注意:本博文最初發布于 2017 年 1 月 25 日,但已進行編輯以反映新的更新。
本文非常簡單地介紹了 CUDA,這是 NVIDIA 的熱門并行計算平臺和編程模型。我在 2013 年寫過一篇文章,名為“ An Easy Introduction to CUDA ”,多年來一直備受歡迎。但是,CUDA 編程變得更加簡單,GPU 也變得更快了,所以現在是時候更新 (甚至更輕松) 介紹了。
CUDA C++ 只是使用 CUDA 創建大規模并行應用程序的多種方法之一。它允許您使用功能強大的 C++ 編程語言來開發由 GPU 上運行的數千個并行線程加速的高性能算法。許多開發者都以這種方式加速了需要大量計算和帶寬的應用程序,包括支持人工智能持續革命 (即 Deep Learning ) 的庫和框架。
您聽說過 CUDA,并且有興趣學習如何在自己的應用中使用 CUDA。如果您是 C++ 程序員,這篇博文應該會為您提供一個良好的開端。為此,您需要一臺配備支持 CUDA 的 GPU ( Windows、WSL 或 64 位 Linux,任何 NVIDIA GPU 都應配備) 的計算機,或配備 GPU 的云實例 ( AWS、Azure、Google Colab 和其他云服務提供商均配備此類 GPU ) 。您還需要安裝免費的 CUDA 工具包 。
我們開始吧!
從簡單開始
我們將從一個簡單的 C++ 程序開始,該程序可添加兩個數組的元素,每個數組包含一百萬個元素。
#include <iostream> #include <math.h> // function to add the elements of two arrays void add( int n, float *x, float *y) { for ( int i = 0; i < n; i++) y[i] = x[i] + y[i]; } int main( void ) { int N = 1<<20; // 1M elements float *x = new float [N]; float *y = new float [N]; // initialize x and y arrays on the host for ( int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; } // Run kernel on 1M elements on the CPU add(N, x, y); // 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 delete [] x; delete [] y; return 0; } |
首先,編譯并運行此 C++ 程序。將上述代碼放入文件并另存為 add.cpp,然后使用 C++ 編譯器進行編譯。我使用的是 Linux,所以我使用的是 g++,但您可以在 Windows 上使用 MSVC (或在 WSL 上使用 g++) 。
> g++ add.cpp -o add |
然后運行:
> ./add Max error: 0.000000 |
(在 Windows 上,您可能希望為可執行文件 add.exe
命名并使用 .\add
運行。)
與預期一樣,它會打印求和中沒有錯誤,然后退出。現在,我想在 GPU 的多個核心上 (并行) 運行這種計算。邁出第一步其實非常簡單。
首先,我只需將 add
函數轉換為 GPU 可以運行的函數,即 CUDA 中的 kernel 函數。為此,我所要做的就是將說明符 __global__
添加到函數中,以告知 CUDA C++ 編譯器這是一個在 GPU 上運行的函數,可以從 CPU 代碼中調用。
// Kernel function to add the elements of two arrays __global__ void add( int n, float *sum, float *x, float *y) { for ( int i = 0; i < n; i++) sum[i] = x[i] + y[i]; } |
此 __global__ function
稱為 CUDA 內核 ,在 GPU 上運行。 在 GPU 上運行的代碼通常稱為設備代碼 ,而 在 CPU 上運行的代碼則是主機代碼 。
CUDA 中的顯存分配
要在 GPU 上進行計算,我需要分配 GPU 可訪問的內存。CUDA 中的 Unified Memory 通過為系統中的所有 GPU 和 CPU 提供可訪問的單一內存空間來簡化這一過程。要在 Unified Memory 中分配數據,請調用 cudaMallocManaged()
,它會返回可從主機 (CPU) 代碼或設備 (GPU) 代碼訪問的指針。要釋放數據,只需將指針傳遞給 cudaFree()
即可。
我只需要將上面代碼中對 new 的調用替換為對 cudaMallocManaged()
的調用,并將對 delete []
的調用替換為對 cudaFree
的調用。
// Allocate Unified Memory -- accessible from CPU or GPU float *x, *y, *sum; cudaMallocManaged(&x, N* sizeof ( float )); cudaMallocManaged(&y, N* sizeof ( float )); ... // Free memory cudaFree(x); cudaFree(y); |
最后,我需要啟動 add()
內核,以在 GPU 上調用它。CUDA 核函數啟動使用三重角度括號語法 <<< >>>
指定。我只需將其添加到參數列表之前的 add
調用中。
add<<<1, 1>>>(N, sum, x, y); |
簡單!我很快就會詳細介紹角括號內的內容;現在您只需要知道這一行啟動一個 GPU 線程來運行 add()
。
還有一件事:我需要 CPU 等到 kernel 完成后再訪問結果 (因為 CUDA kernel 啟動不會阻塞調用 CPU 線程) 。為此,我只需調用 cudaDeviceSynchronize()
,然后再在 CPU 上進行最后一次錯誤檢查。
以下是完整代碼:
#include <iostream> #include <math.h> // Kernel function to add the elements of two arrays __global__ void add( int n, float *x, float *y) { for ( int i = 0; i < n; i++) 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; } // Run kernel on 1M elements on the GPU add<<<1, 1>>>(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; } |
CUDA 文件的文件擴展名為 .cu
。將此代碼保存在名為 add.cu
的文件中,并使用 CUDA C++ 編譯器 nvcc
進行編譯。
> nvcc add.cu -o add_cuda > ./add_cuda Max error: 0.000000 |
這只是第一步,因為正如書面所述,此 kernel 僅適用于單個線程,因為運行此 kernel 的每個線程都將對整個數組執行添加操作。此外,由于多個并行線程會讀寫相同的位置,因此存在 race condition。
注意:在 Windows 上,您需要確保在 Microsoft Visual Studio 的項目配置屬性中將 Platform 設置為 x64。
分析它!
要了解內核的運行時間,一個很好的方法是使用 NSight Systems CLI `nsys` 運行內核。我們可以在命令行中輸入 nsys profile -t cuda --stats=true ./add_cuda
。但是,這會生成詳細的統計數據,在本文中,我們只想了解內核運行所需的時間。因此,我編寫了一個名為 nsys_easy
的包裝器腳本,該腳本僅生成我們需要的輸出,并阻止 nsys 生成會使源目錄混亂的中間文件。 該腳本可在 GitHub 上獲取。 只需下載 nsys_easy
,并將其放在 PATH (甚至當前目錄) 中的某個位置即可。(請注意,我從輸出中刪除了一些統計數據,以便更好地適應本網站的寬度。)
> nsys_easy ./add_cuda Max error: 0 Generating '/tmp/nsys-report-bb25.qdstrm' [1/1] [========================100%] nsys_easy.nsys-rep Generated: /home/nfs/mharris/src/even_easier/nsys_easy.nsys-rep Generating SQLite file nsys_easy.sqlite from nsys_easy.nsys-rep Processing 1259 events: [======================================100%] Processing [nsys_easy.sqlite] with [cuda_gpu_sum.py]... ** CUDA GPU Summary (Kernels/MemOps) (cuda_gpu_sum): Time (%) Total Time (ns) Count Category Operation -------- --------------- ----- ----------- -------------------------- 98.5 75,403,544 1 CUDA_KERNEL add( int , float *, float *) 1.0 768,480 48 MEMORY_OPER [ memcpy Unified H2D] 0.5 352,787 24 MEMORY_OPER [ memcpy Unified D2D] |
CUDA GPU Summary 表顯示了對 add
的一次調用。在 NVIDIA T4 GPU 上大約需要 75 毫秒。讓我們通過并行來提高它的速度。
獲取線程
現在,您已使用一個線程運行核函數來執行一些計算,如何使其并行?密鑰采用 CUDA 的 <<<1, 1>>>
語法。這稱為執行配置,它會告知 CUDA 運行時在 GPU 上啟動要使用的并行線程數。這里有兩個參數,我們先來更改第二個參數:線程塊中的線程數。CUDA GPU 使用線程塊 (大小為 32 的倍數) 運行內核;256 個線程可合理選擇。
add<<<1, 256>>>(N, x, y); |
如果我僅在運行代碼時進行此更改,它將為每個線程執行一次計算,而不是將計算分散到并行線程。為此,我需要修改內核。CUDA C++ 提供關鍵字,允許內核獲取正在運行的線程的索引。具體而言,threadIdx.x
包含其塊內當前線程的索引,blockDim.x
包含塊中的線程數。我只需修改循環,即可使用并行線程在數組中大步移動。
__global__ void add( int n, float *x, float *y) { int index = threadIdx.x; int stride = blockDim.x; for ( int i = index; i < n; i += stride) y[i] = x[i] + y[i]; } |
add
函數沒有太大變化。事實上,將 index
設置為 0 并將 stride
設置為 1 會使其在語義上與第一個版本相同。
將文件另存為 add_block.cu
,然后再次編譯并在 nvprof
中運行。在本文的剩余部分中,我將展示輸出中的相關行。
Time (%) Time (ns) Count Category Operation -------- --------- ----- ----------- -------------------------- 79.0 4,221,011 1 CUDA_KERNEL add(int, float *, float *) |
這是一個很大的加速 (75ms down to 4ms),但并不奇怪,因為執行從一個線程擴展到 256 個線程。讓我們繼續提升性能。
突破限制
CUDA GPU 有許多并行處理器,分為 Streaming Multiprocessors (SM) 。每個 SM 可以運行多個并發線程塊,但每個線程塊都在單個 SM 上運行。例如,基于 Turing GPU 架構的 NVIDIA T4 GPU 具有 40 個 SM 和 2560 個 CUDA 核心,每個 SM 可支持多達 1024 個活動線程。為了充分利用所有這些線程,我應該啟動具有多個線程塊的 kernel。
現在,您可能已經猜到執行配置的第一個參數用于指定線程塊的數量。并行線程塊共同構成了所謂的 grid 。由于要處理 N
元素,且每個線程塊有 256 個線程,因此我只需計算線程塊的數量即可獲得至少 N 個線程。我只需將 N
除以塊大小 (如果 N 不是 tg_ 44 的倍數,請注意進行四舍五入) 。
int blockSize = 256; int numBlocks = (N + blockSize - 1) / blockSize; add<<<numBlocks, blockSize>>>(N, x, y); |

我還需要更新 kernel 代碼,以將整個線程塊網格考慮在內。CUDA 提供 gridDim.x
(包含網格中的塊數量) 和 tg_ 48 (包含網格中當前線程塊的索引) 。圖 1 說明了使用 blockDim.x
、gridDim.x
和 tg_ 51 在 CUDA 中索引為數組 (一維) 的方法。其思路是,每個線程通過計算其塊開頭的偏移量 (塊索引乘以塊大小:blockIdx.x * blockDim.x
) 并在塊 (tg_ 53) 中添加線程索引來獲取其索引。代碼 blockIdx.x * blockDim.x + threadIdx.x
是慣用的 CUDA。
__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]; } |
更新后的內核還將 stride
設置為網格中的線程總數 (blockDim.x * gridDim.x
) 。CUDA 內核中的這類循環通常稱為 grid-stride 循環 。
將文件另存為 add_grid.cu
,然后再次編譯并在 nvprof
中運行。
Time (%) Time (ns) Count Category Operation -------- --------- ----- ----------- ---------------------------- 79.6 4,514,384 1 CUDA_KERNEL add(int, float *, float *) |
這很有趣。這種變化并沒有帶來加速,而且可能會略有減緩。這是為什么?如果將計算增加 40 倍 (SM 數量) 并不會縮短總時間,那么計算就不是瓶頸。
統一內存預取
如果我們從分析器中查看完整的 summary table,就會發現瓶頸:
Time (%) Time (ns) Count Category Operation -------- --------- ------ ----------- -------------------------- 79.6 4,514,384 1 CUDA_KERNEL add(int, float *, float *) 14.2 807,245 64 MEMORY_OPER [CUDA memcpy Unified H2D] 6.2 353,201 24 MEMORY_OPER [CUDA memcpy Unified D2H] |
在這里,我們可以看到有 64 次主機到設備 (H2D) 和 24 次設備到主機 (D2H) “統一”memcpy
運算。但代碼中沒有明確的 memcpy
調用。CUDA 中的 Unified Memory 是虛擬內存。單個虛擬內存頁面可能駐留在系統中任何設備 (GPU 或 CPU) 的內存中,并且這些頁面會按需遷移。此程序首先在 for 循環中初始化 CPU 上的數組,然后啟動由 GPU 讀取和寫入數組的內核。由于內核運行時內存頁均為 CPU 駐留,因此存在多個分頁錯誤,并且硬件會在發生錯誤時將分頁遷移到 GPU 顯存。這會導致內存瓶頸,這也是我們沒有看到加速的原因。
遷移成本高昂,因為 page faults 是單獨發生的,并且 GPU 線程在等待 page migration 時停止。因為我知道內核需要哪些內存 (x 和 y 數組),所以我可以使用 prefetching 來確保數據在內核需要之前位于 GPU 上。我在啟動 kernel 之前使用 cudaMemPrefetchAsync()
函數來執行此操作:
// Prefetch the x and y arrays to the GPU cudaMemPrefetchAsync(x, N* sizeof ( float ), 0, 0); cudaMemPrefetchAsync(y, N* sizeof ( float ), 0, 0); |
使用分析器運行此過程會產生以下輸出。現在,kernel 只需不到 50 微秒!
Time (%) Time (ns) Count Category Operation -------- --------- ----- ----------- -------------------------- 63.2 690,043 4 MEMORY_OPER [CUDA memcpy Unified H2D] 32.4 353,647 24 MEMORY_OPER [CUDA memcpy Unified D2H] 4.4 47,520 1 CUDA_KERNEL add(int, float *, float *) |
總結
一次預取數組的所有頁面比單個頁面錯誤快得多。請注意,此更改可使所有版本的 add
程序受益,因此,讓我們將其添加到所有三個版本中,并在分析器中再次運行它們。這是一個匯總表。
版本 | 時間 | 加速與。單線程 | 帶寬 |
單線程 | 91811206 納秒 | 1x | 137 MB/ 秒 |
單塊 ( 256 個線程) | 2049,034 納秒 | 45 倍 | 6 GB/ 秒 |
多個塊 | 47520 納秒 | 1932 倍 | 265 GB/ 秒 |
數據進入內存后,從單個塊到多個塊的加速與 GPU 上的 SM 數量 (40) 成正比。
如您所見,我們可以在 GPU 上實現非常高的帶寬。加法內核具有很高的帶寬限制 ( 265 GB/s 是 T4 峰值帶寬 ( 320 GB/s) 的 80% 以上) ,但 GPU 在密集矩陣線性代數、 深度學習 、圖像和信號處理、物理模擬等高度計算受限的計算方面也表現出色。
練習
為了讓您繼續前進,您可以自己嘗試以下幾點。請在下面的評論區發布您的體驗。
- 瀏覽 CUDA 工具包文檔 。如果您尚未安裝 CUDA,請查看 快速入門指南 和安裝指南。然后瀏覽 編程指南 和 最佳實踐指南 。我們還提供各種架構的調優指南。
- 在內核中試驗 printf () 。嘗試為部分或全部線程打印
threadIdx.x
和blockIdx.x
的值。它們是否按順序打印?為什么或者為什么不呢? - 打印內核中
threadIdx.y
或threadIdx.z
(或blockIdx.y
) 的值。(同樣適用于blockDim
和gridDim
)。它們為何存在?如何讓它們接受非 0 (dim 為 1) 的值?
從哪里開始?
我希望這篇文章能激發您對 CUDA 的興趣,并希望您有興趣了解更多信息并在自己的計算中應用 CUDA C++。如果您有任何疑問或意見,請隨時使用下面的評論部分與我們聯系。
您可以繼續閱讀一系列較早的介紹性博文:
- 如何在 CUDA C++ 中實現 Performance Metrics
- 如何在 CUDA C++ 中查詢設備屬性并處理錯誤
- 如何在 CUDA C++ 中優化數據傳輸
- 如何在 CUDA C++ 中重疊數據傳輸
- 如何在 CUDA C++ 中高效訪問 Global Memory
- 在 CUDA C++ 中使用 Shared Memory
- CUDA C++ 中的高效矩陣轉置
- CUDA C++ 中的 Finite Difference Methods,第 1 部分
- CUDA C++ 中的 Finite Difference Methods,第 2 部分
- 借助 CUDA 在一個周末加速 Ray Tracing
此外,還有一系列反映上述內容的 CUDA Fortran 帖子,首先是 An Easy Introduction to CUDA Fortran 。
NVIDIA Developer 博客 上有關于 CUDA C++ 和其他 GPU 計算主題的大量其他內容,請四處看看!
如果您喜歡這篇博文并想了解更多信息, NVIDIA DLI 提供了一些深度 CUDA 編程課程。
- 如果您是新手,請參閱 現代 CUDA C++ 中的加速計算入門 ,其中提供專用 GPU 資源、更復雜的編程環境、 NVIDIA Nsight Systems 視覺分析器的使用、數十次交互式練習、詳細演示、8 小時以上的材料,以及獲得 DLI 能力證書的能力。
- 有關 Python 程序員,請參閱 使用 CUDA Python 的加速計算基礎知識 。
- 有關更多中級和高級 CUDA 編程材料,請參閱 NVIDIA DLI 自定進度目錄 的 Accelerated Computing 部分。
?