一、算子融合的核心动机与收益分析
大家好，我是你们的老朋友。今天咱们聊聊算子融合——这个在CUDA优化里绕不开的话题。
说实话，我刚接触GPU编程那会儿，总觉得把每个算子都写成独立的kernel，然后一个一个调，挺清晰的。直到有一次，我做一个简单的卷积+激活+池化流水线，发现性能惨不忍睹。一分析，好家伙，大部分时间都花在了数据搬运和kernel启动上，真正的计算反而没占多少。嗯，从那以后，我就开始认真研究算子融合了。
1.1 为什么需要算子融合？
先问大家一个问题：你写一个神经网络的前向推理，比如 Conv2D → ReLU → BatchNorm → Pooling，你会怎么写？
最直接的做法：四个kernel，依次执行。每个kernel从显存读数据，算完写回去，下一个kernel再读……
你想想看，这中间有多少次全局内存的读写？每次读写都是几百微秒级别的延迟。更别说每个kernel启动还有几十微秒的开销。小模型还好，大模型一跑，这些开销能占到总时间的30%甚至更多。
我个人习惯把这种模式叫做「碎片化计算」——计算本身不慢，但数据在算子之间来回倒腾，把性能拖垮了。
算子融合的核心思想，说白了就是：把多个连续的操作合并成一个kernel执行。这样中间结果不用写回显存，直接在寄存器或共享内存里流转，省掉大量的访存开销。
核心动机总结：
- 减少全局内存访问：中间结果驻留在片上，避免反复读写DRAM
- 降低kernel启动开销：多个小kernel合并成一个大kernel，启动次数从N次降到1次
- 提升并行度与计算密度：融合后的kernel有更多计算可做，能更好地隐藏访存延迟
1.2 访存减少：最直接的收益
咱们来算一笔账。假设一个简单的融合场景：A = B + C 然后 D = A * 2。
不融合的话：
- Kernel 1：读B、C，写A → 3次全局内存访问
- Kernel 2：读A，写D → 2次全局内存访问
- 总共：5次全局内存访问
融合之后：
- Kernel 1：读B、C，直接计算D并写出 → 3次全局内存访问
- 节省了2次访问（A的读写）
看起来不多？但你要知道，在深度学习模型里，这种模式比比皆是。一个ResNet-50有上百个这样的操作组合。我在项目中遇到过，仅仅把Conv+BN+ReLU融合，就能让推理速度提升30%以上。
为什么会这样？因为全局内存的带宽是有限的。现代GPU的HBM带宽虽然高（比如A100有2TB/s），但计算单元更猛。你想想看，一个SM每秒能执行多少TFLOPS的浮点运算？如果数据供应跟不上，计算单元就只能干等。融合之后，数据在片上流转，带宽压力骤减。
避坑指南： 我曾经犯过一个错误——把所有能融合的算子都强行融合到一个kernel里。结果寄存器压力过大，导致每个block能启动的线程数大幅减少，反而性能下降了。融合不是越多越好，要平衡计算密度和资源占用。
1.3 内核启动开销降低：被忽视的隐形成本
很多初学者容易忽略kernel启动的开销。你以为调用一个kernel就是几行代码的事？实际上，GPU driver要做一堆工作：参数校验、命令缓冲区构建、上下文切换……每次启动大概需要5-20微秒。
单个kernel看起来不多，但如果你有100个算子呢？那就是0.5-2毫秒的纯开销。对于推理场景，尤其是实时性要求高的应用（比如自动驾驶、语音交互），这个开销是不能接受的。
我举个例子：
// 不融合：4次kernel启动
conv2d_kernel<<<grid, block>>>(input, weight, temp1);
relu_kernel<<<grid, block>>>(temp1, temp2);
batchnorm_kernel<<<grid, block>>>(temp2, temp3);
pooling_kernel<<<grid, block>>>(temp3, output);
// 融合后：1次kernel启动
fused_conv_relu_bn_pool_kernel<<<grid, block>>>(input, weight, output);
融合后，启动开销从4次降为1次，节省了75%。而且，融合后的kernel内部可以做更精细的调度——比如在卷积计算完成后，立即在寄存器里做ReLU，然后直接做BN，最后做pooling。数据根本不用离开SM。
1.4 并行度提升：让GPU吃饱
还有一个容易被忽视的点：融合能提升并行度。
你想想看，一个单独的ReLU kernel，计算量极小，每个线程就做一次比较和赋值。这种kernel的算术强度（计算量/访存量）很低，GPU的计算单元根本吃不饱。大部分时间都在等数据。
融合之后，一个线程可能要执行：卷积的乘累加 → ReLU的比较 → BN的乘加 → Pooling的取最大值。计算量上去了，算术强度提高了，GPU的SM才能满负荷运转。
我习惯用「计算密度」这个词来衡量。一个融合后的kernel，计算密度往往比单个小kernel高几倍甚至一个数量级。高计算密度意味着更好的延迟隐藏——当一部分线程在等数据时，其他线程可以继续计算，SM的利用率自然就上去了。
收益对比表（以Conv+ReLU+BN融合为例）：
| 指标 | 不融合 | 融合后 | 提升幅度 |
|---|---|---|---|
| 全局内存访问次数 | 6次（读3次+写3次） | 2次（读1次+写1次） | 66.7% |
| kernel启动次数 | 3次 | 1次 | 66.7% |
| 算术强度（FLOP/Byte） | ~2 | ~8 | 300% |
| 端到端延迟 | 100% | ~60% | 40% |
1.5 融合的代价与权衡
当然，融合不是银弹。我前面也提到了，过度融合会导致寄存器压力过大。另外，融合后的kernel代码复杂度会上升，调试难度也更大。
我曾经在一个项目里，把7个算子融合到一起，结果编译时间从几秒变成了几分钟，而且生成的SASS代码体积巨大，指令缓存都装不下了。最后不得不拆成两个融合kernel，性能反而更好。
所以，我的建议是：先分析，再融合。用NVIDIA Nsight Compute或者nvprof看看每个算子的访存和计算占比，找到真正的瓶颈。一般来说，访存密集型的连续算子（比如逐元素操作）是最值得融合的。计算密集型的算子（比如大卷积核的Conv）融合收益相对较小。
注意事项： 融合时要特别留意数据依赖关系。如果两个算子之间有数据依赖，但依赖关系复杂（比如需要全局同步），强行融合可能导致错误或性能下降。这种情况下，可以考虑使用CUDA Graph或者stream来优化，而不是硬融合。
1.6 知识体系总览
为了让大家更直观地理解算子融合的核心逻辑，我画了一张图：
这张图很直观地展示了融合前后的差异。左边是碎片化的多kernel执行，每次都要读写全局内存；右边是融合后的一体化执行，中间结果在寄存器里流转，只有首尾两次全局访问。
好了，关于算子融合的核心动机与收益，今天就聊到这里。记住一句话：融合的本质，是用片上带宽换全局带宽，用计算密度换启动开销。理解了这个，你就能在实战中做出更好的优化决策。
想要更多资料，联系可可