參考：blog.sina.com.cn/s/blog_80ce3a550101lntp.html

GPU的硬體結構中與CUDA相關的幾個概念：thread block grid warp sp sm
streaming processor(sp): 最基本的處理單元，streaming processor 最後具體的指令和任務都是在sp上處理的。GPU進行平行計算，也就是很多個sp同時做處理。現在SP的術語已經有點弱化了，而是直接使用thread來代替。一個SP對應一個thread。
Warp：warp是SM排程和執行的基礎概念，同時也是一個硬體概念，注意到Warp實際上是一個和硬體相關的概念，通常一個SM中的SP(thread)會分成幾個warp(也就是SP在SM中是進行分組的，物理上進行的分組)，每一個WARP中在Tegra中是32個thread.這個WARP中的32個thread(sp)是一起工作的，執行相同的指令，如果沒有這麼多thread需要工作，那麼這個WARP中的一些thread(sp)是不工作的。
每一個執行緒都有自己的暫存器記憶體和local memory，一個warp中的執行緒是同時執行的，也就是當進行平行計算時，執行緒數儘量為32的倍數，如果執行緒數不上32的倍數的話；假如是1，則warp會生成一個掩碼，當一個指令控制器對一個warp單位的執行緒傳送指令時，32個執行緒中只有一個執行緒在真正執行，其他31個 程序會進入靜默狀態。

streaming multiprocessor(sm):多個sp加上其他的一些資源組成一個sm, streaming multiprocessor. 其他資源也就是儲存資源，共享記憶體，寄儲器等。可見，一個SM中的所有SP是先分成warp的，是共享同一個memory和instruction unit。
每個SM通過使用兩個特殊函式(Special Function Unit,SFU)單元進行超越函式和屬性插值函式（根據頂點屬性來對畫素進行插值）計算。SFU用來執行超越函式、插值以及其他特殊運算
在 G80/G92 的架構下，總共會有 128 個 SP，以 8 個 SP 為一組，組成 16 個 SM，再以兩個 SM 為一個 TPC，共分成 8 個 TPC 來運作。而在新一代的 GT200 裡，SP 則是增加到 240 個，還是以 8 個 SP 組成一個 SM，但是改成以 3 個 SM 組成一個 TPC，共 10 組 TPC。
在Tegra系列中，一個GPU中通常只有2個SM，每一個SM中包含4個WARP，每一個warp中有32個thread(SP),因此，一個SM中有128個SP。
Stream：流（Stream）是一系列順序執行的命令，流之間相對無序或併發的執行他們的命令。

軟體概念：
thread–>block–>grid：在利用cuda進行程式設計時，一個grid分為多個block，而一個block分為多個thread。其中任務劃分到是否影響最後的執行效果。劃分的依據是任務特性和GPU本身的硬體特性。GRID,BLOCK,THREAD是軟體概念,而非硬體的概念。

從硬體角度講，一個GPU由多個SM組成（當然還有其他部分），一個SM包含有多個SP（以及還有暫存器資源，shared memory資源，L1cache，scheduler，SPU，LD/ST單元等等），1.x硬體，一個SM包含8個SP，2.0是32個，2.1是48個，3.0和3.5是192個。以及SP目前也稱為CUDA CORE，而SM目前也稱為MP，在KEPLER架構（SM3.0和3.5）下也稱為SMX。

從軟體角度講，CUDA因為是SIMT的形式，GRID，block，thread是thread的組織形式。最小的邏輯單位是一個thread，最小的硬體執行單位是thread warp（簡稱warp），若干個thread（典型值是128~512個）組成一個block，block被載入到SM上執行，多個block組成整體的GRID。

這裡為什麼要有一箇中間的層次block呢？這是因為CUDA通過這個概念，提供了細粒度的通訊手段，因為block是載入在SM上執行的，所以可以利用SM提供的shared memory和__syncthreads()功能實現執行緒同步和通訊，這帶來了很多好處。而block之間，除了結束kernel之外是無法同步的，一般也不保證執行先後順序，這是因為CUDA程式要保證在不同規模（不同SM數量）的GPU上都可以執行，必須具備規模的可擴充套件性，因此block之間不能有依賴。

從上面的表述中可以總結：
在GPU中最小的硬體單元是SP(這個術語通常使用thread來代替),而硬體上一個SM中的所有SP在物理上是分成了幾個WARP(每一個warp包含一些thread),warp中的SP是可以同時工作的，但是執行相同的指令，也就是說取指令單元取一條指令同時發射給WARP中的所有的SP(假設SP都需要工作，否則有些是idle的).可見，在硬體上一個SM->WARPS->threads(sp).
對於軟體thread組織來看，因為一個SM中是分WARP的，而一個WARP包含一定數目（比如Tegra 32個)的sp(thread),因此最好按照這個數目來組織thread，否則硬體該warp上有些SP是不工作的。

這就是CUDA的兩級並行結構。

總而言之，一個kernel對應一個GRID，該GRID又包含若干個block，block內包含若干個thread。GRID跑在GPU上的時候，可能是獨佔一個GPU的，也可能是多個kernel併發佔用一個GPU的（需要fermi及更新的GPU架構支援）。

block是resident在SM上的，一個SM可能有一個或多個resident blocks，需要具體根據資源佔用分析。

thread以warp為單位被SM的scheduler 發射到SP或者其他單元，如SFU，LD/ST unit執行相關操作，需要等待的warp會被切出（依然是resident 狀態），以空出執行單元給其他warps。