我想了解银行冲突如何发生.
如果我在全局内存中有一个大小为256的数组,并且我在一个Block中有256个线程,并且我想将该数组复制到共享内存.因此每个线程复制一个元素.
如果我在全局内存中有一个大小为256的数组,并且我在一个Block中有256个线程,并且我想将该数组复制到共享内存.因此每个线程复制一个元素.
shared_a[threadIdx.x]=global_a[threadIdx.x]
这个简单的行动是否导致银行冲突？
现在假设数组的大小大于线程数,所以我现在正在使用它来将全局内存复制到共享内存中：
tid = threadIdx.x;
for(int i=0;tid+i<N;i+=blockDim.x)
shared_a[tid+i]=global_a[tid+i];
上述代码是否导致银行冲突？
解决方法
检查这个的最好方法是使用“Compute Visual Profiler”来分析您的代码;这与CUDA工具包一起提供.此外,GPU Gems 3年还有一个很棒的部分 – “39.2.3避免银行冲突”.
“当同一个warp中的多个线程访问同一个bank时,会发生bank冲突,除非warp的所有线程在相同的32位字中访问相同的地址” – 首先有16个内存条,每个4字节宽.所以基本上,如果你有一个共享内存库中同一个4bytes的半经编读取内存的线程,你将有银行冲突和序列化等.
好,所以你的第一个例子：
首先让我们假设你的数组是例如int类型(一个32位字).您的代码将这些内存保存到共享内存中,跨越第K个线程保存到第K个内存条的任何半角.所以例如前半线的线程0将保存在第一个存储区中的shared_a [0],线程1将保存到shared_a [1],每个半线具有16个线程,这些映射到16个4字节的存储区.在下一半的warp中,第一个线程现在将其值保存在再次存放在第一个内存中的shared_a [16]中.所以如果你使用一个4字节这样的int,float等等,那么你的第一个例子就不会导致银行冲突.如果您使用1字节的字,如char,则前半部经线0,1,2和3将所有值都保存到共享存储器的第一个存储区,这将导致存储库冲突.
第二个例子：
再次,这一切都取决于你使用的单词的大小,但是对于这个例子,我将使用一个4字节的单词.所以看上半场经线：
线程数= 32
N = 64
线程0：将写入0,31,63
线程1：将写入1,32
半个warp的所有线程并发执行,因此对共享内存的写入不应该导致银行冲突.我必须仔细检查一下.
希望这有帮助,抱歉的巨大答复！
