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

Question

我有两个关于__syncwarp()CUDA的问题：

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

Answer 1

如果我理解正确的话，CUDA 中的扭曲是以 SIMD 方式执行的。这是否意味着扭曲中的所有线程始终是同步的？

不会。可能存在扭曲级别执行分歧（通常是分支，但也可能是其他事情，例如扭曲洗牌、投票和谓词执行），由指令重放或执行屏蔽处理。请注意，在“现代”CUDA 中，隐式扭曲同步编程不再安全，因此扭曲级别同步不仅是可取的，而且是强制性的。

如果是这样，__syncwarp() 到底做什么，为什么有必要？

因为可能存在扭曲级别的执行分歧，这就是在发散扭曲内实现同步的方式。

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

通过确保始终使用计算的扭曲边界（或合适的线程掩码）显式执行分割。