намеренно вызывает конфликты банков для общей памяти на устройстве CUDA

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

#include <cuda_runtime.h>
#include <stdio.h>

#define nblc 13
#define nthr 1024

//------------------------@device--------------------

__device__ int inwarpD[nblc];

__global__ void kernel(){
__shared__ int mywarp;

mywarp=0;
for (int i=0;i<5;i++) mywarp += (10000*threadIdx.x+1);
__syncthreads();

inwarpD[blockIdx.x]=mywarp;
}
//------------------------@host-----------------------

int main(int argc, char **argv){
int inwarpH[nblc];
cudaSetDevice(2);

kernel<<<nblc, nthr>>>();

cudaMemcpyFromSymbol(inwarpH, inwarpD, nblc*sizeof(int), 0, cudaMemcpyDeviceToHost);

for (int i=0;i<nblc;i++) printf("%i : %i\n",i, inwarpH[i]);
}

и запустил его на графическом процессоре K80. Поскольку несколько потоков имеют доступ к одной и той же переменной общей памяти, я ожидал, что эта переменная будет обновляться 5 * nth раз, хотя и не в одном и том же цикле из-за конфликта банков. Однако выходные данные показывают, что общая переменная mywarp была обновлена ​​только 5 раз. Для каждого блока разные потоки выполняли эту задачу:

0 : 35150005
1 : 38350005
2 : 44750005
3 : 38350005
4 : 51150005
5 : 38350005
6 : 38350005
7 : 38350005
8 : 51150005
9 : 44750005
10 : 51150005
11 : 38350005
12 : 38350005

Вместо этого я ожидал

 523776*10000+5*1024=5237765120

для каждого блока. Может кто-нибудь любезно объяснить мне, где мое понимание общей памяти терпит неудачу. Я также хотел бы знать, как возможно, чтобы все потоки в одном блоке обращались (обновляли) к одной и той же общей переменной. Я знаю, что это невозможно в том же цикле MP. Сериализация мне подходит, потому что это будет редкое событие.


shared-memory gpu cuda bank-conflict
person yarchik    schedule 29.05.2015    source источник

Ответы (1)


arrow_upward
1
arrow_downward

Давайте пройдемся по ptx, который он генерирует.

//Declare some registers
.reg .s32       %r<5>;
.reg .s64       %rd<4>;

// demoted variable
.shared .align 4 .u32 _Z6kernelv$__cuda_local_var_35411_30_non_const_mywarp;

//load tid in register r1
mov.u32         %r1, %tid.x;

//multiple tid*5000+5 and store in r2
mad.lo.s32      %r2, %r1, 50000, 5;

//store result in shared memory
st.shared.u32   [_Z6kernelv$__cuda_local_var_35411_30_non_const_mywarp], %r2;

///synchronize
bar.sync        0;

//load from shared memory and store in r3
ld.shared.u32   %r3, [_Z6kernelv$__cuda_local_var_35411_30_non_const_mywarp];

mov.u32         %r4, %ctaid.x;
mul.wide.u32    %rd1, %r4, 4;
mov.u64         %rd2, inwarpD;
add.s64         %rd3, %rd2, %rd1;

//store r3 in global memory
st.global.u32   [%rd3], %r3;
ret;

Так что в основном

for (int i=0;i<5;i++)
    mywarp += (10000*threadIdx.x+1);

оптимизируется до

mywarp=50000*threadIdx.x+5

так что вы не испытываете банковский конфликт. Вы испытываете состояние гонки.

person Christian Sarofeen    schedule 29.05.2015
comment
Ты прав, я не знаю, о чем я думал вчера. Спасибо за внимание. - person Christian Sarofeen; 30.05.2015
comment
Спасибо за ваш анализ. Некоторые вещи мне до сих пор не ясны: i) когда я добавляю атрибут volatile (как вы предложили в своем первом ответе), он немного меняет вывод, например. последняя цифра не 5, а иногда 7,8. ii) Верно ли, что состояние гонки разрешается, позволяя только одному потоку изменять переменную mywarp? iii) Если я хочу, чтобы все потоки участвовали, мне нужен atomicAdd()? - person yarchik; 01.06.2015
comment
Это зависит именно от того, что вам нужно. Если вам нужна простая редукция (суммирование по потокам), вы можете оставить данные в регистре, а затем использовать метод редукции в разделяемой памяти. Если вам нужно, чтобы каждая деформация уникальным образом обновляла одно значение, тогда да, вам нужно использовать atomicAdd. - person Christian Sarofeen; 01.06.2015