---
title: NVIDIA GB10 Grace Blackwell Architecture
url: https://www.emergentmind.com/topics/nvidia-gb10-grace-blackwell
type: topic
---

# NVIDIA GB10 Grace Blackwell Architecture

NVIDIA GB10, codenamed "Grace Blackwell," is a unified CPU-GPU architecture noted for its dual-die design, advanced low-precision tensor core support, high-bandwidth HBM3e memory system, and novel execution and memory subsystem optimizations. GB10 targets large-scale machine learning, high-performance computing (HPC), and exascale workloads by introducing significant microarchitectural improvements over previous NVIDIA Hopper (H100) generation hardware, as substantiated through empirical microbenchmarking and performance modeling [2512.02189][2507.10789][2601.16032].

## 1. Microarchitectural Overview

NVIDIA GB10 features a dual-chip (dual-die) design: two identical GPU “slices,” each integrating 74 SMs (148 SMs total, organized into 8 GPCs) are interconnected via NV-HBI in a unified address space spanning 192 GB HBM3e, served by 8 memory stacks (4 per die). Each die contains four L2 partitions, and the design leverages 208 billion transistors. GB10 is fabricated at TSMC 3 nm, positioning it at the forefront of density scaling and process node efficiency [2512.02189][2601.16032].

Critical hardware components and their architectural context include:

- **5th-Generation Tensor Cores:** Support scalar and matrix operations in FP64, FP32, FP16, BF16, FP8 (E4M3), FP6 (E3M2/E2M3), FP4 (E2M1), INT8, and INT4. The `tcgen05.mma` instruction provides per-thread matrix-multiply-accumulate (MMA), decoupled from warp-level synchronization. Single-instruction latency for all supported tile sizes is approximately 11 cycles, compared to up to 128 cycles for prior Hopper wgmma, significantly increasing throughput for both dense and sparse GEMM patterns [2512.02189][2507.10789].
- **Tensor Memory (TMEM):** Each SM integrates 256 KB TMEM, implemented as a 2D array (512 columns × 128 lanes of 32-bit cells), providing 16 TB/s read and 8 TB/s write bandwidth per SM, additive to L1/shared memory performance.
- **Decompression Engine (DE):** Hardware supports multiple codecs (LZ4, Snappy, Zstandard, GZIP, Cascaded, Bitcomp, ANS), achieving input throughput of 173 GB/s (uncompressed) and output throughput of 539 GB/s (<1.3 ms for 100 MB workloads) [2512.02189].

A typical dataflow involves HBM3e global memory → L2 cache → L1/shared memory → SM, with accelerated transfers to and from TMEM (`tcgen05.cp`, `tcgen05.ld/st`). The DE allows on-the-fly tensor decompression with minimal latency overhead.

## 2. Memory Hierarchy and Subsystem Innovations

The GB10 memory hierarchy is optimized to mitigate bottlenecks in deep learning and scientific applications:

- **Global Memory:** 192 GB HBM3e, with measured peak bandwidth up to 8.2 TB/s (read) and 1.6 TB/s (write) for GDDR7 configurations, and LPDDR5X with ~301 GB/s raw bandwidth for SoC variants [2512.02189][2507.10789][2601.16032].
- **L2 Cache:** Unified 65 MB L2 for GPU, 24 MiB for SoC variant. L2 hit latency is measured at 358 cycles (light load) rising to 508 cycles under saturation (e.g., 45 MB active footprint). L2 behavior dominates streaming-attention and tile-based matrix workloads [2507.10789][2601.16032].
- **L1/Texture and Shared Memory:** 128 KB per SM (GPU), up to 192 KiB per SM (SoC), with hit latency ≈30–40 cycles. GB10 makes L1/share configurable via CUDA attributes, with a tunable default shared memory [2507.10789].
- **Tensor Memory (TMEM):** Reduced memory access latency for TMEM-misses vs. global memory by 58% (from 1000 cycles to 420 cycles). L1/arbitration-style TMEM hits occur at ≈50 cycles [2512.02189].
- **Impact on Algorithm Design:** The increased TMEM and L2 sizes raise the optimal tile sizes for matrix work from 32×32 (Hopper) to 64×64 (GB10), maximizing TMEM usage and reducing global memory writes up to 12 TB/s per SM in attention-style kernels [2512.02189][2601.16032].

## 3. Tensor Core Throughput and Mixed-Precision Capabilities

GB10’s 5th-generation tensor cores drive unprecedented mixed-precision throughput and energy efficiency:

| Precision | Theoretical Peak Throughput (TFLOPS) | Relative to H200 (where supported)  |
|-----------|--------------------------------------|-------------------------------------|
| FP4       | 7702.5                              | N/A (new on Blackwell)              |
| FP6       | 5134.8                              | N/A (new on Blackwell)              |
| FP8       | 3851.4                              | 1.27×                               |
| FP16      | 1929.2                              | 1.27×                               |
| BF16      | 1926.8                              | 1.27×                               |
| FP32      | 481.2                               | 1.27×                               |
| FP64      | 44.8                                | —                                   |

- **Mixed-precision throughput gains**: 1.56× over H200 (Hopper) [2512.02189].
- **Energy efficiency gain:** 42% (η_B200 / η_H200 = 1.42) [2512.02189].
- **Microbenchmark results:** FP4/FP6/FP8 microbenchmark throughput exceeds 11 TFLOP/s per SM; FP4 and FP6 provide ≈20% lower power and ≈15% higher throughput vs. FP8 for inference-oriented kernels [2507.10789].
- **Latency:** Single-instruction latency of 11 cycles for m64×64×16 FP16 operations versus 32 cycles (Hopper wgmma) [2512.02189].

FP4 and FP6 are natively supported in hardware, with `QMMA` or `OMMA` kernel mapping and are recommended for inference where modest losses in model accuracy or perplexity are acceptable. The architecture supports flexible PTX decomposition, decoupling data movement and computation at fine granularity.

## 4. Empirical Workload Evaluations

Microbenchmark and end-to-end workload studies across GB10 yield the following key findings [2512.02189][2507.10789][2601.16032]:

| Workload                     | B200 (GB10)         | H200 (Hopper)        | Speedup Factor     |
|------------------------------|---------------------|----------------------|--------------------|
| Dense GEMM (FP64, ≥16k²)     | 36.3 TFLOPS         | 18.9 TFLOPS          | 1.92×              |
| STREAM Triad (Mem BW, >64GB) | 7.48 TB/s           | 4.38 TB/s            | 1.71×              |
| Sparse GEMV + DE             | 5.04 GFLOPS         | 3.20 GFLOPS          | 1.58×              |
| Transformer Inference (FP8, B32, S=2048) | 78,400 tok/s         | 49,200 tok/s         | 1.59×              |
| Mixed-Precision Training (GPT-1.3B) | 14,397 tok/s         | 9,240 tok/s          | 1.56×              |
| Energy Efficiency (GPT-1.3B) | 22.2 tok/s/W        | 15.6 tok/s/W          | 1.42×             |

Transformer models (e.g., Mistral-7B, Mixtral-8×7B) realize 1.52–1.59× lower inference latency at batch sizes of 1 and 32. B200’s support for FP4 yields up to 2.5× inference speedup with a perplexity increase of +8.2% at extreme quantization (FP4 vs. FP8: ΔPPL +8.2%) [2512.02189]. A plausible implication is that, in deployment scenarios where model quality loss is tolerable, FP4 can be preferentially used in non-critical layers to maximize speed.

## 5. Software and Programming and Kernel Optimization

GB10 exposes new PTX instructions for TMEM (`tcgen05.cp`, `tcgen05.ld/st`), enabling asynchronous global↔TMEM copies and flexible tile decomposition, as well as DE primitives for data decompression. Tile sizes optimized at 64×64 (FP8) or T=80 (CuTile) maximize TMEM and shared memory use [2512.02189][2601.16032].

For attention kernels and GEMM:

- **Tile-based Decomposition:** Optimal for both dense and sparse matrix multiplication. Developers are advised to double-buffer TMEM tiles and aggressively overlap data movement with computation to approach 95–99% of peak throughput.
- **DE Kernel Integration:** Enables asynchronous decompression of large weight tensors during compute, with >85% pipeline efficiency for appropriate chunk sizes [2512.02189].
- **Precision Selection:** Use FP8 in critical paths (e.g., QKV projection in transformers), apply FP4 otherwise, while limiting FP4 deployment to layers with acceptable perplexity drift (ΔPPL < 3%) [2512.02189].
- **Warp Scheduler Tuning:** Expose at least 9 independent operations per thread and use ~32 warps per block to saturate tensor core throughput [2507.10789].

Compiler recommendations include scheduling to maintain steady ILP, generate mma.sync instructions for FP4/FP6, and avoid shared memory bank conflicts (stride ≥ 4) that can induce high latency [2507.10789].

## 6. Memory Access Optimization: Attention Kernels and Sawtooth Wavefront

High-performance attention (e.g., FlashAttention) kernels on GB10 are dominated by L2 cache behavior due to the large KV working set characteristic of LLMs [2601.16032]. Standard cyclic KV tile access patterns result in high reuse distances and L2 conflict/capacity misses when working sets exceed L2 capacity (e.g., S ≈ 80k, KV size ≈ 20 MiB for D=64), sharply degrading throughput.

Sawtooth Wavefront Reordering, an access scheduling technique alternating forward and backward KV tile iteration between cooperating thread arrays (CTAs), halves the mean reuse distance. This pattern empirically reduces L2 misses by 50–67% and boosts throughput by up to 60% on GB10 for both CUDA and CuTile implementations.

Table: GB10 FlashAttention Sawtooth Optimization Experimental Results (S=32k, D=64)

| Configuration          | Cyclic L2 Miss (×10⁶) | Sawtooth L2 Miss (×10⁶) | Reduction | Orig TFLOPS | Opt TFLOPS | Speedup |
|------------------------|----------------------|--------------------------|----------|-------------|------------|---------|
| CUDA, B=1              | 108.0                | 54.0                     | 50%      | 1.3         | 2.4        | 1.85×   |
| CuTile non-causal      | 370                  | 120                      | 67%      | 61          | 69         | 1.13×   |
| CuTile causal          | 380                  | 150                      | 60%      | 41          | 66         | 1.61×   |

The benefit is maximized when sequence length is large and working set size exceeds L2 capacity. Slight divergence overhead (~5%) may occur when working set fits L2. The sawtooth ordering approach is architecture-targeted, generalizing to other tiled kernels and potentially to multi-GPU/NUMA schemes [2601.16032].

## 7. Comparative Analysis and Architectural Significance

Relative to prior Hopper architecture (H100/H200), Grace Blackwell (GB10) demonstrates:

- Substantially increased mixed-precision throughput (up to 1.56×) and energy efficiency (up to 1.42×).
- Lower memory miss latency due to enlarged TMEM and L2; optimal tile sizes shift upward, and kernel-level data movement patterns must be retuned for maximum performance [2512.02189].
- A fundamental trade-off: reduced L1 capacity and external GDDR7 bandwidth in some SKUs, but larger L2 cache, higher TMEM bandwidth, and support for new ultra-low-precision datatypes (FP4, FP6) [2507.10789].
- Microbenchmark evidence that GB10’s more conservative and consistent warp scheduling, larger register file, and aggressive support for low-precision datatypes translate to smoother throughput scaling as ILP increases [2507.10789].
- System-level innovations (coherent GPU/CPU SoC, unified memory, improved cooperative thread array control) support a range of workload patterns not efficiently handled on previous architectures [2601.16032].

A plausible implication is that future algorithmic and software optimizations for GB10-class hardware will need to operate at larger tile-and-batch granularities, leverage dynamic precision selection, and exploit advanced memory scheduling.

---

**References**:  
- [2512.02189] Microbenchmarking NVIDIA's Blackwell Architecture: An in-depth Architectural Analysis  
- [2507.10789] Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks  
- [2601.16032] Sawtooth Wavefront Reordering: Enhanced CuTile FlashAttention on NVIDIA GB10

Source: https://www.emergentmind.com/topics/nvidia-gb10-grace-blackwell