rocRAND: random numbers inside kernel function

[LucasSnatiago]

I'm doing a project using ROCm and I'm needing to get some random numbers inside a kernel function.

Really simple example:

__global__ void doSomething() {
    int idx = threadIdx.x;
    int value = roc_rand() % 1000;

    // Do more calculations
}

How can I create this function?
Right now, just as a placeholder I'm creating a big vector on cpu and write every single position with a random number that comes from C rand() function. Then I hipMemcpy it from host to the device. This vector is really big and I'm basically wasting a lot of my vram just to hold this vector.

[Snektron]

Hi, for this you can use the rocRAND device API. For example:

#include <hip/hip_runtime.h>
#include <rocrand/rocrand_kernel.h>

__global__
void test() {
	uint tid = blockDim.x * blockIdx.x + threadIdx.x;
	rocrand_state_xorwow state;
	rocrand_init(123, tid, 0, &state);

	for (int i = 0; i < 3; ++i) {
		const auto value = rocrand(&state);
		printf("thread %d, index %u: %u\n", tid, i, value);
	}
}

int main() {
	test<<<dim3(1), dim3(32)>>>();
	hipDeviceSynchronize();
}

After compiling as follows (assuming you have rocRAND installed in its default installation location):

$ hipcc -otest test.hip -std=c++17 -lrocrand

When executed, it prints something like

thread 0, index 0: 2008776621
thread 1, index 0: 1048137452
thread 2, index 0: 1940051228
[ ... ]
thread 27, index 2: 2838558949
thread 28, index 2: 1634970841
thread 29, index 2: 2010177377
thread 30, index 2: 3146023297
thread 31, index 2: 3830475989

You can find the other rocrand device functions in the rocRAND API docs, under 'Device Functions'.

On a side note, if you're using modulo for integers, beware of modulo bias: The first 2**32 % 1000 values will be sampled more often than the others. See this blog for a better explanation (or simply google for modulo bias).

[LucasSnatiago]

Thank you!

[LucasSnatiago]

I'm completely new to GPGPU. I was looking at the rocRAND documentation and I have no idea how to use it. Is there any plans to add some examples in the rocRAND API docs, just like your example?

[Snektron]

Yep agreed, those those should be expanded a little.

[Naraenda]

Tracking a) new example request and b) rocRAND documentation improvements internally.