新增 SIMT 协作组 grid_group-asc-devkit-AtomGit

已合并

新增 SIMT 协作组 grid_group #4795

louyuxuan创建于 20 天前从louyuxuan/asc-devkit： master 合入到 cann/asc-devkit： master

新增 SIMT 协作组 grid_group #4795

已合并

louyuxuan创建于 20 天前

从louyuxuan/asc-devkit master合入到cann/asc-devkit master

louyuxuan成员

20 天前

描述

Ascend C SIMT 协作组支持 grid_group，提供跨block 的同步能力，具体支持API：

this_grid()
bool is_valid() const
void sync() const
static unsigned long long thread_rank()
static unsigned long long block_rank()
static unsigned long long num_threads()
static unsigned long long num_blocks()
static dim3 dim_blocks()
static dim3 block_index()
static unsigned long long size()
static dim3 group_dim()

关联的Issue

#1391

测试

SIMT 二级冒烟验证
编写算子测试 grid_group 所有 api

文档更新

docs/zh/api/README.md
docs/zh/api/SIMT-API/SIMD_SIMT_hybrid_programming_intro/api_list.md
docs/zh/api/SIMT-API/SIMT_programming_intro/api_list.md
docs/zh/api/SIMT-API/cooperative_groups/cooperative_groups.md
docs/zh/api/SIMT-API/cooperative_groups/cooperative_groups_intro.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/block_index.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/block_rank.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/dim_blocks.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/grid_group.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/grid_group_constructor.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/grid_group_intro.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/group_dim.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/is_valid.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/num_blocks.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/num_threads.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/size.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/sync.md
docs/zh/api/SIMT-API/cooperative_groups/grid_group/thread_rank.md
docs/zh/api/SIMT-API/cooperative_groups/sync.md

类型标签

🐛 fix: Bug 修复
✨ feat: 新功能
⚡ perf: 性能优化
♻️ refactor: 代码重构
🧪 test: 新增或修改测试
📝 docs: 文档更新
🔧 ci: CI/CD 配置修改
↩️ revert: 回退
🧹 chore: 其他，请具体描述

Pull Request已成功合入, 合并人@CANN-robot

（感谢 louyuxuan 的贡献）

L louyuxuan成员

20 天前创建了 pull request，commit aca52eef

L louyuxuan成员

20 天前关联了issue：[Requirement|需求建议]: Ascend C SIMT 协作组支持 grid_group

atomgit-bot

20 天前评论：

变更摘要

本 PR 新增 SIMT 协作组 grid_group（关联 Issue #1391），支持对整张 grid 上的所有线程块进行跨块同步。核心是在 cooperative_groups.h 中声明 grid_group 类并新增 group_type::grid_group_type 枚举值，在 cooperative_groups_impl.h 中实现其全部接口：grid_group::sync() 基于参数基址前偏移 32 字节的全局同步计数器（高 32 位记录栅栏阶段号、低 32 位统计到达的 block 数，配合 asc_atomic_add 与自旋等待实现可重复的跨 block 屏障），并提供 num_blocks()、num_threads()、size()、block_rank()、thread_rank()、dim_blocks()、group_dim()、block_index()、is_valid() 等查询接口以及入口函数 this_grid()；同时让 thread_group 基类的 size()、thread_rank()、sync() 多态分发支持 grid_group_type。测试侧新增覆盖 grid_group 各接口及同步计数器行为（单 block 不触碰计数器、block 0/非 0 释放、重复屏障推进阶段号）的用例，并在编译桩中补充 __cce_simt_get_para_base_imp() 的实现以支撑测试。

主要改动

新增 grid_group 类与枚举值: 在 cooperative_groups.h 中声明 grid_group（继承 thread_group，声明 this_grid() 友元），并在 group_type 中新增 grid_group_type；cooperative_groups_impl.h 中新增 this_grid() 入口及各成员函数实现，thread_group 基类的 size()、thread_rank()、sync() 增加 grid_group_type 分派分支。
实现跨 block 同步屏障 grid_group::sync(): 通过 details::get_grid_sync_counter() 获取参数基址前 4×8 字节处的 64 位同步计数器，利用低 32 位到达计数与高 32 位阶段号（grid_sync_counter_phase_shift 等常量）配合 asc_atomic_add 原子操作和阶段自旋等待实现可重复的 grid 级 sync()，并在同步前后以 asc_threadfence()/asc_syncthreads() 保证内存可见性。
新增 grid 维度与秩查询接口: 实现 grid_group::num_blocks()（由 gridDim 计算）、num_threads()、size()、block_rank()、thread_rank()（按 block 序展开的全局线程秩）、dim_blocks()/group_dim()/block_index() 等静态接口，并新增 is_valid()。
新增测试用例覆盖 grid_group: 在 test_operator_simt_c_cooperative_groups.cpp 中新增 this_grid() 类型/有效性、基类分发、3D 维度与秩计算，以及单 block 不写计数器、block 0/非 0 释放计数器、重复同步推进阶段号等同步行为测试。
补充测试桩支持: 在 simt_compiler_stub.h/cpp 中新增 __cce_simt_get_para_base_imp() 声明与实现（返回内部静态数组偏移地址），为测试中的同步计数器访问提供支撑。