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

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

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

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

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

CPU 用于查詢(xún)網(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ù)包工作流場(chǎng)景。

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

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

您可以清楚地識(shí)別兩種不同的流：

Data flow ：通過(guò) 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ù)移動(dòng)（發(fā)送或接收數(shù)據(jù)包）。它可以通過(guò) GPUDirect RDMA 技術(shù)實(shí)現(xiàn)，該技術(shù)使用 PCI Express 總線接口的標(biāo)準(zhǔn)功能，在 NVIDIA GPU 和第三方對(duì)等設(shè)備（如網(wǎng)卡）之間實(shí)現(xiàn)直接數(shù)據(jù)路徑。

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

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

要在 Linux 系統(tǒng)上啟用 GPUDirect RDMA ，需要 nvidia-peermem 模塊（在 CUDA 11.4 及更高版本中提供）。圖 3 顯示了最大化 GPUDirect RDMA 內(nèi)部吞吐量的理想系統(tǒng)拓?fù)洌涸?GPU 和 NIC 之間使用專(zhuān)用 PCIe 交換機(jī)，而不是通過(guò)與其他組件共享的系統(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)和同步活動(dòng)的主要參與者，用于喚醒 NIC ，將數(shù)據(jù)包接收到 GPU 內(nèi)存中，并通知 CUDA 工作負(fù)載有新數(shù)據(jù)包可供處理。

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

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

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

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

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

方法 1

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

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

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

方法 2

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

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

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

方法 3

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

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

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

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

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

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

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

方法 4

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

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

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

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

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

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

DPDK 和 GPUdev

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

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

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

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

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

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

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

通過(guò)通用層公開(kāi) GPU 特定于驅(qū)動(dòng)程序的功能。

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

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

DPDK 在 mempool 中接收數(shù)據(jù)包，這是一個(gè)連續(xù)的內(nèi)存塊。通過(guò)以下指令序列，您可以啟用 GPUDirect RDMA 在 GPU 內(nèi)存中分配 mempool ，并將其注冊(cè)到設(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)

對(duì)于控制流，要啟用 CPU 和 GPU 之間的通知機(jī)制，可以使用gpudev通信列表：在 CPU 內(nèi)存和 CUDA 內(nèi)核之間的共享內(nèi)存結(jié)構(gòu)。列表中的每一項(xiàng)都可以保存接收到的數(shù)據(jù)包的地址（mbufs），以及一個(gè)用于更新處理該項(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++; }

為簡(jiǎn)單起見(jiàn)，假設(shè)應(yīng)用程序遵循 CUDA 持久內(nèi)核場(chǎng)景， CUDA 內(nèi)核上的輪詢(xún)端看起來(lái)類(lè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庫(kù)的實(shí)際示例，l2fwd-nv示例代碼已在 /NVIDIA/l2fwd-nv GitHub repo 上發(fā)布。這是使用 GPU 功能增強(qiáng)的普通 DPDK l2fwd示例的擴(kuò)展。應(yīng)用程序布局是接收數(shù)據(jù)包，交換每個(gè)數(shù)據(jù)包的 MAC 地址（源和目的地），并傳輸修改后的數(shù)據(jù)包。

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

CPU 僅限

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

CUDA 持久內(nèi)核

CUDA 圖

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

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

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

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

圖 13 顯示，當(dāng)為數(shù)據(jù)包使用 CPU 或 GPU 內(nèi)存時(shí)，峰值 I / O 吞吐量是相同的，因此使用其中一個(gè)不會(huì)帶來(lái)固有的損失。這里的數(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ù)包處理方法之間的差異

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

結(jié)論

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

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

關(guān)于作者

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

審核編輯：郭婷