add best-mxfp4 adv-api-asc-devkit-AtomGit

Mexamples/01_simd_cpp_api/04_best_practices/00_vector_compute_practices/add_high_performance/CMakeLists.txt+0-7

	@@ -37,13 +37,6 @@ target_compile_definitions(demo PRIVATE
37	NPU_ARCH=${NPU_ARCH_ID}	37	NPU_ARCH=${NPU_ARCH_ID}
38	)	38	)
39		39
40	-# ======================================================================================
41	-# NPU 编译选项配置
42	-#
43	-# 说明：
44	-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
45	-# ======================================================================================
46	-
47	target_compile_options(demo PRIVATE	40	target_compile_options(demo PRIVATE
48	$<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>	41	$<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
49	)	42	)

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/CMakeLists.txt+44-0

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

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/README.md+272-0

@@ -0,0 +1,272 @@
	1	+# MxFP4 Matmul 性能调优样例
	2	+
	3	+## 概述
	4	+
	5	+本样例以 MxFP4 矩阵乘法为例，介绍基于 Ascend C `Matmul` 高阶 API 的 MxMatmul 性能调优方法。样例包含两个场景（Case 1-2），均采用常量化 tiling，使用模板常量化 `MatmulApiStaticTiling`（静态 tiling）替代 runtime tiling 拷贝与计算。
	6	+
	7	+优化路径：
	8	+- Case 1: 多核 MDL 常量化 tiling（scaleA/B 与 A/B 同步搬运）
	9	+- Case 2: 多核 MDL 常量化 tiling（在GM→L1搬运中，scaleA/B 相对 A/B 按照倍数多倍搬运，`mxTypePara`）
	10	+
	11	+## 支持的产品
	12	+- Ascend 950PR / Ascend 950DT
	13	+
	14	+## 目录结构介绍
	15	+
	16	+```
	17	+├── matmul_mxfp4_high_performance
	18	+│ ├── scripts
	19	+│ │ ├── gen_data.py // 输入数据和真值数据生成脚本文件
	20	+│ │ └── verify_result.py // 真值对比文件
	21	+│ ├── CMakeLists.txt // 编译工程文件
	22	+│ ├── data_utils.h // 数据读入写出函数
	23	+│ ├── matmul_mx.asc // Ascend C样例实现（包含2个优化case）
	24	+│ ├── matmul_mx.h // 样例头文件（静态Tiling模板与kernel实现）
	25	+```
	26	+
	27	+## 样例描述
	28	+
	29	+### 样例功能
	30	+
	31	+ 样例实现的是固定 shape 为 8192×8192 的 MxFP4 矩阵乘法（带 scale 量化系数输入）。
	32	+
	33	+### MxMatmul 简介
	34	+
	35	+MxMatmul（Matrix Multiply with Scale）是一种带量化缩放系数的矩阵乘法，是 Ascend C Matmul API
	36	+ 在 MX（Mixed-Precision）量化场景下的扩展能力。与基础 Matmul 相比，MxMatmul 的核心区别在于引入了scale 输入。
	37	+
	38	+#### 计算公式
	39	+
	40	+$$
	41	+C = (A \otimes \text{scaleA}) \times (B \otimes \text{scaleB})
	42	+$$
	43	+
	44	+其中 $\otimes$ 表示广播乘（broadcast），左/右矩阵在左/右量化系数矩阵做乘积时，K 方向上每 32 个元素共享一个量化因子。
	45	+
	46	+#### 参数说明
	47	+
	48	+\| 输入 \| 名称 \| 形状 \| 数据类型 \| 数据排布类型 \| 说明 \|
	49	+\|------\|------\|------\|----------\|--------------\|------\|
	50	+\| A \| 左矩阵 \| [8192, 8192] \| `fp4x2_e1m2_t` \| `ND` \| MX FP4 左矩阵 \|
	51	+\| scaleA \| 左量化系数矩阵 \| [8192, 256] \| `fp8_e8m0_t` \| `ND` \| A 矩阵的缩放因子矩阵，K 方向每 32 个元素共享一个缩放因子 \|
	52	+\| B \| 右矩阵 \| [8192, 8192] \| `fp4x2_e1m2_t` \| `ND` \| MX FP4 右矩阵 \|
	53	+\| scaleB \| 右量化系数矩阵 \| [256, 8192] \| `fp8_e8m0_t` \| `ND`（K方向连续2Byte重排） \| B 矩阵的缩放因子矩阵，K 方向每 32 个元素共享一个缩放因子 \|
	54	+\| C \| 输出 \| [8192, 8192] \| `bfloat16_t` \| `ND` \| 计算结果 \|
	55	+
	56	+ <img src="figure/MxMatmul.png">
	W W watermelon_milk4月20日图片显示有问题 cellur_z 4月20日评论： 4月20日评论：
	57	+
	58	+#### 四路输入说明
	59	+
	60	+- 样例中 `sK = ceil(K / 64) * 2`，当 `K=8192` 时，`sK=256`
	61	+- 因此 `scaleA` 形状为 `[M, sK] = [8192, 256]`，`scaleB` 形状为 `[sK, N] = [256, 8192]`
	62	+- 等价含义是：K 方向每 32 个元素共享一个 scale（数据脚本按 `i // 32` 广播）
	63	+- `scale` 的 ND 需要特别说明：`scaleA` 按常规 row-major `[M, sK]` 写入；`scaleB` 的写盘顺序等价于 `[sK/2, N, 2]`，即先在 K 方向连续 `2 Byte`，再沿 N 方向推进，四路输入的 `ND` 排布如下图所示：
	64	+ <img src="figure/NDformat.png">
	65	+
	66	+- 四路输入的搬运如下图所示：
	67	+
	68	+ <img src="figure/InputOfMxMatmul.png">
	69	+
	70	+## 样例实现
	71	+
	72	+### 实现要点
	73	+
	74	+本样例将 tiling 参数统一在 `matmul_mx.h` 中编译期确定，通过模板常量 `CONSTANT_CFG` 传给 `MatmulImpl`：
	75	+
	76	+```cpp
	77	+constexpr static auto CONSTANT_CFG = GetMxConstantCFG<aType, bType, cType, EnableScaleCache>();
	78	+AscendC::Matmul<aType, bType, cType, cType, CONSTANT_CFG,
	79	+ AscendC::MatmulCallBackFunc<nullptr, nullptr, nullptr>,
	80	+ AscendC::Impl::Detail::MatmulWithScalePolicy>
	81	+ matmulObj;
	CaiBinbin4月20日变量名单独一行，是不是格式有问题？ cellur_z 4月20日评论： 4月20日评论：
	82	+REGIST_MATMUL_OBJ(pipe, GetSysWorkSpacePtr(), matmulObj, (TCubeTiling*)nullptr);
	83	+```
	84	+
	85	+说明：
	86	+- kernel 侧不做 `TCubeTiling` 运行时拷贝/计算。
	87	+- `SCENARIO_NUM` 只决定模板实例：`case1 -> MatmulKernel<false>`，`case2 -> MatmulKernel<true>`。
	88	+
	89	+### Case1 与 Case2 差异
	90	+
	91	+两种场景均为常量化 tiling，L1 参数一致：`depthA1/depthB1=4`、`stepKa/stepKb=2`、`stepM/stepN=1`、`dbL0A/dbL0B=2`。
	92	+唯一差异是 `mxTypePara`：
	93	+
	94	+\| 场景 \| `mxTypePara` \| 语义 \|
	95	+\|------\|--------------\|------\|
	96	+\| Case 1 (`SCENARIO_NUM=1`) \| `CASE1_MX_TYPE_PARA = 0x01010101` \| scaleA/B 与 A/B 同步搬运 \|
	97	+\| Case 2 (`SCENARIO_NUM=2`) \| `CASE2_MX_TYPE_PARA = 0x01010404` \| scaleA/B 在 K 方向相对 A/B 多搬运 \|
	98	+
	99	+`mxTypePara` 定义：
	100	+
	101	+- 在MxMatmul中，可以通过设置mxTypePara来控制Scale矩阵和矩阵A、B在L1中加载的比例。
	102	+- MX Scale 缩放因子：`scaleFactorKa=4` 表示 scaleA 数据在 K 方向的加载比例为 A 矩阵的 4 倍；`scaleFactorKb=4` 表示 scaleB 数据在 K 方向的加载比例为 B 矩阵的 4 倍
	103	+
	104	+ - mxTypePara：组合参数，在 MxMatmul 场景使用，表示 scaleA/scaleB 载入 L1 的大小与 A/B 矩阵载入 L1 大小的倍数：
	105	+ - bit [0:6] `scaleFactorKa`：scaleA 与 A 矩阵在 K 方向载入数据量的比例系数，范围 [1, 127]
	106	+ - bit [8:14] `scaleFactorKb`：scaleB 与 B 矩阵在 K 方向载入数据量的比例系数，范围 [1, 127]
	107	+ - bit [16:22] `scaleFactorM`：scaleA 与 A 矩阵在 M 方向载入数据量的比例系数，范围 [1, 127]
	108	+ - bit [24:30] `scaleFactorN`：scaleB 与 B 矩阵在 N 方向载入数据量的比例系数，范围 [1, 127]
	109	+ - 使用约束：
	110	+ - 仅当 Ka 方向全载时（`baseK * stepKa * scaleFactorKa >= singleCoreK`），才能设置 `scaleFactorM > 1`
	111	+ - 仅当 Kb 方向全载时（`baseK * stepKb * scaleFactorKb >= singleCoreK`），才能设置 `scaleFactorN > 1`
	112	+ - scaleA、scaleB 在 M、N、K 方向的载入数据量不能超过实际大小
	113	+ - 该参数仅在 MDL 模式下生效
	114	+
	115	+### 参数设置与搬运数据量计算
	116	+
	117	+以下统计搬运路径为 GM->L1，按当前固定参数计算：
	118	+
	119	+\| 参数 \| 值 \|
	120	+\|------\|----\|
	121	+\| `M=N=K` \| `8192 / 8192 / 8192` \|
	122	+\| `singleCoreM` \| `2048` \|
	123	+\| `singleCoreN` \| `1024` \|
	124	+\| `singleCoreK` \| `8192` \|
	125	+\| `baseM=baseN=baseK` \| `256` \|
	126	+\| `stepKa=stepKb` \| `2 / 2` \|
	127	+\| `scaleFactorKa=scaleFactorKb (case1)` \| `1 / 1` \|
	128	+\| `scaleFactorKa=scaleFactorKb (case2)` \| `4 / 4` \|
	129	+\| 数据类型 \| A/B: `fp4x2` (`0.5 Byte/elem`), scale: `fp8` (`1 Byte/elem`) \|
	130	+
	131	+说明：
	132	+
	133	+A/B 的 base 块大小：`baseM * baseK * 0.5 = 256 * 256 * 0.5 = 32,768 B = 32 KB`, scaleA/scaleB 的 base 块大小：`256 * (256/32) * 1 = 2,048 B = 2 KB` 。
	134	+
	135	+case1 单次 GM→L1 搬运量：
	136	+
	137	+- A：`stepM * stepKa = 1 * 2 = 2` 个 base 块，字节量 `2 * 32 = 64 KB`
	138	+- B：`stepN * stepKb = 1 * 2 = 2` 个 base 块，字节量 `2 * 32 = 64 KB`
	139	+- scaleA：`stepM * stepKa * scaleFactorKa = 1 * 2 * 1 = 2` 个 base 块，`4 KB`
	140	+- scaleB：`stepN * stepKb * scaleFactorKb = 1 * 2 * 1 = 2` 个 base 块，`4 KB`
	141	+- 合计：`64 + 64 + 4 + 4 = 136 KB`
	142	+
	143	+> 说明：`dbL0A/dbL0B=2` 表示 L1→L0 间采用 double buffer（L0 计算当前份的同时，下一份已从 L1 就位），因此 L1 需要同时容纳 `136 × 2 = 272 KB` 数据，但 GM→L1 每次 MTE2 搬运量仍为 `136 KB`。
	144	+
	145	+case2 单次 GM→L1 搬运量：
	146	+- A/B 与 case1 相同：各 `64 KB`（合计 `128 KB`）
	147	+- scaleA：`1 * 2 * 4 = 8` 个 base 块，`16 KB`
	148	+- scaleB：`1 * 2 * 4 = 8` 个 base 块，`16 KB`
	149	+- 合计 MTE2 每次：`64 + 64 + 16 + 16 = 160 KB`
	150	+
	151	+> 同理，由于 `dbL0A/dbL0B=2`，L1 驻留总量为 `160 × 2 = 320 KB`。
	152	+
	153	+Case1/Case2 在 scale 侧的差异主要体现在“单次搬运粒度与搬运次数”：
	154	+
	155	+- Case1：每次少量搬运，scale 搬运次数约为 `16` 次（`8192/512`）
	156	+- Case2：每次多量搬运，scale 搬运次数约为 `4` 次（`8192/2048`）
	157	+- 在该 shape 下两者 scale 理论总字节量相同，但 Case2 的批次更少、复用窗口更大，更利于降低 MTE2 耗时。
	158	+
	159	+
	160	+
	161	+## 性能对比总结
	162	+
	163	+### Ascend 950PR芯片性能数据
	164	+\| Case version \| Task Duration(μs) \| Block Num \| aicore_time(μs) \| aic_mac_time(μs) \| aic_mac_ratio \| aic_scalar_time(μs) \| aic_scalar_ratio \| aic_mte1_time(μs) \| aic_mte1_ratio \| aic_mte2_time(μs) \| aic_mte2_ratio \| aic_fixpipe_time(μs) \| aic_fixpipe_ratio \|
	165	+\|------\|------------------\|-----------\|----------------\|-----------------\|---------------\|-------------------\|-----------------\|------------------\|----------------\|------------------\|----------------\|--------------------\|-------------------\|
	166	+\| Case 1 \| 750.219 \| 32 \| 749.13 \| 660.15 \| 0.881 \| 258.354 \| 0.345 \| 437.64 \| 0.584 \| 753.906 \| 0.982 \| 33.257 \| 0.044 \|
	167	+\| Case 2 \| 693.283 \| 32 \| 692.34 \| 641.444 \| 0.926 \| 241.563 \| 0.349 \| 428.914 \| 0.62 \| 612.536 \| 0.885 \| 33.965 \| 0.049 \|
	168	+
	169	+### Case 2 收益（相对 Case 1）
	170	+
	171	+两种场景均使用常量化 tiling + 模板常量化。Case 2 相对 Case 1 的收益主要来自 `mxTypePara` 带来的 scale 多搬运能力。
	172	+
	173	+- 端到端时延：`750.219 -> 693.283 μs`，减少 `56.936 μs`，收益 `7.59%`。
	174	+- MTE2 绝对耗时：`753.906 -> 612.536 μs`，减少 `141.370 μs`，收益 `18.75%`。
	175	+- MTE2 占比：`0.982 -> 0.885`，下降 `9.7%` 。
	176	+- MAC 占比：`0.881 -> 0.926`，提升 `4.5%` 。
	177	+
	178	+调优 Tips：
	179	+> MX Matmul 的关键差异在于 `scale` 与 A/B 的搬运可解耦；当 `aic_mte2_ratio` 偏高时，优先通过 `mxTypePara` 调整 `scale` 搬运比例以提升 L1 复用、减少重复 GM->L1 搬运。
	180	+
	181	+
	182	+### 理论性能对比
	183	+本样例的性能数据在Ascend 950PR上运行得到，该处理器的主频为1.65GHz，对于MX-FP4的数据类型，每cycle处理16×64×16次乘加运算。Cube理论运算时间为
	184	+$$
	185	+T_{\text{theory}} = \frac{M \times N \times K}{16 \times 64 \times 16 \times 1.65 \times 10^9 \times \text{核数}} = \frac{8192 \times 8192 \times 8192}{4096 \times 1.65 \times 10^9 \times 32} = 635.5 μs
	186	+$$
	187	+Case 1/Case 2 的 `aic_mac_time` 分别为 `660.150 μs` / `641.444 μs`，相对理论值 `635.5 μs`：
	188	+- Case 1 误差：`(660.150 - 635.5) / 635.5 = 3.88%`
	189	+- Case 2 误差：`(641.444 - 635.5) / 635.5 = 0.94%`
	190	+
	191	+可以看到，Case 2 已达到理论性能峰值的 `99.06%`。
	192	+
	193	+## 编译运行
	194	+
	195	+### 编译执行
	196	+
	197	+在本样例根目录下执行如下步骤，编译并执行样例：
	198	+
	199	+- 配置环境变量
	200	+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
	201	+ - 默认路径，root用户安装CANN软件包
	202	+ ```bash
	203	+ source /usr/local/Ascend/cann/set_env.sh
	204	+ ```
	205	+
	206	+ - 默认路径，非root用户安装CANN软件包
	207	+ ```bash
	208	+ source $HOME/Ascend/cann/set_env.sh
	209	+ ```
	210	+
	211	+ - 指定路径install_path，安装CANN软件包
	212	+ ```bash
	213	+ source ${install_path}/cann/set_env.sh
	214	+ ```
	215	+
	216	+- 样例执行
	217	+ ```bash
	218	+ SCENARIO_NUM=2
	219	+ mkdir -p build && cd build; # 创建并进入 build 目录
	220	+ cmake .. -DSCENARIO_NUM=$SCENARIO_NUM -DCMAKE_ASC_RUN_MODE=npu -DCMAKE_ASC_ARCHITECTURES=dav-3510; make -j;
	221	+ python3 ../scripts/gen_data.py
	222	+ ./demo
	223	+ python3 ../scripts/verify_result.py ./output/output.bin ./output/golden.bin
	224	+ ```
	225	+
	226	+ 使用NPU仿真模式时，设置 `-DCMAKE_ASC_RUN_MODE=sim` 即可
	227	+ ```bash
	228	+ cmake .. -DSCENARIO_NUM=$SCENARIO_NUM -DCMAKE_ASC_RUN_MODE=npu -DCMAKE_ASC_ARCHITECTURES=dav-3510; make -j; # npu模式
	229	+ cmake .. -DSCENARIO_NUM=$SCENARIO_NUM -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510; make -j; # npu仿真模式
	230	+ ```
	231	+
	232	+ 编译选项说明：
	233	+
	234	+ \| 参数 \| 可选值 \| 说明 \|
	235	+ \|------\|--------\|------\|
	236	+ \| `SCENARIO_NUM` \| `1` / `2` \| 1: 常量化tiling + scale同步搬运；2: 常量化tiling + scale多搬运 \|
	237	+ \| `CMAKE_ASC_RUN_MODE` \| `npu`（默认）/ `sim` \| 运行模式：NPU运行、NPU仿真 \|
	238	+ \| `CMAKE_ASC_ARCHITECTURES` \| `dav-3510` \| 目标SoC架构（本样例仅支持3510） \|
	239	+
	240	+ > 注意：切换 `CMAKE_ASC_RUN_MODE` / `CMAKE_ASC_ARCHITECTURES` / `SCENARIO_NUM` 前需清理cmake缓存，可在build目录下执行 `rm CMakeCache.txt` 后重新 cmake。
	241	+
	242	+
	243	+ 执行结果如下，说明精度对比成功。
	244	+ ```bash
	245	+ test pass!
	246	+ ```
	247	+
	248	+### 性能分析
	249	+
	250	+使用 `msprof` 工具获取详细性能数据：
	251	+
	252	+```bash
	253	+msprof ./demo # 分析性能
	254	+```
	255	+
	256	+当前目录下会生成PROF_前缀的文件夹，`mindstudio_profiler_output`目录保存Host和各个Device的性能数据汇总，性能数据分析推荐查看该目录下文件
	257	+
	258	+```bash
	259	+PROF_xxxx_XXXXXX
	260	+├── device_{id}
	261	+└── host
	262	+└── mindstudio_profiler_log
	263	+└── mindstudio_profiler_output # 保存Host和各个Device的性能数据汇总
	264	+ ├── msprof_*.json
	265	+ ├── xx_*.csv
	266	+ └── README.txt
	267	+```
	268	+查看具体的性能分析结果：
	269	+```
	270	+# 查看Task Duration 以及各项数据
	271	+cat ./PROF_/mindstudio_profiler_output/op_summary_.csv
	272	+```

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/data_utils.h+93-0

@@ -0,0 +1,93 @@
	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 Utility functions for reading and writing binary files
	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 ReadFile(const std::string& filePath, size_t& fileSize, void* buffer, size_t bufferSize)
	26	+{
	27	+ struct stat sBuf;
	28	+ int fileStatus = stat(filePath.data(), &sBuf);
	29	+ if (fileStatus == -1) {
	30	+ ERROR_LOG("failed to get file");
	31	+ return false;
	32	+ }
	33	+ if (S_ISREG(sBuf.st_mode) == 0) {
	34	+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
	35	+ return false;
	36	+ }
	37	+
	38	+ std::ifstream file;
	39	+ file.open(filePath, std::ios::binary);
	40	+ if (!file.is_open()) {
	41	+ ERROR_LOG("Open file failed. path = %s", filePath.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 > bufferSize) {
	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	+ fileSize = size;
	60	+ file.close();
	61	+ return true;
	62	+}
	63	+
	64	+/**
	65	+ * @brief Write data to file
	66	+ * @param [in] filePath: file path
	67	+ * @param [in] buffer: data to write to file
	68	+ * @param [in] size: size to write
	69	+ * @return write result
	70	+ */
	71	+bool WriteFile(const std::string& filePath, const void* buffer, size_t size)
	72	+{
	73	+ if (buffer == nullptr) {
	74	+ ERROR_LOG("Write file failed. buffer is nullptr");
	75	+ return false;
	76	+ }
	77	+
	78	+ int fd = open(filePath.c_str(), O_RDWR \| O_CREAT \| O_TRUNC, S_IRUSR \| S_IWRITE);
	79	+ if (fd < 0) {
	80	+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
	81	+ return false;
	82	+ }
	83	+
	84	+ size_t writeSize = write(fd, buffer, size);
	85	+ (void)close(fd);
	86	+ if (writeSize != size) {
	87	+ ERROR_LOG("Write file Failed.");
	88	+ return false;
	89	+ }
	90	+
	91	+ return true;
	92	+}
	93	+#endif // DATA_UTILS_H

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/figure/InputOfMxMatmul.png+3-0

@@ -0,0 +1,3 @@
	1	+version https://git-lfs.github.com/spec/v1
	2	+oid sha256:81ba9d71949f7c4e44ef5c6a3eef943939d8cc5c4664a85527ee1066f68e35b6
	3	+size 36705

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/figure/MxMatmul.png+3-0

@@ -0,0 +1,3 @@
	1	+version https://git-lfs.github.com/spec/v1
	2	+oid sha256:b68b428c3d3d945db29ccb39a3fb4e5a1802956cd41035343a5c8a38b5c858cd
	3	+size 54842

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/figure/NDformat.png+3-0

@@ -0,0 +1,3 @@
	1	+version https://git-lfs.github.com/spec/v1
	2	+oid sha256:2b33a31b6f64a9dd69d7080601c67f1242187f377afc9d93887e22bdfc49257d
	3	+size 71964

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/matmul_mx.asc+157-0

@@ -0,0 +1,157 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under 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 matmul.asc
	13	+ * \brief Matmul_mx sample with case1~case2
	14	+ */
	15	+
	16	+#include "matmul_mx.h"
	17	+#include "data_utils.h"
	18	+#include "tiling/platform/platform_ascendc.h"
	19	+#include "acl/acl.h"
	20	+
	21	+#include <iostream>
	22	+
	23	+#ifndef SCENARIO_NUM
	24	+#define SCENARIO_NUM 1
	25	+#endif
	26	+
	27	+template <bool EnableScaleCache>
	28	+__global__ __cube__ void matmul_custom_mdl_static(
	29	+ GM_ADDR a, GM_ADDR b, GM_ADDR as, GM_ADDR bs, GM_ADDR c, GM_ADDR workspace)
	30	+{
	31	+ (void)workspace;
	32	+ AscendC::TPipe pipe;
	33	+ MatmulKernel<EnableScaleCache> matmulKernel;
	34	+ matmulKernel.Init(a, b, as, bs, c);
	35	+ matmulKernel.Process(&pipe);
	36	+}
	37	+
	38	+int32_t main(int32_t argc, char* argv[])
	39	+{
	40	+ constexpr uint32_t kM = cfg::M;
	41	+ constexpr uint32_t kN = cfg::N;
	42	+ constexpr uint32_t kK = cfg::K;
	43	+ constexpr uint32_t kScaleK = cfg::SCALE_K;
	44	+ constexpr uint32_t kCoreNum = cfg::CORE_NUM;
	45	+
	46	+ printf("Usage: %s (compile-time SCENARIO_NUM=%u)\n", argv[0], (uint32_t)SCENARIO_NUM);
	47	+ printf("Running scenario %u...\n", (uint32_t)SCENARIO_NUM);
	48	+
	49	+ const char* modeName = "";
	50	+ uint32_t numBlocks = kCoreNum;
	51	+ auto ascendc = platform_ascendc::PlatformAscendCManager::GetInstance();
	52	+
	53	+ if constexpr (SCENARIO_NUM == 1) {
	54	+ modeName = "Multi Core MDL Manual Tiling (Scale Sync)";
	55	+ numBlocks = kCoreNum;
	56	+ } else if constexpr (SCENARIO_NUM == 2) {
	57	+ modeName = "Multi Core MDL Manual Tiling (Scale Extra Cache)";
	58	+ numBlocks = kCoreNum;
	59	+ } else {
	60	+ std::cout << "Invalid SCENARIO_NUM: " << SCENARIO_NUM << std::endl;
	61	+ return -1;
	62	+ }
	63	+
	64	+ std::cout << "Running mode: " << modeName << std::endl;
	65	+
	66	+ if (ascendc == nullptr) {
	67	+ std::cout << "Failed to get platform instance." << std::endl;
	68	+ return -1;
	69	+ }
	70	+
	71	+ size_t aFileSize = static_cast<size_t>(kM * kK) * sizeof(uint8_t) / 2;
	72	+ size_t bFileSize = static_cast<size_t>(kK * kN) * sizeof(uint8_t) / 2;
	73	+ size_t asFileSize = static_cast<size_t>(kM * kScaleK) * sizeof(uint8_t);
	74	+ size_t bsFileSize = static_cast<size_t>(kScaleK * kN) * sizeof(uint8_t);
	75	+ size_t cFileSize = static_cast<size_t>(kM * kN) * sizeof(uint16_t);
	76	+
	77	+ size_t userWorkspaceSize = 0;
	78	+ size_t systemWorkspaceSize = static_cast<size_t>(ascendc->GetLibApiWorkSpaceSize());
	79	+ size_t workspaceSize = userWorkspaceSize + systemWorkspaceSize;
	80	+ int32_t deviceId = 0;
	81	+ aclrtStream stream = nullptr;
	82	+ aclrtContext context;
	83	+
	84	+ aclInit(nullptr);
	85	+ aclrtSetDevice(deviceId);
	86	+ aclrtCreateContext(&context, deviceId);
	87	+ aclrtCreateStream(&stream);
	88	+
	89	+ uint8_t* aHost;
	90	+ uint8_t* aDevice;
	91	+ aclrtMallocHost((void**)(&aHost), aFileSize);
	92	+ aclrtMalloc((void**)&aDevice, aFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	93	+ ReadFile("./input/x1_gm.bin", aFileSize, aHost, aFileSize);
	94	+ aclrtMemcpy(aDevice, aFileSize, aHost, aFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	95	+
	96	+ uint8_t* bHost;
	97	+ uint8_t* bDevice;
	98	+ aclrtMallocHost((void**)(&bHost), bFileSize);
	99	+ aclrtMalloc((void**)&bDevice, bFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	100	+ ReadFile("./input/x2_gm.bin", bFileSize, bHost, bFileSize);
	101	+ aclrtMemcpy(bDevice, bFileSize, bHost, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	102	+
	103	+ uint8_t* asHost;
	104	+ uint8_t* asDevice;
	105	+ aclrtMallocHost((void**)(&asHost), asFileSize);
	106	+ aclrtMalloc((void**)&asDevice, asFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	107	+ ReadFile("./input/x1_scale_gm.bin", asFileSize, asHost, asFileSize);
	108	+ aclrtMemcpy(asDevice, asFileSize, asHost, asFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	109	+
	110	+ uint8_t* bsHost;
	111	+ uint8_t* bsDevice;
	112	+ aclrtMallocHost((void**)(&bsHost), bsFileSize);
	113	+ aclrtMalloc((void**)&bsDevice, bsFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	114	+ ReadFile("./input/x2_scale_gm.bin", bsFileSize, bsHost, bsFileSize);
	115	+ aclrtMemcpy(bsDevice, bsFileSize, bsHost, bsFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	116	+
	117	+ uint8_t* cHost;
	118	+ uint8_t* cDevice;
	119	+ aclrtMallocHost((void**)(&cHost), cFileSize);
	120	+ aclrtMalloc((void**)&cDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	121	+
	122	+ uint8_t* workspaceDevice;
	123	+ aclrtMalloc((void**)&workspaceDevice, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	124	+
	125	+ constexpr auto selectedKernel = []() {
	126	+ if constexpr (SCENARIO_NUM == 1) {
	127	+ return matmul_custom_mdl_static<false>;
	128	+ } else if constexpr (SCENARIO_NUM == 2) {
	129	+ return matmul_custom_mdl_static<true>;
	130	+ }
	131	+ }();
	132	+
	133	+ selectedKernel<<<numBlocks, nullptr, stream>>>(aDevice, bDevice, asDevice, bsDevice, cDevice, workspaceDevice);
	134	+
	135	+ aclrtSynchronizeStream(stream);
	136	+ aclrtMemcpy(cHost, cFileSize, cDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST);
	137	+ WriteFile("./output/output.bin", cHost, cFileSize);
	138	+
	139	+ aclrtFree(aDevice);
	140	+ aclrtFreeHost(aHost);
	141	+ aclrtFree(bDevice);
	142	+ aclrtFreeHost(bHost);
	143	+ aclrtFree(asDevice);
	144	+ aclrtFreeHost(asHost);
	145	+ aclrtFree(bsDevice);
	146	+ aclrtFreeHost(bsHost);
	147	+ aclrtFree(cDevice);
	148	+ aclrtFreeHost(cHost);
	149	+ aclrtFree(workspaceDevice);
	150	+
	151	+ aclrtDestroyStream(stream);
	152	+ aclrtDestroyContext(context);
	153	+ aclrtResetDevice(deviceId);
	154	+ aclFinalize();
	155	+
	156	+ return 0;
	157	+}

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/matmul_mx.h+219-0

@@ -0,0 +1,219 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under 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 matmul.h
	13	+ * \brief Header file for matmul_mx case1~case2 implementation
	14	+ */
	15	+
	16	+#ifndef MATMULMX_H
	17	+#define MATMULMX_H
	18	+
	19	+#include "kernel_tiling/kernel_tiling.h"
	20	+#include "tiling/tiling_api.h"
	21	+#include "kernel_operator.h"
	22	+#define SPLIT_CORE_CUBE
	23	+#define ASCENDC_CUBE_ONLY
	24	+#include "lib/matmul_intf.h"
	25	+
	26	+struct ScaleFactors {
	27	+ uint32_t ka;
	28	+ uint32_t kb;
	29	+ uint32_t m;
	30	+ uint32_t n;
	31	+};
	32	+
	33	+struct MatmulMxStaticCfg {
	34	+ static constexpr uint32_t M = 8192;
	35	+ static constexpr uint32_t N = 8192;
	36	+ static constexpr uint32_t K = 8192;
	37	+
	38	+ static constexpr bool IS_TRANS_A = false;
	39	+ static constexpr bool IS_TRANS_B = false;
	40	+
	41	+ static constexpr uint32_t BASE_M = 256;
	42	+ static constexpr uint32_t BASE_N = 256;
	43	+ static constexpr uint32_t BASE_K = 256;
	44	+
	45	+ static constexpr uint32_t SINGLE_M = 2048;
	46	+ static constexpr uint32_t SINGLE_N = 1024;
	47	+ static constexpr uint32_t SINGLE_K = 8192;
	48	+
	49	+ static constexpr uint32_t SCALE_CEIL_NUMBER = 64;
	50	+ static constexpr uint32_t SCALE_NUMBER = 2;
	51	+ static constexpr uint32_t SCALE_K = (K + SCALE_CEIL_NUMBER - 1) / SCALE_CEIL_NUMBER * SCALE_NUMBER;
	52	+
	53	+ static constexpr uint32_t CORE_BLOCKS_M = M / SINGLE_M;
	54	+ static constexpr uint32_t CORE_BLOCKS_N = N / SINGLE_N;
	55	+ static constexpr uint32_t CORE_NUM = CORE_BLOCKS_M * CORE_BLOCKS_N;
	56	+
	57	+ static constexpr int DEPTH_A1 = 4;
	58	+ static constexpr int DEPTH_B1 = 4;
	59	+ static constexpr int STEP_KA = 2;
	60	+ static constexpr int STEP_KB = 2;
	61	+ static constexpr int STEP_M = 1;
	62	+ static constexpr int STEP_N = 1;
	63	+ static constexpr int DB_L0A = 2;
	64	+ static constexpr int DB_L0B = 2;
	65	+
	66	+ static constexpr ScaleFactors CASE1_SCALE = {1U, 1U, 1U, 1U};
	67	+ static constexpr ScaleFactors CASE2_SCALE = {4U, 4U, 1U, 1U};
	68	+};
	69	+
	70	+using cfg = MatmulMxStaticCfg;
	71	+
	72	+static_assert(cfg::M % cfg::SINGLE_M == 0, "M must be divisible by SINGLE_M");
	73	+static_assert(cfg::N % cfg::SINGLE_N == 0, "N must be divisible by SINGLE_N");
	74	+static_assert(cfg::K == cfg::SINGLE_K, "K must equal SINGLE_K");
	75	+
	76	+
	77	+// mxTypePara bit layout (MatmulApiStaticTiling/TCubeTiling):
	78	+// [0:6] scaleFactorKa
	79	+// [8:14] scaleFactorKb
	80	+// [16:22] scaleFactorM
	81	+// [24:30] scaleFactorN
	82	+constexpr uint32_t BuildMxTypePara(const ScaleFactors& scale)
	83	+{
	84	+ return ((scale.ka & 0x7FU) << 0) \| ((scale.kb & 0x7FU) << 8) \| ((scale.m & 0x7FU) << 16) \|
	85	+ ((scale.n & 0x7FU) << 24);
	86	+}
	87	+
	88	+constexpr uint32_t CASE1_MX_TYPE_PARA = BuildMxTypePara(cfg::CASE1_SCALE);
	89	+constexpr uint32_t CASE2_MX_TYPE_PARA = BuildMxTypePara(cfg::CASE2_SCALE);
	90	+
	91	+constexpr MatmulShapeParams SHAPE_PARAMS = {cfg::SINGLE_M, cfg::SINGLE_N, cfg::SINGLE_K,
	92	+ cfg::BASE_M, cfg::BASE_N, cfg::BASE_K};
	93	+
	94	+template <typename AType, typename BType, typename CType, bool EnableScaleCache>
	95	+__aicore__ inline constexpr MatmulApiStaticTiling GetMxConstantCFG()
	96	+{
	97	+ MatmulConfig mmCFG = GetMMConfig<MatmulConfigMode::CONFIG_MDL>(SHAPE_PARAMS);
	98	+ auto constantCFG = AscendC::GetMatmulApiTiling<AType, BType, CType, CType>(mmCFG);
	99	+
	100	+ constantCFG.depthA1 = cfg::DEPTH_A1;
	101	+ constantCFG.depthB1 = cfg::DEPTH_B1;
	102	+ constantCFG.stepKa = cfg::STEP_KA;
	103	+ constantCFG.stepKb = cfg::STEP_KB;
	104	+ constantCFG.stepM = cfg::STEP_M;
	105	+ constantCFG.stepN = cfg::STEP_N;
	106	+ constantCFG.dbL0A = cfg::DB_L0A;
	107	+ constantCFG.dbL0B = cfg::DB_L0B;
	108	+ constantCFG.mxTypePara = EnableScaleCache ? CASE2_MX_TYPE_PARA : CASE1_MX_TYPE_PARA;
	109	+ return constantCFG;
	110	+}
	111	+
	112	+template <bool EnableScaleCache>
	113	+class MatmulKernel {
	114	+public:
	115	+ __aicore__ inline MatmulKernel() {};
	116	+
	117	+ using aType =
	118	+ AscendC::MatmulTypeWithScale<AscendC::TPosition::GM, AscendC::TPosition::GM, CubeFormat::ND, fp4x2_e1m2_t,
	119	+ false>;
	W W watermelon_milk4月20日把scale转不转置的信息表达全或者写注释说明
	120	+ using bType =
	121	+ AscendC::MatmulTypeWithScale<AscendC::TPosition::GM, AscendC::TPosition::GM, CubeFormat::ND, fp4x2_e1m2_t,
	122	+ false>;
	123	+ using cType = AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, bfloat16_t>;
	124	+
	125	+ constexpr static auto CONSTANT_CFG = GetMxConstantCFG<aType, bType, cType, EnableScaleCache>();
	126	+
	127	+ AscendC::Matmul<aType, bType, cType, cType, CONSTANT_CFG,
	128	+ AscendC::MatmulCallBackFunc<nullptr, nullptr, nullptr>,
	129	+ AscendC::Impl::Detail::MatmulWithScalePolicy>
	130	+ matmulObj;
	131	+
	132	+ __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR as, GM_ADDR bs, GM_ADDR c)
	133	+ {
	134	+ aGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ fp4x2_e1m2_t>(a), kM kK);
	135	+ bGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ fp4x2_e1m2_t>(b), kK kN);
	136	+ cGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ bfloat16_t>(c), kM kN);
	137	+ asGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ AscendC::fp8_e8m0_t>(as), kM kScaleK);
	138	+ bsGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ AscendC::fp8_e8m0_t>(bs), kScaleK kN);
	139	+
	140	+ int32_t offsetA = 0;
	141	+ int32_t offsetB = 0;
	142	+ int32_t offsetC = 0;
	143	+ int32_t offsetAscale = 0;
	144	+ int32_t offsetBscale = 0;
	145	+ CalcOffset(AscendC::GetBlockIdx(), offsetA, offsetB, offsetAscale, offsetBscale, offsetC);
	146	+
	147	+ aGlobal = aGlobal[offsetA];
	148	+ bGlobal = bGlobal[offsetB];
	149	+ cGlobal = cGlobal[offsetC];
	150	+ asGlobal = asGlobal[offsetAscale];
	151	+ bsGlobal = bsGlobal[offsetBscale];
	152	+
	153	+ if (GetSysWorkSpacePtr() == nullptr) {
	154	+ return;
	155	+ }
	156	+ }
	157	+
	158	+ __aicore__ inline void Process(AscendC::TPipe* pipe)
	159	+ {
	160	+ if (AscendC::GetBlockIdx() >= kCoreNum) {
	161	+ return;
	162	+ }
	163	+
	164	+ REGIST_MATMUL_OBJ(pipe, GetSysWorkSpacePtr(), matmulObj, (TCubeTiling*)nullptr);
	165	+ matmulObj.SetOrgShape(kM, kN, kK);
	166	+
	167	+ matmulObj.SetTensorA(aGlobal, kIsTransA);
	168	+ matmulObj.SetTensorB(bGlobal, kIsTransB);
	169	+ matmulObj.SetTensorScaleA(asGlobal, false);
	170	+ matmulObj.SetTensorScaleB(bsGlobal, false);
	171	+
	172	+ matmulObj.IterateAll(cGlobal);
	173	+ matmulObj.End();
	174	+ }
	175	+
	176	+private:
	177	+ static constexpr uint32_t kM = cfg::M;
	178	+ static constexpr uint32_t kN = cfg::N;
	179	+ static constexpr uint32_t kK = cfg::K;
	180	+ static constexpr uint32_t kScaleK = cfg::SCALE_K;
	181	+ static constexpr uint32_t kSingleM = cfg::SINGLE_M;
	182	+ static constexpr uint32_t kSingleN = cfg::SINGLE_N;
	183	+ static constexpr uint32_t kScaleNumber = cfg::SCALE_NUMBER;
	184	+ static constexpr uint32_t kCoreBlocksM = cfg::CORE_BLOCKS_M;
	185	+ static constexpr uint32_t kCoreNum = cfg::CORE_NUM;
	186	+ static constexpr bool kIsTransA = cfg::IS_TRANS_A;
	187	+ static constexpr bool kIsTransB = cfg::IS_TRANS_B;
	188	+
	189	+ __aicore__ inline void CalcOffset(
	190	+ int32_t blockIdx, int32_t& offsetA, int32_t& offsetB, int32_t& offsetAscale, int32_t& offsetBscale,
	191	+ int32_t& offsetC)
	192	+ {
	193	+ auto mCoreIndex = blockIdx % kCoreBlocksM;
	194	+ auto nCoreIndex = blockIdx / kCoreBlocksM;
	195	+
	196	+ offsetA = mCoreIndex * kK * kSingleM;
	197	+ if (kIsTransA) {
	198	+ offsetA = mCoreIndex * kSingleM;
	199	+ }
	200	+
	201	+ offsetB = nCoreIndex * kSingleN;
	202	+ if (kIsTransB) {
	203	+ offsetB = nCoreIndex * kK * kSingleN;
	204	+ }
	205	+
	206	+ offsetAscale = mCoreIndex * kScaleK * kSingleM;
	207	+ offsetBscale = nCoreIndex * kSingleN * kScaleNumber;
	208	+
	209	+ offsetC = mCoreIndex * kN * kSingleM + nCoreIndex * kSingleN;
	210	+ }
	211	+
	212	+ AscendC::GlobalTensor<fp4x2_e1m2_t> aGlobal;
	213	+ AscendC::GlobalTensor<fp4x2_e1m2_t> bGlobal;
	214	+ AscendC::GlobalTensor<bfloat16_t> cGlobal;
	215	+ AscendC::GlobalTensor<AscendC::fp8_e8m0_t> asGlobal;
	216	+ AscendC::GlobalTensor<AscendC::fp8_e8m0_t> bsGlobal;
	217	+};
	218	+
	219	+#endif

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/scripts/gen_data.py+74-0

@@ -0,0 +1,74 @@
	1	+#!/usr/bin/python3
	2	+# coding=utf-8
	3	+
	4	+# ----------------------------------------------------------------------------------------------------------
	5	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	6	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	7	+# CANN Open Software License Agreement Version 2.0 (the "License").
	8	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	9	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	10	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	11	+# See LICENSE in the root of the software repository for the full text of the License.
	12	+# ----------------------------------------------------------------------------------------------------------
	13	+
	14	+
	15	+import os
	16	+import numpy as np
	17	+import ml_dtypes
	18	+import en_dtypes
	19	+import math
	20	+
	21	+bfloat16 = ml_dtypes.bfloat16
	22	+fp4_e1m2x2 = en_dtypes.float4_e1m2
	23	+fp4_e2m1x2 = en_dtypes.float4_e2m1
	24	+
	25	+def pack_two_fp4(scale_matrix):
	26	+ scale_matrix_row = scale_matrix.shape[0]
	27	+ scale_matrix_col = scale_matrix.shape[1]
	28	+ scale_matrix_bin = scale_matrix.flatten()
	29	+ scale_matrix_high = scale_matrix_bin[::2].view(np.uint8)
	30	+ scale_matrix_low = scale_matrix_bin[1::2].view(np.uint8)
	31	+ low_bits = (scale_matrix_low & 0x0F) << 4
	32	+ high_bits = scale_matrix_high & 0x0F
	33	+ combined = low_bits \| high_bits
	34	+ scale_matrix_bin = combined.reshape(scale_matrix_row, scale_matrix_col // 2)
	35	+ return scale_matrix_bin
	36	+
	37	+def gen_golden_data():
	38	+ m, n, k = 8192, 8192, 8192
	39	+ sk = (int)(np.ceil(k / 64) * 2)
	40	+
	41	+ os.makedirs("input", exist_ok=True)
	42	+ os.makedirs("output", exist_ok=True)
	43	+
	44	+ x1_gm = np.random.randint(-1, 2, [m, k]).astype(fp4_e1m2x2)
	45	+ x2_gm = np.random.randint(-1, 2, [k, n]).astype(fp4_e1m2x2)
	46	+
	47	+ x1_scale_gm = np.random.randint(127, 130, [m, sk]).astype(np.uint8)
	48	+ x2_scale_gm = np.random.randint(127, 130, [sk, n]).astype(np.uint8)
	49	+
	50	+ ###################### compute ########################
	51	+ x1_mx = 2**(x1_scale_gm.astype(np.float64) - 127)
	52	+ x2_mx = 2**(x2_scale_gm.astype(np.float64) - 127)
	53	+ x1_full = np.zeros([m, k], dtype=np.float64)
	54	+ x2_full = np.zeros([k, n], dtype=np.float64)
	55	+
	56	+ for i in range(x1_gm.shape[1]):
	57	+ x1_full[:, i] = x1_gm[:, i] * x1_mx[:, i // 32]
	58	+ x2_full[i, :] = x2_gm[i, :] * x2_mx[i // 32, :]
	59	+
	60	+ golden = np.matmul(x1_full.astype(np.float64), x2_full.astype(np.float64)).astype(bfloat16)
	61	+
	62	+ x2_scale_gm = x2_scale_gm.reshape(int(sk / 2), 2, n).transpose(0, 2, 1)
	63	+ x1_gm_packed = pack_two_fp4(x1_gm)
	64	+ x2_gm_packed = pack_two_fp4(x2_gm)
	65	+ x1_gm_packed.tofile("./input/x1_gm.bin")
	66	+ x2_gm_packed.tofile("./input/x2_gm.bin")
	67	+ x1_scale_gm.tofile("./input/x1_scale_gm.bin")
	68	+ x2_scale_gm.tofile("./input/x2_scale_gm.bin")
	69	+ golden.tofile("./output/golden.bin")
	70	+
	71	+
	72	+if __name__ == "__main__":
	73	+ gen_golden_data()
	74	+

Aexamples/01_simd_cpp_api/04_best_practices/01_matrix_compute_practices/matmul_mxfp4_high_performance/scripts/verify_result.py+71-0

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