add mmad best practice-asc-devkit-AtomGit

已合并

add mmad best practice #1648

zc1110创建于 4月16日从zhengchao0613/asc-devkit_4553： master 合入到 cann/asc-devkit： master

add mmad best practice #1648

已合并

zc1110创建于 4月16日

从zhengchao0613/asc-devkit_4553 master合入到cann/asc-devkit master

共 7 个文件变更+1353-2

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_basic_API_high_performance/CMakeLists.txt+43-0

@@ -0,0 +1,43 @@
	1	+# ----------------------------------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------------------------------------
	10	+
	11	+
	12	+cmake_minimum_required(VERSION 3.16)
	13	+
	14	+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
	15	+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
	16	+set(SCENARIO_NUM 1 CACHE STRING "Scenario number to compile (1-2)")
	17	+
	18	+find_package(ASC REQUIRED)
	19	+
	20	+project(kernel_samples LANGUAGES ASC CXX)
	21	+
	22	+add_executable(demo
	23	+ mmad.asc
	24	+)
	25	+
	26	+if(CMAKE_ASC_ARCHITECTURES STREQUAL "dav-2201")
	27	+ set(NPU_ARCH_ID 2201)
	28	+elseif(CMAKE_ASC_ARCHITECTURES STREQUAL "dav-3510")
	29	+ set(NPU_ARCH_ID 3510)
	30	+else()
	31	+ message(FATAL_ERROR "Unsupported CMAKE_ASC_ARCHITECTURES=${CMAKE_ASC_ARCHITECTURES}, only dav-2201 and dav-3510 are supported")
	32	+endif()
	33	+
	34	+message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}")
	35	+
	36	+target_compile_definitions(demo PRIVATE
	37	+ NPU_ARCH=${NPU_ARCH_ID}
	38	+ SCENARIO_NUM=${SCENARIO_NUM}
	39	+)
	40	+
	41	+target_compile_options(demo PRIVATE
	42	+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
	43	+)

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_basic_API_high_performance/README.md+472-0

@@ -0,0 +1,472 @@
	1	+# Matmul 基础API最佳实践样例
	2	+
	3	+## 概述
	4	+
	5	+本样例基于静态Tensor编程范式，通过 L1/L0 双缓冲机制、大包搬运、细粒度流水同步、UnitFlag、L2Cache等多种优化手段，实现高性能矩阵乘法运算。本样例基于基础API实现，与高阶API版本所使用的优化方法相同，样例目标为基于静态Tensor编程范式展示调优实现细节。
	6	+## 支持的产品
	7	+
	8	+- Ascend 950PR/Ascend 950DT
	9	+- Atlas A3 训练系列产品/Atlas A3 推理系列产品
	10	+- Atlas A2 训练系列产品/Atlas A2 推理系列产品
	11	+
	12	+## 目录结构介绍
	13	+
	14	+```
	15	+├── matmul_basic_API_high_performance
	16	+│ └── scripts
	17	+│ ├── gen_data.py // 输入数据和真值数据生成脚本文件
	18	+│ └── verify_result.py // 真值对比文件
	19	+│ ├── CMakeLists.txt // 编译工程文件
	20	+│ ├── data_utils.h // 数据读入写出函数
	21	+│ └── mmad.asc // Ascend C样例实现
	22	+```
	23	+
	24	+## 样例描述
	25	+
	26	+ Matmul计算公式：
	27	+ $$
	28	+ C = A * B
	29	+ $$
	30	+ - A、B为源操作数，A为左矩阵，形状为[M, K]；B为右矩阵，形状为[K, N]
	31	+ - C为目的操作数，存放矩阵乘结果的矩阵，形状为[M, N]
	32	+
	33	+- 样例规格：
	34	+
	35	+
	36	+
	37	+ <table>
	CaiBinbin4月20日表格格式有问题
	38	+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="5" align="center">Matmul</td></tr>
	39	+ </tr>
	40	+ <tr><td rowspan="3" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td><td align="center">isTrans</td></tr>
	41	+ <tr><td align="center">A</td><td align="center">[M, K]</td><td align="center">float16</td><td align="center">ND</td><td align="center">false</td></tr>
	42	+ <tr><td align="center">B</td><td align="center">[K, N]</td><td align="center">float16</td><td align="center">ND</td><td align="center">true</td></tr>
	43	+ </tr>
	44	+ </tr>
	45	+ <tr><td rowspan="1" align="center">样例输出</td><td align="center">C</td><td align="center">[M, N]</td><td align="center">float16</td><td align="center">ND</td><td align="center">-</td></tr>
	46	+ </tr>
	47	+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="5" align="center">mmad_custom</td></tr>
	48	+ </table>
	49	+## 样例实现
	50	+
	51	+### 性能指标说明
	52	+
	53	+\| 指标 \| 说明 \|
	54	+\|------\|------\|
	55	+\| Task Duration(μs) \| 整个任务执行的总时间，算子执行时间以该参数为准 \|
	56	+\| Block Num \| 使用的核数（Block数量） \|
	57	+\| aicore_time(μs) \| AI Core 的平均执行时间 \|
	58	+\| aic_mac_time(μs) \| Cube 计算单元的执行时间 \|
	59	+\| aic_mac_ratio \| Cube 计算单元的时间占比，反映计算单元利用率 \|
	60	+\| aic_mte1_time(μs) \| MTE1（L1 到 L0A/L0B 搬运）的执行时间 \|
	61	+\| aic_mte1_ratio \| MTE1 的时间占比，反映 L1 到 L0 的数据搬运压力 \|
	62	+\| aic_mte2_time(μs) \| MTE2（GM 到 L1 搬运）的执行时间 \|
	63	+\| aic_mte2_ratio \| MTE2 的时间占比，反映 GM 到 L1 的数据加载压力 \|
	64	+\| aic_fixpipe_time(μs) \| FixPipe（L0C 到 GM 搬运）的执行时间 \|
	65	+\| aic_fixpipe_ratio \| FixPipe 的时间占比，反映结果写回的访存压力 \|
	66	+
	67	+
	68	+### 数据流路径：
	69	+```
	70	+GM ──(MTE2, DataCopy)──> L1 ──(MTE1, LoadData)──> L0A/L0B ──(Cube, Mmad)──> L0C ──(Fixpipe)──> GM
	71	+ DataCopyInA/B DataLoadA/B Compute CopyOut
	72	+```
	73	+
	74	+### 核心特性
	75	+
	76	+#### 1. L1/L0 双缓冲 Ping-Pong 布局
	77	+
	78	+L1 和 L0 均采用 Ping-Pong 双缓冲，使得 DataCopyIn（GM→L1）、DataLoad（L1→L0）与 Compute 三个阶段形成三级流水，每个阶段处理不同缓冲区的数据，互不阻塞。
	79	+
	80	+```
	81	+时间 ──────────────────────────────────────────────────────────────>
	82	+
	83	+MTE2: \|─ A1 Ping ──\|─ A1 Pong ──\|─ A1 Ping ──\| ...
	84	+MTE1: \|─ A2 Ping ──\|─ A2 Pong ──\|─ A2 Ping ──\| ...
	85	+Cube: \|─ Mmad ─────\|─ Mmad ─────\| ...
	86	+Fixpipe: \|─ CopyOut ──\|(unitflag)
	87	+```
	88	+
	89	+
	90	+L1 双缓冲布局：A1 占 L1 前半（0~256KB），B1 占 L1 后半（256~512KB），各自再分为 Ping/Pong 两块：
	91	+
	92	+```
	93	+L1 (512KB):
	94	+├── A1 Ping: [0, 128KB)
	95	+├── A1 Pong: [128KB, 256KB)
	96	+├── B1 Ping: [256KB, 384KB)
	97	+└── B1 Pong: [384KB, 512KB)
	98	+```
	99	+
	100	+L0 双缓冲布局：A2/B2 各自独立64KB空间，各自再分为 Ping/Pong：
	101	+
	102	+```
	103	+L0A/L0B (64KB):
	104	+├── A2 Ping: [0, 16KB)
	105	+├── A2 Pong: [32KB, 48KB)
	106	+├── B2 Ping: [0, 32KB)
	107	+└── B2 Pong: [32KB, 64KB)
	108	+```
	109	+
	110	+```cpp
	111	+// A1: L1 Ping/Pong
	112	+AscendC::LocalTensor<half> a1LocalPing(AscendC::TPosition::A1, 0, a1PingpongSize);
	113	+AscendC::LocalTensor<half> a1LocalPong(AscendC::TPosition::A1, a1PingpongSize * sizeof(half), a1PingpongSize);
	114	+// A2: L0 Ping/Pong
	115	+AscendC::LocalTensor<half> a2LocalPing(AscendC::TPosition::A2, 0, a2PingpongSize);
	116	+AscendC::LocalTensor<half> a2LocalPong(AscendC::TPosition::A2, L0_PINGPONG_BYTES, a2PingpongSize);
	117	+```
	118	+
	119	+#### 2. 大包搬运
	120	+
	121	+通过 `stepKa`/`stepKb` 参数将多个基本块打包为一次 DataCopyIn 操作（称为"大包"），减少 MTE2 搬运次数。例如 `stepKa=8` 表示一次将 8 个 baseM * baseK 块从 GM 搬入 L1。
	122	+
	123	+
	124	+```cpp
	125	+// DataCopyInA: 一次搬入 stepKa 个 baseK 块
	126	+AscendC::Nd2NzParams nd2nzParams;
	127	+nd2nzParams.nValue = curM;
	128	+nd2nzParams.dValue = baseK * stepKa; // 大包包含 stepKa 个 baseM * baseK
	129	+```
	130	+
	131	+#### 3. 细粒度流水同步
	132	+
	133	+使用四类硬件事件标志实现精确的流水线同步，分为正向同步（数据就绪通知）和反向同步（缓冲区释放通知）：
	134	+
	135	+\| 事件类型 \| 方向 \| 用途 \| flag 编号 \|
	136	+\|---------\|------\|------\|----------\|
	137	+\| MTE2_MTE1 \| 正向 \| L1 数据就绪通知，DataCopyIn 通知 DataLoad 可读取 \| 0/1: A1 Ping/Pong; 2/3: B1 Ping/Pong \|
	138	+\| MTE1_MTE2 \| 反向 \| L1 缓冲区释放通知，DataLoad 通知 DataCopyIn 可写入 \| 同上 \|
	139	+\| MTE1_M \| 正向 \| L0 数据就绪通知，DataLoad 通知 Compute 可计算 \| mte1DBFlag (0/1 交替) \|
	140	+\| M_MTE1 \| 反向 \| L0 缓冲区释放通知，Compute 通知 DataLoad 可写入 \| mte1DBFlag (0/1 交替) \|
	141	+
	142	+反向同步需预置：由于反向同步是"消费方 SetFlag → 生产方 WaitFlag"，首次使用前必须预置 SetFlag，否则首次 WaitFlag 会死锁：
	143	+
	144	+```cpp
	145	+// 初始化：预置反向同步 flag，防止首次 WaitFlag 死锁
	146	+AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(0); // A1 Ping 可写
	147	+AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(1); // A1 Pong 可写
	148	+AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(2); // B1 Ping 可写
	149	+AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(3); // B1 Pong 可写
	150	+AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(0); // L0 Ping 可写
	151	+AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(1); // L0 Pong 可写
	152	+```
	153	+
	154	+大包粒度的正向同步：DataLoad 从 L1 大包中按K方向逐个base块读取数据，只需在读取大包首个base块时等待数据就绪，后续base块与首个在同一个大包内，无需重复等待：
	155	+
	156	+```cpp
	157	+// 正向同步：仅大包首元素需等待数据就绪
	158	+if (kOffsetInChunkA == 0) {
	159	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(a1ReadIdx);
	160	+}
	161	+// 反向同步：仅大包末元素消费完后通知 DataCopyIn 可覆盖
	162	+if ((kOffsetInChunkA + 1) == stepKa) {
	163	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(a1ReadIdx);
	164	+}
	165	+```
	166	+
	167	+#### 4. LoadData3D 替代 LoadData2D——减少指令队列占用
	168	+
	169	+在 Atlas A2/A3 架构上，本样例使用 `LoadData3DParamsV2`（即 LoadData3D）替代 `LoadData2DParams`（即 LoadData2D）完成 L1→L0 的数据搬运。这是一个关键的指令队列优化。
	170	+
	171	+问题背景：MTE1 指令队列深度为 32。使用 LoadData2D 时，由于单条 LoadData2D 指令搬运粒度有限，搬运一个 baseM×baseK 的切片需要用 for 循环发射多条 LoadData2D 指令。例如 baseM=128、baseK=64 时，最少需要发射 `baseK/16 = 4` 条 LoadData2D 指令。
	172	+
	173	+> 注意：Atlas A5 芯片提供了 `LoadData2DParamsV2` 接口，单条指令即可完成搬运，无需 LoadData3D。因此本样例通过 `__NPU_ARCH__` 条件编译区分两种架构的 LoadData 实现。
	174	+
	175	+#### 5. 多核并行切分
	176	+
	177	+按 M/N 方向均匀切分矩阵到多核并行计算，4×6 切分策略（M 方向 4 块、N 方向 6 块，共 24 核）满足地址 512B 对齐，并减少同地址访问冲突：
	178	+
	179	+```cpp
	180	+constexpr uint32_t mIter = AscendC::DivCeil(M, singleCoreM);
	181	+uint32_t mIterIdx = get_block_idx() % mIter;
	182	+uint32_t nIterIdx = get_block_idx() / mIter;
	183	+```
	184	+
	185	+#### 6. 常量 Tiling
	186	+
	187	+所有 Tiling 参数（baseM/baseK/baseN、singleCoreM/K/N、stepKa/stepKb）通过模板参数在编译期确定，运行时无需 Scalar 动态计算，减少 Scalar 开销：
	188	+
	189	+```cpp
	190	+template <uint32_t M, uint32_t K, uint32_t N, uint32_t baseM, uint32_t baseK, uint32_t baseN,
	191	+ uint32_t singleCoreM, uint32_t singleCoreK, uint32_t singleCoreN,
	192	+ uint32_t stepKa, uint32_t stepKb>
	193	+class KernelMmad { ... };
	194	+```
	195	+
	196	+#### 7. UnitFlag 优化
	197	+
	198	+开启 UnitFlag 后，MMAD 和 FIXPIPE 实现细粒度（512B）流水并行，而非指令级同步。每当 Cube 完成一个 512B 数据结果的计算，FIXPIPE 立即搬出该数据，Cube 计算与结果写回流水重叠：
	199	+
	200	+```cpp
	201	+mmadParams.unitFlag = (kBlockIdx != kLoopCount - 1) ? 2 : 3; // 开启 UnitFlag
	202	+```
	203	+
	204	+- `unitFlag = 2`：中间 K 块，MMAD 计算结果不立即写回，而是等待下一个 512B 完成后流水搬出
	205	+- `unitFlag = 3`：最后一个 K 块，通知 FIXPIPE 将所有剩余结果写回 GM
	206	+
	207	+#### 8. DataCopyIn 优先搬运 B 矩阵
	208	+
	209	+当 `stepKa > stepKb` 时，B 矩阵每 `stepKb` 个 baseK 就需要切换到下一个 L1 缓冲区（Pong），而 A 矩阵需要 `stepKa` 个 baseK 才切换。因此 B 的数据需求更紧迫。如果先搬运 A，MTE2 流水被 A 占用，B 的搬运要等 A 搬完才能开始，导致需要 B 数据时 B 尚未就绪。
	210	+
	211	+本样例在 Compute 之后触发 DataCopyIn 时，先 B 后 A，优先搬运更紧迫的 B 数据：
	212	+
	213	+```
	214	+k=0: Compute → DataCopyIn(B1 Pong) → DataCopyIn(A1 Pong)
	215	+ ↑ MTE2 先搬 B ↑ 再搬 A
	216	+k=stepKb: 需要 B1 Pong → 已就绪 ✓（B 已有 stepKb 轮时间搬运）
	217	+k=stepKa: 需要 A1 Pong → 已就绪 ✓（A 有 stepKa 轮时间搬运，更充裕）
	218	+```
	219	+
	220	+#### 9. L2Cache 优化
	221	+
	222	+L2Cache 是 AI Core 共享的外部缓存，纯读带宽约为 GM 的 3 到 4 倍。若数据无法命中 L2Cache，需要访问 GM，带宽利用效率较低，导致 MTE2 成为性能瓶颈。
	223	+
	224	+L2Cache 切分的具体实现与[高阶 API Matmul 样例](../matmul_high_performance/README.md)中的 Case 6 一致，核心思路均为将 A 矩阵 M 轴切分，使 B 矩阵跨轮驻留 L2Cache。本样例通过 `ProcessL2Cache()` 方法实现双重外层循环，按 `outerMIdx` 分轮调度，每轮内 24 核并行计算各自的子块：
	225	+
	226	+```cpp
	227	+// ProcessL2Cache: 按 M 方向分轮，每轮 24 核覆盖 mIterPerRound 份 M 子块
	228	+constexpr uint32_t mIterPerRound = AscendC::DivCeil(M, singleCoreM * 2);
	229	+constexpr uint32_t outerMLoopCount = AscendC::DivCeil(mIterTotal, mIterPerRound);
	230	+
	231	+for (uint32_t outerMIdx = 0; outerMIdx < outerMLoopCount; outerMIdx++) {
	232	+ uint32_t mIterIdx = AscendC::GetBlockIdx() % mIterPerRound + outerMIdx * mIterPerRound;
	233	+ uint32_t nIterIdx = AscendC::GetBlockIdx() / mIterPerRound;
	234	+ if (mIterIdx >= mIterTotal \|\| nIterIdx >= nIterTotal) continue;
	235	+ InitComputeParamsL2Cache(mIterIdx, nIterIdx);
	236	+ ProcessLoop(...);
	237	+}
	238	+```
	239	+
	240	+#### 10. K 方向主循环完整流程
	241	+
	242	+以 stepKa=8、stepKb=4 为例，展示一个完整 (mBlockIdx, nBlockIdx) 子块内 K 方向循环的详细执行流程：
	243	+
	244	+```
	245	+预处理：
	246	+ SetFlag(MTE1_MTE2, 0/1/2/3) // 预置反向同步：L1 Ping/Pong 均可写
	247	+ SetFlag(M_MTE1, 0/1) // 预置反向同步：L0 Ping/Pong 均可写
	248	+ DataCopyIn(A1 Ping, k=0) // 搬入首个 A 大包到 Ping
	249	+ DataCopyIn(B1 Ping, k=0) // 搬入首个 B 大包到 Ping
	250	+ SetFlag(MTE2_MTE1, 0/2) // 通知 A1/B1 Ping 数据就绪
	251	+
	252	+K循环 kBlockIdx = 0, 1, ..., kLoopCount-1:
	253	+ ┌─ a1ReadIdx = (kBlockIdx / stepKa) % 2 // 当前读 L1 A 的 Ping/Pong
	254	+ │ b1ReadIdx = (kBlockIdx / stepKb) % 2 // 当前读 L1 B 的 Ping/Pong
	255	+ │ kOffsetInChunkA = kBlockIdx % stepKa // 当前 baseK 在 A 大包内的偏移
	256	+ │ kOffsetInChunkB = kBlockIdx % stepKb // 当前 baseK 在 B 大包内的偏移
	257	+ │
	258	+ │ WaitFlag(M_MTE1, mte1DBFlag) // 等待上一轮 Compute 释放 L0
	259	+ │ if (kOffsetInChunkA == 0)
	260	+ │ WaitFlag(MTE2_MTE1, a1ReadIdx) // 等待 L1 A 大包数据就绪（仅首元素）
	261	+ │ if (kOffsetInChunkB == 0)
	262	+ │ WaitFlag(MTE2_MTE1, b1ReadIdx + 2) // 等待 L1 B 大包数据就绪（仅首元素）
	263	+ │
	264	+ │ DataLoadA(A1 → A2) // L1 → L0
	265	+ │ DataLoadB(B1 → B2) // L1 → L0
	266	+ │
	267	+ │ if (kOffsetInChunkA + 1 == stepKa)
	268	+ │ SetFlag(MTE1_MTE2, a1ReadIdx) // A 大包末元素：通知 L1 A 可写
	269	+ │ if (kOffsetInChunkB + 1 == stepKb)
	270	+ │ SetFlag(MTE1_MTE2, b1ReadIdx + 2) // B 大包末元素：通知 L1 B 可写
	271	+ │
	272	+ │ Compute(Mmad) // Cube 计算（M 指令）
	273	+ │ SetFlag(M_MTE1, mte1DBFlag) // 通知 L0 可覆盖
	274	+ │ mte1DBFlag ^= 1 // 切换 L0 Ping/Pong
	275	+ │
	276	+ │ // DataCopyIn 放在 Compute 之后，先 B 后 A
	277	+ │ if (B大包末元素 && 还有B数据):
	278	+ │ WaitFlag(MTE1_MTE2, b1WriteIdx + 2) // 等待 L1 B 缓冲区可写
	279	+ │ DataCopyInB(下一个 B 大包) // GM → L1（MTE2 指令）
	280	+ │ SetFlag(MTE2_MTE1, b1WriteIdx + 2) // 通知 L1 B 数据就绪
	281	+ │ if (A大包末元素 && 还有A数据):
	282	+ │ WaitFlag(MTE1_MTE2, a1WriteIdx) // 等待 L1 A 缓冲区可写
	283	+ │ DataCopyInA(下一个 A 大包) // GM → L1（MTE2 指令）
	284	+ │ SetFlag(MTE2_MTE1, a1WriteIdx) // 通知 L1 A 数据就绪
	285	+ └─
	286	+```
	287	+
	288	+时序示意（stepKa=8, stepKb=4）：
	289	+
	290	+```
	291	+预取: DataCopyIn(A1 Ping) + DataCopyIn(B1 Ping) ← K循环前
	292	+
	293	+k=0: WaitFlag(A1Ping, B1Ping) → DataLoad → Compute → DataCopyIn(B1Pong) → DataCopyIn(A1Pong)
	294	+k=1: DataLoad → Compute
	295	+k=2: DataLoad → Compute
	296	+k=3: DataLoad(释放B1Ping) → Compute → DataCopyIn(搬运B1Ping)
	297	+
	298	+k=4: WaitFlag(B1Pong) → DataLoad(B1Pong) → Compute
	299	+k=5: DataLoad → Compute
	300	+k=6: DataLoad → Compute
	301	+k=7: DataLoad(释放A1Ping B1Pong) → Compute → DataCopyIn(B1Pong) → DataCopyIn(A1Ping)
	302	+k=8: WaitFlag(A1Pong已就绪✓) → WaitFlag(B1Ping已就绪✓) → DataLoad → Compute
	303	+...
	304	+```
	305	+
	306	+
	307	+### 性能数据分析
	308	+
	309	+#### Atlas A2训练系列芯片性能数据
	310	+- Scenario 1：不使能 L2Cache 切分，singleCoreM=2048, singleCoreN=1536，24 核一轮全覆盖
	311	+- Scenario 2：使能 L2Cache 切分，singleCoreM=1024, singleCoreN=1536，24 核分 2 轮计算
	312	+
	313	+\| Scenario \| Task Duration(μs) \| Block Num \| aicore_time(μs) \| aic_mac_time(μs) \| aic_mac_ratio \| aic_scalar_time(μs) \| aic_scalar_ratio \| aic_mte1_time(μs) \| aic_mte1_ratio \| aic_mte2_time(μs) \| aic_mte2_ratio \| aic_fixpipe_time(μs) \| aic_fixpipe_ratio \|
	314	+\|------\|------------------\|-----------\|----------------\|-----------------\|---------------\|-------------------\|-----------------\|------------------\|----------------\|------------------\|----------------\|--------------------\|-------------------\|
	315	+\| Scenario 1 \| 4121.16 \| 24 \| 3670.7 \| 3081.664 \| 0.84 \| 337.343 \| 0.092 \| 2538.348 \| 0.692 \| 3552.248 \| 0.968 \| 160.405 \| 0.044 \|
	316	+\| Scenario 2 \| 4081.64 \| 24 \| 3636.85 \| 3082.158 \| 0.847 \| 345.139 \| 0.095 \| 2553.064 \| 0.702 \| 3487.068 \| 0.959 \| 161.812 \| 0.044 \|
	317	+
	318	+除去启动开销，已达成该芯片 84.7% 的峰值算力。
	319	+
	320	+使能 L2Cache 切分后，aic_mte2_time 从 3552.248μs 降低到 3487.068μs，降低了 1.84%。当前切分策略较简单，用户可进一步优化 L2Cache 切分策略以提高 MTE2 带宽。
	321	+
	322	+
	323	+#### Ascend 950PR芯片性能数据
	324	+
	325	+- Scenario 1：不使能 L2Cache 切分，singleCoreM=2048, singleCoreN=1024，32 核一轮全覆盖
	326	+- Scenario 2：使能 L2Cache 切分，singleCoreM=1024, singleCoreN=1024，32 核分 2 轮计算
	327	+
	328	+\| Scenario \| Task Duration(μs) \| Block Num \| aicore_time(μs) \| aic_mac_time(μs) \| aic_mac_ratio \| aic_scalar_time(μs) \| aic_scalar_ratio \| aic_mte1_time(μs) \| aic_mte1_ratio \| aic_mte2_time(μs) \| aic_mte2_ratio \| aic_fixpipe_time(μs) \| aic_fixpipe_ratio \|
	329	+\|------\|------------------\|-----------\|----------------\|-----------------\|---------------\|-------------------\|-----------------\|------------------\|----------------\|------------------\|----------------\|--------------------\|-------------------\|
	330	+\| Scenario 1 \| 2572.047 \| 32 \| 2571.44 \| 2564.813 \| 0.997 \| 144.604 \| 0.056 \| 828.001 \| 0.322 \| 1874.267 \| 0.729 \| 221.997 \| 0.086 \|
	331	+\| Scenario 2 \| 2574.492 \| 32 \| 2573.39 \| 2564.147 \| 0.996 \| 104.845 \| 0.041 \| 819.207 \| 0.318 \| 1892.742 \| 0.736 \| 223.129 \| 0.087 \|
	332	+
	333	+已达成该芯片 99.7% 的峰值算力。
	334	+
	335	+使能 L2Cache 切分后在 Ascend 950PR 芯片上没有明显效果，原因如下：L2Cache 优化的目标是缓解 MTE2 bound，但当前瓶颈为 Cube 计算而非数据搬运，因此减少 MTE2 耗时无法提升整体性能；此外 L2Cache 切分将计算分为 2 轮调度，引入了额外的 Scalar 开销和调度开销，导致 Scenario 2 的 Task Duration 略高于 Scenario 1。同时，Scenario 2 的 aic_mte2_time（1892.742μs）反而略高于 Scenario 1（1874.267μs），这是因为当样例处于 Cube bound 时，MTE2 流水线被 Cube 计算阻塞，profiler 采集的 aic_mte2_time 包含了流水线等待时间而非纯粹的数据搬运时间，L2Cache 优化虽然减少了实际数据访问延迟，但被 Cube 计算瓶颈掩盖，无法在 MTE2 指标上体现。
	336	+
	337	+
	338	+### 理论性能分析
	339	+
	340	+#### Cube 计算性能分析
	341	+
	342	+Atlas A2 训练系列芯片：样例参数 M=N=K=8192，baseM=128，baseN=256，baseK=64，该计算芯片主频为 1.85GHz，每 cycle 处理 16×16×16 次乘加运算。
	343	+
	344	+$$cube\_time = \frac{M \times N \times K}{16 \times 16 \times 16 \times core\_num \times cube\_freq} = \frac{8192 \times 8192 \times 8192}{16 \times 16 \times 16 \times 24 \times 1850} = 3022.92\mu s$$
	345	+
	346	+Cube 计算耗时误差：
	347	+
	348	+$$误差 = \frac{aic\_mac\_time - cube\_time}{cube\_time} = \frac{3082.158 - 3022.92}{3022.92} = 1.95\%$$
	349	+
	350	+
	351	+Ascend 950PR 芯片：样例参数 M=N=K=8192，baseM=256，baseN=256，baseK=64，该处理器主频为 1.65GHz，每 cycle 处理 16×16×16 次乘加运算。
	352	+
	353	+$$cube\_time = \frac{M \times N \times K}{16 \times 16 \times 16 \times core\_num \times cube\_freq} = \frac{8192 \times 8192 \times 8192}{16 \times 16 \times 16 \times 32 \times 1650} = 2542.00\mu s$$
	354	+
	355	+Cube 计算耗时误差：
	356	+
	357	+$$误差 = \frac{aic\_mac\_time - cube\_time}{cube\_time} = \frac{2564.813 - 2542.00}{2542.00} = 0.90\%$$
	358	+
	359	+
	360	+#### MTE2 带宽分析
	361	+
	362	+读入数据总量：
	363	+
	364	+Atlas A2 训练系列芯片（baseM=128，baseN=256）：
	365	+
	366	+$$读入数据总量 = \left(\frac{N}{baseN} \times M \times K + \frac{M}{baseM} \times K \times N\right) \times sizeof(half) = (32 \times 8192 \times 8192 + 64 \times 8192 \times 8192) \times 2B = 12GB$$
	367	+
	368	+Ascend 950PR 芯片（baseM=256，baseN=256）：
	369	+
	370	+$$读入数据总量 = \left(\frac{N}{baseN} \times M \times K + \frac{M}{baseM} \times K \times N\right) \times sizeof(half) = (32 \times 8192 \times 8192 + 32 \times 8192 \times 8192) \times 2B = 8GB$$
	371	+
	372	+MTE2 理论耗时：
	373	+
	374	+Atlas A2 训练系列芯片：L2Cache 峰值带宽约 5TB/s，HBM 带宽约 1.8TB/s。首次从 HBM 读入数据，后续从 L2Cache 读取。
	375	+
	376	+$$第一次从HBM读入的数据总量 = M \times K \times sizeof(half) + K \times N \times sizeof(half) = 256MB$$
	377	+
	378	+$$MTE2理论耗时 = \frac{HBM读入数据总量}{1.8TB/s} + \frac{L2Cache读入数据总量}{5TB/s}$$
	379	+
	380	+MTE2 耗时误差：
	381	+
	382	+$$MTE2耗时误差 = \frac{3487.068 - 2672.44}{2672.44} = 30.48\%$$
	383	+
	384	+当前MTE2耗时与理论值相差较大，因为实际芯片L2Cache大小为192MB，当前L2Cache切分策略较简单；另一方面当MTE2搬运场景为ND2NZ（GM数据Layout为ND，搬运到L1时需做ND→NZ格式转换）时，L2Cache带宽会降低。用户可进一步优化L2Cache切分策略以提高MTE2带宽。
	385	+
	386	+Ascend 950PR 芯片：L2Cache 峰值带宽约 5TB/s，HBM 带宽约 1.6TB/s。
	387	+
	388	+$$MTE2理论耗时 = \frac{HBM读入数据总量}{1.6TB/s} + \frac{L2Cache读入数据总量}{5TB/s}$$
	389	+
	390	+MTE2 耗时误差：
	391	+
	392	+$$MTE2耗时误差 = \frac{1874.267 - 1832.10}{1832.10} = 2.30\%$$
	393	+
	394	+相比于 Atlas A2 训练系列芯片，Ascend 950PR 芯片数据搬运更为高效，MTE2 带宽利用率更高。
	395	+
	396	+
	397	+
	398	+## 编译运行
	399	+
	400	+在本样例根目录下执行如下步骤，编译并执行样例。
	401	+
	402	+### 配置环境变量
	403	+
	404	+请根据当前环境上 CANN 开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
	405	+
	406	+- 默认路径，root 用户安装 CANN 软件包
	407	+ ```bash
	408	+ source /usr/local/Ascend/cann/set_env.sh
	409	+ ```
	410	+
	411	+- 默认路径，非 root 用户安装 CANN 软件包
	412	+ ```bash
	413	+ source $HOME/Ascend/cann/set_env.sh
	414	+ ```
	415	+
	416	+- 指定路径 install_path，安装 CANN 软件包
	417	+ ```bash
	418	+ source ${install_path}/cann/set_env.sh
	419	+ ```
	420	+
	421	+### 样例执行
	422	+
	423	+```bash
	424	+SCENARIO=1 # 选择执行场景（1为不使能L2Cache切分，2为使能L2Cache切分）
	425	+mkdir -p build && cd build; # 创建并进入 build 目录
	426	+cmake -DSCENARIO_NUM=$SCENARIO -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程（默认npu模式）
	427	+python3 ../scripts/gen_data.py # 生成测试输入数据
	428	+./demo # 执行编译生成的可执行程序
	429	+python3 ../scripts/verify_result.py output/output.bin output/golden.bin
	430	+```
	431	+
	432	+使用 NPU仿真模式时，添加 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
	433	+
	434	+示例：
	435	+```bash
	436	+cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU 仿真模式
	437	+```
	438	+
	439	+> 注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
	440	+
	441	+### 编译选项说明
	442	+
	443	+\| 参数 \| 说明 \| 可选值 \| 默认值 \|
	444	+\|------\|------\|---------\|--------\|
	445	+\| `SCENARIO_NUM` \| `1` / `2` \| 1: 不使能L2Cache切分；2: 使能L2Cache切分 \| `1` \|
	446	+\| `CMAKE_ASC_RUN_MODE` \| 运行模式 \| `npu`、`sim` \| `npu` \|
	447	+\| `CMAKE_ASC_ARCHITECTURES` \| NPU 硬件架构 \| `dav-2201`、`dav-3510` \| `dav-2201` \|
	cellur_z4月20日编译选项是不是还有一个SCENARIO_NUM
	448	+
	449	+执行结果如下，说明精度对比成功。
	450	+```bash
	451	+test pass!
	452	+```
	453	+
	454	+### 性能分析
	455	+
	456	+使用 `msprof` 工具获取详细性能数据：
	457	+
	458	+```bash
	459	+msprof ./demo # 分析样例性能
	460	+```
	461	+
	462	+当前目录下会生成 PROF_ 前缀的文件夹，`mindstudio_profiler_output` 目录保存 Host 和各个 Device 的性能数据汇总，性能数据分析推荐查看该目录下文件：
	463	+```bash
	464	+PROF_xxxx_XXXXXX
	465	+├── device_{id}
	466	+└── host
	467	+└── mindstudio_profiler_log
	468	+└── mindstudio_profiler_output # 保存 Host 和各个 Device 的性能数据汇总
	469	+ ├── msprof_*.json
	470	+ ├── xx_*.csv
	471	+ └── README.txt
	472	+```

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_basic_API_high_performance/data_utils.h+94-0

@@ -0,0 +1,94 @@
	1	+/**
	2	+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+* CANN Open Software License Agreement Version 2.0 (the "License").
	5	+* Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+* See LICENSE in the root of the software repository for the full text of the License.
	9	+*/
	10	+
	11	+
	12	+/* !
	13	+ * \file data_utils.h
	14	+ * \brief Utility functions for reading and writing binary files
	15	+ */
	16	+
	17	+#ifndef DATA_UTILS_H
	18	+#define DATA_UTILS_H
	19	+#include <fcntl.h>
	20	+#include <sys/stat.h>
	21	+#include <unistd.h>
	22	+#include <fstream>
	23	+
	24	+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
	25	+
	26	+bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize)
	27	+{
	28	+ struct stat sBuf;
	29	+ int fileStatus = stat(filePath.data(), &sBuf);
	30	+ if (fileStatus == -1) {
	31	+ ERROR_LOG("failed to get file");
	32	+ return false;
	33	+ }
	34	+ if (S_ISREG(sBuf.st_mode) == 0) {
	35	+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
	36	+ return false;
	37	+ }
	38	+
	39	+ std::ifstream file;
	40	+ file.open(filePath, std::ios::binary);
	41	+ if (!file.is_open()) {
	42	+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
	43	+ return false;
	44	+ }
	45	+
	46	+ std::filebuf *buf = file.rdbuf();
	47	+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
	48	+ if (size == 0) {
	49	+ ERROR_LOG("file size is 0");
	50	+ file.close();
	51	+ return false;
	52	+ }
	53	+ if (size > bufferSize) {
	54	+ ERROR_LOG("file size is larger than buffer size");
	55	+ file.close();
	56	+ return false;
	57	+ }
	58	+ buf->pubseekpos(0, std::ios::in);
	59	+ buf->sgetn(static_cast<char *>(buffer), size);
	60	+ fileSize = size;
	61	+ file.close();
	62	+ return true;
	63	+}
	64	+
	65	+/**
	66	+ * @brief Write data to file
	67	+ * @param [in] filePath: file path
	68	+ * @param [in] buffer: data to write to file
	69	+ * @param [in] size: size to write
	70	+ * @return write result
	71	+ */
	72	+bool WriteFile(const std::string &filePath, const void *buffer, size_t size)
	73	+{
	74	+ if (buffer == nullptr) {
	75	+ ERROR_LOG("Write file failed. buffer is nullptr");
	76	+ return false;
	77	+ }
	78	+
	79	+ int fd = open(filePath.c_str(), O_RDWR \| O_CREAT \| O_TRUNC, S_IRUSR \| S_IWRITE);
	80	+ if (fd < 0) {
	81	+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
	82	+ return false;
	83	+ }
	84	+
	85	+ size_t writeSize = write(fd, buffer, size);
	86	+ (void)close(fd);
	87	+ if (writeSize != size) {
	88	+ ERROR_LOG("Write file Failed.");
	89	+ return false;
	90	+ }
	91	+
	92	+ return true;
	93	+}
	94	+#endif // DATA_UTILS_H

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_basic_API_high_performance/mmad.asc+643-0

@@ -0,0 +1,643 @@
	1	+/**
	2	+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+* CANN Open Software License Agreement Version 2.0 (the "License").
	5	+* Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+* See LICENSE in the root of the software repository for the full text of the License.
	9	+*/
	10	+
	11	+
	12	+/* !
	13	+ * \file mmad.asc
	14	+ * \brief matmul_basic_API_high_performance
	15	+ */
	16	+
	17	+#include "acl/acl.h"
	18	+#include "kernel_operator.h"
	19	+#include "data_utils.h"
	20	+
	21	+#ifdef ASCENDC_CPU_DEBUG
	22	+#include "cpu_debug_launch.h"
	23	+#endif
	24	+
	25	+// half type, cube block: [16, 16]
	26	+constexpr uint32_t CUBE_BLOCK = 16;
	27	+constexpr uint32_t CUBE_BLOCK_SIZE = 16 * 16;
	28	+constexpr uint32_t L0_PINGPONG_BYTES = 32 * 1024;
	29	+constexpr uint32_t L1_PINGPONG_BYTES = 256 * 1024;
	30	+constexpr bool IS_B_TRANSPOSE = true;
	31	+constexpr uint32_t scenarioNum = SCENARIO_NUM;
	32	+
	33	+
	34	+template <uint32_t M, uint32_t K, uint32_t N, uint32_t baseM, uint32_t baseK, uint32_t baseN,
	35	+ uint32_t singleCoreM, uint32_t singleCoreK, uint32_t singleCoreN,
	36	+ uint32_t stepKa, uint32_t stepKb>
	37	+class KernelMmad {
	38	+
	39	+public:
	40	+ __aicore__ inline KernelMmad()
	41	+ {
	42	+
	43	+ }
	44	+ __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR c)
	45	+ {
	46	+ aGMOri.SetGlobalBuffer((__gm__ half *)a);
	47	+ bGMOri.SetGlobalBuffer((__gm__ half *)b);
	48	+ cGMOri.SetGlobalBuffer((__gm__ half *)c);
	49	+ }
	50	+
	51	+ __aicore__ inline void Process()
	52	+ {
	53	+ InitComputeParams();
	54	+
	55	+ // ============================================================
	56	+ // 1. Buffer 分配：L1 / L0 双缓冲 Ping-Pong 布局
	57	+ // ============================================================
	58	+ // L1: A1 Ping/Pong (0~256KB), B1 Ping/Pong (256~512KB)
	59	+ // L0: A2/B2 Ping/Pong (0~32KB / 32KB~64KB)
	60	+ uint32_t a1PingpongSize = baseM * baseK * stepKa;
	61	+ uint32_t b1PingpongSize = baseK * baseN * stepKb;
	62	+ uint32_t a2PingpongSize = baseM * baseK;
	63	+ uint32_t b2PingpongSize = baseK * baseN;
	64	+
	65	+ // A1: L1 Ping/Pong
	66	+ AscendC::LocalTensor<half> a1LocalPing(AscendC::TPosition::A1, 0, a1PingpongSize);
	67	+ AscendC::LocalTensor<half> a1LocalPong(AscendC::TPosition::A1, a1PingpongSize * sizeof(half), a1PingpongSize);
	68	+ // A2: L0 Ping/Pong
	69	+ AscendC::LocalTensor<half> a2LocalPing(AscendC::TPosition::A2, 0, a2PingpongSize);
	70	+ AscendC::LocalTensor<half> a2LocalPong(AscendC::TPosition::A2, L0_PINGPONG_BYTES, a2PingpongSize);
	71	+
	72	+ // B1: L1 Ping/Pong
	73	+ AscendC::LocalTensor<half> b1LocalPing(AscendC::TPosition::B1, L1_PINGPONG_BYTES, b1PingpongSize);
	74	+ AscendC::LocalTensor<half> b1LocalPong(AscendC::TPosition::B1, L1_PINGPONG_BYTES + b1PingpongSize * sizeof(half), b1PingpongSize);
	75	+ // B2: L0 Ping/Pong
	76	+ AscendC::LocalTensor<half> b2LocalPing(AscendC::TPosition::B2, 0, b2PingpongSize);
	77	+ AscendC::LocalTensor<half> b2LocalPong(AscendC::TPosition::B2, L0_PINGPONG_BYTES, b2PingpongSize);
	78	+ // CO1: L0C
	79	+ AscendC::LocalTensor<float> cLocal(AscendC::TPosition::CO1, 0, baseM * baseN);
	80	+
	81	+ ProcessLoop(a1LocalPing, a1LocalPong, a2LocalPing, a2LocalPong,
	82	+ b1LocalPing, b1LocalPong, b2LocalPing, b2LocalPong, cLocal);
	83	+ }
	84	+
	85	+ // ============================================================
	86	+ // L2 Cache 优化：双重外层循环（M 维度 + N 维度分批）
	87	+ // ============================================================
	88	+ __aicore__ inline void ProcessL2Cache()
	89	+ {
	90	+
	91	+ // ---- Step 1: Buffer 分配：L1 / L0 双缓冲 Ping-Pong 布局
	92	+ uint32_t a1PingpongSize = baseM * baseK * stepKa;
	93	+ uint32_t b1PingpongSize = baseK * baseN * stepKb;
	94	+ uint32_t a2PingpongSize = baseM * baseK;
	95	+ uint32_t b2PingpongSize = baseK * baseN;
	96	+
	97	+ AscendC::LocalTensor<half> a1LocalPing(AscendC::TPosition::A1, 0, a1PingpongSize);
	98	+ AscendC::LocalTensor<half> a1LocalPong(AscendC::TPosition::A1, a1PingpongSize * sizeof(half), a1PingpongSize);
	99	+ AscendC::LocalTensor<half> a2LocalPing(AscendC::TPosition::A2, 0, a2PingpongSize);
	100	+ AscendC::LocalTensor<half> a2LocalPong(AscendC::TPosition::A2, L0_PINGPONG_BYTES, a2PingpongSize);
	101	+
	102	+ AscendC::LocalTensor<half> b1LocalPing(AscendC::TPosition::B1, L1_PINGPONG_BYTES, b1PingpongSize);
	103	+ AscendC::LocalTensor<half> b1LocalPong(AscendC::TPosition::B1, L1_PINGPONG_BYTES + b1PingpongSize * sizeof(half), b1PingpongSize);
	104	+ AscendC::LocalTensor<half> b2LocalPing(AscendC::TPosition::B2, 0, b2PingpongSize);
	105	+ AscendC::LocalTensor<half> b2LocalPong(AscendC::TPosition::B2, L0_PINGPONG_BYTES, b2PingpongSize);
	106	+ AscendC::LocalTensor<float> cLocal(AscendC::TPosition::CO1, 0, baseM * baseN);
	107	+
	108	+ // ---- Step 2: 计算分批参数 ----
	109	+ // mIterTotal/nIterTotal: M/N 方向总共需要多少个 singleCore 大小的子块
	110	+ constexpr uint32_t mIterTotal = AscendC::DivCeil(M, singleCoreM);
	111	+ constexpr uint32_t nIterTotal = AscendC::DivCeil(N, singleCoreN);
	112	+ // 每轮 numBlocks 核覆盖的 M 份数
	113	+ constexpr uint32_t mIterPerRound = AscendC::DivCeil(M, singleCoreM * 2); // 2 轮 M
	114	+ // 外层循环轮数
	115	+ constexpr uint32_t outerMLoopCount = AscendC::DivCeil(mIterTotal, mIterPerRound);
	116	+
	117	+ // ---- Step 3: outerMIdx/outerNIdx 外层循环 ----
	118	+ for (uint32_t outerMIdx = 0; outerMIdx < outerMLoopCount; outerMIdx++) {
	119	+ // 根据核索引 blockIdx 计算当前核在本轮负责的子块坐标
	120	+ // mIterIdx = blockIdx 在 M 方向的偏移 + 本轮的 M 起始偏移
	121	+ // nIterIdx = blockIdx 在 N 方向的偏移
	122	+ uint32_t mIterIdx = AscendC::GetBlockIdx() % mIterPerRound + outerMIdx * mIterPerRound;
	123	+ uint32_t nIterIdx = AscendC::GetBlockIdx() / mIterPerRound;
	124	+
	125	+ // 边界保护：当 M/N 不能被整除时，最后一轮部分核会超出子块范围，直接跳过
	126	+ if (mIterIdx >= mIterTotal \|\| nIterIdx >= nIterTotal) continue;
	127	+
	128	+ // 根据子块坐标计算 GM 偏移、循环次数、尾块参数等
	129	+ InitComputeParamsL2Cache(mIterIdx, nIterIdx);
	130	+
	131	+ // 执行当前子块的完整 N/M/K 循环计算
	132	+ ProcessLoop(a1LocalPing, a1LocalPong, a2LocalPing, a2LocalPong,
	133	+ b1LocalPing, b1LocalPong, b2LocalPing, b2LocalPong, cLocal);
	134	+
	135	+ // 每轮外层循环后重置双缓冲 flag
	136	+ mte1DBFlag = 0;
	137	+ }
	138	+ }
	139	+
	140	+private:
	141	+ // N/M/K 主循环：执行当前核的完整计算
	142	+ __aicore__ inline void ProcessLoop(
	143	+ AscendC::LocalTensor<half> &a1LocalPing, AscendC::LocalTensor<half> &a1LocalPong,
	144	+ AscendC::LocalTensor<half> &a2LocalPing, AscendC::LocalTensor<half> &a2LocalPong,
	145	+ AscendC::LocalTensor<half> &b1LocalPing, AscendC::LocalTensor<half> &b1LocalPong,
	146	+ AscendC::LocalTensor<half> &b2LocalPing, AscendC::LocalTensor<half> &b2LocalPong,
	147	+ AscendC::LocalTensor<float> &cLocal)
	148	+ {
	149	+ // ============================================================
	150	+ // 1. DataCopyIn 进度跟踪变量
	151	+ // ============================================================
	152	+ uint32_t a1NextKChunkIdx = 0; // A 下一次 DataCopyIn 要搬的 K 方向大包索引
	153	+ uint32_t b1NextKChunkIdx = 0; // B 下一次 DataCopyIn 要搬的 K 方向大包索引
	154	+ uint8_t a1CopyInIdx = 0; // A 下一次 DataCopyIn 写入的 L1 buffer 索引 (0=Ping, 1=Pong)
	155	+ uint8_t b1CopyInIdx = 0; // B 下一次 DataCopyIn 写入的 L1 buffer 索引 (0=Ping, 1=Pong)
	156	+
	157	+ // ============================================================
	158	+ // 2. 同步标志初始化：反向同步需预置，防止首次 WaitFlag 死锁
	159	+ // ============================================================
	160	+ // MTE1_MTE2 (反向): DataCopyIn 可覆盖通知
	161	+ // flag 0: A1 Ping 可写 \| flag 1: A1 Pong 可写
	162	+ // flag 2: B1 Ping 可写 \| flag 3: B1 Pong 可写
	163	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(0);
	164	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(1);
	165	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(2);
	166	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(3);
	167	+ // M_MTE1 (反向): L0 双缓冲，DataLoad 可覆盖通知
	168	+ AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(0);
	169	+ AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(1);
	170	+
	171	+ // ============================================================
	172	+ // 3. N/M 外层循环 + K 主循环
	173	+ // ============================================================
	174	+ for (uint32_t nBlockIdx = 0; nBlockIdx < nLoopCount; nBlockIdx++) {
	175	+ for (uint32_t mBlockIdx = 0; mBlockIdx < mLoopCount; mBlockIdx++) {
	176	+ // 每个 (mBlockIdx, nBlockIdx) 子块重置 DataCopyIn 进度
	177	+ a1NextKChunkIdx = 0;
	178	+ b1NextKChunkIdx = 0;
	179	+ a1CopyInIdx = 0;
	180	+ b1CopyInIdx = 0;
	181	+
	182	+ // ---- 搬入 A1 Ping 和 B1 Ping 的首个大包 ----
	183	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(0); // A1 Ping 可写
	184	+ DataCopyInA(a1LocalPing, a1NextKChunkIdx, mBlockIdx);
	185	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(0); // A1 Ping 数据就绪
	186	+ a1NextKChunkIdx += stepKa;
	187	+ a1CopyInIdx ^= 1; // 下次写 Pong
	188	+
	189	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(2); // B1 Ping 可写
	190	+ DataCopyInB(b1LocalPing, b1NextKChunkIdx, nBlockIdx);
	191	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(2); // B1 Ping 数据就绪
	192	+ b1NextKChunkIdx += stepKb;
	193	+ b1CopyInIdx ^= 1; // 下次写 Pong
	194	+
	195	+ // ---- K 方向主循环 ----
	196	+ for (uint32_t kBlockIdx = 0; kBlockIdx < kLoopCount; kBlockIdx++) {
	197	+ // 确定当前 kBlockIdx 对应的 L1 读取缓冲区（Ping/Pong）
	198	+ uint32_t a1ReadIdx = (kBlockIdx / stepKa) % 2;
	199	+ uint32_t b1ReadIdx = (kBlockIdx / stepKb) % 2;
	200	+ uint32_t kOffsetInChunkA = kBlockIdx % stepKa;
	201	+ uint32_t kOffsetInChunkB = kBlockIdx % stepKb;
	202	+
	203	+ AscendC::LocalTensor<half> a1ReadBuf = (a1ReadIdx == 0) ? a1LocalPing : a1LocalPong;
	204	+ AscendC::LocalTensor<half> b1ReadBuf = (b1ReadIdx == 0) ? b1LocalPing : b1LocalPong;
	205	+
	206	+ // L0 双缓冲选择
	207	+ AscendC::LocalTensor<half> a2Local = (mte1DBFlag == 0) ? a2LocalPing : a2LocalPong;
	208	+ AscendC::LocalTensor<half> b2Local = (mte1DBFlag == 0) ? b2LocalPing : b2LocalPong;
	209	+
	210	+ // ---- 反向同步：等待上一轮 Compute 释放 L0 缓冲区 ----
	211	+ AscendC::WaitFlag<AscendC::HardEvent::M_MTE1>(mte1DBFlag);
	212	+
	213	+ // ---- 正向同步：等待 L1 大包数据就绪（仅大包中首元素需等待） ----
	214	+ if (kOffsetInChunkA == 0) {
	215	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(a1ReadIdx);
	216	+ }
	217	+ if (kOffsetInChunkB == 0) {
	218	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(b1ReadIdx + 2);
	219	+ }
	220	+
	221	+ // ---- DataLoad: L1 → L0 ----
	222	+ DataLoadA(a1ReadBuf, a2Local, mBlockIdx, kOffsetInChunkA);
	223	+ DataLoadB(b1ReadBuf, b2Local, nBlockIdx, kOffsetInChunkB);
	224	+
	225	+ // ---- 反向同步：当前 L1 大包消费完毕，通知 DataCopyIn 可覆盖 ----
	226	+ if ((kOffsetInChunkA + 1) == stepKa) {
	227	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(a1ReadIdx);
	228	+ }
	229	+ if ((kOffsetInChunkB + 1) == stepKb) {
	230	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_MTE2>(b1ReadIdx + 2);
	231	+ }
	232	+
	233	+ // ---- Compute: Mmad 矩阵乘累加 ----
	234	+ Compute(cLocal, a2Local, b2Local, kBlockIdx, mBlockIdx, nBlockIdx);
	235	+
	236	+ // ---- 搬入下一个 L1 大包，使 DataCopyIn 与 Compute 流水重叠 ----
	237	+ // 触发条件：
	238	+ // (1) kBlockIdx == 0：首个 K 块计算时，A1/B1 Pong 尚未使用，可直接搬入
	239	+ // (2) 当前 L1 大包最后一个 baseK 被消费完，该缓冲区可被覆盖
	240	+ if (((kBlockIdx == 0) \|\| ((kOffsetInChunkB + 1) == stepKb)) && b1NextKChunkIdx < kLoopCount) {
	241	+ AscendC::LocalTensor<half> b1WriteBuf = (b1CopyInIdx == 0) ? b1LocalPing : b1LocalPong;
	242	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(b1CopyInIdx + 2);
	243	+ DataCopyInB(b1WriteBuf, b1NextKChunkIdx, nBlockIdx);
	244	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(b1CopyInIdx + 2);
	245	+ b1NextKChunkIdx += stepKb;
	246	+ b1CopyInIdx ^= 1;
	247	+ }
	248	+ if (((kBlockIdx == 0) \|\| ((kOffsetInChunkA + 1) == stepKa)) && a1NextKChunkIdx < kLoopCount) {
	249	+ AscendC::LocalTensor<half> a1WriteBuf = (a1CopyInIdx == 0) ? a1LocalPing : a1LocalPong;
	250	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(a1CopyInIdx);
	251	+ DataCopyInA(a1WriteBuf, a1NextKChunkIdx, mBlockIdx);
	252	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(a1CopyInIdx);
	253	+ a1NextKChunkIdx += stepKa;
	254	+ a1CopyInIdx ^= 1;
	255	+ }
	256	+ }
	257	+ // ---- CopyOut: CO1 → GM ----
	258	+ CopyOut(cLocal, mBlockIdx, nBlockIdx);
	259	+ }
	260	+ }
	261	+
	262	+ // ============================================================
	263	+ // 4. 等待所有同步完成
	264	+ // ============================================================
	265	+ AscendC::WaitFlag<AscendC::HardEvent::M_MTE1>(0);
	266	+ AscendC::WaitFlag<AscendC::HardEvent::M_MTE1>(1);
	267	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(0);
	268	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(1);
	269	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(2);
	270	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_MTE2>(3);
	271	+ }
	272	+
	273	+ __aicore__ inline void InitComputeParams()
	274	+ {
	275	+ // ---- 1. 计算当前核在 M/N 方向的迭代索引，确定 GM 起始偏移 ----
	276	+ constexpr uint32_t mIter = AscendC::DivCeil(M, singleCoreM);
	277	+ uint32_t mIterIdx = AscendC::GetBlockIdx() % mIter;
	278	+ uint32_t nIterIdx = AscendC::GetBlockIdx() / mIter;
	279	+
	280	+ uint64_t gmOffsetA = mIterIdx * singleCoreM * K;
	281	+ uint64_t gmOffsetB = IS_B_TRANSPOSE ? nIterIdx * K * singleCoreN : nIterIdx * singleCoreN;
	282	+ uint64_t gmOffsetC = mIterIdx * singleCoreM * N + nIterIdx * singleCoreN;
	283	+ aGM = aGMOri[gmOffsetA];
	284	+ bGM = bGMOri[gmOffsetB];
	285	+ cGM = cGMOri[gmOffsetC];
	286	+
	287	+ // ---- 2. 计算当前核实际的 M/N 维度大小（最后一块可能不满 singleCore） ----
	288	+ actualSingleCoreM = M - mIterIdx * singleCoreM;
	289	+ actualSingleCoreM = actualSingleCoreM < singleCoreM ? actualSingleCoreM : singleCoreM;
	290	+ actualSingleCoreN = N - nIterIdx * singleCoreN;
	291	+ actualSingleCoreN = actualSingleCoreN < singleCoreN ? actualSingleCoreN : singleCoreN;
	292	+
	293	+ // ---- 3. 计算 K/M/N 三个维度的循环次数 ----
	294	+ kLoopCount = AscendC::DivCeil(singleCoreK, baseK);
	295	+ mLoopCount = AscendC::DivCeil(actualSingleCoreM, baseM);
	296	+ nLoopCount = AscendC::DivCeil(actualSingleCoreN, baseN);
	297	+
	298	+ // ---- 4. 计算 N 方向分块参数：完整 baseN 块数 + 尾块大小及对齐 ----
	299	+ baseNCount = actualSingleCoreN / baseN;
	300	+ tailN = actualSingleCoreN % baseN;
	301	+ tailNAlign = AscendC::DivCeil(tailN, CUBE_BLOCK) * CUBE_BLOCK;
	302	+
	303	+ // ---- 5. 计算 M 方向分块参数：完整 baseM 块数 + 尾块大小及对齐 ----
	304	+ baseMCount = actualSingleCoreM / baseM;
	305	+ tailM = actualSingleCoreM % baseM;
	306	+ tailMAlign = AscendC::DivCeil(tailM, CUBE_BLOCK) * CUBE_BLOCK;
	307	+ }
	308	+
	309	+ // 根据子块索引设置当前子块的全部计算参数
	310	+ // mIterIdx: 当前核在 M 方向的子块索引
	311	+ // nIterIdx: 当前核在 N 方向的子块索引
	312	+ __aicore__ inline void InitComputeParamsL2Cache(uint32_t mIterIdx, uint32_t nIterIdx)
	313	+ {
	314	+ // ---- 计算当前子块在 GM 中的起始偏移 ----
	315	+ // aGM/bGM/cGM 指向当前子块起始位置，后续 DataCopyIn/CopyOut 基于该偏移寻址
	316	+ uint64_t gmOffsetA = mIterIdx * singleCoreM * K;
	317	+ uint64_t gmOffsetB = IS_B_TRANSPOSE ? nIterIdx * K * singleCoreN : nIterIdx * singleCoreN;
	318	+ uint64_t gmOffsetC = mIterIdx * singleCoreM * N + nIterIdx * singleCoreN;
	319	+ aGM = aGMOri[gmOffsetA];
	320	+ bGM = bGMOri[gmOffsetB];
	321	+ cGM = cGMOri[gmOffsetC];
	322	+
	323	+ actualSingleCoreM = M - mIterIdx * singleCoreM;
	324	+ actualSingleCoreM = actualSingleCoreM < singleCoreM ? actualSingleCoreM : singleCoreM;
	325	+ actualSingleCoreN = N - nIterIdx * singleCoreN;
	326	+ actualSingleCoreN = actualSingleCoreN < singleCoreN ? actualSingleCoreN : singleCoreN;
	327	+
	328	+ kLoopCount = AscendC::DivCeil(singleCoreK, baseK);
	329	+ mLoopCount = AscendC::DivCeil(actualSingleCoreM, baseM);
	330	+ nLoopCount = AscendC::DivCeil(actualSingleCoreN, baseN);
	331	+ baseNCount = actualSingleCoreN / baseN;
	332	+ tailN = actualSingleCoreN % baseN;
	333	+ tailNAlign = AscendC::DivCeil(tailN, CUBE_BLOCK) * CUBE_BLOCK;
	334	+
	335	+ baseMCount = actualSingleCoreM / baseM;
	336	+ tailM = actualSingleCoreM % baseM;
	337	+ tailMAlign = AscendC::DivCeil(tailM, CUBE_BLOCK) * CUBE_BLOCK;
	338	+ }
	339	+
	340	+ // GM → A1: 将 A 矩阵的一个大包 (stepKa 个 baseK) 搬入 L1
	341	+ // kChunkIdx: K 方向大包索引，用于计算 GM 读取偏移
	342	+ // mBlockIdx: M 方向 baseM 块索引，用于计算 GM 读取偏移及判断尾块
	343	+ __aicore__ inline void DataCopyInA(AscendC::LocalTensor<half> a1Local, uint32_t kChunkIdx, uint32_t mBlockIdx)
	344	+ {
	345	+ // 尾块使用 tailM，否则使用 baseM
	346	+ uint32_t curM = (mBlockIdx != baseMCount) ? baseM : tailM;
	347	+ AscendC::Nd2NzParams nd2nzParams;
	348	+ nd2nzParams.ndNum = 1;
	349	+ nd2nzParams.nValue = curM;
	350	+ nd2nzParams.dValue = baseK * stepKa;
	351	+ nd2nzParams.srcNdMatrixStride = 0;
	352	+ nd2nzParams.srcDValue = K;
	353	+ nd2nzParams.dstNzC0Stride = baseM;
	354	+ nd2nzParams.dstNzNStride = 1;
	355	+ nd2nzParams.dstNzMatrixStride = 0;
	356	+ AscendC::DataCopy(a1Local, aGM[kChunkIdx * baseK + mBlockIdx * K * baseM], nd2nzParams);
	357	+ }
	358	+
	359	+ // GM → B1: 将 B 矩阵的一个大包 (stepKb 个 baseK) 搬入 L1
	360	+ // kChunkIdx: K 方向大包索引，用于计算 GM 读取偏移
	361	+ // nBlockIdx: N 方向 baseN 块索引，用于计算 GM 读取偏移及判断尾块
	362	+ __aicore__ inline void DataCopyInB(AscendC::LocalTensor<half> b1Local, uint32_t kChunkIdx, uint32_t nBlockIdx)
	363	+ {
	364	+ // 尾块使用 tailN，否则使用 baseN
	365	+ uint32_t curN = (nBlockIdx != baseNCount) ? baseN : tailN;
	366	+ AscendC::Nd2NzParams nd2nzParams;
	367	+ if constexpr (!IS_B_TRANSPOSE) {
	368	+ // B 非转置
	369	+ nd2nzParams.ndNum = 1;
	370	+ nd2nzParams.nValue = baseK * stepKb;
	371	+ nd2nzParams.dValue = curN;
	372	+ nd2nzParams.srcNdMatrixStride = 0;
	373	+ nd2nzParams.srcDValue = N;
	374	+ nd2nzParams.dstNzC0Stride = baseK * stepKb;
	375	+ nd2nzParams.dstNzNStride = 1;
	376	+ nd2nzParams.dstNzMatrixStride = 0;
	377	+ AscendC::DataCopy(b1Local, bGM[kChunkIdx * baseK * N + nBlockIdx * baseN], nd2nzParams);
	378	+ } else {
	379	+ // B 转置
	380	+ nd2nzParams.ndNum = 1;
	381	+ nd2nzParams.nValue = curN;
	382	+ nd2nzParams.dValue = baseK * stepKb;
	383	+ nd2nzParams.srcNdMatrixStride = 0;
	384	+ nd2nzParams.srcDValue = K;
	385	+ nd2nzParams.dstNzC0Stride = baseN;
	386	+ nd2nzParams.dstNzNStride = 1;
	387	+ nd2nzParams.dstNzMatrixStride = 0;
	388	+ AscendC::DataCopy(b1Local, bGM[kChunkIdx * baseK + nBlockIdx * baseN * K], nd2nzParams);
	389	+ }
	390	+ }
	391	+
	392	+ // A1 → A2: 将 L1 中的一个 baseM * baseK 切片搬入 L0
	393	+ // mBlockIdx: M 方向 baseM 块索引，用于判断尾块对齐
	394	+ // kOffsetInChunkA: 当前 base块在大包内的偏移序号 (0 ~ stepKa-1)
	395	+ __aicore__ inline void DataLoadA(AscendC::LocalTensor<half> a1Local, AscendC::LocalTensor<half> a2Local,
	396	+ uint32_t mBlockIdx, uint32_t kOffsetInChunkA)
	397	+ {
	398	+ uint32_t srcAddr = kOffsetInChunkA * baseK * baseM;
	399	+ // 尾块使用对齐后的 tailMAlign，保证 LoadData 地址对齐
	400	+ uint32_t curMAlign = (mBlockIdx != baseMCount) ? baseM : tailMAlign;
	401	+#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 2201 \|\| __NPU_ARCH__ == 2202)
	402	+ AscendC::LoadData3DParamsV2<half> loadDataParams;
	403	+ loadDataParams.l1H = 1;
	404	+ loadDataParams.l1W = baseM;
	405	+ loadDataParams.channelSize = baseK;
	406	+ loadDataParams.kExtension = baseK;
	407	+ loadDataParams.mExtension = curMAlign;
	408	+ loadDataParams.mStartPt = 0;
	409	+ loadDataParams.kStartPt = 0;
	410	+ AscendC::LoadData(a2Local, a1Local[srcAddr], loadDataParams);
	411	+#elif defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
	412	+ AscendC::LoadData2DParamsV2 loadDataParams;
	413	+ loadDataParams.mStartPosition = 0;
	414	+ loadDataParams.kStartPosition = 0;
	415	+ loadDataParams.mStep = AscendC::DivCeil(curMAlign, CUBE_BLOCK);
	416	+ loadDataParams.kStep = AscendC::DivCeil(baseK, CUBE_BLOCK);
	417	+ loadDataParams.srcStride = AscendC::DivCeil(baseM, CUBE_BLOCK);
	418	+ loadDataParams.dstStride = AscendC::DivCeil(curMAlign, CUBE_BLOCK);
	419	+ loadDataParams.sid = 0;
	420	+ loadDataParams.ifTranspose = false;
	421	+ AscendC::LoadData(a2Local, a1Local[srcAddr], loadDataParams);
	422	+#endif
	423	+ }
	424	+
	425	+ // B1 → B2: 将 L1 中的一个 baseK * baseN 切片搬入 L0
	426	+ // nBlockIdx: N 方向 baseN 块索引，用于判断尾块对齐
	427	+ // kOffsetInChunkB: 当前 base块在大包内的偏移序号 (0 ~ stepKb-1)
	428	+ __aicore__ inline void DataLoadB(AscendC::LocalTensor<half> b1Local, AscendC::LocalTensor<half> b2Local,
	429	+ uint32_t nBlockIdx, uint32_t kOffsetInChunkB)
	430	+ {
	431	+ uint32_t srcAddr = kOffsetInChunkB * baseK * (IS_B_TRANSPOSE ? baseN : CUBE_BLOCK);
	432	+ // 尾块使用对齐后的 tailNAlign，保证 LoadData 地址对齐
	433	+ uint32_t curNAlign = (nBlockIdx != baseNCount) ? baseN : tailNAlign;
	434	+#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 2201 \|\| __NPU_ARCH__ == 2202)
	435	+ if constexpr (!IS_B_TRANSPOSE) {
	436	+ // B 非转置: 使用 LoadData3DV2 完成 [K, N] → [N, K] 转置搬运
	437	+ AscendC::LoadData3DParamsV2<half> loadDataParams;
	438	+ loadDataParams.l1H = 1;
	439	+ loadDataParams.l1W = baseK * stepKb;
	440	+ loadDataParams.channelSize = baseN;
	441	+ loadDataParams.kExtension = curNAlign;
	442	+ loadDataParams.mExtension = baseK;
	443	+ loadDataParams.mStartPt = kOffsetInChunkB * baseK;
	444	+ loadDataParams.kStartPt = 0;
	445	+ loadDataParams.strideW = 1;
	446	+ loadDataParams.strideH = 1;
	447	+ loadDataParams.filterW = 1;
	448	+ loadDataParams.filterH = 1;
	449	+ loadDataParams.dilationFilterW = 1;
	450	+ loadDataParams.dilationFilterH = 1;
	451	+ loadDataParams.filterSizeW = false;
	452	+ loadDataParams.filterSizeH = false;
	453	+ loadDataParams.enTranspose = true;
	454	+ loadDataParams.fMatrixCtrl = false;
	455	+ AscendC::LoadData(b2Local, b1Local, loadDataParams);
	456	+ } else {
	457	+ // B 转置: 按 CUBE_BLOCK 粒度逐块搬入 L0，无需转置
	458	+ AscendC::LoadData2DParams loadDataParams;
	459	+ uint32_t dstOffset = curNAlign * CUBE_BLOCK;
	460	+ uint32_t srcOffset = baseN * CUBE_BLOCK;
	461	+ loadDataParams.repeatTimes = AscendC::DivCeil(curNAlign, CUBE_BLOCK);
	462	+ loadDataParams.srcStride = 1;
	463	+ loadDataParams.dstGap = 0;
	464	+ loadDataParams.ifTranspose = false;
	465	+ for (int i = 0; i < AscendC::DivCeil(baseK, CUBE_BLOCK); ++i) {
	466	+ AscendC::LoadData(b2Local[i * dstOffset], b1Local[srcAddr + i * srcOffset], loadDataParams);
	467	+ }
	468	+ }
	469	+#elif defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
	470	+ if constexpr (!IS_B_TRANSPOSE) {
	471	+ // B 非转置: 使用 LoadData2D V2 完成 [K, N] → [N, K] 转置搬运
	472	+ AscendC::LoadData2DParamsV2 loadDataParams;
	473	+ loadDataParams.mStartPosition = 0;
	474	+ loadDataParams.kStartPosition = 0;
	475	+ loadDataParams.mStep = AscendC::DivCeil(baseK, CUBE_BLOCK);
	476	+ loadDataParams.kStep = AscendC::DivCeil(curNAlign * sizeof(half), 32);
	477	+ loadDataParams.srcStride = AscendC::DivCeil(baseK * stepKb, CUBE_BLOCK);
	478	+ loadDataParams.dstStride = AscendC::DivCeil(curNAlign, CUBE_BLOCK);
	479	+ loadDataParams.ifTranspose = true;
	480	+ AscendC::LoadData(b2Local, b1Local[srcAddr], loadDataParams);
	481	+ } else {
	482	+ // B 转置: 无需转置，直接按块搬运
	483	+ AscendC::LoadData2DParamsV2 loadDataParams;
	484	+ loadDataParams.mStartPosition = 0;
	485	+ loadDataParams.kStartPosition = 0;
	486	+ loadDataParams.mStep = AscendC::DivCeil(curNAlign, CUBE_BLOCK);
	487	+ loadDataParams.kStep = AscendC::DivCeil(baseK * sizeof(half), 32);
	488	+ loadDataParams.srcStride = AscendC::DivCeil(baseN, CUBE_BLOCK);
	489	+ loadDataParams.dstStride = AscendC::DivCeil(curNAlign, CUBE_BLOCK);
	490	+ loadDataParams.ifTranspose = false;
	491	+ AscendC::LoadData(b2Local, b1Local[srcAddr], loadDataParams);
	492	+ }
	493	+#endif
	494	+ }
	495	+
	496	+ // A2 * B2 → CO1: 执行一次 Mmad 矩阵乘累加,A矩阵大小baseM * baseK,B矩阵baseK * baseN
	497	+ // kBlockIdx: K 方向 baseK 块索引，用于判断是否首次累加
	498	+ // mBlockIdx/nBlockIdx: M/N 方向块索引，用于判断尾块实际尺寸
	499	+ __aicore__ inline void Compute(AscendC::LocalTensor<float> cLocal,
	500	+ AscendC::LocalTensor<half> a2Local, AscendC::LocalTensor<half> b2Local,
	501	+ uint32_t kBlockIdx, uint32_t mBlockIdx, uint32_t nBlockIdx)
	502	+ {
	503	+ // MTE1_M 正向同步：等待 MTE1 完成 L1 → L0 搬运
	504	+ AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(mte1DBFlag);
	505	+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(mte1DBFlag);
	506	+ // 尾块使用 tailM/tailN，否则使用 baseM/baseN
	507	+ uint32_t curM = (mBlockIdx != baseMCount) ? baseM : tailM;
	508	+ uint32_t curN = (nBlockIdx != baseNCount) ? baseN : tailN;
	509	+ AscendC::MmadParams mmadParams;
	510	+ mmadParams.m = curM;
	511	+ mmadParams.n = curN;
	512	+ mmadParams.k = baseK;
	513	+ mmadParams.cmatrixInitVal = (kBlockIdx == 0);
	514	+ mmadParams.unitFlag = (kBlockIdx != kLoopCount - 1) ? 2 : 3; //开启UnitFlag
	515	+ AscendC::Mmad(cLocal, a2Local, b2Local, mmadParams);
	516	+ // M_MTE1 反向同步：通知下一轮 DataLoad 可以覆盖 L0 缓冲区
	517	+ AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(mte1DBFlag);
	518	+ mte1DBFlag ^= 1;
	519	+ }
	520	+
	521	+ // CO1 → GM: 将计算结果从 L0C 写回 Global Memory
	522	+ // mBlockIdx/nBlockIdx: M/N 方向块索引，用于计算 GM 写入偏移及判断尾块尺寸
	523	+ __aicore__ inline void CopyOut(AscendC::LocalTensor<float> cLocal, uint32_t mBlockIdx, uint32_t nBlockIdx)
	524	+ {
	525	+ // 尾块使用对齐后的 tailMAlign 作为 srcStride，使用 tailM/tailN 作为实际行列数
	526	+ uint32_t curMAlign = (mBlockIdx != baseMCount) ? baseM : tailMAlign;
	527	+ uint32_t curM = (mBlockIdx != baseMCount) ? baseM : tailM;
	528	+ uint32_t curN = (nBlockIdx != baseNCount) ? baseN : tailN;
	529	+ AscendC::FixpipeParamsV220 fixpipeParams;
	530	+ fixpipeParams.nSize = curN;
	531	+ fixpipeParams.mSize = curM;
	532	+ fixpipeParams.srcStride = curMAlign;
	533	+ fixpipeParams.dstStride = N;
	534	+ fixpipeParams.quantPre = QuantMode_t::F322F16;
	535	+ fixpipeParams.ndNum = 1;
	536	+ fixpipeParams.srcNdStride = 0;
	537	+ fixpipeParams.dstNdStride = 0;
	538	+ fixpipeParams.unitFlag = 3;
	539	+ AscendC::Fixpipe(cGM[mBlockIdx * baseM * N + nBlockIdx * baseN], cLocal, fixpipeParams);
	540	+ }
	541	+
	542	+private:
	543	+ AscendC::GlobalTensor<half> aGM;
	544	+ AscendC::GlobalTensor<half> bGM;
	545	+ AscendC::GlobalTensor<half> cGM;
	546	+ AscendC::GlobalTensor<half> aGMOri;
	547	+ AscendC::GlobalTensor<half> bGMOri;
	548	+ AscendC::GlobalTensor<half> cGMOri;
	549	+ uint32_t actualSingleCoreM, actualSingleCoreN;
	550	+ uint32_t mLoopCount, nLoopCount, kLoopCount;
	551	+ uint32_t baseMCount, baseNCount;
	552	+ uint32_t tailM, tailN;
	553	+ uint32_t tailMAlign, tailNAlign;
	554	+ uint8_t mte1DBFlag = 0;
	555	+};
	556	+
	557	+template <uint32_t M, uint32_t K, uint32_t N, uint32_t baseM, uint32_t baseK, uint32_t baseN,
	558	+ uint32_t singleCoreM, uint32_t singleCoreK, uint32_t singleCoreN,
	559	+ uint32_t stepKa, uint32_t stepKb>
	560	+__global__ __cube__ void mmad_custom(GM_ADDR a, GM_ADDR b, GM_ADDR c)
	561	+{
	562	+ AscendC::InitSocState();
	563	+ KernelMmad<M, K, N, baseM, baseK, baseN, singleCoreM, singleCoreK, singleCoreN,
	564	+ stepKa, stepKb> op;
	565	+ op.Init(a, b, c);
	566	+ if constexpr (scenarioNum == 1) {
	567	+ op.Process();
	568	+ } else if constexpr (scenarioNum == 2) {
	569	+ op.ProcessL2Cache();
	570	+ }
	571	+}
	572	+
	573	+int32_t main(int32_t argc, char *argv[])
	574	+{
	575	+ constexpr uint32_t M = 8192;
	576	+ constexpr uint32_t K = 8192;
	577	+ constexpr uint32_t N = 8192;
	578	+ constexpr uint32_t baseK = 64;
	579	+ constexpr uint32_t baseN = 256;
	580	+ constexpr uint32_t singleCoreK = 8192;
	581	+ constexpr uint32_t stepKb = 4;
	582	+ constexpr uint32_t singleCoreM = (scenarioNum == 1) ? 2048 : 1024;
	583	+
	584	+#if (NPU_ARCH == 2201)
	585	+ constexpr uint32_t baseM = 128;
	586	+ constexpr uint32_t stepKa = 8;
	587	+ constexpr uint32_t singleCoreN = 1536;
	588	+ constexpr uint32_t numBlocks = 24;
	589	+#elif (NPU_ARCH == 3510)
	590	+ constexpr uint32_t baseM = 256;
	591	+ constexpr uint32_t stepKa = 4;
	592	+ constexpr uint32_t singleCoreN = 1024;
	593	+ constexpr uint32_t numBlocks = 32;
	594	+#endif
	595	+ size_t aFileSize = M * K * sizeof(int16_t);
	596	+ size_t bFileSize = K * N * sizeof(int16_t);
	597	+ size_t cFileSize = M * N * sizeof(int16_t);
	598	+
	599	+ aclInit(nullptr);
	600	+ int32_t deviceId = 0;
	601	+ aclrtSetDevice(deviceId);
	602	+ aclrtStream stream = nullptr;
	603	+ aclrtCreateStream(&stream);
	604	+
	605	+ uint8_t *aHost;
	606	+ uint8_t *aDevice;
	607	+ aclrtMallocHost((void **)(&aHost), aFileSize);
	608	+ aclrtMalloc((void **)&aDevice, aFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	609	+ ReadFile("./input/x1_gm.bin", aFileSize, aHost, aFileSize);
	610	+ aclrtMemcpy(aDevice, aFileSize, aHost, aFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	611	+
	612	+ uint8_t *bHost;
	613	+ uint8_t *bDevice;
	614	+ aclrtMallocHost((void **)(&bHost), bFileSize);
	615	+ aclrtMalloc((void **)&bDevice, bFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	616	+ ReadFile("./input/x2_gm.bin", bFileSize, bHost, bFileSize);
	617	+ aclrtMemcpy(bDevice, bFileSize, bHost, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	618	+
	619	+ uint8_t *cHost;
	620	+ uint8_t *cDevice;
	621	+ aclrtMallocHost((void **)(&cHost), cFileSize);
	622	+ aclrtMalloc((void **)&cDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	623	+
	624	+ mmad_custom<M, K, N, baseM, baseK, baseN, singleCoreM, singleCoreK, singleCoreN,
	625	+ stepKa, stepKb>
	626	+ <<<numBlocks, nullptr, stream>>>(aDevice, bDevice, cDevice);
	627	+ aclrtSynchronizeStream(stream);
	628	+
	629	+ aclrtMemcpy(cHost, cFileSize, cDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST);
	630	+ WriteFile("./output/output.bin", cHost, cFileSize);
	631	+
	632	+ aclrtFree(aDevice);
	633	+ aclrtFreeHost(aHost);
	634	+ aclrtFree(bDevice);
	635	+ aclrtFreeHost(bHost);
	636	+ aclrtFree(cDevice);
	637	+ aclrtFreeHost(cHost);
	638	+
	639	+ aclrtDestroyStream(stream);
	640	+ aclrtResetDevice(deviceId);
	641	+ aclFinalize();
	642	+ return 0;
	643	+}

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_basic_API_high_performance/scripts/gen_data.py+41-0

@@ -0,0 +1,41 @@
	1	+#!/usr/bin/python3
	2	+# coding=utf-8
	3	+
	4	+# ----------------------------------------------------------------------------------------------------------
	5	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	6	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	7	+# CANN Open Software License Agreement Version 2.0 (the "License").
	8	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	9	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	10	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	11	+# See LICENSE in the root of the software repository for the full text of the License.
	12	+# ----------------------------------------------------------------------------------------------------------
	13	+
	14	+
	15	+import os
	16	+import numpy as np
	17	+
	18	+def gen_golden_data():
	19	+ m, n, k, is_bias = 8192, 8192, 8192, False
	CANN-robot4月16日代码逻辑和结构: 硬编码了矩阵维度(m, n, k)和is_bias标志。这使得代码缺乏灵活性，无法适应不同规模的测试需求。如果需要生成不同大小的测试数据，必须修改源代码。问题类型: 代码逻辑和结构文件路径: `examples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/mmad_high_performance/scripts/gen_data.py` 行号: 20 问题代码: m, n, k, is_bias = 8192, 8192, 8192, False 修改建议: 考虑将这些参数改为函数参数，或者从配置文件/命令行参数中读取。例如：def gen_golden_data(m=8192, n=8192, k=8192, is_bias=False): --- 此评论由代码审查工具自动生成
	20	+
	21	+ os.makedirs("input", exist_ok=True)
	22	+ os.makedirs("output", exist_ok=True)
	23	+
	24	+ x1_gm = np.random.uniform(-1, 1, [m, k]).astype(np.float16)
	25	+ x2_gm = np.random.uniform(-1, 1, [k, n]).astype(np.float16)
	26	+ if is_bias:
	27	+ bias_gm = np.random.uniform(-10, 10, [n]).reshape([n]).astype(np.float16)
	28	+ golden = np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32)).astype(np.float32) + bias_gm
	29	+ bias_gm.tofile("./input/bias_gm.bin")
	30	+ else:
	31	+ golden = np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32)).astype(np.float32)
	32	+ golden = golden.astype(np.float16)
	33	+ x1_gm.tofile("./input/x1_gm.bin")
	CANN-robot4月16日可维护性和扩展性: 文件路径硬编码为相对路径'./input/'和'./output/'。这限制了代码的灵活性，如果脚本在其他目录运行或需要输出到不同位置，将无法正常工作。问题类型: 可维护性和扩展性文件路径: `examples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/mmad_high_performance/scripts/gen_data.py` 行号: 34 问题代码: x1_gm.tofile("./input/x1_gm.bin") 修改建议: 将输出目录路径作为函数参数或配置变量，或者至少使用os.path.join来构建路径以提高可移植性。例如：output_dir = os.path.join(os.getcwd(), 'output') --- 此评论由代码审查工具自动生成
	34	+ # x2_gm transpose to match B matrix transpose
	35	+ x2_gm = x2_gm.transpose()
	36	+ x2_gm.tofile("./input/x2_gm.bin")
	37	+ golden.tofile("./output/golden.bin")
	38	+
	39	+
	40	+if __name__ == "__main__":
	41	+ gen_golden_data()

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_basic_API_high_performance/scripts/verify_result.py+58-0

@@ -0,0 +1,58 @@
	1	+#!/usr/bin/python3
	2	+# coding=utf-8
	3	+
	4	+# ----------------------------------------------------------------------------------------------------------
	5	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	6	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	7	+# CANN Open Software License Agreement Version 2.0 (the "License").
	8	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	9	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	10	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	11	+# See LICENSE in the root of the software repository for the full text of the License.
	12	+# ----------------------------------------------------------------------------------------------------------
	13	+
	14	+
	15	+import sys
	16	+import numpy as np
	17	+
	18	+
	19	+# for float32
	20	+relative_tol = 1e-3
	21	+absolute_tol = 1e-3
	22	+error_tol = 1e-3
	23	+
	24	+
	25	+def verify_result(output, golden):
	26	+ output = np.fromfile(output, dtype=np.float16).reshape(-1)
	27	+ golden = np.fromfile(golden, dtype=np.float16).reshape(-1)
	28	+ different_element_results = np.isclose(output,
	29	+ golden,
	30	+ rtol=relative_tol,
	31	+ atol=absolute_tol,
	32	+ equal_nan=True)
	33	+ different_element_indexes = np.where(different_element_results == False)[0]
	34	+ for index in range(len(different_element_indexes)):
	35	+ real_index = different_element_indexes[index]
	36	+ golden_data = golden[real_index]
	37	+ output_data = output[real_index]
	38	+ print(
	39	+ "data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" %
	40	+ (real_index, golden_data, output_data,
	41	+ abs(output_data - golden_data) / golden_data))
	42	+ if index == 100:
	43	+ break
	44	+ error_ratio = float(different_element_indexes.size) / golden.size
	45	+ print("error ratio: %.4f, tolerance: %.4f" % (error_ratio, error_tol))
	46	+ return error_ratio <= error_tol
	47	+
	48	+
	49	+if __name__ == '__main__':
	50	+ try:
	51	+ res = verify_result(sys.argv[1], sys.argv[2])
	52	+ if not res:
	53	+ raise ValueError("[ERROR] result error")
	54	+ else:
	55	+ print("test pass!")
	56	+ except Exception as e:
	57	+ print(e)
	58	+ sys.exit(1)

Mexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_high_performance/README.md+2-2

	@@ -478,7 +478,7 @@ MTE2理论耗时：
478	$$MTE2理论耗时 =\frac{HBM读入数据总量}{1.8TB/s} +\frac{L2Cache读入数据总量}{5TB/s} = 2672.44\mu s$$	478	$$MTE2理论耗时 =\frac{HBM读入数据总量}{1.8TB/s} +\frac{L2Cache读入数据总量}{5TB/s} = 2672.44\mu s$$
479		479
480	Case 8 MTE2耗时误差：	480	Case 8 MTE2耗时误差：
481	-$$MTE2耗时误差 = \frac{{3435.584\mu s} - {2672.44\mu s}}{{2672.44\mu s}} = 20.77\%$$	481	+$$MTE2耗时误差 = \frac{{3435.584\mu s} - {2672.44\mu s}}{{2672.44\mu s}} = 28.55\%$$
482		482
483	当前MTE2耗时与理论值相差较大，因为实际芯片L2Cache大小为192MB，当前L2Cache切分策略较简单；另一方面当MTE2搬运场景为ND2NZ（GM数据Layout为ND，搬运到L1时需做ND→NZ格式转换）时，L2Cache带宽会降低。用户可进一步优化L2Cache切分策略以提高MTE2带宽。	483	当前MTE2耗时与理论值相差较大，因为实际芯片L2Cache大小为192MB，当前L2Cache切分策略较简单；另一方面当MTE2搬运场景为ND2NZ（GM数据Layout为ND，搬运到L1时需做ND→NZ格式转换）时，L2Cache带宽会降低。用户可进一步优化L2Cache切分策略以提高MTE2带宽。
484		484
	@@ -603,4 +603,4 @@ PROF_xxxx_XXXXXX
603	├── msprof_*.json	603	├── msprof_*.json
604	├── xx_*.csv	604	├── xx_*.csv
605	└── README.txt	605	└── README.txt
606	-```	606	+```