CANN asc-devkit Matmul IBShareB 直调样例解析：基于 L1 Buffer 的 B 矩阵复用与纯 Cube 模式实现

CANN asc-devkit Matmul IBShareB 直调样例解析：基于 L1 Buffer 的 B 矩阵复用与纯 Cube 模式实现

发布时间：2026/10/17 2:29:15
CANN asc-devkit Matmul IBShareB 直调样例解析：基于 L1 Buffer 的 B 矩阵复用与纯 Cube 模式实现 人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本篇文章围绕 asc-devkit 仓库中 matmul_ibshareB 样例 展开讲解如何在调用 Matmul 高阶 API 时开启 IBShare 功能让同一 AIC 对应的两个 AIV 在迭代计算中复用 L1 Buffer 上完全相同的 B 矩阵数据从而减少数据搬运开销。读完本文你将掌握 IBShareB 的两种实现路径默认 IBShare 模板与纯 Cube 模式、内核声明与调用要点、CMake 编译选项的完整用法并能从源码层面理解CFG_IBSHARE_NORM模板配置的真实含义与生效机制。一、IBShare 特性与样例定位1.1 什么是 IBShareIBShareInner Block Sharing是 Matmul 高阶 API 提供的一项优化特性当计算过程中多个计算单元同一 AIC 对应的两个 AIV在每次迭代时用到的 A 矩阵或 B 矩阵在 L1 Buffer 上的数据完全一致时可以复用这块 L1 Buffer 数据避免重复搬运从而降低 GM → L1 的数据传输开销。本样例MatmulIBShareBCustom聚焦仅 B 矩阵复用的场景调用 Matmul 高阶 API 时开启 B 矩阵的 IBShare 参数模板参数中对应isIBShare true计算过程中同一 AIC 对应的两个 AIV 在每次迭代时使用的 B 矩阵 L1 Buffer 数据保持一致。1.2 两种场景SCENARIO_NUM样例通过编译选项SCENARIO_NUM文档中描述为通过环境变量选择场景实际由 CMake 传给编译器的宏定义实现选择两种实现方式scenarioNum场景类型1默认默认实现使能 IBShare 模板CFG_IBSHARE_NORM2使能纯 Cube 模式 IBShareB1.3 约束条件仅 A 矩阵或仅 B 矩阵单独使能 IBShare 的场景要求复用的矩阵必须在 L1 Buffer 上全载。也就是说被复用的那侧矩阵需要完整放入 L1 Buffer才能保证多个 AIV 复用时读到的数据一致。1.4 支持的产品与 CANN 版本产品CANN 软件版本Ascend 950PR 950DT 系列产品 CANN 9.1.0Atlas A3 系列产品 CANN 9.0.0Atlas A2 系列产品 CANN 9.0.0二、目录结构与样例规格2.1 目录结构├── matmul_ibshareB │ ├── scripts │ │ ├── gen_data.py // 输入数据和真值数据生成脚本 │ │ └── verify_result.py // 真值对比脚本 │ ├── CMakeLists.txt // 编译工程文件 │ ├── data_utils.h // 数据读入写出函数 │ ├── matmul_ibshareB.asc // Ascend C 样例实现 调用样例 │ └── README.md // 样例说明文档2.2 样例规格本样例的矩阵规模为M 64N 256K 384项说明样例类型OpTypeMatmulIBShareBCustom输入 ashape [M, K]数据类型 halfND 格式isTrans false输入 bshape [K, N]数据类型 halfND 格式isTrans false输出 cshape [M, N]数据类型 floatND 格式核函数名matmul_ibshareb_custom数据生成脚本 gen_data.py 中对应配置为MatmulGenData(64, 256, 384, 1, False, False, False, float16_float32)即 batch1、A/B 均不转置、无 bias、A/B 为 fp16、C 为 fp32输入数据用np.random.uniform(-1, 1)随机生成并转为 fp16真值由 TensorFlow 的tf.matmul计算得到。三、内核实现两种 IBShareB 使能方式核心实现位于 matmul_ibshareB.asc。内核关键步骤是创建 Matmul 对象并设置 B 矩阵的 IBShare 参数为 true具体通过MatmulType模板的最后一个布尔模板参数体现见源码第 73 行using B_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, BType, false, LayoutMode::NONE, true;MatmulType的模板参数依次为位置GM、格式ND、数据类型、是否转置false、布局模式NONE以及最后的isIBSharetrue——正是这个true使 B 矩阵进入 IBShare 复用路径。3.1 方式一默认实现使能 IBShare 模板在matmul_ibshareB.asc中当SCENARIO_NUM ! 2时源码第 80 行Matmul 对象使用CFG_IBSHARE_NORM模板创建#include lib/matmul_intf.h using A_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, AType; using B_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, BType, false, LayoutMode::NONE, true; using C_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, CType; using BIAS_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, BiasType; AscendC::MatmulA_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, CFG_IBSHARE_NORM matmulObj;该方式为非纯 Cube 模式AICCube 单元与 AIV向量单元按比例协同计算内核入口使用__aicore__源码第 178 行调用时SetDim设置的是AIV:AIC 组合的核数。3.2 方式二纯 Cube 模式 IBShareB纯 Cube 模式通过宏ASCENDC_CUBE_ONLY开启。该宏必须在#include lib/matmul_intf.h之前定义。在matmul_ibshareB.asc中该宏正是由编译选项驱动的条件定义源码第 17-19 行#if SCENARIO_NUM 2 #define ASCENDC_CUBE_ONLY #endif #include lib/matmul_intf.h使用纯 Cube 模式时Matmul 对象改用常规配置CFG_NORM创建源码第 78 行#define ASCENDC_CUBE_ONLY // 设置 ASCENDC_CUBE_ONLY 宏 #include lib/matmul_intf.h using A_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, AType; using B_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, BType, false, LayoutMode::NONE, true; using C_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, CType; using BIAS_TYPE AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, BiasType; AscendC::MatmulA_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, CFG_NORM matmulObj;需要说明方式二下 B 矩阵的 IBShare 使能仍然由B_TYPE中isIBShare true承担CFG_NORM只是不启用 IBShare 归一化模板由于ASCENDC_CUBE_ONLY排除了 AIV 参与纯 Cube 模式下计算完全由 AIC 完成SetDim设置的是AIC 核数。3.3 核函数入口声明默认实现方式一内核函数入口为extern C __global__ __aicore__ void matmul_ibshareb_custom(...)源码第 178 行。纯 Cube 模式方式二内核函数入口需显式指定__cube__源码第 174 行extern C __global__ __cube__ void matmul_ibshareb_custom(...)。两种入口的参数一致A 矩阵地址、B 矩阵地址、bias 地址、C 矩阵地址、workspace 地址、tiling 地址。3.4 内核处理流程源码级MatmulKernel模板类源码第 63 行起的Process()展示了完整的调用链通过AscendC::GetBlockIdx()判断当前核是否超出usedCoreNum超出则直接返回源码第 128 行matmulObj.SetTensorA(aGlobal, isTransA)、matmulObj.SetTensorB(bGlobal, isTransB)绑定 GM 侧张量若tiling.isBias为真则matmulObj.SetBias(biasGlobal)matmulObj.IterateAll(cGlobal)完成全部迭代计算matmulObj.End()收尾。此外Init()中通过CalcOffset()源码第 141 行起按blockIdx对 M/N 维度进行核心划分mSingleBlocks (M singleCoreM - 1) / singleCoreMmCoreIndex blockIdx % mSingleBlocks、nCoreIndex blockIdx / mSingleBlocks据此计算 A、B、C、bias 在 GM 中的偏移。而REGIST_MATMUL_OBJ(pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, tiling)源码第 189 行将 Matmul 对象注册到 TPipe 流水workspace 指针用于承载 Matmul 内部所需的临时空间。3.5 调用实现 内核调用符使用内核调用符调用核函数两种方式的网格grid规模不同matmul_ibshareb_customtilingData.usedCoreNum / MIX_RATIO, 0, stream(x1, x2, bias, y, workspaceDevice, tilingDevice); // 方式一非纯 Cube 模式SetDim 设置为 AIV:AIC 组合的核数 matmul_ibshareb_customtilingData.usedCoreNum, 0, stream(x1, x2, bias, y, workspaceDevice, tilingDevice); // 方式二纯 Cube 模式SetDim 设置为 AIC 核数源码中定义constexpr int32_t MIX_RATIO 2; // AIC:AIV1:2源码第 29 行即在默认的 AIC/AIV 混合模式下总核数 AIC 核数 ×1 MIX_RATIO因此网格大小取usedCoreNum / MIX_RATIO纯 Cube 模式下则直接使用usedCoreNum即 AIC 核数。主程序中两种场景的 launch 由#if SCENARIO_NUM 2分支分别编译源码第 280-285 行。3.6 Tiling 生成源码级GenerateTiling()源码第 196 行起展示了 Matmul 直调场景下 tiling 的完整生成流程通过platform_ascendc::PlatformAscendCManager::GetInstance()获取平台信息CPU 调试模式下按__NPU_ARCH__分别指定 Ascend910B1 / Ascend950PR_9589 实例创建matmul_tiling::MultiCoreMatmulTiling cubeTiling并SetDim(USED_CORE_NUM)样例中USED_CORE_NUM 2源码第 33 行SetAType/SetBType分别声明 Afp16、ND、非转置、Bfp16、ND、非转置SetCType/SetBiasType声明 C 与 bias 为 fp32SetOrgShape/SetShape设置 M、N、KEnableBias(false)、SetBufferSpace(-1, -1, -1)使用默认 buffer 空间最后GetTiling(tilingData)产出TCubeTiling失败时打印 Generate tiling failed. 并返回空。主程序MatmulOp()中还会通过aclrtMalloc为 workspace 与 tiling 分配设备内存aclrtMemcpy将 tiling 从 HOST 拷贝到 DEVICE 后启动内核。四、编译与运行4.1 配置环境变量请根据当前环境上 CANN 开发套件包的安装方式 配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。4.2 样例执行在样例根目录examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_ibshareB下执行SCENARIO_NUM1 # 默认实现使能 IBShare 模板 mkdir -p build cd build; # 创建并进入 build 目录 cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 -DSCENARIO_NUM$SCENARIO_NUM ..;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_MODEcpu或-DCMAKE_ASC_RUN_MODEsim参数SCENARIO_NUM1 cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 -DSCENARIO_NUM$SCENARIO_NUM ..;make -j; # CPU 调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 -DSCENARIO_NUM$SCENARIO_NUM ..;make -j; # NPU 仿真模式注意切换编译模式或切换 SCENARIO_NUM前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新运行 cmake。4.3 编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201、dav-3510NPU 架构dav-2201 对应 Atlas A2 系列产品与 Atlas A3 系列产品dav-3510 对应 Ascend 950PR 950DT 系列产品SCENARIO_NUM1默认、2场景编号1默认实现使能 IBShare 模板2使能纯 Cube 模式 IBShareB这些选项在 CMakeLists.txt 中均有对应实现CMAKE_ASC_RUN_MODE与CMAKE_ASC_ARCHITECTURES以 CACHE STRING 形式定义SCENARIO_NUM同样以 CACHE STRING 声明且非法值非 1 或 2会触发FATAL_ERRORtarget_compile_definitions(demo PRIVATE SCENARIO_NUM${SCENARIO_NUM})将宏传给编译器从而驱动matmul_ibshareB.asc中的#if分支。可执行程序demo由单个源文件matmul_ibshareB.asc构建链接tiling_api、register、platform、m、dl等库编译选项通过--npu-arch${CMAKE_ASC_ARCHITECTURES}指定目标 NPU 架构。4.4 数据生成与精度验证gen_data.py在./input、./output目录下生成x1_gm.binA 矩阵fp16、x2_gm.binB 矩阵fp16与golden.bin真值fp32。真值由 TensorFlowtf.matmul在 fp32 下计算并做 ±65504 饱和截断due_overflow避免 fp16 溢出影响对比。verify_result.py使用np.isclose按相对容差rtol1e-3、绝对容差atol1e-3逐元素比对统计错误比例并与阈值1e-3比较对比通过时打印test pass!否则打印前 100 个差异元素索引与错误比例并退出码置 1。4.5 执行结果执行结果如下说明精度对比成功test pass!五、从源码理解 CFG_IBSHARE_NORMIBShare 模板的底层机制为帮助读者知其所以然这里结合仓库源码揭示CFG_IBSHARE_NORM的真实定义与生效路径。5.1 模板常量定义在 include/adv_api/matmul/tiling.h 中constexpr MatmulConfig CFG_NORM GetNormalConfig(); constexpr MatmulConfig CFG_IBSHARE_NORM GetIBShareNormConfig();GetIBShareNormConfig()tiling.h返回的配置核心是doIBShareNorm true同时保留归一化融合相关能力enUnitFlag true、enableReuse true、enableUBReuse true等。Matmul模板的最后一个模板参数MM_CFG默认即为CFG_NORM见 include/adv_api/matmul/matmul_intf.h因此方式一显式传入CFG_IBSHARE_NORM等价于把 Matmul 对象切换到 IBShare 归一化优化版本。5.2 模式判定与分发在 impl/adv_api/detail/matmul/utils/matmul_config_utils.h 中定义了MatmulMode枚举enum class MatmulMode { NORMAL 0, // Basic version MULTI_DATA_LOAD, // Multiple data preloading BASIC_BLOCK, // Block computation using basic blocks IBSHARE_NORM, // Inner block sharing and normalization fusion optimization };GetMatmulMode()matmul_config_utils.h按doNorm → doBasicBlock → doMDL → doIBShareNorm的优先级判定当CFG_IBSHARE_NORM的doIBShareNorm true时模式落入MatmulMode::IBSHARE_NORM从而触发 IBShare 归一化这一独立实现分支。5.3 参数特化与实现入口在 impl/adv_api/detail/matmul/utils/matmul_param.h 中MatmulParams针对GetMatmulMode(CFG_IBSHARE_NORM)即MatmulMode::IBSHARE_NORM做了模板特化为该模式提供独立的参数结构定义支撑内核运行时的分派。这也印证了IBShare 不是简单的开关而是一条独立的模板实例化与实现路径——这正是方式一与方式二CFG_NORM 纯 Cube在SCENARIO_NUM编译分支下并行存在的原因。六、与 IBShareAB 样例的对比同目录下的 matmul_ibshareAB 样例展示的是 A、B 矩阵同时复用的场景A、B 均使能 IBShare 时不对 K 列切分计算均不使能时按 K 列切分计算通过对比两种场景运行时间可评估特性收益。而本样例IBShareB是仅 B 复用的最小化验证二者共享相同的MatmulType..., isIBShare模板机制与SCENARIO_NUM编译选择方式适合作为理解 IBShare 全貌的递进素材。若要对比不使能 IBShare的基线行为可直接参考 matmul 基础样例 或matmul_ibshareAB的 SCENARIO_NUM2 分支。七、总结本文以 asc-devkit 的matmul_ibshareB直调样例为线索完整覆盖了特性理解IBShareB 使同一 AIC 下两个 AIV 复用 L1 Buffer 上的 B 矩阵数据降低数据搬运开销约束是复用矩阵需在 L1 全载两种实现路径默认CFG_IBSHARE_NORM模板混合 AIC/AIV与ASCENDC_CUBE_ONLY纯 Cube 模式CFG_NORM二者通过SCENARIO_NUM编译分支切换内核入口分别对应__aicore__与__cube__调用网格大小分别为usedCoreNum / MIX_RATIO与usedCoreNum端到端编译运行环境变量配置、cmake/make构建、数据生成、内核执行与精度验证的完整命令链以及CMAKE_ASC_RUN_MODE、CMAKE_ASC_ARCHITECTURES、SCENARIO_NUM三个编译选项的取值与含义源码级原理从 tiling.h 的配置常量、matmul_config_utils.h 的模式判定到 matmul_param.h 的参数特化还原了 IBShare 归一化路径的完整分发链路。对需要在大规模矩阵乘法中减少 GM→L1 搬运、提升迭代效率的开发者而言本样例是一份可直接复制的IBShare 最小可用范式也适合在此基础上继续对比matmul_ibshareAB、matmul_a2b2_share等进阶样例构建完整的 Matmul 数据复用优化知识体系。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐asc-devkit Matmul 高阶 API 实战IBShareB 直调样例原理与实现asc devkit Matmul 高阶 API 实战IBShareB 直调样例原理与实现 导读 本篇技术指南以 asc devkit 仓库中的 matmul人工智能深度学习算子库CANNAscendasc-devkit Batch Matmul 批量矩阵乘直调样例解析基于 IterateBatch 高阶 API 的实现与调优实践asc devkit Batch Matmul 批量矩阵乘直调样例解析基于 IterateBatch 高阶 API 的实现与调优实践 本篇文章以 asc de人工智能深度学习算子库CANNAscendBatch Matmul 复用 Bias 矩阵直调样例深度解析基于 CANN AscendC 的批量矩阵乘 Bias 复用实现Batch Matmul 复用 Bias 矩阵直调样例深度解析基于 CANN AscendC 的批量矩阵乘 Bias 复用实现 批量矩阵乘Batch Mat人工智能深度学习算子库CANNAscend上一篇Composio E2E 测试工具链解析在隔离 Docker 环境中运行 composio/core 与 CLI 端到端测试下一篇PaddleOCR PP-StructureV2 智能文档分析系统全解析版面分析、表格识别、关键信息抽取与版面恢复实战创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考