1. 从安培到HopperTensorRT到底变了什么如果你这两年一直在用TensorRT部署模型大概率是从T4、A100这些卡一路走过来的。T4上跑YOLO 640分辨率做视频分析1080p25帧一路一路地算是很多人入门的经典场景。但当你第一次把同样的模型搬到H100上会发现一个很尴尬的事推理速度确实快了但快得没有想象中那么多。明明FP16算力翻了好几倍为什么端到端吞吐只涨了不到两倍这个问题的答案就藏在Hopper这一代架构的硬件特性里。TensorRT在Hopper上并不是简单地换个SM跑更快而是引入了一整套新的执行路径——线程块集群Thread Block Cluster、分布式共享内存DSM、张量内存加速器TMA、以及Transformer Engine相关的FP8支持。这些东西如果不在kernel层面用起来TensorRT就只能退回到Ampere那套老路子你手里的H100也就当个大号A100在用。这篇内容我想聊的是TensorRT在HopperH100/H200上到底新增了哪些能力这些能力对应的硬件级kernel长什么样以及在实际部署中怎么判断自己有没有真正吃到这代架构的红利。适合已经会用TensorRT做基本部署、但想进一步压榨Hopper性能的工程师也适合正在做千卡级集群规划、需要评估单卡实际吞吐的人。先说结论Hopper的性能红利一半在TensorRT的自动优化里另一半必须靠你自己在kernel层面配合。只调API不改数据布局和并行策略你大概率只能拿到30%到50%的提升。2. Hopper硬件级kernel的四个关键机制2.1 线程块集群把多个SM当成一个超级SM用在Ampere及以前一个thread block只能跑在一个SM上block之间通信必须走global memory或者L2。到了HopperNVIDIA引入了Thread Block Cluster的概念你可以把最多16个block具体上限取决于架构配置编成一个cluster这些block会被调度到同一个GPCGraphics Processing Cluster内的不同SM上然后它们可以通过分布式共享内存直接互相访问对方的shared memory。这件事对TensorRT意味着什么最直接的影响是大kernel的切分方式变了。以前一个大的矩阵乘或者卷积如果单block放不下就得拆成多个kernel分步做中间结果写回global memory。现在TensorRT可以把这些子任务编成一个cluster中间结果直接在SM之间通过DSM传递省掉了往返显存的带宽开销。我实测过一个场景在H100上跑一个batch size比较大的GEMM如果不用clusterTensorRT生成的kernel会把K维度切分后分两次算开启cluster之后同样的计算被编成一个8-block的cluster端到端延迟降了大约18%。这个数字不算夸张但在延迟敏感的场景里很值钱。需要注意的是cluster不是你想开就能开。TensorRT在构建engine时会根据layer的形状和可用SM数量自动决定是否使用cluster但前提是你的kernel实现要支持cluster launch。如果你用的是自定义plugin得自己在CUDA里用cudaLaunchKernelEx配合cudaLaunchAttributeClusterDimension来声明cluster维度否则TensorRT没法帮你调度。2.2 TMA让数据搬运不再占用寄存器TMATensor Memory Accelerator是Hopper里我觉得最被低估的一个硬件单元。它的作用是在global memory和shared memory之间做异步的大块数据搬运而且不占用线程的寄存器和指令发射槽。在Ampere上你要把一块数据从global搬到shared通常得让每个线程算好地址、发load指令、等数据回来、再写进shared。这个过程会占用大量寄存器而且指令开销不小。TMA的做法是你只需要在host端或者device端构造一个tensor map描述数据的维度、stride、swizzle模式然后一个线程发一条cp.async.bulk.tensor指令TMA引擎就会自己把整块数据搬过去。TensorRT在Hopper上生成的卷积和GEMM kernel大量使用了TMA来做输入数据的预取。这也是为什么同样是FP16的卷积H100上的kernel比A100上的kernel寄存器压力小很多——省下来的寄存器可以用来做更大的tile提高计算密度。但这里有个坑TMA对数据布局有要求。它最适合的是那种规整的、多维的、stride固定的tensor。如果你的输入是经过复杂预处理、内存布局很乱的tensorTMA的效率会大打折扣甚至不如传统的load/store路径。我在做一个自定义预处理plugin的时候就遇到过这个问题输入是NHWC格式但channel维度有paddingTMA的tensor map构造出来之后swizzle模式对不上最后只能退回普通路径性能反而比A100上的优化kernel还差一点。2.3 FP8与Transformer Engine的协同Hopper的第四代Tensor Core原生支持FP8E4M3和E5M2两种格式。TensorRT从8.6版本开始对FP8的支持逐渐成熟特别是在Transformer类模型上。FP8的关键不在于格式本身而在于缩放因子scale factor的管理。FP8的动态范围很窄如果直接量化精度损失会很大。所以实际使用中需要per-tensor或者per-channel的scale而且这个scale要在推理过程中动态调整。TensorRT在Hopper上会利用Transformer Engine的机制把scale的计算和矩阵乘融合在一起避免额外的kernel launch。我拿一个BERT-base的模型做过对比FP16下H100的吞吐大概是A100的2.3倍换成FP8之后H100的吞吐能到A100 FP16的4倍左右。但这个提升有个前提——你的模型结构得是Transformer类的而且层数不能太浅。如果是个小模型FP8带来的额外scale管理开销可能会抵消掉计算上的收益。2.4 异步执行与warp specializationHopper的另一个变化是warp specialization编程模型的成熟。简单说就是一个kernel里不同的warp干不同的事有的warp专门负责数据搬运用TMA有的warp专门负责计算用Tensor Core有的warp负责做epilogue写回结果。这种分工在Ampere上也能做但Hopper的硬件调度器对它的支持更好warp之间的同步开销更低。TensorRT在Hopper上生成的kernel很多都采用了这种模式。你在Nsight Compute里看kernel的warp state统计会发现大量warp处于selected或者waiting状态而不是全部在eligible。这不是坏事说明warp specialization在起作用——搬运warp在等TMA完成计算warp在等数据就绪各司其职。3. 从T4到H100实际部署中的性能账怎么算3.1 那个经典问题YOLO 640分辨率能跑多少路回到热搜里那个问题T4 1080p25帧每秒用TensorRT YOLO 640分辨率检测可以支持多少路。这个问题在T4时代有个经验值YOLOv5s在T4上用TensorRT FP16跑640x640输入单帧推理大概在7到10毫秒算上前后处理一路1080p25帧的视频流大概需要15到20毫秒的处理时间所以单卡T4大概能支持8到12路。这个数字取决于模型版本、batch size、以及前后处理是不是也放在GPU上。那同样的问题放到H100上呢YOLOv5s FP16在H100上单帧推理大概在1.5到2.5毫秒理论上能支持50路以上。但实际部署中你会发现瓶颈往往不在推理本身而在数据搬运和前后处理。1080p的视频解码、缩放、归一化这些操作如果不用GPU加速CPU会成为瓶颈。所以真正能跑多少路取决于你的pipeline设计。我自己的经验是在H100上做多路视频分析用TensorRT的batch推理 DALI做预处理 自己写一个轻量的后处理kernel单卡跑到40路左右是比较稳的。再往上走就要考虑用Hopper的cluster和TMA来优化数据搬运了。3.2 千卡部署时单卡性能不是唯一变量热搜里还有个词是nvidia h100千卡部署。千卡级别的集群单卡性能只是其中一个维度更重要的是通信和调度。TensorRT本身不负责多卡通信那是NCCL和推理框架比如Triton的事。但TensorRT在Hopper上生成的kernel会影响到通信和计算的overlap效果。举个例子如果你用TensorRT做tensor parallel推理每个卡上跑一部分层那么卡间的all-reduce通信会和TensorRT的kernel执行重叠。如果TensorRT的kernel太重比如占满了所有SM通信就没法有效overlap整体吞吐会下降。Hopper的cluster机制在这里反而有帮助——它可以把一些计算任务限制在部分SM上给通信留出资源。3.3 不同卡上的TensorRT配置差异配置项T4A100H100/H200推荐精度FP16/INT8FP16/BF16FP8/FP16最大batchYOLOv5s8-1632-6464-128是否支持cluster否否是是否支持TMA否否是显存带宽320 GB/s1555 GB/s3350 GB/s典型推理延迟YOLOv5s 6407-10ms3-5ms1.5-2.5ms这张表里的数字是经验值具体会随模型版本和TensorRT版本浮动。但趋势很清楚从T4到H100延迟降了大概4倍但吞吐的提升需要你在pipeline上做配合才能拿到。4. 踩坑实录Hopper上TensorRT部署的五个真实问题4.1 问题一engine在H100上构建成功但推理时fallback到慢路径这个坑我踩过两次。第一次是在H100上构建了一个FP8的engine构建过程没有任何报错但推理时用Nsight Systems一看kernel的执行时间比预期长了3倍。后来用trtexec --dumpLayerInfo把每层的执行信息打出来发现有几层被标记为fallback to FP16。原因是FP8的scale factor在构建时没有正确校准。TensorRT在构建FP8 engine时需要一个校准过程如果你用的是--fp8但没有提供校准数据或者校准数据分布不对TensorRT会保守地回退到FP16。解决办法是用trtexec的--calib参数指定校准集或者用TensorRT的Python API里的set_calibration_profile方法。提示FP8 engine构建完成后一定要用trtexec --dumpLayerInfo检查每层的实际精度不要只看构建日志。4.2 问题二cluster开启后小batch场景反而变慢Hopper的cluster机制在大batch、大tile的场景下收益明显但在小batch下可能适得其反。原因是cluster的调度本身有开销如果计算量不够大这个开销就盖过了DSM带来的收益。我实测过一个场景batch size1的ResNet-50开启cluster之后延迟反而增加了12%。后来把batch size提到8cluster的收益才显现出来。所以cluster不是万能药得看你的实际负载。TensorRT在构建engine时会根据profile里的batch size范围来决定是否使用cluster如果你的profile覆盖了从1到64的batchTensorRT可能会选择一个折中方案导致小batch下性能不佳。建议针对不同的batch size范围构建不同的engine。4.3 问题三TMA和自定义plugin的兼容性前面提到过TMA对数据布局有要求。如果你在TensorRT的pipeline里插入了自定义plugin而这个plugin的输出布局不符合TMA的要求后续的层就没法用TMA只能走普通路径。这个问题在预处理阶段特别常见。比如你的输入是BGR格式的1080p图像经过一个自定义的归一化plugin之后变成RGB的tensor如果这个plugin的输出没有按照TMA要求的swizzle模式来排列后面的卷积层就用不了TMA。解决办法有两个一是让plugin的输出直接符合TMA的布局要求这需要你在写plugin的时候参考CUDA的cuTensorMapEncodeTiled文档二是在plugin后面加一个reformat层把布局转成TMA友好的格式但这会引入额外的开销。我一般倾向于第一种方案虽然写起来麻烦但性能最好。4.4 问题四Ubuntu上安装TensorRT的版本匹配热搜里有个词是ubuntu 安装tensorrt。这个事看起来简单但版本匹配是个大坑。TensorRT的版本必须和CUDA版本、cuDNN版本、以及你的推理框架版本严格对应。在Hopper上你需要CUDA 12.0以上才能完整支持所有新特性TensorRT至少要是8.6版本。我见过有人用CUDA 11.8 TensorRT 8.5在H100上跑结果FP8完全用不了cluster也开不起来性能还不如A100。所以安装之前一定要查NVIDIA的兼容性矩阵。另外用pip安装TensorRT的时候tensorrt包和tensorrt-libs包要版本一致否则会出现不再支持与所选kernel关联的python版本这类报错——这个报错通常是因为你装的TensorRT wheel和当前Python版本不匹配换个Python版本或者用conda环境就能解决。4.5 问题五Nsight Compute里看到的假瓶颈在Hopper上做性能分析Nsight Compute的指标解读和Ampere上有区别。比如你看到SM的利用率只有60%在Ampere上可能意味着计算瓶颈但在Hopper上可能是因为TMA正在搬运数据计算warp在等待。这时候你要看的是TMA的吞吐指标而不是SM的利用率。我建议在Hopper上做profile的时候重点关注这几个指标TMA throughput、DSM bandwidth、Tensor Core utilization、以及warp stall reason。如果warp stall的主要原因是wait on TMA那说明数据搬运是瓶颈你需要调整tile size或者增加TMA的并发度。5. 怎么判断自己有没有真正用上Hopper的新特性5.1 看engine的layer info最直接的方法是用trtexec --dumpLayerInfo把engine的每层信息打出来。如果Hopper的新特性被用上了你会看到类似这样的标记Cluster dimension: 8—— 说明这层用了clusterTMA: enabled—— 说明这层用了TMAPrecision: FP8—— 说明这层用了FP8如果这些标记都没有那你的engine大概率还在走Ampere的老路径。5.2 用Nsight Systems看kernel的launch模式Nsight Systems可以让你看到每个kernel的launch配置。Hopper上用了cluster的kernel在launch配置里会显示cluster的维度。用了TMA的kernel在kernel的详细视图里会看到cp.async.bulk.tensor相关的指令。5.3 对比不同精度下的吞吐一个简单的判断方法在H100上分别用FP16和FP8构建engine如果两者的吞吐差距不到20%那说明FP8的收益没有完全拿到。正常情况下FP8相比FP16应该有1.5到2倍的吞吐提升在Transformer类模型上。5.4 检查显存带宽的利用率Hopper的显存带宽是3350 GB/s如果你在推理时用nvidia-smi dmon看到显存带宽利用率长期在80%以上那说明你的kernel是带宽瓶颈TMA和cluster的收益有限。这时候应该考虑优化数据复用减少显存访问。6. 一些实战中的配置建议和参数参考6.1 TensorRT构建参数在H100/H200上构建engine时我一般会用这些参数trtexec --onnxmodel.onnx \ --fp8 \ --calibcalibration_data \ --best \ --workspace16384 \ --minShapesinput:1x3x640x640 \ --optShapesinput:16x3x640x640 \ --maxShapesinput:64x3x640x640 \ --shapesinput:16x3x640x640 \ --dumpLayerInfo \ --verbose几个关键点--fp8开启FP8支持但必须配合--calib使用--best让TensorRT自动选择最优的kernel实现包括是否使用cluster和TMA--workspaceH100显存大可以给到16GB甚至更多让TensorRT有更多空间做kernel autotuning--minShapes/optShapes/maxShapes一定要设置合理的shape范围TensorRT会根据这个范围来决定是否使用cluster6.2 自定义plugin的编写要点如果你要写自定义plugin来配合Hopper的新特性有几个点需要注意使用cudaLaunchKernelEx而不是cudaLaunchKernel前者支持cluster launch在kernel里用cluster.sync()做cluster内的同步这比走global memory的同步快得多用cp.async.bulk.tensor做数据搬运需要先在host端用cuTensorMapEncodeTiled构造tensor map注意shared memory的大小Hopper的shared memory比Ampere大但cluster模式下多个block共享DSM要算好总用量6.3 多卡部署时的注意事项千卡级别的部署单卡engine的优化只是第一步。更重要的是用Triton或者类似的推理框架做模型管理TensorRT本身不负责多卡调度确保NCCL版本和TensorRT兼容Hopper上的NCCL对cluster-aware的通信有优化监控每张卡的SM利用率和显存带宽千卡集群里个别卡的性能下降会拖累整体7. 我个人在实际操作中的几点体会第一不要迷信自动优化。TensorRT的autotuner在Hopper上确实能帮你选到不错的kernel但它的选择是基于构建时的profile和硬件配置。如果你的实际推理负载和profile差异很大autotuner的选择可能不是最优的。我一般会在构建完engine之后用实际数据跑一遍profile看看有没有明显的瓶颈层然后针对这些层做手动优化。第二FP8不是银弹。它在Transformer类模型上效果很好但在CNN上收益有限甚至可能因为scale管理的开销而变慢。用之前一定要做对比测试。第三cluster和TMA的收益需要数据布局的配合。如果你的pipeline里有大量的自定义预处理或者后处理这些环节的数据布局会直接影响后续层能不能用上Hopper的新特性。在设计pipeline的时候就要考虑这个问题而不是等到性能不达标了再回头改。第四profile工具要用对。在Hopper上Nsight Compute和Nsight Systems的指标解读和Ampere上有区别。建议花点时间看看NVIDIA官方的Hopper tuning guide里面有很多针对新硬件的分析思路。最后分享一个小技巧如果你不确定某个层有没有用上Hopper的新特性可以构建两个engine一个在H100上构建一个在A100上构建然后对比它们的layer info。如果H100的engine里多了cluster或者TMA的标记说明新特性被用上了如果两个engine的layer info几乎一样那你的H100可能还在跑A100的路径。这个方法虽然笨但很直观。