1. 知識準備
1.1 中央處理器（CPU）
中央處理器（CPU，Central Processing Unit）是一塊超大規(guī)模的集成電路，是一臺計算機的運算核心（Core）和控制核心（ Control Unit）。它的功能主要是解釋計算機指令以及處理計算機軟件中的數(shù)據(jù)。
中央處理器主要包括運算器（算術(shù)邏輯運算單元，ALU，Arithmetic Logic Unit）和高速緩沖存儲器（Cache）及實現(xiàn)它們之間聯(lián)系的數(shù)據(jù)（Data）、控制及狀態(tài)的總線（Bus）。它與內(nèi)部存儲器（Memory）和輸入/輸出（I/O）設(shè)備合稱為電子計算機三大核心部件。
CPU的結(jié)構(gòu)主要包括運算器（ALU, Arithmetic and Logic Unit）、控制單元（CU, Control Unit）、寄存器（Register）、高速緩存器（Cache）和它們之間通訊的數(shù)據(jù)、控制及狀態(tài)的總線。
簡單來說就是：計算單元、控制單元和存儲單元，架構(gòu)如下圖所示：
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">CPU微架構(gòu)示意圖</figcaption>
什么？架構(gòu)記不??？來，我們換種表示方法：
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">CPU微架構(gòu)示意圖（改）</figcaption>
嗯，大概就是這個意思。
從字面上我們也很好理解，計算單元主要執(zhí)行算術(shù)運算、移位等操作以及地址運算和轉(zhuǎn)換；存儲單元主要用于保存運算中產(chǎn)生的數(shù)據(jù)以及指令等；控制單元則對指令譯碼，并且發(fā)出為完成每條指令所要執(zhí)行的各個操作的控制信號。
所以一條指令在CPU中執(zhí)行的過程是這樣的：讀取到指令后，通過指令總線送到控制器（黃色區(qū)域）中進行譯碼，并發(fā)出相應(yīng)的操作控制信號；然后運算器（綠色區(qū)域）按照操作指令對數(shù)據(jù)進行計算，并通過數(shù)據(jù)總線將得到的數(shù)據(jù)存入數(shù)據(jù)緩存器（大塊橙色區(qū)域）。過程如下圖所示：
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">CPU執(zhí)行指令圖</figcaption>
是不是有點兒復雜？沒關(guān)系，這張圖完全不用記住，我們只需要知道，CPU遵循的是馮諾依曼架構(gòu)，其核心就是：存儲程序，順序執(zhí)行。
講到這里，有沒有看出問題，沒錯——在這個結(jié)構(gòu)圖中，負責計算的綠色區(qū)域占的面積似乎太小了，而橙色區(qū)域的緩存Cache和黃色區(qū)域的控制單元占據(jù)了大量空間。
高中化學有句老生常談的話叫：結(jié)構(gòu)決定性質(zhì)，放在這里也非常適用。
因為CPU的架構(gòu)中需要大量的空間去放置存儲單元（橙色部分）和控制單元（黃色部分），相比之下計算單元（綠色部分）只占據(jù)了很小的一部分，所以它在大規(guī)模并行計算能力上極受限制，而更擅長于邏輯控制。
另外，因為遵循馮諾依曼架構(gòu)（存儲程序，順序執(zhí)行），CPU就像是個一板一眼的管家，人們吩咐的事情它總是一步一步來做。但是隨著人們對更大規(guī)模與更快處理速度的需求的增加，這位管家漸漸變得有些力不從心。
于是，大家就想，能不能把多個處理器放在同一塊芯片上，讓它們一起來做事，這樣效率不就提高了嗎？
沒錯，GPU便由此誕生了。
1.2 顯卡
顯卡（Video card，Graphics card）全稱顯示接口卡，又稱顯示適配器，是計算機最基本配置、最重要的配件之一。顯卡作為電腦主機里的一個重要組成部分，是電腦進行數(shù)模信號轉(zhuǎn)換的設(shè)備，承擔輸出顯示圖形的任務(wù)。顯卡接在電腦主板上，它將電腦的數(shù)字信號轉(zhuǎn)換成模擬信號讓顯示器顯示出來，同時顯卡還是有圖像處理能力，可協(xié)助CPU工作，提高整體的運行速度。對于從事專業(yè)圖形設(shè)計的人來說顯卡非常重要。 民用和軍用顯卡圖形芯片供應(yīng)商主要包括AMD(超微半導體)和Nvidia(英偉達)2家。現(xiàn)在的top500計算機，都包含顯卡計算核心。在科學計算中，顯卡被稱為顯示加速卡。
為什么GPU特別擅長處理圖像數(shù)據(jù)呢？這是因為圖像上的每一個像素點都有被處理的需要，而且每個像素點處理的過程和方式都十分相似，也就成了GPU的天然溫床。
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">GPU微架構(gòu)示意圖</figcaption>
從架構(gòu)圖我們就能很明顯的看出，GPU的構(gòu)成相對簡單，有數(shù)量眾多的計算單元和超長的流水線，特別適合處理大量的類型統(tǒng)一的數(shù)據(jù)。
再把CPU和GPU兩者放在一張圖上看下對比，就非常一目了然了。
GPU的工作大部分都計算量大，但沒什么技術(shù)含量，而且要重復很多很多次。
但GPU無法單獨工作，必須由CPU進行控制調(diào)用才能工作。CPU可單獨作用，處理復雜的邏輯運算和不同的數(shù)據(jù)類型，但當需要大量的處理類型統(tǒng)一的數(shù)據(jù)時，則可調(diào)用GPU進行并行計算。
借用知乎上某大佬的說法，就像你有個工作需要計算幾億次一百以內(nèi)加減乘除一樣，最好的辦法就是雇上幾十個小學生一起算，一人算一部分，反正這些計算也沒什么技術(shù)含量，純粹體力活而已；而CPU就像老教授，積分微分都會算，就是工資高，一個老教授資頂二十個小學生，你要是富士康你雇哪個？
注：GPU中有很多的運算器ALU和很少的緩存cache，緩存的目的不是保存后面需要訪問的數(shù)據(jù)的，這點和CPU不同，而是為線程thread提高服務(wù)的。如果有很多線程需要訪問同一個相同的數(shù)據(jù)，緩存會合并這些訪問，然后再去訪問dram。
可愛的你如果對CUDA硬件有更多的興趣，可移步NVIDIA中文官網(wǎng)進一步學習。
1.3 內(nèi)存
內(nèi)存是計算機中重要的部件之一，它是與CPU進行溝通的橋梁。計算機中所有程序的運行都是在內(nèi)存中進行的，因此內(nèi)存的性能對計算機的影響非常大。內(nèi)存(Memory)也被稱為內(nèi)存儲器，其作用是用于暫時存放CPU中的運算數(shù)據(jù)，以及與硬盤等外部存儲器交換的數(shù)據(jù)。只要計算機在運行中，CPU就會把需要運算的數(shù)據(jù)調(diào)到內(nèi)存中進行運算，當運算完成后CPU再將結(jié)果傳送出來，內(nèi)存的運行也決定了計算機的穩(wěn)定運行。 內(nèi)存是由內(nèi)存芯片、電路板、金手指等部分組成的。
1.4 顯存
顯存，也被叫做幀緩存，它的作用是用來存儲顯卡芯片處理過或者即將提取的渲染數(shù)據(jù)。如同計算機的內(nèi)存一樣，顯存是用來存儲要處理的圖形信息的部件。
1.5 顯卡、顯卡驅(qū)動、CUDA之間的關(guān)系
顯卡：（GPU）主流是NVIDIA的GPU，深度學習本身需要大量計算。GPU的并行計算能力，在過去幾年里恰當?shù)貪M足了深度學習的需求。AMD的GPU基本沒有什么支持，可以不用考慮。
驅(qū)動：沒有顯卡驅(qū)動，就不能識別GPU硬件，不能調(diào)用其計算資源。但是呢，NVIDIA在Linux上的驅(qū)動安裝特別麻煩，尤其對于新手簡直就是噩夢。得屏蔽第三方顯卡驅(qū)動。下面會給出教程。
CUDA：是NVIDIA推出的只能用于自家GPU的并行計算框架。只有安裝這個框架才能夠進行復雜的并行計算。主流的深度學習框架也都是基于CUDA進行GPU并行加速的，幾乎無一例外。還有一個叫做cudnn，是針對深度卷積神經(jīng)網(wǎng)絡(luò)的加速庫。
查看顯卡驅(qū)動信息(以實驗室服務(wù)器為例)
ssh ubuntu@192.168.1.158
輸入服務(wù)器密碼登陸
然后，進入cuda
cd /usr/local/cuda-8.0/samples/1_Utilities/deviceQuery
運行其中的可執(zhí)行文件
./deviceQuery
得到如下信息
./deviceQuery Starting... CUDA Device Query (Runtime API) version (CUDART static linking) Detected 4 CUDA Capable device(s) Device 0: "GeForce GTX 1080 Ti" CUDA Driver Version / Runtime Version 9.0 / 8.0 CUDA Capability Major/Minor version number: 6.1 Total amount of global memory: 11171 MBytes (11713708032 bytes) (28) Multiprocessors, (128) CUDA Cores/MP: 3584 CUDA Cores GPU Max Clock rate: 1620 MHz (1.62 GHz) Memory Clock rate: 5505 Mhz Memory Bus Width: 352-bit L2 Cache Size: 2883584 bytes Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 2 copy engine(s) Run time limit on kernels: No Integrated GPU sharing Host Memory: No Support host page-locked memory mapping: Yes Alignment requirement for Surfaces: Yes Device has ECC support: Disabled Device supports Unified Addressing (UVA): Yes Device PCI Domain ID / Bus ID / location ID: 0 / 2 / 0 Compute Mode: < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) > Device 1: "GeForce GTX 1080 Ti" CUDA Driver Version / Runtime Version 9.0 / 8.0 CUDA Capability Major/Minor version number: 6.1 Total amount of global memory: 11172 MBytes (11715084288 bytes) (28) Multiprocessors, (128) CUDA Cores/MP: 3584 CUDA Cores GPU Max Clock rate: 1620 MHz (1.62 GHz) Memory Clock rate: 5505 Mhz Memory Bus Width: 352-bit L2 Cache Size: 2883584 bytes Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 2 copy engine(s) Run time limit on kernels: No Integrated GPU sharing Host Memory: No Support host page-locked memory mapping: Yes Alignment requirement for Surfaces: Yes Device has ECC support: Disabled Device supports Unified Addressing (UVA): Yes Device PCI Domain ID / Bus ID / location ID: 0 / 3 / 0 Compute Mode: < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) > Device 2: "GeForce GTX 1080 Ti" CUDA Driver Version / Runtime Version 9.0 / 8.0 CUDA Capability Major/Minor version number: 6.1 Total amount of global memory: 11172 MBytes (11715084288 bytes) (28) Multiprocessors, (128) CUDA Cores/MP: 3584 CUDA Cores GPU Max Clock rate: 1620 MHz (1.62 GHz) Memory Clock rate: 5505 Mhz Memory Bus Width: 352-bit L2 Cache Size: 2883584 bytes Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 2 copy engine(s) Run time limit on kernels: No Integrated GPU sharing Host Memory: No Support host page-locked memory mapping: Yes Alignment requirement for Surfaces: Yes Device has ECC support: Disabled Device supports Unified Addressing (UVA): Yes Device PCI Domain ID / Bus ID / location ID: 0 / 130 / 0 Compute Mode: < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) > Device 3: "GeForce GTX 1080 Ti" CUDA Driver Version / Runtime Version 9.0 / 8.0 CUDA Capability Major/Minor version number: 6.1 Total amount of global memory: 11172 MBytes (11715084288 bytes) (28) Multiprocessors, (128) CUDA Cores/MP: 3584 CUDA Cores GPU Max Clock rate: 1620 MHz (1.62 GHz) Memory Clock rate: 5505 Mhz Memory Bus Width: 352-bit L2 Cache Size: 2883584 bytes Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 2 copy engine(s) Run time limit on kernels: No Integrated GPU sharing Host Memory: No Support host page-locked memory mapping: Yes Alignment requirement for Surfaces: Yes Device has ECC support: Disabled Device supports Unified Addressing (UVA): Yes Device PCI Domain ID / Bus ID / location ID: 0 / 131 / 0 Compute Mode: < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >> Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU1) : Yes> Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU2) : No> Peer access from GeForce GTX 1080 Ti (GPU0) -> GeForce GTX 1080 Ti (GPU3) : No> Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU0) : Yes> Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU2) : No> Peer access from GeForce GTX 1080 Ti (GPU1) -> GeForce GTX 1080 Ti (GPU3) : No> Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU0) : No> Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU1) : No> Peer access from GeForce GTX 1080 Ti (GPU2) -> GeForce GTX 1080 Ti (GPU3) : Yes> Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU0) : No> Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU1) : No> Peer access from GeForce GTX 1080 Ti (GPU3) -> GeForce GTX 1080 Ti (GPU2) : Yes deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 9.0, CUDA Runtime Version = 8.0, NumDevs = 4, Device0 = GeForce GTX 1080 Ti, Device1 = GeForce GTX 1080 Ti, Device2 = GeForce GTX 1080 Ti, Device3 = GeForce GTX 1080 TiResult = PASS
大家可以在自己PC或者工作機上嘗試一下。
再啰嗦兩句
GPU就是用很多簡單的計算單元去完成大量的計算任務(wù)，純粹的人海戰(zhàn)術(shù)。這種策略基于一個前提，就是小學生A和小學生B的工作沒有什么依賴性，是互相獨立的。
但有一點需要強調(diào)，雖然GPU是為了圖像處理而生的，但是我們通過前面的介紹可以發(fā)現(xiàn)，它在結(jié)構(gòu)上并沒有專門為圖像服務(wù)的部件，只是對CPU的結(jié)構(gòu)進行了優(yōu)化與調(diào)整，所以現(xiàn)在GPU不僅可以在圖像處理領(lǐng)域大顯身手，它還被用來科學計算、密碼破解、數(shù)值分析，海量數(shù)據(jù)處理（排序，Map-Reduce等），金融分析等需要大規(guī)模并行計算的領(lǐng)域。
所以GPU也可以認為是一種較通用的芯片。
2. CUDA軟件構(gòu)架
CUDA是一種新的操作GPU計算的硬件和軟件架構(gòu)，它將GPU視作一個數(shù)據(jù)并行計算設(shè)備，而且無需把這些計算映射到圖形API。操作系統(tǒng)的多任務(wù)機制可以同時管理CUDA訪問GPU和圖形程序的運行庫，其計算特性支持利用CUDA直觀地編寫GPU核心程序。目前Tesla架構(gòu)具有在筆記本電腦、臺式機、工作站和服務(wù)器上的廣泛可用性，配以C/C++語言的編程環(huán)境和CUDA軟件，使這種架構(gòu)得以成為最優(yōu)秀的超級計算平臺。
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">CUDA軟件層次結(jié)構(gòu)</figcaption>
CUDA在軟件方面組成有：一個CUDA庫、一個應(yīng)用程序編程接口（API）及其運行庫(Runtime)、兩個較高級別的通用數(shù)學庫，即CUFFT和CUBLAS。CUDA改進了DRAM的讀寫靈活性，使得GPU與CPU的機制相吻合。另一方面，CUDA提供了片上（on-chip）共享內(nèi)存，使得線程之間可以共享數(shù)據(jù)。應(yīng)用程序可以利用共享內(nèi)存來減少DRAM的數(shù)據(jù)傳送，更少的依賴DRAM的內(nèi)存帶寬。
3. 編程模型
CUDA程序構(gòu)架分為兩部分：Host和Device。一般而言，Host指的是CPU，Device指的是GPU。在CUDA程序構(gòu)架中，主程序還是由CPU來執(zhí)行，而當遇到數(shù)據(jù)并行處理的部分，CUDA 就會將程序編譯成GPU能執(zhí)行的程序，并傳送到GPU。而這個程序在CUDA里稱做核（kernel）。CUDA允許程序員定義稱為核的C語言函數(shù)，從而擴展了C語言，在調(diào)用此類函數(shù)時，它將由N個不同的CUDA線程并行執(zhí)行N次，這與普通的C語言函數(shù)只執(zhí)行一次的方式不同。執(zhí)行核的每個線程都會被分配一個獨特的線程ID，可通過內(nèi)置的threadIdx變量在內(nèi)核中訪問此ID。在 CUDA 程序中，主程序在調(diào)用任何GPU內(nèi)核之前，必須對核進行執(zhí)行配置，即確定線程塊數(shù)和每個線程塊中的線程數(shù)以及共享內(nèi)存大小。
3.1 線程層次結(jié)構(gòu)
在GPU中要執(zhí)行的線程，根據(jù)最有效的數(shù)據(jù)共享來創(chuàng)建塊(Block)，其類型有一維、二維或三維。在同一個塊內(nèi)的線程可彼此協(xié)作，通過一些共享存儲器來共享數(shù)據(jù)，并同步其執(zhí)行來協(xié)調(diào)存儲器訪問。一個塊中的所有線程都必須位于同一個處理器核心中。因而，一個處理器核心的有限存儲器資源制約了每個塊的線程數(shù)量。在早期的NVIDIA 架構(gòu)中，一個線程塊最多可以包含 512個線程，而在后期出現(xiàn)的一些設(shè)備中則最多可支持1024個線程。一般GPU程序線程數(shù)目是很多的，所以不能把所有的線程都塞到同一個塊里。但一個內(nèi)核可由多個大小相同的線程塊同時執(zhí)行，因而線程總數(shù)應(yīng)等于每個塊的線程數(shù)乘以塊的數(shù)量。這些同樣維度和大小的塊將組織為一個一維或二維線程塊網(wǎng)格(Grid)。具體框架如下圖所示。
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">線程塊網(wǎng)格</figcaption>
NOTICE：
線程(Thread)
一般通過GPU的一個核進行處理。（可以表示成一維，二維，三維，具體下面再細說）。
線程塊(Block)
- 由多個線程組成（可以表示成一維，二維，三維，具體下面再細說）。
- 各block是并行執(zhí)行的，block間無法通信，也沒有執(zhí)行順序。
- 注意線程塊的數(shù)量限制為不超過65535（硬件限制）。
線程格(Grid)
由多個線程塊組成（可以表示成一維，二維，三維，具體下面再細說）。
線程束
在CUDA架構(gòu)中，線程束是指一個包含32個線程的集合，這個線程集合被“編織在一起”并且“步調(diào)一致”的形式執(zhí)行。在程序中的每一行，線程束中的每個線程都將在不同數(shù)據(jù)上執(zhí)行相同的命令。
從硬件上看
SP：最基本的處理單元，streaming processor，也稱為CUDA core。最后具體的指令和任務(wù)都是在SP上處理的。GPU進行并行計算，也就是很多個SP同時做處理。
SM：多個SP加上其他的一些資源組成一個streaming multiprocessor。也叫GPU大核，其他資源如：warp scheduler，register，shared memory等。SM可以看做GPU的心臟（對比CPU核心），register和shared memory是SM的稀缺資源。CUDA將這些資源分配給所有駐留在SM中的threads。因此，這些有限的資源就使每個SM中active warps有非常嚴格的限制，也就限制了并行能力。
從軟件上看
thread：一個CUDA的并行程序會被以許多個threads來執(zhí)行。
block：數(shù)個threads會被群組成一個block，同一個block中的threads可以同步，也可以通過shared memory通信。
grid：多個blocks則會再構(gòu)成grid。
warp：GPU執(zhí)行程序時的調(diào)度單位，目前cuda的warp的大小為32，同在一個warp的線程，以不同數(shù)據(jù)資源執(zhí)行相同的指令,這就是所謂 SIMT。
3.2 存儲器層次結(jié)構(gòu)
CUDA設(shè)備擁有多個獨立的存儲空間，其中包括：全局存儲器、本地存儲器、共享存儲器、常量存儲器、紋理存儲器和寄存器，如圖
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">CUDA設(shè)備上的存儲器</figcaption>
NOTICE：
主機(Host)
將CPU及系統(tǒng)的內(nèi)存（內(nèi)存條）稱為主機。
設(shè)備(Device)
將GPU及GPU本身的顯示內(nèi)存稱為設(shè)備。
動態(tài)隨機存取存儲器(DRAM)
DRAM（Dynamic Random Access Memory），即動態(tài)隨機存取存儲器，最為常見的系統(tǒng)內(nèi)存。DRAM只能將數(shù)據(jù)保持很短的時間。為了保持數(shù)據(jù)，DRAM使用電容存儲，所以必須隔一段時間刷新（refresh）一次，如果存儲單元沒有被刷新，存儲的信息就會丟失。 （關(guān)機就會丟失數(shù)據(jù)）
CUDA線程可在執(zhí)行過程中訪問多個存儲器空間的數(shù)據(jù)，如下圖所示其中：
- 每個線程都有一個私有的本地存儲器。
- 每個線程塊都有一個共享存儲器，該存儲器對于塊內(nèi)的所有線程都是可見的，并且與塊具有相同的生命周期。
- 所有線程都可訪問相同的全局存儲器。
- 此外還有兩個只讀的存儲器空間，可由所有線程訪問，這兩個空間是常量存儲器空間和紋理存儲器空間。全局、固定和紋理存儲器空間經(jīng)過優(yōu)化，適于不同的存儲器用途。紋理存儲器也為某些特殊的數(shù)據(jù)格式提供了不同的尋址模式以及數(shù)據(jù)過濾，方便Host對流數(shù)據(jù)的快速存取。
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">存儲器的應(yīng)用層次</figcaption>
3.3 主機（Host）和設(shè)備（Device）
如下圖所示，CUDA假設(shè)線程可在物理上獨立的設(shè)備上執(zhí)行，此類設(shè)備作為運行C語言程序的主機的協(xié)處理器操作。內(nèi)核在GPU上執(zhí)行，而C語言程序的其他部分在CPU上執(zhí)行（即串行代碼在主機上執(zhí)行，而并行代碼在設(shè)備上執(zhí)行）。此外，CUDA還假設(shè)主機和設(shè)備均維護自己的DRAM，分別稱為主機存儲器和設(shè)備存儲器。因而，一個程序通過調(diào)用CUDA運行庫來管理對內(nèi)核可見的全局、固定和紋理存儲器空間。這種管理包括設(shè)備存儲器的分配和取消分配，還包括主機和設(shè)備存儲器之間的數(shù)據(jù)傳輸。
4. CUDA軟硬件
4.1 CUDA術(shù)語
由于CUDA中存在許多概念和術(shù)語，諸如SM、block、SP等多個概念不容易理解，將其與CPU的一些概念進行比較，如下表所示。
| CPU | GPU | 層次 |
|---|---|---|
| 算術(shù)邏輯和控制單元 | 流處理器(SM) | 硬件 |
| 算術(shù)單元 | 批量處理器(SP) | 硬件 |
| 進程 | Block | 軟件 |
| 線程 | thread | 軟件 |
| 調(diào)度單位 | Warp | 軟件 |
4.2 硬件利用率
當為一個GPU分配一個內(nèi)核函數(shù)，我們關(guān)心的是如何才能充分利用GPU的計算能力，但由于不同的硬件有不同的計算能力，SM一次最多能容納的線程數(shù)也不盡相同，SM一次最多能容納的線程數(shù)量主要與底層硬件的計算能力有關(guān)，如下表顯示了在不同的計算能力的設(shè)備上，每個線程塊上開啟不同數(shù)量的線程時設(shè)備的利用率。
| 計算能力 每個線程塊的線程數(shù) | 1.0 | 1.1 | 1.2 | 1.3 | 2.0 | 2.1 | 3.0 |
| 64 | 67 | 50 | 50 | 50 | 33 | 33 | 50 |
| 96 | 100 | 100 | 75 | 75 | 50 | 50 | 75 |
| 128 | 100 | 100 | 100 | 100 | 67 | 67 | 100 |
| 192 | 100 | 100 | 94 | 94 | 100 | 100 | 94 |
| 96 | 100 | 100 | 100 | 100 | 100 | 100 | 10 |
| ... ... | ... ... | | | | | | |
查看顯卡利用率 （以實驗室服務(wù)器為例）
輸入以下命令
nvidia-smi
Thu Aug 23 21:06:36 2018 +-----------------------------------------------------------------------------+| NVIDIA-SMI 384.130 Driver Version: 384.130 ||-------------------------------+----------------------+----------------------+| GPU Name Persistence-M| Bus-Id Disp.A | Volatile Uncorr. ECC || Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. ||===============================+======================+======================|| 0 GeForce GTX 108... Off | 00000000:02:00.0 Off | N/A || 29% 41C P0 58W / 250W | 0MiB / 11171MiB | 0% Default |+-------------------------------+----------------------+----------------------+| 1 GeForce GTX 108... Off | 00000000:03:00.0 Off | N/A || 33% 47C P0 57W / 250W | 0MiB / 11172MiB | 0% Default |+-------------------------------+----------------------+----------------------+| 2 GeForce GTX 108... Off | 00000000:82:00.0 Off | N/A || 36% 49C P0 59W / 250W | 0MiB / 11172MiB | 0% Default |+-------------------------------+----------------------+----------------------+| 3 GeForce GTX 108... Off | 00000000:83:00.0 Off | N/A || 33% 46C P0 51W / 250W | 0MiB / 11172MiB | 1% Default |+-------------------------------+----------------------+----------------------+ +-----------------------------------------------------------------------------+| Processes: GPU Memory || GPU PID Type Process name Usage ||=============================================================================|| No running processes found |+-----------------------------------------------------------------------------+
5. 并行計算
5.1 并發(fā)性
CUDA將問題分解成線程塊的網(wǎng)格，每塊包含多個線程。快可以按任意順序執(zhí)行。不過在某個時間點上，只有一部分塊處于執(zhí)行中。一旦被調(diào)用到GUP包含的N個“流處理器簇（SM）”中的一個上執(zhí)行，一個塊必須從開始到結(jié)束。網(wǎng)格中的塊可以被分配到任意一個有空閑槽的SM上。起初，可以采用“輪詢調(diào)度”策略，以確保分配到每一個SM上的塊數(shù)基本相同。對絕大多數(shù)內(nèi)核程序而言，分塊的數(shù)量應(yīng)該是GPU中物理SM數(shù)量的八倍或更多倍。
以一個軍隊比喻，假設(shè)有一支由士兵（線程）組成的部隊（網(wǎng)格）。部隊被分成若干個連（塊），每個連隊由一位連長來指揮。按照32名士兵一個班（一個線程束），連隊又進一步分成若干個班，每個班由一個班長來指揮。
<figcaption style="box-sizing: border-box; outline: 0px; display: block; text-align: center; margin: 8px; color: rgb(153, 153, 153); font-size: 14px; overflow-wrap: break-word;">基于GPU的線程視圖</figcaption>
要執(zhí)行某個操作，總司令（內(nèi)核程序/ 主機程序）必須提供操作名稱及相應(yīng)的數(shù)據(jù)。每個士兵（線程）只處理分配給他的問題中的一小塊。在連長（負責一個塊）或班長（負責一個束）的控制下，束與束之間的線程或者一個束內(nèi)部的線程之間，要經(jīng)常地交換數(shù)據(jù)。但是，連隊（塊）之間的協(xié)同就得由總司令（內(nèi)核函數(shù)/ 主機程序）來控制。
5.2 局部性
對于GPU程序設(shè)計，程序員必須處理局部性。對于一個給定的工作，他需要事先思考需要哪些工具或零件（即存儲地址或數(shù)據(jù)結(jié)構(gòu)），然后一次性地把他們從硬件倉庫（全局內(nèi)存）可能把與這些數(shù)據(jù)相關(guān)的不同工作都執(zhí)行了，避免發(fā)生“取來--存回--為了下一個工作再取”。
5.3 緩存一致性
GPU與CPU在緩存上的一個重要差別就是“緩存一致性”問題。對于“緩存一致”的系統(tǒng)，一個內(nèi)存的寫操作需要通知所有核的各個級別的緩存。因此，無論何時，所有的處理器核看到的內(nèi)存視圖是完全一樣的。隨著處理器中核數(shù)量的增多，這個“通知”的開銷迅速增大，使得“緩存一致性”成為限制一個處理器中核數(shù)量不能太多的一重要因素?！熬彺嬉恢隆毕到y(tǒng)中最壞的情況是，一個內(nèi)存操作會強迫每個核的緩存都進行更新，進而每個核都要對相鄰的內(nèi)存單元寫操作。
相比之下，非“緩存一致”系統(tǒng)不會自動地更新其他核的緩存。它需要由程序員寫清楚每個處理器核輸出的各自不同的目標區(qū)域。從程序的視角看，這支持一個核僅負責一個輸出或者一個小的輸出集。通常，CPU遵循“緩存一致性”原則，而GPU則不是。故GPU能夠擴展到一個芯片內(nèi)具有大數(shù)量的核心（流處理器簇）。
5.4 弗林分類法
根據(jù)弗林分類法，計算機的結(jié)構(gòu)類型有：
SIMD--單指令，多數(shù)據(jù)
MIMD--多指令，多數(shù)據(jù)
SISD--單指令，單數(shù)據(jù)
MISD--多指令，單數(shù)據(jù)
5.5 分條 / 分塊
CUDA提供的簡單二維網(wǎng)格模型。對于很多問題，這樣的模型就足夠了。如果在一個塊內(nèi)，你的工作是線性分布的，那么你可以很好地將其他分解成CUDA塊。由于在一個SM內(nèi)，最多可以分配16個塊，而在一個GPU內(nèi)有16個（有些是32個）SM，所以問題分成256個甚至更多的塊都可以。實際上，我們更傾向于把一個塊內(nèi)的元素總數(shù)限制為128、256、或者512，這樣有助于在一個典型的數(shù)據(jù)集內(nèi)劃分出更多數(shù)量的塊。
5.6 快速傅氏變換(FFT)
FFT： FFT（Fast Fourier Transformation）是離散傅氏變換（DFT）的快速算法。即為快速傅氏變換。它是根據(jù)離散傅氏變換的奇、偶、虛、實等特性，對離散傅立葉變換的算法進行改進獲得的。
由于不是剛需，這里不展開講。好奇的你可以點擊樓下時光機，通過下面的教程進行學習。
FFT（最詳細最通俗的入門手冊）
5.7 CUDA計算能力的含義
體現(xiàn)GPU計算能力的兩個重要特征：
1)CUDA核的個數(shù)；
2)存儲器大小。
描述GPU性能的兩個重要指標: ：
1)計算性能峰值；
2)存儲器帶寬。
參考
1.CUDA計算能力的含義
2.CUDA GPUs
6. 實踐
6.1 Ubuntu 系統(tǒng)下環(huán)境搭建
6.1.1 系統(tǒng)要求
要搭建 CUDA 環(huán)境，我們需要自己的計算機滿足以下這三個條件：
1. 有至少一顆支持 CUDA 的 GPU（我的是GeForece GT 650M）
2. 有滿足版本要求的 gcc 編譯器和鏈接工具
3. 有 NVIDIA 提供的 CUDA 工具包(點擊神奇的小鏈接下載)
6.1.2 準備工作
下面，我們一步一步來驗證自己的系統(tǒng)是否滿足安裝要求。
Step 1: 驗證計算機是否擁有至少一顆支持 CUDA 的 GPU
打開終端（Ctrl + Alt + T），鍵入以下命令：
lspci | grep -i nvidia
可以看到以下內(nèi)容（結(jié)果因人而異，與具體的GPU有關(guān)）
看到這個就說明至少有一顆支持 CUDA 的 GPU，可以進入下一步了。
Step 2: 驗證一下自己操作系統(tǒng)的版本
鍵入命令：
lsb_release -a
No LSB modules are available.Distributor ID: UbuntuDescription: Ubuntu 16.04.4 LTSRelease: 16.04Codename: xenial
更多信息請移步Ubuntu查看版本信息
Step 3: 驗證 gcc 編譯器的版本
鍵入命令：
gcc --version
或者
gcc -v
得到如下信息
gcc (Ubuntu 5.4.0-6ubuntu1~16.04.10) 5.4.0 20160609Copyright (C) 2015 Free Software Foundation, Inc.This is free software; see the source for copying conditions. There is NOwarranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
Step 4: 驗證系統(tǒng)內(nèi)核版本
鍵入命令：
uname -r
得到如下信息
對照官方提供的對各種 Linux 發(fā)行版的安裝要求進行安裝
6.1.3 搭建 CUDA 環(huán)境
Step 1: 安裝 CUDA 工具包
在前面幾項驗證都順利通過以后就來到最關(guān)鍵的一步。首先下載對應(yīng)自己系統(tǒng)版本的 CUDA 工具包(以CUDA Toolkit 9.2 為例)，然后進入到安裝包所在目錄：
sudo dpkg -i cuda-repo-ubuntu1604-9-2-local_9.2.148-1_amd64.deb
sudo apt-key add /var/cuda-repo-<version>/7fa2af80.pub
sudo apt-get update
sudo apt-get install cuda
NOTICE:
Other installation options are available in the form of meta-packages. For example, to install all the library packages, replace "cuda" with the "cuda-libraries-9-2" meta package. For more information on all the available meta packages click here.
此時靜靜地等待安裝完成。不出意外，一段時間后安裝完成了。
Step 2: 設(shè)置環(huán)境變量
首先在 PATH 變量中加入 /usr/local/cuda-9.2/bin，在Terminal中執(zhí)行：
export PATH=/usr/local/cuda-9.2/bin:$PATH
然后在 LD_LIBRARY_PATH 變量中添加 /usr/local/cuda-9.2/lib64，執(zhí)行：
export LD_LIBRARY_PATH=/usr/local/cuda-9.2/lib64:$LD_LIBRARY_PATH
Step 3: 驗證環(huán)境搭建是否成功
首先執(zhí)行命令：
nvcc -V
關(guān)于測試...聰明的你一定想起來了，我們前面是講過怎么做的。
對，沒錯，就在1.5小節(jié)，話不多說，自行上翻吧。
看到通過測試，到這里，64位 Ubuntu 16.04 系統(tǒng)下 CUDA 環(huán)境搭建就完成了。
6.2 CUDA編程
6.2.1 核函數(shù)
1. 在GPU上執(zhí)行的函數(shù)通常稱為核函數(shù)。
2. 一般通過標識符__global__修飾，調(diào)用通過<<<參數(shù)1,參數(shù)2>>>，用于說明內(nèi)核函數(shù)中的線程數(shù)量，以及線程是如何組織的。
3. 以線程格（Grid）的形式組織，每個線程格由若干個線程塊（block）組成，而每個線程塊又由若干個線程（thread）組成。
4.是以block為單位執(zhí)行的。
5. 叧能在主機端代碼中調(diào)用。
6. 調(diào)用時必須聲明內(nèi)核函數(shù)的執(zhí)行參數(shù)。
7. 在編程時，必須先為kernel函數(shù)中用到的數(shù)組或變量分配好足夠的空間，再調(diào)用kernel函數(shù)，否則在GPU計算時會發(fā)生錯誤，例如越界或報錯，甚至導致藍屏和死機。
看完基本知識，裝好CUDA以后，就可以開始寫第一個CUDA程序了：
#include <cuda_runtime.h> int main(){printf("Hello world!\n");}
慢著，這個程序和C有什么區(qū)別?用到顯卡了嗎?
答：沒有區(qū)別，沒用顯卡。如果你非要用顯卡干點什么事情的話，可以改成這個樣子：
/* * @file_name HelloWorld.cu 后綴名稱.cu */ #include <stdio.h>#include <cuda_runtime.h> //頭文件 //核函數(shù)聲明，前面的關(guān)鍵字__global____global__ void kernel( void ) {} int main( void ) { //核函數(shù)的調(diào)用，注意<<<1,1>>>，第一個1，代表線程格里只有一個線程塊；第二個1，代表一個線程塊里只有一個線程。 kernel<<<1,1>>>(); printf( "Hello, World!\n" ); return 0;}
6.2.2 dim3結(jié)構(gòu)類型
- dim3是基于uint3定義的矢量類型，相當亍由3個unsigned int型組成的結(jié)構(gòu)體。uint3類型有三個數(shù)據(jù)成員
unsigned int x;unsigned int y;unsigned int z; - 可使用于一維、二維或三維的索引來標識線程，構(gòu)成一維、二維或三維線程塊。
-
dim3結(jié)構(gòu)類型變量用在核函數(shù)調(diào)用的<<<,>>>中。 - 相關(guān)的幾個內(nèi)置變量
4.1.threadIdx，顧名思義獲取線程thread的ID索引；如果線程是一維的那么就取threadIdx.x，二維的還可以多取到一個值threadIdx.y，以此類推到三維threadIdx.z。
4.2.blockIdx，線程塊的ID索引；同樣有blockIdx.x，blockIdx.y，blockIdx.z。
4.3.blockDim，線程塊的維度，同樣有blockDim.x，blockDim.y，blockDim.z。
4.4.gridDim，線程格的維度，同樣有gridDim.x，gridDim.y，gridDim.z。 - 對于一維的
block，線程的threadID=threadIdx.x。 - 對于大小為
（blockDim.x, blockDim.y）的 二維block，線程的threadID=threadIdx.x+threadIdx.y*blockDim.x。 - 對于大小為
（blockDim.x, blockDim.y, blockDim.z）的 三維block，線程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。 - 對于計算線程索引偏移增量為已啟動線程的總數(shù)。如
stride = blockDim.x * gridDim.x; threadId += stride。
6.2.3 函數(shù)修飾符
1.__global__，表明被修飾的函數(shù)在設(shè)備上執(zhí)行，但在主機上調(diào)用。
2.__device__，表明被修飾的函數(shù)在設(shè)備上執(zhí)行，但只能在其他__device__函數(shù)或者__global__函數(shù)中調(diào)用。
6.2.4 常用的GPU內(nèi)存函數(shù)
cudaMalloc()
1. 函數(shù)原型： cudaError_t cudaMalloc (void **devPtr, size_t size)。
2. 函數(shù)用處：與C語言中的malloc函數(shù)一樣，只是此函數(shù)在GPU的內(nèi)存你分配內(nèi)存。
3. 注意事項：
3.1. 可以將cudaMalloc()分配的指針傳遞給在設(shè)備上執(zhí)行的函數(shù)；
3.2. 可以在設(shè)備代碼中使用cudaMalloc()分配的指針進行設(shè)備內(nèi)存讀寫操作；
3.3. 可以將cudaMalloc()分配的指針傳遞給在主機上執(zhí)行的函數(shù)；
3.4. 不可以在主機代碼中使用cudaMalloc()分配的指針進行主機內(nèi)存讀寫操作（即不能進行解引用）。
cudaMemcpy()
1. 函數(shù)原型：cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)。
2. 函數(shù)作用：與c語言中的memcpy函數(shù)一樣，只是此函數(shù)可以在主機內(nèi)存和GPU內(nèi)存之間互相拷貝數(shù)據(jù)。
3. 函數(shù)參數(shù)：cudaMemcpyKind kind表示數(shù)據(jù)拷貝方向，如果kind賦值為cudaMemcpyDeviceToHost表示數(shù)據(jù)從設(shè)備內(nèi)存拷貝到主機內(nèi)存。
4. 與C中的memcpy()一樣，以同步方式執(zhí)行，即當函數(shù)返回時，復制操作就已經(jīng)完成了，并且在輸出緩沖區(qū)中包含了復制進去的內(nèi)容。
5. 相應(yīng)的有個異步方式執(zhí)行的函數(shù)cudaMemcpyAsync()，這個函數(shù)詳解請看下面的流一節(jié)有關(guān)內(nèi)容。
cudaFree()
1. 函數(shù)原型：cudaError_t cudaFree ( void* devPtr )。
2. 函數(shù)作用：與c語言中的free()函數(shù)一樣，只是此函數(shù)釋放的是cudaMalloc()分配的內(nèi)存。
下面實例用于解釋上面三個函數(shù)
#include <stdio.h>#include <cuda_runtime.h>__global__ void add( int a, int b, int *c ) { *c = a + b;}int main( void ) { int c; int *dev_c; //cudaMalloc() cudaMalloc( (void**)&dev_c, sizeof(int) ); //核函數(shù)執(zhí)行 add<<<1,1>>>( 2, 7, dev_c ); //cudaMemcpy() cudaMemcpy( &c, dev_c, sizeof(int),cudaMemcpyDeviceToHost ) ; printf( "2 + 7 = %d\n", c ); //cudaFree() cudaFree( dev_c ); return 0;}
6.2.5 GPU內(nèi)存分類
全局內(nèi)存
通俗意義上的設(shè)備內(nèi)存。
共享內(nèi)存
1. 位置：設(shè)備內(nèi)存。
2. 形式：關(guān)鍵字__shared__添加到變量聲明中。如__shared__ float cache[10]。
3. 目的：對于GPU上啟動的每個線程塊，CUDAC編譯器都將創(chuàng)建該共享變量的一個副本。線程塊中的每個線程都共享這塊內(nèi)存，但線程卻無法看到也不能修改其他線程塊的變量副本。這樣使得一個線程塊中的多個線程能夠在計算上通信和協(xié)作。
常量內(nèi)存
1. 位置：設(shè)備內(nèi)存
2. 形式：關(guān)鍵字__constant__添加到變量聲明中。如__constant__ float s[10];。
3. 目的：為了提升性能。常量內(nèi)存采取了不同于標準全局內(nèi)存的處理方式。在某些情況下，用常量內(nèi)存替換全局內(nèi)存能有效地減少內(nèi)存帶寬。
4. 特點：常量內(nèi)存用于保存在核函數(shù)執(zhí)行期間不會發(fā)生變化的數(shù)據(jù)。變量的訪問限制為只讀。NVIDIA硬件提供了64KB的常量內(nèi)存。不再需要cudaMalloc()或者cudaFree(),而是在編譯時，靜態(tài)地分配空間。
5. 要求：當我們需要拷貝數(shù)據(jù)到常量內(nèi)存中應(yīng)該使用cudaMemcpyToSymbol()，而cudaMemcpy()會復制到全局內(nèi)存。
6. 性能提升的原因：
6.1. 對常量內(nèi)存的單次讀操作可以廣播到其他的“鄰近”線程。這將節(jié)約15次讀取操作。（為什么是15，因為“鄰近”指半個線程束，一個線程束包含32個線程的集合。）
6.2. 常量內(nèi)存的數(shù)據(jù)將緩存起來，因此對相同地址的連續(xù)讀操作將不會產(chǎn)生額外的內(nèi)存通信量。
紋理內(nèi)存
1. 位置：設(shè)備內(nèi)存
2. 目的：能夠減少對內(nèi)存的請求并提供高效的內(nèi)存帶寬。是專門為那些在內(nèi)存訪問模式中存在大量空間局部性的圖形應(yīng)用程序設(shè)計，意味著一個線程讀取的位置可能與鄰近線程讀取的位置“非常接近”。如下圖：
3. 紋理變量（引用）必須聲明為文件作用域內(nèi)的全局變量。
4. 形式：分為一維紋理內(nèi)存 和 二維紋理內(nèi)存。
4.1. 一維紋理內(nèi)存
4.1.1. 用texture<類型>類型聲明，如texture<float> texIn。
4.1.2. 通過cudaBindTexture()綁定到紋理內(nèi)存中。
4.1.3. 通過tex1Dfetch()來讀取紋理內(nèi)存中的數(shù)據(jù)。
4.1.4. 通過cudaUnbindTexture()取消綁定紋理內(nèi)存。
4.2. 二維紋理內(nèi)存
4.2.1. 用texture<類型,數(shù)字>類型聲明，如texture<float，2> texIn。
4.2.2. 通過cudaBindTexture2D()綁定到紋理內(nèi)存中。
4.2.3. 通過tex2D()來讀取紋理內(nèi)存中的數(shù)據(jù)。
4.2.4. 通過cudaUnbindTexture()取消綁定紋理內(nèi)存。
固定內(nèi)存
1. 位置：主機內(nèi)存。
2. 概念：也稱為頁鎖定內(nèi)存或者不可分頁內(nèi)存，操作系統(tǒng)將不會對這塊內(nèi)存分頁并交換到磁盤上，從而確保了該內(nèi)存始終駐留在物理內(nèi)存中。因此操作系統(tǒng)能夠安全地使某個應(yīng)用程序訪問該內(nèi)存的物理地址，因為這塊內(nèi)存將不會破壞或者重新定位。
3. 目的：提高訪問速度。由于GPU知道主機內(nèi)存的物理地址，因此可以通過“直接內(nèi)存訪問DMA（Direct Memory Access)技術(shù)來在GPU和主機之間復制數(shù)據(jù)。由于DMA在執(zhí)行復制時無需CPU介入。因此DMA復制過程中使用固定內(nèi)存是非常重要的。
4. 缺點：使用固定內(nèi)存，將失去虛擬內(nèi)存的所有功能；系統(tǒng)將更快的耗盡內(nèi)存。
5. 建議：對cudaMemcpy()函數(shù)調(diào)用中的源內(nèi)存或者目標內(nèi)存，才使用固定內(nèi)存，并且在不再需要使用它們時立即釋放。
6. 形式：通過cudaHostAlloc()函數(shù)來分配；通過cudaFreeHost()釋放。
7. 只能以異步方式對固定內(nèi)存進行復制操作。
原子性
1. 概念：如果操作的執(zhí)行過程不能分解為更小的部分，我們將滿足這種條件限制的操作稱為原子操作。
2. 形式：函數(shù)調(diào)用，如atomicAdd（addr,y)將生成一個原子的操作序列，這個操作序列包括讀取地址addr處的值，將y增加到這個值，以及將結(jié)果保存回地址addr。
6.2.6 常用線程操作函數(shù)
同步方法__syncthreads()，這個函數(shù)的調(diào)用，將確保線程塊中的每個線程都執(zhí)行完__syscthreads()前面的語句后，才會執(zhí)行下一條語句。
使用事件來測量性能
1. 用途：為了測量GPU在某個任務(wù)上花費的時間。CUDA中的事件本質(zhì)上是一個GPU時間戳。由于事件是直接在GPU上實現(xiàn)的。因此不適用于對同時包含設(shè)備代碼和主機代碼的混合代碼設(shè)計。
2. 形式：首先創(chuàng)建一個事件，然后記錄事件，再計算兩個事件之差，最后銷毀事件。如：
cudaEvent_t start, stop;cudaEventCreate( &start );cudaEventCreate( &stop );cudaEventRecord( start, 0 );//do somethingcudaEventRecord( stop, 0 );float elapsedTime;cudaEventElapsedTime( &elapsedTime,start, stop );cudaEventDestroy( start );cudaEventDestroy( stop )；
6.2.7 流
- 扯一扯：并發(fā)重點在于一個極短時間段內(nèi)運行多個不同的任務(wù)；并行重點在于同時運行一個任務(wù)。
- 任務(wù)并行性：是指并行執(zhí)行兩個或多個不同的任務(wù)，而不是在大量數(shù)據(jù)上執(zhí)行同一個任務(wù)。
- 概念：CUDA流表示一個GPU操作隊列，并且該隊列中的操作將以指定的順序執(zhí)行。我們可以在流中添加一些操作，如核函數(shù)啟動，內(nèi)存復制以及事件的啟動和結(jié)束等。這些操作的添加到流的順序也是它們的執(zhí)行順序?？梢詫⒚總€流視為GPU上的一個任務(wù)，并且這些任務(wù)可以并行執(zhí)行。
- 硬件前提：必須是支持設(shè)備重疊功能的GPU。支持設(shè)備重疊功能，即在執(zhí)行一個核函數(shù)的同時，還能在設(shè)備與主機之間執(zhí)行復制操作。
- 聲明與創(chuàng)建：聲明
cudaStream_t stream;，創(chuàng)建cudaSteamCreate(&stream);。 -
cudaMemcpyAsync()：前面在cudaMemcpy()中提到過，這是一個以異步方式執(zhí)行的函數(shù)。在調(diào)用cudaMemcpyAsync()時，只是放置一個請求，表示在流中執(zhí)行一次內(nèi)存復制操作，這個流是通過參數(shù)stream來指定的。當函數(shù)返回時，我們無法確保復制操作是否已經(jīng)啟動，更無法保證它是否已經(jīng)結(jié)束。我們能夠得到的保證是，復制操作肯定會當下一個被放入流中的操作之前執(zhí)行。傳遞給此函數(shù)的主機內(nèi)存指針必須是通過cudaHostAlloc()分配好的內(nèi)存。（流中要求固定內(nèi)存） - 流同步：通過
cudaStreamSynchronize()來協(xié)調(diào)。 - 流銷毀：在退出應(yīng)用程序之前，需要銷毀對GPU操作進行排隊的流，調(diào)用
cudaStreamDestroy()。 - 針對多個流：
9.1. 記得對流進行同步操作。
9.2. 將操作放入流的隊列時，應(yīng)采用寬度優(yōu)先方式，而非深度優(yōu)先的方式，換句話說，不是首先添加第0個流的所有操作，再依次添加后面的第1，2,…個流。而是交替進行添加，比如將a的復制操作添加到第0個流中，接著把a的復制操作添加到第1個流中，再繼續(xù)其他的類似交替添加的行為。
9.3. 要牢牢記住操作放入流中的隊列中的順序影響到CUDA驅(qū)動程序調(diào)度這些操作和流以及執(zhí)行的方式。
TIPS:
- 當線程塊的數(shù)量為GPU中處理數(shù)量的2倍時，將達到最優(yōu)性能。
- 核函數(shù)執(zhí)行的第一個計算就是計算輸入數(shù)據(jù)的偏移。每個線程的起始偏移都是0到線程數(shù)量減1之間的某個值。然后，對偏移的增量為已啟動線程的總數(shù)。
6.2.8 這是一個栗子
我們嘗試用一個程序來比較cuda/c在GPU/CPU的運行效率，來不及了，快上車。
這是一個CUDA程序，請保存文件名為“文件名.cu”,在你的PC或者服務(wù)器上運行。
#include "cuda_runtime.h"#include "device_launch_parameters.h" #include <stdio.h>#include <time.h> #define N (1024*1024)#define M (10000)#define THREADS_PER_BLOCK 1024 void serial_add(double *a, double *b, double *c, int n, int m){ for(int index=0;index<n;index++) { for(int j=0;j<m;j++) { c[index] = a[index]*a[index] + b[index]*b[index]; } }} __global__ void vector_add(double *a, double *b, double *c){ int index = blockIdx.x * blockDim.x + threadIdx.x; for(int j=0;j<M;j++) { c[index] = a[index]*a[index] + b[index]*b[index]; }} int main(){ clock_t start,end; double *a, *b, *c; int size = N * sizeof( double ); a = (double *)malloc( size ); b = (double *)malloc( size ); c = (double *)malloc( size ); for( int i = 0; i < N; i++ ) { a[i] = b[i] = i; c[i] = 0; } start = clock(); serial_add(a, b, c, N, M); printf( "c[%d] = %f\n",0,c[0] ); printf( "c[%d] = %f\n",N-1, c[N-1] ); end = clock(); float time1 = ((float)(end-start))/CLOCKS_PER_SEC; printf("CPU: %f seconds\n",time1); start = clock(); double *d_a, *d_b, *d_c; cudaMalloc( (void **) &d_a, size ); cudaMalloc( (void **) &d_b, size ); cudaMalloc( (void **) &d_c, size ); cudaMemcpy( d_a, a, size, cudaMemcpyHostToDevice ); cudaMemcpy( d_b, b, size, cudaMemcpyHostToDevice ); vector_add<<< (N + (THREADS_PER_BLOCK-1)) / THREADS_PER_BLOCK, THREADS_PER_BLOCK >>>( d_a, d_b, d_c ); cudaMemcpy( c, d_c, size, cudaMemcpyDeviceToHost ); printf( "c[%d] = %f\n",0,c[0] ); printf( "c[%d] = %f\n",N-1, c[N-1] ); free(a); free(b); free(c); cudaFree( d_a ); cudaFree( d_b ); cudaFree( d_c ); end = clock(); float time2 = ((float)(end-start))/CLOCKS_PER_SEC; printf("CUDA: %f seconds, Speedup: %f\n",time2, time1/time2); return 0;}
效率對比
我們通過修改count的值并且加大循環(huán)次數(shù)來觀察變量的效率的差別。
運行結(jié)果：
可見在數(shù)據(jù)量大的情況下效率還是相當不錯的。
7. GPU or FPGA
GPU優(yōu)勢
1.從峰值性能來說，GPU（10Tflops)遠遠高于FPGA(<1TFlops);
2.GPU相對于FPGA還有一個優(yōu)勢就是內(nèi)存接口, GPU的內(nèi)存接口（傳統(tǒng)的GDDR5，最近更是用上了HBM和HBM2）的帶寬遠好于FPGA的傳統(tǒng)DDR接口（大約帶寬高4-5倍）;
3.功耗方面，雖然GPU的功耗遠大于FPGA的功耗，但是如果要比較功耗應(yīng)該比較在執(zhí)行效率相同時需要的功耗。如果FPGA的架構(gòu)優(yōu)化能做到很好以致于一塊FPGA的平均性能能夠接近一塊GPU，那么FPGA方案的總功耗遠小于GPU，散熱問題可以大大減輕。反之，如果需要二十塊FPGA才能實現(xiàn)一塊GPU的平均性能，那么FPGA在功耗方面并沒有優(yōu)勢。
4.FPGA缺點有三點:
第一，基本單元的計算能力有限。為了實現(xiàn)可重構(gòu)特性，F(xiàn)PGA 內(nèi)部有大量極細粒度的基本單元，但是每個單元的計算能力（主要依靠LUT 查找表）都遠遠低于CPU 和GPU 中的ALU模塊。
第二，速度和功耗相對專用定制芯片（ASIC）仍然存在不小差距。
第三，F(xiàn)PGA 價格較為昂貴，在規(guī)模放量的情況下單塊FPGA 的成本要遠高于專用定制芯片。最后誰能勝出, 完全取決于FPGA架構(gòu)優(yōu)化能否彌補峰值性能的劣勢。
5.個人更推薦: CPU+FPGA的組合模式; 其中FPGA用于整形計算，cpu進行浮點計算和調(diào)度，此組合的擁有更高的單位功耗性能和更低的時延。最后更想GPU穩(wěn)定開放,發(fā)揮其長處, 達到真正的物美價廉!
FPGA優(yōu)勢
人工智能目前仍處于早期階段，未來人工智能的主戰(zhàn)場是在推理環(huán)節(jié)，遠沒有爆發(fā)。未來勝負尚未可知，各家技術(shù)路線都有機會勝出。目前英偉達的GPU在訓練場景中占據(jù)著絕對領(lǐng)導地位，但是在未來，專注于推理環(huán)節(jié)的FPGA必將會發(fā)揮巨大的價值。
FPGA和GPU內(nèi)都有大量的計算單元，因此它們的計算能力都很強。在進行神經(jīng)網(wǎng)絡(luò)運算的時候，兩者的速度會比CPU快很多。但是GPU由于架構(gòu)固定，硬件原生支持的指令也就固定了，而FPGA則是可編程的。其可編程性是關(guān)鍵，因為它讓軟件與終端應(yīng)用公司能夠提供與其競爭對手不同的解決方案，并且能夠靈活地針對自己所用的算法修改電路。
在平均性能方面，GPU遜于FPGA，F(xiàn)PGA可以根據(jù)特定的應(yīng)用去編程硬件，例如如果應(yīng)用里面的加法運算非常多就可以把大量的邏輯資源去實現(xiàn)加法器，而GPU一旦設(shè)計完就不能改動了，所以不能根據(jù)應(yīng)用去調(diào)整硬件資源。
目前機器學習大多使用SIMD架構(gòu)，即只需一條指令可以平行處理大量數(shù)據(jù)，因此用GPU很適合。但是有些應(yīng)用是MISD，即單一數(shù)據(jù)需要用許多條指令平行處理，這種情況下用FPGA做一個MISD的架構(gòu)就會比GPU有優(yōu)勢。 所以，對于平均性能，看的就是FPGA加速器架構(gòu)上的優(yōu)勢是否能彌補運行速度上的劣勢。如果FPGA上的架構(gòu)優(yōu)化可以帶來相比GPU架構(gòu)兩到三個數(shù)量級的優(yōu)勢，那么FPGA在平均性能上會好于GPU。
在功耗能效比方面，同樣由于FPGA的靈活性，在架構(gòu)優(yōu)化到很好時，一塊FPGA的平均性能能夠接近一塊GPU，那么FPGA方案的總功耗遠小于GPU，散熱問題可以大大減輕。 能效比的比較也是類似，能效指的是完成程序執(zhí)行消耗的能量，而能量消耗等于功耗乘以程序的執(zhí)行時間。雖然GPU的功耗遠大于FPGA的功耗，但是如果FPGA執(zhí)行相同程序需要的時間比GPU長幾十倍，那FPGA在能效比上就沒有優(yōu)勢了；反之如果FPGA上實現(xiàn)的硬件架構(gòu)優(yōu)化得很適合特定的機器學習應(yīng)用，執(zhí)行算法所需的時間僅僅是GPU的幾倍或甚至于接近GPU，那么FPGA的能效比就會比GPU強。
在峰值性能比方面，雖然GPU的峰值性能（10Tflops）遠大于FPGA的峰值性能（<1Tflops），但針對特定的場景來講吞吐量并不比GPU差。
8. 深度學習的三種硬件方案：ASIC，F(xiàn)PGA，GPU
8.1 對深度學習硬件平臺的要求
要想明白“深度學習”需要怎樣的硬件，必須了解深度學習的工作原理。首先在表層上，我們有一個巨大的數(shù)據(jù)集，并選定了一種深度學習模型。每個模型都有一些內(nèi)部參數(shù)需要調(diào)整，以便學習數(shù)據(jù)。而這種參數(shù)調(diào)整實際上可以歸結(jié)為優(yōu)化問題，在調(diào)整這些參數(shù)時，就相當于在優(yōu)化特定的約束條件。
[圖片上傳失敗...(image-cb45d0-1639707754355)]
矩陣相乘（Matrix Multiplication）——幾乎所有的深度學習模型都包含這一運算，它的計算十分密集。
卷積（Convolution）——這是另一個常用的運算，占用了模型中大部分的每秒浮點運算（浮點／秒）。
循環(huán)層（Recurrent Layers ）——模型中的反饋層，并且基本上是前兩個運算的組合。
All Reduce——這是一個在優(yōu)化前對學習到的參數(shù)進行傳遞或解析的運算序列。在跨硬件分布的深度學習網(wǎng)絡(luò)上執(zhí)行同步優(yōu)化時（如AlphaGo的例子），這一操作尤其有效。
除此之外，深度學習的硬件加速器需要具備數(shù)據(jù)級別和流程化的并行性、多線程和高內(nèi)存帶寬等特性。 另外，由于數(shù)據(jù)的訓練時間很長，所以硬件架構(gòu)必須低功耗。 因此，效能功耗比（Performance per Watt）是硬件架構(gòu)的評估標準之一。
CNN在應(yīng)用中，一般采用GPU加速，請解釋為什么GPU可以有加速效果，主要加速算法的哪一個部分？
這里默認gpu加速是指NVIDIA的CUDA加速。CPU是中央處理單元，gpu是圖形處理單元，gpu由上千個流處理器(core)作為運算器。執(zhí)行采用單指令多線程(SIMT)模式。相比于單核CPU（向量機）流水線式的串行操作，雖然gpu單個core計算能力很弱，但是通過大量線程進行同時計算，在數(shù)據(jù)量很大是會活動較為可觀的加速效果。
具體到cnn，利用gpu加速主要是在conv（卷積）過程上。conv過程同理可以像以上的向量加法一樣通過cuda實現(xiàn)并行化。具體的方法很多，不過最好的還是利用fft（快速傅里葉變換）進行快速卷積。NVIDIA提供了cufft庫實現(xiàn)fft，復數(shù)乘法則可以使用cublas庫里的對應(yīng)的level3的cublasCgemm函數(shù)。
GPU加速的基本準則就是“人多力量大”。CNN說到底主要問題就是計算量大，但是卻可以比較有效的拆分成并行問題。隨便拿一個層的filter來舉例子，假設(shè)某一層有n個filter，每一個需要對上一層輸入過來的map進行卷積操作。那么，這個卷積操作并不需要按照線性的流程去做，每個濾波器互相之間并不影響，可以大家同時做，然后大家生成了n張新的譜之后再繼續(xù)接下來的操作。既然可以并行，那么同一時間處理單元越多，理論上速度優(yōu)勢就會越大。所以，處理問題就變得很簡單粗暴，就像NV那樣，暴力增加顯卡單元數(shù)（當然，顯卡的架構(gòu)、內(nèi)部數(shù)據(jù)的傳輸速率、算法的優(yōu)化等等也都很重要）。
GPU主要是針對圖形顯示及渲染等技術(shù)的出眾，而其中的根本是因為處理矩陣算法能力的強大，剛好CNN中涉及大量的卷積，也就是矩陣乘法等，所以在這方面具有優(yōu)勢。
機器學習的算法一定得經(jīng)過gpu加速嗎？
不一定。只有需要大量浮點數(shù)計算，例如矩陣乘法，才需要GPU加速。 用CNN對圖像進行分類就是一個需要大量浮點數(shù)計算的典型案例，通常需要GPU加速
對于ASIC、FPGA、分布式計算，這里不再展開講，有興趣的小伙伴可以，自行學習。不過....說不定某天博主心情好，就會梳理一下這幾種硬件方案在端到端上應(yīng)用的區(qū)別了。
菜鳥入門教程就到這里了，聰明的你一定不滿足這個入門教程，如有興趣進一步學習CUDA編程，可移步NVIDIA官方的課程平臺CUDA ZONE（PS：中文網(wǎng)站，英文課程）
