OPTIMIZING CUDA APPLICATIONS FOR NVIDIA HOPPER ARCHITECTURE

Guillaume Thomas-Collignon, Vishal Mehta
DevTech Compute, GTC 2022

霍珀架构 (Hopper Architecture)
- H100 GPU 关键特性
- H100 流多处理器 (Streaming Multiprocessor) 关键特性
即将讨论的新特性
线程和内存层级结构 (Thread and Memory Hierarchy)
- Hopper之前的CUDA线程和内存层级结构
  - 线程块间的协作
- Hopper中引入线程块集群 (Thread Block Clusters)
- 集群中所有线程的加速同步
- 分布式共享内存操作
- 使用集群启动CUDA内核
  - 通过集群启动CUDA内核
  - 使用CUDA可扩展内核启动API
- 示例：共享内存直方图
- 示例：分布式共享内存直方图
  - 分布式共享内存直方图实现 (代码片段)
  - 分布式共享内存直方图实现 (续)
- 直方图性能
异步SIMT编程模型 (Asynchronous SIMT Programming Model)
- 张量内存加速器单元 (TMA)
- 异步屏障 (Ampere 模型)
- 异步事务屏障 (Hopper 新特性)
  - 异步 SIMT 编程与异步事务屏障
- cuda::memcpy_async 全局 <-> 共享
- Stencil 代码示例
  - 无 TMA
  - 使用 TMA
  - TMA 示例：性能提升
- TMA 示例：3D TTI 逆时偏移 (Reverse Time Migration)
- 集群中的生产者/消费者编程
  - 异步 SIMT 编程：共享内存之间的 cuda::memcpy_async
  - 消费者等待生产者数据
  - 生产者使用消费者屏障发送异步内存拷贝
  - 更多异步内存拷贝及生产者线程抵达
  - 所有线程和数据抵达，屏障解除
生产者/消费者用例：Longstaff Schwartz 定价模型
- Longstaff Schwartz 定价模型介绍
- 算法和计算方法
- 当前实现
- H100 上的集群实现
- H100 上的集群实现，采用单边通信
- H100 上的性能表现
结论与主要要点

霍珀架构 (Hopper Architecture)

H100 GPU 关键特性

H100 GPU引入了多项关键特性，旨在提升性能和功能。

主要特性包括：
* 第2代多实例GPU (Multi-Instance GPU)
* 机密计算 (Confidential Computing)
* PCIe Gen5
* 更大的50 MB L2缓存
* 80GB HBM3，3 TB/s带宽
* 132个流多处理器 (SMs)
* 第4代张量核心 (Tensor Core)
* 线程块集群 (Thread Block Clusters)
* 第4代NVLink，总带宽900 GB/s

更多信息可参考：
* Inside the NVIDIA Hopper Architecture (S42663)
* CUDA: New Features and Beyond (S41486)

H100 流多处理器 (Streaming Multiprocessor) 关键特性

H100 SMs的设计带来了显著的改进。

主要特性包括：
* 256 KB的组合L1缓存/共享内存，每个SM比A100增加33%。
* 新的线程块集群 (Thread Block Clusters) 和分布式共享内存 (Distributed Shared Memory)。
* 新的张量内存加速器 (Tensor Memory Accelerator) 和异步事务屏障 (Asynchronous Transaction Barriers)。
* 第4代张量核心，每个时钟周期性能提升2倍。

即将讨论的新特性

本次演讲将深入探讨以下Hopper架构中的新特性：

线程层级结构 (Thread Hierarchy)
- 新线程块集群 (New Thread Block Clusters)
- 分布式共享内存 (Distributed Shared Memory)
异步SIMT编程模型 (Asynchronous SIMT Programming Model)
- 张量内存加速器 (Tensor Memory Accelerator) (cuda::memcpy_async)
- 异步事务屏障 (Asynchronous Transaction Barriers)

线程和内存层级结构 (THREAD AND MEMORY HIERARCHY)

在CUDA编程模型中，线程和内存的层级结构对于理解和优化CUDA应用程序至关重要。

Hopper之前的CUDA线程和内存层级结构

在Hopper架构之前，CUDA的线程和内存层级结构有以下特点：

线程块中的所有线程都可以使用共享内存 (Shared memory) 进行协作。
线程块中的所有线程都保证在单个流多处理器 (SM) 上协同调度。
线程块中的线程可以使用cooperative_groups::this_thread_block.sync() 或 __syncthreads() 进行同步/通信。
也可以使用 cuda::barrier<thread_scope_block>::arrive() 和 ::wait()。
线程块中的线程还可以执行集合操作，如 cooperative_groups::reduce()。

线程块间的协作

所有线程块共享全局内存 (global memory) 以进行协作。
独立的线程块可以乱序调度，以提高占用率 (occupancy) 和 GPU 利用率。
所有线程块要在GPU上同步，需要CUDA协同启动 (cooperative launch)。

Hopper中引入线程块集群 (Thread Block Clusters)

Hopper架构引入了线程块集群，为CUDA编程模型带来了新的可选层级结构。

线程块集群引入了CUDA编程模型中一个新的可选层级结构。
集群中的线程块保证在GPU处理集群 (GPC) 中的多个SM上协同调度。
集群内的所有共享内存组成了分布式共享内存 (Distributed shared memory)。

集群中所有线程的加速同步

Hopper架构为集群内的线程提供了硬件加速同步机制。

namespace cg = cooperative_groups;
auto block = cg::this_thread_block();
cg::cluster_group cluster = cg::this_cluster();

<...>

集群同步在硬件中得到加速。
H100每个集群最多支持16个线程块或16384个线程。

cluster.sync();

免责声明：预发布CUDA API，可能会有变动。

分布式共享内存操作

线程块集群内的所有块都可以使用分布式共享内存进行协作。

线程块可以在彼此的共享内存上进行读、写和原子操作。

// All blocks in the cluster have the variable smem
__shared__ int smem;
namespace cg = cooperative_groups;
cg::cluster_group cluster = cg::this_cluster();
unsigned int BlockRank = cluster.block_rank();
int cluster_size = cluster.dim_blocks().x;

免责声明：预发布CUDA API，可能会有变动。

// Get a pointer to peer smem variable based on
// pointer from current block
int *remote_smem = cluster.map_shared_rank(&smem,
                                        (BlockRank + 1) % cluster_size);

if (threadIdx.x == 0)
    *remote_smem = 10; // Store to remote memory

cluster.sync(); // Sync to ensure store is done

分布式共享内存远程存储 (Page 15)
* 免责声明：预发布CUDA API，可能会有变动。

线程和内存层次结构：使用集群启动CUDA内核

通过集群启动CUDA内核

编译时线程块集群大小标注: 允许在编译时指定集群尺寸来标注内核。
内核启动方式: 沿用经典的 <<< >>> 方式启动内核。
- 例如，在编译时，一个内核在X维度和Y维度各包含2个线程块。
- 要求线程块的数量必须是4的倍数。
- 示例代码: global_ void cluster_dims_(2, 2, 1) clusterKernel() { ... }

使用CUDA可扩展内核启动API

启动方式: 通过可扩展启动API启动。
配置对象: 使用 cudaLaunchConfig_t 配置对象，通过 attribute[0].Id = cudaLaunchAttributeClusterDimension; 和 attribute[0].val.clusterDim.x = 2; 等设置集群维度。
内核启动函数: cudaLaunchKernelEx(&config, (void*)clusterKernel, params);

示例：共享内存直方图

传统CUDA直方图计算: 通常在共享内存中计算，然后通过全局内存进行规约。

示例：分布式共享内存直方图

问题: 对于大型直方图，共享内存容量可能不足。例如，300KB或75K个整数直方图桶。
解决方案: 引入分布式共享内存。
处理流程: 每个线程块集群负责N/2个直方图桶，最终在全局内存中进行规约。

分布式共享内存直方图实现 (代码片段)

DSMEM指针初始化: extern __shared__ int smem[];，cg::cluster_group cluster = cg::this_cluster();，unsigned int cluster_size = cluster.dim_blocks().x;
共享内存直方图初始化为零: 循环遍历 smem[i] = 0;。
集群同步: cluster.sync();

分布式共享内存直方图实现 (续)

加载输入数据并找到直方图桶ID:
- int dst_block_rank = (int)(binid / bins_per_block);
- int dst_offset = binid % bins_per_block;
- atomicAdd(&sh_hist[dst_block_rank] + dst_offset, 1);
集群同步: cluster.sync();
执行全局内存规约: // Perform global memory reductions

直方图性能

性能提升: 在H100集群上，使用分布式共享内存的直方图计算比单H100快1.7倍。
示例: 75K个直方图桶 (300KB) 在分布式共享内存中处理，2个线程块集群 -> 每个线程块处理37.5K个 (150KB)。

异步SIMT编程模型 (ASYNCHRONOUS SIMT PROGRAMMING)

介绍

张量内存加速器单元 (TMA)

功能: cuda::memcpy_async
硬件加速: 1D到5D张量内存拷贝。
- 全局内存 -> 共享内存
- 共享内存 -> 全局内存
- 支持全局内存中的元素级规约
硬件加速: 1D或元素级存储和规约。
- 集群中的共享内存 -> 共享内存
关键机制: 使用异步事务屏障来信号数据传输完成。

异步屏障 (Ampere 模型)

目的: 允许独立工作重叠执行。
异步屏障 (Ampere) 模型: 生产者/消费者模型。
- 生产数据 -> 屏障 -> 消费数据
屏障分为两步:
- Arrive (到达): 线程完成数据生产。
- Wait (等待): 线程准备开始消费数据。
Arrive是非阻塞的: 等待时，提前到达的线程可以执行独立工作。
异步屏障状态: 保存在共享内存中。
Arrive/Wait 分裂屏障: 工作流包含计算值并存储结果 -> 到达 -> 独立工作 -> 等待 -> 处理结果。

异步事务屏障 (Hopper 新特性)

扩展异步屏障: 增加了“数据到达跟踪”功能（Hopper新特性）。
Async事务屏障 (Hopper新特性):
- 线程增加 Arrival_count。
- 异步存储到共享内存增加 Transaction_count。
等待机制: 只有当 Arrival_count 和 Transaction_count 都达到预期值时才阻塞等待。
关键作用: cuda::memcpy_async 单边数据交换的构建块。
工作流: 计算值并存储结果 -> 到达/事务 -> 独立工作 -> 等待 -> 处理结果。

异步 SIMT 编程与异步事务屏障

本节介绍异步 SIMT 编程，特别是 cuda::memcpy_async Global <-> Shared 的应用。它强调使用异步事务屏障（Asynchronous Transaction Barrier）来发出完成信号，实现完全异步的线程操作。

上图展示了 H100 SM 中异步事务屏障的工作流程：
1. Arrive (到达)：一个事务屏障被初始化，设置 block.size() 和 sizeof(int) * block.size()。
2. Transaction (事务)：一个事务被启动，例如通过 cuda::memcpy_async 进行数据拷贝。
3. Wait (等待)：等待所有事务完成。
图中的代码片段展示了如何初始化合作线程组、设置屏障，并异步执行 memcpy 操作，最后通过 barrier.wait() 等待完成。

`cuda::memcpy_async` 全局 <-> 共享

完全异步: 线程完全异步，并使用异步事务屏障来信号完成。
初始化: 屏障和块大小的初始化。
- block.size() 定义预期计数。
- init(&abarrier, block.size()); 初始化屏障。
GPU组件交互: 异步事务屏障与H100 SM中的共享内存、寄存器、L1缓存、TMA单元以及GPU内存交互。

集体发出异步拷贝操作:
- cuda::memcpy_async(block, smem, gmem, block.size() * sizeof(int), // int/thread barrier);
- 这表示从GPU内存通过TMA单元将数据拷贝到共享内存。

屏障到达:
- auto token = barrier.arrive(); 线程完成数据发送后，调用 arrive()。

独立工作:
- 发出 arrive() 后，线程可以立即执行独立工作，而无需等待数据传输完成。

Stencil 代码示例

无 TMA

Stencil 代码是一种常见的 GPU 计算模式，其典型步骤包括：

使用一个 2D 线程块。
将 2D 瓦片复制到共享内存。
将 Halo 区域加载到共享内存。
使用 __syncthreads() 进行同步。
计算并写入结果。

上图展示了典型的 Stencil 算子结构，包括中心（Center）和四周的 Halo 区域（X-, X+, Y-, Y+）。

在没有使用 TMA（Tensor Memory Access）的情况下加载 Stencil 数据，需要大量的测试来检查要加载到共享内存中的数据是否存在，如果不存在则用零填充。

上图的 C++ 代码示例显示了在加载 tile 和 halo 区域时，需要通过 if (idx < nx && sy < ny) 等条件判断来处理边界情况，这导致了代码的复杂性。

使用 TMA

通过使用张量描述符（Tensor Descriptor）的 2D TMA 加载，可以简化 Stencil 代码。
张量描述符提供：
* 维度 NX, NY
* Y 维度步长 LDIMX >= NX
* 共享内存框大小
* 基指针

每个块可以使用以下方式加载数据：
* 张量描述符
* (X,Y) 偏移

上图解释了张量描述符如何定义数据区域，包括 NX, NY 尺寸和 LDIMX 步长，以及如何通过 Box size X 和 Box size Y 来定义加载范围。

TMA 的优势：
* TMA 可以通过一条指令加载整个共享内存区域。
* 单个线程可以向共享内存加载数十 KB 的数据。
* 如果数据不存在，TMA 会自动进行零填充。
* 使用标准的 cuda::barrier 同步，实现完全异步。

上图对比了传统 Stencil 区域（左）与 TMA 实现的单次加载（右），表明 TMA 能将复杂的加载操作简化为一次。

通过 TMA 简化数据加载。
与 Page 33 所示的复杂边界检查代码相比，使用 TMA 可以显著简化代码。Page 37 通过红叉强调了这部分复杂代码的移除。

使用 TMA 后，不再需要大量的边界检查代码，因为 TMA 提供了自动零填充。

上图中的伪代码显示，复杂的条件判断和手动加载被简化为一行 memcpy_async 调用，同时图示强调了 TMA 的“自动零填充”功能。

TMA 示例：性能提升

通过计算每个线程 4 x 4 个点而不是一个点，可以改进 Stencil 代码的性能。

上图展示了 Stencil 算子从单个中心点扩展到 4 x 4 区域的计算方式。

TMA 对 Stencil 代码性能的影响：
* 没有 TMA： 需要加载更多数据并进行适当的边界检查，导致大量的测试和加载指令。
* 使用 TMA： 张量描述符只需进行微小改动（更大的框尺寸），可以将重点放在 Stencil 计算本身。

全局内存、共享内存和 TMA 过滤实现器的相对性能：

上图对比了不同过滤半径下，TMA、共享内存 (Smem) 和全局内存 (Gmem) 的性能（相对 Smem 为 1.0x，更高表示更好）：
* 小半径过滤器 (2)：
* TMA: 1.04x
* Smem: 1.0x
* Gmem: 1.06x

大半径过滤器 (8)：
- TMA: 1.11x
- Smem: 1.0x
- Gmem: 0.67x

结论： TMA 代码更简单，指令更少，性能更高。

TMA 示例：3D TTI 逆时偏移 (Reverse Time Migration)

3D TTI 逆时偏移是一个计算密集型、寄存器占用高的地震应用。它已经通过 cuda::memcpy_async (GTC'21 A31115) 进行了加速。

进一步使用 TMA：

用张量描述符替换所有指针。
移除越界检查。
使用 TMA 异步加载和存储所有数据。

上图展示了使用 TMA 的 3D TTI 逆时偏移循环流程：包括 TMA 加载模型、等待、计算 Y/Z 驱动、同步、TMA 加载下一数据、计算其他驱动、等待、TMA 写入结果等步骤。

使用 TMA 带来的改进：
* 寄存器使用率降低 25%。
* 占用率提高 66%。
* 性能提升 1.65x。

集群中的生产者/消费者编程

Hopper 架构引入了新的异步事务屏障，支持“数据到达跟踪”。

上图展示了新的异步事务屏障流程：
1. 计算值并存储结果：线程递增 Arrival_count。
2. Arrive (到达) & Transaction (事务)：异步存储到共享内存，递增 Transaction_count。
3. 独立工作 (Independent Work)。
4. Wait (等待)：等待直到 Arrival_count 和 Transaction_count 都达到预期值。
5. 处理结果 (Process results)。

这是 cuda::memcpy_async 生产者/消费者模型中单向数据交换的构建块。

`cuda::memcpy_async` 共享内存 <-> 共享内存

在使用 cuda::memcpy_async Shared <-> Shared 进行生产者/消费者编程时，需要用生产者和消费者线程计数来初始化屏障。

上图展示了生产者（线程块 0，2 个线程）和消费者（线程块 1，2 个线程）通过共享内存进行交互的设置。屏障被初始化，例如“预期到达计数 = 4”，而“预期事务计数”在开始时为 0。

消费者等待生产者数据 (Page 46)

消费者线程抵达并等待生产者提供数据。图中展示了线程块0（生产者，2个线程）和线程块1（消费者，2个线程）的初始状态。消费者线程在各自的共享内存中调用 bar.arrive() 获取令牌，并随后调用 barrier.wait(token) 等待屏障解除。
异步SIMT编程：消费者等待数据 (Page 46)

生产者使用消费者屏障发送异步内存拷贝 (Page 47)

生产者线程使用 memcpy_async 向消费者的共享内存发起异步内存拷贝操作，同时更新消费者屏障的状态。
异步SIMT编程：生产者发送异步内存拷贝 (Page 47)

生产者线程1抵达 (Page 49)

生产者线程1也执行了 memcpy_async 和 barrier_arrive 操作，使得预期事务计数达到4。
异步SIMT编程：生产者1抵达 (Page 49)

所有线程和数据抵达，屏障解除 (Page 50)

当所有线程和数据都抵达后，屏障解除阻塞。图中显示预期抵达计数和预期事务计数均达到0，表示屏障已成功同步。
异步SIMT编程：所有线程和数据抵达，屏障解除 (Page 50)

生产者/消费者用例：Longstaff Schwartz 定价模型 (PRODUCER/CONSUMER USE CASE: LONGSTAFF SCHWARTZ PRICING MODEL)

Longstaff Schwartz 定价模型介绍

这是一种定量金融中的定价技术。

Longstaff-Schwartz 定价模型
金融风险计算涉及蒙特卡洛模拟。
Longstaff-Schwartz 技术：用于美式期权和其他可在任何时间行使的金融合约的估值。

算法和计算方法 (Page 53)

Longstaff Schwartz 定价模型将问题空间可视化为一个三维结构，包含“资产”、“路径”和“到期时间”维度。
Longstaff Schwartz 定价模型算法视图 (Page 53)

算法和计算方法 (Page 54)

对于每个资产，模型会沿着时间维度进行逆向传播（Propagation (reverse in time)），计算并更新现金流（Cash flows）。这个过程通过迭代完成。
Longstaff Schwartz 定价模型计算流程 (Page 54)

算法和计算方法 (Page 55)

Longstaff Schwartz 定价模型的计算流程包括以下步骤：
1. 初始化 (Initialization)
2. 跨N路径的归约 (Reduction across N paths)
3. 伪逆计算 (Pseudo-inverse computation)
4. 现金流更新 (Cash flow update)
5. 现值 (Present value)
其中，步骤2-4在一个时间步长的迭代中重复进行。
Longstaff Schwartz 定价模型完整算法流程 (Page 55)

当前实现 (Current Implementation)

当前实现存在内存限制：
* 仅 25,000 条路径可以放入共享内存，这可能太低了。
* 对于金融领域一个相关的用例，如果需要处理 100,000 条路径，当前的实现会将现金流存储在全局内存中。
* 评估 1260 个时间步，将启动 3780 个核函数（kernels）。

Longstaff Schwartz 定价模型当前实现瓶颈 (Page 56)
Longstaff Schwartz 定价模型全局内存及核函数启动 (Page 57)

H100 上的集群实现

在 H100 上的 Longstaff Schwartz 定价模型的集群实现中，处理 100,000 条路径时，每个现金流数组为 800KB。这无法在一个块的共享内存中容纳，但可以容纳在 4 个块中。

Longstaff Schwartz 定价模型H100集群实现 (Page 58)

整个集群处理流程如下：
* Thread Block 0, 1, 2, 3 各自执行对其 N/4 路径的归约 (Reduction across N/4 paths)。
* 归约结果汇集到 Thread Block 0 的线程 0 (Reduce to Thread 0, Block 0)。
* 通过 cluster.sync() 进行集群同步。
* 执行伪逆计算并分散值 (Pseudo-Inverse computation. Scatter values)。
* 再次通过 cluster.sync() 进行集群同步。
* 最终，各个线程块更新现金流 (Cash flow update)。

Longstaff Schwartz 定价模型H100集群同步 (Page 59)
Longstaff Schwartz 定价模型H100集群伪逆计算 (Page 60)
Longstaff Schwartz 定价模型 - H100 集群实现概览 Page 61

在评估 1260 个时间步时，这种方法只需启动 1 个内核，而传统方法可能需要 3780 个内核。

Longstaff Schwartz 定价模型 - H100 集群实现（1260 时间步） Page 62

H100 上的集群实现，采用单边通信

通过使用单边屏障通信（one-sided barrier communication）而非完整的集群同步，可以优化流程。

优化后的流程如下：
* Thread Block 0, 1, 2, 3 各自执行对其 N/4 路径的归约。
* 归约结果汇集到 Thread Block 0 的线程 0。
* Thread Block 0 执行 barrier.wait()。
* 在 1260 个时间步的循环中：
* Thread Block 0 执行伪逆计算并分散值。
* 所有线程块执行 barrier.wait()。
* 所有线程块更新现金流。

Longstaff Schwartz 定价模型 - H100 单边通信集群实现 Page 63

H100 上的性能表现

H100 上更大的分布式共享内存容量使得编写高效内核成为可能。
现在可以在共享内存中实现线程块级别的生产者-消费者模型，而无需全局内存同步。
硬件加速的集群同步易于使用且高效。

通过对 Longstaff Schwartz 定价模型的性能测试，结果显示：
* 未采用集群的 H100 (H100 w/o clusters) 作为基准。
* 采用集群并使用 cluster.sync() (H100 w/ Cluster + cluster.sync()) 的吞吐量达到基准的 2.7 倍。
* 采用集群并使用单边屏障 (H100 w/ Cluster + Barrier) 的吞吐量进一步提升 10%。

Longstaff Schwartz 定价模型在 H100 上的性能 Page 65

HOPPER 架构

结论与主要要点

Hopper 架构的关键特性包括：

线程块集群 (Thread Block Cluster)
- 利用 GPU 硬件层级结构。
- 更大的分布式共享内存，在 H100 上高达 3.6MB。
- 降低了块间通信的延迟。
异步 SIMT 编程模型 (Asynchronous SIMT Programming Model)
- cuda::memcpy_async 在所有方向上都得到硬件加速。
  - A100 仅支持全局内存到共享内存的传输。
  - 新的架构支持共享内存到全局内存以及归约操作。
  - 支持共享内存之间的传输。
- 在 CUDA 内核级别实现内存操作和计算的重叠。
- TMA (Tensor Memory Accelerator) 简化了高性能代码的编写。

另请参阅：

Inside the NVIDIA Hopper Architecture (S42663)
CUDA: New Features and Beyond (S41486)