• <xmp id="om0om">
  • <table id="om0om"><noscript id="om0om"></noscript></table>
  • 數據科學

    CUDA 入門教程:更簡單的介紹 (更新版)

    注意:本博文最初發布于 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);

    圖 1。CUDA 內核中的 Grid、Block 和 Thread 索引 (一維) 。

    我還需要更新 kernel 代碼,以將整個線程塊網格考慮在內。CUDA 提供 gridDim.x (包含網格中的塊數量) 和 tg_ 48 (包含網格中當前線程塊的索引) 。圖 1 說明了使用 blockDim.xgridDim.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 在密集矩陣線性代數、 深度學習 、圖像和信號處理、物理模擬等高度計算受限的計算方面也表現出色。

    練習

    為了讓您繼續前進,您可以自己嘗試以下幾點。請在下面的評論區發布您的體驗。

    1. 瀏覽 CUDA 工具包文檔 。如果您尚未安裝 CUDA,請查看 快速入門指南 和安裝指南。然后瀏覽 編程指南 最佳實踐指南 。我們還提供各種架構的調優指南。
    2. 在內核中試驗 printf () 。嘗試為部分或全部線程打印 threadIdx.xblockIdx.x 的值。它們是否按順序打印?為什么或者為什么不呢?
    3. 打印內核中 threadIdx.ythreadIdx.z (或 blockIdx.y) 的值。(同樣適用于 blockDimgridDim)。它們為何存在?如何讓它們接受非 0 (dim 為 1) 的值?

    從哪里開始?

    我希望這篇文章能激發您對 CUDA 的興趣,并希望您有興趣了解更多信息并在自己的計算中應用 CUDA C++。如果您有任何疑問或意見,請隨時使用下面的評論部分與我們聯系。

    您可以繼續閱讀一系列較早的介紹性博文:

    此外,還有一系列反映上述內容的 CUDA Fortran 帖子,首先是 An Easy Introduction to CUDA Fortran

    NVIDIA Developer 博客 上有關于 CUDA C++ 和其他 GPU 計算主題的大量其他內容,請四處看看!

    如果您喜歡這篇博文并想了解更多信息, NVIDIA DLI 提供了一些深度 CUDA 編程課程。

    ?

    +6

    標簽

    人人超碰97caoporen国产