C.1. Introduction

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

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

C.2. What’s New in CUDA 11.0

• 使用網(wǎng)格范圍的組不再需要單獨編譯，并且同步該組的速度現(xiàn)在提高了30%。此外，我們在最新的 Windows 平臺上啟用了協(xié)作啟動，并在 MPS 下運行時增加了對它們的支持。
• 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)在可以在類型中編碼其父級，這允許對發(fā)出的代碼進(jìn)行更好的編譯時優(yōu)化。
• 接口更改：grid_group必須在聲明時使用this_grid()構(gòu)造。默認(rèn)構(gòu)造函數(shù)被刪除。

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

C.3. Programming Model Concept

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

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

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

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

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

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

// 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功能并且您的主機編譯器的默認(rèn)不是 C++11 或更高版本，那么您必須添加--std=c++11到命令行。

C.3.1. Composition Example

為了說明組的概念，此示例嘗試執(zhí)行塊范圍的求和。 以前，編寫此代碼時對實現(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()屏障，但是，對于可能想要使用sum(...)的開發(fā)人員來說，這個約束是隱藏的。 對于合作組，更好的編寫方式是：

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

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

C.4.1.1. Thread Block Group

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

class thread_block;

thread_block g = this_thread_block();

公開成員函數(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

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

class grid_group;
grid_group g = this_grid();

公開成員函數(shù)：

C.4.1.3. Multi Grid Group

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

class multi_grid_group;

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

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

公開成員函數(shù)：

C.4.2. Explicit Groups

C.4.2.1. Thread Block Tile

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

template 
class thread_block_tile;

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

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

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

ParentT是從其中劃分該組的父類型。 它是自動推斷的，但是 void 的值會將此信息存儲在組句柄中而不是類型中。

公開成員函數(shù)：

注意：

shfl、shfl_up、shfl_down 和 shfl_xor函數(shù)在使用 C++11 或更高版本編譯時接受任何類型的對象。 這意味著只要滿足以下約束，就可以對非整數(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

開發(fā)人員可能擁有他們之前對 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保留少量內(nèi)存。 這可以使用必須駐留在共享或全局內(nèi)存中的cooperative_groups::experimental::block_tile_memory結(jié)構(gòu)模板來完成。

template 
struct block_tile_memory;

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

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

然后這個block_tile_memory需要被傳遞到cooperative_groups::experimental::this_thread_block，允許將生成的thread_block劃分為大小大于32的tile。this_thread_block接受block_tile_memory參數(shù)的重載是一個集體操作，必須與所有線程一起調(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);

    // ...
}

公開成員函數(shù):

C.4.2.2. Coalesced Groups

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

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

class coalesced_group;

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

coalesced_group active = coalesced_threads();

公開成員函數(shù):

注意：shfl、shfl_up 和 shfl_down函數(shù)在使用 C++11 或更高版本編譯時接受任何類型的對象。 這意味著只要滿足以下約束，就可以對非整數(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

通常，開發(fā)人員需要使用當(dāng)前活動的線程集。 不對存在的線程做任何假設(shè)，而是開發(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。

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

Codegen 要求：計算能力 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);

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

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”);

那么該語句將由塊中的每四個線程打?。好總€ tile4 組中排名為 0 的線程，它們對應(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方法是一種集體操作，它將父組劃分為一維子組，線程在這些子組中合并。 該實現(xiàn)將評估條件標(biāo)簽并將具有相同標(biāo)簽值的線程分配到同一組中。

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

注意：此功能仍在評估中，將來可能會略有變化。

Codegen 要求：計算能力 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()方法是一種集體操作，它將父組劃分為一維子組，線程在其中合并。 該實現(xiàn)將評估predicate并將具有相同值的線程分配到同一組中。 這是labeled_partition()的一種特殊形式，其中label只能是0 或1。

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

注意：此功能仍在評估中，將來可能會略有變化。

Codegen 要求：計算能力 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)有的組類型，因為它們都支持同步。 如果組是grid_group或multi_grid_group，則內(nèi)核必須已使用適當(dāng)?shù)膮f(xié)作啟動 API 啟動。

C.6.2. Data Transfer

C.6.2.1. memcpy_async

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

必須等待所有未完成的請求可能會失去一些靈活性（但會變得簡單）。為了有效地重疊數(shù)據(jù)傳輸和執(zhí)行，重要的是能夠在等待和操作請求N時啟動N+1 memcpy_async請求。為此，請使用memcpy_async并使用基于集體階段的wait_priorAPI 等待它.有關(guān)詳細(xì)信息，請參閱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，則兩者必須指定相同的對齊方式。

勘誤表

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

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

Codegen 要求：最低計算能力 3.5，異步計算能力 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請求（在等待的情況下）或第一個NumStages（在 wait_prior 的情況下）完成。

Codegen 要求：最低計算能力 3.5，異步計算能力 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對傳入的組中指定的每個線程提供的數(shù)據(jù)執(zhí)行歸約操作。這利用硬件加速（在計算 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
• 對給定的函數(shù)對象具有合適的算術(shù)或比較運算符。

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

Codegen 要求：計算能力 3.5 最低，計算能力 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ù)對象的原型

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

功能說明：

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