SM_120 Microarchitecture Measured
Cycle-level characterization of an NVIDIA Blackwell GPU, synthesized from inline-PTX microbenchmarks.
Microarchitecture Diagram
Simplified SM block diagram and memory hierarchy, annotated with measured latency/throughput values. Peak issue throughput comes from the SMSP count and saturated instruction benchmarks.
Topology
Device limits and on-chip resources reported by the CUDA driver and measured microbenchmarks.
Pipeline Depth & ILP
ILP sweep for FFMA and IADD3: latency at ILP=1 approximates pipeline depth; saturation reveals issue throughput.
Co-Issue
Wall-clock cycle measurement of FMA-only, IMAD-only, and interleaved FMA+IMAD dependent chains. The overlap ratio estimates independent pipe dispatch.
Tensor Memory Accelerator (TMA)
Latency of a 1D cp.async.bulk.tensor global-to-shared load (16-byte tile) measured with an mbarrier.
Clock
Memory Hierarchy
Latencies are measured with dependency-chain or pointer-chase kernels; bandwidths use streaming copies.
Memory Subsystem Probes
Cache-line detection, coalescing behavior, and async global-to-shared copy latency.
Instructions
Latency is the dependent-chain cost per instruction; throughput is the saturated chip-level rate.
More Instructions
Logic, shift, permute, integer division, 64-bit integer, and type-conversion latencies.
Instruction Throughput
Saturated instruction throughput per SM, measured with ILP=4 independent accumulators and reported in instructions per cycle.
Tensor Core
MMA latency/throughput for various shapes and data types via inline PTX.
Tensor Core More Shapes
Additional MMA shapes such as FP64 DMMA.
Atomics
Barriers & Synchronization
Warp Shuffle
Warp Collectives & Divergence
Latency of warp-level collectives and cost of branch divergence.