CUDA：对全局内存的连续字节的半扭曲访问是否保证被合并？答案

【问题标题】：CUDA: are half-warp accesses to consecutive bytes of the global memory guaranteed to be coalesced?CUDA：对全局内存的连续字节的半扭曲访问是否保证被合并？
【发布时间】：2013-01-31 20:11:28
【问题描述】：

我有一个计算能力为 1.3 的 GPU。根据文档，当相同半扭曲的线程根据字长从相同的 32-64 或 128 字节内存段访问字节时，这些内存访问会合并为一个。

但是，在使用 cudaMallocPitch() 分配的二维数组的情况下，当同一个 half-warp 的线程访问连续字节时，是否保证这些字节驻留在同一个内存段？

CUDA coalesced access to global memory 有一个类似的问题，但不包括计算能力 1.3 GPU 和 2D 数组。

【问题讨论】：

标签： c optimization cuda

【解决方案1】：

是的 - cudaMallocPitch() 主要用于确保合并行为从一行持续到下一行。合并的标准是 per-warp，因此它们的粒度要细得多，并且与二维数组的一行中的连续字节有关。

请注意，操作数大小必须至少为 32 位，否则合并根本不起作用。

【讨论】：

"请注意，操作数大小必须至少为 32 位，否则合并根本不起作用。"这句话适用于计算能力 1.0 和 1.1，但不适用于较新的 GPU。
Tera，我想这取决于您如何定义“合并”。在每个级别的 CUDA 硬件上，8 位和 16 位访问都慢得多。在我的 GK104 上，8 位访问的差异为 2.4 倍，16 位访问的差异为 40%（100GB/s 与 64 位访问的 142GB/s）。 github.com/ArchaeaSoftware/cudahandbook/blob/master/memory/…
当然 8 位和 16 位访问要慢得多，因为在 CC2.0+ 上它们浪费了部分缓存行（缓存太小，无法保证 100% 重用）而在 CC 上1.2/1.3 它们没有有效地使用内存突发（并且在 8 位访问的情况下会浪费一半的内存带宽，因为要传输的数据量小于最小事务宽度）。尽管如此，half-warp 的访问仍然会合并为一个事务（前提是满足地址要求）。