cuda - CUDAポインタ演算により、結合されていないメモリアクセスが発生しますか?

Question

ポインターからポインターへの操作が必要な CUDA カーネルを使用しています。カーネルは基本的に、多数の非常に小さいリダクションを実行します。リダクションのサイズは Nptrs=3 ～ 4 であるため、シリアルで実行するのが最適です。カーネルの 2 つの実装を次に示します。

__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 領域は、算術演算を行ったポインターにアクセスしているのに対し、他の領域は角かっこ演算子を使用してアクセスしていることに注意してください。