I'm trying to port some MonteCarlo code to CUDA (this is my first CUDA experience).

The CPU code often calls a function which generates a random number:

float generate_random_number() {...}

Since all the existing code base runs fine on the GPU, that is basically the only thing to adjust.

My strategy is to put a curandState for each thread in global memory. Is this reasonable? Why/why not? Is the following code ok? Can I make it faster using shared memory or so?

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "curand_kernel.h"

__device__ curandState* random_generator_states;
__global__ void set_pointer(curandState* pointer) {
    random_generator_states = pointer;
__global__ void initialize_device(const unsigned int seed) {
    int idx = blockDim.x*blockIdx.x + threadIdx.x;
    curand_init(seed, idx, 0, &random_generator_states[idx]);
__device__ float generate_random_number() {
    int idx = blockDim.x*blockIdx.x + threadIdx.x;
    curandState localState = random_generator_states[idx];
    float RANDOM = curand_uniform(&localState);
    random_generator_states[idx] = localState;
    return RANDOM;

int main() {
    cudaError_t cudaStatus;
    size_t number_of_threads = 128;
    size_t number_of_blocks = 128;
    size_t blocks_times_threads = number_of_blocks * number_of_threads;
    size_t size = blocks_times_threads * sizeof(float);

    unsigned int seed = 0;

    // initialize memory for random generator states
    curandState* random_generator_states_pointer;
    cudaStatus = cudaMalloc(&random_generator_states_pointer, blocks_times_threads * sizeof(curandState));
    // initialize values for random generator states
    set_pointer<<< 1, 1>>>(random_generator_states_pointer);
    initialize_device<<< number_of_blocks, number_of_threads >>>(seed);

    // free GPU

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Browse other questions tagged or ask your own question.