Практический пример использования разделяемой памяти GPU

Question

У меня есть такой массив:

data[16] = {10,1,8,-1,0,-2,3,5,-2,-3,2,7,0,11,0,2}

Я хочу вычислить сокращение этого массива, используя разделяемую память на GPU G80.

Ядро, упомянутое в документе NVIDIA, выглядит так:

__global__ void reduce1(int *g_idata, int *g_odata) {
extern __shared__ int sdata[];

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
sdata[tid] = g_idata[i];
__syncthreads();

// here the reduction :

for (unsigned int s=1; s < blockDim.x; s *= 2) {
int index = 2 * s * tid;
if (index < blockDim.x) {
sdata[index] += sdata[index + s];
}
__syncthreads();
}

Автор статьи сказал, что в этом методе существует проблема банковского конфликта. Я пытался понять, но я не мог понять, почему? Я знаю определение банковского конфликта и широковещательного доступа, но все еще не могу этого понять.

Банковские конфликты

Answer 1

Процессор G80 - это очень старый графический процессор с поддержкой CUDA, в первом поколении графических процессоров CUDA, с вычислительной способностью 1,0. Эти устройства больше не поддерживаются последними версиями CUDA (после 6.5), поэтому онлайн-документация больше не содержит необходимой информации для понимания структуры банка в этих устройствах.

Поэтому я извлеку необходимую информацию для устройств cc 1.x из руководства по программированию CUDA 6.5 C здесь:

G.3.3. Общая память

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

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

В этих устройствах разделяемая память имеет структуру из 16 банков, так что каждый банк имеет "ширину" 32 бита или 4 байта. Каждый банк имеет такую ​​же ширину, как int или же float количество, например. Поэтому давайте представим первые 32 4-байтовых количества, которые могут быть сохранены в этом виде совместно используемой памяти, и их соответствующие банки (используя f вместо sdata для имени массива):

extern __shared__ int f[];

index: f[0] f[1] f[2] f[3] ... f[15] f[16] f[17] f[18] f[19] ... f[31]
bank:    0    1    2    3  ...   15     0     1     2     3  ...   15

Первые 16 int количества в разделяемой памяти принадлежат банкам от 0 до 15, а следующие 16 int количества в разделяемой памяти также принадлежат банкам от 0 до 15 (и так далее, если у нас было больше данных в нашем int массив).

Теперь давайте посмотрим на строки кода, которые вызовут конфликт в банке:

for (unsigned int s=1; s < blockDim.x; s *= 2) {
int index = 2 * s * tid;
if (index < blockDim.x) {
sdata[index] += sdata[index + s];
}

Давайте рассмотрим первый проход через вышеуказанный цикл, где s это 1. Это значит index является 2*1*tid так что для каждой темы, index просто удвоить значение threadIdx.x:

threadIdx.x: 0 1 2 3 4  5  6  7  8  9 10 11 ...
 index:      0 2 4 6 8 10 12 14 16 18 20 22 ...
 bank:       0 2 4 6 8 10 12 14  0  2  4  6 ...

Итак, для этой операции чтения:

+= sdata[index + s]

у нас есть:

threadIdx.x: 0 1 2 3 4  5  6  7  8  9 10 11 ...
 index:      0 2 4 6 8 10 12 14 16 18 20 22 ...
 index + s:  1 3 5 7 9 11 13 15 17 19 21 23 ...
 bank:       1 3 5 7 9 11 13 15  1  3  5  7 ...

Итак, в первых 16 потоках у нас есть два потока, которые хотят читать из банка 1, два, которые хотят читать из банка 3, два, которые хотят читать из банка 5, и т. Д. Таким образом, этот цикл чтения сталкивается с двусторонними конфликтами банков через первую группу с 16 потоками. Обратите внимание, что другие операции чтения и записи в той же строке кода аналогично конфликтуют между банками:

sdata[index] +=

как это будет читать, а затем записывать в банки 0, 2, 4 и т. д. дважды на группу из 16 потоков.

Примечание для тех, кто может читать этот пример: как написано, он относится только к устройствам cc 1.x. Методология демонстрации конфликтов банков на cc 2.x и более новых устройствах может быть схожей, но специфика различается из-за различий в выполнении деформации и того факта, что эти более новые устройства имеют структуру банка с 32 путями, а не банк с 16 путями состав.