Star731
1029
代码
代码
Issues57
Pull Requests177
Actions
项目讨论
Wiki
分析
Star731
1029
已合并
新增C API矩阵搬入样例 #5460
UFOunder创建于 9 天前
新增C API矩阵搬入样例 #5460
已合并
UFOunder创建于 9 天前
从UFOunder/asc-devkit_3107master合入到cann/asc-devkitmaster
共 18 个文件变更+1854-0
Mexamples/02_simd_c_api/03_c_api/03_matrix_compute/README.md+1-0
@@ -11,5 +11,6 @@
11 11 
12| 目录名称 | 功能描述 | 支持的产品 |12| 目录名称 | 功能描述 | 支持的产品 |
13|--------------------------------------------------------------------------------------------| ---------------------------------------------------- | --- |13|--------------------------------------------------------------------------------------------| ---------------------------------------------------- | --- |
14+| [load_data_l12l0_950](./load_data_l12l0_950) | 本样例展示如何调用`asc_copy_l12l0a`/`asc_copy_l12l0b`及其`_transpose`接口，将A/B矩阵从L1搬运到L0A/L0B Buffer。 | Ascend 950PR/Ascend 950DT |
14| [mmad](./mmad) | 本样例以 int8_t 和 bfloat16 两种输入数据类型为例，演示如何通过C API实现矩阵乘法（C = A × B + Bias） | Ascend 950PR/Ascend 950DT |15| [mmad](./mmad) | 本样例以 int8_t 和 bfloat16 两种输入数据类型为例，演示如何通过C API实现矩阵乘法（C = A × B + Bias） | Ascend 950PR/Ascend 950DT |
15| [mmad_mx](./mmad_mx) | 本样例演示如何通过C API实现带量化功能的MX矩阵乘法。A、B矩阵使用FP4或FP8数据类型，ScaleA、ScaleB使用fp8_e8m0_t数据类型；包含C矩阵初始化为0、从BiasTable Buffer初始化以及从L0C累加三种场景，并介绍K轴对齐和FP8尾部清零 | Ascend 950PR/Ascend 950DT |16| [mmad_mx](./mmad_mx) | 本样例演示如何通过C API实现带量化功能的MX矩阵乘法。A、B矩阵使用FP4或FP8数据类型，ScaleA、ScaleB使用fp8_e8m0_t数据类型；包含C矩阵初始化为0、从BiasTable Buffer初始化以及从L0C累加三种场景，并介绍K轴对齐和FP8尾部清零 | Ascend 950PR/Ascend 950DT |
Mexamples/02_simd_c_api/03_c_api/03_matrix_compute/README_en.md+1-0
@@ -11,5 +11,6 @@ This sample set introduces typical usage of different features of matrix computa
11 11 
12| Directory Name | Function Description | Supported Products |12| Directory Name | Function Description | Supported Products |
13|--------------------------------------------------------------------------------------------| ---------------------------------------------------- | --- |13|--------------------------------------------------------------------------------------------| ---------------------------------------------------- | --- |
14+| [load_data_l12l0_950](./load_data_l12l0_950) | This example demonstrates how to call `asc_copy_l12l0a` / `asc_copy_l12l0b` and their `_transpose` variants to transfer A / B matrices from L1 to L0A / L0B Buffer | Ascend 950PR/Ascend 950DT |
14| [mmad](./mmad) | This example uses int8_t and bfloat16 as two input data types to demonstrate how to implement matrix multiplication (C = A x B + Bias) through the C_API | Ascend 950PR/Ascend 950DT |15| [mmad](./mmad) | This example uses int8_t and bfloat16 as two input data types to demonstrate how to implement matrix multiplication (C = A x B + Bias) through the C_API | Ascend 950PR/Ascend 950DT |
15| [mmad_mx](./mmad_mx) | This example demonstrates MX matrix multiplication with quantization through the C API. Matrices A and B use FP4 or FP8 data types, while ScaleA and ScaleB use fp8_e8m0_t. It covers three scenarios: initializing C to zero, initializing C from the BiasTable Buffer, and accumulating from L0C. It also describes K-axis alignment and FP8 tail clearing | Ascend 950PR/Ascend 950DT |16| [mmad_mx](./mmad_mx) | This example demonstrates MX matrix multiplication with quantization through the C API. Matrices A and B use FP4 or FP8 data types, while ScaleA and ScaleB use fp8_e8m0_t. It covers three scenarios: initializing C to zero, initializing C from the BiasTable Buffer, and accumulating from L0C. It also describes K-axis alignment and FP8 tail clearing | Ascend 950PR/Ascend 950DT |
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/CMakeLists.txt+39-0
@@ -0,0 +1,39 @@
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+cmake_minimum_required(VERSION 3.16)
12+ 
13+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
14+ 
15+if(NOT CMAKE_ASC_ARCHITECTURES STREQUAL "dav-3510")
16+ message(FATAL_ERROR "load_data_l12l0_950 only supports dav-3510, but got ${CMAKE_ASC_ARCHITECTURES}")
17+endif()
18+ 
19+find_package(ASC REQUIRED)
20+ 
21+project(kernel_samples LANGUAGES ASC CXX)
22+ 
23+add_executable(demo
24+ load_data_l12l0_950.asc
25+)
26+ 
27+set(SCENARIO_NUM "1" CACHE STRING "Scene number (1-7)")
28+if(NOT SCENARIO_NUM MATCHES "^[1-7]$")
29+ message(FATAL_ERROR "SCENARIO_NUM must be 1, 2, 3, 4, 5, 6 or 7, but got ${SCENARIO_NUM}")
30+endif()
31+ 
32+message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}")
33+target_compile_definitions(demo PRIVATE
34+ SCENARIO_NUM=${SCENARIO_NUM}
35+)
36+ 
37+target_compile_options(demo PRIVATE
38+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
39+)
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/README.md+642-0
@@ -0,0 +1,642 @@
1+# load_data_l12l0_950样例
2+ 
3+## 概述
4+ 
5+本样例展示如何调用`asc_copy_l12l0a` / `asc_copy_l12l0b`及其`_transpose`接口，将A / B矩阵从L1搬运到L0A / L0B Buffer。
6+ 
7+样例覆盖int8_t、half、float三种输入数据类型，以及A / B矩阵转置和不转置输入场景。L0A数据排布为Nz格式，L0B数据排布为Zn格式。
8+ 
9+## 本样例支持的产品及CANN软件版本
10+ 
11+| 产品 | CANN软件版本 |
12+|------|-------------|
13+| Ascend 950PR/Ascend 950DT | >= CANN 9.2.0 |
14+ 
15+## 目录结构介绍
16+ 
17+```
18+├── load_data_l12l0_950
19+│ ├── figures // 图示
20+│ ├── scripts
21+│ │ ├── gen_data.py // 输入数据和真值数据生成脚本
22+│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本
23+│ ├── CMakeLists.txt // 编译工程文件
24+│ ├── data_utils.h // 数据读入写出函数
25+│ ├── load_data_l12l0_950.asc // SIMD C API样例实现 & 调用样例
26+│ └── README.md // 样例说明文档
27+```
28+ 
29+## 样例描述
30+ 
31+### 1. 整体流程
32+ 
33+样例整体流程如下：
34+ 
35+```text
36+GM(ND) ----------> L1(Nz) -> L0A(Nz)/L0B(Zn) -> L0C(Nz) -> GM(ND)
37+ │ │ │ │
38+asc_copy_gm2l1_nd2nz asc_copy_l12l0* asc_mmad asc_copy_l0c2gm
39+```
40+ 
41+**步骤详解**：
42+ 
43+1. **GM -> L1**：使用`asc_set_gm2l1_nz_para`配置Nz排布参数，使用`asc_copy_gm2l1_nd2nz`实现ND到Nz格式转换。
44+2. **L1 -> L0A / L0B**：使用`asc_copy_l12l0a` / `asc_copy_l12l0b`进行非转置搬运，使用对应的`_transpose`接口进行转置搬运。
45+3. **矩阵乘**：使用`asc_mmad`接口执行矩阵乘法。
46+4. **L0C -> GM**：使用`asc_set_l0c_copy_nz_para`配置Nz到ND格式转换参数，使用`asc_copy_l0c2gm`接口搬出结果。
47+ 
48+本样例固定矩阵乘规格为`[m, n, k] = [40, 50, 70]`，输出C矩阵为`[40, 50]`，ND格式。核函数通过`scenario_num`选择不同场景。
49+ 
50+### 2. 场景关键差异
51+ 
52+各场景的输入shape、转置标志和L1 -> L0搬运方式见[表1](#表1)。
53+ 
54+<a name="表1"></a>
55+<table border="2" align="center">
56+<caption style="font-weight: normal;">
57+ <span style="font-weight: bold; font-size: 1.2em;">📌 表1：scenario_num不同取值的含义</span></caption>
58+ <tr>
59+ <td ><span style="font-weight: bold;">scenario_num</span></td>
60+ <td><span style="font-weight: bold;">输入数据类型</span></td>
61+ <td><span style="font-weight: bold;">输出数据类型</span></td>
62+ <td><span style="font-weight: bold;">输入shape</span></td>
63+ <td><span style="font-weight: bold;">is_a_transpose</span></td>
64+ <td><span style="font-weight: bold;">is_b_transpose</span></td>
65+ <td><span style="font-weight: bold;">是否多搬多算</span></td>
66+ <td><span style="font-weight: bold;">L1 -> L0搬运方式</span></td>
67+ </tr>
68+ <tr>
69+ <td><span style="font-weight: bold;">1</span></td>
70+ <td rowspan="2" >int8_t</td>
71+ <td rowspan="2" >int32_t</td>
72+ <td>A [40, 70]<br>B [50, 70]</td>
73+ <td>false</td>
74+ <td>true</td>
75+ <td>否</td>
76+ <td><code>asc_copy_l12l0a</code> / <code>asc_copy_l12l0b</code></td>
77+ </tr>
78+ <tr>
79+ <td><span style="font-weight: bold;">2</span></td>
80+ <td>A [70, 40]<br>B [70, 50]</td>
81+ <td>true</td>
82+ <td>false</td>
83+ <td>是</td>
84+ <td><code>asc_copy_l12l0a_transpose</code> / <code>asc_copy_l12l0b_transpose</code></td>
85+ </tr>
86+ <tr>
87+ <td><span style="font-weight: bold;">3</span></td>
88+ <td rowspan="2" >half</td>
89+ <td rowspan="2" >float</td>
90+ <td>A [40, 70]<br>B [50, 70]</td>
91+ <td>false</td>
92+ <td>true</td>
93+ <td>否</td>
94+ <td><code>asc_copy_l12l0a</code> / <code>asc_copy_l12l0b</code></td>
95+ </tr>
96+ <tr>
97+ <td><span style="font-weight: bold;">4</span></td>
98+ <td>A [70, 40]<br>B [70, 50]</td>
99+ <td>true</td>
100+ <td>false</td>
101+ <td>否</td>
102+ <td><code>asc_copy_l12l0a_transpose</code> / <code>asc_copy_l12l0b_transpose</code></td>
103+ </tr>
104+ <tr>
105+ <td><span style="font-weight: bold;">5</span></td>
106+ <td rowspan="2" >float</td>
107+ <td rowspan="2" >float</td>
108+ <td>A [40, 70]<br>B [50, 70]</td>
109+ <td>false</td>
110+ <td>true</td>
111+ <td>否</td>
112+ <td><code>asc_copy_l12l0a</code> / <code>asc_copy_l12l0b</code></td>
113+ </tr>
114+ <tr>
115+ <td><span style="font-weight: bold;">6</span></td>
116+ <td>A [70, 40]<br>B [70, 50]</td>
117+ <td>true</td>
118+ <td>false</td>
119+ <td>否</td>
120+ <td><code>asc_copy_l12l0a_transpose</code> / <code>asc_copy_l12l0b_transpose</code></td>
121+ </tr>
122+ <tr>
123+ <td><span style="font-weight: bold;">7</span></td>
124+ <td rowspan="1" >int8_t</td>
125+ <td rowspan="1" >int32_t</td>
126+ <td>A [70, 40]<br>B [70, 50]</td>
127+ <td>true</td>
128+ <td>false</td>
129+ <td>否</td>
130+ <td>for循环 + <code>asc_copy_l12l0a_transpose</code> / <code>asc_copy_l12l0b_transpose</code></td>
131+ </tr>
132+</table>
133+ 
134+下面按搬运模式说明场景关键差异。
135+ 
136+**C API不转置搬运：场景1 / 3 / 5**
137+ 
138+- A矩阵不转置输入`[m, k]`，B矩阵转置输入`[n, k]`。
139+- L1 -> L0A / L0B均不需要转置，分别调用`asc_copy_l12l0a` / `asc_copy_l12l0b`。
140+ 
141+**C API转置搬运，单次调用：场景2 / 4 / 6**
142+ 
143+- A矩阵转置输入`[k, m]`，B矩阵不转置输入`[k, n]`。
144+- L1 -> L0A / L0B均需要转置，分别调用`asc_copy_l12l0a_transpose` / `asc_copy_l12l0b_transpose`。
145+- 场景2中，int8_t转置单次搬运会导致m方向多搬多算脏数据分形，`asc_copy_l0c2gm`通过参数`m_size = m`只搬出有效区域。
146+ 
147+**C API转置搬运，for循环调用：场景7**
148+ 
149+- 场景7与场景2同为int8_t转置输入。
150+- A矩阵采用for循环调用`asc_copy_l12l0a_transpose`，在写入L0A时跳过m方向尾部脏数据分形，使矩阵计算时没有额外脏数据分形参与。
151+ 
152+### 3. 参数说明
153+ 
154+后续代码和参数说明会反复使用分形和对齐相关变量，本节先集中定义这些概念。
155+ 
156+- `fractal_shape`：小分形的shape。B8 / B16 / B32输入数据类型的shape为`[16, 32 / sizeof(T)]`，其中`T`表示输入数据类型。本样例涉及的数据类型分形相关信息见[表2](#表2)。
157+- `fractal_size`：1个小分形包含的元素个数，具体见[表2](#表2)。
158+- `fractal_num`：当从L1 -> L0A / L0B需要转置时，`asc_copy_l12l0a_transpose` / `asc_copy_l12l0b_transpose`会按正方形矩阵转置。需要多个连续小分形合并为一个正方形矩阵时，`fractal_num`表示该正方形矩阵包含的小分形个数，具体见[表2](#表2)。
159+ 
160+<a name="表2"></a>
161+<table border="2" align="center">
162+<caption style="font-weight: normal;">
163+ <span style="font-weight: bold; font-size: 1.2em;">📌 表2：不同数据类型分形相关信息</span></caption>
164+ <tr>
165+ <td></td>
166+ <td align="center"><span style="font-weight: bold;">fractal_shape</span></td>
167+ <td align="center"><span style="font-weight: bold;">fractal_size</span></td>
168+ <td align="center"><span style="font-weight: bold;">fractal_num</span></td>
169+ </tr>
170+ <tr>
171+ <td align="center"><span style="font-weight: bold;">B8</span></td>
172+ <td align="center">[16, 32]</td>
173+ <td align="center">512</td>
174+ <td align="center">2</td>
175+ </tr>
176+ <tr>
177+ <td align="center"><span style="font-weight: bold;">B16</span></td>
178+ <td align="center">[16, 16]</td>
179+ <td align="center">256</td>
180+ <td align="center">1</td>
181+ </tr>
182+ <tr>
183+ <td align="center"><span style="font-weight: bold;">B32</span></td>
184+ <td align="center">[16, 8]</td>
185+ <td align="center">128</td>
186+ <td align="center">2</td>
187+ </tr>
188+</table>
189+ 
190+- `ceil_align`：向上对齐操作。例如`m = 30`时，`ceil_align(30, 16) = 32`，表示将m轴向16对齐，对齐后m轴长度为32。
191+ 
192+ ```cpp
193+ #define ceil_div(value, align) (((value) + (align) - 1) / (align))
194+ #define ceil_align(value, align) (ceil_div((value), (align)) * (align))
195+ ```
196+ 
197+- `ceil_div`：向上取整除法，一般用于求解向上对齐后的循环次数。
198+- `m_align_value`：m轴向`m_align_value`对齐。例如`m_align_value = 32`，代表m轴对齐到32；依次类推还有`n_align_value`、`ka_align_value`、`kb_align_value`。
199+- `m_align_l1`和`m_align_l0`：A矩阵分别在L1和L0A上时，m轴对齐后的值。依次类推还有`n_align_l1`、`n_align_l0`、`ka_align_l1`、`ka_align_l0`、`kb_align_l1`、`kb_align_l0`。
200+ 
201+### 4. 对齐要求
202+ 
203+A / B矩阵在L1和L0上各轴的对齐要求不同，后续配置`asc_copy_l12l0a` / `asc_copy_l12l0b`、`asc_mmad`和`asc_copy_l0c2gm`参数时会用到这些值。[表3](#表3)、[表4](#表4)总结scenario_num取1-6时A / B矩阵在L1和L0上的对齐要求；[表5](#表5)、[表6](#表6)总结scenario_num取7时的对齐要求。
204+ 
205+<a name="表3"></a>
206+<table border="2" align="center">
207+<caption style="font-weight: normal;">
208+ <span style="font-weight: bold; font-size: 1.2em;">📌 表3：A、B矩阵在L1上各个轴对齐要求（L1排布格式为Nz）</span></caption>
209+ <tr>
210+ <td></td>
211+ <td align="center"><span style="font-weight: bold;">B8 (fractal_num=2)</span></td>
212+ <td align="center"><span style="font-weight: bold;">B16 (fractal_num=1)</span></td>
213+ <td align="center"><span style="font-weight: bold;">B32 (fractal_num=2)</span></td>
214+ </tr>
215+ <tr>
216+ <td rowspan="2" align="center"><span style="font-weight: bold;">A矩阵不转置输入[m, k]</span></td>
217+ <td colspan="3" align="center">m_align_value = fractal_shape[0]</td>
218+ </tr>
219+ <tr>
220+ <td colspan="3" align="center" >ka_align_value = fractal_shape[1]</td>
221+ </tr>
222+ <tr>
223+ <td rowspan="2" align="center"><span style="font-weight: bold;">A矩阵转置输入[k, m]</span></td>
224+ <td colspan="2" align="center">ka_align_value = fractal_shape[0] * fractal_num</td>
225+ <td colspan="1" align="center">ka_align_value = fractal_shape[0]</td>
226+ </tr>
227+ <tr>
228+ <td colspan="2" align="center" >m_align_value = fractal_shape[1]</td>
229+ <td colspan="1" align="center" >m_align_value = fractal_shape[1] * fractal_num</td>
230+ </tr>
231+ <tr>
232+ <td rowspan="2" align="center"><span style="font-weight: bold;">B矩阵不转置输入[k, n]</span></td>
233+ <td colspan="2" align="center">kb_align_value = fractal_shape[0] * fractal_num</td>
234+ <td colspan="1" align="center">kb_align_value = fractal_shape[0]</td>
235+ </tr>
236+ <tr>
237+ <td colspan="2" align="center" >n_align_value = fractal_shape[1]</td>
238+ <td colspan="1" align="center" >n_align_value = fractal_shape[1] * fractal_num</td>
239+ </tr>
240+ <tr>
241+ <td rowspan="2" align="center"><span style="font-weight: bold;">B矩阵转置输入[n, k]</span></td>
242+ <td colspan="3" align="center">n_align_value = fractal_shape[0]</td>
243+ </tr>
244+ <tr>
245+ <td colspan="3" align="center" >kb_align_value = fractal_shape[1]</td>
246+ </tr>
247+</table>
248+ 
249+<a name="表4"></a>
250+<table border="2" align="center">
251+<caption style="font-weight: normal;">
252+ <span style="font-weight: bold; font-size: 1.2em;">📌 表4：A、B矩阵在L0上各个轴对齐要求</span></caption>
253+ <tr>
254+ <td></td>
255+ <td align="center"><span style="font-weight: bold;">B8 (fractal_num=2)</span></td>
256+ <td align="center"><span style="font-weight: bold;">B16 (fractal_num=1)</span></td>
257+ <td align="center"><span style="font-weight: bold;">B32 (fractal_num=2)</span></td>
258+ </tr>
259+ <tr>
260+ <td rowspan="2" align="center"><span style="font-weight: bold;">A矩阵不转置输入[m, k]，L1 -> L0A不需要转置</span></td>
261+ <td colspan="3" align="center">m_align_value = fractal_shape[0]</td>
262+ </tr>
263+ <tr>
264+ <td colspan="3" align="center" >ka_align_value = fractal_shape[1]</td>
265+ </tr>
266+ <tr>
267+ <td rowspan="2" align="center"><span style="font-weight: bold;">A矩阵转置输入[k, m]，L1 -> L0A需要转置</span></td>
268+ <td colspan="2" align="center">ka_align_value = fractal_shape[1]</td>
269+ <td >ka_align_value = fractal_shape[1] * fractal_num</td>
270+ </tr>
271+ <tr>
272+ <td colspan="2" align="center" >m_align_value = fractal_shape[0] * fractal_num</td>
273+ <td align="center" >m_align_value = fractal_shape[0]</td>
274+ </tr>
275+ <tr>
276+ <td rowspan="2" align="center"><span style="font-weight: bold;">B矩阵不转置输入[k, n]，L1 -> L0B需要转置</span></td>
277+ <td colspan="2" align="center">kb_align_value = fractal_shape[1]</td>
278+ <td align="center">kb_align_value = fractal_shape[1] * fractal_num</td>
279+ </tr>
280+ <tr>
281+ <td colspan="2" align="center">n_align_value = fractal_shape[0] * fractal_num</td>
282+ <td align="center" >n_align_value = fractal_shape[0]</td>
283+ </tr>
284+ <tr>
285+ <td rowspan="2" align="center"><span style="font-weight: bold;">B矩阵转置输入[n, k]，L1 -> L0B不需要转置</span></td>
286+ <td colspan="3" align="center">n_align_value = fractal_shape[0]</td>
287+ </tr>
288+ <tr>
289+ <td colspan="3" align="center" >kb_align_value = fractal_shape[1]</td>
290+ </tr>
291+</table>
292+ 
293+当scenario_num=7时，A矩阵用for循环 + `asc_copy_l12l0a_transpose`实现L1 -> L0A搬运，L0A只写入有效数据按分形对齐后的区域。
294+ 
295+<a name="表5"></a>
296+<table border="2" align="center">
297+<caption style="font-weight: normal;">
298+ <span style="font-weight: bold; font-size: 1.2em;">📌 表5：scenario_num=7，A、B矩阵在L1上各个轴对齐要求</span></caption>
299+ <tr>
300+ <td align="center" ></td>
301+ <td align="center" ><span style="font-weight: bold;">int8_t (fractal_num=2)</span></td>
302+ </tr>
303+ <tr>
304+ <td rowspan="2"><span style="font-weight: bold;">A矩阵转置输入[k, m]</span></td>
305+ <td align="center" >ka_align_value = fractal_shape[0] * fractal_num</td>
306+ </tr>
307+ <tr>
308+ <td align="center" >m_align_value = fractal_shape[1]</td>
309+ </tr>
310+ <tr>
311+ <td rowspan="2"><span style="font-weight: bold;">B矩阵不转置输入[k, n]</span></td>
312+ <td align="center" >kb_align_value = fractal_shape[0] * fractal_num</td>
313+ </tr>
314+ <tr>
315+ <td align="center" >n_align_value = fractal_shape[1]</td>
316+ </tr>
317+</table>
318+ 
319+<a name="表6"></a>
320+<table border="2" align="center">
321+<caption style="font-weight: normal;">
322+ <span style="font-weight: bold; font-size: 1.2em;">📌 表6：scenario_num=7，A、B矩阵在L0上各个轴对齐要求</span></caption>
323+ <tr>
324+ <td align="center" ></td>
325+ <td align="center" ><span style="font-weight: bold;">int8_t (fractal_num=2)</span></td>
326+ </tr>
327+ <tr>
328+ <td rowspan="2"><span style="font-weight: bold;">A矩阵转置输入[k, m]，L1 -> L0A需要转置</span></td>
329+ <td align="center" >m_align_value = fractal_shape[0]</td>
330+ </tr>
331+ <tr>
332+ <td align="center" >ka_align_value = fractal_shape[1]</td>
333+ </tr>
334+ <tr>
335+ <td rowspan="2"><span style="font-weight: bold;">B矩阵不转置输入[k, n]，L1 -> L0B需要转置</span></td>
336+ <td align="center" >kb_align_value = fractal_shape[1]</td>
337+ </tr>
338+ <tr>
339+ <td align="center" >n_align_value = fractal_shape[0] * fractal_num</td>
340+ </tr>
341+</table>
342+ 
343+ 
344+### 5. L1 -> L0 C API
345+ 
346+A / B矩阵在L1上均为Nz排布，在L0A、L0B上分别为Nz、Zn排布。L1 -> L0阶段根据目标存储位置和是否转置，调用`asc_copy_l12l0a`、`asc_copy_l12l0b`、`asc_copy_l12l0a_transpose`或`asc_copy_l12l0b_transpose`。
347+ 
348+阅读本章时可先按[表1](#表1)确定场景，再按[表3](#表3)-[表6](#表6)选择对齐值。上述C API重点关注以下参数：
349+ 
350+- 接口名：无`_transpose`后缀的接口执行非转置搬运，带`_transpose`后缀的接口执行转置搬运。
351+- `m_step`、`k_step`：分别表示本次搬运覆盖的row方向、col方向小分形个数。转置搬运时，B8 / B32需要按方型分形组合搬运，相关步长需满足`fractal_num`要求。
352+- `src_stride`、`dst_stride`：分别表示L1源矩阵和L0目标矩阵row方向相邻小分形之间的间隔。
353+- `m_start_position`、`k_start_position`：表示本次搬运在L1源矩阵中的起始小分形位置。场景7中通过循环修改`m_start_position`，分段搬运A矩阵。
354+ 
355+#### 5.1. A矩阵 L1 -> L0A 不转置
356+ 
357+L1 -> L0A不转置时，B8 / B16 / B32三种数据类型的参数配置基本相同，仅`fractal_shape`不同，可参考[表2](#表2)。以int8_t为例展示图示。
358+ 
359+<div align="center">
360+ <img src="figures/B8_A_l1_l0A_copy_l12l0.png" width="800"><br>
361+ 图1: int8_t数据类型下，L1 -> L0A不转置，调用asc_copy_l12l0a数据排布示意图
362+</div>
363+ 
364+参数配置要点：
365+ 
366+- `m_step = ceil_div(m_align_l1, fractal_shape[0])`，表示row方向搬运m轴对应的小分形个数。
367+- `k_step = ceil_div(ka_align_l1, fractal_shape[1])`，表示col方向搬运k轴对应的小分形个数。
368+- 调用`asc_copy_l12l0a`，表示只完成L1 Nz到L0A Nz的排布搬运，不做转置。
369+ 
370+```cpp
371+m_align_l1 = ceil_align(m, fractal_shape[0]); // 48
372+ka_align_l1 = ceil_align(k, fractal_shape[1]); // 96
373+m_align_l0 = ceil_align(m, fractal_shape[0]); // 48
374+ka_align_l0 = ceil_align(k, fractal_shape[1]); // 96
375+asc_copy_l12l0a(a_l0, a_l1, 0, 0,
376+ ceil_div(m_align_l1, fractal_shape[0]), // m_step = 3
377+ ceil_div(ka_align_l1, fractal_shape[1]), // k_step = 3
378+ ceil_div(m_align_l1, fractal_shape[0]), // src_stride = 3
379+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 3
380+```
381+ 
382+#### 5.2. A矩阵 L1 -> L0A 转置
383+ 
384+L1 -> L0A需要转置时调用`asc_copy_l12l0a_transpose`。该接口会按方型分形组合完成小分形转置和L0A Nz排布写入；不同数据类型的方型分形组合方式不同，下面按B8 / B16 / B32分别说明。
385+ 
386+##### 5.2.1. B8输入数据类型
387+ 
388+B8输入数据类型分形为16 * 32，L1 -> L0转置时会按照方型将row方向的2个16 * 32分形拼成1个32 * 32方型进行转置。下面以int8_t为例说明单次调用和for循环调用两种方式。
389+ 
390+**单次调用**
391+ 
392+调用一次`asc_copy_l12l0a_transpose`完成L1 -> L0A搬运和转置的图示如下：
393+ 
394+<div align="center">
395+ <img src="figures/B8_A_l1_l0A_trans_copy_l12l0.png" width="800"><br>
396+ 图2: int8_t数据类型下，L1 -> L0A转置，调用一次asc_copy_l12l0a_transpose数据排布示意图
397+</div>
398+ 
399+本样例`m = 40`，B8转置搬运时`m_align_l1 = ceil_align(m, fractal_shape[1]) = 64`，`m_align_l1 - m = 24 > 16`。单次调用`asc_copy_l12l0a_transpose`时，m方向会多搬1个无效分形；`asc_mmad`计算时需要设置`m = ceil_align(m, fractal_shape[0] * fractal_num)`让该分形参与计算，最后由`asc_copy_l0c2gm`通过`m_size = m`只搬出有效区域。
400+ 
401+```cpp
402+ka_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 96
403+m_align_l1 = ceil_align(m, fractal_shape[1]); // 64
404+m_align_l0 = ceil_align(m, fractal_shape[0] * fractal_num); // 64
405+ka_align_l0 = ceil_align(k, fractal_shape[1]); // 96
406+asc_copy_l12l0a_transpose(a_l0, a_l1, 0, 0,
407+ ceil_div(ka_align_l1, fractal_shape[0]), // m_step = 6
408+ ceil_div(m_align_l1, fractal_shape[1]), // k_step = 2
409+ ceil_div(ka_align_l1, fractal_shape[0]), // src_stride = 6
410+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 4
411+```
412+ 
413+**for循环调用**
414+ 
415+for循环调用多次`asc_copy_l12l0a_transpose`完成L1 -> L0A搬运和转置的图示如下：
416+ 
417+<div align="center">
418+ <img src="figures/B8_A_l1_l0A_trans_for_copy_l12l0.png" width="800"><br>
419+ 图3: int8_t数据类型下，L1 -> L0A转置，for循环调用多次asc_copy_l12l0a_transpose数据排布示意图
420+</div>
421+ 
422+for循环调用时，沿L1 row方向（A矩阵k轴）分段搬运，每次搬运k轴方向2个分形、m轴方向`ceil_div(m_align_l0, fractal_shape[1])`个分形。`dst_stride`按m方向有效数据向`fractal_shape[0]`对齐配置，写入L0A时跳过转置多读的m方向脏数据分形，使`asc_mmad`计算时m方向没有额外脏数据分形参与。
423+ 
424+```cpp
425+ka_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 96
426+m_align_l1 = ceil_align(m, fractal_shape[1]); // 64
427+m_align_l0 = ceil_align(m, fractal_shape[0]); // 48
428+ka_align_l0 = ceil_align(k, fractal_shape[1]); // 96
429+// 输入为int8类型，A矩阵[k,m]转置输入，L1 -> L0A需要转置
430+// for循环调用asc_copy_l12l0a_transpose，以k轴方向做for循环，每次循环在L1的k方向搬运2个分形，在L0A上跳过m方向尾脏数据分形，m方向多搬运数据不超过1个分形
431+uint16_t a_loop_num = ceil_div(ka_align_l0, fractal_shape[0] * fractal_num); // 3
432+uint32_t dst_offset = 0;
433+for (uint16_t loopIdx = 0; loopIdx < a_loop_num; ++loopIdx) {
434+ asc_copy_l12l0a_transpose(a_l0 + dst_offset, a_l1,
435+ int8_m_step_align * loopIdx, 0, // m_start_position, k_start_position
436+ int8_m_step_align, // m_step = 2
437+ ceil_div(m_align_l0, fractal_shape[1]), // k_step = 2
438+ ceil_div(ka_align_l1, fractal_shape[0]), // src_stride = 6
439+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 3
440+ dst_offset += ceil_align(m_align_l0, fractal_shape[0]) * fractal_shape[1]; // 1536
441+}
442+```
443+ 
444+##### 5.2.2. B16输入数据类型
445+ 
446+B16输入数据类型分形为16 * 16，一个分形即一个方型，L1 -> L0转置时会按照小分形进行转置，调用一次`asc_copy_l12l0a_transpose`即可完成L1 -> L0A数据搬运及转置。本样例以half为例展示图示。
447+ 
448+<div align="center">
449+ <img src="figures/B16_A_l1_l0A_trans_copy_l12l0.png" width="800"><br>
450+ 图4: half数据类型下，L1 -> L0A转置，调用一次asc_copy_l12l0a_transpose数据排布示意图
451+</div>
452+ 
453+参数配置要点：
454+ 
455+- `m_step = ceil_div(ka_align_l1, fractal_shape[0])`，表示row方向搬运k轴对应的小分形个数。
456+- `k_step = ceil_div(m_align_l1, fractal_shape[1])`，表示col方向搬运m轴对应的小分形个数。
457+- 调用`asc_copy_l12l0a_transpose`，表示搬运到L0A时完成转置。
458+ 
459+```cpp
460+ka_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 80
461+m_align_l1 = ceil_align(m, fractal_shape[1]); // 48
462+m_align_l0 = ceil_align(m, fractal_shape[0] * fractal_num); // 48
463+ka_align_l0 = ceil_align(k, fractal_shape[1]); // 80
464+asc_copy_l12l0a_transpose(a_l0, a_l1, 0, 0,
465+ ceil_div(ka_align_l1, fractal_shape[0]), // m_step = 5
466+ ceil_div(m_align_l1, fractal_shape[1]), // k_step = 3
467+ ceil_div(ka_align_l1, fractal_shape[0]), // src_stride = 5
468+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 3
469+```
470+ 
471+##### 5.2.3. B32输入数据类型
472+ 
473+B32输入数据类型分形为16 * 8，L1 -> L0转置时会按照方型将col方向的2个16 * 8分形拼成1个16 * 16方型进行转置。本样例以float为例展示图示。
474+ 
475+<div align="center">
476+ <img src="figures/B32_A_l1_l0A_trans_copy_l12l0.png" width="800"><br>
477+ 图5: float数据类型下，L1 -> L0A转置，调用一次asc_copy_l12l0a_transpose数据排布示意图
478+</div>
479+ 
480+本样例`m = 40`。B32转置搬运时需要将col方向2个分形拼成方型，**`k_step`必须是2的倍数**，因此L1上col方向（m方向）会多读1个无效分形。写入L0A后，k方向尾部会多1个无效分形；由于L0A为Nz排布，`asc_mmad`配置`k = matrix_k`即可只让有效k轴数据参与计算。
481+ 
482+```cpp
483+ka_align_l1 = ceil_align(k, fractal_shape[0]); // 80
484+m_align_l1 = ceil_align(m, fractal_shape[1] * fractal_num); // 48
485+m_align_l0 = ceil_align(m, fractal_shape[0]); // 48
486+ka_align_l0 = ceil_align(k, fractal_shape[1] * fractal_num); // 80
487+asc_copy_l12l0a_transpose(a_l0, a_l1, 0, 0,
488+ ceil_div(ka_align_l1, fractal_shape[0]), // m_step = 5
489+ ceil_div(m_align_l1, fractal_shape[1]), // k_step = 6
490+ ceil_div(ka_align_l1, fractal_shape[0]), // src_stride = 5
491+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 3
492+```
493+ 
494+#### 5.3. B矩阵 L1 -> L0B 不转置
495+ 
496+L1 -> L0B不转置时，B8 / B16 / B32三种数据类型的参数配置基本相同，仅`fractal_shape`不同，可参考[表2](#表2)。以float为例展示图示。
497+ 
498+<div align="center">
499+ <img src="figures/B32_B_l1_l0B_copy_l12l0.png" width="1100"><br>
500+ 图6: float数据类型下，L1 -> L0B不转置，调用一次asc_copy_l12l0b数据排布示意图
501+</div>
502+ 
503+参数配置要点：
504+ 
505+- `m_step = ceil_div(n_align_l1, fractal_shape[0])`，表示row方向搬运n轴对应的小分形个数。
506+- `k_step = ceil_div(kb_align_l1, fractal_shape[1])`，表示col方向搬运k轴对应的小分形个数。
507+- 调用`asc_copy_l12l0b`，表示只完成L1 Nz到L0B Zn的排布搬运，不做转置。
508+ 
509+```cpp
510+n_align_l1 = ceil_align(n, fractal_shape[0]); // 64
511+kb_align_l1 = ceil_align(k, fractal_shape[1]); // 72
512+kb_align_l0 = ceil_align(k, fractal_shape[1]); // 72
513+n_align_l0 = ceil_align(n, fractal_shape[0]); // 64
514+asc_copy_l12l0b(b_l0, b_l1, 0, 0,
515+ ceil_div(n_align_l1, fractal_shape[0]), // m_step = 4
516+ ceil_div(kb_align_l1, fractal_shape[1]), // k_step = 9
517+ ceil_div(n_align_l1, fractal_shape[0]), // src_stride = 4
518+ ceil_div(n_align_l0, fractal_shape[0])); // dst_stride = 4
519+```
520+ 
521+#### 5.4. B矩阵 L1 -> L0B 转置
522+ 
523+L1 -> L0B需要转置时调用`asc_copy_l12l0b_transpose`。该接口会按方型分形组合完成小分形转置和L0B Zn排布写入；不同数据类型的方型分形组合方式不同，下面按B8 / B16 / B32分别说明。
524+ 
525+##### 5.4.1. B8输入数据类型
526+ 
527+B8输入数据类型分形为16 * 32，L1 -> L0转置时会按照方型将row方向的2个16 * 32的分形拼成1个32 * 32的方型进行转置。本样例以int8_t为例展示图示。
528+ 
529+<div align="center">
530+ <img src="figures/B8_B_l1_l0B_trans_copy_l12l0.png" width="1100"><br>
531+ 图7: int8_t数据类型下，L1 -> L0B转置，调用一次asc_copy_l12l0b_transpose数据排布示意图
532+</div>
533+ 
534+本样例`k = 70`。B8转置搬运时需要将row方向2个分形拼成方型，**`m_step`必须是2的倍数**，因此L1上row方向（k方向）会多读1个无效分形。
535+ 
536+```cpp
537+kb_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 96
538+n_align_l1 = ceil_align(n, fractal_shape[1]); // 64
539+kb_align_l0 = ceil_align(k, fractal_shape[1]); // 96
540+n_align_l0 = ceil_align(n, fractal_shape[0] * fractal_num); // 64
541+asc_copy_l12l0b_transpose(b_l0, b_l1, 0, 0,
542+ ceil_div(kb_align_l1, fractal_shape[0]), // m_step = 6
543+ ceil_div(n_align_l1, fractal_shape[1]), // k_step = 2
544+ ceil_div(kb_align_l1, fractal_shape[0]), // src_stride = 6
545+ ceil_div(n_align_l0, fractal_shape[0])); // dst_stride = 4
546+```
547+ 
548+##### 5.4.2. B16输入数据类型
549+ 
550+B16输入数据类型分形为16 * 16，一个分形即一个方型，L1 -> L0转置时会按照小分形进行转置，调用一次`asc_copy_l12l0b_transpose`即可完成L1 -> L0B数据搬运及转置。本样例以half为例展示图示。
551+ 
552+<div align="center">
553+ <img src="figures/B16_B_l1_l0B_trans_copy_l12l0.png" width="800"><br>
554+ 图8: half数据类型下，L1 -> L0B转置，调用一次asc_copy_l12l0b_transpose数据排布示意图
555+</div>
556+ 
557+参数配置要点：
558+ 
559+- `m_step = ceil_div(kb_align_l1, fractal_shape[0])`，表示row方向搬运k轴对应的小分形个数。
560+- `k_step = ceil_div(n_align_l1, fractal_shape[1])`，表示col方向搬运n轴对应的小分形个数。
561+- 调用`asc_copy_l12l0b_transpose`，表示搬运到L0B时完成转置。
562+ 
563+```cpp
564+kb_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 80
565+n_align_l1 = ceil_align(n, fractal_shape[1]); // 64
566+kb_align_l0 = ceil_align(k, fractal_shape[1]); // 80
567+n_align_l0 = ceil_align(n, fractal_shape[0] * fractal_num); // 64
568+asc_copy_l12l0b_transpose(b_l0, b_l1, 0, 0,
569+ ceil_div(kb_align_l1, fractal_shape[0]), // m_step = 5
570+ ceil_div(n_align_l1, fractal_shape[1]), // k_step = 4
571+ ceil_div(kb_align_l1, fractal_shape[0]), // src_stride = 5
572+ ceil_div(n_align_l0, fractal_shape[0])); // dst_stride = 4
573+```
574+ 
575+##### 5.4.3. B32输入数据类型
576+ 
577+B32输入数据类型分形为16 * 8，L1 -> L0转置时会按照方型将col方向的2个16 * 8的分形拼成1个16 * 16的方型进行转置。本样例以float为例展示图示。
578+ 
579+<div align="center">
580+ <img src="figures/B32_B_l1_l0B_trans_copy_l12l0.png" width="1100"><br>
581+ 图9: float数据类型下，L1 -> L0B转置，调用一次asc_copy_l12l0b_transpose数据排布示意图
582+</div>
583+ 
584+本样例`n = 50`。B32转置搬运时需要将col方向2个分形拼成方型，**`k_step`必须是2的倍数**，因此L1上col方向（n方向）会多读1个无效分形。写入L0B后，k方向尾部会多1个无效分形；由于L0B为Zn排布，`asc_mmad`配置`k = matrix_k`即可只让有效k轴数据参与计算。
585+ 
586+```cpp
587+kb_align_l1 = ceil_align(k, fractal_shape[0]); // 80
588+n_align_l1 = ceil_align(n, fractal_shape[1] * fractal_num); // 64
589+kb_align_l0 = ceil_align(k, fractal_shape[1] * fractal_num); // 80
590+n_align_l0 = ceil_align(n, fractal_shape[0]); // 64
591+asc_copy_l12l0b_transpose(b_l0, b_l1, 0, 0,
592+ ceil_div(kb_align_l1, fractal_shape[0]), // m_step = 5
593+ ceil_div(n_align_l1, fractal_shape[1]), // k_step = 8
594+ ceil_div(kb_align_l1, fractal_shape[0]), // src_stride = 5
595+ ceil_div(n_align_l0, fractal_shape[0])); // dst_stride = 4
596+```
597+ 
598+## 编译运行
599+ 
600+在本样例根目录下执行如下步骤，编译并执行样例。
601+- 配置环境变量
602+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/zh/quick_start.md#prepare&install)，配置环境变量。
603+ ```bash
604+ source ${install_path}/cann/set_env.sh
605+ ```
606+ 
607+ > **说明：** `${install_path}` 为CANN包安装目录，未指定安装目录时默认安装至 `/usr/local/Ascend` 下。
608+- 样例执行
609+ 
610+ 在本样例目录下执行如下命令。
611+ ```bash
612+ SCENARIO_NUM=1
613+ mkdir -p build && cd build;
614+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO_NUM ..;make -j;
615+ python3 ../scripts/gen_data.py -scenario_num=$SCENARIO_NUM
616+ ./demo
617+ python3 ../scripts/verify_result.py -scenario_num=$SCENARIO_NUM output/output.bin output/golden.bin
618+ ```
619+ 使用NPU仿真模式时，添加 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
620+ 
621+ 示例如下：
622+ 
623+ ```bash
624+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO_NUM ..;make -j; # NPU仿真模式
625+ ```
626+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
627+ 
628+- 编译选项说明
629+ 
630+ | 选项 | 可选值 | 说明 |
631+ |------|--------|------|
632+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`sim` | 运行模式：NPU 运行、NPU仿真 |
633+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构：Ascend 950PR/Ascend 950DT |
634+ | `SCENARIO_NUM` | `1`-`7` | 场景编号：不同数据类型和转置组合 |
635+ 
636+- 执行结果
637+ 
638+ 执行结果如下，说明精度对比成功。
639+ 
640+ ```bash
641+ test pass!
642+ ```
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/README_en.md+641-0
@@ -0,0 +1,641 @@
1+# load_data_l12l0_950 Example
2+ 
3+## Overview
4+ 
5+This example demonstrates how to call `asc_copy_l12l0a` / `asc_copy_l12l0b` and their `_transpose` variants to transfer A / B matrices from L1 to L0A / L0B Buffer.
6+ 
7+The example covers three input data types: int8_t, half, and float, as well as A / B matrix transposed and non-transposed input scenarios. L0A data layout is Nz format, L0B data layout is Zn format.
8+ 
9+## Supported Products and CANN Versions
10+ 
11+| Product | CANN Version |
12+|---------|-------------|
13+| Ascend 950PR/Ascend 950DT | >= CANN 9.2.0 |
14+ 
15+## Directory Structure
16+ 
17+```
18+├── load_data_l12l0_950
19+│ ├── figures // Illustrations
20+│ ├── scripts
21+│ │ ├── gen_data.py // Script for generating input data and ground truth data
22+│ │ └── verify_result.py // Script for verifying whether output data matches ground truth data
23+│ ├── CMakeLists.txt // Build project file
24+│ ├── data_utils.h // Data read/write functions
25+│ ├── load_data_l12l0_950.asc // SIMD C API implementation and invocation example
26+│ └── README.md // Example documentation
27+```
28+ 
29+## Example Description
30+ 
31+### 1. Overall Flow
32+ 
33+The overall flow of the example is as follows:
34+ 
35+```text
36+GM(ND) -> L1(Nz) -> L0A(Nz)/L0B(Zn) -> L0C(Nz) -> GM(ND)
37+ │ │ │ │
38+asc_copy_gm2l1_nd2nz asc_copy_l12l0* asc_mmad asc_copy_l0c2gm
39+```
40+ 
41+**Step details**:
42+ 
43+1. **GM -> L1**: Use `asc_set_gm2l1_nz_para` to configure the Nz layout, then use `asc_copy_gm2l1_nd2nz` to convert ND to Nz.
44+2. **L1 -> L0A / L0B**: Use `asc_copy_l12l0a` / `asc_copy_l12l0b` for non-transpose transfer and their `_transpose` variants for transpose transfer.
45+3. **Matrix multiplication**: Use `asc_mmad` to perform matrix multiplication.
46+4. **L0C -> GM**: Use `asc_set_l0c_copy_nz_para` to configure Nz-to-ND conversion, then use `asc_copy_l0c2gm` to transfer results out.
47+ 
48+This example uses a fixed matrix multiplication specification of `[m, n, k] = [40, 50, 70]`, with output C matrix `[40, 50]` in ND format. The kernel function selects different scenarios through `scenario_num`.
49+ 
50+### 2. Key Scenario Differences
51+ 
52+The input shape, transpose flags, and L1 -> L0 transfer methods for each scenario are shown in [Table 1](#table1).
53+ 
54+<a name="table1"></a>
55+<table border="2" align="center">
56+<caption style="font-weight: normal;">
57+ <span style="font-weight: bold; font-size: 1.2em;">Table 1: Meaning of Different scenario_num Values</span></caption>
58+ <tr>
59+ <td><span style="font-weight: bold;">scenario_num</span></td>
60+ <td><span style="font-weight: bold;">Input Data Type</span></td>
61+ <td><span style="font-weight: bold;">Output Data Type</span></td>
62+ <td><span style="font-weight: bold;">Input Shape</span></td>
63+ <td><span style="font-weight: bold;">is_a_transpose</span></td>
64+ <td><span style="font-weight: bold;">is_b_transpose</span></td>
65+ <td><span style="font-weight: bold">Extra Dirty Fractals</span></td>
66+ <td><span style="font-weight: bold;">L1 -> L0 Transfer Method</span></td>
67+ </tr>
68+ <tr>
69+ <td><span style="font-weight: bold;">1</span></td>
70+ <td rowspan="2">int8_t</td>
71+ <td rowspan="2">int32_t</td>
72+ <td>A [40, 70]<br>B [50, 70]</td>
73+ <td>false</td>
74+ <td>true</td>
75+ <td>No</td>
76+ <td><code>asc_copy_l12l0a</code> / <code>asc_copy_l12l0b</code></td>
77+ </tr>
78+ <tr>
79+ <td><span style="font-weight: bold;">2</span></td>
80+ <td>A [70, 40]<br>B [70, 50]</td>
81+ <td>true</td>
82+ <td>false</td>
83+ <td>Yes</td>
84+ <td><code>asc_copy_l12l0a_transpose</code> / <code>asc_copy_l12l0b_transpose</code></td>
85+ </tr>
86+ <tr>
87+ <td><span style="font-weight: bold;">3</span></td>
88+ <td rowspan="2">half</td>
89+ <td rowspan="2">float</td>
90+ <td>A [40, 70]<br>B [50, 70]</td>
91+ <td>false</td>
92+ <td>true</td>
93+ <td>No</td>
94+ <td><code>asc_copy_l12l0a</code> / <code>asc_copy_l12l0b</code></td>
95+ </tr>
96+ <tr>
97+ <td><span style="font-weight: bold;">4</span></td>
98+ <td>A [70, 40]<br>B [70, 50]</td>
99+ <td>true</td>
100+ <td>false</td>
101+ <td>No</td>
102+ <td><code>asc_copy_l12l0a_transpose</code> / <code>asc_copy_l12l0b_transpose</code></td>
103+ </tr>
104+ <tr>
105+ <td><span style="font-weight: bold;">5</span></td>
106+ <td rowspan="2">float</td>
107+ <td rowspan="2">float</td>
108+ <td>A [40, 70]<br>B [50, 70]</td>
109+ <td>false</td>
110+ <td>true</td>
111+ <td>No</td>
112+ <td><code>asc_copy_l12l0a</code> / <code>asc_copy_l12l0b</code></td>
113+ </tr>
114+ <tr>
115+ <td><span style="font-weight: bold;">6</span></td>
116+ <td>A [70, 40]<br>B [70, 50]</td>
117+ <td>true</td>
118+ <td>false</td>
119+ <td>No</td>
120+ <td><code>asc_copy_l12l0a_transpose</code> / <code>asc_copy_l12l0b_transpose</code></td>
121+ </tr>
122+ <tr>
123+ <td><span style="font-weight: bold;">7</span></td>
124+ <td rowspan="1">int8_t</td>
125+ <td rowspan="1">int32_t</td>
126+ <td>A [70, 40]<br>B [70, 50]</td>
127+ <td>true</td>
128+ <td>false</td>
129+ <td>No</td>
130+ <td>for loop + <code>asc_copy_l12l0a_transpose</code>; <code>asc_copy_l12l0b_transpose</code></td>
131+ </tr>
132+</table>
133+ 
134+Key differences by transfer mode:
135+ 
136+**C API non-transpose transfer: Scenarios 1 / 3 / 5**
137+ 
138+- A matrix non-transposed input `[m, k]`, B matrix transposed input `[n, k]`.
139+- L1 -> L0A / L0B both do not need transpose. Call `asc_copy_l12l0a` / `asc_copy_l12l0b`, respectively.
140+ 
141+**C API transpose transfer, single call: Scenarios 2 / 4 / 6**
142+ 
143+- A matrix transposed input `[k, m]`, B matrix non-transposed input `[k, n]`.
144+- L1 -> L0A / L0B both need transpose. Call `asc_copy_l12l0a_transpose` / `asc_copy_l12l0b_transpose`, respectively.
145+- In scenario 2, a single int8_t transpose transfer causes extra dirty data fractals in the m direction. `asc_copy_l0c2gm` transfers out only the valid region with `m_size = m`.
146+ 
147+**C API transpose transfer, for loop call: Scenario 7**
148+ 
149+- Scenario 7 is also int8_t transposed input, same as scenario 2.
150+- A matrix uses a for loop to call `asc_copy_l12l0a_transpose`, skipping dirty data fractals at the tail of the m direction when writing to L0A, so no extra dirty data fractals participate in matrix computation.
151+ 
152+### 3. Parameter Description
153+ 
154+The subsequent code and parameter descriptions repeatedly use fractal and alignment related variables. This section defines these concepts first.
155+ 
156+- `fractal_shape`: Small fractal shape. B8 / B16 / B32 input data types have shape `[16, 32 / sizeof(T)]`, where `T` represents the input data type. Fractal information for data types involved in this example is shown in [Table 2](#table2).
157+- `fractal_size`: Number of elements in 1 small fractal, see [Table 2](#table2).
158+- `fractal_num`: When L1 -> L0A / L0B requires transpose, `asc_copy_l12l0a_transpose` / `asc_copy_l12l0b_transpose` transposes by square matrix. When multiple consecutive small fractals need to merge into a square matrix, `fractal_num` represents the number of small fractals in that square matrix, see [Table 2](#table2).
159+ 
160+<a name="table2"></a>
161+<table border="2" align="center">
162+<caption style="font-weight: normal;">
163+ <span style="font-weight: bold; font-size: 1.2em;">Table 2: Fractal Information for Different Data Types</span></caption>
164+ <tr>
165+ <td></td>
166+ <td align="center"><span style="font-weight: bold;">fractal_shape</span></td>
167+ <td align="center"><span style="font-weight: bold;">fractal_size</span></td>
168+ <td align="center"><span style="font-weight: bold;">fractal_num</span></td>
169+ </tr>
170+ <tr>
171+ <td align="center"><span style="font-weight: bold;">B8</span></td>
172+ <td align="center">[16, 32]</td>
173+ <td align="center">512</td>
174+ <td align="center">2</td>
175+ </tr>
176+ <tr>
177+ <td align="center"><span style="font-weight: bold;">B16</span></td>
178+ <td align="center">[16, 16]</td>
179+ <td align="center">256</td>
180+ <td align="center">1</td>
181+ </tr>
182+ <tr>
183+ <td align="center"><span style="font-weight: bold;">B32</span></td>
184+ <td align="center">[16, 8]</td>
185+ <td align="center">128</td>
186+ <td align="center">2</td>
187+ </tr>
188+</table>
189+ 
190+- `ceil_align`: Ceiling alignment operation. For example, when `m = 30`, `ceil_align(30, 16) = 32`, meaning that the m-axis is aligned to 16 and its aligned length is 32.
191+ 
192+ ```cpp
193+ #define ceil_div(value, align) (((value) + (align) - 1) / (align))
194+ #define ceil_align(value, align) (ceil_div((value), (align)) * (align))
195+ ```
196+ 
197+- `ceil_div`: Ceiling division, generally used to compute loop counts after ceiling alignment.
198+- `m_align_value`: m axis aligns to `m_align_value`. For example, when `m_align_value = 32`, the m axis is aligned to 32. Similarly `n_align_value`, `ka_align_value`, `kb_align_value`.
199+- `m_align_l1` and `m_align_l0`: Aligned values of m axis when A matrix is on L1 and L0A respectively. Similarly `n_align_l1`, `n_align_l0`, `ka_align_l1`, `ka_align_l0`, `kb_align_l1`, `kb_align_l0`.
200+ 
201+### 4. Alignment Requirements
202+ 
203+A / B matrices have different alignment requirements on each axis on L1 and L0. These values are used when configuring `asc_copy_l12l0a` / `asc_copy_l12l0b`, `asc_mmad`, and `asc_copy_l0c2gm` parameters. [Table 3](#table3) and [Table 4](#table4) summarize alignment requirements for A / B matrices on L1 and L0 when scenario_num is 1-6; [Table 5](#table5) and [Table 6](#table6) summarize alignment requirements when scenario_num is 7.
204+ 
205+<a name="table3"></a>
206+<table border="2" align="center">
207+<caption style="font-weight: normal;">
208+ <span style="font-weight: bold; font-size: 1.2em;">Table 3: Alignment Requirements for A and B Matrices on L1 (L1 Layout Format is Nz)</span></caption>
209+ <tr>
210+ <td></td>
211+ <td align="center"><span style="font-weight: bold;">B8 (fractal_num=2)</span></td>
212+ <td align="center"><span style="font-weight: bold;">B16 (fractal_num=1)</span></td>
213+ <td align="center"><span style="font-weight: bold;">B32 (fractal_num=2)</span></td>
214+ </tr>
215+ <tr>
216+ <td rowspan="2" align="center"><span style="font-weight: bold;">A matrix non-transposed input [m, k]</span></td>
217+ <td colspan="3" align="center">m_align_value = fractal_shape[0]</td>
218+ </tr>
219+ <tr>
220+ <td colspan="3" align="center" >ka_align_value = fractal_shape[1]</td>
221+ </tr>
222+ <tr>
223+ <td rowspan="2" align="center"><span style="font-weight: bold;">A matrix transposed input [k, m]</span></td>
224+ <td colspan="2" align="center">ka_align_value = fractal_shape[0] * fractal_num</td>
225+ <td colspan="1" align="center">ka_align_value = fractal_shape[0]</td>
226+ </tr>
227+ <tr>
228+ <td colspan="2" align="center" >m_align_value = fractal_shape[1]</td>
229+ <td colspan="1" align="center" >m_align_value = fractal_shape[1] * fractal_num</td>
230+ </tr>
231+ <tr>
232+ <td rowspan="2" align="center"><span style="font-weight: bold;">B matrix non-transposed input [k, n]</span></td>
233+ <td colspan="2" align="center">kb_align_value = fractal_shape[0] * fractal_num</td>
234+ <td colspan="1" align="center">kb_align_value = fractal_shape[0]</td>
235+ </tr>
236+ <tr>
237+ <td colspan="2" align="center" >n_align_value = fractal_shape[1]</td>
238+ <td colspan="1" align="center" >n_align_value = fractal_shape[1] * fractal_num</td>
239+ </tr>
240+ <tr>
241+ <td rowspan="2" align="center"><span style="font-weight: bold;">B matrix transposed input [n, k]</span></td>
242+ <td colspan="3" align="center">n_align_value = fractal_shape[0]</td>
243+ </tr>
244+ <tr>
245+ <td colspan="3" align="center" >kb_align_value = fractal_shape[1]</td>
246+ </tr>
247+</table>
248+ 
249+<a name="table4"></a>
250+<table border="2" align="center">
251+<caption style="font-weight: normal;">
252+ <span style="font-weight: bold; font-size: 1.2em;">Table 4: Alignment Requirements for A and B Matrices on L0</span></caption>
253+ <tr>
254+ <td></td>
255+ <td align="center"><span style="font-weight: bold;">B8 (fractal_num=2)</span></td>
256+ <td align="center"><span style="font-weight: bold;">B16 (fractal_num=1)</span></td>
257+ <td align="center"><span style="font-weight: bold;">B32 (fractal_num=2)</span></td>
258+ </tr>
259+ <tr>
260+ <td rowspan="2" align="center"><span style="font-weight: bold;">A matrix non-transposed input [m, k], L1 -> L0A no transpose needed</span></td>
261+ <td colspan="3" align="center">m_align_value = fractal_shape[0]</td>
262+ </tr>
263+ <tr>
264+ <td colspan="3" align="center" >ka_align_value = fractal_shape[1]</td>
265+ </tr>
266+ <tr>
267+ <td rowspan="2" align="center"><span style="font-weight: bold;">A matrix transposed input [k, m], L1 -> L0A transpose needed</span></td>
268+ <td colspan="2" align="center">ka_align_value = fractal_shape[1]</td>
269+ <td >ka_align_value = fractal_shape[1] * fractal_num</td>
270+ </tr>
271+ <tr>
272+ <td colspan="2" align="center" >m_align_value = fractal_shape[0] * fractal_num</td>
273+ <td align="center" >m_align_value = fractal_shape[0]</td>
274+ </tr>
275+ <tr>
276+ <td rowspan="2" align="center"><span style="font-weight: bold;">B matrix non-transposed input [k, n], L1 -> L0B transpose needed</span></td>
277+ <td colspan="2" align="center">kb_align_value = fractal_shape[1]</td>
278+ <td align="center">kb_align_value = fractal_shape[1] * fractal_num</td>
279+ </tr>
280+ <tr>
281+ <td colspan="2" align="center">n_align_value = fractal_shape[0] * fractal_num</td>
282+ <td align="center" >n_align_value = fractal_shape[0]</td>
283+ </tr>
284+ <tr>
285+ <td rowspan="2" align="center"><span style="font-weight: bold;">B matrix transposed input [n, k], L1 -> L0B no transpose needed</span></td>
286+ <td colspan="3" align="center">n_align_value = fractal_shape[0]</td>
287+ </tr>
288+ <tr>
289+ <td colspan="3" align="center" >kb_align_value = fractal_shape[1]</td>
290+ </tr>
291+</table>
292+ 
293+When scenario_num=7, A matrix uses a for loop with `asc_copy_l12l0a_transpose` for L1 -> L0A transfer, and L0A only writes the region aligned by valid data fractals.
294+ 
295+<a name="table5"></a>
296+<table border="2" align="center">
297+<caption style="font-weight: normal;">
298+ <span style="font-weight: bold; font-size: 1.2em;">Table 5: Alignment Requirements for A and B Matrices on L1 when scenario_num=7</span></caption>
299+ <tr>
300+ <td align="center" ></td>
301+ <td align="center" ><span style="font-weight: bold;">int8_t (fractal_num=2)</span></td>
302+ </tr>
303+ <tr>
304+ <td rowspan="2"><span style="font-weight: bold;">A matrix transposed input [k, m]</span></td>
305+ <td align="center" >ka_align_value = fractal_shape[0] * fractal_num</td>
306+ </tr>
307+ <tr>
308+ <td align="center" >m_align_value = fractal_shape[1]</td>
309+ </tr>
310+ <tr>
311+ <td rowspan="2"><span style="font-weight: bold;">B matrix non-transposed input [k, n]</span></td>
312+ <td align="center" >kb_align_value = fractal_shape[0] * fractal_num</td>
313+ </tr>
314+ <tr>
315+ <td align="center" >n_align_value = fractal_shape[1]</td>
316+ </tr>
317+</table>
318+ 
319+<a name="table6"></a>
320+<table border="2" align="center">
321+<caption style="font-weight: normal;">
322+ <span style="font-weight: bold; font-size: 1.2em;">Table 6: Alignment Requirements for A and B Matrices on L0 when scenario_num=7</span></caption>
323+ <tr>
324+ <td align="center" ></td>
325+ <td align="center" ><span style="font-weight: bold;">int8_t (fractal_num=2)</span></td>
326+ </tr>
327+ <tr>
328+ <td rowspan="2"><span style="font-weight: bold;">A matrix transposed input [k, m], L1 -> L0A transpose needed</span></td>
329+ <td align="center" >m_align_value = fractal_shape[0]</td>
330+ </tr>
331+ <tr>
332+ <td align="center" >ka_align_value = fractal_shape[1]</td>
333+ </tr>
334+ <tr>
335+ <td rowspan="2"><span style="font-weight: bold;">B matrix non-transposed input [k, n], L1 -> L0B transpose needed</span></td>
336+ <td align="center" >kb_align_value = fractal_shape[1]</td>
337+ </tr>
338+ <tr>
339+ <td align="center" >n_align_value = fractal_shape[0] * fractal_num</td>
340+ </tr>
341+</table>
342+ 
343+### 5. L1 -> L0 C API
344+ 
345+A / B matrices on L1 are all in Nz layout, and on L0A and L0B they are in Nz and Zn layouts respectively. Based on the destination and whether transpose is required, the L1 -> L0 phase calls `asc_copy_l12l0a`, `asc_copy_l12l0b`, `asc_copy_l12l0a_transpose`, or `asc_copy_l12l0b_transpose`.
346+ 
347+When reading this section, first determine the scenario from [Table 1](#table1), then select alignment values from [Table 3](#table3)-[Table 6](#table6). The key C API parameters are:
348+ 
349+- API name: APIs without the `_transpose` suffix perform non-transpose transfer; APIs with the suffix perform transpose transfer.
350+- `m_step`, `k_step`: Represent the number of small fractals covered in row and col directions respectively. During transpose transfer, B8 / B32 need to transfer in square fractal combinations, related step sizes must satisfy `fractal_num` requirements.
351+- `src_stride`, `dst_stride`: Represent intervals between adjacent small fractals in the row direction of the L1 source matrix and L0 destination matrix respectively.
352+- `m_start_position`, `k_start_position`: Represent the starting small fractal position in the L1 source matrix for this transfer. In scenario 7, `m_start_position` is modified in the loop to transfer A matrix in segments.
353+ 
354+#### 5.1. A Matrix L1 -> L0A Non-Transpose
355+ 
356+When L1 -> L0A does not require transpose, the parameter configurations for B8 / B16 / B32 data types are essentially the same, differing only in `fractal_shape`, see [Table 2](#table2). The following figure uses int8_t as an example.
357+ 
358+<div align="center">
359+ <img src="figures/B8_A_l1_l0A_copy_l12l0.png" width="800"><br>
360+ Figure 1: int8_t data type, L1 -> L0A non-transpose, asc_copy_l12l0a data layout diagram
361+</div>
362+ 
363+Parameter configuration key points:
364+ 
365+- `m_step = ceil_div(m_align_l1, fractal_shape[0])`, representing the number of small fractals in the row direction for the m axis.
366+- `k_step = ceil_div(ka_align_l1, fractal_shape[1])`, representing the number of small fractals in the col direction for the k axis.
367+- Call `asc_copy_l12l0a` to perform only the L1 Nz to L0A Nz layout transfer without transpose.
368+ 
369+```cpp
370+m_align_l1 = ceil_align(m, fractal_shape[0]); // 48
371+ka_align_l1 = ceil_align(k, fractal_shape[1]); // 96
372+m_align_l0 = ceil_align(m, fractal_shape[0]); // 48
373+ka_align_l0 = ceil_align(k, fractal_shape[1]); // 96
374+asc_copy_l12l0a(a_l0, a_l1, 0, 0,
375+ ceil_div(m_align_l1, fractal_shape[0]), // m_step = 3
376+ ceil_div(ka_align_l1, fractal_shape[1]), // k_step = 3
377+ ceil_div(m_align_l1, fractal_shape[0]), // src_stride = 3
378+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 3
379+```
380+ 
381+#### 5.2. A Matrix L1 -> L0A Transpose
382+ 
383+When L1 -> L0A requires transpose, call `asc_copy_l12l0a_transpose`. It completes small fractal transpose and L0A Nz layout writing via square fractal combinations; different data types have different square fractal combination methods, described below for B8 / B16 / B32.
384+ 
385+##### 5.2.1. B8 Input Data Type
386+ 
387+B8 input data type fractal is 16 * 32. During L1 -> L0 transpose, 2 fractals of 16 * 32 in the row direction are combined into a 32 * 32 square for transposing. The following uses int8_t as an example to illustrate single call and for loop call methods.
388+ 
389+**Single Call**
390+ 
391+The following figure shows calling `asc_copy_l12l0a_transpose` once to complete L1 -> L0A transfer and transpose:
392+ 
393+<div align="center">
394+ <img src="figures/B8_A_l1_l0A_trans_copy_l12l0.png" width="800"><br>
395+ Figure 2: int8_t data type, L1 -> L0A transpose, single asc_copy_l12l0a_transpose call data layout diagram
396+</div>
397+ 
398+In this example `m = 40`. During B8 transpose transfer, `m_align_l1 = ceil_align(m, fractal_shape[1]) = 64`, `m_align_l1 - m = 24 > 16`. When calling `asc_copy_l12l0a_transpose` once, 1 extra invalid fractal is transferred in the m direction; `asc_mmad` uses `m = ceil_align(m, fractal_shape[0] * fractal_num)` to include that fractal in computation, and `asc_copy_l0c2gm` transfers out only the valid region with `m_size = m`.
399+ 
400+```cpp
401+ka_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 96
402+m_align_l1 = ceil_align(m, fractal_shape[1]); // 64
403+m_align_l0 = ceil_align(m, fractal_shape[0] * fractal_num); // 64
404+ka_align_l0 = ceil_align(k, fractal_shape[1]); // 96
405+asc_copy_l12l0a_transpose(a_l0, a_l1, 0, 0,
406+ ceil_div(ka_align_l1, fractal_shape[0]), // m_step = 6
407+ ceil_div(m_align_l1, fractal_shape[1]), // k_step = 2
408+ ceil_div(ka_align_l1, fractal_shape[0]), // src_stride = 6
409+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 4
410+```
411+ 
412+**For Loop Call**
413+ 
414+The following figure shows calling `asc_copy_l12l0a_transpose` multiple times in a for loop to complete L1 -> L0A transfer and transpose:
415+ 
416+<div align="center">
417+ <img src="figures/B8_A_l1_l0A_trans_for_copy_l12l0.png" width="800"><br>
418+ Figure 3: int8_t data type, L1 -> L0A transpose, for loop calling asc_copy_l12l0a_transpose multiple times
419+</div>
420+ 
421+When using for loop, transfer is done in segments along the L1 row direction (A matrix k axis), transferring 2 fractals in the k axis direction and `ceil_div(m_align_l0, fractal_shape[1])` fractals in the m axis direction each time. `dst_stride` is configured with m direction valid data aligned to `fractal_shape[0]`, skipping dirty data fractals at the tail of m direction when writing to L0A due to transpose over-read, so no extra dirty data fractals participate in `asc_mmad` computation in the m direction.
422+ 
423+```cpp
424+ka_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 96
425+m_align_l1 = ceil_align(m, fractal_shape[1]); // 64
426+m_align_l0 = ceil_align(m, fractal_shape[0]); // 48
427+ka_align_l0 = ceil_align(k, fractal_shape[1]); // 96
428+// Input is int8 type, A matrix [k,m] transposed input, L1 -> L0A requires transpose
429+// The loop calls asc_copy_l12l0a_transpose along the k axis, transferring 2 fractals each time.
430+uint16_t a_loop_num = ceil_div(ka_align_l0, fractal_shape[0] * fractal_num); // 3
431+uint32_t dst_offset = 0;
432+for (uint16_t loopIdx = 0; loopIdx < a_loop_num; ++loopIdx) {
433+ asc_copy_l12l0a_transpose(a_l0 + dst_offset, a_l1,
434+ int8_m_step_align * loopIdx, 0, // m_start_position, k_start_position
435+ int8_m_step_align, // m_step = 2
436+ ceil_div(m_align_l0, fractal_shape[1]), // k_step = 2
437+ ceil_div(ka_align_l1, fractal_shape[0]), // src_stride = 6
438+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 3
439+ dst_offset += ceil_align(m_align_l0, fractal_shape[0]) * fractal_shape[1]; // 1536
440+}
441+```
442+ 
443+##### 5.2.2. B16 Input Data Type
444+ 
445+B16 input data type fractal is 16 * 16, one fractal is already a square. During L1 -> L0 transpose, it transposes by small fractal, and a single `asc_copy_l12l0a_transpose` call completes L1 -> L0A data transfer and transpose. This example uses half as the data type.
446+ 
447+<div align="center">
448+ <img src="figures/B16_A_l1_l0A_trans_copy_l12l0.png" width="800"><br>
449+ Figure 4: half data type, L1 -> L0A transpose, single asc_copy_l12l0a_transpose call data layout diagram
450+</div>
451+ 
452+Parameter configuration key points:
453+ 
454+- `m_step = ceil_div(ka_align_l1, fractal_shape[0])`, representing the number of small fractals in the row direction for the k axis.
455+- `k_step = ceil_div(m_align_l1, fractal_shape[1])`, representing the number of small fractals in the col direction for the m axis.
456+- Call `asc_copy_l12l0a_transpose` to complete transpose when transferring to L0A.
457+ 
458+```cpp
459+ka_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 80
460+m_align_l1 = ceil_align(m, fractal_shape[1]); // 48
461+m_align_l0 = ceil_align(m, fractal_shape[0] * fractal_num); // 48
462+ka_align_l0 = ceil_align(k, fractal_shape[1]); // 80
463+asc_copy_l12l0a_transpose(a_l0, a_l1, 0, 0,
464+ ceil_div(ka_align_l1, fractal_shape[0]), // m_step = 5
465+ ceil_div(m_align_l1, fractal_shape[1]), // k_step = 3
466+ ceil_div(ka_align_l1, fractal_shape[0]), // src_stride = 5
467+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 3
468+```
469+ 
470+##### 5.2.3. B32 Input Data Type
471+ 
472+B32 input data type fractal is 16 * 8. During L1 -> L0 transpose, 2 fractals of 16 * 8 in the col direction are combined into a 16 * 16 square for transposing. This example uses float as the data type.
473+ 
474+<div align="center">
475+ <img src="figures/B32_A_l1_l0A_trans_copy_l12l0.png" width="800"><br>
476+ Figure 5: float data type, L1 -> L0A transpose, single asc_copy_l12l0a_transpose call data layout diagram
477+</div>
478+ 
479+In this example `m = 40`. B32 transpose transfer requires combining 2 fractals in the col direction into a square, **`k_step` must be a multiple of 2**, so L1 col direction (m direction) reads 1 extra invalid fractal. After writing to L0A, there is 1 extra invalid fractal at the tail of the k direction; since L0A is in Nz layout, configuring `k = matrix_k` in `asc_mmad` ensures only valid k axis data participates in computation.
480+ 
481+```cpp
482+ka_align_l1 = ceil_align(k, fractal_shape[0]); // 80
483+m_align_l1 = ceil_align(m, fractal_shape[1] * fractal_num); // 48
484+m_align_l0 = ceil_align(m, fractal_shape[0]); // 48
485+ka_align_l0 = ceil_align(k, fractal_shape[1] * fractal_num); // 80
486+asc_copy_l12l0a_transpose(a_l0, a_l1, 0, 0,
487+ ceil_div(ka_align_l1, fractal_shape[0]), // m_step = 5
488+ ceil_div(m_align_l1, fractal_shape[1]), // k_step = 6
489+ ceil_div(ka_align_l1, fractal_shape[0]), // src_stride = 5
490+ ceil_div(m_align_l0, fractal_shape[0])); // dst_stride = 3
491+```
492+ 
493+#### 5.3. B Matrix L1 -> L0B Non-Transpose
494+ 
495+When L1 -> L0B does not require transpose, the parameter configurations for B8 / B16 / B32 data types are essentially the same, differing only in `fractal_shape`, see [Table 2](#table2). The following figure uses float as an example.
496+ 
497+<div align="center">
498+ <img src="figures/B32_B_l1_l0B_copy_l12l0.png" width="1100"><br>
499+ Figure 6: float data type, L1 -> L0B non-transpose, single asc_copy_l12l0b call data layout diagram
500+</div>
501+ 
502+Parameter configuration key points:
503+ 
504+- `m_step = ceil_div(n_align_l1, fractal_shape[0])`, representing the number of small fractals in the row direction for the n axis.
505+- `k_step = ceil_div(kb_align_l1, fractal_shape[1])`, representing the number of small fractals in the col direction for the k axis.
506+- Call `asc_copy_l12l0b` to perform only the L1 Nz to L0B Zn layout transfer without transpose.
507+ 
508+```cpp
509+n_align_l1 = ceil_align(n, fractal_shape[0]); // 64
510+kb_align_l1 = ceil_align(k, fractal_shape[1]); // 72
511+kb_align_l0 = ceil_align(k, fractal_shape[1]); // 72
512+n_align_l0 = ceil_align(n, fractal_shape[0]); // 64
513+asc_copy_l12l0b(b_l0, b_l1, 0, 0,
514+ ceil_div(n_align_l1, fractal_shape[0]), // m_step = 4
515+ ceil_div(kb_align_l1, fractal_shape[1]), // k_step = 9
516+ ceil_div(n_align_l1, fractal_shape[0]), // src_stride = 4
517+ ceil_div(n_align_l0, fractal_shape[0])); // dst_stride = 4
518+```
519+ 
520+#### 5.4. B Matrix L1 -> L0B Transpose
521+ 
522+When L1 -> L0B requires transpose, call `asc_copy_l12l0b_transpose`. It completes small fractal transpose and L0B Zn layout writing via square fractal combinations; different data types have different square fractal combination methods, described below for B8 / B16 / B32.
523+ 
524+##### 5.4.1. B8 Input Data Type
525+ 
526+B8 input data type fractal is 16 * 32. During L1 -> L0 transpose, 2 fractals of 16 * 32 in the row direction are combined into a 32 * 32 square for transposing. This example uses int8_t as the data type.
527+ 
528+<div align="center">
529+ <img src="figures/B8_B_l1_l0B_trans_copy_l12l0.png" width="1100"><br>
530+ Figure 7: int8_t data type, L1 -> L0B transpose, single asc_copy_l12l0b_transpose call data layout diagram
531+</div>
532+ 
533+In this example `k = 70`. B8 transpose transfer requires combining 2 fractals in the row direction into a square, **`m_step` must be a multiple of 2**, so L1 row direction (k direction) reads 1 extra invalid fractal.
534+ 
535+```cpp
536+kb_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 96
537+n_align_l1 = ceil_align(n, fractal_shape[1]); // 64
538+kb_align_l0 = ceil_align(k, fractal_shape[1]); // 96
539+n_align_l0 = ceil_align(n, fractal_shape[0] * fractal_num); // 64
540+asc_copy_l12l0b_transpose(b_l0, b_l1, 0, 0,
541+ ceil_div(kb_align_l1, fractal_shape[0]), // m_step = 6
542+ ceil_div(n_align_l1, fractal_shape[1]), // k_step = 2
543+ ceil_div(kb_align_l1, fractal_shape[0]), // src_stride = 6
544+ ceil_div(n_align_l0, fractal_shape[0])); // dst_stride = 4
545+```
546+ 
547+##### 5.4.2. B16 Input Data Type
548+ 
549+B16 input data type fractal is 16 * 16, one fractal is already a square. During L1 -> L0 transpose, it transposes by small fractal, and a single `asc_copy_l12l0b_transpose` call completes L1 -> L0B data transfer and transpose. This example uses half as the data type.
550+ 
551+<div align="center">
552+ <img src="figures/B16_B_l1_l0B_trans_copy_l12l0.png" width="800"><br>
553+ Figure 8: half data type, L1 -> L0B transpose, single asc_copy_l12l0b_transpose call data layout diagram
554+</div>
555+ 
556+Parameter configuration key points:
557+ 
558+- `m_step = ceil_div(kb_align_l1, fractal_shape[0])`, representing the number of small fractals in the row direction for the k axis.
559+- `k_step = ceil_div(n_align_l1, fractal_shape[1])`, representing the number of small fractals in the col direction for the n axis.
560+- Call `asc_copy_l12l0b_transpose` to complete transpose when transferring to L0B.
561+ 
562+```cpp
563+kb_align_l1 = ceil_align(k, fractal_shape[0] * fractal_num); // 80
564+n_align_l1 = ceil_align(n, fractal_shape[1]); // 64
565+kb_align_l0 = ceil_align(k, fractal_shape[1]); // 80
566+n_align_l0 = ceil_align(n, fractal_shape[0] * fractal_num); // 64
567+asc_copy_l12l0b_transpose(b_l0, b_l1, 0, 0,
568+ ceil_div(kb_align_l1, fractal_shape[0]), // m_step = 5
569+ ceil_div(n_align_l1, fractal_shape[1]), // k_step = 4
570+ ceil_div(kb_align_l1, fractal_shape[0]), // src_stride = 5
571+ ceil_div(n_align_l0, fractal_shape[0])); // dst_stride = 4
572+```
573+ 
574+##### 5.4.3. B32 Input Data Type
575+ 
576+B32 input data type fractal is 16 * 8. During L1 -> L0 transpose, 2 fractals of 16 * 8 in the col direction are combined into a 16 * 16 square for transposing. This example uses float as the data type.
577+ 
578+<div align="center">
579+ <img src="figures/B32_B_l1_l0B_trans_copy_l12l0.png" width="1100"><br>
580+ Figure 9: float data type, L1 -> L0B transpose, single asc_copy_l12l0b_transpose call data layout diagram
581+</div>
582+ 
583+In this example `n = 50`. B32 transpose transfer requires combining 2 fractals in the col direction into a square, **`k_step` must be a multiple of 2**, so L1 col direction (n direction) reads 1 extra invalid fractal. After writing to L0B, there is 1 extra invalid fractal at the tail of the k direction; since L0B is in Zn layout, configuring `k = matrix_k` in `asc_mmad` ensures only valid k axis data participates in computation.
584+ 
585+```cpp
586+kb_align_l1 = ceil_align(k, fractal_shape[0]); // 80
587+n_align_l1 = ceil_align(n, fractal_shape[1] * fractal_num); // 64
588+kb_align_l0 = ceil_align(k, fractal_shape[1] * fractal_num); // 80
589+n_align_l0 = ceil_align(n, fractal_shape[0]); // 64
590+asc_copy_l12l0b_transpose(b_l0, b_l1, 0, 0,
591+ ceil_div(kb_align_l1, fractal_shape[0]), // m_step = 5
592+ ceil_div(n_align_l1, fractal_shape[1]), // k_step = 8
593+ ceil_div(kb_align_l1, fractal_shape[0]), // src_stride = 5
594+ ceil_div(n_align_l0, fractal_shape[0])); // dst_stride = 4
595+```
596+ 
597+## Build and Run
598+ 
599+Run the following steps in the root directory of this example to build and run the example.
600+- Configure environment variables
601+ Configure environment variables according to the [installation method](../../../../../docs/en/quick_start.md#prepare&install) of the CANN development kit in the current environment.
602+ ```bash
603+ source ${install_path}/cann/set_env.sh
604+ ```
605+ 
606+ > **Note:** `${install_path}` is the CANN package installation directory. When no installation directory is specified, the default installation path is `/usr/local/Ascend`.
607+- Run the example
608+ 
609+ Run the following commands in the example directory.
610+ ```bash
611+ SCENARIO_NUM=1
612+ mkdir -p build && cd build;
613+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO_NUM ..;make -j;
614+ python3 ../scripts/gen_data.py -scenario_num=$SCENARIO_NUM
615+ ./demo
616+ python3 ../scripts/verify_result.py -scenario_num=$SCENARIO_NUM output/output.bin output/golden.bin
617+ ```
618+ To use NPU simulation mode, add the `-DCMAKE_ASC_RUN_MODE=sim` parameter.
619+ 
620+ Examples:
621+ 
622+ ```bash
623+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO_NUM ..;make -j; # NPU simulation mode
624+ ```
625+ > **Notice:** Clear the cmake cache before switching build modes. Run `rm CMakeCache.txt` in the build directory and then re-run cmake.
626+ 
627+- Build option description
628+ 
629+ | Option | Values | Description |
630+ |--------|--------|-------------|
631+ | `CMAKE_ASC_RUN_MODE` | `npu` (default), `sim` | Run mode: NPU execution, NPU simulation |
632+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU architecture: Ascend 950PR/Ascend 950DT |
633+ | `SCENARIO_NUM` | `1`-`7` | Scenario number: different data types and transpose combinations |
634+ 
635+- Execution result
636+ 
637+ The following execution result indicates that the accuracy comparison is successful.
638+ 
639+ ```bash
640+ test pass!
641+ ```
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/data_utils.h+86-0
@@ -0,0 +1,86 @@
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+ * \file data_utils.h
13+ * \brief Data read/write utility functions
14+ */
15+ 
16+#ifndef DATA_UTILS_H
17+#define DATA_UTILS_H
18+#include <fcntl.h>
19+#include <sys/stat.h>
20+#include <unistd.h>
21+#include <fstream>
22+ 
23+#define error_log(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
24+ 
25+bool read_file(const std::string& file_path, size_t& file_size, void* buffer, size_t buffer_size)
26+{
27+ struct stat stat_buffer;
28+ int file_status = stat(file_path.data(), &stat_buffer);
29+ if (file_status == -1) {
30+ error_log("failed to get file");
31+ return false;
32+ }
33+ if (S_ISREG(stat_buffer.st_mode) == 0) {
34+ error_log("%s is not a file, please enter a file", file_path.c_str());
35+ return false;
36+ }
37+ 
38+ std::ifstream file;
39+ file.open(file_path, std::ios::binary);
40+ if (!file.is_open()) {
41+ error_log("Open file failed. path = %s", file_path.c_str());
42+ return false;
43+ }
44+ 
45+ std::filebuf* buf = file.rdbuf();
46+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
47+ if (size == 0) {
48+ error_log("file size is 0");
49+ file.close();
50+ return false;
51+ }
52+ if (size > buffer_size) {
53+ error_log("file size is larger than buffer size");
54+ file.close();
55+ return false;
56+ }
57+ buf->pubseekpos(0, std::ios::in);
58+ buf->sgetn(static_cast<char*>(buffer), size);
59+ file_size = size;
60+ file.close();
61+ return true;
62+}
63+ 
64+bool write_file(const std::string& file_path, const void* buffer, size_t size)
65+{
66+ if (buffer == nullptr) {
67+ error_log("Write file failed. buffer is nullptr");
68+ return false;
69+ }
70+ 
71+ int fd = open(file_path.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE);
72+ if (fd < 0) {
73+ error_log("Open file failed. path = %s", file_path.c_str());
74+ return false;
75+ }
76+ 
77+ size_t write_size = write(fd, buffer, size);
78+ (void)close(fd);
79+ if (write_size != size) {
80+ error_log("Write file Failed.");
81+ return false;
82+ }
83+ 
84+ return true;
85+}
86+#endif // DATA_UTILS_H
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B16_A_l1_l0A_trans_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:9fa393bfa0083e677fc9782665cb73bdb3fdccb6634754fc62c3c870dc7b1e8f
3+size 197594
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B16_B_l1_l0B_trans_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:fc128e78cc92307c4b74b0c8cad2a5ca38367de3df88ac29ac36d2aa8cc15ef1
3+size 202980
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B32_A_l1_l0A_trans_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:f95db214c166fdef321fd138a6dcbaa46c6fe44bb06552a686961e90cc833e99
3+size 323371
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B32_B_l1_l0B_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:6095b86e71fd5109340f80a41466b3b5c97bb2cc596eb0e10c15fbb9a6a92a02
3+size 321604
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B32_B_l1_l0B_trans_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:46d20bc93d586aad3d2b3324bc2982d9907c64007287cc2165496f3ada4c3297
3+size 343225
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B8_A_l1_l0A_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:a5b87a7485c73d1dc1e561e81dbf24cfa9f33c825659e6ea9a75be0a710bbe3c
3+size 125666
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B8_A_l1_l0A_trans_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:ee40c8878e917ed97f0f4590f3179e58a230f216e763f07ce9968d122d44c4de
3+size 209475
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B8_A_l1_l0A_trans_for_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:f691c0d35e072195901ce02d8483e77780989042bedb9a4a5e8747a1cde8a630
3+size 216855
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/figures/B8_B_l1_l0B_trans_copy_l12l0.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:546d49f977bb23f6d88b2efeb7f68be752fc5f2ec4dd4eb777fc2df15089d9ab
3+size 189283
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/load_data_l12l0_950.asc+251-0
@@ -0,0 +1,251 @@
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+ * \file load_data_l12l0_950.asc
13+ * \brief 使用C API将Nz格式矩阵从L1 Buffer搬运到L0A/L0B Buffer，并完成矩阵乘计算。
14+ */
15+ 
16+#include "acl/acl.h"
17+#include "c_api/asc_simd.h"
18+#include "data_utils.h"
19+ 
20+#define ceil_div(value, alignment) (((value) + (alignment) - 1) / (alignment))
21+#define ceil_align(value, alignment) (ceil_div((value), (alignment)) * (alignment))
22+ 
23+constexpr uint32_t scenario_num = SCENARIO_NUM;
24+constexpr uint32_t matrix_m = 40;
25+constexpr uint32_t matrix_n = 50;
26+constexpr uint32_t matrix_k = 70;
27+ 
28+// 场景1、2、7使用B8输入，场景3、4使用B16输入，场景5、6使用B32输入。
29+#if SCENARIO_NUM == 1 || SCENARIO_NUM == 2 || SCENARIO_NUM == 7
30+typedef int8_t input_t;
31+typedef int32_t output_t;
32+#elif SCENARIO_NUM == 3 || SCENARIO_NUM == 4
33+typedef half input_t;
34+typedef float output_t;
35+#else
36+typedef float input_t;
37+typedef float output_t;
38+#endif
39+ 
40+// A转置输入时shape为[k, m]，否则为[m, k]；B转置输入时shape为[n, k]，否则为[k, n]。
41+constexpr bool is_a_transpose = scenario_num == 2 || scenario_num == 4 || scenario_num == 6 || scenario_num == 7;
42+constexpr bool is_b_transpose = !is_a_transpose;
43+ 
44+// B8、B16、B32的小分形shape分别为[16, 32]、[16, 16]、[16, 8]。
45+// 转置搬运时，fractal_num表示拼成一个方型所需的连续小分形个数。
46+constexpr uint32_t fractal_m = 16;
47+constexpr uint32_t c0_size = 32 / sizeof(input_t);
48+constexpr bool is_b16 = sizeof(input_t) == 2;
49+constexpr bool is_b32 = sizeof(input_t) == 4;
50+constexpr uint32_t fractal_num = is_b16 ? 1 : 2;
51+// B8转置搬运的m_step必须是2的倍数。
52+constexpr uint32_t int8_m_step_align = 2;
53+ 
54+// A矩阵在L1上为Nz排布。非转置输入按[m, k]对齐，转置输入按[k, m]对齐。
55+// B32转置时，col方向需要将2个小分形拼成方型；B8转置时，row方向需要拼接2个小分形。
56+constexpr uint32_t m_align_l1 =
57+ is_a_transpose ? ceil_align(matrix_m, c0_size*(is_b32 ? fractal_num : 1)) : ceil_align(matrix_m, fractal_m);
58+constexpr uint32_t ka_align_l1 =
59+ is_a_transpose ? ceil_align(matrix_k, fractal_m*(is_b32 ? 1 : fractal_num)) : ceil_align(matrix_k, c0_size);
60+// 场景2单次转置搬运会多搬一个无效分形，L0A的m轴需要按方型对齐。
61+// 场景7通过for循环跳过该无效分形，因此L0A的m轴只需按fractal_m对齐。
62+constexpr uint32_t m_align_l0_single_call =
63+ is_a_transpose ? ceil_align(matrix_m, fractal_m*(is_b32 ? 1 : fractal_num)) : ceil_align(matrix_m, fractal_m);
64+constexpr uint32_t m_align_l0 = scenario_num == 7 ? ceil_align(matrix_m, fractal_m) : m_align_l0_single_call;
65+constexpr uint32_t ka_align_l0 =
66+ is_a_transpose ? ceil_align(matrix_k, c0_size*(is_b32 ? fractal_num : 1)) : ceil_align(matrix_k, c0_size);
67+ 
68+// B矩阵在L1上为Nz排布，在L0B上为Zn排布。各轴根据是否转置和方型分形要求进行对齐。
69+constexpr uint32_t kb_align_l1 =
70+ is_b_transpose ? ceil_align(matrix_k, c0_size) : ceil_align(matrix_k, fractal_m*(is_b32 ? 1 : fractal_num));
71+constexpr uint32_t n_align_l1 =
72+ is_b_transpose ? ceil_align(matrix_n, fractal_m) : ceil_align(matrix_n, c0_size*(is_b32 ? fractal_num : 1));
73+constexpr uint32_t kb_align_l0 =
74+ is_b_transpose ? ceil_align(matrix_k, c0_size) : ceil_align(matrix_k, c0_size*(is_b32 ? fractal_num : 1));
75+constexpr uint32_t n_align_l0 =
76+ is_b_transpose ? ceil_align(matrix_n, fractal_m) : ceil_align(matrix_n, fractal_m*(is_b32 ? 1 : fractal_num));
77+ 
78+// Buffer数组按各轴对齐后的元素个数申请。
79+constexpr uint32_t a_l1_size = m_align_l1 * ka_align_l1;
80+constexpr uint32_t b_l1_size = kb_align_l1 * n_align_l1;
81+constexpr uint32_t a_l0_size = m_align_l0 * ka_align_l0;
82+constexpr uint32_t b_l0_size = kb_align_l0 * n_align_l0;
83+// 场景2的多搬分形参与asc_mmad计算，因此L0C的m轴按m_align_l0申请；其他场景按有效区域对齐。
84+constexpr uint32_t c_l0_size =
85+ scenario_num == 2 ? m_align_l0 * n_align_l0 : ceil_align(matrix_m, fractal_m) * ceil_align(matrix_n, fractal_m);
86+ 
87+// GM(ND) -> L1(Nz)参数。dst_nz_c0_stride表示Nz矩阵中相邻Z分形的间隔，单位为32字节。
88+constexpr uint16_t gm2l1_matrix_num = 1;
89+constexpr uint16_t gm2l1_dst_nz_n_stride = 1;
90+constexpr uint16_t gm2l1_dst_nz_matrix_stride = 0;
91+constexpr uint64_t gm2l1_loop4_src_stride = 0;
92+constexpr bool gm2l1_enable_small_c0 = false;
93+constexpr uint16_t a_nd_rows = is_a_transpose ? matrix_k : matrix_m;
94+constexpr uint32_t a_nd_cols = is_a_transpose ? matrix_m : matrix_k;
95+constexpr uint16_t a_nz_c0_stride = is_a_transpose ? ka_align_l1 : m_align_l1;
96+constexpr uint16_t b_nd_rows = is_b_transpose ? matrix_n : matrix_k;
97+constexpr uint32_t b_nd_cols = is_b_transpose ? matrix_k : matrix_n;
98+constexpr uint16_t b_nz_c0_stride = is_b_transpose ? n_align_l1 : kb_align_l1;
99+ 
100+// A矩阵L1 -> L0A参数。m_step、k_step表示row、col方向搬运的小分形个数。
101+// src_stride、dst_stride表示源、目的矩阵col方向相邻小分形的间隔，单位为512字节。
102+constexpr uint8_t a_m_step = scenario_num == 7 ?
103+ int8_m_step_align :
104+ (is_a_transpose ? ceil_div(ka_align_l1, fractal_m) : ceil_div(m_align_l1, fractal_m));
105+constexpr uint8_t a_k_step = scenario_num == 7 ?
106+ ceil_div(m_align_l0, c0_size) :
107+ (is_a_transpose ? ceil_div(m_align_l1, c0_size) : ceil_div(ka_align_l1, c0_size));
108+constexpr int16_t a_src_stride = is_a_transpose ? ceil_div(ka_align_l1, fractal_m) : ceil_div(m_align_l1, fractal_m);
109+constexpr uint16_t a_dst_stride = ceil_div(m_align_l0, fractal_m);
110+// 场景7沿L1 row方向（A矩阵k轴）分段搬运，每次搬运int8_m_step_align个小分形。
111+constexpr uint32_t a_loop_num = ceil_div(ka_align_l0, fractal_m* fractal_num);
112+// a_dst_offset为每次循环在L0A上的元素偏移，用于跳过m方向尾部的无效分形。
113+constexpr uint32_t a_dst_offset = ceil_align(m_align_l0, fractal_m) * c0_size;
114+ 
115+// B矩阵L1 -> L0B参数，含义与A矩阵对应参数一致。
116+constexpr uint8_t b_m_step = is_b_transpose ? ceil_div(n_align_l1, fractal_m) : ceil_div(kb_align_l1, fractal_m);
117+constexpr uint8_t b_k_step = is_b_transpose ? ceil_div(kb_align_l1, c0_size) : ceil_div(n_align_l1, c0_size);
118+constexpr int16_t b_src_stride = is_b_transpose ? ceil_div(n_align_l1, fractal_m) : ceil_div(kb_align_l1, fractal_m);
119+constexpr uint16_t b_dst_stride = ceil_div(n_align_l0, fractal_m);
120+ 
121+// 场景2需要让多搬的m方向无效分形参与计算；最终仅通过m_size搬出matrix_m行有效数据。
122+constexpr uint16_t mmad_m = scenario_num == 2 ? ceil_align(matrix_m, fractal_m* fractal_num) : matrix_m;
123+constexpr bool mmad_disable_gemv = false;
124+constexpr bool mmad_c_matrix_source = false;
125+constexpr bool mmad_c_matrix_init_val = true;
126+ 
127+// L0C(Nz) -> GM(ND)参数。仅搬出一个矩阵，相邻矩阵stride设置为0。
128+constexpr uint16_t l0c_matrix_num = 1;
129+constexpr uint16_t l0c_src_nz_matrix_stride = 0;
130+constexpr uint32_t l0c_dst_matrix_stride = 0;
131+constexpr uint16_t l0c_src_stride =
132+ scenario_num == 2 ? ceil_align(matrix_m, fractal_m* fractal_num) : ceil_align(matrix_m, fractal_m);
133+constexpr bool l0c_enable_channel_split = false;
134+constexpr bool l0c_enable_nz2nd = true;
135+constexpr bool l0c_enable_nz2dn = false;
136+constexpr bool l0c_enable_clip_relu_pre = false;
137+ 
138+__global__ __cube__ void load_data_l12l0_950_custom(__gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* c)
139+{
140+ asc_init();
141+ 
142+ __cbuf__ input_t a_l1[a_l1_size];
143+ __cbuf__ input_t b_l1[b_l1_size];
144+ __ca__ input_t a_l0[a_l0_size];
145+ __cb__ input_t b_l0[b_l0_size];
146+ __cc__ output_t c_l0[c_l0_size];
147+ 
148+ // GM(ND) -> L1(Nz)：分别根据A、B的输入shape配置ND矩阵行列数和目的Nz排布。
149+ asc_set_gm2l1_nz_para(gm2l1_matrix_num, gm2l1_dst_nz_n_stride, a_nz_c0_stride, gm2l1_dst_nz_matrix_stride);
150+ asc_copy_gm2l1_nd2nz(
151+ a_l1, (__gm__ input_t*)a, a_nd_cols * sizeof(input_t), asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, a_nd_rows,
152+ a_nd_cols, gm2l1_loop4_src_stride, gm2l1_enable_small_c0);
153+ 
154+ asc_set_gm2l1_nz_para(gm2l1_matrix_num, gm2l1_dst_nz_n_stride, b_nz_c0_stride, gm2l1_dst_nz_matrix_stride);
155+ asc_copy_gm2l1_nd2nz(
156+ b_l1, (__gm__ input_t*)b, b_nd_cols * sizeof(input_t), asc_load_l2_cache_mode::NORMAL_FIRST_VICTIM, b_nd_rows,
157+ b_nd_cols, gm2l1_loop4_src_stride, gm2l1_enable_small_c0);
158+ 
159+ asc_sync_notify(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
160+ asc_sync_wait(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
161+ 
162+ // L1(Nz) -> L0A(Nz)/L0B(Zn).
163+#if SCENARIO_NUM == 7
164+ for (uint32_t loop_idx = 0; loop_idx < a_loop_num; ++loop_idx) {
165+ asc_copy_l12l0a_transpose(
166+ a_l0 + loop_idx * a_dst_offset, a_l1, int8_m_step_align * loop_idx, 0, a_m_step, a_k_step, a_src_stride,
167+ a_dst_stride);
168+ }
169+#elif SCENARIO_NUM == 2 || SCENARIO_NUM == 4 || SCENARIO_NUM == 6
170+ asc_copy_l12l0a_transpose(a_l0, a_l1, 0, 0, a_m_step, a_k_step, a_src_stride, a_dst_stride);
171+#else
172+ asc_copy_l12l0a(a_l0, a_l1, 0, 0, a_m_step, a_k_step, a_src_stride, a_dst_stride);
173+#endif
174+ 
175+#if SCENARIO_NUM == 1 || SCENARIO_NUM == 3 || SCENARIO_NUM == 5
176+ asc_copy_l12l0b(b_l0, b_l1, 0, 0, b_m_step, b_k_step, b_src_stride, b_dst_stride);
177+#else
178+ asc_copy_l12l0b_transpose(b_l0, b_l1, 0, 0, b_m_step, b_k_step, b_src_stride, b_dst_stride);
179+#endif
180+ 
181+ asc_sync_notify(PIPE_MTE1, PIPE_M, EVENT_ID0);
182+ asc_sync_wait(PIPE_MTE1, PIPE_M, EVENT_ID0);
183+ 
184+ asc_mmad(
185+ c_l0, a_l0, b_l0, mmad_m, matrix_k, matrix_n, asc_unit_flag_mode::DISABLE, mmad_disable_gemv,
186+ mmad_c_matrix_source, mmad_c_matrix_init_val);
187+ 
188+ asc_sync_notify(PIPE_M, PIPE_FIX, EVENT_ID0);
189+ asc_sync_wait(PIPE_M, PIPE_FIX, EVENT_ID0);
190+ 
191+ // L0C(Nz) -> GM(ND).
192+ asc_set_l0c_copy_nz_para(l0c_matrix_num, l0c_src_nz_matrix_stride, l0c_dst_matrix_stride);
193+ asc_copy_l0c2gm(
194+ (__gm__ output_t*)c, c_l0, matrix_n, matrix_m, matrix_n, l0c_src_stride,
195+ asc_store_l2_cache_mode::NORMAL_FIRST_VICTIM, asc_unit_flag_mode::DISABLE, asc_quant_mode::NoQuant,
196+ asc_relu_pre_mode::NONE, l0c_enable_channel_split, l0c_enable_nz2nd, l0c_enable_nz2dn,
197+ l0c_enable_clip_relu_pre);
198+ asc_sync_pipe(PIPE_ALL);
199+}
200+ 
201+int32_t main(int32_t argc, char* argv[])
202+{
203+ (void)argc;
204+ (void)argv;
205+ 
206+ size_t a_file_size = matrix_m * matrix_k * sizeof(input_t);
207+ size_t b_file_size = matrix_n * matrix_k * sizeof(input_t);
208+ size_t c_file_size = matrix_m * matrix_n * sizeof(output_t);
209+ 
210+ aclInit(nullptr);
211+ int32_t device_id = 0;
212+ aclrtSetDevice(device_id);
213+ aclrtStream stream = nullptr;
214+ aclrtCreateStream(&stream);
215+ 
216+ uint8_t* a_host = nullptr;
217+ uint8_t* a_device = nullptr;
218+ aclrtMallocHost((void**)(&a_host), a_file_size);
219+ aclrtMalloc((void**)&a_device, a_file_size, ACL_MEM_MALLOC_HUGE_FIRST);
220+ read_file("./input/x1_gm.bin", a_file_size, a_host, a_file_size);
221+ aclrtMemcpy(a_device, a_file_size, a_host, a_file_size, ACL_MEMCPY_HOST_TO_DEVICE);
222+ 
223+ uint8_t* b_host = nullptr;
224+ uint8_t* b_device = nullptr;
225+ aclrtMallocHost((void**)(&b_host), b_file_size);
226+ aclrtMalloc((void**)&b_device, b_file_size, ACL_MEM_MALLOC_HUGE_FIRST);
227+ read_file("./input/x2_gm.bin", b_file_size, b_host, b_file_size);
228+ aclrtMemcpy(b_device, b_file_size, b_host, b_file_size, ACL_MEMCPY_HOST_TO_DEVICE);
229+ 
230+ uint8_t* c_host = nullptr;
231+ uint8_t* c_device = nullptr;
232+ aclrtMallocHost((void**)(&c_host), c_file_size);
233+ aclrtMalloc((void**)&c_device, c_file_size, ACL_MEM_MALLOC_HUGE_FIRST);
234+ 
235+ load_data_l12l0_950_custom<<<1, 0, stream>>>(a_device, b_device, c_device);
236+ aclrtSynchronizeStream(stream);
237+ 
238+ aclrtMemcpy(c_host, c_file_size, c_device, c_file_size, ACL_MEMCPY_DEVICE_TO_HOST);
239+ write_file("./output/output.bin", c_host, c_file_size);
240+ 
241+ aclrtFree(a_device);
242+ aclrtFreeHost(a_host);
243+ aclrtFree(b_device);
244+ aclrtFreeHost(b_host);
245+ aclrtFree(c_device);
246+ aclrtFreeHost(c_host);
247+ aclrtDestroyStream(stream);
248+ aclrtResetDevice(device_id);
249+ aclFinalize();
250+ return 0;
251+}
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/scripts/gen_data.py+99-0
@@ -0,0 +1,99 @@
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+import os
15+import argparse
16+import numpy as np
17+ 
18+np.random.seed(9)
19+ 
20+ 
21+def gen_golden_data(scenario_num=1):
22+ m = 40
23+ n = 50
24+ k = 70
25+ 
26+ x1_gm = None
27+ x2_gm = None
28+ golden = None
29+ 
30+ if scenario_num == 1:
31+ # int8_t, A不转置B转置: A[m,k], B[n,k]
32+ x1_gm = np.random.uniform(1, 10, [m, k]).astype(np.int8)
33+ x2_gm = np.random.uniform(1, 10, [k, n]).astype(np.int8)
34+ golden = (np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32))).astype(
35+ np.int32
36+ )
37+ x2_gm = x2_gm.transpose()
38+ elif scenario_num == 2 or scenario_num == 7:
39+ # int8_t, A转置B不转置: A[k,m], B[k,n]
40+ x1_gm = np.random.uniform(1, 10, [m, k]).astype(np.int8)
41+ x2_gm = np.random.uniform(1, 10, [k, n]).astype(np.int8)
42+ golden = (np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32))).astype(
43+ np.int32
44+ )
45+ x1_gm = x1_gm.transpose()
46+ elif scenario_num == 3:
47+ # half, A不转置B转置: A[m,k], B[n,k]
48+ x1_gm = np.random.uniform(1, 10, [m, k]).astype(np.float16)
49+ x2_gm = np.random.uniform(1, 10, [k, n]).astype(np.float16)
50+ golden = (np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32))).astype(
51+ np.float32
52+ )
53+ x2_gm = x2_gm.transpose()
54+ elif scenario_num == 4:
55+ # half, A转置B不转置: A[k,m], B[k,n]
56+ x1_gm = np.random.uniform(1, 10, [m, k]).astype(np.float16)
57+ x2_gm = np.random.uniform(1, 10, [k, n]).astype(np.float16)
58+ golden = (np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32))).astype(
59+ np.float32
60+ )
61+ x1_gm = x1_gm.transpose()
62+ elif scenario_num == 5:
63+ # float, A不转置B转置: A[m,k], B[n,k]
64+ x1_gm = np.random.uniform(1, 10, [m, k]).astype(np.float32)
65+ x2_gm = np.random.uniform(1, 10, [k, n]).astype(np.float32)
66+ golden = (np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32))).astype(
67+ np.float32
68+ )
69+ x2_gm = x2_gm.transpose()
70+ elif scenario_num == 6:
71+ # float, A转置B不转置: A[k,m], B[k,n]
72+ x1_gm = np.random.uniform(1, 10, [m, k]).astype(np.float32)
73+ x2_gm = np.random.uniform(1, 10, [k, n]).astype(np.float32)
74+ golden = (np.matmul(x1_gm.astype(np.float32), x2_gm.astype(np.float32))).astype(
75+ np.float32
76+ )
77+ x1_gm = x1_gm.transpose()
78+ 
79+ if scenario_num <= 2 or scenario_num == 7:
80+ golden = golden.astype(np.int32)
81+ else:
82+ golden = golden.astype(np.float32)
83+ 
84+ print("x1_gm == ", x1_gm)
85+ print("x2_gm == ", x2_gm)
86+ 
87+ os.system("mkdir -p input")
88+ os.system("mkdir -p output")
89+ 
90+ x1_gm.tofile("./input/x1_gm.bin")
91+ x2_gm.tofile("./input/x2_gm.bin")
92+ golden.tofile("./output/golden.bin")
93+ 
94+ 
95+if __name__ == "__main__":
96+ parser = argparse.ArgumentParser()
97+ parser.add_argument("-scenario_num", type=int, default=1, choices=range(1, 8))
98+ args = parser.parse_args()
99+ gen_golden_data(args.scenario_num)
Aexamples/02_simd_c_api/03_c_api/03_matrix_compute/load_data_l12l0_950/scripts/verify_result.py+67-0
@@ -0,0 +1,67 @@
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+import sys
15+import argparse
16+import numpy as np
17+ 
18+relative_tol = 1e-6
19+absolute_tol = 1e-9
20+error_tol = 1e-4
21+ 
22+ 
23+def verify_result(scenario_num, output, golden):
24+ if scenario_num <= 2 or scenario_num == 7:
25+ output_type = np.int32
26+ else:
27+ output_type = np.float32
28+ output = np.fromfile(output, dtype=output_type).reshape(-1)
29+ golden = np.fromfile(golden, dtype=output_type).reshape(-1)
30+ different_element_results = np.isclose(
31+ output, golden, rtol=relative_tol, atol=absolute_tol, equal_nan=True
32+ )
33+ different_element_indexes = np.where(np.logical_not(different_element_results))[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+ % (
41+ real_index,
42+ golden_data,
43+ output_data,
44+ abs(output_data - golden_data) / golden_data,
45+ )
46+ )
47+ if index == 100:
48+ break
49+ error_ratio = float(different_element_indexes.size) / golden.size
50+ return error_ratio <= error_tol
51+ 
52+ 
53+if __name__ == "__main__":
54+ parser = argparse.ArgumentParser()
55+ parser.add_argument("-scenario_num", type=int, default=1, choices=range(1, 8))
56+ parser.add_argument("output", type=str)
57+ parser.add_argument("golden", type=str)
58+ args = parser.parse_args()
59+ try:
60+ res = verify_result(args.scenario_num, args.output, args.golden)
61+ if not res:
62+ raise ValueError("[ERROR] result error")
63+ else:
64+ print("test pass!")
65+ except Exception as e:
66+ print(e)
67+ sys.exit(1)