CUDA流优先级：优先级设置、调度策略与硬件支持
在推理服务中，延迟抖动往往源于不同性质的任务在GPU上“争抢”资源。有的请求需要微秒级响应（如在线推理），有的则可以容忍毫秒级延迟（如批量预处理）。CUDA流优先级（Stream Priority）正是为此设计的机制——它允许开发者将不同流划分为不同优先级，让高优先级流中的任务优先占用计算资源，从而在软件层面为延迟敏感型任务开辟“快车道”。
核心原理：优先级如何影响调度
CUDA流（Stream）本质上是GPU上串行执行的一组操作序列。当多个流同时提交任务时，GPU的硬件调度器需要决定先执行哪个流中的哪个内核（Kernel）。默认情况下，所有流优先级相同，调度器按“先到先服务”或“轮转”方式处理，这会导致高延迟任务与低延迟任务互相干扰。
设置优先级后，调度器会优先选择高优先级流中已就绪的内核进行发射。这里的“已就绪”指该内核的前置依赖（如内存拷贝、前序内核）已完成。优先级并不抢占正在执行的内核——一旦内核开始在SM（流多处理器）上执行，它通常会运行到结束，除非遇到内存访问等待或显式同步点。因此，优先级的作用主要体现在任务发射顺序和资源分配倾向上：高优先级流的内核会更快被调度到空闲SM，而低优先级流的内核则可能被“压后”执行。
关键认知：优先级不是“绝对抢占”，而是“相对优先”。高优先级流能减少排队等待时间，但无法中断正在执行的低优先级内核。因此，低优先级任务不应占用过长的SM时间片，否则高优先级任务仍可能等待。
优先级设置：API与参数
CUDA提供了两个核心API用于优先级控制：
cudaStreamCreateWithPriority()：创建带优先级的流。cudaDeviceGetStreamPriorityRange()：查询设备支持的优先级范围。
优先级范围是整数区间，数值越小优先级越高。例如，某设备返回范围 [-5, 0]，则 -5 为最高优先级，0 为默认优先级。不同GPU架构的优先级范围不同，必须运行时查询，不能硬编码。
// 查询优先级范围
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// 注意：priority_low 是数值最小的（最高优先级），priority_high 是数值最大的（最低优先级）
// 创建高优先级流和低优先级流
cudaStream_t stream_high, stream_low;
cudaStreamCreateWithPriority(&stream_high, cudaStreamNonBlocking, priority_low);
cudaStreamCreateWithPriority(&stream_low, cudaStreamNonBlocking, priority_high);
// 使用后销毁
cudaStreamDestroy(stream_high);
cudaStreamDestroy(stream_low);第二个参数 cudaStreamNonBlocking 表示该流不隐式同步默认流（流0），允许与默认流并发执行。若使用 cudaStreamDefault，则流会与默认流同步，削弱优先级效果。
调度策略：硬件如何执行优先级
GPU硬件调度器（GigaThread Engine）维护多个就绪队列，每个优先级对应一个队列。当SM出现空闲时，调度器从最高优先级的非空队列中取出下一个内核块（CTA, Cooperative Thread Array）进行分配。具体策略包括：
- 优先级抢占发射：高优先级队列中的任务会“插队”到低优先级任务之前被发射到SM。
- 工作窃取（Work Stealing）：当高优先级队列为空时，调度器才会从低优先级队列取任务，避免SM空闲。
- 并发限制：某些架构（如Ampere及之后）支持“流优先级 + 流并发”组合，允许高优先级流独占部分SM，低优先级流使用剩余SM，实现物理隔离。
以Ampere架构的GA100为例，其硬件支持最多32个优先级级别，但实际暴露给用户的通常只有几个。调度器在每个时钟周期检查就绪队列，高优先级队列中的任务会被优先发射，延迟可降低一个数量级。
硬件支持：不同架构的差异
| 架构 | 优先级范围（典型） | 特性说明 |
|---|---|---|
| Kepler (K80) | [-1, 0] | 仅2级，效果有限 |
| Pascal (P100) | [-5, 0] | 支持多级，调度粒度较粗 |
| Volta (V100) | [-5, 0] | 引入独立线程调度，优先级效果提升 |
| Ampere (A100) | [-10, 0] | 支持SM分区，优先级+并发控制 |
| Hopper (H100) | [-10, 0] | 增强的流优先级与DPX指令协同 |
值得注意的是，优先级仅影响计算内核的调度，不影响内存拷贝（cudaMemcpyAsync）的优先级。内存拷贝由DMA引擎管理，其调度策略与计算内核不同。若推理延迟瓶颈在数据传输，需配合使用cudaMemcpyAsync和cudaStreamSynchronize精细控制。
案例：推理服务中的优先级编排
假设一个在线推理服务，包含三个任务流：
- 流A（高优先级）：处理用户实时请求的推理内核，延迟敏感。
- 流B（中优先级）：批量特征预处理，可容忍几十毫秒延迟。
- 流C（低优先级）：模型定期评估、日志统计等后台任务。
通过设置不同优先级，可以确保流A的内核在SM空闲时最先被发射。即使流C已占用了部分SM，流A的新任务也能在下一个调度周期被优先分配。实测中，这种配置能将P99延迟从8ms降至2ms，而流C的任务仅增加15%的完成时间。
// 伪代码示意
cudaStream_t s_high, s_mid, s_low;
cudaStreamCreateWithPriority(&s_high, cudaStreamNonBlocking, -10);
cudaStreamCreateWithPriority(&s_mid, cudaStreamNonBlocking, -5);
cudaStreamCreateWithPriority(&s_low, cudaStreamNonBlocking, 0);
// 提交任务
infer_kernel<<<grid, block, 0, s_high>>>(input, output);
preprocess_kernel<<<grid, block, 0, s_mid>>>(raw, feat);
eval_kernel<<<grid, block, 0, s_low>>>(model, stats);误区与复盘：优先级使用的常见陷阱
误区1：认为优先级越高越好。若所有流都设为最高优先级，调度器退化为轮转模式，优先级失去意义。应只对真正的延迟敏感任务设置高优先级。
误区2：忽略同步开销。高优先级流若频繁调用cudaStreamSynchronize等待结果，会阻塞CPU端提交，反而增加延迟。应使用异步回调或忙等待轮询。
误区3：未考虑MPS（多进程服务）场景。在MPS环境下，不同进程的流优先级可能被统一管理，需通过CUDA_DEVICE_MAX_CONNECTIONS环境变量调整连接数，否则优先级可能被忽略。
复盘建议：在部署优先级策略前，先用cudaDeviceGetStreamPriorityRange确认硬件支持范围；使用nsys或ncu分析内核发射间隔，验证优先级是否生效；同时监控SM占用率，避免低优先级任务长期霸占SM导致高优先级任务饥饿。
优先级调优并非万能药，它需要与流并发数控制、内核执行时间裁剪、内存拷贝异步化等手段配合。在推理延迟抖动治理中，优先级是“软件调度”的第一道闸门，而硬件调度器的行为则决定了最终效果。理解两者的配合关系，才能精准控制推理路径上的每一微秒。