我的第一份CUDA代碼

來源丨h(huán)ttps://zhuanlan.zhihu.com/p/507678214編輯丨極市平臺1. 前言

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

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

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

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

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

首先是一份簡單的C++代碼，主要的運(yùn)行函數(shù)為add函數(shù)，該函數(shù)實(shí)現(xiàn)功能為30M次的for循環(huán)，每次循環(huán)進(jìn)行一次加法。

// 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;
}

編譯以及運(yùn)行代碼：

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

不出意外的話，你應(yīng)該得到下面的結(jié)果：

第一行表示add函數(shù)的運(yùn)行時間，第二行表示每個for循環(huán)里的計(jì)算是否與預(yù)期結(jié)果一致。

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

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

1.首先需要做的是將add函數(shù)變?yōu)镚PU可運(yùn)行函數(shù)，在CUDA中稱為kernel，為此，僅需將變量聲明符添加到函數(shù)中，告訴 CUDA C++ 編譯器這是一個在 GPU 上運(yùn)行并且可以從 CPU 代碼中調(diào)用的函數(shù)。

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

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

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

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

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

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

cudaFree(x);
cudaFree(y);

其實(shí)就相當(dāng)于C++中的new跟delete。

3. add函數(shù)在GPU端運(yùn)行之后，CPU需要等待cuda上的代碼運(yùn)行完畢，才能對數(shù)據(jù)進(jìn)行讀取，因?yàn)镃UDA內(nèi)核啟動時并未對CPU的線程進(jìn)行固定，需要使用cudaDeviceSynchronize()函數(shù)進(jìn)行同步。

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

// add.cu
#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
// __global__ 變量聲明符，作用是將add函數(shù)變成可以在GPU上運(yùn)行的函數(shù)
// __global__ 函數(shù)被稱為kernel，
// 在 GPU 上運(yùn)行的代碼通常稱為設(shè)備代碼（device code），而在 CPU 上運(yùn)行的代碼是主機(jī)代碼（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
  // 內(nèi)存分配，在GPU或者CPU上統(tǒng)一分配內(nèi)存
  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, 執(zhí)行配置
  add<<<1, 1>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  // CPU需要等待cuda上的代碼運(yùn)行完畢，才能對數(shù)據(jù)進(jìn)行讀取
  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對程序進(jìn)行編譯并運(yùn)行：

nvcc add.cu -o add_cuda 
./add_cuda

或者使用nvprof進(jìn)行速度測試：

nvprof ./add_cuda

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

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

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

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

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

修改的部分也比較好理解，不贅述了，接下來運(yùn)行看看結(jié)果：

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

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

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

當(dāng)然因?yàn)樽兂闪硕鄠€`block`，所以此時add函數(shù)需要再改一下：

__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的計(jì)算可以參考上面GPU結(jié)構(gòu)圖以及下面的圖（圖取自An Even Easier Introduction to CUDA | NVIDIA Technical Blog），自行推算，較好理解。

搞定之后再編譯運(yùn)行一下：

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

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

本文代碼：

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

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

參考文獻(xiàn)

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

本文僅做學(xué)術(shù)分享，如有侵權(quán)，請聯(lián)系刪文。