我之前的介紹文章，“ 更容易介紹 CUDA C ++ ”介紹了 CUDA 編程的基本知識(shí)，它演示了如何編寫一個(gè)簡(jiǎn)單的程序，在內(nèi)存中分配兩個(gè)可供 GPU 訪問(wèn)的數(shù)字?jǐn)?shù)組，然后將它們加在 GPU 上。為此，我向您介紹了統(tǒng)一內(nèi)存，這使得分配和訪問(wèn)系統(tǒng)中任何處理器上運(yùn)行的代碼都可以使用的數(shù)據(jù)變得非常容易， CPU 或 GPU 。

我以幾個(gè)簡(jiǎn)單的“練習(xí)”結(jié)束了這篇文章，其中一個(gè)練習(xí)鼓勵(lì)您運(yùn)行最近基于 Pascal 的 GPU ，看看會(huì)發(fā)生什么。（我希望讀者能嘗試一下并對(duì)結(jié)果發(fā)表評(píng)論，你們中的一些人也這樣做了?。?。我建議這樣做有兩個(gè)原因。首先，因?yàn)?PascalMIG 如 NVIDIA Titan X 和 NVIDIA Tesla P100 是第一個(gè)包含頁(yè) GPUs 定額引擎的 GPUs ，它是統(tǒng)一內(nèi)存頁(yè)錯(cuò)誤處理和 MIG 比率的硬件支持。第二個(gè)原因是它提供了一個(gè)很好的機(jī)會(huì)來(lái)學(xué)習(xí)更多的統(tǒng)一內(nèi)存。

快 GPU ，快內(nèi)存…對(duì)嗎？

正確的！但讓我們看看。首先，我將重新打印在兩個(gè) NVIDIA 開(kāi)普勒 GPUs 上運(yùn)行的結(jié)果（一個(gè)在我的筆記本電腦上，一個(gè)在服務(wù)器上）。

現(xiàn)在讓我們嘗試在一個(gè)非?？斓?Tesla P100 加速器上運(yùn)行，它基于 pascalgp100GPU 。

> nvprof ./add_grid ... Time(%) Time Calls Avg Min Max Name 100.00% 2.1192ms 1 2.1192ms 2.1192ms 2.1192ms add(int, float*, float*)

嗯，這低于 6gb / s ：比在我的筆記本電腦基于開(kāi)普勒的 GeForceGPU 上運(yùn)行慢。不過(guò)，別灰心，我們可以解決這個(gè)問(wèn)題的。為了理解這一點(diǎn)，我將告訴你更多關(guān)于統(tǒng)一內(nèi)存的信息。

下面是要添加的完整代碼，以供參考_網(wǎng)格. cu 從上次開(kāi)始。

#include  #include  // CUDA kernel to add elements of two arrays __global__ void add(int n, float *x, float *y) { int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; for (int i = index; i < n; i += stride) y[i] = x[i] + y[i]; } int main(void) { int N = 1<<20; float *x, *y; // Allocate Unified Memory -- accessible from CPU or GPU cudaMallocManaged(&x, N*sizeof(float)); cudaMallocManaged(&y, N*sizeof(float)); // initialize x and y arrays on the host for (int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; } // Launch kernel on 1M elements on the GPU int blockSize = 256; int numBlocks = (N + blockSize - 1) / blockSize; add<<>>(N, x, y); // Wait for GPU to finish before accessing on host cudaDeviceSynchronize(); // Check for errors (all values should be 3.0f) float maxError = 0.0f; for (int i = 0; i < N; i++) maxError = fmax(maxError, fabs(y[i]-3.0f)); std::cout << "Max error: " << maxError << std::endl; // Free memory cudaFree(x); cudaFree(y); return 0; }

對(duì) 27-19 行的內(nèi)存進(jìn)行初始化。

什么是統(tǒng)一內(nèi)存？

統(tǒng)一內(nèi)存是可從系統(tǒng)中的任何處理器訪問(wèn)的單個(gè)內(nèi)存地址空間（請(qǐng)參見(jiàn)圖 1 ）。這種硬件/軟件技術(shù)允許應(yīng)用程序分配可以從 CPU s 或 GPUs 上運(yùn)行的代碼讀取或?qū)懭氲臄?shù)據(jù)。分配統(tǒng)一內(nèi)存非常簡(jiǎn)單，只需將對(duì)malloc()或new的調(diào)用替換為對(duì)cudaMallocManaged()的調(diào)用，這是一個(gè)分配函數(shù)，返回可從任何處理器訪問(wèn)的指針（以下為ptr）。

cudaError_t cudaMallocManaged(void** ptr, size_t size);

當(dāng)在 CPU 或 GPU 上運(yùn)行的代碼訪問(wèn)以這種方式分配的數(shù)據(jù)（通常稱為 CUDA 管理 數(shù)據(jù)）， CUDA 系統(tǒng)軟件和/或硬件負(fù)責(zé)將 MIG 額定內(nèi)存頁(yè)分配給訪問(wèn)處理器的內(nèi)存。這里重要的一點(diǎn)是， PascalGPU 體系結(jié)構(gòu)是第一個(gè)通過(guò)頁(yè)面 MIG 比率引擎對(duì)虛擬內(nèi)存頁(yè)錯(cuò)誤處理和 MIG 比率提供硬件支持的架構(gòu)。基于更舊的 kezbr 架構(gòu)和更為統(tǒng)一的 kezbr 形式的支持。

當(dāng)我打電話給cudaMallocManaged()時(shí)，開(kāi)普勒會(huì)發(fā)生什么？

在具有 pre-PascalGPUs 的系統(tǒng)上，如 Tesla K80 ，調(diào)用 cudaMallocManaged（） 會(huì)分配 size 字節(jié)的托管內(nèi)存 在 GPU 設(shè)備上 ，該內(nèi)存在調(diào)用 1 時(shí)處于活動(dòng)狀態(tài)。在內(nèi)部，驅(qū)動(dòng)程序還為分配覆蓋的所有頁(yè)面設(shè)置頁(yè)表?xiàng)l目，以便系統(tǒng)知道這些頁(yè)駐留在 GPU 上。

所以，在我們的例子中，在 Tesla K80GPU （開(kāi)普勒架構(gòu)）上運(yùn)行， x 和 y 最初都完全駐留在 GPU 內(nèi)存中。然后在第 6 行開(kāi)始的循環(huán)中， CPU 逐步遍歷兩個(gè)數(shù)組，分別將它們的元素初始化為 1.0f 和 2.0f 。由于這些頁(yè)最初駐留在設(shè)備存儲(chǔ)器中，所以它寫入的每個(gè)數(shù)組頁(yè)的 CPU 上都會(huì)發(fā)生一個(gè)頁(yè)錯(cuò)誤， GPU 驅(qū)動(dòng)程序 MIG 會(huì)將設(shè)備內(nèi)存中的頁(yè)面分配給 CPU 內(nèi)存。循環(huán)之后，兩個(gè)數(shù)組的所有頁(yè)都駐留在 CPU 內(nèi)存中。

在初始化 CPU 上的數(shù)據(jù)之后，程序啟動(dòng) add（） 內(nèi)核，將 x 的元素添加到 y 的元素中。

add<<<1, 256>>>(N, x, y);

在 pre-PascalGPUs 上，啟動(dòng)一個(gè)內(nèi)核后， CUDA 運(yùn)行時(shí)必須 MIG 將以前 MIG 額定為主機(jī)內(nèi)存或另一個(gè) GPU 的所有頁(yè)面重新評(píng)級(jí)到運(yùn)行內(nèi)核 2 的設(shè)備內(nèi)存。由于這些舊的 GPUs 不能出現(xiàn)分頁(yè)錯(cuò)誤，所有數(shù)據(jù)都必須駐留在 GPU 以防萬(wàn)一 上，內(nèi)核訪問(wèn)它（即使它不會(huì)訪問(wèn)）。這意味著每次啟動(dòng)內(nèi)核時(shí)都可能存在 MIG 定額開(kāi)銷。

當(dāng)我在 K80 或 macbookpro 上運(yùn)行程序時(shí)，就會(huì)發(fā)生這種情況。但是請(qǐng)注意，探查器顯示的內(nèi)核運(yùn)行時(shí)間與 MIG 定額時(shí)間是分開(kāi)的，因?yàn)?MIG 定額發(fā)生在內(nèi)核運(yùn)行之前。

==15638== Profiling application: ./add_grid ==15638== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00% 93.471us 1 93.471us 93.471us 93.471us add(int, float*, float*) ==15638== Unified Memory profiling result: Device "Tesla K80 (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 6 1.3333MB 896.00KB 2.0000MB 8.000000MB 1.154720ms Host To Device 102 120.47KB 4.0000KB 0.9961MB 12.00000MB 1.895040ms Device To Host Total CPU Page faults: 51

當(dāng)我調(diào)用cudaMallocManaged()時(shí)， Pascal 上會(huì)發(fā)生什么？

在 Pascal 和更高版本的 GPUs 上， cudaMallocManaged（） 返回時(shí)可能不會(huì)物理分配托管內(nèi)存；它只能在訪問(wèn)（或預(yù)?。r(shí)填充。換言之，在 GPU 或 CPU 訪問(wèn)頁(yè)和頁(yè)表項(xiàng)之前，可能無(wú)法創(chuàng)建它們。頁(yè)面可以在任何時(shí)候?qū)θ魏翁幚砥鞯膬?nèi)存進(jìn)行 cudaMemPrefetchAsync（） 速率，驅(qū)動(dòng)程序使用啟發(fā)式來(lái)維護(hù)數(shù)據(jù)的局部性并防止過(guò)多的頁(yè)面錯(cuò)誤 3 。（注意：應(yīng)用程序可以使用 cudaMemAdvise（） 指導(dǎo)驅(qū)動(dòng)程序，并使用 MIG 顯式地 MIG 對(duì)內(nèi)存進(jìn)行速率調(diào)整，如 這篇博文描述了 ）。

與 pre-PascalGPUs 不同， Tesla P100 支持硬件頁(yè)錯(cuò)誤和 MIG 比率。所以在這種情況下，運(yùn)行庫(kù)在運(yùn)行內(nèi)核之前不會(huì)自動(dòng)將 全部的 頁(yè)面復(fù)制回 GPU 。內(nèi)核在沒(méi)有任何 MIG 定額開(kāi)銷的情況下啟動(dòng)，當(dāng)它訪問(wèn)任何缺失的頁(yè)時(shí)， GPU 會(huì)暫停訪問(wèn)線程的執(zhí)行，頁(yè)面 MIG 定額引擎 MIG 會(huì)在恢復(fù)線程之前對(duì)設(shè)備的頁(yè)面進(jìn)行評(píng)級(jí)。

這意味著當(dāng)我在 Tesla P100 （ 2 。 1192ms ）上運(yùn)行程序時(shí)， MIG 定額的成本包含在內(nèi)核運(yùn)行時(shí)中。在這個(gè)內(nèi)核中，數(shù)組中的每一頁(yè)都由 CPU 寫入，然后由 GPU 上的 CUDA 內(nèi)核訪問(wèn)，導(dǎo)致內(nèi)核等待大量的頁(yè) MIG 配額。這就是為什么分析器在像 Tesla P100 這樣的 PascalGPU 上測(cè)量的內(nèi)核時(shí)間更長(zhǎng)。讓我們看看 P100 上程序的完整 nvprof 輸出。

==19278== Profiling application: ./add_grid ==19278== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00% 2.1192ms 1 2.1192ms 2.1192ms 2.1192ms add(int, float*, float*) ==19278== Unified Memory profiling result: Device "Tesla P100-PCIE-16GB (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 146 56.109KB 4.0000KB 988.00KB 8.000000MB 860.5760us Host To Device 24 170.67KB 4.0000KB 0.9961MB 4.000000MB 339.5520us Device To Host 12 - - - - 1.067526ms GPU Page fault groups Total CPU Page faults: 36

如您所見(jiàn)，存在許多主機(jī)到設(shè)備頁(yè)面錯(cuò)誤，降低了 CUDA 內(nèi)核的吞吐量。

我該怎么辦？

在實(shí)際應(yīng)用中， GPU 可能會(huì)在數(shù)據(jù)上執(zhí)行更多的計(jì)算（可能多次），而不需要 CPU 來(lái)接觸它。這個(gè)簡(jiǎn)單代碼中的 MIG 定額開(kāi)銷是由于 CPU 初始化數(shù)據(jù)， GPU 只使用一次。有幾種不同的方法可以消除或更改 MIG 比率開(kāi)銷，從而更準(zhǔn)確地測(cè)量 vector add 內(nèi)核的性能。

將數(shù)據(jù)初始化移動(dòng)到另一個(gè) CUDA 內(nèi)核中的 GPU 。

多次運(yùn)行內(nèi)核，查看平均和最小運(yùn)行時(shí)間。

在運(yùn)行內(nèi)核之前，將數(shù)據(jù)預(yù)取到 GPU 內(nèi)存。

我們來(lái)看看這三種方法。

初始化內(nèi)核中的數(shù)據(jù)

如果我們將初始化從 CPU 移到 GPU ，則add內(nèi)核不會(huì)出現(xiàn)頁(yè)面錯(cuò)誤。這里有一個(gè)簡(jiǎn)單的 CUDA C ++內(nèi)核來(lái)初始化數(shù)據(jù)。我們可以用啟動(dòng)這個(gè)內(nèi)核來(lái)替換初始化x和y的主機(jī)代碼。

__global__ void init(int n, float *x, float *y) { int index = threadIdx.x + blockIdx.x * blockDim.x; int stride = blockDim.x * gridDim.x; for (int i = index; i < n; i += stride) { x[i] = 1.0f; y[i] = 2.0f; } }

當(dāng)我這樣做時(shí)，我在 Tesla P100GPU 的配置文件中看到兩個(gè)內(nèi)核：

==44292== Profiling application: ./add_grid_init ==44292== Profiling result: Time(%) Time Calls Avg Min Max Name 98.06% 1.3018ms 1 1.3018ms 1.3018ms 1.3018ms init(int, float*, float*) 1.94% 25.792us 1 25.792us 25.792us 25.792us add(int, float*, float*) ==44292== Unified Memory profiling result: Device "Tesla P100-PCIE-16GB (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 24 170.67KB 4.0000KB 0.9961MB 4.000000MB 344.2880us Device To Host 16 - - - - 551.9940us GPU Page fault groups Total CPU Page faults: 12

add內(nèi)核現(xiàn)在運(yùn)行得更快： 25 . 8us ，相當(dāng)于接近 500gb / s 。

帶寬=字節(jié)/秒=（ 3 * 4194304 字節(jié)* 1e-9 字節(jié)/ GB ）/ 25 . 8e-6s = 488 [UNK] GB / s

（要了解如何計(jì)算理論帶寬和實(shí)現(xiàn)的帶寬，請(qǐng)參閱這個(gè)帖子。）仍然存在設(shè)備到主機(jī)頁(yè)錯(cuò)誤，但這是由于在程序末尾檢查 CPU 結(jié)果的循環(huán)造成的。

運(yùn)行多次

另一種方法是只運(yùn)行內(nèi)核多次，并查看探查器中的平均時(shí)間。為此，我需要修改錯(cuò)誤檢查代碼，以便正確報(bào)告結(jié)果。以下是在 Tesla P100 上 100 次運(yùn)行內(nèi)核的結(jié)果：

==48760== Profiling application: ./add_grid_many ==48760== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00% 4.5526ms 100 45.526us 24.479us 2.0616ms add(int, float*, float*) ==48760== Unified Memory profiling result: Device "Tesla P100-PCIE-16GB (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 174 47.080KB 4.0000KB 0.9844MB 8.000000MB 829.2480us Host To Device 24 170.67KB 4.0000KB 0.9961MB 4.000000MB 339.7760us Device To Host 14 - - - - 1.008684ms GPU Page fault groups Total CPU Page faults: 36

最短的內(nèi)核運(yùn)行時(shí)間只有 24 . 5 微秒，這意味著它可以獲得超過(guò) 500GB / s 的內(nèi)存帶寬。我還包括了來(lái)自nvprof的統(tǒng)一內(nèi)存分析輸出，它顯示了從主機(jī)到設(shè)備總共 8MB 的頁(yè)面錯(cuò)誤，對(duì)應(yīng)于第一次運(yùn)行add時(shí)通過(guò)頁(yè)面錯(cuò)誤復(fù)制到設(shè)備上的兩個(gè) 4MB 數(shù)組（x和y）。

預(yù)取

第三種方法是在初始化后使用統(tǒng)一內(nèi)存預(yù)取將數(shù)據(jù)移動(dòng)到 GPU 。 CUDA 為此提供了cudaMemPrefetchAsync()。我可以在內(nèi)核啟動(dòng)之前添加以下代碼。

 // Prefetch the data to the GPU int device = -1; cudaGetDevice(&device); cudaMemPrefetchAsync(x, N*sizeof(float), device, NULL); cudaMemPrefetchAsync(y, N*sizeof(float), device, NULL); // Run kernel on 1M elements on the GPU int blockSize = 256; int numBlocks = (N + blockSize - 1) / blockSize; saxpy<<>>(N, 1.0f, x, y);

現(xiàn)在當(dāng)我在 Tesla P100 上評(píng)測(cè)時(shí)，我得到以下輸出。

==50360== Profiling application: ./add_grid_prefetch ==50360== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00% 26.112us 1 26.112us 26.112us 26.112us add(int, float*, float*) ==50360== Unified Memory profiling result: Device "Tesla P100-PCIE-16GB (0)" Count Avg Size Min Size Max Size Total Size Total Time Name 4 2.0000MB 2.0000MB 2.0000MB 8.000000MB 689.0560us Host To Device 24 170.67KB 4.0000KB 0.9961MB 4.000000MB 346.5600us Device To Host Total CPU Page faults: 36

在這里，您可以看到內(nèi)核只運(yùn)行了一次，運(yùn)行時(shí)間為 26 。 1us ，與前面顯示的 100 次運(yùn)行中最快的一次相似。您還可以看到，不再報(bào)告任何 GPU 頁(yè)錯(cuò)誤，主機(jī)到設(shè)備的傳輸顯示為四個(gè) 2MB 的傳輸，這要?dú)w功于預(yù)取。

現(xiàn)在我們已經(jīng)讓它在 P100 上運(yùn)行得很快，讓我們將它添加到上次的結(jié)果表中。

關(guān)于并發(fā)性的注記

請(qǐng)記住，您的系統(tǒng)有多個(gè)處理器同時(shí)運(yùn)行 CUDA 應(yīng)用程序的部分：一個(gè)或多個(gè) CPU 和一個(gè)或多個(gè) GPUs 。即使在我們這個(gè)簡(jiǎn)單的例子中，也有一個(gè) CPU 線程和一個(gè) GPU 執(zhí)行上下文，因此在訪問(wèn)任何一個(gè)處理器上的托管分配時(shí)都要小心，以確保沒(méi)有競(jìng)爭(zhēng)條件。

從計(jì)算能力低于 6 。 0 的 CPU 和 GPUs 同時(shí)訪問(wèn)托管內(nèi)存是不可能的。這是因?yàn)?pre-Pascal GPUs 缺少硬件頁(yè)面錯(cuò)誤，所以不能保證一致性。在這些 GPUs 上，內(nèi)核運(yùn)行時(shí)從 CPU 訪問(wèn)將導(dǎo)致分段錯(cuò)誤。

在 Pascal 和更高版本的 GPUs 上， CPU 和 GPU 可以同時(shí)訪問(wèn)托管內(nèi)存，因?yàn)樗鼈兌伎梢蕴幚眄?yè)錯(cuò)誤；但是，由應(yīng)用程序開(kāi)發(fā)人員來(lái)確保不存在由同時(shí)訪問(wèn)引起的爭(zhēng)用條件。

在我們的簡(jiǎn)單示例中，我們?cè)趦?nèi)核啟動(dòng)后調(diào)用了 cudaDeviceSynchronize（） 。這可以確保內(nèi)核在 CPU 嘗試從托管內(nèi)存指針讀取結(jié)果之前運(yùn)行到完成。否則， CPU 可能會(huì)讀取無(wú)效數(shù)據(jù)（在 Pascal 和更高版本上），或獲得分段錯(cuò)誤（在 pre-Pascal GPUs ）。

Pascal 及更高版本上統(tǒng)一內(nèi)存的好處 GPUs

從 PascalGPU 體系結(jié)構(gòu)開(kāi)始，通過(guò) 49 位虛擬尋址和按需分頁(yè) GPU 比率，統(tǒng)一內(nèi)存功能得到了顯著改善。 49 位虛擬地址足以使 GPUs 訪問(wèn)整個(gè)系統(tǒng)內(nèi)存加上系統(tǒng)中所有 GPUs 的內(nèi)存。頁(yè)面 MIG 比率引擎允許 GPU 線程在非駐留內(nèi)存訪問(wèn)時(shí)出現(xiàn)故障，因此系統(tǒng)可以根據(jù)需要從系統(tǒng)中的任何位置對(duì) MIG 的內(nèi)存中的頁(yè)面進(jìn)行 MIG 分級(jí)，以實(shí)現(xiàn)高效處理。

允許使用統(tǒng)一內(nèi)存 cudaMallocManaged（） 對(duì)統(tǒng)一內(nèi)存進(jìn)行分配。無(wú)論是在一個(gè) GPU 上運(yùn)行還是在多個(gè) GPU 上運(yùn)行，它都不會(huì)對(duì)應(yīng)用程序進(jìn)行任何修改。

另外， Pascal 和 VoltaGPUs 支持系統(tǒng)范圍的原子內(nèi)存操作。這意味著您可以對(duì)系統(tǒng)中任何地方的多個(gè) GPUs 值進(jìn)行原子操作。這對(duì)于編寫高效的 multi-GPU 協(xié)作算法非常有用。

請(qǐng)求分頁(yè)對(duì)于以稀疏模式訪問(wèn)數(shù)據(jù)的應(yīng)用程序尤其有利。在某些應(yīng)用程序中，不知道特定處理器將訪問(wèn)哪些特定內(nèi)存地址。如果沒(méi)有硬件頁(yè)面錯(cuò)誤，應(yīng)用程序只能預(yù)加載整個(gè)陣列，或者承受設(shè)備外訪問(wèn)的高延遲成本（也稱為“零拷貝”）。但是頁(yè)面錯(cuò)誤意味著只有內(nèi)核訪問(wèn)的頁(yè)面需要被 MIG 評(píng)級(jí)。

關(guān)于作者

Mark Harris 是 NVIDIA 杰出的工程師，致力于 RAPIDS 。 Mark 擁有超過(guò) 20 年的 GPUs 軟件開(kāi)發(fā)經(jīng)驗(yàn)，從圖形和游戲到基于物理的模擬，到并行算法和高性能計(jì)算。當(dāng)他還是北卡羅來(lái)納大學(xué)的博士生時(shí)，他意識(shí)到了一種新生的趨勢(shì)，并為此創(chuàng)造了一個(gè)名字： GPGPU （圖形處理單元上的通用計(jì)算）。

審核編輯：郭婷