1、算子融合概述：什么是算子融合、为什么需要算子融合、算子融合的收益分析

1.1 什么是算子融合？

算子融合，说白了就是把多个连续的计算操作合并成一个。你想想看，在深度学习模型里，我们经常看到 Conv + BN + ReLU 这种组合。传统做法是三个算子挨个执行，每个都要读写一次数据。融合之后，这三个操作变成一个大的 kernel，数据从显存读一次，算完所有结果再写回去。

我习惯把算子融合比作「流水线上的工序合并」。假设你要给零件喷漆、烘干、打磨，分开做的话，每个工序之间都要搬运零件。合并成一个工位后，零件不用挪窝，一气呵成。硬件加速的本质逻辑，就是这个道理。

核心定义：算子融合是将多个计算密集且数据依赖连续的算子，合并为一个复合算子，在单次 kernel 启动中完成全部计算。

1.2 为什么需要算子融合？

这个问题我经常被问到。很多刚入行的同学觉得，分开算也没什么不好，反正 GPU 算得快。但实际部署过就知道，瓶颈往往不在计算本身，而在数据搬运。

我曾经在一个边缘端项目里踩过坑。模型跑在 Jetson 上，Conv + BN + ReLU 分开实现，帧率死活上不去。用 profiler 一查，计算单元利用率不到 30%，大部分时间花在读写显存上。后来做了融合，帧率直接翻倍。嗯，从那以后我再也不敢轻视算子融合了。

具体来说，需要算子融合的原因有三点：

  • 访存带宽瓶颈：现代 GPU 算力增长远超显存带宽增长。分开执行时，中间结果反复读写，带宽被白白浪费。
  • kernel launch 开销：每次启动一个 kernel，CPU 都要下发指令、同步。算子多了，启动开销累积起来很可观。
  • 计算密度不足：小算子往往计算量小、访存量大，无法充分利用硬件计算单元。

个人经验：我建议你在做部署优化时，先用 profiler 看看算子的「计算-访存比」。如果这个比值小于 10，基本就是融合的候选对象。

1.3 算子融合的收益分析

收益分析，我们得从两个维度来看：访存优化和计算密度提升。这两个是硬件加速的命门。

1.3.1 访存优化

访存优化是算子融合最直接的收益。我举个例子：

假设一个 Conv 输出 4MB 数据，BN 读这 4MB 再写 4MB，ReLU 再读再写。分开算的话，总访存量是 4 + 4 + 4 + 4 = 16MB（读两次、写两次）。融合后，Conv 算完直接进 BN 和 ReLU，只读一次输入、写一次输出，访存量降到 4 + 4 = 8MB。省了一半。

执行方式 读操作 写操作 总访存量
分开执行 Conv读输入 + BN读Conv输出 + ReLU读BN输出 Conv写输出 + BN写输出 + ReLU写输出 16MB
融合执行 融合算子读输入 融合算子写输出 8MB

你想想看，在带宽受限的硬件上，访存量减半意味着什么？意味着同样的带宽可以喂给计算单元更多数据，吞吐量自然就上去了。

1.3.2 计算密度提升

计算密度，指的是每个访存操作对应的计算量。融合后，中间数据不用写回显存，直接在寄存器或共享内存里流转，计算单元可以持续工作。

我记得有一次优化一个 MobileNet 模型，里面有很多 Depthwise Conv + Pointwise Conv 的组合。分开算的时候，计算密度只有 2-3 FLOPs/Byte，融合后提升到 15+ FLOPs/Byte。硬件利用率从 20% 飙到 70%。

关键指标：计算密度 = 总计算量 (FLOPs) / 总访存量 (Bytes)。融合的目标就是把这个比值尽可能做大。

1.3.3 其他隐性收益

  • 减少 kernel launch 次数：融合后一次启动搞定，CPU 调度开销降低。
  • 减少中间显存占用：不需要为每个中间结果分配独立 buffer，显存压力减小。
  • 利于算子库优化：融合后的复合算子可以针对特定硬件手写汇编优化，比通用库更高效。

注意：算子融合不是万能的。如果融合后的算子计算量过大，导致寄存器溢出或共享内存不足，反而会降低性能。我曾经在融合一个超大 Conv 时踩过这个坑，性能不升反降。所以融合前一定要评估硬件资源上限。

1.4 算子融合的知识体系

下面这张图是我梳理的算子融合核心知识框架，涵盖了融合类型、收益维度、硬件影响和实现方法。你可以把它当作本章的思维导图。

算子融合知识体系 算子融合 融合类型 垂直融合（连续算子） 水平融合（并行算子） 跨层融合（ResNet残差） 收益维度 访存优化（减半） 计算密度提升 减少kernel启动开销 硬件影响 寄存器/共享内存利用 计算单元利用率提升 减少显存带宽压力 实现方法 手动融合（手写CUDA） 编译器自动融合 TVM/TensorRT 核心思路：通过减少访存次数、提升计算密度，让硬件计算单元满负荷工作 融合前务必评估硬件资源（寄存器、共享内存、带宽）

1.5 一个简单的融合示例

光说不练假把式。我们看一个 Conv + ReLU 融合的伪代码，感受一下差异。

// 分开执行
__global__ void conv_kernel(...) { /* 计算卷积，写结果到global memory */ }
__global__ void relu_kernel(float* data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx] = fmaxf(0.0f, data[idx]);
}
// 启动两次kernel
conv_kernel<<<grid, block>>>(input, output);
relu_kernel<<<grid, block>>>(output);

// 融合执行
__global__ void fused_conv_relu_kernel(...) {
    // 计算卷积，结果直接进ReLU，不写回global memory
    float conv_out = compute_conv(...);
    float relu_out = fmaxf(0.0f, conv_out);
    // 只写一次最终结果
    output[idx] = relu_out;
}
// 只启动一次kernel
fused_conv_relu_kernel<<<grid, block>>>(input, output);

你看，融合后的代码更紧凑，数据在寄存器里流转，不用经过显存。这就是性能提升的根源。

我的建议：刚开始做算子融合时，先从 Conv + BN + ReLU 这种经典组合入手。这类融合在 TensorRT、TVM 等框架里已经高度优化，你可以先跑通再自己手写，对比性能差异。

好了，算子融合的基本概念和收益就讲到这里。记住三个关键词：访存优化、计算密度、kernel 启动开销。后面的章节我们会深入具体的融合策略和硬件适配技巧。

← 上一章 📖 返回目录 下一章 →