Cooperative Groups
Cooperative Groups（協同組）是CUDA 9.0引入的一個新概念，主要用於跨線程塊（block）的同步。為使用Cooperative Groups，我們需要包含頭文件#include <cooperative_groups.h>，同時需要cooperative_groups命名空間。
簡介
在CUDA 9.0之前，CUDA僅支持線程塊內的同步，CUDA提供了2個原語操作：__syncthreads()函數用於同步同一線程塊內的所有線程，以及__syncwarp(unsigned mask=0xffffffff)函數用於同步線程束內的線程。
附1：
由於__syncthreads()函數要求整個線程塊內的所有線程都得到達該同步點方能繼續執行，也就是說同一線程塊的if條件必須都相同，否則程序將會被掛起或產生意想不到的結果。為避免此問題，CUDA提供了如下三個函數用於評估if條件的預測值：
int __syncthreads_count(int predicate);
該函數在__syncthreads()函數基礎上增加了返回predicate值非0的線程的數目。
int __syncthreads_and(int predicate);該函數在
__syncthreads()函數基礎上，當且僅當塊內所有線程predicate值非0時返回一個非0值。
int __syncthreads_or(int predicate);該函數在
__syncthreads()函數基礎上，當且僅當塊內存在任意一個線程predicate值非0時返回一個非0值。
顯然，線程塊級的同步並不能滿足開發者的需求，在某些時候，開發者需要跨線程塊同步，針對此問題，CUDA 9.0推出了Cooperative Groups機制，用於線程塊內和跨線程塊的同步。該機制為開發者提供了自定義線程組的方式，並提供了相應的同步函數，同時還包括一個新的kernel啟動API（cudaLaunchCooperativeKernel），該API保證了Cooperative Groups同步的安全性。
塊內組
thread_block
Cooperative Groups引入了一個新的數據結構：thread_block，即線程塊。thread_block可以通過this_thread_block()進行獲取並初始化：
thread_block g = this_thread_block();
thread_block繼承自更廣義的線程組數據結構：thread_group 。thread_group 提供了如下函數：
void sync(); //同步組內的所有線程，這里g.sync()等價於__syncthreads()
unsigned size(); //獲取組內的線程數目
unsigned thread_rank(); //獲取線程的組內索引值（[0,size]）
bool is_valid(); //判斷本組是否違背了任何APIconstraints（API限制）
thread_block則提供如下特定線程塊函數：
dim3 group_index(); //網格grid內3維索引（block索引）
dim3 thread_index(); //塊block內3維索引（線程索引）
注意以上所有操作組內所有線程都要確保執行到，否則行為未定義。
相比__syncthreads()函數，使用g.sync()的好處在於避免了隱式同步隱患：
__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);
}
此時，當開發者調用他人編寫的sum函數時，不一定能發現sum中存在着同步，但當我們顯式傳參時情況就不一樣了：
__device__ int sum(const thread_group& g, int *x, int n)
{
// ...
g.sync()
return total;
}
__global__ void parallel_kernel(float *x)
{
// ...
// Entire thread block must call sum
sum(this_thread_block(), x, n);
}
tiled_partition
tiled_partition()函數用於將一個線程塊分解為多個小的協同線程組（tiled subgroups），比如說：
thread_block wholeBlock = this_thread_block(); //獲取線程塊
以下函數將線程塊分解為若干個大小為32的小線程組：
thread_group tile32 = tiled_partition(wholeBlock, 32);
甚至可以更深一步，將tile32分解為更小的若干個大小為4的小線程組：
thread_group tile4 = tiled_partition(tile32, 4);
注意：小線程組大小僅支持2的冪數且不大於32，也就是僅限於2,4,8,16,32
現在，我們通過如下操作就可以讓線程0,4,8,12,...（相對於wholeBlock 的索引）打印"Hello World"：
if (tile4.thread_rank() == 0) printf("Hello World\n");
Thread Block Tiles
CUDA還提供了thread_block_tile<>模版使得小線程組大小在編譯期就可以得到：
thread_block wholeBlock = this_thread_block(); //獲取線程塊
thread_block_tile<32> tile32 = tiled_partition<32>(wholeBlock);
thread_block_tile<4> tile4 = tiled_partition<4>(tile32);
Thread Block Tiles提供了如下成員函數用於協同同步：
.shfl() //等價__shfl_sync
.shfl_down() //等價__shfl_down_sync
.shfl_up() //等價__shfl_up_sync
.shfl_xor() //等價__shfl_xor_sync
.any() //等價__any_sync
.all() //等價__all_sync
.ballot() //等價__ballot_sync
.match_any() //等價__match_any_sync
.match_all() //等價__match_all_sync
注意相比通過tiled_partition()函數傳參動態設置線程組大小，通過tiled_partition<>模版靜態設置線程組大小使得開發者可以使用如上這些線程束同步函數，前者不能。
附2：
__shfl_sync系列指令（俗稱洗牌指令）用於在線程束中獲取指定線程的變量值，該操作會在mask（一般取0xffffffff，每個bit位代表每個線程id）指定的那些線程中同時執行（同一mask中的線程必須執行相同指令），每次移動4字節或8字節的數據，但若指定線程為非活躍線程，則結果未知。具體功能如下：T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize);
__shfl_sync指令返回索引為srcLane線程的var變量值，其中srcLane大小為[0,width)，類似的，width的值必須是2的冪數且不大於32。T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);
__shfl_up_sync指令返回索引為當前線程索引減去delta的值的線程的var值，若減去后的值小於0則不做任何操作（保持不變）。T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);
__shfl_down_sync指令返回索引為當前線程索引加上delta的值的線程的var值，若加后的值大於width則不做任何操作（保持不變）。T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=warpSize);
__shfl_xor_sync指令返回索引為當前線程索引按位異或laneMask后的值的線程的var值。注意若width值小於warpSize值，此時后面的線程可以訪問前面的線程組的值（獲取成功），但前面的線程不能訪問后面線程組的值（保持不變）。
附3：
__any_sync系列指令（俗稱投票指令）對線程束中的參與線程（同樣由mask指定）比較預測值predicate是否非零，並向所有參與的活躍線程廣播比較結果：int __all_sync(unsigned mask, int predicate);當線程束中所有參與線程的預測值predicate非零時返回一個非零值。
int __any_sync(unsigned mask, int predicate);當線程束中存在任意一個參與線程的預測值predicate非零時返回一個非零值。
unsigned __ballot_sync(unsigned mask, int predicate);若線程束中的第N個線程活躍且其預測值predicate非零時，設定返回值的第N個bit為1，否則為0。
unsigned __activemask();返回線程束內活躍線程組成的掩碼。若線程束中的第N個線程為活躍線程，則設定第N個bit為1，否則為0（注意已退出線程也是非活躍線程）。該指令不執行同步。
附4：
__match_any_sync系列指令對線程束的參與線程（同樣由mask指定）比較value值，並向所有參與線程廣播比較結果：unsigned int __match_any_sync(unsigned mask, T value);返回value值相同的那些線程組成的掩碼。
unsigned int __match_all_sync(unsigned mask, T value, int *pred);返回mask值若所有參與線程的value值都相同，否則返回0。此外前者的預測值pred還將被設定為true，否則為false。
Coalesced Groups
若同一線程束（warp）內的線程出現條件分化（通常由if語句導致），那么程序將序列化運行：既在執行某分支線程時停止其它分支線程的執行，直到所有分支執行完畢。我們稱正執行的活躍線程為coalesced thread，線程束內所有活躍線程組成的線程組即為coalesced groups，其可以通過coalesced_threads函數獲取：
coalesced_group active = coalesced_threads();
coalesced_group也是一類thread_group。
網格級同步
相比塊內組，Cooperative Groups最強大的能力在於跨線程塊同步，在CUDA 9.0之前，不同線程塊僅能在kernel執行結束時同步，現在開發者可以通過grid_group 結構執行網格級同步：
grid_group grid = this_grid();
grid.sync();
注意不同於傳統的<<<...>>>執行配置，網格級同步必須通過cudaLaunchCooperativeKernel API配置並啟動kernel：
cudaError_t cudaLaunchCooperativeKernel(
const T *func, //kernel函數指針
dim3 gridDim,
dim3 blockDim,
void **args, //kernel參數數組
size_t sharedMem = 0,
cudaStream_t stream = 0
)
注意為保證所有協同線程塊能安全的常駐GPU，gridDim和blockDim的值需要慎重考慮，開發者可以通過計算SM的最大活躍線程塊數目來最大化並行率：
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&numBlocksPerSm,
my_kernel,
numThreads,
0
);
// initialize, then launch
cudaLaunchCooperativeKernel(
(void*)my_kernel,
deviceProp.multiProcessorCount*numBlocksPerSm,
numThreads,
args
);
Cooperative Launch目前不支持任務搶占和調度，若一次啟動的block數超過了設備駐留的極限，則報錯too many blocks in cooperative launch cudaLaunchCooperativeKernel，此時你需要檢查一下啟動block數、使用的共享內存大小、使用的寄存器大小。相關問題見(https://bbs.gpuworld.cn/index.php?topic=73127.0)
除特殊的啟動函數外，網格同步還需要在編譯時開啟-rdc=true參數。
該功能僅支持計算能力6.0及以上的設備，在不確定GPU是否支持網格同步時，開發者可以通過如下方式查詢：
int pi=0;
cuDevice dev;
cuDeviceGet(&dev,0) // get handle to device 0
cuDeviceGetAttribute(&pi, CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH, dev);
當pi值為1時表明設備0支持網格級同步。
多設備同步
類似網格級同步，多設備同步通過multi_grid_group 結構執行：
multi_grid_group multi_grid = this_multi_grid();
multi_grid.sync();
並通過cudaLaunchCooperativeKernelMultiDevice API配置並啟動kernel：
cudaError_t cudaLaunchCooperativeKernelMultiDevice(
CUDA_LAUNCH_PARAMS *launchParamsList,
unsigned int numDevices,
unsigned int flags = 0
);
其中CUDA_LAUNCH_PARAMS結構體定義如下：
typedef struct CUDA_LAUNCH_PARAMS_st {
CUfunction function;
unsigned int gridDimX;
unsigned int gridDimY;
unsigned int gridDimZ;
unsigned int blockDimX;
unsigned int blockDimY;
unsigned int blockDimZ;
unsigned int sharedMemBytes;
CUstream hStream;
void **kernelParams;
} CUDA_LAUNCH_PARAMS;
當開發者使用該API需要注意如下幾點：
該API將確保一個launch操作是原子的，例如當API調用成功時，相應數目的線程塊在所有指定設備上launch成功。
對於所有設備，該API調用的kernel函數必須是相同的。
同一設備上的
launchParamsList參數必須是相同的。所有設備的計算能力必須是相同的（major and minor versions）。
對於所有設備，配置的網格大小（
gridDim）、塊大小（blockDim）和每個網格的共享內存大小必須是相同的。自定義的
__device__,__constant__,__managed__全局變量在每個設備上都是獨立實例化的，因此需要開發者對該類變量賦初值。
類似的，該功能僅支持計算能力6.0及以上設備，可以通過CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH 查詢。