基础 MatMul 开发指南【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills适用场景普通矩阵乘 C A × Bfp16/bf16/fp32可选 MMAD bias 输入。路径blaze libraryBlaze::Gemm命名空间非 blaze_custom验证状态已通过blaze_basic_matmul_probe工程穿刺覆盖 fp16/bf16/fp32、ND/NZ、transpose、bias、shape 组合共 480 cases。普通开发模板按用户需求固定 dtype仅保留 ND/NZ 与 transpose 的组合分发。§1 场景背景数学定义C[M,N] A[M,K] × B[K,N]带 bias 时C[i,j] Σ(k0..K-1) A[i,k] × B[k,j] bias[j]输入输出张量shapedtype说明AtransAfalse:(M,K)transAtrue:(K,M)fp16/bf16/fp32左输入矩阵BtransBfalse:(K,N)transBtrue:(N,K)fp16/bf16/fp32右输入矩阵Bias(N)fp32可选按列加到输出C(M,N)同 A/B输出矩阵默认 NDdtype 固定原则普通开发场景通常由用户需求指定 dtype不需要在一个算子工程中同时支持 fp16/bf16/fp32 runtime dispatch。dtype 在 kernel 类型别名和 host 侧元素字节数中固定。用户需求 dtypeAType/BType/CTypeIsFp32NZ C0fp16halffalse16bf16bfloat16_tfalse16fp32floattrue8§2 组件选择基础 MatMul 单算子使用 ops-tensor 的 blaze library Basic MatMul 组件链。组件选择来源路径blaze libraryop_kernel/include/blaze/KernelBlaze::Gemm::Kernel::GemmUniversalblaze/gemm/kernel/kernel_universal.hkernel_matmul_basic.hDispatchPolicyBlaze::Gemm::MatmulMultiBlockBasicblaze/gemm/policy/dispatch_policy.hBlockMmadBlaze::Gemm::Block::BlockMmadblaze/gemm/block/block_mmad.hblock_mmad_matmul_basic.hSchedulerBlaze::Gemm::Block::BlockSchedulerMatmulBasicblaze/gemm/block/block_scheduler_matmul_basic.hEpilogueBlaze::Gemm::Block::BlockEpilogueEmptyblaze/epilogue/block/block_epilogue_empty.h能力边界默认使用 SWAT 流式路径。Basic Kernel 只使用BlockEpilogueEmpty不承载 vector epilogue 融合。如果需求包含 ReLU/Add/Cast/Scale 等后处理路由到fusion-matmul-development.md。Full-load、StreamK、4-buffer 不属于本 skill 的默认路径。§3 Kernel 组装代码Include 顺序部分 ops-tensor Basic MatMul 头文件中使用了非限定名Shape、Coord、SetMMLayoutTransform。在 include blaze 头文件前引入对应 using避免编译期找不到符号。#include kernel_operator.h using AscendC::Coord; using AscendC::SetMMLayoutTransform; using AscendC::Shape; #include blaze/gemm/policy/dispatch_policy.h #include blaze/gemm/block/block_mmad.h #include blaze/gemm/block/block_mmad_matmul_basic.h #include blaze/gemm/block/block_scheduler_matmul_basic.h #include blaze/gemm/kernel/kernel_universal.h #include blaze/gemm/kernel/kernel_matmul_basic.h #include blaze/epilogue/block/block_epilogue_empty.hTilingData#include op_tiling/matmul/blaze_matmul_tiling_data.h enum class CubeFormat : uint32_t { ND 0, NZ 1, };MatmulTilingData由MatmulTilingSwat生成可映射到 BlockMmad 和 Scheduler 参数。Grouped MatMul 也复用该结构不增加 group 字段。Kernel 入口函数template bool TransA, bool TransB, CubeFormat FormatA, CubeFormat FormatB __global__ __aicore__ __cube__ void matmul_kernel( GM_ADDR dA, GM_ADDR dB, GM_ADDR dBias, GM_ADDR dC, const MatmulTilingData tilingData) { // [MODIFY] 根据用户需求固定 dtype。bf16 改为 bfloat16_tfp32 改为 float。 using AType half; using BType half; using CType half; using BiasType float; using LayoutA AscendC::Std::conditional_t (FormatA CubeFormat::NZ), AscendC::Std::conditional_tTransA, AscendC::Te::ZNLayoutPtn, AscendC::Te::NZLayoutPtn, AscendC::Std::conditional_tTransA, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutB AscendC::Std::conditional_t (FormatB CubeFormat::NZ), AscendC::Std::conditional_tTransB, AscendC::Te::ZNLayoutPtn, AscendC::Te::NZLayoutPtn, AscendC::Std::conditional_tTransB, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutC AscendC::Te::NDExtLayoutPtn; using LayoutBias AscendC::Te::NDExtLayoutPtn; constexpr uint64_t FUSED_OP_TYPE 0; constexpr bool IS_ND_FORMAT (FormatA CubeFormat::ND FormatB CubeFormat::ND); constexpr bool IS_FP32 AscendC::Std::is_same_vAType, float; using ProblemShape AscendC::Te::Shapeint64_t, int64_t, int64_t, int64_t; using DispatchPolicy Blaze::Gemm::MatmulMultiBlockBasicNO_FULL_LOAD_MODE, FUSED_OP_TYPE; using BlockMmad Blaze::Gemm::Block::BlockMmad DispatchPolicy, AType, LayoutA, BType, LayoutB, CType, LayoutC, BiasType, LayoutBias; using BlockScheduler Blaze::Gemm::Block::BlockSchedulerMatmulBasic ProblemShape, NO_FULL_LOAD_MODE, IS_FP32, IS_ND_FORMAT; using BlockEpilogue Blaze::Gemm::Block::BlockEpilogueEmpty; using KernelImpl Blaze::Gemm::Kernel::GemmUniversal ProblemShape, BlockMmad, BlockEpilogue, BlockScheduler; using Params typename KernelImpl::Params; using BlockMmadParams typename BlockMmad::Params; using BlockSchedulerParams typename BlockScheduler::Params; using BlockEpilogueParams typename BlockEpilogue::Params; ProblemShape problemShape{ static_castint64_t(tilingData.m), static_castint64_t(tilingData.n), static_castint64_t(tilingData.k), 1L}; BlockMmadParams mmParams; mmParams.aGmAddr dA; mmParams.bGmAddr dB; mmParams.cGmAddr dC; mmParams.biasGmAddr hasBias ? dBias : nullptr; mmParams.ml1 tilingData.mL1; mmParams.nl1 tilingData.nL1; mmParams.kl1 tilingData.kL1; mmParams.ml0 tilingData.baseM; mmParams.nl0 tilingData.baseN; mmParams.kl0 tilingData.baseK; mmParams.l1Stages tilingData.l1BufferNum; mmParams.l0cStages static_castuint16_t(tilingData.l0cDB); BlockSchedulerParams schParams; schParams.mL1 tilingData.mL1; schParams.nL1 tilingData.nL1; schParams.kL1 tilingData.kL1; schParams.baseM tilingData.baseM; schParams.baseN tilingData.baseN; schParams.baseK tilingData.baseK; schParams.mTailCnt tilingData.mTailCnt; schParams.nTailCnt tilingData.nTailCnt; schParams.mBaseTailSplitCnt tilingData.mBaseTailSplitCnt; schParams.nBaseTailSplitCnt tilingData.nBaseTailSplitCnt; schParams.mTailMain tilingData.mTailMain; schParams.nTailMain tilingData.nTailMain; schParams.isHf32 0; schParams.l2CacheMode Blaze::Gemm::L2_CACHE_DEFAULT; schParams.sliceM 0; schParams.srcNdStride 1; schParams.innerBatch 1; BlockEpilogueParams epilogueParams; Params params; params.problemShape problemShape; params.mmadParams mmParams; params.epilogueParams epilogueParams; params.schParams schParams; KernelImpl kernel; kernel(params); }TilingData → Params 映射Params 字段TilingData 来源说明problemShape{m, n, k, 1}batch 固定为 1mmadParams.aGmAddrdAA GM 地址mmadParams.bGmAddrdBB GM 地址mmadParams.cGmAddrdCC GM 地址mmadParams.biasGmAddrhasBias ? dBias : nullptr可选 fp32 biashasBias是 launcher/runtime 参数不来自TilingDatammadParams.ml1/nl1/kl1mL1/nL1/kL1L1 tile 尺寸mmadParams.ml0/nl0/kl0baseM/baseN/baseKL0 tile 尺寸mmadParams.l1Stagesl1BufferNum推荐直接映射当前 SWAT 的 L1 buffer 数mmadParams.l0cStagesl0cDB推荐直接映射当前 SWAT 的 L0C buffer 数schParams.mL1/nL1/kL1mL1/nL1/kL1Scheduler tile 尺寸schParams.baseM/baseN/baseKbaseM/baseN/baseKScheduler L0 粒度schParams.mTailCnt/nTailCntmTailCnt/nTailCnt尾块 split 参数在普通 SWAT 中语义等价于 tail split factor默认值为1/1epilogueParamsBlockEpilogueEmpty::Params{}空后处理§4 Layout 与 Launcher 分发普通开发模板固定 dtype仅分发transA × transB × formatA × formatB 16 种组合推荐使用分层分发避免手写 16 个扁平分支。template bool TransA, bool TransB, CubeFormat FormatA void LaunchByBLayout( aclrtStream stream, GM_ADDR dA, GM_ADDR dB, GM_ADDR dBias, GM_ADDR dC, const MatmulTilingData tilingData, CubeFormat formatB) { if (formatB CubeFormat::NZ) { matmul_kernelTransA, TransB, FormatA, CubeFormat::NZ tilingData.usedCoreNum, nullptr, stream(dA, dB, dBias, dC, tilingData); } else { matmul_kernelTransA, TransB, FormatA, CubeFormat::ND tilingData.usedCoreNum, nullptr, stream(dA, dB, dBias, dC, tilingData); } } template bool TransA, bool TransB void LaunchByALayout( aclrtStream stream, GM_ADDR dA, GM_ADDR dB, GM_ADDR dBias, GM_ADDR dC, const MatmulTilingData tilingData, CubeFormat formatA, CubeFormat formatB) { if (formatA CubeFormat::NZ) { LaunchByBLayoutTransA, TransB, CubeFormat::NZ(stream, dA, dB, dBias, dC, tilingData, formatB); } else { LaunchByBLayoutTransA, TransB, CubeFormat::ND(stream, dA, dB, dBias, dC, tilingData, formatB); } }完整链路LaunchByTransA → LaunchByTransB → LaunchByALayout → LaunchByBLayout§5 Tiling 参数Basic MatMul 使用assets/op_tiling/matmul/MatmulTilingSwat不要手写启发式 tiling。MatmulTilingData tilingData; MatmulTilingSwat tiling; tiling.GetTilingData(m, n, k, inputElemBytes, tilingData, transA, transB, isANz, isBNz, hasBias);关键字段字段含义Kernel 端使用usedCoreNum启动核数usedCoreNum, ...m/n/k问题规模ProblemShapemL1/nL1/kL1L1 tile 尺寸BlockMmadParamsBlockSchedulerParamsbaseM/baseN/baseKL0 tile 尺寸BlockMmadParamsBlockSchedulerParamsmTailCnt/nTailCnt尾块切分份数禁用尾块二次切分时固定为1/1BlockSchedulerParamsmBaseTailSplitCnt/nBaseTailSplitCnt尾块合并数量BlockSchedulerParams§6 Bias 输入blaze library Basic MatMul 的 bias 不需要手写 L1/L0 copy 流水直接通过BlockMmad::Params::biasGmAddr传入。项约定Bias dtypefloatBias shape(N)Bias layoutAscendC::Te::NDExtLayoutPtn无 biasbiasGmAddr nullptrusing BiasType float; using LayoutBias AscendC::Te::NDExtLayoutPtn; mmParams.biasGmAddr hasBias ? dBias : nullptr;如果需求是 bias activation / bias add / bias cast 等 vector 后处理不属于 Basic Kernel 范围应切换到 CV 融合场景。§7 常见陷阱#陷阱症状解决1未在 blaze include 前引入Shape/Coord/SetMMLayoutTransform编译报未声明标识符在 include 前添加using AscendC::...2usedCoreNum1但 M/N 多 tile输出部分列/行全 0usedCoreNum按tileM × tileN与硬件核数取 min3把验证 harness 的 dtype dispatch 带入交付模板代码复杂、实例化膨胀普通开发按用户需求固定 dtype4NZ size 按 ND 计算H2D/D2H 越界或精度异常按ceil(cols/C0) * ceil(rows/16) * 16 * C0计算5fp32 仍使用 C016NZ 数据错位fp32 C08fp16/bf16 C0166Basic Kernel 承载 vector epilogue设计不匹配Basic Kernel 只配BlockEpilogueEmpty融合需求走 fusion 文档7在本 skill 中切换 full-load/StreamK/4-buffer缺少 tiling 与 block 契约默认只使用 SWAT§8 验证建议普通 MatMul 单算子建议至少验证典型 shape小 shape、方阵、非方阵、多 tile、非 64/128 规则 shape。layoutND/ND、ND/NZ、NZ/ND、NZ/NZ。transposefalse/false、true/false、false/true、true/true。bias无 bias / 有 bias。精度阈值建议dtypertolatolfp161e-31e-3bf162e-22e-2fp321e-41e-4下一步→references/development/step3-launcher.md编写 Launcher【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN/cannbot-skills：基础矩阵乘法开发指南
基础 MatMul 开发指南【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills适用场景普通矩阵乘 C A × Bfp16/bf16/fp32可选 MMAD bias 输入。路径blaze libraryBlaze::Gemm命名空间非 blaze_custom验证状态已通过blaze_basic_matmul_probe工程穿刺覆盖 fp16/bf16/fp32、ND/NZ、transpose、bias、shape 组合共 480 cases。普通开发模板按用户需求固定 dtype仅保留 ND/NZ 与 transpose 的组合分发。§1 场景背景数学定义C[M,N] A[M,K] × B[K,N]带 bias 时C[i,j] Σ(k0..K-1) A[i,k] × B[k,j] bias[j]输入输出张量shapedtype说明AtransAfalse:(M,K)transAtrue:(K,M)fp16/bf16/fp32左输入矩阵BtransBfalse:(K,N)transBtrue:(N,K)fp16/bf16/fp32右输入矩阵Bias(N)fp32可选按列加到输出C(M,N)同 A/B输出矩阵默认 NDdtype 固定原则普通开发场景通常由用户需求指定 dtype不需要在一个算子工程中同时支持 fp16/bf16/fp32 runtime dispatch。dtype 在 kernel 类型别名和 host 侧元素字节数中固定。用户需求 dtypeAType/BType/CTypeIsFp32NZ C0fp16halffalse16bf16bfloat16_tfalse16fp32floattrue8§2 组件选择基础 MatMul 单算子使用 ops-tensor 的 blaze library Basic MatMul 组件链。组件选择来源路径blaze libraryop_kernel/include/blaze/KernelBlaze::Gemm::Kernel::GemmUniversalblaze/gemm/kernel/kernel_universal.hkernel_matmul_basic.hDispatchPolicyBlaze::Gemm::MatmulMultiBlockBasicblaze/gemm/policy/dispatch_policy.hBlockMmadBlaze::Gemm::Block::BlockMmadblaze/gemm/block/block_mmad.hblock_mmad_matmul_basic.hSchedulerBlaze::Gemm::Block::BlockSchedulerMatmulBasicblaze/gemm/block/block_scheduler_matmul_basic.hEpilogueBlaze::Gemm::Block::BlockEpilogueEmptyblaze/epilogue/block/block_epilogue_empty.h能力边界默认使用 SWAT 流式路径。Basic Kernel 只使用BlockEpilogueEmpty不承载 vector epilogue 融合。如果需求包含 ReLU/Add/Cast/Scale 等后处理路由到fusion-matmul-development.md。Full-load、StreamK、4-buffer 不属于本 skill 的默认路径。§3 Kernel 组装代码Include 顺序部分 ops-tensor Basic MatMul 头文件中使用了非限定名Shape、Coord、SetMMLayoutTransform。在 include blaze 头文件前引入对应 using避免编译期找不到符号。#include kernel_operator.h using AscendC::Coord; using AscendC::SetMMLayoutTransform; using AscendC::Shape; #include blaze/gemm/policy/dispatch_policy.h #include blaze/gemm/block/block_mmad.h #include blaze/gemm/block/block_mmad_matmul_basic.h #include blaze/gemm/block/block_scheduler_matmul_basic.h #include blaze/gemm/kernel/kernel_universal.h #include blaze/gemm/kernel/kernel_matmul_basic.h #include blaze/epilogue/block/block_epilogue_empty.hTilingData#include op_tiling/matmul/blaze_matmul_tiling_data.h enum class CubeFormat : uint32_t { ND 0, NZ 1, };MatmulTilingData由MatmulTilingSwat生成可映射到 BlockMmad 和 Scheduler 参数。Grouped MatMul 也复用该结构不增加 group 字段。Kernel 入口函数template bool TransA, bool TransB, CubeFormat FormatA, CubeFormat FormatB __global__ __aicore__ __cube__ void matmul_kernel( GM_ADDR dA, GM_ADDR dB, GM_ADDR dBias, GM_ADDR dC, const MatmulTilingData tilingData) { // [MODIFY] 根据用户需求固定 dtype。bf16 改为 bfloat16_tfp32 改为 float。 using AType half; using BType half; using CType half; using BiasType float; using LayoutA AscendC::Std::conditional_t (FormatA CubeFormat::NZ), AscendC::Std::conditional_tTransA, AscendC::Te::ZNLayoutPtn, AscendC::Te::NZLayoutPtn, AscendC::Std::conditional_tTransA, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutB AscendC::Std::conditional_t (FormatB CubeFormat::NZ), AscendC::Std::conditional_tTransB, AscendC::Te::ZNLayoutPtn, AscendC::Te::NZLayoutPtn, AscendC::Std::conditional_tTransB, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutC AscendC::Te::NDExtLayoutPtn; using LayoutBias AscendC::Te::NDExtLayoutPtn; constexpr uint64_t FUSED_OP_TYPE 0; constexpr bool IS_ND_FORMAT (FormatA CubeFormat::ND FormatB CubeFormat::ND); constexpr bool IS_FP32 AscendC::Std::is_same_vAType, float; using ProblemShape AscendC::Te::Shapeint64_t, int64_t, int64_t, int64_t; using DispatchPolicy Blaze::Gemm::MatmulMultiBlockBasicNO_FULL_LOAD_MODE, FUSED_OP_TYPE; using BlockMmad Blaze::Gemm::Block::BlockMmad DispatchPolicy, AType, LayoutA, BType, LayoutB, CType, LayoutC, BiasType, LayoutBias; using BlockScheduler Blaze::Gemm::Block::BlockSchedulerMatmulBasic ProblemShape, NO_FULL_LOAD_MODE, IS_FP32, IS_ND_FORMAT; using BlockEpilogue Blaze::Gemm::Block::BlockEpilogueEmpty; using KernelImpl Blaze::Gemm::Kernel::GemmUniversal ProblemShape, BlockMmad, BlockEpilogue, BlockScheduler; using Params typename KernelImpl::Params; using BlockMmadParams typename BlockMmad::Params; using BlockSchedulerParams typename BlockScheduler::Params; using BlockEpilogueParams typename BlockEpilogue::Params; ProblemShape problemShape{ static_castint64_t(tilingData.m), static_castint64_t(tilingData.n), static_castint64_t(tilingData.k), 1L}; BlockMmadParams mmParams; mmParams.aGmAddr dA; mmParams.bGmAddr dB; mmParams.cGmAddr dC; mmParams.biasGmAddr hasBias ? dBias : nullptr; mmParams.ml1 tilingData.mL1; mmParams.nl1 tilingData.nL1; mmParams.kl1 tilingData.kL1; mmParams.ml0 tilingData.baseM; mmParams.nl0 tilingData.baseN; mmParams.kl0 tilingData.baseK; mmParams.l1Stages tilingData.l1BufferNum; mmParams.l0cStages static_castuint16_t(tilingData.l0cDB); BlockSchedulerParams schParams; schParams.mL1 tilingData.mL1; schParams.nL1 tilingData.nL1; schParams.kL1 tilingData.kL1; schParams.baseM tilingData.baseM; schParams.baseN tilingData.baseN; schParams.baseK tilingData.baseK; schParams.mTailCnt tilingData.mTailCnt; schParams.nTailCnt tilingData.nTailCnt; schParams.mBaseTailSplitCnt tilingData.mBaseTailSplitCnt; schParams.nBaseTailSplitCnt tilingData.nBaseTailSplitCnt; schParams.mTailMain tilingData.mTailMain; schParams.nTailMain tilingData.nTailMain; schParams.isHf32 0; schParams.l2CacheMode Blaze::Gemm::L2_CACHE_DEFAULT; schParams.sliceM 0; schParams.srcNdStride 1; schParams.innerBatch 1; BlockEpilogueParams epilogueParams; Params params; params.problemShape problemShape; params.mmadParams mmParams; params.epilogueParams epilogueParams; params.schParams schParams; KernelImpl kernel; kernel(params); }TilingData → Params 映射Params 字段TilingData 来源说明problemShape{m, n, k, 1}batch 固定为 1mmadParams.aGmAddrdAA GM 地址mmadParams.bGmAddrdBB GM 地址mmadParams.cGmAddrdCC GM 地址mmadParams.biasGmAddrhasBias ? dBias : nullptr可选 fp32 biashasBias是 launcher/runtime 参数不来自TilingDatammadParams.ml1/nl1/kl1mL1/nL1/kL1L1 tile 尺寸mmadParams.ml0/nl0/kl0baseM/baseN/baseKL0 tile 尺寸mmadParams.l1Stagesl1BufferNum推荐直接映射当前 SWAT 的 L1 buffer 数mmadParams.l0cStagesl0cDB推荐直接映射当前 SWAT 的 L0C buffer 数schParams.mL1/nL1/kL1mL1/nL1/kL1Scheduler tile 尺寸schParams.baseM/baseN/baseKbaseM/baseN/baseKScheduler L0 粒度schParams.mTailCnt/nTailCntmTailCnt/nTailCnt尾块 split 参数在普通 SWAT 中语义等价于 tail split factor默认值为1/1epilogueParamsBlockEpilogueEmpty::Params{}空后处理§4 Layout 与 Launcher 分发普通开发模板固定 dtype仅分发transA × transB × formatA × formatB 16 种组合推荐使用分层分发避免手写 16 个扁平分支。template bool TransA, bool TransB, CubeFormat FormatA void LaunchByBLayout( aclrtStream stream, GM_ADDR dA, GM_ADDR dB, GM_ADDR dBias, GM_ADDR dC, const MatmulTilingData tilingData, CubeFormat formatB) { if (formatB CubeFormat::NZ) { matmul_kernelTransA, TransB, FormatA, CubeFormat::NZ tilingData.usedCoreNum, nullptr, stream(dA, dB, dBias, dC, tilingData); } else { matmul_kernelTransA, TransB, FormatA, CubeFormat::ND tilingData.usedCoreNum, nullptr, stream(dA, dB, dBias, dC, tilingData); } } template bool TransA, bool TransB void LaunchByALayout( aclrtStream stream, GM_ADDR dA, GM_ADDR dB, GM_ADDR dBias, GM_ADDR dC, const MatmulTilingData tilingData, CubeFormat formatA, CubeFormat formatB) { if (formatA CubeFormat::NZ) { LaunchByBLayoutTransA, TransB, CubeFormat::NZ(stream, dA, dB, dBias, dC, tilingData, formatB); } else { LaunchByBLayoutTransA, TransB, CubeFormat::ND(stream, dA, dB, dBias, dC, tilingData, formatB); } }完整链路LaunchByTransA → LaunchByTransB → LaunchByALayout → LaunchByBLayout§5 Tiling 参数Basic MatMul 使用assets/op_tiling/matmul/MatmulTilingSwat不要手写启发式 tiling。MatmulTilingData tilingData; MatmulTilingSwat tiling; tiling.GetTilingData(m, n, k, inputElemBytes, tilingData, transA, transB, isANz, isBNz, hasBias);关键字段字段含义Kernel 端使用usedCoreNum启动核数usedCoreNum, ...m/n/k问题规模ProblemShapemL1/nL1/kL1L1 tile 尺寸BlockMmadParamsBlockSchedulerParamsbaseM/baseN/baseKL0 tile 尺寸BlockMmadParamsBlockSchedulerParamsmTailCnt/nTailCnt尾块切分份数禁用尾块二次切分时固定为1/1BlockSchedulerParamsmBaseTailSplitCnt/nBaseTailSplitCnt尾块合并数量BlockSchedulerParams§6 Bias 输入blaze library Basic MatMul 的 bias 不需要手写 L1/L0 copy 流水直接通过BlockMmad::Params::biasGmAddr传入。项约定Bias dtypefloatBias shape(N)Bias layoutAscendC::Te::NDExtLayoutPtn无 biasbiasGmAddr nullptrusing BiasType float; using LayoutBias AscendC::Te::NDExtLayoutPtn; mmParams.biasGmAddr hasBias ? dBias : nullptr;如果需求是 bias activation / bias add / bias cast 等 vector 后处理不属于 Basic Kernel 范围应切换到 CV 融合场景。§7 常见陷阱#陷阱症状解决1未在 blaze include 前引入Shape/Coord/SetMMLayoutTransform编译报未声明标识符在 include 前添加using AscendC::...2usedCoreNum1但 M/N 多 tile输出部分列/行全 0usedCoreNum按tileM × tileN与硬件核数取 min3把验证 harness 的 dtype dispatch 带入交付模板代码复杂、实例化膨胀普通开发按用户需求固定 dtype4NZ size 按 ND 计算H2D/D2H 越界或精度异常按ceil(cols/C0) * ceil(rows/16) * 16 * C0计算5fp32 仍使用 C016NZ 数据错位fp32 C08fp16/bf16 C0166Basic Kernel 承载 vector epilogue设计不匹配Basic Kernel 只配BlockEpilogueEmpty融合需求走 fusion 文档7在本 skill 中切换 full-load/StreamK/4-buffer缺少 tiling 与 block 契约默认只使用 SWAT§8 验证建议普通 MatMul 单算子建议至少验证典型 shape小 shape、方阵、非方阵、多 tile、非 64/128 规则 shape。layoutND/ND、ND/NZ、NZ/ND、NZ/NZ。transposefalse/false、true/false、false/true、true/true。bias无 bias / 有 bias。精度阈值建议dtypertolatolfp161e-31e-3bf162e-22e-2fp321e-41e-4下一步→references/development/step3-launcher.md编写 Launcher【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
