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%，一旦计算了启动开销，在打开ECC的M2090上，在Nptrs=3，Nx=60000时实现了约85GBps的有效带宽(~160GBps将是最佳的)。

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

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

Answer 1

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

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

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

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