CUDA：Compute Unified Device Architecture，是由NVIDIA所推出的一種集成技術(shù)，允許使用標準C來進行GPU代碼編程，最終轉(zhuǎn)為PTX匯編代碼。
CPU與GPU
GPU可以看作是CPU的協(xié)助處理器，使用GPU實際指的是基于CPU+GPU的異構(gòu)計算架構(gòu)。通過PCle總線連接，CPU端成為Host端，GPU端稱為Device端。
GPU適合數(shù)據(jù)并行的計算密集型任務(wù)，如大型矩陣運算，而CPU的運算核心較少，但是其可以實現(xiàn)復雜的邏輯運算，因此其適合控制密集型任務(wù)。此外，CPU上的線程是重量級的，上下文切換開銷大，但是GPU由于存在很多核心，其線程是輕量級的。因此，基于CPU+GPU的異構(gòu)計算平臺可以優(yōu)勢互補，CPU負責處理邏輯復雜的串行程序，而GPU重點處理數(shù)據(jù)密集型的并行計算程序。
CUDA程序執(zhí)行流程
1、分配host內(nèi)存，并進行數(shù)據(jù)初始化；
2、分配device內(nèi)存，并從host將數(shù)據(jù)拷貝到device上；
3、調(diào)用CUDA kernel在device上完成指定的運算；（kernel是在device上線程中并行執(zhí)行的函數(shù)，核函數(shù)用__global__符號聲明，在調(diào)用時需要用<<<grid, block>>>來指定kernel要執(zhí)行的線程數(shù)量，在CUDA中，每一個線程都要執(zhí)行核函數(shù)，并且每個線程會分配一個唯一的線程號thread ID，這個ID值可以通過核函數(shù)的內(nèi)置變量threadIdx來獲得。）
4、將device上的運算結(jié)果拷貝到host上；
5、釋放device和host上分配的內(nèi)存。
CUDA函數(shù)類型限定詞
? ? ? ? 1）__global__：在device上執(zhí)行，從host中調(diào)用（一些特定的GPU也可以從device上調(diào)用），返回類型必須是void，不支持可變參數(shù)參數(shù)，不能成為類成員函數(shù)。注意用__global__定義的kernel是異步的，這意味著host不會等待kernel執(zhí)行完就執(zhí)行下一步。
? ? ? ? 2）__device__：在device上執(zhí)行，單僅可以從device中調(diào)用，不可以和__global__同時用。
? ? ? ? 3）__host__：在host上執(zhí)行，僅可以從host上調(diào)用，一般省略不寫，不可以和__global__同時用，但可和__device__，此時函數(shù)會在device和host都編譯。
Kernel線程結(jié)構(gòu)
kernel在device上執(zhí)行時實際上是啟動很多線程，一個kernel所啟動的所有線程稱為一個網(wǎng)格（grid），同一個網(wǎng)格上的線程共享相同的全局內(nèi)存空間，grid是線程結(jié)構(gòu)的第一層次，而網(wǎng)格又可以分為很多線程塊（block），一個線程塊里面包含很多線程，這是第二個層次。grid和block都是定義為dim3類型的變量，dim3可以看成是包含三個無符號整數(shù)（x，y，z）成員的結(jié)構(gòu)體變量，在定義時，缺省值初始化為1。因此grid和block可以靈活地定義為1-dim，2-dim以及3-dim結(jié)構(gòu)，對于圖中結(jié)構(gòu)（主要水平方向為x軸），定義的grid和block如下所示，kernel在調(diào)用時也必須通過執(zhí)行配置<<<grid, block>>>來指定kernel所使用的線程數(shù)及結(jié)構(gòu)。
一個線程需要兩個內(nèi)置的坐標變量（blockIdx，threadIdx）來唯一標識，它們都是dim3類型變量，其中blockIdx指明線程所在grid中的位置，而threaIdx指明線程所在block中的位置。以Thread(1, 1)為例：
? ? ? ? threadIdx.x = 1? ? ? ? ? ?threadIdx.y = 1
????????blockIdx.x = 1? ? ? ? ? ? ?blockIdx.y =1
測試GPU硬件配置
CUDA內(nèi)存模型
每個線程有自己的私有本地內(nèi)存（Local Memory），而每個線程塊有包含共享內(nèi)存（Shared Memory）,可以被線程塊中所有線程共享，其生命周期與線程塊一致。此外，所有的線程都可以訪問全局內(nèi)存（Global Memory）。還可以訪問一些只讀內(nèi)存塊：常量內(nèi)存（Constant Memory）和紋理內(nèi)存（Texture Memory）。
分配托管內(nèi)存?
實例：矩陣加法
利用上圖2-dim結(jié)構(gòu)實現(xiàn)兩個矩陣的加法，每個線程負責處理每個位置的兩個元素相加，代碼如下所示。線程塊大小為(16, 16)，然后將N*N大小的矩陣均分為不同的線程塊來執(zhí)行加法運算。
實例：矩陣乘法
設(shè)輸入矩陣為?
?和?
?，要得到?
?。實現(xiàn)思路是每個線程計算?
?的一個元素值?
?，對于矩陣運算，應(yīng)該選用grid和block為2-D的。首先定義矩陣的結(jié)構(gòu)體：
nvprof工具
通過命令nvprof cuda9.exe可以得到kernel的運行情況，
