4. 典型融合模式：Conv+BN+ReLU、Conv+Add+ReLU、MatMul+Add

好，咱们今天聊点实在的。算子融合这个事儿，说白了就是把几个小操作捏成一个，减少显存读写，提升计算效率。我这些年调过不少模型，踩过不少坑，今天挑三个最经典的融合模式跟你掰扯掰扯。

4.1 Conv+BN+ReLU 融合

这个组合，你肯定见过。几乎每个CNN里都有它的身影。为什么要把它们仨揉在一起？因为BN和ReLU都是element-wise操作，计算量不大，但访存开销不小。

我个人习惯把BN的公式先展开。BN在推理阶段，其实就是一个线性变换：

BN(x) = γ * (x - μ) / √(σ² + ε) + β
      = (γ / √(σ² + ε)) * x + (β - γ * μ / √(σ² + ε))
      = α * x + b

你看，推理时BN就变成了一个简单的乘加。那Conv本身也是线性变换，两个线性变换叠加，还是一个线性变换。所以我们可以把BN的系数直接吸收进Conv的权重里。

具体怎么做？

// 假设 Conv 权重为 W，偏置为 bias
// BN 参数：gamma, beta, mean, var, eps

// 计算融合后的权重和偏置
alpha = gamma / sqrt(var + eps)
W_fused = W * alpha
bias_fused = bias * alpha + (beta - gamma * mean / sqrt(var + eps))

// 融合后的 Conv 输出直接过 ReLU
output = ReLU(Conv(input, W_fused, bias_fused))

核心收益：省掉一次BN的全局访存，省掉一次ReLU的独立kernel launch。显存带宽节省约20%-30%。

我的经验：我在项目中遇到过一个问题——训练时BN的running mean和var是不断更新的，但推理时是固定的。所以这个融合只能在推理阶段做。训练阶段别乱搞，否则梯度传不过去。

4.2 Conv+Add+ReLU 融合

这个模式在ResNet里特别常见。就是残差连接：主路径的Conv输出，加上旁路的输入，再过ReLU。

你想想看，如果不融合，流程是这样的：

Conv计算 → 写回显存
从显存读Conv结果，读旁路输入 → Add → 写回显存
从显存读Add结果 → ReLU → 写回显存

三次显存读写，太浪费了。融合后怎么做？

// 伪代码：Conv+Add+ReLU 融合 kernel
__global__ void fused_conv_add_relu(float* input, float* skip, float* output, 
                                     float* weight, float* bias, int H, int W, int C) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx >= H * W * C) return;
    
    // 计算 Conv 的一个输出元素
    float conv_val = 0.0f;
    for (int kh = 0; kh < K; kh++)
        for (int kw = 0; kw < K; kw++)
            conv_val += input[...] * weight[...];
    conv_val += bias[oc];
    
    // 直接加旁路，然后 ReLU，一次写回
    float add_val = conv_val + skip[idx];
    output[idx] = add_val > 0 ? add_val : 0;
}

嗯，这里要注意：旁路输入和Conv输出的shape必须一致。如果通道数不同，需要先做1x1 Conv对齐，那就不能简单融合了。

我曾经踩过的坑：有一次我图省事，直接把Add和ReLU写在一个kernel里，但Conv是单独调用的。结果发现性能反而下降了。为什么？因为Conv输出写回显存后，Add+ReLU kernel又要读回来，显存带宽成了瓶颈。后来我把Conv也塞进去，三合一，性能才上去。

4.3 MatMul+Add 融合

这个模式在Transformer里太常见了。QKV投影、FFN层，到处都是MatMul后面跟一个bias add。

说白了，MatMul就是矩阵乘法，Add就是加偏置。这两个操作天然可以融合——在MatMul的最后一个累加循环里，直接把偏置加上去。

// 传统做法：两步走
C = MatMul(A, B);  // MxN 矩阵
D = C + bias;      // 广播加法

// 融合做法：一步到位
// 在 MatMul 的 inner loop 里，计算完 C[i][j] 后直接加 bias[j]
for (int i = 0; i < M; i++) {
    for (int j = 0; j < N; j++) {
        float sum = 0.0f;
        for (int k = 0; k < K; k++) {
            sum += A[i * K + k] * B[k * N + j];
        }
        C[i * N + j] = sum + bias[j];  // 融合点
    }
}

关键点：这个融合能省掉一次N维向量的全局写和一次全局读。对于大矩阵（比如M=4096, N=4096），省掉的显存带宽相当可观。

我个人习惯在实现时，把bias加载到寄存器里，避免重复从显存读。如果bias比较小（比如N=4096），可以全部塞进shared memory甚至寄存器。

4.4 三种融合模式的对比

融合模式	适用场景	省掉的访存次数	典型加速比
Conv+BN+ReLU	CNN推理	2次全局写 + 2次全局读	1.2x - 1.5x
Conv+Add+ReLU	ResNet类网络	2次全局写 + 1次全局读	1.3x - 1.6x
MatMul+Add	Transformer/MLP	1次全局写 + 1次全局读	1.1x - 1.3x

你可能会问：为什么MatMul+Add的加速比看起来没那么高？因为MatMul本身是计算密集型，访存占比相对小。而Conv+BN+ReLU里，BN和ReLU都是访存密集型，融合后收益更明显。

避坑指南：我曾经在优化一个BERT模型时，把所有MatMul后面的Add都融合了，结果发现收益不到5%。后来一分析，瓶颈在Softmax和LayerNorm上。所以融合要挑对地方，别盲目。

4.5 知识体系图

下面这张图，我把三种融合模式的核心逻辑画出来了。你看一眼就能明白它们之间的关系。

这三种融合模式，说白了都是同一个思路：把多个小操作合并成一个，减少中间结果的显存读写。区别只在于具体的操作类型和融合方式。

好了，今天就聊到这儿。这三种模式你掌握了，大部分模型的推理优化就够用了。下次遇到新的融合场景，你也可以用同样的思路去分析——哪些操作是线性的？哪些是element-wise的？能不能合并？想清楚这几个问题，融合方案自然就有了。

📚 课程目录

4. 典型融合模式：Conv+BN+ReLU、Conv+Add+ReLU、MatMul+Add

4.1 Conv+BN+ReLU 融合

4.2 Conv+Add+ReLU 融合

4.3 MatMul+Add 融合

4.4 三种融合模式的对比

4.5 知识体系图