第二十三课：LayerNorm 和 RMSNorm 如何在 GPU 上实现

导语：这一课讲 LayerNorm 与 RMSNorm；随后我会把课程收束到少数几个最必要章节，避免继续无限展开。

LayerNorm 是 Transformer 中最常见的算子之一。

对于一行长度为 $N$ 的输入：

$$x_0,x_1,\ldots,x_{N-1}$$

LayerNorm 的计算过程为：

$$\mu=\frac{1}{N}\sum_{i=0}^{N-1}x_i$$ $$\sigma^2= \frac{1}{N} \sum_{i=0}^{N-1}(x_i-\mu)^2$$ $$\hat{x}_i= \frac{x_i-\mu} {\sqrt{\sigma^2+\epsilon}}$$

最后进行仿射变换：

$$y_i=\gamma_i\hat{x}_i+\beta_i$$

其中：

• $\mu$：均值；
• $\sigma^2$：方差；
• $\epsilon$：防止除零；
• $\gamma_i$：可学习缩放参数；
• $\beta_i$：可学习偏置参数。

---

一、LayerNorm 在 Transformer 中对哪个维度归一化

假设输入张量为：

$$X\in\mathbb{R}^{B\times S\times H}$$

其中：

• $B$：Batch Size；
• $S$：Sequence Length；
• $H$：Hidden Size。

LayerNorm 通常沿最后一个维度 $H$ 归一化。

也就是说：

每个 Token 的 H 个特征
单独计算一组均值和方差

例如：

Token 0：[x₀, x₁, ..., xH-1] → 一次 LayerNorm
Token 1：[x₀, x₁, ..., xH-1] → 一次 LayerNorm
Token 2：[x₀, x₁, ..., xH-1] → 一次 LayerNorm

不同 Token 之间通常相互独立。

因此 GPU 上的基本映射通常是：

一个 Warp 或一个 Block
处理一个 Token 的一行特征

---

二、LayerNorm 为什么包含 Reduction

计算均值需要：

$$\sum_i x_i$$

计算方差需要：

$$\sum_i(x_i-\mu)^2$$

所以每一行至少包含两个全行统计量：

Sum Reduction
Square Sum 或 Variance Reduction

最终还要把统计结果广播给所有线程：

所有线程获得 μ
所有线程获得 1 / sqrt(σ² + ε)

然后才能生成各自负责的输出元素。

---

三、最直接的三遍实现

一种朴素实现可以分成三次遍历。

第一遍：计算均值

$$\mu=\frac{1}{N}\sum_i x_i$$

第二遍：计算方差

$$\sigma^2= \frac{1}{N} \sum_i(x_i-\mu)^2$$

第三遍：归一化并输出

$$y_i= \gamma_i \frac{x_i-\mu}{\sqrt{\sigma^2+\epsilon}} +\beta_i$$

数据流为：

第一次读取 x
    ↓
计算均值

第二次读取 x
    ↓
计算方差

第三次读取 x
    ↓
读取 γ、β
    ↓
写出 y

如果输入来自 Global Memory，则一行输入至少被读取三遍。

---

四、两遍 LayerNorm

可以同时计算：

$$S_1=\sum_i x_i$$

以及：

$$S_2=\sum_i x_i^2$$

然后：

$$\mu=\frac{S_1}{N}$$ $$\sigma^2= \frac{S_2}{N}-\mu^2$$

这样第一遍读取输入时，同时得到均值和平方和。

第二遍再归一化并写出。

数据流变为：

第一遍：
读取 x
同时计算 Σx 和 Σx²

第二遍：
读取 x
读取 γ、β
写出 y

输入只需要读取两遍。

---

五、平方和公式的数值问题

公式：

$$\sigma^2= E[x^2]-E[x]^2$$

数学上完全正确。

但在浮点数中可能发生严重的消减误差。

例如所有输入都接近：

$$10000$$

但彼此只相差很小。

此时：

$$E[x^2]$$

和：

$$E[x]^2$$

都是非常大的数。

两者相减后才得到很小的方差：

一个很大的数
减去另一个非常接近的大数
        ↓
有效精度大量丢失

结果甚至可能因为舍入误差得到一个微小负数。

因此高质量实现通常会考虑更稳定的方差算法。

---

六、Welford 在线方差算法

Welford 算法可以逐个处理元素，并维护：

• 当前元素数量 $n$；
• 当前均值 $\mu$；
• 当前平方偏差和 $M_2$。

初始状态：

$$n=0,\quad \mu=0,\quad M_2=0$$

加入新元素 $x$：

$$n'=n+1$$ $$\delta=x-\mu$$ $$\mu'=\mu+\frac{\delta}{n'}$$ $$\delta_2=x-\mu'$$ $$M_2'=M_2+\delta\delta_2$$

最终方差：

$$\sigma^2=\frac{M_2}{N}$$

LayerNorm 一般使用总体方差，因此除以 $N$，而不是统计学样本方差中的 $N-1$。

---

七、Welford 为什么更稳定

它不是先计算两个巨大且相近的量：

$$E[x^2]$$

和：

$$E[x]^2$$

再相减。

而是逐步累计：

每个值相对当前均值的偏差

所以在输入数值较大、方差较小时，通常具有更好的数值稳定性。

代价是：

• 每个元素需要更多运算；
• 更新存在一定依赖链；
• 并行归约比普通求和更复杂。

---

八、Welford 状态如何合并

GPU 中不同线程分别处理不同元素，因此每个线程会得到自己的局部 Welford 状态：

$$(n_a,\mu_a,M_{2a})$$

另一组状态为：

$$(n_b,\mu_b,M_{2b})$$

令：

$$\delta=\mu_b-\mu_a$$

合并后的数量：

$$n=n_a+n_b$$

合并后的均值：

$$\mu= \mu_a+ \delta\frac{n_b}{n}$$

合并后的平方偏差和：

$$M_2= M_{2a}+M_{2b} + \delta^2\frac{n_an_b}{n}$$

因此，Welford 状态也可以进行树形归约。

---

九、Warp 内 Welford 归约

假设一个 Warp 处理一行。

每个线程先处理一个或多个元素，得到：

Lane 0：状态 A₀
Lane 1：状态 A₁
...
Lane 31：状态 A₃₁

然后通过 Warp Shuffle 交换：

• Count；
• Mean；
• $M_2$。

每一步将两个 Welford 状态合并：

Offset 16
Offset 8
Offset 4
Offset 2
Offset 1

最终 Lane 0 得到整行状态。

之后再广播：

$$\mu$$

和：

$$\frac{1}{\sqrt{\sigma^2+\epsilon}}$$

给整个 Warp。

---

十、短 Hidden Size 的 Warp LayerNorm

如果 Hidden Size 较小，例如：

$$H\leq 1024$$

可以考虑让一个 Warp 处理一行。

例如 $H=128$：

Lane 0：处理 0、32、64、96
Lane 1：处理 1、33、65、97
...

每个 Lane 只需保存4个输入。

如果寄存器足够，可以在第一次读取时保存全部负责的数据：

Global Memory 读取一次
        ↓
寄存器保存输入
        ↓
计算均值和方差
        ↓
直接归一化并写出

这样不需要第二次读取输入。

---

十一、寄存器缓存输入的代价

假设每个线程负责16个 FP32 元素：

$$16\times32\text{ bit}$$

至少需要16个数据寄存器，此外还需要：

• 局部统计寄存器；
• 地址寄存器；
• $\gamma,\beta$；
• 中间计算变量。

因此寄存器使用可能迅速增加。

结果可能是：

输入只读取一次
但每线程寄存器大幅增加
        ↓
Occupancy 下降

这不一定是坏事。

如果减少一次完整 Global Memory 读取的收益更大，较低 Occupancy 仍可能更快。

---

十二、Block LayerNorm

当 Hidden Size 较大时，一个 Warp 可能不足。

可以使用一个 Block 处理一行：

Block：256 Threads
Hidden Size：4096
每线程处理约16个元素

归约过程：

每个线程计算局部统计
        ↓
每个 Warp 内归约
        ↓
Lane 0 写入 Shared Memory
        ↓
__syncthreads()
        ↓
第一个 Warp 归约所有 Warp 状态
        ↓
广播最终均值和方差

这与前一课的 Block Reduction 类似，只是归约对象从一个标量变成了 Welford 三元组。

---

十三、两遍 Block LayerNorm 的典型流程

第一遍：

各线程读取 x
        ↓
计算局部 Sum/Square Sum 或 Welford 状态
        ↓
Block Reduction
        ↓
得到均值和方差

第二遍：

各线程重新读取负责的 x
        ↓
读取 γ、β
        ↓
归一化
        ↓
写出 y

如果输入不能全部保存在寄存器或 Shared Memory 中，这是常见方案。

---

十四、使用 Shared Memory 保存整行是否合适

理论上可以：

读取输入到 Shared Memory
        ↓
计算统计量
        ↓
直接从 Shared Memory 读取并归一化

这样可以避免第二次访问 Global Memory。

但假设：

$$H=8192$$

FP32 一行需要：

$$8192\times4=32\text{ KB}$$

一个 Block 就需要32 KB Shared Memory。

如果还需要其他缓冲区，资源使用会进一步增加。

这会限制：

• 每个 SM 的驻留 Block 数；
• Occupancy；
• 可支持的最大 Hidden Size。

因此实际实现需要根据 Hidden Size 在以下方案间选择：

寄存器缓存
Shared Memory 缓存
重新读取 Global Memory

---

十五、LayerNorm 为什么经常是内存受限

对每个元素，通常需要：

• 读取 $x_i$；
• 读取 $\gamma_i$；
• 读取 $\beta_i$；
• 写出 $y_i$。

如果输入需要两遍读取，还要额外读取一次 $x_i$。

以 FP16 输入输出、FP32 累加为例，主要数据流大致包括：

读取 x
读取 γ
读取 β
写 y

每元素的运算只有：

• 减法；
• 乘法；
• 加法；
• 少量统计分摊；
• 一次整体 rsqrt。

计算量相对数据流量较低。

所以大规模 LayerNorm 通常主要受：

• 显存带宽；
• 读写次数；
• Reduction 同步；

限制。

---

十六、为什么减少一次 Global Memory 读取很重要

假设传统两遍实现读取输入两次：

$$2H$$

个元素。

如果把输入缓存在寄存器中，只需要读取：

$$H$$

个元素。

在带宽受限 Kernel 中，减少一次完整输入读取可能带来明显收益。

因此现代优化经常围绕：

输入是否能够在片上保存到统计量计算完成

展开。

---

十七、什么是 RMSNorm

RMSNorm 可以看作简化版 LayerNorm。

它不减去均值，只使用均方根进行归一化：

$$r= \sqrt{ \frac{1}{N} \sum_{i=0}^{N-1}x_i^2 +\epsilon }$$

输出为：

$$y_i= \gamma_i\frac{x_i}{r}$$

通常没有 $\beta$，但具体模型实现可能有所不同。

---

十八、RMSNorm 与 LayerNorm 的区别

LayerNorm：

$$y_i= \gamma_i \frac{x_i-\mu} {\sqrt{\sigma^2+\epsilon}} +\beta_i$$

RMSNorm：

$$y_i= \gamma_i \frac{x_i} {\sqrt{ \frac{1}{N}\sum_jx_j^2+\epsilon }}$$

RMSNorm 不计算：

均值 μ
中心化 x - μ

只需要：

$$\sum_i x_i^2$$

因此归约状态更简单。

---

十九、RMSNorm 为什么更适合硬件实现

RMSNorm 只需要：

平方
求和
除以 N
加 ε
开平方倒数
乘缩放参数

不需要：

• 均值 Reduction；
• Welford 三元组；
• 中心化减法；
• Bias 参数读取。

因此可能具有：

• 更少运算；
• 更简单归约；
• 更少寄存器；
• 更简单 ASIC 数据通路；
• 更少参数带宽。

但模型是否使用 RMSNorm 取决于网络设计，不能任意把训练好的 LayerNorm 替换成 RMSNorm。

---

二十、Warp RMSNorm 的实现结构

每个 Lane 负责若干元素。

第一步，计算局部平方和：

$$s_t=\sum_{i\in t}x_i^2$$

第二步，Warp Reduction：

$$s=\sum_ts_t$$

第三步：

$$r^{-1} = \operatorname{rsqrt} \left( \frac{s}{N}+\epsilon \right)$$

第四步，每个线程输出：

$$y_i=x_i\cdot r^{-1}\cdot\gamma_i$$

如果输入保存在寄存器中，整个输入只需读取一次。

---

二十一、为什么通常使用 FP32 累加

即使输入是：

• FP16；
• BF16；
• FP8；

均值、平方和和方差通常使用 FP32 计算或累加。

原因是：

• Hidden Size 可能很大；
• 平方会扩大数值范围；
• 低精度累加误差较大；
• 方差对消减误差敏感；
• rsqrt 的输入需要稳定。

典型数据路径是：

FP16/BF16 输入
      ↓ 转换
FP32 统计与归一化
      ↓
转换回 FP16/BF16 输出

---

二十二、rsqrt 为什么常见

归一化需要：

$$\frac{1}{\sqrt{\sigma^2+\epsilon}}$$

可以直接计算：

float inv_std = rsqrtf(var + epsilon);

然后每个元素执行：

y = (x - mean) * inv_std;

而不是每个元素单独执行除法和平方根。

这样只需要：

• 每行一次 rsqrt；
• 每元素一次乘法。

---

二十三、Scale 和 Bias 融合

普通执行可能分成：

Kernel 1：
计算归一化结果 z

Kernel 2：
y = γz + β

中间结果 $z$ 需要：

• 写入显存；
• 再读取。

更合理的是直接融合：

$$y_i= \gamma_i\hat{x}_i+\beta_i$$

在同一个 Kernel 中完成。

这样：

归一化结果保存在寄存器
        ↓
直接乘 γ、加 β
        ↓
写最终输出

避免中间张量。

---

二十四、残差连接融合

Transformer 中常见：

$$z=x+\operatorname{Sublayer}(x)$$

然后：

$$y=\operatorname{LayerNorm}(z)$$

朴素实现：

Kernel 1：
z = x + residual
写 z 到 HBM

Kernel 2：
读取 z
执行 LayerNorm
写 y

融合后：

读取 x
读取 residual
        ↓
寄存器中相加
        ↓
直接参与 LayerNorm Reduction
        ↓
写最终 y

可以减少一次中间张量写入和读取。

---

二十五、Residual + Dropout + LayerNorm 融合

训练中还可能有：

$$z=x+\operatorname{Dropout}(f(x))$$

然后：

$$y=\operatorname{LayerNorm}(z)$$

可以融合：

读取主分支
读取残差
生成或重建 Dropout Mask
执行残差加法
进行 LayerNorm
写最终结果

这样能显著减少显存流量。

但融合会增加：

• Kernel 代码复杂度；
• 寄存器使用；
• 随机数状态管理；
• 数值验证难度。

---

二十六、Pre-Norm 和 Post-Norm

Transformer 中有两种常见结构。

Pre-Norm

$$y=x+F(\operatorname{Norm}(x))$$

先归一化，再进入子层。

Post-Norm

$$y=\operatorname{Norm}(x+F(x))$$

先残差相加，再归一化。

它们对算子融合的数据流不同。

Post-Norm 更自然地融合：

Residual Add + LayerNorm

Pre-Norm 则可能融合：

LayerNorm + 后续线性层输入准备

具体融合策略取决于模型计算图。

---

二十七、LayerNorm 与后续 GEMM 融合的限制

LayerNorm 输出通常进入线性层：

$$Y=\operatorname{LayerNorm}(X)W$$

理论上可以不把归一化结果完整写回 HBM，而是直接送入 GEMM。

但实现困难，因为：

• LayerNorm 按行 Reduction；
• GEMM 按 Tile 组织；
• 两者线程映射不同；
• GEMM 需要复用输入；
• LayerNorm 必须先得到整行统计量；
• 融合会增加寄存器和 Shared Memory 压力。

因此这类深度融合通常需要专门设计的数据流，而不是简单合并两段代码。

---

二十八、LayerNorm 的反向传播

设：

$$\hat{x}_i= \frac{x_i-\mu}{\sqrt{\sigma^2+\epsilon}}$$

前向：

$$y_i=\gamma_i\hat{x}_i+\beta_i$$

反向传播需要计算：

• $dx_i$；
• $d\gamma_i$；
• $d\beta_i$。

其中：

$$d\beta_i=dy_i$$

对 Batch/Token 维度还需要进一步归约。

$$d\gamma_i=dy_i\hat{x}_i$$

而 $dx$ 需要行内统计：

$$\sum_i dy_i\gamma_i$$

以及：

$$\sum_i dy_i\gamma_i\hat{x}_i$$

因此反向传播同样包含多个 Reduction。

---

二十九、LayerNorm 反向的常见形式

令：

$$g_i=dy_i\gamma_i$$

则可以写成：

$$dx_i= \frac{r}{N} \left[ Ng_i - \sum_jg_j - \hat{x}_i \sum_jg_j\hat{x}_j \right]$$

其中：

$$r=\frac{1}{\sqrt{\sigma^2+\epsilon}}$$

因此每行需要计算两个统计量：

$$S_1=\sum_jg_j$$ $$S_2=\sum_jg_j\hat{x}_j$$

然后广播给整行线程。

它仍然是一个典型的 Reduction + Elementwise Kernel。

---

三十、参数梯度为什么是另一方向的归约

$\gamma$ 和 $\beta$ 的形状通常为：

$$[H]$$

它们在所有 Batch 和 Token 之间共享。

因此：

$$d\gamma_i= \sum_{\text{rows}} dy_i\hat{x}_i$$ $$d\beta_i= \sum_{\text{rows}} dy_i$$

这里是沿行方向归约：

相同 Hidden Index
跨所有 Token 求和

它与前向的“行内归约”方向不同。

因此反向 Kernel 的数据布局和并行映射会更复杂。

---

三十一、为什么参数梯度可能需要 Atomic

不同 Block 可能同时计算同一个 $\gamma_i$ 和 $\beta_i$ 的部分梯度。

可以采用：

每个 Block 得到局部 dγ、dβ
        ↓
写部分和
        ↓
第二级归约

也可以：

每个 Block 对全局 dγ、dβ
执行 Atomic Add

如果 Hidden Size 较大、每个地址竞争程度可控，原子方案可能有效。

如果竞争严重，则多级归约更适合。

---

三十二、Nsight Compute 中的典型表现

LayerNorm 或 RMSNorm 常表现为：

Memory Throughput 较高
Compute Throughput 较低或中等
Reduction 相关 Stall
Long Scoreboard
部分 Short Scoreboard
较高寄存器使用

如果输入保存在寄存器中，可能看到：

Global Memory 流量下降
寄存器数量增加
Occupancy 下降
但 Duration 减少

这是一个典型的“低 Occupancy 但高效率”案例。

---

三十三、如何判断优化方向

情况一：显存带宽已经接近上限

应考虑：

• 融合残差；
• 融合 Bias/Scale；
• 减少输入重复读取；
• 使用低精度输入输出；
• 避免中间张量；
• 使用向量化读写。

情况二：带宽不高，Long Scoreboard 很高

可能是：

• Grid 太小；
• 每行只有一个 Warp，Warp 数不足；
• 地址不连续；
• 读取粒度过小；
• 输入行长度不均匀。

情况三：Short Scoreboard 或 Barrier 高

可能是：

• Shared Memory Reduction 开销；
• Bank Conflict；
• Block 过大；
• 同步过多。

情况四：寄存器过高并发生 Spill

应考虑：

• 减少每线程缓存元素数；
• 分两遍读取；
• 调整 Block Size；
• 减少循环展开；
• 使用 Shared Memory 替代部分寄存器。

---

三十四、ASIC 中的 LayerNorm 数据通路

一个简化的 LayerNorm 加速器可以包含：

输入流
  ↓
累加 Sum
累加 Square Sum 或 Welford
  ↓
计算 Mean / Variance
  ↓
RSQRT 单元
  ↓
输入回放
  ↓
减均值
乘 inv_std
乘 γ
加 β
  ↓
输出

核心问题是：

统计量只有在整行输入处理完成后才能得到，但归一化又需要再次使用整行输入。

---

三十五、ASIC 如何处理输入回放

主要有三种方案。

方案一：片上缓存整行

输入写入 SRAM
同时计算统计量
        ↓
统计量完成
        ↓
从 SRAM 读回输入
执行归一化

优点：

• 不需要再次访问外部存储；
• 数据流清晰。

缺点：

• 需要足够片上 SRAM；
• Hidden Size 越大，缓冲越大。

---

方案二：从外部存储重新读取

第一遍从 DRAM 读取并统计
第二遍重新从 DRAM 读取并归一化

优点：

• 片上存储需求小。

缺点：

• 外部带宽加倍；
• 能耗较高；
• 延迟较大。

---

方案三：流式分块与中间缓冲

将一行分块处理，并在片上保存部分数据或压缩状态。

但普通 LayerNorm 仍需要整行统计量，因此最终通常需要：

• 保存输入；
• 或重新读取。

这与 FlashAttention 的 Online Softmax 不完全一样，因为 LayerNorm 的每个输出也依赖最终全行均值和方差。

---

三十六、ASIC 中的 Welford 单元

一个 Welford 单元需要维护：

Count
Mean
M2

每输入一个元素，执行：

减法
除法或乘倒数
乘法
加法

为了提高吞吐率，可以：

• 多路并行处理；
• 每路得到局部状态；
• 使用树形合并单元；
• 在定点或浮点中实现；
• 对固定行长预计算倒数。

但 Welford 合并比普通求和树复杂，面积和延迟更高。

---

三十七、ASIC 中 RMSNorm 更简单

RMSNorm 只需累计：

$$\sum_i x_i^2$$

数据路径为：

x
 ↓
乘法 x²
 ↓
累加
 ↓
乘 1/N
 ↓
加 ε
 ↓
RSQRT
 ↓
与缓存输入和 γ 相乘

不需要维护均值，也不需要做中心化。

因此 RMSNorm 对专用硬件更友好。

---

三十八、定点 LayerNorm 的主要困难

定点化时需要考虑：

• 输入平方后的位宽增长；
• Sum 的累加位宽；
• Square Sum 的累加位宽；
• 均值除法；
• 方差可能因舍入变成负数；
• sqrt 和倒数实现；
• $\gamma,\beta$ 的量化尺度；
• 输出 Requantization。

假设输入是 INT8：

$$x\in[-128,127]$$

平方最大约：

$$127^2=16129$$

若累加4096个元素：

$$16129\times4096 \approx6.6\times10^7$$

至少需要约26位无符号表示，实际还要保留符号、精度和安全余量。

所以不能用INT8累加器直接计算方差。

---

三十九、方差负值保护

使用：

$$\sigma^2= E[x^2]-E[x]^2$$

时，定点或浮点误差可能得到：

$$\sigma^2=-10^{-7}$$

数学上方差不应为负。

实现中通常需要：

$$\sigma^2\leftarrow\max(\sigma^2,0)$$

然后再计算：

$$\operatorname{rsqrt}(\sigma^2+\epsilon)$$

否则可能出现非法平方根。

---

四十、LayerNorm 的完整优化思路

可以总结为：

朴素三遍读取
      ↓
Sum + Square Sum 合并为第一遍
      ↓
使用稳定的 Welford
      ↓
Warp/Block 分层归约
      ↓
短行在寄存器缓存输入
      ↓
长行选择 Shared Memory 或重新读取
      ↓
融合 γ、β
      ↓
融合 Residual / Bias / Dropout
      ↓
减少中间 HBM 流量

RMSNorm 则可以进一步简化为：

平方和归约
   ↓
RSQRT
   ↓
Scale

---

本课核心结论

第一，LayerNorm 沿每个 Token 的 Hidden Dimension 计算均值和方差。

第二，LayerNorm 包含至少两个统计量：

$$\sum x_i$$

以及方差相关归约。

第三，朴素三遍实现会多次读取输入，显存流量较大。

第四，可以同时计算 Sum 和 Square Sum，将其减少为两遍，但：

$$E[x^2]-E[x]^2$$

可能存在数值消减问题。

第五，Welford 算法通过维护 Count、Mean 和 $M_2$，提供更稳定的在线方差计算。

第六，Welford 状态可以合并，因此适合 Warp 和 Block 树形归约。

第七，短行可以由一个 Warp 处理，并将输入保存在寄存器中。

第八，长行通常由一个 Block 处理，通过 Warp Shuffle、Shared Memory 和 Block Reduction 获得统计量。

第九，寄存器缓存可以减少输入重复读取，但可能降低 Occupancy。

第十，LayerNorm 通常是内存受限算子，因此减少显存读写比增加算术吞吐率更重要。

第十一，RMSNorm 不计算均值，只进行平方和归约，因此实现更简单。

第十二，低精度输入通常仍使用 FP32 进行统计和归一化。

第十三，Residual、Bias、Scale、Dropout 与 Norm 的融合可以避免中间张量写回 HBM。

第十四，LayerNorm 反向传播同样需要多个行内归约和跨行参数梯度归约。

第十五，ASIC 实现的核心问题是如何在统计量完成后重新获得整行输入。

---

后续课程如何收束

目前已经覆盖了 GPU 架构学习中的绝大多数主干内容：

CPU 与 GPU 区别
SIMT、Thread、Warp、Block、Grid
SM 内部结构
分支发散与同步
寄存器、Shared Memory、Cache、HBM
Occupancy 与调度
Roofline 与性能瓶颈
矩阵乘法、卷积、Tensor Core
低精度计算
多 GPU、NVLink、NCCL
图形管线与 RT Core
完整 GPU 芯片层级
NVIDIA 架构演进
PTX、SASS 与编译器
Nsight Compute
Reduction、Softmax、FlashAttention
LayerNorm、RMSNorm

继续无限扩展会进入很多专业分支。为了形成一个完整闭环，再讲下面四课基本就足够了。

第二十四课：Transformer 中的 MLP、GEMM 与算子融合

包括：

Linear 层如何映射为 GEMM
GELU、SiLU 与 SwiGLU
Bias + Activation Fusion
GEMM Epilogue
为什么 MLP 占大量计算
Tensor Parallel 下的列并行和行并行

第二十五课：稀疏、不规则访存和图计算

包括：

稀疏矩阵为什么不一定快
CSR/CSC 格式
Gather/Scatter
负载不均衡
原子冲突
结构化稀疏与非结构化稀疏
GPU 不擅长什么任务

第二十六课：如何从零设计一个简化 SIMT GPU

包括：

指令集
Thread Context
Warp PC
Active Mask
Scoreboard
Warp Scheduler
SIMD Datapath
LSU
Shared Memory
Block Dispatcher
最小可运行 GPU RTL 框架

第二十七课：GPU 架构学习总结与分析方法

最终整理为：

看到一个 GPU 架构图该怎么看
看到一个 Kernel 该怎么分析
看到一个性能报告该怎么定位
看到一个算子该怎么做硬件映射
GPU、NPU、ASIC 如何选择
后续应阅读哪些文档和论文

讲完第二十七课，整个基础到进阶课程就可以结束。下一课进入 MLP、GEMM Epilogue 和算子融合。