5

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

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

示例代码：

shared _voltatile_ sdata[16];

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

2 回答 2

8

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

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

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

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

于 2012-04-18T10:00:52.677 回答
1

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

于 2012-04-18T09:15:55.870 回答