NVIDIA H100
Scroll to dissect Hopper architecture
Hopper Architecture
NVIDIA H100
Scroll to dissect
Transistors
80 billion
Die area
814 mm2
Process
TSMC 4N
Active SMs
132
GH100 / 814 mm2
HBM3
HBM3
HBM3
HBM3
HBM3
OFF
GH100 / Hopper GPU die
host / NVLink / PCIe interface
MC1
MC2
MC3
MC4
MC5
MC6
MC7
MC8
MC9
MC10
GPC 1
9 TPC / 18 SMTPC 1SM1SM2
TPC 2SM1SM2
TPC 3SM1SM2
TPC 4SM1SM2
TPC 5SM1SM2
TPC 6SM1SM2
TPC 7SM1SM2
TPC 8SM1SM2
TPC 9SM1SM2
GPC 2
9 TPC / 18 SMTPC 1SM1SM2
TPC 2SM1SM2
TPC 3SM1SM2
TPC 4SM1SM2
TPC 5SM1SM2
TPC 6SM1SM2
TPC 7SM1SM2
TPC 8SM1SM2
TPC 9SM1SM2
GPC 3
8 TPC / 16 SMTPC 1SM1SM2
TPC 2SM1SM2
TPC 3SM1SM2
TPC 4SM1SM2
TPC 5SM1SM2
TPC 6SM1SM2
TPC 7SM1SM2
TPC 8SM1SM2
GPC 4
8 TPC / 16 SMTPC 1SM1SM2
TPC 2SM1SM2
TPC 3SM1SM2
TPC 4SM1SM2
TPC 5SM1SM2
TPC 6SM1SM2
TPC 7SM1SM2
TPC 8SM1SM2
50 MB shared L2 cache
GPC 5
8 TPC / 16 SMTPC 1SM1SM2
TPC 2SM1SM2
TPC 3SM1SM2
TPC 4SM1SM2
TPC 5SM1SM2
TPC 6SM1SM2
TPC 7SM1SM2
TPC 8SM1SM2
GPC 6
8 TPC / 16 SMTPC 1SM1SM2
TPC 2SM1SM2
TPC 3SM1SM2
TPC 4SM1SM2
TPC 5SM1SM2
TPC 6SM1SM2
TPC 7SM1SM2
TPC 8SM1SM2
GPC 7
8 TPC / 16 SMTPC 1SM1SM2
TPC 2SM1SM2
TPC 3SM1SM2
TPC 4SM1SM2
TPC 5SM1SM2
TPC 6SM1SM2
TPC 7SM1SM2
TPC 8SM1SM2
GPC 8
8 TPC / 16 SMTPC 1SM1SM2
TPC 2SM1SM2
TPC 3SM1SM2
TPC 4SM1SM2
TPC 5SM1SM2
TPC 6SM1SM2
TPC 7SM1SM2
TPC 8SM1SM2
H100 Streaming Multiprocessor
4 partitions / 128 FP32 / 64 INT32 / 64 FP64 / 4 Tensor
SM-wide L1 Instruction Cache
Partition 132 FP32 / 16 INT / 16 FP64
L0 Instruction Cache
Warp Scheduler
Dispatch Unit
Register File
16,384 x 32-bit
16,384 x 32-bit
FP64
INT32
FP32
Tensor Core
SFU
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
Partition 232 FP32 / 16 INT / 16 FP64
L0 Instruction Cache
Warp Scheduler
Dispatch Unit
Register File
16,384 x 32-bit
16,384 x 32-bit
FP64
INT32
FP32
Tensor Core
SFU
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
Partition 332 FP32 / 16 INT / 16 FP64
L0 Instruction Cache
Warp Scheduler
Dispatch Unit
Register File
16,384 x 32-bit
16,384 x 32-bit
FP64
INT32
FP32
Tensor Core
SFU
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
Partition 432 FP32 / 16 INT / 16 FP64
L0 Instruction Cache
Warp Scheduler
Dispatch Unit
Register File
16,384 x 32-bit
16,384 x 32-bit
FP64
INT32
FP32
Tensor Core
SFU
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
LD/ST
256 KB L1 Data Cache / Shared Memory
TEX
TEX
TEX
TEX
CUDA / Triton to H100 hardware
GPU > 8 GPC > 66 TPC > 132 SM
PyTorch
z = x + yTriton launch
add_kernel[grid](...)Program
pid = tl.program_id(0)Data
load -> add -> storesoftware / launch model
PyTorch operationtensor expression
Launch gridmany programs / blocks
Program / CUDA blockone data tile
Warp32 CUDA threads
CUDA thread / Triton lanecompiler-mapped values
mapping
operationkernel launch
griddistributed across GPU
program / blockscheduled to one SM
warpissued by scheduler
thread / laneexecution + register state
hardware target / behavior
Compiled GPU kernelmachine instructions
H100 GPU132 enabled SMs
One SMup to 32 resident blocks
Warp scheduler4 per SM
Execution lanesLD/ST + FP32 for vector add
memory hierarchy
lower latency / smaller capacity to higher latency / larger capacityRegistersthread state
65,536 32-bit registers per SM01
Shared memory / L1one SM
256 KB per SM combined L1 data cache and shared memory02
L2 cacheall SMs
50 MB03
HBM3device memory
80 GB HBM3 on H100 SXM / 3.35 TB/s04
Hopper architecture
NVIDIA H100
Scroll to dissect the package, compute hierarchy, one SM, and the CUDA/Triton execution model.
Transistors
80 billion
Process
TSMC 4N
SMs
132
Layer 1 / Package
GH100 die with HBM3
HBM1
HBM2
HBM3
HBM4
HBM5
SITE
GH100814 mm2
Memory
80 GB
Type
HBM3
Bandwidth
3.35 TB/s
Layer 2 / Physical compute
GPU → GPC → TPC → SM
GPC 1
9 TPC
18 SM
GPC 2
9 TPC
18 SM
GPC 3
8 TPC
16 SM
GPC 4
8 TPC
16 SM
GPC 5
8 TPC
16 SM
GPC 6
8 TPC
16 SM
GPC 7
8 TPC
16 SM
GPC 8
8 TPC
16 SM
1 GPC
8 or 9 TPC
1 TPC
2 SM
1 SM
execution
Shared L2 cache50 MB
H100 exposes 8 GPCs, 66 TPCs, and 132 SMs. CUDA schedules blocks to SMs, not directly to GPCs or TPCs.
Layer 3 / One streaming multiprocessor
H100 SM: four scheduler partitions
SM-wide L1 instruction cache
Partition 11 scheduler
Warp scheduler + dispatch
Register file: 16,384 × 32-bit
32 FP32
16 INT32
16 FP64
1 Tensor
Partition 21 scheduler
Warp scheduler + dispatch
Register file: 16,384 × 32-bit
32 FP32
16 INT32
16 FP64
1 Tensor
Partition 31 scheduler
Warp scheduler + dispatch
Register file: 16,384 × 32-bit
32 FP32
16 INT32
16 FP64
1 Tensor
Partition 41 scheduler
Warp scheduler + dispatch
Register file: 16,384 × 32-bit
32 FP32
16 INT32
16 FP64
1 Tensor
256 KB combined L1 data cache / shared memory
Per SM: 128 FP32, 64 INT32, 64 FP64 and 4 Tensor Cores.
Layer 4 / Programming model
Software → hardware mapping
PyTorch
z = x + y
launches or compiles a GPU kernel
Grid
many programs / blocks
distributed across 132 SMs
Program / block
one data tile
resides on one SM
Warp
32 CUDA threads
issued by a warp scheduler
Thread / lane
load → add → store
uses registers and execution lanes
Registers
thread
Shared/L1
SM
L2
GPU
HBM3
device