Optimizing Memory Bandwidth and Latency on Hopper + Blackwell

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

议程 (Agenda)

硬件趋势 (Hardware Trend)
理解内存带宽 (Understanding Memory Bandwidth)
最大化内存带宽 (Maximizing Memory Bandwidth)
针对小问题规模最小化延迟 (Minimizing Latency for small problem sizes)

硬件发展趋势

硬件正在发生什么？

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

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

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

每SM带宽增加带来的影响

随着每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)

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

自动扶梯的比喻

利特尔法则: 系统中的平均单元数 = 平均到达率 * 平均驻留时间
扶梯规格:
- 每级台阶站1人
- 共20级台阶
- 每2秒到达1人
- 峰值到达率 = 0.5 人/秒
- 驻留时间 = 40 秒
问题: 当扶梯上只有1个人（in-flight）时，实现的吞吐量是多少？
计算: 吞吐量 = 人数 / 驻留时间 = 1 / 40 = 0.025 人/秒
- 这个吞吐量远低于0.5人/秒的峰值吞吐量，说明系统（扶梯）未被充分利用。

利特尔法则在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架构的在途字节数需求

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

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

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

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

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

能否产生足够的在途字节数？
估算每SM的在途字节数:
# loads / thread * # bytes / load * # threads / block * # blocks / SM
= 2 * 4 * 256 * 8 = 16 KiB (假设100%占用率)

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

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

__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];
}

估算每SM的在途字节数:
# loads / thread * # bytes / load * # threads / block * # blocks / SM
= 3 * 4 * 256 * 8 = 24 KiB (假设100%占用率)

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

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

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

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

可用工具 (Tools at our disposal)

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

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

考虑一个典型的循环内核。在循环展开前，每次迭代包含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];
    }
}

每个线程的在途字节数估算: # loads / thread * # bytes / load
在这个例子中，每次迭代有load a和load b两个加载操作，总计8字节的在途数据。

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

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

展开后，等效于在一个迭代中处理两个元素，指令变为：
- load a[i1], load b[i1]
- load a[i2], load b[i2]
- mul a[i1], b[i1]
- store c[i1]
- mul a[i2], b[i2]
- store c[i2]
每个线程的在途字节数估算: # loads / thread * # bytes / load
现在每个（逻辑）迭代有4次加载操作，总计16字节的在途数据，有效地将bytes-in-flight翻倍。

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

向量化全局访问
- 需要对齐的数据。
- 宽度为64位或128位。
启用向量化的方法
1. 显式使用向量数据类型，例如 float2, float4。
2. 隐式通过将指针强制转换为向量指针。
3. 需要适当的对齐。

性能比较

这些技术的效果如何？

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

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

使用寄存器的弊端

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

所有先前的技术都以增加寄存器使用量为代价来增加在途字节数（bytes-in-flight）。
- 在途字节需要由寄存器来支持。
- 这可能导致寄存器溢出到本地内存（register spilling）。
新一代GPU需要更高水平的ILP/DLP（即更多寄存器）来饱和内存带宽。
用于计算的寄存器数量不足成为一个挑战。
- 对于寄存器密集的内核（kernel），这是一个难题。
- 可能导致低占用率（occupancy）和寄存器溢出。

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

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

异步内存拷贝机制

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

普通加载 (Normal Loads)
- 来源: 全局内存
- 目的地: 寄存器
- 大小: 每线程 4, 8, 16B
- 预期用途: 多个线程并行
- 可用性: 所有CUDA GPU
异步加载 (Async Loads)
- 来源: 全局内存
- 目的地: 共享内存
- 大小: 每线程 4, 8, 16B
- 预期用途: 多个线程并行
- 引入架构: Ampere
异步批量加载 (Async Bulk Loads)
- 来源: 全局内存
- 目的地: 共享内存
- 大小: 16B - 100+ KB
- 预期用途: 每个线程块一个线程
- 引入架构: Hopper

异步加载

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

异步数据拷贝跳过寄存器，直接进入共享内存。
- 为计算释放更多寄存器。
- 减少L1缓存流量。

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

同步拷贝: 全局内存 → L2缓存 → L1缓存 → 寄存器 → 记分板
异步拷贝 (L1旁路): 全局内存 → L2缓存 → 共享内存 → 记分板 (绕过了L1和寄存器)
异步拷贝 (L1访问): 全局内存 → L2缓存 → L1缓存 → 共享内存 → 记分板 (绕过了寄存器)

使用异步加载的示例

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

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

异步批量加载 (又名 TMA)

一次性加载大量数据。

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

异步拷贝 (L1旁路)
- 多个线程执行拷贝。
- 在线程作用域的流水线（thread-scope pipeline）中完成。
- 数据路径：全局内存 → L2 → 共享内存 → 记分板。
异步批量拷贝
- 1个线程发起拷贝。
- 在共享内存屏障（shared memory barrier）处完成。
- 数据路径：全局内存 → L2 → 共享内存 → 共享内存屏障。

异步批量加载示例

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

完成机制: 共享内存屏障 (shared memory barrier)。
代码流程:
1. 初始化 (INIT): 单个线程 (如 threadIdx.x == 0) 初始化一个 cuda::barrier。
2. 同步所有线程。
3. 触发 (FIRE): 单个线程启动 cuda::memcpy_async 来执行批量拷贝。
4. 等待完成 (WAIT FOR COMPLETION): 所有线程调用 bar.wait() 等待拷贝操作完成。
注意: 当源/目标地址是16字节对齐且大小是16的倍数时，将使用TMA；否则，将回退到同步拷贝。

异步加载总结

使用异步加载的概览

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

为获得最佳性能，全局内存（gmem）和共享内存（smem）都应始终首选128字节对齐。
*Async Bulk Tensor Loads 本次未涵盖。

优化指南

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

开始: 在途字节数（bytes-in-flight）是否足够？
- 是：无需操作。
- 否：进入下一步。
数据加载到何处？
- 寄存器（REG）：进行循环展开/向量化。
- 共享内存（SMEM）：进入下一步。
数据是否对齐？
- 4或8字节对齐：使用异步加载（Async Loads）。
- 16字节对齐：进入下一步。
数据块（tile）的大小是多少？
- < 1 KiB：使用异步加载（Async Loads）。
- > 1KiB 且 < 2KiB：可选择批量或非批量异步加载。
- > 2 KiB：使用异步批量加载（Async Bulk Loads）。

关键要点

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

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

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

小规模问题的挑战

对于一个简单的float4向量加法内核，当问题规模足够大时：
- 每个SM有足够的在途字节（64 KB）。
- 能够达到理论峰值带宽（SoL BW）。
但在较小的问题规模下无法达到SoL带宽。

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

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

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

我们能做什么？

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

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

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

目标：将性能曲线左移

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

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

哪些延迟？

内存延迟由硬件固定。
块（Block）启动延迟不是关键：
- 注：1 波（wave） = 148 个 SM * 每个 SM 64 KB = ~10 MB。
内核（Kernel）启动延迟会影响所有问题规模，是优化的重点。

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

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

问题设置

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

__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

使用 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% 的性能提升。

CUDA 编程指南: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#creating-a-graph-using-stream-capture
GTC: CUDA Techniques to Maximize Concurrency and System Utilization

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

2. Programmatic Dependent Launch (PDL)

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 的优势：
- 允许内核更早地启动：
- 在前一个内核的全局内存存储变得可见之前。

启用更多预取：
- 内核参数、常量库等。
为了保证正确性：
- 内核必须执行 cudaGridDependencySynchronize() 以使前一个内核的存储操作变得可见。
CUDA 编程指南: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#api-description
GTC: CUDA Techniques to Maximize Concurrency and System Utilization

结合 PDL 的性能

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

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

3. Programmatic Dependent Launch + TriggerProgrammaticLaunchCompletion (Early Exit)

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)：
- 下一个内核在前一个内核的所有块都退出后才启动。

现在 (Now)：
- 下一个内核在前一个内核的所有块都执行了 cudaTriggerProgrammaticLaunchCompletion 后就启动。
- 这使得下一个内核能够更早地启动。

结合 Early Exit 的性能

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

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

组合技术的总体影响

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

加速比 (左图)：
- 对于小规模问题，组合技术可实现高达 3倍的加速。
- 随着问题规模的增大，这种加速效果逐渐减弱。
带宽曲线 (右图)：
- 综合效果使得带宽曲线显著左移，在某些点上实现了 75% 甚至 125% 的性能提升。

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

小规模问题优化总结

硬件加速主要针对大规模问题实现。
对于小规模问题：
- 需要降低内核启动延迟。
- 使用的技术包括：
  - CUDA Graphs
  - Programmatic Dependent Launch (PDL)
  - Early Exit (提前退出)
软件优化在小规模问题上实现了高达 3倍的加速。

PaperCache

Optimizing Memory Bandwidth and Latency on Hopper + Blackwell

Optimizing Memory Bandwidth and Latency on Hopper + Blackwell

目录

议程 (Agenda)

硬件发展趋势

硬件正在发生什么？

每SM带宽增加带来的影响

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

自动扶梯的比喻

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

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

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

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

可用工具 (Tools at our disposal)

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

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

性能比较

使用寄存器的弊端

异步内存拷贝机制

异步加载

使用异步加载的示例

异步批量加载 (又名 TMA)

异步批量加载示例

异步加载总结

使用异步加载的概览

优化指南

关键要点

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

小规模问题的挑战

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

目标：将性能曲线左移

问题设置

1. CUDA Graphs

2. Programmatic Dependent Launch (PDL)

3. Programmatic Dependent Launch + TriggerProgrammaticLaunchCompletion (Early Exit)

组合技术的总体影响

小规模问题优化总结

Optimizing Memory Bandwidth and Latency on Hopper + Blackwell

目录

议程 (Agenda)

硬件发展趋势

硬件正在发生什么？

每SM带宽增加带来的影响

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

自动扶梯的比喻

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

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

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

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

可用工具 (Tools at our disposal)

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

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

性能比较

使用寄存器的弊端

异步内存拷贝机制

异步加载

使用异步加载的示例

异步批量加载 (又名 TMA)

异步批量加载示例

异步加载总结

使用异步加载的概览

优化指南

关键要点

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

小规模问题的挑战

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

目标：将性能曲线左移

问题设置

1. CUDA Graphs

2. Programmatic Dependent Launch (PDL)

3. Programmatic Dependent Launch + TriggerProgrammaticLaunchCompletion (Early Exit)

组合技术的总体影响

小规模问题优化总结

登录

注册

忘记密码

重发验证邮件