（1）線程束究竟是什麼？

概念上講，線程束是SM用SIMD方式所同時處理的工作粒度。當一個線程塊被指定到一個SM上，線程塊中的所有線程被分成了線程束，所以優化工作負載要考慮到線程束。或者這樣說，我們從軟體角度看線程，線程是獨立的且並行執行，對硬體來說，卻是被組織成32個線程為一堆的線程束來執行指令。因此，邏輯上我們每個線程塊可以是一維，二維，三維線程，但是物理上都是以32個線程為單位的一維物理布局。

杰倫：線程塊中線程數應該是32的倍數，不然最後那個沒有被填滿的線程束，會和其他填滿的線程束一樣消耗SM的資源，比如寶貴的寄存器。

（2）條件分支與線程束的關係

有一點非常重要，在一個線程束中所有線程在每個周期必須執行相同的指令，也就是說一個線程束中的所有線程，在面臨條件分支的時候，集體選擇同一條件執行，如果該線程束中有16個線程是選擇a條件執行，另外16個是選擇b條件執行，則實際的工作情況是：16個選擇a在執行的時候將會禁用另外16個線程執行b，也就是串列執行不同的條件分支，這個對性能的影響是很惡劣的。舉個栗子：

__global__ void kernel(float *c)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
float a =0,b=0;
if(index%2==0)
a = 100.0f;
else
b = 200.0f;
c[index] = a+b;
}

改進版（以線程束為單位，不再是以線程IDX奇偶，而是線程束的奇偶來分）

__global__ void kernel(float *c)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
float a =0,b=0;
if(（index/wrapSize)%2==0)//線程束為偶數執行，這樣可行的先決條件是，有一半線程執行a，有一半線程執行b
a = 100.0f;
else
b = 200.0f;
c[index] = a+b;
}

杰倫：在內核函數中，盡量避免數據分化，解決方法就是在條件分支處以線程束為單位。

（3）增大並行性來提高性能

實驗表明，當我們的blockDim.x是32的倍數，能獲增加載入吞吐量，而塊數較多時（當然不能超過硬體限制），可以有更多活躍的線程束，這直接影響著可實現佔用率，從而影響性能。但是，並不是說可實現佔用率越高，性能就越好，在實驗的時候要注意調整塊大小。

杰倫：當我們想通過設置塊的大小，網格的大小來獲得最佳性能時，不能單單考慮到可實現佔用率，載入吞吐量，要在指標之間尋求一個好的平衡，原則就是，塊數偏多時會增加並行性，塊第一維（第一維調整對載入吞吐量有絕對的影響）一定要是32的倍數。

（4）CUDA內存模型

無論是CPU還是GPU都遵循時間局部性和空間局部性，這就使得CPU和GPU內存結構設計都是低延遲低容量的層次結構，如cpu中（寄存器、緩存、主存、磁碟存儲器，容量越大，延遲越高）。在CUDA的內存層次中，排序為（寄存器，共享內存，本地內存，常量內存，紋理內存，全局內存），如圖所示：

GPU內存結構
CPU內存結構
  1. 寄存器：低容量低延遲，線程級的內存，為一個線程私有，生命周期與核函數一樣，核函數中聲明的沒有其他修飾符的自變數，通常是存儲在寄存器中，值得注意的是，在核函數中使用較少的寄存器將使在SM上有更多的常駐線程塊，每個SM上並發線程塊越多，使用率和性能越高，所以這也是為什麼要避免半個線程束這種不活躍的線程束的存在，原因就在於浪費了SM的寄存器。
  2. 本地內存：寄存器的備胎，寄存器中變數的溢出地，有限定條件的自變數（沒理解到這是什麼變數）但是當作為寄存器數據的溢出內存用的時候，卻與全局內存在一個地方，於是高延遲低帶寬。一般來說，對於計算能力在2.0及以上的GPU來說，本地內存是SM的一級緩存，設備的二級緩存。
  3. 共享內存：標記：__share__ 高帶寬低延遲，線程塊級，類似於CPU中的一級緩存，但他可編程。生命周期與線程塊一致，但需注意，片上內存64KB中一部分分配給共享內存，一部分分配給GPU一級緩存（可通過cudaFuncSetCacheConfig()手動配置），我的共享內存是48KB。共享內存是線程間通訊的必要內存，但是要注意搭配線程同步。在《CUDA by Example》--chapter05 code中有講過共享內存的具體用法。
  4. 常量內存：標記：__constant__ ，在全局空間和核函數之外進行聲明，對所有線程可見，初始化：cudaMemcpyToSymbol(); 適用範圍：當所有線程從相同的內存地址中讀取數據時，常量內存表現良好。具體用法，見書《CUDA by Example》--chapter06 。
  5. 紋理內存：。。。具體使用還是看《CUDA by Example》--chapter07吧！
  6. 全局內存：用得最多，實際意義上的顯存，也就預示著它是高容量高延遲。我的顯卡是N卡1060，顯示全局內存為2G，我們使用cudaMalloc()與cudaFree()就是在動態得分配全局內存。由於不存在塊間同步，因此不同線程塊並發得修改全局內存中的同一位置可以會出現未定義行為。
  7. 注意，紋理，常量，全局內存都是可以被所有線程訪問，且都是在主機上配置，但是紋理和常量是只讀訪問。

（5）

推薦閱讀：

相关文章