Star727
1027
代码
代码
Issues60
Pull Requests204
Actions
项目讨论
Wiki
分析
Star727
1027
已合并
新增基础api mmad_mx 样例 #2336
wx369创建于 5月20日
新增基础api mmad_mx 样例 #2336
已合并
wx369创建于 5月20日
从wx369/asc-devkitmaster合入到cann/asc-devkitmaster
共 13 个文件变更+1102-8
Mdocs/api/SIMD-API/高阶API/量化操作/AscendQuant.md+1-1
@@ -853,7 +853,7 @@ struct AscendQuantParam {
853 AscendC::AscendQuant<srcType>(dstLocal, srcLocal, sharedTmpBuffer, scaleLocal, offsetLocal, scaleCount, offsetCount, calCount);853 AscendC::AscendQuant<srcType>(dstLocal, srcLocal, sharedTmpBuffer, scaleLocal, offsetLocal, scaleCount, offsetCount, calCount);
854 854
855 // dstTensor非固定数据类型855 // dstTensor非固定数据类型
856- AscendC::AscendQuant<srcType>(dstLocal, srcLocal, scaleLocal, offsetLocal, scaleCount, offsetCount, calCount);856+ AscendC::AscendQuant<dstType, srcType>(dstLocal, srcLocal, scaleLocal, offsetLocal, scaleCount, offsetCount, calCount);
857 ```857 ```
858 858 
859 结果示例如下：859 结果示例如下：
Mexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/README.md+1-0
@@ -20,3 +20,4 @@
20| [mmad_unitflag](./mmad_unitflag) | 本样例介绍了如何在调用Mmad指令时，使用unitFlag功能。 | Ascend 950PR/Ascend 950DT<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |20| [mmad_unitflag](./mmad_unitflag) | 本样例介绍了如何在调用Mmad指令时，使用unitFlag功能。 | Ascend 950PR/Ascend 950DT<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |
21| [mmad_gemv](./mmad_gemv) | 本样例介绍了Gemv（M=1）模式下的矩阵乘法 | Ascend 950PR/Ascend 950DT<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |21| [mmad_gemv](./mmad_gemv) | 本样例介绍了Gemv（M=1）模式下的矩阵乘法 | Ascend 950PR/Ascend 950DT<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |
22| [mmad_with_sparse](./mmad_with_sparse) | 本样例介绍基础API MmadWithSparse调用 | Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |22| [mmad_with_sparse](./mmad_with_sparse) | 本样例介绍基础API MmadWithSparse调用 | Atlas A3 训练系列产品/Atlas A3 推理系列产品<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品 |
23+| [mmad_mx](./mmad_mx) | 样例介绍了输入为ND格式，A、B矩阵为FP4（fp4x2_e1m2_t、fp4x2_e2m1_t）和FP8（fp8_e4m3fn_t、fp8_e5m2_t）数据类型，量化矩阵scaleA、scaleB为fp8_e8m0_t数据类型，左矩阵和左量化矩阵不转置、右矩阵和右量化矩阵转置共 4 种包含缩放功能的矩阵乘（简称：MX矩阵乘法） |
A
Aai_xin5月26日

同步修改英文目录

likedislike
Mexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/README_en.md+2-1
@@ -22,4 +22,5 @@ This sample set introduces typical usage of different features of matrix computa
22| [mmad](./mmad) | This sample introduces matrix multiplication with ND format input, B4/B8/B16/B32 input data types (specifically using int4_t/int8_t/bfloat16/float as examples), demonstrating how to implement matrix multiplication computation (C = A x B + Bias) using Mmad instruction | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |22| [mmad](./mmad) | This sample introduces matrix multiplication with ND format input, B4/B8/B16/B32 input data types (specifically using int4_t/int8_t/bfloat16/float as examples), demonstrating how to implement matrix multiplication computation (C = A x B + Bias) using Mmad instruction | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |
23| [mmad_unitflag](./mmad_unitflag) | This sample introduces how to use unitFlag functionality when calling Mmad instruction | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |23| [mmad_unitflag](./mmad_unitflag) | This sample introduces how to use unitFlag functionality when calling Mmad instruction | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |
24| [mmad_gemv](./mmad_gemv) | This sample introduces matrix multiplication in Gemv (M=1) mode | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |24| [mmad_gemv](./mmad_gemv) | This sample introduces matrix multiplication in Gemv (M=1) mode | Ascend 950PR/Ascend 950DT<br>Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |
25-| [mmad_with_sparse](./mmad_with_sparse) | This sample introduces basic API MmadWithSparse invocation | Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |25+| [mmad_with_sparse](./mmad_with_sparse) | This sample introduces basic API MmadWithSparse invocation | Atlas A3 Training Series Products/Atlas A3 Inference Series Products<br>Atlas A2 Training Series Products/Atlas A2 Inference Series Products |
26+| [mmad_mx](./mmad_mx) | This sample introduces 4 types of matrix multiplication with scaling features (referred to as MX matrix multiplication) with ND format input, where A and B matrices are of FP4 (fp4x2_e1m2_t, fp4x2_e2m1_t) and FP8 (fp8_e4m3fn_t, fp8_e5m2_t) data types, and scaling matrices scaleA and scaleB are of fp8_e8m0_t data type. It covers scenarios with non-transposed left matrix and left scaling matrix, and transposed right matrix and right scaling matrix. |
Mexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad/README.md+3-3
@@ -10,7 +10,7 @@
10## 目录结构介绍10## 目录结构介绍
11```11```
12├── mmad12├── mmad
13-│ ├── figures // 本文中的插图文件13+│ ├── figures // 图示
14│ ├── scripts14│ ├── scripts
15│ │ ├── gen_data.py // 输入数据和真值数据生成脚本15│ │ ├── gen_data.py // 输入数据和真值数据生成脚本
16│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本16│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本
@@ -159,12 +159,12 @@
159图2：float类型，A转置，K轴实际对齐与Mmad指令默认要求不一致159图2：float类型，A转置，K轴实际对齐与Mmad指令默认要求不一致
160</p>160</p>
161 161 
162-**场景4 int4b_t输入，int32_t输出，C矩阵初始值为0**162+**场景4 int4b_t输入，int32_t输出，C矩阵初始值为0，仅支持Atlas A3 训练系列产品/Atlas A3 推理系列产品、Atlas A2 训练系列产品/Atlas A2 推理系列产品**
163 163 
164- 输入：A不转置 [30, 70] int4b_t类型，ND格式；B转置 [40, 70] int4b_t类型，ND格式；不带Bias164- 输入：A不转置 [30, 70] int4b_t类型，ND格式；B转置 [40, 70] int4b_t类型，ND格式；不带Bias
165- 输出：C [30, 40] int32_t类型，ND格式165- 输出：C [30, 40] int32_t类型，ND格式
166- 实现：使用`Mmad`实现矩阵乘法运算，通过参数：`mmadParams.cmatrixInitVal = true`，设置C矩阵初始值为0166- 实现：使用`Mmad`实现矩阵乘法运算，通过参数：`mmadParams.cmatrixInitVal = true`，设置C矩阵初始值为0
167-- 说明：该场景仅支持Atlas A3 训练系列产品/Atlas A3 推理系列产品/Atlas A2 训练系列产品/Atlas A2 推理系列产品，且不支持使用传入biasTensor的方式加Bias（场景3）。167+- 说明：该场景不支持使用传入biasTensor的方式加Bias（场景3）。
168 168 
169### 矩阵乘法（Mmad）169### 矩阵乘法（Mmad）
170 170 
Mexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad/mmad.asc+2-2
@@ -89,7 +89,7 @@ public:
89 {89 {
90 AscendC::LocalTensor<U> a1Local(AscendC::TPosition::A1, a1Addr, aSizeAlignL1);90 AscendC::LocalTensor<U> a1Local(AscendC::TPosition::A1, a1Addr, aSizeAlignL1);
91 AscendC::LocalTensor<U> b1Local(AscendC::TPosition::B1, b1Addr, bSizeAlignL1);91 AscendC::LocalTensor<U> b1Local(AscendC::TPosition::B1, b1Addr, bSizeAlignL1);
92- AscendC::LocalTensor<T> bias1Local(AscendC::TPosition::C1, c1Addr, biasSizeAlign);92+ AscendC::LocalTensor<T> bias1Local(AscendC::TPosition::C1, c1Addr, biasSizeAlign / sizeof(T));
W
Wwatermelon_milk5月25日

这里为什么要除sizeof(T)

likedislike
W
wx369
5月25日 评论：
93 93 
94 // GM -> L194 // GM -> L1
95 CopyIn(a1Local, b1Local, bias1Local);95 CopyIn(a1Local, b1Local, bias1Local);
@@ -108,7 +108,7 @@ public:
108 SplitB(b1Local, b2Local);108 SplitB(b1Local, b2Local);
109 }109 }
110 // L1 -> BT110 // L1 -> BT
111- AscendC::LocalTensor<T> bias2Local(AscendC::TPosition::C2, ADDR_0, biasSizeAlign);111+ AscendC::LocalTensor<T> bias2Local(AscendC::TPosition::C2, ADDR_0, biasSizeAlign / sizeof(T));
112 if constexpr (HAS_BIAS_TENSOR) {112 if constexpr (HAS_BIAS_TENSOR) {
113 SplitBias(bias1Local, bias2Local);113 SplitBias(bias1Local, bias2Local);
114 }114 }
Mexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad/scripts/verify_result.py+1-1
@@ -39,7 +39,7 @@ def verify_result(scenario_num, output, golden):
39 real_index = different_element_indexes[index]39 real_index = different_element_indexes[index]
40 golden_data = golden[real_index]40 golden_data = golden[real_index]
41 output_data = output[real_index]41 output_data = output[real_index]
42- if scenario_num > 2:42+ if scenario_num in [2, 3]:
43 print("data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" %43 print("data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" %
44 (real_index, golden_data, output_data, abs(output_data - golden_data) / golden_data))44 (real_index, golden_data, output_data, abs(output_data - golden_data) / golden_data))
45 else:45 else:
Aexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad_mx/CMakeLists.txt+36-0
@@ -0,0 +1,36 @@
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+find_package(ASC REQUIRED)
16+ 
17+project(kernel_samples LANGUAGES ASC CXX)
18+ 
19+add_executable(demo
20+ mmad_mx.asc
21+)
22+ 
23+set(VALID_SCENARIOS "1" "2" "3" "4")
24+if(SCENARIO_NUM IN_LIST VALID_SCENARIOS)
25+ message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}")
26+else()
27+ message(FATAL_ERROR "Invalid SCENARIO_NUM: '${SCENARIO_NUM}'. Allowed values are: 1, 2, 3, 4.")
28+endif()
29+ 
30+target_compile_definitions(demo PRIVATE
31+ SCENARIO_NUM=${SCENARIO_NUM}
32+)
33+ 
34+target_compile_options(demo PRIVATE
35+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
36+)
Aexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad_mx/README.md+206-0
@@ -0,0 +1,206 @@
1+# Mmad MX样例
2+## 概述
3+ 
4+本样例介绍了输入为ND格式，A、B矩阵为FP4（fp4x2_e1m2_t、fp4x2_e2m1_t）和FP8（fp8_e4m3fn_t、fp8_e5m2_t）数据类型，缩放矩阵scaleA、scaleB为fp8_e8m0_t数据类型，左矩阵和左缩放矩阵不转置、右矩阵和右缩放矩阵转置共 4 种包含缩放功能的矩阵乘（简称：MX矩阵乘法）。
5+ 
6+## 支持的产品
7+ 
8+- Ascend 950PR/Ascend 950DT
9+ 
10+## 目录结构介绍
11+```
12+├── mmad_mx
13+│ ├── figures // 图示
14+│ ├── scripts
15+│ │ ├── gen_data.py // 输入数据和真值数据生成脚本
16+│ │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本
17+│ ├── CMakeLists.txt // 编译工程文件
18+│ ├── data_utils.h // 数据读入写出函数
19+│ ├── mmad_mx.asc // Ascend C样例实现 & 调用样例
20+│ └── README.md // 样例说明文档
21+```
22+ 
23+## 样例描述
24+一次完整的MX矩阵乘法涉及的数据搬运过程包括：GM-->L1、L1-->L0A/L0AMX/L0B/L0BMX、L1-->BT（BiasTable Buffer）、L0C-->GM，其中不同存储单元的数据排布格式，如下表1所示：
25+ 
26+<table border="2">
A
Aai_xin5月26日

表格使用markdown简易语法

likedislike
27+<caption>表1：不同存储单元的数据排布格式</caption>
28+ <tr>
29+ <td >存储单元</td>
30+ <td>数据排布格式</td>
31+ </tr>
32+ <tr>
33+ <td>GM</td>
34+ <td>输入A、B矩阵和输出C矩阵为ND排列。</td>
35+ </tr>
36+ <tr>
37+ <td>L1</td>
38+ <td>A、B矩阵为Nz排列。</td>
39+ </tr>
40+ <tr>
41+ <td>L0A、L0AMX</td>
42+ <td>A、ScaleA矩阵分别为Nz、Zz排列。</td>
43+ </tr>
44+ <tr>
45+ <td>L0B、L0BMX</td>
46+ <td>B、ScaleB矩阵分别为Zn、Nn排列。</td>
47+ </tr>
48+ <tr>
49+ <td>BT（BiasTable Buffer）</td>
50+ <td>Bias是shape为[N]的一维Tensor。</td>
51+ </tr>
52+ <tr>
53+ <td>L0C</td>
54+ <td>C矩阵为Nz排列。</td>
55+ </tr>
56+</table>
57+ 
58+MX矩阵乘法计算公式：$$ C = (ScaleA ⊗ A) x (ScaleB ⊗ B) + Bias $$
W
Wwatermelon_milk5月25日

公式前面$$ 是什么意思？

likedislike
W
wx369
5月25日 评论：
59+其中A、ScaleA、B、ScaleB、Bias、C矩阵的需要满足的shape分别为[M,K]、[M,ScaleK]、[K,N]、[ScaleK, N]、[N]和[M,N]，ScaleK = CeilAlign(CeilDiv(K, 32), 2)。
60+ 
61+程序中scenarioNum参数不同取值对应的场景，如下表2所示：
62+ 
63+<table border="2">
64+<caption>表2：scenarioNum不同取值的含义</caption>
65+ <tr>
66+ <td >scenarioNum</td>
67+ <td>A Type</td>
68+ <td>ScaleA Type</td>
69+ <td>B Type</td>
70+ <td>ScaleB Type</td>
71+ <td>Bias Type</td>
72+ <td>C Type</td>
73+ <td>Bias说明</td>
74+ </tr>
75+ <tr>
76+ <td>1</td>
77+ <td>fp4x2_e1m2_t</td>
78+ <td>fp8_e8m0_t</td>
79+ <td>fp4x2_e2m1_t</td>
80+ <td>fp8_e8m0_t</td>
81+ <td>float</td>
82+ <td>float</td>
83+ <td>不带Bias，C矩阵的初始值为0</td>
84+ </tr>
85+ <tr>
86+ <td>2</td>
87+ <td>fp4x2_e2m1_t</td>
88+ <td>fp8_e8m0_t</td>
89+ <td>fp4x2_e1m2_t</td>
90+ <td>fp8_e8m0_t</td>
91+ <td>float</td>
92+ <td>float</td>
93+ <td>带Bias，矩阵计算时不传入biasTensor，C矩阵的初始值来源于C2（BiasTable Buffer），<br>从C2的起始位置0开始读取</td>
94+ </tr>
95+ <tr>
96+ <td>3</td>
97+ <td>fp8_e4m3fn_t</td>
98+ <td>fp8_e8m0_t</td>
99+ <td>fp8_e5m2_t</td>
100+ <td>fp8_e8m0_t</td>
101+ <td>float</td>
102+ <td>float</td>
103+ <td>不带Bias，C矩阵累加来源于CO1初始值</td>
104+ </tr>
105+ <tr>
106+ <td>4</td>
107+ <td>fp8_e5m2_t</td>
108+ <td>fp8_e8m0_t</td>
109+ <td>fp8_e4m3fn_t</td>
110+ <td>fp8_e8m0_t</td>
111+ <td>float</td>
112+ <td>float</td>
113+ <td>带Bias，矩阵计算时传入biasTensor，C矩阵初始值来源于传入的biasTensor，<br>从传入biasTensor的起始位置开始读取</td>
114+ </tr>
115+</table>
116+ 
117+### 场景详细说明
118+ 本样例通过编译参数 `SCENARIO_NUM` 选择不同的场景，所有场景基于相同的矩阵乘规格：[M, N, K] = [40, 50, 70]，核函数名为 `KernelMmadMXCustom`。
119+ 
120+**场景1 C矩阵初始值为0**
121+ 
122+- 输入：A不转置 [40, 70] fp4x2_e1m2_t类型，ND格式；ScaleA不转置 [40, 4] fp8_e8m0_t类型，ND格式；B转置 [50, 70] fp4x2_e2m1_t类型，ND格式；ScaleB转置 [50, 4] fp8_e8m0_t类型，ND格式；不带Bias
123+- 输出：C [40, 50] float类型，ND格式
124+- 实现：使用`Mmad`实现矩阵乘法运算，通过参数：`mmadParams.cmatrixInitVal = true`，设置C矩阵初始值为0
W
Wwatermelon_milk5月25日

新增了显示表示mx使用的mmad接口，可以确认下是否需要更换为最新的接口

likedislike
125+ 
126+**场景2 C矩阵初始值来源于C2**
127+- 输入：A不转置 [40, 70] fp4x2_e2m1_t类型，ND格式；ScaleA不转置 [40, 4] fp8_e8m0_t类型，ND格式；B转置 [50, 70] fp4x2_e1m2_t类型，ND格式；ScaleB转置 [50, 4] fp8_e8m0_t类型，ND格式；带Bias，C矩阵初始值来源于C2
128+- 输出：C [40, 50] float类型，ND格式
129+- 实现：使用`Mmad`实现矩阵乘法运算，不传入biasTensor，通过参数：`mmadParams.cmatrixInitVal = false、mmadParams.cmatrixSource = true`，设置C矩阵初始值来源于C2
130+ 
131+**场景3 C矩阵初始值来源于CO1**
132+- 输入：A不转置 [40, 70] fp8_e4m3fn_t类型，ND格式；ScaleA不转置 [40, 4] fp8_e8m0_t类型，ND格式；B转置 [50, 70] fp8_e5m2_t类型，ND格式；ScaleB转置 [50, 4] fp8_e8m0_t类型，ND格式；不带Bias，C矩阵初始值来源于CO1
133+- 输出：C [40, 50] float类型，ND格式
134+- 实现：使用`Mmad`实现矩阵乘法运算，通过参数：`mmadParams.cmatrixInitVal = false、mmadParams.cmatrixSource = false`，设置C矩阵初始值来源于CO1
135+ 
136+**场景4 C矩阵初始值来源于biasTensor**
137+- 输入：A不转置 [40, 70] fp8_e5m2_t类型，ND格式；ScaleA不转置 [40, 4] fp8_e8m0_t类型，ND格式；B转置 [50, 70] fp8_e4m3fn_t类型，ND格式；ScaleB转置 [50, 4] fp8_e8m0_t类型，ND格式；带Bias，C矩阵初始值来源于传入的biasTensor
138+- 输出：C [40, 50] float类型，ND格式
139+- 实现：使用`Mmad`实现矩阵乘法运算，传入biasTensor，该场景下`mmadParams.cmatrixSource`参数无效
140+ 
141+### 矩阵乘法（Mmad）
142+ 
143+`Mmad`指令会自动完成左、右矩阵与对应scale矩阵的广播乘法，k方向上每32个元素共享一个缩放因子。以A、B矩阵均为FP4类型为例，下图展示了A、ScaleA、B、ScaleB的分型排布格式和缩放功能原理：
144+ 
145+<p align="center">
146+ <img src="figures/mmad-mx.png" width="1000">
WB
Wwatermelon_milk5月25日

图有问题，B和scaleB之前没有x，A和scaleA之间有？

likedislike
Bbluesky9015月25日

我发现这个图最后L0C里的描述有点错误，可能需要调整一下这个png图。L0C里 出现了L0B Layout。

likedislike
147+</p>
148+ 
149+<p align="center">
150+图1：带有缩放功能的矩阵乘示意图
151+</p>
152+ 
153+需要注意的是，mmadParams.k取alignK = CeilAlign(k, 64) = 128，而非原始的k=70，这是因为MX矩阵乘法中Mmad指令要求k方向向64对齐。
154+ 
155+## 编译运行
156+在本样例根目录下执行如下步骤，编译并执行算子。
157+- 配置环境变量
158+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
159+ - 默认路径，root用户安装CANN软件包
160+ ```bash
161+ source /usr/local/Ascend/cann/set_env.sh
162+ ```
163+ 
164+ - 默认路径，非root用户安装CANN软件包
165+ ```bash
166+ source $HOME/Ascend/cann/set_env.sh
167+ ```
168+ 
169+ - 指定路径install_path，安装CANN软件包
170+ ```bash
171+ source ${install_path}/cann/set_env.sh
172+ ```
173+
174+- 样例执行
175+ ```bash
176+ SCENARIO=1
177+ mkdir -p build && cd build; # 创建并进入build目录
178+ cmake -DSCENARIO_NUM=$SCENARIO ..;make -j; # 编译工程，默认npu模式
179+ python3 ../scripts/gen_data.py -scenarioNum=$SCENARIO # 生成测试输入数据
180+ ./demo # 执行编译生成的可执行程序，执行样例
181+ python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确
182+ ```
183+ 
184+ 使用 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
185+ 
186+ 示例如下：
187+ ```bash
188+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DSCENARIO_NUM=$SCENARIO ..;make -j; # CPU仿真模式
189+ cmake -DCMAKE_ASC_RUN_MODE=sim -DSCENARIO_NUM=$SCENARIO ..;make -j; # NPU仿真模式
190+ ```
191+ 
192+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
193+ 
194+- 编译选项说明
195+ 
196+ | 选项 | 可选值 | 说明 |
197+ |------|--------|------|
198+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`sim` 、`cpu` | 运行模式：NPU运行、NPU仿真、CPU仿真 |
199+ | `SCENARIO_NUM` | `1`（默认）、`2`、`3`、`4` | 场景编号，对应场景说明中的4种场景 |
200+ 
201+- 执行结果
202+ 
203+ 执行结果如下，说明精度对比成功。
204+ ```bash
205+ test pass!
206+ ```
Aexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad_mx/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 Data read/write utility functions
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/02_features/03_basic_api/01_matrix_compute/mmad_mx/figures/mmad-mx.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:24464d4659767558d89183c5a4d570a9911a12a3752f286d0db17f2558564b60
3+size 83386
Aexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad_mx/mmad_mx.asc+492-0
@@ -0,0 +1,492 @@
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_mx.asc
14+ * \brief 本文介绍了包含缩放功能的Mmad，以及4种不同的应用场景（包括不同的bias和C矩阵初始化方式）。
15+ */
16+#ifdef ASCENDC_CPU_DEBUG
17+#include "cpu_debug_launch.h"
18+#endif
19+#include "acl/acl.h"
20+#include "kernel_operator.h"
21+#include "data_utils.h"
22+ 
23+constexpr uint32_t scenarioNum = SCENARIO_NUM;
24+constexpr bool HAS_BIAS_TENSOR = (scenarioNum == 2 || scenarioNum == 4);
25+constexpr uint32_t SCALE_BASE_FACTOR = 64;
26+constexpr uint32_t SCALE_EVEN_NUMBER = 2;
27+constexpr uint32_t SCALE_CEIL_NUMBER = 32;
28+constexpr uint32_t FP4_M_STEP_ALIGN = 4;
29+constexpr uint32_t FP8_M_STEP_ALIGN = 2;
30+constexpr uint32_t FP8_C0SIZE = 32;
31+ 
32+template <class T, class TA, class TB, class TAL0Type, class TBL0Type, uint32_t m, uint32_t n, uint32_t k>
33+class KernelMmadMX {
34+public:
35+ __aicore__ inline KernelMmadMX()
36+ {
37+ fractalShape[0] = 16;
38+ if constexpr (AscendC::IsSameType<TA, fp4x2_e1m2_t>::value || AscendC::IsSameType<TA, fp4x2_e2m1_t>::value) {
39+ fractalShape[1] = 64;
40+ fractalNum = 4;
41+ packedK = CeilDiv(k, 2);
42+ } else {
43+ fractalShape[1] = 32;
44+ fractalNum = 2;
45+ packedK = k;
46+ }
47+ scaleK = CeilDiv(k, SCALE_BASE_FACTOR) * SCALE_EVEN_NUMBER;
48+ alignK = CeilAlign(k, SCALE_BASE_FACTOR);
49+ 
50+ // A矩阵：L1上排布格式为Nz
51+ mAlignL1 = CeilAlign(m, fractalShape[0]);
52+ kaAlignL1 = alignK;
53+ aSizeAlignL1 = mAlignL1 * kaAlignL1;
54+ // A矩阵：L0上排布格式为Nz
55+ mAlignL0 = CeilAlign(m, fractalShape[0]);
56+ kaAlignL0 = alignK;
57+ aSizeAlignL0 = mAlignL0 * kaAlignL0;
58+ 
59+ // B矩阵：L1上排布格式为Nz
60+ nAlignL1 = CeilAlign(n, fractalShape[0]);
61+ kbAlignL1 = alignK;
62+ bSizeAlignL1 = nAlignL1 * kbAlignL1;
63+ // B矩阵：L0上排布格式为Zn
64+ kbAlignL0 = alignK;
65+ nAlignL0 = CeilAlign(n, fractalShape[0]);
66+ bSizeAlignL0 = kbAlignL0 * nAlignL0;
67+ 
68+ // scaleA矩阵：L1/L0上排布格式为Zz
69+ scaleMAlignL1 = CeilAlign(m, fractalShape[0]);
70+ scaleASizeAlignL1 = scaleMAlignL1 * scaleK;
71+ 
72+ // scaleB矩阵：L1/L0上排布格式为Nn
73+ scaleNAlignL1 = CeilAlign(n, fractalShape[0]);
74+ scaleBSizeAlignL1 = scaleK * scaleNAlignL1;
75+ 
76+ biasSizeAlign = CeilAlign(CeilDiv(n * sizeof(T), 32), 2) * 32;
77+ cSizeAlignL0 = mAlignL0 * nAlignL0;
78+ 
79+ a1Addr = 0;
80+ scaleA1Addr = 128 * 1024;
B
Bbluesky9015月25日

既然ADDR_0用constexpr定义，下面的咋不用constexpr，建议统一一下，或者全部用数字+注释方式 或者全部用constexpr命名自注释

likedislike
81+ b1Addr = 192 * 1024;
82+ scaleB1Addr = 320 * 1024;
83+ c1Addr = 384 * 1024;
84+ }
85+ __aicore__ inline void Init(GM_ADDR a, GM_ADDR scaleA, GM_ADDR b, GM_ADDR scaleB, GM_ADDR bias, GM_ADDR c)
86+ {
87+ aGM.SetGlobalBuffer((__gm__ TA *)a);
88+ scaleAGM.SetGlobalBuffer((__gm__ fp8_e8m0_t *)scaleA);
89+ bGM.SetGlobalBuffer((__gm__ TB *)b);
90+ scaleBGM.SetGlobalBuffer((__gm__ fp8_e8m0_t *)scaleB);
91+ biasGM.SetGlobalBuffer((__gm__ T *)bias);
92+ cGM.SetGlobalBuffer((__gm__ T *)c);
93+ }
94+ __aicore__ inline void Process()
95+ {
96+ AscendC::LocalTensor<TA> a1Local(AscendC::TPosition::A1, a1Addr, aSizeAlignL1);
97+ AscendC::LocalTensor<fp8_e8m0_t> scaleA1Local(AscendC::TPosition::A1, scaleA1Addr, scaleASizeAlignL1);
98+ AscendC::LocalTensor<TB> b1Local(AscendC::TPosition::B1, b1Addr, bSizeAlignL1);
99+ AscendC::LocalTensor<fp8_e8m0_t> scaleB1Local(AscendC::TPosition::B1, scaleB1Addr, scaleBSizeAlignL1);
100+ AscendC::LocalTensor<T> bias1Local(AscendC::TPosition::C1, c1Addr, biasSizeAlign / sizeof(T));
101+ 
102+ // GM -> L1
103+ CopyInA1(a1Local);
104+ CopyInScaleA1(scaleA1Local);
105+ CopyInB1(b1Local);
106+ CopyInScaleB1(scaleB1Local);
107+ if constexpr (HAS_BIAS_TENSOR) {
108+ CopyInBias(bias1Local);
109+ }
110+ 
111+ AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID0);
112+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE1>(EVENT_ID0);
113+
114+ // L1 -> L0A
115+ AscendC::LocalTensor<TAL0Type> a2Local(AscendC::TPosition::A2, 0, aSizeAlignL0);
116+ SplitA(a1Local, scaleA1Local, a2Local);
117+ 
118+ // L1 -> L0B
119+ AscendC::LocalTensor<TBL0Type> b2Local(AscendC::TPosition::B2, 0, bSizeAlignL0);
120+ SplitB(b1Local, scaleB1Local, b2Local);
121+ 
122+ // L1 -> BT
123+ AscendC::LocalTensor<T> bias2Local(AscendC::TPosition::C2, 0, biasSizeAlign / sizeof(T));
124+ if constexpr (HAS_BIAS_TENSOR) {
125+ SplitBias(bias1Local, bias2Local);
126+ }
127+ 
128+ AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID0);
129+ AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(EVENT_ID0);
130+ 
131+ // Mmad
132+ AscendC::LocalTensor<T> co1Local(AscendC::TPosition::CO1, 0, cSizeAlignL0);
133+ Compute(a2Local, b2Local, bias2Local, co1Local);
134+ 
135+ AscendC::SetFlag<AscendC::HardEvent::M_FIX>(EVENT_ID0);
136+ AscendC::WaitFlag<AscendC::HardEvent::M_FIX>(EVENT_ID0);
137+ 
138+ // L0C -> GM
139+ CopyOut(co1Local);
140+ }
141+ 
142+private:
143+ // 向上整除
144+ __aicore__ inline uint32_t CeilDiv(uint32_t numerator, uint32_t denominator)
145+ {
146+ return (numerator + denominator - 1) / denominator;
147+ }
148+ // 向上对齐
149+ __aicore__ inline uint32_t CeilAlign(uint32_t numerator, uint32_t denominator)
150+ {
151+ return (numerator + denominator - 1) / denominator * denominator;
152+ }
153+ 
154+ __aicore__ inline void CopyInA1(AscendC::LocalTensor<TA> a1Local)
155+ {
156+ AscendC::Nd2NzParams nd2nzA1Params;
157+ nd2nzA1Params.ndNum = 1;
158+ nd2nzA1Params.nValue = m;
159+ nd2nzA1Params.dValue = packedK;
160+ nd2nzA1Params.srcNdMatrixStride = 0;
161+ nd2nzA1Params.srcDValue = packedK;
162+ nd2nzA1Params.dstNzC0Stride = mAlignL1;
163+ nd2nzA1Params.dstNzNStride = 1;
164+ nd2nzA1Params.dstNzMatrixStride = 0;
165+ 
166+ AscendC::DataCopy(a1Local, aGM, nd2nzA1Params);
167+ }
168+ 
169+ __aicore__ inline void CopyInScaleA1(AscendC::LocalTensor<fp8_e8m0_t> scaleA1Local)
170+ {
171+ AscendC::GlobalTensor<half> scaleAGMB16;
172+ scaleAGMB16.SetGlobalBuffer((__gm__ half *)(scaleAGM.GetPhyAddr()), m * scaleK / 2);
173+ auto scaleA1LocalB16 = scaleA1Local.ReinterpretCast<half>();
174+ 
175+ AscendC::Dn2NzParams dn2nzParams;
176+ dn2nzParams.dnNum = 1;
177+ dn2nzParams.dValue = m;
178+ dn2nzParams.nValue = scaleK / 2;
179+ dn2nzParams.srcDnMatrixStride = 0;
180+ dn2nzParams.srcDValue = scaleK / 2;
181+ dn2nzParams.dstNzC0Stride = scaleK / 2;
182+ dn2nzParams.dstNzNStride = 1;
183+ dn2nzParams.dstNzMatrixStride = 0;
184+ 
185+ AscendC::DataCopy(scaleA1LocalB16, scaleAGMB16, dn2nzParams);
186+ }
187+ 
188+ __aicore__ inline void CopyInB1(AscendC::LocalTensor<TB> b1Local)
189+ {
190+ AscendC::Nd2NzParams nd2nzB1Params;
191+ nd2nzB1Params.ndNum = 1;
192+ nd2nzB1Params.nValue = n;
193+ nd2nzB1Params.dValue = packedK;
194+ nd2nzB1Params.srcNdMatrixStride = 0;
195+ nd2nzB1Params.srcDValue = packedK;
196+ nd2nzB1Params.dstNzC0Stride = nAlignL1;
197+ nd2nzB1Params.dstNzNStride = 1;
198+ nd2nzB1Params.dstNzMatrixStride = 0;
199+ 
200+ AscendC::DataCopy(b1Local, bGM, nd2nzB1Params);
201+ }
202+ 
203+ __aicore__ inline void CopyInScaleB1(AscendC::LocalTensor<fp8_e8m0_t> scaleB1Local)
204+ {
205+ AscendC::GlobalTensor<half> scaleBGMB16;
206+ scaleBGMB16.SetGlobalBuffer((__gm__ half *)(scaleBGM.GetPhyAddr()), n * scaleK / 2);
207+ auto scaleB1LocalB16 = scaleB1Local.ReinterpretCast<half>();
208+ 
209+ AscendC::Dn2NzParams dn2nzParams;
210+ dn2nzParams.dnNum = 1;
211+ dn2nzParams.dValue = n;
212+ dn2nzParams.nValue = scaleK / 2;
213+ dn2nzParams.srcDnMatrixStride = 0;
214+ dn2nzParams.srcDValue = scaleK / 2;
215+ dn2nzParams.dstNzC0Stride = scaleK / 2;
216+ dn2nzParams.dstNzNStride = 1;
217+ dn2nzParams.dstNzMatrixStride = 0;
218+ 
219+ AscendC::DataCopy(scaleB1LocalB16, scaleBGMB16, dn2nzParams);
220+ }
221+ 
222+ __aicore__ inline void CopyInBias(AscendC::LocalTensor<T> c1Local)
223+ {
224+ // 由于N可能非32B对齐，“基础数据搬运”无法处理非对齐数据，且“DataCopyPad”不支持GM到L1的搬运，
225+ // 因此，使用“随路转换ND2Nz搬运”实现Bias的GM到L1搬运，确保搬运到L1数据32B对齐。
226+ AscendC::Nd2NzParams nd2nzC1Params;
227+ nd2nzC1Params.ndNum = 1;
228+ nd2nzC1Params.nValue = 1;
229+ nd2nzC1Params.dValue = n;
230+ nd2nzC1Params.srcNdMatrixStride = 0;
231+ nd2nzC1Params.srcDValue = n;
232+ nd2nzC1Params.dstNzC0Stride = 1;
233+ nd2nzC1Params.dstNzNStride = 1;
234+ nd2nzC1Params.dstNzMatrixStride = 0;
235+ 
236+ AscendC::DataCopy(c1Local, biasGM, nd2nzC1Params);
237+ }
238+ 
239+ // A1 -> A2: 调用一次LoadData，同时将A矩阵搬入L0A、scaleA矩阵搬入L0A_MX
240+ // LoadData2DParamsV2控制A矩阵搬运，LoadData2DMxParams控制scaleA矩阵搬运
241+ __aicore__ inline void SplitA(
242+ AscendC::LocalTensor<TA> a1Local,
243+ AscendC::LocalTensor<fp8_e8m0_t> scaleA1Local,
244+ AscendC::LocalTensor<TAL0Type> a2Local)
245+ {
246+ AscendC::LoadData2DParamsV2 loadDataParams;
247+ loadDataParams.sid = 0;
248+ loadDataParams.mStartPosition = 0;
249+ loadDataParams.kStartPosition = 0;
250+ loadDataParams.mStep = CeilDiv(mAlignL1, fractalShape[0]);
251+ loadDataParams.kStep = CeilDiv(kaAlignL1, fractalShape[1]);
252+ loadDataParams.srcStride = CeilDiv(mAlignL1, fractalShape[0]);
253+ loadDataParams.dstStride = CeilDiv(mAlignL0, fractalShape[0]);
254+ loadDataParams.ifTranspose = false;
255+ 
256+ AscendC::LoadData2DMxParams loadMxDataParams;
257+ loadMxDataParams.xStartPosition = 0;
258+ loadMxDataParams.yStartPosition = 0;
259+ loadMxDataParams.xStep = CeilDiv(scaleMAlignL1, fractalShape[0]);
260+ loadMxDataParams.yStep = scaleK / SCALE_EVEN_NUMBER;
261+ loadMxDataParams.srcStride = scaleK / SCALE_EVEN_NUMBER;
262+ loadMxDataParams.dstStride = scaleK / SCALE_EVEN_NUMBER;
263+ 
264+ AscendC::LoadData(a2Local, a1Local, scaleA1Local, loadDataParams, loadMxDataParams);
265+ }
266+ 
267+ // B1 -> B2: 调用一次LoadData，同时将B矩阵搬入L0B、scaleB矩阵搬入L0B_MX
268+ // LoadData2DParamsV2控制B矩阵搬运，LoadData2DMxParams控制scaleB矩阵搬运
269+ __aicore__ inline void SplitB(
270+ AscendC::LocalTensor<TB> b1Local,
271+ AscendC::LocalTensor<fp8_e8m0_t> scaleB1Local,
272+ AscendC::LocalTensor<TBL0Type> b2Local)
273+ {
274+ AscendC::LoadData2DParamsV2 loadDataParams;
275+ loadDataParams.sid = 0;
276+ loadDataParams.mStartPosition = 0;
277+ loadDataParams.kStartPosition = 0;
278+ loadDataParams.mStep = CeilDiv(nAlignL1, fractalShape[0]);
279+ loadDataParams.kStep = CeilDiv(kbAlignL1, fractalShape[1]);
280+ loadDataParams.srcStride = CeilDiv(nAlignL1, fractalShape[0]);
281+ loadDataParams.dstStride = CeilDiv(nAlignL0, fractalShape[0]);
282+ loadDataParams.ifTranspose = false;
283+ 
284+ AscendC::LoadData2DMxParams loadMxDataParams;
285+ loadMxDataParams.xStartPosition = 0;
286+ loadMxDataParams.yStartPosition = 0;
287+ loadMxDataParams.xStep = CeilDiv(scaleNAlignL1, fractalShape[0]);
288+ loadMxDataParams.yStep = scaleK / SCALE_EVEN_NUMBER;
289+ loadMxDataParams.srcStride = scaleK / SCALE_EVEN_NUMBER;
290+ loadMxDataParams.dstStride = scaleK / SCALE_EVEN_NUMBER;
291+ 
292+ AscendC::LoadData(b2Local, b1Local, scaleB1Local, loadDataParams, loadMxDataParams);
293+ }
294+ 
295+ __aicore__ inline void SplitBias(AscendC::LocalTensor<T>& bias1Local, AscendC::LocalTensor<T>& bias2Local)
296+ {
297+ AscendC::DataCopyParams c12c2Params = {1, static_cast<uint16_t>(biasSizeAlign / 32), 0, 0};
298+ AscendC::DataCopy(bias2Local, bias1Local, c12c2Params);
299+ }
300+ 
301+ __aicore__ inline void Compute(AscendC::LocalTensor<TAL0Type> a2Local, AscendC::LocalTensor<TBL0Type> b2Local,
302+ AscendC::LocalTensor<T>& bias2Local, AscendC::LocalTensor<T>& co1Local)
303+ {
304+ AscendC::MmadParams mmadParams;
305+ mmadParams.m = m;
306+ mmadParams.n = n;
307+ mmadParams.k = alignK;
308+ 
309+ if constexpr (SCENARIO_NUM == 1) {
310+ // 不带Bias场景
311+ mmadParams.cmatrixInitVal = true; // C矩阵初始值为0
312+ AscendC::Mmad(co1Local, a2Local, b2Local, mmadParams);
313+ } else if constexpr (SCENARIO_NUM == 2) {
314+ // 带Bias且不传入biasTensor，C矩阵的初始值来源于C2
315+ mmadParams.cmatrixInitVal = false;
316+ mmadParams.cmatrixSource = true;
317+ AscendC::Mmad(co1Local, a2Local, b2Local, mmadParams);
318+ } else if constexpr (SCENARIO_NUM == 3) {
319+ // 不带Bias，C矩阵累加来源于CO1初始值
320+ // 第一次mmad计算，CO1的计算结果作为下一次mmad计算C矩阵的初始值
321+ AscendC::Mmad(co1Local, a2Local, b2Local, mmadParams);
322+ mmadParams.cmatrixInitVal = false;
323+ mmadParams.cmatrixSource = false; // C矩阵的初始值来源于CO1
324+ AscendC::Mmad(co1Local, a2Local, b2Local, mmadParams);
325+ } else {
326+ // 带Bias且传入biasTensor的场景，该场景下cmatrixSource参数无效
327+ mmadParams.cmatrixInitVal = false;
328+ AscendC::Mmad(co1Local, a2Local, b2Local, bias2Local, mmadParams);
329+ }
330+ }
331+ __aicore__ inline void CopyOut(AscendC::LocalTensor<T>& co1Local)
332+ {
333+ AscendC::FixpipeParamsArch3510<AscendC::CO2Layout::ROW_MAJOR> fixpipeParams;
334+ fixpipeParams.mSize = m;
335+ fixpipeParams.nSize = n;
336+ fixpipeParams.srcStride = mAlignL0;
337+ fixpipeParams.dstStride = n;
338+ 
339+ AscendC::Fixpipe<T, T, AscendC::CFG_ROW_MAJOR>(cGM, co1Local, fixpipeParams);
340+ }
341+ 
342+private:
343+ AscendC::GlobalTensor<TA> aGM;
344+ AscendC::GlobalTensor<fp8_e8m0_t> scaleAGM;
345+ AscendC::GlobalTensor<TB> bGM;
346+ AscendC::GlobalTensor<fp8_e8m0_t> scaleBGM;
347+ AscendC::GlobalTensor<T> biasGM;
348+ AscendC::GlobalTensor<T> cGM;
349+ 
350+ uint32_t mAlignL1 = m, kaAlignL1 = k, nAlignL1 = n, kbAlignL1 = k;
351+ uint32_t mAlignL0 = m, kaAlignL0 = k, nAlignL0 = n, kbAlignL0 = k;
352+ uint32_t scaleMAlignL1 = m, scaleNAlignL1 = n;
353+ uint32_t alignK = k;
354+ 
355+ uint32_t aSizeAlignL1, bSizeAlignL1, scaleASizeAlignL1, scaleBSizeAlignL1, biasSizeAlign;
356+ uint32_t aSizeAlignL0, bSizeAlignL0, scaleASizeAlignL0, scaleBSizeAlignL0, cSizeAlignL0;
357+ uint32_t fractalShape[2] = {0, 0};
358+ uint32_t fractalNum = 0;
359+ uint32_t packedK = k;
360+ uint32_t scaleK = k;
361+ 
362+ uint64_t a1Addr, scaleA1Addr, b1Addr, scaleB1Addr, c1Addr;
363+};
364+ 
365+template<typename KernelType>
366+__aicore__ inline void RunKernel(GM_ADDR a, GM_ADDR scaleA, GM_ADDR b, GM_ADDR scaleB, GM_ADDR bias, GM_ADDR c)
367+{
368+ AscendC::InitSocState();
369+ KernelType op;
370+ op.Init(a, scaleA, b, scaleB, bias, c);
371+ op.Process();
372+ AscendC::PipeBarrier<PIPE_ALL>();
373+}
374+ 
375+template <uint32_t m, uint32_t n, uint32_t k>
376+__global__ __cube__ void KernelMmadMXCustom(GM_ADDR a, GM_ADDR scaleA, GM_ADDR b, GM_ADDR scaleB, GM_ADDR bias, GM_ADDR c)
377+{
378+ if constexpr (scenarioNum == 1) {
379+ RunKernel<KernelMmadMX<float, fp4x2_e1m2_t, fp4x2_e2m1_t, fp4x2_e1m2_t, fp4x2_e2m1_t, m, n, k>>(a, scaleA, b, scaleB, bias, c);
380+ } else if constexpr (scenarioNum == 2) {
381+ RunKernel<KernelMmadMX<float, fp4x2_e2m1_t, fp4x2_e1m2_t, fp4x2_e2m1_t, fp4x2_e1m2_t, m, n, k>>(a, scaleA, b, scaleB, bias, c);
382+ } else if constexpr (scenarioNum == 3) {
383+ RunKernel<KernelMmadMX<float, fp8_e4m3fn_t, fp8_e5m2_t, AscendC::mx_fp8_e4m3_t, AscendC::mx_fp8_e5m2_t, m, n, k>>(a, scaleA, b, scaleB, bias, c);
B
Bbluesky9015月25日

AscendC::mx_fp8_e4m3_t, AscendC::mx_fp8_e5m2_t 这两个类型看下能不能封装到kernel里面去，仅在调用mmad指令的时候做一下类型转换，因为后续有规划要新增一个mmad_mx的基础接口。不会使用这几个mx前缀的类型。

likedislike
384+ } else if constexpr (scenarioNum == 4) {
385+ RunKernel<KernelMmadMX<float, fp8_e5m2_t, fp8_e4m3fn_t, AscendC::mx_fp8_e5m2_t, AscendC::mx_fp8_e4m3_t, m, n, k>>(a, scaleA, b, scaleB, bias, c);
386+ }
387+}
388+ 
389+int32_t main(int32_t argc, char *argv[])
390+{
391+ constexpr uint32_t m = 40;
392+ constexpr uint32_t n = 50;
393+ constexpr uint32_t k = 70;
394+ constexpr uint32_t scaleK = (k + SCALE_BASE_FACTOR - 1) / SCALE_BASE_FACTOR * SCALE_EVEN_NUMBER;
395+ 
396+ size_t aFileSize = 0;
397+ size_t scaleAFileSize = 0;
398+ size_t bFileSize = 0;
399+ size_t scaleBFileSize = 0;
400+ size_t biasFileSize = 0;
401+ size_t cFileSize = 0;
402+ 
403+ if constexpr (scenarioNum == 1 || scenarioNum == 2) {
404+ // 场景1和2: A/B 矩阵为 FP4 类型 (两个元素打包成一个字节)，需要除以2
405+ aFileSize = m * k * sizeof(uint8_t) / 2;
406+ bFileSize = n * k * sizeof(uint8_t) / 2;
407+ } else {
408+ // 场景3和4: A/B 矩阵为 FP8 类型，占用一个字节
409+ aFileSize = m * k * sizeof(uint8_t);
410+ bFileSize = n * k * sizeof(uint8_t);
411+ }
412+ scaleAFileSize = m * scaleK * sizeof(uint8_t); // fp8_e8m0_t
413+ scaleBFileSize = n * scaleK * sizeof(uint8_t); // fp8_e8m0_t
414+ biasFileSize = n * sizeof(float);
415+ cFileSize = m * n * sizeof(float);
416+ 
417+ uint32_t numBlocks = 1;
418+ 
419+ aclInit(nullptr);
420+ int32_t deviceId = 0;
421+ aclrtSetDevice(deviceId);
422+ aclrtStream stream = nullptr;
423+ aclrtCreateStream(&stream);
424+ 
425+ uint8_t *aHost;
426+ uint8_t *aDevice;
427+ aclrtMallocHost((void **)(&aHost), aFileSize);
428+ aclrtMalloc((void **)&aDevice, aFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
429+ ReadFile("./input/x1_gm.bin", aFileSize, aHost, aFileSize);
430+ aclrtMemcpy(aDevice, aFileSize, aHost, aFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
431+ 
432+ uint8_t *scaleAHost;
433+ uint8_t *scaleADevice;
434+ aclrtMallocHost((void **)(&scaleAHost), scaleAFileSize);
435+ aclrtMalloc((void **)&scaleADevice, scaleAFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
436+ ReadFile("./input/x1_scale_gm.bin", scaleAFileSize, scaleAHost, scaleAFileSize);
437+ aclrtMemcpy(scaleADevice, scaleAFileSize, scaleAHost, scaleAFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
438+ 
439+ uint8_t *bHost;
440+ uint8_t *bDevice;
441+ aclrtMallocHost((void **)(&bHost), bFileSize);
442+ aclrtMalloc((void **)&bDevice, bFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
443+ ReadFile("./input/x2_gm.bin", bFileSize, bHost, bFileSize);
444+ aclrtMemcpy(bDevice, bFileSize, bHost, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
445+ 
446+ uint8_t *scaleBHost;
447+ uint8_t *scaleBDevice;
448+ aclrtMallocHost((void **)(&scaleBHost), scaleBFileSize);
449+ aclrtMalloc((void **)&scaleBDevice, scaleBFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
450+ ReadFile("./input/x2_scale_gm.bin", scaleBFileSize, scaleBHost, scaleBFileSize);
451+ aclrtMemcpy(scaleBDevice, scaleBFileSize, scaleBHost, scaleBFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
452+ 
453+ uint8_t *biasHost;
454+ uint8_t *biasDevice;
455+ if constexpr (HAS_BIAS_TENSOR) {
456+ aclrtMallocHost((void **)(&biasHost), biasFileSize);
457+ aclrtMalloc((void **)&biasDevice, biasFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
458+ ReadFile("./input/bias_gm.bin", biasFileSize, biasHost, biasFileSize);
459+ aclrtMemcpy(biasDevice, biasFileSize, biasHost, biasFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
460+ }
461+ 
462+ uint8_t *cHost;
463+ uint8_t *cDevice;
464+ aclrtMallocHost((void **)(&cHost), cFileSize);
465+ aclrtMalloc((void **)&cDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
466+ 
467+ KernelMmadMXCustom<m, n, k><<<numBlocks, nullptr, stream>>>(aDevice, scaleADevice, bDevice, scaleBDevice, biasDevice, cDevice);
468+ aclrtSynchronizeStream(stream);
469+ 
470+ aclrtMemcpy(cHost, cFileSize, cDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST);
471+ WriteFile("./output/output.bin", cHost, cFileSize);
472+ 
473+ aclrtFree(aDevice);
474+ aclrtFreeHost(aHost);
475+ aclrtFree(scaleADevice);
476+ aclrtFreeHost(scaleAHost);
477+ aclrtFree(bDevice);
478+ aclrtFreeHost(bHost);
479+ aclrtFree(scaleBDevice);
480+ aclrtFreeHost(scaleBHost);
481+ if constexpr (HAS_BIAS_TENSOR) {
482+ aclrtFree(biasDevice);
483+ aclrtFreeHost(biasHost);
484+ }
485+ aclrtFree(cDevice);
486+ aclrtFreeHost(cHost);
487+ 
488+ aclrtDestroyStream(stream);
489+ aclrtResetDevice(deviceId);
490+ aclFinalize();
491+ return 0;
492+}
Aexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad_mx/scripts/gen_data.py+192-0
@@ -0,0 +1,192 @@
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 sys
16+import numpy as np
17+ 
18+try:
19+ import ml_dtypes
20+ bfloat16 = ml_dtypes.bfloat16
21+ fp8_e4m3 = ml_dtypes.float8_e4m3fn
22+ fp8_e5m2 = ml_dtypes.float8_e5m2
23+except ImportError:
24+ bfloat16 = np.float16
25+ fp8_e4m3 = np.uint8
26+ fp8_e5m2 = np.uint8
27+ 
28+try:
29+ import en_dtypes
30+ fp4_e1m2 = en_dtypes.float4_e1m2
31+ fp4_e2m1 = en_dtypes.float4_e2m1
32+except ImportError:
33+ fp4_e1m2 = np.uint8
34+ fp4_e2m1 = np.uint8
35+ 
36+ 
37+def pack_fp4_to_fp4x2(fp4_data):
38+ row = fp4_data.shape[0]
39+ col = fp4_data.shape[1]
40+ fp4_flat = fp4_data.flatten()
41+ fp4_high = fp4_flat[::2].view(np.uint8)
42+ fp4_low = fp4_flat[1::2].view(np.uint8)
43+ low_bits = (fp4_low & 0x0F) << 4
44+ high_bits = fp4_high & 0x0F
45+ combined = low_bits | high_bits
46+ packed = combined.reshape(row, col // 2)
47+ return packed
48+ 
49+ 
50+def mx_decompress(fp_data, scale_data, block_size=32):
51+ scale_exp = scale_data.astype(np.float32) - 127.0
52+ scale_factor = np.power(2.0, scale_exp)
53+
54+ result = np.zeros(fp_data.shape, dtype=np.float32)
55+
56+ for i in range(fp_data.shape[1]):
57+ block_idx = i // block_size
58+ result[:, i] = fp_data[:, i].astype(np.float32) * scale_factor[:, block_idx]
59+
60+ return result
61+ 
62+ 
63+def mx_decompress_b(fp_data, scale_data, block_size=32):
64+ scale_exp = scale_data.astype(np.float32) - 127.0
65+ scale_factor = np.power(2.0, scale_exp)
66+
67+ result = np.zeros(fp_data.shape, dtype=np.float32)
68+
69+ for row in range(fp_data.shape[0]):
70+ for col in range(fp_data.shape[1]):
71+ block_idx = row // block_size
72+ result[row, col] = fp_data[row, col].astype(np.float32) * scale_factor[block_idx, col]
73+
74+ return result
75+ 
76+ 
77+def layout_scale_a_trans(scale_data):
78+ m, scale_k = scale_data.shape
79+ return scale_data.reshape(m, scale_k // 2, 2).transpose(1, 0, 2).copy()
80+ 
81+ 
82+def layout_scale_b_notrans(scale_data):
83+ scale_k, n = scale_data.shape
84+ return scale_data.reshape(scale_k // 2, 2, n).transpose(0, 2, 1).copy()
85+ 
86+ 
87+def gen_golden_data_fp4(scenario_num, m, n, k):
88+ scale_ceil_number = 32
89+ scale_align_number = 2
90+ scale_k_unaligned = (k + scale_ceil_number - 1) // scale_ceil_number
91+ sk = ((scale_k_unaligned + scale_align_number - 1) // scale_align_number) * scale_align_number
92+ 
93+ if scenario_num == 1:
94+ a_dtype = fp4_e1m2
95+ b_dtype = fp4_e2m1
96+ elif scenario_num == 2:
97+ a_dtype = fp4_e2m1
98+ b_dtype = fp4_e1m2
99+ else:
100+ raise ValueError(f"Invalid scenario_num {scenario_num} for FP4")
101+ 
102+ os.makedirs("input", exist_ok=True)
103+ os.makedirs("output", exist_ok=True)
104+ 
105+ x1_gm = np.random.uniform(-2, 2, [m, k]).astype(a_dtype)
106+ x2_gm = np.random.uniform(-2, 2, [k, n]).astype(b_dtype)
107+
108+ x1_scale_gm = np.random.randint(127, 130, [m, sk]).astype(np.uint8)
109+ x2_scale_gm = np.random.randint(127, 130, [sk, n]).astype(np.uint8)
110+ 
111+ x1_full = mx_decompress(x1_gm.astype(np.float32), x1_scale_gm, 32)
112+ x2_full = mx_decompress_b(x2_gm.astype(np.float32), x2_scale_gm, 32)
113+ 
114+ golden = np.matmul(x1_full.astype(np.float64), x2_full.astype(np.float64)).astype(np.float32)
115+ if scenario_num == 2:
116+ bias_gm = np.random.randint(-10, 10, [n]).astype(np.float32)
117+ golden = golden + bias_gm
118+ bias_gm.tofile("./input/bias_gm.bin")
119+ 
120+ print("B/scaleB transpose")
121+ x2_gm = x2_gm.transpose()
122+ x2_scale_gm = x2_scale_gm.transpose().copy()
123+ 
124+ x1_packed = pack_fp4_to_fp4x2(x1_gm.view(np.uint8))
125+ x2_packed = pack_fp4_to_fp4x2(x2_gm.view(np.uint8))
126+ 
127+ x1_packed.tofile("./input/x1_gm.bin")
128+ x2_packed.tofile("./input/x2_gm.bin")
129+ x1_scale_gm.tofile("./input/x1_scale_gm.bin")
130+ x2_scale_gm.tofile("./input/x2_scale_gm.bin")
131+ golden.tofile("./output/golden.bin")
132+ 
133+ 
134+def gen_golden_data_fp8(scenario_num, m, n, k):
135+ scale_ceil_number = 32
136+ scale_align_number = 2
137+ scale_k_unaligned = (k + scale_ceil_number - 1) // scale_ceil_number
138+ sk = ((scale_k_unaligned + scale_align_number - 1) // scale_align_number) * scale_align_number
139+ 
140+ if scenario_num == 3:
141+ a_dtype = fp8_e4m3
142+ b_dtype = fp8_e5m2
143+ elif scenario_num == 4:
144+ a_dtype = fp8_e5m2
145+ b_dtype = fp8_e4m3
146+ else:
147+ raise ValueError(f"Invalid scenario_num {scenario_num} for FP8")
148+ 
149+ os.makedirs("input", exist_ok=True)
150+ os.makedirs("output", exist_ok=True)
151+ 
152+ x1_gm = np.random.uniform(-10, 10, [m, k]).astype(a_dtype)
153+ x2_gm = np.random.uniform(-10, 10, [k, n]).astype(b_dtype)
154+
155+ x1_scale_gm = np.random.randint(127, 130, [m, sk]).astype(np.uint8)
156+ x2_scale_gm = np.random.randint(127, 130, [sk, n]).astype(np.uint8)
157+ 
158+ x1_full = mx_decompress(x1_gm.astype(np.float32), x1_scale_gm, 32)
159+ x2_full = mx_decompress_b(x2_gm.astype(np.float32), x2_scale_gm, 32)
160+ 
161+ golden = np.matmul(x1_full.astype(np.float64), x2_full.astype(np.float64)).astype(np.float32)
162+ if scenario_num == 3:
163+ golden = np.matmul(x1_full.astype(np.float64), x2_full.astype(np.float64)).astype(np.float32) * 2
164+ else:
165+ bias_gm = np.random.randint(-10, 10, [n]).astype(np.float32)
166+ golden = golden + bias_gm
167+ bias_gm.tofile("./input/bias_gm.bin")
168+ 
169+ print("B/scaleB transpose")
170+ x2_gm = x2_gm.transpose()
171+ x2_scale_gm = x2_scale_gm.transpose().copy()
172+ 
173+ x1_gm.tofile("./input/x1_gm.bin")
174+ x2_gm.tofile("./input/x2_gm.bin")
175+ x1_scale_gm.tofile("./input/x1_scale_gm.bin")
176+ x2_scale_gm.tofile("./input/x2_scale_gm.bin")
177+ golden.tofile("./output/golden.bin")
178+ 
179+ 
180+if __name__ == "__main__":
181+ m, n, k = 40, 50, 70
182+ 
183+ scenario_num = 1
184+ if len(sys.argv) > 1:
185+ scenario_num = int(sys.argv[1].split("=")[1])
186+
187+ if scenario_num in [1, 2]:
188+ gen_golden_data_fp4(scenario_num, m, n, k)
189+ elif scenario_num in [3, 4]:
190+ gen_golden_data_fp8(scenario_num, m, n, k)
191+ else:
192+ raise ValueError(f"Invalid scenario_num {scenario_num}")
Aexamples/01_simd_cpp_api/02_features/03_basic_api/01_matrix_compute/mmad_mx/scripts/verify_result.py+69-0
@@ -0,0 +1,69 @@
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+from pathlib import Path
17+import numpy as np
18+ 
19+ 
20+def read_result_file(file_path, label, expected_size):
21+ path = Path(file_path)
22+ if not path.is_file():
23+ raise FileNotFoundError(
24+ f"{label} file not found: {file_path}. "
25+ f"Current directory: {Path.cwd()}. "
26+ "Please run gen_data.py and ./demo in the build directory before verification."
27+ )
28+ 
29+ data = np.fromfile(path, dtype=np.float32)
30+ if data.size != expected_size:
31+ raise ValueError(f"{label} file size error: expected {expected_size} float32 values, got {data.size}")
32+ return data
33+ 
34+ 
35+def verify_result(output_file, golden_file):
36+ m, n = 40, 50
37+ expected_size = m * n
38+ 
39+ output_data = read_result_file(output_file, "output", expected_size)
40+ golden_data = read_result_file(golden_file, "golden", expected_size)
41+ 
42+ output_data = output_data.reshape(m, n)
43+ golden_data = golden_data.reshape(m, n)
44+ 
45+ diff = np.abs(output_data - golden_data)
46+ max_diff = np.max(diff)
47+ mean_diff = np.mean(diff)
48+ 
49+ threshold = 0.01
50+ if max_diff < threshold:
51+ print("test pass!")
52+ return True
53+ else:
54+ print(f"test failed! max_diff={max_diff}, mean_diff={mean_diff}")
55+ return False
56+ 
57+ 
58+if __name__ == "__main__":
59+ parser = argparse.ArgumentParser()
60+ parser.add_argument("output", type=str, nargs="?", default="output/output.bin")
61+ parser.add_argument("golden", type=str, nargs="?", default="output/golden.bin")
62+ args = parser.parse_args()
63+ 
64+ try:
65+ if not verify_result(args.output, args.golden):
66+ sys.exit(1)
67+ except Exception as e:
68+ print(e)
69+ sys.exit(1)