為什么深度學習模型在GPU上運行更快？
引言
當前，提到深度學習，我們很自然地會想到利用GPU來提升運算效率。GPU最初是為了加速圖像渲染和2D、3D圖形處理而設計的。但它們強大的并行處理能力，使得它們在深度學習等更廣泛的領域中也發揮了重要作用。
深度學習模型開始采用GPU是在2000年代中期到晚期，到了2012年，隨著AlexNet的誕生，這種使用變得極為普遍。AlexNet是由Alex Krizhevsky、Ilya Sutskever和Geoffrey Hinton共同設計的卷積神經網絡，它在2012年的ImageNet大規模視覺識別挑戰賽（ILSVRC）中獲勝。這一勝利不僅證明了深度神經網絡在圖像分類上的巨大潛力，也展示了使用GPU進行大型模型訓練的優勢。
自那以后，使用GPU進行深度學習模型訓練變得日益流行，這也催生了PyTorch和TensorFlow等框架的誕生。如今，我們只需在PyTorch中簡單地寫上.to("cuda")，即可將數據傳輸至GPU，期待訓練過程能夠更快。但深度學習算法是如何在實際中利用GPU的計算能力的呢？讓我們一探究竟。
深度學習架構，如神經網絡、卷積神經網絡（CNNs）、循環神經網絡（RNNs）和變換器（transformers），本質上是通過矩陣加法、矩陣乘法以及對矩陣應用函數等數學運算構建的。如果我們能夠優化這些運算，就能提升深度學習模型的效率。
讓我們從基礎開始。設想你需要將兩個向量A和B相加得到向量C，即C = A + B。
圖片
在 C 中的一個簡單實現是：
void AddTwoVectors(flaot A[], float B[], float C[]) {
for (int i = 0; i < N; i++) {
C[i] = A[i] + B[i];
}
}你可能會注意到，計算機需要逐個遍歷向量中的元素，每次迭代都依次將一對元素相加。這些加法操作是獨立進行的，即對第i個元素對的加法并不依賴于其他任何元素對。那么，如果我們能夠同時進行這些操作，一次性并行地完成所有元素對的加法，又會如何呢？
一種簡單的解決方案是利用CPU的多線程功能，來并行處理所有的計算任務。但是，在處理深度學習模型時，我們面對的是包含數百萬元素的大型向量。一般CPU能夠同時處理的線程數量大約只有十幾個。
這時，GPU的優勢就顯現出來了！現代GPU能夠同時執行數百萬的線程，極大地提升了對這些龐大向量進行數學運算的效率。
GPU 與 CPU 比較
雖然CPU在單個操作的速度上可能超過GPU，但GPU的真正優勢在于其強大的并行處理功能。這背后的原因在于兩者設計初衷的差異。CPU的設計宗旨是盡可能快速地完成一系列操作序列，它能夠同時處理的線程數量有限，大約只有幾十個；相比之下，GPU的設計宗旨是為了能夠同時執行數百萬條線程，即便這意味著犧牲了單個線程的執行速度。
舉個例子，我們可以把CPU比作一輛法拉利跑車，而GPU則相當于一輛大巴。如果你只需要運送一個人，那么法拉利（CPU）無疑是更佳的選擇。但如果你的任務是運送一群人，盡管法拉利（CPU）每次運送的速度更快，但大巴（GPU）卻能夠一次性將所有人送達，這樣一次性完成運輸的速度，要比法拉利多次往返運送要快得多。所以，CPU更適合執行順序串行操作，而GPU則更擅長處理并行操作。
為了實現更強的并行處理功能，GPU在設計時將更多的晶體管資源用于執行數據處理任務，而不是像CPU那樣，將大量晶體管用于數據緩存和流程控制，這樣做是為了提升單線程的處理速度和復雜指令的執行效率。
下面的圖表展示了CPU和GPU在芯片資源分配上的差異。
圖片
CPU配備了功能強大的核心和更為復雜的緩存內存結構（為此投入了大量的晶體管資源）。這樣的設計讓CPU在處理順序任務時更為迅速。而GPU則側重于擁有眾多核心，以此來達到更高的并行處理水平。
既然我們已經掌握了這些基礎概念，那么在實際應用中，我們該如何發揮這些并行計算的優勢呢？
CUDA簡介
當您啟動某個深度學習模型時，您可能會傾向于選擇像PyTorch或TensorFlow這樣的流行Python庫。但這些庫的底層實際上是在運行C/C++代碼，這是眾所周知的事實。此外，正如我們之前所討論的，您可能會利用GPU來提升處理速度。這就引入了CUDA的概念！CUDA，即Compute Unified Architecture，是NVIDIA為其GPU開發的一個平臺，用于執行通用計算任務。因此，DirectX被游戲引擎用于圖形計算，而CUDA則允許開發者將NVIDIA的GPU計算能力整合到他們的應用程序中，不僅限于圖形渲染。
為了實現這一點，CUDA提供了一個基于C/C++的簡潔接口（CUDA C/C++），它能夠訪問GPU的虛擬指令集和一些特定操作，比如在CPU和GPU之間傳輸數據。
在我們深入之前，先來理解一些基本的CUDA編程概念和術語：
- host：指CPU及其內存；
- device：指GPU及其內存；
- kernel：指在設備（GPU）上執行的函數；
在用CUDA編寫的簡單代碼中，程序在host（CPU）上運行，將數據發送至device（GPU），并啟動kernel（函數）在device（GPU）上執行。這些kernel由多個線程并行執行。執行完畢后，結果會從device（GPU）傳回host（CPU）。
現在，讓我們回到添加兩個向量的問題上：
#include <stdio.h>
void AddTwoVectors(flaot A[], float B[], float C[]) {
for (int i = 0; i < N; i++) {
C[i] = A[i] + B[i];
}
}
int main() {
...
AddTwoVectors(A, B, C);
...
}在CUDA C/C++編程環境中，開發者能夠創建被稱為kernels的C/C++函數，這些函數一旦被觸發，就能由N個不同的CUDA線程同時執行N次。
定義一個kernel時，我們用__global__關鍵字來聲明，而執行這個kernel的CUDA線程數量可以通過特殊的<<<...>>>標記來設置：
#include <stdio.h>
// Kernel definition
__global__ void AddTwoVectors(float A[], float B[], float C[]) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main() {
...
// Kernel invocation with N threads
AddTwoVectors<<<1, N>>>(A, B, C);
...
}每個執行核心（thread）在運行核心函數（kernel）時，都會被分配一個獨一無二的核心標識符 threadIdx，這個標識符可以在核心函數內部通過內建變量來獲取。上述代碼實現了兩個大小為N的向量A和B的相加操作，并將相加結果存放到向量C中。你會注意到，與傳統的順序循環處理每一對元素相加的方式不同，CUDA技術允許我們通過并行使用N個核心來同時完成所有這些操作。
但在我們實際運行這段代碼之前，還需要進行一些調整。需要牢記的是，核心函數是在設備（GPU）上執行的。這意味著它使用的所有數據都應當存儲在GPU的內存中。我們可以通過調用CUDA提供的一系列內建函數來完成這一數據的遷移：
#include <stdio.h>
// Kernel definition
__global__ void AddTwoVectors(float A[], float B[], float C[]) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main() {
int N = 1000; // Size of the vectors
float A[N], B[N], C[N]; // Arrays for vectors A, B, and C
...
float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C
// Allocate memory on the device for vectors A, B, and C
cudaMalloc((void **)&d_A, N * sizeof(float));
cudaMalloc((void **)&d_B, N * sizeof(float));
cudaMalloc((void **)&d_C, N * sizeof(float));
// Copy vectors A and B from host to device
cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);
// Kernel invocation with N threads
AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);
// Copy vector C from device to host
cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
}我們不能將變量A、B和C直接傳入核心函數，而應該使用指針。在CUDA編程中，你無法在核心函數調用（標記為<<<...>>>）中直接使用主機上的數組（比如示例中的A、B和C）。核心函數是在設備內存中運行的，因此你需要將設備指針（d_A、d_B和d_C）傳入核心函數，以便它能夠進行操作。
除此之外，我們還需要通過調用cudaMalloc函數在設備上分配內存，并利用cudaMemcpy函數在主機內存和設備內存之間傳輸數據。
現在，我們可以在代碼的最后添加向量A和B的初始化步驟，并在結束時刷新CUDA內存。
#include <stdio.h>
// Kernel definition
__global__ void AddTwoVectors(float A[], float B[], float C[]) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main() {
int N = 1000; // Size of the vectors
float A[N], B[N], C[N]; // Arrays for vectors A, B, and C
// Initialize vectors A and B
for (int i = 0; i < N; ++i) {
A[i] = 1;
B[i] = 3;
}
float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C
// Allocate memory on the device for vectors A, B, and C
cudaMalloc((void **)&d_A, N * sizeof(float));
cudaMalloc((void **)&d_B, N * sizeof(float));
cudaMalloc((void **)&d_C, N * sizeof(float));
// Copy vectors A and B from host to device
cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);
// Kernel invocation with N threads
AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);
// Copy vector C from device to host
cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
}此外，我們在核心函數調用之后，需要加入 cudaDeviceSynchronize(); 這個調用。這個函數的作用是確保主機線程與設備之間的同步。調用此函數后，主機線程會暫停，直到設備上所有先前發出的CUDA命令都執行完畢才會繼續。
此外，重要的是要加入一些CUDA錯誤檢查機制，以便我們能夠發現GPU上的錯誤。如果我們忽略了這些檢查，代碼會持續執行主機線程（即CPU的線程），這將使得發現與CUDA相關的錯誤變得困難。
以下是這兩種技術的實現方法：
#include <stdio.h>
// Kernel definition
__global__ void AddTwoVectors(float A[], float B[], float C[]) {
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main() {
int N = 1000; // Size of the vectors
float A[N], B[N], C[N]; // Arrays for vectors A, B, and C
// Initialize vectors A and B
for (int i = 0; i < N; ++i) {
A[i] = 1;
B[i] = 3;
}
float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C
// Allocate memory on the device for vectors A, B, and C
cudaMalloc((void **)&d_A, N * sizeof(float));
cudaMalloc((void **)&d_B, N * sizeof(float));
cudaMalloc((void **)&d_C, N * sizeof(float));
// Copy vectors A and B from host to device
cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);
// Kernel invocation with N threads
AddTwoVectors<<<1, N>>>(d_A, d_B, d_C);
// Check for error
cudaError_t error = cudaGetLastError();
if(error != cudaSuccess) {
printf("CUDA error: %s\n", cudaGetErrorString(error));
exit(-1);
}
// Waits untill all CUDA threads are executed
cudaDeviceSynchronize();
// Copy vector C from device to host
cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
}為了編譯和執行CUDA程序，首先得保證你的計算機上已經安裝了CUDA工具集。接著，你可以利用NVIDIA的CUDA編譯器nvcc來編譯你的代碼。如果你的計算機不具備GPU，你可以考慮使用Google Colab平臺。你只需在“Runtime”菜單下的“Notebook settings”選項中選擇相應的GPU，然后將你的代碼保存為example.cu文件，并執行它。
%%shell
nvcc example.cu -o compiled_example # compile
./compiled_example # run
# you can also run the code with bug detection sanitizer
compute-sanitizer --tool memcheck ./compiled_example不過，我們的代碼優化還有待提升。例如，上述示例中的向量大小僅為N = 1000。這個數值偏小，不足以完全體現GPU的并行處理優勢。在深度學習問題中，我們經常要處理包含數百萬參數的大型向量。如果我們嘗試將N設置為500000，并像之前的例子那樣以<<<1, 500000>>>的方式調用核心函數，會遇到錯誤。因此，為了優化代碼并執行這樣的操作，我們首先需要理解CUDA編程中的一個關鍵概念：線程的層級結構。
線程層次結構
核心函數的調用是通過<<<number_of_blocks, threads_per_block>>>這樣的標記來完成的。比如，在我們之前的例子中，我們執行了1個包含N個CUDA線程的區塊。但是，每個區塊支持的線程數是有上限的。這是因為區塊內的所有線程都需要位于同一個流式多處理器核心上，并且需要共享該核心的內存資源。
你可以通過以下代碼片段來查詢這個上限值：
int device;
cudaDeviceProp props;
cudaGetDevice(&device);
cudaGetDeviceProperties(&props, device);
printf("Maximum threads per block: %d\n", props.maxThreadsPerBlock);在Colab平臺的當前GPU配置中，單個線程塊最多可以包含1024個線程。因此，為了在示例中處理大型向量，我們需要更多的線程塊來執行更多的線程。同時，這些線程塊被進一步組織成更大的結構——網格，就像下面展示的那樣：
現在，可以使用以下方式訪問線程 ID：
int i = blockIdx.x * blockDim.x + threadIdx.x;所以，我們的腳本變成：
#include <stdio.h>
// Kernel definition
__global__ void AddTwoVectors(float A[], float B[], float C[], int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) // To avoid exceeding array limit
C[i] = A[i] + B[i];
}
int main() {
int N = 500000; // Size of the vectors
int threads_per_block;
int device;
cudaDeviceProp props;
cudaGetDevice(&device);
cudaGetDeviceProperties(&props, device);
threads_per_block = props.maxThreadsPerBlock;
printf("Maximum threads per block: %d\n", threads_per_block); // 1024
float A[N], B[N], C[N]; // Arrays for vectors A, B, and C
// Initialize vectors A and B
for (int i = 0; i < N; ++i) {
A[i] = 1;
B[i] = 3;
}
float *d_A, *d_B, *d_C; // Device pointers for vectors A, B, and C
// Allocate memory on the device for vectors A, B, and C
cudaMalloc((void **)&d_A, N * sizeof(float));
cudaMalloc((void **)&d_B, N * sizeof(float));
cudaMalloc((void **)&d_C, N * sizeof(float));
// Copy vectors A and B from host to device
cudaMemcpy(d_A, A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, B, N * sizeof(float), cudaMemcpyHostToDevice);
// Kernel invocation with multiple blocks and threads_per_block threads per block
int number_of_blocks = (N + threads_per_block - 1) / threads_per_block;
AddTwoVectors<<<number_of_blocks, threads_per_block>>>(d_A, d_B, d_C, N);
// Check for error
cudaError_t error = cudaGetLastError();
if (error != cudaSuccess) {
printf("CUDA error: %s\n", cudaGetErrorString(error));
exit(-1);
}
// Wait until all CUDA threads are executed
cudaDeviceSynchronize();
// Copy vector C from device to host
cudaMemcpy(C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
}性能對比
下面對不同向量大小的兩個向量相加運算的 CPU 和 GPU 計算進行了比較。
顯而易見，GPU處理的性能優勢在處理大規模向量N時才會明顯體現出來。此外，需要記住的是，這里的時間比較僅針對核心函數的執行時間，并未包括在主機和設備間傳輸數據所需的時間。雖然在大多數情況下，數據傳輸時間可能并不顯著，但在我們只進行簡單加法操作的情況下，這部分時間卻相對較長。因此，我們必須意識到，GPU在處理那些既計算密集又高度可并行化的計算任務時，才能真正發揮其性能優勢。
多維線程
明白了，我們現在掌握了如何提升基本數組操作效率的方法。但在深度學習模型的實踐中，我們更多地需要處理矩陣和張量的操作。回顧我們之前的示例，我們僅使用了一維區塊，每個區塊包含N個線程。實際上，我們可以執行更高維度的區塊（最多可至三維）。因此，如果你需要進行矩陣運算，可以方便地設置一個NxM的線程區塊。在這種情況下，可以通過row = threadIdx.x和col = threadIdx.y來獲取矩陣的行和列索引。此外，為了簡化操作，可以使用dim3數據類型來指定區塊的數量和每個區塊中的線程數。
以下示例展示了如何實現兩個矩陣的相加操作。
#include <stdio.h>
// Kernel definition
__global__ void AddTwoMatrices(float A[N][N], float B[N][N], float C[N][N]) {
int i = threadIdx.x;
int j = threadIdx.y;
C[i][j] = A[i][j] + B[i][j];
}
int main() {
...
// Kernel invocation with 1 block of NxN threads
dim3 threads_per_block(N, N);
AddTwoMatrices<<<1, threads_per_block>>>(A, B, C);
...
}您還可以擴展此示例以處理多個塊：
#include <stdio.h>
// Kernel definition
__global__ void AddTwoMatrices(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 with 1 block of NxN threads
dim3 threads_per_block(32, 32);
dim3 number_of_blocks((N + threads_per_block.x - 1) ∕ threads_per_block.x, (N + threads_per_block.y - 1) ∕ threads_per_block.y);
AddTwoMatrices<<<number_of_blocks, threads_per_block>>>(A, B, C);
...
}您可以按照這個示例的思路，進一步擴展到處理三維數據的操作。
既然您已經掌握了多維數據的操作方式，接下來要學習另一個既重要又簡單的概念：在核心函數內部如何調用函數。這通常是通過__device__關鍵字來實現的。使用__device__關鍵字定義的函數可以直接在設備（即GPU）上調用。這意味著，這些函數只能在__global__核心函數或其他__device__函數中被調用。以下示例展示了如何在向量上應用sigmoid函數——這是深度學習模型中非常普遍的一種操作。
#include <math.h>
// Sigmoid function
__device__ float sigmoid(float x) {
return 1 / (1 + expf(-x));
}
// Kernel definition for applying sigmoid function to a vector
__global__ void sigmoidActivation(float input[], float output[]) {
int i = threadIdx.x;
output[i] = sigmoid(input[i]);
}明白了CUDA編程的基礎關鍵概念后，您就可以著手編寫CUDA核心函數了。對于深度學習模型，它們通常包含一系列矩陣和張量操作，比如求和、乘法、卷積、歸一化等操作。以矩陣乘法為例，一個簡單的算法可以通過以下方式實現并行處理：
// GPU version
__global__ void matMul(float A[M][N], float B[N][P], float C[M][P]) {
int row = blockIdx.x * blockDim.x + threadIdx.x;
int col = blockIdx.y * blockDim.y + threadIdx.y;
if (row < M && col < P) {
float C_value = 0;
for (int i = 0; i < N; i++) {
C_value += A[row][i] * B[i][col];
}
C[row][col] = C_value;
}
}現在將其與下面兩個矩陣乘法的普通 CPU 實現進行比較：
// CPU version
void matMul(float A[M][N], float B[N][P], float C[M][P]) {
for (int row = 0; row < M; row++) {
for (int col = 0; col < P; col++) {
float C_value = 0;
for (int i = 0; i < N; i++) {
C_value += A[row][i] * B[i][col];
}
C[row][col] = C_value;
}
}
}您可以注意到，在 GPU 版本上，我們的循環更少，從而可以更快地處理操作。下面是CPU和GPU在NxN矩陣乘法上的性能比較：
正如您所觀察到的，隨著矩陣大小的增加，矩陣乘法運算的 GPU 處理性能提升甚至更高。
現在，考慮一個基本的神經網絡，它主要涉及 y = σ(Wx + b) 操作，如下所示：
這些操作主要包括矩陣乘法、矩陣加法以及將函數應用于數組，所有這些操作您都已經熟悉了并行化技術。因此，您現在能夠從頭開始實現在 GPU 上運行的您自己的神經網絡！
總結
本文[1]我們探討了提升深度學習模型性能的GPU處理基礎知識。PyTorch和TensorFlow等庫應用了包含優化內存訪問、批量處理等更高級概念的優化技術（它們使用了在CUDA基礎上構建的庫，比如cuBLAS和cuDNN）。希望本文能夠幫助你理解當你執行.to("cuda")并利用GPU運行深度學習模型時，背后所發生的機制。
