[SOLVED] The type of random number generator in cuRAND kernels

The type of random number generator in cuRAND kernels

A typical example of random number generation in CUDA or pyCUDA is reported in the question How to generate random number inside pyCUDA kernel?, namely

#include <curand_kernel.h>

const int nstates = %(NGENERATORS)s;
__device__ curandState_t* states[nstates];

__global__ void initkernel(int seed)
{
    int tidx = threadIdx.x + blockIdx.x * blockDim.x;

    if (tidx < nstates) {
        curandState_t* s = new curandState_t;
        if (s != 0) {
            curand_init(seed, tidx, 0, s);
        }

        states[tidx] = s;
    }
}

__global__ void randfillkernel(float *values, int N)
{
    int tidx = threadIdx.x + blockIdx.x * blockDim.x;

    if (tidx < nstates) {
        curandState_t s = *states[tidx];
        for(int i=tidx; i < N; i += blockDim.x * gridDim.x) {
            values[i] = curand_uniform(&s);
        }
        *states[tidx] = s;
    }
}

Using this classical example, what is the random number generator that is activated (XORWOW, MTGP32, others)?

How is it possible to change the random number generator from within the kernel?

Solution

The default generator in the curand device API is XORWOW, as defined by

typedef struct curandStateXORWOW curandState_t;

in the device API header. You can change to another generator by substituting another state type to the curandInit call. Note that some generators require different arguments to the curandInit routine compared to the default.