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

Question

Я всегда работал с линейной разделяемой памятью (загрузка, хранение, доступ к соседям), но я сделал простой тест в 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);

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

Answer 1

Да, разделяемая память организована в порядке следования строк, как вы и ожидали.Таким образом, ваш [16] [16] массив хранится построчно, примерно так:

       bank0 .... bank15
row 0  [ 0   .... 15  ]
    1  [ 16  .... 31  ]
    2  [ 32  .... 47  ]
    3  [ 48  .... 63  ]
    4  [ 64  .... 79  ]
    5  [ 80  .... 95  ]
    6  [ 96  .... 111 ]
    7  [ 112 .... 127 ]
    8  [ 128 .... 143 ]
    9  [ 144 .... 159 ]
    10 [ 160 .... 175 ]
    11 [ 176 .... 191 ]
    12 [ 192 .... 207 ]
    13 [ 208 .... 223 ]
    14 [ 224 .... 239 ]
    15 [ 240 .... 255 ]
       col 0 .... col 15

Поскольку на оборудовании pre-Fermi есть 16 32-битных банков разделяемой памяти, каждая целочисленная запись в каждом столбце отображается наодин общий банк памяти.Так как это взаимодействует с выбранной вами схемой индексации?

Следует иметь в виду, что потоки в блоке нумеруются в эквиваленте основного порядка столбцов (технически измерение x структуры является самым быстрымизменяющийся, сопровождаемый y, сопровождаемый z).Поэтому, когда вы используете эту схему индексации:

shData[threadIdx.x][threadIdx.y]

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

shData[threadIdx.y][threadIdx.x]

потоки в одном и том же Half-Warp будут читать из одной и той же строки, что подразумевает чтение из каждого из 16 различных банков общей памяти, никаких конфликтов не возникает.