加速 Pony AV 傳感器數據處理流水線

就像人類依靠眼睛看東西一樣，自動駕駛汽車使用傳感器收集信息。這些傳感器收集大量數據，這需要高效的車載數據處理，以便車輛對道路上的情況做出快速反應。這種能力對于自動駕駛汽車的安全至關重要，對于讓虛擬駕駛員更智能也至關重要。

由于需要冗余和多樣的傳感器和計算系統，設計和優化處理管道是一項挑戰。在這篇文章中，我們將介紹Pony.ai 的演變。人工智能的車載傳感器數據處理管道。

Pony.ai 能的傳感器設置包括多個攝像頭、激光雷達和雷達。上游模塊同步傳感器，將數據封裝成消息片段，并將其發送給下游模塊，這些模塊使用它們來分割、分類和檢測對象，等等

每種類型的傳感器數據可能有多個模塊，用戶算法可以是傳統的，也可以是基于神經網絡的。

整個管道必須以最高的效率運行。乘客的安全是我們的首要任務。傳感器數據處理系統在兩個方面影響安全。

首先，安全性的決定因素之一是自動駕駛系統處理傳感器數據的速度。如果感知和定位算法以數百毫秒的延遲獲得傳感器數據，那么車輛做出的決定就太晚了。

第二，整個硬件/軟件系統必須可靠，才能長期成功。如果自動駕駛汽車在制造幾個月后開始出現問題，消費者將永遠不會想要購買或乘坐該車。這在大規模生產階段至關重要。

高效處理傳感器數據

緩解傳感器處理管道中的瓶頸需要采用多方面的方法，同時考慮傳感器、 GPU 體系結構和 GPU 內存。

從傳感器到 GPU

當 Pony.ai 成立時，我們最初的傳感器設置由現成的組件組成。我們為攝像頭使用了基于 USB 和以太網的型號，攝像頭直接連接到車載計算機， CPU 負責從 USB /以太網接口讀取數據。

通過以太網/ USB 傳輸攝像頭數據提供了更高的延遲，但需要 CPU 個周期。

雖然功能良好，但設計存在一個根本問題。 USB 和以太網攝像頭接口（ GigE 攝像頭）占用 CPU 。隨著越來越多的高分辨率攝像頭的加入， CPU 很快變得不知所措，無法執行所有的 I / O 操作。這種設計很難在保持足夠低的延遲的同時實現可擴展性。

我們通過為攝像頭和激光雷達添加基于 FPGA 的傳感器網關解決了這個問題。

添加 FPGA 作為傳感器網關 可以降低 CPU I / O 成本，但 PCIe 上的 DMA 具有更高的帶寬和更低的延遲，因此它可以支持更多攝像頭。

FPGA 處理攝像頭觸發和同步邏輯，以提供更好的傳感器融合。當一個或多個攝像頭數據包準備就緒時，觸發 DMA 傳輸，以通過 PCIe 總線將數據從 FPGA 復制到主存儲器。 DMA 引擎在 FPGA 上執行此操作， CPU 被釋放。它不僅可以打開 CPU 的 I / O 資源，還可以減少數據傳輸延遲，從而實現更具可擴展性的傳感器設置。

由于在 GPU 上運行的許多神經網絡模型都使用相機數據，因此在通過 DMA 將其從 FPGA 傳輸到 CPU 后，仍然必須將其復制到 GPU 內存中。因此，某個地方需要 CUDA HostToDevice內存拷貝， FHD 相機圖像的單幀需要約 1.5 毫秒。

然而，我們希望進一步減少這種延遲。理想情況下，攝像機數據應直接傳輸到 GPU 內存中，而無需通過 CPU 路由。

使用 FPGA- GPU RDMA ，我們添加了一個 PCIe 交換機以獲得最佳性能。該解決方案消除了 CPU- GPU 數據拷貝。我們還將 NVIDIA GPU 直接 RDMA 技術集成到 Xilinx 的 XDMA 驅動程序中，在 PCIe Gen3 x8 上提供約 6 GB / s 的 FPGA- GPU 帶寬。

我們通過使用 NVIDIA GPU 直接 RDMA 實現了這一目標。 GPU 直接 RDMA 使我們能夠通過 PCIe 條（定義 PCIe 地址空間線性窗口的基址寄存器）預先分配可供 PCIe 對等方訪問的 CUDA 內存塊。

它還為第三方設備驅動程序提供了一系列內核空間 API ，以獲取 GPU 內存物理地址。這些 API 有助于第三方設備中的 DMA 引擎直接向 GPU 內存發送和讀取數據，就像它向主存發送和讀取數據一樣。

GPU 直接 RDMA 通過消除 CPU 到 – GPU 拷貝來減少延遲，并在 PCIe Gen3 x8 下實現最高帶寬~ 6 GB / s ，其理論限制為 8 GB / s 。

跨 GPU 縮放

由于計算工作量的增加，我們需要不止一個 GPU 。隨著越來越多的 GPU 添加到系統中， GPU 之間的通信也可能成為瓶頸。通過暫存緩沖區通過 CPU 會增加 CPU 成本并限制總體帶寬。

我們添加了一個 PCIe 交換機，它提供了盡可能最好的對等傳輸性能。在我們的測量中，對等通信可以達到 PCIe 線速度，從而在多個 GPU 之間提供更好的可擴展性。

將計算轉移到專用硬件

我們還將以前在 CUDA 內核上運行的任務卸載到專用硬件上，以加速傳感器數據處理。

例如，將 FHD 攝像頭圖像編碼為 JPEG 字符串時， NvJPEG 庫在帶有 RTX5000 GPU 的單個 CPU 線程上需要約 4ms 的時間。 NvJPEG 可能會消耗 CPU 和 GPU 資源，因為它的某些階段，比如哈夫曼編碼，可能完全在 CPU 上。

GPU 上的 JPEG 編碼會消耗 GPU 和 CPU 的資源。 NvJPEG GPU 編碼仍有一些相位在 CPU 上計算（哈夫曼編碼和解碼）。一個 CPU 線程加上 NVIDIA RTX 5000 需要約 4ms 才能將一幅 FHD 圖像編碼為 JPEG 。這也會干擾其他正在運行的 GPU 任務。

我們采用 NVIDIA 車載視頻編解碼器來減輕 CPU 和 GPU （ CUDA 部分）進行圖像編碼和解碼的負擔。該編解碼器在 GPU 的專用部分使用編碼器。它是 GPU 的一部分，但與用于運行內核和深度學習模型的其他 CUDA 資源不沖突。

我們還通過使用 NVIDIA GPU 上的專用硬件視頻編碼器，將圖像壓縮格式從 JPEG 遷移到了 HEVC （ H.265 ）。我們實現了編碼速度的提高，并將 CPU 和 GPU 資源釋放出來用于其他任務。

在 GPU 上完全編碼 FHD 圖像需要約 3 毫秒，而不會影響其 CUDA 性能。性能是在僅 I 幀模式下測量的，這確保了幀間一致的質量和壓縮偽影。

NVIDIA 視頻編解碼器在 GPU 芯片的專用分區中使用編碼器，不消耗 CUDA 內核或 CPU 。 NVENC 支持 H264 / H265 。將一幅 FHD 圖像編碼到 HEVC 需要約 3 毫秒，因此 GPU 和 CPU 可以自由地執行其他任務。我們使用 I-frame-only 模式來確保每個幀具有相同的質量和相同類型的工件。

關于 – GPU 數據流

另一個關鍵主題是將相機幀作為消息發送到下游模塊的效率。

我們使用谷歌的 protobuf 來定義一條消息。以CameraFrame消息為例。相機規格和屬性是消息中的基本類型。由于 protobuf 的限制，實際有效載荷攝像頭數據必須定義為主系統內存中的字節字段。

下面代碼示例中的消息是 proto 。由于protobuf的限制，數據必須在主存中。

message CameraFrame {
 optional string device_name = 1;
 optional int32 width = 2;
 optional int32 height = 3;
 optional int32 pixel_format = 4;
 optional bytes data = 5;
};

需要 CUDA H2D 副本，才能使 DL 型號獲得攝像頭數據。

我們使用發布者 – 訂閱者模型，在模塊之間傳遞零拷貝消息以共享信息。此CameraFrame消息的許多訂戶模塊使用攝像頭數據進行深入學習推斷。

在最初的設計中，當這樣的模塊收到消息時，它必須調用 CUDA HostToDevice內存副本，以便在推斷之前將相機數據傳輸到 GPU 。

每個模塊都必須對相同的數據進行 H2D 復制，這很耗時！下面的代碼示例顯示在 CPU 上傳遞的零拷貝消息，但不在 GPU 上傳遞。

message CameraFrame {
 optional string device_name = 1;
 optional int32 width = 2;
 optional int32 height = 3;
 optional int32 pixel_format = 4;
 optional bytes data = 5;
};

每個模塊都必須進行 CUDA HostToDevice復制，這是冗余的，而且會消耗資源。盡管零拷貝消息傳遞框架在 CPU 上運行良好，但它涉及大量 CPU- GPU 數據拷貝。

我們使用 protobuf codegen插件來啟用 GPU 內存中的數據字段。下面的代碼示例顯示了在 GPU 上傳遞的零拷貝消息。GPUData字段位于 GPU 內存中。

message CameraFrame {
 optional string device_name = 1;
 optional int32 width = 2;
 optional int32 height = 3;
 optional int32 pixel_format = 4;
 optional GpuData data = 5;
};

我們通過 protobuf 的插件 API 向 protobuf 代碼生成器中添加了一種新類型的數據GpuData字段，從而解決了這個問題。GpuData支持標準的resize操作，就像 CPU 內存bytes字段一樣。然而，它的物理數據存儲在 GPU 上。

當用戶模塊收到消息時，它們可以檢索 GPU 數據指針以供直接使用。因此，我們在整個管道中實現了完全零拷貝。

改進 GPU 內存分配

當我們調用GpuData原型的resize函數時，它調用 CUDA cudaMalloc。當GpuData原型消息被銷毀時，它會調用cudaFree。

這兩個 API 操作并不便宜，因為它們必須修改 GPU 的內存映射。每次通話可能需要約 0.1 毫秒。

由于在攝像機不間斷地生成數據時，該協議被廣泛使用，因此我們應該優化 GPU 協議的alloc/free成本。

我們實現了一個固定插槽大小的 GPU 內存來解決這個問題。想法很簡單：我們維護一堆預先分配的 GPU 內存插槽，這些插槽與我們想要的相機數據幀緩沖區大小相匹配。每次調用alloc時，我們從堆棧中取出一個插槽。每次調用free時，插槽都會返回到池中。通過重新使用 GPU 內存，alloc/free時間接近于零。

camera_frame.mutable_gpu_data()->Resize(size);
ptr = pool->Alloc();

如果我們想支持不同分辨率的攝像頭呢？使用這個固定大小的內存池，我們必須始終分配最大的內存池大小，或者使用不同的插槽大小初始化多個內存池。兩者都會降低效率。

CUDA 11.2 中的新功能解決了這個問題。它正式支持cudaMemPool，可以預先分配，然后用于cudaMalloc和free。與我們之前的實現相比，它有助于任何分配大小。這大大提高了靈活性，但性能成本很低（每次分配約 2 個）。

camera_frame.mutable_gpu_data()->Resize(size);
pool->cudaMallocFromPoolAsync(&ptr, pool, ...);

在這兩種方法中，當內存池溢出時，resize調用會返回到傳統的cudaMalloc和free。

YUV 顏色空間中更干凈的數據流

通過之前對硬件設計和系統軟件架構的所有優化，我們實現了高效的數據流。下一步是優化數據格式本身。

我們的系統用于在 RGB 顏色空間中處理相機數據。然而，我們相機的 ISP 輸出在 YUV 顏色空間中，并且在 GPU 上執行從 YUV 到 RGB 的轉換，這需要約 0.3 毫秒。此外，一些感知組件不需要顏色信息。向它們提供 RGB 顏色像素是浪費。

由于這些原因，我們從 RGB 相機幀遷移到 YUV 幀。我們選擇使用 YUV420 像素格式，因為人類視覺對色度信息的敏感度不如對亮度信息的敏感度。

通過采用 YUV420 像素格式，我們節省了 GPU 內存消耗的一半。這也使我們能夠僅向感知組件發送 Y 通道，感知組件不需要色度信息，與 RGB 相比，節省了 GPU 內存消耗的三分之二。

在 GPU 上處理激光雷達數據

除了相機數據，我們還處理激光雷達數據，這是更稀疏的，主要是在 GPU 上。考慮到不同類型的激光雷達，處理起來更困難。我們在處理激光雷達數據時進行了幾次優化：

• 由于激光雷達掃描數據包含大量的物理信息，我們使用 GPU 友好的陣列結構而不是結構陣列來描述點云，使 GPU 內存訪問模式更緊密而不是分散。
• 當一些字段必須在 CPU 和 GPU 之間交換時，我們將它們保存在頁面鎖定內存中，以加速傳輸。
• NVIDIA CUB 庫廣泛用于我們的處理管道，特別是掃描/選擇操作。

GPU 上的激光雷達數據處理管道產生以下結果：

• 一種 GPU 友好的陣列數據布局結構。
• 頁面鎖定內存以加速 CPU- GPU 傳輸。
• NVIDIA CUB 庫快速例程，過濾速度快約 58% 。

通過所有這些優化，我們在關鍵路徑中將整個管道延遲減少了約 4 毫秒。

總體時間表

通過所有這些優化，我們可以使用內部時間線可視化工具查看系統跟蹤。

總體時間線顯示了我們今天對 GPU 的依賴程度。雖然這兩個 GPU 在大約 80% 的時間內都在使用，但 GPU0 和 GPU1 的工作負載并不理想。對于 GPU 0 ，它在整個 perception module 迭代中被大量使用。對于 GPU 1 ，它在迭代的中間有更多的空閑間隙。

未來，我們將專注于進一步提高 GPU 效率。

生產準備

在開發的早期， FPGA 使我們能夠輕松地在基于硬件的傳感器數據處理中試驗我們的想法。隨著我們的傳感器數據處理器變得越來越成熟，我們一直在研究使用片上系統（ SoC ）來提供緊湊、可靠、可生產的傳感器數據處理器的可能性。

我們發現汽車級 NVIDIA DRIVE Orin SoC 完全符合我們的要求。它是 ASIL 等級，非常適合在生產車輛上運行。

從 FPGA 遷移到 NVIDIA DRIVE Orin

在開發的早期， FPGA 使我們能夠輕松地在基于硬件的傳感器數據處理中試驗我們的想法。

隨著我們的傳感器數據處理器變得越來越成熟，我們一直在研究使用片上系統（ SoC ）來提供緊湊、可靠、可生產的傳感器數據處理器的可能性。

我們發現汽車級 NVIDIA DRIVE Orin SoC 完全符合我們的要求。它是 ASIL 等級，非常適合在生產車輛上運行。盡管其體積小、成本低，但它可以連接到各種汽車級傳感器，并有效地處理大規模傳感器數據。

我們將使用 NVIDIA Orin 來處理所有傳感器信號處理、同步、數據包收集以及相機幀編碼。我們估計，這種設計，結合其他架構優化，將節省約 70% 的 BOM 總成本。

Orin SoC 取代 FPGA 作為傳感器網關，支持 10 多個攝像頭和激光雷達，這是汽車級的，成本降低約 70% 。

在與 NVIDIA 的合作中，我們確保 Orin CPU- GPU 組件之間的所有通信都通過 PCIe 總線，并通過 NvStreams 支持 DMA 。

• 對于計算密集型 DL 工作， NVIDIA Orin SoC 使用 NvStream 將傳感器數據傳輸到離散的 GPU 進行處理。
• 對于非 GPU 工作， NVIDIA Orin SoC 使用 NvStream 將數據傳輸到主機 CPU 進行處理。

2 / 3 級計算平臺應用程序

L4 和 L2 / L3 之間的差距可以縮小，因為 Orin SoC 提供高達 250 個頂級 DL 性能。模型的執行性能與 NVIDIA RTX 5000 類似，因此 L4 場景中使用的模型可以被刪減，以合理的速度在 Orin 上運行。靈活的體系結構設計可以在 L4 和 L2 / L3 解決方案之間共享資源。

這種設計的一個顯著優點是，它有潛力被用作 L2 / L3 計算平臺。

NVIDIA Orin 每秒提供 254 萬億次運算，其計算能力可能與我們目前的 4 級自主車輛計算平臺上使用的 RTX5000 離散式 GPU 類似。然而，要充分發揮 NVIDIA Orin SoC 的潛力，需要進行多項優化，例如：

• 結構稀疏網絡
• 對于核心
• 跨多個 NVIDIA Orin SOC 擴展

結論

Pony 傳感器數據處理管道的演變證明了我們朝著高效數據處理管道和增強系統可靠性的方向發展的系統方法，這有助于實現更高的安全目標。這種方法背后的簡單想法是：

• 使數據流簡單流暢。數據應該以最小化轉換開銷的格式直接傳輸到將被使用的位置。
• 為計算密集型任務使用專用硬件，并為其他任務節省通用計算資源。
• 4 級和 2 級系統之間的資源共享提高了可靠性并節省了工程成本。

硬件和軟件不能單獨通過共同努力來實現。我們認為，這對于滿足快速增長的計算需求和生產預期至關重要。

致謝

這篇文章包括 Pony.ai 傳感器數據管道多年的發展。這項榮譽屬于多個團隊和工程師，他們一直在為開發這種高效的傳感器數據管道做出貢獻。

?