有没有办法将线程显式映射到 CUDA 中的特定扭曲？

Posted 2023-04-15

【中文标题】有没有办法将线程显式映射到 CUDA 中的特定扭曲？

【英文标题】：Is there a way to explicitly map a thread to a specific warp in CUDA?

【发布时间】：2017-08-18 00:15:06

【问题描述】：

比如说，动态分析是在一个 CUDA 程序上完成的，因此某些线程最好在同一个 warp 中。

例如，假设我们有 1024 个 cuda 线程和 32 个 warp 大小。经过动态分析，我们发现线程 989、243、819、...、42（总共列出了 32 个线程）应该在同一个经。我们确定它们应该在同一个 warp 上，因为它们在代码执行方面几乎没有分歧——（在执行 CUDA 程序的动态分析时，它们可能不一定在同一个 warp 上）。

有没有办法控制线程在 CUDA 中扭曲调度？如果没有，是否有另一种 GPU 编程语言可以提供这种显式的扭曲调度。如果没有，可以做些什么（甚至可能是解决这个问题的非常低级的方法）？我希望至少有最后一个问题的答案，因为这可能是 CUDA 的实现方式——除非在硬件级别完成扭曲调度，这将是不幸的。谢谢！

【问题讨论】：

见***.com/questions/21535471/…

【参考方案1】：

不，您不能选择将线程分配给经纱。对此声明的支持覆盖here。

但是，使线程具有特定行为特征的是您编写的线程代码，而不是任何锁定在硬件中的东西。无论是数据访问模式，还是通过控制流的特定路径，都由程序员控制。

CUDA 线程发挥其特性的主要方式之一是通过生成全局唯一线程 ID，这是任何 CUDA 代码的典型样板，例如：

int idx=threadIdx.x+blockDim.x*blockIdx.x;

为每个线程创建一个规范的、全局唯一的一维线程索引。

但没有什么特别的理由必须这样。我也可以轻松做到：

int private_idx = threadIdx.x+blockDim.x*blockIdx.x;
int idx = desired_idx[private_idx];

然后可以按您想要的任何顺序对线程进行编号。如果您的 desired_idx 数组按照您的建议包含一组数字：

989, 243, 819, ..., 42

然后那些相邻的线程将采取与该顺序一致的行为。

【讨论】：

这是一个很好的解决方案——谢谢！此外，您链接到的堆栈溢出线程很有见地。特别是，warp 调度不是动态的——线程 0-31 总是映射到同一个 warp。正如您所提到的，我们可以强制线程 0-31 具有 desired_index。