Блог

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

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

Да, разделяемая память упорядочена в порядке следования строк, как вы и ожидали. Итак, ваш массив [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 различных банков общей памяти, никаких конфликтов не возникает.