[rocm-libraries] ROCm/rocm-libraries#12616 (commit dda0b49)

perf(rocrand): add __restrict__ to pointers to preserve
 philox4_32_10 f64 benchmark performance (#12616)

## Motivation

rocRAND device benchmark (rocrand-device) shows a ~13.35% performance
drop on MI300X (gfx942) for the normal distribution with the
philox4_32_10 engine and f64 type when comparing the WW38 staging
compiler build against the WW38 base build.

The solution is to add __restrict__ on the states and data parameters of
generate_kernel() in benchmark_device_api.hpp. This avoids LDS usage and
an increase in VGPRs for this benchmark that limits the blocksize to 128
instead of 256, resulting in uneven work distribution.

JIRA ID : AIPRIMS-262

## Technical Details

Adds __restrict__ on the states and data parameters of generate_kernel()
in benchmark_device_api.hpp.

## Test Plan

Unit tests pass with no performance regressions.

## Test Result

Unit tests pass with no performance regressions.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Author: amd-jmahovsky

CHANGELOG.md

@@ -3,6 +3,12 @@
 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 5.2.0 for ROCm 10.2.0
+
+### Optimized
+
+* Added `__restrict__` qualifiers to the device API benchmark, enabling better compiler optimization.
+
 ## rocRAND 5.1.0 for ROCm 10.1.0
 
 ### Optimized

CMakeLists.txt

@@ -183,7 +183,7 @@ if(BUILD_ADDRESS_SANITIZER AND BUILD_SHARED_LIBS)
 endif()
 
 # Set version variables
-rocm_setup_version( VERSION "5.1.0" )
+rocm_setup_version( VERSION "5.2.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}")

benchmark/benchmark_device_api.hpp

@@ -335,7 +335,7 @@ struct unrolled
 
 template<typename EngineState, typename T, typename Generator>
 __global__ __launch_bounds__(get_max_block_size<EngineState>())
-void generate_kernel(EngineState* states,
+void generate_kernel(EngineState* __restrict__ states,
                      T* __restrict__ data,
                      const size_t size,
                      Generator    generator)
@@ -397,7 +397,7 @@ struct runner
 
 template<typename T, typename Generator>
 __global__ __launch_bounds__(RAND_DEFAULT_MAX_BLOCK_SIZE)
-void generate_kernel(rand_state_mtgp32_t* states,
+void generate_kernel(rand_state_mtgp32_t* __restrict__ states,
                      T* __restrict__ data,
                      const size_t size,
                      Generator    generator)
@@ -587,7 +587,7 @@ void init_scrambled_sobol_kernel(EngineState* states,
 // generate_kernel for the normal and scrambled sobol generators
 template<typename EngineState, typename T, typename Generator>
 __global__ __launch_bounds__(RAND_DEFAULT_MAX_BLOCK_SIZE)
-void generate_sobol_kernel(EngineState* states,
+void generate_sobol_kernel(EngineState* __restrict__ states,
                            T* __restrict__ data,
                            const size_t size,
                            Generator    generator)