同步线程同步

Question

除了同步线程块中的翘曲的__syncthreads()函数之外，还有另一个叫做__syncwarp()的函数。这个函数到底是做什么的？

cuda编程指南说，

将导致执行线程等待，直到掩码中指定的所有翘曲通道在恢复执行之前执行__syncwarp() (具有相同的掩码)。所有在掩码中指定的非退出线程都必须使用相同的掩码执行相应的__syncwarp()，否则结果是未定义的。 执行__syncwarp()可以确保参与屏障的线程之间的内存排序。因此，希望通过内存进行通信的warp中的线程可以存储到内存，执行__syncwarp()，然后安全地读取翘曲中其他线程存储的值。

那么，这是否意味着该函数确保了掩码所包含的翘曲中线程的同步？如果是这样的话，我们是否需要在同一翘曲中的线程中进行这样的同步，因为它们都确保在锁步中执行？

Answer 1

这一特性可在CUDA 9上使用，是的，它同步了经纱内的所有线程，对于发散的经纱非常有用。这对于Volta体系结构非常有用，在Volta体系结构中，翘曲中的线程可以单独调度。