Random Generator in OpenCL?

Question

Since there is no built-in random function in OpenCL (not that I know of, please correct me if this is not true). To generate a random list that put into kernel will not work for the purpose of my work. It has to be a random generator running on GPU (kernel). I intended to write my own function that generate random number in the range from 0 to 1. The code below is what i have run on CPU, and it seem to work well.

array_of_random_numbers = [0.0] * N
seed = 19073486328125
multiplier = 19073486328125
adder = 0
random_number = 19073486328125
modulus = 2**48.0
RAND_MAX = 2**48.0

for i in range( 0 , (N-1) ):
    random_number = (multiplier * random_number + adder ) % (modulus)
    print(multiplier * random_number)
    array_of_random_numbers[i] = (random_number / RAND_MAX)

However, I have hard time migrate the code to kernel since I cannot set the random_number in a loop and let it change over iterations.

kernel = """__kernel void get_new_rand(__global float* c)
{      
      random_number = (19073486328125 * 19073486328125 + 0) % (281474976710656);
      c[thread_id] = (random_number / 281474976710656.0);
}"""

Is there a way I can write the random generator on kernel?

Thank you in advance!

I intended to write my own function that generate random number in the range from 0 to 1. The code below is what i have run on CPU, and it seem to work well. However, I have hard time migrate the code to kernel since I cannot set the random_number in a loop and let it change over iterations.

Answer 1

The simplest approach is to use a linear congruence generator function (like you already have), that takes in a seed value and outputs a pseudo-random number that can be normalized to the range [0,1[. The remaining problem is how to get the seed.

Solution: you can pass a seed value as a global parameter from the host side, and change the seed value for every call of the kernel on the host side. This does not have any performance impact. Finally, to get different seeds for the different GPU threads, add the global ID to the seed passed over from the host before calling the LCG function.

This way you don't need any array of numbers stored. Also you have full control over the seed value on the host side, and it all remains fully deterministic.