CUDA系列学习（一）An Introduction to GPU and CUDA
本文從軟硬件層面講一下CUDA的結構，應用，邏輯和接口。分為以下章節：
（一）、GPU與CPU
（二）、CUDA硬件層面
（三）、CUDA安裝
（四）、CUDA 結構與接口
???????????? 4.1 Kernels
?????????????4.2 Thread，Block, Grid
???????????? 4.3 Memory
???????????? 4.4 Execution
（五）、碼HelloWorld——數組求和
希望感興趣的同學可以一起討論。
(一)、GPU與CPU
對于浮點數操作能力，CPU與GPU的能力相差在GPU更適用于計算強度高，多并行的計算中。因此，GPU擁有更多晶體管，而不是像CPU一樣的數據Cache和流程控制器。這樣的設計是因為多并行計算的時候每個數據單元執行相同程序，不需要那么繁瑣的流程控制，而更需要高計算能力，這也不需要大cache。
（二）、CUDA硬件層面：
Nvidia于2006年引入CUDA，一個GPU內嵌通用并行計算平臺。CUDA支持C, C++, Fortran, Java, Python等語言。
那么一個多線程CUDA程序如何執行的呢？
GPU建立在一組多處理器（SMX，Streaming Multiprocessors）附近。
一個SMX的配置：
- 192 cores（都是SIMT cores（Single Instruction Multiple Threads） and 64k registers（如下圖所示）
???????? GPU中的SIMT對應于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數量不同），幾個例子：
每個multi-thread程序的execution kernel instance（kernel定義見下一節,instance指block）在一個SMX上執行，一個多線程程序會分配到blocks of threads（每個block中負責一部分線程）中獨立執行。所以GPU中的處理器越多執行越快（因為如果SMX不夠給每個kernel instance分配一個，就要幾個kernel搶一個SMX了）。具體來講，如果SMX上有足夠寄存器和內存（后面會講到，shared memory），就多個kernel instance在一個SMX上執行，否則放到隊列里等。
圖：表示不同SM數帶來的執行速度差異。
GPU工作原理：首先通過主接口讀取中央處理器指令，GigaThread引擎從系統內存中獲取特定的數據并拷貝到顯存中，為顯存控制器提供數據存取所需的高帶寬。GigaThread引擎隨后為各個SMX創建和分派線程塊（warp, 詳細介紹見SIMT架構或者CUDA系列學習（二）），SMX則將多個Warp調度到各CUDA核心以及其他執行單元。在圖形流水線出現工作超載的時候，GigaThread引擎還負責進行工作的重新分配。
（三）、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 結構與接口
4.1 Kernels
CUDA C 中可通過定義kernel，每次被調用就在N個CUDA thread中并行執行。
kernel的定義：
- 聲明 __global__
- 配置kernel_routine<<<gridDim, Blockdim>>>(args)
?????????? 其中gridDim和Blockdim變量可以是int或dim3（<=3維）類型的變量。gridDim表示每個grid中block結構(the number of instances(blocks) of the kernel)，Blockdim表示每個block中thread結構。那么。。thread，block，grid又是啥？往下看。。。見4.2節
- 每個執行該kernel的thread都會通過被分配到一個unique thread ID，就是built-in變量：threadIdx
4.2 Thread，Block，Grid
很多threads組成1維，2維or3維的thread block. 為了標記thread在block中的位置（index），我們可以用上面講的threadIdx。threadIdx是一個維度<=3的vector。還可以用thread index（一個標量）表示這個位置。
thread的index與threadIdx的關系：
| | 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）。
因為一個block內的所有threads會在同一處理器內核上共享內存資源，所以block內有多少threads是有限制的。目前GPU限制每個 block最多有1024個threads。但是一個kernel可以在多個相同shape的block上執行，效果等效于在一個有N*#thread per block個thread的block上執行。
Block又被組織成grid。同樣，grid中block也可以被組織成1維，2維or3維。一個grid中的block數量由系統中處理器個數或待處理的數據量決定。
和threadIdx類似，對于block有built-in變量blockDim（block dimension）和blockIdx（block index）。
回過頭來看4.1中的configureation，舉個栗子，假設A，B，C都是大小[N][N]的二維矩陣，kernel MatAdd目的將A，B對應位置元素加和給C的對應位置。
聲明：
[cpp]?view plain?copy ?
這里threadsPerBlock(16,16)一般是標配。例子中，我們假定grid中block足夠多，保證N/threadsPerBlock不會超限。
4.3 Memory
前面提到了Block中的threads共享內存，那么怎樣同步呢？在kernel中調用內部__synthreads()函數，其作用是block內的所有threads必須全部執行完，程序才能繼續往下走。那么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程序執行的時候，GPU就像一個獨立的設備一樣，kernel部分由GPU執行，其余部分CPU執行。于是memory就被分為host memory（for CPU）& device memory（for GPU）。因此，一個程序需要在CUDA運行時管理device memory的分配，釋放和device & host memory之間的data transfer。
4.4 Execution
從執行角度看，程序經過了以下步驟：
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
總結：每個kernel放在一個grid上執行，1個kernel有多個instance，每個instance在一個block上執行，每個block只能在一個SM上執行，如果block數>SM數，多個block搶SM用。kernel的一個instance在SMX上通過一組進程來執行。如下圖所示：
總結：
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每個變量放在哪塊內存
這里我們以數組對應元素相加為例，看下Code ：
[cpp]?view plain?copy ?
編譯：
????????? nvcc -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc/ Matadd.cu?
運行結果：
????????? 3
OK，大功告成。
這里注意kernel部分的code，所有變量都必須是device variable，即需要通過cudaMalloc分配過memory的。之前我忘記將A，B數組cudaMemcpy到dA,dB，而直接傳入MatAdd kernel就出現了運行一次過后卡住的問題。
參考：
1.?CUDA C Programming Guide
2.?An Introduction to CUDA
3.?CUDA 安裝與配置
4.?CUDA調試工具——CUDA GDB
5.?GPU工作方式
6.?Fermi 架構白皮書（GPU繼承了Fermi的很多架構特點）
7.?GTX460架構
from:?http://blog.csdn.net/abcjennifer/article/details/42436727
總結
以上是生活随笔為你收集整理的CUDA系列学习（一）An Introduction to GPU and CUDA的全部內容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： GMM-HMM语音识别模型 原理篇
- 下一篇： CUDA系列学习（三）GPU设计与结构Q