echo "######alloc swap######"
if [ ! -e /swapfile ];then
    sudo fallocate -l 4G /swapfile
    sudo chmod 600 /swapfile
    sudo mkswap /swapfile
    sudo /bin/sh -c 'echo  "/swapfile \t none \t swap \t defaults \t 0 \t 0" >> /etc/fstab'
    sudo swapon -a
fi

13.2. Optimizing TensorRT Performance

以下部分重點(diǎn)介紹 GPU 上的一般推理流程和一些提高性能的一般策略。這些想法適用于大多數(shù) CUDA 程序員，但對(duì)于來(lái)自其他背景的開(kāi)發(fā)人員可能并不那么明顯。

13.2.1. Batching

最重要的優(yōu)化是使用批處理并行計(jì)算盡可能多的結(jié)果。在 TensorRT 中，批次是可以統(tǒng)一處理的輸入的集合。批次中的每個(gè)實(shí)例都具有相同的形狀，并以完全相同的方式流經(jīng)網(wǎng)絡(luò)。因此，每個(gè)實(shí)例都可以簡(jiǎn)單地并行計(jì)算。

網(wǎng)絡(luò)的每一層都有計(jì)算前向推理所需的一定數(shù)量的開(kāi)銷(xiāo)和同步。通過(guò)并行計(jì)算更多結(jié)果，這種開(kāi)銷(xiāo)可以更有效地得到回報(bào)。此外，許多層的性能受到輸入中最小維度的限制。如果批量大小為 1 或較小，則此大小通常可能是性能限制維度。例如，具有V個(gè)輸入和K個(gè)輸出的完全連接層可以針對(duì)一個(gè)批次實(shí)例實(shí)現(xiàn)為1xV矩陣與VxK權(quán)重矩陣的矩陣乘法。如果對(duì)N個(gè)實(shí)例進(jìn)行批處理，則這將變?yōu)镹xV乘以VxK矩陣。向量矩陣乘法器變成矩陣矩陣乘法器，效率更高。

更大的批量大小幾乎總是在 GPU 上更有效。非常大的批次，例如N 》 2^16 ，有時(shí)可能需要擴(kuò)展索引計(jì)算，因此應(yīng)盡可能避免。但通常，增加批量大小會(huì)提高總吞吐量。此外，當(dāng)網(wǎng)絡(luò)包含 MatrixMultiply 層或完全連接層時(shí)，如果硬件支持，由于使用了 Tensor Cores，32 的倍數(shù)的批大小往往對(duì) FP16 和 INT8 推理具有最佳性能。

由于應(yīng)用程序的組織，有時(shí)無(wú)法進(jìn)行批處理推理工作。在一些常見(jiàn)的應(yīng)用程序中，例如根據(jù)請(qǐng)求進(jìn)行推理的服務(wù)器，可以實(shí)現(xiàn)機(jī)會(huì)批處理。對(duì)于每個(gè)傳入的請(qǐng)求，等待時(shí)間T 。如果在此期間有其他請(qǐng)求進(jìn)來(lái)，請(qǐng)將它們一起批處理。否則，繼續(xù)進(jìn)行單實(shí)例推理。這種類(lèi)型的策略為每個(gè)請(qǐng)求增加了固定的延遲，但可以將系統(tǒng)的最大吞吐量提高幾個(gè)數(shù)量級(jí)。

使用批處理

如果在創(chuàng)建網(wǎng)絡(luò)時(shí)使用顯式批處理模式，則批處理維度是張量維度的一部分，您可以通過(guò)添加優(yōu)化配置文件來(lái)指定批處理大小和批處理大小的范圍以?xún)?yōu)化引擎。有關(guān)更多詳細(xì)信息，請(qǐng)參閱使用動(dòng)態(tài)形狀部分。

如果在創(chuàng)建網(wǎng)絡(luò)時(shí)使用隱式批處理模式，則IExecutionContext：：execute （ Python 中的IExecutionContext.execute ）和IExecutionContext：：enqueue （ Python 中的IExecutionContext.execute_async ）方法采用批處理大小參數(shù)。在使用IBuilder：：setMaxBatchSize （ Python中的 Builder.max_batch_size ）構(gòu)建優(yōu)化網(wǎng)絡(luò)時(shí)，還應(yīng)該為構(gòu)建器設(shè)置最大批量大小。當(dāng)調(diào)用IExecutionContext：：execute或enqueue時(shí)，作為綁定參數(shù)傳遞的綁定是按張量組織的，而不是按實(shí)例組織的。換句話(huà)說(shuō)，一個(gè)輸入實(shí)例的數(shù)據(jù)沒(méi)有組合到一個(gè)連續(xù)的內(nèi)存區(qū)域中。相反，每個(gè)張量綁定都是該張量的實(shí)例數(shù)據(jù)數(shù)組。

另一個(gè)考慮因素是構(gòu)建優(yōu)化的網(wǎng)絡(luò)會(huì)針對(duì)給定的最大批量大小進(jìn)行優(yōu)化。最終結(jié)果將針對(duì)最大批量大小進(jìn)行調(diào)整，但對(duì)于任何較小的批量大小仍然可以正常工作。可以運(yùn)行多個(gè)構(gòu)建操作來(lái)為不同的批量大小創(chuàng)建多個(gè)優(yōu)化引擎，然后在運(yùn)行時(shí)根據(jù)實(shí)際批量大小選擇要使用的引擎。

13.2.2. Streaming

一般來(lái)說(shuō)，CUDA 編程流是一種組織異步工作的方式。放入流中的異步命令保證按順序運(yùn)行，但相對(duì)于其他流可能會(huì)亂序執(zhí)行。特別是，兩個(gè)流中的異步命令可以被調(diào)度為同時(shí)運(yùn)行（受硬件限制）。

在 TensorRT 和推理的上下文中，優(yōu)化的最終網(wǎng)絡(luò)的每一層都需要在 GPU 上工作。但是，并非所有層都能夠充分利用硬件的計(jì)算能力。在單獨(dú)的流中安排請(qǐng)求允許在硬件可用時(shí)立即安排工作，而無(wú)需進(jìn)行不必要的同步。即使只有一些層可以重疊，整體性能也會(huì)提高。

使用流式傳輸

識(shí)別獨(dú)立的推理批次。

為網(wǎng)絡(luò)創(chuàng)建一個(gè)引擎。

cudaStreamCreate為每個(gè)獨(dú)立批次創(chuàng)建一個(gè) CUDA 流，并為每個(gè)獨(dú)立批次創(chuàng)建一個(gè)IExecutionContext 。

IExecutionContext：：enqueue從適當(dāng)?shù)腎ExecutionContext請(qǐng)求異步結(jié)果并傳入適當(dāng)?shù)牧鱽?lái)啟動(dòng)推理工作。

在所有工作啟動(dòng)后，與所有流同步以等待結(jié)果。執(zhí)行上下文和流可以重用于以后的獨(dú)立工作批次。

多個(gè)流 運(yùn)行多個(gè)并發(fā)流通常會(huì)導(dǎo)致多個(gè)流同時(shí)共享計(jì)算資源的情況。這意味著與優(yōu)化 TensorRT 引擎時(shí)相比，推理期間網(wǎng)絡(luò)可用的計(jì)算資源可能更少。這種資源可用性的差異可能會(huì)導(dǎo)致 TensorRT 選擇一個(gè)對(duì)于實(shí)際運(yùn)行時(shí)條件不是最佳的內(nèi)核。為了減輕這種影響，您可以在引擎創(chuàng)建期間限制可用計(jì)算資源的數(shù)量，使其更接近實(shí)際運(yùn)行時(shí)條件。這種方法通常以延遲為代價(jià)來(lái)提高吞吐量。有關(guān)更多信息，請(qǐng)參閱限制計(jì)算資源。

也可以將多個(gè)主機(jī)線(xiàn)程與流一起使用。一種常見(jiàn)的模式是將傳入的請(qǐng)求分派到等待工作線(xiàn)程池中。在這種情況下，工作線(xiàn)程池將每個(gè)都有一個(gè)執(zhí)行上下文和 CUDA 流。當(dāng)工作變得可用時(shí)，每個(gè)線(xiàn)程將在自己的流中請(qǐng)求工作。每個(gè)線(xiàn)程將與其流同步以等待結(jié)果，而不會(huì)阻塞其他工作線(xiàn)程。

13.2.3. CUDA Graphs

CUDA 圖是一種表示內(nèi)核序列（或更一般地是圖）的方式，其調(diào)度方式允許由 CUDA 優(yōu)化。當(dāng)您的應(yīng)用程序性能對(duì)將內(nèi)核排入隊(duì)列所花費(fèi)的 CPU 時(shí)間敏感時(shí)，這可能特別有用。 TensorRT 的enqueuev2（）方法支持對(duì)不需要 CPU 交互的模型進(jìn)行 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);

不支持圖的模型包括帶有循環(huán)或條件的模型。在這種情況下， cudaStreamEndCapture（）將返回cudaErrorStreamCapture*錯(cuò)誤，表示圖捕獲失敗，但上下文可以繼續(xù)用于沒(méi)有 CUDA 圖的正常推理。

捕獲圖時(shí)，重要的是要考慮在存在動(dòng)態(tài)形狀時(shí)使用的兩階段執(zhí)行策略。

更新模型的內(nèi)部狀態(tài)以考慮輸入大小的任何變化

將工作流式傳輸?shù)?GPU

對(duì)于在構(gòu)建時(shí)輸入大小固定的模型，第一階段不需要每次調(diào)用工作。否則，如果自上次調(diào)用以來(lái)輸入大小發(fā)生了變化，則可能需要進(jìn)行一些工作來(lái)更新派生屬性。

第一階段的工作不是為捕獲而設(shè)計(jì)的，即使捕獲成功也可能會(huì)增加模型執(zhí)行時(shí)間。因此，在更改輸入的形狀或形狀張量的值后，調(diào)用enqueueV2（）一次以在捕獲圖形之前刷新延遲更新。

使用 TensorRT 捕獲的圖特定于捕獲它們的輸入大小，以及執(zhí)行上下文的狀態(tài)。修改捕獲圖表的上下文將導(dǎo)致執(zhí)行圖表時(shí)未定義的行為 – 特別是，如果應(yīng)用程序通過(guò)createExecutionContextWithoutDeviceMemory（）為激活提供自己的內(nèi)存，則內(nèi)存地址也會(huì)作為圖表的一部分被捕獲。綁定位置也被捕獲為圖表的一部分。

trtexec允許您檢查構(gòu)建的 TensorRT 引擎是否與 CUDA 圖形捕獲兼容。有關(guān)詳細(xì)信息，請(qǐng)參閱trtexec部分。

13.2.4. Enabling Fusion

13.2.4.1. Layer Fusion

TensorRT 嘗試在構(gòu)建階段在網(wǎng)絡(luò)中執(zhí)行許多不同類(lèi)型的優(yōu)化。在第一階段，盡可能將層融合在一起。融合將網(wǎng)絡(luò)轉(zhuǎn)換為更簡(jiǎn)單的形式，但保持相同的整體行為。在內(nèi)部，許多層實(shí)現(xiàn)具有在創(chuàng)建網(wǎng)絡(luò)時(shí)無(wú)法直接訪(fǎng)問(wèn)的額外參數(shù)和選項(xiàng)。相反，融合優(yōu)化步驟檢測(cè)支持的操作模式，并將多個(gè)層融合到一個(gè)具有內(nèi)部選項(xiàng)集的層中。

考慮卷積后跟 ReLU 激活的常見(jiàn)情況。要?jiǎng)?chuàng)建具有這些操作的網(wǎng)絡(luò)，需要使用 addConvolution 添加卷積層，然后使用addActivation和kRELU的ActivationType添加激活層。未優(yōu)化的圖將包含用于卷積和激活的單獨(dú)層。卷積的內(nèi)部實(shí)現(xiàn)支持直接從卷積核一步計(jì)算輸出上的 ReLU 函數(shù)，而無(wú)需第二次內(nèi)核調(diào)用。融合優(yōu)化步驟將檢測(cè) ReLU 之后的卷積，驗(yàn)證實(shí)現(xiàn)是否支持這些操作，然后將它們?nèi)诤系揭粚印?/p>

為了調(diào)查哪些融合已經(jīng)發(fā)生或沒(méi)有發(fā)生，構(gòu)建器將其操作記錄到構(gòu)建期間提供的記錄器對(duì)象。優(yōu)化步驟在kINFO日志級(jí)別。要查看這些消息，請(qǐng)確保將它們記錄在ILogger回調(diào)中。

融合通常通過(guò)創(chuàng)建一個(gè)新層來(lái)處理，該層的名稱(chēng)包含被融合的兩個(gè)層的名稱(chēng)。例如，在 MNIST 中，名為 ip1 的全連接層（InnerProduct）與名為relu1的 ReLU 激活層融合，以創(chuàng)建名為ip1 + relu1的新層。

13.2.4.2. Types Of Fusions

以下列表描述了支持的融合類(lèi)型。

支持的層融合

ReLU ReLU Activation

執(zhí)行 ReLU 的激活層，然后執(zhí)行 ReLU 的激活將被單個(gè)激活層替換。

Convolution and ReLU Activation 卷積層可以是任何類(lèi)型，并且對(duì)值沒(méi)有限制。激活層必須是 ReLU 類(lèi)型。

Convolution and GELU Activation

輸入輸出精度要一致；它們都是 FP16 或 INT8。激活層必須是 GELU 類(lèi)型。 TensorRT 應(yīng)該在具有 CUDA 10.0 或更高版本的 Turing 或更高版本的設(shè)備上運(yùn)行。

Convolution and Clip Activation

卷積層可以是任何類(lèi)型，并且對(duì)值沒(méi)有限制。激活層必須是Clip類(lèi)型。

Scale and Activation

Scale 層后跟一個(gè) Activation 層可以融合成一個(gè) Activation 層。

Convolution And ElementWise Operation

卷積層后跟 ElementWise 層中的簡(jiǎn)單求和、最小值或最大值可以融合到卷積層中。總和不得使用廣播，除非廣播跨越批量大小。

Padding and Convolution/Deconvolution

如果所有填充大小都是非負(fù)的，則可以將后跟卷積或反卷積的填充融合到單個(gè)卷積/反卷積層中。

Shuffle and Reduce

一個(gè)沒(méi)有 reshape 的 Shuffle 層，然后是一個(gè) Reduce 層，可以融合成一個(gè) Reduce 層。 Shuffle 層可以執(zhí)行排列，但不能執(zhí)行任何重塑操作。 Reduce 層必須有一組keepDimensions維度。

Shuffle and Shuffle

每個(gè) Shuffle 層由轉(zhuǎn)置、重塑和第二個(gè)轉(zhuǎn)置組成。一個(gè) Shuffle 層后跟另一個(gè) Shuffle 層可以被單個(gè) Shuffle 替換（或什么都沒(méi)有）。如果兩個(gè) Shuffle 層都執(zhí)行 reshape 操作，則只有當(dāng)?shù)谝粋€(gè) shuffle 的第二個(gè)轉(zhuǎn)置是第二個(gè) shuffle 的第一個(gè)轉(zhuǎn)置的逆時(shí)，才允許這種融合。

Scale 可以擦除添加0 、乘以1或計(jì)算 1 的冪的Scale 層。

Convolution and Scale

卷積層后跟kUNIFORM或kCHANNEL的 Scale 層融合為單個(gè)卷積。如果秤具有非恒定功率參數(shù)，則禁用此融合。

Reduce

執(zhí)行平均池化的 Reduce 層將被 Pooling 層取代。 Reduce 層必須有一個(gè)keepDimensions集，使用kAVG操作在批處理之前從CHW輸入格式減少H和W維度。

Convolution and Pooling 卷積層和池化層必須具有相同的精度。卷積層可能已經(jīng)具有來(lái)自先前融合的融合激活操作。

Depthwise Separable Convolution 帶有激活的深度卷積，然后是帶有激活的卷積，有時(shí)可能會(huì)融合到單個(gè)優(yōu)化的 DepSepConvolution 層中。兩個(gè)卷積的精度必須為 INT8，并且設(shè)備的計(jì)算能力必須為 7.2 或更高版本。

SoftMax and Log

如果 SoftMax 尚未與先前的日志操作融合，則可以將其融合為單個(gè) Softmax 層。

SoftMax 和 TopK

可以融合成單層。 SoftMax 可能包含也可能不包含 Log 操作。

FullyConnected

FullyConnected 層將被轉(zhuǎn)換為 Convolution 層，所有用于卷積的融合都會(huì)生效。

Supported Reduction Operation Fusions

GELU

一組表示以下方程的 Unary 層和 ElementWise 層可以融合到單個(gè) GELU 歸約操作中。

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

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

L1Norm

一個(gè)一元層kABS操作和一個(gè) Reduce 層kSUM操作可以融合成一個(gè) L1Norm 歸約操作。

Sum of Squares

具有相同輸入（平方運(yùn)算）的乘積 ElementWise 層后跟kSUM 約簡(jiǎn)可以融合為單個(gè)平方和約簡(jiǎn)運(yùn)算。

L2Norm

kSQRT UnaryOperation之后的平方和運(yùn)算可以融合到單個(gè) L2Norm 歸約運(yùn)算中。

LogSum

一個(gè)縮減層kSUM后跟一個(gè)kLOG UnaryOperation 可以融合成一個(gè)單一的 LogSum 縮減操作。

LogSumExp

一個(gè)一元kEXP ElementWise操作后跟一個(gè) LogSum 融合可以融合成一個(gè)單一的 LogSumExp 約簡(jiǎn)。

13.2.4.3. PointWise Fusion

多個(gè)相鄰的 PointWise 層可以融合到一個(gè) PointWise 層中，以提高性能。

支持以下類(lèi)型的 PointWise 層，但有一些限制：

Activation

每個(gè)ActivationType 。

Constant

僅具有單個(gè)值的常量（大小 == 1）。

ElementWise

每個(gè)ElementWiseOperation 。

PointWise

PointWise本身也是一個(gè) PointWise 層。

Scale

僅支持ScaleMode：：kUNIFORM 。

Unary

每個(gè)UnaryOperation 。

融合的 PointWise 層的大小不是無(wú)限的，因此，某些 PointWise 層可能無(wú)法融合。

Fusion 創(chuàng)建一個(gè)新層，其名稱(chēng)由融合的兩個(gè)層組成。例如，名為add1的 ElementWise 層與名為relu1的 ReLU 激活層融合，新層名稱(chēng)為： fusedPointwiseNode（add1， relu1） 。

13.2.4.4. Q/DQ Fusion

從 QAT 工具（如NVIDIA 的 PyTorch 量化工具包）生成的量化 INT8 圖由具有比例和零點(diǎn)的onnx：：QuantizeLinear和onnx：：DequantizeLinear節(jié)點(diǎn)對(duì) （Q/DQ） 組成。從 TensorRT 7.0 開(kāi)始，要求zero_point為0 。

Q/DQ 節(jié)點(diǎn)幫助將 FP32 值轉(zhuǎn)換為 INT8，反之亦然。這樣的圖在 FP32 精度上仍然會(huì)有權(quán)重和偏差。

權(quán)重之后是 Q/DQ 節(jié)點(diǎn)對(duì)，以便在需要時(shí)可以對(duì)它們進(jìn)行量化/去量化。偏置量化是使用來(lái)自激活和權(quán)重的尺度執(zhí)行的，因此偏置輸入不需要額外的 Q/DQ 節(jié)點(diǎn)對(duì)。偏差量化的假設(shè)是$S_weights * S_input = S_bias$ 。

與 Q/DQ 節(jié)點(diǎn)相關(guān)的融合包括量化/去量化權(quán)重，在不改變模型數(shù)學(xué)等價(jià)性的情況下對(duì) Q/DQ 節(jié)點(diǎn)進(jìn)行交換，以及擦除冗余 Q/DQ 節(jié)點(diǎn)。應(yīng)用 Q/DQ 融合后，其余的構(gòu)建器優(yōu)化將應(yīng)用于圖。

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

如果我們有一個(gè)

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

（ ［DQ， DQ］ 》 Node 》 Q ） 序列，然后融合到量化節(jié)點(diǎn) （ QNode ）。