第 4 章 内存层级与代价表

在 H100 上，每从 HBM 读 1 字节，你必须做够 295 次浮点运算才不会被带宽卡住。 这一章要做的事，就是把这句话背后的五层内存、每一层的延迟与带宽，一条条摆到台面上。

4.1 一张图看清五层内存

打开任何一份 GPU 性能调优指南，第一张图大概率是这样：

                Hopper GPU 内存层级（H100 SXM5）
                                                                    延迟        带宽
     ┌───────────────────────────────────────────┐
     │  Registers · 256 KB / SM × 132 SM = 33 MB │              随指令读出  不单独计
     │  (但每线程私有, 不可共享)                 │
     └────────────────────┬──────────────────────┘
                          │
     ┌────────────────────▼──────────────────────┐
     │  Shared Memory (SMEM) · 228 KB / SM       │              ~29 cycles  ~230 GB/s/SM
     │  软件管理, block 内共享                   │                          ~30 TB/s 全卡
     └────────────────────┬──────────────────────┘
                          │
     ┌────────────────────▼──────────────────────┐
     │  L1 Cache · 28 KB / SM (与 SMEM 共享)     │              ~41 cycles  ~230 GB/s/SM
     │  硬件管理, 自动命中                       │
     └────────────────────┬──────────────────────┘
                          │
     ┌────────────────────▼──────────────────────┐
     │  L2 Cache · 50 MB 全卡共享                │              ~263 cycles ~8 TB/s
     │  硬件管理, 自动命中, set-associative      │
     └────────────────────┬──────────────────────┘
                          │
     ┌────────────────────▼──────────────────────┐
     │  HBM3 Memory · 80 GB                      │              ~479 cycles 3.35 TB/s
     │  HBM 物理是 5 stack × 16 GB               │
     └───────────────────────────────────────────┘

数据来源：容量与 HBM 带宽取自 NVIDIA H100 Whitepaper 与 Hopper Tuning Guide（H100 SXM5）；延迟取自第三方 microbenchmark Luo et al., Benchmarking and Dissecting the Nvidia Hopper GPU Architecture（2024）Table IV，测的是 H800，不是 H100 SXM5，只能当量级用。带宽为单 SM 或全卡聚合。其中 HBM 的 3.35 TB/s 是官方规格；SMEM 的 ~30 TB/s 是按 32 bank × 4 B/周期 × 约 1.8 GHz × 132 SM 推算的理论值（Luo et al. 实测单 SM 127.9 B/clk，与 128 B/clk 吻合）；L2 的 ~8 TB/s 官方从未公布，是按 Luo et al. 在 H800 上测得的约 4.5 KB/clk 折算的估计，随访问模式与频率波动。

把这张图记住——这个专栏后面每一章都会回到它。它是 LLM 算子优化的"地图"。

几个值得反复强调的事实：

1. 从 HBM 读比从 SMEM 读慢一个数量级：按上表，延迟约 479 对 29 cycles（约 16 倍），带宽约 3.35 对 30 TB/s（约 9 倍）。这就是为什么"把数据搬到 SMEM"在 GPU 上是头等大事。
2. 寄存器操作数随指令直接读出，不单独付访存延迟；一次 HBM 访问却要等几百 cycle。这就是为什么 GEMM 一定要把 K 维度拆开，让数据在寄存器里反复用。
3. L1 在 LLM 算子里作用有限：重要的 tile 数据都被显式放到 SMEM 里了，L1 主要承接其余的零散 global 访问和 spill。所以 LLM kernel 基本都把 carveout 顶满：SMEM 228 KB / L1 28 KB。注意这不是"默认值"——H100 支持的 SMEM 容量是 0、8、16、32、64、100、132、164、196、228 KB 这几档（Hopper Tuning Guide），驱动按 kernel 需要的 SMEM 量从中挑，也可以用 cudaFuncAttributePreferredSharedMemoryCarveout 自己指定。
4. L2 是隐藏的关键层：50 MB 听起来不大，但放一层里不少中间激活已经够了。前一个 kernel 写出的数据若还留在 L2，相邻的下一个 kernel 读它就能命中 L2 而不回 HBM，得到一点"准 fused"的效果。

4.2 寄存器：零延迟，零开销，但稀缺

寄存器是 GPU 上最快的存储——访问延迟基本是 0 cycle（与算术指令融合在一起）。但寄存器有几个硬约束：

4.2.1 寄存器是线程私有的

__global__ void kernel() {
    float x = 1.0f;       // x 是寄存器, 线程 0 的 x 跟线程 1 的 x 是两个独立寄存器
    int   y = threadIdx.x;
    // 线程间想交换 x: 不能直接, 必须用 shfl 或 SMEM
    int peer_y = __shfl_sync(0xFFFFFFFF, y, 0);  // 拿到 lane 0 的 y
}

寄存器不能在线程间共享。要交换数据必须用 warp shuffle 或 SMEM。

4.2.2 寄存器数量的硬约束

Hopper 上每个 SM 总共 65536 个 32-bit 寄存器。如果每个线程要 N 个寄存器，那 SM 上能 active 的线程数最多是 65536 / N。这个约束直接决定 occupancy：

每线程寄存器数      |    SM 上最大 active 线程数    |    Occupancy
─────────────────────────────────────────────────────────────
   16              |      4096                    |    100% (上限是 2048)
   32              |      2048                    |    100%
   64              |      1024                    |    50%
   128             |       512                    |    25%
   255             |       256                    |    12.5% (255 是每线程寄存器上限)

所以"用更多寄存器存中间结果"和"让更多 warp 活跃以隐藏延迟"是直接竞争的。第 12 章手写 GEMM 时会非常具体地讨论这个权衡。

4.2.3 Register Spill：性能杀手

如果一个线程需要的寄存器超过单线程上限（255 个 32-bit），或者超过 __launch_bounds__ / -maxrregcount 给定的配额，编译器会把"装不下的寄存器"溢出到 local memory。Local memory 的地址空间落在 HBM 上（每线程私有、按 warp 交错排布），但它走 L1/L2 缓存——所以 spill 不必然等于一次 HBM 往返：命中 L1 时只是几十 cycle，可一旦工作集撑破 L1，就真的退化成几百 cycle 的片外访问。更要命的是 spill 会同时吃掉 L1 容量和访存指令带宽，把本该给 SMEM tile 的资源挤走。

判断是否发生 spill 最简单的方法是 nvcc 加 -Xptxas -v：

ptxas info    : Function properties for _Z6kernelPf
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 80 registers, used 0 barriers

（输出格式以 nvcc 13.4 为准，不同 CUDA 版本可能略有不同。）

spill stores/loads = 0 是好的。如果非零，说明发生了 spill。严肃的 LLM kernel 要把 spill 压到零，第 15 章（FA2 前向的寄存器预算）、第 20 章（从 SASS 看 spill）和第 21 章（陷阱 1）会详细讲怎么避免。

4.3 Shared Memory：黄金 228 KB

SMEM 是 LLM kernel 优化的核心舞台。它在物理上是 SM 内部的一块 SRAM，访问延迟约 29 cycles（Luo et al. 2024 在 H800 上实测）。带宽怎么算？每个 SM 有 32 个 bank、每 bank 每周期 4 字节，也就是 128 字节/周期/SM；按 ~1.8 GHz 折算约 230 GB/s 每 SM，全卡 132 个 SM 合计约 30 TB/s——差不多是 HBM3 那 3.35 TB/s 的 9 倍。关键在于它是每 SM 独享的：HBM 那 3.35 TB/s 要被 132 个 SM 分掉，SMEM 不用。

SMEM 的两个关键特性需要特别理解：32 banks 和 bank conflict。

4.3.1 SMEM 被切成 32 个 bank

SMEM 物理上不是一整块 SRAM，而是 32 个 bank——每个 bank 可以独立地处理读/写请求。一个 warp（32 线程）同时访问 SMEM 时，硬件并行地让 32 个 bank 各自服务一个线程。

                       SMEM (228 KB)
┌──────────────────────────────────────────────────────────────┐
│ bank 0  bank 1  bank 2 ... bank 30  bank 31                  │
│  4 B    4 B    4 B   ...   4 B     4 B   <- 第 1 个 32-bit 字 │
│  4 B    4 B    4 B   ...   4 B     4 B   <- 第 2 个 32-bit 字 │
│  ...                                                         │
└──────────────────────────────────────────────────────────────┘

地址 → bank 映射:  bank_id = (addr / 4) % 32

地址映射意味着：

• smem[0], smem[1], ... smem[31] 分别落在 bank 0..31。
• smem[32], smem[33], ... smem[63] 也分别落在 bank 0..31。
• smem[i] 和 smem[i + 32k] 落在同一个 bank。

4.3.2 Bank Conflict 的代价

如果一个 warp 内 N 个线程访问同一个 bank 的不同地址，硬件只能串行处理——这叫 N-way bank conflict，性能是无冲突的 1/N。

// 反例: 32-way bank conflict
__shared__ float smem[32 * 32];
int tid = threadIdx.x;
float v = smem[tid * 32];  // 所有 32 线程访问 bank 0 -> 32-way conflict!
                           // 性能是无冲突的 1/32

// 正例: 无冲突
__shared__ float smem[32 * 32];
int tid = threadIdx.x;
float v = smem[tid];       // 32 线程访问 32 个不同的 bank -> 无冲突

// 也是正例: 广播 (broadcast)
__shared__ float smem[32];
int tid = threadIdx.x;
float v = smem[0];         // 所有 32 线程都访问 bank 0 同一个地址 -> 硬件检测到 broadcast,
                           // 0 cycle 额外开销

特殊情况：所有线程访问同一个地址会触发硬件 broadcast，零额外开销。但只要有一部分线程访问同 bank 不同地址，就是 conflict。

4.3.3 矩阵转置的经典 bank conflict

最经典的 bank conflict 例子是矩阵转置：

// 原始矩阵 (32 行 × 32 列), 转置后写入 SMEM
__shared__ float smem[32][32];
int tid = threadIdx.x;
int row = blockIdx.y;
// 把 mat[row][tid] 写到 smem[tid][row] (转置)
smem[tid][row] = mat[row * 32 + tid];

// 注意 smem[tid][row]:
// tid=0, row=0: smem[0][0]    -> bank 0
// tid=1, row=0: smem[1][0]    -> 偏移 128 字节 -> bank 0 (32-way conflict!)
// tid=2, row=0: smem[2][0]    -> 偏移 256 字节 -> bank 0
// ...
// 32 个线程全部撞到 bank 0!

解决方法是 padding：

__shared__ float smem[32][33];  // 多一列 (33 而不是 32)
smem[tid][row] = mat[row * 32 + tid];

// tid=0, row=0: smem[0][0]    -> bank 0
// tid=1, row=0: smem[1][0]    -> 偏移 132 字节 -> bank 1
// tid=2, row=0: smem[2][0]    -> 偏移 264 字节 -> bank 2
// ... 完美错开!

这种 padding 叫 "+1 trick"。它浪费了一点 SMEM 容量（每 32 列加 1 列），但完全消除 bank conflict，单看这一步 SMEM 访问，理论上能差 32 倍。

4.3.4 Hopper 的 swizzle：硬件加速的反 conflict

Hopper 引入了硬件 swizzle：TMA 指令可以指定一种 swizzle 模式，硬件在写 SMEM 时自动按 swizzle 模式重排地址，无需程序员手动 padding。

TMA descriptor（CUtensorMap）里的 CUtensorMapSwizzle 有四档：

• CU_TENSOR_MAP_SWIZZLE_NONE：不 swizzle。
• CU_TENSOR_MAP_SWIZZLE_32B：在 32 字节跨度内以 16 字节块为单位错位。
• CU_TENSOR_MAP_SWIZZLE_64B：在 64 字节跨度内错位。
• CU_TENSOR_MAP_SWIZZLE_128B：在 128 字节跨度内错位（GEMM/FA 里最常用）。

（较新的 CUDA 里 128B 档又细分出以 32 B / 64 B 块为单位错位的变体：CUDA 13.4 的 cuda.h 列有 CU_TENSOR_MAP_SWIZZLE_128B_ATOM_32B、_128B_ATOM_32B_FLIP_8B、_128B_ATOM_64B 三个；CUTLASS 的映射表按 CUDA 12.7 起映射其中 _ATOM_32B / _ATOM_64B 两个：cutlass-4.7.0/include/cute/arch/copy_sm90_desc.hpp:240。别和 CUtensorMapInterleave 搞混——那是另一个字段，描述的是张量在 global memory 里的 interleave 布局，不是 SMEM 的 bank 错位。）

128B swizzle 的一个单元正好是 64 个 BF16 元素——GEMM 里 K=64 的一行 tile 恰好一整块，硬件按行错开 bank，程序员一行 padding 都不用写。

第 12 章会用 ldmatrix + swizzle 给 Tensor Core 喂 fragment，第 17 章会用 TMA + swizzle 给 WGMMA 喂 tile。

4.4 L2 Cache：被忽视的中间层

L2 cache 50 MB，对所有 SM 共享。它是连接"片上 SMEM"和"片外 HBM"的中间层。

L2 的几个关键特性：

4.4.1 L2 是 set-associative 的

L2 cache line 大小 128 字节，set-associative 设计。同一个地址只能映射到 L2 的某一组 way 上，如果这组 way 都被占用，新数据来时会驱逐旧数据。

实践影响：两个数据流在 L2 上"打架"——比如一个 stream 的 KV cache 和另一个 stream 的 weight 都映射到同一组 way，互相驱逐对方，导致 L2 命中率塌陷。

4.4.2 L2 持久化（L2 Persistence）

CUDA 11+ 提供了 L2 persistence API，让程序员显式标记某些数据"应该常驻 L2"：

// 先划出 L2 的 persisting 区 (上限是 cudaDeviceProp::persistingL2CacheMaxSize)
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, 4 * 1024 * 1024);  // 4 MB

// 再把某段热数据地址标记为 persistent
cudaStreamAttrValue attr = {};
attr.accessPolicyWindow.base_ptr = hot_table_ptr;   // 比如 RoPE cos/sin 表
attr.accessPolicyWindow.num_bytes = 4 * 1024 * 1024;  // 不能超过 accessPolicyMaxWindowSize
attr.accessPolicyWindow.hitRatio = 1.0;  // 窗口大于 set-aside 时应调低, 否则互相驱逐
attr.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
attr.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;

cudaStreamSetAttribute(stream,
    cudaStreamAttributeAccessPolicyWindow, &attr);

不先用 cudaDeviceSetLimit 划出 set-aside，persisting 标记就无处落脚。这对反复访问的小热数据（比如 RoPE 的 cos/sin 表、layer-wise scale 表）有用。但对冷数据（KV cache 这种大几个 GB）没用——它放不下 L2。

4.4.3 L2 的 streaming 写

LLM 算子里很多写操作是 write-only（输出张量），写完之后这次 kernel 不会再读。这种数据在 L2 里久留没意义——只会驱逐有用的数据。CUDA 提供 streaming store，让它在 cache 里优先被驱逐：

// PTX: st.global.cs (cs = cache streaming, 写进 cache 但标记为优先驱逐)
asm("st.global.cs.f32 [%0], %1;" :: "l"(addr), "f"(val));

或者 CUDA C++ 的 hint：

__ldca(addr);       // ld.global.ca — cache at all levels (L1+L2)，这才是**默认**行为
__ldcg(addr);       // ld.global.cg — 只在 L2 缓存，跳过 L1
__ldcs(addr);       // ld.global.cs — streaming，仍然进 L1/L2，但标记为"优先驱逐"
__stcs(addr, val);  // st.global.cs — 同理，写进去但标记为优先驱逐，尽量不挤掉热数据

注意 .cs 的语义常被误传成"绕过 L2"——PTX ISA 的原文是 evict-first 优先级，数据照样进 cache，只是排在驱逐队列最前面。真正"不经过 L1"的是 .cg。

这些 hint 在 GEMM epilogue（写最终结果）和 attention 的 mask 写入这种"用一次就扔"的场景非常有用。

4.5 HBM3：3.35 TB/s 的金矿

HBM3 是 Hopper 与外部世界的接口。80 GB 容量、3.35 TB/s 带宽——听起来很多，但被 132 SM 平摊后，每 SM 大约 25 GB/s。这就是为什么"减少 HBM 访问"是 LLM 算子的第一性原理。

4.5.1 Coalesced Access：访存的第一原则

global 访存以 transaction 为单位——计算能力 6.0 起按 32 字节 sector 计（CUDA C++ Best Practices Guide）。一个 warp（32 线程）发 32 个 4 字节读请求时：

• 如果 32 个请求是连续的 128 字节（地址 0..127）：硬件合并成 4 个 32 字节 sector（恰好覆盖 128 字节），效率 100%。
• 如果 32 个请求散落在 4096 字节范围内（每个间隔 128 字节）：硬件发出 32 次独立的 32 字节 transaction，总线上搬了 1024 字节而只有 128 字节有用——有效带宽是 1/8，事务数是 8 倍（32 对 4）。

flowchart LR
  subgraph C [Coalesced]
    direction LR
    T1[32 线程读地址 0,4,8,...124] --> M1[4 个 32B sector]
    M1 --> Eff1[效率 100%]
  end
  subgraph U [Uncoalesced]
    direction LR
    T2[32 线程读地址 0,128,256,...3968] --> M2[32 次 32B transaction]
    M2 --> Eff2[效率 12.5% 实际占用 1024B 总线 拿到 128B 有效数据]
  end

LLM 算子里 coalesced 的常见保证方式：

• 行优先存储 + 沿最低维度展开线程：让 thread 0..31 访问相邻 32 个元素。
• vectorized load（float4 / int4）：每个线程一次读 16 字节，8 个线程就读完一个 128 字节 cache line，访存指令数降到 1/4。
• TMA：硬件自动保证 coalesced，无需程序员关心。

4.5.2 Vectorized Load：一条指令顶四条

在 Hopper 引入 TMA 之前，最常用的访存优化是 vectorized load：

// 标量 load: 4 次访存
float v0 = arr[i + 0];
float v1 = arr[i + 1];
float v2 = arr[i + 2];
float v3 = arr[i + 3];

// Vectorized load: 1 次访存读 16 字节
float4 v = *reinterpret_cast<float4*>(&arr[i]);
// 或 (走只读数据路径):
// float4 v = __ldg(reinterpret_cast<const float4*>(&arr[i]));

float4 在硬件层面对应一条 LDG.E.128 指令（__ldg 版本是 LDG.E.128.CONSTANT；nvcc 13.4 以 sm_90a 编译所得，不同 CUDA 版本可能不同），单条指令读 16 字节。这有几个好处：

1. 访存指令数降到 1/4：原来 4 条指令变成 1 条。
2. 更易触发 coalesced：8 个线程的 float4 = 128 字节，正好一个 cache line。
3. 更高的内存级并行度：一条指令在 pipeline 里"占用"的资源更少。

但要注意 vectorized load 要求地址 16 字节对齐，否则会触发 misaligned address 错误。好消息是 cudaMalloc 返回的指针本身至少按 256 字节对齐；真正要小心的是自己做指针偏移（ptr + offset）时把对齐破坏掉，以及张量最后一维不是 4 的倍数时的尾部处理。

4.5.3 一个典型案例：INT4 反量化的两种写法

最能说明问题的是这样一个例子：对一个 long×long 的 INT4 量化矩阵做反量化（dequantize）。

// 朴素版本: 每线程读 1 个 INT4 (实际读 1 字节, 用其中 4 位)
int8_t packed = arr[tid];
int8_t a = packed & 0x0F;
int8_t b = (packed >> 4) & 0x0F;
out[tid * 2 + 0] = scale * (float)a;
out[tid * 2 + 1] = scale * (float)b;
// 带宽利用率明显偏低 (定性判断, 未实测)

这版本看起来没毛病——线程是 coalesced 的。但它离 HBM 峰值差得远。

问题在哪？每条指令处理的数据太少。每线程读 1 字节、做几个算术、写 8 字节，每 SM 每周期最多只能发射 4 条 warp 指令（4 个调度器各 1 条），要让在途访存量撑满 HBM，就得发出极多条这样的小访存指令。

优化版：

// Vectorized 版本: 每线程读 16 字节 = 32 个 INT4
uint4 packed = *reinterpret_cast<const uint4*>(&arr[tid * 16]);
const uint32_t w[4] = {packed.x, packed.y, packed.z, packed.w};
float4* out4 = reinterpret_cast<float4*>(&out[tid * 32]);  // out 需 16 字节对齐
// 在寄存器里展开 32 个 INT4, 每 4 个拼成一个 float4 写出
#pragma unroll
for (int j = 0; j < 4; ++j) {          // 4 个 32 位字
    #pragma unroll
    for (int k = 0; k < 2; ++k) {      // 每个字 8 个 INT4 -> 2 个 float4
        uint32_t x = w[j] >> (k * 16);
        out4[j * 2 + k] = make_float4(scale * (float)( x        & 0xF),
                                      scale * (float)((x >> 4)  & 0xF),
                                      scale * (float)((x >> 8)  & 0xF),
                                      scale * (float)((x >> 12) & 0xF));
    }
}
// 可以逼近 HBM 峰值 (定性判断, 未实测)

这个提升来自三处：

1. 指令密度：每条 LDG 处理 16 字节而不是 1 字节。
2. 更高 ILP：编译器可以把 32 次解包并行调度。
3. 写入也 vectorized：32 个 float 共 128 字节，拼成 float4 后用 8 条 STG.E.128 写完（一条 STG 一次最多 16 字节），而不是 32 条标量 STG。注意要显式拼 float4：若逐个写 out[tid * 32 + i]，编译器无法假定 out 16 字节对齐，只会生成 32 条标量 STG.E（nvcc 13.4 以 sm_90a 编译核对：上面的版本是 1 条 LDG.E.128 + 8 条 STG.E.128）。

这个案例很典型——LLM 推理的 dequant kernel 几乎都长这个样。第 9 章会详细讲。

4.6 算术强度与 Roofline 模型

理解了内存层级，就可以引入一个关键的性能分析工具：Roofline 模型。

Arithmetic Intensity（算术强度）= 算术操作数 ÷ 内存访问字节数

对于一个 kernel：

• 如果算术强度低（比如 1 FLOP/byte），它一定被带宽限制。
• 如果算术强度高（比如 1000 FLOPs/byte），它有可能被算力限制。
• 临界点就是：算力峰值 ÷ 带宽峰值。H100 SXM5 的 FP16/BF16 Tensor Core 稠密峰值是 989 TFLOPs，÷ 3.35 TB/s ≈ 295 FLOPs/byte。

flowchart LR
  X[算术强度 FLOPs/byte] --> Y[实际性能 TFLOPs]
  subgraph Roofline
    direction TB
    BW[带宽墙: y = 3.35 × x]
    COMP[算力墙: y = 989]
    BW -.-> COMP
  end

把不同 kernel 标在 roofline 上：

                  实际 TFLOPs/s
                       ▲
              989 ─────┼──────────────────────  算力天花板 (FP16 Tensor Core 稠密)
                       │ ╱          ● GEMM 4096^3 (AI~1300)
                       │╱           ● FA2 长序列 (AI~10^3, K/V 命中 L2)
                       │            ↑ 这两个都在算力墙下, 上限是 989
                  500 ─┤
                       │      ← 临界点 AI = 295 (989 / 3.35)
                       │
                       │
                       │    ● Attention naive (AI≈d/2≈64, 上限 ≈ 214)
                  100 ─┤ ╱
                       │╱
                       │    ● Attention decode (AI≈GQA 组数 g, g=8 时上限 ≈ 27)
                       │
                       │    ● LayerNorm / Element-wise (AI ≲ 1, 上限个位数)
                       └─────────────────────►  算术强度 (FLOPs/byte)
                       1     10    100   1000

                       斜线 = 带宽屋顶 y = 3.35 × AI (TFLOPs/s)

阈值 295 FLOPs/byte 来自 H100 SXM5 FP16 稠密算力 / 带宽 = 989/3.35。decode attention 的估算：每个 KV head 读 K、V 共 $2Ld$ 个 FP16 元素（$4Ld$ 字节），供 g 个 query head 各做 $4Ld$ FLOPs，AI = g——MHA 是 1，GQA 组数 4～8 时是 4～8，KV 用 FP8 再翻倍；batch 不改变它，因为每条序列的 KV 各读各的。naive attention 按 S、P 以 FP16 在 HBM 各写读一次估算，整体 AI 约 d/2（d=128 时约 64）。

Roofline 是 LLM 算子优化的"温度计"。看一个新算子，先估算它的算术强度，就能预估它是带宽 bound 还是算力 bound。

这里要提醒一句：算术强度是对"哪一层内存"算的。以 FA2 的长序列前向为例，它的浮点量是 $4N^2d$；如果每个 Q block 都要重新从 HBM 拉一遍 K/V，HBM 流量是 $4N^2d/B_r$ 字节，算术强度就等于 tile 行数 $B_r$（100 量级），落在临界点 295 左边；但 N=4096、d=128 时一个 head 的 K/V 才 2 MB，同时在算的几个 head 加起来也稳稳落在 50 MB 的 L2 里，真正打到 HBM 的只有读一遍 Q/K/V、写一遍 O，算术强度是 $4N^2d / (8Nd) = N/2$，也就是 2000 量级——远在临界点 295 的右边。这也是 FA2 能在 A100（L2 40 MB）上跑到理论峰值 ~73% 的前提之一。同一个 kernel，若 K/V 每次都回 HBM 就是带宽 bound，K/V 命中 L2 后按 HBM 算就是算力 bound，说清楚"对哪一层"才有意义。

LLM 工作负载的算术强度分布：

算子	算术强度	类型	优化方向
GEMM (M=N=K=4096, FP16)	~1300	算力 bound	Tensor Core 利用
GEMM (small batch)	≈ batch 大小 M（如 10-50）	带宽 bound	减少 HBM 访问
Attention (FA2, 长序列)	~10³	算力 bound	TMA + WGMMA 流水
Attention (decode)	≈ GQA 组数 g（1～8，FP16 KV）	带宽 bound	减少 KV 重复读
LayerNorm	~1	带宽 bound	fused kernel
RMSNorm	~1	带宽 bound	fused kernel
Softmax	~3	带宽 bound	减少 pass 数
Element-wise	0.25	严重带宽 bound	算子融合

（LayerNorm、RMSNorm、Softmax、Element-wise 几行是按每元素读写一遍、做几次运算估出的量级，随 dtype 与实现浮动。）

LLM 推理的一个隐藏事实：绝大多数算子都跑不到算力峰值，它们被带宽卡住。这是为什么在 decode 阶段，Tensor Core 989 TFLOPs 的算力大部分时间是闲置的——HBM 带宽追不上。减少 HBM 访问比"调用更厉害的 Tensor Core 指令"更重要。

4.7 一个完整案例：reduce 的访存优化

为了把这一章所有概念串起来，看一个具体例子：对 1 亿个 float 求和。

朴素版本：

__global__ void reduce_v1(const float* arr, float* out, int N) {
    __shared__ float smem[256];
    int tid = threadIdx.x;
    int gid = blockIdx.x * 256 + tid;
    smem[tid] = (gid < N) ? arr[gid] : 0.0f;
    __syncthreads();
    for (int s = 128; s > 0; s >>= 1) {
        if (tid < s) smem[tid] += smem[tid + s];
        __syncthreads();
    }
    if (tid == 0) atomicAdd(out, smem[0]);
}
// 定性: 离 HBM 峰值有明显差距 (未实测)

离 3.35 TB/s 还很远。问题在哪？

1. 每线程读 1 个 float（4 字节）。指令带宽是瓶颈。
2. 所有 block 的 atomicAdd 都打到同一个地址（原子操作在 L2 执行），彼此串行化。

优化版：

__global__ void reduce_v2(const float* arr, float* out, int N) {
    __shared__ float smem[256];
    int tid = threadIdx.x;
    int gid = blockIdx.x * 256 * 4 + tid * 4;

    // 1. Vectorized load: 每线程读 4 个 float (尾部不足 4 个时逐个读)
    float local_sum = 0.0f;
    if (gid + 4 <= N) {
        float4 v = *reinterpret_cast<const float4*>(&arr[gid]);
        local_sum = v.x + v.y + v.z + v.w;
    } else {
        for (int k = gid; k < N; ++k) local_sum += arr[k];
    }

    // 2. Warp shuffle 归约 (无 SMEM, 无 sync)
    for (int offset = 16; offset > 0; offset >>= 1) {
        local_sum += __shfl_xor_sync(0xFFFFFFFF, local_sum, offset);
    }

    // 3. Block 内汇总 warp partial sum (8 warps -> 1 block)
    if (tid % 32 == 0) smem[tid / 32] = local_sum;
    __syncthreads();
    if (tid < 8) {
        local_sum = smem[tid];
        for (int offset = 4; offset > 0; offset >>= 1) {
            local_sum += __shfl_xor_sync(0xFF, local_sum, offset);
        }
        if (tid == 0) atomicAdd(out, local_sum);
    }
}
// 定性: 可以逼近 HBM 峰值 (未实测)

改的是什么？

1. vectorized load：每线程读 16 字节，cache line 利用率拉满。
2. warp shuffle 替代 SMEM 归约：减少 SMEM 访问和 sync 次数。
3. 每个 block 处理 4 倍数据：atomic 次数降到 1/4。

这个例子的细节会在第 5 章 reduce 那一章完整展开。但要看到的是：同样的算法，不同的访存策略，能差出成倍的带宽利用率。这种差距在 LLM 算子里到处都是。

4.8 这一章的小结与下一章

这一章我们建立了 GPU 访存的精确认知：

1. 五层内存的代价表：register（随指令读出）→ SMEM (~29c) → L1 (~41c) → L2 (~263c) → HBM (~479c)（H800 实测量级）。越往外延迟越高，SMEM 到 HBM 差一个数量级。
2. Coalesced 访存是 HBM 的第一性原则：32 线程访问连续 128 字节 = 4 个 32B sector，完全散落 = 32 个 sector（8 倍）。
3. SMEM 的 32 banks 与 bank conflict：访问同 bank 不同地址 → N-way conflict → 性能 1/N。padding 或 swizzle 可避免。
4. Roofline 决定了一个 kernel 的上限：算术强度低 = 带宽 bound，高 = 算力 bound。LLM 大多算子在 100 以下，都是带宽 bound。
5. Vectorized load 与 streaming store：float4 减少访存指令、提升带宽利用率，__stcs 等 cache hint 减少一次性数据对 L2 热数据的挤占。

第 5 章我们正式开始写 kernel 代码——以 reduce 为例，把第 1-4 章的所有概念落到具体的 kernel 实现上。读者会看到一个 reduce kernel 从朴素写法一步步优化到贴近 HBM 峰值的完整过程，这是后续 GEMM/FA2 优化的"压缩演练"。

本章动手练习：

1. 写一个 kernel 故意制造 32-way bank conflict，再写一个无冲突版本，用 Nsight Compute 测 SMEM 带宽对比。
2. 跑 NVIDIA 官方 cuda-samples 里的 bandwidthTest，记录你 H100 的实际 H2D / D2H / D2D 带宽，拿 D2D 对照 3.35 TB/s 的标称值。
3. 计算一个 7B 模型 (hidden_size=4096, num_layers=32) 推理时一次 forward pass 的总 HBM 访问字节数。验证一下"为什么大模型推理是带宽 bound"。