pythoncudarandom-seedpycuda

Generating single random number in pyCuda kernel


I have seen many ways to generate an array of random numbers. but I want to generate a single random number. Is there any function as rand() in c++. I don't want a series of random numbers. I just need to generate a random number inside the kernel. is there any builtin function to generate random numbers? I have tried the given code below, but it not working.

import numpy as np
import pycuda.autoinit
from pycuda.compiler import SourceModule
from pycuda import gpuarray

code = """
    #include <curand_kernel.h>
       __device__ float getRand()
       {
          curandState_t s;
          curand_init(clock64(), 123456, 0, &s);
          return curand_uniform(&s);
       }

        __global__ void myRand(float *values)
        {
          values[0] = getRand();
        }
"""


mod = SourceModule(code)
myRand = mod.get_function("myRand")
gdata = gpuarray.zeros(2, dtype=np.float32)
myRand(gdata, block=(1,1,1), grid=(1,1,1))
print(gdata)

Errors are like:

/usr/local/cuda/bin/../targets/x86_64-linux/include/curand_poisson.h(548): error: this declaration may not have extern "C" linkage

/usr/local/cuda/bin/../targets/x86_64-linux/include/curand_discrete2.h(69): error: this declaration may not have extern "C" linkage

/usr/local/cuda/bin/../targets/x86_64-linux/include/curand_discrete2.h(78): error: this declaration may not have extern "C" linkage

/usr/local/cuda/bin/../targets/x86_64-linux/include/curand_discrete2.h(86): error: this declaration may not have extern "C" linkage

30 errors detected in the compilation of "kernel.cu".

Solution

  • The basic problem is that, by default, PyCUDA silently applies C linkage to all code compiled in a SourceModule. As the error is showing, cuRand requires C++ linkage, so getRand can't have C linkage.

    You can fix this either by changing these two lines:

    mod = SourceModule(code)
    myRand = mod.get_function("myRand")
    

    to

    mod = SourceModule(code, no_extern_c=True)
    myRand = mod.get_function("_Z6myRandPf")
    

    This disables C linkage, but does mean you need to supply the C++ mangled name to the get_function call. You will need to look at the verbose compiler output or compile the code outside of PyCUDA to get that name (for example Godbolt).

    Alternatively you can modify the code like this:

    import numpy as np
    import pycuda.autoinit
    from pycuda.compiler import SourceModule
    from pycuda import gpuarray
    
    code = """
           #include <curand_kernel.h>
    
           __device__ float getRand()
           {
              curandState_t s;
              curand_init(clock64(), 123456, 0, &s);
              return curand_uniform(&s);
           }
            
            extern "C" {
            __global__ void myRand(float *values)
            {
              values[0] = getRand();
            }
            }
    """
    
    
    mod = SourceModule(code, no_extern_c=True)
    myRand = mod.get_function("myRand")
    gdata = gpuarray.zeros(2, dtype=np.float32)
    myRand(gdata, block=(1,1,1), grid=(1,1,1))
    print(gdata)
    

    This leaves the kernel with C linkage, but doesn't touch the device function which is using cuRand.