GroupedDynamicMxQuant add cublas algorithm and params-ops-nn-AtomGit

已合并

GroupedDynamicMxQuant add cublas algorithm and params #3577

yeeyanghsia创建于 4月7日从yeeyanghsia/ops-nn： GroupedDynamicMxQuant 合入到 cann/ops-nn： master

GroupedDynamicMxQuant add cublas algorithm and params #3577

已合并

yeeyanghsia创建于 4月7日

从yeeyanghsia/ops-nn GroupedDynamicMxQuant合入到cann/ops-nn master

收起

docs / zh
- op_api_list.md
quant / grouped_dynamic_mx_quant
- CMakeLists.txt
- README.md
- docs
  - aclnnGroupedDynamicMxQuantV2.md
- examples / arch35
  - test_aclnn_grouped_dynamic_mx_quant.cpp
  - test_aclnn_grouped_dynamic_mx_quant_v2.cpp
- op_api
  - aclnn_grouped_dynamic_mx_quant.cpp
  - aclnn_grouped_dynamic_mx_quant_v2.cpp
  - aclnn_grouped_dynamic_mx_quant_v2.h
  - grouped_dynamic_mx_quant.cpp
  - grouped_dynamic_mx_quant.h
- op_graph
  - grouped_dynamic_mx_quant_proto.h
- op_host
  - arch35
    - grouped_dynamic_mx_quant_tiling_arch35.cpp
    - grouped_dynamic_mx_quant_tiling_arch35.h
  - config / ascend950
    - grouped_dynamic_mx_quant_binary.json
  - grouped_dynamic_mx_quant_def.cpp
  - grouped_dynamic_mx_quant_infershape.cpp
  - op_api
    - aclnn_grouped_dynamic_mx_quant_v2.h
    - grouped_dynamic_mx_quant.h
- op_kernel
  - arch35
    - grouped_dynamic_mx_quant_common.h
    - grouped_dynamic_mx_quant_not_tail_axis_fp8.h
    - grouped_dynamic_mx_quant_struct.h
    - grouped_dynamic_mx_quant_tilingdata.h
  - grouped_dynamic_mx_quant.cpp
- tests
  - CMakeLists.txt
  - ut
    - CMakeLists.txt
    - op_api
      - CMakeLists.txt
      - test_aclnn_grouped_dynamic_mx_quant.cpp
    - op_host
      - CMakeLists.txt
      - arch35
        test_grouped_dynamic_mx_quant_tiling.cpp
      - test_grouped_dynamic_mx_quant_infershape.cpp

共 31 个文件变更+2203-539

Mdocs/zh/op_api_list.md+1-0

	@@ -99,6 +99,7 @@
99	\| [aclnnDequantSwigluQuantV2](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuantV2.md) \| 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 \| 默认确定性实现 \| 默认确定性实现 \|	99	\| [aclnnDequantSwigluQuantV2](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuantV2.md) \| 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 \| 默认确定性实现 \| 默认确定性实现 \|
100	\| [aclnnSwigluMxQuant](../../quant/swiglu_mx_quant/docs/aclnnSwigluMxQuant.md) \| 在Swish门控线性单元激活函数前后添加DynamicMxQuant操作，实现x的SwigluMxQuant计算。 \| 默认确定性实现 \| 默认确定性实现 \|	100	\| [aclnnSwigluMxQuant](../../quant/swiglu_mx_quant/docs/aclnnSwigluMxQuant.md) \| 在Swish门控线性单元激活函数前后添加DynamicMxQuant操作，实现x的SwigluMxQuant计算。 \| 默认确定性实现 \| 默认确定性实现 \|
101	\| [aclnnGroupedDynamicMxQuant](../../quant/grouped_dynamic_mx_quant/docs/aclnnGroupedDynamicMxQuant.md) \| 根据传入的分组索引的起始值，对传入的数据进行分组的float8的动态量化。 \| - \| 默认确定性实现 \|	101	\| [aclnnGroupedDynamicMxQuant](../../quant/grouped_dynamic_mx_quant/docs/aclnnGroupedDynamicMxQuant.md) \| 根据传入的分组索引的起始值，对传入的数据进行分组的float8的动态量化。 \| - \| 默认确定性实现 \|
		102	+\| [aclnnGroupedDynamicMxQuantV2](../../quant/grouped_dynamic_mx_quant/docs/aclnnGroupedDynamicMxQuantV2.md) \| 根据传入的分组索引的起始值，对传入的数据进行分组的float8的动态量化。根据scaleAlg不同取值，选择不同的量化算法。 \| - \| 默认确定性实现 \|
102	\| [aclnnDynamicDualLevelMxQuant](../../quant/dynamic_dual_level_mx_quant/docs/aclnnDynamicDualLevelMxQuant.md) \| 通过给定的level0BlockSize和level1BlockSize将输入进行两次划分，以数据块为粒度，进行目的数据类型为FLOAT4类的MX二级量化。 \| - \| 默认确定性实现 \|	103	\| [aclnnDynamicDualLevelMxQuant](../../quant/dynamic_dual_level_mx_quant/docs/aclnnDynamicDualLevelMxQuant.md) \| 通过给定的level0BlockSize和level1BlockSize将输入进行两次划分，以数据块为粒度，进行目的数据类型为FLOAT4类的MX二级量化。 \| - \| 默认确定性实现 \|
103	\| [aclnnGroupedDynamicBlockQuant](../../quant/grouped_dynamic_block_quant/docs/aclnnGroupedDynamicBlockQuant.md) \| 根据传入的分组索引的起始值对各个group以基本块的粒度进行量化，并输出量化参数scale。 \| - \| 默认确定性实现 \|	104	\| [aclnnGroupedDynamicBlockQuant](../../quant/grouped_dynamic_block_quant/docs/aclnnGroupedDynamicBlockQuant.md) \| 根据传入的分组索引的起始值对各个group以基本块的粒度进行量化，并输出量化参数scale。 \| - \| 默认确定性实现 \|
104	\| [aclnnDualLevelQuantMatmulWeightNz](../../matmul/dual_level_quant_batch_matmul/docs/aclnnDualLevelQuantMatmulWeightNz.md) \| 完成二级量化mxfp4的矩阵乘计算。 \| - \| 默认确定性实现 \|	105	\| [aclnnDualLevelQuantMatmulWeightNz](../../matmul/dual_level_quant_batch_matmul/docs/aclnnDualLevelQuantMatmulWeightNz.md) \| 完成二级量化mxfp4的矩阵乘计算。 \| - \| 默认确定性实现 \|

Mquant/grouped_dynamic_mx_quant/CMakeLists.txt+1-1

	@@ -1,5 +1,5 @@
1	# -----------------------------------------------------------------------------------------------------------	1	# -----------------------------------------------------------------------------------------------------------
2	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.

Mquant/grouped_dynamic_mx_quant/README.md+85-34

	@@ -13,43 +13,74 @@
13		13
14	## 功能说明	14	## 功能说明
15		15
16	-- 接口功能：根据传入的分组索引的起始值，对传入的数据进行分组的float8的动态量化。	16	+- 算子功能：根据传入的分组索引的起始值，对传入的数据进行分组的float8的动态量化。
17		17
18	- 计算公式：	18	- 计算公式：
19	- - 将输入x在第0维上先按照groupIndex进行分组，每个group内按k = blocksize个数分组，一组k个数 {{x<sub>i</sub>}<sub>i=1</sub><sup>k</sup>} 计算出这组数对应的量化尺度mxscale_pre, {mxscale_pre, {P<sub>i</sub>}<sub>i=1</sub><sup>k</sup>}, 计算公式为下面公式(1)(2)。	19	+ - 场景1，当scaleAlg为0时：
20	- $$	20	+ - 将输入x在第0维上先按照groupIndex进行分组，每个group内按k = blocksize个数分组，一组k个数 {{x<sub>i</sub>}<sub>i=1</sub><sup>k</sup>} 计算出这组数对应的量化尺度mxscale_pre, {mxscale_pre, {P<sub>i</sub>}<sub>i=1</sub><sup>k</sup>}, 计算公式为下面公式(1)(2)。
21	- shared\_exp = floor(log_2(max_i(\|V_i\|))) - emax \tag{1}	21	+
22	- $$	22	+ $$
23	- $$	23	+ shared\_exp = floor(log_2(max_i(\|V_i\|))) - emax \tag{1}
24	- mxscale\_pre = 2^{shared\_exp} \tag{2}	24	+ $$
25	- $$	25	+
26	- - 这组数每一个除以mxscale，根据round_mode转换到对应的dst_type，得到量化结果y, 计算公式为下面公式(3)。	26	+ $$
27	- $$	27	+ mxscale\_pre = 2^{shared\_exp} \tag{2}
28	- P_i = cast\_to\_dst\_type(V_i/mxscale, round\_mode), \space i\space from\space 1\space to\space blocksize \tag{3}	28	+ $$
29	- $$	29	+
30	-	30	+ - 这组数每个数都除以mxscale，根据round_mode转换到对应的dst_type，得到量化结果y，计算公式为下面公式(3)。
31	- 量化后的P<sub>i</sub>按对应的x<sub>i</sub>的位置组成输出y，mxscale_pre按对应的groupIndex分组，分组内第一个维度pad为偶数，组成输出mxscale。	31	+
32	-	32	+ $$
33	- - emax: 对应数据类型的最大正则数的指数位。	33	+ P_i = cast\_to\_dst\_type(V_i/mxscale, round\_mode), \space i\space from\space 1\space to\space blocksize \tag{3}
34	-	34	+ $$
35	- \| DataType \| emax \|	35	+
36	- \| :-----------: \| :--: \|	36	+ - 量化后的 $P_{i}$ 按对应的 $V_{i}$ 的位置组成输出y，mxscale_pre按对应的groupIndex分组，分组内第一个维度pad为偶数，组成输出mxscale。
37	- \| FLOAT8_E4M3FN \| 8 \|	37	+
38	- \| FLOAT8_E5M2 \| 15 \|	38	+ - emax：对应数据类型的最大正则数的指数位。
		39	+
		40	+ \| DataType \| emax \|
		41	+ \| :-----------: \| :--: \|
		42	+ \| FLOAT8_E4M3FN \| 8 \|
		43	+ \| FLOAT8_E5M2 \| 15 \|
		44	+
		45	+ - 场景2，当scaleAlg为1时：
		46	+ - 将长向量按块分，每块长度为k，对每块单独计算一个块缩放因子$S_{fp32}^b$，再把块内所有元素用同一个$S_{fp32}^b$映射到目标低精度类型FP8。如果最后一块不足k个元素，把缺失值视为0，按照完整块处理。
		47	+ - 找到该块中数值的最大绝对值:
		48	+
		49	+ $$
		50	+ Amax(D_{fp32}^b)=max(\{\|d_{i}\|\}_{i=1}^{k})
		51	+ $$
		52	+
		53	+ - 将FP32映射到目标数据类型FP8可表示的范围内，其中$Amax(DType)$是目标精度能表示的最大值。
		54	+
		55	+ $$
		56	+ S_{fp32}^b = \frac{Amax(D_{fp32}^b)}{Amax(DType)}
		57	+ $$
		58	+
		59	+ - 将块缩放因子$S_{fp32}^b$转换为FP8格式下可表示的缩放值$S_{ue8m0}^b$
		60	+ - 从块的浮点缩放因子$S_{fp32}^b$中提取无偏指数$E_{int}^b$和尾数$M_{fixp}^b$
		61	+ - 为保证量化时不溢出，对指数进行向上取整，且在FP8可表示的范围内：
		62	+
		63	+ $$
		64	+ E_{int}^b = \begin{cases} E_{int}^b + 1, & \text{如果} S_{fp32}^b \text{为正规数，且} E_{int}^b < 254 \text{且} M_{fixp}^b > 0 \\ E_{int}^b + 1, & \text{如果} S_{fp32}^b \text{为非正规数，且} M_{fixp}^b > 0.5 \\ E_{int}^b, & \text{否则} \end{cases}
		65	+ $$
		66	+
		67	+ - 计算块缩放因子：$S_{ue8m0}^b=2^{E_{int}^b}$
		68	+ - 计算块转换因子：$R_{fp32}^b=\frac{1}{fp32(S_{ue8m0}^b)}$
		69	+ - 应用到量化的最终步骤，对于每个块内元素，$d^i = DType(d_{fp32}^i \cdot R_{fp32}^n)$，最终输出的量化结果是$\left(S^b, [d^i]_{i=1}^k\right)$，其中$S^b$代表块的缩放因子，这里指$S_{ue8m0}^b$，$[d^i]_{i=1}^k$代表块内量化后的数据。
39		70
40	## 参数说明	71	## 参数说明
41		72
42	-<table style="undefined;table-layout: fixed; width: 980px"><colgroup>	73	+<table style="undefined;table-layout: fixed; width: 1010px"><colgroup>
43	- <col style="width: 100px">	74	+ <col style="width: 130px">
44	<col style="width: 150px">	75	<col style="width: 150px">
45	- <col style="width: 280px">
46	<col style="width: 330px">	76	<col style="width: 330px">
		77	+ <col style="width: 280px">
47	<col style="width: 120px">	78	<col style="width: 120px">
48	</colgroup>	79	</colgroup>
49	<thead>	80	<thead>
50	<tr>	81	<tr>
51	<th>参数名</th>	82	<th>参数名</th>
52	- <th>输入/输出</th>	83	+ <th>输入/输出/属性</th>
53	<th>描述</th>	84	<th>描述</th>
54	<th>数据类型</th>	85	<th>数据类型</th>
55	<th>数据格式</th>	86	<th>数据格式</th>
	@@ -65,42 +96,56 @@
65	<tr>	96	<tr>
66	<td>groupIndex</td>	97	<td>groupIndex</td>
67	<td>输入</td>	98	<td>输入</td>
68	- <td>Device侧的aclTensor，量化分组的起始索引。shape仅支持1维。支持非连续的Tensor，支持空Tensor。</td>	99	+ <td>Device侧的aclTensor，量化分组的起始索引。shape仅支持1维。支持非连续的Tensor，不支持空Tensor。</td>
69	<td>INT32</td>	100	<td>INT32</td>
70	<td>ND</td>	101	<td>ND</td>
71	</tr>	102	</tr>
72	<tr>	103	<tr>
73	<td>roundMode</td>	104	<td>roundMode</td>
74	- <td>输入</td>	105	+ <td>属性</td>
75	<td>host侧的string，公式中的round_mode，数据转换的模式，仅支持"rint"模式。</td>	106	<td>host侧的string，公式中的round_mode，数据转换的模式，仅支持"rint"模式。</td>
76	- <td>-</td>	107	+ <td>STRING</td>
77	<td>-</td>	108	<td>-</td>
78	</tr>	109	</tr>
79	<tr>	110	<tr>
80	<td>dstType</td>	111	<td>dstType</td>
81	- <td>输入</td>	112	+ <td>属性</td>
82	<td>host侧的int64_t，公式中的dst_type，指定数据转换后y的类型，输入范围为{35, 36}，分别对应输出y的数据类型为{35: FLOAT8_E5M2, 36: FLOAT8_E4M3FN}。</td>	113	<td>host侧的int64_t，公式中的dst_type，指定数据转换后y的类型，输入范围为{35, 36}，分别对应输出y的数据类型为{35: FLOAT8_E5M2, 36: FLOAT8_E4M3FN}。</td>
83	- <td>-</td>	114	+ <td>INT64</td>
84	<td>-</td>	115	<td>-</td>
85	</tr>	116	</tr>
86	<tr>	117	<tr>
87	<td>blocksize</td>	118	<td>blocksize</td>
88	- <td>输入</td>	119	+ <td>属性</td>
89	<td>host侧的int64_t，公式中的blocksize，指定每次量化的元素个数，仅支持32。</td>	120	<td>host侧的int64_t，公式中的blocksize，指定每次量化的元素个数，仅支持32。</td>
		121	+ <td>INT64</td>
90	<td>-</td>	122	<td>-</td>
		123	+ </tr>
		124	+ <tr>
		125	+ <td>scaleAlg</td>
		126	+ <td>属性</td>
		127	+ <td>host侧的int64_t，指定mxscale计算时采用的算法，仅支持0和1。</td>
		128	+ <td>INT64</td>
		129	+ <td>-</td>
		130	+ </tr>
		131	+ <tr>
		132	+ <td>dstTypeMax</td>
		133	+ <td>属性</td>
		134	+ <td>host侧的float32，在scale_alg=2时生效。默认值0.0表示max_type为目标数据类型的最大值，若传入其它数值，则需要按照传入的数值计算mxscale。当前支持取值为0.0/6.0-12.0，只支持在FLOAT4_E2M1场景设置该值。</td>
		135	+ <td>FLOAT</td>
91	<td>-</td>	136	<td>-</td>
92	</tr>	137	</tr>
93	<tr>	138	<tr>
94	<td>y</td>	139	<td>y</td>
95	<td>输出</td>	140	<td>输出</td>
96	- <td>Device侧的aclTensor，公式中的输出y，输入x量化后的对应结果。需与dstType对应，shape仅支持2维，支持空Tensor，Shape和输入x一致。</td>	141	+ <td>Device侧的aclTensor，公式中的输出y，输入x量化后的对应结果。需与dstType对应，shape仅支持2维，支持空Tensor，shape和输入x一致。</td>
97	<td>FLOAT8_E4M3FN、FLOAT8_E5M2</td>	142	<td>FLOAT8_E4M3FN、FLOAT8_E5M2</td>
98	<td>ND</td>	143	<td>ND</td>
99	</tr>	144	</tr>
100	<tr>	145	<tr>
101	<td>mxscale</td>	146	<td>mxscale</td>
102	<td>输出</td>	147	<td>输出</td>
103	- <td>Device侧的aclTensor，公式中的mxscale_pre组成的输出mxscale，每个分组对应的量化尺度。需与dstType对应，shape仅支持3维，支持空Tensor，Shape和输入x一致。假设x的shape为 $[m,n]$，groupedIndex的shape为 $[g]$，则mxscale的shape为 $[(m/(blocksize * 2)+g), n, 2]$</td>	148	+ <td>Device侧的aclTensor，公式中的mxscale_pre组成的输出mxscale，每个分组对应的量化尺度。shape仅支持3维，支持空Tensor。假设x的shape为 [m,n]，groupedIndex的shape为 [g]，则mxscale的shape为 [(m/(blocksize * 2)+g), n, 2]。</td>
104	<td>FLOAT8_E8M0</td>	149	<td>FLOAT8_E8M0</td>
105	<td>ND</td>	150	<td>ND</td>
106	</tr>	151	</tr>
	@@ -108,10 +153,16 @@
108		153
109	## 约束说明	154	## 约束说明
110		155
111	-无	156	+- 关于x、groupIndex、y、mxscale的约束说明如下：
		157	+ - groupIndex中的值必须非递减，且不能小于0，最后一个元素必须为x第一个维度的长度。
		158	+ - $rank(mxscale) = rank(x) + 1$。
		159	+ - 假设x的shape为 $[m,n]$，groupedIndex的shape为 $[g]$，则mxscale的shape为 $[(m/(blocksize * 2)+g), n, 2]$。
		160	+ - $mxscale.shape[-1] = 2$。
		161	+ - 输出y的shape与输入x一致。
112		162
113	## 调用说明	163	## 调用说明
114		164
115	\| 调用方式 \| 调用样例 \| 说明 \|	165	\| 调用方式 \| 调用样例 \| 说明 \|
116	\|--------------\|------------------------------------------------------------------------\|--------------------------------------------------------------\|	166	\|--------------\|------------------------------------------------------------------------\|--------------------------------------------------------------\|
117	\| aclnn调用 \| [test_aclnn_grouped_dynamic_mx_quant](./examples/arch35/test_aclnn_grouped_dynamic_mx_quant.cpp) \| 通过[aclnnGroupedDynamicMxQuant](./docs/aclnnGroupedDynamicMxQuant.md)接口方式调用GroupedDynamicMxQuant算子。 \|	167	\| aclnn调用 \| [test_aclnn_grouped_dynamic_mx_quant](./examples/arch35/test_aclnn_grouped_dynamic_mx_quant.cpp) \| 通过[aclnnGroupedDynamicMxQuant](./docs/aclnnGroupedDynamicMxQuant.md)接口方式调用GroupedDynamicMxQuant算子。 \|
		168	+\| aclnn调用 \| [test_aclnn_grouped_dynamic_mx_quant_v2](./examples/arch35/test_aclnn_grouped_dynamic_mx_quant.cpp) \| 通过[aclnnGroupedDynamicMxQuantV2](./docs/aclnnGroupedDynamicMxQuantV2.md)接口方式调用GroupedDynamicMxQuant算子。 \|

Aquant/grouped_dynamic_mx_quant/docs/aclnnGroupedDynamicMxQuantV2.md+531-0

@@ -0,0 +1,531 @@
	1	+# aclnnGroupedDynamicMxQuantV2
	2	+
	3	+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/quant/grouped_dynamic_mx_quant)
	4	+
	5	+## 产品支持情况
	6	+
	7	+\| 产品 \| 是否支持 \|
	8	+\| :----------------------------------------------------------- \| :------: \|
	9	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
	10	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| × \|
	11	+\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| × \|
	12	+\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|
	13	+\| <term>Atlas 推理系列产品</term> \| × \|
	14	+\| <term>Atlas 训练系列产品</term> \| × \|
	15	+
	16	+## 功能说明
	17	+
	18	+- 接口功能：根据传入的分组索引的起始值，对传入的数据进行分组的float8的动态量化。
	19	+
	20	+- 计算公式：
	21	+ - 场景1，当scaleAlg为0时：
	22	+ - 将输入x在第0维上先按照groupIndex进行分组，每个group内按k = blocksize个数分组，一组k个数 {{x<sub>i</sub>}<sub>i=1</sub><sup>k</sup>} 计算出这组数对应的量化尺度mxscale_pre, {mxscale_pre, {P<sub>i</sub>}<sub>i=1</sub><sup>k</sup>}, 计算公式为下面公式(1)(2)。
	23	+
	24	+ $$
	25	+ shared\_exp = floor(log_2(max_i(\|V_i\|))) - emax \tag{1}
	26	+ $$
	27	+
	28	+ $$
	29	+ mxscale\_pre = 2^{shared\_exp} \tag{2}
	30	+ $$
	31	+
	32	+ - 这组数每个数都除以mxscale，根据round_mode转换到对应的dst_type，得到量化结果y，计算公式为下面公式(3)。
	33	+
	34	+ $$
	35	+ P_i = cast\_to\_dst\_type(V_i/mxscale, round\_mode), \space i\space from\space 1\space to\space blocksize \tag{3}
	36	+ $$
	37	+
	38	+ - 量化后的 $P_{i}$ 按对应的 $V_{i}$ 的位置组成输出y，mxscale_pre按对应的groupIndex分组，分组内第一个维度pad为偶数，组成输出mxscale。
	39	+
	40	+ - emax：对应数据类型的最大正则数的指数位。
	41	+
	42	+ \| DataType \| emax \|
	43	+ \| :-----------: \| :--: \|
	44	+ \| FLOAT8_E4M3FN \| 8 \|
	45	+ \| FLOAT8_E5M2 \| 15 \|
	46	+
	47	+ - 场景2，当scaleAlg为1时：
	48	+ - 将长向量按块分，每块长度为k，对每块单独计算一个块缩放因子$S_{fp32}^b$，再把块内所有元素用同一个$S_{fp32}^b$映射到目标低精度类型FP8。如果最后一块不足k个元素，把缺失值视为0，按照完整块处理。
	49	+ - 找到该块中数值的最大绝对值:
	50	+
	51	+ $$
	52	+ Amax(D_{fp32}^b)=max(\{\|d_{i}\|\}_{i=1}^{k})
	53	+ $$
	54	+
	55	+ - 将FP32映射到目标数据类型FP8可表示的范围内，其中$Amax(DType)$是目标精度能表示的最大值。
	56	+
	57	+ $$
	58	+ S_{fp32}^b = \frac{Amax(D_{fp32}^b)}{Amax(DType)}
	59	+ $$
	60	+
	61	+ - 将块缩放因子$S_{fp32}^b$转换为FP8格式下可表示的缩放值$S_{ue8m0}^b$
	62	+ - 从块的浮点缩放因子$S_{fp32}^b$中提取无偏指数$E_{int}^b$和尾数$M_{fixp}^b$
	63	+ - 为保证量化时不溢出，对指数进行向上取整，且在FP8可表示的范围内：
	64	+
	65	+ $$
	66	+ E_{int}^b = \begin{cases} E_{int}^b + 1, & \text{如果} S_{fp32}^b \text{为正规数，且} E_{int}^b < 254 \text{且} M_{fixp}^b > 0 \\ E_{int}^b + 1, & \text{如果} S_{fp32}^b \text{为非正规数，且} M_{fixp}^b > 0.5 \\ E_{int}^b, & \text{否则} \end{cases}
	67	+ $$
	68	+
	69	+ - 计算块缩放因子：$S_{ue8m0}^b=2^{E_{int}^b}$
	70	+ - 计算块转换因子：$R_{fp32}^b=\frac{1}{fp32(S_{ue8m0}^b)}$
	71	+ - 应用到量化的最终步骤，对于每个块内元素，$d^i = DType(d_{fp32}^i \cdot R_{fp32}^n)$，最终输出的量化结果是$\left(S^b, [d^i]_{i=1}^k\right)$，其中$S^b$代表块的缩放因子，这里指$S_{ue8m0}^b$，$[d^i]_{i=1}^k$代表块内量化后的数据。
	72	+
	73	+## 函数原型
	74	+
	75	+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnGroupedDynamicMxQuantV2GetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnGroupedDynamicMxQuantV2”接口执行计算。
	76	+
	77	+```cpp
	78	+aclnnStatus aclnnGroupedDynamicMxQuantV2GetWorkspaceSize(
	79	+ const aclTensor *x,
	80	+ const aclTensor *groupIndex,
	81	+ const char *roundMode,
	82	+ int64_t dstType,
	83	+ int64_t blocksize,
	84	+ int64_t scaleAlg,
	85	+ double dstTypeMax,
	86	+ const aclTensor *y,
	87	+ const aclTensor *mxscale,
	88	+ uint64_t *workspaceSize,
	89	+ aclOpExecutor **executor)
	90	+```
	91	+
	92	+```cpp
	93	+aclnnStatus aclnnGroupedDynamicMxQuantV2(
	94	+ void *workspace,
	95	+ uint64_t workspaceSize,
	96	+ aclOpExecutor *executor,
	97	+ aclrtStream stream)
	98	+```
	99	+
	100	+## aclnnGroupedDynamicMxQuantV2GetWorkspaceSize
	101	+
	102	+- 参数说明：
	103	+ <table style="undefined;table-layout: fixed; width: 1550px"><colgroup>
	104	+ <col style="width: 180px">
	105	+ <col style="width: 120px">
	106	+ <col style="width: 280px">
	107	+ <col style="width: 320px">
	108	+ <col style="width: 250px">
	109	+ <col style="width: 120px">
	110	+ <col style="width: 140px">
	111	+ <col style="width: 140px">
	112	+ </colgroup>
	113	+ <thead>
	114	+ <tr>
	115	+ <th>参数名</th>
	116	+ <th>输入/输出</th>
	117	+ <th>描述</th>
	118	+ <th>使用说明</th>
	119	+ <th>数据类型</th>
	120	+ <th>数据格式</th>
	121	+ <th>维度(shape)</th>
	122	+ <th>非连续Tensor</th>
	123	+ </tr></thead>
	124	+ <tbody>
	125	+ <tr>
	126	+ <td>x (aclTensor*)</td>
	127	+ <td>输入</td>
	128	+ <td>表示算子输入的Tensor。计算公式中的输入x。</td>
	129	+ <td>支持空Tensor。</td>
	130	+ <td>FLOAT16、BFLOAT16</td>
	131	+ <td>ND</td>
	132	+ <td>2</td>
	133	+ <td>√</td>
	134	+ </tr>
	135	+ <tr>
	136	+ <td>groupIndex (aclTensor*)</td>
	137	+ <td>输入</td>
	138	+ <td>量化分组的起始索引。</td>
	139	+ <td><ul><li>不支持空Tensor。</li><li>索引要求大于等于0，且非递减，并且最后一个数需要与x的第一个维度大小相等。</li></ul></td>
	140	+ <td>INT32</td>
	141	+ <td>ND</td>
	142	+ <td>1</td>
	143	+ <td>√</td>
	144	+ </tr>
	145	+ <tr>
	146	+ <td>roundMode（char*）</td>
	147	+ <td>输入</td>
	148	+ <td>公式中的round_mode，数据转换的模式。</td>
	149	+ <td>仅支持"rint"模式。</td>
	150	+ <td>STRING</td>
	151	+ <td>-</td>
	152	+ <td>-</td>
	153	+ <td>-</td>
	154	+ </tr>
	155	+ <tr>
	156	+ <td>dstType (int64_t)</td>
	157	+ <td>输入</td>
	158	+ <td>公式中的dst_type，指定数据转换后y的类型。</td>
	159	+ <td>输入范围为{35, 36}，分别对应输出y的数据类型为{35: FLOAT8_E5M2, 36: FLOAT8_E4M3FN}。</td>
	160	+ <td>INT64</td>
	161	+ <td>-</td>
	162	+ <td>-</td>
	163	+ <td>-</td>
	164	+ </tr>
	165	+ <tr>
	166	+ <td>blocksize (int64_t)</td>
	167	+ <td>输入</td>
	168	+ <td>公式中的blocksize，指定每次量化的元素个数。</td>
	169	+ <td>当前取值仅支持32。</td>
	170	+ <td>INT64</td>
	171	+ <td>-</td>
	172	+ <td>-</td>
	173	+ <td>-</td>
	174	+ </tr>
	175	+ <tr>
	176	+ <td>scaleAlg (int64_t)</td>
	177	+ <td>输入</td>
	178	+ <td>指定mxscale计算时采用的算法。</td>
	179	+ <td>当前取值仅支持0和1。</td>
	180	+ <td>INT64</td>
	181	+ <td>-</td>
	182	+ <td>-</td>
	183	+ <td>-</td>
	184	+ </tr>
	185	+ <tr>
	186	+ <td>dstTypeMax (double)</td>
	187	+ <td>输入</td>
	188	+ <td>预留参数。</td>
	189	+ <td>当前取值仅支持0.0。</td>
	190	+ <td>DOUBLE</td>
	191	+ <td>-</td>
	192	+ <td>-</td>
	193	+ <td>-</td>
	194	+ </tr>
	195	+ <tr>
	196	+ <td>y (aclTensor*)</td>
	197	+ <td>输出</td>
	198	+ <td>表示量化后的输出Tensor。对应公式中的y。</td>
	199	+ <td><ul><li>支持空Tensor。</li><li>shape的维度与x保持一致。</li><li>数据类型支持FLOAT8_E4M3FN、FLOAT8_E5M2，需与dstType对应。</li></ul></td>
	200	+ <td>FLOAT8_E4M3FN、FLOAT8_E5M2</td>
	201	+ <td>ND</td>
	202	+ <td>2</td>
	203	+ <td>√</td>
	204	+ </tr>
	205	+ <tr>
	206	+ <td>mxscale (aclTensor*)</td>
	207	+ <td>输出</td>
	208	+ <td>公式中的mxscale_pre组成的输出mxscale，每个分组对应的量化尺度。</td>
	209	+ <td><ul><li>支持空Tensor。</li><li>假设x的shape为 [m,n]，groupedIndex的shape为 [g]，则mxscale的shape为 [(m/(blocksize∗2)+g),n,2]。</li></ul></td>
	210	+ <td>FLOAT8_E8M0</td>
	211	+ <td>ND</td>
	212	+ <td>3</td>
	213	+ <td>√</td>
	214	+ </tr>
	215	+ <tr>
	216	+ <td>workspaceSize</td>
	217	+ <td>输出</td>
	218	+ <td>返回需要在Device侧申请的workspace大小。</td>
	219	+ <td>-</td>
	220	+ <td>-</td>
	221	+ <td>-</td>
	222	+ <td>-</td>
	223	+ <td>-</td>
	224	+ </tr>
	225	+ <tr>
	226	+ <td>executor</td>
	227	+ <td>输出</td>
	228	+ <td>返回op执行器，包含了算子计算流程。</td>
	229	+ <td>-</td>
	230	+ <td>-</td>
	231	+ <td>-</td>
	232	+ <td>-</td>
	233	+ <td>-</td>
	234	+ </tr>
	235	+ </tbody></table>
	236	+
	237	+- 返回值：
	238	+
	239	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
	240	+
	241	+ 第一段接口完成入参校验，出现以下场景时报错：
	242	+
	243	+ <table style="undefined;table-layout: fixed; width: 1048px"><colgroup>
	244	+ <col style="width: 319px">
	245	+ <col style="width: 108px">
	246	+ <col style="width: 621px">
	247	+ </colgroup>
	248	+ <thead>
	249	+ <tr>
	250	+ <th>返回码</th>
	251	+ <th>错误码</th>
	252	+ <th>描述</th>
	253	+ </tr></thead>
	254	+ <tbody>
	255	+ <tr>
	256	+ <td rowspan="2">ACLNN_ERR_PARAM_NULLPTR</td>
	257	+ <td rowspan="2">161001</td>
	258	+ <td>如果传入参数是必选输入，输出或者必选属性，且是空指针，则返回161001。</td>
	259	+ </tr>
	260	+ <tr>
	261	+ <td>传入的roundMode是空指针。</td>
	262	+ </tr>
	263	+ <tr>
	264	+ <td rowspan="5">ACLNN_ERR_PARAM_INVALID</td>
	265	+ <td rowspan="5">161002</td>
	266	+ <td>x、groupIndex、y、mxscale的数据类型不在支持的范围之内。</td>
	267	+ </tr>
	268	+ <tr>
	269	+ <td>x、y和mxscale的shape不满足校验条件。</td>
	270	+ </tr>
	271	+ <tr>
	272	+ <td>x、groupIndex、y和mxscale的维度不在支持的范围之内。</td>
	273	+ </tr>
	274	+ <tr>
	275	+ <td>roundMode、dstType、blocksize、scaleAlg、dstTypeMax不符合当前支持的值。</td>
	276	+ </tr>
	277	+ <tr>
	278	+ <td>mxscale不支持非连续的Tensor。</td>
	279	+ </tr>
	280	+ <tr>
	281	+ <td rowspan="1">ACLNN_ERR_PARAM_NULLPTR</td>
	282	+ <td rowspan="1">361001</td>
	283	+ <td>当前平台不在支持的平台范围内。</td>
	284	+ </tr>
	285	+ </tbody></table>
	286	+
	287	+## aclnnGroupedDynamicMxQuantV2
	288	+
	289	+- 参数说明：
	290	+
	291	+ <table style="undefined;table-layout: fixed; width: 1033px"><colgroup>
	292	+ <col style="width: 173px">
	293	+ <col style="width: 112px">
	294	+ <col style="width: 748px">
	295	+ </colgroup>
	296	+ <thead>
	297	+ <tr>
	298	+ <th>参数名</th>
	299	+ <th>输入/输出</th>
	300	+ <th>描述</th>
	301	+ </tr></thead>
	302	+ <tbody>
	303	+ <tr>
	304	+ <td>workspace</td>
	305	+ <td>输入</td>
	306	+ <td>在Device侧申请的workspace内存地址。</td>
	307	+ </tr>
	308	+ <tr>
	309	+ <td>workspaceSize</td>
	310	+ <td>输入</td>
	311	+ <td>在Device侧申请的workspace大小，由第一段接口aclnnGroupedDynamicMxQuantV2GetWorkspaceSize获取。</td>
	312	+ </tr>
	313	+ <tr>
	314	+ <td>executor</td>
	315	+ <td>输入</td>
	316	+ <td>op执行器，包含了算子计算流程。</td>
	317	+ </tr>
	318	+ <tr>
	319	+ <td>stream</td>
	320	+ <td>输入</td>
	321	+ <td>指定执行任务的Stream。</td>
	322	+ </tr>
	323	+ </tbody>
	324	+ </table>
	325	+
	326	+- 返回值：
	327	+
	328	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
	329	+
	330	+## 约束说明
	331	+
	332	+- 确定性计算：
	333	+ - aclnnGroupedDynamicMxQuantV2默认确定性实现。
	334	+
	335	+## 调用示例
	336	+
	337	+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
	338	+
	339	+```Cpp
	340	+#include <iostream>
	341	+#include <memory>
	342	+#include <vector>
	343	+
	344	+#include "acl/acl.h"
	345	+#include "aclnnop/aclnn_grouped_dynamic_mx_quant_v2.h"
	346	+
	347	+#define CHECK_RET(cond, return_expr) \
	348	+ do { \
	349	+ if (!(cond)) { \
	350	+ return_expr; \
	351	+ } \
	352	+ } while (0)
	353	+
	354	+#define CHECK_FREE_RET(cond, return_expr) \
	355	+ do { \
	356	+ if (!(cond)) { \
	357	+ Finalize(deviceId, stream); \
	358	+ return_expr; \
	359	+ } \
	360	+ } while (0)
	361	+
	362	+#define LOG_PRINT(message, ...) \
	363	+ do { \
	364	+ printf(message, ##__VA_ARGS__); \
	365	+ } while (0)
	366	+
	367	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
	368	+{
	369	+ int64_t shapeSize = 1;
	370	+ for (auto i : shape) {
	371	+ shapeSize *= i;
	372	+ }
	373	+ return shapeSize;
	374	+}
	375	+
	376	+int Init(int32_t deviceId, aclrtStream* stream)
	377	+{
	378	+ // 固定写法，资源初始化
	379	+ auto ret = aclInit(nullptr);
	380	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	381	+ ret = aclrtSetDevice(deviceId);
	382	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	383	+ ret = aclrtCreateStream(stream);
	384	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	385	+ return 0;
	386	+}
	387	+
	388	+template <typename T>
	389	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void deviceAddr, aclDataType dataType, aclTensor tensor)
	390	+{
	391	+ auto size = GetShapeSize(shape) * sizeof(T);
	392	+ // 调用aclrtMalloc申请device侧内存
	393	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	394	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	395	+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
	396	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	397	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	398	+
	399	+ // 计算连续tensor的strides
	400	+ std::vector<int64_t> strides(shape.size(), 1);
	401	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	402	+ strides[i] = shape[i + 1] * strides[i + 1];
	403	+ }
	404	+
	405	+ // 调用aclCreateTensor接口创建aclTensor
	406	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	407	+ shape.data(), shape.size(), *deviceAddr);
	408	+ return 0;
	409	+}
	410	+
	411	+void Finalize(int32_t deviceId, aclrtStream stream)
	412	+{
	413	+ aclrtDestroyStream(stream);
	414	+ aclrtResetDevice(deviceId);
	415	+ aclFinalize();
	416	+}
	417	+
	418	+int aclnnGroupedDynamicMxQuantV2Test(int32_t deviceId, aclrtStream& stream)
	419	+{
	420	+ auto ret = Init(deviceId, &stream);
	421	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	422	+
	423	+ // 2. 构造输入与输出，需要根据API的接口自定义构造
	424	+ std::vector<int64_t> xShape = {8, 1};
	425	+ std::vector<int64_t> groupedIndexShape = {2};
	426	+ std::vector<int64_t> yOutShape = {8, 1};
	427	+ std::vector<int64_t> mxscaleOutShape = {2, 1, 2};
	428	+ void* xDeviceAddr = nullptr;
	429	+ void* groupedIndexDeviceAddr = nullptr;
	430	+ void* yOutDeviceAddr = nullptr;
	431	+ void* mxscaleOutDeviceAddr = nullptr;
	432	+ aclTensor* x = nullptr;
	433	+ aclTensor* groupedIndex = nullptr;
	434	+ aclTensor* yOut = nullptr;
	435	+ aclTensor* mxscaleOut = nullptr;
	436	+ //对应BF16的值(0, 8, 64, 512)
	437	+ std::vector<uint16_t> xHostData = {{0}, {16640}, {17024}, {17408}, {0}, {16640}, {17024}, {17408}};
	438	+
	439	+ std::vector<uint32_t> groupedIndexHostData = {4, 8};
	440	+ //对应float8_e4m3的值(0, 4, 32, 256)
	441	+ std::vector<uint8_t> yOutHostData = {{0}, {72}, {96}, {120}, {0}, {72}, {96}, {120}};
	442	+ //对应float8_e8m0的值(2)
	443	+ std::vector<std::vector<uint8_t>> mxscaleOutHostData = {{{128, 0}}, {{128, 0}}};
	444	+ const char* roundModeOptional = "rint";
	445	+ int64_t dstType = 36;
	446	+ int64_t blocksize = 32;
	447	+ int64_t scaleAlg = 0;
	448	+ double dstTypeMax = 0.0;
	449	+ // 创建x aclTensor
	450	+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_BF16, &x);
	451	+ std::unique_ptr<aclTensor, aclnnStatus ()(const aclTensor)> xTensorPtr(x, aclDestroyTensor);
	452	+ std::unique_ptr<void, aclError ()(void)> xDeviceAddrPtr(xDeviceAddr, aclrtFree);
	453	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	454	+ // 创建groudedIndex aclTensor
	455	+ ret = CreateAclTensor(groupedIndexHostData, groupedIndexShape, &groupedIndexDeviceAddr, aclDataType::ACL_INT32, &groupedIndex);
	456	+ std::unique_ptr<aclTensor, aclnnStatus ()(const aclTensor)> groupedIndexTensorPtr(groupedIndex, aclDestroyTensor);
	457	+ std::unique_ptr<void, aclError ()(void)> groupedIndexDeviceAddrPtr(groupedIndexDeviceAddr, aclrtFree);
	458	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	459	+ // 创建yOut aclTensor
	460	+ ret = CreateAclTensor(yOutHostData, yOutShape, &yOutDeviceAddr, aclDataType::ACL_FLOAT8_E4M3FN, &yOut);
	461	+ std::unique_ptr<aclTensor, aclnnStatus ()(const aclTensor)> yOutTensorPtr(yOut, aclDestroyTensor);
	462	+ std::unique_ptr<void, aclError ()(void)> yOutDeviceAddrPtr(yOutDeviceAddr, aclrtFree);
	463	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	464	+ // 创建mxscaleOut aclTensor
	465	+ ret = CreateAclTensor(mxscaleOutHostData, mxscaleOutShape, &mxscaleOutDeviceAddr, aclDataType::ACL_FLOAT8_E8M0, &mxscaleOut);
	466	+ std::unique_ptr<aclTensor, aclnnStatus ()(const aclTensor)> mxscaleOutTensorPtr(mxscaleOut, aclDestroyTensor);
	467	+ std::unique_ptr<void, aclError ()(void)> mxscaleOutDeviceAddrPtr(mxscaleOutDeviceAddr, aclrtFree);
	468	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	469	+
	470	+ // 调用CANN算子库API，需要修改为具体的Api名称
	471	+ uint64_t workspaceSize = 0;
	472	+ aclOpExecutor* executor;
	473	+
	474	+ // 调用aclnnGroupedDynamicMxQuantV2第一段接口
	475	+ ret = aclnnGroupedDynamicMxQuantV2GetWorkspaceSize(x, groupedIndex, roundModeOptional, dstType, blocksize, scaleAlg, dstTypeMax,
	476	+ yOut, mxscaleOut, &workspaceSize, &executor);
	477	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuantV2GetWorkspaceSize failed. ERROR: %d\n", ret);
	478	+ return ret);
	479	+ // 根据第一段接口计算出的workspaceSize申请device内存
	480	+ void* workspaceAddr = nullptr;
	481	+ std::unique_ptr<void, aclError ()(void)> workspaceAddrPtr(nullptr, aclrtFree);
	482	+ if (workspaceSize > 0) {
	483	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	484	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
	485	+ workspaceAddrPtr.reset(workspaceAddr);
	486	+ }
	487	+ // 调用aclnnGroupedDynamicMxQuantV2第二段接口
	488	+ ret = aclnnGroupedDynamicMxQuantV2(workspaceAddr, workspaceSize, executor, stream);
	489	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuantV2 failed. ERROR: %d\n", ret); return ret);
	490	+
	491	+ //（固定写法）同步等待任务执行结束
	492	+ ret = aclrtSynchronizeStream(stream);
	493	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	494	+
	495	+ // 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
	496	+ auto size = GetShapeSize(yOutShape);
	497	+ std::vector<uint8_t> yOutData(
	498	+ size, 0); // C语言中无法直接打印fp4的数据，需要用uint8读出来，自行通过二进制转成fp4
	499	+ ret = aclrtMemcpy(yOutData.data(), yOutData.size() * sizeof(yOutData[0]), yOutDeviceAddr,
	500	+ size * sizeof(yOutData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
	501	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy yOut from device to host failed. ERROR: %d\n", ret);
	502	+ return ret);
	503	+ for (int64_t i = 0; i < size; i++) {
	504	+ LOG_PRINT("y[%ld] is: %d\n", i, yOutData[i]);
	505	+ }
	506	+ size = GetShapeSize(mxscaleOutShape);
	507	+ std::vector<uint8_t> mxscaleOutData(
	508	+ size, 0); // C语言中无法直接打印fp8的数据，需要用uint8读出来，自行通过二进制转成fp8
	509	+ ret = aclrtMemcpy(mxscaleOutData.data(), mxscaleOutData.size() * sizeof(mxscaleOutData[0]), mxscaleOutDeviceAddr,
	510	+ size * sizeof(mxscaleOutData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
	511	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy mxscaleOut from device to host failed. ERROR: %d\n", ret);
	512	+ return ret);
	513	+ for (int64_t i = 0; i < size; i++) {
	514	+ LOG_PRINT("mxscaleOut[%ld] is: %d\n", i, mxscaleOutData[i]);
	515	+ }
	516	+ return ACL_SUCCESS;
	517	+}
	518	+
	519	+int main()
	520	+{
	521	+ // 1. （固定写法）device/stream初始化，参考acl API手册
	522	+ // 根据自己的实际device填写deviceId
	523	+ int32_t deviceId = 0;
	524	+ aclrtStream stream;
	525	+ auto ret = aclnnGroupedDynamicMxQuantV2Test(deviceId, stream);
	526	+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuantV2Test failed. ERROR: %d\n", ret); return ret);
	527	+
	528	+ Finalize(deviceId, stream);
	529	+ return 0;
	530	+}
	531	+```

Mquant/grouped_dynamic_mx_quant/examples/arch35/test_aclnn_grouped_dynamic_mx_quant.cpp+3-3

	@@ -145,7 +145,7 @@ int aclnnGroupedDynamicMxQuantTest(int32_t deviceId, aclrtStream& stream)
145	uint64_t workspaceSize = 0;	145	uint64_t workspaceSize = 0;
146	aclOpExecutor* executor;	146	aclOpExecutor* executor;
147		147
148	- // 调用aclnnDynamicMxQuant第一段接口	148	+ // 调用aclnnGroupedDynamicMxQuant第一段接口
149	ret = aclnnGroupedDynamicMxQuantGetWorkspaceSize(x, groupedIndex, roundModeOptional, dstType, blocksize, yOut, mxscaleOut, &workspaceSize, &executor);	149	ret = aclnnGroupedDynamicMxQuantGetWorkspaceSize(x, groupedIndex, roundModeOptional, dstType, blocksize, yOut, mxscaleOut, &workspaceSize, &executor);
150	CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuantGetWorkspaceSize failed. ERROR: %d\n", ret);	150	CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuantGetWorkspaceSize failed. ERROR: %d\n", ret);
151	return ret);	151	return ret);
	@@ -157,9 +157,9 @@ int aclnnGroupedDynamicMxQuantTest(int32_t deviceId, aclrtStream& stream)
157	CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);	157	CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
158	workspaceAddrPtr.reset(workspaceAddr);	158	workspaceAddrPtr.reset(workspaceAddr);
159	}	159	}
160	- // 调用aclnnDynamicMxQuant第二段接口	160	+ // 调用aclnnGroupedDynamicMxQuant第二段接口
161	ret = aclnnGroupedDynamicMxQuant(workspaceAddr, workspaceSize, executor, stream);	161	ret = aclnnGroupedDynamicMxQuant(workspaceAddr, workspaceSize, executor, stream);
162	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDynamicMxQuant failed. ERROR: %d\n", ret); return ret);	162	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuant failed. ERROR: %d\n", ret); return ret);
163		163
164	//（固定写法）同步等待任务执行结束	164	//（固定写法）同步等待任务执行结束
165	ret = aclrtSynchronizeStream(stream);	165	ret = aclrtSynchronizeStream(stream);

Aquant/grouped_dynamic_mx_quant/examples/arch35/test_aclnn_grouped_dynamic_mx_quant_v2.cpp+206-0

@@ -0,0 +1,206 @@
	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 test_aclnn_grouped_dynamic_mx_quant.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include <iostream>
	17	+#include <memory>
	18	+#include <vector>
	19	+
	20	+#include "acl/acl.h"
	21	+#include "aclnnop/aclnn_grouped_dynamic_mx_quant_v2.h"
	22	+
	23	+#define CHECK_RET(cond, return_expr) \
	24	+ do { \
	25	+ if (!(cond)) { \
	26	+ return_expr; \
	27	+ } \
	28	+ } while (0)
	29	+
	30	+#define CHECK_FREE_RET(cond, return_expr) \
	31	+ do { \
	32	+ if (!(cond)) { \
	33	+ Finalize(deviceId, stream); \
	34	+ return_expr; \
	35	+ } \
	36	+ } while (0)
	37	+
	38	+#define LOG_PRINT(message, ...) \
	39	+ do { \
	40	+ printf(message, ##__VA_ARGS__); \
	41	+ } while (0)
	42	+
	43	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
	44	+{
	45	+ int64_t shapeSize = 1;
	46	+ for (auto i : shape) {
	47	+ shapeSize *= i;
	48	+ }
	49	+ return shapeSize;
	50	+}
	51	+
	52	+int Init(int32_t deviceId, aclrtStream* stream)
	53	+{
	54	+ // 固定写法，资源初始化
	55	+ auto ret = aclInit(nullptr);
	56	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	57	+ ret = aclrtSetDevice(deviceId);
	58	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	59	+ ret = aclrtCreateStream(stream);
	60	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	61	+ return 0;
	62	+}
	63	+
	64	+template <typename T>
	65	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void deviceAddr, aclDataType dataType, aclTensor tensor)
	66	+{
	67	+ auto size = GetShapeSize(shape) * sizeof(T);
	68	+ // 调用aclrtMalloc申请device侧内存
	69	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	70	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	71	+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
	72	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	73	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	74	+
	75	+ // 计算连续tensor的strides
	76	+ std::vector<int64_t> strides(shape.size(), 1);
	77	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	78	+ strides[i] = shape[i + 1] * strides[i + 1];
	79	+ }
	80	+
	81	+ // 调用aclCreateTensor接口创建aclTensor
	82	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	83	+ shape.data(), shape.size(), *deviceAddr);
	84	+ return 0;
	85	+}
	86	+
	87	+void Finalize(int32_t deviceId, aclrtStream stream)
	88	+{
	89	+ aclrtDestroyStream(stream);
	90	+ aclrtResetDevice(deviceId);
	91	+ aclFinalize();
	92	+}
	93	+
	94	+int aclnnGroupedDynamicMxQuantV2Test(int32_t deviceId, aclrtStream& stream)
	95	+{
	96	+ auto ret = Init(deviceId, &stream);
	97	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	98	+
	99	+ // 2. 构造输入与输出，需要根据API的接口自定义构造
	100	+ std::vector<int64_t> xShape = {8, 1};
	101	+ std::vector<int64_t> groupedIndexShape = {2};
	102	+ std::vector<int64_t> yOutShape = {8, 1};
	103	+ std::vector<int64_t> mxscaleOutShape = {2, 1, 2};
	104	+ void* xDeviceAddr = nullptr;
	105	+ void* groupedIndexDeviceAddr = nullptr;
	106	+ void* yOutDeviceAddr = nullptr;
	107	+ void* mxscaleOutDeviceAddr = nullptr;
	108	+ aclTensor* x = nullptr;
	109	+ aclTensor* groupedIndex = nullptr;
	110	+ aclTensor* yOut = nullptr;
	111	+ aclTensor* mxscaleOut = nullptr;
	112	+ //对应BF16的值(0, 8, 64, 512)
	113	+ std::vector<uint16_t> xHostData = {{0}, {16640}, {17024}, {17408}, {0}, {16640}, {17024}, {17408}};
	114	+
	115	+ std::vector<uint32_t> groupedIndexHostData = {4, 8};
	116	+ //对应float8_e4m3的值(0, 4, 32, 256)
	117	+ std::vector<uint8_t> yOutHostData = {{0}, {72}, {96}, {120}, {0}, {72}, {96}, {120}};
	118	+ //对应float8_e8m0的值(2)
	119	+ std::vector<std::vector<uint8_t>> mxscaleOutHostData = {{{128, 0}}, {{128, 0}}};
	120	+ const char* roundModeOptional = "rint";
	121	+ int64_t dstType = 36;
	122	+ int64_t blocksize = 32;
	123	+ int64_t scaleAlg = 0;
	124	+ double dstTypeMax = 0.0;
	125	+ // 创建x aclTensor
	126	+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_BF16, &x);
	127	+ std::unique_ptr<aclTensor, aclnnStatus ()(const aclTensor)> xTensorPtr(x, aclDestroyTensor);
	128	+ std::unique_ptr<void, aclError ()(void)> xDeviceAddrPtr(xDeviceAddr, aclrtFree);
	129	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	130	+ // 创建groudedIndex aclTensor
	131	+ ret = CreateAclTensor(groupedIndexHostData, groupedIndexShape, &groupedIndexDeviceAddr, aclDataType::ACL_INT32, &groupedIndex);
	132	+ std::unique_ptr<aclTensor, aclnnStatus ()(const aclTensor)> groupedIndexTensorPtr(groupedIndex, aclDestroyTensor);
	133	+ std::unique_ptr<void, aclError ()(void)> groupedIndexDeviceAddrPtr(groupedIndexDeviceAddr, aclrtFree);
	134	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	135	+ // 创建yOut aclTensor
	136	+ ret = CreateAclTensor(yOutHostData, yOutShape, &yOutDeviceAddr, aclDataType::ACL_FLOAT8_E4M3FN, &yOut);
	137	+ std::unique_ptr<aclTensor, aclnnStatus ()(const aclTensor)> yOutTensorPtr(yOut, aclDestroyTensor);
	138	+ std::unique_ptr<void, aclError ()(void)> yOutDeviceAddrPtr(yOutDeviceAddr, aclrtFree);
	139	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	140	+ // 创建mxscaleOut aclTensor
	141	+ ret = CreateAclTensor(mxscaleOutHostData, mxscaleOutShape, &mxscaleOutDeviceAddr, aclDataType::ACL_FLOAT8_E8M0, &mxscaleOut);
	142	+ std::unique_ptr<aclTensor, aclnnStatus ()(const aclTensor)> mxscaleOutTensorPtr(mxscaleOut, aclDestroyTensor);
	143	+ std::unique_ptr<void, aclError ()(void)> mxscaleOutDeviceAddrPtr(mxscaleOutDeviceAddr, aclrtFree);
	144	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	145	+
	146	+ // 调用CANN算子库API，需要修改为具体的Api名称
	147	+ uint64_t workspaceSize = 0;
	148	+ aclOpExecutor* executor;
	149	+
	150	+ // 调用aclnnGroupedDynamicMxQuantV2第一段接口
	151	+ ret = aclnnGroupedDynamicMxQuantV2GetWorkspaceSize(x, groupedIndex, roundModeOptional, dstType, blocksize, scaleAlg, dstTypeMax,
	152	+ yOut, mxscaleOut, &workspaceSize, &executor);
	153	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuantV2GetWorkspaceSize failed. ERROR: %d\n", ret);
	154	+ return ret);
	155	+ // 根据第一段接口计算出的workspaceSize申请device内存
	156	+ void* workspaceAddr = nullptr;
	157	+ std::unique_ptr<void, aclError ()(void)> workspaceAddrPtr(nullptr, aclrtFree);
	158	+ if (workspaceSize > 0) {
	159	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	160	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
	161	+ workspaceAddrPtr.reset(workspaceAddr);
	162	+ }
	163	+ // 调用aclnnGroupedDynamicMxQuantV2第二段接口
	164	+ ret = aclnnGroupedDynamicMxQuantV2(workspaceAddr, workspaceSize, executor, stream);
	165	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuantV2 failed. ERROR: %d\n", ret); return ret);
	166	+
	167	+ //（固定写法）同步等待任务执行结束
	168	+ ret = aclrtSynchronizeStream(stream);
	169	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	170	+
	171	+ // 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
	172	+ auto size = GetShapeSize(yOutShape);
	173	+ std::vector<uint8_t> yOutData(
	174	+ size, 0); // C语言中无法直接打印fp4的数据，需要用uint8读出来，自行通过二进制转成fp4
	175	+ ret = aclrtMemcpy(yOutData.data(), yOutData.size() * sizeof(yOutData[0]), yOutDeviceAddr,
	176	+ size * sizeof(yOutData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
	177	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy yOut from device to host failed. ERROR: %d\n", ret);
	178	+ return ret);
	179	+ for (int64_t i = 0; i < size; i++) {
	180	+ LOG_PRINT("y[%ld] is: %d\n", i, yOutData[i]);
	181	+ }
	182	+ size = GetShapeSize(mxscaleOutShape);
	183	+ std::vector<uint8_t> mxscaleOutData(
	184	+ size, 0); // C语言中无法直接打印fp8的数据，需要用uint8读出来，自行通过二进制转成fp8
	185	+ ret = aclrtMemcpy(mxscaleOutData.data(), mxscaleOutData.size() * sizeof(mxscaleOutData[0]), mxscaleOutDeviceAddr,
	186	+ size * sizeof(mxscaleOutData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
	187	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy mxscaleOut from device to host failed. ERROR: %d\n", ret);
	188	+ return ret);
	189	+ for (int64_t i = 0; i < size; i++) {
	190	+ LOG_PRINT("mxscaleOut[%ld] is: %d\n", i, mxscaleOutData[i]);
	191	+ }
	192	+ return ACL_SUCCESS;
	193	+}
	194	+
	195	+int main()
	196	+{
	197	+ // 1. （固定写法）device/stream初始化，参考acl API手册
	198	+ // 根据自己的实际device填写deviceId
	199	+ int32_t deviceId = 0;
	200	+ aclrtStream stream;
	201	+ auto ret = aclnnGroupedDynamicMxQuantV2Test(deviceId, stream);
	202	+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGroupedDynamicMxQuantV2Test failed. ERROR: %d\n", ret); return ret);
	203	+
	204	+ Finalize(deviceId, stream);
	205	+ return 0;
	206	+}

Mquant/grouped_dynamic_mx_quant/op_api/aclnn_grouped_dynamic_mx_quant.cpp+4-2

Aquant/grouped_dynamic_mx_quant/op_api/aclnn_grouped_dynamic_mx_quant_v2.cpp+203-0

@@ -0,0 +1,203 @@
	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	+#include <dlfcn.h>
	11	+#include <new>
	12	+#include "aclnn_grouped_dynamic_mx_quant_v2.h"
	13	+#include "grouped_dynamic_mx_quant.h"
	14	+#include "level0/fault_injection.h"
	15	+#include "aclnn_kernels/contiguous.h"
	16	+#include "aclnn_kernels/reshape.h"
	17	+#include "aclnn/aclnn_base.h"
	18	+#include "aclnn_kernels/common/op_error_check.h"
	19	+#include "acl/acl.h"
	20	+#include "opdev/common_types.h"
	21	+#include "opdev/data_type_utils.h"
	22	+#include "opdev/format_utils.h"
	23	+#include "opdev/op_dfx.h"
	24	+#include "opdev/op_executor.h"
	25	+#include "opdev/op_log.h"
	26	+#include "opdev/platform.h"
	27	+#include "opdev/shape_utils.h"
	28	+#include "opdev/tensor_view_utils.h"
	29	+#include "opdev/make_op_executor.h"
	30	+#include "op_api/aclnn_util.h"
	31	+
	32	+using namespace op;
	33	+
	34	+#ifdef __cplusplus
	35	+extern "C" {
	36	+#endif
	37	+static constexpr int64_t X_DIM_NUM = 2;
	38	+static constexpr int64_t NUM_TWO = 2;
	39	+static constexpr int64_t SCALE_DIM_NUM = 3;
	40	+static constexpr uint64_t NUM_ZERO = 0;
	41	+
	42	+static const std::initializer_list<op::DataType> X_DTYPE_SUPPORT_LIST = {op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
	43	+
	44	+static const std::initializer_list<op::DataType> GROUP_INDEX_DTYPE_SUPPORT_LIST = {op::DataType::DT_INT32};
	45	+
	46	+static const std::initializer_list<op::DataType> OUTPUT_DTYPE_SUPPORT_LIST = {
	47	+ op::DataType::DT_FLOAT8_E4M3FN, op::DataType::DT_FLOAT8_E5M2};
	48	+
	49	+static const std::initializer_list<op::DataType> MXSCALE_DTYPE_SUPPORT_LIST = {
	50	+ op::DataType::DT_FLOAT8_E8M0};
	51	+
	52	+static inline bool CheckNotNull(const aclTensor* x, const aclTensor* groupIndex, const char* roundMode, const aclTensor* y, const aclTensor* mxscale) {
	53	+ OP_CHECK_NULL(x, return false);
	54	+ OP_CHECK_NULL(groupIndex, return false);
	55	+ if (roundMode == nullptr) {
	56	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "roundMode cannot be nullptr");
	57	+ return false;
	58	+ }
	59	+ OP_CHECK_NULL(y, return false);
	60	+ OP_CHECK_NULL(mxscale, return false);
	61	+ return true;
	62	+}
	63	+
	64	+static bool CheckShape(const aclTensor* x, const aclTensor* groupIndex, int64_t blocksize, int64_t scaleAlg, double dstTypeMax, const aclTensor* y, const aclTensor* mxscale) {
	CANN-robot4月7日函数参数不匹配: 函数CheckShape的声明参数列表为(const aclTensor* x, const aclTensor* groupIndex, int64_t blocksize, int64_t scaleAlg, double dstTypeMax, const aclTensor* y, const aclTensor* mxscale)，但在第64行定义时，参数列表缺少了scaleAlg和dstTypeMax。这会导致编译错误或运行时参数传递错误。问题类型: 函数参数不匹配文件路径: `quant/grouped_dynamic_mx_quant/op_api/aclnn_grouped_dynamic_mx_quant_v2.cpp` 行号: 64 问题代码: static bool CheckShape(const aclTensor* x, const aclTensor* groupIndex, int64_t blocksize, int64_t scaleAlg, double dstTypeMax, const aclTensor* y, const aclTensor* mxscale) { 修改建议: 修正函数定义，确保参数列表与声明一致。正确的定义应为：static bool CheckShape(const aclTensor* x, const aclTensor* groupIndex, int64_t blocksize, int64_t scaleAlg, double dstTypeMax, const aclTensor* y, const aclTensor* mxscale)。 --- 此评论由代码审查工具自动生成
	65	+ auto xShape = x->GetViewShape();
	66	+ auto groupShape = groupIndex->GetViewShape();
	67	+ auto yShape = y->GetViewShape();
	68	+ auto mxscaleShape = mxscale->GetViewShape();
	69	+ OP_CHECK(xShape.GetDimNum() == X_DIM_NUM,
	70	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "input x Dims is %ld, should be 2D.", xShape.GetDimNum()), return false);
	71	+ OP_CHECK(groupShape.GetDimNum() == 1,
	72	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "input groupIndex Dims is %ld, should be 1D.", groupShape.GetDimNum()), return false);
	73	+ OP_CHECK(yShape.GetDimNum() == X_DIM_NUM,
	74	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "output yShape Dims is %ld, should be 2D.", yShape.GetDimNum()), return false);
	75	+ OP_CHECK(mxscaleShape.GetDimNum() == SCALE_DIM_NUM,
	76	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "output mxscale Dims is %ld, should be 3D.", mxscaleShape.GetDimNum()), return false);
	77	+ OP_CHECK_SHAPE_NOT_EQUAL(y, x, return false);
	78	+ int64_t xDim0 = xShape.GetDim(0);
	79	+ int64_t xDim1 = xShape.GetDim(1);
	80	+ int64_t groupIndexDim0 = groupShape.GetDim(0);
	81	+ int64_t mxscaleDim0 = mxscaleShape.GetDim(0);
	82	+ int64_t mxscaleDim1 = mxscaleShape.GetDim(1);
	83	+ int64_t mxscaleDim2 = mxscaleShape.GetDim(NUM_TWO);
	84	+ int64_t mxscaleDim0Count = (xDim0 / (blocksize * NUM_TWO) + groupIndexDim0);
	85	+ OP_CHECK(mxscaleDim2 == NUM_TWO,
	CANN-robot4月7日错误信息不准确: 在第86行的错误日志中，打印的变量是mxscaleDim1，但实际检查的是mxscaleDim2。这会导致错误信息与实际检查的维度不匹配，误导开发者。问题类型: 错误信息不准确文件路径: `quant/grouped_dynamic_mx_quant/op_api/aclnn_grouped_dynamic_mx_quant_v2.cpp` 行号: 85 问题代码: OP_CHECK(mxscaleDim2 == NUM_TWO, OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mxscale dim2 is %ld, should be 2.", mxscaleDim1), return false); 修改建议: 将错误日志中的变量改为mxscaleDim2，以准确反映检查的维度。修改为：OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mxscale dim2 is %ld, should be 2.", mxscaleDim2)。 --- 此评论由代码审查工具自动生成
	86	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mxscale dim2 is %ld, should be 2.", mxscaleDim1), return false);
	87	+ OP_CHECK(
	88	+ xDim1 == mxscaleDim1,
	89	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mxscale dim1 is %ld, should be same as x dim1 (%ld).", mxscaleDim1, xDim1),
	90	+ return false);
	91	+ OP_CHECK(mxscaleDim0 == mxscaleDim0Count,
	92	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mxscale dim0 is %ld, should be same with mxscaleDim0Count (%ld).", mxscaleDim0, mxscaleDim0Count), return false);
	93	+ return true;
	94	+}
	95	+
	96	+static bool CheckDtypeValid(const aclTensor* x, const aclTensor* groupIndex, const char* roundMode, int64_t dstType,
	97	+ int64_t blocksize, int64_t scaleAlg, double dstTypeMax, const aclTensor* y, const aclTensor* mxscale) {
	98	+ // 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
	99	+ bool IsRegbaseSocVersion = Ops::NN::AclnnUtil::IsRegbase();
	100	+ if (IsRegbaseSocVersion) {
	101	+ OP_CHECK_DTYPE_NOT_SUPPORT(x, X_DTYPE_SUPPORT_LIST, return false);
	102	+ OP_CHECK_DTYPE_NOT_SUPPORT(groupIndex, GROUP_INDEX_DTYPE_SUPPORT_LIST, return false);
	103	+ OP_CHECK_DTYPE_NOT_SUPPORT(y, OUTPUT_DTYPE_SUPPORT_LIST, return false);
	104	+ OP_CHECK_DTYPE_NOT_SUPPORT(mxscale, MXSCALE_DTYPE_SUPPORT_LIST, return false);
	105	+ const std::string mode = std::string(roundMode);
	106	+ OP_CHECK(mode == "rint",
	107	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "expected roundMode equals 'rint', get: %s", mode.c_str()),
	108	+ return false);
	109	+ OP_CHECK(blocksize == 32,
	110	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "blocksize only support '32' now, get: %ld", blocksize),
	111	+ return false);
	112	+ OP_CHECK(scaleAlg == 0 \|\| scaleAlg == 1,
	113	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "scaleAlg only support '0' or '1' now, get: %ld", scaleAlg),
	114	+ return false);
	115	+ OP_CHECK(dstTypeMax == 0.0,
	116	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "dstTypeMax only support '0.0' now, get: %f", dstTypeMax),
	117	+ return false);
	118	+ OP_CHECK(static_cast<int64_t>(y->GetDataType()) == dstType,
	119	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "dstType:%ld(%s) is must be the same as y dtype[%s].",
	120	+ dstType, op::ToString(static_cast<op::DataType>(dstType)).GetString(), op::ToString(y->GetDataType()).GetString()),
	121	+ return false);
	122	+ } else {
	123	+ auto curArch = GetCurrentPlatformInfo().GetCurNpuArch();
	124	+ OP_LOGE(ACLNN_ERR_RUNTIME_ERROR, "support for npuArch %u is not implemented",
	125	+ static_cast<uint32_t>(curArch));
	126	+ return false;
	127	+ }
	128	+ return true;
	129	+}
	130	+
	131	+inline static aclnnStatus CheckParams(const aclTensor* x, const aclTensor* groupIndex, const char* roundMode, int64_t dstType,
	132	+ int64_t blocksize, int64_t scaleAlg, double dstTypeMax, const aclTensor* y, const aclTensor* mxscale) {
	133	+ CHECK_RET(CheckNotNull(x, groupIndex, roundMode, y, mxscale), ACLNN_ERR_PARAM_NULLPTR);
	134	+ CHECK_RET(CheckDtypeValid(x, groupIndex, roundMode, dstType, blocksize, scaleAlg, dstTypeMax, y, mxscale), ACLNN_ERR_PARAM_INVALID);
	135	+ CHECK_RET(CheckShape(x, groupIndex, blocksize, scaleAlg, dstTypeMax, y, mxscale), ACLNN_ERR_PARAM_INVALID);
	136	+ return ACLNN_SUCCESS;
	137	+}
	138	+
	139	+aclnnStatus aclnnGroupedDynamicMxQuantV2GetWorkspaceSize(const aclTensor* x, const aclTensor* groupIndex, const char* roundMode,
	140	+ int64_t dstType, int64_t blocksize, int64_t scaleAlg, double dstTypeMax,
	141	+ const aclTensor* y, const aclTensor* mxscale, uint64_t* workspaceSize,
	142	+ aclOpExecutor** executor) {
	143	+ L2_DFX_PHASE_1(aclnnGroupedDynamicMxQuantV2, DFX_IN(x, groupIndex, roundMode, dstType, blocksize, scaleAlg, dstTypeMax),
	144	+ DFX_OUT(y, mxscale));
	145	+ // 固定写法，创建OpExecutor
	146	+ auto uniqueExecutor = CREATE_EXECUTOR();
	147	+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
	148	+
	149	+ // 固定写法，参数检查
	150	+ auto ret = CheckParams(x, groupIndex, roundMode, dstType, blocksize, scaleAlg, dstTypeMax, y, mxscale);
	151	+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
	152	+
	153	+ // 空Tensor处理
	154	+ if (groupIndex->IsEmpty()) {
	155	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "groupIndex does not support empty values.");
	156	+ return ACLNN_ERR_PARAM_INVALID;
	157	+ }
	158	+
	159	+ if (x->IsEmpty()) {
	160	+ *workspaceSize = NUM_ZERO;
	161	+ uniqueExecutor.ReleaseTo(executor);
	162	+ return ACLNN_SUCCESS;
	163	+ }
	164	+
	165	+ // x如果非连续，需要转连续
	166	+ auto selfContiguous = l0op::Contiguous(x, uniqueExecutor.get());
	167	+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	168	+
	169	+ auto groupIndexContiguous = l0op::Contiguous(groupIndex, uniqueExecutor.get());
	170	+ CHECK_RET(groupIndexContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	171	+
	172	+ auto result = l0op::GroupedDynamicMxQuant(selfContiguous, groupIndexContiguous, roundMode, dstType, blocksize,
	173	+ scaleAlg, dstTypeMax, uniqueExecutor.get());
	174	+ const aclTensor *yOut = std::get<0>(result);
	175	+ const aclTensor *mxscaleOut = std::get<1>(result);
	176	+ // 如果出参y是非连续Tensor，需要把计算完的连续Tensor转非连续
	177	+ auto viewCopyResult0 = l0op::ViewCopy(yOut, y, uniqueExecutor.get());
	178	+ CHECK_RET(viewCopyResult0 != nullptr, ACLNN_ERR_INNER_NULLPTR);
	179	+ if (!IsContiguous(mxscale)) {
	180	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "mxscale must be contiguous.");
	181	+ return ACLNN_ERR_PARAM_INVALID;
	182	+ }
	183	+ auto viewCopyResult1 = l0op::ViewCopy(mxscaleOut, mxscale, uniqueExecutor.get());
	184	+ CHECK_RET(viewCopyResult1 != nullptr, ACLNN_ERR_INNER_NULLPTR);
	185	+
	186	+ // 固定写法，获取计算过程中需要使用的workspace大小
	187	+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
	188	+ uniqueExecutor.ReleaseTo(executor); // 需要把 uniqueExecutor持有executor转移给executor
	189	+ return ACLNN_SUCCESS;
	190	+}
	191	+
	192	+aclnnStatus aclnnGroupedDynamicMxQuantV2(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream) {
	193	+ L2_DFX_PHASE_2(aclnnGroupedDynamicMxQuantV2);
	194	+ auto ret = CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
	195	+ if (ret != ACLNN_SUCCESS) {
	196	+ OP_LOGE(ACLNN_ERR_INNER, "This is an error in GroupedDynamicMxQuant launch aicore");
	197	+ return ACLNN_ERR_INNER;
	198	+ }
	199	+ return ACLNN_SUCCESS;
	200	+}
	201	+#ifdef __cplusplus
	202	+}
	203	+#endif

Aquant/grouped_dynamic_mx_quant/op_api/aclnn_grouped_dynamic_mx_quant_v2.h+58-0

@@ -0,0 +1,58 @@
	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	+#ifndef ACLNN_GROUPED_DYNAMIC_MX_QUANT_V2_H_
	11	+#define ACLNN_GROUPED_DYNAMIC_MX_QUANT_V2_H_
	12	+
	13	+#include "aclnn/acl_meta.h"
	14	+
	15	+#ifdef __cplusplus
	16	+extern "C" {
	17	+#endif
	18	+
	19	+/**
	20	+ * @brief aclnnGroupedDynamicMxQuantV2的第一段接口，根据具体的计算流程，计算workspace大小。
	21	+ * @domain aclnn_ops_infer
	22	+ *
	23	+ * @param [in] x: 待进行GroupedDynamicMxQuant计算的入参。npu device侧的aclTensor，
	24	+ * 数据类型支持float16, bfloat16, 数据格式支持ND，支持非连续的Tensor。
	25	+ * @param [in] groupIndex: npu device侧的aclTensor，数据类型支持int32
	26	+ * @param [in] roundMode: host侧的aclScalar，数据类型string，仅支持 "rint"
	27	+ * @param [in] dstType: host侧的aclScalar, 数据类型int, 输入范围为{35, 36}，分别对应输出y的数据类型为{35: FLOAT8_E5M2, 36: FLOAT8_E4M3FN}
	28	+ * @param [in] blocksize: host侧的aclScalar, 数据类型int，仅支持 "32"
	29	+ * @param [in] scaleAlg: host侧的aclScalar, 数据类型int，仅支持 "0"和 "1"
	30	+ * @param [in] dstTypeMax: host侧的aclScalar, 数据类型double，仅支持 "0.0"和 "6.0-12.0"
	31	+ * @param [in] y: GroupedDynamicMxQuant计算的出参。npu device侧的aclTensor，
	32	+ * 数据类型支持float8_e4m3fn, float8_e5m2, 数据格式支持ND，支持非连续的Tensor。
	33	+ * @param [in] mxscale: GroupedDynamicMxQuant计算的出参。npu device侧的aclTensor，
	34	+ * 数据类型支持float8_e8m0, 数据格式支持ND，不支持非连续的Tensor。
	35	+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
	36	+ * @param [out] executor: 返回op执行器，包含算子计算流程。
	37	+ * @return aclnnStatus: 返回状态码。
	38	+ */
	39	+__attribute__((visibility("default"))) aclnnStatus aclnnGroupedDynamicMxQuantV2GetWorkspaceSize(
	40	+ const aclTensor* x, const aclTensor* groupIndex, const char* roundMode, int64_t dstType, int64_t blocksize, int64_t scaleAlg,
	41	+ double dstTypeMax, const aclTensor* y, const aclTensor* mxscale, uint64_t* workspaceSize, aclOpExecutor** executor);
	42	+
	43	+/**
	44	+ * @brief aclnnGroupedDynamicMxQuantV2的第二段接口，用于执行计算。
	45	+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
	46	+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnGroupQuantV2GetWorkspaceSize获取。
	47	+ * @param [in] executor: op执行器，包含了算子计算流程。
	48	+ * @param [in] stream: acl stream流。
	49	+ * @return aclnnStatus: 返回状态码。
	50	+ */
	51	+__attribute__((visibility("default"))) aclnnStatus aclnnGroupedDynamicMxQuantV2(
	52	+ void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
	53	+
	54	+#ifdef __cplusplus
	55	+}
	56	+#endif
	57	+
	58	+#endif // ACLNN_GROUPED_DYNAMIC_MX_QUANT_V2_H_

Mquant/grouped_dynamic_mx_quant/op_api/grouped_dynamic_mx_quant.cpp+5-4

Mquant/grouped_dynamic_mx_quant/op_api/grouped_dynamic_mx_quant.h+3-3

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -15,8 +15,8 @@
15		15
16	namespace l0op {	16	namespace l0op {
17	std::tuple<aclTensor, aclTensor> GroupedDynamicMxQuant(const aclTensor x, const aclTensor groupIndex,	17	std::tuple<aclTensor, aclTensor> GroupedDynamicMxQuant(const aclTensor x, const aclTensor groupIndex,
18	- const char *roundMode, int64_t dstType,	18	+ const char *roundMode, int64_t dstType, int64_t blocksize,
19	- int64_t blocksize, aclOpExecutor *executor);	19	+ int64_t scaleAlg, float dstTypeMax, aclOpExecutor *executor);
20	} // l0op	20	} // l0op
21		21
22	#endif // OP_API_INC_LEVEL0_GROUP_QUANT_H_	22	#endif // OP_API_INC_LEVEL0_GROUP_QUANT_H_

Mquant/grouped_dynamic_mx_quant/op_graph/grouped_dynamic_mx_quant_proto.h+17-6

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -7,12 +7,13 @@
7	* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	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.	8	* See LICENSE in the root of the software repository for the full text of the License.
9	*/	9	*/
		10	+
10	/*!	11	/*!
11	- * \file quantize_ops.h	12	+ * \file grouped_dynamic_mx_quant_proto.h
12	* \brief	13	* \brief
13	*/	14	*/
14	-#ifndef OPS_BUILT_IN_OP_PROTO_INC_QUANTIZE_OPS_H_	15	+#ifndef QUANTIZE_GROUPED_DYNAMIC_MX_QUANT_PROTO_H
15	-#define OPS_BUILT_IN_OP_PROTO_INC_QUANTIZE_OPS_H_	16	+#define QUANTIZE_GROUPED_DYNAMIC_MX_QUANT_PROTO_H
16	#include "graph/operator_reg.h"	17	#include "graph/operator_reg.h"
17		18
18	namespace ge {	19	namespace ge {
	@@ -29,10 +30,18 @@ namespace ge {
29	* @par Attributes:	30	* @par Attributes:
30	* @li round_mode: An optional string, specifying the quantization rounding mode.	31	* @li round_mode: An optional string, specifying the quantization rounding mode.
31	* Defaults and only supports "rint".	32	* Defaults and only supports "rint".
32	-* @li dst_type: An optional int, specifying the dtype of output y.	33	+* @li dst_type: An optional int, specifying the dtype of output y. Target data type enum value:
		34	+* - 35: DT_FLOAT8_E5M2
		35	+* - 36: DT_FLOAT8_E4M3FN
33	* Defaults to FLOAT8_E5M2, only supports FLOAT8_E4M3FN or FLOAT8_E5M2.	36	* Defaults to FLOAT8_E5M2, only supports FLOAT8_E4M3FN or FLOAT8_E5M2.
34	* @li blocksize: An optional int, specifying the block size of quantization.	37	* @li blocksize: An optional int, specifying the block size of quantization.
35	* Defaults and only supports 32.	38	* Defaults and only supports 32.
		39	+* @li scale_alg: An Optional Int. The algorithm for the scale in quantization.
		40	+* Support MxFP4/MxFP8(OCP Microscaling Formats(Mx) Specification, count 0) or MxFP8(nvidia-cuBLAS , count 1).
		41	+* Defaults to 0.
		42	+* @li dst_type_max: An Optional Float. Max_dtype takes the maximum value of the quant_data_type, or the provided value.
		43	+* Only support in FP4_E2M1 mode, with a valid range of 0 or 6 to 12.
		44	+* Defaults to 0.
36		45
37	* @par Outputs:	46	* @par Outputs:
38	* @li y: An output tensor of type FLOAT8_E4M3FN or FLOAT8_E5M2. It has the same shape and rank as input x.	47	* @li y: An output tensor of type FLOAT8_E4M3FN or FLOAT8_E5M2. It has the same shape and rank as input x.
	@@ -52,7 +61,9 @@ REG_OP(GroupedDynamicMxQuant)
52	.ATTR(round_mode, String, "rint")	61	.ATTR(round_mode, String, "rint")
53	.ATTR(dst_type, Int, DT_FLOAT8_E5M2)	62	.ATTR(dst_type, Int, DT_FLOAT8_E5M2)
54	.ATTR(blocksize, Int, 32)	63	.ATTR(blocksize, Int, 32)
		64	+ .ATTR(scale_alg, Int, 0)
		65	+ .ATTR(dst_type_max, Float, 0.0)
55	.OP_END_FACTORY_REG(GroupedDynamicMxQuant)	66	.OP_END_FACTORY_REG(GroupedDynamicMxQuant)
56	} // namespace ge	67	} // namespace ge
57		68
58	-#endif // OPS_BUILT_IN_OP_PROTO_INC_QUANTIZE_OPS_H_	69	+#endif // QUANTIZE_GROUPED_DYNAMIC_MX_QUANT_PROTO_H

Mquant/grouped_dynamic_mx_quant/op_host/arch35/grouped_dynamic_mx_quant_tiling_arch35.cpp+123-81

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -14,21 +14,34 @@
14	*/	14	*/
15		15
16	#include "grouped_dynamic_mx_quant_tiling_arch35.h"	16	#include "grouped_dynamic_mx_quant_tiling_arch35.h"
		17	+#include "quant/grouped_dynamic_mx_quant/op_kernel/arch35/grouped_dynamic_mx_quant_tilingdata.h"
		18	+#include "quant/grouped_dynamic_mx_quant/op_kernel/arch35/grouped_dynamic_mx_quant_struct.h"
17	#include <cmath>	19	#include <cmath>
18	-#include "op_host/util/platform_util.h"	20	+#include "op_common/op_host/util/platform_util.h"
		21	+#include "platform/platform_info.h"
		22	+#include "tiling/platform/platform_ascendc.h"
		23	+#include "platform/platform_infos_def.h"
		24	+#include "error_util.h"
19	#include "graph/utils/type_utils.h"	25	#include "graph/utils/type_utils.h"
20	#include "util/math_util.h"	26	#include "util/math_util.h"
21	#include "register/op_impl_registry.h"	27	#include "register/op_impl_registry.h"
		28	+#include "register/tilingdata_base.h"
22		29
23	using namespace std;	30	using namespace std;
24	using namespace ge;	31	using namespace ge;
25	using namespace Ops::Base;	32	using namespace Ops::Base;
		33	+using namespace GroupedDynamicMxQuantOp;
26		34
27	namespace optiling {	35	namespace optiling {
28	constexpr int64_t INDEX_ATTR_ROUND_MODE = 0;	36	constexpr int64_t INDEX_ATTR_ROUND_MODE = 0;
29	constexpr int64_t INDEX_ATTR_DST_DTYPE = 1;	37	constexpr int64_t INDEX_ATTR_DST_DTYPE = 1;
30	constexpr int64_t INDEX_ATTR_BLOCK_SIZE = 2;	38	constexpr int64_t INDEX_ATTR_BLOCK_SIZE = 2;
		39	+constexpr int64_t INDEX_ATTR_SCALE_ALG = 3;
		40	+constexpr int64_t INDEX_ATTR_DST_TYPE_MAX = 4;
31	constexpr int64_t BYTES_OF_INPUT_TYPE = 2;	41	constexpr int64_t BYTES_OF_INPUT_TYPE = 2;
		42	+constexpr int64_t DIGIT_ZERO= 0;
		43	+constexpr float DIGIT_ZERO_FLOAT= 0.0;
		44	+constexpr int64_t DIGIT_ONE = 1;
32	constexpr int64_t DIGIT_TWO = 2;	45	constexpr int64_t DIGIT_TWO = 2;
33	constexpr int64_t DIGIT_TEN = 10;	46	constexpr int64_t DIGIT_TEN = 10;
34	constexpr int64_t N_BUFFER = 2;	47	constexpr int64_t N_BUFFER = 2;
	@@ -41,9 +54,24 @@ const std::set<ge::DataType> GROUPIDX_SUPPORT_DTYPE_SET = { ge::DT_INT32 };
41	const std::set<ge::DataType> Y_SUPPORT_DTYPE_SET = { ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2 };	54	const std::set<ge::DataType> Y_SUPPORT_DTYPE_SET = { ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2 };
42	const std::set<ge::DataType> OUTPUT_SUPPORT_DTYPE_SET = { ge::DT_FLOAT8_E8M0 };	55	const std::set<ge::DataType> OUTPUT_SUPPORT_DTYPE_SET = { ge::DT_FLOAT8_E8M0 };
43		56
		57	+inline static ge::graphStatus GroupedDynamicMxQuantSetTilingData(gert::TilingContext* context, GroupedDynamicMxQuantTilingData& tilingData)
		58	+{
		59	+ uint64_t tilingDataSize = sizeof(tilingData);
		60	+ OP_CHECK_NULL_WITH_CONTEXT(context, context->GetRawTilingData());
		61	+ auto rawTilingData = context->GetRawTilingData();
		62	+ errno_t ret = memcpy_s(rawTilingData->GetData(), rawTilingData->GetCapacity(),
		63	+ reinterpret_cast<void *>(&tilingData), tilingDataSize);
		64	+ if (ret != EOK) {
		65	+ OP_LOGE(context->GetNodeName(), "memcpy_s failed, ret = %d", ret);
		66	+ return ge::GRAPH_FAILED;
		67	+ }
		68	+ context->GetRawTilingData()->SetDataSize(tilingDataSize);
		69	+ return ge::GRAPH_SUCCESS;
		70	+}
		71	+
44	static ge::graphStatus GetAttr(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)	72	static ge::graphStatus GetAttr(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)
45	{	73	{
46	- OP_LOGD(context->GetNodeName(), "GetAttr begin.");	74	+ OP_LOGD(context, "GetAttr begin.");
47	auto *attrs = context->GetAttrs();	75	auto *attrs = context->GetAttrs();
48	OP_CHECK_NULL_WITH_CONTEXT(context, attrs);	76	OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
49		77
	@@ -51,8 +79,7 @@ static ge::graphStatus GetAttr(const gert::TilingContext *context, GroupedDynami
51	OP_CHECK_NULL_WITH_CONTEXT(context, attrRoundMode);	79	OP_CHECK_NULL_WITH_CONTEXT(context, attrRoundMode);
52	std::string roundModeStr = attrRoundMode;	80	std::string roundModeStr = attrRoundMode;
53	OP_CHECK_IF((roundModeStr != "rint"),	81	OP_CHECK_IF((roundModeStr != "rint"),
54	- OP_LOGE(context->GetNodeName(),	82	+ OP_LOGE(context, "round_mode only supports rint currently, please check."),
55	- "invalid round_mode:%s; round_mode only supports rint", roundModeStr.c_str()),
56	return ge::GRAPH_FAILED);	83	return ge::GRAPH_FAILED);
57		84
58	auto *attrDstType = attrs->GetAttrPointer<int64_t>(INDEX_ATTR_DST_DTYPE);	85	auto *attrDstType = attrs->GetAttrPointer<int64_t>(INDEX_ATTR_DST_DTYPE);
	@@ -60,7 +87,7 @@ static ge::graphStatus GetAttr(const gert::TilingContext *context, GroupedDynami
60	int checkDstType = static_cast<int>(*attrDstType);	87	int checkDstType = static_cast<int>(*attrDstType);
61	OP_CHECK_IF((tilingParam.outDtype == ge::DT_FLOAT8_E4M3FN && checkDstType != 36) \|\|	88	OP_CHECK_IF((tilingParam.outDtype == ge::DT_FLOAT8_E4M3FN && checkDstType != 36) \|\|
62	(tilingParam.outDtype == ge::DT_FLOAT8_E5M2 && checkDstType != 35),	89	(tilingParam.outDtype == ge::DT_FLOAT8_E5M2 && checkDstType != 35),
63	- OP_LOGE(context->GetNodeName(),	90	+ OP_LOGE(context,
64	"y's data type and dst_type is not corresponded, y's data type: FLOAT8_E4M3FN/FLOAT8_E5M2 correspond to dst_type: 36/35."),	91	"y's data type and dst_type is not corresponded, y's data type: FLOAT8_E4M3FN/FLOAT8_E5M2 correspond to dst_type: 36/35."),
65	return ge::GRAPH_FAILED);	92	return ge::GRAPH_FAILED);
66		93
	@@ -68,50 +95,56 @@ static ge::graphStatus GetAttr(const gert::TilingContext *context, GroupedDynami
68	OP_CHECK_NULL_WITH_CONTEXT(context, attrBlockSize);	95	OP_CHECK_NULL_WITH_CONTEXT(context, attrBlockSize);
69	tilingParam.blockSize = static_cast<int64_t>(*attrBlockSize);	96	tilingParam.blockSize = static_cast<int64_t>(*attrBlockSize);
70	OP_CHECK_IF(tilingParam.blockSize != ATTR_BLOCK_SIZE,	97	OP_CHECK_IF(tilingParam.blockSize != ATTR_BLOCK_SIZE,
71	- OP_LOGE(context->GetNodeName(),	98	+ OP_LOGE(context,
72	"The blocksize only supports 32."),	99	"The blocksize only supports 32."),
73	return ge::GRAPH_FAILED);	100	return ge::GRAPH_FAILED);
		101	+
		102	+ auto *attrScaleAlg = attrs->GetAttrPointer<int64_t>(INDEX_ATTR_SCALE_ALG);
		103	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrScaleAlg);
		104	+ tilingParam.scaleAlg = static_cast<int64_t>(*attrScaleAlg);
		105	+ OP_CHECK_IF(tilingParam.scaleAlg != DIGIT_ZERO && tilingParam.scaleAlg != DIGIT_ONE,
		106	+ OP_LOGE(context, "The scale_alg only supports 0 or 1."),
		107	+ return ge::GRAPH_FAILED);
		108	+
		109	+ auto *attrDstTypeMax = attrs->GetAttrPointer<float>(INDEX_ATTR_DST_TYPE_MAX);
		110	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrDstTypeMax);
		111	+ tilingParam.dstTypeMax = static_cast<float>(*attrDstTypeMax);
		112	+ OP_CHECK_IF(tilingParam.dstTypeMax != DIGIT_ZERO_FLOAT,
		113	+ OP_LOGE(context, "The dst_type_max only supports 0.0."),
		114	+ return ge::GRAPH_FAILED);
74		115
75	return ge::GRAPH_SUCCESS;	116	return ge::GRAPH_SUCCESS;
76	}	117	}
77		118
78	static ge::graphStatus CheckDtype(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)	119	static ge::graphStatus CheckDtype(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)
79	{	120	{
80	- OP_LOGD(context->GetNodeName(), "CheckDtype begin.");	121	+ OP_LOGD(context, "CheckDtype begin.");
81	auto inputXPtr = context->GetInputDesc(0);	122	auto inputXPtr = context->GetInputDesc(0);
82	OP_CHECK_NULL_WITH_CONTEXT(context, inputXPtr);	123	OP_CHECK_NULL_WITH_CONTEXT(context, inputXPtr);
83	tilingParam.inDtype = inputXPtr->GetDataType();	124	tilingParam.inDtype = inputXPtr->GetDataType();
84	OP_CHECK_IF(INPUT_SUPPORT_DTYPE_SET.count(tilingParam.inDtype) == 0,	125	OP_CHECK_IF(INPUT_SUPPORT_DTYPE_SET.count(tilingParam.inDtype) == 0,
85	- OP_LOGE(context->GetNodeName(),	126	+ OP_LOGE(context, "Input x dtype only support FLOAT16/BFLOAT16 currently, please check."),
86	- "Input x's data type is [%s], current only supports FLOAT16/BFLOAT16.",
87	- ge::TypeUtils::DataTypeToSerialString(tilingParam.inDtype).c_str()),
88	return ge::GRAPH_FAILED);	127	return ge::GRAPH_FAILED);
89		128
90	auto groupIndexPtr = context->GetInputDesc(1);	129	auto groupIndexPtr = context->GetInputDesc(1);
91	OP_CHECK_NULL_WITH_CONTEXT(context, groupIndexPtr);	130	OP_CHECK_NULL_WITH_CONTEXT(context, groupIndexPtr);
92	auto groupIndexDtype = groupIndexPtr->GetDataType();	131	auto groupIndexDtype = groupIndexPtr->GetDataType();
93	OP_CHECK_IF(GROUPIDX_SUPPORT_DTYPE_SET.count(groupIndexDtype) == 0,	132	OP_CHECK_IF(GROUPIDX_SUPPORT_DTYPE_SET.count(groupIndexDtype) == 0,
94	- OP_LOGE(context->GetNodeName(),	133	+ OP_LOGE(context->GetNodeName(), "Input group_index only supports Int32 currently, please check."),
95	- "group_index's data type is [%s], current only supports Int32.",
96	- ge::TypeUtils::DataTypeToSerialString(groupIndexDtype).c_str()),
97	return ge::GRAPH_FAILED);	134	return ge::GRAPH_FAILED);
98		135
99	auto outputYPtr = context->GetOutputDesc(0);	136	auto outputYPtr = context->GetOutputDesc(0);
100	OP_CHECK_NULL_WITH_CONTEXT(context, outputYPtr);	137	OP_CHECK_NULL_WITH_CONTEXT(context, outputYPtr);
101	tilingParam.outDtype = outputYPtr->GetDataType();	138	tilingParam.outDtype = outputYPtr->GetDataType();
102	OP_CHECK_IF(Y_SUPPORT_DTYPE_SET.count(tilingParam.outDtype) == 0,	139	OP_CHECK_IF(Y_SUPPORT_DTYPE_SET.count(tilingParam.outDtype) == 0,
103	- OP_LOGE(context->GetNodeName(),	140	+ OP_LOGE(context, "Output y only supports FLOAT8_E4M3FN/FLOAT8_E5M2 currently, please check."),
104	- "Output y's data type is [%s], current only supports FLOAT8_E4M3FN/FLOAT8_E5M2.",
105	- ge::TypeUtils::DataTypeToSerialString(tilingParam.outDtype).c_str()),
106	return ge::GRAPH_FAILED);	141	return ge::GRAPH_FAILED);
107		142
108	auto outputMxScalePtr = context->GetOutputDesc(1);	143	auto outputMxScalePtr = context->GetOutputDesc(1);
109	OP_CHECK_NULL_WITH_CONTEXT(context, outputMxScalePtr);	144	OP_CHECK_NULL_WITH_CONTEXT(context, outputMxScalePtr);
110	auto scaleDtype = outputMxScalePtr->GetDataType();	145	auto scaleDtype = outputMxScalePtr->GetDataType();
111	OP_CHECK_IF(OUTPUT_SUPPORT_DTYPE_SET.count(scaleDtype) == 0,	146	OP_CHECK_IF(OUTPUT_SUPPORT_DTYPE_SET.count(scaleDtype) == 0,
112	- OP_LOGE(context->GetNodeName(),	147	+ OP_LOGE(context, "Output mxscale only supports FLOAT8_E8M0 currently, please check."),
113	- "Input mxscale's data type is [%s], current only supports FLOAT8_E8M0.",
114	- ge::TypeUtils::DataTypeToSerialString(scaleDtype).c_str()),
115	return ge::GRAPH_FAILED);	148	return ge::GRAPH_FAILED);
116		149
117	return ge::GRAPH_SUCCESS;	150	return ge::GRAPH_SUCCESS;
	@@ -119,7 +152,7 @@ static ge::graphStatus CheckDtype(const gert::TilingContext *context, GroupedDyn
119		152
120	static ge::graphStatus CheckShape(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)	153	static ge::graphStatus CheckShape(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)
121	{	154	{
122	- OP_LOGD(context->GetNodeName(), "CheckShape begin.");	155	+ OP_LOGD(context, "CheckShape begin.");
123	auto xShapePtr = context->GetInputShape(0);	156	auto xShapePtr = context->GetInputShape(0);
124	OP_CHECK_NULL_WITH_CONTEXT(context, xShapePtr);	157	OP_CHECK_NULL_WITH_CONTEXT(context, xShapePtr);
125	auto xShape = xShapePtr->GetStorageShape();	158	auto xShape = xShapePtr->GetStorageShape();
	@@ -137,37 +170,32 @@ static ge::graphStatus CheckShape(const gert::TilingContext *context, GroupedDyn
137	auto mxScaleShape = mxScaleShapePtr->GetStorageShape();	170	auto mxScaleShape = mxScaleShapePtr->GetStorageShape();
138		171
139	OP_CHECK_IF(xShape != yShape,	172	OP_CHECK_IF(xShape != yShape,
140	- OP_LOGE(context->GetNodeName(),	173	+ OP_LOGE(context, "The shape of output y must be same with shape of input x."),
141	- "The shape of output y must be same with shape of input x."),
142	return ge::GRAPH_FAILED);	174	return ge::GRAPH_FAILED);
143	OP_CHECK_IF(xShape.GetDimNum() != 2,	175	OP_CHECK_IF(xShape.GetDimNum() != 2,
144	- OP_LOGE(context->GetNodeName(),	176	+ OP_LOGE(context, "The shape of input x must be 2-D."),
145	- "The shape of input x must be 2-D."),
146	return ge::GRAPH_FAILED);	177	return ge::GRAPH_FAILED);
147	OP_CHECK_IF(groupIndexShape.GetDimNum() != 1,	178	OP_CHECK_IF(groupIndexShape.GetDimNum() != 1,
148	- OP_LOGE(context->GetNodeName(),	179	+ OP_LOGE(context, "The shape of input group_index must be 1-D."),
149	- "The shape of input group_index must be 1-D."),
150	return ge::GRAPH_FAILED);	180	return ge::GRAPH_FAILED);
151	OP_CHECK_IF(mxScaleShape.GetDimNum() != SCALE_DIM_NUM,	181	OP_CHECK_IF(mxScaleShape.GetDimNum() != SCALE_DIM_NUM,
152	- OP_LOGE(context->GetNodeName(),	182	+ OP_LOGE(context, "The shape of output mxscale must be 3-D."),
153	- "The shape of output mxscale must be 3-D."),
154	return ge::GRAPH_FAILED);	183	return ge::GRAPH_FAILED);
155		184
156	tilingParam.groupSize = groupIndexShape.GetDim(0);	185	tilingParam.groupSize = groupIndexShape.GetDim(0);
157	tilingParam.preAxisSize = xShape.GetDim(0);	186	tilingParam.preAxisSize = xShape.GetDim(0);
158	tilingParam.postAxisSize = xShape.GetDim(1);	187	tilingParam.postAxisSize = xShape.GetDim(1);
159	OP_CHECK_IF(tilingParam.groupSize == 0,	188	OP_CHECK_IF(tilingParam.groupSize == 0,
160	- OP_LOGE(context->GetNodeName(),	189	+ OP_LOGE(context, "group_index does not support empty tensor."),
161	- "group_index does not support empty tensor."),
162	return ge::GRAPH_FAILED);	190	return ge::GRAPH_FAILED);
163		191
164	- xShape.SetDim(0, tilingParam.preAxisSize/(tilingParam.blockSize *DIGIT_TWO)+tilingParam.groupSize);	192	+ xShape.SetDim(0, tilingParam.preAxisSize / (tilingParam.blockSize * DIGIT_TWO) + tilingParam.groupSize);
165	xShape.SetDim(1, tilingParam.postAxisSize * DIGIT_TWO);	193	xShape.SetDim(1, tilingParam.postAxisSize * DIGIT_TWO);
166	OP_CHECK_IF(	194	OP_CHECK_IF(
167	mxScaleShape[0] != xShape[0] \|\| mxScaleShape[1] != tilingParam.postAxisSize \|\|	195	mxScaleShape[0] != xShape[0] \|\| mxScaleShape[1] != tilingParam.postAxisSize \|\|
168	mxScaleShape[SCALE_DIM_NUM - 1] != DIGIT_TWO,	196	mxScaleShape[SCALE_DIM_NUM - 1] != DIGIT_TWO,
169	OP_LOGE(	197	OP_LOGE(
170	- context->GetNodeName(),	198	+ context,
171	"The shape of output mxscale is incorrect, it should be [x.shape[0] / (2 * "	199	"The shape of output mxscale is incorrect, it should be [x.shape[0] / (2 * "
172	"blocksize) + group_index.shape[0], x.shape[1], 2]."),	200	"blocksize) + group_index.shape[0], x.shape[1], 2]."),
173	return ge::GRAPH_FAILED);	201	return ge::GRAPH_FAILED);
	@@ -176,37 +204,43 @@ static ge::graphStatus CheckShape(const gert::TilingContext *context, GroupedDyn
176		204
177	static ge::graphStatus GetPlatInfo(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)	205	static ge::graphStatus GetPlatInfo(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)
178	{	206	{
179	- OP_LOGD(context->GetNodeName(), "GetPlatInfo begin.");	207	+ OP_LOGD(context, "GetPlatInfo begin.");
180	auto platformInfo = context->GetPlatformInfo();	208	auto platformInfo = context->GetPlatformInfo();
181	OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);	209	OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
182	auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);	210	auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
		211	+
183	tilingParam.totalCoreNum = ascendcPlatform.GetCoreNumAiv();	212	tilingParam.totalCoreNum = ascendcPlatform.GetCoreNumAiv();
184	OP_CHECK_IF((tilingParam.totalCoreNum <= 0),	213	OP_CHECK_IF((tilingParam.totalCoreNum <= 0),
185	- OP_LOGE(context->GetNodeName(), "Failed to get core num."), return ge::GRAPH_FAILED);	214	+ OP_LOGE(context, "Failed to get core num."), return ge::GRAPH_FAILED);
		215	+
186	uint64_t ubSize;	216	uint64_t ubSize;
187	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);	217	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
188	tilingParam.ubSize = static_cast<int64_t>(ubSize);	218	tilingParam.ubSize = static_cast<int64_t>(ubSize);
189	OP_CHECK_IF((tilingParam.ubSize <= 0),	219	OP_CHECK_IF((tilingParam.ubSize <= 0),
190	- OP_LOGE(context->GetNodeName(), "Failed to get ub size."), return ge::GRAPH_FAILED);	220	+ OP_LOGE(context, "Failed to get ub size."), return ge::GRAPH_FAILED);
		221	+
191	tilingParam.vfLen = Ops::Base::GetVRegSize(context);	222	tilingParam.vfLen = Ops::Base::GetVRegSize(context);
		223	+ OP_CHECK_IF((tilingParam.ubSize <= 0),
		224	+ OP_LOGE(context, "Failed to get ub size."), return ge::GRAPH_FAILED);
192	return ge::GRAPH_SUCCESS;	225	return ge::GRAPH_SUCCESS;
193	}	226	}
194		227
195	static ge::graphStatus DoTiling(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)	228	static ge::graphStatus DoTiling(const gert::TilingContext *context, GroupedDynamicMxQuantTilingParam &tilingParam)
196	{	229	{
197	- OP_LOGD(context->GetNodeName(), "DoTiling begin.");	230	+ OP_LOGD(context, "DoTiling begin.");
198	// 计算tilingkey	231	// 计算tilingkey
199	// 十位数为1、2，分别表示输入类型是float16、bfloat16;	232	// 十位数为1、2，分别表示输入类型是float16、bfloat16;
200	- int64_t hundredDigit = tilingParam.inDtype == DT_FLOAT16 ? 1 : DIGIT_TWO;	233	+ int64_t hundredDigit = tilingParam.inDtype == DT_FLOAT16 ? DIGIT_ONE : DIGIT_TWO;
201	// 个位数为1、2，分别表示输出类型是float8_e4m3fn、float8_e5m2	234	// 个位数为1、2，分别表示输出类型是float8_e4m3fn、float8_e5m2
202	- int64_t tenDigit = tilingParam.outDtype == DT_FLOAT8_E4M3FN? 1: 2;	235	+ int64_t tenDigit = tilingParam.outDtype == DT_FLOAT8_E4M3FN ? DIGIT_ONE : DIGIT_TWO;
203	tilingParam.tilingKey = hundredDigit * DIGIT_TEN + tenDigit;	236	tilingParam.tilingKey = hundredDigit * DIGIT_TEN + tenDigit;
204		237
205	// 计算ubFactor	238	// 计算ubFactor
206	const int64_t cacheline = static_cast<int64_t>(tilingParam.vfLen / BYTES_OF_INPUT_TYPE);	239	const int64_t cacheline = static_cast<int64_t>(tilingParam.vfLen / BYTES_OF_INPUT_TYPE);
207	int64_t maxUbAvailable = tilingParam.ubSize / N_BUFFER / EXIST_NODE_NUM;	240	int64_t maxUbAvailable = tilingParam.ubSize / N_BUFFER / EXIST_NODE_NUM;
208	// 按照2倍blocksize对齐，保证e8m0_2可以ub内交织计算	241	// 按照2倍blocksize对齐，保证e8m0_2可以ub内交织计算
209	- tilingParam.maxUbCol = static_cast<int64_t>(maxUbAvailable / static_cast<int64_t>(tilingParam.vfLen) / (tilingParam.blockSizeDIGIT_TWO) (tilingParam.blockSize*DIGIT_TWO));	242	+ tilingParam.maxUbCol = static_cast<int64_t>(maxUbAvailable / static_cast<int64_t>(tilingParam.vfLen) /
		243	+ (tilingParam.blockSize * DIGIT_TWO) * (tilingParam.blockSize * DIGIT_TWO));
210	tilingParam.ubFactor = cacheline;	244	tilingParam.ubFactor = cacheline;
211	tilingParam.uo = CeilDiv(tilingParam.postAxisSize, tilingParam.ubFactor);	245	tilingParam.uo = CeilDiv(tilingParam.postAxisSize, tilingParam.ubFactor);
212	tilingParam.tailUbFactor = tilingParam.postAxisSize - (tilingParam.uo - 1) * tilingParam.ubFactor;	246	tilingParam.tailUbFactor = tilingParam.postAxisSize - (tilingParam.uo - 1) * tilingParam.ubFactor;
	@@ -220,79 +254,87 @@ static ge::graphStatus DoTiling(const gert::TilingContext *context, GroupedDynam
220	return ge::GRAPH_SUCCESS;	254	return ge::GRAPH_SUCCESS;
221	}	255	}
222		256
		257	+inline static ge::graphStatus SetTilingKeyParam(gert::TilingContext *context,
		258	+ const GroupedDynamicMxQuantTilingParam &tilingParam, GroupedDynamicMxQuantTilingData &tilingData)
		259	+{
		260	+ uint64_t mode = 0;
		261	+ int64_t tilingKey = GET_TPL_TILING_KEY(mode);
		262	+ OP_LOGD(context->GetNodeName(), "mode is %ld", mode);
		263	+ context->SetTilingKey(tilingKey);
		264	+
		265	+ return ge::GRAPH_SUCCESS;
		266	+}
		267	+
223	inline static ge::graphStatus SetTilingData(gert::TilingContext *context,	268	inline static ge::graphStatus SetTilingData(gert::TilingContext *context,
224	const GroupedDynamicMxQuantTilingParam &tilingParam, GroupedDynamicMxQuantTilingData &tilingData)	269	const GroupedDynamicMxQuantTilingParam &tilingParam, GroupedDynamicMxQuantTilingData &tilingData)
225	{	270	{
226	- OP_LOGD(context->GetNodeName(), "SetTilingData begin.");	271	+ tilingData.totalCoreNum = tilingParam.totalCoreNum;
227	- tilingData.set_totalCoreNum(tilingParam.totalCoreNum);	272	+ tilingData.usedCoreNum = tilingParam.usedCoreNum;
228	- tilingData.set_usedCoreNum(tilingParam.usedCoreNum);	273	+ tilingData.blockFactor = tilingParam.blockFactor;
229	- tilingData.set_blockFactor(tilingParam.blockFactor);	274	+ tilingData.tailBlockFactor = tilingParam.tailBlockFactor;
230	- tilingData.set_tailBlockFactor(tilingParam.tailBlockFactor);	275	+ tilingData.uo = tilingParam.uo;
231	- tilingData.set_uo(tilingParam.uo);	276	+ tilingData.maxUbCol = tilingParam.maxUbCol;
232	- tilingData.set_maxUbCol(tilingParam.maxUbCol);	277	+ tilingData.ubFactor = tilingParam.ubFactor;
233	- tilingData.set_ubFactor(tilingParam.ubFactor);	278	+ tilingData.tailUbFactor = tilingParam.tailUbFactor;
234	- tilingData.set_tailUbFactor(tilingParam.tailUbFactor);	279	+ tilingData.blockSize = tilingParam.blockSize;
235	- tilingData.set_blockSize(tilingParam.blockSize);	280	+ tilingData.scaleAlg = tilingParam.scaleAlg;
236	- tilingData.set_preAxisSize(tilingParam.preAxisSize);	281	+ tilingData.preAxisSize = tilingParam.preAxisSize;
237	- tilingData.set_postAxisSize(tilingParam.postAxisSize);	282	+ tilingData.postAxisSize = tilingParam.postAxisSize;
		283	+ tilingData.dstTypeMax = tilingParam.dstTypeMax;
238		284
239	- OP_CHECK_IF(tilingData.GetDataSize() > context->GetRawTilingData()->GetCapacity(),
240	- OP_LOGE(context->GetNodeName(), "tiling datasize: %zu is bigger than %zu",
241	- tilingData.GetDataSize(), context->GetRawTilingData()->GetCapacity()),
242	- return ge::GRAPH_FAILED);
243	- tilingData.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
244	- context->GetRawTilingData()->SetDataSize(tilingData.GetDataSize());
245	- context->SetBlockDim(tilingData.get_usedCoreNum());
246	- context->SetTilingKey(tilingParam.tilingKey);
247	- size_t *workspaces = context->GetWorkspaceSizes(1);
248	- OP_CHECK_NULL_WITH_CONTEXT(context, workspaces);
249	- workspaces[0] = WORKSPACE_SIZE;
250	return ge::GRAPH_SUCCESS;	285	return ge::GRAPH_SUCCESS;
251	}	286	}
252		287
253	inline static void PrintTilingData(const gert::TilingContext *context, GroupedDynamicMxQuantTilingData &tilingData)	288	inline static void PrintTilingData(const gert::TilingContext *context, GroupedDynamicMxQuantTilingData &tilingData)
254	{	289	{
255	- OP_LOGI(context->GetNodeName(), "tilingData is totalCoreNum:%ld, usedCoreNum:%ld, ubFactor:%ld, \	290	+ OP_LOGI(context, "tilingData is totalCoreNum:%ld, usedCoreNum:%ld, blockFactor:%ld, tailUbFactor:%ld, uo:%ld, \
256	- tailUbFactor:%ld, blockFactor:%ld, tailBlockFactor:%ld, uo:%ld, maxUbCol:%ld, blockSize:%ld, preAxisSize:%ld, postAxisSize:%ld",	291	+ maxUbCol:%ld, ubFactor:%ld, tailBlockFactor:%ld, blockSize:%ld, scaleAlg:%ld, preAxisSize:%ld, postAxisSize:%ld, \
257	- tilingData.get_totalCoreNum(), tilingData.get_usedCoreNum(), tilingData.get_ubFactor(),	292	+ dstTypeMax:%f",
258	- tilingData.get_tailUbFactor(), tilingData.get_blockFactor(), tilingData.get_tailBlockFactor(),	293	+ tilingData.totalCoreNum, tilingData.usedCoreNum, tilingData.blockFactor, tilingData.tailUbFactor, tilingData.uo,
259	- tilingData.get_uo(), tilingData.get_maxUbCol(), tilingData.get_blockSize(),	294	+ tilingData.maxUbCol, tilingData.ubFactor, tilingData.tailBlockFactor, tilingData.blockSize, tilingData.scaleAlg,
260	- tilingData.get_preAxisSize(), tilingData.get_postAxisSize());	295	+ tilingData.preAxisSize, tilingData.postAxisSize, tilingData.dstTypeMax);
261	}	296	}
262		297
263	ge::graphStatus Tiling4GroupedDynamicMxQuant(gert::TilingContext *context)	298	ge::graphStatus Tiling4GroupedDynamicMxQuant(gert::TilingContext *context)
264	{	299	{
265	- OP_LOGD(context->GetNodeName(), "Tiling4GroupedDynamicMxQuant running begin.");	300	+ OP_LOGD(context, "Tiling4GroupedDynamicMxQuant running begin.");
266		301
267	GroupedDynamicMxQuantTilingParam tilingParam;	302	GroupedDynamicMxQuantTilingParam tilingParam;
268		303
269	OP_CHECK_IF(CheckDtype(context, tilingParam) != ge::GRAPH_SUCCESS,	304	OP_CHECK_IF(CheckDtype(context, tilingParam) != ge::GRAPH_SUCCESS,
270	- OP_LOGE(context->GetNodeName(), "The data type check failed."), return ge::GRAPH_FAILED);	305	+ OP_LOGE(context, "The data type check failed."), return ge::GRAPH_FAILED);
271		306
272	OP_CHECK_IF(GetAttr(context, tilingParam) != ge::GRAPH_SUCCESS,	307	OP_CHECK_IF(GetAttr(context, tilingParam) != ge::GRAPH_SUCCESS,
273	- OP_LOGE(context->GetNodeName(), "The attr get failed."), return ge::GRAPH_FAILED);	308	+ OP_LOGE(context, "The attr get failed."), return ge::GRAPH_FAILED);
274		309
275	OP_CHECK_IF(CheckShape(context, tilingParam) != ge::GRAPH_SUCCESS,	310	OP_CHECK_IF(CheckShape(context, tilingParam) != ge::GRAPH_SUCCESS,
276	- OP_LOGE(context->GetNodeName(), "The shape check failed."), return ge::GRAPH_FAILED);	311	+ OP_LOGE(context, "The shape check failed."), return ge::GRAPH_FAILED);
277		312
278	OP_CHECK_IF(GetPlatInfo(context, tilingParam) != ge::GRAPH_SUCCESS,	313	OP_CHECK_IF(GetPlatInfo(context, tilingParam) != ge::GRAPH_SUCCESS,
279	- OP_LOGE(context->GetNodeName(), "GetPlatInfo failed."), return ge::GRAPH_FAILED);	314	+ OP_LOGE(context, "GetPlatInfo failed."), return ge::GRAPH_FAILED);
280		315
		316	+ GroupedDynamicMxQuantTilingData tilingData = {};
281	OP_CHECK_IF(DoTiling(context, tilingParam) != ge::GRAPH_SUCCESS,	317	OP_CHECK_IF(DoTiling(context, tilingParam) != ge::GRAPH_SUCCESS,
282	- OP_LOGE(context->GetNodeName(), "DoTiling failed."), return ge::GRAPH_FAILED);	318	+ OP_LOGE(context, "DoTiling failed."), return ge::GRAPH_FAILED);
283	-	319	+
284	- GroupedDynamicMxQuantTilingData tilingData;	320	+ SetTilingKeyParam(context, tilingParam, tilingData);
285	- OP_CHECK_IF(SetTilingData(context, tilingParam, tilingData) != ge::GRAPH_SUCCESS,	321	+ SetTilingData(context, tilingParam, tilingData);
286	- OP_LOGE(context->GetNodeName(), "SetContext fail."),	322	+ OP_CHECK_IF(GroupedDynamicMxQuantSetTilingData(context, tilingData) != ge::GRAPH_SUCCESS,
		323	+ OP_LOGE(context, "GroupedDynamicMxQuantSetTilingData set tiling data fail."),
287	return ge::GRAPH_FAILED);	324	return ge::GRAPH_FAILED);
288	-	325	+
		326	+ context->SetBlockDim(tilingData.usedCoreNum);
		327	+ size_t *workspaces = context->GetWorkspaceSizes(1);
		328	+ OP_CHECK_NULL_WITH_CONTEXT(context, workspaces);
		329	+ workspaces[0] = WORKSPACE_SIZE;
289	PrintTilingData(context, tilingData);	330	PrintTilingData(context, tilingData);
		331	+
290	return ge::GRAPH_SUCCESS;	332	return ge::GRAPH_SUCCESS;
291	}	333	}
292		334
293	ge::graphStatus TilingPrepare4GroupedDynamicMxQuant(gert::TilingParseContext *context)	335	ge::graphStatus TilingPrepare4GroupedDynamicMxQuant(gert::TilingParseContext *context)
294	{	336	{
295	- OP_LOGD(context->GetNodeName(), "TilingPrepare4GroupedDynamicMxQuant entering.");	337	+ OP_LOGD(context, "TilingPrepare4GroupedDynamicMxQuant entering.");
296	return ge::GRAPH_SUCCESS;	338	return ge::GRAPH_SUCCESS;
297	}	339	}
298		340

Mquant/grouped_dynamic_mx_quant/op_host/arch35/grouped_dynamic_mx_quant_tiling_arch35.h+27-36

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -16,50 +16,41 @@
16	#ifndef AIR_CXX_RUNTIME_V2_OP_IMPL_GROUPED_DYNAMIC_MX_QUANT_H	16	#ifndef AIR_CXX_RUNTIME_V2_OP_IMPL_GROUPED_DYNAMIC_MX_QUANT_H
17	#define AIR_CXX_RUNTIME_V2_OP_IMPL_GROUPED_DYNAMIC_MX_QUANT_H	17	#define AIR_CXX_RUNTIME_V2_OP_IMPL_GROUPED_DYNAMIC_MX_QUANT_H
18		18
19	-#include "register/tilingdata_base.h"	19	+#include "op_host/tiling_base.h"
		20	+#include "op_host/tiling_templates_registry.h"
		21	+#include "op_host/tiling_util.h"
20	#include "graph/types.h"	22	#include "graph/types.h"
		23	+#include "register/op_impl_registry.h"
		24	+#include "quant/grouped_dynamic_mx_quant/op_kernel/arch35/grouped_dynamic_mx_quant_tilingdata.h"
21		25
		26	+using namespace Ops::NN::Optiling;
22	namespace optiling {	27	namespace optiling {
23	-BEGIN_TILING_DATA_DEF(GroupedDynamicMxQuantTilingData)
24	-TILING_DATA_FIELD_DEF(int64_t, totalCoreNum);
25	-TILING_DATA_FIELD_DEF(int64_t, usedCoreNum); // 实际使用的核数
26	-TILING_DATA_FIELD_DEF(int64_t, blockFactor); // 单核循环次数
27	-TILING_DATA_FIELD_DEF(int64_t, tailBlockFactor); // 尾核循环次数
28	-TILING_DATA_FIELD_DEF(int64_t, uo); // 切分轴上的循环次数
29	-TILING_DATA_FIELD_DEF(int64_t, maxUbCol); // 单次循环要处理的数据大小
30	-TILING_DATA_FIELD_DEF(int64_t, ubFactor); // 单次循环要处理的数据大小
31	-TILING_DATA_FIELD_DEF(int64_t, tailUbFactor); // 尾循环要处理的数据大小
32	-TILING_DATA_FIELD_DEF(int64_t, blockSize); // 进行微缩的数据块大小
33	-TILING_DATA_FIELD_DEF(int64_t, preAxisSize); // 输入row长度
34	-TILING_DATA_FIELD_DEF(int64_t, postAxisSize); // 输入column长度
35	-END_TILING_DATA_DEF;
36	-
37	-REGISTER_TILING_DATA_CLASS(GroupedDynamicMxQuant, GroupedDynamicMxQuantTilingData)
38	-
39	struct GroupedDynamicMxQuantCompileInfo {	28	struct GroupedDynamicMxQuantCompileInfo {
40	int64_t coreNum = 0;	29	int64_t coreNum = 0;
41	int64_t ubSize = 0;	30	int64_t ubSize = 0;
42	};	31	};
43		32
44	struct GroupedDynamicMxQuantTilingParam {	33	struct GroupedDynamicMxQuantTilingParam {
45	- int64_t totalCoreNum { 0 };	34	+ int64_t totalCoreNum = 0;
46	- int64_t usedCoreNum { 0 };	35	+ int64_t usedCoreNum = 0;
47	- int64_t blockFactor { 0 };	36	+ int64_t blockFactor = 0;
48	- int64_t tailBlockFactor { 0 };	37	+ int64_t tailBlockFactor = 0;
49	- int64_t uo { 1 };	38	+ int64_t uo = 1;
50	- int64_t maxUbCol { 1 };	39	+ int64_t maxUbCol = 1;
51	- int64_t ubFactor { 0 };	40	+ int64_t ubFactor = 0;
52	- int64_t tailUbFactor { 0 };	41	+ int64_t tailUbFactor = 0;
53	- int64_t blockSize { 0 };	42	+ int64_t blockSize = 0;
54	- int64_t preAxisSize {0};	43	+ int64_t scaleAlg = 0;
55	- int64_t postAxisSize { 1 };	44	+ int64_t preAxisSize = 0;
56	- bool isTailAxis { false };	45	+ int64_t postAxisSize = 1;
57	- int64_t ubSize { 0 };	46	+ float dstTypeMax = 0.0;
58	- uint32_t vfLen { 0 };	47	+ bool isTailAxis = false;
59	- int64_t tilingKey { 0 };	48	+ int64_t ubSize = 0;
60	- int64_t groupSize {1};	49	+ uint32_t vfLen = 0;
61	- ge::DataType inDtype {ge::DT_FLOAT16};	50	+ int64_t tilingKey = 0;
62	- ge::DataType outDtype {ge::DT_FLOAT8_E4M3FN};	51	+ int64_t groupSize = 1;
		52	+ ge::DataType inDtype = ge::DT_FLOAT16;
		53	+ ge::DataType outDtype = ge::DT_FLOAT8_E4M3FN;
63	};	54	};
64		55
65	} // namespace optiling	56	} // namespace optiling

Mquant/grouped_dynamic_mx_quant/op_host/config/ascend950/grouped_dynamic_mx_quant_binary.json+40-0

	@@ -66,6 +66,16 @@
66	"name": "blocksize",	66	"name": "blocksize",
67	"dtype": "int",	67	"dtype": "int",
68	"value": null	68	"value": null
		69	+ },
		70	+ {
		71	+ "name": "scale_alg",
		72	+ "dtype": "int",
		73	+ "value": null
		74	+ },
		75	+ {
		76	+ "name": "dst_type_max",
		77	+ "dtype": "float32",
		78	+ "value": null
69	}	79	}
70	]	80	]
71	},	81	},
	@@ -134,6 +144,16 @@
134	"name": "blocksize",	144	"name": "blocksize",
135	"dtype": "int",	145	"dtype": "int",
136	"value": null	146	"value": null
		147	+ },
		148	+ {
		149	+ "name": "scale_alg",
		150	+ "dtype": "int",
		151	+ "value": null
		152	+ },
		153	+ {
		154	+ "name": "dst_type_max",
		155	+ "dtype": "float32",
		156	+ "value": null
137	}	157	}
138	]	158	]
139	},	159	},
	@@ -202,6 +222,16 @@
202	"name": "blocksize",	222	"name": "blocksize",
203	"dtype": "int",	223	"dtype": "int",
204	"value": null	224	"value": null
		225	+ },
		226	+ {
		227	+ "name": "scale_alg",
		228	+ "dtype": "int",
		229	+ "value": null
		230	+ },
		231	+ {
		232	+ "name": "dst_type_max",
		233	+ "dtype": "float32",
		234	+ "value": null
205	}	235	}
206	]	236	]
207	},	237	},
	@@ -270,6 +300,16 @@
270	"name": "blocksize",	300	"name": "blocksize",
271	"dtype": "int",	301	"dtype": "int",
272	"value": null	302	"value": null
		303	+ },
		304	+ {
		305	+ "name": "scale_alg",
		306	+ "dtype": "int",
		307	+ "value": null
		308	+ },
		309	+ {
		310	+ "name": "dst_type_max",
		311	+ "dtype": "float32",
		312	+ "value": null
273	}	313	}
274	]	314	]
275	}	315	}

Mquant/grouped_dynamic_mx_quant/op_host/grouped_dynamic_mx_quant_def.cpp+39-17

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -17,36 +17,57 @@
17	#include "register/op_def_registry.h"	17	#include "register/op_def_registry.h"
18		18
19	namespace ops {	19	namespace ops {
20	-constexpr int32_t BLOCK_SIZE = 32;	20	+static constexpr int32_t DEFAULT_DST_TYPE = 35;
		21	+static constexpr int32_t DEFAULT_BLOCK_SIZE = 32;
		22	+static constexpr int32_t DEFAULT_SCALE_ALG = 0;
		23	+static constexpr float DEFAULT_DST_TYPE_MAX_VALUE = 0.0;
		24	+
		25	+static const std::vector<ge::DataType> groupedDynamicMxQuantXDataType = {
		26	+ ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16};
		27	+
		28	+static const std::vector<ge::DataType> groupedDynamicMxQuantGroupIndexDataType = {
		29	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32};
		30	+
		31	+static const std::vector<ge::DataType> groupedDynamicMxQuantYDataType = {
		32	+ ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2,ge::DT_FLOAT8_E5M2};
		33	+
		34	+static const std::vector<ge::DataType> groupedDynamicMxQuantMxScaleDataType = {
		35	+ ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0};
		36	+
		37	+static const std::vector<ge::Format> groupedDynamicMxQuantNDFormat = {
		38	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
		39	+
21	class GroupedDynamicMxQuant : public OpDef {	40	class GroupedDynamicMxQuant : public OpDef {
22	public:	41	public:
23	explicit GroupedDynamicMxQuant(const char* name) : OpDef(name)	42	explicit GroupedDynamicMxQuant(const char* name) : OpDef(name)
24	{	43	{
25	this->Input("x")	44	this->Input("x")
26	.ParamType(REQUIRED)	45	.ParamType(REQUIRED)
27	- .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16})	46	+ .DataType(groupedDynamicMxQuantXDataType)
28	- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	47	+ .Format(groupedDynamicMxQuantNDFormat)
29	- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	48	+ .UnknownShapeFormat(groupedDynamicMxQuantNDFormat)
30	.AutoContiguous();	49	.AutoContiguous();
31	this->Input("group_index")	50	this->Input("group_index")
32	.ParamType(REQUIRED)	51	.ParamType(REQUIRED)
33	- .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})	52	+ .DataType(groupedDynamicMxQuantGroupIndexDataType)
34	- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	53	+ .Format(groupedDynamicMxQuantNDFormat)
35	- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	54	+ .UnknownShapeFormat(groupedDynamicMxQuantNDFormat)
36	.AutoContiguous();	55	.AutoContiguous();
37	this->Output("y")	56	this->Output("y")
38	.ParamType(REQUIRED)	57	.ParamType(REQUIRED)
39	- .DataType({ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2,ge::DT_FLOAT8_E5M2})	58	+ .DataType(groupedDynamicMxQuantYDataType)
40	- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	59	+ .Format(groupedDynamicMxQuantNDFormat)
41	- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});	60	+ .UnknownShapeFormat(groupedDynamicMxQuantNDFormat);
42	this->Output("mxscale")	61	this->Output("mxscale")
43	.ParamType(REQUIRED)	62	.ParamType(REQUIRED)
44	- .DataType({ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0})	63	+ .DataType(groupedDynamicMxQuantMxScaleDataType)
45	- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	64	+ .Format(groupedDynamicMxQuantNDFormat)
46	- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});	65	+ .UnknownShapeFormat(groupedDynamicMxQuantNDFormat);
47	this->Attr("round_mode").AttrType(OPTIONAL).String("rint");	66	this->Attr("round_mode").AttrType(OPTIONAL).String("rint");
48	- this->Attr("dst_type").AttrType(OPTIONAL).Int(ge::DT_FLOAT8_E5M2);	67	+ this->Attr("dst_type").AttrType(OPTIONAL).Int(DEFAULT_DST_TYPE);
49	- this->Attr("blocksize").AttrType(OPTIONAL).Int(BLOCK_SIZE);	68	+ this->Attr("blocksize").AttrType(OPTIONAL).Int(DEFAULT_BLOCK_SIZE);
		69	+ this->Attr("scale_alg").AttrType(OPTIONAL).Int(DEFAULT_SCALE_ALG);
		70	+ this->Attr("dst_type_max").AttrType(OPTIONAL).Float(DEFAULT_DST_TYPE_MAX_VALUE);
50		71
51	OpAICoreConfig aicoreConfig;	72	OpAICoreConfig aicoreConfig;
52	aicoreConfig.DynamicCompileStaticFlag(true)	73	aicoreConfig.DynamicCompileStaticFlag(true)
	@@ -54,7 +75,8 @@ public:
54	.DynamicRankSupportFlag(true)	75	.DynamicRankSupportFlag(true)
55	.DynamicShapeSupportFlag(true)	76	.DynamicShapeSupportFlag(true)
56	.NeedCheckSupportFlag(false)	77	.NeedCheckSupportFlag(false)
57	- .PrecisionReduceFlag(true);	78	+ .PrecisionReduceFlag(true)
		79	+ .ExtendCfgInfo("opFile.value", "grouped_dynamic_mx_quant");
58	this->AICore().AddConfig("ascend950", aicoreConfig);	80	this->AICore().AddConfig("ascend950", aicoreConfig);
59	}	81	}
60	};	82	};

Mquant/grouped_dynamic_mx_quant/op_host/grouped_dynamic_mx_quant_infershape.cpp+69-68

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -9,13 +9,14 @@
9	*/	9	*/
10		10
11	/*!	11	/*!
12	- * \file grouped_dynamic_mx_quant.cc	12	+ * \file grouped_dynamic_mx_quant.cpp
13	* \brief	13	* \brief
14	*/	14	*/
15		15
16	#include "log/log.h"	16	#include "log/log.h"
17	#include "register/op_impl_registry.h"	17	#include "register/op_impl_registry.h"
18	-#include "op_host/util/shape_util.h"	18	+#include "util/shape_util.h"
		19	+#include "util/math_util.h"
19		20
20	using namespace ge;	21	using namespace ge;
21	namespace ops {	22	namespace ops {
	@@ -40,83 +41,83 @@ static std::string Shape2String(const T& shape) {
40	}	41	}
41		42
42	graphStatus InferShapeForGroupedDynamicMxQuant(gert::InferShapeContext* context) {	43	graphStatus InferShapeForGroupedDynamicMxQuant(gert::InferShapeContext* context) {
43	- OP_LOGD(context->GetNodeName(), "Begin to do InferShapeForGroupedDynamicMxQuant");	44	+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeForGroupedDynamicMxQuant");
44	- const gert::Shape* xShape = context->GetInputShape(0);	45	+ const gert::Shape* xShape = context->GetInputShape(0);
45	- OP_CHECK_NULL_WITH_CONTEXT(context, xShape);	46	+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
46	- const gert::Shape* groupIdxShape = context->GetInputShape(1);	47	+ const gert::Shape* groupIdxShape = context->GetInputShape(1);
47	- OP_CHECK_NULL_WITH_CONTEXT(context, groupIdxShape);	48	+ OP_CHECK_NULL_WITH_CONTEXT(context, groupIdxShape);
48		49
49	- gert::Shape* yShape = context->GetOutputShape(0);	50	+ gert::Shape* yShape = context->GetOutputShape(0);
50	- OP_CHECK_NULL_WITH_CONTEXT(context, yShape);	51	+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
51	- yShape = xShape;	52	+ yShape = xShape;
52		53
53	- gert::Shape* scaleShape = context->GetOutputShape(1);	54	+ gert::Shape* scaleShape = context->GetOutputShape(1);
54	- OP_CHECK_NULL_WITH_CONTEXT(context, scaleShape);	55	+ OP_CHECK_NULL_WITH_CONTEXT(context, scaleShape);
55		56
56	- auto attrsPtr = context->GetAttrs();	57	+ auto attrsPtr = context->GetAttrs();
57	- OP_CHECK_NULL_WITH_CONTEXT(context, attrsPtr);	58	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrsPtr);
58	-
59	- const int32_t *blockSize = attrsPtr->GetAttrPointer<int32_t>(INDEX_ATTR_BLOCK_SIZE);
60	- OP_CHECK_NULL_WITH_CONTEXT(context, blockSize);
61	- OP_CHECK_IF(static_cast<int64_t>(*blockSize) != 32,
62	- OP_LOGE(context->GetNodeName(), "blockSize is invalid, must be 32"),
63	- return ge::GRAPH_FAILED);
64	- size_t xShapeSize = xShape->GetDimNum();
65	- size_t groupIdxShapeSize = groupIdxShape->GetDimNum();
66	- OP_CHECK_IF(groupIdxShapeSize != 1,
67	- OP_LOGE(context->GetNodeName(),
68	- "group_index's shape must be 1D, but is %lu", groupIdxShapeSize),
69	- return ge::GRAPH_FAILED);
70	- int64_t groupIdxDim0 = groupIdxShape->GetDim(0);
71	- OP_CHECK_IF(groupIdxDim0 == 0,
72	- OP_LOGE(context->GetNodeName(), "group_index does not support empty tensor"),
73	- return ge::GRAPH_FAILED);
74		59
75	- // dynamic -2 (input x)	60	+ const int32_t *blockSize = attrsPtr->GetAttrPointer<int32_t>(INDEX_ATTR_BLOCK_SIZE);
76	- if (Ops::Base::IsUnknownRank(*xShape)) {	61	+ OP_CHECK_NULL_WITH_CONTEXT(context, blockSize);
77	- OP_LOGD(context->GetNodeName(), "input x is UnknownRank, set outputs' shape to -2");	62	+ OP_CHECK_IF(static_cast<int64_t>(*blockSize) != 32,
78	- scaleShape = xShape;	63	+ OP_LOGE(context->GetNodeName(), "blockSize is invalid, must be 32"),
79	- return ge::GRAPH_SUCCESS;	64	+ return ge::GRAPH_FAILED);
80	- } else {	65	+
81	- OP_CHECK_IF(xShapeSize != 2,	66	+ size_t xShapeSize = xShape->GetDimNum();
		67	+ size_t groupIdxShapeSize = groupIdxShape->GetDimNum();
		68	+ OP_CHECK_IF(groupIdxShapeSize != 1,
		69	+ OP_LOGE(context->GetNodeName(),
		70	+ "group_index's shape must be 1D, but is %lu", groupIdxShapeSize),
		71	+ return ge::GRAPH_FAILED);
		72	+ int64_t groupIdxDim0 = groupIdxShape->GetDim(0);
		73	+ OP_CHECK_IF(groupIdxDim0 == 0,
		74	+ OP_LOGE(context->GetNodeName(), "group_index does not support empty tensor"),
		75	+ return ge::GRAPH_FAILED);
		76	+
		77	+ // dynamic -2 (input x)
		78	+ if (Ops::Base::IsUnknownRank(*xShape)) {
		79	+ OP_LOGD(context->GetNodeName(), "input x is UnknownRank, set outputs' shape to -2");
		80	+ scaleShape = xShape;
		81	+ return ge::GRAPH_SUCCESS;
		82	+ } else {
		83	+ OP_CHECK_IF(xShapeSize != 2,
82	OP_LOGE(context->GetNodeName(),	84	OP_LOGE(context->GetNodeName(),
83	"input x is not UnknownRank, shape must be 2D, but is %lu", xShapeSize),	85	"input x is not UnknownRank, shape must be 2D, but is %lu", xShapeSize),
84	return ge::GRAPH_FAILED);	86	return ge::GRAPH_FAILED);
85	- }	87	+ }
86		88
87	- int64_t dim0Size = (xShape->GetDim(0) / (static_cast<int64_t>(blockSize) 2) + groupIdxDim0); // 不带起始0	89	+ int64_t dim0Size = (xShape->GetDim(0) / (static_cast<int64_t>(blockSize) 2) + groupIdxDim0); // 不带起始0
88	- // dynamic -2 or -1 (groupIdxShape), dynamic -1 (input x)	90	+ // dynamic -2 or -1 (groupIdxShape), dynamic -1 (input x)
89	- if (Ops::Base::IsUnknownRank(*groupIdxShape) \|\| groupIdxDim0 == -1\|\| xShape->GetDim(0) == -1) {	91	+ if (Ops::Base::IsUnknownRank(*groupIdxShape) \|\| groupIdxDim0 == -1 \|\| xShape->GetDim(0) == -1) {
90	- dim0Size = -1;	92	+ dim0Size = -1;
91	- }	93	+ }
92		94
93	- scaleShape->SetDimNum(SCALE_DIM_NUM);	95	+ scaleShape->SetDimNum(SCALE_DIM_NUM);
94	- scaleShape->SetDim(0, dim0Size);	96	+ scaleShape->SetDim(0, dim0Size);
95	- scaleShape->SetDim(1, (xShape->GetDim(1) == -1) ? -1 : xShape->GetDim(1));	97	+ scaleShape->SetDim(1, (xShape->GetDim(1) == -1) ? -1 : xShape->GetDim(1));
96	- scaleShape->SetDim(2, 2); // mxscale's third dimensions only support 2	98	+ scaleShape->SetDim(2, 2); // mxscale's third dimensions only support 2
97	- OP_LOGD(context->GetNodeName(), "mxscale shape is :%s after infershape.", Shape2String(*scaleShape).c_str());	99	+ OP_LOGD(context->GetNodeName(), "mxscale shape is :%s after infershape.", Shape2String(*scaleShape).c_str());
98	- OP_LOGD(context->GetNodeName(), "End to do InferShapeForGroupedDynamicMxQuant");	100	+ OP_LOGD(context->GetNodeName(), "End to do InferShapeForGroupedDynamicMxQuant");
99	- return ge::GRAPH_SUCCESS;	101	+ return ge::GRAPH_SUCCESS;
100	}	102	}
101		103
102	static ge::graphStatus InferDataTypeForGroupedDynamicMxQuant(gert::InferDataTypeContext *context) {	104	static ge::graphStatus InferDataTypeForGroupedDynamicMxQuant(gert::InferDataTypeContext *context) {
103	- OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeForGroupedDynamicMxQuant");	105	+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeForGroupedDynamicMxQuant");
104	- auto attrsPtr = context->GetAttrs();	106	+ auto attrsPtr = context->GetAttrs();
105	- OP_CHECK_NULL_WITH_CONTEXT(context, attrsPtr);	107	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrsPtr);
106	- ge::DataType yDtype = ge::DT_FLOAT8_E5M2;	108	+ ge::DataType yDtype = ge::DT_FLOAT8_E5M2;
107	- const int32_t *pDstDtype = attrsPtr->GetAttrPointer<int32_t>(INDEX_ATTR_DST_TYPE);	109	+ const int32_t *pDstDtype = attrsPtr->GetAttrPointer<int32_t>(INDEX_ATTR_DST_TYPE);
108	- if (pDstDtype != nullptr) {	110	+ if (pDstDtype != nullptr) {
109	- int32_t dstDtype = *pDstDtype;	111	+ int32_t dstDtype = *pDstDtype;
110	- OP_CHECK_IF(dstDtype != DTYPE_FLOAT8_E5M2 && dstDtype != DTYPE_FLOAT8_E4M3FN,	112	+ OP_CHECK_IF(dstDtype != DTYPE_FLOAT8_E5M2 && dstDtype != DTYPE_FLOAT8_E4M3FN,
111	- OP_LOGE(context->GetNodeName(),	113	+ OP_LOGE(context->GetNodeName(), "attr dst_type only support 35(FLOAT8_E5M2) and 36(FLOAT8_E4M3FN)"),
112	- "attr dst_type only support 35(FLOAT8_E5M2) and 36(FLOAT8_E4M3FN)"),	114	+ return ge::GRAPH_FAILED);
113	- return ge::GRAPH_FAILED);	115	+ yDtype = static_cast<ge::DataType>(dstDtype);
114	- yDtype = static_cast<ge::DataType>(dstDtype);	116	+ }
115	- }	117	+ context->SetOutputDataType(0, yDtype);
116	- context->SetOutputDataType(0, yDtype);	118	+ context->SetOutputDataType(1, ge::DT_FLOAT8_E8M0);
117	- context->SetOutputDataType(1, ge::DT_FLOAT8_E8M0);	119	+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeForGroupedDynamicMxQuant");
118	- OP_LOGD(context->GetNodeName(), "End to do InferDataTypeForGroupedDynamicMxQuant");	120	+ return ge::GRAPH_SUCCESS;
119	- return ge::GRAPH_SUCCESS;
120	}	121	}
121		122
122	IMPL_OP_INFERSHAPE(GroupedDynamicMxQuant)	123	IMPL_OP_INFERSHAPE(GroupedDynamicMxQuant)

Aquant/grouped_dynamic_mx_quant/op_host/op_api/aclnn_grouped_dynamic_mx_quant_v2.h+58-0

@@ -0,0 +1,58 @@
	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	+#ifndef ACLNN_GROUPED_DYNAMIC_MX_QUANT_V2_H_
	11	+#define ACLNN_GROUPED_DYNAMIC_MX_QUANT_V2_H_
	12	+
	13	+#include "aclnn/acl_meta.h"
	14	+
	15	+#ifdef __cplusplus
	16	+extern "C" {
	17	+#endif
	18	+
	19	+/**
	20	+ * @brief aclnnGroupedDynamicMxQuantV2的第一段接口，根据具体的计算流程，计算workspace大小。
	21	+ * @domain aclnn_ops_infer
	22	+ *
	23	+ * @param [in] x: 待进行GroupedDynamicMxQuant计算的入参。npu device侧的aclTensor，
	24	+ * 数据类型支持float16, bfloat16, 数据格式支持ND，支持非连续的Tensor。
	25	+ * @param [in] groupIndex: npu device侧的aclTensor，数据类型支持int32
	26	+ * @param [in] roundMode: host侧的aclScalar，数据类型string，仅支持 "rint"
	27	+ * @param [in] dstType: host侧的aclScalar, 数据类型int, 输入范围为{35, 36}，分别对应输出y的数据类型为{35: FLOAT8_E5M2, 36: FLOAT8_E4M3FN}
	28	+ * @param [in] blocksize: host侧的aclScalar, 数据类型int，仅支持 "32"
	29	+ * @param [in] scaleAlg: host侧的aclScalar, 数据类型int，仅支持 "0"和 "1"
	30	+ * @param [in] dstTypeMax: host侧的aclScalar, 数据类型double，仅支持 "0.0"和 "6.0-12.0"
	31	+ * @param [in] y: GroupedDynamicMxQuant计算的出参。npu device侧的aclTensor，
	32	+ * 数据类型支持float8_e4m3fn, float8_e5m2, 数据格式支持ND，支持非连续的Tensor。
	33	+ * @param [in] mxscale: GroupedDynamicMxQuant计算的出参。npu device侧的aclTensor，
	34	+ * 数据类型支持float8_e8m0, 数据格式支持ND，不支持非连续的Tensor。
	35	+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
	36	+ * @param [out] executor: 返回op执行器，包含算子计算流程。
	37	+ * @return aclnnStatus: 返回状态码。
	38	+ */
	39	+__attribute__((visibility("default"))) aclnnStatus aclnnGroupedDynamicMxQuantV2GetWorkspaceSize(
	40	+ const aclTensor* x, const aclTensor* groupIndex, const char* roundMode, int64_t dstType, int64_t blocksize, int64_t scaleAlg,
	41	+ double dstTypeMax, const aclTensor* y, const aclTensor* mxscale, uint64_t* workspaceSize, aclOpExecutor** executor);
	42	+
	43	+/**
	44	+ * @brief aclnnGroupedDynamicMxQuantV2的第二段接口，用于执行计算。
	45	+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
	46	+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnGroupQuantV2GetWorkspaceSize获取。
	47	+ * @param [in] executor: op执行器，包含了算子计算流程。
	48	+ * @param [in] stream: acl stream流。
	49	+ * @return aclnnStatus: 返回状态码。
	50	+ */
	51	+__attribute__((visibility("default"))) aclnnStatus aclnnGroupedDynamicMxQuantV2(void* workspace, uint64_t workspaceSize,
	52	+ aclOpExecutor* executor, aclrtStream stream);
	53	+
	54	+#ifdef __cplusplus
	55	+}
	56	+#endif
	57	+
	58	+#endif // ACLNN_GROUPED_DYNAMIC_MX_QUANT_V2_H_

Mquant/grouped_dynamic_mx_quant/op_host/op_api/grouped_dynamic_mx_quant.h+3-3

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -15,8 +15,8 @@
15		15
16	namespace l0op {	16	namespace l0op {
17	std::tuple<aclTensor, aclTensor> GroupedDynamicMxQuant(const aclTensor x, const aclTensor groupIndex,	17	std::tuple<aclTensor, aclTensor> GroupedDynamicMxQuant(const aclTensor x, const aclTensor groupIndex,
18	- const char *roundMode, int64_t dstType,	18	+ const char *roundMode, int64_t dstType, int64_t blocksize,
19	- int64_t blocksize, aclOpExecutor *executor);	19	+ int64_t scaleAlg, float dstTypeMax, aclOpExecutor *executor);
20	} // l0op	20	} // l0op
21		21
22	#endif // OP_API_INC_LEVEL0_GROUP_QUANT_H_	22	#endif // OP_API_INC_LEVEL0_GROUP_QUANT_H_

Mquant/grouped_dynamic_mx_quant/op_kernel/arch35/grouped_dynamic_mx_quant_common.h+17-1

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -17,6 +17,9 @@
17	#define GRROUPED_DYNAMIC_MX_QUANT_COMMON_H	17	#define GRROUPED_DYNAMIC_MX_QUANT_COMMON_H
18		18
19	#include "kernel_operator.h"	19	#include "kernel_operator.h"
		20	+#include "../inc/platform.h"
		21	+#include "grouped_dynamic_mx_quant_tilingdata.h"
		22	+
20	namespace GroupedDynamicMxQuant {	23	namespace GroupedDynamicMxQuant {
21	template<typename Tp, Tp v>	24	template<typename Tp, Tp v>
22	struct IntegralConstant {	25	struct IntegralConstant {
	@@ -40,5 +43,18 @@ constexpr int16_t SHR_NUM_FOR_BF16 = 7;
40	constexpr uint16_t BF16_EXP_BIAS = 0x7f00;	43	constexpr uint16_t BF16_EXP_BIAS = 0x7f00;
41	constexpr uint16_t FP8_E4M3_MAX_EXP = 0x0400; // elem_emax右移7位(BF16E8M7)	44	constexpr uint16_t FP8_E4M3_MAX_EXP = 0x0400; // elem_emax右移7位(BF16E8M7)
42	constexpr uint16_t FP8_E5M2_MAX_EXP = 0x0780;	45	constexpr uint16_t FP8_E5M2_MAX_EXP = 0x0780;
		46	+
		47	+constexpr uint16_t ABS_FOR_UINT16 = 0x7fff;
		48	+constexpr uint32_t MAN_FOR_FP32 = 0x007fffff;
		49	+constexpr uint32_t MAX_EXP_FOR_FP8_IN_FP32 = 0x000000ff;
		50	+constexpr uint32_t FP32_EXP_BIAS_CUBLAS = 0x00007f00;
		51	+constexpr uint32_t NAN_CUSTOMIZATION_PACK = 0x00007f81;
		52	+constexpr int16_t SHR_NUM_FOR_FP32 = 23;
		53	+constexpr uint32_t MAX_EXP_FOR_FP32 = 0x7f800000;
		54	+constexpr uint32_t NUMBER_ZERO = 0x00000000;
		55	+constexpr uint32_t NUMBER_TWO_FIVE_FOUR = 0x000000fe;
		56	+constexpr uint32_t NUMBER_HALF = 0x00400000;
		57	+constexpr float FP8_E4M3_INV_MAX = 0.002232142857; // 1/448
		58	+constexpr float FP8_E5M2_INV_MAX = 0.000017438616; // 1/57344
43	}	59	}
44	#endif // GRROUPED_DYNAMIC_MX_QUANT_COMMON_H	60	#endif // GRROUPED_DYNAMIC_MX_QUANT_COMMON_H

Mquant/grouped_dynamic_mx_quant/op_kernel/arch35/grouped_dynamic_mx_quant_not_tail_axis_fp8.h+510-198

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -16,12 +16,16 @@
16	#ifndef GROUPED_DYNAMIC_MX_QUANT_NOT_TAIL_AXIS_FP8_H	16	#ifndef GROUPED_DYNAMIC_MX_QUANT_NOT_TAIL_AXIS_FP8_H
17	#define GROUPED_DYNAMIC_MX_QUANT_NOT_TAIL_AXIS_FP8_H	17	#define GROUPED_DYNAMIC_MX_QUANT_NOT_TAIL_AXIS_FP8_H
18	#include "grouped_dynamic_mx_quant_common.h"	18	#include "grouped_dynamic_mx_quant_common.h"
		19	+#include "grouped_dynamic_mx_quant_tilingdata.h"
		20	+#include "grouped_dynamic_mx_quant_struct.h"
		21	+#include "kernel_operator.h"
19	#include "op_kernel/platform_util.h"	22	#include "op_kernel/platform_util.h"
20	#include "op_kernel/math_util.h"	23	#include "op_kernel/math_util.h"
21		24
22	namespace GroupedDynamicMxQuant {	25	namespace GroupedDynamicMxQuant {
23	using namespace AscendC;	26	using namespace AscendC;
24	#define FLOAT_OVERFLOW_MODE_CTRL 60	27	#define FLOAT_OVERFLOW_MODE_CTRL 60
		28	+
25	template <typename T, typename U>	29	template <typename T, typename U>
26	class GroupedDynamicMxQuantBaseFP8 {	30	class GroupedDynamicMxQuantBaseFP8 {
27	public:	31	public:
	@@ -34,8 +38,11 @@ private:
34	__aicore__ inline void SplitPostAxisCompute(int64_t dataLen, int64_t blockCount);	38	__aicore__ inline void SplitPostAxisCompute(int64_t dataLen, int64_t blockCount);
35	__aicore__ inline void CopyIn(int64_t offset, int64_t blockCount, int64_t dataLen);	39	__aicore__ inline void CopyIn(int64_t offset, int64_t blockCount, int64_t dataLen);
36	__aicore__ inline void CopyOut(int64_t xOffset, int64_t scaleOffset, int64_t blockCount, int64_t dataLen);	40	__aicore__ inline void CopyOut(int64_t xOffset, int64_t scaleOffset, int64_t blockCount, int64_t dataLen);
37	- __aicore__ inline void Compute(int64_t dataLen, int64_t blockCount,	41	+ __aicore__ inline void ComputeOCP(int64_t dataLen, int64_t blockCount,
38	__ubuf__ T xAddr, __ubuf__ uint8_t mxScaleAddr, __ubuf__ uint8_t *yAddr);	42	__ubuf__ T xAddr, __ubuf__ uint8_t mxScaleAddr, __ubuf__ uint8_t *yAddr);
		43	+ __aicore__ inline void ComputecuBLAS(int64_t dataLen, int64_t blockCount,
		44	+ __ubuf__ T xAddr, __ubuf__ uint8_t mxScaleAddr, __ubuf__ uint8_t *yAddr);
		45	+
39	private:	46	private:
40	TPipe pipe_;	47	TPipe pipe_;
41	TQue<QuePosition::VECIN, DB_BUFFER> inQueue_;	48	TQue<QuePosition::VECIN, DB_BUFFER> inQueue_;
	@@ -58,7 +65,8 @@ private:
58	int64_t tailUbFactor_ = 0; // n轴cacheline尾块	65	int64_t tailUbFactor_ = 0; // n轴cacheline尾块
59	int64_t preAxisSize_ = 0; // m轴大小	66	int64_t preAxisSize_ = 0; // m轴大小
60	int64_t postAxisSize_ = 0; // n轴大小	67	int64_t postAxisSize_ = 0; // n轴大小
61	- int64_t blockSize_ = 0; // 使用核数，仅支持32	68	+ int64_t blockSize_ = 0; // 量化数据块大小，仅支持32
		69	+ int64_t scaleAlg_ = 0;
62	};	70	};
63		71
64	template <typename T, typename U>	72	template <typename T, typename U>
	@@ -72,6 +80,7 @@ __aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::ParseTilingData(const
72	blockFactor_ = tilingData.blockFactor;	80	blockFactor_ = tilingData.blockFactor;
73	tailBlockFactor_ = tilingData.tailBlockFactor;	81	tailBlockFactor_ = tilingData.tailBlockFactor;
74	blockSize_ = tilingData.blockSize;	82	blockSize_ = tilingData.blockSize;
		83	+ scaleAlg_ = tilingData.scaleAlg;
75	preAxisSize_ = tilingData.preAxisSize;	84	preAxisSize_ = tilingData.preAxisSize;
76	postAxisSize_ = tilingData.postAxisSize;	85	postAxisSize_ = tilingData.postAxisSize;
77	}	86	}
	@@ -83,7 +92,7 @@ __aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::CopyIn(int64_t offset
83	static_cast<uint16_t>(blockCount),	92	static_cast<uint16_t>(blockCount),
84	static_cast<uint32_t>(dataLen * sizeof(T)),	93	static_cast<uint32_t>(dataLen * sizeof(T)),
85	static_cast<uint32_t>((postAxisSize_ - dataLen) * sizeof(T)),	94	static_cast<uint32_t>((postAxisSize_ - dataLen) * sizeof(T)),
86	- static_cast<uint32_t>((dataLen +31)/32*2-(dataLen+15)/16),	95	+ static_cast<uint32_t>((dataLen + 31) / 32 * 2 - (dataLen + 15) / 16),
87	static_cast<uint32_t>(0)	96	static_cast<uint32_t>(0)
88	};	97	};
89		98
	@@ -124,9 +133,10 @@ template <typename T, typename U>
124	__aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::Init(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR mxScale,	133	__aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::Init(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR mxScale,
125	const GroupedDynamicMxQuantTilingData &tilingData)	134	const GroupedDynamicMxQuantTilingData &tilingData)
126	{	135	{
127	- #if (__NPU_ARCH__ == 3510)	136	+#if (__NPU_ARCH__ == 3510)
128	- AscendC::SetCtrlSpr<FLOAT_OVERFLOW_MODE_CTRL,FLOAT_OVERFLOW_MODE_CTRL>(0);	137	+ SetCtrlSpr<FLOAT_OVERFLOW_MODE_CTRL, FLOAT_OVERFLOW_MODE_CTRL>(0);
129	- #endif	138	+#endif
		139	+
130	blockIdx_ = GetBlockIdx();	140	blockIdx_ = GetBlockIdx();
131	ParseTilingData(tilingData);	141	ParseTilingData(tilingData);
132	if (this->blockIdx_ >= this->usedCoreNum_) {	142	if (this->blockIdx_ >= this->usedCoreNum_) {
	@@ -167,20 +177,22 @@ __aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::Process()
167	gIdxValueStart = groupIndexGm_.GetValue(gIdx -1);	177	gIdxValueStart = groupIndexGm_.GetValue(gIdx -1);
168	}	178	}
169	int64_t gIdxValueEnd = groupIndexGm_.GetValue(gIdx);	179	int64_t gIdxValueEnd = groupIndexGm_.GetValue(gIdx);
170	- assert((0 <= gIdxValueEnd && gIdxValueEnd <= preAxisSize_), "groupIndex %lld = %lld out of range[0 %lld]!\n", gIdx, gIdxValueEnd, preAxisSize_);	180	+ assert((0 <= gIdxValueEnd && gIdxValueEnd <= preAxisSize_), "groupIndex %lld = %lld out of range[0 %lld]!\n",
171	- assert((gIdxValueStart <= gIdxValueEnd), "groupIndex %lld = %lld less than previous groupIdx = %lld!\n", gIdx, gIdxValueEnd, gIdxValueStart);	181	+ gIdx, gIdxValueEnd, preAxisSize_);
		182	+ assert((gIdxValueStart <= gIdxValueEnd), "groupIndex %lld = %lld less than previous groupIdx = %lld!\n",
		183	+ gIdx, gIdxValueEnd, gIdxValueStart);
172	int64_t groupSizeIdx = gIdxValueEnd - gIdxValueStart;	184	int64_t groupSizeIdx = gIdxValueEnd - gIdxValueStart;
173	int64_t rowOffset = gIdxValueStart * postAxisSize_ + nIdx * ubFactor_;	185	int64_t rowOffset = gIdxValueStart * postAxisSize_ + nIdx * ubFactor_;
174	- int64_t rowGroupOffset = 2(gIdxValueStart /64 + gIdx) postAxisSize_ + 2 * nIdx * ubFactor_; // e8m0_2	186	+ int64_t rowGroupOffset = 2 * (gIdxValueStart / 64 + gIdx) * postAxisSize_ + 2 * nIdx * ubFactor_; // e8m0_2
175	bool isTailLoopInUbDim = nIdx == uo_ - 1;	187	bool isTailLoopInUbDim = nIdx == uo_ - 1;
176	int64_t dataLen = isTailLoopInUbDim ? tailUbFactor_ : ubFactor_;	188	int64_t dataLen = isTailLoopInUbDim ? tailUbFactor_ : ubFactor_;
177	int64_t inLoopNum = Ops::Base::CeilDiv(groupSizeIdx, maxUbCol_);	189	int64_t inLoopNum = Ops::Base::CeilDiv(groupSizeIdx, maxUbCol_);
178	for (int64_t j = 0; j < inLoopNum; j++){	190	for (int64_t j = 0; j < inLoopNum; j++){
179	int64_t blockCount = (j == inLoopNum - 1)? groupSizeIdx - j * maxUbCol_: maxUbCol_;	191	int64_t blockCount = (j == inLoopNum - 1)? groupSizeIdx - j * maxUbCol_: maxUbCol_;
180	int64_t xGmOffset = rowOffset + j * maxUbCol_ * postAxisSize_;	192	int64_t xGmOffset = rowOffset + j * maxUbCol_ * postAxisSize_;
181	- int64_t scaleGmOffset = rowGroupOffset + j * (maxUbCol_/32) * postAxisSize_;	193	+ int64_t scaleGmOffset = rowGroupOffset + j * (maxUbCol_ / 32) * postAxisSize_;
182	CopyIn(xGmOffset, blockCount, dataLen);	194	CopyIn(xGmOffset, blockCount, dataLen);
183	- SplitPostAxisCompute(blockCount, (dataLen+31)/32*32);	195	+ SplitPostAxisCompute(blockCount, (dataLen + 31) / 32 * 32);
184	CopyOut(xGmOffset, scaleGmOffset, blockCount, dataLen);	196	CopyOut(xGmOffset, scaleGmOffset, blockCount, dataLen);
185	}	197	}
186	}	198	}
	@@ -195,31 +207,47 @@ __aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::SplitPostAxisCompute(
195	LocalTensor<uint8_t> tmpBuff = calcBuf.Get<uint8_t>();	207	LocalTensor<uint8_t> tmpBuff = calcBuf.Get<uint8_t>();
196		208
197	int64_t rowNumFullCount = blockCount / 32;	209	int64_t rowNumFullCount = blockCount / 32;
198	- int64_t rowNumFullAlign32 = (blockCount+31) / 32;	210	+ int64_t rowNumFullAlign32 = (blockCount + 31) / 32;
199	- int64_t rowNumFullAlign64 = (blockCount+63) / 64;	211	+ int64_t rowNumFullAlign64 = (blockCount + 63) / 64;
200	int64_t rowNumResSize = blockCount % 32;	212	int64_t rowNumResSize = blockCount % 32;
201		213
202	int64_t offset = 0;	214	int64_t offset = 0;
203	- for(int64_t i=0; i < rowNumFullCount; i++){	215	+ if (scaleAlg_ == 0) {
204	- offset = i * 32 * dataLen;	216	+ for(int64_t i = 0; i < rowNumFullCount; i++) {
205	- auto xAddr = (__ubuf__ T *)x.GetPhyAddr() + offset;	217	+ offset = i * 32 * dataLen;
206	- auto mxScaleAddr = (__ubuf__ uint8_t )tmpBuff.GetPhyAddr() + i dataLen;	218	+ auto xAddr = (__ubuf__ T *)x.GetPhyAddr() + offset;
207	- auto yAddr = (__ubuf__ uint8_t *)y.GetPhyAddr() + offset;	219	+ auto mxScaleAddr = (__ubuf__ uint8_t )tmpBuff.GetPhyAddr() + i dataLen;
208	- Compute(dataLen, 32, xAddr, mxScaleAddr, yAddr);	220	+ auto yAddr = (__ubuf__ uint8_t *)y.GetPhyAddr() + offset;
		221	+ ComputeOCP(dataLen, 32, xAddr, mxScaleAddr, yAddr);
		222	+ }
		223	+ } else if (scaleAlg_ == 1) {
		224	+ for(int64_t i = 0; i < rowNumFullCount; i++) {
		225	+ offset = i * 32 * dataLen;
		226	+ auto xAddr = (__ubuf__ T *)x.GetPhyAddr() + offset;
		227	+ auto mxScaleAddr = (__ubuf__ uint8_t )tmpBuff.GetPhyAddr() + i dataLen;
		228	+ auto yAddr = (__ubuf__ uint8_t *)y.GetPhyAddr() + offset;
		229	+ ComputecuBLAS(dataLen, 32, xAddr, mxScaleAddr, yAddr);
		230	+ }
209	}	231	}
210	- if (rowNumResSize!=0) {	232	+
		233	+ if (rowNumResSize != 0) {
211	int64_t i = rowNumFullCount;	234	int64_t i = rowNumFullCount;
212	offset = i * 32 * dataLen;	235	offset = i * 32 * dataLen;
213	auto xAddr = (__ubuf__ T *)x.GetPhyAddr() + offset;	236	auto xAddr = (__ubuf__ T *)x.GetPhyAddr() + offset;
214	auto mxScaleAddr = (__ubuf__ uint8_t )tmpBuff.GetPhyAddr() + i dataLen;	237	auto mxScaleAddr = (__ubuf__ uint8_t )tmpBuff.GetPhyAddr() + i dataLen;
215	auto yAddr = (__ubuf__ uint8_t *)y.GetPhyAddr() + offset;	238	auto yAddr = (__ubuf__ uint8_t *)y.GetPhyAddr() + offset;
216	- Compute(dataLen, rowNumResSize, xAddr, mxScaleAddr, yAddr);	239	+ if (scaleAlg_ == 0) {
		240	+ ComputeOCP(dataLen, rowNumResSize, xAddr, mxScaleAddr, yAddr);
		241	+ } else if (scaleAlg_ == 1) {
		242	+ ComputecuBLAS(dataLen, rowNumResSize, xAddr, mxScaleAddr, yAddr);
		243	+ }
217	}	244	}
218	- if (rowNumFullAlign32 % 2 != 0){	245	+ if (rowNumFullAlign32 % 2 != 0) {
219	Duplicate<uint8_t>(tmpBuff[rowNumFullAlign32 * dataLen], (uint8_t)0, 128); // e8m0_2偶数pad 0	246	Duplicate<uint8_t>(tmpBuff[rowNumFullAlign32 * dataLen], (uint8_t)0, 128); // e8m0_2偶数pad 0
220	}	247	}
221	- for(int64_t i=0; i < rowNumFullAlign64; i++){	248	+ for(int64_t i = 0; i < rowNumFullAlign64; i++){
222	- Interleave(mxScale[2idataLen], mxScale[(2i+1)dataLen], tmpBuff[2idataLen], tmpBuff[(2i+1)dataLen], dataLen);	249	+ Interleave(mxScale[2 * i * dataLen], mxScale[(2 * i + 1) * dataLen], tmpBuff[2 * i * dataLen],
		250	+ tmpBuff[(2 * i + 1) * dataLen], dataLen);
223	}	251	}
224	this->mxScaleQueue_.template EnQue(mxScale);	252	this->mxScaleQueue_.template EnQue(mxScale);
225	this->outQueue_.template EnQue(y);	253	this->outQueue_.template EnQue(y);
	@@ -227,7 +255,7 @@ __aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::SplitPostAxisCompute(
227	}	255	}
228		256
229	template <typename T, typename U>	257	template <typename T, typename U>
230	-__aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::Compute(int64_t dataLen, int64_t blockCount,	258	+__aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::ComputeOCP(int64_t dataLen, int64_t blockCount,
231	__ubuf__ T xAddr, __ubuf__ uint8_t mxScaleAddr, __ubuf__ uint8_t *yAddr)	259	__ubuf__ T xAddr, __ubuf__ uint8_t mxScaleAddr, __ubuf__ uint8_t *yAddr)
232	{	260	{
233	constexpr uint32_t vfLen = Ops::Base::GetVRegSize() / sizeof(T); // 寄存器单次处理能处理的长度	261	constexpr uint32_t vfLen = Ops::Base::GetVRegSize() / sizeof(T); // 寄存器单次处理能处理的长度
	@@ -240,219 +268,503 @@ __aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::Compute(int64_t dataL
240	uint32_t tailLoopNum1 = tailVfLen <= vfNum ? 0 : (tailVfLen - vfNum);	268	uint32_t tailLoopNum1 = tailVfLen <= vfNum ? 0 : (tailVfLen - vfNum);
241	int64_t outDataLenAlign = dataLen;	269	int64_t outDataLenAlign = dataLen;
242	uint16_t FP8_BF16_MAX_EXP = 0;	270	uint16_t FP8_BF16_MAX_EXP = 0;
		271	+ int64_t blockCountLoop = blockCount / 2;
243	if constexpr (IsSame<U, fp8_e4m3fn_t>::value){	272	if constexpr (IsSame<U, fp8_e4m3fn_t>::value){
244	FP8_BF16_MAX_EXP = FP8_E4M3_MAX_EXP;	273	FP8_BF16_MAX_EXP = FP8_E4M3_MAX_EXP;
245	} else if constexpr (IsSame<U, fp8_e5m2_t>::value){	274	} else if constexpr (IsSame<U, fp8_e5m2_t>::value){
246	FP8_BF16_MAX_EXP = FP8_E5M2_MAX_EXP;	275	FP8_BF16_MAX_EXP = FP8_E5M2_MAX_EXP;
247	}	276	}
		277	+
248	__VEC_SCOPE__	278	__VEC_SCOPE__
249	{	279	{
250	- AscendC::MicroAPI::RegTensor<T> xRegTensor;	280	+ Reg::RegTensor<T> xRegTensor;
251	- AscendC::MicroAPI::RegTensor<bfloat16_t> xBF16RegTensor;	281	+ Reg::RegTensor<uint16_t> expRegTensor;
252	- AscendC::MicroAPI::RegTensor<uint16_t> expRegTensor;	282	+ Reg::RegTensor<uint16_t> expMaxRegTensor;
253	- AscendC::MicroAPI::RegTensor<uint16_t> expMaxRegTensor;	283	+ Reg::RegTensor<T> xRegTensor0;
254	- AscendC::MicroAPI::RegTensor<uint16_t> maxEleRegTensor;	284	+ Reg::RegTensor<T> xRegTensor1;
255	- AscendC::MicroAPI::RegTensor<uint16_t> fp8MaxExpRegTensor;	285	+ Reg::RegTensor<bfloat16_t> xBF16RegTensor0;
256	- AscendC::MicroAPI::RegTensor<uint16_t> fp8NanRegTensor;	286	+ Reg::RegTensor<bfloat16_t> xBF16RegTensor1;
257	- AscendC::MicroAPI::RegTensor<uint16_t> shareExpRegTensor;	287	+ Reg::RegTensor<uint16_t> expMaxRegTensor0;
258	- AscendC::MicroAPI::RegTensor<uint16_t> mxScaleRegTensor;	288	+ Reg::RegTensor<uint16_t> expMaxRegTensor1;
259	- AscendC::MicroAPI::RegTensor<uint8_t> mxScale;	289	+ Reg::RegTensor<uint16_t> maxEleRegTensor;
260	- AscendC::MicroAPI::RegTensor<uint16_t> reversedShareExpRegTensor;	290	+ Reg::RegTensor<uint16_t> fp8MaxExpRegTensor;
261	- AscendC::MicroAPI::RegTensor<float> reversedShareExpRegTensorFP32Zero;	291	+ Reg::RegTensor<uint16_t> fp8NanRegTensor;
262	- AscendC::MicroAPI::RegTensor<float> reversedShareExpRegTensorFP32One;	292	+ Reg::RegTensor<uint16_t> mxScaleRegTensor;
263	- AscendC::MicroAPI::RegTensor<uint16_t> specialExpRegTensor;	293	+ Reg::RegTensor<uint8_t> mxScale;
264	- AscendC::MicroAPI::RegTensor<uint16_t> biasRegTensor;	294	+ Reg::RegTensor<uint16_t> reversedShareExpRegTensor;
265	- AscendC::MicroAPI::RegTensor<uint16_t> zeroRegTensor;	295	+ Reg::RegTensor<float> reversedShareExpRegTensorFP32Zero;
266	- AscendC::MicroAPI::RegTensor<uint16_t> nanRegTensor;	296	+ Reg::RegTensor<float> reversedShareExpRegTensorFP32One;
267	- AscendC::MicroAPI::RegTensor<uint16_t> yRegTensorZero;	297	+ Reg::RegTensor<uint16_t> specialExpRegTensor;
268	- AscendC::MicroAPI::RegTensor<uint16_t> yRegTensorOne;	298	+ Reg::RegTensor<uint16_t> biasRegTensor;
269	- AscendC::MicroAPI::RegTensor<uint8_t> outZero;	299	+ Reg::RegTensor<uint16_t> zeroRegTensor;
270	- AscendC::MicroAPI::RegTensor<uint8_t> outOne;	300	+ Reg::RegTensor<uint16_t> nanRegTensor;
271	- AscendC::MicroAPI::RegTensor<float> yZero;	301	+ Reg::RegTensor<uint16_t> yRegTensorZero;
272	- AscendC::MicroAPI::RegTensor<float> yOne;	302	+ Reg::RegTensor<uint16_t> yRegTensorOne;
273	- AscendC::MicroAPI::RegTensor<U> yZeroFP8;	303	+ Reg::RegTensor<uint8_t> outZero;
274	- AscendC::MicroAPI::RegTensor<U> yOneFP8;	304	+ Reg::RegTensor<uint8_t> outOne;
275	- AscendC::MicroAPI::RegTensor<bfloat16_t> valueRegTensor;	305	+ Reg::RegTensor<float> yZero;
276	- AscendC::MicroAPI::UnalignReg u0;	306	+ Reg::RegTensor<float> yOne;
277	- AscendC::MicroAPI::UnalignReg u1;	307	+ Reg::RegTensor<U> yZeroFP8;
278	- AscendC::MicroAPI::MaskReg infMask;	308	+ Reg::RegTensor<U> yOneFP8;
279	- AscendC::MicroAPI::MaskReg zeroMask;	309	+ Reg::RegTensor<bfloat16_t> valueRegTensor;
280	- AscendC::MicroAPI::MaskReg invalidDataMask;	310	+
281	- AscendC::MicroAPI::MaskReg specialDataMask;	311	+ Reg::UnalignReg u0;
282	- AscendC::MicroAPI::MaskReg maskAll = AscendC::MicroAPI::CreateMask<uint16_t, AscendC::MicroAPI::MaskPattern::ALL>();	312	+ Reg::MaskReg infMask;
283	- static constexpr AscendC::MicroAPI::CastTrait castTraitZero = {	313	+ Reg::MaskReg zeroMask;
284	- AscendC::MicroAPI::RegLayout::ZERO, AscendC::MicroAPI::SatMode::UNKNOWN,	314	+ Reg::MaskReg invalidDataMask;
285	- AscendC::MicroAPI::MaskMergeMode::ZEROING, RoundMode::UNKNOWN};	315	+ Reg::MaskReg specialDataMask;
286	- static constexpr AscendC::MicroAPI::CastTrait castTraitOne = {	316	+ Reg::MaskReg maskAll = Reg::CreateMask<uint16_t, Reg::MaskPattern::ALL>();
287	- AscendC::MicroAPI::RegLayout::ONE, AscendC::MicroAPI::SatMode::UNKNOWN,	317	+
288	- AscendC::MicroAPI::MaskMergeMode::ZEROING, RoundMode::UNKNOWN};	318	+ static constexpr Reg::CastTrait castTraitZero = {
289	- static constexpr AscendC::MicroAPI::CastTrait castTrait32to8 = {	319	+ Reg::RegLayout::ZERO, Reg::SatMode::UNKNOWN,
290	- AscendC::MicroAPI::RegLayout::ZERO, AscendC::MicroAPI::SatMode::SAT,	320	+ Reg::MaskMergeMode::ZEROING, RoundMode::UNKNOWN};
291	- AscendC::MicroAPI::MaskMergeMode::ZEROING, RoundMode::CAST_RINT};	321	+ static constexpr Reg::CastTrait castTraitOne = {
292	- static constexpr AscendC::MicroAPI::CastTrait castTraitHalf2Bf16 = {	322	+ Reg::RegLayout::ONE, Reg::SatMode::UNKNOWN,
293	- AscendC::MicroAPI::RegLayout::UNKNOWN, AscendC::MicroAPI::SatMode::UNKNOWN,	323	+ Reg::MaskMergeMode::ZEROING, RoundMode::UNKNOWN};
294	- AscendC::MicroAPI::MaskMergeMode::ZEROING, RoundMode::CAST_TRUNC};	324	+ static constexpr Reg::CastTrait castTrait32to8 = {
295	- AscendC::MicroAPI::Duplicate(maxEleRegTensor, MAX_EXP_FOR_BF16);	325	+ Reg::RegLayout::ZERO, Reg::SatMode::SAT,
296	- AscendC::MicroAPI::Duplicate(fp8MaxExpRegTensor, FP8_BF16_MAX_EXP);	326	+ Reg::MaskMergeMode::ZEROING, RoundMode::CAST_RINT};
297	- AscendC::MicroAPI::Duplicate(fp8NanRegTensor, MAX_EXP_FOR_FP8);	327	+ static constexpr Reg::CastTrait castTraitHalf2Bf16 = {
298	- AscendC::MicroAPI::Duplicate(biasRegTensor, BF16_EXP_BIAS);	328	+ Reg::RegLayout::UNKNOWN, Reg::SatMode::UNKNOWN,
299	- AscendC::MicroAPI::Duplicate(zeroRegTensor, 0);	329	+ Reg::MaskMergeMode::ZEROING, RoundMode::CAST_TRUNC};
300	- AscendC::MicroAPI::Duplicate(nanRegTensor, NAN_CUSTOMIZATION);	330	+
301	- AscendC::MicroAPI::Duplicate(specialExpRegTensor, SPECIAL_EXP_THRESHOLD);	331	+ Reg::Duplicate(maxEleRegTensor, MAX_EXP_FOR_BF16);
		332	+ Reg::Duplicate(fp8MaxExpRegTensor, FP8_BF16_MAX_EXP);
		333	+ Reg::Duplicate(fp8NanRegTensor, MAX_EXP_FOR_FP8);
		334	+ Reg::Duplicate(biasRegTensor, BF16_EXP_BIAS);
		335	+ Reg::Duplicate(zeroRegTensor, 0);
		336	+ Reg::Duplicate(nanRegTensor, NAN_CUSTOMIZATION);
		337	+ Reg::Duplicate(specialExpRegTensor, SPECIAL_EXP_THRESHOLD);
		338	+ Reg::Duplicate(expMaxRegTensor, 0);
		339	+
302	for (uint16_t i = 0; i < regLoop; i++) {	340	for (uint16_t i = 0; i < regLoop; i++) {
303	uint32_t pnum = vfLen;	341	uint32_t pnum = vfLen;
304	- AscendC::MicroAPI::MaskReg p0 = AscendC::MicroAPI::UpdateMask<uint16_t>(pnum);	342	+ Reg::MaskReg p0 = Reg::UpdateMask<uint16_t>(pnum);
305	- DataCopy(xRegTensor, xAddr + i * vfLen);	343	+ for (uint16_t j = 0; j <= static_cast<uint16_t>(blockCountLoop); j++) {
306	- if constexpr (IsSame<T, half>::value) {	344	+ Reg::LoadAlign(xRegTensor0, xAddr + j * dataLen + i * vfLen);
307	- AscendC::MicroAPI::Cast<bfloat16_t, T, castTraitHalf2Bf16>(xBF16RegTensor, xRegTensor, p0);	345	+ Reg::LoadAlign(xRegTensor1, xAddr + (static_cast<uint16_t>(blockCount) - j - 1) * dataLen + i * vfLen);
308	- AscendC::MicroAPI::And(expMaxRegTensor, (AscendC::MicroAPI::RegTensor<uint16_t>&)xBF16RegTensor,
309	- maxEleRegTensor, p0);
310	- } else {
311	- AscendC::MicroAPI::And(expMaxRegTensor, (AscendC::MicroAPI::RegTensor<uint16_t>&)xRegTensor,
312	- maxEleRegTensor, p0);
313	- }
314	- for (uint16_t j = 1; j < static_cast<uint16_t>(blockCount); j++) {
315	- DataCopy(xRegTensor, xAddr + j * dataLen + i * vfLen);
316	if constexpr (IsSame<T, half>::value) {	346	if constexpr (IsSame<T, half>::value) {
317	- AscendC::MicroAPI::Cast<bfloat16_t, T, castTraitHalf2Bf16>(xBF16RegTensor, xRegTensor, p0);	347	+ Reg::Cast<bfloat16_t, T, castTraitHalf2Bf16>(xBF16RegTensor0, xRegTensor0, p0);
318	- AscendC::MicroAPI::And(expRegTensor, (AscendC::MicroAPI::RegTensor<uint16_t>&)xBF16RegTensor,	348	+ Reg::Cast<bfloat16_t, T, castTraitHalf2Bf16>(xBF16RegTensor1, xRegTensor1, p0);
319	- maxEleRegTensor, p0);	349	+ Reg::And(expMaxRegTensor0, (Reg::RegTensor<uint16_t>&)xBF16RegTensor0, maxEleRegTensor, p0);
		350	+ Reg::And(expMaxRegTensor1, (Reg::RegTensor<uint16_t>&)xBF16RegTensor1, maxEleRegTensor, p0);
320	} else {	351	} else {
321	- AscendC::MicroAPI::And(expRegTensor, (AscendC::MicroAPI::RegTensor<uint16_t>&)xRegTensor,	352	+ Reg::And(expMaxRegTensor0, (Reg::RegTensor<uint16_t>&)xRegTensor0, maxEleRegTensor, p0);
322	- maxEleRegTensor, p0);	353	+ Reg::And(expMaxRegTensor1, (Reg::RegTensor<uint16_t>&)xRegTensor1, maxEleRegTensor, p0);
323	}	354	}
324	- AscendC::MicroAPI::Max(expMaxRegTensor, expMaxRegTensor, expRegTensor, p0);	355	+ Reg::Max(expRegTensor, expMaxRegTensor0, expMaxRegTensor1, p0);
		356	+ Reg::Max(expMaxRegTensor, expRegTensor, expMaxRegTensor, p0);
325	}	357	}
326	- AscendC::MicroAPI::Compare<uint16_t, CMPMODE::NE>(infMask, expMaxRegTensor, maxEleRegTensor, p0);	358	+ Reg::Compare<uint16_t, CMPMODE::NE>(infMask, expMaxRegTensor, maxEleRegTensor, p0);
327	- AscendC::MicroAPI::Compare<uint16_t, CMPMODE::NE>(zeroMask, expMaxRegTensor, zeroRegTensor, p0);	359	+ Reg::Compare<uint16_t, CMPMODE::NE>(zeroMask, expMaxRegTensor, zeroRegTensor, p0);
328	- AscendC::MicroAPI::Compare<uint16_t, CMPMODE::LE>(invalidDataMask, expMaxRegTensor,	360	+ Reg::Compare<uint16_t, CMPMODE::LE>(invalidDataMask, expMaxRegTensor, fp8MaxExpRegTensor, p0);
329	- fp8MaxExpRegTensor, p0);	361	+ Reg::Select<uint16_t>(expMaxRegTensor, fp8MaxExpRegTensor, expMaxRegTensor, invalidDataMask);
330	- AscendC::MicroAPI::Select<uint16_t>(expMaxRegTensor, fp8MaxExpRegTensor, expMaxRegTensor,	362	+ Reg::Sub(expMaxRegTensor, expMaxRegTensor, fp8MaxExpRegTensor, p0);
331	- invalidDataMask);	363	+ Reg::ShiftRights(mxScaleRegTensor, expMaxRegTensor, SHR_NUM_FOR_BF16, p0);
332	- AscendC::MicroAPI::Sub(expMaxRegTensor, expMaxRegTensor, fp8MaxExpRegTensor, p0);	364	+ Reg::Select<uint16_t>(mxScaleRegTensor, mxScaleRegTensor, fp8NanRegTensor, infMask);
333	- AscendC::MicroAPI::ShiftRights(mxScaleRegTensor, expMaxRegTensor, SHR_NUM_FOR_BF16, p0);	365	+ Reg::Select<uint16_t>(mxScaleRegTensor, mxScaleRegTensor, zeroRegTensor, zeroMask);
334	- AscendC::MicroAPI::Select<uint16_t>(mxScaleRegTensor, mxScaleRegTensor, fp8NanRegTensor, infMask);	366	+ Reg::Pack(mxScale, mxScaleRegTensor);
335	- AscendC::MicroAPI::Select<uint16_t>(mxScaleRegTensor, mxScaleRegTensor, zeroRegTensor, zeroMask);	367	+ Reg::StoreUnAlign(mxScaleAddr, mxScale, u0, vfLen);
336	- AscendC::MicroAPI::Pack(mxScale, mxScaleRegTensor);	368	+ Reg::StoreUnAlignPost(mxScaleAddr, u0, 0);
337	- AscendC::MicroAPI::DataCopyUnAlign(mxScaleAddr, mxScale, u1, vfLen);	369	+
338	- AscendC::MicroAPI::DataCopyUnAlignPost(mxScaleAddr, u1, 0);
339	// 求1/scale	370	// 求1/scale
340	- AscendC::MicroAPI::Compare<uint16_t, CMPMODE::EQ>(specialDataMask, expMaxRegTensor, biasRegTensor, p0);	371	+ Reg::Compare<uint16_t, CMPMODE::EQ>(specialDataMask, expMaxRegTensor, biasRegTensor, p0);
341	- AscendC::MicroAPI::Sub(reversedShareExpRegTensor, biasRegTensor, expMaxRegTensor, p0);	372	+ Reg::Sub(reversedShareExpRegTensor, biasRegTensor, expMaxRegTensor, p0);
342	- AscendC::MicroAPI::Select<uint16_t>(reversedShareExpRegTensor, reversedShareExpRegTensor, nanRegTensor,	373	+ Reg::Select<uint16_t>(reversedShareExpRegTensor, reversedShareExpRegTensor, nanRegTensor, infMask);
343	- infMask);	374	+ Reg::Select<uint16_t>(reversedShareExpRegTensor, reversedShareExpRegTensor, zeroRegTensor, zeroMask);
344	- AscendC::MicroAPI::Select<uint16_t>(reversedShareExpRegTensor, reversedShareExpRegTensor, zeroRegTensor,	375	+ Reg::Select<uint16_t>(reversedShareExpRegTensor, specialExpRegTensor, reversedShareExpRegTensor, specialDataMask);
345	- zeroMask);
346	- AscendC::MicroAPI::Select<uint16_t>(reversedShareExpRegTensor, specialExpRegTensor,
347	- reversedShareExpRegTensor, specialDataMask);
348		376
349	// 求data value	377	// 求data value
		378	+ if constexpr (IsSame<T, half>::value) {
		379	+ Reg::Cast<float, bfloat16_t, castTraitZero>(reversedShareExpRegTensorFP32Zero,
		380	+ (Reg::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p0);
		381	+ Reg::Cast<float, bfloat16_t, castTraitOne>(reversedShareExpRegTensorFP32One,
		382	+ (Reg::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p0);
		383	+ }
350	for (uint16_t j = 0; j < static_cast<uint16_t>(blockCount); j++) {	384	for (uint16_t j = 0; j < static_cast<uint16_t>(blockCount); j++) {
351	- DataCopy(xRegTensor, xAddr + j * dataLen + i * vfLen);	385	+ Reg::LoadAlign(xRegTensor, xAddr + j * dataLen + i * vfLen);
352	if constexpr (IsSame<T, half>::value) {	386	if constexpr (IsSame<T, half>::value) {
353	- AscendC::MicroAPI::Cast<float, T, castTraitZero>(yZero, xRegTensor, p0);	387	+ Reg::Cast<float, T, castTraitZero>(yZero, xRegTensor, p0);
354	- AscendC::MicroAPI::Cast<float, T, castTraitOne>(yOne, xRegTensor, p0);	388	+ Reg::Cast<float, T, castTraitOne>(yOne, xRegTensor, p0);
355	- AscendC::MicroAPI::Cast<float, bfloat16_t, castTraitZero>(reversedShareExpRegTensorFP32Zero, (AscendC::MicroAPI::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p0);	389	+ Reg::Mul(yZero, yZero, reversedShareExpRegTensorFP32Zero, maskAll);
356	- AscendC::MicroAPI::Cast<float, bfloat16_t, castTraitOne>(reversedShareExpRegTensorFP32One, (AscendC::MicroAPI::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p0);	390	+ Reg::Mul(yOne, yOne, reversedShareExpRegTensorFP32One, maskAll);
357	- AscendC::MicroAPI::Mul(yZero, yZero, reversedShareExpRegTensorFP32Zero, maskAll);
358	- AscendC::MicroAPI::Mul(yOne, yOne, reversedShareExpRegTensorFP32One, maskAll);
359	} else {	391	} else {
360	- AscendC::MicroAPI::Mul(valueRegTensor, xRegTensor,	392	+ Reg::Mul(valueRegTensor, xRegTensor,
361	- (AscendC::MicroAPI::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p0);	393	+ (Reg::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p0);
362	- AscendC::MicroAPI::Cast<float, bfloat16_t, castTraitZero>(yZero, valueRegTensor, maskAll);	394	+ Reg::Cast<float, bfloat16_t, castTraitZero>(yZero, valueRegTensor, maskAll);
363	- AscendC::MicroAPI::Cast<float, bfloat16_t, castTraitOne>(yOne, valueRegTensor, maskAll);	395	+ Reg::Cast<float, bfloat16_t, castTraitOne>(yOne, valueRegTensor, maskAll);
364	}	396	}
365	- AscendC::MicroAPI::Interleave(yZero, yOne, yZero, yOne);	397	+ Reg::Interleave(yZero, yOne, yZero, yOne);
366	- AscendC::MicroAPI::Cast<U, float, castTrait32to8>(yZeroFP8, yZero, maskAll);	398	+ Reg::Cast<U, float, castTrait32to8>(yZeroFP8, yZero, maskAll);
367	- AscendC::MicroAPI::Cast<U, float, castTrait32to8>(yOneFP8, yOne, maskAll);	399	+ Reg::Cast<U, float, castTrait32to8>(yOneFP8, yOne, maskAll);
368	- AscendC::MicroAPI::Pack(yRegTensorZero, (AscendC::MicroAPI::RegTensor<uint32_t>&)yZeroFP8);	400	+ Reg::Pack(yRegTensorZero, (Reg::RegTensor<uint32_t>&)yZeroFP8);
369	- AscendC::MicroAPI::Pack(outZero, yRegTensorZero);	401	+ Reg::Pack(outZero, yRegTensorZero);
370	- AscendC::MicroAPI::Pack(yRegTensorOne, (AscendC::MicroAPI::RegTensor<uint32_t>&)yOneFP8);	402	+ Reg::Pack(yRegTensorOne, (Reg::RegTensor<uint32_t>&)yOneFP8);
371	- AscendC::MicroAPI::Pack(outOne, yRegTensorOne);	403	+ Reg::Pack(outOne, yRegTensorOne);
372	auto addr0 = yAddr + (j * outDataLenAlign + i * vfLen);	404	auto addr0 = yAddr + (j * outDataLenAlign + i * vfLen);
373	- AscendC::MicroAPI::DataCopyUnAlign(addr0, outZero, u1, loopNum0);	405	+ Reg::StoreUnAlign(addr0, outZero, u0, loopNum0);
374	- AscendC::MicroAPI::DataCopyUnAlignPost(addr0, u1, 0);	406	+ Reg::StoreUnAlignPost(addr0, u0, 0);
375	auto addr1 = yAddr + (j * outDataLenAlign + i * vfLen) + loopNum0;	407	auto addr1 = yAddr + (j * outDataLenAlign + i * vfLen) + loopNum0;
376	- AscendC::MicroAPI::DataCopyUnAlign(addr1, outOne, u1, loopNum1);	408	+ Reg::StoreUnAlign(addr1, outOne, u0, loopNum1);
377	- AscendC::MicroAPI::DataCopyUnAlignPost(addr1, u1, 0);	409	+ Reg::StoreUnAlignPost(addr1, u0, 0);
378	}	410	}
379	}	411	}
380		412
381	if (tailVfLen != 0) {	413	if (tailVfLen != 0) {
382	uint32_t tailPnum = tailVfLen;	414	uint32_t tailPnum = tailVfLen;
383	- AscendC::MicroAPI::MaskReg p1 = AscendC::MicroAPI::UpdateMask<T>(tailPnum);	415	+ Reg::MaskReg p1 = Reg::UpdateMask<T>(tailPnum);
384	- DataCopy(xRegTensor, xAddr + regLoop * vfLen);	416	+ for (uint16_t k = 0; k <= static_cast<uint16_t>(blockCountLoop); k++) {
385	- if constexpr (IsSame<T, half>::value) {	417	+ Reg::LoadAlign(xRegTensor0, xAddr + k * dataLen + regLoop * vfLen);
386	- AscendC::MicroAPI::Cast<bfloat16_t, T, castTraitHalf2Bf16>(xBF16RegTensor, xRegTensor, p1);	418	+ Reg::LoadAlign(xRegTensor1, xAddr + (static_cast<uint16_t>(blockCount) - k - 1) * dataLen + regLoop * vfLen);
387	- AscendC::MicroAPI::And(expMaxRegTensor, (AscendC::MicroAPI::RegTensor<uint16_t>&)xBF16RegTensor,
388	- maxEleRegTensor, p1);
389	- } else {
390	- AscendC::MicroAPI::And(expMaxRegTensor, (AscendC::MicroAPI::RegTensor<uint16_t>&)xRegTensor,
391	- maxEleRegTensor, p1);
392	- }
393	- for (uint16_t j = 1; j < static_cast<uint16_t>(blockCount); j++) {
394	- DataCopy(xRegTensor, xAddr + regLoop * vfLen + j * dataLen);
395	if constexpr (IsSame<T, half>::value) {	419	if constexpr (IsSame<T, half>::value) {
396	- AscendC::MicroAPI::Cast<bfloat16_t, T, castTraitHalf2Bf16>(xBF16RegTensor, xRegTensor, p1);	420	+ Reg::Cast<bfloat16_t, T, castTraitHalf2Bf16>(xBF16RegTensor0, xRegTensor0, p1);
397	- AscendC::MicroAPI::And(expRegTensor, (AscendC::MicroAPI::RegTensor<uint16_t>&)xBF16RegTensor,	421	+ Reg::Cast<bfloat16_t, T, castTraitHalf2Bf16>(xBF16RegTensor1, xRegTensor1, p1);
398	- maxEleRegTensor, p1);	422	+ Reg::And(expMaxRegTensor0, (Reg::RegTensor<uint16_t>&)xBF16RegTensor0, maxEleRegTensor, p1);
		423	+ Reg::And(expMaxRegTensor1, (Reg::RegTensor<uint16_t>&)xBF16RegTensor1, maxEleRegTensor, p1);
399	} else {	424	} else {
400	- AscendC::MicroAPI::And(expRegTensor, (AscendC::MicroAPI::RegTensor<uint16_t>&)xRegTensor,	425	+ Reg::And(expMaxRegTensor0, (Reg::RegTensor<uint16_t>&)xRegTensor0, maxEleRegTensor, p1);
401	- maxEleRegTensor, p1);	426	+ Reg::And(expMaxRegTensor1, (Reg::RegTensor<uint16_t>&)xRegTensor1, maxEleRegTensor, p1);
402	- }	427	+ }
403	- AscendC::MicroAPI::Max(expMaxRegTensor, expMaxRegTensor, expRegTensor, p1);	428	+ Reg::Max(expRegTensor, expMaxRegTensor0, expMaxRegTensor1, p1);
		429	+ Reg::Max(expMaxRegTensor, expRegTensor, expMaxRegTensor, p1);
404	}	430	}
405	- AscendC::MicroAPI::Compare<uint16_t, CMPMODE::NE>(infMask, expMaxRegTensor, maxEleRegTensor, p1);	431	+ Reg::Compare<uint16_t, CMPMODE::NE>(infMask, expMaxRegTensor, maxEleRegTensor, p1);
406	- AscendC::MicroAPI::Compare<uint16_t, CMPMODE::NE>(zeroMask, expMaxRegTensor, zeroRegTensor, p1);	432	+ Reg::Compare<uint16_t, CMPMODE::NE>(zeroMask, expMaxRegTensor, zeroRegTensor, p1);
407	- AscendC::MicroAPI::Compare<uint16_t, CMPMODE::LE>(invalidDataMask, expMaxRegTensor,	433	+ Reg::Compare<uint16_t, CMPMODE::LE>(invalidDataMask, expMaxRegTensor, fp8MaxExpRegTensor, p1);
408	- fp8MaxExpRegTensor, p1);	434	+ Reg::Select<uint16_t>(expMaxRegTensor, fp8MaxExpRegTensor, expMaxRegTensor, invalidDataMask);
409	- AscendC::MicroAPI::Select<uint16_t>(expMaxRegTensor, fp8MaxExpRegTensor, expMaxRegTensor,	435	+ Reg::Sub(expMaxRegTensor, expMaxRegTensor, fp8MaxExpRegTensor, p1);
410	- invalidDataMask);	436	+ Reg::ShiftRights(mxScaleRegTensor, expMaxRegTensor, SHR_NUM_FOR_BF16, p1);
411	- AscendC::MicroAPI::Sub(expMaxRegTensor, expMaxRegTensor, fp8MaxExpRegTensor, p1);	437	+ Reg::Select<uint16_t>(mxScaleRegTensor, mxScaleRegTensor, fp8NanRegTensor, infMask);
412	- AscendC::MicroAPI::ShiftRights(mxScaleRegTensor, expMaxRegTensor, SHR_NUM_FOR_BF16, p1);	438	+ Reg::Select<uint16_t>(mxScaleRegTensor, mxScaleRegTensor, zeroRegTensor, zeroMask);
413	- AscendC::MicroAPI::Select<uint16_t>(mxScaleRegTensor, mxScaleRegTensor, fp8NanRegTensor, infMask);	439	+ Reg::Pack(mxScale, mxScaleRegTensor);
414	- AscendC::MicroAPI::Select<uint16_t>(mxScaleRegTensor, mxScaleRegTensor, zeroRegTensor, zeroMask);	440	+ Reg::StoreUnAlign(mxScaleAddr, mxScale, u0, tailVfLen);
415	- AscendC::MicroAPI::Pack(mxScale, mxScaleRegTensor);	441	+ Reg::StoreUnAlignPost(mxScaleAddr, u0, 0);
416	- AscendC::MicroAPI::DataCopyUnAlign(mxScaleAddr, mxScale, u1, tailVfLen);	442	+
417	- AscendC::MicroAPI::DataCopyUnAlignPost(mxScaleAddr, u1, 0);
418	// 求1/scale	443	// 求1/scale
419	- AscendC::MicroAPI::Compare<uint16_t, CMPMODE::EQ>(specialDataMask, expMaxRegTensor, biasRegTensor, p1);	444	+ Reg::Compare<uint16_t, CMPMODE::EQ>(specialDataMask, expMaxRegTensor, biasRegTensor, p1);
420	- AscendC::MicroAPI::Sub(reversedShareExpRegTensor, biasRegTensor, expMaxRegTensor, p1);	445	+ Reg::Sub(reversedShareExpRegTensor, biasRegTensor, expMaxRegTensor, p1);
421	- AscendC::MicroAPI::Select<uint16_t>(reversedShareExpRegTensor, specialExpRegTensor,	446	+ Reg::Select<uint16_t>(reversedShareExpRegTensor, specialExpRegTensor,
422	reversedShareExpRegTensor, specialDataMask);	447	reversedShareExpRegTensor, specialDataMask);
423	- AscendC::MicroAPI::Select<uint16_t>(reversedShareExpRegTensor, reversedShareExpRegTensor, nanRegTensor,	448	+ Reg::Select<uint16_t>(reversedShareExpRegTensor, reversedShareExpRegTensor, nanRegTensor,
424	infMask);	449	infMask);
425	- AscendC::MicroAPI::Select<uint16_t>(reversedShareExpRegTensor, reversedShareExpRegTensor, zeroRegTensor,	450	+ Reg::Select<uint16_t>(reversedShareExpRegTensor, reversedShareExpRegTensor, zeroRegTensor,
426	zeroMask);	451	zeroMask);
		452	+
427	// 求data value	453	// 求data value
		454	+ if constexpr (IsSame<T, half>::value) {
		455	+ Reg::Cast<float, bfloat16_t, castTraitZero>(reversedShareExpRegTensorFP32Zero,
		456	+ (Reg::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p1);
		457	+ Reg::Cast<float, bfloat16_t, castTraitOne>(reversedShareExpRegTensorFP32One,
		458	+ (Reg::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p1);
		459	+ }
428	for (uint16_t j = 0; j < static_cast<uint16_t>(blockCount); j++) {	460	for (uint16_t j = 0; j < static_cast<uint16_t>(blockCount); j++) {
429	- DataCopy(xRegTensor, xAddr + regLoop * vfLen + j * dataLen);	461	+ Reg::LoadAlign(xRegTensor, xAddr + regLoop * vfLen + j * dataLen);
430	if constexpr (IsSame<T, half>::value) {	462	if constexpr (IsSame<T, half>::value) {
431	- AscendC::MicroAPI::Cast<float, T, castTraitZero>(yZero, xRegTensor, p1);	463	+ Reg::Cast<float, T, castTraitZero>(yZero, xRegTensor, p1);
432	- AscendC::MicroAPI::Cast<float, T, castTraitOne>(yOne, xRegTensor, p1);	464	+ Reg::Cast<float, T, castTraitOne>(yOne, xRegTensor, p1);
433	- AscendC::MicroAPI::Cast<float, bfloat16_t, castTraitZero>(reversedShareExpRegTensorFP32Zero, (AscendC::MicroAPI::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p1);	465	+ Reg::Mul(yZero, yZero, reversedShareExpRegTensorFP32Zero, maskAll);
434	- AscendC::MicroAPI::Cast<float, bfloat16_t, castTraitOne>(reversedShareExpRegTensorFP32One, (AscendC::MicroAPI::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p1);	466	+ Reg::Mul(yOne, yOne, reversedShareExpRegTensorFP32One, maskAll);
435	- AscendC::MicroAPI::Mul(yZero, yZero, reversedShareExpRegTensorFP32Zero, maskAll);
436	- AscendC::MicroAPI::Mul(yOne, yOne, reversedShareExpRegTensorFP32One, maskAll);
437	} else {	467	} else {
438	- AscendC::MicroAPI::Mul(valueRegTensor, xRegTensor,	468	+ Reg::Mul(valueRegTensor, xRegTensor,
439	- (AscendC::MicroAPI::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p1);	469	+ (Reg::RegTensor<bfloat16_t>&)reversedShareExpRegTensor, p1);
440	- AscendC::MicroAPI::Cast<float, bfloat16_t, castTraitZero>(yZero, valueRegTensor, maskAll);	470	+ Reg::Cast<float, bfloat16_t, castTraitZero>(yZero, valueRegTensor, maskAll);
441	- AscendC::MicroAPI::Cast<float, bfloat16_t, castTraitOne>(yOne, valueRegTensor, maskAll);	471	+ Reg::Cast<float, bfloat16_t, castTraitOne>(yOne, valueRegTensor, maskAll);
442	}	472	}
443	- AscendC::MicroAPI::Interleave(yZero, yOne, yZero, yOne);	473	+ Reg::Interleave(yZero, yOne, yZero, yOne);
444	- AscendC::MicroAPI::Cast<U, float, castTrait32to8>(yZeroFP8, yZero, maskAll);	474	+ Reg::Cast<U, float, castTrait32to8>(yZeroFP8, yZero, maskAll);
445	- AscendC::MicroAPI::Cast<U, float, castTrait32to8>(yOneFP8, yOne, maskAll);	475	+ Reg::Cast<U, float, castTrait32to8>(yOneFP8, yOne, maskAll);
446	- AscendC::MicroAPI::Pack(yRegTensorZero, (AscendC::MicroAPI::RegTensor<uint32_t>&)yZeroFP8);	476	+ Reg::Pack(yRegTensorZero, (Reg::RegTensor<uint32_t>&)yZeroFP8);
447	- AscendC::MicroAPI::Pack(outZero, yRegTensorZero);	477	+ Reg::Pack(outZero, yRegTensorZero);
448	- AscendC::MicroAPI::Pack(yRegTensorOne, (AscendC::MicroAPI::RegTensor<uint32_t>&)yOneFP8);	478	+ Reg::Pack(yRegTensorOne, (Reg::RegTensor<uint32_t>&)yOneFP8);
449	- AscendC::MicroAPI::Pack(outOne, yRegTensorOne);	479	+ Reg::Pack(outOne, yRegTensorOne);
450	auto addr0 = yAddr + (regLoop * vfLen + j * outDataLenAlign);	480	auto addr0 = yAddr + (regLoop * vfLen + j * outDataLenAlign);
451	- AscendC::MicroAPI::DataCopyUnAlign(addr0, outZero, u1, tailLoopNum0);	481	+ Reg::StoreUnAlign(addr0, outZero, u0, tailLoopNum0);
452	- AscendC::MicroAPI::DataCopyUnAlignPost(addr0, u1, 0);	482	+ Reg::StoreUnAlignPost(addr0, u0, 0);
453	auto addr1 = yAddr + (regLoop * vfLen + j * outDataLenAlign) + tailLoopNum0;	483	auto addr1 = yAddr + (regLoop * vfLen + j * outDataLenAlign) + tailLoopNum0;
454	- AscendC::MicroAPI::DataCopyUnAlign(addr1, outOne, u1, tailLoopNum1);	484	+ Reg::StoreUnAlign(addr1, outOne, u0, tailLoopNum1);
455	- AscendC::MicroAPI::DataCopyUnAlignPost(addr1, u1, 0);	485	+ Reg::StoreUnAlignPost(addr1, u0, 0);
		486	+ }
		487	+ }
		488	+ }
		489	+}
		490	+
		491	+template <typename T, typename U>
		492	+__aicore__ inline void GroupedDynamicMxQuantBaseFP8<T, U>::ComputecuBLAS(int64_t dataLen, int64_t blockCount,
		493	+ __ubuf__ T xAddr, __ubuf__ uint8_t mxScaleAddr, __ubuf__ uint8_t *yAddr)
		494	+{
		495	+ constexpr uint32_t vfNum16 = Ops::Base::GetVRegSize() / sizeof(T);
		496	+ constexpr uint32_t vfNum32 = Ops::Base::GetVRegSize() / sizeof(float);
		497	+ uint16_t regLoop = static_cast<uint16_t>(dataLen) / static_cast<uint16_t>(vfNum32);
		498	+ uint16_t tailVfLen = static_cast<uint16_t>(dataLen) % static_cast<uint16_t>(vfNum32);
		499	+ uint32_t singleLoopNum = dataLen <= vfNum32 ? vfNum16 : vfNum32;
		500	+ uint32_t singleTailLoopNum = tailVfLen;
		501	+ int64_t outDataLenAlign = (dataLen + 32 - 1) / 32 * 32;
		502	+ float inv_dtype_max = 0;
		503	+ int64_t blockCountLoop = blockCount / 2;
		504	+ if constexpr (IsSame<U, fp8_e4m3fn_t>::value) {
		505	+ inv_dtype_max = FP8_E4M3_INV_MAX; // 1.0f / 448.0f
		506	+ } else if constexpr (IsSame<U, fp8_e5m2_t>::value) {
		507	+ inv_dtype_max = FP8_E5M2_INV_MAX; // 1.0f / 57344.0f
		508	+ }
		509	+
		510	+ __VEC_SCOPE__
		511	+ {
		512	+ Reg::RegTensor<T> xRegTensor;
		513	+ Reg::RegTensor<T> xRegTensor0;
		514	+ Reg::RegTensor<T> xRegTensor1;
		515	+ Reg::RegTensor<uint16_t> absMaskBF16Reg;
		516	+ Reg::RegTensor<uint16_t> fullZeroReg;
		517	+ Reg::RegTensor<uint32_t> zeroRegTensor32;
		518	+ Reg::RegTensor<uint32_t> manMaskFP32Reg;
		519	+ Reg::RegTensor<uint32_t> fp8NanRegTensor32;
		520	+ Reg::RegTensor<uint32_t> biasRegTensor32;
		521	+ Reg::RegTensor<uint32_t> nanRegTensor32;
		522	+ Reg::RegTensor<uint16_t> expScaleMxYZeroRegTensor;
		523	+ Reg::RegTensor<uint16_t> expMaxRegTensor;
		524	+ Reg::RegTensor<uint16_t> expMaxRegTensor0;
		525	+ Reg::RegTensor<uint16_t> expMaxRegTensor1;
		526	+ Reg::RegTensor<uint16_t> expMaxreShareExpFP16YOneRegTensor;
		527	+ Reg::RegTensor<uint32_t> expMaxAndAddOneFP32RegTensor;
		528	+ Reg::RegTensor<uint32_t> expAndreShareExpFP32RegTensor;
		529	+ Reg::RegTensor<uint32_t> manAndmxScaleFP32RegTensor;
		530	+ Reg::RegTensor<float> yZero;
		531	+ Reg::RegTensor<float> yOne;
		532	+ Reg::RegTensor<float> reversedShareExpRegTensorFP32Zero;
		533	+ Reg::RegTensor<float> reversedShareExpRegTensorFP32One;
		534	+ Reg::RegTensor<bfloat16_t> valueRegTensor;
		535	+ Reg::RegTensor<U> yZeroFP8;
		536	+ Reg::RegTensor<uint8_t> outZeromxScaleFp8;
		537	+
		538	+ Reg::UnalignReg u1;
		539	+ Reg::MaskReg infMask;
		540	+ Reg::MaskReg zeroMask;
		541	+ Reg::MaskReg p0;
		542	+ Reg::MaskReg p1;
		543	+ Reg::MaskReg p2;
		544	+ Reg::MaskReg p3;
		545	+ Reg::MaskReg preMaskScale = Reg::CreateMask<uint32_t>();
		546	+ Reg::MaskReg maskAll = Reg::CreateMask<uint16_t, Reg::MaskPattern::ALL>();
		547	+
		548	+ static constexpr Reg::CastTrait castTraitZero = {
		549	+ Reg::RegLayout::ZERO, Reg::SatMode::UNKNOWN, Reg::MaskMergeMode::ZEROING, RoundMode::UNKNOWN};
		550	+ static constexpr Reg::CastTrait castTraitOne = {
		551	+ Reg::RegLayout::ONE, Reg::SatMode::UNKNOWN, Reg::MaskMergeMode::ZEROING, RoundMode::UNKNOWN};
		552	+ static constexpr Reg::CastTrait castTrait32to8 = {
		553	+ Reg::RegLayout::ZERO, Reg::SatMode::SAT, Reg::MaskMergeMode::ZEROING, RoundMode::CAST_RINT};
		554	+ static constexpr Reg::CastTrait castTraitHalf2Bf16 = {
		555	+ Reg::RegLayout::UNKNOWN, Reg::SatMode::UNKNOWN, Reg::MaskMergeMode::ZEROING, RoundMode::CAST_TRUNC};
		556	+ static constexpr Reg::CastTrait castTraitHalf2Float = {
		557	+ Reg::RegLayout::ZERO, Reg::SatMode::UNKNOWN, Reg::MaskMergeMode::ZEROING, RoundMode::UNKNOWN};
		558	+
		559	+ Reg::Duplicate(absMaskBF16Reg, ABS_FOR_UINT16);
		560	+ Reg::Duplicate(zeroRegTensor32, 0);
		561	+ Reg::Duplicate(manMaskFP32Reg, MAN_FOR_FP32);
		562	+ Reg::Duplicate(fp8NanRegTensor32, MAX_EXP_FOR_FP8_IN_FP32);
		563	+ Reg::Duplicate(biasRegTensor32, FP32_EXP_BIAS_CUBLAS);
		564	+ Reg::Duplicate(nanRegTensor32, NAN_CUSTOMIZATION_PACK);
		565	+ Reg::Duplicate(fullZeroReg, 0);
		566	+
		567	+ uint32_t pnum = vfNum32;
		568	+ p0 = Reg::UpdateMask<T>(pnum);
		569	+ for (uint16_t i = 0; i < regLoop; i++) {
		570	+ Reg::Duplicate(expMaxreShareExpFP16YOneRegTensor, 0);
		571	+ for (uint16_t j = 0; j <= static_cast<uint16_t>(blockCountLoop); j++) {
		572	+ Reg::LoadAlign(xRegTensor0, xAddr + j * dataLen + i * vfNum32);
		573	+ Reg::LoadAlign(xRegTensor1, xAddr + (static_cast<uint16_t>(blockCount) - j - 1) * dataLen + i * vfNum32);
		574	+ Reg::And(expMaxRegTensor0, (Reg::RegTensor<uint16_t>&)xRegTensor0, absMaskBF16Reg, p0);
		575	+ Reg::And(expMaxRegTensor1, (Reg::RegTensor<uint16_t>&)xRegTensor1, absMaskBF16Reg, p0);
		576	+ Reg::Max(expMaxRegTensor, expMaxRegTensor0, expMaxRegTensor1, p0);
		577	+ Reg::Max(expMaxreShareExpFP16YOneRegTensor, expMaxRegTensor, expMaxreShareExpFP16YOneRegTensor, p0);
		578	+ }
		579	+ // 2. BF16→FP32
		580	+ Reg::Interleave(expMaxreShareExpFP16YOneRegTensor, fullZeroReg,
		581	+ expMaxreShareExpFP16YOneRegTensor, fullZeroReg);
		582	+ Reg::Cast<float, T, castTraitHalf2Float>(
		583	+ (Reg::RegTensor<float>&)expMaxAndAddOneFP32RegTensor,
		584	+ (Reg::RegTensor<T>&)expMaxreShareExpFP16YOneRegTensor, preMaskScale);
		585	+ // 3. Inf/Zero校验
		586	+ Reg::CompareScalar<uint32_t, CMPMODE::LT>(
		587	+ infMask, expMaxAndAddOneFP32RegTensor, MAX_EXP_FOR_FP32, preMaskScale);
		588	+ Reg::Compare<uint32_t, CMPMODE::NE>(
		589	+ zeroMask, expMaxAndAddOneFP32RegTensor, zeroRegTensor32, preMaskScale);
		590	+ // 4. S_fp32 = Amax / Amax(DType)
		591	+ Reg::Muls((Reg::RegTensor<float>&)expMaxAndAddOneFP32RegTensor,
		592	+ (Reg::RegTensor<float>&)expMaxAndAddOneFP32RegTensor, inv_dtype_max, preMaskScale);
		593	+ // 5. 提取 Exp 和 Man
		594	+ Reg::ShiftRights(expAndreShareExpFP32RegTensor,
		595	+ expMaxAndAddOneFP32RegTensor, SHR_NUM_FOR_FP32, preMaskScale);
		596	+ Reg::And(manAndmxScaleFP32RegTensor,
		597	+ expMaxAndAddOneFP32RegTensor, manMaskFP32Reg, preMaskScale);
		598	+ // 6. 条件向上取整
		599	+ Reg::CompareScalar<uint32_t, CMPMODE::GT>(
		600	+ p1, expAndreShareExpFP32RegTensor, NUMBER_ZERO, preMaskScale);
		601	+ Reg::CompareScalar<uint32_t, CMPMODE::LT>(
		602	+ p2, expAndreShareExpFP32RegTensor, NUMBER_TWO_FIVE_FOUR, preMaskScale);
		603	+ Reg::CompareScalar<uint32_t, CMPMODE::GT>(
		604	+ p3, manAndmxScaleFP32RegTensor, NUMBER_ZERO, preMaskScale);
		605	+ Reg::MaskAnd(p1, p1, p2, preMaskScale);
		606	+ Reg::MaskAnd(p1, p1, p3, preMaskScale);
		607	+ Reg::CompareScalar<uint32_t, CMPMODE::EQ>(
		608	+ p2, expAndreShareExpFP32RegTensor, NUMBER_ZERO, preMaskScale);
		609	+ Reg::CompareScalar<uint32_t, CMPMODE::GT>(
		610	+ p3, manAndmxScaleFP32RegTensor, NUMBER_HALF, preMaskScale);
		611	+ Reg::MaskAnd(p2, p2, p3, preMaskScale);
		612	+ Reg::MaskXor(p1, p1, p2, preMaskScale);
		613	+ Reg::Adds(expMaxAndAddOneFP32RegTensor, expAndreShareExpFP32RegTensor, 1, preMaskScale);
		614	+ Reg::Select(manAndmxScaleFP32RegTensor,
		615	+ expMaxAndAddOneFP32RegTensor, expAndreShareExpFP32RegTensor, p1);
		616	+ // 7. 异常值处理
		617	+ Reg::Select<uint32_t>(manAndmxScaleFP32RegTensor,
		618	+ manAndmxScaleFP32RegTensor, fp8NanRegTensor32, infMask);
		619	+ Reg::Select<uint32_t>(manAndmxScaleFP32RegTensor,
		620	+ manAndmxScaleFP32RegTensor, zeroRegTensor32, zeroMask);
		621	+ // 8. 输出 mxscale（E8M0）
		622	+ Reg::Pack<uint16_t, uint32_t, Reg::HighLowPart::LOWEST>(
		623	+ expScaleMxYZeroRegTensor, manAndmxScaleFP32RegTensor);
		624	+ Reg::Pack<uint8_t, uint16_t, Reg::HighLowPart::LOWEST>(
		625	+ outZeromxScaleFp8, expScaleMxYZeroRegTensor);
		626	+ Reg::StoreUnAlign(mxScaleAddr, outZeromxScaleFp8, u1, vfNum32);
		627	+ Reg::StoreUnAlignPost(mxScaleAddr, u1, 0);
		628	+ // 9. 求 1/scale
		629	+ Reg::ShiftLefts(manAndmxScaleFP32RegTensor,
		630	+ manAndmxScaleFP32RegTensor, SHR_NUM_FOR_BF16, preMaskScale);
		631	+ Reg::Sub(expAndreShareExpFP32RegTensor,
		632	+ biasRegTensor32, manAndmxScaleFP32RegTensor, preMaskScale);
		633	+ Reg::Select<uint32_t>(expAndreShareExpFP32RegTensor,
		634	+ expAndreShareExpFP32RegTensor, nanRegTensor32, infMask);
		635	+ Reg::Select<uint32_t>(expAndreShareExpFP32RegTensor,
		636	+ expAndreShareExpFP32RegTensor, zeroRegTensor32, zeroMask);
		637	+ Reg::Pack<uint16_t, uint32_t, Reg::HighLowPart::LOWEST>(
		638	+ expMaxreShareExpFP16YOneRegTensor, expAndreShareExpFP32RegTensor);
		639	+ // 10. 量化
		640	+ if constexpr (IsSame<T, half>::value) {
		641	+ Reg::Cast<float, bfloat16_t, castTraitZero>(reversedShareExpRegTensorFP32Zero,
		642	+ (Reg::RegTensor<bfloat16_t>&)expMaxreShareExpFP16YOneRegTensor, p0);
		643	+ Reg::Cast<float, bfloat16_t, castTraitOne>(reversedShareExpRegTensorFP32One,
		644	+ (Reg::RegTensor<bfloat16_t>&)expMaxreShareExpFP16YOneRegTensor, p0);
		645	+ }
		646	+ for (uint16_t j = 0; j < static_cast<uint16_t>(blockCount); j++) {
		647	+ Reg::LoadAlign(xRegTensor, xAddr + j * dataLen + i * vfNum32);
		648	+ if constexpr (IsSame<T, half>::value) {
		649	+ Reg::Cast<float, T, castTraitZero>(yZero, xRegTensor, p0);
		650	+ Reg::Cast<float, T, castTraitOne>(yOne, xRegTensor, p0);
		651	+ Reg::Mul(yZero, yZero, reversedShareExpRegTensorFP32Zero, maskAll);
		652	+ Reg::Mul(yOne, yOne, reversedShareExpRegTensorFP32One, maskAll);
		653	+ } else {
		654	+ Reg::Mul(valueRegTensor, xRegTensor,
		655	+ (Reg::RegTensor<bfloat16_t>&)expMaxreShareExpFP16YOneRegTensor, p0);
		656	+ Reg::Cast<float, bfloat16_t, castTraitZero>(yZero, valueRegTensor, maskAll);
		657	+ Reg::Cast<float, bfloat16_t, castTraitOne>(yOne, valueRegTensor, maskAll);
		658	+ }
		659	+ Reg::Interleave(yZero, yOne, yZero, yOne);
		660	+ Reg::Cast<U, float, castTrait32to8>(yZeroFP8, yZero, maskAll);
		661	+ Reg::Pack(expScaleMxYZeroRegTensor, (Reg::RegTensor<uint32_t>&)yZeroFP8);
		662	+ Reg::Pack(outZeromxScaleFp8, expScaleMxYZeroRegTensor);
		663	+ auto addr = yAddr + (j * outDataLenAlign + i * vfNum32);
		664	+ Reg::StoreUnAlign(addr, outZeromxScaleFp8, u1, singleLoopNum);
		665	+ Reg::StoreUnAlignPost(addr, u1, 0);
		666	+ }
		667	+ }
		668	+
		669	+ if (tailVfLen != 0) {
		670	+ uint32_t tailPnum = tailVfLen;
		671	+ p0 = Reg::UpdateMask<T>(tailPnum);
		672	+ Reg::Duplicate(expMaxreShareExpFP16YOneRegTensor, 0);
		673	+ for (uint16_t k = 0; k <= static_cast<uint16_t>(blockCountLoop); k++) {
		674	+ Reg::LoadAlign(xRegTensor0, xAddr + k * dataLen + regLoop * vfNum32);
		675	+ Reg::LoadAlign(xRegTensor1, xAddr + (static_cast<uint16_t>(blockCount) - k - 1) * dataLen + regLoop * vfNum32);
		676	+ Reg::And(expMaxRegTensor0, (Reg::RegTensor<uint16_t>&)xRegTensor0, absMaskBF16Reg, p0);
		677	+ Reg::And(expMaxRegTensor1, (Reg::RegTensor<uint16_t>&)xRegTensor1, absMaskBF16Reg, p0);
		678	+ Reg::Max(expMaxRegTensor, expMaxRegTensor0, expMaxRegTensor1, p0);
		679	+ Reg::Max(expMaxreShareExpFP16YOneRegTensor, expMaxRegTensor, expMaxreShareExpFP16YOneRegTensor, p0);
		680	+ }
		681	+ // 2. BF16→FP32
		682	+ Reg::Interleave(expMaxreShareExpFP16YOneRegTensor, fullZeroReg,
		683	+ expMaxreShareExpFP16YOneRegTensor, fullZeroReg);
		684	+ Reg::Cast<float, T, castTraitHalf2Float>(
		685	+ (Reg::RegTensor<float>&)expMaxAndAddOneFP32RegTensor,
		686	+ (Reg::RegTensor<T>&)expMaxreShareExpFP16YOneRegTensor, preMaskScale);
		687	+ // 3. Inf/Zero校验
		688	+ Reg::CompareScalar<uint32_t, CMPMODE::LT>(
		689	+ infMask, expMaxAndAddOneFP32RegTensor, MAX_EXP_FOR_FP32, preMaskScale);
		690	+ Reg::Compare<uint32_t, CMPMODE::NE>(
		691	+ zeroMask, expMaxAndAddOneFP32RegTensor, zeroRegTensor32, preMaskScale);
		692	+ // 4. S_fp32 = Amax / Amax(DType)
		693	+ Reg::Muls((Reg::RegTensor<float>&)expMaxAndAddOneFP32RegTensor,
		694	+ (Reg::RegTensor<float>&)expMaxAndAddOneFP32RegTensor, inv_dtype_max, preMaskScale);
		695	+ // 5. 提取 Exp 和 Man
		696	+ Reg::ShiftRights(expAndreShareExpFP32RegTensor,
		697	+ expMaxAndAddOneFP32RegTensor, SHR_NUM_FOR_FP32, preMaskScale);
		698	+ Reg::And(manAndmxScaleFP32RegTensor,
		699	+ expMaxAndAddOneFP32RegTensor, manMaskFP32Reg, preMaskScale);
		700	+ // 6. 条件向上取整
		701	+ Reg::CompareScalar<uint32_t, CMPMODE::GT>(
		702	+ p1, expAndreShareExpFP32RegTensor, NUMBER_ZERO, preMaskScale);
		703	+ Reg::CompareScalar<uint32_t, CMPMODE::LT>(
		704	+ p2, expAndreShareExpFP32RegTensor, NUMBER_TWO_FIVE_FOUR, preMaskScale);
		705	+ Reg::CompareScalar<uint32_t, CMPMODE::GT>(
		706	+ p3, manAndmxScaleFP32RegTensor, NUMBER_ZERO, preMaskScale);
		707	+ Reg::MaskAnd(p1, p1, p2, preMaskScale);
		708	+ Reg::MaskAnd(p1, p1, p3, preMaskScale);
		709	+ Reg::CompareScalar<uint32_t, CMPMODE::EQ>(
		710	+ p2, expAndreShareExpFP32RegTensor, NUMBER_ZERO, preMaskScale);
		711	+ Reg::CompareScalar<uint32_t, CMPMODE::GT>(
		712	+ p3, manAndmxScaleFP32RegTensor, NUMBER_HALF, preMaskScale);
		713	+ Reg::MaskAnd(p2, p2, p3, preMaskScale);
		714	+ Reg::MaskXor(p1, p1, p2, preMaskScale);
		715	+ Reg::Adds(expMaxAndAddOneFP32RegTensor, expAndreShareExpFP32RegTensor, 1, preMaskScale);
		716	+ Reg::Select(manAndmxScaleFP32RegTensor,
		717	+ expMaxAndAddOneFP32RegTensor, expAndreShareExpFP32RegTensor, p1);
		718	+ // 7. 异常值处理
		719	+ Reg::Select<uint32_t>(manAndmxScaleFP32RegTensor,
		720	+ manAndmxScaleFP32RegTensor, fp8NanRegTensor32, infMask);
		721	+ Reg::Select<uint32_t>(manAndmxScaleFP32RegTensor,
		722	+ manAndmxScaleFP32RegTensor, zeroRegTensor32, zeroMask);
		723	+ // 8. 输出 mxscale（E8M0）
		724	+ Reg::Pack<uint16_t, uint32_t, Reg::HighLowPart::LOWEST>(
		725	+ expScaleMxYZeroRegTensor, manAndmxScaleFP32RegTensor);
		726	+ Reg::Pack<uint8_t, uint16_t, Reg::HighLowPart::LOWEST>(
		727	+ outZeromxScaleFp8, expScaleMxYZeroRegTensor);
		728	+ Reg::StoreUnAlign(mxScaleAddr, outZeromxScaleFp8, u1, tailVfLen);
		729	+ Reg::StoreUnAlignPost(mxScaleAddr, u1, 0);
		730	+ // 9. 求 1/scale
		731	+ Reg::ShiftLefts(manAndmxScaleFP32RegTensor,
		732	+ manAndmxScaleFP32RegTensor, SHR_NUM_FOR_BF16, preMaskScale);
		733	+ Reg::Sub(expAndreShareExpFP32RegTensor,
		734	+ biasRegTensor32, manAndmxScaleFP32RegTensor, preMaskScale);
		735	+ Reg::Select<uint32_t>(expAndreShareExpFP32RegTensor,
		736	+ expAndreShareExpFP32RegTensor, nanRegTensor32, infMask);
		737	+ Reg::Select<uint32_t>(expAndreShareExpFP32RegTensor,
		738	+ expAndreShareExpFP32RegTensor, zeroRegTensor32, zeroMask);
		739	+ Reg::Pack<uint16_t, uint32_t, Reg::HighLowPart::LOWEST>(
		740	+ expMaxreShareExpFP16YOneRegTensor, expAndreShareExpFP32RegTensor);
		741	+ // 10. 量化
		742	+ if constexpr (IsSame<T, half>::value) {
		743	+ Reg::Cast<float, bfloat16_t, castTraitZero>(reversedShareExpRegTensorFP32Zero,
		744	+ (Reg::RegTensor<bfloat16_t>&)expMaxreShareExpFP16YOneRegTensor, p0);
		745	+ Reg::Cast<float, bfloat16_t, castTraitOne>(reversedShareExpRegTensorFP32One,
		746	+ (Reg::RegTensor<bfloat16_t>&)expMaxreShareExpFP16YOneRegTensor, p0);
		747	+ }
		748	+ for (uint16_t j = 0; j < static_cast<uint16_t>(blockCount); j++) {
		749	+ Reg::LoadAlign(xRegTensor, xAddr + regLoop * vfNum32 + j * dataLen);
		750	+ if constexpr (IsSame<T, half>::value) {
		751	+ Reg::Cast<float, T, castTraitZero>(yZero, xRegTensor, p0);
		752	+ Reg::Cast<float, T, castTraitOne>(yOne, xRegTensor, p0);
		753	+ Reg::Mul(yZero, yZero, reversedShareExpRegTensorFP32Zero, maskAll);
		754	+ Reg::Mul(yOne, yOne, reversedShareExpRegTensorFP32One, maskAll);
		755	+ } else {
		756	+ Reg::Mul(valueRegTensor, xRegTensor,
		757	+ (Reg::RegTensor<bfloat16_t>&)expMaxreShareExpFP16YOneRegTensor, p0);
		758	+ Reg::Cast<float, bfloat16_t, castTraitZero>(yZero, valueRegTensor, maskAll);
		759	+ Reg::Cast<float, bfloat16_t, castTraitOne>(yOne, valueRegTensor, maskAll);
		760	+ }
		761	+ Reg::Interleave(yZero, yOne, yZero, yOne);
		762	+ Reg::Cast<U, float, castTrait32to8>(yZeroFP8, yZero, maskAll);
		763	+ Reg::Pack(expScaleMxYZeroRegTensor, (Reg::RegTensor<uint32_t>&)yZeroFP8);
		764	+ Reg::Pack(outZeromxScaleFp8, expScaleMxYZeroRegTensor);
		765	+ auto addr = yAddr + (regLoop * vfNum32 + j * outDataLenAlign);
		766	+ Reg::StoreUnAlign(addr, outZeromxScaleFp8, u1, singleTailLoopNum);
		767	+ Reg::StoreUnAlignPost(addr, u1, 0);
456	}	768	}
457	}	769	}
458	}	770	}

Aquant/grouped_dynamic_mx_quant/op_kernel/arch35/grouped_dynamic_mx_quant_struct.h+34-0

@@ -0,0 +1,34 @@
	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 grouped_dynamic_mx_quant_struct.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef _GROUPED_DYNAMIC_MX_QUANT_STRUCT_H_
	17	+#define _GROUPED_DYNAMIC_MX_QUANT_STRUCT_H_
	18	+
	19	+#include "ascendc/host_api/tiling/template_argument.h"
	20	+
	21	+#define TPL_MODE_0 0
	22	+#define TPL_MODE_1 1
	23	+
	24	+namespace GroupedDynamicMxQuantOp {
	25	+ASCENDC_TPL_ARGS_DECL(
	26	+ GroupedDynamicMxQuant,
	27	+ ASCENDC_TPL_UINT_DECL(mode, 1, ASCENDC_TPL_UI_LIST, TPL_MODE_0, TPL_MODE_1));
	28	+
	29	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
	30	+ ASCENDC_TPL_UINT_SEL(mode, ASCENDC_TPL_UI_LIST, TPL_MODE_0, TPL_MODE_1)));
	31	+
	32	+} // namespace GroupedDynamicMxQuantOp
	33	+
	34	+#endif // _GROUPED_DYNAMIC_MX_QUANT_STRUCT_H_

Aquant/grouped_dynamic_mx_quant/op_kernel/arch35/grouped_dynamic_mx_quant_tilingdata.h+31-0

@@ -0,0 +1,31 @@
	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	+#ifndef GROUPED_DYNAMIC_MX_QUANT_TILINGDATA_H
	12	+#define GROUPED_DYNAMIC_MX_QUANT_TILINGDATA_H
	13	+
	14	+#include <cstdint>
	15	+
	16	+struct GroupedDynamicMxQuantTilingData {
	17	+ int64_t totalCoreNum;
	18	+ int64_t usedCoreNum; // 实际使用核数
	19	+ int64_t blockFactor; // 头核处理数据量
	20	+ int64_t tailBlockFactor; // 尾核处理数据量
	21	+ int64_t uo; // 切分轴上的循环次数
	22	+ int64_t maxUbCol; // 单次循环要处理的数据大小
	23	+ int64_t ubFactor; // 单次循环要处理的数据大小
	24	+ int64_t tailUbFactor; // 尾循环要处理的数据大小
	25	+ int64_t blockSize; // 量化数据块大小
	26	+ int64_t scaleAlg; // OCP Microscaling Formats(Mx) Specification实现/cuBLAS实现，默认OCP实现
	27	+ int64_t preAxisSize; // 输入row长度
	28	+ int64_t postAxisSize; // 输入col长度
	29	+ float dstTypeMax; // 目前仅支持0.0
	30	+};
	31	+#endif // GROUPED_DYNAMIC_MX_QUANT_TILINGDATA_H

Mquant/grouped_dynamic_mx_quant/op_kernel/grouped_dynamic_mx_quant.cpp+12-30

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -14,48 +14,30 @@
14	*/	14	*/
15		15
16	#include "arch35/grouped_dynamic_mx_quant_not_tail_axis_fp8.h"	16	#include "arch35/grouped_dynamic_mx_quant_not_tail_axis_fp8.h"
17	-	17	+#include "arch35/grouped_dynamic_mx_quant_tilingdata.h"
18	-#define TILING_KEY_FP16_FP8E4M3FN_QUANT_OTHER_AXIS 11
19	-#define TILING_KEY_BF16_FP8E4M3FN_QUANT_OTHER_AXIS 21
20	-#define TILING_KEY_FP16_FP8E5M2_QUANT_OTHER_AXIS 12
21	-#define TILING_KEY_BF16_FP8E5M2_QUANT_OTHER_AXIS 22
22		18
23	#define FLOAT_OVERFLOW_MODE_CTRL 60	19	#define FLOAT_OVERFLOW_MODE_CTRL 60
24		20
25	-// 十位数为1、2，分别表示输入类型是float16、bfloat16;
26	-// 个位数为1、2，分别表示输出类型是float8_e4m3fn、float8_e5m2
27	-
28	using namespace GroupedDynamicMxQuant;	21	using namespace GroupedDynamicMxQuant;
		22	+using namespace GroupedDynamicMxQuantOp;
29		23
30	-extern "C" __global__ __aicore__ void grouped_dynamic_mx_quant(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR mxScale,	24	+template <uint64_t mode>
		25	+__global__ __aicore__ void grouped_dynamic_mx_quant(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR mxScale,
31	GM_ADDR workspace, GM_ADDR tiling)	26	GM_ADDR workspace, GM_ADDR tiling)
32	{	27	{
33	- GET_TILING_DATA(tilingData, tiling);	28	+ REGISTER_TILING_DEFAULT(GroupedDynamicMxQuantTilingData);
		29	+ GET_TILING_DATA_WITH_STRUCT(GroupedDynamicMxQuantTilingData, tilingData, tiling);
34	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);	30	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
35		31
36	#if (__NPU_ARCH__ == 3510)	32	#if (__NPU_ARCH__ == 3510)
37	- int64_t oriOverflowMode = AscendC::GetCtrlSpr<FLOAT_OVERFLOW_MODE_CTRL,FLOAT_OVERFLOW_MODE_CTRL>();	33	+ int64_t oriOverflowMode = AscendC::GetCtrlSpr<FLOAT_OVERFLOW_MODE_CTRL, FLOAT_OVERFLOW_MODE_CTRL>();
38	#endif	34	#endif
39		35
40	- if (TILING_KEY_IS(TILING_KEY_FP16_FP8E4M3FN_QUANT_OTHER_AXIS)) {	36	+ GroupedDynamicMxQuant::GroupedDynamicMxQuantBaseFP8<DTYPE_X, DTYPE_Y> op;
41	- GroupedDynamicMxQuant::GroupedDynamicMxQuantBaseFP8<half, fp8_e4m3fn_t> op;	37	+ op.Init(x, groupIndex, y, mxScale, tilingData);
42	- op.Init(x, groupIndex, y, mxScale, tilingData);	38	+ op.Process();
43	- op.Process();
44	- } else if (TILING_KEY_IS(TILING_KEY_BF16_FP8E4M3FN_QUANT_OTHER_AXIS)) {
45	- GroupedDynamicMxQuant::GroupedDynamicMxQuantBaseFP8<bfloat16_t, fp8_e4m3fn_t> op;
46	- op.Init(x, groupIndex, y, mxScale, tilingData);
47	- op.Process();
48	- } else if (TILING_KEY_IS(TILING_KEY_FP16_FP8E5M2_QUANT_OTHER_AXIS)) {
49	- GroupedDynamicMxQuant::GroupedDynamicMxQuantBaseFP8<half, fp8_e5m2_t> op;
50	- op.Init(x, groupIndex, y, mxScale, tilingData);
51	- op.Process();
52	- } else if (TILING_KEY_IS(TILING_KEY_BF16_FP8E5M2_QUANT_OTHER_AXIS)) {
53	- GroupedDynamicMxQuant::GroupedDynamicMxQuantBaseFP8<bfloat16_t, fp8_e5m2_t> op;
54	- op.Init(x, groupIndex, y, mxScale, tilingData);
55	- op.Process();
56	- }
57		39
58	#if (__NPU_ARCH__ == 3510)	40	#if (__NPU_ARCH__ == 3510)
59	- AscendC::SetCtrlSpr<FLOAT_OVERFLOW_MODE_CTRL,FLOAT_OVERFLOW_MODE_CTRL>(oriOverflowMode);	41	+ AscendC::SetCtrlSpr<FLOAT_OVERFLOW_MODE_CTRL, FLOAT_OVERFLOW_MODE_CTRL>(oriOverflowMode);
60	#endif	42	#endif
61	}	43	}

Mquant/grouped_dynamic_mx_quant/tests/CMakeLists.txt+1-1

	@@ -1,5 +1,5 @@
1	# -----------------------------------------------------------------------------------------------------------	1	# -----------------------------------------------------------------------------------------------------------
2	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.

Mquant/grouped_dynamic_mx_quant/tests/ut/CMakeLists.txt+1-1

	@@ -1,5 +1,5 @@
1	# -----------------------------------------------------------------------------------------------------------	1	# -----------------------------------------------------------------------------------------------------------
2	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.

Mquant/grouped_dynamic_mx_quant/tests/ut/op_api/CMakeLists.txt+1-1

	@@ -1,5 +1,5 @@
1	# -----------------------------------------------------------------------------------------------------------	1	# -----------------------------------------------------------------------------------------------------------
2	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.

Mquant/grouped_dynamic_mx_quant/tests/ut/op_api/test_aclnn_grouped_dynamic_mx_quant.cpp+1-1

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.

Mquant/grouped_dynamic_mx_quant/tests/ut/op_host/CMakeLists.txt+1-1

	@@ -1,5 +1,5 @@
1	# -----------------------------------------------------------------------------------------------------------	1	# -----------------------------------------------------------------------------------------------------------
2	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.

Mquant/grouped_dynamic_mx_quant/tests/ut/op_host/arch35/test_grouped_dynamic_mx_quant_tiling.cpp+105-42

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -23,9 +23,9 @@
23	#include "exe_graph/runtime/storage_shape.h"	23	#include "exe_graph/runtime/storage_shape.h"
24	#include "platform/platform_infos_def.h"	24	#include "platform/platform_infos_def.h"
25	#include "ut_op_util.h"	25	#include "ut_op_util.h"
26	-#include "../../../../op_host/arch35/grouped_dynamic_mx_quant_tiling_arch35.h"	26	+#include "quant/grouped_dynamic_mx_quant/op_host/arch35/grouped_dynamic_mx_quant_tiling_arch35.h"
27	#include "any_value.h"	27	#include "any_value.h"
28	-	28	+#include "register/op_impl_registry.h"
29		29
30	using namespace std;	30	using namespace std;
31		31
	@@ -40,21 +40,22 @@ protected:
40	}	40	}
41	};	41	};
42		42
43	-static string TilingData2Str(const gert::TilingData* tilingData)	43	+template <typename T>
		44	+static string to_string(void* buf, size_t size)
44	{	45	{
45	- auto data = tilingData->GetData();	46	+ std::string result;
46	- string result;	47	+ const T* data = reinterpret_cast<const T*>(buf);
47	- for (size_t i = 0; i < tilingData->GetDataSize(); i += sizeof(int64_t)) {	48	+ size_t len = size / sizeof(T);
48	- result += std::to_string((reinterpret_cast<const int64_t*>(tilingData->GetData())[i / sizeof(int64_t)]));	49	+ for (size_t i = 0; i < len; i++) {
		50	+ result += std::to_string(data[i]);
49	result += " ";	51	result += " ";
50	}	52	}
51	-
52	return result;	53	return result;
53	}	54	}
54		55
55	-static void ExecuteTestCase(ge::DataType inDtype, ge::DataType outDtype, ge::DataType in2Dtype, ge::DataType out2Dtype, gert::StorageShape shape, gert::StorageShape groupIdxShape,	56	+static void ExecuteTestCase(ge::DataType inDtype, ge::DataType outDtype, ge::DataType in2Dtype, ge::DataType out2Dtype,
56	- gert::StorageShape outShape, gert::StorageShape scaleShape, int64_t blockSize, string expectTilingData,	57	+ gert::StorageShape shape, gert::StorageShape groupIdxShape, gert::StorageShape outShape, gert::StorageShape scaleShape,
57	- ge::graphStatus status = ge::GRAPH_SUCCESS) {	58	+ int64_t blockSize, int64_t scaleAlg, float dstTypeMax, string expectTilingData, ge::graphStatus status = ge::GRAPH_SUCCESS) {
58	string compile_info_string = R"({	59	string compile_info_string = R"({
59	"hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",	60	"hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
60	"Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,	61	"Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
	@@ -112,7 +113,9 @@ static void ExecuteTestCase(ge::DataType inDtype, ge::DataType outDtype, ge::Dat
112	.NodeOutputTd(1, out2Dtype, ge::FORMAT_ND, ge::FORMAT_ND)	113	.NodeOutputTd(1, out2Dtype, ge::FORMAT_ND, ge::FORMAT_ND)
113	.NodeAttrs({{"round_mode", Ops::NN::AnyValue::CreateFrom<string>("rint")},	114	.NodeAttrs({{"round_mode", Ops::NN::AnyValue::CreateFrom<string>("rint")},
114	{"dst_type", Ops::NN::AnyValue::CreateFrom((int64_t)outDtype)},	115	{"dst_type", Ops::NN::AnyValue::CreateFrom((int64_t)outDtype)},
115	- {"blocksize", Ops::NN::AnyValue::CreateFrom(blockSize)}})	116	+ {"blocksize", Ops::NN::AnyValue::CreateFrom(blockSize)},
		117	+ {"scale_alg", Ops::NN::AnyValue::CreateFrom(scaleAlg)},
		118	+ {"dst_type_max", Ops::NN::AnyValue::CreateFrom(dstTypeMax)}})
116	.TilingData(param.get())	119	.TilingData(param.get())
117	.Workspace(ws_size)	120	.Workspace(ws_size)
118	.Build();	121	.Build();
	@@ -133,7 +136,9 @@ static void ExecuteTestCase(ge::DataType inDtype, ge::DataType outDtype, ge::Dat
133	// todo check tiling result	136	// todo check tiling result
134	auto tiling_key = tiling_context->GetTilingKey();	137	auto tiling_key = tiling_context->GetTilingKey();
135	auto block_dim = tiling_context->GetBlockDim();	138	auto block_dim = tiling_context->GetBlockDim();
136	- auto tiling_data_result = TilingData2Str(tiling_context->GetRawTilingData());	139	+
		140	+ auto raw_tiling_data = tiling_context->GetRawTilingData();
		141	+ auto tiling_data_result = to_string<int64_t>(raw_tiling_data->GetData(), raw_tiling_data->GetDataSize());
137	EXPECT_EQ(tiling_data_result, expectTilingData);	142	EXPECT_EQ(tiling_data_result, expectTilingData);
138	}	143	}
139		144
	@@ -142,9 +147,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_bfloat1
142	gert::StorageShape groupIdxShape = {{1}, {1}};	147	gert::StorageShape groupIdxShape = {{1}, {1}};
143	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	148	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
144	int64_t blockSize = 32;	149	int64_t blockSize = 32;
145	- string expectTilingData = "64 1 1 1 1 128 128 128 32 32 128 ";	150	+ int64_t scaleAlg = 0;
		151	+ float dstTypeMax = 0.0;
		152	+ string expectTilingData = "64 1 1 1 1 128 128 128 32 0 32 128 0 ";
146		153
147	- ExecuteTestCase(ge::DT_BF16, ge::DT_FLOAT8_E4M3FN, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData);	154	+ ExecuteTestCase(ge::DT_BF16, ge::DT_FLOAT8_E4M3FN, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
		155	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData);
148	}	156	}
149		157
150	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_float16_fp8e4m3fn) {	158	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_float16_fp8e4m3fn) {
	@@ -152,9 +160,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_float16
152	gert::StorageShape groupIdxShape = {{1}, {1}};	160	gert::StorageShape groupIdxShape = {{1}, {1}};
153	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	161	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
154	int64_t blockSize = 32;	162	int64_t blockSize = 32;
155	- string expectTilingData = "64 1 1 1 1 128 128 128 32 32 128 ";	163	+ int64_t scaleAlg = 0;
		164	+ float dstTypeMax = 0.0;
		165	+ string expectTilingData = "64 1 1 1 1 128 128 128 32 0 32 128 0 ";
156		166
157	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E4M3FN, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData);	167	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E4M3FN, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
		168	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData);
158	}	169	}
159		170
160	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_bfloat16_fp8e5m2) {	171	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_bfloat16_fp8e5m2) {
	@@ -162,9 +173,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_bfloat1
162	gert::StorageShape groupIdxShape = {{1}, {1}};	173	gert::StorageShape groupIdxShape = {{1}, {1}};
163	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	174	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
164	int64_t blockSize = 32;	175	int64_t blockSize = 32;
165	- string expectTilingData = "64 1 1 1 1 128 128 128 32 32 128 ";	176	+ int64_t scaleAlg = 0;
		177	+ float dstTypeMax = 0.0;
		178	+ string expectTilingData = "64 1 1 1 1 128 128 128 32 0 32 128 0 ";
166		179
167	- ExecuteTestCase(ge::DT_BF16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData);	180	+ ExecuteTestCase(ge::DT_BF16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
		181	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData);
168	}	182	}
169		183
170	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_float16_fp8e5m2) {	184	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_float16_fp8e5m2) {
	@@ -172,9 +186,25 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_float16
172	gert::StorageShape groupIdxShape = {{1}, {1}};	186	gert::StorageShape groupIdxShape = {{1}, {1}};
173	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	187	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
174	int64_t blockSize = 32;	188	int64_t blockSize = 32;
175	- string expectTilingData = "64 1 1 1 1 128 128 128 32 32 128 ";	189	+ int64_t scaleAlg = 0;
		190	+ float dstTypeMax = 0.0;
		191	+ string expectTilingData = "64 1 1 1 1 128 128 128 32 0 32 128 0 ";
176		192
177	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData);	193	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
		194	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData);
		195	+}
		196	+
		197	+TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_float16_fp8e5m2_scale_alg_1) {
		198	+ gert::StorageShape shape = {{32, 128}, {32, 128}};
		199	+ gert::StorageShape groupIdxShape = {{1}, {1}};
		200	+ gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
		201	+ int64_t blockSize = 32;
		202	+ int64_t scaleAlg = 1;
		203	+ float dstTypeMax = 0.0;
		204	+ string expectTilingData = "64 1 1 1 1 128 128 128 32 1 32 128 0 ";
		205	+
		206	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
		207	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData);
178	}	208	}
179		209
180	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_inDtype) {	210	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_inDtype) {
	@@ -182,10 +212,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_i
182	gert::StorageShape groupIdxShape = {{1}, {1}};	212	gert::StorageShape groupIdxShape = {{1}, {1}};
183	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	213	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
184	int64_t blockSize = 32;	214	int64_t blockSize = 32;
		215	+ int64_t scaleAlg = 0;
		216	+ float dstTypeMax = 0.0;
185	string expectTilingData = "";	217	string expectTilingData = "";
186		218
187	- ExecuteTestCase(ge::DT_FLOAT, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize,	219	+ ExecuteTestCase(ge::DT_FLOAT, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
188	- expectTilingData, ge::GRAPH_FAILED);	220	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
189	}	221	}
190		222
191	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_inDtype2) {	223	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_inDtype2) {
	@@ -193,10 +225,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_i
193	gert::StorageShape groupIdxShape = {{1}, {1}};	225	gert::StorageShape groupIdxShape = {{1}, {1}};
194	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	226	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
195	int64_t blockSize = 32;	227	int64_t blockSize = 32;
		228	+ int64_t scaleAlg = 0;
		229	+ float dstTypeMax = 0.0;
196	string expectTilingData = "";	230	string expectTilingData = "";
197		231
198	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT16, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize,	232	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT16, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
199	- expectTilingData, ge::GRAPH_FAILED);	233	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
200	}	234	}
201		235
202	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_outDtype) {	236	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_outDtype) {
	@@ -204,10 +238,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_o
204	gert::StorageShape groupIdxShape = {{1}, {1}};	238	gert::StorageShape groupIdxShape = {{1}, {1}};
205	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	239	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
206	int64_t blockSize = 32;	240	int64_t blockSize = 32;
		241	+ int64_t scaleAlg = 0;
		242	+ float dstTypeMax = 0.0;
207	string expectTilingData = "";	243	string expectTilingData = "";
208		244
209	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData,	245	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
210	- ge::GRAPH_FAILED);	246	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
211	}	247	}
212		248
213	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_outDtype2) {	249	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_outDtype2) {
	@@ -215,10 +251,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_o
215	gert::StorageShape groupIdxShape = {{1}, {1}};	251	gert::StorageShape groupIdxShape = {{1}, {1}};
216	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	252	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
217	int64_t blockSize = 32;	253	int64_t blockSize = 32;
		254	+ int64_t scaleAlg = 0;
		255	+ float dstTypeMax = 0.0;
218	string expectTilingData = "";	256	string expectTilingData = "";
219		257
220	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT16, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData,	258	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT16, shape, groupIdxShape, shape,
221	- ge::GRAPH_FAILED);	259	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
222	}	260	}
223		261
224	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_blockSize) {	262	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_blockSize) {
	@@ -226,10 +264,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_b
226	gert::StorageShape groupIdxShape = {{1}, {1}};	264	gert::StorageShape groupIdxShape = {{1}, {1}};
227	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	265	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
228	int64_t blockSize = 64;	266	int64_t blockSize = 64;
		267	+ int64_t scaleAlg = 0;
		268	+ float dstTypeMax = 0.0;
229	string expectTilingData = "";	269	string expectTilingData = "";
230		270
231	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData,	271	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
232	- ge::GRAPH_FAILED);	272	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
233	}	273	}
234		274
235	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_scale_shape) {	275	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_scale_shape) {
	@@ -237,10 +277,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_s
237	gert::StorageShape groupIdxShape = {{1}, {1}};	277	gert::StorageShape groupIdxShape = {{1}, {1}};
238	gert::StorageShape scaleShape = {{1, 128, 3}, {1, 128, 3}};	278	gert::StorageShape scaleShape = {{1, 128, 3}, {1, 128, 3}};
239	int64_t blockSize = 32;	279	int64_t blockSize = 32;
		280	+ int64_t scaleAlg = 0;
		281	+ float dstTypeMax = 0.0;
240	string expectTilingData = "";	282	string expectTilingData = "";
241		283
242	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData,	284	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
243	- ge::GRAPH_FAILED);	285	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData , ge::GRAPH_FAILED);
244	}	286	}
245		287
246		288
	@@ -249,10 +291,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_i
249	gert::StorageShape groupIdxShape = {{1}, {1}};	291	gert::StorageShape groupIdxShape = {{1}, {1}};
250	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	292	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
251	int64_t blockSize = 32;	293	int64_t blockSize = 32;
		294	+ int64_t scaleAlg = 0;
		295	+ float dstTypeMax = 0.0;
252	string expectTilingData = "";	296	string expectTilingData = "";
253		297
254	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData,	298	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
255	- ge::GRAPH_FAILED);	299	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
256	}	300	}
257		301
258	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_input1_dim) {	302	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_input1_dim) {
	@@ -260,10 +304,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_i
260	gert::StorageShape groupIdxShape = {{1, 1}, {1, 1}};	304	gert::StorageShape groupIdxShape = {{1, 1}, {1, 1}};
261	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	305	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
262	int64_t blockSize = 32;	306	int64_t blockSize = 32;
		307	+ int64_t scaleAlg = 0;
		308	+ float dstTypeMax = 0.0;
263	string expectTilingData = "";	309	string expectTilingData = "";
264		310
265	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData,	311	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
266	- ge::GRAPH_FAILED);	312	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
267	}	313	}
268		314
269	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_output0_dim) {	315	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_output0_dim) {
	@@ -272,10 +318,12 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_o
272	gert::StorageShape outShape = {{1, 32, 128}, {1, 32, 128}};	318	gert::StorageShape outShape = {{1, 32, 128}, {1, 32, 128}};
273	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};	319	gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
274	int64_t blockSize = 32;	320	int64_t blockSize = 32;
		321	+ int64_t scaleAlg = 0;
		322	+ float dstTypeMax = 0.0;
275	string expectTilingData = "";	323	string expectTilingData = "";
276		324
277	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, outShape, scaleShape, blockSize, expectTilingData,	325	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, outShape,
278	- ge::GRAPH_FAILED);	326	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
279	}	327	}
280		328
281	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_output1_dim) {	329	TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_output1_dim) {
	@@ -283,8 +331,23 @@ TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_o
283	gert::StorageShape groupIdxShape = {{1}, {1}};	331	gert::StorageShape groupIdxShape = {{1}, {1}};
284	gert::StorageShape scaleShape = {{1, 256}, {1, 256}};	332	gert::StorageShape scaleShape = {{1, 256}, {1, 256}};
285	int64_t blockSize = 32;	333	int64_t blockSize = 32;
		334	+ int64_t scaleAlg = 0;
		335	+ float dstTypeMax = 0.0;
286	string expectTilingData = "";	336	string expectTilingData = "";
287		337
288	- ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape, scaleShape, blockSize, expectTilingData,	338	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
289	- ge::GRAPH_FAILED);	339	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
		340	+}
		341	+
		342	+TEST_F(GroupedDynamicMxQuantTiling, GroupedDynamicMxQuant_tiling_ascendc_error_scale_alg) {
		343	+ gert::StorageShape shape = {{32, 128}, {32, 128}};
		344	+ gert::StorageShape groupIdxShape = {{1}, {1}};
		345	+ gert::StorageShape scaleShape = {{1, 128, 2}, {1, 128, 2}};
		346	+ int64_t blockSize = 32;
		347	+ int64_t scaleAlg = 2;
		348	+ float dstTypeMax = 0.0;
		349	+ string expectTilingData = "";
		350	+
		351	+ ExecuteTestCase(ge::DT_FLOAT16, ge::DT_FLOAT8_E5M2, ge::DT_INT32, ge::DT_FLOAT8_E8M0, shape, groupIdxShape, shape,
		352	+ scaleShape, blockSize, scaleAlg, dstTypeMax, expectTilingData, ge::GRAPH_FAILED);
290	}	353	}

Mquant/grouped_dynamic_mx_quant/tests/ut/op_host/test_grouped_dynamic_mx_quant_infershape.cpp+13-5

	@@ -1,5 +1,5 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	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.	5	* Please refer to the License for details. You may not use this file except in compliance with the License.
	@@ -41,14 +41,18 @@ protected:
41		41
42	TEST_F(GroupedDynamicMxQuant, GroupedDynamicMxQuant_infershape_case_1)	42	TEST_F(GroupedDynamicMxQuant, GroupedDynamicMxQuant_infershape_case_1)
43	{	43	{
44	- constexpr int32_t BLOCK_SIZE = 32;	44	+ constexpr int64_t BLOCK_SIZE = 32;
		45	+ constexpr int64_t SCALE_ALG = 0;
		46	+ constexpr float DST_TYPE_MAX = 0.0;
45	ge::op::GroupedDynamicMxQuant op;	47	ge::op::GroupedDynamicMxQuant op;
46	op.UpdateInputDesc("x", create_desc({32, 128}, ge::DT_BF16));	48	op.UpdateInputDesc("x", create_desc({32, 128}, ge::DT_BF16));
47	op.UpdateInputDesc("group_index", create_desc({1}, ge::DT_INT32));	49	op.UpdateInputDesc("group_index", create_desc({1}, ge::DT_INT32));
48	op.SetAttr("round_mode", "rint");	50	op.SetAttr("round_mode", "rint");
49	op.SetAttr("dst_type", (int64_t)ge::DT_FLOAT8_E4M3FN);	51	op.SetAttr("dst_type", (int64_t)ge::DT_FLOAT8_E4M3FN);
50	op.SetAttr("blocksize", BLOCK_SIZE);	52	op.SetAttr("blocksize", BLOCK_SIZE);
51	- Runtime2TestParam param{{"round_mode", "dst_type", "blocksize"}, {}, {}};	53	+ op.SetAttr("scale_alg", SCALE_ALG);
		54	+ op.SetAttr("dst_type_max", DST_TYPE_MAX);
		55	+ Runtime2TestParam param{{"round_mode", "dst_type", "blocksize", "scale_alg", "dst_type_max"}, {}, {}};
52	EXPECT_EQ(InferShapeTest(op, param), ge::GRAPH_SUCCESS);	56	EXPECT_EQ(InferShapeTest(op, param), ge::GRAPH_SUCCESS);
53	auto outputY = op.GetOutputDesc(0);	57	auto outputY = op.GetOutputDesc(0);
54	std::vector<int64_t> expectedYShape = {32, 128};	58	std::vector<int64_t> expectedYShape = {32, 128};
	@@ -60,13 +64,14 @@ TEST_F(GroupedDynamicMxQuant, GroupedDynamicMxQuant_InferDtype_case_1)
60	ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl("GroupedDynamicMxQuant"), nullptr);	64	ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl("GroupedDynamicMxQuant"), nullptr);
61	auto data_type_func = gert::OpImplRegistry::GetInstance().GetOpImpl("GroupedDynamicMxQuant")->infer_datatype;	65	auto data_type_func = gert::OpImplRegistry::GetInstance().GetOpImpl("GroupedDynamicMxQuant")->infer_datatype;
62		66
63	- constexpr int32_t BLOCK_SIZE = 32;
64	if (data_type_func != nullptr) {	67	if (data_type_func != nullptr) {
65	ge::DataType inDtype = ge::DT_BF16;	68	ge::DataType inDtype = ge::DT_BF16;
66	ge::DataType in2Dtype = ge::DT_INT32;	69	ge::DataType in2Dtype = ge::DT_INT32;
67	ge::DataType outDtype = ge::DT_FLOAT8_E4M3FN;	70	ge::DataType outDtype = ge::DT_FLOAT8_E4M3FN;
68	ge::DataType out2Dtype = ge::DT_FLOAT8_E8M0;	71	ge::DataType out2Dtype = ge::DT_FLOAT8_E8M0;
69	int64_t blockSize = 32;	72	int64_t blockSize = 32;
		73	+ int64_t scaleAlg = 0;
		74	+ float dstTypeMax = 0.0;
70	auto context_holder = gert::InferDataTypeContextFaker()	75	auto context_holder = gert::InferDataTypeContextFaker()
71	.IrInputNum(2)	76	.IrInputNum(2)
72	.NodeIoNum(2, 2)	77	.NodeIoNum(2, 2)
	@@ -76,7 +81,10 @@ TEST_F(GroupedDynamicMxQuant, GroupedDynamicMxQuant_InferDtype_case_1)
76	.NodeOutputTd(1, out2Dtype, ge::FORMAT_ND, ge::FORMAT_ND)	81	.NodeOutputTd(1, out2Dtype, ge::FORMAT_ND, ge::FORMAT_ND)
77	.NodeAttrs({{"round_mode", Ops::NN::AnyValue::CreateFrom<string>("rint")},	82	.NodeAttrs({{"round_mode", Ops::NN::AnyValue::CreateFrom<string>("rint")},
78	{"dst_type", Ops::NN::AnyValue::CreateFrom((int64_t)outDtype)},	83	{"dst_type", Ops::NN::AnyValue::CreateFrom((int64_t)outDtype)},
79	- {"blocksize", Ops::NN::AnyValue::CreateFrom(blockSize)}})	84	+ {"blocksize", Ops::NN::AnyValue::CreateFrom(blockSize)},
		85	+ {"scale_alg", Ops::NN::AnyValue::CreateFrom(scaleAlg)},
		86	+ {"dst_type_max", Ops::NN::AnyValue::CreateFrom(dstTypeMax)}})
		87	+
80	.InputDataTypes({&inDtype, &in2Dtype})	88	.InputDataTypes({&inDtype, &in2Dtype})
81	.OutputDataTypes({&outDtype, &out2Dtype})	89	.OutputDataTypes({&outDtype, &out2Dtype})
82	.Build();	90	.Build();

	@@ -37,7 +37,9 @@ extern "C" {
37	static constexpr int64_t X_DIM_NUM = 2;	37	static constexpr int64_t X_DIM_NUM = 2;
38	static constexpr int64_t NUM_TWO = 2;	38	static constexpr int64_t NUM_TWO = 2;
39	static constexpr int64_t SCALE_DIM_NUM = 3;	39	static constexpr int64_t SCALE_DIM_NUM = 3;
40	-static constexpr uint64_t NUM_ZERO = 0 ;	40	+static constexpr uint64_t NUM_ZERO = 0;
		41	+static constexpr int64_t DEFAULT_SCALE_ALG = 0;
		42	+static constexpr float DEFAULT_DST_TYPE_MAX = 0.0;
41		43
42	static const std::initializer_list<op::DataType> X_DTYPE_SUPPORT_LIST = {op::DataType::DT_FLOAT16, op::DataType::DT_BF16};	44	static const std::initializer_list<op::DataType> X_DTYPE_SUPPORT_LIST = {op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
43		45
	@@ -164,7 +166,7 @@ aclnnStatus aclnnGroupedDynamicMxQuantGetWorkspaceSize(const aclTensor* x, const
164	CHECK_RET(groupIndexContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);	166	CHECK_RET(groupIndexContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
165		167
166	auto result = l0op::GroupedDynamicMxQuant(selfContiguous, groupIndexContiguous, roundMode, dstType, blocksize,	168	auto result = l0op::GroupedDynamicMxQuant(selfContiguous, groupIndexContiguous, roundMode, dstType, blocksize,
167	- uniqueExecutor.get());	169	+ DEFAULT_SCALE_ALG, DEFAULT_DST_TYPE_MAX, uniqueExecutor.get());
168	const aclTensor *yOut = std::get<0>(result);	170	const aclTensor *yOut = std::get<0>(result);
169	const aclTensor *mxscaleOut = std::get<1>(result);	171	const aclTensor *mxscaleOut = std::get<1>(result);
170	// 如果出参y是非连续Tensor，需要把计算完的连续Tensor转非连续	172	// 如果出参y是非连续Tensor，需要把计算完的连续Tensor转非连续