Optimizing Memory Bandwidth and Latency on Hopper + Blackwell

Allard Hendriksen, Sr. Developer Technology Engineer
Beijing Open AI Day, May 2025

目录

  1. 议程 (Agenda
  2. 硬件发展趋势
  3. 理解内存带宽：利特尔法则 (Little's Law
  4. 如何增加在途字节数以提升带宽
  5. 针对小问题规模最小化延迟

议程 (Agenda)

硬件发展趋势

硬件正在发生什么？

Page 4 - 硬件发展趋势图表
Page 4 - 硬件发展趋势图表

上图展示了NVIDIA GPU几代架构的硬件发展趋势：
* 总带宽（GB/s）：从P100到H20，总带宽增长迅速，大约增长了2.2倍。
* SM数量（# SMs）：SM（流式多处理器）的数量增长相对缓慢，大约增长了1.1倍。
* 每SM带宽（Bandwidth per SM (GB/s)）*：由于总带宽增速远超SM数量增速，每个SM可用的带宽正在显著增加，大约增长了2.0倍。

核心问题是：如何充分利用（饱和）带宽？

*任何提供的基准测试数据仅用于技术讨论。

每SM带宽增加带来的影响

Page 5 - 带宽利用率图表
Page 5 - 带宽利用率图表

随着每SM可用带宽的增加，简单的内核（Kernel）越来越难以充分利用硬件的带宽潜力。

如上图所示的简单vectorAdd内核：

__global__ void kernel(float *a, float *b, float *c)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

该图表显示，在从V100到B200*的几代GPU中，虽然绝对带宽（BW，以TB/s计）持续提升，但带宽利用率（BWUtil，占峰值百分比）却在下降。这意味着简单的程序无法产生足够的内存请求来“喂饱”现代GPU。

*任何提供的基准测试数据仅用于技术讨论。

理解内存带宽：利特尔法则 (Little's Law)

利特尔法则是一个用于排队论的普适公式，可以帮助我们理解系统吞吐量。

自动扶梯的比喻

利特尔法则在GPU内存中的应用

Page 8 - 利特尔法则与GPU内存带宽
Page 8 - 利特尔法则与GPU内存带宽

将利特尔法则应用于GPU内存系统：

在途字节数 (bytes-in-flight) = 带宽 (bandwidth) * 平均延迟 (mean latency)

为了饱和DRAM带宽，需要有足够的“在途字节数”。随着每一代GPU的发展，这个需求也在增加：
* 主要原因是带宽的增长。
* 从Hopper到Blackwell架构，所需的在途字节数大约增加了2倍。
* 同时，每SM的带宽也在增加，因此需要为每个SM提供更多的在途字节数来饱和带宽。

上图显示了不同GPU（H100, H200, GB200-NVL）上，峰值带宽百分比与每SM在途字节数（Bytes in flight / SM）的关系。可以看出，要达到接近峰值的带宽，需要更多的在途字节数。

*任何提供的基准测试数据仅用于技术讨论。

不同GPU架构的在途字节数需求

Page 9 - H100, H200, B200在途字节数对比
Page 9 - H100, H200, B200在途字节数对比

上图详细对比了NVIDIA H100、H200和B200三款GPU。
* 对于H100，大约需要32-40 KiB的在途字节数/SM才能接近饱和。
* 对于H200和B200，则需要大约64 KiB的在途字节数/SM才能达到相似的饱和水平。
* 结论是：H200需要的在途字节数比H100多，与B200*大致相同。

*任何提供的基准测试数据仅用于技术讨论。

简单内核能否饱和内存带宽？

Page 10 - 简单内核的在途字节数估算 (1)
Page 10 - 简单内核的在途字节数估算 (1)

我们回头看之前的简单内核：

__global__ void kernel(float *a, float *b, float *c)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

16 KiB的在途数据量对于现代GPU来说是不足的，这解释了为什么简单内核的带宽利用率低。

Page 11 - 简单内核的在途字节数估算 (2)
Page 11 - 简单内核的在途字节数估算 (2)

通过在内核中增加一次加载操作，可以增加在途字节数：

__global__ void kernel(float *a, float *b, float *c, float *d)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    d[i] = a[i] + b[i] + c[i];
}

在途字节数增加到了24 KiB，这有助于提升带宽利用率，但可能仍不足以完全饱和最新架构的GPU。

*任何提供的基准测试数据仅用于技术讨论。

如何增加在途字节数以提升带宽？

有多种方法可以增加在途字节数。

可用工具 (Tools at our disposal)

  1. 线程内更多的独立内存操作: 在单个线程中执行更多的加载/存储指令。
  2. 线程内向量化的内存操作: 使用float2, float4等向量类型一次性读写更多数据。
  3. 异步数据拷贝: 利用硬件特性实现计算与数据传输的重叠。

增加指令级并行度（ILP）：循环展开

Page 14 - 循环展开前
Page 14 - 循环展开前

考虑一个典型的循环内核。在循环展开前，每次迭代包含2次加载操作。
* 代码示例:

__global__
void kernel(int n, 
            const float * __restrict__ a, 
            const float * __restrict__ b, 
                  float * __restrict__ c) 
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = tid; i < n; i += stride) {
        c[i] = a[i] * b[i];
    }
}
Page 15 - 循环展开后
Page 15 - 循环展开后

通过使用 #pragma unroll 2 进行循环展开，编译器会将循环体复制一次，从而增加指令级并行度。
* 代码示例:

#pragma unroll 2
for (int i = tid; i < n; i += stride) {
    c[i] = a[i] * b[i];
}

增加数据级并行度（DLP）：使用向量化加载

Page 16
Page 16

性能比较

这些技术的效果如何？

下图展示了在不同GPU架构上，循环展开（unroll）和向量化（vec）对元素级向量乘法（向量大小4GiB）所带来的带宽提升百分比。可以看出，随着GPU架构的演进（从V100到B200），这些技术带来的性能提升越来越显著。

Page 17
Page 17

注：所有基准测试数据仅供技术讨论之用。

使用寄存器的弊端

提高指令级并行度（ILP，通过循环展开）和数据级并行度（DLP，通过向量化）会增加寄存器压力。

下图显示，为了达到峰值带宽（SoL Bandwidth），新一代GPU（如B200）需要比前代（如H100）使用更高比例的寄存器。例如，在B200上达到SoL带宽所需的寄存器比H100多40%。

Page 18
Page 18

注：所有基准测试数据仅供技术讨论之用。

异步内存拷贝机制

在Ampere和Hopper架构中引入了新的内存拷贝机制。

Page 19
Page 19

异步加载

异步加载可以跳过寄存器，直接将数据加载到共享内存。

下图比较了同步拷贝和两种异步拷贝的数据流路径：

Page 20
Page 20

使用异步加载的示例

异步加载可以像同步操作一样使用。

下面的代码示例展示了如何将一个标准的同步内核（左侧）转换为使用异步加载的内核（右侧）。主要步骤包括：
1. 包含 <cuda/pipeline> 头文件。
2. 定义共享内存缓冲区。
3. 创建一个 cuda::pipeline 对象。
4. 使用 cuda::memcpy_async 启动异步拷贝。
5. 使用 pipe.producer_commit() 提交生产者阶段。
6. 使用 cuda::pipeline_consumer_wait_prior 等待拷贝完成。
7. 使用共享内存中的数据进行计算。

Page 21
Page 21

异步批量加载 (又名 TMA)

一次性加载大量数据。

异步批量加载（又称Tensor Memory Accelerator, TMA）与普通的异步拷贝在机制上有所不同：

Page 22
Page 22

异步批量加载示例

以下是使用 cuda::memcpy_async 和 cuda::barrier 实现异步批量加载的示例。

Page 23
Page 23

异步加载总结

使用异步加载的概览

下表总结了不同加载类型的对齐约束和额外优势。

Page 25
Page 25

优化指南

以下流程图可用于指导选择合适的优化策略。

Page 26
Page 26

  1. 开始: 在途字节数（bytes-in-flight）是否足够？

    • 是：无需操作。
    • 否：进入下一步。

  2. 数据加载到何处？

    • 寄存器（REG）：进行循环展开/向量化。
    • 共享内存（SMEM）：进入下一步。

  3. 数据是否对齐？

    • 4或8字节对齐：使用异步加载（Async Loads）。
    • 16字节对齐：进入下一步。

  4. 数据块（tile）的大小是多少？

    • < 1 KiB：使用异步加载（Async Loads）。
    • > 1KiB 且 < 2KiB：可选择批量或非批量异步加载。
    • > 2 KiB：使用异步批量加载（Async Bulk Loads）。

关键要点

Page 27
Page 27
  1. H100, B200 拥有更高的 每个SM的带宽。
  2. 这需要更多的 在途字节 来饱和带宽。
  3. 通过循环展开/向量化，这需要更多的 寄存器。
  4. 解决方案是：使用 异步加载 将在途字节转移到 共享内存 中，从而释放寄存器。

注：所有基准测试数据仅供技术讨论之用。

针对小问题规模最小化延迟

小规模问题的挑战

下图展示了NVIDIA B200的DRAM带宽随传输字节数的变化。只有当数据量达到约100MB以上时，带宽才能接近峰值。

Page 29
Page 29

注：所有基准测试数据仅供技术讨论之用。

与H20的比较：
- 大规模问题：可以看到预期的约2倍加速。
- 中等规模问题：可以看到最高2倍的加速。
- 小规模问题：看不到加速。

我们能做什么？

下图比较了H20和B200在不同问题规模下的带宽表现。在强扩展区（大规模问题），B200性能约为H20的两倍。但在无扩展区（小规模问题，<1MB），两者性能几乎相同，没有体现出B200的优势。

Page 30
Page 30

注：所有基准测试数据仅供技术讨论之用。

优化策略：减少内核启动延迟

目标：将性能曲线左移

Page 32
Page 32

目标
- 在相同的问题规模下实现更高的带宽。

方法
- 减少总运行时间。
- 减少延迟。

哪些延迟？

图中曲线展示了在 NVIDIA B200 设备上，DRAM 带宽随传输字节数的变化。目标是将此曲线向左移动，意味着在处理较小数据量时也能达到高带宽。蓝点表示一个波次的线程块读取约 10MB 的数据。

任何基准测试数据仅供技术讨论之用。

问题设置

Page 33
Page 33

为了衡量基准性能，使用了一个简单的向量加法内核：

__global__ void k(float4 *a, float4 *b, float4 *c)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

实验设置：
- 将上述内核运行 1000 次。
- 轮换使用 a, b, c 指针，以避免命中 L2 缓存。
- 在不同数据规模下测量带宽。

右图展示了在此设置下测得的基准性能曲线。

任何基准测试数据仅供技术讨论之用。

1. CUDA Graphs

Page 34
Page 34

使用 CUDA Graphs 可以显著减少重复内核启动的开销。其工作流程分为捕获、创建、启动和清理四个阶段：

// Capture
cudaGraph_t g;
cudaGraphCreate(&g, 0);
cudaStream_t stream;
cudaStreamCreate(&stream);
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
for (int i=0; i<1000; ++i)
    kernel<<<grid, block, smem_size, stream>>>(params);
cudaStreamEndCapture(stream, &g);

// Create
cudaGraphExec_t gEx;
cudaGraphInstantiate(&gEx, g, nullptr, nullptr, 0));

// Launch
CUDA_CHECK(cudaGraphLaunch(gEx, stream));
CUDA_CHECK(cudaDeviceSynchronize());

// Cleanup
CUDA_CHECK(cudaStreamDestroy(stream));
CUDA_CHECK(cudaGraphExecDestroy(gEx));

如右图所示，使用 CUDA Graph 后，性能曲线明显左移，实现了约 50% 的性能提升。

任何基准测试数据仅供技术讨论之用。

2. Programmatic Dependent Launch (PDL)

Page 35
Page 35

Programmatic Dependent Launch (PDL) 是一种进一步减少延迟的技术。

内核代码修改：
在内核中添加 cudaGridDependencySynchronize() 以确保数据依赖的正确性。

__global__ void k(float4 *a, float4 *b, float4 *c)
{
    cudaGridDependencySynchronize();
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

启动代码：
使用 cudaLaunchKernelEx 并设置相应的属性来启用 PDL。

// Launch
cudaLaunchConfig_t config = {0};
config.gridDim = grid_dim;
config.blockDim = block_dim;
config.dynamicSmemBytes = smem_size;
config.stream = stream;

cudaLaunchAttribute attr[1];
attr[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attr[0].val.programmaticStreamSerializationAllowed = 1;
config.attrs = attr;
config.numAttrs = 1;

cudaLaunchKernelEx(&config, kernel, param0, param1, ..);

PDL 的优势：
- 允许内核更早地启动：
- 在前一个内核的全局内存存储变得可见之前。

结合 PDL 的性能

Page 36
Page 36

将 PDL 与 CUDA Graph 结合使用，可以进一步将性能曲线左移。如图所示，性能提升从 50% 增加到 70%。

任何基准测试数据仅供技术讨论之用。

3. Programmatic Dependent Launch + TriggerProgrammaticLaunchCompletion (Early Exit)

Page 37
Page 37

cudaTriggerProgrammaticLaunchCompletion 是对 PDL 的进一步增强。

内核代码修改：

__global__ void k(float4 *a, float4 *b, float4 *c)
{
    cudaGridDependencySynchronize();
    cudaTriggerProgrammaticLaunchCompletion();
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

cudaTriggerProgrammaticLaunchCompletion 的作用：
- 在一个块（block）真正退出之前，提前发出该块已退出的信号。
- 一个块中只需一个线程执行此操作即可。

对比：
- 通常情况 (Normally)：
- 下一个内核在前一个内核的所有块都退出后才启动。

结合 Early Exit 的性能

Page 38
Page 38

将 CUDA Graph、PDL 和 Early Exit (提前退出) 三种技术结合，性能得到进一步提升。如图所示，性能提升从 70% 增加到 75%。

任何基准测试数据仅供技术讨论之用。

组合技术的总体影响

Page 39
Page 39

结合上述所有技术，对性能的总体影响如下：

任何基-准测试数据仅供技术讨论之用。

小规模问题优化总结

Page 40
Page 40