Skip to content

Commit

Permalink
Merge pull request #146 from stanleytsang-amd/xnack_on_hmm
Browse files Browse the repository at this point in the history
Cherry-picking HMM unit test support for ROCm 4.3
  • Loading branch information
stanleytsang-amd authored Jun 7, 2021
2 parents 8c41e8b + aaaf755 commit e54e8fd
Show file tree
Hide file tree
Showing 21 changed files with 292 additions and 251 deletions.
2 changes: 1 addition & 1 deletion cmake/Dependencies.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -75,7 +75,7 @@ if(HIP_COMPILER STREQUAL "hcc" OR HIP_COMPILER STREQUAL "clang")
download_project(
PROJ rocprim
GIT_REPOSITORY https://github.com/ROCmSoftwarePlatform/rocPRIM.git
GIT_TAG develop
GIT_TAG release/rocm-rel-4.3
INSTALL_DIR ${CMAKE_CURRENT_BINARY_DIR}/deps/rocprim
CMAKE_ARGS -DBUILD_TEST=OFF -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> -DCMAKE_PREFIX_PATH=/opt/rocm
LOG_DOWNLOAD TRUE
Expand Down
135 changes: 88 additions & 47 deletions test/hipcub/common_test_header.hpp
Original file line number Diff line number Diff line change
@@ -1,47 +1,88 @@
// MIT License
//
// Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.

#include <algorithm>
#include <functional>
#include <iostream>
#include <type_traits>
#include <vector>
#include <utility>
#include <tuple>
#include <random>
#include <limits>
#include <cmath>
#include <cstdlib>

// Google Test
#include <gtest/gtest.h>

// HIP API
#include <hip/hip_runtime.h>

// test_utils.hpp should only be included by this header.
// The following definition is used as guard in test_utils.hpp
// Including test_utils.hpp by itself will cause a compile error.
#define TEST_UTILS_INCLUDE_GAURD
#include "test_utils.hpp"

#define HIP_CHECK(error) ASSERT_EQ(error, hipSuccess)
// MIT License
//
// Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.

#include <algorithm>
#include <functional>
#include <iostream>
#include <type_traits>
#include <vector>
#include <utility>
#include <tuple>
#include <random>
#include <limits>
#include <cmath>
#include <cstdlib>

// Google Test
#include <gtest/gtest.h>

// HIP API
#include <hip/hip_runtime.h>

// test_utils.hpp should only be included by this header.
// The following definition is used as guard in test_utils.hpp
// Including test_utils.hpp by itself will cause a compile error.
#define TEST_UTILS_INCLUDE_GAURD
#include "test_utils.hpp"

#define HIP_CHECK(condition) \
{ \
hipError_t error = condition; \
if(error != hipSuccess){ \
std::cout << "HIP error: " << hipGetErrorString(error) << " line: " << __LINE__ << std::endl; \
exit(error); \
} \
}

namespace test_common_utils
{

bool use_hmm()
{
if (getenv("HIPCUB_USE_HMM") == nullptr)
{
return false;
}

if (strcmp(getenv("HIPCUB_USE_HMM"), "1") == 0)
{
return true;
}
return false;
}

// Helper for HMM allocations: HMM is requested through HIPCUB_USE_HMM environment variable
template <class T>
hipError_t hipMallocHelper(T** devPtr, size_t size)
{
if (use_hmm())
{
return hipMallocManaged((void**)devPtr, size);
}
else
{
return hipMalloc((void**)devPtr, size);
}
return hipSuccess;
}

}
14 changes: 7 additions & 7 deletions test/hipcub/test_hipcub_block_discontinuity.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -237,9 +237,9 @@ TYPED_TEST(HipcubBlockDiscontinuity, FlagHeads)

// Preparing Device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
long long* device_heads;
HIP_CHECK(hipMalloc(&device_heads, heads.size() * sizeof(typename decltype(heads)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_heads, heads.size() * sizeof(typename decltype(heads)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -373,9 +373,9 @@ TYPED_TEST(HipcubBlockDiscontinuity, FlagTails)

// Preparing Device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
long long* device_tails;
HIP_CHECK(hipMalloc(&device_tails, tails.size() * sizeof(typename decltype(tails)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_tails, tails.size() * sizeof(typename decltype(tails)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -535,11 +535,11 @@ TYPED_TEST(HipcubBlockDiscontinuity, FlagHeadsAndTails)

// Preparing Device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
long long* device_heads;
HIP_CHECK(hipMalloc(&device_heads, tails.size() * sizeof(typename decltype(heads)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_heads, tails.size() * sizeof(typename decltype(heads)::value_type)));
long long* device_tails;
HIP_CHECK(hipMalloc(&device_tails, tails.size() * sizeof(typename decltype(tails)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_tails, tails.size() * sizeof(typename decltype(tails)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down
28 changes: 14 additions & 14 deletions test/hipcub/test_hipcub_block_exchange.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -157,9 +157,9 @@ TYPED_TEST(HipcubBlockExchangeTests, BlockedToStriped)

// Preparing device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
output_type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -260,9 +260,9 @@ TYPED_TEST(HipcubBlockExchangeTests, StripedToBlocked)

// Preparing device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
output_type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -381,9 +381,9 @@ TYPED_TEST(HipcubBlockExchangeTests, BlockedToWarpStriped)

// Preparing device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
output_type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -504,9 +504,9 @@ TYPED_TEST(HipcubBlockExchangeTests, WarpStripedToBlocked)

// Preparing device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
output_type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -616,11 +616,11 @@ TYPED_TEST(HipcubBlockExchangeTests, ScatterToBlocked)

// Preparing device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
output_type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
unsigned int* device_ranks;
HIP_CHECK(hipMalloc(&device_ranks, ranks.size() * sizeof(typename decltype(ranks)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_ranks, ranks.size() * sizeof(typename decltype(ranks)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -741,11 +741,11 @@ TYPED_TEST(HipcubBlockExchangeTests, ScatterToStriped)

// Preparing device
type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
output_type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
unsigned int* device_ranks;
HIP_CHECK(hipMalloc(&device_ranks, ranks.size() * sizeof(typename decltype(ranks)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_ranks, ranks.size() * sizeof(typename decltype(ranks)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down
4 changes: 2 additions & 2 deletions test/hipcub/test_hipcub_block_histogram.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -172,9 +172,9 @@ TYPED_TEST(HipcubBlockHistogramInputArrayTests, Histogram)

// Preparing device
T* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(T)));
T* device_output_bin;
HIP_CHECK(hipMalloc(&device_output_bin, output_bin.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output_bin, output_bin.size() * sizeof(T)));

HIP_CHECK(
hipMemcpy(
Expand Down
18 changes: 9 additions & 9 deletions test/hipcub/test_hipcub_block_load_store.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -226,9 +226,9 @@ TYPED_TEST(HipcubBlockLoadStoreClassTests, LoadStoreClass)

// Preparing device
Type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
Type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -331,9 +331,9 @@ TYPED_TEST(HipcubBlockLoadStoreClassTests, LoadStoreClassValid)

// Preparing device
Type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
Type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -446,9 +446,9 @@ TYPED_TEST(HipcubBlockLoadStoreClassTests, LoadStoreClassDefault)

// Preparing device
Type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
Type* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(typename decltype(output)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -611,11 +611,11 @@ TYPED_TEST(HipcubBlockLoadStoreClassTests, LoadStoreDiscardIterator)

// Preparing device
Type* device_input;
HIP_CHECK(hipMalloc(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_input, input.size() * sizeof(typename decltype(input)::value_type)));
Type* device_guarded_elements;
HIP_CHECK(hipMalloc(&device_guarded_elements, guarded_expected.size() * sizeof(typename decltype(unguarded)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_guarded_elements, guarded_expected.size() * sizeof(typename decltype(unguarded)::value_type)));
Type* device_unguarded_elements;
HIP_CHECK(hipMalloc(&device_unguarded_elements, unguarded_expected.size() * sizeof(typename decltype(guarded)::value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_unguarded_elements, unguarded_expected.size() * sizeof(typename decltype(guarded)::value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down
6 changes: 3 additions & 3 deletions test/hipcub/test_hipcub_block_radix_sort.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -241,7 +241,7 @@ TYPED_TEST(HipcubBlockRadixSort, SortKeys)

// Preparing device
key_type* device_keys_output;
HIP_CHECK(hipMalloc(&device_keys_output, keys_output.size() * sizeof(key_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_keys_output, keys_output.size() * sizeof(key_type)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -411,9 +411,9 @@ TYPED_TEST(HipcubBlockRadixSort, SortKeysValues)
}

key_type* device_keys_output;
HIP_CHECK(hipMalloc(&device_keys_output, keys_output.size() * sizeof(key_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_keys_output, keys_output.size() * sizeof(key_type)));
value_type* device_values_output;
HIP_CHECK(hipMalloc(&device_values_output, values_output.size() * sizeof(value_type)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_values_output, values_output.size() * sizeof(value_type)));

HIP_CHECK(
hipMemcpy(
Expand Down
12 changes: 6 additions & 6 deletions test/hipcub/test_hipcub_block_reduce.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -155,9 +155,9 @@ TYPED_TEST(HipcubBlockReduceSingleValueTests, Reduce)

// Preparing device
T* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(T)));
T* device_output_reductions;
HIP_CHECK(hipMalloc(&device_output_reductions, output_reductions.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output_reductions, output_reductions.size() * sizeof(T)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -266,9 +266,9 @@ TYPED_TEST(HipcubBlockReduceSingleValueTests, ReduceValid)

// Preparing device
T* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(T)));
T* device_output_reductions;
HIP_CHECK(hipMalloc(&device_output_reductions, output_reductions.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output_reductions, output_reductions.size() * sizeof(T)));

HIP_CHECK(
hipMemcpy(
Expand Down Expand Up @@ -418,9 +418,9 @@ TYPED_TEST(HipcubBlockReduceInputArrayTests, Reduce)

// Preparing device
T* device_output;
HIP_CHECK(hipMalloc(&device_output, output.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output, output.size() * sizeof(T)));
T* device_output_reductions;
HIP_CHECK(hipMalloc(&device_output_reductions, output_reductions.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&device_output_reductions, output_reductions.size() * sizeof(T)));

HIP_CHECK(
hipMemcpy(
Expand Down
Loading

0 comments on commit e54e8fd

Please sign in to comment.