Skip to content

[rocRAND][Code Coverage] Increase rocRAND code coverage #173

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

Closed
NguyenNhuDi wants to merge 14 commits into from
Closed

[rocRAND][Code Coverage] Increase rocRAND code coverage #173

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
4624c8e
ported over rocrand unit tests
NguyenNhuDi Jun 9, 2025
f88450b
Merge branch 'release-staging/rocm-rel-7.0' into zenguyen-rocrand/inc…
NguyenNhuDi Jun 19, 2025
007e763
implemented host side unit test for log_normal
NguyenNhuDi Jun 20, 2025
2ab8b8f
implemented host side test for uniform distribution
NguyenNhuDi Jun 20, 2025
6824013
implemneted host side test for discrete distribution
NguyenNhuDi Jun 20, 2025
8c95647
added poisson host tests
NguyenNhuDi Jun 23, 2025
3c876e6
implemented suggestions
NguyenNhuDi Jun 23, 2025
036c78f
implemented additional poisson host tests
NguyenNhuDi Jun 24, 2025
016987d
implemented additional host unit tests for rocrand normal
NguyenNhuDi Jun 24, 2025
3ad11a9
fixed some out of bounds error in uniform_distribution tests
NguyenNhuDi Jun 25, 2025
74421cf
added ROCRAND_DETAIL_BM_NOT_IN_STATE tests
NguyenNhuDi Jun 25, 2025
15783b5
Merge branch 'release-staging/rocm-rel-7.0' into zenguyen-rocrand/inc…
NguyenNhuDi Jun 26, 2025
110808e
implemented additional ROCRAND_DETAIL_BM_NOT_IN_STATE tests
NguyenNhuDi Jun 26, 2025
d4528e4
started initial work for mt100037_octo_engine_test
NguyenNhuDi Jun 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
18 changes: 18 additions & 0 deletions projects/rocrand/CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,24 @@ Documentation for rocRAND is available at
* Changed return type for `rocrand_generate_poisson` for `SOBOL64` and `SCRAMBLED_SOBOL64` engines
* Changed unnecessarily large 64-bit data type of constants used for skipping in `MRG32K3A` to 32-bit data type

### Added
* Additional unit tests for `test_log_normal_distribution.cpp`
* Additional unit tests for `test_normal_distribution.cpp`
* Additional unit tests for `test_rocrand_mtgp32_prng.cpp`
* Additional unit tests for `test_rocrand_scrambled_sobol32_qrng.cpp`
* Additional unit tests for `test_rocrand_scrambled_sobol64_qrng.cpp`
* Additional unit tests for `test_rocrand_sobol32_qrng.cpp`
* Additional unit tests for `test_rocrand_sobol64_qrng.cpp`
* Additional unit tests for `test_rocrand_threefry2x32_20_prng.cpp`
* Additional unit tests for `test_rocrand_threefry2x64_20_prng.cpp`
* Additional unit tests for `test_rocrand_threefry4x32_20_prng.cpp`
* Additional unit tests for `test_rocrand_threefry4x64_20_prng.cpp`
* Additional unit tests for `test_uniform_distribution.cpp`
* New unit tests for `include/rocrand/rocrand_discrete.h` in `test_rocrand_discrete.cpp`
* New unit tests for `include/rocrand/rocrand_mrg31k3p.h` in `test_rocrand_mrg31k3p_prng.cpp`
* New unit tests for `include/rocrand/rocrand_mrg32k3a.h` in `test_rocrand_mrg32k3a_prng.cpp`
* New unit tests for `include/rocrand/rocrand_poisson.h` in `test_rocrand_poisson.cpp`

## rocRAND 3.4.0 for ROCm 6.5

### Added
Expand Down
82 changes: 29 additions & 53 deletions projects/rocrand/library/include/rocrand/rocrand_log_normal.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,8 +63,8 @@
* \return Log-normally distributed \p float value
*/
#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
__forceinline__ __device__ __host__ float
rocrand_log_normal(rocrand_state_philox4x32_10* state, float mean, float stddev)
__forceinline__ __device__ __host__
float rocrand_log_normal(rocrand_state_philox4x32_10* state, float mean, float stddev)
{
typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;

Expand Down Expand Up @@ -103,10 +103,7 @@ float2 rocrand_log_normal2(rocrand_state_philox4x32_10* state, float mean, float
auto state2 = rocrand(state);

float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
return float2 {
expf(mean + (stddev * r.x)),
expf(mean + (stddev * r.y))
};
return float2{expf(mean + (stddev * r.x)), expf(mean + (stddev * r.y))};
}

/**
Expand All @@ -127,12 +124,10 @@ __forceinline__ __device__ __host__
float4 rocrand_log_normal4(rocrand_state_philox4x32_10* state, float mean, float stddev)
{
float4 r = rocrand_device::detail::normal_distribution4(rocrand4(state));
return float4 {
expf(mean + (stddev * r.x)),
expf(mean + (stddev * r.y)),
expf(mean + (stddev * r.z)),
expf(mean + (stddev * r.w))
};
return float4{expf(mean + (stddev * r.x)),
expf(mean + (stddev * r.y)),
expf(mean + (stddev * r.z)),
expf(mean + (stddev * r.w))};
}

/**
Expand All @@ -151,8 +146,8 @@ float4 rocrand_log_normal4(rocrand_state_philox4x32_10* state, float mean, float
* \return Log-normally distributed \p double value
*/
#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
__forceinline__ __device__ __host__ double
rocrand_log_normal_double(rocrand_state_philox4x32_10* state, double mean, double stddev)
__forceinline__ __device__ __host__
double rocrand_log_normal_double(rocrand_state_philox4x32_10* state, double mean, double stddev)
{
typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;

Expand Down Expand Up @@ -184,10 +179,7 @@ __forceinline__ __device__ __host__
double2 rocrand_log_normal_double2(rocrand_state_philox4x32_10* state, double mean, double stddev)
{
double2 r = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
return double2 {
exp(mean + (stddev * r.x)),
exp(mean + (stddev * r.y))
};
return double2{exp(mean + (stddev * r.x)), exp(mean + (stddev * r.y))};
}

/**
Expand All @@ -210,9 +202,7 @@ double4 rocrand_log_normal_double4(rocrand_state_philox4x32_10* state, double me
double2 r1, r2;
r1 = rocrand_log_normal_double2(state, mean, stddev);
r2 = rocrand_log_normal_double2(state, mean, stddev);
return double4 {
r1.x, r1.y, r2.x, r2.y
};
return double4{r1.x, r1.y, r2.x, r2.y};
}

/**
Expand All @@ -231,8 +221,8 @@ double4 rocrand_log_normal_double4(rocrand_state_philox4x32_10* state, double me
* \return Log-normally distributed \p float value
*/
#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
__forceinline__ __device__ __host__ float
rocrand_log_normal(rocrand_state_mrg31k3p* state, float mean, float stddev)
__forceinline__ __device__ __host__
float rocrand_log_normal(rocrand_state_mrg31k3p* state, float mean, float stddev)
{
typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;

Expand Down Expand Up @@ -292,8 +282,8 @@ float2 rocrand_log_normal2(rocrand_state_mrg31k3p* state, float mean, float stdd
* \return Log-normally distributed \p double value
*/
#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
__forceinline__ __device__ __host__ double
rocrand_log_normal_double(rocrand_state_mrg31k3p* state, double mean, double stddev)
__forceinline__ __device__ __host__
double rocrand_log_normal_double(rocrand_state_mrg31k3p* state, double mean, double stddev)
{
typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;

Expand Down Expand Up @@ -355,8 +345,8 @@ double2 rocrand_log_normal_double2(rocrand_state_mrg31k3p* state, double mean, d
* \return Log-normally distributed \p float value
*/
#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
__forceinline__ __device__ __host__ float
rocrand_log_normal(rocrand_state_mrg32k3a* state, float mean, float stddev)
__forceinline__ __device__ __host__
float rocrand_log_normal(rocrand_state_mrg32k3a* state, float mean, float stddev)
{
typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;

Expand Down Expand Up @@ -397,10 +387,7 @@ float2 rocrand_log_normal2(rocrand_state_mrg32k3a* state, float mean, float stdd

float2 r
= rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
return float2 {
expf(mean + (stddev * r.x)),
expf(mean + (stddev * r.y))
};
return float2{expf(mean + (stddev * r.x)), expf(mean + (stddev * r.y))};
}

/**
Expand All @@ -419,8 +406,8 @@ float2 rocrand_log_normal2(rocrand_state_mrg32k3a* state, float mean, float stdd
* \return Log-normally distributed \p double value
*/
#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
__forceinline__ __device__ __host__ double
rocrand_log_normal_double(rocrand_state_mrg32k3a* state, double mean, double stddev)
__forceinline__ __device__ __host__
double rocrand_log_normal_double(rocrand_state_mrg32k3a* state, double mean, double stddev)
{
typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;

Expand Down Expand Up @@ -463,10 +450,7 @@ double2 rocrand_log_normal_double2(rocrand_state_mrg32k3a* state, double mean, d
double2 r
= rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
state2);
return double2 {
exp(mean + (stddev * r.x)),
exp(mean + (stddev * r.y))
};
return double2{exp(mean + (stddev * r.x)), exp(mean + (stddev * r.y))};
}

/**
Expand All @@ -485,8 +469,8 @@ double2 rocrand_log_normal_double2(rocrand_state_mrg32k3a* state, double mean, d
* \return Log-normally distributed \p float value
*/
#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
__forceinline__ __device__ __host__ float
rocrand_log_normal(rocrand_state_xorwow* state, float mean, float stddev)
__forceinline__ __device__ __host__
float rocrand_log_normal(rocrand_state_xorwow* state, float mean, float stddev)
{
typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;

Expand Down Expand Up @@ -525,10 +509,7 @@ float2 rocrand_log_normal2(rocrand_state_xorwow* state, float mean, float stddev
auto state2 = rocrand(state);

float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
return float2 {
expf(mean + (stddev * r.x)),
expf(mean + (stddev * r.y))
};
return float2{expf(mean + (stddev * r.x)), expf(mean + (stddev * r.y))};
}

/**
Expand All @@ -547,8 +528,8 @@ float2 rocrand_log_normal2(rocrand_state_xorwow* state, float mean, float stddev
* \return Log-normally distributed \p double value
*/
#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
__forceinline__ __device__ __host__ double
rocrand_log_normal_double(rocrand_state_xorwow* state, double mean, double stddev)
__forceinline__ __device__ __host__
double rocrand_log_normal_double(rocrand_state_xorwow* state, double mean, double stddev)
{
typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;

Expand All @@ -563,8 +544,7 @@ __forceinline__ __device__ __host__ double
auto state4 = rocrand(state);

double2 r = rocrand_device::detail::normal_distribution_double2(
uint4 { state1, state2, state3, state4 }
);
uint4{state1, state2, state3, state4});
bm_helper::save_double(state, r.y);
return exp(mean + (stddev * r.x));
}
Expand Down Expand Up @@ -593,12 +573,8 @@ double2 rocrand_log_normal_double2(rocrand_state_xorwow* state, double mean, dou
auto state4 = rocrand(state);

double2 r = rocrand_device::detail::normal_distribution_double2(
uint4 { state1, state2, state3, state4 }
);
return double2 {
exp(mean + (stddev * r.x)),
exp(mean + (stddev * r.y))
};
uint4{state1, state2, state3, state4});
return double2{exp(mean + (stddev * r.x)), exp(mean + (stddev * r.y))};
}

/**
Expand Down
Loading
Loading