我了解处理 4 字节数据类型时的银行冲突，但我想知道 如果我们使用以下代码遇到任何银行冲突（4 路/8 路？）
__shared__ char shared[];
foo = shared[threadIdx.x];
上面的代码导致一个warp中的4个连续线程访问同一个bank中的同一个字地址。
类似的内存访问模式是否会导致任何 cuda 设备系列的存储库冲突？显然，它只适用于旧卡，但我想确认一下。
我的问题可以进一步概括，如果多个线程访问同一个银行可寻址单元 [8 字节或 4 字节]，但每个线程都需要它的一小部分。硬件会在没有任何银行冲突的情况下处理此类请求吗？谢谢
【问题讨论】：
所有 cc2.0 和更新的 GPU 设备都具有广播机制，这样对于参与扭曲请求的任意数量的线程正在访问给定的 32 位对齐位置或该位置的任何部分（或多个此类线程组，每个组访问给定的 32 位对齐位置或该位置的任何部分），该组中的线程将在单个事务中提供服务，无需序列化。
对 warp 的共享内存请求不会在访问同一 32 位字内的任何地址的两个线程之间产生存储体冲突（即使两个地址位于同一存储体中）：在这种情况下，对于读取访问，这个词被广播到请求线程（多个词可以在一个事务中广播）并且对于写访问，每个地址只由一个线程写入（哪个线程执行写入是未定义的）。
对于支持 8-byte-bank-mode 且处于 8-byte-bank-mode 的设备，上述广播机制概括为 64 位对齐的位置。
请注意，我在这里仔细选择了措辞。假设我在一个 warp 请求中有两个这样的广播组。现在还假设这些组正在寻址两个不同的位置，但同一银行中的两个位置。例如。组 A 以地址 0 为目标，组 B 以地址 1024 为目标。在这种情况下，参与组 A 的所有线程将在单个事务中服务，而参与组 B 的所有线程将在单个事务中服务，但那些两个组将相互序列化。
另一方面，如果组目标地址 0 和组 B 目标地址 8，它们位于不同的组中，因此组 A 中的所有线程和组 B 中的所有线程都可以在单个事务中提供服务，因为允许每个请求进行多个广播。
【讨论】：