1. 先把对象模型捋顺:Builder、Network、Engine、Context 到底谁管谁
“三、深入学习TensorRT,Developer Guide篇(二)”这个标题摆在这儿,不少同行第一反应是——Guide 我早翻过了,API 就那么几个,还能有什么新鲜货。我第一次也是这么想的,直到线上模型换成多输入加动态 batch 之后,一个setInputShape的调用顺序把我卡了整整两天,日志里没有任何报错,输出张量全是零。那次之后我才老老实实把 Developer Guide 里对象模型那几页重新读了一遍,才发现很多坑不是 API 写错,而是脑子里对“谁持有谁、谁在什么时候有效”这件事没建立清晰图景。
这篇内容我想聊的是构建期和运行期这两段路怎么走稳,适合已经能跑通一个 Hello World 级别推理、但一上生产就被动态 shape、精度对齐、显存峰值折磨过的人。TensorRT 的学习曲线陡,陡的地方不在算子怎么写,而在于它的设计哲学是“把能提前算的都在构建期算完”,所以你必须理解哪些东西是编译期常量、哪些必须在运行期决定,否则你会一直在一个错误的抽象层级上 debug。
1.1 三段式生命周期的边界在哪里
TensorRT 的整个工作流可以粗暴切成三段:定义网络、构建引擎、执行推理。第一段由IBuilder、INetworkDefinition和解析器(ONNX Parser 或手搭网络)负责,这一段你描述的是“我要算什么”,跟硬件无关;第二段IBuilder::buildSerializedNetwork把网络编译成针对特定 GPU 架构、特定 TensorRT 版本、特定精度配置的ICudaEngine,这一段是重活,可能耗时几十秒到几分钟;第三段IExecutionContext拿着引擎去跑真实数据,这一段才是每次推理都要走的路径。三段边界清晰,但依赖关系是单向的——引擎一旦构建完成就固化了,你没法在运行期再改层融合策略或者换精度。
这个单向依赖带来一个很实际的约束:Engine 是不可移植的。同一个 ONNX 在 A100 上构建的引擎拿到 T4 上加载会直接报错,跨 TensorRT 大版本(比如 8.x 到 10.x)也基本不通。我在项目里一般会把引擎构建放进部署流水线,按目标卡型和 TRT 版本做矩阵,而不是跟着模型文件一起打进镜像。听起来麻烦,但比线上启动时现场 build 要稳得多——后者会让你的服务冷启动时间从秒级变成分钟级,K8s 的健康检查直接把你打成不健康。
另外要提醒的是,IBuilder和INetworkDefinition的生命周期必须覆盖整个构建过程,很多人写成局部变量提前析构,然后在build阶段拿到野指针,现象是段错误或者构建出明显异常的引擎。Python 侧因为有引用计数相对安全,C++ 侧一定要用SampleUniquePtr这类智能指针管住。
1.2 Logger 与显存池这两个容易被忽略的角色
ILogger看起来只是个日志接口,但它是你排查构建问题的第一手信息来源。默认很多人图省事写个只打印错误的 logger,结果 INT8 校准精度崩了、某个层 fallback 到 GPU 之外、tactic 选择失败,这些信息全在 INFO 和 VERBOSE 级别里。我的做法是构建期挂 VERBOSE logger 把日志落盘,运行期降到 WARNING,一套开关切换。日志里出现Could not find any implementation for node这种字样时,基本可以确定是某个算子当前精度下没有可用 kernel,要么降精度要求,要么换等效写法。
显存池是另一个隐形角色。TensorRT 内部有 workspace 和 activation 两类显存,前者是构建期挑 tactic 用的临时空间,后者是运行期存中间激活值用的。Guide 里明确写了 activation 显存由引擎根据网络结构自动推算,但 workspace 上限需要你通过set_memory_pool_limit显式给。给太小会导致某些高性能 tactic 被跳过——注意这不是报错,是静默降级,你的吞吐可能莫名其妙掉了 30% 却找不到原因。给太大又会在多进程部署时互相挤占。我的经验值是先给到目标卡显存的 15% 到 25%,看构建日志里有没有skipping tactic due to workspace limit之类的提示,再往上加。
2. BuilderConfig:真正决定性能上限的几组旋钮
IBuilderConfig是构建阶段唯一能调的东西,所有性能相关的取舍都压在这一个对象上。Guide 里把它拆成 flags、memory pool、optimization profile、calibrator、timing cache 几大块,但文档是平铺直叙的,你不会知道哪几个参数在实际项目里权重最高。按我的排序,影响从大到小依次是:精度开关、profile 定义、timing cache、workspace 上限、tactic 相关 flag。精度和 profile 是决定性的,后面几个是锦上添花,别本末倒置。
有个常见的误区是把 config 当成“设了就有用”。比如set_flag(BuilderFlag.FP16)只是告诉构建器“允许用 FP16”,它会不会真的用,取决于每一层有没有对应的 FP16 kernel,以及构建器评估下来精度损失是否可接受。你设了 FP16 flag 结果引擎还是 FP32 跑,这种情况在含有个别不支持半精度算子的网络里非常常见,得靠日志确认。
2.1 精度开关不是一开了之
TensorRT 支持 FP32、FP16、INT8、BF16(新架构)以及 FP8(Hopper 之后)。开发阶段我建议按这个顺序推进:先用 FP32 把功能跑通、拿到基准输出,再开 FP16 对比精度和速度,最后才考虑 INT8。INT8 的门槛不在 API 而在校准,ONNX 里没有量化信息的话你得提供 calibrator,用一批有代表性的校准数据跑 min-max 或 entropy 统计。校准集的分布如果和线上差太远,INT8 的精度会掉得很夸张,我见过分类模型 top-1 直接掉 8 个点的。
层级别的精度控制是个实用技巧。Guide 里提到可以用config.set_flag的 deny 机制,或者直接在层上设置 precision。比如你的网络最后有一个对数值范围很敏感的 softmax 或 layernorm,可以把它单独钉在 FP32,其余层走 FP16。Python 里的写法大概是这样:
import tensorrt as trt logger = trt.Logger(trt.Logger.VERBOSE) builder = trt.Builder(logger) network = builder.create_network( 1 << int(trt.NetworkDefinitionCreationFlag.EXPLICIT_BATCH) ) config = builder.create_builder_config() config.set_flag(trt.BuilderFlag.FP16) # 单独锁某一层,防止精度敏感层被降精度 layer = network.get_layer(0) layer.precision = trt.DataType.FLOAT layer.set_output_type(0, trt.DataType.FLOAT) config.set_memory_pool_limit(trt.MemoryPoolType.WORKSPACE, 2 << 30)注意set_output_type这个调用,它比单设layer.precision更“硬”,因为它锁定的是这层输出的数据类型,能有效阻止后续的精度传播把它改掉。这条经验文档里写得比较散,我是被一个 argmax 前的 cast 层坑过才记住的。
2.2 Timing Cache 与 Tactic 选择
构建器给每一层挑 kernel 的过程叫 tactic selection,它会枚举候选实现、实测耗时、选最快的。这个过程的代价是构建时间长,而且每次重建都重来一遍。ITimingCache就是把这个测量结果缓存下来,下次构建同一网络结构时直接复用。多卡批量构建的场景收益非常明显,我们之前一个含 400 多个节点的检测网络,单次构建 6 分钟,开了 timing cache 后续构建压到 40 秒左右。
用法上有个细节要注意:cache 是和 GPU 架构、TensorRT 版本绑定的,跨卡复用会命中失败,set_timing_cache的第二个参数ignore_mismatch设成 false 时它会直接报错,设成 true 则忽略不匹配部分继续构建。我一般设 false,让不匹配暴露出来,免得带着过期的耗时数据挑出次优 tactic。
# 从磁盘读回上次的 timing cache with open("timing.cache", "rb") as f: blob = f.read() cache = config.create_timing_cache(blob) config.set_timing_cache(cache, ignore_mismatch=False)另一组和 tactic 相关的 flag 是BuilderFlag.TF32,在 Ampere 之后的卡上默认可能开启,如果你的模型对精度敏感又不需要那点加速,可以显式关掉。还有BuilderFlag.SPARSE_WEIGHTS,只有权重确实稀疏时打开才有意义,否则构建器会多花时间做无用检查。
2.3 Workspace 显存上限的手工设定
前面提过 workspace 影响 tactic 选择,这里补一个具体做法。TensorRT 10 里用set_memory_pool_limit(trt.MemoryPoolType.WORKSPACE, bytes),8.x 则是config.max_workspace_size。设定时建议按目标卡实际的可用显存倒推,而不是拍脑袋给个 1GB。举个例子,一张 16GB 的卡要跑两个模型实例,每个引擎的 activation 显存估算是 5GB,那 workspace 最多给到 2GB 就得留余量给 CUDA context 和输入输出 buffer。
构建日志里会明确打印每个 tactic 尝试时的 workspace 需求,如果看到大量out of workspace的跳过记录,说明你给少了。反过来如果构建日志显示实际峰值只用了你给的三分之一,那就可以往下调,把省下的显存留给并发实例。这是个需要看日志反复试的过程,别指望一次设对。
3. 动态 Shape:Optimization Profile 的实操细节
静态 shape 的时代已经过去了,现在上线的模型十个里有九个是动态 batch,图像类还有动态分辨率。TensorRT 处理动态 shape 的机制是 optimization profile,你在构建期为每个输入指定 min、opt、max 三组尺寸,构建器会以 opt 为准做 tactic 选择,同时为 min 和 max 预留足够的激活显存。这套机制设计得很聪明,但用错了会出现“显存按 max 分配、性能按 min 跑”的尴尬局面。
我第一次用 profile 的时候犯了个典型错误:min 设成 1,max 设成 64,opt 也设成 64,想着常见 batch 就是 64。结果线上大部分请求是 batch 1 到 8 的小包,推理延迟比预期高了一截。原因是构建器针对 opt=64 挑的 tactic 在小 batch 下有大量并行度浪费。后来把 opt 改成 8,大 batch 的性能略降,但小包延迟改善了近四成,整体 P99 更稳。
3.1 为什么必须给三组尺寸
min 和 max 不只是个范围声明,max 直接决定激活显存的峰值分配。引擎构建完成后,运行期无论你送多大的输入,占用的 activation 显存都是按 max 预留的,所以 max 设得过大等于白白浪费显存。而 min 的作用是保证小于 opt 的输入也有正确的 kernel 覆盖——有些 tactic 在极端小尺寸下的行为和大尺寸不同,构建器需要为 min 单独验证可行性。
三组尺寸的选取逻辑我一般这么定:max 取业务上真实出现过的最大尺寸,宁可留一点余量也别把 max 当成“理论上限”随便填;opt 取线上流量的众数区间,用监控数据统计一下 batch size 分布,取覆盖 60% 以上请求的那个值;min 取业务允许的最小合法值,很多场景就是 1。这三组数一旦定了,引擎构建时就会为它们分别验证并分配资源。
多维输入要特别注意每个维度都要给全。一个有四个维度的图像输入(N, C, H, W),profile 里得写成完整的四元组,你不能只指定 N 让 H、W 走默认,那样构建器会认为后三维是固定的。Python 侧的写法:
profile = builder.create_optimization_profile() profile.set_shape( "input", min=(1, 3, 224, 224), opt=(8, 3, 640, 640), max=(32, 3, 1280, 1280), ) config.add_optimization_profile(profile)多个输入要逐个set_shape,而且不同输入的 batch 维度必须能对上,否则运行期会报维度不匹配。我们做过一个双输入模型,图像和文本特征分开走,文本那条的 batch 维度一开始漏设了,构建出来的引擎每次都按固定 1 处理,直到压测才发现吞吐上不去。
3.2 Profile 的显存峰值与 kOPT 的选法
激活显存的估算公式 Guide 里没给完整推导,但大致的思路是每一层的输出张量大小乘以存活周期内的并发数,再取峰值。这个值构建器会自己算,你能做的是通过 profile 间接影响它。一个实用观察是:max 尺寸翻倍,激活显存大致翻倍甚至更多,因为某些层的中间结果会按 max 分配。
所以 kOPT 的选法直接关系到成本和性能的平衡。如果业务是离线批量推理,opt 贴近 max 没有问题;如果是在线服务,opt 贴近常见小 batch 更划算。我还遇到过一种情况:模型输入是变长序列,padding 到 max 长度会让激活显存爆炸,这时候更优的做法是给几个不同 profile 覆盖不同长度区间,运行时按实际长度选最合适的那个,而不是一个 profile 通吃。
多个 profile 的写法是add_optimization_profile调多次,运行时用set_optimization_profile_async或者按名字切。要注意引擎的 activation 显存是按所有 profile 里最大的那个 peak 预留的,所以 profile 数量多不等于显存翻倍,但 max 尺寸最大的那个 profile 会主导显存占用。
3.3 运行时切换 Profile 的正确姿势
动态 shape 在运行期必须显式设置输入尺寸,不设就是未定义行为。TensorRT 10 的 API 是context.set_input_shape(name, dims),返回 bool,很多人不看返回值,结果实际用的是上一次的形状或者默认值。正确姿势是先set_input_shape,再查一次context.get_tensor_shape(name)确认生效,然后才绑定地址。
context = engine.create_execution_context() ok = context.set_input_shape("input", (4, 3, 640, 640)) assert ok, "input shape rejected, check optimization profile range" # 用实际生效的形状去算输出 buffer 大小 out_shape = context.get_tensor_shape("output")一个高频坑是形状超出 profile 范围。set_input_shape返回 false 时不会有异常抛出,如果你 assert 掉以轻心,后续enqueueV3会用一个错误的形状跑,输出要么全零要么是上一次的残留数据。我现在的习惯是在服务启动时就用 min、opt、max 三组尺寸各跑一遍 warmup,确认返回值和输出 shape 都对,把问题挡在流量进来之前。
还有个小细节:切换 profile 时如果有正在执行的 stream,需要先同步。set_optimization_profile_async是带 stream 参数的异步接口,混用同步和异步接口容易出现 race。我的做法是每个 worker 线程独占一个 context,profile 在初始化阶段就定好,运行期不再切换,用空间换确定性。
4. 运行阶段:ExecutionContext、Stream 与内存搬运
构建期一次性完成,运行期就是每个请求都要走的流水线,这段的优化收益是实打实的。Guide 里运行阶段的核心对象是IExecutionContext,它持有推理过程中的状态数据,包括中间激活值的指针、当前绑定的输入输出地址、当前的形状设置。一个 engine 可以创建多个 context,context 之间互不干扰,这是做并发的关键。
这里有个设计上的约束需要理解:context 不是线程安全的。同一时刻一个 context 只能被一个线程使用,多个线程要并发就得多建几个 context。但 context 的数量不是越多越好,每个 context 都会占用独立的 activation 显存(虽然部分空间可以共享,但 Guide 里明确说默认不共享)。所以并发度和显存是一对需要权衡的量。
4.1 Context 与 Engine 的从属关系
Engine 是只读的、可共享的,创建一次可以长期持有;Context 是有状态的,每次推理都会改写内部状态。这个区别决定了资源管理策略:Engine 全局单例,Context 按线程或按请求池化。我在一个服务里做过对比,每请求新建 context 的开销不小,尤其是动态 shape 场景下每次都要重设形状,池化之后能省掉这部分。
Context 池的实现要注意一个问题:池化复用的 context 会保留上一次的形状和地址绑定,如果下一次请求的输入尺寸不同,必须重新set_input_shape,否则用的还是旧形状。我一般会在池化取出时无条件重设一次形状,即使看起来没变,这个操作本身很便宜,换来的是确定性。
C++ 侧还有createExecutionContextWithoutDeviceMemory这个接口,允许你自己管理 activation 显存,多个 context 共用一块大显存。这个用法在多模型多实例的场景里能省下可观的显存,代价是你要自己做显存分配和生命周期管理,一旦某次推理的显存需求超过你分配的大小,行为是未定义的。除非显存真的很紧张,否则我倾向于用默认的自动管理。
4.2 内存绑定的三种玩法
TensorRT 的输入输出通过“张量地址”绑定,TensorRT 10 里是set_tensor_address(name, ptr)配合enqueueV3。三种常见玩法分别是:每请求 cudaMalloc/cudaFree、预分配固定 buffer、以及用 CUDA 的 pinned memory 做零拷贝。
每请求分配是最容易写但最慢的,cudaMalloc是同步且昂贵的操作,在 QPS 稍微高一点的场景下会直接成为瓶颈。预分配是标准做法,按 max 尺寸分配一次输入输出 buffer,运行期只做数据搬运。
void* d_input = nullptr; void* d_output = nullptr; size_t in_bytes = max_batch * 3 * 1280 * 1280 * sizeof(float); size_t out_bytes = context->getMaxOutputSize("output"); cudaMalloc(&d_input, in_bytes); cudaMalloc(&d_output, out_bytes); context->setTensorAddress("input", d_input); context->setTensorAddress("output", d_output);第三种 pinned memory 用法针对的是 host 到 device 的拷贝开销。普通内存的 H2D 拷贝要经过一次中转,pinned memory 允许 DMA 直接访问,带宽能翻倍。具体做法是 host 侧用cudaHostAlloc分配,配合cudaMemcpyAsync在 stream 上异步拷贝,让拷贝和计算有机会重叠。这个优化在输入数据量大的场景(比如高分辨率图像)收益明显,小张量输入就无所谓了。
还有个容易忽略的点:getMaxOutputSize返回的是按 max profile 算的输出大小,如果你实际用的是小 batch,这个值会偏大,分配出来的 buffer 用不满。有些输出形状依赖输入的模型,更精确的做法是用get_tensor_shape拿到实际形状再算字节数。
4.3 多流并发与 CUDA Graph
单流串行执行是最简单的,但对 GPU 利用率不友好。多流并发的思路是把不同请求丢到不同 stream 上,让 kernel 有机会并行——注意 TensorRT 的 kernel 本身在单次推理里是串行依赖的,但要靠 context 级别的并发来填满 GPU。
并发的前提是每个 stream 配一个独立的 context,并且各自的输入输出 buffer 不重叠。如果多个 context 共用同一块输入 buffer,会出现数据被覆盖的 race,现象是偶发的推理结果错乱,特别难查。
CUDA Graph 是另一个层面的优化,思路是把整条推理流水线的 kernel launch 序列录制成一张图,之后一次 launch 就能重放,省掉大量的 CPU 侧 launch 开销。TensorRT 支持在固定形状下捕获 graph,Guide 里提到用enqueueV3加 stream capture 的方式。适用条件是形状完全固定——动态 shape 下每次的 kernel 序列可能不同,graph 捕获会失败或者重放出错。
cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal); context->enqueueV3(stream); cudaStreamEndCapture(stream, &graph); cudaGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0); // 之后每次推理只需 cudaGraphLaunch(graph_exec, stream);graph 捕获我踩过的坑是地址绑定。capture 的时候 kernel 参数里记录的地址是固定的,如果你后来换了输入 buffer 地址,重放的 graph 还是往旧地址写。所以要么保证 buffer 地址长期不变(预分配就能做到),要么用cudaGraphExecKernelNodeSetParams去更新节点参数,后者麻烦得多。我的建议是只有在形状和数据地址都完全稳定的服务里才上 graph,否则收益不够覆盖维护成本。
5. 常见问题与排查实录
构建期和运行期的问题表现差异很大,构建期多半有明确日志,运行期则是结果不对但没报错,后者更难。我把这几年遇到的典型问题整理成速查表,配合原来的排查思路,遇到相似的可以直接对号入座。
5.1 精度与数值异常
输出全零是最常见的一种。原因排查顺序我一般是:先确认set_input_shape返回值是不是 true,再确认输入 buffer 是不是真的拷进去了(很多人漏了cudaMemcpyAsync或者 stream 没同步),最后确认set_tensor_address绑定的地址和拷数据的地址是同一个。这三个都对了还全零,就要怀疑 profile 范围问题或者引擎本身构建异常。
输出数值抖动大、和 FP32 基准差得多,通常是精度选择的问题。先对比 FP32 引擎的输出,确认是降精度引入的。如果 FP16 下某个层明显异常,把它单独锁回 FP32。INT8 的话要检查校准集,覆盖率不够或者分布偏移都会导致 scale 估偏,表现是某些输入下结果离谱、另一些输入下正常。
还有一种隐蔽的情况是 padding 处理。动态分辨率做 resize 加 padding 的模型,如果 padding 值在预处理和后处理两边不一致,边界区域的数值会持续偏高或者偏低,看起来像精度问题,实际是数据管线问题。
| 现象 | 高概率原因 | 快速验证方式 |
|---|---|---|
| 输出全零 | 输入未正确绑地址或未拷贝 | 打印输入 buffer 前几个值 |
| 输出全零 | 形状设置失败未检查返回值 | assert set_input_shape 的返回值 |
| 数值偏差大 | 敏感层被降精度 | 逐层锁 FP32 对比 |
| 偶发错乱 | 多 context 共用 buffer | 检查每 context 独立 buffer |
| 精度正常但延迟高 | opt 尺寸偏离实际流量 | 改 opt 到众数区间重建 |
| 构建无报错但变慢 | workspace 太小触发 tactic 降级 | 看 VERBOSE 日志的跳过记录 |
5.2 显存与构建失败的定位
构建阶段显存不足的报错通常比较直白,会告诉你需求多少、可用多少。这时候不一定非要加卡,先看是不是 profile 的 max 设得太激进。把 max 减半看看能不能过,能过说明是激活显存峰值的问题,得从模型结构或者 profile 设计上想办法,比如拆分大层、减少同时存活的中间张量。
序列化失败是另一个类型。build_serialized_network返回空指针,日志里可能只有一句含糊的错误。这时候先用build_engine(非序列化接口)试试,看报错是否更具体;再看网络里有没有不支持的操作,ONNX Parser 对某些动态控制流是直接拒绝的。还有一个常见原因是 plugin 没注册,自定义层在构建时会找不到实现。
运行期 OOM 一般是绑定 buffer 超了显存,尤其是按 max 尺寸分配了输入输出还同时开了多个 context。这种时候可以用cudaMemGetInfo在分配前后各打一次,看谁吃掉了显存。另外要注意 TensorRT 的 activation 显存在 engine 加载时就已经占住了,不是按需分配的,所以监控里看到的显存占用在推理开始前就已经很高,别误判成泄漏。
5.3 版本兼容与引擎移植
引擎跨设备不可用这个前面提过,实际操作中还要注意 TensorRT 的小版本差异。同一个大版本下的 minor 更新,官方说引擎格式是兼容的,但我在实际项目里还是遇到过某个 minor 版本构建的引擎在新版本下加载后性能明显下降的情况,没有报错,就是慢。所以升级 TensorRT 之后我一般会重新构建一遍引擎做对比,别直接复用旧的。
ONNX 的 opset 也是个雷区。较新的 opset 里某些算子在新版本 TensorRT 才支持,老版本 parser 会静默转成 plugin 或者直接报不支持。构建日志里搜fallback、not supported这类关键词能提前发现。我的做法是在 CI 里跑一遍构建,把 WARNING 级别以上的日志全部收集起来,有新出现的告警就人工过一遍。
还有个经验是引擎文件要带上版本元信息一起存。文件名里塞上 TensorRT 版本、GPU 架构代号、profile 的关键参数,加载时先校验再加载。这样能避免拿错引擎导致的诡异问题,比事后从日志里反推要省事得多。
最后分享一个我个人一直在用的习惯:每换一个新模型或者新卡型,先写一个最小的验证脚本,只做三件事——用 FP32 构建、跑一组固定输入、把输出和 ONNX Runtime 或 PyTorch 的结果做数值对比。这一步跑通了,再往上加精度、加动态 shape、加并发。很多人喜欢一步到位把配置拉满,出问题之后面对一堆变量无从下手,反而更慢。把变量拆开、一步一步加,是我在 TensorRT 上最省时间的做法。