第1章:Conv+Add+ReLU融合:残差结构融合、Elementwise融合技巧、TensorRT实现分析

1.1 残差结构融合——我踩过最深的坑

残差结构,说白了就是“抄近道”。

你想想看,ResNet为什么能堆到152层?就是因为有了这个shortcut连接。但推理的时候,这个“近道”反而成了性能杀手。

我刚开始做INT8推理优化时,遇到一个ResNet-50模型。量化后精度没问题,但延迟就是降不下来。后来一分析,发现Conv+Add+ReLU这个组合,GPU利用率只有30%多。为什么?因为kernel launch太频繁了。

嗯,这里要注意:残差结构融合的核心,就是把三个操作合并成一个kernel。

融合后的计算流程:

输入A → Conv → 中间结果 → Add(加上输入B) → ReLU → 输出

原本需要3次kernel launch,融合后只需要1次。

我个人习惯把残差融合分成两类:

  • 同维度残差融合:输入A和输入B的shape完全一致,直接相加
  • 跨维度残差融合:输入A经过Conv后通道数变化,需要1x1卷积对齐维度

同维度的好办,直接写一个融合kernel就行。跨维度的就麻烦了——我曾经在项目里遇到一个情况,1x1卷积的量化参数和主分支对不上,结果精度掉了2个点。排查了两天才发现是scale factor没对齐。

避坑指南:

我曾经因为残差分支的量化参数没同步,导致模型精度崩了。记住:残差融合时,两个分支的量化scale必须一致,否则Add操作会引入误差。

1.2 Elementwise融合技巧——不只是加法

Elementwise操作,除了Add,还有Mul、Sub、Div等。但实际推理中,Add占了80%以上。

为什么Elementwise融合这么重要?

因为Elementwise操作是“访存密集型”的。你想想看,计算量不大,但数据搬运量很大。如果不融合,GPU的算力全浪费在等数据上了。

我总结了几种常见的Elementwise融合模式:

融合模式 适用场景 性能提升
Conv+Add+ReLU 残差结构 2-3倍
Conv+Add 偏置项融合 1.5-2倍
Conv+BN+ReLU BN层融合 2-4倍
Conv+Add+BN+ReLU 复杂残差 3-5倍

这里有个技巧:Elementwise融合的顺序很重要。

我建议先做Conv+BN融合,再做Conv+Add融合。因为BN融合会改变权重,如果先做Add融合,BN的参数就不好处理了。

实战经验:

我在优化MobileNetV3时,发现它的残差结构里有个SE模块。SE模块里的Elementwise操作特别多,但每个都很小。我尝试把SE模块里的所有Elementwise操作合并成一个kernel,结果延迟降低了40%。

1.3 TensorRT实现分析——工具怎么用

TensorRT做融合,说白了就是“自动帮你写优化后的kernel”。但自动归自动,你得知道它做了什么。

TensorRT的融合策略,我总结成三步:

  1. 图解析:把网络结构解析成计算图
  2. 模式匹配:找到Conv+Add+ReLU这种模式
  3. kernel生成:生成融合后的CUDA kernel

具体到Conv+Add+ReLU,TensorRT的做法是这样的:

// TensorRT融合后的伪代码
__global__ void fused_conv_add_relu_kernel(
    float* input, float* residual, float* weight, float* bias,
    float* output, int batch, int channels, int height, int width
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < batch * channels * height * width) {
        // 1. 卷积计算
        float conv_out = convolution(input[idx], weight, bias);
        // 2. 残差相加
        float add_out = conv_out + residual[idx];
        // 3. ReLU激活
        output[idx] = max(0.0f, add_out);
    }
}

嗯,代码看着简单,但实际实现要考虑很多细节:

  • 内存对齐:输入输出地址要对齐到128字节
  • 共享内存:卷积的中间结果放共享内存,减少全局内存访问
  • 向量化加载:用float4一次加载4个float,提高带宽利用率

我记得有一次,TensorRT自动融合后,性能反而下降了。排查后发现是输入tensor的shape不是4的倍数,导致向量化加载失效。后来我手动padding了一下,性能就上来了。

注意:

TensorRT的自动融合不是万能的。如果网络结构太复杂(比如有多个分支的残差),它可能匹配不到模式。这时候就需要手动写plugin了。

1.4 知识体系总览

下面这张图,是我做INT8推理优化时总结的融合知识体系。你看一眼,心里就有数了。

Conv+Add+ReLU融合知识体系 残差结构融合 Elementwise融合 TensorRT实现 同维度残差 跨维度残差 Add/Mul/Sub 融合顺序优化 自动模式匹配 手动Plugin 核心目标:减少kernel launch + 提高访存效率 一次融合 = 3次kernel launch → 1次

这张图里,我把融合分成了三个维度。残差结构融合解决的是“怎么把shortcut加进来”,Elementwise融合解决的是“怎么把多个小操作合并”,TensorRT实现解决的是“怎么让工具帮我们做”。

三者缺一不可。我见过很多工程师只关注TensorRT的自动融合,忽略了前两个维度。结果模型部署后,性能就是上不去。

我的建议:

先手工分析网络结构,找到所有Conv+Add+ReLU模式。然后看哪些可以自动融合,哪些需要手动写plugin。最后用profiling工具验证融合效果。

这一步做好了,INT8推理的性能基本就稳了。

可可精品资料库