NVIDIA GPUs 以 SIMT (單指令,多線程)方式執(zhí)行稱(chēng)為 warps 的線程組。許多 CUDA 程序通過(guò)利用 warp 執(zhí)行來(lái)獲得高性能。在這個(gè)博客中,我們將展示如何使用 CUDA 9 中引入的原語(yǔ),使您的 warp 級(jí)編程安全有效。
扭曲級(jí)別基本體
NVIDIA GPUs 和 CUDA 編程模型采用一種稱(chēng)為 SIMT (單指令,多線程)的執(zhí)行模型。 SIMT 擴(kuò)展了計(jì)算機(jī)體系結(jié)構(gòu)的 弗林分類(lèi)學(xué) ,它根據(jù)指令和數(shù)據(jù)流的數(shù)量描述了四類(lèi)體系結(jié)構(gòu)。作為 Flynn 的四個(gè)類(lèi)之一, SIMD (單指令,多數(shù)據(jù))通常用于描述類(lèi)似 GPUs 的體系結(jié)構(gòu)。但是 SIMD 和 SIMT 之間有一個(gè)微妙但重要的區(qū)別。在 SIMD 體系結(jié)構(gòu)中,同一個(gè)指令中有多個(gè)并行操作。 SIMD 通常使用帶有向量寄存器和執(zhí)行單元的處理器來(lái)實(shí)現(xiàn);標(biāo)量線程發(fā)出以 SIMD 方式執(zhí)行的向量指令。在 SIMT 體系結(jié)構(gòu)中,多線程向任意數(shù)據(jù)發(fā)出通用指令,而不是單線程發(fā)出應(yīng)用于數(shù)據(jù)向量的向量指令。
SIMT 對(duì)于可編程性的好處使得 NVIDIA 的 GPU 架構(gòu)師為這種架構(gòu)命名,而不是將其描述為 SIMD 。 NVIDIA GPUs 使用 SIMT 執(zhí)行 32 個(gè)并行線程的 warp ,這使得每個(gè)線程能夠訪問(wèn)自己的寄存器,從不同的地址加載和存儲(chǔ),并遵循不同的控制流路徑。 CUDA 編譯器和 GPU 一起工作,以確保 warp 的線程盡可能頻繁地一起執(zhí)行相同的指令序列,從而最大限度地提高性能。
雖然通過(guò) warp 執(zhí)行獲得的高性能發(fā)生在場(chǎng)景后面,但是許多 CUDA 程序可以通過(guò)顯式 warp 級(jí)編程獲得更高的性能。并行程序通常使用集體通信操作,例如并行縮減和掃描。 CUDA C ++通過(guò)提供扭曲級(jí)基元和合作群集合來(lái)支持這樣的集合運(yùn)算。合作組 collectives ( 在上一篇文章中描述過(guò) )是在本文關(guān)注的 warp 原語(yǔ)之上實(shí)現(xiàn)的。
使用 shfl _ down _ sync ()進(jìn)行扭曲級(jí)別并行減少的一部分。
清單 1 顯示了一個(gè)使用 warp 級(jí)別原語(yǔ)的示例。它使用 __shfl_down_sync() 執(zhí)行樹(shù)縮減來(lái)計(jì)算扭曲中每個(gè)線程持有的 val 變量的總和。在第一個(gè)環(huán)的末尾, val 包含第一個(gè)線程的和。
__match_all_sync
活動(dòng)掩碼查詢(xún):返回一個(gè) 32 位掩碼,指示扭曲中的哪些線程與當(dāng)前正在執(zhí)行的線程處于活動(dòng)狀態(tài)。
__activemask
線程同步:同步扭曲中的線程并提供內(nèi)存邊界。
__syncwarp
請(qǐng)看
同步數(shù)據(jù)交換
每個(gè)“同步數(shù)據(jù)交換”原語(yǔ)在一個(gè) warp 中的一組線程之間執(zhí)行一個(gè)集體操作。例如,清單 2 顯示了其中的三個(gè)。調(diào)用 __shfl_sync() 或 __shfl_down_sync() 的每個(gè)線程都從同一個(gè) warp 中的線程接收數(shù)據(jù),而調(diào)用 __ballot_sync() 的每個(gè)線程都會(huì)接收一個(gè)位掩碼,該掩碼表示 warp 中為謂詞參數(shù)傳遞真值的所有線程。
int __shfl_sync(unsigned mask, int val, int src_line, int width=warpSize); int __shfl_down_sync(unsigned mask, int var, unsigned detla, int width=warpSize); int __ballot_sync(unsigned mask, int predicate);
參與調(diào)用每個(gè)原語(yǔ)的線程集是使用 32 位掩碼指定的,這是這些原語(yǔ)的第一個(gè)參數(shù)。所有參與線程必須同步,集體操作才能正常工作。因此,如果線程尚未同步,這些原語(yǔ)將首先同步線程。
一個(gè)常見(jiàn)的問(wèn)題是“對(duì)于mask
參數(shù),我應(yīng)該使用什么?”. 可以將遮罩視為扭曲中應(yīng)參與集體操作的線程集。這組線程由程序邏輯決定,通??梢酝ㄟ^(guò)程序流中早期的某些分支條件來(lái)計(jì)算。以清單 1 中的縮減代碼為例。假設(shè)我們要計(jì)算一個(gè)數(shù)組input[],
的所有元素的總和,該數(shù)組的大小NUM_ELEMENTS
小于線程塊中的線程數(shù)。我們可以使用清單 3 中的方法。
unsigned mask = __ballot_sync(FULL_MASK, threadIdx.x < NUM_ELEMENTS); if (threadIdx.x < NUM_ELEMENTS) { val = input[threadIdx.x]; for (int offset = 16; offset > 0; offset /= 2) val += __shfl_down_sync(mask, val, offset); … }
代碼使用條件thread.idx.x < NUM_ELEMENTS
來(lái)確定線程是否將參與縮減。__ballot_sync()
用于計(jì)算__shfl_down_sync()
操作的成員掩碼。__ballot_sync()
本身使用FULL_MASK
(0xffffffff
表示 32 個(gè)線程),因?yàn)槲覀兗僭O(shè)所有線程都將執(zhí)行它。
在 Volta 和更高版本的 GPU 架構(gòu)中,數(shù)據(jù)交換原語(yǔ)可以用于線程發(fā)散的分支:在這種分支中, warp 中的一些線程采用不同于其他線程的路徑。清單 4 顯示了一個(gè)示例,其中一個(gè) warp 中的所有線程都從第 0 行的線程獲得val
的值。偶數(shù)和奇數(shù)編號(hào)的線程采用if
語(yǔ)句的不同分支。
if (threadIdx.x % 2) { val += __shfl_sync(FULL_MASK, val, 0); … } else { val += __shfl_sync(FULL_MASK, val, 0); … }
在最新(和將來(lái) )的 Volta 的 GPU 上,您可以運(yùn)行使用 warp 同步原語(yǔ)的庫(kù)函數(shù),而不必?fù)?dān)心函數(shù)是否在線程發(fā)散分支中被調(diào)用。
活動(dòng)掩碼查詢(xún)
__activemask() 返回調(diào)用扭曲中所有當(dāng)前活動(dòng)線程的 32 位 unsigned int 掩碼。換句話(huà)說(shuō),它顯示了在其 warp 中的線程也在執(zhí)行相同的 __activemask() 的調(diào)用線程。這對(duì)于我們稍后解釋的:機(jī)會(huì)扭曲級(jí)編程”技術(shù)以及調(diào)試和理解程序行為非常有用。
但是,正確使用 __activemask() 很重要。清單 5 說(shuō)明了一個(gè)不正確的用法。代碼嘗試執(zhí)行與清單 4 中所示相同的總和縮減,但是它在分支內(nèi)部使用了 __activemask() ,而不是在分支之前使用 __ballot_sync() 來(lái)計(jì)算掩碼。這是不正確的,因?yàn)檫@將導(dǎo)致部分和而不是總和。 CUDA 執(zhí)行模型并不能保證將分支連接在一起的所有線程將一起執(zhí)行 __activemask() 。正如我們將要解釋的那樣,不能保證隱式鎖步驟的執(zhí)行。
// // Incorrect use of __activemask() // if (threadIdx.x < NUM_ELEMENTS) { unsigned mask = __activemask(); val = input[threadIdx.x]; for (int offset = 16; offset > 0; offset /= 2) val += __shfl_down_sync(mask, val, offset); … }
翹曲同步
當(dāng) warp 中的線程需要執(zhí)行比數(shù)據(jù)交換原語(yǔ)提供的更復(fù)雜的通信或集體操作時(shí),可以使用 __syncwarp() 原語(yǔ)來(lái)同步 warp 中的線程。它類(lèi)似于 __syncthreads() 原語(yǔ)(同步線程塊中的所有線程),但粒度更細(xì)。
void __syncwarp(unsigned mask=FULL_MASK);
__syncwarp()
原語(yǔ)使執(zhí)行線程等待,直到mask
中指定的所有線程都執(zhí)行了__syncwarp()
(使用相同的mask
),然后再繼續(xù)執(zhí)行。它還提供了一個(gè)記憶柵欄,允許線程在調(diào)用原語(yǔ)之前和之后通過(guò)內(nèi)存進(jìn)行通信。
清單 6 顯示了一個(gè)在 warp 中的線程之間混亂矩陣元素所有權(quán)的示例。
float val = get_value(…); __shared__ float smem[4][8]; // 0 1 2 3 4 5 6 7 // 8 9 10 11 12 13 14 15 // 16 17 18 19 20 21 22 23 // 24 25 26 27 28 29 30 31 int x1 = threadIdx.x % 8; int y1 = threadIdx.x / 8; // 0 4 8 12 16 20 24 28 // 1 5 10 13 17 21 25 29 // 2 6 11 14 18 22 26 30 // 3 7 12 15 19 23 27 31 int x2= threadIdx.x / 4; int y2 = threadIdx.x % 4; smem[y1][x1] = val; __syncwarp(); val = smem[y2][x2]; use(val);
假設(shè)使用了一維線程塊(即 threadIdx . y 始終為 0 )。在代碼的開(kāi)頭,一個(gè) warp 中的每個(gè)線程都擁有一個(gè) 4 × 8 矩陣的元素,該矩陣具有行主索引。換句話(huà)說(shuō),第 0 車(chē)道擁有[0][0]
車(chē)道,第 1 車(chē)道擁有[0][1]
。每個(gè)線程將其值存儲(chǔ)到共享內(nèi)存中 4 × 8 數(shù)組的相應(yīng)位置。然后使用__syncwarp()
來(lái)確保在每個(gè)線程從數(shù)組中的一個(gè)轉(zhuǎn)置位置讀取數(shù)據(jù)之前,所有線程都完成了存儲(chǔ)。最后, warp 中的每一個(gè)線程都擁有一個(gè)矩陣元素,列主索引為: lane0 擁有[0][0]
, lane1 擁有[1][0]
。
確保__syncwarp()
將共享內(nèi)存讀寫(xiě)分開(kāi),以避免爭(zhēng)用情況。清單 7 演示了共享內(nèi)存中樹(shù)和縮減的錯(cuò)誤用法。在每?jī)蓚€(gè)__syncwarp()
調(diào)用之間有一個(gè)共享內(nèi)存讀取,然后是共享內(nèi)存寫(xiě)入。 CUDA 編程模型不能保證所有的讀操作都會(huì)在所有的寫(xiě)操作之前執(zhí)行,因此存在競(jìng)爭(zhēng)條件。
unsigned tid = threadIdx.x; // Incorrect use of __syncwarp() shmem[tid] += shmem[tid+16]; __syncwarp(); shmem[tid] += shmem[tid+8]; __syncwarp(); shmem[tid] += shmem[tid+4]; __syncwarp(); shmem[tid] += shmem[tid+2]; __syncwarp(); shmem[tid] += shmem[tid+1]; __syncwarp();
清單 8 通過(guò)插入額外的__syncwarp()
調(diào)用修復(fù)了競(jìng)爭(zhēng)條件。 CUDA 編譯器可以在最終生成的代碼中省略一些同步指令,這取決于目標(biāo)體系結(jié)構(gòu)(例如,在預(yù)伏打體系結(jié)構(gòu)上)。
unsigned tid = threadIdx.x; int v = 0; v += shmem[tid+16]; __syncwarp(); shmem[tid] = v; __syncwarp(); v += shmem[tid+8]; __syncwarp(); shmem[tid] = v; __syncwarp(); v += shmem[tid+4]; __syncwarp(); shmem[tid] = v; __syncwarp(); v += shmem[tid+2]; __syncwarp(); shmem[tid] = v; __syncwarp(); v += shmem[tid+1]; __syncwarp(); shmem[tid] = v;
在最新的 Volta (和 future ) GPUs 上,也可以在線程發(fā)散分支中使用 __syncwarp() 來(lái)同步兩個(gè)分支的線程,但是一旦它們從原語(yǔ)返回,線程就會(huì)再次發(fā)散。請(qǐng)參見(jiàn)清單 13 中的示例。
機(jī)會(huì)主義翹曲水平編程
正如我們?cè)谕綌?shù)據(jù)交換一節(jié)中所示,在同步數(shù)據(jù)交換原語(yǔ)中使用的成員關(guān)系 mask 通常是在程序流中的分支條件之前計(jì)算的。在許多情況下,程序需要沿著程序流傳遞掩碼;例如,在函數(shù)內(nèi)部使用扭曲級(jí)原語(yǔ)時(shí),作為函數(shù)參數(shù)。如果要在庫(kù)函數(shù)內(nèi)使用 warp 級(jí)編程,但不能更改函數(shù)接口,則這可能很困難。
有些計(jì)算可以使用碰巧一起執(zhí)行的任何線程。我們可以使用一種稱(chēng)為機(jī)會(huì)主義翹曲級(jí)別編程的技術(shù),如下例所示。
// increment the value at ptr by 1 and return the old value __device__ int atomicAggInc(int *ptr) { int mask = __match_any_sync(__activemask(), (unsigned long long)ptr); int leader = __ffs(mask) – 1; // select a leader int res; if(lane_id() == leader) // leader does the update res = atomicAdd(ptr, __popc(mask)); res = __shfl_sync(mask, res, leader); // get leader’s old value return res + __popc(mask & ((1 << lane_id()) – 1)); //compute old value }
atomicAggInc() 以原子方式將 ptr 指向的值遞增 1 并返回舊值。它使用 atomicAdd() 函數(shù),這可能會(huì)引發(fā)爭(zhēng)用。為了減少爭(zhēng)用, atomicAggInc 用 per-warp atomicAdd() 替換了 per-thread atomicAdd() 操作。第 4 行中的 __activemask() 在 warp 中查找將要執(zhí)行原子操作的線程集。[zx7]的傳入線程具有相同的值,這些線程的[zx7]與[ez3]的值相同。每個(gè)組選擇一個(gè)引導(dǎo)線程(第 5 行),該線程為整個(gè)組執(zhí)行 atomicAdd() (第 8 行)。每個(gè)線程從 atomicAdd() 返回的前導(dǎo)(第 9 行)獲取舊值。第 10 行計(jì)算并返回當(dāng)前線程調(diào)用函數(shù)而不是 atomicAggInc 時(shí)從 atomicInc() 獲得的舊值。
隱式 Warp 同步編程是不安全的
CUDA 版本 9 。 0 之前的工具箱提供了一個(gè)(現(xiàn)在是遺留的) warp 級(jí)別基本體版本。與 CUDA 9 原語(yǔ)相比,傳統(tǒng)原語(yǔ)不接受 mask 參數(shù)。例如, int __any(int predicate) 是 int __any_sync(unsigned mask, int predicate) 的舊版本。
如前所述, mask 參數(shù)指定扭曲中必須參與原語(yǔ)的線程集。如果掩碼指定的線程在執(zhí)行過(guò)程中尚未同步,則新基元將執(zhí)行扭曲線程級(jí)內(nèi)同步。
傳統(tǒng)的 warp 級(jí)別原語(yǔ)不允許程序員指定所需的線程,也不執(zhí)行同步。因此,必須參與翹曲級(jí)別操作的線程不是由 CUDA 程序顯式表示的。這樣一個(gè)程序的正確性取決于隱式 warp 同步行為,這種行為可能從一個(gè)硬件體系結(jié)構(gòu)改變到另一個(gè),從一個(gè) CUDA 工具包版本到另一個(gè)(例如,由于編譯器優(yōu)化的變化),甚至從一個(gè)運(yùn)行時(shí)執(zhí)行到另一個(gè)。這種隱式 warp 同步編程是不安全的,可能無(wú)法正常工作。
例如,在下面的代碼中,假設(shè) warp 中的所有 32 個(gè)線程一起執(zhí)行第 2 行。第 4 行的 if 語(yǔ)句導(dǎo)致線程發(fā)散,奇數(shù)線程在第 5 行調(diào)用 foo() ,偶數(shù)線程在第 8 行調(diào)用 bar() 。
// Assuming all 32 threads in a warp execute line 1 together. assert(__ballot(1) == FULL_MASK); int result; if (thread_id % 2) { result = foo(); } else { result = bar(); } unsigned ballot_result = __ballot(result);
CUDA 編譯器和硬件將嘗試在第 10 行重新聚合線程,以獲得更好的性能。但這一重新收斂是不保證的。因此,ballot_result
可能不包含來(lái)自所有 32 個(gè)線程的投票結(jié)果。
在__ballot()
之前的第 10 行調(diào)用新的__syncwarp()
原語(yǔ),如清單 11 所示,也不能解決這個(gè)問(wèn)題。這又是隱式翹曲同步編程。它假設(shè)同一個(gè)扭曲中的線程一旦同步,將保持同步,直到下一個(gè)線程發(fā)散分支為止。盡管這通常是真的,但在 CUDA 編程模型中并不能保證它。
__syncwarp(); unsigned ballot_result = __ballot(result);
正確的修復(fù)方法是使用清單 12 中的__ballot_sync()
。
unsigned ballot_result = __ballot_sync(FULL_MASK, result);
一個(gè)常見(jiàn)的錯(cuò)誤是假設(shè)在舊的 warp 級(jí)別原語(yǔ)之前和/或之后調(diào)用__syncwarp()
在功能上等同于調(diào)用原語(yǔ)的sync
版本。例如,__syncwarp(); v = __shfl(0); __syncwarp();
與__shfl_sync(FULL_MASK, 0)
相同嗎?答案是否定的,有兩個(gè)原因。首先,如果在線程發(fā)散分支中使用序列,那么__shfl(0)
不會(huì)由所有線程一起執(zhí)行。清單 13 顯示了一個(gè)示例。第 3 行和第 7 行的__syncwarp()
將確保在執(zhí)行第 4 行或第 8 行之前, warp 中的所有線程都會(huì)調(diào)用foo()
。一旦線程離開(kāi)__syncwarp()
,奇數(shù)線程和偶數(shù)線程將再次發(fā)散。因此,第 4 行的__shfl(0)
將得到一個(gè)未定義的值,因?yàn)楫?dāng)?shù)?4 行執(zhí)行時(shí),第 0 行將不活動(dòng)。__shfl_sync(FULL_MASK, 0)
可以在線程發(fā)散的分支中使用,沒(méi)有這個(gè)問(wèn)題。
v = foo(); if (threadIdx.x % 2) { __syncwarp(); v = __shfl(0); // L3 will get undefined result because lane 0 __syncwarp(); // is not active when L3 is executed. L3 and L6 } else { // will execute divergently. __syncwarp(); v = __shfl(0); __syncwarp(); }
第二,即使所有線程一起調(diào)用序列, CUDA 執(zhí)行模型也不能保證線程在離開(kāi)__syncwarp()
后保持收斂,如清單 14 所示。不能保證隱式鎖步驟的執(zhí)行。請(qǐng)記住,線程收斂只在顯式同步的扭曲級(jí)別原語(yǔ)中得到保證。
assert(__activemask() == FULL_MASK); // assume this is true __syncwarp(); assert(__activemask() == FULL_MASK); // this may fail
因?yàn)槭褂盟鼈兛赡軙?huì)導(dǎo)致不安全的程序,所以從 CUDA 9 。 0 開(kāi)始就不推薦使用舊的 warp 級(jí)別原語(yǔ)。
更新舊版曲速級(jí)編程
如果您的程序使用舊的 warp 級(jí)原語(yǔ)或任何形式的隱式 warp 同步編程(例如在沒(méi)有同步的 warp 線程之間通信),您應(yīng)該更新代碼以使用原語(yǔ)的 sync 版本。您可能還需要重新構(gòu)造代碼以使用 Cooperative Groups ,這提供了更高級(jí)別的抽象以及諸如多塊同步等新功能。
使用翹曲級(jí)別原語(yǔ)最棘手的部分是找出要使用的成員掩碼。我們希望以上幾節(jié)能給你一個(gè)好主意,從哪里開(kāi)始,注意什么。以下是建議列表:
不要只使用 FULL_MASK (即對(duì)于 32 個(gè)線程使用 0xffffffff )作為 mask 值。如果不是所有的線程都能根據(jù)程序邏輯到達(dá)原語(yǔ),那么使用 FULL_MASK 可能會(huì)導(dǎo)致程序掛起。
不要只使用 __activemask() 作為掩碼值。 __activemask() 告訴您調(diào)用函數(shù)時(shí)哪些線程會(huì)收斂,這可能與您希望在集合操作中的情況不同。
分析程序邏輯并理解成員資格要求。根據(jù)程序邏輯提前計(jì)算掩碼。
如果您的程序執(zhí)行機(jī)會(huì)主義 warp 同步編程,請(qǐng)使用“ detective ”函數(shù),如 __activemask() 和 __match_all_sync() 來(lái)找到正確的掩碼。
使用 __syncwarp() 來(lái)分離與內(nèi)部扭曲相關(guān)的操作。不要假設(shè)執(zhí)行鎖步。
最后一個(gè)訣竅。如果您現(xiàn)有的 CUDA 程序在 Volta architecture GPUs 上給出了不同的結(jié)果,并且您懷疑差異是由 Volta 新的獨(dú)立線程調(diào)度 引起的,它可能會(huì)改變翹曲同步行為,您可能需要使用 nvcc 選項(xiàng) -arch=compute_60 -code=sm_70 重新編譯程序。這樣的編譯程序選擇使用 Pascal 的線程調(diào)度。當(dāng)有選擇地使用時(shí),它可以幫助更快地確定罪魁禍?zhǔn)啄K,允許您更新代碼以避免隱式 warp 同步編程。
Volta 獨(dú)立的線程調(diào)度允許交叉執(zhí)行來(lái)自不同分支的語(yǔ)句。這使得執(zhí)行細(xì)粒度并行算法成為可能,其中 warp 中的線程可以同步和通信。
關(guān)于作者
Yuan Lin 是 NVIDIA 編譯團(tuán)隊(duì)的首席工程師。他對(duì)所有使程序更高效、編程更高效的技術(shù)感興趣。在加入 NVIDIA 之前,他是 Sun Microsystems 的一名高級(jí)職員工程師。
Vinod Grover 是 CUDA C ++編譯器團(tuán)隊(duì) NVIDIA 的主管。在此之前,他曾在微軟和太陽(yáng)微系統(tǒng)公司擔(dān)任各種研究、工程和管理職務(wù)。
審核編輯:郭婷
-
NVIDIA
+關(guān)注
關(guān)注
14文章
4816瀏覽量
102632 -
gpu
+關(guān)注
關(guān)注
27文章
4632瀏覽量
128442 -
CUDA
+關(guān)注
關(guān)注
0文章
121瀏覽量
13570
發(fā)布評(píng)論請(qǐng)先 登錄
相關(guān)推薦
評(píng)論