Roofline
并发与调度 影响 高 2026-08-07

多 stream 却不并发 —— 逻辑 stream 抢占物理硬件队列

现象
代码里开了很多 CUDA stream，但 Nsight 时间线上 kernel 首尾相接、几乎没有重叠，多流优化完全没有收益
根因
CUDA stream 是逻辑概念，真正下发要经过数量有限的硬件工作队列。逻辑 stream 多于可用队列时会被复用到同一条队列上，彼此产生假依赖而串行
工具
Nsight Systemsnvidia-smi
关注指标
sm__throughput.avg.pct_of_peak_sustained_elapsedgpu__time_duration.sum

这是多流优化里最容易踩、也最容易误判的一个坑：代码看起来完全正确，cudaStreamCreate 开了十几条流，任务之间也确实没有数据依赖，但一抓 profile 就发现 kernel 老老实实排成一条线。

很多人第一反应是「GPU 资源不够了，所以跑不满」。但这里的瓶颈往往不是算力，而是下发通路。

现象长什么样

在 Nsight Systems 的时间线上，把 CUDA HW 那一组行展开，典型表现是：

  • 明明有 Stream 14、Stream 15、Stream 16 等多行
  • 但同一时刻只有一个 kernel 在跑，块与块之间首尾相接
  • SM 利用率不高，GPU 也没有明显空转，就是「一个接一个」

如果 kernel 本身很小，还会看到 kernel 之间有稳定的小间隙 —— 那是排队和下发的开销。

为什么会这样

关键在于 逻辑 stream ≠ 物理队列。

应用侧创建的 stream 是一个软件抽象。驱动要把这些 stream 上的命令，通过若干条 硬件工作队列（hardware work queue / channel）送进 GPU。这些队列的数量是有限的， 默认值比大多数人以为的要小。

当逻辑 stream 的数量超过可用队列数时，多条逻辑 stream 会被映射到同一条硬件队列上。 同一条队列内部是严格 FIFO 的，于是两个逻辑上毫无关系的 kernel，也会因为共享队列而 被迫一前一后 —— 这就是所谓的假依赖（false dependency）。

逻辑 stream 多于硬件队列时被复用，队列内被迫串行，形成假依赖
逻辑 stream 多于硬件队列时被复用，队列内被迫串行，形成假依赖

控制这个数量的是环境变量 CUDA_DEVICE_MAX_CONNECTIONS：

项 值
默认值 8
可设上限 32
生效时机 进程启动时读取，运行中改无效
代价 每条队列占用一些设备资源，并非越大越好

也就是说，默认情况下你开到第 9 条 stream 就已经开始复用了。开 32 条 stream 却只有 8 条通路，并发度自然上不去。

怎么确认是这个原因

先抓一段 profile：

Bash
nsys profile \
  --trace=cuda,nvtx \
  --cuda-memory-usage=false \
  --output=stream_check \
  ./your_app

然后做一个对照实验 —— 这一步是判定的关键，比盯着时间线猜要可靠得多：

Bash
# 基线：默认 8 条通路
CUDA_DEVICE_MAX_CONNECTIONS=8  nsys profile -o q8  ./your_app

# 放开到 32 条
CUDA_DEVICE_MAX_CONNECTIONS=32 nsys profile -o q32 ./your_app

对比两次的总耗时和 kernel 重叠情况：

  • 重叠明显变多、总时间下降 → 基本可以确认瓶颈就在硬件队列复用上
  • 两次几乎一样 → 不是这个问题，去查下面「容易混淆的几种情况」

解法

按性价比排序：

  1. 调大 CUDA_DEVICE_MAX_CONNECTIONS。 最省事的一招，改环境变量即可。注意它必须在 进程启动前设置，且要配合实测 —— 调到 32 不一定最优。

  2. 减少 stream 数量，让它匹配通路数。 与其开 32 条流各自排队，不如开 8 条流把任务 均摊进去。多出来的逻辑流只会增加调度开销，不会带来额外并发。

  3. 合并小 kernel。 如果单个 kernel 只跑几微秒，并发带来的收益本身就有限， 下发开销反而占主导。这种情况优先考虑 kernel 融合或 CUDA Graph。

  4. 检查是不是真的需要并发。 如果 kernel 已经能把 SM 占满，串行执行并不吃亏 —— 此时多流优化的方向本来就是错的。

容易混淆的几种情况

时间线上「不重叠」有好几种原因，别一上来就归到队列复用：

  • 用了 legacy default stream。 默认流会和其它流隐式同步。要么用 cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking)，要么编译时加 --default-stream per-thread。
  • 中间夹了隐式同步操作。 cudaMalloc、cudaFree、同步版 cudaMemcpy、 以及很多 cudaDeviceSynchronize 都会打断流水。
  • 单个 kernel 已经吃满 SM。 资源不够，后来的 kernel 只能等，这属于正常。
  • stream 优先级设置。 cudaStreamCreateWithPriority 会让高优先级流插队。

判断顺序建议：先排掉默认流和隐式同步（这两个最常见且好查），再看 occupancy 是不是 已经打满，最后才做 CUDA_DEVICE_MAX_CONNECTIONS 的对照实验。

参考

  • CUDA C Programming Guide — Streams and Concurrency
  • Nsight Systems User Guide — CUDA HW rows 的解读
streamhardware queueCUDA_DEVICE_MAX_CONNECTIONS假依赖并发串行抢占
← 全部案例