在 cuda 内核中生成随机数答案

【问题标题】：generating random numbers within cuda kernel在 cuda 内核中生成随机数
【发布时间】：2013-01-13 03:35:07
【问题描述】：

我正在编写一个 cuda 程序，我需要在其中生成一个随机变量，该变量将通过遵循正态分布生成。我希望随机变量的值限制在 0 到 8 之间。所以我希望在内核函数中生成随机变量，然后将随机变量结果用于进一步使用。我打算为此目的使用 cuRAND 库。我一直在尝试使用 curand_normal 设备 api 来生成值，但没有任何成功。如果有人可以向我提供内核函数代码，那将非常有帮助。感谢您的所有帮助。

下面提供的代码是我在 gpu 中搜索的 cpu 实现：

  #include "stdafx.h"
    #include <iostream>
    #include <random>

    using namespace std;
    int _tmain(int argc, _TCHAR* argv[])
    {
        const int nrolls=10000;  // number of experiments
        const int nstars=100;    // maximum number of stars to distribute
        int i;
        default_random_engine generator;
        normal_distribution<double> distribution(0.0,3);


       for (i=0;i<=nstars;i++)
       {   int number = distribution(generator);
           printf("%d\n\n",number);
        }


        return 0;
    }

我想补充一点，我不懂 C++，我只是按照我在其他网站上看到的其他代码编写了这个程序。谢谢。

【问题讨论】：

你为什么不发布你尝试过的方法没有成功？这通常是个好主意。 “为我编写代码”类型的问题不太可能得到好的结果。你看过设备 API 示例here 吗？它提供了一个完整的程序，其中一个generate_uniform_kernel 选项应该非常接近您的要求。
嗨！你好！感谢您的支持。我实际上浏览了设备 API 示例并复制了那里给出的几乎一些东西，但仅使用 XORWOW 随机生成器，但结果不是很有说服力，而且只给出了分数。我不确定如何在 0 到 8 的范围内获得正态分布后的随机变量。另一件事是我尝试了 here 中提到的程序，我在顶部设备函数中将 curand_uniform 更改为 curand_normal 并且我我得到了一些结果。
你想要一个离散的均匀分布，取值 (0, 1, 2, 3, 4, 5, 6, 7, 8)（即整数）还是你想要一个连续的值取值介于 0.0 和 8.0 之间的均匀分布（即浮点数。）
我想要从 0 到 8 的离散数，平均值为 0
我不知道如何创建一个具有准确统计数据且占据有限范围的正态分布。你知道如何？你如何处理分布的尾部？我猜你的意思是零是指你只想要以零为中心的正态分布的正半部分？显然，如果范围从零扩展到正值，则结果分布的平均值不能为零。对我来说，这是一个 CUDA 问题并不明显。您是否有代表您正在尝试做的非 CUDA（基于 CPU）的实现？

标签： cuda gpgpu normal-distribution

【解决方案1】：

这是对this code 的改编，它将产生一组近似“正常”分布的随机数，可以取大约 0 到 8 之间的离散值。我不明白 cmets 中要求有一个范围的请求0 到 8，平均值为 0。

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <curand_kernel.h>
#include <math.h>
#define SCALE 2.0
#define SHIFT 4.5
#define DISCRETE
#define BLOCKS 1024
#define THREADS 512

#define CUDA_CALL(x) do { if((x) != cudaSuccess) { \
    printf("Error at %s:%d\n",__FILE__,__LINE__); \
    return EXIT_FAILURE;}} while(0)

__global__ void setup_kernel(curandState *state)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    /* Each thread gets different seed, a different sequence
       number, no offset */
    curand_init(7+id, id, 0, &state[id]);
}



__global__ void generate_normal_kernel(curandState *state,
                                int *result)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    float x;
    /* Copy state to local memory for efficiency */
    curandState localState = state[id];
    /* Generate pseudo-random uniforms */
    for(int n = 0; n < 10; n++) {
        x = (curand_normal(&localState) * SCALE)+SHIFT;
        /* Discretize */
#if defined DISCRETE
        x = truncf(x);
#endif
    }
    /* Copy state back to global memory */
    state[id] = localState;
    /* Store last generated result per thread */
    result[id] = (int) x;
}


int main(int argc, char *argv[])
{
    int i;
    unsigned int total;
    curandState *devStates;
    int *devResults, *hostResults;
    int device;
    struct cudaDeviceProp properties;

    CUDA_CALL(cudaGetDevice(&device));
    CUDA_CALL(cudaGetDeviceProperties(&properties,device));


    /* Allocate space for results on host */
    hostResults = (int *)calloc(THREADS * BLOCKS, sizeof(int));

    /* Allocate space for results on device */
    CUDA_CALL(cudaMalloc((void **)&devResults, BLOCKS * THREADS *
              sizeof(int)));
    /* Set results to 0 */
    CUDA_CALL(cudaMemset(devResults, 0, THREADS * BLOCKS *
              sizeof(int)));

    /* Allocate space for prng states on device */
    CUDA_CALL(cudaMalloc((void **)&devStates, THREADS * BLOCKS *
                  sizeof(curandState)));

    /* Setup prng states */
    setup_kernel<<<BLOCKS, THREADS>>>(devStates);


    /* Generate and use uniform pseudo-random  */
    generate_normal_kernel<<<BLOCKS, THREADS>>>(devStates, devResults);

    /* Copy device memory to host */
    CUDA_CALL(cudaMemcpy(hostResults, devResults, BLOCKS * THREADS *
        sizeof(int), cudaMemcpyDeviceToHost));

    /* Show result */
    if (THREADS*BLOCKS > 20){
      printf("First 20 stored results:\n");
      for (i=0; i<20; i++)
        printf("%d\n", hostResults[i]);
      }

    total = 0;
    for(i = 0; i < BLOCKS * THREADS; i++) {
        total += hostResults[i];
    }
    printf("Results mean = %f\n", (total/(1.0*BLOCKS*THREADS)));



    /* Cleanup */
    CUDA_CALL(cudaFree(devStates));
    CUDA_CALL(cudaFree(devResults));
    free(hostResults);
    return EXIT_SUCCESS;
}

您也可以轻松修改此代码以生成连续值正态分布（浮点数）。

正态分布的两个参数是均值和标准差。这些使用 SHIFT 和 SCALE 参数表示。 SHIFT 将平均值从零移动。 SCALE 修改标准偏差（从 1.0，到 SCALE 指示的任何值）。因此，您可以使用 SHIFT 和 SCALE 参数来获得您想要的分布。请注意，随机数生成器的实值输出的截断会影响统计信息。您可以通过调整 SCALE 或 SHIFT 来进行调整，或者您可以从 truncf() 切换到某种形式的舍入。

你可以编译这个：

nvcc -arch=sm_20 -o uniform uniform.cu

假设您拥有 cc2.0 或更高版本的 GPU。

如果没有，编译时使用：

nvcc -o uniform uniform.cu

在这种情况下 double 被降级为 float 的编译器警告可以忽略。

THREADS 和BLOCKS 是机器范围内的任意选择。您可以修改这些以适合您自己代码的特定启动配置。

【讨论】：