Skip to content

Commit

Permalink
Enabling HMM unit test support (#234)
Browse files Browse the repository at this point in the history
* Converting benchmar_device_radix_sort to HMM.

* Adding initial HMM unit tests.

* Refactoring unit tests so that all can use hipMallocManaged

* Temporarily disabling managed memory check

* Removing remainder of size zero test cases

* Updating CHANGELOG for ROCm 4.3

* Temporarily removing deprecation compiler warning for ROCm 4.3

* Revert "Removing remainder of size zero test cases"

This reverts commit e70b00c.

* Disabling benchmark build by default

* More updates to HMM

* Updating HMM unit tests

* updating changelog for 4.4
  • Loading branch information
stanleytsang-amd authored Jun 11, 2021
1 parent ac4712a commit 027bf32
Show file tree
Hide file tree
Showing 40 changed files with 604 additions and 400 deletions.
17 changes: 8 additions & 9 deletions benchmark/benchmark_device_radix_sort.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -41,15 +41,6 @@
// rocPRIM
#include <rocprim/rocprim.hpp>

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

#ifndef DEFAULT_N
const size_t DEFAULT_N = 1024 * 1024 * 32;
#endif
Expand Down Expand Up @@ -89,8 +80,10 @@ void run_sort_keys_benchmark(benchmark::State& state,

key_type * d_keys_input;
key_type * d_keys_output;

HIP_CHECK(hipMalloc(&d_keys_input, size * sizeof(key_type)));
HIP_CHECK(hipMalloc(&d_keys_output, size * sizeof(key_type)));

HIP_CHECK(
hipMemcpy(
d_keys_input, keys_input->data(),
Expand Down Expand Up @@ -174,8 +167,10 @@ void run_sort_pairs_benchmark(benchmark::State& state,

key_type * d_keys_input;
key_type * d_keys_output;

HIP_CHECK(hipMalloc(&d_keys_input, size * sizeof(key_type)));
HIP_CHECK(hipMalloc(&d_keys_output, size * sizeof(key_type)));

HIP_CHECK(
hipMemcpy(
d_keys_input, keys_input->data(),
Expand All @@ -186,8 +181,10 @@ void run_sort_pairs_benchmark(benchmark::State& state,

value_type * d_values_input;
value_type * d_values_output;

HIP_CHECK(hipMalloc(&d_values_input, size * sizeof(value_type)));
HIP_CHECK(hipMalloc(&d_values_output, size * sizeof(value_type)));

HIP_CHECK(
hipMemcpy(
d_values_input, values_input.data(),
Expand All @@ -207,6 +204,8 @@ void run_sort_pairs_benchmark(benchmark::State& state,
)
);

HIP_CHECK(hipMalloc(&d_temporary_storage, temporary_storage_bytes));

HIP_CHECK(hipMalloc(&d_temporary_storage, temporary_storage_bytes));
HIP_CHECK(hipDeviceSynchronize());

Expand Down
9 changes: 9 additions & 0 deletions benchmark/benchmark_utils.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,15 @@

#include <rocprim/rocprim.hpp>

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

// Support half operators on host side

ROCPRIM_HOST inline
Expand Down
1 change: 0 additions & 1 deletion rocprim/include/rocprim/intrinsics/thread.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,6 @@ BEGIN_ROCPRIM_NAMESPACE
/// It is constant for a device.
/// This function is not supported for the gfx1030 architecture and will be removed in a future release.
/// Please use the new host_warp_size() and device_warp_size() functions.
[[deprecated]]
ROCPRIM_HOST_DEVICE inline
constexpr unsigned int warp_size()
{
Expand Down
183 changes: 107 additions & 76 deletions test/common_test_header.hpp
Original file line number Diff line number Diff line change
@@ -1,76 +1,107 @@
// 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.

#ifndef ROCPRIM_COMMON_TEST_HEADER
#define ROCPRIM_COMMON_TEST_HEADER

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

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

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

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

#include <cstdlib>
#include <string>
#include <cctype>

namespace test_common_utils
{

int obtain_device_from_ctest()
{
static const std::string rg0 = "CTEST_RESOURCE_GROUP_0";
if (std::getenv(rg0.c_str()) != nullptr)
{
std::string amdgpu_target = std::getenv(rg0.c_str());
std::transform(amdgpu_target.cbegin(), amdgpu_target.cend(), amdgpu_target.begin(), ::toupper);
std::string reqs = std::getenv((rg0 + "_" + amdgpu_target).c_str());
return std::atoi(reqs.substr(reqs.find(':') + 1, reqs.find(',') - (reqs.find(':') + 1)).c_str());
}
else
return 0;
}

}

#endif
// 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.

#ifndef ROCPRIM_COMMON_TEST_HEADER
#define ROCPRIM_COMMON_TEST_HEADER

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

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

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

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

#include <cstdlib>
#include <string>
#include <cctype>

namespace test_common_utils
{

int obtain_device_from_ctest()
{
static const std::string rg0 = "CTEST_RESOURCE_GROUP_0";
if (std::getenv(rg0.c_str()) != nullptr)
{
std::string amdgpu_target = std::getenv(rg0.c_str());
std::transform(amdgpu_target.cbegin(), amdgpu_target.cend(), amdgpu_target.begin(), ::toupper);
std::string reqs = std::getenv((rg0 + "_" + amdgpu_target).c_str());
return std::atoi(reqs.substr(reqs.find(':') + 1, reqs.find(',') - (reqs.find(':') + 1)).c_str());
}
else
return 0;
}

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

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

// Helper for HMM allocations: HMM is requested through ROCPRIM_USE_HMM=1 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;
}

}

#endif
6 changes: 3 additions & 3 deletions test/rocprim/test_arg_index_iterator.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -131,8 +131,8 @@ TYPED_TEST(RocprimArgIndexIteratorTests, ReduceArgMinimum)

T * d_input;
key_value * d_output;
HIP_CHECK(hipMalloc(&d_input, input.size() * sizeof(T)));
HIP_CHECK(hipMalloc(&d_output, output.size() * sizeof(key_value)));
HIP_CHECK(test_common_utils::hipMallocHelper(&d_input, input.size() * sizeof(T)));
HIP_CHECK(test_common_utils::hipMallocHelper(&d_output, output.size() * sizeof(key_value)));
HIP_CHECK(
hipMemcpy(
d_input, input.data(),
Expand Down Expand Up @@ -167,7 +167,7 @@ TYPED_TEST(RocprimArgIndexIteratorTests, ReduceArgMinimum)
ASSERT_GT(temp_storage_size_bytes, 0);

// allocate temporary storage
HIP_CHECK(hipMalloc(&d_temp_storage, temp_storage_size_bytes));
HIP_CHECK(test_common_utils::hipMallocHelper(&d_temp_storage, temp_storage_size_bytes));
HIP_CHECK(hipDeviceSynchronize());

// Run
Expand Down
14 changes: 7 additions & 7 deletions test/rocprim/test_block_discontinuity.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -280,9 +280,9 @@ auto test_block_discontinuity()

// 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 @@ -392,9 +392,9 @@ auto test_block_discontinuity()

// 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 @@ -516,11 +516,11 @@ auto test_block_discontinuity()

// 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
Loading

0 comments on commit 027bf32

Please sign in to comment.