- 人工智能
- 深度学习
- 算子库
- CANN
- Ascend
【免费下载链接】asc-devkit
本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。
本文基于 CANN asc-devkit 仓库中的 matmul_k_reorder_load 直调样例，讲解如何在高阶 Matmul API（MDL 模板）中开启 K 轴错峰加载数据特性，缓解多核对 Global Memory 同地址访问冲突，并结合仓库源码解析该特性的配置入口、底层迭代器实现与调优案例，帮助读者在大 K 轴场景下写出性能更优的矩阵乘算子。
一、为什么需要 K 轴错峰加载：多核同地址访问冲突问题
在昇腾 AI 处理器的多核并行执行架构中，Matmul 计算会按 M、N、K 方向对矩阵进行切分，将不同分块分配给不同的 AI Core。此时存在一个典型的性能隐患：
- 如果多个核共享同一份左矩阵（A 矩阵）或右矩阵（B 矩阵），且该矩阵存储在Global Memory（GM）中；
- 那么多个核往往会同时访问同一个 GM 地址来加载矩阵数据，产生同地址访问冲突（Address Access Conflict）；
- 冲突会导致数据搬运（MTE2）效率下降，进而拉低整体算子性能。
开启K 轴错峰加载数据（K-axis staggered data loading，源码中对应enableKdimReorderLoad）后，多核在执行 Matmul 时会尽量在相同时间访问不同的 Global Memory 地址：即沿 K 轴方向切分出的不同核，各自从不同的 GM 起始地址开始错峰搬运数据，从而降低地址访问冲突概率、提升性能。
该特性的适用前提非常明确，来自样例 README与仓库调优文档 matmul_multi_core_k_stagger.md 的约束条件：
- K 轴较大：建议 K 不小于 4096；
- 矩阵非全载：左矩阵或右矩阵的 K 方向数据不能一次性全部搬入并保持在 L1 Buffer 中（即 K 轴非全载）；
- 仅支持 MDL 模板（
doMultiDataLoad，即多数据加载模式的 Matmul 配置）； - 多核执行Matmul 计算；
- A 矩阵或 B 矩阵的内存位置位于 GM。
注意：开启该特性的场景是“左矩阵与右矩阵均非全载”，且核心目标是减少多核搬运阶段（MTE2）的 GM 地址访问冲突，而非改变矩阵乘的数学结果——它属于纯性能优化开关。
二、支持的产品与 CANN 版本
根据样例文档的支持矩阵：
| 产品 | CANN 版本 |
|---|---|
| Ascend 950PR & 950DT 系列产品 | >= CANN 9.1.0 |
| Atlas A3 系列产品 | >= CANN 9.0.0 |
| Atlas A2 系列产品 | >= CANN 9.0.0 |
编译时通过CMAKE_ASC_ARCHITECTURES指定 NPU 架构：dav-2201对应 Atlas A2/A3 系列产品，dav-3510对应 Ascend 950PR & 950DT 系列产品（详见下文编译选项）。
三、样例整体结构与运行效果
matmul_k_reorder_load样例位于examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_k_reorder_load/，采用**直调（Direct Invocation）**方式演示该特性：
matmul_k_reorder_load/ ├── CMakeLists.txt # 构建配置（ASC 语言工程） ├── matmul_k_reorder_load.asc # 算子内核源码（含 tiling 生成与 host 侧启动） ├── data_utils.h # 数据读写工具 ├── scripts/ │ ├── gen_data.py # 生成测试输入数据与 golden 结果 │ └── verify_result.py # 精度校验脚本 ├── README.md / README_en.md # 中英文样例说明样例规格（来自 matmul_k_reorder_load.asc 与调优文档）：
- 左矩阵 A：
M × K = 768 × 6144，float16，ND 格式； - 右矩阵 B：
K × N = 6144 × 2048，float16，ND 格式； - 偏置 bias：
1 × N = 1 × 2048，float32； - 输出 C：
M × N = 768 × 2048，float32。
K = 6144 满足“K 轴不小于 4096”的推荐条件，且 K 方向数据无法全载入 L1 Buffer，是典型的 K 轴错峰加载适用场景。
执行成功后，精度校验脚本输出test pass!，表示算子输出与 golden 数据一致、算法逻辑正确。
四、关键源码解读：如何在代码中开启该特性
4.1 配置入口：GetMDLConfig与enableKdimReorderLoad
样例在 matmul_k_reorder_load.asc 中通过模板参数开启该特性：
// enableKdimReorderLoad set to true static constexpr auto CFG_MDL_REORDER = GetMDLConfig(false, false, 0, false, false, false, false, true, true, false, false, true); AscendC::Matmul< AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, AType>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BType>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, CType>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BiasType>, CFG_MDL_REORDER> matmulObj;GetMDLConfig是 MDL（Multi Data Load）模板的配置构造函数，声明于 include/adv_api/matmul/tiling.h，其完整签名与默认值如下：
__aicore__ constexpr MatmulConfig GetMDLConfig( const bool intrinsicsLimit = false, // 单核输入矩阵内轴元素数 >= 65535 时，是否开启 GM 到 L1 的循环搬运 const bool batchLoop = false, // 是否开启多 batch 输入输出（仅 BatchMatmul） const uint32_t doMTE2Preload = 0, // 是否开启 M/N 方向 preload const bool isVecND2NZ = false, // 是否使用向量指令完成 ND 到 NZ 格式转换 bool isPerTensor = false, // A16W8 场景下 B 矩阵是否按 per-tensor 量化 bool hasAntiQuantOffset = false, // A16W8 场景下 B 矩阵是否使用 offset 系数反量化 const bool enUnitFlag = false, // 是否使能 UnitFlag const bool isMsgReuse = true, // SetSelfDefineData 回调的 dataPtr 是否直接传值复用 const bool enableUBReuse = true, // 是否复用 Unified Buffer 空间 const bool enableL1CacheUB = false, // 是否使能 L1 Buffer 缓存 UB 计算块 const bool enableMixDualMaster = false, // 是否使能 MixDualMaster const bool enableKdimReorderLoad = false,// 是否使能 KdimReorder（K 轴错峰加载）★ const bool enableL1BankConflictOptimise = false) // 是否使能 L1BankConflictOptimise- 样例传入
GetMDLConfig(false, false, 0, false, false, false, false, true, true, false, false, true)，第 12 个参数enableKdimReorderLoad = true，即开启 K 轴错峰加载； - 该参数在 include/adv_api/matmul/matmul_config.h 中默认为
false（关闭），注释为 “false: Disable staggered K-axis data loading (default)”； - 若不想手写 13 个位置参数，也可以采用调优文档 matmul_multi_core_k_stagger.md 中“先取 MDL 默认配置再修改成员”的写法：
constexpr MatmulConfig GetMDLKDimReorderConfig() { auto CFG = CFG_MDL; // CFG_MDL 即默认 MDL 模板，见 include/adv_api/matmul/tiling.h CFG.enableKdimReorderLoad = true; return CFG; } constexpr static MatmulConfig MM_CFG = GetMDLKDimReorderConfig();其中CFG_MDL是仓库预置的默认 MDL 配置常量（constexpr MatmulConfig CFG_MDL = GetMDLConfig();，见 include/adv_api/matmul/tiling.h），默认已开启doMultiDataLoad（多数据加载），这是 K 轴错峰特性的前置模板要求。
4.2 内核类结构：Init / Process / CalcOffset
样例内核MatmulKernel<AType, BType, CType, BiasType>是典型的直调 Matmul 编写范式，完整代码见 matmul_k_reorder_load.asc：
- Init：设置 A/B/C/Bias 的 GlobalTensor 缓冲（
aGlobal.SetGlobalBuffer(...)），并通过CalcOffset(GetBlockIdx(), ...)依据当前核的 blockIdx 计算各自的数据偏移，支持isTransA/isTransB两种转置布局的偏移计算； - Process：处理尾核（tail core）场景——当
tailM < singleCoreM || tailN < singleCoreN时调用matmulObj.SetTail(curSingleCoreM, curSingleCoreN)；随后SetTensorA/SetTensorB（可选SetBias）、IterateAll(cGlobal)完成全 K 迭代，最后End()收尾； - CalcOffset：按
mSingleBlocks = (M + singleCoreM - 1) / singleCoreM把核划分为 M/N 二维网格，mCoreIndex = blockIdx % mSingleBlocks、nCoreIndex = blockIdx / mSingleBlocks，据此计算 A、B、C、Bias 各自的 GM 偏移。
4.3 Host 侧 tiling 生成与启动
直调样例在GenerateTiling（matmul_k_reorder_load.asc）中使用matmul_tiling::MultiCoreMatmulTiling生成 tiling 数据：
tilingApi.SetDim(ascendcPlatform->GetCoreNumAiv()); // 使用全部 AIV 核 tilingApi.SetAType(leftPosition, leftFormat, leftDtype, isTransA); // A: GM/ND/float16 tilingApi.SetBType(rightPosition, rightFormat, rightDtype, isTransB); // B: GM/ND/float16 tilingApi.SetBiasType(biasPosition, biasFormat, biasDtype); // bias: GM/ND/float32 tilingApi.SetCType(resultPosition, resultFormat, resultDtype); // C: GM/ND/float32 tilingApi.SetOrgShape(M, N, K); tilingApi.SetShape(M, N, K); tilingApi.EnableBias(isBias); tilingApi.SetBufferSpace(-1, -1, -1); int64_t res = tilingApi.GetTiling(tilingData); // 生成 Matmul tiling 数据随后main中完成 ACL 初始化（aclInit/aclrtSetDevice/aclrtCreateStream）、Host/Device 内存分配与数据搬入，并以matmul_k_reorder_load_custom<<<numBlocks, 0, stream>>>(...)方式直接 launch 内核（matmul_k_reorder_load.asc）。
五、底层实现原理：K 循环迭代器的错峰机制
开启enableKdimReorderLoad后，Matmul 的 K 循环会走专门实现的迭代器KLoop<..., IsKdimReorderLoad<MM_CFG> && ...>，对应内部头文件 impl/adv_api/detail/matmul/scheduler/iterator/k_loop/k_loop_mdl_reorder.h。
核心机制可以从源码推断如下：
错峰起始索引计算：
InitMinStepKStartIdx()中，根据当前blockIdx以及 A/B 两侧的 K 步长stepKa/stepKb计算每个核的最小 K 步起始外循环索引minStepKStartOuterIdx_：- 当
stepKa > stepKb时，minStepKStartOuterIdx_ = (blockIdx % outerKaIter_) * kaStepFactor_； - 否则取
(blockIdx % outerKbIter_) * kbStepFactor_。
即不同核在 K 维外循环上拥有不同的起点，从而在时间上错开对 GM 不同地址块的访问。
- 当
循环轮转：
OuterStart()从错峰起点开始，OuterNext()通过(outerIdx_ + 1) % outIter_回绕推进，OuterEnd()以“回到起点”作为终止条件，保证每个核仍完整遍历全部 K 分块，只是遍历顺序/相位不同——这就是“错峰”（staggered/peak-shifting）的本质：计算总量与结果不变，访问相位被拉开。配套判定：该专用迭代器通过模板偏特化启用，条件为
DoMatmulMDL(MM_CFG) && IsKdimReorderLoad<MM_CFG> && !IsNeedUB() && !HasScalePosition<A_TYPE>，即必须是 MDL 模板、开启 KdimReorder、且 A 无 scale 位置（详见 k_loop_mdl_reorder.h）。从该模板偏特化条件也可以印证文档中的约束：该特性仅支持 MDL 模板。
仓库中还提供了对应的单测 tests/api/adv_api/matmul/iterator/test_k_loop_mdl_reorder.cpp，以stepKa=2, stepKb=4, baseK=32, singleCoreK=512为例验证：总迭代数GetTotalIter() == 16、FirstOuterIter/LastOuterIter/FirstInnerIter的错峰边界判断、outerIdx最终回绕到 0 等行为，直接佐证了上述错峰循环逻辑的正确性。
六、编译与运行
6.1 配置环境变量
根据当前环境上 CANN 开发套件包的安装方式（参见 docs/en/quick_start.md），配置环境变量：
source ${install_path}/cann/set_env.sh
${install_path}为 CANN 包安装目录；未指定安装目录时默认安装至/usr/local/Ascend。
6.2 构建并运行样例
在样例目录examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_k_reorder_load/下执行：
mkdir -p build && cd build; # 创建并进入 build 目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..; make -j; # 编译工程，默认 npu 模式 python3 ../scripts/gen_data.py # 生成测试输入数据（input/*.bin）与 golden 结果（output/golden.bin） ./demo # 执行编译生成的可执行程序，运行样例 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 校验输出精度，确认算法逻辑正确如需CPU 调试模式或NPU 仿真模式，在 cmake 时追加运行模式参数：
cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..; make -j; # CPU 调试模式 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..; make -j; # NPU 仿真模式注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下执行
rm CMakeCache.txt后重新 cmake。
6.3 编译选项说明
| 选项 | 可选值 | 说明 |
|---|---|---|
CMAKE_ASC_RUN_MODE | npu（默认）、cpu、sim | 运行模式：NPU 运行、CPU 调试、NPU 仿真 |
CMAKE_ASC_ARCHITECTURES | dav-2201（默认）、dav-3510 | NPU 架构：dav-2201对应 Atlas A2/A3 系列产品，dav-3510对应 Ascend 950PR & 950DT 系列产品 |
这两个选项同样定义在样例的 CMakeLists.txt 中（CACHE STRING类型，可通过-D覆盖）；工程通过find_package(ASC REQUIRED)引入 ASC 编译器支持，并以add_executable(demo matmul_k_reorder_load.asc)构建可执行文件。
6.4 数据生成与精度校验
- scripts/gen_data.py：生成
768×6144的 A（float16）、6144×2048的 B（float16）、1×2048的 bias（float32），并用 NumPy 计算C = A×B + bias作为 golden； - scripts/verify_result.py：以
rtol=1e-3、atol=1e-3、误差占比阈值1e-3对比 output 与 golden，全部满足时打印：
test pass!若出现不匹配，脚本会打印前 100 个错误元素的下标、期望值、实际值与相对误差，以及整体 error ratio，便于定位问题。
七、调优视角：该特性带来的性能收益（仓库案例佐证）
仓库调优文档 matmul_multi_core_k_stagger.md 记录了该特性在同规格算子上的真实调优过程，可直接作为该样例收益的参考：
- 优化前瓶颈：矩阵按 M/N 方向切分，单核 shape 为
[128, 512, 6144]、基本块 shape 为[128, 256, 64]；Profiling 显示aic_mte2_ratio平均达 0.93，算子为 MTE2 Bound，MTE2 平均带宽利用率仅 34.4%，整体耗时为 98.72 us——多核同时访问同一 GM 地址是 MTE2 效率低的主因； - 开启 K 轴错峰后（Tiling 参数不变）：MTE2 平均耗时从 90 us 降至 69.87 us、最大耗时从 91.94 us 降至 75.82 us；MTE2 平均带宽利用率提升至 41.7%；算子整体耗时从 98.72 us 降至 85.68 us，性能提升约 13.2%。
（上图为仓库调优文档中“多核访问地址冲突”的示意，下图为开启 K 轴错峰后“缓解地址冲突”的示意，均来自 docs/zh/guide/figures/addr_conflict.png 与 docs/zh/guide/figures/ease_addr_conflict.png。）
需要说明：上述收益数据来自仓库内该调优案例在特定软硬件环境下的测试结果，实际收益与算子 shape、核数、带宽等因素相关，建议在目标产品上使用 msOpProf 工具（算子仿真流水图与上板 Profiling）实测 MTE2 流水与带宽利用率后再决定是否开启。
八、总结
K 轴错峰加载数据（enableKdimReorderLoad）是 asc-devkit 高阶 Matmul API 中一项面向多核 + 大 K 轴 + 矩阵非全载场景的性能优化特性：
- 一句话原理：通过让不同核在 K 维外循环上拥有不同的起始相位（
k_loop_mdl_reorder.h中的minStepKStartOuterIdx_机制），把对 GM 的同步同地址访问变成错峰的不同地址访问，缓解地址冲突、提升 MTE2 搬运效率； - 开启方式：使用 MDL 模板（
GetMDLConfig/CFG_MDL），并将enableKdimReorderLoad置为true后作为AscendC::Matmul的模板参数传入； - 约束条件：K 建议不小于 4096、A/B 非全载、多核执行、A 或 B 位于 GM、仅支持 MDL 模板；
- 参考实现：可直接复用 matmul_k_reorder_load 样例 的直调框架（内核 + tiling + host 启动 + 精度校验），并通过仓库调优文档了解其预期的性能收益与分析方法。
在动手之前，建议同时参考 MatmulConfig.md 与 GetMDLConfig.md 了解完整配置语义，并结合 Profiling 数据确认 MTE2 是否为当前算子瓶颈，再做开关决策。
- 人工智能
- 深度学习
- 算子库
- CANN
- Ascend
【免费下载链接】asc-devkit
本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。
相关推荐
asc-devkit Matmul M/N 轴流水并行实战：matmul_mn_double_buffer 样例深度解析
asc devkit Matmul M/N 轴流水并行实战：matmul_mn_double_buffer 样例深度解析 导读 在昇腾 AI 处理器算子开发中，
人工智能深度学习算子库CANNAscendasc-devkit Matmul 高阶 API 实战：IBShareB 直调样例原理与实现
asc devkit Matmul 高阶 API 实战：IBShareB 直调样例原理与实现 导读 本篇技术指南以 asc devkit 仓库中的 matmul
人工智能深度学习算子库CANNAscendasc-devkit Matmul SplitM 策略直调样例解析：多核切 M 场景下的 Matmul 编程实战
asc devkit Matmul SplitM 策略直调样例解析：多核切 M 场景下的 Matmul 编程实战 导读 本文围绕 asc devkit 仓库中的
人工智能深度学习算子库CANNAscend
创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
-物理系统)
