Как организована 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 разных банков разделяемой памяти, никаких конфликтов не происходит.
