協(xié)作組編程模型的特點(diǎn)及應(yīng)用

C.1. Introduction

Cooperative Groups 是 CUDA 9 中引入的 CUDA 編程模型的擴(kuò)展，用于組織通信線程組。協(xié)作組允許開(kāi)發(fā)人員表達(dá)線程通信的粒度，幫助他們表達(dá)更豐富、更有效的并行分解。

從歷史上看，CUDA 編程模型為同步協(xié)作線程提供了一個(gè)單一、簡(jiǎn)單的構(gòu)造：線程塊的所有線程之間的屏障，如使用__syncthreads()內(nèi)部函數(shù)實(shí)現(xiàn)的那樣。但是，程序員希望以其他粒度定義和同步線程組，以“集體”組范圍功能接口的形式實(shí)現(xiàn)更高的性能、設(shè)計(jì)靈活性和軟件重用。為了表達(dá)更廣泛的并行交互模式，許多面向性能的程序員已經(jīng)求助于編寫自己的臨時(shí)和不安全的原語(yǔ)來(lái)同步單個(gè) warp 中的線程，或者跨運(yùn)行在單個(gè) GPU 上的線程塊集。雖然實(shí)現(xiàn)的性能改進(jìn)通常很有價(jià)值，但這導(dǎo)致了越來(lái)越多的脆弱代碼集合，隨著時(shí)間的推移和跨 GPU 架構(gòu)的不同，這些代碼的編寫、調(diào)整和維護(hù)成本很高。合作組通過(guò)提供安全且面向未來(lái)的機(jī)制來(lái)啟用高性能代碼來(lái)解決這個(gè)問(wèn)題。

C.2. What’s New in CUDA 11.0

• 使用網(wǎng)格范圍的組不再需要單獨(dú)編譯，并且同步該組的速度現(xiàn)在提高了30%。此外，我們?cè)?a href="http://srfitnesspt.com/article/zt/" target="_blank">最新的 Windows 平臺(tái)上啟用了協(xié)作啟動(dòng)，并在 MPS 下運(yùn)行時(shí)增加了對(duì)它們的支持。
• grid_group現(xiàn)在可以轉(zhuǎn)換為thread_group。
• 線程塊切片和合并組的新集合：reduce和memcpy_async。
• 線程塊切片和合并組的新分區(qū)操作：labeled_pa??rtition和binary_partition。
• 新的 API，meta_group_rank和meta_group_size，它們提供有關(guān)導(dǎo)致創(chuàng)建該組的分區(qū)的信息。
• 線程塊tile現(xiàn)在可以在類型中編碼其父級(jí)，這允許對(duì)發(fā)出的代碼進(jìn)行更好的編譯時(shí)優(yōu)化。
• 接口更改：grid_group必須在聲明時(shí)使用this_grid()構(gòu)造。默認(rèn)構(gòu)造函數(shù)被刪除。

注意：在此版本中，我們正朝著要求 C++11 提供新功能的方向發(fā)展。在未來(lái)的版本中，所有現(xiàn)有 API 都需要這樣做。

C.3. Programming Model Concept

協(xié)作組編程模型描述了 CUDA 線程塊內(nèi)和跨線程塊的同步模式。 它為應(yīng)用程序提供了定義它們自己的線程組的方法，以及同步它們的接口。 它還提供了強(qiáng)制執(zhí)行某些限制的新啟動(dòng) API，因此可以保證同步正常工作。 這些原語(yǔ)在 CUDA 內(nèi)啟用了新的協(xié)作并行模式，包括生產(chǎn)者-消費(fèi)者并行、機(jī)會(huì)并行和整個(gè)網(wǎng)格的全局同步。

合作組編程模型由以下元素組成：

• 表示協(xié)作線程組的數(shù)據(jù)類型；
• 獲取由 CUDA 啟動(dòng) API 定義的隱式組的操作（例如，線程塊）；
• 將現(xiàn)有群體劃分為新群體的集體；
• 用于數(shù)據(jù)移動(dòng)和操作的集體算法（例如memcpy_async、reduce、scan）；
• 同步組內(nèi)所有線程的操作；
• 檢查組屬性的操作；
• 公開(kāi)低級(jí)別、特定于組且通常是硬件加速的操作的集合。

協(xié)作組中的主要概念是對(duì)象命名作為其中一部分的線程集的對(duì)象。 這種將組表示為一等程序?qū)ο蟮姆绞礁倪M(jìn)了軟件組合，因?yàn)榧虾瘮?shù)可以接收表示參與線程組的顯式對(duì)象。 該對(duì)象還明確了程序員的意圖，從而消除了不合理的架構(gòu)假設(shè)，這些假設(shè)會(huì)導(dǎo)致代碼脆弱、對(duì)編譯器優(yōu)化的不良限制以及與新一代 GPU 的更好兼容性。

為了編寫高效的代碼，最好使用專門的組（通用會(huì)失去很多編譯時(shí)優(yōu)化），并通過(guò)引用打算以某種協(xié)作方式使用這些線程的函數(shù)來(lái)傳遞這些組對(duì)象。

合作組需要 CUDA 9.0 或更高版本。 要使用合作組，請(qǐng)包含頭文件：

// Primary header is compatible with pre-C++11, collective algorithm headers require C++11
#include 
// Optionally include for memcpy_async() collective
#include 
// Optionally include for reduce() collective
#include 
// Optionally include for inclusive_scan() and exclusive_scan() collectives
#include 

并使用合作組命名空間：

using namespace cooperative_groups;
// Alternatively use an alias to avoid polluting the namespace with collective algorithms
namespace cg = cooperative_groups;

可以使用 nvcc 以正常方式編譯代碼，但是如果您希望使用memcpy_async、reduce或scan功能并且您的主機(jī)編譯器的默認(rèn)不是 C++11 或更高版本，那么您必須添加--std=c++11到命令行。

C.3.1. Composition Example

為了說(shuō)明組的概念，此示例嘗試執(zhí)行塊范圍的求和。 以前，編寫此代碼時(shí)對(duì)實(shí)現(xiàn)存在隱藏的約束：

__device__ int sum(int *x, int n) {
    // ...
    __syncthreads();
    return total;
}

__global__ void parallel_kernel(float *x) {
    // ...
    // Entire thread block must call sum
    sum(x, n);
}

線程塊中的所有線程都必須到達(dá)__syncthreads()屏障，但是，對(duì)于可能想要使用sum(...)的開(kāi)發(fā)人員來(lái)說(shuō)，這個(gè)約束是隱藏的。 對(duì)于合作組，更好的編寫方式是：

__device__ int sum(const thread_block& g, int *x, int n) {
    // ...
    g.sync()
    return total;
}

__global__ void parallel_kernel(...) {
    // ...
    // Entire thread block must call sum
    thread_block tb = this_thread_block();
    sum(tb, x, n);
    // ...
}

C.4. Group Types

C.4.1. Implicit Groups

隱式組代表內(nèi)核的啟動(dòng)配置。不管你的內(nèi)核是如何編寫的，它總是有一定數(shù)量的線程、塊和塊尺寸、單個(gè)網(wǎng)格和網(wǎng)格尺寸。另外，如果使用多設(shè)備協(xié)同啟動(dòng)API，它可以有多個(gè)網(wǎng)格（每個(gè)設(shè)備一個(gè)網(wǎng)格）。這些組為分解為更細(xì)粒度的組提供了起點(diǎn)，這些組通常是硬件加速的，并且更專門針對(duì)開(kāi)發(fā)人員正在解決的問(wèn)題。

盡管您可以在代碼中的任何位置創(chuàng)建隱式組，但這樣做很危險(xiǎn)。為隱式組創(chuàng)建句柄是一項(xiàng)集體操作——組中的所有線程都必須參與。如果組是在并非所有線程都到達(dá)的條件分支中創(chuàng)建的，則可能導(dǎo)致死鎖或數(shù)據(jù)損壞。出于這個(gè)原因，建議您預(yù)先為隱式組創(chuàng)建一個(gè)句柄（盡可能早，在任何分支發(fā)生之前）并在整個(gè)內(nèi)核中使用該句柄。出于同樣的原因，必須在聲明時(shí)初始化組句柄（沒(méi)有默認(rèn)構(gòu)造函數(shù)），并且不鼓勵(lì)復(fù)制構(gòu)造它們。

C.4.1.1. Thread Block Group

任何 CUDA 程序員都已經(jīng)熟悉某一組線程：線程塊。 Cooperative Groups 擴(kuò)展引入了一個(gè)新的數(shù)據(jù)類型thread_block，以在內(nèi)核中明確表示這個(gè)概念。

class thread_block;

thread_block g = this_thread_block();

公開(kāi)成員函數(shù)：

示例:

/// Loading an integer from global into shared memory
__global__ void kernel(int *globalInput) {
    __shared__ int x;
    thread_block g = this_thread_block();
    // Choose a leader in the thread block
    if (g.thread_rank() == 0) {
        // load from global into shared for all threads to work with
        x = (*globalInput);
    }
    // After loading data into shared memory, you want to synchronize
    // if all threads in your thread block need to see it
    g.sync(); // equivalent to __syncthreads();
}

注意：組中的所有線程都必須參與集體操作，否則行為未定義。

相關(guān)：thread_block數(shù)據(jù)類型派生自更通用的thread_group數(shù)據(jù)類型，可用于表示更廣泛的組類。

C.4.1.2. Grid Group

該組對(duì)象表示在單個(gè)網(wǎng)格中啟動(dòng)的所有線程。 除了sync()之外的 API 始終可用，但要能夠跨網(wǎng)格同步，您需要使用協(xié)作啟動(dòng) API。

class grid_group;
grid_group g = this_grid();

公開(kāi)成員函數(shù)：

C.4.1.3. Multi Grid Group

該組對(duì)象表示跨設(shè)備協(xié)作組啟動(dòng)的所有設(shè)備啟動(dòng)的所有線程。 與grid.group不同，所有 API 都要求您使用適當(dāng)?shù)膯?dòng) API。

class multi_grid_group;

通過(guò)一下方式構(gòu)建:

// Kernel must be launched with the cooperative multi-device API
multi_grid_group g = this_multi_grid();

公開(kāi)成員函數(shù)：

C.4.2. Explicit Groups

C.4.2.1. Thread Block Tile

tile組的模板版本，其中模板參數(shù)用于指定tile的大小 – 在編譯時(shí)已知這一點(diǎn)，有可能實(shí)現(xiàn)更優(yōu)化的執(zhí)行。

template 
class thread_block_tile;

通過(guò)以下構(gòu)建:

template 
_CG_QUALIFIER thread_block_tile tiled_partition(const ParentT& g),>

Size必須是 2 的冪且小于或等于 32。

ParentT是從其中劃分該組的父類型。 它是自動(dòng)推斷的，但是 void 的值會(huì)將此信息存儲(chǔ)在組句柄中而不是類型中。

公開(kāi)成員函數(shù)：

注意：

shfl、shfl_up、shfl_down 和 shfl_xor函數(shù)在使用 C++11 或更高版本編譯時(shí)接受任何類型的對(duì)象。 這意味著只要滿足以下約束，就可以對(duì)非整數(shù)類型進(jìn)行shuffle ：

• 符合普通可復(fù)制的條件，即is_trivially_copyable::value == true
• sizeof(T) <= 32

示例:

/// The following code will create two sets of tiled groups, of size 32 and 4 respectively:
/// The latter has the provenance encoded in the type, while the first stores it in the handle
thread_block block = this_thread_block();
thread_block_tile<32> tile32 = tiled_partition<32>(block);
thread_block_tile<4, thread_block> tile4 = tiled_partition<4>(block);

注意：這里使用的是 thread_block_tile 模板化數(shù)據(jù)結(jié)構(gòu)，并且組的大小作為模板參數(shù)而不是參數(shù)傳遞給 tiled_partition 調(diào)用。

C.4.2.1.1. Warp-Synchronous Code Pattern

開(kāi)發(fā)人員可能擁有他們之前對(duì) warp 大小做出隱含假設(shè)并圍繞該數(shù)字進(jìn)行編碼的 warp 同步代碼。 現(xiàn)在這需要明確指定。

__global__ void cooperative_kernel(...) {
    // obtain default "current thread block" group
    thread_block my_block = this_thread_block();

    // subdivide into 32-thread, tiled subgroups
    // Tiled subgroups evenly partition a parent group into
    // adjacent sets of threads - in this case each one warp in size
    auto my_tile = tiled_partition<32>(my_block);

    // This operation will be performed by only the
    // first 32-thread tile of each block
    if (my_tile.meta_group_rank() == 0) {
        // ...
        my_tile.sync();
    }
}

C.4.2.1.2. Single thread group

可以從 this_thread 函數(shù)中獲取代表當(dāng)前線程的組：

thread_block_tile<1> this_thread();

以下memcpy_asyncAPI 使用thread_group將int元素從源復(fù)制到目標(biāo)：

#include 
#include 

cooperative_groups::memcpy_async(cooperative_groups::this_thread(), dest, src, sizeof(int));

可以在使用?cuda::pipeline的單階段異步數(shù)據(jù)拷貝和使用?cuda::pipeline的多階段異步數(shù)據(jù)拷貝部分中找到使用this_thread執(zhí)行異步復(fù)制的更詳細(xì)示例。

C.4.2.1.3. Thread Block Tile of size larger than 32

使用cooperative_groups::experimental命名空間中的新API 可以獲得大小為64、128、256 或512的thread_block_tile。 要使用它，_CG_ABI_EXPERIMENTAL必須在源代碼中定義。 在分區(qū)之前，必須為thread_block_tile保留少量?jī)?nèi)存。 這可以使用必須駐留在共享或全局內(nèi)存中的cooperative_groups::experimental::block_tile_memory結(jié)構(gòu)模板來(lái)完成。

template 
struct block_tile_memory;

TileCommunicationSize確定為集體操作保留多少內(nèi)存。 如果對(duì)大于指定通信大小的大小類型執(zhí)行此類操作，則集合可能涉及多次傳輸并需要更長(zhǎng)的時(shí)間才能完成。

MaxBlockSize指定當(dāng)前線程塊中的最大線程數(shù)。 此參數(shù)可用于最小化僅以較小線程數(shù)啟動(dòng)的內(nèi)核中block_tile_memory的共享內(nèi)存使用量。

然后這個(gè)block_tile_memory需要被傳遞到cooperative_groups::experimental::this_thread_block，允許將生成的thread_block劃分為大小大于32的tile。this_thread_block接受block_tile_memory參數(shù)的重載是一個(gè)集體操作，必須與所有線程一起調(diào)用 線程塊。 返回的線程塊可以使用experimental::tiled_partition函數(shù)模板進(jìn)行分區(qū)，該模板接受與常規(guī)tiled_partition相同的參數(shù)。

#define _CG_ABI_EXPERIMENTAL // enable experimental API

__global__ void cooperative_kernel(...) {
    // reserve shared memory for thread_block_tile usage.
    __shared__ experimental::block_tile_memory<4, 256> shared;
    thread_block thb = experimental::this_thread_block(shared);

    auto tile = experimental::tiled_partition<128>(thb);

    // ...
}

公開(kāi)成員函數(shù):

C.4.2.2. Coalesced Groups

在 CUDA 的 SIMT 架構(gòu)中，在硬件級(jí)別，多處理器以 32 個(gè)一組的線程執(zhí)行線程，稱為 warp。 如果應(yīng)用程序代碼中存在依賴于數(shù)據(jù)的條件分支，使得 warp 中的線程發(fā)散，那么 warp 會(huì)串行執(zhí)行每個(gè)分支，禁用不在該路徑上的線程。 在路徑上保持活動(dòng)的線程稱為合并。 協(xié)作組具有發(fā)現(xiàn)和創(chuàng)建包含所有合并線程的組的功能。

通過(guò)coalesced_threads()構(gòu)造組句柄是伺機(jī)的(opportunistic)。 它在那個(gè)時(shí)間點(diǎn)返回一組活動(dòng)線程，并且不保證返回哪些線程（只要它們是活動(dòng)的）或者它們?cè)谡麄€(gè)執(zhí)行過(guò)程中保持合并（它們將被重新組合在一起以執(zhí)行一個(gè)集合，但之后可以再次發(fā)散）。

class coalesced_group;

通過(guò)以下重構(gòu):

coalesced_group active = coalesced_threads();

公開(kāi)成員函數(shù):

注意：shfl、shfl_up 和 shfl_down函數(shù)在使用 C++11 或更高版本編譯時(shí)接受任何類型的對(duì)象。 這意味著只要滿足以下約束，就可以對(duì)非整數(shù)類型進(jìn)行洗牌：

• 符合普通可復(fù)制的條件，即is_trivially_copyable::value == true
• sizeof(T) <= 32

示例:

/// Consider a situation whereby there is a branch in the
/// code in which only the 2nd, 4th and 8th threads in each warp are
/// active. The coalesced_threads() call, placed in that branch, will create (for each
/// warp) a group, active, that has three threads (with
/// ranks 0-2 inclusive).
__global__ void kernel(int *globalInput) {
    // Lets say globalInput says that threads 2, 4, 8 should handle the data
    if (threadIdx.x == *globalInput) {
        coalesced_group active = coalesced_threads();
        // active contains 0-2 inclusive
        active.sync();
    }
}

C.4.2.2.1. Discovery Pattern

通常，開(kāi)發(fā)人員需要使用當(dāng)前活動(dòng)的線程集。 不對(duì)存在的線程做任何假設(shè)，而是開(kāi)發(fā)人員使用碰巧存在的線程。 這可以在以下“在warp中跨線程聚合原子增量”示例中看到（使用正確的 CUDA 9.0 內(nèi)在函數(shù)集編寫）：

{
    unsigned int writemask = __activemask();
    unsigned int total = __popc(writemask);
    unsigned int prefix = __popc(writemask & __lanemask_lt());
    // Find the lowest-numbered active lane
    int elected_lane = __ffs(writemask) - 1;
    int base_offset = 0;
    if (prefix == 0) {
        base_offset = atomicAdd(p, total);
    }
    base_offset = __shfl_sync(writemask, base_offset, elected_lane);
    int thread_offset = prefix + base_offset;
    return thread_offset;
}

這可以用Cooperative Groups重寫如下：

{
    cg::coalesced_group g = cg::coalesced_threads();
    int prev;
    if (g.thread_rank() == 0) {
        prev = atomicAdd(p, g.num_threads());
    }
    prev = g.thread_rank() + g.shfl(prev, 0);
    return prev;
}

C.5. Group Partitioning

C.5.1. tiled_partition

template 
thread_block_tile tiled_partition(const ParentT& g);

thread_group tiled_partition(const thread_group& parent, unsigned int tilesz);,>

tiled_partition方法是一種集體操作，它將父組劃分為一維、行主序的子組平鋪。 總共將創(chuàng)建((size(parent)/tilesz)子組，因此父組大小必須能被Size整除。允許的父組是thread_block或thread_block_tile。

該實(shí)現(xiàn)可能導(dǎo)致調(diào)用線程在恢復(fù)執(zhí)行之前等待，直到父組的所有成員都調(diào)用了該操作。功能僅限于本地硬件大小，1/2/4/8/16/32和cg::size(parent)必須大于size參數(shù)。cooperative_groups::experimental命名空間的實(shí)驗(yàn)版本支持64/128/256/512大小。

Codegen 要求：計(jì)算能力 3.5 最低，C++11 用于大于 32 的size

示例:

/// The following code will create a 32-thread tile
thread_block block = this_thread_block();
thread_block_tile<32> tile32 = tiled_partition<32>(block);

我們可以將這些組中的每一個(gè)分成更小的組，每個(gè)組的大小為 4 個(gè)線程：

auto tile4 = tiled_partition<4>(tile32);
// or using a general group
// thread_group tile4 = tiled_partition(tile32, 4);

例如，如果我們要包含以下代碼行：

if (tile4.thread_rank()==0) printf(“Hello from tile4 rank 0\n”);

那么該語(yǔ)句將由塊中的每四個(gè)線程打?。好總€(gè) tile4 組中排名為 0 的線程，它們對(duì)應(yīng)于塊組中排名為 0、4、8、12.. 的那些線程。

C.5.2. labeled_partition

coalesced_group labeled_partition(const coalesced_group& g, int label);
template 
coalesced_group labeled_partition(const thread_block_tile& g, int label);

labeled_partition方法是一種集體操作，它將父組劃分為一維子組，線程在這些子組中合并。 該實(shí)現(xiàn)將評(píng)估條件標(biāo)簽并將具有相同標(biāo)簽值的線程分配到同一組中。

該實(shí)現(xiàn)可能會(huì)導(dǎo)致調(diào)用線程在恢復(fù)執(zhí)行之前等待直到父組的所有成員都調(diào)用了該操作。

注意：此功能仍在評(píng)估中，將來(lái)可能會(huì)略有變化。

Codegen 要求：計(jì)算能力 7.0 最低，C++11

C.5.3. binary_partition

coalesced_group binary_partition(const coalesced_group& g, bool pred);
template 
coalesced_group binary_partition(const thread_block_tile& g, bool pred);

binary_partition()方法是一種集體操作，它將父組劃分為一維子組，線程在其中合并。 該實(shí)現(xiàn)將評(píng)估predicate并將具有相同值的線程分配到同一組中。 這是labeled_partition()的一種特殊形式，其中label只能是0 或1。

該實(shí)現(xiàn)可能會(huì)導(dǎo)致調(diào)用線程在恢復(fù)執(zhí)行之前等待直到父組的所有成員都調(diào)用了該操作。

注意：此功能仍在評(píng)估中，將來(lái)可能會(huì)略有變化。

Codegen 要求：計(jì)算能力 7.0 最低，C++11

示例:

/// This example divides a 32-sized tile into a group with odd
/// numbers and a group with even numbers
_global__ void oddEven(int *inputArr) {
    cg::thread_block cta = cg::this_thread_block();
    cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);

    // inputArr contains random integers
    int elem = inputArr[cta.thread_rank()];
    // after this, tile32 is split into 2 groups,
    // a subtile where elem&1 is true and one where its false
    auto subtile = cg::binary_partition(tile32, (elem & 1));
}

C.6. Group Collectives

C.6.1. Synchronization

C.6.1.1. sync

cooperative_groups::sync(T& group);

sync同步組中指定的線程。T可以是任何現(xiàn)有的組類型，因?yàn)樗鼈兌贾С滞健?如果組是grid_group或multi_grid_group，則內(nèi)核必須已使用適當(dāng)?shù)膮f(xié)作啟動(dòng) API 啟動(dòng)。

C.6.2. Data Transfer

C.6.2.1. memcpy_async

memcpy_async是一個(gè)組范圍的集體memcpy，它利用硬件加速支持從全局到共享內(nèi)存的非阻塞內(nèi)存事務(wù)。給定組中命名的一組線程，memcpy_async將通過(guò)單個(gè)管道階段傳輸指定數(shù)量的字節(jié)或輸入類型的元素。此外，為了在使用memcpy_asyncAPI 時(shí)獲得最佳性能，共享內(nèi)存和全局內(nèi)存都需要 16 字節(jié)對(duì)齊。需要注意的是，雖然在一般情況下這是一個(gè)memcpy，但只有當(dāng)源(source)是全局內(nèi)存而目標(biāo)是共享內(nèi)存并且兩者都可以通過(guò) 16、8 或 4 字節(jié)對(duì)齊來(lái)尋址時(shí)，它才是異步的。異步復(fù)制的數(shù)據(jù)只能在調(diào)用wait或wait_prior之后讀取，這表明相應(yīng)階段已完成將數(shù)據(jù)移動(dòng)到共享內(nèi)存。

必須等待所有未完成的請(qǐng)求可能會(huì)失去一些靈活性（但會(huì)變得簡(jiǎn)單）。為了有效地重疊數(shù)據(jù)傳輸和執(zhí)行，重要的是能夠在等待和操作請(qǐng)求N時(shí)啟動(dòng)N+1 memcpy_async請(qǐng)求。為此，請(qǐng)使用memcpy_async并使用基于集體階段的wait_priorAPI 等待它.有關(guān)詳細(xì)信息，請(qǐng)參閱wait 和 wait_prior。

用法1:

template 
void memcpy_async(
  const TyGroup &group,
  TyElem *__restrict__ _dst,
  const TyElem *__restrict__ _src,
  const TyShape &shape
);

執(zhí)行shape字節(jié)的拷貝

用法2:

template 
void memcpy_async(
  const TyGroup &group,
  TyElem *__restrict__ dst,
  const TyDstLayout &dstLayout,
  const TyElem *__restrict__ src,
  const TySrcLayout &srcLayout
);

執(zhí)行min(dstLayout, srcLayout)元素的拷貝。 如果布局的類型為cuda::aligned_size_t，則兩者必須指定相同的對(duì)齊方式。

勘誤表

CUDA 11.1 中引入的具有 src 和 dst 輸入布局的memcpy_asyncAPI 期望布局以元素而不是字節(jié)形式提供。 元素類型是從TyElem推斷出來(lái)的，大小為sizeof(TyElem)。 如果使用cuda::aligned_size_t類型作為布局，指定的元素個(gè)數(shù)乘以sizeof(TyElem)必須是 N 的倍數(shù)，建議使用std::byte或char作為元素類型。

如果副本的指定形狀或布局是cuda::aligned_size_t類型，則將保證至少為min(16, N)。 在這種情況下，dst 和 src 指針都需要與 N 個(gè)字節(jié)對(duì)齊，并且復(fù)制的字節(jié)數(shù)需要是 N 的倍數(shù)。

Codegen 要求：最低計(jì)算能力 3.5，異步計(jì)算能力 8.0，C++11

需要包含collaborative_groups/memcpy_async.h頭文件。

示例:

/// This example streams elementsPerThreadBlock worth of data from global memory
/// into a limited sized shared memory (elementsInShared) block to operate on.
#include 
#include 

namespace cg = cooperative_groups;

__global__ void kernel(int* global_data) {
    cg::thread_block tb = cg::this_thread_block();
    const size_t elementsPerThreadBlock = 16 * 1024;
    const size_t elementsInShared = 128;
    __shared__ int local_smem[elementsInShared];

    size_t copy_count;
    size_t index = 0;
    while (index < elementsPerThreadBlock) {
        cg::memcpy_async(tb, local_smem, elementsInShared, global_data + index, elementsPerThreadBlock - index);
        copy_count = min(elementsInShared, elementsPerThreadBlock - index);
        cg::wait(tb);
        // Work with local_smem
        index += copy_count;
    }
}

C.6.2.2. wait and wait_prior

template 
void wait(TyGroup & group);

template 
void wair_prior(TyGroup & group);

wait和wait_prior集合同步指定的線程和線程塊，直到所有未完成的memcpy_async請(qǐng)求（在等待的情況下）或第一個(gè)NumStages（在 wait_prior 的情況下）完成。

Codegen 要求：最低計(jì)算能力 3.5，異步計(jì)算能力 8.0，C++11

需要包含collaborative_groups/memcpy_async.h 頭文件。

示例:

/// This example streams elementsPerThreadBlock worth of data from global memory
/// into a limited sized shared memory (elementsInShared) block to operate on in
/// multiple (two) stages. As stage N is kicked off, we can wait on and operate on stage N-1.
#include 
#include 

namespace cg = cooperative_groups;

__global__ void kernel(int* global_data) {
    cg::thread_block tb = cg::this_thread_block();
    const size_t elementsPerThreadBlock = 16 * 1024 + 64;
    const size_t elementsInShared = 128;
    __align__(16) __shared__ int local_smem[2][elementsInShared];
    int stage = 0;
    // First kick off an extra request
    size_t copy_count = elementsInShared;
    size_t index = copy_count;
    cg::memcpy_async(tb, local_smem[stage], elementsInShared, global_data, elementsPerThreadBlock - index);
    while (index < elementsPerThreadBlock) {
        // Now we kick off the next request...
        cg::memcpy_async(tb, local_smem[stage ^ 1], elementsInShared, global_data + index, elementsPerThreadBlock - index);
        // ... but we wait on the one before it
        cg::wait_prior<1>(tb);

        // Its now available and we can work with local_smem[stage] here
        // (...)
        //

        // Calculate the amount fo data that was actually copied, for the next iteration.
        copy_count = min(elementsInShared, elementsPerThreadBlock - index);
        index += copy_count;

        // A cg::sync(tb) might be needed here depending on whether
        // the work done with local_smem[stage] can release threads to race ahead or not
        // Wrap to the next stage
        stage ^= 1;
    }
    cg::wait(tb);
    // The last local_smem[stage] can be handled here

C.6.3. Data manipulation

C.6.3.1. reduce

template 
auto reduce(const TyGroup& group, TyArg&& val, TyOp&& op) -> decltype(op(val, val));

reduce對(duì)傳入的組中指定的每個(gè)線程提供的數(shù)據(jù)執(zhí)行歸約操作。這利用硬件加速（在計(jì)算 80 及更高的設(shè)備上）進(jìn)行算術(shù)加法、最小或最大操作以及邏輯 AND、OR、或 XOR，以及在老一代硬件上提供軟件替代支持(fallback)。只有 4B 類型由硬件加速。

group：有效的組類型是coalesced_group和thread_block_tile。

val：滿足以下要求的任何類型：

• 符合普通可復(fù)制的條件，即is_trivially_copyable::value == true
• sizeof(TyArg) <= 32
• 對(duì)給定的函數(shù)對(duì)象具有合適的算術(shù)或比較運(yùn)算符。

op：將提供具有整數(shù)類型的硬件加速的有效函數(shù)對(duì)象是plus()、less()、greater()、bit_and()、bit_xor()、bit_or()。這些必須構(gòu)造，因此需要TyVal模板參數(shù)，即plus()。Reduce還支持可以使用operator()調(diào)用的lambda和其他函數(shù)對(duì)象

Codegen 要求：計(jì)算能力 3.5 最低，計(jì)算能力 8.0 用于硬件加速，C++11。

需要包含collaborative_groups/reduce.h 頭文件。

示例:

#include 
#include 
namespace cg=cooperative_groups;

/// The following example accepts input in *A and outputs a result into *sum
/// It spreads the data within the block, one element per thread
#define blocksz 256
__global__ void block_reduce(const int *A, int *sum) {
    __shared__ int reduction_s[blocksz];

    cg::thread_block cta = cg::this_thread_block();
    cg::thread_block_tile<32> tile = cg::tiled_partition<32>(cta);

    const int tid = cta.thread_rank();
    int beta = A[tid];
    // reduce across the tile
    // cg::plus allows cg::reduce() to know it can use hardware acceleration for addition
    reduction_s[tid] = cg::reduce(tile, beta, cg::plus());
    // synchronize the block so all data is ready
    cg::sync(cta);
    // single leader accumulates the result
    if (cta.thread_rank() == 0) {
        beta = 0;
        for (int i = 0; i < blocksz; i += tile.num_threads()) {
            beta += reduction_s[i];
        }
        sum[blockIdx.x] = beta;
    }

C.6.3.2. Reduce Operators

下面是一些可以用reduce完成的基本操作的函數(shù)對(duì)象的原型

namespace cooperative_groups {
  template 
  struct cg::plus;

  template 
  struct cg::less;

  template 
  struct cg::greater;

  template 
  struct cg::bit_and;

  template 
  struct cg::bit_xor;

  template 
  struct cg::bit_or;
}

Reduce僅限于在編譯時(shí)可用于實(shí)現(xiàn)的信息。 因此，為了利用 CC 8.0 中引入的內(nèi)在函數(shù)，cg::命名空間公開(kāi)了幾個(gè)鏡像硬件的功能對(duì)象。 這些對(duì)象看起來(lái)與 C++ STL 中呈現(xiàn)的對(duì)象相似，除了less/greater。 與 STL 有任何差異的原因在于，這些函數(shù)對(duì)象旨在實(shí)際反映硬件內(nèi)聯(lián)函數(shù)的操作。

功能說(shuō)明：

• cg::plus：接受兩個(gè)值并使用operator +返回兩者之和。
• cg::less: 接受兩個(gè)值并使用operator 返回較小的值。 這不同之處在于返回較低的值而不是布爾值。
• cg::greater：接受兩個(gè)值并使用operator <返回較大的值。 這不同之處在于返回更大的值而不是布爾值。
• cg::bit_and：接受兩個(gè)值并返回operator &的結(jié)果。
• cg::bit_xor：接受兩個(gè)值并返回operator ^的結(jié)果。
• cg::bit_or：接受兩個(gè)值并返回operator |的結(jié)果。

示例:

{
    // cg::plus is specialized within cg::reduce and calls __reduce_add_sync(...) on CC 8.0+
    cg::reduce(tile, (int)val, cg::plus());

    // cg::plus fails to match with an accelerator and instead performs a standard shuffle based reduction
    cg::reduce(tile, (float)val, cg::plus());

    // While individual components of a vector are supported, reduce will not use hardware intrinsics for the following
    // It will also be necessary to define a corresponding operator for vector and any custom types that may be used
    int4 vec = {...};
    cg::reduce(tile, vec, cg::plus())

    // Finally lambdas and other function objects cannot be inspected for dispatch
    // and will instead perform shuffle based reductions using the provided function object.
    cg::reduce(tile, (int)val, [](int l, int r) -> int {return l + r;});
}

template 
auto inclusive_scan(const TyGroup& group, TyVal&& val, TyFn&& op) -> decltype(op(val, val));

template 
TyVal inclusive_scan(const TyGroup& group, TyVal&& val);

template 
auto exclusive_scan(const TyGroup& group, TyVal&& val, TyFn&& op) -> decltype(op(val, val));

template 
TyVal exclusive_scan(const TyGroup& group, TyVal&& val);

#include 
#include 
#include 
namespace cg = cooperative_groups;

__global__ void kernel() {
    auto thread_block = cg::this_thread_block();
    auto tile = cg::tiled_partition<8>(thread_block);
    unsigned int val = cg::inclusive_scan(tile, tile.thread_rank());
    printf("%u: %u\n", tile.thread_rank(), val);
}

/*  prints for each group:
    0: 0
    1: 1
    2: 3
    3: 6
    4: 10
    5: 15
    6: 21
    7: 28
*/

#include 
#include 
namespace cg = cooperative_groups;

// Buffer partitioning is static to make the example easier to follow,
// but any arbitrary dynamic allocation scheme can be implemented by replacing this function.
__device__ int calculate_buffer_space_needed(cg::thread_block_tile<32>& tile) {
    return tile.thread_rank() % 2 + 1;
}

__device__ int my_thread_data(int i) {
    return i;
}

__global__ void kernel() {
    __shared__ int buffer_used;
    extern __shared__ int buffer[];
    auto thread_block = cg::this_thread_block();
    auto tile = cg::tiled_partition<32>(thread_block);

    buffer_used = 0;
    thread_block.sync();

    // each thread calculates buffer size it needs and its offset within the allocation
    int buf_needed = calculate_buffer_space_needed(tile);
    int buf_offset = cg::exclusive_scan(tile, buf_needed);

    // last thread in the tile allocates buffer space with an atomic operation
    int alloc_offset = 0;
    if (tile.thread_rank() == tile.num_threads() - 1) {
        alloc_offset = atomicAdd(&buffer_used, buf_offset + buf_needed);
    }
    // that thread shares the allocation start with other threads in the tile
    alloc_offset = tile.shfl(alloc_offset, tile.num_threads() - 1);
    buf_offset += alloc_offset;

    // each thread fill its part of the buffer with thread specific data
    for (int i = 0 ; i < buf_needed ; ++i) {
        buffer[buf_offset + i] = my_thread_data(i);
    }

    // buffer is {0, 0, 1, 0, 0, 1 ...};
}

grid_group grid = this_grid();
grid.sync();

int device = 0;
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
// initialize, then launch
cudaLaunchCooperativeKernel((void*)my_kernel, deviceProp.multiProcessorCount, numThreads, args);

/// This will launch a grid that can maximally fill the GPU, on the default stream with kernel arguments
int numBlocksPerSm = 0;
 // Number of threads my_kernel will be launched with
int numThreads = 128;
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocksPerSm, my_kernel, numThreads, 0);
// launch
void *kernelArgs[] = { /* add kernel args */ };
dim3 dimBlock(numThreads, 1, 1);
dim3 dimGrid(deviceProp.multiProcessorCount*numBlocksPerSm, 1, 1);
cudaLaunchCooperativeKernel((void*)my_kernel, dimGrid, dimBlock, kernelArgs);

int dev = 0;
int supportsCoopLaunch = 0;
cudaDeviceGetAttribute(&supportsCoopLaunch, cudaDevAttrCooperativeLaunch, dev);

cudaDeviceProp deviceProp;
cudaGetDeviceCount(&numGpus);

// Per device launch parameters
cudaLaunchParams *launchParams = (cudaLaunchParams*)malloc(sizeof(cudaLaunchParams) * numGpus);
cudaStream_t *streams = (cudaStream_t*)malloc(sizeof(cudaStream_t) * numGpus);

// The kernel arguments are copied over during launch
// Its also possible to have individual copies of kernel arguments per device, but
// the signature and name of the function/kernel must be the same.
void *kernelArgs[] = { /* Add kernel arguments */ };

for (int i = 0; i < numGpus; i++) {
    cudaSetDevice(i);
    // Per device stream, but its also possible to use the default NULL stream of each device
    cudaStreamCreate(&streams[i]);
    // Loop over other devices and cudaDeviceEnablePeerAccess to get a faster barrier implementation
}
// Since all devices must be of the same compute capability and have the same launch configuration
// it is sufficient to query device 0 here
cudaGetDeviceProperties(&deviceProp[i], 0);
dim3 dimBlock(numThreads, 1, 1);
dim3 dimGrid(deviceProp.multiProcessorCount, 1, 1);
for (int i = 0; i < numGpus; i++) {
    launchParamsList[i].func = (void*)my_kernel;
    launchParamsList[i].gridDim = dimGrid;
    launchParamsList[i].blockDim = dimBlock;
    launchParamsList[i].sharedMem = 0;
    launchParamsList[i].stream = streams[i];
    launchParamsList[i].args = kernelArgs;
}
cudaLaunchCooperativeKernelMultiDevice(launchParams, numGpus);

multi_grid_group multi_grid = this_multi_grid();
multi_grid.sync();

int dev = 0;
int supportsMdCoopLaunch = 0;
cudaDeviceGetAttribute(&supportsMdCoopLaunch, cudaDevAttrCooperativeMultiDeviceLaunch, dev);