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 typically see speedups without any code changes.
At the streaming-multiprocessor level, the guide summarizes Blackwell as offering similar occupancy to Hopper, with the ceiling 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.
The tuning guide never states what that HBM is worth in bandwidth, which is the number the roofline math of Foundations actually wants. NVIDIA’s Blackwell architecture technical brief does: its HGX system table gives 7.7 TB/s per GPU for HGX B200, against up to 192 GB of HBM3e per GPU, and the same figure for HGX B300 at 270 GB. For a single GPU in a GB200 the brief quotes 8 TB/s of HBM3e. The two documents do not agree on capacity, the tuning guide saying up to 180 GB and the brief up to 192 GB for what both call a B200, which is the ordinary situation with vendor documentation and the reason the comparison table at the end of this chapter says which document each of its B200 figures came from.
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.
Moving a kernel from Hopper to Blackwell#
The tuning guide’s headline claim is deliberately modest and worth taking literally. Blackwell retains and extends the same CUDA programming model, and applications already following the best practices for Ampere and Hopper “should typically see speedups on the Blackwell GPUs without any code changes.” The practices it means are the six high-priority recommendations it reprints from the general guides: parallelize sequential code, minimize host and device transfers, adjust the launch configuration to maximize device utilization, keep global memory accesses coalesced, minimize redundant global accesses, and avoid long sequences of diverged execution within a warp. Not one of them is new, and everything Blackwell-specific in the guide is what remains after they are done.
The first thing that actually breaks on a new architecture is not tuning but compilation. A CUDA binary carries compiled cubins for particular compute capabilities and, optionally, forward-compatible PTX. A cubin “is supported to run on any GPU with the same major revision and same or higher minor revision of compute capability,” so a Hopper cubin does not run on Blackwell at all, while PTX “is supported to run on any GPU with compute capability higher than the compute capability assumed for generation of that PTX” and is JIT-compiled at load. Binaries that include PTX “should work as-is on the Blackwell GPUs”; binaries that ship only cubins “need to be rebuilt.” Testing which one you have requires no Blackwell hardware: setting CUDA_FORCE_PTX_JIT=1 makes the runtime ignore embedded binary code and JIT the PTX instead, and an application that fails under it is one that will fail on the new architecture.
The exception is the case that matters most for kernel libraries. Architecture-conditional targets, written sm_100a and compute_100a, are exactly what a CUTLASS-style kernel uses to reach the tensor core instructions of the next subsection, and the compatibility guide states that binaries using them “are not forward or backward compatible. For example, PTX compiled for compute_90a (Hopper) are not supported on the Blackwell architecture.” The Hopper fast path does not survive even in PTX form. It has to be recompiled against a Blackwell-conditional target, which is why the tcgen05 instructions of the next subsection are a separate kernel family from Hopper’s WGMMA ones rather than a recompilation of them, and why a build covering both architectures lists both explicitly: -gencode=arch=compute_90,code=sm_90 beside -gencode=arch=compute_100,code=sm_100, with a trailing -gencode=arch=compute_100,code=compute_100 to keep PTX in the binary for whatever comes after. The shorthand -arch=sm_XX cannot express this, because it “can only specify a single target cubin architecture at a time.”
Once it compiles, the assumption most likely to be wrong is that there is one Blackwell. The occupancy ceilings above already split by compute capability, and shared memory splits harder: capacity per SM is 228 KB at compute capability 10.0 against 128 KB at 12.0, and the maximum a single thread block may address is 227 KB against 99 KB. A kernel sized to fill an SM on one Blackwell part can be unlaunchable on the other. Two older constants survive intact: CUDA reserves 1 KB of shared memory per thread block, and “static shared memory allocations remain limited to 48 KB, and an explicit opt-in is also required to enable dynamic allocations above this limit.”
The rest of the memory system rewards the same habits it did on Ampere. Blackwell keeps the unified L1 and texture cache, which “acts as a coalescing buffer for memory accesses, gathering up the data requested by the threads of a warp before delivery of that data to the warp,” so the coalescing arithmetic of Foundations is still what decides how many transactions a warp costs. The carveout between that cache and shared memory is a runtime decision made with cudaFuncSetAttribute() and the cudaFuncAttributePreferredSharedMemoryCarveout attribute rather than a compile-time one, and Blackwell also “allows CUDA users to control the persistence of data in L2 cache similar to the NVIDIA Ampere GPU Architecture,” which matters most for the repeatedly read operands the benchmark methodology of the profiling chapter goes to such lengths to flush.
Clusters carry over from Hopper unchanged, but the guide adds behavior that a specification sheet does not show. Distributed shared memory “can be used by an SM simultaneously with L2 cache accesses,” so a kernel communicating between SMs can draw on the combined bandwidth of both paths rather than choosing one. The access rules are those of Foundations, applied to a new address space: accesses to distributed shared memory “should be coalesced and aligned to 32-byte segments, if possible,” and non-unit stride patterns should be avoided, staged through local shared memory instead. Reaching for the nonportable cluster size means setting the cudaFuncAttributeNonPortableClusterSizeAllowed function attribute, and it is not free, because “using larger cluster sizes may reduce the maximum number of active blocks across the GPU.” For any cluster kernel the guide declines to give a number at all and recommends computing occupancy with cudaOccupancyMaxActiveClusters and launching accordingly.
Two smaller behaviors round out the move. NVLink “operates transparently within the existing CUDA model”: transfers between NVLink-connected endpoints are “automatically routed through NVLink, rather than PCIe,” but cudaDeviceEnablePeerAccess() is still required to enable direct transfers at all, with cudaDeviceCanAccessPeer() reporting whether a given pair can. And for code written before Volta that quietly assumes warp synchronicity, the compatibility guide offers an escape hatch rather than a fix: compiling with -gencode=arch=compute_60,code=sm_100 opts the kernel into the Pascal scheduling model on Blackwell hardware. It is a way to get a port running, not a place to leave it.
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.
Those instruction kinds are the software face of a hardware feature the architecture brief gives a name to. Blackwell carries a second-generation Transformer Engine, and NVIDIA describes it as using “advanced dynamic range management algorithms and fine-grain scaling techniques, called micro-tensor scaling, to optimize inference performance, accuracy, and enable FP4 AI.” NVIDIA claims three consequences, each a doubling: the performance of Blackwell’s FP4 tensor core, the parameter bandwidth to HBM, and the size of models supported per GPU. Read against the paragraph above, micro-tensor scaling is what a per-block scale factor is called when you describe it from the hardware side rather than from the instruction encoding, and the doublings are the vendor’s framing of why the mxf4 and mxf4nvf4 kinds exist at all: a 4-bit format is only usable if something restores the dynamic range that four bits cannot hold, and a scale factor per 16 or 32 elements is that something. The same brief states that the Blackwell Ultra GPU provides a 2x speedup over Blackwell GPUs for attention-layer compute, with new instructions to improve the performance of long input sequences, so an attention kernel is the one place to expect the two Blackwell parts to diverge without any change to the source.
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.