如果n < warpSize，它是否合并？
// In kernel
int x;
if (threadId < n)
x = globalMem[threadId];
如果某些N 不能被warpSize 整除，则这种情况出现在循环的最后一次迭代中。我应该运行这些情况并分配只能被warpSize 整除的设备内存还是按原样合并？
【问题讨论】：
标签： c++ cuda memory-access coalescing
如果threadId 的计算正确，如cuda programming guide - thread hierachy 中所述，则此访问将被合并 - threadId = threadIdx.x 就是这种情况。
对于不同的计算架构，内存合并略有不同。更多详情请访问appendix G of cuda programming guide。
通常你可以说全局内存访问是合并的，如果你的线程从你第一个线程访问的元素的地址开始抓取内存中的连续元素。
假设您有一个浮点数组。 float array[]
你的内存访问看起来是这样的
array[threadIdx.x == 0, threadId.x == 1, threadIdx.x == 2, ..., threadIdx.x == 31]
您的访问权限将被合并。
但是如果你以这种方式访问内存（例如交错）
array[threadIdx.x == 0, NONE, threadId.x == 1, NONE, threadIdx.x == 2, ..., NONE, threadIdx.x == 31]
你的访问没有被合并（NONE 表示这个数组元素没有被任何线程访问）
在第一种情况下，您获取 128 个连续字节的内存。在第二种情况下，您获取 256 个字节。对于第二种情况，需要两个扭曲来从全局内存中加载内存，而不是第一种情况下的一个扭曲。但在这两种情况下，以下计算只需要 32 个浮点元素（即 128 个字节）。因此，在这种简单情况下，您的全局加载率将从 1.0 下降到 0.5。
【讨论】：