Cuda 随机数生成

Posted

技术标签:

【中文标题】Cuda 随机数生成【英文标题】:Cuda Random Number Generation 【发布时间】:2013-02-21 06:06:10 【问题描述】:

我想知道通过使用 curand 或其他东西来生成一个 0 到 49k 之间的伪随机数的最佳方法是什么,这对于每个线程都是相同的。

我更喜欢在内核中生成随机数,因为我必须一次生成一个,但大约 10k 次。

我可以使用介于 0.0 和 1.0 之间的浮点数,但我不知道如何使我的 PRN 可用于所有线程,因为大多数帖子和示例都展示了如何为每个线程设置不同的 PRN。

谢谢

【问题讨论】:

【参考方案1】:

可能你只需要研究curand documentation,尤其是device API。为每个线程获取相同序列的关键是为每个线程创建状态(大多数示例都这样做),然后将相同的序列号传递给每个线程的 init 函数。在curand_init中,参数顺序如下:

curand_init(seed, subsequence number, offset, state)

通过为每个 init 调用设置相同的种子,我们为每个线程生成相同的序列。通过将子序列和偏移数设置为相同,我们在该序列中为每个线程选择相同的起始值。

下面是演示代码:

// compile with: nvcc -arch=sm_20 -lcurand -o t89 t89.cu
#include <stdio.h>
#include <curand.h>
#include <curand_kernel.h>

#define SCALE 49000
#define DSIZE 5000
#define nTPB 256

#define cudaCheckErrors(msg) \
    do  \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess)  \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
         \
     while (0)

__device__ float getnextrand(curandState *state)

  return (float)(curand_uniform(state));


__device__ int getnextrandscaled(curandState *state, int scale)

  return (int) scale * getnextrand(state);



__global__ void initCurand(curandState *state, unsigned long seed)
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    curand_init(seed, 0, 0, &state[idx]);


__global__ void testrand(curandState *state, int *a1, int *a2)
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    a1[idx] = getnextrandscaled(&state[idx], SCALE);
    a2[idx] = getnextrandscaled(&state[idx], SCALE);


int main() 

    int *h_a1, *h_a2, *d_a1, *d_a2;
    curandState *devState;

    h_a1 = (int *)malloc(DSIZE*sizeof(int));
    if (h_a1 == 0) printf("malloc fail\n"); return 1;
    h_a2 = (int *)malloc(DSIZE*sizeof(int));
    if (h_a2 == 0) printf("malloc fail\n"); return 1;
    cudaMalloc((void**)&d_a1, DSIZE * sizeof(int));
    cudaMalloc((void**)&d_a2, DSIZE * sizeof(int));
    cudaMalloc((void**)&devState, DSIZE * sizeof(curandState));
    cudaCheckErrors("cudamalloc");



     initCurand<<<(DSIZE+nTPB-1)/nTPB,nTPB>>>(devState, 1);
     cudaDeviceSynchronize();
     cudaCheckErrors("kernels1");
     testrand<<<(DSIZE+nTPB-1)/nTPB,nTPB>>>(devState, d_a1, d_a2);
     cudaDeviceSynchronize();
     cudaCheckErrors("kernels2");
     cudaMemcpy(h_a1, d_a1, DSIZE*sizeof(int), cudaMemcpyDeviceToHost);
     cudaMemcpy(h_a2, d_a2, DSIZE*sizeof(int), cudaMemcpyDeviceToHost);
     cudaCheckErrors("cudamemcpy");
     printf("1st returned random value is %d\n", h_a1[0]);
     printf("2nd returned random value is %d\n", h_a2[0]);

     for (int i=1; i< DSIZE; i++)
       if (h_a1[i] != h_a1[0]) 
         printf("mismatch on 1st value at %d, val = %d\n", i, h_a1[i]);
         return 1;
         
       if (h_a2[i] != h_a2[0]) 
         printf("mismatch on 2nd value at %d, val = %d\n", i, h_a2[i]);
         return 1;
         
       
     printf("thread values match!\n");


【讨论】:

以上是关于Cuda 随机数生成的主要内容,如果未能解决你的问题,请参考以下文章

CUDA 的 Mersenne Twister 用于任意数量的线程

如何在 pyCUDA 内核中生成随机数?

Qt 随机数的生成

MATLAB的简单随机生成函数

详解随机数的生成

Shell进阶脚本-生成随机数( $RANDOM )