人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读在昇腾 AI 处理器的 MIX 混合编程场景中AICAI Cube 核负责 Matmul 计算、AIVAI Vector 核负责向量运算二者通过核间消息协同。本文以 CANN asc-devkit 开源仓库为依托系统讲解如何通过Iteratefalse/IterateAllfalse异步接口减少 AIC 与 AIV 之间的消息交互、消除同步依赖并给出可复制的正反例代码、Workspace 配置要点以及仓库内的完整可运行样例帮助读者在算子开发中直接落地这一核间通信优化手段。一、背景MIX 场景下的核间同步开销在 MIX 场景AIC 与 AIV 混合编程中AIV 核需要调用 Matmul 的Iterate或IterateAll接口通过发送消息到 AIC 核来启动 Matmul 计算。每一次 Matmul 计算片C 分片的产出与取回都伴随着 Cube 核与 Vector 核之间的消息传递与同步等待。该问题在 asc-devkit 文档体系中属于高优先级优化项专门针对流水线调度Pipeline Scheduling场景。其核心矛盾在于同步方式Iteratetrue默认每次调用Iterate都会触发一次消息发送AIV 需要等待 AIC 计算完成后才能继续异步方式Iteratefalse仅第一次调用需要发送消息启动计算后续迭代不再发送消息AIV 无需逐次等待从而减少 Cube 与 Vector 核间交互显著降低核间通信开销。因此MIX 场景官方推荐使用Iteratefalse或IterateAllfalse异步接口前提是使用异步接口时必须设置 Workspace。下图分别展示了同步与异步两种方式下 AIV 向 AIC 发送消息启动 Matmul 计算的差异同步方式消息发送示意图异步方式消息发送示意图二、接口本质sync模板参数决定消息发送行为Iterate与IterateAll都是带模板参数sync的接口默认值为true同步。相关声明位于仓库公开接口文件 matmul.h 与 matmul.h// 计算一个 baseM * baseN 大小的 C 矩阵分片 // tparam sync: 设置为同步或异步模式 // param enPartialSum: 是否将 Iterate 的结果累加到 CO1 数据中 template bool sync true __aicore__ inline bool Iterate(bool enPartialSum false); // 计算 singleCoreM * singleCoreN 大小的完整 C 矩阵 // param waitIterateAll: 异步模式下是否通过 WaitIterateAll 等待 IterateAll 完成 template bool sync true __aicore__ inline void IterateAll( const GlobalTensorDstT gm, uint8_t enAtomic 0, bool enSequentialWrite false, bool waitIterateAll false, bool fakeMsg false);Iterate以baseM × baseN为基本分片粒度逐片计算配合GetTensorC逐片取回结果适合 AIV 侧边算边处理的流水线场景IterateAll一次性计算singleCoreM × singleCoreN的完整 C 矩阵适合整块输出的场景。两者语义一致sync true为同步模式sync false为异步模式。GetTensorC同样支持template bool sync true异步取数见 matmul.h异步场景下通常与Iteratefalse配套使用。2.1 客户端实现异步模式仅发送一次消息在 KFCKernel Function Call客户端实现 matmul_client.h 中Iterate的行为清晰地体现了仅第一次发送消息这一机制template bool sync true __aicore__ inline bool Iterate(bool enPartialSum false) { TRACE_START(TraceId::KFC_CLIENT_POST_MSG); if (unlikely(kfcMsg_.body.isFirstIter)) { cntIter_ 0; cOffset_ 0; curProcess 0; } else { if (cntIter_ mnIter_) { TRACE_STOP(TraceId::KFC_CLIENT_POST_MSG); return false; // 迭代结束 } if constexpr (!sync) { TRACE_STOP(TraceId::KFC_CLIENT_POST_MSG); return true; // 异步模式后续迭代直接返回不发送消息 } } if constexpr (!sync) { // 异步模式约束 ASSERT(cacheWorkspaceAddr ! 0); // 必须已通过 SetWorkspace 配置缓存地址 ASSERT(PhyPosIsUB(C_TYPE::pos)); // 异步模式仅支持输出到 UB } isSyncGetC sync; kfcMsg_.body.sync sync; kfcMsg_.body.cAddr reinterpret_castuint64_t(cacheWorkspaceAddr); ... PostMessageKFC_Enum::MMFUN_ITERATE, false(); // 同步模式每次迭代都会发送消息 SyncCubeWithVecA_TYPE::ibShare, B_TYPE::ibShare(); ... return true; }从源码可以清楚看到异步模式sync false除第一次isFirstIter外后续每次Iterate直接return true不会执行PostMessage即不再向 AIC 发送消息同步模式sync true每次迭代都会执行PostMessageKFC_Enum::MMFUN_ITERATE发送消息并执行SyncCubeWithVec同步这正是每次调用都会触发一次消息发送的代码级根因。2.2 服务器端实现异步模式下 C 分片直出在 KFC 服务器端 matmul_server_impl_3510.h 的Iterate处理逻辑中异步 Workspace 场景会走直出快速路径// 异步且配置了 workspace 时C 分片直接搬出无需逐次切换消息队列 if (unlikely(isIterate body.cAddr ! 0 hasSetWorkspace sync 0)) { mul.GetTensorC(cGlobal[offset], (uint8_t)(enAtomic), enSequentialWrite); offset offsetSize; CrossCoreSetFlagINTRA_MODE, PIPE_FIX( GetIntraFlagId(instID, static_castuint8_t(VEC_WAIT_INTRA_Enum::WAIT_FIXP), mul.GetSubBlockIdx())); continue; }同时GetOffsetSize 中专门注释说明For asynchronous Iterate, the offset must be used for address calculation and the size is baseM x baseN即异步 Iterate 需要按baseM × baseN计算偏移地址C 分片依次落入 Workspace 中不同的偏移位置。服务器端通过hasSetWorkspace即是否调用过SetWorkspace与sync字段共同判定是否走异步直出路径这进一步印证了使用异步接口必须设置 Workspace的硬性要求。三、反例与正例同步 vs 异步代码对比3.1 反例MIX 场景使用 Iterate 的同步方式以下代码在 MIX 场景中使用Iterate()即Iteratetrue同步方式每次循环迭代都会触发一次 AIV→AIC 的消息发送与核间同步TQueBindTPosition::CO2, TPosition::VECIN qVecIn; TQueBindTPosition::VECIN, TPosition::VECOUT qVecOut; mm.SetTensorA(gmA); mm.SetTensorB(gmB); int16_t scalar 2; while(mm.template Iterate()){ auto cInUB qVecIn.AllocTensorfloat(); mm.GetTensorC(cInUB); qVecIn.EnQue(cInUB); cInUB qVecIn.DeQuefloat(); auto cOutUB qVecOut.AllocTensorfloat(); Muls(cOutUB, cInUB, scalar, baseM*baseN); qVecIn.FreeTensor(cInUB); ... }3.2 正例MIX 场景使用 Iterate 的异步方式正例与反例相比仅有两处关键差异① 循环条件改为Iteratefalse② 循环前调用SetWorkspace配置临时空间TQueBindTPosition::CO2, TPosition::VECIN qVecIn; TQueBindTPosition::VECIN, TPosition::VECOUT qVecOut; mm.SetTensorA(gmA); mm.SetTensorB(gmB); mm.SetWorkspace(workspace, size);//其中workspace为临时空间的物理地址size为singleCoreM*singleCoreN大小的矩阵C占用的内存大小singleCoreM*singleCoreN*sizeof(float) int16_t scalar 2; while(mm.template Iteratefalse()){ auto cInUB qVecIn.AllocTensorfloat(); mm.GetTensorC(cInUB); qVecIn.EnQue(cInUB); cInUB qVecIn.DeQuefloat(); auto cOutUB qVecOut.AllocTensorfloat(); Muls(cOutUB, cInUB, scalar, baseM*baseN); qVecIn.FreeTensor(cInUB); ... }两段代码的对比要点总结如下对比维度同步方式Iteratetrue异步方式Iteratefalse消息发送次数每次迭代均发送一次仅第一次发送后续迭代不再发送核间同步每次迭代等待 AIC 计算完成无需逐次等待并行度更高是否必须设置 Workspace否必须SetWorkspace适用输出位置不限输出到 UBC 矩阵在 UB 上核间通信开销高低注意原文档明确指出使用异步接口时必须设置 Workspace否则会触发ASSERT(cacheWorkspaceAddr ! 0)断言失败见 matmul_client.h。四、Workspace 配置详解4.1SetWorkspace的两种调用形式SetWorkspace在公开接口 matmul.h 中提供两种重载// 形式一GM 裸指针 内存大小size 的单位为字节 // param addr: workspace on GM provided by user, GM address type // param size: 内存大小 template class T __aicore__ inline void SetWorkspace(__gm__ const T* addr, int size); // 形式二GlobalTensor推荐使用 // param addr: workspace on GM provided by user, GlobalTensor type template class T __aicore__ inline void SetWorkspace(GlobalTensorT addr);以原文档正例中的SetWorkspace(workspace, size)为例workspace临时空间的物理地址GM 地址sizesingleCoreM * singleCoreN * sizeof(float)即单核 Matmul 计算所需的整块 C 矩阵占用的内存大小字节数。在客户端实现 matmul_client.h 中GlobalTensor形式的SetWorkspace最终转换为裸指针形式并记录缓存地址template class T __aicore__ inline void SetWorkspace(GlobalTensorT addr) { ASSERT(addr.GetSize() 0); SetWorkspace(addr.GetPhyAddr(), addr.GetSize() * sizeof(T)); // size 按字节数传递 } template class T __aicore__ inline void SetWorkspace(__gm__ const T* addr, int size) { ASSERT(addr ! nullptr); cacheWorkspaceAddr reinterpret_castGM_ADDR(const_cast__gm__ T*(addr)); cOffset_ 0; }注意两点size 的单位从GlobalTensor形式的转换实现看addr.GetSize() * sizeof(T)传入的是字节数与原文档singleCoreM*singleCoreN*sizeof(float)的描述一致后续kfcMsg_.body.cAddr即使用该缓存地址作为 C 分片的落盘位置调用时机SetWorkspace必须在Iteratefalse之前调用且每个 Matmul 对象实例只需设置一次cOffset_会随迭代自动累加偏移。4.2 示例工程中的 Workspace 分配方式仓库示例 matmul_async_iterate.asc 展示了完整的 Workspace 获取与设置流程__global__ __aicore__ void matmul_async_iterate_custom( __gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* bias, __gm__ uint8_t* c, __kfc_workspace__ __gm__ uint8_t* workspace, __gm__ uint8_t* tilingGm) { TCubeTiling tiling; CopyTiling(tiling, tilingGm); MatmulKernelhalf, half, float, float matmulKernel; AscendC::TPipe pipe; REGIST_MATMUL_OBJ(pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, tiling); matmulKernel.Init(a, b, bias, c, workspace, tiling); matmulKernel.Process(pipe); }在内核Init中通过AscendC::GetUserWorkspace(workspace)获取用户可用的临时空间并按核偏移切分workspaceGlobal.SetGlobalBuffer( reinterpret_cast__gm__ CType*(AscendC::GetUserWorkspace(workspace)), tiling.M * tiling.N); ... workspaceGlobal workspaceGlobal[AscendC::GetBlockIdx() * tiling.singleCoreM * tiling.singleCoreN];随后在Process中调用matmulObj.SetTensorA(aGlobal, IS_TRANS_A); matmulObj.SetTensorB(bGlobal, IS_TRANS_B); if (IS_BIAS) { matmulObj.SetBias(biasGlobal); } matmulObj.SetWorkspace(workspaceGlobal); matmulObj.template Iteratefalse();即Workspace 由系统通过__kfc_workspace__参数和REGIST_MATMUL_OBJ宏注册每个核分得singleCoreM * singleCoreN大小的独立空间再以GlobalTensor形式传给SetWorkspace。五、IterateAllfalse异步场景与WaitIterateAll与Iterate逐片处理不同IterateAllfalse用于一次性输出整块 C 矩阵。当采用异步模式且输出为连续 GM 时需要配合waitIterateAll true与WaitIterateAll()使用。客户端 matmul_client.h 对WaitIterateAll的约束如下// Only support the mode that the IterateAll is asynchronous and GM output is continuous. // In discontinuous scenarios, the system stops responding. __aicore__ inline void WaitIterateAll() { ASSERT(!isSyncGetC); // Must be asynchronous mode ... }仓库示例 matmul_async_iterate_all.asc 给出了标准用法matmulObj.SetTensorA(aGlobal, IS_TRANS_A); matmulObj.SetTensorB(bGlobal, IS_TRANS_B); if (IS_BIAS) { matmulObj.SetBias(biasGlobal); } matmulObj.template IterateAllfalse(cGlobal, 0, false, true); // 第4个参数 waitIterateAll true matmulObj.WaitIterateAll(); matmulObj.End();要点IterateAllfalse(cGlobal, 0, false, true)的四个参数依次为C 矩阵 GlobalTensor、enAtomic是否使能原子操作0 表示关闭、enSequentialWrite是否顺序写、waitIterateAll是否等待完成异步模式下必须调用WaitIterateAll()等待 AIC 侧全部计算完成再释放/复用后续资源从源码注释看该模式仅支持GM 输出连续的场景若输出地址不连续可能导致系统挂起使用时需确保 C 矩阵在 GM 上连续排布。六、适用前提与注意事项综合原文档说明与仓库源码使用异步接口需要把握以下边界条件必须设置 Workspace异步模式依赖 Workspace 作为 C 分片的缓存中转Iteratefalse路径上有ASSERT(cacheWorkspaceAddr ! 0)硬校验matmul_client.h输出位置约束异步Iterate模式要求 C 矩阵输出到 UBASSERT(PhyPosIsUB(C_TYPE::pos))即C_TYPE的物理位置为 UBIterateAll 异步模式要求 GM 输出连续并配合WaitIterateAll()使用同步/异步的取舍当循环迭代次数较少、或 AIV 侧后续向量计算依赖单次 Matmul 结果强同步时同步模式逻辑更简单当分片迭代次数多、AIV 侧有独立的向量算子流水可并行时异步模式的核间通信收益更明显尾块处理异步模式下服务器端按baseM × baseN计算偏移若存在尾块tail需要像示例中那样通过SetTail(tailM, tailN)处理单核边界形状见 matmul_async_iterate.asc。七、仓库内完整样例与运行方式在 asc-devkit 仓库中与本文主题直接相关的可运行样例位于examples/01_simd_cpp_api/04_advanced_api/00_matmul/下matmul_async_iterateIteratefalseGetTensorCfalse异步直调样例输出到 VECINmatmul_async_iterate_allIterateAllfalseWaitIterateAll()异步整块输出样例matmul_fused融合场景中Iteratefalse异步调用示例batch_matmul_iterate_n_batchBatch Matmul 中SetWorkspace配合异步迭代的示例。以matmul_async_iterate样例为例其 README 给出的编译运行方式如下参考 README.mdsource ${install_path}/cann/set_env.sh # 配置 CANN 环境变量${install_path} 默认 /usr/local/Ascend mkdir -p build cd build; # 创建并进入 build 目录 cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # 编译工程默认 NPU 模式 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行生成的可执行程序 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证结果精度编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2/A3 系列dav-3510 对应 Ascend 950PR/950DT 系列# CPU 调试模式 cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # NPU 仿真模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j;注意切换编译模式前需清理 cmake 缓存如删除 build 目录下的CMakeCache.txt后重新 cmake。样例精度验证通过时终端输出test pass!。八、总结在 MIX 混合编程场景中Matmul 的同步Iterate/IterateAll调用会逐次触发 AIV→AIC 的消息发送与核间同步是算子流水线中的常见性能瓶颈。asc-devkit 提供Iteratefalse与IterateAllfalse异步接口其本质是利用 KFC 消息机制的首次发送、后续直出特性将核间通信次数从每分片一次降为整个计算一次代码改造成本极低仅需将Iterate()改为Iteratefalse或将IterateAll()改为IterateAllfalse并在调用前增加一次SetWorkspace收益明确减少 Cube 与 Vector 核间的消息交互与同步等待降低核间通信开销提升算子整体并行度边界清晰异步模式要求输出到 UB、必须设置 WorkspaceIterateAllfalse要求 GM 输出连续并配合WaitIterateAll()。开发者可直接参照本文正例与仓库内matmul_async_iterate、matmul_async_iterate_all样例完成改造并通过verify_result.py校验精度在保障正确性的前提下获得核间通信优化的收益。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐asc-devkit C_API 异步场景 Add 算子实战基于 asc_copy/asc_add/asc_sync 的矢量算子开发指南asc devkit C_API 异步场景 Add 算子实战基于 asc_copy/asc_add/asc_sync 的矢量算子开发指南 导读 本文以 asc人工智能深度学习算子库CANNAscendasc-devkit Matmul MIX 模式融合编程实践AIC/AIV 分工实现 Matmul LeakyRelu 单核函数融合asc devkit Matmul MIX 模式融合编程实践AIC/AIV 分工实现 Matmul LeakyRelu 单核函数融合 Ascend C 的人工智能深度学习算子库CANNAscendAscend C 核间同步实战基于 IBSet 与 IBWait 的多核数据依赖消除asc-devkit 样例解析Ascend C 核间同步实战基于 IBSet 与 IBWait 的多核数据依赖消除asc devkit 样例解析 本文基于 asc devkit 仓库中人工智能深度学习算子库CANNAscend上一篇TinyUSB USB PD 3.0实现Type-C电力传输协议开发详解下一篇asmttpd 源码架构分析模块化设计的汇编实践创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 
长河君 · 2026/10/17 1:48:17 · 阅读 1.2k asc-devkit 算子优化实战：MIX 场景启用 Iterate/IterateAll 异步接口消除 AIC-AIV 同步依赖
人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读在昇腾 AI 处理器的 MIX 混合编程场景中AICAI Cube 核负责 Matmul 计算、AIVAI Vector 核负责向量运算二者通过核间消息协同。本文以 CANN asc-devkit 开源仓库为依托系统讲解如何通过Iteratefalse/IterateAllfalse异步接口减少 AIC 与 AIV 之间的消息交互、消除同步依赖并给出可复制的正反例代码、Workspace 配置要点以及仓库内的完整可运行样例帮助读者在算子开发中直接落地这一核间通信优化手段。一、背景MIX 场景下的核间同步开销在 MIX 场景AIC 与 AIV 混合编程中AIV 核需要调用 Matmul 的Iterate或IterateAll接口通过发送消息到 AIC 核来启动 Matmul 计算。每一次 Matmul 计算片C 分片的产出与取回都伴随着 Cube 核与 Vector 核之间的消息传递与同步等待。该问题在 asc-devkit 文档体系中属于高优先级优化项专门针对流水线调度Pipeline Scheduling场景。其核心矛盾在于同步方式Iteratetrue默认每次调用Iterate都会触发一次消息发送AIV 需要等待 AIC 计算完成后才能继续异步方式Iteratefalse仅第一次调用需要发送消息启动计算后续迭代不再发送消息AIV 无需逐次等待从而减少 Cube 与 Vector 核间交互显著降低核间通信开销。因此MIX 场景官方推荐使用Iteratefalse或IterateAllfalse异步接口前提是使用异步接口时必须设置 Workspace。下图分别展示了同步与异步两种方式下 AIV 向 AIC 发送消息启动 Matmul 计算的差异同步方式消息发送示意图异步方式消息发送示意图二、接口本质sync模板参数决定消息发送行为Iterate与IterateAll都是带模板参数sync的接口默认值为true同步。相关声明位于仓库公开接口文件 matmul.h 与 matmul.h// 计算一个 baseM * baseN 大小的 C 矩阵分片 // tparam sync: 设置为同步或异步模式 // param enPartialSum: 是否将 Iterate 的结果累加到 CO1 数据中 template bool sync true __aicore__ inline bool Iterate(bool enPartialSum false); // 计算 singleCoreM * singleCoreN 大小的完整 C 矩阵 // param waitIterateAll: 异步模式下是否通过 WaitIterateAll 等待 IterateAll 完成 template bool sync true __aicore__ inline void IterateAll( const GlobalTensorDstT gm, uint8_t enAtomic 0, bool enSequentialWrite false, bool waitIterateAll false, bool fakeMsg false);Iterate以baseM × baseN为基本分片粒度逐片计算配合GetTensorC逐片取回结果适合 AIV 侧边算边处理的流水线场景IterateAll一次性计算singleCoreM × singleCoreN的完整 C 矩阵适合整块输出的场景。两者语义一致sync true为同步模式sync false为异步模式。GetTensorC同样支持template bool sync true异步取数见 matmul.h异步场景下通常与Iteratefalse配套使用。2.1 客户端实现异步模式仅发送一次消息在 KFCKernel Function Call客户端实现 matmul_client.h 中Iterate的行为清晰地体现了仅第一次发送消息这一机制template bool sync true __aicore__ inline bool Iterate(bool enPartialSum false) { TRACE_START(TraceId::KFC_CLIENT_POST_MSG); if (unlikely(kfcMsg_.body.isFirstIter)) { cntIter_ 0; cOffset_ 0; curProcess 0; } else { if (cntIter_ mnIter_) { TRACE_STOP(TraceId::KFC_CLIENT_POST_MSG); return false; // 迭代结束 } if constexpr (!sync) { TRACE_STOP(TraceId::KFC_CLIENT_POST_MSG); return true; // 异步模式后续迭代直接返回不发送消息 } } if constexpr (!sync) { // 异步模式约束 ASSERT(cacheWorkspaceAddr ! 0); // 必须已通过 SetWorkspace 配置缓存地址 ASSERT(PhyPosIsUB(C_TYPE::pos)); // 异步模式仅支持输出到 UB } isSyncGetC sync; kfcMsg_.body.sync sync; kfcMsg_.body.cAddr reinterpret_castuint64_t(cacheWorkspaceAddr); ... PostMessageKFC_Enum::MMFUN_ITERATE, false(); // 同步模式每次迭代都会发送消息 SyncCubeWithVecA_TYPE::ibShare, B_TYPE::ibShare(); ... return true; }从源码可以清楚看到异步模式sync false除第一次isFirstIter外后续每次Iterate直接return true不会执行PostMessage即不再向 AIC 发送消息同步模式sync true每次迭代都会执行PostMessageKFC_Enum::MMFUN_ITERATE发送消息并执行SyncCubeWithVec同步这正是每次调用都会触发一次消息发送的代码级根因。2.2 服务器端实现异步模式下 C 分片直出在 KFC 服务器端 matmul_server_impl_3510.h 的Iterate处理逻辑中异步 Workspace 场景会走直出快速路径// 异步且配置了 workspace 时C 分片直接搬出无需逐次切换消息队列 if (unlikely(isIterate body.cAddr ! 0 hasSetWorkspace sync 0)) { mul.GetTensorC(cGlobal[offset], (uint8_t)(enAtomic), enSequentialWrite); offset offsetSize; CrossCoreSetFlagINTRA_MODE, PIPE_FIX( GetIntraFlagId(instID, static_castuint8_t(VEC_WAIT_INTRA_Enum::WAIT_FIXP), mul.GetSubBlockIdx())); continue; }同时GetOffsetSize 中专门注释说明For asynchronous Iterate, the offset must be used for address calculation and the size is baseM x baseN即异步 Iterate 需要按baseM × baseN计算偏移地址C 分片依次落入 Workspace 中不同的偏移位置。服务器端通过hasSetWorkspace即是否调用过SetWorkspace与sync字段共同判定是否走异步直出路径这进一步印证了使用异步接口必须设置 Workspace的硬性要求。三、反例与正例同步 vs 异步代码对比3.1 反例MIX 场景使用 Iterate 的同步方式以下代码在 MIX 场景中使用Iterate()即Iteratetrue同步方式每次循环迭代都会触发一次 AIV→AIC 的消息发送与核间同步TQueBindTPosition::CO2, TPosition::VECIN qVecIn; TQueBindTPosition::VECIN, TPosition::VECOUT qVecOut; mm.SetTensorA(gmA); mm.SetTensorB(gmB); int16_t scalar 2; while(mm.template Iterate()){ auto cInUB qVecIn.AllocTensorfloat(); mm.GetTensorC(cInUB); qVecIn.EnQue(cInUB); cInUB qVecIn.DeQuefloat(); auto cOutUB qVecOut.AllocTensorfloat(); Muls(cOutUB, cInUB, scalar, baseM*baseN); qVecIn.FreeTensor(cInUB); ... }3.2 正例MIX 场景使用 Iterate 的异步方式正例与反例相比仅有两处关键差异① 循环条件改为Iteratefalse② 循环前调用SetWorkspace配置临时空间TQueBindTPosition::CO2, TPosition::VECIN qVecIn; TQueBindTPosition::VECIN, TPosition::VECOUT qVecOut; mm.SetTensorA(gmA); mm.SetTensorB(gmB); mm.SetWorkspace(workspace, size);//其中workspace为临时空间的物理地址size为singleCoreM*singleCoreN大小的矩阵C占用的内存大小singleCoreM*singleCoreN*sizeof(float) int16_t scalar 2; while(mm.template Iteratefalse()){ auto cInUB qVecIn.AllocTensorfloat(); mm.GetTensorC(cInUB); qVecIn.EnQue(cInUB); cInUB qVecIn.DeQuefloat(); auto cOutUB qVecOut.AllocTensorfloat(); Muls(cOutUB, cInUB, scalar, baseM*baseN); qVecIn.FreeTensor(cInUB); ... }两段代码的对比要点总结如下对比维度同步方式Iteratetrue异步方式Iteratefalse消息发送次数每次迭代均发送一次仅第一次发送后续迭代不再发送核间同步每次迭代等待 AIC 计算完成无需逐次等待并行度更高是否必须设置 Workspace否必须SetWorkspace适用输出位置不限输出到 UBC 矩阵在 UB 上核间通信开销高低注意原文档明确指出使用异步接口时必须设置 Workspace否则会触发ASSERT(cacheWorkspaceAddr ! 0)断言失败见 matmul_client.h。四、Workspace 配置详解4.1SetWorkspace的两种调用形式SetWorkspace在公开接口 matmul.h 中提供两种重载// 形式一GM 裸指针 内存大小size 的单位为字节 // param addr: workspace on GM provided by user, GM address type // param size: 内存大小 template class T __aicore__ inline void SetWorkspace(__gm__ const T* addr, int size); // 形式二GlobalTensor推荐使用 // param addr: workspace on GM provided by user, GlobalTensor type template class T __aicore__ inline void SetWorkspace(GlobalTensorT addr);以原文档正例中的SetWorkspace(workspace, size)为例workspace临时空间的物理地址GM 地址sizesingleCoreM * singleCoreN * sizeof(float)即单核 Matmul 计算所需的整块 C 矩阵占用的内存大小字节数。在客户端实现 matmul_client.h 中GlobalTensor形式的SetWorkspace最终转换为裸指针形式并记录缓存地址template class T __aicore__ inline void SetWorkspace(GlobalTensorT addr) { ASSERT(addr.GetSize() 0); SetWorkspace(addr.GetPhyAddr(), addr.GetSize() * sizeof(T)); // size 按字节数传递 } template class T __aicore__ inline void SetWorkspace(__gm__ const T* addr, int size) { ASSERT(addr ! nullptr); cacheWorkspaceAddr reinterpret_castGM_ADDR(const_cast__gm__ T*(addr)); cOffset_ 0; }注意两点size 的单位从GlobalTensor形式的转换实现看addr.GetSize() * sizeof(T)传入的是字节数与原文档singleCoreM*singleCoreN*sizeof(float)的描述一致后续kfcMsg_.body.cAddr即使用该缓存地址作为 C 分片的落盘位置调用时机SetWorkspace必须在Iteratefalse之前调用且每个 Matmul 对象实例只需设置一次cOffset_会随迭代自动累加偏移。4.2 示例工程中的 Workspace 分配方式仓库示例 matmul_async_iterate.asc 展示了完整的 Workspace 获取与设置流程__global__ __aicore__ void matmul_async_iterate_custom( __gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* bias, __gm__ uint8_t* c, __kfc_workspace__ __gm__ uint8_t* workspace, __gm__ uint8_t* tilingGm) { TCubeTiling tiling; CopyTiling(tiling, tilingGm); MatmulKernelhalf, half, float, float matmulKernel; AscendC::TPipe pipe; REGIST_MATMUL_OBJ(pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, tiling); matmulKernel.Init(a, b, bias, c, workspace, tiling); matmulKernel.Process(pipe); }在内核Init中通过AscendC::GetUserWorkspace(workspace)获取用户可用的临时空间并按核偏移切分workspaceGlobal.SetGlobalBuffer( reinterpret_cast__gm__ CType*(AscendC::GetUserWorkspace(workspace)), tiling.M * tiling.N); ... workspaceGlobal workspaceGlobal[AscendC::GetBlockIdx() * tiling.singleCoreM * tiling.singleCoreN];随后在Process中调用matmulObj.SetTensorA(aGlobal, IS_TRANS_A); matmulObj.SetTensorB(bGlobal, IS_TRANS_B); if (IS_BIAS) { matmulObj.SetBias(biasGlobal); } matmulObj.SetWorkspace(workspaceGlobal); matmulObj.template Iteratefalse();即Workspace 由系统通过__kfc_workspace__参数和REGIST_MATMUL_OBJ宏注册每个核分得singleCoreM * singleCoreN大小的独立空间再以GlobalTensor形式传给SetWorkspace。五、IterateAllfalse异步场景与WaitIterateAll与Iterate逐片处理不同IterateAllfalse用于一次性输出整块 C 矩阵。当采用异步模式且输出为连续 GM 时需要配合waitIterateAll true与WaitIterateAll()使用。客户端 matmul_client.h 对WaitIterateAll的约束如下// Only support the mode that the IterateAll is asynchronous and GM output is continuous. // In discontinuous scenarios, the system stops responding. __aicore__ inline void WaitIterateAll() { ASSERT(!isSyncGetC); // Must be asynchronous mode ... }仓库示例 matmul_async_iterate_all.asc 给出了标准用法matmulObj.SetTensorA(aGlobal, IS_TRANS_A); matmulObj.SetTensorB(bGlobal, IS_TRANS_B); if (IS_BIAS) { matmulObj.SetBias(biasGlobal); } matmulObj.template IterateAllfalse(cGlobal, 0, false, true); // 第4个参数 waitIterateAll true matmulObj.WaitIterateAll(); matmulObj.End();要点IterateAllfalse(cGlobal, 0, false, true)的四个参数依次为C 矩阵 GlobalTensor、enAtomic是否使能原子操作0 表示关闭、enSequentialWrite是否顺序写、waitIterateAll是否等待完成异步模式下必须调用WaitIterateAll()等待 AIC 侧全部计算完成再释放/复用后续资源从源码注释看该模式仅支持GM 输出连续的场景若输出地址不连续可能导致系统挂起使用时需确保 C 矩阵在 GM 上连续排布。六、适用前提与注意事项综合原文档说明与仓库源码使用异步接口需要把握以下边界条件必须设置 Workspace异步模式依赖 Workspace 作为 C 分片的缓存中转Iteratefalse路径上有ASSERT(cacheWorkspaceAddr ! 0)硬校验matmul_client.h输出位置约束异步Iterate模式要求 C 矩阵输出到 UBASSERT(PhyPosIsUB(C_TYPE::pos))即C_TYPE的物理位置为 UBIterateAll 异步模式要求 GM 输出连续并配合WaitIterateAll()使用同步/异步的取舍当循环迭代次数较少、或 AIV 侧后续向量计算依赖单次 Matmul 结果强同步时同步模式逻辑更简单当分片迭代次数多、AIV 侧有独立的向量算子流水可并行时异步模式的核间通信收益更明显尾块处理异步模式下服务器端按baseM × baseN计算偏移若存在尾块tail需要像示例中那样通过SetTail(tailM, tailN)处理单核边界形状见 matmul_async_iterate.asc。七、仓库内完整样例与运行方式在 asc-devkit 仓库中与本文主题直接相关的可运行样例位于examples/01_simd_cpp_api/04_advanced_api/00_matmul/下matmul_async_iterateIteratefalseGetTensorCfalse异步直调样例输出到 VECINmatmul_async_iterate_allIterateAllfalseWaitIterateAll()异步整块输出样例matmul_fused融合场景中Iteratefalse异步调用示例batch_matmul_iterate_n_batchBatch Matmul 中SetWorkspace配合异步迭代的示例。以matmul_async_iterate样例为例其 README 给出的编译运行方式如下参考 README.mdsource ${install_path}/cann/set_env.sh # 配置 CANN 环境变量${install_path} 默认 /usr/local/Ascend mkdir -p build cd build; # 创建并进入 build 目录 cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # 编译工程默认 NPU 模式 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行生成的可执行程序 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证结果精度编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2/A3 系列dav-3510 对应 Ascend 950PR/950DT 系列# CPU 调试模式 cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # NPU 仿真模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j;注意切换编译模式前需清理 cmake 缓存如删除 build 目录下的CMakeCache.txt后重新 cmake。样例精度验证通过时终端输出test pass!。八、总结在 MIX 混合编程场景中Matmul 的同步Iterate/IterateAll调用会逐次触发 AIV→AIC 的消息发送与核间同步是算子流水线中的常见性能瓶颈。asc-devkit 提供Iteratefalse与IterateAllfalse异步接口其本质是利用 KFC 消息机制的首次发送、后续直出特性将核间通信次数从每分片一次降为整个计算一次代码改造成本极低仅需将Iterate()改为Iteratefalse或将IterateAll()改为IterateAllfalse并在调用前增加一次SetWorkspace收益明确减少 Cube 与 Vector 核间的消息交互与同步等待降低核间通信开销提升算子整体并行度边界清晰异步模式要求输出到 UB、必须设置 WorkspaceIterateAllfalse要求 GM 输出连续并配合WaitIterateAll()。开发者可直接参照本文正例与仓库内matmul_async_iterate、matmul_async_iterate_all样例完成改造并通过verify_result.py校验精度在保障正确性的前提下获得核间通信优化的收益。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐asc-devkit C_API 异步场景 Add 算子实战基于 asc_copy/asc_add/asc_sync 的矢量算子开发指南asc devkit C_API 异步场景 Add 算子实战基于 asc_copy/asc_add/asc_sync 的矢量算子开发指南 导读 本文以 asc人工智能深度学习算子库CANNAscendasc-devkit Matmul MIX 模式融合编程实践AIC/AIV 分工实现 Matmul LeakyRelu 单核函数融合asc devkit Matmul MIX 模式融合编程实践AIC/AIV 分工实现 Matmul LeakyRelu 单核函数融合 Ascend C 的人工智能深度学习算子库CANNAscendAscend C 核间同步实战基于 IBSet 与 IBWait 的多核数据依赖消除asc-devkit 样例解析Ascend C 核间同步实战基于 IBSet 与 IBWait 的多核数据依赖消除asc devkit 样例解析 本文基于 asc devkit 仓库中人工智能深度学习算子库CANNAscend上一篇TinyUSB USB PD 3.0实现Type-C电力传输协议开发详解下一篇asmttpd 源码架构分析模块化设计的汇编实践创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考