4、常见融合模式（下）：Add+ReLU融合、MatMul+Add融合、LayerNorm融合
好，咱们接着聊算子融合。上一节我们把 Element-wise 类的融合讲得差不多了，这一节要啃的，是几个更“重”一点的融合模式。说实话，这几个模式在实际工程里出现的频率极高，几乎每个大模型推理框架都会做。我个人习惯把它们称为“黄金搭档”——因为它们实在太常一起出现了。
4.1 Add+ReLU 融合：顺手的事
先看一个最简单的。Add 操作后面紧跟一个 ReLU，这在残差网络里遍地都是。你想想看，一个张量经过卷积或者全连接层，输出结果要和残差分支相加，然后过 ReLU。如果不做融合，GPU 要启动两次 kernel：一次加法，一次 ReLU。
为什么要融合？说白了，就是减少访存。加法操作把结果写回显存，ReLU 又读出来，这中间白白浪费了一次全局内存读写。融合之后，加法做完直接判断正负，结果只写一次。
我在项目中遇到过一种情况：有些框架的 ReLU 实现是 in-place 的，也就是直接在输入张量上修改。这时候如果你强行把 Add 和 ReLU 融合，反而会破坏残差分支的数据。嗯，这里要注意——融合前一定要确认数据流的依赖关系。
融合后的计算逻辑：
// 伪代码：Add+ReLU 融合 kernel
__global__ void fused_add_relu(float* c, const float* a, const float* b, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
float sum = a[idx] + b[idx];
c[idx] = sum > 0 ? sum : 0; // ReLU 就地完成
}
}这个 kernel 看起来平平无奇，但性能提升很可观。我记得有一次在 ResNet-50 上做 profiling，光这一个融合就让整体推理速度提升了 3% 左右。别小看这 3%，在工业级部署里，每一点性能都是钱。
4.2 MatMul+Add 融合：大模型的心脏
接下来这个，可以说是 Transformer 架构里最核心的融合了。MatMul（矩阵乘法）后面跟一个 Add（偏置加法），这就是全连接层的标准形式。你去看任何一个大模型，不管是 BERT 还是 GPT，里面全是这玩意儿。
为什么不分开做？因为矩阵乘法本身已经是计算密集型了，加一个偏置几乎不增加计算量，但如果不融合，就要多一次 kernel 启动和一次显存写入。说白了，融合就是“顺手把偏置加上”，几乎零成本。
我曾经在优化一个 7B 参数的模型时，发现推理框架里 MatMul 和 Add 是分开的 kernel。我当时就纳闷了——这明显是没做融合优化啊。改完之后，单层 MLP 的延迟从 1.2ms 降到了 0.9ms，效果立竿见影。
实现技巧： 很多框架（比如 TensorRT、ONNX Runtime）会把 MatMul+Add 融合成一个 Gemm 算子。在 cuBLAS 里，你可以直接调用 cublasGemmEx 并传入偏置指针，让矩阵乘法在计算过程中就把偏置加上。这样既利用了 Tensor Core 的高吞吐，又省掉了额外的 kernel。
这里有个坑，我必须要提醒你。有些框架的 MatMul 输出是行优先存储，而偏置是列优先的，或者维度不匹配。我曾经因为偏置的 broadcast 规则没搞清楚，调试了一整个下午。所以，做融合之前，先确认好数据布局和 broadcast 语义。
4.3 LayerNorm 融合：精度与性能的博弈
LayerNorm 的融合，比前面两个要复杂一些。LayerNorm 本身包含多个步骤：计算均值、计算方差、归一化、缩放和平移。如果不融合，每一步都要启动一个 kernel，访存开销巨大。
标准的 LayerNorm 融合，就是把所有步骤塞进一个 kernel 里。大致流程是这样的：
- 读取一行数据（一个样本的所有特征）
- 计算均值和方差（一次遍历或两次遍历）
- 用均值和方差做归一化
- 乘以 gamma 并加上 beta
- 写回结果
这样做的好处很明显：数据只从显存读一次，写一次。对于大模型来说，LayerNorm 的输入维度通常很大（比如 4096 或 8192），一次融合能省掉 3-4 次全局内存访问。
注意： LayerNorm 融合对数值精度有影响。因为均值和方差是在同一个 kernel 里计算的，如果使用单精度浮点，累加过程中可能出现精度损失。我建议在训练场景下使用 FP32 做累加，推理场景下可以用 FP16 但要做好数值监控。
我曾经在优化一个 13B 参数的模型时，发现 LayerNorm 融合后，某些层的输出和原始实现有 1e-3 级别的差异。排查下来，是因为方差计算时用了半精度累加，导致数值不稳定。后来改成 FP32 累加，问题就解决了。
4.4 三种融合模式的对比
为了方便你快速对比，我把这三种融合模式的关键信息整理成了表格：
| 融合模式 | 计算强度 | 访存节省 | 常见场景 | 实现难度 |
|---|---|---|---|---|
| Add+ReLU | 低 | 1 次写 + 1 次读 | 残差网络、CNN | 简单 |
| MatMul+Add | 高 | 1 次 kernel 启动 | 全连接层、Transformer | 中等 |
| LayerNorm | 中等 | 3-4 次访存 | 大模型、RNN | 较难 |
从表格里能看出来，Add+ReLU 是最容易上手的，MatMul+Add 是性价比最高的，而 LayerNorm 融合虽然难，但收益也最大。我个人建议，如果你刚开始做算子融合，先从 Add+ReLU 练手，再挑战 MatMul+Add，最后啃 LayerNorm。
4.5 融合后的性能分析
最后，我想用一张流程图来展示这三种融合模式在模型中的位置和关系。这张图能帮你快速理解它们是如何协同工作的。
从这张图能看出来，这三种融合模式在 Transformer 的一层里是串联出现的。你优化了其中一个，整层的性能都会受益。但反过来，如果有一个没做好，就会成为瓶颈。我见过不少团队，把 MatMul+Add 优化得特别好，但 LayerNorm 还是分开的 kernel，结果整体延迟还是下不来。
总结一下：
- Add+ReLU 融合：简单高效，适合入门
- MatMul+Add 融合：大模型标配，性价比之王
- LayerNorm 融合：收益大但坑也多，需要细心调优
好了，这一节的内容就到这里。这三种融合模式，你可以在自己的模型上试试看。先从 Add+ReLU 开始，感受一下融合带来的性能提升，然后再挑战更复杂的。记住，融合不是万能的，但不融合是万万不能的。
想要更多资料，联系可可