CUDA：Compute Unified Device Architecture盆昙，統(tǒng)一計(jì)算設(shè)備架構(gòu)五芝，CUDA?是一種由NVIDIA推出的通用并行計(jì)算架構(gòu)写妥，該架構(gòu)使GPU (Graphics processing unit) 能夠解決復(fù)雜的計(jì)算問題级零。它包含了CUDA指令集架構(gòu)（ISA）以及GPU內(nèi)部的并行計(jì)算引擎舀瓢。
CUDA 即是 NVIDIA 的 GPGPU 模型，它使用C語言為基礎(chǔ)
和CPU的比較
優(yōu)點(diǎn)：
- 更大的內(nèi)存帶寬
- 更多的執(zhí)行單元蚓聘，雖然頻率比CPU低
- 價(jià)格更低
缺點(diǎn)：
- 運(yùn)算單元多腌乡，只適合高度并行化的工作
- 對(duì)于具有高度分支的程序，效率會(huì)比較差
具體來說夜牡，GPU特別適合大量并行的數(shù)據(jù)運(yùn)算（高運(yùn)算密度）与纽。由于對(duì)每個(gè)數(shù)據(jù)進(jìn)行相同的操作，所以對(duì)復(fù)雜的流控制需求較低塘装，并且因?yàn)樘幚碓S多數(shù)據(jù)單元并且具有高運(yùn)算密度急迂，內(nèi)存讀取延時(shí)可以通過運(yùn)算來隱藏，CPU采用的是高速緩存cache來縮減延時(shí) （latency ）
可伸縮Scaleable的編程模式
核心是三個(gè)關(guān)鍵的抽象：
- 線程組的層次結(jié)構(gòu)
- 共享內(nèi)存
- 障礙同步
這種可擴(kuò)展的編程模型允許市場(chǎng)上的各種GPU架構(gòu)蹦肴，GeForce僚碎，Quadro，Tesla等阴幌。
A GPU is built around a scalable array of multithreaded Streaming Multiprocessors (SMs).
Streaming Multiprocessors (SMs)
在一個(gè)multiprocess中勺阐，一個(gè)線程塊的所有線程可以同時(shí)執(zhí)行卷中，多個(gè)線程塊block也可以同時(shí)執(zhí)行，使用SIMT (Single-Instruction, Multiple-Thread) 構(gòu)架皆看，一個(gè)線程內(nèi)部通過指令流水進(jìn)行指令級(jí)別的并行仓坞，通過硬件多線程進(jìn)行線程級(jí)別的并行。
一個(gè)GPU中有多個(gè)SM,每個(gè)SM有多個(gè)core（processor）腰吟，但是只有一個(gè)指令單元无埃，同時(shí)只能夠執(zhí)行完全相同的指令集。
Warp
Multiprocessor執(zhí)行線程塊時(shí)毛雇，把他們以32個(gè)并行線程為一組嫉称，稱為warps,每個(gè)warp由warp scheduler調(diào)度執(zhí)行，warp中的每個(gè)線程在相同的程序地址處開始執(zhí)行灵疮，但他們有自己的指令地址計(jì)數(shù)器和寄存器狀態(tài)织阅。
當(dāng)他們執(zhí)行相同的指令時(shí)可以達(dá)到最大的效率，但當(dāng)因?yàn)橛捎谝揽繑?shù)據(jù)決定的條件分支產(chǎn)生了分歧（warp divergence）震捣，warp連續(xù)執(zhí)行每個(gè)分支路徑荔棉，禁止不相關(guān)的線程，執(zhí)行完成后蒿赢，又回歸到相同的執(zhí)行路徑润樱。不同warps執(zhí)行互相獨(dú)立。
多線程的CUDA程序具有自動(dòng)適應(yīng)性羡棵，它被分解為相互獨(dú)立的線程塊壹若，可以以任意的順序執(zhí)行，不管是并行還是串行皂冰，所以一個(gè)編譯好的CUDA程序可以在任意數(shù)目的多處理器上運(yùn)行店展，只有運(yùn)行的系統(tǒng)需要知道實(shí)際的處理器數(shù)量
Kernels
一種擴(kuò)展的C函數(shù)，調(diào)用時(shí)通過N個(gè)不同的CUDA線程執(zhí)行N次秃流，不像通常C函數(shù)一樣只執(zhí)行一次
__global__說明符來聲明kernel函數(shù)赂蕴，執(zhí)行的次數(shù)通過<<<...>>>執(zhí)行配置語法來說明
每個(gè)線程對(duì)應(yīng)的唯一標(biāo)識(shí)ID通過內(nèi)置的threadIdx變量獲取
// 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);
...
}
Thread Hierarchy
線程被分成成塊的網(wǎng)格。Grid和Block都是三維的
例如：
dim3 dimBlock(8, 8, 8);
dim3 dimGrid(100, 100, 1);
Kernel<<<dimGrid, dimBlock>>>(…);
不過通常我們使用的一維grid和block
Kernel<<<block_count, block_size>>>(…);
每個(gè)block中的thread數(shù)目有限舶胀，因?yàn)樗麄冊(cè)谕粋€(gè)內(nèi)存核心中睡腿，共享有限的內(nèi)存資源，通常最多1024個(gè)線程峻贮。
grid中的block的數(shù)量：最大數(shù)目一般為65535，由被處理的數(shù)據(jù)的大小或者系統(tǒng)中處理器的數(shù)量決定
- 如果超出了任何一個(gè)數(shù)目应闯，GPU會(huì)出錯(cuò)或者產(chǎn)生垃圾數(shù)據(jù)
- GPU編程的部分時(shí)間就是用來處理硬件限制的
- 這個(gè)限制意味著必須要對(duì)Kernel程序分配線程不足進(jìn)行處理
index of thread 和 thread ID的關(guān)系
threadIdx是一個(gè)三維矢量纤控，所以線程可以通過一、二碉纺、三維線程索引index來識(shí)別船万，組成對(duì)應(yīng)維度的線程塊thread block刻撒。
一維索引直接對(duì)應(yīng)線程ID
二維(Dx, Dy)的塊， 線程索引為 (x, y)對(duì)應(yīng)ID為(x + y Dx);
三維(Dx, Dy, Dz)塊, 索引為(x, y, z) 對(duì)應(yīng)ID為 (x + y Dx+ z Dx Dy).
blockIdx 索引變量可以用來標(biāo)識(shí)網(wǎng)格中的塊耿导，blockDim可以獲取維度
同步 synchronization
線程通信和資源共享：
通過同步來相互協(xié)調(diào)声怔，調(diào)用固有函數(shù)__syncthreads()
只在block層次產(chǎn)生作用，類似于C/C++中的 barrier()函數(shù)
原子操作 Atomic Operation
執(zhí)行讀-修改-寫的原子操作舱呻，在全局或共享內(nèi)存空間中
串行操作
atomic<op>(float *address, float val);
op的范圍: Add, Sub, Exch, Min, Max, Inc, Dec, And, Or, Xor
atomicCAS(int *address, int compare, int val)
warp shuffle
CC >= 3.0 在一個(gè)warp內(nèi)的線程交換變量
int __shfl(int var, int srcLane, int width=warpSize);
計(jì)算能力 CC
設(shè)備的計(jì)算能力由版本號(hào)代表醋火，也叫做 SM Version。
這個(gè)版本號(hào)標(biāo)識(shí)了GPU硬件的支持特性箱吕，在應(yīng)用運(yùn)行時(shí)使用芥驳，決定在當(dāng)前GPU上可以實(shí)現(xiàn)的硬件特性和指令。它由主版本號(hào)X和小版本號(hào)Y組成茬高，表示為X.Y
主版本號(hào)表示核心構(gòu)架相同兆旬，副版本號(hào)表示核心構(gòu)架的改進(jìn)，可能增加了新的特性
| The major revision number | Core architecture |
|---|---|
| 5 | Maxwell architecture |
| 3 | Kepler architecture |
| 2 | Fermi architecture |
| 1 | Tesla architecture. |
The Tesla architecture is no longer supported starting with CUDA 7.0.
