CUDA統(tǒng)一內(nèi)存分析

PascalMIG 如 NVIDIA Titan X 和 NVIDIA Tesla P100 是第一個包含頁 GPUs 定額引擎的 GPUs ，它是統(tǒng)一內(nèi)存頁錯誤處理和 MIG 比率的硬件支持。提供了一個很好的機(jī)會來學(xué)習(xí)更多的統(tǒng)一內(nèi)存。

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

正確的！

首先，將重新打印在兩個 NVIDIA 開普勒 GPUs 上運(yùn)行的結(jié)果（一個在筆記本電腦上，一個在服務(wù)器上）。

現(xiàn)在嘗試在一個非常快的 Tesla P100 加速器上運(yùn)行，它基于pascalgp100GPU 。

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

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

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

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

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

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

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

調(diào)用 cudaMallocManaged() 時，開普勒會發(fā)生什么？

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

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

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

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

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

當(dāng)在 K80 或 macbookpro 上運(yùn)行程序時，就會發(fā)生這種情況。只是探查器顯示的內(nèi)核運(yùn)行時間與 MIG 定額時間是分開的，因為 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

調(diào)用 cudaMallocManaged() 時， Pascal 上會發(fā)生什么？

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

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

這意味著在 Tesla P100 （ 2 . 1192ms ）上運(yùn)行程序時， MIG 定額的成本包含在內(nèi)核運(yùn)行中。在這個內(nèi)核中，數(shù)組中的每一頁都由 CPU 寫入，然后由 GPU 上的 CUDA 內(nèi)核訪問，導(dǎo)致內(nèi)核等待大量的頁 MIG 配額。這就是為什么分析器在像 Tesla P100 這樣的 PascalGPU 上測量的內(nèi)核時間更長。看看 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

存在許多主機(jī)到設(shè)備頁面錯誤，降低了 CUDA 內(nèi)核的吞吐量。

怎么辦？

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

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

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

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

看看這三種方法。

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

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

global void init(int n, float x, floaty) { 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; } }

在 Tesla P100GPU 的配置文件中看到兩個內(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

仍然存在設(shè)備到主機(jī)頁錯誤，但這是由于在程序末尾檢查 CPU 結(jié)果的循環(huán)造成的。

運(yùn)行多次

另一種方法是只運(yùn)行內(nèi)核多次，并查看探查器中的平均時間。為此，需要修改錯誤檢查代碼，以便正確報告結(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)行時間只有 24 . 5 微秒，意味著可以獲得超過 500GB / s 的內(nèi)存帶寬。來自 nvprof 的統(tǒng)一內(nèi)存分析輸出，它顯示了從主機(jī)到設(shè)備總共 8MB 的頁面錯誤，對應(yīng)于第一次運(yùn)行 add 時通過頁面錯誤復(fù)制到設(shè)備上的兩個 4MB 數(shù)組（ x 和 y ）。

預(yù)取

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

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

在 Tesla P100 上評測時，得到以下輸出。

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)行時間為 26 . 1us ，與前面顯示的 100 次運(yùn)行中最快的一次相似。還可以看到，不再報告任何 GPU 頁錯誤，主機(jī)到設(shè)備的傳輸顯示為四個 2MB 的傳輸，這要?dú)w功于預(yù)取。

在 P100 上運(yùn)行得很快，將它添加到上次的結(jié)果表中。

并發(fā)性的分析

系統(tǒng)有多個處理器同時運(yùn)行 CUDA 應(yīng)用程序的部分：一個或多個 CPU 和一個或多個 GPUs 。這個簡單的例子中，有一個 CPU 線程和一個 GPU 執(zhí)行上下文，因此在訪問任何一個處理器上的托管分配時都要小心，以確保沒有競爭條件。

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

在 Pascal 和更高版本的 GPUs 上， CPU 和 GPU 可同時訪問托管內(nèi)存，都可以處理頁錯誤；但是，由應(yīng)用程序開發(fā)人員來確保不存在由同時訪問引起的競爭條件。

在內(nèi)核啟動后調(diào)用了 cudaDeviceSynchronize() 。可確保內(nèi)核在 CPU 嘗試從托管內(nèi)存指針讀取結(jié)果之前運(yùn)行到完成。否則， CPU 可能會讀取無效數(shù)據(jù)（在 Pascal 和更高版本上），或獲得分段錯誤（在 pre-Pascal GPUs ）。

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

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

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

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

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

總結(jié)

以上是生活随笔為你收集整理的CUDA统一内存分析的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。