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í)候每個(gè)數(shù)據(jù)單元執(zhí)行相同程序，不需要那么繁瑣的流程控制，而更需要高計(jì)算能力，這也不需要大cache。
（二）、CUDA硬件層面：
Nvidia于2006年引入CUDA，一個(gè)GPU內(nèi)嵌通用并行計(jì)算平臺。CUDA支持C, C++, Fortran, Java, Python等語言。
那么一個(gè)多線程CUDA程序如何執(zhí)行的呢？
GPU建立在一組多處理器（SMX，Streaming Multiprocessors）附近。
一個(gè)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ù)量不同），幾個(gè)例子：
每個(gè)multi-thread程序的execution kernel instance（kernel定義見下一節(jié),instance指block）在一個(gè)SMX上執(zhí)行，一個(gè)多線程程序會(huì)分配到blocks of threads（每個(gè)block中負(fù)責(zé)一部分線程）中獨(dú)立執(zhí)行。所以GPU中的處理器越多執(zhí)行越快（因?yàn)槿绻鸖MX不夠給每個(gè)kernel instance分配一個(gè)，就要幾個(gè)kernel搶一個(gè)SMX了）。具體來講，如果SMX上有足夠寄存器和內(nèi)存（后面會(huì)講到，shared memory），就多個(gè)kernel instance在一個(gè)SMX上執(zhí)行，否則放到隊(duì)列里等。
圖：表示不同SM數(shù)帶來的執(zhí)行速度差異。
GPU工作原理：首先通過主接口讀取中央處理器指令，GigaThread引擎從系統(tǒng)內(nèi)存中獲取特定的數(shù)據(jù)并拷貝到顯存中，為顯存控制器提供數(shù)據(jù)存取所需的高帶寬。GigaThread引擎隨后為各個(gè)SMX創(chuàng)建和分派線程塊（warp, 詳細(xì)介紹見SIMT架構(gòu)或者CUDA系列學(xué)習(xí)（二）），SMX則將多個(gè)Warp調(diào)度到各CUDA核心以及其他執(zhí)行單元。在圖形流水線出現(xiàn)工作超載的時(shí)候，GigaThread引擎還負(fù)責(zé)進(jìn)行工作的重新分配。
（三）、CUDA安裝
裝CUDA主要裝以下3個(gè)組建：
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個(gè)CUDA thread中并行執(zhí)行。
kernel的定義：
- 聲明 __global__
- 配置kernel_routine<<<gridDim, Blockdim>>>(args)
?????????? 其中g(shù)ridDim和Blockdim變量可以是int或dim3（<=3維）類型的變量。gridDim表示每個(gè)grid中block結(jié)構(gòu)(the number of instances(blocks) of the kernel)，Blockdim表示每個(gè)block中thread結(jié)構(gòu)。那么。。thread，block，grid又是啥？往下看。。。見4.2節(jié)
- 每個(gè)執(zhí)行該kernel的thread都會(huì)通過被分配到一個(gè)unique thread ID，就是built-in變量：threadIdx
4.2 Thread，Block，Grid
很多threads組成1維，2維or3維的thread block. 為了標(biāo)記thread在block中的位置（index），我們可以用上面講的threadIdx。threadIdx是一個(gè)維度<=3的vector。還可以用thread index（一個(gè)標(biāo)量）表示這個(gè)位置。
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)">一個(gè)block內(nèi)的所有threads會(huì)在同一處理器內(nèi)核上共享內(nèi)存資源，所以block內(nèi)有多少threads是有限制的。目前GPU限制每個(gè) block最多有1024個(gè)threads。但是一個(gè)kernel可以在多個(gè)相同shape的block上執(zhí)行，效果等效于在一個(gè)有N*#thread per block個(gè)thread的block上執(zhí)行。
Block又被組織成grid。同樣，grid中block也可以被組織成1維，2維or3維。一個(gè)grid中的block數(shù)量由系統(tǒng)中處理器個(gè)數(shù)或待處理的數(shù)據(jù)量決定。
和threadIdx類似，對于block有built-in變量blockDim（block dimension）和blockIdx（block index）。
回過頭來看4.1中的configureation，舉個(gè)栗子，假設(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不會(huì)超限。
4.3 Memory
前面提到了Block中的threads共享內(nèi)存，那么怎樣同步呢？在kernel中調(diào)用內(nèi)部__synthreads()函數(shù)，其作用是block內(nèi)的所有threads必須全部執(zhí)行完，程序才能繼續(xù)往下走。那么thread到底怎樣使用memory呢？
- 每個(gè)thread有private local memory
- 每個(gè)block有shared memory
- 所有thread都能訪問到相同的一塊global memory
- 所有thread都能訪問兩塊read-only memory：constant & texture array(通常放查找表)
其中，global，constant，texture memory伴隨kernel生死。
CUDA程序執(zhí)行的時(shí)候，GPU就像一個(gè)獨(dú)立的設(shè)備一樣，kernel部分由GPU執(zhí)行，其余部分CPU執(zhí)行。于是memory就被分為host memory（for CPU）& device memory（for GPU）。因此，一個(gè)程序需要在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é)：每個(gè)kernel放在一個(gè)grid上執(zhí)行，1個(gè)kernel有多個(gè)instance，每個(gè)instance在一個(gè)block上執(zhí)行，每個(gè)block只能在一個(gè)SM上執(zhí)行，如果block數(shù)>SM數(shù)，多個(gè)block搶SM用。kernel的一個(gè)instance在SMX上通過一組進(jìn)程來執(zhí)行。如下圖所示：
總結(jié)：
CUDA的3個(gè)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每個(gè)變量放在哪塊內(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语音识别模型 原理篇
- 下一篇： CUDA系列学习（三）GPU设计与结构Q