Programming Model and Applications for Grace Hopper Superchip

Vishal Mehta, Mathias Wagner, Devtech Compute
GTC, March 2023

Grace Hopper Superchip 概述
Grace Hopper 架构
- Hopper 架构：H100 GPU 关键特性
- H100 在 HPC 和 AI 领域的领导地位
- GRACE 架构
- NVIDIA Grace CPU 性能和效率
- Grace Hopper Superchip 结构
- Grace Hopper 平台
- Grace Hopper 利用率：弹性计算
Grace Hopper 应用与性能
- 异构计算概述
- 加速系统上的应用分类
  - 完全 GPU 加速
  - 部分 GPU 加速
- 高带宽内存访问与自动数据迁移
  - GPU 访问 GPU 内存带宽
  - CPU 访问 CPU 内存带宽
  - GPU 访问 CPU 内存带宽
- 应用案例分析
  - OpenFoam
  - NAMD Collective Variables
  - CP2K
  - Hash Joins TPC-H Query 4
- 加速系统上的应用原则
内存一致性 (Coherent Memory)
- 传统 X86 + GPU：独立的页表
- Grace Hopper：地址转换服务 (ATS)
- Grace Hopper：统一内存的自动迁移
- Grace Hopper 中的缓存一致性
- 多节点 Grace Hopper 访问路径
Grace Hopper 编程模型与内存管理
- 异构内存的直接访问 (cudaMalloc)
- CUDA 托管内存的直接访问 (cudaMallocManaged)
- 系统分配内存 (传统 vs. Grace Hopper)
- 异构编程示例 (cudaMalloc)
- 异步异构编程 (Grace Hopper 优势)
- CUDA 统一内存或托管内存 (工作流)
  - 内存分配
  - CPU 写入数据
  - GPU 启动核函数
  - 数据预取
- 系统分配内存 (深入探讨)
  - 系统分配器
  - CUDA 内存提示
具有 CPU/GPU 一致性的应用程序实例
- cuCollections 哈希图基准测试
- 通过优化内存访问增加性能
Grace Hopper Superchip 总结

1. Grace Hopper Superchip 概述

Grace Hopper Superchip 专为巨型规模的 AI 和 HPC 而设计，具备以下关键特性：
* 最高加速性能：通过 Grace CPU 和 Hopper GPU 加速实现。
* GPU 可用内存接近 600GB：支持巨型 AI 模型进行训练和推理。
* 最高内存带宽 4 TB/s：采用 LPDDR5X 和 HBM3。
* 900GB/S 相干接口：通过 NVLink-C2C 连接 Grace 和 Hopper。
* 运行完整的 NVIDIA 计算栈：支持 HPC、AI 和 Omniverse 应用。

2. Grace Hopper 架构

Hopper 架构：H100 GPU 关键特性

Hopper 架构，特别是 H100 GPU，提供了多项创新，进一步提升了性能：
* 第二代多实例 GPU (Multi-Instance GPU)
* 机密计算 (Confidential Computing)
* PCIe Gen5
* 更大的 60 MB L2 缓存
* 96GB HBM3，带宽达到 4 TB/s
* 132 个流多处理器 (SMs)
* 第四代 Tensor 核心
* 线程块簇 (Thread Block Clusters)
* 第四代 NVLink，总带宽达到 900 GB/s

Hopper Architecture H100 GPU Key Features

CUDA 编程模型和 Hopper 架构应用优化是后续的深入讨论点。

H100 在 HPC 和 AI 领域的领导地位

NVIDIA H100 在 HPC 和 AI 训练方面展现出显著的性能提升。

HPC GOLDEN SUITE 加速
与 2016 年的 Broadwell 双路处理器相比，H100 实现了 470x 的性能加速。

AI 训练
H100 在各种 AI 模型训练中相较于 A100 提供了更高的性能：
* Mask R-CNN
* GPT-3 (16B 参数)
* DLRM (14TB 嵌入)
* GPT-3 (175B 参数)
* MoE Switch-XXL (395B 参数)

H100 continues NVIDIA's leadership in HPC and AI

GRACE 架构

GRACE 架构基于 ARM Neoverse V2 和 NVIDIA 可扩展相干性结构，提供了强大的 CPU 能力：
* 核心配置：72 个 Neoverse V2 核心，Armv9，SVE2 支持 4x128b。
* 缓存配置：
* 每个核心 64KB i-cache 和 64KB d-cache。
* 每个核心 1 MB L2 缓存。
* 117 MB L3 SCF 缓存。

互连：
- 900 GB/s NVLink C2C。
- 高达 4x PCIe Gen5 x16。
NVIDIA 可扩展相干性结构：
- 3,225.6 GB/s 双向带宽。
- 117MB L3 缓存。
- 支持远程 CPU/GPU 内存的本地缓存。
- 通过相干 NVLINK 支持最多 4 芯片相干性。
内存：LPDDR5X，500 GB/s 带宽，每个 CPU 240 GB。

NVIDIA Grace CPU 性能和效率

Grace CPU 实现了数据中心吞吐量 2 倍的提升。
* CFD (计算流体动力学)：1.9X 相对吞吐量。
* 大数据 (Big Data)：2.0X 相对吞吐量。
* 微服务 (Microservices)：2.3X 相对吞吐量。
这些性能均与下一代 x86 CPU 相比。

Nvidia Grace CPU Delivers Breakthrough Performance and Efficiency

Grace Hopper Superchip 结构

Grace Hopper Superchip 使得 GPU 能够以 CPU 内存速度访问 CPU 内存，实现了硬件一致性。
* CPU 部分：
* CPU LPDDR5X 240 GB。
* Grace CPU。
* 500 GB/s 内存带宽。

GPU 部分：
- GPU HBM3 96 GB。
- Hopper GPU。
- 4000 GB/s 内存带宽。
互联：
- NVLINK C2C，900 GB/s 连接 Grace CPU 和 Hopper GPU。
- 高速 I/O：4x 16x PCIe-5，512 GB/s。
- NVLink 网络：18x NVLINK 4，900 GB/s，可连接多达 256 个 GPU。

Grace Hopper 平台

Grace Hopper 平台通过 NVLink 和 InfiniBand 实现大规模扩展。

顶部互联：NVIDIA Quantum-2 InfiniBand NDR400，提供 100 GB/s 带宽连接 BlueField DPU (0-255)。
Grace Hopper Superchip 阵列：多个 Grace CPU 和 Hopper GPU 通过 NVLINK-C2C (900 GB/s) 和硬件一致性连接。
底部互联：NVIDIA NVLink Switch System。

Grace Hopper 利用率：弹性计算

系统管理员可以组合 CPU 分区和 GPU 实例。NVLink C2C 可以通过 CPU 分区进行分区。
* 左侧图：展示了 NVLINK-C2C 的分区，分为 PARTID 0 和 PARTID 1，每个分区包含 LPDDR、SCC 和 CSN 块。
* 右侧图：展示了 GPU 的多实例 GPU (MIG) 分区，例如 MIG partition 0。

3. Grace Hopper 应用与性能

异构计算概述

异构计算结合了不同类型的处理器，每种处理器专精于不同类型的执行。一个应用程序可以同时利用多个处理器。
应用程序可分为三类：
* 完全 GPU 加速 (Fully GPU Accelerated)
* 部分 GPU 加速 (Partially GPU Accelerated)
* 相干 GPU 加速 (Coherently GPU Accelerated)

加速系统上的应用分类

完全 GPU 加速

在这类应用中，计算几乎完全在 GPU 上完成，数据也驻留在 GPU 内存中。CPU 和数据传输对性能的限制极小或没有限制。

Application on Accelerated Systems Fully GPU Accelerated

部分 GPU 加速

随着 GPU 速度越来越快，应用程序越来越受非 GPU 因素的限制，例如：
* 主要受数据传输 (PCIe) 限制：CPU 和 GPU 之间频繁的数据传输成为瓶颈。
* 主要受 CPU 限制：GPU 可能等待 CPU 完成其任务或准备数据。

Application on Accelerated Systems Partially GPU Accelerated

高带宽内存访问与自动数据迁移

Grace Hopper 架构通过 NVLink C2C 实现了 CPU 和 GPU 之间的高带宽内存访问和自动数据迁移。

GPU 访问 GPU 内存带宽

GPU 访问其自身的 HBM3 内存时，带宽表现如下：
- 峰值：4000 GB/sec
- 实际达到：3732 GB/sec
GPU访问GPU内存

CPU 访问 CPU 内存带宽

CPU 访问其自身的 LPDDR5X 内存时，带宽表现如下：
- 峰值：500 GB/sec
- 实际达到：475 GB/sec
CPU访问CPU内存

GPU 访问 CPU 内存带宽

GPU 通过 NVLink C2C 访问 CPU 的 LPDDR5X 内存时，带宽表现如下：
- 峰值：500 GB/sec
- 实际达到：486 GB/sec
这表明 GPU 能够以接近 CPU 访问自身内存的峰值带宽来访问 CPU 内存，彰显了 Grace Hopper 架构中内存访问的灵活性和高效率。
GPU访问CPU内存

应用案例分析

OpenFoam

OpenFoam 是一个由 OpenCFD 开发的计算流体动力学 (CFD) 工具箱，在汽车及其他工程领域广泛应用。它包含高度可配置的流体流动求解器，并利用 GPU 加速的 AMGX 线性求解器。OpenFoam 当前是部分 GPU 加速，但主要受限于 CPU 性能。

HPC 摩托车问题 (大型)：约 30% 的纯 CPU 执行时间花费在线性求解器中。
Grace Hopper 性能表现：
- 高 CPU 和 GPU 内存带宽提升了计算性能。
- C2C 带宽最大限度地降低了迁移 CPU 矩阵数据的成本。
性能对比 (Openfoam MotorBike L)：
- x86+A100 (基准): 1.0x
- x86+H100: 1.4x
- Grace Hopper: 2.6x

通过 Nsight Systems 对 OpenFoam 进行性能分析显示：

在 Grace Hopper (H100 96GB HBM3) 系统上，相较于 x86 + H100 (H100 80GB HBM3) 系统，CPU 部分的执行速度提高了 2 倍，GPU 部分的执行速度提高了 1.3 倍。这表明 Grace Hopper 在处理 OpenFoam 负载时，CPU 和 GPU 协同工作，显著提升了整体性能。

NAMD Collective Variables

NAMD 是一款分子动力学模拟软件，其 Collective Variables 功能由第三方 Colvars 模块提供（纯 CPU 实现）。此功能是流行的扩展，也用于 Gromacs、LAMMPS 等软件，并且当前只与 NAMD 的 GPU 卸载模式兼容。NAMD 当前是部分 GPU 加速，但主要受限于 CPU 性能。

葡萄糖转运蛋白 3 系统 (143k 原子)：
- 大部分力计算由 GPU 加速。
- 额外的 Colvars 力在 CPU 上计算。
- 主要瓶颈是受限于 CPU 内存带宽的积分步骤。
Grace Hopper 性能表现：
- Grace CPU 内存带宽加速了积分步骤。
- NVLink C2C 带宽避免了数据传输瓶颈。
性能对比 (NAMD + Colvars)：
- x86+A100 (基准): 1.0x
- x86+H100: 1.4x
- Grace Hopper: 2.6x

CP2K

CP2K 是一款量子化学应用，实现了多种方法，但许多尚未 GPU 加速。内存容量限制常常要求将部分数据保留在系统内存中。CP2K 当前是部分 GPU 加速，但主要受限于数据传输。

数据集 "128-H2O" (随机相近似 RPA 方法)：
- PDGEMM 主导 CPU 运行时 -> 使用 GPU 加速的 PDGEMM。
- 在 x86 + 当前 GPU 系统上：PDGEMM 性能受限于 PCIe 数据传输。
- 性能受限于 CPU 内存带宽、数据传输和 MPI 通信。
Grace Hopper 性能表现：
- Grace CPU 内存带宽和 NVLink C2C 极大地加速了 CP2K RPA。
- C2C 加速的传输隐藏在 GPU 加速的 PDGEMM 之后。
性能对比 (CP2K RPA)：
- x86+A100 (基准): 1.0x
- x86+H100: 2.2x
- Grace Hopper: 3.8x

通过 Nsight Systems 对 CP2K RPA 进行性能分析显示：

在 Grace Hopper (H100 96GB HBM3) 系统上，相较于 x86 + H100 (H100 80GB HBM3) 系统，"主机到设备" 的数据传输时间显著缩短，性能提高了 6.2 倍。这表明 Grace Hopper 有效缓解了数据传输瓶颈。

Hash Joins TPC-H Query 4

哈希连接 (Hash Joins) TPC-H Query 4 包含连接操作和分组聚合。该应用是部分 GPU 加速的，但主要受限于数据传输。

在 GPU 上创建哈希表以查找匹配键：
- 步骤 1：基于第一个输入表的键构建哈希表。
  - 直接访问 CPU 内存中的键（零拷贝访问）。
- 步骤 2：针对第二个输入表的键探测哈希表并构建输出结果。
  - 从 CPU 内存复制输入到 GPU 内存 (cudaMemcpy)。
- 步骤 2 中的 PCIe 数据传输是主要限制因素。
Grace Hopper 性能表现：
- NVLink-C2C 显著减少了传输时间。
性能对比 (Hash Join TPC-H Query 4)：
- x86+A100 (基准): 1.0x
- x86+H100: 2.0x
- Grace Hopper: 4.6x

加速系统上的应用原则

加速系统上的应用应充分利用 GPU / CPU 的一致性，并利用所有可用的系统特性。在这些系统中，阶段之间不一定存在清晰的界限。

4. 内存一致性 (Coherent Memory)

传统 X86 + GPU：独立的页表

在传统的 X86 + GPU 架构中，CPU 和 GPU 拥有独立的页表。

CPU 物理内存 (LPDDR5X) 和 GPU 物理内存 (HBM3) 分开管理。
系统页表 (转换 malloc()) 映射到 CPU 物理内存中的 CPU 常驻访问 (PTE A)。
GPU 页表 (转换 GPU 分配) 映射到 GPU 物理内存中的相应区域。
当 CPU 尝试直接访问 GPU 物理内存，或 GPU 尝试直接访问 CPU 物理内存时，由于页表独立，会发生未映射的访问失败。

Grace Hopper：地址转换服务 (ATS)

Grace Hopper 引入了地址转换服务 (ATS)。
- CPU 物理内存 (LPDDR5X, Page A) 和 GPU 物理内存 (HBM3, Page B)。
- GRACE CPU 和 HOPPER GPU 通过 NVLINK C2C 连接。
- 系统页表 (将 CPU malloc() 转换为 CPU 或 GPU 地址) 统一管理内存。
- CPU 常驻访问 (PTE A) 指向 CPU 物理内存中的 Page A。
- 远程访问 (PTE B) 通过 NVLINK C2C 指向 GPU 物理内存中的 Page B。
- GPU 常驻访问直接指向 GPU 物理内存中的 Page B。
ATS 实现了 CPU 和 GPU 共享一个统一的系统页表，使得远程访问成为可能。
Grace Hopper 地址转换服务 (ATS)

Grace Hopper：统一内存的自动迁移

Grace Hopper 实现了统一内存的自动数据迁移。
- 初始时，数据页面 (Page A) 可能位于 CPU 物理内存。
- GRACE CPU 偶尔访问 Page A (Infrequent accesses)。
- 如果 HOPPER GPU 对 Page A 的访问变得频繁 (Frequent accesses)，系统将检测到并自动将 Page A 从 CPU 物理内存迁移到 GPU 物理内存。
- 数据迁移能够提高性能。
Grace Hopper 统一内存的自动迁移

Grace Hopper 中的缓存一致性

Grace Hopper 通过 NVLink C2C 实现了缓存一致性访问。
- CPU 访问 GPU 内存：
- GRACE CPU 核心通过 CPU L3 缓存 (GPU 缓存行) 访问 NVLINK C2C，进而访问 GPU HBM3。
- CPU 通过“签出缓存行”来访问 GPU 内存。

GPU 访问 CPU 内存：
- HOPPER SMs 通过 GPU L2 缓存访问 NVLINK C2C，进而访问 CPU L3 缓存 (CPU 缓存行)，再访问 CPU LPDDR5X。
- GPU “无需签出缓存行”即可一致地访问 CPU 内存。

多节点 Grace Hopper 访问路径

本节介绍了多节点Grace Hopper与内存一致性NVLink的访问路径。

该图展示了两个Grace Hopper节点之间的连接和数据访问路径：

每个节点包含一个GRACE CPU和LPDDR5X内存，以及一个HOPPER GPU和HBM3内存。
GRACE CPU和HOPPER GPU通过NVLINK C2C连接，提供本地CPU到GPU的访问（Local CPU <-> GPU）。
两个HOPPER GPU之间通过NVIDIA NVLINK Network连接，支持GPU到对等GPU（GPU -> Peer GPU）以及GPU到对等CPU（GPU -> Peer CPU）的访问。

5. Grace Hopper 编程模型与内存管理

Grace Hopper超级芯片的内存模型旨在优化异构内存访问。

异构内存的直接访问 (cudaMalloc)

本节比较了PCIe H100和Grace Hopper在cudaMalloc内存模型下的表现。

PCIe H100： CPU页表和CPU内存与GPU页表和GPU内存是独立的。CPU页表中的指针ptr无法直接访问GPU内存。这意味着cudaMalloc分配的GPU内存无法被CPU直接访问。
Grace Hopper： 对于cudaMalloc，Grace Hopper的内存模型与PCIe H100的cudaMalloc行为类似，CPU页表中的指针ptr同样无法直接访问GPU内存。

CUDA 托管内存的直接访问 (cudaMallocManaged)

本节比较了PCIe H100和Grace Hopper在cudaMallocManaged内存模型下的表现。

PCIe H100： 访问托管内存需要故障和迁移。CPU页表中的指针ptr最初无法直接访问GPU内存，需要数据在CPU和GPU之间迁移。
Grace Hopper： 直接访问可以减少故障。更快的迁移。CPU页表中的指针ptr可以直接指向GPU内存。这是Grace Hopper在托管内存方面的一个关键优势。

系统分配内存 (传统 vs. Grace Hopper)

本节比较了PCIe H100和Grace Hopper在系统分配内存方面的处理。

PCIe H100： 需要cudaHostRegister()函数来将主机内存注册为GPU可访问的页锁定内存。CPU页表中的ptr指向CPU内存。
Grace Hopper - ATS： 不需要cudaHostRegister()。访问速度达到NVLink C2C的带宽。CPU页表中的ptr可以直接访问CPU内存。通过共享页表，GPU可以以NVLink C2C的速度直接访问系统分配的内存。

异构编程示例 (cudaMalloc)

这是一个在x86平台和Grace Hopper上均可工作的异构编程示例，使用cudaMalloc。

cudaMalloc在Grace Hopper上仍然保持与x86平台相同的行为，即CPU无法直接访问由cudaMalloc分配的GPU内存。
代码示例流程：
1. CPU使用malloc分配主机内存。
2. GPU使用cudaMalloc分配设备内存。
3. 使用cudaMemcpyHostToDevice将数据从主机复制到设备。
4. 启动CUDA核函数在GPU上执行。
5. 使用cudaMemcpyDeviceToHost将结果从设备复制回主机。
6. 释放CPU和GPU内存。
系统架构图中展示了Grace CPU与Hopper GPU通过900 GB/s的NVLINK C2C连接。

异步异构编程 (Grace Hopper 优势)

本节比较了x86 + Hopper与Grace Hopper在异步异构编程中使用cudaMallocHost和cudaHostRegister的区别。

在x86 + Hopper平台上：
异步异构编程 - x86+Hopper

cudaMallocHost：为cudaMemcpyAsync分配页锁定内存。
cudaHostRegister：为cudaMemcpyAsync页锁定现有主机内存。
图示的执行时间线（Host2Device, Kernel, Device2Host）表明这些操作通常是串行的。

在Grace Hopper平台上：
异步异构编程 - Grace Hopper

在Grace Hopper上，cudaMallocHost和cudaHostRegister 不再需要。
所有内存都可以以C2C带宽进行迁移。
cudaMallocHost：不需要锁定。但cudaMallocHost == malloc() + cudaHostRegister()，意味着它的行为等同于常规分配加注册。
cudaHostRegister：不需要锁定。但cudaHostRegister()可以填充页面。
图示的执行时间线显示了Host2Device、Kernel和Device2Host操作可以重叠执行，实现真正的异步性。

CUDA 统一内存或托管内存 (工作流)

此内存模型通过移动内存页面，使得代码可以在x86和Grace Hopper之间移植。

内存分配

代码示例使用cudaMallocManaged分配统一内存：cudaMallocManaged((void **)&data, N * sizeof(int));。
这行代码在CPU和GPU之间分配了一个可由两者访问的内存区域。

CPU 写入数据

CPU通过data[i] = 1;访问并写入数据。在Unified Memory模型下，当CPU访问时，数据可能被迁移到CPU侧。

GPU 启动核函数

GPU通过kernel<<<1, N>>>(data);启动核函数访问数据。当GPU访问时，数据可能被迁移到GPU侧，或者GPU可以直接访问。

此图进一步强调了GPU在执行核函数时对数据的访问。

数据预取

cudaMemPrefetchAsync(data, N * sizeof(int), 0); 被用于异步预取数据到设备0（通常是GPU）。这可以显式地将数据从主机内存（CPU LPDDR5X）预取到设备内存（GPU HBM3），以优化访问性能。

系统分配内存 (深入探讨)

系统分配器

代码轻松迁移到 GPU。
利用地址转换服务 (ATS)。
内存分配过程：
- 初始状态：CPU 使用 malloc 分配内存，物理页尚未分配。
- 首次接触：当 GPU 内核 kernel<<<1, 128>>>(data, data2); 首次访问 CPU 分配的内存时，物理页才被分配。
- 直接访问：CPU 可以直接访问 GPU 数据，体现了统一内存模型的优势。

CUDA 内存提示

用于调整内存移动。
通过 cudaMemPrefetchAsync API 进行内存预取，将数据从 CPU 内存异步预取到 GPU 内存，以优化性能。
预取数据后，GPU 可以直接访问。
CUDA 内存 API 或 NUMA API 可用于调整内存放置和移动，例如通过 cudaMemLocationTypeHostNUMA 和 CPUNUMAId 精细控制内存位置。

6. 具有 CPU/GPU 一致性的应用程序实例

cuCollections 哈希图基准测试

当数据不在 GPU 内存中时，cuCollections 能够在 CPU-GPU 链接带宽速度下构建和查询哈希表。
在 GPU 上构建哈希表，而数据集驻留在 CPU 上，方便且无需显式移动数据。
从 GPU 到 CPU 内存的一致性访问。
Grace Hopper 性能表现：
- 在 CuCo-find 和 CuCo-Insert 操作中，Grace Hopper 显示出显著优于 x86+A100 和 x86+H100 的性能。
- 特别是在 CuCo-Insert 操作中，Grace Hopper 的性能比 x86+A100 高出超过 3.5 倍。

通过优化内存访问增加性能

展示了通过优化内存访问和 CPU/GPU 一致性实现的性能提升。
右侧的报告显示，在 GPU 光速吞吐量 (GPU Speed Of Light Throughput) 方面，相比左侧报告有 2.7 倍的性能提升。
具体的吞吐量数据：
- 左侧报告（基线）：计算 (SM 吞吐量) 19.31 (55.31%)，内存吞吐量 17.49 (58.08%)，L1/TEX 缓存吞吐量 3.54 (56.37%)，L2 缓存吞吐量 22.91 (54.11%)，DRAM 吞吐量 17.49 (58.08%)。
- 右侧报告（优化后）：计算 (SM 吞吐量) 50.43 (1.51%)，内存吞吐量 39.60 (4.93%)，L1/TEX 缓存吞吐量 25.88 (4.71%)，L2 缓存吞吐量 49.86 (1.28%)，DRAM 吞吐量 39.60 (4.93%)。

7. Grace Hopper Superchip 总结

Grace Hopper Superchip 的主要特点：

平衡且多功能的高性能系统：适用于所有现有工作负载。
新应用的连贯架构：Grace CPU 和 Hopper GPU 紧密结合，提供一致性。
提升开发人员加速应用程序的速度：通过优化和统一的编程模型，提高开发效率。