一、水平融合与垂直融合:两种最基本的算子合并策略

大家好,我是你们的老朋友。今天我们来聊聊算子融合里最基础、也最重要的两个概念——水平融合和垂直融合。

说实话,我刚接触GPU计算那会儿,也分不清这两者的区别。总觉得「把算子合并到一起不就完事了吗?」后来在项目里踩了坑才明白——融合的方向不同,优化的思路完全不同。

1.1 水平融合:把相同类型的算子「打包」

什么叫水平融合?说白了,就是把多个相同类型的算子合并成一个。比如你有三个element-wise的加法操作,分别对不同的张量做运算,那就可以把它们合并成一个大的element-wise操作。

核心思想:减少kernel launch次数,提高GPU利用率。

我举个例子。假设你有三个独立的ReLU激活函数,分别作用在不同的特征图上:

// 融合前:三个独立的ReLU kernel
relu_kernel<<<grid, block>>>(input1, output1, N);
relu_kernel<<<grid, block>>>(input2, output2, N);
relu_kernel<<<grid, block>>>(input3, output3, N);

// 融合后:一个kernel搞定三个ReLU
fused_relu_kernel<<<grid, block>>>(input1, input2, input3, 
                                        output1, output2, output3, N);

你想想看,每次启动kernel都有开销——把参数传到GPU、分配寄存器、调度线程块……这些开销虽然单个不大,但积少成多。我在一个NLP模型里见过,光ReLU就调用了上百次,融合后性能提升了将近30%。

我的经验:水平融合最适合那种「计算量小、调用频繁」的算子。比如激活函数、归一化、简单的数学运算。如果算子本身计算量很大,融合带来的收益反而没那么明显。

1.2 垂直融合:把流水线上的算子「串起来」

垂直融合就完全不一样了。它是把计算流水线上前后连续的算子合并成一个。比如Conv + BN + ReLU,这是CNN里最常见的组合,也是垂直融合的经典案例。

为什么要做垂直融合?核心原因只有一个——减少全局内存的读写。

你想想看,如果不融合,Conv算完要把结果写回显存,BN读出来算完再写回去,ReLU再读再写……这一来一回,带宽全浪费在数据搬运上了。融合之后,中间结果直接留在寄存器或者共享内存里,根本不用碰全局内存。

// 融合前:三步走
conv_kernel<<<grid, block>>>(input, conv_out, ...);
bn_kernel<<<grid, block>>>(conv_out, bn_out, ...);
relu_kernel<<<grid, block>>>(bn_out, output, ...);

// 融合后:一步到位
fused_conv_bn_relu_kernel<<<grid, block>>>(input, output, ...);

我记得有一次优化一个ResNet-50的推理模型,光是把Conv+BN+ReLU融合,端到端延迟就降了40%。当时团队里有人还不太信,觉得「不就是少了几次内存读写吗?」——嗯,事实证明,内存带宽才是GPU最大的瓶颈。

注意:垂直融合不是万能的。如果两个算子之间需要频繁的同步或者依赖关系复杂,强行融合反而会降低并行度。我曾经试过把两个计算量都很大的卷积融合在一起,结果因为寄存器压力太大,性能反而下降了。

1.3 两种策略的对比

维度 水平融合 垂直融合
合并对象 相同类型的算子 流水线上连续的算子
主要收益 减少kernel launch次数 减少全局内存读写
适用场景 element-wise、激活函数等轻量算子 Conv+BN+ReLU、MatMul+Add等组合
潜在风险 寄存器压力增大 并行度下降、依赖复杂
典型收益 10%-30%性能提升 20%-50%性能提升

1.4 典型应用场景

说几个我在实际项目中遇到的例子吧。

场景一:Transformer中的LayerNorm + Add + ReLU

这是典型的垂直融合场景。LayerNorm算完,紧接着做残差连接,然后过ReLU。如果不融合,中间结果要写回显存再读出来。我建议的做法是把这三个操作合并成一个kernel,中间结果全部留在寄存器里。

场景二:多个独立的卷积层

这个适合水平融合。比如在MobileNet里,depthwise卷积后面经常跟着多个pointwise卷积,它们之间没有数据依赖。这时候可以把多个pointwise卷积合并成一个大的矩阵乘法,减少kernel launch次数。

场景三:Conv + BN + ReLU + Pooling

这是垂直融合的经典案例。我曾经在优化一个图像分类模型时,把Conv+BN+ReLU+MaxPooling四个算子融合成一个。注意这里有个坑——Pooling的窗口大小和步长要跟卷积的输出尺寸匹配,否则融合后计算逻辑会变得很复杂。

避坑指南:我曾经犯过一个错误——把两个计算量都很大的卷积强行水平融合。结果因为寄存器不够用,编译器不得不做spill,性能反而下降了15%。所以,融合前一定要评估寄存器压力和共享内存的使用量。

1.5 知识体系图

下面这张图展示了水平融合和垂直融合在整个算子融合体系中的位置,以及它们各自的核心逻辑:

算子融合策略总览 水平融合 垂直融合 合并相同类型的算子 ReLU ReLU ReLU 融合后的ReLU 收益:减少kernel launch次数 典型场景:多个激活函数 多个element-wise操作 合并流水线上的连续算子 Conv BN ReLU 融合后的Conv+BN+ReLU 收益:减少全局内存读写 典型场景:Conv+BN+ReLU MatMul+Add+Activation 核心区别:水平融合减少启动开销,垂直融合减少内存带宽开销

从这张图可以看得很清楚——水平融合是「横向扩展」,把多个同类型的算子打包;垂直融合是「纵向深入」,把流水线上的算子串起来。两种策略各有各的用武之地,实际项目中往往需要结合使用。

好了,这一章的内容就到这里。水平融合和垂直融合是算子融合的基石,理解了它们,后面再讲更复杂的融合策略就容易多了。

可可精品资料库