Как организована 2D-общая память в 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);

Спасибо заранее.

Ответ 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]

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