了解cuda的統一內存

1. CUDA 6中的統一內存

在CUDA 6中，從Kepler GPU架構（計算能力3.0或更高）開始，在64位Windows 7、8和Linux操作系統（內核2.6.18+）上開始支持統一內存.
在這里插入圖片描述
從CUDA 6開始，NVIDIA推出了CUDA平臺歷史上最引人注目的編程模型改進之一 ---- 統一內存。
在當今典型的PC或集群節點中，CPU和GPU的內存在物理上是不同的，并由PCI Express總線分隔。
在CUDA 6之前，程序員必須這樣看待事物。CPU和GPU之間共享的數據必須在兩個內存中分配，并由程序在它們之間明確復制。這給CUDA程序增加了很多復雜性。

在這里插入圖片描述
統一內存創建了一個在CPU和GPU之間共享的托管內存池，彌合了CPU-GPU的鴻溝。
CPU和GPU都可以使用單個指針訪問托管內存。關鍵在于，系統會自動在主機和設備之間遷移統一內存中分配的數據，使其看起來像CPU內存當在CPU上運行代碼時，而在GPU上運行代碼時像是GPU內存。

在這篇文章中，我將向您展示統一內存如何大大簡化GPU加速應用程序中的內存管理。
下圖顯示了一個非常簡單的例子。這兩個代碼都從磁盤加載文件，對其中的字節進行排序，然后在釋放內存之前在CPU上使用排序后的數據。右側的代碼使用CUDA和統一內存在GPU上運行.
兩個代碼的唯一的區別是GPU版本啟動內核函數（并在啟動后進行同步），并使用新的API cudaMallocManaged（）為加載的文件在統一內存中分配空間.

在這里插入圖片描述
如果你以前用過CUDA C/C++編程，你無疑會被右邊代碼的簡潔性所打動。請注意，我們只分配了一次內存，并且我們有一個指向主機和設備都可以訪問的數據的指針。我們可以直接從文件讀取內容到分配的內存，然后將指針直接傳遞給在設備上運行的CUDA內核。然后，在等待內核完成之后，我們可以再次從CPU訪問數據。CUDA運行時隱藏了所有的復雜性，自動將數據遷移到訪問它的地方。

2 統一內存提供了什么

程序員從統一內存中受益的主要方式有兩種。

2.1 更簡單的編程和內存模型

統一內存降低了CUDA平臺上并行編程的門檻,通過使設備內存管理成為一種優化，而不是一種要求
有了統一內存，現在程序員可以直接開發并行CUDA內核，而不會陷入分配和復制設備內存的細節中。
這將使學習CUDA平臺的編程和將現有代碼移植到GPU更簡單。但這不僅僅適用于初學者。
本文后面的示例展示了統一內存如何使復雜的數據結構更容易與設備代碼一起使用，以及它與C++結合時的強大功能。

2.2 通過數據本地化提升性能

通過在CPU和GPU之間按需遷移數據，統一內存可以在GPU上提供本地數據的性能，同時提供全局共享數據的易用性。此功能的復雜性被隱藏在CUDA驅動程序和運行時的保護之下，確保應用程序代碼更易于編寫。遷移的目的是實現每個處理器的全帶寬；250 GB/s的GDDR5內存對于提升開普勒GPU的計算吞吐量至關重要。

一個重要的點是，一個經過精心調優的CUDA程序，它使用流和cudaemcpyAsync來有效地將執行與數據傳輸重疊，可能會比只使用統一內存的CUDA軟件表現得更好。容易理解的是：CUDA運行時永遠不會像程序員那樣了解需要數據的位置和時間！CUDA程序員仍然可以訪問顯式設備內存分配和異步內存副本，以優化數據管理和CPU-GPU并發性。統一內存首先是一種生產力功能，它為并行計算提供了更平滑的入口，而不會剝奪CUDA為高級用戶提供的任何功能。

3 統一內存還是統一虛擬尋址？

CUDA自CUDA 4以來一直支持統一虛擬尋址（UVA），雖然統一內存依賴于UVA，但它們不是一回事。UVA為系統中的所有內存提供了一個單一的虛擬內存地址空間，并允許從GPU代碼訪問指針，無論它們位于系統的哪個位置，無論是設備內存（在相同或不同的GPU上）、主機內存還是片上共享內存。它還允許使用cudaMemcpy，而無需指定輸入和輸出參數的確切位置。UVA支持“零拷貝”內存，即設備代碼可以直接通過PCI Express訪問的固定主機內存，無需memcpy。Zero Copy提供了統一內存的一些便利，但沒有提供任何性能，因為它總是使用PCI Express的低帶寬和高延遲進行訪問。

UVA不會像統一內存那樣自動將數據從一個物理位置遷移到另一個。由于統一內存能夠在主機和設備內存之間自動遷移單個頁面級別的數據，因此需要大量的工程來構建，因為它需要CUDA運行時、設備驅動程序甚至操作系統內核中的新功能。以下示例旨在讓您了解其功能。

3.1 例子：消除深拷貝

統一內存的一個關鍵好處是簡化了異構計算內存模型，因為在訪問GPU內核中的結構化數據時不需要深度副本。將包含指針的數據結構從CPU傳遞到GPU需要進行“深度復制”，如下圖所示。
在這里插入圖片描述
dataElem結構如下

struct dataElem {int prop1;int prop2;char *name;
}

要在設備上使用此結構，我們必須復制結構本身及其數據成員，然后復制結構指向的所有數據，然后更新結構副本中的所有指針。這導致了以下復雜的代碼，只是為了將數據元素傳遞給內核函數。

void launch(dataElem *elem) {dataElem *d_elem;char *d_name;int namelen = strlen(elem->name) + 1;// Allocate storage for struct and namecudaMalloc(&d_elem, sizeof(dataElem));cudaMalloc(&d_name, namelen);// Copy up each piece separately, including new “name” pointer valuecudaMemcpy(d_elem, elem, sizeof(dataElem), cudaMemcpyHostToDevice);cudaMemcpy(d_name, elem->name, namelen, cudaMemcpyHostToDevice);cudaMemcpy(&(d_elem->name), &d_name, sizeof(char*), cudaMemcpyHostToDevice);// Finally we can launch our kernel, but CPU & GPU use different copies of “elem”Kernel<<< ... >>>(d_elem);
}

可以想象，在CPU和GPU代碼之間共享復雜數據結構所需的額外主機端代碼對生產率有重大影響。在統一內存中分配我們的dataElem結構消除了所有多余的設置代碼，只剩下內核啟動，它與宿主代碼在同一指針上運行。這是一個很大的進步！

void launch(dataElem *elem) {kernel<<< ... >>>(elem);
}

但這不僅僅是代碼復雜性的一大改進。統一內存還可以做以前無法想象的事情。讓我們來看另一個例子:

3.2 例子: CPU/GPU之間共享鏈表

在這里插入圖片描述
鏈表是一種非常常見的數據結構，但由于它們本質上是由指針組成的嵌套數據結構，在內存空間之間傳遞它們非常復雜。如果沒有統一內存，CPU和GPU之間共享鏈表是無法管理的。唯一的選擇是在零拷貝內存（固定主機內存）中分配列表，這意味著GPU訪問僅限于PCI express性能。通過在統一內存中分配鏈表數據，設備代碼可以在GPU上正常跟隨指針，并具有設備內存的全部性能。該程序可以維護一個鏈表，可以在主機或設備上添加和刪除列表元素.

將具有現有復雜數據結構的代碼移植到GPU曾經是一項艱巨的任務，但統一內存使這變得更加容易。我預計統一內存將為CUDA程序員帶來巨大的生產力提升。

4. c++中使用統一內存

統一內存在C++數據結構中大放異彩。C++通過使用帶有復制構造函數的類簡化了深度復制問題。復制構造函數是一個函數，它知道如何創建類的對象，為其成員分配空間，并從另一個對象復制它們的值。C++還允許重載new和delete內存管理運算符。這意味著我們可以創建一個基類，我們稱之為Managed，它在重載的new運算符中使用cudaAllocManaged（），如下代碼所示。

class Managed {
public:void *operator new(size_t len) {void *ptr;cudaMallocManaged(&ptr, len);cudaDeviceSynchronize();return ptr;}void operator delete(void *ptr) {cudaDeviceSynchronize();cudaFree(ptr);}
};

然后，我們可以讓String類從Managed類繼承，并實現一個復制構造函數，為復制的字符串分配統一內存。

// Deriving from “Managed” allows pass-by-reference
class String : public Managed {int length;char *data;public:// Unified memory copy constructor allows pass-by-valueString (const String &s) {length = s.length;cudaMallocManaged(&data, length);memcpy(data, s.data, length);}// ...
};

同樣，我們使dataElem類繼承Managed。

// Note “managed” on this class, too.
// C++ now handles our deep copies
class dataElem : public Managed {
public:int prop1;int prop2;String name;
};

通過這些更改，C++類在統一內存中分配存儲，并自動處理深度副本。我們可以像任何C++對象一樣在統一內存中分配dataElem。

dataElem *data = new dataElem;

請注意，您需要確保繼承樹中的每個類都繼承自Managed，否則您的內存映射中會出現漏洞。實際上，您可能需要在CPU和GPU之間共享的所有內容都應該繼承Managed。如果你更喜歡簡單地使用統一內存來處理所有事情，你可以在全局范圍內重載new和delete，但這只有在沒有僅CPU數據的情況下才有意義，否則數據將不必要地遷移。
現在，當我們將對象傳遞給內核函數時，我們有一個選擇；與C++中的正常情況一樣，我們可以按值傳遞或按引用傳遞，如下面的示例代碼所示。

// Pass-by-reference version
__global__ void kernel_by_ref(dataElem &data) { ... }// Pass-by-value version
__global__ void kernel_by_val(dataElem data) { ... }int main(void) {dataElem *data = new dataElem;...// pass data to kernel by referencekernel_by_ref<<<1,1>>>(*data);// pass data to kernel by value -- this will create a copykernel_by_val<<<1,1>>>(*data);
}

得益于統一內存，深度拷貝、按值傳遞和按引用傳遞都能正常工作。這為在GPU上運行C++代碼提供了巨大的價值。
這篇文章中的例子可以在Github上找到。

5. 統一內存的光明未來

CUDA 6中統一內存最令人興奮的事情之一是它只是一個開始。cuda圍繞統一內存計劃了一個漫長的改進和功能路線圖。
統一內存的第一個版本旨在使CUDA編程更容易，特別是對于初學者。從CUDA 6開始，cudaemcpy（）不再是必需的。通過使用cudaAllocManaged（），您可以有一個指向數據的指針，并且可以在CPU和GPU之間共享復雜的C/C++數據結構。這使得編寫CUDA程序變得更加容易，因為您可以直接編寫內核，而不是編寫大量數據管理代碼并維護所有數據的重復主機和設備副本。您仍然可以自由地使用cudaemcpy（）（特別是cudamemppyAsync（））來提高性能，但這不是一項要求，而是一種優化。

CUDA的未來版本可能會通過添加數據預取和遷移提示來提高使用統一內存的應用程序的性能。統一內存還將增加對更多操作系統的支持。nv的下一代GPU架構將帶來一系列硬件改進，以進一步提高性能和靈活性.