Практический пример общей памяти графического процессора

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

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

Я хочу вычислить сокращение этого массива, используя общую память на графическом процессоре 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();
}

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

Конфликты банков


gpu-shared-memory cuda bank-conflict
person sara idrissi    schedule 03.04.2017    source источник
comment
Не было бы никаких конфликтов банков для размера данных 16 на G80, если предположить, что ваш blockDim.x также равен 16. Я совершенно уверен, что автор статьи не имел в виду ваш пример. При размере данных не менее 32 и blockDim.x не менее 32 несложно продемонстрировать, как возникает конфликт банков на G80.   -  person Robert Crovella    schedule 04.04.2017
comment
stackoverflow.com/q/7903566/681865   -  person talonmies    schedule 04.04.2017
comment
Я использовал тот же пример, что и в этой статье бумага(Я использовал тот же пример, что и в этой статье) Я говорю о методе на странице 11(вы можете видеть это на картинке, Я только что добавил в свой вопрос. Пожалуйста, не могли бы вы продемонстрировать, как возникает конфликт банков с 32 элементами? Большое спасибо @Robert Crovellla   -  person sara idrissi    schedule 04.04.2017

Ответы (1)


arrow_upward
2
arrow_downward

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

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

Г.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. структура.

person Robert Crovella    schedule 04.04.2017
comment
Это действительно очень четкое объяснение проблемы, на понимание которой я потратил много времени. Большое спасибо, дорогой мистер Робер @Robert Crovella - person sara idrissi; 05.04.2017