From 2aee15543d3825c91b8387337a0e7a16f9fbe745 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Fri, 25 Sep 2026 15:17:16 +0200 Subject: [PATCH 1/8] Adding two examples for using RNG libraries on the GPU. Structure generated with Claude, and later modified manually. --- examples/README.rst | 6 ++++ examples/cuda/curand.py | 70 +++++++++++++++++++++++++++++++++++++ examples/hip/hiprand.py | 77 +++++++++++++++++++++++++++++++++++++++++ 3 files changed, 153 insertions(+) create mode 100644 examples/cuda/curand.py create mode 100644 examples/hip/hiprand.py diff --git a/examples/README.rst b/examples/README.rst index e2749882d..7f5cedf26 100644 --- a/examples/README.rst +++ b/examples/README.rst @@ -101,3 +101,9 @@ Code Generator -------------- [`CUDA `__] [`OpenCL `__] - use a Python function as a code generator + +Random Number Generation +------------------------- +[`CUDA `__] [`HIP `__] + - use curand/hiprand's device-side API to generate random numbers inside a tuned kernel + - initialize the generator state once from the host using run\_kernel, then reuse it across tuning diff --git a/examples/cuda/curand.py b/examples/cuda/curand.py new file mode 100644 index 000000000..f2075228b --- /dev/null +++ b/examples/cuda/curand.py @@ -0,0 +1,70 @@ +#!/usr/bin/env python +"""Example showing how to tune a CUDA kernel that draws from a curand state +initialized on the host, and updates that state after drawing numbers""" + +import numpy +from kernel_tuner import tune_kernel, run_kernel + +# curandStateXORWOW's fields (unsigned int d, v[5]; int boxmuller_flag, +# boxmuller_flag_double; float boxmuller_extra; double boxmuller_extra_double) +# sum to 44 bytes, padded to 48 for the trailing double's 8-byte alignment. +CURAND_STATE_SIZE = 48 + + +def tune(): + kernel_string = """ + #define DRAWS_PER_THREAD 16 + #include + + extern "C" __global__ void setup_kernel(curandState *state, unsigned long long seed, int n) { + int i = blockIdx.x * block_size_x + threadIdx.x; + if (i < n) { + curand_init(seed, i, 0, &state[i]); + } + } + + __global__ void generate_random(float *output, curandState *state, int n) { + int i = blockIdx.x * block_size_x + threadIdx.x; + if (i < n) { + curandState local_state = state[i]; + + float sum = 0.0f; + #pragma unroll unroll_draws + for (int j = 0; j < DRAWS_PER_THREAD; j++) { + sum += curand_uniform(&local_state); + } + + output[i] = sum; + state[i] = local_state; + } + } + """ + + size = 10_000_000 + n = numpy.int32(size) + seed = numpy.uint64(42) + compiler_options = ["-O3"] + + # curandState is opaque, host-side content is irrelevant, only its size matters + state = numpy.zeros(size * CURAND_STATE_SIZE, dtype=numpy.uint8) + + # initialize the curand state from a separate kernel, using a fixed block size + setup_params = {"block_size_x": 256} + state = run_kernel("setup_kernel", kernel_string, size, [state, seed, n], setup_params, lang="nvcuda", compiler_options=compiler_options)[0] + + output = numpy.zeros(size).astype(numpy.float32) + args = [output, state, n] + + tune_params = dict() + tune_params["block_size_x"] = [32 * i for i in range(33)] + tune_params["unroll_draws"] = [1, 2, 4, 8, 16] + + # note: each benchmarked launch of generate_random advances the curand state + # further along its sequence, since the kernel writes the updated state back + results, env = tune_kernel("generate_random", kernel_string, size, args, tune_params, lang="nvcuda", compiler_options=compiler_options) + + return results + + +if __name__ == "__main__": + tune() diff --git a/examples/hip/hiprand.py b/examples/hip/hiprand.py new file mode 100644 index 000000000..82b28a444 --- /dev/null +++ b/examples/hip/hiprand.py @@ -0,0 +1,77 @@ +#!/usr/bin/env python +"""Example showing how to tune a HIP kernel that draws from a hiprand state +initialized on the host, and updates that state after drawing numbers""" + +import numpy +from kernel_tuner import tune_kernel, run_kernel + +# hiprandStateXORWOW_t (rocrand_state_xorwow) wraps a single xorwow_state: +# unsigned int d, boxmuller_float_state, boxmuller_double_state; (12 bytes) +# float boxmuller_float; (4 bytes) +# double boxmuller_double; (8 bytes, needs 8-byte alignment) +# unsigned int x[5]; (20 bytes) +# which sums to 44 bytes, padded to 48 for struct alignment. Unlike CUDA's +# curandState, those Box-Muller fields are dropped from the state entirely +# when rocRAND is built with ROCRAND_DETAIL_BM_NOT_IN_STATE defined, shrinking +# the struct to 24 bytes -- so this layout can vary by build, not just by +# version. +HIPRAND_STATE_SIZE = 48 + + +def tune(): + kernel_string = """ + #define DRAWS_PER_THREAD 16 + #include + + extern "C" __global__ void setup_kernel(hiprandState *state, unsigned long long seed, int n) { + int i = blockIdx.x * block_size_x + threadIdx.x; + if (i < n) { + hiprand_init(seed, i, 0, &state[i]); + } + } + + __global__ void generate_random(float *output, hiprandState *state, int n) { + int i = blockIdx.x * block_size_x + threadIdx.x; + if (i < n) { + hiprandState local_state = state[i]; + + float sum = 0.0f; + #pragma unroll unroll_draws + for (int j = 0; j < DRAWS_PER_THREAD; j++) { + sum += hiprand_uniform(&local_state); + } + + output[i] = sum; + state[i] = local_state; + } + } + """ + + size = 10_000_000 + n = numpy.int32(size) + seed = numpy.uint64(42) + compiler_options = ["-O3"] + + # hiprandState is opaque, host-side content is irrelevant, only its size matters + state = numpy.zeros(size * HIPRAND_STATE_SIZE, dtype=numpy.uint8) + + # initialize the hiprand state once from a separate kernel, using a fixed block size + setup_params = {"block_size_x": 256} + state = run_kernel("setup_kernel", kernel_string, size, [state, seed, n], setup_params, lang="HIP", compiler_options=compiler_options)[0] + + output = numpy.zeros(size).astype(numpy.float32) + args = [output, state, n] + + tune_params = dict() + tune_params["block_size_x"] = [32 * i for i in range(33)] + tune_params["unroll_draws"] = [1, 2, 4, 8, 16] + + # note: each benchmarked launch of generate_random advances the hiprand state + # further along its sequence, since the kernel writes the updated state back + results, env = tune_kernel("generate_random", kernel_string, size, args, tune_params, lang="HIP", compiler_options=compiler_options) + + return results + + +if __name__ == "__main__": + tune() From 42599ac497994f46c0f35b1d337187c49efbf366 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Fri, 25 Sep 2026 15:26:13 +0200 Subject: [PATCH 2/8] Fixed a couple of bugs. --- examples/cuda/curand.py | 4 ++-- examples/hip/hiprand.py | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/examples/cuda/curand.py b/examples/cuda/curand.py index f2075228b..b2b4c7045 100644 --- a/examples/cuda/curand.py +++ b/examples/cuda/curand.py @@ -43,7 +43,7 @@ def tune(): size = 10_000_000 n = numpy.int32(size) seed = numpy.uint64(42) - compiler_options = ["-O3"] + compiler_options = ["--dopt"] # curandState is opaque, host-side content is irrelevant, only its size matters state = numpy.zeros(size * CURAND_STATE_SIZE, dtype=numpy.uint8) @@ -56,7 +56,7 @@ def tune(): args = [output, state, n] tune_params = dict() - tune_params["block_size_x"] = [32 * i for i in range(33)] + tune_params["block_size_x"] = [32 * i for i in range(1, 33)] tune_params["unroll_draws"] = [1, 2, 4, 8, 16] # note: each benchmarked launch of generate_random advances the curand state diff --git a/examples/hip/hiprand.py b/examples/hip/hiprand.py index 82b28a444..e3fcea764 100644 --- a/examples/hip/hiprand.py +++ b/examples/hip/hiprand.py @@ -63,7 +63,7 @@ def tune(): args = [output, state, n] tune_params = dict() - tune_params["block_size_x"] = [32 * i for i in range(33)] + tune_params["block_size_x"] = [32 * i for i in range(1, 33)] tune_params["unroll_draws"] = [1, 2, 4, 8, 16] # note: each benchmarked launch of generate_random advances the hiprand state From 323bb18b3dec0d36f6853b99a2d9f822714ff737 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Tue, 29 Sep 2026 11:50:47 +0200 Subject: [PATCH 3/8] Use the right cli option for the compiler. --- examples/cuda/curand.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/cuda/curand.py b/examples/cuda/curand.py index b2b4c7045..380dbaa3d 100644 --- a/examples/cuda/curand.py +++ b/examples/cuda/curand.py @@ -43,7 +43,7 @@ def tune(): size = 10_000_000 n = numpy.int32(size) seed = numpy.uint64(42) - compiler_options = ["--dopt"] + compiler_options = ["--dopt=on"] # curandState is opaque, host-side content is irrelevant, only its size matters state = numpy.zeros(size * CURAND_STATE_SIZE, dtype=numpy.uint8) From 86f3b2ab6ef9c84497616674161b59654a3c8dc9 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Tue, 29 Sep 2026 14:52:49 +0200 Subject: [PATCH 4/8] Fix the unroll_draws parameter for the setup. --- examples/cuda/curand.py | 2 +- examples/hip/hiprand.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/examples/cuda/curand.py b/examples/cuda/curand.py index 380dbaa3d..f4de08f8f 100644 --- a/examples/cuda/curand.py +++ b/examples/cuda/curand.py @@ -49,7 +49,7 @@ def tune(): state = numpy.zeros(size * CURAND_STATE_SIZE, dtype=numpy.uint8) # initialize the curand state from a separate kernel, using a fixed block size - setup_params = {"block_size_x": 256} + setup_params = {"block_size_x": 256, "unroll_draws": 1} state = run_kernel("setup_kernel", kernel_string, size, [state, seed, n], setup_params, lang="nvcuda", compiler_options=compiler_options)[0] output = numpy.zeros(size).astype(numpy.float32) diff --git a/examples/hip/hiprand.py b/examples/hip/hiprand.py index e3fcea764..4acd2478e 100644 --- a/examples/hip/hiprand.py +++ b/examples/hip/hiprand.py @@ -56,7 +56,7 @@ def tune(): state = numpy.zeros(size * HIPRAND_STATE_SIZE, dtype=numpy.uint8) # initialize the hiprand state once from a separate kernel, using a fixed block size - setup_params = {"block_size_x": 256} + setup_params = {"block_size_x": 256, "unroll_draws": 1} state = run_kernel("setup_kernel", kernel_string, size, [state, seed, n], setup_params, lang="HIP", compiler_options=compiler_options)[0] output = numpy.zeros(size).astype(numpy.float32) From 25bac2ca5561bcc5653a734cb9ab17f49f14cd73 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Tue, 29 Sep 2026 15:48:31 +0200 Subject: [PATCH 5/8] Wrong header path for hiprand. --- examples/hip/hiprand.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/hip/hiprand.py b/examples/hip/hiprand.py index 4acd2478e..f4b76b681 100644 --- a/examples/hip/hiprand.py +++ b/examples/hip/hiprand.py @@ -21,7 +21,7 @@ def tune(): kernel_string = """ #define DRAWS_PER_THREAD 16 - #include + #include extern "C" __global__ void setup_kernel(hiprandState *state, unsigned long long seed, int n) { int i = blockIdx.x * block_size_x + threadIdx.x; From 16924b0f6dfb87c178f8f7bbba6465039c3ff460 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Tue, 29 Sep 2026 16:26:31 +0200 Subject: [PATCH 6/8] Unfortunately hiprand and hiprtc are not compatible, so Claude ported the example to rocrand as a test. --- examples/hip/hiprand.py | 30 +++++++++++++++++------------- 1 file changed, 17 insertions(+), 13 deletions(-) diff --git a/examples/hip/hiprand.py b/examples/hip/hiprand.py index f4b76b681..2368cba77 100644 --- a/examples/hip/hiprand.py +++ b/examples/hip/hiprand.py @@ -1,11 +1,11 @@ #!/usr/bin/env python -"""Example showing how to tune a HIP kernel that draws from a hiprand state +"""Example showing how to tune a HIP kernel that draws from a rocRAND state initialized on the host, and updates that state after drawing numbers""" import numpy from kernel_tuner import tune_kernel, run_kernel -# hiprandStateXORWOW_t (rocrand_state_xorwow) wraps a single xorwow_state: +# rocrand_state_xorwow wraps a single xorwow_state: # unsigned int d, boxmuller_float_state, boxmuller_double_state; (12 bytes) # float boxmuller_float; (4 bytes) # double boxmuller_double; (8 bytes, needs 8-byte alignment) @@ -15,30 +15,34 @@ # when rocRAND is built with ROCRAND_DETAIL_BM_NOT_IN_STATE defined, shrinking # the struct to 24 bytes -- so this layout can vary by build, not just by # version. -HIPRAND_STATE_SIZE = 48 +ROCRAND_STATE_SIZE = 48 def tune(): + # The hiprand_kernel.h header pulls in the host API and , which + # does not compile with hiprtc (used by Kernel Tuner's HIP backend), so we + # use rocRAND's device headers directly. kernel_string = """ #define DRAWS_PER_THREAD 16 - #include + #include + #include - extern "C" __global__ void setup_kernel(hiprandState *state, unsigned long long seed, int n) { + extern "C" __global__ void setup_kernel(rocrand_state_xorwow *state, unsigned long long seed, int n) { int i = blockIdx.x * block_size_x + threadIdx.x; if (i < n) { - hiprand_init(seed, i, 0, &state[i]); + rocrand_init(seed, i, 0, &state[i]); } } - __global__ void generate_random(float *output, hiprandState *state, int n) { + __global__ void generate_random(float *output, rocrand_state_xorwow *state, int n) { int i = blockIdx.x * block_size_x + threadIdx.x; if (i < n) { - hiprandState local_state = state[i]; + rocrand_state_xorwow local_state = state[i]; float sum = 0.0f; #pragma unroll unroll_draws for (int j = 0; j < DRAWS_PER_THREAD; j++) { - sum += hiprand_uniform(&local_state); + sum += rocrand_uniform(&local_state); } output[i] = sum; @@ -52,10 +56,10 @@ def tune(): seed = numpy.uint64(42) compiler_options = ["-O3"] - # hiprandState is opaque, host-side content is irrelevant, only its size matters - state = numpy.zeros(size * HIPRAND_STATE_SIZE, dtype=numpy.uint8) + # rocrand_state_xorwow is opaque, host-side content is irrelevant, only its size matters + state = numpy.zeros(size * ROCRAND_STATE_SIZE, dtype=numpy.uint8) - # initialize the hiprand state once from a separate kernel, using a fixed block size + # initialize the rocRAND state once from a separate kernel, using a fixed block size setup_params = {"block_size_x": 256, "unroll_draws": 1} state = run_kernel("setup_kernel", kernel_string, size, [state, seed, n], setup_params, lang="HIP", compiler_options=compiler_options)[0] @@ -66,7 +70,7 @@ def tune(): tune_params["block_size_x"] = [32 * i for i in range(1, 33)] tune_params["unroll_draws"] = [1, 2, 4, 8, 16] - # note: each benchmarked launch of generate_random advances the hiprand state + # note: each benchmarked launch of generate_random advances the rocRAND state # further along its sequence, since the kernel writes the updated state back results, env = tune_kernel("generate_random", kernel_string, size, args, tune_params, lang="HIP", compiler_options=compiler_options) From c12d8c0126ad727aea31eca3b90bbf0f33656c6f Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Tue, 29 Sep 2026 16:31:01 +0200 Subject: [PATCH 7/8] Claude generated changes to still try to have this example working. It may be dropped later if impossible with hiprtc. --- examples/hip/hiprand.py | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/examples/hip/hiprand.py b/examples/hip/hiprand.py index 2368cba77..cd5683f02 100644 --- a/examples/hip/hiprand.py +++ b/examples/hip/hiprand.py @@ -25,7 +25,12 @@ def tune(): kernel_string = """ #define DRAWS_PER_THREAD 16 #include - #include + + // rocrand_uniform.h includes the host-side mtgp32 code, so do the (0, 1] conversion by hand + __device__ float uniform(rocrand_state_xorwow *state) { + const float inv = 2.3283064e-10f; // 2^-32 + return inv + rocrand(state) * inv; + } extern "C" __global__ void setup_kernel(rocrand_state_xorwow *state, unsigned long long seed, int n) { int i = blockIdx.x * block_size_x + threadIdx.x; @@ -42,7 +47,7 @@ def tune(): float sum = 0.0f; #pragma unroll unroll_draws for (int j = 0; j < DRAWS_PER_THREAD; j++) { - sum += rocrand_uniform(&local_state); + sum += uniform(&local_state); } output[i] = sum; @@ -54,7 +59,9 @@ def tune(): size = 10_000_000 n = numpy.int32(size) seed = numpy.uint64(42) - compiler_options = ["-O3"] + # rocrand_common.h includes , which clashes with hiprtc's built-in + # types; defining libstdc++'s include guard skips it (gcc's header layout) + compiler_options = ["-O3", "-D_GLIBCXX_MATH_H"] # rocrand_state_xorwow is opaque, host-side content is irrelevant, only its size matters state = numpy.zeros(size * ROCRAND_STATE_SIZE, dtype=numpy.uint8) From eb991f56234e0970ca2bae92edf2f14b4a183de5 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Tue, 29 Sep 2026 16:33:35 +0200 Subject: [PATCH 8/8] Update documentation. --- examples/README.rst | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/examples/README.rst b/examples/README.rst index 7f5cedf26..d0a9736bc 100644 --- a/examples/README.rst +++ b/examples/README.rst @@ -105,5 +105,6 @@ Code Generator Random Number Generation ------------------------- [`CUDA `__] [`HIP `__] - - use curand/hiprand's device-side API to generate random numbers inside a tuned kernel + - use curand's (CUDA) or rocRAND's (HIP) device-side API to generate random numbers inside a tuned kernel - initialize the generator state once from the host using run\_kernel, then reuse it across tuning + - tune the unrolling factor of the loop that draws the random numbers