cooperative_groups网格级同步的本质是启动必须全部参与且原子性完成的跨SM kernel，并用cg::grid_group::synchronize()实现全局屏障；需满足计算能力≥6.0、TCC模式（Tesla卡）、空流三条件，且仅在该kernel内生效。
什么是 cooperative_groups 网格级同步的本质
它不是“让所有 GPU 线程等一个栅栏”，而是允许你启动一个跨多个 SM 的、**必须全部参与且原子性完成**的 kernel，并在其中用 cg::grid_group 调用 synchronize()。前提是：kernel 必须用 cudaLaunchCooperativeKernel 启动，且设备支持（cudaDeviceGetAttribute(&attr, cudaDevAttrCooperativeLaunch, dev) 返回非零），常见于 Tesla A100 / V100 / H100，消费卡（如 RTX 4090）默认禁用或不支持。
启动 cooperative kernel 的三个硬性条件
缺一不可，否则 cudaLaunchCooperativeKernel 直接返回 cudaErrorNotSupported 或 cudaErrorInvalidValue：
- GPU 计算能力 ≥ 6.0（Pascal），但实际需看驱动与模式：Tesla 卡需在 TCC 模式下运行；Windows WDDM 下一律不支持
- Host 端调用前必须设置流为 NULL：
cudaStream_t stream = 0，不能传自定义流 - Kernel 中若使用
cg::grid_group，必须在 kernel 入口第一行声明并获取，例如：__global__ void my_kernel() { namespace cg = cooperative_groups; cg::grid_group grid = cg::this_grid(); // 后续才能调用 grid.synchronize() }
cg::grid_group::synchronize() 的行为边界
它只对当前 cooperative kernel 内所有线程生效，不阻塞 host，也不影响其他 kernel。但它会强制等待：所有 SM 上该 kernel 的所有 block 都执行到该语句，且所有 warp 都到达后才继续。注意以下陷阱：
- 不能在
if (threadIdx.x == 0)这类 divergent 分支里调用——会导致部分 warp 永远不执行，kernel hang 住 - 不能和
__syncthreads()混用在同个 block 内做嵌套同步——无定义行为，常见死锁 - 每个 block 内只能调用一次
grid.synchronize()（多次调用不报错但无意义，且可能触发驱动异常） - 同步开销显著：A100 上单次耗时约 5–10 μs，比
__syncthreads()高两个数量级
典型协作模式：全局归约 + 分布式 barrier
适合需要所有 block 协同完成阶段性计算的场景，比如多 block 共同构建哈希表、分阶段排序、或分布式 SGD 的梯度聚合。下面是最小可运行结构：
立即学习“C++免费学习笔记（深入）”；
__global__ void global_reduce_kernel(float* data, int n) {
namespace cg = cooperative_groups;
cg::grid_group grid = cg::this_grid();
extern __shared__ float sdata[];
int tid = threadIdx.x;
int bid = blockIdx.x;
int offset = bid * blockDim.x;
// Step 1: block 内规约到 shared memory
sdata[tid] = (offset + tid < n) ? data[offset + tid] : 0.f;
__syncthreads();
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) sdata[tid] += sdata[tid + s];
__syncthreads();
}
// Step 2: 只有 block 0 把结果写入 global memory（或其他协调逻辑）
if (bid == 0 && tid == 0) {
data[0] = sdata[0];
}
// Step 3: 所有 block 等待 block 0 完成写入，再进入下一阶段
grid.synchronize();
// Step 4: 此时可安全读取 data[0] 做后续广播/校正
if (bid == 1 && tid == 0) {
float global_sum = data[0];
// do something with global_sum...
}}
真正难的从来不是写这十几行代码，而是确认你的运行环境是否满足 cooperative launch 的全部约束——尤其是驱动模式、设备属性、流参数这三个点，漏查一个，kernel 就静默失败或直接 crash。
