CUDA 6中的統(tǒng)一內(nèi)存模型

CUDA介紹

CUDA（Compute Unified Device Architecture，統(tǒng)一計算設(shè)備架構(gòu)）是由NVIDIA公司于 2006 年所推出的一種并行計算技術(shù)，是該公司對于GPGPU（ General-purpose computing on graphics processing units， 圖形處理單元上的通用計算 ）技術(shù)的正式命名。通過此技術(shù)，用戶可在GPU上進行通用計算，而開發(fā)人員可以使用C語言來為CUDA架構(gòu)編寫程序 。相比CPU，擁有CUDA技術(shù)的GPU成本不高，但計算性能很突出。本文中提到的是2014年發(fā)布的CUDA6， CUDA6最重要的新特性就是支持統(tǒng)一內(nèi)存模型（Unified Memory）。

注：文中經(jīng)常出現(xiàn)“主機和設(shè)備”，本文的“主機”特指CPU、“設(shè)備”特指GPU。

CUDA 6中的統(tǒng)一內(nèi)存模型

NVIDIA在CUDA 6中引入了統(tǒng)一內(nèi)存模型 （ Unified Memory ），這是CUDA歷史上最重要的編程模型改進之一。在當今典型的PC或群集節(jié)點中，CPU和GPU的內(nèi)存在物理上是獨立的，并通過PCI-Express總線相連。在CUDA6之前， 這是程序員最需要注意的地方。CPU和GPU之間共享的數(shù)據(jù)必須在兩個內(nèi)存中都分配，并由程序直接地在兩個內(nèi)存之間來回復(fù)制。這給CUDA編程帶來了很大難度。

統(tǒng)一內(nèi)存模型創(chuàng)建了一個托管內(nèi)存池（a pool of managed memory），該托管內(nèi)存池由CPU和GPU共享，跨越了CPU與GPU之間的鴻溝。CPU和GPU都可以使用單指針訪問托管內(nèi)存。關(guān)鍵是系統(tǒng)會自動地在主機和設(shè)備之間遷移在統(tǒng)一內(nèi)存中分配的數(shù)據(jù)，從而使那些看起來像CPU內(nèi)存中的代碼在CPU上運行，而另一些看起來像GPU內(nèi)存中的代碼在GPU上運行。

在本文中，我將向您展示統(tǒng)一內(nèi)存模型如何顯著簡化GPU加速型應(yīng)用程序中的內(nèi)存管理。下圖顯示了一個非常簡單的示例。兩種代碼都從磁盤加載文件，對其中的字節(jié)進行排序，然后在釋放內(nèi)存之前使用CPU上已排序的數(shù)據(jù)。右側(cè)的代碼使用CUDA和統(tǒng)一內(nèi)存模型在GPU上運行。和左邊代碼唯一的區(qū)別是，右邊代碼由GPU來啟動一個內(nèi)核（并在啟動后進行同步），并使用新的API cudaMallocManaged（） 在統(tǒng)一內(nèi)存模型中為加載的文件分配空間。

如果您曾經(jīng)編程過CUDA C / C++，那么毫無疑問，右側(cè)的代碼會為您帶來震撼。請注意，我們只分配了一次內(nèi)存，并且只有一個指針指向主機和設(shè)備上的可訪問數(shù)據(jù)。我們可以直接地將文件的內(nèi)容讀取到已分配的內(nèi)存，然后就可以將內(nèi)存的指針傳遞給在設(shè)備上運行的CUDA內(nèi)核。然后，在等待內(nèi)核處理完成之后，我們可以再次從CPU訪問數(shù)據(jù)。CUDA運行時隱藏了所有復(fù)雜性，自動將數(shù)據(jù)遷移到訪問它的地方。

統(tǒng)一內(nèi)存模型提供了什么

統(tǒng)一內(nèi)存模型為程序員提供了兩大捷徑

簡化編程、簡化內(nèi)存模型

統(tǒng)一內(nèi)存模型通過使設(shè)備內(nèi)存管理（device memory management）成為一項可選的優(yōu)化，而不是一項硬性的要求，從而降低了CUDA平臺上并行編程的門檻。借助統(tǒng)一內(nèi)存模型，程序員現(xiàn)在可以直接開發(fā)并行的CUDA內(nèi)核，而不必擔心分配和復(fù)制設(shè)備內(nèi)存的細節(jié)。這將降低在CUDA平臺上編程的學(xué)習(xí)成本，也使得將現(xiàn)有代碼移植到GPU的工作變得容易。但這些好處不僅有利于初學(xué)者。我在本文后面的示例中將展示統(tǒng)一內(nèi)存模型如何使復(fù)雜的數(shù)據(jù)結(jié)構(gòu)更易于與設(shè)備代碼一起使用，以及它與C++結(jié)合時的強大威力。

通過數(shù)據(jù)局部性原理提高性能

通過在CPU和GPU之間按需遷移數(shù)據(jù)，統(tǒng)一內(nèi)存模型可以滿足GPU上本地數(shù)據(jù)的性能需求，同時還提供了易于使用的全局共享數(shù)據(jù)。這個功能的復(fù)雜細節(jié)被 CUDA驅(qū)動程序和運行時隱藏了，以確保應(yīng)用程序代碼更易于編寫。遷移的關(guān)鍵是從每個處理器獲得全部帶寬。250 GB / s的GDDR5內(nèi)存對于保證開普勒（ Kepler ）GPU的計算吞吐量至關(guān)重要。

值得注意的是， 一個經(jīng)過精心調(diào)優(yōu)的CUDA程序，即使用流（streams）和 cudaMemcpyAsync來有效地將執(zhí)行命令與數(shù)據(jù)傳輸重疊的程序，會比僅使用統(tǒng)一內(nèi)存模型的CUDA程序更好 。可以理解的是：CUDA運行時從來沒有像程序員那樣提供何處需要數(shù)據(jù)或何時需要數(shù)據(jù)的信息！CUDA程序員仍然可以顯式地訪問設(shè)備內(nèi)存分配和異步內(nèi)存拷貝，以優(yōu)化數(shù)據(jù)管理和CPU-GPU并發(fā)機制 。首先，統(tǒng)一內(nèi)存模型提高了生產(chǎn)力，它為并行計算提供了更順暢的入口，同時它又不影響高級用戶的任何CUDA功能。

統(tǒng)一內(nèi)存模型 vs 統(tǒng)一虛擬尋址？

自CUDA4起，CUDA就支持統(tǒng)一虛擬尋址（UVA），并且盡管統(tǒng)一內(nèi)存模型依賴于UVA，但它們并不是一回事。UVA為 系統(tǒng)中的所有內(nèi)存提供了單個虛擬內(nèi)存地址空間，無論指針位于系統(tǒng)中的何處，無論在設(shè)備內(nèi)存（在相同或不同的GPU上）、主機內(nèi)存、或片上共享存儲器。UVA也允許 cudaMemcpy在不指定輸入和輸出參數(shù)確切位置的情況下使用。UVA啟用“零復(fù)制（Zero-Copy）” 內(nèi)存，“零復(fù)制”內(nèi)存是固定的主機內(nèi)存，可由設(shè)備上的代碼通過PCI-Express總線直接訪問，而無需使用 memcpy。零復(fù)制為統(tǒng)一內(nèi)存模型提供了一些便利，但是卻沒有提高性能，因為它總是通過帶寬低而且延遲高的PCI-Express進行訪問。

UVA不會像統(tǒng)一內(nèi)存模型一樣自動將數(shù)據(jù)從一個物理位置遷移到另一個物理位置。由于統(tǒng)一內(nèi)存模型能夠在主機和設(shè)備內(nèi)存之間的各級頁面自動地遷移數(shù)據(jù)，因此它需要進行大量的工程設(shè)計，因為它需要在CUDA運行時（runtime）、設(shè)備驅(qū)動程序、甚至OS內(nèi)核中添加新功能。以下示例旨在讓您領(lǐng)會到這一點。示例：消除深層副本

統(tǒng)一內(nèi)存模型的主要優(yōu)勢在于，在訪問GPU內(nèi)核中的結(jié)構(gòu)化數(shù)據(jù)時，無需進行深度復(fù)制（deep copies），從而簡化了異構(gòu)計算內(nèi)存模型。如下圖所示，將包含指針的數(shù)據(jù)結(jié)構(gòu)從CPU傳遞到GPU要求進行“深度復(fù)制”。

下面以struct dataElem為例。

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

要在設(shè)備上使用此結(jié)構(gòu)體，我們必須復(fù)制結(jié)構(gòu)體本身及其數(shù)據(jù)成員，然后復(fù)制該結(jié)構(gòu)體指向的所有數(shù)據(jù)，然后更新該結(jié)構(gòu)體。副本中的所有指針。這導(dǎo)致下面的復(fù)雜代碼，這些代碼只是將數(shù)據(jù)元素傳遞給內(nèi)核函數(shù)。

void launch（dataElem *elem） { dataElem *d_elem;char *d_name;

int namelen = strlen（elem-》name） + 1;

// Allocate storage for struct and name cudaMalloc（&d_elem， sizeof（dataElem））; cudaMalloc（&d_name， namelen）;

// Copy up each piece separately， including new “name” pointer value cudaMemcpy（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代碼之間分享復(fù)雜的數(shù)據(jù)結(jié)構(gòu)所需的額外主機端代碼對生產(chǎn)率有嚴重影響。統(tǒng)一內(nèi)存模型中分配我們的“ dataElem”結(jié)構(gòu)可消除所有多余的設(shè)置代碼，這些代碼與主機代碼被相同的指針操作，留給我們的就只有內(nèi)核啟動了。這是一個很大的進步！

void launch（dataElem *elem） { kernel《《《 。.. 》》》（elem）;}

但統(tǒng)一內(nèi)存模型不僅大幅降低了代碼復(fù)雜性。還可以做一些以前無法想象的事情。讓我們看另一個例子。

Example： CPU/GPU Shared Linked Lists

鏈表是一種非常常見的數(shù)據(jù)結(jié)構(gòu)，但是由于它們本質(zhì)上是由指針組成的嵌套數(shù)據(jù)結(jié)構(gòu)，因此在內(nèi)存空間之間傳遞它們非常復(fù)雜。如果沒有統(tǒng)一內(nèi)存模型，則無法在CPU和GPU之間分享鏈表。唯一的選擇是在零拷貝內(nèi)存（被pin住的主機內(nèi)存）中分配鏈表，這意味著GPU的訪問受限于PCI-express性能。通過在統(tǒng)一內(nèi)存模型中分配鏈表數(shù)據(jù)，設(shè)備代碼可以正常使用GPU上的指針，從而發(fā)揮設(shè)備內(nèi)存的全部性能。程序可以維護單鏈表，并且無論在主機或設(shè)備中都可以添加和刪除鏈表元素。

將具有復(fù)雜數(shù)據(jù)結(jié)構(gòu)的代碼移植到GPU上曾經(jīng)是一項艱巨的任務(wù)，但是統(tǒng)一內(nèi)存模型使此操作變得非常容易。我希望統(tǒng)一內(nèi)存模型能夠為CUDA程序員帶來巨大的生產(chǎn)力提升。

Unified Memory with C++

統(tǒng)一內(nèi)存模型確實在C++數(shù)據(jù)結(jié)構(gòu)中大放異彩。C++通過帶有拷貝構(gòu)造函數(shù)（copy constructors）的類來簡化深度復(fù)制問題?？截悩?gòu)造函數(shù)是一個知道如何創(chuàng)建類所對應(yīng)對象的函數(shù)，拷貝構(gòu)造函數(shù)為對象的成員分配空間并從其他對象復(fù)制值過來。C++還允許 new和 delete這倆個內(nèi)存管理運算符被重載。這意味著我們可以創(chuàng)建一個基類，我們將其稱為 Managed，它在重載的 new運算符內(nèi)部使用 cudaMallocManaged（），如以下代碼所示。

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類，并實現(xiàn)一個拷貝構(gòu)造函數(shù)，該拷貝構(gòu)造函數(shù)為需要拷貝的字符串分配統(tǒng)一內(nèi)存。

// Deriving from “Managed” allows pass-by-referenceclass String ： public Managed { int length; char *data;

public:// Unified memory copy constructor allows pass-by-value String （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 copiesclass dataElem ： public Managed {public:int prop1;int prop2; String name;};

通過這些更改，C++的類將在統(tǒng)一內(nèi)存中分配空間，并自動處理深度復(fù)制。我們可以像分配任何C++的對象那樣在統(tǒng)一內(nèi)存中分配一個 dataElem。

dataElem *data = new dataElem;

請注意，您需要確保樹中的每個類都繼承自 Managed，否則您的內(nèi)存映射中會有一個漏洞。實際上，任何你想在CPU和GPU之間分享的內(nèi)容都應(yīng)該繼承 Managed。如果你傾向于對所有程序都簡單地使用統(tǒng)一內(nèi)存模型，你可以在全局重載 new和 delete， 但這只在這種情況下有作用——你的程序中沒有僅被CPU訪問的數(shù)據(jù)（即程序中的所有數(shù)據(jù)都被GPU訪問），因為只有CPU數(shù)據(jù)時沒有必要遷移數(shù)據(jù)。

現(xiàn)在，我們可以選擇將對象傳遞給內(nèi)核函數(shù)了。如在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 reference kernel_by_ref《《《1，1》》》（*data）;

// pass data to kernel by value -- this will create a copy kernel_by_val《《《1，1》》》（*data）;}

多虧了統(tǒng)一內(nèi)存模型，深度復(fù)制、按值傳遞和按引用傳遞都可以正常工作。統(tǒng)一內(nèi)存模型為在GPU上運行C++代碼提供了巨大幫助。

這篇文章的例子可以在Github上找到。

統(tǒng)一內(nèi)存模型的光明前景

CUDA 6中關(guān)于統(tǒng)一內(nèi)存模型的最令人興奮的事情之一就是它僅僅是個開始。我們針對統(tǒng)一內(nèi)存模型有一個包括性能提升與特性的長遠規(guī)劃。我們的第一個發(fā)行版旨在使CUDA編程更容易，尤其是對于初學(xué)者而言。從CUDA 6開始， cudaMemcpy（）不再是必需的。通過使用 cudaMallocManaged（），您可以擁有一個指向數(shù)據(jù)的指針，并且可以在CPU和GPU之間共享復(fù)雜的C / C++數(shù)據(jù)結(jié)構(gòu)。這使編寫CUDA程序變得容易得多，因為您可以直接編寫內(nèi)核，而不是編寫大量數(shù)據(jù)管理代碼并且要維護在主機和設(shè)備之間所有重復(fù)的數(shù)據(jù)。您仍然可以自由使用 cudaMemcpy（）（特別是 cudaMemcpyAsync（））來提高性能，但現(xiàn)在這不是一項要求，而是一項優(yōu)化。

CUDA的未來版本可能會通過添加數(shù)據(jù)預(yù)取和遷移提示來提高使用統(tǒng)一內(nèi)存模型的應(yīng)用程序的性能。我們還將增加對更多操作系統(tǒng)的支持。我們的下一代GPU架構(gòu)將帶來許多硬件改進，以進一步提高性能和靈活性。
責任編輯:pj