Professional CUDA C Programing
代碼下載：http://www.wrox.com/WileyCDA/
Warp資源分配
? 程序計(jì)數(shù)器
?寄存器
? 共享內(nèi)存
每個(gè)warp的上下文都是全部保留在SM上的，所以warp之間的切換沒(méi)有什么消耗窥浪。每個(gè)SM上的寄存器和共享內(nèi)存分配給線(xiàn)程塊漾脂，根據(jù)寄存器的多上和共享存儲(chǔ)器的大小可以決定同時(shí)駐留在一個(gè)SM上的warp數(shù)目和線(xiàn)程塊數(shù)目。
一個(gè)SM上同時(shí)駐留的線(xiàn)程越多笨鸡，則每個(gè)線(xiàn)程占用的寄存器數(shù)量越少坦冠。
一個(gè)SM上駐留的線(xiàn)程塊越多辙浑，每個(gè)線(xiàn)程塊占用的共享顯存越少。
查看GPU資源的一些限制信息：一個(gè)線(xiàn)程塊中最多1000個(gè)線(xiàn)程昔期，GT740只有2個(gè)SM硼一，但是Tesla K80有13個(gè)SM。
Active Warp
當(dāng)計(jì)算資源分配給了該線(xiàn)程塊時(shí)愧哟，該線(xiàn)程塊叫做active block蕊梧，其中包含的warps叫做active warps腮介。active waps分為以下3類(lèi)：
? Selected warp：warp調(diào)度器選中的warp，正在執(zhí)行的warp甘改。
? Stalled warp：還沒(méi)有準(zhǔn)備好執(zhí)行的warp
? Eligible warp：已經(jīng)準(zhǔn)備好執(zhí)行灭抑，但是還沒(méi)有執(zhí)行的warp腾节。（準(zhǔn)備好的限制條件：1.32個(gè)CUDA core可以用來(lái)執(zhí)行；2.當(dāng)前指令的所有參數(shù)都準(zhǔn)備就緒）
延遲隱藏Latency Hiding
GPU設(shè)計(jì)成處理大量輕量的并發(fā)的線(xiàn)程庆冕，最大化實(shí)現(xiàn)吞吐率救湖。
指令分為兩類(lèi)：
? 算術(shù)類(lèi)指令：10-20個(gè)時(shí)鐘周期
? 訪(fǎng)存類(lèi)指令：400-800個(gè)時(shí)鐘周期訪(fǎng)問(wèn)global memory
鞋既？？邑闺？陡舅？沒(méi)有看懂
Number of Required Warps = Latency × Throughput
Bandwidth VS hroughput：帶寬一般是理論上的峰值，吞吐量一般是實(shí)際達(dá)到的值灾炭。帶寬一般指單位時(shí)間內(nèi)數(shù)據(jù)的傳輸多少，吞吐量一般指單位時(shí)間內(nèi)完成的某種操作或計(jì)算田弥，比如說(shuō)單位時(shí)間內(nèi)完成的指令次數(shù)偷厦。
Occupancy
每個(gè)SM：occupancy = active warps/maximum warps
CUDA Toolkit中有一個(gè)幫助用戶(hù)確定grid和block大小的工具：/usr/local/cuda-8.0/tools
?小線(xiàn)程塊：每個(gè)塊的線(xiàn)程太少導(dǎo)致在所有硬件資源完全利用之前燕刻，已經(jīng)達(dá)到了每個(gè)SM最多的warps卵洗。比如一個(gè)線(xiàn)程塊只有10個(gè)thread，那么一個(gè)線(xiàn)程塊就要占用一個(gè)warp籍滴。
?大線(xiàn)程塊：每個(gè)塊太多的線(xiàn)程導(dǎo)致每個(gè)線(xiàn)程可以利用SM的資源更少榴啸。
選擇策略：根據(jù)kernel的計(jì)算量調(diào)整block的size晚岭，并進(jìn)行多次實(shí)驗(yàn)發(fā)現(xiàn)最優(yōu)的grid和block的設(shè)置坦报。
?每個(gè)block中含有的thread是warpSize的整數(shù)倍數(shù)。
?避免一個(gè)block太少的thread潜的，一個(gè)block最少128或256個(gè)線(xiàn)程字管。
?盡量使block的數(shù)目大于GPU的SM的數(shù)目。
同步Synchronization
屏障同步是許多并行編程語(yǔ)言中常見(jiàn)的原語(yǔ)亡呵。 在CUDA的同步可以在兩個(gè)層面上執(zhí)行：
?系統(tǒng)級(jí)別：等待主機(jī)和設(shè)備上的所有工作完成硫戈。
?塊級(jí)別：等待在設(shè)備上的線(xiàn)程塊中的所有線(xiàn)程到達(dá)執(zhí)行中的同一點(diǎn)（同步點(diǎn)）。
由于許多CUDA API調(diào)用和所有內(nèi)核啟動(dòng)都是與主機(jī)異步的汁胆，
cudaDeviceSynchronize可用于阻止主機(jī)應(yīng)用程序嫩码，直到所有CUDA操作（copies，內(nèi)核等）已經(jīng)完成：
cudaError_t cudaDeviceSynchronize(void);
__device__ void __syncthreads(void);
同一個(gè)block中threads要注意避免資源競(jìng)爭(zhēng)释牺，不同的warps的執(zhí)行順序是隨機(jī)的回挽，多個(gè)thread訪(fǎng)問(wèn)同一個(gè)變量要注意read-write，write-read等問(wèn)題祭刚，避免讀臟數(shù)據(jù)等墙牌。不同的block的執(zhí)行順序是隨機(jī)的喜滨。
可擴(kuò)展性Scalability
可擴(kuò)展：當(dāng)計(jì)算量增大時(shí)可以通過(guò)增加CUDA core來(lái)解決。
參考在不同數(shù)量的計(jì)算核心上執(zhí)行相同應(yīng)用程序代碼的能力
作為透明的可擴(kuò)展性棒口。 透明可擴(kuò)展的平臺(tái)拓寬了現(xiàn)有用例
應(yīng)用程序辜膝，并減輕開(kāi)發(fā)人員的負(fù)擔(dān)厂抖，因?yàn)樗鼈兛梢员苊鈱?duì)新的更改或不同的硬件。 可擴(kuò)展性比效率更重要七蜘。 一個(gè)可擴(kuò)展但效率低的系統(tǒng)可以通過(guò)簡(jiǎn)單地添加硬件核心來(lái)處理更大的工作負(fù)載墙懂。 效率很高但不可擴(kuò)展系統(tǒng)可能快速達(dá)到可實(shí)現(xiàn)性能的上限。
Checking Active Warps with nvprof
代碼來(lái)源：http://www.wrox.com/WileyCDA/
第三章sumMatrix.cu
//矩陣大小16384＊16384
// invoke kernel at host side
int dimx = 32;
int dimy = 32;
if(argc > 2)
{
dimx = atoi(argv[1]);
dimy = atoi(argv[2]);
}
dim3 block(dimx, dimy);
dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);
// grid 2D block 2D
__global__ void sumMatrixOnGPU2D(float *A, float *B, float *C, int NX, int NY)
{
unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
unsigned int idx = iy * NX + ix;
if (ix < NX && iy < NY)
{
C[idx] = A[idx] + B[idx];
}
}
$ nvprof --metrics achieved_occupancy ./sumMatrix 32 32
32 32: Achieved Occupancy 0.758286
32 16: Achieved Occupancy 0.777452
16 32: Achieved Occupancy 0.783850
16 32: Achieved Occupancy 0.810251
$ nvprof --metrics gld_throughput ./sumMatrix 32 32
32 32: Global Load Throughput 69.013GB/s
32 16: Global Load Throughput 71.597GB/s
16 32: Global Load Throughput 67.425GB/s
16 32:Global Load Throughput 70.240GB/s
$ nvprof --metrics gld_efficiency ./sumMatrix 32 32
