内存延迟隐藏

高性能处理器经常花费更多时间等待数据，而非处理数据。虽然循环平铺等方法旨在提升数据重用性，确保数据一旦进入缓存便被多次使用，但从全局内存将数据获取到寄存器的最初动作仍会带来明显的延迟开销。如果算术逻辑单元（ALU）在等待内存请求完成时处于空闲状态，硬件就未能充分利用。

内存延迟隐藏是一系列优化方法，旨在使内存操作与计算重叠进行。目标是在处理器忙于执行当前迭代时，发出对未来迭代的内存请求。成功实现后，核的执行时间将从内存时间和计算时间之和转变为两者中的较大值。

内存墙

为认识延迟隐藏的必要性，请考虑硬件中的时间尺度差异。在GPU上，一次浮点乘法可能需要4到6个时钟周期。从全局内存获取一个浮点数可能需要400到800个周期。

在一种简单实现中，执行流程是严格串行的：

1. 发出加载指令。
2. 等待400+周期。
3. 执行数学指令。
4. 重复。

在第2步中，计算单元会停滞。为避免这种情况，编译器和硬件调度器会设法保持大量“在途”内存请求。通过识别独立指令，即不依赖于待定加载的即时结果的指令，编译器可以重新编排代码，以使ALU保持活跃。

指令调度与预取

延迟隐藏的最基本形式在单线程或循环体内部通过指令调度实现。编译器分析中间表示（IR）中的数据依赖关系，并将加载指令尽可能早地移至执行流中。

思考一个简单的向量 (vector)加法循环：

# 简单顺序
for i in range(N):
    a = A[i]      # 加载
    b = B[i]      # 加载
    c = a + b     # 计算 (必须等待 a, b)
    C[i] = c      # 存储

如果编译器展开此循环，它可以将加载操作组合在一起。这种方法称为软件流水线或预取。通过在处理迭代 $i$ 的同时请求迭代 $i+1$ 的数据，我们便能隐藏加载的延迟。

# 流水线/预取顺序 (伪代码)
reg_a = A[0]
reg_b = B[0]

for i in range(N - 1):
    # 立即为下一次迭代发出加载请求
    next_a = load_async(A[i+1])
    next_b = load_async(B[i+1])

    # 计算当前迭代，同时加载操作处于“在途”状态
    c = reg_a + reg_b
    store(C[i], c)

    # 更新寄存器以用于下一次循环
    reg_a = next_a
    reg_b = next_b

# 最终迭代的收尾
c = reg_a + reg_b
store(C[N-1], c)

这种转换改变了循环结构。它引入了序言（在循环开始前加载首批元素）和收尾（在循环结束后处理最终计算）。在稳态循环内部，内存子系统为未来获取数据，同时计算单元处理已有的数据。

流水线可视化

为直观感受其影响，我们可以查看操作的时间线。在串行版本中，内存总线和ALU轮流工作。在流水线版本中，它们同时操作。

串行执行与软件流水线的比较。在流水线版本中，A[1] 的加载与 A[0] 的计算同时进行。

双缓冲

在GPU和TPU等深度学习 (deep learning)加速器中，延迟隐藏常通过双缓冲实现。这特别适用于在全球高带宽内存（HBM）与更快的片上内存（共享内存或暂存器）之间移动数据的情形。

双缓冲为同一数据块分配两个独立的内存区域（缓冲区）。当计算核处理来自缓冲区A的数据时，直接内存访问（DMA）引擎会将下一个数据块加载到缓冲区B中。一旦两项操作完成，角色便会互换：核在缓冲区B上计算，DMA加载到缓冲区A中。

这种方法能有效地将内存传输开销隐藏在计算之后，前提是算术密度足够高。如果计算一个数据块的时间 ($T_{math}$) 大于加载一个数据块的时间 ($T_{load}$)，内存延迟便会完全隐藏。$N$ 个数据块的总时间近似为：

$T_{总计} \approx T_{首次加载} + N \times \max(T_{计算}, T_{加载})$

如果 $T_{math} > T_{load}$，应用程序是计算受限的。如果 $T_{load} > T_{math}$，应用程序仍是内存受限的，但性能仍比串行版本好许多。

异步复制指令

现代硬件架构提供专用指令以方便此模式。例如，NVIDIA GPU引入了 cp.async（异步复制）指令。这些命令会从全局内存发起复制到共享内存，而不阻塞执行线程。

为这类目标生成代码时，ML编译器必须：

1. 分配共享内存： 为单个数据块所需的共享内存空间预留双倍容量。
2. 发出异步加载： 生成特定的内部指令以启动传输。
3. 插入屏障： 在循环中的恰当位置放置同步屏障（wait 或 bar.sync），以确保数据在计算指令尝试读取之前已到达。

这些屏障的正确放置比较困难。如果屏障放置过早，处理器将停滞，抵消了优势。如果放置过晚或被遗漏，核将读取未初始化的内存，从而导致数值错误。

向量 (vector)化与占用率

延迟隐藏也依赖于占用率，即多处理器上同时运行的活跃warp或线程数量。即使采用软件流水线，单个线程最终仍可能因等待某个值而停滞。硬件调度器通过即时切换到另一个准备执行的线程上下文 (context)来减轻此问题。

然而，激进的优化会损害占用率。循环展开和双缓冲会增加每个线程的寄存器和共享内存压力。如果一个核需要过多寄存器来存储预取值，硬件就只能支持更少的活跃线程。

优化循环嵌套时，必须权衡三个方面：

• 流水线深度： 预取多少个阶段（单缓冲、双缓冲或三缓冲）。
• 寄存器使用： 保持这些预取值活跃的存储开销。
• 线程占用率： 硬件切换线程以自然隐藏延迟的能力。

编译器常使用成本模型来确定最佳的展开因子和流水线深度。对于矩阵乘法核，双缓冲是标准基线，但三缓冲（使用3个缓冲区）有时在计算与内存比极高的架构上会带来优势。