Step 3: 编写 Launcher【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills定位Launcher 编写完整参考。CMake 和 run.sh 配置见references/development/step1-setup.md§4-§5。§1 Launcher 职责Launcher 是 host 端 C 程序承担以下职责Tiling 计算调用 Tiling 引擎生成TilingData结构体内存管理分配 host/device 缓冲区管理生命周期Kernel 启动通过gridDim, nullptr, stream语法启动 device 端 kernel结果输出D2H 回拷后将计算结果写入文件三种 Kernel 模式模式修饰符典型场景纯 AICCube__cube__基础 matmul、MX 量化 matmul混合 AIC/AIV__mix__(aicCount, aivCount)matmul vector 后融合Blaze 库直调__cube__复用 Blaze GemmUniversal§2 ACL 会话管理constexpr int32_t deviceId 0; aclrtContext context nullptr; aclrtStream stream nullptr; ACL_CHECK(aclInit(nullptr)); ACL_CHECK(aclrtSetDevice(deviceId)); ACL_CHECK(aclrtCreateContext(context, deviceId)); ACL_CHECK(aclrtCreateStream(stream));退出时显式调用aclrtDestroyStream、aclrtDestroyContext、aclrtResetDevice和aclFinalizedeviceId通常固定为0多卡场景通过npu-smi info确认可用设备运行 launcher 前必须 source CANN 环境变量§3 内存管理Host 内存uint16_t* hA nullptr; CHECK_COND(aclrtMallocHost((void**)hA, sizeA) ACL_SUCCESS, alloc host A failed.); std::unique_ptrvoid, aclError (*)(void*) hostA(hA, aclrtFreeHost);Device 内存GM_ADDR dA nullptr; CHECK_COND(aclrtMalloc((void**)dA, sizeA, ACL_MEM_MALLOC_HUGE_ONLY) ACL_SUCCESS, alloc device A failed.); std::unique_ptrvoid, aclError (*)(void*) deviceA(dA, aclrtFree);缓冲区大小计算ND 格式uint64_t sizeA m * k * ELEM_BYTES;NZ 格式物理排列与 ND 不同需按物理维度计算static uint64_t CalcNzSize(uint64_t dim0, uint64_t dim1, uint64_t c0, uint64_t elemBytes) { uint64_t dim0Blocks (dim0 15) / 16; uint64_t dim1Blocks (dim1 c0 - 1) / c0; return dim1Blocks * dim0Blocks * 16 * c0 * elemBytes; } constexpr uint64_t ELEM_BYTES sizeof(MyDType); // bf16/fp16: 2, fp32: 4, fp8/int8: 1 constexpr uint64_t C0 32 / ELEM_BYTES; // bf16/fp16: 16, fp8/int8: 32 uint64_t sizeA aIsNz ? (transA ? CalcNzSize(k, m, C0, ELEM_BYTES) : CalcNzSize(m, k, C0, ELEM_BYTES)) : m * k * ELEM_BYTES;关键原则所有缓冲区必须用std::unique_ptr包装Host 内存用aclrtFreeHost释放Device 内存用aclrtFree释放Device 内存分配使用ACL_MEM_MALLOC_HUGE_ONLYNZ 格式的 size 计算必须传入elemBytes参数不要硬编码§4 文件 I/O输入读取ExampleIoPaths paths GetExampleIoPaths(); CHECK_COND(ReadExactFile(paths.inputDir /input_a.bin, hA, sizeA), read A failed.); CHECK_COND(ReadExactFile(paths.inputDir /input_b.bin, hB, sizeB), read B failed.);输出写入CHECK_COND(WriteFile(paths.outputDir /output.bin, hC, sizeC), write output failed.);目录约定executable_dir/ ├── input/ # 输入 .bin 文件 │ ├── input_a.bin │ ├── input_b.bin │ ├── input_d.bin # epilogue 额外输入 │ ├── input_scaleA.bin # MX 量化 Scale │ └── input_scaleB.bin └── output/ # 输出 .bin 文件 └── output.bin或 npu_out.bin§5 Layout Dispatchconditional_t 扁平分发推荐模式Kernel 签名含 CubeFormat 模板参数template bool TransA, bool TransB, CubeFormat FormatA, CubeFormat FormatB __global__ __aicore__ __cube__ void my_kernel(...) { 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; // LayoutB 同理 }Host 端分发16 种组合CubeFormat fA (aLayout nz) ? CubeFormat::NZ : CubeFormat::ND; CubeFormat fB (bLayout nz) ? CubeFormat::NZ : CubeFormat::ND; if (!transA !transB fA CubeFormat::ND fB CubeFormat::ND) my_kernelfalse, false, CubeFormat::ND, CubeFormat::ND...(...); else if (!transA !transB fA CubeFormat::ND fB CubeFormat::NZ) my_kernelfalse, false, CubeFormat::ND, CubeFormat::NZ...(...); // ... 共 16 行transA/transB 各 2 值 × fA/fB 各 2 值 else my_kerneltrue, true, CubeFormat::NZ, CubeFormat::NZ...(...);简化分发固定 layout / 固定 trans仅支持 ND 输入时只需分发 transA/transB4 种组合if (!transA !transB) my_kernelfalse, false...(...); else if (!transA transB) my_kernelfalse, true...(...); else if ( transA !transB) my_kerneltrue, false...(...); else my_kerneltrue, true...(...);普通 MatMul 分发范围普通开发模板按用户需求固定 dtype不默认生成 fp16/bf16/fp32 runtime dispatchLauncher 只分发transA × transB × formatA × formatB 16 种组合推荐使用分层函数组织LaunchByTransA → LaunchByTransB → LaunchByALayout → LaunchByBLayout具体示例见references/scenarios/basic-matmul-development.md§4。MX 量化分发MX 量化交付工程通常固定为 MXFP8 或 MXFP4除非用户明确要求通用 demo不默认生成 MXFP8/MXFP4 runtime dispatch。若需要同时支持两类 dtype可在 host 端显式分发到不同模板实例。if (dtype mxfp4) { my_quant_kerneltransA, transB, DT_FLOAT4_E2M1, DT_FLOAT4_E2M1...(...); } else { my_quant_kerneltransA, transB, DT_FLOAT8_E4M3FN, DT_FLOAT8_E4M3FN...(...); }Layout 格式的详细定义和 ND/NZ/ZN 转换方法详见references/fundamentals/blaze-matmul-layout.md。§6 Kernel LaunchkernelNametemplateArgs gridDim, nullptr, stream(kernelArgs...);gridDim启动核数取tilingData.usedCoreNum第二个参数保留固定传nullptrstream从AclRtSession获取的 ACL stream模板参数在 launch 时显式指定不要依赖模板参数推导§7 开发验证流程Launcher 本身只负责到 D2H 回拷 写输出文件。验证步骤编译无错误、无警告运行不崩溃无 segfault、无 hang精度外部脚本与 CPU golden 对比非 launcher 职责精度标准数据类型rtolatolfp161e-31e-3bf162e-22e-2fp321e-41e-4mxfp81e-21e-2mxfp45e-25e-2fp32 建议从1e-4/1e-4作为工程验收阈值开始如需更严精度再单独分析 HF32、累加路径和硬件误差。常见调试手段检查 TilingData 字段是否合理baseM/baseN/baseK 是否为 16 的倍数检查 buffer 大小是否正确NZ 格式需按物理维度计算检查 kernel launch 参数gridDim 是否等于 usedCoreNum§8 Include 文件组织Launcher 的 include 按 4 层分组层来源路径前缀职责标准库C STL...基础类型、容器、异常SDKCANN Toolkitacl/...,kernel_operator.hACL 运行时 AscendC 设备 API项目 Host 侧工程assets/op_tiling/op_tiling/...tiling 引擎项目 Kernel 侧工程op_kernel/op_kernel/...kernel 模板、blaze 组件按需 include 决策表场景Host 侧差异Kernel 侧差异普通 MatMul 单算子Basic MatMul TilingData / Tiling 引擎op_kernel/op_kernel.h内部 includeblaze/gemm/...基础场景使用最小 SDK include 集MX 量化单算子assets/op_tiling/quant_matmul_mx_tiling.hop_kernel/op_kernel.h内部 includeblaze/gemm/...默认先用最小 SDK include 集缺失再最小回补普通 CV 融合fusion tiling epilogue 额外输入blaze_customfusion kernel assets/blaze_custom/epilogue/*.hMX CV 融合QuantMatmulTilingData ScaleA/ScaleB epilogue 额外输入MxMatmulKernelFused blaze library MX Block/Scheduler 自定义 EpilogueGrouped MatMulgrouped tilingblaze_customgroup kernelGrouped CV 融合grouped tiling groupList epilogue 额外输入GroupMatmulKernel..., Epilogue context-based Epilogue§9 完整 Launcher 模板以下是普通 MatMul 单算子最小化 Launcher 骨架。Kernel 侧完整 blaze library 组装代码见references/scenarios/basic-matmul-development.md§3Launcher 只负责 tiling、内存、分发和启动。#include cstdint #include exception #include iostream #include memory #include string #include acl/acl.h #include kernel_operator.h // ---- 项目Host 侧 ---- #include op_tiling/matmul/blaze_matmul_tiling.h #include op_tiling/matmul/blaze_matmul_tiling_data.h // ---- 项目Kernel 侧 ---- #include op_kernel/my_matmul_kernel.h // 内部使用 blaze/gemm/... 组装 extern C void my_matmul_launch( aclrtStream stream, GM_ADDR dA, GM_ADDR dB, GM_ADDR dBias, GM_ADDR dC, const MatmulTilingData tilingData, bool transA, bool transB, CubeFormat formatA, CubeFormat formatB); int main(int argc, char* argv[]) { try { // 参数解析普通开发模板固定 dtype仅解析 shape/layout/trans/bias uint64_t m ParsePositiveUint64(argv[1], m); uint64_t k ParsePositiveUint64(argv[2], k); uint64_t n ParsePositiveUint64(argv[3], n); bool transA ParseBool(argv[4]); bool transB ParseBool(argv[5]); CubeFormat formatA ParseFormat(argv[6]); // nd / nz CubeFormat formatB ParseFormat(argv[7]); bool hasBias ParseBool(argv[8]); // [MODIFY] dtype 由用户需求固定。fp16/bf162fp324NZ C0 分别为 16/8。 constexpr uint64_t ELEM_BYTES sizeof(uint16_t); constexpr uint64_t C0 32 / ELEM_BYTES; uint64_t sizeA CalcInputSize(m, k, transA, formatA, C0, ELEM_BYTES); uint64_t sizeB CalcInputSize(k, n, transB, formatB, C0, ELEM_BYTES); uint64_t sizeBias hasBias ? n * sizeof(float) : 0; uint64_t sizeC m * n * ELEM_BYTES; MatmulTilingData tilingData; MatmulTilingSwat tilingEngine; tilingEngine.GetTilingData(m, n, k, ELEM_BYTES, tilingData, transA, transB, formatA CubeFormat::NZ, formatB CubeFormat::NZ, hasBias); aclrtContext context nullptr; aclrtStream stream nullptr; ACL_CHECK(aclInit(nullptr)); ACL_CHECK(aclrtSetDevice(0)); ACL_CHECK(aclrtCreateContext(context, 0)); ACL_CHECK(aclrtCreateStream(stream)); // 内存分配 文件读取 H2D // ... 见 §3 内存管理、§4 文件 I/O my_matmul_launch(stream, dA, dB, dBias, dC, tilingData, transA, transB, formatA, formatB); // D2H 同步 写输出 // ... 见 §7 验证流程 ACL_CHECK(aclrtSynchronizeStream(stream)); ACL_CHECK(aclrtDestroyStream(stream)); ACL_CHECK(aclrtDestroyContext(context)); ACL_CHECK(aclrtResetDevice(0)); ACL_CHECK(aclFinalize()); return 0; } catch (const std::exception e) { std::cerr e.what() std::endl; return 1; } }MX 量化单算子 Launcher 仍按references/scenarios/mx-matmul-development.md使用QuantMatmulTilingData和 Scale 输入。Grouped 场景使用totalM调用对应非 grouped tiling并额外准备groupList/groupNum作为 grouped kernel 参数它们不写入 tiling data。MX CV 融合转到references/scenarios/fusion-matmul-development.mdhost 侧仍使用QuantMatmulTilingSwat但需要同时准备 A/B/ScaleA/ScaleB、epilogue 额外输入和输出kernel wrapper 使用__mix__(1, 2)路径并按 transA/transB 分发。host_utils目录不是本 skill 的一部分launcher 所需 helper 由工程内自行实现或直接内联。§10 构建与运行CMake 配置详见references/development/step1-setup.md§4。按需查阅references/fundamentals/blaze-matmul-layout.mdLayout 格式详情ND/NZ/ZN 定义、数据生成、LayoutPtn 选择【免费下载链接】cannbot-skillsCANNBot 是面向 CANN 开发的用于提升开发效率的系列智能体本仓库为其提供可复用的 Skills 模块。项目地址: https://gitcode.com/cann/cannbot-skills创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考