掌握如何組織線程是CUDA編程的重要部分。CUDA線程分成Grid和Block兩個層次。
由一個單獨的kernel啟動的所有線程組成一個grid，grid中所有線程共享global memory。一個grid由許多block組成，block由許多線程組成，grid和block都可以是一維二維或者三維，上圖是一個二維grid和二維block。
這里介紹幾個CUDA內置變量：
blockIdx：block的索引，blockIdx.x表示block的x坐標。
threadIdx：線程索引，同理blockIdx。
blockDim：block維度，上圖中blockDim.x=5.
gridDim：grid維度，同理blockDim。
一般會把grid組織成2D，block為3D。grid和block都使用dim3作為聲明，例如：
dim3 block(3);
dim3 grid((nElem+block.x-1)/block.x);
需要注意的是，dim3僅為host端可見，其對應的device端類型為uint3。
啟動CUDA kernel
CUDA kernel的調用格式為：
kernel_name<<<grid, block>>>(argument list);
其中grid和block即為上文中介紹的類型為dim3的變量。通過這兩個變量可以配置一個kernel的線程總和，以及線程的組織形式。例如：
kernel_name<<<4, 8>>>(argumentt list);
該行代碼表明有grid為一維，有4個block，block為一維，每個block有8個線程，故此共有4*8=32個線程。
一些基本的描述：
gridDim.x-線程網絡X維度上線程塊的數量
gridDim.y-線程網絡Y維度上線程塊的數量
blockDim.x-一個線程塊X維度上的線程數量
blockDim.y-一個線程塊Y維度上的線程數量
blockIdx.x-線程網絡X維度上的線程塊索引
blockIdx.y-線程網絡Y維度上的線程塊索引
threadIdx.x-線程塊X維度上的線程索引
threadIdx.y-線程塊Y維度上的線程索引
線程索引
一般，一個矩陣以線性存儲在global memory中的，並以行來實現線性：
在kernel里，線程的唯一索引非常有用，為了確定一個線程的索引，我們以2D為例：
- 線程和block索引
- 矩陣中元素坐標
- 線性global memory 的偏移
首先可以將thread和block索引映射到矩陣坐標：
ix = threadIdx.x + blockIdx.x * blockDim.x
iy = threadIdx.y + blockIdx.y * blockDim.y
之后可以利用上述變量計算線性地址：
idx = iy * nx + ix
上圖展示了block和thread索引，矩陣坐標以及線性地址之間的關系，謹記，相鄰的thread擁有連續的threadIdx.x，也就是索引為(0,0)(1,0)(2,0)(3,0)...的thread連續，而不是(0,0)(0,1)(0,2)(0,3)...連續，跟我們線代里玩矩陣的時候不一樣。
現在可以驗證出下面的關系：
thread_id（2,1）block_id（1,0） coordinate（6,1） global index 14 ival 14
下圖顯示了三者之間的關系：
線程塊分配要求：
CUDA設備上面的SM的數量；dev_prop.multiProcessorCount
每個SM上面SP（流處理器）的數量，真正執行指令的部件；
一個SM（多核流處理器）上最多可以分配的線程數量；
一個SM上分配線程塊的上線；
一個線程塊中的最大線程數量；dev_prop.maxthreadsPerBlock
每個維度允許分配的最大線程數量;x:dev_prop.maxthreadsDim[0] 、y:dev_prop.maxthreadsDim[1]
每個維度允許分配的最大線程塊數量;x:dev_prop.maxGridSize[0] 、y:dev_prop. maxGridSize[1]
Warp單元:SM中的線程調度單元，用來隱藏其它類型的操作延遲，由32個線程組成。
當一個網格啟動時，網格中的線程塊以任意順序分配到SM上，因此不同線程塊上的線程不能同步.
Warp調度
邏輯上，所有thread是並行的，但是，從硬件的角度來說，實際上並不是所有的thread能夠在同一時刻執行，接下來我們將解釋有關warp的一些本質。
同一個warp中的thread可以以任意順序執行，active warps被SM資源限制。當一個warp空閑時，SM就可以調度駐留在該SM中另一個可用warp。在並發的warp之間切換是沒什么消耗的，因為硬件資源早就被分配到所有thread和block，所以該新調度的warp的狀態已經存儲在SM中了。
SM可以看做GPU的心臟，寄存器和共享內存是SM的稀缺資源。CUDA將這些資源分配給所有駐留在SM中的thread。因此，這些有限的資源就使每個SM中active warps有非常嚴格的限制，也就限制了並行能力。所以，掌握部分硬件知識，有助於CUDA性能提升。
warp是SM的基本執行單元。一個warp包含32個並行thread，這32個thread執行於SMIT模式。也就是說所有thread執行同一條指令，並且每個thread會使用各自的data執行該指令。
block可以是一維二維或者三維的，但是，從硬件角度看，所有的thread都被組織成一維，每個thread都有個唯一的ID。
每個block的warp數量可以由下面的公式計算獲得：
一個warp中的線程必然在同一個block中，如果block所含線程數目不是warp大小的整數倍，那么多出的那些thread所在的warp中，會剩余一些inactive的thread，也就是說，即使湊不夠warp整數倍的thread，硬件也會為warp湊足，只不過那些thread是inactive狀態，需要注意的是，即使這部分thread是inactive的，也會消耗SM資源。
每一個塊上面最多可以有1024個線程。
每一個SM（多核流處理器）上面最多有1636個線程。
SM中的線程調度單元又將分配到的塊進行細分，將其中的線程組織成更小的結構，稱為線程束（warp）。所以由32個線程組成的Warp是CUDA程序執行的最小單位，並且同一個warp是嚴格串行的。warp的設計被用於隱藏延遲操作。盡可能充分利用每個線程塊的線程容量能得到足夠多的warp隱藏長延遲操作。