Warp

邏輯上，所有thread是并行的，但是，從硬件的角度來(lái)說(shuō)，實(shí)際上并不是所有的thread能夠在同一時(shí)刻執(zhí)行，接下來(lái)我們將解釋有關(guān)warp的一些本質(zhì)。

Warps and Thread Blocks

warp是SM的基本執(zhí)行單元。一個(gè)warp包含32個(gè)并行thread，這32個(gè)thread執(zhí)行于SMIT模式。也就是說(shuō)所有thread執(zhí)行同一條指令，并且每個(gè)thread會(huì)使用各自的data執(zhí)行該指令。

block可以是一維二維或者三維的，但是，從硬件角度看，所有的thread都被組織成一維，每個(gè)thread都有個(gè)唯一的ID(ID的計(jì)算可以在之前的博文查看)。

每個(gè)block的warp數(shù)量可以由下面的公式計(jì)算獲得：

?

一個(gè)warp中的線(xiàn)程必然在同一個(gè)block中，如果block所含線(xiàn)程數(shù)目不是warp大小的整數(shù)倍，那么多出的那些thread所在的warp中，會(huì)剩余一些inactive的thread，也就是說(shuō)，即使湊不夠warp整數(shù)倍的thread，硬件也會(huì)為warp湊足，只不過(guò)那些thread是inactive狀態(tài)，需要注意的是，即使這部分thread是inactive的，也會(huì)消耗SM資源。

?

Warp Divergence

控制流語(yǔ)句普遍存在于各種編程語(yǔ)言中，GPU支持傳統(tǒng)的，C-style，顯式控制流結(jié)構(gòu)，例如if…else,for,while等等。

CPU有復(fù)雜的硬件設(shè)計(jì)可以很好的做分支預(yù)測(cè)，即預(yù)測(cè)應(yīng)用程序會(huì)走哪個(gè)path。如果預(yù)測(cè)正確，那么CPU只會(huì)有很小的消耗。和CPU對(duì)比來(lái)說(shuō)，GPU就沒(méi)那么復(fù)雜的分支預(yù)測(cè)了（CPU和GPU這方面的差異的原因不是我們關(guān)心的，了解就好，我們關(guān)心的是由這差異引起的問(wèn)題）。

這樣我們的問(wèn)題就來(lái)了，因?yàn)樗型粋€(gè)warp中的thread必須執(zhí)行相同的指令，那么如果這些線(xiàn)程在遇到控制流語(yǔ)句時(shí)，如果進(jìn)入不同的分支，那么同一時(shí)刻除了正在執(zhí)行的分之外，其余分支都被阻塞了，十分影響性能。這類(lèi)問(wèn)題就是warp divergence。

請(qǐng)注意，warp divergence問(wèn)題只會(huì)發(fā)生在同一個(gè)warp中。

下圖展示了warp divergence問(wèn)題：

為了獲得最好的性能，就需要避免同一個(gè)warp存在不同的執(zhí)行路徑。避免該問(wèn)題的方法很多，比如這樣一個(gè)情形，假設(shè)有兩個(gè)分支，分支的決定條件是thread的唯一ID的奇偶性：

__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大小為步調(diào)，然后取奇偶，如下：

__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; }

代碼：

View Code

編譯運(yùn)行：

$ nvcc -O3 -arch=sm_20 simpleDivergence.cu -o simpleDivergence $./simpleDivergence

輸出：

$ ./simpleDivergence using Device 0: Tesla M2070 Data size 64 Execution Configuration (block 64 grid 1) Warmingup elapsed 0.000040 sec mathKernel1 elapsed 0.000016 sec mathKernel2 elapsed 0.000014 sec

我們也可以直接使用nvprof（之后會(huì)詳細(xì)介紹）這個(gè)工具來(lái)度量性能：

$ nvprof --metrics branch_efficiency ./simpleDivergence

輸出為：

Kernel: mathKernel1(void) 1 branch_efficiency Branch Efficiency 100.00% 100.00% 100.00% Kernel: mathKernel2(void) 1 branch_efficiency Branch Efficiency 100.00% 100.00% 100.00%

Branch Efficiency的定義如下：

到這里你應(yīng)該在奇怪為什么二者表現(xiàn)相同呢，實(shí)際上當(dāng)我們的代碼很簡(jiǎn)單，可以被預(yù)測(cè)時(shí)，CUDA的編譯器會(huì)自動(dòng)幫助優(yōu)化我們的代碼。稍微提一下GPU分支預(yù)測(cè)（理解的有點(diǎn)暈，不過(guò)了解下就好），這里，一個(gè)被稱(chēng)為預(yù)測(cè)變量的東西會(huì)被設(shè)置成1或者0，所有分支都會(huì)得到執(zhí)行，但是只有預(yù)測(cè)值為1時(shí)，才會(huì)得到執(zhí)行。當(dāng)條件狀態(tài)少于某一個(gè)閾值時(shí)，編譯器會(huì)將一個(gè)分支指令替換為預(yù)測(cè)指令，因此，現(xiàn)在回到自動(dòng)優(yōu)化問(wèn)題，一份較長(zhǎng)的代碼就會(huì)導(dǎo)致warp divergence了。

可以使用下面的命令強(qiáng)制編譯器不優(yōu)化（貌似不怎么管用）：

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

?

Resource Partitioning

一個(gè)warp的context包括以下三部分：

Program counter

Register

Shared memory

再次重申，在同一個(gè)執(zhí)行context中切換是沒(méi)有消耗的，因?yàn)樵谡麄€(gè)warp的生命期內(nèi)，SM處理的每個(gè)warp的執(zhí)行context都是on-chip的。

每個(gè)SM有一個(gè)32位register集合放在register file中，還有固定數(shù)量的shared memory，這些資源都被thread瓜分了，由于資源是有限的，所以，如果thread比較多，那么每個(gè)thread占用資源就叫少，thread較少，占用資源就較多，這需要根據(jù)自己的要求作出一個(gè)平衡。

資源限制了駐留在SM中blcok的數(shù)量，不同的device，register和shared memory的數(shù)量也不同，就像之前介紹的Fermi和Kepler的差別。如果沒(méi)有足夠的資源，kernel的啟動(dòng)就會(huì)失敗。

當(dāng)一個(gè)block或得到足夠的資源時(shí)，就成為active block。block中的warp就稱(chēng)為active warp。active warp又可以被分為下面三類(lèi)：

Selected warp

Stalled warp

Eligible warp

SM中warp調(diào)度器每個(gè)cycle會(huì)挑選active warp送去執(zhí)行，一個(gè)被選中的warp稱(chēng)為selected warp，沒(méi)被選中，但是已經(jīng)做好準(zhǔn)備被執(zhí)行的稱(chēng)為Eligible warp，沒(méi)準(zhǔn)備好要執(zhí)行的稱(chēng)為Stalled warp。warp適合執(zhí)行需要滿(mǎn)足下面兩個(gè)條件：

32個(gè)CUDA core有空

所有當(dāng)前指令的參數(shù)都準(zhǔn)備就緒

例如，Kepler任何時(shí)刻的active warp數(shù)目必須少于或等于64個(gè)（GPU架構(gòu)篇有介紹）。selected warp數(shù)目必須小于或等于4個(gè)（因?yàn)閟cheduler有4個(gè)？不確定，至于4個(gè)是不是太少則不用擔(dān)心，kernel啟動(dòng)前，會(huì)有一個(gè)warmup操作，可以使用cudaFree()來(lái)實(shí)現(xiàn)）。如果一個(gè)warp阻塞了，調(diào)度器會(huì)挑選一個(gè)Eligible warp準(zhǔn)備去執(zhí)行。

CUDA編程中應(yīng)該重視對(duì)計(jì)算資源的分配：這些資源限制了active warp的數(shù)量。因此，我們必須掌握硬件的一些限制，為了最大化GPU利用率，我們必須最大化active warp的數(shù)目。

Latency Hiding

指令從開(kāi)始到結(jié)束消耗的clock cycle稱(chēng)為指令的latency。當(dāng)每個(gè)cycle都有eligible warp被調(diào)度時(shí)，計(jì)算資源就會(huì)得到充分利用，基于此，我們就可以將每個(gè)指令的latency隱藏于issue其它warp的指令的過(guò)程中。

和CPU編程相比，latency hiding對(duì)GPU非常重要。CPU cores被設(shè)計(jì)成可以最小化一到兩個(gè)thread的latency，但是GPU的thread數(shù)目可不是一個(gè)兩個(gè)那么簡(jiǎn)單。

當(dāng)涉及到指令latency時(shí)，指令可以被區(qū)分為下面兩種：

Arithmetic instruction

Memory instruction

顧名思義，Arithmetic? instruction latency是一個(gè)算數(shù)操作的始末間隔。另一個(gè)則是指load或store的始末間隔。二者的latency大約為：

10-20 cycle for arithmetic operations

400-800 cycles for global memory accesses

下圖是一個(gè)簡(jiǎn)單的執(zhí)行流程，當(dāng)warp0阻塞時(shí)，執(zhí)行其他的warp，當(dāng)warp變?yōu)閑ligible時(shí)從新執(zhí)行。

你可能想要知道怎樣評(píng)估active warps 的數(shù)量來(lái)hide latency。Little’s Law可以提供一個(gè)合理的估計(jì)：

?

對(duì)于Arithmetic operations來(lái)說(shuō)，并行性可以表達(dá)為用來(lái)hide ?Arithmetic latency的操作的數(shù)目。下表顯示了Fermi和Kepler相關(guān)數(shù)據(jù)，這里是以（a + b * c）作為操作的例子。不同的算數(shù)指令，throughput（吞吐）也是不同的。

這里的throughput定義為每個(gè)SM每個(gè)cycle的操作數(shù)目。由于每個(gè)warp執(zhí)行同一種指令，因此每個(gè)warp對(duì)應(yīng)32個(gè)操作。所以，對(duì)于Fermi來(lái)說(shuō)，每個(gè)SM需要640/32=20個(gè)warp來(lái)保持計(jì)算資源的充分利用。這也就意味著，arithmetic operations的并行性可以表達(dá)為操作的數(shù)目或者warp的數(shù)目。二者的關(guān)系也對(duì)應(yīng)了兩種方式來(lái)增加并行性：

Instruction-level Parallelism（ILP）：同一個(gè)thread中更多的獨(dú)立指令

Thread-level Parallelism （TLP）：更多并發(fā)的eligible threads

對(duì)于Memory operations，并行性可以表達(dá)為每個(gè)cycle的byte數(shù)目。

因?yàn)閙emory throughput總是以GB/Sec為單位，我們需要先作相應(yīng)的轉(zhuǎn)化。可以通過(guò)下面的指令來(lái)查看device的memory frequency：

$ nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory"

以Fermi為例，其memory frequency可能是1.566GHz，Kepler的是1.6GHz。那么轉(zhuǎn)化過(guò)程為：

?

乘上這個(gè)92可以得到上圖中的74，這里的數(shù)字是針對(duì)整個(gè)device的，而不是每個(gè)SM。

有了這些數(shù)據(jù)，我們可以做一些計(jì)算了，以Fermi為例，假設(shè)每個(gè)thread的任務(wù)是將一個(gè)float（4 bytes）類(lèi)型的數(shù)據(jù)從global memory移至SM用來(lái)計(jì)算，你應(yīng)該需要大約18500個(gè)thread，也就是579個(gè)warp來(lái)隱藏所有的memory latency。

?

Fermi有16個(gè)SM，所以每個(gè)SM需要579/16=36個(gè)warp來(lái)隱藏memory latency。

?

Occupancy

當(dāng)一個(gè)warp阻塞了，SM會(huì)執(zhí)行另一個(gè)eligible warp。理想情況是，每時(shí)每刻到保證cores被占用。Occupancy就是每個(gè)SM的active warp占最大warp數(shù)目的比例：

?

我們可以使用的device篇提到的方法來(lái)獲取warp最大數(shù)目：

cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device);

然后用maxThreadsPerMultiProcessor來(lái)獲取具體數(shù)值。

grid和block的配置準(zhǔn)則：

• 保證block中thrad數(shù)目是32的倍數(shù)。
• 避免block太小：每個(gè)blcok最少128或256個(gè)thread。
• 根據(jù)kernel需要的資源調(diào)整block。
• 保證block的數(shù)目遠(yuǎn)大于SM的數(shù)目。
• 多做實(shí)驗(yàn)來(lái)挖掘出最好的配置。

Occupancy專(zhuān)注于每個(gè)SM中可以并行的thread或者warp的數(shù)目。不管怎樣，Occupancy不是唯一的性能指標(biāo)，Occupancy達(dá)到當(dāng)某個(gè)值是，再做優(yōu)化就可能不在有效果了，還有許多其它的指標(biāo)需要調(diào)節(jié)，我們會(huì)在之后的博文繼續(xù)探討。

Synchronize

同步是并行編程的一個(gè)普遍的問(wèn)題。在CUDA的世界里，有兩種方式實(shí)現(xiàn)同步：

System-level：等待所有host和device的工作完成

Block-level：等待device中block的所有thread執(zhí)行到某個(gè)點(diǎn)

因?yàn)镃UDA API和host代碼是異步的，cudaDeviceSynchronize可以用來(lái)停住CUP等待CUDA中的操作完成：

cudaError_t cudaDeviceSynchronize(void);

因?yàn)閎lock中的thread執(zhí)行順序不定，CUDA提供了一個(gè)function來(lái)同步block中的thread。

__device__ void __syncthreads(void);

當(dāng)該函數(shù)被調(diào)用，block中的每個(gè)thread都會(huì)等待所有其他thread執(zhí)行到某個(gè)點(diǎn)來(lái)實(shí)現(xiàn)同步。

總結(jié)

以上是生活随笔為你收集整理的cuda warp的全部?jī)?nèi)容，希望文章能夠幫你解決所遇到的問(wèn)題。

如果覺(jué)得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。