我正在 CUDA 中实现并行缩减。
内核有一个__syncthreads 来等待所有线程完成从共享内存中的2 次读取，然后将总和写回共享内存。
我应该使用__threadfence_block 来确保对共享内存的写入对于下一次迭代对所有线程都是可见的，还是使用__syncthreads 中给出的NVIDIA's example ？
【问题讨论】：
标签： parallel-processing cuda synchronization reduction
__syncthreads() 也暗示了内存围栏功能。这在documentation:
等待直到线程块中的所有线程都达到这一点并且这些线程在 __syncthreads() 之前进行的所有全局和共享内存访问对块中的所有线程都是可见的。
所以在这种情况下，除了__syncthreads() 之外，就没有必要使用__threadfence_block()
在通常的一般并行缩减中，您不能用 threadfence 函数代替执行屏障。除了内存防护功能外，还需要执行屏障 (__syncthreads())。在一般情况下，通常需要等待所有线程执行给定的一轮归约，然后再进行下一轮； __threadfence_block() 本身不会强制 warp 等待，而其他 warp 正在执行给定的一轮缩减。
因此__syncthreads()一般是必须的，假设你使用得当，__threadfence_block()一般是不需要的。
__syncthreads() 暗示 __threadfence_block()。
__threadfence_block() 并不暗示__syncthreads()
【讨论】：