12

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

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

4

2 回答 2

16

检查这一点的最佳方法是使用“Compute Visual Profiler”来分析您的代码；这与 CUDA 工具包一起提供。GPU Gems 3中也有一个很棒的部分- “39.2.3 Avoiding Bank Conflicts”。

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

好的，你的第一个例子：

首先让我们假设您的数组例如是int类型（一个 32 位字）。您的代码将这些整数保存到共享内存中，跨过任何半扭曲，第 K 个线程正在保存到第 K 个内存库。因此，例如前半束的线程 0 将保存到shared_a[0]第一个内存组中的哪个，线程 1 将保存到shared_a[1]，每个半束有 16 个线程，这些线程映射到 16 个 4 字节组。在接下来的半个 warp 中，第一个线程现在将其值保存到 shared_a[16] 中，这是第一个再次记忆库。因此，如果您使用 4 字节字，例如 int、float 等，那么您的第一个示例不会导致银行冲突。如果您使用 1 字节的字，例如 char，则在前半部分 warp 线程 0、1、2 和 3 都会将它们的值保存到共享内存的第一组，这将导致组冲突。

第二个例子：

同样，这完全取决于您使用的单词的大小，但对于示例，我将使用 4 字节的单词。所以看看前半个经线：

线程数 = 32

N = 64

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

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

希望这会有所帮助，很抱歉回复了这么多！

于 2010-12-09T09:46:04.367 回答
3

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

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

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

...如果您使用 1 字节字，例如 char，则在前半部分 warp 线程 0、1、2 和 3 都会将它们的值保存到共享内存的第一组，这将导致组冲突。

这对于旧 GPU 可能是正确的，但对于 cc >= 2.x 的最新 GPU，它们不会导致存储库冲突，这实际上是由于广播机制（链接）。以下引用来自CUDA C 编程指南 (v8.0.61) G3.3。共享内存。

对 warp 的共享内存请求不会在访问同一 32 位字中的任何地址的两个线程之间产生存储区冲突（即使两个地址位于同一存储区中）：在这种情况下，对于读取访问，字被广播到请求线程（可以在单个事务中广播多个字）并且对于写访问，每个地址仅由一个线程写入（哪个线程执行写入未定义）。

这尤其意味着，如果按如下方式访问 char 数组，则不会发生存储库冲突，例如：

   extern __shared__ char shared[];
   char data = shared[BaseIndex + tid];
于 2017-04-18T10:21:31.037 回答