Арифметика указателя CUDA вызывает несвязанный доступ к памяти?

Я работаю с ядром 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% быстрее, чем его аналог, достигая, с учетом накладных расходов на запуск, эффективной пропускной способности около 85 Гбит/с при Nptrs=3, Nx=60000 на M2090 с включенным ECC (~160 Гбит/с). будет оптимальным).

Запуск их через nvvp показывает, что ядро ​​​​привязано к пропускной способности. Странно, однако, что строки, которые я пометил //FLAGGED, выделяются профилировщиком как области неоптимального доступа к памяти. Я не понимаю, почему это так, поскольку доступ здесь выглядит объединенным для меня. Почему бы и нет?

Редактировать: я забыл указать на это, но обратите внимание, что области //FLAGGED обращаются к указателям, над которыми я выполнял арифметические действия, тогда как к другим обращались с помощью оператора квадратных скобок.