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

觀看次數： 38,944

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

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

1. CUDA 平行加速工作流程

在 host(CPU) 配置記憶體並初始化資料
在 device(GPU) 配置記憶體
將資料從 host 傳送到 device (CPU->GPU)
在GPU上執行核心程式碼(Kernel)
將資料從 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 這個函式，只需要傳入cudaMemcpyHostToDevice或cudaMemcpyDeviceToHost 即可指定是將資料從哪裡複製到哪裡

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);
    ...
}

詳細的 Index、Dimension 等觀念可以參考 Programming Guide :: CUDA Toolkit Documentation
因為 blockDim * blockIdx + threadIdx 很常用，所以建議可以記起來