Как устроена двумерная общая память в CUDA

Я всегда работал с линейной разделяемой памятью (загрузка, хранение, доступ к соседям), но я сделал простой тест в 2D для изучения конфликтов банков, результаты которого меня запутали.

Следующий код считывает данные из одномерного массива глобальной памяти в общую память и копирует их обратно из общей памяти в глобальную память.

__global__ void update(int* gIn, int* gOut, int w) {

// shared memory space
__shared__ int shData[16][16];
// map from threadIdx/BlockIdx to data position
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
// calculate the global id into the one dimensional array
int gid = x + y * w;

// load shared memory
shData[threadIdx.x][threadIdx.y] = gIn[gid];
// synchronize threads not really needed but keep it for convenience
__syncthreads();
// write data back to global memory
gOut[gid] = shData[threadIdx.x][threadIdx.y];
}

Визуальный профилировщик сообщил о конфликтах в разделяемой памяти . Следующий код избегает этих конфликтов (только показывает различия)

// load shared memory
shData[threadIdx.y][threadIdx.x] = gIn[gid];

// write data back to global memory
gOut[gid] = shData[threadIdx.y][threadIdx.x];

Такое поведение меня смутило, потому что в программировании массивно параллельных процессоров. Практический подход мы можем прочитать:

матричные элементы в C и CUDA помещаются в ячейки с линейной адресацией в соответствии с соглашением о старших строках. То есть элементы строки 0 матрицы сначала размещаются по порядку в последовательных местах.

Связано ли это с массивом разделяемой памяти? или с указателями потоков? Может быть, я что-то упустил?

Конфигурация ядра следующая:

// kernel configuration
dim3 dimBlock  = dim3 ( 16, 16, 1 );
dim3 dimGrid   = dim3 ( 64, 64 );
// Launching a grid of 64x64 blocks with 16x16 threads -> 1048576 threads
update<<<dimGrid, dimBlock>>>(d_input, d_output, 1024);

Заранее спасибо.