元鉴
返回中文阅读流

NVIDIA Developer Blog

使用批处理模式 VC-6 和 NVIDIA Nsight 加速视觉 AI 流水线

在视觉 AI 系统中,模型吞吐量持续提高。周围的流水线阶段必须跟上步伐,包括解码、预处理和 GPU……

中文内容

已翻译official company source英文原文2026-05-26

在视觉 AI 系统中,模型吞吐量持续提升。包括解码、预处理和 GPU 调度在内的周边流水线阶段必须跟上节奏。在上一篇文章《Build High-Performance Vision AI Pipelines with NVIDIA CUDA-Accelerated VC-6》中,这被描述为数据到张量的差距——即 AI 流水线各阶段之间的性能不匹配。

SMPTE VC-6 (ST 2117-1) 编解码器通过一种分层的、基于图块的架构来解决这一差距。图像被编码为可逐步细化的质量等级(Levels of Quality,LoQs),每个等级都会增加增量细节。这使得系统能够仅选择性地检索和解码所需的分辨率、感兴趣区域或颜色平面,并可随机访问可独立解码的帧。流水线可以只检索和解码模型所需的内容。

然而,高效的单图像执行并不会自动转化为高效的扩展。随着批大小增加,瓶颈会从单图像内核效率转移到工作负载编排、启动节奏和 GPU 占用率。

本文重点介绍为扩展 VC-6 解码以支持批量推理和训练工作负载所需的架构变更。由于 NVIDIA Nsight Systems 和 NVIDIA Nsight Compute 能够帮助开发者识别系统级和内核级约束,因此它们被用于重新设计面向批量吞吐的 VC-6 CUDA 实现。结果显示,在输出质量完全相同的情况下,与此前实现相比,每张图像的解码时间最高降低约 85%;批量模式下 LoQ-0(约 4K)的解码时间低于 1 毫秒,较低 LoQ 的解码时间约为 0.2 毫秒。这显著提升了生产级视觉 AI 工作负载的流水线效率。

介绍 VC-6 批量模式实现

新的实现围绕多项架构变更构建,包括批量模式和内核级优化。

批量模式:从 N 个解码器到单一解码器

  • 执行模型重新设计:通过算法变更,使用单个解码器同时解码多张图像
  • 改进并行化:利用新的工作维度(图像),结合现有并行化维度(图块、平面),将初始 VC-6 图块层级工作转移到 GPU。小批量流水线处理

内核级优化

  • 基于 Nsight Compute 的范围解码器内核优化
  • 这些优化使内核速度提升了约 20%

以下各节将深入详细介绍对 VC-6 解码器所做的这些更改。与任何 CUDA 优化一样,计划是先使用 Nsight Systems 这类系统级分析器来识别并修复初始性能瓶颈,然后使用 Nsight Compute 来优化各个内核。

从 N 个解码器转向单个解码器

图 1 的上半部分展示了起点,具体如《Build High-Performance Vision AI Pipelines with NVIDIA CUDA-Accelerated VC-6》中所述。

A comparison between two Nsight Systems screenshots. Above, in “individual async decoders”, each decoder issues its own kernel, resulting in a quite messy overview with many decoders calling many small kernels, resulting in less than 100% lA comparison between two Nsight Systems screenshots. Above, in “individual async decoders”, each decoder issues its own kernel, resulting in a quite messy overview with many decoders calling many small kernels, resulting in less than 100% l
图 1. 在 Nsight Systems 中显示的将多个图像解码批处理到单个 VC-6 解码器中

中间几行显示了大量 CUDA API 使用情况,每一行对应一个单独的解码器实例,每个实例解码一张图像。在 All Streams 中,GPU 上许多小型并发运行的内核以蓝色显示。最上方一行显示设备利用率。浅橙色表示未达到满负载,深橙色表示满负载。在此示例中,即使分派了足够的工作,也很少显示达到满利用率。因此,所分析的算法并非最优。

这种低效可由大量小型内核的执行来解释。每次内核启动都有若干相关开销,例如调度和内核资源管理。在这种情况下,每个内核固定的开销以及每个内核较少的工作量,会导致开销与实际工作之间的比例不利。

要改变这一点,需要将范式从许多小型内核转变为少数较大型内核。

在这个案例中,NVIDIA Nsight 促使执行模型从针对 N 张图像使用 N 个解码器,重新设计为使用单个解码器一次性解码 N 张图像的批次。这种新的执行模型将固定的工作量重新分配到更少的内核中,每个内核承担更多工作。图 1 的下半部分展示了这次重新实现的效果。它仅显示了两条 CUDA API 时间线、少数几个大型内核,以及由深橙色 GPU 利用率所指示的满负载 GPU 利用率。

将更多工作转移到 GPU

在初始实现中,VC-6 图块层级结构的根层级和较窄层级的解码是在 CPU 上执行的。对于单图像解码而言,这些较窄阶段的工作量太小,不足以证明使用 GPU 执行是合理的。在批处理设计中,尽管每张图像的工作量仍然很小,但多张图像的聚合提供了足够的并行性,能够高效利用 GPU。

此外,该算法经过修改,以消除主机端用于处理可变图像尺寸的逻辑。将该逻辑嵌入 GPU 内核后,NVIDIA Nsight 显示,这既减少了同步点和提交延迟,又提高了流水线的流畅性。

图 2 和图 3 展示了在 LoQ-0 和 LoQ-2 下解码图像的利用率和 CPU 开销概览,表明 LoQ-2 存在更严重的低效问题。

Similarly to the comparison in Figure 1, Figure 2 shows two Nsight Systems screenshots, this time showing NVTX ranges in relation to GPU execution statistics like clocks, compute in flight, SMs active, warp occupancy and more.Similarly to the comparison in Figure 1, Figure 2 shows two Nsight Systems screenshots, this time showing NVTX ranges in relation to GPU execution statistics like clocks, compute in flight, SMs active, warp occupancy and more.
图 2. 比较 GPU 利用率:多解码器与批处理模式(全分辨率,LoQ-0),底行显示 CPU 开销
The same comparison as in the previous figures for a lower quality, LoQ-2. It largely shows the same effect, that even for this amount of work, the GPU and the warps are fully utilized in the new setting, as compared to the previous softwarThe same comparison as in the previous figures for a lower quality, LoQ-2. It largely shows the same effect, that even for this amount of work, the GPU and the warps are fully utilized in the new setting, as compared to the previous softwar
图 3. 比较 GPU 利用率:多解码器与批处理模式(四分之一分辨率,LoQ-2),底行显示 CPU 开销

然而,使用批处理模式 VC-6(图 2 和图 3 底部)时,即使是最小的 LoQ,在 GPU 上执行也是可行的,因为多张图像的聚合工作负载可以在 GPU 上高效计算。

小批量流水线处理

新的解码器设计将每个批次拆分为多个小批量。这些小批量会经过一条由 CPU 处理、PCIe 传输和 GPU 解码阶段组成的流水线。一个小批量中的图像会同时驻留在某个流水线阶段中,而各阶段并发运行,并相互隐藏彼此的成本。

An Nsight Systems screenshot showing GPU utilization in relation to two instances of the CUDA API and NVTX. The GPU is fully utilized at all times, whilst the CUDA APIs show NVTX ranges titled “UPLOAD”, “GPU” and “CPU”. This indicates thatAn Nsight Systems screenshot showing GPU utilization in relation to two instances of the CUDA API and NVTX. The GPU is fully utilized at all times, whilst the CUDA APIs show NVTX ranges titled “UPLOAD”, “GPU” and “CPU”. This indicates that
图 4. Nsight Systems 中的小批量流水线处理,展示了 GPU 利用率与 CUDA API 和 NVTX 实例之间的关系

图 4 展示了这种小批量流水线处理。与图 1 类似,CUDA API 调用由 UPLOAD 和 GPU 两个线程发出,主机端资源使用量极低。工作聚合显著减少了 CUDA API 调用、内存操作和同步,同时将内核启动开销在整个批次中摊销。

内核级优化

Nsight Systems 显示,初始优化缓解了 CPU 开销,而进一步改进需要进行内核优化。实现范围解码器的 terminal_decode 内核值得关注。Nsight Compute 突出了此前并非关键的微架构限制。以下算法问题被重点指出:典型的底层低效问题,例如较低的流式 SM 占用率、warp 分歧、非合并内存访问以及寄存器压力。这些洞察对于开发者随后在可能的情况下消除或最大限度减少这些算法问题至关重要。

Nsight Compute 源代码热力图和 Warp Stall Sampling(All Samples)突出显示了每一行源代码所测得的耗时。它们显示,在范围更新逻辑中的整数除法上花费了大量时间(图 5)。由于 GPU 并未针对整数除法进行优化,而准确性又不容妥协,因此这些操作无法优化。

A screenshot of Nsight Compute showing three lines of source code on the left, each associated with warp stall sampling (all samples), and (not-issued samples.) The source code lines containing integer divisions have a large green and purplA screenshot of Nsight Compute showing three lines of source code on the left, each associated with warp stall sampling (all samples), and (not-issued samples.) The source code lines containing integer divisions have a large green and purpl
图 5。范围解码器更新显示,整数除法花费了大量时间

对于解码器查找表(以共享内存上的二分查找实现),Nsight Compute 也显示出显著的短 scoreboard 停顿(图 6)。

An Nsight Compute screenshot, this time source code lines and stall reasons in relation to SASS-instructions, the latter including a scoreboard dependency graphic. The scoreboard dependencies show that some SASS machine instruction lines arAn Nsight Compute screenshot, this time source code lines and stall reasons in relation to SASS-instructions, the latter including a scoreboard dependency graphic. The scoreboard dependencies show that some SASS machine instruction lines ar
图 6. 左侧为短 scoreboard 停顿(粉色),右侧为 LDS(加载共享内存)操作的 scoreboard 依赖关系

这些停顿指向共享内存加载(图 6 中的 LDS),因为对局部数组进行动态索引否则会导致缓慢的本地内存访问。由于查找表的大小是恒定的,可以用局部变量和展开循环来替代这种方法。与二分查找相比,这种穷举搜索允许对可驻留在寄存器中的固定大小数组进行常量索引。将这两项更改同时应用于两个范围解码器,使该内核获得了约 20% 的加速。

Two side-by-side images showing a subset of the memory hierarchy diagram of Nsight Compute. This shows data the kernel consumes, for example directly from global, local and shared memory, and then also from L1/Tex Cache. Two versions are shTwo side-by-side images showing a subset of the memory hierarchy diagram of Nsight Compute. This shows data the kernel consumes, for example directly from global, local and shared memory, and then also from L1/Tex Cache. Two versions are sh
图 7. 修改前(左)和修改后(右)的内核内存使用情况,显示未使用共享内存(底部)和本地内存

图 7 显示了 Nsight Compute 的内存图表。从视觉上看,它证实了修改后既没有使用共享内存(最后一行),也没有使用本地内存(第 2 行)。

其权衡是寄存器使用量增加,从每个线程 48 个寄存器增加到 92 个寄存器。考虑到每个线程 255 个寄存器的限制以及该内核相对较小的网格维度,这里是可以接受的。由于在此阶段并不优先追求每个 SM 的高块驻留率,额外的寄存器压力不会限制总体吞吐量。

另一项优化是用 cub::DeviceSelect 函数调用替换自定义选择例程。这简化了代码,并将当前及未来硬件的维护和优化工作交由 CUB 负责。

性能扩展和更新后的结果

图 8 比较了先前实现与改进实现之间在不同批量大小下的单图像解码时间,使用 UHD-IQA 数据集(可通过 Hugging Face 上的 V-Nova 获取)在四个 LoQ(LoQ-0 约 4K、LoQ-1 约 2K、LoQ-2 约 1K、LoQ-3 约 0.5K)下进行评估。

A bar chart, comparing VC-6 before and after optimization in per-image decode time in relation to batch size from 1-256. Four LoQs are shown, corresponding to even steps between 4K down to ~0.5K. Overall, the optimized version is much fasteA bar chart, comparing VC-6 before and after optimization in per-image decode time in relation to batch size from 1-256. Four LoQs are shown, corresponding to even steps between 4K down to ~0.5K. Overall, the optimized version is much faste
图 8. VC-6 在 NVIDIA L40s(g6e.8xlarge)上的单图像解码时间比较,UHD-IQA 数据集

出现了两种不同的扩展行为:

  • 先前的实现在较小批量大小(约 1–16)之后趋于平台期。增加图像数量并未转化为进一步的单图像收益。相比之下,优化后的 CUDA 实现会随着批量大小增加而持续改进。例如,在较大批量大小下,LoQ-0(约 4K)的解码时间降至每张图像 1 ms 以下。
  • 相对改进在较低 LoQ 下更为明显。较小的单图像工作负载暴露出更多可聚合的独立工作,从而带来更好的 GPU 利用率。在较高批大小下,LoQ-2 解码达到每张图像约 0.2 毫秒,LoQ-3 达到约 0.14 毫秒。

测得的改进包括:

  • 批大小为 1 时(LoQ-0),单图像解码时间降低约 36%
  • 在批大小为 16–32 时,LoQ-2 和 LoQ-3 的单图像解码时间降低约 70–80%
  • 批量大小为 256 时,每张图像解码时间最多降低约 85%

图 9 显示了重新设计的实现在 NVIDIA H100(Hopper)和 NVIDIA B200(Blackwell)GPU 上不同批量大小下的性能。结果表明,性能提升并非特定硅芯片所独有,而是源于改进后的批处理模式。这有效地暴露了足够的并行工作量,以充分利用现代 GPU 架构。

A bar chart, comparing 3 LoQ decode times (per-image decode time in ms) in relation to batch sizes from 1-256. Left shows the result of an H100, and right of a B200 GPU. The scaling is similar, batch-size 1 taking the longest, and the restA bar chart, comparing 3 LoQ decode times (per-image decode time in ms) in relation to batch sizes from 1-256. Left shows the result of an H100, and right of a B200 GPU. The scaling is similar, batch-size 1 taking the longest, and the rest
图 9. VC-6 批处理模式下的每张图像解码时间(NVIDIA H100、NVIDIA B200、UHD-IQA 数据集)

用于视觉 AI 流水线的 VC-6

利用 VC-6 随机访问、仅帧内、LoQ 解码,以及选择性访问感兴趣区域或颜色通道的智能化、定制化解码,可使训练、推理和视频摘要工作流受益。这是未来工作的一个方向。

开始使用 VC-6 解码

扩展 VC-6 解码能力需要的不仅仅是内核调优。Nsight 性能分析揭示了启动节奏、占用率、线程分歧和内存行为方面的结构性限制。通过重新设计 CUDA 执行模型,以暴露更多独立工作并在批处理中分摊开销,新的实现将单张图像解码时间最多降低约 85%,在批处理中实现了 LoQ-0(约 4K)的亚毫秒级解码,并在较低 LoQ 下达到约 0.2 毫秒,同时输出质量保持一致。

随着视觉 AI 工作负载持续扩展,整体流水线效率由每个步骤共同决定,包括解码和预处理阶段。

要开始使用,请查看以下资源:

  • VC-6 示例:VC-6 编码和选择性解码的示例;用于通过 Hugging Face 数据集复现我们结果的基准测试套件
  • VC-6 AI Blueprint:展示视觉 AI 流水线中 VC-6 选择性解码的演示;适用于多个用例的参考集成模式
Like

标签

原文标题

Accelerating Vision AI Pipelines with Batch Mode VC-6 and NVIDIA Nsight