Star908
1826
代码
代码
Issues226
Pull Requests308
Actions
项目讨论
Wiki
分析
Star908
1826
已合并
add forward deformable_offsets #1504
demon_coffee创建于 2月4日
add forward deformable_offsets #1504
已合并
demon_coffee创建于 2月4日
从demon_coffee/ops-nndeformable_offsets合入到cann/ops-nnmaster
共 27 个文件变更+2998-0
Aconv/deformable_offsets/CMakeLists.txt+22-0
@@ -0,0 +1,22 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
W
Wwangyongguang2月10日

2026？

likedislike
demon_coffee
demon_coffee
2月10日 评论：
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+ 
14+if(NOT ENABLE_TEST AND NOT BENCHMARK)
15+ list(REMOVE_ITEM CURRENT_DIRS tests)
16+endif()
17+ 
18+foreach(SUB_DIR ${CURRENT_DIRS})
19+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
20+ add_subdirectory(${SUB_DIR})
21+ endif()
22+endforeach()
Aconv/deformable_offsets/README.md+113-0
@@ -0,0 +1,113 @@
1+# DeformableOffsets
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+ 
11+## 功能说明
12+ 
13+- 算子功能：用于计算变形卷积（Deformable Convolution）输出的函数。通过引入偏移参数offsets，使得卷积核在输入特征图上的位置可以动态调整，从而适配不规则的集合变化。
14+ 
15+## 参数说明
16+ 
17+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
18+ <col style="width: 100px">
19+ <col style="width: 150px">
20+ <col style="width: 280px">
21+ <col style="width: 330px">
22+ <col style="width: 120px">
23+ </colgroup>
24+ <thead>
25+ <tr>
26+ <th>参数名</th>
27+ <th>输入/输出/属性</th>
28+ <th>描述</th>
29+ <th>数据类型</th>
30+ <th>数据格式</th>
31+ </tr></thead>
32+ <tbody>
33+ <tr>
34+ <td>x</td>
35+ <td>输入</td>
36+ <td>输入的坐标。</td>
37+ <td>DT_FLOAT16, DT_FLOAT, DT_BF16</td>
38+ <td>NHWC</td>
39+ </tr>
40+ <tr>
41+ <td>offsets</td>
42+ <td>输入</td>
43+ <td>偏移值。</td>
44+ <td>DT_FLOAT16, DT_FLOAT, DT_BF16</td>
45+ <td>NHWC</td>
46+ </tr>
47+ <tr>
48+ <td>y</td>
49+ <td>输出</td>
50+ <td>变形卷积的输出。</td>
51+ <td>DT_FLOAT16, DT_FLOAT, DT_BF16</td>
52+ <td>NHWC</td>
53+ </tr>
54+ <tr>
55+ <td>strides</td>
56+ <td>属性</td>
57+ <td>指定卷积核在高度和宽度方向上的步幅。</td>
58+ <td>INT64</td>
59+ <td>-</td>
60+ </tr>
61+ <tr>
62+ <td>pads</td>
63+ <td>属性</td>
64+ <td>指定在输入的高度和宽度方向上添加的填充。</td>
65+ <td>INT64</td>
66+ <td>-</td>
67+ </tr>
68+ <tr>
69+ <td>ksize</td>
70+ <td>属性</td>
71+ <td>指定卷积核的大小。</td>
72+ <td>INT64</td>
73+ <td>-</td>
74+ </tr>
75+ <tr>
76+ <td>dilations</td>
77+ <td>属性</td>
78+ <td>指定每个维度上的膨胀因子。</td>
79+ <td>INT64</td>
80+ <td>-</td>
81+ </tr>
82+ <tr>
83+ <td>data_format</td>
84+ <td>属性</td>
85+ <td>指定输入 x 的数据格式。</td>
86+ <td>STRING</td>
87+ <td>-</td>
88+ </tr>
89+ <tr>
90+ <td>deformable_groups</td>
91+ <td>属性</td>
92+ <td>指定输入 x 在 C 轴上的分组数。</td>
93+ <td>INT32</td>
94+ <td>-</td>
95+ </tr>
96+ <tr>
97+ <td>modulated</td>
98+ <td>属性</td>
99+ <td>指定变形卷积的版本。</td>
100+ <td>BOOL</td>
101+ <td>-</td>
102+ </tr>
103+ </tbody></table>
104+ 
105+## 约束说明
106+ 
107+无
108+ 
109+## 调用说明
110+ 
111+| 调用方式 | 样例代码 | 说明 |
112+|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
113+| 图模式 | [test_geir_deformable_offsets](./examples/test_geir_deformable_offsets.cpp) | 通过[算子IR](./op_graph/deformable_offsets_proto.h)构图方式调用DeformableOffsets算子。 |
Aconv/deformable_offsets/examples/test_geir_deformable_offsets.cpp+384-0
@@ -0,0 +1,384 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file test_geir_deformable_offsets.cpp
13+ * \brief
14+ */
15+ 
16+#include <iostream>
17+#include <fstream>
18+#include <string.h>
19+#include <stdint.h>
20+#include <vector>
21+#include <string>
22+#include <map>
23+#include "assert.h"
24+ 
25+#include "graph.h"
26+#include "types.h"
27+#include "tensor.h"
28+#include "ge_error_codes.h"
29+#include "ge_api_types.h"
30+#include "ge_api.h"
31+#include "array_ops.h"
32+#include "ge_ir_build.h"
33+ 
34+#include "experiment_ops.h"
35+#include "nn_other.h"
36+#include "../op_graph/deformable_offsets_proto.h"
37+ 
38+#define FAILED -1
39+#define SUCCESS 0
40+ 
41+using namespace ge;
42+using std::map;
43+using std::string;
44+using std::vector;
45+#define ADD_INPUT(inputIndex, inputName, inputDtype, inputShape) \
46+ vector<int64_t> placeholder##inputIndex##_shape = inputShape; \
47+ auto placeholder##inputIndex = op::Data("placeholder" + inputIndex).set_attr_index(0); \
48+ TensorDesc placeholder##inputIndex##_desc = \
49+ TensorDesc(ge::Shape(placeholder##inputIndex##_shape), ge::FORMAT_NHWC, inputDtype); \
50+ placeholder##inputIndex##_desc.SetPlacement(ge::kPlacementHost); \
51+ placeholder##inputIndex##_desc.SetFormat(ge::FORMAT_NHWC); \
52+ Tensor tensor_placeholder##inputIndex; \
53+ ret = GenOnesData( \
54+ placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, placeholder##inputIndex##_desc, inputDtype, \
55+ 2); \
56+ if (ret != SUCCESS) { \
57+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
58+ return FAILED; \
59+ } \
60+ placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \
61+ input.push_back(tensor_placeholder##inputIndex); \
62+ graph.AddOp(placeholder##inputIndex); \
63+ add1.set_input_##inputName(placeholder##inputIndex); \
64+ inputs.push_back(placeholder##inputIndex)
65+ 
66+#define ADD_CONST_INPUT(inputIndex, inputName, inputDtype, inputShape) \
67+ vector<int64_t> placeholder##inputIndex##_shape = inputShape; \
68+ auto placeholder##inputIndex = op::Const("placeholder" + inputIndex); \
69+ TensorDesc placeholder##inputIndex##_desc = \
70+ TensorDesc(ge::Shape(placeholder##inputIndex##_shape), FORMAT_NHWC, inputDtype); \
71+ placeholder##inputIndex##_desc.SetPlacement(ge::kPlacementHost); \
72+ placeholder##inputIndex##_desc.SetFormat(FORMAT_NHWC); \
73+ Tensor tensor_placeholder##inputIndex; \
74+ ret = GenOnesData( \
75+ placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, placeholder##inputIndex##_desc, inputDtype, \
76+ 2); \
77+ if (ret != SUCCESS) { \
78+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
79+ return FAILED; \
80+ } \
81+ placeholder##inputIndex.SetAttr("value", tensor_placeholder##inputIndex); \
82+ placeholder##inputIndex.update_output_desc_y(placeholder##inputIndex##_desc); \
83+ graph.AddOp(placeholder##inputIndex); \
84+ add1.set_input_##inputName(placeholder##inputIndex); \
85+ add1.update_input_desc_##inputName(placeholder##inputIndex##_desc); \
86+ inputs.push_back(placeholder##inputIndex)
87+ 
88+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
89+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), ge::FORMAT_NHWC, outputDtype); \
90+ add1.update_output_desc_##outputName(outputName##outputIndex##_desc)
91+ 
92+#define LOG_PRINT(message, ...) \
93+ do { \
94+ printf(message, ##__VA_ARGS__); \
95+ } while (0)
96+ 
97+#define ADD_INPUT_ATTR(attrName, attrValue) add1.set_attr_##attrName(attrValue)
98+ 
99+string GetTime()
100+{
101+ time_t timep;
102+ time(&timep);
103+ char tmp[64];
104+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
105+ return tmp;
106+}
107+ 
108+uint32_t GetDataTypeSize(DataType dt)
109+{
110+ uint32_t dilation = 1;
111+ uint32_t oneByte = 1;
112+ uint32_t twoByte = 2;
113+ uint32_t fourByte = 4;
114+ uint32_t eightByte = 8;
115+ 
116+ if (dt == ge::DT_FLOAT) {
117+ dilation = fourByte;
118+ } else if (dt == ge::DT_FLOAT16) {
119+ dilation = twoByte;
120+ } else if (dt == ge::DT_BF16) {
121+ dilation = twoByte;
122+ } else if (dt == ge::DT_INT16) {
123+ dilation = twoByte;
124+ } else if (dt == ge::DT_UINT16) {
125+ dilation = twoByte;
126+ } else if (dt == ge::DT_INT32) {
127+ dilation = fourByte;
128+ } else if (dt == ge::DT_UINT32) {
129+ dilation = fourByte;
130+ } else if (dt == ge::DT_INT64) {
131+ dilation = eightByte;
132+ } else if (dt == ge::DT_UINT64) {
133+ dilation = eightByte;
134+ } else if (dt == ge::DT_INT8) {
135+ dilation = oneByte;
136+ }
137+ return dilation;
138+}
139+ 
140+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value)
141+{
142+ input_tensor_desc.SetRealDimCnt(shapes.size());
143+ size_t size = 1;
144+ for (uint32_t i = 0; i < shapes.size(); i++) {
145+ size *= shapes[i];
146+ }
147+ uint32_t byteSizeFloat32 = 4;
148+ uint32_t data_len = size * byteSizeFloat32;
149+ float* pData = new (std::nothrow) float[size];
150+ 
151+ for (size_t i = 0; i < size; ++i) {
152+ *(pData + i) = value;
153+ }
154+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
155+ return SUCCESS;
156+}
157+ 
158+int32_t GenOnesData(
159+ vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type, int value)
160+{
161+ input_tensor_desc.SetRealDimCnt(shapes.size());
162+ size_t size = 1;
163+ for (uint32_t i = 0; i < shapes.size(); i++) {
164+ size *= shapes[i];
165+ }
166+ uint32_t data_len = size * GetDataTypeSize(data_type);
167+ int32_t* pData = new (std::nothrow) int32_t[data_len];
168+ for (uint32_t i = 0; i < size; ++i) {
169+ *(pData + i) = value;
170+ }
171+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
172+ return SUCCESS;
173+}
174+ 
175+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
176+{
177+ FILE* fp = fopen(bin_file.c_str(), "w");
178+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
179+ fclose(fp);
180+ return SUCCESS;
181+}
182+ 
183+int CreateOppInGraph(
184+ DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, std::vector<Operator>& outputs,
185+ Graph& graph)
186+{
187+ Status ret = SUCCESS;
188+ // 自定义代码：添加单算子定义到图中
189+ auto add1 = op::DeformableOffsets("deformable_offsets");
190+ vector<vector<int64_t>> shapes = {
191+ {1, 7, 11, 256},
192+ {1, 7, 11, 27},
193+ {1, 21, 33, 256},
194+ {1, 1, 1, 1},
195+ {1, 1, 1, 1},
196+ {3, 3},
197+ {1, 1, 1, 1}
198+ };
199+ 
200+ ADD_INPUT(1, x, inDtype, shapes[0]);
201+ ADD_INPUT(2, offsets, inDtype, shapes[1]);
202+ ADD_OUTPUT(3, y, inDtype, shapes[2]);
203+ 
204+ add1.set_attr_strides(shapes[3]);
205+ add1.set_attr_pads(shapes[4]);
206+ add1.set_attr_ksize(shapes[5]);
207+ add1.set_attr_dilations(shapes[6]);
208+ ADD_INPUT_ATTR(data_format, "NHWC");
209+ std::int64_t deformable_groups_ = 1;
210+ add1.set_attr_deformable_groups(deformable_groups_);
211+ bool modulated_ = true;
212+ add1.set_attr_modulated(modulated_);
213+ 
214+ outputs.push_back(add1);
215+ // 添加完毕
216+ return SUCCESS;
217+}
218+ 
219+bool InitEnv()
220+{
221+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
222+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
223+ Status ret = ge::GEInitialize(global_options);
224+ if (ret != SUCCESS) {
225+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
226+ return false;
227+ }
228+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
229+ return true;
230+}
231+ 
232+bool CreateAndConfigGraph(Graph& graph, std::vector<ge::Tensor>& input)
233+{
234+ printf("%s - INFO - [XIR]: Start to CreateAndConfigGraph\n", GetTime().c_str());
235+ std::vector<Operator> inputs{};
236+ std::vector<Operator> outputs{};
237+ 
238+ DataType inDtype = DT_FLOAT;
239+ std::cout << inDtype << std::endl;
240+ 
241+ Status ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
242+ if (ret != SUCCESS) {
243+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
244+ return false;
245+ }
246+ 
247+ if (!inputs.empty() && !outputs.empty()) {
248+ graph.SetInputs(inputs).SetOutputs(outputs);
249+ }
250+ return true;
251+}
252+ 
253+bool AddGraphToSession(ge::Session* session, Graph& graph, uint32_t graph_id)
254+{
255+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
256+ 
257+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
258+ 
259+ std::map<AscendString, AscendString> graph_options = {
260+ 
261+ };
262+ 
263+ Status ret = session->AddGraph(graph_id, graph, graph_options);
264+ if (ret != SUCCESS) {
265+ printf("%s - INFO - [XIR]: Add graph failed\n", GetTime().c_str());
266+ delete session;
267+ ge::GEFinalize();
268+ return false;
269+ }
270+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
271+ 
272+ return true;
273+}
274+ 
275+bool DumpAndRunGraph(
276+ ge::Session* session, Graph& graph, std::vector<ge::Tensor>& input, std::vector<ge::Tensor>& output,
277+ uint32_t graph_id)
278+{
279+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
280+ std::string file_path = "./dump";
281+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
282+ 
283+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
284+ 
285+ Status ret = session->RunGraph(graph_id, input, output);
286+ if (ret != SUCCESS) {
287+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
288+ delete session;
289+ ge::GEFinalize();
290+ return false;
291+ }
292+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
293+ return true;
294+}
295+ 
296+void ProcessInputData(std::vector<ge::Tensor>& input)
297+{
298+ int input_num = input.size();
299+ for (int i = 0; i < input_num; i++) {
300+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
301+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
302+ uint8_t* input_data_i = input[i].GetData();
303+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
304+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
305+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
306+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
307+ }
308+}
309+ 
310+void ProcessOutputData(std::vector<ge::Tensor>& output)
311+{
312+ int output_num = output.size();
313+ for (int i = 0; i < output_num; i++) {
314+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
315+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
316+ uint8_t* output_data_i = output[i].GetData();
317+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
318+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
319+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
320+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
321+ for (int64_t j = 0; j < output_shape; j++) {
322+ LOG_PRINT("result[%ld] is: %u\n", j, output_data_i[j]);
323+ }
324+ }
325+}
326+ 
327+int FinalizeRes()
328+{
329+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
330+ std::string error_str(error_msg.GetString());
331+ std::cout << "Error message: " << error_str << std::endl;
332+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
333+ std::string warning_str(warning_msg.GetString());
334+ std::cout << "Warning message: " << warning_str << std::endl;
335+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
336+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
337+ Status ret = ge::GEFinalize();
338+ if (ret != SUCCESS) {
339+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
340+ return FAILED;
341+ }
342+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
343+ return SUCCESS;
344+}
345+ 
346+int main(int argc, char* argv[])
347+{
348+ // 初始化环境
349+ if (!InitEnv()) {
350+ return FAILED;
351+ }
352+ 
353+ // 创建计算图
354+ const char* graph_name = "tc_ge_irrun_test";
355+ Graph graph(graph_name);
356+ std::vector<ge::Tensor> input;
357+ if (!CreateAndConfigGraph(graph, input)) {
358+ return FAILED;
359+ }
360+ 
361+ // 创建会话并添加图
362+ std::map<AscendString, AscendString> build_options = {
363+ 
364+ };
365+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
366+ ge::Session* session = new Session(build_options);
367+ 
368+ uint32_t graph_id = 0;
369+ if (!AddGraphToSession(session, graph, graph_id)) {
370+ return FAILED;
371+ }
372+ 
373+ // 执行图
374+ std::vector<ge::Tensor> output;
375+ if (!DumpAndRunGraph(session, graph, input, output, graph_id)) {
376+ return FAILED;
377+ }
378+ // 处理输入输出数据
379+ ProcessInputData(input);
380+ ProcessOutputData(output);
381+ 
382+ // 清理资源
383+ return FinalizeRes();
384+}
Aconv/deformable_offsets/op_graph/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) Huawei Technologies Co., Ltd. 2025. All rights reserved.
3+# This file is a part of the CANN Open Software.
4+# Licensed under 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.conv.deformable_offsets.graph_plugin.CMakeLists.txt ")
12+add_graph_plugin_sources()
Aconv/deformable_offsets/op_graph/deformable_offsets_proto.h+63-0
@@ -0,0 +1,63 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file deformable_offsets_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_CONV_DEFORMABLE_OFFSETS_PLUGIN_DEFORMABLE_OFFSETS_PROTO_H_
16+#define OPS_CONV_DEFORMABLE_OFFSETS_PLUGIN_DEFORMABLE_OFFSETS_PROTO_H_
17+ 
18+#include "graph/operator_reg.h"
19+ 
20+namespace ge {
21+/**
22+*@brief Computes the deformed convolution output with the expected input
23+* @par Inputs:
24+* Two inputs:
25+* @li x: A 4D tensor of input image. A tensor of type float16, float32, bfloat16. The format support NHWC.
26+* Shape support 4D.
27+* @li offsets: A tensor of type float16, float32, bfloat16. Deformation offset parameter.
28+* The format support NHWC. Shape support 4D. Has the same format and dtype as "x".
29+ 
30+*@par Attributes:
31+* @li strides: A tuple/list of 4 integers. The stride of the sliding window for
32+* height and width for H/W dimension. Required and no default value.
33+* @li pads: A tuple/list of 4 integers. Padding added to H/W dimension
34+* of the input. Required and no default value.
35+* @li ksize: A tuple/list of 2 integers. Kernel size. Required and no default value.
36+* @li dilations: A tuple/list of 4 integers. The dilation factor for each dimension
37+* of input. Defaults to [1, 1, 1, 1]
38+* @li data_format: An optional string from: "NCHW", "NHWC". The default value "NCHW" is not supported.
39+* Specify the data format of the input x. The format of the attribute
40+* @li deformable_groups: An optional int specify the c-axis grouping number of input x. Defaults to "1".
41+* @li modulated: An optional bool specify version of DeformableConv2D, true means v2, false means v1. Defaults to
42+"true".
43+* Only support true now.
44+ 
45+*@par Outputs:
46+* y: Deformed convolution output. A tensor of type float16, float32, bfloat16. The format support NHWC.
47+* Shape support 4D. Has the same format and dtype as input "x".
48+*/
49+REG_OP(DeformableOffsets)
50+ .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
51+ .INPUT(offsets, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
52+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
53+ .REQUIRED_ATTR(strides, ListInt)
54+ .REQUIRED_ATTR(pads, ListInt)
55+ .REQUIRED_ATTR(ksize, ListInt)
56+ .ATTR(dilations, ListInt, {1, 1, 1, 1})
57+ .ATTR(data_format, String, "NHWC")
58+ .ATTR(deformable_groups, Int, 1)
59+ .ATTR(modulated, Bool, true)
60+ .OP_END_FACTORY_REG(DeformableOffsets)
61+} // namespace ge
62+ 
63+#endif // OPS_CONV_DEFORMABLE_OFFSETS_PLUGIN_DEFORMABLE_OFFSETS_PROTO_H_
Aconv/deformable_offsets/op_host/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+# 设置算子定义时支持的芯片类型
12+set(SUPPORT_COMPUTE_UNIT "ascend950")
13+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14+set(SUPPORT_TILING_DIR "arch35")
15+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE deformable_offsets
16+ ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Aconv/deformable_offsets/op_host/arch35/deformable_offsets_tiling_arch35.cpp+413-0
@@ -0,0 +1,413 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file deformable_offsets_tiling_arch35.cpp
13+ * \brief deformable_offsets_tiling_arch35 impl info
14+ */
15+ 
16+#include "deformable_offsets_tiling_arch35.h"
17+#include <cstdint>
18+#include "log/log.h"
19+#include "util/math_util.h"
20+#include "util/shape_util.h"
21+#include "register/op_def_registry.h"
22+#include "tiling/tiling_api.h"
23+ 
24+namespace optiling {
25+static constexpr int32_t OUTPUT_Y_INDEX = 0;
26+static constexpr int32_t INPUT_X_INDEX = 0;
27+static constexpr int32_t INPUT_OFFSET_INDEX = 1;
28+static constexpr int32_t DIM_NUM_4D = 4;
29+static constexpr int32_t STRIDES_INDEX = 0;
30+static constexpr int32_t PADS_INDEX = 1;
31+static constexpr int32_t KSIZE_INDEX = 2;
32+static constexpr int32_t DILATIONS_INDEX = 3;
33+static constexpr int32_t DATA_FORMAT_INDEX = 4;
34+static constexpr int32_t DEFORMABLE_GROUPS_ATTR_INDEX = 5;
35+static constexpr int32_t MODULATE_INDEX = 6;
36+static constexpr int32_t LIST_INDEX_0 = 0;
37+static constexpr int32_t LIST_INDEX_1 = 1;
38+static constexpr int32_t LIST_INDEX_2 = 2;
39+static constexpr int32_t LIST_INDEX_3 = 3;
40+static constexpr int32_t POINT_WEIGHT_SIZE = 3;
41+static constexpr int32_t POINT_NOT_WEIGHT_SIZE = 2;
42+static constexpr uint32_t MAX_CORE_NUM = 64;
43+static constexpr uint32_t VF_MAX_THREAD_NUM = 512;
44+static constexpr uint32_t EXCEPTED_KERNEL_SIZE = 2;
45+static constexpr uint32_t SIMT_COMMON_TILING_KEY = 1000;
46+static const std::set<ge::DataType> supportDtype = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
47+ 
48+ge::graphStatus CheckDeformableOffsetParams(
49+ gert::TilingContext* context, gert::Shape& inputXShape, gert::Shape& inputOffsetShape, gert::Shape& outputShapeInfo)
50+{
51+ // input data Format
52+ auto inputXdesc = context->GetInputDesc(INPUT_X_INDEX);
53+ OP_CHECK_NULL_WITH_CONTEXT(context, inputXdesc);
54+ auto inputXDataFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(inputXdesc->GetStorageFormat()));
55+ auto inputXDataType = inputXdesc->GetDataType();
56+ OP_CHECK_IF(
57+ supportDtype.count(inputXDataType) == 0, OP_LOGE(context->GetNodeName(), "X not support data type"),
58+ return ge::GRAPH_FAILED);
59+ // get input x shape info
60+ auto inputX = context->GetInputShape(INPUT_X_INDEX);
61+ OP_CHECK_IF(inputX == nullptr, OP_LOGE(context->GetNodeName(), "Get input x failed"), return ge::GRAPH_FAILED);
62+ inputXShape = inputX->GetStorageShape();
63+ int64_t inputXShapeDims = inputXShape.GetDimNum();
64+ 
65+ // get input offset shape info
66+ auto inputOffsetdesc = context->GetInputDesc(INPUT_OFFSET_INDEX);
67+ OP_CHECK_NULL_WITH_CONTEXT(context, inputOffsetdesc);
68+ auto inputYDataType = inputOffsetdesc->GetDataType();
69+ OP_CHECK_IF(
70+ supportDtype.count(inputYDataType) == 0, OP_LOGE(context->GetNodeName(), "offset not support data type"),
71+ return ge::GRAPH_FAILED);
72+ auto inputOffset = context->GetInputShape(INPUT_OFFSET_INDEX);
73+ OP_CHECK_IF(
74+ inputOffset == nullptr, OP_LOGE(context->GetNodeName(), "Get input offset failed"), return ge::GRAPH_FAILED);
75+ inputOffsetShape = inputOffset->GetStorageShape();
76+ int64_t inputOffsetShapeDims = inputOffsetShape.GetDimNum();
77+ 
78+ // get output shape info
79+ auto outputYdesc = context->GetOutputDesc(OUTPUT_Y_INDEX);
80+ OP_CHECK_NULL_WITH_CONTEXT(context, outputYdesc);
81+ auto outputYDataType = outputYdesc->GetDataType();
82+ OP_CHECK_IF(
83+ supportDtype.count(outputYDataType) == 0, OP_LOGE(context->GetNodeName(), "Not support output data type"),
84+ return ge::GRAPH_FAILED);
85+ auto outputShape = context->GetOutputShape(OUTPUT_Y_INDEX);
86+ OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);
87+ outputShapeInfo = outputShape->GetStorageShape();
88+ 
89+ // check shape and format
90+ OP_CHECK_IF(
91+ inputXShapeDims != DIM_NUM_4D || inputOffsetShapeDims != DIM_NUM_4D,
92+ OP_LOGE(context->GetNodeName(), "X offset shape check failed"), return ge::GRAPH_FAILED);
93+ ge::Format inputOffsetDataFormat =
94+ static_cast<ge::Format>(ge::GetPrimaryFormat(inputOffsetdesc->GetStorageFormat()));
95+ OP_CHECK_IF(
96+ inputXDataFormat != ge::FORMAT_NHWC || inputOffsetDataFormat != ge::FORMAT_NHWC,
97+ OP_LOGE(context->GetNodeName(), "X offset format support [NHWC]"), return ge::GRAPH_FAILED);
98+ return ge::GRAPH_SUCCESS;
99+}
100+ 
101+ge::graphStatus UpdateStrideAndDilationByFormat(
102+ gert::TilingContext* context, const std::string format, DeformableOffsetAttr& deformableOffsetAttrInfo,
103+ const int64_t* stridesData, const int64_t* dilatesData)
104+{
105+ if (format == "NCHW") {
106+ deformableOffsetAttrInfo.strideH = stridesData[LIST_INDEX_2];
107+ deformableOffsetAttrInfo.strideW = stridesData[LIST_INDEX_3];
108+ deformableOffsetAttrInfo.dilationH = dilatesData[LIST_INDEX_2];
109+ deformableOffsetAttrInfo.dilationW = dilatesData[LIST_INDEX_3];
110+ OP_CHECK_IF(
111+ stridesData[LIST_INDEX_0] != 1 || stridesData[LIST_INDEX_1] != 1 || dilatesData[LIST_INDEX_0] != 1 ||
112+ dilatesData[LIST_INDEX_1] != 1,
113+ OP_LOGE(context->GetNodeName(), "Stride or dilation data index 0 and 1 value must equal 1"),
114+ return ge::GRAPH_FAILED);
115+ } else if (format == "NHWC") {
116+ deformableOffsetAttrInfo.strideH = stridesData[LIST_INDEX_1];
117+ deformableOffsetAttrInfo.strideW = stridesData[LIST_INDEX_2];
118+ deformableOffsetAttrInfo.dilationH = dilatesData[LIST_INDEX_1];
119+ deformableOffsetAttrInfo.dilationW = dilatesData[LIST_INDEX_2];
120+ OP_CHECK_IF(
121+ stridesData[LIST_INDEX_0] != 1 || stridesData[LIST_INDEX_3] != 1 || dilatesData[LIST_INDEX_0] != 1 ||
122+ dilatesData[LIST_INDEX_3] != 1,
123+ OP_LOGE(context->GetNodeName(), "Stride or dilation data index 0 and 3 value must equal 1"),
124+ return ge::GRAPH_FAILED);
125+ } else {
126+ OP_LOGE(context->GetNodeName(), "dataFormat %s is not support", format.c_str());
127+ return ge::GRAPH_FAILED;
128+ }
129+ OP_CHECK_IF(
130+ deformableOffsetAttrInfo.strideH == 0 || deformableOffsetAttrInfo.strideW == 0,
131+ OP_LOGE(context->GetNodeName(), "StrideH not equal to 0"), return ge::GRAPH_FAILED);
132+ return ge::GRAPH_SUCCESS;
133+}
134+ 
135+ge::graphStatus CheckDeformableOffsetAttrs(gert::TilingContext* context, DeformableOffsetAttr& deformableOffsetAttrInfo)
136+{
137+ // attr info
138+ auto attrs = context->GetAttrs();
139+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
140+ auto strides = attrs->GetAttrPointer<gert::ContinuousVector>(STRIDES_INDEX);
141+ OP_CHECK_NULL_WITH_CONTEXT(context, strides);
142+ OP_CHECK_IF(strides->GetSize() != DIM_NUM_4D, OP_LOGE("DeformableOffsetTiling", "Not stride support len"),
143+ return ge::GRAPH_FAILED);
144+ const int64_t* stridesData = reinterpret_cast<const int64_t*>(strides->GetData());
145+ 
146+ auto pads = attrs->GetAttrPointer<gert::ContinuousVector>(PADS_INDEX);
147+ OP_CHECK_NULL_WITH_CONTEXT(context, pads);
148+ OP_CHECK_IF(pads->GetSize() != DIM_NUM_4D, OP_LOGE("DeformableOffsetTiling", "Not pads support len"),
149+ return ge::GRAPH_FAILED);
150+ const int64_t* padsData = reinterpret_cast<const int64_t*>(pads->GetData());
151+ 
152+ auto ksizes = attrs->GetAttrPointer<gert::ContinuousVector>(KSIZE_INDEX);
153+ OP_CHECK_NULL_WITH_CONTEXT(context, ksizes);
154+ OP_CHECK_IF(ksizes->GetSize() != EXCEPTED_KERNEL_SIZE, OP_LOGE("DeformableOffsetTiling", "Not ksize support len"),
155+ return ge::GRAPH_FAILED);
156+ const int64_t* ksizesData = reinterpret_cast<const int64_t*>(ksizes->GetData());
157+ 
158+ auto dilates = attrs->GetAttrPointer<gert::ContinuousVector>(DILATIONS_INDEX);
159+ OP_CHECK_NULL_WITH_CONTEXT(context, dilates);
160+ OP_CHECK_IF(
161+ dilates->GetSize() != DIM_NUM_4D, OP_LOGE("DeformableOffsetTiling", "Not dilation support len"),
162+ return ge::GRAPH_FAILED);
163+ const int64_t* dilatesData = reinterpret_cast<const int64_t*>(dilates->GetData());
164+ 
165+ auto modulatePtr = attrs->GetAttrPointer<bool>(MODULATE_INDEX);
166+ OP_CHECK_NULL_WITH_CONTEXT(context, modulatePtr);
167+ bool isModulated = *modulatePtr;
168+ 
169+ auto deformableGroupsPtr = attrs->GetAttrPointer<int>(DEFORMABLE_GROUPS_ATTR_INDEX);
170+ OP_CHECK_NULL_WITH_CONTEXT(context, deformableGroupsPtr);
171+ deformableOffsetAttrInfo.deformableGroupsAttr = *deformableGroupsPtr;
172+ const std::string dataFormat = std::string(attrs->GetAttrPointer<char>(DATA_FORMAT_INDEX));
173+ OP_CHECK_IF(
174+ UpdateStrideAndDilationByFormat(context, dataFormat, deformableOffsetAttrInfo, stridesData, dilatesData) !=
175+ ge::GRAPH_SUCCESS,
176+ OP_LOGE(context->GetNodeName(), "get stride failed."), return ge::GRAPH_FAILED);
177+ // get attr value
178+ deformableOffsetAttrInfo.padsHeightUp = padsData[LIST_INDEX_0];
179+ deformableOffsetAttrInfo.padsHeightDown = padsData[LIST_INDEX_1];
180+ deformableOffsetAttrInfo.padsWidthLeft = padsData[LIST_INDEX_2];
181+ deformableOffsetAttrInfo.padsWidthRight = padsData[LIST_INDEX_3];
182+ deformableOffsetAttrInfo.dimKh = ksizesData[LIST_INDEX_0];
183+ deformableOffsetAttrInfo.dimKw = ksizesData[LIST_INDEX_1];
184+ deformableOffsetAttrInfo.offsetValueDim =
185+ static_cast<uint32_t>(isModulated ? POINT_WEIGHT_SIZE : POINT_NOT_WEIGHT_SIZE);
186+ OP_CHECK_IF(
187+ !isModulated, OP_LOGE(context->GetNodeName(), "Deformable modulate attr only support true."),
188+ return ge::GRAPH_FAILED);
189+ return ge::GRAPH_SUCCESS;
190+}
191+ 
192+ge::graphStatus CheckOffsetArgs(
193+ gert::TilingContext* context, DeformableOffsetsOffset& deformableOffsetsOffset, const gert::Shape inputOffsetShape,
194+ const gert::Shape outputShapeInfo, DeformableOffsetAttr deformableOffsetAttrInfo) {
195+ // input offset info
196+ uint32_t offsetWidth = inputOffsetShape.GetDim(LIST_INDEX_2);
197+ uint32_t offsetHeight = inputOffsetShape.GetDim(LIST_INDEX_1);
198+ uint32_t offsetChannel = inputOffsetShape.GetDim(LIST_INDEX_3);
199+ deformableOffsetsOffset.offsetBatchStride = offsetHeight * offsetWidth * offsetChannel;
200+ deformableOffsetsOffset.deformableGroups = static_cast<uint32_t>(offsetChannel /
201+ (deformableOffsetAttrInfo.dimKh * deformableOffsetAttrInfo.dimKw * deformableOffsetAttrInfo.offsetValueDim));
202+ OP_CHECK_IF(deformableOffsetsOffset.imgChannel % deformableOffsetsOffset.deformableGroups != 0,
203+ OP_LOGE(context->GetNodeName(), "Img channel(%u) not equal to %d", deformableOffsetsOffset.imgChannel,
204+ deformableOffsetsOffset.deformableGroups),
205+ return ge::GRAPH_FAILED);
206+ OP_CHECK_IF(
207+ deformableOffsetAttrInfo.deformableGroupsAttr != deformableOffsetsOffset.deformableGroups,
208+ OP_LOGE(
209+ context->GetNodeName(), "Deformable groups attr is %u is invalid when deformableGroups is %u",
210+ deformableOffsetAttrInfo.deformableGroupsAttr, deformableOffsetsOffset.deformableGroups),
211+ return ge::GRAPH_FAILED);
212+ 
213+ uint32_t outputHeight = outputShapeInfo.GetDim(LIST_INDEX_1);
214+ uint32_t outputWidth = outputShapeInfo.GetDim(LIST_INDEX_2);
215+ deformableOffsetsOffset.outputBatchStride = outputHeight * outputWidth * deformableOffsetsOffset.imgChannel;
216+ uint32_t offsetHeightSame = static_cast<uint32_t>(offsetHeight * deformableOffsetAttrInfo.dimKh);
217+ uint32_t offsetWidthSame = static_cast<uint32_t>(offsetWidth * deformableOffsetAttrInfo.dimKw);
218+ 
219+ OP_CHECK_IF(outputShapeInfo.GetDim(LIST_INDEX_1) != offsetHeightSame,
220+ OP_LOGE(context->GetNodeName(), "Img out height not equal to offsetHeight muti dimkh"),
221+ return ge::GRAPH_FAILED);
222+ 
223+ OP_CHECK_IF(outputShapeInfo.GetDim(LIST_INDEX_2) != offsetWidthSame,
224+ OP_LOGE(context->GetNodeName(), "Img out width not equal to offsetWidth muti dimkw"), return ge::GRAPH_FAILED);
225+ 
226+ OP_CHECK_IF(deformableOffsetsOffset.imgOutHeight != inputOffsetShape.GetDim(LIST_INDEX_1),
227+ OP_LOGE(context->GetNodeName(), "Img out height not equal to offsetHeight"), return ge::GRAPH_FAILED);
228+ 
229+ OP_CHECK_IF(deformableOffsetsOffset.imgOutWidth != inputOffsetShape.GetDim(LIST_INDEX_2),
230+ OP_LOGE(context->GetNodeName(), "Img out width not equal to offsetWidth"), return ge::GRAPH_FAILED);
231+ 
232+ OP_CHECK_IF(deformableOffsetsOffset.imgBatchNum != inputOffsetShape.GetDim(LIST_INDEX_0),
233+ OP_LOGE(context->GetNodeName(), "Offset batchNum not equal to input"), return ge::GRAPH_FAILED);
234+ 
235+ OP_CHECK_IF(deformableOffsetsOffset.imgBatchNum != outputShapeInfo.GetDim(LIST_INDEX_0),
236+ OP_LOGE(context->GetNodeName(), "Output batchNum not equal to input"), return ge::GRAPH_FAILED);
237+ // offset NHC (N, H, W, offsetValueDim, group=1, k_h, k_w)
238+ deformableOffsetsOffset.offsetKernelElementStride = static_cast<uint32_t>(
239+ deformableOffsetsOffset.deformableGroups * deformableOffsetAttrInfo.dimKh * deformableOffsetAttrInfo.dimKw);
240+ deformableOffsetsOffset.offsetPointStride =
241+ deformableOffsetAttrInfo.offsetValueDim * deformableOffsetsOffset.offsetKernelElementStride;
242+ deformableOffsetsOffset.offsetWidthStride = offsetWidth * deformableOffsetsOffset.offsetPointStride;
243+ return ge::GRAPH_SUCCESS;
244+}
245+ 
246+ge::graphStatus CalDeformableOffsetsOffset(
247+ gert::TilingContext* context, gert::Shape& inputOffsetShape, gert::Shape& outputShapeInfo,
248+ DeformableOffsetAttr& deformableOffsetAttrInfo, DeformableOffsetsOffset& deformableOffsetsOffset)
249+{
250+ deformableOffsetsOffset.imgWidthStride = deformableOffsetsOffset.imgWidth * deformableOffsetsOffset.imgChannel;
251+ deformableOffsetsOffset.imgBatchStride =
252+ deformableOffsetsOffset.imgHeight * deformableOffsetsOffset.imgWidth * deformableOffsetsOffset.imgChannel;
253+ // conved img
254+ deformableOffsetsOffset.imgOutHeight = static_cast<uint32_t>(
255+ (deformableOffsetsOffset.imgHeight + deformableOffsetAttrInfo.padsHeightUp +
256+ deformableOffsetAttrInfo.padsHeightDown -
257+ (deformableOffsetAttrInfo.dilationH * (deformableOffsetAttrInfo.dimKh - 1) + 1)) /
258+ deformableOffsetAttrInfo.strideH +
259+ 1);
260+ deformableOffsetsOffset.imgOutWidth = static_cast<uint32_t>(
261+ (deformableOffsetsOffset.imgWidth + deformableOffsetAttrInfo.padsWidthLeft +
262+ deformableOffsetAttrInfo.padsWidthRight -
263+ (deformableOffsetAttrInfo.dilationW * (deformableOffsetAttrInfo.dimKw - 1) + 1)) /
264+ deformableOffsetAttrInfo.strideW +
265+ 1);
266+ 
267+ OP_CHECK_IF(
268+ deformableOffsetsOffset.imgOutHeight <= 0, OP_LOGE(context->GetNodeName(), "ImgOutHeight must greate 0"),
269+ return ge::GRAPH_FAILED);
270+ OP_CHECK_IF(
271+ deformableOffsetsOffset.imgOutWidth <= 0, OP_LOGE(context->GetNodeName(), "ImgOutWidth must greate 0"),
272+ return ge::GRAPH_FAILED);
273+ OP_CHECK_IF(
274+ deformableOffsetsOffset.imgChannel <= 0, OP_LOGE(context->GetNodeName(), "Img channel must greate 0"),
275+ return ge::GRAPH_FAILED);
276+ 
277+ OP_CHECK_IF(
278+ CheckOffsetArgs(
279+ context, deformableOffsetsOffset, inputOffsetShape, outputShapeInfo, deformableOffsetAttrInfo) !=
280+ ge::GRAPH_SUCCESS,
281+ OP_LOGE(context->GetNodeName(), "check args failed."), return ge::GRAPH_FAILED);
282+ // output NHWC
283+ // real is (N, H, K_h, W, k_w, c)
284+ deformableOffsetsOffset.outputPointWidthStride =
285+ static_cast<uint32_t>(deformableOffsetAttrInfo.dimKw * deformableOffsetsOffset.imgChannel);
286+ deformableOffsetsOffset.outputWidthStride =
287+ deformableOffsetsOffset.imgOutWidth * deformableOffsetsOffset.outputPointWidthStride;
288+ deformableOffsetsOffset.outputKernelWidthStride =
289+ static_cast<uint32_t>(deformableOffsetAttrInfo.dimKh * deformableOffsetsOffset.outputWidthStride);
290+ deformableOffsetsOffset.numKernels = deformableOffsetsOffset.imgOutWidth * deformableOffsetsOffset.imgOutHeight *
291+ deformableOffsetsOffset.imgChannel * deformableOffsetsOffset.imgBatchNum;
292+ deformableOffsetsOffset.blockDimValue = static_cast<uint32_t>(
293+ (deformableOffsetsOffset.numKernels / VF_MAX_THREAD_NUM) +
294+ ((deformableOffsetsOffset.numKernels % VF_MAX_THREAD_NUM) == 0 ? 0 : 1));
295+ return ge::GRAPH_SUCCESS;
296+}
297+ 
298+void SetDeformableOffsetsTilingData(
299+ gert::TilingContext* context, DeformableOffsetsTilingDataSimt& deformableOffsetTilingData,
300+ const DeformableOffsetsOffset deformableOffsetsOffset, const DeformableOffsetAttr deformableOffsetAttrInfo)
301+{
302+ deformableOffsetTilingData.set_blockNum(deformableOffsetsOffset.blockDimValue);
303+ deformableOffsetTilingData.set_strideHeight(deformableOffsetAttrInfo.strideH);
304+ deformableOffsetTilingData.set_strideWidth(deformableOffsetAttrInfo.strideW);
305+ deformableOffsetTilingData.set_dilationHeight(deformableOffsetAttrInfo.dilationH);
306+ deformableOffsetTilingData.set_dilationWidth(deformableOffsetAttrInfo.dilationW);
307+ deformableOffsetTilingData.set_padsHeight(deformableOffsetAttrInfo.padsHeightUp);
308+ deformableOffsetTilingData.set_padsWidth(deformableOffsetAttrInfo.padsWidthLeft);
309+ deformableOffsetTilingData.set_dimKHeight(deformableOffsetAttrInfo.dimKh);
310+ deformableOffsetTilingData.set_dimKWidth(deformableOffsetAttrInfo.dimKw);
311+ deformableOffsetTilingData.set_imgChannel(deformableOffsetsOffset.imgChannel);
312+ deformableOffsetTilingData.set_imgWidth(deformableOffsetsOffset.imgWidth);
313+ deformableOffsetTilingData.set_imgHeight(deformableOffsetsOffset.imgHeight);
314+ deformableOffsetTilingData.set_imgWidthStride(deformableOffsetsOffset.imgWidthStride);
315+ deformableOffsetTilingData.set_imgOutHeight(deformableOffsetsOffset.imgOutHeight);
316+ deformableOffsetTilingData.set_imgOutWidth(deformableOffsetsOffset.imgOutWidth);
317+ deformableOffsetTilingData.set_offsetKernelElementStride(deformableOffsetsOffset.offsetKernelElementStride);
318+ deformableOffsetTilingData.set_offsetPointStride(deformableOffsetsOffset.offsetPointStride);
319+ deformableOffsetTilingData.set_offsetWidthStride(deformableOffsetsOffset.offsetWidthStride);
320+ deformableOffsetTilingData.set_offsetValueDim(deformableOffsetAttrInfo.offsetValueDim);
321+ deformableOffsetTilingData.set_deformableGroups(deformableOffsetsOffset.deformableGroups);
322+ deformableOffsetTilingData.set_outputPointWidthStride(deformableOffsetsOffset.outputPointWidthStride);
323+ deformableOffsetTilingData.set_outputWidthStride(deformableOffsetsOffset.outputWidthStride);
324+ deformableOffsetTilingData.set_outputKernelWidthStride(deformableOffsetsOffset.outputKernelWidthStride);
325+ deformableOffsetTilingData.set_numKernels(deformableOffsetsOffset.numKernels);
326+ deformableOffsetTilingData.set_imgBatchStride(deformableOffsetsOffset.imgBatchStride);
327+ deformableOffsetTilingData.set_offsetBatchStride(deformableOffsetsOffset.offsetBatchStride);
328+ deformableOffsetTilingData.set_outputBatchStride(deformableOffsetsOffset.outputBatchStride);
329+ deformableOffsetTilingData.set_imgBatchNum(deformableOffsetsOffset.imgBatchNum);
330+ deformableOffsetTilingData.SaveToBuffer(
331+ context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
332+}
333+ 
334+ge::graphStatus DeformableOffsetTiling(gert::TilingContext* context, int32_t maxCoreNum)
335+{
336+ OP_LOGD("DeformableOffsetTilingForAscendC", "DeformableOffsetTiling start");
337+ DeformableOffsetsTilingDataSimt deformableOffsetTilingData;
338+ gert::Shape inputXShape, inputOffsetShape, outputShapeInfo;
339+ 
340+ OP_CHECK_IF(
341+ CheckDeformableOffsetParams(context, inputXShape, inputOffsetShape, outputShapeInfo) != ge::GRAPH_SUCCESS,
342+ OP_LOGE(context->GetNodeName(), "check params failed."), return ge::GRAPH_FAILED);
343+ DeformableOffsetAttr deformableOffsetAttrInfo;
344+ OP_CHECK_IF(
345+ CheckDeformableOffsetAttrs(context, deformableOffsetAttrInfo) != ge::GRAPH_SUCCESS,
346+ OP_LOGE(context->GetNodeName(), "check attr failed."), return ge::GRAPH_FAILED);
347+ 
348+ DeformableOffsetsOffset deformableOffsetsOffset;
349+ // input x info
350+ deformableOffsetsOffset.imgBatchNum = inputXShape.GetDim(LIST_INDEX_0);
351+ deformableOffsetsOffset.imgChannel = inputXShape.GetDim(LIST_INDEX_3);
352+ deformableOffsetsOffset.imgWidth = inputXShape.GetDim(LIST_INDEX_2);
353+ deformableOffsetsOffset.imgHeight = inputXShape.GetDim(LIST_INDEX_1);
354+ OP_CHECK_IF(
355+ CalDeformableOffsetsOffset(
356+ context, inputOffsetShape, outputShapeInfo, deformableOffsetAttrInfo, deformableOffsetsOffset) !=
357+ ge::GRAPH_SUCCESS,
358+ OP_LOGE(context->GetNodeName(), "get offsets failed."), return ge::GRAPH_FAILED);
359+ deformableOffsetsOffset.blockDimValue =
360+ std::min(deformableOffsetsOffset.blockDimValue, static_cast<uint32_t>(maxCoreNum));
361+ context->SetTilingKey(SIMT_COMMON_TILING_KEY);
362+ context->SetBlockDim(deformableOffsetsOffset.blockDimValue);
363+ SetDeformableOffsetsTilingData(
364+ context, deformableOffsetTilingData, deformableOffsetsOffset, deformableOffsetAttrInfo);
365+ context->GetRawTilingData()->SetDataSize(deformableOffsetTilingData.GetDataSize());
366+ 
367+ size_t usrSize = 0;
368+ size_t sysWorkSpaceSize = 16 * 1024 * 1024;
369+ 
370+ size_t* userWorkSpaceSize = context->GetWorkspaceSizes(1);
371+ userWorkSpaceSize[0] = usrSize + sysWorkSpaceSize;
372+ OP_LOGD("DeformableOffsetTilingForAscendC", "DeformableOffsetTiling end");
373+ return ge::GRAPH_SUCCESS;
374+}
375+ 
376+ge::graphStatus DeformableOffsetTilingSimt(gert::TilingContext* context, int32_t maxCoreNum)
377+{
378+ return DeformableOffsetTiling(context, maxCoreNum);
379+}
380+ 
381+ge::graphStatus Tiling4DeformableOffsets(gert::TilingContext* context)
382+{
383+ OP_LOGI(context->GetNodeName(), "Tiling4DeformableOffsets running.");
384+ const TilingPrepareForDeformableOffsetsCompileInfo* compileInfo =
385+ context->GetCompileInfo<TilingPrepareForDeformableOffsetsCompileInfo>();
386+ 
387+ int32_t maxCoreNum = static_cast<int32_t>(compileInfo->coreNum);
388+ OP_CHECK_IF(
389+ DeformableOffsetTilingSimt(context, maxCoreNum) != ge::GRAPH_SUCCESS,
390+ OP_LOGE(context->GetNodeName(), "The simd tiling function failed"), return ge::GRAPH_FAILED);
391+ return ge::GRAPH_SUCCESS;
392+}
393+ 
394+ge::graphStatus Tiling4PrepareDeformableOffsets(gert::TilingParseContext* context)
395+{
396+ OP_LOGD(context->GetNodeName(), "begin to do TilingPrepare4DeformableOffsets.");
397+ 
398+ auto compileInfo = context->GetCompiledInfo<TilingPrepareForDeformableOffsetsCompileInfo>();
399+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
400+ auto platformInfo = context->GetPlatformInfo();
401+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
402+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
403+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
404+ OP_CHECK_IF(
405+ (compileInfo->coreNum <= 0), OP_LOGE(context->GetNodeName(), "The core num is invalid."),
406+ return ge::GRAPH_FAILED);
407+ return ge::GRAPH_SUCCESS;
408+}
409+ 
410+IMPL_OP_OPTILING(DeformableOffsets)
411+ .Tiling(Tiling4DeformableOffsets)
412+ .TilingParse<TilingPrepareForDeformableOffsetsCompileInfo>(Tiling4PrepareDeformableOffsets);
413+} // namespace optiling
Aconv/deformable_offsets/op_host/arch35/deformable_offsets_tiling_arch35.h+106-0
@@ -0,0 +1,106 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file deformable_offsets_tiling_arch35.h
13+ * \brief deformable_offsets_tiling_arch35 info
14+ */
15+ 
16+#ifndef DEFORMABLE_OFFSETS_TILING_ARCH35_H
17+#define DEFORMABLE_OFFSETS_TILING_ARCH35_H
18+ 
19+#include <cstdint>
20+ 
21+#include "util/shape_util.h"
22+#include "register/tilingdata_base.h"
23+#include "op_common/op_host/util/opbase_export.h"
24+#include "register/op_impl_registry.h"
25+ 
26+namespace optiling {
27+BEGIN_TILING_DATA_DEF(DeformableOffsetsTilingDataSimt);
28+TILING_DATA_FIELD_DEF(uint32_t, blockNum);
29+TILING_DATA_FIELD_DEF(uint32_t, strideHeight);
30+TILING_DATA_FIELD_DEF(uint32_t, strideWidth);
31+TILING_DATA_FIELD_DEF(uint32_t, dilationHeight);
32+TILING_DATA_FIELD_DEF(uint32_t, dilationWidth);
33+TILING_DATA_FIELD_DEF(uint32_t, padsHeight);
34+TILING_DATA_FIELD_DEF(uint32_t, padsWidth);
35+TILING_DATA_FIELD_DEF(uint32_t, dimKHeight);
36+TILING_DATA_FIELD_DEF(uint32_t, dimKWidth);
37+TILING_DATA_FIELD_DEF(uint32_t, imgChannel);
38+TILING_DATA_FIELD_DEF(uint32_t, imgWidth);
39+TILING_DATA_FIELD_DEF(uint32_t, imgHeight);
40+TILING_DATA_FIELD_DEF(uint32_t, imgWidthStride);
41+TILING_DATA_FIELD_DEF(uint32_t, imgOutHeight);
42+TILING_DATA_FIELD_DEF(uint32_t, imgOutWidth);
43+TILING_DATA_FIELD_DEF(uint32_t, offsetKernelElementStride);
44+TILING_DATA_FIELD_DEF(uint32_t, offsetPointStride);
45+TILING_DATA_FIELD_DEF(uint32_t, offsetWidthStride);
46+TILING_DATA_FIELD_DEF(uint32_t, offsetValueDim);
47+TILING_DATA_FIELD_DEF(uint32_t, deformableGroups);
48+TILING_DATA_FIELD_DEF(uint32_t, outputPointWidthStride);
49+TILING_DATA_FIELD_DEF(uint32_t, outputWidthStride);
50+TILING_DATA_FIELD_DEF(uint32_t, outputKernelWidthStride);
51+TILING_DATA_FIELD_DEF(uint32_t, numKernels);
52+TILING_DATA_FIELD_DEF(uint32_t, imgBatchStride);
53+TILING_DATA_FIELD_DEF(uint32_t, offsetBatchStride);
54+TILING_DATA_FIELD_DEF(uint32_t, outputBatchStride);
55+TILING_DATA_FIELD_DEF(uint32_t, imgBatchNum);
56+END_TILING_DATA_DEF;
57+ 
58+struct TilingPrepareForDeformableOffsetsCompileInfo {
59+ int64_t coreNum;
60+ int64_t ubSize;
61+};
62+ 
63+struct DeformableOffsetAttr {
64+ int64_t strideH;
65+ int64_t strideW;
66+ int64_t dilationH;
67+ int64_t dilationW;
68+ int64_t padsHeightUp;
69+ int64_t padsHeightDown;
70+ int64_t padsWidthLeft;
71+ int64_t padsWidthRight;
72+ int64_t dimKh;
73+ int64_t dimKw;
74+ uint32_t deformableGroupsAttr;
75+ uint32_t offsetValueDim;
76+};
77+ 
78+struct DeformableOffsetsOffset {
79+ uint32_t imgBatchNum;
80+ uint32_t imgChannel;
81+ uint32_t imgWidth;
82+ uint32_t imgHeight;
83+ uint32_t imgWidthStride;
84+ uint32_t imgBatchStride;
85+ uint32_t imgOutHeight;
86+ uint32_t imgOutWidth;
87+ uint32_t offsetBatchStride;
88+ uint32_t deformableGroups;
89+ uint32_t offsetKernelElementStride;
90+ uint32_t offsetPointStride;
91+ uint32_t offsetWidthStride;
92+ 
93+ uint32_t outputBatchStride;
94+ uint32_t outputPointWidthStride;
95+ uint32_t outputWidthStride;
96+ uint32_t outputKernelWidthStride;
97+ 
98+ uint32_t numKernels;
99+ uint32_t blockDimValue;
100+};
101+ 
102+REGISTER_TILING_DATA_CLASS(DeformableOffsets, DeformableOffsetsTilingDataSimt)
103+ 
104+ge::graphStatus DeformableOffsetTilingSimt(gert::TilingContext* context, int32_t maxCoreNum);
105+} // namespace optiling
106+#endif
Aconv/deformable_offsets/op_host/config/ascend950/deformable_offsets_binary.json+467-0
@@ -0,0 +1,467 @@
1+{
2+ "op_type": "DeformableOffsets",
3+ "op_list": [
4+ {
5+ "bin_filename": "DeformableOffsets_0dbbb7ea43b1978c622d505927bce5d3",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "bfloat16",
11+ "format": "NHWC",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatFixed"
17+ },
18+ {
19+ "name": "offsets",
20+ "index": 1,
21+ "dtype": "bfloat16",
22+ "format": "NHWC",
23+ "paramType": "required",
24+ "shape": [
25+ -2
26+ ],
27+ "format_match_mode": "FormatFixed"
28+ }
29+ ],
30+ "outputs": [
31+ {
32+ "name": "y",
33+ "index": 0,
34+ "dtype": "bfloat16",
35+ "format": "NHWC",
36+ "paramType": "required",
37+ "shape": [
38+ -2
39+ ],
40+ "format_match_mode": "FormatFixed"
41+ }
42+ ],
43+ "attrs": [
44+ {
45+ "name": "strides",
46+ "dtype": "list_int",
47+ "value": null
48+ },
49+ {
50+ "name": "pads",
51+ "dtype": "list_int",
52+ "value": null
53+ },
54+ {
55+ "name": "ksize",
56+ "dtype": "list_int",
57+ "value": null
58+ },
59+ {
60+ "name": "dilations",
61+ "dtype": "list_int",
62+ "value": null
63+ },
64+ {
65+ "name": "data_format",
66+ "dtype": "string",
67+ "value": "NCHW"
68+ },
69+ {
70+ "name": "deformable_groups",
71+ "dtype": "int",
72+ "value": null
73+ },
74+ {
75+ "name": "modulated",
76+ "dtype": "bool",
77+ "value": true
78+ }
79+ ]
80+ },
81+ {
82+ "bin_filename": "DeformableOffsets_74099a9777b318fbe6da27158ecf991c",
83+ "inputs": [
84+ {
85+ "name": "x",
86+ "index": 0,
87+ "dtype": "bfloat16",
88+ "format": "NHWC",
89+ "paramType": "required",
90+ "shape": [
91+ -2
92+ ],
93+ "format_match_mode": "FormatFixed"
94+ },
95+ {
96+ "name": "offsets",
97+ "index": 1,
98+ "dtype": "bfloat16",
99+ "format": "NHWC",
100+ "paramType": "required",
101+ "shape": [
102+ -2
103+ ],
104+ "format_match_mode": "FormatFixed"
105+ }
106+ ],
107+ "outputs": [
108+ {
109+ "name": "y",
110+ "index": 0,
111+ "dtype": "bfloat16",
112+ "format": "NHWC",
113+ "paramType": "required",
114+ "shape": [
115+ -2
116+ ],
117+ "format_match_mode": "FormatFixed"
118+ }
119+ ],
120+ "attrs": [
121+ {
122+ "name": "strides",
123+ "dtype": "list_int",
124+ "value": null
125+ },
126+ {
127+ "name": "pads",
128+ "dtype": "list_int",
129+ "value": null
130+ },
131+ {
132+ "name": "ksize",
133+ "dtype": "list_int",
134+ "value": null
135+ },
136+ {
137+ "name": "dilations",
138+ "dtype": "list_int",
139+ "value": null
140+ },
141+ {
142+ "name": "data_format",
143+ "dtype": "string",
144+ "value": "NHWC"
145+ },
146+ {
147+ "name": "deformable_groups",
148+ "dtype": "int",
149+ "value": null
150+ },
151+ {
152+ "name": "modulated",
153+ "dtype": "bool",
154+ "value": true
155+ }
156+ ]
157+ },
158+ {
159+ "bin_filename": "DeformableOffsets_bb470772b7a6c739609b57a9acf7d265",
160+ "inputs": [
161+ {
162+ "name": "x",
163+ "index": 0,
164+ "dtype": "float16",
165+ "format": "NHWC",
166+ "paramType": "required",
167+ "shape": [
168+ -2
169+ ],
170+ "format_match_mode": "FormatFixed"
171+ },
172+ {
173+ "name": "offsets",
174+ "index": 1,
175+ "dtype": "float16",
176+ "format": "NHWC",
177+ "paramType": "required",
178+ "shape": [
179+ -2
180+ ],
181+ "format_match_mode": "FormatFixed"
182+ }
183+ ],
184+ "outputs": [
185+ {
186+ "name": "y",
187+ "index": 0,
188+ "dtype": "float16",
189+ "format": "NHWC",
190+ "paramType": "required",
191+ "shape": [
192+ -2
193+ ],
194+ "format_match_mode": "FormatFixed"
195+ }
196+ ],
197+ "attrs": [
198+ {
199+ "name": "strides",
200+ "dtype": "list_int",
201+ "value": null
202+ },
203+ {
204+ "name": "pads",
205+ "dtype": "list_int",
206+ "value": null
207+ },
208+ {
209+ "name": "ksize",
210+ "dtype": "list_int",
211+ "value": null
212+ },
213+ {
214+ "name": "dilations",
215+ "dtype": "list_int",
216+ "value": null
217+ },
218+ {
219+ "name": "data_format",
220+ "dtype": "string",
221+ "value": "NCHW"
222+ },
223+ {
224+ "name": "deformable_groups",
225+ "dtype": "int",
226+ "value": null
227+ },
228+ {
229+ "name": "modulated",
230+ "dtype": "bool",
231+ "value": true
232+ }
233+ ]
234+ },
235+ {
236+ "bin_filename": "DeformableOffsets_a54b7ddfee7ac692290eb7c13650d2dd",
237+ "inputs": [
238+ {
239+ "name": "x",
240+ "index": 0,
241+ "dtype": "float16",
242+ "format": "NHWC",
243+ "paramType": "required",
244+ "shape": [
245+ -2
246+ ],
247+ "format_match_mode": "FormatFixed"
248+ },
249+ {
250+ "name": "offsets",
251+ "index": 1,
252+ "dtype": "float16",
253+ "format": "NHWC",
254+ "paramType": "required",
255+ "shape": [
256+ -2
257+ ],
258+ "format_match_mode": "FormatFixed"
259+ }
260+ ],
261+ "outputs": [
262+ {
263+ "name": "y",
264+ "index": 0,
265+ "dtype": "float16",
266+ "format": "NHWC",
267+ "paramType": "required",
268+ "shape": [
269+ -2
270+ ],
271+ "format_match_mode": "FormatFixed"
272+ }
273+ ],
274+ "attrs": [
275+ {
276+ "name": "strides",
277+ "dtype": "list_int",
278+ "value": null
279+ },
280+ {
281+ "name": "pads",
282+ "dtype": "list_int",
283+ "value": null
284+ },
285+ {
286+ "name": "ksize",
287+ "dtype": "list_int",
288+ "value": null
289+ },
290+ {
291+ "name": "dilations",
292+ "dtype": "list_int",
293+ "value": null
294+ },
295+ {
296+ "name": "data_format",
297+ "dtype": "string",
298+ "value": "NHWC"
299+ },
300+ {
301+ "name": "deformable_groups",
302+ "dtype": "int",
303+ "value": null
304+ },
305+ {
306+ "name": "modulated",
307+ "dtype": "bool",
308+ "value": true
309+ }
310+ ]
311+ },
312+ {
313+ "bin_filename": "DeformableOffsets_5f976e94bb4de68c4617f442e9cb719c",
314+ "inputs": [
315+ {
316+ "name": "x",
317+ "index": 0,
318+ "dtype": "float32",
319+ "format": "NHWC",
320+ "paramType": "required",
321+ "shape": [
322+ -2
323+ ],
324+ "format_match_mode": "FormatFixed"
325+ },
326+ {
327+ "name": "offsets",
328+ "index": 1,
329+ "dtype": "float32",
330+ "format": "NHWC",
331+ "paramType": "required",
332+ "shape": [
333+ -2
334+ ],
335+ "format_match_mode": "FormatFixed"
336+ }
337+ ],
338+ "outputs": [
339+ {
340+ "name": "y",
341+ "index": 0,
342+ "dtype": "float32",
343+ "format": "NHWC",
344+ "paramType": "required",
345+ "shape": [
346+ -2
347+ ],
348+ "format_match_mode": "FormatFixed"
349+ }
350+ ],
351+ "attrs": [
352+ {
353+ "name": "strides",
354+ "dtype": "list_int",
355+ "value": null
356+ },
357+ {
358+ "name": "pads",
359+ "dtype": "list_int",
360+ "value": null
361+ },
362+ {
363+ "name": "ksize",
364+ "dtype": "list_int",
365+ "value": null
366+ },
367+ {
368+ "name": "dilations",
369+ "dtype": "list_int",
370+ "value": null
371+ },
372+ {
373+ "name": "data_format",
374+ "dtype": "string",
375+ "value": "NCHW"
376+ },
377+ {
378+ "name": "deformable_groups",
379+ "dtype": "int",
380+ "value": null
381+ },
382+ {
383+ "name": "modulated",
384+ "dtype": "bool",
385+ "value": true
386+ }
387+ ]
388+ },
389+ {
390+ "bin_filename": "DeformableOffsets_e0f33dbf4cacedb8080e0a61cbc85716",
391+ "inputs": [
392+ {
393+ "name": "x",
394+ "index": 0,
395+ "dtype": "float32",
396+ "format": "NHWC",
397+ "paramType": "required",
398+ "shape": [
399+ -2
400+ ],
401+ "format_match_mode": "FormatFixed"
402+ },
403+ {
404+ "name": "offsets",
405+ "index": 1,
406+ "dtype": "float32",
407+ "format": "NHWC",
408+ "paramType": "required",
409+ "shape": [
410+ -2
411+ ],
412+ "format_match_mode": "FormatFixed"
413+ }
414+ ],
415+ "outputs": [
416+ {
417+ "name": "y",
418+ "index": 0,
419+ "dtype": "float32",
420+ "format": "NHWC",
421+ "paramType": "required",
422+ "shape": [
423+ -2
424+ ],
425+ "format_match_mode": "FormatFixed"
426+ }
427+ ],
428+ "attrs": [
429+ {
430+ "name": "strides",
431+ "dtype": "list_int",
432+ "value": null
433+ },
434+ {
435+ "name": "pads",
436+ "dtype": "list_int",
437+ "value": null
438+ },
439+ {
440+ "name": "ksize",
441+ "dtype": "list_int",
442+ "value": null
443+ },
444+ {
445+ "name": "dilations",
446+ "dtype": "list_int",
447+ "value": null
448+ },
449+ {
450+ "name": "data_format",
451+ "dtype": "string",
452+ "value": "NHWC"
453+ },
454+ {
455+ "name": "deformable_groups",
456+ "dtype": "int",
457+ "value": null
458+ },
459+ {
460+ "name": "modulated",
461+ "dtype": "bool",
462+ "value": true
463+ }
464+ ]
465+ }
466+ ]
467+ }
Aconv/deformable_offsets/op_host/config/ascend950/deformable_offsets_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+[DeformableOffsets]
13+default=0
Aconv/deformable_offsets/op_host/deformable_offsets_def.cpp+56-0
@@ -0,0 +1,56 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file deformable_offsets_def.cpp
13+ * \brief deformable_offsets_def op_host
14+ */
15+ 
16+#include <cstdint>
17+#include "register/op_def_registry.h"
18+ 
19+namespace ops {
20+class DeformableOffsets : public OpDef {
21+public:
22+ explicit DeformableOffsets(const char* name) : OpDef(name)
23+ {
24+ this->Input("x")
25+ .ParamType(REQUIRED)
26+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
27+ .Format({ge::FORMAT_NHWC, ge::FORMAT_NHWC, ge::FORMAT_NHWC})
28+ .UnknownShapeFormat({ge::FORMAT_NHWC, ge::FORMAT_NHWC, ge::FORMAT_NHWC});
29+ this->Input("offsets")
30+ .ParamType(REQUIRED)
31+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
32+ .Format({ge::FORMAT_NHWC, ge::FORMAT_NHWC, ge::FORMAT_NHWC})
33+ .UnknownShapeFormat({ge::FORMAT_NHWC, ge::FORMAT_NHWC, ge::FORMAT_NHWC});
34+ this->Output("y")
35+ .ParamType(REQUIRED)
36+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
37+ .Format({ge::FORMAT_NHWC, ge::FORMAT_NHWC, ge::FORMAT_NHWC})
38+ .UnknownShapeFormat({ge::FORMAT_NHWC, ge::FORMAT_NHWC, ge::FORMAT_NHWC});
39+ this->Attr("strides").AttrType(REQUIRED).ListInt();
40+ this->Attr("pads").AttrType(REQUIRED).ListInt();
41+ this->Attr("ksize").AttrType(REQUIRED).ListInt();
42+ this->Attr("dilations").AttrType(OPTIONAL).ListInt({1, 1, 1, 1});
43+ this->Attr("data_format").AttrType(OPTIONAL).String("NHWC");
44+ this->Attr("deformable_groups").AttrType(OPTIONAL).Int(1);
45+ this->Attr("modulated").AttrType(OPTIONAL).Bool(true);
46+ 
47+ OpAICoreConfig aicoreConfig;
48+ aicoreConfig.DynamicCompileStaticFlag(true)
49+ .DynamicRankSupportFlag(true)
50+ .DynamicShapeSupportFlag(true)
51+ .ExtendCfgInfo("opFile.value", "deformable_offsets_apt");
52+ this->AICore().AddConfig("ascend950", aicoreConfig);
53+ }
54+};
55+OP_ADD(DeformableOffsets);
56+} // namespace ops
Aconv/deformable_offsets/op_host/deformable_offsets_infershape.cpp+164-0
@@ -0,0 +1,164 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file deformable_offsets_infershape.cpp
13+ * \brief
14+ */
15+ 
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "util/shape_util.h"
19+#include "op_common/op_host/util/opbase_export.h"
20+#include "register/op_impl_registry.h"
21+#include "runtime/storage_shape.h"
22+ 
23+using namespace ge;
24+namespace ops {
25+constexpr size_t kDimNum = 4U;
26+constexpr size_t kDilationsSize = 4U;
27+constexpr size_t kStridesSize = 4U;
28+constexpr size_t kKSizeSize = 2U;
29+constexpr size_t kPadsSize = 4U;
30+constexpr size_t kNCHWN = 0U;
31+constexpr size_t kNCHWC = 1U;
32+constexpr size_t kNCHWH = 2U;
33+constexpr size_t kNCHWW = 3U;
34+constexpr size_t kNHWCN = 0U;
35+constexpr size_t kNHWCH = 1U;
36+constexpr size_t kNHWCW = 2U;
37+constexpr size_t kNHWCC = 3U;
38+ 
39+static ge::graphStatus DeformableOffsetsInferShape(gert::InferShapeContext* context)
40+{
41+ auto attrs = context->GetAttrs();
42+ auto stridesPtr = attrs->GetListInt(0);
43+ OP_CHECK_NULL_WITH_CONTEXT(context, stridesPtr);
44+ auto padsPtr = attrs->GetListInt(1);
45+ OP_CHECK_NULL_WITH_CONTEXT(context, padsPtr);
46+ auto ksizePtr = attrs->GetListInt(2); // 2 is for ksize
47+ OP_CHECK_NULL_WITH_CONTEXT(context, ksizePtr);
48+ auto dilationsPtr = attrs->GetListInt(3); // 3 is for dilations
49+ OP_CHECK_NULL_WITH_CONTEXT(context, dilationsPtr);
50+ auto dataFormat = attrs->GetStr(4); // 4 is for dataFormat
51+ OP_CHECK_NULL_WITH_CONTEXT(context, dataFormat);
52+ auto dilations = dilationsPtr->GetData();
53+ OP_CHECK_NULL_WITH_CONTEXT(context, dilations);
54+ if (dilationsPtr->GetSize() != kDilationsSize) {
55+ OP_LOGE(context->GetNodeName(), "dilations list size should be 4, but got %zu", dilationsPtr->GetSize());
56+ return ge::GRAPH_FAILED;
57+ }
58+ auto strides = stridesPtr->GetData();
59+ OP_CHECK_NULL_WITH_CONTEXT(context, strides);
60+ if (stridesPtr->GetSize() != kStridesSize) {
61+ OP_LOGE(context->GetNodeName(), "strides list size should be 4, but got %zu", stridesPtr->GetSize());
62+ return ge::GRAPH_FAILED;
63+ }
64+ int64_t dilationsH;
65+ int64_t dilationsW;
66+ int64_t strideH;
67+ int64_t strideW;
68+ // 这里的data_format 仅仅表示属性的format，已经不能表示input/output的format
69+ if (strcmp(dataFormat, "NCHW") == 0) {
70+ dilationsH = dilations[kNCHWH];
71+ dilationsW = dilations[kNCHWW];
72+ strideH = strides[kNCHWH];
73+ strideW = strides[kNCHWW];
74+ } else if (strcmp(dataFormat, "NHWC") == 0) {
75+ dilationsH = dilations[kNHWCH];
76+ dilationsW = dilations[kNHWCW];
77+ strideH = strides[kNHWCH];
78+ strideW = strides[kNHWCW];
79+ } else {
80+ OP_LOGE(context->GetNodeName(), "dataFormat attr only support NCHW or NHWC, but got %s", dataFormat);
81+ return ge::GRAPH_FAILED;
82+ }
83+ 
84+ if ((strideH <= 0) || (strideW <= 0)) {
85+ OP_LOGE(
86+ context->GetNodeName(), "stride should be greater than 0, strideH [%ld], strideW [%ld]", strideH, strideW);
87+ return ge::GRAPH_FAILED;
88+ }
89+ 
90+ OP_CHECK_IF(
91+ (ksizePtr->GetSize() != kKSizeSize), OP_LOGE(context->GetNodeName(), "kSize list size should be 2"),
92+ return ge::GRAPH_FAILED);
93+ auto ksize = ksizePtr->GetData();
94+ OP_CHECK_NULL_WITH_CONTEXT(context, ksize);
95+ auto ksizeH = ksize[0];
96+ auto ksizeW = ksize[1];
97+ 
98+ auto dilKsizeH = (ksizeH - 1) * dilationsH + 1;
99+ auto dilKsizeW = (ksizeW - 1) * dilationsW + 1;
100+ 
101+ const gert::Shape* xShape = context->GetInputShape(0);
102+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
103+ OP_CHECK_IF(
104+ (xShape->GetDimNum() != kDimNum), OP_LOGE(context->GetNodeName(), "x rank should be 4D"),
105+ return ge::GRAPH_FAILED);
106+ const gert::Shape* offsetShape = context->GetInputShape(1);
107+ OP_CHECK_NULL_WITH_CONTEXT(context, offsetShape);
108+ OP_CHECK_IF(
109+ (offsetShape->GetDimNum() != kDimNum), OP_LOGE(context->GetNodeName(), "offset rank should be 4D"),
110+ return ge::GRAPH_FAILED);
111+ auto posH = strchr(dataFormat, 'H') - dataFormat;
112+ auto posW = strchr(dataFormat, 'W') - dataFormat;
113+ 
114+ auto xH = xShape->GetDim(posH);
115+ auto xW = xShape->GetDim(posW);
116+ auto offsetH = offsetShape->GetDim(posH);
117+ auto offsetW = offsetShape->GetDim(posW);
118+ auto pads = padsPtr->GetData();
119+ OP_CHECK_NULL_WITH_CONTEXT(context, pads);
120+ OP_CHECK_IF(
121+ (padsPtr->GetSize() != kPadsSize), OP_LOGE(context->GetNodeName(), "pads list size should be 4"),
122+ return ge::GRAPH_FAILED);
123+ auto padU = pads[0];
124+ auto padD = pads[1];
125+ auto padL = pads[2];
126+ auto padR = pads[3];
127+ 
128+ auto convOutH = (xH + padU + padD - dilKsizeH) / strideH + 1;
129+ auto convOutW = (xW + padL + padR - dilKsizeW) / strideW + 1;
130+ 
131+ if ((convOutH != offsetH) || (convOutW != offsetW)) {
132+ OP_LOGE(
133+ context->GetNodeName(),
134+ "Input_offsets h/w should be same as h/w after convolution, but now offset: [h:%ld, w:%ld]. conv_out: "
135+ "[h:%ld, w:%ld].",
136+ offsetH, offsetW, convOutH, convOutW);
137+ return ge::GRAPH_FAILED;
138+ }
139+ 
140+ auto outputShape = context->GetOutputShape(0);
141+ OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);
142+ *outputShape = *xShape;
143+ outputShape->SetDim(posH, offsetH * ksizeH);
144+ outputShape->SetDim(posW, offsetW * ksizeW);
145+ OP_LOGD(
146+ context->GetNodeName(), "x shape is %s, offset shape is %s, output shape is %s, dataFormat is %s",
147+ Ops::Base::ToString(*xShape).c_str(), Ops::Base::ToString(*offsetShape).c_str(),
148+ Ops::Base::ToString(*outputShape).c_str(), dataFormat);
149+ return ge::GRAPH_SUCCESS;
150+}
151+ 
152+static ge::graphStatus DeformableOffsetsInferDataType(gert::InferDataTypeContext* context)
153+{
154+ OP_LOGD(context->GetNodeName(), "DeformableOffsetsInferDataType begin");
155+ auto inputXDtype = context->GetInputDataType(0);
156+ context->SetOutputDataType(0, inputXDtype);
157+ OP_LOGD(context->GetNodeName(), "DeformableOffsetsInferDataType end");
158+ return ge::GRAPH_SUCCESS;
159+}
160+ 
161+IMPL_OP_INFERSHAPE(DeformableOffsets)
162+ .InferShape(DeformableOffsetsInferShape)
163+ .InferDataType(DeformableOffsetsInferDataType);
164+} // namespace ops
Aconv/deformable_offsets/op_kernel/arch35/deformable_offsets.h+201-0
@@ -0,0 +1,201 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file deformable_offsets.h
13+ * \brief deformable_offsets kernel info
14+ */
15+#ifndef DEFORMABLE_OFFSET_H
16+#define DEFORMABLE_OFFSET_H
17+#include "kernel_operator.h"
18+namespace DeformableOffsets {
19+using namespace AscendC;
20+const uint32_t WIDTH_OFFSET_INDEX = 0;
21+const uint32_t HEIGHT_OFFSET_INDEX = 1;
22+const uint32_t POINT_WEIGHT_OFFSET_INDEX = 2;
23+const uint32_t VF_MAX_THREAD_NUM = 512;
24+const uint32_t OFFSET_DIM_VALUE = 3;
25+template <typename T>
26+class DeformableOffset {
27+public:
28+ __aicore__ inline DeformableOffset()
29+ {}
30+ __aicore__ inline void Init(
31+ GM_ADDR x, GM_ADDR offsets, GM_ADDR y, GM_ADDR workspace,
32+ const DeformableOffsetsTilingDataSimt* __restrict tilingData);
33+ __aicore__ inline void Process();
34+ 
35+private:
36+ GlobalTensor<T> inputImgGm_;
37+ GlobalTensor<T> offsetsGm_;
38+ GlobalTensor<T> yGm_;
39+ uint32_t blockId_ = GetBlockIdx();
40+ const DeformableOffsetsTilingDataSimt* tiling_;
41+};
42+ 
43+template <typename T>
44+__aicore__ inline void DeformableOffset<T>::Init(
45+ GM_ADDR x, GM_ADDR offsets, GM_ADDR y, GM_ADDR workspace,
46+ const DeformableOffsetsTilingDataSimt* __restrict tilingData)
47+{
48+ inputImgGm_.SetGlobalBuffer((__gm__ T*)(x));
49+ offsetsGm_.SetGlobalBuffer((__gm__ T*)(offsets));
50+ yGm_.SetGlobalBuffer((__gm__ T*)(y));
51+ 
52+ tiling_ = tilingData;
53+}
54+ 
55+__aicore__ __attribute__((always_inline)) inline float GetFloorValue(float x)
56+{
57+ return __floorf(x);
58+}
59+ 
60+template <typename T>
61+__aicore__ __attribute__((always_inline)) inline T GetInputPointValue(
62+ __gm__ T* inputImgGmAddr, int32_t inputHeight, int32_t inputWidth, uint32_t channelIndex,
63+ uint32_t inputDataBatchOffset, uint32_t imgHeight, uint32_t imgWidth, uint32_t imgWidthStride, uint32_t imgChannel)
64+{
65+ if (inputHeight >= 0 && inputWidth >= 0 && inputHeight < imgHeight && inputWidth < imgWidth) {
66+ return inputImgGmAddr
67+ [inputDataBatchOffset + inputHeight * imgWidthStride + inputWidth * imgChannel + channelIndex];
68+ }
69+ return static_cast<T>(0.0);
70+}
71+ 
72+template <typename T>
73+__aicore__ __attribute__((always_inline)) inline T DeformableOffsetBilinear(
74+ __gm__ T* inputImgGmAddr, float pointHeight, float pointWidth, uint32_t channelIndex, T offsetPointWeight,
75+ uint32_t inputDataBatchOffset, uint32_t imgHeight, uint32_t imgWidth, uint32_t imgWidthStride, uint32_t imgChannel)
76+{
77+ float heightFloor = GetFloorValue(pointHeight);
78+ float widthFloor = GetFloorValue(pointWidth);
79+ 
80+ float heightFloorDelta = pointHeight - heightFloor;
81+ float widthFloorDelta = pointWidth - widthFloor;
82+ // pointLeftUp
83+ float inputValue = static_cast<float>(GetInputPointValue(
84+ (__gm__ T*)inputImgGmAddr, heightFloor, widthFloor, channelIndex, inputDataBatchOffset, imgHeight, imgWidth,
85+ imgWidthStride, imgChannel));
86+ float inputWeight = (1.0f - heightFloorDelta) * (1.0f - widthFloorDelta);
87+ float bilinearValue = (inputValue * inputWeight);
88+ 
89+ // pointRightUp
90+ inputValue = static_cast<float>(GetInputPointValue(
91+ (__gm__ T*)inputImgGmAddr, heightFloor, (widthFloor + 1), channelIndex, inputDataBatchOffset, imgHeight,
92+ imgWidth, imgWidthStride, imgChannel));
93+ inputWeight = (1.0f - heightFloorDelta) * widthFloorDelta;
94+ bilinearValue += (inputValue * inputWeight);
95+ 
96+ // pointLeftBottom
97+ inputValue = static_cast<float>(GetInputPointValue(
98+ (__gm__ T*)inputImgGmAddr, (heightFloor + 1), widthFloor, channelIndex, inputDataBatchOffset, imgHeight,
99+ imgWidth, imgWidthStride, imgChannel));
100+ inputWeight = heightFloorDelta * (1.0f - widthFloorDelta);
101+ bilinearValue += (inputValue * inputWeight);
102+ 
103+ // pointRightBottom
104+ inputValue = static_cast<float>(GetInputPointValue(
105+ (__gm__ T*)inputImgGmAddr, (heightFloor + 1), (widthFloor + 1), channelIndex, inputDataBatchOffset, imgHeight,
106+ imgWidth, imgWidthStride, imgChannel));
107+ inputWeight = heightFloorDelta * widthFloorDelta;
108+ bilinearValue += (inputValue * inputWeight);
109+ 
110+ return static_cast<T>(bilinearValue * static_cast<float>(offsetPointWeight));
111+}
112+ 
113+// LAUNCH_BOUND
114+template <typename T>
115+__simt_vf__ LAUNCH_BOUND(VF_MAX_THREAD_NUM) __aicore__ void ComputeDeformableOffset(
116+ __gm__ T* inputImgGmAddr, __gm__ T* offsetsGmAddr, __gm__ T* yGmAddr, uint32_t blockNumber, uint32_t numKernels,
117+ uint32_t imgOutWidth, uint32_t imgChannel, uint32_t imgHeight, uint32_t imgWidth, uint32_t strideH,
118+ uint32_t strideW, uint32_t dilationH, uint32_t dilationW, uint32_t padsH, uint32_t padsW, uint32_t dimKh,
119+ uint32_t dimKw, uint32_t outputPointWidthStride, uint32_t outputWidthStride, uint32_t outputKernelWidthStride,
120+ uint32_t outputBatchStride, uint32_t offsetBatchStride, uint32_t offsetKernelElementStride,
121+ uint32_t offsetPointStride, uint32_t offsetWidthStride, uint32_t imgBatchStride, uint32_t imgWidthStride,
122+ uint32_t groups, uint32_t outImgSize, uint32_t shiftB_, uint32_t mB_, uint32_t shiftH_, uint32_t mH_,
123+ uint32_t shiftW_, uint32_t mW_, uint32_t shiftC_, uint32_t mC_, uint32_t blockId_)
124+{
125+ uint32_t offsetGroupKernelStride = dimKh * dimKw;
126+ uint32_t heightOffset = HEIGHT_OFFSET_INDEX * offsetKernelElementStride;
127+ uint32_t widthOffset = WIDTH_OFFSET_INDEX * offsetKernelElementStride;
128+ uint32_t weightOffset = POINT_WEIGHT_OFFSET_INDEX * offsetKernelElementStride;
129+ 
130+ for (uint32_t index = blockId_ * VF_MAX_THREAD_NUM + Simt::GetThreadIdx(); index < numKernels;
131+ index += (blockNumber * VF_MAX_THREAD_NUM)) {
132+ // output info (N H K_h W K_w, groups, groupC)
133+ uint32_t batchNum, heightCol, widthCol, channelIndex, groupsIndex;
134+ // fast division, addr/factor
135+ batchNum = Simt::UintDiv(index, mB_, shiftB_);
136+ uint32_t remain = index - batchNum * outImgSize;
137+ 
138+ heightCol = Simt::UintDiv(remain, mH_, shiftH_);
139+ remain = remain - heightCol * (imgOutWidth * imgChannel);
140+ 
141+ widthCol = Simt::UintDiv(remain, mW_, shiftW_);
142+ channelIndex = remain - widthCol * imgChannel;
143+ 
144+ groupsIndex = Simt::UintDiv(channelIndex, mC_, shiftC_);
145+ 
146+ uint32_t newIndex = batchNum * outputBatchStride;
147+ int32_t heightInput = heightCol * strideH - padsH;
148+ int32_t widthInput = widthCol * strideW - padsW;
149+ 
150+ uint32_t outputOffset = newIndex + heightCol * outputKernelWidthStride + widthCol * outputPointWidthStride;
151+ uint32_t newOffsetIndex = batchNum * offsetBatchStride;
152+ uint32_t newInputIndex = batchNum * imgBatchStride;
153+ 
154+ uint32_t offsetBaseAdrr = newOffsetIndex + heightCol * offsetWidthStride + widthCol * offsetPointStride +
155+ groupsIndex * offsetGroupKernelStride;
156+ for (int32_t i = 0; i < dimKh; i++) {
157+ for (int32_t j = 0; j < dimKw; j++) {
158+ uint32_t offsetAdrr = offsetBaseAdrr + (i * dimKw + j);
159+ // offset height info
160+ uint32_t offsetValueIndex = offsetAdrr + heightOffset;
161+ float pointHeight = static_cast<float>(heightInput) + static_cast<float>(i * dilationH) +
162+ static_cast<float>(offsetsGmAddr[offsetValueIndex]);
163+ // offset width info
164+ offsetValueIndex = offsetAdrr + widthOffset;
165+ float pointWidth = static_cast<float>(widthInput) + static_cast<float>(j * dilationW) +
166+ static_cast<float>(offsetsGmAddr[offsetValueIndex]);
167+ // offset weight info
168+ offsetValueIndex = offsetAdrr + weightOffset;
169+ T bilinearValue = DeformableOffsetBilinear(
170+ (__gm__ T*)(inputImgGmAddr), pointHeight, pointWidth, channelIndex, offsetsGmAddr[offsetValueIndex],
171+ newInputIndex, imgHeight, imgWidth, imgWidthStride, imgChannel);
172+ // data layout (n, h, k_h, w, k_w, c)
173+ yGmAddr[outputOffset + i * outputWidthStride + j * imgChannel + channelIndex] = bilinearValue;
174+ }
175+ }
176+ }
177+}
178+ 
179+template <typename T>
180+__aicore__ inline void DeformableOffset<T>::Process()
181+{
182+ uint32_t outImgSize = tiling_->imgOutWidth * tiling_->imgOutHeight * tiling_->imgChannel;
183+ uint32_t shiftB_, mB_, shiftH_, mH_, shiftW_, mW_, shiftC_, mC_;
184+ GetUintDivMagicAndShift(mB_, shiftB_, outImgSize);
185+ GetUintDivMagicAndShift(mH_, shiftH_, tiling_->imgOutWidth * tiling_->imgChannel);
186+ GetUintDivMagicAndShift(mW_, shiftW_, tiling_->imgChannel);
187+ GetUintDivMagicAndShift(mC_, shiftC_, tiling_->imgChannel / tiling_->deformableGroups);
188+ Simt::VF_CALL<ComputeDeformableOffset<T>>(
189+ Simt::Dim3{VF_MAX_THREAD_NUM, 1, 1}, (__gm__ T*)(inputImgGm_.GetPhyAddr()),
190+ (__gm__ T*)(offsetsGm_.GetPhyAddr()), (__gm__ T*)(yGm_.GetPhyAddr()), tiling_->blockNum, tiling_->numKernels,
191+ tiling_->imgOutWidth, tiling_->imgChannel, tiling_->imgHeight, tiling_->imgWidth, tiling_->strideHeight,
192+ tiling_->strideWidth, tiling_->dilationHeight, tiling_->dilationWidth, tiling_->padsHeight, tiling_->padsWidth,
193+ tiling_->dimKHeight, tiling_->dimKWidth, tiling_->outputPointWidthStride, tiling_->outputWidthStride,
194+ tiling_->outputKernelWidthStride, tiling_->outputBatchStride, tiling_->offsetBatchStride,
195+ tiling_->offsetKernelElementStride, tiling_->offsetPointStride, tiling_->offsetWidthStride,
196+ tiling_->imgBatchStride, tiling_->imgWidthStride, tiling_->deformableGroups, outImgSize, shiftB_, mB_, shiftH_,
197+ mH_, shiftW_, mW_, shiftC_, mC_, blockId_);
198+}
199+ 
200+} // namespace DeformableOffsets
201+#endif // DEFORMABLE_OFFSETS_H
Aconv/deformable_offsets/op_kernel/deformable_offsets_apt.cpp+33-0
@@ -0,0 +1,33 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file deformable_offsets.cpp
13+ * \brief deformable_offsets kernel main
14+ */
15+ 
16+#include "arch35/deformable_offsets.h"
17+ 
18+#define TILING_SIMT_COMMON_KEY 1000
19+ 
20+using namespace DeformableOffsets;
21+using namespace AscendC;
22+ 
23+extern "C" __global__ __aicore__ void deformable_offsets(
24+ GM_ADDR x, GM_ADDR offsets, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
25+{
26+ GET_TILING_DATA(tilingData, tiling);
27+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
28+ if (TILING_KEY_IS(TILING_SIMT_COMMON_KEY)) {
29+ DeformableOffset<DTYPE_X> deformableOffsetObject;
30+ deformableOffsetObject.Init(x, offsets, y, workspace, &tilingData);
31+ deformableOffsetObject.Process();
32+ }
33+}
Aconv/deformable_offsets/tests/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aconv/deformable_offsets/tests/ut/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aconv/deformable_offsets/tests/ut/op_host/CMakeLists.txt+14-0
@@ -0,0 +1,14 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_HOST_UT)
13+ add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
Aconv/deformable_offsets/tests/ut/op_host/arch35/test_deformable_offsets_tiling.cpp+68-0
@@ -0,0 +1,68 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file test_deformable_offsets_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "../../../op_host/arch35/deformable_offsets_tiling_arch35.h"
17+#include <iostream>
18+#include <gtest/gtest.h>
19+#include "tiling_context_faker.h"
20+#include "tiling_case_executor.h"
21+ 
22+using namespace std;
23+using namespace ge;
24+ 
25+class DeformableOffsetsTiling : public testing::Test {
26+protected:
27+ static void SetUpTestCase()
28+ {
29+ std::cout << "DeformableOffsetsTiling SetUp" << std::endl;
30+ }
31+ 
32+ static void TearDownTestCase()
33+ {
34+ std::cout << "DeformableOffsetsTiling TearDown" << std::endl;
35+ }
36+};
37+ 
38+TEST_F(DeformableOffsetsTiling, deformable_offsets_test_0)
39+{
40+ optiling::TilingPrepareForDeformableOffsetsCompileInfo compileInfo = {64, 245760};
41+ gert::TilingContextPara tilingContextPara(
42+ "DeformableOffsets",
43+ {
44+ {{{1, 7, 11, 256}, {1, 7, 11, 256}}, ge::DT_FLOAT, ge::FORMAT_NHWC},
45+ {{{1, 7, 11, 27}, {1, 7, 11, 27}}, ge::DT_FLOAT, ge::FORMAT_NHWC},
46+ },
47+ {
48+ {{{1, 21, 33, 256}, {1, 21, 33, 256}}, ge::DT_FLOAT, ge::FORMAT_NHWC},
49+ },
50+ {
51+ gert::TilingContextPara::OpAttr(
52+ "strides", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({1, 1, 1, 1})),
53+ gert::TilingContextPara::OpAttr("pads", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({1, 1, 1, 1})),
54+ gert::TilingContextPara::OpAttr("ksize", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({3, 3})),
55+ gert::TilingContextPara::OpAttr(
56+ "dilations", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({1, 1, 1, 1})),
57+ gert::TilingContextPara::OpAttr("data_format", Ops::NN::AnyValue::CreateFrom<string>("NHWC")),
58+ gert::TilingContextPara::OpAttr("deformable_groups", Ops::NN::AnyValue::CreateFrom<int64_t>(1)),
59+ gert::TilingContextPara::OpAttr("modulated", Ops::NN::AnyValue::CreateFrom<bool>(true)),
60+ },
61+ &compileInfo);
62+ uint64_t expectTilingKey = 1000;
63+ string expectTilingData =
64+ "4294967335 4294967297 4294967297 12884901889 1099511627779 30064771083 30064773888 38654705675 1275605286939 "
65+ "4294967299 36283883717376 84662395364096 8929237028096 4295144704 ";
66+ std::vector<size_t> expectWorkspaces = {16777216};
67+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
68+}
Aconv/deformable_offsets/tests/ut/op_host/test_deformable_offsets_infershape.cpp+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+ 
11+/*!
12+ * \file test_deformable_offsets_infershape.cpp
13+ * \brief
14+ */
15+ 
16+#include <iostream>
17+#include <gtest/gtest.h>
18+#include "infershape_context_faker.h"
19+#include "infershape_case_executor.h"
20+ 
21+class DeformableOffsetsInfershape : public testing::Test {
22+protected:
23+ static void SetUpTestCase()
24+ {
25+ std::cout << "DeformableOffsetsInfershape SetUp" << std::endl;
26+ }
27+ 
28+ static void TearDownTestCase()
29+ {
30+ std::cout << "DeformableOffsetsInfershape TearDown" << std::endl;
31+ }
32+};
33+ 
34+TEST_F(DeformableOffsetsInfershape, deformable_offsets_infer_shape_test)
35+{
36+ gert::InfershapeContextPara infershapeContextPara(
37+ "DeformableOffsets",
38+ {
39+ {{{4, 16, 64, 64}, {4, 16, 64, 64}}, ge::DT_FLOAT16, ge::FORMAT_NCHW},
40+ {{{4, 216, 64, 64}, {4, 216, 64, 64}}, ge::DT_FLOAT16, ge::FORMAT_NCHW},
41+ },
42+ {
43+ {{{4, 32, 192, 192}, {4, 32, 192, 192}}, ge::DT_FLOAT16, ge::FORMAT_NCHW},
44+ },
45+ {
46+ {"strides", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({1, 1, 1, 1})},
47+ {"pads", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({1, 1, 1, 1})},
48+ {"ksize", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({3, 3})},
49+ {"dilations", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({1, 1, 1, 1})},
50+ {"data_format", Ops::NN::AnyValue::CreateFrom<std::string>("NCHW")},
51+ {"deformable_groups", Ops::NN::AnyValue::CreateFrom<int64_t>(8)},
52+ {"modulated", Ops::NN::AnyValue::CreateFrom<bool>(true)},
53+ });
54+ std::vector<std::vector<int64_t>> expectOutputShape = {
55+ {4, 16, 192, 192},
56+ };
57+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
58+}
Mdocs/zh/op_list.md+10-0
@@ -764,6 +764,16 @@
764 <td>AI Core</td>764 <td>AI Core</td>
765 <td>实现卷积功能，支持2D卷积，同时支持可变形卷积、分组卷积。</td>765 <td>实现卷积功能，支持2D卷积，同时支持可变形卷积、分组卷积。</td>
766 </tr>766 </tr>
767+ <tr>
768+ <td>conv</td>
769+ <td><a href="../../conv/deformable_offsets/README.md">deformable_offsets</a></td>
770+ <td>✓</td>
771+ <td>✓</td>
772+ <td>✓</td>
773+ <td>✓</td>
774+ <td>AI Core</td>
775+ <td>用于计算变形卷积（Deformable Convolution）输出的函数。通过引入偏移参数offsets，使得卷积核在输入特征图上的位置可以动态调整，从而适配不规则的集合变化。</td>
776+ </tr>
767 <tr>777 <tr>
768 <td>foreach</td>778 <td>foreach</td>
769 <td><a href="../../foreach/foreach_abs/README.md">foreach_abs</a></td>779 <td><a href="../../foreach/foreach_abs/README.md">foreach_abs</a></td>
Atests/ut/common/infershape_case_executor.cpp+111-0
@@ -0,0 +1,111 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "infershape_case_executor.h"
12+#include <gtest/gtest.h>
13+#include "base/registry/op_impl_space_registry_v2.h"
14+ 
15+#define DO_INFERSHAPE(infershapeContextPara) \
16+ auto contextFaker = gert::InferShapeContextFaker(); \
17+ /* 1. input/output information */ \
18+ size_t inputNum = infershapeContextPara.inputTensorDesc_.size(); \
19+ size_t outputNum = infershapeContextPara.outputTensorDesc_.size(); \
20+ if (infershapeContextPara.inputInstanceNum_.size() != 0 || infershapeContextPara.outputInstanceNum_.size() != 0) { \
21+ contextFaker.IrInstanceNum(infershapeContextPara.inputInstanceNum_, infershapeContextPara.outputInstanceNum_); \
22+ } else { \
23+ contextFaker.NodeIoNum(inputNum, outputNum); \
24+ } \
25+ std::vector<gert::Tensor *> inputTensors = {}; \
26+ std::vector<std::unique_ptr<gert::Tensor>> inputTensorsKeepAlive = {}; \
27+ for (size_t index = 0; index < inputNum; index++) { \
28+ std::unique_ptr<gert::Tensor> curTensor = std::make_unique<gert::Tensor>( \
29+ infershapeContextPara.inputTensorDesc_[index].shape_, \
30+ gert::StorageFormat(infershapeContextPara.inputTensorDesc_[index].format_, \
31+ infershapeContextPara.inputTensorDesc_[index].format_, \
32+ gert::ExpandDimsType()), \
33+ gert::TensorPlacement::kOnHost, \
34+ infershapeContextPara.inputTensorDesc_[index].dtype_, \
35+ infershapeContextPara.inputTensorDesc_[index].isConst_ ? \
36+ infershapeContextPara.inputTensorDesc_[index].constValue_: \
37+ nullptr); \
38+ inputTensors.push_back(curTensor.get()); \
39+ inputTensorsKeepAlive.push_back(std::move(curTensor)); \
40+ } \
41+ for (size_t index = 0; index < outputNum; index++) { \
42+ contextFaker.NodeOutputTd(index, \
43+ infershapeContextPara.outputTensorDesc_[index].dtype_, \
44+ infershapeContextPara.outputTensorDesc_[index].format_, \
45+ infershapeContextPara.outputTensorDesc_[index].format_); \
46+ } \
47+ contextFaker.InputTensors(inputTensors); \
48+ for (auto& attrInfo : infershapeContextPara.attrs_) { \
49+ switch (attrInfo.attr_.type_) { \
50+ case Ops::NN::AnyValue::ValueType::VT_BOOL: { \
51+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<bool*>(attrInfo.attr_.valuePtr_.get())); \
52+ break;} \
53+ case Ops::NN::AnyValue::ValueType::VT_INT: { \
54+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<int64_t*>(attrInfo.attr_.valuePtr_.get())); \
55+ break;} \
56+ case Ops::NN::AnyValue::ValueType::VT_FLOAT: { \
57+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<float*>(attrInfo.attr_.valuePtr_.get())); \
58+ break;} \
59+ case Ops::NN::AnyValue::ValueType::VT_STRING: { \
60+ contextFaker.Attr(attrInfo.attrName_, ge::AscendString(reinterpret_cast<std::string*>(attrInfo.attr_.valuePtr_.get())->c_str()));\
61+ break;} \
62+ case Ops::NN::AnyValue::ValueType::VT_LIST_BOOL: { \
63+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<std::vector<bool>*>(attrInfo.attr_.valuePtr_.get()));\
64+ break;} \
65+ case Ops::NN::AnyValue::ValueType::VT_LIST_INT: { \
66+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<std::vector<int64_t>*>(attrInfo.attr_.valuePtr_.get()));\
67+ break;} \
68+ case Ops::NN::AnyValue::ValueType::VT_LIST_LIST_INT: { \
69+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<std::vector<std::vector<int64_t>>*>(attrInfo.attr_.valuePtr_.get()));\
70+ break;} \
71+ case Ops::NN::AnyValue::ValueType::VT_LIST_FLOAT: { \
72+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<std::vector<float>*>(attrInfo.attr_.valuePtr_.get()));\
73+ break;} \
74+ default: \
75+ std::cout << "[ERROR]" << __FILE__ << ":" << __LINE__ << "The ValueType " << attrInfo.attr_.type_ << "is not supported!" << std::endl;\
76+ } \
77+ } \
78+ auto contextHolder = contextFaker.SetOpType(infershapeContextPara.opName_.c_str()).Build(); \
79+ /* 2. get infershape func */ \
80+ auto spaceRegistry = gert::DefaultOpImplSpaceRegistryV2::GetInstance().GetSpaceRegistry(); \
81+ auto infershapeFunc = spaceRegistry->GetOpImpl(infershapeContextPara.opName_.c_str())->infer_shape; \
82+ /* 3. check infershape func */ \
83+ auto infershapeRet = infershapeFunc(contextHolder.GetContext());
84+ 
85+static std::vector<int64_t> ToVector(const gert::Shape& shape) {
86+ size_t shapeSize = shape.GetDimNum();
87+ std::vector<int64_t> shapeVec(shapeSize, 0);
88+ 
89+ for (size_t i = 0; i < shapeSize; i++) {
90+ shapeVec[i] = shape.GetDim(i);
91+ }
92+ return shapeVec;
93+}
94+ 
95+void ExecuteTestCase(gert::InfershapeContextPara& infershapeContextPara,
96+ ge::graphStatus expectResult,
97+ const std::vector<std::vector<int64_t>>& expectOutputShape)
98+{
99+ DO_INFERSHAPE(infershapeContextPara);
100+ 
101+ // check infershape func
102+ EXPECT_EQ(infershapeRet, expectResult);
103+ if (expectResult == ge::GRAPH_FAILED) {
104+ return;
105+ }
106+ 
107+ // check output shape
108+ for (int i = 0; i < expectOutputShape.size(); i++) {
109+ EXPECT_EQ(ToVector(*contextHolder.GetContext()->GetOutputShape(i)), expectOutputShape[i]);
110+ }
111+}
Atests/ut/common/infershape_case_executor.h+20-0
@@ -0,0 +1,20 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OPS_NN_TESTS_UT_COMMON_INFERSHAPE_CASE_EXECUTOR_H
12+#define OPS_NN_TESTS_UT_COMMON_INFERSHAPE_CASE_EXECUTOR_H
13+ 
14+#include "infershape_context_faker.h"
15+ 
16+void ExecuteTestCase(gert::InfershapeContextPara& infershapeContextPara,
17+ ge::graphStatus expectResult = ge::GRAPH_FAILED,
18+ const std::vector<std::vector<int64_t>>& expectOutputShape = {});
19+ 
20+#endif // OPS_NN_TESTS_UT_COMMON_INFERSHAPE_CASE_EXECUTOR_H
Atests/ut/common/infershape_context_faker.cpp+63-0
@@ -0,0 +1,63 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "infershape_context_faker.h"
12+ 
13+namespace gert {
14+ 
15+InferShapeContextFaker& InferShapeContextFaker::SetOpType(const std::string opType)
16+{
17+ OpInferShapeContextBuilder::OpType(opType.c_str()).OpName(opType.c_str());
18+ return *this;
19+}
20+ 
21+InferShapeContextFaker& InferShapeContextFaker::NodeIoNum(size_t inputNum, size_t outputNum)
22+{
23+ OpInferShapeContextBuilder::IONum(inputNum, outputNum);
24+ return *this;
25+}
26+ 
27+InferShapeContextFaker& InferShapeContextFaker::IrInstanceNum(const std::vector<uint32_t>& inputInstanceNum,
28+ const std::vector<uint32_t>& outputInstanceNum)
29+{
30+ OpInferShapeContextBuilder::IOInstanceNum(inputInstanceNum, outputInstanceNum);
31+ return *this;
32+}
33+ 
34+InferShapeContextFaker& InferShapeContextFaker::NodeInputTd(int32_t index, ge::DataType dtype, ge::Format originFormat,
35+ ge::Format storageFormat)
36+{
37+ return *this;
38+}
39+ 
40+InferShapeContextFaker& InferShapeContextFaker::NodeOutputTd(int32_t index, ge::DataType dtype, ge::Format originFormat,
41+ ge::Format storageFormat)
42+{
43+ OpInferShapeContextBuilder::OutputTensorDesc(index, dtype, originFormat, storageFormat);
44+ return *this;
45+}
46+ 
47+InferShapeContextFaker& InferShapeContextFaker::InputTensors(const std::vector<Tensor *>& inputTensors)
48+{
49+ OpInferShapeContextBuilder::InputTensors(inputTensors);
50+ return *this;
51+}
52+ 
53+InferShapeContextFaker& InferShapeContextFaker::OutputShapes(const std::vector<StorageShape *>& outputShapes)
54+{
55+ return *this;
56+}
57+ 
58+ContextHolder<InferShapeContext> InferShapeContextFaker::Build()
59+{
60+ return OpInferShapeContextBuilder::Build();
61+}
62+ 
63+} // namespace gert
Atests/ut/common/infershape_context_faker.h+138-0
@@ -0,0 +1,138 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OPS_NN_TESTS_UT_COMMON_INFERSHAPE_CONTEXT_FAKER_H
12+#define OPS_NN_TESTS_UT_COMMON_INFERSHAPE_CONTEXT_FAKER_H
13+ 
14+#include "op_infer_shape_context_builder.h"
15+#include "any_value.h"
16+ 
17+namespace gert {
18+ 
19+ class InfershapeContextPara {
20+public:
21+ class TensorDescription {
22+ public:
23+ TensorDescription(const gert::StorageShape& shape, ge::DataType dtype, ge::Format format, bool isConst = false,
24+ void* constValue = nullptr) :
25+ shape_(shape), dtype_(dtype), format_(format), isConst_(isConst), constValue_(constValue) {}
26+ public:
27+ gert::StorageShape shape_;
28+ ge::DataType dtype_ = ge::DT_FLOAT;
29+ ge::Format format_ = ge::FORMAT_ND;
30+ bool isConst_ = false;
31+ void* constValue_ = nullptr;
32+ };
33+ 
34+ class OpAttr {
35+ public:
36+ OpAttr(const std::string& attrName, const Ops::NN::AnyValue& attr) : attrName_(attrName), attr_(attr) {}
37+ public:
38+ std::string attrName_;
39+ Ops::NN::AnyValue attr_;
40+ };
41+public:
42+ InfershapeContextPara(const std::string& opName,
43+ const std::vector<TensorDescription>& inputTensorDesc,
44+ const std::vector<TensorDescription>& outputTensorDesc,
45+ const std::vector<OpAttr>& attrs,
46+ const std::vector<uint32_t>& inputInstanceNum = {},
47+ const std::vector<uint32_t>& outputInstanceNum = {}) :
48+ opName_(opName),
49+ inputTensorDesc_(inputTensorDesc),
50+ outputTensorDesc_(outputTensorDesc),
51+ attrs_(attrs),
52+ inputInstanceNum_(inputInstanceNum),
53+ outputInstanceNum_(outputInstanceNum) {}
54+ 
55+ InfershapeContextPara(const std::string& opName,
56+ const std::vector<TensorDescription>& inputTensorDesc,
57+ const std::vector<TensorDescription>& outputTensorDesc,
58+ const std::vector<uint32_t>& inputInstanceNum = {},
59+ const std::vector<uint32_t>& outputInstanceNum = {}) :
60+ opName_(opName),
61+ inputTensorDesc_(inputTensorDesc),
62+ outputTensorDesc_(outputTensorDesc),
63+ inputInstanceNum_(inputInstanceNum),
64+ outputInstanceNum_(outputInstanceNum) {}
65+ 
66+public:
67+ std::string opName_;
68+ std::vector<uint32_t> inputInstanceNum_;
69+ std::vector<uint32_t> outputInstanceNum_;
70+ std::vector<TensorDescription> inputTensorDesc_;
71+ std::vector<TensorDescription> outputTensorDesc_;
72+ std::vector<OpAttr> attrs_;
73+};
74+ 
75+class InferShapeContextFaker : public OpInferShapeContextBuilder {
76+public:
77+ InferShapeContextFaker& SetOpType(const std::string opType);
78+ 
79+ /* only one can be choosed from IrInstanceNum */
80+ InferShapeContextFaker& NodeIoNum(size_t inputNum, size_t outputNum);
81+ 
82+ /* can be used for dynamic inputs/outputs
83+ * only one can be choosed from NodeIoNum */
84+ InferShapeContextFaker& IrInstanceNum(const std::vector<uint32_t>& inputInstanceNum,
85+ const std::vector<uint32_t>& outputInstanceNum);
86+ 
87+ InferShapeContextFaker& NodeInputTd(int32_t index, ge::DataType dtype, ge::Format originFormat,
88+ ge::Format storageFormat);
89+ 
90+ InferShapeContextFaker& NodeOutputTd(int32_t index, ge::DataType dtype, ge::Format originFormat,
91+ ge::Format storageFormat);
92+ 
93+ InferShapeContextFaker& Attr(const std::string& attrName, bool attr) {
94+ OpInferShapeContextBuilder::AppendAttr(attr);
95+ return *this;
96+ }
97+ InferShapeContextFaker& Attr(const std::string& attrName, int64_t attr) {
98+ OpInferShapeContextBuilder::AppendAttr(attr);
99+ return *this;
100+ }
101+ InferShapeContextFaker& Attr(const std::string& attrName, float attr) {
102+ OpInferShapeContextBuilder::AppendAttr(attr);
103+ return *this;
104+ }
105+ InferShapeContextFaker& Attr(const std::string& attrName, const ge::AscendString& attr) {
106+ OpInferShapeContextBuilder::AppendAttr(attr);
107+ return *this;
108+ }
109+ InferShapeContextFaker& Attr(const std::string& attrName, const std::vector<bool>& attr) {
110+ OpInferShapeContextBuilder::AppendAttr(attr);
111+ return *this;
112+ }
113+ InferShapeContextFaker& Attr(const std::string& attrName, const std::vector<int64_t>& attr) {
114+ OpInferShapeContextBuilder::AppendAttr(attr);
115+ return *this;
116+ }
117+ InferShapeContextFaker& Attr(const std::string& attrName, const std::vector<float>& attr) {
118+ OpInferShapeContextBuilder::AppendAttr(attr);
119+ return *this;
120+ }
121+ InferShapeContextFaker& Attr(const std::string& attrName, const std::vector<ge::AscendString>& attr) {
122+ OpInferShapeContextBuilder::AppendAttr(attr);
123+ return *this;
124+ }
125+ InferShapeContextFaker& Attr(const std::string& attrName, const std::vector<std::vector<int64_t>>& attr) {
126+ OpInferShapeContextBuilder::AppendAttr(attr);
127+ return *this;
128+ }
129+ 
130+ InferShapeContextFaker& InputTensors(const std::vector<Tensor *>& inputTensors);
131+ 
132+ InferShapeContextFaker& OutputShapes(const std::vector<StorageShape *>& outputShapes);
133+ 
134+ ContextHolder<InferShapeContext> Build();
135+};
136+ 
137+} // namespace gert
138+#endif // OPS_NN_TESTS_UT_COMMON_INFERSHAPE_CONTEXT_FAKER_H
Atests/ut/common/tiling_case_executor.cpp+267-0
@@ -0,0 +1,267 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "tiling_case_executor.h"
12+#include <gtest/gtest.h>
13+#include <nlohmann/json.hpp>
14+#include "platform/platform_infos_def.h"
15+#include "base/registry/op_impl_space_registry_v2.h"
16+ 
17+#define DO_TILING(tilingContextPara) \
18+ auto contextFaker = gert::TilingContextFaker(); \
19+ /* 1. input/output information */ \
20+ size_t inputNum = tilingContextPara.inputTensorDesc_.size(); \
21+ size_t outputNum = tilingContextPara.outputTensorDesc_.size(); \
22+ if (tilingContextPara.inputInstanceNum_.size() != 0 || tilingContextPara.outputInstanceNum_.size() != 0) { \
23+ contextFaker.IrInstanceNum(tilingContextPara.inputInstanceNum_, tilingContextPara.outputInstanceNum_); \
24+ } else { \
25+ contextFaker.NodeIoNum(inputNum, outputNum); \
26+ } \
27+ std::vector<gert::Tensor *> inputTensors = {}; \
28+ std::vector<gert::Tensor *> outputTensors = {}; \
29+ std::vector<std::unique_ptr<gert::Tensor>> inputTensorsKeepAlive = {}; \
30+ std::vector<std::unique_ptr<gert::Tensor>> outputTensorsKeepAlive = {}; \
31+ for (size_t index = 0; index < inputNum; index++) { \
32+ std::unique_ptr<gert::Tensor> curTensor = std::make_unique<gert::Tensor>( \
33+ tilingContextPara.inputTensorDesc_[index].shape_, \
34+ gert::StorageFormat(tilingContextPara.inputTensorDesc_[index].format_, \
35+ tilingContextPara.inputTensorDesc_[index].format_, \
36+ gert::ExpandDimsType()), \
37+ gert::TensorPlacement::kOnHost, \
38+ tilingContextPara.inputTensorDesc_[index].dtype_, \
39+ tilingContextPara.inputTensorDesc_[index].isConst_ ? \
40+ tilingContextPara.inputTensorDesc_[index].constValue_: \
41+ nullptr); \
42+ inputTensors.push_back(curTensor.get()); \
43+ inputTensorsKeepAlive.push_back(std::move(curTensor)); \
44+ } \
45+ for (size_t index = 0; index < outputNum; index++) { \
46+ std::unique_ptr<gert::Tensor> curTensor = std::make_unique<gert::Tensor>( \
47+ tilingContextPara.outputTensorDesc_[index].shape_, \
48+ gert::StorageFormat(tilingContextPara.outputTensorDesc_[index].format_, \
49+ tilingContextPara.outputTensorDesc_[index].format_, \
50+ gert::ExpandDimsType()), \
51+ gert::TensorPlacement::kOnHost, \
52+ tilingContextPara.outputTensorDesc_[index].dtype_, \
53+ tilingContextPara.outputTensorDesc_[index].isConst_ ? \
54+ tilingContextPara.outputTensorDesc_[index].constValue_: \
55+ nullptr); \
56+ outputTensors.push_back(curTensor.get()); \
57+ outputTensorsKeepAlive.push_back(std::move(curTensor)); \
58+ } \
59+ contextFaker.InputTensors(inputTensors).OutputTensors(outputTensors); \
60+ for (auto& attrInfo : tilingContextPara.attrs_) { \
61+ switch (attrInfo.attr_.type_) { \
62+ case Ops::NN::AnyValue::ValueType::VT_BOOL: { \
63+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<bool*>(attrInfo.attr_.valuePtr_.get())); \
64+ break;} \
65+ case Ops::NN::AnyValue::ValueType::VT_INT: { \
66+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<int64_t*>(attrInfo.attr_.valuePtr_.get())); \
67+ break;} \
68+ case Ops::NN::AnyValue::ValueType::VT_FLOAT: { \
69+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<float*>(attrInfo.attr_.valuePtr_.get())); \
70+ break;} \
71+ case Ops::NN::AnyValue::ValueType::VT_STRING: { \
72+ contextFaker.Attr(attrInfo.attrName_, ge::AscendString(reinterpret_cast<std::string*>(attrInfo.attr_.valuePtr_.get())->c_str()));\
73+ break;} \
74+ case Ops::NN::AnyValue::ValueType::VT_LIST_BOOL: { \
75+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<std::vector<bool>*>(attrInfo.attr_.valuePtr_.get()));\
76+ break;} \
77+ case Ops::NN::AnyValue::ValueType::VT_LIST_INT: { \
78+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<std::vector<int64_t>*>(attrInfo.attr_.valuePtr_.get()));\
79+ break;} \
80+ case Ops::NN::AnyValue::ValueType::VT_LIST_LIST_INT: { \
81+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<std::vector<std::vector<int64_t>>*>(attrInfo.attr_.valuePtr_.get()));\
82+ break;} \
83+ case Ops::NN::AnyValue::ValueType::VT_LIST_FLOAT: { \
84+ contextFaker.Attr(attrInfo.attrName_, *reinterpret_cast<std::vector<float>*>(attrInfo.attr_.valuePtr_.get()));\
85+ break;} \
86+ default: \
87+ std::cout << "[ERROR]" << __FILE__ << ":" << __LINE__ << "The ValueType " << attrInfo.attr_.type_ << "is not supported!" << std::endl;\
88+ } \
89+ } \
90+ /* 2. base information */ \
91+ fe::PlatFormInfos platformInfo; \
92+ platformInfo.Init(); \
93+ auto tilingData = gert::TilingData::CreateCap(tilingContextPara.tilingDataSize_); \
94+ auto workspace = gert::ContinuousVector::Create<size_t>(4096); \
95+ auto contextHolder = contextFaker.SetOpType(tilingContextPara.opName_.c_str()) \
96+ .CompileInfo(tilingContextPara.compileInfo_) \
97+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo)) \
98+ .TilingData(tilingData.get()) \
99+ .Workspace(reinterpret_cast<gert::ContinuousVector *>(workspace.get())) \
100+ .Build(); \
101+ string compileInfoStringPrefix = R"({"hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1", "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false, "UB_SIZE": )";\
102+ string compileInfoStringMiddle = R"(, "L2_SIZE": 33554432, "L1_SIZE": 524288, "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": )";\
103+ string compileInfoStringSuffix = R"(} })";\
104+ string compileInfoString = compileInfoStringPrefix + \
105+ std::to_string(tilingContextPara.ubSize_) + \
106+ compileInfoStringMiddle + \
107+ std::to_string(tilingContextPara.coreNum_) + \
108+ compileInfoStringSuffix; \
109+ map<string, string> socToArch = { \
110+ {"Ascend310P", "2002"}, \
111+ {"Ascend910B", "2201"}, \
112+ {"Ascend910_93", "2201"}, \
113+ {"Ascend950", "3510"}, \
114+ {"Ascend910", "1001"} \
115+ }; \
116+ map<string, string> socInfos; \
117+ map<string, string> aicoreSpec; \
118+ map<string, string> intrinsics; \
119+ map<string, string> socversions = { \
120+ {"NpuArch", socToArch[tilingContextPara.socVersion_]}, \
121+ {"Short_SoC_version", tilingContextPara.socVersion_} \
122+ }; \
123+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics); \
124+ auto tilingContext = contextHolder.GetContext(); \
125+ tilingContext->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos); \
126+ tilingContext->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec); \
127+ tilingContext->GetPlatformInfo()->SetCoreNumByCoreType("AICore"); \
128+ tilingContext->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics); \
129+ tilingContext->GetPlatformInfo()->SetPlatformRes("version", socversions); \
130+ /* 3. get tiling func */ \
131+ auto spaceRegistry = gert::DefaultOpImplSpaceRegistryV2::GetInstance().GetSpaceRegistry(); \
132+ auto tilingFunc = spaceRegistry->GetOpImpl(tilingContextPara.opName_.c_str())->tiling; \
133+ /* 4. check tiling func */ \
134+ auto tilingRet = tilingFunc(tilingContext);
135+ 
136+template <typename T>
137+static string to_string(void* buf, size_t size) {
138+ string result;
139+ const T* data = reinterpret_cast<const T*>(buf);
140+ size_t len = size / sizeof(T);
141+ for (size_t i = 0; i < len; i++) {
142+ result += std::to_string(data[i]);
143+ result += " ";
144+ }
145+ return result;
146+}
147+ 
148+static void GetPlatFormInfos(const char* compileInfoStr, map<string, string>& socInfos, map<string, string>& aicoreSpec,
149+ map<string, string>& intrinsics) {
150+ string default_hardward_info = R"({
151+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1", "Intrinsic_fix_pipe_l0c2out": false,
152+ "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true,
153+ "Intrinsic_data_move_out2l1_nd2nz": false, "UB_SIZE": 262144, "L2_SIZE": 33554432,
154+ "L1_SIZE": 1048576, "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 262144,
155+ "CORE_NUM": 32}})";
156+ nlohmann::json compileInfoJson = nlohmann::json::parse(compileInfoStr);
157+ if (compileInfoJson.type() != nlohmann::json::value_t::object) {
158+ compileInfoJson = nlohmann::json::parse(default_hardward_info.c_str());
159+ }
160+ 
161+ map<string, string> socInfoKeys = {{"ai_core_cnt", "CORE_NUM"},
162+ {"l2_size", "L2_SIZE"},
163+ {"cube_core_cnt", "cube_core_cnt"},
164+ {"vector_core_cnt", "vector_core_cnt"},
165+ {"core_type_list", "core_type_list"}};
166+ socInfos["core_type_list"] = "AICore";
167+ 
168+ for (auto &t : socInfoKeys) {
169+ if (compileInfoJson.contains("hardware_info") && compileInfoJson["hardware_info"].contains(t.second)) {
170+ auto &objJson = compileInfoJson["hardware_info"][t.second];
171+ if (objJson.is_number_integer()) {
172+ socInfos[t.first] = to_string(compileInfoJson["hardware_info"][t.second].get<uint32_t>());
173+ } else if (objJson.is_string()) {
174+ socInfos[t.first] = objJson;
175+ }
176+ }
177+ }
178+ map<string, string> aicoreSpecKeys = {{"ub_size", "UB_SIZE"},
179+ {"l0_a_size", "L0A_SIZE"},
180+ {"l0_b_size", "L0B_SIZE"},
181+ {"l0_c_size", "L0C_SIZE"},
182+ {"l1_size", "L1_SIZE"},
183+ {"bt_size", "BT_SIZE"},
184+ {"load3d_constraints", "load3d_constraints"}};
185+ aicoreSpec["cube_freq"] = "cube_freq";
186+ for (auto &t : aicoreSpecKeys) {
187+ if (compileInfoJson.contains("hardware_info") && compileInfoJson["hardware_info"].contains(t.second)) {
188+ if (t.second == "load3d_constraints") {
189+ aicoreSpec[t.first] = compileInfoJson["hardware_info"][t.second].get<string>();
190+ } else {
191+ aicoreSpec[t.first] = to_string(compileInfoJson["hardware_info"][t.second].get<uint32_t>());
192+ }
193+ }
194+ }
195+ 
196+ std::string intrinsicsKeys[] = {"Intrinsic_data_move_l12ub", "Intrinsic_data_move_l0c2ub",
197+ "Intrinsic_fix_pipe_l0c2out", "Intrinsic_data_move_out2l1_nd2nz",
198+ "Intrinsic_matmul_ub_to_ub", "Intrinsic_conv_ub_to_ub",
199+ "Intrinsic_data_move_l12bt"};
200+ for (string key : intrinsicsKeys) {
201+ if (compileInfoJson.contains("hardware_info") && compileInfoJson["hardware_info"].contains(key) &&
202+ compileInfoJson["hardware_info"][key].get<bool>()) {
203+ intrinsics[key] = "float16";
204+ if (key.find("Intrinsic_data_move_l12bt") != string::npos) {
205+ intrinsics[key] = "bf16";
206+ }
207+ }
208+ }
209+}
210+ 
211+void ExecuteTestCase(const gert::TilingContextPara& tilingContextPara,
212+ ge::graphStatus expectResult,
213+ uint64_t expectTilingKey,
214+ const string& expectTilingData,
215+ const std::vector<size_t>& expectWorkspaces)
216+{
217+ DO_TILING(tilingContextPara);
218+ 
219+ // check tiling func
220+ EXPECT_EQ(tilingRet, expectResult);
221+ if (expectResult == ge::GRAPH_FAILED) {
222+ return;
223+ }
224+ 
225+ // check workspace
226+ size_t workspaceCount = tilingContext->GetWorkspaceNum();
227+ if (workspaceCount > 0) {
228+ auto workspaceSizes = tilingContext->GetWorkspaceSizes(workspaceCount);
229+ for (size_t i = 0; i < workspaceCount; i++) {
230+ ASSERT_EQ(workspaceSizes[i], expectWorkspaces[i]);
231+ }
232+ }
233+ 
234+ // check tiling key
235+ auto tilingKeyResult = tilingContext->GetTilingKey();
236+ ASSERT_EQ(tilingKeyResult, expectTilingKey);
237+ 
238+ // check tiling data
239+ auto rawTilingData = tilingContext->GetRawTilingData();
240+ auto tilingDataResult = to_string<int64_t>(rawTilingData->GetData(), rawTilingData->GetDataSize());
241+ EXPECT_EQ(tilingDataResult, expectTilingData);
242+}
243+ 
244+bool ExecuteTiling(const gert::TilingContextPara& tilingContextPara, TilingInfo& tilingInfo)
245+{
246+ DO_TILING(tilingContextPara);
247+ 
248+ if (tilingRet != ge::GRAPH_SUCCESS) {
249+ return false;
250+ }
251+ 
252+ tilingInfo.tilingKey = tilingContext->GetTilingKey();
253+ tilingInfo.blockNum = tilingContext->GetBlockDim();
254+ size_t workspaceCount = tilingContext->GetWorkspaceNum();
255+ if (workspaceCount > 0) {
256+ auto workSpaceSizes = tilingContext->GetWorkspaceSizes(workspaceCount);
257+ for (size_t i = 0; i < workspaceCount; i++) {
258+ tilingInfo.workspaceSizes.push_back(workSpaceSizes[i]);
259+ }
260+ }
261+ auto rawTilingData = tilingContext->GetRawTilingData();
262+ tilingInfo.tilingData = std::make_unique<uint8_t[]>(rawTilingData->GetDataSize());
263+ tilingInfo.tilingDataSize = rawTilingData->GetDataSize();
264+ std::memcpy(tilingInfo.tilingData.get(), rawTilingData->GetData(), rawTilingData->GetDataSize());
265+ 
266+ return true;
267+}
Atests/ut/common/tiling_case_executor.h+34-0
@@ -0,0 +1,34 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OPS_NN_TESTS_UT_COMMON_TILING_CASE_EXECUTOR_H
12+#define OPS_NN_TESTS_UT_COMMON_TILING_CASE_EXECUTOR_H
13+ 
14+#include "tiling_context_faker.h"
15+ 
16+using namespace std;
17+ 
18+struct TilingInfo {
19+ int64_t tilingKey = -1;
20+ std::vector<int64_t> workspaceSizes;
21+ std::unique_ptr<uint8_t[]> tilingData;
22+ size_t tilingDataSize = 0;
23+ size_t blockNum = 0;
24+};
25+ 
26+void ExecuteTestCase(const gert::TilingContextPara& tilingContextPara,
27+ ge::graphStatus expectResult = ge::GRAPH_FAILED,
28+ uint64_t expectTilingKey = 0,
29+ const string& expectTilingData = "",
30+ const std::vector<size_t>& expectWorkspaces = {});
31+ 
32+bool ExecuteTiling(const gert::TilingContextPara& tilingContextPara, TilingInfo& tilingInfo);
33+ 
34+#endif // OPS_NN_TESTS_UT_COMMON_TILING_CASE_EXECUTOR_H
Mtests/ut/common/tiling_context_faker.h+120-0
@@ -18,6 +18,126 @@
18#include "any_value.h"18#include "any_value.h"
19 19 
20namespace gert {20namespace gert {
21+ 
22+class TilingContextPara {
23+public:
24+ class TensorDescription {
25+ public:
26+ TensorDescription(const gert::StorageShape& shape,
27+ ge::DataType dtype,
28+ ge::Format format,
29+ bool isConst = false,
30+ void* constValue = nullptr) :
31+ shape_(shape), dtype_(dtype), format_(format), isConst_(isConst), constValue_(constValue) {}
32+ public:
33+ gert::StorageShape shape_;
34+ ge::DataType dtype_ = ge::DT_FLOAT;
35+ ge::Format format_ = ge::FORMAT_ND;
36+ bool isConst_ = false;
37+ void* constValue_ = nullptr;
38+ };
39+ 
40+ class OpAttr {
41+ public:
42+ OpAttr(const std::string& attrName, const Ops::NN::AnyValue& attr) : attrName_(attrName), attr_(attr) {}
43+ public:
44+ std::string attrName_;
45+ Ops::NN::AnyValue attr_;
46+ };
47+public:
48+ TilingContextPara(const std::string& opName,
49+ const std::vector<TensorDescription>& inputTensorDesc,
50+ const std::vector<TensorDescription>& outputTensorDesc,
51+ const std::vector<OpAttr>& attrs,
52+ void* compileInfo = nullptr,
53+ std::string socVersion = "Ascend910b",
54+ uint64_t coreNum = 64,
55+ uint64_t ubSize = 262144,
56+ uint64_t tilingDataSize = 4096) :
57+ opName_(opName),
58+ inputTensorDesc_(inputTensorDesc),
59+ outputTensorDesc_(outputTensorDesc),
60+ attrs_(attrs),
61+ compileInfo_(compileInfo),
62+ socVersion_(socVersion),
63+ coreNum_(coreNum),
64+ ubSize_(ubSize),
65+ tilingDataSize_(tilingDataSize) {}
66+ 
67+ TilingContextPara(const std::string& opName,
68+ const std::vector<TensorDescription>& inputTensorDesc,
69+ const std::vector<TensorDescription>& outputTensorDesc,
70+ void* compileInfo = nullptr,
71+ std::string socVersion = "Ascend910b",
72+ uint64_t coreNum = 64,
73+ uint64_t ubSize = 262144,
74+ uint64_t tilingDataSize = 4096) :
75+ opName_(opName),
76+ inputTensorDesc_(inputTensorDesc),
77+ outputTensorDesc_(outputTensorDesc),
78+ compileInfo_(compileInfo),
79+ socVersion_(socVersion),
80+ coreNum_(coreNum),
81+ ubSize_(ubSize),
82+ tilingDataSize_(tilingDataSize) {}
83+ 
84+ TilingContextPara(const std::string& opName,
85+ const std::vector<TensorDescription>& inputTensorDesc,
86+ const std::vector<TensorDescription>& outputTensorDesc,
87+ const std::vector<OpAttr>& attrs,
88+ const std::vector<uint32_t>& inputInstanceNum,
89+ const std::vector<uint32_t>& outputInstanceNum,
90+ void* compileInfo = nullptr,
91+ std::string socVersion = "Ascend910b",
92+ uint64_t coreNum = 64,
93+ uint64_t ubSize = 262144,
94+ uint64_t tilingDataSize = 4096) :
95+ opName_(opName),
96+ inputTensorDesc_(inputTensorDesc),
97+ outputTensorDesc_(outputTensorDesc),
98+ attrs_(attrs),
99+ inputInstanceNum_(inputInstanceNum),
100+ outputInstanceNum_(outputInstanceNum),
101+ compileInfo_(compileInfo),
102+ socVersion_(socVersion),
103+ coreNum_(coreNum),
104+ ubSize_(ubSize),
105+ tilingDataSize_(tilingDataSize) {}
106+ 
107+ TilingContextPara(const std::string& opName,
108+ const std::vector<TensorDescription>& inputTensorDesc,
109+ const std::vector<TensorDescription>& outputTensorDesc,
110+ const std::vector<uint32_t>& inputInstanceNum,
111+ const std::vector<uint32_t>& outputInstanceNum,
112+ void* compileInfo = nullptr,
113+ std::string socVersion = "Ascend910b",
114+ uint64_t coreNum = 64,
115+ uint64_t ubSize = 262144,
116+ uint64_t tilingDataSize = 4096) :
117+ opName_(opName),
118+ inputTensorDesc_(inputTensorDesc),
119+ outputTensorDesc_(outputTensorDesc),
120+ inputInstanceNum_(inputInstanceNum),
121+ outputInstanceNum_(outputInstanceNum),
122+ compileInfo_(compileInfo),
123+ socVersion_(socVersion),
124+ coreNum_(coreNum),
125+ ubSize_(ubSize),
126+ tilingDataSize_(tilingDataSize) {}
127+public:
128+ std::string opName_;
129+ std::vector<uint32_t> inputInstanceNum_;
130+ std::vector<uint32_t> outputInstanceNum_;
131+ std::vector<TensorDescription> inputTensorDesc_;
132+ std::vector<TensorDescription> outputTensorDesc_;
133+ std::vector<OpAttr> attrs_;
134+ uint64_t coreNum_ = 64;
135+ uint64_t ubSize_ = 262144;
136+ uint64_t tilingDataSize_ = 4096;
137+ std::string socVersion_ = "Ascend910b";
138+ void* compileInfo_ = nullptr;
139+};
140+ 
21class TilingContextFaker : public OpTilingContextBuilder, public KernelRunContextHolder {141class TilingContextFaker : public OpTilingContextBuilder, public KernelRunContextHolder {
22public:142public:
23 TilingContextFaker() = default;143 TilingContextFaker() = default;