【平行運算】CUDA教學(一) 概念介紹

cuda

現在大部份做深度學習的工程師視 CUDA、GPU加速為黑盒子,但若你能了解 CUDA,這就是你比 90% 的人還要厲害的關鍵。

CUDA (Compute Unified Device Architecture,統一計算架構) 是 NVIDIA 研發的平行運算平台及編程模型,可利用繪圖處理單元 (GPU) 的能力大幅提升運算效能。目前最流行的深度學習,在訓練類神經網路時因為牽涉到大量的運算,也是使用 CUDA 等平行運算技術來進行加速。目前主流的 Tensorflow、Pytorch 等深度學習框架也大量了使用 CUDA。也正是因為 GPU 的運算效能,才讓深度學習可以有今天的表現。因此了解 GPU、CUDA、平行運算的技術是非常重要的!

1. CUDA 平行加速工作流程

  1. 在 host(CPU) 配置記憶體並初始化資料
  2. 在 device(GPU) 配置記憶體
  3. 將資料從 host 傳送到 device (CPU->GPU)
  4. 在GPU上執行核心程式碼(Kernel)
  5. 將資料從 device 傳送回 host (GPU->CPU)

上面提到了幾個關鍵字 : 配置記憶體資料傳輸核心程式碼 Kernel,以下將一一詳細介紹

2. 配置記憶體

記憶體的配置是相當簡單,在 CPU 上使用 malloc (memory allocation)來配置記憶體,在 GPU 使用 cudaMalloc (cuda memory allocation) 來配置記憶體。

CPU 記憶體:

  • malloc
  • free
float *Array_CPU;

Array_CPU = (float*)malloc(Array_Size);

free(Array_CPU);

GPU 記憶體:

  • cudaMalloc
  • cudaFree
float *Array_GPU;

cudaMalloc(&Array_GPU, Array_Size);

cudaFree(Array_GPU);

3. 資料傳輸

無論是從 CPU 到 GPU 都是 cudaMemcpy 這個函式,只需要傳入cudaMemcpyHostToDevicecudaMemcpyDeviceToHost 即可指定是將資料從哪裡複製到哪裡

CPU 到 GPU (Host to Device):

  • cudaMemcpy( )
  • cudaMemcpyHostToDevice
cudaMemcpy(Array_GPU, Array_CPU, Array_Size, cudaMemcpyHostToDevice)

GPU 到 CPU (Device to Host):

  • cudaMemcpy( )
  • cudaMemcpyDeviceToHost
cudaMemcpy(Array_CPU, Array_GPU, Array_Size, cudaMemcpyDeviceToHost)

4. 核心程式碼 (Kernel)

  • CUDA 讓我們可以自己定義核心程式碼(kernel),當他被呼叫的時候,會被每個 thread 都執行一次, N 個 thread 執行 N 次
  • Kernel function 在定義時要使用 global
  • Kernel function 在呼叫時要使用 <<<…>>> 並指定 block 數目,和每個 block 裡面有多少 threads
GPU_Kernel <<< blocksPerGrid, threadsPerBlock >>> (N, Array_GPU);

以下的程式碼用了 1 個 block, N 個 threads 來執行向量的加法,向量A + 向量B 並存至 C

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

5. 執行緒層級 (Thread Hierachy)

在 CUDA 中,最基本的運算單元是 thread,很多 thread 組成一個 block,很多 block 組成一個 gird ( block 也可以稱作 thread block,由 thread 組成的 block)

所以由大到小是:

  • Grid
  • Block
  • Thread

GBT、GBT、GBT (幫助記憶)

  • 要知道每一個 thread 獨特的座標以指派運算任務,我們會需要 block 的 index 和 thread 的 index。這些 index 可以是 1D、2D 或 3D,想像成三維座標就可以了
  • 要知道每一個 thread 獨特的座標以指派運算任務,我們會需要 block 的大小。假設一個 block 裡面有 4 x 3 的 threads (像上圖),則 blockDim.x = 4, blockDim.y = 3
//在一個 grid 裡面,一個 thread 的座標(x, y) 為:
x = blockDim.x * blockIdx.x + threadIdx.x
y = blockDim.y * blockIdx.y + threadIdx.y

以下這張圖很清楚的介紹了在 1D 的情況下,thread_id 的計算方式:

(圖片來源:An Even Easier Introduction to CUDA)

以下的程式碼執行了矩陣的加法,使用了大小為 16 x 16 的 block (每個block裡有256個thread),並且使用了 numBlcoks 個 block,這邊 block的數目為動態計算,可以針對給訂的資料來決定要使用多少個 block

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

6. 參考資料

留言討論區