企业官网建设流程全解析

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

导读

Matmul API 默认的 MIX 模式依赖消息机制，由 AIV 通过 KFC 消息驱动 AIC 执行矩阵计算；而双主模式（MixDualMaster）让 AIC 与 AIV 各自独立运行代码，不再依赖消息驱动，从而降低通信开销、提升性能。本文以 asc-devkit 仓库中的matmul_mixdualmaster样例为主体，讲解双主模式的开启条件、enableMixDualMaster配置项的底层实现、样例关键代码（含 IBSHARE 参数启用、尾块 Tail 处理、多核偏移计算、Tiling 生成），并给出完整的编译、运行与精度验证流程。读完本文，你将掌握在 AscendC/CANN 工程中如何编写并调测一个双主模式 Matmul 算子。

双主模式是什么：与 MIX 模式的区别

在 Ascend 昇腾 AI 处理器的 Cube（AIC）与 Vector（AIV）异构架构中，Matmul 算子通常以 MIX 模式运行，即“矩阵计算 + 矢量计算”混合模式。默认的 MIX 模式下，AIV 作为主控，通过消息（KFC message）机制驱动 AIC 执行矩阵乘计算，AIC 与 AIV 之间存在消息同步与分发开销。

双主模式（MixDualMaster）则改变这一协作方式：AIC 和 AIV 独立运行各自代码，不依赖消息驱动，双方并行推进计算，从而减少消息传递带来的等待与调度开销，获得更优性能。样例根目录 README 明确指出：使能双主模式的 Matmul 样例中，AIC 和 AIV 独立运行代码，不依赖消息驱动，用于提升性能。

这一设计理念在源码中也有直接印证：模板配置项enableMixDualMaster默认值为false，定义于 include/adv_api/matmul/matmul_config.h#L184；而在双主模式使能后，KFC 注册与通信路径会切换为无消息依赖的独立执行方式（参见 impl/adv_api/detail/kfc/kfc_register_obj.h 中围绕asEnableMixDualMaster的分支处理）。

开启双主模式的条件

双主模式并非所有场景都可用，只有满足以下条件之一时才能开启（对应样例 README.md）：

• 核函数类型为 MIX，且 AIC 核数 : AIV 核数 = 1:1；
• 核函数类型为 MIX，且 AIC 核数 : AIV 核数 = 1:2，且 A 矩阵和 B 矩阵同时使能 IBSHARE 参数。

注意：双主模式下 A/B/Bias 矩阵只支持从 Global Memory（GM）搬入；计算结果只支持通过IterateAll接口输出到 GlobalTensor 或 LocalTensor，不能调用GetTensorC、Iterate、IterateBatch、SetWorkspace等接口。同一算子中所有 Matmul 对象的enableMixDualMaster取值必须保持一致。这些约束在 docs/zh/api/SIMD-API/adv_api/cube_compute/Matmul_Kernel/MatmulConfig.md 中有完整说明，并在 impl/adv_api/detail/matmul/kfc/matmul_server_aux.h 的if constexpr编译期断言中落地（例如SetTensorA localTensor not support when enableMixDualMaster is enabled）。

样例支持的产品及 CANN 版本

样例工程结构

样例位于 examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_mixdualmaster 目录，核心文件如下：

关键代码解析：如何使能双主模式

1. 通过 MatmulConfig 使能enableMixDualMaster

样例通过自定义MatmulConfig开启双主模式（matmul_mixdualmaster.asc#L42-L47）：

__aicore__ inline constexpr MatmulConfig GetCustomCFG() { auto mmCfg = CFG_NORM; // 以默认 Norm 模板配置为基底 mmCfg.enableMixDualMaster = true; // 使能双主模式 return mmCfg; }

其中CFG_NORM是库提供的默认配置常量，定义于 include/adv_api/matmul/tiling.h#L398；enableMixDualMaster在 include/adv_api/matmul/matmul_config.h#L184 中默认值为false。将模板第五个参数指定为CUSTOM_CFG即可让整个 Matmul 对象运行在双主模式：

constexpr static MatmulConfig CUSTOM_CFG = GetCustomCFG(); AscendC::Matmul< AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, AType, IS_TRANS_A, LayoutMode::NONE, true>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BType, IS_TRANS_B, LayoutMode::NONE, true>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, CType>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BiasType>, CUSTOM_CFG> matmulObj;

代码注释也明确说明：// Enable IBSHARE parameters for matrix A and matrix B——MatmulType末尾的模板参数true即为 A、B 矩阵的 IBSHARE 使能位，这与 README 中“1:2 核数比需 A/B 同时使能 IBSHARE”的条件相呼应。IBSHARE 的语义是复用同一份 A 或 B 矩阵数据于 L1 Buffer，与之配套的还有doIBShareNorm配置与CFG_IBSHARE_NORM模板常量（include/adv_api/matmul/tiling.h#L401）。

2. 注册 Matmul 对象并完成初始化

REGIST_MATMUL_OBJ宏将 Matmul 对象注册到 TPipe 上，并绑定系统 Workspace 与 Tiling 数据（matmul_mixdualmaster.asc#L152-L164）：

__global__ __cube__ void matmul_mixdualmaster_custom( __gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* bias, __gm__ uint8_t* c, __gm__ uint8_t* workspace, __gm__ uint8_t* tilingGm) { TCubeTiling tiling; CopyTiling(&tiling, tilingGm); AscendC::TPipe pipe; MatmulKernel<half, half, float, float> matmulKernel; REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, &tiling); matmulKernel.Init(a, b, bias, c, tiling); matmulKernel.Process(); }

CopyTiling将 GM 侧 Tiling 数据按 32 位字逐一搬运到TCubeTiling结构体，供设备侧使用。REGIST_MATMUL_OBJ的底层实现位于 impl/adv_api/detail/kfc/kfc_register_obj.h，其中会根据enableMixDualMaster的编译期常量切换 KFC 客户端构造方式（如KfcCommClient的初始化与消息路径选择），这正是“不依赖消息驱动”的实现入口之一。

3. 多核偏移计算与尾块处理

样例按 M/N 两个维度把矩阵切分到多个核：CalcOffset以mSingleBlocks作为 M 方向分核数，blockIdx % mSingleBlocks得到 M 方向核号、blockIdx / mSingleBlocks得到 N 方向核号，再据此计算 A/B/C/Bias 的 GM 偏移（matmul_mixdualmaster.asc#L130-L150）：

auto mSingleBlocks = (tiling.M + tiling.singleCoreM - 1) / tiling.singleCoreM; // split M into mSingleBlocks cores mCoreIndex = blockIdx % mSingleBlocks; nCoreIndex = blockIdx / mSingleBlocks; offsetA = mCoreIndex * tiling.Ka * tiling.singleCoreM; if (IS_TRANS_A) { offsetA = mCoreIndex * tiling.singleCoreM; } offsetB = nCoreIndex * tiling.singleCoreN; if (IS_TRANS_B) { offsetB = nCoreIndex * tiling.Kb * tiling.singleCoreN; } offsetC = mCoreIndex * tiling.N * tiling.singleCoreM + nCoreIndex * tiling.singleCoreN; offsetBias = nCoreIndex * tiling.singleCoreN;

Process中先按usedCoreNum过滤超发核，再针对 M/N 方向不足一个单核分片的情况调用SetTail(tailM, tailN)设置尾块，随后依次SetTensorA、SetTensorB、SetBias（tiling.isBias为真时），最后以IterateAll(cGlobal)+End()完成计算（matmul_mixdualmaster.asc#L106-L128）。注意，双主模式下计算入口固定为IterateAll，这与前文所述 API 约束一致。

4. Tiling 生成

宿主侧通过matmul_tiling::MultiCoreMatmulTiling生成 Tiling：SetDim使用GetCoreNumAic()（即 AIC 核数）作为维度，SetAType/SetBType指定 GM 侧 ND 布局的 FP16 输入，SetCType/SetBiasType指定 FP32 输出与偏置，SetOrgShape/SetShape设置 M/N/K（12288×256×128），EnableBias使能偏置，SetBufferSpace(-1, -1, -1)表示使用默认 Buffer 空间，最后GetTiling生成并SaveToBuffer序列化（matmul_mixdualmaster.asc#L166-L192）。

5. 数据规模与精度语义

样例采用half（FP16）输入、float（FP32）输出：A 为 12288×128 的 FP16 矩阵，B 为 128×256 的 FP16 矩阵，Bias 为长度 256 的 FP32 向量，输出 C 为 12288×256 的 FP32 矩阵（matmul_mixdualmaster.asc#L24-L29）。gen_data.py中 golden 的计算方式为：np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32)).astype(np.float32) + bias_gm，即先把 FP16 输入提升到 FP32 后做精确矩阵乘再加偏置，作为比对基准。

编译运行与验证

1. 配置环境变量

根据当前环境上 CANN 开发套件包的安装方式配置环境变量：

source ${install_path}/cann/set_env.sh

说明：${install_path}为 CANN 包安装目录，未指定安装目录时默认安装至/usr/local/Ascend下。

2. 样例执行

在样例根目录examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_mixdualmaster下依次执行：

mkdir -p build && cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译生成的可执行程序，执行样例 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确

使用CPU 调试或NPU 仿真模式时，添加-DCMAKE_ASC_RUN_MODE=cpu或-DCMAKE_ASC_RUN_MODE=sim参数即可，例如：

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

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

3. 编译选项说明

这两个选项在 CMakeLists.txt 中以 CACHE STRING 形式定义，并通过find_package(ASC REQUIRED)引入 ASC 编译器工具链；target_compile_options中的--npu-arch=${CMAKE_ASC_ARCHITECTURES}将架构透传给 ASC 编译器。可执行程序链接tiling_api、register、platform等库。

4. 执行结果

执行成功后输出以下内容，说明精度对比通过：

test pass!

该判定来自 scripts/verify_result.py：其使用np.isclose以相对容差rtol=1e-3、绝对容差atol=1e-3比对输出与 golden，统计误差率error_ratio，当误差率不超过1e-3时打印test pass!并返回成功；若超过则打印前 100 个不一致元素的下标、期望值、实际值与相对误差，便于定位精度问题。

运行机理补充：源码视角的双主模式实现

从源码结构可以进一步理解双主模式在库内部的处理路径：

• 配置声明：MatmulConfig中enableMixDualMaster默认false（include/adv_api/matmul/matmul_config.h#L184），并通过GetNormalConfig()等构造函数在 include/adv_api/matmul/tiling.h 中生成CFG_NORM等预置常量。
• 编译期分流：Matmul 客户端与服务端大量使用if constexpr (ToMatmulConfig(MM_CFG).enableMixDualMaster)做编译期分支（include/adv_api/matmul/matmul_client.h、impl/adv_api/detail/matmul/kfc/matmul_server_aux.h），在双主模式下编译期剔除消息驱动相关代码路径，同时对SetTensorA(localTensor)、SetWorkspace、Iterate、GetTensorC、IterateBatch等不兼容接口直接以ASSERT形式拒绝（例如"IterateAll localTensor not support when enableMixDualMaster is enabled"）。
• 共享对象判定：IsSharedObj(MatmulConfig)返回!mmCFG.enableInit || mmCFG.enableMixDualMaster（impl/adv_api/detail/matmul/utils/matmul_config_utils.h#L176），即双主模式对象按共享对象管理，进一步印证其独立并行的运行模型。
• 注册与一致性校验：KFC 注册层（impl/adv_api/detail/kfc/kfc_register_obj.h）会汇总所有 Matmul 对象的enableMixDualMasterValue，并通过static_assert（"enableMixDualMaster must be consistent for all cube objects."）保证同一算子内取值一致。

总结与实践建议

• 何时使用双主模式：当算子的 Matmul 计算具备 MIX 核函数类型，且 AIC:AIV 核数比为 1:1，或 1:2 且 A/B 同时开启 IBSHARE 时，可通过enableMixDualMaster = true让 AIC/AIV 独立运行，摆脱消息驱动的同步开销，换取性能提升。
• 代码改动量小：仅需在MatmulConfig中打开一个开关、将自定义配置作为模板参数传入，并把计算入口固定为IterateAll，其余数据搬运、偏移计算、Tiling 生成流程与普通 Matmul 样例一致。
• 注意约束：A/B/Bias 仅支持 GM 输入、结果只能经IterateAll输出、同一算子内所有 Matmul 对象配置必须一致；切换CMAKE_ASC_RUN_MODE前需清理 CMake 缓存。
• 验证闭环：gen_data.py生成输入与 golden，verify_result.py以 1e-3 容差完成精度校验，test pass!即为最终成功标志。

如需进一步了解enableMixDualMaster与其他 Matmul 模板参数（如scheduleType、enableDoubleCache、isA2B2Shared）的配合约束，可查阅 MatmulConfig 参数说明；相关不兼容接口的逐条说明散见于对应 API 文档（如 AsyncGetTensorC、GetBatchC 等）。

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