CUDA 內存統一分析
關于CUDA 編程的基本知識，如何編寫一個簡單的程序，在內存中分配兩個可供 GPU 訪問的數字數組，然后將它們加在 GPU 上。
本文介紹內存統一，這使得分配和訪問系統中任何處理器上運行的代碼都可以使用的數據變得非常容易， CPU 或 GPU 。

圖 1 .內存統一是可從系統中的任何處理器訪問的單個內存地址空間。
以幾個簡單的“練習”介紹，其中一個練習，運行最近基于 Pascal 的 GPU ，看看會發生什么。
建議這樣做有兩個原因。首先，因為 PascalMIG 如 NVIDIA Titan X 和 NVIDIA Tesla P100 是第一個包含頁 GPUs 定額引擎的 GPUs ，它是內存統一頁錯誤處理和 MIG 比率的硬件支持。第二個原因是提供了一個很好的機會來學習更多的內存統一。
快 GPU ，快內存…對嗎？
正確的！首先，我將重新打印在兩個 NVIDIA 開普勒 GPUs 上運行的結果（一個在筆記本電腦上，一個在服務器上）。

現在嘗試在一個非常快的 Tesla P100 加速器上運行，它基于 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 ：比在筆記本電腦基于開普勒的 GeForceGPU 上運行慢。不過，別灰心，可以解決這個問題的。為了理解這一點，將介紹更多關于內存統一的信息。
下面是要添加的完整代碼，以供參考_網格. cu 從上次開始。
#include #include <math.h> // 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, Nsizeof(float)); cudaMallocManaged(&y, Nsizeof(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<<<numBlocks, blockSize>>>(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; }
對 27-19 行的內存進行初始化。
什么是內存統一？
內存統一是可從系統中的任何處理器訪問的單個內存地址空間（請參見圖 1 ）。這種硬件/軟件技術允許應用程序分配可以從 CPU s 或 GPUs 上運行的代碼讀取或寫入的數據。分配內存統一非常簡單，只需將對 malloc() 或 new 的調用替換為對 cudaMallocManaged() 的調用，這是一個分配函數，返回可從任何處理器訪問的指針（以下為 ptr ）。
cudaError_t cudaMallocManaged(void ptr, size_t size);
當在 CPU 或 GPU 上運行的代碼訪問以這種方式分配的數據（通常稱為 CUDA 管理 數據）， CUDA 系統軟件和/或硬件負責將 MIG 額定內存頁分配給訪問處理器的內存。這里重要的一點是， PascalGPU 體系結構是第一個通過頁面 MIG 比率引擎對虛擬內存頁錯誤處理和 MIG 比率提供硬件支持的架構。基于更老的 kezbr 架構和更為統一的 kezbr 形式的支持。
調用 cudaMallocManaged() 時，開普勒會發生什么？

在具有 pre-PascalGPUs 的系統上，如 Tesla K80 ，調用 cudaMallocManaged() 會分配 size 字節的托管內存 在 GPU 設備上 ，該內存在調用 1 時處于活動狀態。在內部，驅動程序還為分配覆蓋的所有頁面設置頁表條目，以便系統理解這些頁駐留在 GPU 上。
所以，在 Tesla K80GPU （開普勒架構）上運行， x 和 y 最初都完全駐留在 GPU 內存中。然后在第 6 行開始的循環中， CPU 逐步遍歷兩個數組，分別將它們的元素初始化為 1.0f 和 2.0f 。由于這些頁最初駐留在設備存儲器中，所以寫入的每個數組頁的 CPU 上都會發生一個頁錯誤， GPU 驅動程序 MIG 會將設備內存中的頁面分配給 CPU 內存。循環之后，兩個數組的所有頁都駐留在 CPU 內存中。
在初始化 CPU 上的數據之后，程序啟動 add() 內核，將 x 的元素添加到 y 的元素中。
add<<<1, 256>>>(N, x, y);
在 pre-PascalGPUs 上，啟動一個內核后， CUDA 運行時必須 MIG，將以前 MIG 額定為主機內存或另一個 GPU 的所有頁面重新評級到運行內核 2 的設備內存。由于這些老的 GPUs 不能出現分頁錯誤，所有數據都必須駐留在 GPU 以防萬一 上，內核訪問它（即使它不會訪問）。這意味著每次啟動內核時都可能存在 MIG 定額開銷。
在 K80 或 macbookpro 上運行程序時，就會發生這種情況。注意，探查器顯示的內核運行時間與 MIG 定額時間是分開的，因為 MIG 定額發生在內核運行之前。
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
調用 cudaMallocManaged() 時， Pascal 上會發生什么？
在 Pascal 和更高版本的 GPUs 上， cudaMallocManaged() 返回時可能不會物理分配托管內存；它只能在訪問（或預取）時填充。換言之，在 GPU 或 CPU 訪問頁和頁表項之前，可能無法創建它們。頁面可以在任何時候對任何處理器的內存進行 cudaMemPrefetchAsync() 速率，驅動程序使用啟發式來維護數據的局部性并防止過多的頁面錯誤 3 。（注意：應用程序可以使用 cudaMemAdvise() 指示驅動程序，并使用 MIG 顯式地 MIG 對內存進行速率調整，如 這篇博文描述了 ）。
與 pre-PascalGPUs 不同， Tesla P100 支持硬件頁錯誤和 MIG 比率。所以在這種情況下，運行庫在運行內核之前不會自動將 全部的 頁面復制回 GPU 。內核在沒有任何 MIG 定額開銷的情況下啟動，當訪問任何缺失的頁時， GPU 會暫停訪問線程的執行，頁面 MIG 定額引擎 MIG 會在恢復線程之前對設備的頁面進行評級。
這意味著在 Tesla P100 （ 2 . 1192ms ）上運行程序時， MIG 定額的成本包含在內核運行時中。在這個內核中，數組中的每一頁都由 CPU 寫入，然后由 GPU 上的 CUDA 內核訪問，導致內核等待大量的頁 MIG 配額。這就是為什么分析器在像 Tesla P100 這樣的 PascalGPU 上測量的內核時間更長。讓我們看看 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
存在許多主機到設備頁面錯誤，降低了 CUDA 內核的吞吐量。
該怎么辦？
在實際應用中， GPU 可能會在數據上執行更多的計算（可能多次），而不需要 CPU 來接觸它。這個簡單代碼中的 MIG 定額開銷是由于 CPU 初始化數據， GPU 只使用一次。有幾種不同的方法可以消除或更改 MIG 比率開銷，從而更準確地測量 vector add 內核的性能。

1. 將數據初始化移動到另一個 CUDA 內核中的 GPU 。
2. 多次運行內核，查看平均和最小運行時間。
3. 在運行內核之前，將數據預取到 GPU 內存。
   來看看這三種方法。
   初始化內核中的數據
   如果將初始化從 CPU 移到 GPU ，則 add 內核不會出現頁面錯誤。這里有一個簡單的 CUDA C ++內核來初始化數據。可以用啟動這個內核來替換初始化 x 和 y 的主機代碼。
   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; } }
   這樣做時，在 Tesla P100GPU 的配置文件中看到兩個內核：
   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 內核現在運行得更快： 25 . 8us ，相當于接近 500gb / s 。
   帶寬=字節/秒=（ 3 * 4194304 字節* 1e-9 字節/ GB ）/ 25 . 8e-6s = 488 [UNK] GB / s
   仍然存在設備到主機頁錯誤，但這是由于在程序末尾檢查 CPU 結果的循環造成的。
   運行多次
   另一種方法是只運行內核多次，并查看探查器中的平均時間。為此，需要修改錯誤檢查代碼，以便正確報告結果。以下是在 Tesla P100 上 100 次運行內核的結果：
   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
   最短的內核運行時間只有 24 . 5 微秒，這意味著它可以獲得超過 500GB / s 的內存帶寬。還包括了來自 nvprof 的內存統一分析輸出，它顯示了從主機到設備總共 8MB 的頁面錯誤，對應于第一次運行 add 時通過頁面錯誤復制到設備上的兩個 4MB 數組（ x 和 y ）。
   預取
   第三種方法是在初始化后使用內存統一預取將數據移動到 GPU 。 CUDA 為此提供了 cudaMemPrefetchAsync() 。可以在內核啟動之前添加以下代碼。
   // 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
   在這里，您可以看到內核只運行了一次，運行時間為 26 . 1us ，與前面顯示的 100 次運行中最快的一次相似。您還可以看到，不再報告任何 GPU 頁錯誤，主機到設備的傳輸顯示為四個 2MB 的傳輸，這要歸功于預取。
   現在已經讓它在 P100 上運行得很快，將它添加到上次的結果表中。
   

關于并發性的注記
請記住，系統有多個處理器同時運行 CUDA 應用程序的部分：一個或多個 CPU 和一個或多個 GPUs 。即使在這個簡單的例子中，也有一個 CPU 線程和一個 GPU 執行上下文，因此在訪問任何一個處理器上的托管分配時都要小心，以確保沒有競爭條件。
從計算能力低于 6 . 0 的 CPU 和 GPUs 同時訪問托管內存是不可能的。這是因為 pre-Pascal GPUs 缺少硬件頁面錯誤，所以不能保證一致性。在這些 GPUs 上，內核運行時從 CPU 訪問將導致分段錯誤。
在 Pascal 和更高版本的 GPUs 上， CPU 和 GPU 可以同時訪問托管內存，因為它們都可以處理頁錯誤；但是，由應用程序開發人員來確保不存在由同時訪問引起的爭用條件。
在簡單示例中，在內核啟動后調用了 cudaDeviceSynchronize() 。這可以確保內核在 CPU 嘗試從托管內存指針讀取結果之前運行到完成。否則， CPU 可能會讀取無效數據（在 Pascal 和更高版本上），或獲得分段錯誤（在 pre-Pascal GPUs ）。
Pascal 及更高版本上內存統一的好處 GPUs
從 PascalGPU 體系結構開始，通過 49 位虛擬尋址和按需分頁 GPU 比率，內存統一功能得到了顯著改善。 49 位虛擬地址足以使 GPUs 訪問整個系統內存加上系統中所有 GPUs 的內存。頁面 MIG 比率引擎允許 GPU 線程在非駐留內存訪問時出現故障，因此系統可以根據需要從系統中的任何位置對 MIG 的內存中的頁面進行 MIG 分級，以實現高效處理。
允許使用內存統一 cudaMallocManaged() 對內存統一進行分配。無論是在一個 GPU 上運行還是在多個 GPU 上運行，它都不會對應用程序進行任何修改。
另外， Pascal 和 VoltaGPUs 支持系統范圍的原子內存操作。這意味著您可以對系統中任何地方的多個 GPUs 值進行原子操作。這對于編寫高效的 multi-GPU 協作算法非常有用。
請求分頁對于以稀疏模式訪問數據的應用程序尤其有利。在某些應用程序中，不知道特定處理器將訪問哪些特定內存地址。如果沒有硬件頁面錯誤，應用程序只能預加載整個陣列，或者承受設備外訪問的高延遲成本（也稱為“零拷貝”）。但是頁面錯誤意味著只有內核訪問的頁面需要被 MIG 評級。
下一步？
本文幫助繼續學習 CUDA 編程，并且有興趣學習更多，并在計算中應用 CUDA C ++。
有關內存統一預取和使用提示（ cudaMemAdvise() ）的更多信息，請參閱文章
在 Pascal 上使用內存統一超出 GPU 內存限制 。如果想了解使用 cudaMemcpy 和 cudaMemcpy 在 CUDA 中進行顯式內存管理的信息，請參閱文章 CUDA C / C ++的簡單介紹 。
計劃用更多的 CUDA 編程材料來跟進本文，可以繼續閱讀一系列比較老的介紹性文章。
? 如何在 CUDA C ++中實現性能度量
? 如何查詢 CUDA C ++中的設備屬性和處理錯誤
? 如何優化 CUDA C ++中的數據傳輸
? 如何在 CUDA C ++中重疊數據傳輸
? 如何在 CUDA C ++中高效訪問全局內存
? 在 CUDA C ++中使用共享內存
? CUDA C ++中的一種高效矩陣轉置
? CUDA C ++中的有限差分方法，第 1 部分
? CUDA C ++中的有限差分方法，第 2 部分
還有一系列的設備。
從技術上講，這是一種簡化。在帶有 pre-Pascal GPUs 的 multi-GPU 系統上，如果某些 GPUs 禁用了對等訪問，則將分配內存，使其最初駐留在 CPU 上。
嚴格地說，可以使用 cudaStreamAttachMemAsync() 將分配的可見性限制到特定的 CUDA 流。這允許驅動程序 MIG 只對附加到啟動內核的流的頁面進行評級。默認情況下，托管分配附加到所有流，因此任何內核啟動都會觸發 MIG 配額。 請閱讀 CUDA 編程指南中的更多內容 。
設備屬性 concurrentManagedAccess 說明 GPU 是否支持硬件頁 MIG 比率以及它所啟用的并發訪問功能。值為 1 表示支持。目前，它只在運行 64 位 Linux 的 Pascal 和更新的 GPUs 上受支持。

總結

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

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。