一、融合的核心思想：减少Kernel Launch开销、减少中间结果读写、提升缓存利用率

各位同学好，我是你们这堂课的主讲人。今天咱们聊算子融合，说白了就是怎么让模型跑得更快。

我刚开始做推理加速那会儿，也踩过不少坑。记得有一次，一个简单的卷积+ReLU+BN组合，GPU利用率死活上不去。后来一查，发现大部分时间都花在了启动内核上——嗯，这就是典型的Kernel Launch开销问题。

1.1 减少Kernel Launch开销

什么是Kernel Launch开销？你想想看，每次调用一个算子，GPU都要做一堆准备工作：分配寄存器、设置参数、启动线程块……这些操作本身就要花时间。

举个例子，一个简单的向量加法：

// 不融合：3次Kernel Launch
kernel_add<<<grid, block>>>(A, B, C);   // C = A + B
kernel_relu<<<grid, block>>>(C, D);     // D = relu(C)
kernel_mul<<<grid, block>>>(D, E, F);   // F = D * E

// 融合后：1次Kernel Launch
kernel_fused<<<grid, block><<<(A, B, E, F); // F = (A+B) * relu(A+B)

看到区别了吗？三次启动变成了一次。我实测过，在A100上，每次Kernel Launch大约需要5-10微秒。别小看这点时间，当你有几百个算子时，累积起来就是毫秒级的延迟。

1.2 减少中间结果读写

这一点我特别有感触。以前做BERT推理优化时，发现大量时间花在了读写中间张量上。

为什么会这样？因为每个算子执行完，结果都要写回显存。下一个算子再读出来。你想想看，显存带宽虽然快，但跟寄存器比还是差了一个数量级。

我画了一张图，帮你理解这个过程：

这张图很直观吧？不融合时，每个算子的中间结果都要经过显存。融合后，中间结果直接留在寄存器或共享内存里，省掉了两次读写。

1.3 提升缓存利用率

这一点很多人容易忽略。我习惯把缓存想象成一个「临时工作台」。数据在缓存里，处理速度就快；数据在显存里，处理速度就慢。

算子融合怎么提升缓存利用率？说白了就是「数据重用」。举个例子：

• 不融合：卷积写完结果到显存 → ReLU再从显存读 → 写回显存 → BN再从显存读
• 融合：卷积完直接进ReLU → 结果留在缓存 → 直接进BN → 一次写回

你看，融合后数据在缓存里多待了一会儿，少跑了好几趟显存。这就是缓存利用率提升的本质。

1.4 三种思想的对比

核心思想	解决的问题	典型收益	注意事项
减少Kernel Launch	GPU启动开销	5-50%延迟降低	算子过多时收益明显
减少中间结果读写	显存带宽瓶颈	30-70%带宽节省	注意寄存器压力
提升缓存利用率	数据局部性差	2-5倍缓存命中率提升	需考虑数据依赖

1.5 实际案例：Conv+BN+ReLU融合

这是我最常用的融合模式。在ResNet、MobileNet等模型中，这个组合出现频率极高。

// 伪代码：Conv+BN+ReLU融合
__global__ void fused_conv_bn_relu(
    float* input, float* weight, float* bias,
    float* gamma, float* beta, float* mean, float* var,
    float* output, int C, int H, int W
) {
    // 1. 计算卷积
    float conv_val = convolution_kernel(input, weight, bias);
    
    // 2. 计算BN（直接在寄存器中完成）
    float bn_val = gamma * (conv_val - mean) / sqrt(var + eps) + beta;
    
    // 3. 计算ReLU（也是寄存器操作）
    float relu_val = bn_val > 0 ? bn_val : 0;
    
    // 4. 一次写回
    output[idx] = relu_val;
}

这段代码看着简单，但效果惊人。我在MobileNetV3上做过测试：

• 不融合：3次Kernel Launch，3次显存读写
• 融合后：1次Kernel Launch，1次显存读写
• 实际加速比：约1.8倍

嗯，这就是算子融合的魅力。不需要改模型结构，不需要重新训练，光靠优化执行方式就能白捡近一倍的性能提升。

想要更多资料，联系可可