Introduction

I have two GPUs in my workstation, running voice conversation, a lip-sync avatar, an LLM, image generation, and video generation all at once. A third card is arriving, so I was working out which workload belongs on which card.

Along the way, a plausible-sounding hypothesis came up.

LLM decoding is memory-bandwidth bound. Diffusion models are compute bound. They consume different resources — so shouldn't they coexist on the same card without fighting?

NVIDIA has a feature called MPS (Multi-Process Service) that looks like it exists precisely to enable this. So I tried it.

The hypothesis was wrong. And the process by which I concluded it was wrong was itself wrong — twice. It took three rounds of measurement before I had something that actually explained the behavior. This is all three rounds.

---

Act 1: The spatial hypothesis

Measuring the roofline of both cards

I had no baseline, so I measured. RTX PRO 6000 Blackwell Max-Q (96GB) and RTX PRO 4000 Blackwell (24GB).

	6000 Max-Q	4000	Ratio
SM count	188	70	2.69x
Memory bandwidth	1469 GB/s	553 GB/s	2.66x
bf16 GEMM (sustained)	228 TFLOPS	103 TFLOPS	2.22x
Ridge point	155.5 FLOP/byte	185.9 FLOP/byte	—
Throttle over 4s	-2.0%	-0.9%	—

Both cards run identical clocks — 14001 MHz memory, 3090 MHz SM. The differences come down to bus width, SM count, and power budget. The bandwidth ratio (2.66x) exceeds the compute ratio (2.22x), and the 6000 has the lower ridge point, meaning it is the relatively better-fed card on the memory side.

I suspected the Max-Q's 300W limit would eat into compute, but a 4-second GEMM only throttled 2.0%.

Why the hypothesis looked good

The ridge point is where a workload flips from bandwidth-bound to compute-bound. On the 6000, that's 155.5 FLOP/byte.

LLM decoding sits far below it. Producing one token means reading every weight matrix in every layer and multiplying each by a single thin activation vector. Every weight is read once, used once, discarded. With no reuse, arithmetic intensity is 2B/b where b is bytes per parameter and B is batch size. At NVFP4 (0.5 byte/param) and batch 1, that's 4 FLOP/byte — 1/39th of the ridge point.

Applying the measured bandwidth to a 27B NVFP4 model (~13.5GB):

13.5 GB / 1469 GB/s = 9.2 ms/token  →  theoretical ceiling ≈ 109 tok/s

An LLM's speed is set not by compute but by how fast you can pull weights out of VRAM. The tensor cores sit mostly idle.

Diffusion is on the opposite side. Generating a 2048² image reuses loaded data thousands of times. Arithmetic intensity runs in the hundreds to thousands. Bandwidth is wide open while tensor cores saturate.

The resources cleanly divide. So they should overlap — that was the hypothesis.

But first: a GPU doesn't run kernels from different processes simultaneously

A GPU works on one CUDA context per process, and kernels from different contexts do not run concurrently on the same SMs. The GPU switches contexts and processes them in turn. Time-slicing.

Without MPS — time-slicing. Only one side runs at any moment
time →
LLM        ████░░░░████░░░░████░░░░
Image gen  ░░░░████░░░░████░░░░████
SM usage    lo  hi  lo  hi  lo  hi   ← tensor cores idle during LLM's turn
BW usage    hi  lo  hi  lo  hi  lo   ← bandwidth idle during image's turn

MPS changes this. A daemon becomes a proxy that funnels every process's GPU work into a single shared context, so kernels from different processes can genuinely run at the same time. In theory.

Building an instrument

Pitting the real things against each other immediately tells you nothing about why. So I built two proxies that reproduce not the real kernels but the real positions on the roofline.

Bandwidth side (LLM decode proxy) — huge matrix times a thin vector. Each call reads 537MB and computes only 537 MFLOP. Arithmetic intensity 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)

Compute side (diffusion proxy) — square times square. Each call touches 1.6GB and computes 8.8 TFLOP. Arithmetic intensity 5461 FLOP/byte, 35x above the ridge point.

a = torch.randn(16384, 16384, dtype=torch.bfloat16, device="cuda")
while ...:
    torch.mm(a, b, out=c)

Naively launching both mixes in "time when one finished and the other had the GPU to itself," so I hand both processes a shared UNIX timestamp, warm up, meet at a barrier, and measure over exactly the same 5-second window.

And the instrument itself is the normalized sum. Divide each condition's throughput by its solo value, then add the two.

sum ≈ 1.0  →  zero-sum. One side's gain is the other's loss = no overlap
sum >  1.0  →  genuine overlap. Idle resources got filled

It collapses "are they overlapping, or queueing?" into a single number.

Result: time-slicing is zero-sum, and MPS starves the bandwidth side

Solo baselines: 1435 GB/s and 232 TFLOPS.

	Bandwidth side	Compute side	Sum
No MPS	44%	51%	96%
MPS, no cap	2%	98%	100%

Without MPS the sum lands exactly on zero-sum. Textbook, which confirmed the instrument wasn't broken.

Then MPS made the bandwidth side 21x worse than time-slicing.

Why: "doesn't use compute" and "doesn't need SMs" are different things

This was the hole in the hypothesis. Reading memory still requires threads resident on an SM to issue the load instructions. Not using tensor cores is not the same as not needing SMs.

The compute-side GEMM is a monster kernel that runs 37.9ms per launch and fills all 188 SMs with thread blocks. And since there is no preemption at block granularity, a block that lands runs to completion. The bandwidth side's blocks just queue behind it and never get dispatched.

Under time-slicing, at least your turn came around. Under MPS you get to "run alongside" — but no seat ever frees up.

An SM cap brings it back, at the price of a permanent tax

MPS has a knob, CUDA_MPS_ACTIVE_THREAD_PERCENTAGE, that limits each client's share of SMs. Squeezing the compute side brings the bandwidth side back.

Compute-side cap	Bandwidth side	Compute side	Sum
none	2%	98%	100%
50%	61%	43%	104%
25%	85%	22%	107%
10%	93%	7%	100%

Measuring solo, with no competitor, showed the cap is not arbitration during contention but a flat tax. At cap 25% you get 46% even when nothing else is running.

cap	Bandwidth side (solo)	Compute side (solo)
100%	100%	100%
50%	93%	75%
25%	61%	46%

Note the asymmetry: the bandwidth side keeps 93% on half the SMs. Half is enough to saturate memory.

Capping both sides still doesn't partition the SMs

From the solo-cap numbers you can predict what perfect partitioning would give. Bandwidth side at 50% (93% solo) plus compute side at 50% (75% solo) should sum to 168%.

Setting	If partitioned	Measured
BW 50 / CMP 50	168%	107%
BW 50 / CMP 25	139%	102%
BW 75 / CMP 25	~143%	109%
BW 25 / CMP 75	~149%	99%

Nowhere close. However you set it, the sum pins between 99% and 109% — the ceiling won't move off roughly 107%. A cap is an upper bound on how many SMs you may use, not a reservation. It can take away; it cannot guarantee.

And here I almost published a conclusion

Sharing one GPU between two saturating workloads has a hardware ceiling. Time-slicing is zero-sum, MPS tops out around 107% however you configure it, MIG isn't supported on this card generation. Software cannot get past it.

It was clean. The numbers lined up. And it was wrong.

---

Act 2: Doubting the instrument

What nagged at me was "a cap is a bound, not a reservation." So what happens if you build a genuinely disjoint set of SMs?

CUDA exposes %smid, a special register telling a thread which SM it's running on. Have blocks that land outside your assigned range return immediately, and you get real partitioning without any library support.

__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;    // not my SM — step off
  ...
}

Learning from Act 1's mistakes, I rebuilt it: no more persistent blocks squatting until a deadline (real kernels finish and relaunch), a 16GiB buffer to overflow the 128MiB L2 for certain, and an atomic work queue so masked-out blocks don't cause data to go unread.

Result: partitioning does nothing. But the sum was 139%

BW side, solo (all SMs)         1650 GB/s
BW side, solo (lower half only) 1651 GB/s   ← half the SMs fully saturate bandwidth (100%)
CMP side, solo (upper half)     53.9 TFLOPS (56%)
  → if partitioning worked perfectly, the sum should be 157%

Concurrent / shared SMs         sum 139%
Concurrent / partitioned SMs    sum 138%   ← no change

Explicitly partitioning the SMs did not raise the sum. The bandwidth side already saturates memory with half the SMs, so there was never any room for partitioning to help. For this pair, the hardware's greedy scheduler was already near-optimal.

But that wasn't what caught my eye. The sum is 139% — nothing like MPS's 107% ceiling, nothing like time-slicing's 96%.

Same kernels, different process topology

Act 1 used cuBLAS; Act 2 used custom kernels. Different kernels can't be compared. So I re-measured holding the kernels fixed and varying only the topology.

Configuration	Bandwidth side	Compute side	Sum
2 processes / MPS OFF	48%	48%	96%
2 processes / MPS ON	66%	66%	132%
1 process, 2 streams / shared SMs	91%	48%	139%
1 process, 2 streams / partitioned SMs	95%	44%	138%

MPS took it from 96% to 132%. It works fine. Act 1's "the ceiling is +7%" was false as a general claim.

The culprit was kernel granularity

Two differences.

Launch duration differed by 7x.

• cuBLAS GEMM: 37.9 ms per launch
• My custom kernel: 5.2 ms per launch

With no preemption at block granularity, waiting for a seat to free is the only option. That 7x translates directly into how easy it is to get a word in.

And I had left occupancy headroom. Measured:

Max threads per SM         : 1536
My custom kernel           : 20 registers/thread, 0 B shared memory
  → could pack 6 blocks/SM (100% occupancy)
  → but I was only launching 4 blocks/SM
  → 33% of the SM's thread slots sat empty

That empty 33% was exactly where the other client's blocks landed. cuBLAS is the opposite: tuned with large shared-memory tiles and heavy register use to consume all available occupancy, leaving no room for anyone else.

In other words, my proxy was too lightweight to stand in for a diffusion model. Its position on the roofline was right, but the way it occupied SMs was completely different. Matching the roofline alone does not make a valid proxy.

---

Act 3: Checking against the real thing

If the proxy can't be trusted, measure the real thing. This time I wrote the prediction down first.

HiDream's heavy layers (attention and linear) ultimately call cuBLAS / cuBLASLt, which are exactly the long, occupancy-consuming kind of kernel. So it should land pessimistic. But a diffusion step isn't one giant GEMM — it's hundreds of kernels in sequence: layernorms, elementwise ops, softmax, small projections, with occupancy dipping between them. So it should land between 107% (cuBLAS-like) and 132% (fine-grained), toward the lower end.

I isolated ThinkCap-27B (NVFP4, vLLM) and HiDream-O1-Image from production and emptied the GPU. Images stayed at the production 2048² with only the step count cut from 50 to 10 — steps are just repetitions of the same kernels, so this shortens the experiment while preserving granularity.

Results

Baselines: LLM 71.83 tok/s (TTFT 62.0ms), image 8.12 s/image.

Condition	LLM		TTFT	Image		Sum
MPS OFF (status quo)	33.18 tok/s	46.2%	84.5 ms	15.05 s	54.0%	100.1%
MPS ON	18.53	25.8%	126.1 ms	9.13	88.9%	114.7%
MPS ON + image SM 25%	52.80	73.5%	67.2 ms	36.86	22.0%	95.5%

114.7% — the lower end of the predicted range.

For what it's worth, the 71.83 tok/s baseline is 66% of the 109 tok/s ceiling predicted from bandwidth back in Act 1. Given MTP overhead and KV cache reads, that's a reasonable figure — the roofline model itself is still alive.

Three configurations, and "best" differs by metric in every case

• Highest sum is MPS ON (114.7%) — but the LLM falls to 25.8%
• Lowest sum is MPS + cap (95.5%, below time-slicing) — but the LLM is best at 73.5%
• Time-slicing sits in the middle on everything

What determines conversational experience is LLM token rate and TTFT. On that axis alone:

Today (time-slicing)     : 46.2% / TTFT 84.5ms
MPS + image cap 25%      : 73.5% / TTFT 67.2ms   ← nearly back to the unloaded 62.0ms

Conversation runs 1.6x faster during image generation, and TTFT returns almost to idle. The cost is images going from 15.1s to 36.9s (2.4x slower). Images are async and queued, so the trade is clearly favorable.

The configuration with the lowest aggregate is the best one for the product. Aggregate throughput was the wrong currency for this problem.

---

Conclusion: not space, but time

Three rounds of measurement, and here's what survived.

The spatial argument — which resource you consume, which SM you sit on — never mattered once.

The roofline analysis I started from was sound, and it is correct: about how fast a workload runs alone. About whether two workloads can share, it predicted almost nothing. The data said so repeatedly.

• The bandwidth side saturated 100% of memory bandwidth on half the SMs (so there was never spatial room to divide)
• Explicitly disjoint SM sets moved the sum from 139% to 138% (so spatial separation has no value)
• Meanwhile, changing kernel duration from 37.9ms to 5.2ms moved it from 107% to 132%

What mattered throughout was how much you can interrupt. And that is a product of two things.

	Question	Determined by
Turnover	When does a seat free up?	Duration of one kernel launch (no block-level preemption, so you wait it out)
Vacancy	Is there room when one does?	Occupancy your neighbor left unused

Both are about admission — and have nothing to do with which execution units you use once you're in.

The interesting part: the one spatial knob that worked (the SM cap) was also working in the time domain. It doesn't give anyone dedicated SMs; it shrinks the greedy client's footprint so that seats free up more often. That's why pure spatial partitioning is worthless while a cap as an upper bound works.

What this means in practice

There is nothing you can do inside someone else's code. Shortening kernels and leaving occupancy headroom are things you can design when you wrote the kernel. You cannot touch the insides of vLLM and cuBLAS. Making third-party stacks coexist is, in principle, solvable only by physical separation.

Which brings it back to my own product. We do real-time voice roleplay, and the conversation loop is speech detection → ASR → LLM → TTS → lip-sync avatar. Avatar frame generation carries a hard deadline.

Through this model, it's clear why the avatar breaks under co-tenancy. Every 400ms deadline, the avatar runs a dependency chain: encode → 9 denoise steps → decode. Each link in that chain enters the admission lottery separately. It's not one throughput hit — it has to win 11 lotteries back to back. So even at an average of 1.3x slower, the worst block misses its deadline.

The answer is physical separation. Isolate the real-time path (voice and avatar) onto a dedicated card, and put the LLM and image/video generation elsewhere. The third card exists to finish drawing that line.

I'm not deploying MPS + cap. An LLM at 73.5% is tempting, but the daemon becomes a single point of failure, there's no track record for running vLLM under MPS for long stretches, and above all, the problem disappears when the card arrives. I'm filing it as the move for when co-tenancy on one card becomes unavoidable.

What three rounds taught me

• Validating an instrument only proves it isn't broken. In Act 1, seeing the no-MPS case land exactly on 96% (zero-sum) made me trust the instrument. It genuinely wasn't broken. It just wasn't representative. Those are two different properties
• When building a proxy, matching its position on the roofline is not enough. You have to match how it occupies the machine, or it's useless for co-tenancy experiments
• Writing the prediction before measuring lets you claim the model has predictive power when it lands. Hitting 114.7% in Act 3 means that next time I add a workload, I can reason about it on paper before running anything. That was the real payoff of three rounds

Limits of this experiment

• I only measured pairs of two workloads with opposite characteristics. Whether three-way mixes, or workloads with heavy CPU-side synchronization, behave the same way is untested
• Video generation (LTX-2) spends nearly half of its 88 seconds loading from disk with the GPU idle, so this model predicts it is actually a good neighbor. Not measured
• I never measured a capped image workload running solo. In the synthetic case the cap was an unconditional tax, so if that holds for the real thing, images stay slow even at 3am with no load. Required follow-up before any deployment
• Windows were 5 to 45 seconds, so thermal steady state is not covered