CUDA __syncthreads()在warp中的用法

Question

如果块中的所有线程绝对需要在代码中的同一点,如果正在启动的线程数等于warp中的线程数,我们是否需要__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];

Answer 1

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

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

• 请记住,编译器会假设它可以优化,提供内部线程语义保持正确,包括将存储延迟到可以将数据保存在寄存器中的内存.__syncthreads()充当屏障,因此确保在其他线程读取数据之前将数据写入共享内存.使用volatile导致编译器执行内存写入而不是保留在寄存器中,但是这有一些风险并且更像是一个hack(意味着我不知道将来会如何影响它)
  • 从技术上讲,您应始终使用__syncthreads()符合CUDA编程模型
• 经线尺寸一直是32,但您可以:
  • 在编译时使用warpSize设备代码中的特殊变量(在CUDA编程指南中,在"内置变量"下,4.1版本中的B.4节中记录)
  • 在运行时使用cudaDeviceProp结构的warpSize字段(在CUDA参考手册中有记录)

请注意,某些SDK示例(特别是缩减和扫描)使用此扭曲同步技术.