From 2aee15543d3825c91b8387337a0e7a16f9fbe745 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Fri, 25 Sep 2026 15:17:16 +0200 Subject: [PATCH 01/15] 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 02/15] 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 03/15] 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 04/15] 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 05/15] 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 06/15] 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 07/15] 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 08/15] 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 From 1c70bb2699923b803e919104eb72f3381cd56018 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Thu, 1 Oct 2026 15:26:42 +0200 Subject: [PATCH 09/15] Formatting. --- examples/cuda/curand.py | 14 ++++++++++++-- examples/hip/hiprand.py | 14 ++++++++++++-- 2 files changed, 24 insertions(+), 4 deletions(-) diff --git a/examples/cuda/curand.py b/examples/cuda/curand.py index f4de08f8f..5e2b44d45 100644 --- a/examples/cuda/curand.py +++ b/examples/cuda/curand.py @@ -50,7 +50,15 @@ def tune(): # initialize the curand state 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="nvcuda", compiler_options=compiler_options)[0] + 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] @@ -61,7 +69,9 @@ def tune(): # 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) + results, env = tune_kernel( + "generate_random", kernel_string, size, args, tune_params, lang="nvcuda", compiler_options=compiler_options + ) return results diff --git a/examples/hip/hiprand.py b/examples/hip/hiprand.py index cd5683f02..3abd93eb4 100644 --- a/examples/hip/hiprand.py +++ b/examples/hip/hiprand.py @@ -68,7 +68,15 @@ def tune(): # 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] + 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] @@ -79,7 +87,9 @@ def tune(): # 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) + results, env = tune_kernel( + "generate_random", kernel_string, size, args, tune_params, lang="HIP", compiler_options=compiler_options + ) return results From 21d176fa3812d9a383e60bb5536946fc32222475 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Thu, 15 Oct 2026 14:52:25 +0200 Subject: [PATCH 10/15] Adding an example that uses the host API of cuRAND. --- examples/cuda/curand.py | 3 +- examples/cuda/curand_host.py | 63 ++++++++++++++++++++++++++++++++++++ 2 files changed, 65 insertions(+), 1 deletion(-) create mode 100644 examples/cuda/curand_host.py diff --git a/examples/cuda/curand.py b/examples/cuda/curand.py index 5e2b44d45..335a9465d 100644 --- a/examples/cuda/curand.py +++ b/examples/cuda/curand.py @@ -1,6 +1,7 @@ #!/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""" +initialized on the device in a separate kernel, and updates that state after +drawing numbers""" import numpy from kernel_tuner import tune_kernel, run_kernel diff --git a/examples/cuda/curand_host.py b/examples/cuda/curand_host.py new file mode 100644 index 000000000..98d7fe6f2 --- /dev/null +++ b/examples/cuda/curand_host.py @@ -0,0 +1,63 @@ +#!/usr/bin/env python +"""Example showing how to tune a CUDA kernel that draws numbers generated +with the cuRAND host API""" + +import numpy +from kernel_tuner import tune_kernel +from cuda.bindings import driver +from nvmath.bindings import curand + +DRAWS_PER_THREAD = 16 + + +def tune(): + kernel_string = """ + __global__ void draw_random(float *output, float *random, int n) { + int i = blockIdx.x * block_size_x + threadIdx.x; + if (i < n) { + float sum = 0.0f; + #pragma unroll unroll_draws + for (int j = 0; j < DRAWS_PER_THREAD; j++) { + sum += random[(j * n) + i]; + } + + output[i] = sum; + } + } + """ + + size = 10_000_000 + compiler_options = ["--dopt=on"] + + output = numpy.zeros(size).astype(numpy.float32) + random = numpy.zeros(size * DRAWS_PER_THREAD).astype(numpy.float32) + n = numpy.int32(size) + args = [output, random, n] + + # use cuRAND to generate the random numbers on the GPU via the host API + err, stream = driver.cuStreamCreate(driver.CUstream_flags.CU_STREAM_NON_BLOCKING.value) + err, device_random = driver.cuMemAlloc(random.nbytes) + gen = curand.create_generator(curand.RngType.PSEUDO_PHILOX4_32_10) + curand.set_pseudo_random_generator_seed(gen, 42) + curand.set_stream(gen, int(stream)) + curand.generate_uniform(gen, int(device_random), size * DRAWS_PER_THREAD) + driver.cuStreamSynchronize(stream) + driver.cuMemcpyDtoH(random.ctypes.data, device_random, random.nbytes) + driver.cuMemFree(device_random) + curand.destroy_generator(gen) + driver.cuStreamDestroy(stream) + + tune_params = dict() + tune_params["DRAWS_PER_THREAD"] = [DRAWS_PER_THREAD] + tune_params["block_size_x"] = [32 * i for i in range(1, 33)] + tune_params["unroll_draws"] = [1, 2, 4, 8, 16] + + results, env = tune_kernel( + "draw_random", kernel_string, size, args, tune_params, lang="nvcuda", compiler_options=compiler_options + ) + + return results + + +if __name__ == "__main__": + tune() From 36bdf5a6fe46599504769c798d9d8a610fbc4606 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Thu, 15 Oct 2026 15:22:14 +0200 Subject: [PATCH 11/15] Fixing a bug in the example. --- examples/cuda/curand_host.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/examples/cuda/curand_host.py b/examples/cuda/curand_host.py index 98d7fe6f2..e75616fb2 100644 --- a/examples/cuda/curand_host.py +++ b/examples/cuda/curand_host.py @@ -35,6 +35,10 @@ def tune(): args = [output, random, n] # use cuRAND to generate the random numbers on the GPU via the host API + driver.cuInit(0) + driver.cuDeviceGet(0) + err, ctx = driver.cuDevicePrimaryCtxRetain(dev) + driver.cuCtxSetCurrent(ctx) err, stream = driver.cuStreamCreate(driver.CUstream_flags.CU_STREAM_NON_BLOCKING.value) err, device_random = driver.cuMemAlloc(random.nbytes) gen = curand.create_generator(curand.RngType.PSEUDO_PHILOX4_32_10) From 7c0a2fb703ddb044771f40105d77807a7e08fd37 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Thu, 15 Oct 2026 15:27:39 +0200 Subject: [PATCH 12/15] Fixed the last bug, runs on DAS6. --- examples/cuda/curand_host.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/examples/cuda/curand_host.py b/examples/cuda/curand_host.py index e75616fb2..95a39ea91 100644 --- a/examples/cuda/curand_host.py +++ b/examples/cuda/curand_host.py @@ -36,7 +36,7 @@ def tune(): # use cuRAND to generate the random numbers on the GPU via the host API driver.cuInit(0) - driver.cuDeviceGet(0) + err, dev = driver.cuDeviceGet(0) err, ctx = driver.cuDevicePrimaryCtxRetain(dev) driver.cuCtxSetCurrent(ctx) err, stream = driver.cuStreamCreate(driver.CUstream_flags.CU_STREAM_NON_BLOCKING.value) From 67731eb75ef9d94443da42acd4286fa1c69f6f64 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Fri, 16 Oct 2026 09:53:44 +0200 Subject: [PATCH 13/15] Correctness checks. --- examples/cuda/curand_host.py | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/examples/cuda/curand_host.py b/examples/cuda/curand_host.py index 95a39ea91..235496068 100644 --- a/examples/cuda/curand_host.py +++ b/examples/cuda/curand_host.py @@ -50,6 +50,12 @@ def tune(): driver.cuMemFree(device_random) curand.destroy_generator(gen) driver.cuStreamDestroy(stream) + driver.cuDevicePrimaryCtxRelease(dev) + + # Check that something was actually generated :) + assert abs(random.mean() - 0.5) < 1e-3 + + expected = random.reshape(DRAWS_PER_THREAD, size).sum(axis=0, dtype=numpy.float32) tune_params = dict() tune_params["DRAWS_PER_THREAD"] = [DRAWS_PER_THREAD] @@ -57,7 +63,7 @@ def tune(): tune_params["unroll_draws"] = [1, 2, 4, 8, 16] results, env = tune_kernel( - "draw_random", kernel_string, size, args, tune_params, lang="nvcuda", compiler_options=compiler_options + "draw_random", kernel_string, size, args, tune_params, lang="nvcuda", compiler_options=compiler_options, answer=[expected, None, None] ) return results From 3edaccc553753c15d408bf3a34615cb3b06af178 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Fri, 16 Oct 2026 10:04:32 +0200 Subject: [PATCH 14/15] Format. --- examples/cuda/curand_host.py | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/examples/cuda/curand_host.py b/examples/cuda/curand_host.py index 235496068..43772bdb8 100644 --- a/examples/cuda/curand_host.py +++ b/examples/cuda/curand_host.py @@ -63,7 +63,14 @@ def tune(): tune_params["unroll_draws"] = [1, 2, 4, 8, 16] results, env = tune_kernel( - "draw_random", kernel_string, size, args, tune_params, lang="nvcuda", compiler_options=compiler_options, answer=[expected, None, None] + "draw_random", + kernel_string, + size, + args, + tune_params, + lang="nvcuda", + compiler_options=compiler_options, + answer=[expected, None, None], ) return results From 4d9da2255873e4ffa147f07203cf5c5e5a045984 Mon Sep 17 00:00:00 2001 From: Alessio Sclocco Date: Fri, 16 Oct 2026 10:31:22 +0200 Subject: [PATCH 15/15] Adding an hip equivalent version for the rand host api. --- examples/hip/hiprand_host.py | 75 ++++++++++++++++++++++++++++++++++++ 1 file changed, 75 insertions(+) create mode 100644 examples/hip/hiprand_host.py diff --git a/examples/hip/hiprand_host.py b/examples/hip/hiprand_host.py new file mode 100644 index 000000000..6f53eee05 --- /dev/null +++ b/examples/hip/hiprand_host.py @@ -0,0 +1,75 @@ +#!/usr/bin/env python +"""Example showing how to tune a HIP kernel that draws numbers generated +with the hipRAND host API""" + +import numpy +from kernel_tuner import tune_kernel +from hip import hip, hiprand + +DRAWS_PER_THREAD = 16 + + +def tune(): + kernel_string = """ + __global__ void draw_random(float *output, float *random, int n) { + int i = blockIdx.x * block_size_x + threadIdx.x; + if (i < n) { + float sum = 0.0f; + #pragma unroll unroll_draws + for (int j = 0; j < DRAWS_PER_THREAD; j++) { + sum += random[(j * n) + i]; + } + + output[i] = sum; + } + } + """ + + size = 10_000_000 + + output = numpy.zeros(size).astype(numpy.float32) + random = numpy.zeros(size * DRAWS_PER_THREAD).astype(numpy.float32) + n = numpy.int32(size) + args = [output, random, n] + + # use hipRAND to generate the random numbers on the GPU via the host API + hip.hipInit(0) + hip.hipSetDevice(0) + err, stream = hip.hipStreamCreate() + err, device_random = hip.hipMalloc(random.nbytes) + err, gen = hiprand.hiprandCreateGenerator(hiprand.hiprandRngType.HIPRAND_RNG_PSEUDO_PHILOX4_32_10) + hiprand.hiprandSetPseudoRandomGeneratorSeed(gen, 42) + hiprand.hiprandSetStream(gen, stream) + hiprand.hiprandGenerateUniform(gen, device_random, random.size) + hip.hipStreamSynchronize(stream) + hip.hipMemcpy(random, device_random, random.nbytes, hip.hipMemcpyKind.hipMemcpyDeviceToHost) + hip.hipFree(device_random) + hiprand.hiprandDestroyGenerator(gen) + hip.hipStreamDestroy(stream) + + # Check that something was actually generated + assert abs(random.mean() - 0.5) < 1e-3 + + expected = random.reshape(DRAWS_PER_THREAD, size).sum(axis=0, dtype=numpy.float32) + + tune_params = dict() + tune_params["DRAWS_PER_THREAD"] = [DRAWS_PER_THREAD] + tune_params["block_size_x"] = [32 * i for i in range(1, 33)] + tune_params["unroll_draws"] = [1, 2, 4, 8, 16] + + results, env = tune_kernel( + "draw_random", + kernel_string, + size, + args, + tune_params, + lang="HIP", + answer=[expected, None, None], + atol=1e-3, + ) + + return results + + +if __name__ == "__main__": + tune()