CUDA：curand_uniform 全局内存对齐

Question

分析一个项目，我注意到对curand_uniform 的调用与global memory access 存在问题。例如使用内核创建的random number generator，如下所示：

__device__ curandState randGPU_d_state[200000];

__global__ void
initCurand(const unsigned long seed)
{
  int i = blockIdx.x * blockDimx. + threadIdx.x;
  if (i < 200000)
    curand_init(seed, i, 0, &randGPU_d_state[i]);
}

稍后在后续内核中通过以下方式访问，其中threadIdx.x < 200000：

float temp = curand_uniform(&randGPU_d_state[threadIdx.x]);

在将'Global Memory Access Pattern' 分析为'Global Load L2 Transactions/Access = 31.8, Ideal Transactions/Access = 8[ 12000 L2 transactions for 377 total executions ] ' 时，导致NVIDIA Visual Profiler 抛出此行。

事实上，我在同一行收到了 7 个这样的警告。

此外，如果我改用curand_normal，NVIDIA Visual Profiler 还会警告curand_normal.h 的第 310、312、313、315 和 316 行存在问题，Ideal Transactions/Access 的错误比率为 4 of 8。

我相信我正在访问合并的状态（虽然我没有将内存的细节打结在其中，但仍然访问了合并的状态变量），因此，为什么要预设这些不良比率？

Answer 1

您对合并内存访问的假设是不正确的。如果你运行这样的东西：

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

__device__ curandState randGPU_d_state[200000];

__global__ void
initCurand()
{
    printf("%ld\n", sizeof(randGPU_d_state[0]));
}

int main()
{
    initCurand<<<1,1>>>();
    cudaDeviceReset();
    return 0;
}

您会看到它为 sizeof(curandState) 打印 48。我不知道可以完全合并访问这么大的类型数组。