在 CUDA 中访问全局内存很慢答案

【问题标题】：Accessing global memory in CUDA is slow在 CUDA 中访问全局内存很慢
【发布时间】：2016-09-21 21:42:02
【问题描述】：

我有一个 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 介入，并没有发现访问冲突。

这是我启动内核的方式（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 矩阵的循环。我会说这是相当大量的计算。

【问题讨论】：

你确定是全局内存访问需要“这么长时间”（你没有指定时间）吗？通过在全局数组中注释存储，在将 val 设置为 0 并调用您的“计算”后，永远不会使用它。编译器可能会认为做这项工作是无用的，并且可能会在编译您的源代码时忽略这一点。
@Taro 真的吗？那么我怎么知道 val 是否真的被计算出是否调用了更少的全局内存写入？对不起，我没有计时，但这么长至少意味着几十分钟。我什至没有完成它。
是的。这取决于很多因素，但对于大多数编译器来说，这是一个“经典”优化。在此处查看我的答案，了解如何查看生成的 PTX/SASS 代码并将其与源代码进行比较：stackoverflow.com/a/36477199/6172231
这取决于您在SomeComputationOnVal 中所做的工作量。在你的最后一个例子中，由于 val 甚至没有被使用并且 SomeComputationOnVal 没有参数（可能没有边界效应），它被优化了。所以你的最后一个样本没有测量SomeComputationOnVal。顺便说一句，你能提供那个函数的内容吗？
@FlorentDUGUET 当然。查看更新。

标签： memory cuda gpu gpgpu

【解决方案1】：

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

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

因此情况二和三不进行计算非常快。使用 coesced 线程在 gpu 内存上写入 64 MB 非常快（毫秒范围）。

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

【讨论】：