我的第一份CUDA代碼

來源丨https://zhuanlan.zhihu.com/p/507678214編輯丨極市平臺1. 前言

這是一份簡單的CUDA編程入門，主要參考英偉達的官方文檔進行學習，本人也是剛開始學習，如有表述錯誤，還請指出。官方文檔鏈接如下：

https://developer.nvidia.com/blog/even-easier-introduction-cuda/

本文先從一份簡單的C++代碼開始，然后逐步介紹如何將C++代碼轉換為CUDA代碼，以及對轉換前后程序的運行時間進行對比，本文代碼放在我的github中，有需要可以自取。

https://github.com/xcyuyuyu/My-First-CUDA-Code

本文所使用的CPU為i7-4790，GPU為GTX 1080，那就開始吧。

首先是一份簡單的C++代碼，主要的運行函數為add函數，該函數實現功能為30M次的for循環，每次循環進行一次加法。

// add.cpp
#include <iostream>
#include <math.h>
#include <sys/time.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<<25; // 30M 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;
  }

  struct timeval t1,t2;
  double timeuse;
  gettimeofday(&t1,NULL);
  // Run kernel on 30M elements on the CPU
  add(N, x, y);
  gettimeofday(&t2,NULL);
  timeuse = (t2.tv_sec - t1.tv_sec) + (double)(t2.tv_usec - t1.tv_usec)/1000.0;

  std::cout << "add(int, float*, float*) time: " << timeuse << "ms" << std::endl;
  // 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;
}

編譯以及運行代碼：

g++ add.cpp -o add
./add

不出意外的話，你應該得到下面的結果：

第一行表示add函數的運行時間，第二行表示每個for循環里的計算是否與預期結果一致。

這個簡單的C++代碼在CPU端運行，運行時間為85ms，接下來介紹如何將主要運算的add函數遷移至GPU端。

將C++代碼改為CUDA代碼，目的是將add函數的計算過程遷移至GPU端，利用GPU的并行性加速運算，需要修改的地方主要有3處：

1.首先需要做的是將add函數變為GPU可運行函數，在CUDA中稱為kernel，為此，僅需將變量聲明符添加到函數中，告訴 CUDA C++ 編譯器這是一個在 GPU 上運行并且可以從 CPU 代碼中調用的函數。

__global__ 
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    y[i] = x[i] + y[i];
}

那么修改后的add函數的調用也比較簡單，僅需要在add函數名后面加上三角括號語法<<<i,j>>>指定CUDA內核啟動即可，<<<i,j>>>稱為執行配置（execution configuration），用于配置程序運行時的線程，后續會講到，目前先將其設置為<<<i,j>>>：

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

2. 那么為了在GPU進行計算，需要在GPU上分配可訪問的內存。CUDA中通過Unified Memory（統一內存）機制來提供可同時供GPU和CPU訪問的內存，使用cudaMallocManaged()函數進行分配：

cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));

同時，在程序最后使用cudaFree()進行內存釋放：

cudaFree(x);
cudaFree(y);

其實就相當于C++中的new跟delete。

3. add函數在GPU端運行之后，CPU需要等待cuda上的代碼運行完畢，才能對數據進行讀取，因為CUDA內核啟動時并未對CPU的線程進行固定，需要使用cudaDeviceSynchronize()函數進行同步。

4. 整體的程序如下所示：

// add.cu
#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
// __global__ 變量聲明符，作用是將add函數變成可以在GPU上運行的函數
// __global__ 函數被稱為kernel，
// 在 GPU 上運行的代碼通常稱為設備代碼（device code），而在 CPU 上運行的代碼是主機代碼（host code）。
__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<<25;
  float *x, *y;

  // Allocate Unified Memory – accessible from CPU or GPU
  // 內存分配，在GPU或者CPU上統一分配內存
  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
  // execution configuration, 執行配置
  add<<<1, 1>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  // CPU需要等待cuda上的代碼運行完畢，才能對數據進行讀取
  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;
}

使用nvcc對程序進行編譯并運行：

nvcc add.cu -o add_cuda 
./add_cuda

或者使用nvprof進行速度測試：

nvprof ./add_cuda

不出意外的話，你會得到以下輸出：

框出來的就是add函數在GPU端的運行時間，為4s。沒錯，就是比CPU端85ms還要慢，那還學個錘子。

好的回過頭看看，問題出現在這個執行配置 <<<i,j>>> 上。不急，先看一下一個簡單的GPU結構示意圖，按照層次從大到小可將GPU按照 grid -> block -> thread劃分，其中最小單元是thread，并行的本質就是將程序的計算模塊拆分成多個小模塊扔給每個thread并行計算。

再看一下前面執行配置 `<<<i,j>>>` 的含義，`<<<i,j>>>` 應該寫成 `<<<numBlocks, blockSize>>>` ，即表示函數運行時使用的block數量以及每個block的大小，前面我們將其設置為`<<<1,1>>>` ，說明程序是單線程運行的，那當然慢了~~。下面我們以單個block為例，將其改為`<<<1,256>>>`，add函數也需要適當修改：

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x; // threadIdx.x表示當前在第幾個thread上運行
  int stride = blockDim.x; // blockDim.x表示每個block的大小
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

修改的部分也比較好理解，不贅述了，接下來運行看看結果：

你看，開始加速了吧，4s加速到了77ms。

那么，`<<<numBlocks, blockSize>>>` 的兩個參數應該怎么設置好呢。首先，CUDA GPU 使用大小為 32 的倍數的線程塊運行內核，因此 `blockSize` 的大小應該設置為32的倍數，例如128、256、512等。確定 `blockSize` 之后，可以根據for循環的總個數`N`確定 `numBlock` 的大?。ㄗ⒁馑纳嵛迦氲恼`差）：

int numBlock = (N + blockSize - 1) / blockSize;

當然因為變成了多個`block`，所以此時add函數需要再改一下：

__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];
}

這里index跟stride的計算可以參考上面GPU結構圖以及下面的圖（圖取自An Even Easier Introduction to CUDA | NVIDIA Technical Blog），自行推算，較好理解。

搞定之后再編譯運行一下：

看看，又加速了不是，通過提升并行度而加速，相比于CPU端（85ms）加速了接近一倍左右。

以上僅是一份簡單的CUDA入門代碼，看起來還算比較簡單，不過繼續深入肯定有更多的坑，期待后面有時間繼續學習。

本文代碼：

GitHub - xcyuyuyu/My-First-CUDA-Code: The introduction to cuda, a simple and easy cuda project

https://github.com/xcyuyuyu/My-First-CUDA-Code

參考文獻

[1] An Even Easier Introduction to CUDA | NVIDIA Technical Blog（https://developer.nvidia.com/blog/even-easier-introduction-cuda/）

本文僅做學術分享，如有侵權，請聯系刪文。