GPU共享内存库冲突

Question

我试图了解银行冲突是如何发生的。
如果我在全局内存中有一个大小256的数组，并且在一个块中有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

检查此问题的最佳方法是使用“计算视觉分析器”介绍您的代码；这附带了CUDA工具包。还有一个很棒的部分 GPU宝石3 因此 - “ 39.2.3避免银行冲突”。

"当同一扭曲中的多个线程访问同一银行时，除非经纱的所有线程在同一32位单词中访问相同地址，否则会发生银行冲突“ - 第一件事，每个4个宽度都有16个记忆库。因此，如果您有 半扭曲中的任何线程 从共享内存库中的相同4个比较读取内存，您将有银行冲突和序列化等。

好的，你的第一个示例:

首先，假设您的数组是说例如类型 int (一个32位单词）。您的代码将这些INT保存到共享内存中，在任何一半的扭曲中，KTH线程都保存到KTH内存库。因此，例如，上半年扭曲的线程0将保存到 shared_a[0] 在第一个内存库中，线程1将保存到 shared_a[1], ，每半扭曲都有16个线程，这些地图是16个4 Byte Banks。在接下来的扭曲中，第一个线程现在将其值保存到共享_a [16]中 第一的 再次记忆库。因此，如果您使用这样的int，float等的4个单词，那么您的第一个示例将不会导致银行冲突。如果您使用1个字节词（例如char），则在上半年扭曲线程0、1、2和3中，所有的值都将其值将其值保存到共享内存的第一岸，这将导致银行冲突。

第二个示例:

同样，这完全取决于您使用的单词的大小，但是在示例中，我将使用一个4 Byte Word。因此，看上半场扭曲：

线程数= 32

n = 64

线程0：将写入0、31、63线程1：将写入1、32

一半扭曲的所有线程同时执行，因此写入共享内存的写入不应引起银行冲突。不过，我必须仔细检查一下。

希望这会有所帮助，对不起，回答很大！

Answer 2

在这两种情况下，线程都有连续地址访问共享内存。这取决于共享内存的元素大小，但是通过线程扭曲对共享内存的连续访问不会导致“小”元素大小的银行冲突。

分析 此代码 使用Nvidia Visual Profiler显示，对于小于32的元素大小和4（4、8、12，...，28）的倍数，对共享内存的连续访问不会导致银行冲突。但是，元素大小为32，导致银行冲突。

---

Ljdawson的答案包含一些过时的信息：

...如果您使用1个字节词，例如char，在上半年扭曲线程0、1、2和3中，所有的值都将其值保存到共享内存的第一岸，这将导致银行冲突。

对于旧的GPU来说，这可能是正确的，但是对于CC> = 2.x最近的GPU，它们不会导致银行冲突，这是由于广播机制的有效性（关联）。下引用来自 CUDA C编程指南（V8.0.61）G3.3。共享内存.

扭曲的共享内存请求不会在两个线程之间产生银行冲突，这些线程访问同一32位单词中的任何地址（即使两个地址落在同一银行中）：在这种情况下，对于阅读访问，被广播到请求线程（可以在单个事务中广播多个单词），对于写入访问，每个地址仅由一个线程之一（线程执行的写入不确定）编写。

特别是，这意味着，如果访问以下方式，则没有银行冲突，例如：

   extern __shared__ char shared[];
   char data = shared[BaseIndex + tid];