parallel-processing - CUDA __syncthreads() 在 warp 中的使用

Question

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

注意:没有额外的线程或块，内核只有一个扭曲。

示例代码:

shared _voltatile_ sdata[16];

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

Answer 1

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

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

请记住，编译器将假定它可以优化提供线程内语义保持正确，包括延迟存储到内存中，数据可以保存在寄存器中。 __syncthreads()作为一个屏障，因此确保在其他线程读取数据之前将数据写入共享内存。使用 volatile导致编译器执行内存写入而不是保留在寄存器中，但是这有一些风险并且更像是一个黑客(意味着我不知道这将如何在 future 受到影响)

从技术上讲，您应该始终使用 __syncthreads()符合 CUDA 编程模型

经纱尺寸一直是 32，但您可以:

在编译时使用特殊变量 warpSize在设备代码中(记录在 CUDA Programming Guide 中，在“内置变量”下，4.1 版本的 B.4 部分)

在运行时使用 cudaDeviceProp 结构的 warpSize 字段(记录在 CUDA Reference Manual 中)

请注意，某些 SDK 示例(特别是缩减和扫描)使用了这种扭曲同步技术。