• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend

导读

本篇文章围绕 asc-devkit 仓库中的 Matmul L0 缓存（L0 Cache）直调样例展开，讲解如何在 AscendC 算子中利用 Matmul API 的 L0 缓存特性，通过让 A 矩阵数据在 A2 缓冲区驻留复用、与不同 B2 数据依次完成乘加运算，从而减少 MTE1 的重复搬运。读完本文，你将掌握 L0 缓存的使能条件（全部常量化 Tiling 配置、singleCoreM=baseM且singleCoreK=baseK）、样例工程的编译运行方式（NPU/CPU/仿真三种模式），以及该特性在 Matmul 库内部的自动判定与缓存命中实现原理。

概述：什么是 Matmul L0 缓存

以左矩阵 A 为例，Matmul 的 Cube 计算流水线中，A 矩阵数据需要先由 MTE2 从 GM 搬运到 L1（A1 缓冲区），再经 MTE1 从 L1 搬运到 L0（A2 缓冲区）后才能参与乘加运算。所谓L0 缓存，指的是：当 A1 中数据搬运到 A2 后，该份数据保持在 A2 中不释放，后续循环中缓存数据依次与 B2 中不同的数据完成乘加运算，从而避免同一份 A 数据被重复执行 MTE1 搬运。其核心收益即文档所述：减少 MTE1 重复搬运。

说明：L0 缓存没有对外开关，是否开启由 Matmul API 内部根据用户配置的 shape 信息自动推导，开发者只需按照使能条件配置 Tiling 即可。

使能条件：L0 缓存的开启前提

根据 matmul_l0cache 样例文档，当前仅支持L0A 缓存，使能场景要求如下：

• 需要配置全部常量化 Tiling 场景：singleCoreM/singleCoreN/singleCoreK和baseM/baseN/baseK均需通过常量化接口设置；
• singleCoreM = baseM同时singleCoreK = baseK。

从源码看自动判定逻辑

Matmul 库内部在编译期通过IsL0ACache/IsL0Cache常量表达式函数自动判断是否启用 L0 缓存，实现在 matmul_utils.h：

template <typename T, const auto& MM_CFG> __aicore__ inline constexpr bool IsL0ACache() { return (ToMatmulConfig(MM_CFG).singleCoreK <= ToMatmulConfig(MM_CFG).basicK) && (ToMatmulConfig(MM_CFG).singleCoreM <= ToMatmulConfig(MM_CFG).basicM); }

IsL0Cache在IsL0ACache之外还叠加了多项前提检查，包括：不支持含 scale 位置（MX 量化）的 A 类型、不支持 Load2dV2 特性、需要开启 Norm/MDL 调度、不允许intraBlockPartSum、A/B 布局需为LayoutMode::NONE、不允许isA2B2Shared等，且singleCoreM等六个 shape 参数均需大于 0，最终要求IsL0ACache && IsL0BCache同时成立。

可见文档中“singleCoreM=baseM且singleCoreK=baseK”正是命中IsL0ACache判定的核心 shape 条件，其余条件则由 Matmul API 的默认模板（CONFIG_NORM）与使用约束共同保证。

缓存命中机制：TBufPoolL0 与调度器的 Hit 判断

L0 缓存的落地实现位于缓存池专用特化类 tbuf_pool_l0_cache.h，该类通过模板特化条件IsL0Cache<A_TYPE, MM_CFG>()与常规缓存池（tbuf_pool_l0_common.h、tbuf_pool_l0_v200.h）分流：

• 维护l0aPingCachedPos_/l0aPongCachedPos_（以及对应的 l0b 两组）缓存位置标记，记录当前驻留在 A2/B2 中的数据块位置；
• Hit<TPosition::A2>(pos)判断请求的数据块是否已在缓存中命中，命中时返回 true 并切换 ping/pong 标志；
• ResetCache()通过WaitFlag/SetFlag<HardEvent::M_MTE1>事件同步解决多对象间的同步问题，并将缓存位置重置为 -1。

调度器侧（Norm 调度 scheduler_norm.h）的SplitA/SplitB配合Hit实现“命中即复用”：

// posA 由 MLoop 内层迭代序号与 KLoop 迭代序号共同计算 if (!(MATMUL_MODULE(TBufPoolL0)->template Hit<TPosition::A2>(posA))) { a2 = MATMUL_MODULE(TBufPoolL0)->template GetBuffer<TPosition::A2, TransAT>(); MATMUL_MODULE(LoadToA2)->Load(...); // 未命中：执行 MTE1 搬运 } else { a2 = MATMUL_MODULE(TBufPoolL0)->template GetBuffer<TPosition::A2, TransAT>(); // 命中：跳过 MTE1 搬运，直接复用 A2 中驻留的数据 }

当singleCoreM == baseM时 M 方向单核内无需切分，同一份 A 数据块在 K 方向外层循环中可反复命中，这正是“减少 MTE1 重复搬运”的底层来源。Hit中注释singleMK == baseMK, no need to pingpong也印证了该设计。

样例工程结构

matmul_l0cache 样例位于 examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_l0cache，目录结构如下：

算子实现剖析：matmul_l0cache.asc

1. Shape 与常量 Tiling 配置

算子以常量方式定义矩阵规模与单核/基础分块参数（matmul_l0cache.asc）：

constexpr uint32_t M = 2560; // 原始矩阵 M 维 constexpr uint32_t N = 2048; // 原始矩阵 N 维 constexpr uint32_t K = 128; // 原始矩阵 K 维 constexpr uint32_t NUM_BLOCKS = 20; // 参与计算的核数 constexpr bool IS_TRANS_A = false; constexpr bool IS_TRANS_B = false; constexpr bool IS_BIAS = true; constexpr static int32_t SINGLE_M = 256; // 单核 M 维 shape constexpr static int32_t SINGLE_N = 1024; // 单核 N 维 shape constexpr static int32_t SINGLE_K = 128; // 单核 K 维 shape constexpr static int32_t BASE_M = 256; // 基础分块 M，BASEM * BASEK * sizeof(typeA) <= L0A size constexpr static int32_t BASE_N = 128; // 基础分块 N，BASEN * BASEK * sizeof(typeB) <= L0B size constexpr static int32_t BASE_K = 128; // 基础分块 K，BASEM * BASEN * sizeof(typeC) <= L0C size

注意这里SINGLE_M == BASE_M、SINGLE_K == BASE_K，恰好满足文档给出的 L0A 缓存使能条件。注释中还给出了 L0 容量约束：baseM * baseK * sizeof(typeA) <= L0A size、baseN * baseK * sizeof(typeB) <= L0B size、baseM * baseN * sizeof(typeC) <= L0C size，对应源码 matmul_shape_tiling.h 中StaticTilingCheck的static_assert检查。

2. 常量化接口组装 Matmul 对象

通过MatmulShapeParams结构体（定义见 matmul_config.h，字段包括singleCoreM/singleCoreN/singleCoreK与basicM/basicN/basicK）与GetMMConfig<MatmulConfigMode::CONFIG_NORM>、GetMatmulApiTiling（定义见 constant_tiling.h）完成编译期常量 Tiling 推导：

constexpr static MatmulShapeParams shapeParams = {SINGLE_M, SINGLE_N, SINGLE_K, BASE_M, BASE_N, BASE_K}; constexpr static MatmulConfig CUSTOM_CFG = GetMMConfig<MatmulConfigMode::CONFIG_NORM>(shapeParams); constexpr static auto CONSTANT_CFG = AscendC::GetMatmulApiTiling<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE>(CUSTOM_CFG); AscendC::Matmul<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, CONSTANT_CFG> matmulObj;

GetMatmulApiTiling在编译期把singleCoreM/singleCoreN/singleCoreK写入MatmulApiStaticTiling的M/N/Ka/Kb及baseM/baseN/baseK，实现“全部常量化 Tiling”。A/B/C/Bias 类型均通过AscendC::MatmulType<TPosition::GM, CubeFormat::ND, T, IS_TRANS_A>声明，本样例为half（A、B）×float（C、Bias）。

3. Kernel 主体流程

MatmulKernel的Init中通过SetGlobalBuffer绑定 GM 上的 A/B/C/Bias 张量，并调用matmulObj.SetOrgShape(tiling.M, tiling.N, tiling.Ka, tiling.Kb)设置原始 shape；Process中先根据GetBlockIdx()判断尾块（tail）并调用SetTail，随后依次执行：

matmulObj.SetTensorA(aGlobal, IS_TRANS_A); matmulObj.SetTensorB(bGlobal, IS_TRANS_B); if (tiling.isBias) { matmulObj.SetBias(biasGlobal); } matmulObj.IterateAll(cGlobal); matmulObj.End();

多核切分逻辑在CalcOffset中完成：按mSingleBlocks = (M + singleCoreM - 1) / singleCoreM把 M 维切分到多个核，mCoreIndex = blockIdx % mSingleBlocks、nCoreIndex = blockIdx / mSingleBlocks，并据此计算 A/B/C/Bias 的偏移。核入口matmul_l0cache_custom通过CopyTiling从 GM 拷贝 Tiling，并用REGIST_MATMUL_OBJ注册 Matmul 对象。

4. Tiling 生成

GenerateTiling使用matmul_tiling::MultiCoreMatmulTiling生成运行时 Tiling：SetDim(NUM_BLOCKS)设置核数，SetAType/SetBType/SetCType/SetBiasType设置各矩阵类型（A/B 为DT_FLOAT16，C/Bias 为DT_FLOAT），SetOrgShape与SetShape(M, N, K)设置矩阵规模，SetSingleShape(SINGLE_M, SINGLE_N, SINGLE_K)设置单核 shape，EnableBias(IS_BIAS)使能 Bias，SetBufferSpace(-1, -1, -1)使用默认缓冲空间，最后GetTiling得到 Tiling 数据并保存到 buffer。

5. ACL 运行框架

main函数按ASCENDC_CPU_DEBUG与__NPU_ARCH__选择平台实例（CPU 调试时dav-2201对应 "Ascend910B1"，dav-3510对应 "Ascend950PR_9589"），完成aclInit→aclrtSetDevice→aclrtCreateContext→aclrtCreateStream初始化后，通过aclrtMallocHost/aclrtMalloc申请主机与设备内存，读取./input/x1_gm.bin、./input/x2_gm.bin、./input/bias_gm.bin并拷贝到设备，以matmul_l0cache_custom<<<NUM_BLOCKS, 0, stream>>>(...)启动 20 核 Kernel，同步后将结果写回./output/output.bin，最后按序释放资源。

编译运行：三种模式完整步骤

配置环境变量

根据当前环境上 CANN 开发套件包的安装方式，配置环境变量：

source ${install_path}/cann/set_env.sh

说明：${install_path}为 CANN 包安装目录，未指定安装目录时默认安装至/usr/local/Ascend下。

NPU 模式编译与执行

在样例根目录examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_l0cache下执行：

mkdir -p build && cd build; # 创建并进入 build 目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认 npu 模式 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译生成的可执行程序，执行样例 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确

CPU 调试与 NPU 仿真模式

使用 CPU 调试或 NPU 仿真模式时，添加-DCMAKE_ASC_RUN_MODE=cpu或-DCMAKE_ASC_RUN_MODE=sim参数即可：

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。

编译选项说明

这两个选项在 CMakeLists.txt 中以CACHE STRING形式声明，编译时通过--npu-arch=${CMAKE_ASC_ARCHITECTURES}传递给 ASC 编译器，并链接tiling_api、register、platform、m、dl库。

数据生成与精度校验

生成输入与 golden 数据

gen_data.py 以与算子一致的m, n, k = 2560, 2048, 128生成数据：

• A 矩阵x1_gm：np.random.uniform(-1, 1, [m, k])转float16；
• B 矩阵x2_gm：np.random.uniform(-1, 1, [k, n])转float16；
• Bias：np.random.uniform(-10, 10, [n])转float32；
• golden：np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32)) + bias_gm，即标准矩阵乘加 Bias 的结果。

数据分别写入./input/x1_gm.bin、./input/x2_gm.bin、./input/bias_gm.bin与./output/golden.bin。

结果校验

verify_result.py 通过np.isclose对比输出与 golden，容差设置为：

RELATIVE_TOL = 1e-3 # 相对误差容限 ABSOLUTE_TOL = 1e-3 # 绝对误差容限 ERROR_TOL = 1e-3 # 允许的错误元素比例上限

错误元素比例不超过ERROR_TOL时打印并返回成功，否则逐条打印前 100 个差异元素的索引、期望值、实际值与相对偏差。执行成功后终端输出：

test pass!

支持的产品与 CANN 版本

小结

matmul_l0cache 样例完整展示了 AscendC Matmul L0 缓存特性的使用路径：通过MatmulShapeParams+GetMMConfig<CONFIG_NORM>+GetMatmulApiTiling实现全部常量化 Tiling，并使singleCoreM=baseM、singleCoreK=baseK触发库内IsL0ACache编译期判定，最终由TBufPoolL0缓存池与 Norm 调度器的Hit命中机制在运行时复用 A2 驻留数据、跳过重复的 MTE1 搬运。对于希望优化 Matmul 算子 MTE1 搬运开销的开发者，可参照本样例的 shape 配置与编译运行流程，并结合本文给出的源码依据深入理解其判定与调度细节。

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend