asc-devkit GroupBarrier 样例详解：基于 Ascend C 的多 AIV 分组屏障同步实现与实战

人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本篇文章围绕 asc-devkit 仓库中的 GroupBarrier 样例 展开深入讲解如何基于 Ascend C 的GroupBarrier同步原语实现两组存在依赖关系的 AIVAI Vector Core之间的分组同步Arrive 组 AIV 完成计算后Wait 组 AIV 才能基于其结果继续执行。读完本文你将掌握GroupBarrier的接口用法、底层 GM workspace 布局与原子计数原理、完整可编译运行的样例代码以及 NPU 实机与仿真两种模式的编译运行方法。样例概述什么是 GroupBarrier在多核多 AIV并行计算场景中经常出现一组核的计算结果被另一组核依赖的流水式依赖关系。GroupBarrier正是为此设计的分组屏障同步原语Arrive 组先执行计算的 AIV 集合。完成自己的任务后调用Arrive通知其他核Wait 组依赖 Arrive 组结果的 AIV 集合。通过调用Wait阻塞等待直到 Arrive 组全部完成。本样例启用 8 个 AIV 核其中 2 个 AIV 作为 Arrive 组通过原子累加把指定数值写入 Global MemoryGM后调用Arrive其余 6 个 AIV 作为 Wait 组先调用Wait等待 Arrive 组完成写入再从 GM 读出该数值并借助printf打印验证。注意本样例仅适用于基于 TPipe 和 TQue 的编程模式即 SIMD向量编程范式不适用于 SIMT 等纯标量编程模式。在同步控制 API 样例集中GroupBarrier 样例 与 SyncAll、CrossCoreSetFlag/CrossCoreWaitFlag 同属于核间同步类别区别在于 GroupBarrier 面向的是两组核之间存在单向数据依赖的场景粒度是分组而非全局。支持的产品与 CANN 版本产品CANN 软件版本Ascend 950PR950DT系列产品 CANN 9.1.0Atlas A3系列产品 CANN 9.0.0Atlas A2系列产品 CANN 9.0.0从源码实现看GroupBarrier的接口头文件按架构做了条件编译include/basic_api/core_mng/roc/kernel_operator_group_barrier_intf.h 中__NPU_ARCH__ 2201时引入 dav_c220 的实现、__NPU_ARCH__ 3510时引入 dav_3510 的实现分别对应 Atlas A2/A3 系列与 Ascend 950PR950DT 系列与样例声明支持的硬件范围一致。样例规格与目录结构输入输出规格类型nameshapedata typeformat样例输入barGm[3072]uint8_tND样例输出z[8]int32_tND核函数名kernel_group_barrier其中barGm同时承担两个角色既是 Arrive 组向 GM 写入数值的目标缓冲区也是GroupBarrier使用的组同步消息 workspace下文原理部分会详述其内部布局。目录结构examples/01_simd_cpp_api/03_basic_api/05_sync_control/group_barrier/ ├── CMakeLists.txt // 编译工程文件 ├── group_barrier.asc // Ascend C 样例实现 调用样例 ├── README.md // 样例说明文档 └── README_en.md // 英文版样例说明实现原理GM workspace 与原子计数接口原型GroupBarrier的公开接口定义在 include/basic_api/core_mng/roc/kernel_operator_group_barrier_intf.htemplate PipeMode pipeMode class GroupBarrier { public: __aicore__ inline GroupBarrier(GM_ADDR groupWorkspace, uint32_t arriveSizeIn, uint32_t waitSizeIn); __aicore__ inline void Arrive(uint32_t arriveIndex); // stuck in while loop until all aiv has arrived, then update wait counter __aicore__ inline void Wait(uint32_t waitIndex); __aicore__ inline uint64_t GetWorkspaceLen(); };模板参数PipeMode发送组同步消息时使用的执行单元仅支持PipeMode::MTE3_MODE构造函数中有ASCENDC_DEBUG_ASSERT断言校验groupWorkspace该 GroupBarrier 在 GM 上的起始地址用于存储组同步消息必须 512B 对齐且需要用户自行管理地址对齐与清零arriveSizeIn/waitSizeInArrive 组与 Wait 组的 AIV 个数均需大于 0 且不超过GetBlockNum()即当前核函数实际启动的核数。Workspace 内存布局关键常量定义在 impl/basic_api/dav_c220/core_mng/roc/kernel_operator_cube_group_info.hconstexpr uint32_t BARRIER_SIZE 64; // 1 barrier for 1 64B due to dcci constexpr uint32_t BARRIER_MAX_AIV 50; // at most 50 aiv constexpr uint16_t CACHE_LINE_LEN 512; // cacheline length is 512B构造时见 impl/basic_api/dav_c220/core_mng/roc/kernel_operator_group_barrier_impl.h 与 impl/basic_api/dav_3510/core_mng/roc/kernel_operator_group_barrier_impl.hbarrierInfoArrive groupWorkspace前 64B存放 arrive 计数barrierInfoWait groupWorkspace BARRIER_SIZE紧接着的 64B存放 wait 计数每个 Wait 组 / Arrive 组的核使用CACHE_LINE_LEN(512B) / sizeof(BarrierInfo)为步长错位访问自己的计数器条目配合dccidata cache clean invalidate指令做缓存一致性维护。GetWorkspaceLen()的返回值即所需 GM 消息空间大小return (arriveSize waitSize) ? arriveSize * CACHE_LINE_LEN : waitSize * CACHE_LINE_LEN;例如 3 个 Arrive、6 个 Wait则GetWorkspaceLen()返回6 * 512 3072字节这也解释了本样例输入barGm为什么是[3072]它正是为 6 个 Wait 组 AIV 预留的 3072B 消息空间。Arrive 与 Wait 的底层行为Arrive通过SetAtomicAddImplint32_t()开启原子累加将 UB 中预设为 1 的计数块以DataCopyUB2GMImpl写入 GM 中的 arrive 计数器每核一个 32B block、GM 侧按 512B 间隔排布随后关闭原子模式SetAtomicNoneImpl。如果counter 1多轮使用场景会先检查上一轮所有 Wait 是否都已结束确保轮次间不串扰Wait进入自旋等待循环dcci失效缓存后反复读取barrierInfoArrive-head直到其值 waitCounter * arriveSize即 Arrive 组本轮全部到达随后同样用原子加更新 wait 计数器。因此从实现层面看GroupBarrier 本质上是GM 上的原子计数器 dcci 缓存一致性指令 while 自旋的组合而非硬件专用屏障指令这也决定了它支持在循环中反复使用但循环次数受多核通信效率限制API 文档约束为不超过 1,048,575 次。样例实现代码解析完整实现位于 group_barrier.asc核心逻辑在KernelGroupBarrier::Process()中constexpr int32_t ARRIVE_NUM 2; // arriver num constexpr int32_t WAIT_NUM 6; // wait num constexpr int32_t SIZE_NUM 1; // tensor size constexpr int32_t TEMP ARRIVE_NUM * WAIT_NUM; // initial number __aicore__ inline void Process() { // apply GroupBarrier AscendC::GroupBarrierAscendC::PipeMode::MTE3_MODE barA(barWorkspace, ARRIVE_NUM, WAIT_NUM); auto id AscendC::GetBlockIdx(); if (id 0) { AscendC::InitOutputint32_t(dstGlobal, totalLength, 0); // init output zero } AscendC::SyncAll(); if (id 0 id ARRIVE_NUM) // arrive realize { AscendC::LocalTensorint32_t srcLocal inQueue.AllocTensorint32_t(); AscendC::DisableDmaAtomic(); srcLocal.SetValue(0, TEMP); // set initial number AscendC::SetAtomicAddint32_t(); // add initial number AscendC::DataCopy(dstGlobal, srcLocal, totalLength); AscendC::DisableDmaAtomic(); inQueue.FreeTensor(srcLocal); barA.Arrive(id); // arrive end } else if (id ARRIVE_NUM id ARRIVE_NUM WAIT_NUM) // wait realize { barA.Wait(id - ARRIVE_NUM); // wait end int32_t val dstGlobal.GetValue(0); AscendC::printf([Block (%lu/%d)]: OUTPUT %d\n, id - ARRIVE_NUM, WAIT_NUM, val); } }关键步骤拆解创建 GroupBarrier以传入的barWorkspace即barGm为消息空间声明 Arrive 组 2 核、Wait 组 6 核写初值Block 0 先把输出dstGlobal初始化为 0随后SyncAll()保证所有核看到已初始化的状态Arrive 组Block 0、1在inQueue中分配 LocalTensor写入初值TEMP ARRIVE_NUM * WAIT_NUM 12开启原子累加后通过DataCopy写入dstGlobal。由于两个 Arrive 核都对同一地址执行原子加最终 GM 中的值为12 12 24写完后调用barA.Arrive(id)通知 Wait 组Wait 组Block 2~7调用barA.Wait(id - ARRIVE_NUM)阻塞直到 Arrive 组两个核都完成写入随后dstGlobal.GetValue(0)读到 24用printf打印。值得注意的是样例通过AscendC::SetAtomicAddint32_t()与DisableDmaAtomic()配合把Arrive 组多个核向同一 GM 地址做原子累加这一过程完整展示了出来——这正是TEMP取ARRIVE_NUM * WAIT_NUM的原因使最终结果 24 能够同时验证原子加生效与Wait 同步正确两层语义。核函数入口与主机侧调用使用 ACL 运行时 APIaclrtMalloc/aclrtMemcpy/aclrtLaunchKernel风格extern C __global__ __vector__ void kernel_group_barrier(__gm__ uint8_t* barGm, __gm__ uint8_t* dstGm) { AscendC::TPipe pipe; KernelGroupBarrier op; op.Init(barGm, dstGm, pipe); op.Process(); }主机侧以kernel_group_barriernumBlocks, 0, stream(xDevice, zDevice)形式启动 8 个 Block其中numBlocks 8与 Arrive(2) Wait(6) 的划分严格对应。编译运行1. 配置环境变量请根据当前环境上 CANN 开发套件包的安装方式配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。2. 编译与执行在样例目录下执行mkdir -p build cd build; # 创建并进入build目录 cmake .. -DCMAKE_ASC_ARCHITECTURESdav-2201;make -j; # 编译工程 ./demo # 执行编译生成的可执行程序执行样例使用 NPU 仿真模式时添加-DCMAKE_ASC_RUN_MODEsim参数即可cmake .. -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201;make -j; # NPU仿真模式注意切换编译模式前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新 cmake。编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、sim运行模式NPU 运行、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2系列产品和 Atlas A3系列产品dav-3510 对应 Ascend 950PR950DT系列产品编译工程文件 CMakeLists.txt 通过find_package(ASC REQUIRED)引入 Ascend C 编译工具链以project(kernel_samples LANGUAGES ASC CXX)声明 ASC/C 双语言并将--npu-arch${CMAKE_ASC_ARCHITECTURES}作为 ASC 源文件的编译选项通过add_executable(demo group_barrier.asc)生成名为demo的可执行程序。执行结果执行成功时输出如下6 个 Wait 组 AIV 均读到由 Arrive 组原子累加写入的数值 24[Block (0/6)]: OUTPUT 24 [Block (1/6)]: OUTPUT 24 [Block (2/6)]: OUTPUT 24 [Block (3/6)]: OUTPUT 24 [Block (4/6)]: OUTPUT 24 [Block (5/6)]: OUTPUT 24使用约束与注意事项结合 API 文档docs/zh/api/SIMD-API/basic_api/cube_group_mgmt_ISASI/GroupBarrier与源码实现使用GroupBarrier时有以下约束仅支持PipeMode::MTE3_MODE构造函数的ASCENDC_DEBUG_ASSERT会在其他模式下直接报错workspace 需用户自管groupWorkspace必须 512B 对齐、且由用户负责清零所需大小可通过GetWorkspaceLen()获取UB 空间占用使用该接口时Unified BufferUB空间末尾的 1600B 会被占用对应源码中UB_START_ADDR TOTAL_UB_SIZE - ONE_BLK_SIZE * BARRIER_MAX_AIV预留的屏障计数区不可与REGIST_MATMUL_OBJ混用两者都涉及 AIC 与 AIV 的核间通信管理同时使用可能导致通信消息错误等异常循环使用上限Arrive/Wait支持在循环中使用但受多核间通信效率限制循环最大次数不超过 1,048,575 次序号范围arriveIndex取值范围为[0, arriveSize - 1]waitIndex取值范围为[0, waitSize - 1]AIV 数量上限从常量BARRIER_MAX_AIV 50可见单个 GroupBarrier 支持的 AIV 规模存在上限最多 50 个 AIV设计分组时应预留余量。源码佐证接口、实现与测试闭环若要深入验证上述行为可以在仓库中找到完整的三层证据链接口层include/basic_api/core_mng/roc/kernel_operator_group_barrier_intf.h 声明了GroupBarrier模板类的全部公开方法实现层按架构分目录实现如 dav_c220 实现、dav_3510 实现、dav_9201 实现其中可见原子加、dcci、自旋等待等关键细节配套常量定义在 kernel_operator_cube_group_info.h测试层test_cube_group_barrier.cpp 中的GroupBarrierSingleAivArriveCompletes用例验证了单核 Arrive 场景下GetWorkspaceLen()返回CACHE_LINE_LEN(512)并覆盖了 Cube 组消息收发、消息状态流转等相邻能力可作为功能正确性的参照。此外GroupBarrier的官方 API 使用说明创建 → Arrive → Wait 的完整调用范式参见 GroupBarrier使用说明构造函数、Arrive、Wait、GetWorkspaceLen 的逐参数说明分别见 构造函数、Arrive、Wait、GetWorkspaceLen。小结GroupBarrier 样例是理解 Ascend C 多核分组同步的极佳起点它用最简的写数-等待-打印逻辑完整呈现了 Arrive 组与 Wait 组之间的依赖同步过程同时覆盖了原子累加写 GM、workspace 布局、编译选项与仿真运行等实战要素。在实际算子开发中凡是前一组 AIV 的计算结果被后一组 AIV 消费的流水场景例如分阶段流水、生产者-消费者形态的多核协同计算都可以参照本样例将GroupBarrierMTE3_MODE嵌入基于 TPipe/TQue 的 SIMD 编程模型实现正确且高效的核间同步。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐asc-devkit PhiloxRandom 随机数生成样例详解基于 Ascend C 高阶 API 的端到端实现asc devkit PhiloxRandom 随机数生成样例详解基于 Ascend C 高阶 API 的端到端实现 导读 本文以 asc devkit 开源人工智能深度学习算子库CANNAscendAscend C 算子样例解析基于 Interleave/DeInterleave 实现元素交织与解交织asc-devkitAscend C 算子样例解析基于 Interleave/DeInterleave 实现元素交织与解交织asc devkit 本篇文章围绕 asc dev人工智能深度学习算子库CANNAscendasc-devkit RegBase 向量计算实战基于 SIMD C API 的 add 与 GELU 样例详解asc devkit RegBase 向量计算实战基于 SIMD C API 的 add 与 GELU 样例详解 本文基于 asc devkitCANN人工智能深度学习算子库CANNAscend上一篇从零复刻 Mango手把手教你用 Kotlin MVP 写一个 Dribbble 客户端下一篇Qwen3-VL-30B-A3B-Thinking-FP840亿参数开启多模态AI普惠时代创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考