CV 融合 MatMul 开发指南【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills适用场景MatMul Vector Epilogue 融合即C epilogue(A x B, extra...)。CV 的开发分两层先确定 C 部分的 MatMul 主体和 L0C2UB 能力再设计 V 部分的 Epilogue。不同 MatMul 场景的具体组件不同但开发方法一致。§1 CV 融合模型CV 融合是 Cube Vector 协同模型AIC 做 MatMulFixpipe 将 L0C 写入 UBAIC 通过 CrossCore flag 通知 AIVAIV 从 UB 读取 MatMul 结果加载额外输入执行 Vector Epilogue 并写回 GMAIV 完成后再通知 AIC 释放 UB。融合 kernel 使用__mix__(1, 2)1 AIC 2 AIV 协作。AIV 不提前返回不能沿用纯 MatMul__cube__if (ASCEND_IS_AIV) return的模板。§2 设计流程总览层次设计目标关键问题C 部分选择 MatMul 主体组件当前 MatMul 场景是什么对应 Block 是否具备 L0C2UB 能力V 部分设计 EpilogueVector 公式适合 MemBase 还是 RegBase剩余 UB 如何分配同步依赖是否完整Tiling复用 C 部分 tilingCV 不新增独立 vector tiling engineEpilogue 只消费 Cube tiling 后的剩余 UBLauncher组织输入和参数MatMul 输入、Epilogue 额外输入、TilingData、layout/trans 分发是否一致§3 C 部分设计3.1 选择 MatMul 场景C 部分场景KernelBlockMmadSchedulerTiling说明普通 MatMul VKernel::MatmulKernelFusedBlock::BlockMmadMatmulSwatSchedulerMatmulTilingSwat完整blaze_custom路径MXFP8/MXFP4 MatMul VKernel::MxMatmulKernelFusedBlaze::Gemm::Block::BlockMmadMXBlockSchedulerQuantBatchMatmulV3QuantMatmulTilingSwat受控组合态MX C 部分复用 blaze libraryEpilogue 由 custom bridge 承载Grouped MatMul VKernel::GroupMatmulKernel..., EpilogueBlock::BlockMmadGroupMatmulBlockSchedulerSplitM复用对应非 groupedMatmulTilingSwatEpilogue 必须使用 context-based view纯GemmUniversal是纯 MatMul 路径不承载自定义 Vector Epilogue。需要 Vector Epilogue 时必须选择能把 L0C 结果写到 UB 并与 AIV 同步的 Kernel / Block 组合。3.2 L0C2UB 能力判断CV 不是纯 MatMul 后面硬接一段 Vector。C 部分必须提供以下契约判断项要求TensorC 位置BlockMmad 支持 TensorC LocationUBL0C→UB 搬运使用CopyL0C2UBCopyL0C2UBSplitMTraitDUAL_DST_SPLIT_MTrait 选择__mix__(1,2)标准 CV 必须用CopyL0C2UBSplitMTraitCopyL0C2UBNonSplitTrait仅用于单 AIV 调试SplitMDUAL_DST_SPLIT_M将 L0C 的 M 行对半切分各 AIV 从各自 UB offset 0 读取半份数据UB 读取不需要 sub-block 偏移UB 布局Epilogue 必须知道 UB 中 MatMul 结果的 dtype、nAlign、rows 和起始位置同步Kernel 层负责 AIC/AIV CrossCore set/waitEpilogue 完成后必须回复 AIC3.3 最小组装骨架#include kernel/matmul_kernel_fused.h #include block/matmul_block_mmad.h #include block/block_scheduler_policy.h #include policy/dispatch_policy.h #include epilogue/my_epilogue.h template bool TransA, bool TransB __global__ __aicore__ __mix__(1, 2) void fused_matmul_kernel( GM_ADDR dA, GM_ADDR dB, GM_ADDR dExtra, GM_ADDR dC, const MatmulTilingData tilingData) { using AType bfloat16_t; using BType bfloat16_t; using CType bfloat16_t; using LayoutA AscendC::Std::conditional_tTransA, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutB AscendC::Std::conditional_tTransB, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutC AscendC::Te::NDExtLayoutPtn; using ProblemShape MatmulShape; using DispatchPolicy MatmulMultiBlockPolicyNO_FULL_LOAD_MODE; using BlockScheduler MatmulSwatSchedulerNO_FULL_LOAD_MODE; using BlockMmad Block::BlockMmadDispatchPolicy, AType, LayoutA, BType, LayoutB, CType, LayoutC; using EpilogueOp MyEpilogue; using KernelImpl Kernel::MatmulKernelFusedProblemShape, BlockMmad, BlockScheduler, EpilogueOp; ProblemShape problemShape{tilingData.m, tilingData.n, tilingData.k, 1L}; typename BlockMmad::Params mmParams{dA, dB, dC}; typename KernelImpl::L1Params l1Params{static_castuint64_t(tilingData.kL1)}; typename KernelImpl::BlockSchedulerParams schParams{ tilingData.baseM, tilingData.baseN, tilingData.mTailCnt, tilingData.nTailCnt, tilingData.mBaseTailSplitCnt, tilingData.nBaseTailSplitCnt, tilingData.mTailMain, tilingData.nTailMain}; typename KernelImpl::QBMMTiling qbmmParams{ tilingData.baseM, tilingData.baseN, tilingData.baseK, tilingData.l0cDB}; // mTailCnt/nTailCnt 在普通 SWAT 中的真实语义是尾块切分份数传给 // blaze_custom scheduler 时分别对应 mTailTile/nTailTile默认值为 1/1。 typename EpilogueOp::Params epilogueParams{dExtra, dC}; typename KernelImpl::Params params{problemShape, mmParams, l1Params, schParams, qbmmParams, epilogueParams}; KernelImpl kernel; kernel(params); }§4 Epilogue 设计Epilogue 负责消费 UB 中的 MatMul 结果、加载额外输入、执行 Vector 公式并写回 GM。4.1 MemBase vs RegBaseEpilogue 类型适用场景特点MemBase只有一个简单 vector 操作且有明确可用的 AscendC LocalTensor API如Mul/Add/Div/Cast中间值通常写 UB代码简单RegBaseGELU / SwiGLU / LayerNorm / 多输入 scale / 多中间值公式链中间值在 RegTensor 内流转减少 UB 占用和读写默认选择 RegBase。只有当公式足够简单、API 明确且 UB 空间满足时才选择 MemBase。详细设计通用入口references/modules/blaze-custom/development/epilogue-dev-guide.mdMemBasereferences/modules/blaze-custom/development/epilogue-membase-design.mdRegBasereferences/modules/blaze-custom/development/epilogue-regbase-design.md4.2 SplitM 与 row-dependent 输入CopyL0C2UBSplitMTraitDUAL_DST_SPLIT_M将 L0C 的 M 行对半切分到两个 AIV SubBlock。GM offset需要 sub-block 偏移origM blockShapeM halfM ceilDiv(origM, GetTaskRation()) localRows (origM is odd) ? (halfM - GetSubBlockIdx()) : halfM tileM0 gmOffset / N tileN0 gmOffset % N subM0 tileM0 GetSubBlockIdx() * halfM stageM0 subM0 stageRowOffset rowDependentInputOffset stageM0 outputOffset subM0 * N tileN0UB 读取不需要 sub-block 偏移DUAL_DST_SPLIT_M硬件自动分片每个 AIV 从各自 UB offset 0 读取半份数据srcAddr cLocal_.GetPhyAddr() // 不加 GetSubBlockIdx() * halfM * nAlign rowSrc srcAddr row * nAlign关键GM 侧 offset 需要加 sub-block 偏移UB 侧不需要。从 GM 公式推断 UB 也需要偏移是常见错误。localRows0 边界当curM1奇数halfM1时 V1 的localRows0early return 即可CV 同步由 kernel 层处理。详见references/modules/blaze-custom/development/epilogue-dev-guide.md§4。4.3 Epilogue 设计输出每个 Epilogue 应明确项内容Params所有额外输入和输出的 GM_ADDRUB 分区cLocal_、extra input、tmp、output staging 的字节范围公式伪代码根据用户公式列出 vector API 调用链同步伪代码CrossCore 之外的 MTE2/V/MTE3 正向与反向依赖Offset 公式SplitM、stage、tail、row-dependent input 和 output offsetAPI 来源MemBase 查ascendc-api-best-practicesRegBase 查ascendc-regbase-best-practice§5 Tiling 与 UB 剩余空间关系CV 融合不新增独立 vector tiling engine。Tiling 选择见references/tiling/tiling-selection.md。设计关系Cube tiling - 决定 baseM/baseN/baseK/kL1 和 L0C2UB 的 matmul result 区域 - 计算 matmulAreaBytes - 剩余 UB 交给 Epilogue 设计 extra/tmp/output buffer - 若剩余 UB 不足回调 C 部分调整 baseM/baseN 或切换 RegBase要点先保 L0C2UB 的 MatMul 结果完整落 UB。V 部分不独立切 tile只在剩余 UB 中设计stageRows/stageSize。复杂公式链优先用 RegBase 降低 UB 中间值数量。§6 Launcher 与参数组织Launcher 需要同时组织 C 部分输入、Epilogue 额外输入和 TilingData。内容要求GM_ADDR 顺序输入在前、输出在后、tilingData 最后C 部分输入A/B 和可能存在的 MX scale / groupListV 部分输入Bias / scale / residual / activation 参数等gridDim使用tilingData.usedCoreNumtrans / format 分发与对应 MatMul 场景一致Grouped CV 使用totalM调用对应非 grouped SWAT tilinggroupList/groupNum独立传入 grouped kernel不进入 tiling data。§7 常见陷阱#现象根因修复F1AIC hangAIV epilogue 崩溃或未回复 CrossCore flag检查 epilogue 末尾写回与 kernel 层CrossCoreSetFlagF2AIV 输出全零未等 AIC 信号就读 UB确认 AIV 侧CrossCoreWaitFlag后才读 UBF3大 shape FAIL多 tile flag 轮转错误保持count / COUNT_ID_MAX % COUNT_FLAG轮转协议F4UB 越界epilogue 额外 buffer 未从 UB 预算扣除计算remainBytes UB_SIZE - matmulAreaBytes - epilogueBytesF5tail tile 错乱nAlign 用 Init 时 baseN 固定计算per-call 从blockShapeN计算 nAlignF6误用__cube__复制纯 MatMul 模板CV 必须使用__mix__(1, 2)F7SplitM 精度错位GM 侧Epilogue 未按 SubBlock 计算 rows / gmOffset / extra input offset使用GetTaskRation()/GetSubBlockIdx()F8纯GemmUniversal做融合无法接入自定义 Vector Epilogue走本文件的 CV 组装路径F9SplitM 精度错位UB 侧UB 读取 cLocal_ 时加了 sub-block 偏移SplitM 已硬件分片从 UB offset 0 读取F10UB buffer 重叠matmulAreaBytes行步长用 16 而非 nAlignL0C行步长 nAlignL0C不是 L0C cube 边长 16F11tail 场景输出错乱UB→GM DataCopyPadsrcStride传 bytes 而非 32B 单位UB 侧 stride 传 0nAlign 保证 32B 对齐F12多 stage 数据竞争用MTE3_MTE2自 Set 自 Wait无实际同步效果用V_MTE2MTE3_V分别保护不同 bufferF13首轮 hang 或尾轮 flag 泄漏缺少 Init 预发射或析构排空Init 预发射所有反向 SetFlag析构排空所有 WaitFlag【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANNBot技能：C+V融合MatMul开发指南
CV 融合 MatMul 开发指南【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills适用场景MatMul Vector Epilogue 融合即C epilogue(A x B, extra...)。CV 的开发分两层先确定 C 部分的 MatMul 主体和 L0C2UB 能力再设计 V 部分的 Epilogue。不同 MatMul 场景的具体组件不同但开发方法一致。§1 CV 融合模型CV 融合是 Cube Vector 协同模型AIC 做 MatMulFixpipe 将 L0C 写入 UBAIC 通过 CrossCore flag 通知 AIVAIV 从 UB 读取 MatMul 结果加载额外输入执行 Vector Epilogue 并写回 GMAIV 完成后再通知 AIC 释放 UB。融合 kernel 使用__mix__(1, 2)1 AIC 2 AIV 协作。AIV 不提前返回不能沿用纯 MatMul__cube__if (ASCEND_IS_AIV) return的模板。§2 设计流程总览层次设计目标关键问题C 部分选择 MatMul 主体组件当前 MatMul 场景是什么对应 Block 是否具备 L0C2UB 能力V 部分设计 EpilogueVector 公式适合 MemBase 还是 RegBase剩余 UB 如何分配同步依赖是否完整Tiling复用 C 部分 tilingCV 不新增独立 vector tiling engineEpilogue 只消费 Cube tiling 后的剩余 UBLauncher组织输入和参数MatMul 输入、Epilogue 额外输入、TilingData、layout/trans 分发是否一致§3 C 部分设计3.1 选择 MatMul 场景C 部分场景KernelBlockMmadSchedulerTiling说明普通 MatMul VKernel::MatmulKernelFusedBlock::BlockMmadMatmulSwatSchedulerMatmulTilingSwat完整blaze_custom路径MXFP8/MXFP4 MatMul VKernel::MxMatmulKernelFusedBlaze::Gemm::Block::BlockMmadMXBlockSchedulerQuantBatchMatmulV3QuantMatmulTilingSwat受控组合态MX C 部分复用 blaze libraryEpilogue 由 custom bridge 承载Grouped MatMul VKernel::GroupMatmulKernel..., EpilogueBlock::BlockMmadGroupMatmulBlockSchedulerSplitM复用对应非 groupedMatmulTilingSwatEpilogue 必须使用 context-based view纯GemmUniversal是纯 MatMul 路径不承载自定义 Vector Epilogue。需要 Vector Epilogue 时必须选择能把 L0C 结果写到 UB 并与 AIV 同步的 Kernel / Block 组合。3.2 L0C2UB 能力判断CV 不是纯 MatMul 后面硬接一段 Vector。C 部分必须提供以下契约判断项要求TensorC 位置BlockMmad 支持 TensorC LocationUBL0C→UB 搬运使用CopyL0C2UBCopyL0C2UBSplitMTraitDUAL_DST_SPLIT_MTrait 选择__mix__(1,2)标准 CV 必须用CopyL0C2UBSplitMTraitCopyL0C2UBNonSplitTrait仅用于单 AIV 调试SplitMDUAL_DST_SPLIT_M将 L0C 的 M 行对半切分各 AIV 从各自 UB offset 0 读取半份数据UB 读取不需要 sub-block 偏移UB 布局Epilogue 必须知道 UB 中 MatMul 结果的 dtype、nAlign、rows 和起始位置同步Kernel 层负责 AIC/AIV CrossCore set/waitEpilogue 完成后必须回复 AIC3.3 最小组装骨架#include kernel/matmul_kernel_fused.h #include block/matmul_block_mmad.h #include block/block_scheduler_policy.h #include policy/dispatch_policy.h #include epilogue/my_epilogue.h template bool TransA, bool TransB __global__ __aicore__ __mix__(1, 2) void fused_matmul_kernel( GM_ADDR dA, GM_ADDR dB, GM_ADDR dExtra, GM_ADDR dC, const MatmulTilingData tilingData) { using AType bfloat16_t; using BType bfloat16_t; using CType bfloat16_t; using LayoutA AscendC::Std::conditional_tTransA, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutB AscendC::Std::conditional_tTransB, AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn; using LayoutC AscendC::Te::NDExtLayoutPtn; using ProblemShape MatmulShape; using DispatchPolicy MatmulMultiBlockPolicyNO_FULL_LOAD_MODE; using BlockScheduler MatmulSwatSchedulerNO_FULL_LOAD_MODE; using BlockMmad Block::BlockMmadDispatchPolicy, AType, LayoutA, BType, LayoutB, CType, LayoutC; using EpilogueOp MyEpilogue; using KernelImpl Kernel::MatmulKernelFusedProblemShape, BlockMmad, BlockScheduler, EpilogueOp; ProblemShape problemShape{tilingData.m, tilingData.n, tilingData.k, 1L}; typename BlockMmad::Params mmParams{dA, dB, dC}; typename KernelImpl::L1Params l1Params{static_castuint64_t(tilingData.kL1)}; typename KernelImpl::BlockSchedulerParams schParams{ tilingData.baseM, tilingData.baseN, tilingData.mTailCnt, tilingData.nTailCnt, tilingData.mBaseTailSplitCnt, tilingData.nBaseTailSplitCnt, tilingData.mTailMain, tilingData.nTailMain}; typename KernelImpl::QBMMTiling qbmmParams{ tilingData.baseM, tilingData.baseN, tilingData.baseK, tilingData.l0cDB}; // mTailCnt/nTailCnt 在普通 SWAT 中的真实语义是尾块切分份数传给 // blaze_custom scheduler 时分别对应 mTailTile/nTailTile默认值为 1/1。 typename EpilogueOp::Params epilogueParams{dExtra, dC}; typename KernelImpl::Params params{problemShape, mmParams, l1Params, schParams, qbmmParams, epilogueParams}; KernelImpl kernel; kernel(params); }§4 Epilogue 设计Epilogue 负责消费 UB 中的 MatMul 结果、加载额外输入、执行 Vector 公式并写回 GM。4.1 MemBase vs RegBaseEpilogue 类型适用场景特点MemBase只有一个简单 vector 操作且有明确可用的 AscendC LocalTensor API如Mul/Add/Div/Cast中间值通常写 UB代码简单RegBaseGELU / SwiGLU / LayerNorm / 多输入 scale / 多中间值公式链中间值在 RegTensor 内流转减少 UB 占用和读写默认选择 RegBase。只有当公式足够简单、API 明确且 UB 空间满足时才选择 MemBase。详细设计通用入口references/modules/blaze-custom/development/epilogue-dev-guide.mdMemBasereferences/modules/blaze-custom/development/epilogue-membase-design.mdRegBasereferences/modules/blaze-custom/development/epilogue-regbase-design.md4.2 SplitM 与 row-dependent 输入CopyL0C2UBSplitMTraitDUAL_DST_SPLIT_M将 L0C 的 M 行对半切分到两个 AIV SubBlock。GM offset需要 sub-block 偏移origM blockShapeM halfM ceilDiv(origM, GetTaskRation()) localRows (origM is odd) ? (halfM - GetSubBlockIdx()) : halfM tileM0 gmOffset / N tileN0 gmOffset % N subM0 tileM0 GetSubBlockIdx() * halfM stageM0 subM0 stageRowOffset rowDependentInputOffset stageM0 outputOffset subM0 * N tileN0UB 读取不需要 sub-block 偏移DUAL_DST_SPLIT_M硬件自动分片每个 AIV 从各自 UB offset 0 读取半份数据srcAddr cLocal_.GetPhyAddr() // 不加 GetSubBlockIdx() * halfM * nAlign rowSrc srcAddr row * nAlign关键GM 侧 offset 需要加 sub-block 偏移UB 侧不需要。从 GM 公式推断 UB 也需要偏移是常见错误。localRows0 边界当curM1奇数halfM1时 V1 的localRows0early return 即可CV 同步由 kernel 层处理。详见references/modules/blaze-custom/development/epilogue-dev-guide.md§4。4.3 Epilogue 设计输出每个 Epilogue 应明确项内容Params所有额外输入和输出的 GM_ADDRUB 分区cLocal_、extra input、tmp、output staging 的字节范围公式伪代码根据用户公式列出 vector API 调用链同步伪代码CrossCore 之外的 MTE2/V/MTE3 正向与反向依赖Offset 公式SplitM、stage、tail、row-dependent input 和 output offsetAPI 来源MemBase 查ascendc-api-best-practicesRegBase 查ascendc-regbase-best-practice§5 Tiling 与 UB 剩余空间关系CV 融合不新增独立 vector tiling engine。Tiling 选择见references/tiling/tiling-selection.md。设计关系Cube tiling - 决定 baseM/baseN/baseK/kL1 和 L0C2UB 的 matmul result 区域 - 计算 matmulAreaBytes - 剩余 UB 交给 Epilogue 设计 extra/tmp/output buffer - 若剩余 UB 不足回调 C 部分调整 baseM/baseN 或切换 RegBase要点先保 L0C2UB 的 MatMul 结果完整落 UB。V 部分不独立切 tile只在剩余 UB 中设计stageRows/stageSize。复杂公式链优先用 RegBase 降低 UB 中间值数量。§6 Launcher 与参数组织Launcher 需要同时组织 C 部分输入、Epilogue 额外输入和 TilingData。内容要求GM_ADDR 顺序输入在前、输出在后、tilingData 最后C 部分输入A/B 和可能存在的 MX scale / groupListV 部分输入Bias / scale / residual / activation 参数等gridDim使用tilingData.usedCoreNumtrans / format 分发与对应 MatMul 场景一致Grouped CV 使用totalM调用对应非 grouped SWAT tilinggroupList/groupNum独立传入 grouped kernel不进入 tiling data。§7 常见陷阱#现象根因修复F1AIC hangAIV epilogue 崩溃或未回复 CrossCore flag检查 epilogue 末尾写回与 kernel 层CrossCoreSetFlagF2AIV 输出全零未等 AIC 信号就读 UB确认 AIV 侧CrossCoreWaitFlag后才读 UBF3大 shape FAIL多 tile flag 轮转错误保持count / COUNT_ID_MAX % COUNT_FLAG轮转协议F4UB 越界epilogue 额外 buffer 未从 UB 预算扣除计算remainBytes UB_SIZE - matmulAreaBytes - epilogueBytesF5tail tile 错乱nAlign 用 Init 时 baseN 固定计算per-call 从blockShapeN计算 nAlignF6误用__cube__复制纯 MatMul 模板CV 必须使用__mix__(1, 2)F7SplitM 精度错位GM 侧Epilogue 未按 SubBlock 计算 rows / gmOffset / extra input offset使用GetTaskRation()/GetSubBlockIdx()F8纯GemmUniversal做融合无法接入自定义 Vector Epilogue走本文件的 CV 组装路径F9SplitM 精度错位UB 侧UB 读取 cLocal_ 时加了 sub-block 偏移SplitM 已硬件分片从 UB offset 0 读取F10UB buffer 重叠matmulAreaBytes行步长用 16 而非 nAlignL0C行步长 nAlignL0C不是 L0C cube 边长 16F11tail 场景输出错乱UB→GM DataCopyPadsrcStride传 bytes 而非 32B 单位UB 侧 stride 传 0nAlign 保证 32B 对齐F12多 stage 数据竞争用MTE3_MTE2自 Set 自 Wait无实际同步效果用V_MTE2MTE3_V分别保护不同 bufferF13首轮 hang 或尾轮 flag 泄漏缺少 Init 预发射或析构排空Init 预发射所有反向 SetFlag析构排空所有 WaitFlag【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
