# What a consumer RTX 5090 is missing next to a datacenter GPU
> The 5090 and the B200 are both Blackwell, but the consumer chip is missing whole capabilities. Here is what is gone, what it costs, and why you cannot just port data centre code.
Source: https://rfriedmann.de/blog/what-the-5090-lacks-vs-datacenter/
Published: 2026-05-26 · Track: log · Level: Advanced
NVIDIA's RTX 5090 and its data centre sibling, the B200, are both "Blackwell".
The name suggests a family resemblance, and on a spec sheet you will see the same
terms: FP4 tensor cores, fifth-generation this, Blackwell that. Spend a week
writing kernels for the consumer chip, though, and you find that it is missing
whole capabilities the data centre part takes for granted.
This is a tour of that gap, why it is real silicon and not a software lock, and
what it actually costs you when running large language models.
## The big one: the matrix engine cannot work in the background
Modern data centre GPUs have a feature that consumer Blackwell does not. On a B200,
the tensor cores (the units that do the giant matrix multiplies at the heart of
every model) can run **asynchronously**: you kick off a matmul, the GPU goes off
and does other useful work, and you come back later to collect the result. A
whole family of instructions exists to make this smooth, with names like
`tcgen05`, `TMEM`, and `wgmma`. They're how the fastest attention kernels
keep every part of the chip busy at once.
It helps to split this into two halves. The async-*copy* engine, TMA
(`cp.async.bulk.tensor`, the Tensor Memory Accelerator that has existed since
Hopper) plus plain `cp.async`, the 5090 has. The async-*compute* path, the part
that lets the tensor cores consume directly from that pipeline (`tcgen05`,
`TMEM`, and Hopper's `wgmma`), it does not. So the 5090 can stream data in the
background, but its matrix engine cannot run in the background.
The result is that the RTX 5090 has none of the async-compute instructions. As
imp's own kernel notes put it, on this chip "the MMA always blocks the issuing
warp": the matrix instruction stops and waits for its own result before anything
else can happen. The async copy can prefetch the next tile, but it cannot hide
the matrix multiply itself, which stays synchronous.
Datacenter Blackwell - B200 (sm_100)
tcgen05 / TMEM (async matrix engine)
wgmma (warp-group matrix multiply)
TMA (async copy)
FP4 tensor cores
Consumer Blackwell - RTX 5090 (sm_120a)
tcgen05 / TMEM
wgmma
TMA (async copy, cp.async.bulk.tensor)
FP4 tensor cores (mma.sync, blocking)
cp.async, HMMA m16n8k16
So the 5090 is not a weak chip. It has real FP4 tensor cores and the fast
asynchronous-copy plumbing to feed them. It is simply shaped differently: it can do
the heavy maths, it just cannot do it in the background while it gets on with
something else.
## The other gap: memory
Compute is only half the story. The other half is how fast you can feed it.
Memory bandwidth - the gap you can't tune away
[diagram omitted — see the page for the chart]
The 5090's 1.79 TB/s (32 GB GDDR7) is documented in imp's sm120.md. The B200 figure (~8 TB/s HBM3e) is NVIDIA's published spec, shown for scale.
Generating text is mostly a memory problem, not a maths problem. For a typical
decode step, imp's notes measure the ratio of memory traffic to actual
arithmetic at roughly **28 to 1**: the chip spends almost all its time moving
weights around and almost none of it calculating. "Compute is free", as the docs
bluntly put it. That has a reassuring side and an awkward one. The
reassuring side: for single-stream decoding, the 5090's 1.79 TB/s is enough to be
genuinely fast. The awkward side: when you do hit a job that is bound by raw
bandwidth or capacity, the data centre card has ~4.5x the memory bandwidth, and
~6x the capacity (192 vs 32 GB), and simply pulls away, with no clever code that
closes that.
There is a quieter difference too: the 5090 has a 98 MB L2 cache, and data centre
Blackwell pairs a large L2 of its own with far more HBM bandwidth and capacity on
top. (Neither part has a gaming-style L3.) Small on a slide, but it shapes which
optimisations are even worth trying.
## Why you cannot just port the data centre code
Put those two gaps together and you get the practical conclusion: the best kernel
for a B200 is not the best kernel for a 5090. They are different designs, not the
same design at different speeds.
B200: FlashAttention-4 style
Three hardware pipelines overlapping
Async matrix engine hides its own latency
Built on tcgen05 / TMEM / TMA
5090: what actually wins
FlashAttention-2, output kept in registers
cp.async double-buffering to hide loads
Softmax overlapped on a side unit
This is not a guess. Someone (well, Claude Code, under my direction) built the
full data centre-style approach for the 5090's main matrix multiply: the async
TMA loads, the warp-specialised producer/consumer design, the entire method. It
was correct and it ran **slower**: in a single internal measurement, 509 against
629 TFLOP/s for the simpler design.
The elaborate machinery starves the chip of parallelism to feed its one clever loader.
On this hardware, the textbook data centre technique is the wrong tool, and that
result is checked into the repo so nobody has to rediscover it.
The summary, side by side:
| | RTX 5090 (sm_120a) | B200 (sm_100) |
|---|---|---|
| Async matrix engine | no | yes |
| Warp-group GEMM (wgmma) | no | yes |
| Warp-specialized async-compute GEMM | no | yes |
| TMA async copy (cp.async.bulk.tensor) | yes | yes |
| FP4 tensor cores | yes (blocking) | yes (async) |
| Usable shared memory / SM | ~99 KB | more |
| L2 cache | 98 MB | large L2 (no gaming-style L3 either side) |
| Memory | 32 GB GDDR7, 1.79 TB/s | ~192 GB HBM3e, ~8 TB/s |
## So what is the 5090, then?
Not a baby B200. It is its own machine, and the honest framing is that you gain
some and lose some. You lose the async matrix engine, the bandwidth, the
capacity, and with them any hope of matching data centre throughput on the heavy
parallel jobs. You also lose NVLink (the 5090 talks over PCIe only, so multi-GPU
scale-out is far weaker) and most of the FP64 throughput (cut to roughly 1:64),
which rules it out for double-precision HPC work. In exchange you get a card
that, with the right kind of kernel,
decodes large models at speeds nobody else delivers on consumer hardware, for a
small fraction of the price and power.
The gap to a data centre card is real, and on the hardest kernels it is a ceiling
made of silicon, not code. But "differently shaped" is not the same as "slower at
everything", and the [whole game](/blog/optimizing-kernels-consumer-blackwell/)
is finding the things this chip is quietly very good at.