Skip to content

Commit

Permalink
Merge pull request #179 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
stanleytsang-amd authored Jun 7, 2021
2 parents d5e2be9 + 0e2fd40 commit 16989af
Show file tree
Hide file tree
Showing 20 changed files with 184 additions and 140 deletions.
52 changes: 52 additions & 0 deletions test/test_common.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -21,9 +21,19 @@
#ifndef TEST_COMMON_HPP_
#define TEST_COMMON_HPP_

#include <cstdlib>

#define HIP_CHECK(state) ASSERT_EQ(state, hipSuccess)
#define ROCRAND_CHECK(state) ASSERT_EQ(state, ROCRAND_STATUS_SUCCESS)

#define HIP_CHECK_NON_VOID(condition) \
{ \
hipError_t error = condition; \
if(error != hipSuccess){ \
std::cout << "HIP error: " << error << " line: " << __LINE__ << std::endl; \
exit(error); \
} \
}
const rocrand_rng_type rng_types[] = {
ROCRAND_RNG_PSEUDO_PHILOX4_32_10,
ROCRAND_RNG_PSEUDO_MRG32K3A,
Expand All @@ -32,4 +42,46 @@ const rocrand_rng_type rng_types[] = {
ROCRAND_RNG_QUASI_SOBOL32
};

bool supports_hmm()
{
hipDeviceProp_t device_prop;
int device_id;
HIP_CHECK_NON_VOID(hipGetDevice(&device_id));
HIP_CHECK_NON_VOID(hipGetDeviceProperties(&device_prop, device_id));
if (device_prop.managedMemory == 1) return true;

return false;
}

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

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

// Helper for HMM allocations: if HMM is requested through
// setting environment variable ROCRAND_USE_HMM=1
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 // TEST_COMMON_HPP_
20 changes: 11 additions & 9 deletions test/test_hiprand_api.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -24,9 +24,11 @@
#include <hip/hip_runtime.h>
#include <hiprand.h>

#define HIP_CHECK(x) ASSERT_EQ(x, hipSuccess)
#include "test_common.hpp"
#define HIPRAND_CHECK(state) ASSERT_EQ(state, HIPRAND_STATUS_SUCCESS)



template<hiprandRngType_t rng_type>
void hiprand_generate_test_func()
{
Expand All @@ -36,7 +38,7 @@ void hiprand_generate_test_func()
const size_t output_size = 8192;
unsigned int * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(unsigned int))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -101,7 +103,7 @@ void hiprand_generate_uniform_test_func()
const size_t output_size = 8192;
float * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(float))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -166,7 +168,7 @@ void hiprand_generate_uniform_double_test_func()
const size_t output_size = 8192;
double * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(double))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -231,7 +233,7 @@ void hiprand_generate_normal_test_func()
const size_t output_size = 8192;
float * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(float))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -304,7 +306,7 @@ void hiprand_generate_normal_double_test_func()
const size_t output_size = 8192;
double * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(double))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -377,7 +379,7 @@ void hiprand_generate_lognormal_test_func()
const size_t output_size = 8192;
float * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(float))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -454,7 +456,7 @@ void hiprand_generate_lognormal_double_test_func()
const size_t output_size = 8192;
double * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(double))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -532,7 +534,7 @@ void hiprand_generate_poisson_test_func()
const size_t output_size = 8192;
unsigned int * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(unsigned int))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down
12 changes: 6 additions & 6 deletions test/test_hiprand_cpp_wrapper.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
#include <hip/hip_runtime.h>
#include <hiprand.hpp>

#define HIP_CHECK(x) ASSERT_EQ(x, hipSuccess)
#include "test_common.hpp"

TEST(hiprand_cpp_wrapper, hiprand_error)
{
Expand Down Expand Up @@ -230,7 +230,7 @@ void hiprand_uniform_int_dist_template()
const size_t output_size = 8192;
IntType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(IntType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -287,7 +287,7 @@ void hiprand_uniform_real_dist_template()
const size_t output_size = 8192;
RealType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(RealType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -363,7 +363,7 @@ void hiprand_normal_dist_template()
const size_t output_size = 8192;
RealType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(RealType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -461,7 +461,7 @@ void hiprand_lognormal_dist_template()
const size_t output_size = 8192;
RealType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(RealType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -568,7 +568,7 @@ void hiprand_poisson_dist_template(const double lambda)
const size_t output_size = 8192;
IntType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(IntType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down
22 changes: 11 additions & 11 deletions test/test_hiprand_kernel.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -26,12 +26,12 @@

#include <hip/hip_runtime.h>


#define QUALIFIERS __forceinline__ __host__ __device__
#include <hiprand_kernel.h>
#include <hiprand.h>

#define HIP_CHECK(x) ASSERT_EQ(x, hipSuccess)
#define HIPRAND_CHECK(state) ASSERT_EQ(state, HIPRAND_STATUS_SUCCESS)
#include "test_common.hpp"

template <class GeneratorState>
__global__
Expand Down Expand Up @@ -235,7 +235,7 @@ void hiprand_kernel_h_hiprand_init_test()

const size_t states_size = 256;
state_type * states;
HIP_CHECK(hipMalloc((void **)&states, states_size * sizeof(state_type)));
HIP_CHECK(hipMallocHelper((void **)&states, states_size * sizeof(state_type)));
HIP_CHECK(hipDeviceSynchronize());

hipLaunchKernelGGL(
Expand Down Expand Up @@ -283,7 +283,7 @@ TEST(hiprand_kernel_h_philox4x32_10, hiprand_init_nvcc)

const size_t states_size = 256;
state_type * states;
HIP_CHECK(hipMalloc((void **)&states, states_size * sizeof(state_type)));
HIP_CHECK(hipMallocHelper((void **)&states, states_size * sizeof(state_type)));
HIP_CHECK(hipDeviceSynchronize());

hipLaunchKernelGGL(
Expand Down Expand Up @@ -338,7 +338,7 @@ TEST(hiprand_kernel_h_philox4x32_10, hiprand_skip_nvcc)

const size_t states_size = 256;
state_type * states;
HIP_CHECK(hipMalloc((void **)&states, states_size * sizeof(state_type)));
HIP_CHECK(hipMallocHelper((void **)&states, states_size * sizeof(state_type)));
HIP_CHECK(hipDeviceSynchronize());

hipLaunchKernelGGL(
Expand Down Expand Up @@ -380,7 +380,7 @@ void hiprand_kernel_h_hiprand_test()

const size_t output_size = 8192;
unsigned int * output;
HIP_CHECK(hipMalloc((void **)&output, output_size * sizeof(unsigned int)));
HIP_CHECK(hipMallocHelper((void **)&output, output_size * sizeof(unsigned int)));
HIP_CHECK(hipDeviceSynchronize());

hipLaunchKernelGGL(
Expand Down Expand Up @@ -441,7 +441,7 @@ void hiprand_kernel_h_hiprand_uniform_test()

const size_t output_size = 8192;
float * output;
HIP_CHECK(hipMalloc((void **)&output, output_size * sizeof(float)));
HIP_CHECK(hipMallocHelper((void **)&output, output_size * sizeof(float)));
HIP_CHECK(hipDeviceSynchronize());

hipLaunchKernelGGL(
Expand Down Expand Up @@ -502,7 +502,7 @@ void hiprand_kernel_h_hiprand_normal_test()

const size_t output_size = 8192;
float * output;
HIP_CHECK(hipMalloc((void **)&output, output_size * sizeof(float)));
HIP_CHECK(hipMallocHelper((void **)&output, output_size * sizeof(float)));
HIP_CHECK(hipDeviceSynchronize());

hipLaunchKernelGGL(
Expand Down Expand Up @@ -571,7 +571,7 @@ void hiprand_kernel_h_hiprand_log_normal_test()

const size_t output_size = 8192;
float * output;
HIP_CHECK(hipMalloc((void **)&output, output_size * sizeof(float)));
HIP_CHECK(hipMallocHelper((void **)&output, output_size * sizeof(float)));
HIP_CHECK(hipDeviceSynchronize());

hipLaunchKernelGGL(
Expand Down Expand Up @@ -644,7 +644,7 @@ void hiprand_kernel_h_hiprand_poisson_test(double lambda)

const size_t output_size = 8192;
unsigned int * output;
HIP_CHECK(hipMalloc((void **)&output, output_size * sizeof(unsigned int)));
HIP_CHECK(hipMallocHelper((void **)&output, output_size * sizeof(unsigned int)));
HIP_CHECK(hipDeviceSynchronize());

hipLaunchKernelGGL(
Expand Down Expand Up @@ -690,7 +690,7 @@ void hiprand_kernel_h_hiprand_discrete_test(double lambda)

const size_t output_size = 8192;
unsigned int * output;
HIP_CHECK(hipMalloc((void **)&output, output_size * sizeof(unsigned int)));
HIP_CHECK(hipMallocHelper((void **)&output, output_size * sizeof(unsigned int)));
HIP_CHECK(hipDeviceSynchronize());

hiprandDiscreteDistribution_t discrete_distribution;
Expand Down
12 changes: 6 additions & 6 deletions test/test_rocrand_cpp_wrapper.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
#include <hip/hip_runtime.h>
#include <rocrand.hpp>

#define HIP_CHECK(x) ASSERT_EQ(x, hipSuccess)
#include "test_common.hpp"

TEST(rocrand_cpp_wrapper, rocrand_error)
{
Expand Down Expand Up @@ -230,7 +230,7 @@ void rocrand_uniform_int_dist_template()
const size_t output_size = 8192;
IntType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(IntType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -287,7 +287,7 @@ void rocrand_uniform_real_dist_template()
const size_t output_size = 8192;
RealType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(RealType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -363,7 +363,7 @@ void rocrand_normal_dist_template()
const size_t output_size = 8192;
RealType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(RealType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -461,7 +461,7 @@ void rocrand_lognormal_dist_template()
const size_t output_size = 8192;
RealType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(RealType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down Expand Up @@ -568,7 +568,7 @@ void rocrand_poisson_dist_template(const double lambda)
const size_t output_size = 8192;
IntType * output;
HIP_CHECK(
hipMalloc((void **)&output,
hipMallocHelper((void **)&output,
output_size * sizeof(IntType))
);
HIP_CHECK(hipDeviceSynchronize());
Expand Down
6 changes: 3 additions & 3 deletions test/test_rocrand_generate.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ TEST_P(rocrand_generate_tests, int_test)

const size_t size = 12563;
unsigned int * data;
HIP_CHECK(hipMalloc((void **)&data, size * sizeof(unsigned int)));
HIP_CHECK(hipMallocHelper((void **)&data, size * sizeof(unsigned int)));
HIP_CHECK(hipDeviceSynchronize());

// Any sizes
Expand Down Expand Up @@ -80,7 +80,7 @@ TEST_P(rocrand_generate_tests, char_test)

const size_t size = 12563;
unsigned char * data;
HIP_CHECK(hipMalloc((void **)&data, size * sizeof(unsigned char)));
HIP_CHECK(hipMallocHelper((void **)&data, size * sizeof(unsigned char)));
HIP_CHECK(hipDeviceSynchronize());

// Any sizes
Expand Down Expand Up @@ -118,7 +118,7 @@ TEST_P(rocrand_generate_tests, short_test)

const size_t size = 12563;
unsigned short * data;
HIP_CHECK(hipMalloc((void **)&data, size * sizeof(unsigned short)));
HIP_CHECK(hipMallocHelper((void **)&data, size * sizeof(unsigned short)));
HIP_CHECK(hipDeviceSynchronize());

// Any sizes
Expand Down
Loading

0 comments on commit 16989af

Please sign in to comment.