GPU共享内存库冲突

Question

我想了解银行冲突是如何发生的.
如果我在全局内存中有一个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];

以上代码会导致银行冲突吗？

Answer 1

检查这个的最好方法是使用"Compute Visual Profiler"来分析您的代码; 这附带了CUDA工具包.此外,GPU Gems 3中有一个很棒的部分- "39.2.3避免银行冲突".

" 当同一个warp中的多个线程访问同一个bank时,除非warp的所有线程在同一个32位字内访问相同的地址,否则会发生bank冲突 " - 首先有16个内存bank,每个4bytes宽.基本上,如果你在共享内存库中的相同4字节的半经线读取内存中有任何线程,你将会发生银行冲突和序列化等.

好的第一个例子:

首先假设您的数组是例如int类型(32位字).您的代码将这些内容保存到共享内存中,跨越Kth线程保存到第K个内存库的任何一半扭曲.因此,例如前半个warp的线程0将保存到shared_a[0]第一个内存库中,线程1将保存到shared_a[1],每个half warp有16个线程这些映射到16个4byte bank.在下一半warp中,第一个线程现在将其值保存到shared_a [16]中,这是第一个记忆库再次.因此,如果您使用4byte字,如int,float等,那么您的第一个示例将不会导致银行冲突.如果你使用一个1字节的字,如char,在前半部分经线0,1,2和3都将它们的值保存到第一个共享存储器组,这将导致存储体冲突.

第二个例子:

同样,这将取决于您正在使用的单词的大小,但对于示例,我将使用4字节的单词.所以看着上半场经线:

线程数= 32

N = 64

线程0:将写入0,31,63线程1:将写入1,32

半warp中的所有线程同时执行,因此对共享内存的写入不应导致存储体冲突.我不得不仔细检查这一个.

希望这有帮助,对不起的巨大回复!