自7多年前推出至今， CUDA 統(tǒng)一內(nèi)存編程模型一直在開發(fā)人員中廣受歡迎。統(tǒng)一內(nèi)存為 GPU 應(yīng)用程序的原型設(shè)計(jì)提供了一個(gè)簡單的接口，而無需在主機(jī)和設(shè)備之間手動 MIG 評級內(nèi)存。

從 NVIDIA Pascal 更容易擴(kuò)展到更大的問題規(guī)模體系結(jié)構(gòu)開始，支持統(tǒng)一內(nèi)存的應(yīng)用程序可以使用系統(tǒng) CPU 中所有可用的 CPU 和 GPU 內(nèi)存。有關(guān)使用統(tǒng)一內(nèi)存開始 GPU 計(jì)算的更多信息，請參閱 CUDA 更簡單的介紹。

您是否希望使用大型數(shù)據(jù)集無縫運(yùn)行應(yīng)用程序，同時(shí)保持內(nèi)存管理簡單？統(tǒng)一內(nèi)存可用于使虛擬內(nèi)存分配大于可用 GPU 內(nèi)存。在發(fā)生超額訂閱時(shí)， GPU 自動開始將內(nèi)存頁逐出到系統(tǒng)內(nèi)存，以便為活動的在用虛擬內(nèi)存地址騰出空間。

但是，應(yīng)用程序性能在很大程度上取決于內(nèi)存訪問模式、數(shù)據(jù)駐留和運(yùn)行的系統(tǒng)。在過去幾年中，我們發(fā)表了幾篇關(guān)于使用統(tǒng)一內(nèi)存實(shí)現(xiàn) GPU 內(nèi)存超額訂閱的文章。我們通過各種編程技術(shù)（如預(yù)取和內(nèi)存使用提示）為您的應(yīng)用程序?qū)崿F(xiàn)更高的性能提供了幫助。

在這篇文章中，我們深入研究了一個(gè)微基準(zhǔn)測試的性能特征，它強(qiáng)調(diào)了超額訂閱場景中不同的內(nèi)存訪問模式。它可以幫助您分解并了解統(tǒng)一內(nèi)存的所有性能方面：什么時(shí)候適合，什么時(shí)候不適合，以及您可以做些什么。正如您將從我們的結(jié)果中看到的，根據(jù)平臺、超額訂閱因素和內(nèi)存提示，性能可能會變化 100 倍。我們希望這篇文章能讓您更清楚地知道何時(shí)以及如何在應(yīng)用程序中使用統(tǒng)一內(nèi)存！

基準(zhǔn)設(shè)置和訪問模式

要評估統(tǒng)一內(nèi)存超額訂閱性能，可以使用分配和讀取內(nèi)存的簡單程序。使用cudaMallocManaged分配一大塊連續(xù)內(nèi)存，然后在 GPU 上訪問該內(nèi)存，并測量有效的內(nèi)核內(nèi)存帶寬。不同的統(tǒng)一內(nèi)存性能提示，如cudaMemPrefetchAsync和cudaMemAdvise修改分配的統(tǒng)一內(nèi)存。我們將在本文后面討論它們對性能的影響。

我們定義了一個(gè)名為“ oversubscription factor ”的參數(shù)，它控制分配給測試的可用 GPU 內(nèi)存的分?jǐn)?shù)。

值為 1.0 表示 GPU 上的所有可用內(nèi)存都已分配。

小于 1.0 的值表示 GPU 未被超額認(rèn)購

大于 1.0 的值可以解釋為給定 GPU 的超額認(rèn)購量。例如，具有 32 GB 內(nèi)存的 GPU 的超額訂閱因子值為 1.5 意味著使用統(tǒng)一內(nèi)存分配了 48 GB 內(nèi)存。

我們在微基準(zhǔn)測試中測試了三種內(nèi)存訪問內(nèi)核：網(wǎng)格步長、塊邊和隨機(jī)每扭曲。網(wǎng)格跨步和塊跨步是許多 CUDA 應(yīng)用程序中最常見的順序訪問模式。然而，非結(jié)構(gòu)化或隨機(jī)訪問在新興的 CUDA 工作負(fù)載中也非常流行，如圖形應(yīng)用程序、哈希表和推薦系統(tǒng)中的嵌入。我們決定測試這三個(gè)。

網(wǎng)格步長

每個(gè)線程塊在循環(huán)迭代中訪問相鄰內(nèi)存區(qū)域中的元素，然后進(jìn)行網(wǎng)格跨步（blockDim.x * gridDim.x）。

圖 1 網(wǎng)格訪問模式

template
 __global__ void read_thread(data_type *ptr, const size_t size)
 {
     size_t n = size / sizeof(data_type);
     data_type accum = 0;
  
     for(size_t tid = threadIdx.x + blockIdx.x * blockDim.x; tid < n; tid += blockDim.x * gridDim.x)
         accum += ptr[tid];
  
     if (threadIdx.x == 0)
       ptr[0] = accum;
 } 

擋步

每個(gè)線程塊訪問一大塊連續(xù)內(nèi)存，這是根據(jù)分配的總內(nèi)存大小確定的。在任何給定的時(shí)間， SM 上的駐留塊都可以訪問不同的內(nèi)存頁，因?yàn)榉峙浣o每個(gè)塊的內(nèi)存域很大。

圖 2 塊跨步訪問模式

template
 __global__ void read_thread_blockCont(data_type *ptr, const size_t size)
 {
   size_t n = size / sizeof(data_type);
   data_type accum = 0;
  
   size_t elements_per_block = ((n + (gridDim.x - 1)) / gridDim.x) + 1;
   size_t startIdx = elements_per_block * blockIdx.x;
  
   for (size_t rid = threadIdx.x; rid < elements_per_block; rid += blockDim.x) {
     if ((rid + startIdx) < n)
       accum += ptr[rid + startIdx];
   }
  
   if (threadIdx.x == 0)
     ptr[0] = accum;
 } 

隨機(jī)翹曲

在此訪問模式中，對于 warp 的每個(gè)循環(huán)迭代，選擇一個(gè)隨機(jī)頁面，然后訪問一個(gè)連續(xù)的 128B （ 4B 的 32 個(gè)元素）區(qū)域。這將導(dǎo)致線程塊的每個(gè)扭曲跨所有線程塊訪問隨機(jī)頁面。扭曲的循環(huán)計(jì)數(shù)由扭曲的總數(shù)和分配的總內(nèi)存決定。

圖 3 隨機(jī)扭曲訪問模式，扭曲的每個(gè)循環(huán)迭代選擇一個(gè)隨機(jī)頁面并訪問頁面中的隨機(jī) 128B 區(qū)域

內(nèi)核使用線程塊和網(wǎng)格參數(shù)啟動，以實(shí)現(xiàn) 100% 的占用率。內(nèi)核的所有塊始終駐留在 GPU 上。

硬件設(shè)置

本文中的基準(zhǔn)測試使用了以下三種不同硬件設(shè)置中的一種 GPU 。

我們研究了不同的內(nèi)存駐留技術(shù)，以提高這些訪問模式的超額訂閱性能。從根本上說，我們試圖消除統(tǒng)一內(nèi)存頁錯(cuò)誤，并找到最佳的數(shù)據(jù)分區(qū)策略，以獲得基準(zhǔn)測試的最佳讀取帶寬。在本文中，我們將討論以下內(nèi)存模式：

按需 MIG 定額

Zero-copy

CPU 和 GPU 之間的數(shù)據(jù)分區(qū)

在下面的部分中，我們將深入到性能分析和所有優(yōu)化的解釋中。我們還討論了哪些工作負(fù)載能夠與統(tǒng)一內(nèi)存一起很好地解決超額訂閱問題。

基線實(shí)施：按需 MIG 定額

在此測試用例中，使用cudaMallocManaged執(zhí)行內(nèi)存分配，然后按照以下方式在系統(tǒng)（ CPU ）內(nèi)存上填充頁面：

cudaMallocManaged(&uvm_alloc_ptr, allocation_size);
 // all the pages are initialized on CPU
  
 for (int i = 0; i < num_elements; i++)
     uvm_alloc_ptr[i] = 0.0f;

然后，執(zhí)行 GPU 內(nèi)核，并測量內(nèi)核的性能：

read_thread<<>>((float*)uvm_alloc_ptr, allocation_size);

我們使用了上一節(jié)中描述的三種訪問模式之一。這是使用統(tǒng)一內(nèi)存進(jìn)行超額訂閱的最簡單方法，因?yàn)?a target="_blank">程序員不需要提示。

在內(nèi)核調(diào)用時(shí)， GPU 嘗試訪問駐留在主機(jī)上的虛擬內(nèi)存地址。這會觸發(fā)一個(gè)頁面錯(cuò)誤事件，導(dǎo)致通過 CPU – GPU 互連將內(nèi)存頁面 MIG 分配到 GPU 內(nèi)存。內(nèi)核性能受生成的頁面錯(cuò)誤模式和 CPU – GPU 互連速度的影響。

頁面錯(cuò)誤模式是動態(tài)的，因?yàn)樗Q于流式多處理器上塊和扭曲的調(diào)度。然后是 GPU 線程發(fā)出的內(nèi)存加載指令。

圖 4 grid stride ` read _ thread `內(nèi)核執(zhí)行的 NVIDIA NSight 系統(tǒng)時(shí)間線視圖。內(nèi)存行上顯示的 HtoD 和 DtoH 傳輸是由于 MIG 定量和從 GPU 從頁面錯(cuò)誤中逐出造成的。

圖 5 顯示了如何在空 GPU 和超額訂閱 GPU 上處理頁面錯(cuò)誤。在超額訂閱時(shí)，首先將內(nèi)存頁從 GPU 內(nèi)存移出到系統(tǒng)內(nèi)存，然后將請求的內(nèi)存從 CPU 轉(zhuǎn)移到 GPU 。

圖 5 頁面錯(cuò)誤服務(wù)和數(shù)據(jù)逐出機(jī)制。

圖 6 顯示了使用 Power9 CPU 在 V100 、 A100 和 V100 上通過不同訪問模式獲得的內(nèi)存帶寬。

圖 6 基線內(nèi)存分配的讀取帶寬

順序存取分析

訪問模式和不同平臺之間頁面故障驅(qū)動的內(nèi)存讀取帶寬的差異可以通過以下因素來解釋：

訪問模式的影響：傳統(tǒng)上，已知網(wǎng)格跨步訪問模式在訪問 GPU 駐留內(nèi)存時(shí)可實(shí)現(xiàn)最大內(nèi)存帶寬。這里，由于該模式生成的頁面錯(cuò)誤通信量，塊跨步訪問模式實(shí)現(xiàn)了更高的內(nèi)存帶寬。還值得注意的是， Power9 CPU 上的默認(rèn)系統(tǒng)內(nèi)存頁大小為 64 KB ，而 x86 系統(tǒng)上為 4 KB 。這有助于在觸發(fā)頁面錯(cuò)誤事件時(shí)，統(tǒng)一內(nèi)存錯(cuò)誤 MIG 將較大的內(nèi)存塊從 CPU 移動到 GPU 。

對 GPU 體系結(jié)構(gòu)和互連的敏感性： DGX A100 在 CPU 和 GPU 之間具有更快的 PCIe Gen4 互連。這可能是 A100 實(shí)現(xiàn)更高帶寬的原因。然而，互連帶寬并不是飽和的。更高帶寬的主要因素是 A100 GPU 和 108 個(gè)流式多處理器可以產(chǎn)生更多的頁面錯(cuò)誤，因?yàn)?GPU 上有更多的活動線程塊。 P9 測試也證實(shí)了這一理解，盡管 GPU – CPU 之間的 NVLink 連接理論峰值帶寬為 75 GB / s ，但讀取帶寬低于 A100 。

Tip：在這篇文章的實(shí)驗(yàn)中，我們發(fā)現(xiàn)流式網(wǎng)格和塊跨步內(nèi)核訪問模式對線程塊大小和塊內(nèi)同步不敏感。但是，為了使用討論的其他優(yōu)化方法獲得更好的性能，我們在一個(gè)塊中使用了 128 個(gè)線程，在每個(gè)循環(huán)展開時(shí)進(jìn)行塊內(nèi)同步。這確保了塊的所有扭曲有效地使用 SM 的地址轉(zhuǎn)換單元。要了解塊內(nèi)同步的內(nèi)核設(shè)計(jì)，請參閱本文發(fā)布的源代碼。嘗試使用不同塊大小的同步和不同步變體。

隨機(jī)存取分析

在 x86 平臺的超額訂閱域中，由于許多頁面錯(cuò)誤以及由此產(chǎn)生的從 GPU 到 GPU 的內(nèi)存 MIG 比率，隨機(jī)扭曲訪問模式僅產(chǎn)生幾百 KB / s 的讀取帶寬。由于訪問是隨機(jī)的，因此使用了 MIG 額定內(nèi)存的一小部分。額定為 MIG 的內(nèi)存可能最終被逐出回 CPU ，以便為其他內(nèi)存片段騰出空間。

但是，在 Power9 系統(tǒng)上啟用了訪問計(jì)數(shù)器，從而從 GPU 進(jìn)行 CPU 映射內(nèi)存訪問，并且并非所有訪問的內(nèi)存片段都立即被 MIG 評級為 GPU 。這導(dǎo)致了一致的內(nèi)存讀取帶寬，與 x86 系統(tǒng)相比，內(nèi)存抖動更少。

優(yōu)化 1 ：直接訪問系統(tǒng)內(nèi)存（零拷貝）

除了通過互連將內(nèi)存頁從系統(tǒng)內(nèi)存移動到 GPU 內(nèi)存之外，您還可以直接從 GPU 訪問固定系統(tǒng)內(nèi)存。這種內(nèi)存分配方法也稱為零拷貝內(nèi)存。

可使用 CUDA API 調(diào)用cudaMallocHost或通過將虛擬地址范圍的首選位置設(shè)置為 CPU ，從統(tǒng)一內(nèi)存接口分配固定系統(tǒng)內(nèi)存。

cudaMemAdvise（uvm_alloc_ptr， allocation_size， cudaMemAdviseSetPreferredLocation， cudaCpuDeviceId）;

cudaMemAdvise（uvm_alloc_ptr， allocation_size， cudaMemAdviseSetAccessedBy， current_gpu_device）;

圖 7 grid stride ` read _ thread `內(nèi)核直接訪問固定系統(tǒng)內(nèi)存的 NVIDIA NSight 系統(tǒng)時(shí)間線視圖。沒有任何頁面錯(cuò)誤事件或任何方向的內(nèi)存?zhèn)鬏敗?/p>

圖 8 零拷貝內(nèi)存的數(shù)據(jù)訪問路徑

圖 9 顯示了讀內(nèi)核實(shí)現(xiàn)的內(nèi)存帶寬。在 x86 平臺上， A100 GPU 可以實(shí)現(xiàn)比 V100 更高的帶寬，因?yàn)?DGX A100 上 CPU 和 GPU 之間的 PCIe Gen4 互連速度更快。類似地， Power9 系統(tǒng)通過網(wǎng)格跨步訪問模式實(shí)現(xiàn)接近互連帶寬的峰值帶寬。 A100 GPU 上的網(wǎng)格跨步帶寬模式會隨著過度訂閱而降低，因?yàn)?GPU MMU 地址轉(zhuǎn)換未命中會增加加載指令的延遲。

圖 9 零拷貝內(nèi)存的內(nèi)存讀取帶寬

對于所有測試的系統(tǒng)，隨機(jī)扭曲訪問在超額訂閱域中產(chǎn)生 3-4 GB / s 的恒定帶寬。這比前面介紹的故障驅(qū)動場景要好得多。

收獲

從數(shù)據(jù)中可以清楚地看出，零拷貝方法實(shí)現(xiàn)了比基線更高的帶寬。當(dāng)您希望避免從 CPU 和 GPU 取消映射和映射內(nèi)存時(shí)，固定系統(tǒng)內(nèi)存是有利的。如果應(yīng)用程序只使用分配的數(shù)據(jù)一次，那么使用零拷貝內(nèi)存直接訪問更好。但是，如果應(yīng)用程序中存在數(shù)據(jù)重用，則根據(jù)訪問模式和重用情況，對 GPU 的錯(cuò)誤和 MIG 評級數(shù)據(jù)可以產(chǎn)生更高的聚合帶寬。

優(yōu)化 2 ：在 CPU – GPU 之間進(jìn)行數(shù)據(jù)分區(qū)的直接內(nèi)存訪問

對于前面解釋的故障驅(qū)動 MIG 比率， GPU MMU 系統(tǒng)在 GPU 上達(dá)到所需的內(nèi)存范圍之前會出現(xiàn)額外的暫停開銷。為了克服這一開銷，您可以在 CPU 和 GPU 之間分配內(nèi)存，并將內(nèi)存從 GPU 映射到 CPU ，以便于無故障內(nèi)存訪問。

在 CPU 和 GPU 之間分配內(nèi)存有幾種方法：

為內(nèi)存分配設(shè)置了SetAccessedBy統(tǒng)一內(nèi)存提示的cudaMemPrefetchAsync API 調(diào)用。

CPU 和 GPU 之間的手動混合內(nèi)存分配，帶有手動預(yù)取和使用SetPreferredLocation和SetAccessedBy提示。

我們發(fā)現(xiàn)，這兩種方法在許多訪問模式和體系結(jié)構(gòu)組合中表現(xiàn)相似，只有少數(shù)例外。在本節(jié)中，我們主要討論手動頁面分發(fā)。您可以在unified-memory-oversubscription GitHub repo 中查找這兩者的代碼。

圖 10 分配到 GPU 和 CPU 內(nèi)存的頁的內(nèi)存訪問路徑

在混合內(nèi)存分發(fā)中，很少有內(nèi)存頁可以固定到 CPU ，并使用cudaMemAdvise API 調(diào)用將setAccessedBy提示設(shè)置為 GPU 設(shè)備顯式映射內(nèi)存。在我們的測試用例中，我們以循環(huán)方式將多余的內(nèi)存頁映射到 CPU ，其中到 CPU 的映射取決于 GPU 的超額訂閱量。例如，在超額訂閱因子值為 1 。 5 時(shí)，每三個(gè)頁面映射到 CPU 。超額認(rèn)購系數(shù)為 2 。 0 時(shí)，每隔一頁將映射到 CPU 。

在我們的實(shí)驗(yàn)中，內(nèi)存頁設(shè)置為 2MB ，這是 GPU MMU 可以操作的最大頁大小。

圖 11 分布在 CPU 和 GPU 的 2MB 頁面。 Y 軸使用對數(shù)刻度。

對于小于 1 。 0 的超額訂閱值，所有內(nèi)存頁都駐留在 GPU 上。與超額認(rèn)購率大于 1 。 0 的情況相比，您可以看到更高的帶寬。對于大于 1 。 0 的超額訂閱值，基本 HBM 內(nèi)存帶寬和 CPU – GPU 互連速度等因素控制最終內(nèi)存讀取帶寬。

Tip：在 Power9 系統(tǒng)上進(jìn)行測試時(shí)，我們遇到了顯式大容量內(nèi)存預(yù)取的有趣行為（選項(xiàng) a ）。因?yàn)樵?P9 系統(tǒng)上啟用了訪問計(jì)數(shù)器，所以移出的內(nèi)存并不總是固定在 GPU 上，統(tǒng)一內(nèi)存驅(qū)動程序可以啟動從 CPU 到 GPU 的數(shù)據(jù) MIG 分配。這將導(dǎo)致從 GPU 逐出，并且該循環(huán)將在內(nèi)核的整個(gè)生命周期內(nèi)持續(xù)。這個(gè)過程會對流塊和網(wǎng)格步長內(nèi)核產(chǎn)生負(fù)面影響，并且它們比手動頁面分發(fā)獲得的帶寬更低。

解決方案：單一 GPU 超額認(rèn)購

在使用統(tǒng)一內(nèi)存的 GPU 超額訂閱的三種不同內(nèi)存分配策略中，給定應(yīng)用程序分配方法的最佳選擇取決于內(nèi)存訪問模式和 GPU 內(nèi)存的重用。

當(dāng)您在故障和固定系統(tǒng)內(nèi)存分配之間進(jìn)行選擇時(shí)，后者在所有平臺和 GPU 上的性能始終更好。如果內(nèi)存子區(qū)域的 GPU 駐留從總體應(yīng)用程序速度中受益，那么 GPU 和 CPU 之間的內(nèi)存頁分配是一種更好的分配策略。

嘗試統(tǒng)一內(nèi)存優(yōu)化

在這篇文章中，我們回顧了一個(gè)具有一些常見訪問模式的基準(zhǔn)測試，并分析了從 x86 到 P9 ，以及 V100 和 A100 GPU s 的各種平臺上的性能。您可以使用這些數(shù)據(jù)作為參考來進(jìn)行預(yù)測，并考慮在代碼中使用統(tǒng)一內(nèi)存是否有益。我們還介紹了多種數(shù)據(jù)分布模式和統(tǒng)一內(nèi)存模式，它們有時(shí)會帶來顯著的性能優(yōu)勢。有關(guān)更多信息，請參閱 GitHub 上的unified-memory-oversubscription微基準(zhǔn)源代碼。

在上一篇文章中，我們證明了基于統(tǒng)一內(nèi)存的超額訂閱對大數(shù)據(jù)分析和大深度學(xué)習(xí)模型特別有效。請嘗試在代碼中使用統(tǒng)一內(nèi)存進(jìn)行超額訂閱，并讓我們知道它如何幫助您提高應(yīng)用程序性能。

關(guān)于作者

Chirayu Garg 是 NVIDIA 的高級人工智能開發(fā)技術(shù)工程師。他致力于加速 GPU 上的深度學(xué)習(xí)和機(jī)器學(xué)習(xí)應(yīng)用程序。此前，他為 NVIDIA 的游戲流媒體服務(wù)開發(fā)了視頻和圖像處理算法。

審核編輯：郭婷