经线中的 CUDA __syncthreads() 使用答案

【问题标题】：CUDA __syncthreads() usage within a warp经线中的 CUDA __syncthreads() 使用
【发布时间】：2012-04-18 07:56:55
【问题描述】：

如果一个块中的所有线程绝对需要在代码中的同一点，如果正在启动的线程数等于线程数中的线程数，我们是否需要 __syncthreads 函数？

注意：没有额外的线程或块，只是内核的一个 warp。

示例代码：

shared _voltatile_ sdata[16];

int index = some_number_between_0_and_15;
sdata[tid] = some_number;
output[tid] = x ^ y ^ z ^ sdata[index];

【问题讨论】：

我认为这在很大程度上是一个无关紧要的问题，因为如果您在每个块中运行 warpSize (32) 个线程，您最终的性能可能会低于运行您的CPU 上的算法。
@RogerDahl 我不确定我是否完全理解，你是说在 warp 中使用 32 个线程总是比任何内核的 CPU 都慢？
我没有说“总是”，因为人们可能会想出一组我的陈述不适用的特殊情况。问题是，如果您将threads per block 限制为warpSize，您最终可能会得到极低的占用率（受max blocks per multiprocessor 的限制），这意味着性能低下。此外，当您可以简单地增加 threads per block 并添加 __syncthreads() 以获得更好的性能时，人为地将 threads per block 保持在低水平是没有任何意义的。
因此，如您所见，使用限制为 16 个线程的算法，您可以做的任何优化都是徒劳的。在 CPU 上运行它会更好。
嘿。 SO 告诉我避免在 cmets 中进行扩展讨论，并希望我将其移至聊天。因此，如果您在阅读本文后仍有疑问，请继续并提出一个新问题。引用http.developer.nvidia.com/GPUGems3/gpugems3_ch36.html:但是，因为CBC模式需要前一个步骤的密文来处理下一步，所以在前一个块被加密之前不可能开始加密。所以我们不能指望在这种模式的加密阶段进行并行处理。

【解决方案1】：

更新了有关使用 volatile 的更多信息

假设您希望所有线程处于同一点，因为它们正在将其他线程写入的数据读取到共享内存中，如果您正在启动单个 warp（在每个块中），那么您知道所有线程都在一起执行。从表面上看，这意味着您可以省略__syncthreads()，这种做法被称为“warp 同步编程”。但是，有一些事情需要注意。

请记住，如果线程内语义保持正确，编译器会假定它可以进行优化，包括延迟将数据存储到可以将数据保存在寄存器中的内存中。 __syncthreads() 充当屏障，因此确保在其他线程读取数据之前将数据写入共享内存。使用volatile 会导致编译器执行内存写入而不是保存在寄存器中，但是这有一些风险并且更像是一种黑客攻击（这意味着我不知道这在未来会受到怎样的影响）
- 从技术上讲，您应该始终使用 __syncthreads() 以符合 CUDA 编程模型
经线大小一直是 32，但您可以：
- 在编译时在设备代码中使用特殊变量warpSize（记录在CUDA Programming Guide，在“内置变量”下，4.1 版本的 B.4 节中）
- 在运行时使用 cudaDeviceProp 结构的 warpSize 字段（记录在 CUDA Reference Manual 中）

请注意，一些 SDK 示例（尤其是缩减和扫描）使用这种扭曲同步技术。

【讨论】：

当你说“如果核心数量小于经线大小，这适用”，这是否意味着 C1060 是不可能的？另外，如果我在内核中只启动 16 个线程呢？
C1060 还是可以的。您真正关心的是对某个位置的写入，然后是读取（或 WAR/WAW），并且您知道整个 warp 将在 warp 中的任何线程开始后续读取之前完成写入。如果您启动的线程少于 32 个，则同样的行为适用，尽管效率较低，但它仍然只是一个扭曲。请注意，我强烈建议您考虑启动更大的块，每个块只有一个 warp，您会发现无法覆盖大多数延迟。
我添加了一个我正在使用的类似代码的示例。线程将读取在前一行代码中写入的数据。因为对共享数据也进行了一些按位运算，我可以将按位运算分离出来并在读取之前放置它们吗？ warp 中的所有线程是否仍会执行相同的代码行？

【解决方案2】：

你仍然需要__syncthreads()，即使warp 是并行执行的。硬件中的实际执行可能不是并行的，因为一个 SM（Stream Multiprocessor）内的内核数量可能少于 32 个。例如，GT200 架构在每个 SM 中有 8 个内核，因此您永远无法确定所有线程都在代码中的相同点。

【讨论】：

__syncthreads() 可能会对一些密集型代码（例如缩减和扫描）产生性能影响。当您想在线程之间共享数据（例如通过共享内存）时，就会出现此问题。在这些情况下，如果您有一个存储后跟一个加载，您知道整个经线将在任何线程开始加载之前执行存储，即使 LS 单元的数量（CUDA 核心的数量与此无关）小于经纱尺寸。编译器优化还有其他一些注意事项！
-1：CUDA 模型保证一个 warp 中的所有线程都是同步的。编程指南中的引述：因为 warp 一次执行一条公共指令，所以 warp 中的线程是隐式同步的，有时可以使用它来省略 __syncthreads() 以获得更好的性能。
较新的 Kepler_Tuning_Guide.pdf 第 1.4.8 节贬值通过消除 __syncthreads() 来利用 Warp 同步编程。