4. 算子融合基础策略：Horizontal Fusion 与 Vertical Fusion

聊到算子融合，很多同学第一反应就是「把几个算子拼成一个」。没错，但怎么拼、拼哪些，这里头门道不少。我做了这么多年端侧推理引擎，踩过最多的坑就是融合策略选错了，结果推理速度反而更慢。

今天咱们就聚焦 SwiGLU 这个结构，把 Horizontal Fusion（水平融合）和 Vertical Fusion（垂直融合）这两个基础策略彻底讲透。

4.1 先搞清楚：什么是水平融合？什么是垂直融合？

说白了，这两种融合的区别就是「往宽了拼」还是「往深了拼」。

• 水平融合（Horizontal Fusion）：把同一层里、输入输出形状相似的算子合并。比如两个并行的全连接层，它们吃同样的输入，各自算完再拼起来。这种融合减少的是 kernel launch 开销和显存带宽占用。
• 垂直融合（Vertical Fusion）：把前后串联的算子合并。比如一个全连接后面接一个激活函数，再后面接一个 LayerNorm。这种融合减少的是中间结果的读写开销。

你想想看，SwiGLU 这个结构天然就适合这两种融合。为什么？因为它既有并行的分支（两个全连接层），又有串联的激活和门控操作。

4.2 SwiGLU 的结构回顾

先画个图，帮大家回忆一下 SwiGLU 的计算流程。我习惯用 SVG 画这种结构图，清晰直观。

从图上可以清楚看到：水平融合的目标是那两个并行的全连接层 W1 和 W2；垂直融合的目标是 W1→SiLU 这条链，以及 W2→直接传递→逐元素相乘 这条链。

4.3 水平融合在 SwiGLU 上的应用

我在项目中遇到过这样一个场景：一个 7B 模型的 SwiGLU 层，两个全连接层的权重矩阵形状一模一样，都是 [hidden_size, intermediate_size]。当时我一看，这不就是典型的水平融合机会吗？

水平融合的具体做法是这样的：

// 融合前：两个独立的全连接计算
float* x = ...;  // 输入，形状 [batch, hidden_size]
float* w1 = ...; // 权重1，形状 [hidden_size, intermediate_size]
float* w2 = ...; // 权重2，形状 [hidden_size, intermediate_size]

// 分别计算
float* out1 = matmul(x, w1);  // 形状 [batch, intermediate_size]
float* out2 = matmul(x, w2);  // 形状 [batch, intermediate_size]

// 融合后：合并为一个更大的矩阵乘法
float* w_combined = concat(w1, w2, axis=1);  // 形状 [hidden_size, 2 * intermediate_size]
float* out_combined = matmul(x, w_combined);  // 形状 [batch, 2 * intermediate_size]

// 然后切分
float* out1 = out_combined[:, :intermediate_size];
float* out2 = out_combined[:, intermediate_size:];

嗯，这里要注意：水平融合不是简单的「把两个矩阵拼起来就完事」。它有几个关键收益：

• 减少 kernel launch 次数：从两次 matmul 变成一次，GPU 上的调度开销直接减半
• 提高计算密度：合并后的矩阵更大，更容易达到 GPU 的峰值算力
• 减少显存访问：输入 x 只需要从显存读一次，而不是两次

4.4 垂直融合在 SwiGLU 上的应用

垂直融合的思路更直接——把前后串联的算子「粘」在一起。在 SwiGLU 里，最典型的垂直融合是：全连接层 + SiLU 激活函数。

我曾经犯过一个错误：把全连接和 SiLU 分开实现，结果中间结果写回显存又读出来，白白浪费了带宽。后来改成融合实现，速度直接翻了一倍。

垂直融合的代码实现大致是这样：

// 融合前：分开计算
float* out = matmul(x, w);  // 全连接
for (int i = 0; i < n; i++) {
    out[i] = out[i] / (1.0f + expf(-out[i]));  // SiLU
}

// 融合后：在矩阵乘法 kernel 内部直接做激活
// 伪代码，实际是在 CUDA kernel 里实现的
__global__ void fused_matmul_silu(float* x, float* w, float* out, int M, int N, int K) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    
    float sum = 0.0f;
    for (int k = 0; k < K; k++) {
        sum += x[row * K + k] * w[k * N + col];
    }
    
    // 直接在 kernel 内部做 SiLU 激活
    out[row * N + col] = sum / (1.0f + expf(-sum));
}

除了全连接+SiLU，SwiGLU 里还有另一个垂直融合机会：逐元素相乘 + 后面的线性层。不过这个融合稍微复杂一些，需要处理两个输入分支的对齐问题。

4.5 水平融合 vs 垂直融合：怎么选？

很多同学会问：这两种融合策略，到底哪个更重要？我的答案是：看你的瓶颈在哪里。

场景	推荐策略	原因
显存带宽受限	垂直融合	减少中间结果的读写，直接降低带宽压力
计算单元利用率低	水平融合	增大计算粒度，提高计算密度
小 batch 推理	水平融合优先	kernel launch 开销占比大，减少 launch 次数收益明显
大 batch 推理	垂直融合优先	计算本身已经饱和，减少带宽浪费更关键
端侧设备（手机、IoT）	两者结合	端侧资源有限，需要同时优化计算和带宽

说白了，没有银弹。我在做端侧推理引擎时，通常的做法是：先做水平融合，再做垂直融合。因为水平融合的改动相对独立，不容易引入 bug；垂直融合涉及算子内部的逻辑修改，需要更谨慎的测试。

4.6 实战中的避坑指南

最后分享几个我踩过的坑，希望能帮大家少走弯路：

• 不要盲目融合：我曾经把 SwiGLU 里的所有算子都塞进一个 kernel，结果编译时间从 10 秒变成 10 分钟，而且性能没提升。融合要「精准打击」，不是「大包大揽」。
• 注意数据类型对齐：水平融合时，两个分支的权重精度可能不同（比如一个 FP16、一个 INT8）。强行融合会导致精度损失，需要做量化对齐。
• 端侧设备要特别小心寄存器压力：手机 GPU 的寄存器数量远少于桌面 GPU，垂直融合时很容易 spill。我建议在端侧做垂直融合时，先做 profilling，看看寄存器使用情况。
• 测试覆盖要全：融合后的 kernel 逻辑更复杂，边界条件更多。我习惯写一个「融合前 vs 融合后」的数值对比测试，确保每个元素的误差在 1e-5 以内。

---

公众号:蓝海资料掘金营，微信deep3321