本小節來自《大規模並行處理器編程實戰》第四節，該書是很好的從內部原理結構上來講述了CUDA的，對於理解CUDA很有幫助，借以博客的形式去繁取間，肯定會加入自己個人理解，所以有錯誤之處還望指正。
一、塊索引與線程索引
CUDA是細粒度的，數據並行的輕量級線程，在啟動一個CUDA的一個Kernel函數的時候，就會創建一個線程網格grid，該網格中的所有線程都是執行該kernel函數的，對於kernel的函數調用形式：kernel<<<dimGrid, dimBlk >>>（argument list）；可以看出，一個函數代表一個網格，里面是由塊構成，而每個塊中就有許多的輕量級線程了。塊是cuda的資源調度單位，塊之間交互相較於塊內較為困難：
如上圖所示，一個網格有許多塊，每個塊有差不多3維的線程形式，其中線程中的具體的索引就是對應着具體的線程計算，這可以用在矩陣操作中。所以一個kernel對應一個網格，那么什么對應每個塊和塊中的線程呢？是blockIdx （blockIdx.x ，blockIdx.y，blockIdx.z）和threadIdx（threadIdx.x ，threadIdx.y，threadIdx.z），這兩個變量是c語言版本的dim3結構其中有3個無符號整數類型的字段：x、y、z，由CUDA內置的變量，對應的值由運行時系統在運行的時候分配，前者就是用來索引網格中塊的位置；后者就是用來索引塊中的線程的位置。
一個塊中的線程是有上限的（在http://blog.csdn.net/shouhuxianjian/article/details/42427285 的最后有所提及），所以如果只用塊中的索引來矩陣計算勢必是不夠的，這里提下，上面kernel的函數形式中dimGrid, dimBlk 都是三維形式，前者是用來說明這個網格中塊在x,y,z 三個維度上的分布；而后者是用來說明每個塊中線程在x,y,z 三個維度上的分布。但是根據書中p52頁所說，網格是一個二維的塊數組，所以網格維度參數中第三個字段z通常是默認為1的，（個人：這里說下，這是這本書說的，那時候還挺早的，現在不一定，因為按照cuda的屬性返回值看出來不同的顯卡有不同的信息，有興趣的可以自己測試下）：
例子代碼如下：
dim3 dimGrid（32,32,1）; dim3 dimBlk(16,16,16);//這是錯的，因為16^3大於512（早期顯卡），也大於1024（780ti），所以這里需要注意分配 addCUDA<<<dimGrid, dimBlk>>>(參數列表)；
上面的kernel函數自己寫的，這里是為了說明<<<>>>中的兩個參數，不同維度上最大值請先調取自己顯卡的信息核實再決定，這個函數就是告訴運行時系統啟動kernel函數了，通過blockIdx和threadIdx,兩個的組合就能生成成千上萬，上百萬的線程用來計算。
上圖就是一個矩陣乘法的例子，行向量乘以列向量作為生成向量的一個元素。如果矩陣很大，那么只靠threadIdx來是無法索引的，可以通過下圖這種將整個矩陣划分成不同的子塊矩陣來計算，剛好就對應了塊索引和塊中的元素：
下面給出一個例子，因為暫時看到的資料和教程還沒有涉及到2維或者3維的操作，就先按照行主排序或者列主排序來將矩陣划分成1維的向量來計算，不過對於矩陣的元素索引還是遵從2維形式的。
__global__ void MatrixMulKernel(float *Md,float *Pd,int Width){
int Row = blockIdx.y*TILE_WIDTH+threadIdx.y;//計算矩陣的第幾行
int Col = blockIdx.x*TILE_WIDTH+threadIdx.x;//計算矩陣的第幾列
float Pvalue= 0；//臨時變量
for(int k = 0;k<Witdh;++k){
Pvalue += Md[Row*Width+k]*Nd[k*Witdh +Col];//循環計算行向量乘以列向量
Pd[Row*Width +Col] = Pvalue;//將值賦值給結果矩陣
}
}
//執行配置參數
dim3 dimGrid(Width/TILE_WIDTH,Width/TILE_WIDTH);
dim3 dimBlk(TILE_WIDTH,TILE_WIDTH);
//啟動核函數，其中WIDTH為矩陣的維度，這里默認方形矩陣，TILE_WIDTH是每個子方形矩陣的維度
MatrixMulKernel<<<dimGrid,dimBlk>>>(Md,Nd,Pd,Width)；
二、同步與透明可擴展性
CUDA中有一個用於一個塊中所有線程的柵欄同步函數 __syncthreads()來用於協調這個塊中的所有線程，該函數是用來進行塊中的線程等待的，即假設一個塊中有500個線程，kernel中有工作a，工作b，前250個線程完成了工作a，可是剩下250個線程還沒完成，這時候使用了該函數，那么前250個線程就需要等待后續的250個線程完成工作a之后在一起整個500個線程接着完成工作b。該函數具有唯一性，不會被異步，假設在 if -else 語句中:
if（argument） __syncthreads()； else __syncthreads()； end
這種情況下無需擔心，因為該函數的意思是要么塊中所有的線程都按照其中包含這個函數的路徑執行，要么都不執行，不然一半停留在if后面，一半停留在else后面，互相等待，豈不是成了死鎖了。
正式因為該函數基於塊，所以一個塊中相互之間的線程等待時間還能夠忍受，如果是基於塊的相互等待那么就不能忍受了。正是這種相同塊需要柵欄同步，不同塊無需柵欄同步，所以不同的顯卡上程序才有可伸縮性：
如上圖，程序剛好划分成8個塊，前者設備每次執行2個塊，后者每次可以執行4個塊，但是都能很靈活的划分任務進行執行。
三、線程分配
啟動kernel函數之后，運行時系統會生成相應的線程網格，然后以塊為單位把這些線程分配給這些執行資源。這些執行資源組織成多核流處理器（streaming multiprocessor，SM）。書中這里接着說CUDA運行時系統會自動減少每個SM中分配的塊的數量，這個情況是發生在任何一種或者多種資源都不能滿足一個SM中所有的塊運行的時候。也就是在之前的人為分配好塊和線程的參數的時候，在執行的時候運行時系統發現每個sm無法使用滿運行的所有的塊，那么就自動讓那幾個塊休息不工作（個人：結合CUDA2.2講述的存儲器等的限制，應該是假設一個SM最大支持768個線程，而每個塊最大支持256個線程，所以SM算下來最少可以支持3個塊，但是如果在kernel中設置好了參數，但是發現由於存儲器等限制，那么就減少塊數，本來是3塊的，現在一個SM只能駐留2個塊了）。
這里引用《cuda並行程序設計》6.3章節中一段話：“GPU與CPU不同，GPU不使用寄存器重命名的機制，而是致力於為每個線程都分配真實的寄存器，因此當需要上下文切換時，所需要的操作就是將指向當前寄存器組的選擇器或者指針進行更新，指向下一個要執行的線程束的寄存器組，因此是零開銷。...如果一個內核函數的每個線程需要的寄存器過多，則每個SM中GPU能夠調度的線程塊的數量就會受到限制，因此總的可以執行的線程數量也會受到限制。開啟的線程數量過少會造成硬件無法被充分利用，性能急劇下降，但是開啟過多又意味着資源可能短缺，調度到SM上的線程塊數量會減少...由於所使用硬件不同，每個SM可供所有線程使用的寄存器空間大小也不同，分分別有8KB，16KB，32KB以及64KB。牢記，每個線程中每個變量都會占用一個寄存器，舉個例子，每個SM擁有32KB的寄存器空間，如果每個線程塊有256個線程，則對於每個變量是32位，需要占用4個Bytes的情況下，每個線程可以使用32×1024/4/256=32個寄存器。而每個SM也有所謂最大可使用上限的寄存器數量，如果當前線程塊上的寄存器數目是允許的最大值時，每個SM只會處理一個線程塊”。
一旦一個塊分配給了一個SM，那么該塊就會被以大小為32的線程warp所划分，warp的個數的多少（個人：英文版是size，覺得應該說的是分配給每個SM中划分的warp的個數，而每個warp是固定的32個線程）是在具體實現的時候指定的，對於程序員來說這個是透明的，因為無需我們關心。這個參數也可以在不同的顯卡的屬性信息中得知，在SM中，warp才是線程調度的單位，而不是單個的線程，這個概念不屬於CUDA的規范，但是卻有助於理解和優化在特定CUDA設備上運行的程序的性能。通常cuda會以半個線程束做一次調度，此時可以將一半的線程束的取數據操作合並成一次連續取數據操作，且指令取一次，並廣播給這整個線程束（cuda並行程序設計5.5章節：“由於硬件每次只能為一個線程束獲取一條指令...在指令執行層，硬件的調度是基於半個線程束，而不是整個線程束”）。塊中相鄰的Warp的threadIdx的值是連續的，假設第一個waro包含線程0-31，那么第二個就包含線程32-63，以此類推。
上圖就是一個SM，這個sm的上面就是三個不同的塊，和每個塊划分成不同的warp，每個warp中有32個線程。只要給定塊的大小和每個sm中塊的數量，就能計算每個sm中駐留的warp的數量了，上圖中假設每個塊有256個線程，那么每個塊就有256/32 = 8個warp了，如果每個sm中只有3個塊，那么每個sm就有8×3 = 24個warp，假設在G80類型的顯卡中每個sm最多駐留768個線程，每個sm最多駐留768/32 = 24個warp。但是如果說每個SM中只有8個sp（streaming processors）（個人：一個sp執行一個warp還是一個線程？應該是執行一個線程，不過是流水形式的串行執行），那么為什么一個sm中會有這么多的warp。這里就在於CUDA的處理器需要高效的執行長延時操作，比如訪問全局存儲器（時間比訪問寄存器長）。如果warp中線程執行一條指令需要等待前面啟動的長延時操作的結果（就是該warp需要從全局存儲器中提取數值計算），那么就不選擇該warp，而是選擇另一個不需要等待結果的駐留的warp（這個warp已經得到了自己需要的結果，所以已經無需等待了，可以直接執行了），當多個warp准備執行的時候，采用優先機制選擇一個warp執行，這種機制不產生延時的線程先執行，這就是所謂的延時隱藏（latency hiding）。執行這種處於就緒狀態的warp不會帶來多余的時間開銷，被稱之為“零開銷線程調度（zero-overhead thread scheduling）”。
所以我們在寫網格維度和塊的維度的時候，最好考慮到這個，將每個塊中的線程的數量設置成32的倍數，而且盡可能的先滿足一個塊中線程的上限，因為塊內線程調度相比較與塊間調度來說，時間開銷更小，更能加速程序運行，也就是更少的粗粒度，更多的細粒度。（該想法是錯的，見下面博客）
順帶看了下其他人博客http://blog.csdn.net/mysniper11/article/details/8269776 的相關理解。其中的：
“每一個塊內線程數應該首先是32的倍數，因為這樣的話可以適應每一個warp包含32個線程的要求，每一個warp中串行執行，這就要求每一個線程中不可以有過多的循環或者需要的資源過多。但是每一個塊中如果線程數過多，可能由於線程中參數過多帶來存儲器要求過大，從而使SM處理的效率更低。所以，在函數不是很復雜的情況下，可以適當的增加線程數目，線程中不要加入循環。在函數比較復雜的情況下，每一個塊中分配32或是64個線程比較合適。每一個SM同時處理一個塊，只有在粗粒度層面上以及細粒度層面上均達到平衡，才能使得GPU的利用到達最大。“解答了上面錯誤想法的地方。