CUDA Techniques to Maximize Memory Bandwidth and Hide Latency [S72683]

Athena Elafrou, Sr. Developer Technology Engineer
Allard Hendriksen, Sr. Developer Technology Engineer
GTC, March 17th 2025

议程 (Agenda
"问题在于内存，笨蛋！" ("It's the Memory, Stupid
GPU 内存层次结构 (GPU Memory Hierarchy
最大化内存带宽 (Maximizing Memory Bandwidth
异步数据拷贝 (Asynchronous Data Copies
内存模型 (Memory model
低延迟集群同步 (Low-Latency Cluster Synchronization
CUDA 开发者会议 (CUDA Developer Sessions

议程 (Agenda)

"问题在于内存，笨蛋！" ("It's the Memory, Stupid!")
最大化内存带宽 (Maximizing Memory Bandwidth)
理解内存模型 (Understanding the Memory Model)
低延迟集群同步 (Low-Latency Cluster Synchronization)

"问题在于内存，笨蛋！" ("It's the Memory, Stupid!")

"It's the Memory, Stupid!"
- Richard Sites, Multiprocessor Report, 1996

</blockquote>

GPU 内存层次结构 (GPU Memory Hierarchy)

硬件趋势 (Hardware trends)

下图展示了典型的GPU内存层次结构，以及各级别内存的发展趋势。

每个GPU架构的演进趋势如下：
- 更多的流多处理器 (SMs) 以执行计算。
- 每个SM拥有更大的L1缓存/共享内存。
- 更大的L2缓存和更高的带宽。
- 更大的DRAM和更高的带宽。

在接下来的几页中，我们将分别关注L1/共享内存、L2和DRAM的发展。

首先，关注每个SM的L1缓存/共享内存。

硬件趋势 - 共享内存 (Hardware trends - Shared memory)

从Kepler到Hopper Blackwell架构，共享内存（Shared Memory）的大小在持续增加。增加共享内存可以减少访问全局内存的往返次数。

Kepler: 48 KiB
P100: 64 KiB
V100: 96 KiB
A100: 164 KiB
Hopper Blackwell: 224 KiB

Hopper Blackwell架构引入了分布式共享内存（Distributed Shared Memory）。

分布式共享内存 (Distributed shared memory)

Hopper Blackwell架构中的分布式共享内存允许线程块集群（Thread Block Cluster）之间高效地共享数据。这引出了一个关键问题：如何在线程块集群内高效地同步和交换数据？

硬件趋势 - L2缓存

接下来，关注L2缓存。每一代GPU架构都带来了更大的L2缓存和更高的带宽。

硬件趋势 - DRAM

最后，关注DRAM。每一代GPU架构都配备了更大的DRAM和更高的带宽。

硬件趋势 - DRAM 带宽

下图展示了从P100到B200架构，DRAM总带宽、SM数量以及每个SM的带宽的变化趋势。

总带宽 (Bandwidth GB/s): 增长迅速，大约提升了2.2倍。
SM数量 (# SMs): 增长缓慢，大约提升了1.1倍。
每个SM的带宽 (Bandwidth per SM GB/s): 持续增长，大约提升了2.0倍。

这引出了一个核心问题：如何饱和带宽 (How to saturate bandwidth)?

最大化内存带宽 (Maximizing Memory Bandwidth)

利特尔定律 (Little's Law)

利特尔定律指出：系统中的平均单元数 = 平均到达率 * 平均停留时间。

我们可以用一个自动扶梯的例子来类比：
- 扶梯规格:
- 每级台阶1人
- 高度为20级台阶
- 每2秒到达一级新台阶

计算得出:
- 峰值到达率 = 0.5 人/秒
- 停留时间 = 40 秒

问题: 如果扶梯上只有1个人（in-flight），实现的吞吐量是多少？
吞吐量 = #人数 / 停留时间 = 0.025 人/秒

问题: 我们需要多少人同时在扶梯上（in-flight）才能最大化吞吐量？
并发数 = 峰值到达率 * 停留时间 = 0.5 人/秒 * 40 秒 = 20 人

利特尔定律应用于GPU内存

将利特尔定律应用于GPU内存，公式为：
在途字节数 (bytes-in-flight) = 带宽 (bandwidth) * 平均延迟 (mean latency)

其中，在途字节数由软件控制，而平均延迟由硬件决定。

趋势:
- 每一代GPU，为了饱和DRAM带宽所需的在途字节数都在增加。
- 这主要是由于带宽的增加（例如，从Hopper到Blackwell大约增加2倍）。
- 同时，每个SM的带宽也在增加。
结论:
- 需要更多的在途字节数/SM来饱和带宽。
- 这需要我们相应地调整软件。

下图显示了峰值带宽（占理论峰值百分比）与每个SM的在途字节数之间的关系。可以看出，为了达到高带宽利用率，需要有足够多的在途字节数（例如，H200需要约48 KiB，GB200-NVL需要约64 KiB）。

利特尔定律应用于GPU内存的性能图 (Page 14)
*图中的点代表使用不同操作、线程块维度、数据类型和并行加载数量的类STREAM负载。

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

我们能否用简单的内核来饱和内存带宽？考虑以下简单的向量加法内核：

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

我们能产生足够的在途字节数吗？
估算每个SM的在途字节数：
estimated bytes-in-flight / SM = # loads / thread * # bytes / load * # threads / block * # blocks / SM
= 2 * 4 * 256 * 8 = 16 KiB
（假设达到100%占用率）

下图显示，尽管GPU的理论带宽（BW）不断提升，但对于这个简单内核，带宽利用率（BWUtil）却在下降，因为16 KiB的在途字节数不足以饱和新一代GPU的内存带宽。

我们再考虑一个稍复杂的内核: d[i] = a[i] + b[i] + c[i];。

在途字节估算 (3次加载):
- 估计的在途字节/SM = (# loads / thread) * (# bytes / load) * (# blocks / SM) * (# threads / block)
- = 3 * 4 * 8 * 256 = 24 KiB
- 此计算基于 100% 的占用率 (occupancy)。
在途字节估算 (3次加载 + 1次存储):
- 如果我们将存储操作也考虑在内，在途字节数会增加。
- 估计的在途字节/SM = (# memory ops / thread) * (# bytes / op) * (# blocks / SM) * (# threads / block)
- = 4 * 4 * 8 * 256 = 32 KiB
- 同样，此计算基于 100% 的占用率。

图表显示，从 V100 到 B200，尽管原始带宽（TB/s）持续增加，但对于这些简单的内核，峰值带宽利用率（BWUtil %）却在下降。这表明随着硬件的发展，简单内核越来越难以充分利用可用的内存带宽，即使我们增加了操作数。

增加在途字节 (Increasing Bytes-in-Flight)

有三种主要技术可用于增加在途字节（bytes-in-flight）：

线程内更多的独立内存操作 (指令级并行，Instruction-level parallelism)。
线程内向量化的内存操作 (数据级并行，data-level parallelism)。
异步数据拷贝 (Asynchronous data copies)。

增加指令级并行 (Increasing ILP)

以下示例展示了通过增加指令级并行（ILP）来增加在途字节。

基础情况: 一个简单的 c[i] = a[i] * b[i] 内核。
指令序列: load a, load b, mul a, b, store c。
在途字节估算: 每个线程有两个加载操作 (# loads / thread = 2)，产生 8 字节的在途数据。

循环展开

使用循环展开（Loop unrolling）可以增加 ILP。

方法: 在内核中使用 #pragma unroll 2 指令。
效果: 编译器将循环体展开，每个线程在一次迭代中处理两个元素。这使得加载操作（load a[i1], load b[i1], load a[i2], load b[i2]）可以独立地并发执行。
在途字节估算: 每个线程的加载操作数增加到 4，在途数据量增加到 16 字节。

循环展开的问题

循环展开并非总是能按预期工作。

问题: 编译器可能无法生成我们期望的代码，因为它不能安全地假设 i1 和 i2 的内存地址不重叠（指针别名问题）。
结果: 如果编译器不能排除别名，它可能会串行化内存操作，例如先完成 i1 的所有操作（加载、计算、存储），再开始 i2 的操作。这会阻碍 ILP 的提升。

手动循环展开

手动循环展开是一种获得最佳性能的解决方案。

方法: 使用模板和内层循环来显式地手动展开循环，从而向编译器明确指令的独立性。
效果: 这种方法可以帮助编译器生成更优的指令调度，将多个加载操作聚合在一起执行。
在途字节估算: 每个线程的加载操作数增加到 6，在途数据量增加到 24 字节。

增加数据级并行 (Increasing DLP)

通过向量化加载可以增加数据级并行（DLP）。

基线情况: 一个 warp（32 个线程）执行连续、对齐的 float 类型内存访问。这种访问模式通过单条指令获取 1 个缓存行（cache line）。

向量化加载

float2 访问: 一个 warp 执行连续、对齐的 float2 类型内存访问。这种模式下，单条指令可以获取 2 个缓存行。

float4 访问: 一个 warp 执行连续、对齐的 float4 类型内存访问。单条指令可以获取 4 个缓存行，进一步增加了数据级并行度。

实现方法

要求:
- 向量化的全局和共享内存访问需要数据对齐。
- 访问宽度通常为 64 位或 128 位。
实现途径:
1. 显式: 使用向量数据类型，如 float2、float4。
2. 隐式: 将指针强制转换为向量指针（需要确保正确的内存对齐）。
代码示例: 展示了使用 float2 类型的指针和数据访问的内核。这种方法可以产生 16 字节的在途数据。

性能比较

比较不同技术（循环展开和向量化）的有效性。

实验设置:
- 逐元素向量乘法。
- 向量大小为 4GiB。
比较的技术: unroll2（展开2次）、unroll4（展开4次）、vec2（float2向量化）、vec4（float4向量化）。
结果:
- 图表显示了在 V100、A100、H100、H200 和 B200 GPU 上的带宽提升百分比。
- 所有技术都带来了性能提升，尤其是在较新的 GPU 架构（H200, B200）上更为显著。
- 在最新的 GPU 上，unroll4 和 vec4 通常能提供最大的性能提升。

寄存器使用

增加 ILP 和 DLP 带来了一个权衡：寄存器压力增加。

核心问题:
- 所有提高在途字节的技术都是以增加寄存器使用量为代价的。
- 在途字节需要由寄存器来支持。
- 这可能导致寄存器溢出（spilling）到局部内存，从而降低性能。
对新一代 GPU 的影响:
- 新的 GPU 需要更高水平的 ILP/DLP（即更多寄存器）来饱和内存带宽。
- 这给计算密集型内核留下的可用寄存器更少，可能导致占用率低或寄存器溢出。
图表分析:
- 图表显示峰值带宽（% SOL）与寄存器使用率（% of SM）的关系。
- 对于 H100、H200 等新架构，只有当寄存器使用率超过一定阈值（例如 30%-40%）时，才能接近峰值带宽。

异步数据拷贝 (Asynchronous Data Copies)

异步数据拷贝是一种可以绕过寄存器直接将数据拷贝到共享内存的技术。

机制: 异步数据拷贝跳过寄存器，直接将数据从全局内存移动到共享内存。
优势:
- 释放更多寄存器用于计算。
- 减少 L1 缓存流量。
- 减少 MIO（内存输入/输出）压力（指令更少）。
流程对比:
- 同步拷贝: 全局内存 -> L2 -> L1 -> 寄存器 -> 共享内存。
- 异步拷贝: 全局内存 -> L2 -> 共享内存（绕过 L1 和寄存器）。

允许内存传输与计算重叠

异步数据拷贝可以与计算操作重叠执行，从而隐藏内存延迟并增加在途字节。

基线情况 (Vanilla): 传统的串行执行模式：加载(i) -> 计算(i) -> 加载(i+1) -> 计算(i+1)... 在这种模式下，加载和计算之间没有重叠。

数据预取（距离为1）: 通过将下一次迭代（i=1）的数据加载（load）与当前迭代（i=0）的计算（comp）重叠，可以隐藏内存延迟。

数据预取（距离为2）: 通过增加预取距离，可以进一步增加重叠，隐藏更长的延迟。例如，预取距离为2时，当前迭代的计算可以与未来两次迭代的数据加载重叠。

启用生产者-消费者模式

异步数据拷贝天然支持生产者-消费者模式。生产者线程（producers）负责从全局内存（GMEM）加载数据到共享内存（SMEM），而消费者线程（consumers）则从共享内存中读取数据进行计算。这个过程可以流水线化，以实现高效的数据处理。

使用场景

CUDA 核函数通常采用一种模式：
- 从全局内存读取数据到共享内存。
- 对共享内存中的数据执行计算。
- 可能会将结果写回全局内存。
对于采用这种模式的核函数，仅仅切换到异步拷贝对延迟的改善可能很小。
- 除非有机会进行批量加载。
- 例如，在条件代码中有多个加载操作，可参考 [S62192] "Advanced Performance Optimization in CUDA"。
对于可以为未来迭代预取数据的迭代式核函数，其收益可能非常显著。
- 特别是对于低占用率、计算密集型的核函数。

概述

下表总结了不同内存空间之间的异步数据拷贝及其完成机制。

LDGSTS指令

LDGSTS 是 smem[sidx] = gmem[gidx] 的异步版本，支持一次性拷贝4、8或16字节。

两种模式:
- L1 BYPASS: 访问会使L1缓存无效或绕过L1。要求数据类型大小（sizeof(datatype)）和对齐方式为16字节。
- L1 ACCESS: 访问通过L1缓存。要求数据类型大小和对齐方式为4或8字节。
- 对于大多数GPU，如果满足要求，编译器将使用L1 BYPASS模式。

LDGSTS APIs

实现LDGSTS功能可以通过以下几组API：

Primitives API <cuda_pipeline.h>
- 共享内存屏障完成: __pipeline_memcpy_async(), __pipeline_arrive_on() 结合 __mbarrier_*()
- 线程局部完成: __pipeline_memcpy_async(), __pipeline_memcpy_commit() 和 __pipeline_memcpy_wait_prior()
libcudacxx API <cuda/barrier> 或 <cuda/pipeline>
- 共享内存屏障完成: cuda::memcpy_async() 结合 cuda::barrier
- 线程局部完成: cuda::memcpy_async() 结合 cuda::pipeline
Cooperative groups API <cooperative_groups/memcpy_async.h>
- cooperative_groups::memcpy_async() 结合 cooperative_groups::wait() 或 cooperative_groups::wait_prior()

代码示例

切换到异步拷贝以进行批处理计算

Primitives API 示例

以下代码展示了如何将一个标准的同步数据加载计算循环转换为使用 Primitives API 的异步版本。异步版本通过在计算前发起异步内存拷贝并等待其完成，实现了计算与数据传输的重叠。

Libcudacxx API 示例

使用 libcudacxx API 可以实现类似的功能。该API支持大于16字节的拷贝，并使用 cuda::aligned_size_t 帮助编译器进行优化。

数据预取

使用 `cuda::pipeline` (1/2): 序言 (Prologue)

数据预取通常分为两部分：序言和主循环。序言部分负责为第一次迭代预取数据。为了避免线程束分化（warp entanglement），producer_acquire() 和 producer_commit() 应该在收敛的代码路径中调用。

使用 `cuda::pipeline` (2/2): 主循环 (Main loop)

主循环中，在处理当前阶段的数据之前，会为下一个迭代预取数据。cuda::pipeline_consumer_wait_prior<1>(pipe) 用于等待当前阶段的数据拷贝完成。计算完成后，通过 pipe.consumer_release() 释放已使用的阶段。

多阶段数据预取

序言 (1/2)

通过使用多级缓存（Multi-Stage Buffering），可以隐藏更高的内存延迟。使用编译时常量 NUM_STAGES 可以确保编译器消除内部的簿记指令。预取距离等于 NUM_STAGES - 1。序言部分会加载所有流水线阶段的数据。

主循环 (2/2)

在主循环中，等待操作 cuda::pipeline_consumer_wait_prior<NUM_STAGES - 1>(pipe) 会一直等到除了最近的 NUM_STAGES - 1 个阶段外所有数据都加载完毕。然后进行计算，释放已消耗的阶段，并为 NUM_STAGES 次迭代之后的数据发起新的预取。

使用生产者-消费者模式 (1/2): 序言

这种模式下，可以指定一部分线程（例如，memcpy_threads）专门用于内存拷贝。每个线程拷贝16字节可以启用L1 BYPASS模式，以获得更好的性能。

使用生产者-消费者模式 (2/2): 主循环

在主循环中，计算步骤前后需要同步（__syncthreads()），以确保所有线程在计算开始前都能访问到最新的数据，并在计算结束后再进行下一次数据预取，避免数据覆盖。

我需要多少个阶段 (How Many Stages Do I Need?)

调整内核以达到目标在途字节数 (Tuning our kernel to reach a target bytes-in-flight)

为了确定流水线所需的阶段数量，可以使用以下公式来计算每个SM（Streaming Multiprocessor）的在途字节数（bytes in flight）：

公式分解如下：

#bytes in flight / SM: 每个SM的在途字节数。这是一个固定的目标值，对于Hopper架构 >32 KiB，对于Blackwell架构 >40 KiB。
#blocks / SM: 每个SM的线程块数量。这是一个占用率目标，与不使用分阶段加载时相同。
#threads / block: 每个线程块的线程数。固定为256。
#loads / thread: 每个线程的加载次数。这是一个可调参数，等于 2 * 阶段数。
#bytes / load: 每次加载的字节数。固定为4。

根据这个公式，对于Hopper架构，我们需要2个阶段；对于Blackwell架构，需要3个阶段。

性能分析：简单计算内核

对于一个简单的计算任务 compute(a, b) = a * b，在NVIDIA H100上的性能表现如下：

Vanilla (基准): 原始内核使用带有禁用展开的网格步长循环（grid-stride loop）。
2 stages: 使用2个阶段的流水线加载，相较于基准版本，时间从2.288ms减少到2.148ms，带宽利用率从84.00%提升到89.49%，获得了1.065倍的加速。在途字节数从16KiB增加到32KiB。
3 stages / 4 stages: 进一步增加阶段数，在途字节数相应增加，但执行时间和加速比没有明显改善。

性能分析工具NVIDIA Nsight Compute显示，该内核的主要瓶颈是Stall Long Scoreboard，即等待长延迟操作（如内存加载）完成。

性能分析：复杂计算内核

对于一个计算延迟更高的任务 compute(a, b) = sqrt(sqrt(a) / sqrt(b))，在NVIDIA H100上的性能表现如下：

当计算延迟增加时，尽管长计分板停滞（long scoreboard stalls）仍然是主要瓶颈，但增加在途字节数可以产生显著的影响。
- 2 stages: 相较于基准版本，获得了1.305倍的显著加速，带宽利用率从68.62%大幅提升至89.56%。
- 3 stages: 性能略有下降，但仍比基准版本快1.281倍。

Nsight Compute的分析再次确认Stall Long Scoreboard是主要的性能瓶颈。

张量内存加速器 (Tensor Memory Accelerator - TMA)

TMA是一种用于批量拷贝的高效异步数据传输机制。
- 两个编程模型:
- 一维连续数组的批量异步拷贝 (TMA 1D)。
- 多维数组的批量异步拷贝 (TMA ND)。
- 在 [S62192]: "Advanced Performance Optimization in CUDA" 中有广泛介绍。

编程模型是线程束统一的 (warp uniform):
- 从每个线程束的单个线程调用TMA操作更高效。
- 如果每个线程束中有多个线程处于活动状态，编译器将生成一个剥离循环（peeling loop）来顺序执行TMA操作。

下图展示了TMA在全局内存和共享内存之间传输数据的过程。
Page 51 - TMA架构示意图

TMA 1D (UBLKCP)

UBLKCP (Unified Bulk Copy) 是TMA一维拷贝的实现。

完成机制: 共享内存屏障 (Shared memory barrier)。
对齐要求:
- 源/目标指针必须是16字节对齐。
- 拷贝大小必须是16字节的倍数。
API: libcudacxx <cuda/ptx>
- cuda::memcpy_async() 结合 cuda::barrier 使用。
- cuda::device::memcpy_async_tx() 结合 cuda::barrier 使用。
- PTX汇编在cuda::ptx命名空间中公开，允许更细粒度的屏障同步。
也在Thrust中启用: Thrust::transform (CCCL)。

TMA 1D编程模式 (GMEM to SMEM)

异步拷贝的典型编程模式分为三个阶段：初始化(INIT)、触发(FIRE)和等待完成(WAIT FOR COMPLETION)。

使用cuda::memcpy_async:
- 使用单个线程启动异步拷贝。如果源/目标指针是16字节对齐且大小是16的倍数，则使用TMA；否则，它会回退到同步拷贝。
- 注意: 如果不满足大小和对齐要求，行为是未定义的。此功能仅在Hopper+架构上有效。
使用PTX内联汇编:
- 提供更底层的控制。代码结构类似，但使用ptx::cp_async_bulk等指令。
- 编译器问题: 编译器不知道 if (threadIdx.x == 0) 这个条件对于整个线程束是恒定的，因此可能会为单个线程生成一个剥离循环，影响效率。
- 解决方案: 使用cooperative_groups::invoke_one来明确告知编译器，在线程束中只有一个活动的线程将执行TMA操作，从而避免生成不必要的代码。

下图展示了如何将一个标准的批处理计算内核重构为使用异步拷贝的模式，从而实现计算和数据传输的重叠。

零开销异步拷贝 (Zero-Effort Async Copies)

使用 `Thrust::transform`

Thrust库提供了一种更简单的方式来使用异步拷贝，几乎不需要手动管理。

代码转换: 一个标准的CUDA内核可以被一个thrust::transform调用替代。
启用TMA: 通过调用cuda::proclaim_copyable_arguments，可以告知Thrust lambda函数的参数数据可以被拷贝到共享内存，从而启用TMA。
自动调优: Thrust::transform会根据lambda函数的特性在内部进行自动调优，以最大化在途字节数。

数据预取 (Data Prefetching)

使用 TMA 1D

以下代码片段展示了使用一维张量内存加速器 (TMA 1D) 进行数据预取。

序言: 内核开始时，初始化共享内存屏障。在首次迭代中，由一个单独的线程发起异步内存拷贝操作。这里使用了 cooperative_groups 来协调线程束执行 cuda::device::memcpy_async_tx。拷贝操作的完成通过共享内存屏障进行管理。

主循环: 在一个循环中，单个线程为下一次迭代发起异步数据拷贝。在计算当前阶段的数据之前，代码使用 cuda::ptx::mbarrier_try_wait_parity 等待当前阶段的数据准备就绪。数据到达后，线程执行计算。在进入下一次迭代之前，使用 __syncthreads() 同步块内的所有线程。

使用异步拷贝 (Using Asynchronous Copies)

下表总结了不同异步拷贝机制的对齐约束和额外优势。

LDG+STS: 没有对齐约束。
LDGSTS: 需要 4、8 或 16 字节对齐。其优势在于可以批量处理拷贝，从而增加在途字节数（bytes-in-flight）。
TMA 1D: 需要 16 字节对齐。其优势在于批量拷贝可以减少指令数量。
TMA ND: 共享内存（SMEM）需 128 字节对齐，全局内存（GMEM）需 16 字节对齐，并且步长（strides）必须是 16 字节的倍数。额外优势包括支持共享内存的交错（swizzling）能力和越界处理能力。

建议:
- 优先选择 TMA 来拷贝较大数据块。
- TMA 指令的延迟高于 LDGSTS，因此需要更多数据来分摊其成本。

优化指南 (Optimization Guidelines)

该流程图为选择合适的内存优化策略提供了指导。

检查在途字节数 (bytes-in-flight):
- 如果已经足够，则无需优化。
- 如果不足，则考虑预取。
选择预取目标:
- 寄存器 (REG): 进行循环展开或向量化。
- 共享内存 (SMEM): 检查数据对齐。
基于对齐和数据块大小选择指令:
- 4 或 8 字节对齐: 使用 LDGSTS。
- 16 字节对齐: 根据数据块（tile）大小决定：
  - < 1 KiB: 使用 LDGSTS。
  - > 1KiB 且 < 2KiB: 使用 LDGSTS 或 TMA。
  - > 2KiB: 使用 TMA。

关键要点 (Key Takeaways)

要点 #1:

最新的 GPU 每个 SM（Streaming Multiprocessor）拥有更高的可用带宽。
根据利特尔法则（Little's Law），可实现的带宽取决于在途字节数（bytes-in-flight）。
为了充分利用带宽，可能需要通过软件层面的更改来增加在途字节数。

要点 #2:
- CUDA 提供了异步数据拷贝机制（LDGSTS 和 TMA），这些机制不占用额外的寄存器。
- 使用这些特性编写内核会更复杂。
- 在某些情况下，我们可以利用库来“免费”启用 TMA。

内存模型 (Memory model)

什么是内存模型？

内存模型决定了从内存加载（load）操作可以返回的值。
它是用户与编译器/硬件/语言之间的一个契约。
对于在并行系统中线程间的通信和同步至重要。

四个主题

本节将涵盖内存模型的四个主题，分别对应不同的 GPU 架构演进：
- 单线程 (Single thread)
- 多线程 (Multi-thread): Volta 架构
- 异步线程 (Async thread): Ampere 架构
- 异步代理 (Async proxy): Hopper 架构

单线程加载和存储

对于单个线程：
- 存储（store）操作对执行该存储的线程是可见的。
- 对同一地址的加载和存储操作会保持其顺序。这被称为同地址排序 (same-address ordering)。
- 如下图代码所示，对 val 的写入和读取操作不会被重排，因此断言 assert(val == 42) 总是成立。
- 问题: 这种排序规则是否总是成立？是否存在例外？

非一致性：常量缓存

同地址排序并不总是成立。
- 对于常规的加载和存储，缓存是保持一致的。
- 但在某些情况下会存在非一致性 (non-coherence)，常量缓存 (constant caches) 就是一个例子。

常量缓存的工作方式:
- 常量缓存与 L2 缓存有链接，但此链接独立于 L1 缓存。
- L1 缓存和常量缓存之间没有通信，因此它们之间的数据不是相互一致的。

以下代码展示了常量缓存可能导致的非一致性问题。
- 一个 __constant__ 变量 val 被修改。
- 即使 val 在全局内存（通过 L1/L2 路径）中被更新为 42，后续对 val 的加载操作可能会命中常量缓存，从而返回一个过时的值（stale value）。
- 这可能导致 assert(val == 42) 失败。
- 问题: 在多线程并行的情况下，排序是如何工作的？

多线程内存顺序: Relaxed / Release / Acquire

"内存顺序指定了内存访问（包括常规的非原子访问）如何围绕一个原子操作进行排序。"

以下是四种内存顺序的比较：

顺序类型	描述
Sequentially consistent	- 加载和存储不能在原子操作之前或之后移动。 - 在单线程内保持同地址排序。 - 易于编程，但对硬件而言速度较慢。
Acquire	- 加载和存储不能移动到 `acquire` 操作之前。 - 在单线程内保持同地址排序。
Release	- 加载和存储不能移动到 `release` 操作之后。 - 在单线程内保持同地址排序。
Relaxed	- 加载和存储可以在原子操作之前或之后移动。 - 在单线程内保持同地址排序。

顺序一致 (Sequentially consistent)

这是最强的内存顺序。
任何加载和存储操作都不能跨越此原子操作。之前的访问必须在它之前完成，之后的访问必须在它之后开始。
如下图所示，Prior load 和 Later load 都不能被重排到 a.load 的另一侧。

获取 (Acquire)

Acquire 语义创建了一个单向的屏障。
它确保在 acquire 操作之后的加载和存储不会被重排到该操作之前。
这常用于确保在一个线程中读取由另一个线程“发布”的数据时，能看到所有相关写入。

释放 (Release)

Release 语义也创建了一个单向的屏障。
它确保在 release 操作之前的加载和存储不会被重排到该操作之后。
这常用于“发布”数据，确保所有相关的写入在 release 操作完成前对其他线程可见。

松散 (Relaxed)

加载和存储可以在 relaxed 操作之前或之后移动。
在单个线程内保留相同地址的排序。

内存模型：线程作用域

CUDA C++ 作用域：thread, block, device, system

作用域定义了哪些线程可以观察到当前线程的加载和存储操作。

cuda::thread_scope_thread (线程): 只有本地线程可以观察到加载和存储。
cuda::thread_scope_block (线程块): 线程块中的其他线程可以观察到该线程的加载和存储。
cuda::thread_scope_device (GPU设备): 设备(GPU)中的其他线程可以观察到该线程的加载和存储。
cuda::thread_scope_system (系统): 系统中的其他线程（CPU、其他GPU、其他节点）可以观察到该线程的加载和存储。

CUDA PTX 作用域：block, cluster, device, system

PTX（并行线程执行）指令集体系结构有其自身的作用域定义。

Thread (线程): 只有本地线程可以观察到加载和存储。
.cta (线程块 - Thread Block): 线程块中的其他线程可以观察到该线程的加载和存储。
.cluster (线程块集群 - Thread block Cluster): 线程块集群中的其他线程可以观察到该线程的加载和存储。
.gpu (GPU设备): 设备(GPU)中的其他线程可以观察到该线程的加载和存储。
.sys (系统): 系统中的其他线程（CPU、其他GPU、其他节点）可以观察到该线程的加载和存储。

内存模型：GPU内存层级与作用域

作用域与内存层级的关系

每个作用域都有一个关联的一致性点（point of coherency），它决定了在该作用域内，内存操作在何处变得对其他线程可见。

Block scope (块作用域): 一致性点是 L1 缓存。同一线程块内的线程通过L1缓存（或共享内存）实现数据同步。
Cluster scope (集群作用域): 一致性点是 L2 缓存。同一集群内的线程块通过L2缓存同步。
Device scope (设备作用域): 一致性点是 L2 缓存。同一GPU设备上的所有线程通过L2缓存同步。
System scope (系统作用域): 一致性点是 L2 缓存 + 连接的缓存。跨设备（例如，通过NVLINK连接的GPU或通过PCIe连接的CPU）的同步点，涉及各自的L2缓存以及它们之间的连接。

内存模型：多线程示例

块内松散原子操作示例 (Relaxed block example)

场景: 一个生产者线程（block0_thread0）和一个消费者线程（block0_thread1）在同一个线程块内，通过共享内存（SMEM）中的变量 val 进行通信。
代码:
- 生产者使用 val.store(42, cuda::memory_order_relaxed) 写入值。
- 消费者在一个循环中通过 tmp = atom_v.load(cuda::memory_order_relaxed) 读取 val，直到值不为-1。
结果: 此操作成功。因为两个线程在同一个块中，它们的作用域是 thread_scope_block，一致性点是L1缓存/共享内存。消费者最终会读取到生产者写入的值 42。

跨块松散原子操作的失败示例 (A not so relaxed device example)

场景: 生产者（block0_thread0）和消费者（block1_thread0）位于不同的线程块。它们通过L2/DRAM中的变量val通信，但代码中指定的作用域仍为thread_scope_block。
潜在失败: 消费者对val的加载操作可能会命中其本地的L1缓存，而L1缓存中可能没有生产者写入的新值。由于作用域被限制在块级别，没有机制强制更新或使消费者的L1缓存无效。这可能导致消费者永远无法读取到新值，assert(tmp == 42) 失败。

跨块松散原子操作的正确示例 (A relaxed device example)

场景: 与前一个示例相同，但作用域已更正为 thread_scope_device。
代码: 生产者和消费者都使用 cuda::atomic_ref<int, cuda::thread_scope_device> val。
结果: 此操作成功。通过将作用域提升到设备级别，加载和存储的一致性点变为L2/DRAM。这确保了生产者的写入操作对设备上的所有其他线程（包括消费者）都是可见的。

使用标志位的松散原子操作失败示例

当需要同步多个值时，仅使用松散原子操作可能会引入问题。

L1缓存导致的问题
$Page 87$
- 场景: 生产者需要更新一个值 val 和一个标志 flag。消费者等待 flag 被设置，然后读取 val。所有操作都是 relaxed。
- 潜在失败: 消费者可能成功读取到 flag 的更新值（假设该读取操作命中了L2），但在读取 val 时，却命中了其本地L1缓存中的旧值。这导致断言失败。
乱序观察导致的问题
- 场景: 与上一个类似。
- 潜在失败: 即使 val 和 flag 的读写都命中了L2/DRAM，relaxed 内存顺序也不保证操作的顺序。消费者可能会观察到 flag 的更新先于 val 的更新，即使在生产者的代码中 val 的更新在 flag 之前。这同样会导致消费者读取到旧的 val 值。

Release-Acquire 模式

场景: 使用 Release-Acquire 模式来解决上述带标志位的同步问题。
代码:
- 生产者在更新 val 之后，使用 flag.store(1, cuda::memory_order_release) 来设置标志。
- 消费者使用 while (flag.load(cuda::memory_order_acquire) == -1) 来等待标志。
结果: 此操作成功。release 操作确保在它之前的所有内存写入（如此处的 val = 42）对其他线程可见。acquire 操作确保在它之后的所有内存读取都能看到由匹配的 release 操作同步的数据。这保证了 val 和 flag 的更新被消费者按顺序观察到。

Relaxed 与 Release-Acquire 的比较

Relaxed (松散):
- 更快: 只是在一致性点对缓存进行单次存储或加载。
- 无序: 不提供相对于其他读写的排序保证。
- 用途: 适用于两个线程交换单个值的情况。
Release + Acquire (释放+获取):
- 更慢: 需要将数据刷新到一致性点和/或使缓存无效。
- 有序: 提供相对于其他读写的排序保证。
- 用途: 适用于多个线程交换多个值的情况。

内存模型：异步线程 (Async threads)

异步线程：一个激励性示例

PTX 指令 st.async 的作用是将一个值存储到集群中另一个块的分布式共享内存（Distributed Shared Memory）中。一旦存储完成，它会更新另一个块的共享内存中的一个共享内存屏障（shared memory barrier）。

PTX 指令: ptx::st_async(remote_addr, 42, remote_bar);

然而，这种异步操作存在一个问题：后续的加载（load）或存储（store）操作可能会提前执行，从而违反了同一地址的顺序性（same-address ordering）。如下图所示，对 remote_addr 的加载操作可能会在 st.async 存储操作完成之前执行，导致数据竞争。

解决方案与模型

PTX 指令 st.async 相对于其后的加载或存储操作，不遵守同一地址顺序。这个问题的解决方法是：

该指令被标记为 "async"（异步）。
它的行为模式如同它派生了一个独立的线程来执行存储和 mbarrier 更新操作。
因此，多个线程读取相同的值会构成数据竞争（data race）。
适用于多线程的常规推理方式在这里同样适用。

内存模型：异步代理 (Async proxies)

代理（Proxies）代表了这样一种情况：从单个线程到单个物理内存位置存在多条不同的路径，而这些路径之间没有一致性/窥探（coherence/snooping）机制。

通用代理 (Generic proxy): 所有常规的加载和存储都通过此代理进行。
异步代理 (Async proxy): 一条由 TMA 单元、张量核心（tensor cores）以及其他一些指令使用的不同路径。
在通用代理的加载/存储和异步代理的加载/存储之间，不存在同一地址顺序。这意味着常规的存储操作可能会在异步代理的加载操作之前“竞争”执行。

异步代理：共享内存示例

以下代码展示了一个数据竞争的例子。对共享内存 smem 的存储是通过通用代理进行的，而后续使用 ptx::cp_async_bulk 从共享内存复制到全局内存的操作（本质上是对 smem 的加载）是通过异步代理（TMA）进行的。这可能导致从 smem 的加载操作在对它的存储操作之前执行。

#include <cuda/ptx>
namespace ptx = cuda::ptx;

__device__ float4 gmem;

__global__ void kernel() {
  __shared__ float4 smem;

  // Store value to shared memory
  // (generic proxy)
  smem = {42., 42., 42., 42.};

  // Copy from shared to global memory
  // (async proxy)
  ptx::cp_async_bulk(
    ptx::space_global, ptx::space_shared,
    gmem, &smem, sizeof(smem)
  );
}

为了解决这个问题，需要在两个代理操作之间插入一个栅栏（fence）。ptx::fence_proxy_async 指令可以确保代理之间的加载和存储顺序。

// ...
  // Store value to shared memory
  // (generic proxy)
  smem = {42., 42., 42., 42.};
  // Fence between proxies
  ptx::fence_proxy_async(ptx::space_shared);
  // Copy from shared to global memory
  // (async proxy)
  ptx::cp_async_bulk(/*...*/);
// ...

异步代理：自动跨代理栅栏

在某些情况下，栅栏是自动插入的。例如，当一个异步代理操作（如 cp.async.bulk）之后跟着一个等待屏障的操作（如 mbarrier_try_wait），屏障的等待操作会隐式地创建一个跨代理的栅栏。这确保了在屏障状态翻转之前，所有先前的内存操作都已完成，从而保证了后续通用代理加载操作的顺序性。

异步线程与异步代理指令总结

下表总结了哪些指令属于异步线程模型，哪些属于异步代理模型。

异步线程 (Async thread)：与其他通用加载/存储操作之间存在同一地址顺序，但与其他异步线程操作之间存在数据竞争。
异步代理 (Async proxy)：与其他通用加载/存储操作之间没有同一地址顺序，与其他异步代理操作之间也没有同一地址顺序。

低延迟集群同步 (Low-Latency Cluster Synchronization)

关键点

对于一个集群（Cluster）而言，其一致性点（point of coherency）是 L2 缓存。
任何具有集群作用域（cluster scope）的释放-获取（release-acquire）模式都需要一次到 L2 的往返通信，这非常耗时。
为了降低延迟，需要避免到 L2 的往返通信。

两种在集群中同步线程的方式

协作组 (Cooperative groups):
- 使用 cluster::sync()。
- 必须由集群中的所有线程执行。
- 总是使其之前的加载/存储对集群中的其他线程可见。
- 速度慢：需要到 L2 的往返通信。
CUDA PTX:
- 使用 ptx::barrier_cluster_arrive 和 ptx::barrier_cluster_wait 分离到达和等待阶段。
- arrive 必须由所有线程执行。
- 可以选择性地使加载/存储可见：sem_release 使其可见，而 sem_relaxed 则不。

屏障初始化：简单方式

共享内存屏障是另一种同步机制。它们在使用前必须被初始化。初始化后的屏障必须对集群中的其他线程可见。简单的方法是使用 cluster::sync()，但这很慢。

屏障初始化：快速方式

为了避免 L2 的往返通信，可以采取以下措施：

执行一次宽松的（relaxed）集群到达 ptx::barrier_cluster_arrive(ptx::sem_relaxed)。
为了保持正确性，仅对 mbarrier 的初始化操作设置栅栏 ptx::fence_mbarrier_init(ptx::sem_release, ptx::scope_cluster)。
栅栏和屏障的集群等待（barrier_cluster_wait）共同构成了一个释放-获取模式，从而在避免L2往返通信的同时保证了正确性。

数据通信

同步通信 (Synchronous communication):
- 存储到远程分布式共享内存。
- 同步整个集群（例如使用 cluster::sync()）。
异步通信 (Asynchronous communication):
- 异步存储到远程共享内存。
- 在本地屏障上等待。这种方式延迟更低。

下面是一个使用 cluster::sync() 进行数据通信的简单基准测试示例。

集群中有2个块。
在每次迭代中，它们通信一个整数。
代码在一个循环中执行远程写入，然后调用 cluster::sync() 来等待其他集群成员。

使用 st.async() 通信数据

以下代码展示了使用 st.async() 和内存屏障（mbarrier）实现低延迟集群同步的 PTX 代码示例。

代码逻辑解析：
- for 循环：在多次迭代中执行通信。
- // Send value：使用 ptx::st_async 异步发送一个值（42）到远程内存地址 remote_val。
- // Arrive on local barrier：在本地内存屏障上执行 arrive 操作，并期望一次传输（expect_tx）。
- // Wait for value from other cluster：在一个 while 循环中，使用 ptx::mbarrier_try_wait 尝试在屏障上等待，直到从另一个集群接收到值。
- // Wait for other block to have received our value：使用 ptx::barrier_cluster_arrive 和 ptx::barrier_cluster_wait 来确保其他块已经收到了我们发送的值，这是一个集群范围的同步。

性能对比

此页面展示了异步与同步通信性能的基准测试对比。

简单基准测试 (Simple benchmark)：
- 集群中有 2 个块。
- 在每次迭代中，它们通信一个整数。
- 测试平台为 H100。

结果 (Result)：
- 同步（Synchronous）通信的性能为 1.3M 次迭代/秒。
- 异步（Asynchronous）通信的性能为 7M 次迭代/秒。
- 异步版本的速度比同步版本快 5倍以上。

结论 (Conclusion)：
- 避免在热循环（hot loops）中使用 cluster::sync()。
- 使用 st.async 来获得显著的加速。

CUDA 开发者会议 (CUDA Developer Sessions)

此页面列出了一系列与 CUDA 相关的开发者会议，涵盖了从入门到高级优化的多个主题。