c++ - CUDA:二维网格中的线程 ID 分配

Question

假设我有一个带有 2D 网格的内核调用，如下所示:

dim3 dimGrid(x, y); // not important what the actual values are
dim3 dimBlock(blockSize, blockSize);
myKernel <<< dimGrid, dimBlock >>>();

现在我读到多维网格只是为了简化编程——底层硬件只会使用一维线性缓存内存(除非你使用纹理内存，但这与这里无关)。

我的问题是:在 warp 调度期间，线程将以什么顺序分配给网格索引？它们将水平分配(“迭代”x，然后 y)还是垂直分配(“迭代”y，然后 x)？这可能与改进内存合并有关，具体取决于我在内核中访问内存的方式。

为了说得更清楚，假设以下内容表示应用于我的(假想的)网格的线程 ID，具有“水平”分布:

[ 0  1  2  3 ]
[ 4  5  6  7 ]
[ 8  9 10 11 ]
[ ...        ]

“垂直”分布是:

[ 0  4  8 .. ]
[ 1  5  9 .. ]
[ 2  6 10 .. ]
[ 3  7 11 .. ]

我希望您能看到这对合并有何影响:对于每个变体，都会有一种特定的最佳方式来访问我的设备内存缓冲区。

不幸的是，我还没有找到任何关于这方面的详细信息..

Answer 1

水平和垂直是任意的。但是线程确实具有明确定义的 x、y 和 z 维度。线程按 x、y、z 的顺序分组到 warp 中。因此，一个 16x16 线程 block 在第一个 32 线程 warp 中将具有以下顺序的线程:

warp lane:线程 ID (x,y,z)

• 0: 0,0,0
• 1: 1,0,0
• 2: 2,0,0
• 3: 3,0,0
• ...
• 15: 15,0,0
• 16: 0,1,0
• 17: 1,1,0
• 18: 2,1,0
• 19:3,1,0
• ...
• 31: 15,1,0

以上也是您将拥有的 16,2 线程 block 尺寸的确切图案，它将填充一个经线。

对于没有完全填充 warp 的维度，例如 8,3(或 8,29，这将使最终的 warp 仅部分充满事件线程)，分配顺序遵循 the same pattern .正如您在上面看到的，当您将线程分配给扭曲时，快速变化的维度是 x 维度。下一个变化最快的维度是 y，然后是 z。

programming guide还解释了如何按顺序对线程进行编号。当线程按此顺序编号时，前 32 个线程属于第一个 warp，接下来的 32 个线程属于下一个 warp，依此类推。

如果在上述编号顺序中只有部分补充(即少于 32 个)线程可用于特定的最后一个 warp，那么最后一个 warp(仅)将由少于 32 个事件线程组成。