CUDA 编程上手指南：CUDA C 编程及 GPU 基本知识
作者丨科技猛獸
編輯丨極市平臺
本文原創首發于極市平臺，轉載請獲得授權并標明出處。
推薦大家關注極市平臺公眾號，每天都會更新最新的計算機視覺論文解讀、綜述盤點、調參攻略、面試經驗等干貨~
目錄
1 CPU 和 GPU 的基礎知識
2 CUDA 編程的重要概念
3 并行計算向量相加
4 實踐
4.1 向量相加 CUDA 代碼
4.2 實踐向量相加
5 給大家的一點參考資料
1 CPU 和 GPU 的基礎知識
提到處理器結構，有2個指標是經常要考慮的：延遲和吞吐量。所謂延遲，是指從發出指令到最終返回結果中間經歷的時間間隔。而所謂吞吐量，就是單位之間內處理的指令的條數。
下圖1是 CPU 的示意圖。從圖中可以看出 CPU 的幾個特點：
所以綜合以上三點，CPU 在設計時的導向就是減少指令的時延，我們稱之為延遲導向設計，如下圖3所示。
下圖2是 GPU 的示意圖，它與之前 CPU 的示意圖相比有著非常大的不同。從圖中可以看出 GPU 的幾個特點 (注意紫色和黃色的區域分別是緩存單元和控制單元)：
所以，GPU 在設計過程中以一個原則為核心：增加簡單指令的吞吐。因此，我們稱 GPU 為吞吐導向設計，，如下圖3所示。
那么究竟在什么情況下使用 CPU，什么情況下使用 GPU 呢？
CPU 在連續計算部分，延遲優先，CPU 比 GPU ，單條復雜指令延遲快10倍以上。
GPU 在并行計算部分，吞吐優先，GPU 比 CPU ，單位時間內執行指令數量10倍以上。
適合 GPU 的問題：
2 CUDA 編程的重要概念
CUDA (Compute Unified Device Architecture)，由英偉達公司2007年開始推出，初衷是為 GPU 增加一個易用的編程接口，讓開發者無需學習復雜的著色語言或者圖形處理原語。
OpenCL (Open Computing Languge) 是2008年發布的異構平臺并行編程的開放標準，也是一個編程框架。OpenCL 相比 CUDA，支持的平臺更多，除了 GPU 還支持 CPU、DSP、FPGA 等設備。
下面我們將以 CUDA 為例，介紹 GPU 編程的基本思想和基本操作。
首先主機端 (host) 和設備端 (device)，主機端一般指我們的 CPU，設備端一般指我們的 GPU。
一個 CUDA 程序，我們可以把它分成3個部分：
第1部分是： 從主機 (host) 端申請 device memory，把要拷貝的內容從 host memory 拷貝到申請的 device memory 里面。
第2部分是： 設備端的核函數對拷貝進來的東西進行計算，來得到和實現運算的結果，圖4中的 Kernel 就是指在 GPU 上運行的函數。
第3部分是： 把結果從 device memory 拷貝到申請的 host memory 里面，并且釋放設備端的顯存和內存。
CUDA 編程中的內存模型
這里就引出了一個非常重要的概念就是 CUDA 編程中的內存模型。
從硬件的角度來講：
CUDA 內存模型的最基本的單位就是 SP (線程處理器)。每個線程處理器 (SP) 都用自己的 registers (寄存器) 和 local memory (局部內存)。寄存器和局部內存只能被自己訪問，不同的線程處理器之間呢是彼此獨立的。
由多個線程處理器 (SP) 和一塊共享內存所構成的就是 SM (多核處理器) (灰色部分)。多核處理器里邊的多個線程處理器是互相并行的，是不互相影響的。每個多核處理器 (SM) 內都有自己的 shared memory (共享內存)，shared memory 可以被線程塊內所有線程訪問。
再往上，由這個 SM (多核處理器) 和一塊全局內存，就構成了 GPU。一個 GPU 的所有 SM 共有一塊 global memory (全局內存)，不同線程塊的線程都可使用。
上面這段話可以表述為：每個 thread 都有自己的一份 register 和 local memory 的空間。同一個 block 中的每個 thread 則有共享的一份 share memory。此外，所有的 thread (包括不同 block 的 thread) 都共享一份 global memory。不同的 grid 則有各自的 global memory。
從軟件的角度來講：
如下圖6所示，所謂線程塊內存模型在軟件側的一個最基本的執行單位，所以我們從這里開始梳理。線程塊就是線程的組合體，它具有如下這些特點：
如下圖7所示的線程塊就是由256個線程組成的，它執行的任務就是一個最基本的向量相加的一個操作。在線程塊內，這256個線程的計算是彼此互相獨立的，并行的。下面的這個 [i]，就是如何確定每個線程的索引 (在顯存中的位置)。在計算完以后 (圖中彎箭頭的頭部)，會設置一個時鐘，將這256個線程的計算結果進行同步。
以上就是一個256位向量的加的操作的并行處理方法，得到最終的向量加的結果。
所謂網格 (grid)，其實就是線程塊的組合體，如下圖8所示。
CUDA 核函數由線程網格 (數組) 執行。每個線程都有一個索引，用于計算內存地址和做出控制決策。在計算完以后 (圖中所有彎箭頭的頭部)，會設置一個時鐘，將這N個線程塊的計算結果進行同步。
線程塊 id & 線程 id：定位獨立線程的門牌號
核函數需要確定每個線程在顯存中的位置，我們之前提到 CUDA 的核函數是要在設備端來進行計算和處理的，在執行核函數時需要訪問到每個線程的 registers (寄存器) 和 local memory (局部內存)。在這個過程中需要確定每一個線程在顯存上的位置。所以我們需要像圖9那樣使用線程塊的 index 和線程的 index 來確定線程在顯存上的位置。
如圖9所示，圖9中的線程塊索引是2維的，每個網格都由2×2個線程塊組成；線程索引是3維的，每個線程塊都由2×4×2個線程組成，所以代碼應該是：
圖10中：M=N=2，P,Q,S=2,4,2。
每個線程x的那一維應該是線程塊的索引×線程塊的x維度大小+線程的索引。(設備端線程x的那一維的索引)。
每個線程y的那一維應該是線程塊的索引×線程塊的y維度大小+線程的索引。(設備端線程y的那一維的索引)。
線程束 (warp)
前面我們提到，如圖11所示的每一行由1個控制單元加上若干計算單元所組成，這些所有的計算單元執行的控制指令是一個。這其實就是個非常典型的 “單指令多數據流機制”。
單指令多數據流機制是說：執行的指令是一條，只不過不同的計算單元使用的數據是不一樣的。而上面這一行，我們就稱之為一個線程束 (warp)。
所以，SM 采用的 SIMT (Single-Instruction, Multiple-Thread，單指令多線程) 架構，warp (線程束) 是最基本的執行單元。一個 warp 包含32個并行 thread，這些 thread 以不同數據資源執行相同的指令。一個 warp 只包含一條指令，所以：warp 本質上是線程在 GPU 上運行的最小單元。
由于warp的大小為32，所以block所含的thread的大小一般要設置為32的倍數。
當一個 kernel 被執行時，grid 中的線程塊被分配到 SM (多核處理器) 上，一個線程塊的 thread 只能在一個SM 上調度，SM 一般可以調度多個線程塊，大量的 thread 可能被分到不同的 SM 上。每個 thread 擁有它自己的程序計數器和狀態寄存器，并且用該線程自己的數據執行指令，這就是所謂的 Single Instruction Multiple Thread (SIMT)，如圖12所示。
3 并行計算向量相加
下面我們就用一個實際的例子來看看 CUDA 編程具體是如何操作的。例子就是兩個長度為N的張量相加，如下圖13所示。
在 CPU 中完成相加的操作很簡單：
// Compute vector sum C = A+B void vecAdd(float* A, float* B, float* C, int n) { for (i= 0, i< n, i++) C[i] = A[i] + B[i]; } int main() { // Memory allocation for A_h, B_h, and C_h // I/O to read A_hand B_h, N elements … vecAdd(A_h, B_h, C_h, N); }要在 GPU 中完成這一操作，首先我們想一下它是否適合使用 GPU，我們當時總結了四個特點：
所以，向量相家的任務適合在 GPU 上編程。
再回顧下 GPU 運算步驟，如圖4所示：
一個 CUDA 程序，我們可以把它分成3個部分：
第1部分是： 從主機 (host) 端申請 device memory，把要拷貝的內容從 host memory 拷貝到申請的 device memory 里面。
第2部分是： 設備端的核函數對拷貝進來的東西進行計算，來得到和實現運算的結果，圖4中的 Kernel 就是指在 GPU 上運行的函數。
第3部分是： 把結果從 device memory 拷貝到申請的 host memory 里面，并且釋放設備端的顯存和內存。
如下：
#include <cuda.h> void vecAdd(float* A, float* B, float* C, int n) { int size = n* sizeof(float); float* A_d, B_d, C_d; … 1. // Allocate device memory for A, B, and C // copy A and B to device memory 2. // Kernel launch code –to have the device // to perform the actual vector addition 3. // copy C from the device memory // Free device vectors }下面我們把這些內容細化到函數。
設備端代碼:
主機端代碼：
內存是插在主板上的內存插槽上的內存條，而顯存是獨立顯卡上焊在顯卡上的內存芯片。
申請顯存的函數 cudaMalloc()：
在主機端完成顯存的申請，得到相應的指針。
釋放顯存的函數 cudaFree( )：
將指向顯存的指針釋放掉。
內存和顯存之間互相拷貝的函數 cudaMemcpy( )：
參數含義是：終點的指針，起點的指針，拷貝的大小，模式 (主機端到設備端，設備端到主機端，設備端之間的拷貝)
以上三個函數是 CUDA 幫我們寫好的，如果調用的話需要先：
# include cuda.h下面就是具體的 C++ 代碼實現：
申請內存的大小是 n *sizeof(float)，定義3個指針 A_d，B_d，C_d。
cudaMalloc 函數需要傳入 1. 指針的指針 (指向申請得到的顯存的指針)。2. 申請顯存的大小。 所以分別傳入 &A_d 和 size。同理后面依次傳入 &B_d 和 size，&C_d 和 size。
cudaMemcpy 函數需要傳入 1. 終點的指針。2. 起點的指針。3. 拷貝的大小。4. 模式。 所以分別傳入 A_d, A, size, cudaMemcpyHostToDevice。同理后面依次傳入 B_d, B, size, cudaMemcpyHostToDevice 和 C, C_d, size, cudaMemcpyHostToDevice。
最后把設備端申請的顯存都釋放掉。cudaFree 函數需要傳入設備端申請顯存的指針，即 A_d，B_d，C_d。
下面我們進入最重要的部分，即：如何自己書寫一個 kernel 函數。
核函數調用的注意事項
CUDA 編程的標識符號
不同的表示符號對應著不同的工作地點和被調用地點。核函數使用 __global__ 標識，必須返回 void。__device__ & __host__ 可以一起用。
下面，按照我們剛才的對核函數的介紹，我們展示了向量相加的代碼。
代碼講解：
首先，看到 __global__ 標識，返回的是 void，就意味著 vecAddKernel 函數是一個在 host 端調用，在 device 端執行的核函數。它的三個參數就是我們之前申請好的指向三段顯存的指針。
通過 int i= threadIdx.x+ blockDim.x* blockIdx.x; (線程的索引，線程塊的索引，線程塊維度的大小) 來計算好要訪問的線程的索引的位置。
那么如何在主機端調用呢？我們使用尖括號**<<<網格 grid 維度，線程塊 block 維度>>>**來包括：線程塊數 ceil(n/256) 和一個線程塊的線程數256。
第1步主機端 __host__ 修飾：申請顯存，內存。顯存，內存的互相拷貝。內存，顯存釋放。比如圖19中申請的網格是 ceil(n/256) 維的代表一個網格有 ceil(n/256) 個線程塊；線程塊是256維的，代表一個線程塊有256個線程。
第2步設備端 __global__ 修飾：計算索引絕對位置，并行計算。
詳細地講，核函數只能在主機端調用，調用時必須申明執行參數。調用形式如下：
Kernel<<<Dg,Db, Ns, S>>>(param list);<<<>>> 運算符內是核函數的執行參數，告訴編譯器運行時如何啟動核函數，用于說明內核函數中的線程數量，以及線程是如何組織的。
<<<>>> 運算符對 kernel 函數完整的執行配置參數形式是 <<<Dg, Db, Ns, S>>>
最后我們簡單介紹下 CUDA 編程如何執行編譯的過程。因為我們之前在 CPU 上編程，使用 g++ 或 gcc 進行編譯，再通過 link 生成可執行程序。那么在 GPU 端，編譯器就是 NVCC (NVIDIA Cuda compiler driver)。
通常我們會把和 GPU 相關的頭文件放在 .h 文件里，把設備端執行的程序 (__global__ 定義的函數) 放在 .cu 文件里，這些程序我們用 NVCC 來進行編譯。主機端的程序放在 .h 和 .cpp 里面，這些程序我們可以繼續用 g++ 或 gcc 來進行編譯。
通常我們有這幾種編譯的方法：
CUDA 中 threadIdx，blockIdx，blockDim，gridDim 的使用
下面這張圖21比較清晰的表示的幾個概念的關系：
cuda 通過<<< >>>符號來分配索引線程的方式，我知道的一共有15種索引方式。
4 實踐
4.1 向量相加 CUDA 代碼
這一節我們通過一個實例直觀感受下 CUDA 并經計算究竟能使這些計算簡單，并行度高的操作加速多少。
我們先看一下 CPU 執行向量相加的代碼：
#include <iostream> #include <cstdlib> #include <sys/time.h>using namespace std;void vecAdd(float* A, float* B, float* C, int n) {for (int i = 0; i < n; i++) {C[i] = A[i] + B[i];} }int main(int argc, char *argv[]) {int n = atoi(argv[1]);cout << n << endl;size_t size = n * sizeof(float);// host memeryfloat *a = (float *)malloc(size);float *b = (float *)malloc(size);float *c = (float *)malloc(size);for (int i = 0; i < n; i++) {float af = rand() / double(RAND_MAX);float bf = rand() / double(RAND_MAX);a[i] = af;b[i] = bf;}struct timeval t1, t2;gettimeofday(&t1, NULL);vecAdd(a, b, c, n);gettimeofday(&t2, NULL);//for (int i = 0; i < 10; i++) // cout << vecA[i] << " " << vecB[i] << " " << vecC[i] << endl;double timeuse = (t2.tv_sec - t1.tv_sec) + (double)(t2.tv_usec - t1.tv_usec)/1000000.0;cout << timeuse << endl;free(a);free(b);free(c);return 0; }注釋：
float*a =(float*)malloc(size); 分配一段內存，使用指針 a 指向它。
for 循環產生一些隨機數，并放在分配的內存里面。
vecAdd(float* A,float* B,float* C,int n) 要輸入指向3段內存的指針名，也就是 a, b, c。
gettimeofday 函數來得到精確時間。它的精度可以達到微妙，是C標準庫的函數。
最后的 free 函數把申請的3段內存釋放掉。
編譯：
g++ -O3 main_cpu.cpp -o VectorSumCPU我們再看一下 CUDA 執行向量相加的代碼：
#include <iostream> #include <cstdlib> #include <sys/time.h> #include <cuda_runtime.h>using namespace std;__global__ void vecAddKernel(float* A_d, float* B_d, float* C_d, int n) {int i = threadIdx.x + blockDim.x * blockIdx.x;if (i < n) C_d[i] = A_d[i] + B_d[i]; }int main(int argc, char *argv[]) {int n = atoi(argv[1]);cout << n << endl;size_t size = n * sizeof(float);// host memeryfloat *a = (float *)malloc(size);float *b = (float *)malloc(size);float *c = (float *)malloc(size);for (int i = 0; i < n; i++) {float af = rand() / double(RAND_MAX);float bf = rand() / double(RAND_MAX);a[i] = af;b[i] = bf;}float *da = NULL;float *db = NULL;float *dc = NULL;cudaMalloc((void **)&da, size);cudaMalloc((void **)&db, size);cudaMalloc((void **)&dc, size);cudaMemcpy(da,a,size,cudaMemcpyHostToDevice);cudaMemcpy(db,b,size,cudaMemcpyHostToDevice);cudaMemcpy(dc,c,size,cudaMemcpyHostToDevice);struct timeval t1, t2;int threadPerBlock = 256;int blockPerGrid = (n + threadPerBlock - 1)/threadPerBlock;printf("threadPerBlock: %d \nblockPerGrid: %d \n",threadPerBlock,blockPerGrid);gettimeofday(&t1, NULL);vecAddKernel <<< blockPerGrid, threadPerBlock >>> (da, db, dc, n);gettimeofday(&t2, NULL);cudaMemcpy(c,dc,size,cudaMemcpyDeviceToHost);//for (int i = 0; i < 10; i++) // cout << vecA[i] << " " << vecB[i] << " " << vecC[i] << endl;double timeuse = (t2.tv_sec - t1.tv_sec) + (double)(t2.tv_usec - t1.tv_usec)/1000000.0;cout << timeuse << endl;cudaFree(da);cudaFree(db);cudaFree(dc);free(a);free(b);free(c);return 0; }注釋：
首先要用 __global__ 來修飾。
vecAdd(float* A,float* B,float* C,int n) 要輸入指向3段顯存的指針名，也就是 d_a, d_b, d_c。
float*da =NULL; 定義空指針。
cudaMalloc((void**)&da, size); 申請顯存，da 指向申請的顯存，注意 cudaMalloc 函數傳入指針的指針 (指向申請得到的顯存的指針)。
cudaMemcpy(da,a,size,cudaMemcpyHostToDevice) 把內存的東西拷貝到顯存，也就是把 a, b, c 里面的東西拷貝到 d_a, d_b, d_c 中。
int threadPerBlock =256; int blockPerGrid =(n + threadPerBlock -1)/threadPerBlock; 計算線程塊和網格的數量。
vecAddKernel <<< blockPerGrid, threadPerBlock >>> (da, db, dc, n); 調用核函數。
gettimeofday 函數來得到精確時間。它的精度可以達到微妙，是C標準庫的函數。
最后的 free 函數把申請的3段內存釋放掉。
編譯：
/usr/local/cuda/bin/nvcc main_gpu.cu -o VectorSumGPU4.2 實踐向量相加
編譯之后得到可執行文件 VectorSumCPU 和 VectorSumGPU 之后，我們可以執行一下比較下運行時間 (注意要在 linux 下運行)：
在 CPU 下，執行1000000000次加需要4.18秒。
./VectorSumCPU 1000000000 1000000000 4.18261在 GPU 下，執行1000000000次加只需要1.6e-05秒，哇。
(base) wjh19@iccv:~/mage/CUDA/db$ ./VectorSumGPU 1000000000 1000000000 threadPerBlock: 256 blockPerGrid: 3906250 1.6e-05GPU 對于計算簡單，并行度高的計算果然可以大幅提速！！！
在 CPU 下，執行1000次加需要1e-06秒。
(base) wjh19@iccv:~/mage/CUDA/db$ ./VectorSumCPU 1000 1000 1e-06在 GPU 下，執行1000次加需要1.3e-05秒。
(base) wjh19@iccv:~/mage/CUDA/db$ ./VectorSumGPU 1000 1000 threadPerBlock: 256 blockPerGrid: 4 1.3e-05GPU 對于少量計算效率反倒不如 CPU。
參考
2. D. Kirk and W. Hwu, “Programming Massively Parallel Processors –A Hands-on Approach, Second Edition”
3. CUDA by example, Sanders and Kandrot
4. Nvidia CUDA C Programming Guide：
5. CS/EE217 GPU Architecture andProgramming
總結
以上是生活随笔為你收集整理的CUDA 编程上手指南：CUDA C 编程及 GPU 基本知识的全部內容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： 专访 Swin Transformer
- 下一篇： 深入思考：算法工程师的落地能力具体指什么