I’ve been going through blogs and papers regarding B200’s design and its tradeoffs. Here I will summarize what I’ve learned so far. B200 is NVIDIA’s datacenter Blackwell GPU: 148 SMs across two dies connected by NV-HBI, with 183 GB HBM3e at 8 TB/s and a 1000W TDP.

Blackwell Arch

Quick comparison with the previous Hopper arch. B200 unifies the fp32 and int32 cores while introducing TMEM. It also introduces new 4bit Tensor Cores for int4/fp4.

Component Hopper (H100 SXM5) Blackwell (B200)
SMs 132 148
GPCs 8 8 (4 per die)
FP32/INT32 128 FP32 + 64 INT32 (separate) 128 unified (FP32 or INT32 per cycle)
FP64 64 dedicated ~64 dedicated per SM
Tensor Cores 4 (4th-gen) 4 (5th-gen)
TMEM — 256 KB
L1/SMEM per SM 256 KB unified 256 KB unified
L2 cache 50 MB ~126 MB
HBM 80 GB @ 3.35 TB/s 192 GB raw @ 8 TB/s
Native FP4/FP6 No (FP8 only) Yes

On Hopper the 64 int32 cores can also execute fp32 instructions, giving Hopper a peak of 192 FP32 FMA/clock per SM. By unifying to 128, Blackwell actually reduces scalar FP32 throughput per SM by 33%. Blackwell is explicitly designed to sacrifice scalar compute density to make room for TMEM and larger Tensor Cores.

Register Pressure and the TMEM Migration

There has been a slow migration of the accumulator from registers to where the matrix multiply data actually lives, closer to Tensor Cores.

Architecture Matrix A Matrix B Accumulator D Key Instruction
Volta (2017) Register File Register File Register File mma.sync
Ampere (2020) Register File / SMEM SMEM Register File mma.sync + cp.async
Hopper (2022) SMEM (via TMA) SMEM (via TMA) RMEM (special registers) wgmma.mma_async
Blackwell (2024) TMEM / SMEM SMEM TMEM tcgen05.mma

TMEM (Tensor Memory) is a new on-SM scratchpad that’s private to the Tensor Cores. TMEM doesn’t support the typical load/store operations per thread. TMA moves data from global memory into shared memory, then tcgen05.mma reads operands from SMEM and writes the accumulator to TMEM. The only way to read it back into registers is tcgen05.ld.

TMEM

Each Tensor Core generation moved from more threads participating to fewer, making the MMA progressively less intrusive on the scalar pipeline. On Volta and Ampere mma.sync is required for all 32 warp threads to participate which stalls the entire warp until the MMA is completed (no overlap). Hopper’s wgmma.mma_async cuts this to a warpgroup (128 threads) and makes it asynchronous so the warpgroup could issue an MMA and then proceed to other work while it ran in the background.

Blackwell’s tcgen05.mma takes this to its logical endpoint. A single thread issues the MMA on behalf of the entire CTA. The other 127 threads never touch the MMA path, they’re free to issue TMA loads, manage barriers or compute addresses concurrently. This is “full decoupling”. The Tensor Cores are now independent coprocessors driven by one thread with its accumulator in TMEM and completely separate from the scalar register file.

Register Pressure

I compiled two CUTLASS tutorial GEMM kernels on the B200:

  • Ampere path (sgemm_sm80.cu): uses mma.sync with a 128×64 tile, FP16 inputs and an FP16 accumulator held in registers. Compiled with nvcc -arch=sm_100 (backward-compatible on B200).
  • Blackwell path (01_mma_sm100.cu): uses tcgen05.mma with a 128×256 tile, FP16 inputs and an FP32 accumulator held in TMEM. Compiled with nvcc -arch=sm_100a (the a variant enables tcgen05/TMEM).

Both were compiled with -ptxas-options=-v which makes the PTX assembler print register allocation for every function.

Register Comparison

I observed that the Blackwell kernel hits the 255 register ceiling and spills to local memory. The question is, if TMEM removes the accumulator from the register file then why does it need more registers and not fewer?

Since ncu isn’t available on a RunPod cloud instance, I used two approaches to dissect this - kernel deletion and SASS disassembly. First I removed the epilogue from the original kernel that included tcgen05.ld, type conversion, and GMEM store. The resulting kernel used only 32 registers with 0 bytes of spill traffic. To further confirm this, I disassembled the full kernel with cuobjdump –dump-sass and located every LDL (spill reload) instruction in the instruction stream, 0 in the MMA compute region, 547 (100%) in the epilogue. This confirms the spill traffic comes entirely from the epilogue, not the Tensor Core compute path.

The point of TMEM isn’t fewer registers in total. It’s to keep the compute phase’s registers free so the MMA pipeline stays fed. The epilogue can tolerate spills because it can be chunked but the compute phase can’t.

The SFU Bottleneck

B200 Tensor Cores are enormously fast with 2.25 PFLOPS in FP16 and 9 PFLOPS in FP4. But attention isn’t just matmul. After QK^T you run softmax(QK^T / √d) which calls exp() on every element. The exp() function runs on the SFU (Special Function Unit), a separate hardware block that uses lookup table based interpolation.

The SFU has roughly 4–8 units per SM, compared to 32+ FP32 ALUs. It’s much lower throughput than the Tensor Core. On B200, Nvidia shrank the SM count to ~80 per die (to fit TMEM, DSMEM, and L2 changes) but the SFU per-SM throughput stayed roughly the same. The result is that we got ~4x faster Tensor Core but SFU didn’t scale.

SFU Ceiling

I ran a micro-benchmark that does nothing but __expf() in a tight loop. No memory access and no reduction, just pure SFU throughput:

SFU vs Tensor Core

Unit Throughput Source
B200 SFU (raw __expf ceiling) 1.4 TOPS Measured (100K iterations, 148 SMs, stable across run counts)
B200 Tensor Core FP16 (dense) 2,250 TOPS Spec sheet (2.25 PFLOPS)
B200 Tensor Core FP4 (dense) 9,000 TOPS Spec sheet (9.0 PFLOPS)

The ratio: ~1,600x. The SFU is three orders of magnitude slower than the Tensor Core. This is the structural imbalance on B200. You can matmul at 2.25 PFLOPS but you can only expf at 1.4 TOPS.

Softmax in Practice

A standalone softmax kernel (one block per row, __expf() for the exponential and shared-memory reduction for the sum) shows how this plays out in a real workload:

Softmax Time

The throughput plateaus at ~88B elements/sec. That’s the SFU wall. Adding more rows (larger seq_len) doesn’t increase throughput, it just increases time linearly. The SFU is saturated.

Cross-Die Memory Latency

The Architecture

B200 is two GPU dies in one package. Each die has its own L2 cache slice (~63MB) and its own HBM3e stacks (96GB). The two dies are connected by NV-HBI at 10 TB/s. An SM on die A can access die B’s L2 and HBM, but the request must traverse the NV-HBI interconnect. CUDA 13 exposes no CTA memory affinity API meaning you can’t tell the runtime “schedule this thread block near this data”.

L2: The Cross-Die Penalty is Visible

Latency vs Offset

SemiAnalysis measured this directly. They ran an L2 sized pointer chase 126MB sized to fill L2. With all 148 SMs traversing the same shuffled chain and recording per load latency with clock64(). By comparing latency vectors between every SM pair they built a 148×148 distance matrix. Two clear clusters emerged which were separated by 300 cycles. SMs on the same die see similar L2 latency patterns while those on the different die see inverse patterns. Local L2 hit was 150 cycles and remote L2 hit (cross-die) was around 450 cycles. The 300 cycle NV-HBI crossing is a clean 3× signal at L2 depths.

HBM: The Penalty Disappears

I measured the same question at HBM depth. I created a pointer chase kernel that traversed a 150GB buffer. This is 82% of the total VRAM spanning both dies. Each read’s address depends on the previous read’s value, preventing pipelining or prefetching. Each block ran 2000 dependent reads and recorded total cycles. I divided by 2000 to get per read latency. The result was uniform across all offsets. 1231–1327 cycles per read (1.08× ratio) with no bimodal distribution.

Latency Histogram

Latency vs Offset

Buffer Size Min cycles/read Max cycles/read Ratio
10 GB 447 484 1.08x
100 GB 1,230 1,322 1.07x
150 GB (spans both dies) 1,231 1,327 1.08x

No visible cross-die penalty. All blocks show uniform latency regardless of where in the 150GB buffer they read. The same ~300 cycle NV-HBI crossing exists at both levels. In my opinion the penalty is there at HBM but gets buried in noise.

  • L2 hit: ~150 cycles base. Adding ~300 cycles for NV-HBI produces a clean 3× signal (450 cycles) that sticks out above the low L2 jitter.
  • HBM access: ~1,000+ cycles base. A 300 cycle cross-die penalty is small relative to the base latency and its natural DRAM variance

Wrapping Up

Blackwell brings some serious firepower. A fully decoupled Tensor Core with its own memory, 2.25 PFLOPS FP16, 9 PFLOPS FP4 and an interconnect that hides the two die penalty at HBM. But the tradeoffs are real. The SFU is 1600× slower than the Tensor Core, scalar FP32 throughput dropped 33% and the epilogue is where register pressure actually lives. I’ll be working more with the B200 and sharing findings along the way.

References