CUDA:在2D网格中分配线程ID

CUDA: Thread ID assignment in 2D grid

本文关键字：分配 线程 ID 网格 2D CUDA      更新时间：2023-10-16

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

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

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

我的问题是:在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 .. ]

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

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

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

warp lane: thread 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的线程块，可以填充一次经纱。

对于没有完全填满经线的维度，例如8,3(或8,29，这将使最后的经线仅部分满活动线程)，赋值顺序遵循相同的模式。当您将线程分配给warp时，快速变化的维度是x维度，如上所示。下一个变化最快的维度是y，然后是z。

编程指南还解释了如何按顺序为线程编号。当线程按此顺序编号时，前32个线程属于第一个经线，后32个线程属于下一个经线，以此类推。

如果在上述编号顺序中，特定的最后一次经络只有部分补充(即少于32个)可用的线程，则最后一次经络(仅)将由少于32个活动线程组成。

相关文章：