在開始使用 TensorRT 進行任何優化工作之前,必須確定應該測量什么。沒有衡量標準,就不可能取得可靠的進展或衡量是否取得了成功
Latency
網絡推理的性能度量是從輸入呈現給網絡到輸出可用所經過的時間。這是單個推理的網絡延遲。較低的延遲更好。在某些應用中,低延遲是一項關鍵的安全要求。在其他應用程序中,延遲作為服務質量問題對用戶來說是直接可見的。對于批量處理,延遲可能根本不重要。
Throughput
另一個性能測量是在固定的時間單位內可以完成多少推理。這是網絡的吞吐量。吞吐量越高越好。更高的吞吐量表明更有效地利用固定計算資源。對于批量處理,所花費的總時間將由網絡的吞吐量決定。
查看延遲和吞吐量的另一種方法是確定最大延遲并在該延遲下測量吞吐量。像這樣的服務質量測量可以是用戶體驗和系統效率之間的合理折衷。
在測量延遲和吞吐量之前,您需要選擇開始和停止計時的確切點。根據網絡和應用程序,選擇不同的點可能是有意義的。
在很多應用中,都有一個處理流水線,整個系統的性能可以通過整個處理流水線的延遲和吞吐量來衡量。由于預處理和后處理步驟在很大程度上取決于特定應用程序,因此本節僅考慮網絡推理的延遲和吞吐量。
13.1.1. Wall-clock Timing
經過時間(計算開始和結束之間經過的時間)可用于測量應用程序的整體吞吐量和延遲,以及將推理時間置于更大系統的上下文中。 C++11 在標準庫中提供了高精度計時器。例如,std::chrono::system_clock表示系統范圍的經過時間,而std::chrono::high_resolution_clock以可用的最高精度測量時間。
以下示例代碼片段顯示了測量網絡推理主機時間:
#includeauto startTime = std::chrono::high_resolution_clock::now(); context->enqueueV2(&buffers[0], stream, nullptr); cudaStreamSynchronize(stream); auto endTime = std::chrono::high_resolution_clock::now(); float totalTime = std::chrono::duration (endTime - startTime).count();
如果設備上一次只發生一個推理,那么這可能是一種簡單的方法來分析各種操作所花費的時間。推理通常是異步的,因此請確保添加顯式 CUDA 流或設備同步以等待結果可用。
13.1.2. CUDA Events
僅在主機上計時的一個問題是它需要主機/設備同步。優化的應用程序可能會在設備上并行運行許多推理,并具有重疊的數據移動。此外,同步本身給定時測量增加了一些噪聲。 為了幫助解決這些問題,CUDA 提供了一個事件 API 。此 API 允許您將事件放入 CUDA 流中,這些事件將在遇到事件時由 GPU 打上時間戳。然后,時間戳的差異可以告訴您不同操作花費了多長時間。
以下示例代碼片段顯示了計算兩個 CUDA 事件之間的時間:
cudaEvent_t start, end; cudaEventCreate(&start); cudaEventCreate(&end); cudaEventRecord(start, stream); context->enqueueV2(&buffers[0], stream, nullptr); cudaEventRecord(end, stream); cudaEventSynchronize(end); float totalTime; cudaEventElapsedTime(&totalTime, start, end);
13.1.3. Built-In TensorRT Profiling
深入挖掘推理性能需要在優化網絡中進行更細粒度的時序測量。 TensorRT 有一個Profiler ( C++ , Python ) 接口,您可以實現該接口以便讓 TensorRT 將分析信息傳遞給您的應用程序。調用時,網絡將以分析模式運行。完成推理后,將調用您的類的分析器對象以報告網絡中每一層的時間。這些時序可用于定位瓶頸、比較序列化引擎的不同版本以及調試性能問題。
分析信息可以從常規推理enqueueV2()啟動或 CUDA 圖啟動中收集。有關詳細信息,請參閱IExecutionContext::setProfiler()和IExecutionContext::reportToProfiler() ( C++ 、 Python )。
循環內的層編譯為單個單片層,因此,這些層的單獨時序不可用。
公共示例代碼 ( common.h ) 中提供了一個展示如何使用IProfiler接口的示例,然后在位于 GitHub 存儲庫中的sampleNMT中使用。
您還可以使用trtexec在給定輸入網絡或計劃文件的情況下使用 TensorRT 分析網絡。有關詳細信息,請參閱trtexec部分。
13.1.4. CUDA Profiling Tools
推薦的 CUDA 分析器是NVIDIA Nsight? Systems 。一些 CUDA 開發人員可能更熟悉 nvprof 和 nvvp,但是,這些已被棄用。在任何情況下,這些分析器都可以用于任何 CUDA 程序,以報告有關在執行期間啟動的內核、主機和設備之間的數據移動以及使用的 CUDA API 調用的時序信息。
Nsight Systems 可以通過多種方式配置,以僅報告程序執行的一部分的時序信息,或者也可以將傳統的 CPU 采樣配置文件信息與 GPU 信息一起報告。
僅分析推理階段
分析 TensorRT 應用程序時,您應該僅在構建引擎后啟用分析。在構建階段,所有可能的策略都被嘗試和計時。分析這部分執行將不會顯示任何有意義的性能測量,并將包括所有可能的內核,而不是實際選擇用于推理的內核。限制分析范圍的一種方法是:
第一階段:構建應用程序,然后在一個階段序列化引擎。
第二階段:加載序列化引擎并在第二階段運行推理并僅對第二階段進行分析。
如果應用程序無法序列化引擎,或者應用程序必須連續運行兩個階段,您還可以在第二階段周圍添加cudaProfilerStart() / cudaProfilerStop() CUDA API,并在 Nsight Systems 命令中添加-c cudaProfilerApi標志以僅配置文件cudaProfilerStart()和cudaProfilerStop()之間的部分。
在 Nsight Systems 中使用 NVTX 跟蹤 啟用NVIDIA 工具擴展 SDK (NVTX)跟蹤允許 Nsight Compute 和 Nsight Systems 收集由 TensorRT 應用程序生成的數據。 NVTX 是一個基于 C 的 API,用于標記應用程序中的事件和范圍。
將內核名稱解碼回原始網絡中的層可能很復雜。因此,TensorRT 使用 NVTX 為每一層標記一個范圍,然后允許 CUDA 分析器將每一層與調用來實現它的內核相關聯。在 TensorRT 中,NVTX 有助于將運行時引擎層的執行與 CUDA內核調用相關聯。 Nsight Systems 支持在時間軸上收集和可視化這些事件和范圍。 Nsight Compute 還支持在應用程序掛起時收集和顯示給定線程中所有活動 NVTX 域和范圍的狀態。
在 TensorRT 中,每一層都可以啟動一個或多個內核來執行其操作。啟動的確切內核取決于優化的網絡和存在的硬件。根據構建器的選擇,可能會有多個額外的操作對穿插在層計算中的數據進行重新排序;這些重新格式化操作可以作為設備到設備的內存副本或自定義內核來實現。
例如,以下屏幕截圖來自 Nsight Systems。
控制 NVTX 跟蹤中的詳細程度
默認情況下,TensorRT 僅在 NVTX 標記中顯示層名稱,而用戶可以在構建引擎時通過設置IBuilderConfig中的 ProfilingVerbosity 來控制細節級別。例如,要禁用 NVTX 跟蹤,請將 ProfilingVerbosity 設置為kNONE :
C++
builderConfig-》setProfilingVerbosity(ProfilingVerbosity::kNONE);
Python
builder_config.profiling_verbosity = trt.ProfilingVerbosity.NONE
另一方面,您可以通過將ProfilingVerbosity設置為kDETAILED來選擇允許 TensorRT 在 NVTX 標記中打印更詳細的層信息,包括輸入和輸出尺寸、操作、參數、順序編號等:
C++
builderConfig-》setProfilingVerbosity(ProfilingVerbosity::kDETAILED);
Python
builder_config.profiling_verbosity = trt.ProfilingVerbosity.DETAILED
trtexec運行 Nsight 系統 以下是使用trtexec工具收集 Nsight Systems 配置文件的命令示例:
trtexec --onnx=foo.onnx --profilingVerbosity=detailed --saveEngine=foo.plan
nsys profile -o foo_profile trtexec --loadEngine=foo.plan --warmUp=0 --duration=0 --iterations=50
第一個命令構建引擎并將其序列化為foo.plan ,第二個命令使用foo.plan運行推理并生成一個foo_profile.qdrep文件,然后可以在 Nsight Systems GUI 界面中打開該文件以進行可視化。
--profilingVerbosity=detailed標志允許 TensorRT 在 NVTX 標記中顯示更詳細的層信息,而--warmUp =0 --duration=0 --iterations=50標志允許您控制要運行的推理迭代次數。默認情況下, trtexec運行推理三秒鐘,這可能會導致輸出 qdrep 文件非常大。
13.1.5. Tracking Memory
跟蹤內存使用情況與執行性能一樣重要。通常,設備上的內存比主機上的內存更受限制。為了跟蹤設備內存,推薦的機制是創建一個簡單的自定義 GPU 分配器,它在內部保留一些統計信息,然后使用常規 CUDA 內存分配函數cudaMalloc和cudaFree 。
可以為構建器IBuilder設置自定義 GPU 分配器以進行網絡優化,并在使用IGpuAllocator API反序列化引擎時為IRuntime 設置。自定義分配器的一個想法是跟蹤當前分配的內存量,并將帶有時間戳和其他信息的分配事件推送到分配事件的全局列表中。查看分配事件列表可以分析一段時間內的內存使用情況。
在移動平臺上,GPU 內存和 CPU 內存共享系統內存。在內存大小非常有限的設備上,如 Nano,系統內存可能會因大型網絡而耗盡;甚至所需的 GPU 內存也小于系統內存。在這種情況下,增加系統交換大小可以解決一些問題。一個示例腳本是:
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
以下部分重點介紹 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 )。
支持權重的 Q/DQ 節點對需要加權節點支持多個輸入。因此,我們支持添加第二個輸入(用于權重張量)和第三個輸入(用于偏置張量)。可以使用setInput(index, tensor) API 為卷積、反卷積和全連接層設置其他輸入,其中 index = 2 用于權重張量, index = 3 用于偏置張量。
在與加權節點融合期間,我們會將 FP32 權重量化為 INT8,并將其與相應的加權節點融合。類似地,FP32 偏差將被量化為 INT32 并融合。
使用非加權節點融合 Q/DQ
如果我們有一個DequantizeLinear 》 Node 》 QuantizeLinear ( DQ 》 Node 》 Q ) 序列,那么它將融合到量化節點 ( QNode )。
Commutate Q/DQ nodes
DequantizeLinear commutation is allowed when $Φ (DQ (x)) == DQ (Φ (x))$ 。 QuantizeLinear commutation is allowed when $Q (Φ (x)) == Φ (Q (x))$ 。
此外,交換邏輯還考慮了可用的內核實現,從而保證了數學等價性。
Insert missing Q/DQ nodes
如果一個節點缺少一個 Q/DQ 節點對,并且$max (abs (Φ (x))) == max (abs (x))$ ; (例如,MaxPool),將插入缺少的 Q/DQ 對以運行更多具有 INT8 精度的節點。
Erase redundant Q/DQ nodes
有可能在應用所有優化之后,該圖仍然有 Q/DQ 節點對,它們本身就是一個空操作。 Q/DQ 節點擦除融合將刪除此類冗余對。
13.2.5. Limiting Compute Resources
當減少的數量更好地代表運行時的預期條件時,限制在引擎創建期間可用于 TensorRT 的計算資源量是有益的。例如,當期望 GPU 與 TensorRT 引擎并行執行額外工作時,或者當期望引擎在資源較少的不同 GPU 上運行時(請注意,推薦的方法是在 GPU 上構建引擎,即將用于推理,但這可能并不總是可行的)。
您可以通過以下步驟限制可用計算資源的數量:
啟動 CUDA MPS 控制守護進程。nvidia-cuda-mps-control -d
CUDA_MPS_ACTIVE_THREAD_PERCENTAGE環境變量一起使用的計算資源量。例如,導出 CUDA_MPS_ACTIVE_THREAD_PERCENTAGE=50 。
構建網絡引擎。
停止 CUDA MPS 控制守護程序。echo quit | nvidia-cuda-mps-control
生成的引擎針對減少的計算核心數量(本例中為 50%)進行了優化,并在推理期間使用類似條件時提供更好的吞吐量。鼓勵您嘗試不同數量的流和不同的 MPS 值,以確定網絡的最佳性能。
有關nvidia-cuda-mps-control 的更多詳細信息,請參閱nvidia-cuda-mps-control文檔和此處的相關 GPU 要求。
13.3. Optimizing Layer Performance
以下描述詳細說明了如何優化列出的層。
Concatenation Layer
如果使用隱式批處理維度,連接層的主要考慮是如果多個輸出連接在一起,它們不能跨批處理維度廣播,必須顯式復制。大多數層支持跨批次維度的廣播以避免不必要地復制數據,但如果輸出與其他張量連接,這將被禁用。
Gather Layer 請使用0軸。 Gather 層沒有可用的融合。
Reduce Layer
要從 Reduce 層獲得最大性能,請在最后一個維度上執行歸約(尾部歸約)。這允許最佳內存通過順序內存位置讀取/寫入模式。如果進行常見的歸約操作,請盡可能以將融合為單個操作的方式表達歸約。
RNN Layer
如果可能,請選擇使用較新的 RNNv2 接口而不是傳統的 RNN 接口。較新的接口支持可變序列長度和可變批量大小,以及具有更一致的接口。為了獲得最佳性能,更大的批量大小更好。通常,大小為 64 的倍數可獲得最高性能。雙向 RNN 模式由于增加了依賴性而阻止了波前傳播,因此,它往往更慢。
此外,新引入的基于 ILoop 的 API 提供了一種更靈活的機制,可以在循環中使用通用層,而不受限于一小組預定義的 RNNv2 接口。 ILoop 循環實現了一組豐富的自動循環優化,包括循環融合、展開和循環不變的代碼運動,僅舉幾例。例如,當同一 MatrixMultiply 或 FullyConnected 層的多個實例正確組合以在沿序列維度展開循環后最大化機器利用率時,通常會獲得顯著的性能提升。如果您可以避免 MatrixMultiply 或 FullyConnected 層沿序列維度具有循環數據依賴性,則此方法效果最佳。
Shuffle
如果輸入張量僅用于 shuffle 層,并且該層的輸入和輸出張量不是網絡的輸入和輸出張量,則省略相當于對基礎數據的身份操作的 shuffle 操作。 TensorRT 不會為此類操作執行額外的內核或內存副本。
TopK
要從 TopK 層中獲得最大性能,請使用較小的K值來減少數據的最后一維,以實現最佳的順序內存訪問。通過使用 Shuffle 層來重塑數據,然后適當地重新解釋索引值,可以一次模擬沿多個維度的縮減。
有關層的更多信息,請參閱TensorRT 層。
13.4. Optimizing for Tensor Cores
Tensor Core 是在 NVIDIA GPU 上提供高性能推理的關鍵技術。在 TensorRT 中,所有計算密集型層(MatrixMultiply、FullyConnected、Convolution 和 Deconvolution)都支持 Tensor Core 操作。
如果輸入/輸出張量維度與某個最小粒度對齊,則張量核心層往往會獲得更好的性能:
在卷積和反卷積層中,對齊要求是輸入/輸出通道維度
在 MatrixMultiply 和 FullyConnected 層中,對齊要求是在 MatrixMultiply 中的矩陣維度K和N上,即M x K乘以K x N
下表捕獲了建議的張量維度對齊,以獲得更好的張量核心性能。
在不滿足這些要求的情況下使用 Tensor Core 實現時,TensorRT 會隱式地將張量填充到最接近的對齊倍數,將模型定義中的維度向上舍入,以在不增加計算或內存流量的情況下允許模型中的額外容量。
TensorRT 總是對層使用最快的實現,因此在某些情況下,即使可用,也可能不使用 Tensor Core 實現。
13.5. Optimizing Plugins
TensorRT 提供了一種注冊執行層操作的自定義插件的機制。插件創建者注冊后,您可以在序列化/反序列化過程中查找注冊表找到創建者并將相應的插件對象添加到網絡中。
加載插件庫后,所有 TensorRT 插件都會自動注冊。有關自定義插件的更多信息,請參閱使用自定義層擴展 TensorRT 。
插件的性能取決于執行插件操作的 CUDA 代碼。適用標準CUDA 最佳實踐。在開發插件時,從執行插件操作并驗證正確性的簡單獨立 CUDA 應用程序開始會很有幫助。然后可以通過性能測量、更多單元測試和替代實現來擴展插件程序。代碼運行并優化后,可以作為插件集成到 TensorRT 中。 為了盡可能獲得最佳性能,在插件中支持盡可能多的格式非常重要。這消除了在網絡執行期間對內部重新格式化操作的需要。有關示例,請參閱使用自定義層擴展 TensorRT部分。
13.6. Optimizing Python Performance
使用 Python API 時,大多數相同的性能注意事項都適用。在構建引擎時,構建器優化階段通常會成為性能瓶頸;不是 API 調用來構建網絡。 Python API 和 C++ API 的推理時間應該幾乎相同。
在 Python API 中設置輸入緩沖區涉及使用pycuda或其他 CUDA Python 庫(如cupy )將數據從主機傳輸到設備內存。其工作原理的詳細信息將取決于主機數據的來源。在內部, pycuda支持允許有效訪問內存區域的Python 緩沖區協議。這意味著,如果輸入數據在numpy數組或其他也支持緩沖區協議的類型中以合適的格式可用,則可以有效地訪問和傳輸到 GPU。為了獲得更好的性能,請確保您使用pycuda分配一個頁面鎖定的緩沖區,并在那里寫入最終的預處理輸入。
有關使用 Python API 的更多信息,請參閱Python API 。
13.7. Improving Model Accuracy
TensorRT 可以根據構建器配置以 FP32、FP16 或 INT8 精度執行層。默認情況下,TensorRT 選擇以可實現最佳性能的精度運行層。有時這可能會導致準確性下降。通常,以更高的精度運行層有助于提高準確性,但會影響一些性能。
我們可以采取幾個步驟來提高模型的準確性:
驗證層輸出:
使用Polygraphy轉儲層輸出并驗證沒有 NaN 或 Inf。 --validate選項可以檢查 NaN 和 Infs 。此外,我們可以將層輸出與來自例如 ONNX 運行時的黃金值進行比較。
對于 FP16,模型可能需要重新訓練以確保中間層輸出可以以 FP16 精度表示,而不會出現上溢/下溢。
對于 INT8,考慮使用更具代表性的校準數據集重新校準。如果您的模型來自 PyTorch,除了 TensorRT 中的 PTQ,我們還提供了 NVIDIA 的 Quantization Toolkit for PyTorch for QAT 框架中的 QAT。您可以嘗試這兩種方法并選擇更準確的方法。
操縱層精度:
有時以特定精度運行層會導致輸出不正確。這可能是由于固有的層約束(例如,LayerNorm 輸出不應為 INT8)、模型約束(輸出發散導致準確性差)或報告TensorRT 錯誤。
您可以控制層執行精度和輸出精度。
一個實驗性的調試精度工具可以幫助自動找到以高精度運行的層。
使用算法選擇和可重現構建來禁用不穩定的策略:
當構建+運行與構建+運行之間的準確性發生變化時,可能是由于為層選擇了錯誤的策略。
使用算法選擇器從好的和壞的運行中轉儲策略。將算法選擇器配置為僅允許策略的子集(即僅允許來自良好運行的策略等)。 c. 您可以使用Polygraphy自動執行此過程。
每次運行變化的準確性不應改變;一旦為特定 GPU 構建了引擎,它應該會在多次運行中產生位準確的輸出。如果沒有,請提交TensorRT 錯誤。
關于作者
Ken He 是 NVIDIA 企業級開發者社區經理 & 高級講師,擁有多年的 GPU 和人工智能開發經驗。自 2017 年加入 NVIDIA 開發者社區以來,完成過上百場培訓,幫助上萬個開發者了解人工智能和 GPU 編程開發。在計算機視覺,高性能計算領域完成過多個獨立項目。并且,在機器人和無人機領域,有過豐富的研發經驗。對于圖像識別,目標的檢測與跟蹤完成過多種解決方案。曾經參與 GPU 版氣象模式GRAPES,是其主要研發者。
審核編輯:郭婷
-
NVIDIA
+關注
關注
14文章
5026瀏覽量
103284 -
計時器
+關注
關注
1文章
426瀏覽量
32758
發布評論請先 登錄
相關推薦
評論