CUDA 中的扭曲改组是什么以及它为何有用？

Question

来自 CUDA 编程指南：

[Warp shuffle 函数] 在 warp 内的线程之间交换变量。

我知道这是共享内存的替代方案，因此它被用于扭曲中的线程来“交换”或共享值。但它背后的直觉是什么（它是如何工作的）？与使用共享内存相比，它有什么好处？

Answer 1

warp shuffle 是关于线程间通信的。在 warp shuffle 存在之前，在线程块中的线程之间交换数据的最直接、最有效的机制是使用共享内存，就像您在典型的共享扫描式缩减中所做的那样。

NVIDIA 没有提供 warp shuffle 的详细设计（它是如何工作的？），但在行为层面上，它允许在 warp 中从一个线程到另一个线程直接交换寄存器数据（例如线程局部变量），使用相当灵活的源/目标描述。为了使 shuffle 操作合理，源线程和目标线程都必须参与。编程指南中给出了额外的行为描述，并且 SO 标签上有许多问题cuda对此进行了讨论。

一个例子可以是：

int r = __shfl_sync(0xffffffff, value,     0); 
    ^                 ^         ^          ^
destination           bit mask  source     source lane
variable              for       variable
                      threads 
                      which 
                      must 
                      participate

r在上述 shuffle 操作之后， warp 中每个线程的变量将包含valuewarp 中线程 0 所持有的数量。线程 0 已将其“广播”value给其他线程。

Warp Shuffle 的主要优点是：

1. 更少的指令/步骤/操作/（以及更低的延迟）：而通过共享内存将数据项从线程 B 中的寄存器传送到线程 A 中的寄存器需要至少 2 个步骤（共享加载指令和共享存储指令，并且可能还有一个同步步骤），通过 warp shuffle 进行的相同通信需要单个操作/指令。

2. 减少“共同压力”：这有两个方面。首先，可以减少所使用的共享存储器的总量。Warp shuffle 本身不需要共享内存，甚至成熟的 1024 个线程减少也只需要 32 个共享内存存储元素。由于共享内存是一种宝贵的资源（大小非常有限），并且共享使用也可能是占用的限制因素，因此这在两个方面都带来了好处：更多的共享内存可用于算法的其他部分，并且如果共享使用是关心占用，那么这可以在很大程度上避免共享使用（对于可以通过随机操作处理的通信模式）。其次，共享压力可以表现为共享内存的吞吐量。例如，根据共享内存减少的实现方式，共享内存吞吐量可能是性能的限制因素。因此，如果我们可以将部分或全部共享内存事务转移到备用路径/单元，则可能会减少共享事务压力，如果这是一个限制器，这将提高代码性能。

如上所述，使用 warp shuffle 的另一个可能的好处可能是减少线程块同步指令（例如__syncthreads()）的需要/使用，这在 CUDA 编程中通常是一件好事。

尽管我主要关注减少激励示例，但 shuffle 操作可用于构建其他类型的简洁操作，例如前缀和。