fix(AccumulateNV2): P1缺陷整改+broadcast收回+fp16性能优化+golden精度标准落位-ops-math-AtomGit

Mmath/accumulate_nv2/README.md+11-8

	@@ -1,4 +1,4 @@
1	-# AccumulateNv2	1	+# AccumulateNV2
2		2
3	## 产品支持情况	3	## 产品支持情况
4		4
	@@ -39,8 +39,8 @@ $$y = \sum_{i=0}^{N-1} x_i$$
39	<tbody>	39	<tbody>
40	<tr>	40	<tr>
41	<td>x</td>	41	<td>x</td>
42	- <td>输入（动态输入）</td>	42	+ <td>输入</td>
43	- <td>需要求和的输入tensor列表，公式中的x_i。各tensor的shape需满足broadcast关系。</td>	43	+ <td>需要求和的输入tensor列表，公式中的x_i。图模式直连时各输入shape需完全相同（算子本体不支持broadcast）；通过aclnnSum接口调用时各tensor的shape可不同，由接口层完成对齐。</td>
44	<td>FLOAT16、FLOAT、INT8、INT32、UINT8</td>	44	<td>FLOAT16、FLOAT、INT8、INT32、UINT8</td>
45	<td>ND</td>	45	<td>ND</td>
46	</tr>	46	</tr>
	@@ -53,8 +53,8 @@ $$y = \sum_{i=0}^{N-1} x_i$$
53	</tr>	53	</tr>
54	<tr>	54	<tr>
55	<td>N</td>	55	<td>N</td>
56	- <td>属性（必填）</td>	56	+ <td>属性</td>
57	- <td>输入tensor列表的大小。</td>	57	+ <td>输入tensor列表的大小。N需大于等于1，且必须与实际输入Tensor的数量一致；不一致时在图编译阶段报错。</td>
58	<td>INT</td>	58	<td>INT</td>
59	<td>-</td>	59	<td>-</td>
60	</tr>	60	</tr>
	@@ -62,11 +62,14 @@ $$y = \sum_{i=0}^{N-1} x_i$$
62		62
63	## 约束说明	63	## 约束说明
64		64
65	-- 输入tensors列表中各tensor的shape需满足broadcast关系，broadcast后的shape需与输出y的shape一致。	65	+- 通过aclnnSum接口调用时，输入tensors列表中各tensor的shape需满足broadcast关系，由实现层完成对齐；通过图模式（GE/GEIR）直连算子时，各输入shape需完全相同且与输出y一致（算子本体不支持broadcast），shape不一致时在图编译阶段报错。
66	-- 输入tensors与输出的数据类型需相同。	66	+- 输入tensors与输出的数据类型需相同，且tensors列表内各tensor的数据类型需一致；不一致时在图编译阶段报错。
		67	+- 输入维度不超过8维：aclnnSum接口第一段校验拦截（ACLNN_ERR_PARAM_INVALID 161002）；图模式直连时在tiling阶段拦截。
		68	+- 图模式（GE/GEIR）直连时输入数量需在[1, 40]范围内（tiling阶段拦截，超限报错）；aclnnSum接口内部按16个一组分组累加，不受该上限约束。
67		69
68	## 调用说明	70	## 调用说明
69		71
70	\| 调用方式 \| 调用样例 \| 说明 \|	72	\| 调用方式 \| 调用样例 \| 说明 \|
71	\|--------------\|------------------------------------------------------------------------\|--------------------------------------------------------------\|	73	\|--------------\|------------------------------------------------------------------------\|--------------------------------------------------------------\|
72	-\| aclnn调用 \| [test_aclnn_sum](./examples/test_aclnn_sum.cpp) \| 通过[aclnnSum](./docs/aclnnSum.md)接口方式调用AccumulateNv2算子。 \|	74	+\| aclnn调用 \| [test_aclnn_sum](./examples/test_aclnn_sum.cpp) \| 通过[aclnnSum](./docs/aclnnSum.md)接口方式调用AccumulateNV2算子。 \|
		75	+\| 图模式（GEIR） \| [accumulate_nv2_proto.h](./op_graph/accumulate_nv2_proto.h) \| 通过[AccumulateNV2算子IR定义](./op_graph/accumulate_nv2_proto.h)在GE图中调用（各输入shape需完全相同）。 \|

Mmath/accumulate_nv2/docs/aclnnSum.md+19-6

	@@ -83,7 +83,7 @@ aclnnStatus aclnnSum(
83	<td>tensors（aclTensorList*）</td>	83	<td>tensors（aclTensorList*）</td>
84	<td>输入</td>	84	<td>输入</td>
85	<td>需要求和的输入tensor列表，公式中的tensors<sub>i</sub>表示列表中的第i个tensor。</td>	85	<td>需要求和的输入tensor列表，公式中的tensors<sub>i</sub>表示列表中的第i个tensor。</td>
86	- <td>需要与out数据类型相同。tensors中各tensor的shape需要与out满足<a href="../../../docs/zh/context/broadcast_relationship.md" target="_blank">broadcast关系</a>。</td>	86	+ <td>需要与out数据类型相同。tensors中各tensor的shape需要与out满足<a href="../../../docs/zh/context/broadcast_relationship.md" target="_blank">broadcast关系</a>。支持空Tensor（shape含0轴）。</td>
87	<td>FLOAT16、FLOAT、INT8、INT32、UINT8</td>	87	<td>FLOAT16、FLOAT、INT8、INT32、UINT8</td>
88	<td>ND</td>	88	<td>ND</td>
89	<td>不大于8</td>	89	<td>不大于8</td>
	@@ -93,7 +93,7 @@ aclnnStatus aclnnSum(
93	<td>out（aclTensor*）</td>	93	<td>out（aclTensor*）</td>
94	<td>输出</td>	94	<td>输出</td>
95	<td>求和结果tensor，公式中的out。</td>	95	<td>求和结果tensor，公式中的out。</td>
96	- <td>需要与tensors数据类型相同。shape需要与tensors中各tensor的shape满足<a href="../../../docs/zh/context/broadcast_relationship.md" target="_blank">broadcast关系</a>。</td>	96	+ <td>需要与tensors数据类型相同。shape需要与tensors中各tensor的shape满足<a href="../../../docs/zh/context/broadcast_relationship.md" target="_blank">broadcast关系</a>。支持空Tensor（shape含0轴，此时输出为空Tensor）。</td>
97	<td>FLOAT、FLOAT16、INT8、INT32、UINT8</td>	97	<td>FLOAT、FLOAT16、INT8、INT32、UINT8</td>
98	<td>ND</td>	98	<td>ND</td>
99	<td>-</td>	99	<td>-</td>
	@@ -140,21 +140,30 @@ aclnnStatus aclnnSum(
140	</tr></thead>	140	</tr></thead>
141	<tbody>	141	<tbody>
142	<tr>	142	<tr>
143	- <td>ACLNN_ERR_PARAM_NULLPTR</td>	143	+ <td rowspan="2">ACLNN_ERR_PARAM_NULLPTR</td>
144	- <td>161001</td>	144	+ <td rowspan="2">161001</td>
145	<td>传入的tensors或out是空指针时。</td>	145	<td>传入的tensors或out是空指针时。</td>
146	</tr>	146	</tr>
147	<tr>	147	<tr>
148	- <td rowspan="3">ACLNN_ERR_PARAM_INVALID</td>	148	+ <td>tensors列表中存在空指针元素时。</td>
149	- <td rowspan="3">161002</td>	149	+ </tr>
		150	+ <tr>
		151	+ <td rowspan="5">ACLNN_ERR_PARAM_INVALID</td>
		152	+ <td rowspan="5">161002</td>
150	<td>tensors列表中tensor或out的数据类型不在支持的范围之内。</td>	153	<td>tensors列表中tensor或out的数据类型不在支持的范围之内。</td>
151	</tr>	154	</tr>
152	<tr>	155	<tr>
153	<td>tensors列表和out的数据类型不一致。</td>	156	<td>tensors列表和out的数据类型不一致。</td>
154	</tr>	157	</tr>
		158	+ <tr>
		159	+ <td>传入空的TensorList（tensors->Size()==0）。</td>
		160	+ </tr>
155	<tr>	161	<tr>
156	<td>tensors和out的shape不满足broadcast规则，或者broadcast后的shape与out不一致。</td>	162	<td>tensors和out的shape不满足broadcast规则，或者broadcast后的shape与out不一致。</td>
157	</tr>	163	</tr>
		164	+ <tr>
		165	+ <td>tensors列表或out的维度大于8。</td>
		166	+ </tr>
158	</tbody>	167	</tbody>
159	</table>	168	</table>
160		169
	@@ -203,6 +212,10 @@ aclnnStatus aclnnSum(
203		212
204	## 约束说明	213	## 约束说明
205		214
		215	+- NaN 输入传播：任一输入含 NaN 时，对应输出位置为 NaN。
		216	+- ±Inf 求和：同号 Inf 相加为该符号 Inf；异号 Inf 相加为 NaN（IEEE 754 语义）。
		217	+- 动态 shape 标记（图模式直连路径）：-1（未知维度）在输入间逐维比对时按字面值处理（各输入的 -1 位置须一致）；-2（未知 rank）作为首输入时输出透传未知 rank，后续输入跳过比对。
		218	+
206	- 确定性计算：	219	- 确定性计算：
207	- aclnnSum默认确定性实现。	220	- aclnnSum默认确定性实现。
208		221

Mmath/accumulate_nv2/framework/CMakeLists.txt+11-12

	@@ -1,12 +1,11 @@
1	-# ----------------------------------------------------------------------------	1	+# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify.	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	-# This file is a part of the CANN Open Software.	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
6	-# 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	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of	8	+# See LICENSE in the root of the software repository for the full text of the License.
9	-# the software repository for the full text of the License.	9	+# ----------------------------------------------------------------------------
10	-# ----------------------------------------------------------------------------	10	+
11	-	11	+add_onnx_plugin_sources()
12	-add_onnx_plugin_sources()

Mmath/accumulate_nv2/framework/sum_onnx_plugin.cpp+10-15

	@@ -30,18 +30,13 @@ static Status ParseParamsSum(const Message* op_src, ge::Operator& op_dest)
30	}	30	}
31		31
32	REGISTER_CUSTOM_OP("AccumulateNV2")	32	REGISTER_CUSTOM_OP("AccumulateNV2")
33	- .FrameworkType(ONNX)	33	+ .FrameworkType(ONNX)
34	- .OriginOpType({ge::AscendString("ai.onnx::8::Sum"),	34	+ .OriginOpType({ge::AscendString("ai.onnx::8::Sum"), ge::AscendString("ai.onnx::9::Sum"),
35	- ge::AscendString("ai.onnx::9::Sum"),	35	+ ge::AscendString("ai.onnx::10::Sum"), ge::AscendString("ai.onnx::11::Sum"),
36	- ge::AscendString("ai.onnx::10::Sum"),	36	+ ge::AscendString("ai.onnx::12::Sum"), ge::AscendString("ai.onnx::13::Sum"),
37	- ge::AscendString("ai.onnx::11::Sum"),	37	+ ge::AscendString("ai.onnx::14::Sum"), ge::AscendString("ai.onnx::15::Sum"),
38	- ge::AscendString("ai.onnx::12::Sum"),	38	+ ge::AscendString("ai.onnx::16::Sum"), ge::AscendString("ai.onnx::17::Sum"),
39	- ge::AscendString("ai.onnx::13::Sum"),	39	+ ge::AscendString("ai.onnx::18::Sum")})
40	- ge::AscendString("ai.onnx::14::Sum"),	40	+ .ParseParamsFn(ParseParamsSum)
41	- ge::AscendString("ai.onnx::15::Sum"),	41	+ .ImplyType(ImplyType::TVM);
42	- ge::AscendString("ai.onnx::16::Sum"),	42	+} // namespace domi
43	- ge::AscendString("ai.onnx::17::Sum"),
44	- ge::AscendString("ai.onnx::18::Sum")})
45	- .ParseParamsFn(ParseParamsSum)
46	- .ImplyType(ImplyType::TVM);
47	-} // namespace domi

Mmath/accumulate_nv2/op_api/accumulate_nv2.cpp+10-9

	@@ -19,18 +19,19 @@
19	#include "opdev/op_executor.h"	19	#include "opdev/op_executor.h"
20	#include "opdev/shape_utils.h"	20	#include "opdev/shape_utils.h"
21		21
22	-
23	using namespace op;	22	using namespace op;
24		23
25	namespace l0op {	24	namespace l0op {
26	OP_TYPE_REGISTER(AccumulateNV2);	25	OP_TYPE_REGISTER(AccumulateNV2);
27		26
28	-const aclTensor AccumulateNV2(const aclTensorList tensors, aclOpExecutor *executor) {	27	+const aclTensor* AccumulateNV2(const aclTensorList* tensors, aclOpExecutor* executor)
29	- L0_DFX(AccumulateNV2, tensors);	28	+{
30	- auto out = executor->AllocTensor((tensors)[0]->GetViewShape(), (tensors)[0]->GetDataType());	29	+ L0_DFX(AccumulateNV2, tensors);
31	- auto ret = ADD_TO_LAUNCHER_LIST_AICORE(AccumulateNV2, OP_INPUT(tensors), OP_OUTPUT(out));	30	+ auto out = executor->AllocTensor((tensors)[0]->GetViewShape(), (tensors)[0]->GetDataType());
32	- OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "AccumulateNV2AiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),	31	+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(AccumulateNV2, OP_INPUT(tensors), OP_OUTPUT(out));
33	- return nullptr);	32	+ OP_CHECK(ret == ACLNN_SUCCESS,
34	- return out;	33	+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "AccumulateNV2AiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),
		34	+ return nullptr);
		35	+ return out;
35	}	36	}
36	-} // namespace l0op	37	+} // namespace l0op

Mmath/accumulate_nv2/op_api/accumulate_nv2.h+2-2

	@@ -18,7 +18,7 @@
18	#include "opdev/op_executor.h"	18	#include "opdev/op_executor.h"
19		19
20	namespace l0op {	20	namespace l0op {
21	-const aclTensor AccumulateNV2(const aclTensorList tensors, aclOpExecutor *executor);	21	+const aclTensor* AccumulateNV2(const aclTensorList* tensors, aclOpExecutor* executor);
22	}	22	}
23		23
24	-#endif // OP_API_INC_LEVEL0_OP_ACCUMULATE_NV2_H_	24	+#endif // OP_API_INC_LEVEL0_OP_ACCUMULATE_NV2_H_

Mmath/accumulate_nv2/op_api/aclnn_sum.cpp+8-0

	@@ -58,6 +58,14 @@ static bool CheckDtypeValid(const aclTensorList* tensors, const aclTensor* out)
58	}	58	}
59	}	59	}
60	OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST, return false);	60	OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST, return false);
		61	+ for (uint64_t i = 0; i < tensors->Size(); i++) {
		62	+ if ((*tensors)[i]->GetDataType() != out->GetDataType()) {
		63	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "tensors[%lu] dtype %s is not the same as out dtype %s.", i,
		64	+ op::ToString((*tensors)[i]->GetDataType()).GetString(),
		65	+ op::ToString(out->GetDataType()).GetString());
		66	+ return false;
		67	+ }
		68	+ }
61	return true;	69	return true;
62	}	70	}
63		71

Mmath/accumulate_nv2/op_graph/accumulate_nv2_graph_infer.cpp+49-57

	@@ -16,6 +16,7 @@
16	#include "graph/operator_reg.h"	16	#include "graph/operator_reg.h"
17	#include "register/op_impl_registry.h"	17	#include "register/op_impl_registry.h"
18	#include "exe_graph/runtime/infer_datatype_context.h"	18	#include "exe_graph/runtime/infer_datatype_context.h"
		19	+#include "log/log.h"
19	#include "accumulate_nv2_graph_infer.h"	20	#include "accumulate_nv2_graph_infer.h"
20		21
21	namespace ge {	22	namespace ge {
	@@ -23,59 +24,8 @@ namespace {
23		24
24	bool IsUnknownRank(const std::vector<int64_t>& dims) { return dims == UNKNOWN_RANK; }	25	bool IsUnknownRank(const std::vector<int64_t>& dims) { return dims == UNKNOWN_RANK; }
25		26
26	-graphStatus MergeBroadcastDim(int64_t lhs, int64_t rhs, int64_t& output)
27	-{
28	- if (lhs == rhs) {
29	- // This branch intentionally handles 0 and avoids the legacy 0 % 0
30	- // SIGFPE in the built-in AccumulateNV2 V1 infer implementation.
31	- output = lhs;
32	- return GRAPH_SUCCESS;
33	- }
34	- if (lhs == 1) {
35	- output = rhs;
36	- return GRAPH_SUCCESS;
37	- }
38	- if (rhs == 1) {
39	- output = lhs;
40	- return GRAPH_SUCCESS;
41	- }
42	- if (lhs == UNKNOWN_DIM \|\| rhs == UNKNOWN_DIM) {
43	- // A known non-one dimension constrains the broadcast result. If both
44	- // dimensions are unknown, keep the result unknown.
45	- output = (lhs == UNKNOWN_DIM && rhs == UNKNOWN_DIM) ? UNKNOWN_DIM : (lhs == UNKNOWN_DIM ? rhs : lhs);
46	- return GRAPH_SUCCESS;
47	- }
48	- return GRAPH_FAILED;
49	-}
50	-
51	} // namespace	27	} // namespace
52		28
53	-namespace accumulate_nv2 {
54	-
55	-graphStatus MergeBroadcastShape(const std::vector<int64_t>& lhs, const std::vector<int64_t>& rhs,
56	- std::vector<int64_t>& output)
57	-{
58	- if (IsUnknownRank(lhs) \|\| IsUnknownRank(rhs)) {
59	- output = UNKNOWN_RANK;
60	- return GRAPH_SUCCESS;
61	- }
62	-
63	- const size_t outputRank = std::max(lhs.size(), rhs.size());
64	- output.assign(outputRank, 1);
65	- const size_t lhsOffset = outputRank - lhs.size();
66	- const size_t rhsOffset = outputRank - rhs.size();
67	- for (size_t i = 0; i < outputRank; ++i) {
68	- const int64_t lhsDim = (i < lhsOffset) ? 1 : lhs[i - lhsOffset];
69	- const int64_t rhsDim = (i < rhsOffset) ? 1 : rhs[i - rhsOffset];
70	- if (MergeBroadcastDim(lhsDim, rhsDim, output[i]) != GRAPH_SUCCESS) {
71	- return GRAPH_FAILED;
72	- }
73	- }
74	- return GRAPH_SUCCESS;
75	-}
76	-
77	-} // namespace accumulate_nv2
78	-
79	IMPLEMT_COMMON_INFERFUNC(AccumulateNV2LegacyV1InferShape)	29	IMPLEMT_COMMON_INFERFUNC(AccumulateNV2LegacyV1InferShape)
80	{	30	{
81	int64_t attrN = 0;	31	int64_t attrN = 0;
	@@ -94,21 +44,47 @@ IMPLEMT_COMMON_INFERFUNC(AccumulateNV2LegacyV1InferShape)
94	const TensorDesc firstInputDesc = op.GetDynamicInputDesc("x", 0U);	44	const TensorDesc firstInputDesc = op.GetDynamicInputDesc("x", 0U);
95	const DataType dataType = firstInputDesc.GetDataType();	45	const DataType dataType = firstInputDesc.GetDataType();
96	std::vector<int64_t> outputDims = firstInputDesc.GetShape().GetDims();	46	std::vector<int64_t> outputDims = firstInputDesc.GetShape().GetDims();
		47	+ // The operator itself does not support broadcast: every input must share the
		48	+ // first input's shape. aclnnSum aligns tensors with BroadcastTo in op_api
		49	+ // before composing the inner operator, so aligned calls still pass.
		50	+ // Dynamic-shape markers keep their inference-time semantics: -2 (unknown
		51	+ // rank) inputs are runtime-determined and skipped; -1 (unknown dim) is a
		52	+ // wildcard that resolves against the first concrete value seen for that
		53	+ // dim (the output keeps the resolved dims), and two different concrete
		54	+ // values for the same dim fail.
		55	+ const bool firstUnknownRank = IsUnknownRank(outputDims);
		56	+ std::vector<int64_t> resolvedDims = outputDims;
97	for (size_t i = 1; i < inputNum; ++i) {	57	for (size_t i = 1; i < inputNum; ++i) {
98	const TensorDesc inputDesc = op.GetDynamicInputDesc("x", static_cast<uint32_t>(i));	58	const TensorDesc inputDesc = op.GetDynamicInputDesc("x", static_cast<uint32_t>(i));
99	if (inputDesc.GetDataType() != dataType) {	59	if (inputDesc.GetDataType() != dataType) {
100	return GRAPH_FAILED;	60	return GRAPH_FAILED;
101	}	61	}
102	- std::vector<int64_t> mergedDims;	62	+ if (firstUnknownRank) {
103	- if (accumulate_nv2::MergeBroadcastShape(outputDims, inputDesc.GetShape().GetDims(), mergedDims) !=	63	+ continue;
104	- GRAPH_SUCCESS) {	64	+ }
		65	+ const std::vector<int64_t> inputDims = inputDesc.GetShape().GetDims();
		66	+ if (IsUnknownRank(inputDims)) {
		67	+ continue;
		68	+ }
		69	+ if (inputDims.size() != outputDims.size()) {
105	return GRAPH_FAILED;	70	return GRAPH_FAILED;
106	}	71	}
107	- outputDims = std::move(mergedDims);	72	+ for (size_t d = 0; d < outputDims.size(); ++d) {
		73	+ if (inputDims[d] == -1) {
		74	+ continue;
		75	+ }
		76	+ if (resolvedDims[d] == -1) {
		77	+ resolvedDims[d] = inputDims[d];
		78	+ continue;
		79	+ }
		80	+ if (inputDims[d] != resolvedDims[d]) {
		81	+ return GRAPH_FAILED;
		82	+ }
		83	+ }
108	}	84	}
109		85
110	TensorDesc outputDesc = op.GetOutputDesc(0U);	86	TensorDesc outputDesc = op.GetOutputDesc(0U);
111	- outputDesc.SetShape(Shape(outputDims));	87	+ outputDesc.SetShape(Shape(resolvedDims));
112	outputDesc.SetDataType(dataType);	88	outputDesc.SetDataType(dataType);
113	return op.UpdateOutputDesc(0U, outputDesc);	89	return op.UpdateOutputDesc(0U, outputDesc);
114	}	90	}
	@@ -124,7 +100,23 @@ namespace ops {
124		100
125	static ge::graphStatus InferDataType4AccumulateNV2(gert::InferDataTypeContext* context)	101	static ge::graphStatus InferDataType4AccumulateNV2(gert::InferDataTypeContext* context)
126	{	102	{
127	- context->SetOutputDataType(0, context->GetInputDataType(0));	103	+ const size_t inputNum = context->GetComputeNodeInputNum();
		104	+ if (inputNum == 0U) {
		105	+ OP_LOGE("AccumulateNV2", "input count must not be zero");
		106	+ return ge::GRAPH_FAILED;
		107	+ }
		108	+ const ge::DataType firstDtype = context->GetInputDataType(0);
		109	+ for (size_t i = 1U; i < inputNum; ++i) {
		110	+ const ge::DataType dtype = context->GetInputDataType(i);
		111	+ if (dtype != firstDtype) {
		112	+ OP_LOGE("AccumulateNV2",
		113	+ "all input dtypes must be the same as the first input, but input[%zu] dtype %d differs from "
		114	+ "input[0] dtype %d",
		115	+ i, static_cast<int32_t>(dtype), static_cast<int32_t>(firstDtype));
		116	+ return ge::GRAPH_FAILED;
		117	+ }
		118	+ }
		119	+ context->SetOutputDataType(0, firstDtype);
128	return ge::GRAPH_SUCCESS;	120	return ge::GRAPH_SUCCESS;
129	}	121	}
130		122

Mmath/accumulate_nv2/op_graph/accumulate_nv2_graph_infer.h+3-7

	@@ -11,18 +11,14 @@
11	#ifndef ACCUMULATE_NV2_GRAPH_INFER_H_	11	#ifndef ACCUMULATE_NV2_GRAPH_INFER_H_
12	#define ACCUMULATE_NV2_GRAPH_INFER_H_	12	#define ACCUMULATE_NV2_GRAPH_INFER_H_
13		13
14	-#include <cstdint>
15	-#include <vector>
16	-
17	#include "graph/ge_error_codes.h"	14	#include "graph/ge_error_codes.h"
18		15
19	namespace ge {	16	namespace ge {
20	-namespace accumulate_nv2 {
21		17
22	-graphStatus MergeBroadcastShape(const std::vector<int64_t>& lhs, const std::vector<int64_t>& rhs,	18	+// V1 legacy InferShape registered via COMMON_INFER_FUNC_REG in the .cpp file.
23	- std::vector<int64_t>& output);	19	+// The operator itself does not support broadcast (all inputs must share the
		20	+// first input's shape); aclnnSum aligns tensors in op_api before composing.
24		21
25	-} // namespace accumulate_nv2
26	} // namespace ge	22	} // namespace ge
27		23
28	#endif // ACCUMULATE_NV2_GRAPH_INFER_H_	24	#endif // ACCUMULATE_NV2_GRAPH_INFER_H_

Mmath/accumulate_nv2/op_graph/accumulate_nv2_proto.h+6-6

	@@ -40,11 +40,11 @@ since minimum temporary storage is proportional to the output size rather than t
40	*	40	*
41	*/	41	*/
42	REG_OP(AccumulateNV2)	42	REG_OP(AccumulateNV2)
43	- .DYNAMIC_INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_INT32, DT_INT8, DT_UINT8}))	43	+ .DYNAMIC_INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_INT32, DT_INT8, DT_UINT8}))
44	- .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT, DT_INT32, DT_INT8, DT_UINT8}))	44	+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT, DT_INT32, DT_INT8, DT_UINT8}))
45	- .REQUIRED_ATTR(N, Int)	45	+ .REQUIRED_ATTR(N, Int)
46	- .OP_END_FACTORY_REG(AccumulateNV2)	46	+ .OP_END_FACTORY_REG(AccumulateNV2)
47		47
48	-} // namespace ge	48	+} // namespace ge
49		49
50	-#endif // ARG_MAX_V2_PROTO_H_	50	+#endif // ARG_MAX_V2_PROTO_H_

Mmath/accumulate_nv2/op_host/accumulate_nv2_def.cpp+2-4

	@@ -29,9 +29,7 @@ public:
29	.Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	29	.Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
30	.UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	30	.UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
31	.AutoContiguous();	31	.AutoContiguous();
32	- this->Attr("N")	32	+ this->Attr("N").AttrType(REQUIRED).Int();
33	- .AttrType(REQUIRED)
34	- .Int();
35		33
36	OpAICoreConfig aiCoreConfig;	34	OpAICoreConfig aiCoreConfig;
37	aiCoreConfig.DynamicCompileStaticFlag(true)	35	aiCoreConfig.DynamicCompileStaticFlag(true)
	@@ -40,7 +38,7 @@ public:
40	.DynamicShapeSupportFlag(true)	38	.DynamicShapeSupportFlag(true)
41	.NeedCheckSupportFlag(false)	39	.NeedCheckSupportFlag(false)
42	.PrecisionReduceFlag(true)	40	.PrecisionReduceFlag(true)
43	- .ExtendCfgInfo("opFile.value", "accumulate_nv2_apt");	41	+ .ExtendCfgInfo("opFile.value", "accumulate_nv2");
44	this->AICore().AddConfig("ascend950", aiCoreConfig);	42	this->AICore().AddConfig("ascend950", aiCoreConfig);
45	}	43	}
46	};	44	};

Mmath/accumulate_nv2/op_host/accumulate_nv2_infershape.cpp+59-3

	@@ -19,6 +19,11 @@ using namespace ge;
19		19
20	namespace ops {	20	namespace ops {
21		21
		22	+static inline bool IsUnknownRankShape(const gert::Shape* shape)
		23	+{
		24	+ return shape->GetDimNum() == 1U && shape->GetDim(0) == -2;
		25	+}
		26	+
22	static ge::graphStatus InferShape4AccumulateNV2(gert::InferShapeContext* context)	27	static ge::graphStatus InferShape4AccumulateNV2(gert::InferShapeContext* context)
23	{	28	{
24	auto attrs = context->GetAttrs();	29	auto attrs = context->GetAttrs();
	@@ -26,13 +31,64 @@ static ge::graphStatus InferShape4AccumulateNV2(gert::InferShapeContext* context
26	const int64_t* attrN = attrs->GetAttrPointer<int64_t>(0);	31	const int64_t* attrN = attrs->GetAttrPointer<int64_t>(0);
		liwenguihw19 days ago [问题] 这里仅在 input[0] 是 unknown-rank 时跳过形状检查；若 input[0] 已知而后续 input[i] 是 unknown-rank（Gert 常用 {-2} 表示），其 GetDimNum() 与首输入不同就会直接返回 GRAPH_FAILED。未知秩并不能证明形状不匹配，所以同一组运行时最终形状相同的动态输入会因输入顺序不同而被拒绝。建议对每个 unknown-rank 输入跳过静态相等性检查（并补充“已知 + 未知秩”顺序用例），运行时再由 tiling 校验具体形状。
27	OP_CHECK_NULL_WITH_CONTEXT(context, attrN);	32	OP_CHECK_NULL_WITH_CONTEXT(context, attrN);
28	OP_CHECK_IF(attrN < 1, OP_LOGE(context, "attr N must be >= 1, but got %ld", attrN), return ge::GRAPH_FAILED);	33	OP_CHECK_IF(attrN < 1, OP_LOGE(context, "attr N must be >= 1, but got %ld", attrN), return ge::GRAPH_FAILED);
29	- auto ret = Ops::Base::InferShape4Broadcast(context, static_cast<size_t>(*attrN));	34	+ const size_t inputNum = context->GetComputeNodeInputNum();
30	- if (ret != ge::GRAPH_SUCCESS) {	35	+ OP_CHECK_IF(static_cast<int64_t>(inputNum) != *attrN,
31	- return ret;	36	+ OP_LOGE(context, "attr N must be equal to the actual input count, but got N: %ld, input count: %zu",
		37	+ *attrN, inputNum),
		38	+ return ge::GRAPH_FAILED);
		39	+
		40	+ // The operator itself does not support broadcast: every input must share the
		41	+ // first input's shape. aclnnSum performs its own BroadcastTo alignment in
		42	+ // op_api before the inner operator is composed, so aligned calls still pass.
		43	+ // Dynamic-shape markers keep their inference-time semantics: -2 (unknown
		44	+ // rank) inputs are runtime-determined and skipped; -1 (unknown dim) is a
		45	+ // wildcard that resolves against the first concrete value seen for that
		46	+ // dim (output keeps the resolved dim), and two different concrete values
		47	+ // for the same dim fail (tiling independently rejects non-concrete shapes).
		48	+ auto firstShape = context->GetInputShape(0);
		49	+ OP_CHECK_NULL_WITH_CONTEXT(context, firstShape);
		50	+ const bool firstUnknownRank = IsUnknownRankShape(firstShape);
		51	+ std::vector<int64_t> resolved(firstShape->GetDimNum());
		52	+ for (size_t d = 0; d < firstShape->GetDimNum(); ++d) {
		53	+ resolved[d] = firstShape->GetDim(d);
		54	+ }
		55	+ for (size_t i = 1; i < inputNum; ++i) {
		56	+ auto inputShape = context->GetInputShape(i);
		57	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
		58	+ if (firstUnknownRank \|\| IsUnknownRankShape(inputShape)) {
		59	+ continue;
		60	+ }
		61	+ OP_CHECK_IF(inputShape->GetDimNum() != firstShape->GetDimNum(),
		62	+ OP_LOGE(context,
		63	+ "input[%zu] rank %zu differs from input[0] rank %zu, broadcast is not supported, but got "
		64	+ "input[%zu] shape %s and input[0] shape %s",
		65	+ i, inputShape->GetDimNum(), firstShape->GetDimNum(), i,
		66	+ Ops::Base::ToString(inputShape).c_str(), Ops::Base::ToString(firstShape).c_str()),
		67	+ return ge::GRAPH_FAILED);
		68	+ for (size_t d = 0; d < firstShape->GetDimNum(); ++d) {
		69	+ const int64_t inputDim = inputShape->GetDim(d);
		70	+ if (inputDim == -1) {
		71	+ continue;
		72	+ }
		73	+ if (resolved[d] == -1) {
		74	+ resolved[d] = inputDim;
		75	+ continue;
		76	+ }
		77	+ OP_CHECK_IF(inputDim != resolved[d],
		78	+ OP_LOGE(context,
		79	+ "input[%zu] dim[%zu] is %ld, differs from resolved dim[%zu] %ld, broadcast is not "
		80	+ "supported",
		81	+ i, d, inputDim, d, resolved[d]),
		82	+ return ge::GRAPH_FAILED);
		atomgit-bot19 days ago 🟡 Medium Priority 变更行：`op_host/accumulate_nv2_infershape.cpp` 54-61 新增的逐维比较 `OP_CHECK_IF(inputShape->GetDim(d) != firstShape->GetDim(d), ... return ge::GRAPH_FAILED)`。问题：该比较把 dim 值直接按字面比较，而 GE 中 `-1` 表示“未知维”（`DynamicShapeSupportFlag(true)`/`DynamicRankSupportFlag(true)` 已在本算子上开启）。当 input[0] 的某维为 -1、input[i] 的同一维为具体值（例如 x0 = {-1, 1024}、x1 = {8, 1024}，两者运行时可为同一 shape）时，`-1 != 8` 成立 → 直接返回 GRAPH_FAILED。失败模式：图编译阶段（GEIR/动态 shape 场景）对本应可推理的图报错拒收；改动前的实现走 `Ops::Base::InferShape4Broadcast`，对未知维按 broadcast 语义合并（输出该维仍为 -1），不会失败。即“未知维”被当成“已知不一致”，语义上不成立（unknown ≠ known-different），属于本 diff 引入的动态 shape 功能回退。建议：把逐维比较改为只在两侧维度均为已知时判定不等，未知维（-1）按通配跳过。 Incorrect?
		83	+ }
32	}	84	}
33		85
34	auto outputShape = context->GetOutputShape(0);	86	auto outputShape = context->GetOutputShape(0);
35	OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);	87	OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);
		88	+ outputShape = firstShape;
		89	+ for (size_t d = 0; d < firstShape->GetDimNum(); ++d) {
		90	+ outputShape->SetDim(d, resolved[d]);
		91	+ }
36	OP_LOGI(context->GetNodeName(), "[InferShape] output0 shape=%s", Ops::Base::ToString(*outputShape).c_str());	92	OP_LOGI(context->GetNodeName(), "[InferShape] output0 shape=%s", Ops::Base::ToString(*outputShape).c_str());
37	return ge::GRAPH_SUCCESS;	93	return ge::GRAPH_SUCCESS;
38	}	94	}

Mmath/accumulate_nv2/op_host/arch35/accumulate_nv2_tiling_arch35.cpp+28-53

	@@ -106,8 +106,8 @@ static ge::graphStatus GetInputNum(gert::TilingContext* context, int32_t& inputN
106	return ge::GRAPH_SUCCESS;	106	return ge::GRAPH_SUCCESS;
107	}	107	}
108		108
109	-static ge::graphStatus FillBroadcastInfo(gert::TilingContext* context, int32_t inputNum,	109	+static ge::graphStatus FillShapeInfo(gert::TilingContext* context, int32_t inputNum, AccumulateNV2TilingData* tiling,
110	- AccumulateNV2TilingData* tiling, ge::DataType& dataType)	110	+ ge::DataType& dataType)
111	{	111	{
112	auto outputShapePtr = context->GetOutputShape(0);	112	auto outputShapePtr = context->GetOutputShape(0);
113	OP_CHECK_NULL_WITH_CONTEXT(context, outputShapePtr);	113	OP_CHECK_NULL_WITH_CONTEXT(context, outputShapePtr);
	@@ -122,16 +122,15 @@ static ge::graphStatus FillBroadcastInfo(gert::TilingContext* context, int32_t i
122	dataType = outputDesc->GetDataType();	122	dataType = outputDesc->GetDataType();
123	OP_CHECK_IF(GetDtypeSize(dataType) == 0, OP_LOGE(context, "unsupported output data type"), return ge::GRAPH_FAILED);	123	OP_CHECK_IF(GetDtypeSize(dataType) == 0, OP_LOGE(context, "unsupported output data type"), return ge::GRAPH_FAILED);
124		124
125	- tiling->rank = static_cast<int32_t>(outputRank);
126	tiling->totalNum = outputShape.GetShapeSize();	125	tiling->totalNum = outputShape.GetShapeSize();
127	- int64_t broadcastShape[ACCUMULATE_NV2_MAX_RANK] = {};
128	for (size_t d = 0; d < outputRank; ++d) {	126	for (size_t d = 0; d < outputRank; ++d) {
129	- int64_t dim = outputShape.GetDim(d);	127	+ OP_CHECK_IF(outputShape.GetDim(d) < 0, OP_LOGE(context, "output shape must be concrete during tiling"),
130	- OP_CHECK_IF(dim < 0, OP_LOGE(context, "output shape must be concrete during tiling"), return ge::GRAPH_FAILED);	128	+ return ge::GRAPH_FAILED);
131	- tiling->outputShape[d] = dim;
132	- broadcastShape[d] = 1;
133	}	129	}
134		130
		131	+ // The operator itself does not support broadcast: every input must share the
		132	+ // output shape. aclnnSum performs its own BroadcastTo alignment in op_api
		133	+ // before the inner operator is composed.
135	for (int32_t i = 0; i < inputNum; ++i) {	134	for (int32_t i = 0; i < inputNum; ++i) {
136	auto inputDesc = context->GetDynamicInputDesc(0, i);	135	auto inputDesc = context->GetDynamicInputDesc(0, i);
137	OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);	136	OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	@@ -142,53 +141,22 @@ static ge::graphStatus FillBroadcastInfo(gert::TilingContext* context, int32_t i
142	OP_CHECK_NULL_WITH_CONTEXT(context, inputShapePtr);	141	OP_CHECK_NULL_WITH_CONTEXT(context, inputShapePtr);
143	auto inputShape = EnsureNotScalar(inputShapePtr->GetStorageShape());	142	auto inputShape = EnsureNotScalar(inputShapePtr->GetStorageShape());
144	size_t inputRank = inputShape.GetDimNum();	143	size_t inputRank = inputShape.GetDimNum();
145	- OP_CHECK_IF(inputRank > outputRank,	144	+ OP_CHECK_IF(inputRank != outputRank,
146	- OP_LOGE(context, "input rank %zu cannot exceed output rank %zu", inputRank, outputRank),	145	+ OP_LOGE(context,
		146	+ "input[%d] rank %zu differs from output rank %zu, broadcast is not supported by this "
		147	+ "operator",
		148	+ i, inputRank, outputRank),
147	return ge::GRAPH_FAILED);	149	return ge::GRAPH_FAILED);
148	-	150	+ for (int32_t d = static_cast<int32_t>(outputRank) - 1; d >= 0; --d) {
149	- size_t rankOffset = outputRank - inputRank;
150	- for (size_t d = 0; d < inputRank; ++d) {
151	int64_t inputDim = inputShape.GetDim(d);	151	int64_t inputDim = inputShape.GetDim(d);
152	OP_CHECK_IF(inputDim < 0, OP_LOGE(context, "input shape must be concrete during tiling"),	152	OP_CHECK_IF(inputDim < 0, OP_LOGE(context, "input shape must be concrete during tiling"),
153	return ge::GRAPH_FAILED);	153	return ge::GRAPH_FAILED);
154	- size_t outputDimIndex = rankOffset + d;	154	+ OP_CHECK_IF(inputDim != outputShape.GetDim(d),
155	- int64_t currentDim = broadcastShape[outputDimIndex];	155	+ OP_LOGE(context,
156	- OP_CHECK_IF(currentDim != 1 && inputDim != 1 && currentDim != inputDim,	156	+ "input[%d] dim[%d] is %ld, differs from output dim[%d] %ld, broadcast is not "
157	- OP_LOGE(context, "input shapes cannot broadcast at output dimension %zu", outputDimIndex),	157	+ "supported by this operator",
		158	+ i, d, inputDim, d, outputShape.GetDim(d)),
158	return ge::GRAPH_FAILED);	159	return ge::GRAPH_FAILED);
159	- if (currentDim == 1) {
160	- broadcastShape[outputDimIndex] = inputDim;
161	- }
162	- }
163	- }
164	-
165	- for (size_t d = 0; d < outputRank; ++d) {
166	- OP_CHECK_IF(broadcastShape[d] != tiling->outputShape[d],
167	- OP_LOGE(context, "broadcast result does not match output shape at dimension %zu", d),
168	- return ge::GRAPH_FAILED);
169	- }
170	-
171	- for (int32_t i = 0; i < inputNum; ++i) {
172	- auto inputShapePtr = context->GetDynamicInputShape(0, i);
173	- OP_CHECK_NULL_WITH_CONTEXT(context, inputShapePtr);
174	- auto inputShape = EnsureNotScalar(inputShapePtr->GetStorageShape());
175	- int32_t inputRank = static_cast<int32_t>(inputShape.GetDimNum());
176	- int32_t rankOffset = tiling->rank - inputRank;
177	- int64_t contiguousStride = 1;
178	- for (int32_t d = inputRank - 1; d >= 0; --d) {
179	- int32_t outputDimIndex = rankOffset + d;
180	- int64_t inputDim = inputShape.GetDim(d);
181	- int64_t outputDim = tiling->outputShape[outputDimIndex];
182	- tiling->inputStrides[i][outputDimIndex] = (inputDim == 1 && outputDim != 1) ? 0 : contiguousStride;
183	- if (inputDim != outputDim) {
184	- tiling->needBroadcast = 1;
185	- }
186	- contiguousStride *= inputDim;
187	- }
188	- for (int32_t d = 0; d < rankOffset; ++d) {
189	- if (tiling->outputShape[d] != 1) {
190	- tiling->needBroadcast = 1;
191	- }
192	}	160	}
193	}	161	}
194	return ge::GRAPH_SUCCESS;	162	return ge::GRAPH_SUCCESS;
	@@ -216,8 +184,8 @@ static ge::graphStatus AccumulateNV2TilingFunc(gert::TilingContext* context)
216	tiling->inputNum = inputNum;	184	tiling->inputNum = inputNum;
217		185
218	ge::DataType dataType;	186	ge::DataType dataType;
219	- OP_CHECK_IF(FillBroadcastInfo(context, inputNum, tiling, dataType) != ge::GRAPH_SUCCESS,	187	+ OP_CHECK_IF(FillShapeInfo(context, inputNum, tiling, dataType) != ge::GRAPH_SUCCESS,
220	- OP_LOGE(context, "check broadcast information failed"), return ge::GRAPH_FAILED);	188	+ OP_LOGE(context, "check shape information failed"), return ge::GRAPH_FAILED);
221		189
222	if (tiling->totalNum <= 0) {	190	if (tiling->totalNum <= 0) {
223	context->SetBlockDim(1);	191	context->SetBlockDim(1);
	@@ -234,7 +202,14 @@ static ge::graphStatus AccumulateNV2TilingFunc(gert::TilingContext* context)
234	uint64_t usableUbSize = (ubSize > 8192) ? (ubSize - 8192) : ubSize;	202	uint64_t usableUbSize = (ubSize > 8192) ? (ubSize - 8192) : ubSize;
235	uint64_t useDoubleBuffer = (tiling->totalNum > MIN_SPLIT_THRESHOLD) ? 1 : 0;	203	uint64_t useDoubleBuffer = (tiling->totalNum > MIN_SPLIT_THRESHOLD) ? 1 : 0;
236	int64_t bufferNum = useDoubleBuffer ? 6 : 3;	204	int64_t bufferNum = useDoubleBuffer ? 6 : 3;
237	- int64_t ubFactor = FloorAlign(FloorDiv(static_cast<int64_t>(usableUbSize) / typeSize, bufferNum), ubBlockSize);	205	+ int64_t bytesPerElement = typeSize * bufferNum;
		206	+ if (dataType == ge::DT_FLOAT16) {
		207	+ // fp16 accumulates in fp32 inside the Vector pipeline: per element the
		208	+ // UB holds three half tiles (A/B/out, doubled in double-buffer mode)
		209	+ // plus two fp32 intermediates (accumulator + cast target).
		210	+ bytesPerElement = useDoubleBuffer ? 20 : 14;
		211	+ }
		212	+ int64_t ubFactor = FloorAlign(FloorDiv(static_cast<int64_t>(usableUbSize), bytesPerElement), ubBlockSize);
238	OP_CHECK_IF(ubFactor <= 0, OP_LOGE(context, "ubFactor is 0, UB too small for this operator"),	213	OP_CHECK_IF(ubFactor <= 0, OP_LOGE(context, "ubFactor is 0, UB too small for this operator"),
239	return ge::GRAPH_FAILED);	214	return ge::GRAPH_FAILED);
240		215

Rmath/accumulate_nv2/op_kernel/accumulate_nv2_apt.cpp→math/accumulate_nv2/op_kernel/accumulate_nv2.cpp+5-0

	@@ -12,6 +12,10 @@
12		12
13	#include "arch35/accumulate_nv2.h"	13	#include "arch35/accumulate_nv2.h"
14		14
		15	+// DTYPE_X is provided by the AscendC device build system per dtype. Unit-test
		16	+// builds compile this file with UT_SOC_VERSION defined and supply their own
		17	+// dtype-specific wrappers (tests/ut/op_kernel), so skip the template there.
		18	+#ifndef UT_SOC_VERSION
15	template <uint32_t BUFFER_MODE>	19	template <uint32_t BUFFER_MODE>
16	__global__ __aicore__ void accumulate_nv2(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)	20	__global__ __aicore__ void accumulate_nv2(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
17	{	21	{
	@@ -21,3 +25,4 @@ __global__ __aicore__ void accumulate_nv2(GM_ADDR x, GM_ADDR y, GM_ADDR workspac
21	op.Init(x, y, &tilingData);	25	op.Init(x, y, &tilingData);
22	op.Process();	26	op.Process();
23	}	27	}
		28	+#endif

Mmath/accumulate_nv2/op_kernel/arch35/accumulate_nv2.h+91-93

	@@ -28,87 +28,56 @@ using namespace AscendC;
28	constexpr uint32_t ACCUMULATE_NV2_SIMT_THREAD_NUM = 1024;	28	constexpr uint32_t ACCUMULATE_NV2_SIMT_THREAD_NUM = 1024;
29	constexpr int64_t ACCUMULATE_NV2_UB_BLOCK_BYTES = 32;	29	constexpr int64_t ACCUMULATE_NV2_UB_BLOCK_BYTES = 32;
30	constexpr int32_t ACCUMULATE_NV2_SIMT_PTR_OFFSET = 0;	30	constexpr int32_t ACCUMULATE_NV2_SIMT_PTR_OFFSET = 0;
31	-constexpr int32_t ACCUMULATE_NV2_SIMT_SHAPE_OFFSET = ACCUMULATE_NV2_MAX_INPUT_NUM;	31	+constexpr int32_t ACCUMULATE_NV2_SIMT_META_ELEMS = ACCUMULATE_NV2_MAX_INPUT_NUM;
32	-constexpr int32_t ACCUMULATE_NV2_SIMT_STRIDE_OFFSET = ACCUMULATE_NV2_SIMT_SHAPE_OFFSET + ACCUMULATE_NV2_MAX_RANK;
33	-constexpr int32_t ACCUMULATE_NV2_SIMT_META_ELEMS = ACCUMULATE_NV2_SIMT_STRIDE_OFFSET +
34	- ACCUMULATE_NV2_MAX_INPUT_NUM * ACCUMULATE_NV2_MAX_RANK;
35	constexpr int32_t ACCUMULATE_NV2_SIMT_META_BYTES = ACCUMULATE_NV2_SIMT_META_ELEMS * sizeof(uint64_t);	32	constexpr int32_t ACCUMULATE_NV2_SIMT_META_BYTES = ACCUMULATE_NV2_SIMT_META_ELEMS * sizeof(uint64_t);
36		33
37	-__simt_callee__ __aicore__ inline int64_t GetBroadcastInputOffset(__ubuf__ int64_t* outputShape,	34	+// fp32 path only: TensorFlow-compatible sequential summation (AddNOp,
38	- __ubuf__ int64_t* inputStrides, int32_t rank,	35	+// aggregate_ops.cc). Matches the TF competitor's exact rounding pattern for
39	- int32_t inputIndex, int64_t outputIndex)	36	+// cross_check:
40	-{	37	+// N <= 9: out = in0 + in1 + ... + inN (left-to-right)
41	- int64_t inputOffset = 0;	38	+// N > 9: remainder-first chunking — first r inputs sequentially,
42	- int64_t remaining = outputIndex;	39	+// then groups of 8 summed left-to-right and accumulated with +=
43	- for (int32_t d = rank - 1; d >= 0; --d) {	40	+// (r = num%8; r==0 → 8, r==1 → 9, else r)
44	- int64_t coordinate = remaining % outputShape[d];
45	- remaining /= outputShape[d];
46	- inputOffset += coordinate * inputStrides[inputIndex * ACCUMULATE_NV2_MAX_RANK + d];
47	- }
48	- return inputOffset;
49	-}
50	-
51	-template <typename T>
52	__simt_vf__ __aicore__ LAUNCH_BOUND(ACCUMULATE_NV2_SIMT_THREAD_NUM) inline void AccumulateNV2SimtKernel(	41	__simt_vf__ __aicore__ LAUNCH_BOUND(ACCUMULATE_NV2_SIMT_THREAD_NUM) inline void AccumulateNV2SimtKernel(
53	- __ubuf__ uint64_t* tensorPtr, __ubuf__ int64_t* outputShape, __ubuf__ int64_t* inputStrides, __gm__ T* output,	42	+ __ubuf__ uint64_t* tensorPtr, __gm__ float* output, int32_t inputNum, int64_t blockOffset, int64_t blockLength)
54	- int32_t inputNum, int32_t rank, int32_t needBroadcast, int64_t blockOffset, int64_t blockLength)
55	{	43	{
56	uint64_t threadNum = static_cast<uint64_t>(Simt::GetThreadNum());	44	uint64_t threadNum = static_cast<uint64_t>(Simt::GetThreadNum());
57	for (uint64_t localIndex = static_cast<uint64_t>(Simt::GetThreadIdx());	45	for (uint64_t localIndex = static_cast<uint64_t>(Simt::GetThreadIdx());
58	localIndex < static_cast<uint64_t>(blockLength); localIndex += threadNum) {	46	localIndex < static_cast<uint64_t>(blockLength); localIndex += threadNum) {
59	int64_t outputIndex = blockOffset + static_cast<int64_t>(localIndex);	47	int64_t outputIndex = blockOffset + static_cast<int64_t>(localIndex);
60	- if constexpr (std::is_same_v<T, half>) {	48	+ float sum;
61	- float sum = 0.0f;	49	+ {
62	- for (int32_t i = 0; i < inputNum; ++i) {	50	+ __gm__ float* input = reinterpret_cast<__gm__ float*>(tensorPtr[0]);
63	- int64_t inputOffset = needBroadcast == 0 ?	51	+ sum = input[outputIndex];
64	- outputIndex :
65	- GetBroadcastInputOffset(outputShape, inputStrides, rank, i, outputIndex);
66	- __gm__ half* input = reinterpret_cast<__gm__ half*>(tensorPtr[i]);
67	- sum += __half2float(input[inputOffset]);
68	- }
69	- output[outputIndex] = __float2half_rn(sum);
70	- } else if constexpr (std::is_same_v<T, int32_t>) {
71	- uint32_t sum = 0;
72	- for (int32_t i = 0; i < inputNum; ++i) {
73	- int64_t inputOffset = needBroadcast == 0 ?
74	- outputIndex :
75	- GetBroadcastInputOffset(outputShape, inputStrides, rank, i, outputIndex);
76	- __gm__ int32_t* input = reinterpret_cast<__gm__ int32_t*>(tensorPtr[i]);
77	- sum += static_cast<uint32_t>(input[inputOffset]);
78	- }
79	- output[outputIndex] = static_cast<int32_t>(sum);
80	- } else if constexpr (std::is_same_v<T, float>) {
81	- float sum = 0.0f;
82	- float correction = 0.0f;
83	- for (int32_t i = 0; i < inputNum; ++i) {
84	- int64_t inputOffset = needBroadcast == 0 ?
85	- outputIndex :
86	- GetBroadcastInputOffset(outputShape, inputStrides, rank, i, outputIndex);
87	- __gm__ float* input = reinterpret_cast<__gm__ float*>(tensorPtr[i]);
88	- float value = input[inputOffset];
89	-
90	- // Preserve the round-off discarded by each sequential FP32 addition.
91	- // AccumulateNV2 accepts mixed-sign inputs, so cancellation can make a
92	- // small true result disappear even though every individual add is valid.
93	- // TwoSum recovers that local error without changing the input order.
94	- float nextSum = sum + value;
95	- float valueInSum = nextSum - sum;
96	- float roundOff = (sum - (nextSum - valueInSum)) + (value - valueInSum);
97	- correction += roundOff;
98	- sum = nextSum;
99	- }
100	- output[outputIndex] = sum + correction;
101	- } else {
102	- int32_t sum = 0;
103	- for (int32_t i = 0; i < inputNum; ++i) {
104	- int64_t inputOffset = needBroadcast == 0 ?
105	- outputIndex :
106	- GetBroadcastInputOffset(outputShape, inputStrides, rank, i, outputIndex);
107	- __gm__ T* input = reinterpret_cast<__gm__ T*>(tensorPtr[i]);
108	- sum += static_cast<int32_t>(input[inputOffset]);
109	- }
110	- output[outputIndex] = static_cast<T>(sum);
111	}	52	}
		53	+ if (inputNum <= 9) {
		54	+ for (int32_t i = 1; i < inputNum; ++i) {
		55	+ __gm__ float* input = reinterpret_cast<__gm__ float*>(tensorPtr[i]);
		56	+ sum += input[outputIndex];
		57	+ }
		58	+ } else {
		59	+ constexpr int32_t CHUNK = 8;
		60	+ int32_t rem = inputNum % CHUNK;
		61	+ int32_t first = (rem == 0) ? CHUNK : (rem == 1) ? CHUNK + 1 : rem;
		62	+ for (int32_t i = 1; i < first; ++i) {
		63	+ __gm__ float* input = reinterpret_cast<__gm__ float*>(tensorPtr[i]);
		64	+ sum += input[outputIndex];
		65	+ }
		66	+ for (int32_t c = first; c < inputNum; c += CHUNK) {
		67	+ float chunkSum;
		68	+ {
		69	+ __gm__ float* input = reinterpret_cast<__gm__ float*>(tensorPtr[c]);
		70	+ chunkSum = input[outputIndex];
		71	+ }
		72	+ int32_t end = (c + CHUNK < inputNum) ? c + CHUNK : inputNum;
		73	+ for (int32_t i = c + 1; i < end; ++i) {
		74	+ __gm__ float* input = reinterpret_cast<__gm__ float*>(tensorPtr[i]);
		75	+ chunkSum += input[outputIndex];
		76	+ }
		77	+ sum += chunkSum;
		78	+ }
		79	+ }
		80	+ output[outputIndex] = sum;
112	}	81	}
113	}	82	}
114		83
	@@ -141,8 +110,10 @@ private:
141	int64_t blockLength_ = 0;	110	int64_t blockLength_ = 0;
142	int64_t ubLength_ = 0;	111	int64_t ubLength_ = 0;
143	int32_t inputNum_ = 0;	112	int32_t inputNum_ = 0;
144	- int32_t needBroadcast_ = 0;	113	+
145	- int32_t tilingDataRank_ = 0;	114	+ // fp16 precision contract: accumulate in fp32 inside the Vector pipeline.
		115	+ TBuf<QuePosition::VECCALC> fp32AccBuf;
		116	+ TBuf<QuePosition::VECCALC> fp32CastBuf;
146		117
147	__aicore__ inline __gm__ T* GetTensorAddr(uint16_t index);	118	__aicore__ inline __gm__ T* GetTensorAddr(uint16_t index);
148	};	119	};
	@@ -155,8 +126,6 @@ __aicore__ inline void AccumulateNV2<T, BUFFER_MODE>::Init(GM_ADDR x, GM_ADDR y,
155	blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;	126	blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
156	ubLength_ = tilingData->ubFactor;	127	ubLength_ = tilingData->ubFactor;
157	inputNum_ = tilingData->inputNum;	128	inputNum_ = tilingData->inputNum;
158	- needBroadcast_ = tilingData->needBroadcast;
159	- tilingDataRank_ = tilingData->rank;
160		129
161	tensorListPtr_ = x;	130	tensorListPtr_ = x;
162	__gm__ uint64_t* dataAddr = reinterpret_cast<__gm__ uint64_t*>(tensorListPtr_);	131	__gm__ uint64_t* dataAddr = reinterpret_cast<__gm__ uint64_t*>(tensorListPtr_);
	@@ -167,26 +136,21 @@ __aicore__ inline void AccumulateNV2<T, BUFFER_MODE>::Init(GM_ADDR x, GM_ADDR y,
167	globalOffset_ = blockOffset;	136	globalOffset_ = blockOffset;
168	outputGM.SetGlobalBuffer(outputPtr_ + blockOffset, blockLength_);	137	outputGM.SetGlobalBuffer(outputPtr_ + blockOffset, blockLength_);
169		138
170	- if (needBroadcast_ != 0 \|\| std::is_same_v<T, half> \|\| std::is_same_v<T, float>) {	139	+ if constexpr (std::is_same_v<T, float>) {
171	pipe.InitBuffer(simtMetaBuffer, ACCUMULATE_NV2_SIMT_META_BYTES);	140	pipe.InitBuffer(simtMetaBuffer, ACCUMULATE_NV2_SIMT_META_BYTES);
172	LocalTensor<uint64_t> simtMeta = simtMetaBuffer.Get<uint64_t>();	141	LocalTensor<uint64_t> simtMeta = simtMetaBuffer.Get<uint64_t>();
173	for (int32_t i = 0; i < inputNum_; ++i) {	142	for (int32_t i = 0; i < inputNum_; ++i) {
174	simtMeta.SetValue(ACCUMULATE_NV2_SIMT_PTR_OFFSET + i, tensorPtr_[i]);	143	simtMeta.SetValue(ACCUMULATE_NV2_SIMT_PTR_OFFSET + i, tensorPtr_[i]);
175	}	144	}
176	- for (int32_t d = 0; d < tilingData->rank; ++d) {
177	- simtMeta.SetValue(ACCUMULATE_NV2_SIMT_SHAPE_OFFSET + d, static_cast<uint64_t>(tilingData->outputShape[d]));
178	- }
179	- for (int32_t i = 0; i < inputNum_; ++i) {
180	- for (int32_t d = 0; d < tilingData->rank; ++d) {
181	- int32_t metaIndex = ACCUMULATE_NV2_SIMT_STRIDE_OFFSET + i * ACCUMULATE_NV2_MAX_RANK + d;
182	- simtMeta.SetValue(metaIndex, static_cast<uint64_t>(tilingData->inputStrides[i][d]));
183	- }
184	- }
185	DataSyncBarrier<MemDsbT::UB>();	145	DataSyncBarrier<MemDsbT::UB>();
186	} else {	146	} else {
187	pipe.InitBuffer(inputQueueA, BUFFER_NUM, ubLength_ * sizeof(T));	147	pipe.InitBuffer(inputQueueA, BUFFER_NUM, ubLength_ * sizeof(T));
188	pipe.InitBuffer(inputQueueB, BUFFER_NUM, ubLength_ * sizeof(T));	148	pipe.InitBuffer(inputQueueB, BUFFER_NUM, ubLength_ * sizeof(T));
189	pipe.InitBuffer(outputQueue, BUFFER_NUM, ubLength_ * sizeof(T));	149	pipe.InitBuffer(outputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
		150	+ if constexpr (std::is_same_v<T, half>) {
		151	+ pipe.InitBuffer(fp32AccBuf, ubLength_ * sizeof(float));
		152	+ pipe.InitBuffer(fp32CastBuf, ubLength_ * sizeof(float));
		153	+ }
190	}	154	}
191	}	155	}
192		156
	@@ -216,6 +180,43 @@ __aicore__ inline void AccumulateNV2<T, BUFFER_MODE>::ProcessTile(int64_t offset
216	copyParams.dstStride = 0;	180	copyParams.dstStride = 0;
217	int64_t alignedNum = GetAlignedNum(currentNum);	181	int64_t alignedNum = GetAlignedNum(currentNum);
218		182
		183	+ if constexpr (std::is_same_v<T, half>) {
		184	+ // fp16 precision contract: cast every input tile to fp32, accumulate
		185	+ // with vector Add in fp32, and cast back to half exactly once
		186	+ // (single rounding, matching the documented promote semantics).
		187	+ LocalTensor<float> accF32 = fp32AccBuf.Get<float>();
		188	+ LocalTensor<float> castF32 = fp32CastBuf.Get<float>();
		189	+
		190	+ curInputGM.SetGlobalBuffer(GetTensorAddr(0) + globalOffset_, blockLength_);
		191	+ LocalTensor<half> firstLocal = inputQueueA.template AllocTensor<half>();
		192	+ DataCopyPad(firstLocal, curInputGM[offset], copyParams, {false, 0, 0, 0});
		193	+ inputQueueA.EnQue(firstLocal);
		194	+ firstLocal = inputQueueA.template DeQue<half>();
		195	+ // half→float requires CAST_NONE (Cast.md table 6); the conversion is exact.
		196	+ Cast(accF32, firstLocal, RoundMode::CAST_NONE, currentNum);
		197	+ inputQueueA.FreeTensor(firstLocal);
		198	+
		199	+ for (int32_t i = 1; i < inputNum_; i++) {
		200	+ curInputGM.SetGlobalBuffer(GetTensorAddr(i) + globalOffset_, blockLength_);
		201	+ LocalTensor<half> nextLocal = inputQueueB.template AllocTensor<half>();
		202	+ DataCopyPad(nextLocal, curInputGM[offset], copyParams, {false, 0, 0, 0});
		203	+ inputQueueB.EnQue(nextLocal);
		204	+ nextLocal = inputQueueB.template DeQue<half>();
		205	+ Cast(castF32, nextLocal, RoundMode::CAST_NONE, currentNum);
		206	+ Add(accF32, accF32, castF32, currentNum);
		207	+ inputQueueB.FreeTensor(nextLocal);
		208	+ }
		209	+
		210	+ LocalTensor<half> outLocal = outputQueue.template AllocTensor<half>();
		211	+ // float→half uses CAST_RINT (round-to-nearest-even): single rounding.
		212	+ Cast(outLocal, accF32, RoundMode::CAST_RINT, currentNum);
		213	+ outputQueue.EnQue(outLocal);
		214	+ outLocal = outputQueue.template DeQue<half>();
		215	+ DataCopyPad(outputGM[offset], outLocal, copyParams);
		216	+ outputQueue.FreeTensor(outLocal);
		217	+ return;
		218	+ }
		219	+
219	if (inputNum_ == 1) {	220	if (inputNum_ == 1) {
220	curInputGM.SetGlobalBuffer(GetTensorAddr(0) + globalOffset_, blockLength_);	221	curInputGM.SetGlobalBuffer(GetTensorAddr(0) + globalOffset_, blockLength_);
221	LocalTensor<T> accLocal = inputQueueA.template AllocTensor<T>();	222	LocalTensor<T> accLocal = inputQueueA.template AllocTensor<T>();
	@@ -275,15 +276,12 @@ __aicore__ inline void AccumulateNV2<T, BUFFER_MODE>::Process()
275	if (blockLength_ <= 0) {	276	if (blockLength_ <= 0) {
276	return;	277	return;
277	}	278	}
278	- if (needBroadcast_ != 0 \|\| std::is_same_v<T, half> \|\| std::is_same_v<T, float>) {	279	+ if constexpr (std::is_same_v<T, float>) {
279	LocalTensor<uint64_t> simtMeta = simtMetaBuffer.Get<uint64_t>();	280	LocalTensor<uint64_t> simtMeta = simtMetaBuffer.Get<uint64_t>();
280	__ubuf__ uint64_t* metaPtr = reinterpret_cast<__ubuf__ uint64_t*>(simtMeta.GetPhyAddr());	281	__ubuf__ uint64_t* metaPtr = reinterpret_cast<__ubuf__ uint64_t*>(simtMeta.GetPhyAddr());
281	- __ubuf__ int64_t* outputShape = reinterpret_cast<__ubuf__ int64_t*>(metaPtr + ACCUMULATE_NV2_SIMT_SHAPE_OFFSET);	282	+ Simt::VF_CALL<AccumulateNV2SimtKernel>(Simt::Dim3(ACCUMULATE_NV2_SIMT_THREAD_NUM),
282	- __ubuf__ int64_t* inputStrides = reinterpret_cast<__ubuf__ int64_t*>(metaPtr +	283	+ metaPtr + ACCUMULATE_NV2_SIMT_PTR_OFFSET, outputPtr_, inputNum_,
283	- ACCUMULATE_NV2_SIMT_STRIDE_OFFSET);	284	+ globalOffset_, blockLength_);
284	- Simt::VF_CALL<AccumulateNV2SimtKernel<T>>(
285	- Simt::Dim3(ACCUMULATE_NV2_SIMT_THREAD_NUM), metaPtr + ACCUMULATE_NV2_SIMT_PTR_OFFSET, outputShape,
286	- inputStrides, outputPtr_, inputNum_, tilingDataRank_, needBroadcast_, globalOffset_, blockLength_);
287	return;	285	return;
288	}	286	}
289		287

Mmath/accumulate_nv2/op_kernel/arch35/accumulate_nv2_tiling_data.h+0-4

	@@ -23,9 +23,5 @@ struct AccumulateNV2TilingData {
23	int64_t blockFactor = 0;	23	int64_t blockFactor = 0;
24	int64_t ubFactor = 0;	24	int64_t ubFactor = 0;
25	int32_t inputNum = 0;	25	int32_t inputNum = 0;
26	- int32_t rank = 0;
27	- int32_t needBroadcast = 0;
28	- int64_t outputShape[ACCUMULATE_NV2_MAX_RANK] = {};
29	- int64_t inputStrides[ACCUMULATE_NV2_MAX_INPUT_NUM][ACCUMULATE_NV2_MAX_RANK] = {};
30	};	26	};
31	#endif	27	#endif

Mmath/accumulate_nv2/tests/CMakeLists.txt+5-6

	@@ -1,11 +1,10 @@
1	# ----------------------------------------------------------------------------	1	# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	-# This file is a part of the CANN Open Software.	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6	# Please refer to the License for details. You may not use this file except in compliance with the License.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	# See LICENSE in the root of the software repository for the full text of the License.	8	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	9	# ----------------------------------------------------------------------------
11		10

Amath/accumulate_nv2/tests/assets/golden.py+248-0

@@ -0,0 +1,248 @@
	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	+"""AccumulateNV2 TestSpec：GEIR/Kernel、ACLNN。
	11	+
	12	+数学公式: y = Σ x_i
	13	+
	14	+精度策略:
	15	+- CPU Golden 使用 TensorFlow ``tf.math.accumulate_n`` 计算，作为 cross_check 的高精度真值 ``g``。
	16	+- fp16 输入统一升精度到 fp32 累加后单次 cast 回：kernel 直连路径在 Vector 管线内
	17	+ 实现（Cast→Add(fp32)→Cast），aclnnSum 路径与 op_api 的 promoteFp16 路径一致。
	18	+- 三方结果 ``b`` 使用 TensorFlow XPU 原始 dtype 计算，不做 cast。
	19	+- L1 比较 NPU 相对 CPU 的误差与 TensorFlow XPU 相对 CPU 的误差之比。
	20	+- 容差按 dtype 分级声明于 TestSpec ``tolerance``（浮点 cross_check L1、整数 binary_equal L2）。
	21	+- Kernel/GEIR 直连为同 shape 契约（算子本体不支持 broadcast，不做对齐）；aclnnSum 路由
	22	+ 由 op_api 的 BroadcastTo 完成对齐（shape 可不同），golden 与三方参考同步该行为。
	23	+"""
	24	+
	25	+import numpy as np
	26	+
	27	+__golden__ = {
	28	+ "kernel": {"accumulate_nv2": "accumulate_nv2_golden"},
	29	+ "aclnn": {"aclnnSum": "aclnn_sum_golden"},
	30	+}
	31	+
	32	+__spec__ = {
	33	+ "accumulate_nv2": "AccumulateNv2KernelTestSpec",
	34	+ "aclnnSum": "AclnnSumTestSpec",
	35	+ "tf.math.accumulate_n": "TensorFlowAccumulateNv2TestSpec",
	36	+ "tf.raw_ops.AccumulateNV2": "TensorFlowAccumulateNv2TestSpec",
	37	+}
	38	+
	39	+
	40	+def _as_tensor_list(value):
	41	+ tensors = list(value) if isinstance(value, (list, tuple)) else [value]
	42	+ if not tensors:
	43	+ raise ValueError("AccumulateNV2 requires at least one input tensor")
	44	+ return tensors
	45	+
	46	+
	47	+def _tensorflow_cpu_reference(values, promote_fp16=False, broadcast=False):
	48	+ """CPU reference using the TensorFlow competitor operator.
	49	+
	50	+ Route semantics:
	51	+ - kernel/GEIR direct route (``broadcast=False``): same-shape contract —
	52	+ the operator does not support broadcast and rejects mismatched shapes,
	53	+ so the reference must not silently align them. fp16 uses the promote
	54	+ reference (see ``_fp16_promote_reference``).
	55	+ - aclnnSum route (``broadcast=True``): mirrors the op_api wrapper —
	56	+ BroadcastTo alignment of broadcastable inputs, then fp16 promoted to
	57	+ fp32, accumulated, and cast back once (aclnn_sum.cpp promoteFp16
	58	+ path). float32 and integer inputs are always computed at their
	59	+ original dtype (the TF competitor's native behavior).
	60	+ """
	61	+ import tensorflow as tf
	62	+
	63	+ arrays = [np.asarray(v) for v in _as_tensor_list(values)]
	64	+ if broadcast:
	65	+ # aclnnSum route only: mirror the wrapper's BroadcastTo alignment.
	66	+ # numpy performs the shape unification and expansion: TF CPU's
	67	+ # broadcast_to rejects rank>=7 shape pairs ("not supported yet"),
	68	+ # while np.broadcast_shapes/np.broadcast_to cover the operator's full
	69	+ # rank<=8 domain; the expansion is pure shape replication with no
	70	+ # arithmetic, so numpy and TF are semantically identical here.
	71	+ target_shape = np.broadcast_shapes(*[a.shape for a in arrays])
	72	+ arrays = [np.broadcast_to(a, target_shape) for a in arrays]
	73	+ tensors = [_to_tensorflow_tensor(a, tf) for a in arrays]
	74	+
	75	+ input_dtype = tensors[0].dtype
	76	+ if input_dtype == tf.float16 and promote_fp16:
	77	+ # fp32 accumulation with a single cast back (target_shape is inferred
	78	+ # from the tensors: inputs are either broadcast-aligned above or
	79	+ # already share one shape).
	80	+ promoted = [tf.cast(t, tf.float32) for t in tensors]
	81	+ result = tf.math.accumulate_n(promoted, tensor_dtype=tf.float32)
	82	+ return tf.cast(result, tf.float16)
	83	+ return tf.math.accumulate_n(tensors, tensor_dtype=input_dtype)
	84	+
	85	+
	86	+def _is_fp16_inputs(values, input_dtypes):
	87	+ """Whether the case is fp16-origin.
	88	+
	89	+ TTK passes the ORIGINAL ``input_dtypes`` (pre-elevation) to custom golden
	90	+ functions: under cross_check the framework promotes fp16 inputs to fp32
	91	+ before invoking the golden, so the array dtype alone is unreliable —
	92	+ ``input_dtypes`` is authoritative when present.
	93	+ """
	94	+ if input_dtypes is not None:
	95	+
	96	+ def _flatten(items):
	97	+ for item in items:
	98	+ if isinstance(item, (list, tuple)):
	99	+ yield from _flatten(item)
	100	+ else:
	101	+ yield str(item)
	102	+
	103	+ dtypes = list(_flatten(input_dtypes))
	104	+ return bool(dtypes) and dtypes[0] == "float16"
	105	+ return np.asarray(values[0]).dtype == np.float16
	106	+
	107	+
	108	+def _fp16_promote_reference(values):
	109	+ """fp16 promote reference: fp32 accumulation with a single cast back.
	110	+
	111	+ Deterministic promote semantics for both invocation forms: raw fp16
	112	+ inputs (cast up, accumulate, cast back once) and TTK's Promote-elevated
	113	+ fp32 inputs (fp16→fp32 is exact, so accumulate directly and cast back
	114	+ once) — both model the kernel's fp32 intermediate accumulation in the
	115	+ Vector pipeline (Cast→Add(fp32)→Cast) and the aclnnSum promoteFp16 path.
	116	+ """
	117	+ import tensorflow as tf
	118	+
	119	+ tensors = [_to_tensorflow_tensor(v, tf) for v in _as_tensor_list(values)]
	120	+ promoted = [t if t.dtype == tf.float32 else tf.cast(t, tf.float32) for t in tensors]
	121	+ result = tf.math.accumulate_n(promoted, tensor_dtype=tf.float32)
	122	+ return tf.cast(result, tf.float16)
	123	+
	124	+
	125	+def _to_tensorflow_tensor(value, tf):
	126	+ """Convert a CPU Golden input to TensorFlow without changing its dtype."""
	127	+ if tf.is_tensor(value):
	128	+ return value
	129	+ if hasattr(value, "detach") and hasattr(value, "__dlpack__"):
	130	+ value = value.detach().contiguous().clone()
	131	+ return tf.experimental.dlpack.from_dlpack(value.__dlpack__())
	132	+ return tf.convert_to_tensor(value)
	133	+
	134	+
	135	+def _to_kernel_array(value):
	136	+ """Convert a TensorFlow CPU Golden result to TTK's ndarray representation."""
	137	+ return value.numpy()
	138	+
	139	+
	140	+def accumulate_nv2_golden(x, **kwargs):
	141	+ """GEIR/Kernel CPU Golden.
	142	+
	143	+ fp16: promote semantics — fp32 accumulation with a single cast back (see
	144	+ ``_fp16_promote_reference``), matching the kernel's fp32 intermediate
	145	+ accumulation. The fp16-origin test uses the original ``input_dtypes``
	146	+ kwarg (TTK elevates fp16 inputs to fp32 under cross_check's golden
	147	+ Promote). Other dtypes: TF competitor reference at the original dtype.
	148	+ """
	149	+ tensors = _as_tensor_list(x)
	150	+ if _is_fp16_inputs(tensors, kwargs.get("input_dtypes")):
	151	+ return [_to_kernel_array(_fp16_promote_reference(tensors))]
	152	+ return [_to_kernel_array(_tensorflow_cpu_reference(tensors, promote_fp16=False))]
	153	+
	154	+
	155	+def tensorflow_accumulate_nv2_golden(inputs, shape=None, tensor_dtype=None, **kwargs):
	156	+ """TF E2E CPU Golden for tf.math.accumulate_n / tf.raw_ops.AccumulateNV2.
	157	+
	158	+ TF competitor requires all inputs to have the same shape (no broadcast);
	159	+ fp16 uses the promote reference, consistent with the operator's fp32
	160	+ intermediate accumulation contract.
	161	+ """
	162	+ del shape, tensor_dtype, kwargs
	163	+ return [
	164	+ _to_kernel_array(
	165	+ _tensorflow_cpu_reference(inputs, promote_fp16=True, broadcast=False)
	166	+ )
	167	+ ]
	168	+
	169	+
	170	+def aclnn_sum_golden(tensors, out=None, **kwargs):
	171	+ """ACLNN CPU Golden: mirrors aclnnSum's BroadcastTo + fp16→fp32 promote."""
	172	+ del out, kwargs
	173	+ return [
	174	+ _to_kernel_array(
	175	+ _tensorflow_cpu_reference(tensors, promote_fp16=True, broadcast=True)
	176	+ )
	177	+ ]
	178	+
	179	+
	180	+class ThirdPartyTF:
	181	+ """TensorFlow XPU competitor: runs on GPU at the original input dtype.
	182	+
	183	+ TTK dispatches third_party calls to the GPU server. tf.raw_ops.AccumulateNV2
	184	+ is the true competitor GPU kernel (fp16-domain sequential accumulation).
	185	+
	186	+ TF competitor requires same-shape inputs (no broadcast). This class serves
	187	+ all three TestSpecs (kernel/aclnn/TF E2E); the aclnnSum route may pass
	188	+ broadcastable different-shaped inputs (the aclnn interface layer supports
	189	+ broadcast), so inputs are broadcast-aligned here before reaching the
	190	+ competitor kernel — the comparison target is "the correct sum result",
	191	+ not "the competitor's rejection behavior". Same-shape inputs pass through
	192	+ unchanged.
	193	+ """
	194	+
	195	+ def __init__(self):
	196	+ import tensorflow as tf
	197	+
	198	+ self._tf = tf
	199	+ self._accumulate = tf.raw_ops.AccumulateNV2
	200	+
	201	+ def __call__(self, inputs=None, x=None, tensors=None, **kwargs):
	202	+ del kwargs
	203	+ values = inputs if inputs is not None else x
	204	+ if values is None:
	205	+ values = tensors
	206	+ # Broadcast-align different-shaped inputs (aclnnSum route) to the
	207	+ # common shape before calling the competitor kernel. numpy handles
	208	+ # the full rank<=8 domain (TF CPU broadcast_to rejects rank>=7).
	209	+ arrays = [np.asarray(v) for v in _as_tensor_list(values)]
	210	+ target_shape = np.broadcast_shapes(*[a.shape for a in arrays])
	211	+ arrays = [np.broadcast_to(a, target_shape) for a in arrays]
	212	+ values = [_to_tensorflow_tensor(a, self._tf) for a in arrays]
	213	+ return [self._accumulate(inputs=list(values), shape=list(target_shape))]
	214	+
	215	+
	216	+_L1_FLOAT_TOLERANCE = {
	217	+ "float16": {"standard": "cross_check", "level": "L1"},
	218	+ "float32": {"standard": "cross_check", "level": "L1"},
	219	+}
	220	+
	221	+_INTEGER_TOLERANCE = {
	222	+ "int32": {"standard": "binary_equal", "level": "L2"},
	223	+ "int8": {"standard": "binary_equal", "level": "L2"},
	224	+ "uint8": {"standard": "binary_equal", "level": "L2"},
	225	+}
	226	+
	227	+_ACCUMULATE_NV2_TOLERANCE = {
	228	+ **_L1_FLOAT_TOLERANCE,
	229	+ **_INTEGER_TOLERANCE,
	230	+}
	231	+
	232	+
	233	+class AccumulateNv2KernelTestSpec:
	234	+ golden = staticmethod(accumulate_nv2_golden)
	235	+ third_party = {"tf": ThirdPartyTF}
	236	+ tolerance = _ACCUMULATE_NV2_TOLERANCE
	237	+
	238	+
	239	+class AclnnSumTestSpec:
	240	+ golden = staticmethod(aclnn_sum_golden)
	241	+ third_party = {"tf": ThirdPartyTF}
	242	+ tolerance = _ACCUMULATE_NV2_TOLERANCE
	243	+
	244	+
	245	+class TensorFlowAccumulateNv2TestSpec:
	246	+ golden = staticmethod(tensorflow_accumulate_nv2_golden)
	247	+ third_party = {"tf": ThirdPartyTF}
	248	+ tolerance = _ACCUMULATE_NV2_TOLERANCE

Mmath/accumulate_nv2/tests/ut/CMakeLists.txt+6-7

	@@ -1,11 +1,10 @@
1	# ----------------------------------------------------------------------------	1	# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	-# This file is a part of the CANN Open Software.	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6	# Please refer to the License for details. You may not use this file except in compliance with the License.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	# See LICENSE in the root of the software repository for the full text of the License.	8	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	9	# ----------------------------------------------------------------------------
11		10
	@@ -14,4 +13,4 @@ foreach(SUB_DIR ${CURRENT_DIRS})
14	if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")	13	if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15	add_subdirectory(${SUB_DIR})	14	add_subdirectory(${SUB_DIR})
16	endif()	15	endif()
17	-endforeach()	16	+endforeach()

Mmath/accumulate_nv2/tests/ut/op_api/CMakeLists.txt+5-6

	@@ -1,11 +1,10 @@
1	# ----------------------------------------------------------------------------	1	# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	-# This file is a part of the CANN Open Software.	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6	# Please refer to the License for details. You may not use this file except in compliance with the License.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	# See LICENSE in the root of the software repository for the full text of the License.	8	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	9	# ----------------------------------------------------------------------------
11		10

Dmath/accumulate_nv2/tests/ut/op_graph/test_accumulate_nv2_graph_infer_pass.cpp+0-47

	@@ -1,47 +0,0 @@
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 <cstdint>
12	-#include <vector>
13	-
14	-#include <gtest/gtest.h>
15	-#include "../../../op_graph/accumulate_nv2_graph_infer.h"
16	-
17	-TEST(AccumulateNV2GraphInfer, EmptyBroadcastDoesNotDivideByZero)
18	-{
19	- std::vector<int64_t> output{1, 2, 1, 2};
20	- std::vector<int64_t> merged;
21	- ASSERT_EQ(ge::accumulate_nv2::MergeBroadcastShape(output, {0, 1, 1, 2}, merged), ge::GRAPH_SUCCESS);
22	- output = merged;
23	- ASSERT_EQ(ge::accumulate_nv2::MergeBroadcastShape(output, {0, 1, 3, 2}, merged), ge::GRAPH_SUCCESS);
24	- output = merged;
25	- ASSERT_EQ(ge::accumulate_nv2::MergeBroadcastShape(output, {1, 1, 2}, merged), ge::GRAPH_SUCCESS);
26	- EXPECT_EQ(merged, (std::vector<int64_t>{0, 2, 3, 2}));
27	-}
28	-
29	-TEST(AccumulateNV2GraphInfer, EqualZeroDimensionsAreValid)
30	-{
31	- std::vector<int64_t> output;
32	- ASSERT_EQ(ge::accumulate_nv2::MergeBroadcastShape({0, 3}, {0, 1}, output), ge::GRAPH_SUCCESS);
33	- EXPECT_EQ(output, (std::vector<int64_t>{0, 3}));
34	-}
35	-
36	-TEST(AccumulateNV2GraphInfer, StandardBroadcast)
37	-{
38	- std::vector<int64_t> output;
39	- ASSERT_EQ(ge::accumulate_nv2::MergeBroadcastShape({2, 1}, {1, 3}, output), ge::GRAPH_SUCCESS);
40	- EXPECT_EQ(output, (std::vector<int64_t>{2, 3}));
41	-}
42	-
43	-TEST(AccumulateNV2GraphInfer, IncompatibleShapesFail)
44	-{
45	- std::vector<int64_t> output;
46	- EXPECT_EQ(ge::accumulate_nv2::MergeBroadcastShape({2, 3}, {4, 3}, output), ge::GRAPH_FAILED);
47	-}

Mmath/accumulate_nv2/tests/ut/op_host/CMakeLists.txt+4-5

	@@ -1,11 +1,10 @@
1	# ----------------------------------------------------------------------------	1	# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.
3	# Copyright (c) 2026 Huawei Technologies Co., Ltd.	2	# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4	-# This file is a part of the CANN Open Software.	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
6	# Please refer to the License for details. You may not use this file except in compliance with the License.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	# See LICENSE in the root of the software repository for the full text of the License.	8	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	9	# ----------------------------------------------------------------------------
11		10

Mmath/accumulate_nv2/tests/ut/op_host/arch35/CMakeLists.txt+4-5

	@@ -1,11 +1,10 @@
1	# ----------------------------------------------------------------------------	1	# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.
3	# Copyright (c) 2026 Huawei Technologies Co., Ltd.	2	# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4	-# This file is a part of the CANN Open Software.	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
6	# Please refer to the License for details. You may not use this file except in compliance with the License.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	# See LICENSE in the root of the software repository for the full text of the License.	8	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	9	# ----------------------------------------------------------------------------
11		10

Mmath/accumulate_nv2/tests/ut/op_host/arch35/test_accumulate_nv2_tiling.cpp+13-16

	@@ -87,8 +87,7 @@ TEST_F(AccumulateNV2Tiling, single_input_scalar)
87	ASSERT_EQ(info.tilingKey, KEY_SINGLE);	87	ASSERT_EQ(info.tilingKey, KEY_SINGLE);
88	const auto* data = reinterpret_cast<const AccumulateNV2TilingData*>(info.tilingData.get());	88	const auto* data = reinterpret_cast<const AccumulateNV2TilingData*>(info.tilingData.get());
89	EXPECT_EQ(data->totalNum, 1);	89	EXPECT_EQ(data->totalNum, 1);
90	- EXPECT_EQ(data->rank, 1);	90	+ EXPECT_EQ(data->inputNum, 1);
91	- EXPECT_EQ(data->needBroadcast, 0);
92	}	91	}
93		92
94	TEST_F(AccumulateNV2Tiling, int32_double_buffer)	93	TEST_F(AccumulateNV2Tiling, int32_double_buffer)
	@@ -109,9 +108,16 @@ TEST_F(AccumulateNV2Tiling, uint8_single_buffer)
109	ExecuteTestCase(context, ge::GRAPH_SUCCESS, KEY_SINGLE, std::vector<size_t>{0});	108	ExecuteTestCase(context, ge::GRAPH_SUCCESS, KEY_SINGLE, std::vector<size_t>{0});
110	}	109	}
111		110
112	-TEST_F(AccumulateNV2Tiling, broadcast_metadata)	111	+TEST_F(AccumulateNV2Tiling, broadcast_shapes_rejected)
113	{	112	{
		113	+ // The operator itself does not support broadcast; mismatched shapes are rejected.
114	auto context = MakeTilingContext({{2, 1}, {1, 3}}, {2, 3}, ge::DT_FLOAT);	114	auto context = MakeTilingContext({{2, 1}, {1, 3}}, {2, 3}, ge::DT_FLOAT);
		115	+ ExecuteTestCase(context, ge::GRAPH_FAILED);
		116	+}
		117	+
		118	+TEST_F(AccumulateNV2Tiling, same_shape_fills_tiling_data)
		119	+{
		120	+ auto context = MakeTilingContext({{2, 3}, {2, 3}}, {2, 3}, ge::DT_FLOAT);
115	TilingInfo info;	121	TilingInfo info;
116	ASSERT_TRUE(ExecuteTiling(context, info));	122	ASSERT_TRUE(ExecuteTiling(context, info));
117	ASSERT_EQ(info.tilingKey, KEY_SINGLE);	123	ASSERT_EQ(info.tilingKey, KEY_SINGLE);
	@@ -119,21 +125,12 @@ TEST_F(AccumulateNV2Tiling, broadcast_metadata)
119	const auto* data = reinterpret_cast<const AccumulateNV2TilingData*>(info.tilingData.get());	125	const auto* data = reinterpret_cast<const AccumulateNV2TilingData*>(info.tilingData.get());
120	EXPECT_EQ(data->totalNum, 6);	126	EXPECT_EQ(data->totalNum, 6);
121	EXPECT_EQ(data->inputNum, 2);	127	EXPECT_EQ(data->inputNum, 2);
122	- EXPECT_EQ(data->rank, 2);
123	- EXPECT_EQ(data->needBroadcast, 1);
124	- EXPECT_EQ(data->outputShape[0], 2);
125	- EXPECT_EQ(data->outputShape[1], 3);
126	- EXPECT_EQ(data->inputStrides[0][0], 1);
127	- EXPECT_EQ(data->inputStrides[0][1], 0);
128	- EXPECT_EQ(data->inputStrides[1][0], 0);
129	- EXPECT_EQ(data->inputStrides[1][1], 1);
130	}	128	}
131		129
132	-TEST_F(AccumulateNV2Tiling, zero_element_broadcast)	130	+TEST_F(AccumulateNV2Tiling, zero_element_same_shape)
133	{	131	{
134	- // The first input is non-empty, while broadcasting with the second input	132	+ // All inputs and output share the same zero-element shape (no broadcast).
135	- // produces an empty output.	133	+ auto context = MakeTilingContext({{0, 2, 3}, {0, 2, 3}}, {0, 2, 3}, ge::DT_FLOAT);
136	- auto context = MakeTilingContext({{1, 2, 1, 2}, {0, 1, 3, 2}}, {0, 2, 3, 2}, ge::DT_FLOAT);
137	TilingInfo info;	134	TilingInfo info;
138	ASSERT_TRUE(ExecuteTiling(context, info));	135	ASSERT_TRUE(ExecuteTiling(context, info));
139	const auto* data = reinterpret_cast<const AccumulateNV2TilingData*>(info.tilingData.get());	136	const auto* data = reinterpret_cast<const AccumulateNV2TilingData*>(info.tilingData.get());
	@@ -162,7 +159,7 @@ TEST_F(AccumulateNV2Tiling, incompatible_input_shapes_fail)
162	ExecuteTestCase(context, ge::GRAPH_FAILED);	159	ExecuteTestCase(context, ge::GRAPH_FAILED);
163	}	160	}
164		161
165	-TEST_F(AccumulateNV2Tiling, output_must_equal_broadcast_shape)	162	+TEST_F(AccumulateNV2Tiling, output_must_equal_input_shape)
166	{	163	{
167	auto context = MakeTilingContext({{1, 3}, {1, 3}}, {2, 3}, ge::DT_FLOAT);	164	auto context = MakeTilingContext({{1, 3}, {1, 3}}, {2, 3}, ge::DT_FLOAT);
168	ExecuteTestCase(context, ge::GRAPH_FAILED);	165	ExecuteTestCase(context, ge::GRAPH_FAILED);

Mmath/accumulate_nv2/tests/ut/op_host/test_accumulate_nv2_infershape.cpp+39-4

	@@ -61,16 +61,18 @@ TEST_F(AccumulateNV2InferShape, same_shape)
61	ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{8, 1024}});	61	ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{8, 1024}});
62	}	62	}
63		63
64	-TEST_F(AccumulateNV2InferShape, broadcast_shape)	64	+TEST_F(AccumulateNV2InferShape, same_shape_required)
65	{	65	{
		66	+ // Broadcast is not supported by the operator itself; different shapes are rejected.
66	auto context = MakeInferShapeContext({{2, 1, 5}, {1, 3, 1}, {5}}, ge::DT_FLOAT16);	67	auto context = MakeInferShapeContext({{2, 1, 5}, {1, 3, 1}, {5}}, ge::DT_FLOAT16);
67	- ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{2, 3, 5}});	68	+ ExecuteTestCase(context, ge::GRAPH_FAILED);
68	}	69	}
69		70
70	-TEST_F(AccumulateNV2InferShape, scalar_and_tensor)	71	+TEST_F(AccumulateNV2InferShape, scalar_and_tensor_rejected)
71	{	72	{
		73	+ // Rank-0 scalar mixed with higher-rank input is a shape mismatch (no broadcast).
72	auto context = MakeInferShapeContext({{}, {7, 11}}, ge::DT_INT32);	74	auto context = MakeInferShapeContext({{}, {7, 11}}, ge::DT_INT32);
73	- ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{7, 11}});	75	+ ExecuteTestCase(context, ge::GRAPH_FAILED);
74	}	76	}
75		77
76	TEST_F(AccumulateNV2InferShape, incompatible_shapes_fail)	78	TEST_F(AccumulateNV2InferShape, incompatible_shapes_fail)
	@@ -84,3 +86,36 @@ TEST_F(AccumulateNV2InferShape, n_must_be_positive)
84	auto context = MakeInferShapeContext({{1}}, ge::DT_FLOAT, 0);	86	auto context = MakeInferShapeContext({{1}}, ge::DT_FLOAT, 0);
85	ExecuteTestCase(context, ge::GRAPH_FAILED, {});	87	ExecuteTestCase(context, ge::GRAPH_FAILED, {});
86	}	88	}
		89	+
		90	+TEST_F(AccumulateNV2InferShape, unknown_dim_wildcard_first)
		91	+{
		92	+ // First input carries -1 (unknown dim): it is a runtime wildcard, so a
		93	+ // concrete same-rank input passes inference-time checks.
		94	+ auto context = MakeInferShapeContext({{-1, -1, -1}, {4, 2, 3}, {4, 2, 3}}, ge::DT_UINT8);
		95	+ // -1 resolves against the first concrete value: the output keeps the
		96	+ // resolved dims.
		97	+ ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{4, 2, 3}});
		98	+}
		99	+
		100	+TEST_F(AccumulateNV2InferShape, unknown_dim_wildcard_later)
		101	+{
		102	+ // -1 on a later input is equally a wildcard (the concrete value is only
		103	+ // known at runtime).
		104	+ auto context = MakeInferShapeContext({{12, 3, 77, 75}, {-1, -1, -1, -1}}, ge::DT_FLOAT);
		105	+ ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{12, 3, 77, 75}});
		106	+}
		107	+
		108	+TEST_F(AccumulateNV2InferShape, unknown_rank_later_input_skipped)
		109	+{
		110	+ // A -2 (unknown rank) input is runtime-determined; its comparison is
		111	+ // skipped even when it is not the first input.
		112	+ auto context = MakeInferShapeContext({{}, {}, {-2}, {}}, ge::DT_FLOAT16);
		113	+ ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{}});
		114	+}
		115	+
		116	+TEST_F(AccumulateNV2InferShape, unknown_dim_does_not_mask_real_mismatch)
		117	+{
		118	+ // -1 is a wildcard, but two different concrete dims still fail.
		119	+ auto context = MakeInferShapeContext({{-1, 5}, {4, 5}, {3, 5}}, ge::DT_FLOAT);
		120	+ ExecuteTestCase(context, ge::GRAPH_FAILED);
		121	+}

Mmath/accumulate_nv2/tests/ut/op_kernel/CMakeLists.txt+4-5

	@@ -1,11 +1,10 @@
1	# ----------------------------------------------------------------------------	1	# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.
3	# Copyright (c) 2026 Huawei Technologies Co., Ltd.	2	# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4	-# This file is a part of the CANN Open Software.	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
6	# Please refer to the License for details. You may not use this file except in compliance with the License.	5	# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	# See LICENSE in the root of the software repository for the full text of the License.	8	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	9	# ----------------------------------------------------------------------------
11		10

Mmath/accumulate_nv2/tests/ut/op_kernel/test_accumulate_nv2.cpp+50-67

	@@ -262,9 +262,10 @@ TEST_F(AccumulateNV2KernelTest, test_float32_2inputs_double_buffer)
262	AscendC::GmFree(tiling);	262	AscendC::GmFree(tiling);
263	}	263	}
264		264
265	-// float32 uses compensated sequential accumulation so a small term is not	265	+// float32 uses TensorFlow-compatible sequential summation: 1e8 + 1 rounds back
266	-// lost permanently when later inputs cancel a much larger partial sum.	266	+// to 1e8 (ulp(1e8) = 8), the -1e8 term then cancels the partial sum to 0 and
267	-TEST_F(AccumulateNV2KernelTest, test_float32_compensated_cancellation)	267	+// only the trailing 3 survives.
		268	+TEST_F(AccumulateNV2KernelTest, test_float32_sequential_cancellation)
268	{	269	{
269	constexpr int32_t INPUT_NUM = 4;	270	constexpr int32_t INPUT_NUM = 4;
270	constexpr int64_t TOTAL_NUM = 32;	271	constexpr int64_t TOTAL_NUM = 32;
	@@ -289,11 +290,6 @@ TEST_F(AccumulateNV2KernelTest, test_float32_compensated_cancellation)
289	td->blockFactor = TOTAL_NUM;	290	td->blockFactor = TOTAL_NUM;
290	td->ubFactor = TOTAL_NUM;	291	td->ubFactor = TOTAL_NUM;
291	td->inputNum = INPUT_NUM;	292	td->inputNum = INPUT_NUM;
292	- td->rank = 1;
293	- td->outputShape[0] = TOTAL_NUM;
294	- for (int32_t i = 0; i < INPUT_NUM; ++i) {
295	- td->inputStrides[i][0] = 1;
296	- }
297		293
298	ICPU_SET_TILING_KEY(0UL);	294	ICPU_SET_TILING_KEY(0UL);
299	AscendC::SetKernelMode(KernelMode::AIV_MODE);	295	AscendC::SetKernelMode(KernelMode::AIV_MODE);
	@@ -301,7 +297,7 @@ TEST_F(AccumulateNV2KernelTest, test_float32_compensated_cancellation)
301		297
302	auto* result = reinterpret_cast<float*>(y);	298	auto* result = reinterpret_cast<float*>(y);
303	for (int64_t j = 0; j < TOTAL_NUM; ++j) {	299	for (int64_t j = 0; j < TOTAL_NUM; ++j) {
304	- EXPECT_FLOAT_EQ(result[j], 4.0f);	300	+ EXPECT_FLOAT_EQ(result[j], 3.0f);
305	}	301	}
306		302
307	builder.Free();	303	builder.Free();
	@@ -476,59 +472,8 @@ TEST_F(AccumulateNV2KernelTest, test_float32_4inputs_unaligned_tail)
476	AscendC::GmFree(tiling);	472	AscendC::GmFree(tiling);
477	}	473	}
478		474
479	-// float32 broadcast: [2, 1] + [1, 3] => [2, 3]	475	+// FP16 inputs are promoted to FP32 before accumulation and cast back once:
480	-TEST_F(AccumulateNV2KernelTest, test_float32_broadcast)	476	+// 2048 + 1 + 1 = 2050 survives the single rounding exactly.
481	-{
482	- constexpr int32_t INPUT_NUM = 2;
483	- constexpr int64_t TOTAL_NUM = 6;
484	-
485	- TensorListBuilder builder;
486	- float* input0 = (float)AscendC::GmAlloc(2 sizeof(float));
487	- input0[0] = 1.0f;
488	- input0[1] = 2.0f;
489	- builder.tensors.push_back((uint8_t*)input0);
490	- float* input1 = (float)AscendC::GmAlloc(3 sizeof(float));
491	- input1[0] = 10.0f;
492	- input1[1] = 20.0f;
493	- input1[2] = 30.0f;
494	- builder.tensors.push_back((uint8_t*)input1);
495	-
496	- uint8_t* x = builder.Build(INPUT_NUM);
497	- uint8_t* y = (uint8_t)AscendC::GmAlloc(TOTAL_NUM sizeof(float));
498	- memset(y, 0, TOTAL_NUM * sizeof(float));
499	- uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(1024);
500	- uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(AccumulateNV2TilingData));
501	- auto* td = reinterpret_cast<AccumulateNV2TilingData*>(tiling);
502	- td->totalNum = TOTAL_NUM;
503	- td->blockFactor = TOTAL_NUM;
504	- td->ubFactor = 32;
505	- td->inputNum = INPUT_NUM;
506	- td->rank = 2;
507	- td->needBroadcast = 1;
508	- td->outputShape[0] = 2;
509	- td->outputShape[1] = 3;
510	- td->inputStrides[0][0] = 1;
511	- td->inputStrides[0][1] = 0;
512	- td->inputStrides[1][0] = 0;
513	- td->inputStrides[1][1] = 1;
514	-
515	- ICPU_SET_TILING_KEY(0UL);
516	- AscendC::SetKernelMode(KernelMode::AIV_MODE);
517	- ICPU_RUN_KF(kernel_accumulate_nv2_float_single, 1, x, y, workspace, tiling);
518	-
519	- const float expected[TOTAL_NUM] = {11.0f, 21.0f, 31.0f, 12.0f, 22.0f, 32.0f};
520	- auto* result = reinterpret_cast<float*>(y);
521	- for (int64_t j = 0; j < TOTAL_NUM; ++j) {
522	- EXPECT_FLOAT_EQ(result[j], expected[j]);
523	- }
524	-
525	- builder.Free();
526	- AscendC::GmFree(y);
527	- AscendC::GmFree(workspace);
528	- AscendC::GmFree(tiling);
529	-}
530	-
531	-// FP16 inputs are converted to FP32 before accumulation and cast back once.
532	TEST_F(AccumulateNV2KernelTest, test_float16_promote_before_accumulate)	477	TEST_F(AccumulateNV2KernelTest, test_float16_promote_before_accumulate)
533	{	478	{
534	constexpr int32_t INPUT_NUM = 3;	479	constexpr int32_t INPUT_NUM = 3;
	@@ -553,11 +498,6 @@ TEST_F(AccumulateNV2KernelTest, test_float16_promote_before_accumulate)
553	td->blockFactor = TOTAL_NUM;	498	td->blockFactor = TOTAL_NUM;
554	td->ubFactor = 32;	499	td->ubFactor = 32;
555	td->inputNum = INPUT_NUM;	500	td->inputNum = INPUT_NUM;
556	- td->rank = 1;
557	- td->outputShape[0] = TOTAL_NUM;
558	- for (int32_t i = 0; i < INPUT_NUM; ++i) {
559	- td->inputStrides[i][0] = 1;
560	- }
561		501
562	ICPU_SET_TILING_KEY(0UL);	502	ICPU_SET_TILING_KEY(0UL);
563	AscendC::SetKernelMode(KernelMode::AIV_MODE);	503	AscendC::SetKernelMode(KernelMode::AIV_MODE);
	@@ -573,3 +513,46 @@ TEST_F(AccumulateNV2KernelTest, test_float16_promote_before_accumulate)
573	AscendC::GmFree(workspace);	513	AscendC::GmFree(workspace);
574	AscendC::GmFree(tiling);	514	AscendC::GmFree(tiling);
575	}	515	}
		516	+
		517	+// fp16, 4 inputs of 15 elements: verifies the Cast-based promote pipeline on a
		518	+// non-32-byte-aligned tail (currentNum not a multiple of the vector block).
		519	+TEST_F(AccumulateNV2KernelTest, test_float16_4inputs_unaligned_tail)
		520	+{
		521	+ constexpr int32_t INPUT_NUM = 4;
		522	+ constexpr int64_t TOTAL_NUM = 15;
		523	+ const float values[INPUT_NUM] = {1.5f, 2.25f, 3.0f, 4.0f};
		524	+
		525	+ TensorListBuilder builder;
		526	+ for (int32_t i = 0; i < INPUT_NUM; ++i) {
		527	+ half* tensor = (half)AscendC::GmAlloc(TOTAL_NUM sizeof(half));
		528	+ for (int64_t j = 0; j < TOTAL_NUM; ++j) {
		529	+ tensor[j] = static_cast<half>(values[i]);
		530	+ }
		531	+ builder.tensors.push_back((uint8_t*)tensor);
		532	+ }
		533	+ uint8_t* x = builder.Build(INPUT_NUM);
		534	+ uint8_t* y = (uint8_t)AscendC::GmAlloc(TOTAL_NUM sizeof(half));
		535	+ memset(y, 0, TOTAL_NUM * sizeof(half));
		536	+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(1024);
		537	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(AccumulateNV2TilingData));
		538	+ auto* td = reinterpret_cast<AccumulateNV2TilingData*>(tiling);
		539	+ memset(td, 0, sizeof(AccumulateNV2TilingData));
		540	+ td->totalNum = TOTAL_NUM;
		541	+ td->blockFactor = TOTAL_NUM;
		542	+ td->ubFactor = 8;
		543	+ td->inputNum = INPUT_NUM;
		544	+
		545	+ ICPU_SET_TILING_KEY(0UL);
		546	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
		547	+ ICPU_RUN_KF(kernel_accumulate_nv2_half_single, 1, x, y, workspace, tiling);
		548	+
		549	+ auto* result = reinterpret_cast<half*>(y);
		550	+ for (int64_t j = 0; j < TOTAL_NUM; ++j) {
		551	+ EXPECT_FLOAT_EQ(static_cast<float>(result[j]), 10.75f);
		552	+ }
		553	+
		554	+ builder.Free();
		555	+ AscendC::GmFree(y);
		556	+ AscendC::GmFree(workspace);
		557	+ AscendC::GmFree(tiling);
		558	+}