第1章：融合的本质——减少Kernel Launch次数与减少中间结果读写

大家好，我是你们的老朋友。今天咱们聊聊算子融合最核心的东西。

很多人刚开始接触算子融合时，总觉得这是个很高深的技术。其实没那么复杂。说白了，融合的本质就两件事：减少Kernel Launch次数和减少中间结果读写。把多个算子合并成一个Kernel，就这么简单。

1.1 为什么Kernel Launch这么贵？

先说说Kernel Launch。我刚开始做GPU编程时，犯过一个低级错误——写了个循环，每次迭代都启动一个小Kernel。结果呢？性能惨不忍睹。

为什么会这样？因为每次启动Kernel，GPU都要做一堆准备工作：

• 把指令从CPU传到GPU
• 分配寄存器资源
• 设置线程块和网格
• 等待前一个Kernel完成

这些开销，少说也要几微秒。你想想看，如果你的Kernel本身只跑几微秒，那启动开销就占了半壁江山。这不划算。

1.2 中间结果读写——被忽视的性能杀手

另一个问题是中间结果。我记得有个项目，做的是图像处理流水线：先做卷积，再做池化，最后做激活。每个算子都输出一个中间张量，写到全局内存里。

你猜怎么着？这些中间结果占用了大量带宽。而且，写一次读一次，来回折腾。我算过一笔账：

操作	带宽消耗	延迟
卷积写中间结果	高	高
池化读中间结果	高	高
池化写输出	中	中
激活读输出	中	中

看到没？中间结果读写占了总带宽的大头。如果能把这三个算子融合成一个，中间结果直接留在寄存器或共享内存里，根本不用出去。

1.3 融合的两种基本模式

算子融合不是乱来的。我总结下来，主要有两种模式：

1.3.1 水平融合

把多个独立的算子合并成一个。比如两个卷积层，如果它们没有数据依赖，可以并行执行。这种融合主要减少Kernel Launch次数。

1.3.2 垂直融合

把流水线上的算子合并。比如Conv + ReLU + Pooling，它们有先后顺序。这种融合主要减少中间结果读写。

嗯，这里要注意：垂直融合的效果通常比水平融合更好。因为中间结果读写是真正的性能瓶颈。

1.4 一个简单的例子

咱们看个代码示例。假设你有两个算子：

// 算子1：加法
__global__ void add_kernel(float* a, float* b, float* c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) c[idx] = a[idx] + b[idx];
}

// 算子2：ReLU
__global__ void relu_kernel(float* c, float* d, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) d[idx] = max(0.0f, c[idx]);
}

// 启动两次Kernel
add_kernel<<<grid, block>>>(a, b, c, n);
relu_kernel<<<grid, block>>>(c, d, n);

融合之后：

// 融合Kernel：加法 + ReLU
__global__ void fused_add_relu_kernel(float* a, float* b, float* d, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        float tmp = a[idx] + b[idx];  // 中间结果在寄存器里
        d[idx] = max(0.0f, tmp);      // 直接输出
    }
}

// 只启动一次Kernel
fused_add_relu_kernel<<<grid, block>>>(a, b, d, n);

看到了吗？融合后，中间结果c根本不需要分配全局内存。直接在寄存器里算完就输出。这省了多少带宽啊。

1.5 融合的收益量化

咱们用数据说话。假设一个网络有10个算子，每个算子启动开销10微秒，中间结果读写100微秒：

• 不融合：10次启动 × 10微秒 + 9次中间读写 × 100微秒 = 100 + 900 = 1000微秒
• 完全融合：1次启动 × 10微秒 + 0次中间读写 = 10微秒

收益接近100倍！当然，这是理想情况。实际中，受限于寄存器数量和计算复杂度，不可能完全融合。但方向是明确的。

1.6 本章小结

好了，咱们总结一下：

• Kernel Launch次数：每次启动都有固定开销，融合能减少启动次数
• 中间结果读写：这是真正的性能杀手，融合能避免不必要的内存访问
• 两种模式：水平融合（减少启动）和垂直融合（减少读写），后者效果更明显
• 注意平衡：别过度融合，小心寄存器压力

我个人觉得，理解这两点，你就抓住了算子融合的命脉。后面的章节，咱们会深入具体的融合技术和实现细节。

想要更多资料，联系可可