企业官网建设流程全解析

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend

导读

本篇文章围绕 asc-devkit 仓库中的 GroupBarrier 样例 展开，深入讲解如何基于 Ascend C 的GroupBarrier同步原语，实现两组存在依赖关系的 AIV（AI 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 Memory（GM）后调用Arrive；其余 6 个 AIV 作为 Wait 组，先调用Wait等待 Arrive 组完成写入，再从 GM 读出该数值并借助printf打印验证。

注意：本样例仅适用于基于 TPipe 和 TQue 的编程模式，即 SIMD（向量）编程范式，不适用于 SIMT 等纯标量编程模式。

在同步控制 API 样例集中，GroupBarrier 样例 与 SyncAll、CrossCoreSetFlag/CrossCoreWaitFlag 同属于核间同步类别，区别在于 GroupBarrier 面向的是"两组核之间存在单向数据依赖"的场景，粒度是分组而非全局。

支持的产品与 CANN 版本

从源码实现看，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 950PR&950DT 系列，与样例声明支持的硬件范围一致。

样例规格与目录结构

输入输出规格

其中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.h：

template <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/waitSizeIn：Arrive 组与 Wait 组的 AIV 个数，均需大于 0 且不超过GetBlockNum()（即当前核函数实际启动的核数）。

Workspace 内存布局

关键常量定义在 impl/basic_api/dav_c220/core_mng/roc/kernel_operator_cube_group_info.h：

constexpr 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.h）：

• barrierInfoArrive = groupWorkspace（前 64B，存放 arrive 计数）；
• barrierInfoWait = groupWorkspace + BARRIER_SIZE（紧接着的 64B，存放 wait 计数）；
• 每个 Wait 组 / Arrive 组的核使用CACHE_LINE_LEN(512B) / sizeof(BarrierInfo)为步长错位访问自己的计数器条目，配合dcci（data 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：通过SetAtomicAddImpl<int32_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::GroupBarrier<AscendC::PipeMode::MTE3_MODE> barA(barWorkspace, ARRIVE_NUM, WAIT_NUM); auto id = AscendC::GetBlockIdx(); if (id == 0) { AscendC::InitOutput<int32_t>(dstGlobal, totalLength, 0); // init output zero } AscendC::SyncAll(); if (id >= 0 && id < ARRIVE_NUM) // arrive realize { AscendC::LocalTensor<int32_t> srcLocal = inQueue.AllocTensor<int32_t>(); AscendC::DisableDmaAtomic(); srcLocal.SetValue(0, TEMP); // set initial number AscendC::SetAtomicAdd<int32_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); } }

关键步骤拆解：

1. 创建 GroupBarrier：以传入的barWorkspace（即barGm）为消息空间，声明 Arrive 组 2 核、Wait 组 6 核；
2. 写初值：Block 0 先把输出dstGlobal初始化为 0，随后SyncAll()保证所有核看到已初始化的状态；
3. Arrive 组（Block 0、1）：在inQueue中分配 LocalTensor，写入初值TEMP = ARRIVE_NUM * WAIT_NUM = 12，开启原子累加后通过DataCopy写入dstGlobal。由于两个 Arrive 核都对同一地址执行原子加，最终 GM 中的值为12 + 12 = 24；写完后调用barA.Arrive(id)通知 Wait 组；
4. Wait 组（Block 2~7）：调用barA.Wait(id - ARRIVE_NUM)阻塞，直到 Arrive 组两个核都完成写入；随后dstGlobal.GetValue(0)读到 24，用printf打印。

值得注意的是，样例通过AscendC::SetAtomicAdd<int32_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_barrier<<<numBlocks, 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_ARCHITECTURES=dav-2201;make -j; # 编译工程 ./demo # 执行编译生成的可执行程序，执行样例

使用 NPU 仿真模式时，添加-DCMAKE_ASC_RUN_MODE=sim参数即可：

cmake .. -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201;make -j; # NPU仿真模式

注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下执行rm CMakeCache.txt后重新 cmake。

编译选项说明

编译工程文件 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 Buffer（UB）空间末尾的 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），设计分组时应预留余量。

源码佐证：接口、实现与测试闭环

若要深入验证上述行为，可以在仓库中找到完整的三层证据链：

1. 接口层：include/basic_api/core_mng/roc/kernel_operator_group_barrier_intf.h 声明了GroupBarrier模板类的全部公开方法；
2. 实现层：按架构分目录实现，如 dav_c220 实现、dav_3510 实现、dav_9201 实现，其中可见原子加、dcci、自旋等待等关键细节；配套常量定义在 kernel_operator_cube_group_info.h；
3. 测试层：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 消费"的流水场景（例如分阶段流水、生产者-消费者形态的多核协同计算），都可以参照本样例将GroupBarrier<MTE3_MODE>嵌入基于 TPipe/TQue 的 SIMD 编程模型，实现正确且高效的核间同步。

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend