Blaze Matmul 工程模板指南【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills定位提供基于 Blaze 库开发 matmul 类算子的参考骨架示例明确目录结构、编码风格与规范。适用架构DAV_3510工程模板位置ascendc-direct-invoke-templateskill 的references/matmul_blaze_template/§1 Blaze 编程概述1.1 Blaze 是什么Blaze 是 Ascend C 的高层 matmul 编程库封装了底层 tensor_api 和硬件指令提供可组合的组件BlockMmad、BlockScheduler、Kernel使开发者能够以声明式方式构建高性能矩阵乘法算子。1.2 四层抽象Launcherhost 端 ↓ 选择组件 传递参数 Kerneldevice 端入口 ↓ 组装类型链 BlockMmad BlockScheduler核心组件 ↓ 调用 tensor_api硬件指令封装Launcherhost 端 C 程序负责 tiling 计算、内存管理、kernel 启动Kerneldevice 端入口函数__global__ __aicore__ __cube__组装类型链驱动计算BlockMmad数据搬运 MMAD 计算的核心组件管理 L1/L0 流水BlockSchedulertile 分配到多核的调度器Serpentine 遍历DispatchPolicy流水策略NO_FULL_LOAD / A_FULL_LOADtensor_api硬件指令封装Mmad、CopyGM2L1 等一般不直接修改1.3 关键组件职责组件职责来源Launcherhost 端程序负责 tiling、内存管理、kernel 启动用户开发Kerneldevice 端入口函数组装类型链驱动计算用户开发BlockMmad数据搬运 MMAD 计算L1/L0 流水Blaze 库或 blaze_customBlockSchedulertile 分配到多核Serpentine 遍历Blaze 库或 blaze_customDispatchPolicy流水策略NO_FULL_LOAD / A_FULL_LOADBlaze 库或 blaze_customProblemShape问题规模定义m, n, k, batchtensor_api§2 工程目录结构2.1 整体目录树matmul_blaze_template/ ├── test_matmul_blaze.cpp # C 直调 launcherhost 端入口 ├── run.sh # 一键编译运行验证 ├── CMakeLists.txt # 构建配置 ├── op_kernel/ # Kernel 层device 端 │ ├── matmul_blaze_kernel.h # Kernel 模板函数 │ ├── matmul_blaze_kernel.cpp # Wrapperextern C 包装 │ └── include/ │ ├── blaze/ # [拉取] Blaze 库 │ ├── tensor_api/ # [拉取] tensor_api │ └── blaze_custom/ # [可选] 自定义 Blaze 扩展模块 ├── op_tiling/ # Tiling 层host 端 │ ├── matmul_tiling_data.h # TilingData POD 结构体 │ ├── matmul_tiling_stub.h # Mock tiling 实现仅用于示例 │ └── matmul_constant.h # Tiling 相关常量 ├── op_extension/ # [可选] PyTorch 接入层 │ ├── matmul_blaze_torch.cpp # Torch extension 实现 │ ├── register.cpp # TORCH_LIBRARY 注册 │ └── ops.h # 函数声明 ├── common/ # 公共工具 │ ├── acl_utils.h # ACL 会话 RAII │ ├── common_utils.h # 参数解析、CeilDiv 等 │ └── io_utils.h # 文件 I/O └── scripts/ # 测试脚本 ├── gen_data_mxfp8.py # mxfp8 数据生成 golden 计算 ├── gen_data_mxfp4.py # mxfp4 数据生成 golden 计算 └── verify_result.py # 精度验证2.2 关键目录说明op_kernel/include/依赖管理blaze/和tensor_api/是外部依赖从 ops-tensor 仓库拉取git clone --depth 1 https://gitcode.com/cann/ops-tensor.git /tmp/ops-tensor cp -r /tmp/ops-tensor/include/blaze op_kernel/include/ cp -r /tmp/ops-tensor/include/tensor_api op_kernel/include/ rm -rf /tmp/ops-tensorrun.sh首次运行时会自动检测并拉取。blaze_custom/自定义扩展本模板提供两条开发路径Blaze 库路径本模板示例所有组件来自blaze/gemm/适用于标准 matmul 场景blaze_custom 路径自定义模块在blaze_custom/开发适用于需要特殊功能的场景两条路径的组件接口不同不能混用。当前blaze_custom/为空目录.gitkeep占位。op_tiling/边界说明当前matmul_tiling_stub.h是 mock 实现返回固定参数baseM128, baseN128, baseK128仅用于示例演示。正式开发时必须根据BlockScheduler::Params结构体定义开发完整的 tiling 引擎确保TilingData字段与Params字段映射正确。op_extension/PyTorch 接入可选提供torch.ops.npu.matmul_blaze()接口使用TORCH_LIBRARY_FRAGMENT注册。通过 wrapper 函数matmul_blaze_kernel.cpp调用 kernel。本模板使用的组件本模板使用 Blaze 库路径组件来自op_kernel/include/blaze/gemm/KernelGemmUniversal通用 kernel 模板BlockMmadBlockMmad9 个模板参数含 BiasType LayoutBiasBlockSchedulerBlockSchedulerQuantBatchMatmulV3量化 batch matmul 调度器DispatchPolicyMatmulWithScaleMxMX 量化策略参考op_kernel/matmul_blaze_kernel.h第 31-42 行2.3 文件依赖关系test_matmul_blaze.cpp (launcher) ↓ includes op_kernel/matmul_blaze_kernel.h ↓ includes blaze/gemm/kernel/kernel_qbmm_mx.h ↓ includes blaze/gemm/block/block_mmad_qbmm_mx.h blaze/gemm/block/block_scheduler_qbmm.h tensor_api/tensor.h§3 开发流程概览Step 1: 工程初始化复制模板目录重命名为目标算子名修改 CMakeLists.txt目标名、include 路径、编译选项--npu-archdav-3510拉取 blaze tensor_api 依赖run.sh首次运行自动完成确认op_kernel/include/blaze/和tensor_api/目录存在参考CMakeLists.txtStep 1.5: 根据 DESIGN.md 改造样例工程结构matmul_blaze_template/是 Blaze 样例工程不是目标算子的最终交付结构。复制模板后必须先根据DESIGN.md的“工程目录目标设计”完成目录和文件改造再进入 Kernel / Tiling / 测试逻辑开发。改造原则原则要求以 Blaze 样例结构为基础目标工程继承根目录 launcher、op_kernel/、op_tiling/、op_extension/、common/、scripts/等结构不引入 Add/Vector 模板的op_host/目录先重命名再开发样例业务文件必须重命名为目标算子语义或明确删除不要在样例文件旁边叠加一套新文件后保留旧文件内容必须改写算子名、kernel 名、launch wrapper、PyTorch 注册名、dtype、shape、输入输出、转置语义、tiling、golden、verify 均需与DESIGN.md一致CMake 必须收敛add_executable/add_library只纳入目标算子源文件未纳入 target 的样例业务源文件不应继续留在工程中外部依赖例外op_kernel/include/blaze/和op_kernel/include/tensor_api/是外部依赖可保留其多 dtype / 多场景代码自定义扩展位置需要扩展 Blaze 时只在op_kernel/include/blaze_custom/下新增或改造模块常见样例文件改造映射样例文件/目录常见处理说明test_matmul_blaze.cpp重命名并改写为{operator_name}.cpp或按设计删除根目录直调 launchertest_matmul_blaze_torch.py重命名并改写为test_{operator_name}_torch.py或按设计删除Blaze 模板的 PyTorch 测试脚本在根目录op_kernel/matmul_blaze_kernel.h重命名并改写为op_kernel/{operator_name}_kernel.h或按设计删除Kernel 类型链与入口声明op_kernel/matmul_blaze_kernel.cpp重命名并改写为op_kernel/{operator_name}_kernel.cpp或按设计删除Kernel wrapper / 分发入口op_extension/matmul_blaze_torch.cpp重命名并改写为op_extension/{operator_name}_torch.cpp或按设计删除PyTorch extension hostop_extension/register.cpp/ops.h改写注册名和函数声明或按设计删除TORCH_LIBRARY 接入op_tiling/matmul_tiling_data.h重命名并改写为{operator_name}tiling data字段必须匹配目标 Kernel Paramsop_tiling/matmul_tiling_stub.h改写为正式 tiling或删除mock tiling 不能作为正式实现scripts/gen_data_mxfp8.py/gen_data_mxfp4.py非 MX/FP4 算子删除或改写为目标gen_data.py不得保留与目标 dtype 无关的业务脚本scripts/verify_result.py改写后保留精度标准、输入输出文件、Golden 路径需与目标算子一致CV 融合目录要求普通 MatMul Vector 融合使用blaze_custom路线。自定义 Epilogue 是 Blaze Custom 组件必须放在op_kernel/include/blaze_custom/epilogue/{epilogue_name}.hKernel 中通过以下方式引用#include blaze_custom/epilogue/{epilogue_name}.h不要把融合 Epilogue 放在op_kernel/顶层。op_kernel/顶层只放目标 Kernel 入口、Wrapper 或按设计需要保留的 kernel 相关文件。目标目录结构由 DESIGN.md 决定不同 Blaze 场景的目录结构不同Developer 必须按DESIGN.md中基于组件选择、Epilogue 设计和 Tiling 设计推导出的目标目录执行纯 MatMul可只保留 blaze library 路径和目标 kernel/tiling/launcher。普通 CV需要blaze_custom/kernel、block、policy、utils、epilogue中与设计相关的模块。MX CV按受控组合路径保留 MX 相关 scale / Quant tiling 文件不得套普通 CV 路径重写 MX 反量化。GroupMatMul按 grouped scheduler / context-based epilogue 的设计保留对应模块。Step 2: 定义 KernelKernel 签名设计修饰符选择纯 matmul 用__cube__融合用__mix__(aicCount, aivCount)GM_ADDR参数顺序输入在前、输出在后、tilingData 最后模板参数conditional_t推导模式推荐或bool enum模式类型链组装从ProblemShape出发逐步推导完整类型链ProblemShape → DispatchPolicy → BlockScheduler → BlockMmad → Kernel → ParamsBlaze 库路径所有组件来自blaze/gemm/不能混用 blaze_custom每个组件暴露::Params类型从中提取字段填充参数结构体不要硬编码 Params 结构始终从组件类型推导TilingData 对齐TilingData字段必须与BlockScheduler::Params对齐当前模板使用 mock tiling正式开发时需替换参考op_kernel/matmul_blaze_kernel.hStep 3: 编写 LauncherHost boilerplateACL 会话初始化AclRtSessionRAII内存管理aclrtMallocHoststd::unique_ptr文件 I/OReadExactFile/WriteFileLayout Dispatchconditional_t 扁平 if/else 分发每种 transA/transB 组合一行Kernel LaunchgridDim, nullptr, stream语法gridDim取tilingData.usedCoreNum参考test_matmul_blaze.cpp改造点改造点文件说明算子名称test_matmul_blaze.cpp、CMakeLists.txt、run.sh三处保持一致数据类型matmul_blaze_kernel.h修改AType/BType/CType第 19-24 行Tiling 引擎op_tiling/替换 mock 实现为正式 tiling数据生成scripts/gen_data_*.py根据 dtype 调整精度验证scripts/verify_result.py根据 dtype 调整容差常见变种模式A_FULL_LOADA 矩阵常驻 L1跨 N-tile 复用修改DispatchPolicy模板参数CV 融合matmul vector 后处理GELU/SwiGLU 等使用__mix__修饰符MX 量化mxfp8/mxfp4 输入bf16 输出本模板示例Group matmul多个小 matmul 批量执行使用GroupMatmulBlockScheduler自定义扩展当标准 Blaze 组件无法满足需求时在blaze_custom/开发自定义模块§4 编译、运行与验证4.1 首次使用bash run.sh 16 128 16384 false true自动拉取 blaze tensor_api → 编译 → 生成数据 → 运行 → 验证4.2 常用命令bash run.sh --skip-build 256 256 256 # 跳过编译 bash run.sh --clean # 清理临时文件build/、blaze/、tensor_api/、input/、output/4.3 验证标准数据类型rtolatolbf16 / fp161e-31e-3fp321e-51e-5mxfp81e-21e-2mxfp45e-25e-24.4 验证流程编译无错误、无警告运行不崩溃无 segfault、无 hang精度verify_result.py与 CPU golden 对比输出[PASS]或[FAIL]§5 常见误区5.1 依赖管理⚠️不要修改op_kernel/include/blaze/和tensor_api/的内容这些是外部依赖修改后会在下次拉取时被覆盖。如需扩展 Blaze 功能在blaze_custom/目录开发自定义模块。5.2 Tiling 对齐⚠️TilingData字段必须与BlockScheduler::Params对齐TilingData是 host 端填充的 POD 结构体Params是 device 端使用的结构体。两者的字段映射关系必须正确否则会导致 kernel 行为异常。⚠️Mock tiling 仅用于示例正式开发时必须替换matmul_tiling_stub.h返回固定参数不考虑实际矩阵规模和硬件约束。正式开发时需根据BlockScheduler的要求开发完整的 tiling 引擎。5.3 路径混用⚠️Blaze 库路径和 blaze_custom 路径不能混用Blaze 库组件BlockMmad、BlockScheduler、Kernel的接口与 blaze_custom 不同。选择一条路径后所有组件都必须来自同一路径。5.4 Params 字段顺序⚠️Params聚合初始化时字段顺序必须与声明完全一致错位会导致excess elements in scalar initializer编译错误。参考matmul_blaze_kernel.h第 84 行的params初始化。§6 模板文件索引文件用途op_kernel/matmul_blaze_kernel.hKernel 模板函数参考类型链组装 Params 填充op_kernel/matmul_blaze_kernel.cppWrapper 函数参考extern C 包装 trans 分发test_matmul_blaze.cppLauncher 参考ACL 内存 I/O dispatch launchop_tiling/matmul_tiling_data.hTilingData POD 结构参考CMakeLists.txt构建配置参考run.sh编译运行验证流程参考scripts/gen_data_mxfp8.pyMXFP8 数据生成 golden 计算参考scripts/gen_data_mxfp4.pyMXFP4 数据生成 golden 计算参考scripts/verify_result.py精度验证参考【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考