Cooperative Groups
Cooperative Groups 用 C++ 对象表达线程组协作，适合需要明确线程组粒度的同步、规约、扫描和通用异步搬运场景。本文将 GPU 执行函数称为内核（kernel）。
基本使用
#include <cooperative_groups.h>
namespace cg = cooperative_groups;
__global__ void group_kernel(float* data) {
auto block = cg::this_thread_block();
block.sync();
}
规约、扫描和异步搬运需要额外的公开头文件：
#include <cooperative_groups/reduce.h>
#include <cooperative_groups/scan.h>
#include <cooperative_groups/memcpy_async.h>
使用约束
- 集体操作必须由参与线程组以一致控制流调用；
sync、reduce、scan和memcpy_async不应被不一致分支破坏；- 异步搬运完成前不能读取目标 shared memory；
- grid/cluster 级能力需要额外确认启动方式和设备能力，不作为普通 kernel 的默认路径。
通用异步搬运
对于通用的 global-to-shared 搬运，可以使用 cg::memcpy_async 和对应的线程组等待。它不需要 TME descriptor。需要 tensor layout、tile descriptor 或 TME barrier 流程时，转到 TME 与异步数据搬运。
选择建议
| 场景 | 入口 |
|---|---|
| block/tile 协作同步 | thread_block 或 thread_block_tile |
| group 规约/扫描 | reduce.h 或 scan.h |
| 通用异步搬运 | memcpy_async.h 和 group wait |
| Tensor tile 搬运 | TME 页面 |
