NVIDIA DOCA GPUNetIO是 NVIDIA DOCA SDK 中的一個庫，專門為實時內聯 GPU 數據包處理而設計。它結合了GPUDirect RDMA和GPUDirect Async等技術，能夠創建以 GPU 為中心的應用程序，其中 CUDA 內核可以直接與網絡接口卡（NIC）通信，用于發送和接收數據包，繞過 CPU 并將其排除在關鍵路徑之外。
DOCA GPUNetIO 的核心原理和用途已在前幾篇文章《Inline GPU Packet Processing with NVIDIA DOCA GPUNetIO》和《Realizing the Power of Real-Time Network Processing with NVIDIA DOCA GPUNetIO》以及DOCA GPUNetIO 編程指南中進行了討論。
此前，DOCA GPUNetIO與DOCA Ethernet和DOCA Flow一起，僅限于處理以太網傳輸層上的數據包傳輸。隨著 DOCA 2.7 的推出，現在有一組擴展的 API使 DOCA GPUNetIO 能夠直接從 GPU CUDA 內核使用 RoCE 或 InfiniBand 傳輸層支持 RDMA 通信。
這篇文章探討了由 GPU CUDA 內核和 DOCA GPUNetIO 控制的新的遠程直接內存訪問（RDMA）功能，并對其與性能測試（perftest）微基準的性能進行了比較。
注意，RDMA 首字母縮寫描述了一種協議，該協議允許從一臺計算機的存儲器到另一臺計算機存儲器的遠程直接存儲器訪問，而不涉及任何一臺計算機中的操作系統。操作示例包括 RDMA 寫入和 RDMA 讀取。不能將其與GPUDirect RDMA混淆，后者與 RDMA 協議無關。GPUDirect RDMA 是 NVIDIA 在 GPUDirect 技術家族中啟用的技術之一，使網卡能夠繞過 CPU 內存副本和操作系統例程，直接發送或接收訪問 GPU 內存的數據。GPUDirect RDMA 可以由任何使用以太網、InfiniBand 或 RoCE 的網絡框架啟用。
具有 GPUNetIO 的 RDMA GPU 數據路徑
RDMA 提供了在兩個主機的主內存之間的直接訪問，而不涉及操作系統、緩存或存儲。這使得數據傳輸具有高吞吐量、低延遲和低 CPU 利用率。這是通過注冊并共享本地內存區域，以便遠程主機知道如何訪問它。
兩個對等方需要通過 RDMA 交換數據的應用程序通常遵循三個基本步驟：
- 步驟 1–本地配置：每個對等端在本地創建 RDMA 隊列和內存緩沖區，以便與其他對等端共享這些資源。
- 步驟 2–交換信息： 使用帶外（OOB）機制（例如，Linux 套接字），對等端交換有關 RDMA 隊列和要遠程訪問的內存緩沖區的信息。
- 步驟 3–數據路徑：兩個對等方使用遠程內存地址執行 RDMA 讀、寫、發送和接收，以交換數據。
DOCA RDMA 庫按照上面列出的三個步驟通過 InfiniBand 或 RoCE 實現 RDMA 通信，所有這些步驟都是用 CPU 執行的。隨著新GPUNetIO RDMA功能的引入，應用程序可以在 GPU 上執行步驟 3，使用 CUDA 內核管理 RDMA 應用程序的數據路徑，而步驟 1 和 2 保持不變，因為它們與 GPU 數據路徑無關。
將 RDMA 數據路徑移動到 GPU 上的好處與以太網用例中的好處相同。在數據處理發生在 GPU 上的網絡應用程序中，將網絡通信從 CPU 卸載到 GPU，使其能夠成為應用程序的主控制器，消除與 CPU 交互所需的額外延遲，知道數據何時準備就緒以及數據位于何處，這也釋放了 CPU 周期。此外，GPU 可以同時并行管理多個 RDMA 隊列，例如，每個 CUDA 塊可以在不同的 RDMA 隊列上發布 RDMA 操作。
IB Verbs 和 DOCA GPUNetIO 性能測試
在 DOCA 2.7 中，引入了一個新的 DOCA GPUNetIO RDMA 客戶機-服務器代碼示例，以顯示新 API 的使用情況并評估其正確性。這篇文章分析了 GPUNetIO RDMA 函數與 IB Verbs RDMA 函數之間的性能比較，重現了眾所周知的 perftest 套件中的一個微基準。
簡而言之，perftest 是一組微基準點，用于使用基本的 RDMA 操作測量 RDMA 帶寬（BW）和兩個對等點（服務器和客戶端）之間的延遲盡管網絡控制部分發生在 CPU 中，但可以通過啟用 GPUDirect RDMA 并指定--use_cuda標志來指定數據是否駐留在 GPU 內存中。
一般來說，RDMA 寫單向 BW 基準測試（即 ib_write_bw）在每個 RDMA 隊列上發布一個針對相同大小消息的寫請求列表，用于固定迭代次數，并命令 NIC 執行發布的寫操作，這就是所謂的“按門鈴”程序。為了確保所有寫入都已發出，在進入下一次迭代之前，它輪詢完成隊列，等待每個寫入都已正確執行的確認。然后，對于每個消息大小，可以檢索發布和輪詢所花費的總時間，并以 MB/s 為單位計算 BW。
圖 1 顯示了 IB 謂詞ib_write_bw性能測試主循環。在每次迭代中，CPU 發布一個 RDMA 寫入請求列表，命令 NIC 執行這些請求（按門鈴），然后等待完成后移動到下一次迭代。啟用 CUDA 標志后，要寫入的數據包將從 GPU 內存本地獲取，而不是從 CPU 內存。
ib_write_bw性能測試主回路實驗是用 DOCA 庫復制ib_write_bw微基準標記，使用 DOCA RDMA 作為 CPU 上的控制路徑以建立客戶端-服務器連接，并使用 DOCA GPUNetIO RDMA 作為數據路徑，在 CUDA 內核內發布寫入。這種比較并不完全一致，因為 perftest 使用 GPUDirect RDMA 來傳輸數據，但網絡通信由 CPU 控制，而 DOCA GPUNetIO 同時使用 GPUDirect RDMA 和 GPUDirect Async 來控制網絡通信和來自 GPU 的數據傳輸。目標是證明 DOCA GPUNetIO RDMA 性能與 IB Verbs 性能測試相當，后者被視為基線。
為了重現ib_write_bw數據路徑并測量針對每個消息大小發布 RDMA 寫入操作所花費的時間，CPU 記錄一個 CUDA 事件，啟動rdma_write_bw CUDA 內核，然后記錄第二個 CUDA 事件。這應該可以很好地近似 CUDA 內核使用 DOCA GPUNetIO 函數發布 RDMA 寫入所用的時間（以毫秒為單位），如下面的代碼段 1 所示。
Int msg_sizes[MAX_MSG] = {....};for (int msg_idx = 0; msg_idx < MAX_MSG; msg_idx++) {?????????do_warmup();?????????cuEventRecord(start_event, stream);?????????rdma_write_bw<<<num_queue, msg_per_size, 0, stream>>>(msg_sizes[msg_idx], …);?????????cuEventRecord(end_event, stream);?????????cuEventSynchronize(end_event);?????????cuEventElapsedTime(&total_ms, start_event, end_event);?????????calculate_result(total_ms, msg_sizes[msg_idx], …)} |
在下面的代碼段 2 中，CUDA 內核rdma_write_bw，按照弱模式使用 DOCA GPUNetIO 設備函數，對于給定數量的迭代，并行發布一系列 RDMA 寫入，每個 CUDA 塊中的 CUDA 線程發布一個寫操作，按照弱模式。
__global__ void rdma_write_bw(struct doca_gpu_dev_rdma *rdma_gpu,???????????????const int num_iter, const size_t msg_size,???????????????const struct doca_gpu_buf_arr *server_local_buf_arr,???????????????const struct doca_gpu_buf_arr *server_remote_buf_arr){???struct doca_gpu_buf *remote_buf;???struct doca_gpu_buf *local_buf;???uint32_t curr_position;???uint32_t mask_max_position;?
???doca_gpu_dev_buf_get_buf(server_local_buf_arr, threadIdx.x, &local_buf);???doca_gpu_dev_buf_get_buf(server_remote_buf_arr, threadIdx.x, &remote_buf);?
???for (int iter_idx = 0; iter_idx < num_iter; iter_idx++) {???????doca_gpu_dev_rdma_get_info(rdma_gpu, &curr_position, &mask_max_position);???????doca_gpu_dev_rdma_write_weak(rdma_gpu,???????????????????remote_buf, 0,???????????????????local_buf, 0,???????????????????msg_size, 0,???????????????????DOCA_GPU_RDMA_WRITE_FLAG_NONE,???????????????????(curr_position + threadIdx.x) & mask_max_position);???????/* Wait all CUDA threads to post their RDMA Write */???????__syncthreads();?
???????if (threadIdx.x == 0) {???????????/* Only 1 CUDA thread can commit the writes in the queue to execute them */???????????doca_gpu_dev_rdma_commit_weak(rdma_gpu, blockDim.x);????????????????/* Only 1 CUDA thread can flush the RDMA queue waiting for the actual execution of the writes */???????doca_gpu_dev_rdma_flush(rdma_gpu);???????}???????__syncthreads();???}?
???return;} |
圖 2 描述了代碼段 2。在每次迭代時，GPU CUDA 內核并行發布一個 RDMA 寫入請求列表，每個 CUDA 塊中的 CUDA 線程一個。在同步所有 CUDA 線程后，只有線程 0 命令 NIC 執行寫入并等待完成，然后刷新隊列，最后再進行下一次迭代。
為了比較性能，為 IB Verbs perftest 和 DOCA GPUNetIO perftest 設置了相同的參數：1 個 RDMA 隊列，2048 次迭代，每次迭代執行 512 次 RDMA 寫入，并測試消息大小從 64 字節到 4096 字節。
已在 Dell R750 機器上執行基準測試，該機器配備NVIDIA H100 GPU和ConnectX-7網卡（RoCE 模式），通過系統 PCIe 總線連接（無專用 PCIe 交換機）。如下圖所示，perftest 的兩種實現所花費的總時間是完全可比較的（圖 3），以及以 MB/s 為單位報告的峰值 BW（圖 4）。由于代碼中不同邏輯的性質，時間和 BW 是用不同的方法來測量的，IB Verbs perftest 使用系統時鐘，而 DOCA GPUNetIO perftest 則依賴于 CUDA 事件，后者可能具有不同的內部時間測量開銷。
請注意，像 perftest 這樣的應用程序并不是顯示 GPU 利用率優勢的最佳工具，因為可實現的并行化量非常低。DOCA GPUNetIO 性能測試 RDMA 寫入是以并行方式發布在隊列中的（512 個寫入，每個寫入由不同的 CUDA 線程執行），發布所花費的時間微不足道，約 4 微秒。大部分性能測試時間是由 NIC 實際執行 RDMA 寫入、通過網絡發送數據和返回正反饋所花費。
這個實驗可以被認為是成功的，因為它證明了使用 DOCA GPUNetIO RDMA API 與使用常規 IB Verbs 相比不會引入任何相關開銷，并且在運行相同類型的工作負載和工作流時可以達到性能目標。ISV 開發人員和最終用戶可以使用 DOCA GPUNetIO RDMA，獲得 GPUDirect 異步技術將通信控制卸載到 GPU 的好處。
這種架構選擇提供了以下好處：
- 一個可擴展性更強的應用程序，能夠同時并行管理多個 RDMA 隊列（通常每個 CUDA 塊一個隊列）。
- 利用 GPU 提供的高度并行性的能力，該 GPU 擁有多個 CUDA 線程在不同數據上并行工作，并能以盡可能低的延遲在同一隊列上發布 RDMA 操作。
- ?
- 更少的內部總線事務（例如 PCIe），因為不需要同步 GPU 上的工作與 CPU 活動。CPU 不再負責發送或接收 GPU 必須處理的數據。
準備好深入了解 DOCA GPUNetIO 了嗎？查看官方 DOCA GPUNetIO 程序員指南 以獲取全面的見解和技術細節。
要了解 DOCA GPUNetIO 的更多功能，請參閱使用 NVIDIA DOCA GPUNetIO 的內聯 GPU 數據包處理。要了解 DOCA GPUNetIO 如何徹底改變 GPU 通信，請參閱使用 NVIDIA DOCA GPUNetIO 實現實時網絡處理的強大功能。
?
