在不同范围内的 Cuda 内核中生成随机数

Question

我正在尝试在 cuda 内核中生成随机数随机数。我希望从均匀分布和整数形式中生成随机数，从 1 到 8。每个线程的随机数都不同。可以生成随机数的范围也因一个线程而异。一个线程中范围的最大值可能低至 2，或者在另一个线程中可能高达 8，但不会高于该值。因此，我在下面提供了一个示例，说明我希望如何生成数字：

In thread#1 --> maximum of the range is 2 and so the random number should be between 1 and 2
In thread#2 --> maximum of the range is 6  and so the random number should be between 1 and 6
In thread#3 --> maximum of the range is 5 and so the random number should be between 1 and 5

等等……

Answer 1

编辑：我已经编辑了我的答案，以修复其他答案 (@tudorturcu) 和 cmets 中指出的一些缺陷。

1. 使用 CURAND 生成 uniform distribution 在 0.0 和 1.0 之间。注意：包含 1.0，排除 0.0
2. 然后将其乘以所需范围（最大值 - 最小值 值 + 0.999999)。
3. 然后添加偏移量（+ 最小值）。
4. 然后截断为整数。

您的设备代码中有这样的内容：

int idx = threadIdx.x+blockDim.x*blockIdx.x;
// assume have already set up curand and generated state for each thread...
// assume ranges vary by thread index
float myrandf = curand_uniform(&(my_curandstate[idx]));
myrandf *= (max_rand_int[idx] - min_rand_int[idx] + 0.999999);
myrandf += min_rand_int[idx];
int myrand = (int)truncf(myrandf);

你应该：

#include <math.h>

为truncf

这是一个完整的例子：

$ cat t527.cu
#include <stdio.h>
#include <curand.h>
#include <curand_kernel.h>
#include <math.h>
#include <assert.h>
#define MIN 2
#define MAX 7
#define ITER 10000000

__global__ void setup_kernel(curandState *state){

  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  curand_init(1234, idx, 0, &state[idx]);
}

__global__ void generate_kernel(curandState *my_curandstate, const unsigned int n, const unsigned *max_rand_int, const unsigned *min_rand_int,  unsigned int *result){

  int idx = threadIdx.x + blockDim.x*blockIdx.x;

  int count = 0;
  while (count < n){
    float myrandf = curand_uniform(my_curandstate+idx);
    myrandf *= (max_rand_int[idx] - min_rand_int[idx]+0.999999);
    myrandf += min_rand_int[idx];
    int myrand = (int)truncf(myrandf);

    assert(myrand <= max_rand_int[idx]);
    assert(myrand >= min_rand_int[idx]);
    result[myrand-min_rand_int[idx]]++;
    count++;}
}

int main(){

  curandState *d_state;
  cudaMalloc(&d_state, sizeof(curandState));
  unsigned *d_result, *h_result;
  unsigned *d_max_rand_int, *h_max_rand_int, *d_min_rand_int, *h_min_rand_int;
  cudaMalloc(&d_result, (MAX-MIN+1) * sizeof(unsigned));
  h_result = (unsigned *)malloc((MAX-MIN+1)*sizeof(unsigned));
  cudaMalloc(&d_max_rand_int, sizeof(unsigned));
  h_max_rand_int = (unsigned *)malloc(sizeof(unsigned));
  cudaMalloc(&d_min_rand_int, sizeof(unsigned));
  h_min_rand_int = (unsigned *)malloc(sizeof(unsigned));
  cudaMemset(d_result, 0, (MAX-MIN+1)*sizeof(unsigned));
  setup_kernel<<<1,1>>>(d_state);

  *h_max_rand_int = MAX;
  *h_min_rand_int = MIN;
  cudaMemcpy(d_max_rand_int, h_max_rand_int, sizeof(unsigned), cudaMemcpyHostToDevice);
  cudaMemcpy(d_min_rand_int, h_min_rand_int, sizeof(unsigned), cudaMemcpyHostToDevice);
  generate_kernel<<<1,1>>>(d_state, ITER, d_max_rand_int, d_min_rand_int, d_result);
  cudaMemcpy(h_result, d_result, (MAX-MIN+1) * sizeof(unsigned), cudaMemcpyDeviceToHost);
  printf("Bin:    Count: \n");
  for (int i = MIN; i <= MAX; i++)
    printf("%d    %d\n", i, h_result[i-MIN]);

  return 0;
}


$ nvcc -arch=sm_20 -o t527 t527.cu -lcurand
$ cuda-memcheck ./t527
========= CUDA-MEMCHECK
Bin:    Count:
2    1665496
3    1668130
4    1667644
5    1667435
6    1665026
7    1666269
========= ERROR SUMMARY: 0 errors
$

Answer 2

@Robert 的示例没有生成完美均匀分布（尽管生成了范围内的所有数字并且所有生成的数字都在范围内）。最小值和最大值都有 0.5 的概率在该范围内的其余数字中被选中。

在第 2 步，您应该乘以范围内的值数：（最大值 - 最小值 + 0.999999）。 *

在第 3 步，偏移量应为 (+ 最小值) 而不是 (+ 最小值 + 0.5)。

第 1 步和第 4 步保持不变。

*正如@Kamil Czerski 所说，1.0 包含在发行版中。添加 1.0 而不是 0.99999 有时会导致数字超出所需范围。