https://mp.weixin.qq.com/s/KgK3ertk9XVTxWhynv2AgA
本系列是為了彌補教程和實際應用之間的空白，幫助大家理解 CUDA 編程並最終熟練使用 CUDA 編程。你不需要具備 OpenGL 或者 DirectX 的知識，也不需要有計算及圖形學的背景。
目錄
1 CPU 和 GPU 的基礎知識
2 CUDA 編程的重要概念
3 並行計算向量相加
4 實踐
4.1 向量相加 CUDA 代碼
4.2 實踐向量相加
5 給大家的一點參考資料
1 CPU 和 GPU 的基礎知識
提到處理器結構，有2個指標是經常要考慮的：延遲和吞吐量。所謂延遲，是指從發出指令到最終返回結果中間經歷的時間間隔。而所謂吞吐量，就是單位之間內處理的指令的條數。
下圖1是 CPU 的示意圖。從圖中可以看出 CPU 的幾個特點：
- CPU 中包含了多級高速的緩存結構。 因為我們知道處理運算的速度遠高於訪問存儲的速度，那么奔着空間換時間的思想，設計了多級高速的緩存結構，將經常訪問的內容放到低級緩存中，將不經常訪問的內容放到高級緩存中，從而提升了指令訪問存儲的速度。
- CPU 中包含了很多控制單元。 具體有2種，一個是分支預測機制，另一個是流水線前傳機制。
- CPU 的運算單元 (Core) 強大，整型浮點型復雜運算速度快。
圖1：CPU 的示意圖
所以綜合以上三點，CPU 在設計時的導向就是減少指令的時延，我們稱之為延遲導向設計，如下圖3所示。
下圖2是 GPU 的示意圖，它與之前 CPU 的示意圖相比有着非常大的不同。從圖中可以看出 GPU 的幾個特點 (注意紫色和黃色的區域分別是緩存單元和控制單元)：
- GPU 中雖有緩存結構但是數量少。 因為要減少指令訪問緩存的次數。
- GPU 中控制單元非常簡單。 控制單元中也沒有分支預測機制和數據轉發機制。對於復雜的指令運算就會比較慢。
- GPU 的運算單元 (Core) 非常多，采用長延時流水線以實現高吞吐量。 每一行的運算單元的控制器只有一個，意味着每一行的運算單元使用的指令是相同的，不同的是它們的數據內容。那么這種整齊划一的運算方式使得 GPU 對於那些控制簡單但運算高效的指令的效率顯著增加。
圖2：GPU 的示意圖
所以，GPU 在設計過程中以一個原則為核心：增加簡單指令的吞吐。因此，我們稱 GPU 為吞吐導向設計，，如下圖3所示。
圖3：CPU 是延遲導向設計，GPU 是吞吐導向設計
那么究竟在什么情況下使用 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 里面，並且釋放設備端的顯存和內存。
圖4：一個 CUDA 程序可以分成3個部分
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。
圖5：CUDA 內存模型，硬件角度
從軟件的角度來講：
- 線程處理器 (SP) 對應線程 (thread)。
- 多核處理器 (SM) 對應線程塊 (thread block)。
- 設備端 (device) 對應線程塊組合體 (grid)。
圖6：CUDA 內存模型，軟件角度
如下圖6所示，所謂線程塊內存模型在軟件側的一個最基本的執行單位，所以我們從這里開始梳理。線程塊就是線程的組合體，它具有如下這些特點：
- 塊內的線程通過共享內存、原子操作和屏障同步進行協作 (shared memory, atomic operations and barrier synchronization)
- 不同塊中的線程不能協作。
如下圖7所示的線程塊就是由256個線程組成的，它執行的任務就是一個最基本的向量相加的一個操作。在線程塊內，這256個線程的計算是彼此互相獨立的，並行的。下面的這個 [i]，就是如何確定每個線程的索引 (在顯存中的位置)。在計算完以后 (圖中彎箭頭的頭部)，會設置一個時鍾，將這256個線程的計算結果進行同步。
圖7：一個256個線程組成的線程塊
以上就是一個256位向量的加的操作的並行處理方法，得到最終的向量加的結果。
所謂網格 (grid)，其實就是線程塊的組合體，如下圖8所示。
- 網格 (grid) 內的線程塊是彼此互相獨立，互不影響的。
- 全局內存可以由所有的線程塊進行訪問。
CUDA 核函數由線程網格 (數組) 執行。每個線程都有一個索引，用於計算內存地址和做出控制決策。在計算完以后 (圖中所有彎箭頭的頭部)，會設置一個時鍾，將這N個線程塊的計算結果進行同步。
圖8：網格就是線程塊的組合體
線程塊 id & 線程 id：定位獨立線程的門牌號
核函數需要確定每個線程在顯存中的位置，我們之前提到 CUDA 的核函數是要在設備端來進行計算和處理的，在執行核函數時需要訪問到每個線程的 registers (寄存器) 和 local memory (局部內存)。在這個過程中需要確定每一個線程在顯存上的位置。所以我們需要像圖9那樣使用線程塊的 index 和線程的 index 來確定線程在顯存上的位置。
圖9：使用線程塊的 index 和線程的 index 來確定線程在顯存上的位置
如圖9所示，圖9中的線程塊索引是2維的，每個網格都由2×2個線程塊組成；線程索引是3維的，每個線程塊都由2×4×2個線程組成，所以代碼應該是：
圖10：線程Id計算
圖10中：M=N=2，P,Q,S=2,4,2。
每個線程x的那一維應該是線程塊的索引×線程塊的x維度大小+線程的索引。(設備端線程x的那一維的索引)。
每個線程y的那一維應該是線程塊的索引×線程塊的y維度大小+線程的索引。(設備端線程y的那一維的索引)。
線程束 (warp)
前面我們提到，如圖11所示的每一行由1個控制單元加上若干計算單元所組成，這些所有的計算單元執行的控制指令是一個。這其實就是個非常典型的 "單指令多數據流機制"。
圖11：一個線程束 (warp)：采用單指令多數據流機制
單指令多數據流機制是說：執行的指令是一條，只不過不同的計算單元使用的數據是不一樣的。而上面這一行，我們就稱之為一個線程束 (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所示。
圖12：Single Instruction Multiple Thread (SIMT)
3 並行計算向量相加
下面我們就用一個實際的例子來看看 CUDA 編程具體是如何操作的。例子就是兩個長度為N的張量相加，如下圖13所示。
圖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，我們當時總結了四個特點：
- 訪問內存次數少，滿足。
- 控制指令簡單，無復雜分枝預測，跳轉指令，滿足。
- 計算指令簡單，滿足，是簡單的加法操作。
- 並行度高，滿足，不同的 [i] 之間不互相影響。
所以，向量相家的任務適合在 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
}
下面我們把這些內容細化到函數。
設備端代碼:
- 讀寫線程寄存器
- 讀寫 Grid 中全局內存
- 讀寫 block 中共享內存
主機端代碼：
- 申請顯存，內存
- Grid 中全局內存拷貝轉移 (顯存，內存互相拷貝)
- 內存，顯存釋放
內存是插在主板上的內存插槽上的內存條，而顯存是獨立顯卡上焊在顯卡上的內存芯片。
申請顯存的函數 cudaMalloc()：
在主機端完成顯存的申請，得到相應的指針。
圖14：申請顯存的函數 cudaMalloc()
釋放顯存的函數 cudaFree( )：
將指向顯存的指針釋放掉。
圖15：釋放顯存的函數 cudaFree( )
內存和顯存之間互相拷貝的函數 cudaMemcpy( )：
參數含義是：終點的指針，起點的指針，拷貝的大小，模式 (主機端到設備端，設備端到主機端，設備端之間的拷貝)
圖16：內存和顯存之間互相拷貝的函數 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。
void vecAdd(float* A, float* B, float* C, int n)
{
int size = n * sizeof(float);
float* A_d, *B_d, *C_d;
1. // Transfer A and B to device memory
cudaMalloc((void **) &A_d, size);
cudaMemcpy(A_d, A, size, cudaMemcpyHostToDevice);
cudaMalloc((void **) &B_d, size);
cudaMemcpy(B_d, B, size, cudaMemcpyHostToDevice);
// Allocate device memory for
cudaMalloc((void **) &C_d, size);
2. // Kernel invocation code –to be shown later
…
3. // Transfer C from device to host
cudaMemcpy(C, C_d, size, cudaMemcpyDeviceToHost);
// Free device memory for A, B, C
cudaFree(A_d); cudaFree(B_d); cudaFree(C_d);
}
下面我們進入最重要的部分，即：如何自己書寫一個 kernel 函數。
核函數調用的注意事項
- 在 GPU 上執行的函數。
- 一般通過標識符 __global__ 修飾。
- 調用通過<<<參數1,參數2>>>，用於說明內核函數中的線程數量，以及線程是如何組織的。
- 以網格 (Grid) 的形式組織，每個線程格由若干個線程塊 (block) 組成，而每個線程塊又由若干個線程 (thread) 組成。
- 調用時必須聲明內核函數的執行參數。
- 在編程時，必須先為 kernel 函數中用到的數組或變量分配好足夠的空間，再調用 kernel 函數，否則在 GPU 計算時會發生錯誤。
CUDA 編程的標識符號
不同的表示符號對應着不同的工作地點和被調用地點。核函數使用 __global__ 標識，必須返回 void。__device__ & __host__ 可以一起用。
圖17：CUDA 編程的標識符號
下面，按照我們剛才的對核函數的介紹，我們展示了向量相加的代碼。
代碼講解：
首先，看到 __global__ 標識，返回的是 void，就意味着 vecAddKernel 函數是一個在 host 端調用，在 device 端執行的核函數。它的三個參數就是我們之前申請好的指向三段顯存的指針。
通過 int i= threadIdx.x+ blockDim.x* blockIdx.x; (線程的索引，線程塊的索引，線程塊維度的大小) 來計算好要訪問的線程的索引的位置。
那么如何在主機端調用呢？我們使用尖括號**<<<網格 grid 維度，線程塊 block 維度>>>**來包括：線程塊數 ceil(n/256) 和一個線程塊的線程數256。
圖18：向量相加的代碼
第1步主機端 __host__ 修飾：申請顯存，內存。顯存，內存的互相拷貝。內存，顯存釋放。比如圖19中申請的網格是 ceil(n/256) 維的代表一個網格有 ceil(n/256) 個線程塊；線程塊是256維的，代表一個線程塊有256個線程。
第2步設備端 __global__ 修飾：計算索引絕對位置，並行計算。
圖19：主機端和設備端代碼
詳細地講，核函數只能在主機端調用，調用時必須申明執行參數。調用形式如下：
Kernel<<<Dg,Db, Ns, S>>>(param list);
<<<>>> 運算符內是核函數的執行參數，告訴編譯器運行時如何啟動核函數，用於說明內核函數中的線程數量，以及線程是如何組織的。
<<<>>> 運算符對 kernel 函數完整的執行配置參數形式是 <<<Dg, Db, Ns, S>>>
- 參數 Dg 用於定義整個 grid 的維度和尺寸，即一個 grid 有多少個 block。為 dim3 類型。Dim3 Dg(Dg.x, Dg.y, 1) 表示grid中每行有 Dg.x 個 block，每列有 Dg.y 個 block，第三維恆為1(目前一個核函數只有一個grid)。整個 grid 中共有 Dg.x*Dg.y 個 block，其中 Dg.x 和 Dg.y 最大值為65535。
- 參數 Db 用於定義一個 block 的維度和尺寸，即一個 block 有多少個 thread。為 dim3 類型。Dim3 Db(Db.x, Db.y, Db.z) 表示整個 block 中每行有 Db.x 個 thread，每列有 Db.y 個 thread，高度為 Db.z。Db.x 和 Db.y 最大值為512，Db.z 最大值為62。一個 block 中共有 Db.x*Db.y*Db.z 個 thread。計算能力為1.0,1.1的硬件該乘積的最大值為768，計算能力為1.2,1.3的硬件支持的最大值為1024。
- 參數 Ns 是一個可選參數，用於設置每個 block 除了靜態分配的 shared Memory 以外，最多能動態分配的shared memory 大小，單位為 byte。不需要動態分配時該值為0或省略不寫。
- 參數 S 是一個 cudaStream_t 類型的可選參數，初始值為零，表示該核函數處在哪個流之中。
最后我們簡單介紹下 CUDA 編程如何執行編譯的過程。因為我們之前在 CPU 上編程，使用 g++ 或 gcc 進行編譯，再通過 link 生成可執行程序。那么在 GPU 端，編譯器就是 NVCC (NVIDIA Cuda compiler driver)。
通常我們會把和 GPU 相關的頭文件放在 .h 文件里，把設備端執行的程序 (__global__ 定義的函數) 放在 .cu 文件里，這些程序我們用 NVCC 來進行編譯。主機端的程序放在 .h 和 .cpp 里面，這些程序我們可以繼續用 g++ 或 gcc 來進行編譯。
通常我們有這幾種編譯的方法：
- 逐個文件編譯 (GPU 和 CPU 的程序都編譯成 .o 文件。最后把它們匯總在一起，並 link 為一個可執行文件 .exe)，但是這只適用於文件數較少的情況，當文件數較多時，這種辦法就顯得比較復雜。
- 使用 cmake 方式編譯，寫一個 cmake.txt，下文有介紹。
圖20：CUDA 編程如何執行編譯的過程
CUDA 中 threadIdx，blockIdx，blockDim，gridDim 的使用
- threadIdx是一個uint3類型，表示一個線程的索引。
- blockIdx是一個uint3類型，表示一個線程塊的索引，一個線程塊中通常有多個線程。
- blockDim是一個dim3類型，表示線程塊的大小。
- gridDim是一個dim3類型，表示網格的大小，一個網格中通常有多個線程塊。
下面這張圖21比較清晰的表示的幾個概念的關系：
圖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 memery
float *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 memery
float *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 VectorSumGPU
4.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-05
GPU 對於計算簡單，並行度高的計算果然可以大幅提速！！！
在 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-05
GPU 對於少量計算效率反倒不如 CPU。
參考
1.深藍學院課程講解：https://www.shenlanxueyuan.com/course/410
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：https://docs.nvidia.com/cuda/cuda-c-programming-guide/
5. CS/EE217 GPU Architecture andProgramming
如果覺得有用，就請分享到朋友圈吧！
