NVIDIA TensorRT的最佳性能实践

13.2. Optimizing TensorRT Performance

以下部分重点介绍 GPU 上的一般推理流程和一些提高性能的一般策略。这些想法适用于大多数 CUDA 程序员，但对于来自其他背景的开发人员可能并不那么明显。

13.2.1. Batching

最重要的优化是使用批处理并行计算尽可能多的结果。在 TensorRT 中，批次是可以统一处理的输入的集合。批次中的每个实例都具有相同的形状，并以完全相同的方式流经网络。因此，每个实例都可以简单地并行计算。

网络的每一层都有计算前向推理所需的一定数量的开销和同步。通过并行计算更多结果，这种开销可以更有效地得到回报。此外，许多层的性能受到输入中最小维度的限制。如果批量大小为 1 或较小，则此大小通常可能是性能限制维度。例如，具有V个输入和K个输出的完全连接层可以针对一个批次实例实现为1xV矩阵与VxK权重矩阵的矩阵乘法。如果对N个实例进行批处理，则这将变为NxV乘以VxK矩阵。向量矩阵乘法器变成矩阵矩阵乘法器，效率更高。

更大的批量大小几乎总是在 GPU 上更有效。非常大的批次，例如N 》 2^16 ，有时可能需要扩展索引计算，因此应尽可能避免。但通常，增加批量大小会提高总吞吐量。此外，当网络包含 MatrixMultiply 层或完全连接层时，如果硬件支持，由于使用了 Tensor Cores，32 的倍数的批大小往往对 FP16 和 INT8 推理具有最佳性能。

由于应用程序的组织，有时无法进行批处理推理工作。在一些常见的应用程序中，例如根据请求进行推理的服务器，可以实现机会批处理。对于每个传入的请求，等待时间T 。如果在此期间有其他请求进来，请将它们一起批处理。否则，继续进行单实例推理。这种类型的策略为每个请求增加了固定的延迟，但可以将系统的最大吞吐量提高几个数量级。

使用批处理

如果在创建网络时使用显式批处理模式，则批处理维度是张量维度的一部分，您可以通过添加优化配置文件来指定批处理大小和批处理大小的范围以优化引擎。有关更多详细信息，请参阅使用动态形状部分。

如果在创建网络时使用隐式批处理模式，则IExecutionContext：：execute （ Python 中的IExecutionContext.execute ）和IExecutionContext：：enqueue （ Python 中的IExecutionContext.execute_async ）方法采用批处理大小参数。在使用IBuilder：：setMaxBatchSize （ Python中的 Builder.max_batch_size ）构建优化网络时，还应该为构建器设置最大批量大小。当调用IExecutionContext：：execute或enqueue时，作为绑定参数传递的绑定是按张量组织的，而不是按实例组织的。换句话说，一个输入实例的数据没有组合到一个连续的内存区域中。相反，每个张量绑定都是该张量的实例数据数组。

另一个考虑因素是构建优化的网络会针对给定的最大批量大小进行优化。最终结果将针对最大批量大小进行调整，但对于任何较小的批量大小仍然可以正常工作。可以运行多个构建操作来为不同的批量大小创建多个优化引擎，然后在运行时根据实际批量大小选择要使用的引擎。

13.2.2. Streaming

一般来说，CUDA 编程流是一种组织异步工作的方式。放入流中的异步命令保证按顺序运行，但相对于其他流可能会乱序执行。特别是，两个流中的异步命令可以被调度为同时运行（受硬件限制）。

在 TensorRT 和推理的上下文中，优化的最终网络的每一层都需要在 GPU 上工作。但是，并非所有层都能够充分利用硬件的计算能力。在单独的流中安排请求允许在硬件可用时立即安排工作，而无需进行不必要的同步。即使只有一些层可以重叠，整体性能也会提高。

使用流式传输

识别独立的推理批次。

为网络创建一个引擎。

cudaStreamCreate为每个独立批次创建一个 CUDA 流，并为每个独立批次创建一个IExecutionContext 。

IExecutionContext：：enqueue从适当的IExecutionContext请求异步结果并传入适当的流来启动推理工作。

在所有工作启动后，与所有流同步以等待结果。执行上下文和流可以重用于以后的独立工作批次。

多个流 运行多个并发流通常会导致多个流同时共享计算资源的情况。这意味着与优化 TensorRT 引擎时相比，推理期间网络可用的计算资源可能更少。这种资源可用性的差异可能会导致 TensorRT 选择一个对于实际运行时条件不是最佳的内核。为了减轻这种影响，您可以在引擎创建期间限制可用计算资源的数量，使其更接近实际运行时条件。这种方法通常以延迟为代价来提高吞吐量。有关更多信息，请参阅限制计算资源。

也可以将多个主机线程与流一起使用。一种常见的模式是将传入的请求分派到等待工作线程池中。在这种情况下，工作线程池将每个都有一个执行上下文和 CUDA 流。当工作变得可用时，每个线程将在自己的流中请求工作。每个线程将与其流同步以等待结果，而不会阻塞其他工作线程。

13.2.3. CUDA Graphs

CUDA 图是一种表示内核序列（或更一般地是图）的方式，其调度方式允许由 CUDA 优化。当您的应用程序性能对将内核排入队列所花费的 CPU 时间敏感时，这可能特别有用。 TensorRT 的enqueuev2（）方法支持对不需要 CPU 交互的模型进行 CUDA 图捕获。例如：

C++

// Capture a CUDA graph instance
cudaGraph_t graph;
cudaGraphExec_t instance;
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
context->enqueueV2(buffers, stream, nullptr);
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);

// To run inferences:
cudaGraphLaunch(instance, stream);
cudaStreamSynchronize(stream);

不支持图的模型包括带有循环或条件的模型。在这种情况下， cudaStreamEndCapture（）将返回cudaErrorStreamCapture*错误，表示图捕获失败，但上下文可以继续用于没有 CUDA 图的正常推理。

捕获图时，重要的是要考虑在存在动态形状时使用的两阶段执行策略。

更新模型的内部状态以考虑输入大小的任何变化

将工作流式传输到 GPU

对于在构建时输入大小固定的模型，第一阶段不需要每次调用工作。否则，如果自上次调用以来输入大小发生了变化，则可能需要进行一些工作来更新派生属性。

第一阶段的工作不是为捕获而设计的，即使捕获成功也可能会增加模型执行时间。因此，在更改输入的形状或形状张量的值后，调用enqueueV2（）一次以在捕获图形之前刷新延迟更新。

使用 TensorRT 捕获的图特定于捕获它们的输入大小，以及执行上下文的状态。修改捕获图表的上下文将导致执行图表时未定义的行为 – 特别是，如果应用程序通过createExecutionContextWithoutDeviceMemory（）为激活提供自己的内存，则内存地址也会作为图表的一部分被捕获。绑定位置也被捕获为图表的一部分。

trtexec允许您检查构建的 TensorRT 引擎是否与 CUDA 图形捕获兼容。有关详细信息，请参阅trtexec部分。

13.2.4. Enabling Fusion

13.2.4.1. Layer Fusion

TensorRT 尝试在构建阶段在网络中执行许多不同类型的优化。在第一阶段，尽可能将层融合在一起。融合将网络转换为更简单的形式，但保持相同的整体行为。在内部，许多层实现具有在创建网络时无法直接访问的额外参数和选项。相反，融合优化步骤检测支持的操作模式，并将多个层融合到一个具有内部选项集的层中。

考虑卷积后跟 ReLU 激活的常见情况。要创建具有这些操作的网络，需要使用 addConvolution 添加卷积层，然后使用addActivation和kRELU的ActivationType添加激活层。未优化的图将包含用于卷积和激活的单独层。卷积的内部实现支持直接从卷积核一步计算输出上的 ReLU 函数，而无需第二次内核调用。融合优化步骤将检测 ReLU 之后的卷积，验证实现是否支持这些操作，然后将它们融合到一层。

为了调查哪些融合已经发生或没有发生，构建器将其操作记录到构建期间提供的记录器对象。优化步骤在kINFO日志级别。要查看这些消息，请确保将它们记录在ILogger回调中。

融合通常通过创建一个新层来处理，该层的名称包含被融合的两个层的名称。例如，在 MNIST 中，名为 ip1 的全连接层（InnerProduct）与名为relu1的 ReLU 激活层融合，以创建名为ip1 + relu1的新层。

13.2.4.2. Types Of Fusions

以下列表描述了支持的融合类型。

支持的层融合

ReLU ReLU Activation

执行 ReLU 的激活层，然后执行 ReLU 的激活将被单个激活层替换。

Convolution and ReLU Activation 卷积层可以是任何类型，并且对值没有限制。激活层必须是 ReLU 类型。

Convolution and GELU Activation

输入输出精度要一致；它们都是 FP16 或 INT8。激活层必须是 GELU 类型。 TensorRT 应该在具有 CUDA 10.0 或更高版本的 Turing 或更高版本的设备上运行。

Convolution and Clip Activation

卷积层可以是任何类型，并且对值没有限制。激活层必须是Clip类型。

Scale and Activation

Scale 层后跟一个 Activation 层可以融合成一个 Activation 层。

Convolution And ElementWise Operation

卷积层后跟 ElementWise 层中的简单求和、最小值或最大值可以融合到卷积层中。总和不得使用广播，除非广播跨越批量大小。

Padding and Convolution/Deconvolution

如果所有填充大小都是非负的，则可以将后跟卷积或反卷积的填充融合到单个卷积/反卷积层中。

Shuffle and Reduce

一个没有 reshape 的 Shuffle 层，然后是一个 Reduce 层，可以融合成一个 Reduce 层。 Shuffle 层可以执行排列，但不能执行任何重塑操作。 Reduce 层必须有一组keepDimensions维度。

Shuffle and Shuffle

每个 Shuffle 层由转置、重塑和第二个转置组成。一个 Shuffle 层后跟另一个 Shuffle 层可以被单个 Shuffle 替换（或什么都没有）。如果两个 Shuffle 层都执行 reshape 操作，则只有当第一个 shuffle 的第二个转置是第二个 shuffle 的第一个转置的逆时，才允许这种融合。

Scale 可以擦除添加0 、乘以1或计算 1 的幂的Scale 层。

Convolution and Scale

卷积层后跟kUNIFORM或kCHANNEL的 Scale 层融合为单个卷积。如果秤具有非恒定功率参数，则禁用此融合。

Reduce

执行平均池化的 Reduce 层将被 Pooling 层取代。 Reduce 层必须有一个keepDimensions集，使用kAVG操作在批处理之前从CHW输入格式减少H和W维度。

Convolution and Pooling 卷积层和池化层必须具有相同的精度。卷积层可能已经具有来自先前融合的融合激活操作。

Depthwise Separable Convolution 带有激活的深度卷积，然后是带有激活的卷积，有时可能会融合到单个优化的 DepSepConvolution 层中。两个卷积的精度必须为 INT8，并且设备的计算能力必须为 7.2 或更高版本。

SoftMax and Log

如果 SoftMax 尚未与先前的日志操作融合，则可以将其融合为单个 Softmax 层。

SoftMax 和 TopK

可以融合成单层。 SoftMax 可能包含也可能不包含 Log 操作。

FullyConnected

FullyConnected 层将被转换为 Convolution 层，所有用于卷积的融合都会生效。

Supported Reduction Operation Fusions

GELU

一组表示以下方程的 Unary 层和 ElementWise 层可以融合到单个 GELU 归约操作中。

$0.5x × （1+tanh （2/π （x+0.044715x^3）））$

或替代表示： $0.5x × （1+erf （x/\sqrt{2}））$

L1Norm

一个一元层kABS操作和一个 Reduce 层kSUM操作可以融合成一个 L1Norm 归约操作。

Sum of Squares

具有相同输入（平方运算）的乘积 ElementWise 层后跟kSUM 约简可以融合为单个平方和约简运算。

L2Norm

kSQRT UnaryOperation之后的平方和运算可以融合到单个 L2Norm 归约运算中。

LogSum

一个缩减层kSUM后跟一个kLOG UnaryOperation 可以融合成一个单一的 LogSum 缩减操作。

LogSumExp

一个一元kEXP ElementWise操作后跟一个 LogSum 融合可以融合成一个单一的 LogSumExp 约简。

13.2.4.3. PointWise Fusion

多个相邻的 PointWise 层可以融合到一个 PointWise 层中，以提高性能。

支持以下类型的 PointWise 层，但有一些限制：

Activation

每个ActivationType 。

Constant

仅具有单个值的常量（大小 == 1）。

ElementWise

每个ElementWiseOperation 。

PointWise

PointWise本身也是一个 PointWise 层。

Scale

仅支持ScaleMode：：kUNIFORM 。

Unary

每个UnaryOperation 。

融合的 PointWise 层的大小不是无限的，因此，某些 PointWise 层可能无法融合。

Fusion 创建一个新层，其名称由融合的两个层组成。例如，名为add1的 ElementWise 层与名为relu1的 ReLU 激活层融合，新层名称为： fusedPointwiseNode（add1， relu1） 。

13.2.4.4. Q/DQ Fusion

从 QAT 工具（如NVIDIA 的 PyTorch 量化工具包）生成的量化 INT8 图由具有比例和零点的onnx：：QuantizeLinear和onnx：：DequantizeLinear节点对 （Q/DQ） 组成。从 TensorRT 7.0 开始，要求zero_point为0 。

Q/DQ 节点帮助将 FP32 值转换为 INT8，反之亦然。这样的图在 FP32 精度上仍然会有权重和偏差。

权重之后是 Q/DQ 节点对，以便在需要时可以对它们进行量化/去量化。偏置量化是使用来自激活和权重的尺度执行的，因此偏置输入不需要额外的 Q/DQ 节点对。偏差量化的假设是$S_weights * S_input = S_bias$ 。

与 Q/DQ 节点相关的融合包括量化/去量化权重，在不改变模型数学等价性的情况下对 Q/DQ 节点进行交换，以及擦除冗余 Q/DQ 节点。应用 Q/DQ 融合后，其余的构建器优化将应用于图。

Fuse Q/DQ with weighted node （Conv， FC， Deconv）

如果我们有一个

[DequantizeLinear (Activations), DequantizeLinear (weights)] > Node >
        QuantizeLinear

（ ［DQ， DQ］ 》 Node 》 Q ） 序列，然后融合到量化节点 （ QNode ）。