故意造成CUDA设备上共享内存的银行冲突

Question

对我来说，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 : %in",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周期是不可能的。序列化对我来说很好，因为它将是一个罕见的事件。

Answer 1

让我们看一下它生成的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

所以你没有遇到银行冲突。