c++ - GPU 共享内存库冲突

Question

我想了解银行冲突是如何发生的。
如果我在全局内存中有一个大小为 256 的数组，并且我在单个 block 中有 256 个线程，并且我想将该数组复制到共享内存。因此每个线程都复制一个元素。

shared_a[threadIdx.x]=global_a[threadIdx.x]

这个简单的 Action 会导致银行冲突吗？

假设现在数组的大小大于线程数，所以我现在使用它来将全局内存复制到共享内存:

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 个内存组，每个 4 字节宽。所以本质上，如果你有 half warp 中的任何线程 从共享内存库中的相同 4 字节读取内存，你将遇到内存库冲突和序列化等问题。

好的，你的第一个例子:

首先假设您的数组是 int 类型的示例(32 位字)。您的代码将这些 int 保存到共享内存中，在第 K 个线程保存到第 K 个内存库的任何半个 warp 中。因此，例如，前半部分 warp 的线程 0 将保存到第一个内存库中的 shared_a[0]，线程 1 将保存到 shared_a[1]，每个half warp 有 16 个线程，这些线程映射到 16 个 4 字节组。在下半个 warp 中，第一个线程现在将其值再次保存到 first 内存库中的 shared_a[16] 中。因此，如果您使用 4 字节的字，如 int、float 等，那么您的第一个示例将不会导致银行冲突。如果使用 char 等单字节字，在前半部分 warp 线程 0、1、2 和 3 会将它们的值全部保存到共享内存的第一个 bank，这将导致 bank 冲突。

第二个例子:

同样，这完全取决于您使用的单词的大小，但对于示例，我将使用 4 字节的单词。所以看看前半部分扭曲:

线程数 = 32

N = 64

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

half warp 中的所有线程并发执行，因此对共享内存的写入不应导致存储体冲突。不过，我必须仔细检查一下。

希望这对您有所帮助，对于大量回复感到抱歉!