CUDA系列学习（一）An Introduction to GPU and CUDA
本文從軟硬件層面講一下CUDA的結(jié)構(gòu)，應(yīng)用，邏輯和接口。分為以下章節(jié)：
（一）、GPU與CPU
（二）、CUDA硬件層面
（三）、CUDA安裝
（四）、CUDA 結(jié)構(gòu)與接口
???????????? 4.1 Kernels
?????????????4.2 Thread，Block, Grid
???????????? 4.3 Memory
???????????? 4.4 Execution
（五）、碼HelloWorld——數(shù)組求和
希望感興趣的同學(xué)可以一起討論。
(一)、GPU與CPU
對于浮點(diǎn)數(shù)操作能力，CPU與GPU的能力相差在GPU更適用于計(jì)算強(qiáng)度高，多并行的計(jì)算中。因此，GPU擁有更多晶體管，而不是像CPU一樣的數(shù)據(jù)Cache和流程控制器。這樣的設(shè)計(jì)是因?yàn)槎嗖⑿杏?jì)算的時(shí)候每個數(shù)據(jù)單元執(zhí)行相同程序，不需要那么繁瑣的流程控制，而更需要高計(jì)算能力，這也不需要大cache。
（二）、CUDA硬件層面：
Nvidia于2006年引入CUDA，一個GPU內(nèi)嵌通用并行計(jì)算平臺。CUDA支持C, C++, Fortran, Java, Python等語言。
那么一個多線程CUDA程序如何執(zhí)行的呢？
GPU建立在一組多處理器（SMX，Streaming Multiprocessors）附近。
一個SMX的配置：
- 192 cores（都是SIMT cores（Single Instruction Multiple Threads） and 64k registers（如下圖所示）
???????? GPU中的SIMT對應(yīng)于CPU中的SIMD（Single Instruction Multiple Data）
- 64KB of shared memory / L1 cache
- 8KB cache for constants
- 48KB texture cache for read-only arrays
- up to 2K threads per SMX
不同顯卡有不同配置（即SMX數(shù)量不同），幾個例子：
每個multi-thread程序的execution kernel instance（kernel定義見下一節(jié),instance指block）在一個SMX上執(zhí)行，一個多線程程序會分配到blocks of threads（每個block中負(fù)責(zé)一部分線程）中獨(dú)立執(zhí)行。所以GPU中的處理器越多執(zhí)行越快（因?yàn)槿绻鸖MX不夠給每個kernel instance分配一個，就要幾個kernel搶一個SMX了）。具體來講，如果SMX上有足夠寄存器和內(nèi)存（后面會講到，shared memory），就多個kernel instance在一個SMX上執(zhí)行，否則放到隊(duì)列里等。
圖：表示不同SM數(shù)帶來的執(zhí)行速度差異。
GPU工作原理：首先通過主接口讀取中央處理器指令，GigaThread引擎從系統(tǒng)內(nèi)存中獲取特定的數(shù)據(jù)并拷貝到顯存中，為顯存控制器提供數(shù)據(jù)存取所需的高帶寬。GigaThread引擎隨后為各個SMX創(chuàng)建和分派線程塊（warp, 詳細(xì)介紹見SIMT架構(gòu)或者CUDA系列學(xué)習(xí)（二）），SMX則將多個Warp調(diào)度到各CUDA核心以及其他執(zhí)行單元。在圖形流水線出現(xiàn)工作超載的時(shí)候，GigaThread引擎還負(fù)責(zé)進(jìn)行工作的重新分配。
（三）、CUDA安裝
裝CUDA主要裝以下3個組建：
1. driver
- low-level software that controls the graphics card
2. toolkit
- nvcc CUDA compiler
- profiling and debugging tools
- several libraries
3. SDK
- lots of demonstration examples
- some error-checking utilities
- not officially supported by NVIDIA
- almost no documentation
詳情請見CUDA 安裝與配置
（四）、CUDA 結(jié)構(gòu)與接口
4.1 Kernels
CUDA C 中可通過定義kernel，每次被調(diào)用就在N個CUDA thread中并行執(zhí)行。
kernel的定義：
- 聲明 __global__
- 配置kernel_routine<<<gridDim, Blockdim>>>(args)
?????????? 其中g(shù)ridDim和Blockdim變量可以是int或dim3（<=3維）類型的變量。gridDim表示每個grid中block結(jié)構(gòu)(the number of instances(blocks) of the kernel)，Blockdim表示每個block中thread結(jié)構(gòu)。那么。。thread，block，grid又是啥？往下看。。。見4.2節(jié)
- 每個執(zhí)行該kernel的thread都會通過被分配到一個unique thread ID，就是built-in變量：threadIdx
4.2 Thread，Block，Grid
很多threads組成1維，2維or3維的thread block. 為了標(biāo)記thread在block中的位置（index），我們可以用上面講的threadIdx。threadIdx是一個維度<=3的vector。還可以用thread index（一個標(biāo)量）表示這個位置。
thread的index與threadIdx的關(guān)系：
| | Thread index |
| 1 | T |
| 2 | T.x + T.y * Dx |
| 3 | T.x+T.y*Dx+z*Dx*Dy |
其中T表示變量threadIdx。（Dx, Dy, Dz）為block的size（每一維有多少threads）。
因?yàn)?span style="color:rgb(255,0,0)">一個block內(nèi)的所有threads會在同一處理器內(nèi)核上共享內(nèi)存資源，所以block內(nèi)有多少threads是有限制的。目前GPU限制每個 block最多有1024個threads。但是一個kernel可以在多個相同shape的block上執(zhí)行，效果等效于在一個有N*#thread per block個thread的block上執(zhí)行。
Block又被組織成grid。同樣，grid中block也可以被組織成1維，2維or3維。一個grid中的block數(shù)量由系統(tǒng)中處理器個數(shù)或待處理的數(shù)據(jù)量決定。
和threadIdx類似，對于block有built-in變量blockDim（block dimension）和blockIdx（block index）。
回過頭來看4.1中的configureation，舉個栗子，假設(shè)A，B，C都是大小[N][N]的二維矩陣，kernel MatAdd目的將A，B對應(yīng)位置元素加和給C的對應(yīng)位置。
聲明：
[cpp]?view plain?copy ?
這里threadsPerBlock(16,16)一般是標(biāo)配。例子中，我們假定grid中block足夠多，保證N/threadsPerBlock不會超限。
4.3 Memory
前面提到了Block中的threads共享內(nèi)存，那么怎樣同步呢？在kernel中調(diào)用內(nèi)部__synthreads()函數(shù)，其作用是block內(nèi)的所有threads必須全部執(zhí)行完，程序才能繼續(xù)往下走。那么thread到底怎樣使用memory呢？
- 每個thread有private local memory
- 每個block有shared memory
- 所有thread都能訪問到相同的一塊global memory
- 所有thread都能訪問兩塊read-only memory：constant & texture array(通常放查找表)
其中，global，constant，texture memory伴隨kernel生死。
CUDA程序執(zhí)行的時(shí)候，GPU就像一個獨(dú)立的設(shè)備一樣，kernel部分由GPU執(zhí)行，其余部分CPU執(zhí)行。于是memory就被分為host memory（for CPU）& device memory（for GPU）。因此，一個程序需要在CUDA運(yùn)行時(shí)管理device memory的分配，釋放和device & host memory之間的data transfer。
4.4 Execution
從執(zhí)行角度看，程序經(jīng)過了以下步驟：
1. initialises card
2. allocates memory in host and on device
3. copies data from host to device memory
4. launches multiple instances of execution “kernel” on device
5. copies data from device memory to host
6. repeats 3-5 as needed
7. de-allocates all memory and terminates
總結(jié)：每個kernel放在一個grid上執(zhí)行，1個kernel有多個instance，每個instance在一個block上執(zhí)行，每個block只能在一個SM上執(zhí)行，如果block數(shù)>SM數(shù)，多個block搶SM用。kernel的一個instance在SMX上通過一組進(jìn)程來執(zhí)行。如下圖所示：
總結(jié)：
CUDA的3個key abstraction：thread groups, shared memories, 和barrier synchronization
CUDA中的built-in變量：gridDim, blockDim, blockIdx（block在grid中的index）, threadIdx, warpSize（threads的warp size）
（五）、碼HelloWorld
- kernel code很像MPI，從單線程的角度coding
- 需要think about每個變量放在哪塊內(nèi)存
這里我們以數(shù)組對應(yīng)元素相加為例，看下Code ：
[cpp]?view plain?copy ?
編譯：
????????? nvcc -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc/ Matadd.cu?
運(yùn)行結(jié)果：
????????? 3
OK，大功告成。
這里注意kernel部分的code，所有變量都必須是device variable，即需要通過cudaMalloc分配過memory的。之前我忘記將A，B數(shù)組cudaMemcpy到dA,dB，而直接傳入MatAdd kernel就出現(xiàn)了運(yùn)行一次過后卡住的問題。
參考：
1.?CUDA C Programming Guide
2.?An Introduction to CUDA
3.?CUDA 安裝與配置
4.?CUDA調(diào)試工具——CUDA GDB
5.?GPU工作方式
6.?Fermi 架構(gòu)白皮書（GPU繼承了Fermi的很多架構(gòu)特點(diǎn)）
7.?GTX460架構(gòu)
from:?http://blog.csdn.net/abcjennifer/article/details/42436727
總結(jié)
以上是生活随笔為你收集整理的CUDA系列学习（一）An Introduction to GPU and CUDA的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： GMM-HMM语音识别模型 原理篇
- 下一篇： 循环神经网络教程Recurrent Ne