Get the most performance from Grace Hopper

Akshay Subramaniam, Devtech Compute | GTC 2025

目录

• NVIDIA Grace Hopper Superchip
• NVIDIA GB200 Superchip
• Grace Hopper Superchip 架构

• CUDA工具以充分利用Grace Hopper

  • 内存分配器

    • 仅设备内存管理 (cudaMalloc
    • CUDA统一内存 (cudaMallocManaged
    • 系统分配内存 (malloc/mmap

  • Grace Hopper 统一内存

• Grace Hopper 优化

  • 统一内存的自动迁移 (CPU -> GPU
  • CUDA内存调优API
  • 用于显式控制的内存调节 API

  • Grace Hopper 的缓存一致性

    • 跨 NVLink C2C 的原子操作
    • 用于生产者/消费者模式的并发队列

  • Grace Hopper 优化性能

    • NVLink-C2C 性能

• 为 Grace-Hopper 做好准备

  • 只需重编译和运行
  • 加速现有应用

• Grace Hopper 带来的开发者效率提升

• 在 Grace Hopper 上优化应用程序

  • 关键优化技术

  • 监督式微调 (Supervised Fine Tuning, SFT

    • 混合专家模型 (Mixture of Experts - MoE
    • 技术：激活检查点 (Activation Checkpointing
    • 技术：异步激活卸载 (Asynchronous activation offloading
    • 在 Megatron-LM 中异步卸载到 CPU 内存
    • 性能分析与评测 (Profiling and Performance

  • 可微流体动力学 (Differentiable Fluid Dynamics

    • XLB
    • LBM: 数据与计算
    • 使用 XLB 进行核外计算 (Out of Core Computing
    • 核外算法 (Out-of-Core Algorithm
    • 使用 XLB 进行核外计算的详细流程
    • 异步内存传输
    • 使用核外内存的开销成本

  • 数据库查询

    • x86 与 Grace Hopper 架构对比
    • 数据库查询性能

• 结论

  • 关键要点 (Key Takeaways

• CUDA 开发者会议

NVIDIA Grace Hopper Superchip

Soul是全新的NVLink-C2C CPU↔GPU互联技术。

• 内存一致性：易于使用

  • 所有线程——GPU和CPU——访问系统内存：C++ new、malloc、mmaped文件、原子操作等。
  • 快速自动页迁移。
  • 线程缓存对等内存 → 减少迁移。

• 高带宽：900 GB/s (与原生NVLink 4相同)

  • GPU读取或写入本地/对等LPDDR5X内存时达到峰值带宽。

• 低延迟：GPU→HBM延迟

  • GPU读取或写入LPDDR5X内存时具有“HBM3延迟”。

对于系统中的所有线程，内存的体验在行为、延迟和带宽方面都符合预期。

NVIDIA GB200 Superchip

这是Grace Hopper的增强版（on steroids）。

• 下一代超级芯片

  • 2个Blackwell GPU + 1个Grace CPU
  • 第二代Transformer引擎，支持FP4/FP6 Tensor Cores
  • 解压引擎，速度高达800 GB/s
  • 第五代NVLink
  • 多节点NVLink

• 与Grace Hopper相同的编程模型

  • 为Grace Hopper开发的代码在GB200上同样有效。

更多信息请参考 LLM Inference Performance and Optimization on NVIDIA GB200 NVL72 [S72503]。

Grace Hopper Superchip 架构

GPU可以以CPU内存的速度访问CPU内存。

该架构图展示了NVIDIA Grace Hopper Superchip的组件和连接：

• Grace CPU 连接高达 512 GB/s 的 LPDDR5X 内存 (480 GB)。
• Hopper GPU 连接高达 4.9 TB/s 的 HBM3 (96GB) 或 HBM3e (144 GB) 内存。
• CPU和GPU通过 NVLINK C2C 连接，带宽为 900 GB/s。
• 系统通过 高速I/O (4x 16x PCIe-5, 512 GB/s) 和 NVLINK NETWORK (18x NVLINK 4, 900 GB/s) 与外部连接，支持最多256个GPU。
• 整个系统实现了硬件一致性。

CUDA工具以充分利用Grace Hopper

内存分配器

仅设备内存管理 (cudaMalloc)

在 PCIe Hopper 和 Grace Hopper 架构上，cudaMalloc 的行为是相同的。通过 cudaMalloc 分配的GPU内存，CPU无法直接访问。CPU指针 ptr 无法解引用以访问GPU内存中的数据。

cudaMalloc 的行为继续与x86平台保持一致，CPU无法访问其分配的内存。如代码所示，数据必须通过 cudaMemcpyHostToDevice 从CPU内存显式复制到GPU内存，并在计算后通过 cudaMemcpyDeviceToHost 复制回来。

CUDA统一内存 (cudaMallocManaged)

在 PCIe Hopper 和 Grace Hopper 架构上，使用 cudaMallocManaged 分配的内存，系统通过页错误（Page fault）和迁移（migrate）机制工作。内存会在首次被访问时（on first touch）被物理分配和放置。CPU和GPU都可以访问指向该内存的指针，系统会自动处理数据迁移。

cudaMallocManaged 使得代码在x86和Grace Hopper之间具有可移植性，通过直接访问或移动内存页实现。如代码所示，使用 cudaMallocManaged 后，无需显式的 cudaMemcpy 调用。CPU可以直接初始化数据，然后GPU内核可以直接使用该数据指针。

系统分配内存 (malloc/mmap)

• PCIe Hopper：GPU无法访问CPU内存，反之亦然。CPU通过malloc()分配的内存对GPU不可见，GPU内存对CPU也不可见。（注：异构内存管理HMM可以通过页错误和迁移来模拟此行为）。
• Grace Hopper：CPU和GPU可以通过C2C互联直接访问对方的内存。malloc()分配的内存，无论是CPU首次接触还是GPU首次接触，都可以在两者之间共享和直接访问。

在Grace Hopper上，使用标准 malloc 分配的内存，其物理页面在首次被访问时分配。如代码所示，当GPU内核 kernel<<<1, 128>>>(data, data2) 首次访问 data 和 data2 时，这些内存页将被物理分配。

GPU内核执行完毕后，CPU可以直接访问由GPU写入的数据。如代码中的 for 循环所示，CPU可以直接从 data 和 data2 中读取GPU的计算结果，无需任何显式的数据拷贝。

Grace Hopper 统一内存

Grace Hopper的统一内存通过 地址转换服务 (Address Translation Service, ATS) 实现。系统页表 (System Page Table) 能够将CPU malloc() 分配的内存地址或GPU内存地址进行转换。
- CPU物理内存中的页面A (Page A)，可以被CPU本地访问 (CPU-resident access)，也可以被GPU远程访问 (Remote accesses)。
- GPU物理内存中的页面B (Page B)，可以被GPU本地访问 (GPU-resident access)。
这种机制使得CPU和GPU可以无缝访问彼此的内存空间。

Grace Hopper 优化

统一内存的自动迁移 (CPU -> GPU)

当一个位于CPU物理内存中的页面（Page A）被GPU频繁访问时，系统会自动将该页面迁移到GPU物理内存中，以提高性能。不频繁的访问则会继续通过NVLINK C2C进行远程访问。

CUDA内存调优API

在Grace Hopper上，CUDA内存调优API可用于 cudaMallocManaged 和系统分配内存 (malloc)。

• Grace Hopper 批量传输 (Bulk Transfers)

  • 使用 cudaMemPrefetchAsync API可以将CPU内存中的数据预取到GPU内存中。这是一种显式的批量数据传输，适用于 cudaMallocManaged 或 malloc 分配的、在首次接触时分配的内存。

• Grace Hopper 直接访问 (Direct Access)

  • 使用 cudaMemAdvise API并附带 cudaPreferredLocation 参数，可以向驱动程序建议数据的首选位置。这使得数据可以保留在CPU内存中，同时允许GPU通过NVLINK C2C进行高性能的直接访问。

用于显式控制的内存调节 API

Grace Hopper 平台提供了用于显式内存控制的调节 API。开发者可以精细地管理数据在 CPU 内存和 GPU 内存之间的位置和移动，以实现性能最优化。

• 内存分配与建议:
  • cudaMemAdvise_v2 和 cudaMemPrefetchAsync_v2 等 API 允许开发者指定内存的首选位置（Host 或 Device）并异步预取数据。
  • 示例代码分析:
    1. data1 被分配在 CPU（Grace）内存中。代码通过 cudaMemAdviseSetPreferredLocation 明确建议其位置为主机端。
    2. data2 被分配在 GPU（Hopper）内存中。代码通过 cudaMemAdviseSetPreferredLocation 明确建议其位置为设备端，并使用 cudaMemPrefetchAsync_v2 将数据异步预取到 GPU。
    3. 下图展示了数据从 CPU 内存（蓝色箭头）和 GPU 内存（绿色箭头）流向 Hopper GPU 进行计算的过程。

• 统一内存访问:
  • 在数据布局完成后，GPU Kernel 和 CPU 代码都可以直接访问这些内存区域。
  • Grace Hopper 的 NVLink C2C 互连技术支持 CPU 和 GPU 之间双向直接访问，无需显式的内存拷贝。
  • 示例代码分析:
    1. kernel<<<1, 128>>>(data, data2) 在 GPU 上启动，可以直接访问位于 CPU 内存的 data1 和位于 GPU 内存的 data2。
    2. GPU Kernel 执行完毕后，CPU 上的 for 循环可以直接打印 data1 和 data2 的内容，展示了 CPU 对 GPU 内存的直接访问能力。

Grace Hopper 的缓存一致性

Grace Hopper 通过 NVLink C2C 实现了缓存一致性访问，极大地简化了 CPU 和 GPU 之间的编程模型。

• CPU 访问 GPU 内存: 当 CPU 访问 GPU 的 HBM3 内存时，它会通过 "检出 (checking out)" 缓存行的方式，将数据拉取到 CPU 的 L3 缓存中。这确保了 CPU 能够高效地访问 GPU 数据。
• GPU 访问 CPU 内存: 当 GPU 访问 CPU 的 LPDDR5X 内存时，它可以实现一致性访问，而无需 "检出" 缓存行。数据可以直接流向 Hopper 的 SMs，同时保持与 CPU 缓存的同步。

跨 NVLink C2C 的原子操作

缓存一致性使得 CPU 和 GPU 能够对共享内存地址执行原子操作。

• 示例代码:
  • 代码展示了 GPU Kernel 和 CPU 代码如何通过 cuda::atomic_ref 对一个位于统一内存中的整数进行原子性的 fetch_add 操作。
  • cuda::thread_scope_system 作用域确保了原子操作在整个系统（包括 CPU 和 GPU）中都是可见和一致的。

• 执行流程:
  1. 在 main 函数中，CPU 初始化一个系统范围的原子变量 atomic_data 的值为 0。
  2. 随后，启动 GPU Kernel，该 Kernel 会对 atomic_data 执行原子增加操作。
  3. cudaDeviceSynchronize() 确保所有 GPU 操作完成后，CPU 可以安全地读取最终结果。这展示了 CPU 与 GPU 在共享数据结构上的紧密协作。

用于生产者/消费者模式的并发队列

缓存一致性非常适用于实现高效的生产者/消费者模式，其中 CPU 和 GPU 可以分别扮演生产者和消费者的角色，操作一个共享的并发工作队列。

• 队列结构:

  • 队列由一系列工作项组成，分为已完成、进行中和待处理三部分。
  • Producer available work pointer 指向下一个可供生产者添加工作的位置。
  • Consumer reserved work pointer 指向下一个可供消费者提取工作的位置。

• 消费者（GPU）:

  • GPU 上的设备函数 load_work_queue 使用 cuda::atomic_ref 和 memory_order_acquire 内存顺序，安全地从共享队列中加载工作索引，确保能观察到生产者发布的所有工作。

• 生产者（CPU）:
  • CPU 负责生成工作项并将其添加到队列中。
  • 通过对 a_total_work_items 执行 fetch_add 原子操作，并使用 cuda::memory_order_acq_rel 内存顺序，CPU 可以原子性地更新工作指针，并确保其写入的数据对 GPU 消费者可见。

• 消费者（GPU）拉取工作:
  • GPU 上的 Kernel 线程块通过对工作索引 _work_idx 执行 fetch_add 原子操作来预留新的工作项。
  • 这种方式允许多个 GPU 线程块并发地、无锁地从队列中获取任务，实现了高效的任务分发。

Grace Hopper 优化性能

NVLink-C2C 性能

Grace Hopper 的 NVLink-C2C 互连显著提升了 CPU 和 GPU 之间的数据传输带宽。

• cudaMemcpy 性能对比:
  • Pinned Memory: Grace Hopper 的 Host-to-Device (H2D) 带宽达到 381.7 GB/s，远高于 x86+H100 系统的 55.67 GB/s。
  • Pageable Memory: 最显著的提升在于可分页内存。传统 x86 系统中，可分页内存的传输性能极差（H2D 仅 9.5 GB/s）。而在 Grace Hopper 上，由于统一内存架构和高带宽互连，可分页内存的传输性能（H2D 381 GB/s）几乎与固定内存持平。这极大地简化了内存管理并提升了易用性。

为 Grace-Hopper 做好准备

只需重编译和运行

• 现有应用程序通常无需修改代码。
• 只需为 ARM Neoverse-V2 (Grace) 和 sm_90 (Hopper) 架构重新编译应用程序。
• 即可从平台无处不在的更高带宽中获益。

加速现有应用

• 移植应用程序比以往任何时候都更容易。
• 支持广泛的编程模型和语言。
• 硬件缓存一致性简化了编程。
• 即使是部分移植的应用，也能通过 Grace CPU 和 C2C 获得整体加速。
• 提供丰富的 NVIDIA 和第三方工具。
• 平衡的架构减少了阿姆达尔定律带来的瓶颈。

Grace Hopper 带来的开发者效率提升

Grace Hopper 旨在加速通往加速计算的路径，显著提升开发者效率和应用性能。

• 性能与工程投入曲线对比:
  • x86+H100 (深青色线): 从纯 CPU 开始，当引入 GPU 时，由于 PCIe 瓶颈，性能可能暂时下降，这会降低开发者的积极性。需要大量的工程努力才能跨越瓶颈，实现性能的显著提升。
  • Grace Hopper (绿色线): 从纯 CPU 开始，由于 C2C 的高带宽和低延迟，可以立即获得增量式的性能加速。这鼓励开发者进一步优化，通过 "生产力加速" 阶段，更快地达到 "为正确的目标使用正确的芯片" 的高效状态。

在 Grace Hopper 上优化应用程序

关键优化技术

Grace Hopper 提供了两种主要的 CPU-GPU 交互和优化模式：

1. 异步批量 cudaMemcpy: 将 CPU 内存视为 GPU 内存的扩展。通过 NVLink C2C 进行异步卸载，适用于需要批量传输大量数据的场景。
2. 直接内存访问 (Direct Memory Access): 通过 NVLink C2C 对 CPU 内存进行显式的高带宽和缓存访问。适用于需要 GPU 对 CPU 内存进行细粒度、低延迟访问的场景。

监督式微调 (Supervised Fine Tuning, SFT)

本节内容感谢 Matthias Langer (DevTech Compute APAC)。

监督式微调是调整大型语言模型（LLM）以适应特定任务或领域知识的关键步骤。

• 流程:
  1. 预训练 (Pre-training): 使用网页规模的庞大数据集训练一个通用的预训练 LLM。
  2. 微调 (Fine-tuning): 使用特定领域的数据或自定义知识库，对预训练的 LLM 进行进一步训练，使其适应特定任务。
  3. 结果: 得到一个针对特定领域进行优化的微调后 LLM。

Grace Hopper 的统一内存架构为处理大规模模型和数据集的微调过程提供了便利，简化了数据在 CPU 和 GPU 之间的流动。

混合专家模型 (Mixture of Experts - MoE)

• 每个 Transformer 层由于包含多个前馈（Feed Forward）专家层，因此是内存密集型的。
• 持续预训练（Continual Pre-training）涉及使用新的数据集来增强已训练的模型。

下图展示了在 LLM 模型中，一个 Transformer 层的结构，该层采用了混合专家（MoE）架构。输入经过层归一化（Layer Norm）和注意力（Attention）模块后，再次进行层归一化，然后通过 MoE 路由（MoE Routing）机制，将计算分配给 N 个专家网络中的一个或几个。最后，将专家们的输出进行加权求和（Weighted Sum），并与原始输入进行残差连接。

技术：激活检查点 (Activation Checkpointing)

该技术旨在通过重计算激活值来更有效地利用内存。

• 中间张量（Intermediate tensors）具有很高的内存需求。
• 重计算（Re-computation）的内存开销较低，但会导致冗余计算和更高的能耗。

在反向传播（Backward pass）过程中，为了节省内存，不会存储所有前向传播（Forward pass）的中间激活值，而是在需要时重新计算它们。

技术：异步激活卸载 (Asynchronous activation offloading)

该技术将激活值异步卸载到 CPU 内存中，以减少 GPU 内存的占用。

前向传播过程：
在模型的前向传播过程中，每一层计算产生的中间张量（Intermediate Tensors）会被生成。

随着前向传播的进行，产生的中间张量（特别是较小的张量）被异步地从 GPU 内存卸载（Offload）到 CPU 内存中。

这个卸载过程在模型的每一层都会发生，从而持续释放 GPU 内存。

反向传播过程：
在反向传播开始时，当计算需要前向传播过程中产生的中间张量时，这些张量会从 CPU 内存中被召回（Recall）到 GPU 内存。

随着反向传播的继续，每一层所需的中间张量都会按需从 CPU 内存中召回，直到反向传播完成。

在 Megatron-LM 中异步卸载到 CPU 内存

• 在 Megatron-LM 中的实现会暴露用于卸载优化的运行时参数。

代码示例展示了如何使用 cpu_offload_context 来管理激活值的卸载。

self.cpu_offload_context = CPUOffloadContext(args)

with self.cpu_offload_context as ctx:
    for index in range(self.num_layers):
        layer = self.get_layer(index)
        hidden_states = layer(hidden_states, **forward_kwargs)
        hidden_states = ctx(hidden_states, self.num_layers)

相关配置参数包括：--cpu-offload-activations, --cpu-offload-async, --cpu-offload-size-threshold 67108864。

下表分析了在一个约 108B 参数的 MoE MegatronLM 模型（配置为 8 个 H100 GPU 系统，2 个专家，15 层，batch size 为 8）中临时张量的分布情况。可以看出，尽管大于 1000MB 的大型张量数量不多，但它们占据了总内存的绝大部分（54.1% + 12.0%）。

性能分析与评测 (Profiling and Performance)

• 与 x86 平台相比，在 Grace Hopper 平台上的激活卸载效率更高。

下图的性能分析显示：
* 在 Grace Hopper 平台上，卸载操作（D2H 和 H2D）可以被计算（Compute）有效隐藏。
* 在 x86 平台上，卸载操作主要受数据传输（data transfer）的限制，成为性能瓶颈。

在 8 个 NVLink 连接的 GPU 上对 108B MoE 模型进行监督式微调的性能对比显示，Grace Hopper 平台上的卸载（offload）方案比重计算（recompute）方案快 25%。

可微流体动力学 (Differentiable Fluid Dynamics)

感谢：Łukasz Wawrzyniak, Oliver Hennigh

XLB

来自 Autodesk Research 的加速格点玻尔兹曼框架 (Accelerated Lattice Boltzmann Framework)

• XLB 是一个由 Autodesk Research 开源的、在 Python 中实现的可微、可扩展的计算流体动力学（CFD）库。

特点:
* 基于格点玻尔兹曼方法（Lattice Boltzmann Method）。
* 支持多种后端：
* JAX by Google，使用数组编程。
* Warp by NVIDIA，使用核函数编程。
* Neon。

• 可在多核 CPU、GPU 和 TPU 上进行大规模并行计算。
• 对机器学习（ML）友好，允许与深度学习（DL）模型集成。
• 支持核外计算（out-of-core computations），即流式计算（streaming）。

LBM: 数据与计算

• 格点玻尔兹曼方法 (Lattice Boltzmann Method - LBM)
  • 基于模板（Stencil）的方法/局部通信。
  • 对 GPU 友好，通常受内存带宽限制。
  • 仿真规模受限于 GPU 内存。
  • 通常为碰撞（collision）和迁移（streaming）步骤使用单个融合核函数（fused kernel）。

LBM 的计算过程主要分为两个步骤：

1. 碰撞步骤 (Collision Step): 在每个格点上，粒子分布函数根据碰撞模型进行更新，以模拟流体粒子间的相互作用。
2. 迁移步骤 (Streaming Step): 更新后的粒子分布函数沿着离散速度方向传播到相邻的格点。

使用 XLB 进行核外计算 (Out of Core Computing)

利用 GH/GB 架构

• 仿真数据以瓦片（tiles）的形式存储在 CPU 内存中。
• 瓦片具有期望的填充厚度（padding thickness），这允许每个传输到 GPU 的瓦片可以执行多个时间步（timestep）的计算。
• 排队（Queueing）和多流（multiple streams）技术允许异步内存复制。

下图展示了该过程：来自完整仿真域（Lattice Boltzmann Lid Driven Cavity）的瓦片被存储在主机内存（Host Memory）中，然后被传输到设备内存（Device Memory）中进行多时间步的更新。

核外算法 (Out-of-Core Algorithm)

使用统一内存和双缓冲技术

该技术实现了异步重叠的内存传输。

• 核外 LBM (Out-of-Core LBM)

  • 仿真规模受限于 GPU 内存？可以使用主机 CPU 内存！
  • 可以实现 5-10 倍更大的仿真规模！
  • 使用缓冲和异步内存传输，以最小化性能影响。
    • 只要内存传输足够快。

• 适用于其他模板类型的方法

  • 显式有限体积法 (Explicit Finite Volume)
  • 时域有限差分法 (Finite Difference Time Domain)
  • 细胞内粒子法 (Particle in Cell)
  • 直接模拟蒙特卡洛法 (Direct Simulation Monte Carlo)

下图展示了使用双缓冲（Buffer A, Buffer B）技术，将计算（Compute）和 CPU 到 GPU 的数据传输进行重叠，从而提高效率。

使用 XLB 进行核外计算的详细流程

利用 GH/GB 架构

下图详细描述了核外计算的 7 个步骤：

1. 分区 (Partitioning): 整个仿真域被分割成瓦片（tiles），存储在 CPU RAM 中。
2. 数据移动: 瓦片从 CPU 内存移动到 GPU 内存。
3. 时间步计算 #1: 在 GPU 上执行第一个时间步的计算。
4. 时间步计算 #2: 继续在 GPU 上执行后续时间步的计算。
5. 时间步计算 #N: 瓦片具有参数化的光环（halos），允许在 GPU 上执行多个时间步。
6. 光环更新 (Halo Update): 光环区域在 GPU 上进行局部更新。
7. 异步复制: 排队和多流技术允许异步内存复制，将更新后的瓦片写回 CPU RAM，同时加载新的瓦片。

异步内存传输

下图展示了在 Grace Hopper (GH) 上使用核外 (Out of Core) 内存的格子玻尔兹曼方法 (LBM) 的典型性能分析剖面图。其开销极小，约等于两次内存复制的成本。

从图中可以看出：
- 每个数据块更新包含多个时间步长：计算内核（蓝色和紫色条）与内存传输操作交错进行。
- CPU与GPU之间的异步内存来回传输：内存复制操作 Memcpy DtoH (Device to Host) 和 Memcpy HtoD (Host to Device) 与 CUDA 内核的执行是重叠的，这表明了异步执行的特性，有效隐藏了数据传输的延迟。

使用核外内存的开销成本

使用核外内存虽然能够处理更大规模的问题，但也会带来一定的性能开销。

上图展示了格子玻尔兹曼方法（BGK D3Q19）在纯 GPU 内存和使用核外内存两种配置下的最大问题规模和性能对比。

• 最大尺寸对比 (LBM Max Size)：

  • GPU 内存: 可处理 1.1 十亿（Billion）单元格。
  • 核外内存: 可处理 5.8 十亿单元格。

• 性能对比 (LBM Performance)：

  • GPU 内存: 性能为 21027.4 MLLUPS (Million Lattice-Site Updates Per Second)。
  • 核外内存: 性能为 14569.5 MLLUPS。

要点总结：

• 使用核外内存可以运行规模显著增大的模拟，例如在顶盖驱动流（Lid Driven Cavity, LDC）问题中，规模可扩大约 5倍。
• 使用核外内存会产生开销。在LDC问题上，观察到的开销约为 31%。
• 这项工作仍在进行中，理论最优（“光速”）开销有望降低到 20% 以下。

数据库查询

本节内容感谢 Rui Lan, Eyal Soha, Nikolay Sakharnykh (DevTech Compute NALA) 的贡献。

x86 与 Grace Hopper 架构对比

在处理数据库查询等需要大量CPU-GPU数据交互的应用时，传统架构的瓶颈和 Grace Hopper 的优势得以体现。

• 在 x86 平台上，PCIe 成为系统瓶颈。
• Grace Hopper 凭借其 NVLink C2C 的高带宽，能够高效支持数据库查询。

如上图所示：
- 左侧 (x86+H100)：CPU 与 Hopper GPU 通过 PCIe 总线连接。当输入表位于CPU内存，而哈希表/输出位于GPU内存时，PCIe 的带宽（~64GB/s）成为性能瓶颈。
- 右侧 (NVIDIA Grace Hopper Superchip)：Grace CPU 与 Hopper GPU 通过带宽高-达 900 GB/s 的 NVLink C2C 互连。这使得 GPU 可以高速访问 CPU 的 LPDDR5X 内存（高达 512 GB/s），CPU 也可以高速访问 GPU 的 HBM3 内存（高达 4.9 TB/s），从而消除了数据传输瓶颈。

数据库查询性能

下图展示了在不同数据库查询任务中，Grace Hopper 相对于 x86 平台的性能优势。该测试采用的技术是“高带宽直接CPU访问”。

• 在聚合查询 (Aggregation Query)、连接查询+压缩 (Join Query + compression)、连接查询 (Join Query) 和复杂连接查询 (Complex Join Query) 等多种场景下，Grace Hopper 的吞吐量（GB/s）都显著高于 x86+A100 和 x86+H100 的组合。
• 这再次证明了 PCIe 在 x86 平台上的瓶颈效应，以及 Grace Hopper 的 NVLink C2C 在数据密集型应用中的巨大价值。

结论

关键要点 (Key Takeaways)

本次报告的关键要点可总结为以下三方面：

1. 通过异步卸载实现更高的 GPU 吞吐量：Grace CPU 可以高效地将计算任务异步卸载到 Hopper GPU，最大化 GPU 的利用率。
2. 通过统一内存为系统内存提供高带宽缓存：NVLink C2C 实现了 CPU 和 GPU 内存的统一访问，为系统内存提供了高带宽的缓存机制，支持 CPU 和 GPU 对彼此内存的高效远程访问。
3. 优化的 AI 软件栈与开发者生产力：NVIDIA 提供了从硬件（Grace CPU + Hopper GPU）到 CUDA，再到上层训练与推理框架的完整优化软件栈，极大地提升了开发者的生产力。

CUDA 开发者会议

以下是推荐的 CUDA 开发者会议，涵盖了从入门到高级优化的各类主题：

• 通用 CUDA (General CUDA)
• CUDA Python
• CUDA C++
• 开发者工具 (Developer Tools)
• 与专家交流 (Connect with the Experts)
• 多 GPU 编程 (Multi-GPU Programming)
• 性能优化 (Performance Optimization)

更多详情请访问：nvidia.com/gtc/sessions/cuda-tuesday