random - curand_uniform 不确定？

Question

我想以一种确定的方式在 CUDA 设备上生成伪随机数，如果我运行该程序两次，我希望得到完全相同的结果，因为该程序使用的是硬编码的种子。按照 nvidia 提供的示例：https ://docs.nvidia.com/cuda/curand/device-api-overview.html#device-api-example 我希望完全符合所描述的行为。

但我确实得到了不同的结果，多次运行完全相同的代码。正如我所描述的，有没有办法以确定的方式获得伪随机数？

以下示例代码显示了我的问题：

#include <iostream>

#include <cuda.h>
#include <curand_kernel.h>

__global__ void setup_kernel(curandState *state)
{
  auto id = threadIdx.x + blockIdx.x * blockDim.x;
  curand_init(123456, id, 0, &state[id]);
}

__global__ void draw_numbers(curandState *state, float* results)
{
  auto id = threadIdx.x + blockIdx.x * blockDim.x;
  // Copy state
  curandState localState = state[id % 1024];
  // Generate random number
  results[id] = curand_uniform(&localState);
  // Copy back state
  state[id % 1024] = localState;
}

int main(int argc, char* argv[])
{
  // Setup
  curandState* dStates;
  cudaMalloc((void **) &dStates, sizeof(curandState) * 1024);
  setup_kernel<<<1024, 1>>>(dStates);

  // Random numbers
  float* devResults;
  cudaMalloc((void **) &devResults, sizeof(float) * 16 * 1024);
  float *hostResults = (float*) calloc(16 * 1024, sizeof(float));

  // Call draw random numbers
  draw_numbers<<<1024, 16>>>(dStates, devResults);

  // Copy results
  cudaMemcpy(hostResults, devResults, 16 * 1024 * sizeof(float), cudaMemcpyDeviceToHost);

  // Output number 12345
  ::std::cout << "12345 is: " << hostResults[12345] << ::std::endl;

  return 0;
}

编译和运行代码会在我的机器上产生不同的输出：

$ nvcc -std=c++11 curand.cu && ./a.out && ./a.out && ./a.out 
12345 is: 0.8059
12345 is: 0.53454
12345 is: 0.382981

正如我所说，在此示例中，我希望输出相同的三倍。

标签： randomcudarandom-seed

Answer 1

curand_uniform确实取决于提供它的状态。

感谢 Robert Crovella 的评论，我现在看到错误在于依赖于线程执行顺序。只是不重用状态会导致相同的“随机”数字，当 draw_numbers 内核被多次调用时，这对我来说也不是一个选项。

我的猜测是，就我而言，最好的解决方案是只启动 1024 个线程（与设置的 curandState 一样多）并在每个线程中生成多个随机数（在我的示例中为 16 个/线程）。这样，我在程序内的多次调用中收到不同的随机数，但每次程序启动时都收到相同的数字。

更新代码：

#include <iostream>

#include <cuda.h>
#include <curand_kernel.h>

__global__ void setup_kernel(curandState *state)
{
  auto id = threadIdx.x + blockIdx.x * blockDim.x;
  curand_init(123456, id, 0, &state[id]);
}

__global__ void draw_numbers(curandState *state, float* results, int runs)
{
  auto id = threadIdx.x + blockIdx.x * blockDim.x;

  // Copy state
  curandState localState = state[id];

  // Generate random numbers
  for (int i = 0; i < runs; ++i)
  {
    results[id + i * 1024] = curand_uniform(&localState);
  }

  // Copy back state
  state[id] = localState;
}

int main(int argc, char* argv[])
{
  // Setup
  curandState* dStates;
  cudaMalloc((void **) &dStates, sizeof(curandState) * 1024);
  setup_kernel<<<1024, 1>>>(dStates);

  // Random numbers
  float* devResults;
  cudaMalloc((void **) &devResults, sizeof(float) * 16 * 1024);
  float *hostResults = (float*) calloc(16 * 1024, sizeof(float));

  // Call draw random numbers
  draw_numbers<<<16, 64>>>(dStates, devResults, 16);
  // Copy results
  cudaMemcpy(hostResults, devResults, 16 * 1024 * sizeof(float), cudaMemcpyDeviceToHost);

  // Output number 12345
  ::std::cout << "12345 is " << hostResults[12345];

  // Call draw random numbers (again)
  draw_numbers<<<16, 64>>>(dStates, devResults, 16);
  // Copy results
  cudaMemcpy(hostResults, devResults, 16 * 1024 * sizeof(float), cudaMemcpyDeviceToHost);

  // Output number 12345 again
  ::std::cout << " and " << hostResults[12345] << ::std::endl;

  return 0;
}

产生以下输出：

$ nvcc -std=c++11 curand.cu && ./a.out && ./a.out && ./a.out
12345 is 0.164181 and 0.295907
12345 is 0.164181 and 0.295907
12345 is 0.164181 and 0.295907

这正是我的用例。