Step 2: Kernel 设计总纲【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills定位告诉 Agent Kernel 入口函数怎么设计的通用方法论。具体场景的组装代码见references/scenarios/目录。§1 Kernel 签名设计修饰符选择场景修饰符说明纯 matmul无 vector 后处理__cube__AIC 独占 Cube Corematmul epilogue 融合__mix__(aicCount, aivCount)AIC AIV 混合执行GM_ADDR 参数约定参数顺序输入在前、输出在后、tilingData 最后。额外输入Scale、Bias插在输出之前。// 基础 matmul __global__ __aicore__ __cube__ void my_kernel( GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, const MyTilingData tilingData); // MX 量化 matmul额外 Scale 输入 __global__ __aicore__ __cube__ void mx_kernel( GM_ADDR dA, GM_ADDR dB, GM_ADDR dScaleA, GM_ADDR dScaleB, GM_ADDR dC, const QuantMatmulTilingData tilingData); // 融合 matmul额外 Bias 输入 __global__ __aicore__ __mix__(1, 2) void fused_kernel( GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, GM_ADDR dBias, const MyTilingData tilingData);模板参数设计模式适用场景优点conditional_t推导多种 layout/trans 组合编译期推导host 端扁平分发boolenum量化算子dtype 运行时可变灵活支持 runtime dispatch推荐模式模板参数bool TransA, bool TransB Kernel 内部conditional_t推导 Layouttemplate bool TransA, bool TransB __global__ __aicore__ __cube__ void my_kernel(GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, ...) { using LayoutA AscendC::Std::conditional_tTransA, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutB AscendC::Std::conditional_tTransB, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; // ... 后续组装 }NZ 格式需要额外的CubeFormat模板参数详见references/scenarios/basic-matmul-development.md§3。§2 类型链推导从ProblemShape出发逐步推导完整类型链ProblemShape问题规模m, n, k, batch ↓ DispatchPolicy流水策略SWAT / MatmulWithScaleMx ↓ BlockScheduler块调度器 ↓ BlockMmad数据搬运 MMAD 计算 ↓ Kernel顶层 kernel 类 ↓ Params参数结构体两条开发路径路径适用场景组件来源blaze 库路径普通 MatMul 单算子、MX 量化 MatMul、A8W8 等标准库已覆盖场景op_kernel/include/blaze/gemm/blaze_custom 路径普通 CV 融合、Group MatMul、Grouped CV、自定义 Block/Scheduler/Epilogue 扩展op_kernel/include/blaze_custom/MX CV 受控组合态MXFP8/MXFP4 MatMul Vector EpilogueKernel::MxMatmulKernelFused blaze library MX Block/Scheduler 自定义 Epilogue混用禁令默认情况下blaze_custom 模块和 blaze 库模块不能在同一 kernel 入口函数中任意混合使用。blaze_custom 使用Kernel::、Block::命名空间blaze 库使用Blaze::Gemm::Kernel::、Blaze::Gemm::Block::命名空间两者的BlockMmad模板参数数量、Params 结构和 SFINAE 机制不同普通 MatMul 单算子和纯 MX 量化 MatMul 均使用 blaze 库全套路径。基础 MatMul 场景的 CMake include 应优先采用最小可编译集合普通 CV、Grouped CV 和标准 blaze library 暂不能覆盖的扩展场景使用 blaze_custom 路径。MX CV 是唯一明确设计的受控例外Kernel::MxMatmulKernelFused专门桥接 blaze library MX Block/Scheduler 与自定义 Epilogue不视为任意混用。关键原则选择 blaze 库路径时所有组件都必须来自blaze/gemm/不能混用 blaze_customMX CV 只能通过MxMatmulKernelFused这个受控 bridge 例外接入自定义 Epilogue每个组件都暴露::Params类型从中提取字段填充参数结构体不要硬编码 Params 结构始终从组件类型推导§3 组件选择决策树场景路径KernelBlockMmadSchedulerPolicy详见基础 MatMul 单算子blaze 库GemmUniversalBlaze::Gemm::Block::BlockMmad(Basic)BlockSchedulerMatmulBasicMatmulMultiBlockBasicreferences/scenarios/basic-matmul-development.mdMX 量化单算子blaze 库GemmUniversalBlaze::Gemm::Block::BlockMmad(ScaleMx)BlockSchedulerQuantBatchMatmulV3MatmulWithScaleMxreferences/scenarios/mx-matmul-development.mdGroup MatMulblaze_customGroupMatmulKernelBlock::BlockMmadGroupMatmulBlockSchedulerSplitMMatmulMultiBlockPolicyreferences/scenarios/group-matmul-development.md普通 MatMul Vectorblaze_customMatmulKernelFusedBlock::BlockMmadMatmulSwatSchedulerMatmulMultiBlockPolicyreferences/scenarios/fusion-matmul-development.mdMXFP8/MXFP4 MatMul Vector受控组合态MxMatmulKernelFusedBlaze::Gemm::Block::BlockMmad(ScaleMx)BlockSchedulerQuantBatchMatmulV3MatmulWithScaleMxreferences/scenarios/fusion-matmul-development.mdGrouped MatMul Vectorblaze_customGroupMatmulKernel..., EpilogueBlock::BlockMmadGroupMatmulBlockSchedulerSplitMMatmulMultiBlockPolicyreferences/scenarios/fusion-matmul-development.mdLOAD 模式本 skill 默认只提供 SWAT 流式路径。Full-load、StreamK、4-buffer 不作为默认开发路径也不提供对应 tiling engine。各模块的详细能力支持的 dtype/trans/format和参数说明查阅references/modules/目录blaze_custom 模块references/modules/blaze-custom/blaze 库模块references/modules/blaze-library/blaze-modules-index.md§4 TilingData 约定TilingData 是 host 端 Tiling 引擎计算、device 端 Kernel 消费的 POD 结构体。设计规范字段必须与BlockScheduler::Params对齐#pragma pack(push, 8)alignas(8)确保 8 字节对齐#ifndef __CCE_AICORE__保护cstdintinclude现有可用 Tiling 引擎Blaze 路径下的 Tiling 选择统一由references/tiling/tiling-selection.md维护。本文只说明 TilingData 约定和字段含义不展开 tiling 算法细节。场景Tiling 引擎TilingData来源基础 MatMul 单算子MatmulTilingSwatMatmulTilingDataassets/op_tiling/matmul/MX 量化 matmulQuantMatmulTilingSwatDTypeA, DTypeBQuantMatmulTilingDataassets/op_tiling/mx/Grouped MatMul / Grouped CV复用对应非 grouped SWAT tilingMatmulTilingData或QuantMatmulTilingDataassets/op_tiling/matmul/或assets/op_tiling/mx/CV 融合场景epilogue 组装不影响 tiling 引擎选择。融合场景的 tiling 引擎与其基础 matmul 场景一致bf16/fp16 matmul epilogue → 仍按融合场景的 blaze_custom tiling 方案处理MXFP8/MXFP4 量化 matmul epilogue → 同 MX 量化 matmulQuantMatmulTilingSwatGrouped matmul epilogue → 使用{totalM,N,K}调用对应非 grouped SWAT tilinggroupNum/groupList独立传给 grouped kernel不进入 tiling data注意CV 融合的 V 部分不新增独立 tiling engine。Vector 侧只在 Cube tiling 产出的剩余 UB 中规划 extra input / tmp / output staging 和stageRows/stageSize。详见references/tiling/tiling-selection.mdreferences/scenarios/fusion-matmul-development.mdreferences/modules/blaze-custom/development/epilogue-dev-guide.md关键字段字段含义Launcher 使用位置usedCoreNum实际使用核数gridDimm / n / k问题规模kernel 端ProblemShapebaseM / baseN / baseKL0 切分颗粒kernel 端BlockSchedulerParams/QBMMTiling各 Scheduler 的 TilingData 字段含义详见references/modules/blaze-custom/scheduler-modules.md。注意本文档不介绍 tiling 算法原理。Tiling 引擎作为现成模块使用只需了解输入输出和字段含义具体选择入口见references/tiling/tiling-selection.md。§5 Wrapper 函数Wrapper 函数是extern C包装层接收运行时 trans/format 参数通过 if/else 分发到对应模板实例化的 kernel 入口extern C void my_op_launch( aclrtStream stream, GM_ADDR dA, GM_ADDR dB, GM_ADDR dC, const MyTilingData tilingData, bool transA, bool transB) { if (transA transB) { my_kerneltrue, truetilingData.usedCoreNum, nullptr, stream(dA, dB, dC, tilingData); } else if (transA !transB) { my_kerneltrue, falsetilingData.usedCoreNum, nullptr, stream(dA, dB, dC, tilingData); } else if (!transA transB) { my_kernelfalse, truetilingData.usedCoreNum, nullptr, stream(dA, dB, dC, tilingData); } else { my_kernelfalse, falsetilingData.usedCoreNum, nullptr, stream(dA, dB, dC, tilingData); } }设计要点extern C确保 C 链接供 launcher 通过函数声明调用每种 trans 组合一行编译期实例化正确版本gridDim, nullptr, stream中 gridDim 取tilingData.usedCoreNumNZ 格式需要额外的CubeFormat模板参数分发16 种组合详见references/scenarios/basic-matmul-development.md§3 和references/development/step3-launcher.md§5。下一步→references/development/step3-launcher.md编写 Launcher按需查阅references/scenarios/*.md具体场景的组装代码references/modules/blaze-custom/*.mdblaze_custom 模块能力查阅references/modules/blaze-library/blaze-modules-index.mdblaze 库模块索引【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考