Star873
1771
代码
代码
Issues224
Pull Requests304
Actions
项目讨论
Wiki
分析
Star873
1771
已合并
FakeQuantAffineCachemask算子zero_point参数支持float16/32 #780
jiangjiawei创建于 1月17日
FakeQuantAffineCachemask算子zero_point参数支持float16/32 #780
已合并
jiangjiawei创建于 1月17日
从jiangjiawei/ops-nnsync_fake_quant合入到cann/ops-nnmaster
共 14 个文件变更+949-81
Mquant/fake_quant_affine_cachemask/README.md+5-1
@@ -4,8 +4,12 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| ---- | :----:|6| ---- | :----:|
7+|Ascend 950PR/Ascend 950DT|×|
7|Atlas A3 训练系列产品/Atlas A3 推理系列产品|√|8|Atlas A3 训练系列产品/Atlas A3 推理系列产品|√|
8-|Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件|√|9+|Atlas A2 训练系列产品/Atlas A2 推理系列产品|√|
10+|Atlas 200I/500 A2推理产品|×|
11+|Atlas 推理系列产品|×|
12+|Atlas 训练系列产品|×|
9 13 
10## 功能说明14## 功能说明
11 15 
Mquant/fake_quant_affine_cachemask/docs/aclnnFakeQuantPerChannelAffineCachemask.md+235-30
@@ -4,17 +4,26 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |6| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | × |
7| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |8| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |9| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | × |
9 13 
10## 功能说明14## 功能说明
11 15 
12-- 算子功能：对于输入数据self，使用scale和zero_point对输入self在指定轴axis上进行伪量化处理，并根据quant_min和quant_max对伪量化输出进行值域更新，最终返回结果out及对应位置掩码mask。16+- 接口功能：对于输入数据self，使用scale和zero_point对输入self在指定轴axis上进行伪量化处理，并根据quant_min和quant_max对伪量化输出进行值域更新，最终返回结果out及对应位置掩码mask。
13- 计算公式：根据算子功能先计算临时变量qval，再计算得出out和mask。17- 计算公式：根据算子功能先计算临时变量qval，再计算得出out和mask。
14 18 
19+ - zero\_point的数据类型为int32
15 $$20 $$
16 qval = Round(std::nearby\_int(self / scale) + zero\_point)21 qval = Round(std::nearby\_int(self / scale) + zero\_point)
17 $$22 $$
23+ - zero\_point的数据类型为float16或float32
24+ $$
25+ qval = std::lrintf(self / scale + zero\_point)
26+ $$
18 27 
19 $$28 $$
20 out = (Min(quant\_max, Max(quant\_min, qval)) - zero\_point) * scale29 out = (Min(quant\_max, Max(quant\_min, qval)) - zero\_point) * scale
@@ -28,47 +37,242 @@
28 37 
29每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnFakeQuantPerChannelAffineCachemaskGetWorkspaceSize”接口获取入参并根据流程计算所需workspace大小，再调用“aclnnFakeQuantPerChannelAffineCachemask”接口执行计算。38每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnFakeQuantPerChannelAffineCachemaskGetWorkspaceSize”接口获取入参并根据流程计算所需workspace大小，再调用“aclnnFakeQuantPerChannelAffineCachemask”接口执行计算。
30 39 
31-- `aclnnStatus aclnnFakeQuantPerChannelAffineCachemaskGetWorkspaceSize(const aclTensor* self, const aclTensor* scale, const aclTensor* zeroPoint, int64_t axis, int64_t quantMin, int64_t quantMax, aclTensor* out, aclTensor* mask, uint64_t* workspaceSize, aclOpExecutor** executor)`40+```Cpp
32-- `aclnnStatus aclnnFakeQuantPerChannelAffineCachemask(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)`41+aclnnStatus aclnnFakeQuantPerChannelAffineCachemaskGetWorkspaceSize(
42+ const aclTensor *self,
43+ const aclTensor *scale,
44+ const aclTensor *zeroPoint,
45+ int64_t axis,
46+ int64_t quantMin,
47+ int64_t quantMax,
48+ aclTensor *out,
49+ aclTensor *mask,
50+ uint64_t *workspaceSize,
51+ aclOpExecutor **executor)
52+```
53+```Cpp
54+aclnnStatus aclnnFakeQuantPerChannelAffineCachemask(
55+ void *workspace,
56+ uint64_t workspaceSize,
57+ aclOpExecutor *executor,
58+ aclrtStream stream)
59+```
33 60 
34## aclnnFakeQuantPerChannelAffineCachemaskGetWorkspaceSize61## aclnnFakeQuantPerChannelAffineCachemaskGetWorkspaceSize
35 62 
36- **参数说明：**63- **参数说明：**
37 64 
38- - self(aclTensor\*, 计算输入)：公式中的`self`，Device侧的aclTensor，数据类型支持FLOAT16、FLOAT32。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。65+ <table style="undefined;table-layout: fixed; width: 1478px"><colgroup>
39- - scale(aclTensor\*, 计算输入)：公式中的`scale`，Device侧的aclTensor，表示输入伪量化的缩放系数。数据类型支持FLOAT16、FLOAT32，shape只支持1维，size需要与zeroPoint一致且需要与self在axis轴的size一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。66+ <col style="width: 149px">
40- - zeroPoint(aclTensor\*, 计算输入)：公式中的`zero_point`，Device侧的aclTensor，表示输入伪量化的零基准参数。数据类型支持INT32，shape只支持1维，size需要与scale一致且需要与self在axis轴的size一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。67+ <col style="width: 121px">
41- - axis(int64_t, 计算输入)：Host侧的整型，表示计算维度，范围为[-self.dim(), self.dim()-1]。68+ <col style="width: 304px">
42- - quantMin(int64_t, 计算输入)：Host侧的整型，表示输入数据伪量化后的最小值，需要小于等于quantMax。69+ <col style="width: 253px">
43- - quantMax(int64_t, 计算输入)：Host侧的整型，表示输入数据伪量化后的最大值，需要大于等于quantMin。70+ <col style="width: 222px">
44- - out(aclTensor\*, 计算输出)：公式中的`out`，Device侧的aclTensor，数据类型支持FLOAT16、FLOAT32，shape需要和计算输入`self`一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。71+ <col style="width: 148px">
45- - mask(aclTensor\*, 计算输出)：公式中的`mask`，Device侧的aclTensor，数据类型支持BOOL，shape需要和计算输入`self`一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。72+ <col style="width: 135px">
46- - workspaceSize(uint64_t\*, 出参)：返回需要在Device侧申请的workspace大小。73+ <col style="width: 146px">
47- - executor(aclOpExecutor\*\*, 出参)：返回op执行器，包含了算子计算流程。74+ </colgroup>
75+ <thead>
76+ <tr>
77+ <th>参数名</th>
78+ <th>输入/输出</th>
79+ <th>描述</th>
80+ <th>使用说明</th>
81+ <th>数据类型</th>
82+ <th>数据格式</th>
83+ <th>维度(shape)</th>
84+ <th>非连续Tensor</th>
85+ </tr></thead>
86+ <tbody>
87+ <tr>
88+ <td>self</td>
89+ <td>输入</td>
90+ <td>公式中的self。</td>
91+ <td>-</td>
92+ <td>FLOAT16、FLOAT32</td>
93+ <td>ND</td>
94+ <td>0-8</td>
95+ <td>√</td>
96+ </tr>
97+ <tr>
98+ <td>scale</td>
99+ <td>输入</td>
100+ <td>公式中的scale，表示输入伪量化的缩放系数。</td>
101+ <td>size需要与zeroPoint一致且需要与self在axis轴的size一致。</td>
102+ <td>FLOAT16、FLOAT32</td>
103+ <td>ND</td>
104+ <td>1</td>
105+ <td>√</td>
106+ </tr>
107+ <tr>
108+ <td>zeroPoint</td>
109+ <td>输入</td>
110+ <td>公式中的zero_point，表示输入伪量化的零基准参数。</td>
111+ <td>size需要与scale一致且需要与self在axis轴的size一致。</td>
112+ <td>INT32</td>
113+ <td>ND</td>
114+ <td>1</td>
115+ <td>√</td>
116+ </tr>
117+ <tr>
118+ <td>axis</td>
119+ <td>输入</td>
120+ <td>表示计算维度。</td>
121+ <td>范围为[-self.dim(), self.dim()-1]。</td>
122+ <td>INT64</td>
123+ <td>-</td>
124+ <td>-</td>
125+ <td>-</td>
126+ </tr>
127+ <tr>
128+ <td>quantMin</td>
129+ <td>输入</td>
130+ <td>表示输入数据伪量化后的最小值。</td>
131+ <td>需要小于等于quantMax。</td>
132+ <td>INT64</td>
133+ <td>-</td>
134+ <td>-</td>
135+ <td>-</td>
136+ </tr>
137+ <tr>
138+ <td>quantMax</td>
139+ <td>输入</td>
140+ <td>表示输入数据伪量化后的最大值。</td>
141+ <td>需要大于等于quantMin。</td>
142+ <td>INT64</td>
143+ <td>-</td>
144+ <td>-</td>
145+ <td>-</td>
146+ </tr>
147+ <tr>
148+ <td>out</td>
149+ <td>输出</td>
150+ <td>公式中的out。</td>
151+ <td>shape需要和`self`一致。</td>
152+ <td>FLOAT16、FLOAT32</td>
153+ <td>ND</td>
154+ <td>0-8</td>
155+ <td>√</td>
156+ </tr>
157+ <tr>
158+ <td>mask</td>
159+ <td>输出</td>
160+ <td>公式中的mask。</td>
161+ <td>shape需要和`self`一致。</td>
162+ <td>BOOL</td>
163+ <td>ND</td>
164+ <td>0-8</td>
165+ <td>√</td>
166+ </tr>
167+ <tr>
168+ <td>workspaceSize</td>
169+ <td>输出</td>
170+ <td>返回需要在Device侧申请的workspace大小。</td>
171+ <td>-</td>
172+ <td>-</td>
173+ <td>-</td>
174+ <td>-</td>
175+ <td>-</td>
176+ </tr>
177+ <tr>
178+ <td>executor</td>
179+ <td>输出</td>
180+ <td>返回op执行器，包含了算子计算流程。</td>
181+ <td>-</td>
182+ <td>-</td>
183+ <td>-</td>
184+ <td>-</td>
185+ <td>-</td>
186+ </tr>
187+ </tbody></table>
48- **返回值：**188- **返回值：**
49 189 
50 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。190 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
51 191 
52- ```
53 第一段接口完成入参校验，出现以下场景时报错：192 第一段接口完成入参校验，出现以下场景时报错：
54- 返回161001 (ACLNN_ERR_PARAM_NULLPTR): 1. 传入的self、scale、zeroPoint、out或mask是空指针。193+ 
55- 返回161002 (ACLNN_ERR_PARAM_INVALID): 1. self、scale、zeroPoint、out或mask的数据类型不在支持的范围之内。194+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
56- 2. out和mask的shape与self不一致。195+ <col style="width: 267px">
57- 3. scale或zeroPoint的shape不是1维。196+ <col style="width: 124px">
58- 4. scale和zeroPoint的size大小不一致。197+ <col style="width: 775px">
59- 5. axis指定的维度不在合法范围[-self.dim(), self.dim()-1]内。198+ </colgroup>
60- 6. scale或zeroPoint的size与self在axis轴上的size大小不一致。199+ <thead>
61- 7. quantMin大于quantMax。200+ <tr>
62- ```201+ <th>返回码</th>
202+ <th>错误码</th>
203+ <th>描述</th>
204+ </tr></thead>
205+ <tbody>
206+ <tr>
207+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
208+ <td>161001</td>
209+ <td>传入的self、scale、zeroPoint、out或mask是空指针。</td>
210+ </tr>
211+ <tr>
212+ <td rowspan="7">ACLNN_ERR_PARAM_INVALID</td>
213+ <td rowspan="7">161002</td>
214+ <td>self、scale、zeroPoint、out或mask的数据类型不在支持的范围之内。</td>
215+ </tr>
216+ <tr>
217+ <td>out和mask的shape与self不一致。</td>
218+ </tr>
219+ <tr>
220+ <td>scale或zeroPoint的shape不是1维。</td>
221+ </tr>
222+ <tr>
223+ <td>scale和zeroPoint的size大小不一致。</td>
224+ </tr>
225+ <tr>
226+ <td>axis指定的维度不在合法范围[-self.dim(), self.dim()-1]内。</td>
227+ </tr>
228+ <tr>
229+ <td>scale或zeroPoint的size与self在axis轴上的size大小不一致。</td>
230+ </tr>
231+ <tr>
232+ <td>quantMin大于quantMax。</td>
233+ </tr>
234+ </tbody>
235+ </table>
63 236 
64## aclnnFakeQuantPerChannelAffineCachemask237## aclnnFakeQuantPerChannelAffineCachemask
65 238 
66- **参数说明：**239- **参数说明：**
240+
241+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
242+ <col style="width: 173px">
243+ <col style="width: 133px">
244+ <col style="width: 860px">
245+ </colgroup>
246+ <thead>
247+ <tr>
248+ <th>参数名</th>
249+ <th>输入/输出</th>
250+ <th>描述</th>
251+ </tr></thead>
252+ <tbody>
253+ <tr>
254+ <td>workspace</td>
255+ <td>输入</td>
256+ <td>在Device侧申请的workspace内存地址。</td>
257+ </tr>
258+ <tr>
259+ <td>workspaceSize</td>
260+ <td>输入</td>
261+ <td>在Device侧申请的workspace大小，由第一段接口aclnnFakeQuantPerChannelAffineCachemaskGetWorkspaceSize获取。</td>
262+ </tr>
263+ <tr>
264+ <td>executor</td>
265+ <td>输入</td>
266+ <td>op执行器，包含了算子计算流程。</td>
267+ </tr>
268+ <tr>
269+ <td>stream</td>
270+ <td>输入</td>
271+ <td>指定执行任务的Stream。</td>
272+ </tr>
273+ </tbody>
274+ </table>
67 275 
68- - workspace(void\*, 入参)：在Device侧申请的workspace内存地址。
69- - workspaceSize(uint64_t, 入参)：在Device侧申请的workspace大小，由第一段接口aclnnFakeQuantPerChannelAffineCachemaskGetWorkspaceSize获取。
70- - executor(aclOpExecutor\*, 入参)：op执行器，包含了算子计算流程。
71- - stream(aclrtStream, 入参)：指定执行任务的Stream。
72- **返回值：**276- **返回值：**
73 277 
74 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。278 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
@@ -79,6 +283,7 @@
79 283 
80- 当前新算子FakeQuantPerChannelAffineCachemask不支持zero_point的float32和float16输入，故先在aclnn接口内部拦截，待算子支持后放开该限制。284- 当前新算子FakeQuantPerChannelAffineCachemask不支持zero_point的float32和float16输入，故先在aclnn接口内部拦截，待算子支持后放开该限制。
81 285 
286+ 
82## 调用示例287## 调用示例
83 288 
84示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。289示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
@@ -150,7 +355,7 @@ int main() {
150 aclrtStream stream;355 aclrtStream stream;
151 auto ret = Init(deviceId, &stream);356 auto ret = Init(deviceId, &stream);
152 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);357 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
153- 358+
154 // 2. 构造输入与输出，需要根据API的接口自定义构造359 // 2. 构造输入与输出，需要根据API的接口自定义构造
155 std::vector<int64_t> selfShape = {1};360 std::vector<int64_t> selfShape = {1};
156 std::vector<int64_t> scaleShape = {1};361 std::vector<int64_t> scaleShape = {1};
@@ -202,11 +407,11 @@ int main() {
202 // 调用aclnnFakeQuantPerChannelAffineCachemask第二段接口407 // 调用aclnnFakeQuantPerChannelAffineCachemask第二段接口
203 ret = aclnnFakeQuantPerChannelAffineCachemask(workspaceAddr, workspaceSize, executor, stream);408 ret = aclnnFakeQuantPerChannelAffineCachemask(workspaceAddr, workspaceSize, executor, stream);
204 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFakeQuantPerChannelAffineCachemask failed. ERROR: %d\n", ret); return ret);409 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFakeQuantPerChannelAffineCachemask failed. ERROR: %d\n", ret); return ret);
205- 410+
206 // 4. （固定写法）同步等待任务执行结束411 // 4. （固定写法）同步等待任务执行结束
207 ret = aclrtSynchronizeStream(stream);412 ret = aclrtSynchronizeStream(stream);
208 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);413 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
209- 414+
210 // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改415 // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
211 auto size = GetShapeSize(outShape);416 auto size = GetShapeSize(outShape);
212 std::vector<float> resultData(size, 0);417 std::vector<float> resultData(size, 0);
Mquant/fake_quant_affine_cachemask/docs/aclnnFakeQuantPerTensorAffineCachemask.md+224-27
@@ -4,12 +4,16 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |6| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | × |
7| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |8| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |9| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | × |
9 13 
10## 功能说明14## 功能说明
11 15 
12-- 算子功能：16+- 接口功能：
13 - fake_quant_enabled >= 1: 对于输入数据self，使用scale和zero_point对输入self进行伪量化处理，并根据quant_min和quant_max对伪量化输出进行值域更新，最终返回结果out及对应位置掩码mask。17 - fake_quant_enabled >= 1: 对于输入数据self，使用scale和zero_point对输入self进行伪量化处理，并根据quant_min和quant_max对伪量化输出进行值域更新，最终返回结果out及对应位置掩码mask。
14 - fake_quant_enabled < 1: 返回结果out为self.clone()对象，掩码mask为全True。18 - fake_quant_enabled < 1: 返回结果out为self.clone()对象，掩码mask为全True。
15- 计算公式：在fake_quant_enabled >= 1的情况下，根据算子功能先计算临时变量qval，再计算得出out和mask。19- 计算公式：在fake_quant_enabled >= 1的情况下，根据算子功能先计算临时变量qval，再计算得出out和mask。
@@ -30,41 +34,234 @@
30 34 
31每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnFakeQuantPerTensorAffineCachemaskGetWorkspaceSize”接口获取入参并根据流程计算所需workspace大小，再调用“aclnnFakeQuantPerTensorAffineCachemask”接口执行计算。35每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnFakeQuantPerTensorAffineCachemaskGetWorkspaceSize”接口获取入参并根据流程计算所需workspace大小，再调用“aclnnFakeQuantPerTensorAffineCachemask”接口执行计算。
32 36 
33-* `aclnnStatus aclnnFakeQuantPerTensorAffineCachemaskGetWorkspaceSize(const aclTensor* self, const aclTensor* scale, const aclTensor* zeroPoint, float fakeQuantEnabled, int64_t quantMin, int64_t quantMax, aclTensor* out, aclTensor* mask, uint64_t* workspaceSize, aclOpExecutor** executor)`37+```Cpp
34-* `aclnnStatus aclnnFakeQuantPerTensorAffineCachemask(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)`38+aclnnStatus aclnnFakeQuantPerTensorAffineCachemaskGetWorkspaceSize(
39+ const aclTensor *self,
40+ const aclTensor *scale,
41+ const aclTensor *zeroPoint,
42+ float fakeQuantEnabled,
43+ int64_t quantMin,
44+ int64_t quantMax,
45+ aclTensor *out,
46+ aclTensor *mask,
47+ uint64_t *workspaceSize,
48+ aclOpExecutor **executor)
49+```
50+```Cpp
51+aclnnStatus aclnnFakeQuantPerTensorAffineCachemask(
52+ void *workspace,
53+ uint64_t workspaceSize,
54+ aclOpExecutor *executor,
55+ aclrtStream stream)
56+```
35 57 
36## aclnnFakeQuantPerTensorAffineCachemaskGetWorkspaceSize58## aclnnFakeQuantPerTensorAffineCachemaskGetWorkspaceSize
37 59 
38- **参数说明：**60- **参数说明：**
39- - self(aclTensor*, 计算输入)：公式中的`self`，Device侧的aclTensor，数据类型支持FLOAT16、FLOAT32。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。61+ 
40- - scale(aclTensor*, 计算输入)：公式中的`scale`，Device侧的aclTensor，表示输入伪量化的缩放系数。数据类型支持FLOAT16、FLOAT32，shape支持1维，size大小为1。62+ <table style="undefined;table-layout: fixed; width: 1478px"><colgroup>
41- - zeroPoint(aclTensor*, 计算输入)：公式中的`zero_point`，Device侧的aclTensor，表示输入伪量化的零基准参数。数据类型支持INT32，shape支持1维，size大小为1。63+ <col style="width: 149px">
42- - fakeQuantEnabled(float, 计算输入)：Host侧的浮点型，表示是否进行伪量化计算。64+ <col style="width: 121px">
43- - quantMin(int64_t, 计算输入)：Host侧的整型，表示输入数据伪量化后的最小值，需要小于等于quantMax。65+ <col style="width: 304px">
44- - quantMax(int64_t, 计算输入)：Host侧的整型，表示输入数据伪量化后的最大值，需要大于等于quantMin。66+ <col style="width: 253px">
45- - out(aclTensor\*, 计算输出)：Device侧的aclTensor，数据类型支持FLOAT16、FLOAT32，shape需要和计算输入`self`一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。67+ <col style="width: 222px">
46- - mask(aclTensor\*, 计算输出)：Device侧的aclTensor，数据类型支持BOOL，shape需要和计算输入`self`一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。68+ <col style="width: 148px">
47- - workspaceSize(uint64_t\*, 出参)：返回需要在Device侧申请的workspace大小。69+ <col style="width: 135px">
48- - executor(aclOpExecutor\*\*, 出参)：返回op执行器，包含了算子计算流程。70+ <col style="width: 146px">
71+ </colgroup>
72+ <thead>
73+ <tr>
74+ <th>参数名</th>
75+ <th>输入/输出</th>
76+ <th>描述</th>
77+ <th>使用说明</th>
78+ <th>数据类型</th>
79+ <th>数据格式</th>
80+ <th>维度(shape)</th>
81+ <th>非连续Tensor</th>
82+ </tr></thead>
83+ <tbody>
84+ <tr>
85+ <td>self</td>
86+ <td>输入</td>
87+ <td>公式中的self。</td>
88+ <td>-</td>
89+ <td>FLOAT16、FLOAT32</td>
90+ <td>ND</td>
91+ <td>0-8</td>
92+ <td>√</td>
93+ </tr>
94+ <tr>
95+ <td>scale</td>
96+ <td>输入</td>
97+ <td>公式中的scale，表示输入伪量化的缩放系数。</td>
98+ <td>size大小为1。</td>
99+ <td>FLOAT16、FLOAT32</td>
100+ <td>ND</td>
101+ <td>1</td>
102+ <td>√</td>
103+ </tr>
104+ <tr>
105+ <td>zeroPoint</td>
106+ <td>输入</td>
107+ <td>公式中的zero_point，表示输入伪量化的零基准参数。</td>
108+ <td>size大小为1。</td>
109+ <td>INT32</td>
110+ <td>ND</td>
111+ <td>1</td>
112+ <td>√</td>
113+ </tr>
114+ <tr>
115+ <td>fakeQuantEnabled</td>
116+ <td>输入</td>
117+ <td>表示是否进行伪量化计算。</td>
118+ <td>-</td>
119+ <td>FLOAT32</td>
120+ <td>-</td>
121+ <td>-</td>
122+ <td>-</td>
123+ </tr>
124+ <tr>
125+ <td>quantMin</td>
126+ <td>输入</td>
127+ <td>表示输入数据伪量化后的最小值。</td>
128+ <td>需要小于等于quantMax。</td>
129+ <td>INT64</td>
130+ <td>-</td>
131+ <td>-</td>
132+ <td>-</td>
133+ </tr>
134+ <tr>
135+ <td>quantMax</td>
136+ <td>输入</td>
137+ <td>表示输入数据伪量化后的最大值。</td>
138+ <td>需要大于等于quantMin。</td>
139+ <td>INT64</td>
140+ <td>-</td>
141+ <td>-</td>
142+ <td>-</td>
143+ </tr>
144+ <tr>
145+ <td>out</td>
146+ <td>输出</td>
147+ <td>公式中的out。</td>
148+ <td>shape需要和`self`一致。</td>
149+ <td>FLOAT16、FLOAT32</td>
150+ <td>ND</td>
151+ <td>0-8</td>
152+ <td>√</td>
153+ </tr>
154+ <tr>
155+ <td>mask</td>
156+ <td>输出</td>
157+ <td>公式中的mask。</td>
158+ <td>shape需要和`self`一致。</td>
159+ <td>BOOL</td>
160+ <td>ND</td>
161+ <td>0-8</td>
162+ <td>√</td>
163+ </tr>
164+ <tr>
165+ <td>workspaceSize</td>
166+ <td>输出</td>
167+ <td>返回需要在Device侧申请的workspace大小。</td>
168+ <td>-</td>
169+ <td>-</td>
170+ <td>-</td>
171+ <td>-</td>
172+ <td>-</td>
173+ </tr>
174+ <tr>
175+ <td>executor</td>
176+ <td>输出</td>
177+ <td>返回op执行器，包含了算子计算流程。</td>
178+ <td>-</td>
179+ <td>-</td>
180+ <td>-</td>
181+ <td>-</td>
182+ <td>-</td>
183+ </tr>
184+ </tbody></table>
185+ 
49- **返回值：**186- **返回值：**
50 187 
51 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。188 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
52- ```189+ 
53 第一段接口完成入参校验，出现以下场景时报错：190 第一段接口完成入参校验，出现以下场景时报错：
54- 返回161001 (ACLNN_ERR_PARAM_NULLPTR): 1. 传入的self、scale、zeroPoint、out或mask是空指针。191+ 
55- 返回161002 (ACLNN_ERR_PARAM_INVALID): 1. self、scale、zeroPoint、out或mask的数据类型不在支持的范围之内。192+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
56- 2. scale或zeroPoint的size大小不是1。193+ <col style="width: 267px">
57- 3. out和mask的shape与self不一致194+ <col style="width: 124px">
58- 4. quantMin大于quantMax。195+ <col style="width: 775px">
59- ```196+ </colgroup>
197+ <thead>
198+ <tr>
199+ <th>返回码</th>
200+ <th>错误码</th>
201+ <th>描述</th>
202+ </tr></thead>
203+ <tbody>
204+ <tr>
205+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
206+ <td>161001</td>
207+ <td>传入的self、scale、zeroPoint、out或mask是空指针。</td>
208+ </tr>
209+ <tr>
210+ <td rowspan="4">ACLNN_ERR_PARAM_INVALID</td>
211+ <td rowspan="4">161002</td>
212+ <td>self、scale、zeroPoint、out或mask的数据类型不在支持的范围之内。</td>
213+ </tr>
214+ <tr>
215+ <td>scale或zeroPoint的size大小不是1。</td>
216+ </tr>
217+ <tr>
218+ <td>out和mask的shape与self不一致</td>
219+ </tr>
220+ <tr>
221+ <td>quantMin大于quantMax。</td>
222+ </tr>
223+ </tbody>
224+ </table>
60 225 
61## aclnnFakeQuantPerTensorAffineCachemask226## aclnnFakeQuantPerTensorAffineCachemask
62 227 
63- **参数说明：**228- **参数说明：**
64- - workspace(void\*, 入参)：在Device侧申请的workspace内存地址。229+
65- - workspaceSize(uint64_t, 入参)：在Device侧申请的workspace大小，由第一段接口aclnnFakeQuantPerTensorAffineCachemaskGetWorkspaceSize获取。230+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
66- - executor(aclOpExecutor\*, 入参)：op执行器，包含了算子计算流程。231+ <col style="width: 173px">
67- - stream(aclrtStream, 入参)：指定执行任务的Stream。232+ <col style="width: 133px">
233+ <col style="width: 860px">
234+ </colgroup>
235+ <thead>
236+ <tr>
237+ <th>参数名</th>
238+ <th>输入/输出</th>
239+ <th>描述</th>
240+ </tr></thead>
241+ <tbody>
242+ <tr>
243+ <td>workspace</td>
244+ <td>输入</td>
245+ <td>在Device侧申请的workspace内存地址。</td>
246+ </tr>
247+ <tr>
248+ <td>workspaceSize</td>
249+ <td>输入</td>
250+ <td>在Device侧申请的workspace大小，由第一段接口aclnnFakeQuantPerTensorAffineCachemaskGetWorkspaceSize获取。</td>
251+ </tr>
252+ <tr>
253+ <td>executor</td>
254+ <td>输入</td>
255+ <td>op执行器，包含了算子计算流程。</td>
256+ </tr>
257+ <tr>
258+ <td>stream</td>
259+ <td>输入</td>
260+ <td>指定执行任务的Stream。</td>
261+ </tr>
262+ </tbody>
263+ </table>
264+
68- **返回值：**265- **返回值：**
69 266 
70 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。267 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
@@ -144,7 +341,7 @@ int main() {
144 aclrtStream stream;341 aclrtStream stream;
145 auto ret = Init(deviceId, &stream);342 auto ret = Init(deviceId, &stream);
146 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);343 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
147- 344+
148 // 2. 构造输入与输出，需要根据API的接口自定义构造345 // 2. 构造输入与输出，需要根据API的接口自定义构造
149 std::vector<int64_t> selfShape = {1};346 std::vector<int64_t> selfShape = {1};
150 std::vector<int64_t> scaleShape = {1};347 std::vector<int64_t> scaleShape = {1};
@@ -196,11 +393,11 @@ int main() {
196 // 调用aclnnFakeQuantPerTensorAffineCachemask第二段接口393 // 调用aclnnFakeQuantPerTensorAffineCachemask第二段接口
197 ret = aclnnFakeQuantPerTensorAffineCachemask(workspaceAddr, workspaceSize, executor, stream);394 ret = aclnnFakeQuantPerTensorAffineCachemask(workspaceAddr, workspaceSize, executor, stream);
198 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFakeQuantPerTensorAffineCachemask failed. ERROR: %d\n", ret); return ret);395 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFakeQuantPerTensorAffineCachemask failed. ERROR: %d\n", ret); return ret);
199- 396+
200 // 4. （固定写法）同步等待任务执行结束397 // 4. （固定写法）同步等待任务执行结束
201 ret = aclrtSynchronizeStream(stream);398 ret = aclrtSynchronizeStream(stream);
202 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);399 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
203- 400+
204 // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改401 // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
205 auto size = GetShapeSize(outShape);402 auto size = GetShapeSize(outShape);
206 std::vector<float> resultData(size, 0);403 std::vector<float> resultData(size, 0);
Mquant/fake_quant_affine_cachemask/op_host/config/ascend910_93/fake_quant_affine_cachemask_binary.json+148-0
@@ -148,6 +148,154 @@
148 "value": 1148 "value": 1
149 }149 }
150 ]150 ]
151+ },
152+ {
153+ "bin_filename": "FakeQuantAffineCachemask_fp32_fp32",
154+ "inputs": [
155+ {
156+ "name": "x",
157+ "index": 0,
158+ "dtype": "float32",
159+ "format": "ND",
160+ "paramType": "required",
161+ "shape": [
162+ -2
163+ ]
164+ },
165+ {
166+ "name": "scale",
167+ "index": 1,
168+ "dtype": "float32",
169+ "format": "ND",
170+ "paramType": "required",
171+ "shape": [
172+ -2
173+ ]
174+ },
175+ {
176+ "name": "zero_point",
177+ "index": 2,
178+ "dtype": "float32",
179+ "format": "ND",
180+ "paramType": "required",
181+ "shape": [
182+ -2
183+ ]
184+ }
185+ ],
186+ "outputs": [
187+ {
188+ "name": "y",
189+ "index": 0,
190+ "dtype": "float32",
191+ "format": "ND",
192+ "paramType": "required",
193+ "shape": [
194+ -2
195+ ]
196+ },
197+ {
198+ "name": "mask",
199+ "index": 0,
200+ "dtype": "bool",
201+ "format": "ND",
202+ "paramType": "required",
203+ "shape": [
204+ -2
205+ ]
206+ }
207+ ],
208+ "attrs": [
209+ {
210+ "name": "axis",
211+ "dtype": "int",
212+ "value": 1
213+ },
214+ {
215+ "name": "quant_min",
216+ "dtype": "int",
217+ "value": 1
218+ },
219+ {
220+ "name": "quant_max",
221+ "dtype": "int",
222+ "value": 1
223+ }
224+ ]
225+ },
226+ {
227+ "bin_filename": "FakeQuantAffineCachemask_fp16_fp16",
228+ "inputs": [
229+ {
230+ "name": "x",
231+ "index": 0,
232+ "dtype": "float16",
233+ "format": "ND",
234+ "paramType": "required",
235+ "shape": [
236+ -2
237+ ]
238+ },
239+ {
240+ "name": "scale",
241+ "index": 1,
242+ "dtype": "float16",
243+ "format": "ND",
244+ "paramType": "required",
245+ "shape": [
246+ -2
247+ ]
248+ },
249+ {
250+ "name": "zero_point",
251+ "index": 2,
252+ "dtype": "float16",
253+ "format": "ND",
254+ "paramType": "required",
255+ "shape": [
256+ -2
257+ ]
258+ }
259+ ],
260+ "outputs": [
261+ {
262+ "name": "y",
263+ "index": 0,
264+ "dtype": "float16",
265+ "format": "ND",
266+ "paramType": "required",
267+ "shape": [
268+ -2
269+ ]
270+ },
271+ {
272+ "name": "mask",
273+ "index": 0,
274+ "dtype": "bool",
275+ "format": "ND",
276+ "paramType": "required",
277+ "shape": [
278+ -2
279+ ]
280+ }
281+ ],
282+ "attrs": [
283+ {
284+ "name": "axis",
285+ "dtype": "int",
286+ "value": 1
287+ },
288+ {
289+ "name": "quant_min",
290+ "dtype": "int",
291+ "value": 1
292+ },
293+ {
294+ "name": "quant_max",
295+ "dtype": "int",
296+ "value": 1
297+ }
298+ ]
151 }299 }
152 ]300 ]
153 }301 }
Mquant/fake_quant_affine_cachemask/op_host/config/ascend910b/fake_quant_affine_cachemask_binary.json+148-0
@@ -148,6 +148,154 @@
148 "value": 1148 "value": 1
149 }149 }
150 ]150 ]
151+ },
152+ {
153+ "bin_filename": "FakeQuantAffineCachemask_fp32_fp32",
154+ "inputs": [
155+ {
156+ "name": "x",
157+ "index": 0,
158+ "dtype": "float32",
159+ "format": "ND",
160+ "paramType": "required",
161+ "shape": [
162+ -2
163+ ]
164+ },
165+ {
166+ "name": "scale",
167+ "index": 1,
168+ "dtype": "float32",
169+ "format": "ND",
170+ "paramType": "required",
171+ "shape": [
172+ -2
173+ ]
174+ },
175+ {
176+ "name": "zero_point",
177+ "index": 2,
178+ "dtype": "float32",
179+ "format": "ND",
180+ "paramType": "required",
181+ "shape": [
182+ -2
183+ ]
184+ }
185+ ],
186+ "outputs": [
187+ {
188+ "name": "y",
189+ "index": 0,
190+ "dtype": "float32",
191+ "format": "ND",
192+ "paramType": "required",
193+ "shape": [
194+ -2
195+ ]
196+ },
197+ {
198+ "name": "mask",
199+ "index": 0,
200+ "dtype": "bool",
201+ "format": "ND",
202+ "paramType": "required",
203+ "shape": [
204+ -2
205+ ]
206+ }
207+ ],
208+ "attrs": [
209+ {
210+ "name": "axis",
211+ "dtype": "int",
212+ "value": 1
213+ },
214+ {
215+ "name": "quant_min",
216+ "dtype": "int",
217+ "value": 1
218+ },
219+ {
220+ "name": "quant_max",
221+ "dtype": "int",
222+ "value": 1
223+ }
224+ ]
225+ },
226+ {
227+ "bin_filename": "FakeQuantAffineCachemask_fp16_fp16",
228+ "inputs": [
229+ {
230+ "name": "x",
231+ "index": 0,
232+ "dtype": "float16",
233+ "format": "ND",
234+ "paramType": "required",
235+ "shape": [
236+ -2
237+ ]
238+ },
239+ {
240+ "name": "scale",
241+ "index": 1,
242+ "dtype": "float16",
243+ "format": "ND",
244+ "paramType": "required",
245+ "shape": [
246+ -2
247+ ]
248+ },
249+ {
250+ "name": "zero_point",
251+ "index": 2,
252+ "dtype": "float16",
253+ "format": "ND",
254+ "paramType": "required",
255+ "shape": [
256+ -2
257+ ]
258+ }
259+ ],
260+ "outputs": [
261+ {
262+ "name": "y",
263+ "index": 0,
264+ "dtype": "float16",
265+ "format": "ND",
266+ "paramType": "required",
267+ "shape": [
268+ -2
269+ ]
270+ },
271+ {
272+ "name": "mask",
273+ "index": 0,
274+ "dtype": "bool",
275+ "format": "ND",
276+ "paramType": "required",
277+ "shape": [
278+ -2
279+ ]
280+ }
281+ ],
282+ "attrs": [
283+ {
284+ "name": "axis",
285+ "dtype": "int",
286+ "value": 1
287+ },
288+ {
289+ "name": "quant_min",
290+ "dtype": "int",
291+ "value": 1
292+ },
293+ {
294+ "name": "quant_max",
295+ "dtype": "int",
296+ "value": 1
297+ }
298+ ]
151 }299 }
152 ]300 ]
153 }301 }
Mquant/fake_quant_affine_cachemask/op_host/op_api/aclnn_fake_quant_per_channel_affine_cachemask.cpp+4-1
@@ -37,6 +37,9 @@ extern "C" {
37static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {37static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {
38 op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT};38 op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT};
39 39 
40+static const std::initializer_list<op::DataType> ZERO_POINT_DTYPE_SUPPORT_LIST = {
41+ op::DataType::DT_INT32, op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT};
42+ 
40static bool CheckNotNull(43static bool CheckNotNull(
41 const aclTensor* self, const aclTensor* scale, const aclTensor* zeroPoint, const aclTensor* out,44 const aclTensor* self, const aclTensor* scale, const aclTensor* zeroPoint, const aclTensor* out,
42 const aclTensor* mask)45 const aclTensor* mask)
@@ -55,7 +58,7 @@ static bool CheckDtypeValid(
55{58{
56 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST, return false);59 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST, return false);
57 OP_CHECK_DTYPE_NOT_SUPPORT(scale, DTYPE_SUPPORT_LIST, return false);60 OP_CHECK_DTYPE_NOT_SUPPORT(scale, DTYPE_SUPPORT_LIST, return false);
58- OP_CHECK_DTYPE_NOT_MATCH(zeroPoint, op::DataType::DT_INT32, return false);61+ OP_CHECK_DTYPE_NOT_SUPPORT(zeroPoint, ZERO_POINT_DTYPE_SUPPORT_LIST, return false);
59 62 
60 op::DataType promoteType = op::PromoteType(self->GetDataType(), scale->GetDataType());63 op::DataType promoteType = op::PromoteType(self->GetDataType(), scale->GetDataType());
61 if (!CanCast(DataType(out->GetDataType()), promoteType)) {64 if (!CanCast(DataType(out->GetDataType()), promoteType)) {
Mquant/fake_quant_affine_cachemask/op_kernel/fake_quant_affine_cachemask_fp16.h+20-10
@@ -34,7 +34,8 @@ public:
34 34 
35 xGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(x) + this->offset, this->blockLength);35 xGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(x) + this->offset, this->blockLength);
36 scaleGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(scale) + this->scaleOffset, this->circleNum);36 scaleGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(scale) + this->scaleOffset, this->circleNum);
37- zeroGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(zero_point) + this->scaleOffset, this->circleNum);37+ zeroGm.SetGlobalBuffer(
CANN-robot
CANN-robotCANN-robot1月17日

数组越界与内存安全: 在设置zeroGm的GlobalBuffer时，使用reinterpret_cast<__gm__ DTYPE_ZERO_POINT*>(zero_point) + this->scaleOffset进行指针偏移。scaleOffset是用于scale张量的偏移量，但这里被错误地用于zero_point张量。zero_point张量可能有自己的内存布局和偏移量，直接复用scaleOffset可能导致访问越界或数据错位。从代码逻辑看，scale和zero_point应该是成对使用的参数，但它们的偏移量可能不同，特别是在某些数据排布场景下。

问题类型: 数组越界与内存安全 文件路径: quant/fake_quant_affine_cachemask/op_kernel/fake_quant_affine_cachemask_fp16.h 行号: 37 问题代码:

zeroGm.SetGlobalBuffer(
    reinterpret_cast<__gm__ DTYPE_ZERO_POINT*>(zero_point) + this->scaleOffset, this->circleNum);

修改建议:

应该为zero_point张量定义独立的偏移量成员变量（如`zeroOffset`），并在tiling数据中传递该值。修改为：`reinterpret_cast<__gm__ DTYPE_ZERO_POINT*>(zero_point) + this->zeroOffset, this->circleNum`，并确保基类`FakeQuantAffineCachemaskBase`中有相应的`zeroOffset`成员。

此评论由代码审查工具自动生成

likedislike
38+ reinterpret_cast<__gm__ DTYPE_ZERO_POINT*>(zero_point) + this->scaleOffset, this->circleNum);
38 yGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(y) + this->offset, this->blockLength);39 yGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(y) + this->offset, this->blockLength);
39 maskGm.SetGlobalBuffer(reinterpret_cast<__gm__ uint8_t*>(mask) + this->offset, this->blockLength);40 maskGm.SetGlobalBuffer(reinterpret_cast<__gm__ uint8_t*>(mask) + this->offset, this->blockLength);
40 41 
@@ -101,13 +102,21 @@ private:
101 selectTemp = selectBuf.Get<yType>();102 selectTemp = selectBuf.Get<yType>();
102 maskTemp = maskBuf.Get<uint8_t>();103 maskTemp = maskBuf.Get<uint8_t>();
103 104 
104- // tmp = x / scale + zero_point
105 Cast(curTemp, xLocal, RoundMode::CAST_NONE, calCount);105 Cast(curTemp, xLocal, RoundMode::CAST_NONE, calCount);
106 Muls(curTemp, curTemp, static_cast<float>(1.0f / scaleValue), calCount);106 Muls(curTemp, curTemp, static_cast<float>(1.0f / scaleValue), calCount);
107- Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);107+ if constexpr (IsSameType<DTYPE_ZERO_POINT, int32_t>::value) {
CANN-robot
CANN-robotCANN-robot1月17日

逻辑运算与副作用: 在条件编译分支中，当DTYPE_ZERO_POINT为int32_t时，执行Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount)将float转换为int32_t，然后立即又Cast(curTemp, curInt32Temp, RoundMode::CAST_NONE, calCount)转换回float。这个操作实际上是在进行四舍五入取整，但紧接着又执行Adds(curTemp, curTemp, static_cast<float>(zeroPointValue), calCount)。这里存在精度损失风险：先取整再加zero_point，与先加zero_point再取整的结果可能不同。从注释看，算法应该是tmp = x / scale + zero_point，但这里的实现实际上是tmp = round(x / scale) + zero_point。

问题类型: 逻辑运算与副作用 文件路径: quant/fake_quant_affine_cachemask/op_kernel/fake_quant_affine_cachemask_fp16.h 行号: 107 问题代码:

if constexpr (IsSameType<DTYPE_ZERO_POINT, int32_t>::value) {
    // tmp = x / scale + zero_point
    Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);
    Cast(curTemp, curInt32Temp, RoundMode::CAST_NONE, calCount);
    Adds(curTemp, curTemp, static_cast<float>(zeroPointValue), calCount);
    Cast(curHf16Temp, curTemp, RoundMode::CAST_RINT, calCount);
}

修改建议:

如果算法意图确实是`tmp = x / scale + zero_point`，应该先执行加法再取整。修改为：`Adds(curTemp, curTemp, static_cast<float>(zeroPointValue), calCount); Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);`。如果意图是`tmp = round(x / scale) + zero_point`，则注释应相应更新以准确反映实现。

此评论由代码审查工具自动生成

likedislike
108- Cast(curTemp, curInt32Temp, RoundMode::CAST_NONE, calCount);108+ // tmp = x / scale + zero_point
109- Adds(curTemp, curTemp, static_cast<float>(zeroPointValue), calCount);109+ Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);
110- Cast(curHf16Temp, curTemp, RoundMode::CAST_RINT, calCount);110+ Cast(curTemp, curInt32Temp, RoundMode::CAST_NONE, calCount);
111+ Adds(curTemp, curTemp, static_cast<float>(zeroPointValue), calCount);
112+ Cast(curHf16Temp, curTemp, RoundMode::CAST_RINT, calCount);
113+ } else {
114+ // tmp = lrintf(x / scale + zero_point)
115+ Adds(curTemp, curTemp, static_cast<float>(zeroPointValue), calCount);
116+ Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);
117+ Cast(curTemp, curInt32Temp, RoundMode::CAST_NONE, calCount);
118+ Cast(curHf16Temp, curTemp, RoundMode::CAST_RINT, calCount);
119+ }
111 PipeBarrier<PIPE_ALL>();120 PipeBarrier<PIPE_ALL>();
112 121 
113 // maskTemp = (round(tmp) >= quant_min) & (round(tmp) <= quant_max)122 // maskTemp = (round(tmp) >= quant_min) & (round(tmp) <= quant_max)
@@ -130,8 +139,9 @@ private:
130 Cast(curTemp, curInt32Temp, RoundMode::CAST_ROUND, calCount);139 Cast(curTemp, curInt32Temp, RoundMode::CAST_ROUND, calCount);
131 Compare(maskTemp, xLocal, xLocal, CMPMODE::EQ, calCount);140 Compare(maskTemp, xLocal, xLocal, CMPMODE::EQ, calCount);
132 Select(141 Select(
133- curTemp, maskTemp, curTemp, 0.0f, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->mask, repeatTimes, repeatParams);142+ curTemp, maskTemp, curTemp, 0.0f, SELMODE::VSEL_TENSOR_SCALAR_MODE,
134- Adds(curTemp, curTemp, static_cast<float>(-1 * zeroPointValue), calCount);143+ this->mask, repeatTimes, repeatParams);
144+ Adds(curTemp, curTemp, static_cast<float>(-1 * static_cast<float>(zeroPointValue)), calCount);
135 Muls(curTemp, curTemp, static_cast<float>(scaleValue), calCount);145 Muls(curTemp, curTemp, static_cast<float>(scaleValue), calCount);
136 Cast(yLocal, curTemp, RoundMode::CAST_RINT, calCount);146 Cast(yLocal, curTemp, RoundMode::CAST_RINT, calCount);
137 147 
@@ -143,12 +153,12 @@ private:
143private:153private:
144 TPipe pipe;154 TPipe pipe;
145 155 
146- int32_t zeroPointValue = 1;156+ DTYPE_ZERO_POINT zeroPointValue = 1;
147 uint8_t repeatTimes = 0;157 uint8_t repeatTimes = 0;
148 float scaleValue = 1.0;158 float scaleValue = 1.0;
149 159 
150 GlobalTensor<yType> xGm, scaleGm, yGm;160 GlobalTensor<yType> xGm, scaleGm, yGm;
151- GlobalTensor<int32_t> zeroGm;161+ GlobalTensor<DTYPE_ZERO_POINT> zeroGm;
152 GlobalTensor<uint8_t> maskGm;162 GlobalTensor<uint8_t> maskGm;
153 163 
154 TQue<QuePosition::VECIN, BUFFER_NUM> inQueueData;164 TQue<QuePosition::VECIN, BUFFER_NUM> inQueueData;
Mquant/fake_quant_affine_cachemask/op_kernel/fake_quant_affine_cachemask_fp32.h+19-11
@@ -34,7 +34,8 @@ public:
34 34 
35 xGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(x) + this->offset, this->blockLength);35 xGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(x) + this->offset, this->blockLength);
36 scaleGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(scale) + this->scaleOffset, this->circleNum);36 scaleGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(scale) + this->scaleOffset, this->circleNum);
37- zeroGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(zero_point) + this->scaleOffset, this->circleNum);37+ zeroGm.SetGlobalBuffer(
38+ reinterpret_cast<__gm__ DTYPE_ZERO_POINT*>(zero_point) + this->scaleOffset, this->circleNum);
38 yGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(y) + this->offset, this->blockLength);39 yGm.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(y) + this->offset, this->blockLength);
39 maskGm.SetGlobalBuffer(reinterpret_cast<__gm__ uint8_t*>(mask) + this->offset, this->blockLength);40 maskGm.SetGlobalBuffer(reinterpret_cast<__gm__ uint8_t*>(mask) + this->offset, this->blockLength);
40 41 
@@ -101,13 +102,19 @@ private:
101 selectTemp = selectBuf.Get<yType>();102 selectTemp = selectBuf.Get<yType>();
102 maskTemp = maskBuf.Get<uint8_t>();103 maskTemp = maskBuf.Get<uint8_t>();
103 104 
104- // tmp = x / scale + zero_point
105 Muls(curTemp, xLocal, static_cast<yType>(1.0f / scaleValue), calCount);105 Muls(curTemp, xLocal, static_cast<yType>(1.0f / scaleValue), calCount);
106- Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);106+ if constexpr (IsSameType<DTYPE_ZERO_POINT, int32_t>::value) {
107- Cast(curTemp, curInt32Temp, RoundMode::CAST_NONE, calCount);107+ // tmp = x / scale + zero_point
108- Cast(yLocal, curInt32Temp, RoundMode::CAST_NONE, calCount);108+ Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);
109- Adds(curTemp, curTemp, static_cast<yType>(zeroPointValue), calCount);109+ Cast(curTemp, curInt32Temp, RoundMode::CAST_NONE, calCount);
110- Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);110+ Adds(curTemp, curTemp, static_cast<yType>(zeroPointValue), calCount);
111+ Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);
112+ } else {
113+ // tmp = lrintf(x / scale + zero_point)
114+ Adds(curTemp, curTemp, static_cast<yType>(zeroPointValue), calCount);
115+ Cast(curInt32Temp, curTemp, RoundMode::CAST_RINT, calCount);
116+ Cast(curTemp, curInt32Temp, RoundMode::CAST_NONE, calCount);
117+ }
111 PipeBarrier<PIPE_ALL>();118 PipeBarrier<PIPE_ALL>();
112 119 
113 // maskTemp = (round(curTemp) >= quant_min) & (round(curTemp) <= quant_max)120 // maskTemp = (round(curTemp) >= quant_min) & (round(curTemp) <= quant_max)
@@ -133,8 +140,9 @@ private:
133 Cast(curTemp, curInt32Temp, RoundMode::CAST_ROUND, calCount);140 Cast(curTemp, curInt32Temp, RoundMode::CAST_ROUND, calCount);
134 Compare(maskTemp, xLocal, xLocal, CMPMODE::EQ, calCount);141 Compare(maskTemp, xLocal, xLocal, CMPMODE::EQ, calCount);
135 Select(142 Select(
136- curTemp, maskTemp, curTemp, 0.0f, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->mask, repeatTimes, repeatParams);143+ curTemp, maskTemp, curTemp, 0.0f, SELMODE::VSEL_TENSOR_SCALAR_MODE,
137- Adds(curTemp, curTemp, static_cast<yType>(-1 * zeroPointValue), calCount);144+ this->mask, repeatTimes, repeatParams);
145+ Adds(curTemp, curTemp, static_cast<yType>(-1 * static_cast<float>(zeroPointValue)), calCount);
138 Muls(yLocal, curTemp, static_cast<yType>(scaleValue), calCount);146 Muls(yLocal, curTemp, static_cast<yType>(scaleValue), calCount);
139 PipeBarrier<PIPE_ALL>();147 PipeBarrier<PIPE_ALL>();
140 148 
@@ -146,12 +154,12 @@ private:
146private:154private:
147 TPipe pipe;155 TPipe pipe;
148 156 
149- int32_t zeroPointValue = 1;157+ DTYPE_ZERO_POINT zeroPointValue = 1;
CANN-robot
CANN-robotCANN-robot1月17日

变量初始化: 第157行zeroPointValue初始化为1，但DTYPE_ZERO_POINT可能是int32_t或浮点类型。对于浮点类型，初始化为1.0更合适。虽然不影响运行时（因为会在Process()中被重新赋值），但为了代码清晰性和避免潜在的类型转换问题，建议根据类型进行初始化。

问题类型: 变量初始化 文件路径: quant/fake_quant_affine_cachemask/op_kernel/fake_quant_affine_cachemask_fp32.h 行号: 157 问题代码:

DTYPE_ZERO_POINT zeroPointValue = 1;

修改建议:

使用条件初始化：zeroPointValue = static_cast<DTYPE_ZERO_POINT>(1); 或者利用if constexpr在构造函数中初始化。

此评论由代码审查工具自动生成

likedislike
150 float scaleValue = 1.0;158 float scaleValue = 1.0;
151 uint8_t repeatTimes = 0;159 uint8_t repeatTimes = 0;
152 160 
153 GlobalTensor<yType> xGm, scaleGm, yGm;161 GlobalTensor<yType> xGm, scaleGm, yGm;
154- GlobalTensor<int32_t> zeroGm;162+ GlobalTensor<DTYPE_ZERO_POINT> zeroGm;
155 GlobalTensor<uint8_t> maskGm;163 GlobalTensor<uint8_t> maskGm;
156 164 
157 TQue<QuePosition::VECIN, BUFFER_NUM> inQueueData;165 TQue<QuePosition::VECIN, BUFFER_NUM> inQueueData;
Aquant/fake_quant_affine_cachemask/tests/st/aclnnFakeQuantPerChannelAffineCachemask/all_aclnnFakeQuantPerChannelAffineCachemask.json+1-0
Aquant/fake_quant_affine_cachemask/tests/st/aclnnFakeQuantPerChannelAffineCachemask/executor_aclnnFakeQuantPerChannelAffineCachemask.py+69-0
@@ -0,0 +1,69 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+ 
13+import torch
14+from atk.configs.dataset_config import InputDataset
15+from atk.configs.results_config import TaskResult
16+from atk.tasks.api_execute import register
17+from atk.tasks.api_execute.base_api import BaseApi
18+from atk.tasks.dataset.base_dataset import OpsDataset
19+from atk.configs.base_config import TaskType
20+ 
21+ 
22+@register("aclnn_fake_quant_per_channel_affine_cachemask")
23+class MethodTorchNnFakeQuantAffineCachemaskApi(BaseApi):
24+ def __init__(self, task_result: TaskResult):
25+ super(MethodTorchNnFakeQuantAffineCachemaskApi, self).__init__(task_result)
26+ OpsDataset.seed_everything()
27+ 
28+ def init_by_input_data(self, input_data: InputDataset):
29+ """
30+ 该接口可实现部门场景下api的初始化需要依赖于当前的输入数据，且不希望计入耗时，
31+ 可以在此接口实现
32+ :param input_data:
33+ :return:
34+ """
35+ # device处理 pyaclnn仅在此阶段区分 后续调用仍走cpu获取输出格式，此处转device仅构造输入
36+ if self.device == "gpu":
37+ device = f"cuda:{self.device_id}"
38+ elif self.device in ("npu", "pyaclnn"):
39+ device = f"npu:{self.device_id}"
40+ else:
41+ device = "cpu"
42+
43+ if self.device == "cpu":
44+ # 标杆的scale仅支持fp32
45+ input_data.kwargs['scale'] = input_data.kwargs['scale'].to(torch.float32)
46+ input_data.kwargs['self'] = input_data.kwargs['self'].to(device)
47+ input_data.kwargs['scale'] = input_data.kwargs['scale'].to(device)
48+ input_data.kwargs['zeroPoint'] = input_data.kwargs['zeroPoint'].to(device)
49+
50+ def __call__(self, input_data: InputDataset, with_output: bool = False):
51+ 
52+ func = eval(self.api_name)
53+ output = mask = None
54+ if with_output:
55+ output, mask = func(input_data.kwargs['self'],
56+ input_data.kwargs['scale'],
57+ input_data.kwargs['zeroPoint'],
58+ input_data.kwargs['axis'],
59+ input_data.kwargs['quantMin'],
60+ input_data.kwargs['quantMax'])
61+ else:
62+ func(input_data.kwargs['self'],
63+ input_data.kwargs['scale'],
64+ input_data.kwargs['zeroPoint'],
65+ input_data.kwargs['axis'],
66+ input_data.kwargs['quantMin'],
67+ input_data.kwargs['quantMax'])
68+ 
69+ return output, mask
Aquant/fake_quant_affine_cachemask/tests/st/aclnnFakeQuantPerTensorAffineCachemask/all_aclnnFakeQuantPerTensorAffineCachemask.json+1-0
Aquant/fake_quant_affine_cachemask/tests/st/aclnnFakeQuantPerTensorAffineCachemask/executor_aclnnFakeQuantPerTensorAffineCachemask.py+72-0
@@ -0,0 +1,72 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+ 
13+import torch
14+from atk.configs.dataset_config import InputDataset
15+from atk.configs.results_config import TaskResult
16+from atk.tasks.api_execute import register
17+from atk.tasks.api_execute.base_api import BaseApi
18+from atk.tasks.dataset.base_dataset import OpsDataset
19+ 
20+ 
21+@register("aclnn_fake_quant_per_tensor_affine_cachemask")
22+class MethodTorchNnFakeQuantAffineCachemaskApi(BaseApi):
23+ def __init__(self, task_result: TaskResult):
24+ super(MethodTorchNnFakeQuantAffineCachemaskApi, self).__init__(task_result)
25+ OpsDataset.seed_everything()
26+ 
27+ def init_by_input_data(self, input_data: InputDataset):
28+ """
29+ 该接口可实现部门场景下api的初始化需要依赖于当前的输入数据，且不希望计入耗时，
30+ 可以在此接口实现
31+ :param input_data:
32+ :return:
33+ """
34+ # 确保quantMin <= quantMax
35+ if input_data.kwargs['quantMin'] is None or input_data.kwargs['quantMax'] is None:
36+ return
37+ if input_data.kwargs['quantMin'] > input_data.kwargs['quantMax']:
38+ input_data.kwargs['quantMin'], input_data.kwargs['quantMax'] = \
39+ input_data.kwargs['quantMax'], input_data.kwargs['quantMin']
40+ 
41+ def __call__(self, input_data: InputDataset, with_output: bool = False):
42+ func = eval(self.api_name)
43+
44+ if self.device == "gpu":
45+ device = f"cuda:{self.device_id}"
46+ elif self.device == "npu":
47+ device = f"{self.device}:{self.device_id}"
48+ else:
49+ device = "cpu"
50+ input_data.kwargs['self'] = input_data.kwargs['self'].to(device)
51+ input_data.kwargs['scale'] = input_data.kwargs['scale'].to(device)
52+ input_data.kwargs['zeroPoint'] = input_data.kwargs['zeroPoint'].to(device)
53+ input_data.kwargs['fakeQuantEnbled'] = torch.tensor(input_data.kwargs['fakeQuantEnbled'],
54+ dtype=torch.bool).to(device)
55+ 
56+ output = mask = None
57+ if with_output:
58+ output, mask = func(input_data.kwargs['self'],
59+ input_data.kwargs['scale'],
60+ input_data.kwargs['zeroPoint'],
61+ input_data.kwargs['fakeQuantEnbled'],
62+ input_data.kwargs['quantMin'],
63+ input_data.kwargs['quantMax'])
64+ else:
65+ func(input_data.kwargs['self'],
66+ input_data.kwargs['scale'],
67+ input_data.kwargs['zeroPoint'],
68+ input_data.kwargs['fakeQuantEnbled'],
69+ input_data.kwargs['quantMin'],
70+ input_data.kwargs['quantMax'])
71+ 
72+ return output, mask
Mquant/fake_quant_affine_cachemask/tests/ut/op_kernel/CMakeLists.txt+1-1
@@ -9,5 +9,5 @@
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)11if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
12- AddOpTestCase(fake_quant_affine_cachemask "ascend910B1" "")12+ AddOpTestCase(fake_quant_affine_cachemask "ascend910B1" "-DDTYPE_ZERO_POINT=float")
13endif()13endif()
Mquant/fake_quant_affine_cachemask/tests/ut/op_kernel/test_fake_quant_affine_cachemask.cpp+2-0
@@ -22,6 +22,8 @@
22 22 
23#include <cstdint>23#include <cstdint>
24 24 
25+#define DTYPE_ZERO_POINT DT_FLOAT
26+ 
25using namespace std;27using namespace std;
26 28 
27extern "C" __global__ __aicore__ void fake_quant_affine_cachemask(29extern "C" __global__ __aicore__ void fake_quant_affine_cachemask(