2D / 3D CUDA块如何划分成线程束？

Question

如果我使用一个具有以下块维度的网格启动我的内核：

dim3 block_dims(16,16);

网格块现在如何被划分为warp？这样的块的前两行是否形成一个warp，或者是前两列，还是任意顺序？

假设GPU计算能力为2.0。

- Gabriel

Answer 1

线程在块内按顺序编号，因此threadIdx.x变化最快，然后是threadIdx.y第二快，threadIdx.z变化最慢。这与多维数组中的列主序相同。warp按此顺序连续构建线程。因此，2d块的计算为：

unsigned int tid = threadIdx.x + threadIdx.y * blockDim.x;
unsigned int warpid = tid / warpSize;

这个问题在编程指南和PTX指南中都有涉及。

- talonmies

---

注意，“列优先顺序”假定dim3是一个数组，而不是一个结构体。更精确的描述是，.x是维度中变化最快的，.y是第二快变化的，.z变化最慢。如何将.x, .y和.z与内存中的行、列、切片、偏移量、树级别或任何其他寻址方式相对应，取决于您自己。 - harrism

Answer 2

为了阐明 @talonmies 的答案，我们通过 'Visual Studio WarpWatch' 窗口展示了两个连续的warp（dim3 block_dims(16,16); 和WarpSize=32）：

第一次warp的图片如下：，第二次warp的图片如下：。

- Mohsen