对于我来说,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周期内这是不可能的。串行化对我来说没问题,因为这将是一个罕见的事件。