Star893
1805
代码
代码
Issues248
Pull Requests300
Actions
项目讨论
Wiki
分析
Star893
1805
已合并
新增IFMR（Input Feature Map Reconstruction）算子 #2153
fujun19创建于 2月27日
新增IFMR（Input Feature Map Reconstruction）算子 #2153
已合并
fujun19创建于 2月27日
从fujun19/ops-nnifmr合入到cann/ops-nnmaster
共 23 个文件变更+2816-0
Mdocs/zh/op_list.md+10-0
@@ -3063,6 +3063,16 @@
3063 <td>✓</td>3063 <td>✓</td>
3064 <td>AI Core</td>3064 <td>AI Core</td>
3065 <td>完成量化计算参数scale数据类型的转换，将FLOAT32的数据类型转换为硬件需要的UINT64类型。</td>3065 <td>完成量化计算参数scale数据类型的转换，将FLOAT32的数据类型转换为硬件需要的UINT64类型。</td>
3066+ </tr>
3067+ <tr>
3068+ <td>quant</td>
3069+ <td><a href="../../quant/ifmr/README.md">ifmr</a></td>
3070+ <td>✓</td>
3071+ <td>✓</td>
3072+ <td>✗</td>
3073+ <td>✓</td>
3074+ <td>AI Core</td>
3075+ <td>输入特征图重建的量化方法。</td>
3066 </tr>3076 </tr>
3067 <tr>3077 <tr>
3068 <td>rnn</td>3078 <td>rnn</td>
Aquant/ifmr/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start ops.quant.ifmr.CMakeLists.txt ")
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+if(NOT ENABLE_TEST AND NOT BENCHMARK)
14+ list(REMOVE_ITEM CURRENT_DIRS tests)
15+endif()
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/ifmr/README.md+130-0
@@ -0,0 +1,130 @@
1+# IFMR
2+ 
3+## 产品支持情况
C
Cchenjiao3月4日

缺少交付件op_list.md

likedislike
4+ 
5+| 产品 | 是否支持 |
6+| ---- | :----:|
7+|Ascend 950PR/Ascend 950DT|√|
W
Wwulinyu3月3日

这里的产品型号 确定是这样的吗 下面是A2 A3 上面是Ascend xxx?

likedislike
F
fujun19
3月3日 评论：
8+|Atlas A3 训练系列产品/Atlas A3 推理系列产品|×|
9+|Atlas A2 训练系列产品/Atlas A2 推理系列产品|×|
10+|Atlas 200I/500 A2 推理产品|×|
11+|Atlas 推理系列产品|×|
12+|Atlas 训练系列产品|×|
13+ 
14+## 功能说明
15+ 
16+输入特征图重建的量化方法。
17+ 
18+## 参数说明
19+ 
20+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
21+ <col style="width: 150px">
22+ <col style="width: 150px">
23+ <col style="width: 280px">
24+ <col style="width: 330px">
25+ <col style="width: 120px">
26+ </colgroup>
27+ <thead>
28+ <tr>
29+ <th>参数名</th>
30+ <th>输入/输出/属性</th>
31+ <th>描述</th>
32+ <th>数据类型</th>
33+ <th>数据格式</th>
34+ </tr></thead>
35+ <tbody>
36+ <tr>
37+ <td>data</td>
38+ <td>输入</td>
39+ <td>特征图的张量。</td>
40+ <td>FLOAT、FLOAT16</td>
41+ <td>ND</td>
42+ </tr>
43+ <tr>
44+ <td>data_min</td>
45+ <td>输入</td>
46+ <td>特征图最小值的张量。</td>
47+ <td>FLOAT、FLOAT16</td>
48+ <td>ND</td>
49+ </tr>
50+ <tr>
51+ <td>data_max</td>
52+ <td>输入</td>
53+ <td>特征图最大值的张量。</td>
54+ <td>FLOAT、FLOAT16</td>
55+ <td>ND</td>
56+ </tr>
57+ <tr>
58+ <td>cumsum</td>
59+ <td>输入</td>
60+ <td>数据的累积求和的张量。</td>
61+ <td>INT32</td>
62+ <td>ND</td>
63+ </tr>
64+ <tr>
65+ <td>min_percentile</td>
66+ <td>属性</td>
67+ <td>最小初始化百分位数。</td>
68+ <td>FLOAT</td>
69+ <td>-</td>
70+ </tr>
71+ <tr>
72+ <td>max_percentile</td>
73+ <td>属性</td>
74+ <td>最大初始化百分位数。</td>
75+ <td>FLOAT</td>
76+ <td>-</td>
77+ </tr>
78+ <tr>
79+ <td>search_range</td>
80+ <td>属性</td>
81+ <td>搜索范围。</td>
82+ <td>LISTFLOAT</td>
83+ <td>-</td>
84+ </tr>
85+ <tr>
86+ <td>search_step</td>
87+ <td>属性</td>
88+ <td>搜索步长。</td>
89+ <td>FLOAT</td>
90+ <td>-</td>
91+ </tr>
92+ <tr>
93+ <td>with_offset</td>
94+ <td>属性</td>
95+ <td>是否使用偏移量。</td>
96+ <td>BOOL</td>
97+ <td>-</td>
98+ </tr>
99+ <tr>
100+ <td>quant_bits</td>
101+ <td>属性</td>
102+ <td>量化位数，可选属性，默认值为8。</td>
103+ <td>INT</td>
104+ <td>-</td>
105+ </tr>
106+ <tr>
107+ <td>scale</td>
108+ <td>输出</td>
109+ <td>最优尺度。</td>
110+ <td>FLOAT</td>
111+ <td>ND</td>
112+ </tr>
113+ <tr>
114+ <td>offset</td>
115+ <td>输出</td>
116+ <td>最优偏移量。</td>
117+ <td>FLOAT</td>
118+ <td>ND</td>
119+ </tr>
120+ </tbody></table>
121+ 
122+## 约束说明
123+ 
124+无
125+ 
126+## 调用说明
127+ 
128+| 调用方式 | 样例代码 | 说明 |
129+| ---------------- | --------------------------- | --------------------------------------------------- |
130+| 图模式 | [test_geir_ifmr](./examples/test_geir_ifmr.cpp) | 通过[算子IR](./op_graph/ifmr_proto.h)构图方式调用IFMR算子。 |
Aquant/ifmr/examples/test_geir_ifmr.cpp+332-0
@@ -0,0 +1,332 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <fstream>
13+#include <string.h>
14+#include <stdint.h>
15+#include <vector>
16+#include <string>
17+#include <map>
18+#include "assert.h"
19+ 
20+#include "graph.h"
21+#include "types.h"
22+#include "tensor.h"
23+#include "ge_error_codes.h"
24+#include "ge_api_types.h"
25+#include "ge_api.h"
26+#include "array_ops.h"
27+#include "ge_ir_build.h"
28+ 
29+#include "experiment_ops.h"
30+#include "nn_other.h"
31+#include "../op_graph/ifmr_proto.h"
32+ 
33+#define FAILED -1
34+#define SUCCESS 0
35+ 
36+using namespace ge;
37+using std::map;
38+using std::string;
39+using std::vector;
40+ 
41+const int B = 24; // Batch size
42+const int S = 1; // Sequence length
43+const int H = 128; // Hidden size
44+ 
45+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
46+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
47+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
48+ TensorDesc placeholder##intputIndex##_desc = \
49+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
50+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
51+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
52+ Tensor tensor_placeholder##intputIndex; \
53+ ret = GenOnesData(placeholder##intputIndex##_shape, \
54+ tensor_placeholder##intputIndex, \
55+ placeholder##intputIndex##_desc, \
56+ intputDtype, \
57+ 2); \
58+ if (ret != SUCCESS) { \
59+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
60+ return FAILED; \
61+ } \
62+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
63+ input.push_back(tensor_placeholder##intputIndex); \
64+ graph.AddOp(placeholder##intputIndex); \
65+ ifmr.set_input_##intputName(placeholder##intputIndex); \
66+ inputs.push_back(placeholder##intputIndex);
67+ 
68+#define ADD_INPUT_ATTR(attrName, attrValue) \
69+ ifmr.set_attr_##attrName(attrValue);
70+ 
71+#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape) \
72+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
73+ auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \
74+ TensorDesc placeholder##intputIndex##_desc = \
75+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
76+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
77+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
78+ Tensor tensor_placeholder##intputIndex; \
79+ ret = GenOnesData(placeholder##intputIndex##_shape, \
80+ tensor_placeholder##intputIndex, \
81+ placeholder##intputIndex##_desc, \
82+ intputDtype, \
83+ 2); \
84+ if (ret != SUCCESS) { \
85+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
86+ return FAILED; \
87+ } \
88+ placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \
89+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
90+ graph.AddOp(placeholder##intputIndex); \
91+ ifmr.set_input_##intputName(placeholder##intputIndex); \
92+ ifmr.update_input_desc_##intputName(placeholder##intputIndex##_desc); \
93+ inputs.push_back(placeholder##intputIndex);
94+ 
95+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
96+ TensorDesc outputName##outputIndex##_desc = \
97+ TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
98+ ifmr.update_output_desc_##outputName(outputName##outputIndex##_desc);
99+ 
100+#define LOG_PRINT(message, ...) \
101+ do { \
102+ printf(message, ##__VA_ARGS__); \
103+ } while (0)
104+ 
105+string GetTime()
106+{
107+ time_t timep;
108+ time(&timep);
109+ char tmp[64];
110+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
111+ return tmp;
112+}
113+ 
114+uint32_t GetDataTypeSize(DataType dt)
115+{
116+ uint32_t dilation = 1;
117+ uint32_t oneByte = 1;
118+ uint32_t twoByte = 2;
119+ uint32_t fourByte = 4;
120+ uint32_t eightByte = 8;
121+ 
122+ if (dt == ge::DT_FLOAT) {
123+ dilation = fourByte;
124+ } else if (dt == ge::DT_FLOAT16) {
125+ dilation = twoByte;
126+ } else if (dt == ge::DT_BF16) {
127+ dilation = twoByte;
128+ } else if (dt == ge::DT_INT16) {
129+ dilation = twoByte;
130+ } else if (dt == ge::DT_UINT16) {
131+ dilation = twoByte;
132+ } else if (dt == ge::DT_INT32) {
133+ dilation = fourByte;
134+ } else if (dt == ge::DT_UINT32) {
135+ dilation = fourByte;
136+ } else if (dt == ge::DT_INT64) {
137+ dilation = eightByte;
138+ } else if (dt == ge::DT_UINT64) {
139+ dilation = eightByte;
140+ } else if (dt == ge::DT_INT8) {
141+ dilation = oneByte;
142+ }
143+ return dilation;
144+}
145+ 
146+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, float value)
147+{
148+ input_tensor_desc.SetRealDimCnt(shapes.size());
149+ size_t size = 1;
150+ for (uint32_t i = 0; i < shapes.size(); i++) {
151+ size *= shapes[i];
152+ }
153+ uint32_t byteSizeFloat32 = 4;
154+ uint32_t data_len = size * byteSizeFloat32;
155+ float *pData = new (std::nothrow) float[size];
156+ 
157+ for (size_t i = 0; i < size; ++i) {
158+ *(pData + i) = value;
159+ }
160+ input_tensor = Tensor(input_tensor_desc, (uint8_t *)pData, data_len);
161+ return SUCCESS;
162+}
163+ 
164+int32_t GenOnesData(
165+ vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, DataType data_type, int value)
166+{
167+ input_tensor_desc.SetRealDimCnt(shapes.size());
168+ size_t size = 1;
169+ for (uint32_t i = 0; i < shapes.size(); i++) {
170+ size *= shapes[i];
171+ }
172+ uint32_t data_len = size * GetDataTypeSize(data_type);
173+ int32_t *pData = new (std::nothrow) int32_t[data_len];
174+ for (uint32_t i = 0; i < size; ++i) {
175+ *(pData + i) = value;
176+ }
177+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t *>(pData), data_len);
178+ return SUCCESS;
179+}
180+ 
181+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t *inputData)
182+{
183+ FILE *fp;
184+ fp = fopen(bin_file.c_str(), "w");
185+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
186+ fclose(fp);
187+ return SUCCESS;
188+}
189+ 
190+int CreateOppInGraph(std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
191+ std::vector<Operator> &outputs, Graph &graph)
192+{
193+ Status ret = SUCCESS;
194+ // 自定义代码：添加单算子定义到图中
195+ auto ifmr = op::IFMR("test_geir_ifmr");
196+ 
197+ // shape定义
198+ std::vector<int64_t> data = {1024};
199+ std::vector<int64_t> data_min = {1};
200+ std::vector<int64_t> data_max = {1};
201+ std::vector<int64_t> cumsum = {1024};
202+ 
203+ std::vector<int64_t> scale = {1};
204+ std::vector<int64_t> offset = {1};
205+ 
206+ // 添加输入（顺序严格匹配 proto.h）
207+ ADD_INPUT(1, data, DT_FLOAT16, data);
208+ ADD_INPUT(2, data_min, DT_FLOAT16, data_min);
209+ ADD_INPUT(3, data_max, DT_FLOAT16, data_max);
210+ ADD_INPUT(4, cumsum, DT_INT32, cumsum);
211+ 
212+ // 添加必选属性（顺序严格匹配 proto.h）
213+ ADD_INPUT_ATTR(min_percentile, 0.9);
214+ ADD_INPUT_ATTR(max_percentile, 0.9);
215+ ADD_INPUT_ATTR(search_range, {0.7, 1.3});
216+ ADD_INPUT_ATTR(search_step, 0.1);
217+ ADD_INPUT_ATTR(with_offset, true);
218+ // 添加输出（顺序严格匹配 proto.h）
219+ ADD_OUTPUT(1, scale, DT_FLOAT, var_out_shape);
220+ ADD_OUTPUT(2, offset, DT_FLOAT, scale_out_shape);
221+ 
222+ outputs.push_back(ifmr);
223+ // 添加完毕
224+ return SUCCESS;
225+}
226+ 
227+int main(int argc, char *argv[])
228+{
229+ const char *graph_name = "tc_ge_irrun_test";
230+ Graph graph(graph_name);
231+ std::vector<ge::Tensor> input;
232+ 
233+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
234+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
235+ Status ret = ge::GEInitialize(global_options);
236+ if (ret != SUCCESS) {
237+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
238+ return FAILED;
239+ }
240+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
241+ 
242+ std::vector<Operator> inputs{};
243+ std::vector<Operator> outputs{};
244+ 
245+ std::cout << argv[1] << std::endl;
246+ char *endptr;
247+ 
248+ ret = CreateOppInGraph(input, inputs, outputs, graph);
249+ if (ret != SUCCESS) {
250+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
251+ return FAILED;
252+ }
253+ 
254+ if (!inputs.empty() && !outputs.empty()) {
255+ graph.SetInputs(inputs).SetOutputs(outputs);
256+ }
257+ 
258+ std::map<AscendString, AscendString> build_options = {
259+ 
260+ };
261+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
262+ ge::Session *session = new Session(build_options);
263+ 
264+ if (session == nullptr) {
265+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
266+ return FAILED;
267+ }
268+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
269+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
270+ 
271+ std::map<AscendString, AscendString> graph_options = {
272+ 
273+ };
274+ uint32_t graph_id = 0;
275+ ret = session->AddGraph(graph_id, graph, graph_options);
276+ 
277+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
278+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
279+ std::string file_path = "./dump";
280+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
281+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
282+ std::vector<ge::Tensor> output;
283+ ret = session->RunGraph(graph_id, input, output);
284+ if (ret != SUCCESS) {
285+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
286+ delete session;
287+ GEFinalize();
288+ return FAILED;
289+ }
290+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
291+ 
292+ int input_num = input.size();
293+ for (int i = 0; i < input_num; i++) {
294+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
295+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
296+ uint8_t *input_data_i = input[i].GetData();
297+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
298+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
299+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
300+ WriteDataToFile((const char *)input_file.c_str(), data_size, input_data_i);
301+ }
302+ 
303+ int output_num = output.size();
304+ for (int i = 0; i < output_num; i++) {
305+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
306+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
307+ uint8_t *output_data_i = output[i].GetData();
308+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
309+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
310+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
311+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
312+ int32_t *result = (int32_t*)output_data_i;
313+ for (int64_t j = 0; j < output_shape; j++) {
314+ LOG_PRINT("result[%ld] is: %d\n", j, result[j]);
315+ }
316+ }
317+ 
318+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
319+ std::string error_str(error_msg.GetString());
320+ std::cout << "Error message: " << error_str << std::endl;
321+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
322+ std::string warning_str(warning_msg.GetString());
323+ std::cout << "Warning message: " << warning_str << std::endl;
324+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
325+ ret = ge::GEFinalize();
326+ if (ret != SUCCESS) {
327+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
328+ return FAILED;
329+ }
330+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
331+ return SUCCESS;
332+}
Aquant/ifmr/op_graph/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+#/
9+ 
10+message(STATUS "=== Debug: start ifmr.op_host.CMakeLists.txt ")
11+add_graph_plugin_sources()
Aquant/ifmr/op_graph/ifmr_proto.h+58-0
@@ -0,0 +1,58 @@
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+#ifndef IFMR_PROTO_H
11+#define IFMR_PROTO_H
12+ 
13+#include "graph/operator_reg.h"
14+#include "graph/operator.h"
15+ 
16+namespace ge {
17+/**
18+* @brief IFMR(Input Feature Map Reconstruction).
19+ 
20+* @par Inputs:
21+* @li data: A Tensor of feature map.
22+* @li data_min: A Tensor of min value of feature map.
23+* @li data_max: A Tensor of max value of feature map.
24+* @li cumsum: A Tensor of cumsum bin of data . \n
25+ 
26+* @par Attributes:
27+* @li min_percentile: Min init percentile.
28+* @li max_percentile: Max init percentile.
29+* @li search_range: Search range.
30+* @li search_step: Step size of searching.
31+* @li with_offset: Whether using offset.
32+* @li quant_bits: Bits of quant, an optional attr, default value is 8. \n
33+ 
34+* @par Outputs:
35+* @li scale: Optimal scale.
36+* @li offset: Optimal offset. \n
37+ 
38+* @par Third-party framework compatibility
39+* Compatible with mindspore.
40+*/
41+ 
42+REG_OP(IFMR)
43+ .INPUT(data, TensorType({DT_FLOAT16, DT_FLOAT}))
44+ .INPUT(data_min, TensorType({DT_FLOAT16, DT_FLOAT}))
45+ .INPUT(data_max, TensorType({DT_FLOAT16, DT_FLOAT}))
46+ .INPUT(cumsum, TensorType({DT_INT32}))
47+ .OUTPUT(scale, TensorType({DT_FLOAT}))
48+ .OUTPUT(offset, TensorType({DT_FLOAT}))
49+ .REQUIRED_ATTR(min_percentile, Float)
50+ .REQUIRED_ATTR(max_percentile, Float)
51+ .REQUIRED_ATTR(search_range, ListFloat)
52+ .REQUIRED_ATTR(search_step, Float)
53+ .REQUIRED_ATTR(with_offset, Bool)
54+ .ATTR(quant_bits, Int, 8)
55+ .OP_END_FACTORY_REG(IFMR)
56+} // namespace ge
57+ 
58+#endif // IFMR_PROTO_H
Aquant/ifmr/op_host/CMakeLists.txt+13-0
@@ -0,0 +1,13 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start ops.quant.ifmr.op_host.CMakeLists.txt ")
12+ 
13+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE ifmr ACLNNTYPE aclnn)
Aquant/ifmr/op_host/config/ascend950/ifmr_binary.json+402-0
@@ -0,0 +1,402 @@
1+{
2+ "op_type": "IFMR",
3+ "op_list": [
4+ {
5+ "bin_filename": "IFMR_float16_with_offset",
6+ "inputs": [
7+ {
8+ "name": "data",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "data_min",
19+ "index": 1,
20+ "dtype": "float16",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ },
27+ {
28+ "name": "data_max",
29+ "index": 2,
30+ "dtype": "float16",
31+ "format": "ND",
32+ "paramType": "required",
33+ "shape": [
34+ -2
35+ ]
36+ },
37+ {
38+ "name": "cumsum",
39+ "index": 3,
40+ "dtype": "int32",
41+ "format": "ND",
42+ "paramType": "required",
43+ "shape": [
44+ -2
45+ ]
46+ }
47+ ],
48+ "outputs": [
49+ {
50+ "name": "scale",
51+ "index": 0,
52+ "dtype": "float32",
53+ "format": "ND",
54+ "paramType": "required",
55+ "shape": [
56+ -2
57+ ]
58+ },
59+ {
60+ "name": "offset",
61+ "index": 1,
62+ "dtype": "float32",
63+ "format": "ND",
64+ "paramType": "required",
65+ "shape": [
66+ -2
67+ ]
68+ }
69+ ],
70+ "attrs": [
71+ {
72+ "name": "min_percentile",
73+ "dtype": "float32",
74+ "value": -1.0
75+ },
76+ {
77+ "name": "max_percentile",
78+ "dtype": "float32",
79+ "value": -1.0
80+ },
81+ {
82+ "name": "search_range",
83+ "dtype": "list_float",
84+ "value": [-1.0, -1.0]
85+ },
86+ {
87+ "name": "search_step",
88+ "dtype": "float",
89+ "value": -1.0
90+ },
91+ {
92+ "name": "with_offset",
93+ "dtype": "bool",
94+ "value": true
95+ },
96+ {
97+ "name": "quant_bits",
98+ "dtype": "int",
99+ "value": -1
100+ }
101+ ]
102+ },
103+ {
104+ "bin_filename": "IFMR_float16_without_offset",
105+ "inputs": [
106+ {
107+ "name": "data",
108+ "index": 0,
109+ "dtype": "float16",
110+ "format": "ND",
111+ "paramType": "required",
112+ "shape": [
113+ -2
114+ ]
115+ },
116+ {
117+ "name": "data_min",
118+ "index": 1,
119+ "dtype": "float16",
120+ "format": "ND",
121+ "paramType": "required",
122+ "shape": [
123+ -2
124+ ]
125+ },
126+ {
127+ "name": "data_max",
128+ "index": 2,
129+ "dtype": "float16",
130+ "format": "ND",
131+ "paramType": "required",
132+ "shape": [
133+ -2
134+ ]
135+ },
136+ {
137+ "name": "cumsum",
138+ "index": 3,
139+ "dtype": "int32",
140+ "format": "ND",
141+ "paramType": "required",
142+ "shape": [
143+ -2
144+ ]
145+ }
146+ ],
147+ "outputs": [
148+ {
149+ "name": "scale",
150+ "index": 0,
151+ "dtype": "float32",
152+ "format": "ND",
153+ "paramType": "required",
154+ "shape": [
155+ -2
156+ ]
157+ },
158+ {
159+ "name": "offset",
160+ "index": 1,
161+ "dtype": "float32",
162+ "format": "ND",
163+ "paramType": "required",
164+ "shape": [
165+ -2
166+ ]
167+ }
168+ ],
169+ "attrs": [
170+ {
171+ "name": "min_percentile",
172+ "dtype": "float32",
173+ "value": -1.0
174+ },
175+ {
176+ "name": "max_percentile",
177+ "dtype": "float32",
178+ "value": -1.0
179+ },
180+ {
181+ "name": "search_range",
182+ "dtype": "list_float",
183+ "value": [-1.0, -1.0]
184+ },
185+ {
186+ "name": "search_step",
187+ "dtype": "float",
188+ "value": -1.0
189+ },
190+ {
191+ "name": "with_offset",
192+ "dtype": "bool",
193+ "value": false
194+ },
195+ {
196+ "name": "quant_bits",
197+ "dtype": "int",
198+ "value": -1
199+ }
200+ ]
201+ },
202+ {
203+ "bin_filename": "IFMR_float32_with_offset",
204+ "inputs": [
205+ {
206+ "name": "data",
207+ "index": 0,
208+ "dtype": "float32",
209+ "format": "ND",
210+ "paramType": "required",
211+ "shape": [
212+ -2
213+ ]
214+ },
215+ {
216+ "name": "data_min",
217+ "index": 1,
218+ "dtype": "float32",
219+ "format": "ND",
220+ "paramType": "required",
221+ "shape": [
222+ -2
223+ ]
224+ },
225+ {
226+ "name": "data_max",
227+ "index": 2,
228+ "dtype": "float32",
229+ "format": "ND",
230+ "paramType": "required",
231+ "shape": [
232+ -2
233+ ]
234+ },
235+ {
236+ "name": "cumsum",
237+ "index": 3,
238+ "dtype": "int32",
239+ "format": "ND",
240+ "paramType": "required",
241+ "shape": [
242+ -2
243+ ]
244+ }
245+ ],
246+ "outputs": [
247+ {
248+ "name": "scale",
249+ "index": 0,
250+ "dtype": "float32",
251+ "format": "ND",
252+ "paramType": "required",
253+ "shape": [
254+ -2
255+ ]
256+ },
257+ {
258+ "name": "offset",
259+ "index": 1,
260+ "dtype": "float32",
261+ "format": "ND",
262+ "paramType": "required",
263+ "shape": [
264+ -2
265+ ]
266+ }
267+ ],
268+ "attrs": [
269+ {
270+ "name": "min_percentile",
271+ "dtype": "float32",
272+ "value": -1.0
273+ },
274+ {
275+ "name": "max_percentile",
276+ "dtype": "float32",
277+ "value": -1.0
278+ },
279+ {
280+ "name": "search_range",
281+ "dtype": "list_float",
282+ "value": [-1.0, -1.0]
283+ },
284+ {
285+ "name": "search_step",
286+ "dtype": "float",
287+ "value": -1.0
288+ },
289+ {
290+ "name": "with_offset",
291+ "dtype": "bool",
292+ "value": true
293+ },
294+ {
295+ "name": "quant_bits",
296+ "dtype": "int",
297+ "value": -1
298+ }
299+ ]
300+ },
301+ {
302+ "bin_filename": "IFMR_float32_without_offset",
303+ "inputs": [
304+ {
305+ "name": "data",
306+ "index": 0,
307+ "dtype": "float32",
308+ "format": "ND",
309+ "paramType": "required",
310+ "shape": [
311+ -2
312+ ]
313+ },
314+ {
315+ "name": "data_min",
316+ "index": 1,
317+ "dtype": "float32",
318+ "format": "ND",
319+ "paramType": "required",
320+ "shape": [
321+ -2
322+ ]
323+ },
324+ {
325+ "name": "data_max",
326+ "index": 2,
327+ "dtype": "float32",
328+ "format": "ND",
329+ "paramType": "required",
330+ "shape": [
331+ -2
332+ ]
333+ },
334+ {
335+ "name": "cumsum",
336+ "index": 3,
337+ "dtype": "int32",
338+ "format": "ND",
339+ "paramType": "required",
340+ "shape": [
341+ -2
342+ ]
343+ }
344+ ],
345+ "outputs": [
346+ {
347+ "name": "scale",
348+ "index": 0,
349+ "dtype": "float32",
350+ "format": "ND",
351+ "paramType": "required",
352+ "shape": [
353+ -2
354+ ]
355+ },
356+ {
357+ "name": "offset",
358+ "index": 1,
359+ "dtype": "float32",
360+ "format": "ND",
361+ "paramType": "required",
362+ "shape": [
363+ -2
364+ ]
365+ }
366+ ],
367+ "attrs": [
368+ {
369+ "name": "min_percentile",
370+ "dtype": "float32",
371+ "value": -1.0
372+ },
373+ {
374+ "name": "max_percentile",
375+ "dtype": "float32",
376+ "value": -1.0
377+ },
378+ {
379+ "name": "search_range",
380+ "dtype": "list_float",
381+ "value": [-1.0, -1.0]
382+ },
383+ {
384+ "name": "search_step",
385+ "dtype": "float",
386+ "value": -1.0
387+ },
388+ {
389+ "name": "with_offset",
390+ "dtype": "bool",
391+ "value": false
392+ },
393+ {
394+ "name": "quant_bits",
395+ "dtype": "int",
396+ "value": -1
397+ }
398+ ]
399+ }
400+ ]
401+ }
402+
Aquant/ifmr/op_host/config/ascend950/ifmr_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascnedxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[IFMR]
13+simplified_key_mode=0
Aquant/ifmr/op_host/ifmr_def.cpp+68-0
@@ -0,0 +1,68 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ifmr.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class IFMR : public OpDef {
19+public:
20+ explicit IFMR(const char* name) : OpDef(name)
21+ {
22+ this->Input("data")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
27+ this->Input("data_min")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
32+ this->Input("data_max")
33+ .ParamType(REQUIRED)
34+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
35+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
36+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
37+ this->Input("cumsum")
38+ .ParamType(REQUIRED)
39+ .DataType({ge::DT_INT32, ge::DT_INT32})
40+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
41+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
42+ this->Output("scale")
43+ .ParamType(REQUIRED)
44+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT})
45+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
46+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
47+ this->Output("offset")
48+ .ParamType(REQUIRED)
49+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT})
50+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
51+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
52+ this->Attr("min_percentile").AttrType(REQUIRED).Float();
53+ this->Attr("max_percentile").AttrType(REQUIRED).Float();
54+ this->Attr("search_range").AttrType(REQUIRED).ListFloat();
55+ this->Attr("search_step").AttrType(REQUIRED).Float();
56+ this->Attr("with_offset").AttrType(REQUIRED).Bool();
57+ this->Attr("quant_bits").AttrType(OPTIONAL).Int();
58+ 
59+ OpAICoreConfig aicoreConfig;
60+ aicoreConfig.DynamicCompileStaticFlag(true)
61+ .DynamicRankSupportFlag(true)
62+ .DynamicShapeSupportFlag(true)
63+ .NeedCheckSupportFlag(false);
64+ this->AICore().AddConfig("ascend950", aicoreConfig);
65+ }
66+};
67+OP_ADD(IFMR); // 添加算子信息库
68+} // namespace ops
Aquant/ifmr/op_host/ifmr_infershape.cpp+38-0
@@ -0,0 +1,38 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ifmr_infershape.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+ 
20+namespace ops {
21+static constexpr int64_t IDX_0 = 0;
22+ 
23+static ge::graphStatus InferShapeIfmr(gert::InferShapeContext* context)
24+{
25+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeIfmr");
26+ auto outShape = context->GetOutputShape(0);
27+ OP_CHECK_NULL_WITH_CONTEXT(context, outShape);
28+ outShape->SetDimNum(0);
29+ outShape->AppendDim(1);
30+ outShape = context->GetOutputShape(1);
31+ outShape->SetDimNum(0);
32+ outShape->AppendDim(1);
33+ OP_LOGD(context->GetNodeName(), "End to do InferShapeIfmr");
34+ return GRAPH_SUCCESS;
35+}
36+ 
37+IMPL_OP_INFERSHAPE(IFMR).InferShape(InferShapeIfmr);
38+} // namespace ops
Aquant/ifmr/op_host/ifmr_tiling.cpp+297-0
@@ -0,0 +1,297 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ifmr_tiling.cpp
13+ * \brief
14+ */
15+#include "ifmr_tiling.h"
16+#include "log/log.h"
17+#include "error_util.h"
18+#include "op_host/tiling_util.h"
19+#include "op_host/tiling_templates_registry.h"
20+#include "ifmr/op_kernel/ifmr_tiling_data.h"
21+ 
22+namespace optiling {
23+ 
24+struct IfmrCompileInfo {
25+ int32_t dataNum;
26+};
27+ 
28+ge::graphStatus IfmrTiling::CheckIfmrTilingAttrs(void)
29+{
30+ if (attrs_.minPercentile <= K_PERCENTILE_LOW_BOUND || attrs_.minPercentile > K_PERCENTILE_UPPER_BOUND) {
31+ OP_LOGE(nodeName_, "The minPercentile must be greater than 0.5 "\
32+ "and less than or equal to 1.0");
33+ return ge::GRAPH_FAILED;
34+ }
35+ if (attrs_.maxPercentile <= K_PERCENTILE_LOW_BOUND || attrs_.maxPercentile > K_PERCENTILE_UPPER_BOUND) {
36+ OP_LOGE(nodeName_, "The maxPercentile must be greater than 0.5 "\
37+ "and less than or equal to 1.0");
38+ return ge::GRAPH_FAILED;
39+ }
40+ if (attrs_.searchRange[0] <= 0) {
41+ OP_LOGE(nodeName_, "search_start must be greater than zero.");
42+ return ge::GRAPH_FAILED;
43+ }
44+ if (attrs_.searchRange[0] >= attrs_.searchRange[1]) {
45+ OP_LOGE(nodeName_, "search_start must be less than search_end.");
46+ return ge::GRAPH_FAILED;
47+ }
48+ if (attrs_.searchStep <= 0) {
49+ OP_LOGE(nodeName_, "The searchStep must be greater than zero.");
50+ return ge::GRAPH_FAILED;
51+ }
52+ if (std::round((attrs_.searchRange[1] - attrs_.searchRange[0]) / attrs_.searchStep) + 1 > MAX_STEP_NUMS) {
53+ OP_LOGE(nodeName_, "step size should be equal or less than 4096");
54+ return ge::GRAPH_FAILED;
55+ }
56+ if (std::find(SUPPORTED_QUANT_BITS, (SUPPORTED_QUANT_BITS + SUPPORT_QUANT_BITS_NUM), attrs_.quantBits) ==
57+ (SUPPORTED_QUANT_BITS + SUPPORT_QUANT_BITS_NUM)) {
58+ OP_LOGE(nodeName_, "quant bits only support 8 or 16");
59+ return ge::GRAPH_FAILED;
60+ }
61+ return ge::GRAPH_SUCCESS;
62+}
63+ge::graphStatus IfmrTiling::GetIfmrTilingAttrInfo(void)
64+{
65+ OP_LOGD(nodeName_, "[IFMR] GetIfmrTilingAttrInfo start running");
66+ OP_CHECK_NULL_WITH_CONTEXT(context_, context_->GetAttrs());
67+ auto minPercentilPtr = context_->GetAttrs()->GetFloat(ATTR_MIN_PERCENTILE_INDEX);
68+ auto maxPercentilPtr = context_->GetAttrs()->GetFloat(ATTR_MAX_PERCENTILE_INDEX);
69+ auto searchRangePtr = context_->GetAttrs()->GetListFloat(ATTR_SEARCH_RANGE_INDEX);
70+ auto searchStepPtr = context_->GetAttrs()->GetFloat(ATTR_SEARCH_STEP_INDEX);
71+ auto withOffsetPtr = context_->GetAttrs()->GetBool(ATTR_WITH_OFFSET_INDEX);
72+ auto quantBitsPtr = context_->GetAttrs()->GetInt(ATTR_QUANT_BITS_INDEX);
73+ OP_CHECK_NULL_WITH_CONTEXT(context_, minPercentilPtr);
74+ OP_CHECK_NULL_WITH_CONTEXT(context_, maxPercentilPtr);
75+ OP_CHECK_NULL_WITH_CONTEXT(context_, searchRangePtr);
76+ OP_CHECK_NULL_WITH_CONTEXT(context_, searchStepPtr);
77+ OP_CHECK_NULL_WITH_CONTEXT(context_, withOffsetPtr);
78+ OP_CHECK_NULL_WITH_CONTEXT(context_, quantBitsPtr);
79+ attrs_.minPercentile = *minPercentilPtr;
80+ attrs_.maxPercentile = *maxPercentilPtr;
81+ if (searchRangePtr->GetSize() != ATTR_SEARCH_RANGE_SIZE) {
82+ OP_LOGE(nodeName_, "[IFMR] step_range dim should be 2");
83+ return ge::GRAPH_FAILED;
84+ }
85+ attrs_.searchRange[0] = searchRangePtr->GetData()[0];
86+ attrs_.searchRange[1] = searchRangePtr->GetData()[1];
87+ attrs_.searchStep = *searchStepPtr;
88+ attrs_.withOffset = *withOffsetPtr;
89+ attrs_.quantBits = *quantBitsPtr;
90+ return CheckIfmrTilingAttrs();
91+}
92+ge::graphStatus IfmrTiling::CheckIfmrTilingInputDataShape(uint32_t inputIndex, std::string inputName)
93+{
94+ auto dataShapePtr = context_->GetInputShape(inputIndex);
95+ OP_CHECK_NULL_WITH_CONTEXT(context_, dataShapePtr);
96+ auto dataShape = dataShapePtr->GetStorageShape();
97+ auto dataDataDim = dataShape.GetDimNum();
98+ if (dataDataDim != 1 || dataShapePtr->GetStorageShape().GetDim(0) != 1) {
99+ OP_LOGE(nodeName_, "The shape of %s must be [1]!", inputName.c_str());
100+ return ge::GRAPH_FAILED;
101+ }
102+ return ge::GRAPH_SUCCESS;
103+}
104+ge::graphStatus IfmrTiling::CheckIfmrTilingInputDtype(void)
105+{
106+ // check input dtype
107+ auto dataDesc = context_->GetInputDesc(DATA_INPUT_INDEX);
108+ auto dataDtype = dataDesc->GetDataType();
109+ if (dataDtype != ge::DataType::DT_FLOAT && \
110+ dataDtype != ge::DataType::DT_FLOAT16) {
111+ OP_LOGE(nodeName_, "Input data support only DT_FLOAT and DT_FLOAT16!");
112+ return ge::GRAPH_FAILED;
113+ }
114+ auto dataMinDesc = context_->GetInputDesc(DATA_MIN_INPUT_INDEX);
115+ auto dataMinDtype = dataMinDesc->GetDataType();
116+ auto dataMaxDesc = context_->GetInputDesc(DATA_MAX_INPUT_INDEX);
117+ auto dataMaxDtype = dataMaxDesc->GetDataType();
118+ if (dataDtype != dataMaxDtype || dataDtype != dataMinDtype) {
119+ OP_LOGE(nodeName_, "Input data, data_max, data_min must have same data type!");
120+ return ge::GRAPH_FAILED;
121+ }
122+ auto cumsumDesc = context_->GetInputDesc(CUMSUM_INPUT_INDEX);
123+ if (cumsumDesc->GetDataType() != ge::DataType::DT_INT32) {
124+ OP_LOGE(nodeName_, "Input cumsum support only DT_INT32!");
125+ return ge::GRAPH_FAILED;
126+ }
127+ return ge::GRAPH_SUCCESS;
128+}
129+ge::graphStatus IfmrTiling::GetDataLength(void)
130+{
131+ auto dataShapePtr = context_->GetInputShape(DATA_INPUT_INDEX);
132+ OP_CHECK_NULL_WITH_CONTEXT(context_, dataShapePtr);
133+ auto dataShape = dataShapePtr->GetStorageShape();
134+ auto dataDataDim = dataShape.GetDimNum();
135+ uint64_t dataLength = 1;
136+ for (uint32_t i = 0; i < dataDataDim; i++) {
137+ int64_t shape = dataShapePtr->GetStorageShape().GetDim(i);
138+ if (shape <= 0) {
139+ OP_LOGE(nodeName_, "The input shape should be greater than 0!");
140+ return ge::GRAPH_FAILED;
141+ }
142+ uint64_t preDataLength = dataLength;
143+ dataLength *= shape;
144+ // DataLength overflow validity check
145+ if ((dataLength > SHAPE_SIZE_LIMIT) || ((shape != 0) && (dataLength / shape != preDataLength))) {
146+ OP_LOGE(nodeName_, "Excessive amount of input_data(more than 2^31)!");
147+ return ge::GRAPH_FAILED;
148+ }
149+ }
150+ attrs_.dataLength = dataLength;
151+ return ge::GRAPH_SUCCESS;
152+}
153+ge::graphStatus IfmrTiling::GetIfmrTilingInputInfo(void)
154+{
155+ // check input shape
156+ if (GetDataLength() != ge::GRAPH_SUCCESS) {
157+ return ge::GRAPH_FAILED;
158+ }
159+ auto cumsumShapePtr = context_->GetInputShape(CUMSUM_INPUT_INDEX);
160+ OP_CHECK_NULL_WITH_CONTEXT(context_, cumsumShapePtr);
161+ auto cumSumShape = cumsumShapePtr->GetStorageShape();
162+ auto cumsumDataDim = cumSumShape.GetDimNum();
163+ if (cumsumDataDim != 1) {
164+ OP_LOGE(nodeName_, "The shape of input_cumsum must be (x,)!");
165+ return ge::GRAPH_FAILED;
166+ }
167+ uint32_t cumsumLength = cumsumShapePtr->GetStorageShape().GetDim(0);
168+ if (cumsumLength > MAX_CUMSUM_LENGTH || cumsumLength == 0) {
169+ OP_LOGE(nodeName_, "Excessive amount of input_cumsum(more than 8192) or cumsum_len is 0!");
170+ return ge::GRAPH_FAILED;
171+ }
172+ attrs_.cumsumLength = cumsumLength;
173+ if (CheckIfmrTilingInputDataShape(DATA_MIN_INPUT_INDEX, "input_min") != ge::GRAPH_SUCCESS) {
174+ return ge::GRAPH_FAILED;
175+ }
176+ if (CheckIfmrTilingInputDataShape(DATA_MAX_INPUT_INDEX, "input_max") != ge::GRAPH_SUCCESS) {
177+ return ge::GRAPH_FAILED;
178+ }
179+ if (CheckIfmrTilingInputDtype() != ge::GRAPH_SUCCESS) {
180+ return ge::GRAPH_FAILED;
181+ }
182+ uint64_t ubSizePlatform;
183+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context_->GetPlatformInfo());
184+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatform);
185+ if (ubSizePlatform <= UB_SIZE_RESERVE) {
186+ OP_LOGE(nodeName_, "UB size is not enough!");
187+ return ge::GRAPH_FAILED;
188+ }
189+ return ge::GRAPH_SUCCESS;
190+}
191+ge::graphStatus IfmrTiling::CheckIfmrTilingOutputInfo(void)
192+{
193+ auto scaleDesc = context_->GetOutputDesc(0);
194+ auto offsetDesc = context_->GetOutputDesc(1);
195+ OP_TILING_CHECK(scaleDesc == nullptr, VECTOR_INNER_ERR_REPORT_TILIING(nodeName_, "scaleDesc cannot be nullptr!"),
196+ return ge::GRAPH_FAILED);
197+ OP_TILING_CHECK(offsetDesc == nullptr, VECTOR_INNER_ERR_REPORT_TILIING(nodeName_, "offsetDesc cannot be nullptr!"),
198+ return ge::GRAPH_FAILED);
199+ if (scaleDesc->GetDataType() != ge::DataType::DT_FLOAT) {
200+ OP_LOGE(nodeName_, "Output scale support only DT_FLOAT!");
201+ return ge::GRAPH_FAILED;
202+ }
203+ if (offsetDesc->GetDataType() != ge::DataType::DT_FLOAT) {
204+ OP_LOGE(nodeName_, "Output offset support only DT_FLOAT!");
205+ return ge::GRAPH_FAILED;
206+ }
207+ auto scaleShapePtr = context_->GetOutputShape(0);
208+ OP_CHECK_NULL_WITH_CONTEXT(context_, scaleShapePtr);
209+ auto scaleShape = scaleShapePtr->GetStorageShape();
210+ auto scaleDataDim = scaleShape.GetDimNum();
211+ if (scaleDataDim != 1 || scaleShapePtr->GetStorageShape().GetDim(0) != 1) {
212+ OP_LOGE(nodeName_, "The shape of scale must be [1]!");
213+ return ge::GRAPH_FAILED;
214+ }
215+ auto offsetShapePtr = context_->GetOutputShape(1);
216+ OP_CHECK_NULL_WITH_CONTEXT(context_, offsetShapePtr);
217+ auto offsetShape = offsetShapePtr->GetStorageShape();
218+ auto offsetDataDim = offsetShape.GetDimNum();
219+ if (offsetDataDim != 1 || offsetShapePtr->GetStorageShape().GetDim(0) != 1) {
220+ OP_LOGE(nodeName_, "The shape of offset must be [1]!");
221+ return ge::GRAPH_FAILED;
222+ }
223+ return ge::GRAPH_SUCCESS;
224+}
225+void IfmrTiling::SetIfmrTiling(void)
226+{
227+ OP_LOGD(nodeName_, "[IFMR] SetIfmrTiling start running");
228+ IfmrTilingData *tilingData = context_->GetTilingData<IfmrTilingData>();
229+ (void)memset_s(tilingData, sizeof(IfmrTilingData), 0, sizeof(IfmrTilingData));
230+ tilingData->minPercentile = attrs_.minPercentile;
231+ tilingData->maxPercentile = attrs_.maxPercentile;
232+ tilingData->searchRange[0] = attrs_.searchRange[0];
233+ tilingData->searchRange[1] = attrs_.searchRange[1];
234+ tilingData->searchStep = attrs_.searchStep;
235+ tilingData->withOffset = attrs_.withOffset;
236+ tilingData->quantBits = attrs_.quantBits;
237+ tilingData->dataLength = attrs_.dataLength;
238+ tilingData->cumsumLength = attrs_.cumsumLength;
239+ OP_LOGI("IFMR", "nodeName: %s, minPercentile: %lf, maxPercentile: %lf, searchRange: [%lf, %lf],\
240+ searchStep: %lf, withOffset: %d, quantBits: %d, dataLength: %d, cumsumLength: %d",
241+ nodeName_, tilingData->minPercentile, tilingData->maxPercentile,
242+ tilingData->searchRange[0], tilingData->searchRange[1], tilingData->searchStep,
243+ tilingData->withOffset, tilingData->quantBits,
244+ tilingData->dataLength, tilingData->cumsumLength);
245+ return;
246+}
247+void IfmrTiling::PostTiling(void)
248+{
249+ context_->SetTilingKey(0);
250+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context_->GetPlatformInfo());
251+ context_->SetBlockDim(ascendcPlatform.GetCoreNumAiv());
252+ size_t* workspaces = context_->GetWorkspaceSizes(1);
253+ size_t workspaceSize = static_cast<size_t>(20) * 1024 * 1024; // 16M for AscendC framework, 4M reserved for ifmr op
254+ workspaces[0] = workspaceSize;
255+ OP_LOGD(nodeName_, "[IFMR] PostTiling run completed");
256+}
257+ 
258+// tiling 分发入口
259+ge::graphStatus IfmrTiling::IfmrTilingFunc(void)
260+{
261+ OP_LOGD(nodeName_, "[IFMR] RunIfmrTiling start running");
262+ 
263+ auto ret = GetIfmrTilingAttrInfo();
264+ if (ret != ge::GRAPH_SUCCESS) {
265+ return ret;
266+ }
267+ ret = GetIfmrTilingInputInfo();
268+ if (ret != ge::GRAPH_SUCCESS) {
269+ return ret;
270+ }
271+ ret = CheckIfmrTilingOutputInfo();
272+ if (ret != ge::GRAPH_SUCCESS) {
273+ return ret;
274+ }
275+ SetIfmrTiling();
276+ PostTiling();
277+ return ret;
278+}
279+ 
280+static ge::graphStatus TilingForIfmr(gert::TilingContext* context)
281+{
282+ OP_TILING_CHECK(context == nullptr, VECTOR_INNER_ERR_REPORT_TILIING("IFMR", "context should not be nullptr."),
283+ return ge::GRAPH_FAILED);
284+ IfmrTiling tiling(context);
285+ return tiling.IfmrTilingFunc();
286+}
287+ 
288+static ge::graphStatus TilingParseForIfmr(gert::TilingParseContext* context)
289+{
290+ // The operator does not need compile_info, return success
291+ (void)context;
292+ return ge::GRAPH_SUCCESS;
293+}
294+ 
295+// tiling注册入口.
296+IMPL_OP_OPTILING(IFMR).Tiling(TilingForIfmr).TilingParse<IfmrCompileInfo>(TilingParseForIfmr);
297+} // namespace optiling
Aquant/ifmr/op_host/ifmr_tiling.h+86-0
@@ -0,0 +1,86 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ifmr_tiling.h
13+ * \brief
14+ */
15+ 
16+#ifndef IFMR_TILING_H
17+#define IFMR_TILING_H
18+ 
19+#include <cstdint>
20+#include <string>
21+#include "register/tilingdata_base.h"
22+#include "tiling/tiling_api.h"
23+#include "op_host/tiling_base.h"
24+#include "quant/ifmr/op_kernel/ifmr_tiling_data.h"
25+ 
26+namespace optiling {
27+ constexpr int ATTR_MIN_PERCENTILE_INDEX = 0;
28+ constexpr int ATTR_MAX_PERCENTILE_INDEX = 1;
29+ constexpr int ATTR_SEARCH_RANGE_INDEX = 2;
30+ constexpr int ATTR_SEARCH_STEP_INDEX = 3;
31+ constexpr int ATTR_WITH_OFFSET_INDEX = 4;
32+ constexpr int ATTR_QUANT_BITS_INDEX = 5;
33+ 
34+ constexpr int ATTR_SEARCH_RANGE_SIZE = 2;
35+ 
36+ constexpr float K_PERCENTILE_LOW_BOUND = 0.5f;
37+ constexpr float K_PERCENTILE_UPPER_BOUND = 1.0f;
38+ 
39+ constexpr uint32_t DATA_INPUT_INDEX = 0;
40+ constexpr uint32_t DATA_MIN_INPUT_INDEX = 1;
41+ constexpr uint32_t DATA_MAX_INPUT_INDEX = 2;
42+ constexpr uint32_t CUMSUM_INPUT_INDEX = 3;
43+ constexpr uint32_t SHAPE_SIZE_LIMIT = 2147483648; //pow(2, 31);
44+ constexpr uint32_t MAX_CUMSUM_LENGTH = 8192;
45+ constexpr uint32_t MAX_STEP_NUMS = 4096;
46+ constexpr uint32_t SUPPORTED_QUANT_BITS[2] = {8, 16};
47+ constexpr uint32_t SUPPORT_QUANT_BITS_NUM = 2;
48+ constexpr uint64_t UB_SIZE_RESERVE = 133280;
49+ 
50+ struct IfmrAttrs {
51+ float minPercentile;
52+ float maxPercentile;
53+ float searchRange[2];
54+ float searchStep;
55+ bool withOffset;
56+ int quantBits;
57+ uint32_t dataLength;
58+ uint32_t cumsumLength;
59+ };
60+ 
61+ class IfmrTiling {
62+ public:
63+ explicit IfmrTiling(gert::TilingContext* context) : context_(context), nodeName_(context->GetNodeName()) {}
64+ ~IfmrTiling() {}
65+ ge::graphStatus IfmrTilingFunc(void);
66+ 
67+ protected:
68+ gert::TilingContext* context_ = nullptr;
69+ const ge::char_t* nodeName_;
70+ 
71+ private:
72+ ge::graphStatus GetIfmrTilingAttrInfo(void);
73+ ge::graphStatus CheckIfmrTilingAttrs(void);
74+ ge::graphStatus CheckIfmrTilingInputDtype(void);
75+ ge::graphStatus CheckIfmrTilingInputDataShape(uint32_t inputIndex, std::string inputName);
76+ ge::graphStatus GetIfmrTilingInputInfo(void);
77+ ge::graphStatus GetDataLength(void);
78+ ge::graphStatus CheckIfmrTilingOutputInfo(void);
79+ void SetIfmrTiling(void);
80+ void PostTiling(void);
81+ 
82+ IfmrAttrs attrs_;
83+ };
84+ 
85+} // namespace optiling
86+#endif // IFMR_TILING_H
Aquant/ifmr/op_kernel/ifmr.cpp+42-0
@@ -0,0 +1,42 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ifmr.cpp
13+ * \brief
14+ */
15+ 
16+#include "ifmr.h"
17+ 
18+using namespace AscendC;
19+
20+extern "C" __global__ __aicore__ void ifmr(
21+ GM_ADDR data,
22+ GM_ADDR data_min,
23+ GM_ADDR data_max,
24+ GM_ADDR cumsum,
25+ GM_ADDR scale,
26+ GM_ADDR offset,
27+ GM_ADDR workspace,
28+ GM_ADDR tiling) {
29+ GM_ADDR user_workspace = AscendC::GetUserWorkspace(workspace);
30+ REGISTER_TILING_DEFAULT(IfmrTilingData);
31+ GET_TILING_DATA_WITH_STRUCT(IfmrTilingData, tilingData, tiling);
32+
33+#if defined(DTYPE_DATA)
34+ if (TILING_KEY_IS(0)) {
35+ Ifmr::KernelIfmr<DTYPE_DATA> op;
36+ op.Init(data, data_min, data_max, cumsum, scale, offset, workspace, tilingData);
37+ op.Process();
38+ }
39+#endif
40+
41+ return;
42+}
Aquant/ifmr/op_kernel/ifmr.h+507-0
@@ -0,0 +1,507 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ifmr.h
13+ * \brief
14+ */
15+#ifndef IFMR_H
16+#define IFMR_H
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "ifmr_tiling_data.h"
21+ 
22+using namespace AscendC;
23+namespace Ifmr {
24+ 
25+constexpr uint64_t ALIGN_LENGTH = 32;
26+constexpr uint64_t RESERVED_LENGTH = 160;
27+constexpr float EPSILON = 1.192092896e-7f;
28+ 
29+struct IfmrAttrs {
30+ float minPercentile;
31+ float maxPercentile;
32+ float searchRange[2];
33+ float searchStep;
34+ bool withOffset;
35+ uint32_t quantBits;
36+ uint32_t dataNums;
37+ uint32_t cumsumNums;
38+ uint32_t stepNums;
39+ uint32_t stepAlignedBytes;
40+};
41+ 
42+template <typename T>
43+class KernelIfmr {
44+public:
45+ __aicore__ inline KernelIfmr() {}
46+ __aicore__ inline void ParseTilingData(const IfmrTilingData& tilingData) {
47+ attrs_.minPercentile = tilingData.minPercentile;
48+ attrs_.maxPercentile = tilingData.maxPercentile;
49+ attrs_.searchRange[0] = tilingData.searchRange[0];
50+ attrs_.searchRange[1] = tilingData.searchRange[1];
51+ attrs_.searchStep = tilingData.searchStep;
52+ attrs_.withOffset = tilingData.withOffset;
53+ attrs_.quantBits = tilingData.quantBits;
54+ attrs_.dataNums = tilingData.dataLength;
55+ attrs_.cumsumNums = tilingData.cumsumLength;
56+ attrs_.stepNums = Round((attrs_.searchRange[1] - attrs_.searchRange[0]) / attrs_.searchStep) + 1;
57+ attrs_.stepAlignedBytes = AlignB(attrs_.stepNums * sizeof(float));
58+ }
59+ 
60+ __aicore__ inline void Init(GM_ADDR data, GM_ADDR data_min, GM_ADDR data_max, GM_ADDR cumsum,
61+ GM_ADDR scale, GM_ADDR offset, GM_ADDR workSpace, const IfmrTilingData& tiling)
62+ {
63+ ParseTilingData(tiling);
64+ dataGm_.SetGlobalBuffer((__gm__ T *)data, AlignB(attrs_.dataNums * sizeof(T)) / sizeof(T));
65+ dataMinGm_.SetGlobalBuffer((__gm__ T *)data_min, ALIGN_LENGTH / sizeof(T));
66+ dataMaxGm_.SetGlobalBuffer((__gm__ T *)data_max, ALIGN_LENGTH / sizeof(T));
67+ cumsumGm_.SetGlobalBuffer((__gm__ int32_t*)cumsum,
68+ AlignB(attrs_.cumsumNums * sizeof(int32_t)) / sizeof(int32_t));
69+ scaleGm_.SetGlobalBuffer((__gm__ float *)scale, 1);
70+ offsetGm_.SetGlobalBuffer((__gm__ float *)offset, 1);
71+ 
72+ pipe_.InitBuffer(dataMaxTQue_, 1, ALIGN_LENGTH);
73+ pipe_.InitBuffer(dataMinTQue_, 1, ALIGN_LENGTH);
74+ uint64_t cumsumAlignedSize = AlignB(attrs_.cumsumNums * sizeof(int32_t));
75+ pipe_.InitBuffer(cumsumTQue_, 1, cumsumAlignedSize);
76+ pipe_.InitBuffer(maxListTQue_, 1, attrs_.stepNums * sizeof(float));
77+ pipe_.InitBuffer(minListTQue_, 1, attrs_.stepNums * sizeof(float));
78+ 
79+ pipe_.InitBuffer(scaleListQueue_, 1, attrs_.stepAlignedBytes);
80+ pipe_.InitBuffer(offsetListQueue_, 1, attrs_.stepAlignedBytes);
81+ 
82+ clipMax_ = (1 << (attrs_.quantBits - 1)) - 1;
83+ clipMin_ = -1 << (attrs_.quantBits - 1);
84+ coreNum_ = GetBlockNum();
85+ ASSERT(coreNum_ != 0 && "block dim can not be zero!");
86+ coreIdx_ = GetBlockIdx();
87+ dataEachBlock_ = ALIGN_LENGTH / sizeof(T);
88+ eachCoreDataNum_ = attrs_.dataNums / dataEachBlock_ / coreNum_ * dataEachBlock_;
89+ lastCoreDataNum_ = attrs_.dataNums - (eachCoreDataNum_ * (coreNum_ - 1));
90+ uint32_t ubFactor;
91+ if constexpr (std::is_same<T, half>::value) {
92+ ubFactor = 4; // max((orignal(1) + cast_to_f32(2)), (f32(2) + quant(2))) = 4
93+ } else {
94+ ubFactor = 2; // orignal(1) + quant(1) = 2
95+ }
96+ // 6 means maxList+minList+scale+offset+lossOut+lossIn
97+ eachLoopDataNum_ = (TOTAL_UB_SIZE - attrs_.stepAlignedBytes * 6 - cumsumAlignedSize -
98+ coreNum_ * ALIGN_LENGTH - RESERVED_LENGTH) / sizeof(T) / ubFactor / dataEachBlock_ * dataEachBlock_;
99+ eachCoreLossNum_ = attrs_.stepAlignedBytes / sizeof(float);
100+ mseLossGm_.SetGlobalBuffer((__gm__ float *)workSpace, eachCoreLossNum_ * coreNum_);
101+ pipe_.InitBuffer(dataQueue_, 1, eachLoopDataNum_ * sizeof(float));
102+ pipe_.InitBuffer(lossOutQueue_, 1, eachCoreLossNum_ * sizeof(float));
103+ pipe_.InitBuffer(lossInQueue_, 1, eachCoreLossNum_ * sizeof(float));
104+ syncGm_.SetGlobalBuffer((__gm__ int32_t*)(workSpace) + eachCoreLossNum_ * coreNum_,
105+ coreNum_ * ALIGN_LENGTH / sizeof(int32_t));
106+ pipe_.InitBuffer(syncGmQue_, 1, coreNum_ * ALIGN_LENGTH);
107+ pipe_.InitBuffer(tempBuf_, eachLoopDataNum_ * sizeof(float));
108+ }
109+ 
110+ __aicore__ inline int32_t Round(float x)
111+ {
112+ return static_cast<int32_t>(x + 0.5f);
113+ }
114+ 
115+ __aicore__ inline uint64_t AlignB(uint64_t a, uint64_t b = ALIGN_LENGTH)
116+ {
117+ if (b == 0) {
118+ return a;
119+ }
120+ return ((a + b - 1) / b) * b;
121+ }
122+ 
123+ __aicore__ inline void CopyIn()
124+ {
125+ LocalTensor<T> dataMax = dataMaxTQue_.AllocTensor<T>();
126+ LocalTensor<T> dataMin = dataMinTQue_.AllocTensor<T>();
127+ LocalTensor<int32_t> cumsum = cumsumTQue_.AllocTensor<int32_t>();
128+ 
129+ uint64_t cumsumAlignedSize = AlignB(attrs_.cumsumNums * sizeof(int32_t));
130+ DataCopy<int32_t>(cumsum, cumsumGm_, cumsumAlignedSize / sizeof(uint32_t));
131+ DataCopy<T>(dataMax, dataMaxGm_, ALIGN_LENGTH / sizeof(T));
132+ DataCopy<T>(dataMin, dataMinGm_, ALIGN_LENGTH / sizeof(T));
133+ 
134+ dataMaxTQue_.EnQue(dataMax);
135+ dataMinTQue_.EnQue(dataMin);
136+ cumsumTQue_.EnQue(cumsum);
137+ }
138+ 
139+ __aicore__ inline void Process()
140+ {
141+ CopyIn();
142+ event_t eventIdMTE2ToV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));
143+ SetFlag<HardEvent::MTE2_V>(eventIdMTE2ToV);
144+ WaitFlag<HardEvent::MTE2_V>(eventIdMTE2ToV);
145+ 
146+ LocalTensor<T> dataMax = dataMaxTQue_.DeQue<T>();
147+ LocalTensor<T> dataMin = dataMinTQue_.DeQue<T>();
148+ LocalTensor<int32_t> cumsum = cumsumTQue_.DeQue<int32_t>();
149+ LocalTensor<float> maxList = maxListTQue_.AllocTensor<float>();
150+ LocalTensor<float> minList = minListTQue_.AllocTensor<float>();
151+ GenerateMaxMinCandidates(maxList, minList, cumsum, dataMax, dataMin);
152+ 
153+ dataMaxTQue_.FreeTensor(dataMax);
154+ dataMinTQue_.FreeTensor(dataMin);
155+ cumsumTQue_.FreeTensor(cumsum);
156+ 
157+ LocalTensor<float> scaleList = scaleListQueue_.AllocTensor<float>();
158+ LocalTensor<float> offsetList = offsetListQueue_.AllocTensor<float>();
159+ CalScaleOffset(scaleList, offsetList, maxList, minList);
160+ maxListTQue_.FreeTensor(maxList);
161+ minListTQue_.FreeTensor(minList);
162+ 
163+ ComputeMSELoss(scaleList, offsetList);
164+ if (coreNum_ > 1) {
165+ SyncAll();
166+ }
167+ 
168+ if (coreIdx_ == 0) {
169+ ReduceAndOutput(scaleList, offsetList);
170+ }
171+ 
172+ scaleListQueue_.FreeTensor(scaleList);
173+ offsetListQueue_.FreeTensor(offsetList);
174+ }
175+ 
176+private:
177+ TPipe pipe_;
178+ GlobalTensor<T> dataGm_;
179+ GlobalTensor<T> dataMinGm_;
180+ GlobalTensor<T> dataMaxGm_;
181+ GlobalTensor<int32_t> cumsumGm_;
182+ 
183+ GlobalTensor<float> scaleGm_;
184+ GlobalTensor<float> offsetGm_;
185+ 
186+ TQue<TPosition::VECIN, 1> dataMaxTQue_;
187+ TQue<TPosition::VECIN, 1> dataMinTQue_;
188+ TQue<TPosition::VECIN, 1> cumsumTQue_;
189+ TQue<TPosition::VECOUT, 1> maxListTQue_;
190+ TQue<TPosition::VECOUT, 1> minListTQue_;
191+ TQue<TPosition::VECOUT, 1> scaleListQueue_;
192+ TQue<TPosition::VECOUT, 1> offsetListQueue_;
193+ 
194+ GlobalTensor<int32_t> syncGm_;
195+ TQue<TPosition::VECOUT, 1> syncGmQue_;
196+ GlobalTensor<float> mseLossGm_;
197+ TQue<TPosition::VECIN, 1> dataQueue_;
198+ TQue<TPosition::VECOUT, 1> lossOutQueue_;
199+ TQue<TPosition::VECIN, 1> lossInQueue_;
200+ TBuf<TPosition::VECCALC> tempBuf_;
201+ 
202+ int64_t coreNum_{0};
203+ int64_t coreIdx_{0};
204+ uint32_t dataEachBlock_{0};
205+ uint32_t eachCoreDataNum_{0};
206+ uint32_t lastCoreDataNum_{0};
207+ uint32_t eachLoopDataNum_{0};
208+ uint32_t eachCoreLossNum_{0};
209+ float clipMax_{0.0f};
210+ float clipMin_{0.0f};
211+ 
212+ IfmrAttrs attrs_;
213+ 
214+ __aicore__ inline void CalScaleOffset(LocalTensor<float>& scaleList, LocalTensor<float>& offsetList,
215+ const LocalTensor<float>& maxList, const LocalTensor<float>& minList);
216+ 
217+ __aicore__ inline void GetMinMaxIndex(LocalTensor<float>& minIndex, LocalTensor<float>& maxIndex,
218+ const LocalTensor<int32_t>& cumsumLocal);
219+ 
220+ __aicore__ inline void CalcMinInitMaxInit(LocalTensor<float>& minIndex, LocalTensor<float>& maxIndex,
221+ const LocalTensor<T>& dataMin, const LocalTensor<T>& dataMax);
222+ 
223+ __aicore__ inline void GenerateMaxMinCandidates(LocalTensor<float>& maxList, LocalTensor<float>& minList,
224+ const LocalTensor<int32_t>& cumsumLocal, const LocalTensor<T>& dataMax, const LocalTensor<T>& dataMin);
225+ 
226+ __aicore__ inline void GetInputData(uint32_t loopIdx, uint32_t calNum);
227+ 
228+ __aicore__ inline void ComputeEachMSELoss(const LocalTensor<float>& scaleList,
229+ const LocalTensor<float>& offsetList, LocalTensor<float>& eachCoreLoss, uint32_t calNum);
230+ 
231+ __aicore__ inline void ComputeMSELoss(const LocalTensor<float>& scaleList, const LocalTensor<float>& offsetList);
232+ 
233+ __aicore__ inline void ReduceAndOutput(const LocalTensor<float>& scaleList, const LocalTensor<float>& offsetList);
234+};
235+ 
236+template <typename T>
237+__aicore__ inline void KernelIfmr<T>::CalScaleOffset(LocalTensor<float>& scaleList, LocalTensor<float>& offsetList,
238+ const LocalTensor<float>& maxList, const LocalTensor<float>& minList)
239+{
240+ LocalTensor<float> tmp = tempBuf_.Get<float>(attrs_.stepNums);
241+ if (attrs_.withOffset) {
242+ Sub(tmp, maxList, minList, attrs_.stepNums);
243+ float divisor = 1 / static_cast<float>((1ULL << attrs_.quantBits) - 1);
244+ 
245+ PipeBarrier<PIPE_ALL>();
246+ Muls(scaleList, tmp, divisor, attrs_.stepNums);
247+ PipeBarrier<PIPE_ALL>();
L
Llouyuxuan3月3日

尽量避免使用 PIPE_ALL，对性能影响大，应该可以明确需要同步哪类流水

likedislike
F
fujun19
3月3日 评论：
248+ 
249+ Div(tmp, minList, scaleList, attrs_.stepNums);
250+ PipeBarrier<PIPE_ALL>();
251+ Cast(tmp, tmp, RoundMode::CAST_RINT, attrs_.stepNums);
252+ PipeBarrier<PIPE_ALL>();
253+ Muls(tmp, tmp, -1.0f, attrs_.stepNums);
254+ PipeBarrier<PIPE_ALL>();
255+ Adds(offsetList, tmp, clipMin_, attrs_.stepNums);
256+ } else {
257+ float divisor = 1 / clipMax_;
258+ Muls(scaleList, maxList, divisor, attrs_.stepNums);
259+ PipeBarrier<PIPE_ALL>();
260+ Duplicate(offsetList, 0.0f, attrs_.stepNums);
261+ }
262+ PipeBarrier<PIPE_ALL>();
263+ return;
264+}
265+ 
266+template <typename T>
267+__aicore__ inline void KernelIfmr<T>::GetMinMaxIndex(LocalTensor<float>& minIndex, LocalTensor<float>& maxIndex,
268+ const LocalTensor<int32_t>& cumsumLocal)
269+{
270+ LocalTensor<float> cdf = tempBuf_.Get<float>(attrs_.cumsumNums);
271+ Cast(cdf, cumsumLocal, RoundMode::CAST_NONE, attrs_.cumsumNums);
272+ PipeBarrier<PIPE_ALL>();
273+ 
274+ Muls(cdf, cdf , 1 / static_cast<float>(attrs_.dataNums), attrs_.cumsumNums);
275+ PipeBarrier<PIPE_ALL>();
276+ 
277+ int32_t compareResultSize = AlignB(attrs_.cumsumNums, AscendCUtils::GetBitSize(sizeof(uint8_t))) /
278+ AscendCUtils::GetBitSize(sizeof(uint8_t));
279+ LocalTensor<uint8_t> compareResult = tempBuf_.GetWithOffset<uint8_t>(compareResultSize,
280+ AlignB(attrs_.cumsumNums * sizeof(float)));
281+ CompareScalar(compareResult, cdf, attrs_.maxPercentile, CMPMODE::LT, attrs_.cumsumNums);
282+ 
283+ LocalTensor<float> selectResult = tempBuf_.GetWithOffset<float>(attrs_.cumsumNums,
284+ AlignB(attrs_.cumsumNums * sizeof(float)) + AlignB(compareResultSize * sizeof(uint8_t)));
285+ Duplicate(selectResult, 1.0f, attrs_.cumsumNums);
286+ PipeBarrier<PIPE_ALL>();
287+ 
288+ Select(selectResult, compareResult, selectResult, 0.0f,
289+ SELMODE::VSEL_TENSOR_SCALAR_MODE, attrs_.cumsumNums);
290+ PipeBarrier<PIPE_ALL>();
291+ 
292+ // get max index by reduce sum max percentile comp result
293+ TBuf<TPosition::VECCALC> workSpaceTbuf;
294+ pipe_.InitBuffer(workSpaceTbuf, 1 * sizeof(float));
295+ LocalTensor<float> workSpaceTensor = workSpaceTbuf.Get<float>();
296+ 
297+ ReduceSum(maxIndex, selectResult, workSpaceTensor, attrs_.cumsumNums);
298+ PipeBarrier<PIPE_ALL>();
299+ 
300+ // get min index
301+ CompareScalar(compareResult, cdf, 1 - attrs_.minPercentile, CMPMODE::LT, attrs_.cumsumNums);
302+ Duplicate(selectResult, 1.0f, attrs_.cumsumNums);
303+ PipeBarrier<PIPE_ALL>();
304+ 
305+ Select(selectResult, compareResult, selectResult, 0.0f,
306+ SELMODE::VSEL_TENSOR_SCALAR_MODE, attrs_.cumsumNums);
307+ PipeBarrier<PIPE_ALL>();
308+ ReduceSum(minIndex, selectResult, workSpaceTensor, attrs_.cumsumNums);
309+ PipeBarrier<PIPE_ALL>();
310+}
311+ 
312+template <typename T>
313+__aicore__ inline void KernelIfmr<T>::CalcMinInitMaxInit(LocalTensor<float>& minIndex, LocalTensor<float>& maxIndex,
314+ const LocalTensor<T>& dataMin, const LocalTensor<T>& dataMax)
315+{
316+ TBuf<TPosition::VECCALC> dataMaxFloat32Buf;
317+ TBuf<TPosition::VECCALC> dataMinFloat32Buf;
318+ 
319+ pipe_.InitBuffer(dataMaxFloat32Buf, sizeof(float));
320+ pipe_.InitBuffer(dataMinFloat32Buf, sizeof(float));
321+ 
322+ // calc max init
323+ LocalTensor<float> dataMaxFloat32 = dataMaxFloat32Buf.Get<float>();
324+ LocalTensor<float> dataMinFloat32 = dataMinFloat32Buf.Get<float>();
325+ if constexpr (std::is_same<T, half>::value) {
326+ Cast(dataMaxFloat32, dataMax, RoundMode::CAST_NONE, 1);
327+ Cast(dataMinFloat32, dataMin, RoundMode::CAST_NONE, 1);
328+ } else {
329+ Adds(dataMaxFloat32, dataMax, 0.0f, 1);
330+ Adds(dataMinFloat32, dataMin, 0.0f, 1);
331+ }
332+ PipeBarrier<PIPE_ALL>();
333+ 
334+ Muls(maxIndex, maxIndex, 1 / static_cast<float>(attrs_.cumsumNums), 1);
335+ Sub(dataMaxFloat32, dataMaxFloat32, dataMinFloat32, 1);
336+ PipeBarrier<PIPE_ALL>();
337+ Mul(maxIndex, maxIndex, dataMaxFloat32, 1);
338+ Add(maxIndex, maxIndex, dataMinFloat32, 1);
339+ 
340+ // calc min init
341+ Muls(minIndex, minIndex, 1 / static_cast<float>(attrs_.cumsumNums), 1);
342+ Mul(minIndex, minIndex, dataMaxFloat32, 1);
343+ Add(minIndex, minIndex, dataMinFloat32, 1);
344+ PipeBarrier<PIPE_ALL>();
345+}
346+ 
347+template <typename T>
348+__aicore__ inline void KernelIfmr<T>::GenerateMaxMinCandidates(LocalTensor<float>& maxList,
349+ LocalTensor<float>& minList, const LocalTensor<int32_t>& cumsumLocal,
350+ const LocalTensor<T>& dataMax, const LocalTensor<T>& dataMin)
351+{
352+ // apply local tensor
353+ TQue<TPosition::VECOUT, 1> maxIndexQueue;
354+ TQue<TPosition::VECOUT, 1> minIndexQueue;
355+ 
356+ // init buffer for vars in calc
357+ pipe_.InitBuffer(minIndexQueue, 1, sizeof(float));
358+ pipe_.InitBuffer(maxIndexQueue, 1, sizeof(float));
359+ 
360+ LocalTensor<float> minIndex = minIndexQueue.AllocTensor<float>();
361+ LocalTensor<float> maxIndex = maxIndexQueue.AllocTensor<float>();
362+ GetMinMaxIndex(minIndex, maxIndex, cumsumLocal);
363+ maxIndexQueue.EnQue(maxIndex);
364+ minIndexQueue.EnQue(minIndex);
365+ 
366+ maxIndex = maxIndexQueue.DeQue<float>();
367+ minIndex = minIndexQueue.DeQue<float>();
368+ CalcMinInitMaxInit(minIndex, maxIndex, dataMin, dataMax);
369+ 
370+ if (attrs_.withOffset) {
371+ Maxs(maxIndex, maxIndex, 0.0f, 1);
372+ Mins(minIndex, minIndex, 0.0f, 1);
373+ PipeBarrier<PIPE_ALL>();
374+ 
375+ Duplicate(minList, minIndex.GetValue(0), attrs_.stepNums);
376+ PipeBarrier<PIPE_ALL>();
377+ } else {
378+ Abs(minIndex, minIndex, 1);
379+ PipeBarrier<PIPE_ALL>();
380+ Max(maxIndex, maxIndex, minIndex, 1);
381+ PipeBarrier<PIPE_ALL>();
382+ }
383+ 
384+ float maxInit = maxIndex.GetValue(0);
385+ PipeBarrier<PIPE_ALL>();
386+ minIndexQueue.FreeTensor(minIndex);
387+ maxIndexQueue.FreeTensor(maxIndex);
388+ for (int32_t i = 0; i < attrs_.stepNums; i++) {
389+ maxList.SetValue(i, (attrs_.searchRange[0] + attrs_.searchStep * i) * maxInit);
390+ }
391+ PipeBarrier<PIPE_ALL>();
392+ return;
393+}
394+ 
395+template <typename T>
396+__aicore__ inline void KernelIfmr<T>::GetInputData(uint32_t loopIdx, uint32_t calNum)
397+{
398+ LocalTensor<float> data = dataQueue_.AllocTensor<float>();
399+ if constexpr (std::is_same<T, half>::value) {
400+ LocalTensor<half> tempData = tempBuf_.Get<half>(calNum);
401+ DataCopy(tempData, dataGm_[eachCoreDataNum_ * coreIdx_ + eachLoopDataNum_ * loopIdx], calNum);
402+ event_t eventID = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));
403+ SetFlag<HardEvent::MTE2_V>(eventID);
404+ WaitFlag<HardEvent::MTE2_V>(eventID);
405+ Cast(data, tempData, RoundMode::CAST_NONE, calNum);
406+ } else {
407+ DataCopy(data, dataGm_[eachCoreDataNum_ * coreIdx_ + eachLoopDataNum_ * loopIdx], calNum);
408+ }
409+ dataQueue_.EnQue(data);
410+}
411+ 
412+template <typename T>
413+__aicore__ inline void KernelIfmr<T>::ComputeEachMSELoss(const LocalTensor<float>& scaleList,
414+ const LocalTensor<float>& offsetList, LocalTensor<float>& eachCoreLoss, uint32_t calNum)
415+{
416+ LocalTensor<float> data = dataQueue_.DeQue<float>();
417+ LocalTensor<float> tempData = tempBuf_.Get<float>(calNum);
418+ for (uint32_t i = 0; i < attrs_.stepNums; i++) {
419+ // quant
420+ Muls(tempData, data, 1 / scaleList.GetValue(i), calNum);
421+ Adds(tempData, tempData, offsetList.GetValue(i), calNum);
422+ // round
423+ Cast(tempData, tempData, RoundMode::CAST_RINT, calNum);
424+ // clip
425+ Mins(tempData, tempData, clipMax_, calNum);
426+ Maxs(tempData, tempData, clipMin_, calNum);
427+ // dequant
428+ Adds(tempData, tempData, -offsetList.GetValue(i), calNum);
429+ Muls(tempData, tempData, scaleList.GetValue(i), calNum);
430+ // calculate mse loss
431+ Sub(tempData, tempData, data, calNum);
432+ Mul(tempData, tempData, tempData, calNum);
433+ ReduceSum(tempData, tempData, tempData, calNum);
434+ event_t eventID = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_S));
435+ SetFlag<HardEvent::V_S>(eventID);
436+ WaitFlag<HardEvent::V_S>(eventID);
437+ eachCoreLoss.SetValue(i, eachCoreLoss.GetValue(i) + tempData.GetValue(0));
438+ }
439+ dataQueue_.FreeTensor(data);
440+ PipeBarrier<PIPE_ALL>();
441+}
442+ 
443+template <typename T>
444+__aicore__ inline void KernelIfmr<T>::ComputeMSELoss(const LocalTensor<float>& scaleList,
445+ const LocalTensor<float>& offsetList)
446+{
447+ LocalTensor<float> eachCoreLoss = lossOutQueue_.AllocTensor<float>();
448+ Duplicate(eachCoreLoss, 0.0f, eachCoreLossNum_);
449+ uint32_t inputNum = eachCoreDataNum_;
450+ if (coreIdx_ == (coreNum_ - 1)) {
451+ inputNum = lastCoreDataNum_;
452+ }
453+ uint32_t loopNum = inputNum / eachLoopDataNum_;
454+ for (uint32_t i = 0; i < loopNum; i++) {
455+ GetInputData(i, eachLoopDataNum_);
456+ ComputeEachMSELoss(scaleList, offsetList, eachCoreLoss, eachLoopDataNum_);
457+ }
458+ uint32_t leftDataNum = inputNum % eachLoopDataNum_;
459+ if (leftDataNum > 0) {
460+ GetInputData(loopNum, AlignB(leftDataNum, dataEachBlock_));
461+ ComputeEachMSELoss(scaleList, offsetList, eachCoreLoss, leftDataNum);
462+ }
463+ DataCopy(mseLossGm_[eachCoreLossNum_ * coreIdx_], eachCoreLoss, eachCoreLossNum_);
464+ lossOutQueue_.EnQue(eachCoreLoss);
465+ lossOutQueue_.DeQue<float>();
466+ lossOutQueue_.FreeTensor(eachCoreLoss);
467+ PipeBarrier<PIPE_ALL>();
468+}
469+ 
470+template <typename T>
471+__aicore__ inline void KernelIfmr<T>::ReduceAndOutput(const LocalTensor<float>& scaleList,
472+ const LocalTensor<float>& offsetList)
473+{
474+ LocalTensor<float> localLoss = lossInQueue_.AllocTensor<float>();
475+ LocalTensor<float> tempLoss = tempBuf_.Get<float>(eachCoreLossNum_);
476+ Duplicate(tempLoss, 0.0f, eachCoreLossNum_);
477+ for (uint32_t i = 0; i < coreNum_; i++) {
478+ event_t eventID = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE2));
479+ SetFlag<HardEvent::V_MTE2>(eventID);
480+ WaitFlag<HardEvent::V_MTE2>(eventID);
481+ DataCopy(localLoss, mseLossGm_[eachCoreLossNum_ * i], eachCoreLossNum_);
482+ eventID = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));
483+ SetFlag<HardEvent::MTE2_V>(eventID);
484+ WaitFlag<HardEvent::MTE2_V>(eventID);
485+ Add(tempLoss, tempLoss, localLoss, attrs_.stepNums);
486+ }
487+ ReduceMin(tempLoss, tempLoss, tempLoss, attrs_.stepNums, true);
488+ event_t eventID = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_S));
489+ SetFlag<HardEvent::V_S>(eventID);
490+ WaitFlag<HardEvent::V_S>(eventID);
491+ float tempMinLossIndex = tempLoss.GetValue(1);
492+ uint32_t minIndex = *reinterpret_cast<uint32_t*>(&tempMinLossIndex);
493+ float scale = scaleList.GetValue(minIndex);
494+ float offset = offsetList.GetValue(minIndex);
495+ if (scale < EPSILON) {
496+ scale = 1.0f;
497+ offset = attrs_.withOffset ? clipMin_ : 0.0f;
498+ }
499+ scaleGm_.SetValue(0, scale);
500+ offsetGm_.SetValue(0, offset);
501+ DataCacheCleanAndInvalid<float, AscendC::CacheLine::SINGLE_CACHE_LINE, AscendC::DcciDst::CACHELINE_OUT>(scaleGm_);
502+ DataCacheCleanAndInvalid<float, AscendC::CacheLine::SINGLE_CACHE_LINE, AscendC::DcciDst::CACHELINE_OUT>(offsetGm_);
503+ lossInQueue_.FreeTensor(localLoss);
504+ PipeBarrier<PIPE_ALL>();
505+}
506+}
507+#endif // IFMR_H
Aquant/ifmr/op_kernel/ifmr_tiling_data.h+32-0
@@ -0,0 +1,32 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file ifmr_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef IFMR_TILLING_DATA_H
17+#define IFMR_TILLING_DATA_H
18+ 
19+#include <cstdint>
20+ 
21+struct IfmrTilingData {
22+ float minPercentile;
23+ float maxPercentile;
24+ float searchRange[2];
25+ float searchStep;
26+ bool withOffset;
27+ uint32_t quantBits;
28+ uint32_t dataLength;
29+ uint32_t cumsumLength;
30+};
31+ 
32+#endif
Aquant/ifmr/tests/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start ops.quant.ifmr.tests.CMakeLists.txt ")
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
14+foreach(SUB_DIR ${CURRENT_DIRS})
15+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16+ add_subdirectory(${SUB_DIR})
17+ endif()
18+endforeach()
Aquant/ifmr/tests/op_host/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
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()
Aquant/ifmr/tests/op_host/test_ifmr_tiling.cpp+428-0
@@ -0,0 +1,428 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+
13+#include <iostream>
14+#include <map>
15+#include <sstream>
16+#include <string>
17+#include <vector>
18+#include <stdio.h>
19+ 
20+#include "log/log.h"
21+#include "exe_graph/runtime/storage_format.h"
22+#include "exe_graph/runtime/storage_shape.h"
23+#include "kernel_run_context_facker.h"
24+#include "test_cube_util.h"
25+#include "platform/platform_infos_def.h"
26+#include "ut_op_util.h"
27+#include "quant/ifmr/op_kernel/ifmr_tiling_data.h"
28+ 
29+struct IFMRCompileInfo {
30+ int32_t dataNum;
31+};
32+ 
33+class IFMRTilingTest : public testing::Test {
34+protected:
35+ static void SetUpTestCase() {
36+ std::cout << "IFMRTilingTest SetUp" << std::endl;
37+ }
38+
39+ static void TearDownTestCase() {
40+ std::cout << "IFMRTilingTest TearDown" << std::endl;
41+ }
42+};
43+
44+struct IFMRAttrs {
45+ float minPercentile;
46+ float maxPercentile;
47+ std::vector<float> searchRange;
48+ float searchStep;
49+ bool withOffset;
50+ uint32_t quantBits;
51+};
52+
53+ge::graphStatus IFMRTestCase(vector<vector<int64_t>> input_shapes, vector<vector<int64_t>> output_shapes,
54+ IFMRAttrs attrs, ge::DataType dataType, IfmrTilingData &tilingParam) {
55+ std::string opType("IFMR");
56+ if (gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()) == nullptr) {
57+ std::cout << "gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()) is nullptr" << std::endl;
58+ return ge::GRAPH_FAILED;
59+ }
60+
61+ gert::StorageShape data = {{input_shapes[0][0]}, {input_shapes[0][0]}};
62+ gert::StorageShape dataMin = {{input_shapes[1][0]}, {input_shapes[1][0]}};
63+ gert::StorageShape dataMax = {{input_shapes[2][0]}, {input_shapes[2][0]}};
64+ gert::StorageShape cumsum = {{input_shapes[3][0]}, {input_shapes[3][0]}};
65+ gert::StorageShape scale = {{output_shapes[0][0]}, {output_shapes[0][0]}};
66+ gert::StorageShape offset = {{output_shapes[0][0]}, {output_shapes[0][0]}};
67+
68+ if (gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()) == nullptr) {
69+ std::cout << "gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()) is nullptr" << std::endl;
70+ return ge::GRAPH_FAILED;
71+ }
72+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
73+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
74+ 
75+ uint64_t L1_SIZE = 524288;
76+ uint64_t L0a_SIZE = 65536;
77+ uint64_t L0b_SIZE = 65536;
78+ uint64_t L0c_SIZE = 262144;
79+ uint64_t aicoreNum = 32;
80+ string compile_info_string = R"({
81+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
82+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
83+ "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
84+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
85+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 262144,
86+ "CORE_NUM": 32}
87+ })";
88+ map<string, string> socInfos;
89+ map<string, string> aicoreSpec;
90+ map<string, string> intrinsics;
91+ GetPlatFormInfos(compile_info_string.c_str(), socInfos, aicoreSpec, intrinsics);
92+ map<string, string> socVersionInfos = {{"Short_SoC_version", "Ascend950"}};
93+ fe::PlatFormInfos platformInfo;
94+ platformInfo.Init();
95+ auto tilingDataPtr = gert::TilingData::CreateCap(4096);
96+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
97+ auto wsSize = reinterpret_cast<gert::ContinuousVector *>(workspaceSizeHoler.get());
98+ if (tilingDataPtr == nullptr) {
99+ std::cout << "tilingDataPtr is nullptr" << std::endl;
100+ return ge::GRAPH_FAILED;
101+ }
102+
103+ std::vector<void*> inputShapeRef = {&data, &dataMin, &dataMax, &cumsum};
104+ std::vector<void*> outputShapesRef = {&scale, &offset};
105+
106+ // compile info
107+ IFMRCompileInfo compileInfo;
108+ // tilingParseFunc simulate
109+ auto kernelHolder = gert::KernelRunContextFaker()
110+ .KernelIONum(2, 1)
111+ .Inputs({const_cast<char *>(compile_info_string.c_str()), reinterpret_cast<void *>(&platformInfo)})
112+ .Outputs({&compileInfo})
113+ .Build();
114+ if (kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init() == false) {
115+ std::cout << "kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init() is nullptr" << std::endl;
116+ return ge::GRAPH_FAILED;
117+ }
118+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
119+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
120+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
121+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
122+ 
123+ auto holder = gert::TilingContextFaker().SetOpType(opType)
124+ .NodeIoNum(4, 2)
125+ .IrInstanceNum({1, 1, 1, 1})
126+ .InputShapes(inputShapeRef)
127+ .OutputShapes(outputShapesRef)
128+ .CompileInfo(&compileInfo)
129+ .PlatformInfo(reinterpret_cast<char *>(&platformInfo))
130+ .NodeInputTd(0, dataType, ge::FORMAT_ND, ge::FORMAT_ND)
131+ .NodeInputTd(1, dataType, ge::FORMAT_ND, ge::FORMAT_ND)
132+ .NodeInputTd(2, dataType, ge::FORMAT_ND, ge::FORMAT_ND)
133+ .NodeInputTd(3, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
134+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
135+ .NodeOutputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
136+ .NodeAttrs({
137+ {"min_percentile", Ops::NN::AnyValue::CreateFrom<float>(attrs.minPercentile)},
138+ {"max_percentile", Ops::NN::AnyValue::CreateFrom<float>(attrs.maxPercentile)},
139+ {"search_range", Ops::NN::AnyValue::CreateFrom<std::vector<float>>(attrs.searchRange)},
140+ {"search_step", Ops::NN::AnyValue::CreateFrom<float>(attrs.searchStep)},
141+ {"with_offset", Ops::NN::AnyValue::CreateFrom<bool>(attrs.withOffset)},
142+ {"quant_bits", Ops::NN::AnyValue::CreateFrom<int64_t>(attrs.quantBits)}
143+ })
144+ .TilingData(tilingDataPtr.get())
145+ .Workspace(wsSize)
146+ .Build();
147+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
148+ if (tilingContext->GetPlatformInfo() == nullptr) {
149+ std::cout << "tilingContext->GetPlatformInfo() is nullptr" << std::endl;
150+ return ge::GRAPH_FAILED;
151+ }
152+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("version", socVersionInfos);
153+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
154+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
155+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
156+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
157+ auto ret = tilingFunc(tilingContext);
158+ auto buf = (IfmrTilingData*)tilingContext->GetTilingData<IfmrTilingData>();
159+ tilingParam = *buf;
160+ return ret;
161+}
162+
163+TEST_F(IFMRTilingTest, run_IFMR_case_0) {
164+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
165+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
166+ IFMRAttrs attrs;
167+ attrs.minPercentile = 0.9;
168+ attrs.maxPercentile = 0.9;
169+ attrs.searchRange = {0.7, 1.3};
170+ attrs.searchStep = 0.1;
171+ attrs.withOffset = true;
172+ attrs.quantBits = 8;
173+ ge::DataType dataType = ge::DT_FLOAT;
174+
175+ IfmrTilingData tilingdata;
176+
177+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
178+ EXPECT_EQ(ret, ge::GRAPH_SUCCESS);
179+
180+ EXPECT_EQ(attrs.minPercentile, tilingdata.minPercentile);
181+ EXPECT_EQ(attrs.maxPercentile, tilingdata.maxPercentile);
182+ EXPECT_EQ(attrs.searchRange[0], tilingdata.searchRange[0]);
183+ EXPECT_EQ(attrs.searchRange[1], tilingdata.searchRange[1]);
184+ EXPECT_EQ(attrs.searchStep, tilingdata.searchStep);
185+ EXPECT_EQ((uint32_t)attrs.withOffset, tilingdata.withOffset);
186+ EXPECT_EQ(attrs.quantBits, tilingdata.quantBits);
187+}
188+
189+TEST_F(IFMRTilingTest, run_IFMR_case_1) {
190+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
191+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
192+ IFMRAttrs attrs;
193+ attrs.minPercentile = 0.9;
194+ attrs.maxPercentile = 0.9;
195+ attrs.searchRange = {0.7, 1.3};
196+ attrs.searchStep = 0.1;
197+ attrs.withOffset = true;
198+ attrs.quantBits = 8;
199+ ge::DataType dataType = ge::DT_FLOAT16;
200+
201+ IfmrTilingData tilingdata;
202+
203+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
204+ EXPECT_EQ(ret, ge::GRAPH_SUCCESS);
205+
206+ EXPECT_EQ(attrs.minPercentile, tilingdata.minPercentile);
207+ EXPECT_EQ(attrs.maxPercentile, tilingdata.maxPercentile);
208+ EXPECT_EQ(attrs.searchRange[0], tilingdata.searchRange[0]);
209+ EXPECT_EQ(attrs.searchRange[1], tilingdata.searchRange[1]);
210+ EXPECT_EQ(attrs.searchStep, tilingdata.searchStep);
211+ EXPECT_EQ((uint32_t)attrs.withOffset, tilingdata.withOffset);
212+ EXPECT_EQ(attrs.quantBits, tilingdata.quantBits);
213+}
214+
215+TEST_F(IFMRTilingTest, run_IFMR_case_2) {
216+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
217+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
218+ IFMRAttrs attrs;
219+ attrs.minPercentile = 0.9;
220+ attrs.maxPercentile = 0.9;
221+ attrs.searchRange = {0.7, 1.3};
222+ attrs.searchStep = 0.1;
223+ attrs.withOffset = false;
224+ attrs.quantBits = 8;
225+ ge::DataType dataType = ge::DT_FLOAT16;
226+
227+ IfmrTilingData tilingdata;
228+
229+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
230+ EXPECT_EQ(ret, ge::GRAPH_SUCCESS);
231+
232+ EXPECT_EQ(attrs.minPercentile, tilingdata.minPercentile);
233+ EXPECT_EQ(attrs.maxPercentile, tilingdata.maxPercentile);
234+ EXPECT_EQ(attrs.searchRange[0], tilingdata.searchRange[0]);
235+ EXPECT_EQ(attrs.searchRange[1], tilingdata.searchRange[1]);
236+ EXPECT_EQ(attrs.searchStep, tilingdata.searchStep);
237+ EXPECT_EQ((uint32_t)attrs.withOffset, tilingdata.withOffset);
238+ EXPECT_EQ(attrs.quantBits, tilingdata.quantBits);
239+}
240+
241+TEST_F(IFMRTilingTest, run_IFMR_case_3) {
242+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
243+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
244+ IFMRAttrs attrs;
245+ attrs.minPercentile = 0.9;
246+ attrs.maxPercentile = 0.9;
247+ attrs.searchRange = {0.7, 1.3};
248+ attrs.searchStep = 0.1;
249+ attrs.withOffset = false;
250+ attrs.quantBits = 8;
251+ ge::DataType dataType = ge::DT_FLOAT;
252+
253+ IfmrTilingData tilingdata;
254+
255+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
256+ EXPECT_EQ(ret, ge::GRAPH_SUCCESS);
257+
258+ EXPECT_EQ(attrs.minPercentile, tilingdata.minPercentile);
259+ EXPECT_EQ(attrs.maxPercentile, tilingdata.maxPercentile);
260+ EXPECT_EQ(attrs.searchRange[0], tilingdata.searchRange[0]);
261+ EXPECT_EQ(attrs.searchRange[1], tilingdata.searchRange[1]);
262+ EXPECT_EQ(attrs.searchStep, tilingdata.searchStep);
263+ EXPECT_EQ((uint32_t)attrs.withOffset, tilingdata.withOffset);
264+ EXPECT_EQ(attrs.quantBits, tilingdata.quantBits);
265+}
266+
267+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_minpercentile_01) {
268+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
269+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
270+ IFMRAttrs attrs;
271+ attrs.minPercentile = 0.1;
272+ attrs.maxPercentile = 0.9;
273+ attrs.searchRange = {0.7, 1.3};
274+ attrs.searchStep = 0.1;
275+ attrs.withOffset = false;
276+ attrs.quantBits = 8;
277+ ge::DataType dataType = ge::DT_FLOAT;
278+
279+ IfmrTilingData tilingdata;
280+
281+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
282+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
283+}
284+
285+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_minpercentile_02) {
286+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
287+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
288+ IFMRAttrs attrs;
289+ attrs.minPercentile = 1.1;
290+ attrs.maxPercentile = 0.9;
291+ attrs.searchRange = {0.7, 1.3};
292+ attrs.searchStep = 0.1;
293+ attrs.withOffset = false;
294+ attrs.quantBits = 8;
295+ ge::DataType dataType = ge::DT_FLOAT;
296+
297+ IfmrTilingData tilingdata;
298+
299+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
300+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
301+}
302+
303+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_maxpercentile_01) {
304+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
305+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
306+ IFMRAttrs attrs;
307+ attrs.minPercentile = 0.9;
308+ attrs.maxPercentile = 0.1;
309+ attrs.searchRange = {0.7, 1.3};
310+ attrs.searchStep = 0.1;
311+ attrs.withOffset = false;
312+ attrs.quantBits = 8;
313+ ge::DataType dataType = ge::DT_FLOAT;
314+
315+ IfmrTilingData tilingdata;
316+
317+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
318+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
319+}
320+
321+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_maxpercentile_02) {
322+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
323+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
324+ IFMRAttrs attrs;
325+ attrs.minPercentile = 0.9;
326+ attrs.maxPercentile = 1.1;
327+ attrs.searchRange = {0.7, 1.3};
328+ attrs.searchStep = 0.1;
329+ attrs.withOffset = false;
330+ attrs.quantBits = 8;
331+ ge::DataType dataType = ge::DT_FLOAT;
332+
333+ IfmrTilingData tilingdata;
334+
335+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
336+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
337+}
338+
339+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_searchRange_01) {
340+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
341+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
342+ IFMRAttrs attrs;
343+ attrs.minPercentile = 0.9;
344+ attrs.maxPercentile = 0.9;
345+ attrs.searchRange = {0, 1};
346+ attrs.searchStep = 0.1;
347+ attrs.withOffset = false;
348+ attrs.quantBits = 8;
349+ ge::DataType dataType = ge::DT_FLOAT;
350+
351+ IfmrTilingData tilingdata;
352+
353+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
354+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
355+}
356+
357+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_searchRange_02) {
358+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
359+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
360+ IFMRAttrs attrs;
361+ attrs.minPercentile = 0.9;
362+ attrs.maxPercentile = 0.9;
363+ attrs.searchRange = {1.1, 1.1};
364+ attrs.searchStep = 0.1;
365+ attrs.withOffset = false;
366+ attrs.quantBits = 8;
367+ ge::DataType dataType = ge::DT_FLOAT;
368+
369+ IfmrTilingData tilingdata;
370+
371+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
372+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
373+}
374+
375+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_searchStep_01) {
376+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
377+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
378+ IFMRAttrs attrs;
379+ attrs.minPercentile = 0.9;
380+ attrs.maxPercentile = 0.9;
381+ attrs.searchRange = {1, 8193};
382+ attrs.searchStep = 1;
383+ attrs.withOffset = false;
384+ attrs.quantBits = 8;
385+ ge::DataType dataType = ge::DT_FLOAT;
386+
387+ IfmrTilingData tilingdata;
388+
389+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
390+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
391+}
392+
393+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_quantBits_01) {
394+ vector<vector<int64_t>> input_shapes = {{1024}, {1}, {1}, {1024}};
395+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
396+ IFMRAttrs attrs;
397+ attrs.minPercentile = 0.9;
398+ attrs.maxPercentile = 0.9;
399+ attrs.searchRange = {0.7, 1.3};
400+ attrs.searchStep = 0.1;
401+ attrs.withOffset = false;
402+ attrs.quantBits = 9;
403+ ge::DataType dataType = ge::DT_FLOAT;
404+
405+ IfmrTilingData tilingdata;
406+
407+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
408+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
409+}
410+ 
411+
412+TEST_F(IFMRTilingTest, run_IFMR_case_invalid_data_length) {
413+ vector<vector<int64_t>> input_shapes = {{0}, {0}, {0}, {0}};
414+ vector<vector<int64_t>> output_shapes = {{1}, {1}};
415+ IFMRAttrs attrs;
416+ attrs.minPercentile = 0.9;
417+ attrs.maxPercentile = 0.9;
418+ attrs.searchRange = {0.7, 1.3};
419+ attrs.searchStep = 0.1;
420+ attrs.withOffset = false;
421+ attrs.quantBits = 8;
422+ ge::DataType dataType = ge::DT_FLOAT;
423+
424+ IfmrTilingData tilingdata;
425+
426+ auto ret = IFMRTestCase(input_shapes, output_shapes, attrs, dataType, tilingdata);
427+ EXPECT_EQ(ret, ge::GRAPH_FAILED);
428+}
Aquant/ifmr/tests/op_kernel/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+# 使用AddOpTestCase
12+# param1：算子名称，以kernel方式命名
13+# param2：soc版本，多个以分号分隔，例如："ascend950pr_9599;AscendB1"
14+# param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔，例如："-DDTYPE_X=float -DDTYPE_Y=float"
15+# param4：该算子依赖的所有tiling源码文件
16+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
17+ AddOpTestCase(ifmr "ascend950" "")
18+endif()
Aquant/ifmr/tests/op_kernel/ifmr_data/gen_data.py+130-0
@@ -0,0 +1,130 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+"""
11+gen_data.py
12+"""
13+import sys
14+from functools import reduce
15+ 
16+import numpy as np
17+import torch
18+ 
19+ 
20+class IfmrAlgNpu():
21+ def __init__(self, bins_num=512,
22+ min_percentile=0.999999,
23+ max_percentile=0.999999,
24+ search_range=[0.7, 1.3],
25+ search_step=0.01,
26+ with_offset=True):
27+ self.ifmr_param = {
28+ "bins_num": bins_num,
29+ "min_percentile": min_percentile,
30+ "max_percentile": max_percentile,
31+ "search_range": search_range,
32+ "search_step": search_step,
33+ "with_offset": with_offset
34+ }
35+ 
36+ def forward(self, data):
37+ bins_num = self.ifmr_param["bins_num"]
38+ min_percentile = self.ifmr_param["min_percentile"]
39+ max_percentile = self.ifmr_param["max_percentile"]
40+ search_range = self.ifmr_param["search_range"]
41+ search_step = self.ifmr_param["search_step"]
42+ with_offset = self.ifmr_param["with_offset"]
43+ 
44+ # process
45+ data = data.astype(np.float32)
46+ data_shape = data.shape
47+ data_type = data.dtype
48+ # 数据预处理
49+ data_max = np.max(data)
50+ data_min = np.min(data)
51+ if data_min > 0:
52+ data_min = 0
53+ if data_max < 0:
54+ data_max = 0
55+ data_num = reduce(lambda x, y: x * y, data_shape)
56+ data_max = np.array([data_max], dtype=data_type)
57+ data_min = np.array([data_min], dtype=data_type)
58+ # 计算累加和
59+ bins, threshold = np.histogram(data, bins_num)
60+ cumsum = np.cumsum(bins).astype(np.int32)
61+ cdf = cumsum / data_num
62+ 
63+ print('---------ifmr debug------------')
64+ print('min', data_min)
65+ print('max', data_max)
66+ print('cdf', cumsum)
67+ print('-------------------------------')
68+ # 生成 scale&offset
69+ max_index = np.where(cdf > max_percentile, 0, 1).sum()
70+ min_index = np.where(cdf > 1 - min_percentile, 0, 1).sum()
71+ max_init = max_index / bins_num * (data_max - data_min) + data_min
72+ min_init = min_index / bins_num * (data_max - data_min) + data_min
73+ step = np.arange(search_range[0], search_range[1], search_step)
74+ max_list = max_init * step
75+ min_list = min_init * np.ones(step.shape)
76+ scale = (max_list - min_list) / 255
77+ 
78+ offset = np.round(min_list / scale)
79+ offset = -(offset + 128)
80+ 
81+ # 找到最优
82+ data_list = data.flatten()
83+ loss_list = np.zeros(len(step))
84+ for i in range(len(step)):
85+ quant_data_list = np.round(data_list / scale[i]) + offset[i]
86+ np.clip(quant_data_list, -128, 127, out=quant_data_list)
87+ quant_data_list = (quant_data_list - offset[i]) * scale[i]
88+ loss = np.sum(np.square(quant_data_list - data_list))
89+ loss_list[i] = loss
90+ 
91+ index = np.unravel_index(np.argmin(loss_list), loss_list.shape)
92+ best_scale = scale[index]
93+ best_offset = offset[index]
94+ 
95+ return best_scale, best_offset
96+ 
97+ 
98+def gen_golden_data_simple(shape, min, max,
99+ bins_num=512, min_percentile=0.999999, max_percentile=0.999999,
100+ search_range=[0.7, 1.3], search_step=0.01, with_offset=True):
101+ 
102+ bins_num = int(bins_num)
103+ min = float(min)
104+ max = float(max)
105+ shape = int(shape)
106+ 
107+ if (min == max):
108+ input_np = np.random.uniform(0, 1024, size=shape).astype(float)
109+ else:
110+ input_np = np.random.uniform(min, max, size=shape).astype(float)
111+ 
112+ input_x = torch.from_numpy(input_np).reshape(shape)
113+ input_x_fp32 = input_x.to(torch.float32)
114+ hist = torch.histc(input_x_fp32, bins_num)
115+ cumsum = torch.cumsum(hist, 0)
116+ ifmr_mod = IfmrAlgNpu(bins_num, min_percentile,
117+ max_percentile, search_range, search_step, with_offset)
118+ golen_scale, golden_offset = ifmr_mod.forward(input_np)
119+ # save inputs
120+ input_x_fp32.numpy().tofile("./inputs.bin")
121+ torch.max(input_x_fp32).numpy().tofile("./inputs_max.bin")
122+ torch.min(input_x_fp32).numpy().tofile("./inputs_min.bin")
123+ cumsum.detach().numpy().tofile("./cumsum.bin")
124+ # save outputs
125+ golen_scale.tofile("./golen_scale.bin")
126+ golden_offset.tofile("./golden_offset.bin")
127+ 
128+ 
129+if __name__ == "__main__":
130+ gen_golden_data_simple(*sys.argv[1:])
Aquant/ifmr/tests/op_kernel/test_ifmr.cpp+97-0
@@ -0,0 +1,97 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+#include "test_ifmr.h"
15+ 
16+#ifdef __CCE_KT_TEST__
17+#include "tikicpulib.h"
18+#include "../data_utils.h"
19+#include "string.h"
20+#include <iostream>
21+#include <string>
22+#endif
23+#include <cstdint>
24+#include "../../../op_kernel/ifmr.cpp"
25+#include "../../../op_kernel/ifmr_tiling_data.h"
26+ 
27+using namespace std;
28+ 
29+extern "C" __global__ __aicore__ void ifmr(GM_ADDR self, GM_ADDR min, GM_ADDR max, GM_ADDR binsCount, GM_ADDR workspace, GM_ADDR tiling);
30+class ifmr_test : public testing::Test {
31+ protected:
32+ static void SetUpTestCase() {
33+ cout << "ifmr_test SetUp\n" << endl;
34+ }
35+ static void TearDownTestCase() {
36+ cout << "ifmr_test TearDown\n" << endl;
37+ }
38+};
39+ 
40+IfmrTilingData* FakeGetTilingData(uint64_t tilingKey, uint8_t *tiling, uint32_t blockDim) {
41+ IfmrTilingData* tilingData = reinterpret_cast<IfmrTilingData*>(tiling);
42+ 
43+ tilingData->minPercentile = 0.999999;
44+ tilingData->maxPercentile = 0.999999;
45+ tilingData->searchRange[0] = 0.7;
46+ tilingData->searchRange[1] = 1.3;
47+ tilingData->searchStep = 0.01;
48+ tilingData->withOffset = true;
49+ tilingData->quantBits = 8;
50+ tilingData->dataLength = 320000;
51+ tilingData->cumsumLength = 512;
52+ 
53+ return tilingData;
54+}
55+ 
56+TEST_F(ifmr_test, test_case_0) {
57+ int64_t totalLength = 320000;
58+ int64_t bins = 512;
59+ 
60+ // inputs
61+ size_t inputs_size = totalLength * sizeof(float);
62+ size_t min_size = sizeof(float);
63+ size_t max_size = sizeof(float);
64+ size_t cumsum_size = bins * sizeof(float);
65+ size_t tiling_data_size = sizeof(IfmrTilingData);
66+ 
67+ uint8_t *inputs = (uint8_t*)AscendC::GmAlloc(inputs_size);
68+ uint8_t *min = (uint8_t*)AscendC::GmAlloc(min_size);
69+ uint8_t *max = (uint8_t*)AscendC::GmAlloc(max_size);
70+ uint8_t *cumsum = (uint8_t*)AscendC::GmAlloc(cumsum_size);
71+ uint8_t *workspace = (uint8_t *)AscendC::GmAlloc(1024 * 16 * 1024);
72+ uint8_t *tiling = (uint8_t *)AscendC::GmAlloc(tiling_data_size);
73+ uint32_t blockDim = 1; //cpu模拟使用单核
74+ system("cp -r ../ifmr_data ./");
75+ system("chmod -R 755 ./ifmr_data/");
76+ system("cd ./ifmr_data/ && rm -rf ./*bin");
77+ system("cd ./ifmr_data/ && python3 gen_data.py 320000 0 1 512");
78+ 
79+ char * path_ = get_current_dir_name();
80+ string path(path_);
81+ ReadFile(path + "/ifmr_data/inputs.bin", inputs_size, inputs, inputs_size);
82+ ReadFile(path + "/ifmr_data/inputs_min.bin", min_size, min, min_size);
83+ ReadFile(path + "/ifmr_data/inputs_max.bin", max_size, max, max_size);
84+ ReadFile(path + "/ifmr_data/cumsum.bin", cumsum_size, cumsum, cumsum_size);
85+ uint64_t tilingKey = 0;
86+ auto tilingData = FakeGetTilingData(tilingKey, tiling, blockDim);
87+ ICPU_SET_TILING_KEY(tilingKey);
88+ ICPU_RUN_KF(ifmr, blockDim, inputs, min, max, cumsum, workspace, (uint8_t*)(tilingData));
89+ 
90+ AscendC::GmFree(inputs);
91+ AscendC::GmFree(min);
92+ AscendC::GmFree(max);
93+ AscendC::GmFree(cumsum);
94+ AscendC::GmFree(workspace);
95+ AscendC::GmFree(tiling);
96+ free(path_);
97+}
Aquant/ifmr/tests/op_kernel/test_ifmr.h+51-0
@@ -0,0 +1,51 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef TEST_IFMR_H
12+#define TEST_IFMR_H
13+ 
14+#include "kernel_tiling/kernel_tiling.h"
15+struct IfmrTilingData {
16+ float minPercentile;
17+ float maxPercentile;
18+ float searchRange[2];
19+ float searchStep;
20+ bool withOffset;
21+ int quantBits;
22+ int dataLength;
23+ int cumsumLength;
24+};
25+ 
26+#define DTYPE_X int64_t
27+ 
28+#pragma pack(1)
29+ 
30+#pragma pack()
31+ 
32+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
33+ __ubuf__ tilingStruct* tilingDataPointer = \
34+ reinterpret_cast<__ubuf__ tilingStruct*>((__ubuf__ uint8_t*)(tilingPointer));
35+ 
36+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
37+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
38+ 
39+#define GET_TILING_DATA(tilingData, tilingPointer) \
40+ IfmrTilingData tilingData; \
41+ INIT_TILING_DATA(IfmrTilingData, tilingDataPointer, tilingPointer); \
42+ (tilingData).minPercentile = tilingDataPointer->minPercentile; \
43+ (tilingData).maxPercentile = tilingDataPointer->maxPercentile; \
44+ (tilingData).searchRange[0] = tilingDataPointer->searchRange[0]; \
45+ (tilingData).searchRange[1] = tilingDataPointer->searchRange[1]; \
46+ (tilingData).searchStep = tilingDataPointer->searchStep; \
47+ (tilingData).withOffset = tilingDataPointer->withOffset; \
48+ (tilingData).quantBits = tilingDataPointer->quantBits; \
49+ (tilingData).dataLength = tilingDataPointer->dataLength; \
50+ (tilingData).cumsumLength = tilingDataPointer->cumsumLength;
51+#endif // TEST_IFMR_H