为什么在CUDA中有翘曲级同步原语？

Question

我对CUDA中的__syncwarp()有两个问题：

1. 如果我理解正确，将在SIMD功能中执行CUDA中的扭曲。难道不意味着扭曲中的所有线程都始终同步吗？如果是这样，__syncwarp()到底是做什么的，为什么有必要？
2. 假设我们启动了一个内核，其块大小为1024，其中一个块中的线程被分成32个线程的组。每个线程都通过共享内存与该组中的其他线程通信，但不与该组外的任何线程通信。在这样的内核中，我可以看到比__syncthreads()更细粒度的同步可能有用，但是由于将块分割成的翘曲可能与组不匹配，使用__syncwarp()时如何保证正确性？

Answer 1

如果我理解正确，将在SIMD功能中执行CUDA中的扭曲。难道这并不意味着扭曲中的所有线程总是同步的吗？

没有可能存在warp级执行差异（通常是分支，但也可能是warp shuffle，表决和谓词执行等其他事物），由指令重播或执行屏蔽处理。请注意，在“现代” CUDA中，隐式翘曲同步编程为no longer safe，因此翘曲级同步不仅可取，而且是必需的。]

如果是这样，__ syncwarp（）的确切作用是什么，为什么有必要？

因为存在可以

扭曲级别执行差异，所以这是在离散扭曲中实现同步的方式。

假设我们启动了一个内核，其块大小为1024，其中一个块中的线程被分成32个线程的组。每个线程都通过共享内存与该组中的其他线程通信，但不与该组外的任何线程通信。在这样的内核中，我可以看到比__syncthreads（）更细粒度的同步可能有用，但是由于将块拆分为warp可能与组不匹配，因此在使用__syncwarp（）时如何保证正确性？

通过确保始终使用计算的经纱边界（或适当的线程掩码）显式执行拆分。