抢占式调度：抢占点设计、内核级抢占、用户级抢占
抢占式调度，说白了就是「正在跑的任务，说停就得停」。
我早年做GPU驱动时，觉得这玩意儿挺简单——不就是中断一下嘛。后来被坑了几次才明白，抢占点的设计才是真正的艺术。停早了，数据乱掉；停晚了，抢占没意义。
抢占点的核心设计原则
抢占点，就是GPU执行过程中允许被打断的位置。不是所有地方都能停的。
核心原则：抢占点必须保证状态一致性。说白了，就是停下来的那一刻，所有关键数据都是「干净」的。
我在项目中遇到过一种情况：某个kernel正在写全局内存，突然被抢占了。结果呢？写了一半的数据被另一个任务读走了，整个计算结果全乱套。嗯，从那以后我设计抢占点时，都会加一条规则——内存操作必须完整。
抢占点通常分两类：
- 自然抢占点：指令流中天然存在的边界，比如线程块结束、同步点、内存屏障。
- 强制抢占点：人为插入的检查点，比如每N条指令检查一次抢占标志。
你想想看，自然抢占点开销小，但等待时间长。强制抢占点响应快，但会拖慢正常执行。怎么选？看场景。
内核级抢占
内核级抢占，是在GPU驱动层面实现的。它不关心上层应用在干嘛，只管「这个kernel能不能停」。
我习惯把内核级抢占分成两种模式：
| 模式 | 原理 | 延迟 | 适用场景 |
|---|---|---|---|
| 粗粒度抢占 | 等当前线程块执行完再停 | 高（毫秒级） | 批处理任务、离线计算 |
| 细粒度抢占 | 在指令级别打断 | 低（微秒级） | 实时渲染、交互式应用 |
粗粒度抢占实现简单，但响应慢。细粒度抢占响应快，但需要硬件支持——你得在指令流水线里插入检查点。
注意：细粒度抢占对硬件设计要求极高。我曾经在一个项目里尝试在ALU指令之间插入抢占点，结果流水线深度增加了30%，功耗直接炸了。后来我们改成了只在内存访问指令后检查抢占标志，才平衡了性能和响应速度。
内核级抢占的典型流程：
// 伪代码：内核级抢占处理
void kernel_preempt_handler(uint32_t kernel_id) {
// 1. 保存当前kernel的上下文
save_kernel_context(kernel_id);
// 2. 标记kernel为暂停状态
set_kernel_state(kernel_id, KERNEL_PAUSED);
// 3. 调度下一个kernel
uint32_t next_kernel = schedule_next();
// 4. 恢复下一个kernel的上下文
restore_kernel_context(next_kernel);
// 5. 开始执行
launch_kernel(next_kernel);
}这段代码看着简单，但实际实现时坑很多。比如保存上下文这一步——GPU的寄存器文件动辄几百KB，全保存？太慢。只保存关键寄存器？你得知道哪些是「关键」的。我建议的做法是：只保存那些在抢占点之后还会被用到的寄存器。这需要编译器配合，在编译时标注寄存器的生命周期。
用户级抢占
用户级抢占，是让应用程序自己决定什么时候可以被抢。听起来有点反直觉——都抢占了还让用户决定？但实际场景中，这很有用。
举个例子：一个AI训练任务，正在执行一个很长的卷积kernel。如果系统强行抢占，这个kernel的中间结果可能全丢了。但如果让用户在kernel里主动插入yield点，那就可以优雅地保存中间结果，再让出GPU。
我的经验：用户级抢占最适合那些「知道自己什么时候安全」的应用。比如深度学习框架，可以在每个layer执行完后检查一下抢占标志。我曾经帮一个团队优化过，他们在每个batch的forward/backward之间插入yield点，抢占延迟从5ms降到了200μs。
用户级抢占的API设计通常长这样：
// 用户代码中主动让出GPU
__global__ void my_kernel() {
for (int i = 0; i < N; i++) {
// 做一些计算
compute_partial(i);
// 检查是否需要让出
if (should_yield()) {
// 保存当前进度
save_progress(i);
// 让出GPU
cudaYield();
// 恢复后继续
restore_progress();
}
}
}这里有个关键问题：should_yield() 怎么实现？我见过两种做法：
- 轮询标志位：在全局内存里放一个标志，驱动想抢占时就置位。用户代码定期检查。简单，但每次检查都要访存，开销不小。
- 硬件中断：驱动发一个中断信号，用户代码在特定位置处理。响应快，但需要硬件支持，而且中断处理本身也有开销。
我个人更倾向于轮询标志位，因为它的行为可预测。中断方式虽然响应快，但处理时机不确定，容易引入微妙的bug。
抢占点设计的实战经验
说了这么多理论，来点实际的。我总结了几条抢占点设计的「避坑指南」：
避坑指南：
- 不要在共享内存操作中间设抢占点——共享内存是线程块内共享的，抢占后恢复时，其他线程可能已经把数据改了。
- 原子操作前后不要设抢占点——原子操作本身是不可分割的，你在中间插一脚，原子性就没了。
- 抢占点要避开屏障同步——__syncthreads() 前后设抢占点，容易导致死锁。我曾经踩过这个坑，debug了两天才找到原因。
还有一个容易被忽略的点：抢占点的数量要适中。太多，正常执行被拖慢；太少，抢占响应太慢。我一般建议每100-1000条指令设一个检查点，具体看你的延迟要求。
内核级 vs 用户级：怎么选？
这个问题没有标准答案。我个人的经验是：
- 如果你做的是通用GPU调度系统，用内核级抢占。它不依赖应用配合，适用范围广。
- 如果你做的是特定领域的加速库（比如AI推理、渲染），用用户级抢占。应用开发者更了解自己的安全点，可以做到更优雅的抢占。
- 如果两者都想要？那就混合使用。内核级做兜底，用户级做优化。我在一个项目里就是这么干的——用户级抢占处理90%的情况，内核级抢占处理剩下的10%。效果还不错。
一个小技巧：用户级抢占的yield点，可以设计成「可选的」。也就是说，应用可以不处理yield，这时候由内核级接管。这样既给了应用灵活性，又保证了系统的可控性。
抢占点设计的可视化
下面这张图展示了抢占点在GPU执行流中的位置：
从图上可以看得很清楚：内核级抢占点在线程块边界或指令之间，用户级抢占点则由应用自己控制。两种方式各有优劣，关键看你的场景需要多快的响应、多高的效率。
好了，关于抢占点设计、内核级抢占和用户级抢占，我就讲这么多。记住一点：没有完美的抢占策略，只有适合你场景的方案。多试、多调、多踩坑，慢慢就有感觉了。
想要更多资料，联系可可