我对来自 NVidia 的 2 个文件感到困惑。 “CUDA 最佳实践”描述了共享内存是按组组织的,通常在 32 位模式下,每 4 个字节是一个组(这就是我的理解)。然而 Parallel Prefix Sum (Scan) with CUDA详细介绍了由于组冲突而应如何将填充添加到扫描算法中。
对我来说问题是,这个算法的基本类型是浮点数,大小是 4 个字节。因此,每个 float 都是一家银行,不存在银行冲突。
那么我的理解是否正确——即,如果您使用 4*N 字节类型,则不必担心银行冲突,因为根据定义,不会有任何冲突?如果不是,我应该如何理解(何时使用填充)?
最佳答案
您可能对 this webinar 感兴趣来自 NVIDIA CUDA webinar page来自 this webinar 的幻灯片 35-45 中也描述了包括存储体在内的共享内存。 .
通常,当两个不同的线程尝试访问(从同一内核指令)共享内存中的低 4 位(cc2.0 之前的设备)或 5 位(cc2.0 和更新的设备)的地址相同。当确实发生 bank 冲突时,共享内存系统会序列化对同一 bank 中的位置的访问,从而降低性能。对于某些访问模式,填充尝试避免这种情况。请注意,对于 cc2.0 和更新版本,如果所有位都相同(即相同的位置),这不会导致库冲突。
从图片上看,我们可以这样看:
__shared__ int A[2048];
int my;
my = A[0]; // A[0] is in bank 0
my = A[1]; // A[1] is in bank 1
my = A[2]; // A[2] is in bank 2
...
my = A[31]; // A[31] is in bank 31 (cc2.0 or newer device)
my = A[32]; // A[32] is in bank 0
my = A[33]; // A[33] is in bank 1
现在,如果我们在 warp 中跨线程访问共享内存,我们可能会遇到 bank 冲突:
my = A[threadIdx.x]; // no bank conflicts or serialization - handled in one trans.
my = A[threadIdx.x*2]; // 2-way bank conflicts - will cause 2 level serialization
my = A[threadIdx.x*32]; // 32-way bank conflicts - will cause 32 level serialization
让我们仔细看看上面的 2-way bank 冲突。由于我们正在乘以
threadIdx.x
通过 2,线程 0 访问存储区 0 中的位置 0,但线程 16 访问也在存储区 0 中的位置 32,从而产生存储区冲突。对于上面的 32 路示例,所有地址都对应于 bank 0。因此,必须发生 32 个到共享内存的事务才能满足此请求,因为它们都是序列化的。所以回答这个问题,如果我知道我的访问模式会是这样的,例如:
my = A[threadIdx.x*32];
然后我可能想要填充我的数据存储,以便
A[32]
是一个虚拟/焊盘位置,正如 A[64]
, A[96]
等等。然后我可以像这样获取相同的数据:
my = A[threadIdx.x*33];
并在没有银行冲突的情况下获取我的数据。
希望这可以帮助。
关于cuda - 什么时候真正需要填充共享内存?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/15056842/