Star1337
1454
代码
代码
Issues117
Pull Requests261
Actions
项目讨论
Wiki
分析
Star1337
1454
已合并
Pad/Tile算子支持下一代实现 #964
Guoqh创建于 1月28日
Pad/Tile算子支持下一代实现 #964
已合并
Guoqh创建于 1月28日
从guoqiuhao/ops-mathmigrate合入到cann/ops-mathmaster
共 38 个文件变更+2156-233
Mconversion/pad/CMakeLists.txt+6-10
@@ -1,5 +1,5 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.2+# Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
5# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -8,12 +8,8 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11+# 设置算子定义时支持的芯片类型
daiwei18
daiwei18daiwei181月28日

op_graph/fusion_pass/.gitkeep是新增的空文件，建议删除

likedislike
12-if(NOT ENABLE_TEST AND NOT BENCHMARK)12+set(SUPPORT_COMPUTE_UNIT "ascend950")
13- list(REMOVE_ITEM CURRENT_DIRS tests)13+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14-endif()14+set(SUPPORT_TILING_DIR "arch35")
15-foreach(SUB_DIR ${CURRENT_DIRS})15+add_all_modules_sources(OPTYPE pad ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES pad_v3)
16- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17- add_subdirectory(${SUB_DIR})
18- endif()
19-endforeach()
Rconversion/pad/op_host/op_api/pad.cpp→conversion/pad/op_api/pad.cpp+0-0
文件重命名但无更改。
Aconversion/pad/op_graph/pad_proto.h+48-0
@@ -0,0 +1,48 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file pad_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_OP_PROTO_INC_PAD_OPS_H_
16+#define OPS_OP_PROTO_INC_PAD_OPS_H_
17+ 
18+#include "graph/operator_reg.h"
19+ 
20+namespace ge {
21+/**
22+* @brief Pad a tensor.
23+ 
24+* @par Inputs:
25+* Two inputs, including:
26+* @li x: A Tensor. Must be one of the following types: bfloat16, float16,
27+* float32, double, int32, uint8, int16, int8, complex64, int64, qint8,
28+* quint8, qint32, qint16, quint16, uint16, complex128, uint32, uint64, bool. Supported format list ["ND"].
29+* @li paddings: A Tensor of type int32 or int64. Supported format list ["ND"]. \n
30+ 
31+* @par Outputs:
32+* y: A Tensor of the same type as "x". Supported format list ["ND"]. \n
33+ 
34+* @li Due to different architectures, the calculation results of this operator
35+* on NPU and CPU may be inconsistent. \n
36+ 
37+* @par Third-party framework compatibility:
38+* Compatible with TensorFlow operator Pad.
39+*/
40+REG_OP(Pad)
41+ .INPUT(x, TensorType({TensorType::BasicType(), DT_BOOL}))
42+ .INPUT(paddings, TensorType::IndexNumberType())
43+ .OUTPUT(y, TensorType({TensorType::BasicType(), DT_BOOL}))
44+ .OP_END_FACTORY_REG(Pad)
45+ 
46+} // namespace ge
47+ 
48+#endif
Dconversion/pad/op_host/CMakeLists.txt+0-11
@@ -1,11 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# ----------------------------------------------------------------------------
10- 
11-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
Aconversion/pad/op_host/arch35/pad_tiling_arch35.cpp+34-0
@@ -0,0 +1,34 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file pad_tiling_arch35.cpp
13+ * \brief
14+ */
15+#include "pad_tiling_arch35.h"
16+#include "conversion/pad_v3/op_host/arch35/pad_v3_tiling_arch35.h"
17+#include "log/log.h"
18+ 
19+namespace optiling {
20+ 
21+static ge::graphStatus Tiling4Pad(gert::TilingContext* context) {
22+ PadACTiling tilingObject(context);
23+ return tilingObject.DoTiling();
24+ 
25+}
26+ 
27+static ge::graphStatus TilingPrepare4Pad(gert::TilingParseContext* context) {
28+ OP_LOGD(context->GetNodeName(), "TilingPreparePad entering.");
29+ return ge::GRAPH_SUCCESS;
30+}
31+ 
32+// register tiling interface of the Pad op.
33+IMPL_OP_OPTILING(Pad).Tiling(Tiling4Pad).TilingParse<PadCompileInfo>(TilingPrepare4Pad);
34+} // namespace optiling
Aconversion/pad/op_host/arch35/pad_tiling_arch35.h+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file pad_tiling_arch35.h
13+ * \brief
14+ */
15+ 
16+#ifndef OPS_PAD_TILING_H_
17+#define OPS_PAD_TILING_H_
18+ 
19+#include <cstdint>
20+ 
21+namespace optiling {
22+struct PadCompileInfo {
23+ int64_t core_num;
24+ int64_t ub_size;
25+};
26+ 
27+} // namespace optiling
28+ 
29+#endif // OPS_PAD_TILING_H_
Aconversion/pad/op_host/config/ascend950/pad_binary.json+341-0
@@ -0,0 +1,341 @@
1+{
2+ "op_type": "Pad",
3+ "op_list": [
4+ {
5+ "bin_filename": "Pad_int8_int32",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "int8",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "dtype_match_mode": "DtypeByte",
17+ "format_match_mode": "FormatAgnostic"
18+ },
19+ {
20+ "name": "paddings",
21+ "index": 1,
22+ "dtype": "int32",
23+ "format": "ND",
24+ "paramType": "required",
25+ "shape": [
26+ -2
27+ ],
28+ "format_match_mode": "FormatAgnostic"
29+ }
30+ ],
31+ "outputs": [
32+ {
33+ "name": "y",
34+ "index": 0,
35+ "dtype": "int8",
36+ "format": "ND",
37+ "paramType": "required",
38+ "shape": [
39+ -2
40+ ],
41+ "dtype_match_mode": "DtypeByte",
42+ "format_match_mode": "FormatAgnostic"
43+ }
44+ ]
45+ },
46+ {
47+ "bin_filename": "Pad_int8_int64",
48+ "inputs": [
49+ {
50+ "name": "x",
51+ "index": 0,
52+ "dtype": "int8",
53+ "format": "ND",
54+ "paramType": "required",
55+ "shape": [
56+ -2
57+ ],
58+ "dtype_match_mode": "DtypeByte",
59+ "format_match_mode": "FormatAgnostic"
60+ },
61+ {
62+ "name": "paddings",
63+ "index": 1,
64+ "dtype": "int64",
65+ "format": "ND",
66+ "paramType": "required",
67+ "shape": [
68+ -2
69+ ],
70+ "format_match_mode": "FormatAgnostic"
71+ }
72+ ],
73+ "outputs": [
74+ {
75+ "name": "y",
76+ "index": 0,
77+ "dtype": "int8",
78+ "format": "ND",
79+ "paramType": "required",
80+ "shape": [
81+ -2
82+ ],
83+ "dtype_match_mode": "DtypeByte",
84+ "format_match_mode": "FormatAgnostic"
85+ }
86+ ]
87+ },
88+ {
89+ "bin_filename": "Pad_float16_int32",
90+ "inputs": [
91+ {
92+ "name": "x",
93+ "index": 0,
94+ "dtype": "float16",
95+ "format": "ND",
96+ "paramType": "required",
97+ "shape": [
98+ -2
99+ ],
100+ "dtype_match_mode": "DtypeByte",
101+ "format_match_mode": "FormatAgnostic"
102+ },
103+ {
104+ "name": "paddings",
105+ "index": 1,
106+ "dtype": "int32",
107+ "format": "ND",
108+ "paramType": "required",
109+ "shape": [
110+ -2
111+ ],
112+ "format_match_mode": "FormatAgnostic"
113+ }
114+ ],
115+ "outputs": [
116+ {
117+ "name": "y",
118+ "index": 0,
119+ "dtype": "float16",
120+ "format": "ND",
121+ "paramType": "required",
122+ "shape": [
123+ -2
124+ ],
125+ "dtype_match_mode": "DtypeByte",
126+ "format_match_mode": "FormatAgnostic"
127+ }
128+ ]
129+ },
130+ {
131+ "bin_filename": "Pad_float16_int64",
132+ "inputs": [
133+ {
134+ "name": "x",
135+ "index": 0,
136+ "dtype": "float16",
137+ "format": "ND",
138+ "paramType": "required",
139+ "shape": [
140+ -2
141+ ],
142+ "dtype_match_mode": "DtypeByte",
143+ "format_match_mode": "FormatAgnostic"
144+ },
145+ {
146+ "name": "paddings",
147+ "index": 1,
148+ "dtype": "int64",
149+ "format": "ND",
150+ "paramType": "required",
151+ "shape": [
152+ -2
153+ ],
154+ "format_match_mode": "FormatAgnostic"
155+ }
156+ ],
157+ "outputs": [
158+ {
159+ "name": "y",
160+ "index": 0,
161+ "dtype": "float16",
162+ "format": "ND",
163+ "paramType": "required",
164+ "shape": [
165+ -2
166+ ],
167+ "dtype_match_mode": "DtypeByte",
168+ "format_match_mode": "FormatAgnostic"
169+ }
170+ ]
171+ },
172+ {
173+ "bin_filename": "Pad_float32_int32",
174+ "inputs": [
175+ {
176+ "name": "x",
177+ "index": 0,
178+ "dtype": "float32",
179+ "format": "ND",
180+ "paramType": "required",
181+ "shape": [
182+ -2
183+ ],
184+ "dtype_match_mode": "DtypeByte",
185+ "format_match_mode": "FormatAgnostic"
186+ },
187+ {
188+ "name": "paddings",
189+ "index": 1,
190+ "dtype": "int32",
191+ "format": "ND",
192+ "paramType": "required",
193+ "shape": [
194+ -2
195+ ],
196+ "format_match_mode": "FormatAgnostic"
197+ }
198+ ],
199+ "outputs": [
200+ {
201+ "name": "y",
202+ "index": 0,
203+ "dtype": "float32",
204+ "format": "ND",
205+ "paramType": "required",
206+ "shape": [
207+ -2
208+ ],
209+ "dtype_match_mode": "DtypeByte",
210+ "format_match_mode": "FormatAgnostic"
211+ }
212+ ]
213+ },
214+ {
215+ "bin_filename": "Pad_float32_int64",
216+ "inputs": [
217+ {
218+ "name": "x",
219+ "index": 0,
220+ "dtype": "float32",
221+ "format": "ND",
222+ "paramType": "required",
223+ "shape": [
224+ -2
225+ ],
226+ "dtype_match_mode": "DtypeByte",
227+ "format_match_mode": "FormatAgnostic"
228+ },
229+ {
230+ "name": "paddings",
231+ "index": 1,
232+ "dtype": "int64",
233+ "format": "ND",
234+ "paramType": "required",
235+ "shape": [
236+ -2
237+ ],
238+ "format_match_mode": "FormatAgnostic"
239+ }
240+ ],
241+ "outputs": [
242+ {
243+ "name": "y",
244+ "index": 0,
245+ "dtype": "float32",
246+ "format": "ND",
247+ "paramType": "required",
248+ "shape": [
249+ -2
250+ ],
251+ "dtype_match_mode": "DtypeByte",
252+ "format_match_mode": "FormatAgnostic"
253+ }
254+ ]
255+ },
256+ {
257+ "bin_filename": "Pad_int64_int32",
258+ "inputs": [
259+ {
260+ "name": "x",
261+ "index": 0,
262+ "dtype": "int64",
263+ "format": "ND",
264+ "paramType": "required",
265+ "shape": [
266+ -2
267+ ],
268+ "dtype_match_mode": "DtypeByte",
269+ "format_match_mode": "FormatAgnostic"
270+ },
271+ {
272+ "name": "paddings",
273+ "index": 1,
274+ "dtype": "int32",
275+ "format": "ND",
276+ "paramType": "required",
277+ "shape": [
278+ -2
279+ ],
280+ "format_match_mode": "FormatAgnostic"
281+ }
282+ ],
283+ "outputs": [
284+ {
285+ "name": "y",
286+ "index": 0,
287+ "dtype": "int64",
288+ "format": "ND",
289+ "paramType": "required",
290+ "shape": [
291+ -2
292+ ],
293+ "dtype_match_mode": "DtypeByte",
294+ "format_match_mode": "FormatAgnostic"
295+ }
296+ ]
297+ },
298+ {
299+ "bin_filename": "Pad_int64_int64",
300+ "inputs": [
301+ {
302+ "name": "x",
303+ "index": 0,
304+ "dtype": "int64",
305+ "format": "ND",
306+ "paramType": "required",
307+ "shape": [
308+ -2
309+ ],
310+ "dtype_match_mode": "DtypeByte",
311+ "format_match_mode": "FormatAgnostic"
312+ },
313+ {
314+ "name": "paddings",
315+ "index": 1,
316+ "dtype": "int64",
317+ "format": "ND",
318+ "paramType": "required",
319+ "shape": [
320+ -2
321+ ],
322+ "format_match_mode": "FormatAgnostic"
323+ }
324+ ],
325+ "outputs": [
326+ {
327+ "name": "y",
328+ "index": 0,
329+ "dtype": "int64",
330+ "format": "ND",
331+ "paramType": "required",
332+ "shape": [
333+ -2
334+ ],
335+ "dtype_match_mode": "DtypeByte",
336+ "format_match_mode": "FormatAgnostic"
337+ }
338+ ]
339+ }
340+ ]
341+}
Aconversion/pad/op_host/config/ascend950/pad_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+[Pad]
13+default=0
Aconversion/pad/op_host/pad_def.cpp+62-0
@@ -0,0 +1,62 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file pad_def.cpp
13+ * \brief Pad ophost
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace {
18+static const std::vector<ge::Format> format = {
19+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
20+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
21+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
22+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND
23+};
24+static const std::vector<ge::DataType> valueDataType = {
25+ ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
26+ ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_DOUBLE, ge::DT_BOOL,
27+ ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
28+ ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_DOUBLE, ge::DT_BOOL
29+};
30+ 
31+static const std::vector<ge::DataType> padDataType = {
32+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
33+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
34+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
35+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32
36+};
37+}
38+ 
39+namespace ops {
40+class Pad : public OpDef {
41+ public:
42+ explicit Pad(const char* name) : OpDef(name) {
43+ this->Input("x").ParamType(REQUIRED).DataType(valueDataType).Format(format);
44+ this->Input("paddings")
45+ .ParamType(REQUIRED)
46+ .ValueDepend(OPTIONAL)
47+ .DataType(padDataType)
48+ .Format(format);
49+ this->Output("y").ParamType(REQUIRED).DataType(valueDataType).Format(format);
50+ 
51+ OpAICoreConfig aicore_config;
52+ aicore_config.DynamicCompileStaticFlag(true)
53+ .DynamicRankSupportFlag(true)
54+ .DynamicShapeSupportFlag(true)
55+ .NeedCheckSupportFlag(false)
56+ .ExtendCfgInfo("opFile.value", "pad_apt");
57+ this->AICore().AddConfig("ascend950", aicore_config);
58+ }
59+};
60+ 
61+OP_ADD(Pad);
62+} // namespace ops
Aconversion/pad/op_host/pad_infershape.cpp+131-0
@@ -0,0 +1,131 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file pad_infershape.cpp
13+ * \brief
14+ */
15+ 
16+#include "register/op_impl_registry.h"
17+#include "op_api/op_util.h"
18+#include "log/log.h"
19+#include "util/shape_util.h"
20+ 
21+using namespace ge;
22+using namespace Ops::Base;
23+ 
24+namespace ops {
25+static constexpr size_t PAD_IN_IDX_X = 0;
26+static constexpr size_t PAD_IN_IDX_PADDINGS = 1;
27+static constexpr size_t PAD_OUT_IDX_Y = 0;
28+static constexpr size_t INT_DATA_2 = 2;
29+static constexpr int64_t UNKNOWN_DIM_VALUE_ = -1L;
30+ 
31+template <typename T>
32+static bool PadInfershape(const gert::InferShapeContext* context, const gert::Shape* x_shape, const T* paddings_value,
33+ const size_t paddings_size, gert::Shape* y_shape) {
34+ OP_LOGD(context->GetNodeName(), "Begin to do PadInfershape");
35+ OP_LOGD(context->GetNodeName(), "input x = %s", ToString(*x_shape).c_str());
36+ size_t input_dim_size = x_shape->GetDimNum();
37+ OP_CHECK_IF(input_dim_size == 0,
38+ OP_LOGE(context->GetNodeName(), "input shape cannot empty"),
39+ return false);
40+ if (input_dim_size * INT_DATA_2 != paddings_size) {
41+ OP_LOGE(context->GetNodeName(), "the paddings num must be twice of the input x rank."
42+ "but paddings num is %zu, input x rank is %zu", paddings_size, input_dim_size);
43+ return false;
44+ }
45+ y_shape->SetDimNum(input_dim_size);
46+ int64_t dim_value = UNKNOWN_DIM_VALUE_;
47+ for (size_t i = 0; i < input_dim_size; ++i) {
48+ dim_value = x_shape->GetDim(i) == UNKNOWN_DIM_VALUE_
49+ ? UNKNOWN_DIM_VALUE_
50+ : x_shape->GetDim(i) + paddings_value[INT_DATA_2 * i] + paddings_value[INT_DATA_2 * i + 1];
51+ if (x_shape->GetDim(i) != UNKNOWN_DIM_VALUE_ && dim_value < 0) {
52+ OP_LOGE(
53+ context->GetNodeName(),
54+ "The output shape at index %zu is %ld, but output shape CANNOT contain negative values. x_shape at "
55+ "index %zu: %ld, corresponding pad_front: %ld, corresponding pad_end: %ld.",
56+ i, dim_value, i, x_shape->GetDim(i), static_cast<int64_t>(paddings_value[INT_DATA_2 * i]),
57+ static_cast<int64_t>(paddings_value[INT_DATA_2 * i + 1]));
58+ return false;
59+ }
60+ y_shape->SetDim(i, dim_value);
61+ }
62+ OP_LOGD(context->GetNodeName(), "output y = %s", ToString(*y_shape).c_str());
63+ OP_LOGD(context->GetNodeName(), "End to do PadInfershape");
64+ 
65+ return true;
66+}
67+ 
68+template <typename T>
69+ge::graphStatus PadInfershapeWithTensor(const gert::InferShapeContext* context, const gert::Shape* x_shape,
70+ const gert::Tensor* paddings_tensor, gert::Shape* y_shape) {
71+ const T* paddings_value = paddings_tensor->GetData<T>();
72+ const size_t paddings_num = paddings_tensor->GetShapeSize();
73+ OP_CHECK_IF(!PadInfershape<T>(context, x_shape, paddings_value, paddings_num, y_shape),
74+ OP_LOGE(context->GetNodeName(), "do PadInfershape failed"),
75+ return ge::GRAPH_FAILED);
76+ 
77+ return ge::GRAPH_SUCCESS;
78+}
79+ 
80+static ge::graphStatus SetAllUnknownDim(const int64_t rank, gert::Shape* output_shape)
81+{
82+ OP_CHECK_IF(
83+ output_shape == nullptr, OP_LOGD("SetAllUnknownDim", "the output_shape is nullptr, return unsuccess"),
84+ return ge::GRAPH_FAILED);
85+ 
86+ output_shape->SetDimNum(rank);
87+ for (int64_t i = 0; i < rank; ++i) {
88+ output_shape->SetDim(i, UNKNOWN_DIM_VALUE_);
89+ }
90+ OP_LOGD("SetAllUnknownDim", "set all dim = -1, output = %s", Ops::Base::ToString(*output_shape).c_str());
91+ 
92+ return ge::GRAPH_SUCCESS;
93+}
94+ 
95+static ge::graphStatus InferShapeForPad(gert::InferShapeContext* context) {
96+ const gert::Shape* x_shape = context->GetInputShape(PAD_IN_IDX_X);
97+ OP_CHECK_NULL_WITH_CONTEXT(context, x_shape);
98+ gert::Shape* y_shape = context->GetOutputShape(PAD_OUT_IDX_Y);
99+ OP_CHECK_NULL_WITH_CONTEXT(context, y_shape);
100+ const gert::Tensor* paddings_tensor = context->GetInputTensor(PAD_IN_IDX_PADDINGS);
101+ OP_CHECK_NULL_WITH_CONTEXT(context, paddings_tensor);
102+ 
103+ // if x_shape is unknown rank [-2] that means cannot know how many ranks,
104+ // which make output unknown rank.
105+ if (IsUnknownRank(*x_shape)) {
106+ SetUnknownRank(*y_shape);
107+ return ge::GRAPH_SUCCESS;
108+ }
109+ 
110+ if (IsConstTensor(paddings_tensor)) {
111+ ge::DataType paddings_dtype = paddings_tensor->GetDataType();
112+ switch (paddings_dtype) {
113+ case ge::DT_INT32: {
114+ return PadInfershapeWithTensor<int32_t>(context, x_shape, paddings_tensor, y_shape);
115+ }
116+ case ge::DT_INT64: {
117+ return PadInfershapeWithTensor<int64_t>(context, x_shape, paddings_tensor, y_shape);
118+ }
119+ default:
120+ OP_LOGE_WITH_INVALID_INPUT_DTYPE(
121+ context->GetNodeName(),
122+ "paddings", Ops::Base::ToString(paddings_dtype).c_str(), "[int32, int64]");
123+ return ge::GRAPH_FAILED;
124+ }
125+ } else {
126+ return SetAllUnknownDim(x_shape->GetDimNum(), y_shape);
127+ }
128+}
129+ 
130+IMPL_OP_INFERSHAPE(Pad).InferShape(InferShapeForPad).InputsDataDependency({PAD_IN_IDX_PADDINGS});
131+} // namespace ops
Aconversion/pad/op_kernel/pad_apt.cpp+106-0
@@ -0,0 +1,106 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file pad_apt.cpp
13+ * \brief pad kernel
14+ */
15+#include "../pad_v3/arch35/pad_constant.h"
16+#include "../pad_v3/arch35/pad_slice.h"
17+ 
18+using namespace PadV3;
19+ 
20+#define CONSTANT_SLICE_BRANCH 10000
21+#define CONSTANT_SIMT_BRANCH 20000
22+#define CONSTANT_SIMT_BIG_SIZE_BRANCH 20001
23+#define CONSTANT_CUT_LAST_DIM_BRANCH 30010
24+#define CONSTANT_BIG_LAST_DIM_BRANCH_DIM2 30021
25+#define CONSTANT_BIG_LAST_DIM_BRANCH_DIM3 30031
26+#define CONSTANT_BIG_LAST_DIM_BRANCH_DIM4 30041
27+#define CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM2 30022
28+#define CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM3 30032
29+#define CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM4 30042
30+#define CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM2 30023
31+#define CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM3 30033
32+#define CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM4 30043
33+ 
34+#define PAD_SLICE_KEY_MOVE_ALIGN 10100
35+#define PAD_SLICE_KEY_MOVE_ALIGN_LAST_DIM 10101
36+#define PAD_SLICE_KEY_NDDMA 10102
37+#define PAD_SLICE_KEY_NDDMA_LAST_DIM 10103
38+#define PAD_SLICE_KEY_MOVE_ALIGN_TWO_DIM 10150
39+#define PAD_SLICE_KEY_SIMT 10200
40+#define PAD_SLICE_KEY_MOVE_ALIGN_GATHER 10300
41+#define PAD_SLICE_KEY_MOVE_UNALIGN_GATHER 10301
42+#define PAD_SLICE_KEY_TWO_DIM_SMALL_SHAPE 10400
43+ 
44+extern "C" __global__ __aicore__ void pad(GM_ADDR x, GM_ADDR paddings, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
45+{
46+ if (workspace == nullptr) {
47+ return;
48+ }
49+ SetSysWorkspace(workspace);
50+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
51+ REGISTER_TILING_DEFAULT(SliceFakeTilingData);
52+ if (TILING_KEY_IS(CONSTANT_CUT_LAST_DIM_BRANCH)) { // 30000
53+ PadV3::LaunchKernelPadWithHugeWidth<DTYPE_X>(x, paddings, y, tiling);
54+ } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM2)) { // 30021
55+ PadV3::LaunchKernelPadWithNormalWidth<DTYPE_X, CONSTANT_BIG_LAST_DIM_BRANCH_DIM2>(x, paddings, y, tiling);
56+ } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM3)) { // 30031
57+ PadV3::LaunchKernelPadWithNormalWidth<DTYPE_X, CONSTANT_BIG_LAST_DIM_BRANCH_DIM3>(x, paddings, y, tiling);
58+ } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM4)) { // 30041
59+ PadV3::LaunchKernelPadWithNormalWidth<DTYPE_X, CONSTANT_BIG_LAST_DIM_BRANCH_DIM4>(x, paddings, y, tiling);
60+ } else if (
61+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM2) ||
62+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM3) ||
63+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM4)) { // 30002
64+ PadV3::LaunchKernelPadGather<DTYPE_X>(x, paddings, y, tiling);
65+ } else if (
66+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM2) ||
67+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM3) ||
68+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM4)) { // 30002
69+ PadV3::LaunchKernelPadScatter<DTYPE_X>(x, paddings, y, tiling);
70+ } else if (TILING_KEY_IS(CONSTANT_SIMT_BRANCH)) { // 20000
71+ PadV3::LaunchKernelPadSimt<DTYPE_X>(x, paddings, y, tiling);
72+ } else if (TILING_KEY_IS(CONSTANT_SIMT_BIG_SIZE_BRANCH)) { // 20001
73+ PadV3::LaunchKernelPadSimtHuge<DTYPE_X>(x, paddings, y, tiling);
74+ } else {
75+ TPipe pipe;
76+ __gm__ uint8_t* offsets = nullptr;
77+ __gm__ uint8_t* size = nullptr;
78+ if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN)) {
79+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignTilingData, tilingData, tiling);
80+ PadSliceMoveAlignProcess(x, offsets, size, y, &tilingData, &pipe);
81+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_NDDMA)) {
82+ GET_TILING_DATA_WITH_STRUCT(SliceNDDMATilingData, tilingData, tiling);
83+ PadSliceNDDMAProcess(x, offsets, size, y, &tilingData, &pipe);
84+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_LAST_DIM)) {
85+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLastDimTilingData, tilingData, tiling);
86+ PadSliceMoveAlignLastDimProcess(x, offsets, size, y, &tilingData, &pipe);
87+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_NDDMA_LAST_DIM)) {
88+ GET_TILING_DATA_WITH_STRUCT(SliceNDDMALastDimTilingData, tilingData, tiling);
89+ PadSliceNDDMALastDimProcess(x, offsets, size, y, &tilingData, &pipe);
90+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_TWO_DIM)) {
91+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);
92+ PadSliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);
93+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_SIMT)) {
94+ // 空tenseor处理
95+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_GATHER)) {
96+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);
97+ PadSliceMoveAlignGatherProcess(x, offsets, size, y, &tilingData, &pipe);
98+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_UNALIGN_GATHER)) {
99+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);
100+ PadSliceMoveAlignDataCopyUnalignProcess(x, offsets, size, y, &tilingData, &pipe);
101+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_TWO_DIM_SMALL_SHAPE)) {
102+ GET_TILING_DATA_WITH_STRUCT(SliceTwoDimSmallSapeTilingData, tilingData, tiling);
103+ PadSliceTwoDimSmallShapeProcess(x, offsets, size, y, &tilingData, &pipe);
104+ }
105+ }
106+}
Dconversion/pad/tests/CMakeLists.txt+0-16
@@ -1,16 +0,0 @@
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()
Dconversion/pad/tests/ut/CMakeLists.txt+0-16
@@ -1,16 +0,0 @@
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()
Rconversion/pad/tests/ut/op_host/op_api/test_pad.cpp→conversion/pad/tests/ut/op_api/test_pad.cpp+0-0
文件重命名但无更改。
Dconversion/pad/tests/ut/op_host/CMakeLists.txt+0-20
@@ -1,20 +0,0 @@
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()
17- 
18-if(TILING_UT OR PROTO_UT OR OP_API_UT OR (UT_TEST_ALL AND NOT AICPU_ONLY))
19- # add_modules_llt_sources(HOSTNAME ${OPTEST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
20-endif()
Aconversion/pad/tests/ut/op_host/arch35/test_pad_tiling.cpp+179-0
@@ -0,0 +1,179 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file test_pad_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "../../../../op_host/arch35/pad_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+class PadTilingTest : public testing::Test {
25+ protected:
26+ static void SetUpTestCase() {
27+ std::cout << "PadTilingTest SetUp" << std::endl;
28+ }
29+ 
30+ static void TearDownTestCase() {
31+ std::cout << "PadTilingTest TearDown" << std::endl;
32+ }
33+};
34+ 
35+TEST_F(PadTilingTest, Pad_AC_tiling_NDDMA_20000) // NDDMA_20000
36+{
37+ optiling::PadCompileInfo compileInfo;
38+ compileInfo.core_num = 64;
39+ compileInfo.ub_size = 245760; // 240 * 1024
40+ std::vector<int32_t> paddingsValue = {4, 1};
41+ gert::StorageShape xShape = {{55}, {55}};
42+ gert::StorageShape paddingsShape = {{1, 2}, {1, 2}};
43+ gert::StorageShape yShape = {{60}, {60}};
44+ gert::TilingContextPara tilingContextPara(
45+ "Pad",
46+ {
47+ { xShape, ge::DT_FLOAT, ge::FORMAT_ND },
48+ { paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data() }
49+ },
50+ {
51+ { yShape, ge::DT_FLOAT, ge::FORMAT_ND }
52+ },
53+ &compileInfo);
54+ uint64_t expectedTilingKey = 20000;
55+ std::vector<size_t> expectedWorkspaces = { 16777216 };
56+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectedTilingKey, expectedWorkspaces);
57+}
58+ 
59+TEST_F(PadTilingTest, Pad_AC_tiling_NDDMA_30010) // NDDMA_30010 ----
60+{
61+ optiling::PadCompileInfo compileInfo;
62+ compileInfo.core_num = 64;
63+ compileInfo.ub_size = 245760; // 240 * 1024
64+ std::vector<int32_t> paddingsValue = {9, 17};
65+ gert::StorageShape xShape = {{1991203}, {1991203}};
66+ gert::StorageShape paddingsShape = {{1, 2}, {1, 2}};
67+ gert::StorageShape yShape = {{1991203}, {1991203}};
68+ gert::TilingContextPara tilingContextPara(
69+ "Pad",
70+ {
71+ { xShape, ge::DT_FLOAT, ge::FORMAT_ND },
72+ { paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data() }
73+ },
74+ {
75+ { yShape, ge::DT_FLOAT, ge::FORMAT_ND }
76+ },
77+ &compileInfo);
78+ uint64_t expectedTilingKey = 30010;
79+ std::vector<size_t> expectedWorkspaces = { 16777216 };
80+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectedTilingKey, expectedWorkspaces);
81+}
82+ 
83+TEST_F(PadTilingTest, Pad_AC_tiling_NDDMA_30021) // NDDMA_30021
84+{
85+ optiling::PadCompileInfo compileInfo;
86+ compileInfo.core_num = 64;
87+ compileInfo.ub_size = 245760; // 240 * 1024
88+ std::vector<int32_t> paddingsValue = {20, 25, 10, 18};
89+ gert::StorageShape xShape = {{1239, 1025}, {1239, 1025}};
90+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
91+ gert::StorageShape yShape = {{1284, 1053}, {1284, 1053}};
92+ gert::TilingContextPara tilingContextPara(
93+ "Pad",
94+ {
95+ { xShape, ge::DT_BOOL, ge::FORMAT_ND },
96+ { paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data() }
97+ },
98+ {
99+ { yShape, ge::DT_BOOL, ge::FORMAT_ND }
100+ },
101+ &compileInfo);
102+ uint64_t expectedTilingKey = 30021;
103+ std::vector<size_t> expectedWorkspaces = { 16777216 };
104+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectedTilingKey, expectedWorkspaces);
105+}
106+ 
107+TEST_F(PadTilingTest, Pad_AC_tiling_NDDMA_30023) // NDDMA_30023 ---
108+{
109+ optiling::PadCompileInfo compileInfo;
110+ compileInfo.core_num = 64;
111+ compileInfo.ub_size = 245760; // 240 * 1024
112+ std::vector<int32_t> paddingsValue = {15, 22, 11, 1, 22, 5};
113+ gert::StorageShape xShape = {{109, 7078, 4}, {109, 7078, 4}};
114+ gert::StorageShape paddingsShape = {{3, 2}, {3, 2}};
115+ gert::StorageShape yShape = {{146, 7090, 31}, {146, 7090, 31}};
116+ gert::TilingContextPara tilingContextPara(
117+ "Pad",
118+ {
119+ { xShape, ge::DT_INT8, ge::FORMAT_ND },
120+ { paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data() }
121+ },
122+ {
123+ { yShape, ge::DT_INT8, ge::FORMAT_ND }
124+ },
125+ &compileInfo);
126+ uint64_t expectedTilingKey = 30023;
127+ std::vector<size_t> expectedWorkspaces = { 16777216 };
128+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectedTilingKey, expectedWorkspaces);
129+}
130+ 
131+TEST_F(PadTilingTest, Pad_AC_tiling_NDDMA_30031) // NDDMA_30031 ---
132+{
133+ optiling::PadCompileInfo compileInfo;
134+ compileInfo.core_num = 64;
135+ compileInfo.ub_size = 245760; // 240 * 1024
136+ std::vector<int32_t> paddingsValue = {5, 3, 4, 4, 4, 4, 5, 4};
137+ gert::StorageShape xShape = {{70, 73, 74, 69}, {70, 73, 74, 69}};
138+ gert::StorageShape paddingsShape = {{4, 2}, {4, 2}};
139+ gert::StorageShape yShape = {{78, 81, 82, 78}, {78, 81, 82, 78}};
140+ gert::TilingContextPara tilingContextPara(
141+ "Pad",
142+ {
143+ { xShape, ge::DT_FLOAT16, ge::FORMAT_ND },
144+ { paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data() }
145+ },
146+ {
147+ { yShape, ge::DT_FLOAT16, ge::FORMAT_ND }
148+ },
149+ &compileInfo);
150+ uint64_t expectedTilingKey = 30031;
151+ std::vector<size_t> expectedWorkspaces = { 16777216 };
152+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectedTilingKey, expectedWorkspaces);
153+}
154+ 
155+TEST_F(PadTilingTest, Pad_AC_tiling_NDDMA_30032) // NDDMA_30033 ---
156+{
157+ optiling::PadCompileInfo compileInfo;
158+ compileInfo.core_num = 64;
159+ compileInfo.ub_size = 245760; // 240 * 1024
160+ std::vector<int32_t> paddingsValue = {24, 14, 6, 17, 6, 7, 15, 17};
161+ gert::StorageShape xShape = {{4, 45, 42, 21}, {4, 45, 42, 21}};
162+ gert::StorageShape paddingsShape = {{4, 2}, {4, 2}};
163+ gert::StorageShape yShape = {{42, 68, 55, 53}, {42, 68, 55, 53}};
164+ gert::TilingContextPara tilingContextPara(
165+ "Pad",
166+ {
167+ { xShape, ge::DT_BOOL, ge::FORMAT_ND },
168+ { paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data() }
169+ },
170+ {
171+ { yShape, ge::DT_BOOL, ge::FORMAT_ND }
172+ },
173+ &compileInfo);
174+ uint64_t expectedTilingKey = 30033;
175+ std::vector<size_t> expectedWorkspaces = { 16777216 };
176+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectedTilingKey, expectedWorkspaces);
177+}
178+ 
179+ 
Dconversion/pad/tests/ut/op_host/op_api/CMakeLists.txt+0-11
@@ -1,11 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# ----------------------------------------------------------------------------
10- 
11-# add_modules_llt_sources(HOSTNAME ${OPTEST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
Aconversion/pad/tests/ut/op_host/test_pad_infershape.cpp+89-0
@@ -0,0 +1,89 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+#include <iostream>
13+#include "infershape_context_faker.h"
14+#include "infershape_case_executor.h"
15+ 
16+class Pad : public testing::Test {
17+ protected:
18+ static void SetUpTestCase() {
19+ setenv("ASCEND_SLOG_PRINT_TO_STDOUT", "1", true);
20+ std::cout << "Pad SetUp" << std::endl;
21+ }
22+ 
23+ static void TearDownTestCase() {
24+ unsetenv("ASCEND_SLOG_PRINT_TO_STDOUT");
25+ std::cout << "Pad TearDown" << std::endl;
26+ }
27+};
28+ 
29+using namespace ge;
30+ 
31+TEST_F(Pad, Pad_const_infer_1) {
32+ std::vector<int32_t> values = {1, 1, 2, 2};
33+ gert::InfershapeContextPara infershapeContextPara("Pad",
34+ {
35+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
36+ {{{2, 2}, {2, 2}}, ge::DT_INT32, ge::FORMAT_ND, true, values.data()}
37+ },
38+ {
39+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
40+ }
41+ );
42+ std::vector<std::vector<int64_t>> expectOutputShape = {{5, 9}};
43+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
44+}
45+ 
46+TEST_F(Pad, Pad_const_infer_2_error_paddings_num) {
47+ std::vector<int32_t> values = {1, 1, 2, 2, 0, 0};
48+ gert::InfershapeContextPara infershapeContextPara("Pad",
49+ {
50+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
51+ {{{3, 2}, {3, 2}}, ge::DT_INT64, ge::FORMAT_ND, true, values.data()}
52+ },
53+ {
54+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
55+ }
56+ );
57+ std::vector<std::vector<int64_t>> expectOutputShape = {{5, 9}};
58+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
59+}
60+ 
61+TEST_F(Pad, Pad_const_infer_3_error_paddings_dtype) {
62+ std::vector<int32_t> values = {1, 1, 2, 2};
63+ gert::InfershapeContextPara infershapeContextPara("Pad",
64+ {
65+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
66+ {{{2, 2}, {2, 2}}, ge::DT_UINT64, ge::FORMAT_ND, true, values.data()}
67+ },
68+ {
69+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
70+ }
71+ );
72+ std::vector<std::vector<int64_t>> expectOutputShape = {{5, 9}};
73+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
74+}
75+ 
76+TEST_F(Pad, Pad_const_infer_4_error_paddings_num) {
77+ std::vector<int32_t> values = {1, 1, 2, 2, 0, 0};
78+ gert::InfershapeContextPara infershapeContextPara("Pad",
79+ {
80+ {{{3, 5}, {3, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
81+ {{{3, 2}, {3, 2}}, ge::DT_INT64, ge::FORMAT_ND, true, values.data()}
82+ },
83+ {
84+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
85+ }
86+ );
87+ std::vector<std::vector<int64_t>> expectOutputShape = {{5, 9}};
88+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
89+}
Mexperimental/CMakeLists.txt+1-1
@@ -23,7 +23,7 @@ set(COMMON_L0_SRCS
23 ${CMAKE_SOURCE_DIR}/conversion/tensor_move/op_api/tensor_move.cpp23 ${CMAKE_SOURCE_DIR}/conversion/tensor_move/op_api/tensor_move.cpp
24 ${CMAKE_SOURCE_DIR}/conversion/trans_data/op_host/op_api/transdata.cpp24 ${CMAKE_SOURCE_DIR}/conversion/trans_data/op_host/op_api/transdata.cpp
25 ${CMAKE_SOURCE_DIR}/conversion/transpose/op_api/transpose.cpp25 ${CMAKE_SOURCE_DIR}/conversion/transpose/op_api/transpose.cpp
26- ${CMAKE_SOURCE_DIR}/conversion/pad/op_host/op_api/pad.cpp26+ ${CMAKE_SOURCE_DIR}/conversion/pad/op_api/pad.cpp
27 ${CMAKE_SOURCE_DIR}/conversion/fill/op_api/fill.cpp27 ${CMAKE_SOURCE_DIR}/conversion/fill/op_api/fill.cpp
28 )28 )
29 29 
Mmath/tile/CMakeLists.txt+4-4
@@ -1,5 +1,5 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.2+# Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
5# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -9,7 +9,7 @@
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12-set(SUPPORT_COMPUTE_UNIT "ascend310p" "ascend910_93" "ascend910b" "ascend950")12+set(SUPPORT_COMPUTE_UNIT "ascend950")
13# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch35")14+set(SUPPORT_TILING_DIR "arch35")
15-add_all_modules_sources(OPTYPE mem_set_v2 ACLNNTYPE aclnn_exclude DISABLE_IN_OPP TRUE)15+add_all_modules_sources(OPTYPE tile ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES broadcast_to)
Mmath/tile/README.md+126-1
@@ -1,3 +1,128 @@
1# Tile1# Tile
2 2 
3-本目录仅包含Tile算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | √ |
12+| <term>Atlas 训练系列产品</term> | √ |
13+ 
14+## 功能说明
15+ 
16+接口功能：对输入tensor沿着repeats中对每个维度指定的复制次数进行复制。示例：
17+假设输入Tensor为[[a,b],[c,d],[e,f]]，即shape为[3,2]，repeats为(2,4)，则生成的Tensor的shape为[6,8]，值如下所示：
18+ 
19+```
20+>>> x = torch.tensor([[a,b],[c,d],[e,f]])
21+>>> x.repeat(2,4)
22+tensor([[a,b,a,b,a,b,a,b],
23+ [c,d,c,d,c,d,c,d],
24+ [e,f,e,f,e,f,e,f],
25+ [a,b,a,b,a,b,a,b],
26+ [c,d,c,d,c,d,c,d],
27+ [e,f,e,f,e,f,e,f],
28+ ])
29+```
30+当repeats为(2,4,2)时，即repeats的元素个数大于Tensor中的维度，则输出Tensor等效为如下操作：先将输入Tensor的shape扩张到和repeats个数相同的维度：[1,3,2]，而后按照对应维度和repeats的值进行扩张，即输出Tensor的shape为[2,12,4]，结果如下：
31+```
32+>>> x.repeat(2,4,2)
33+tensor([[[a,b,a,b],
34+ [c,d,c,d],
35+ [e,f,e,f],
36+ [a,b,a,b],
37+ [c,d,c,d],
38+ [e,f,e,f],
39+ [a,b,a,b],
40+ [c,d,c,d],
41+ [e,f,e,f],
42+ [a,b,a,b],
43+ [c,d,c,d],
44+ [e,f,e,f]],
45+ 
46+ [[a,b,a,b],
47+ [c,d,c,d],
48+ [e,f,e,f],
49+ [a,b,a,b],
50+ [c,d,c,d],
51+ [e,f,e,f],
52+ [a,b,a,b],
53+ [c,d,c,d],
54+ [e,f,e,f],
55+ [a,b,a,b],
56+ [c,d,c,d],
57+ [e,f,e,f]]])
58+```
59+计算时需要满足以下条件：
60+repeats中参数个数不能少于输入Tensor的维度。
61+repeats中的值必须大于等于0。
62+ 
63+## 参数说明
64+ 
65+ <table style="undefined;table-layout: fixed; width: 1528px"><colgroup>
66+ <col style="width: 132px">
67+ <col style="width: 120px">
68+ <col style="width: 256px">
69+ <col style="width: 253px">
70+ <col style="width: 333px">
71+ <col style="width: 126px">
72+ <col style="width: 160px">
73+ <col style="width: 145px">
74+ </colgroup>
75+ <thead>
76+ <tr>
77+ <th>参数名</th>
78+ <th>输入/输出</th>
79+ <th>描述</th>
80+ <th>使用说明</th>
81+ <th>数据类型</th>
82+ <th>数据格式</th>
83+ <th>维度（shape）</th>
84+ <th>非连续张量Tensor</th>
85+ </tr></thead>
86+ <tbody>
87+ <tr>
88+ <td>self</td>
89+ <td>输入</td>
90+ <td>-</td>
91+ <td>-</td>
92+ <td>FLOAT、DOUBLE、FLOAT16、COMPLEX64、COMPLEX128、UINT8、INT8、INT16、INT32、INT64、UINT16、UINT32、UINT64、BOOL、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN</td>
93+ <td>ND</td>
94+ <td>≤8</td>
95+ <td>√</td>
96+ </tr>
97+ <tr>
98+ <td>repeats</td>
99+ <td>输入</td>
100+ <td>-</td>
101+ <td>表示沿每个维度重复输入tensor的次数，参数个数不大于8, 当前不支持对超过4个维度同时做repeat的场景，详细约束请见约束说明。</td>
102+ <td>INT64</td>
103+ <td>-</td>
104+ <td>-</td>
105+ <td>-</td>
106+ </tr>
107+ <tr>
108+ <td>out</td>
109+ <td>输出</td>
110+ <td>-</td>
111+ <td>-</td>
112+ <td>与self一致</td>
113+ <td>ND</td>
114+ <td>≤8</td>
115+ <td>√</td>
116+ </tr>
117+ </tbody></table>
118+ 
119+## 约束说明
120+ 
121+- 确定性计算：
122+ - 默认确定性实现。
123+ 
124+## 调用示例
125+ 
126+| 调用方式 | 样例代码 | 说明 |
127+| ---------------- | --------------------------- | --------------------------------------------------- |
128+| aclnn接口 | [test_aclnn_repeat](./examples/test_aclnn_repeat.cpp) | 通过[aclnnRepeat](docs/aclnnRepeat.md)接口方式调用Tile算子。 |
Amath/tile/op_graph/tile_proto.h+54-0
@@ -0,0 +1,54 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_OP_PROTO_INC_TILE_OPS_H_
16+#define OPS_OP_PROTO_INC_TILE_OPS_H_
17+ 
18+#include "graph/operator_reg.h"
19+ 
20+namespace ge {
21+/**
22+* @brief Constructs a tensor by tiling a given tensor .
23+ 
24+* @par Inputs:
25+* Two inputs, including:
26+* @li x: A Tensor.
27+* Must be one of the following types: DT_FLOAT, DT_FLOAT16, DT_DOUBLE, DT_COMPLEX64, DT_COMPLEX128,
28+ DT_INT8, DT_UINT8, DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_INT64, DT_UINT64,
29+ DT_QINT8, DT_QUINT8, DT_QINT16, DT_QUINT16, DT_QINT32, DT_BF16, DT_BOOL,DT_HIFLOAT8, DT_FLOAT8_E5M2,
30+ DT_FLOAT8_E4M3FN
31+* @li multiples: A 1D Tensor of type int32 or int64.
32+* The length must be the same as the number of dimensions in "input"
33+ 
34+* @par Outputs:
35+* y: A Tensor. Has the same type as "x" . \n
36+ 
37+* @see TileD()
38+ 
39+* @par Third-party framework compatibility
40+* Compatible with the TensorFlow operator Tile.
41+*/
42+REG_OP(Tile)
43+ .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16, DT_DOUBLE, DT_COMPLEX64, DT_COMPLEX128,
44+ DT_INT8, DT_UINT8, DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_INT64, DT_UINT64,
45+ DT_QINT8, DT_QUINT8, DT_QINT16, DT_QUINT16, DT_QINT32, DT_BF16, DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN}))
46+ .INPUT(multiples, TensorType::IndexNumberType())
47+ .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16, DT_DOUBLE, DT_COMPLEX64, DT_COMPLEX128,
48+ DT_INT8, DT_UINT8, DT_INT16, DT_UINT16, DT_INT32, DT_UINT32, DT_INT64, DT_UINT64,
49+ DT_QINT8, DT_QUINT8, DT_QINT16, DT_QUINT16, DT_QINT32, DT_BF16, DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN}))
50+ .OP_END_FACTORY_REG(Tile)
51+ 
52+} // namespace ge
53+ 
54+#endif
Dmath/tile/op_host/CMakeLists.txt+0-11
@@ -1,11 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# ----------------------------------------------------------------------------
10- 
11-add_modules_sources()
Amath/tile/op_host/arch35/tile_tiling_arch35.cpp+191-0
@@ -0,0 +1,191 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_tiling_arch35.cpp
13+ * \brief calc tiling for tile
14+ */
15+#include "tile_tiling_arch35.h"
16+#include "log/log.h"
17+#include "register/op_impl_registry.h"
18+#include "tiling/tiling_api.h"
19+#include "op_host/tiling_base.h"
20+#include "op_host/tiling_util.h"
21+#include "util/platform_util.h"
22+ 
23+namespace optiling {
24+constexpr size_t TILE_MAX_DIM_NUM = 0x8;
25+ 
26+template <typename T>
27+static std::string Shape2String(const T& shape)
28+{
29+ std::ostringstream oss;
30+ oss << "[";
31+ if (shape.GetDimNum() > 0) {
32+ for (size_t i = 0; i < shape.GetDimNum() - 1; ++i) {
33+ oss << shape.GetDim(i) << ", ";
34+ }
35+ oss << shape.GetDim(shape.GetDimNum() - 1);
36+ }
37+ oss << "]";
38+ return oss.str();
39+}
40+ 
41+static ge::graphStatus CheckTileRule(const gert::TilingContext* context, const gert::Shape& inShape,
42+ const gert::Shape& outShape)
43+{
44+ auto outDimNum = outShape.GetDimNum();
45+ OP_CHECK_IF(
46+ inShape.GetDimNum() != outDimNum,
47+ OP_LOGE(context->GetNodeName(), "The input shape dims are different with output's!"),
48+ return ge::GRAPH_FAILED);
49+ 
50+ for (size_t i = 0; i < outDimNum; i++) {
51+ if (outShape[i] < inShape[i] || outShape[i] % inShape[i] != 0) {
52+ return ge::GRAPH_FAILED;
53+ }
54+ }
55+ return ge::GRAPH_SUCCESS;
56+}
57+ 
58+static void ConvertTileAxis2BroadcastToAxis(gert::Shape& inShape, gert::Shape& outShape)
59+{
60+ gert::Shape newInShape;
61+ gert::Shape newOutShape;
62+ auto DimNum = inShape.GetDimNum();
63+ for (size_t i = 0; i < DimNum; i++) {
64+ newInShape.AppendDim(1);
65+ newInShape.AppendDim(inShape[i]);
66+ newOutShape.AppendDim(outShape[i] / inShape[i]);
CANN-robot
CANN-robotCANN-robot1月28日

逻辑运算与副作用: 在ConvertTileAxis2BroadcastToAxis函数中，计算outShape[i] / inShape[i]时，存在整数除零的潜在风险。虽然CheckTileRule中检查了outShape[i] % inShape[i] != 0，但如果inShape[i]为0，那么取模运算本身就会导致除零错误（在C++中，整数除零是未定义行为）。因此，需要确保inShape[i]不为0。

问题类型: 逻辑运算与副作用 文件路径: math/tile/op_host/arch35/tile_tiling_arch35.cpp 行号: 66 问题代码:

        newOutShape.AppendDim(outShape[i] / inShape[i]);

修改建议:

在CheckTileRule中，添加对inShape[i] == 0的检查。如果inShape[i]为0，那么outShape[i]也必须为0（因为outShape[i] % 0无意义），但通常形状维度为0是无效的。建议在CheckTileRule的循环中，先检查if (inShape[i] == 0) return ge::GRAPH_FAILED;。

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

likedislike
67+ newOutShape.AppendDim(inShape[i]);
68+ }
69+ inShape = newInShape;
70+ outShape = newOutShape;
71+}
72+ 
73+static ge::graphStatus GetShapeInfo(const gert::TilingContext* context, gert::Shape& inShape, gert::Shape& outShape)
74+{
75+ auto xStorage = context->GetInputShape(0);
76+ OP_CHECK_NULL_WITH_CONTEXT(context, xStorage);
77+ inShape = Ops::Math::OpTiling::EnsureNotScalar(xStorage->GetStorageShape());
78+ auto yStorage = context->GetOutputShape(0);
79+ OP_CHECK_NULL_WITH_CONTEXT(context, yStorage);
80+ outShape = Ops::Math::OpTiling::EnsureNotScalar(yStorage->GetStorageShape());
81+ 
82+ auto outDimNum = outShape.GetDimNum();
83+ OP_CHECK_IF(inShape.GetDimNum() > outDimNum,
84+ OP_LOGE(context->GetNodeName(),
85+ "The input shape has more dimensions than output shape!"),
86+ return ge::GRAPH_FAILED);
87+ OP_CHECK_IF(
88+ outDimNum > TILE_MAX_DIM_NUM,
89+ OP_LOGE(context->GetNodeName(), "Not support the dim num: %lu yet!", outDimNum),
90+ return ge::GRAPH_FAILED);
91+ brcto::AdjustShapesToSameDimNum(inShape, outDimNum);
92+ OP_CHECK_IF(inShape.GetShapeSize() == 0 || outShape.GetShapeSize() == 0,
93+ OP_LOGE(context->GetNodeName(), "The input or output shape is empty!"),
94+ return ge::GRAPH_FAILED);
95+ OP_CHECK_IF(CheckTileRule(context, inShape, outShape) != ge::GRAPH_SUCCESS,
96+ OP_LOGE(context->GetNodeName(),
97+ "The input and output shapes mismatch the broadcast rule!"),
98+ return ge::GRAPH_FAILED);
99+ ConvertTileAxis2BroadcastToAxis(inShape, outShape);
100+ 
101+ OP_LOGI(context->GetNodeName(), "The broadcastTo input and output after convertion is: %s and %s", Shape2String(inShape).c_str(),
102+ Shape2String(outShape).c_str());
103+ 
104+ OP_CHECK_IF(brcto::DeleteOneSizeAxis(context, inShape, outShape) != ge::GRAPH_SUCCESS,
105+ OP_LOGE(context->GetNodeName(), "Failed to delete one size axes!"),
106+ return ge::GRAPH_FAILED);
107+ OP_LOGI(context->GetNodeName(), "The reshaped input and output is: %s and %s", Shape2String(inShape).c_str(),
108+ Shape2String(outShape).c_str());
109+ 
110+ OP_CHECK_IF(brcto::MergeAxis(context, inShape, outShape) != ge::GRAPH_SUCCESS,
111+ OP_LOGE(context->GetNodeName(), "Failed to merge axes!"),
112+ return ge::GRAPH_FAILED);
113+ OP_LOGI(context->GetNodeName(), "The merged input and output is: %s and %s", Shape2String(inShape).c_str(),
114+ Shape2String(outShape).c_str());
115+ 
116+ return ge::GRAPH_SUCCESS;
117+}
118+ 
119+ge::graphStatus Tiling4TileAscendC(gert::TilingContext* context, const gert::Shape* inShapePtr,
120+ const gert::Shape* outShapePtr)
121+{
122+ OP_CHECK_NULL_WITH_CONTEXT(context, inShapePtr);
123+ OP_CHECK_NULL_WITH_CONTEXT(context, outShapePtr);
124+ brcto::BroadcastToTilingAscendC brcToTiling(context, inShapePtr, outShapePtr);
125+ 
126+ OP_CHECK_IF((brcToTiling.GetHardwareInfo<TileCompileInfo>() != ge::GRAPH_SUCCESS),
127+ OP_LOGE(context->GetNodeName(),
128+ "BroadcastToTilingAscendC failed to get hardware info."),
129+ return ge::GRAPH_FAILED);
130+ 
131+ return brcToTiling.DoTiling();
132+}
133+ 
134+static ge::graphStatus Tiling4Tile(gert::TilingContext* context) {
135+ auto compile_info = reinterpret_cast<const TileCompileInfo*>(context->GetCompileInfo());
136+ OP_CHECK_NULL_WITH_CONTEXT(context, compile_info);
137+ 
138+ gert::Shape inShape;
139+ gert::Shape outShape;
140+ 
141+ // convert tile input shape to broadcastTo input, then use broadcastTo template
142+ OP_CHECK_IF(GetShapeInfo(context, inShape, outShape) != ge::GRAPH_SUCCESS,
143+ OP_LOGE(context->GetNodeName(), "Get input or output shape was failed!"),
144+ return ge::GRAPH_FAILED);
145+ return Tiling4TileAscendC(context, &inShape, &outShape);
146+
147+}
148+ 
149+static ge::graphStatus TilingPrepare4Tile(gert::TilingParseContext* context) {
150+ OP_LOGD(context->GetNodeName(), "Enter TilingPrepare4Tile.");
151+ 
152+ auto compileInfo = context->GetCompiledInfo<TileCompileInfo>();
153+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
154+ auto platformInfo = context->GetPlatformInfo();
155+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
156+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
157+ 
158+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
159+ OP_CHECK_IF((compileInfo->coreNum <= 0),
160+ OP_LOGE(context->GetNodeName(), "The core num is negative."),
161+ return ge::GRAPH_FAILED);
162+ 
163+ uint64_t ubSize = 0;
164+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
165+ compileInfo->ubSize = static_cast<int64_t>(ubSize);
166+ OP_CHECK_IF((compileInfo->ubSize <= 0),
167+ OP_LOGE(context->GetNodeName(), "Failed to get ub size."),
168+ return ge::GRAPH_FAILED);
169+ 
170+ compileInfo->clSize = Ops::Base::GetCacheLineSize(context);
171+ OP_CHECK_IF((compileInfo->clSize <= 0),
172+ OP_LOGE(context->GetNodeName(), "Failed to get cache line size."),
173+ return ge::GRAPH_FAILED);
174+ 
175+ compileInfo->blockSize = Ops::Base::GetUbBlockSize(context);
176+ OP_CHECK_IF((compileInfo->blockSize <= 0),
177+ OP_LOGE(context->GetNodeName(), "Failed to get block size."),
178+ return ge::GRAPH_FAILED);
179+ 
180+ compileInfo->vRegSize = Ops::Base::GetVRegSize(context);
181+ OP_CHECK_IF((compileInfo->vRegSize <= 0),
182+ OP_LOGE(context->GetNodeName(), "Failed to get vReg size."),
183+ return ge::GRAPH_FAILED);
184+ 
185+ OP_LOGD(context->GetNodeName(), "Exit TilingPrepare4Tile.");
186+ return ge::GRAPH_SUCCESS;
187+}
188+ 
189+// register tiling interface of the Tile op.
190+IMPL_OP_OPTILING(Tile).Tiling(Tiling4Tile).TilingParse<TileCompileInfo>(TilingPrepare4Tile);
191+} // namespace optiling
Amath/tile/op_host/arch35/tile_tiling_arch35.h+33-0
@@ -0,0 +1,33 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_tiling_arch35.h
13+ * \brief head file of tile tiling
14+ */
15+ 
16+#ifndef OPS_BUILT_IN_OP_TILING_RUNTIME_TILE_H_
17+#define OPS_BUILT_IN_OP_TILING_RUNTIME_TILE_H_
18+#include "conversion/broadcast_to/op_host/arch35/broadcast_to_tiling_base.h"
19+#include "register/tilingdata_base.h"
20+ 
21+namespace optiling {
22+ 
23+REGISTER_TILING_DATA_CLASS(Tile, BroadcastToTilingData);
24+ 
25+struct TileCompileInfo {
26+ int64_t coreNum;
27+ int64_t ubSize;
28+ uint32_t clSize;
29+ uint32_t vRegSize;
30+ int64_t blockSize;
31+};
32+} // namespace optiling
33+#endif // OPS_BUILT_IN_OP_TILING_RUNTIME_TILE_H_
Amath/tile/op_host/config/ascend950/tile_binary.json+341-0
@@ -0,0 +1,341 @@
1+{
2+"op_type": "Tile",
3+"op_list": [
4+ {
5+ "bin_filename": "Tile_5ce316c9f935f0585864b2d372bcc322111",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "bool",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic",
17+ "dtype_match_mode": "DtypeByte"
18+ },
19+ {
20+ "name": "multiples",
21+ "index": 1,
22+ "dtype": "int32",
23+ "format": "ND",
24+ "paramType": "required",
25+ "shape": [
26+ -2
27+ ],
28+ "format_match_mode": "FormatAgnostic"
29+ }
30+ ],
31+ "outputs": [
32+ {
33+ "name": "y",
34+ "index": 0,
35+ "dtype": "bool",
36+ "format": "ND",
37+ "paramType": "required",
38+ "shape": [
39+ -2
40+ ],
41+ "format_match_mode": "FormatAgnostic",
42+ "dtype_match_mode": "DtypeByte"
43+ }
44+ ]
45+ },
46+ {
47+ "bin_filename": "Tile_be837fe498fad076d01c0ce8dfd3d01422",
48+ "inputs": [
49+ {
50+ "name": "x",
51+ "index": 0,
52+ "dtype": "float16",
53+ "format": "ND",
54+ "paramType": "required",
55+ "shape": [
56+ -2
57+ ],
58+ "format_match_mode": "FormatAgnostic",
59+ "dtype_match_mode": "DtypeByte"
60+ },
61+ {
62+ "name": "multiples",
63+ "index": 1,
64+ "dtype": "int32",
65+ "format": "ND",
66+ "paramType": "required",
67+ "shape": [
68+ -2
69+ ],
70+ "format_match_mode": "FormatAgnostic"
71+ }
72+ ],
73+ "outputs": [
74+ {
75+ "name": "y",
76+ "index": 0,
77+ "dtype": "float16",
78+ "format": "ND",
79+ "paramType": "required",
80+ "shape": [
81+ -2
82+ ],
83+ "format_match_mode": "FormatAgnostic",
84+ "dtype_match_mode": "DtypeByte"
85+ }
86+ ]
87+ },
88+ {
89+ "bin_filename": "Tile_7538f6ca8b4e71ac38750eb26a539cca33",
90+ "inputs": [
91+ {
92+ "name": "x",
93+ "index": 0,
94+ "dtype": "float32",
95+ "format": "ND",
96+ "paramType": "required",
97+ "shape": [
98+ -2
99+ ],
100+ "format_match_mode": "FormatAgnostic",
101+ "dtype_match_mode": "DtypeByte"
102+ },
103+ {
104+ "name": "multiples",
105+ "index": 1,
106+ "dtype": "int32",
107+ "format": "ND",
108+ "paramType": "required",
109+ "shape": [
110+ -2
111+ ],
112+ "format_match_mode": "FormatAgnostic"
113+ }
114+ ],
115+ "outputs": [
116+ {
117+ "name": "y",
118+ "index": 0,
119+ "dtype": "float32",
120+ "format": "ND",
121+ "paramType": "required",
122+ "shape": [
123+ -2
124+ ],
125+ "format_match_mode": "FormatAgnostic",
126+ "dtype_match_mode": "DtypeByte"
127+ }
128+ ]
129+ },
130+ {
131+ "bin_filename": "Tile_7538f6ca8b4e71ac38750eb26a539cca44",
132+ "inputs": [
133+ {
134+ "name": "x",
135+ "index": 0,
136+ "dtype": "int64",
137+ "format": "ND",
138+ "paramType": "required",
139+ "shape": [
140+ -2
141+ ],
142+ "format_match_mode": "FormatAgnostic",
143+ "dtype_match_mode": "DtypeByte"
144+ },
145+ {
146+ "name": "multiples",
147+ "index": 1,
148+ "dtype": "int32",
149+ "format": "ND",
150+ "paramType": "required",
151+ "shape": [
152+ -2
153+ ],
154+ "format_match_mode": "FormatAgnostic"
155+ }
156+ ],
157+ "outputs": [
158+ {
159+ "name": "y",
160+ "index": 0,
161+ "dtype": "int64",
162+ "format": "ND",
163+ "paramType": "required",
164+ "shape": [
165+ -2
166+ ],
167+ "format_match_mode": "FormatAgnostic",
168+ "dtype_match_mode": "DtypeByte"
169+ }
170+ ]
171+ },
172+ {
173+ "bin_filename": "Tile_5ce316c9f935f0585864b2d372bcc32255",
174+ "inputs": [
175+ {
176+ "name": "x",
177+ "index": 0,
178+ "dtype": "bool",
179+ "format": "ND",
180+ "paramType": "required",
181+ "shape": [
182+ -2
183+ ],
184+ "format_match_mode": "FormatAgnostic",
185+ "dtype_match_mode": "DtypeByte"
186+ },
187+ {
188+ "name": "multiples",
189+ "index": 1,
190+ "dtype": "int64",
191+ "format": "ND",
192+ "paramType": "required",
193+ "shape": [
194+ -2
195+ ],
196+ "format_match_mode": "FormatAgnostic"
197+ }
198+ ],
199+ "outputs": [
200+ {
201+ "name": "y",
202+ "index": 0,
203+ "dtype": "bool",
204+ "format": "ND",
205+ "paramType": "required",
206+ "shape": [
207+ -2
208+ ],
209+ "format_match_mode": "FormatAgnostic",
210+ "dtype_match_mode": "DtypeByte"
211+ }
212+ ]
213+ },
214+ {
215+ "bin_filename": "Tile_be837fe498fad076d01c0ce8dfd3d01466",
216+ "inputs": [
217+ {
218+ "name": "x",
219+ "index": 0,
220+ "dtype": "float16",
221+ "format": "ND",
222+ "paramType": "required",
223+ "shape": [
224+ -2
225+ ],
226+ "format_match_mode": "FormatAgnostic",
227+ "dtype_match_mode": "DtypeByte"
228+ },
229+ {
230+ "name": "multiples",
231+ "index": 1,
232+ "dtype": "int64",
233+ "format": "ND",
234+ "paramType": "required",
235+ "shape": [
236+ -2
237+ ],
238+ "format_match_mode": "FormatAgnostic"
239+ }
240+ ],
241+ "outputs": [
242+ {
243+ "name": "y",
244+ "index": 0,
245+ "dtype": "float16",
246+ "format": "ND",
247+ "paramType": "required",
248+ "shape": [
249+ -2
250+ ],
251+ "format_match_mode": "FormatAgnostic",
252+ "dtype_match_mode": "DtypeByte"
253+ }
254+ ]
255+ },
256+ {
257+ "bin_filename": "Tile_7538f6ca8b4e71ac38750eb26a539cca77",
258+ "inputs": [
259+ {
260+ "name": "x",
261+ "index": 0,
262+ "dtype": "float32",
263+ "format": "ND",
264+ "paramType": "required",
265+ "shape": [
266+ -2
267+ ],
268+ "format_match_mode": "FormatAgnostic",
269+ "dtype_match_mode": "DtypeByte"
270+ },
271+ {
272+ "name": "multiples",
273+ "index": 1,
274+ "dtype": "int64",
275+ "format": "ND",
276+ "paramType": "required",
277+ "shape": [
278+ -2
279+ ],
280+ "format_match_mode": "FormatAgnostic"
281+ }
282+ ],
283+ "outputs": [
284+ {
285+ "name": "y",
286+ "index": 0,
287+ "dtype": "float32",
288+ "format": "ND",
289+ "paramType": "required",
290+ "shape": [
291+ -2
292+ ],
293+ "format_match_mode": "FormatAgnostic",
294+ "dtype_match_mode": "DtypeByte"
295+ }
296+ ]
297+ },
298+ {
299+ "bin_filename": "Tile_7538f6ca8b4e71ac38750eb26a539cca88",
300+ "inputs": [
301+ {
302+ "name": "x",
303+ "index": 0,
304+ "dtype": "int64",
305+ "format": "ND",
306+ "paramType": "required",
307+ "shape": [
308+ -2
309+ ],
310+ "format_match_mode": "FormatAgnostic",
311+ "dtype_match_mode": "DtypeByte"
312+ },
313+ {
314+ "name": "multiples",
315+ "index": 1,
316+ "dtype": "int64",
317+ "format": "ND",
318+ "paramType": "required",
319+ "shape": [
320+ -2
321+ ],
322+ "format_match_mode": "FormatAgnostic"
323+ }
324+ ],
325+ "outputs": [
326+ {
327+ "name": "y",
328+ "index": 0,
329+ "dtype": "int64",
330+ "format": "ND",
331+ "paramType": "required",
332+ "shape": [
333+ -2
334+ ],
335+ "format_match_mode": "FormatAgnostic",
336+ "dtype_match_mode": "DtypeByte"
337+ }
338+ ]
339+ }
340+]
341+}
Amath/tile/op_host/config/ascend950/tile_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+[Tile]
13+default=0
Amath/tile/op_host/tile_def.cpp+77-0
@@ -0,0 +1,77 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_def.cpp
13+ * \brief op config of tile
14+ */
15+ 
16+#include <cstdint>
17+#include "register/op_def_registry.h"
18+ 
19+namespace ops
20+{
21+// remain same as aic-ascend910b-ops-info.ini, and add 3 more aicore type
22+static const std::vector<ge::DataType> xDType = {
23+ ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_INT32, ge::DT_INT64,
24+ ge::DT_BOOL, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8,
25+ ge::DT_INT16, ge::DT_UINT16, ge::DT_UINT32, ge::DT_UINT64,
26+ ge::DT_COMPLEX64, ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E4M3FN,
27+ ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_INT32, ge::DT_INT64,
28+ ge::DT_BOOL, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8,
29+ ge::DT_INT16, ge::DT_UINT16, ge::DT_UINT32, ge::DT_UINT64,
30+ ge::DT_COMPLEX64, ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E4M3FN};
31+static const std::vector<ge::Format> xFormat = {
32+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
33+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
34+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
35+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
36+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
37+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
38+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
39+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
40+static const std::vector<ge::DataType> constDType = {
41+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
42+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
43+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
44+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
45+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
46+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
47+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
48+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64};
49+ 
50+class Tile : public OpDef
51+{
52+public:
53+ explicit Tile(const char* name) : OpDef(name)
54+ {
55+ this->Input("x").ParamType(REQUIRED).DataType(xDType).Format(xFormat).UnknownShapeFormat(xFormat);
56+ this->Input("multiples")
57+ .ParamType(REQUIRED)
58+ .DataType(constDType)
59+ .Format(xFormat)
60+ .UnknownShapeFormat(xFormat)
61+ .ValueDepend(OPTIONAL);
62+ this->Output("y").ParamType(REQUIRED).DataType(xDType).Format(xFormat).UnknownShapeFormat(xFormat);
63+ 
64+ OpAICoreConfig aicoreConfig;
65+ aicoreConfig.DynamicCompileStaticFlag(true)
66+ .DynamicFormatFlag(false)
67+ .DynamicRankSupportFlag(true)
68+ .DynamicShapeSupportFlag(true)
69+ .NeedCheckSupportFlag(false)
70+ .ExtendCfgInfo("opFile.value", "tile_apt");
71+ this->AICore().AddConfig("ascend950", aicoreConfig);
72+ this->AICore().AddConfig("mc62cm12a", aicoreConfig);
73+ }
74+};
75+ 
76+OP_ADD(Tile);
77+} // namespace ops
Amath/tile/op_host/tile_infershape.cpp+101-0
@@ -0,0 +1,101 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_infershape.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "op_api/op_util.h"
zhanw_coding
zhanw_codingzhanw_coding1月31日

op_host 下不要引 op_api 下的头文件，而且此处也没用到

likedislike
17+#include "log/log.h"
18+#include "util/shape_util.h"
19+ 
20+using namespace ge;
21+using namespace Ops::Base;
22+ 
23+namespace ops {
24+constexpr size_t ATTR_INDEX_MULTIPLES = 0;
25+constexpr size_t TILE_IN_IDX = 0;
26+constexpr size_t TILE_OUT_IDX = 0;
27+static constexpr size_t MAXDIMNUM = 8;
28+constexpr size_t INPUT_INDEX_MULTIPLES = 1;
29+ 
30+template <typename T>
31+ge::graphStatus TileInferShapeCommon(gert::InferShapeContext* context, const T* multiples_data, size_t multiples_len) {
32+ auto in_shape = context->GetInputShape(TILE_IN_IDX);
33+ OP_CHECK_NULL_WITH_CONTEXT(context, in_shape);
34+ auto out_shape = context->GetOutputShape(TILE_OUT_IDX);
35+ OP_CHECK_NULL_WITH_CONTEXT(context, out_shape);
36+ auto in_shape_len = in_shape->GetDimNum();
37+ OP_CHECK_IF(multiples_len > MAXDIMNUM,
38+ OP_LOGE(context->GetNodeName(), "the tile multiples len is more than MaxDimNum 8"),
39+ return ge::GRAPH_FAILED);
40+ // align shape for input
41+ gert::Shape in_shape_new;
42+ if (in_shape_len < multiples_len) {
43+ OP_LOGI(context->GetNodeName(), "The tile multiples len is more than the input len.");
44+ int32_t len_diff = multiples_len - in_shape_len;
45+ for (int32_t i = 0; i < len_diff; i++) {
46+ in_shape_new.AppendDim(1);
47+ }
48+ for (size_t i = 0; i < in_shape_len; i++) {
49+ in_shape_new.AppendDim(in_shape->GetDim(i));
50+ }
51+ in_shape_len = multiples_len;
52+ } else {
53+ OP_LOGI(context->GetNodeName(), "The tile multiples len is less or equal than the input len.");
54+ in_shape_new = *in_shape;
55+ }
56+ // in shape == [], out shape = []
57+ if (in_shape_len == 0) {
58+ OP_LOGI(context->GetNodeName(), "input shape is [], output shape is [].");
59+ *out_shape = *in_shape;
60+ return GRAPH_SUCCESS;
61+ }
62+ // calculate output shape dim value
63+ out_shape->SetDimNum(in_shape_len);
64+ for (uint64_t i = 0; i < in_shape_len; i++) {
65+ if (in_shape_new[i] >= 0) {
66+ int32_t multiples_index = multiples_len - in_shape_len + i;
67+ out_shape->SetDim(i, in_shape_new[i] * (multiples_index >= 0 ? multiples_data[i] : 1));
68+ } else {
69+ OP_LOGE(context->GetNodeName(), "Runtime infershape illegal input dim:%lu, value is %ld", i, in_shape_new[i]);
70+ return ge::GRAPH_FAILED;
71+ }
72+ }
73+ return GRAPH_SUCCESS;
74+}
75+ 
76+static ge::graphStatus InferShape4Tile(gert::InferShapeContext* context) {
77+ auto x_shape = context->GetInputShape(0);
78+ OP_CHECK_NULL_WITH_CONTEXT(context, x_shape);
79+ auto multiples_tensor = context->GetInputTensor(INPUT_INDEX_MULTIPLES);
80+ OP_CHECK_NULL_WITH_CONTEXT(context, multiples_tensor);
81+ auto out_shape = context->GetOutputShape(0);
82+ OP_CHECK_NULL_WITH_CONTEXT(context, out_shape);
83+ auto multiples_size = static_cast<size_t>(multiples_tensor->GetShapeSize());
84+ OP_LOGD(context->GetNodeName(), "multiples_size is %zu", multiples_size);
85+ DataType data_type = multiples_tensor->GetDataType();
86+ OP_CHECK_IF((data_type != DT_INT32) && (data_type != DT_INT64),
87+ OP_LOGE(
88+ context->GetNodeName(),
89+ "multiples's dtype %s must be in (int32,int64)!", ToString(data_type).c_str()),
90+ return ge::GRAPH_FAILED);
91+ 
92+ if (data_type == DT_INT32) {
93+ const int32_t* multiples_data = multiples_tensor->GetData<int32_t>();
94+ return TileInferShapeCommon<int32_t>(context, multiples_data, multiples_size);
95+ }
96+ const int64_t* multiples_data = multiples_tensor->GetData<int64_t>();
97+ return TileInferShapeCommon<int64_t>(context, multiples_data, multiples_size);
98+}
99+ 
100+IMPL_OP_INFERSHAPE(Tile).InferShape(InferShape4Tile).InputsDataDependency({INPUT_INDEX_MULTIPLES});
101+} // namespace ops
Amath/tile/op_kernel/tile_apt.cpp+23-0
@@ -0,0 +1,23 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file tile_apt.cpp
13+ * \brief kernel of op tile
14+ */
15+ 
16+#include "../broadcast_to/arch35/broadcast_to.h"
17+ 
18+extern "C" __global__ __aicore__ void tile(GM_ADDR x, GM_ADDR shape, GM_ADDR y, GM_ADDR workspace,
19+ GM_ADDR tiling)
20+ 
21+{
22+ broadcast_to_impl(x, shape, y, workspace, tiling);
23+}
Dmath/tile/tests/CMakeLists.txt+0-16
@@ -1,16 +0,0 @@
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()
Dmath/tile/tests/ut/CMakeLists.txt+0-16
@@ -1,16 +0,0 @@
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()
Rmath/tile/tests/ut/op_host/op_api/test_aclnn_repeat.cpp→math/tile/tests/ut/op_api/test_aclnn_repeat.cpp+1-71
@@ -12,7 +12,7 @@
12 12 
13#include "gtest/gtest.h"13#include "gtest/gtest.h"
14#include "platform/platform_info.h"14#include "platform/platform_info.h"
15-#include "aclnn_repeat.h"15+#include "../../../op_api/aclnn_repeat.h"
16#include "op_api_ut_common/op_api_ut.h"16#include "op_api_ut_common/op_api_ut.h"
17#include "op_api_ut_common/scalar_desc.h"17#include "op_api_ut_common/scalar_desc.h"
18#include "op_api_ut_common/tensor_desc.h"18#include "op_api_ut_common/tensor_desc.h"
@@ -67,76 +67,6 @@ protected:
67///// 检查dtype /////67///// 检查dtype /////
68///////////////////////////////////////68///////////////////////////////////////
69 69 
70-TEST_F(l2_repeat_test, l2_repeat_test_dtype)
71-{
72- vector<aclDataType> ValidList = {ACL_INT8, ACL_INT16, ACL_INT32, ACL_INT64, ACL_UINT8,
73- ACL_DOUBLE, ACL_FLOAT16, ACL_FLOAT, ACL_BOOL, ACL_COMPLEX64,
74- ACL_COMPLEX128, ACL_UINT32, ACL_UINT64};
75- 
76- int64_t length = ValidList.size();
77- vector<int64_t> input_dim = {1, 2};
78- vector<int64_t> repeats_dim = {2, 3, 2, 3};
79- vector<int64_t> result_dim = {2, 3, 2, 6};
80- 
81- for (int i = 0; i < length; i++) {
82- auto inputDesc = TensorDesc(input_dim, ValidList[i], ACL_FORMAT_ND).ValueRange(-2, 2);
83- auto repeatsDesc = IntArrayDesc(repeats_dim);
84- auto outDesc = TensorDesc(result_dim, ValidList[i], ACL_FORMAT_ND).Precision(0.0001, 0.0001);
85- 
86- auto ut = OP_API_UT(aclnnRepeat, INPUT(inputDesc, repeatsDesc), OUTPUT(outDesc));
87- uint64_t workspaceSize = 0;
88- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
89- EXPECT_EQ(aclRet, ACL_SUCCESS);
90- // ut.TestPrecision();
91- }
92-}
93- 
94-TEST_F(l2_repeat_test, ascend910B2_test_dtype)
95-{
96- vector<aclDataType> ValidList = {ACL_INT8, ACL_INT16, ACL_INT32, ACL_INT64, ACL_UINT8, ACL_DOUBLE,
97- ACL_FLOAT16, ACL_FLOAT, ACL_BOOL, ACL_COMPLEX64, ACL_COMPLEX128};
98- 
99- int64_t length = ValidList.size();
100- vector<int64_t> input_dim = {1, 2};
101- vector<int64_t> repeats_dim = {2, 3, 2, 3};
102- vector<int64_t> result_dim = {2, 3, 2, 6};
103- 
104- for (int i = 0; i < length; i++) {
105- auto inputDesc = TensorDesc(input_dim, ValidList[i], ACL_FORMAT_ND).ValueRange(-2, 2);
106- auto repeatsDesc = IntArrayDesc(repeats_dim);
107- auto outDesc = TensorDesc(result_dim, ValidList[i], ACL_FORMAT_ND).Precision(0.0001, 0.0001);
108- 
109- auto ut = OP_API_UT(aclnnRepeat, INPUT(inputDesc, repeatsDesc), OUTPUT(outDesc));
110- uint64_t workspaceSize = 0;
111- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
112- EXPECT_EQ(aclRet, ACL_SUCCESS);
113- ut.TestPrecision();
114- }
115-}
116- 
117-TEST_F(l2_repeat_test, ascend310P1_test_dtype)
118-{
119- vector<aclDataType> ValidList = {ACL_INT8, ACL_INT16, ACL_INT32, ACL_INT64, ACL_UINT8, ACL_DOUBLE,
120- ACL_FLOAT16, ACL_FLOAT, ACL_BOOL, ACL_COMPLEX64, ACL_COMPLEX128};
121- 
122- int64_t length = ValidList.size();
123- vector<int64_t> input_dim = {1, 2};
124- vector<int64_t> repeats_dim = {2, 3, 2, 3};
125- vector<int64_t> result_dim = {2, 3, 2, 6};
126- 
127- for (int i = 0; i < length; i++) {
128- auto inputDesc = TensorDesc(input_dim, ValidList[i], ACL_FORMAT_ND).ValueRange(-2, 2);
129- auto repeatsDesc = IntArrayDesc(repeats_dim);
130- auto outDesc = TensorDesc(result_dim, ValidList[i], ACL_FORMAT_ND).Precision(0.0001, 0.0001);
131- 
132- auto ut = OP_API_UT(aclnnRepeat, INPUT(inputDesc, repeatsDesc), OUTPUT(outDesc));
133- uint64_t workspaceSize = 0;
134- aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
135- EXPECT_EQ(aclRet, ACL_SUCCESS);
136- ut.TestPrecision();
137- }
138-}
139- 
140TEST_F(l2_repeat_test, l2_repeat_test_dtype_bf16_error)70TEST_F(l2_repeat_test, l2_repeat_test_dtype_bf16_error)
141{71{
142 test_run({1, 2}, ACL_BF16, ACL_FORMAT_ND, {-10, 10}, {3, 2, 3}, {3, 2, 6}, ACL_BF16, ACL_FORMAT_ND);72 test_run({1, 2}, ACL_BF16, ACL_FORMAT_ND, {-10, 10}, {3, 2, 3}, {3, 2, 6}, ACL_BF16, ACL_FORMAT_ND);
Dmath/tile/tests/ut/op_host/CMakeLists.txt+0-19
@@ -1,19 +0,0 @@
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()
17- 
18-if(OP_API_UT OR (UT_TEST_ALL AND NOT AICPU_ONLY))
19-endif()
Amath/tile/tests/ut/op_host/arch35/test_tile_tiling_arch35.cpp+51-0
@@ -0,0 +1,51 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file test_tile_tiling_arch35.cpp
13+ * \brief tile tiling ut test
14+ */
15+ 
16+#include "../../../../op_host/arch35/tile_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 ge;
23+ 
24+class TileTiling : public testing::Test
25+{
26+protected:
27+ static void SetUpTestCase() { std::cout << "TileTiling SetUp" << std::endl; }
28+ 
29+ static void TearDownTestCase() { std::cout << "TileTiling TearDown" << std::endl; }
30+};
31+ 
32+TEST_F(TileTiling, TileTiling_001)
33+{
34+ optiling::TileCompileInfo compileInfo;
35+ compileInfo.coreNum = 64;
36+ compileInfo.ubSize = 245760; // 240 * 1024
37+ compileInfo.clSize = 256;
38+ compileInfo.vRegSize = 256;
39+ compileInfo.blockSize = 32;
40+ gert::StorageShape xShape = {{2, 3, 1, 1000, 1, 1}, {2, 3, 1, 1000, 1, 1}};
41+ gert::StorageShape shape = {{2, 3, 1, 1, 2, 2}, {2, 3, 1, 1, 2, 2}};
42+ gert::StorageShape yShape = {{4, 9, 1, 1000, 2, 2}, {4, 9, 1, 1000, 2, 2}};
43+ gert::TilingContextPara tilingContextPara(
44+ "Tile",
45+ {{ xShape, ge::DT_FLOAT, ge::FORMAT_ND }, { shape, ge::DT_INT32, ge::FORMAT_ND }},
46+ {{ yShape, ge::DT_FLOAT, ge::FORMAT_ND }},
47+ &compileInfo);
48+ uint64_t expectedTilingKey = 11000;
49+ std::vector<size_t> expectedWorkspaces = { 16777216 };
50+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectedTilingKey, expectedWorkspaces);
51+}
Dmath/tile/tests/ut/op_host/op_api/CMakeLists.txt+0-10
@@ -1,10 +0,0 @@
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- 
Amath/tile/tests/ut/op_host/test_tile_infershape.cpp+100-0
@@ -0,0 +1,100 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+#include <iostream>
13+#include "infershape_context_faker.h"
14+#include "infershape_case_executor.h"
15+ 
16+class TileTest : public testing::Test {
17+ protected:
18+ static void SetUpTestCase() {
19+ std::cout << "TileTest SetUp" << std::endl;
20+ }
21+ 
22+ static void TearDownTestCase() {
23+ std::cout << "TileTest TearDown" << std::endl;
24+ }
25+};
26+ 
27+TEST_F(TileTest, tile_infershape_test_0) {
28+ std::vector<int32_t> values = {8, 9};
29+ gert::InfershapeContextPara infershapeContextPara("Tile",
30+ {
31+ {{{2, 1}, {2, 1}}, ge::DT_FLOAT, ge::FORMAT_ND},
32+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND, true, values.data()}
33+ },
34+ {
35+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
36+ }
37+ );
38+ std::vector<std::vector<int64_t>> expectOutputShape = {{16, 9}};
39+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
40+}
41+ 
42+TEST_F(TileTest, tile_infershape_test_1) {
43+ std::vector<int32_t> values = {9};
44+ gert::InfershapeContextPara infershapeContextPara("Tile",
45+ {
46+ {{{1}, {1}}, ge::DT_FLOAT, ge::FORMAT_ND},
47+ {{{1}, {1}}, ge::DT_INT32, ge::FORMAT_ND, true, values.data()}
48+ },
49+ {
50+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
51+ }
52+ );
53+ std::vector<std::vector<int64_t>> expectOutputShape = {{9}};
54+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
55+}
56+ 
57+TEST_F(TileTest, tile_infershape_test_2) {
58+ std::vector<int32_t> values = {9};
59+ gert::InfershapeContextPara infershapeContextPara("Tile",
60+ {
61+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
62+ {{{1}, {1}}, ge::DT_INT32, ge::FORMAT_ND, true, values.data()}
63+ },
64+ {
65+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
66+ }
67+ );
68+ std::vector<std::vector<int64_t>> expectOutputShape = {{9}};
69+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
70+}
71+ 
72+TEST_F(TileTest, tile_infershape_test_3) {
73+ std::vector<int32_t> values = {};
74+ gert::InfershapeContextPara infershapeContextPara("Tile",
75+ {
76+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
77+ {{{0}, {0}}, ge::DT_INT32, ge::FORMAT_ND, true, values.data()}
78+ },
79+ {
80+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
81+ }
82+ );
83+ std::vector<std::vector<int64_t>> expectOutputShape = {{}};
84+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
85+}
86+ 
87+TEST_F(TileTest, tile_infershape_test_4) {
88+ std::vector<int32_t> values = {0};
89+ gert::InfershapeContextPara infershapeContextPara("Tile",
90+ {
91+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
92+ {{{1}, {1}}, ge::DT_INT32, ge::FORMAT_ND, true, values.data()}
93+ },
94+ {
95+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
96+ }
97+ );
98+ std::vector<std::vector<int64_t>> expectOutputShape = {{0}};
99+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
100+}
Mscripts/ci/ascend950/ops_math_operator_list.yaml+2-0
@@ -55,6 +55,8 @@ operator_group_1:
55 - grouped_bias_add_grad55 - grouped_bias_add_grad
56 - histogram_v256 - histogram_v2
57 - invert57 - invert
58+ - pad
59+ - tile
58 - is_close60 - is_close
59 - is_finite61 - is_finite
60 - is_inf62 - is_inf