为什么 CUDA 中有 warp 级同步原语？
Why is there a warp-level synchronization primitive in CUDA?
我有两个关于 CUDA 中 __syncwarp() 的问题：
- 如果我没理解错的话，CUDA 中的 warp 是在 SIMD fasion 中执行的。这是否意味着 warp 中的所有线程总是同步的？如果是这样，
__syncwarp() 究竟做了什么，为什么有必要？
- 假设我们启动了一个块大小为 1024 的内核，其中一个块中的线程被分为每组 32 个线程。每个线程通过共享内存与其组中的其他线程通信，但不与其组外的任何线程通信。在这样的内核中，我可以看到比
__syncthreads() 更细粒度的同步可能是多么有用，但是由于块被分割成的扭曲可能与组不匹配，所以在使用 [=10= 时如何保证正确性]?
If I understand correctly, a warp in CUDA is executed in an SIMD fasion. Does that not imply that all threads in a warp are always synchronized?
没有。可以存在 warp 级执行分歧（通常是分支，但也可以是其他事情，例如 warp 洗牌、投票和预测执行），由指令重放或执行屏蔽处理。请注意，在 "modern" CUDA 中，隐式 warp 同步编程是 no longer safe，因此 warp 级别同步不仅是可取的，而且是强制性的。
If so, what exactly does __syncwarp() do, and why is it necessary?
因为可以 warp 级执行发散，这就是发散 warp 内同步的实现方式。
Say we have a kernel launched with a block size of 1024, where the threads within a block are divided into groups of 32 threads each. Each thread communicates with other threads in it's group via shared memory, but does not communicate with any thread outside it's group. In such a kernel, I can see how a more granular synchronization than __syncthreads() may be useful, but since the warps the block is split into may not match with the groups, how would one guarantee correctness when using __syncwarp()?
通过确保始终使用计算出的扭曲边界（或合适的线程掩码）明确执行拆分。
我有两个关于 CUDA 中 __syncwarp() 的问题：
- 如果我没理解错的话，CUDA 中的 warp 是在 SIMD fasion 中执行的。这是否意味着 warp 中的所有线程总是同步的？如果是这样，
__syncwarp()究竟做了什么，为什么有必要？ - 假设我们启动了一个块大小为 1024 的内核，其中一个块中的线程被分为每组 32 个线程。每个线程通过共享内存与其组中的其他线程通信，但不与其组外的任何线程通信。在这样的内核中，我可以看到比
__syncthreads()更细粒度的同步可能是多么有用，但是由于块被分割成的扭曲可能与组不匹配，所以在使用 [=10= 时如何保证正确性]?
If I understand correctly, a warp in CUDA is executed in an SIMD fasion. Does that not imply that all threads in a warp are always synchronized?
没有。可以存在 warp 级执行分歧（通常是分支，但也可以是其他事情，例如 warp 洗牌、投票和预测执行），由指令重放或执行屏蔽处理。请注意，在 "modern" CUDA 中，隐式 warp 同步编程是 no longer safe，因此 warp 级别同步不仅是可取的，而且是强制性的。
If so, what exactly does __syncwarp() do, and why is it necessary?
因为可以 warp 级执行发散，这就是发散 warp 内同步的实现方式。
Say we have a kernel launched with a block size of 1024, where the threads within a block are divided into groups of 32 threads each. Each thread communicates with other threads in it's group via shared memory, but does not communicate with any thread outside it's group. In such a kernel, I can see how a more granular synchronization than __syncthreads() may be useful, but since the warps the block is split into may not match with the groups, how would one guarantee correctness when using __syncwarp()?
通过确保始终使用计算出的扭曲边界（或合适的线程掩码）明确执行拆分。