常见融合模式:从理论到实战
算子融合,说白了就是把多个连续的小算子合并成一个大的计算核。为什么要这么做?因为每次调用一个算子,都有启动开销、显存读写开销。合并之后,中间结果不用写回显存,直接在寄存器里流转,速度能快好几倍。
我个人习惯把融合分成两类:一类是计算密集型融合,比如 Conv+BatchNorm;另一类是访存密集型融合,比如 Conv+ReLU。前者省计算,后者省带宽。今天咱们就把最常见的四种模式挨个过一遍。
1. Conv + BatchNorm 融合
这个模式在推理场景中几乎是必做的优化。BatchNorm 在训练时是独立算子,但推理时它的参数是固定的——均值、方差、缩放因子、偏移量都是常量。那为什么不把这些常量直接吸收进卷积的权重和偏置里呢?
公式推导其实很简单。卷积输出是 y = W * x + b,BatchNorm 推理公式是 z = γ * (y - μ) / √(σ² + ε) + β。把 y 代入,合并同类项,就得到新的权重 W' = γ * W / √(σ² + ε),新的偏置 b' = γ * (b - μ) / √(σ² + ε) + β。
核心要点:融合后的卷积层直接输出归一化后的结果,BatchNorm 算子从计算图中消失。推理速度能提升 10%~30%,具体取决于卷积核大小和通道数。
我在项目中遇到过一个问题:有些框架的 BatchNorm 实现里,ε 的默认值很小(比如 1e-5),但某些硬件对除零特别敏感。有一次在 FPGA 上部署时,因为 σ² 恰好接近 0,导致除零异常。后来我养成了一个习惯——融合前先检查方差,如果某个通道的方差几乎为 0,就给它一个保底值。
// 伪代码:Conv + BN 融合
void fuse_conv_bn(float* conv_weight, float* conv_bias,
float* bn_mean, float* bn_var,
float* bn_gamma, float* bn_beta,
float eps, int out_channels, int kernel_size) {
for (int c = 0; c < out_channels; c++) {
float scale = bn_gamma[c] / sqrtf(bn_var[c] + eps);
// 融合权重
for (int i = 0; i < kernel_size; i++) {
conv_weight[c * kernel_size + i] *= scale;
}
// 融合偏置
conv_bias[c] = scale * (conv_bias[c] - bn_mean[c]) + bn_beta[c];
}
}
我的小技巧:融合后的权重精度会略微变化。如果模型对精度极其敏感(比如医学影像),建议融合后做一次精度验证,误差通常在 1e-5 量级,基本无感。
2. Conv + ReLU 融合
这个融合更简单,但收益同样可观。ReLU 本质上就是个 max(0, x) 操作,没有可学习的参数。融合的思路是:在卷积的输出阶段,每个元素计算完后直接判断是否小于 0,如果是就直接写 0,否则写原值。
你想想看,如果不融合,卷积的输出要先写进显存,然后 ReLU 再读出来,处理完再写回去。这一来一回,显存带宽就浪费了。融合后,ReLU 的读写完全消失,数据在寄存器里就处理完了。
性能收益:Conv+ReLU 融合通常能减少约 30%~50% 的显存读写量。对于带宽受限的模型(比如 MobileNet 系列),这个优化效果非常明显。
我曾经在优化一个实时视频处理模型时,发现 Conv+ReLU 融合后,单帧处理时间从 12ms 降到了 9ms。嗯,就是这 3ms 的差距,让帧率从 83fps 跳到了 111fps,刚好跨过了 100fps 的门槛。
// CUDA 示例:Conv + ReLU 融合 kernel
__global__ void fused_conv_relu_kernel(
float* input, float* weight, float* bias,
float* output, int H, int W, int C, int K) {
int tx = threadIdx.x, ty = threadIdx.y;
int row = blockIdx.y * blockDim.y + ty;
int col = blockIdx.x * blockDim.x + tx;
if (row >= H || col >= W) return;
float sum = bias[blockIdx.z];
// 卷积计算(简化)
for (int c = 0; c < C; c++) {
sum += input[row * W * C + col * C + c] * weight[blockIdx.z * C + c];
}
// 融合 ReLU:直接在这里做激活
output[row * W * K + col * K + blockIdx.z] = (sum > 0) ? sum : 0.0f;
}
注意:不是所有激活函数都适合融合。像 ReLU、LeakyReLU 这种 element-wise 且无跨通道依赖的,融合起来很轻松。但像 Softmax、Sigmoid 这种涉及全局统计或指数运算的,融合收益有限,有时反而会降低灵活性。
3. Conv + Add + ReLU 融合
这个模式在残差网络(ResNet)里特别常见。一个残差块的结构通常是:Conv → BN → ReLU → Conv → BN → Add(加上 shortcut)→ ReLU。你看,最后三步是 Conv+BN → Add → ReLU,这三者可以融合成一个算子。
融合的关键在于 Add 操作。Add 是两个张量逐元素相加,本身没有参数。所以我们可以把卷积的输出直接和 shortcut 相加,然后过 ReLU。这样,卷积的输出不用单独写显存,Add 的结果也不用单独写显存,全部在一个 kernel 里完成。
| 融合方式 | 减少的 kernel 数 | 减少的显存读写 | 典型加速比 |
|---|---|---|---|
| Conv+BN | 1 | ~40% | 1.1x ~ 1.3x |
| Conv+ReLU | 1 | ~50% | 1.2x ~ 1.5x |
| Conv+Add+ReLU | 2 | ~60% | 1.3x ~ 1.8x |
| MatMul+Add | 1 | ~30% | 1.1x ~ 1.4x |
我记得有一次在优化 ResNet-50 时,把残差块里的 Conv+BN+Add+ReLU 全部融合后,整个模型的推理速度提升了 22%。而且精度完全没变——因为融合是数学等价的。
实现要点:融合 Conv+Add+ReLU 时,要注意 Add 的输入对齐。如果 shortcut 是 1x1 卷积降维后的结果,需要确保两个分支的 shape 完全一致。我习惯在融合前加一个 assert 检查。
4. MatMul + Add 融合
这个模式在 Transformer 和全连接层里非常普遍。说白了就是 y = x * W + b,也就是全连接层本身。很多框架把 MatMul 和 Add 拆成两个算子,但本质上它们是一个线性变换。
融合的思路很简单:把偏置 b 直接加到 MatMul 的输出上。在 GPU 上,这通常通过一个 fused kernel 实现——矩阵乘法计算完每个元素后,立即加上对应的偏置值,然后写回显存。
为什么这个融合重要?因为 Transformer 模型里有大量的 QKV 投影、FFN 层,每个都包含 MatMul+Add。如果不融合,每个 都要多一次 kernel launch 和一次显存读写。对于大模型来说,这个开销累积起来非常可观。
// 伪代码:MatMul + Add 融合
void fused_matmul_add(float* A, float* B, float* bias,
float* C, int M, int N, int K) {
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];
}
// 融合 Add:加上偏置
C[i * N + j] = sum + bias[j];
}
}
}
避坑指南:我曾经在优化一个 模型时,发现 MatMul+Add 融合后精度出现了微小偏差。排查了半天,发现是偏置的广播方式有问题——bias 是 [N] 维向量,但 MatMul 的输出是 [M, N] 矩阵,偏置需要沿着 M 维度广播。如果 kernel 实现里广播逻辑写错了,结果就会出错。所以,融合后一定要做数值对比测试。
融合策略的选择原则
说了这么多,到底什么时候该融合,什么时候不该?我总结了几条经验:
- 优先融合访存密集型的组合:比如 Conv+ReLU、Conv+Add,这些操作本身计算量小,但访存开销大,融合收益最明显。
- 计算密集型的融合要谨慎:比如 Conv+BN,虽然收益也不错,但 BN 的参数吸收会改变权重数值,对量化模型可能有影响。
- 考虑硬件特性:在 GPU 上,融合 kernel 的寄存器压力不能太大。如果融合后的 kernel 寄存器溢出,性能反而会下降。我一般在融合后做一次 occupancy 分析。
- 不要过度融合:把整个网络融合成一个 kernel 听起来很美好,但实际上不可能。融合的粒度要适中,一般以 3~5 个算子为一组比较合适。
总结一下:算子融合是推理优化的基本功。Conv+BN 融合省计算,Conv+ReLU 融合省带宽,Conv+Add+ReLU 融合是残差网络的标配,MatMul+Add 融合是 Transformer 的必备技能。把这四种模式吃透,你就能应对 80% 以上的融合场景。
嗯,这四种模式在实际项目中我几乎每天都在用。下次当你看到一个推理模型跑得慢时,别急着上量化或者剪枝——先看看算子融合做了没有。很多时候,融合带来的收益比你想的要大得多。