Professional CUDA C Programing
代碼下載：http:www.wrox.com/go/procudac
本章的主要內(nèi)容：
?了解warp執(zhí)行的本質(zhì)
?將更多的并行性暴露給GPU
?掌握網(wǎng)格和塊配置的設(shè)置方法
?學(xué)習(xí)各種CUDA性能指標(biāo)和事件
?探測(cè)動(dòng)態(tài)并行和嵌套執(zhí)行
GPU的硬件結(jié)構(gòu)
GPU是由Streaming Multiprocessors (SM)組成的，每個(gè)SM如下：
? CUDA Cores
? Shared Memory/L1 Cache
? Register File
? Load/Store Units
? Special Function Units
? Warp Scheduler
GPU中的每個(gè)SM都支持?jǐn)?shù)百個(gè)線程的并發(fā)執(zhí)行，通常是每個(gè)GPU有多個(gè)SM，所以有可能有數(shù)千個(gè)線程并發(fā)執(zhí)行。
CUDA采用了SIMT單指令多線程執(zhí)行，一個(gè)指令32個(gè)線程執(zhí)行，32個(gè)線程組織成warp。一個(gè)warp中的線程同一時(shí)刻執(zhí)行同一個(gè)指令。每個(gè)線程有自己的指令技術(shù)計(jì)數(shù)器和寄存器，在自己的數(shù)據(jù)上執(zhí)行指令。
SIMT 和 SIMD最大的差異：
? 每個(gè)線程有自己獨(dú)立的指令寄存器
? 每個(gè)線程有自己獨(dú)立的寄存器狀態(tài)
? 每個(gè)線程有獨(dú)立的執(zhí)行路徑
一個(gè)線程塊只能分配到一個(gè)SM上執(zhí)行，一個(gè)SM可以同時(shí)允許多個(gè)線程塊。
共享存儲(chǔ)器和寄存器都是SM上珍貴的資源，共享存儲(chǔ)器按線程塊進(jìn)行劃分，同一個(gè)線程塊中的線程可以通過(guò)共享內(nèi)存互相通信，在邏輯上同一個(gè)線程塊中的所有線程同時(shí)執(zhí)行，但是在物理上，同一個(gè)線程塊中的所有線程并不是同時(shí)執(zhí)行的，所以同一個(gè)線程塊中的線程并不是同時(shí)執(zhí)行結(jié)束的。While all threads in a thread block run logically in parallel, not all threads can execute physically at the same time. As a result, different threads in a thread block may make progressat a different pace.
共享內(nèi)存可能會(huì)導(dǎo)致線程之間的競(jìng)爭(zhēng)：多個(gè)線程同時(shí)訪問(wèn)某個(gè)數(shù)據(jù)。CUDA提供了線程塊內(nèi)的同步，保證同一個(gè)線程塊中的線程在下一步執(zhí)行前都完成了上一步的執(zhí)行。但是線程塊之間無(wú)法同步。
在SM1中warp1正在執(zhí)行，但是warp1需要從device中讀取數(shù)據(jù)，此時(shí)SM1將調(diào)用warp2繼續(xù)執(zhí)行，warp1和warp2之間的轉(zhuǎn)換開(kāi)銷不大（SM的資源為所有線程共享），由于warp間并發(fā)的執(zhí)行提高了SM的利用率。（一個(gè)SM中真正執(zhí)行的warp數(shù)目和GPU的資源有關(guān)）
Fermi Architecture
Fermi有16個(gè)SM，每個(gè)SM有32個(gè)CUDA core（一個(gè)warp32個(gè)線程），每個(gè)CUDA core有ALU和FPU。當(dāng)一個(gè)線程塊分配到一個(gè)SM上時(shí)，線程塊被組織成warps，SM上的warp調(diào)度器選擇合適的warp執(zhí)行。
對(duì)于計(jì)算能力2.0以上的Fermi結(jié)構(gòu)，一個(gè)SM最多同時(shí)處理48個(gè)warps。
Fermi的兩個(gè)關(guān)鍵點(diǎn)：
? 可以通過(guò)CUDA runtime API 設(shè)置共享內(nèi)存和L1cache
? 支持并發(fā)的內(nèi)核執(zhí)行：多個(gè)小的kernel可以并發(fā)執(zhí)行，最多16個(gè)kernels同時(shí)在設(shè)備上運(yùn)行。
Kepler Architecture
? 15個(gè)SM
? 每個(gè)SM：192 單精度CUDA core，64個(gè)雙精度計(jì)算單元，32個(gè)特殊功能計(jì)算單元，32個(gè)load／store計(jì)算單元。4個(gè)warp調(diào)度器，8個(gè)指令分配器。
? 計(jì)算能力3.5每個(gè)SM一次可以調(diào)度64個(gè)warps駐留在SM上。
? 動(dòng)態(tài)并行性。一個(gè)kernel可以創(chuàng)建其它的kernel
? Hyper-Q。Hyper-Q在CPU和GPU之間增加了更多同步的硬件連接，從而實(shí)現(xiàn)了CPU核心同時(shí)在GPU上運(yùn)行更多任務(wù)。 因此，可以增加GPU 使用率。 費(fèi)米GPU依靠單一硬件工作隊(duì)列將任務(wù)從CPU傳遞到GPU，這可能導(dǎo)致單個(gè)任務(wù)阻止所有其他任務(wù)落后于隊(duì)伍中取得進(jìn)展。 開(kāi)普勒Hyper-Q消除了這個(gè)限制。Kepler GPU在主機(jī)和主機(jī)之間提供32個(gè)硬件工作隊(duì)列GPU。 Hyper-Q可以在GPU上實(shí)現(xiàn)更多的并發(fā)性，最大限度地提高GPU的利用率。。
性能優(yōu)化
? 時(shí)間復(fù)雜度、空間復(fù)雜度
? 特殊指令的使用
? 調(diào)用函數(shù)的頻率
性能優(yōu)化的必要性：
?簡(jiǎn)單的內(nèi)核實(shí)現(xiàn)通常不會(huì)產(chǎn)生最佳性能。 性能調(diào)優(yōu)工具可以幫助您查找代碼中的關(guān)鍵區(qū)域，這些區(qū)域是性能瓶頸。
?CUDA中的SM資源在多個(gè)駐留線程塊中分分配。此分配可能會(huì)導(dǎo)致一些資源成為性能限制。 Profiling工具可以幫助您深入了解如何利用計(jì)算資源。
?CUDA提供了硬件架構(gòu)的抽象，使您能夠控制線程并發(fā)性 。Profiling工具可以幫助您測(cè)量，可視化和指導(dǎo)您的優(yōu)化。
nvvp：可視化性能分析工具
nvprof：命令行性能那分析工具
**注意：＊＊
1.很多性能指標(biāo)都是針對(duì)的每個(gè)SM并不是整個(gè)GPU。
2.運(yùn)行一次可能只會(huì)得到某些參數(shù)，多次運(yùn)行可以收集完整。
3.多次運(yùn)行的結(jié)果可能會(huì)不同。
考慮的因素：
1.存儲(chǔ)器帶寬
2.計(jì)算資源
3.指令和存儲(chǔ)的時(shí)延
Warp的執(zhí)行方式
當(dāng)創(chuàng)建了一個(gè)kernel時(shí)，從邏輯上理解為kernel中的所有線程都在并行，但是從硬件物理?xiàng)l件上看同一時(shí)刻并不是所有的線程都在執(zhí)行。。
Warp和線程塊
warp是SM上的基本執(zhí)行單元。warp一定是同一個(gè)block中的，如果一個(gè)block中的threads不足32個(gè)，則補(bǔ)足成為32個(gè)構(gòu)成一個(gè)warp。
如圖所示，本來(lái)只需要80個(gè)線程，但是實(shí)際上仍然需要32＊3=96個(gè)threads，盡管最后一個(gè)warp的16個(gè)線程沒(méi)有使用，但是仍然會(huì)消耗SM上的資源，比如共享存儲(chǔ)器、寄存器。
Warp分支
定義：一個(gè)warp中的線程執(zhí)行不同的指令，叫做warp分支。
如果warp發(fā)生分支，則需要順序執(zhí)行每個(gè)分支路徑。
在一個(gè)warp中所有線程都必須具有兩個(gè)分支if...else....一個(gè)warp中如果有線程的條件為true，則執(zhí)行if子句，其它為false的線程將等待if執(zhí)行完成。然后執(zhí)行else語(yǔ)句，當(dāng)條件為true的線程則等待else執(zhí)行完成。
為了獲得更高的性能，盡量避免warp分支，warp是32個(gè)連續(xù)的線程，在算法允許的情況下，可以將數(shù)據(jù)分割，使同一個(gè)warp避免分支。
Example
實(shí)現(xiàn)偶數(shù)的線程計(jì)算結(jié)果為100，奇數(shù)線程的計(jì)算結(jié)果為200.
// set up data size
int size = 64;
int blocksize = 64;
//線程分支
__global__ void mathKernel1(float *c)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float ia, ib;
ia = ib = 0.0f;
if (tid % 2 == 0)
{
ia = 100.0f;
}
else
{
ib = 200.0f;
}
c[tid] = ia + ib;
}
//沒(méi)有warp分支，設(shè)備利用率更高，計(jì)算結(jié)果相同，但是順序不同。
__global__ void mathKernel2(float *c)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float ia, ib;
ia = ib = 0.0f;
if ((tid / warpSize) % 2 == 0)
{
ia = 100.0f;
}
else
{
ib = 200.0f;
}
c[tid] = ia + ib;
}
分支效率：
???不知道為什么，我的電腦運(yùn)行結(jié)果很奇怪Tesla K80，反而是kernel1運(yùn)行時(shí)間更短，kernel2運(yùn)行時(shí)間更長(zhǎng)。
warmingup:不分支
mathKernel1：分支
mathKernel2：不分支
mathKernel3：分支
mathKernel4：不分支
以前的nvprof計(jì)算warp分支的效率，但是我的CUDA8.0已經(jīng)提示沒(méi)有該metrics了和events。
$ nvprof --metrics branch_efficiency
$ nvprof --events branch,divergent_branch
