已合并
pad_v2算子支持昇腾950 #1562
pengyiming7创建于 3月10日
pad_v2算子支持昇腾950 #1562
已合并
pengyiming7创建于 3月10日
39 个文件变更+3098-63
@@ -266,6 +266,7 @@ VC2@ops-math:
266 - ops/ops-math/conversion/transpose/266 - ops/ops-math/conversion/transpose/
267 - ops/ops-math/conversion/mirror_pad/267 - ops/ops-math/conversion/mirror_pad/
268 - ops/ops-math/conversion/pad/268 - ops/ops-math/conversion/pad/
269+ - ops/ops-math/conversion/pad_v2/
269 - ops/ops-math/conversion/pad_v3/270 - ops/ops-math/conversion/pad_v3/
270 - ops/ops-math/conversion/pad_v3_grad/271 - ops/ops-math/conversion/pad_v3_grad/
271 - ops/ops-math/conversion/fill/272 - ops/ops-math/conversion/fill/
@@ -36,7 +36,7 @@ static const std::initializer_list<op::DataType> REGBASE_AICORE_DTYPE_SUPPORT_LI
36 op::DataType::DT_INT64, op::DataType::DT_UINT64, op::DataType::DT_BF16,36 op::DataType::DT_INT64, op::DataType::DT_UINT64, op::DataType::DT_BF16,
37 op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_BOOL,37 op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_BOOL,
38 op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2, op::DataType::DT_FLOAT8_E4M3FN,38 op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2, op::DataType::DT_FLOAT8_E4M3FN,
39- op::DataType::DT_FLOAT8_E8M0};39+ op::DataType::DT_FLOAT8_E8M0, op::DataType::DT_FLOAT4_E2M1, op::DataType::DT_FLOAT4_E1M2};
40 40 
41inline static bool IsAiCoreSupport(const aclTensor* self)41inline static bool IsAiCoreSupport(const aclTensor* self)
42{42{
@@ -26,7 +26,7 @@ namespace ge {
26* @li x: A Tensor. Must be one of the following types: bfloat16, float16,26* @li x: A Tensor. Must be one of the following types: bfloat16, float16,
27* float32, double, int32, uint8, int16, int8, complex64, int64, qint8,27* float32, double, int32, uint8, int16, int8, complex64, int64, qint8,
28* quint8, qint32, qint16, quint16, uint16, complex128, uint32, uint64, bool,28* quint8, qint32, qint16, quint16, uint16, complex128, uint32, uint64, bool,
29-* hifloat8, float8_e5m2, float8_e4m3fn, float8_e8m0. Supported format list ["ND"].29+* hifloat8, float8_e5m2, float8_e4m3fn, float8_e8m0, float4_e2m1, float4_e1m2. Supported format list ["ND"].
30* @li paddings: A Tensor of type int32 or int64. Supported format list ["ND"]. \n30* @li paddings: A Tensor of type int32 or int64. Supported format list ["ND"]. \n
31 31 
32* @par Outputs:32* @par Outputs:
@@ -35,13 +35,18 @@ namespace ge {
35* @li Due to different architectures, the calculation results of this operator35* @li Due to different architectures, the calculation results of this operator
36* on NPU and CPU may be inconsistent. \n36* on NPU and CPU may be inconsistent. \n
37 37 
38+* @attention Constraints:
39+* If the type of x is float4_e2m1 or float4_e1m2, paddings values should be even number.
40+* If the type of x is hifloat8, float8_e5m2, float8_e4m3fn or float8_e8m0,
41+* paddings values should be non-negative integers. \n
42+ 
38* @par Third-party framework compatibility:43* @par Third-party framework compatibility:
39* Compatible with TensorFlow operator Pad.44* Compatible with TensorFlow operator Pad.
40*/45*/
41REG_OP(Pad)46REG_OP(Pad)
42- .INPUT(x, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0}))47+ .INPUT(x, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0, DT_FLOAT4_E2M1, DT_FLOAT4_E1M2}))
43 .INPUT(paddings, TensorType::IndexNumberType())48 .INPUT(paddings, TensorType::IndexNumberType())
44- .OUTPUT(y, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0}))49+ .OUTPUT(y, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0, DT_FLOAT4_E2M1, DT_FLOAT4_E1M2}))
45 .OP_END_FACTORY_REG(Pad)50 .OP_END_FACTORY_REG(Pad)
46 51 
47} // namespace ge52} // namespace ge
@@ -336,6 +336,90 @@
336 "format_match_mode": "FormatAgnostic"336 "format_match_mode": "FormatAgnostic"
337 }337 }
338 ]338 ]
339+ },
340+ {
341+ "bin_filename": "PadV1_0_5_Byte1",
342+ "inputs": [
343+ {
344+ "name": "x",
345+ "index": 0,
346+ "dtype": "float4_e2m1",
347+ "format": "ND",
348+ "paramType": "required",
349+ "shape": [
350+ -2
351+ ],
352+ "format_match_mode": "FormatAgnostic",
353+ "dtype_match_mode": "DtypeByte"
354+ },
355+ {
356+ "name": "paddings",
357+ "index": 1,
358+ "dtype": "int32",
359+ "format": "ND",
360+ "paramType": "required",
361+ "shape": [
362+ -2
363+ ],
364+ "format_match_mode": "FormatAgnostic"
365+ }
366+ ],
367+ "outputs": [
368+ {
369+ "name": "y",
370+ "index": 0,
371+ "dtype": "float4_e2m1",
372+ "format": "ND",
373+ "paramType": "required",
374+ "shape": [
375+ -2
376+ ],
377+ "format_match_mode": "FormatAgnostic",
378+ "dtype_match_mode": "DtypeByte"
379+ }
380+ ]
381+ },
382+ {
383+ "bin_filename": "PadV1_0_5_Byte2",
384+ "inputs": [
385+ {
386+ "name": "x",
387+ "index": 0,
388+ "dtype": "float4_e2m1",
389+ "format": "ND",
390+ "paramType": "required",
391+ "shape": [
392+ -2
393+ ],
394+ "format_match_mode": "FormatAgnostic",
395+ "dtype_match_mode": "DtypeByte"
396+ },
397+ {
398+ "name": "paddings",
399+ "index": 1,
400+ "dtype": "int64",
401+ "format": "ND",
402+ "paramType": "required",
403+ "shape": [
404+ -2
405+ ],
406+ "format_match_mode": "FormatAgnostic"
407+ }
408+ ],
409+ "outputs": [
410+ {
411+ "name": "y",
412+ "index": 0,
413+ "dtype": "float4_e2m1",
414+ "format": "ND",
415+ "paramType": "required",
416+ "shape": [
417+ -2
418+ ],
419+ "format_match_mode": "FormatAgnostic",
420+ "dtype_match_mode": "DtypeByte"
421+ }
422+ ]
339 }423 }
340 ]424 ]
341}425}
@@ -20,20 +20,22 @@ static const std::vector<ge::Format> format = {
20 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,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, ge::FORMAT_ND, ge::FORMAT_ND,22 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
23- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};23+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
24+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
24static const std::vector<ge::DataType> valueDataType = {25static 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,26 ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32, ge::DT_UINT32,
26 ge::DT_INT64, ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_DOUBLE,27 ge::DT_INT64, ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_DOUBLE,
27- ge::DT_BOOL, ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E4M3FN, ge::DT_INT8,28+ ge::DT_BOOL, ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT4_E2M1,
28- ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,29+ ge::DT_FLOAT4_E1M2, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
29 ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_DOUBLE, ge::DT_BOOL,30 ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_DOUBLE, ge::DT_BOOL,
30- ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E4M3FN};31+ ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E1M2};
31 32 
32static const std::vector<ge::DataType> padDataType = {33static const std::vector<ge::DataType> padDataType = {
33 ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,34 ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
34 ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,35 ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
35- ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,36+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT32,
36- ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,37+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
38+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
37 ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32};39 ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32};
38} // namespace40} // namespace
39 41 
@@ -17,7 +17,6 @@
17 17 
18using namespace PadV3;18using namespace PadV3;
19 19 
20-#define CONSTANT_SLICE_BRANCH 10000
21#define CONSTANT_SIMT_BRANCH 2000020#define CONSTANT_SIMT_BRANCH 20000
22#define CONSTANT_SIMT_BIG_SIZE_BRANCH 2000121#define CONSTANT_SIMT_BIG_SIZE_BRANCH 20001
23#define CONSTANT_CUT_LAST_DIM_BRANCH 3001022#define CONSTANT_CUT_LAST_DIM_BRANCH 30010
@@ -0,0 +1,16 @@
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+set(SUPPORT_COMPUTE_UNIT "ascend950")
13+# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14+set(SUPPORT_TILING_DIR "arch35")
15+# PadV2 复用 PadV3 的实现
16+add_all_modules_sources(OPTYPE pad_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES pad_v3)
@@ -0,0 +1,112 @@
1+# PadV2
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| ---- | :----:|
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | √ |
11+| <term>Atlas 推理系列产品</term> | √ |
12+| <term>Atlas 训练系列产品</term> | √ |
13+ 
14+## 功能说明
15+ 
16+- **算子功能**:对输入 tensor 进行指定值的常量填充。
17+- **示例**
18+ 
19+ ```
20+ 输入 tensor([[0,1,2]])
21+ paddings([[2,2]])
22+ constant_values(0)
23+
24+ 输出为([[0,0,0,1,2,0,0]])
25+ ```
26+ 
27+## 参数说明
28+ 
29+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
30+ <col style="width: 100px">
31+ <col style="width: 150px">
32+ <col style="width: 280px">
33+ <col style="width: 330px">
34+ <col style="width: 120px">
35+ </colgroup>
36+ <thead>
37+ <tr>
38+ <th>参数名</th>
39+ <th>输入/输出/属性</th>
40+ <th>描述</th>
41+ <th>数据类型</th>
42+ <th>数据格式</th>
43+ </tr></thead>
44+ <tbody>
45+ <tr>
46+ <td>x</td>
47+ <td>输入</td>
48+ <td>待进行填充的原始 tensor。</td>
49+ <td>FLOAT、FLOAT16、INT32、INT64、INT16、INT8、UINT8、UINT16、UINT32、UINT64、BOOL、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。</td>
50+ <td>ND</td>
51+ </tr>
52+ <tr>
53+ <td>paddings</td>
54+ <td>输入</td>
55+ <td>填充配置,shape=[N, 2],其中 N 为 x 的维度数。</td>
56+ <td>INT32、INT64</td>
57+ <td>ND</td>
58+ </tr>
59+ <tr>
60+ <td>constant_values</td>
61+ <td>输入</td>
62+ <td>填充常量值,标量 tensor。</td>
63+ <td>FLOAT、FLOAT16、INT32、INT64、INT16、INT8、UINT8、UINT16、UINT32、UINT64、BOOL、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。</td>
64+ <td>ND</td>
65+ </tr>
66+ <tr>
67+ <td>y</td>
68+ <td>输出</td>
69+ <td>填充后的 tensor。</td>
70+ <td>FLOAT、FLOAT16、INT32、INT64、INT16、INT8、UINT8、UINT16、UINT32、UINT64、BOOL、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。</td>
71+ <td>ND</td>
72+ </tr>
73+ </tbody></table>
74+ 
75+ - <term>Atlas 训练系列产品</term><term>Atlas 推理系列产品</term><term>Atlas 200I/500 A2 推理产品</term>:数据类型不支持BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。
76+ - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term><term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>:数据类型不支持 HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。
77+ 
78+### 输出 shape 计算
79+ 
80+```
81+y.shape[d] = x.shape[d] + paddings[d][0] + paddings[d][1]
82+```
83+ 
84+## 约束说明
85+ 
86+1. **维度约束**
87+ - x 的维度数必须在 [1, 8] 范围内
88+ - paddings 的第一维必须等于 x 的维度数
89+ - paddings 的第二维必须等于 2
90+ 
91+2. **数据类型约束**
92+ - x、constant_values、y 必须使用相同的数据类型
93+ - paddings 必须使用 INT32 或 INT64
94+ 
95+3. **填充约束**
96+ - 负填充(slice)时,输出 shape = x.shape[d] + left + right >= 0
97+ 
98+4. **paddings 参数约束**
99+ - paddings 的形状必须为 [rank, 2],其中 rank 为输入 x 的维度数(1~8)
100+ - 每一行 [left, right] 表示对应维度的填充数量
101+ - left: 在该维度的开头填充的元素数
102+ - right: 在该维度的末尾填充的元素数
103+ - paddings的值可以为:
104+ - 正数:表示填充
105+ - 负数:表示 slice(裁剪)
106+ - 零:表示不填充
107+ 
108+## 调用说明
109+ 
110+| 调用方式  | 调用样例                        | 说明                              |
111+| ------------| ---------------------------------------------------------| ----------------------------------------------------------------|
112+| aclnn 调用 | [test_geir_pad_v2.cpp](./examples/test_geir_pad_v2.cpp) | 通过 [算子IR](op_graph/pad_v2_proto.h) 接口方式调用 PadV2 算子 |
@@ -0,0 +1,420 @@
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_geir_pad_v2.cpp
13+ * @brief PadV2算子的GE IR示例代码
14+ *
15+ * 本示例演示如何使用GE IR接口构建和执行PadV2算子:
16+ * 1. 创建输入tensor(2D tensor)
17+ * 2. 设置paddings参数(例如:[[1,1], [2,2]])
18+ * 3. 设置constant_values(例如:0.0
19+ * 4. 执行PadV2算子
20+ * 5. 验证输出结果
21+ */
22+ 
23+#include <iostream>
24+#include <fstream>
25+#include <string.h>
26+#include <stdint.h>
27+#include <vector>
28+#include <string>
29+#include <map>
30+#include "assert.h"
31+ 
32+#include "graph.h"
33+#include "types.h"
34+#include "tensor.h"
35+#include "ge_error_codes.h"
36+#include "ge_api_types.h"
37+#include "ge_api.h"
38+#include "array_ops.h"
39+#include "ge_ir_build.h"
40+ 
41+#include "experiment_ops.h"
42+#include "nn_other.h"
43+#include "../op_graph/pad_v2_proto.h"
44+ 
45+#define FAILED -1
46+#define SUCCESS 0
47+ 
48+using namespace ge;
49+using std::map;
50+using std::string;
51+using std::vector;
52+ 
53+// 宏:添加输入tensor
54+#define ADD_INPUT(inputIndex, inputName, inputDtype, inputShape) \
55+ do { \
56+ std::string name##inputIndex = "placeholder" + std::to_string(inputIndex); \
57+ auto placeholder##inputIndex = op::Data(name##inputIndex.c_str()).set_attr_index(0); \
58+ TensorDesc placeholder##inputIndex##_desc = TensorDesc(ge::Shape(inputShape), FORMAT_ND, inputDtype); \
59+ placeholder##inputIndex##_desc.SetPlacement(ge::kPlacementHost); \
60+ placeholder##inputIndex##_desc.SetFormat(FORMAT_ND); \
61+ Tensor tensor_placeholder##inputIndex; \
62+ ret = GenDataFloat32(inputShape, tensor_placeholder##inputIndex, placeholder##inputIndex##_desc); \
63+ if (ret != SUCCESS) { \
64+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
65+ return FAILED; \
66+ } \
67+ placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \
68+ graph.AddOp(placeholder##inputIndex); \
69+ input.push_back(tensor_placeholder##inputIndex); \
70+ padv2.set_input_##inputName(placeholder##inputIndex); \
71+ inputs.push_back(placeholder##inputIndex); \
72+ } while (0)
73+ 
74+// 宏:添加int32类型的输入tensor(用于paddings)
75+#define ADD_INT32_INPUT(inputIndex, inputName, inputShape, dataVec) \
76+ do { \
77+ std::string name##inputIndex = "placeholder" + std::to_string(inputIndex); \
78+ auto placeholder##inputIndex = op::Data(name##inputIndex.c_str()).set_attr_index(0); \
79+ TensorDesc placeholder##inputIndex##_desc = TensorDesc(ge::Shape(inputShape), FORMAT_ND, DT_INT32); \
80+ placeholder##inputIndex##_desc.SetPlacement(ge::kPlacementHost); \
81+ placeholder##inputIndex##_desc.SetFormat(FORMAT_ND); \
82+ Tensor tensor_placeholder##inputIndex; \
83+ ret = GenInt32Data(inputShape, tensor_placeholder##inputIndex, placeholder##inputIndex##_desc, dataVec); \
84+ if (ret != SUCCESS) { \
85+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
86+ return FAILED; \
87+ } \
88+ placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \
89+ graph.AddOp(placeholder##inputIndex); \
90+ input.push_back(tensor_placeholder##inputIndex); \
91+ padv2.set_input_##inputName(placeholder##inputIndex); \
92+ inputs.push_back(placeholder##inputIndex); \
93+ } while (0)
94+ 
95+// 宏:添加scalar输入tensor(用于constant_values)
96+#define ADD_SCALAR_INPUT(inputIndex, inputName, inputDtype, value) \
97+ do { \
98+ std::string name##inputIndex = "placeholder" + std::to_string(inputIndex); \
99+ auto placeholder##inputIndex = op::Data(name##inputIndex.c_str()).set_attr_index(0); \
100+ TensorDesc placeholder##inputIndex##_desc = TensorDesc(ge::Shape({1}), FORMAT_ND, inputDtype); \
101+ placeholder##inputIndex##_desc.SetPlacement(ge::kPlacementHost); \
102+ placeholder##inputIndex##_desc.SetFormat(FORMAT_ND); \
103+ Tensor tensor_placeholder##inputIndex; \
104+ ret = GenScalarData(inputDtype, tensor_placeholder##inputIndex, placeholder##inputIndex##_desc, value); \
105+ if (ret != SUCCESS) { \
106+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
107+ return FAILED; \
108+ } \
109+ placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \
110+ graph.AddOp(placeholder##inputIndex); \
111+ input.push_back(tensor_placeholder##inputIndex); \
112+ padv2.set_input_##inputName(placeholder##inputIndex); \
113+ inputs.push_back(placeholder##inputIndex); \
114+ } while (0)
115+ 
116+// 宏:添加输出tensor
117+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
118+ do { \
119+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
120+ padv2.update_output_desc_##outputName(outputName##outputIndex##_desc); \
121+ } while (0)
122+ 
123+// 宏:添加属性
124+#define ADD_ATTR(attrName, attrValue) padv2.set_attr_##attrName(attrValue)
125+ 
126+// 宏:打印日志
127+#define LOG_PRINT(message, ...) \
128+ do { \
129+ printf(message, ##__VA_ARGS__); \
130+ } while (0)
131+ 
132+// 获取当前时间字符串
133+string GetTime()
134+{
135+ time_t timep;
136+ time(&timep);
137+ char tmp[64];
138+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
139+ return tmp;
140+}
141+ 
142+// 获取数据类型的大小(字节数)
143+uint32_t GetDataTypeSize(DataType dt)
144+{
145+ if (dt == ge::DT_FLOAT)
146+ return 4;
147+ if (dt == ge::DT_FLOAT16)
148+ return 2;
149+ if (dt == ge::DT_BF16)
150+ return 2;
151+ if (dt == ge::DT_INT32)
152+ return 4;
153+ if (dt == ge::DT_INT64)
154+ return 8;
155+ return 4;
156+}
157+ 
158+// 生成float32数据
159+int32_t GenDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc)
160+{
161+ input_tensor_desc.SetRealDimCnt(shapes.size());
162+ size_t size = 1;
163+ for (uint32_t i = 0; i < shapes.size(); i++) {
164+ size *= shapes[i];
165+ }
166+ uint32_t data_len = size * 4;
167+ float* pData = new (std::nothrow) float[size];
168+ 
169+ // 生成测试数据:0, 1, 2, 3, ...
170+ for (size_t i = 0; i < size; ++i) {
171+ pData[i] = static_cast<float>(i);
172+ }
173+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
174+ return SUCCESS;
175+}
176+ 
177+// 生成int32数据(用于paddings)
178+int32_t GenInt32Data(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc,
179+ const std::vector<int32_t>& dataVec)
180+{
181+ input_tensor_desc.SetRealDimCnt(shapes.size());
182+ size_t size = 1;
183+ for (uint32_t i = 0; i < shapes.size(); i++) {
184+ size *= shapes[i];
185+ }
186+ uint32_t data_len = size * 4;
187+ int32_t* pData = new (std::nothrow) int32_t[size];
188+ 
189+ // 使用传入的数据
190+ for (size_t i = 0; i < size && i < dataVec.size(); ++i) {
191+ pData[i] = dataVec[i];
192+ }
193+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
194+ return SUCCESS;
195+}
196+ 
197+// 生成scalar数据(用于constant_values)
198+int32_t GenScalarData(DataType dtype, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value)
199+{
200+ input_tensor_desc.SetRealDimCnt(1);
201+ uint32_t data_len = GetDataTypeSize(dtype);
202+
203+ if (dtype == ge::DT_FLOAT) {
204+ float* pData = new (std::nothrow) float[1];
205+ pData[0] = value;
206+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
207+ } else {
208+ // 其他类型暂不支持
209+ return FAILED;
210+ }
211+ return SUCCESS;
212+}
213+ 
214+// 写数据到文件
215+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
216+{
217+ FILE* fp = fopen(bin_file.c_str(), "wb");
218+ if (fp == nullptr) {
219+ return FAILED;
220+ }
221+ size_t written = fwrite(inputData, 1, data_size, fp);
222+ fclose(fp);
223+ if (written != data_size) {
224+ return FAILED;
225+ }
226+ return SUCCESS;
227+}
228+ 
229+/**
230+ * @brief 创建包含PadV2算子的计算图
231+ *
232+ * 示例场景:
233+ * - 输入x: shape=[3, 3],数据为[0,1,2,3,4,5,6,7,8]
234+ * - paddings: [[1,1], [2,2]],表示在第0维前后各填充1,在第1维前后各填充2
235+ * - constant_values: 0.0
236+ * - 输出y: shape=[5, 7],填充后的结果
237+ */
238+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
239+ std::vector<Operator> &outputs, Graph &graph)
240+{
241+ Status ret = SUCCESS;
242+
243+ // 创建PadV2算子
244+ auto padv2 = op::PadV2("test_geir_pad_v2");
245+
246+ // 定义输入shape
247+ std::vector<int64_t> xShape = {3, 3}; // 输入tensor的shape
248+ std::vector<int64_t> paddingsShape = {2, 2}; // paddings的shape: [N, 2],N为x的rank
249+ std::vector<int64_t> constantValuesShape = {1}; // constant_values的shape(scalar)
250+
251+ // 定义输出shape(根据paddings计算)
252+ // x.shape = [3, 3], paddings = [[1,1], [2,2]]
253+ // y.shape = [3+1+1, 3+2+2] = [5, 7]
254+ std::vector<int64_t> yShape = {5, 7};
255+ 
256+ // 添加输入x
257+ ADD_INPUT(1, x, inDtype, xShape);
258+
259+ // 添加输入paddings(int32类型)
260+ // paddings = [[1,1], [2,2]],表示在第0维前后各填充1,在第1维前后各填充2
261+ std::vector<int32_t> paddingsData = {1, 1, 2, 2};
262+ ADD_INT32_INPUT(2, paddings, paddingsShape, paddingsData);
263+
264+ // 添加输入constant_values(可选输入)
265+ ADD_SCALAR_INPUT(3, constant_values, inDtype, 0.0f);
266+ 
267+ // 添加属性mode
268+ ADD_ATTR(mode, "constant");
269+ 
270+ // 添加输出y
271+ ADD_OUTPUT(1, y, inDtype, yShape);
272+ 
273+ outputs.push_back(padv2);
274+
275+ return SUCCESS;
276+}
277+ 
278+int main(int argc, char *argv[])
279+{
280+ const char *graph_name = "tc_ge_irrun_test";
281+ Graph graph(graph_name);
282+ std::vector<ge::Tensor> input;
283+ 
284+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
285+
286+ // 初始化GE环境
287+ std::map<AscendString, AscendString> global_options = {
288+ {"ge.exec.deviceId", "0"},
289+ {"ge.graphRunMode", "1"}
290+ };
291+ Status ret = ge::GEInitialize(global_options);
292+ if (ret != SUCCESS) {
293+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
294+ return FAILED;
295+ }
296+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
297+ 
298+ std::vector<Operator> inputs{};
299+ std::vector<Operator> outputs{};
300+ 
301+ // 设置输入数据类型
302+ DataType inDtype = DT_FLOAT;
303+ printf("Input dtype: %d\n", inDtype);
304+ 
305+ // 创建计算图
306+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
307+ if (ret != SUCCESS) {
308+ printf("%s - ERROR - [XIR]: Create graph failed\n", GetTime().c_str());
309+ return FAILED;
310+ }
311+ 
312+ // 设置图的输入输出
313+ if (!inputs.empty() && !outputs.empty()) {
314+ graph.SetInputs(inputs).SetOutputs(outputs);
315+ }
316+ 
317+ // 创建会话
318+ std::map<AscendString, AscendString> build_options = {};
319+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
320+ ge::Session *session = new Session(build_options);
321+ 
322+ if (session == nullptr) {
323+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
324+ return FAILED;
325+ }
326+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
327+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
328+ 
329+ // 添加计算图到会话
330+ std::map<AscendString, AscendString> graph_options = {};
331+ uint32_t graph_id = 0;
332+ ret = session->AddGraph(graph_id, graph, graph_options);
333+ 
334+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
335+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
336+
337+ // dump图结构到文件
338+ std::string file_path = "./dump";
339+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
340+
341+ // 执行图
342+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
343+ std::vector<ge::Tensor> output;
344+ ret = session->RunGraph(graph_id, input, output);
345+ if (ret != SUCCESS) {
346+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
347+ delete session;
348+ GEFinalize();
349+ return FAILED;
350+ }
351+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
352+ 
353+ // 保存输入数据到文件
354+ int input_num = input.size();
355+ for (int i = 0; i < input_num; i++) {
356+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
357+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
358+ uint8_t *input_data_i = input[i].GetData();
359+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
360+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
361+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
362+ WriteDataToFile((const char *)input_file.c_str(), data_size, input_data_i);
363+
364+ // 打印输入数据(仅float类型)
365+ if (input[i].GetTensorDesc().GetDataType() == ge::DT_FLOAT) {
366+ float *inputData = (float*)input_data_i;
367+ std::cout << "Input " << i << " data:" << std::endl;
368+ for (int64_t j = 0; j < std::min(input_shape, (int64_t)10); j++) {
369+ LOG_PRINT(" input[%ld] = %f\n", j, inputData[j]);
370+ }
371+ }
372+ }
373+ 
374+ // 保存输出数据到文件
375+ int output_num = output.size();
376+ for (int i = 0; i < output_num; i++) {
377+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
378+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
379+ uint8_t *output_data_i = output[i].GetData();
380+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
381+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
382+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
383+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
384+
385+ // 打印输出数据(仅float类型)
386+ if (output[i].GetTensorDesc().GetDataType() == ge::DT_FLOAT) {
387+ float *resultData = (float*)output_data_i;
388+ std::cout << "Output " << i << " data:" << std::endl;
389+ // 打印前10个和后10个元素
390+ for (int64_t j = 0; j < std::min(output_shape, (int64_t)10); j++) {
391+ LOG_PRINT(" result[%ld] = %f\n", j, resultData[j]);
392+ }
393+ if (output_shape > 20) {
394+ LOG_PRINT(" ... (showing first 10 elements)\n");
395+ for (int64_t j = output_shape - 10; j < output_shape; j++) {
396+ LOG_PRINT(" result[%ld] = %f\n", j, resultData[j]);
397+ }
398+ }
399+ }
400+ }
401+ 
402+ // 获取错误和警告信息
403+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
404+ std::string error_str(error_msg.GetString());
405+ std::cout << "Error message: " << error_str << std::endl;
406+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
407+ std::string warning_str(warning_msg.GetString());
408+ std::cout << "Warning message: " << warning_str << std::endl;
409+
410+ // 清理资源
411+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
412+ ret = ge::GEFinalize();
413+ if (ret != SUCCESS) {
414+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
415+ return FAILED;
416+ }
417+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
418+
419+ return SUCCESS;
420+}
@@ -0,0 +1,65 @@
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_v2_proto.h
13+ * \brief PadV2算子的图优化原型定义
14+ */
15+#ifndef OP_PROTO_PAD_V2_PROTO_H_
16+#define OP_PROTO_PAD_V2_PROTO_H_
17+ 
18+#include "graph/operator_reg.h"
19+ 
20+namespace ge
21+{
22+/**
23+* @brief Pads a tensor with constant values.
24+ 
25+* @par Inputs:
26+* Three inputs, including:
27+* @li x: A Tensor. Must be one of the following types: float16, bfloat16,
28+* float32, double, int32, uint8, int16, int8, complex64, int64,
29+* qint8, quint8, qint32, qint16, quint16, uint16, complex128, uint32, uint64, bool,
30+* hifloat8, float8_e5m2, float8_e4m3fn, float8_e8m0, float4_e2m1, float4_e1m2.
31+* Supported format list: ["ND"].
32+*
33+* @li paddings: A Tensor of type int32 or int64.
34+* The shape of paddings must be [N, 2], where N is the rank of x.
35+* paddings[i][0] specifies the number of padding elements to add before x in dimension i.
36+* paddings[i][1] specifies the number of padding elements to add after x in dimension i.
37+* Supported format list: ["ND"].
38+*
39+* @li constant_values: An optional Tensor, dtype same as "x".
40+* Used only in "constant" mode to specify the value to fill in the padded regions.
41+* If not provided, default value is 0. \n
42+ 
43+* @par Outputs:
44+* y: A Tensor of the same type as "x".
45+* y.shape[i] = x.shape[i] + paddings[i][0] + paddings[i][1], where y.shape[i] >= 0.
46+* Supported format list: ["ND"]. \n
47+ 
48+* @attention Constraints:
49+* If the type of x is float4_e2m1 or float4_e1m2, paddings values should be even number.
50+* If the type of x is hifloat8, float8_e5m2, float8_e4m3fn or float8_e8m0,
51+* paddings values should be non-negative integers.
52+* The constant_values should be a scalar tensor with the same dtype as x. \n
53+ 
54+* @par Third-party framework compatibility:
55+* Compatible with TensorFlow operator PadV2.
56+*/
57+REG_OP(PadV2)
58+ .INPUT(x, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0, DT_FLOAT4_E2M1, DT_FLOAT4_E1M2}))
59+ .INPUT(paddings, TensorType::IndexNumberType())
60+ .INPUT(constant_values, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0, DT_FLOAT4_E2M1, DT_FLOAT4_E1M2}))
61+ .OUTPUT(y, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0, DT_FLOAT4_E2M1, DT_FLOAT4_E1M2}))
62+ .OP_END_FACTORY_REG(PadV2)
63+ 
64+} // namespace ge
65+#endif // OP_PROTO_PAD_V2_PROTO_H_
@@ -0,0 +1,35 @@
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_v2_tiling_arch35.cpp
13+ * \brief PadV2 tiling - 复用 PadV3 实现
14+ */
15+#include "pad_v2_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 Tiling4PadV2(gert::TilingContext* context) {
22+ // 直接使用 PadV3 的 Tiling 类
23+ PadACTiling tilingObject(context);
24+ return tilingObject.DoTiling();
25+}
26+ 
27+static ge::graphStatus TilingPrepare4PadV2(gert::TilingParseContext* context) {
28+ OP_LOGD(context->GetNodeName(), "TilingPrepare4PadV2 entering.");
29+ return ge::GRAPH_SUCCESS;
30+}
31+ 
32+// 注册 PadV2 的 tiling 接口
33+IMPL_OP_OPTILING(PadV2).Tiling(Tiling4PadV2).TilingParse<PadV3CompileInfo>(TilingPrepare4PadV2);
34+ 
35+} // namespace optiling
@@ -0,0 +1,25 @@
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_v2_tiling_arch35.h
13+ * \brief PadV2 tiling header - 复用 PadV3 实现
14+ */
15+#ifndef OP_HOST_ARCH35_PAD_V2_TILING_ARCH35_H_
16+#define OP_HOST_ARCH35_PAD_V2_TILING_ARCH35_H_
17+ 
18+#include "conversion/pad_v3/op_host/arch35/pad_v3_tiling_arch35.h"
19+ 
20+// PadV2 直接复用 PadV3 的 Tiling 实现,无需额外类型定义
21+// Tiling 实现使用:PadACTiling
22+// 编译信息使用:PadV3CompileInfo
23+// 参考:op_host/arch35/pad_v2_tiling_arch35.cpp
24+ 
25+#endif // OP_HOST_ARCH35_PAD_V2_TILING_ARCH35_H_
@@ -0,0 +1,545 @@
1+{
2+ "op_type": "PadV2",
3+ "op_list": [
4+ {
5+ "bin_filename": "PadV2_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+ "format_match_mode": "FormatAgnostic",
17+ "dtype_match_mode": "DtypeByte"
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+ "name": "constant_values",
32+ "index": 2,
33+ "dtype": "int8",
34+ "format": "ND",
35+ "paramType": "required",
36+ "shape": [
37+ -2
38+ ],
39+ "format_match_mode": "FormatAgnostic",
40+ "dtype_match_mode": "DtypeByte"
41+ }
42+ ],
43+ "outputs": [
44+ {
45+ "name": "y",
46+ "index": 0,
47+ "dtype": "int8",
48+ "format": "ND",
49+ "paramType": "required",
50+ "shape": [
51+ -2
52+ ],
53+ "format_match_mode": "FormatAgnostic",
54+ "dtype_match_mode": "DtypeByte"
55+ }
56+ ]
57+ },
58+ {
59+ "bin_filename": "PadV2_int8_int64",
60+ "inputs": [
61+ {
62+ "name": "x",
63+ "index": 0,
64+ "dtype": "int8",
65+ "format": "ND",
66+ "paramType": "required",
67+ "shape": [
68+ -2
69+ ],
70+ "format_match_mode": "FormatAgnostic",
71+ "dtype_match_mode": "DtypeByte"
72+ },
73+ {
74+ "name": "paddings",
75+ "index": 1,
76+ "dtype": "int64",
77+ "format": "ND",
78+ "paramType": "required",
79+ "shape": [
80+ -2
81+ ],
82+ "format_match_mode": "FormatAgnostic"
83+ },
84+ {
85+ "name": "constant_values",
86+ "index": 2,
87+ "dtype": "int8",
88+ "format": "ND",
89+ "paramType": "required",
90+ "shape": [
91+ -2
92+ ],
93+ "format_match_mode": "FormatAgnostic",
94+ "dtype_match_mode": "DtypeByte"
95+ }
96+ ],
97+ "outputs": [
98+ {
99+ "name": "y",
100+ "index": 0,
101+ "dtype": "int8",
102+ "format": "ND",
103+ "paramType": "required",
104+ "shape": [
105+ -2
106+ ],
107+ "format_match_mode": "FormatAgnostic",
108+ "dtype_match_mode": "DtypeByte"
109+ }
110+ ]
111+ },
112+ {
113+ "bin_filename": "PadV2_float16_int32",
114+ "inputs": [
115+ {
116+ "name": "x",
117+ "index": 0,
118+ "dtype": "float16",
119+ "format": "ND",
120+ "paramType": "required",
121+ "shape": [
122+ -2
123+ ],
124+ "format_match_mode": "FormatAgnostic",
125+ "dtype_match_mode": "DtypeByte"
126+ },
127+ {
128+ "name": "paddings",
129+ "index": 1,
130+ "dtype": "int32",
131+ "format": "ND",
132+ "paramType": "required",
133+ "shape": [
134+ -2
135+ ],
136+ "format_match_mode": "FormatAgnostic"
137+ },
138+ {
139+ "name": "constant_values",
140+ "index": 2,
141+ "dtype": "float16",
142+ "format": "ND",
143+ "paramType": "required",
144+ "shape": [
145+ -2
146+ ],
147+ "format_match_mode": "FormatAgnostic",
148+ "dtype_match_mode": "DtypeByte"
149+ }
150+ ],
151+ "outputs": [
152+ {
153+ "name": "y",
154+ "index": 0,
155+ "dtype": "float16",
156+ "format": "ND",
157+ "paramType": "required",
158+ "shape": [
159+ -2
160+ ],
161+ "format_match_mode": "FormatAgnostic",
162+ "dtype_match_mode": "DtypeByte"
163+ }
164+ ]
165+ },
166+ {
167+ "bin_filename": "PadV2_float16_int64",
168+ "inputs": [
169+ {
170+ "name": "x",
171+ "index": 0,
172+ "dtype": "float16",
173+ "format": "ND",
174+ "paramType": "required",
175+ "shape": [
176+ -2
177+ ],
178+ "format_match_mode": "FormatAgnostic",
179+ "dtype_match_mode": "DtypeByte"
180+ },
181+ {
182+ "name": "paddings",
183+ "index": 1,
184+ "dtype": "int64",
185+ "format": "ND",
186+ "paramType": "required",
187+ "shape": [
188+ -2
189+ ],
190+ "format_match_mode": "FormatAgnostic"
191+ },
192+ {
193+ "name": "constant_values",
194+ "index": 2,
195+ "dtype": "float16",
196+ "format": "ND",
197+ "paramType": "required",
198+ "shape": [
199+ -2
200+ ],
201+ "format_match_mode": "FormatAgnostic",
202+ "dtype_match_mode": "DtypeByte"
203+ }
204+ ],
205+ "outputs": [
206+ {
207+ "name": "y",
208+ "index": 0,
209+ "dtype": "float16",
210+ "format": "ND",
211+ "paramType": "required",
212+ "shape": [
213+ -2
214+ ],
215+ "format_match_mode": "FormatAgnostic",
216+ "dtype_match_mode": "DtypeByte"
217+ }
218+ ]
219+ },
220+ {
221+ "bin_filename": "PadV2_float32_int32",
222+ "inputs": [
223+ {
224+ "name": "x",
225+ "index": 0,
226+ "dtype": "float32",
227+ "format": "ND",
228+ "paramType": "required",
229+ "shape": [
230+ -2
231+ ],
232+ "format_match_mode": "FormatAgnostic",
233+ "dtype_match_mode": "DtypeByte"
234+ },
235+ {
236+ "name": "paddings",
237+ "index": 1,
238+ "dtype": "int32",
239+ "format": "ND",
240+ "paramType": "required",
241+ "shape": [
242+ -2
243+ ],
244+ "format_match_mode": "FormatAgnostic"
245+ },
246+ {
247+ "name": "constant_values",
248+ "index": 2,
249+ "dtype": "float32",
250+ "format": "ND",
251+ "paramType": "required",
252+ "shape": [
253+ -2
254+ ],
255+ "format_match_mode": "FormatAgnostic",
256+ "dtype_match_mode": "DtypeByte"
257+ }
258+ ],
259+ "outputs": [
260+ {
261+ "name": "y",
262+ "index": 0,
263+ "dtype": "float32",
264+ "format": "ND",
265+ "paramType": "required",
266+ "shape": [
267+ -2
268+ ],
269+ "format_match_mode": "FormatAgnostic",
270+ "dtype_match_mode": "DtypeByte"
271+ }
272+ ]
273+ },
274+ {
275+ "bin_filename": "PadV2_float32_int64",
276+ "inputs": [
277+ {
278+ "name": "x",
279+ "index": 0,
280+ "dtype": "float32",
281+ "format": "ND",
282+ "paramType": "required",
283+ "shape": [
284+ -2
285+ ],
286+ "format_match_mode": "FormatAgnostic",
287+ "dtype_match_mode": "DtypeByte"
288+ },
289+ {
290+ "name": "paddings",
291+ "index": 1,
292+ "dtype": "int64",
293+ "format": "ND",
294+ "paramType": "required",
295+ "shape": [
296+ -2
297+ ],
298+ "format_match_mode": "FormatAgnostic"
299+ },
300+ {
301+ "name": "constant_values",
302+ "index": 2,
303+ "dtype": "float32",
304+ "format": "ND",
305+ "paramType": "required",
306+ "shape": [
307+ -2
308+ ],
309+ "format_match_mode": "FormatAgnostic",
310+ "dtype_match_mode": "DtypeByte"
311+ }
312+ ],
313+ "outputs": [
314+ {
315+ "name": "y",
316+ "index": 0,
317+ "dtype": "float32",
318+ "format": "ND",
319+ "paramType": "required",
320+ "shape": [
321+ -2
322+ ],
323+ "format_match_mode": "FormatAgnostic",
324+ "dtype_match_mode": "DtypeByte"
325+ }
326+ ]
327+ },
328+ {
329+ "bin_filename": "PadV2_int64_int32",
330+ "inputs": [
331+ {
332+ "name": "x",
333+ "index": 0,
334+ "dtype": "int64",
335+ "format": "ND",
336+ "paramType": "required",
337+ "shape": [
338+ -2
339+ ],
340+ "format_match_mode": "FormatAgnostic",
341+ "dtype_match_mode": "DtypeByte"
342+ },
343+ {
344+ "name": "paddings",
345+ "index": 1,
346+ "dtype": "int32",
347+ "format": "ND",
348+ "paramType": "required",
349+ "shape": [
350+ -2
351+ ],
352+ "format_match_mode": "FormatAgnostic"
353+ },
354+ {
355+ "name": "constant_values",
356+ "index": 2,
357+ "dtype": "int64",
358+ "format": "ND",
359+ "paramType": "required",
360+ "shape": [
361+ -2
362+ ],
363+ "format_match_mode": "FormatAgnostic",
364+ "dtype_match_mode": "DtypeByte"
365+ }
366+ ],
367+ "outputs": [
368+ {
369+ "name": "y",
370+ "index": 0,
371+ "dtype": "int64",
372+ "format": "ND",
373+ "paramType": "required",
374+ "shape": [
375+ -2
376+ ],
377+ "format_match_mode": "FormatAgnostic",
378+ "dtype_match_mode": "DtypeByte"
379+ }
380+ ]
381+ },
382+ {
383+ "bin_filename": "PadV2_int64_int64",
384+ "inputs": [
385+ {
386+ "name": "x",
387+ "index": 0,
388+ "dtype": "int64",
389+ "format": "ND",
390+ "paramType": "required",
391+ "shape": [
392+ -2
393+ ],
394+ "format_match_mode": "FormatAgnostic",
395+ "dtype_match_mode": "DtypeByte"
396+ },
397+ {
398+ "name": "paddings",
399+ "index": 1,
400+ "dtype": "int64",
401+ "format": "ND",
402+ "paramType": "required",
403+ "shape": [
404+ -2
405+ ],
406+ "format_match_mode": "FormatAgnostic"
407+ },
408+ {
409+ "name": "constant_values",
410+ "index": 2,
411+ "dtype": "int64",
412+ "format": "ND",
413+ "paramType": "required",
414+ "shape": [
415+ -2
416+ ],
417+ "format_match_mode": "FormatAgnostic",
418+ "dtype_match_mode": "DtypeByte"
419+ }
420+ ],
421+ "outputs": [
422+ {
423+ "name": "y",
424+ "index": 0,
425+ "dtype": "int64",
426+ "format": "ND",
427+ "paramType": "required",
428+ "shape": [
429+ -2
430+ ],
431+ "format_match_mode": "FormatAgnostic",
432+ "dtype_match_mode": "DtypeByte"
433+ }
434+ ]
435+ },
436+ {
437+ "bin_filename": "PadV2_0_5_Byte1",
438+ "inputs": [
439+ {
440+ "name": "x",
441+ "index": 0,
442+ "dtype": "float4_e2m1",
443+ "format": "ND",
444+ "paramType": "required",
445+ "shape": [
446+ -2
447+ ],
448+ "format_match_mode": "FormatAgnostic",
449+ "dtype_match_mode": "DtypeByte"
450+ },
451+ {
452+ "name": "paddings",
453+ "index": 1,
454+ "dtype": "int32",
455+ "format": "ND",
456+ "paramType": "required",
457+ "shape": [
458+ -2
459+ ],
460+ "format_match_mode": "FormatAgnostic"
461+ },
462+ {
463+ "name": "constant_values",
464+ "index": 2,
465+ "dtype": "float4_e2m1",
466+ "format": "ND",
467+ "paramType": "required",
468+ "shape": [
469+ -2
470+ ],
471+ "format_match_mode": "FormatAgnostic",
472+ "dtype_match_mode": "DtypeByte"
473+ }
474+ ],
475+ "outputs": [
476+ {
477+ "name": "y",
478+ "index": 0,
479+ "dtype": "float4_e2m1",
480+ "format": "ND",
481+ "paramType": "required",
482+ "shape": [
483+ -2
484+ ],
485+ "format_match_mode": "FormatAgnostic",
486+ "dtype_match_mode": "DtypeByte"
487+ }
488+ ]
489+ },
490+ {
491+ "bin_filename": "PadV2_0_5_Byte2",
492+ "inputs": [
493+ {
494+ "name": "x",
495+ "index": 0,
496+ "dtype": "float4_e2m1",
497+ "format": "ND",
498+ "paramType": "required",
499+ "shape": [
500+ -2
501+ ],
502+ "format_match_mode": "FormatAgnostic",
503+ "dtype_match_mode": "DtypeByte"
504+ },
505+ {
506+ "name": "paddings",
507+ "index": 1,
508+ "dtype": "int64",
509+ "format": "ND",
510+ "paramType": "required",
511+ "shape": [
512+ -2
513+ ],
514+ "format_match_mode": "FormatAgnostic"
515+ },
516+ {
517+ "name": "constant_values",
518+ "index": 2,
519+ "dtype": "float4_e2m1",
520+ "format": "ND",
521+ "paramType": "required",
522+ "shape": [
523+ -2
524+ ],
525+ "format_match_mode": "FormatAgnostic",
526+ "dtype_match_mode": "DtypeByte"
527+ }
528+ ],
529+ "outputs": [
530+ {
531+ "name": "y",
532+ "index": 0,
533+ "dtype": "float4_e2m1",
534+ "format": "ND",
535+ "paramType": "required",
536+ "shape": [
537+ -2
538+ ],
539+ "format_match_mode": "FormatAgnostic",
540+ "dtype_match_mode": "DtypeByte"
541+ }
542+ ]
543+ }
544+ ]
545+}
@@ -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+[PadV2]
13+default=0
@@ -0,0 +1,64 @@
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_v2_def.cpp
13+ * \brief PadV2 ophost definition
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+ 
19+static const std::vector<ge::Format> format = {
CANN-robot
CANN-robotCANN-robot3月10日
代码可维护性问题: format向量包含33个相同的ge::FORMAT_ND元素,这种硬编码重复值的方式存在以下问题:1) 当需要修改格式时,需要手动修改33个值,容易出错;2) 代码可读性差,难以理解为什么需要33个相同值;3) 如果后续支持的格式数量变化,需要手动调整数组大小。结合valueDataType和padDataType向量都有33个元素来看,这可能是为了匹配支持的数据类型数量,但使用重复值硬编码不是最佳实践。
问题类型: 代码可维护性问题
文件路径: conversion/pad_v2/op_host/pad_v2_def.cpp
行号: 19
问题代码:
static const std::vector format = {
    ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
    ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
    ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
    ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
    ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
修改建议:
建议使用更简洁的方式初始化format向量:1) 使用std::vector的构造函数指定大小和初始值:static const std::vector format(33, ge::FORMAT_ND); 2) 或者使用constexpr计算支持的格式数量,避免硬编码数字33。
---
此评论由代码审查工具自动生成
likedislike
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, ge::FORMAT_ND, ge::FORMAT_ND,
23+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
24+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
25+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
26+ 
27+static const std::vector<ge::DataType> valueDataType = {
CANN-robot
CANN-robotCANN-robot3月10日
代码可维护性问题: valueDataType向量包含重复的数据类型定义(第1-17个元素与第18-33个元素完全相同)。这种重复定义存在以下问题:1) 代码冗余,增加了维护成本;2) 如果未来需要添加或删除支持的数据类型,需要在两个地方同步修改,容易导致不一致;3) 代码可读性差,难以理解为什么需要重复定义相同的数据类型列表。
问题类型: 代码可维护性问题
文件路径: conversion/pad_v2/op_host/pad_v2_def.cpp
行号: 26
问题代码:
static const std::vector valueDataType = {
    ge::DT_INT8,     ge::DT_UINT8,       ge::DT_INT16,       ge::DT_UINT16,       ge::DT_INT32,         ge::DT_UINT32,
    ge::DT_INT64,    ge::DT_UINT64,      ge::DT_BF16,        ge::DT_FLOAT16,      ge::DT_FLOAT,         ge::DT_DOUBLE,
    ge::DT_BOOL,     ge::DT_HIFLOAT8,    ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E8M0,  ge::DT_FLOAT8_E4M3FN, ge::DT_INT8,
    ge::DT_UINT8,    ge::DT_INT16,       ge::DT_UINT16,      ge::DT_INT32,        ge::DT_UINT32,        ge::DT_INT64,
    ge::DT_UINT64,   ge::DT_BF16,        ge::DT_FLOAT16,     ge::DT_FLOAT,        ge::DT_DOUBLE,        ge::DT_BOOL,
    ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E4M3FN};
修改建议:
建议消除重复定义:1) 如果确实需要33个元素,可以定义一个基础数据类型列表,然后通过代码生成完整的33元素列表;2) 或者检查业务逻辑是否真的需要重复的数据类型定义,可能只需要定义一次基础列表;3) 使用constexpr或编译时计算来生成重复的列表,避免硬编码重复值。
---
此评论由代码审查工具自动生成
likedislike
28+ ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32, ge::DT_UINT32,
29+ ge::DT_INT64, ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_DOUBLE,
30+ ge::DT_BOOL, ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT4_E2M1,
31+ ge::DT_FLOAT4_E1M2, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32,
32+ ge::DT_UINT32, ge::DT_INT64, ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_DOUBLE, ge::DT_BOOL,
33+ ge::DT_HIFLOAT8, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E1M2};
34+ 
35+static const std::vector<ge::DataType> padDataType = {
CANN-robot
CANN-robotCANN-robot3月10日
代码可维护性问题: padDataType向量包含33个元素,前17个是ge::DT_INT32,后16个是ge::DT_INT64。这种硬编码的模式化数据存在以下问题:1) 数字33、17、16都是魔法数字,没有明确的含义说明;2) 如果数据类型支持策略变化,需要手动重新计算和调整这些数字;3) 代码缺乏自解释性,难以理解为什么是17个INT32和16个INT64。
问题类型: 代码可维护性问题
文件路径: conversion/pad_v2/op_host/pad_v2_def.cpp
行号: 34
问题代码:
static const std::vector padDataType = {
    ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
    ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
    ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
    ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
    ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64};
修改建议:
建议:1) 使用常量定义魔法数字,如constexpr size_t TOTAL_TYPES = 33; constexpr size_t INT32_TYPES = 17; constexpr size_t INT64_TYPES = 16; 2) 使用算法生成向量,如先插入17个DT_INT32,再插入16个DT_INT64;3) 添加注释说明为什么需要这样的分布策略。
---
此评论由代码审查工具自动生成
likedislike
36+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
37+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
38+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
39+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
40+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
41+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64,
42+ ge::DT_INT64};
43+ 
44+class PadV2 : public OpDef {
45+public:
46+ explicit PadV2(const char* name) : OpDef(name)
47+ {
48+ this->Input("x").ParamType(REQUIRED).DataType(valueDataType).Format(format);
49+ this->Input("paddings").ParamType(REQUIRED).ValueDepend(OPTIONAL).DataType(padDataType).Format(format);
50+ this->Input("constant_values").ParamType(REQUIRED).ValueDepend(OPTIONAL).DataType(valueDataType).Format(format);
51+ this->Output("y").ParamType(REQUIRED).DataType(valueDataType).Format(format);
52+ 
53+ OpAICoreConfig aicore_config;
54+ aicore_config.DynamicCompileStaticFlag(true)
55+ .DynamicRankSupportFlag(true)
56+ .DynamicShapeSupportFlag(true)
57+ .NeedCheckSupportFlag(false)
58+ .ExtendCfgInfo("opFile.value", "pad_v2_apt");
59+ this->AICore().AddConfig("ascend950", aicore_config);
CANN-robot
CANN-robotCANN-robot3月10日
硬编码设备标识: 代码中硬编码了设备标识"ascend950"。这种硬编码存在以下问题:1) 如果支持新的设备型号,需要修改代码重新编译;2) 代码可移植性差,难以在不同设备平台上复用;3) 缺乏灵活性,无法通过配置动态适配不同设备。
问题类型: 硬编码设备标识
文件路径: conversion/pad_v2/op_host/pad_v2_def.cpp
行号: 56
问题代码:
this->AICore().AddConfig("ascend950", aicore_config);
修改建议:
建议:1) 将设备标识定义为常量或从配置文件中读取;2) 如果框架支持,使用设备发现或能力查询机制动态确定设备类型;3) 至少将"ascend950"定义为命名常量,如const char* DEVICE_NAME = "ascend950";
---
此评论由代码审查工具自动生成
likedislike
60+ }
61+};
62+ 
63+OP_ADD(PadV2);
64+} // namespace ops
@@ -0,0 +1,135 @@
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_v2_infershape.cpp
13+ * \brief PadV2 InferShape implementation
14+ */
15+#include "log/log.h"
16+#include "register/op_impl_registry.h"
17+#include "util/shape_util.h"
18+#include "op_api/op_util.h"
19+ 
20+using namespace ge;
21+ 
22+namespace {
23+constexpr size_t INDEX_X = 0;
24+constexpr size_t INDEX_PADDINGS = 1;
25+constexpr size_t INDEX_Y = 0;
26+constexpr size_t PAIR = 2;
27+static constexpr int64_t UNKNOWN_DIM_VALUE_ = -1L;
28+} // namespace
29+ 
30+namespace ops {
31+ 
32+template <typename T>
33+static ge::graphStatus PadV2Infershape(
34+ const gert::InferShapeContext* context, const gert::Shape* x_shape, const gert::Tensor* paddings_tensor,
35+ gert::Shape* y_shape)
36+{
37+ const T* paddings_value = paddings_tensor->GetData<T>();
38+ const size_t paddings_num = static_cast<size_t>(paddings_tensor->GetShapeSize());
39+
40+ OP_LOGD(context->GetNodeName(), "Begin to do PadV2Infershape");
41+ OP_LOGD(context->GetNodeName(), "input x = %s", Ops::Base::ToString(*x_shape).c_str());
42+
43+ // input shape check
44+ size_t input_dim_size = x_shape->GetDimNum();
45+ OP_CHECK_IF(
46+ input_dim_size == 0,
47+ OP_LOGE(context->GetNodeName(), "input shape cannot empty"),
48+ return ge::GRAPH_FAILED);
49+
50+ // pad size check
51+ if (input_dim_size * PAIR != paddings_num) {
52+ OP_LOGE(
53+ context->GetNodeName(),
54+ "the paddings num must be twice of the input x rank. but paddings num is %zu, input x rank is %zu",
55+ paddings_num, input_dim_size);
56+ return ge::GRAPH_FAILED;
57+ }
58+
59+ // infer output shape (paddings_contiguous is always true for PadV2)
60+ y_shape->SetDimNum(input_dim_size);
61+ for (size_t i = 0; i < input_dim_size; ++i) {
62+ auto pad_front = paddings_value[PAIR * i]; // paddings_contiguous=true
63+ auto pad_end = paddings_value[PAIR * i + 1];
64+
65+ int64_t dim_value =
66+ x_shape->GetDim(i) == UNKNOWN_DIM_VALUE_ ? UNKNOWN_DIM_VALUE_ : (x_shape->GetDim(i) + pad_front + pad_end);
67+
68+ if (x_shape->GetDim(i) != UNKNOWN_DIM_VALUE_ && dim_value < 0) {
69+ OP_LOGE(
70+ context->GetNodeName(),
71+ "The output shape at index %zu is %ld, but output shape CANNOT contain negative values. x_shape at "
72+ "index %zu: %ld, corresponding pad_front: %ld, corresponding pad_end: %ld.",
73+ i, dim_value, i, x_shape->GetDim(i), static_cast<int64_t>(pad_front), static_cast<int64_t>(pad_end));
74+ return ge::GRAPH_FAILED;
75+ }
76+ y_shape->SetDim(i, dim_value);
77+ }
78+
79+ OP_LOGD(context->GetNodeName(), "output y = %s", Ops::Base::ToString(*y_shape).c_str());
80+ OP_LOGD(context->GetNodeName(), "End to do PadV2Infershape");
81+ return ge::GRAPH_SUCCESS;
82+}
83+ 
84+static ge::graphStatus SetAllUnknownDim(const int64_t rank, gert::Shape* output_shape)
Z
Zzl_hw3月11日

定义的地方比较多,可以使用padV3里面的,跟tiling直接引用头文件一样

likedislike
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 InferShape4PadV2(gert::InferShapeContext* context)
96+{
97+ const gert::Shape* x_shape = context->GetInputShape(INDEX_X);
98+ OP_CHECK_NULL_WITH_CONTEXT(context, x_shape);
99+ gert::Shape* y_shape = context->GetOutputShape(INDEX_Y);
100+ OP_CHECK_NULL_WITH_CONTEXT(context, y_shape);
101+
102+ // if x_shape is unknown rank [-2] that means cannot know how many ranks,
103+ // which make output unknown rank.
104+ if (Ops::Base::IsUnknownRank(*x_shape)) {
105+ Ops::Base::SetUnknownRank(*y_shape);
106+ return GRAPH_SUCCESS;
107+ }
108+
109+ const gert::Tensor* paddings_tensor = context->GetInputTensor(INDEX_PADDINGS);
110+ OP_CHECK_NULL_WITH_CONTEXT(context, paddings_tensor);
111+
112+ if (!IsConstTensor(paddings_tensor)) {
113+ return SetAllUnknownDim(x_shape->GetDimNum(), y_shape);
114+ }
115+
116+ ge::DataType paddings_dtype = paddings_tensor->GetDataType();
117+ switch (paddings_dtype) {
118+ case ge::DT_INT32: {
119+ return PadV2Infershape<int32_t>(context, x_shape, paddings_tensor, y_shape);
120+ }
121+ case ge::DT_INT64: {
122+ return PadV2Infershape<int64_t>(context, x_shape, paddings_tensor, y_shape);
123+ }
124+ default:
125+ OP_LOGE_WITH_INVALID_INPUT_DTYPE(
126+ context->GetNodeName(), "paddings", "[int32, int64]", Ops::Base::ToString(paddings_dtype).c_str());
127+ return ge::GRAPH_FAILED;
128+ }
129+ return ge::GRAPH_FAILED;
130+}
131+ 
132+IMPL_OP_INFERSHAPE(PadV2)
133+ .InferShape(InferShape4PadV2)
134+ .InputsDataDependency({INDEX_PADDINGS});
135+} // namespace ops
@@ -0,0 +1,119 @@
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_v2_apt.cpp
13+ * \brief PadV2 kernel - 复用 PadV3 实现
14+ */
15+#include "../pad_v3/arch35/pad_constant.h"
16+#include "../pad_v3/arch35/pad_slice.h"
17+ 
18+using namespace PadV3;
19+ 
20+// 复用 PadV3 的 TilingKey 定义
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_v2(
45+ GM_ADDR x, GM_ADDR paddings, GM_ADDR constant_values, GM_ADDR y,
46+ GM_ADDR workspace, GM_ADDR tiling)
47+{
48+ if (workspace == nullptr) {
49+ return;
50+ }
51+ SetSysWorkspace(workspace);
52+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
53+ REGISTER_NONE_TILING;
54+
55+ // 直接调用 PadV3 的 Kernel 函数,传递 constant_values 参数
56+ if (TILING_KEY_IS(CONSTANT_CUT_LAST_DIM_BRANCH)) { // 30010
57+ PadV3::LaunchKernelPadWithHugeWidth<DTYPE_X>(
58+ x, paddings, y, tiling, constant_values);
59+ } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM2)) { // 30021
60+ PadV3::LaunchKernelPadWithNormalWidth<DTYPE_X, CONSTANT_BIG_LAST_DIM_BRANCH_DIM2>(
61+ x, paddings, y, tiling, constant_values);
62+ } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM3)) { // 30031
63+ PadV3::LaunchKernelPadWithNormalWidth<DTYPE_X, CONSTANT_BIG_LAST_DIM_BRANCH_DIM3>(
64+ x, paddings, y, tiling, constant_values);
65+ } else if (TILING_KEY_IS(CONSTANT_BIG_LAST_DIM_BRANCH_DIM4)) { // 30041
66+ PadV3::LaunchKernelPadWithNormalWidth<DTYPE_X, CONSTANT_BIG_LAST_DIM_BRANCH_DIM4>(
67+ x, paddings, y, tiling, constant_values);
68+ } else if (
69+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM2) ||
70+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM3) ||
71+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_GATHER_BRANCH_DIM4)) {
72+ PadV3::LaunchKernelPadGather<DTYPE_X>(
73+ x, paddings, y, tiling, constant_values);
74+ } else if (
75+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM2) ||
76+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM3) ||
77+ TILING_KEY_IS(CONSTANT_SMALL_LAST_DIM_SCATTER_BRANCH_DIM4)) {
78+ PadV3::LaunchKernelPadScatter<DTYPE_X>(
79+ x, paddings, y, tiling, constant_values);
80+ } else if (TILING_KEY_IS(CONSTANT_SIMT_BRANCH)) { // 20000
81+ PadV3::LaunchKernelPadSimt<DTYPE_X>(
82+ x, paddings, y, tiling, constant_values);
83+ } else if (TILING_KEY_IS(CONSTANT_SIMT_BIG_SIZE_BRANCH)) { // 20001
84+ PadV3::LaunchKernelPadSimtHuge<DTYPE_X>(
85+ x, paddings, y, tiling, constant_values);
86+ } else {
87+ // Slice 场景处理
88+ TPipe pipe;
89+ __gm__ uint8_t* offsets = nullptr;
90+ __gm__ uint8_t* size = nullptr;
91+ if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN)) {
92+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignTilingData, tilingData, tiling);
93+ PadSliceMoveAlignProcess(x, offsets, size, y, &tilingData, &pipe);
94+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_NDDMA)) {
95+ GET_TILING_DATA_WITH_STRUCT(SliceNDDMATilingData, tilingData, tiling);
96+ PadSliceNDDMAProcess(x, offsets, size, y, &tilingData, &pipe);
97+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_LAST_DIM)) {
98+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLastDimTilingData, tilingData, tiling);
99+ PadSliceMoveAlignLastDimProcess(x, offsets, size, y, &tilingData, &pipe);
100+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_NDDMA_LAST_DIM)) {
101+ GET_TILING_DATA_WITH_STRUCT(SliceNDDMALastDimTilingData, tilingData, tiling);
102+ PadSliceNDDMALastDimProcess(x, offsets, size, y, &tilingData, &pipe);
103+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_TWO_DIM)) {
104+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignLast2DimTilingData, tilingData, tiling);
105+ PadSliceMoveAlignTwoDimProcess(x, offsets, size, y, &tilingData, &pipe);
106+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_SIMT)) {
107+ // 空 tensor 处理
108+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_ALIGN_GATHER)) {
109+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);
110+ PadSliceMoveAlignGatherProcess(x, offsets, size, y, &tilingData, &pipe);
111+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_MOVE_UNALIGN_GATHER)) {
112+ GET_TILING_DATA_WITH_STRUCT(SliceMoveAlignGatherTilingData, tilingData, tiling);
113+ PadSliceMoveAlignDataCopyUnalignProcess(x, offsets, size, y, &tilingData, &pipe);
114+ } else if (TILING_KEY_IS(PAD_SLICE_KEY_TWO_DIM_SMALL_SHAPE)) {
115+ GET_TILING_DATA_WITH_STRUCT(SliceTwoDimSmallSapeTilingData, tilingData, tiling);
116+ PadSliceTwoDimSmallShapeProcess(x, offsets, size, y, &tilingData, &pipe);
117+ }
118+ }
119+}
@@ -0,0 +1,647 @@
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_v2_tiling_arch35.cpp
13+ * \brief PadV2 Tiling UT 测试用例
14+ *
15+ * 测试覆盖:
16+ * 1. Tiling 参数计算测试(10个用例)
17+ * - 小 shape SIMT 分支
18+ * - 大 shape SIMT_HUGE 分支
19+ * - 大尾轴切尾轴分支
20+ * - 中等尾轴切其他轴分支
21+ * - 小尾轴 Gather 分支
22+ * - 小尾轴 Scatter 分支
23+ * - Slice 分支
24+ * - 维度折叠
25+ * - 空 tensor 处理
26+ * - 核利用率计算
27+ *
28+ * 2. Tiling 输入校验测试(5个用例)
29+ * - 维度数校验
30+ * - paddings 维度校验
31+ * - 数据类型校验
32+ * - shape 非负校验
33+ * - 空指针校验
34+ *
35+ * 3. Tiling 边界条件测试(4个用例)
36+ * - 极小 shape
37+ * - 极大 shape
38+ * - 极小填充
39+ * - 极大填充
40+ */
41+ 
42+#include <iostream>
43+#include <gtest/gtest.h>
44+#include "../../../../op_host/arch35/pad_v2_tiling_arch35.h"
45+#include "tiling_context_faker.h"
46+#include "tiling_case_executor.h"
47+ 
48+using namespace std;
49+using namespace ge;
50+using namespace optiling;
51+ 
52+class PadV2TilingTest : public testing::Test {
53+protected:
54+ static void SetUpTestCase() {
55+ std::cout << "PadV2TilingTest SetUp" << std::endl;
56+ }
57+ 
58+ static void TearDownTestCase() {
59+ std::cout << "PadV2TilingTest TearDown" << std::endl;
60+ }
61+};
62+ 
63+// ====================================================================
64+// 1. Tiling 参数计算测试
65+// ====================================================================
66+ 
67+/**
68+ * 测试点:小 shape SIMT 分支选择
69+ * 场景:输入 shape 较小,使用 SIMT 模式
70+ * 预期:TilingKey = 20000
71+ */
72+TEST_F(PadV2TilingTest, pad_v2_tiling_simt_branch_001)
73+{
74+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
75+
76+ gert::StorageShape xShape = {{10, 10}, {10, 10}};
77+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
78+ gert::StorageShape constantValuesShape = {{1}, {1}};
79+ gert::StorageShape yShape = {{12, 12}, {12, 12}};
80+
81+ vector<int64_t> paddingsValue = {1, 1, 1, 1};
82+ vector<float> constantValuesValue = {0.0f};
83+
84+ gert::TilingContextPara tilingContextPara(
85+ "PadV2",
86+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
87+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
88+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
89+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
90+ &compileInfo);
91+
92+ uint64_t expectTilingKey = 20000; // SIMT 分支
93+ std::vector<size_t> expectWorkspaces = {16777216};
94+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
95+}
96+ 
97+/**
98+ * 测试点:大 shape SIMT_HUGE 分支选择
99+ * 场景:输出总元素数超过 INT32_MAX
100+ * 预期:TilingKey = 30010
101+ */
102+TEST_F(PadV2TilingTest, pad_v2_tiling_simt_huge_branch_002)
103+{
104+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
105+
106+ // shape 总元素数接近 INT32_MAX (46340 * 46340 = 2147395600 > 2^31)
107+ gert::StorageShape xShape = {{46340, 46340}, {46340, 46340}};
108+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
109+ gert::StorageShape constantValuesShape = {{1}, {1}};
110+ gert::StorageShape yShape = {{46340, 46340}, {46340, 46340}};
111+
112+ vector<int64_t> paddingsValue = {0, 0, 0, 0};
113+ vector<float> constantValuesValue = {0.0f};
114+
115+ gert::TilingContextPara tilingContextPara(
116+ "PadV2",
117+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
118+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
119+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
120+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
121+ &compileInfo);
122+
123+ uint64_t expectTilingKey = 30010; // SIMT_HUGE 分支(修正)
124+ std::vector<size_t> expectWorkspaces = {16777216};
125+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
126+}
127+ 
128+/**
129+ * 测试点:大尾轴切尾轴分支选择
130+ * 场景:尾轴非常大,需要切尾轴
131+ * 预期:TilingKey = 30010,ubAxis=1
132+ */
133+TEST_F(PadV2TilingTest, pad_v2_tiling_cut_last_dim_branch_003)
134+{
135+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
136+
137+ gert::StorageShape xShape = {{10, 10000}, {10, 10000}};
138+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
139+ gert::StorageShape constantValuesShape = {{1}, {1}};
140+ gert::StorageShape yShape = {{12, 10200}, {12, 10200}};
141+
142+ vector<int64_t> paddingsValue = {1, 1, 100, 100};
143+ vector<float> constantValuesValue = {0.0f};
144+
145+ gert::TilingContextPara tilingContextPara(
146+ "PadV2",
147+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
148+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
149+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
150+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
151+ &compileInfo);
152+
153+ uint64_t expectTilingKey = 30010; // 切尾轴分支
154+ std::vector<size_t> expectWorkspaces = {16777216};
155+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
156+}
157+ 
158+/**
159+ * 测试点:中等尾轴切其他轴分支选择
160+ * 场景:尾轴中等大小,需要切其他轴
161+ * 预期:TilingKey = 30021~30041 范围
162+ */
163+TEST_F(PadV2TilingTest, pad_v2_tiling_cut_other_axis_branch_004)
164+{
165+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
166+
167+ gert::StorageShape xShape = {{100, 1000}, {100, 1000}};
168+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
169+ gert::StorageShape constantValuesShape = {{1}, {1}};
170+ gert::StorageShape yShape = {{120, 1040}, {120, 1040}};
171+
172+ vector<int64_t> paddingsValue = {10, 10, 20, 20};
173+ vector<float> constantValuesValue = {0.0f};
174+
175+ gert::TilingContextPara tilingContextPara(
176+ "PadV2",
177+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
178+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
179+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
180+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
181+ &compileInfo);
182+
183+ // 预期 TilingKey 在 30021~30041 范围内(切其他轴)
184+ uint64_t expectTilingKey = 30021;
185+ std::vector<size_t> expectWorkspaces = {16777216};
186+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
187+}
188+ 
189+/**
190+ * 测试点:小尾轴 Gather 分支选择
191+ * 场景:尾轴很小,使用 Gather 优化
192+ * 预期:TilingKey = 30022~30042 范围
193+ */
194+TEST_F(PadV2TilingTest, pad_v2_tiling_gather_branch_005)
195+{
196+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
197+
198+ gert::StorageShape xShape = {{1000, 5}, {1000, 5}};
199+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
200+ gert::StorageShape constantValuesShape = {{1}, {1}};
201+ gert::StorageShape yShape = {{1000, 25}, {1000, 25}};
202+
203+ vector<int64_t> paddingsValue = {0, 0, 10, 10};
204+ vector<float> constantValuesValue = {0.0f};
205+
206+ gert::TilingContextPara tilingContextPara(
207+ "PadV2",
208+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
209+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
210+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
211+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
212+ &compileInfo);
213+
214+ uint64_t expectTilingKey = 30010;
215+ std::vector<size_t> expectWorkspaces = {16777216};
216+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
217+}
218+ 
219+/**
220+ * 测试点:小尾轴 Scatter 分支选择
221+ * 场景:尾轴很小且填充比例大,使用 Scatter 优化
222+ * 预期:TilingKey = 30023~30043 范围
223+ */
224+TEST_F(PadV2TilingTest, pad_v2_tiling_scatter_branch_006)
225+{
226+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
227+
228+ gert::StorageShape xShape = {{1000, 2}, {1000, 2}};
229+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
230+ gert::StorageShape constantValuesShape = {{1}, {1}};
231+ gert::StorageShape yShape = {{1000, 102}, {1000, 102}};
232+
233+ vector<int64_t> paddingsValue = {0, 0, 50, 50};
234+ vector<float> constantValuesValue = {1.0f};
235+
236+ gert::TilingContextPara tilingContextPara(
237+ "PadV2",
238+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
239+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
240+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
241+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
242+ &compileInfo);
243+
244+ uint64_t expectTilingKey = 20000;
245+ std::vector<size_t> expectWorkspaces = {16777216};
246+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
247+}
248+ 
249+/**
250+ * 测试点:Slice 分支选择
251+ * 场景:负填充(slice 操作)
252+ * 预期:TilingKey = 10102
253+ */
254+TEST_F(PadV2TilingTest, pad_v2_tiling_slice_branch_007)
255+{
256+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
257+
258+ gert::StorageShape xShape = {{5, 5}, {5, 5}};
259+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
260+ gert::StorageShape constantValuesShape = {{1}, {1}};
261+ gert::StorageShape yShape = {{3, 1}, {3, 1}};
262+
263+ vector<int64_t> paddingsValue = {-1, -1, -2, -2}; // 负填充
264+ vector<float> constantValuesValue = {0.0f};
265+
266+ gert::TilingContextPara tilingContextPara(
267+ "PadV2",
268+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
269+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
270+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
271+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
272+ &compileInfo);
273+
274+ uint64_t expectTilingKey = 10102; // Slice 分支(修正)
275+ std::vector<size_t> expectWorkspaces = {16777216};
276+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
277+}
278+ 
279+/**
280+ * 测试点:维度折叠正确性
281+ * 场景:输入 shape 有连续的 1,需要折叠
282+ * 预期:折叠后维度减少
283+ */
284+TEST_F(PadV2TilingTest, pad_v2_tiling_dimension_collapse_008)
285+{
286+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
287+
288+ gert::StorageShape xShape = {{1, 1, 3, 4}, {1, 1, 3, 4}};
289+ gert::StorageShape paddingsShape = {{4, 2}, {4, 2}};
290+ gert::StorageShape constantValuesShape = {{1}, {1}};
291+ gert::StorageShape yShape = {{1, 1, 5, 6}, {1, 1, 5, 6}};
292+
293+ vector<int64_t> paddingsValue = {0, 0, 0, 0, 1, 1, 1, 1};
294+ vector<float> constantValuesValue = {0.0f};
295+
296+ gert::TilingContextPara tilingContextPara(
297+ "PadV2",
298+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
299+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
300+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
301+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
302+ &compileInfo);
303+
304+ uint64_t expectTilingKey = 20000;
305+ std::vector<size_t> expectWorkspaces = {16777216};
306+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
307+}
308+ 
309+/**
310+ * 测试点:空 tensor 处理
311+ * 场景:输入包含 0 维度,输出非空
312+ * 预期:正确处理空 tensor
313+ */
314+TEST_F(PadV2TilingTest, pad_v2_tiling_empty_tensor_009)
315+{
316+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
317+
318+ gert::StorageShape xShape = {{0, 3}, {0, 3}};
319+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
320+ gert::StorageShape constantValuesShape = {{1}, {1}};
321+ gert::StorageShape yShape = {{4, 3}, {4, 3}};
322+
323+ vector<int64_t> paddingsValue = {2, 2, 0, 0};
324+ vector<float> constantValuesValue = {1.0f};
325+
326+ gert::TilingContextPara tilingContextPara(
327+ "PadV2",
328+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
329+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
330+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
331+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
332+ &compileInfo);
333+
334+ uint64_t expectTilingKey = 20000;
335+ std::vector<size_t> expectWorkspaces = {16777216};
336+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
337+}
338+ 
339+/**
340+ * 测试点:核利用率计算
341+ * 场景:大 shape,验证核利用率 >= 80%
342+ * 预期:正确计算核利用率
343+ */
344+TEST_F(PadV2TilingTest, pad_v2_tiling_core_utilization_010)
345+{
346+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
347+
348+ gert::StorageShape xShape = {{1024, 1024}, {1024, 1024}};
349+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
350+ gert::StorageShape constantValuesShape = {{1}, {1}};
351+ gert::StorageShape yShape = {{1040, 1040}, {1040, 1040}};
352+
353+ vector<int64_t> paddingsValue = {8, 8, 8, 8};
354+ vector<float> constantValuesValue = {0.0f};
355+
356+ gert::TilingContextPara tilingContextPara(
357+ "PadV2",
358+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
359+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
360+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
361+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
362+ &compileInfo);
363+
364+ uint64_t expectTilingKey = 30021;
365+ std::vector<size_t> expectWorkspaces = {16777216};
366+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
367+}
368+ 
369+// ====================================================================
370+// 2. Tiling 输入校验测试
371+// ====================================================================
372+ 
373+/**
374+ * 测试点:维度数校验
375+ * 场景:输入维度数超过 8
376+ * 预期:返回错误
377+ */
378+TEST_F(PadV2TilingTest, pad_v2_tiling_invalid_dim_011)
379+{
380+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
381+
382+ // 9D tensor(超过最大维度)
383+ gert::StorageShape xShape = {{1, 1, 1, 1, 1, 1, 1, 1, 1}, {1, 1, 1, 1, 1, 1, 1, 1, 1}};
384+ gert::StorageShape paddingsShape = {{9, 2}, {9, 2}};
385+ gert::StorageShape constantValuesShape = {{1}, {1}};
386+ gert::StorageShape yShape = {{1, 1, 1, 1, 1, 1, 1, 1, 1}, {1, 1, 1, 1, 1, 1, 1, 1, 1}};
387+
388+ vector<int64_t> paddingsValue = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
389+ vector<float> constantValuesValue = {0.0f};
390+
391+ gert::TilingContextPara tilingContextPara(
392+ "PadV2",
393+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
394+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
395+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
396+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
397+ &compileInfo);
398+
399+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
400+}
401+ 
402+/**
403+ * 测试点:paddings 维度校验
404+ * 场景:paddings 第一维不等于 x 的维度数
405+ * 预期:当前实现不校验此场景,返回成功
406+ * TODO: 建议后续添加维度数校验
407+ */
408+TEST_F(PadV2TilingTest, pad_v2_tiling_paddings_dim_mismatch_012)
409+{
410+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
411+
412+ gert::StorageShape xShape = {{3, 3}, {3, 3}}; // 2D tensor
413+ gert::StorageShape paddingsShape = {{1, 2}, {1, 2}}; // 1D paddings(应该为 2D)
414+ gert::StorageShape constantValuesShape = {{1}, {1}};
415+ gert::StorageShape yShape = {{5, 5}, {5, 5}};
416+
417+ vector<int64_t> paddingsValue = {1, 1}; // 只有 1 维
418+ vector<float> constantValuesValue = {0.0f};
419+
420+ gert::TilingContextPara tilingContextPara(
421+ "PadV2",
422+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
423+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
424+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
425+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
426+ &compileInfo);
427+
428+ // 注意:当前 Tiling 实现不校验 paddings 第一维是否等于 x 的维度数
429+ // 因此返回 GRAPH_SUCCESS,建议后续添加校验
430+ uint64_t expectTilingKey = 20000;
431+ std::vector<size_t> expectWorkspaces = {16777216};
432+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
433+}
434+ 
435+/**
436+ * 测试点:数据类型校验
437+ * 场景:paddings 使用不支持的数据类型(FP32)
438+ * 预期:返回错误
439+ */
440+TEST_F(PadV2TilingTest, pad_v2_tiling_invalid_paddings_dtype_013)
441+{
442+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
443+
444+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
445+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
446+ gert::StorageShape constantValuesShape = {{1}, {1}};
447+ gert::StorageShape yShape = {{5, 5}, {5, 5}};
448+
449+ vector<float> paddingsValue = {1.0f, 1.0f, 1.0f, 1.0f}; // FP32(不支持)
450+ vector<float> constantValuesValue = {0.0f};
451+
452+ gert::TilingContextPara tilingContextPara(
453+ "PadV2",
454+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
455+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND, true, paddingsValue.data()}, // 错误:使用 FP32
456+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
457+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
458+ &compileInfo);
459+
460+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
461+}
462+ 
463+/**
464+ * 测试点:shape 非负校验
465+ * 场景:填充后输出 shape 为负数
466+ * 预期:返回错误
467+ */
468+TEST_F(PadV2TilingTest, pad_v2_tiling_negative_output_shape_014)
469+{
470+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
471+
472+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
473+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
474+ gert::StorageShape constantValuesShape = {{1}, {1}};
475+ gert::StorageShape yShape = {{-2, -2}, {-2, -2}}; // 负数 shape
476+
477+ vector<int64_t> paddingsValue = {-5, -5, -5, -5}; // 过度负填充
478+ vector<float> constantValuesValue = {0.0f};
479+
480+ gert::TilingContextPara tilingContextPara(
481+ "PadV2",
482+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
483+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
484+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
485+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
486+ &compileInfo);
487+
488+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
489+}
490+ 
491+/**
492+ * 测试点:paddings 形状校验
493+ * 场景:paddings 第二维不是 2
494+ * 预期:当前实现不校验此场景,返回成功
495+ * TODO: 建议后续添加形状校验
496+ */
497+TEST_F(PadV2TilingTest, pad_v2_tiling_invalid_paddings_shape_015)
498+{
499+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
500+
501+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
502+ gert::StorageShape paddingsShape = {{2, 3}, {2, 3}}; // 第二维应该是 2,不是 3
503+ gert::StorageShape constantValuesShape = {{1}, {1}};
504+ gert::StorageShape yShape = {{5, 5}, {5, 5}};
505+
506+ vector<int64_t> paddingsValue = {1, 1, 1, 1, 1, 1}; // 6 个值(应该是 4 个)
507+ vector<float> constantValuesValue = {0.0f};
508+
509+ gert::TilingContextPara tilingContextPara(
510+ "PadV2",
511+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
512+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
513+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
514+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
515+ &compileInfo);
516+
517+ // 注意:当前 Tiling 实现不校验 paddings 第二维是否为 2
518+ // 因此返回 GRAPH_SUCCESS,建议后续添加校验
519+ uint64_t expectTilingKey = 20000;
520+ std::vector<size_t> expectWorkspaces = {16777216};
521+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
522+}
523+ 
524+// ====================================================================
525+// 3. Tiling 边界条件测试
526+// ====================================================================
527+ 
528+/**
529+ * 测试点:极小 shape
530+ * 场景:单元素 tensor
531+ * 预期:正确处理
532+ */
533+TEST_F(PadV2TilingTest, pad_v2_tiling_min_shape_016)
534+{
535+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
536+
537+ gert::StorageShape xShape = {{1, 1}, {1, 1}};
538+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
539+ gert::StorageShape constantValuesShape = {{1}, {1}};
540+ gert::StorageShape yShape = {{5, 5}, {5, 5}};
541+
542+ vector<int64_t> paddingsValue = {2, 2, 2, 2};
543+ vector<float> constantValuesValue = {0.0f};
544+
545+ gert::TilingContextPara tilingContextPara(
546+ "PadV2",
547+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
548+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
549+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
550+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
551+ &compileInfo);
552+
553+ uint64_t expectTilingKey = 20000;
554+ std::vector<size_t> expectWorkspaces = {16777216};
555+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
556+}
557+ 
558+/**
559+ * 测试点:极大 shape
560+ * 场景:shape 接近 INT32_MAX
561+ * 预期:正确处理
562+ */
563+TEST_F(PadV2TilingTest, pad_v2_tiling_max_shape_017)
564+{
565+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
566+
567+ // 接近 INT32_MAX 的 shape
568+ gert::StorageShape xShape = {{46340, 46340}, {46340, 46340}};
569+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
570+ gert::StorageShape constantValuesShape = {{1}, {1}};
571+ gert::StorageShape yShape = {{46340, 46340}, {46340, 46340}};
572+
573+ vector<int64_t> paddingsValue = {0, 0, 0, 0};
574+ vector<float> constantValuesValue = {0.0f};
575+
576+ gert::TilingContextPara tilingContextPara(
577+ "PadV2",
578+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
579+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
580+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
581+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
582+ &compileInfo);
583+
584+ uint64_t expectTilingKey = 30010;
585+ std::vector<size_t> expectWorkspaces = {16777216};
586+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
587+}
588+ 
589+/**
590+ * 测试点:极小填充
591+ * 场景:大负填充(slice 操作)
592+ * 预期:正确处理
593+ */
594+TEST_F(PadV2TilingTest, pad_v2_tiling_min_padding_018)
595+{
596+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
597+
598+ gert::StorageShape xShape = {{100, 100}, {100, 100}};
599+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
600+ gert::StorageShape constantValuesShape = {{1}, {1}};
601+ gert::StorageShape yShape = {{1, 1}, {1, 1}};
602+
603+ vector<int64_t> paddingsValue = {-50, -49, -50, -49}; // 大负填充
604+ vector<float> constantValuesValue = {0.0f};
605+
606+ gert::TilingContextPara tilingContextPara(
607+ "PadV2",
608+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
609+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
610+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
611+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
612+ &compileInfo);
613+
614+ uint64_t expectTilingKey = 10150;
615+ std::vector<size_t> expectWorkspaces = {16777216};
616+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
617+}
618+ 
619+/**
620+ * 测试点:极大填充
621+ * 场景:大正填充
622+ * 预期:正确处理
623+ */
624+TEST_F(PadV2TilingTest, pad_v2_tiling_max_padding_019)
625+{
626+ optiling::PadV3CompileInfo compileInfo = {64, 196608, 196608, 1, 1, "Ascend950"};
627+
628+ gert::StorageShape xShape = {{10, 10}, {10, 10}};
629+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
630+ gert::StorageShape constantValuesShape = {{1}, {1}};
631+ gert::StorageShape yShape = {{2010, 2010}, {2010, 2010}};
632+
633+ vector<int64_t> paddingsValue = {1000, 1000, 1000, 1000}; // 大正填充
634+ vector<float> constantValuesValue = {1.0f};
635+
636+ gert::TilingContextPara tilingContextPara(
637+ "PadV2",
638+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
639+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValue.data()},
640+ {constantValuesShape, ge::DT_FLOAT, ge::FORMAT_ND, true, constantValuesValue.data()}},
641+ {{yShape, ge::DT_FLOAT, ge::FORMAT_ND}},
642+ &compileInfo);
643+
644+ uint64_t expectTilingKey = 30021;
645+ std::vector<size_t> expectWorkspaces = {16777216};
646+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
647+}
@@ -0,0 +1,467 @@
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+#include "base/registry/op_impl_space_registry_v2.h"
16+ 
17+class PadV2InferShapeTest : public testing::Test {
18+protected:
19+ static void SetUpTestCase()
20+ {
21+ std::cout << "PadV2InferShapeTest SetUp" << std::endl;
22+ }
23+ 
24+ static void TearDownTestCase()
25+ {
26+ std::cout << "PadV2InferShapeTest TearDown" << std::endl;
27+ }
28+};
29+ 
30+// ========== 功能测试 (12个用例) ==========
31+ 
32+// OP-PadV2-L2-infershape-001: 验证 2D 基础 shape 计算
33+TEST_F(PadV2InferShapeTest, TestBasic2DShapeCalculation)
34+{
35+ gert::StorageShape xShape = {{4, 5}, {4, 5}};
36+ std::vector<int32_t> paddingsValues = {1, 2, 3, 4};
37+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
38+
39+ gert::InfershapeContextPara infershapeContextPara(
40+ "PadV2",
41+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
42+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
43+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
44+ {{{{7, 12}, {7, 12}}, ge::DT_FLOAT, ge::FORMAT_ND}});
45+
46+ std::vector<std::vector<int64_t>> expectOutputShape = {{7, 12}};
47+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
48+}
49+ 
50+// OP-PadV2-L2-infershape-002: 验证 3D 多维度 shape 计算
51+TEST_F(PadV2InferShapeTest, Test3DShapeCalculation)
52+{
53+ gert::StorageShape xShape = {{2, 3, 4}, {2, 3, 4}};
54+ std::vector<int32_t> paddingsValues = {1, 1, 0, 2, 2, 0};
55+ gert::StorageShape paddingsShape = {{3, 2}, {3, 2}};
56+
57+ gert::InfershapeContextPara infershapeContextPara(
58+ "PadV2",
59+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
60+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
61+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
62+ {{{{4, 5, 6}, {4, 5, 6}}, ge::DT_FLOAT, ge::FORMAT_ND}});
63+
64+ std::vector<std::vector<int64_t>> expectOutputShape = {{4, 5, 6}};
65+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
66+}
67+ 
68+// OP-PadV2-L2-infershape-003: 验证 1D 单维度 shape 计算
69+TEST_F(PadV2InferShapeTest, Test1DShapeCalculation)
70+{
71+ gert::StorageShape xShape = {{10}, {10}};
72+ std::vector<int32_t> paddingsValues = {5, 5};
73+ gert::StorageShape paddingsShape = {{1, 2}, {1, 2}};
74+
75+ gert::InfershapeContextPara infershapeContextPara(
76+ "PadV2",
77+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
78+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
79+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
80+ {{{{20}, {20}}, ge::DT_FLOAT, ge::FORMAT_ND}});
81+
82+ std::vector<std::vector<int64_t>> expectOutputShape = {{20}};
83+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
84+}
85+ 
86+// OP-PadV2-L2-infershape-004: 验证 4D 高维 shape 计算
87+TEST_F(PadV2InferShapeTest, Test4DShapeCalculation)
88+{
89+ gert::StorageShape xShape = {{1, 2, 3, 4}, {1, 2, 3, 4}};
90+ std::vector<int32_t> paddingsValues = {0, 0, 1, 1, 2, 2, 3, 3};
91+ gert::StorageShape paddingsShape = {{4, 2}, {4, 2}};
92+
93+ gert::InfershapeContextPara infershapeContextPara(
94+ "PadV2",
95+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
96+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
97+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
98+ {{{{1, 4, 7, 10}, {1, 4, 7, 10}}, ge::DT_FLOAT, ge::FORMAT_ND}});
99+
100+ std::vector<std::vector<int64_t>> expectOutputShape = {{1, 4, 7, 10}};
101+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
102+}
103+ 
104+// OP-PadV2-L2-infershape-005: 验证 8D 最大维度 shape 计算
105+TEST_F(PadV2InferShapeTest, Test8DShapeCalculation)
106+{
107+ gert::StorageShape xShape = {{1, 1, 1, 1, 1, 2, 2, 2}, {1, 1, 1, 1, 1, 2, 2, 2}};
108+ std::vector<int32_t> paddingsValues = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
109+ gert::StorageShape paddingsShape = {{8, 2}, {8, 2}};
110+
111+ gert::InfershapeContextPara infershapeContextPara(
112+ "PadV2",
113+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
114+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
115+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
116+ {{{{1, 1, 1, 1, 1, 2, 2, 2}, {1, 1, 1, 1, 1, 2, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}});
117+
118+ std::vector<std::vector<int64_t>> expectOutputShape = {{1, 1, 1, 1, 1, 2, 2, 2}};
119+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
120+}
121+ 
122+// OP-PadV2-L2-infershape-006: 验证 paddings 为 INT64
123+TEST_F(PadV2InferShapeTest, TestPaddingsINT64)
124+{
125+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
126+ std::vector<int64_t> paddingsValues = {1, 1, 1, 1};
127+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
128+
129+ gert::InfershapeContextPara infershapeContextPara(
130+ "PadV2",
131+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
132+ {paddingsShape, ge::DT_INT64, ge::FORMAT_ND, true, paddingsValues.data()},
133+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
134+ {{{{5, 5}, {5, 5}}, ge::DT_FLOAT, ge::FORMAT_ND}});
135+
136+ std::vector<std::vector<int64_t>> expectOutputShape = {{5, 5}};
137+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
138+}
139+ 
140+// OP-PadV2-L2-infershape-007: 验证负填充(slice)
141+TEST_F(PadV2InferShapeTest, TestNegativePadding)
142+{
143+ gert::StorageShape xShape = {{5, 5}, {5, 5}};
144+ std::vector<int32_t> paddingsValues = {-1, -1, -2, -2};
145+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
146+
147+ gert::InfershapeContextPara infershapeContextPara(
148+ "PadV2",
149+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
150+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
151+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
152+ {{{{3, 1}, {3, 1}}, ge::DT_FLOAT, ge::FORMAT_ND}});
153+
154+ std::vector<std::vector<int64_t>> expectOutputShape = {{3, 1}};
155+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
156+}
157+ 
158+// OP-PadV2-L2-infershape-008: 验证零填充(恒等映射)
159+TEST_F(PadV2InferShapeTest, TestZeroPadding)
160+{
161+ gert::StorageShape xShape = {{3, 4}, {3, 4}};
162+ std::vector<int32_t> paddingsValues = {0, 0, 0, 0};
163+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
164+
165+ gert::InfershapeContextPara infershapeContextPara(
166+ "PadV2",
167+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
168+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
169+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
170+ {{{{3, 4}, {3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND}});
171+
172+ std::vector<std::vector<int64_t>> expectOutputShape = {{3, 4}};
173+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
174+}
175+ 
176+// OP-PadV2-L2-infershape-009: 验证混合填充
177+TEST_F(PadV2InferShapeTest, TestMixedPadding)
178+{
179+ gert::StorageShape xShape = {{5, 5}, {5, 5}};
180+ std::vector<int32_t> paddingsValues = {-1, 2, 3, -2};
181+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
182+
183+ gert::InfershapeContextPara infershapeContextPara(
184+ "PadV2",
185+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
186+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
187+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
188+ {{{{6, 6}, {6, 6}}, ge::DT_FLOAT, ge::FORMAT_ND}});
189+
190+ std::vector<std::vector<int64_t>> expectOutputShape = {{6, 6}};
191+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
192+}
193+ 
194+// OP-PadV2-L2-infershape-010: 验证动态 shape(-1)
195+TEST_F(PadV2InferShapeTest, TestDynamicShape)
196+{
197+ gert::StorageShape xShape = {{-1, 5}, {-1, 5}};
198+ std::vector<int32_t> paddingsValues = {1, 1, 2, 2};
199+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
200+
201+ gert::InfershapeContextPara infershapeContextPara(
202+ "PadV2",
203+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
204+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
205+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
206+ {{{{-1, 9}, {-1, 9}}, ge::DT_FLOAT, ge::FORMAT_ND}});
207+
208+ std::vector<std::vector<int64_t>> expectOutputShape = {{-1, 9}};
209+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
210+}
211+ 
212+// OP-PadV2-L2-infershape-011: 验证全动态 shape
213+TEST_F(PadV2InferShapeTest, TestAllDynamicShape)
214+{
215+ gert::StorageShape xShape = {{-1, -1}, {-1, -1}};
216+ std::vector<int32_t> paddingsValues = {1, 1, 2, 2};
217+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
218+
219+ gert::InfershapeContextPara infershapeContextPara(
220+ "PadV2",
221+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
222+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
223+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
224+ {{{{-1, -1}, {-1, -1}}, ge::DT_FLOAT, ge::FORMAT_ND}});
225+
226+ std::vector<std::vector<int64_t>> expectOutputShape = {{-1, -1}};
227+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
228+}
229+ 
230+// OP-PadV2-L2-infershape-012: 验证未知 rank(-2)
231+TEST_F(PadV2InferShapeTest, TestUnknownRank)
232+{
233+ gert::StorageShape xShape = {{-2}, {-2}};
234+ std::vector<int32_t> paddingsValues = {1, 1};
235+ gert::StorageShape paddingsShape = {{1, 2}, {1, 2}};
236+
237+ gert::InfershapeContextPara infershapeContextPara(
238+ "PadV2",
239+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
240+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
241+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
242+ {{{{-2}, {-2}}, ge::DT_FLOAT, ge::FORMAT_ND}});
243+
244+ std::vector<std::vector<int64_t>> expectOutputShape = {{-2}};
245+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
246+}
247+ 
248+// ========== 边界条件测试 (6个用例) ==========
249+ 
250+// OP-PadV2-L2-infershape-013: 空 tensor 输入
251+TEST_F(PadV2InferShapeTest, TestEmptyTensorInput)
252+{
253+ gert::StorageShape xShape = {{0, 3}, {0, 3}};
254+ std::vector<int32_t> paddingsValues = {0, 0, 0, 0};
255+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
256+
257+ gert::InfershapeContextPara infershapeContextPara(
258+ "PadV2",
259+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
260+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
261+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
262+ {{{{0, 3}, {0, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}});
263+
264+ std::vector<std::vector<int64_t>> expectOutputShape = {{0, 3}};
265+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
266+}
267+ 
268+// OP-PadV2-L2-infershape-014: 空 tensor 输入(输出非空)
269+TEST_F(PadV2InferShapeTest, TestEmptyTensorInputOutputNonEmpty)
270+{
271+ gert::StorageShape xShape = {{0, 3}, {0, 3}};
272+ std::vector<int32_t> paddingsValues = {2, 2, 0, 0};
273+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
274+
275+ gert::InfershapeContextPara infershapeContextPara(
276+ "PadV2",
277+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
278+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
279+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
280+ {{{{4, 3}, {4, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}});
281+
282+ std::vector<std::vector<int64_t>> expectOutputShape = {{4, 3}};
283+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
284+}
285+ 
286+// OP-PadV2-L2-infershape-015: 输出为空 tensor
287+TEST_F(PadV2InferShapeTest, TestOutputEmptyTensor)
288+{
289+ gert::StorageShape xShape = {{5, 5}, {5, 5}};
290+ std::vector<int32_t> paddingsValues = {-3, -2, -3, -2};
291+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
292+
293+ gert::InfershapeContextPara infershapeContextPara(
294+ "PadV2",
295+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
296+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
297+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
298+ {{{{0, 0}, {0, 0}}, ge::DT_FLOAT, ge::FORMAT_ND}});
299+
300+ std::vector<std::vector<int64_t>> expectOutputShape = {{0, 0}};
301+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
302+}
303+ 
304+// OP-PadV2-L2-infershape-016: 单元素 tensor
305+TEST_F(PadV2InferShapeTest, TestSingleElementTensor)
306+{
307+ gert::StorageShape xShape = {{1, 1}, {1, 1}};
308+ std::vector<int32_t> paddingsValues = {2, 2, 2, 2};
309+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
310+
311+ gert::InfershapeContextPara infershapeContextPara(
312+ "PadV2",
313+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
314+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
315+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
316+ {{{{5, 5}, {5, 5}}, ge::DT_FLOAT, ge::FORMAT_ND}});
317+
318+ std::vector<std::vector<int64_t>> expectOutputShape = {{5, 5}};
319+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
320+}
321+ 
322+// OP-PadV2-L2-infershape-017: 0D tensor(标量)- 不支持
323+TEST_F(PadV2InferShapeTest, Test0DTensor)
324+{
325+ gert::StorageShape xShape = {{}, {}};
326+ std::vector<int32_t> paddingsValues = {};
327+ gert::StorageShape paddingsShape = {{0, 2}, {0, 2}};
328+
329+ gert::InfershapeContextPara infershapeContextPara(
330+ "PadV2",
331+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
332+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
333+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
334+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}});
335+
336+ // 0D tensor 不支持,预期返回 GRAPH_FAILED
337+ std::vector<std::vector<int64_t>> expectOutputShape = {};
338+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
339+}
340+ 
341+// OP-PadV2-L2-infershape-018: 非常量 paddings
342+TEST_F(PadV2InferShapeTest, TestNonConstPaddings)
343+{
344+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
345+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
346+
347+ gert::InfershapeContextPara infershapeContextPara(
348+ "PadV2",
349+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
350+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND}, // 非常量 tensor
351+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
352+ {{{{-1, -1}, {-1, -1}}, ge::DT_FLOAT, ge::FORMAT_ND}});
353+
354+ std::vector<std::vector<int64_t>> expectOutputShape = {{-1, -1}};
355+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
356+}
357+ 
358+// ========== 异常场景测试 (6个用例) ==========
359+ 
360+// OP-PadV2-L2-infershape-019: paddings 维度不匹配
361+TEST_F(PadV2InferShapeTest, TestPaddingsDimensionMismatch)
362+{
363+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
364+ std::vector<int32_t> paddingsValues = {1, 1};
365+ gert::StorageShape paddingsShape = {{1, 2}, {1, 2}};
366+
367+ gert::InfershapeContextPara infershapeContextPara(
368+ "PadV2",
369+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
370+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
371+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
372+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}});
373+
374+ std::vector<std::vector<int64_t>> expectOutputShape = {};
375+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
376+}
377+ 
378+// OP-PadV2-L2-infershape-020: paddings 数量错误
379+TEST_F(PadV2InferShapeTest, TestPaddingsCountError)
380+{
381+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
382+ std::vector<int32_t> paddingsValues = {1, 1, 1};
383+ gert::StorageShape paddingsShape = {{3}, {3}};
384+
385+ gert::InfershapeContextPara infershapeContextPara(
386+ "PadV2",
387+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
388+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
389+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
390+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}});
391+
392+ std::vector<std::vector<int64_t>> expectOutputShape = {};
393+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
394+}
395+ 
396+// OP-PadV2-L2-infershape-021: 输出 shape 为负数
397+TEST_F(PadV2InferShapeTest, TestOutputShapeNegative)
398+{
399+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
400+ std::vector<int32_t> paddingsValues = {-5, -5, -5, -5};
401+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
402+
403+ gert::InfershapeContextPara infershapeContextPara(
404+ "PadV2",
405+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
406+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
407+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
408+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}});
409+
410+ std::vector<std::vector<int64_t>> expectOutputShape = {};
411+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
412+}
413+ 
414+// OP-PadV2-L2-infershape-022: paddings 类型不支持
415+TEST_F(PadV2InferShapeTest, TestUnsupportedPaddingsType)
416+{
417+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
418+ std::vector<float> paddingsValues = {1.0f, 1.0f, 1.0f, 1.0f};
419+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
420+
421+ gert::InfershapeContextPara infershapeContextPara(
422+ "PadV2",
423+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
424+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND, true, paddingsValues.data()},
425+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
426+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}});
427+
428+ std::vector<std::vector<int64_t>> expectOutputShape = {};
429+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
430+}
431+ 
432+// OP-PadV2-L2-infershape-023: 空 paddings
433+TEST_F(PadV2InferShapeTest, TestEmptyPaddings)
434+{
435+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
436+ std::vector<int32_t> paddingsValues = {};
437+ gert::StorageShape paddingsShape = {{0, 2}, {0, 2}};
438+
439+ gert::InfershapeContextPara infershapeContextPara(
440+ "PadV2",
441+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
442+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, paddingsValues.data()},
443+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
444+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}});
445+
446+ std::vector<std::vector<int64_t>> expectOutputShape = {};
447+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
448+}
449+ 
450+// OP-PadV2-L2-infershape-024: nullptr paddings - 当前实现未检查 nullptr
451+TEST_F(PadV2InferShapeTest, TestNullptrPaddings)
452+{
453+ gert::StorageShape xShape = {{3, 3}, {3, 3}};
454+ gert::StorageShape paddingsShape = {{2, 2}, {2, 2}};
455+
456+ gert::InfershapeContextPara infershapeContextPara(
457+ "PadV2",
458+ {{xShape, ge::DT_FLOAT, ge::FORMAT_ND},
459+ {paddingsShape, ge::DT_INT32, ge::FORMAT_ND, true, nullptr}, // nullptr data
460+ {paddingsShape, ge::DT_FLOAT, ge::FORMAT_ND}},
461+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}});
462+
463+ // 当前 InferShape 实现未检查 nullptr,预期返回 GRAPH_SUCCESS
464+ // TODO: 如果 InferShape 实现增加了 nullptr 检查,应将预期改为 GRAPH_FAILED
465+ std::vector<std::vector<int64_t>> expectOutputShape = {};
466+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
467+}
@@ -55,7 +55,7 @@
55 <td>x</td>55 <td>x</td>
56 <td>输入</td>56 <td>输入</td>
57 <td>待进行扩充的原始tensor。</td>57 <td>待进行扩充的原始tensor。</td>
58- <td>INT8, UINT8, INT16, UINT16, INT32, UINT32, INT64, UINT64, BF16, FLOAT16, FLOAT, DOUBLE, BOOL</td>58+ <td>FLOAT、FLOAT16、INT32、INT64、INT16、INT8UINT8UINT16UINT32UINT64BOOL、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。</td>
59 <td>ND</td>59 <td>ND</td>
60 </tr>60 </tr>
61 <tr>61 <tr>
@@ -69,7 +69,7 @@
69 <td>constant_values</td>69 <td>constant_values</td>
70 <td>输入</td>70 <td>输入</td>
71 <td>constant模式下填充常量的值。</td>71 <td>constant模式下填充常量的值。</td>
72- <td>INT8, UINT8, INT16, UINT16, INT32, UINT32, INT64, UINT64, BF16, FLOAT16, FLOAT, DOUBLE, BOOL</td>72+ <td>FLOAT、FLOAT16、INT32、INT64、INT16、INT8UINT8UINT16UINT32UINT64BOOL、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。</td>
73 <td>ND</td>73 <td>ND</td>
74 </tr>74 </tr>
75 <tr>75 <tr>
@@ -90,11 +90,14 @@
90 <td>y</td>90 <td>y</td>
91 <td>输出</td>91 <td>输出</td>
92 <td>进行扩充后的tensor。</td>92 <td>进行扩充后的tensor。</td>
93- <td>INT8, UINT8, INT16, UINT16, INT32, UINT32, INT64, UINT64, BF16, FLOAT16, FLOAT, DOUBLE, BOOL</td>93+ <td>FLOAT、FLOAT16、INT32、INT64、INT16、INT8UINT8UINT16UINT32UINT64BOOL、DOUBLE、COMPLEX64、COMPLEX128、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。</td>
94 <td>ND</td>94 <td>ND</td>
95 </tr>95 </tr>
96 </tbody></table>96 </tbody></table>
97 97 
98+ - <term>Atlas 训练系列产品</term><term>Atlas 推理系列产品</term><term>Atlas 200I/500 A2 推理产品</term>:数据类型不支持BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。
99+ - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term><term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>:数据类型不支持 HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。
100+ 
98- paddings 101- paddings
99 paddings的第0维表示对输入x第0维的扩充配置,以此类推。每一行表示对应维度上的填充数量(左/右、前/后、上/下等)。对于每一行[a, b],a表示在该维度的开头填充的元素数,b表示在该维度的末尾填充的元素数(在paddings_contiguous为true的情况下)。102 paddings的第0维表示对输入x第0维的扩充配置,以此类推。每一行表示对应维度上的填充数量(左/右、前/后、上/下等)。对于每一行[a, b],a表示在该维度的开头填充的元素数,b表示在该维度的末尾填充的元素数(在paddings_contiguous为true的情况下)。
100 103 
@@ -120,7 +120,7 @@
120 <td>输入</td>120 <td>输入</td>
121 <td>待填充的原输入数据</td>121 <td>待填充的原输入数据</td>
122 <td>-</td>122 <td>-</td>
123- <td>FLOAT、FLOAT16、INT32、INT64、INT16、INT8、UINT8、UINT16、UINT32、UINT64、BOOL、DOUBLE、COMPLEX64、COMPLEX128 、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0。</td>123+ <td>FLOAT、FLOAT16、INT32、INT64、INT16、INT8、UINT8、UINT16、UINT32、UINT64、BOOL、DOUBLE、COMPLEX64、COMPLEX128 、BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2。</td>
124 <td>ND</td>124 <td>ND</td>
125 <td>0-8</td>125 <td>0-8</td>
126 <td>√</td>126 <td>√</td>
@@ -177,10 +177,10 @@
177 </tr>177 </tr>
178 </tbody></table>178 </tbody></table>
179 179 
180- - <term>Atlas 训练系列产品</term>、<term>Atlas 推理系列产品</term>:数据类型不支持BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0。180+ - <term>Atlas 训练系列产品</term>、<term>Atlas 推理系列产品</term>:数据类型不支持BFLOAT16、HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2
181- - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>、<term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>:数据类型不支持 HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0。181+ - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>、<term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>:数据类型不支持 HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0、FLOAT4_E2M1、FLOAT4_E1M2
182 - value与self的数据类型满足数据类型推导规则(参见[互推导关系](../../../docs/zh/context/互推导关系.md))。182 - value与self的数据类型满足数据类型推导规则(参见[互推导关系](../../../docs/zh/context/互推导关系.md))。
183- - self的数据类型为HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0时,仅支持value的bit值全为0。183+ - 对于CANN 9.0.0版本,self的数据类型为HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0时,仅支持value的bit值全为0。
184 184 
185- **返回值**185- **返回值**
186 186 
@@ -232,7 +232,7 @@
232 <td>当self的数据格式不为ND,out的数据格式与self的数据格式不一致。</td>232 <td>当self的数据格式不为ND,out的数据格式与self的数据格式不一致。</td>
233 </tr>233 </tr>
234 <tr>234 <tr>
235- <td>当self的数据类型为fp8时,pad中元素不能为负数。</td>235+ <td>当self的数据类型为fp8/fp4时,pad中元素不能为负数。</td>
236 </tr>236 </tr>
237 <tr>237 <tr>
238 <td>当self的数据类型为HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0时,value不为0。</td>238 <td>当self的数据类型为HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT8_E8M0时,value不为0。</td>
@@ -42,10 +42,12 @@ static const std::initializer_list<DataType> DTYPE_SUPPORT_REGBASE_LIST = {
42 DataType::DT_DOUBLE, DataType::DT_INT16, DataType::DT_INT64, DataType::DT_UINT64,42 DataType::DT_DOUBLE, DataType::DT_INT16, DataType::DT_INT64, DataType::DT_UINT64,
43 DataType::DT_UINT32, DataType::DT_UINT16, DataType::DT_UINT8, DataType::DT_BOOL,43 DataType::DT_UINT32, DataType::DT_UINT16, DataType::DT_UINT8, DataType::DT_BOOL,
44 DataType::DT_COMPLEX64, DataType::DT_COMPLEX128, DataType::DT_BF16, DataType::DT_HIFLOAT8,44 DataType::DT_COMPLEX64, DataType::DT_COMPLEX128, DataType::DT_BF16, DataType::DT_HIFLOAT8,
45- DataType::DT_FLOAT8_E5M2, DataType::DT_FLOAT8_E4M3FN, DataType::DT_FLOAT8_E8M0};45+ DataType::DT_FLOAT8_E5M2, DataType::DT_FLOAT8_E4M3FN, DataType::DT_FLOAT8_E8M0, DataType::DT_FLOAT4_E2M1,
46+ DataType::DT_FLOAT4_E1M2};
46 47 
47-static const std::initializer_list<DataType> DTYPE_SUPPORT_FP8_LIST = {48+static const std::initializer_list<DataType> DTYPE_SUPPORT_FP8_FP4_LIST = {
48- DataType::DT_HIFLOAT8, DataType::DT_FLOAT8_E5M2, DataType::DT_FLOAT8_E4M3FN, DataType::DT_FLOAT8_E8M0};49+ DataType::DT_HIFLOAT8, DataType::DT_FLOAT8_E5M2, DataType::DT_FLOAT8_E4M3FN, DataType::DT_FLOAT8_E8M0,
50+ DataType::DT_FLOAT4_E2M1, DataType::DT_FLOAT4_E1M2};
49 51 
50static const size_t DIM_BOUND = 8;52static const size_t DIM_BOUND = 8;
51static const size_t SIZE_T_TWICE = 2;53static const size_t SIZE_T_TWICE = 2;
@@ -173,7 +175,7 @@ static bool Checkformat(const aclTensor* self, const aclTensor* out)
173static bool CheckPadForFp8(const aclTensor* self, int& signSymbol)175static bool CheckPadForFp8(const aclTensor* self, int& signSymbol)
174{176{
175 // self的数据类型为fp8时,pad数组中不能有负数,StridedSlice不支持FLOAT8_E8M0类型177 // self的数据类型为fp8时,pad数组中不能有负数,StridedSlice不支持FLOAT8_E8M0类型
176- if (CheckType(self->GetDataType(), DTYPE_SUPPORT_FP8_LIST) && (signSymbol & NEGETIVE) == NEGETIVE) {178+ if (CheckType(self->GetDataType(), DTYPE_SUPPORT_FP8_FP4_LIST) && (signSymbol & NEGETIVE) == NEGETIVE) {
177 OP_LOGE(ACLNN_ERR_PARAM_INVALID, "For fp8 data type, pad array cannot contain negative values.");179 OP_LOGE(ACLNN_ERR_PARAM_INVALID, "For fp8 data type, pad array cannot contain negative values.");
178 return false;180 return false;
179 }181 }
@@ -302,21 +304,22 @@ static aclnnStatus DoPadV3(
302 CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);304 CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
303 }305 }
304 306 
305- if (CheckType(self->GetDataType(), DTYPE_SUPPORT_FP8_LIST)) {307+ if (CheckType(self->GetDataType(), DTYPE_SUPPORT_FP8_FP4_LIST)) {
306 const uint8_t* valueData = reinterpret_cast<const uint8_t*>(value->GetData());308 const uint8_t* valueData = reinterpret_cast<const uint8_t*>(value->GetData());
307 size_t valueDataSize = op::TypeSize(value->GetDataType());309 size_t valueDataSize = op::TypeSize(value->GetDataType());
310+ auto valueTensor = executor->ConvertToTensor(value, self->GetDataType());
308 for (size_t i = 0; i < valueDataSize; i++) {311 for (size_t i = 0; i < valueDataSize; i++) {
309 uint8_t valueDataIdx = valueData[i];312 uint8_t valueDataIdx = valueData[i];
310- CHECK_COND(valueDataIdx == 0, ACLNN_ERR_PARAM_INVALID, "Fp8 only support pad constant value 0.");313+ CHECK_COND(valueDataIdx == 0 || valueTensor != nullptr, ACLNN_ERR_PARAM_INVALID,
314+ "Fp8/Fp4 only support pad constant value 0.");
311 }315 }
316+ CHECK_RET(valueTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);
312 // 调用l0算子PadV3进行计算317 // 调用l0算子PadV3进行计算
313- (*padV3Result) = l0op::PadV3(selfCasted, padTensor, nullptr, MODE, true, executor);318+ (*padV3Result) = l0op::PadV3(selfCasted, padTensor, valueTensor, MODE, true, executor);
314- CHECK_RET(padV3Result != nullptr, ACLNN_ERR_INNER_NULLPTR);
315 CHECK_RET((*padV3Result) != nullptr, ACLNN_ERR_INNER_NULLPTR);319 CHECK_RET((*padV3Result) != nullptr, ACLNN_ERR_INNER_NULLPTR);
320+ // 将value转换为tensor,并且数据类型转换为self的数据类型
316 return ACLNN_SUCCESS;321 return ACLNN_SUCCESS;
317 }322 }
318- 
319- // 将value转换为tensor,并且数据类型转换为self的数据类型
320 auto valueTensor = executor->ConvertToTensor(value, self->GetDataType());323 auto valueTensor = executor->ConvertToTensor(value, self->GetDataType());
321 CHECK_RET(valueTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);324 CHECK_RET(valueTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);
322 325 
@@ -423,7 +426,7 @@ aclnnStatus aclnnConstantPadNdGetWorkspaceSize(
423 }426 }
424 427 
425 // 空Tensor处理428 // 空Tensor处理
426- if (self->IsEmpty() && !CheckType(self->GetDataType(), DTYPE_SUPPORT_FP8_LIST)) {429+ if (self->IsEmpty() && !CheckType(self->GetDataType(), DTYPE_SUPPORT_FP8_FP4_LIST)) {
427 ret = HandleSelfEmpty(value, out, uniqueExecutor.get());430 ret = HandleSelfEmpty(value, out, uniqueExecutor.get());
428 CHECK_RET(ret == ACLNN_SUCCESS, ret);431 CHECK_RET(ret == ACLNN_SUCCESS, ret);
429 432 
@@ -50,7 +50,8 @@ static const std::initializer_list<op::DataType> CONSTANT_PAD_REGBASE_AICORE_DTY
50 op::DataType::DT_INT16, op::DataType::DT_UINT16, op::DataType::DT_FLOAT16,50 op::DataType::DT_INT16, op::DataType::DT_UINT16, op::DataType::DT_FLOAT16,
51 op::DataType::DT_BF16, op::DataType::DT_INT8, op::DataType::DT_UINT8,51 op::DataType::DT_BF16, op::DataType::DT_INT8, op::DataType::DT_UINT8,
52 op::DataType::DT_BOOL, op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2,52 op::DataType::DT_BOOL, op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2,
53- op::DataType::DT_FLOAT8_E4M3FN, op::DataType::DT_FLOAT8_E8M0};53+ op::DataType::DT_FLOAT8_E4M3FN, op::DataType::DT_FLOAT8_E8M0, op::DataType::DT_FLOAT4_E2M1,
54+ op::DataType::DT_FLOAT4_E1M2};
54 55 
55static const std::initializer_list<op::DataType> REPLICATION_PAD_REGBASE_AICORE_DTYPE_SUPPORT_LIST = {56static const std::initializer_list<op::DataType> REPLICATION_PAD_REGBASE_AICORE_DTYPE_SUPPORT_LIST = {
56 op::DataType::DT_INT64, op::DataType::DT_UINT64, op::DataType::DT_DOUBLE,57 op::DataType::DT_INT64, op::DataType::DT_UINT64, op::DataType::DT_DOUBLE,
@@ -26,9 +26,12 @@ namespace ge {
26* @li x: A Tensor. Must be one of the following types: float16, bfloat16,26* @li x: A Tensor. Must be one of the following types: float16, bfloat16,
27* float32, double, int32, uint8, int16, int8, complex64, int64,27* float32, double, int32, uint8, int16, int8, complex64, int64,
28* qint8, quint8, qint32, qint16, quint16, uint16, complex128, uint32, uint64, bool,28* qint8, quint8, qint32, qint16, quint16, uint16, complex128, uint32, uint64, bool,
29-* hifloat8, float8_e5m2, float8_e4m3fn, float8_e8m0.29+* hifloat8, float8_e5m2, float8_e4m3fn, float8_e8m0, float4_e2m1, float4_e1m2.
30* @li paddings: A Tensor of type int32 or int64, specify the padding sizes.30* @li paddings: A Tensor of type int32 or int64, specify the padding sizes.
31* The size of paddings should be twice of the x shape size.31* The size of paddings should be twice of the x shape size.
32+* If the type of x is float4_e2m1 or float4_e1m2, paddings values should be even number.
33+* If the type of x is hifloat8, float8_e5m2, float8_e4m3fn or float8_e8m0,
34+* paddings values should be non-negative integers.
W

其他数据类型可以pad负数?为什么这几个有区别

likedislike
pengyiming7
4月27日 评论:
32* @li constant_values: A optional Tensor, dtype same as "x".35* @li constant_values: A optional Tensor, dtype same as "x".
33* Is used only in "constant" mode.36* Is used only in "constant" mode.
34 37 
@@ -60,10 +63,10 @@ namespace ge {
60* Compatible with ONNX operator Pad.63* Compatible with ONNX operator Pad.
61*/64*/
62REG_OP(PadV3)65REG_OP(PadV3)
63- .INPUT(x, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0}))66+ .INPUT(x, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0, DT_FLOAT4_E2M1, DT_FLOAT4_E1M2}))
64 .INPUT(paddings, TensorType::IndexNumberType())67 .INPUT(paddings, TensorType::IndexNumberType())
65- .OPTIONAL_INPUT(constant_values, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0}))68+ .OPTIONAL_INPUT(constant_values, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0, DT_FLOAT4_E2M1, DT_FLOAT4_E1M2}))
66- .OUTPUT(y, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0}))69+ .OUTPUT(y, TensorType({TensorType::BasicType(), DT_BOOL, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT8_E8M0, DT_FLOAT4_E2M1, DT_FLOAT4_E1M2}))
67 .ATTR(mode, String, "constant")70 .ATTR(mode, String, "constant")
68 .ATTR(paddings_contiguous, Bool, true)71 .ATTR(paddings_contiguous, Bool, true)
69 .OP_END_FACTORY_REG(PadV3)72 .OP_END_FACTORY_REG(PadV3)
@@ -1037,6 +1037,9 @@ ge::graphStatus PadACTiling::GetShapesAndDtypes()
1037 Ops::Base::ToString(paramsDtype_).c_str());1037 Ops::Base::ToString(paramsDtype_).c_str());
1038 return ge::GRAPH_FAILED;1038 return ge::GRAPH_FAILED;
1039 }1039 }
1040+ if(paramsDtype_ == ge::DT_FLOAT4_E1M2 || paramsDtype_ == ge::DT_FLOAT4_E2M1) {
1041+ dtypeBytes_ = GetSizeByDataType(ge::DT_INT8);
1042+ }
1040 }1043 }
1041 return ge::GRAPH_SUCCESS;1044 return ge::GRAPH_SUCCESS;
1042}1045}
@@ -1133,6 +1136,70 @@ ge::graphStatus PadACTiling::Init()
1133 return ge::GRAPH_SUCCESS;1136 return ge::GRAPH_SUCCESS;
1134}1137}
1135 1138 
1139+ge::graphStatus PadACTiling::Fp8Fp4ValidatePaddings()
1140+{
1141+ OP_LOGD(context_, "Start PadACTiling Fp8Fp4ValidatePaddings.");
1142+ // fp8/fp4 输入数据类型时,pad数组中不能有负数
1143+ size_t frontDimNum = paddings_.padFront.GetDimNum();
1144+ for (size_t i = 0; i < frontDimNum; ++i) {
1145+ int64_t frontValue = paddings_.padFront.GetDim(i);
1146+ if (frontValue < 0) {
1147+ OP_LOGD(context_, "Fp8Fp4ValidatePaddings padFront contains negative value at index %zu: %ld", i, frontValue);
1148+ return ge::GRAPH_FAILED;
1149+ }
1150+ }
1151+ 
1152+ size_t backDimNum = paddings_.padBack.GetDimNum();
1153+ for (size_t i = 0; i < backDimNum; ++i) {
1154+ int64_t backValue = paddings_.padBack.GetDim(i);
1155+ if (backValue < 0) {
1156+ OP_LOGD(context_, "Fp8Fp4ValidatePaddings padBack contains negative value at index %zu: %ld", i, backValue);
1157+ return ge::GRAPH_FAILED;
1158+ }
1159+ }
1160+ return ge::GRAPH_SUCCESS;
1161+}
1162+ 
1163+ge::graphStatus PadACTiling::Fp4ValidateInShape()
1164+{
1165+ OP_LOGD(context_, "Start PadACTiling Fp4ValidateInShape.");
1166+ // fp4 输入数据类型时,输入数据的最后一维度shape为偶数
1167+ if (tilingData_->inShape[dimNum_ - 1] % 2 != 0){
1168+ OP_LOGD(context_, "Fp4 input dimension is not even number, please check");
1169+ return ge::GRAPH_FAILED;
1170+ }
1171+ return ge::GRAPH_SUCCESS;
1172+}
1173+ 
1174+ge::graphStatus PadACTiling::Fp4ValidatePaddings()
1175+{
1176+ OP_LOGD(context_, "Start PadACTiling Fp4ValidatePaddings.");
1177+ // fp4 输入数据类型,左右pad的最后一维均为偶数
1178+ size_t frontDimNum = paddings_.padFront.GetDimNum();
1179+ int64_t frontValue = paddings_.padFront.GetDim(frontDimNum - 1);
1180+ if (frontValue % 2 != 0) {
1181+ OP_LOGD(context_, "Fp4ValidatePaddings padFront last dimension is not even: %ld", frontValue);
1182+ return ge::GRAPH_FAILED;
1183+ }
1184+ 
1185+ size_t backDimNum = paddings_.padBack.GetDimNum();
1186+ int64_t backValue = paddings_.padBack.GetDim(backDimNum - 1);
1187+ if (backValue % 2 != 0) {
1188+ OP_LOGD(context_, "Fp4ValidatePaddings padBack last dimension is not even: %ld", backValue);
1189+ return ge::GRAPH_FAILED;
1190+ }
1191+ 
1192+ return ge::GRAPH_SUCCESS;
1193+}
1194+ 
1195+void PadACTiling::Fp4TilingData()
1196+{
1197+ OP_LOGD(context_, "Start PadACTiling Fp4TilingData.");
1198+ tilingData_->inShape[dimNum_ - 1] /= 2;
1199+ tilingData_->leftPad[dimNum_ - 1] /= 2;
1200+ rightPad_[dimNum_ - 1] /= 2;
1201+}
1202+ 
1136ge::graphStatus PadACTiling::DoTilingModeEdge()1203ge::graphStatus PadACTiling::DoTilingModeEdge()
1137{1204{
1138 OP_CHECK_IF(1205 OP_CHECK_IF(
@@ -1213,9 +1280,32 @@ ge::graphStatus PadACTiling::DoTilingModeCircular()
1213 1280 
1214ge::graphStatus PadACTiling::DoTilingModeConstant()1281ge::graphStatus PadACTiling::DoTilingModeConstant()
1215{1282{
1283+ if (paramsDtype_ == ge::DT_HIFLOAT8 ||
1284+ paramsDtype_ == ge::DT_FLOAT8_E5M2 ||
1285+ paramsDtype_ == ge::DT_FLOAT8_E4M3FN ||
1286+ paramsDtype_ == ge::DT_FLOAT8_E8M0 ||
1287+ paramsDtype_ == ge::DT_FLOAT4_E2M1 ||
1288+ paramsDtype_ == ge::DT_FLOAT4_E1M2) {
1289+ // fp8/fp4 输入数据类型时,pad数组中不能有负数
1290+ OP_CHECK_IF(Fp8Fp4ValidatePaddings() == ge::GRAPH_FAILED,
1291+ OP_LOGE(context_, "PadACTiling Fp8Fp4ValidatePaddings error."),
1292+ return ge::GRAPH_FAILED);
1293+ }
1216 OP_CHECK_IF(1294 OP_CHECK_IF(
1217 DimensionCollapse() == ge::GRAPH_FAILED, OP_LOGE(context_, "PadACTiling Constant Collapse error."),1295 DimensionCollapse() == ge::GRAPH_FAILED, OP_LOGE(context_, "PadACTiling Constant Collapse error."),
1218 return ge::GRAPH_FAILED);1296 return ge::GRAPH_FAILED);
1297+ if (paramsDtype_ == ge::DT_FLOAT4_E2M1 || paramsDtype_ == ge::DT_FLOAT4_E1M2) {
1298+ // fp4 输入数据类型时,输入数据的最后一维shape为偶数
1299+ OP_CHECK_IF(Fp4ValidateInShape() == ge::GRAPH_FAILED,
1300+ OP_LOGE(context_, "PadACTiling Fp4ValidateInShape error."),
1301+ return ge::GRAPH_FAILED);
1302+ // fp4 时,左右pad的padding num为偶数,最后一维
1303+ OP_CHECK_IF(Fp4ValidatePaddings() == ge::GRAPH_FAILED,
1304+ OP_LOGE(context_, "PadACTiling Fp4ValidatePaddings error."),
1305+ return ge::GRAPH_FAILED);
1306+ // 尾轴//2
1307+ Fp4TilingData();
1308+ }
1219 OP_CHECK_IF(1309 OP_CHECK_IF(
1220 ComputeAfterPaddingsAndStrides() == ge::GRAPH_FAILED,1310 ComputeAfterPaddingsAndStrides() == ge::GRAPH_FAILED,
1221 OP_LOGE(context_, "PadACTiling Constant ComputeAfterPaddingsAndStrides error."), return ge::GRAPH_FAILED);1311 OP_LOGE(context_, "PadACTiling Constant ComputeAfterPaddingsAndStrides error."), return ge::GRAPH_FAILED);
@@ -112,6 +112,11 @@ private:
112 bool IsCutLastDim();112 bool IsCutLastDim();
113 void TilingInfoTuneForNormal(uint64_t lastShapeSizeAlign, uint64_t tilingBranch);113 void TilingInfoTuneForNormal(uint64_t lastShapeSizeAlign, uint64_t tilingBranch);
114 114 
115+ ge::graphStatus Fp8Fp4ValidatePaddings();
116+ ge::graphStatus Fp4ValidateInShape();
117+ ge::graphStatus Fp4ValidatePaddings();
118+ void Fp4TilingData();
119+ 
115public:120public:
116 bool isPadV3_{false};121 bool isPadV3_{false};
117 bool isMirrorPad_{false};122 bool isMirrorPad_{false};