何时需要在CUDA中使用二维线程？

Question

我想知道在CUDA中何时应该使用 x 和 y 坐标来控制线程？我看过一些代码，当它们有嵌套循环时，他们会使用 x 和 y 坐标。是否有一般规则？谢谢

- dibid

---

我能想到最常见的原因是，如果你有一个二维数据集，而你的CUDA线程策略为每个输出（或输入）点分配一个线程。那么，2D线程索引是表达代码的一种方便方式。另一个原因是对于cc2.x设备，第一个网格维度被限制为65535。所以启动非常大的1D网格变得不切实际，而你可以使用2D网格代替。有许多问题可以在这个标签上搜索，比如这个可能会给你一些思路。 - Robert Crovella

谢谢Robert，你真的很有帮助。 - dibid

Answer 1

标题中的问题的答案很简单：永远不需要。您永远真正需要2D坐标。

然而，它们实际上存在有几个原因。其中一个主要原因是它简化了某些问题的建模。特别是那些GPU擅长的问题，或者出于“历史”原因一直在使用的问题。我在想像图像处理或矩阵运算这样的问题。当您可以清楚地说出时，编写图像处理或矩阵乘法CUDA内核要直观得多：

int pixelX = threadIdx.x + blockIdx.x * blockDim.x;
int pixelY = threadIdx.y + blockIdx.y * blockDim.y;

从那时起，只需处理简单的像素坐标。 这实际上使索引麻烦更加明显，特别是在涉及共享内存时，例如在矩阵乘法期间，您想要从较大矩阵中切割一组行+列，以将其复制到本地内存中。如果您只有1D索引并且必须处理偏移和步长，那么这将容易出错。

CUDA实际上不仅支持2D，还支持3D内核，这可能源于3D纹理经常用于诸如体积渲染之类的事物，这也是可以通过GPU大大加速的内容（包含关键字“volume ray casting”的Web搜索将带您浏览一些漂亮的演示）。

（旁注：在OpenCL中，这个想法甚至已被泛化。虽然CUDA只支持1D、2D和3D内核，在OpenCL中，您只有N D内核，其中N 明确给出为 work_dim 参数）

_{（另一个旁注：我非常确定还有更多与GPU硬件架构或视频存储器缓存相关的低级技术原因，其中2D内核的局部性可能很容易被利用，并对整体性能有益-但我不熟悉这方面的内容，所以现在只是一个猜测）}

- Marco13

---

谢谢，我已经完全明白了 :) - dibid

我想补充一下，使用二维或三维坐标有时可以稍微提高性能。虽然你也可以在每个线程内将一维坐标转换为二维坐标，但如果你不使用2的幂，则需要较慢的取模或除法运算。使用二维或三维块，并直接从硬件寄存器读取x和y可以避免这些除法/取模操作。 - Jan Lucas