CUDA性能優化----warp深度解析

時間 2019-12-18

原文原文鏈接

本文轉自：http://blog.163.com/wujiaxing009@126/blog/static/71988399201701224540201/編程

一、引言

CUDA性能優化----sp, sm, thread, block, grid, warp概念中提到：邏輯上，CUDA中全部thread是並行的，可是，從硬件的角度來講，實際上並非全部的thread可以在同一時刻執行，接下來咱們將深刻學習和了解有關warp的一些本質。

二、Warps and Thread Blocks

warp是SM的基本執行單元。一個warp包含32個並行thread，這32個thread執行於SIMT模式。也就是說全部thread執行同一條指令，而且每一個thread會使用各自的data執行該指令。

block能夠是1D、2D或者3D的，可是，從硬件角度看，全部的thread都被組織成一維的，每一個thread都有個惟一的ID。每一個block的warp數量能夠由下面的公式計算得到：

一個warp中的線程必然在同一個block中，若是block所含線程數目不是warp大小的整數倍，那麼多出的那些thread所在的warp中，會剩餘一些 inactive的thread，也就是說，即便湊不夠warp整數倍的thread，硬件也會爲warp湊足，只不過那些thread是inactive狀態，須要注意的是，即便這部分thread是inactive的，也會消耗SM資源，這點是編程時應避免的。

三、Warp Divergence(warp分歧)

控制流語句廣泛存在於各類編程語言中，GPU支持傳統的、C-style的顯式控制流結構，例如if…else,for,while等等。

CPU有複雜的硬件設計能夠很好的作分支預測，即預測應用程序會走哪一個path分支。若是預測正確，那麼CPU只會有很小的消耗。和CPU對比來講，GPU就沒那麼複雜的分支預測了。

這樣問題就來了，由於全部同一個warp中的thread必須執行相同的指令，那麼若是這些線程在遇到控制流語句時，若是進入不一樣的分支，那麼同一時刻除了正在執行的分支外，其他分支都被阻塞了，十分影響性能。這類問題就是warp divergence。

注意，warp divergence問題只會發生在同一個warp中。下圖展現了warp divergence問題：

爲了得到最好的性能，就須要避免同一個warp存在不一樣的執行路徑。避免該問題的方法不少，好比這樣一個情形，假設有兩個分支，分支的決定條件是thread的惟一ID的奇偶性，kernel函數以下(simpleWarpDivergence.cu)：

__global__ void mathKernel1(float *c) 性能優化

{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float a, b;
a = b = 0.0f;
if (tid % 2 == 0)
a = 100.0f;
else
b = 200.0f;
c[tid] = a + b;
}架構

一種方法是，將條件改成以warp大小爲步調，而後取奇偶，代碼以下：併發

__global__ void mathKernel2(void) 異步

{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float a, b;
a = b = 0.0f;
if ((tid / warpSize) % 2 == 0)
a = 100.0f;
else
b = 200.0f;
c[tid] = a + b;
}編程語言

經過測試發現兩個kernel函數性能相近，到這裏你應該在奇怪爲何兩者表現相同呢，其實是由於當咱們的代碼很簡單，能夠被預測時，CUDA的編譯器會自動幫助優化咱們的代碼。（稍微提一下GPU分支預測，這裏一個被稱爲預測變量的東西會被設置成1或者0，全部分支都會執行，可是隻有預測變量值爲1時，該分支纔會獲得執行。當條件狀態少於某一個閾值時，編譯器會將一個分支指令替換爲預測指令。）所以，如今回到自動優化問題，一段較長的代碼就可能會致使warp divergence問題了。

可使用下面的命令強制編譯器不作優化：

$ nvcc -g -G -arch=sm_20 simpleWarpDivergence.cu -o simpleWarpDivergence

四、Resource Partitioning(資源劃分)

一個warp的context包括如下三部分：

Program counter
Register
Shared memory

再次重申，在同一個執行context中切換是沒有消耗的，由於在整個warp的生命期內，SM處理的每一個warp的執行context都是「on-chip」的。

每一個SM有一個32位register集合放在register file中，還有固定數量的shared memory，這些資源都被thread瓜分了，因爲資源是有限的，因此，若是thread數量比較多，那麼每一個thread佔用資源就比較少，反之若是thread數量較少，每一個thread佔用資源就較多，這須要根據本身的需求做出一個平衡。

資源限制了駐留在SM中blcok的數量，不一樣的GPU，register和shared memory的數量也不一樣，就像Fermi和Kepler架構的差異。若是沒有足夠的資源，kernel的啓動就會失敗。下圖是計算能力爲2.x和3.x的device參數對比：

當一個block得到到足夠的資源時，就成爲 active block。block中的warp就稱爲 active warp。active warp又能夠被分爲下面三類：

Selected warp
Stalled warp
Eligible warp

SM中 warp調度器每一個cycle會挑選active warp送去執行，一個被選中的warp稱爲 Selected warp，沒被選中，可是已經作好準備被執行的稱爲 Eligible warp，沒準備好要被執行的稱爲 Stalled warp。warp適合執行須要知足下面兩個條件：

32個CUDA core有空
全部當前指令的參數都準備就緒

例如，Kepler架構GPU任什麼時候刻的active warp數目必須少於或等於64個。selected warp數目必須小於或等於4個（由於scheduler有4個？不肯定，至於4個是否是太少則不用擔憂，kernel啓動前，會有一個 warmup操做，可使用cudaFree()來實現）。若是一個warp阻塞了，調度器會挑選一個Eligible warp準備去執行。

CUDA編程中應該重視對計算資源的分配：這些資源限制了active warp的數量。所以，咱們必須掌握硬件的一些限制，爲了最大化GPU利用率，咱們必須最大化active warp的數目。

五、Latency Hiding(延遲隱藏)

指令從開始到結束消耗的clock cycle稱爲指令的latency。當每一個cycle都有eligible warp被調度時，計算資源就會獲得充分利用，基於此，咱們就能夠將每一個指令的latency隱藏於issue其它warp的指令的過程當中。

和CPU編程相比， latency hiding對GPU很是重要。CPU cores被設計成能夠最小化一到兩個thread的latency，可是GPU的thread數目可不是一個兩個那麼簡單。

當涉及到指令latency時，指令能夠被區分爲下面兩種：

Arithmetic instruction
Memory instruction

顧名思義，Arithmetic instruction latency是一個算術操做的始末間隔。另外一個則是指load或store的始末間隔。

兩者的latency大約爲：

10-20 cycle for arithmetic operations
400-800 cycles for global memory accesses

下圖是一個簡單的執行流程，當warp0阻塞時，執行其餘的warp，當warp變爲eligible時從新執行。

你可能想要知道怎樣評估active warps 的數量來hide latency。Little’s Law能夠提供一個合理的估計：

對於Arithmetic operations來講，並行性能夠表達爲用來hide Arithmetic latency的操做的數目。下表顯示了Fermi和Kepler架構的相關數據，這裏是以(a + b * c)做爲操做的例子。不一樣的算術指令，throughput（吞吐）也是不一樣的。

這裏的throughput定義爲每一個SM每一個cycle的操做數目。因爲每一個warp執行同一種指令，所以每一個warp對應32個操做。因此，對於Fermi來講，每一個SM須要640/32=20個warp來保持計算資源的充分利用。這也就意味着，arithmetic operations的並行性能夠表達爲操做的數目或者warp的數目。兩者的關係也對應了兩種方式來增長並行性：