CUDA——"從入門到放棄"

開篇一張圖旬蟋，后面聽我編

1. 知識準(zhǔn)備

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)如下圖所示：

CPU微架構(gòu)示意圖

什么？架構(gòu)記不赘饰睢含潘？來，我們換種表示方法：

CPU微架構(gòu)示意圖（改）

嗯线婚，大概就是這個意思遏弱。

從字面上我們也很好理解，計算單元主要執(zhí)行算術(shù)運算塞弊、移位等操作以及地址運算和轉(zhuǎn)換腾窝；存儲單元主要用于保存運算中產(chǎn)生的數(shù)據(jù)以及指令等；控制單元則對指令譯碼居砖，并且發(fā)出為完成每條指令所要執(zhí)行的各個操作的控制信號虹脯。

所以一條指令在CPU中執(zhí)行的過程是這樣的：讀取到指令后，通過指令總線送到控制器（黃色區(qū)域）中進(jìn)行譯碼奏候，并發(fā)出相應(yīng)的操作控制信號循集；然后運算器（綠色區(qū)域）按照操作指令對數(shù)據(jù)進(jìn)行計算，并通過數(shù)據(jù)總線將得到的數(shù)據(jù)存入數(shù)據(jù)緩存器（大塊橙色區(qū)域）蔗草。過程如下圖所示：

CPU執(zhí)行指令圖

是不是有點兒復(fù)雜咒彤？沒關(guān)系疆柔，這張圖完全不用記住，我們只需要知道镶柱，CPU遵循的是馮諾依曼架構(gòu)旷档，其核心就是：存儲程序，順序執(zhí)行歇拆。

講到這里鞋屈，有沒有看出問題，沒錯——在這個結(jié)構(gòu)圖中故觅，負(fù)責(zé)計算的綠色區(qū)域占的面積似乎太小了厂庇，而橙色區(qū)域的緩存Cache和黃色區(qū)域的控制單元占據(jù)了大量空間。

高中化學(xué)有句老生常談的話叫：結(jié)構(gòu)決定性質(zhì)输吏，放在這里也非常適用权旷。

因為CPU的架構(gòu)中需要大量的空間去放置存儲單元（橙色部分）和控制單元（黃色部分），相比之下計算單元（綠色部分）只占據(jù)了很小的一部分贯溅，所以它在大規(guī)模并行計算能力上極受限制拄氯，而更擅長于邏輯控制。

另外它浅，因為遵循馮諾依曼架構(gòu)（存儲程序坤邪，順序執(zhí)行），CPU就像是個一板一眼的管家罚缕，人們吩咐的事情它總是一步一步來做筹吐。但是隨著人們對更大規(guī)模與更快處理速度的需求的增加苇倡，這位管家漸漸變得有些力不從心断医。

于是廷雅，大家就想，能不能把多個處理器放在同一塊芯片上腌乡，讓它們一起來做事盟劫，這樣效率不就提高了嗎？

沒錯与纽，GPU便由此誕生了侣签。

1.2 顯卡

顯卡（Video card，Graphics card）全稱顯示接口卡急迂，又稱顯示適配器影所，是計算機最基本配置、最重要的配件之一僚碎。顯卡作為電腦主機里的一個重要組成部分猴娩，是電腦進(jìn)行數(shù)模信號轉(zhuǎn)換的設(shè)備，承擔(dān)輸出顯示圖形的任務(wù)。顯卡接在電腦主板上卷中，它將電腦的數(shù)字信號轉(zhuǎn)換成模擬信號讓顯示器顯示出來矛双，同時顯卡還是有圖像處理能力，可協(xié)助CPU工作蟆豫，提高整體的運行速度议忽。對于從事專業(yè)圖形設(shè)計的人來說顯卡非常重要。民用和軍用顯卡圖形芯片供應(yīng)商主要包括AMD(超微半導(dǎo)體)和Nvidia(英偉達(dá))2家∈酰現(xiàn)在的top500計算機栈幸，都包含顯卡計算核心。在科學(xué)計算中嫉称，顯卡被稱為顯示加速卡侦镇。

為什么GPU特別擅長處理圖像數(shù)據(jù)呢灵疮？這是因為圖像上的每一個像素點都有被處理的需要织阅，而且每個像素點處理的過程和方式都十分相似，也就成了GPU的天然溫床震捣。

GPU微架構(gòu)示意圖

從架構(gòu)圖我們就能很明顯的看出荔棉，GPU的構(gòu)成相對簡單，有數(shù)量眾多的計算單元和超長的流水線蒿赢，特別適合處理大量的類型統(tǒng)一的數(shù)據(jù)润樱。

再把CPU和GPU兩者放在一張圖上看下對比，就非常一目了然了羡棵。

GPU的工作大部分都計算量大壹若，但沒什么技術(shù)含量，而且要重復(fù)很多很多次皂冰。

但GPU無法單獨工作店展，必須由CPU進(jìn)行控制調(diào)用才能工作。CPU可單獨作用秃流，處理復(fù)雜的邏輯運算和不同的數(shù)據(jù)類型赂蕴，但當(dāng)需要大量的處理類型統(tǒng)一的數(shù)據(jù)時，則可調(diào)用GPU進(jìn)行并行計算舶胀。

借用知乎上某大佬的說法概说，就像你有個工作需要計算幾億次一百以內(nèi)加減乘除一樣，最好的辦法就是雇上幾十個小學(xué)生一起算嚣伐，一人算一部分糖赔，反正這些計算也沒什么技術(shù)含量，純粹體力活而已轩端；而CPU就像老教授挂捻，積分微分都會算，就是工資高，一個老教授資頂二十個小學(xué)生刻撒，你要是富士康你雇哪個骨田？

注：GPU中有很多的運算器ALU和很少的緩存cache，緩存的目的不是保存后面需要訪問的數(shù)據(jù)的声怔，這點和CPU不同态贤，而是為線程thread提高服務(wù)的。如果有很多線程需要訪問同一個相同的數(shù)據(jù)醋火，緩存會合并這些訪問悠汽，然后再去訪問dram。

可愛的你如果對CUDA硬件有更多的興趣芥驳，可移步NVIDIA中文官網(wǎng)進(jìn)一步學(xué)習(xí)柿冲。

1.3 內(nèi)存

內(nèi)存是計算機中重要的部件之一，它是與CPU進(jìn)行溝通的橋梁兆旬。計算機中所有程序的運行都是在內(nèi)存中進(jìn)行的假抄，因此內(nèi)存的性能對計算機的影響非常大。內(nèi)存(Memory)也被稱為內(nèi)存儲器丽猬，其作用是用于暫時存放CPU中的運算數(shù)據(jù)宿饱，以及與硬盤等外部存儲器交換的數(shù)據(jù)。只要計算機在運行中脚祟，CPU就會把需要運算的數(shù)據(jù)調(diào)到內(nèi)存中進(jìn)行運算谬以，當(dāng)運算完成后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，深度學(xué)習(xí)本身需要大量計算骤星。GPU的并行計算能力经瓷，在過去幾年里恰當(dāng)?shù)貪M足了深度學(xué)習(xí)的需求。AMD的GPU基本沒有什么支持洞难，可以不用考慮舆吮。

驅(qū)動：沒有顯卡驅(qū)動，就不能識別GPU硬件，不能調(diào)用其計算資源色冀。但是呢潭袱，NVIDIA在Linux上的驅(qū)動安裝特別麻煩，尤其對于新手簡直就是噩夢锋恬。得屏蔽第三方顯卡驅(qū)動屯换。下面會給出教程。

CUDA：是NVIDIA推出的只能用于自家GPU的并行計算框架与学。只有安裝這個框架才能夠進(jìn)行復(fù)雜的并行計算彤悔。主流的深度學(xué)習(xí)框架也都是基于CUDA進(jìn)行GPU并行加速的，幾乎無一例外索守。還有一個叫做cudnn晕窑，是針對深度卷積神經(jīng)網(wǎng)絡(luò)的加速庫。

查看顯卡驅(qū)動信息(以實驗室服務(wù)器為例)

ssh ubuntu@192.168.1.158

輸入服務(wù)器密碼登陸
然后卵佛，進(jìn)入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 Ti
Result = PASS

大家可以在自己PC或者工作機上嘗試一下杨赤。

再啰嗦兩句

GPU就是用很多簡單的計算單元去完成大量的計算任務(wù)，純粹的人海戰(zhàn)術(shù)级遭。這種策略基于一個前提望拖，就是小學(xué)生A和小學(xué)生B的工作沒有什么依賴性渺尘，是互相獨立的挫鸽。

但有一點需要強調(diào)，雖然GPU是為了圖像處理而生的鸥跟，但是我們通過前面的介紹可以發(fā)現(xiàn)丢郊，它在結(jié)構(gòu)上并沒有專門為圖像服務(wù)的部件，只是對CPU的結(jié)構(gòu)進(jìn)行了優(yōu)化與調(diào)整医咨，所以現(xiàn)在GPU不僅可以在圖像處理領(lǐng)域大顯身手枫匾，它還被用來科學(xué)計算、密碼破解拟淮、數(shù)值分析干茉，海量數(shù)據(jù)處理（排序，Map-Reduce等）很泊，金融分析等需要大規(guī)模并行計算的領(lǐng)域角虫。

所以GPU也可以認(rèn)為是一種較通用的芯片。

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)秀的超級計算平臺。

CUDA軟件層次結(jié)構(gòu)

CUDA在軟件方面組成有：一個CUDA庫死讹、一個應(yīng)用程序編程接口（API）及其運行庫(Runtime)携狭、兩個較高級別的通用數(shù)學(xué)庫，即CUFFT和CUBLAS回俐。CUDA改進(jìn)了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í)行粱年，而當(dāng)遇到數(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)核之前，必須對核進(jìn)行執(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個線程胳岂。一般GPGPU程序線程數(shù)目是很多的，所以不能把所有的線程都塞到同一個塊里舔稀。但一個內(nèi)核可由多個大小相同的線程塊同時執(zhí)行乳丰，因而線程總數(shù)應(yīng)等于每個塊的線程數(shù)乘以塊的數(shù)量。這些同樣維度和大小的塊將組織為一個一維或二維線程塊網(wǎng)格(Grid)内贮。具體框架如下圖所示产园。

線程塊網(wǎng)格

NOTICE：

線程(Thread)
一般通過GPU的一個核進(jìn)行處理。（可以表示成一維夜郁，二維什燕，三維，具體下面再細(xì)說）竞端。
線程塊(Block)

由多個線程組成（可以表示成一維屎即，二維，三維婶熬，具體下面再細(xì)說）剑勾。

各block是并行執(zhí)行的埃撵，block間無法通信赵颅，也沒有執(zhí)行順序。

注意線程塊的數(shù)量限制為不超過65535（硬件限制）暂刘。

線程格(Grid)
由多個線程塊組成（可以表示成一維饺谬，二維，三維谣拣，具體下面再細(xì)說）募寨。
線程束
在CUDA架構(gòu)中，線程束是指一個包含32個線程的集合森缠，這個線程集合被“編織在一起”并且“步調(diào)一致”的形式執(zhí)行拔鹰。在程序中的每一行，線程束中的每個線程都將在不同數(shù)據(jù)上執(zhí)行相同的命令贵涵。

從硬件上看

SP：最基本的處理單元列肢，streaming processor恰画，也稱為CUDA core。最后具體的指令和任務(wù)都是在SP上處理的瓷马。GPU進(jìn)行并行計算拴还，也就是很多個SP同時做處理。
SM：多個SP加上其他的一些資源組成一個streaming multiprocessor欧聘。也叫GPU大核片林，其他資源如：warp scheduler，register怀骤，shared memory等费封。SM可以看做GPU的心臟（對比CPU核心），register和shared memory是SM的稀缺資源蒋伦。CUDA將這些資源分配給所有駐留在SM中的threads孝偎。因此，這些有限的資源就使每個SM中active warps有非常嚴(yán)格的限制凉敲，也就限制了并行能力衣盾。

從軟件上看

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è)備擁有多個獨立的存儲空間，其中包括：全局存儲器独柑、本地存儲器迈窟、共享存儲器、常量存儲器忌栅、紋理存儲器和寄存器车酣，如圖

CUDA設(shè)備上的存儲器

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ù)保持很短的時間娘摔。為了保持?jǐn)?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ù)的快速存取脾歇。

存儲器的應(yīng)用層次

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è)備均維護(hù)自己的DRAM膘魄，分別稱為主機存儲器和設(shè)備存儲器乌逐。因而，一個程序通過調(diào)用CUDA運行庫來管理對內(nèi)核可見的全局创葡、固定和紋理存儲器空間浙踢。這種管理包括設(shè)備存儲器的分配和取消分配，還包括主機和設(shè)備存儲器之間的數(shù)據(jù)傳輸蹈丸。

4. CUDA軟硬件

4.1 CUDA術(shù)語

由于CUDA中存在許多概念和術(shù)語成黄，諸如SM、block逻杖、SP等多個概念不容易理解，將其與CPU的一些概念進(jìn)行比較思瘟，如下表所示荸百。

CPU	GPU	層次
算術(shù)邏輯和控制單元	流處理器(SM)	硬件
算術(shù)單元	批量處理器(SP)	硬件
進(jìn)程	Block	軟件
線程	thread	軟件
調(diào)度單位	Warp	軟件

4.2 硬件利用率

當(dāng)為一個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	100
···	···

查看顯卡利用率 （以實驗室服務(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名士兵一個班（一個線程束），連隊又進(jìn)一步分成若干個班司顿，每個班由一個班長來指揮芒粹。

基于GPU的線程視圖

要執(zhí)行某個操作，總司令（內(nèi)核程序/ 主機程序）必須提供操作名稱及相應(yīng)的數(shù)據(jù)大溜。每個士兵（線程）只處理分配給他的問題中的一小塊化漆。在連長（負(fù)責(zé)一個塊）或班長（負(fù)責(zé)一個束）的控制下，束與束之間的線程或者一個束內(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ù)量不能太多的一重要因素≡呶啵“緩存一致”系統(tǒng)中最壞的情況是滤蝠，一個內(nèi)存操作會強迫每個核的緩存都進(jìn)行更新，進(jìn)而每個核都要對相鄰的內(nèi)存單元寫操作授嘀。

相比之下物咳，非“緩存一致”系統(tǒng)不會自動地更新其他核的緩存。它需要由程序員寫清楚每個處理器核輸出的各自不同的目標(biāo)區(qū)域蹄皱。從程序的視角看览闰，這支持一個核僅負(fù)責(zé)一個輸出或者一個小的輸出集。通常巷折，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ù)離散傅氏變換的奇、偶波附、虛艺晴、實等特性，對離散傅立葉變換的算法進(jìn)行改進(jìn)獲得的掸屡。

由于不是剛需封寞，這里不展開講。好奇的你可以點擊樓下時光機仅财，通過下面的教程進(jìn)行學(xué)習(xí)狈究。
FFT（最詳細(xì)最通俗的入門手冊）

5.7 CUDA計算能力的含義

體現(xiàn)GPU計算能力的兩個重要特征：
1)CUDA核的個數(shù)；
2)存儲器大小盏求。
描述GPU性能的兩個重要指標(biāo): ：
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 準(zhǔn)備工作

下面磅废，我們一步一步來驗證自己的系統(tǒng)是否滿足安裝要求。
Step 1: 驗證計算機是否擁有至少一顆支持 CUDA 的 GPU
打開終端（Ctrl + Alt + T）荆烈，鍵入以下命令：

lspci | grep -i nvidia

可以看到以下內(nèi)容（結(jié)果因人而異拯勉，與具體的GPU有關(guān)）

看到這個就說明至少有一顆支持 CUDA 的 GPU，可以進(jìn)入下一步了憔购。

Step 2: 驗證一下自己操作系統(tǒng)的版本
鍵入命令：

lsb_release -a

No LSB modules are available.
Distributor ID: Ubuntu
Description:    Ubuntu 16.04.4 LTS
Release:    16.04
Codename:   xenial

更多信息請移步Ubuntu查看版本信息

Step 3: 驗證 gcc 編譯器的版本
鍵入命令：

gcc --version

或者

gcc -v

得到如下信息

gcc (Ubuntu 5.4.0-6ubuntu1~16.04.10) 5.4.0 20160609
Copyright (C) 2015 Free Software Foundation, Inc.
This is free software; see the source for copying conditions.  There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

Step 4: 驗證系統(tǒng)內(nèi)核版本
鍵入命令：

uname -r

得到如下信息

對照官方提供的對各種 Linux 發(fā)行版的安裝要求進(jìn)行安裝

6.1.3 搭建 CUDA 環(huán)境

Step 1: 安裝 CUDA 工具包
在前面幾項驗證都順利通過以后就來到最關(guān)鍵的一步宫峦。首先下載對應(yīng)自己系統(tǒng)版本的 CUDA 工具包(以CUDA Toolkit 9.2 為例)，然后進(jìn)入到安裝包所在目錄：

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. 一般通過標(biāo)識符__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ā)生錯誤顾患，例如越界或報錯，甚至導(dǎo)致藍(lán)屏和死機个唧。

看完基本知識江解，裝好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定義的矢量類型，相當(dāng)亍由3個unsigned int型組成的結(jié)構(gòu)體稀蟋。uint3類型有三個數(shù)據(jù)成員unsigned int x; unsigned int y; unsigned int z;
可使用于一維煌张、二維或三維的索引來標(biāo)識線程，構(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。
1. 對于大小為（blockDim.x, blockDim.y, blockDim.z）的三維 block胃夏，線程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y轴或。
2. 對于計算線程索引偏移增量為已啟動線程的總數(shù)。如stride = blockDim.x * gridDim.x; threadId += stride仰禀。

6.2.3 函數(shù)修飾符

1.__global__照雁，表明被修飾的函數(shù)在設(shè)備上執(zhí)行，但在主機上調(diào)用答恶。

__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()分配的指針進(jìn)行設(shè)備內(nèi)存讀寫操作；
3.3. 可以將cudaMalloc()分配的指針傳遞給在主機上執(zhí)行的函數(shù)；
3.4. 不可以在主機代碼中使用cudaMalloc()分配的指針進(jìn)行主機內(nèi)存讀寫操作（即不能進(jìn)行解引用）悟狱。

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í)行得问，即當(dāng)函數(shù)返回時，復(fù)制操作就已經(jīng)完成了软免，并且在輸出緩沖區(qū)中包含了復(fù)制進(jìn)去的內(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上啟動的每個線程塊，CUDA C編譯器都將創(chuàng)建該共享變量的一個副本艘希。線程塊中的每個線程都共享這塊內(nèi)存硼身，但線程卻無法看到也不能修改其他線程塊的變量副本。這樣使得一個線程塊中的多個線程能夠在計算上通信和協(xié)作覆享。

常量內(nèi)存
1. 位置：設(shè)備內(nèi)存
2. 形式：關(guān)鍵字__constant__添加到變量聲明中佳遂。如__constant__ float s[10];。
3. 目的：為了提升性能撒顿。常量內(nèi)存采取了不同于標(biāo)準(zhǔn)全局內(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. 要求：當(dāng)我們需要拷貝數(shù)據(jù)到常量內(nèi)存中應(yīng)該使用cudaMemcpyToSymbol()唧席，而cudaMemcpy()會復(fù)制到全局內(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和主機之間復(fù)制數(shù)據(jù)兔辅。由于DMA在執(zhí)行復(fù)制時無需CPU介入腊敲。因此DMA復(fù)制過程中使用固定內(nèi)存是非常重要的。
4. 缺點：使用固定內(nèi)存维苔，將失去虛擬內(nèi)存的所有功能碰辅；系統(tǒng)將更快的耗盡內(nèi)存。
5. 建議：對cudaMemcpy()函數(shù)調(diào)用中的源內(nèi)存或者目標(biāo)內(nèi)存介时，才使用固定內(nèi)存没宾，并且在不再需要使用它們時立即釋放忍法。
6. 形式：通過cudaHostAlloc()函數(shù)來分配；通過cudaFreeHost()釋放榕吼。
7. 只能以異步方式對固定內(nèi)存進(jìn)行復(fù)制操作饿序。

原子性
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 something
cudaEventRecord( 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)存復(fù)制以及事件的啟動和結(jié)束等堂氯。這些操作的添加到流的順序也是它們的執(zhí)行順序蔑担。可以將每個流視為GPU上的一個任務(wù)咽白，并且這些任務(wù)可以并行執(zhí)行啤握。
硬件前提：必須是支持設(shè)備重疊功能的GPU。支持設(shè)備重疊功能晶框，即在執(zhí)行一個核函數(shù)的同時排抬，還能在設(shè)備與主機之間執(zhí)行復(fù)制操作懂从。
聲明與創(chuàng)建：聲明cudaStream_t stream;，創(chuàng)建cudaSteamCreate(&stream);蹲蒲。
cudaMemcpyAsync()：前面在cudaMemcpy()中提到過番甩，這是一個以異步方式執(zhí)行的函數(shù)。在調(diào)用cudaMemcpyAsync()時届搁，只是放置一個請求缘薛，表示在流中執(zhí)行一次內(nèi)存復(fù)制操作，這個流是通過參數(shù)stream來指定的卡睦。當(dāng)函數(shù)返回時宴胧，我們無法確保復(fù)制操作是否已經(jīng)啟動，更無法保證它是否已經(jīng)結(jié)束表锻。我們能夠得到的保證是恕齐，復(fù)制操作肯定會當(dāng)下一個被放入流中的操作之前執(zhí)行。傳遞給此函數(shù)的主機內(nèi)存指針必須是通過cudaHostAlloc()分配好的內(nèi)存瞬逊。（流中要求固定內(nèi)存）
流同步：通過cudaStreamSynchronize()來協(xié)調(diào)显歧。
流銷毀：在退出應(yīng)用程序之前，需要銷毀對GPU操作進(jìn)行排隊的流码耐，調(diào)用cudaStreamDestroy()追迟。
針對多個流：
9.1. 記得對流進(jìn)行同步操作。
9.2. 將操作放入流的隊列時骚腥，應(yīng)采用寬度優(yōu)先方式，而非深度優(yōu)先的方式瓶逃，換句話說束铭，不是首先添加第0個流的所有操作，再依次添加后面的第1厢绝，2,…個流契沫。而是交替進(jìn)行添加，比如將a的復(fù)制操作添加到第0個流中昔汉，接著把a的復(fù)制操作添加到第1個流中懈万，再繼續(xù)其他的類似交替添加的行為。
9.3. 要牢牢記住操作放入流中的隊列中的順序影響到CUDA驅(qū)動程序調(diào)度這些操作和流以及執(zhí)行的方式靶病。

TIPS:

當(dāng)線程塊的數(shù)量為GPU中處理數(shù)量的2倍時会通，將達(dá)到最優(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ù)量大的情況下效率還是相當(dāng)不錯的。

7. GPU or FPGA

GPU優(yōu)勢
1.從峰值性能來說郊闯，GPU（10Tflops)遠(yuǎn)遠(yuǎn)高于FPGA(<1TFlops);

2.GPU相對于FPGA還有一個優(yōu)勢就是內(nèi)存接口, GPU的內(nèi)存接口（傳統(tǒng)的GDDR5妻献，最近更是用上了HBM和HBM2）的帶寬遠(yuǎn)好于FPGA的傳統(tǒng)DDR接口（大約帶寬高4-5倍）;

3.功耗方面，雖然GPU的功耗遠(yuǎn)大于FPGA的功耗虚婿，但是如果要比較功耗應(yīng)該比較在執(zhí)行效率相同時需要的功耗旋奢。如果FPGA的架構(gòu)優(yōu)化能做到很好以致于一塊FPGA的平均性能能夠接近一塊GPU，那么FPGA方案的總功耗遠(yuǎn)小于GPU然痊，散熱問題可以大大減輕至朗。反之，如果需要二十塊FPGA才能實現(xiàn)一塊GPU的平均性能剧浸，那么FPGA在功耗方面并沒有優(yōu)勢锹引。

4.FPGA缺點有三點:
第一，基本單元的計算能力有限唆香。為了實現(xiàn)可重構(gòu)特性嫌变，F(xiàn)PGA 內(nèi)部有大量極細(xì)粒度的基本單元，但是每個單元的計算能力（主要依靠LUT 查找表）都遠(yuǎn)遠(yuǎn)低于CPU 和GPU 中的ALU模塊躬它。
第二腾啥，速度和功耗相對專用定制芯片（ASIC）仍然存在不小差距。
第三冯吓，F(xiàn)PGA 價格較為昂貴倘待，在規(guī)模放量的情況下單塊FPGA 的成本要遠(yuǎn)高于專用定制芯片。最后誰能勝出, 完全取決于FPGA架構(gòu)優(yōu)化能否彌補峰值性能的劣勢组贺。

5.個人更推薦: CPU+FPGA的組合模式; 其中FPGA用于整形計算凸舵，cpu進(jìn)行浮點計算和調(diào)度，此組合的擁有更高的單位功耗性能和更低的時延失尖。最后更想GPU穩(wěn)定開放,發(fā)揮其長處, 達(dá)到真正的物美價廉!

FPGA優(yōu)勢
人工智能目前仍處于早期階段啊奄，未來人工智能的主戰(zhàn)場是在推理環(huán)節(jié)，遠(yuǎn)沒有爆發(fā)掀潮。未來勝負(fù)尚未可知菇夸，各家技術(shù)路線都有機會勝出。目前英偉達(dá)的GPU在訓(xùn)練場景中占據(jù)著絕對領(lǐng)導(dǎo)地位胧辽，但是在未來峻仇，專注于推理環(huán)節(jié)的FPGA必將會發(fā)揮巨大的價值。

FPGA和GPU內(nèi)都有大量的計算單元邑商，因此它們的計算能力都很強摄咆。在進(jìn)行神經(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)整硬件資源。
目前機器學(xué)習(xí)大多使用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方案的總功耗遠(yuǎn)小于GPU蚂维，散熱問題可以大大減輕。能效比的比較也是類似路狮，能效指的是完成程序執(zhí)行消耗的能量虫啥，而能量消耗等于功耗乘以程序的執(zhí)行時間。雖然GPU的功耗遠(yuǎn)大于FPGA的功耗奄妨，但是如果FPGA執(zhí)行相同程序需要的時間比GPU長幾十倍涂籽，那FPGA在能效比上就沒有優(yōu)勢了；反之如果FPGA上實現(xiàn)的硬件架構(gòu)優(yōu)化得很適合特定的機器學(xué)習(xí)應(yīng)用砸抛，執(zhí)行算法所需的時間僅僅是GPU的幾倍或甚至于接近GPU评雌，那么FPGA的能效比就會比GPU強树枫。

在峰值性能比方面，雖然GPU的峰值性能（10Tflops）遠(yuǎn)大于FPGA的峰值性能（<1Tflops）景东，但針對特定的場景來講吞吐量并不比GPU差砂轻。

8. 深度學(xué)習(xí)的三種硬件方案：ASIC，F(xiàn)PGA斤吐，GPU

8.1 對深度學(xué)習(xí)硬件平臺的要求

要想明白“深度學(xué)習(xí)”需要怎樣的硬件搔涝，必須了解深度學(xué)習(xí)的工作原理。首先在表層上和措，我們有一個巨大的數(shù)據(jù)集庄呈，并選定了一種深度學(xué)習(xí)模型。每個模型都有一些內(nèi)部參數(shù)需要調(diào)整派阱，以便學(xué)習(xí)數(shù)據(jù)诬留。而這種參數(shù)調(diào)整實際上可以歸結(jié)為優(yōu)化問題，在調(diào)整這些參數(shù)時颁褂，就相當(dāng)于在優(yōu)化特定的約束條件故响。

矩陣相乘（Matrix Multiplication）——幾乎所有的深度學(xué)習(xí)模型都包含這一運算，它的計算十分密集颁独。
卷積（Convolution）——這是另一個常用的運算彩届，占用了模型中大部分的每秒浮點運算（浮點／秒）。
循環(huán)層（Recurrent Layers ）——模型中的反饋層誓酒，并且基本上是前兩個運算的組合樟蠕。
All Reduce——這是一個在優(yōu)化前對學(xué)習(xí)到的參數(shù)進(jìn)行傳遞或解析的運算序列。在跨硬件分布的深度學(xué)習(xí)網(wǎng)絡(luò)上執(zhí)行同步優(yōu)化時（如AlphaGo的例子）靠柑，這一操作尤其有效寨辩。

除此之外，深度學(xué)習(xí)的硬件加速器需要具備數(shù)據(jù)級別和流程化的并行性歼冰、多線程和高內(nèi)存帶寬等特性靡狞。另外，由于數(shù)據(jù)的訓(xùn)練時間很長隔嫡，所以硬件架構(gòu)必須低功耗甸怕。因此，效能功耗比（Performance per Watt）是硬件架構(gòu)的評估標(biāo)準(zhǔn)之一腮恩。

CNN在應(yīng)用中梢杭，一般采用GPU加速，請解釋為什么GPU可以有加速效果秸滴，主要加速算法的哪一個部分武契？

這里默認(rèn)gpu加速是指NVIDIA的CUDA加速。CPU是中央處理單元，gpu是圖形處理單元咒唆，gpu由上千個流處理器(core)作為運算器届垫。執(zhí)行采用單指令多線程(SIMT)模式。相比于單核CPU（向量機）流水線式的串行操作钧排，雖然gpu單個core計算能力很弱敦腔，但是通過大量線程進(jìn)行同時計算，在數(shù)據(jù)量很大是會活動較為可觀的加速效果恨溜。

具體到cnn符衔，利用gpu加速主要是在conv（卷積）過程上。conv過程同理可以像以上的向量加法一樣通過cuda實現(xiàn)并行化糟袁。具體的方法很多判族，不過最好的還是利用fft（快速傅里葉變換）進(jìn)行快速卷積。NVIDIA提供了cufft庫實現(xiàn)fft项戴，復(fù)數(shù)乘法則可以使用cublas庫里的對應(yīng)的level3的cublasCgemm函數(shù)形帮。

GPU加速的基本準(zhǔn)則就是“人多力量大”。CNN說到底主要問題就是計算量大周叮，但是卻可以比較有效的拆分成并行問題辩撑。隨便拿一個層的filter來舉例子，假設(shè)某一層有n個filter仿耽，每一個需要對上一層輸入過來的map進(jìn)行卷積操作合冀。那么，這個卷積操作并不需要按照線性的流程去做项贺，每個濾波器互相之間并不影響君躺，可以大家同時做，然后大家生成了n張新的譜之后再繼續(xù)接下來的操作开缎。既然可以并行棕叫，那么同一時間處理單元越多，理論上速度優(yōu)勢就會越大奕删。所以俺泣，處理問題就變得很簡單粗暴，就像NV那樣完残，暴力增加顯卡單元數(shù)（當(dāng)然砌滞，顯卡的架構(gòu)、內(nèi)部數(shù)據(jù)的傳輸速率坏怪、算法的優(yōu)化等等也都很重要）。

GPU主要是針對圖形顯示及渲染等技術(shù)的出眾绊茧，而其中的根本是因為處理矩陣算法能力的強大铝宵，剛好CNN中涉及大量的卷積，也就是矩陣乘法等，所以在這方面具有優(yōu)勢鹏秋。

機器學(xué)習(xí)的算法一定得經(jīng)過gpu加速嗎尊蚁？

不一定。只有需要大量浮點數(shù)計算侣夷，例如矩陣乘法横朋，才需要GPU加速。用CNN對圖像進(jìn)行分類就是一個需要大量浮點數(shù)計算的典型案例百拓，通常需要GPU加速

對于ASIC琴锭、FPGA、分布式計算衙传，這里不再展開講决帖，有興趣的小伙伴可以，自行學(xué)習(xí)蓖捶。不過....說不定某天博主心情好地回，就會梳理一下這幾種硬件方案在端到端上應(yīng)用的區(qū)別了。

菜鳥入門教程就到這里了俊鱼，聰明的你一定不滿足這個入門教程刻像，如有興趣進(jìn)一步學(xué)習(xí)CUDA編程，可移步NVIDIA官方的課程平臺CUDA ZONE（PS：中文網(wǎng)站并闲，英文課程）

歡迎交流 ?????
Author：He_Yu
Email：heyu.nwpu@gmail.com

最后編輯于：2018.08.24 15:11:46

?著作權(quán)歸作者所有,轉(zhuǎn)載或內(nèi)容合作請聯(lián)系作者

人面猴
序言：七十年代末细睡，一起剝皮案震驚了整個濱河市，隨后出現(xiàn)的幾起案子焙蚓，更是在濱河造成了極大的恐慌纹冤，老刑警劉巖，帶你破解...
沈念sama閱讀 206,214評論 6贊 481
死咒
序言：濱河連續(xù)發(fā)生了三起死亡事件购公，死亡現(xiàn)場離奇詭異萌京，居然都是意外死亡，警方通過查閱死者的電腦和手機宏浩，發(fā)現(xiàn)死者居然都...
沈念sama閱讀 88,307評論 2贊 382
救了他兩次的神仙讓他今天三更去死
文/潘曉璐我一進(jìn)店門知残，熙熙樓的掌柜王于貴愁眉苦臉地迎上來，“玉大人比庄，你說我怎么就攤上這事求妹。” “怎么了佳窑？”我有些...
開封第一講書人閱讀 152,543評論 0贊 341
道士緝兇錄：失蹤的賣姜人
文/不壞的土叔我叫張陵制恍，是天一觀的道長。經(jīng)常有香客問我神凑，道長净神，這世上最難降的妖魔是什么何吝？我笑而不...
開封第一講書人閱讀 55,221評論 1贊 279
?港島之戀（遺憾婚禮）
正文為了忘掉前任，我火速辦了婚禮鹃唯，結(jié)果婚禮上爱榕，老公的妹妹穿的比我還像新娘。我一直安慰自己坡慌，他們只是感情好黔酥，可當(dāng)我...
茶點故事閱讀 64,224評論 5贊 371
惡毒庶女頂嫁案：這布局不是一般人想出來的
文/花漫我一把揭開白布。她就那樣靜靜地躺著洪橘，像睡著了一般跪者。火紅的嫁衣襯著肌膚如雪。梳的紋絲不亂的頭發(fā)上梨树，一...
開封第一講書人閱讀 49,007評論 1贊 284
城市分裂傳說
那天坑夯，我揣著相機與錄音，去河邊找鬼抡四。笑死柜蜈，一個胖子當(dāng)著我的面吹牛，可吹牛的內(nèi)容都是我干的指巡。我是一名探鬼主播淑履，決...
沈念sama閱讀 38,313評論 3贊 399
雙鴛鴦連環(huán)套：你想象不到人心有多黑
文/蒼蘭香墨我猛地睜開眼，長吁一口氣：“原來是場噩夢啊……” “哼藻雪！你這毒婦竟也來了秘噪？” 一聲冷哼從身側(cè)響起，我...
開封第一講書人閱讀 36,956評論 0贊 259
萬榮殺人案實錄
序言：老撾萬榮一對情侶失蹤勉耀，失蹤者是張志新（化名）和其女友劉穎指煎，沒想到半個月后，有當(dāng)?shù)厝嗽跇淞掷锇l(fā)現(xiàn)了一具尸體便斥，經(jīng)...
沈念sama閱讀 43,441評論 1贊 300
?護(hù)林員之死
正文獨居荒郊野嶺守林人離奇死亡至壤，尸身上長有42處帶血的膿包…… 初始之章·張勛以下內(nèi)容為張勛視角年9月15日...
茶點故事閱讀 35,925評論 2贊 323
?白月光啟示錄
正文我和宋清朗相戀三年，在試婚紗的時候發(fā)現(xiàn)自己被綠了枢纠。大學(xué)時的朋友給我發(fā)了我未婚夫和他白月光在一起吃飯的照片像街。...
茶點故事閱讀 38,018評論 1贊 333
活死人
序言：一個原本活蹦亂跳的男人離奇死亡，死狀恐怖晋渺，靈堂內(nèi)的尸體忽然破棺而出镰绎，到底是詐尸還是另有隱情，我是刑警寧澤木西，帶...
沈念sama閱讀 33,685評論 4贊 322
?日本核電站爆炸內(nèi)幕
正文年R本政府宣布畴栖，位于F島的核電站，受9級特大地震影響八千，放射性物質(zhì)發(fā)生泄漏驶臊。R本人自食惡果不足惜挪挤，卻給世界環(huán)境...
茶點故事閱讀 39,234評論 3贊 307
男人毒藥：我在死后第九天來索命
文/蒙蒙一、第九天我趴在偏房一處隱蔽的房頂上張望关翎。院中可真熱鬧，春花似錦鸠信、人聲如沸纵寝。這莊子的主人今日做“春日...
開封第一講書人閱讀 30,240評論 0贊 19
一樁弒父案星立，背后竟有這般陰謀
文/蒼蘭香墨我抬頭看了看天上的太陽爽茴。三九已至，卻和暖如春绰垂，著一層夾襖步出監(jiān)牢的瞬間室奏，已是汗流浹背。一陣腳步聲響...
開封第一講書人閱讀 31,464評論 1贊 261
情欲美人皮
我被黑心中介騙來泰國打工劲装，沒想到剛下飛機就差點兒被人妖公主榨干…… 1. 我叫王不留胧沫，地道東北人。一個月前我還...
沈念sama閱讀 45,467評論 2贊 352
代替公主和親
正文我出身青樓占业，卻偏偏與公主長得像绒怨，于是被迫代替她去往敵國和親。傳聞我的和親對象是個殘疾皇子谦疾，可洞房花燭夜當(dāng)晚...
茶點故事閱讀 42,762評論 2贊 345

CUDA——"從入門到放棄"

1. 知識準(zhǔn)備

1.1 中央處理器（CPU）

1.2 顯卡

1.3 內(nèi)存

1.4 顯存

1.5 顯卡淮野、顯卡驅(qū)動捧书、CUDA之間的關(guān)系

再啰嗦兩句

2. CUDA軟件構(gòu)架

3. 編程模型

3.1 線程層次結(jié)構(gòu)

3.2 存儲器層次結(jié)構(gòu)

3.3 主機（Host）和設(shè)備（Device）

4. CUDA軟硬件

4.1 CUDA術(shù)語

4.2 硬件利用率

5. 并行計算

5.1 并發(fā)性

5.2 局部性

5.3 緩存一致性

5.4 弗林分類法

5.5 分條 / 分塊

5.6 快速傅氏變換(FFT)

5.7 CUDA計算能力的含義

6. 實踐

6.1 Ubuntu 系統(tǒng)下環(huán)境搭建

6.1.1 系統(tǒng)要求

6.1.2 準(zhǔn)備工作

6.1.3 搭建 CUDA 環(huán)境

6.2 CUDA編程

6.2.1 核函數(shù)

6.2.2 dim3結(jié)構(gòu)類型

6.2.3 函數(shù)修飾符

6.2.4 常用的GPU內(nèi)存函數(shù)

6.2.5 GPU內(nèi)存分類

6.2.6 常用線程操作函數(shù)

6.2.7 流

6.2.8 這是一個栗子

7. GPU or FPGA

8. 深度學(xué)習(xí)的三種硬件方案：ASIC，F(xiàn)PGA斤吐，GPU

8.1 對深度學(xué)習(xí)硬件平臺的要求

推薦閱讀更多精彩內(nèi)容