第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的融合策略,我总结成三步:
- 图解析:把网络结构解析成计算图
- 模式匹配:找到Conv+Add+ReLU这种模式
- 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推理优化时总结的融合知识体系。你看一眼,心里就有数了。
这张图里,我把融合分成了三个维度。残差结构融合解决的是“怎么把shortcut加进来”,Elementwise融合解决的是“怎么把多个小操作合并”,TensorRT实现解决的是“怎么让工具帮我们做”。
三者缺一不可。我见过很多工程师只关注TensorRT的自动融合,忽略了前两个维度。结果模型部署后,性能就是上不去。
我的建议:
先手工分析网络结构,找到所有Conv+Add+ReLU模式。然后看哪些可以自动融合,哪些需要手动写plugin。最后用profiling工具验证融合效果。
这一步做好了,INT8推理的性能基本就稳了。