CUDA CURAND 是否容易受到数据竞争的影响？答案

【问题标题】：Is CUDA CURAND susceptible to data races?CUDA CURAND 是否容易受到数据竞争的影响？
【发布时间】：2017-09-12 01:49:22
【问题描述】：

我从我的 Setup() 内核生成 1 个 256 个线程块，以设置具有 256 个 CURAND 状态的数组 RNGstates：

__global__ void Setup(curandState *RNGstates, long seed) {
    int tid = threadIdx.x;
    curand_init(seed, tid, 0, &RNGstates[tid]);
}

现在，我从我的 Generate() 内核中生成 1000 个 256 个线程的块，用 256,000 个随机数填充数组 result。但是，我只使用 RNGstates 的 256 个状态，这样每个状态将被 1000 个线程访问（每个块一个）：

__global__ void Generate(curandState *RNGstates, float *result) {
    int tid = blockIdx.x*blockDim.x + threadIdx.x;
    float rnd = curand_uniform(&RNGstates[threadIdx.x]);
    result[tid] = rnd;
}

我知道调用curand_uniform() 会以某种方式更新状态，所以我推测正在发生一些写操作。

那么当映射到 256 个 CURAND 状态中的每一个的 1000 个线程尝试通过 curand_uniform() 隐式更新状态时，我是否应该担心发生数据争用？这会影响我的随机数的质量（例如获得频繁的重复值）吗？

非常感谢。

【问题讨论】：

标签： random cuda gpgpu race-condition

【解决方案1】：

我认为共享状态肯定会影响质量。重复值是共享状态的最佳情况。数据竞赛可能会彻底毁掉各州。

您可以为每个线程保留一个状态。

当使用 1000 个块时，您的案例需要 256,000 个状态。代码应该是这样的

__global__ void Setup(curandState *RNGstates, long seed) {
  int tid = blockIdx.x*blockDim.x + threadIdx.x;
  curand_init(seed, tid, 0, &RNGstates[tid]);
}

和

__global__ void Generate(curandState *RNGstates, float *result) {
  int tid = blockIdx.x*blockDim.x + threadIdx.x;
  float rnd = curand_uniform(&RNGstates[tid]);
  result[tid] = rnd;
}

为了减少多个块的内存需求，您可以将#block 限制为一个较小的数字，并为每个线程生成多个随机数，而不是每个线程生成 1 个随机数。

__global__ void generate_uniform_kernel(curandState *state, 
                                unsigned int *result)
{
    int id = threadIdx.x + blockIdx.x * 64;
    unsigned int count = 0;
    float x;
    /* Copy state to local memory for efficiency */
    curandState localState = state[id];
    /* Generate pseudo-random uniforms */
    for(int n = 0; n < 10000; n++) {
        x = curand_uniform(&localState);
        /* Check if > .5 */
        if(x > .5) {
            count++;
        }
    }
    /* Copy state back to global memory */
    state[id] = localState;
    /* Store results */
    result[id] += count;
}

有关如何处理多个块的完整示例，请参阅 cuRAND 参考手册中的 Device API Examples 部分。

【讨论】：

是的，但理想情况下，如果可以的话，我希望尽量少使用内存，那么 256 个州仍然可以安全工作吗？
@MiloChen 不安全。为了节省内存，您可以限制#blocks。我已经更新了我的答案。

【解决方案2】：

您也可以使用 curandStateMtgp32_t，每个块只需要一个（如果每个块最多有 256 个线程）http://docs.nvidia.com/cuda/curand/device-api-overview.html#bit-generation-1

【讨论】：