4. 水平融合：将多个相同形状的算子合并执行

水平融合，说白了就是把多个长得一模一样的算子，合并成一个来执行。

我刚开始接触这个概念时，觉得这有什么好讲的？不就是把几个加法拼一起吗？后来在项目中踩过坑才明白——这里面的门道，远比想象中深。

4.1 什么是水平融合

先看个场景。你的模型里有三个独立的卷积层，输入输出形状完全一样：

# 三个独立的卷积
x1 = Conv2D(64, 3, padding='same')(input)
x2 = Conv2D(64, 3, padding='same')(input)
x3 = Conv2D(64, 3, padding='same')(input)

# 然后合并
output = x1 + x2 + x3

每个卷积都要读一次 input，写一次结果。三次读写，三次 kernel 启动。内存带宽全浪费在重复搬运上了。

水平融合的做法是：

# 合并成一个算子
output = FusedConvSum(input, weights=[w1, w2, w3])

一次读取 input，三个卷积同时算，结果直接累加。读写次数从 6 次降到 2 次。

核心收益：减少内存访问次数，提升计算密度。

4.2 适用条件

不是所有算子都能水平融合。我总结了几条硬性条件：

输入形状必须一致——这是前提，否则没法合并
算子类型可以不同——比如 Conv + BatchNorm + ReLU 也能水平融合
输出需要合并——通常是求和、取最大、拼接等操作
没有数据依赖——算子之间不能互相等结果

举个例子，ResNet 里的残差分支就特别适合水平融合。主路径和捷径的卷积形状一样，最后做加法。我优化过一个 ResNet-50 模型，光这一项就省了 15% 的推理时间。

4.3 实现方式

水平融合的实现，我习惯分三步走：

算子识别——遍历计算图，找出所有形状匹配、可并行的算子组
内存规划——为合并后的算子分配共享的输入/输出缓冲区
代码生成——生成一个 kernel，内部循环执行多个算子

这里给个简化版的伪代码：

// 水平融合 kernel 伪代码
__global__ void fused_kernel(float* input, float* output,
                             float* w1, float* w2, float* w3) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    
    // 一次读取
    float x = input[idx];
    
    // 三个卷积并行计算
    float y1 = conv_single(x, w1);
    float y2 = conv_single(x, w2);
    float y3 = conv_single(x, w3);
    
    // 结果累加，一次写入
    output[idx] = y1 + y2 + y3;
}

你看，输入只读一次，输出只写一次。中间的计算全部在寄存器里完成。

小技巧：如果算子数量太多（比如超过 8 个），寄存器可能会不够用。我一般建议控制在 4-6 个，性能收益最明显。

4.4 性能分析

水平融合的效果，跟算子的计算强度有关。我画了张图来说明：

从图上能看出两个规律：

计算密集型算子（比如卷积）收益更大，因为计算时间占比高，融合后能掩盖内存延迟
带宽密集型算子（比如逐元素操作）收益相对小，但融合后也能省掉不少内存搬运
算子数量越多，收益越趋于饱和——到 6 个以上时，提升就不明显了

注意：水平融合不是万能的。我曾经在一个模型里强行融合了 12 个算子，结果寄存器溢出，性能反而下降了 8%。适可而止，别贪多。

4.5 实战案例

说个我实际优化过的例子。一个目标检测模型，里面有 4 个并行的 3x3 卷积，输入都是 128x128x64 的特征图，最后做 element-wise 求和。

优化前：

4 次 kernel 启动，每次读 2MB、写 2MB
总内存访问：16MB
延迟：2.3ms

水平融合后：

1 次 kernel 启动，读 2MB、写 2MB
总内存访问：4MB
延迟：1.1ms

性能提升 52%。而且代码量还减少了，维护起来更省心。

我的习惯：每次写新模型时，先看看有没有形状相同的并行分支。有的话，直接设计成融合算子。别等优化阶段再改，那时候改起来更麻烦。

4.6 实现注意事项

水平融合看着简单，但有几个坑得注意：

问题	表现	解决方案
寄存器溢出	融合算子过多，寄存器不够用，数据被迫写到 local memory	限制融合数量 ≤ 6，或使用 shared memory 做中转
线程束发散	不同分支的计算路径不同，导致 warp 内线程执行不同指令	确保所有分支的计算模式一致，或拆分成多个 warp
精度问题	融合后计算顺序改变，浮点累加结果可能有微小差异	使用更高精度累加（如 float32 累加 float16）
调试困难	融合后的 kernel 难以单步调试	保留原始版本做对比验证，用数值差分检查

嗯，这里要特别提一下精度问题。我之前有个语音模型，融合后精度差了 0.3%。查了半天才发现是 float16 累加时溢出了。换成 float32 累加就解决了。

水平融合是个好工具，但别滥用。我见过有人把整个网络都塞进一个 kernel，结果编译都过不了。合适的才是最好的。

想要更多资料，联系可可

📚 课程目录

4. 水平融合：将多个相同形状的算子合并执行

4.1 什么是水平融合

4.2 适用条件

4.3 实现方式

4.4 性能分析

4.5 实战案例

4.6 实现注意事项