"Our GPUs are at 100% utilization — we need more GPUs."
That sentence has justified more hardware spend than any other in this industry, and it is routinely said about fleets doing a fraction of the work they could. The problem is that three different meters all sound like "how hard is the hardware working," and the two you can see easily saturate while the one that matters starves.
The three meters
1. Utilization — is the chip claimed? nvidia-smi "GPU utilization" is the
fraction of time at least one kernel is resident. Any kernel. One SM doing inefficient
work all day reads 100%. It answers "is anyone in the building?"
2. Occupancy — is the scheduler packed? Of the warp slots each SM can hold, how many hold resident warps. It measures how well threads are packed onto the scheduler — not what they do. Warps that all stall on memory, or all grind slow arithmetic, still count. Every seat filled; nothing says anyone is eating.
3. MFU / MBU — is the silicon working? Achieved FLOPs against peak FLOPs (compute-bound work), achieved bytes/s against peak bandwidth (memory-bound work). The only meters that measure work done.
The arithmetic of "dark"
An H100's arithmetic lives overwhelmingly in one place: CUDA cores deliver ~67 TFLOPS
of FP32; the tensor cores deliver ~989 TFLOPS of dense BF16 — ~15× more, and most
of the die you paid for. Now run a competent plain-CUDA matmul: FFMA opcodes on CUDA
cores, tuned, full occupancy, nvidia-smi pinned at 100%. Its ceiling is 67 of 989
TFLOPS ≈ 7%. The other ~93% of the chip is dark — powered, idle, unreachable —
because tensor cores are only addressable through specific instructions
(mma/wgmma/tcgen05) that this kernel never issues. They aren't slow; they are
unaddressed. Claimed: 100%. Packed: 100%. Working: 7%.
This is not exotic. It is what "we use the GPU" quietly means whenever kernels predate the current tensor-core generation, whenever a framework falls back to unfused ops, whenever "engine A is 10× faster than engine B on identical silicon" — which usually decodes to A's kernels speak the tensor-core instruction and B's don't. A software fact wearing a hardware costume.
The full ladder — silicon to solutions
The three meters are just the bottom rungs. Keep climbing and each rung exposes waste the rung below certified as healthy. Every rung here has been measured somewhere in this practice's work:
| Rung | Question | The lie it catches |
|---|---|---|
| Utilization | claimed? | one busy SM reads 100% |
| Occupancy | packed? | busy-waiting warps read 100% |
| MFU / MBU | silicon working? | PENNANT: good occupancy, 97–99% of cycles stalled on memory — on both vendors |
| Phase fit | right silicon working? | decode barely touches tensor cores — prefill is MFU's problem, decode is MBU's |
| Goodput | users served within SLA? | our serving eval: dashboards green, zero errors, goodput −92% from one config default |
| Solution rate | problems actually solved? | an agent fleet's wall-clock is mostly waiting on tools — perfect serving of an idle loop |
| $/solved-task | more solutions, faster, cheaper? | our agent bake-off's honest unit: ~$0.004/solved on a self-hosted 32B — the number a buyer feels |
Two field results anchor the ladder's span. At the bottom: PENNANT's stall study found an H100 and an MI300X leaving 97–99.7% of compute idle on memory with identical signatures — the waste was the algorithm's, invisible to utilization and occupancy alike. In the middle: our inference evaluation watched raw throughput hold flat and green while goodput collapsed 20.8 → 1.8 tok/s the moment a second concurrent user existed — then recovered 15× from a single serving flag. No hardware changed hands in either story.
The real metric: solutions per dollar per hour
Here is the reorientation the ladder forces. Nobody buys FLOPs. Nobody even buys tokens. They buy solved problems — a bug fixed, a simulation run, a question answered within the time it was worth answering. So the efficiency question is not "is the hardware busy?" but:
Solutions ÷ (dollars × hours) — maximized. More solutions, faster, cheaper.
Every rung of the ladder is a multiplicative factor in that number, and the cheap wins are almost never at the bottom:
- Issuing the right opcodes (using the silicon you own): up to ~15× on matmul-heavy work, free with a mature engine.
- Matching phase to substrate (prefill/compute, decode/bandwidth): the disaggregation dividend.
- Fixing serving config before buying hardware: our measured 15× goodput from one flag — the cheapest capacity in the industry.
- Making the verifier the router: escalate only failures — all-strong quality at near-cheap cost (the V-cycle result).
- Attacking the tool-wait: for agents, faster tests beat faster models.
Notice what buying more GPUs does to this list: it addresses none of it. Scaling out multiplies whatever efficiency you already have — including the waste. A fleet at 7% working-efficiency scaled 2× is twice the money for the same darkness. The order of operations is: light the silicon, fit the phase, fix the config, route by the verifier, then scale.
Takeaway
Claimed is a lobby metric. Packed is a seating chart. Working starts at MFU/MBU and only becomes real at goodput and $/solved-task. If your dashboards top out at utilization, you are optimizing the lobby — and the honest question to put to any fleet, before any purchase order, is the one this practice keeps measuring: how many solutions per dollar per hour, and which rung of the ladder is eating them?
Related: The general absorbs the specific · Do agents run on CPU or GPU? · Measuring training goodput · A V-cycle of models.