NVIDIA GPU Architecture History: From Tesla and Fermi to Blackwell

NVIDIA's GPU architecture went from 128 simple scalar cores in the 2006 G80 to dual-die Blackwell GPUs with 208 billion transistors. Each generation moved one of four levers: more parallel threads, faster memory, narrower number formats for matrix math, and smarter data movement. This post walks through every generation from Tesla to Blackwell, with figures cited to NVIDIA's own documents.

We read each NVIDIA architecture whitepaper, the CUDA Programming Guide, the architecture tuning guides and NVIDIA's own product pages and press releases on 25 October 2026. Every number here is a vendor-stated peak unless we say otherwise, and peak is not sustained: real kernels reach a fraction of these figures, and how large a fraction depends on the code. Where we did arithmetic ourselves, we checked it in Python and say so. The CUDA code samples follow NVIDIA's documentation but were not compiled or run, because this post was written on a Mac without an NVIDIA GPU.

Key takeaways

• Tesla (G80, 2006) unified the vertex and pixel shaders into one pool of scalar processors and made that pool programmable in C through CUDA. Everything since builds on its model of threads, blocks and warps of 32.
• Fermi (2009) made the GPU a credible compute device: real caches, ECC, IEEE 754-2008 fused multiply-add and fast double precision.
• Kepler, Maxwell and Pascal traded between raw width, efficiency and memory. Kepler added dynamic parallelism and Hyper-Q, Maxwell rebuilt the SM for performance per watt, and Pascal brought HBM2, NVLink and page-faulting unified memory.
• Volta (2017) added Tensor Cores and independent thread scheduling. From here on, the headline numbers are matrix-math numbers.
• Turing, Ampere and Hopper kept widening the matrix path (INT8, TF32, BF16, structured sparsity, FP8) and attacked data movement with asynchronous copies, the Tensor Memory Accelerator and thread block clusters.
• Blackwell (2024) joins two reticle-limited dies into one GPU, adds FP4 and microscaling formats, and scales NVLink to 72 GPU racks. Rubin is announced and, per NVIDIA, in production; we cover only what NVIDIA has published.
• The ratio of compute to memory bandwidth grew from about 4 to over 600 FLOPs per byte across these generations (our arithmetic from NVIDIA's figures), which is why modern GPU programming is mostly about feeding the math units.
• The software around these GPUs is a supply chain. Drivers, the CUDA toolkit, cuDNN, NCCL and container images all carry versions, support windows and vulnerabilities that you have to track.

How to read a GPU generation

Before the history, four terms that every section uses.

A streaming multiprocessor (SM) is the unit NVIDIA replicates to build a GPU. Each SM has its own scheduler, register file, shared memory and execution units. A big GPU is many SMs plus an L2 cache, memory controllers and interconnect.

A warp is a group of 32 threads that the SM schedules together. CUDA has kept the warp size at 32 from the first CUDA guide in 2007 to the current CUDA 13 guide. Threads are grouped into thread blocks, blocks into a grid, and since Hopper, blocks can also be grouped into clusters.

Compute capability (CC) is NVIDIA's version number for what an SM can do. It has a major number (the architecture family) and a minor number (variants within the family). The compiler targets it with flags such as sm_80 or sm_90a.

Peak versus sustained. A peak FLOP rate is units times operations per clock times clock speed. Sustained throughput depends on whether the kernel can keep those units fed.

Each generation pulls one or more of four levers:

Timeline: NVIDIA compute architectures with dates from primary sources

The dates below are the dates printed on, or attached to, the NVIDIA documents we fetched. They are document dates, not always first shipment dates.

For Kepler, the whitepaper edition we could fetch is the GK110/GK210 edition carrying a 2014 copyright. We did not find a primary NVIDIA page in our fetches that states the first GK110 ship date, so we leave that out rather than guess.

Compute capability by generation

This is the table people search for most. Product mappings come from NVIDIA's current and legacy compute capability pages; architecture labels for older parts come from the whitepapers.

Two details matter for anyone who builds or ships CUDA software. First, CC 9.0 introduced architecture-specific targets with an a suffix (for example compute_90a), and CC 10.0 introduced family-specific targets with an f suffix. NVIDIA's CUDA 13 guide says code compiled for an a target runs only on that exact compute capability, while compute_100f runs on 10.0, 10.3 and 10.7 devices. Second, NVIDIA's CUDA 13 compute capability tables start at 7.5, and the cuFFT 13.0 release notes list removed support for Maxwell, Pascal and Volta. A toolkit upgrade can therefore drop a GPU you still run.

The per-SM limits are where the generations really differ for a programmer. The table collects the resident limits that drive occupancy.

The register file per SM stopped growing at 64K 32-bit registers (256 KB) after Kepler GK110, except for GK210. What grew instead was the number of SMs, shared memory per SM and everything around the SM.

Tesla: the G80 and the unified shader (2006)

Before G80, graphics chips had separate hardware for vertex shading and pixel shading. If a frame was heavy on pixels, the vertex units sat idle, and the reverse. NVIDIA's GeForce 8800 technical brief, dated 8 November 2006, describes the G80 as a "fully unified architecture" in which one pool of processors runs every shader type.

The brief gives the core numbers. The GeForce 8800 GTX has 128 stream processors running at 1.35 GHz, while the rest of the core runs at 575 MHz. The brief argues for 128 scalar processors over 32 four-component vector units, because scalar units stay busy on code that does not fill all four lanes of a vector. That choice, one thread per lane with scalar instructions, is what made the chip programmable as a general parallel processor.

Memory came from six partitions, each with a 64-bit interface, for a 384-bit bus. The 8800 GTX used GDDR3 at a 900 MHz default clock and 768 MB of frame buffer. The brief states 86.4 GB/s of bandwidth, and the arithmetic checks out:

bus width      384 bits = 48 bytes per transfer
data rate      900 MHz clock x 2 transfers per clock (double data rate)
bandwidth      48 B x 1.8e9 transfers/s = 86.4e9 B/s = 86.4 GB/s

CUDA 1.0 and the first compute capability

The CUDA Programming Guide 1.0 (23 June 2007) describes the device as "a set of SIMD multiprocessors with on-chip shared memory". It lists 16 multiprocessors for the GeForce 8800 GTX and the Tesla C870, so each multiprocessor held 8 of the 128 stream processors. Appendix A gives the limits for compute capability 1.x:

• 8,192 registers per multiprocessor;
• 16 KB of shared memory per multiprocessor, in 16 banks;
• 64 KB of constant memory with an 8 KB cache working set per multiprocessor;
• at most 8 resident blocks, 24 warps and 768 threads per multiprocessor;
• at most 512 threads per block;
• a warp size of 32.

Compute capability 1.1 added atomic functions. That guide used the word SIMD; NVIDIA's later documents describe the same execution model as SIMT, single instruction multiple thread, because each lane is programmed as its own thread with its own registers.

G80 streaming multiprocessor (simplified, from the CUDA 1.0 guide and 8800 brief)

+------------------------------------------------+
| Instruction unit (one warp of 32 issued        |
| over several clocks on 8 processors)           |
+------------------------------------------------+
| SP | SP | SP | SP | SP | SP | SP | SP          |  8 scalar stream processors
+------------------------------------------------+
| Registers: 8,192 x 32-bit                      |
| Shared memory: 16 KB, 16 banks                 |
| Constant cache 8 KB | Texture cache 8 KB (1D)  |
+------------------------------------------------+
x 16 multiprocessors = 128 stream processors on GeForce 8800 GTX

Worked number: the G80 compute to bandwidth ratio

NVIDIA's brief gives clock and core count but our sources did not give a single peak FLOP figure we could cite, so here is the arithmetic, labelled as ours. If each of the 128 processors completes one multiply-add (two floating point operations) per clock at 1.35 GHz:

128 x 2 FLOP x 1.35e9 /s = 345.6 GFLOP/s (multiply-add only, our arithmetic)
345.6 GFLOP/s / 86.4 GB/s = 4.0 FLOP per byte

Keep that 4 FLOPs per byte in mind. By Blackwell it is over 600 for FP8 matrix math.

Why it mattered

The G80 made a cheap consumer graphics card into a parallel computer with a C dialect. A thread hierarchy of grids, blocks and threads, a fast scratchpad (shared memory) and barrier synchronization inside a block were all present in CUDA 1.0. They are still the core of CUDA today. A minimal CUDA kernel from 2007 still looks like this:

// SAXPY: y = a*x + y. Follows the CUDA programming model; not executed on this Mac.
__global__ void saxpy(int n, float a, const float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

// Host side: launch enough 256-thread blocks to cover n elements.
// saxpy<<<(n + 255) / 256, 256>>>(n, 2.0f, d_x, d_y);

Fermi: caches, ECC and real double precision (2009)

The Fermi Compute Architecture whitepaper (v1.1, 2009) opens with a list of what GPU computing users had asked for: faster double precision, ECC, a true cache hierarchy, more shared memory, faster context switching and faster atomics. Fermi answered almost every item.

The first Fermi GPU, GF100, had 3.0 billion transistors and up to 512 CUDA cores in 16 SMs of 32 cores each. Six 64-bit memory partitions formed a 384-bit interface to up to 6 GB of GDDR5. The whitepaper lists the SM changes:

• 32 CUDA cores per SM, four times the GT200 SM;
• 8 times the peak double precision rate of GT200, with up to 16 double precision fused multiply-adds per SM per clock;
• a dual warp scheduler that picks two warps and issues one instruction from each;
• 64 KB of on-chip memory per SM, configurable as 48 KB shared memory with 16 KB L1 cache, or the reverse;
• a 768 KB unified L2 cache shared by all SMs (the GT200 had no L2 for compute);
• full IEEE 754-2008 single and double precision, with fused multiply-add for both;
• a unified address space, which let CUDA support C++ pointers properly;
• ECC on memory, using single-error-correct, double-error-detect (SECDED) codes;
• up to 16 kernels running concurrently, and roughly 10 times faster application context switching.

Fermi SM (GF100), from the Fermi whitepaper

+--------------------------------------------------------------+
| Instruction cache                                            |
| Warp scheduler 0 + dispatch | Warp scheduler 1 + dispatch    |
| Register file: 32,768 x 32-bit                               |
+--------------------------------------------------------------+
| 32 CUDA cores (FP32 and integer), arranged as 2 x 16         |
| 16 load/store units | 4 special function units (SFU)         |
+--------------------------------------------------------------+
| Interconnect network                                         |
| 64 KB shared memory / L1 cache (48/16 or 16/48)              |
| Uniform cache | Texture cache                                |
+--------------------------------------------------------------+
x 16 SMs, around a common 768 KB L2

Why fused multiply-add mattered

A separate multiply then add rounds twice. A fused multiply-add computes a*b + c with one rounding at the end. The whitepaper points out that this improves accuracy in the addition step and that Fermi supports FMA in both single and double precision, matching IEEE 754-2008. ECC was the other half: you cannot run thousands of GPUs in a data center if a single flipped bit silently corrupts a simulation.

The occupancy model gets bigger

The Kepler whitepaper's comparison table gives the Fermi limits: 1,536 threads, 48 warps and 8 blocks per SM, 32,768 registers per SM, and at most 63 registers per thread. Doubling resident warps from 24 to 48 gave the scheduler more warps to switch to while others waited on memory, which is how GPUs hide latency.

Worked number: Fermi double precision per clock

Our sources give Fermi's FP64 rate per clock, not a product clock, so we show the per clock figure only:

16 FP64 FMA per SM per clock x 2 FLOP per FMA x 16 SMs = 512 FP64 FLOP per clock

Multiply by a product's shader clock to get its FP64 peak. We did not find a primary NVIDIA source in our fetches for a specific Fermi product clock, so we stop there.

Kepler: the SMX, dynamic parallelism and Hyper-Q

The Kepler GK110/GK210 whitepaper describes a 7.1 billion transistor GPU with up to 15 SMX units and six 64-bit memory controllers. Products could ship with 13, 14 or 15 SMX enabled. GK110 is compute capability 3.5 and GK210, used in the Tesla K80, is 3.7.

Kepler made the SM very wide. Each SMX has 192 single precision CUDA cores, 64 double precision units, 32 special function units and 32 load/store units, fed by four warp schedulers. The whitepaper describes several features that changed how CUDA programs could be structured:

• Dynamic parallelism. A kernel can launch new kernels on the GPU without going back to the CPU. The new Grid Management Unit tracks these launches.
• Hyper-Q. 32 simultaneous, hardware-managed connections between host and GPU, up from Fermi's single connection, so many CPU processes or streams can feed one GPU without false serialization.
• 255 registers per thread (up from 63), through a new instruction encoding, which helps register-hungry kernels such as double precision linear algebra.
• Shuffle instructions that let threads in a warp read each other's registers directly, without a round trip through shared memory.
• A 48 KB read-only data cache per SMX for data known to be constant during a kernel.
• GPUDirect RDMA, letting a network card or other device read and write GPU memory without staging through system memory.

Worked number: Tesla K40 peak, from first principles

The Volta whitepaper's comparison table lists the Tesla K40 (GK180, a GK110 variant) with 15 SMs, 2,880 FP32 cores, 960 FP64 cores and a boost clock of 810 or 875 MHz, and states 5 TFLOPS FP32 and 1.7 TFLOPS FP64. The Pascal whitepaper's table gives 5,040 and 1,680 GFLOPS. Our arithmetic at 875 MHz:

FP32: 2,880 cores x 2 FLOP x 0.875e9 = 5.04e12 FLOP/s = 5.04 TFLOPS
FP64:   960 cores x 2 FLOP x 0.875e9 = 1.68e12 FLOP/s = 1.68 TFLOPS

This is how every peak in this post is built: units, times two for a fused multiply-add, times clock. When a vendor number and your arithmetic disagree, the usual reason is a different clock (base versus boost) or sparsity.

Dynamic parallelism in code

Dynamic parallelism is still in the CUDA 13 guide. It needs relocatable device code (-rdc=true) at compile time. A parent kernel can launch a child grid sized by data it discovers. The sample below follows the guide's pattern. Note that recent CUDA versions removed device-side cudaDeviceSynchronize(); the parent does not wait on the child inside the kernel.

// Compile with: nvcc -arch=sm_90 -rdc=true dp.cu -lcudadevrt
// Follows the CUDA Programming Guide; not executed on this Mac.
__global__ void child(float* data, int offset, int count) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < count) data[offset + i] *= 2.0f;
}

__global__ void parent(float* data, const int* segment_start,
                       const int* segment_len, int segments) {
    int s = blockIdx.x * blockDim.x + threadIdx.x;
    if (s < segments) {
        int len = segment_len[s];
        if (len > 0) {
            // Each parent thread launches a grid sized to its own segment.
            child<<<(len + 127) / 128, 128>>>(data, segment_start[s], len);
        }
    }
}

Why it mattered

Kepler arrived as GPUs were becoming the default engine for neural network training: NVIDIA's Pascal whitepaper notes that the 2012 ImageNet entry now called AlexNet "used the parallel processing performance of GPUs". Kepler's features were about feeding the GPU from many sources (Hyper-Q, GPUDirect) and letting it schedule itself (dynamic parallelism). The very wide 192-core SMX also showed a limit: wide SMs are hard to keep busy, which is what Maxwell went after next.

Maxwell: the SMM and performance per watt (2014)

NVIDIA's Maxwell tuning guide says the first Maxwell GPU, GM107, was designed for power-limited notebooks and small form factor PCs. The second-generation GM204 followed. NVIDIA's developer blog of 18 September 2014 introduced GM204, which powers the GeForce GTX 980 and 970, with 16 SMs and 2,048 CUDA cores at a 165 W TDP.

The new SM, called SMM, went the other way from Kepler. Instead of one wide 192-core array, it is split into four quadrants, each with 32 cores and its own warp scheduler that can issue two instructions per clock. The blog states this gives more than 40 percent higher delivered performance per CUDA core than Kepler and, overall, twice the efficiency of Kepler GK104.

Maxwell SMM, quadrant layout (from the Maxwell blog and tuning guide)

+----------------------+----------------------+
| Warp scheduler       | Warp scheduler       |
| 32 cores, own regs   | 32 cores, own regs   |
+----------------------+----------------------+
| Warp scheduler       | Warp scheduler       |
| 32 cores, own regs   | 32 cores, own regs   |
+----------------------+----------------------+
| 96 KB dedicated shared memory (GM204)       |
| Unified L1 / texture cache                  |
+---------------------------------------------+

Three changes mattered for compute programmers:

1. Dedicated shared memory. Fermi and Kepler split one array between L1 and shared memory. The tuning guide says Maxwell gives shared memory its own space, and the blog puts it at 64 KB per SMM on GM107 and 96 KB on GM204. L1 merged with the texture cache instead.
2. Native shared memory atomics. Kepler implemented shared memory atomics with a lock, update and unlock sequence. GM204 adds native 32-bit integer atomics and 32-bit and 64-bit compare-and-swap in shared memory, which speeds up histograms and similar code.
3. A larger L2. GM204's L2 is 2 MB, four times GK104's 512 KB, so working sets between those sizes stop hitting DRAM.

The blog also states that GM204 supports up to 32 active thread blocks per SM, which helps kernels with small blocks.

The trade you can see in the data center table

The Volta whitepaper's comparison table shows the price of Maxwell's efficiency focus in the data center. The Tesla M40 (GM200) has 24 SMs and 3,072 FP32 cores for 6.8 TFLOPS FP32, but only 96 FP64 cores for 0.21 TFLOPS FP64. The older K40 had 1.7 TFLOPS FP64. Maxwell's big chip was a single precision part. For deep learning that was fine. For double precision simulation codes, the K40 and K80 stayed in service until Pascal.

Pascal: HBM2, NVLink and page-faulting unified memory (2016)

The GP100 whitepaper (May 2016) describes a 15.3 billion transistor GPU on a 16 nm FinFET process. A full GP100 has six GPCs, 60 SMs in 30 TPCs and eight 512-bit memory controllers. The Tesla P100 uses 56 SMs.

Pascal went back to smaller SMs: 64 FP32 cores and 32 FP64 cores per SM, with 64 KB of shared memory and the same 256 KB register file per SM as before. Fewer cores per register file means more registers and shared memory per thread. The P100's headline figures, from the whitepaper:

• 5.3 TFLOPS FP64, 10.6 TFLOPS FP32 and 21.2 TFLOPS FP16;
• 16 GB of HBM2 on a 4,096-bit interface;
• 4 MB of L2;
• 300 W TDP.

Our arithmetic at the 1,480 MHz boost clock in the whitepaper's table reproduces those numbers: 3,584 FP32 cores give 10.61 TFLOPS, 1,792 FP64 cores give 5.30 TFLOPS, and paired FP16 (two half precision operations in one instruction, as the whitepaper notes) doubles FP32 to 21.2 TFLOPS.

HBM2 and CoWoS

The whitepaper calls P100 the world's first GPU with HBM2 and says HBM2 delivers three times the memory bandwidth of the Maxwell GM200. HBM stacks DRAM dies vertically and places them next to the GPU on a silicon interposer, which NVIDIA's paper describes using TSMC's CoWoS (chip on wafer on substrate) packaging. HBM2 also brings native ECC support. Every NVIDIA data center flagship since has used HBM. For how that package is built, see our post on how a GPU is made, from wafer to HBM package.

NVLink, first generation

NVLink was Pascal's answer to PCIe as a bottleneck between GPUs. The whitepaper's appendix describes each link as eight differential pairs per direction at up to 20 Gb/s, giving 40 GB/s of bidirectional bandwidth per link. P100 has four links, which is 160 GB/s bidirectional per GPU (our arithmetic: 4 times 40). The whitepaper also describes CPU to GPU NVLink with IBM POWER8 processors.

Unified memory with page faulting

CUDA 6 had introduced unified memory, but before Pascal the driver had to migrate whole allocations before a kernel ran. The whitepaper says GP100 is the first NVIDIA GPU with hardware page faulting, and pairs it with 49-bit virtual addressing, enough for 512 TB, covering the CPU's address space plus the GPU's. Pages now migrate on demand when a kernel touches them, and oversubscription (allocating more than the GPU's memory) becomes possible.

// Unified memory with on-demand page migration (Pascal and later).
// Follows the CUDA Programming Guide; not executed on this Mac.
#include <cstdio>
__global__ void scale(float* v, size_t n, float s) {
    size_t i = blockIdx.x * (size_t)blockDim.x + threadIdx.x;
    if (i < n) v[i] *= s;
}

int main() {
    size_t n = 1ull << 28;                    // 268,435,456 floats = 1 GiB
    float* v = nullptr;
    cudaMallocManaged(&v, n * sizeof(float)); // one pointer valid on CPU and GPU
    for (size_t i = 0; i < n; ++i) v[i] = 1.0f;          // pages start on the CPU
    scale<<<(unsigned)((n + 255) / 256), 256>>>(v, n, 3.0f); // GPU faults pages in
    cudaDeviceSynchronize();
    printf("%f\n", v[n - 1]);                 // CPU touch migrates the page back
    cudaFree(v);
    return 0;
}

On-demand faulting is convenient but each fault costs time. Production code usually adds explicit prefetch hints; the prefetch API's signature has changed between CUDA releases, so check the runtime API reference for your toolkit version before copying an older sample.

Compute preemption

Pascal also added compute preemption at instruction-level granularity, according to the whitepaper, where Kepler and Maxwell could only preempt at thread block boundaries. That means a long running kernel no longer blocks the display or other work, and debuggers can stop a kernel safely.

Volta: Tensor Cores and independent thread scheduling (2017)

The Volta whitepaper (August 2017) describes GV100: 21.1 billion transistors on an 815 mm² die, built on a TSMC 12 nm FFN process customized for NVIDIA. A full GV100 has 84 SMs; the Tesla V100 uses 80.

The SM split into four

Each Volta SM is divided into four processing blocks. The whitepaper lists each block as having 16 FP32 cores, 8 FP64 cores, 16 INT32 cores, two Tensor Cores, an L0 instruction cache, one warp scheduler, one dispatch unit and a 64 KB register file. Separate INT32 cores mean address arithmetic no longer competes with floating point for the same pipe.

Volta SM (GV100), from the Volta whitepaper

+------------------+------------------+------------------+------------------+
| L0 i-cache       | L0 i-cache       | L0 i-cache       | L0 i-cache       |
| Warp scheduler   | Warp scheduler   | Warp scheduler   | Warp scheduler   |
| Regs 64 KB       | Regs 64 KB       | Regs 64 KB       | Regs 64 KB       |
| 16 FP32 16 INT32 | 16 FP32 16 INT32 | 16 FP32 16 INT32 | 16 FP32 16 INT32 |
| 8 FP64           | 8 FP64           | 8 FP64           | 8 FP64           |
| 2 Tensor Cores   | 2 Tensor Cores   | 2 Tensor Cores   | 2 Tensor Cores   |
+------------------+------------------+------------------+------------------+
| 128 KB combined L1 data cache and shared memory (up to 96 KB shared)      |
+---------------------------------------------------------------------------+

What a Tensor Core does

The whitepaper defines the operation: each Tensor Core performs D = A x B + C on 4x4 matrices, where A and B are FP16 and C and D can be FP16 or FP32. One 4x4x4 multiply is 64 fused multiply-adds. With 8 Tensor Cores per SM and 640 on the V100, the whitepaper states 125 Tensor TFLOPS. Our check, at the 1,530 MHz boost clock in its table:

640 Tensor Cores x 64 FMA per clock x 2 FLOP x 1.53e9 = 125.3e12 FLOP/s = 125 TFLOPS
V100 FP32 for comparison: 5,120 cores x 2 x 1.53e9 = 15.7 TFLOPS

That is eight times the FP32 rate from the same chip, as long as your problem is matrix multiplication in FP16 with FP32 accumulation. Deep learning is mostly that, which is why Tensor Cores reshaped the product line.

In CUDA C++, the warp matrix (WMMA) API exposes Tensor Cores at warp level. All 32 threads of a warp cooperate on one 16x16x16 tile:

// Warp-level matrix multiply with WMMA (Volta and later).
// Follows the CUDA Programming Guide; not executed on this Mac.
#include <mma.h>
#include <cuda_fp16.h>
using namespace nvcuda;

__global__ void wmma_tile(const half* a, const half* b, float* c) {
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> fa;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> fb;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> fc;

    wmma::fill_fragment(fc, 0.0f);
    wmma::load_matrix_sync(fa, a, 16);   // leading dimension 16
    wmma::load_matrix_sync(fb, b, 16);
    wmma::mma_sync(fc, fa, fb, fc);      // fc = fa x fb + fc on Tensor Cores
    wmma::store_matrix_sync(c, fc, 16, wmma::mem_row_major);
}
// Launch with one warp: wmma_tile<<<1, 32>>>(d_a, d_b, d_c);

Later generations moved the fastest Tensor Core paths to warpgroup instructions (wgmma in the PTX ISA, for Hopper) and to tcgen05 instructions (for Blackwell), which NVIDIA recommends using through libraries such as cuBLAS and CUTLASS rather than by hand.

Independent thread scheduling

The other big Volta change is less visible but changes correctness. Before Volta, a warp shared one program counter. When threads in a warp took different branches, the hardware ran one path, then the other, with a mask, and the paths could not interleave. The whitepaper describes Volta as keeping per-thread scheduling state, including a program counter and call stack per thread, so the scheduler can interleave divergent paths and threads can wait on each other inside a warp.

The practical consequence: code that assumed all 32 threads of a warp run in lockstep ("warp-synchronous" code) can break. The whitepaper introduces __syncwarp() in CUDA 9 to make convergence explicit. Every warp-level primitive since (__shfl_sync, __ballot_sync) takes a mask for this reason.

Memory and links

V100 has 16 GB of HBM2 with 900 GB/s of peak bandwidth, a 6 MB L2, and NVLink 2.0 with six links totalling 300 GB/s. Shared memory and L1 merged again into a 128 KB array per SM, with up to 96 KB usable as shared memory.

Turing: RT Cores, INT8 and INT4, GDDR6 (2018)

Turing is mostly a graphics architecture, but it shaped inference hardware. The Turing whitepaper describes TU102 with 18.6 billion transistors on TSMC 12 nm FFN, six GPCs, 36 TPCs and 72 SMs, twelve 32-bit GDDR6 controllers (384 bits) and 6 MB of L2. Each SM has 64 FP32 cores, 8 Tensor Cores, one RT Core and 96 KB of L1 and shared memory. Turing is compute capability 7.5.

Four changes stand out:

1. RT Cores accelerate bounding volume hierarchy traversal and ray-triangle intersection. The whitepaper lists 10 giga rays per second for the RTX 2080 Ti and Quadro RTX 6000.
2. INT8 and INT4 Tensor Core modes for inference. The Quadro RTX 6000 is listed at 130.5 FP16 Tensor TFLOPS, 261 INT8 TOPS and 522 INT4 TOPS: halve the bits, double the rate.
3. Concurrent integer and floating point execution. The whitepaper says that across many workloads NVIDIA sees about 36 integer instructions for every 100 floating point instructions, and Turing runs them on a separate pipe at the same time.
4. GDDR6, first on Turing per the whitepaper.

Worked numbers: Turing memory and the integer pipe

The whitepaper's table gives the memory figures, and they follow from bus width times per-pin rate:

RTX 2080 Ti:       352-bit x 14 Gbps / 8 = 616 GB/s
Quadro RTX 6000:   384-bit x 14 Gbps / 8 = 672 GB/s

On the integer pipe: if 36 integer instructions per 100 floating point instructions had to share one pipe, the work takes 136 issue slots; on separate pipes it can take as few as 100. That puts the upper bound on the speedup from this change at 1.36 times for that instruction mix (our arithmetic). Real gains are lower because of dependencies and memory stalls.

Turing also brought Tensor Cores to the T4 data center inference card, which NVIDIA's Ampere whitepaper notes as the 2018 move of Tensor Cores into mainstream inference.

Ampere: TF32, sparsity, MIG and asynchronous copy (2020)

The Ampere whitepaper (2020) describes GA100 with 54.2 billion transistors on an 826 mm² die, built on TSMC's 7 nm N7 process. A full GA100 has 8 GPCs and 128 SMs; the A100 Tensor Core GPU ships with 108 SMs, 6,912 FP32 cores and 432 third-generation Tensor Cores (four per SM).

The A100's peak figures, as the whitepaper lists them, show how far the matrix path had pulled away from the scalar one:

Our arithmetic reproduces the dense figures at the 1,410 MHz boost clock the whitepaper uses. FP32 is 6,912 cores times 2 times 1.41 GHz, or 19.49 TFLOPS. For FP16 Tensor Cores, the whitepaper says the four Tensor Cores of an A100 SM have twice the raw FMA rate of the eight V100 Tensor Cores in a V100 SM. Eight V100 Tensor Cores at 64 FMA each is 512 FMA per SM per clock, so each A100 Tensor Core does 256: 432 times 256 times 2 times 1.41 GHz is 311.9 TFLOPS.

TF32 and BF16

TF32 is the format that made Tensor Cores useful for code written in FP32. The whitepaper defines it as 1 sign bit, an 8-bit exponent (the same range as FP32) and a 10-bit mantissa (the same precision as FP16), 19 bits in all. Tensor Cores read FP32 inputs, round them to TF32 internally and accumulate in FP32, so frameworks can turn it on without changing the data type in user code.

BF16 has an 8-bit exponent and a 7-bit mantissa. It keeps FP32's range with less precision than FP16, which avoids most of the loss scaling FP16 training needs. On A100, BF16 Tensor Core math runs at the FP16 rate.

Format bit layouts (sign | exponent | mantissa), from NVIDIA whitepapers
FP32   1 | 8 | 23
TF32   1 | 8 | 10    Ampere, internal Tensor Core format
BF16   1 | 8 | 7
FP16   1 | 5 | 10
FP8    1 | 4 | 3     E4M3, Hopper
FP8    1 | 5 | 2     E5M2, Hopper
FP4    1 | 2 | 1     E2M1, used by NVFP4 on Blackwell

Fine-grained structured sparsity

Ampere Tensor Cores can skip zeros when a weight matrix follows a 2:4 pattern: in every group of four values, at most two are non-zero. The hardware stores only the non-zero values plus small indices, and the whitepaper states twice the math throughput on those operations. The catch is that you have to prune and usually fine-tune the model to fit the pattern; the whitepaper reports accuracy results for several networks after doing so. Every "with sparsity" number since then, on Hopper and Blackwell too, is this 2x factor. When you compare generations, compare dense to dense.

Multi-Instance GPU

MIG lets one A100 be split into as many as seven GPU instances, each with its own SMs, L2 slice and memory controllers, so separate jobs get isolated performance. For inference fleets running many small models this matters more than raw FLOPs.

Memory, L2 and asynchronous copy

A100 has 40 GB of HBM2 at 1,555 GB/s, which the whitepaper calls a 73 percent increase over V100, and a 40 MB L2, nearly seven times V100's. The combined L1 and shared memory array grew to 192 KB per SM, with up to 164 KB as shared memory (CC 8.0). Third-generation NVLink doubled the link count to twelve, for 600 GB/s per GPU.

The quiet but important change is asynchronous copy. Before Ampere, loading a tile into shared memory meant each thread loaded from global memory into a register, then stored it to shared memory. The whitepaper describes a new instruction that copies from global memory straight into shared memory, skipping the register file, and runs in the background while threads keep computing. Hardware asynchronous barriers track when the data has arrived.

// Asynchronous global-to-shared copy with cooperative groups (CC 8.0 and later
// have hardware acceleration). Assumes n is a multiple of blockDim.x.
// Follows the CUDA Programming Guide; not executed on this Mac.
#include <cooperative_groups.h>
#include <cooperative_groups/memcpy_async.h>
namespace cg = cooperative_groups;

__global__ void stage_and_scale(const float* in, float* out, float s) {
    extern __shared__ float tile[];                 // dynamic shared memory
    cg::thread_block block = cg::this_thread_block();
    size_t base = (size_t)blockIdx.x * blockDim.x;

    // The whole block issues one collective copy; no per-thread register staging.
    cg::memcpy_async(block, tile, in + base, sizeof(float) * blockDim.x);
    cg::wait(block);                                // block until the copy lands

    out[base + threadIdx.x] = tile[threadIdx.x] * s;
}
// Launch: stage_and_scale<<<n / 256, 256, 256 * sizeof(float)>>>(d_in, d_out, 2.0f);

The payoff comes from double buffering: copy tile k+1 while computing on tile k, as fast matrix multiply kernels do.

Ampere outside the data center

GA102 and its siblings in the RTX 30 series and A40 are compute capability 8.6, with 1,536 threads and 100 KB of shared memory per SM rather than GA100's 2,048 and 164 KB. Ada (CC 8.9), used in the L4, L40 and RTX 40 series, keeps those SM limits and, according to the CUDA 13 guide's Tensor Core table, adds FP8 Tensor Core inputs. Code tuned for an A100 does not always fit the smaller SM of an RTX card, which is a common source of surprise.

Hopper: FP8, the Transformer Engine, TMA and clusters (2022)

NVIDIA introduced Hopper on 22 March 2022. Its technical blog describes GH100: 80 billion transistors on an 814 mm² die, built on TSMC 4N. A full GH100 has 8 GPCs, 72 TPCs, 144 SMs, six HBM stacks and a 60 MB L2. The H100 SXM5 ships with 132 SMs, 16,896 FP32 cores, 528 fourth-generation Tensor Cores, 80 GB of HBM3 in five stacks and a 50 MB L2. The PCIe version has 114 SMs. Hopper is compute capability 9.0.

Preliminary versus final numbers

The March 2022 blog's tables carry NVIDIA's own warning that the figures were preliminary, with the boost clock "not finalized". The H100 product page lists the shipping figures, and they are higher than the preliminary ones:

The product page lists Tensor Core figures "with sparsity", so the dense numbers in the right column are our division by two. The FP32 figure implies a boost clock near 1.98 GHz (67 TFLOPS divided by 16,896 cores times 2), which is our arithmetic, not a clock NVIDIA states on the pages we read. When you read old blog posts about H100, check which set of numbers they used.

FP8 and the Transformer Engine

Hopper's fourth-generation Tensor Cores add two FP8 formats. The blog defines E4M3 (4 exponent bits, 3 mantissa bits) and E5M2 (5 exponent bits, 2 mantissa bits). In the blog's words, E4M3 suits computations that need less dynamic range and more precision, while E5M2 gives wider dynamic range with less precision. The blog states FP8 halves storage and doubles throughput compared with FP16 or BF16, and per SM, H100 Tensor Cores deliver twice the A100 rate on the same type and four times using FP8 against A100's 16-bit formats.

FP8 is too narrow to use naively, so NVIDIA paired it with the Transformer Engine: software plus hardware that, layer by layer, tracks the range of each tensor and picks scaling factors and the FP8 or 16-bit format to keep accuracy. In practice you meet it through NVIDIA's Transformer Engine library and frameworks, not by writing FP8 scale management yourself.

The Tensor Memory Accelerator

As Tensor Cores got faster, the threads that loaded data for them became the bottleneck. The Hopper tuning guide describes the TMA as an asynchronous copy engine that moves 1D to 5D tensors between global and shared memory, in both directions, and between the shared memories of SMs in the same cluster. A single thread issues a large copy and the rest of the block keeps working. The blog puts it plainly: only a small number of CUDA threads are needed to drive the full memory bandwidth of H100 with TMA.

Thread block clusters and distributed shared memory

Hopper added a level to the CUDA thread hierarchy for the first time since 2007. A cluster is a group of thread blocks that the hardware guarantees to schedule at the same time on SMs within one GPC. Blocks in a cluster can read, write and run atomics on each other's shared memory, which NVIDIA calls distributed shared memory. The tuning guide gives a portable maximum cluster size of 8 blocks, with H100 allowing 16 when a kernel opts in to a non-portable size.

CUDA thread hierarchy, Tesla to Hopper

2007 (CC 1.0):  grid  ->  block  ->  warp (32)  ->  thread
2022 (CC 9.0):  grid  ->  cluster (up to 8, or 16 opt-in)  ->  block  ->  warp  ->  thread
                            |
                            +-- blocks share each other's shared memory (DSMEM)

// Thread block cluster with distributed shared memory (CC 9.0 and later).
// Each block counts values into its own shared histogram, then adds its counts
// into the partner block's histogram. Follows the CUDA Programming Guide;
// not executed on this Mac.
#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void __cluster_dims__(2, 1, 1)
pair_histogram(const unsigned char* data, size_t n, unsigned int* out) {
    __shared__ unsigned int bins[256];
    cg::cluster_group cluster = cg::this_cluster();
    unsigned int rank = cluster.block_rank();

    for (int b = threadIdx.x; b < 256; b += blockDim.x) bins[b] = 0;
    cluster.sync();                       // all blocks in the cluster initialised

    size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) atomicAdd(&bins[data[i]], 1u);
    __syncthreads();

    // Map the partner block's shared array into this block's address space.
    unsigned int* partner = cluster.map_shared_rank(bins, rank ^ 1);
    if (rank == 1)
        for (int b = threadIdx.x; b < 256; b += blockDim.x)
            atomicAdd(&partner[b], bins[b]);   // remote shared-memory atomic
    cluster.sync();                       // keep shared memory alive until done

    if (rank == 0)
        for (int b = threadIdx.x; b < 256; b += blockDim.x)
            atomicAdd(&out[b], bins[b]);
}

The final cluster.sync() matters: a block's shared memory disappears when the block exits, so the cluster must not let a block finish while a partner may still be writing to it.

DPX, memory and the rest

The blog says Hopper's DPX instructions speed up dynamic programming algorithms, such as Smith-Waterman for genome alignment and Floyd-Warshall for route finding, by up to 7 times over A100. The CUDA 13 guide's feature table lists DPX as native on CC 9.0 and emulated with multiple instructions on other compute capabilities, which is a good example of a feature that is architecture-specific rather than inherited.

Other Hopper changes from the blog and tuning guide:

• 256 KB of combined L1 and shared memory per SM, with up to 228 KB as shared memory (227 KB per block, since CUDA reserves 1 KB per block). That is 39 percent more than A100's 164 KB.
• The first GPU with HBM3, per NVIDIA.
• Fourth-generation NVLink with 18 links and 900 GB/s per GPU, and the NVLink Switch System that the blog says can connect up to 256 GPUs.
• Second-generation MIG and NVIDIA Confidential Computing, which NVIDIA describes as the first native confidential computing GPU.

Hopper is also the first generation where NVIDIA introduced architecture-specific compiler targets. The wgmma warpgroup Tensor Core instructions in the PTX ISA require sm_90a, and a binary built only for sm_90a will not run on a Blackwell GPU. That is a deliberate break from the earlier rule that newer GPUs run older code.

Blackwell: two dies, FP4 and microscaling, NVLink 5 (2024)

NVIDIA announced Blackwell on 18 March 2024. Its press release and architecture page state that Blackwell GPUs pack 208 billion transistors, are built on a custom TSMC 4NP process, and use two reticle-limited dies joined by a 10 TB/s chip-to-chip link so that software sees one unified GPU. NVIDIA's Blackwell Ultra blog names that link NV-HBI (NVIDIA High-Bandwidth Interface) and says the two dies share a fully coherent L2.

Why two dies? A lithography scanner can only expose a field of limited size in one shot, the reticle limit. GH100 at 814 mm² was already close to it. To keep adding transistors on a similar process node, NVIDIA had to go to multiple dies and make the link between them fast enough that programmers do not have to care.

What NVIDIA has and has not published

NVIDIA's Blackwell Ultra blog (22 August 2025) gives a comparison table for Hopper, Blackwell and Blackwell Ultra GPUs. It is the most complete primary source we found:

For Blackwell Ultra the blog gives the SM layout: up to 160 SMs in eight GPCs across the two dies, 640 fifth-generation Tensor Cores (four per SM) and 256 KB of Tensor Memory (TMEM) per SM for warp-synchronous storage of intermediate results close to the Tensor Cores. NVIDIA notes that SM count and HBM capacity vary by product. For the original B200 we did not find an SM count in the NVIDIA sources we fetched, so we do not state one. The Blackwell tuning guide does give a 126 MB L2 for the GB200 GPU, 228 KB of shared memory per SM for CC 10.0, and the same 256 KB combined L1, texture and shared memory array as Hopper.

Second-generation Transformer Engine and microscaling formats

Blackwell's Tensor Cores add FP6 and FP4 alongside FP8. NVIDIA's architecture page says the second-generation Transformer Engine uses fine-grain micro-tensor scaling and supports community-defined microscaling formats. The idea of microscaling is simple: a 4-bit number has so little range that one scale factor per tensor is not enough, so each small block of values gets its own scale.

NVIDIA's own format, NVFP4, is described in the Blackwell Ultra blog as FP4 (E2M1) values with two levels of scaling: an FP8 E4M3 scale for each block of 16 values, plus one FP32 scale for the whole tensor.

Worked numbers: what NVFP4 costs in memory

NVFP4 block: 16 values x 4 bits + one 8-bit scale = 72 bits
72 bits / 16 values = 4.5 bits per value (ignoring the single per-tensor FP32 scale)

Weights of a 70-billion-parameter model (our arithmetic, weights only):
  FP16 / BF16   70e9 x 2 bytes      = 140 GB
  FP8           70e9 x 1 byte       =  70 GB
  NVFP4         70e9 x 4.5 / 8 B    = 39.4 GB

So NVFP4 cuts weight memory by about 3.6 times against BF16, not 4, because of the scales. Whether a given model keeps acceptable accuracy at FP4 is an empirical question for that model; NVIDIA's documents describe the format and the hardware, not your accuracy.

NVLink 5 and the NVL72 rack

Fifth-generation NVLink gives each Blackwell GPU 1.8 TB/s of bidirectional bandwidth, which the Ultra blog breaks down as 18 links at 100 GB/s. The GB200 NVL72 rack connects 72 Blackwell GPUs and 36 Grace CPUs in one liquid-cooled NVLink domain with 130 TB/s of NVLink bandwidth, and the Ultra blog says NVLink 5 supports up to 576 GPUs in a non-blocking fabric. Each GB200 superchip pairs two Blackwell GPUs with one Grace CPU over NVLink-C2C at 900 GB/s.

NVLink bandwidth per GPU, bidirectional (links x per-link rate, from NVIDIA papers)

Pascal P100    4 links x  40 GB/s =   160 GB/s
Volta V100     6 links x  50 GB/s =   300 GB/s
Ampere A100   12 links x  50 GB/s =   600 GB/s
Hopper H100   18 links x  50 GB/s =   900 GB/s
Blackwell     18 links x 100 GB/s = 1,800 GB/s
Rubin (announced)                  = 3,600 GB/s per GPU, per NVIDIA's January 2026 release

Other Blackwell engines

The architecture page and launch release list three more engines. A RAS engine for reliability, availability and serviceability diagnostics. A decompression engine for formats such as LZ4, Snappy and Deflate, aimed at database and analytics work. And confidential computing that NVIDIA describes as the first TEE-I/O capable GPU, with inline protection over NVLink. Blackwell Ultra adds doubled special function unit throughput for the exponentials in attention softmax, which the blog says gives up to 2 times faster attention-layer compute than Blackwell.

Blackwell has more than one compute capability

Under the one Blackwell name, NVIDIA ships parts with two different major compute capabilities. (Hopper and Ada already split data center and consumer parts across 9.0 and 8.9, but those carried different architecture names.) B200 and GB200 are CC 10.0, B300 and GB300 are 10.3, RTX 50 series and RTX PRO Blackwell cards are 12.0, and the GB10 in DGX Spark is 12.1. The CUDA 13 guide shows how different they are at SM level: CC 10.0 has 64 resident warps and 228 KB of shared memory per SM, while CC 12.x has 48 warps and 100 KB. The compute_100f family target covers 10.0, 10.3 and 10.7 but not 12.x. If you ship one container image for both a B200 cluster and RTX Blackwell workstations, you need code for both families.

Rubin: what NVIDIA has published so far

NVIDIA's press release of 5 January 2026 announced the Rubin platform as six chips: the Vera CPU, the Rubin GPU, the NVLink 6 switch, the ConnectX-9 SuperNIC, the BlueField-4 DPU and the Spectrum-6 Ethernet switch. For the GPU, the release states:

• a third-generation Transformer Engine with hardware-accelerated adaptive compression;
• 50 petaflops of NVFP4 compute for inference;
• sixth-generation NVLink with 3.6 TB/s per GPU, and 260 TB/s across a Vera Rubin NVL72 rack of 72 Rubin GPUs and 36 Vera CPUs.

The release said Rubin was in full production with partner products expected in the second half of 2026, and NVIDIA's Rubin platform page now describes it as ramping into full production. We did not find a Rubin architecture whitepaper, a per-SM description, a dense versus sparse breakdown of the 50 petaflops figure, or a compute capability for Rubin in the NVIDIA sources we fetched. Treat any such numbers you see elsewhere as unconfirmed until NVIDIA publishes them.

Across the generations: the numbers side by side

This table collects NVIDIA's stated figures for each generation's data center flagship. All are vendor peaks; "derived" marks our own arithmetic from NVIDIA's numbers.

FP32 peak grew about 13 times from K40 to H100 (5.04 to 67 TFLOPS), while low-precision matrix throughput grew far more, because each generation added a narrower format. FP64 was cut when a chip targeted deep learning only (M40) and kept when it targeted HPC (P100 onward).

Worked numbers: the roofline ridge point by generation

The roofline model says a kernel is limited by memory when its arithmetic intensity (FLOPs per byte moved from DRAM) is below the ridge point, peak FLOPs divided by peak bandwidth, and by compute above it. Using NVIDIA's dense peaks and bandwidths, checked in Python:

A concrete reading: a vector add in FP32 reads two floats and writes one, 12 bytes for 1 FLOP, an intensity of about 0.083. On any GPU in the table it is memory bound by a factor of hundreds. A large FP8 matrix multiply can have an intensity in the thousands, so it can approach the ridge. Everything between, which includes attention, normalisation and most of what real models run, is a fight to raise intensity with tiling, fusion and caching. That is why Ampere's async copy, Hopper's TMA and clusters, and Blackwell's TMEM exist.

Worked numbers: occupancy then and now

Occupancy is resident warps divided by the maximum resident warps per SM. It is limited by whichever of registers, shared memory, block slots or warp slots runs out first. Here is the same style of kernel on CC 1.0 and CC 9.0, using the limits in the tables above. These ignore register allocation granularity, so treat them as upper bounds and use cudaOccupancyMaxActiveBlocksPerMultiprocessor or Nsight Compute for exact figures.

CC 1.0 (G80): 8,192 registers, 24 warps, 8 blocks per SM
  256-thread blocks, 16 registers per thread
    by registers: 8,192 / (16 x 256) = 2 blocks
    by warps:     24 / 8             = 3 blocks
    resident:     2 blocks = 16 warps -> 16 / 24 = 67% occupancy
  same kernel at 20 registers per thread
    by registers: 8,192 / (20 x 256) = 1 block -> 8 / 24 = 33% occupancy

CC 9.0 (H100): 65,536 registers, 64 warps, 32 blocks, 228 KB shared per SM
  256-thread blocks, 64 registers, 48 KB shared memory per block (+1 KB reserved)
    by registers: 65,536 / (64 x 256) = 4 blocks
    by shared:    228 / 49            = 4 blocks
    resident:     4 blocks = 32 warps -> 32 / 64 = 50% occupancy
  cut registers to 32: registers allow 8 blocks, shared memory still caps at 4 -> 50%
  then cut shared memory to 32 KB: 228 / 33 = 6 blocks = 48 warps -> 75%

The lesson carries across every generation: find the binding limit before you tune anything else. On G80 four extra registers per thread halved occupancy. On H100 the same kernel can be bound by shared memory instead, and lowering registers does nothing until you fix that.

What stayed the same: SIMT, PTX and the compatibility contract

Through eleven architectures, three things barely changed. The warp is 32 threads. The programming model is SIMT with blocks that share a scratchpad. And CUDA programs are compiled to two forms: machine code (SASS) for specific SM versions, and PTX, a virtual instruction set that the driver can compile just in time for GPUs that did not exist when the program was built. The PTX ISA documentation we fetched is version 9.4 and covers instruction families from mma.sync (Volta onward) through wgmma.mma_async (Hopper) to tcgen05.mma (Blackwell).

A typical build that runs on several generations and stays forward compatible looks like this:

# Fat binary: SASS for each GPU you ship to, plus PTX for future GPUs.
# Follows the nvcc documentation; not executed on this Mac.
nvcc app.cu -o app \
  -gencode arch=compute_75,code=sm_75 \
  -gencode arch=compute_80,code=sm_80 \
  -gencode arch=compute_90,code=sm_90 \
  -gencode arch=compute_100,code=sm_100 \
  -gencode arch=compute_120,code=sm_120 \
  -gencode arch=compute_120,code=compute_120

# See what is inside the binary
cuobjdump -lelf app     # SASS images by SM version
cuobjdump -lptx app     # embedded PTX

# See what the machine has
nvidia-smi -q | grep -i -E "product name|driver version|cuda version"

The a and f targets are the exception to forward compatibility. NVIDIA's CUDA 13 guide is explicit that architecture-specific code runs only on its exact compute capability, and family-specific code only within its family. A library built with sm_90a kernels for H100 needs a separate build for Blackwell. This is one of many reasons GPU software has to be versioned and tracked carefully.

Why GPU history matters for your software supply chain

Every generation in this post arrived with software: a driver branch, a CUDA toolkit that added its compute capability, and libraries such as cuDNN, cuBLAS and NCCL compiled for it. Older generations leave the same way. The cuFFT 13.0 release notes list removed support for Maxwell, Pascal and Volta, and the CUDA 13 release notes say the NVIDIA driver is no longer bundled with the toolkit on Windows from CUDA 13.1 and on Linux from CUDA 13.4. NVIDIA publishes security bulletins for its GPU drivers and software on its product security page and recommends following their guidance on driver and package updates.

So the GPU stack in a real deployment is a dependency graph with many independent version numbers:

Here is where Safeguard fits, honestly. Safeguard does not design, program, tune or speed up GPUs, and nothing in this post depends on it. What Safeguard does is software supply chain work on the layers you build and ship. It generates and analyses SBOMs in SPDX and CycloneDX, scans repositories, container images (including their layers) and existing SBOM files (safeguard scan --dir ., safeguard scan --image <name>, safeguard scan --sbom <file>), matches components against vulnerability databases, ranks findings with reachability analysis, and shows which of your projects contain a given component and version. For a GPU team that means you can answer "which of our inference images still ship cuDNN version X" or "which projects pin an old CUDA base image" in one query, and gate new builds with policies. It does not scan host kernel drivers or GPU firmware outside the images you give it; those still come from NVIDIA's bulletins and your platform team.

Common mistakes

• Comparing sparse to dense. Since Ampere, many headline Tensor Core figures assume 2:4 sparsity. Halve them before comparing with a dense number, and do not assume your model is sparse.
• Comparing different precisions across generations. An FP8 or FP4 figure on a new GPU and an FP16 figure on an old one are not the same work. Compare like with like.
• Using preliminary launch numbers. Hopper's launch blog figures were preliminary and differ from the product page. Use the current product page or datasheet and say which you used.
• Treating peak as sustained. A peak is units times clock. Your kernel's throughput depends on its arithmetic intensity, occupancy and memory pattern. Measure with a profiler.
• Assuming warp-synchronous execution. Since Volta, threads in a warp can diverge and interleave. Use __syncwarp() and the _sync warp primitives with correct masks.
• Shipping only architecture-specific binaries. Code built only for sm_90a or sm_100a will not run on the next GPU. Include PTX or a family target if you need forward compatibility.
• Assuming a data center tuning carries to workstation parts. CC 8.0 and 8.6, or 10.0 and 12.0, have different shared memory and warp limits. Check occupancy on each target.
• Upgrading the toolkit without checking hardware support. CUDA 13 era libraries dropped Maxwell, Pascal and Volta. Inventory your fleet's compute capabilities before you upgrade base images.
• Forgetting the driver. A new CUDA runtime in an image needs a compatible host driver, and the driver is now installed separately from the toolkit. Track both versions.

FAQ

What are NVIDIA's GPU architectures in order?

From NVIDIA's own documents: Tesla (G80, 2006), Fermi (2009), Kepler, Maxwell (2014), Pascal (2016), Volta (2017), Turing (2018), Ampere (2020), Ada (consumer and inference parts, CC 8.9), Hopper (2022), Blackwell (2024) and Rubin (announced January 2026). The years are the dates on the NVIDIA whitepapers, blogs and press releases cited in this post.

What is the compute capability of H100, A100 and B200?

H100, H200 and GH200 are 9.0. A100 and A30 are 8.0. B200 and GB200 are 10.0, B300 and GB300 are 10.3, and RTX 50 series cards are 12.0. NVIDIA's compute capability page lists every current part.

What did Tensor Cores change?

They added units that perform small matrix multiply-accumulate operations in one step. On V100 that gave 125 TFLOPS of FP16 matrix math next to 15.7 TFLOPS of FP32, an eightfold gap. Each later generation added narrower formats: INT8 and INT4 on Turing, TF32 and BF16 on Ampere, FP8 on Hopper and FP4 on Blackwell.

What is the difference between Hopper and Blackwell?

Hopper (H100) is one 80 billion transistor die on TSMC 4N with FP8 Tensor Cores, 3.35 TB/s of HBM3 and 900 GB/s NVLink. Blackwell joins two reticle-limited dies (208 billion transistors, TSMC 4NP) into one GPU, adds FP6 and FP4 with microscaling, raises HBM3E bandwidth to 8 TB/s and NVLink to 1.8 TB/s per GPU. NVIDIA's table lists FP8 at 2 dense petaflops for Hopper and 5 for Blackwell.

Is the warp size still 32?

Yes. The CUDA 1.0 guide in 2007 and the CUDA 13 guide both give a warp size of 32 for every compute capability they cover.

What is a thread block cluster?

A Hopper feature that groups thread blocks so the hardware schedules them together and lets them access each other's shared memory. The portable maximum is 8 blocks, and H100 and B200 allow 16 if the kernel opts in.

What is NVFP4?

NVIDIA's 4-bit floating point format for Blackwell. Each value is FP4 (E2M1), each block of 16 values carries an FP8 E4M3 scale, and each tensor has an FP32 scale. That works out to about 4.5 bits per value.

Will my old CUDA binary run on a new GPU?

If it includes PTX, the driver can compile it for newer GPUs. If it contains only machine code for older SM versions, or only architecture-specific a targets, it will not. Check with cuobjdump, and rebuild with an embedded PTX target if needed.

Does Rubin have a published compute capability?

Not in the NVIDIA sources we read on 25 October 2026. NVIDIA's January 2026 release gives NVFP4 and NVLink figures, but we did not find a Rubin whitepaper or compute capability entry. Check NVIDIA's compute capability page for updates.

Can Safeguard make my CUDA code faster?

No. Safeguard does not program or tune GPUs. It helps you inventory and patch the software around them: container images, CUDA and cuDNN libraries in those images, Python packages and your repositories.

GPU software ages fast: each generation brings new driver branches, CUDA toolkits, cuDNN and NCCL builds and framework wheels, and old images keep running long after the hardware they targeted was retired. Those images are full of open source components with their own vulnerabilities, so the habits below apply to a GPU inference or training stack as much as to a web service, even though none of them touch the GPU itself.

More in this series

Other posts in our GPUs, CPUs, QPUs and lithography series:

• AMD GPU Architecture Explained: GCN, RDNA, CDNA, Infinity Fabric and HIP
• CPU vs GPU vs QPU Explained: How Each Works and Why They Coexist
• CUDA Programming Model Explained: Grids, Blocks, Warps and Memory
• GPU Memory and Interconnects: HBM, GDDR, NVLink, PCIe, CXL
• CUDA vs HIP vs SYCL vs OpenCL vs Metal vs Vulkan vs Triton Compared
• GPU Software Supply Chain: Drivers, CUDA, cuDNN, NCCL and Your SBOM
• How a QPU Works: Qubit Chips, Fridges, Control Pulses and Readout
• How CPUs Work: Pipelines, Out-of-Order, Caches, Branches, SIMD
• How EUV Lithography Machines Work: Tin Plasma, Mirrors and High-NA
• How GPUs Work: Graphics Pipeline, Shader Cores, Ray Tracing and Compute
• Inside an NVIDIA SM: Warp Schedulers, Registers, Tensor Cores
• Intel GPU Architecture: Xe Cores, XMX, Ponte Vecchio and oneAPI SYCL
• NVIDIA vs AMD vs Intel vs Apple GPUs: Architectures Compared
• Optimizing a CUDA Matrix Multiply Step by Step, From Naive to Tensor Cores
• Photolithography Explained: Contact Printing to DUV Immersion and Overlay
• Tile-Based GPUs Explained: Apple, Mali, Adreno, PowerVR and TBDR

Fix it without waiting for upstream

Whatever the vulnerability, most of the risk sits in two gaps. The first is the time between a fix being published and that fix actually running in your production builds. The second, and worse, is the time before any fix exists at all, while a maintainer works out a patch and a release schedule. Waiting through both gaps is a habit, not a requirement. Safeguard gives teams three ways to close them.

1. Autonomous remediation: pull requests, not tickets

When a fixed version exists, Safeguard's Griffin models open a tested, policy gated pull request that moves the affected dependency to the safe version. You can run it on one repository or in bulk across every project that carries the vulnerable version. Remediation covers eight categories, including dependencies, licenses, security posture and findings. Reachability analysis ranks the vulnerable code paths your application actually calls first, so the first pull requests you merge are the ones that matter most. Autonomous remediation ships switched off, with every category set to Manual, so your team decides how much runs on its own and what still needs a human approval.

2. Gold Open Source: the safe version in seconds, free

gold.safeguard.sh is free and needs no login. Search any CVE or package to see affected and fixed versions, CVSS, EPSS, CISA KEV status and a recommended safe version, or use the "Am I affected?" checker. The same data is available through a free read API, README status badges and RSS threat feeds. The safeguard-sh/gold-scan-action GitHub Action can fail a build on a known exploited vulnerability, on any known vulnerability, or never, so a version you have just removed cannot quietly come back in next week's merge.

3. Gold Registry: start from components that are already fixed

registry.safeguard.sh provides hardened, drop in replacements for upstream packages and container images. Every Gold artifact is held continuously to zero exploitable critical or high CVEs, zero malware and a completed license check, with SLSA L3 or higher provenance, a Sigstore signature and a public evidence page you can hand to an auditor. An artifact loses its Gold badge if it regresses. Teams can point their own package directory at the registry so hardened equivalents resolve without changing how developers install anything.

When upstream has no fix yet, or will not ship one, Safeguard first works to get the fix accepted upstream. If the issue is severe or upstream will not move, Safeguard builds a fork, fixes the vulnerability there and publishes a hardened Gold build, so you are not stuck on a maintainer's release schedule. You can also request a hardened, compatibility tested build of a specific package. This fork and patch flow is in early access today, so ask which packages it covers before you plan around it.

One honest limit: when the vulnerable code lives in a vendor appliance, a hosted service or an operating system runtime you do not build yourself, the fix still has to come from that vendor or distribution. The steps above apply to the open source dependencies and container images you build and ship, which is where most teams carry the largest and least tracked share of their exposure.

See how Safeguard would help

The first demo is a real Safeguard CLI session. You install it, run safeguard to open the interactive terminal and ask it to fix the lodash CVE in a repository. It reads the finding and the advisory, asks before any command that changes something, edits package.json to lodash 4.17.21, runs npm install and npm test with your approval, prints a summary and pushes nothing. The same flow applies to a Python serving repository or a Dockerfile that pins an old CUDA base image: the assistant proposes the change, you approve it, your tests decide.

Click through it at your own pace, or open it full screen.

The second demo sets up the Safeguard local runner on your own machine, approves it with its device code and then runs scans from the terminal. For GPU teams this is how you would scan a training or inference image where it is built, for example with safeguard scan --image <name>, and keep the results with your build.

Click through it at your own pace, or open it full screen.

The third demo connects Claude to the Safeguard MCP server with OAuth and shows how you choose which Safeguard tools Claude can use. Once connected, you can ask in plain language which of your projects contain a given library version, such as a particular cuDNN or NCCL build.

Click through it at your own pace, or open it full screen.

To try it on your own images and repositories, download the Safeguard CLI and follow the CLI documentation. For more background on the hardware side, see our posts on AI chip architecture, what an AI accelerator is and how it works and estimating GPU cost per million tokens when self-hosting.

Quick check

Sources

• NVIDIA: GeForce 8800 GPU Architecture Technical Brief (November 2006)
• NVIDIA: CUDA Programming Guide version 1.0 (June 2007)
• NVIDIA: Fermi Compute Architecture whitepaper v1.1 (2009)
• NVIDIA: Kepler GK110/GK210 architecture whitepaper
• NVIDIA Technical Blog: Maxwell, the most advanced CUDA GPU ever made (18 September 2014)
• NVIDIA: Maxwell tuning guide
• NVIDIA: Pascal GP100 architecture whitepaper (2016)
• NVIDIA: Pascal tuning guide
• NVIDIA: Volta GV100 architecture whitepaper (August 2017)
• NVIDIA: Turing architecture whitepaper (2018)
• NVIDIA: Ampere GA100 architecture whitepaper (2020)
• NVIDIA Technical Blog: NVIDIA Hopper architecture in depth (22 March 2022)
• NVIDIA: Hopper tuning guide
• NVIDIA: H100 Tensor Core GPU product page and specifications
• NVIDIA Newsroom: NVIDIA Blackwell platform arrives to power a new era of computing (18 March 2024)
• NVIDIA: Blackwell architecture page
• NVIDIA Technical Blog: Inside NVIDIA Blackwell Ultra (22 August 2025)
• NVIDIA: Blackwell tuning guide
• NVIDIA: GB200 NVL72 product page
• NVIDIA: NVLink and NVLink Switch page
• NVIDIA Newsroom: NVIDIA kicks off the next generation of AI with Rubin (5 January 2026)
• NVIDIA: Vera Rubin platform page
• NVIDIA: CUDA Programming Guide v13.4.2, compute capabilities appendix
• NVIDIA: CUDA C++ Programming Guide 12.0 (archive)
• NVIDIA: PTX ISA 9.4 documentation
• NVIDIA: CUDA GPU compute capability list
• NVIDIA: Legacy CUDA GPU compute capability list
• NVIDIA: CUDA Toolkit 13.4 Update 2 release notes
• NVIDIA: Product security and security bulletins

Last reviewed: 25 October 2026. NVIDIA updates product pages, compute capability tables and Rubin details often; check the linked NVIDIA documents for current figures before you rely on them.