故意导致 CUDA 设备上共享内存的银行冲突
Posted
技术标签:
【中文标题】故意导致 CUDA 设备上共享内存的银行冲突【英文标题】:purposely causing bank conflicts for shared memory on CUDA device 【发布时间】:2015-08-12 15:47:48 【问题描述】: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*nhr 次,尽管由于银行冲突而不是在同一个周期。但是,输出表明 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】:让我们来看看它生成的 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
因此您不会遇到银行冲突。您遇到了竞争状况。
【讨论】:
你说得对,我不知道我昨天在想什么。感谢您的提醒。 感谢您的分析。有些事情我仍然不清楚:i)当我放置 volatile 属性(如您在第一个答案中建议的那样)时,它会稍微改变输出,例如最后一位数字不是 5,而是 7,8 有时。 ii) 是否通过只允许一个线程修改 mywarp 变量来解决竞争条件? iii) 如果我希望所有线程都参与,我需要 atomicAdd()? 这完全取决于您需要什么。如果您需要简单的归约(跨线程求和),您可以将数据留在寄存器中,然后在共享内存中使用归约方法。如果您需要每个 warp 唯一地更新单个值,那么是的,您需要使用 atomicAdd。以上是关于故意导致 CUDA 设备上共享内存的银行冲突的主要内容,如果未能解决你的问题,请参考以下文章