首页 文章

共享内存库与char数组冲突

提问于
浏览 次
2

我理解处理4字节数据类型时的银行冲突，但我想知道我们是否与以下代码发生任何银行冲突（4路/ 8路？）

__shared__ char shared[];
foo = shared[threadIdx.x];

上面的代码导致warp中的4个连续线程访问同一个bank中的相同字地址 .

类似的内存访问模式会导致任何cuda设备系列的银行冲突吗？显然，它只适用于旧卡，但我想确认一下 .

我的问题可以进一步推广，如果多个线程访问同一个bank可寻址单元[8字节或4字节]但每个都需要它的一小部分会怎样 . 硬件是否会处理此类请求而不存在任何银行冲突？谢谢

cuda gpgpu

1 回答

  • 1

    所有cc2.0和更新的GPU设备都具有广播机制，使得参与warp请求的任意数量的线程正在访问给定的32位对齐位置或该位置的任何部分（或多组此类线程，每组访问给定的32位对齐位置或该位置的任何部分），该组中的线程将在单个事务中进行服务，而不进行序列化 .

    来自the documentation：

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

    对于支持8字节组模式且处于8字节组模式的设备，上述广播机制推广到64位对齐位置 .

    请注意，我在这里仔细选择了我的措辞 . 假设我在一个warp请求中有两个这样的广播组 . 现在还假设这些组正在寻址两个不同的位置，但同一个银行中有两个位置 . 例如 . 组A的目标地址为0，组B的目标地址为1024.在这种情况下，参与组A的所有线程将在单个事务中进行服务，参与组B的所有线程将在单个事务中进行服务，但那些两组将相互序列化 .

    另一方面，如果组目标地址0和组B目标地址8，则它们位于不同的组中，因此组A中的所有线程都可以在单个事务中被服务，因为多个广播允许每个请求 .

    回复于

相关问题