这是 CUDA 的 线程组织和协作扩展，属于 高级并行编程模型，它在普通的 block / warp 层次之上引入了 更灵活的线程分组（thread grouping）机制。

背景

在传统 CUDA 模型里：

• Warp：32 个线程的基本执行单元，硬件级别同步。

• Block：多个 warp 组成，block 内可以使用 __syncthreads() 同步。

• Grid：多个 block 组成，可以通过 global memory 通信。

问题：

• __syncthreads() 只能同步同一个 block 的线程。

• 跨 block 的同步非常受限，导致一些算法（如大规模归约、scan、softmax）难以优化。

Cooperative Groups (CG)

CUDA 7.0+ 引入 Cooperative Groups API，提供 灵活的线程分组和同步接口。

分组类型	描述	同步方式
thread_block	对应传统 block	sync()
coalesced_group	warp 内线程分组，按物理执行顺序	warp 内同步，可用 sync()
grid_group	整个 grid 的所有线程（仅限支持 cooperative kernel）	sync()，跨 block
tiled_partition	block 内的任意线程划分成 tile	tile 内同步

常用api

需要包含头文件 #include <cooperative_groups.h>

cooperative_groups::this_thread_block

namespace cooperative_groups {
    thread_block this_thread_block() noexcept;
}

返回类型：cooperative_groups::thread_block

它直接从当前线程上下文获取所在的 block 信息

• 返回一个 thread_block 对象，表示当前线程所在的 block。

• 你可以通过这个对象调用 block 内的各种方法

thread_block

• 类型定义在 cooperative_groups 命名空间

• 对应 CUDA 一个 block 的封装对象

• 提供 block 内线程信息访问和同步接口

• 面向对象的方式比直接用 threadIdx / __syncthreads() 更清晰

成员	类型/返回值	功能说明
dim_threads	dim3	Block 内线程的维度信息，对应原生 blockDim
num_threads()	int	Block 内线程总数，与 size() 相同，等于 blockDim.x * blockDim.y * blockDim.z
size()	int	Block 内线程总数，等同 num_threads()
group_dim()	dim3	线程组（block）在每个维度上的大小，也就是 block 的维度（x,y,z）
group_index()	dim3	Block 的索引（x,y,z），对应原生 blockIdx
id	dim3	Block 的索引，等同 group_index()（内部别名）
thread_rank()	int	Block 内线程的线性索引（0 ~ block 内总线程数-1），常用于一维化访问
thread_index()	dim3	Block 内线程在每个维度的索引（x,y,z），对应原生 threadIdx
thread_scope	枚举	描述线程所在的粒度级别（block/warp/tile 等），通常内部使用
_group_id	内部实现	用于标识该 block 在线程组层次中的 ID，一般不直接使用

成员	返回值	功能说明
sync()	void	Block 内所有线程同步，等效于 __syncthreads()
get_type()	枚举	返回线程组类型（thread_block 对应 block），通常内部使用

• 常用成员：

  • 线程编号/索引：thread_rank() / thread_index()

  • Block 大小信息：size() / num_threads() / dim_threads

  • Block 索引：group_index() / id

  • 同步：sync()

• 内部/实验性：

  • _group_id、thread_scope、get_type() 不推荐直接使用，主要给 CG 内部或 cluster 扩展用

cooperative_groups::coalesced_threads

coalesced_group coalesced_threads();

• 返回值：

  • 类型：cooperative_groups::coalesced_group

  • 表示 warp 内连续执行线程的一个封装对象，提供 warp 内通信、同步等接口。

coalesced_group

coalesced_group 是一个 warp 封装类，把所有底层 warp intrinsics (__all_sync, __any_sync, __shfl_sync, __ballot_sync 等)
封成了更易用、更可泛化的 C++ 接口。

成员	类型	功能 / 说明
sync()	void	warp 内线程同步（本质上通常是 no-op，因为 warp 内天然同步，但为了泛型编程提供统一接口）
thread_rank()	int	warp 内线程的线性编号（0~warpSize-1）
thread_index()	int	同 thread_rank()，便于统一接口
size()	int	warp 内线程数量（一般 = warpSize = 32）
thread_scope	enum	表示线程的作用域，warp 级别
_group_id	int	内部实现使用的 group ID，通常不直接访问

bool all(int predicate) const

功能：
判断 group 中所有线程的 predicate 是否都为真。
如果全部为真 → 返回 true，否则 false。

bool any(int predicate) const

功能：
判断 group 内是否 有任意线程 满足条件。

unsigned ballot(bool predicate) const

功能：
收集 group 内每个线程的 predicate 状态，返回一个 bitmask。

const char* get_type() const

功能：
返回 group 的类型描述字符串（调试用途）。

示例输出：

• "coalesced_group"

• "thread_block_tile" 等。

unsigned match_all(T value, bool &pred) const

功能：
判断 group 内所有线程的值是否都相等。
如果都相等，则 pred = true 并返回该值所在的掩码。

unsigned match_any(T value) const

功能：
返回与当前线程拥有相同 value 的线程掩码。

int meta_group_rank() const

功能：
返回当前 group 在上一级 group 中的编号。
例如：block 内多个 warp，每个 warp 的 meta_group_rank() 表示它是第几个 warp。

返回值：

• group 在上级 group（meta group）中的 rank。

int meta_group_size() const

功能：
返回上一级 group 中包含多少个子 group。
例如：一个 block 有 8 个 warp → 返回 8。

int num_threads() const

功能：
返回 group 内线程总数。
几乎等价于 size()。

template <typename T> T shfl(T var, int srcLane) const

功能：
在 group 内线程间直接交换变量。
当前线程会获得 srcLane 指定线程的 var。

template <typename T> T shfl_down(T var, unsigned delta) const

功能：
从当前线程的下方（编号更大的）线程获取数据。
通常用于规约。

template <typename T> T shfl_up(T var, unsigned delta) const

功能：
与上面的反向：从当前线程的上方（编号更小的）线程获取数据。

tiled_partition

template <unsigned int Size, typename ParentT>
thread_block_tile<Size, ParentT> tiled_partition(const ParentT& g)

把一个已有的线程组（比如 thread_block 或 coalesced_group）划分成若干个 大小为 TileSize 的子 group（tile），
每个 tile 里包含的线程数固定为 TileSize。

例如：

• block 有 128 个线程；

• 你调用 tiled_partition<32>(block);；

• 则这个 block 被划分为 4 个 tile（每个 32 个线程，对应 1 个 warp）。

每个线程在编译时只会知道自己属于哪个 tile，并能在 tile 内通信（shuffle、sync、reduce 等）。

api与coalesced_group基本类似

“群体算法”（Collective Algorithms）

这些算法本质上是“在一个线程组（group）内进行的并行操作”，
与 Thrust 的 STL 风格算法类似，但作用范围局限在 GPU 内部的一个 group 上。

#include <cooperative_groups/reduce.h>
#include <cooperative_groups/scan.h>
#include <cooperative_groups/memcpy_async.h>

分类	典型函数	功能
规约（Reduction）	reduce()	对 group 内所有线程的值进行求和、最值、乘积等聚合
前缀和（Scan）	inclusive_scan() / exclusive_scan()	对 group 内的值进行扫描（累加）
异步拷贝	memcpy_async() / wait()	在 group 内异步加载全局内存
广播 / shuffle	shfl() / broadcast()	在 group 内传递数据
同步	group.sync()	同步 group 内线程

reduce()

#include <cooperative_groups/reduce.h>
template <typename Group, typename T, typename BinaryOp>
__device__ T reduce(const Group& g, T value, BinaryOp op);

作用：
在指定线程组 g 内对每个线程的 value 进行聚合（如求和、求最大值等），返回结果（每个线程都返回同样的值）。

#include <cooperative_groups.h>
#include <cooperative_groups/reduce.h>
namespace cg = cooperative_groups;

__global__ void kernel(float* data) {
    cg::thread_block block = cg::this_thread_block();
    auto tile32 = cg::tiled_partition<32>(block);

    float x = data[block.thread_rank()];
    float sum = cg::reduce(tile32, x, cg::plus<float>());

    if (tile32.thread_rank() == 0)
        printf("Tile sum = %f\n", sum);
}

exclusive_scan() / inclusive_scan()

#include <cooperative_groups/scan.h>
template <typename Group, typename T, typename BinaryOp>
__device__ T exclusive_scan(const Group& g, T value, BinaryOp op);

template <typename Group, typename T, typename BinaryOp>
__device__ T inclusive_scan(const Group& g, T value, BinaryOp op);

作用：
对 group 内的所有线程执行前缀操作（前缀和、前缀最大值等）。

#include <cooperative_groups/scan.h>
namespace cg = cooperative_groups;

__global__ void prefix_sum(float *data) {
    cg::thread_block block = cg::this_thread_block();
    auto tile8 = cg::tiled_partition<8>(block);

    float x = data[block.thread_rank()];
    float prefix = cg::inclusive_scan(tile8, x, cg::plus<float>());

    data[block.thread_rank()] = prefix;
}

memcpy_async() / wait()

#include <cooperative_groups/memcpy_async.h>
template <typename Group>
__device__ void memcpy_async(const Group& g, void* dst, const void* src, size_t size);

template <typename Group>
__device__ void wait(const Group& g);

作用：
在 group 内发起异步数据传输（比如从 global memory 到 shared memory），
通常用于 Ampere 及更高架构的 GPU 上，用来隐藏内存访问延迟。

#include <cooperative_groups/memcpy_async.h>
namespace cg = cooperative_groups;

__global__ void async_copy_kernel(float* dst, const float* src) {
    extern __shared__ float smem[];

    cg::thread_block block = cg::this_thread_block();

    cg::memcpy_async(block, smem, src, 256 * sizeof(float)); // 异步拷贝
    cg::wait(block); // 等待拷贝完成

    // 使用 smem 中的数据
    smem[block.thread_rank()] *= 2;
}

常用的辅助函数和运算符

运算符	类型	作用
cg::plus<T>()	二元函数对象	加法
cg::maximum<T>()	二元函数对象	求最大值
cg::minimum<T>()	二元函数对象	求最小值
cg::multiplies<T>()	二元函数对象	乘法

这些类似于 thrust::plus / std::plus，用在 reduce、scan 等算法里。

规约碎碎念

到目前为止，已经见过CUB库的规约，Thrust库的规约和cooperative_groups库的规约了

CUB 库特点

• 极致性能（由 NVIDIA 工程师手写 PTX 优化）

• 支持泛型操作（不仅是 sum，可自定义 ReduceOp）

• 支持 warp、block、device 三层规约

• 自动处理边界、bank conflict、同步

Thrust 库特点

• 类似 C++ STL 算法接口

• 代码简洁

• 性能比手写略低（但仍然不错）

• 适合快速原型、科研、教学

cooperative_groups 库特点

• 可以定义 warp/block/tile***/多block group

• 支持 同步（sync）、规约（reduce）、广播（broadcast）

• 比 warp shuffle 更抽象、可移植

跨 block 同步

cooperative_groups 默认并不支持跨 block 的同步或规约。
要想真正实现 “跨 block” 的 group（grid 级同步），需要显式启用“cooperative launch”机制。

group 类型	说明	是否支持跨 block
thread_block	当前 block 内线程组	❌ 否
thread_block_tile<N>	block 内的 warp/tile 分组	❌ 否
grid_group	整个 grid 的所有线程组	✅ 是
multi_grid_group	多 grid（多 kernel）级组	✅ 是（rare）

跨 block 同步必须启用 cooperative launch

普通的 kernel 启动方式：

kernel<<<gridDim, blockDim>>>(...);

这种 无法保证 所有 block 同时在 GPU 上执行，因此不能安全地跨 block 同步。

要启用 跨 block 同步（即 grid_group），必须满足两个条件：

✅ 条件 1：使用 cooperative launch API

void *kernelArgs[] = { &input, &output };
cudaLaunchCooperativeKernel((void*)kernel, gridDim, blockDim, kernelArgs);

cudaError_t cudaLaunchCooperativeKernel(const void *func, dim3 gridDim, dim3 blockDim,
 void **args, size_t sharedMem, cudaStream_t stream)

✅ 条件 2：硬件支持

• Pascal (SM 6.0) 及以上 GPU 支持 cooperative launch。

• 驱动和设备属性必须满足：

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, device);
if (!prop.cooperativeLaunch) { /* 不支持跨block同步 */ }

示例代码

#ifndef __CUDACC__
#define __CUDACC__
#endif
#include <cuda_runtime.h>
#include <iostream>
#include <cooperative_groups.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

namespace cg = cooperative_groups;

__global__ void kernel(float *data, int N) {
    cg::grid_group grid = cg::this_grid();  // 获取整个 grid 的 group

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N)
        data[idx] += 1.0f;

    grid.sync();  // 所有 block 同步

    if (idx < N)
        data[idx] *= 2.0f;
}

int main() {
    int N = 64;
    int blocksize = 32;
    int gridsize  = 2;

    thrust::device_vector<float> d_vec(N, 1.0f);
    thrust::host_vector<float> h_vec(N);

    float *d_ptr = thrust::raw_pointer_cast(d_vec.data());

    // ✅ 检查设备是否支持 cooperative launch
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    if (!prop.cooperativeLaunch) {
        std::cerr << "Device does NOT support cooperative launch!" << std::endl;
        return 1;
    }

    // ✅ 检查 grid 是否能完全驻留
    int maxBlocksPerSM;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &maxBlocksPerSM, kernel, blocksize, 0);

    int maxBlocks = maxBlocksPerSM * prop.multiProcessorCount;
    if (gridsize > maxBlocks) {
        std::cerr << "Grid too large for cooperative launch! "
                     "Max allowed: " << maxBlocks << std::endl;
        return 1;
    }

    // ✅ 参数数组
    void *args[] = { &d_ptr, &N };

    // ✅ 启动 cooperative kernel
    cudaError_t err = cudaLaunchCooperativeKernel(
        (void*)kernel, gridsize, blocksize, args, 0, nullptr);

    if (err != cudaSuccess) {
        std::cerr << "Launch failed: " << cudaGetErrorString(err) << std::endl;
        return 1;
    }

    cudaDeviceSynchronize();

    h_vec = d_vec;
    for (int i = 0; i < h_vec.size(); i++)
        std::cout << h_vec[i] << " ";
    std::cout << std::endl;
}