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