NVIDIA Hopper Architecture

Updated 5 July 2025

NVIDIA Hopper Architecture is a next-gen GPU platform that integrates advanced tensor cores, enhanced memory systems, and secure computing features to accelerate AI and HPC workloads.
It introduces innovative features like FP8 tensor cores, asynchronous WGMMA instructions, and the Tensor Memory Accelerator to boost throughput and efficiency for large-scale data processing.
Unified memory integration via NVLink and built-in confidential computing ensure seamless performance and high data security in scientific and industrial applications.

The NVIDIA Hopper Architecture is a GPU platform designed to accelerate AI, high-performance computing (HPC), and data-intensive workloads. As the successor to the NVIDIA Ampere series, Hopper introduces advances in hardware functionality, memory hierarchy, confidential computing, and programming models. Key features include fourth-generation tensor cores with FP8 support, the Tensor Memory Accelerator (TMA), advanced distributed shared memory mechanisms, and secure computing features integrated at the hardware and system-software levels. Hopper models such as the H100 and GH200 (Grace Hopper Superchip) have been extensively benchmarked in academic and industrial research for their effects on throughput, programmability, and reliability across scientific and AI applications.

1. Microarchitecture and Core Features

The Hopper microarchitecture (SM90) integrates several technological innovations targeting both AI and general-purpose workloads:

Tensor Cores with FP8 Support: Hopper’s fourth-generation tensor cores support a variety of precisions, including FP8 (E5M2, E4M3), FP16, BF16, TF32, and INT8. FP8 support is specifically geared to the demands of LLM training and inference, often doubling throughput over FP16 for sufficiently large matrix operations. These tensor cores execute both traditional synchronous matrix–multiply–accumulate (MMA) instructions and the new asynchronous warp-group-level instructions known as wgmma, introducing greater concurrency and efficiency, especially for large matrices (2402.13499, 2501.12084).
Asynchronous Warpgroup Matrix-Multiply-Accumulate (WGMMA): The wgmma instructions cater to 4-warp (128-thread) warpgroups and allow overlapping memory movement with computation. Wgmma supports a wide range of matrix shapes and enables direct operations on larger tiles—translating to high utilization on large AI workloads (2501.12084).
Tensor Memory Accelerator (TMA): The TMA is a dedicated asynchronous copy engine for orchestrating bulk data transfers between global and shared memory, allowing copy operations (such as prefetching weights or activation tiles) to overlap with computation. TMA exposes interfaces for asynchronous operations within CUDA and frameworks like CUTLASS (2312.11918, 2501.12084).
Dynamic Programming Acceleration (DPX Instructions): Hopper introduces hardware instructions to support dynamic programming primitives (e.g., min/max, fused addition, and ReLU). These were shown to accelerate key use cases (such as the Smith–Waterman algorithm) by up to 13× for 16-bit arithmetic, compared to prior generation GPUs relying on software emulation (2402.13499, 2501.12084).
Distributed Shared Memory (DSM): DSM enables low-latency, cluster-local SM-to-SM communication via CUDA-exposed mechanisms, reducing the latency of inter-block or inter-SM data movement by as much as 32% over L2 cache, and supporting data access patterns not efficiently serviced by global memory (2402.13499).
Partitioned L2 and Enhanced Bandwidth: Hopper increases L2 cache to 50 MB (compared to Ampere’s 40 MB), reconfigures L2 as a partitioned cache, and adds improvements to memory throughput—L2 bandwidth is 2.2–2.6× higher than on previous architectures in benchmarks, and HBM bandwidths approach 3.6 TB/s in Grace Hopper systems (2501.12084).

2. Unified Memory and Heterogeneous Integration

Hopper is featured prominently in the GH200 "Grace Hopper" Superchip, which integrates a 72-core ARM-based Grace CPU and a Hopper GPU via NVLink-C2C interconnect:

Unified Memory Architecture (UMA): The CPU and GPU share a single, cache-coherent address space across LPDDR5X for the CPU and HBM for the GPU. This integration allows both processors to access all system memory without explicit data copies or buffer handoffs. Applications allocate memory with standard allocation calls—and those pointers are valid for BLAS calls on either CPU or GPU (2404.13195, 2408.11556, 2501.00279).
NVLink C2C Interconnect: Each direction supports up to 450 GB/s. While intra-node HBM access can approach 4 TB/s, efficient data placement remains crucial; traversing NVLink twice (such as DDR–DDR transfers) halves the effective bandwidth (2408.11556).
Fine-Grained Page Migration and NUMA Affinity: Techniques such as the Device First-Use data migration policy (an editor's term, inspired by OpenMP First-Touch) ensure that BLAS data moves to HBM at first use and remains resident, minimizing transfer overhead for iterative and block-structured scientific codes (2404.13195, 2501.00279).
Programming Flexibility: UMA enables tools that intercept and offload BLAS/LAPACK routines dynamically (e.g., via LD_PRELOAD and dynamic binary instrumentation), providing automatic GPU acceleration for legacy codes without code changes, so long as access patterns allow reusing migrated data (2404.13195, 2501.00279).

3. Programming Models and Kernel Design

Hopper’s innovation in hardware requires corresponding advancements in software and programming paradigms:

Task-oriented, Data-centric Models: Newer task-programming environments such as Cypress exploit Hopper’s asynchronous TMA and Tensor Core units by decoupling high-level sequential semantics from the low-level hardware orchestration. Programmers specify computations as sequential tasks on logical tensors, and a mapping specification directs the Cypress compiler to insert asynchronous copies, pipelining, and warpgroup specialization automatically. This bridging enables performance within 0.88–1.06× that of expert-tuned cuBLAS GEMM and 0.80–0.98× of state-of-the-art Flash Attention kernels (2504.07004).
Optimizing Hybrid Kernels: Practical attention mechanisms (e.g., FlashAttention-2) exploit kernel fusion, overlapping TMA-driven tile copy with double-buffered, pipelined GEMM calls (via WGMMA), and thread-level softmax computation. Tile sizes are tuned (e.g., 64×128, 128×64) to balance register pressure, shared memory, and throughput; excessive tile size can cause register spills counteracting the benefit of larger accumulators. Experimental results show up to 50% FLOP/s improvement over Ampere-based implementations (2312.11918).
Evaluation Across Programming Models: In stencil and classical HPC workloads, CUDA remains the dominant model for absolute throughput; however, directive-based models like OpenACC and OpenMP are increasingly competitive, especially when augmented with asynchronous execution strategies. On Hopper, optimized OpenACC can outperform OpenMP by ~33%, and tuned CUDA still outpaces even the best OpenACC by around 2.1×, highlighting the persistent benefits of vendor-specific tuning (2404.04441).

4. Numerical Properties and Portability

Hopper’s matrix accelerators exhibit unique (and often undocumented) numerical behaviors compared to past NVIDIA (V100, A100) and competing AMD matrix engines:

Block-FMA Width and Extra Precision: The Hopper H100 employs a fused-multiply-add (FMA) accumulation width of at least 16, with two or more internal guard bits during accumulation, in contrast to A100’s width of 8. This expands accumulation fidelity but can introduce non-trivial differences in output for numerically sensitive routines. Rounding modes may deviate from standard RTN–TE (round-to-nearest ties-to-even), at times using "truncate" (chopping). This has practical ramifications: porting algorithms, such as mixed-precision iterative refinement, may yield divergent results—sometimes by orders of magnitude—across architectures (2403.00232).
Feature-Targeted Testing: Empirically designed tests (e.g., T_blk_fma_width, T₃_bits_fin_rnd) are essential to identify platform-specific behavior—developers must validate their computation’s reproducibility if cross-GPU portability is a requirement (2403.00232).

5. Confidential Computing and Security Model

NVIDIA Hopper is the first GPU platform to extend confidential computing from CPUs to the GPU domain, introducing hardware-enforced GPU Confidential Computing (GPU-CC):

Trusted Execution Environments (TEE): Hopper GPUs create secure enclaves, the Confidential Processing Region (CPR), that isolate code and data via encryption/integrity checks, secure boot, and attestation protocols. Secure channels are established for all communication, with hardware-supported key derivation and access controls (2507.02770, 2409.03992).
System Components: Security management is handled via the Foundation Security Processor (FSP), GPU System Processor (GSP), and SEC2 engine (all RISC-V cores), orchestrating secure boot, attestation, device authentication, and confidential memory management. Copy Engines enforce encrypted DMA transfers in and out of the CPR using keys established via SPDM-derived protocols (2507.02770).
Usability and Integration: End-user applications run unmodified—integration is managed by the system stack (drivers, CUDA runtime, user-mode, and kernel-mode drivers modify communication channels "under the hood") (2507.02770). Even when GPU-CC is enabled, performance overhead for representative LLM inference tasks is minimal (generally <7%, and often near zero for large models with long sequences); the dominant penalty arises from encrypted PCIe transfers, not per-token computation (2409.03992).
Security Caveats: The architecture’s complexity and proprietary nature present analysis challenges. A small fraction of register fields may remain accessible in CC mode. While critical paths are encrypted, some metadata in RPC remains in plaintext, creating minimal side-channel risk. Timing channels and signature key exposure in memory scrubbing have been observed; all reported to NVIDIA PSIRT (2507.02770).

6. Application Benchmarks and Scientific Impact

Benchmarks across domains demonstrate practical benefits:

AI Workloads: FP8 tensor cores deliver up to 2× throughput over FP16 for large LLMs, particularly when leveraging new libraries such as Transformer Engine and the dedicated wgmma instruction set. Application-level benchmarks in image classification, object detection (notably on transformer-based detectors), and LLM inference show substantial throughput increases over the A100 and even distributed multi-A100 setups for certain workloads (2410.16487, 2501.12084).
Scientific Computing and BLAS-Heavy Codes: Automatic BLAS offloading via unified memory (with data migration policies like Device First-Use) enables up to 3× speedup for iterative, matrix-heavy codes in quantum physics. Rewriting legacy codes becomes less burdensome, with tools intercepting BLAS symbols and directing operations to the GPU transparently (2404.13195, 2501.00279).
Real-Time Streaming: In radio astronomy applications such as binary pulsar detection, Hopper-based GPU implementations (Pulscan on AstroAccelerate) yield speedups of 70×–250× over CPU, easily sustaining real-time throughput and full-device pipeline execution (2406.15186).
Energy Efficiency: Hopper achieves higher compute-per-watt efficiency in selected AI and DP workloads, reducing overall time to solution and enhancing the energy profile for exascale applications (2501.12084, 2404.04441).

7. Significance and Future Trajectories

The Hopper architecture embodies several key trends in contemporary accelerated computing: integration of heterogeneous processing elements (CPU+GPU), flexible precision tensor operations, hardware-level security models, and fine-grained memory management. While the feature set delivers strong performance gains (20–58% over Ampere in common kernels, 2–3× in unified memory BLAS operations), effective utilization requires adaptive kernel design and awareness of architectural subtleties—especially in numerical accuracy and code portability. The move toward transparent, dynamic offloading and security by default—without user-level change—signals a continued evolution toward sustainable, performant, and secure scientific computing and AI infrastructure.

Table: Key Novel Features and Their Performance Impact

Feature	Description	Performance/Flexibility Impact
FP8 Tensor Cores	Low precision (E5M2/E4M3) arithmetic in tensor cores	Up to 2× throughput vs. FP16 on LLMs
WGMMA Instructions	Async, warpgroup matrix operations	95%+ of peak, better pipeline utilization
TMA (Tensor Memory Accelerator)	Async copy between global/shared memory	Overlapped copy/compute, hides latency
Distributed Shared Memory (DSM)	Cluster-local SM-to-SM communication	32% lower latency vs L2, efficient inter-SM
Unified Memory via NVLink C2C	CPU+GPU single address space, cache-coherent	Enables seamless auto-offload, efficient data
Confidential Computing (GPU-CC/TEE)	Secure enclaves, key management, attestation	<7% overhead for LLM inference
Device First-Use Policy (Editor term)	NUMA-aware auto-migration of data to GPU memory	Enables iterative reuse, minim. migration

This architecture, with its comprehensive combination of hardware, memory, and security features, establishes a new baseline for future AI, HPC, and scientific applications, setting both a high-performance and high-integrity standard for GPU-accelerated platforms.