0

我正在使用必须对指针进行操作的 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 的 M2090 上实现约 85GBps 有效带宽，ECC 开启（~160GBps将是最佳的）。

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

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

4

1 回答 1

2

要理解这种行为，我们需要知道到目前为止所有 CUDA GPU 都按顺序执行指令。在发出从内存中加载操作数的指令后，其他独立指令仍会继续执行。但是，一旦遇到依赖于内存中操作数的指令，对该指令流的所有进一步操作都会停止，直到操作数变得可用。

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

在“SLOW”示例中，从内存加载和加法混合在一起，因此每个循环操作都会产生多次全局内存延迟。

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

于 2017-06-11T07:39:23.510 回答