CUDA 指针运算导致未合并的内存访问？

Question

我正在使用必须对指针指针进行操作的 CUDA 内核。内核基本上执行大量非常小的缩减，最好连续进行，因为缩减的大小为 Nptrs=3-4。下面是内核的两个实现：

__global__
void kernel_RaiseIndexSLOW(double*__restrict__*__restrict__ A0,
        const double*__restrict__*__restrict__ B0,
        const double*__restrict__*__restrict__ C0,
        const int Nptrs, const int Nx){
      const int i = blockIdx.y;
      const int j = blockIdx.z;
      const int idx = blockIdx.x*blockDim.x + threadIdx.x;
      if(i<Nptrs) {
         if(j<Nptrs) {
           for (int x = idx; x < Nx; x += blockDim.x*gridDim.x){
              A0gpu[i+3*j][x] = B0gpu[i][x]*C0gpu[3*j][x]
                       +B0gpu[i+3][x]*C0gpu[1+3*j][x]
                       +B0gpu[i+6][x]*C0gpu[2+3*j][x];               
           }
         }
       }
 }

__global__
void kernel_RaiseIndexsepderef(double*__restrict__*__restrict__  A0gpu, 
               const double*__restrict__*__restrict__ B0gpu,
               const double*__restrict__*__restrict__ C0gpu,
               const int Nptrs, const int Nx){
const int i = blockIdx.y;
const int j = blockIdx.z;
const int idx = blockIdx.x*blockDim.x + threadIdx.x;
if(i<Nptrs) {
  if(j<Nptrs){
    double*__restrict__ A0ptr = A0gpu[i+3*j];
    const double*__restrict__ B0ptr0 = B0gpu[i];
    const double*__restrict__ C0ptr0 = C0gpu[3*j];
    const double*__restrict__ B0ptr1 = B0ptr0+3;
    const double*__restrict__ B0ptr2 = B0ptr0+6;
    const double*__restrict__ C0ptr1 = C0ptr0+1;
    const double*__restrict__ C0ptr2 = C0ptr0+2;

    for (int x = idx; x < Nx; x +=blockDim.x *gridDim.x){
      double d2 = C0ptr0[x];
      double d4 = C0ptr1[x]; //FLAGGED
      double d6 = C0ptr2[x]; //FLAGGED
      double d1 = B0ptr0[x];
      double d3 = B0ptr1[x]; //FLAGGED
      double d5 = B0ptr2[x]; //FLAGGED
      A0ptr[x] = d1*d2 + d3*d4 + d5*d6;

    }
   }                                                                        
  }
 }

如名称所示，内核 "sepderef" 的执行速度比其对应版本快约 40%，一旦计入启动开销，在 Nptrs=3、Nx=60000 时实现了大约 85GBps 的有效带宽开启 ECC 的 M2090（~160GBps 最佳）。

运行这些通过 nvvp 表明内核是带宽限制的。然而，奇怪的是，我标记为 //FLAGGED 的行被分析器突出显示为次优内存访问区域。我不明白为什么会这样，因为这里的访问对我来说看起来是合并的。为什么不是？

编辑：我忘了指出这一点，但请注意，//FLAGGED 区域正在访问我已对其进行算术运算的指针，而其他区域是使用方括号运算符访问的。

Answer 1

要理解这一行为，需要了解目前所有 CUDA GPU 都执行指令 in-order。从内存中加载操作数的指令发出后，其他独立的指令仍然继续执行。但是，一旦遇到依赖于内存中操作数的指令，所有对该指令流的进一步操作都会停止，直到操作数可用。

在您的 "sepderef" 示例中，您在对它们求和之前从内存中加载所有操作数，这意味着每次循环迭代可能只发生一次全局内存延迟（每次循环迭代有六次加载，但是它们都可以重叠。只有循环的第一个添加会停止，直到它的操作数可用。停止后，所有其他添加的操作数将很容易或很快可用。

在"SLOW"的例子中，从内存加载和加法是混合的，所以每次循环操作都会产生多次全局内存延迟。

您可能想知道为什么编译器不会在计算之前自动重新排序加载指令。 CUDA 编译器过去非常积极地执行此操作，在操作数等待使用的地方扩展额外的寄存器。然而，CUDA 8.0 在这方面似乎远没有那么积极，更多地坚持源代码中的指令顺序。这为程序员提供了更好的机会以性能最佳的方式构建代码 where the compiler's instruction scheduling was suboptimal。同时，即使以前的编译器版本正确，它也给程序员显式调度指令带来了更多负担。

CUDA 指针运算导致未合并的内存访问？

CUDA pointer arithmetic causes uncoalesced memory access?

cuda

nvvp