在 CUDA 中访问全局内存很慢

Question

我有一个 CUDA 内核对局部变量（在寄存器中）进行一些计算，计算完成后，它的值被写入全局数组 p:

__global__ void dd( float* p, int dimX, int dimY, int dimZ  )
{
    int 
        i = blockIdx.x*blockDim.x + threadIdx.x,
        j = blockIdx.y*blockDim.y + threadIdx.y,
        k = blockIdx.z*blockDim.z + threadIdx.z,
        idx = j*dimX*dimY + j*dimX +i;   

    if (i >= dimX || j >= dimY || k >= dimZ)
    {
        return;
    }   

    float val = 0;

    val = SomeComputationOnVal();

     p[idx ]=  val;
    __syncthreads();    

}

不幸的是，这个函数执行起来很慢。

但是，如果我这样做，它运行得非常快：

  __global__ void dd(   float* p, int dimX, int dimY, int dimZ  )
    {
        int 
            i = blockIdx.x*blockDim.x + threadIdx.x,
            j = blockIdx.y*blockDim.y + threadIdx.y,
            k = blockIdx.z*blockDim.z + threadIdx.z,
            idx = j*dimX*dimY + j*dimX +i;   

        if (i >= dimX || j >= dimY || k >= dimZ)
        {
            return;
        }   

        float val = 0;

        //val = SomeComputationOnVal();

         p[idx ]=  val;
        __syncthreads();    

    }

如果我这样做它也运行得非常快:

__global__ void dd( float* p, int dimX, int dimY, int dimZ  )
{
    int 
        i = blockIdx.x*blockDim.x + threadIdx.x,
        j = blockIdx.y*blockDim.y + threadIdx.y,
        k = blockIdx.z*blockDim.z + threadIdx.z,
        idx = j*dimX*dimY + j*dimX +i;   

    if (i >= dimX || j >= dimY || k >= dimZ)
    {
        return;
    }   

    float val = 0;

    val = SomeComputationOnVal();

  //   p[idx ]=  val;
    __syncthreads();    

}

所以我很困惑，不知道如何解决这个问题。我已经使用 NSight step in，并没有发现访问冲突。

这是我启动内核的方式 (dimX:924; dimY: 16: dimZ: 1120):

dim3 
      blockSize(8,16,2),
      gridSize(dimX/blockSize.x+1,dimY/blockSize.y, dimZ/blockSize.z);
float* dev_p;       cudaMalloc((void**)&dev_p, dimX*dimY*dimZ*sizeof(float));

dd<<<gridSize, blockSize>>>(     dev_p,dimX,dimY,dimZ);

谁能指点一下吗？因为这对我来说没有多大意义。 val 的所有计算都很快，最后一步是将 val 移动到 p。 p从不参与计算，只出现一次。那为什么这么慢？

计算基本上是 512 X 512 矩阵的循环。我会说这是相当多的计算量。

Answer 1

您在 SomeComputationOnVal 中执行的计算非常昂贵。每个线程至少读取 1MB 的数据，这些数据不在缓存中（或者最好在 L2 中读取一小部分，如果 k 在小范围内变化），总计运行大约 16 TB 的数据。即使在高端 gpu 上，至少也需要大约 2 分钟才能达到运行。更不用说所有可能减慢速度的事情了。

你的函数不会在全局内存中写入任何数据，也没有边界效应。如果您不使用输出，编译器可能会决定优化方法调用。

所以案例二和案例三不做计算的速度非常快。在 gpu 内存上写入 64 MB，并发线程非常快（毫秒范围）。

您可以验证生成的 ptx 以查看代码是否得到优化。在 nvcc 中使用 --keep 选项并搜索 ptx 文件。

在 CUDA 中访问全局内存很慢

Accessing global memory in CUDA is slow

memory

cuda

gpu

gpgpu