Skip to content

[rocRAND][hipRAND] Update error handling in rocRAND/hipRAND #1227

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
stanleytsang-amd merged 3 commits into from
Aug 28, 2025
Merged

[rocRAND][hipRAND] Update error handling in rocRAND/hipRAND #1227

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
51a0f47
Update error handling in rocRAND/hipRAND
umfranzw Aug 15, 2025
08f3237
Merge branch 'develop' into users/umfranzw/update_error_handling
stanleytsang-amd Aug 22, 2025
24c5690
Correct HIP_CHECK name collision
umfranzw Aug 27, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 7 additions & 0 deletions projects/hiprand/CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,13 @@
Documentation for hipRAND is available at
[https://rocm.docs.amd.com/projects/hipRAND/en/latest/](https://rocm.docs.amd.com/projects/hipRAND/en/latest/).

## hipRAND 3.1.0 for ROCm 7.1

### Resolved issues

* Updated error handling for several hipRAND unit tests to accomodate the new hipGetLastError behaviour that was introduced in ROCm 7.0.
As of ROCm 7.0, the internal error state is cleared on each call to `hipGetLastError` rather than on every HIP API call.

## hipRAND 3.0.0 for ROCm 7.0

### Added
Expand Down
2 changes: 1 addition & 1 deletion projects/hiprand/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,7 @@ endif()
# hipRAND project
#
project(hipRAND CXX)
set(hipRAND_VERSION "3.0.0")
set(hipRAND_VERSION "3.1.0")

# Build options
option(BUILD_ADDRESS_SANITIZER "Build with address sanitizer enabled" OFF)
Expand Down
16 changes: 15 additions & 1 deletion projects/hiprand/test/test_common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,21 @@

#include <cstdlib>

#define HIP_CHECK(state) ASSERT_EQ(state, hipSuccess)
// GoogleTest-compatible HIP_CHECK macro. FAIL is called to log the Google Test trace.
// The lambda is invoked immediately as assertions that generate a fatal failure can
// only be used in void-returning functions.
#define HIP_CHECK(condition) \
{ \
hipError_t error = condition; \
if(error != hipSuccess) \
{ \
[error]() \
{ FAIL() << "HIP error " << error << ": " << hipGetErrorString(error); }(); \
exit(error); \
} \
}

#define HIPRAND_CHECK(state) ASSERT_EQ(state, HIPRAND_STATUS_SUCCESS)

#define HIP_CHECK_NON_VOID(condition) \
{ \
Expand Down
4 changes: 1 addition & 3 deletions projects/hiprand/test/test_hiprand_api.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,8 +26,6 @@

#include <gtest/gtest.h>

#define HIPRAND_CHECK(state) ASSERT_EQ(state, HIPRAND_STATUS_SUCCESS)

constexpr hiprandRngType_t hiprand_rng_types[] = {HIPRAND_RNG_PSEUDO_XORWOW,
HIPRAND_RNG_PSEUDO_MRG32K3A,
HIPRAND_RNG_PSEUDO_MTGP32,
Expand Down Expand Up @@ -570,7 +568,7 @@ INSTANTIATE_TEST_SUITE_P(hiprand, hiprand_ordering, ::testing::ValuesIn(hiprand_
void hiprand_generate_uniform_host_test_func(hiprandRngType_t rng_type)
{
hiprandGenerator_t generator = 0;
HIP_CHECK(hiprandCreateGeneratorHost(&generator, rng_type));
HIPRAND_CHECK(hiprandCreateGeneratorHost(&generator, rng_type));

constexpr size_t output_size = 8192;
std::vector<float> output_host(output_size);
Expand Down
32 changes: 16 additions & 16 deletions projects/hiprand/test/test_hiprand_kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,7 +95,7 @@ void load_sobol_vectors_to_gpu(const unsigned int dimensions,

constexpr unsigned int vec_size = sizeof(T) * 8;

HIP_CHECK(get_sobol_directions<S>(&h_directions));
HIPRAND_CHECK(get_sobol_directions<S>(&h_directions));

HIP_CHECK(hipMallocHelper(reinterpret_cast<void**>(direction_vectors),
sizeof(T) * dimensions * vec_size));
Expand All @@ -114,7 +114,7 @@ void load_scrambled_sobol_constants_and_vectors_to_gpu(const unsigned int d

const T* h_constants;

HIP_CHECK(get_sobol_constants(&h_constants));
HIPRAND_CHECK(get_sobol_constants(&h_constants));

HIP_CHECK(
hipMallocHelper(reinterpret_cast<void**>(scramble_constants), sizeof(T) * dimensions));
Expand Down Expand Up @@ -765,10 +765,10 @@ void hiprand_kernel_h_hiprand_mtgp_init_test()
mtgp32_kernel_params_t* k;
HIP_CHECK(hipMallocHelper((void**)&k, states_size * sizeof(mtgp32_kernel_params_t)));

HIP_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));
HIPRAND_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));

unsigned long long seed = 0;
HIP_CHECK(
HIPRAND_CHECK(
hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, k, states_size, seed));

HIP_CHECK(hipFree(states));
Expand Down Expand Up @@ -997,10 +997,10 @@ void hiprand_kernel_h_hiprand_mtgp_test()
mtgp32_kernel_params_t* k;
HIP_CHECK(hipMallocHelper((void**)&k, states_size * sizeof(mtgp32_kernel_params_t)));

HIP_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));
HIPRAND_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));

unsigned long long seed = 0;
HIP_CHECK(
HIPRAND_CHECK(
hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, k, states_size, seed));

const size_t output_size = 8192;
Expand Down Expand Up @@ -1155,10 +1155,10 @@ void hiprand_kernel_h_hiprand_uniform_mtgp_test()
mtgp32_kernel_params_t* k;
HIP_CHECK(hipMallocHelper((void**)&k, states_size * sizeof(mtgp32_kernel_params_t)));

HIP_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));
HIPRAND_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));

unsigned long long seed = 0;
HIP_CHECK(
HIPRAND_CHECK(
hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, k, states_size, seed));

const size_t output_size = 8192;
Expand Down Expand Up @@ -1328,10 +1328,10 @@ void hiprand_kernel_h_hiprand_normal_mtgp_test()
mtgp32_kernel_params_t* k;
HIP_CHECK(hipMallocHelper((void**)&k, states_size * sizeof(mtgp32_kernel_params_t)));

HIP_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));
HIPRAND_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));

unsigned long long seed = 0;
HIP_CHECK(
HIPRAND_CHECK(
hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, k, states_size, seed));

const size_t output_size = 8192;
Expand Down Expand Up @@ -1522,10 +1522,10 @@ void hiprand_kernel_h_hiprand_log_normal_mtgp_test()
mtgp32_kernel_params_t* k;
HIP_CHECK(hipMallocHelper((void**)&k, states_size * sizeof(mtgp32_kernel_params_t)));

HIP_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));
HIPRAND_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));

unsigned long long seed = 0;
HIP_CHECK(
HIPRAND_CHECK(
hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, k, states_size, seed));

const size_t output_size = 8192;
Expand Down Expand Up @@ -1829,10 +1829,10 @@ void hiprand_kernel_h_hiprand_poisson_mtgp_test(double lambda)
mtgp32_kernel_params_t* k;
HIP_CHECK(hipMallocHelper((void**)&k, states_size * sizeof(mtgp32_kernel_params_t)));

HIP_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));
HIPRAND_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));

unsigned long long seed = 0;
HIP_CHECK(
HIPRAND_CHECK(
hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, k, states_size, seed));

const size_t output_size = 8192;
Expand Down Expand Up @@ -1990,10 +1990,10 @@ void hiprand_kernel_h_hiprand_discrete_mtgp_test(double lambda)
mtgp32_kernel_params_t* k;
HIP_CHECK(hipMallocHelper((void**)&k, states_size * sizeof(mtgp32_kernel_params_t)));

HIP_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));
HIPRAND_CHECK(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, k));

unsigned long long seed = 0;
HIP_CHECK(
HIPRAND_CHECK(
hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, k, states_size, seed));

const size_t output_size = 8192;
Expand Down
7 changes: 7 additions & 0 deletions projects/rocrand/CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,13 @@
Documentation for rocRAND is available at
[https://rocm.docs.amd.com/projects/rocRAND/en/latest/](https://rocm.docs.amd.com/projects/rocRAND/en/latest/)

## rocRAND 4.1.0 for ROCm 7.1

### Resolved issues

* Updated error handling for several rocRAND unit tests to accomodate the new hipGetLastError behaviour that was introduced in ROCm 7.0.
As of ROCm 7.0, the internal error state is cleared on each call to `hipGetLastError` rather than on every HIP API call.

## rocRAND 4.0.0 for ROCm 7.0

### Added
Expand Down
2 changes: 1 addition & 1 deletion projects/rocrand/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -178,7 +178,7 @@ if(BUILD_ADDRESS_SANITIZER AND BUILD_SHARED_LIBS)
endif()

# Set version variables
rocm_setup_version( VERSION "4.0.0" )
rocm_setup_version( VERSION "4.1.0" )
set ( rocrand_VERSION ${rocRAND_VERSION} )
# Old-style version number used within the library's API. rocrand_get_version should be modified.
math(EXPR rocrand_VERSION_NUMBER "${rocRAND_VERSION_MAJOR} * 100000 + ${rocRAND_VERSION_MINOR} * 100 + ${rocRAND_VERSION_PATCH}")
Expand Down
4 changes: 2 additions & 2 deletions projects/rocrand/test/internal/test_rocrand_config_dispatch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -121,11 +121,11 @@ TEST(rocrand_config_dispatch_tests, get_config_on_host_and_device)
HIP_CHECK(hipMallocHelper(&d_grid_size, sizeof(*d_grid_size)));

rocrand_impl::host::generator_config config{};
const hipError_t error = rocrand_impl::host::get_generator_config<dummy_rng_type, T>(
const hipError_t err = rocrand_impl::host::get_generator_config<dummy_rng_type, T>(
stream,
ROCRAND_ORDERING_PSEUDO_DEFAULT,
config);
HIP_CHECK(error);
HIP_CHECK(err);

hipLaunchKernelGGL(write_config<T>,
dim3(config.blocks),
Expand Down
14 changes: 13 additions & 1 deletion projects/rocrand/test/internal/test_rocrand_generator_type.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,19 @@

#include <vector>

#define HIP_CHECK(state) ASSERT_EQ(state, hipSuccess)
// GoogleTest-compatible HIP_CHECK macro. FAIL is called to log the Google Test trace.
// The lambda is invoked immediately as assertions that generate a fatal failure can
// only be used in void-returning functions.
#define HIP_CHECK(condition) \
{ \
hipError_t error = condition; \
if(error != hipSuccess) \
{ \
[error]() \
{ FAIL() << "HIP error " << error << ": " << hipGetErrorString(error); }(); \
exit(error); \
} \
}

using namespace rocrand_impl::host;

Expand Down
14 changes: 13 additions & 1 deletion projects/rocrand/test/test_common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,19 @@
#include <iostream>
#include <vector>

#define HIP_CHECK(state) ASSERT_EQ(state, hipSuccess)
// GoogleTest-compatible HIP_CHECK macro. FAIL is called to log the Google Test trace.
// The lambda is invoked immediately as assertions that generate a fatal failure can
// only be used in void-returning functions.
#define HIP_CHECK(condition) \
{ \
hipError_t error = condition; \
if(error != hipSuccess) \
{ \
[error]() \
{ FAIL() << "HIP error " << error << ": " << hipGetErrorString(error); }(); \
exit(error); \
} \
}

#define HIP_CHECK_NON_VOID(condition) \
{ \
Expand Down
15 changes: 14 additions & 1 deletion projects/rocrand/test/test_rocrand_kernel_sobol64.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,20 @@

#include <rocrand/rocrand_kernel.h>

#define HIP_CHECK(state) ASSERT_EQ(state, hipSuccess)
// GoogleTest-compatible HIP_CHECK macro. FAIL is called to log the Google Test trace.
// The lambda is invoked immediately as assertions that generate a fatal failure can
// only be used in void-returning functions.
#define HIP_CHECK(condition) \
{ \
hipError_t error = condition; \
if(error != hipSuccess) \
{ \
[error]() \
{ FAIL() << "HIP error " << error << ": " << hipGetErrorString(error); }(); \
exit(error); \
} \
}

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

template<class GeneratorState>
Expand Down
Loading