我正在尝试了解银行冲突是如何发生的。
如果我在全局内存中有一个大小为 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];
上面的代码会导致银行冲突吗？
【问题讨论】：
标签： c++ cuda gpgpu bank-conflict
检查这一点的最佳方法是使用“Compute Visual Profiler”来分析您的代码；这与 CUDA 工具包一起提供。 GPU Gems 3 中还有一个很棒的部分 - “39.2.3 避免银行冲突”。
“当同一个 warp 中的多个线程访问同一个 bank 时，除非 warp 的所有线程访问同一个 32 位字内的同一个地址，否则会发生 bank 冲突” - 首先有16 个内存库，每个 4 字节宽。所以本质上，如果你有 任何线程在半扭曲中从共享内存银行中的相同 4 字节读取内存，你将有银行冲突和序列化等。
好的，所以你的第一个例子：
首先假设您的数组是 int 类型（32 位字）。您的代码将这些整数保存到共享内存中，跨过任何半扭曲，第 K 个线程正在保存到第 K 个内存库。因此，例如前半个 warp 的线程 0 将保存到第一个内存库中的 shared_a[0]，线程 1 将保存到 shared_a[1]，每个半个 warp 有 16 个线程，这些线程映射到 16 个 4 字节组。在接下来的半个 warp 中，第一个线程现在将其值保存到 shared_a[16] 中，它再次位于 first 内存库中。因此，如果您使用 4 字节字，例如 int、float 等，那么您的第一个示例不会导致银行冲突。如果使用 char 等 1 字节字，则在前半部分 warp 线程 0、1、2 和 3 都会将它们的值保存到共享内存的第一个 bank 中，这将导致 bank 冲突。
第二个例子：
同样，这完全取决于您使用的单词的大小，但对于示例，我将使用 4 字节的单词。所以看看前半段经线：
线程数 = 32
N = 64
线程 0：将写入 0、31、63 线程 1：将写入 1, 32
half warp 中的所有线程同时执行，因此写入共享内存不应导致银行冲突。不过，我必须仔细检查一下。
希望这会有所帮助，对于大量的回复感到抱歉！
【讨论】：
for (int i = tid ; i < N ; i += blockDim.x) shared_a[i] = global_a[i];
在这两种情况下，线程都使用连续地址访问共享内存。它取决于共享内存的元素大小，但线程束对共享内存的连续访问不会导致“小”元素大小的存储库冲突。
使用 NVIDIA Visual Profiler 分析 this code 表明，对于小于 32 和 4 的倍数（4、8、12、...、28）的元素大小，连续访问共享内存不会导致存储库冲突。但是，元素大小为 32 会导致存储库冲突。
Ljdawson 的回答包含一些过时的信息：
...如果您使用 1 字节字，例如 char，则在前半部分 warp 线程 0、1、2 和 3 都会将它们的值保存到共享内存的第一组，这将导致组冲突。
这对于旧 GPU 可能是正确的，但对于 cc >= 2.x 的最新 GPU，它们不会导致存储库冲突，这实际上是由于广播机制 (link)。以下引用来自CUDA C PROGRAMMING GUIDE (v8.0.61) G3.3. Shared Memory。
对 warp 的共享内存请求不会在访问同一 32 位字内的任何地址的两个线程之间产生存储体冲突（即使两个地址位于同一存储体中）：在这种情况下，对于读取访问，这个词被广播到请求线程（多个词可以在一个事务中广播）并且对于写访问，每个地址只由一个线程写入（哪个线程执行写入是未定义的）。
这意味着，特别是，如果按以下方式访问 char 数组，则不会发生存储库冲突，例如：
extern __shared__ char shared[]; char data = shared[BaseIndex + tid];
【讨论】：