Star1371
2131
代码
代码
Issues796
Pull Requests717
Actions
项目讨论
Wiki
分析
Star1371
2131
已合并
【9.2.0beta2】【cherry-pick】DequantSituQuant和SituMxQuant算子 #8958
east_yang创建于 8月20日
【9.2.0beta2】【cherry-pick】DequantSituQuant和SituMxQuant算子 #8958
已合并
east_yang创建于 8月20日
从east_yang/ops-nn_9393cherry-pick-mr-8798-1787210837614-auto合入到cann/ops-nn9.2.0-beta.2
共 63 个文件变更+9752-0
Mdocs/zh/menu_aclnn_api.md+2-0
@@ -83,6 +83,7 @@
83- [aclnnDeformableConv2d](../../conv/deformable_conv2d/docs/aclnnDeformableConv2d.md)83- [aclnnDeformableConv2d](../../conv/deformable_conv2d/docs/aclnnDeformableConv2d.md)
84- [aclnnDeformableConv2dBackward](../../conv/convolution_backward/docs/aclnnDeformableConv2dBackward.md)84- [aclnnDeformableConv2dBackward](../../conv/convolution_backward/docs/aclnnDeformableConv2dBackward.md)
85- [aclnnDequantBias](../../quant/dequant_bias/docs/aclnnDequantBias.md)85- [aclnnDequantBias](../../quant/dequant_bias/docs/aclnnDequantBias.md)
86+- [aclnnDequantSituQuant](../../quant/dequant_situ_quant/docs/aclnnDequantSituQuant.md)
86- [aclnnDequantSwigluQuant](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuant.md)87- [aclnnDequantSwigluQuant](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuant.md)
87- [aclnnDequantSwigluQuantV2](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuantV2.md)88- [aclnnDequantSwigluQuantV2](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuantV2.md)
88- [aclnnDualLevelQuantMatmulWeightNz](../../matmul/dual_level_quant_batch_matmul/docs/aclnnDualLevelQuantMatmulWeightNz.md)89- [aclnnDualLevelQuantMatmulWeightNz](../../matmul/dual_level_quant_batch_matmul/docs/aclnnDualLevelQuantMatmulWeightNz.md)
@@ -405,6 +406,7 @@
405- [aclnnSwigluGroupQuant](../../activation/swiglu_group_quant/docs/aclnnSwigluGroupQuant.md)406- [aclnnSwigluGroupQuant](../../activation/swiglu_group_quant/docs/aclnnSwigluGroupQuant.md)
406- [aclnnSwigluGroupQuantGrad](../../activation/swiglu_group_quant_grad/docs/aclnnSwigluGroupQuantGrad.md)407- [aclnnSwigluGroupQuantGrad](../../activation/swiglu_group_quant_grad/docs/aclnnSwigluGroupQuantGrad.md)
407- [aclnnSwigluMxQuant](../../quant/swiglu_mx_quant/docs/aclnnSwigluMxQuant.md)408- [aclnnSwigluMxQuant](../../quant/swiglu_mx_quant/docs/aclnnSwigluMxQuant.md)
409+- [aclnnSituMxQuant](../../quant/situ_mx_quant/docs/aclnnSituMxQuant.md)
408- [aclnnSwigluMxQuantWithDualAxis](../../quant/swiglu_mx_quant_with_dual_axis/docs/aclnnSwigluMxQuantWithDualAxis.md)410- [aclnnSwigluMxQuantWithDualAxis](../../quant/swiglu_mx_quant_with_dual_axis/docs/aclnnSwigluMxQuantWithDualAxis.md)
409- [aclnnSwish](../../activation/swish/docs/aclnnSwish.md)411- [aclnnSwish](../../activation/swish/docs/aclnnSwish.md)
410- [aclnnSwishBackward](../../activation/swish_grad/docs/aclnnSwishBackward.md)412- [aclnnSwishBackward](../../activation/swish_grad/docs/aclnnSwishBackward.md)
Mdocs/zh/menu_torch_api.md+2-0
@@ -13,3 +13,5 @@
13- [flat_quant](../../quant/flat_quant/docs/torchapi_flat_quant.md)13- [flat_quant](../../quant/flat_quant/docs/torchapi_flat_quant.md)
14- [matmul_emu_split_weight](../../matmul/matmul_emu_split_weight/docs/torchapi_matmul_emu_split_weight.md)14- [matmul_emu_split_weight](../../matmul/matmul_emu_split_weight/docs/torchapi_matmul_emu_split_weight.md)
15- [mx_to_block_mx_quant](../../quant/mx_to_block_mx_quant/docs/torchapi_mx_to_block_mx_quant.md)15- [mx_to_block_mx_quant](../../quant/mx_to_block_mx_quant/docs/torchapi_mx_to_block_mx_quant.md)
16+- [dequant_situ_quant](../../quant/dequant_situ_quant/docs/torchapi_dequant_situ_quant.md)
17+- [situ_mx_quant](../../quant/situ_mx_quant/docs/torchapi_situ_mx_quant.md)
Mdocs/zh/op_api_list.md+2-0
@@ -113,10 +113,12 @@
113| [aclnnDeformableConv2d](../../conv/deformable_conv2d/docs/aclnnDeformableConv2d.md) | 实现卷积功能，支持2D卷积，同时支持可变形卷积、分组卷积。 | 默认确定性实现 | 默认确定性实现 |113| [aclnnDeformableConv2d](../../conv/deformable_conv2d/docs/aclnnDeformableConv2d.md) | 实现卷积功能，支持2D卷积，同时支持可变形卷积、分组卷积。 | 默认确定性实现 | 默认确定性实现 |
114| [aclnnDeformableConv2dBackward](../../conv/convolution_backward/docs/aclnnDeformableConv2dBackward.md) | 实现卷积反向功能，支持2D卷积反向，同时支持可变形卷积、分组卷积反向。 | - | 默认确定性实现 |114| [aclnnDeformableConv2dBackward](../../conv/convolution_backward/docs/aclnnDeformableConv2dBackward.md) | 实现卷积反向功能，支持2D卷积反向，同时支持可变形卷积、分组卷积反向。 | - | 默认确定性实现 |
115| [aclnnDequantBias](../../quant/dequant_bias/docs/aclnnDequantBias.md) | 对输入x反量化操作，将输入的int32的数据转化为FLOAT16/BFLOAT16输出。 | 默认确定性实现 | - |115| [aclnnDequantBias](../../quant/dequant_bias/docs/aclnnDequantBias.md) | 对输入x反量化操作，将输入的int32的数据转化为FLOAT16/BFLOAT16输出。 | 默认确定性实现 | - |
116+| [aclnnDequantSituQuant](../../quant/dequant_situ_quant/docs/aclnnDequantSituQuant.md) | 在Situ激活函数前后添加dequant和quant操作，实现x的DequantSituQuant计算。 | 默认确定性实现 | - |
116| [aclnnDequantSwigluQuant](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuant.md) | 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 | 默认确定性实现 | 默认确定性实现 |117| [aclnnDequantSwigluQuant](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuant.md) | 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 | 默认确定性实现 | 默认确定性实现 |
117| [aclnnDequantSwigluQuantV2](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuantV2.md) | 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 | 默认确定性实现 | 默认确定性实现 |118| [aclnnDequantSwigluQuantV2](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuantV2.md) | 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 | 默认确定性实现 | 默认确定性实现 |
118| [aclnnSwigluGroup](../../activation/swiglu_group/docs/aclnnSwigluGroup.md) | 执行SwiGLU激活，实现x的SwiGlu计算，输出与x相同数据类型的激活结果。 | - | 默认确定性实现 |119| [aclnnSwigluGroup](../../activation/swiglu_group/docs/aclnnSwigluGroup.md) | 执行SwiGLU激活，实现x的SwiGlu计算，输出与x相同数据类型的激活结果。 | - | 默认确定性实现 |
119| [aclnnSwigluMxQuant](../../quant/swiglu_mx_quant/docs/aclnnSwigluMxQuant.md) | 在Swish门控线性单元激活函数后添加DynamicMxQuant操作，实现x的SwigluMxQuant计算。 | - | 默认确定性实现 |120| [aclnnSwigluMxQuant](../../quant/swiglu_mx_quant/docs/aclnnSwigluMxQuant.md) | 在Swish门控线性单元激活函数后添加DynamicMxQuant操作，实现x的SwigluMxQuant计算。 | - | 默认确定性实现 |
121+| [aclnnSituMxQuant](../../quant/situ_mx_quant/docs/aclnnSituMxQuant.md) | 将Situ激活函数与动态MX量化融合，实现x的SituMxQuant计算。 | - | 默认确定性实现 |
120| [aclnnSwigluMxQuantWithDualAxis](../../quant/swiglu_mx_quant_with_dual_axis/docs/aclnnSwigluMxQuantWithDualAxis.md) | 在Swish门控线性单元激活函数后添加GroupDynamicMxQuant操作，根据传入的分组索引的起始值，对swiglu后的数据进行分组的-1轴和-2轴同时动态量化。 | - | 默认确定性实现 |122| [aclnnSwigluMxQuantWithDualAxis](../../quant/swiglu_mx_quant_with_dual_axis/docs/aclnnSwigluMxQuantWithDualAxis.md) | 在Swish门控线性单元激活函数后添加GroupDynamicMxQuant操作，根据传入的分组索引的起始值，对swiglu后的数据进行分组的-1轴和-2轴同时动态量化。 | - | 默认确定性实现 |
121| [aclnnSwigluBackwardMxQuantWithDualAxis](../../quant/swiglu_backward_mx_quant_with_dual_axis/docs/aclnnSwigluBackwardMxQuantWithDualAxis.md) | 在Swish门控线性单元激活的反向计算函数后添加GroupDynamicMxQuant操作，根据传入的分组索引的起始值，对swiglu后的数据进行分组的-1轴和-2轴同时动态量化。 | - | 默认确定性实现 |123| [aclnnSwigluBackwardMxQuantWithDualAxis](../../quant/swiglu_backward_mx_quant_with_dual_axis/docs/aclnnSwigluBackwardMxQuantWithDualAxis.md) | 在Swish门控线性单元激活的反向计算函数后添加GroupDynamicMxQuant操作，根据传入的分组索引的起始值，对swiglu后的数据进行分组的-1轴和-2轴同时动态量化。 | - | 默认确定性实现 |
122| [aclnnGroupedDynamicMxQuant](../../quant/grouped_dynamic_mx_quant/docs/aclnnGroupedDynamicMxQuant.md) | 根据传入的分组索引的起始值，对传入的数据进行分组的float8的动态量化。 | - | 默认确定性实现 |124| [aclnnGroupedDynamicMxQuant](../../quant/grouped_dynamic_mx_quant/docs/aclnnGroupedDynamicMxQuant.md) | 根据传入的分组索引的起始值，对传入的数据进行分组的float8的动态量化。 | - | 默认确定性实现 |
Mdocs/zh/op_list.md+20-0
@@ -4448,6 +4448,16 @@
4448 <td>AI Core</td>4448 <td>AI Core</td>
4449 <td>对输入x反量化操作，将输入的int32的数据转化为FLOAT16/BFLOAT16输出。</td>4449 <td>对输入x反量化操作，将输入的int32的数据转化为FLOAT16/BFLOAT16输出。</td>
4450 </tr>4450 </tr>
4451+ <tr>
4452+ <td>quant</td>
4453+ <td><a href="../../quant/dequant_situ_quant/README.md">dequant_situ_quant</a></td>
4454+ <td>✓</td>
4455+ <td>✓</td>
4456+ <td>✓</td>
4457+ <td>✓</td>
4458+ <td>AI Core</td>
4459+ <td>在Situ激活函数前后添加dequant和quant操作，实现x的DequantSituQuant计算。</td>
4460+ </tr>
4451 <tr>4461 <tr>
4452 <td>quant</td>4462 <td>quant</td>
4453 <td><a href="../../quant/dequant_swiglu_quant/README.md">dequant_swiglu_quant</a></td>4463 <td><a href="../../quant/dequant_swiglu_quant/README.md">dequant_swiglu_quant</a></td>
@@ -4728,6 +4738,16 @@
4728 <td>AI Core</td>4738 <td>AI Core</td>
4729 <td>在SwiGlu激活函数后添加mx_quant操作，实现输入x的SwigluMxQuant计算，支持FP8或FP4量化输出。</td>4739 <td>在SwiGlu激活函数后添加mx_quant操作，实现输入x的SwigluMxQuant计算，支持FP8或FP4量化输出。</td>
4730 </tr>4740 </tr>
4741+ <tr>
4742+ <td>quant</td>
4743+ <td><a href="../../quant/situ_mx_quant/README.md">situ_mx_quant</a></td>
4744+ <td>✓</td>
4745+ <td>✓</td>
4746+ <td>✓</td>
4747+ <td>✓</td>
4748+ <td>AI Core</td>
4749+ <td>将Situ激活函数与动态MX量化融合，实现输入x的SituMxQuant计算，支持FP8量化输出。</td>
4750+ </tr>
4731 <tr>4751 <tr>
4732 <td>quant</td>4752 <td>quant</td>
4733 <td><a href="../../quant/swiglu_mx_quant_with_dual_axis/README.md">swiglu_mx_quant_with_dual_axis</a></td>4753 <td><a href="../../quant/swiglu_mx_quant_with_dual_axis/README.md">swiglu_mx_quant_with_dual_axis</a></td>
Mdocs/zh/torch_api_list.md+2-0
@@ -36,3 +36,5 @@
36|[flat_quant](../../quant/flat_quant/docs/torchapi_flat_quant.md)| 为矩阵x依次进行两次克罗内克积小矩阵乘法，然后针对矩阵乘的结果进行量化处理。 |默认支持确定性计算|默认支持确定性计算|36|[flat_quant](../../quant/flat_quant/docs/torchapi_flat_quant.md)| 为矩阵x依次进行两次克罗内克积小矩阵乘法，然后针对矩阵乘的结果进行量化处理。 |默认支持确定性计算|默认支持确定性计算|
37|[matmul_emu_split_weight](../../matmul/matmul_emu_split_weight/docs/torchapi_matmul_emu_split_weight.md)| 使用双路BF16 GEMM融合模拟FP32精度矩阵乘法。离线将FP32权重拆分为高位BF16与低位残差BF16，推理阶段执行两次BF16 GEMM并做线性组合，激活值全程保持BF16，两路矩阵乘均运行在Cube上。 |-|默认支持确定性计算|37|[matmul_emu_split_weight](../../matmul/matmul_emu_split_weight/docs/torchapi_matmul_emu_split_weight.md)| 使用双路BF16 GEMM融合模拟FP32精度矩阵乘法。离线将FP32权重拆分为高位BF16与低位残差BF16，推理阶段执行两次BF16 GEMM并做线性组合，激活值全程保持BF16，两路矩阵乘均运行在Cube上。 |-|默认支持确定性计算|
38|[mx_to_block_mx_quant](../../quant/mx_to_block_mx_quant/docs/torchapi_mx_to_block_mx_quant.md)| 将调用 npu_dynamic_mx_quant 量化得到的 FLOAT4 的 Tensor 结合 FLOAT8_E8M0 缩放系数，转换为 FLOAT8 分块量化格式，同时输出 -1 轴和 -2 轴方向的量化尺度。|-| 默认确定性实现 |38|[mx_to_block_mx_quant](../../quant/mx_to_block_mx_quant/docs/torchapi_mx_to_block_mx_quant.md)| 将调用 npu_dynamic_mx_quant 量化得到的 FLOAT4 的 Tensor 结合 FLOAT8_E8M0 缩放系数，转换为 FLOAT8 分块量化格式，同时输出 -1 轴和 -2 轴方向的量化尺度。|-| 默认确定性实现 |
39+|[dequant_situ_quant](../../quant/dequant_situ_quant/docs/torchapi_dequant_situ_quant.md)|在Situ激活函数前后添加dequant和quant操作，实现x的DequantSituQuant计算。|默认支持确定性计算|-|
40+|[situ_mx_quant](../../quant/situ_mx_quant/docs/torchapi_situ_mx_quant.md)|将Situ激活函数与动态MX量化融合，实现x的SituMxQuant计算，支持FP8量化输出。|-|默认支持确定性计算|
Aquant/dequant_situ_quant/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST AND NOT BENCHMARK)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+list(REMOVE_ITEM CURRENT_DIRS examples)
16+foreach(SUB_DIR ${CURRENT_DIRS})
17+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18+ add_subdirectory(${SUB_DIR})
19+ endif()
20+endforeach()
Aquant/dequant_situ_quant/README.md+204-0
@@ -0,0 +1,204 @@
1+# DequantSituQuant
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | × |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | × |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：在Situ激活函数前后添加dequant和quant操作，实现x的DequantSituQuant计算。
17+ 
18+- 计算公式：
19+ 
20+ 1. 根据输入数据类型x的不同，反量化路径不同：
21+ 
22+ - INT8路径：
23+ 
24+ $$
25+ dequantOut_i = cast\_to\_float(x_i) \times weight\_scale_i + bias_i
26+ $$
27+ 
28+ - INT32路径：
29+ 
30+ $$
31+ dequantOut_i = cast\_to\_float(x_i) \times weight\_scale_i \times activation\_scale_i + bias_i
32+ $$
33+ 
34+ - BF16/FLOAT16路径（预反量化）：
35+ 
36+ $$
37+ dequantOut_i = cast\_to\_float(x_i)
38+ $$
39+ 
40+ 2. Situ激活：
41+ 
42+ $$
43+ situ_a = \beta \times \tanh(gate / \beta) \times sigmoid(gate)
44+ $$
45+ 
46+ 当linear_beta > 0时：
47+ 
48+ $$
49+ up = linear\_beta \times \tanh(up / linear\_beta)
50+ $$
51+ 
52+ $$
53+ situOut = situ_a \times up
54+ $$
55+ 
56+ 其中，当activate_left为true时，gate取dequantOut的前半部分，up取后半部分；当activate_left为false时，gate取dequantOut的后半部分，up取前半部分。
57+ 
58+ 3. 量化：
59+ 
60+ - static模式：
61+ 
62+ $$
63+ out_i = trunc(situOut_i / quant\_scale_i + quant\_offset_i)
64+ $$
65+ 
66+ - dynamic模式：
67+ 
68+ $$
69+ scale_i = absmax(situOut_i) / 127
70+ $$
71+ 
72+ $$
73+ out_i = trunc(situOut_i / scale_i)
74+ $$
75+ 
76+## 参数说明
77+ 
78+<table style="undefined;table-layout: fixed; width: 951px"><colgroup>
79+ <col style="width: 121px">
80+ <col style="width: 144px">
81+ <col style="width: 313px">
82+ <col style="width: 257px">
83+ <col style="width: 116px">
84+ </colgroup>
85+ <thead>
86+ <tr>
87+ <th>参数名</th>
88+ <th>输入/输出/属性</th>
89+ <th>描述</th>
90+ <th>数据类型</th>
91+ <th>数据格式</th>
92+ </tr></thead>
93+ <tbody>
94+ <tr>
95+ <td>x</td>
96+ <td>输入</td>
97+ <td>输入待处理的数据，公式中的x。输入不支持包含±inf或nan。</td>
98+ <td>INT8、INT32、BFLOAT16、FLOAT16</td>
99+ <td>ND</td>
100+ </tr>
101+ <tr>
102+ <td>weight_scale</td>
103+ <td>输入</td>
104+ <td>反量化的weight scale，公式中的weight_scale。输入不支持包含±inf或nan。</td>
105+ <td>FLOAT</td>
106+ <td>ND</td>
107+ </tr>
108+ <tr>
109+ <td>activation_scale</td>
110+ <td>输入</td>
111+ <td>反量化的activation scale，公式中的activation_scale。输入不支持包含±inf或nan。</td>
112+ <td>FLOAT</td>
113+ <td>ND</td>
114+ </tr>
115+ <tr>
116+ <td>bias</td>
117+ <td>输入</td>
118+ <td>反量化的bias，公式中的bias。输入不支持包含±inf或nan。</td>
119+ <td>FLOAT</td>
120+ <td>ND</td>
121+ </tr>
122+ <tr>
123+ <td>quant_scale</td>
124+ <td>输入</td>
125+ <td>量化的scale，公式中的quant_scale。输入不支持包含±inf或nan。</td>
126+ <td>FLOAT</td>
127+ <td>ND</td>
128+ </tr>
129+ <tr>
130+ <td>quant_offset</td>
131+ <td>输入</td>
132+ <td>量化的offset，公式中的quant_offset。输入不支持包含±inf或nan。</td>
133+ <td>FLOAT</td>
134+ <td>ND</td>
135+ </tr>
136+ <tr>
137+ <td>group_index</td>
138+ <td>输入</td>
139+ <td>MoE分组需要的group_index。输入不支持包含±inf或nan。</td>
140+ <td>INT64</td>
141+ <td>ND</td>
142+ </tr>
143+ <tr>
144+ <td>beta</td>
145+ <td>属性</td>
146+ <td>Situ激活的beta参数，公式中的β。不能为0。默认4.0。</td>
147+ <td>FLOAT</td>
148+ <td>-</td>
149+ </tr>
150+ <tr>
151+ <td>linear_beta</td>
152+ <td>属性</td>
153+ <td>Situ激活的linear_beta参数，公式中的linear_beta。当值≤0时不启用。默认25.0。</td>
154+ <td>FLOAT</td>
155+ <td>-</td>
156+ </tr>
157+ <tr>
158+ <td>activate_left</td>
159+ <td>属性</td>
160+ <td>表示gate取dequantOut的前半部分还是后半部分，公式中的activate_left。默认true。</td>
161+ <td>BOOL</td>
162+ <td>-</td>
163+ </tr>
164+ <tr>
165+ <td>quant_type</td>
166+ <td>属性</td>
167+ <td>表示量化模式，对应公式中的static/dynamic模式。支持"static"和"dynamic"。默认"dynamic"。</td>
168+ <td>STRING</td>
169+ <td>-</td>
170+ </tr>
171+ <tr>
172+ <td>y</td>
173+ <td>输出</td>
174+ <td>量化后的输出，公式中的out。</td>
175+ <td>INT8</td>
176+ <td>ND</td>
177+ </tr>
178+ <tr>
179+ <td>y_scale</td>
180+ <td>输出</td>
181+ <td>动态量化的scale，公式中的scale（仅dynamic模式有意义）。</td>
182+ <td>FLOAT</td>
183+ <td>ND</td>
184+ </tr>
185+ </tbody></table>
186+ 
187+## 约束说明
188+ 
189+- x的最后一维需要是2的倍数。
190+- 当x的数据类型为INT8时，x维度≥2维；当x的数据类型为INT32/BF16/FLOAT16时，x维度为2维。
191+- beta参数不能为0。
192+- INT8路径：必须提供weight_scale，禁止activation_scale和group_index。
193+- INT32：必须提供weight_scale和activation_scale，禁止quant_scale和quant_offset。
194+- BFLOAT16/FLOAT16：所有可选输入均不使用（预反量化模式）。
195+- 当quant_type为static时，quant_scale必须提供。
196+- 当quant_type为dynamic时，quant_scale可选（作为smoothScale使用）。
197+- 算子支持的输入张量的内存大小有上限，校验公式：weight_scale张量内存大小+bias张量内存大小+quant_scale张量内存大小+quant_offset张量内存大小 + （activation_scale张量内存大小 + y_scale张量内存大小）/40 + x张量最后一维H内存大小 * 10 < 192KB。
198+ 
199+## 调用说明
200+ 
201+| 调用方式 | 调用样例 | 说明 |
202+|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|
203+| aclnn调用 | [test_aclnn_dequant_situ_quant](./examples/test_aclnn_dequant_situ_quant.cpp) | 通过[aclnnDequantSituQuant](./docs/aclnnDequantSituQuant.md)接口方式调用DequantSituQuant算子。 |
204+| 图模式调用 | - | 通过[算子IR](./op_graph/dequant_situ_quant_proto.h)构图方式调用DequantSituQuant算子。 |
Aquant/dequant_situ_quant/docs/aclnnDequantSituQuant.md+573-0
@@ -0,0 +1,573 @@
1+# aclnnDequantSituQuant
2+ 
3+## 产品支持情况
4+ 
5+<!-- npu="950" id1 -->
6+- <term>Ascend 950PR/Ascend 950DT</term>：不支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>：支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- <term>Atlas 200I/500 A2 推理产品</term>：不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- <term>Atlas 推理系列产品</term>：不支持
19+<!-- end id5 -->
20+<!-- npu="910" id6 -->
21+- <term>Atlas 训练系列产品</term>：不支持
22+<!-- end id6 -->
23+ 
24+## 功能说明
25+ 
26+- 接口功能：在Situ激活函数前后添加dequant和quant操作，实现x的DequantSituQuant计算。
27+ 
28+- 计算公式：
29+ 
30+1. 根据输入数据类型x的不同，反量化路径不同：
31+ 
32+- INT8路径
33+ 
34+ $$
35+ dequantOut_i = cast\_to\_float(x_i) \times weight\_scale_i + bias_i
36+ $$
37+ 
38+- INT32路径
39+ 
40+ $$
41+ dequantOut_i = cast\_to\_float(x_i) \times weight\_scale_i \times activation\_scale_i + bias_i
42+ $$
43+ 
44+- BF16/FLOAT16路径（预反量化）
45+ 
46+ $$
47+ dequantOut_i = cast\_to\_float(x_i)
48+ $$
49+ 
50+1. Situ激活
51+ 
52+$$
53+situ_a = \beta \times \tanh(gate / \beta) \times sigmoid(gate)
54+$$
55+ 
56+当linear_beta > 0时：
57+ 
58+$$
59+up = linear\_beta \times \tanh(up / linear\_beta)
60+$$
61+ 
62+$$
63+situOut = situ_a \times up
64+$$
65+ 
66+ 其中，当activate_left为true时，gate取dequantOut的前半部分，up取后半部分；当activate_left为false时，gate取dequantOut的后半部分，up取前半部分。
67+ 
68+2. 量化
69+ 
70+- quant_type = static
71+ 
72+ $$
73+ y_i = trunc(situOut_i / quant\_scale_i + quant\_offset_i)
74+ $$
75+ 
76+- quant_type = dynamic
77+ 
78+ $$
79+ scale_i = absmax(situOut_i) / 127
80+ $$
81+ 
82+ $$
83+ y_i = trunc(situOut_i / scale_i)
84+ $$
85+ 
86+## 函数原型
87+ 
88+每个算子分为[两段式接口](../../../docs/zh/context/two_phase_api.md)，必须先调用"aclnnDequantSituQuantGetWorkspaceSize"接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用"aclnnDequantSituQuant"接口执行计算。
89+ 
90+```Cpp
91+aclnnStatus aclnnDequantSituQuantGetWorkspaceSize(
92+ const aclTensor* x,
93+ const aclTensor* weightScaleOptional,
94+ const aclTensor* activationScaleOptional,
95+ const aclTensor* biasOptional,
96+ const aclTensor* quantScaleOptional,
97+ const aclTensor* quantOffsetOptional,
98+ const aclTensor* groupIndexOptional,
99+ float beta,
100+ float linearBeta,
Aquant/dequant_situ_quant/docs/torchapi_dequant_situ_quant.md+350-0
@@ -0,0 +1,350 @@
1+# dequant_situ_quant
2+ 
3+## 产品支持情况
4+ 
5+<!-- npu="950" id1 -->
6+- <term>Ascend 950PR/Ascend 950DT</term>：不支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>：支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- <term>Atlas 200I/500 A2 推理产品</term>：不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- <term>Atlas 推理系列产品</term>：不支持
19+<!-- end id5 -->
20+<!-- npu="910" id6 -->
21+- <term>Atlas 训练系列产品</term>：不支持
22+<!-- end id6 -->
23+ 
24+## 功能说明
25+ 
26+- 接口功能：对输入x执行反量化(Dequant)、Situ激活函数和量化(Quant)的融合计算。底层封装aclnnDequantSituQuant。
27+ 
28+- 计算公式：
29+ 
30+ 1. 根据输入数据类型x的不同，反量化路径不同：
31+ 
32+ - INT8路径：
33+ 
34+ $$
35+ dequantOut_i = cast\_to\_float(x_i) \times weight\_scale_i + bias_i
36+ $$
37+ 
38+ - INT32路径：
39+ 
40+ $$
41+ dequantOut_i = cast\_to\_float(x_i) \times weight\_scale_i \times activation\_scale_i + bias_i
42+ $$
43+ 
44+ - BFLOAT16/FLOAT16路径（预反量化）：
45+ 
46+ $$
47+ dequantOut_i = cast\_to\_float(x_i)
48+ $$
49+ 2. Situ激活：
50+ 
51+ $$
52+ situ_a = \beta \times \tanh(gate / \beta) \times sigmoid(gate)
53+ $$
54+ 
55+ 当linear_beta > 0时：
56+ 
57+ $$
58+ up = linear\_beta \times \tanh(up / linear\_beta)
59+ $$
60+ 
61+ $$
62+ situOut = situ_a \times up
63+ $$
64+ 
65+ 其中，当activate_left为true时，gate取dequantOut的前半部分，up取后半部分；当activate_left为false时，gate取dequantOut的后半部分，up取前半部分。
66+ 3. 量化：
67+ 
68+ - static模式：
69+ 
70+ $$
71+ y_i = trunc(situOut_i / quant\_scale_i + quant\_offset_i)
72+ $$
73+ 
74+ - dynamic模式：
75+ 
76+ $$
77+ scale_i = absmax(situOut_i) / 127
78+ $$
79+ 
80+ $$
81+ y_i = trunc(situOut_i / scale_i)
82+ $$
83+ 
84+## 函数原型
85+ 
86+```python
87+cann_ops_nn.dequant_situ_quant(
88+ x,
89+ *,
90+ weight_scale=None,
91+ activation_scale=None,
92+ bias=None,
93+ quant_scale=None,
94+ quant_offset=None,
95+ group_index=None,
96+ beta=4.0,
97+ linear_beta=25.0,
98+ activate_left=True,
99+ quant_type="dynamic",
100+) -> (Tensor, Tensor)
101+```
102+ 
103+## 参数说明
104+ 
105+<table style="undefined;table-layout: fixed; width:1625px"><colgroup>
106+<col style="width: 147px">
107+<col style="width: 132px">
108+<col style="width: 132px">
109+<col style="width: 480px">
110+<col style="width: 330px">
111+<col style="width: 280px">
112+</colgroup>
113+<thead>
114+<tr>
115+ <th>参数名</th>
116+ <th>参数类型</th>
117+ <th>可选/必选</th>
118+ <th>描述</th>
119+ <th>数据类型</th>
120+ <th>维度(shape)</th>
121+</tr>
122+</thead>
123+<tbody>
124+ <tr>
125+ <td>x</td>
126+ <td>Tensor</td>
127+ <td>必选</td>
128+ <td>输入数据，对应公式中的x。不支持非连续，数据格式支持ND。</td>
129+ <td>int8、int32、bfloat16、float16</td>
130+ <td>INT8: 2-8维；其他: 2维，最后一维为2H</td>
131+ </tr>
132+ <tr>
133+ <td>weight_scale</td>
134+ <td>Tensor</td>
135+ <td>可选</td>
136+ <td>反量化weight scale，对应公式中的weight_scale。不支持非连续，数据格式支持ND。</td>
137+ <td>float32</td>
138+ <td>1维，shape为(2H,)或(1,)</td>
139+ </tr>
140+ <tr>
141+ <td>activation_scale</td>
142+ <td>Tensor</td>
143+ <td>可选</td>
144+ <td>反量化activation scale，对应公式中的activation_scale。不支持非连续，数据格式支持ND。</td>
145+ <td>float32</td>
146+ <td>1维，shape为(1,)</td>
147+ </tr>
148+ <tr>
149+ <td>bias</td>
150+ <td>Tensor</td>
151+ <td>可选</td>
152+ <td>反量化bias，对应公式中的bias。不支持非连续，数据格式支持ND。</td>
153+ <td>float32</td>
154+ <td>1维，shape为(2H,)或(1,)</td>
155+ </tr>
156+ <tr>
157+ <td>quant_scale</td>
158+ <td>Tensor</td>
159+ <td>可选</td>
160+ <td>量化scale，对应公式中的quant_scale。不支持非连续，数据格式支持ND。</td>
161+ <td>float32</td>
162+ <td>1维，shape为(H,)或(1,)</td>
163+ </tr>
164+ <tr>
165+ <td>quant_offset</td>
166+ <td>Tensor</td>
167+ <td>可选</td>
168+ <td>量化offset，对应公式中的quant_offset。不支持非连续，数据格式支持ND。</td>
169+ <td>float32</td>
170+ <td>1维，shape为(H,)或(1,)</td>
171+ </tr>
172+ <tr>
173+ <td>group_index</td>
174+ <td>Tensor</td>
175+ <td>可选</td>
176+ <td>MoE group索引。不支持非连续，数据格式支持ND。</td>
177+ <td>int64</td>
178+ <td>1维，shape为(K,)</td>
179+ </tr>
180+ <tr>
181+ <td>beta</td>
182+ <td>float</td>
183+ <td>可选</td>
184+ <td>Situ激活的beta参数，对应公式中的β。默认4.0。</td>
185+ <td>-</td>
186+ <td>-</td>
187+ </tr>
188+ <tr>
189+ <td>linear_beta</td>
190+ <td>float</td>
191+ <td>可选</td>
192+ <td>Situ激活的linear_beta参数，对应公式中的linear_beta。默认25.0。</td>
193+ <td>-</td>
194+ <td>-</td>
195+ </tr>
196+ <tr>
197+ <td>activate_left</td>
198+ <td>bool</td>
199+ <td>可选</td>
200+ <td>表示gate取dequantOut的前半部分还是后半部分，对应公式中的activate_left。默认true。</td>
201+ <td>-</td>
202+ <td>-</td>
203+ </tr>
204+ <tr>
205+ <td>quant_type</td>
206+ <td>str</td>
207+ <td>可选</td>
208+ <td>量化模式，对应公式中的static/dynamic模式。支持"static"和"dynamic"。默认"dynamic"。</td>
209+ <td>-</td>
210+ <td>-</td>
211+ </tr>
212+</tbody>
213+</table>
214+ 
215+## 返回值说明
216+ 
217+<table style="undefined;table-layout: fixed; width:1625px"><colgroup>
218+<col style="width: 147px">
219+<col style="width: 132px">
220+<col style="width: 132px">
221+<col style="width: 480px">
222+<col style="width: 330px">
223+<col style="width: 280px">
224+</colgroup>
225+<thead>
226+<tr>
227+ <th>输出名</th>
228+ <th>输出类型</th>
229+ <th>可选/必选</th>
230+ <th>描述</th>
231+ <th>数据类型</th>
232+ <th>维度(shape)</th>
233+</tr>
234+</thead>
235+<tbody>
236+ <tr>
237+ <td>y</td>
238+ <td>Tensor</td>
239+ <td>必选</td>
240+ <td>量化输出，对应公式中的y。不支持非连续，数据格式支持ND。</td>
241+ <td>int8</td>
242+ <td>INT8: x.shape[:-1]+[H]；其他: [M, H]</td>
243+ </tr>
244+ <tr>
245+ <td>y_scale</td>
246+ <td>Tensor</td>
247+ <td>必选</td>
248+ <td>动态量化scale，对应公式中的scale（仅dynamic模式有意义）。不支持非连续，数据格式支持ND。</td>
249+ <td>float32</td>
250+ <td>INT8: x.shape[:-1]；其他: [M]</td>
251+ </tr>
252+</tbody>
253+</table>
254+ 
255+其中H = x.shape[-1] / 2。
256+ 
257+## 约束说明
258+ 
259+- x及可选Tensor均需为NPU Tensor；可选Tensor可以传None。
260+- x的最后一维必须为偶数。当x的数据类型为INT8时，维度≥2维；当x的数据类型为INT32/BF16/FLOAT16时，维度为2维。
261+- beta参数不能为0。
262+- 当quant_type为"static"时，quant_scale必须提供。
263+- 支持空Tensor（当x.size(0) == 0时直接返回空输出）。
264+- 关于quant_type的约束说明如下：
265+ 
266+ | quant_type | 含义 | quant_scale | quant_offset |
267+ | ---- | ---- | ---- | ---- |
268+ | "static" | 静态量化 | 必选 | 可选 |
269+ | "dynamic" | 动态量化 | 可选（作为smoothScale） | 不生效 |
270+ 
271+- 各数据类型的输入约束如下：
272+ 
273+ | x数据类型 | 必选输入 | 可选输入 | 禁止输入 |
274+ | ---- | ---- | ---- | ---- |
275+ | INT8 | x, weight_scale | bias, quant_scale, quant_offset | activation_scale, group_index |
276+ | INT32 | x, weight_scale, activation_scale | bias, group_index | quant_scale, quant_offset |
277+ | BF16 | x | — | weight_scale, activation_scale, bias, quant_scale, quant_offset, group_index |
278+ | FLOAT16 | x | — | weight_scale, activation_scale, bias, quant_scale, quant_offset, group_index |
279+ 
280+- 该接口支持训练、推理场景下使用。
281+- 该接口支持单算子模式和TorchAir图模式调用。
282+ 
283+## 确定性计算
284+ 
285+默认支持确定性计算。
286+ 
287+## 调用示例
288+ 
289+- 单算子模式调用（eager）
290+ 
291+ ```python
292+ import torch
293+ import torch_npu
294+ import cann_ops_nn
295+ 
296+ # INT8, dynamic量化
297+ x = torch.randint(-127, 127, (16, 64), dtype=torch.int8).npu()
298+ weight_scale = torch.full((64,), 0.1, dtype=torch.float32).npu()
299+ 
300+ y, y_scale = cann_ops_nn.dequant_situ_quant(
301+ x, weight_scale=weight_scale,
302+ beta=4.0, linear_beta=25.0, activate_left=True, quant_type="dynamic",
303+ )
304+ print("y:============", y.shape, y.cpu())
305+ print("y_scale:============", y_scale.shape, y_scale.cpu())
306+ 
307+ # BF16, 预反量化
308+ x_bf16 = torch.randn(32, 128, dtype=torch.bfloat16).npu()
309+ y_bf16, y_scale_bf16 = cann_ops_nn.dequant_situ_quant(
310+ x_bf16, beta=4.0, linear_beta=25.0, activate_left=True, quant_type="dynamic",
311+ )
312+ print("y_bf16:============", y_bf16.shape, y_bf16.cpu())
313+ ```
314+ 
315+- 图模式（torchair）调用
316+ 
317+ ```python
318+ import torch
319+ import torch_npu
320+ import torchair
321+ import cann_ops_nn
322+ 
323+ 
324+ class NetModel(torch.nn.Module):
325+ def __init__(self):
326+ super().__init__()
327+ 
328+ def forward(self, x, weight_scale, quant_scale):
329+ return cann_ops_nn.dequant_situ_quant(
330+ x, weight_scale=weight_scale, quant_scale=quant_scale,
331+ beta=4.0, linear_beta=25.0, activate_left=True, quant_type="static",
332+ )
333+ 
334+ 
335+ def dequant_situ_quant_test():
336+ x = torch.randint(-127, 127, (16, 64), dtype=torch.int8).npu()
337+ weight_scale = torch.full((64,), 0.1, dtype=torch.float32).npu()
338+ quant_scale = torch.tensor([1.0], dtype=torch.float32).npu()
339+ model = NetModel()
340+ config = torchair.CompilerConfig()
341+ npu_backend = torchair.get_npu_backend(compiler_config=config)
342+ model = torch.compile(model, backend=npu_backend, dynamic=True, fullgraph=True)
343+ y, y_scale = model(x, weight_scale, quant_scale)
344+ print("y:============", y.shape, y.cpu())
345+ print("y_scale:============", y_scale.shape, y_scale.cpu())
346+ 
347+ 
348+ if __name__ == "__main__":
349+ dequant_situ_quant_test()
350+ ```
Aquant/dequant_situ_quant/examples/test_aclnn_dequant_situ_quant.cpp+451-0
@@ -0,0 +1,451 @@
1+/**
2+ * Copyright (c) 2025 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+#include <iostream>
12+#include <vector>
13+#include <cmath>
14+#include <cstring>
15+#include <algorithm>
16+#include "acl/acl.h"
17+#include "aclnnop/aclnn_dequant_situ_quant.h"
18+ 
19+#define CHECK_RET(cond, return_expr) \
20+ do { \
21+ if (!(cond)) { \
22+ return_expr; \
23+ } \
24+ } while (0)
25+ 
26+#define LOG_PRINT(message, ...) \
27+ do { \
28+ printf(message, ##__VA_ARGS__); \
29+ } while (0)
30+ 
31+int64_t GetShapeSize(const std::vector<int64_t>& shape)
32+{
33+ int64_t shapeSize = 1;
34+ for (auto i : shape) {
35+ shapeSize *= i;
36+ }
37+ return shapeSize;
38+}
39+ 
40+template <typename T>
41+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
42+ aclDataType dataType, aclTensor** tensor)
43+{
44+ auto size = GetShapeSize(shape) * sizeof(T);
45+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
46+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
47+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
48+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
49+ 
50+ std::vector<int64_t> strides(shape.size(), 1);
51+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
52+ strides[i] = shape[i + 1] * strides[i + 1];
53+ }
54+ 
55+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
56+ shape.data(), shape.size(), *deviceAddr);
57+ return 0;
58+}
59+ 
60+// Truncate towards zero (matching kernel's CAST_TRUNC for fp16→int8)
61+int8_t TruncToInt8(float val)
62+{
63+ int32_t intVal = static_cast<int32_t>(val); // truncation towards zero
64+ intVal = std::max(-128, std::min(127, intVal));
65+ return static_cast<int8_t>(intVal);
66+}
67+ 
68+// CPU reference: DequantSituQuant static mode
69+void CpuStaticQuant(const std::vector<int8_t>& x, const std::vector<float>& dequantScale,
70+ const std::vector<float>& quantScale, const std::vector<float>& quantOffset, int64_t rowLen,
71+ int64_t inDimy, int64_t outDimy, double beta, double linearBeta, bool activateLeft,
72+ std::vector<int8_t>& yOut)
73+{
74+ for (int64_t r = 0; r < rowLen; r++) {
75+ for (int64_t c = 0; c < outDimy; c++) {
76+ int64_t gateIdx = r * inDimy + (activateLeft ? c : c + outDimy);
77+ int64_t upIdx = r * inDimy + (activateLeft ? c + outDimy : c);
78+ 
79+ float gate = static_cast<float>(x[gateIdx]) * dequantScale[gateIdx % dequantScale.size()];
80+ float up = static_cast<float>(x[upIdx]) * dequantScale[upIdx % dequantScale.size()];
81+ 
82+ float invBeta = 1.0f / static_cast<float>(beta);
83+ float tanhResult = std::tanh(gate * invBeta) * static_cast<float>(beta);
84+ float sigmoidResult = 1.0f / (1.0f + std::exp(-gate));
85+ float situA = tanhResult * sigmoidResult;
86+ 
87+ if (linearBeta > 0.0) {
88+ float invLb = 1.0f / static_cast<float>(linearBeta);
89+ up = std::tanh(up * invLb) * static_cast<float>(linearBeta);
90+ }
91+ 
92+ float situOut = situA * up;
93+ 
94+ float qs = quantScale.size() == 1 ? quantScale[0] : quantScale[c];
95+ float qo = quantOffset.empty() ? 0.0f : (quantOffset.size() == 1 ? quantOffset[0] : quantOffset[c]);
96+ float quantized = situOut / qs + qo;
97+ 
98+ yOut[r * outDimy + c] = TruncToInt8(quantized);
99+ }
100+ }
101+}
102+ 
103+// CPU reference: DequantSituQuant dynamic mode
104+void CpuDynamicQuant(const std::vector<int8_t>& x, const std::vector<float>& dequantScale, int64_t rowLen,
105+ int64_t inDimy, int64_t outDimy, double beta, double linearBeta, bool activateLeft,
106+ std::vector<int8_t>& yOut, std::vector<float>& yScaleOut)
107+{
108+ std::vector<float> situOutBuf(rowLen * outDimy, 0.0f);
109+ 
110+ for (int64_t r = 0; r < rowLen; r++) {
111+ for (int64_t c = 0; c < outDimy; c++) {
112+ int64_t gateIdx = r * inDimy + (activateLeft ? c : c + outDimy);
113+ int64_t upIdx = r * inDimy + (activateLeft ? c + outDimy : c);
114+ 
115+ float gate = static_cast<float>(x[gateIdx]) * dequantScale[gateIdx % dequantScale.size()];
116+ float up = static_cast<float>(x[upIdx]) * dequantScale[upIdx % dequantScale.size()];
117+ 
118+ float invBeta = 1.0f / static_cast<float>(beta);
119+ float tanhResult = std::tanh(gate * invBeta) * static_cast<float>(beta);
120+ float sigmoidResult = 1.0f / (1.0f + std::exp(-gate));
121+ float situA = tanhResult * sigmoidResult;
122+ 
123+ if (linearBeta > 0.0) {
124+ float invLb = 1.0f / static_cast<float>(linearBeta);
125+ up = std::tanh(up * invLb) * static_cast<float>(linearBeta);
126+ }
127+ 
128+ situOutBuf[r * outDimy + c] = situA * up;
129+ }
130+ 
131+ float absMax = 0.0f;
132+ for (int64_t c = 0; c < outDimy; c++) {
133+ absMax = std::max(absMax, std::abs(situOutBuf[r * outDimy + c]));
134+ }
135+ float scale = absMax / 127.0f;
136+ if (scale == 0.0f)
137+ scale = 1.0f;
138+ float invScale = 1.0f / scale;
139+ yScaleOut[r] = scale;
140+ 
141+ for (int64_t c = 0; c < outDimy; c++) {
142+ yOut[r * outDimy + c] = TruncToInt8(situOutBuf[r * outDimy + c] * invScale);
143+ }
144+ }
145+}
146+ 
147+// ============ Test 1: static quant, scalar quant_scale, no bias ============
148+int TestStaticQuantScalar(int32_t deviceId, aclrtStream stream)
149+{
150+ LOG_PRINT("\n========== Test 1: Static Quant (scalar scale, no bias) ==========\n");
151+ 
152+ int64_t rowLen = 16;
153+ int64_t inDimy = 64;
154+ int64_t outDimy = 32;
155+ double beta = 1.0;
156+ double linearBeta = 0.0;
157+ bool activateLeft = false;
158+ 
159+ std::vector<int64_t> xShape = {rowLen, inDimy};
160+ std::vector<int64_t> dequantScaleShape = {inDimy};
161+ std::vector<int64_t> quantScaleShape = {1};
162+ std::vector<int64_t> quantOffsetShape = {1};
163+ std::vector<int64_t> yShape = {rowLen, outDimy};
164+ std::vector<int64_t> yScaleOutShape = {rowLen};
165+ 
166+ auto xSize = GetShapeSize(xShape);
167+ auto dsSize = GetShapeSize(dequantScaleShape);
168+ auto qsSize = GetShapeSize(quantScaleShape);
169+ auto qoSize = GetShapeSize(quantOffsetShape);
170+ auto ySize = GetShapeSize(yShape);
171+ auto yScaleSize = GetShapeSize(yScaleOutShape);
172+ 
173+ std::vector<int8_t> xHostData(xSize);
174+ for (int64_t i = 0; i < xSize; i++) {
175+ xHostData[i] = static_cast<int8_t>((i * 7 + 3) % 100);
176+ }
177+ std::vector<float> dequantScaleHostData(dsSize, 0.1f);
178+ std::vector<float> quantScaleHostData(qsSize, 1.0f);
179+ std::vector<float> quantOffsetHostData(qoSize, 0.0f);
180+ std::vector<int8_t> yHostData(ySize, 0);
181+ std::vector<float> yScaleOutHostData(yScaleSize, 0.0f);
182+ 
183+ void* xDeviceAddr = nullptr;
184+ void* dsDeviceAddr = nullptr;
185+ void* qsDeviceAddr = nullptr;
186+ void* qoDeviceAddr = nullptr;
187+ void* yDeviceAddr = nullptr;
188+ void* yScaleDeviceAddr = nullptr;
189+ aclTensor* x = nullptr;
190+ aclTensor* dequantScale = nullptr;
191+ aclTensor* quantScale = nullptr;
192+ aclTensor* quantOffset = nullptr;
193+ aclTensor* y = nullptr;
194+ aclTensor* yScaleOut = nullptr;
195+ 
196+ auto ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_INT8, &x);
197+ CHECK_RET(ret == ACL_SUCCESS, return ret);
198+ ret = CreateAclTensor(dequantScaleHostData, dequantScaleShape, &dsDeviceAddr, aclDataType::ACL_FLOAT,
199+ &dequantScale);
200+ CHECK_RET(ret == ACL_SUCCESS, return ret);
201+ ret = CreateAclTensor(quantScaleHostData, quantScaleShape, &qsDeviceAddr, aclDataType::ACL_FLOAT, &quantScale);
202+ CHECK_RET(ret == ACL_SUCCESS, return ret);
203+ ret = CreateAclTensor(quantOffsetHostData, quantOffsetShape, &qoDeviceAddr, aclDataType::ACL_FLOAT, &quantOffset);
204+ CHECK_RET(ret == ACL_SUCCESS, return ret);
205+ ret = CreateAclTensor(yHostData, yShape, &yDeviceAddr, aclDataType::ACL_INT8, &y);
206+ CHECK_RET(ret == ACL_SUCCESS, return ret);
207+ ret = CreateAclTensor(yScaleOutHostData, yScaleOutShape, &yScaleDeviceAddr, aclDataType::ACL_FLOAT, &yScaleOut);
208+ CHECK_RET(ret == ACL_SUCCESS, return ret);
209+ 
210+ uint64_t workspaceSize = 0;
211+ aclOpExecutor* executor;
212+ ret = aclnnDequantSituQuantGetWorkspaceSize(x, dequantScale, nullptr, nullptr, quantScale, quantOffset, nullptr,
213+ beta, linearBeta, activateLeft, const_cast<char*>("static"), y,
214+ yScaleOut, &workspaceSize, &executor);
215+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("GetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
216+ 
217+ void* workspaceAddr = nullptr;
218+ if (workspaceSize > 0) {
219+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
220+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
221+ }
222+ 
223+ ret = aclnnDequantSituQuant(workspaceAddr, workspaceSize, executor, stream);
224+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDequantSituQuant failed. ERROR: %d\n", ret); return ret);
225+ 
226+ ret = aclrtSynchronizeStream(stream);
227+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
228+ 
229+ std::vector<int8_t> npuYResult(ySize, 0);
230+ ret = aclrtMemcpy(npuYResult.data(), npuYResult.size() * sizeof(int8_t), yDeviceAddr, ySize * sizeof(int8_t),
231+ ACL_MEMCPY_DEVICE_TO_HOST);
232+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy y result failed. ERROR: %d\n", ret); return ret);
233+ 
234+ std::vector<int8_t> cpuYResult(ySize, 0);
235+ CpuStaticQuant(xHostData, dequantScaleHostData, quantScaleHostData, quantOffsetHostData, rowLen, inDimy, outDimy,
236+ beta, linearBeta, activateLeft, cpuYResult);
237+ 
238+ int64_t matchCount = 0;
239+ int64_t closeCount = 0;
240+ int64_t mismatchCount = 0;
241+ int maxDiff = 0;
242+ for (int64_t i = 0; i < ySize; i++) {
243+ int diff = std::abs(static_cast<int>(npuYResult[i]) - static_cast<int>(cpuYResult[i]));
244+ if (diff == 0)
245+ matchCount++;
246+ if (diff <= 1)
247+ closeCount++;
248+ else
249+ mismatchCount++;
250+ maxDiff = std::max(maxDiff, diff);
251+ }
252+ 
253+ LOG_PRINT("Static quant: total=%ld, exact_match=%ld, within_1=%ld, mismatch=%ld, max_diff=%d\n", ySize, matchCount,
254+ closeCount, mismatchCount, maxDiff);
255+ LOG_PRINT("NPU y (first 16): ");
256+ for (int64_t i = 0; i < std::min<int64_t>(16, ySize); i++)
257+ LOG_PRINT("%d ", npuYResult[i]);
258+ LOG_PRINT("\nCPU y (first 16): ");
259+ for (int64_t i = 0; i < std::min<int64_t>(16, ySize); i++)
260+ LOG_PRINT("%d ", cpuYResult[i]);
261+ LOG_PRINT("\n");
262+ 
263+ aclDestroyTensor(x);
264+ aclDestroyTensor(dequantScale);
265+ aclDestroyTensor(quantScale);
266+ aclDestroyTensor(quantOffset);
267+ aclDestroyTensor(y);
268+ aclDestroyTensor(yScaleOut);
269+ aclrtFree(xDeviceAddr);
270+ aclrtFree(dsDeviceAddr);
271+ aclrtFree(qsDeviceAddr);
272+ aclrtFree(qoDeviceAddr);
273+ aclrtFree(yDeviceAddr);
274+ aclrtFree(yScaleDeviceAddr);
275+ if (workspaceSize > 0)
276+ aclrtFree(workspaceAddr);
277+ 
278+ bool passed = (maxDiff <= 2);
279+ LOG_PRINT("Test 1 %s (max_diff=%d, tolerance=2)\n", passed ? "PASSED" : "FAILED", maxDiff);
280+ return passed ? 0 : 1;
281+}
282+ 
283+// ============ Test 2: dynamic quant, no bias, no smooth ============
284+int TestDynamicQuant(int32_t deviceId, aclrtStream stream)
285+{
286+ LOG_PRINT("\n========== Test 2: Dynamic Quant (no bias, no smooth) ==========\n");
287+ 
288+ int64_t rowLen = 32;
289+ int64_t inDimy = 128;
290+ int64_t outDimy = 64;
291+ double beta = 1.0;
292+ double linearBeta = 0.0;
293+ bool activateLeft = false;
294+ 
295+ std::vector<int64_t> xShape = {rowLen, inDimy};
296+ std::vector<int64_t> dequantScaleShape = {inDimy};
297+ std::vector<int64_t> yShape = {rowLen, outDimy};
298+ std::vector<int64_t> yScaleOutShape = {rowLen};
299+ 
300+ auto xSize = GetShapeSize(xShape);
301+ auto dsSize = GetShapeSize(dequantScaleShape);
302+ auto ySize = GetShapeSize(yShape);
303+ auto yScaleSize = GetShapeSize(yScaleOutShape);
304+ 
305+ std::vector<int8_t> xHostData(xSize);
306+ for (int64_t i = 0; i < xSize; i++) {
307+ xHostData[i] = static_cast<int8_t>((i * 7 + 3) % 100);
308+ }
309+ std::vector<float> dequantScaleHostData(dsSize, 0.1f);
310+ std::vector<int8_t> yHostData(ySize, 0);
311+ std::vector<float> yScaleOutHostData(yScaleSize, 0.0f);
312+ 
313+ void* xDeviceAddr = nullptr;
314+ void* dsDeviceAddr = nullptr;
315+ void* yDeviceAddr = nullptr;
316+ void* yScaleDeviceAddr = nullptr;
317+ aclTensor* x = nullptr;
318+ aclTensor* dequantScale = nullptr;
319+ aclTensor* y = nullptr;
320+ aclTensor* yScaleOut = nullptr;
321+ 
322+ auto ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_INT8, &x);
323+ CHECK_RET(ret == ACL_SUCCESS, return ret);
324+ ret = CreateAclTensor(dequantScaleHostData, dequantScaleShape, &dsDeviceAddr, aclDataType::ACL_FLOAT,
325+ &dequantScale);
326+ CHECK_RET(ret == ACL_SUCCESS, return ret);
327+ ret = CreateAclTensor(yHostData, yShape, &yDeviceAddr, aclDataType::ACL_INT8, &y);
328+ CHECK_RET(ret == ACL_SUCCESS, return ret);
329+ ret = CreateAclTensor(yScaleOutHostData, yScaleOutShape, &yScaleDeviceAddr, aclDataType::ACL_FLOAT, &yScaleOut);
330+ CHECK_RET(ret == ACL_SUCCESS, return ret);
331+ 
332+ uint64_t workspaceSize = 0;
333+ aclOpExecutor* executor;
334+ ret = aclnnDequantSituQuantGetWorkspaceSize(x, dequantScale, nullptr, nullptr, nullptr, nullptr, nullptr, beta,
335+ linearBeta, activateLeft, const_cast<char*>("dynamic"), y, yScaleOut,
336+ &workspaceSize, &executor);
337+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("GetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
338+ 
339+ void* workspaceAddr = nullptr;
340+ if (workspaceSize > 0) {
341+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
342+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
343+ }
344+ 
345+ ret = aclnnDequantSituQuant(workspaceAddr, workspaceSize, executor, stream);
346+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDequantSituQuant failed. ERROR: %d\n", ret); return ret);
347+ 
348+ ret = aclrtSynchronizeStream(stream);
349+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
350+ 
351+ std::vector<int8_t> npuYResult(ySize, 0);
352+ ret = aclrtMemcpy(npuYResult.data(), npuYResult.size() * sizeof(int8_t), yDeviceAddr, ySize * sizeof(int8_t),
353+ ACL_MEMCPY_DEVICE_TO_HOST);
354+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy y result failed. ERROR: %d\n", ret); return ret);
355+ 
356+ std::vector<float> npuScaleResult(yScaleSize, 0.0f);
357+ ret = aclrtMemcpy(npuScaleResult.data(), npuScaleResult.size() * sizeof(float), yScaleDeviceAddr,
358+ yScaleSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
359+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy scale result failed. ERROR: %d\n", ret); return ret);
360+ 
361+ std::vector<int8_t> cpuYResult(ySize, 0);
362+ std::vector<float> cpuScaleResult(yScaleSize, 0.0f);
363+ CpuDynamicQuant(xHostData, dequantScaleHostData, rowLen, inDimy, outDimy, beta, linearBeta, activateLeft,
364+ cpuYResult, cpuScaleResult);
365+ 
366+ int64_t matchCount = 0;
367+ int64_t closeCount = 0;
368+ int maxDiff = 0;
369+ for (int64_t i = 0; i < ySize; i++) {
370+ int diff = std::abs(static_cast<int>(npuYResult[i]) - static_cast<int>(cpuYResult[i]));
371+ if (diff == 0)
372+ matchCount++;
373+ if (diff <= 1)
374+ closeCount++;
375+ maxDiff = std::max(maxDiff, diff);
376+ }
377+ 
378+ int64_t scaleMatch = 0;
379+ float maxScaleDiff = 0.0f;
380+ for (int64_t i = 0; i < yScaleSize; i++) {
381+ float diff = std::abs(npuScaleResult[i] - cpuScaleResult[i]);
382+ if (diff < 1e-3f)
383+ scaleMatch++;
384+ maxScaleDiff = std::max(maxScaleDiff, diff);
385+ }
386+ 
387+ LOG_PRINT("Dynamic quant y: total=%ld, exact_match=%ld, within_1=%ld, max_diff=%d\n", ySize, matchCount, closeCount,
388+ maxDiff);
389+ LOG_PRINT("Dynamic quant scale: total=%ld, match=%ld, max_diff=%f\n", yScaleSize, scaleMatch, maxScaleDiff);
390+ LOG_PRINT("NPU y (first 16): ");
391+ for (int64_t i = 0; i < std::min<int64_t>(16, ySize); i++)
392+ LOG_PRINT("%d ", npuYResult[i]);
393+ LOG_PRINT("\nCPU y (first 16): ");
394+ for (int64_t i = 0; i < std::min<int64_t>(16, ySize); i++)
395+ LOG_PRINT("%d ", cpuYResult[i]);
396+ LOG_PRINT("\n");
397+ LOG_PRINT("NPU scale (first 5): ");
398+ for (int64_t i = 0; i < std::min<int64_t>(5, yScaleSize); i++)
399+ LOG_PRINT("%f ", npuScaleResult[i]);
400+ LOG_PRINT("\nCPU scale (first 5): ");
401+ for (int64_t i = 0; i < std::min<int64_t>(5, yScaleSize); i++)
402+ LOG_PRINT("%f ", cpuScaleResult[i]);
403+ LOG_PRINT("\n");
404+ 
405+ aclDestroyTensor(x);
406+ aclDestroyTensor(dequantScale);
407+ aclDestroyTensor(y);
408+ aclDestroyTensor(yScaleOut);
409+ aclrtFree(xDeviceAddr);
410+ aclrtFree(dsDeviceAddr);
411+ aclrtFree(yDeviceAddr);
412+ aclrtFree(yScaleDeviceAddr);
413+ if (workspaceSize > 0)
414+ aclrtFree(workspaceAddr);
415+ 
416+ bool passed = (maxDiff <= 2) && (maxScaleDiff < 1e-2f);
417+ LOG_PRINT("Test 2 %s (max_y_diff=%d, max_scale_diff=%f)\n", passed ? "PASSED" : "FAILED", maxDiff, maxScaleDiff);
418+ return passed ? 0 : 1;
419+}
420+ 
421+int main()
422+{
423+ LOG_PRINT("=== DequantSituQuant Precision Test ===\n");
424+ 
425+ int32_t deviceId = 0;
426+ aclrtStream stream;
427+ auto ret = aclInit(nullptr);
428+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
429+ ret = aclrtSetDevice(deviceId);
430+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
431+ ret = aclrtCreateStream(&stream);
432+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
433+ 
434+ int ret1 = TestStaticQuantScalar(deviceId, stream);
435+ int ret2 = TestDynamicQuant(deviceId, stream);
436+ 
437+ LOG_PRINT("\n=== Summary ===\n");
438+ LOG_PRINT("Test 1 (Static Quant): %s\n", ret1 == 0 ? "PASSED" : "FAILED");
439+ LOG_PRINT("Test 2 (Dynamic Quant): %s\n", ret2 == 0 ? "PASSED" : "FAILED");
440+ 
441+ aclrtDestroyStream(stream);
442+ aclrtResetDevice(deviceId);
443+ aclFinalize();
444+ 
445+ if (ret1 == 0 && ret2 == 0) {
446+ LOG_PRINT("\nAll tests PASSED!\n");
447+ } else {
448+ LOG_PRINT("\nSome tests FAILED!\n");
449+ }
450+ return (ret1 == 0 && ret2 == 0) ? 0 : 1;
451+}
Aquant/dequant_situ_quant/op_graph/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 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+add_graph_plugin_sources()
Aquant/dequant_situ_quant/op_graph/dequant_situ_quant_proto.h+76-0
@@ -0,0 +1,76 @@
1+/**
2+ * Copyright (c) 2025 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 dequant_situ_quant_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_QUANT_DEQUANT_SITU_QUANT_PROTO_H_
16+#define OPS_QUANT_DEQUANT_SITU_QUANT_PROTO_H_
17+#include "graph/operator_reg.h"
18+ 
19+namespace ge {
20+ 
21+/**
22+* @brief Combine Dequant + Situ + Quant.
23+ 
24+* @par Inputs:
25+* Seven inputs:
26+* @li x: Required tensor. INT8 for per-channel dequant path; INT32 for MoE grouped-matmul accumulator;
27+* BF16 for pre-dequantized path; FLOAT16 for pre-dequantized path. Shape is (N..., H) for INT8
28+* (dim > 1, H even) or [rows, width] for INT32/BF16/FLOAT16.
29+* @li weight_scale: Optional FP32. Per-channel or per-expert dequantization scale. Required for INT32 x.
30+* @li activation_scale: Optional FP32. Per-row dequantization scale (one value per row). Required for INT32 x.
31+* @li bias: Optional FP32. Dequantization bias, same shape as weight_scale.
32+* @li quant_scale: Optional FP32. Static quant scale or dynamic smooth scale.
33+* @li quant_offset: Optional FP32. Static quant offset.
34+* @li group_index: Optional INT64 [experts]. Per-expert consecutive routed row counts for MoE.
35+ 
36+* @par Outputs:
37+* @li y: INT8 tensor. Last dim is x last dim / 2.
38+* @li y_scale: FP32 tensor. Per-row dynamic quant scale (meaningless for static mode).
39+ 
40+* @par Attributes:
41+* @li beta: Float. The beta parameter for Situ activation. Default is 4.0.
42+* @li linear_beta: Float. The linear_beta parameter for Situ activation. When value <= 0, the linear_beta
43+* transformation is not applied. Default is 25.0.
44+* @li activate_left: Bool. Whether gate is the left half (true) or right half (false). Default is true.
45+* @li quant_type: String. The quant type to use: 'static' or 'dynamic', default is 'dynamic'.
46+ 
47+* @attention Constraints:
48+* @li The last dimension of x must be even.
49+* @li INT8 path: weight_scale required, activation_scale/group_index must be absent, x dim > 1.
50+* @li INT32 path: weight_scale and activation_scale required, quant_scale/quant_offset must be absent,
51+* x rank == 2, quant_type must be dynamic.
52+* @li BF16 path: weight_scale/activation_scale/bias/group_index must be absent, x rank == 2,
53+* quant_type must be dynamic.
54+* @li FLOAT16 path: weight_scale/activation_scale/bias/group_index must be absent, x rank == 2,
55+* quant_type must be dynamic.
56+* @li When quant_type is 'static', quant_scale must be provided.
57+* @li When quant_type is 'dynamic', quant_scale is optional (used as smoothScale).
58+*/
59+REG_OP(DequantSituQuant)
60+ .INPUT(x, TensorType({DT_INT8, DT_INT32, DT_BF16, DT_FLOAT16}))
61+ .OPTIONAL_INPUT(weight_scale, TensorType({DT_FLOAT, DT_FLOAT, DT_FLOAT, DT_FLOAT}))
62+ .OPTIONAL_INPUT(activation_scale, TensorType({DT_FLOAT, DT_FLOAT, DT_FLOAT, DT_FLOAT}))
63+ .OPTIONAL_INPUT(bias, TensorType({DT_FLOAT, DT_FLOAT, DT_FLOAT, DT_FLOAT}))
64+ .OPTIONAL_INPUT(quant_scale, TensorType({DT_FLOAT, DT_FLOAT, DT_FLOAT, DT_FLOAT}))
65+ .OPTIONAL_INPUT(quant_offset, TensorType({DT_FLOAT, DT_FLOAT, DT_FLOAT, DT_FLOAT}))
66+ .OPTIONAL_INPUT(group_index, TensorType({DT_INT64, DT_INT64, DT_INT64, DT_INT64}))
67+ .OUTPUT(y, TensorType({DT_INT8, DT_INT8, DT_INT8, DT_INT8}))
68+ .OUTPUT(y_scale, TensorType({DT_FLOAT, DT_FLOAT, DT_FLOAT, DT_FLOAT}))
69+ .ATTR(beta, Float, 4.0)
70+ .ATTR(linear_beta, Float, 25.0)
71+ .ATTR(activate_left, Bool, true)
72+ .ATTR(quant_type, String, "dynamic")
73+ .OP_END_FACTORY_REG(DequantSituQuant)
74+} // namespace ge
75+ 
76+#endif // OPS_QUANT_DEQUANT_SITU_QUANT_PROTO_H_
Aquant/dequant_situ_quant/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 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+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE dequant_situ_quant ACLNNTYPE aclnn)
Aquant/dequant_situ_quant/op_host/dequant_situ_quant_def.cpp+81-0
@@ -0,0 +1,81 @@
1+/**
2+ * Copyright (c) 2025 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 dequant_situ_quant_def.cpp
13+ * \brief
14+ */
15+ 
16+#include <cstdint>
17+#include "register/op_def_registry.h"
18+ 
19+namespace ops {
20+class DequantSituQuant : public OpDef {
21+public:
22+ explicit DequantSituQuant(const char* name) : OpDef(name)
23+ {
24+ this->Input("x")
25+ .ParamType(REQUIRED)
26+ .DataType({ge::DT_INT8, ge::DT_INT32, ge::DT_BF16, ge::DT_FLOAT16})
27+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
28+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
29+ this->Input("weight_scale")
30+ .ParamType(OPTIONAL)
31+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
32+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
33+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
34+ this->Input("activation_scale")
35+ .ParamType(OPTIONAL)
36+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
37+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
38+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
39+ this->Input("bias")
40+ .ParamType(OPTIONAL)
41+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
42+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
43+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
44+ this->Input("quant_scale")
45+ .ParamType(OPTIONAL)
46+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
47+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
48+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
49+ this->Input("quant_offset")
50+ .ParamType(OPTIONAL)
51+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
52+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
53+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
54+ this->Input("group_index")
55+ .ParamType(OPTIONAL)
56+ .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
57+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
58+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
59+ this->Output("y")
60+ .ParamType(REQUIRED)
61+ .DataType({ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8})
62+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
63+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
64+ this->Output("y_scale")
65+ .ParamType(REQUIRED)
66+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
67+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
68+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
69+ 
70+ this->Attr("beta").AttrType(OPTIONAL).Float(4.0);
71+ this->Attr("linear_beta").AttrType(OPTIONAL).Float(25.0);
72+ this->Attr("activate_left").AttrType(OPTIONAL).Bool(true);
73+ this->Attr("quant_type").AttrType(OPTIONAL).String("dynamic");
74+ 
75+ this->AICore().AddConfig("ascend910b");
76+ this->AICore().AddConfig("ascend910_93");
77+ }
78+};
79+ 
80+OP_ADD(DequantSituQuant);
81+} // namespace ops
Aquant/dequant_situ_quant/op_host/dequant_situ_quant_infershape.cpp+90-0
@@ -0,0 +1,90 @@
1+/**
2+ * Copyright (c) 2025 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 dequant_situ_quant_infershape.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "graph/utils/type_utils.h"
17+#include "util/shape_util.h"
18+#include "log/log.h"
19+#include "util/math_util.h"
20+ 
21+using namespace ge;
22+namespace ops {
23+constexpr size_t INPUT_IDX_X = 0;
24+constexpr size_t OUTPUT_IDX_Y = 0;
25+constexpr size_t OUTPUT_IDX_Y_SCALE = 1;
26+constexpr int64_t CONST_UNKNOW_SHAPE = -1;
27+constexpr int64_t NUM_TWO = 2;
28+ 
29+graphStatus InferShape4DequantSituQuant(gert::InferShapeContext* context)
30+{
31+ OP_LOGD(context, "Begin to do InferShape4DequantSituQuant.");
32+ 
33+ const gert::Shape* xShape = context->GetInputShape(INPUT_IDX_X);
34+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
35+ gert::Shape* yShape = context->GetOutputShape(OUTPUT_IDX_Y);
36+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
37+ gert::Shape* yScaleShape = context->GetOutputShape(OUTPUT_IDX_Y_SCALE);
38+ OP_CHECK_NULL_WITH_CONTEXT(context, yScaleShape);
39+ 
40+ *yShape = *xShape;
41+ OP_CHECK_IF(Ops::Base::IsUnknownRank(*xShape),
42+ OP_LOGD(context, "End to do InferShape4DequantSituQuant, inputx is [-2]."), return GRAPH_SUCCESS);
43+ 
44+ int64_t xShapeRank = static_cast<int64_t>(xShape->GetDimNum());
45+ 
46+ auto inputDesc = context->GetInputDesc(INPUT_IDX_X);
47+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
48+ ge::DataType xDataType = inputDesc->GetDataType();
49+ 
50+ bool isInt8 = (xDataType == ge::DT_INT8);
51+ if (isInt8) {
52+ OP_CHECK_IF(xShapeRank <= 1, OP_LOGE(context, "x shape rank must > 1 for int8, but is %ld", xShapeRank),
53+ return ge::GRAPH_FAILED);
54+ } else {
55+ OP_CHECK_IF(xShapeRank != 2,
56+ OP_LOGE(context, "x shape rank must be 2 for int32/bfloat16/float16, but is %ld", xShapeRank),
57+ return ge::GRAPH_FAILED);
58+ }
59+ 
60+ int64_t lastDim = xShape->GetDim(xShapeRank - 1);
61+ int64_t outLastDim = lastDim == CONST_UNKNOW_SHAPE ? CONST_UNKNOW_SHAPE : lastDim / NUM_TWO;
62+ OP_CHECK_IF((lastDim != CONST_UNKNOW_SHAPE) && (lastDim % NUM_TWO != 0),
63+ OP_LOGE(context, "The last dim of x must be even number, but is %ld", lastDim),
64+ return ge::GRAPH_FAILED);
65+ 
66+ yShape->SetDim(xShapeRank - 1, outLastDim);
67+ 
68+ *yScaleShape = *yShape;
69+ if (isInt8) {
70+ yScaleShape->SetDimNum(xShapeRank - 1);
71+ } else {
72+ yScaleShape->SetDimNum(1);
73+ yScaleShape->SetDim(0, xShape->GetDim(0));
74+ }
75+ 
76+ OP_LOGD(context, "End to do InferShape4DequantSituQuant");
77+ return GRAPH_SUCCESS;
78+}
79+ 
80+graphStatus InferDtype4DequantSituQuant(gert::InferDataTypeContext* context)
81+{
82+ OP_LOGD(context, "InferDtype4DequantSituQuant enter");
83+ context->SetOutputDataType(OUTPUT_IDX_Y, DT_INT8);
84+ context->SetOutputDataType(OUTPUT_IDX_Y_SCALE, DT_FLOAT);
85+ OP_LOGD(context, "InferDtype4DequantSituQuant end");
86+ return GRAPH_SUCCESS;
87+}
88+ 
89+IMPL_OP_INFERSHAPE(DequantSituQuant).InferShape(InferShape4DequantSituQuant).InferDataType(InferDtype4DequantSituQuant);
90+} // namespace ops
Aquant/dequant_situ_quant/op_host/dequant_situ_quant_tiling.cpp+686-0
@@ -0,0 +1,686 @@
1+/**
2+ * Copyright (c) 2025 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 dequant_situ_quant_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include <cmath>
17+#include <cstdint>
18+#include <algorithm>
19+#include <limits>
20+#include <string>
21+#include "tiling/tiling_api.h"
22+#include "dequant_situ_quant_tiling.h"
23+#include "op_host/tiling_util.h"
24+#include "op_host/tiling_templates_registry.h"
25+ 
26+#define CHECK_FAIL(cont, cond, ...) \
27+ do { \
28+ if (cond) { \
29+ OP_LOGE(cont->GetNodeName(), ##__VA_ARGS__); \
30+ return ge::GRAPH_FAILED; \
31+ } \
32+ } while (0)
33+ 
34+namespace optiling {
35+using Ops::NN::Optiling::TilingBaseClass;
36+using Ops::NN::Optiling::TilingRegistry;
37+ 
38+constexpr uint32_t UB_RESERVED_BUFF = 1024;
39+constexpr uint32_t ALIGN_UINT_IN_CACHE_32B = 32;
40+constexpr uint32_t PACK_UINT_IN_CACHE_512B = 512;
41+constexpr uint32_t DEFAULT_BUFFER_NUM = 2;
42+constexpr uint32_t MAX_CORE_NUMBER = 64;
43+constexpr uint32_t PERFORMANCE_COL_LEN = 1536;
44+constexpr uint32_t PERFORMANCE_ROW_LEN = 128;
45+constexpr uint32_t MIN_CORE = 12;
46+constexpr uint64_t USER_WORKSPACE = 16777216; // 16 * 1024 * 1024
47+constexpr uint64_t BLOCK_BYTES = 32;
48+ 
49+constexpr size_t INDEX_IN_X = 0;
50+constexpr size_t INDEX_IN_WEIGHT_SCALE = 1;
51+constexpr size_t INDEX_IN_ACTIVATION_SCALE = 2;
52+constexpr size_t INDEX_IN_BIAS = 3;
53+constexpr size_t INDEX_IN_QUANT_SCALE = 4;
54+constexpr size_t INDEX_IN_QUANT_OFFSET = 5;
55+constexpr size_t INDEX_IN_GROUP_INDEX = 6;
56+constexpr int64_t NUM_TWO = 2;
57+constexpr int64_t QUANT_TYPE_STATIC = 0;
58+constexpr int64_t QUANT_TYPE_DYNAMIC = 1;
59+ 
60+void DequantSituQuantTiling::Reset()
61+{
62+ opName = nullptr;
63+ totalCore = 0;
64+ totalUsedCoreNum = 0;
65+ inputDTypeLen = 1;
66+ ubMinBlockLen = 32;
67+ cacheLineLen = 512;
68+ maxTileLen = 0;
69+ optBaseRowLen = 0;
70+ optBaseColLen = 0;
71+ workspaceSize_ = 0;
72+ hasDequantBias = false;
73+ hasQuantScale = false;
74+ hasQuantOffset = false;
75+ quantIsOne = false;
76+ activateLeft = 0;
77+ quantType = 0;
78+ beta = 4.0f;
79+ linearBeta = 25.0f;
80+ quantScaleShapeSize = 0;
81+ inDimx = 0;
82+ inDimy = 0;
83+ outDimy = 0;
84+ xDtype_ = ge::DT_INT8;
85+ isPreDequantized_ = false;
86+ hasWeightScale_ = false;
87+ hasActivationScale_ = false;
88+ hasGroupIndex_ = false;
89+ expertNum_ = 1;
90+}
91+ 
92+bool DequantSituQuantTiling::IsCapable()
93+{
94+ if (Ops::NN::OpTiling::IsRegbaseSocVersion(context_)) {
95+ return false;
96+ }
97+ return true;
98+}
99+ 
100+ge::graphStatus DequantSituQuantTiling::GetPlatformInfo()
Aquant/dequant_situ_quant/op_host/dequant_situ_quant_tiling.h+150-0
@@ -0,0 +1,150 @@
1+/**
2+ * Copyright (c) 2025 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 dequant_situ_quant_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef DEQUANT_SITU_QUANT_TILING_H
17+#define DEQUANT_SITU_QUANT_TILING_H
18+ 
19+#include <vector>
20+#include <iostream>
21+#include "register/op_impl_registry.h"
22+#include "util/math_util.h"
23+#include "log/log.h"
24+#include "tiling/platform/platform_ascendc.h"
25+#include "platform/platform_infos_def.h"
26+#include "register/tilingdata_base.h"
27+#include "tiling/tiling_api.h"
28+#include "../op_graph/dequant_situ_quant_proto.h"
29+#include "op_host/tiling_base.h"
30+#include "op_host/tiling_templates_registry.h"
31+ 
32+using namespace Ops::NN::Optiling;
33+namespace optiling {
34+ 
35+BEGIN_TILING_DATA_DEF(DequantSituQuantTilingData)
36+TILING_DATA_FIELD_DEF(uint32_t, is32BAligned);
37+TILING_DATA_FIELD_DEF(uint32_t, isDoubleBuffer);
38+TILING_DATA_FIELD_DEF(uint64_t, rowLen);
39+TILING_DATA_FIELD_DEF(uint64_t, colLen);
40+TILING_DATA_FIELD_DEF(uint32_t, baseRowLen);
41+TILING_DATA_FIELD_DEF(uint32_t, baseColLen);
42+TILING_DATA_FIELD_DEF(uint32_t, activateLeft);
43+TILING_DATA_FIELD_DEF(uint32_t, dequantBiasIsEmpty);
44+TILING_DATA_FIELD_DEF(uint32_t, quantScaleIsEmpty);
45+TILING_DATA_FIELD_DEF(uint32_t, quantOffsetIsEmpty);
46+TILING_DATA_FIELD_DEF(uint32_t, quantIsOne);
47+TILING_DATA_FIELD_DEF(uint32_t, usedCoreNum);
48+TILING_DATA_FIELD_DEF(int64_t, quantType);
49+TILING_DATA_FIELD_DEF(float, beta);
50+TILING_DATA_FIELD_DEF(float, linearBeta);
51+TILING_DATA_FIELD_DEF(uint32_t, expertNum);
52+TILING_DATA_FIELD_DEF(uint32_t, hasGroupIndex);
53+TILING_DATA_FIELD_DEF(uint32_t, hasActivationScale);
54+TILING_DATA_FIELD_DEF(uint32_t, isPreDequantized);
55+TILING_DATA_FIELD_DEF(uint32_t, inputWidth);
56+TILING_DATA_FIELD_DEF(uint32_t, outputWidth);
57+END_TILING_DATA_DEF;
58+ 
59+REGISTER_TILING_DATA_CLASS(DequantSituQuant, DequantSituQuantTilingData)
60+ 
61+struct DequantSituQuantCompileInfo {
62+ uint64_t coreNum = 0;
63+ uint64_t ubSize = 0;
64+};
65+ 
66+constexpr int64_t DSQ_STATIC_QUANT_ONE = 10000;
67+constexpr int64_t DSQ_STATIC_QUANT_VEC = 10001;
68+constexpr int64_t DSQ_STATIC_QUANT_ONE_BIAS = 10002;
69+constexpr int64_t DSQ_STATIC_QUANT_VEC_BIAS = 10003;
70+constexpr int64_t DSQ_DYNAMIC_QUANT_NO_SMOOTH = 20000;
71+constexpr int64_t DSQ_DYNAMIC_QUANT_SMOOTH = 20001;
72+constexpr int64_t DSQ_DYNAMIC_QUANT_NO_SMOOTH_BIAS = 20002;
73+constexpr int64_t DSQ_DYNAMIC_QUANT_SMOOTH_BIAS = 20003;
74+constexpr int64_t DSQ_INT32_DYNAMIC = 30000;
75+constexpr int64_t DSQ_BF16_DYNAMIC = 40000;
76+constexpr int64_t DSQ_FP16_DYNAMIC = 50000;
77+ 
78+class DequantSituQuantTiling : public Ops::NN::Optiling::TilingBaseClass {
79+public:
80+ explicit DequantSituQuantTiling(gert::TilingContext* cont) : TilingBaseClass(cont) { Reset(); }
81+ ~DequantSituQuantTiling() override = default;
82+ 
83+ void Reset(gert::TilingContext* cont) override
84+ {
85+ TilingBaseClass::Reset(cont);
86+ Reset();
87+ }
88+ 
89+protected:
90+ bool IsCapable() override;
91+ ge::graphStatus GetPlatformInfo() override;
92+ ge::graphStatus GetShapeAttrsInfo() override;
93+ ge::graphStatus DoOpTiling() override;
94+ ge::graphStatus DoLibApiTiling() override;
95+ uint64_t GetTilingKey() const override;
96+ ge::graphStatus GetWorkspaceSize() override;
97+ ge::graphStatus PostTiling() override;
98+ void Reset();
99+ 
100+private:
101+ void ShowTilingData();
102+ ge::graphStatus GetShapeAttrsInfoInner();
103+ ge::graphStatus CheckInputShapes();
104+ ge::graphStatus CheckInputShapesInt8(int64_t dimNum, int64_t inDimy, int64_t outDimy);
105+ ge::graphStatus CheckInputShapesInt32(int64_t inDimy, int64_t outDimy);
106+ ge::graphStatus CheckInputShapesBF16();
107+ bool SetAttrs(const gert::RuntimeAttrs* attrs);
108+ bool CalcTiling(const uint32_t totalCores, const uint64_t ubSize);
109+ bool CalcUbMaxTileLen(const uint64_t ubSize, uint32_t& maxTileLen);
110+ bool CalcOptBaseShape(uint32_t maxTileLen);
111+ ge::graphStatus ValidateInt32Contract();
112+ 
113+ const char* opName = "";
114+ uint32_t totalCore = 0;
115+ uint32_t totalUsedCoreNum = 0;
116+ uint32_t inputDTypeLen = 1;
117+ uint32_t ubMinBlockLen = 32;
118+ uint32_t cacheLineLen = 512;
119+ uint32_t maxTileLen = 0;
120+ uint32_t optBaseRowLen = 0;
121+ uint32_t optBaseColLen = 0;
122+ uint64_t workspaceSize_ = 0;
123+ 
124+ bool hasDequantBias = false;
125+ bool hasQuantScale = false;
126+ bool hasQuantOffset = false;
127+ bool quantIsOne = false;
128+ uint32_t activateLeft = 0;
129+ int64_t quantType = 0;
130+ float beta = 4.0f;
131+ float linearBeta = 25.0f;
132+ uint64_t quantScaleShapeSize = 0;
133+ 
134+ int64_t inDimx = 0;
135+ int64_t inDimy = 0;
136+ int64_t outDimy = 0;
137+ 
138+ ge::DataType xDtype_ = ge::DT_INT8;
139+ bool isPreDequantized_ = false;
140+ bool hasWeightScale_ = false;
141+ bool hasActivationScale_ = false;
142+ bool hasGroupIndex_ = false;
143+ uint32_t expertNum_ = 1;
144+ 
145+ DequantSituQuantTilingData tilingData;
146+ platform_ascendc::SocVersion socVersion = platform_ascendc::SocVersion::ASCEND910B;
147+};
148+ 
149+} // namespace optiling
150+#endif // DEQUANT_SITU_QUANT_TILING_H
Aquant/dequant_situ_quant/op_kernel/dequant_situ_quant.cpp+113-0
@@ -0,0 +1,113 @@
1+/**
2+ * Copyright (c) 2025 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 dequant_situ_quant.cpp
13+ * \brief Kernel entry point for DequantSituQuant operator.
14+ */
15+ 
16+#include "kernel_operator.h"
17+#include "kernel_tiling/kernel_tiling.h"
18+#include "dequant_situ_quant.h"
19+ 
20+using namespace AscendC;
21+ 
22+#define DSQ_STATIC_QUANT_ONE 10000
23+#define DSQ_STATIC_QUANT_VEC 10001
24+#define DSQ_STATIC_QUANT_ONE_BIAS 10002
25+#define DSQ_STATIC_QUANT_VEC_BIAS 10003
26+#define DSQ_DYNAMIC_QUANT_NO_SMOOTH 20000
27+#define DSQ_DYNAMIC_QUANT_SMOOTH 20001
28+#define DSQ_DYNAMIC_QUANT_NO_SMOOTH_BIAS 20002
29+#define DSQ_DYNAMIC_QUANT_SMOOTH_BIAS 20003
30+#define DSQ_INT32_DYNAMIC 30000
31+#define DSQ_BF16_DYNAMIC 40000
32+#define DSQ_FP16_DYNAMIC 50000
33+ 
34+extern "C" __global__ __aicore__ void dequant_situ_quant(GM_ADDR xGM, GM_ADDR weightScaleGM, GM_ADDR activationScaleGM,
35+ GM_ADDR biasGM, GM_ADDR quantScaleGM, GM_ADDR quantOffsetGM,
36+ GM_ADDR groupIndexGM, GM_ADDR yGM, GM_ADDR yScaleGM,
37+ GM_ADDR workspace, GM_ADDR tiling)
38+{
39+#if (ORIG_DTYPE_X == DT_INT8)
40+ if (workspace == nullptr) {
41+ return;
42+ }
43+ GM_ADDR userspace = GetUserWorkspace(workspace);
44+ if (userspace == nullptr) {
45+ return;
46+ }
47+ 
48+ TPipe pipe;
49+ GET_TILING_DATA_WITH_STRUCT(DequantSituQuantTilingData, tilingDataIn, tiling);
50+ const DequantSituQuantTilingData* __restrict__ tilingData = &tilingDataIn;
51+ 
52+ if (TILING_KEY_IS(DSQ_STATIC_QUANT_ONE)) {
53+ DequantSituQuantOps::DequantSituQuantKernel<false> op(&pipe);
54+ op.Init(xGM, weightScaleGM, biasGM, quantScaleGM, quantOffsetGM, yGM, yScaleGM, userspace, tilingData);
55+ op.Process();
56+ } else if (TILING_KEY_IS(DSQ_STATIC_QUANT_VEC)) {
57+ DequantSituQuantOps::DequantSituQuantKernel<false> op(&pipe);
58+ op.Init(xGM, weightScaleGM, biasGM, quantScaleGM, quantOffsetGM, yGM, yScaleGM, userspace, tilingData);
59+ op.Process();
60+ } else if (TILING_KEY_IS(DSQ_STATIC_QUANT_ONE_BIAS)) {
61+ DequantSituQuantOps::DequantSituQuantKernel<true> op(&pipe);
62+ op.Init(xGM, weightScaleGM, biasGM, quantScaleGM, quantOffsetGM, yGM, yScaleGM, userspace, tilingData);
63+ op.Process();
64+ } else if (TILING_KEY_IS(DSQ_STATIC_QUANT_VEC_BIAS)) {
65+ DequantSituQuantOps::DequantSituQuantKernel<true> op(&pipe);
66+ op.Init(xGM, weightScaleGM, biasGM, quantScaleGM, quantOffsetGM, yGM, yScaleGM, userspace, tilingData);
67+ op.Process();
68+ } else if (TILING_KEY_IS(DSQ_DYNAMIC_QUANT_NO_SMOOTH)) {
69+ DequantSituQuantOps::DequantSituQuantKernel<false> op(&pipe);
70+ op.Init(xGM, weightScaleGM, biasGM, quantScaleGM, quantOffsetGM, yGM, yScaleGM, userspace, tilingData);
71+ op.Process();
72+ } else if (TILING_KEY_IS(DSQ_DYNAMIC_QUANT_SMOOTH)) {
73+ DequantSituQuantOps::DequantSituQuantKernel<false> op(&pipe);
74+ op.Init(xGM, weightScaleGM, biasGM, quantScaleGM, quantOffsetGM, yGM, yScaleGM, userspace, tilingData);
75+ op.Process();
76+ } else if (TILING_KEY_IS(DSQ_DYNAMIC_QUANT_NO_SMOOTH_BIAS)) {
77+ DequantSituQuantOps::DequantSituQuantKernel<true> op(&pipe);
78+ op.Init(xGM, weightScaleGM, biasGM, quantScaleGM, quantOffsetGM, yGM, yScaleGM, userspace, tilingData);
79+ op.Process();
80+ } else if (TILING_KEY_IS(DSQ_DYNAMIC_QUANT_SMOOTH_BIAS)) {
81+ DequantSituQuantOps::DequantSituQuantKernel<true> op(&pipe);
82+ op.Init(xGM, weightScaleGM, biasGM, quantScaleGM, quantOffsetGM, yGM, yScaleGM, userspace, tilingData);
83+ op.Process();
84+ }
85+#elif (ORIG_DTYPE_X == DT_INT32)
86+ if (TILING_KEY_IS(DSQ_INT32_DYNAMIC)) {
87+ TPipe pipe;
88+ GET_TILING_DATA_WITH_STRUCT(DequantSituQuantTilingData, tilingDataIn, tiling);
89+ const DequantSituQuantTilingData* __restrict__ tilingData = &tilingDataIn;
90+ DequantSituQuantOps::DequantSituQuantK3Kernel<int32_t> op(&pipe);
91+ op.Init(xGM, weightScaleGM, activationScaleGM, biasGM, groupIndexGM, yGM, yScaleGM, tilingData);
92+ op.Process();
93+ }
94+#elif (ORIG_DTYPE_X == DT_BF16)
95+ if (TILING_KEY_IS(DSQ_BF16_DYNAMIC)) {
96+ TPipe pipe;
97+ GET_TILING_DATA_WITH_STRUCT(DequantSituQuantTilingData, tilingDataIn, tiling);
98+ const DequantSituQuantTilingData* __restrict__ tilingData = &tilingDataIn;
99+ DequantSituQuantOps::DequantSituQuantK3Kernel<bfloat16_t> op(&pipe);
100+ op.Init(xGM, weightScaleGM, activationScaleGM, biasGM, groupIndexGM, yGM, yScaleGM, tilingData);
101+ op.Process();
102+ }
103+#elif (ORIG_DTYPE_X == DT_FLOAT16)
104+ if (TILING_KEY_IS(DSQ_FP16_DYNAMIC)) {
105+ TPipe pipe;
106+ GET_TILING_DATA_WITH_STRUCT(DequantSituQuantTilingData, tilingDataIn, tiling);
107+ const DequantSituQuantTilingData* __restrict__ tilingData = &tilingDataIn;
108+ DequantSituQuantOps::DequantSituQuantK3Kernel<half> op(&pipe);
109+ op.Init(xGM, weightScaleGM, activationScaleGM, biasGM, groupIndexGM, yGM, yScaleGM, tilingData);
110+ op.Process();
111+ }
112+#endif
113+}
Aquant/dequant_situ_quant/tests/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(ENABLE_TEST)
13+ foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+ endforeach()
18+endif()
Aquant/dequant_situ_quant/tests/ut/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aquant/dequant_situ_quant/tests/ut/op_host/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_HOST_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
15+endif()
16+ 
17+if(UT_TEST_ALL OR OP_API_UT)
18+ add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
19+endif()
Atorch_extension/cann_ops_nn/ops/quant/__init__.py+14-0
Atorch_extension/cann_ops_nn/ops/quant/dequant_situ_quant/__init__.py+14-0
Atorch_extension/cann_ops_nn/ops/quant/situ_mx_quant/__init__.py+14-0