Warp & Lane
- Warp 是 GPU 中最基本的调度单位。
- Warp 是一个包含 32 个线程的执行单元。这些线程被称为 Lane。 Warp 中的所有 Lane 同时执行相同的指令。这被称为单指令多线程 (SIMT) 执行模型。
- 每个 Lane 都是一个独立的线程，拥有自己的数据和状态。它们可以访问自己的寄存器、私有内存和全局内存。
- 每个 Lane 都有一个唯一的 Lane ID，从 0 到 31 。可以使用内建函数
__Laneid()或threadIdx.x & 31来获取 Lane ID。 - CUDA 提供了一组 Warp shuffle 指令，允许 Lane 之间进行数据交换。这些指令包括
__shfl_sync()、__shfl_up_sync()、__shfl_down_sync()和__shfl_xor_sync()。 这些指令可以高效地进行 Warp 内部的数据通信，而无需访问共享内存。
可以通过使用 Warp shuffle 指令和 Lane ID 来间接地控制 Lane 的行为。不能直接对 Lane 进行编程。
Warp 之间没有（传统意义上的）上下文切换
有 Warp 调度器。所以 Warp 需要调度。
CUDA 并非以传统操作系统意义上的“上下文切换”方式在线程之间切换。 GPU 的 SM 会同时执行多个 Warp，并通过调度器动态地选择哪个 Warp 执行下一条指令。这更像是一种指令级并行，而不是线程级上下文切换。 没有显式的“切换”动作，而是并发执行。CUDA 利用 SIMT 架构，通过并发执行多个 Warp 来隐藏内存延迟和指令执行延迟，而不是通过频繁的线程上下文切换。
Warp 调度器会检查哪些 Warp 的指令已经准备好执行（例如，没有数据依赖性或内存访问冲突）。然后，调度器会从这些 “eligible” 的 Warp 中选择一个，并发出它的下一条指令。这个切换过程几乎没有开销，因为它不需要像传统上下文切换那样保存和加载线程状态。
所以没有传统意义上的上下文切换，没有保存和加载线程状态的开销（零开销），是零开销的指令流的交织执行。 ***
虽然理想情况下，我们希望所有 Warp 都能在同一时间一起执行，但由于 SIMT 架构的限制、延迟、资源冲突和依赖关系等因素，事实是 GPU 中有 Warp 调度器，它必须对 Warp 的执行进行调度，以实现最佳的性能。 ***
资源（寄存器、共享内存等）并非一直分配给每个线程直到它完成。 一个 SM 的资源是有限的。当一个 Warp完成执行后，其占用的资源会被释放，供其他 Warp 使用。 虽然单个线程的寄存器在 Warp 执行期间保持不变，但 SM 整体的资源分配是动态的，由调度器管理。
Warp 内部线程同一时刻执行相同的指令
SM 通过 Warp 调度器快速切换 Warp 来实现并发，而 Warp 内部的线程在理想情况下是同时执行的（SIMT）。 ***
@ Warp 内部的 reduce 为什么效率高
Warp 内部的 reduce 操作通常不需要额外的同步指令（如
__syncthreads()），因为 Warp 内的线程是同步执行的。在 Warp 内部进行 reduce 操作时，线程可以访问连续的内存地址，这符合GPU的内存访问优化原则。连续的内存访问可以进一步减少访存延迟。
而且通过 Warp shuffle 机制，
可以减少 bank conflict。Warp 不操作 Shared memory，直接访问寄存器的，但是寄存器也有bank，所以可能会有bank conflict。寄存器的bank conflict 很少被提及。
@ 每个 Warp 代表一个指令流，调度器负责在这些指令流之间进行动态切换？为什么要切换，不能在同一个时间一起执行吗
硬件就是这么设计的，有 Warp 调度器。没有为啥。
Warp Scheduler
负责选择要执行的 Warp。Warp Scheduler 根据一定的策略（例如 Round-Robin 或优先级）从 ready 的 Warp 队列中选择一个 Warp。
Dispatch Unit
主要作用是将 Warp Scheduler 选择的 Warp 的指令发送到 SM 中的各个执行单元。 负责将 Warp Scheduler 选择的 Warp 的指令发送到各个执行单元。
Warp 中线程 ID
dim3 blockDim(32, 32) 的情况下， Warp 的线程 ID 是如何分布的。
线程块内的线程会被划分成多个 Warp。划分的顺序是先按照 x 维度，再按照 y 维度。 *** x维度变化最快。
第一个 Warp 的线程 ID：threadIdx.x 从 0 到 31，threadIdx.y 都为 0 第二个 Warp 的线程 ID：threadIdx.x 从 0 到 31，threadIdx.y 都为 1 第三个 Warp 的线程 ID：threadIdx.x 从 0 到 31，threadIdx.y 都为 2 … 第 32 个 Warp 的线程 ID：threadIdx.x 从 0 到 31，threadIdx.y 都为 31
Divergence
GPU 采用 SIMT 架构，这意味着一个线程束中的所有线程在同一时刻执行相同的指令，当一个 Warp 中的线程执行不同的指令时，就会发生 divergence。比如 Warp 中的线程执行循环的次数不同时。假设一个线程束有8个线程，它们的循环次数分别是：
线程ID: 0, 1, 2, 3, 4, 5, 6, 7 循环次数: 10, 12, 15, 10, 18, 20, 12, 15
在这种情况下，GPU会按照以下步骤执行：
循环体执行10次：所有8个线程都处于活跃状态。
循环体执行2次：线程1和线程6处于活跃状态。
循环体执行3次：线程2和线程7处于活跃状态。
循环体执行5次：线程4处于活跃状态。
循环体执行2次：线程5处于活跃状态。
总执行次数：10 + 2 + 3 + 5 + 2 = 22 次,
@ 发散为什么会导致性能下降
- GPU 需要多次执行相同的指令，而每次只有部分线程处于活跃状态。当没有发散时，所有线程相同的指令执行一次。
- 线程束发散会导致 GPU 执行额外的指令，从而降低性能。活跃线程的比例越低，性能损失越大。
线程束发散会导致 GPU 执行的额外指令
主要包括以下几种：
1. 分支指令（Branch Instructions）
当线程束中的线程执行到分支语句（例如 if 语句或循环）时，如果某些线程满足条件，而另一些线程不满足条件，GPU 会生成额外的分支指令。这些分支指令用于控制线程的执行路径，使得不同的线程执行不同的代码块。
2. 屏蔽指令（Masking Instructions）
当线程束发散发生时，GPU 会使用屏蔽指令来屏蔽不满足条件的线程，使得它们不执行当前的代码块。屏蔽指令会增加指令的数量，并可能导致性能下降。
3. 线程同步指令（Thread Synchronization Instructions）
当线程束发散发生时，GPU 需要在不同的分支执行完毕后，将线程束中的所有线程重新同步。线程同步指令会增加指令的数量，并可能导致性能下降。
4. Predicated Execution 指令
虽然 predicated execution 可以减少线程束发散，但它也会增加指令的数量。Predicated execution 使用 @p 修饰符来控制指令的执行，这意味着每条指令都需要判断条件是否满足，这会增加指令的数量。