在CUDA设备上故意引起共享内存银行冲突

3

对于我来说,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 GPU上运行了它。由于多个线程都可以访问同一个共享内存变量,我预计这个变量会被更新5*nthr次,尽管由于银行冲突的原因,不会在同一周期内进行更新。然而,输出表明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周期内这是不可能的。串行化对我来说没问题,因为这将是一个罕见的事件。

1个回答

2
让我们一起浏览它生成的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;

基本上,这个标签是HTML代码。
for (int i=0;i<5;i++)
    mywarp += (10000*threadIdx.x+1);

“正在优化到”
mywarp=50000*threadIdx.x+5

所以你不是遇到了银行冲突,而是遇到了竞态条件。


你说得对,我昨天想得有些偏了。谢谢你提醒我。 - Christian Sarofeen
感谢您的分析。有些事情对我仍然不太清楚:i)当我添加volatile属性(正如您在第一个答案中建议的那样)时,它会稍微改变输出,例如最后一位数字有时不是5,而是7、8。ii)允许只有一个线程修改mywarp变量是否可以解决竞争条件问题?iii)如果我希望所有线程都参与,我需要使用atomicAdd()吗? - yarchik
这取决于你需要什么。如果你只需要简单的规约(跨线程求和),你可以将数据保留在寄存器中,然后使用共享内存中的规约方法。如果你需要每个warp唯一地更新一个值,那么是的,你需要使用atomicAdd。 - Christian Sarofeen

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接