前言
我这台机器上装了两块 GPU，语音对话、口型同步数字人、LLM、图像生成、视频生成全都跑在上面。第三块卡快到了，所以我在规划哪个负载该放在哪块卡上。
过程中冒出来一个看着挺有道理的假设。
LLM 的解码是显存带宽受限的，扩散模型是算力受限的。既然吃的资源不一样，放在同一块卡上不就应该互不干扰吗？
NVIDIA 有个叫 MPS（Multi-Process Service）的功能，看上去正是为此而生。于是我试了。
先说结论：假设是错的。而且我判定它「错」的那个过程，本身也错了两次。 测了三遍，才拿到一个真正能解释现象的模型。下面是这三遍的完整记录。
第一幕：空间的假设
先测两块卡的 Roofline
没有基准就没法讨论，所以先测。RTX PRO 6000 Blackwell Max-Q（96GB）和 RTX PRO 4000 Blackwell（24GB）。
| 6000 Max-Q | 4000 | 比值 | |
|---|---|---|---|
| SM 数量 | 188 | 70 | 2.69x |
| 显存带宽 | 1469 GB/s | 553 GB/s | 2.66x |
| bf16 GEMM（持续） | 228 TFLOPS | 103 TFLOPS | 2.22x |
| 拐点（ridge point） | 155.5 FLOP/byte | 185.9 FLOP/byte | — |
| 4 秒内的降频 | -2.0% | -0.9% | — |
有意思的是，两块卡的显存频率（14001MHz）和 SM 频率（3090MHz）完全一致，差别只在位宽、SM 数量和功耗墙。带宽比（2.66x）大于算力比（2.22x），6000 的拐点也更低，说明它是相对「喂得更饱」的那块卡。
我本来怀疑 Max-Q 的 300W 限制会砍算力，但 4 秒的 GEMM 只掉了 2.0%。
假设的依据
拐点就是「以这个运算强度为界，带宽受限和算力受限互换」的位置。6000 是 155.5 FLOP/byte。
LLM 解码远在它下方。 吐一个 token，要把每一层的每个权重矩阵从头读一遍，各自乘上「当前这一个 token 的细向量」。权重读一次、用一次、扔掉，完全没有复用，所以运算强度只有 2B/b（b 是每参数字节数，B 是 batch）。NVFP4（0.5 byte/param）在 batch 1 下就是 4 FLOP/byte，是拐点的 1/39。
把实测带宽套到 27B 的 NVFP4 模型（约 13.5GB）上：
13.5 GB / 1469 GB/s = 9.2 ms/token → 理论上限约 109 tok/s
LLM 的速度不由算力决定，而由「能多快把权重从显存里抽出来」决定。 Tensor Core 基本在闲着。
扩散模型在另一侧。 2048² 的图像生成会把读进来的数据复用成千上万次，运算强度在几百到几千。带宽空得很，只有 Tensor Core 满负荷。
资源使用的位置分得很干净。所以应该能重叠——这就是当初的假设。
但首先：GPU 根本不会同时跑不同进程的 kernel
GPU 是「一进程 = 一个 CUDA context」，而且不同 context 的 kernel 不会在同一批 SM 上同时运行。GPU 是按 context 整体切换、轮流处理的。也就是分时。
没有 MPS —— 分时。任何时刻只有一边在跑
时间 →
LLM ████░░░░████░░░░████░░░░
图像生成 ░░░░████░░░░████░░░░████
SM 占用 低 高 低 高 低 高 ← LLM 的回合里 Tensor Core 在闲置
带宽占用 高 低 高 低 高 低 ← 图像的回合里带宽在闲置
MPS 改变了这一点。守护进程作为代理，把所有进程的 GPU 工作汇集到一个共享 context 里，于是不同进程的 kernel 就能真正同时跑了——理论上。
造一个测量仪
直接让真家伙对打，什么都说明不了。所以我做了两个代理，复现的不是真实 kernel，而是真实的「Roofline 上的位置」。
带宽侧（LLM 解码的代理）——巨大矩阵乘细向量。每次读 537MB，却只算 537 MFLOP。运算强度 1.0 FLOP/byte。
w = torch.randn(16384, 16384, dtype=torch.bfloat16, device="cuda") # 537 MB
x = torch.randn(16384, 1, dtype=torch.bfloat16, device="cuda") # batch 1
while ...:
torch.mm(w, x, out=y)
算力侧（扩散模型的代理）——方阵乘方阵。每次碰 1.6GB，算 8.8 TFLOP。运算强度 5461 FLOP/byte，在拐点上方 35 倍。
a = torch.randn(16384, 16384, dtype=torch.bfloat16, device="cuda")
while ...:
torch.mm(a, b, out=c)
如果只是朴素地同时启动，会混进「一边先跑完、另一边独占 GPU」的时间。所以给两个进程一个共同的 UNIX 时刻，各自预热后在屏障处会合，在完全相同的 5 秒窗口内计量。
而这个实验真正的仪器，是「归一化之后的和」。 把各条件的吞吐除以单独运行时的值，然后把两边相加。
和 ≈ 1.0 → 零和。一边的所得就是另一边的所失 = 没有重叠
和 > 1.0 → 真的重叠了。原本闲置的资源被填上了
把「到底是在重叠，还是在排队」压缩成一个数字。
结果：分时是零和，而 MPS 把带宽侧饿死了
单独运行的基准是带宽侧 1435 GB/s、算力侧 232 TFLOPS。
| 带宽侧 | 算力侧 | 和 | |
|---|---|---|---|
| 无 MPS | 44% | 51% | 96% |
| MPS / 不设上限 | 2% | 98% | 100% |
无 MPS 时和正好落在零和上。跟理论完全吻合，也确认了仪器没坏。
然后开了 MPS，带宽侧比分时还差 21 倍。
为什么：「不用算力」和「不需要 SM」是两回事
这就是假设的漏洞。读显存同样需要线程驻留在 SM 上去发出 load 指令。 「不用 Tensor Core」和「不需要 SM」完全是两码事。
算力侧的 GEMM 是单发 37.9ms 的怪物 kernel，会用 thread block 把 188 个 SM 全部塞满。而且在 block 粒度上没有抢占，一旦上去就跑到底。带宽侧的 GEMV 只能排在队尾，根本发射不出去。
分时的时候，至少还轮得到自己。开了 MPS 之后，换来的是「可以并排跑」，代价是「座位永远不空」。
加上 SM 上限就能回来，但那是一笔常年缴的税
MPS 有个 CUDA_MPS_ACTIVE_THREAD_PERCENTAGE，可以按客户端限制 SM 的份额。掐住算力侧，带宽侧就回来了。
| 算力侧的 cap | 带宽侧 | 算力侧 | 和 |
|---|---|---|---|
| 不限 | 2% | 98% | 100% |
| 50% | 61% | 43% | 104% |
| 25% | 85% | 22% | 107% |
| 10% | 93% | 7% | 100% |
再让它单独运行（没有对手）去测，就发现 cap 不是竞争时的仲裁，而是固定税。cap 25% 时，就算旁边一个人都没有，也只有 46%。
| cap | 带宽侧（单独） | 算力侧（单独） |
|---|---|---|
| 100% | 100% | 100% |
| 50% | 93% | 75% |
| 25% | 61% | 46% |
这里有个关键的不对称：带宽侧砍掉一半 SM 仍有 93%。 一半就足够把带宽喂饱。
两边都加 cap，SM 依然不会被分割
根据单独 cap 的实测，可以推出「如果真的被分割」应该是多少。带宽侧 50%（单独 93%）加算力侧 50%（单独 75%），和应该是 168%。
| 设置 | 若真分割 | 实测 |
|---|---|---|
| 带宽 50 / 算力 50 | 168% | 107% |
| 带宽 50 / 算力 25 | 139% | 102% |
| 带宽 75 / 算力 25 | ~143% | 109% |
| 带宽 25 / 算力 75 | ~149% | 99% |
差得太远。不管怎么调，和都贴在 99〜109% 之间，天花板始终在 107% 左右不动。cap 是「你最多能用多少 SM」的上限，而不是「给你留着」。它能拿走，但保证不了。
我在这里差点就下了结论
把一块 GPU 分给两个都能吃满的负载，存在硬件层面的天花板。分时是零和，MPS 无论怎么配都在 107% 左右，MIG 在这代卡上不支持。软件层面越不过去。
结论很漂亮，数字也齐。而且是错的。
第二幕：怀疑自己的仪器
让我别扭的是「cap 是上限而不是预留」这一句。那如果真的做出彼此不相交的 SM 集合呢？
CUDA 有个 %smid 特殊寄存器，能读到自己跑在哪个 SM 上。让落在自己范围之外的 block 立刻 return，就能不依赖任何库做出真正的分割。
__device__ __forceinline__ unsigned smid() {
unsigned r; asm volatile("mov.u32 %0, %%smid;" : "=r"(r)); return r;
}
__global__ void bw_pass(...) {
unsigned s = smid();
if (s < lo || s >= hi) return; // 不是我的 SM，直接下车
...
}
吸取第一幕的教训重做了：放弃常驻 block（占着直到截止时刻），改成像真实 kernel 那样反复 launch；缓冲区加到 16GiB，确保溢出 128MiB 的 L2；用原子工作队列，保证被 mask 掉的 block 不会导致数据漏读。
结果：分割没用。但和有 139%
带宽侧 单独（全部 SM） 1650 GB/s
带宽侧 单独（只用下半 SM） 1651 GB/s ← 一半 SM 就把带宽完全喂饱（100%）
算力侧 单独（只用上半 SM） 53.9 TFLOPS（56%）
→ 若分割完全生效，和应该是 157%
同时 / SM 共享 和 139%
同时 / SM 分割 和 138% ← 没有变化
显式分割 SM，和并没有提高。 带宽侧用一半 SM 就已经饱和，所以从一开始就没有可分割的余地。对这一对负载来说，硬件的贪心调度器本来就已经接近最优。
但真正扎眼的不是这个。和有 139%。 跟 MPS 的天花板 107%、分时的 96% 完全不是一个世界。
用同一套 kernel，只改进程结构
第一幕用的是 cuBLAS，第二幕用的是自写 kernel。kernel 不同就没法比较。于是固定 kernel、只改结构重测了一遍。
| 结构 | 带宽侧 | 算力侧 | 和 |
|---|---|---|---|
| 2 进程 / MPS OFF | 48% | 48% | 96% |
| 2 进程 / MPS ON | 66% | 66% | 132% |
| 1 进程 2 stream / SM 共享 | 91% | 48% | 139% |
| 1 进程 2 stream / SM 分割 | 95% | 44% | 138% |
MPS 把 96% 提到了 132%，它是正常起作用的。 第一幕那句「天花板只有 +7%」，作为一般命题是假的。
真凶是 kernel 的粒度
差别有两处。
单发的长度差了 7 倍。
- cuBLAS GEMM：37.9 ms/发
- 自写 kernel：5.2 ms/发
既然 block 粒度上没有抢占，就只能「等它空出来」。这 7 倍直接变成了插队的难易度。
而且我在占用率上留了空隙。 实测：
每个 SM 的最大线程数 : 1536
自写 kernel : 每线程 20 个寄存器，shared memory 0 B
→ 理论上能塞进 6 block/SM（占用率 100%）
→ 但我实际只 launch 了 4 block/SM
→ SM 的线程槽空着 33%
那空出来的 33%，正是对方的 block 挤进去的地方。cuBLAS 恰好相反，它用大块 shared memory tile 和大量寄存器，被调优成把占用率吃干净，结构上不给别人留缝。
换句话说，我的代理作为扩散模型的替身太轻了。 运算强度这个 Roofline 上的位置是对的，但占用 SM 的方式完全不同。只对齐 Roofline，做不出合格的代理。
第三幕：拿真家伙来对答案
代理不可信，那就测真的。不过这次我先把预测写下来再测。
HiDream 的重层（attention 和 linear）归根结底调的是 cuBLAS / cuBLASLt，那正是「又长又吃满占用率」的那类 kernel，所以应该偏悲观。但扩散的一步不是一发巨型 GEMM，而是 layernorm、逐元素运算、softmax、小投影等几百个 kernel 排成的队列，其间占用率会掉下来。所以应该落在 107%（cuBLAS 级）和 132%（细粒度级）之间，偏下端。
我把 ThinkCap-27B（NVFP4, vLLM）和 HiDream-O1-Image 从生产环境隔离出来，清空 GPU 后测量。图像保持生产同样的 2048²，只把步数从 50 降到 10——步数只是同一批 kernel 的重复次数，所以能在保持粒度的前提下缩短实验时间。
结果
基准：LLM 71.83 tok/s（TTFT 62.0ms）、图像 8.12 秒/张。
| 条件 | LLM | TTFT | 图像 | 和 | ||
|---|---|---|---|---|---|---|
| MPS OFF（现状） | 33.18 tok/s | 46.2% | 84.5 ms | 15.05 秒 | 54.0% | 100.1% |
| MPS ON | 18.53 | 25.8% | 126.1 ms | 9.13 | 88.9% | 114.7% |
| MPS ON + 图像 SM 25% | 52.80 | 73.5% | 67.2 ms | 36.86 | 22.0% | 95.5% |
114.7%——落在预测区间的下端。
顺带一提，基准的 71.83 tok/s 是第一幕里根据带宽推出的上限 109 tok/s 的 66%。考虑到 MTP 的开销和 KV cache 的读取，这个数字是合理的，说明 Roofline 模型本身依然成立。
三种配置，而且「最优」按指标看每次都不一样
- 总和最大的是 MPS ON（114.7%），但 LLM 掉到 25.8%
- 总和最小的是 MPS + cap（95.5%，比分时还低），但 LLM 最好，有 73.5%
- 分时在所有指标上都居中
决定对话体验的是 LLM 的 token 速度和 TTFT。只看这一条轴：
现在（分时） : 46.2% / TTFT 84.5ms
MPS + 图像 cap 25% : 73.5% / TTFT 67.2ms ← 几乎回到空载时的 62.0ms
图像生成期间的对话速度快了 1.6 倍，TTFT 几乎回到空载水平。代价是图像从 15.1 秒变成 36.9 秒（慢 2.4 倍）。图像是异步的、可以排队，所以这笔交易明显划算。
总和最小的那个配置，对产品来说反而是最优的。 吞吐总和在这个问题里是错误的货币。
结论：不是空间，是时间
测了三遍，最后剩下的是这个。
关于空间的讨论——吃哪种资源、放在哪个 SM——从头到尾一次都没起作用。
作为出发点的 Roofline 分析是讲得通的，而且是正确的——关于单独运行时的速度。而关于两者能否共存，它几乎什么都没预测出来。数据反复在说这件事。
- 带宽侧用一半 SM 就把带宽跑满 100%（也就是说，一开始就没有可分的空间）
- 显式做出不相交的 SM 集合，和只从 139% 变成 138%（也就是说，空间分离没有价值）
- 而另一边，只是把 kernel 的长度从 37.9ms 改成 5.2ms，就从 107% 动到了 132%
一直在起作用的，是「能插进去多少」。而它是两件事的乘积。
| 问题 | 由什么决定 | |
|---|---|---|
| 周转 | 座位什么时候空出来 | 单发 kernel 的长度（block 粒度没有抢占，只能等它跑完） |
| 空位 | 空出来时有没有地方进 | 邻居没用掉的占用率 |
两者都是能不能入场的问题，跟进去之后用哪个运算单元毫无关系。
有意思的是，唯一起作用的那个空间旋钮（SM cap），其实也是在时间轴上起作用的。 它并没有给谁专属的 SM，而是把贪心客户端的足迹缩小，从而提高座位空出来的频率。所以纯粹的空间分割没有价值，而作为上限的 cap 却有效。
落到实务上
在别人的代码里无计可施。 把 kernel 写短、在占用率上留余地，这些是自己写的 kernel 才能设计的东西。vLLM 和 cuBLAS 的内部动不了。让第三方技术栈共存这个问题，原理上只有物理隔离一条解。
然后就回到自己的产品。我们做的是实时语音角色扮演，对话回路是「语音检测 → 语音识别 → LLM → 语音合成 → 口型同步数字人」。其中数字人的帧生成带着硬性截止时间。
用这个模型看，数字人在共存下崩掉的原因就一目了然。数字人在每个 400ms 的截止周期里，有一条 encode → 9 次去噪 → decode 的依赖链。链上的每一环都要单独参加一次入场抽签。 这不是一次吞吐下降，而是要连续中签 11 次。所以哪怕平均只慢 1.3 倍，最差的那个 block 也会错过截止。
答案就是物理隔离。把实时路径（语音和数字人）隔离到专用卡上，LLM 和图像/视频生成放到另一块。快到的第三块卡，就是用来把这条线画完的。
MPS + cap 不上线。LLM 73.5% 很诱人，但守护进程会变成单点故障，vLLM 在 MPS 下长时间运行也没有实绩，更重要的是卡一到问题本身就消失了。把它记录为「不得不在一块卡上共存时的手牌」。
测三遍学到的东西
- 验证仪器只能证明它「没坏」。 第一幕里看到无 MPS 那一栏正好是 96%（零和），我就信了这个仪器。它确实没坏，只是不具代表性。这是两件不同的事
- 做代理时，只对齐 Roofline 上的位置是不够的。 不把占用方式也对齐，就不能拿来做共存实验
- 先写下预测再测量，命中时才能说「这个模型有预测力」。 第三幕押中了 114.7%，意味着下次再往上加负载时，动手之前就能在纸上判断。这才是三遍下来真正的收获
这个实验的边界
- 测的只是性格两极的一对负载。三个以上，或者 CPU 侧同步很重的负载，是否同样成立没有验证
- 视频生成（LTX-2）的 88 秒里有近一半是在从磁盘加载、GPU 空闲，所以按这个模型推断它反而是好邻居。没测
- 没有测「加了 cap 的图像单独运行」时的速度。 合成负载里 cap 是无条件税，如果真实情况相同，那么深夜无负载时图像也会一直慢。要上线的话这是必做的补测
- 窗口是 5〜45 秒，没有覆盖热平衡