NVIDIA GPUs 以 SIMT （單指令，多線程）方式執(zhí)行稱為 warps 的線程組。許多 CUDA 程序通過(guò)利用 warp 執(zhí)行來(lái)獲得高性能。在這個(gè)博客中，我們將展示如何使用 CUDA 9 中引入的原語(yǔ)，使您的 warp 級(jí)編程安全有效。

扭曲級(jí)別基本體

NVIDIA GPUs 和 CUDA 編程模型采用一種稱為 SIMT （單指令，多線程）的執(zhí)行模型。 SIMT 擴(kuò)展了計(jì)算機(jī)體系結(jié)構(gòu)的 弗林分類學(xué) ，它根據(jù)指令和數(shù)據(jù)流的數(shù)量描述了四類體系結(jié)構(gòu)。作為 Flynn 的四個(gè)類之一， SIMD （單指令，多數(shù)據(jù)）通常用于描述類似 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í)行樹縮減來(lái)計(jì)算扭曲中每個(gè)線程持有的 val 變量的總和。在第一個(gè)環(huán)的末尾， val 包含第一個(gè)線程的和。

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)掩碼查詢

__activemask（） 返回調(diào)用扭曲中所有當(dāng)前活動(dòng)線程的 32 位 unsigned int 掩碼。換句話說(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 中的線程。它類似于 __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 矩陣的元素，該矩陣具有行主索引。換句話說(shuō)，第 0 車道擁有[0][0]車道，第 1 車道擁有[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)存讀寫分開(kāi)，以避免爭(zhēng)用情況。清單 7 演示了共享內(nèi)存中樹和縮減的錯(cuò)誤用法。在每?jī)蓚€(gè)__syncwarp()調(diào)用之間有一個(gè)共享內(nèi)存讀取，然后是共享內(nèi)存寫入。 CUDA 編程模型不能保證所有的讀操作都會(huì)在所有的寫操作之前執(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í)行的任何線程。我們可以使用一種稱為機(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ù)。

審核編輯：郭婷