使用 GPU 對網(wǎng)絡(luò)數(shù)據(jù)包進(jìn)行內(nèi)聯(lián)處理是一種數(shù)據(jù)包分析技術(shù)，可用于許多不同的應(yīng)用領(lǐng)域：信號處理、網(wǎng)絡(luò)安全、信息收集、輸入重建等。

這些應(yīng)用程序類型的主要要求是盡快將接收到的數(shù)據(jù)包移動到 GPU 內(nèi)存中，以觸發(fā)負(fù)責(zé)對其執(zhí)行并行處理的 CUDA 內(nèi)核。

總體思路是創(chuàng)建一個連續(xù)的異步管道，能夠?qū)?shù)據(jù)包從網(wǎng)卡直接接收到 GPU 內(nèi)存中。您還可以使用 CUDA 內(nèi)核來處理傳入的數(shù)據(jù)包，而無需同步 GPU 和 CPU 。

有效的應(yīng)用程序工作流包括使用無鎖通信機(jī)制在以下播放器組件之間創(chuàng)建一個協(xié)調(diào)的連續(xù)異步管道：

network controller 向 GPU 內(nèi)存提供接收到的網(wǎng)絡(luò)數(shù)據(jù)包

CPU 用于查詢網(wǎng)絡(luò)控制器以獲取有關(guān)接收到的數(shù)據(jù)包的信息

GPU 用于接收數(shù)據(jù)包信息并直接將其處理到 GPU 內(nèi)存中

圖 1 顯示了使用 NVIDIA GPU 和 ConnectX 網(wǎng)卡的加速內(nèi)聯(lián)數(shù)據(jù)包處理應(yīng)用程序的典型數(shù)據(jù)包工作流場景。

圖 1 。典型的內(nèi)聯(lián)數(shù)據(jù)包處理工作流場景

在這種情況下，避免延遲是至關(guān)重要的。不同組件之間的通信越優(yōu)化，系統(tǒng)的響應(yīng)速度就越快，吞吐量也就越高。每一步都必須在所需資源可用時以內(nèi)聯(lián)方式進(jìn)行，而不會阻塞任何其他等待的組件。

您可以清楚地識別兩種不同的流：

Data flow ：通過 PCIe 總線在網(wǎng)卡和 GPU 之間交換優(yōu)化的數(shù)據(jù)（網(wǎng)絡(luò)數(shù)據(jù)包）。

Control flow ： CPU 協(xié)調(diào) GPU 和網(wǎng)卡。

數(shù)據(jù)流

關(guān)鍵是優(yōu)化網(wǎng)絡(luò)控制器和 GPU 之間的數(shù)據(jù)移動（發(fā)送或接收數(shù)據(jù)包）。它可以通過 GPUDirect RDMA 技術(shù)實(shí)現(xiàn)，該技術(shù)使用 PCI Express 總線接口的標(biāo)準(zhǔn)功能，在 NVIDIA GPU 和第三方對等設(shè)備（如網(wǎng)卡）之間實(shí)現(xiàn)直接數(shù)據(jù)路徑。

GPUDirect RDMA 依賴于 NVIDIA GPU 在 PCI Express 基址寄存器（ BAR ）區(qū)域上公開部分設(shè)備內(nèi)存的能力。有關(guān)更多信息，請參閱 CUDA 工具包文檔中的 使用 GPUDirect-RDMA 開發(fā) Linux 內(nèi)核模塊 。 在現(xiàn)代服務(wù)器平臺上對 GPUDirect-RDMA 進(jìn)行基準(zhǔn)測試 文章對使用不同系統(tǒng)拓?fù)涞臉?biāo)準(zhǔn) IB 謂詞執(zhí)行網(wǎng)絡(luò)操作（發(fā)送和接收）時的 GPUDirect RDMA 帶寬和延遲進(jìn)行了更深入的分析。

圖 2 。 NVIDIA GPUDirect RDMA 使用 PCI Express 的標(biāo)準(zhǔn)功能，為 GPU 和第三方對等設(shè)備之間的數(shù)據(jù)交換提供了直接路徑

要在 Linux 系統(tǒng)上啟用 GPUDirect RDMA ，需要 nvidia-peermem 模塊（在 CUDA 11.4 及更高版本中提供）。圖 3 顯示了最大化 GPUDirect RDMA 內(nèi)部吞吐量的理想系統(tǒng)拓?fù)洌涸?GPU 和 NIC 之間使用專用 PCIe 交換機(jī)，而不是通過與其他組件共享的系統(tǒng) PCIe 連接。

圖 3 。理想的拓?fù)浣Y(jié)構(gòu)，最大限度地提高網(wǎng)絡(luò)控制器和 GPU 之間的內(nèi)部數(shù)據(jù)吞吐量

控制流

CPU 是網(wǎng)絡(luò)控制器和 GPU 之間協(xié)調(diào)和同步活動的主要參與者，用于喚醒 NIC ，將數(shù)據(jù)包接收到 GPU 內(nèi)存中，并通知 CUDA 工作負(fù)載有新數(shù)據(jù)包可供處理。

在處理 GPU 時，強(qiáng)調(diào) CPU 和 GPU 之間的異步非常重要。例如，假設(shè)一個簡單的應(yīng)用程序在主循環(huán)中執(zhí)行以下三個步驟：

接收數(shù)據(jù)包。

處理數(shù)據(jù)包。

發(fā)回修改過的數(shù)據(jù)包。

在本文中，我將介紹在這種應(yīng)用程序中實(shí)現(xiàn)控制流的四種不同方法，包括優(yōu)缺點(diǎn)。

方法 1

圖 4 顯示了最簡單但最不有效的方法：單個 CPU 線程負(fù)責(zé)接收數(shù)據(jù)包，啟動 CUDA 內(nèi)核來處理它們，等待 CUDA 內(nèi)核完成，并將修改后的數(shù)據(jù)包發(fā)送回網(wǎng)絡(luò)控制器。

圖 4 。單個 CPU 將數(shù)據(jù)包傳遞到 CUDA 內(nèi)核并等待完成以執(zhí)行下一步的工作流

如果數(shù)據(jù)包處理不是那么密集，那么這種方法的性能可能會比只使用 CPU 處理數(shù)據(jù)包而不使用 GPU 更差。例如，您可能具有高度的并行性來解決數(shù)據(jù)包上的一個困難且耗時的算法。

方法 2

在這種方法中，應(yīng)用程序?qū)?CPU 工作負(fù)載分成兩個 CPU 線程：一個用于接收數(shù)據(jù)包并啟動 GPU 處理，另一個用于等待 GPU 處理完成并通過網(wǎng)絡(luò)傳輸修改后的數(shù)據(jù)包（圖 5 ）。

圖 5 。拆分 CPU 線程以通過 GPU 處理數(shù)據(jù)包

這種方法的一個缺點(diǎn)是，每次累積數(shù)據(jù)包的突發(fā)都會啟動一個新的 CUDA 內(nèi)核。 CPU 必須為每次迭代支付 CUDA 內(nèi)核啟動延遲。如果 GPU 被淹沒，數(shù)據(jù)包處理可能不會立即執(zhí)行，從而導(dǎo)致延遲。

方法 3

圖 6 顯示了第三種方法，它涉及使用 CUDA 持久內(nèi)核。

圖 6 。使用持久 CUDA 內(nèi)核進(jìn)行內(nèi)聯(lián)數(shù)據(jù)包處理。

CUDA 持久內(nèi)核是一個預(yù)啟動的內(nèi)核，它正忙著等待來自 CPU 的通知：新數(shù)據(jù)包已經(jīng)到達(dá)并準(zhǔn)備好進(jìn)行處理。當(dāng)數(shù)據(jù)包準(zhǔn)備好后，內(nèi)核通知第二個 CPU 線程它可以向前發(fā)送數(shù)據(jù)包。

實(shí)現(xiàn)此通知系統(tǒng)的最簡單方法是使用忙等待標(biāo)志更新機(jī)制在 CPU 和 GPU 之間共享一些內(nèi)存。雖然 GPUDirect RDMA 旨在從第三方設(shè)備直接訪問 GPU 內(nèi)存，但您可以使用這些 API 創(chuàng)建 GPU 內(nèi)存的完全有效的 CPU 映射。 CPU 驅(qū)動的拷貝的優(yōu)點(diǎn)是所涉及的開銷小?，F(xiàn)在可以通過 GDRCopy 庫啟用此功能。

直接映射 GPU 內(nèi)存進(jìn)行信令，可以從 CPU 修改內(nèi)存，并在輪詢期間降低 GPU 的延遲成本。您也可以將該標(biāo)志放在從 GPU 可見的 CPU 固定內(nèi)存中，但 CUDA 內(nèi)核輪詢 CPU 內(nèi)存標(biāo)志將消耗更多 PCIe 帶寬并增加總體延遲。

這種快速解決方案的問題在于它有風(fēng)險，而且 CUDA 編程模型不支持它。 GPU 內(nèi)核不能被搶占。如果寫得不正確，持久內(nèi)核可能會永遠(yuǎn)循環(huán)。此外，長期運(yùn)行的持久內(nèi)核可能會失去與其他 CUDA 內(nèi)核、 CPU 活動、內(nèi)存分配狀態(tài)等的同步。

它還擁有 GPU 資源（例如，流式多處理器），如果 GPU 真的忙于其他任務(wù)，這可能不是最好的選擇。如果您使用 CUDA 持久內(nèi)核，那么您的應(yīng)用程序必須具有良好的處理能力。

方法 4

最后一種方法是前兩種方法的混合解決方案：使用 CUDA 流內(nèi)存操作 要等待或更新通知標(biāo)志，請?jiān)?CUDA 流上預(yù)啟動一個 CUDA 內(nèi)核，每接收一組數(shù)據(jù)包。

圖 7 。使用模型組合的內(nèi)聯(lián)數(shù)據(jù)包處理的混合方法

這種方法的不同之處在于 GPU HW （使用cuStreamWaitValue）輪詢內(nèi)存標(biāo)志，而不是阻塞 GPU 流式多處理器，并且只有在數(shù)據(jù)包準(zhǔn)備就緒時才會觸發(fā)數(shù)據(jù)包的處理內(nèi)核。

類似地，當(dāng)處理內(nèi)核結(jié)束時，cuStreamWriteValue通知負(fù)責(zé)發(fā)送數(shù)據(jù)包的 CPU 線程數(shù)據(jù)包已被處理。

這種方法的缺點(diǎn)是，應(yīng)用程序必須不時地用cuStreamWriteValue+cuStreamWaitValue內(nèi)核+ CUDA 的新序列重新填充 GPU ，以避免在空流沒有準(zhǔn)備好處理更多數(shù)據(jù)包的情況下浪費(fèi)執(zhí)行時間。這里的 CUDA 圖是在流上重新發(fā)布的好方法。

不同的方法適用于不同的應(yīng)用程序模式。

DPDK 和 GPUdev

數(shù)據(jù)平面開發(fā)工具包 （ DPDK ）是一組庫，用于幫助加速在各種 CPU 體系結(jié)構(gòu)和不同設(shè)備上運(yùn)行的數(shù)據(jù)包處理工作負(fù)載。

在 DPDK 21.11 中， NVIDIA 引入了一個名為 GPUdev 的新庫，以在 DPDK 的上下文中引入 GPU 的概念，并增強(qiáng) CPU 、網(wǎng)卡和 GPU 之間的對話。 GPUdev 在 DPDK 22.03 中擴(kuò)展了更多功能。

圖書館的目標(biāo)如下：

介紹從 DPDK 通用庫管理的 GPU 設(shè)備的概念。

實(shí)現(xiàn)基本的 GPU 內(nèi)存交互，隱藏特定于 GPU 的實(shí)現(xiàn)細(xì)節(jié)。

減少網(wǎng)卡、 GPU 設(shè)備和 CPU 之間的間隙，增強(qiáng)通信。

將 DPDK 集成簡化為 GPU 應(yīng)用程序。

通過通用層公開 GPU 特定于驅(qū)動程序的功能。

對于特定于 NVIDIA 的 GPU ， GPUdev 庫功能通過 CUDA 驅(qū)動程序 DPDK 庫 。要為 NVIDIA GPU 啟用所有g(shù)pudev可用功能， DPDK 必須構(gòu)建在具有 CUDA 庫和 GDRCopy 的系統(tǒng)上。

有了這個新庫提供的功能，您可以輕松地通過 GPU 實(shí)現(xiàn)內(nèi)聯(lián)數(shù)據(jù)包處理，同時處理數(shù)據(jù)流和控制流。

DPDK 在 mempool 中接收數(shù)據(jù)包，這是一個連續(xù)的內(nèi)存塊。通過以下指令序列，您可以啟用 GPUDirect RDMA 在 GPU 內(nèi)存中分配 mempool ，并將其注冊到設(shè)備網(wǎng)絡(luò)中。

struct rte_pktmbuf_extmem gpu_mem; gpu_mem.buf_ptr = rte_gpu_mem_alloc(gpu_id, gpu_mem.buf_len, alignment)); /* Make the GPU memory visible to DPDK */ rte_extmem_register(gpu_mem.buf_ptr, gpu_mem.buf_len, NULL, gpu_mem.buf_iova, NV_GPU_PAGE_SIZE); /* Create DMA mappings on the NIC */ rte_dev_dma_map(rte_eth_devices[PORT_ID].device, gpu_mem.buf_ptr, gpu_mem.buf_iova, gpu_mem.buf_len)); /* Create the actual mempool */ struct rte_mempool *mpool = rte_pktmbuf_pool_create_extbuf(... , &gpu_mem, ...);

圖 8 顯示了 mempool 的結(jié)構(gòu)：

圖 8 用于內(nèi)聯(lián)數(shù)據(jù)包處理的 mempool 結(jié)構(gòu)

對于控制流，要啟用 CPU 和 GPU 之間的通知機(jī)制，可以使用gpudev通信列表：在 CPU 內(nèi)存和 CUDA 內(nèi)核之間的共享內(nèi)存結(jié)構(gòu)。列表中的每一項(xiàng)都可以保存接收到的數(shù)據(jù)包的地址（mbufs），以及一個用于更新處理該項(xiàng)狀態(tài)的標(biāo)志（數(shù)據(jù)包就緒、處理完成等）。

struct rte_gpu_comm_list { /** DPDK GPU ID that will use the communication list. */ uint16_t dev_id; /** List of mbufs populated by the CPU with a set of mbufs. */ struct rte_mbuf **mbufs; /** List of packets populated by the CPU with a set of mbufs info. */ struct rte_gpu_comm_pkt *pkt_list; /** Number of packets in the list. */ uint32_t num_pkts; /** Status of the packets’ list. CPU pointer. */ enum rte_gpu_comm_list_status *status_h; /** Status of the packets’ list. GPU pointer. */ enum rte_gpu_comm_list_status *status_d;
};

偽代碼示例：

struct rte_mbuf * rx_mbufs[MAX_MBUFS]; int item_index = 0; struct rte_gpu_comm_list *comm_list = rte_gpu_comm_create_list(gpu_id, NUM_ITEMS); while(exit_condition) { ... // Receive and accumulate enough packets nb_rx += rte_eth_rx_burst(port_id, queue_id, &(rx_mbufs[0]), rx_pkts); // Populate next item in the communication list. rte_gpu_comm_populate_list_pkts(&(p_v->comm_list[index]), rx_mbufs, nb_rx); ... index++; }

為簡單起見，假設(shè)應(yīng)用程序遵循 CUDA 持久內(nèi)核場景， CUDA 內(nèi)核上的輪詢端看起來類似于以下代碼示例：

__global__ void cuda_persistent_kernel(struct rte_gpu_comm_list *comm_list, int comm_list_entries) { int item_index = 0; uint32_t wait_status; /* GPU kernel keeps checking exit condition as it can’t be preempted. */ while (!exit_condition()) { wait_status = RTE_GPU_VOLATILE(comm_list[item_index].status_d[0]); if (wait_status != RTE_GPU_COMM_LIST_READY) continue; if (threadIdx.x < comm_list[item_index]->num_pkts) { /* Each CUDA thread processes a different packet. */ packet_processing(comm_list[item_index]->addr, comm_list[item_index]->size, ..); } __syncthreads(); /* Notify packets in the items have been processed */ if (threadIdx.x == 0) { RTE_GPU_VOLATILE(comm_list[item_index].status_d[0]) = RTE_GPU_COMM_LIST_DONE; __threadfence_system(); } /* Wait for new packets on the next communication list entry. */ item_index = (item_index+1) % comm_list_entries; } }

圖 10 使用 DPDK 的內(nèi)聯(lián)數(shù)據(jù)包處理用例 gpudev 在空中 5G 軟件中

l2fwd nv 應(yīng)用程序

為了提供如何實(shí)現(xiàn)內(nèi)聯(lián)數(shù)據(jù)包處理和使用 DPDK gpudev庫的實(shí)際示例，l2fwd-nv示例代碼已在 /NVIDIA/l2fwd-nv GitHub repo 上發(fā)布。這是使用 GPU 功能增強(qiáng)的普通 DPDK l2fwd示例的擴(kuò)展。應(yīng)用程序布局是接收數(shù)據(jù)包，交換每個數(shù)據(jù)包的 MAC 地址（源和目的地），并傳輸修改后的數(shù)據(jù)包。

L2fwd-nv為本文討論的所有方法提供了一個實(shí)現(xiàn)示例，以供比較：

CPU 僅限

CUDA 每組數(shù)據(jù)包的內(nèi)核數(shù)

CUDA 持久內(nèi)核

CUDA 圖

例如，圖 11 顯示了帶有 DPDK gpudev對象的 CUDA 持久內(nèi)核的時間線。

圖 11 使用 DPDK 的 CUDA 持久內(nèi)核的時間線示例gpudev objects

為了測量l2fwd-nv相對于 DPDK testpmd數(shù)據(jù)包生成器的性能，圖 12 中使用了兩個與 CPU 背靠背連接的千兆字節(jié)服務(wù)器： Intel Xeon Gold 6240R 、 PCIe gen3 專用交換機(jī)、 Ubuntu 20.04 、 MOFED 5.4 和 CUDA 11.4 。

圖 12 測試 l2fwd nv 性能的兩個千兆字節(jié)服務(wù)器配置

圖 13 顯示，當(dāng)為數(shù)據(jù)包使用 CPU 或 GPU 內(nèi)存時，峰值 I / O 吞吐量是相同的，因此使用其中一個不會帶來固有的損失。這里的數(shù)據(jù)包被轉(zhuǎn)發(fā)而不被修改。

圖 13 峰值 I / O 吞吐量是相同的

為了突出不同 GPU 數(shù)據(jù)包處理方法之間的差異，圖 14 顯示了方法 2 （ CUDA 內(nèi)核/數(shù)據(jù)包集）和方法 3 （ CUDA 持久內(nèi)核）之間的吞吐量比較。這兩種方法都將數(shù)據(jù)包大小保持在 1024 字節(jié)，在觸發(fā) GPU 工作以交換數(shù)據(jù)包的 MAC 地址之前，改變累積數(shù)據(jù)包的數(shù)量。

圖 14 GPU 數(shù)據(jù)包處理方法之間的差異

對于這兩種方法，每次迭代 16 個數(shù)據(jù)包會導(dǎo)致控制平面中的交互過多，并且無法達(dá)到峰值吞吐量。由于每次迭代 32 個數(shù)據(jù)包，持久化內(nèi)核可以跟上峰值吞吐量，而每次迭代的單個啟動仍然有太多的控制平面開銷。對于每次迭代 64 和 128 個數(shù)據(jù)包，這兩種方法都能夠達(dá)到峰值 I / O 吞吐量。這里的吞吐量測量不是零丟失數(shù)據(jù)包。

結(jié)論

在本文中，我討論了使用 GPU 優(yōu)化內(nèi)聯(lián)數(shù)據(jù)包處理的幾種方法。根據(jù)應(yīng)用程序的需要，您可以應(yīng)用多個工作流模型，以減少延遲，從而提高性能。 DPDK gpudev 庫還有助于簡化您的編碼工作，以在最短的時間內(nèi)獲得最佳結(jié)果。

其他需要考慮的因素，取決于應(yīng)用程序，包括在觸發(fā)數(shù)據(jù)包處理之前，在接收端積累足夠的數(shù)據(jù)包需要花費(fèi)多少時間，有多少線程可用于盡可能多地增強(qiáng)不同任務(wù)之間的并行性，以及內(nèi)核在執(zhí)行中應(yīng)該持續(xù)多長時間。

關(guān)于作者

Elena Agostini與意大利國家研究委員會（ National Research Council of Italy ）合作，獲得了羅馬大學(xué)（ University of Rome “ La Sapienza ”）計(jì)算機(jī)科學(xué)與工程博士學(xué)位。她目前是 NVIDIA 的高級軟件工程師。她的研究興趣包括高性能互連、 GPUDirect 技術(shù)、網(wǎng)絡(luò)協(xié)議、快速數(shù)據(jù)包處理和 DOCA 。 Elena 目前的重點(diǎn)是應(yīng)用于 Aeror 的 NVIDA GPUDirect 技術(shù)，這是一組 SDK ，可實(shí)現(xiàn) GPU 加速、軟件定義的 5G 無線 RAN 。她也是 DPDK 的撰稿人。

審核編輯：郭婷