NVIDIA Blackwell
NVIDIA Blackwell-architecture GPUs pack 208 billion transistors, manufactured on a custom TSMC 4NP process, split across two reticle-limited dies. Those dies are connected by a 10 terabytes per second chip-to-chip interconnect and presented to software as a single unified GPU rather than two separate devices the driver has to stitch together. The programming model underneath does not change: the Blackwell Tuning Guide describes the architecture as retaining and extending the same CUDA model used by Ampere and Hopper, so code already following those architectures’ best practices should see speedups without any changes.
At the streaming-multiprocessor level, occupancy limits move in both directions depending on which Blackwell you have. The maximum number of concurrent warps per SM is 64 for compute capability 10.0 devices and 48 for compute capability 12.0 devices; every SM carries a 64K-entry, 32-bit register file with up to 255 registers addressable per thread, and up to 32 thread blocks can be resident on one SM at once. Hopper’s thread block clusters carry over unchanged: a thread block can reach into another block’s shared memory within its cluster through distributed shared memory. Every Blackwell GPU supports a portable cluster size of up to 8, and the B200 additionally allows a nonportable cluster size of 16 for applications that opt in.
The memory system centers on HBM3 and HBM3e: the B200 GPU supports up to 180 GB of it, and the GB200 GPU raises L2 cache capacity to 126 MB. On-chip, the combined L1 data cache, texture cache, and shared memory still tops out at 256 KB per SM, the same ceiling as Hopper, with the shared-memory carveout selectable at runtime across the same 0, 8, 16, 32, 64, 100, 132, 164, 196, and 228 KB increments on both the H100 and the B200.
Scaling out relies on fifth-generation NVLink, which can connect up to 576 GPUs into a single domain. Inside a rack, the NVLink Switch Chip delivers 130 TB/s of aggregate GPU bandwidth across one 72-GPU NVL72 domain, and multiple NVL72 racks can be joined over the same 1.8 TB/s switch interconnect; NVIDIA states that a full NVL72 system supports nine times the GPU throughput of a single eight-GPU server. The link to the host side of the system is separate again: Blackwell reaches into Grace CPU memory over a dedicated connection running at 900 GB/s of bidirectional bandwidth. See NVIDIA Blackwell documentation in the reading list for the full tuning guide and architecture materials.
What CUTLASS targets on Blackwell#
The tuning guide describes the envelope; NVIDIA’s CUTLASS documentation describes what a matmul kernel actually executes inside it. Blackwell SM100 introduces the tcgen05.mma family: seven new matrix-multiply instructions that CUTLASS states are 2x to 4x faster than the Hopper architecture’s WGMMA instructions. They cover the legacy tensor core types (tf32, fp16, bf16, and 8-bit integers) alongside new 4-, 6-, and 8-bit floating point types, with and without scale factors, and each instruction comes in a cta_group::1 and a cta_group::2 variant. The 2-SM variants split one MMA tile across a pair of SMs, with tile shapes up to 256x256 in which each SM produces half of the output, and a kernel that selects a 2-SM instruction must launch with a cluster whose first dimension is a multiple of 2.
The block scaled instruction kinds are where Blackwell meets the same OCP microscaling standard that AMD’s CDNA 4 adopts in the next section. The mxf8f6f4, mxf4, and mxf4nvf4 kinds compute D = C + (A x SFA) * (B x SFB), where every 16 or 32 elements of A and B along the reduction dimension share one scale factor, so an M x K operand carries a much smaller matrix of scale factors beside it. CUTLASS’s mx_float8, mx_float6, and mx_float4 types pair the data with a ue8m0 scale factor over 32-element blocks and follow the OCP specification; nv_float4 pairs 4-bit data with a ue4m3 scale factor over 16-element blocks and is not OCP compliant. Per CUTLASS’s own throughput table, the mxf4 kinds reach 4x the throughput of Hopper’s FP8 tensor core.
None of this is reached through free-form assembly. CUTLASS exposes the instructions through its collective builder interface, where a kernel names element types, layouts, alignments, an MMA tile shape, a cluster shape, and a dispatch policy such as KernelTmaWarpSpecialized2SmMxf8f6f4Sm100, and the documentation’s tables enumerate exactly which combinations are valid. The constraints tighten as the types narrow: a float4_t operand must be 128-element aligned to target the f8f6f4 instruction kind, and the block scaled mxf4 kinds accept only the TN layout, row-major A against column-major B.
Scheduling gets its own hardware assist. CUTLASS GEMMs are persistent kernels: workers stay resident on the GPU and loop over output tiles, hiding prologue and epilogue costs, but a static tile schedule balances poorly when part of the GPU is occupied by another kernel. Blackwell adds cluster launch control: the kernel launches a grid with as many thread blocks as there are output tiles, and each resident worker asks the hardware for the next unstarted block coordinate via the clusterlaunchcontrol.try_cancel instruction. Every coordinate is guaranteed to be either launched as a worker or handed to an existing one, so tiles flow to whatever capacity actually exists, a dynamic schedule with the launch semantics of an ordinary grid.