使用cuRAND在函数内生成随机数,无需预先分配

Question

是否可以在设备功能中生成随机数而无需预先分配所有状态？我想在"实时"中生成和使用它们.我需要将它们用于蒙特卡罗模拟最适合此目的的是什么？下面生成的数字是单精度是否可以使它们具有双精度？

#include <iostream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <curand_kernel.h>

__global__ void cudaRand(float *d_out, unsigned long seed)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    curandState state;
    curand_init(seed, i, 0, &state);
    d_out[i] = curand_uniform(&state);
}

int main(int argc, char** argv)
{
    size_t N = 1 << 4;
    float *v = new float[N];

    float *d_out;
    cudaMalloc((void**)&d_out, N * sizeof(float));

    // generate random numbers
    cudaRand << < 1, N >> > (d_out, time(NULL));

    cudaMemcpy(v, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);

    for (size_t i = 0; i < N; i++)
    {
        printf("out: %f\n", v[i]);
    }

    cudaFree(d_out);
    delete[] v;

    return 0;
}

---

UPDATE

#include <iostream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <curand_kernel.h>
#include <ctime>

__global__ void cudaRand(double *d_out)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    curandState state;
    curand_init((unsigned long long)clock() + i, 0, 0, &state);

    d_out[i] = curand_uniform_double(&state);
}

int main(int argc, char** argv)
{
    size_t N = 1 << 4;
    double *h_v = new double[N];

    double *d_out;
    cudaMalloc((void**)&d_out, N * sizeof(double));

    // generate random numbers
    cudaRand << < 1, N >> > (d_out);

    cudaMemcpy(h_v, d_out, N * sizeof(double), cudaMemcpyDeviceToHost);

    for (size_t i = 0; i < N; i++)
        printf("out: %f\n", h_v[i]);

    cudaFree(d_out);
    delete[] h_v;

    return 0;
}

Answer 1

我是如何处理过去类似的情况,内部__device__/ __global__功能:

int tId = threadIdx.x + (blockIdx.x * blockDim.x);
curandState state;
curand_init((unsigned long long)clock() + tId, 0, 0, &state);

double rand1 = curand_uniform_double(&state);
double rand2 = curand_uniform_double(&state);

所以只是curand_uniform_double用于生成随机双打,而且我相信你不希望所有线程使用相同的种子,这就是我试图通过使用clock() + tId而实现的.这样,在两个线程中的任何一个中具有相同rand1/ 的几率rand2接近于零.

编辑:

但是,根据以下评论,提出的方法可能 会导致偏见的结果:

• JackOLantern向我指出了这部分文档:

  用不同种子产生的序列通常不具有统计学相关的值,但是种子的一些选择可以给出统计学相关的序列.

• 还有一个致力于如何提高性能的devtalk线程curand_init,其中提出的加速语言初始化的解决方案是:

  你可以做的一件事是为每个线程使用不同的种子,固定的子序列为0,偏移量为0.

  但同一张海报后来说明:

  缺点是你在线程之间失去了一些很好的数学属性.在从种子初始化生成器状态的哈希函数与生成器的周期性之间可能存在不良交互.如果发生这种情况,您可能会获得两个线程,其中一些种子具有高度相关 我不知道这样的问题,即使它们确实存在,也很可能很少见.

因此,基本上由您决定是否需要更好的性能(如我所做)或1000%无偏见的结果.如果这是你想要的,那么JackOLantern提出的解决方案是正确的,即初始化curand为:

curand_init((unsigned long long)clock(), tId, 0, &state)

然而,使用not 0for value offset和subsequence参数会降低性能.有关这些参数的更多信息,您可以查看此SO线程以及curand文档.

我看到JackOLantern在评论中说:

我想说从同一个内核调用curand_init和curand_uniform_double是不可推荐的,原因有两个......第二,curand_init初始化伪随机数生成器并设置它的所有参数,所以我担心你的方法会有点慢.

我在几篇页面的论文中处理了这个问题,尝试了各种方法在每个线程中获取不同的随机数,并在每个线程中创建curandState结果对我来说是最可行的解决方案.我需要在每个线程中生成~10个随机数,其中我尝试过:

• 开发我自己的简单随机数生成器(线性同余生成器),其初始化基本上是免费的,但是,在生成数字时性能受到很大影响,因此最终curandState在每个线程中都证明是优越的,
• 预先分配curandStates并重新使用它们 - 这是内存繁重的,当我减少预分配状态的数量时,我必须使用非零值offset/ subsequence参数curand_uniform_double来消除偏差,导致生成数字时性能下降.

因此使得深入分析后,我决定确实调用curand_init并curand_uniform_double在每个线程.唯一的问题是这些状态占用的注册表数量,所以我必须小心块大小不超过每个块可用的最大注册数.

多数民众赞成我提供的解决方案,我终于能够测试,它在我的机器/ GPU上工作得很好.我在上面的问题中运行UPDATE部分的代码,并在控制台中正确显示了16个不同的随机数.因此,我建议您在执行内核后正确执行错误检查,以查看内部出了什么问题.这个SO主题很好地涵盖了这个主题.