Star1296
1410
代码
代码
Issues108
Pull Requests213
Actions
项目讨论
Wiki
分析
Star1296
1410
已合并
fix: TensorRedirect 交付件规范对齐(issue #4553)，并删除 AddV2 不可达的混合 dtype 组合 #4505
raoliang_sac创建于 19 天前
fix: TensorRedirect 交付件规范对齐(issue #4553)，并删除 AddV2 不可达的混合 dtype 组合 #4505
已合并
raoliang_sac创建于 19 天前
从raoliang_sac/ops-mathfix/review-and-issue-4553合入到cann/ops-mathmaster
共 17 个文件变更+327-373
Aconversion/tensor_redirect/op_graph/tensor_redirect_graph_infer.cpp+35-0
@@ -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 tensor_redirect_graph_infer.cpp
13+ * \brief tensor_redirect operator graph infer resource
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+namespace ops {
20+static constexpr size_t TENSOR_REDIRECT_IN_X_IDX = 0;
21+static constexpr size_t TENSOR_REDIRECT_OUT_Y_IDX = 0;
22+ 
23+// output_x 与 x 同 dtype（same-as-input）
24+static ge::graphStatus InferDataType4TensorRedirect(gert::InferDataTypeContext* context)
25+{
26+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataType4TensorRedirect");
27+ auto ret = context->SetOutputDataType(TENSOR_REDIRECT_OUT_Y_IDX,
28+ context->GetInputDataType(TENSOR_REDIRECT_IN_X_IDX));
29+ OP_LOGD(context->GetNodeName(), "End to do InferDataType4TensorRedirect");
30+ return ret;
31+}
32+ 
33+IMPL_OP(TensorRedirect).InferDataType(InferDataType4TensorRedirect);
34+ 
35+} // namespace ops
Mconversion/tensor_redirect/op_graph/tensor_redirect_proto.h+3-0
@@ -30,12 +30,15 @@ namespace ge {
30* output_x: A ND Tensor. Has the same dtype and format as "x". \n30* output_x: A ND Tensor. Has the same dtype and format as "x". \n
31 31 
32*/32*/
33+#ifndef OPS_PROTO_DEF_TENSORREDIRECT
34+#define OPS_PROTO_DEF_TENSORREDIRECT
33REG_OP(TensorRedirect)35REG_OP(TensorRedirect)
34 .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_INT8, DT_INT32, DT_UINT8, DT_INT64, DT_INT16, DT_UINT16, DT_UINT64,36 .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_INT8, DT_INT32, DT_UINT8, DT_INT64, DT_INT16, DT_UINT16, DT_UINT64,
35 DT_UINT32, DT_BF16}))37 DT_UINT32, DT_BF16}))
36 .OUTPUT(output_x, TensorType({DT_FLOAT16, DT_FLOAT, DT_INT8, DT_INT32, DT_UINT8, DT_INT64, DT_INT16, DT_UINT16,38 .OUTPUT(output_x, TensorType({DT_FLOAT16, DT_FLOAT, DT_INT8, DT_INT32, DT_UINT8, DT_INT64, DT_INT16, DT_UINT16,
37 DT_UINT64, DT_UINT32, DT_BF16}))39 DT_UINT64, DT_UINT32, DT_BF16}))
38 .OP_END_FACTORY_REG(TensorRedirect)40 .OP_END_FACTORY_REG(TensorRedirect)
41+#endif
39 42 
40} // namespace ge43} // namespace ge
41#endif // OP_PROTO_TENSOR_REDIRECT_H_44#endif // OP_PROTO_TENSOR_REDIRECT_H_
Mconversion/tensor_redirect/op_host/arch35/tensor_redirect_tiling_arch35.cpp+18-0
@@ -90,6 +90,16 @@ static ge::graphStatus CheckTensorRedirectShape(const gert::TilingContext* conte
90 "rank must be within [1, 8]"),90 "rank must be within [1, 8]"),
91 return ge::GRAPH_FAILED);91 return ge::GRAPH_FAILED);
92 92 
93+ // CheckDim: Tiling 是外部输入边界，concrete shape 的每一维必须非负；
94+ // -1/-2 等动态占位符或非法负值不能进入切分计算
95+ for (size_t i = 0; i < xRank; ++i) {
96+ OP_CHECK_IF(
97+ xShape.GetDim(i) < 0,
98+ OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON("TensorRedirect", "x", std::to_string(xShape.GetDim(i)).c_str(),
99+ "every dim of the concrete shape must be non-negative"),
100+ return ge::GRAPH_FAILED);
101+ }
102+ 
93 auto yShapePtr = context->GetOutputShape(INDEX_OUTPUT_X);103 auto yShapePtr = context->GetOutputShape(INDEX_OUTPUT_X);
94 OP_CHECK_NULL_WITH_CONTEXT(context, yShapePtr);104 OP_CHECK_NULL_WITH_CONTEXT(context, yShapePtr);
95 auto yShape = yShapePtr->GetStorageShape();105 auto yShape = yShapePtr->GetStorageShape();
@@ -228,6 +238,14 @@ static ge::graphStatus Tiling4TensorRedirect(gert::TilingContext* context)
228 // 1D 线性展平，不解释 stride/rank238 // 1D 线性展平，不解释 stride/rank
229 int64_t numel = xShapePtr->GetStorageShape().GetShapeSize();239 int64_t numel = xShapePtr->GetStorageShape().GetShapeSize();
230 240 
241+ // 溢出防护：GetShapeSize() 在维度乘积溢出 int64_t 时返回 kInvalidDimValue，不会自行报错。
242+ // 必须在 numel == 0 判断之前拦截，否则负的 numel 会穿透到 DoTiling 产生 usedCoreNum == 0。
243+ OP_CHECK_IF(numel < 0,
244+ OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON("TensorRedirect", "x",
245+ Ops::Base::ToString(xShapePtr->GetStorageShape()).c_str(),
246+ "the product of all dims overflows int64_t"),
247+ return ge::GRAPH_FAILED);
248+ 
231 // 空 Tensor 早返回249 // 空 Tensor 早返回
232 if (numel == 0) {250 if (numel == 0) {
233 OP_LOGD(context, "TensorRedirect: empty tensor, skip kernel computation.");251 OP_LOGD(context, "TensorRedirect: empty tensor, skip kernel computation.");
Mconversion/tensor_redirect/op_host/tensor_redirect_def.cpp+3-1
@@ -41,7 +41,9 @@ public:
41 .DynamicRankSupportFlag(true) // dynamic rank41 .DynamicRankSupportFlag(true) // dynamic rank
42 .DynamicShapeSupportFlag(true) // 动态 shape42 .DynamicShapeSupportFlag(true) // 动态 shape
43 .NeedCheckSupportFlag(false)43 .NeedCheckSupportFlag(false)
44- .PrecisionReduceFlag(true) // 恒等拷贝无计算，不改变 bit-exact 口径44+ // 恒等拷贝语义要求逐 bit 相等（含负零/NaN payload/非规格化数），
45+ // 不允许 allow_mix_precision 下把 FP32 节点降为 FP16
46+ .PrecisionReduceFlag(false)
45 .ExtendCfgInfo("op.pattern", "formatAgnostic")47 .ExtendCfgInfo("op.pattern", "formatAgnostic")
46 .ExtendCfgInfo("opFile.value", "tensor_redirect_apt");48 .ExtendCfgInfo("opFile.value", "tensor_redirect_apt");
47 this->AICore().AddConfig("ascend950", aicoreConfig);49 this->AICore().AddConfig("ascend950", aicoreConfig);
Mconversion/tensor_redirect/op_host/tensor_redirect_infershape.cpp+3-13
@@ -17,17 +17,7 @@
17 17 
18using namespace ge;18using namespace ge;
19namespace ops {19namespace ops {
20-static constexpr size_t TENSOR_REDIRECT_IN_X_IDX = 0;20+// output_x 与 x 同 shape（same-as-input）
21-static constexpr size_t TENSOR_REDIRECT_OUT_Y_IDX = 0;21+// InferDataType 仅图场景使用，交付在 op_graph/tensor_redirect_graph_infer.cpp
22- 22+IMPL_OP_INFERSHAPE(TensorRedirect).InferShape(Ops::Base::InferShape4Elewise);
23-static ge::graphStatus InferDataType4TensorRedirect(gert::InferDataTypeContext* context)
24-{
25- context->SetOutputDataType(TENSOR_REDIRECT_OUT_Y_IDX, context->GetInputDataType(TENSOR_REDIRECT_IN_X_IDX));
26- return ge::GRAPH_SUCCESS;
27-}
28- 
29-// output_x 与 x 同 shape/dtype（same-as-input）
30-IMPL_OP_INFERSHAPE(TensorRedirect)
31- .InferShape(Ops::Base::InferShape4Elewise)
32- .InferDataType(InferDataType4TensorRedirect);
33} // namespace ops23} // namespace ops
Mconversion/tensor_redirect/tests/ut/op_host/arch35/test_tensor_redirect_tiling_arch35.cpp+31-0
@@ -13,6 +13,7 @@
13 * \brief TensorRedirect op_host Tiling UT13 * \brief TensorRedirect op_host Tiling UT
14 */14 */
15 15 
16+#include <cstdint>
16#include <cstring>17#include <cstring>
17#include <vector>18#include <vector>
18 19 
@@ -286,6 +287,36 @@ TEST_F(TensorRedirectTilingTest, tiling_check_shape_mismatch_rank_same_numel_fai
286 ExecuteTestCase(para, ge::GRAPH_FAILED, EXPECT_TILING_KEY, std::vector<size_t>{});287 ExecuteTestCase(para, ge::GRAPH_FAILED, EXPECT_TILING_KEY, std::vector<size_t>{});
287}288}
288 289 
290+// concrete shape 含 -1（动态占位符穿透到 Tiling）→ 必须拦截，
291+// 否则 GetShapeSize() 返回负 numel，绕过 numel==0 守卫并算出 usedCoreNum==0
292+TEST_F(TensorRedirectTilingTest, tiling_check_negative_dim_minus1_failed)
293+{
294+ gert::StorageShape shape = {{-1, 32}, {-1, 32}};
295+ ExecuteTestCase(MakePara(shape, ge::DT_FLOAT16), ge::GRAPH_FAILED, EXPECT_TILING_KEY, std::vector<size_t>{});
296+}
297+ 
298+// concrete shape 含 -2（unknown rank 占位符）→ 必须拦截
299+TEST_F(TensorRedirectTilingTest, tiling_check_negative_dim_minus2_failed)
300+{
301+ gert::StorageShape shape = {{-2}, {-2}};
302+ ExecuteTestCase(MakePara(shape, ge::DT_FLOAT16), ge::GRAPH_FAILED, EXPECT_TILING_KEY, std::vector<size_t>{});
303+}
304+ 
305+// 负维出现在非首维 → 逐维校验须覆盖全部维度，而不只是 dim0
306+TEST_F(TensorRedirectTilingTest, tiling_check_negative_dim_at_last_axis_failed)
307+{
308+ gert::StorageShape shape = {{4, 8, -1}, {4, 8, -1}};
309+ ExecuteTestCase(MakePara(shape, ge::DT_FLOAT16), ge::GRAPH_FAILED, EXPECT_TILING_KEY, std::vector<size_t>{});
310+}
311+ 
312+// 维度乘积溢出 int64_t：GetShapeSize() 返回 kInvalidDimValue(INT64_MIN)，
313+// 该哨兵不是错误码，必须由 Tiling 显式拦截
314+TEST_F(TensorRedirectTilingTest, tiling_check_shape_size_overflow_failed)
315+{
316+ gert::StorageShape shape = {{INT64_MAX, 2}, {INT64_MAX, 2}};
317+ ExecuteTestCase(MakePara(shape, ge::DT_FLOAT16), ge::GRAPH_FAILED, EXPECT_TILING_KEY, std::vector<size_t>{});
318+}
319+ 
289// 三、多核切分核心路径320// 三、多核切分核心路径
290 321 
291// #4：[1048577] 触发提核优化322// #4：[1048577] 触发提核优化
Mdocs/zh/op_list.md+1-1
@@ -114,7 +114,7 @@
114 <td>×</td>114 <td>×</td>
115 <td>√</td>115 <td>√</td>
116 <td>AI Core</td>116 <td>AI Core</td>
117- <td>对两个输入张量执行逐元素加法，兼容TensorFlow AddV2语义，bool输入按logical_or处理，支持广播。</td>117+ <td>对两个输入张量执行逐元素加法，兼容TensorFlow AddV2语义，支持广播与空Tensor。</td>
118 </tr>118 </tr>
119 <tr>119 <tr>
120 <td>math</td>120 <td>math</td>
Mmath/add_v2/README.md+4-11
@@ -54,32 +54,25 @@ x1 shape (3, 1), x2 shape (1, 4) -> y shape (3, 4)
54 <tr>54 <tr>
55 <td>x2</td>55 <td>x2</td>
56 <td>输入</td>56 <td>输入</td>
57- <td>加法运算的第二个输入张量，公式中的x2_i。</td>57+ <td>加法运算的第二个输入张量，公式中的x2_i。数据类型需与x1一致。</td>
58 <td>BFLOAT16、FLOAT16、FLOAT、INT32、INT16、UINT8、INT8、INT64、COMPLEX64</td>58 <td>BFLOAT16、FLOAT16、FLOAT、INT32、INT16、UINT8、INT8、INT64、COMPLEX64</td>
59 <td>ND</td>59 <td>ND</td>
60 </tr>60 </tr>
61 <tr>61 <tr>
62 <td>y</td>62 <td>y</td>
63 <td>输出</td>63 <td>输出</td>
64- <td>加法运算的输出张量，公式中的y_i。同类型输入时输出类型与输入一致；混合精度输入时输出为类型提升后的结果。</td>64+ <td>加法运算的输出张量，公式中的y_i。数据类型与x1一致。</td>
65 <td>BFLOAT16、FLOAT16、FLOAT、INT32、INT16、UINT8、INT8、INT64、COMPLEX64</td>65 <td>BFLOAT16、FLOAT16、FLOAT、INT32、INT16、UINT8、INT8、INT64、COMPLEX64</td>
66 <td>ND</td>66 <td>ND</td>
67 </tr>67 </tr>
68 </tbody></table>68 </tbody></table>
69 69 
70-混合精度组合（x1 dtype, x2 dtype -> y dtype）：
71- 
72-| x1 | x2 | y |
73-|-------------|-------------|--------|
74-| FLOAT16 | FLOAT | FLOAT |
75-| FLOAT | FLOAT16 | FLOAT |
76-| BFLOAT16 | FLOAT | FLOAT |
77-| FLOAT | BFLOAT16 | FLOAT |
78- 
79## 约束说明70## 约束说明
80 71 
81- 输入x1和x2的shape需满足广播规则。72- 输入x1和x2的shape需满足广播规则。
73+- 输入x1和x2的数据类型需相同，不支持混合数据类型输入。
82- 输入数据类型需在支持列表内，不支持DOUBLE、COMPLEX128、BOOL、COMPLEX32。74- 输入数据类型需在支持列表内，不支持DOUBLE、COMPLEX128、BOOL、COMPLEX32。
75+- 支持空Tensor。当输出y的元素个数为0时，算子不下发有效计算，直接返回成功。
83 76 
84## 调用说明77## 调用说明
85 78 
Mmath/add_v2/op_graph/add_v2_graph_infer.cpp+16-106
@@ -17,119 +17,29 @@
17 17 
18using namespace ge;18using namespace ge;
19namespace ops {19namespace ops {
20-namespace promote_type_for_add_v2_detail {20+static constexpr size_t ADD_V2_IN_X1_IDX = 0;
21-constexpr auto u1 = ge::DataType::DT_UINT8;21+static constexpr size_t ADD_V2_IN_X2_IDX = 1;
22-constexpr auto i1 = ge::DataType::DT_INT8;22+static constexpr size_t ADD_V2_OUT_Y_IDX = 0;
23-constexpr auto i2 = ge::DataType::DT_INT16;
24-constexpr auto i4 = ge::DataType::DT_INT32;
25-constexpr auto i8 = ge::DataType::DT_INT64;
26-constexpr auto f2 = ge::DataType::DT_FLOAT16;
27-constexpr auto f4 = ge::DataType::DT_FLOAT;
28-constexpr auto f8 = ge::DataType::DT_DOUBLE;
29-constexpr auto c2 = ge::DataType::DT_COMPLEX32;
30-constexpr auto c4 = ge::DataType::DT_COMPLEX64;
31-constexpr auto c8 = ge::DataType::DT_COMPLEX128;
32-constexpr auto b1 = ge::DataType::DT_BOOL;
33-constexpr auto bf = ge::DataType::DT_BF16;
34-constexpr auto ud = ge::DataType::DT_UNDEFINED;
35-// @formatter:off
36-static constexpr ge::DataType
37- kPromoteTypesLookup[static_cast<int>(ge::DataType::DT_MAX)][static_cast<int>(ge::DataType::DT_MAX)] = {
38- /* f4 f2 i1 i4 u1 xx i2 u2 u4 i8 u8 f8 b1 sv d1 D1 c4 c8 q1 q2 q4 Q1 Q2 rs sr
39- du va bf, ud t4 T1 t2 T2 c2*/
40- /* f4 0 */ {f4, f4, f4, f4, f4, ud, f4, ud, ud, f4, ud, f8, f4, ud, ud, ud, c4,
41- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, f4, ud, ud, ud, ud, ud, c4},
42- /* f2 1 */ {f4, f2, f2, f2, f2, ud, f2, ud, ud, f2, ud, f8, f2, ud, ud, ud, c4,
43- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, f4, ud, ud, ud, ud, ud, c2},
44- /* i1 2 */ {f4, f2, i1, i4, i2, ud, i2, ud, ud, i8, ud, f8, i1, ud, ud, ud, c4,
45- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, bf, ud, ud, ud, ud, ud, c2},
46- /* i4 3 */ {f4, f2, i4, i4, i4, ud, i4, ud, ud, i8, ud, f8, i4, ud, ud, ud, c4,
47- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, bf, ud, ud, ud, ud, ud, c2},
48- /* u1 4 */ {f4, f2, i2, i4, u1, ud, i2, ud, ud, i8, ud, f8, u1, ud, ud, ud, c4,
49- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, bf, ud, ud, ud, ud, ud, c2},
50- /* xx 5 */ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
51- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
52- /* i2 6 */ {f4, f2, i2, i4, i2, ud, i2, ud, ud, i8, ud, f8, i2, ud, ud, ud, c4,
53- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, bf, ud, ud, ud, ud, ud, c2},
54- /* u2 7 */ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, c4,
55- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
56- /* u4 8 */ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, c4,
57- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
58- /* i8 9 */ {f4, f2, i8, i8, i8, ud, i8, ud, ud, i8, ud, f8, i8, ud, ud, ud, c4,
59- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, bf, ud, ud, ud, ud, ud, c2},
60- /* u8 10*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, c8,
61- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
62- /* f8 11*/ {f8, f8, f8, f8, f8, ud, f8, ud, ud, f8, ud, f8, f8, ud, ud, ud, c8,
63- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, f8, ud, ud, ud, ud, ud, c8},
64- /* b1 12*/ {f4, f2, i1, i4, u1, ud, i2, ud, ud, i8, ud, f8, b1, ud, ud, ud, c4,
65- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, bf, ud, ud, ud, ud, ud, c2},
66- /* sv 13*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
67- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
68- /* d1 14*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
69- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
70- /* D1 15*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
71- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
72- /* c4 16*/ {c4, c4, c4, c4, c4, ud, c4, ud, ud, c4, ud, c8, c4, ud, ud, ud, c4,
73- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, c4, ud, ud, ud, ud, ud, c4},
74- /* c8 17*/ {c8, c8, c8, c8, c8, ud, c8, ud, ud, c8, ud, c8, c8, ud, ud, ud, c8,
75- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, c8, ud, ud, ud, ud, ud, c8},
76- /* q1 18*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
77- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
78- /* q2 19*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
79- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
80- /* q4 20*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
81- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
82- /* Q1 21*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
83- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
84- /* Q2 22*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
85- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
86- /* rs 23*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
87- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
88- /* sr 24*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
89- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
90- /* du 25*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
91- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
92- /* va 26*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
93- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
94- /* bf 27*/ {f4, f4, bf, bf, bf, ud, bf, ud, ud, bf, ud, f8, bf, ud, ud, ud, c4,
95- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, bf, ud, ud, ud, ud, ud, c4},
96- /* ud 28*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
97- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
98- /* t4 29*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
99- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
100- /* T1 30*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
101- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
102- /* t2 31*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
103- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
104- /* T2 32*/ {ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud,
105- ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud, ud},
106- /* c2 33*/ {c4, c2, c2, c2, c2, ud, c2, ud, ud, c2, ud, c8, c2, ud, ud, ud, c4,
107- c8, ud, ud, ud, ud, ud, ud, ud, ud, ud, c4, ud, ud, ud, ud, ud, c2},
108-};
109-// @formatter:on
110- 
111-inline ge::DataType PromoteType(ge::DataType type_a, ge::DataType type_b)
112-{
113- if (type_a < 0 || type_b < 0 || type_a >= ge::DataType::DT_MAX || type_b >= ge::DataType::DT_MAX) {
114- return ge::DataType::DT_UNDEFINED;
115- }
116- if (type_a == type_b) {
117- return type_a;
118- }
119- return kPromoteTypesLookup[static_cast<int>(type_a)][static_cast<int>(type_b)];
120-}
121-} // namespace promote_type_for_add_v2_detail
122 23 
24+// canonical AddV2 契约（canndev op_graph 原型与 CheckTwoInputDtypeSame Verifier）
25+// 要求 x1/x2 同 dtype，y 与 x1 同 dtype；本算子不注册异类型组合，故此处不做类型提升。
123static ge::graphStatus InferDataTypeAddV2(gert::InferDataTypeContext* context)26static ge::graphStatus InferDataTypeAddV2(gert::InferDataTypeContext* context)
124{27{
125 OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeAddV2");28 OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeAddV2");
126 29 
127- const ge::DataType x1DataType = context->GetInputDataType(0);30+ const ge::DataType x1DataType = context->GetInputDataType(ADD_V2_IN_X1_IDX);
128- const ge::DataType x2DataType = context->GetInputDataType(1);31+ const ge::DataType x2DataType = context->GetInputDataType(ADD_V2_IN_X2_IDX);
129- context->SetOutputDataType(0, promote_type_for_add_v2_detail::PromoteType(x1DataType, x2DataType));32+ OP_CHECK_IF(
33+ x1DataType != x2DataType,
34+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(
35+ "AddV2", "x1 and x2", (Ops::Base::ToString(x1DataType) + " and " + Ops::Base::ToString(x2DataType)).c_str(),
36+ "x1 and x2 must have the same dtype"),
37+ return ge::GRAPH_FAILED);
38+ 
39+ auto ret = context->SetOutputDataType(ADD_V2_OUT_Y_IDX, x1DataType);
130 40 
131 OP_LOGD(context->GetNodeName(), "End to do InferDataTypeAddV2");41 OP_LOGD(context->GetNodeName(), "End to do InferDataTypeAddV2");
132- return GRAPH_SUCCESS;42+ return ret;
133}43}
134 44 
135IMPL_OP(AddV2).InferDataType(InferDataTypeAddV2);45IMPL_OP(AddV2).InferDataType(InferDataTypeAddV2);
Mmath/add_v2/op_host/add_v2_def.cpp+9-6
@@ -16,33 +16,36 @@
16 16 
17namespace ops {17namespace ops {
18namespace {18namespace {
19-#define ADD_V2_FORMAT_LIST \19+#define ADD_V2_FORMAT_LIST \
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, \
21- 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}
22} // namespace22} // namespace
23 23 
24class AddV2 : public OpDef {24class AddV2 : public OpDef {
25public:25public:
26 explicit AddV2(const char* name) : OpDef(name)26 explicit AddV2(const char* name) : OpDef(name)
27 {27 {
28+ // 9 组同 dtype 组合。canonical AddV2 Verifier（canndev
29+ // elewise_calculation_ops.cc 的 CheckTwoInputDtypeSame）要求 x1/x2 同
30+ // dtype，异类型组合在 GE 图通路上不可达，故不注册。
28 this->Input("x1")31 this->Input("x1")
29 .ParamType(REQUIRED)32 .ParamType(REQUIRED)
30 .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT16, ge::DT_UINT8, ge::DT_INT8,33 .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT16, ge::DT_UINT8, ge::DT_INT8,
31- ge::DT_INT64, ge::DT_COMPLEX64, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT})34+ ge::DT_INT64, ge::DT_COMPLEX64})
32 .Format(ADD_V2_FORMAT_LIST)35 .Format(ADD_V2_FORMAT_LIST)
33 .UnknownShapeFormat(ADD_V2_FORMAT_LIST)36 .UnknownShapeFormat(ADD_V2_FORMAT_LIST)
34 .AutoContiguous();37 .AutoContiguous();
35 this->Input("x2")38 this->Input("x2")
36 .ParamType(REQUIRED)39 .ParamType(REQUIRED)
37 .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT16, ge::DT_UINT8, ge::DT_INT8,40 .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT16, ge::DT_UINT8, ge::DT_INT8,
38- ge::DT_INT64, ge::DT_COMPLEX64, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})41+ ge::DT_INT64, ge::DT_COMPLEX64})
39 .Format(ADD_V2_FORMAT_LIST)42 .Format(ADD_V2_FORMAT_LIST)
40 .UnknownShapeFormat(ADD_V2_FORMAT_LIST)43 .UnknownShapeFormat(ADD_V2_FORMAT_LIST)
41 .AutoContiguous();44 .AutoContiguous();
42 this->Output("y")45 this->Output("y")
43 .ParamType(REQUIRED)46 .ParamType(REQUIRED)
44 .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT16, ge::DT_UINT8, ge::DT_INT8,47 .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT16, ge::DT_UINT8, ge::DT_INT8,
45- ge::DT_INT64, ge::DT_COMPLEX64, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})48+ ge::DT_INT64, ge::DT_COMPLEX64})
46 .Format(ADD_V2_FORMAT_LIST)49 .Format(ADD_V2_FORMAT_LIST)
47 .UnknownShapeFormat(ADD_V2_FORMAT_LIST)50 .UnknownShapeFormat(ADD_V2_FORMAT_LIST)
48 .AutoContiguous();51 .AutoContiguous();
Mmath/add_v2/op_host/arch35/add_v2_tiling_arch35.cpp+77-36
@@ -35,6 +35,11 @@ namespace optiling {
35 35 
36constexpr int64_t ASCEND_WORKSPACE = 16777216; // 16M36constexpr int64_t ASCEND_WORKSPACE = 16777216; // 16M
37 37 
38+// 自定义模板分支：空 Tensor 走 schMode 999 + userDef 1，见 add_v2_struct_arch35.h
39+constexpr uint64_t ADD_V2_SCH_MODE_EMPTY = 999;
40+constexpr uint64_t ADD_V2_USER_DEF_NORMAL = 0;
41+constexpr uint64_t ADD_V2_USER_DEF_EMPTY = 1;
42+ 
38class AddV2TilingArch35 {43class AddV2TilingArch35 {
39public:44public:
40 explicit AddV2TilingArch35(gert::TilingContext* context) : tilingContext(context) {};45 explicit AddV2TilingArch35(gert::TilingContext* context) : tilingContext(context) {};
@@ -43,13 +48,41 @@ public:
43protected:48protected:
44 ge::graphStatus CalcDtype();49 ge::graphStatus CalcDtype();
45 ge::graphStatus CheckShape() const;50 ge::graphStatus CheckShape() const;
46- bool IsMixedDtype(const ge::DataType& d0, const ge::DataType& d1) const;51+ ge::graphStatus CheckDtype() const;
52+ ge::graphStatus SetWorkspace() const;
53+ ge::graphStatus HandleEmptyTensor() const;
47 54 
48private:55private:
49 ge::DataType inputDtype = ge::DT_UNDEFINED;56 ge::DataType inputDtype = ge::DT_UNDEFINED;
50 gert::TilingContext* tilingContext;57 gert::TilingContext* tilingContext;
51};58};
52 59 
60+ge::graphStatus AddV2TilingArch35::SetWorkspace() const
61+{
62+ size_t* currentWorkspace = tilingContext->GetWorkspaceSizes(1);
63+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, currentWorkspace);
64+ currentWorkspace[0] = static_cast<uint64_t>(ASCEND_WORKSPACE);
65+ return ge::GRAPH_SUCCESS;
66+}
67+ 
68+// 空 Tensor 早返回。ATVOSS 的 BroadcastBaseTiling 在合轴之后会显式拒绝 0 元素
69+// （broadcast_tiling.h: "tensor check is empty, check failed"），不能落到 DoTiling，
70+// 因此这里自己出一份 tiling：blockDim = 1，tilingKey 选自定义分支，kernel 侧直接返回。
71+ge::graphStatus AddV2TilingArch35::HandleEmptyTensor() const
72+{
73+ OP_LOGD(tilingContext, "AddV2: empty tensor, skip kernel computation.");
74+ OP_CHECK_IF(SetWorkspace() != ge::GRAPH_SUCCESS,
75+ OP_LOGE(tilingContext, "AddV2: set workspace failed (empty tensor)."), return ge::GRAPH_FAILED);
76+ 
77+ auto* tilingData = tilingContext->GetTilingData<AddV2EmptyTilingData>();
78+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, tilingData);
79+ tilingData->numel = 0;
80+ 
81+ tilingContext->SetBlockDim(1);
82+ tilingContext->SetTilingKey(GET_TPL_TILING_KEY(ADD_V2_SCH_MODE_EMPTY, ADD_V2_USER_DEF_EMPTY));
83+ return ge::GRAPH_SUCCESS;
84+}
85+ 
53ge::graphStatus AddV2TilingArch35::CalcDtype()86ge::graphStatus AddV2TilingArch35::CalcDtype()
54{87{
55 auto inputDesc = tilingContext->GetInputDesc(0);88 auto inputDesc = tilingContext->GetInputDesc(0);
@@ -69,10 +102,20 @@ ge::graphStatus AddV2TilingArch35::CheckShape() const
69 return ge::GRAPH_SUCCESS;102 return ge::GRAPH_SUCCESS;
70}103}
71 104 
72-bool AddV2TilingArch35::IsMixedDtype(const ge::DataType& d0, const ge::DataType& d1) const105+// 仅注册同 dtype 组合（canonical AddV2 Verifier 亦要求 x1/x2 同 dtype）
106+ge::graphStatus AddV2TilingArch35::CheckDtype() const
73{107{
74- return (d0 == ge::DT_FLOAT16 && d1 == ge::DT_FLOAT) || (d0 == ge::DT_FLOAT && d1 == ge::DT_FLOAT16) ||108+ auto input1Desc = tilingContext->GetInputDesc(1);
75- (d0 == ge::DT_BF16 && d1 == ge::DT_FLOAT) || (d0 == ge::DT_FLOAT && d1 == ge::DT_BF16);109+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, input1Desc);
110+ ge::DataType input1Dtype = input1Desc->GetDataType();
111+ OP_CHECK_IF(input1Dtype != this->inputDtype,
112+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(tilingContext->GetNodeName(), "x1 and x2",
113+ (ge::TypeUtils::DataTypeToSerialString(this->inputDtype) +
114+ " and " + ge::TypeUtils::DataTypeToSerialString(input1Dtype))
115+ .c_str(),
116+ "x1 and x2 must have the same dtype"),
117+ return ge::GRAPH_FAILED);
118+ return ge::GRAPH_SUCCESS;
76}119}
77 120 
78ge::graphStatus AddV2TilingArch35::RunTiling()121ge::graphStatus AddV2TilingArch35::RunTiling()
@@ -87,61 +130,60 @@ ge::graphStatus AddV2TilingArch35::RunTiling()
87 "input shape check failed"),130 "input shape check failed"),
88 return ge::GRAPH_FAILED);131 return ge::GRAPH_FAILED);
89 132 
90- auto input1Desc = tilingContext->GetInputDesc(1);133+ OP_CHECK_IF(CheckDtype() == ge::GRAPH_FAILED,
91- OP_CHECK_NULL_WITH_CONTEXT(tilingContext, input1Desc);134+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(tilingContext->GetNodeName(), "input dtype", "invalid",
92- ge::DataType input1Dtype = input1Desc->GetDataType();135+ "input dtype check failed"),
93- bool isMixedDtype = IsMixedDtype(this->inputDtype, input1Dtype);136+ return ge::GRAPH_FAILED);
137+ 
138+ // y 的元素个数为 0 即空 Tensor（广播规则下 0 只能与 0 或 1 相配，输出空 <=> 有输入为空）。
139+ auto outputY = tilingContext->GetOutputShape(0);
140+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outputY);
141+ int64_t numel = outputY->GetStorageShape().GetShapeSize();
142+ // GetShapeSize() 在维度乘积溢出 int64_t 时返回 kInvalidDimValue（负数）而不报错，
143+ // 必须在 numel == 0 判断之前拦掉，否则负的 numel 会被当成非空穿透到 DoTiling。
144+ OP_CHECK_IF(numel < 0,
145+ OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(tilingContext->GetNodeName(), "y",
146+ Ops::Base::ToString(outputY->GetStorageShape()).c_str(),
147+ "the product of all dims overflows int64_t"),
148+ return ge::GRAPH_FAILED);
149+ if (numel == 0) {
150+ return HandleEmptyTensor();
151+ }
94 152 
95 ge::graphStatus ret = ge::GRAPH_FAILED;153 ge::graphStatus ret = ge::GRAPH_FAILED;
96 uint64_t tilingKey = 0;154 uint64_t tilingKey = 0;
97- if (isMixedDtype && input1Dtype == ge::DT_FLOAT && this->inputDtype == ge::DT_FLOAT16) {155+ if (this->inputDtype == ge::DT_FLOAT16) {
98- BroadcastBaseTiling<AddMixDtypeCompute<half, float>::OpDag> brcBaseTiling(tilingContext);
99- ret = brcBaseTiling.DoTiling();
100- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
101- } else if (isMixedDtype && input1Dtype == ge::DT_FLOAT && this->inputDtype == ge::DT_BF16) {
102- BroadcastBaseTiling<AddMixDtypeCompute<bfloat16_t, float>::OpDag> brcBaseTiling(tilingContext);
103- ret = brcBaseTiling.DoTiling();
104- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
105- } else if (isMixedDtype && this->inputDtype == ge::DT_FLOAT && input1Dtype == ge::DT_FLOAT16) {
106- BroadcastBaseTiling<AddMixDtypeCompute<float, half>::OpDag> brcBaseTiling(tilingContext);
107- ret = brcBaseTiling.DoTiling();
108- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
109- } else if (isMixedDtype && this->inputDtype == ge::DT_FLOAT && input1Dtype == ge::DT_BF16) {
110- BroadcastBaseTiling<AddMixDtypeCompute<float, bfloat16_t>::OpDag> brcBaseTiling(tilingContext);
111- ret = brcBaseTiling.DoTiling();
112- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
113- } else if (this->inputDtype == ge::DT_FLOAT16) {
114 BroadcastBaseTiling<AddWithCastCompute<half>::OpDag> brcBaseTiling(tilingContext);156 BroadcastBaseTiling<AddWithCastCompute<half>::OpDag> brcBaseTiling(tilingContext);
115 ret = brcBaseTiling.DoTiling();157 ret = brcBaseTiling.DoTiling();
116- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());158+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), ADD_V2_USER_DEF_NORMAL);
117 } else if (this->inputDtype == ge::DT_BF16) {159 } else if (this->inputDtype == ge::DT_BF16) {
118 BroadcastBaseTiling<AddWithCastCompute<bfloat16_t>::OpDag> brcBaseTiling(tilingContext);160 BroadcastBaseTiling<AddWithCastCompute<bfloat16_t>::OpDag> brcBaseTiling(tilingContext);
119 ret = brcBaseTiling.DoTiling();161 ret = brcBaseTiling.DoTiling();
120- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());162+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), ADD_V2_USER_DEF_NORMAL);
121 } else if (this->inputDtype == ge::DT_FLOAT) {163 } else if (this->inputDtype == ge::DT_FLOAT) {
122 BroadcastBaseTiling<AddWithCastCompute<float>::OpDag> brcBaseTiling(tilingContext);164 BroadcastBaseTiling<AddWithCastCompute<float>::OpDag> brcBaseTiling(tilingContext);
123 ret = brcBaseTiling.DoTiling();165 ret = brcBaseTiling.DoTiling();
124- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());166+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), ADD_V2_USER_DEF_NORMAL);
125 } else if (this->inputDtype == ge::DT_INT64 || this->inputDtype == ge::DT_COMPLEX64) {167 } else if (this->inputDtype == ge::DT_INT64 || this->inputDtype == ge::DT_COMPLEX64) {
126 BroadcastBaseTiling<AddWithoutCastCompute<int64_t>::OpDag> brcBaseTiling(tilingContext);168 BroadcastBaseTiling<AddWithoutCastCompute<int64_t>::OpDag> brcBaseTiling(tilingContext);
127 ret = brcBaseTiling.DoTiling();169 ret = brcBaseTiling.DoTiling();
128- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());170+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), ADD_V2_USER_DEF_NORMAL);
129 } else if (this->inputDtype == ge::DT_UINT8) {171 } else if (this->inputDtype == ge::DT_UINT8) {
130 BroadcastBaseTiling<AddWithoutCastCompute<uint8_t>::OpDag> brcBaseTiling(tilingContext);172 BroadcastBaseTiling<AddWithoutCastCompute<uint8_t>::OpDag> brcBaseTiling(tilingContext);
131 ret = brcBaseTiling.DoTiling();173 ret = brcBaseTiling.DoTiling();
132- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());174+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), ADD_V2_USER_DEF_NORMAL);
133 } else if (this->inputDtype == ge::DT_INT8) {175 } else if (this->inputDtype == ge::DT_INT8) {
134 BroadcastBaseTiling<AddWithoutCastCompute<int8_t>::OpDag> brcBaseTiling(tilingContext);176 BroadcastBaseTiling<AddWithoutCastCompute<int8_t>::OpDag> brcBaseTiling(tilingContext);
135 ret = brcBaseTiling.DoTiling();177 ret = brcBaseTiling.DoTiling();
136- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());178+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), ADD_V2_USER_DEF_NORMAL);
137 } else if (this->inputDtype == ge::DT_INT32) {179 } else if (this->inputDtype == ge::DT_INT32) {
138 BroadcastBaseTiling<AddWithoutCastCompute<int32_t>::OpDag> brcBaseTiling(tilingContext);180 BroadcastBaseTiling<AddWithoutCastCompute<int32_t>::OpDag> brcBaseTiling(tilingContext);
139 ret = brcBaseTiling.DoTiling();181 ret = brcBaseTiling.DoTiling();
140- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());182+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), ADD_V2_USER_DEF_NORMAL);
141 } else if (this->inputDtype == ge::DT_INT16) {183 } else if (this->inputDtype == ge::DT_INT16) {
142 BroadcastBaseTiling<AddWithoutCastCompute<int16_t>::OpDag> brcBaseTiling(tilingContext);184 BroadcastBaseTiling<AddWithoutCastCompute<int16_t>::OpDag> brcBaseTiling(tilingContext);
143 ret = brcBaseTiling.DoTiling();185 ret = brcBaseTiling.DoTiling();
144- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());186+ tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), ADD_V2_USER_DEF_NORMAL);
145 } else {187 } else {
146 OP_LOGE_FOR_INVALID_DTYPE(tilingContext->GetNodeName(), "x1",188 OP_LOGE_FOR_INVALID_DTYPE(tilingContext->GetNodeName(), "x1",
147 ge::TypeUtils::DataTypeToSerialString(this->inputDtype),189 ge::TypeUtils::DataTypeToSerialString(this->inputDtype),
@@ -153,9 +195,8 @@ ge::graphStatus AddV2TilingArch35::RunTiling()
153 "broadcastBaseTiling failed"),195 "broadcastBaseTiling failed"),
154 return ge::GRAPH_FAILED);196 return ge::GRAPH_FAILED);
155 197 
156- size_t* currentWorkspace = tilingContext->GetWorkspaceSizes(1);198+ OP_CHECK_IF(SetWorkspace() != ge::GRAPH_SUCCESS, OP_LOGE(tilingContext, "AddV2: set workspace failed."),
157- OP_CHECK_NULL_WITH_CONTEXT(tilingContext, currentWorkspace);199+ return ge::GRAPH_FAILED);
158- currentWorkspace[0] = static_cast<uint64_t>(ASCEND_WORKSPACE);
159 200 
160 OP_LOGD(tilingContext, "[TilingData] : tilingKey=%lu", tilingKey);201 OP_LOGD(tilingContext, "[TilingData] : tilingKey=%lu", tilingKey);
161 tilingContext->SetTilingKey(tilingKey);202 tilingContext->SetTilingKey(tilingKey);
Mmath/add_v2/op_host/config/ascend950/add_v2_binary.json+0-160
@@ -41,46 +41,6 @@
41 }41 }
42 ]42 ]
43 },43 },
44- {
45- "bin_filename": "AddV2_BFLOAT16_FLOAT32",
46- "inputs": [
47- {
48- "name": "x1",
49- "index": 0,
50- "dtype": "bfloat16",
51- "format": "ND",
52- "paramType": "required",
53- "shape": [
54- -2
55- ],
56- "format_match_mode": "FormatAgnostic"
57- },
58- {
59- "name": "x2",
60- "index": 1,
61- "dtype": "float32",
62- "format": "ND",
63- "paramType": "required",
64- "shape": [
65- -2
66- ],
67- "format_match_mode": "FormatAgnostic"
68- }
69- ],
70- "outputs": [
71- {
72- "name": "y",
73- "index": 0,
74- "dtype": "float32",
75- "format": "ND",
76- "paramType": "required",
77- "shape": [
78- -2
79- ],
80- "format_match_mode": "FormatAgnostic"
81- }
82- ]
83- },
84 {44 {
85 "bin_filename": "AddV2_COMPLEX64",45 "bin_filename": "AddV2_COMPLEX64",
86 "inputs": [46 "inputs": [
@@ -161,126 +121,6 @@
161 }121 }
162 ]122 ]
163 },123 },
164- {
165- "bin_filename": "AddV2_FLOAT16_FLOAT32",
166- "inputs": [
167- {
168- "name": "x1",
169- "index": 0,
170- "dtype": "float16",
171- "format": "ND",
172- "paramType": "required",
173- "shape": [
174- -2
175- ],
176- "format_match_mode": "FormatAgnostic"
177- },
178- {
179- "name": "x2",
180- "index": 1,
181- "dtype": "float32",
182- "format": "ND",
183- "paramType": "required",
184- "shape": [
185- -2
186- ],
187- "format_match_mode": "FormatAgnostic"
188- }
189- ],
190- "outputs": [
191- {
192- "name": "y",
193- "index": 0,
194- "dtype": "float32",
195- "format": "ND",
196- "paramType": "required",
197- "shape": [
198- -2
199- ],
200- "format_match_mode": "FormatAgnostic"
201- }
202- ]
203- },
204- {
205- "bin_filename": "AddV2_FLOAT32_BFLOAT16",
206- "inputs": [
207- {
208- "name": "x1",
209- "index": 0,
210- "dtype": "float32",
211- "format": "ND",
212- "paramType": "required",
213- "shape": [
214- -2
215- ],
216- "format_match_mode": "FormatAgnostic"
217- },
218- {
219- "name": "x2",
220- "index": 1,
221- "dtype": "bfloat16",
222- "format": "ND",
223- "paramType": "required",
224- "shape": [
225- -2
226- ],
227- "format_match_mode": "FormatAgnostic"
228- }
229- ],
230- "outputs": [
231- {
232- "name": "y",
233- "index": 0,
234- "dtype": "float32",
235- "format": "ND",
236- "paramType": "required",
237- "shape": [
238- -2
239- ],
240- "format_match_mode": "FormatAgnostic"
241- }
242- ]
243- },
244- {
245- "bin_filename": "AddV2_FLOAT32_FLOAT16",
246- "inputs": [
247- {
248- "name": "x1",
249- "index": 0,
250- "dtype": "float32",
251- "format": "ND",
252- "paramType": "required",
253- "shape": [
254- -2
255- ],
256- "format_match_mode": "FormatAgnostic"
257- },
258- {
259- "name": "x2",
260- "index": 1,
261- "dtype": "float16",
262- "format": "ND",
263- "paramType": "required",
264- "shape": [
265- -2
266- ],
267- "format_match_mode": "FormatAgnostic"
268- }
269- ],
270- "outputs": [
271- {
272- "name": "y",
273- "index": 0,
274- "dtype": "float32",
275- "format": "ND",
276- "paramType": "required",
277- "shape": [
278- -2
279- ],
280- "format_match_mode": "FormatAgnostic"
281- }
282- ]
283- },
284 {124 {
285 "bin_filename": "AddV2_FLOAT32",125 "bin_filename": "AddV2_FLOAT32",
286 "inputs": [126 "inputs": [
Mmath/add_v2/op_kernel/arch35/add_v2.cpp+21-20
@@ -22,28 +22,29 @@ using namespace Ops::Base;
22using namespace AscendC;22using namespace AscendC;
23using namespace AddV2Op;23using namespace AddV2Op;
24 24 
25-template <uint64_t schMode>25+template <uint64_t schMode, uint64_t userDef>
26__global__ __aicore__ void add_v2(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)26__global__ __aicore__ void add_v2(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
27{27{
28- constexpr bool isMixDtype = (std::is_same<DTYPE_X1, half>::value && std::is_same<DTYPE_X2, float>::value) ||28+ if constexpr (userDef == 1) {
29- (std::is_same<DTYPE_X1, float>::value && std::is_same<DTYPE_X2, half>::value) ||29+ // 空 Tensor：y 的元素个数为 0，没有任何数据需要搬运或写回，直接返回。
30- (std::is_same<DTYPE_X1, bfloat16_t>::value && std::is_same<DTYPE_X2, float>::value) ||30+ // Tiling 侧已把 blockDim 设为 1，这里只是把这一个核空转掉。
31- (std::is_same<DTYPE_X1, float>::value && std::is_same<DTYPE_X2, bfloat16_t>::value);31+ // 这一次 GET_TILING_DATA_WITH_STRUCT 不能省：opc 靠它反推本模板实例的 tiling
32- constexpr bool isNeedCast = (std::is_same<DTYPE_X1, half>::value || std::is_same<DTYPE_X1, bfloat16_t>::value ||32+ // 结构体大小，分支里不引用任何结构体会导致 tiling_struct_size 未定义而编译失败。
33- std::is_same<DTYPE_X1, float>::value) &&33+ GET_TILING_DATA_WITH_STRUCT(AddV2EmptyTilingData, emptyTilingData, tiling);
34- (!isMixDtype);34+ return;
35- 
36- if constexpr (isMixDtype) {
37- using OpDag = AddMixDtypeCompute<DTYPE_X1, DTYPE_X2>::OpDag;
38- BroadcastSch<schMode, OpDag> sch(tiling);
39- sch.Process(x1, x2, y);
40- } else if constexpr (isNeedCast) {
41- using OpDag = AddWithCastCompute<DTYPE_X1>::OpDag;
42- BroadcastSch<schMode, OpDag> sch(tiling);
43- sch.Process(x1, x2, y);
44 } else {35 } else {
45- using OpDag = AddWithoutCastCompute<DTYPE_X1>::OpDag;36+ // 仅注册同 dtype 组合，DTYPE_X1 与 DTYPE_X2 恒等
46- BroadcastSch<schMode, OpDag> sch(tiling);37+ constexpr bool isNeedCast = std::is_same<DTYPE_X1, half>::value || std::is_same<DTYPE_X1, bfloat16_t>::value ||
47- sch.Process(x1, x2, y);38+ std::is_same<DTYPE_X1, float>::value;
39+ 
40+ if constexpr (isNeedCast) {
41+ using OpDag = AddWithCastCompute<DTYPE_X1>::OpDag;
42+ BroadcastSch<schMode, OpDag> sch(tiling);
43+ sch.Process(x1, x2, y);
44+ } else {
45+ using OpDag = AddWithoutCastCompute<DTYPE_X1>::OpDag;
46+ BroadcastSch<schMode, OpDag> sch(tiling);
47+ sch.Process(x1, x2, y);
48+ }
48 }49 }
49}50}
Mmath/add_v2/op_kernel/arch35/add_v2_dag.h+0-13
@@ -50,19 +50,6 @@ struct AddWithCastCompute {
50 using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;50 using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
51 using OpDag = DAGSch<Outputs, void, MemCfg>;51 using OpDag = DAGSch<Outputs, void, MemCfg>;
52};52};
53- 
54-template <typename T1, typename T2>
55-struct AddMixDtypeCompute {
56- using OpInputX1 = Bind<Vec::CopyInBrc<T1>, Placeholder::In0<T1>>;
57- using OpInputX2 = Bind<Vec::CopyInBrc<T2>, Placeholder::In1<T2>>;
58- using OpCastX1 = Bind<Vec::Cast<float, T1, CAST_NONE_MODE>, OpInputX1>;
59- using OpCastX2 = Bind<Vec::Cast<float, T2, CAST_NONE_MODE>, OpInputX2>;
60- using OpAddRes = Bind<Vec::Add<float>, OpCastX1, OpCastX2>;
61- using OpCopyOut = Bind<Vec::CopyOut<float>, Placeholder::Out0<float>, OpAddRes>;
62- using Outputs = Elems<OpCopyOut>;
63- using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
64- using OpDag = DAGSch<Outputs, void, MemCfg>;
65-};
66} // namespace AddV2Op53} // namespace AddV2Op
67 54 
68#endif // ADD_V2_DAG_H55#endif // ADD_V2_DAG_H
Mmath/add_v2/op_kernel/arch35/add_v2_struct_arch35.h+19-2
@@ -19,8 +19,25 @@
19#include "atvoss/broadcast/broadcast_base_struct.h"19#include "atvoss/broadcast/broadcast_base_struct.h"
20 20 
21namespace AddV2Op {21namespace AddV2Op {
22-ASCENDC_TPL_ARGS_DECL(AddV2, BRC_TEMP_SCH_MODE_KEY_DECL(schMode));22+// 空 Tensor 分支的 tiling data。内容其实用不上（无元素可搬），但**必须存在**：
23-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(BRC_TEMP_SCH_MODE_KEY_SEL(schMode)));23+// opc 依赖 kernel 里的 GET_TILING_DATA_WITH_STRUCT 反推每个模板实例的 tiling 结构体
24+// 大小，自定义分支若完全不引用任何结构体，会报
25+// "UnboundLocalError: cannot access local variable 'tiling_struct_size'"。
26+struct AddV2EmptyTilingData {
27+ int64_t numel; // 恒为 0，仅作占位与调试
28+};
29+ 
30+// userDef = 0：常规 broadcast 通路，schMode 由 BroadcastBaseTiling 决定；
31+// userDef = 1：空 Tensor 通路，走自定义 schMode 999。ATVOSS 的 BroadcastBaseTiling
32+// 在合轴后显式拒绝 0 元素（broadcast_tiling.h 的 "tensor check is empty"），
33+// 所以空 Tensor 不能复用常规 schMode，必须另开一条自定义模板分支。
34+// 这里的写法与 math/select 的 SIMT 自定义分支一致。
35+ASCENDC_TPL_ARGS_DECL(AddV2, BRC_TEMP_SCH_MODE_KEY_DECL(schMode),
36+ ASCENDC_TPL_UINT_DECL(userDef, 8, ASCENDC_TPL_UI_LIST, 0, 1));
37+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(BRC_TEMP_SCH_MODE_KEY_SEL(schMode),
38+ ASCENDC_TPL_UINT_SEL(userDef, ASCENDC_TPL_UI_LIST, 0)),
39+ ASCENDC_TPL_ARGS_SEL(BRC_TEMP_SCH_CUSTOM_MODE_KEY_SEL(schMode),
40+ ASCENDC_TPL_UINT_SEL(userDef, ASCENDC_TPL_UI_LIST, 1)));
24} // namespace AddV2Op41} // namespace AddV2Op
25 42 
26#endif // ADD_V2_STRUCT_ARCH35_H43#endif // ADD_V2_STRUCT_ARCH35_H
Mmath/add_v2/tests/assets/golden.py+7-4
@@ -25,16 +25,19 @@ def add_v2_golden(x1, x2, **kwargs):
25 input_formats, output_formats, input_ori_formats, output_ori_formats,25 input_formats, output_formats, input_ori_formats, output_ori_formats,
26 input_dtypes, output_dtypes.26 input_dtypes, output_dtypes.
27 """27 """
28+ # 仅注册同 dtype 组合，x1/x2 dtype 恒等，输出 dtype 与 x1 一致
28 dtype = x1.dtype29 dtype = x1.dtype
30+ if str(x2.dtype) != str(dtype):
31+ raise ValueError(
32+ f"add_v2 only supports identical input dtypes, got x1={dtype}, x2={x2.dtype}"
33+ )
34+ # torch 无原生 bfloat16 numpy 视图，先升 float32 计算再还原
29 if "bfloat16" in str(dtype):35 if "bfloat16" in str(dtype):
30 x1 = x1.astype("float32")36 x1 = x1.astype("float32")
31 x2 = x2.astype("float32")37 x2 = x2.astype("float32")
32 x = torch.from_numpy(x1)38 x = torch.from_numpy(x1)
33 y = torch.from_numpy(x2)39 y = torch.from_numpy(x2)
34- if "bool" in str(dtype):40+ res = torch.add(x, y).numpy()
35- res = torch.logical_or(x, y).numpy()
36- else:
37- res = torch.add(x, y).numpy()
38 if "bfloat16" in str(dtype):41 if "bfloat16" in str(dtype):
39 res = res.astype(dtype)42 res = res.astype(dtype)
40 43 
Mmath/add_v2/tests/ut/op_host/arch35/test_add_v2_tiling.cpp+80-0
@@ -22,6 +22,13 @@
22using namespace std;22using namespace std;
23using namespace ge;23using namespace ge;
24 24 
25+// 空 Tensor 分支的 tiling key = 65550 = 0x1000E：
26+// 低 16 位 0x000E = 14，是 schMode 999 在 BRC_TEMP_SCH_MODE_KEY_DECL 取值表
27+// (1,2,101,102,103,104,109,201,202,301,302,303,304,305,999) 里的序号；
28+// 第 16 位 = userDef = 1。
29+// 常规通路是 userDef = 0，所以高位为 0，原有用例的 key 不受影响（仍为 8）。
30+static constexpr uint64_t ADD_V2_UT_EMPTY_TILING_KEY = 65550;
31+ 
25class AddV2Tiling : public testing::Test {32class AddV2Tiling : public testing::Test {
26protected:33protected:
27 static void SetUpTestCase() { std::cout << "AddV2Tiling SetUp" << std::endl; }34 static void SetUpTestCase() { std::cout << "AddV2Tiling SetUp" << std::endl; }
@@ -128,3 +135,76 @@ TEST_F(AddV2Tiling, add_v2_tiling_invalid_dtype)
128 &compileInfo);135 &compileInfo);
129 ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);136 ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
130}137}
138+ 
139+// ── 空 Tensor ──────────────────────────────────────────────────────────────
140+// ATVOSS 的 BroadcastBaseTiling 在合轴后显式拒绝 0 元素，空 Tensor 必须走
141+// 自定义模板分支（schMode 999 + userDef 1），blockDim = 1，kernel 侧直接返回。
142+TEST_F(AddV2Tiling, add_v2_tiling_empty_1d)
143+{
144+ optiling::AddV2CompileInfoArch35 compileInfo = {64, 245760};
145+ gert::TilingContextPara tilingContextPara("AddV2",
146+ {
147+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
148+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
149+ },
150+ {
151+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
152+ },
153+ &compileInfo);
154+ uint64_t expectTilingKey = ADD_V2_UT_EMPTY_TILING_KEY;
155+ std::vector<size_t> expectWorkspaces = {16777216};
156+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
157+}
158+ 
159+TEST_F(AddV2Tiling, add_v2_tiling_empty_2d)
160+{
161+ optiling::AddV2CompileInfoArch35 compileInfo = {64, 245760};
162+ gert::TilingContextPara tilingContextPara("AddV2",
163+ {
164+ {{{0, 3}, {0, 3}}, ge::DT_FLOAT16, ge::FORMAT_ND},
165+ {{{0, 3}, {0, 3}}, ge::DT_FLOAT16, ge::FORMAT_ND},
166+ },
167+ {
168+ {{{0, 3}, {0, 3}}, ge::DT_FLOAT16, ge::FORMAT_ND},
169+ },
170+ &compileInfo);
171+ uint64_t expectTilingKey = ADD_V2_UT_EMPTY_TILING_KEY;
172+ std::vector<size_t> expectWorkspaces = {16777216};
173+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
174+}
175+ 
176+// 空张量 + 广播：x1 空、x2 是标量，输出仍为空
177+TEST_F(AddV2Tiling, add_v2_tiling_empty_broadcast)
178+{
179+ optiling::AddV2CompileInfoArch35 compileInfo = {64, 245760};
180+ gert::TilingContextPara tilingContextPara("AddV2",
181+ {
182+ {{{0, 3}, {0, 3}}, ge::DT_INT32, ge::FORMAT_ND},
183+ {{{1, 3}, {1, 3}}, ge::DT_INT32, ge::FORMAT_ND},
184+ },
185+ {
186+ {{{0, 3}, {0, 3}}, ge::DT_INT32, ge::FORMAT_ND},
187+ },
188+ &compileInfo);
189+ uint64_t expectTilingKey = ADD_V2_UT_EMPTY_TILING_KEY;
190+ std::vector<size_t> expectWorkspaces = {16777216};
191+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
192+}
193+ 
194+// 高维中间维为 0
195+TEST_F(AddV2Tiling, add_v2_tiling_empty_highrank)
196+{
197+ optiling::AddV2CompileInfoArch35 compileInfo = {64, 245760};
198+ gert::TilingContextPara tilingContextPara("AddV2",
199+ {
200+ {{{2, 0, 4}, {2, 0, 4}}, ge::DT_INT8, ge::FORMAT_ND},
201+ {{{2, 0, 4}, {2, 0, 4}}, ge::DT_INT8, ge::FORMAT_ND},
202+ },
203+ {
204+ {{{2, 0, 4}, {2, 0, 4}}, ge::DT_INT8, ge::FORMAT_ND},
205+ },
206+ &compileInfo);
207+ uint64_t expectTilingKey = ADD_V2_UT_EMPTY_TILING_KEY;
208+ std::vector<size_t> expectWorkspaces = {16777216};
209+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
210+}