GPU 调度这块，我几年前刚开始认真写 CUDA kernel 的时候，最直观的感受是：明明照着官方文档把线程块大小设成了 256， occupancy 算出来也还行，但性能就是上不去。后来把 NVIDIA 的任务调度模型真正啃了一遍，才明白瓶颈到底卡在哪。这篇是这个系列的第二章，前面聊完整体调度分析的大框架，这次把 NVIDIA 的任务调度模型单独拉出来拆开揉碎讲清楚。
先说清楚这篇要解决什么问题。很多人把“调度”理解成操作系统里那种进程/线程切换，觉得 GPU 调度就是一堆线程排队上核心。这个理解错得很离谱。GPU 的调度是硬件级的、以线程束（warp）为最小单位的、基于状态机轮转的调度方式，它压根不走操作系统内核，也不需要什么上下文切换的软件开销。不理解这一层，你写 kernel 的时候连“为什么 block 大小设成 128 有时比 256 好”这种问题都答不上来。
这篇文章的受众，我默认是对 CUDA 编程有一定基础、但没系统研究过 GPU 内部机制的开发者。如果你手头有 NVIDIA 的显卡，跟着思路把 profiling 数据打开对照着看，理解速度会快很多。下面先把最关键的几张“硬件底图”铺开。
1. 先从硬件底图说起：GPC、TPC 与 SM 的关系
1.1 一张图理解 GPU 的物理结构分层
NVIDIA GPU 的物理结构，从上到下大概可以切成这么几层：整个芯片叫 GPU，芯片内部被划分成多个GPC（Graphics Processing Cluster，图形处理集群），每个 GPC 内部包含若干个TPC（Texture Processing Cluster，纹理处理集群），TPC 再往下拆就是SM（Streaming Multiprocessor，流式多处理器）。
每一层各管各的事：
- GPC：可以理解成芯片里的“大园区”，它内部有独立的几何引擎、光栅化单元、以及若干个 SM。图形渲染任务里的三角形装配、光栅化这类活，就是在 GPC 这个级别分工的。
- TPC：一个 TPC 通常包含两个 SM（老架构里也有一个 TPC 只挂一个 SM 的情况）。TPC 主要起的是“分组管理”作用，让纹理单元和 SM 之间的数据通路更紧凑。
- SM：这才是真正的“计算工人”，所有 CUDA 核心、Tensor Core、共享内存、寄存器堆、调度器，全都住在 SM 里面。SM 是任务调度真正发生的地方。
拿 Ampere 架构的 GA102 核心来举例：它一共有 7 个 GPC、42 个 TPC、84 个 SM，每个 SM 里有 128 个 CUDA 核心，整卡总共就有 10752 个 CUDA 核心。这个数字看着唬人，但你要记住，这些 CUDA 核心并不是像 CPU 那样被“单个指令流”驱动，它们是靠 SM 内部调度器喂数据才能干活的。
这里插一句我踩过的坑：早期我误以为一个 CUDA 核心对应一个线程，那我开一万个线程就能让一万个核心满负荷跑。实际上线程和核心的关系是通过线程束和调度器间接建立的，线程数超过核心数不代表就能并行执行，得看调度器能喂多快。
1.2 SM 内部的关键部件和调度相关寄存器
SM 内部藏着几样和调度直接相关的“家当”，我按重要性排个序：
- 线程束调度器（Warp Scheduler）：每个 SM 里有 4 个（Volta 之后普遍是 4 个），这是调度的“大脑”，负责从就绪的线程束里挑一个发指令。
- 指令分发单元（Dispatch Unit）：每个调度器旁边配一个，负责把取到的指令真正发给执行单元。
- 寄存器堆（Register File）：每个 SM 里有 65536 个 32 位寄存器（64K 是 Ampere 常见的配置）。这些寄存器是给线程存局部变量的，总容量固定，所以线程越多、每个线程能分到的寄存器就越少。
- 共享内存（Shared Memory）：常见 128KB 或 164KB 每 SM，用于线程块内部通信。
- Warp 状态表：SM 内部维护一张表，记录每个 warp 当前处于什么状态（就绪、等待、阻塞等），这是调度器做决策的依据。
这些部件之间的关系，我习惯用“餐厅后厨”来类比。GPC 是餐饮集团的一个门店，TPC 是店里的几个档口，SM 就是一个灶台，warp 就是灶台上的炒锅。炒锅（warp）里装着一批菜（线程），灶台师傅（调度器）按照火候情况决定先翻炒哪口锅。每个灶台不止一口锅，师傅也不止一个，关键是谁能让锅不闲着。
2. 任务调度的最小单位：warp 的概念与意义
2.1 为什么偏偏是 32 个线程
NVIDIA 的硬件线程调度单位，不是单个线程，而是warp（线程束），一个 warp 恰好包含 32 个线程。这 32 个线程在硬件层面是“锁步”（lockstep）执行的，也就是说，同一个 warp 里的所有线程在同一时刻执行同一条指令，只是处理的数据不同。
为什么是 32 不是 16 也不是 64？这里没有特别神秘的答案，主要是硬件设计上的折中。32 个线程一条指令，意味着指令 fetch 和 decode 的开销被 32 份工作分摊了，这个倍数刚好能填满现代 GPU 执行单元的多周期流水线。历史上 Fermi 时代就定了 32 这个数，后续架构一直沿用，说明这个设计经过了充分的实践验证。
你可以这样想：CPU 一个核一次就处理一个线程的指令，GPU 一个调度器一次处理 32 个线程的指令，这就是 GPU 能靠“少控制、多并行”换取吞吐量的根本原因。代价是，warp 内如果出现分支发散，性能会断崖式下跌。
2.2 分支发散为什么是性能杀手
想象一个 warp 里有 32 个线程，代码里写了一个 if-else 分支，一半线程走 if，一半走 else。硬件没法让这 32 个线程同时执行不同指令，于是调度器只好先让走 if 的线程执行，再让走 else 的线程执行，最终这个 warp 需要串行执行两个分支的指令。
这意味着什么？本来一个周期能完成的指令，现在要两个周期；如果一个分支里嵌套了更多分支，耗时还要成倍上涨。写 kernel 的时候，最怕的不是分支本身，而是同一个 warp 里出现“有的线程走 A、有的线程走 B”的分叉。
代码层面怎么避免？核心思想是让分支判断基于 warp 内统一的值，而不是基于线程独有的数据。比如可以根据threadIdx.x / 32来分块处理，让同一 warp 内的线程走同一条路径；或者把分支逻辑改成算术运算，用掩码来“选择”结果，而不是真写 if。
实操心得：我早期写归约（reduction）核函数的时候，用了
if (tid % 2 == 0)这种写法做折叠操作，结果性能比预期慢了一倍还多。后来改成让每个线程固定处理连续两段数据，再配合 warp shuffle，性能直接翻倍。分支发散这个坑，真的是你不亲自拿 profiling 数据对比，光靠感觉完全发现不了。
2.3 线程束与线程块的换算关系
一个线程块（block）由若干 warp 组成，block 内的线程数必须是 warp 大小的整数倍吗？硬件不强制，但实践里强烈建议这样做。因为如果你把 block 大小设成 100，硬件会把它补成 4 个 warp（128 个线程），其中 28 个线程是“空转”的，它们不干任何活，但照样占用调度资源和寄存器。
所以你会看到业界几乎所有 kernel 的 block 大小都习惯性设成 32 的倍数：128、256、512 都有，但很少看见 100、300 这种数字。把 block 设成 warp 大小的整数倍，是最基本的入门礼仪。
3. SM 调度器的工作机制：从取指到发射的完整链路
3.1 四个调度器如何分工协作
前面提到 Ampere 架构的 SM 里有 4 个 warp scheduler，每个 scheduler 配一条独立的指令发射通路。这意味着一个 SM 在同一个时钟周期内，最多能从 4 个不同的 warp 里各取一条指令发射出去。
但这不意味着一个周期只能发射 4 条指令。NVIDIA 的调度器支持dual-issue，也就是一个调度器在一个周期里可以发射两条指令，前提是这两条指令互不依赖、并且执行单元有空闲。所以理论上一个 SM 一个周期最多可以发射 8 条指令（4 个调度器 × 2 条）。
调度器怎么决定发射谁的指令？底层逻辑是每个 warp 在每个周期都有“状态位”，要么是“就绪可发射”（eligible），要么是“等待某种资源”（stalled）。调度器从 eligible 的 warp 里挑一个（具体策略各家 NVIDIA 架构里略有不同，常见的是轮转和优先级结合的方案）。
3.2 指令发射的延迟隐藏原理
GPU 的调度器做一次“挑 warp 并发射指令”的动作，在硬件层面只要几个时钟周期。但它真正的精髓，在于用大量 ready 的 warp 来掩盖长延迟操作。比如从全局内存取数，延迟可能有几百个周期，CPU 的做法是停下等，GPU 的做法是：这个 warp 等着取数，调度器立刻切换到另一个就绪的 warp 继续发射指令。
这个机制叫latency hiding（延迟隐藏）。它要求 SM 里同时有足够多“蓄势待发”的 warp，否则一旦所有 warp 都在等内存返回，SM 就“空转”了，也就是所谓的occupancy 不足。
来算一笔账。一个 SM 有 4 个调度器，每个调度器每个周期最多发射 2 条指令，也就是每个周期最多需要 8 个 warp 供它轮转。假设每个 warp 的平均停顿周期数是 200（等内存），那 SM 里至少要同时有 8 × 200 = 1600 个 warp 在轮转，才能把每周期 8 条发射管线喂满。而实际上 SM 里能容纳的 warp 总数是有限的——Ampere 架构最多 64 个 warp（2048 线程）per SM。这个远小于 1600，所以现实是 GPU 只能做到“尽量隐藏延迟”，不可能完全隐藏。
这就是为什么occupancy 高不代表性能一定好，但 occupancy 太低性能一定好不了。如果你的 kernel 每个线程用了太多寄存器，导致 SM 里能同时常驻的 warp 数减少，延迟隐藏能力就下降，内存延迟就会暴露出来。
3.3 调度器怎么知道你“准备好了”
这个细节容易被忽略。每个 warp 的每条指令在进入执行流水线之前，要经过所谓的scoreboard（记分板）机制。记分板记录每个 warp 的每条指令的输入操作数是否已经准备好。比如全局内存 load 指令的结果还没返回时，这条指令会卡在 scoreboard 阶段，warp 被标记为 stalled。等到数据写回寄存器了，scoreboard 更新状态，warp 变成 eligible，调度器才会把它纳入候选。
这个机制和 CPU 里的动态调度有点类似，但 GPU 的实现是分布在各个 warp 状态里的，所有跟踪逻辑都是硬件电路，没有任何软件参与。这也是为什么 GPU 调度能做到纳秒级别，而操作系统进程调度要微秒甚至毫秒级别的原因。
4. 线程块到 SM 的映射策略
4.1 线程块如何分配到 SM
前几节讲的是 warp 在 SM 内部的调度，还有一个重要问题是：线程块怎么分配到不同的 SM 上。这个分配是硬件完成的，由GigaThread Engine（全局调度器）来负责。当你 launch 一个 kernel 时，CUDA runtime 把整个 grid（网格）的任务描述交给 GigaThread Engine，它会按顺序把一个一个线程块（block）分发到各个 SM 上。
分配的依据是 SM 的资源容量：寄存器和共享内存还够不够容纳一个新的 block？如果够，就分配；如果不够，就等当前在跑的 block 里有线程退出、资源释放了，再分发新的 block 进来。这就是为什么你 launch 一个超大 grid（比如一百万个 block）时，GPU 并不会一次全收下，而是“边跑边补”。
4.2 并发上限：怎么算一个 SM 能装下多少 block
这里给一个具体的计算例子。假设我在 Ampere 架构的 A100 上写 kernel，每个 SM 最多 2048 个线程、32 个 block（软件限制）、65536 个寄存器、164KB 共享内存。我的 block 设成 256 线程，每线程用 32 个寄存器：
- 按线程上限：2048 ÷ 256 = 8 个 block
- 按寄存器上限：65536 ÷ (256 × 32) = 8 个 block
- 按 block 上限：32 个 block
三个限制取最小值，所以这个配置下每个 SM 只能装 8 个 block，换算成 warp 数就是 8 block × 8 warp/block = 64 个 warp，刚好顶到硬件上限，occupancy 100%。
如果把每线程寄存器数改成 64：
- 按寄存器上限：65536 ÷ (256 × 64) = 4 个 block
这时候每个 SM 只能装 4 个 block，总共 32 个 warp，occupancy 掉到 50%。如果内核本身没有太多寄存器压力，只是编译器默认给了很宽的寄存器窗口，那降低 occupancy 就得不偿失。你可以用__launch_bounds__(256, 8)提示编译器限制每线程寄存器数，让它去 spill 一些非常用变量，从而换取更高的 occupancy。
4.3 block 大小和调度效率的微妙关系
block 设得越大，单个 block 包含的 warp 越多，调度器在 block 内部可轮转的 warp 也越多。block 设得越小，block 间切换更灵活，更容易做到负载均衡。但是 block 太小也有问题：假设一个 block 只有 32 线程（1 个 warp），那 SM 要管 64 个 block（如果 occupancy 100%），block 调度的簿记开销会变大，而且每个 block 能用的共享内存量也少。
我个人的经验是，常规计算型 kernel 用 256 线程/block 是最不容易出错的起点，因为它在延迟隐藏效率和块级负载均衡之间取得了一个比较均衡的位置。需要调优的时候，再往 128 或 512 两个方向测试，结合 profiling 数据做决定，而不是拍脑袋硬改。
5. 从 GPU 到 CUDA 软件栈的任务下发链路
5.1 CPU 侧如何把任务交给 GPU
你写 CUDA 代码的时候，kernel launch 在 CPU 侧只做一件事：把任务描述扔进命令缓冲区，然后立即返回（异步行为）。GPU 驱动和硬件之间通过命令处理器（Command Processor，或叫 Host Interface）沟通，驱动把核函数入口地址、grid/block 尺寸、参数列表都写进一个叫control buffer的内存区域，然后写一个门铃寄存器（doorbell register）告诉 GPU“有新任务了”。
GPU 侧的工作调度器（Work Distributor）收到信号后，开始读取控制缓冲区里的任务描述，把 grid 拆成 block，通过 GigaThread Engine 分发到各个 SM。整个过程都是硬件帮衬的，CPU 只负责“下单”，GPU 自己决定“怎么做、做多快”。
实操心得：如果你发现 kernel launch 的开销特别大（用 Nsight Systems 能看到 launch 耗时几个微秒以上），通常不是 GPU 慢，而是 CPU 侧驱动栈和命令缓冲区的开销。这时候可以试试 CUDA Graph，把多个 kernel 的依赖关系预先构建成一张图，一次 launch 整张图，能把启动开销摊薄到极致。这一点对短小 kernel 频繁调用的场景尤其有效。
5.2 任务调度模型视角下的流（Stream）与事件（Event）
CUDA 的Stream在任务调度模型里扮演什么角色？你可以把它理解成一条“任务流水线”，stream 之间的任务是可以并行的（前提是硬件资源足够），同一个 stream 内部的任务保持顺序执行。
底层来看，每个 stream 对应一条独立的命令队列，GPU 的多个队列之间可以并发取任务。但要注意，“可以并行”不等于“一定并行”，如果两个 stream 的任务都要用到同一个 SM 上的全部资源，那它们实际上还是串行的。
事件（Event）则用于跨 stream 同步：一个 stream 里可以cudaStreamWaitEvent(stream, event)，强制这个 stream 等某个事件完成才开始后序任务。这本质上是把依赖关系告诉硬件调度器，让它别做“优化排序”把有先后依赖的任务搞乱。
5.3 MPS、多进程场景下的调度差异
多进程共享 GPU 的情况，和单进程多 stream 的调度模型有很大不同。MPS（Multi-Process Service）允许来自不同进程的 kernel 同时被调度到同一个 SM 上，这相当于把 GPU 的调度单位从“进程”进一步细粒度化了。
普通模式下，多个进程要时间片轮转地独占整个 GPU；MPS 模式下，不同进程的 block 可以被交错调度到同一个 SM 上，利用率显著提升，但代价是单个 kernel 的延迟可能变高，因为要和别人共享执行单元。如果你在做推理服务，下游可能有多个模型实例，MPS 几乎是把 GPU 利用率拉满的必备技能。
6. 任务调度模型的性能关键参数与调优策略
6.1 Occupancy 的计算与解读
Occupancy 的定义：一个 SM 上实际活跃的 warp 数占硬件最大容量的比例。它是衡量“资源利用充分度”最直观的指标。用 CUDA 自带的 occupancy calculator（或cudaOccupancyMaxActiveBlocksPerMultiprocessorAPI）可以直接查到给定 block 大小和寄存器/共享内存用量下的理论 occupancy。
但我要给一句忠告：不要把 occupancy 当唯一指标。我见过很多 kernel，occupancy 从 50% 提到 100%，性能反而下降了。原因是高 occupancy 意味着每个线程能用的寄存器更少，编译器不得不把局部变量吐到本地内存，本地内存虽然走 L1/L2 缓存，但比寄存器慢一两个数量级。这时候低 occupancy 但高寄存器命中率的配置反而更快。
调优的正确路径应该是：用 profiling 工具看 stall 原因（memory dependency、execution dependency、barrier 等），是“因为等内存而空转”就加 occupancy，是“计算单元太挤”反而要降。
6.2 常见 stall 原因与排查思路
Nsight Compute 的 profiler 里，你能看到每个 kernel 的 warp stall 周期分布。几个常见的 stall reason 对应的调优方向：
- Long Scoreboard：等待全局内存或共享内存数据返回。核心思路是提高内存访问局部性、用
__ldg走只读缓存、加大 block 规模以提供更多独立 warp。 - Barrier：在等待同一个 block 内其他 warp 到达 barrier（如
__syncthreads()）。这说明 block 内同步太频繁或负载不均，可以考虑减少同步次数、让每个线程做更多独立工作。 - MIO Throttle：访问共享内存或特殊指令（如 shuffle）触发指令队列拥塞。可以降低共享内存访问密度，把部分数据挪到寄存器。
- Not Selected：warp 已就绪但调度器选了其他 warp。这个通常无害，说明调度器有得选，不算性能瓶颈。
实操心得：把
Nsight Compute的 Sampling Data 打开，按 stall reason 排序，基本上一眼就能定位 kernel 的短板。这种“用数据说话”的方式，比反复猜“是不是 block 大小不对”高效得多。我在调一个流体模拟 kernel 的时候，一直以为瓶颈是全局内存带宽，结果 profiling 一看是 barrier 同步太多，改完同步策略性能提升 40%。
6.3 调优实例：从 60% occupancy 到 90% plus
分享一个实际的调优案例。当时写一个金融领域的 Monte Carlo 模拟内核，基础版本参数是 block=256、每个线程 96 个寄存器，occupancy 只有 37.5%（8 个 block 上限被寄存器限制卡死），实测带宽不到 40%。
第一步：用__launch_bounds__(256, 6)限制每线程寄存器数，编译器把一些冷变量 spill 到本地内存，寄存器降到 40 个。occupancy 提升到 75%，但因为 spill 增加了内存流量，性能只提升了 15%。
第二步：改代码结构，把大多数 spill 的变量改成float精度（原来是double），并把部分数组改成手工分块放入共享内存。这一步将 spill 消除大半，occupancy 达到 87.5%，性能又提升了 32%。
第三步：调整 block 到 128，每个 SM 的 block 数从 14 个升到 28 个（受 block 上限限制），实际活跃 warp 数不变但 block 调度更均衡，最终性能比初版提升了 60% 以上。
这个例子说明，occupancy 只是表象，真正的坑在资源使用的细粒度上。你要同时盯着寄存器、共享内存、spill 三条线，才知道到底是谁限制了性能。
7. 常见问题与排查技巧实录
7.1 同一个 kernel 在不同 GPU 上表现天差地别
换卡之后某些 kernel 效率骤降，是特别常见的事。原因通常是不同架构的 SM 资源参数不同：寄存器文件大小、共享内存大小、warp 调度器数量、甚至 warp 大小（基本都是 32）都可能变化。老卡上压出来的魔数参数（比如 block=512、每线程 64 寄存器）直接搬到新卡上，很可能把新卡的 occupancy 压得很低。
排查思路：跑一遍deviceQuery拿设备参数，再重新走一遍 occupancy 计算。遇到这种情况，代码里的硬编码参数应该抽成宏或参数对象，按设备属性动态计算。
7.2 Kernel launch 失败但没报错
一种让人挠头的情况是：kernel 调用之后，后面的cudaMemcpy报“invalid argument”或直接卡死。常见原因是 block 数超过了硬件的 grid 上限（x 方向一般是 2^31 - 1，y/z 方向有更严格限制），或者共享内存申请量超过了默认上限导致 launch 失败，但错误码被忽略了。
排查方法很简单：每次 kernel launch 后检查返回值，或者设置cudaDeviceSetLimit调大共享内存上限。这个坑我踩了不止一次，奉劝大家 launch kernel 后立刻cudaGetLastError()，成本几乎为零，能救你几个小时。
7.3 分支发散比想象中隐蔽
有些发散不是 if-else 那种明显的写法，而是循环的退出条件不同。同一个 warp 里 32 个线程执行一个 while 循环，循环次数有的线程是 10 次、有的是 20 次，那这个 warp 实际上要执行 20 次迭代，前 10 次所有线程都活跃，后 10 次只有一部分线程活跃。这种发散在 profiling 里不容易一眼看出来，但计算量被白白拉高了。
这种场景的优化方式是：把循环拉平（flatten）或改成针对固定迭代次数的展开循环；或者预计算每个线程的“终止位置”，用掩码跳过多余迭代。实在无法解决时，考虑用“分桶”的思路，让相同循环次数的线程分到同一 warp。
7.4 cudaDeviceSynchronize 卡住不动
这个几乎都会碰到。它本身不是调度模型的锅，但理解了调度就知道为什么它危险：cudaDeviceSynchronize会让 CPU 阻塞等待 GPU 上所有已提交任务完成，如果 GPU 上一个 kernel 是死循环或等待一个永远不满足的条件（典型是 kernel 里访问未正确初始化的全局内存），CPU 就会一直挂着。
排查思路：先用 cuda-gdb 或 nsight 把卡住的 kernel 找出来，然后检查 kernel 内是否有循环条件永远为真、共享内存索引是否越界、以及是否误用了threadIdx和blockIdx的组合。不要小看这个，我见过有人在if (blockIdx.x > gridDim.x)这种永远是假的条件里加了死循环，排查了两小时才找到。
7.5 多 stream 并行时性能不升反降
多 stream 并行没提速，甚至更慢了，大概率是 kernel 太小、CPU 侧 launch 开销占了大头，或者是多个 stream 的 kernel 太大，同时挤占同一个 SM 的资源。还有一种情况是 stream 之间没有足够独立的数据，导致它们互相争抢 L2 缓存。
排查思路：
- 用 Nsight Systems 看 stream 的时间轴，确认各 stream 是不是真的在并行。
- 用小 kernel 调大 stream 数测试，找到“并行收益阈值”。
- 检查 L2 缓存命中率，命中率下降严重说明 stream 间的数据复用率太低，这时候可以考虑单 stream 顺序执行。
8. 需要纠正的几个错误认知
8.1 线程数超过核心数就等于充分利用
这是最常见的误解。核心数是固定的，但 SM 里能驻留的线程数远大于物理核心数，靠的就是调度器在 warps 之间来回切换。线程多只能说明“池子大”，但池子里的线程如果都在等同一个资源（比如同一块显存），利用率照样上不去。真正应该关注的是“就绪的 warp 数”而不是“总线程数”。
8.2 高 occupancy 一定等价于高性能
前文已经反复提到了，这里再强调一次。高 occupancy 可以掩盖延迟，但代价是资源受限、寄存器溢出、甚至 cache 争抢加剧。低 occupancy 的内核如果每个线程的局部性特别好、指令级并行度高，完全可能比高 occupancy 版本更快。指标服务于瓶颈，而不是反过来。
8.3 GPU 调度和 CPU 调度可以走同一套思路
GPU 调度是硬件的、无操作系统介入的、极度高频的（纳秒级）；CPU 调度是软件/内核的、有优先级策略的、相对低频的（微秒以上）。你把 CPU 多线程的调优手段（比如锁、条件变量、线程池栈大小调优）搬进 CUDA kernel，基本全是反模式。GPU 上正确的“并发”方式是开足够多的独立线程、避免同步、减少资源占用。
9. 后续还能深挖的方向
这一篇说到底，是从“模型”层面把 NVIDIA 任务调度的骨架梳理了一遍。再往下挖，还有几块硬骨头值得单独写：
- Volta 之后的独立线程调度：从 Volta 开始，NVIDIA 引入了独立线程调度（Independent Thread Scheduling），warp 内不再是完全锁步，这让以前的“隐式同步假设”失效，也带来了一些新的同步语义（比如
__syncwarp的引入）。这背后有一套更精细的程序计数器（PC）管理机制，值得单独讲。 - 持久线程（Persistent Threads）与手工负载均衡：在一些大规模计算场景里，让线程数量等于 SM 容量而非数据量，然后靠循环自己拉取任务，能做到更精细的负载均衡，这算是对调度模型的“手动接管”。
- 用户态任务图（CUDA Graphs）与流序调度：面向低延迟推理场景，把 kernel 启动开销压到接近零的做法，调度模型会变成一张图，而不是一条链。
这些内容等后续有时间我再慢慢写。最后再分享一个实际的建议：手头有 NVIDIA GPU 的朋友，别光看这篇文章，拿 Nsight Compute 跑一个自己以前的 kernel，把 warp state 那一栏翻出来，对照着这一篇的概念去看，你会发现很多以前觉得“玄学”的性能问题，其实都是调度模型里写得很明白的事情。
