计算机组成原理 Computer Organization & Architecture

阶段总览

1. 数据表示与浮点精度

1.1 整数、补码与溢出：为什么 AI 里到处是整数

AI 系统里真正耗资源的不只是矩阵乘法，还有大量的整数运算：token id、位置索引、attention mask、分页表、张量步长（stride）、量化缩放因子。这些都用定长整数表示，而定长整数会溢出——这不是理论问题，是真实事故来源（比如用 int32 索引一个超大张量时静默回绕）。

python# 定长整数的溢出是「静默回绕」，不是报错——这是最危险的一类 bug
import numpy as np

a = np.int32(2**31 - 1)
print(a + 1)                    # -2147483648  ← 没有异常，直接变成负数

# 真实场景：算 token 总数时溢出，导致分配了错误的 buffer 大小
n_seq, n_len = 70000, 40000
total = np.int32(n_seq) * np.int32(n_len)      # 期望 2.8e9
print(int(total), '应为', n_seq * n_len)        # 溢出成负数或错误值

# 正确做法：索引与计数一律用 int64；只有确定不会溢出的热路径才降位宽
total = np.int64(n_seq) * np.int64(n_len)

python# 把 int8 量化的缩放关系算清楚（AI 里最重要的整数运算之一）
import numpy as np

# 对称量化：把浮点区间 [-a, a] 线性映射到 int8 的 [-127, 127]
def sym_quant(x, a):                 # a = 该张量的最大绝对值
    s = a / 127.0
    q = np.clip(np.round(x / s), -127, 127).astype(np.int8)
    return q, s

w = np.array([0.0, 0.01, -0.02, 0.35, -1.0], dtype=np.float32)
q, s = sym_quant(w, a=float(np.abs(w).max()))
print('scale =', round(float(s), 6))        # ≈ 0.007874  (= 1/127)
print('int8  =', q.tolist())                # [0, 1, -3, 44, -127]
print('反量化=', (q.astype(np.float32) * s).round(4).tolist())
#   → [0.0, 0.0079, -0.0236, 0.3465, -1.0]，最大误差 ≈ 0.0035 = scale/2
# 关键结论：量化误差上限 = scale/2，而 scale 由「该张量最大绝对值」决定
#   → 逐通道 / 逐张量量化优于全局量化，因为局部最大绝对值更小、scale 更细

• 补码的价值：把减法变成加法（a − b = a + (~b) + 1），硬件只需一套加法器。这也是 int8 对称范围取 −128 ~ 127 的来由：−128 没有对应正数，反量化时要用 −127 兜底。
• 溢出的三种行为：C/汇编里是有符号溢出（未定义行为，编译器甚至可能把溢出判断优化掉）；NumPy/PyTorch 是静默回绕；Python 原生 int 是任意精度、永不上溢。跨语言迁移时这是最常见的事故源。
• 经验数字：token id 空间 1e5~2e5 用 int32 足够；当 batch × seq 逼近 2^31（约 21 亿 token）时索引必须升到 int64；而 GPU kernel 内用 int32 索引通常更快（省寄存器、省带宽），所以框架常「先检查 shape 保证不溢出，再降位宽」。

1.2 IEEE 754 与低精度数值：bf16 / fp16 / fp8 的真实差别

这是本阶段最直接影响 AI 工作的一节。浮点数的位宽分配决定了它「能表示多大的数」和「能表示多精细的数」——这两件事是互相挤占的。理解了这个取舍，你才能回答「为什么 bf16 训练稳、fp16 训练容易炸」「为什么 fp8 要分 E4M3 和 E5M2 两种」。

读这张表的正确方式：bf16 和 fp16 都是 16 位，但 bf16 把位宽给了指数（范围大、精度低），fp16 给了尾数（精度高、范围小）。fp16 的最大值是 65504，训练中只要出现稍大的激活值就会溢出成 inf，进而污染梯度——这就是「fp16 训练需要 loss scaling」的原因。bf16 因为指数位和 fp32 一样宽，范围完全够用，所以开箱即用、不需要 loss scaling，代价是精度低一些（对训练几乎无感，因为梯度更新本身是统计量）。

pythonimport numpy as np, torch

# ---- 实验 1：十进制小数的二进制表示误差（为什么 0.1+0.2 != 0.3）----
print(0.1 + 0.2 == 0.3)                     # False
print(f'{0.1:.20f}')                        # 0.10000000000000000555
print(f'{0.1+0.2:.20f}')                    # 0.30000000000000004441

# 0.1 的二进制是无限循环小数 0.0001100110011...，23 位尾数只能截断 → 必然有误差

# ---- 实验 2：累加误差（为什么大规模求和要用 Kahan / pairwise）----
xs = np.full(10_000_000, 0.1, dtype=np.float32)
print(np.float32(0.1) * 10_000_000)          # 1000000.0（先乘，误差小）
print(xs.sum())                              # 1087937.0 ← 逐项累加，误差爆炸！

# 原因：当累加和远大于单个 0.1 时，0.1 的高位被吃掉，低位直接丢失
# 解决：用 float64 累加，或分块 pairwise 求和（torch.sum 内部已做）

# ---- 实验 3：fp16 溢出与 bf16 的稳定性 ----
x = torch.tensor([1e5], dtype=torch.float16)
print(x)                                     # tensor([inf], dtype=torch.float16)  ← 溢出
y = torch.tensor([1e5], dtype=torch.bfloat16)
print(y)                                     # tensor([99840.], dtype=torch.bfloat16)  ← 只是不精确

# 这也是为什么混合精度训练的惯用做法是：
#   前向/反向用 bf16（范围安全）→ 主权重副本保持 fp32（精度安全）

• ε（机器精度）比你想的更重要：fp32 的 ε ≈ 1.19e−7。任何「两个大数相减得到小数」的运算（如 a - b 当 |a|≈|b|）都会出现灾难性抵消，有效位数直接掉光。softmax 里减最大值就是为了避免这个。
• 累加顺序会改变结果：并行归约（GPU 上）的累加顺序不确定，所以同一份代码两次运行可能有 1e-6 级差异。调试时不要以为「结果不一样 = 有 bug」。
• 量化不是简单截断：int8 量化需要 scale（和 zero-point），把浮点区间线性映射到 −128~127。对称量化（weight 常用）只有 scale；非对称量化（激活常用）还要 zero-point。
• bf16 不适合存小数值：bf16 的尾数只有 7 位，表示 (1.0, 1.01) 之间的数只能落到很稀疏的点上。所以优化器状态和主权重必须用 fp32，否则细小的更新会被直接舍入掉。
• 降低精度换到的是带宽和显存，不总是算力：省显存和带宽是确定的（位宽减半）；算力是否翻倍取决于硬件是否有对应的加速单元（Tensor Core 支持 bf16/fp8，但普通 CUDA core 不一定）。

python# FP8 训练为什么必须「逐张量缩放」：E4M3 只有 3 位尾数、最大值 448
import numpy as np

E4M3_MAX = 448.0
def to_fp8_e4m3(x):
    amax = float(np.abs(x).max()) + 1e-12
    scale = E4M3_MAX / amax              # 把最大绝对值顶到 448，吃满动态范围
    q = np.clip(x * scale, -E4M3_MAX, E4M3_MAX)
    return q, scale

x = np.array([0.001, 0.5, 8.0, -300.0], dtype=np.float32)
q, s = to_fp8_e4m3(x)
print('scale =', round(float(s), 4))         # 448/300 ≈ 1.4933
print('缩放后=', np.round(q, 2).tolist())    # [0.0, 0.75, 11.95, -448.0]
# 不缩放直接存 → 0.001 与 0.5 会一起被压成 0（E4M3 最小正规数 ~0.002）
# 代价：额外保存 scale（每张量 4 字节，可忽略），反向要能还原
#   → FP8 训练常配「延迟缩放」或逐块缩放（Blackwell 的 NVFP4）
# 相对误差量级：E4M3 有效十进制位 ≈ 1.3 → 单次量化相对误差 ~2^-4 ≈ 6%
# 所以 FP8 只用在前向 / GEMM，主权重与累加仍需 fp32 / bf16

python# 用位运算拆开一个 float32 的符号 / 指数 / 尾数（把公式变成可见的比特）
import struct

def dissect(x):
    u = struct.unpack('<I', struct.pack('<f', x))[0]
    sign = u >> 31
    exp  = (u >> 23) & 0xFF
    frac = u & 0x7FFFFF
    return sign, exp, frac, exp - 127

print(dissect(0.15625))   # (0, 124, 0x200000, -3) → 0.15625 = 1.25 × 2^-3，精确
print(dissect(0.1))       # (0, 123, 0x4CCCCD, -4) → 尾数被截断，故不精确
# 判据：一个十进制小数能被 float32 精确表示，当且仅当它是「整数 × 2^-k」形式
#   分母只有因子 2 的小数（0.5 / 0.25 / 0.15625 …）才精确；
#   0.1 / 0.2 / 0.3 的分母含因子 5，二进制是无限循环 → 必然截断

1.3 动手练习与自测

1. 把十进制 0.15625 写成 IEEE 754 单精度二进制，给出符号位、指数位（含偏移）、尾数位，并说明它为什么能被精确表示。
2. 训练任务：7B 模型、长上下文 32k、要求范围安全且开箱即用。权重与激活选 bf16 还是 fp16？主权重与优化器状态呢？
3. 对取值区间 [−0.5, 0.5] 的张量做 int8 对称量化，scale 与最大量化误差分别是多少？
4. 用「灾难性抵消」解释 softmax 为什么必须先减去最大值。
5. 把 attention mask 从 bool 改成 float32，会带来什么后果（显存与数值两方面）？

• 参考 1：0.15625 = 1.25 × 2^−3 → 符号 0；指数 127−3=124=0b01111100；尾数 0.25=0b010…0（23 位）。判据：分母只含因子 2 的小数才精确。
• 参考 2：权重与激活用 bf16（指数位同 fp32、范围安全、无需 loss scaling）；主权重副本与 Adam 状态必须是 fp32，否则细小更新会被舍入掉。
• 参考 3：a=0.5 → scale = 0.5/127 ≈ 0.003937；最大量化误差 = scale/2 ≈ 0.00197。
• 参考 4：若 logits 很大，exp(x) 会溢出；而 大数之间相减会出现抵消，有效位骤减。减去 max 后指数都 ≤ 0，既防溢出又保精度，数学上结果不变。
• 参考 5：显存上 [1,32,s,s] 的 float mask 是 bool 的 4 倍；数值上 float mask 常用大负数（如 −1e9）填充，一旦参与乘法/加法会污染梯度、更易产生 NaN，同时白白消耗带宽。

2. 指令执行与 CPU 微架构

2.1 指令集与五级流水线

CPU 提高性能的两条路：提高频率（受功耗墙限制，基本走到头了）和提高每周期执行的指令数（IPC）。流水线是把一条指令拆成多个阶段、让多条指令重叠执行——这是 IPC 的第一层来源。

理想情况下五级流水线每拍完成一条指令（IPC=1）。但只要有冒险（hazard）就会插入气泡（bubble），IPC 掉下来。三类冒险：结构冒险（硬件资源冲突）、数据冒险（后一条指令要用前一条的结果）、控制冒险（分支导致不知道下一条取哪）。

c// 数据冒险的经典例子：load-use
int a = arr[i];      // load：结果要到 MEM/WB 阶段才可用
int b = a + 1;       // 紧接着就要用 a → 必须停顿 1~2 拍（除非有转发/forwarding）

// 编译器为何要「指令调度」：把无关指令插进停顿槽
// 差：                             好（编译器重排后）：
//   load  r1, [r0]                   load  r1, [r0]
//   add   r2, r1, #1   <停顿>        load  r3, [r4]      ← 无关指令填充
//   store [r5], r2                   add   r2, r1, #1
//                                    store [r5], r2

// 这也解释了一个反直觉现象：
// 两个循环版本在数学上等价，但一个快 3 倍——差别就在「依赖链长度」和「访存局部性」

python# 用 timeit 测不同运算的「延迟量级」——比背表格更直观
import timeit
setup = 'x=1.234; y=5.678; xs=[i*0.1 for i in range(1000)]'

for stmt, name in [
    ('x*y', '浮点乘'),
    ('x/y', '浮点除'),
    ('sum(xs)', '1000 元素求和'),
    ('[v*2 for v in xs]', 'Python 逐元素循环'),
]:
    t = timeit.timeit(stmt, setup=setup, number=10000) / 10000
    print(name, round(t * 1e9, 1), 'ns/次')

# 典型量级（x86 笔记本）：浮点乘 ~30~50 ns；浮点除约为乘的 1.5~3 倍
# 注意：在解释器里，「一次除法 ≈ 几十拍」的差异被淹没了。
# 要看到流水线 / 冒险的真实代价，必须去 C / Rust / CUDA 里测。
# 结论：语言抽象层会掩盖微架构差异——这也是 GPU 调优必须用 ncu 而非 Python 计时

• 前递 / 转发（forwarding）能消掉大部分数据冒险：ALU 结果可以旁路给下一条指令，不必等写回。真正无法消除的是 load-use 冒险（要先从内存取数）——这是「指针追逐」类代码慢的根本原因。
• 关键路径决定 IPC 上限：一段代码里最长的依赖链长度除以并行宽度，就是理论最少周期数。「把依赖链打断」往往比「多用几个核」更有效。
• 在 LLM 解码里，token 之间的依赖是天然的串行链：第 t+1 个 token 依赖第 t 个的输出，无法并行。所以解码只能靠 batch（并行服务不同请求）来摊薄，而不是靠更多线程算同一个请求。

2.2 分支预测、乱序执行与超标量

现代 CPU 的实际 IPC 能到 4~6，靠的是超标量（一拍发射多条）+ 乱序执行（按数据就绪顺序执行）+ 推测执行（猜分支方向先跑）。这三者的共同代价是：一旦猜错，要回滚。

• 分支预测：现代预测器准确率能到 95%+，但分支密集且随机的代码会让它失效——每次失败要清空流水线（几十拍）。所以「有序数据上跑得快、随机数据上跑得慢」在 CPU 上是常见现象。
• 乱序执行（OoO）：CPU 用重排序缓冲区（ROB）在窗口内乱序执行、顺序提交。窗口大小有限（几百条 μop），所以依赖链太长时仍然会卡住。
• 推测执行与安全：Spectre / Meltdown 这类侧信道攻击就是利用推测执行的副作用。这是「性能与安全的取舍」的经典案例。
• 对 AI 数据管线的启示：如果你的 tokenize / 清洗代码里有大量 if 类型 == X 的分支，考虑改成查表 / 位运算 / 分支无关（branchless）写法，或者干脆交给向量化框架。

python# 一个能自己跑出结论的实验：分支预测的影响
import random, time

def sum_branchy(xs, threshold):
    s = 0
    for x in xs:
        if x >= threshold:      # 数据随机时，这个分支预测失败率约 50%
            s += x
    return s

n = 20_000_000
sorted_data = sorted(random.random() for _ in range(n))
random_data = [random.random() for _ in range(n)]

for name, data in [('有序数据', sorted_data), ('随机数据', random_data)]:
    t0 = time.perf_counter(); sum_branchy(data, 0.5); dt = time.perf_counter() - t0
    print(f'{name}: {dt:.2f}s')

# 在编译型语言（C / Rust / Java）上差异会非常明显（可达 2~4 倍）；
# 在 CPython 上被解释器开销掩盖，但换成 NumPy 向量化后差异就消失了——
# 因为向量化根本没有分支。这就是「用数据布局换分支」的最简单例子。

text// 为什么「分支无关（branchless）」写法更快 —— 用选择运算替代分支
// 有分支：数据随机时预测失败率约 50%
int clamp_branchy(int x) {
    if (x < 0) return 0;
    if (x > 255) return 255;
    return x;
}
// 无分支：用位运算，预测器无压力
int clamp_branchless(int x) {
    x = x & ~(x >> 31);                       // 负数清零：算术右移得到全 1 掩码
    return x ^ ((x ^ 255) & -((x > 255)));    // 上限裁剪
}
// 经验：分支密集且数据随机时，branchless 可快 1.5~3x；
//       但数据有序时预测器本来就准，branchless 反而可能略慢。
// 结论：先看数据分布，再决定要不要 branchless。

1. 画出一条 load-use 数据冒险的 5 级流水线时序，标出需要插入的停顿拍数，并说明前递为什么消不掉它。
2. 对 if (x > 0) 则累加、否则累减的代码：数据随机与数据有序两种情况下，预测器表现差多少？如何改写成 branchless？
3. 为什么「IPC 从 3 提到 4」和「频率从 3GHz 提到 4GHz」对吞吐是等效的？各自的代价是什么？
4. LLM 解码时 token 之间为什么不能并行？这对应流水线的什么概念？
5. 若推理服务 90% 时间在 GPU 算、10% 在 CPU 预处理，用 Amdahl 说明优化 CPU 预处理最多能省多少。

• 判据 1：load 的结果要到 MEM/WB 才可用，紧接的 add 至少停 1~2 拍；前递只能旁路 ALU/EX 结果，无法旁路「还没从内存拿回来」的数据。
• 判据 2：随机数据预测失败率约 50%、每次清空流水线几十拍；有序数据预测几乎全中。改写：用条件传送（三元运算）或 s += x * sign。
• 判据 3：吞吐 ≈ 频率 × IPC，二者相乘所以等效；提频受功耗/散热限制（电压墙），提 IPC 受并行度、依赖链与面积成本限制。
• 判据 4：自回归解码是串行依赖链（token t+1 依赖 t 的输出），无法在同一请求内并行——对应流水线的数据/控制依赖，只能靠 batch 掩盖。
• 判据 5：Amdahl 下即使 CPU 部分无限快，上限是 1/0.9 ≈ 1.11x，即最多省约 10%；所以应先优化占 90% 的 GPU 侧。

3. 存储层次与 Cache：局部性就是性能

3.1 层级结构与访问延迟：数量级差异才是重点

这节的核心不是记住具体数字，而是记住数量级。CPU 一拍约 0.3 ns，而一次主存访问是 100 ns 量级——相差 300 倍。这意味着「算 100 次加减」可能比「从主存取一个数」还快。整个存储层次的存在意义，就是让绝大多数访问落在快的层。

注意最后两行：从主存到 SSD 是一次 1000 倍的跳跃，从 SSD 到网络是又一次 1000 倍的跳跃。这就解释了一个架构原则——能缓存的绝不重算，能批量的绝不逐条。你在 Hamauls Orion 里设计的检索结果缓存、前缀缓存、语义缓存，收益量级就是「省掉一次下游 100 ms 的调用」。

bash# 亲手测一遍延迟差异（Linux）——比看表格印象深得多
# 顺序访问 vs 随机访问同一块内存
$ python - <<'PY'
import time, random
import numpy as np

N = 64 * 1024 * 1024          # 64M 个 int64 = 512MB，远超 L3
a = np.zeros(N, dtype=np.int64)

# 顺序遍历：硬件预取友好
t0 = time.perf_counter()
for i in range(0, N, 16): a[i] += 1
print('顺序步进访问:', round(time.perf_counter() - t0, 3), 's')

# 随机访问：预取完全失效，每次都可能 miss 到主存
idx = list(range(0, N, 16)); random.shuffle(idx)
t0 = time.perf_counter()
for i in idx: a[i] += 1
print('随机访问:    ', round(time.perf_counter() - t0, 3), 's')
PY
# 典型结果：随机访问慢 5~20 倍（在 Python 里被解释开销稀释，C/Rust 里差距更大）
# 但真正的指导意义在于：这就是 embedding 查表、稀疏索引、图遍历慢的根本原因

python# 用「指针追逐」测各级 Cache 的延迟台阶（经典 cache latency benchmark）
import time
import numpy as np

def chase(n, iters=3_000_000):
    # 构造随机置换的访问链，硬件预取器无法命中
    idx = np.random.permutation(n)
    a = np.zeros(n, dtype=np.int64)
    a[idx[:-1]] = idx[1:]
    a[idx[-1]] = idx[0]
    p = 0
    t0 = time.perf_counter()
    for _ in range(iters):
        p = int(a[p])            # 串行依赖：每次都要等上一次访存返回
    return (time.perf_counter() - t0) / iters * 1e9

for kb in [32, 256, 2048, 8192, 65536]:       # 工作集 (KB)
    n = kb * 1024 // 8
    print(f'{kb:6d} KB ->', round(chase(n), 1), 'ns/次')
# 典型台阶：32KB ≈ 1 ns(L1)、256KB ≈ 4 ns(L2)、8MB ≈ 20 ns(L3)、64MB ≈ 90 ns(DRAM)
# 关键：p 是串行依赖，测的是「延迟」而非「带宽」——
#      这与吞吐型负载（顺序读大数组）看到的瓶颈完全不同，别混为一谈

3.2 Cache 的组织方式：line、set、映射与写策略

• Cache line（块）通常是 64 字节：这是取数据的最小单位。所以哪怕你只要 1 个字节，也会把整条 64 字节搬进来——这既是机会（空间局部性受益）也是浪费（步长刚好等于 line 大小且只用一个元素时，利用率只有 1/64）。
• 三种映射方式：直接映射（一个主存块只能放一个位置，冲突多）、组相联（N 路，主流，如 8 路 / 16 路）、全相联（任意位置，硬件贵，只用于小容量如 TLB）。
• 替换策略：LRU 及其近似（伪 LRU、随机）。这与你第 3 阶段要手写的 LRU 缓存是同一个思想——缓存设计是「用一小块快存储假装成一大块慢存储」的通用技巧，CPU 用它、数据库用它、你的向量索引也要用它。
• 写策略：write-through（写穿，简单但流量大）vs write-back（写回，脏位标记，延迟写出）。常识：现代 CPU 的 L1/L2 基本都是 write-back + write-allocate。
• 三种 miss：compulsory（首次访问，无法避免）、capacity（容量不够，只能换更大的 Cache 或改算法分块）、conflict（映射冲突，可以通过 padding 缓解）。性能调优时要能判断你遇的是哪一种。

python# 用循环顺序把「局部性」变成可测量的数字
import numpy as np, time

N = 2048
A = np.random.rand(N, N).astype(np.float64)
B = np.random.rand(N, N).astype(np.float64)

# ⚠️ 注意：NumPy 内部是优化过的，要看到差异得用下面「分块」的思路
# 这里演示的是原理；真正的对比实验建议用 C 或 Numba 写朴素三重循环

# 朴素三重循环里，访问顺序决定了 cache 命中率：
#   for i: for j: for k:  C[i][j] += A[i][k]*B[k][j]
#   → A 行优先连续（好），B 按列访问（每次跨 N 个元素 = 跨 N/8 条 cache line，极差）
# 交换循环顺序（i-k-j）后 C 和 B 都变成连续访问，实测可快 3~8 倍

# 分块（tiling）：把大矩阵切成能塞进 L1/L2 的小块，让每个元素被复用时都还在 Cache 里
def blocked_matmul(A, B, BS=64):
    N = A.shape[0]
    C = np.zeros((N, N))
    for i0 in range(0, N, BS):
        for k0 in range(0, N, BS):
            for j0 in range(0, N, BS):
                # 小块内做矩阵乘：三个小块都在 Cache 里
                C[i0:i0+BS, j0:j0+BS] += A[i0:i0+BS, k0:k0+BS] @ B[k0:k0+BS, j0:j0+BS]
    return C

# 分块大小 BS 的经验：让 3 个 BS×BS 的块（A块+B块+C块）× 8 字节 ≤ L2 容量
# BS=64 → 3*64*64*8 = 98KB ≈ L2 大小；BS 太大反而退化

c// 分块矩阵乘：一个能直接看到 cache 收益的经典实验（gcc -O2 -march=native）
// naive: 三重循环 i-j-k，B 按列访问，cache line 利用率极低
for (i=0;i<N;i++) for (j=0;j<N;j++) { float s=0;
    for (k=0;k<N;k++) s += A[i*N+k]*B[k*N+j];      // B 每次跨 N 个元素
    C[i*N+j]=s; }

// blocked: 切成 BS×BS 的小块，三个块都能常驻 L1/L2
for (i0=0;i0<N;i0+=BS) for (k0=0;k0<N;k0+=BS) for (j0=0;j0<N;j0+=BS)
  for (i=i0;i<i0+BS;i++) for (k=k0;k<k0+BS;k++){ float a=A[i*N+k];
      for (j=j0;j<j0+BS;j++) C[i*N+j] += a*B[k*N+j]; }

// 实测（N=2048，单核）：naive ≈ 3.2s，BS=64 分块 ≈ 0.4s → 约 8x
// 经验：BS 取 32/64/128，取决于 L1(32~48KB)/L2 大小；太大 → 往返主存，退化
// 这就是「同样 2N^3 次浮点运算，耗时能差一个数量级」的来源

text// Cache 地址划分：给定参数就能算出 set 数（面试常考）
// 例：64B line、8 路组相联、总容量 32KB 的 L1D
//   set 数 = 32KB / (64B × 8 路) = 64 个 set
//   地址位划分（48 位虚拟 / 物理地址）：
//     [ offset 6 位 ] + [ set index 6 位 ] + [ tag 其余高位 ]
//   6 位 offset → 64B；6 位 index → 64 个 set；tag 用于命中比对
// 推论：步长恰为 4096B（64 set × 64B）的访问会全落在同一个 set → 冲突 miss
//   这就是「矩阵按列步进访问慢」在硬件层的解释（可通过 padding 打散）

1. 一个 64B line、8 路组相联、容量 32KB 的 L1D 共有多少 set？地址如何划分为 tag / set / offset？
2. 顺序访问 vs 随机访问一个 512MB 数组，实测差多少倍？给出归因。
3. 矩阵转置为什么慢？如何用 padding 消除冲突 miss？
4. AoS 与 SoA 各适合什么场景？解释「先 SoA 筛选、再 AoS 聚合」的理由。
5. 分块的块大小 BS 怎么定？给出经验公式并说明 BS 过大的后果。

• 判据 1：set 数 = 32×1024 / (64×8) = 64；offset 6 位、set index 6 位、其余为 tag。
• 判据 2：随机路径慢 5~20 倍（C/Rust 更明显）；顺序访问触发硬件预取（可能 4~8 条 line/次）并吃满带宽，随机访问每次都可能 miss 到 DRAM 且预取失效。
• 判据 3：按列访问步长 = 行宽，若行宽是 4096B 的整数倍就会全部落在同一 set 造成 conflict miss；把行宽 padding 到 33/65 个元素（非 2 的幂）即可打散。
• 判据 4：只扫一列 → SoA（连续流式、带宽 100%）；要取整条记录 → AoS（一次 line 拿全）。所以「先 SoA 排序/筛选，再对少量命中结果转 AoS 聚合」最优。
• 判据 5：让 A块+B块+C块（3×BS²×sizeof）之和 ≤ L2、或让 A 行块+B 行块 ≤ L1；BS 过大会超出 Cache 容量，退化成反复往返主存，反而更慢。

3.3 为 AI 优化数据布局：AoS vs SoA 与对齐

在 AI 的预处理与检索热路径里，数据结构的内存布局往往比算法复杂度更决定性能。经典对比是 AoS（Array of Structs，结构体数组）与 SoA（Struct of Arrays，数组的结构体）。

pythonimport numpy as np

# AoS：每个样本的所有字段连续存放 —— 单个样本的完整信息一次拿到（cache 友好）
#   [ (id0,len0,score0), (id1,len1,score1), ... ]
aos = np.zeros(1_000_000, dtype=[('id', 'i8'), ('len', 'i4'), ('score', 'f4')])

# SoA：每个字段各自连续 —— 只读某一个字段时，带宽利用率 100%
#   ids[...], lens[...], scores[...]
soa = {'id': np.zeros(1_000_000, 'i8'),
       'len': np.zeros(1_000_000, 'i4'),
       'score': np.zeros(1_000_000, 'f4')}

# 场景决定选择：
#   「按 score 排序取 top-k」 → SoA 快得多（只扫 scores，连续流式读取）
#   「取出第 i 条样本的所有字段」 → AoS 快得多（一次 cache line 拿全）
#   实践中常见做法：热路径用 SoA 做筛选/排序，最后对选中的少量结果再按 AoS 聚合

# 另一个高频技巧：结构体填充（padding）避免 cache line 跨行与 false sharing
# 多线程下，两个线程各自写同一 cache line 的不同变量 → 缓存行来回失效（false sharing）
# 解决：把每个线程的热变量按 64 字节对齐隔离

• 对齐（alignment）：现代 SIMD 指令要求 32/64 字节对齐才能发挥全部吞吐。NumPy / PyTorch 张量默认已对齐，但当你做 x[:, ::2] 这类非连续切片时，底层会退化成逐元素拷贝——这就是为什么「切片后变慢」很常见。
• 连续性是硬约束：很多算子（reshape、view）要求张量是内存连续的。.contiguous() 会触发一次真实拷贝，代价可能是几十毫秒。在热路径里频繁调用它是最常见的隐形性能杀手。
• layout 选择（NCHW vs NHWC）：图像任务里 NHWC 通常对卷积更友好（通道维连续，便于向量化）；而某些推理引擎在特定硬件上偏好 NCHW。这不是数学问题，是访存问题。
• 在 Hamauls Orion 里怎么落地：M2 里程碑要求你对检索候选集做一次布局改造——把「每行一个 dict」改成「列式数组」，并测出排序阶段的加速比。这个改动通常能带来数倍收益，且完全不需要换算法。

python# AoS vs SoA 的真实排序加速比（NumPy 一行搞定）
import numpy as np, time

n = 5_000_000
rng = np.random.default_rng(0)
ids = rng.integers(0, 1_000_000, n)
scores = rng.random(n).astype(np.float32)

# SoA：只扫 score 一列（连续流式），再按 index 取 id
t0 = time.perf_counter()
order = np.argsort(scores)[::-1]
top = order[:100]
dt_soa = time.perf_counter() - t0

print('SoA top-100 耗时:', round(dt_soa, 3), 's')
print('命中 id:', ids[top][:5].tolist())

# AoS 版（record array）做同样的事通常慢 2~5x：
#   因为每个候选都要把整条记录（含未用到的字段）搬进 cache line
# 经验结论：筛选/排序阶段用 SoA，聚合输出阶段再转 AoS ——
#   这是向量检索、推荐召回、日志分析里通用的布局策略

4. 内存墙与 Roofline：判断瓶颈的定量框架

4.1 算术强度与 Roofline 模型

这是本阶段最有用的一个分析工具，也是最该背下来的一个公式。它回答的问题是：这个算子到底是算力受限还是带宽受限？以及理论上还能快多少？

text算术强度 (Arithmetic Intensity, AI) = 浮点运算次数 / 访问字节数     单位: FLOP/Byte

性能上限 = min( 峰值算力 ,  算术强度 × 内存带宽 )

             log(性能 FLOP/s)
                    ▲
                    │            ╱  ← 带宽受限区（斜率 = 内存带宽）
                    │          ╱
                    │        ╱
            峰值算力 ├───────●────────────  ← compute-bound 天花板
                    │      ╱  ↑
                    │    ╱  拐点 = 峰值算力 / 内存带宽
                    └───┴──────────────────► log(算术强度)
                    带宽受限         算力受限

例子（H100 数量级）：峰值 bf16 算力 ≈ 1000 TFLOPS，HBM 带宽 ≈ 3.35 TB/s
   → 拐点算术强度 = 1000e12 / 3.35e12 ≈ 299 FLOP/Byte
   → 算术强度 < 299 的算子就是 memory-bound，再堆算力也没用

把常见算子的算术强度算一遍，你就理解了 GPU 利用率为什么提不上去：

python# 给算子画 roofline（无需 GPU 也能算上界）
def roofline(flops, bytes_moved, peak_tflops, bw_gbs):
    ai = flops / bytes_moved                       # FLOP/Byte
    mem_bound = ai * bw_gbs * 1e9 / 1e12           # 折算成 TFLOPS
    kind = 'memory' if mem_bound < peak_tflops else 'compute'
    return ai, min(peak_tflops, mem_bound), kind

# 以 H100(SXM, bf16) 为例：峰值 ≈ 989 TFLOPS，HBM3 ≈ 3.35 TB/s，拐点 ≈ 295
PEAK, BW = 989.0, 3350.0
cases = [
    ('向量加      ', 2*4e9,     3*4e9*4),          # 1e9 元素，读2写1，fp32
    ('LayerNorm   ', 5*4e9,     2*4e9*4),
    ('GEMV(4096^2)', 2*4096**2, 4096**2*4 + 4096*4),
    ('GEMM(4096^2)', 2*4096**3, 3*4096**2*4),
]
for name, f, by in cases:
    ai, cap, kind = roofline(f, by, PEAK, BW)
    print(name, 'AI=', round(ai, 2), ' 上限=', round(cap, 1), 'TFLOPS ', kind)
# 预期：向量加 / LayerNorm / GEMV 全是 memory-bound；GEMM(4096) AI≈1365 → compute-bound
# 记住判据：AI < 295（H100）就是带宽受限，堆算力无效

4.2 内存墙：为什么算力涨得比带宽快

过去几十年，CPU 算力大约每 2 年翻倍，而内存带宽与延迟的改善速率慢得多——这个差距被称为内存墙（Memory Wall）。后果是：越来越多的程序从「算得慢」变成「等数据」。AI 是这个趋势的最极端案例。

• PCIe 是一道极常见的隐形墙：HBM 有 3 TB/s，而 PCIe 只有 64 GB/s —— 差 50 倍。所以把数据从 CPU 内存搬到 GPU 显存这一步，往往就是数据管线的真实瓶颈。优化方向：提前搬运（prefetch）、在 GPU 侧做增广、用 pinned memory + 异步拷贝、合并小传输。
• Roofline 的实用读法：先算出你这个算子的算术强度，再和拐点比。远离拐点的一侧就是你的约束；如果是 memory-bound，优化方向是减少访存（融合算子、分块、量化、增大 batch 摊薄），而不是「多写几个 CUDA core 并行」。
• 算子融合（fusion）为什么有效：把「LayerNorm → 加法 → GELU」三个 memory-bound 算子融合成一个，访存从 6 遍降到 2 遍，理论上直接快 3 倍。这是 torch.compile / TensorRT / 手写 kernel 的主要收益来源。
• KV Cache 是显存的主要消费者之一：公式 2 × layers × heads × head_dim × seq_len × batch × dtype_bytes。对 70B 模型、4096 序列、bf16，单条请求的 KV Cache 就约 1.25 GB。长上下文场景下它比权重还吃显存——这正是 PagedAttention / GQA / KV 量化要解决的问题。

python# Hamauls Orion M2 的核心交付：一份能自己算出「跑不跑得动」的容量模型
def vram_budget(params_b: float, bytes_per_param: int = 2,
                 n_layers: int = 80, n_kv_heads: int = 8, head_dim: int = 128,
                 seq: int = 4096, batch: int = 1, kv_bytes: int = 2):
    """返回各部分的显存占用（GB）。参数按十亿计，返回按 GB。"""
    GB = 1024 ** 3

    weights = params_b * 1e9 * bytes_per_param / GB

    # KV Cache: 2(K和V) × 层数 × KV头数 × head_dim × 序列长 × 批大小 × 位宽
    kv = 2 * n_layers * n_kv_heads * head_dim * seq * batch * kv_bytes / GB

    return {
        'weights_GB': round(weights, 2),
        'kv_cache_GB': round(kv, 2),
        'kv_per_token_MB': round(2 * n_layers * n_kv_heads * head_dim * kv_bytes / 1024**2, 3),
        'total_GB': round(weights + kv, 2),
    }

print(vram_budget(70, n_layers=80, n_kv_heads=8, head_dim=128, seq=4096, batch=1))
# → weights 130.4GB（bf16），kv_cache 0.625GB，单条 4096 上下文合计 ≈ 131GB
# 结论：70B 的 bf16 权重连一张 80GB 卡都放不下 → 必须量化或张量并行
# 练习：把 kv_bytes 改成 1（fp8 KV）看省了多少；把 seq 改成 32768 再看 kv 涨了多少倍

python# 测你机器的真实内存带宽（达到标称的百分之几？）
import numpy as np, time

def bandwidth_gbs(n_read, n_write, dt):
    return (n_read + n_write) / dt / 1e9

N = 200_000_000                       # 200M float32 = 800MB，远超 L3
a = np.ones(N, dtype=np.float32)
b = np.ones(N, dtype=np.float32)

# 顺序 copy：读 N + 写 N
t0 = time.perf_counter(); c = a + b; dt = time.perf_counter() - t0
print('copy 带宽 :', round(bandwidth_gbs(4*N, 4*N, dt), 1), 'GB/s')

# 归约：读 N（写出极少）
t0 = time.perf_counter(); s = a.sum(); dt = time.perf_counter() - t0
print('reduce 带宽:', round(bandwidth_gbs(4*N, 0, dt), 1), 'GB/s')
# 典型结果：单通道 DDR5 约 30~45 GB/s；双通道 60~90；GPU HBM 2000~4000
# 达到标称 60%+ 就算不错；差距说明访存模式、NUMA 或没打满通道

1. 算 H100 的 roofline 拐点（bf16 989 TFLOPS / 3.35 TB/s），说明 70B 解码的算术强度为何远低于它。
2. 一个逐元素 GELU 算子的算术强度约 0.25，它受限于什么？给出三种优化手段。
3. 推导 70B、80 层、8 个 KV 头、head_dim 128、4096 序列、bf16 的单条请求 KV Cache 大小。
4. 为什么算子融合能把 3 个 memory-bound 算子的时间降到约 1/3？
5. FlashAttention 把 attention 的 HBM 访存从 O(s²) 降到什么量级？靠什么机制？

• 判据 1：拐点 = 989e12 / 3.35e12 ≈ 295 FLOP/Byte；70B 解码每步读 140GB 权重、只做 140 GFLOP → AI ≈ 1，远小于 295，故是带宽受限。
• 判据 2：GELU 是逐元素算子，AI ≈ 0.25 < 拐点 → memory-bound。手段：① 与前后算子融合；② 降精度（bf16/fp8）减少字节；③ 增大 batch 摊薄权重读取。
• 判据 3：2( K/V ) × 80 × 8 × 128 × 4096 × 2 bytes = 1,342,177,280 B ≈ 1.25 GB；与权重同量级，长上下文下更吃显存。
• 判据 4：融合把「读+写」的次数从每个算子各一遍（6 遍）降到 2 遍（1 读 1 写），访存减少 3 倍 → 时间约 1/3（带宽受限时近似线性）。
• 判据 5：降到约 O(s²·d/BC) 且中间矩阵不再落 HBM；机制是分块（tiling）+ 在线 softmax（online rescaling）+ 把 S、P 留在 shared memory / 寄存器。

5. 并行体系结构：从 ILP 到 SIMD 到多核

5.1 三种并行与 Amdahl 定律

textAmdahl 定律：加速比 = 1 / ( (1-p) + p/s )

  p = 可并行部分占比，s = 该部分的加速倍数

  例子：如果 95% 的代码可以并行、这 95% 部分做到无限快
        → 最大加速比 = 1/(0.05) = 20 倍（而不是无限）

  这是「先找瓶颈，再并行」的数学依据：
  优化一个只占 5% 时间的部分，即使做到极致也只省 5%。
  → 所以必须先 profile，再动手。

python# 强扩展 vs 弱扩展：为什么「加卡」有时几乎不涨吞吐
def amdahl(p, s):
    return 1 / ((1 - p) + p / s)

for p in [0.5, 0.9, 0.95, 0.99]:
    vals = [round(amdahl(p, s), 2) for s in (2, 8, 100, 1e9)]
    print(f'p={p:<5} 加速 s=2/8/100/∞ → {vals}')
# 关键数字：p=0.95 时即使并行部分无限快，上限也只有 20x；p=0.99 → 100x
# → 分布式训练里通信 / 同步 / 数据加载这些「串行部分」才是扩展性天花板
# 强扩展：固定问题规模、加资源 → 受 Amdahl 限制，收益递减
# 弱扩展：问题规模随资源同比例增长 → 更接近线性（Gustafson），更像 LLM 训练

5.2 SIMD 与向量化：为什么 NumPy 快

SIMD（Single Instruction Multiple Data）让一条指令同时处理 8/16/32 个数据元素。这是 NumPy、PyTorch 底层、以及一切「向量化改写」的硬件基础。理解了 SIMD，你就理解了为什么「用 for 循环写 Python」和「用 NumPy 写」能差 100 倍。

pythonimport numpy as np, time

n = 10_000_000
xs = np.random.rand(n).astype(np.float32)

# ① 纯 Python 循环：每轮有解释器开销，且完全没有 SIMD
t0 = time.perf_counter()
out = [x * 2.0 + 1.0 for x in xs]
print('Python 循环: ', round(time.perf_counter() - t0, 3), 's')

# ② NumPy 向量化：底层是 SIMD 指令 + 循环零解释开销
t0 = time.perf_counter()
out = xs * 2.0 + 1.0
print('NumPy 向量化:', round(time.perf_counter() - t0, 3), 's')

# 典型差距 30~80 倍。注意：这个差距的两个来源是
#   (a) 消除了逐元素的解释器开销（大头）
#   (b) SIMD 让一条指令处理多个元素（几倍）
# 很多教程只强调 (b)，其实 (a) 往往贡献更大。

• 向量化的前提是「无循环依赖」：每轮迭代不能依赖上一轮的结果。有依赖时必须换算法（前缀和、扫描），或用专门实现。
• 注意内存布局：向量化读连续内存效率最高。xs[::2] 这类跨步访问会让 SIMD 退化甚至回退到标量。
• 在 GPU 上这叫 coalescing：一个 warp 的 32 个线程访问连续地址时可以合并成一次事务；散乱访问则要发多次请求，带宽利用率暴跌。这是同一个「连续性」原则在 GPU 上的名字。
• 广播语义要看清：(n,1) * (1,m) 会先物化成 (n,m)，内存可能直接爆掉。大矩阵场景优先用 einsum 或显式分块，避免隐式广播的中间张量。

c// SIMD 加速比从哪来：一条指令处理 8 / 16 个 float（AVX2 / AVX-512）
#include <immintrin.h>
void add_scalar(float* a, float* b, float* c, int n) {
    for (int i=0;i<n;i++) c[i] = a[i] + b[i];          // 每拍 1 个元素
}
void add_avx2(float* a, float* b, float* c, int n) {
    for (int i=0;i+8<=n;i+=8)
        _mm256_storeu_ps(c+i, _mm256_add_ps(_mm256_loadu_ps(a+i),
                                            _mm256_loadu_ps(b+i)));  // 8 个 / 条
}
// 理论加速 8x（AVX2）或 16x（AVX-512）；实测多为 4~7x
//   因为向量加已受内存带宽限制（见 roofline），数据能常驻 L1 才接近理论值
// 结论：SIMD 提高的是「算力」，算子 bandwidth-bound 时收益被带宽封顶

1. 用 Amdahl 计算：p=0.9、并行部分加速 10x 时的总加速比。
2. 写出 AVX-512 一条指令能处理多少个 fp32 / fp64，以及对应的理论加速比。
3. 对比数据并行、张量并行、流水并行的通信量级，并说明为什么 TP 通常限制在单机 NVLink 内。
4. 弱扩展（Gustafson）为什么比强扩展（Amdahl）更接近 LLM 训练的现实？
5. 向量化改写的两个前提条件是什么？跨步访问为什么会让它退化？

• 判据 1：1.1 / (1−0.9 + 0.9/10) = 1/(0.1+0.09) = 5.26x。
• 判据 2：AVX-512 有 512 位宽 → 16 个 fp32 或 8 个 fp64；理论加速分别 16x / 8x（实测受带宽限制）。
• 判据 3：DP 每步 all-reduce 约 2×参数量（通信量大）；TP 每层两次 all-reduce（最大）；PP 只在层边界传激活（最小）。TP 通信最密，需要 NVLink 级带宽，跨机 IB 撑不住，故通常限单机。
• 判据 4：LLM 训练的 batch/序列/模型规模随卡数同步增长，问题规模不固定 → 属于弱扩展场景，趋近线性；Amdahl 的悲观上界不适用。
• 判据 5：前提是「无循环依赖（可重排）」与「访问连续（可合并）」。跨步 xs[::2] 时相邻车道取不到相邻内存，无法打包成一条向量指令，退化成标量。

6. GPU 体系结构：为什么它适合 AI

6.1 SIMT 执行模型：warp、分支发散与延迟隐藏

CPU 追求低延迟（让单条指令流尽快跑完），所以把面积花在分支预测、乱序执行、大 Cache 上。GPU 追求高吞吐（单位时间完成尽量多的运算），所以把绝大多数面积花在运算单元上，用海量线程切换来隐藏延迟。这是两种完全不同的设计哲学。

• warp 是调度的最小单位：通常 32 个线程一组，锁步（lockstep）执行同一条指令。
• 分支发散（divergence）：如果 warp 内 32 个线程走了不同的分支，硬件会把每条路径各执行一遍，用 mask 屏蔽不参与的线程。最坏情况 32 倍效率损失。所以 GPU 代码要尽量避免数据相关的分支，或者用「查表 / 选择运算」替代分支。
• 延迟隐藏的关键是 occupancy（占用率）：SM 上同时驻留的 warp 越多，越能在某些 warp 等内存时切换到其他 warp 干活。occupancy 太低（比如寄存器用太多、shared memory 占太多）就直接暴露延迟，性能腰斩。
• 但 occupancy 不是越高越好：很多高性能 kernel（如 FlashAttention）故意用高寄存器数换更少的访存，occupancy 只有 25% 却更快。目标是「足够隐藏延迟」，不是「占满」。

python# 用 PyTorch 观察 GPU 行为（无需写 CUDA）
import torch

device = 'cuda'
x = torch.randn(4096, 4096, device=device)

# ① 连续 vs 非连续：同一个数学运算，速度可以差数倍
a = torch.randn(8192, 8192, device=device)
contig = a.clone()
strided = a.t()                       # 转置视图：非连续
print('contiguous:', contig.sum().item())
print('strided:   ', strided.sum().item())
# 用 torch.profiler 看 kernel 时间，通常 strided 版本慢 2~10 倍
# → 需要时显式 .contiguous()，但要清楚它触发了一次拷贝

# ② 元素级算子融合：3 个 memory-bound 算子 vs 1 个融合算子
def unfused(x):
    h = torch.nn.functional.layer_norm(x, x.shape[-1:])
    return torch.nn.functional.gelu(h + 1.0)

fused = torch.compile(unfused, mode='max-autotune')
# 对比两者在 profiler 里的 kernel 数量与总时间
# 未融合：layernorm / add / gelu 至少 3 次全量读写
# 融合后：1 次读 + 1 次写 → 理论快 3 倍，实测通常 2~3 倍

# ③ 用 profiler 看真实的 kernel 级耗时（这是所有优化的起点）
from torch.profiler import profile, ProfilerActivity
with profile(activities=[ProfilerActivity.CUDA]) as prof:
    for _ in range(20):
        unfused(x)
print(prof.key_averages().table(sort_by='cuda_time_total', row_limit=10))

python# 量化「合并访存」的代价（用 PyTorch 索引表达）
import torch, time
dev = 'cuda'
x = torch.randn(1_000_000, device=dev)

idx_contig = torch.arange(0, 1_000_000, device=dev)     # 连续 → 合并访存
idx_random = torch.randperm(1_000_000, device=dev)      # 散乱 → 每次多发事务

def bench(fn, iters=50):
    fn(); torch.cuda.synchronize()
    t0 = time.perf_counter()
    for _ in range(iters): fn()
    torch.cuda.synchronize()
    return (time.perf_counter() - t0) / iters * 1e3     # ms

print('连续 gather:', round(bench(lambda: x[idx_contig]), 3), 'ms')
print('随机 gather:', round(bench(lambda: x[idx_random]), 3), 'ms')
# 典型：随机 gather 慢 3~10x —— 这就是 embedding 查表、稀疏注意力的成本来源
# 启示：能重排成连续访问的先重排（这正是 MoE 推理里 token 重排的作用）

6.2 GPU 存储层次与三个致命陷阱

• 陷阱一：合并访存（coalescing）。一个 warp 的 32 个线程若访问连续地址，合并为极少事务；若地址散乱，则每次都要独立事务，带宽利用率可能掉到 1/32。「让相邻线程读相邻内存」是 GPU 编程的第一原则。
• 陷阱二：bank conflict。Shared memory 分成 32 个 bank，同一 warp 内两个线程访问同一 bank 的不同地址就要串行化。矩阵转置是经典受害者——解决方式是用 padding 让行宽变成 33 而不是 32。
• 陷阱三：CPU↔GPU 同步拷贝。tensor.cpu() / .item() / print(tensor) 都会触发同步，让流水线断掉。训练循环里写 .item() 取 loss 是常见性能杀手——正确做法是累积到 GPU 上的标量，每隔 N 步同步一次。
• 显存碎片：PyTorch 的 caching allocator 缓解了大部分问题，但「逐步增大张量」的模式仍会导致碎片化 OOM。PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True 是常见的止血手段。

pythonimport torch

# 反例：训练循环里每步都同步 —— 把异步的 GPU 流水线打断
for step, batch in enumerate(loader):
    loss = model(batch)
    train_log.append(loss.item())        # ❌ 每步都同步，GPU 空转等待
    loss.backward(); opt.step(); opt.zero_grad()

# 正例：累积到 GPU 上，定期同步；用 pinned memory + 异步搬运重叠 IO
loss_acc = torch.zeros((), device='cuda')
for step, batch in enumerate(loader):
    batch = batch.to('cuda', non_blocking=True)      # 异步拷贝
    loss = model(batch)
    loss_acc += loss.detach()                        # 全程留在 GPU
    loss.backward(); opt.step(); opt.zero_grad()
    if step % 50 == 0:
        train_log.append((loss_acc / 50).item())      # ✅ 每 50 步同步一次
        loss_acc.zero_()

# 数据加载用 pinned memory + 多 worker，让 H2D 拷贝与计算重叠
loader = torch.utils.data.DataLoader(
    dataset, batch_size=64, num_workers=8, pin_memory=True, persistent_workers=True)

python# pinned vs pageable 内存：一个常被忽略却常占大头的差异
import torch, time

n = 256 * 1024 * 1024 // 4                    # 256MB 张量
cpu_plain = torch.empty(n)                    # 可分页（默认）
cpu_pin   = torch.empty(n, pin_memory=True)   # 锁页

def h2d(t):
    torch.cuda.synchronize(); t0 = time.perf_counter()
    t.to('cuda'); torch.cuda.synchronize()
    return time.perf_counter() - t0

print('pageable H2D:', round(h2d(cpu_plain)*1e3, 2), 'ms')
print('pinned   H2D:', round(h2d(cpu_pin)*1e3, 2), 'ms')
# 典型：pinned 比 pageable 快 2~3x，且 pinned 才能真正异步（non_blocking）
# PCIe 5.0 x16 单向上限 ~64 GB/s → 256MB 的理论下界约 4 ms
# 若 H2D 时间是理论值的 3~5 倍，多半没开 pinned 或没做重叠

6.3 Tensor Core、混合精度与硬件加速单元

Tensor Core 是专门为矩阵乘累加（MMA）设计的硬件单元。它能在若干周期内完成一个 16×16×16 的矩阵乘累加——这是「AI 算力爆炸」的直接来源。理解它的约束条件，才能解释为什么某些 shape 跑得飞快、某些却慢得离谱。

• 它只做特定形状的 MMA：Tensor Core 对矩阵维度有对齐要求（常见是 8 或 16 的倍数）。这就是为什么「把 hidden size 调到 16 的倍数」能让训练快不少——不是玄学，是硬件约束。
• 它支持的低精度比 CUDA core 多：bf16 / fp16 / tf32 / fp8 / int8 都有对应通路。所以混合精度不只是省显存，更是换到更快的计算通路。
• tf32 是「零成本提速」的典型：保持 fp32 的指数范围，只把尾数砍到 10 位。在 Ampere 及以后默认开启（PyTorch 里 torch.backends.cuda.matmul.allow_tf32 = True），矩阵乘能快 2~3 倍，而多数任务精度损失可忽略。
• 算力数字要看清位宽标注：厂商标称的「1000 TFLOPS」通常是 bf16 / fp8 稀疏下的数字。实际 fp32 密集算力可能只有几十分之一。看算力参数一定要对齐「位宽 + 是否稀疏」，否则会严重高估。

python# 不同精度下的 GEMM 吞吐：Tensor Core 的低精度通路有多快
import torch, time

def gemm_tflops(dtype, n=8192, iters=50):
    a = torch.randn(n, n, device='cuda', dtype=dtype)
    b = torch.randn(n, n, device='cuda', dtype=dtype)
    for _ in range(3): a @ b                    # 预热（含 cudnn autotune）
    torch.cuda.synchronize(); t0 = time.perf_counter()
    for _ in range(iters): a @ b
    torch.cuda.synchronize()
    dt = (time.perf_counter() - t0) / iters
    return 2 * n**3 / dt / 1e12

torch.backends.cuda.matmul.allow_tf32 = False
print('fp32 :', round(gemm_tflops(torch.float32), 1), 'TFLOPS')   # ~30~60（H100）
torch.backends.cuda.matmul.allow_tf32 = True
print('tf32 :', round(gemm_tflops(torch.float32), 1), 'TFLOPS')   # 开启后 ~400+
print('bf16 :', round(gemm_tflops(torch.bfloat16), 1), 'TFLOPS')  # ~600~900
# 结论：同一张卡，bf16 密集算力约是 fp32 的 10~16 倍；
#   但这是 GEMM（compute-bound）；换成一个 memory-bound 的逐元素算子，
#   bf16 只省带宽，算力优势几乎体现不出来

1. 一个 warp 的 32 线程走 2 条分支路径，最坏效率是多少？如何避免？
2. 为什么降低 register 用量会提高 occupancy？但高 occupancy 为什么不一定更快？
3. coalescing 是什么？一个 warp 访问步长为 32 个 float 的地址，带宽利用率约多少？
4. Tensor Core 对矩阵维度有何要求？由此解释「hidden 取 16 的倍数更快」。
5. 为什么训练循环里调用 .item() 会拖慢？给出正确写法。

• 判据 1：两条路径串行执行，每条路径上另一半线程被 mask 掉 → 最坏效率 ≤ 50%（实际上是 1/路径数）。避免：用 select/where 替代分支、让同 warp 处理同质数据、把分支提到 warp 以上粒度。
• 判据 2：寄存器是 occupancy 的硬约束之一，用量减少 → 每 SM 能驻留更多 warp → 更能隐藏延迟。但 FlashAttention 类 kernel 故意高寄存器换少访存，occupancy 30% 也更快；目标是「够隐藏延迟」。
• 判据 3：coalescing = warp 内相邻线程访问相邻地址，合并为极少内存事务。步长 32 float = 128 字节，每个线程落进不同 cache line，32 条 line 只用 4 字节 → 利用率约 4/128 ≈ 3%。
• 判据 4：Tensor Core 的 MMA 要求维度是 8 或 16 的倍数；hidden 不是 16 的倍数时需要 padding 到对齐或退化，导致有效算力下降。同类还有 head_dim、序列分块、卷积通道的对齐。
• 判据 5：.item() 触发 device→host 同步拷贝，会阻塞异步流水、让 GPU 空转等 CPU。正确写法：把 loss 累积到 GPU 上的标量（loss_acc += loss.detach()），每 N 步再同步一次。

7. 性能剖析与容量规划（里程碑 M2）

7.1 度量方法：不要猜，要测

性能工作最大的浪费是「凭直觉优化」。正确的顺序永远是：测量 → 定位瓶颈 → 建立模型 → 预测收益 → 实施 → 再测量。跳过前两步的优化，通常是在优化不重要的地方。

bash# 一套可直接抄的性能排查流程
# 1) 先看宏观：GPU 到底有没有在干活？利用率低说明瓶颈在数据侧或同步点
nvidia-smi dmon -s pucm -d 1

# 2) 再看整体耗时构成
python -m torch.utils.bottleneck train.py          # 同时给出 CPU 与 CUDA 视角

# 3) 定位到 kernel 级：谁最耗时、有多少个小 kernel 可以融合
python -c "
import torch; from torch.profiler import profile, ProfilerActivity
with profile(activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
             record_shapes=True, profile_memory=True) as prof:
    run_one_epoch()
print(prof.key_averages().table(sort_by='cuda_time_total', row_limit=20))
prof.export_chrome_trace('trace.json')   # 用 chrome://tracing 打开看时间线
"

# 4) 用带宽利用率判断是不是 memory-bound（关键指标）
#    achieved_bandwidth / peak_bandwidth > 80% 且算力利用率低 → 就是带宽瓶颈
#    此时优化方向：减少访存（融合/分块/量化），而不是加并行度

# 5) 记录「改前 vs 改后」：没有对比数字的优化等于没做
#    每次只改一个变量，否则归因不了

python# 一个反例：不预热 + 单次测量会骗你
import torch, time, statistics

def bad_bench(fn):
    t0 = time.perf_counter(); fn(); return (time.perf_counter()-t0)*1e3

def good_bench(fn, warmup=10, iters=100):
    for _ in range(warmup): fn()        # 预热：CUDA module 加载 / autotune / Triton 编译
    torch.cuda.synchronize()
    ts = []
    for _ in range(iters):
        torch.cuda.synchronize(); t0 = time.perf_counter()
        fn(); torch.cuda.synchronize()
        ts.append((time.perf_counter()-t0)*1e3)
    return statistics.median(ts), max(ts)

x = torch.randn(4096, 4096, device='cuda')
f = lambda: x @ x
print('单次(无预热):', round(bad_bench(f), 2), 'ms')   # 常被首次 launch 拖成 10x
med, worst = good_bench(f)
print('中位数:', round(med, 2), 'ms  P100:', round(worst, 2), 'ms')
# 规矩：预热 + 多次 + 中位数/P95 + 固定变量，否则量到的多半是噪声

7.2 建立容量模型：从「试试看」到「我算得出来」

这一节是本阶段的产出目标。有了前六节的知识，你现在可以先算再跑：在写第一行代码之前，就估算出显存够不够、吞吐大概多少、瓶颈会在哪。这个能力在面试与工程评审中极其值钱——它把讨论从「我觉得」变成「按公式应该是」。

python# Hamauls Orion 容量模型骨架：把「能不能跑」变成可计算的函数
def capacity_model(
    params_b: float,           # 参数量（十亿）
    layers: int, hidden: int, n_kv_heads: int, head_dim: int,
    seq: int, batch: int,
    weight_dtype_bytes: int = 2,     # bf16
    kv_dtype_bytes: int = 2,
    training: bool = False,
    gpu_mem_gb: float = 80, hbm_bw_gbs: float = 3350,
):
    GB = 1024 ** 3
    m = {'weights_GB': params_b * 1e9 * weight_dtype_bytes / GB}
    m['kv_cache_GB'] = (2 * layers * n_kv_heads * head_dim * seq * batch
                        * kv_dtype_bytes) / GB

    if training:
        m['grads_GB'] = m['weights_GB']                          # 梯度 ≈ 权重
        m['optimizer_GB'] = params_b * 1e9 * 8 / GB              # Adam: fp32 m + v
        m['activations_GB'] = (batch * seq * hidden * layers * 2) / GB  # 粗略，与 batch×seq 成正比
    else:
        # 解码吞吐：每 token 至少要读一遍权重 → 理论 token/s 上限
        bytes_per_decode_step = m['weights_GB'] * GB + m['kv_cache_GB'] * GB
        m['decode_tokens_per_s_ceiling'] = round(hbm_bw_gbs * 1e9 / bytes_per_decode_step, 1)
        # 预填充阶段更接近 compute-bound，规模约为 2×参数量×token 数
        m['prefill_tflops_for_seq'] = round(2 * params_b * seq / 1000, 2)

    m['total_GB'] = round(sum(v for k, v in m.items()
                              if k.endswith('_GB') and 'per' not in k), 2)
    m['fits_in_one_gpu'] = m['total_GB'] <= gpu_mem_gb
    return m

# 跑几个典型配置，建立数量级直觉
for name, kw in [
    ('7B 推理 4k×1',  dict(params_b=7,  layers=32, hidden=4096, n_kv_heads=8,  head_dim=128, seq=4096,  batch=1)),
    ('7B 推理 4k×32', dict(params_b=7,  layers=32, hidden=4096, n_kv_heads=8,  head_dim=128, seq=4096,  batch=32)),
    ('70B 推理 8k×1', dict(params_b=70, layers=80, hidden=8192, n_kv_heads=8,  head_dim=128, seq=8192,  batch=1)),
    ('7B 训练 2k×8',  dict(params_b=7,  layers=32, hidden=4096, n_kv_heads=8,  head_dim=128, seq=2048,  batch=8, training=True)),
]:
    r = capacity_model(**kw)
    print(f"{name:16s} 总显存 {r['total_GB']:7.2f} GB  单卡可容: {r['fits_in_one_gpu']}"
          + (f"  解码上限 {r['decode_tokens_per_s_ceiling']} tok/s" if 'decode_tokens_per_s_ceiling' in r else ''))

• 先用模型估算，再用实测校准：模型给你方向和量级，实测给你修正系数。两者差距很大时，说明你漏了某个因素（通常是激活值、碎片、或框架开销）。
• 关键结论要能一句话说清：比如「70B bf16 单卡 80G 放不下，需要 2 卡张量并行或 INT4 量化」「batch 从 1 加到 32 时 KV Cache 涨 32 倍，是长上下文场景的显存主因」。
• 把模型写进代码而不是文档：Hamauls Orion 里 capacity_model() 是一个真实函数，带单测，参数变了会自动重算。文档会过期，代码不会。

python# 用实测反推「实际可用带宽」，再预测别的算子 —— 从公式到工程
def predict_decode_tok_s(params_b, weight_bytes, kv_gb, bw_gbs, eff=0.75):
    # eff：经验效率，多数 memory-bound 算子只用到标称带宽的 70~85%
    w_gb = params_b * 1e9 * weight_bytes / 1e9     # 权重体积 (GB)
    per_step = (w_gb + kv_gb) * 1e9                # 每步需读的字节
    return bw_gbs * eff * 1e9 / per_step

# 7B 模型在 H100（3.35 TB/s）上的解码上限：
print(round(predict_decode_tok_s(7, 2, 0.5, 3350), 0), 'tok/s')     # ~110 量级
print(round(predict_decode_tok_s(7, 1, 0.5, 3350), 0), 'tok/s')     # FP8 → 约 2x
print(round(predict_decode_tok_s(7, 0.5, 0.5, 3350), 0), 'tok/s')   # INT4 → 约 3x
# 结论：预测量级很准、绝对值有偏差 —— 用实测校准系数，是工程容量的正确用法

1. 基准测试三纪律是什么？各为什么重要？
2. 用 roofline 判断你的算子是否 memory-bound，需要哪两个实测指标？
3. 写出 7B 模型全参微调的显存估算（按 16 字节/参数），并说明 LoRA 为什么能救它。
4. 一次优化改动后如何证明收益不是噪声？至少列出三条要求。
5. 先算再跑：预测 batch=32、seq=8k 时 70B（bf16）能否放进 192GB 的 B200。

• 判据 1：预热（规避 JIT/autotune 干扰）、多次取分位数（压噪声）、固定变量（可归因）；缺任一条结论都不可信。
• 判据 2：需要 achieved bandwidth ÷ peak bandwidth 与算力利用率两个数；带宽利用率高而算力利用率低 → memory-bound。
• 判据 3：7B × 16~20 B ≈ 112~140 GB（权重 2 + 梯度 2 + Adam 8 + 主权重 4 + 激活/缓冲）；LoRA 只训练低秩增量、冻结主干，优化器状态与梯度从「全参」降到「增量」量级，可把显存压到几十 GB。
• 判据 4：多次迭代取中位数、报告 P95、同一硬件与数据、单变量改动、给出改前改后数字；最好做显著性（多次运行的方差）。
• 判据 5：权重 70×2 = 140 GB；KV = 2×80×8×128×8192×32×2 B ≈ 80 GB；合计 ≈ 220 GB > 192 GB → 放不下，需 FP8/INT4 量化或张量并行。

项目里程碑

阶段练习项目

常见误区

• 把「位宽减半」等同于「速度翻倍」：省显存和带宽是确定的，算力是否翻倍取决于硬件是否有对应通路。
• 在训练循环里写 `.item()` 或 `print(tensor)`，无意中让 GPU 每步都同步等待。
• 只优化占比 5% 的代码路径（Amdahl 定律的直接体现），却忽略了真正的热点。
• 看到「慢」就加并行度，但代码其实是 memory-bound —— 加并行只会让带宽更拥挤。
• 用单次跑分下结论，不做预热、不取分位数、不固定其他变量。
• 把 bf16 当成「便宜的 fp32」用来存优化器状态和主权重，导致细小更新被舍入掉、训练停滞。
• 忽略非连续张量与隐式广播带来的拷贝，让访存量凭空翻几倍。
• 只看厂商标称的峰值算力（往往是 fp8 稀疏），严重高估机器真实能力。

面试高频问题速答

为什么大模型推理的解码阶段通常不是算力瓶颈？

因为解码是 memory-bound：每生成一个 token，都要把全部权重（以及全部 KV Cache）从 HBM 读一遍，但只做 2×参数量 次浮点运算，算术强度只有约 1 FLOP/Byte，远低于 roofline 拐点（H100 上约 299）。所以解码速度 ≈ 显存带宽 ÷ 每步需读取的字节数。推论：① 权重量化（INT8/FP8/INT4）能近似线性提速；② 增大 batch 能摊薄权重读取成本、提高吞吐（但会拉高延迟）；③ 投机解码用少量额外算力换更少的权重读取轮次。

bf16 和 fp16 都是 16 位，为什么训练一般选 bf16？

位宽分配不同：bf16 用 8 位指数 + 7 位尾数，动态范围和 fp32 一样（±3.4e38）；fp16 用 5 位指数 + 10 位尾数，最大值只有 65504，训练中稍大的激活或梯度就会溢出成 inf，因此必须配 loss scaling。bf16 范围安全、开箱即用，代价是尾数精度低（约 2.4 位十进制），但梯度更新本身是统计量，对这点精度损失不敏感。注意无论用哪种，主权重与优化器状态都应保持 fp32。

什么是 roofline 模型？怎么用它指导优化？

roofline 用「算术强度 = FLOP / 访存字节」为横轴、「性能 = FLOP/s」为纵轴，画出一条由两条线构成的上界：低强度区受内存带宽限制（斜率为带宽），高强度区受峰值算力限制（水平线），拐点 = 峰值算力 ÷ 带宽。用法：先算出目标算子的算术强度，判断它落在哪一侧。若在带宽侧，优化方向是减少访存（算子融合、分块、量化、增大 batch）；若在算力侧，才考虑更好的矩阵形状对齐、低精度加速、以及分布式切分。

Cache 为什么能起作用？什么情况下会失效？

靠局部性：时间局部性（刚访问的数据很可能再被访问）与空间局部性（相邻地址很可能被一起访问，配合 64 字节 cache line 预取）。失效情形有三类 miss：compulsory（首次访问不可避免）、capacity（工作集大于缓存容量）、conflict（映射冲突）。对应手段分别是：无解 / 分块（tiling）缩小工作集 / padding 打散冲突。判断代码是否 cache 友好，看它的访存步长与工作集大小是否匹配缓存容量。

GPU 的 warp 分支发散是什么？如何避免？

一个 warp（通常 32 线程）锁步执行同一条指令。若 warp 内线程走了不同分支，硬件会把每条路径各执行一遍并用 mask 屏蔽不参与的线程，最坏情况效率降到 1/32。避免方式：① 用「选择运算」替代分支（如 `where` / 位运算）；② 让同一 warp 内线程处理同质数据（数据重排 / 分桶）；③ 把分支提到 warp 粒度之上（不同 warp 走不同路径没关系）。这也是为什么排序后再做分支判断会更快——数据变得同质了。

训练一个 7B 模型大概需要多少显存？怎么算？

粗略经验是参数量 × 16~20 字节。拆开算（bf16 混合精度）：权重 2 B/参数、梯度 2 B/参数（bf16）或 4 B（fp32）、Adam 的 m 和 v 各 4 B（fp32）、再加上主权重副本 4 B —— 仅优化器与权重相关部分就约 16 B/参数，即 7B × 16 B ≈ 112 GB；再叠加激活值（与 batch×seq×hidden×layers 成正比，可用梯度检查点大幅压缩）与通信缓冲。所以 7B 全参微调在单张 80G 卡上跑不动，需要张量并行/ZeRO 或改用 LoRA / QLoRA。

为什么把 hidden size 设成 16 的倍数会更快？

因为 Tensor Core 的 MMA 指令对矩阵维度有对齐要求（常见 8 或 16）。维度不是 16 的倍数时，要么无法使用 Tensor Core 通路、要么需要 padding 到对齐后计算，导致有效算力下降。同理还有：head_dim 对齐、attention 的序列长度分块对齐、以及卷积的通道对齐。这不是玄学参数调优，而是硬件约束的直接后果。

学习资源