Inside the H100
From package to tensor core: a scroll-driven journey through NVIDIA's Hopper architecture, one silicon layer at a time.
The DGX H100 Node
8× H100 SXM5 · 4× NVSwitch · 900 GB/s bidirectional per GPU
DGX H100 and the eight-GPU HGX H100 baseboard connect eight through four over fourth-generation NVLink. Each GPU exposes up to 900 GB/s of aggregate bidirectional NVLink bandwidth and the node contains 640 GB of HBM3. Workloads may use tensor, pipeline, expert, or data parallelism; the ring later in this guide is one illustrative collective, not a promise about every training step.
This view is specifically the eight-GPU DGX H100 / HGX H100 NVSwitch configuration.
Depending on the parallel strategy, workloads may use all-reduce, reduce-scatter, all-gather, all-to-all, or point-to-point traffic.
The node contains 640 GB of HBM3; 7.2 TB/s is summed bidirectional endpoint bandwidth, while fabric bisection is 3.6 TB/s.
The Package
H100 SXM5 · Up to 700 W (configurable) · 80 GB HBM3 · 3.35 TB/s
The H100 SXM5 module is where the die, HBM stacks, power delivery, and interconnect come together on a single board. The shipping SXM5 configuration exposes and ten 512-bit memory controllers for 80 GB of memory at 3.35 TB/s. The sixth HBM site in the full GH100 floorplan is useful context when you zoom into the die. provides up to 900 GB/s of aggregate bidirectional GPU-to-fabric bandwidth, 1.5× that of A100. handles host connectivity at 128 GB/s bidirectional.
The SXM5 module joins GH100 silicon, HBM stacks, NVLink, PCIe, and power delivery into one accelerator module.
Shipping H100 SXM5 exposes five active HBM3 stacks; the full GH100 floorplan explains the sixth site.
H100 SXM5 is configurable up to a 700 W maximum TDP, so operating power and clocks depend on the deployed configuration.
The Die
GH100 Floorplan · 814mm² · TSMC 4N
Under the heat spreader sits an 814 mm² monolithic GH100 die. NVIDIA's full-chip organization contains , 60 MB of L2 cache, twelve 512-bit memory controllers, and 144 SMs. H100 SXM5 exposes a documented subset: 132 SMs, , and . The public material establishes those counts, but it does not disclose why each individual unit is disabled or the exact physical routing shown by this teaching model.
The full GH100 die arranges eight GPCs around shared L2 slices and memory-controller edges.
NVIDIA names the GigaThread Engine in the public diagram, but does not publish the detailed internal placement policy modeled here.
The view separates the full 144-SM GH100 organization from the documented 132-SM H100 SXM5 configuration.
Inside a GPC
Full-GH100 maximum · 9 TPCs · 18 SMs
In the full GH100 organization, each Graphics Processing Cluster contains and each TPC contains . That maximum organization yields 18 SMs per full GPC; H100 SXM5 enables 66 TPCs across the die, so the SXM5 device is not eight identical full GPCs. Graphics blocks in the scene are illustrative context rather than a published per-GPC inventory. NVIDIA states that only two TPCs in H100 SXM5 and H100 PCIe are graphics-capable. Thread block clusters are what make the GPC a level a programmer schedules against: the blocks of a cluster are co-scheduled within one GPC and can reach each other's shared memory over the SM-to-SM network .
A full-GH100 GPC contains nine TPCs with two SMs each; enabled H100 SXM5 GPCs need not all contain that maximum.
The GPC is a documented hierarchy level, while the exact internal wires and graphics-block inventory in this scene are illustrative.
One full GPC accounts for 18 SMs, 2,304 CUDA cores, and 72 Tensor Cores in the GH100 floorplan.
The Streaming Multiprocessor
4 Processing Blocks · 128 FP32 Cores · 4 Tensor Cores
The SM is where CUDA thread blocks execute. NVIDIA’s public Hopper diagram divides an SM into four processing blocks and publishes plus four Tensor Cores per SM. The scene places , , and execution resources into a readable quadrant layout, but it is not a literal route map. The 256 KB L1/texture/shared-memory pool provides cache and software-managed staging, with up to 228 KB configurable as shared memory per SM.
// eligible == next instruction's operands ready && its pipe free
warp = select_from(resident_warps, where: eligible);
instruction = fetch(warp.pc);
dispatch(instruction, warp.active_lanes);
warp.pc = next_pc(warp.pc); // SASS instructions are 16 bytes
A Hopper SM contains four processing blocks with schedulers and published CUDA-lane and Tensor Core resources.
The 256 KB combined L1/texture/shared pool supports caching and software-managed staging for tiled kernels.
TMA and asynchronous barriers let a kernel overlap future-tile movement with eligible matrix work.
The Tensor Core
D = A × B + C · FP8 to FP64
Tensor Cores execute matrix multiply-accumulate operations on tile fragments. Hopper's fourth-generation Tensor Cores add FP8 support and asynchronous warp-group MMA (WGMMA) . A WGMMA operation uses four contiguous warps: B is described in , A may come from registers or a shared-memory descriptor, and accumulators remain in registers. Published TFLOPS figures are device peak specifications; sustained application throughput depends on data movement, instruction mix, and utilization.
fence_generic_proxy_writes_to_shared(); // only if B was written by st.shared
fence_wgmma_operands();
issue_async_mma(D_regs, A_regs_or_smem, B_smem);
commit_wgmma_group();
wait_for_wgmma_group(0);
Tensor Cores execute matrix multiply-accumulate on fragments: D = A × B + C.
Hopper adds FP8 Tensor Core paths with FP16 or FP32 accumulator forms, while Transformer Engine manages FP8 scaling and casts in supported software flows.
Four warps can cooperate on one larger operation, with shared-memory descriptors reducing register pressure.
Execution Fundamentals
SIMT · Warps · Occupancy · Memory · Pipeline
Understanding GPU performance requires thinking about five interlocking concepts. A scheduler issues a warp instruction to the currently active lanes of 32 threads ; since Volta, independent thread scheduling also maintains per-thread execution state. The thread hierarchy groups warps into blocks and, on Hopper, blocks into thread block clusters before grids — a cluster's blocks are co-scheduled on one GPC and can address each other's shared memory directly (Chapter 03). Resource limits bound occupancy , while data locality, dependencies, and eligible warps determine whether latency is hidden in a particular kernel.
A scheduler issues a warp instruction to active lanes; independent thread scheduling retains per-thread execution state within the warp.
Registers, shared memory, block size, and launch limits bound how many warps can be resident; dependencies determine which resident warps are eligible to issue.
Coalescing, cache level, and stalls usually determine whether peak math units can stay fed.
Inside a CUDA Core
Warp Issue · Scalar Arithmetic · Dependencies
"CUDA core" is NVIDIA's lane-count term for scalar arithmetic capability, not a published gate-level block diagram. At the programming-model level, a warp scheduler selects an eligible warp, reads its source operands, issues an instruction to the appropriate execution pipeline, and later makes the destination available to dependent instructions. An FMA computes A × B + C with one final rounding. NVIDIA does not publish the exact GH100 operand-routing, result-bus, register-file-cache, or physical pipeline implementation shown in simplified teaching diagrams.
warp = select_eligible_warp();
operands = read_sources(warp.instruction);
result = fused_multiply_add(operands); // one final rounding
make_destination_ready(result);
A CUDA core is a pipelined lane for scalar arithmetic, fed by warp-level issue rather than individual thread control.
Register operands and dependencies are architectural concepts; GH100 physical collectors and routing are not publicly documented.
Dependent instructions become eligible only after their inputs are ready, but the exact internal result path is intentionally not asserted here.
See how kernels map to PTX and SASS in the PTX Playground
A Conceptual FMA Datapath
Conventional arithmetic model · Not disclosed GH100 circuitry
IEEE 754 defines the observable FMA result: compute A × B + C as if with unbounded precision, then round once. A conventional hardware implementation may recode the multiplier, compress partial products with carry-save adders, align the addend, perform a final carry-propagate addition, normalize, and round. The diagram uses that textbook sequence to explain the arithmetic; it does not claim that GH100 uses a particular Booth radix, Wallace topology, carry-lookahead design, number of physical stages, or cycle latency. FP32 has 23 stored fraction bits and 24 bits of precision.
exact_product = exact(A) * exact(B);
exact_sum = exact_product + exact(C);
result = round_once(exact_sum, mode=RNE);
A conventional multiplier may recode and compress partial products; this is a teaching model, not a GH100 disclosure.
A conventional fused datapath aligns and combines C before the single final rounding required by IEEE 754.
Normalize and round stages turn raw bits back into an IEEE 754 result with one final rounding.
The Transistor
TSMC 4N customized for NVIDIA · 80 billion transistors
At the bottom of the stack is a FinFET-family manufacturing process. NVIDIA identifies GH100 as an 80-billion-transistor, 814 mm² design manufactured on TSMC 4N, a process customized for NVIDIA. Public NVIDIA and TSMC material does not disclose GH100's gate length, contacted-gate pitch, fin pitch, SRAM cell, detailed metal stack, or material choices. The scene therefore explains generic CMOS and FinFET concepts only; its geometry and layer spacing are not a physical reconstruction or measurement of GH100.
FinFETs provide generic physical context; dimensions in this scene are not measured or disclosed GH100 geometry.
CMOS inverters and NAND gates are the cells from which schedulers, caches, and datapaths are built.
BEOL layers route signals and power, but GH100 layer counts, pitches, and material choices are not public.
Scale positions are illustrative teaching magnitudes
You have descended roughly eight orders of magnitude.
Click a graduation to climb back up
H100 SXM5 Quick Reference
- Architecture
- Hopper (GH100)
- Process
- TSMC 4N
- Transistors
- 80 billion
- Die Size
- 814 mm²
- SMs (H100 SXM5)
- 132
- CUDA Cores
- 16,896
- Tensor Cores
- 528 (4th gen)
- FP8 Tensor
- 1,979 TFLOPS dense / 3,958 with 2:4 sparsity
- FP16/BF16 Tensor
- 989 TFLOPS dense / 1,979 with 2:4 sparsity
- TF32 Tensor
- 495 TFLOPS dense / 989 with 1:2 sparsity
- FP64 Tensor
- 67 TFLOPS (scalar FP64 ~34 TFLOPS)
- FP32
- 67 TFLOPS
- Memory
- 80 GB HBM3
- Memory BW
- 3.35 TB/s
- L2 Cache
- 50 MB
- Interconnect
- NVLink 4.0 · 900 GB/s aggregate bidirectional
- PCIe
- Gen5 x16 · 128 GB/s aggregate bidirectional
- TDP
- Up to 700 W (configurable)