cuda初步认识
特此聲明：這個內容我是轉別人的
我只摘錄一些我需要的東西，若是想看原文的，請點擊下面的鏈接
原文：http://hi.baidu.com/coolrainbow/item/de05efc83151671a50505878
?
?
?
1 硬件架構
CUDA編程中，習慣稱CPU為Host，GPU為Device。
?
?
2 并行模型
Thread：并行基本單位
Block：相互合作的一組線程。可以彼此同步，快速交換數據，最多可以512個線程
Grid：一組Block，有共享全局內存
Kernel：在GPU上執行的程序，一個Kernel對應一個Grid
Block和Thread都有各自的ID，記作blockIdx（1D，2D），threadIdx（1D，2D，3D）
Block和Thread還有Dim，即blockDim與threadDim. 他們都有三個分量x，y，z
線程同步：void __syncthreads(); 可以同步一個Block內的所有線程
?
?
3 存儲層次
per-thread register 1 cycle
per-thread local memory slow
per-block shared memory 1 cycle
per-grid global memory 500 cycle,not cached!!
constant and texture memories 500 cycle, but cached and read-only
分配內存：cudaMalloc，cudaFree，它們分配的是global memory
Hose-Device數據交換：cudaMemcpy
?
4 變量類型
__device__：GPU的global memory空間，grid中所有線程可訪問
__constant__：GPU的constant memory空間，grid中所有線程可訪問
__shared__：GPU上的thread block空間，block中所有線程可訪問
local：位于SM內，僅本thread可訪問
在編程中，可以在變量名前面加上這些前綴以區分。
?
?
5 數據類型
內建矢量：int1，int2，int3，int4，float1，float2， float3，float4 ...
紋理類型：texture<Type, Dim, ReadMode>texRef;
內建dim3類型：定義grid和block的組織方法。
例如：
dim3 dimGrid(2, 2);
dim3 dimBlock(4, 2, 2);
kernelFoo<<<dimGrid, dimBlock>>>(argument);
?
?
6 CUDA函數定義
__device__：執行于Device，僅能從Device調用。
限制：不能用&取地址；不支持遞歸；不支持static variable；不支持可變長度參數
__global__ void： 執行于Device，僅能從Host調用。此類函數必須返回void
__host__：執行于Host，僅能從Host調用
在執行kernel函數時，必須提供execution configuration，即<<<....>>>的部分。
例如：
__global__ void KernelFunc(...);
dim3 DimGrid(100, 50); // 5000 thread blocks
dim3 DimBlock(4, 8, 8); // 256 threads per block
size_t SharedMemBytes = 64; // 64 bytes of shared memory
KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>(...);
?
?
7 CUDA包含一些數學函數，如sin，pow等。每一個函數包含有兩個版本，例如正弦函數sin，一個普通版本sin，另一個不精確但速度極快的__sin版本。
?
?
8 內置變量
gridDim, blockIdx, blockDim, threadIdx, wrapsize. 這些內置變量不允許賦值的
?
?
9 編寫程序
目前CUDA僅能良好的支持C，在編寫含有CUDA代碼的程序時，首先要導入頭文件cuda_runtime_api.h。文件名后綴為.cu，使用nvcc編譯器編譯。本來想在這里給出些源碼的，但是源碼教程，以后單獨開一個文章在說吧。
?
?
?
?
這部分是一些枯燥的硬件知識的總結，但是對優化CUDA程序有著至關重要的作用，在后面的文章里，我將盡量結合實例來講解這些東西
?
1 GPU硬件
i GPU一個最小單元稱為Streaming Processor(SP)，全流水線單事件無序微處理器，包含兩個ALU和一個FPU，多組寄存器文件（register file，很多寄存器的組合），這個SP沒有cache。事實上，現代GPU就是一組SP的array，即SPA。每一個SP執行一個thread
?
ii 多個SP組成Streaming Multiprocessor(SM)。每一個SM執行一個block。每個SM包含
8個SP；
2個special function unit(SFU)：這里面有4個FPU可以進行超越函數和插值計算
MultiThreading Issue Unit：分發線程指令
具有指令和常量緩存。
包含shared memory
?
?
iii Texture Processor Cluster(TPC) ：包含某些其他單元的一組SM
?
?
2 Single-Program Multiple-Data （SPMD）模型
?
i CPU以順序結構執行代碼，GPU以threads blocks組織并發執行的代碼，即無數個threads同時執行
ii 回顧一下CUDA的概念：
一個kernel程序執行在一個grid of threads blocks之中
一個threads block是一批相互合作的threads：可以用過__syncthreads同步；通過shared memory共享變量，不同block的不能同步。
iii Threads block聲明：
可以包含有1到512個并發線程，具有唯一的blockID，可以是1,2,3D
同一個block中的線程執行同一個程序，不同的操作數，可以同步，每個線程具有唯一的ID
?
3 線程硬件原理
i GPU通過Global block scheduler來調度block，根據硬件架構分配block到某一個SM。每個SM最多分配8個block，每個SM最多可接受768個thread（可以是一個block包含512個thread，也可以是3個block每個包含256個thread（3*256=768！））。同一個SM上面的block的尺寸必須相同。每個線程的調度與ID由該SM管理。
ii SM滿負載工作效率最高！考慮某個Block，其尺寸可以為8*8,16*16,32*32
8*8：每個block有64個線程，由于每個SM最多處理768個線程，因此需要768/64=12個block。但是由于SM最多8個block，因此一個SM實際執行的線程為8*64=512個線程。
16*16：每個block有256個線程，SM可以同時接受三個block，3*256=768，滿負載：）
32*32：每個block有1024個線程，SM無法處理！
iii Block是獨立執行的，每個Block內的threads是可協同的。
iv 每個線程由SM中的一個SP執行。當然，由于SM中僅有8個SP，768個線程是以warp為單位執行的，每個warp包含32個線程，這是基于線程指令的流水線特性完成的。Warp是SM基本調度單位，實際上，一個Warp是一個32路SIMD指令。基本單位是half-warp。
如，SM滿負載工作有768個線程，則共有768/32=24個warp，每一瞬時，只有一組warp在SM中執行。
Warp全部線程是執行同一個指令，每個指令需要4個clock cycle，通過復雜的機制執行。
v 一個thread的一生：
Grid在GPU上啟動；block被分配到SM上；SM把線程組織為warp；SM調度執行warp；執行結束后釋放資源；block繼續被分配....
4 線程存儲模型
i Register and local memory：線程私有，對程序員透明。
每個SM中有8192個register，分配給某些block，block內部的thread只能使用分配的寄存器。線程數多，每個線程使用的寄存器就少了。
ii shared memory：block內共享，動態分配。如__shared__ float region[N]。
shared memory 存儲器是被劃分為16個小單元，與half-warp長度相同，稱為bank，每個bank可以提供自己的地址服務。連續的32位word映射到連續的bank。
對同一bank的同時訪問稱為bank conflict。盡量減少這種情形。
?
iii Global memory：沒有緩存！容易稱為性能瓶頸，是優化的關鍵！
一個half-warp里面的16個線程對global memory的訪問可以被coalesce成整塊內存的訪問，如果：
數據長度為4,8或16bytes；地址連續；起始地址對齊；第N個線程訪問第N個數據。
Coalesce可以大大提升性能
?
uncoalesced
Coalesced方法：如果所有線程讀取同一地址，不妨使用constant memory；如果為不規則讀取可以使用texture內存
如果使用了某種結構體，其大小不是4 8 16的倍數，可以通過__align(X)強制對齊，X=4 8 16
總結
- 上一篇： CUDA从入门到精通（四）：加深对设备的
- 下一篇： 计算机用户组连接打印机,在组策略中使用脚