【社区任务】 Equal 算子 A2/A3 实现-ops-math-AtomGit

Mexperimental/math/equal/README.md+112-40

	@@ -1,44 +1,116 @@
1	# Equal	1	# Equal
2		2
		3	+## 产品支持情况
		4	+
		5	+\| 产品 \| 是否支持 \|
		6	+\| :--- \| :---: \|
		7	+\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|
		8	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|
		9	+
		10	+## 功能说明
		11	+
		12	+- 算子功能：对两个输入张量进行逐元素相等比较，并根据广播后的形状输出BOOL张量。
		13	+- 计算公式：
		14	+
		15	+$$
		16	+out_i = (self_i == other_i)
		17	+$$
		18	+
		19	+- 浮点比较采用数值语义：`+0.0`与`-0.0`相等，`NaN`与任意值（包括`NaN`）均不相等，同号无穷大相等。
		20	+
		21	+## 参数说明
		22	+
		23	+<table style="table-layout: fixed; width: 1494px"><colgroup>
		24	+ <col style="width: 146px">
		25	+ <col style="width: 110px">
		26	+ <col style="width: 301px">
		27	+ <col style="width: 219px">
		28	+ <col style="width: 328px">
		29	+ <col style="width: 101px">
		30	+ <col style="width: 143px">
		31	+ <col style="width: 146px">
		32	+ </colgroup>
		33	+ <thead>
		34	+ <tr>
		35	+ <th>参数名</th>
		36	+ <th>输入/输出</th>
		37	+ <th>描述</th>
		38	+ <th>使用说明</th>
		39	+ <th>数据类型</th>
		40	+ <th>数据格式</th>
		41	+ <th>维度(shape)</th>
		42	+ <th>非连续Tensor</th>
		43	+ </tr>
		44	+ </thead>
		45	+ <tbody>
		46	+ <tr>
		47	+ <td>self</td>
		48	+ <td>输入</td>
		49	+ <td>待比较的第一个输入张量。</td>
		50	+ <td>数据类型与other一致；shape与other满足broadcast关系。</td>
		51	+ <td>FLOAT16、FLOAT、BFLOAT16、INT8、UINT8、INT16、INT32、UINT32、INT64、BOOL</td>
		52	+ <td>ND</td>
		53	+ <td>0-8</td>
		54	+ <td>√</td>
		55	+ </tr>
		56	+ <tr>
		57	+ <td>other</td>
		58	+ <td>输入</td>
		59	+ <td>待比较的第二个输入张量。</td>
		60	+ <td>数据类型与self一致；shape与self满足broadcast关系。</td>
		61	+ <td>FLOAT16、FLOAT、BFLOAT16、INT8、UINT8、INT16、INT32、UINT32、INT64、BOOL</td>
		62	+ <td>ND</td>
		63	+ <td>0-8</td>
		64	+ <td>√</td>
		65	+ </tr>
		66	+ <tr>
		67	+ <td>out</td>
		68	+ <td>输出</td>
		69	+ <td>逐元素比较结果。</td>
		70	+ <td>shape为self与other广播后的公共shape。</td>
		71	+ <td>BOOL</td>
		72	+ <td>ND</td>
		73	+ <td>0-8</td>
		74	+ <td>√</td>
		75	+ </tr>
		76	+ <tr>
		77	+ <td>workspaceSize</td>
		78	+ <td>输出</td>
		79	+ <td>返回需要在Device侧申请的workspace大小。</td>
		80	+ <td>-</td>
		81	+ <td>-</td>
		82	+ <td>-</td>
		83	+ <td>-</td>
		84	+ <td>-</td>
		85	+ </tr>
		86	+ <tr>
		87	+ <td>executor</td>
		88	+ <td>输出</td>
		89	+ <td>返回op执行器，包含算子计算流程。</td>
		90	+ <td>-</td>
		91	+ <td>-</td>
		92	+ <td>-</td>
		93	+ <td>-</td>
		94	+ <td>-</td>
		95	+ </tr>
		96	+ </tbody>
		97	+</table>
		98	+
		99	+## 约束说明
		100	+
		101	+- self与other的数据类型必须一致。
		102	+- self与other的shape必须满足broadcast关系。
		103	+- out的数据类型必须为BOOL，shape必须与输入广播后的公共shape一致。
		104	+
		105	+## 调用说明
		106	+
		107	+\| 调用方式 \| 调用样例 \| 说明 \|
		108	+\| :--- \| :--- \| :--- \|
		109	+\| aclnn调用 \| [test_aclnn_eq_tensor.cpp](./examples/test_aclnn_eq_tensor.cpp) \| 通过[aclnnEqTensor/aclnnInplaceEqTensor](./docs/aclnnEqTensor&aclnnInplaceEqTensor.md)接口调用Equal算子。 \|
		110	+\| aclnn调用 \| [test_aclnn_eq_scalar.cpp](./examples/test_aclnn_eq_scalar.cpp) \| 通过[aclnnEqScalar/aclnnInplaceEqScalar](./docs/aclnnEqScalar&aclnnInplaceEqScalar.md)接口调用Equal算子。 \|
		111	+
3	## 贡献说明	112	## 贡献说明
4		113
5	-\| 贡献者 \| 贡献方 \| 贡献算子 \| 贡献时间 \| 贡献内容 \|	114	+\| 贡献者 \| 贡献方 \| 贡献算子 \| 贡献时间 \| 贡献内容 \|
6	-\|-------\|------------------\|-------\|-----------\|-----------\|	115	+\| :--- \| :--- \| :--- \| :--- \| :--- \|
7	-\| skywang2 \| 个人开发者 \| Equal \| 2026/6/21 \| 新增Equal算子 \|	116	+\| gcw_rESKWmgp \| 个人开发者 \| Equal \| 2026/08/27 \| 完成A2/A3 Ascend C实现、数据类型适配、广播支持及性能优化。 \|
8	-
9	-## 支持的产品型号
10	-
11	-- Atlas A2 训练系列产品
12	-
13	-产品形态详细说明请参见[昇腾产品形态说明](http://www.hiascend.com/document/redirect/CannCommunityProductForm)
14	-
15	-## 算子描述
16	-
17	-- 功能描述
18	-
19	- `Equal`算子将输入的两个向量数据进行各对应位置的判等运算，返回结果向量。
20	-
21	-- 原型信息
22	-
23	- <table>
24	- <tr><th align="center">算子类型(OpType)</th><th colspan="4" align="center">Equal</th></tr>
25	- <tr><td align="center"> </td><td align="center">name</td><td align="center">Type</td><td align="center">data type</td><td align="center">format</td></tr>
26	- <tr><td rowspan="2" align="center">算子输入</td><td align="center">x1</td><td align="center">tensor</td><td align="center">float32,float16,bfloat16,uint8,int8,uint32,int32</td><td align="center">ND</td></tr>
27	- <tr><td align="center">x2</td><td align="center">tensor</td><td align="center">float32,float16,bfloat16,uint8,int8,uint32,int32</td><td align="center">ND</td></tr>
28	- <tr><td align="center">算子输出</td><td align="center">y</td><td align="center">tensor</td><td align="center">bool</td><td align="center">ND</td></tr>
29	- <tr><td align="center">核函数名</td><td colspan="4" align="center">equal</td></tr>
30	- </table>
31	-
32	-## 约束与限制
33	-
34	-- x,y,out的数据类型仅支持float32,float16,bfloat16,uint8,int8,uint32,int32，数据格式仅支持ND
35	-
36	-### 运行验证
37	-
38	-测试命令调用方式：[build.sh](../../../docs/zh/invocation/quick_op_invocation.md)
39	-<table>
40	- <tr><th>目录</th><th>描述</th></tr>
41	- <tr>
42	- <td><a href="./examples/test_aclnn_equal.cpp">test_aclnn_equal.cpp</a></td><td>通过aclnn调用的方式调用Equal算子。</td>
43	- </tr>
44	-</table>

Mexperimental/math/equal/docs/aclnnEqScalar&aclnnInplaceEqScalar.md+4-3

	@@ -5,6 +5,7 @@
5	\| 产品 \| 是否支持 \|	5	\| 产品 \| 是否支持 \|
6	\| :----------------------------------------------------------- \| :------: \|	6	\| :----------------------------------------------------------- \| :------: \|
7	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|	7	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|
		8	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|
8		9
9	## 功能说明	10	## 功能说明
10		11
	@@ -79,7 +80,7 @@ aclnnStatus aclnnInplaceEqScalar(
79	<th>使用说明</th>	80	<th>使用说明</th>
80	<th>数据类型</th>	81	<th>数据类型</th>
81	<th>数据格式</th>	82	<th>数据格式</th>
82	- <th>维度(shape)</th>	83	+ <th>维度（shape）</th>
83	<th>非连续Tensor</th>	84	<th>非连续Tensor</th>
84	</tr></thead>	85	</tr></thead>
85	<tbody>	86	<tbody>
	@@ -141,7 +142,7 @@ aclnnStatus aclnnInplaceEqScalar(
141	- self、other不支持UINT32数据类型。	142	- self、other不支持UINT32数据类型。
142	- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：	143	- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：
143	- self与other满足[互推导关系](../../../../docs/zh/context/deduction_relationship.md)。	144	- self与other满足[互推导关系](../../../../docs/zh/context/deduction_relationship.md)。
144	- - self、other不支持UINT32、UINT64数据类型	145	+ - self、other不支持UINT32、UINT64数据类型。
145	- out不支持UINT64、UINT32、UINT16数据类型。	146	- out不支持UINT64、UINT32、UINT16数据类型。
146	- <term>Atlas 训练系列产品</term>：	147	- <term>Atlas 训练系列产品</term>：
147	- self与other满足[互推导关系](../../../../docs/zh/context/deduction_relationship.md)。	148	- self与other满足[互推导关系](../../../../docs/zh/context/deduction_relationship.md)。
	@@ -253,7 +254,7 @@ aclnnStatus aclnnInplaceEqScalar(
253	<th>使用说明</th>	254	<th>使用说明</th>
254	<th>数据类型</th>	255	<th>数据类型</th>
255	<th>数据格式</th>	256	<th>数据格式</th>
256	- <th>维度(shape)</th>	257	+ <th>维度（shape）</th>
257	<th>非连续Tensor</th>	258	<th>非连续Tensor</th>
258	</tr></thead>	259	</tr></thead>
259	<tbody>	260	<tbody>

Mexperimental/math/equal/docs/aclnnEqTensor&aclnnInplaceEqTensor.md+8-7

	@@ -5,10 +5,11 @@
5	\| 产品 \| 是否支持 \|	5	\| 产品 \| 是否支持 \|
6	\| :----------------------------------------------------------- \| :------: \|	6	\| :----------------------------------------------------------- \| :------: \|
7	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|	7	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|
		8	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|
8		9
9	## 功能说明	10	## 功能说明
10		11
11	-- 接口功能：计算两个Tensor中的元素是否相等，返回一个Tensor，self=other的为True(1.)，否则为False(0.)。	12	+- 接口功能：计算两个Tensor中的元素是否相等，返回一个Tensor，self=other的为True（1），否则为False（0）。
12	- 计算表达式：	13	- 计算表达式：
13		14
14	$$	15	$$
	@@ -79,7 +80,7 @@ aclnnStatus aclnnInplaceEqTensor(
79	<th>使用说明</th>	80	<th>使用说明</th>
80	<th>数据类型</th>	81	<th>数据类型</th>
81	<th>数据格式</th>	82	<th>数据格式</th>
82	- <th>维度(shape)</th>	83	+ <th>维度（shape）</th>
83	<th>非连续Tensor</th>	84	<th>非连续Tensor</th>
84	</tr></thead>	85	</tr></thead>
85	<tbody>	86	<tbody>
	@@ -147,7 +148,7 @@ aclnnStatus aclnnInplaceEqTensor(
147		148
148	- 返回值：	149	- 返回值：
149		150
150	- aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn_return_code.md)。	151	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn_return_code.md)。
151		152
152	第一段接口完成入参校验，出现以下场景时报错：	153	第一段接口完成入参校验，出现以下场景时报错：
153		154
	@@ -229,7 +230,7 @@ aclnnStatus aclnnInplaceEqTensor(
229		230
230	- 返回值：	231	- 返回值：
231		232
232	- aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn_return_code.md)。	233	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn_return_code.md)。
233		234
234	## aclnnInplaceEqTensorGetWorkspaceSize	235	## aclnnInplaceEqTensorGetWorkspaceSize
235		236
	@@ -253,7 +254,7 @@ aclnnStatus aclnnInplaceEqTensor(
253	<th>使用说明</th>	254	<th>使用说明</th>
254	<th>数据类型</th>	255	<th>数据类型</th>
255	<th>数据格式</th>	256	<th>数据格式</th>
256	- <th>维度(shape)</th>	257	+ <th>维度（shape）</th>
257	<th>非连续Tensor</th>	258	<th>非连续Tensor</th>
258	</tr></thead>	259	</tr></thead>
259	<tbody>	260	<tbody>
	@@ -309,7 +310,7 @@ aclnnStatus aclnnInplaceEqTensor(
309		310
310	- 返回值：	311	- 返回值：
311		312
312	- aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn_return_code.md)。	313	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn_return_code.md)。
313		314
314	第一段接口完成入参校验，出现以下场景时报错：	315	第一段接口完成入参校验，出现以下场景时报错：
315		316
	@@ -394,7 +395,7 @@ aclnnStatus aclnnInplaceEqTensor(
394		395
395	- 返回值：	396	- 返回值：
396		397
397	- aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn_return_code.md)。	398	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn_return_code.md)。
398		399
399	## 约束说明	400	## 约束说明
400		401

Mexperimental/math/equal/examples/test_aclnn_equal.cpp+33-31

	@@ -1,12 +1,12 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	#include <iostream>	11	#include <iostream>
12	#include <vector>	12	#include <vector>
	@@ -27,11 +27,11 @@
27		27
28	int64_t GetShapeSize(const std::vector<int64_t>& shape)	28	int64_t GetShapeSize(const std::vector<int64_t>& shape)
29	{	29	{
30	- int64_t shape_size = 1;	30	+ int64_t shapeSize = 1;
31	for (auto i : shape) {	31	for (auto i : shape) {
32	- shape_size *= i;	32	+ shapeSize *= i;
33	}	33	}
34	- return shape_size;	34	+ return shapeSize;
35	}	35	}
36		36
37	int Init(int32_t deviceId, aclrtStream* stream)	37	int Init(int32_t deviceId, aclrtStream* stream)
	@@ -47,9 +47,8 @@ int Init(int32_t deviceId, aclrtStream* stream)
47	}	47	}
48		48
49	template <typename T>	49	template <typename T>
50	-int CreateAclTensor(	50	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
51	- const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,	51	+ aclDataType dataType, aclTensor** tensor)
52	- aclTensor** tensor)
53	{	52	{
54	auto size = GetShapeSize(shape) * sizeof(T);	53	auto size = GetShapeSize(shape) * sizeof(T);
55	// 调用aclrtMalloc申请device侧内存	54	// 调用aclrtMalloc申请device侧内存
	@@ -67,9 +66,14 @@ int CreateAclTensor(
67	}	66	}
68		67
69	// 调用aclCreateTensor接口创建aclTensor	68	// 调用aclCreateTensor接口创建aclTensor
70	- *tensor = aclCreateTensor(	69	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
71	- shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),	70	+ shape.data(), shape.size(), *deviceAddr);
72	- *deviceAddr);	71	+ if (*tensor == nullptr) {
		72	+ LOG_PRINT("aclCreateTensor failed.\n");
		73	+ aclrtFree(*deviceAddr);
		74	+ *deviceAddr = nullptr;
		75	+ return ACL_ERROR_FAILURE;
		76	+ }
73	return 0;	77	return 0;
74	}	78	}
75		79
	@@ -80,10 +84,9 @@ aclError InitAcl(int32_t deviceId, aclrtStream* stream)
80	return ACL_SUCCESS;	84	return ACL_SUCCESS;
81	}	85	}
82		86
83	-aclError CreateInputs(	87	+aclError CreateInputs(const std::vector<int64_t>& selfShape, const std::vector<int64_t>& otherShape,
84	- std::vector<int64_t>& selfShape, std::vector<int64_t>& otherShape, std::vector<int64_t>& outShape,	88	+ const std::vector<int64_t>& outShape, void selfDeviceAddr, void otherDeviceAddr,
85	- void selfDeviceAddr, void otherDeviceAddr, void outDeviceAddr, aclTensor self, aclTensor** other,	89	+ void outDeviceAddr, aclTensor self, aclTensor other, aclTensor out)
86	- aclTensor** out)
87	{	90	{
88	std::vector<double> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};	91	std::vector<double> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
89	std::vector<double> otherHostData = {0, 1, 9, 9, 9, 9, 9, 9};	92	std::vector<double> otherHostData = {0, 1, 9, 9, 9, 9, 9, 9};
	@@ -101,9 +104,9 @@ aclError CreateInputs(
101	return ACL_SUCCESS;	104	return ACL_SUCCESS;
102	}	105	}
103		106
104	-aclError ExecOpApi(	107	+aclError ExecOpApi(const aclTensor* self, const aclTensor* other, aclTensor* out, const void* outDeviceAddr,
105	- aclTensor* self, aclTensor* other, aclTensor* out, void** workspaceAddrOut, uint64_t& workspaceSize,	108	+ const std::vector<int64_t>& outShape, aclrtStream stream, void** workspaceAddrOut,
106	- void* outDeviceAddr, std::vector<int64_t>& outShape, aclrtStream stream)	109	+ uint64_t& workspaceSize)
107	{	110	{
108	aclOpExecutor* executor;	111	aclOpExecutor* executor;
109		112
	@@ -126,9 +129,8 @@ aclError ExecOpApi(
126	auto size = GetShapeSize(outShape);	129	auto size = GetShapeSize(outShape);
127	std::vector<char> resultData(size, 0);	130	std::vector<char> resultData(size, 0);
128		131
129	- ret = aclrtMemcpy(	132	+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(char),
130	- resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(char),	133	+ ACL_MEMCPY_DEVICE_TO_HOST);
131	- ACL_MEMCPY_DEVICE_TO_HOST);
132	CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);	134	CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
133		135
134	for (int64_t i = 0; i < size; i++) {	136	for (int64_t i = 0; i < size; i++) {
	@@ -158,14 +160,14 @@ int main()
158	aclTensor* other = nullptr;	160	aclTensor* other = nullptr;
159	aclTensor* out = nullptr;	161	aclTensor* out = nullptr;
160		162
161	- ret = CreateInputs(	163	+ ret = CreateInputs(selfShape, otherShape, outShape, &selfDeviceAddr, &otherDeviceAddr, &outDeviceAddr, &self,
162	- selfShape, otherShape, outShape, &selfDeviceAddr, &otherDeviceAddr, &outDeviceAddr, &self, &other, &out);	164	+ &other, &out);
163	CHECK_RET(ret == ACL_SUCCESS, return ret);	165	CHECK_RET(ret == ACL_SUCCESS, return ret);
164		166
165	uint64_t workspaceSize = 0;	167	uint64_t workspaceSize = 0;
166	void* workspaceAddr = nullptr;	168	void* workspaceAddr = nullptr;
167		169
168	- ret = ExecOpApi(self, other, out, &workspaceAddr, workspaceSize, outDeviceAddr, outShape, stream);	170	+ ret = ExecOpApi(self, other, out, outDeviceAddr, outShape, stream, &workspaceAddr, workspaceSize);
169	CHECK_RET(ret == ACL_SUCCESS, return ret);	171	CHECK_RET(ret == ACL_SUCCESS, return ret);
170		172
171	// 释放资源	173	// 释放资源
	@@ -185,4 +187,4 @@ int main()
185	aclrtResetDevice(deviceId);	187	aclrtResetDevice(deviceId);
186	aclFinalize();	188	aclFinalize();
187	return 0;	189	return 0;
188	-}	190	+}

Mexperimental/math/equal/op_api/aclnn_eq_scalar.cpp+19-44

	@@ -1,12 +1,12 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10	#include "aclnn_eq_scalar.h"	10	#include "aclnn_eq_scalar.h"
11	#include "equal.h"	11	#include "equal.h"
12	#include "aclnn_kernels/contiguous.h"	12	#include "aclnn_kernels/contiguous.h"
	@@ -32,31 +32,6 @@ static const std::initializer_list<op::DataType> DTYPE_SUPPORT_910B_LIST = {
32	op::DataType::DT_INT16, op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_BOOL,	32	op::DataType::DT_INT16, op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_BOOL,
33	op::DataType::DT_DOUBLE, op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128, op::DataType::DT_BF16};	33	op::DataType::DT_DOUBLE, op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128, op::DataType::DT_BF16};
34		34
35	-static inline double GetCastedDouble(const aclTensor* self, const aclScalar* other)
36	-{
37	- double res = 0;
38	- switch (self->GetDataType()) {
39	- case DataType::DT_FLOAT:
40	- res = static_cast<double>(other->ToFloat());
41	- break;
42	- case DataType::DT_FLOAT16:
43	- res = static_cast<double>(other->ToFp16());
44	- break;
45	- default:
46	- res = other->ToDouble();
47	- break;
48	- }
49	- return res;
50	-}
51	-
52	-static inline bool IsDoubleEqual(double a, double b)
53	-{
54	- if (std::abs(a - b) <= std::numeric_limits<float>::epsilon()) {
55	- return true;
56	- }
57	- return false;
58	-}
59	-
60	static op::DataType PromoteTypeScalar(op::DataType selfDtype, op::DataType otherDtype)	35	static op::DataType PromoteTypeScalar(op::DataType selfDtype, op::DataType otherDtype)
61	{	36	{
62	if (IsComplexType(selfDtype)) {	37	if (IsComplexType(selfDtype)) {
	@@ -111,9 +86,8 @@ static bool CheckPromoteType(const aclTensor* self, const aclScalar* other, cons
111	op::DataType promoteType;	86	op::DataType promoteType;
112	promoteType = op::PromoteType(self->GetDataType(), other->GetDataType());	87	promoteType = op::PromoteType(self->GetDataType(), other->GetDataType());
113	if (promoteType == DataType::DT_UNDEFINED) {	88	if (promoteType == DataType::DT_UNDEFINED) {
114	- OP_LOGE(	89	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self dtype %s and other dtype %s cannot be promoted to a common dtype.",
115	- ACLNN_ERR_PARAM_INVALID, "Self dtype %s and other dtype %s can not promote dtype.",	90	+ op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
116	- op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
117	return false;	91	return false;
118	}	92	}
119		93
	@@ -140,24 +114,25 @@ static bool CheckShape(const aclTensor* self, const aclTensor* out)
140		114
141	static aclnnStatus CheckParams(const aclTensor* self, const aclScalar* other, const aclTensor* out)	115	static aclnnStatus CheckParams(const aclTensor* self, const aclScalar* other, const aclTensor* out)
142	{	116	{
143	- // 1. 检查参数是否为空指针	117	+ // 检查参数是否为空指针
144	CHECK_RET(CheckNotNull(self, other, out), ACLNN_ERR_PARAM_NULLPTR);	118	CHECK_RET(CheckNotNull(self, other, out), ACLNN_ERR_PARAM_NULLPTR);
145		119
146	- // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验	120	+ // 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
147	CHECK_RET(CheckDtypeValid(self, other, out), ACLNN_ERR_PARAM_INVALID);	121	CHECK_RET(CheckDtypeValid(self, other, out), ACLNN_ERR_PARAM_INVALID);
148		122
149	- // 3. 检查self的数据类型能否转换为输出数据类型	123	+ // 检查self的数据类型能否转换为输出数据类型
150	CHECK_RET(CheckPromoteType(self, other, out), ACLNN_ERR_PARAM_INVALID);	124	CHECK_RET(CheckPromoteType(self, other, out), ACLNN_ERR_PARAM_INVALID);
151		125
152	- // 4. 检查shape是否满足约束	126	+ // 检查shape是否满足约束
153	CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID);	127	CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID);
154		128
155	return ACLNN_SUCCESS;	129	return ACLNN_SUCCESS;
156	}	130	}
157		131
158	-aclnnStatus aclnnEqScalarGetWorkspaceSize(	132	+aclnnStatus aclnnEqScalarGetWorkspaceSize(const aclTensor* self, const aclScalar* other, aclTensor* out,
159	- const aclTensor* self, const aclScalar* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)	133	+ uint64_t* workspaceSize, aclOpExecutor** executor)
160	{	134	{
		135	+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
161	L2_DFX_PHASE_1(aclnnEqScalar, DFX_IN(self, other), DFX_OUT(out));	136	L2_DFX_PHASE_1(aclnnEqScalar, DFX_IN(self, other), DFX_OUT(out));
162	// 固定写法，创建OpExecutor	137	// 固定写法，创建OpExecutor
163	auto uniqueExecutor = CREATE_EXECUTOR();	138	auto uniqueExecutor = CREATE_EXECUTOR();
	@@ -204,8 +179,8 @@ aclnnStatus aclnnEqScalarGetWorkspaceSize(
204	return ACLNN_SUCCESS;	179	return ACLNN_SUCCESS;
205	}	180	}
206		181
207	-aclnnStatus aclnnInplaceEqScalarGetWorkspaceSize(	182	+aclnnStatus aclnnInplaceEqScalarGetWorkspaceSize(const aclTensor* selfRef, const aclScalar* other,
208	- const aclTensor* selfRef, const aclScalar* other, uint64_t* workspaceSize, aclOpExecutor** executor)	183	+ uint64_t* workspaceSize, aclOpExecutor** executor)
209	{	184	{
210	auto out = const_cast<aclTensor*>(selfRef);	185	auto out = const_cast<aclTensor*>(selfRef);
211	return aclnnEqScalarGetWorkspaceSize(selfRef, other, out, workspaceSize, executor);	186	return aclnnEqScalarGetWorkspaceSize(selfRef, other, out, workspaceSize, executor);

Mexperimental/math/equal/op_api/aclnn_eq_scalar.h+40-32

	@@ -1,12 +1,12 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	#ifndef OP_API_INC_LEVEL2_ACLNN_EQ_SCALAR_H_	11	#ifndef OP_API_INC_LEVEL2_ACLNN_EQ_SCALAR_H_
12	#define OP_API_INC_LEVEL2_ACLNN_EQ_SCALAR_H_	12	#define OP_API_INC_LEVEL2_ACLNN_EQ_SCALAR_H_
	@@ -21,7 +21,8 @@ extern "C" {
21	/**	21	/**
22	* @brief aclnnEqScalar的第一段接口，根据具体的计算流程，计算workspace大小。	22	* @brief aclnnEqScalar的第一段接口，根据具体的计算流程，计算workspace大小。
23	* @domain aclnn_math	23	* @domain aclnn_math
24	- * 计算self中的元素的值与other的值是否相等，将self每个元素与other的值的比较结果写入out中：	24	+ * 计算self中的元素的值与other的值是否相等，
		25	+ * 将self每个元素与other的值的比较结果写入out中：
25	*	26	*
26	* $$ out_i = (self_i == \mathit{other} ) ? [True] : [False] $$	27	* $$ out_i = (self_i == \mathit{other} ) ? [True] : [False] $$
27	*	28	*
	@@ -35,66 +36,73 @@ extern "C" {
35	* D -->F1([l0op::ViewCopy])--> J[(out)]	36	* D -->F1([l0op::ViewCopy])--> J[(out)]
36	* ```	37	* ```
37	* @param [in] self: npu device侧的aclTensor，	38	* @param [in] self: npu device侧的aclTensor，
38	- * 数据类型支持FLOAT16, FLOAT, INT64, UINT64, INT32, INT8, UINT8, BOOL, UINT32, BFLOAT16, DOUBLE, INT16, COMPLEX64,	39	+ * 数据类型支持FLOAT16, FLOAT, INT64, UINT64, INT32, INT8, UINT8, BOOL, UINT32, BFLOAT16, DOUBLE,
		40	+ * INT16, COMPLEX64,
39	* COMPLEX128，数据格式支持ND，支持非连续的Tensor。	41	* COMPLEX128，数据格式支持ND，支持非连续的Tensor。
40	- * @param [in] other: host侧的aclScalar，数据类型支持FLOAT16, FLOAT, INT64, UINT64, INT32, INT8, UINT8, BOOL, UINT32,	42	+ * @param [in] other: host侧的aclScalar，
		43	+ * 数据类型支持FLOAT16, FLOAT, INT64, UINT64, INT32, INT8, UINT8, BOOL, UINT32,
41	* BFLOAT16, DOUBLE, INT16, COMPLEX64, COMPLEX128	44	* BFLOAT16, DOUBLE, INT16, COMPLEX64, COMPLEX128
42	* @param [in] out: npu device侧的aclTensor，数据类型为bool，shape与self的shape一致，数据格式支持ND	45	* @param [in] out: npu device侧的aclTensor，数据类型为bool，shape与self的shape一致，数据格式支持ND
43	- * @param [out] workspace_size: 返回用户需要在npu device侧申请的workspace大小。	46	+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
44	* @param [out] executor: 返回op执行器，包含算子计算流程。	47	* @param [out] executor: 返回op执行器，包含算子计算流程。
45	* @return aclnnStatus: 返回状态码。	48	* @return aclnnStatus: 返回状态码。
46	*/	49	*/
47	-ACLNN_API aclnnStatus aclnnEqScalarGetWorkspaceSize(	50	+ACLNN_API aclnnStatus aclnnEqScalarGetWorkspaceSize(const aclTensor* self, const aclScalar* other, aclTensor* out,
48	- const aclTensor* self, const aclScalar* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);	51	+ uint64_t* workspaceSize, aclOpExecutor** executor);
49		52
50	/**	53	/**
51	* @brief aclnnEqScalar的第二段接口，用于执行计算。	54	* @brief aclnnEqScalar的第二段接口，用于执行计算。
52	* @param [in] workspace: 在npu device侧申请的workspace内存起址。	55	* @param [in] workspace: 在npu device侧申请的workspace内存起址。
53	- * @param [in] workspace_size: 在npu device侧申请的workspace大小，由第一段接口aclnnEqScalarGetWorkspaceSize获取。	56	+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，
		57	+ * 由第一段接口aclnnEqScalarGetWorkspaceSize获取。
54	* @param [in] stream: acl stream流。	58	* @param [in] stream: acl stream流。
55	* @param [in] executor: op执行器，包含了算子计算流程。	59	* @param [in] executor: op执行器，包含了算子计算流程。
56	* @return aclnnStatus: 返回状态码。	60	* @return aclnnStatus: 返回状态码。
57	*/	61	*/
58	-ACLNN_API aclnnStatus	62	+ACLNN_API aclnnStatus aclnnEqScalar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
59	-aclnnEqScalar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream);	63	+ const aclrtStream stream);
60		64
61	/**	65	/**
62	* @brief aclnnInplaceEqScalar的第一段接口，根据具体的计算流程，计算workspace大小。	66	* @brief aclnnInplaceEqScalar的第一段接口，根据具体的计算流程，计算workspace大小。
63	* @domain aclnn_math	67	* @domain aclnn_math
64	- * 计算self中的元素的值与other的值是否相等，将self每个元素与other的值的比较结果原地返回到self中：	68	+ * 计算self中的元素的值与other的值是否相等，
		69	+ * 将self每个元素与other的值的比较结果原地返回到self中：
65	*	70	*
66	* $$ self_i = (self_i == \mathit{other} ) ? [True] : [False] $$	71	* $$ self_i = (self_i == \mathit{other} ) ? [True] : [False] $$
67	*	72	*
68	* 计算图：	73	* 计算图：
69	-*```mermaid	74	+ * ```mermaid
70	-* graph LR	75	+ * graph LR
71	-* A[(SelfRef)] -->B([l0op::Contiguous])	76	+ * A[(SelfRef)] -->B([l0op::Contiguous])
72	-* B -->C1([l0op::Cast])-->D([l0op::Equal])	77	+ * B -->C1([l0op::Cast])-->D([l0op::Equal])
73	-* F((other)) -->D	78	+ * F((other)) -->D
74	-* D -->F1([l0op::ViewCopy])--> J[(selfRef)]	79	+ * D -->F1([l0op::ViewCopy])--> J[(selfRef)]
75	-```	80	+ * ```
76	* @param [in] selfRef: npu device侧的aclTensor，	81	* @param [in] selfRef: npu device侧的aclTensor，
77	- * 数据类型支持FLOAT16, FLOAT, INT64, UINT64, INT32, INT8, UINT8, BOOL, UINT32, BFLOAT16, DOUBLE, INT16, COMPLEX64,	82	+ * 数据类型支持FLOAT16, FLOAT, INT64, UINT64, INT32, INT8, UINT8, BOOL, UINT32, BFLOAT16, DOUBLE,
		83	+ * INT16, COMPLEX64,
78	* COMPLEX128，数据格式支持ND，支持非连续的Tensor。	84	* COMPLEX128，数据格式支持ND，支持非连续的Tensor。
79	- * @param [in] other: host侧的aclScalar，数据类型支持FLOAT16, FLOAT, INT64, UINT64, INT32, INT8, UINT8, BOOL, UINT32,	85	+ * @param [in] other: host侧的aclScalar，
		86	+ * 数据类型支持FLOAT16, FLOAT, INT64, UINT64, INT32, INT8, UINT8, BOOL, UINT32,
80	* BFLOAT16, DOUBLE, INT16, COMPLEX64, COMPLEX128	87	* BFLOAT16, DOUBLE, INT16, COMPLEX64, COMPLEX128
81	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。	88	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
82	* @param [out] executor: 返回op执行器，包含算子计算流程。	89	* @param [out] executor: 返回op执行器，包含算子计算流程。
83	* @return aclnnStatus: 返回状态码。	90	* @return aclnnStatus: 返回状态码。
84	*/	91	*/
85	-ACLNN_API aclnnStatus aclnnInplaceEqScalarGetWorkspaceSize(	92	+ACLNN_API aclnnStatus aclnnInplaceEqScalarGetWorkspaceSize(const aclTensor* selfRef, const aclScalar* other,
86	- const aclTensor* selfRef, const aclScalar* other, uint64_t* workspaceSize, aclOpExecutor** executor);	93	+ uint64_t* workspaceSize, aclOpExecutor** executor);
87		94
88	/**	95	/**
89	* @brief aclnnInplaceEqScalar的第二段接口，用于执行计算。	96	* @brief aclnnInplaceEqScalar的第二段接口，用于执行计算。
90	* @param [in] workspace: 在npu device侧申请的workspace内存起址。	97	* @param [in] workspace: 在npu device侧申请的workspace内存起址。
91	- * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnEqScalarGetWorkspaceSize获取。	98	+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，
		99	+ * 由第一段接口aclnnEqScalarGetWorkspaceSize获取。
92	* @param [in] stream: acl stream流。	100	* @param [in] stream: acl stream流。
93	* @param [in] executor: op执行器，包含了算子计算流程。	101	* @param [in] executor: op执行器，包含了算子计算流程。
94	* @return aclnnStatus: 返回状态码。	102	* @return aclnnStatus: 返回状态码。
95	*/	103	*/
96	-ACLNN_API aclnnStatus	104	+ACLNN_API aclnnStatus aclnnInplaceEqScalar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
97	-aclnnInplaceEqScalar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);	105	+ aclrtStream stream);
98		106
99	#ifdef __cplusplus	107	#ifdef __cplusplus
100	}	108	}

Mexperimental/math/equal/op_api/aclnn_eq_tensor.cpp+40-57

	@@ -1,12 +1,12 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10	#include "aclnn_eq_tensor.h"	10	#include "aclnn_eq_tensor.h"
11	#include "equal.h"	11	#include "equal.h"
12	#include "aclnn_kernels/cast.h"	12	#include "aclnn_kernels/cast.h"
	@@ -14,16 +14,17 @@
14	#include "aclnn/aclnn_base.h"	14	#include "aclnn/aclnn_base.h"
15	#include "opdev/common_types.h"	15	#include "opdev/common_types.h"
16	#include "opdev/data_type_utils.h"	16	#include "opdev/data_type_utils.h"
		17	+#include "opdev/make_op_executor.h"
17	#include "opdev/shape_utils.h"	18	#include "opdev/shape_utils.h"
18	#include "opdev/format_utils.h"	19	#include "opdev/format_utils.h"
19	#include "opdev/op_dfx.h"	20	#include "opdev/op_dfx.h"
20	#include "opdev/op_executor.h"	21	#include "opdev/op_executor.h"
21	#include "opdev/op_log.h"	22	#include "opdev/op_log.h"
22	#include "opdev/tensor_view_utils.h"	23	#include "opdev/tensor_view_utils.h"
23	-#include "opdev/platform.h"
24	#include "aclnn_kernels/common/op_error_check.h"	24	#include "aclnn_kernels/common/op_error_check.h"
25		25
26	using namespace op;	26	using namespace op;
		27	+
27	#ifdef __cplusplus	28	#ifdef __cplusplus
28	extern "C" {	29	extern "C" {
29	#endif	30	#endif
	@@ -36,22 +37,7 @@ static bool CheckNotNull(const aclTensor* self, const aclTensor* other, const ac
36	return true;	37	return true;
37	}	38	}
38		39
39	-/* Equal 算子的完整计算流程如下:	40	+// Equal 由自研 AICore kernel 同时处理连续输入和广播输入。
40	- * self other
41	- * \| \|
42	- * \ /
43	- * Contiguous(workspace_0) Contiguous(workspace_2)
44	- * \ /
45	- * Cast(workspace_1) Cast(workspace_3)
46	- * \ /
47	- * Equal(workspace_4)
48	- * \|
49	- * Cast(workspace_5)
50	- * \|
51	- * ViewCopy
52	- * \|
53	- * result
54	- */
55		41
56	// 根据API定义，需要列出所能支持的所有dtype (1971)	42	// 根据API定义，需要列出所能支持的所有dtype (1971)
57	static const std::initializer_list<op::DataType> DTYPE_SUPPORT_910B_LIST = {	43	static const std::initializer_list<op::DataType> DTYPE_SUPPORT_910B_LIST = {
	@@ -65,26 +51,27 @@ static const size_t DIM_SUPPORT_MAX = 8;
65	static bool CheckDtypeValid(const aclTensor* self, const aclTensor* other, const aclTensor* out)	51	static bool CheckDtypeValid(const aclTensor* self, const aclTensor* other, const aclTensor* out)
66	{	52	{
67	const std::initializer_list<op::DataType> inputSupportList = DTYPE_SUPPORT_910B_LIST;	53	const std::initializer_list<op::DataType> inputSupportList = DTYPE_SUPPORT_910B_LIST;
		54	+ // Check each input before type promotion so unsupported dtypes cannot enter a supported promotion path.
		55	+ OP_CHECK_DTYPE_NOT_SUPPORT(self, inputSupportList, return false);
		56	+ OP_CHECK_DTYPE_NOT_SUPPORT(other, inputSupportList, return false);
		57	+
68	// 检查out的数据类型是否在equal算子的支持列表内	58	// 检查out的数据类型是否在equal算子的支持列表内
69	OP_CHECK_DTYPE_NOT_SUPPORT(out, inputSupportList, return false);	59	OP_CHECK_DTYPE_NOT_SUPPORT(out, inputSupportList, return false);
70		60
71	// 检查self和other能否做数据类型推导	61	// 检查self和other能否做数据类型推导
72	op::DataType promoteType = op::PromoteType(self->GetDataType(), other->GetDataType());	62	op::DataType promoteType = op::PromoteType(self->GetDataType(), other->GetDataType());
73	if (promoteType == DataType::DT_UNDEFINED) {	63	if (promoteType == DataType::DT_UNDEFINED) {
74	- OP_LOGE(	64	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Self dtype %s and other dtype %s cannot be promoted to a common dtype.",
75	- ACLNN_ERR_PARAM_INVALID, "Self dtype %s and other dtype %s can not promote dtype.",	65	+ op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
76	- op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
77	return false;	66	return false;
78	}	67	}
79		68
80	// 检查promoteType的数据类型是否在equal算子的支持列表内	69	// 检查promoteType的数据类型是否在equal算子的支持列表内
81	if (!CheckType(promoteType, inputSupportList)) {	70	if (!CheckType(promoteType, inputSupportList)) {
82	- OP_LOGE(	71	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
83	- ACLNN_ERR_PARAM_INVALID,	72	+ "Promoted dtype %s for self dtype %s and other dtype %s is not in the supported dtype list [%s].",
84	- "Self dtype %s and other dtype %s get promoteType dtype %s should be in "	73	+ op::ToString(promoteType).GetString(), op::ToString(self->GetDataType()).GetString(),
85	- "dtype support list [%s].",	74	+ op::ToString(other->GetDataType()).GetString(), op::ToString(inputSupportList).GetString());
86	- op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString(),
87	- op::ToString(promoteType).GetString(), op::ToString(inputSupportList).GetString());
88	return false;	75	return false;
89	}	76	}
90		77
	@@ -104,9 +91,8 @@ static bool CheckShape(const aclTensor* self, const aclTensor* other, const aclT
104	OP_CHECK_BROADCAST_AND_INFER_SHAPE(self, other, outShape, return false);	91	OP_CHECK_BROADCAST_AND_INFER_SHAPE(self, other, outShape, return false);
105		92
106	if (outShape != out->GetViewShape()) {	93	if (outShape != out->GetViewShape()) {
107	- OP_LOGE(	94	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Broadcast shape %s does not match the output shape %s.",
108	- ACLNN_ERR_PARAM_INVALID, "BroadcastShape %s is not equal out's shape %s.",	95	+ op::ToString(outShape).GetString(), op::ToString(out->GetViewShape()).GetString());
109	- op::ToString(outShape).GetString(), op::ToString(out->GetViewShape()).GetString());
110	return false;	96	return false;
111	}	97	}
112	return true;	98	return true;
	@@ -114,25 +100,23 @@ static bool CheckShape(const aclTensor* self, const aclTensor* other, const aclT
114		100
115	static aclnnStatus CheckParams(const aclTensor* self, const aclTensor* other, const aclTensor* out)	101	static aclnnStatus CheckParams(const aclTensor* self, const aclTensor* other, const aclTensor* out)
116	{	102	{
117	- // 1. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
118	CHECK_RET(CheckDtypeValid(self, other, out), ACLNN_ERR_PARAM_INVALID);	103	CHECK_RET(CheckDtypeValid(self, other, out), ACLNN_ERR_PARAM_INVALID);
119	-
120	- // 2. 检查双输入是否能broadcast,检查boradcast后的输出与out是否一致
121	CHECK_RET(CheckShape(self, other, out), ACLNN_ERR_PARAM_INVALID);	104	CHECK_RET(CheckShape(self, other, out), ACLNN_ERR_PARAM_INVALID);
122		105
123	return ACLNN_SUCCESS;	106	return ACLNN_SUCCESS;
124	}	107	}
125		108
126	-aclnnStatus aclnnEqTensorGetWorkspaceSize(	109	+aclnnStatus aclnnEqTensorGetWorkspaceSize(const aclTensor* self, const aclTensor* other, aclTensor* out,
127	- const aclTensor* self, const aclTensor* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)	110	+ uint64_t* workspaceSize, aclOpExecutor** executor)
128	{	111	{
		112	+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
129	L2_DFX_PHASE_1(aclnnEqTensor, DFX_IN(self, other), DFX_OUT(out));	113	L2_DFX_PHASE_1(aclnnEqTensor, DFX_IN(self, other), DFX_OUT(out));
130		114
131	// 固定写法，创建OpExecutor	115	// 固定写法，创建OpExecutor
132	auto uniqueExecutor = CREATE_EXECUTOR();	116	auto uniqueExecutor = CREATE_EXECUTOR();
133	CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);	117	CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
134		118
135	- // 1. 检查三个入参参数是否为空指针	119	+ // 检查三个入参是否为空指针
136	CHECK_RET(CheckNotNull(self, other, out), ACLNN_ERR_PARAM_NULLPTR);	120	CHECK_RET(CheckNotNull(self, other, out), ACLNN_ERR_PARAM_NULLPTR);
137		121
138	// 固定写法，参数检查	122	// 固定写法，参数检查
	@@ -164,7 +148,6 @@ aclnnStatus aclnnEqTensorGetWorkspaceSize(
164	auto otherCasted = l0op::Cast(otherContiguous, promoteType, uniqueExecutor.get());	148	auto otherCasted = l0op::Cast(otherContiguous, promoteType, uniqueExecutor.get());
165	CHECK_RET(otherCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);	149	CHECK_RET(otherCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
166		150
167	- // 调用Equal算子kernel
168	const aclTensor* equalOpOut = l0op::Equal(selfCasted, otherCasted, uniqueExecutor.get());	151	const aclTensor* equalOpOut = l0op::Equal(selfCasted, otherCasted, uniqueExecutor.get());
169	CHECK_RET(equalOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);	152	CHECK_RET(equalOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
170		153
	@@ -184,27 +167,27 @@ aclnnStatus aclnnEqTensorGetWorkspaceSize(
184		167
185	static inline aclnnStatus CheckInplace(const aclTensor* selfRef, const aclTensor* other)	168	static inline aclnnStatus CheckInplace(const aclTensor* selfRef, const aclTensor* other)
186	{	169	{
187	- OP_CHECK_NULL(selfRef, return false);	170	+ OP_CHECK_NULL(selfRef, return ACLNN_ERR_PARAM_NULLPTR);
188	- OP_CHECK_NULL(other, return false);	171	+ OP_CHECK_NULL(other, return ACLNN_ERR_PARAM_NULLPTR);
		172	+ OP_CHECK(selfRef->IsEmpty() == other->IsEmpty(),
		173	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "selfRef and other must either both be empty or both be non-empty."),
		174	+ return ACLNN_ERR_PARAM_INVALID);
189		175
190	op::Shape broadcastShape;	176	op::Shape broadcastShape;
191	OP_CHECK(	177	OP_CHECK(
192	BroadcastInferShape(selfRef->GetViewShape(), other->GetViewShape(), broadcastShape),	178	BroadcastInferShape(selfRef->GetViewShape(), other->GetViewShape(), broadcastShape),
193	- OP_LOGE(	179	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Shape of selfRef and other can't broadcast, got %s, %s.",
194	- ACLNN_ERR_PARAM_INVALID, "Shape of selfRef and other can't broadcast, got %s, %s.",	180	+ op::ToString(selfRef->GetViewShape()).GetString(), op::ToString(other->GetViewShape()).GetString()),
195	- op::ToString(selfRef->GetViewShape()).GetString(), op::ToString(other->GetViewShape()).GetString()),
196	- return ACLNN_ERR_PARAM_INVALID);
197	- OP_CHECK(
198	- selfRef->GetViewShape() == broadcastShape,
199	- OP_LOGE(
200	- ACLNN_ERR_PARAM_NULLPTR, "Expected shape of selfRef should be %s, but got %s.",
201	- op::ToString(broadcastShape).GetString(), op::ToString(selfRef->GetViewShape()).GetString()),
202	return ACLNN_ERR_PARAM_INVALID);	181	return ACLNN_ERR_PARAM_INVALID);
		182	+ OP_CHECK(selfRef->GetViewShape() == broadcastShape,
		183	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Expected shape of selfRef should be %s, but got %s.",
		184	+ op::ToString(broadcastShape).GetString(), op::ToString(selfRef->GetViewShape()).GetString()),
		185	+ return ACLNN_ERR_PARAM_INVALID);
203	return ACLNN_SUCCESS;	186	return ACLNN_SUCCESS;
204	}	187	}
205		188
206	-aclnnStatus aclnnInplaceEqTensorGetWorkspaceSize(	189	+aclnnStatus aclnnInplaceEqTensorGetWorkspaceSize(const aclTensor* selfRef, const aclTensor* other,
207	- const aclTensor* selfRef, const aclTensor* other, uint64_t* workspaceSize, aclOpExecutor** executor)	190	+ uint64_t* workspaceSize, aclOpExecutor** executor)
208	{	191	{
209	auto ret = CheckInplace(selfRef, other);	192	auto ret = CheckInplace(selfRef, other);
210	CHECK_RET(ret == ACLNN_SUCCESS, ret);	193	CHECK_RET(ret == ACLNN_SUCCESS, ret);

Mexperimental/math/equal/op_api/aclnn_eq_tensor.h+32-23

	@@ -1,18 +1,21 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	#ifndef OP_API_INC_EQTENSOR_H_	11	#ifndef OP_API_INC_EQTENSOR_H_
12	#define OP_API_INC_EQTENSOR_H_	12	#define OP_API_INC_EQTENSOR_H_
13		13
14	#include "aclnn/aclnn_base.h"	14	#include "aclnn/aclnn_base.h"
15	-#include "aclnn_util.h"	15	+
		16	+#ifndef ACLNN_API
		17	+#define ACLNN_API __attribute__((visibility("default")))
		18	+#endif
16		19
17	#ifdef __cplusplus	20	#ifdef __cplusplus
18	extern "C" {	21	extern "C" {
	@@ -48,56 +51,62 @@ extern "C" {
48	* @param [out] executor: 返回op执行器，包含算子计算流程。	51	* @param [out] executor: 返回op执行器，包含算子计算流程。
49	* @return aclnnStatus: 返回状态码。	52	* @return aclnnStatus: 返回状态码。
50	*/	53	*/
51	-ACLNN_API aclnnStatus aclnnEqTensorGetWorkspaceSize(	54	+ACLNN_API aclnnStatus aclnnEqTensorGetWorkspaceSize(const aclTensor* self, const aclTensor* other, aclTensor* out,
52	- const aclTensor* self, const aclTensor* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);	55	+ uint64_t* workspaceSize, aclOpExecutor** executor);
53		56
54	/**	57	/**
55	* @brief aclnnEqTensor的第二段接口，用于执行计算。	58	* @brief aclnnEqTensor的第二段接口，用于执行计算。
56	* @param [in] workspace: 在npu device侧申请的workspace内存起址。	59	* @param [in] workspace: 在npu device侧申请的workspace内存起址。
57	- * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnEqTensorGetWorkspaceSize获取。	60	+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，
		61	+ * 由第一段接口aclnnEqTensorGetWorkspaceSize获取。
58	* @param [in] stream: acl stream流。	62	* @param [in] stream: acl stream流。
59	* @param [in] executor: op执行器，包含了算子计算流程。	63	* @param [in] executor: op执行器，包含了算子计算流程。
60	* @return aclnnStatus: 返回状态码。	64	* @return aclnnStatus: 返回状态码。
61	*/	65	*/
62	-ACLNN_API aclnnStatus	66	+ACLNN_API aclnnStatus aclnnEqTensor(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
63	-aclnnEqTensor(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);	67	+ aclrtStream stream);
64		68
65	/**	69	/**
66	* @brief aclnnInplaceEqTensor的第一段接口，根据具体的计算流程，计算workspace大小。	70	* @brief aclnnInplaceEqTensor的第一段接口，根据具体的计算流程，计算workspace大小。
67	* @domain aclnn_math	71	* @domain aclnn_math
68	*	72	*
69	- * 算子功能：计算selfRef中元素的值与other的值是否相等，将selfRef每个元素与other的值的比较结果原地返回到selfRef中；	73	+ * 算子功能：计算selfRef中元素的值与other的值是否相等，
		74	+ * 将selfRef每个元素与other的值的比较结果原地返回到selfRef中；
70	*	75	*
71	* 计算公式：$$ selfRef = (selfRef == other) ? [True] : [False] $$	76	* 计算公式：$$ selfRef = (selfRef == other) ? [True] : [False] $$
72	*	77	*
73	* @param [in] selfRef: npu device侧的aclTensor，	78	* @param [in] selfRef: npu device侧的aclTensor，
74	- * 数据类型支持FLOAT16,FLOAT,INT64,INT32,INT8,UINT8,BOOL,BFLOAT16(仅1971支持),DOUBLE,INT16,COMPLEX64,COMPLEX128，	79	+ * 数据类型支持FLOAT16,FLOAT,INT64,INT32,INT8,UINT8,BOOL,BFLOAT16(仅1971支持),DOUBLE,INT16,
		80	+ * COMPLEX64,COMPLEX128，
75	* 数据格式支持ND，支持非连续的Tensor。	81	* 数据格式支持ND，支持非连续的Tensor。
76	* @param [in] other: npu device侧的aclTensor，	82	* @param [in] other: npu device侧的aclTensor，
77	- * 数据类型支持FLOAT16,FLOAT,INT64,INT32,INT8,UINT8,BOOL,BFLOAT16(仅1971支持) ,DOUBLE,INT16,COMPLEX64,COMPLEX128，	83	+ * 数据类型支持FLOAT16,FLOAT,INT64,INT32,INT8,UINT8,BOOL,BFLOAT16(仅1971支持),DOUBLE,INT16,
		84	+ * COMPLEX64,COMPLEX128，
78	* 数据格式支持ND，支持非连续的Tensor。	85	* 数据格式支持ND，支持非连续的Tensor。
79	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。	86	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
80	* @param [out] executor: 返回op执行器，包含算子计算流程。	87	* @param [out] executor: 返回op执行器，包含算子计算流程。
81	* @return aclnnStatus: 返回状态码。	88	* @return aclnnStatus: 返回状态码。
82	*/	89	*/
83	-ACLNN_API aclnnStatus aclnnInplaceEqTensorGetWorkspaceSize(	90	+ACLNN_API aclnnStatus aclnnInplaceEqTensorGetWorkspaceSize(const aclTensor* selfRef, const aclTensor* other,
84	- const aclTensor* selfRef, const aclTensor* other, uint64_t* workspaceSize, aclOpExecutor** executor);	91	+ uint64_t* workspaceSize, aclOpExecutor** executor);
85		92
86	/**	93	/**
87	* @brief aclnnInplaceEqTensor的第二段接口，用于执行计算。	94	* @brief aclnnInplaceEqTensor的第二段接口，用于执行计算。
88	*	95	*
89	- * 算子功能：计算selfRef中元素的值与other的值是否相等，将selfRef每个元素与other的值的比较结果原地返回到selfRef中；	96	+ * 算子功能：计算selfRef中元素的值与other的值是否相等，
		97	+ * 将selfRef每个元素与other的值的比较结果原地返回到selfRef中；
90	*	98	*
91	* 计算公式：$$ selfRef = (selfRef == other) ? [True] : [False] $$	99	* 计算公式：$$ selfRef = (selfRef == other) ? [True] : [False] $$
92	*	100	*
93	* @param [in] workspace: 在npu device侧申请的workspace内存起址。	101	* @param [in] workspace: 在npu device侧申请的workspace内存起址。
94	- * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnInplaceAddGetWorkspaceSize获取。	102	+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，
		103	+ * 由第一段接口aclnnInplaceEqTensorGetWorkspaceSize获取。
95	* @param [in] executor: op执行器，包含了算子计算流程。	104	* @param [in] executor: op执行器，包含了算子计算流程。
96	* @param [in] stream: acl stream流。	105	* @param [in] stream: acl stream流。
97	* @return aclnnStatus: 返回状态码。	106	* @return aclnnStatus: 返回状态码。
98	*/	107	*/
99	-ACLNN_API aclnnStatus	108	+ACLNN_API aclnnStatus aclnnInplaceEqTensor(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
100	-aclnnInplaceEqTensor(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);	109	+ aclrtStream stream);
101		110
102	#ifdef __cplusplus	111	#ifdef __cplusplus
103	}	112	}

Mexperimental/math/equal/op_api/equal.cpp+54-43

	@@ -1,12 +1,12 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10	#include "equal.h"	10	#include "equal.h"
11	#include "opdev/aicpu/aicpu_task.h"	11	#include "opdev/aicpu/aicpu_task.h"
12	#include "opdev/make_op_executor.h"	12	#include "opdev/make_op_executor.h"
	@@ -23,52 +23,64 @@ namespace l0op {
23		23
24	OP_TYPE_REGISTER(Equal);	24	OP_TYPE_REGISTER(Equal);
25		25
26	-// A2	26	+static inline const std::initializer_list<op::DataType>& GetAiCoreDtypeSupportList()
27	-static const std::initializer_list<op::DataType> AICORE910B_DTYPE_SUPPORT_LIST = {	27	+{
28	- op::DataType::DT_FLOAT, op::DataType::DT_INT32, op::DataType::DT_FLOAT16,	28	+ // Direct A2/A3 Ascend C kernel dtypes.
29	- op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_BF16};	29	+ static const std::initializer_list<op::DataType> dtypeSupportList = {
		30	+ op::DataType::DT_FLOAT, op::DataType::DT_INT32, op::DataType::DT_UINT32, op::DataType::DT_INT64,
		31	+ op::DataType::DT_FLOAT16, op::DataType::DT_INT16, op::DataType::DT_INT8, op::DataType::DT_UINT8,
		32	+ op::DataType::DT_BF16, op::DataType::DT_BOOL};
		33	+ return dtypeSupportList;
		34	+}
30		35
31	-// AICPU TF	36	+static inline const std::initializer_list<op::DataType>& GetAiCpuTfDtypeSupportList()
32	-static const std::initializer_list<op::DataType> AICPU_TF_DTYPE_SUPPORT_LIST = {	37	+{
33	- op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128};	38	+ // Public ACLNN compatibility types that are not implemented by the direct Ascend C kernel.
		39	+ static const std::initializer_list<op::DataType> dtypeSupportList = {op::DataType::DT_COMPLEX64,
		40	+ op::DataType::DT_COMPLEX128};
		41	+ return dtypeSupportList;
		42	+}
34		43
35	// 根据芯片类型、dtype判断算子是否支持走aicore	44	// 根据芯片类型、dtype判断算子是否支持走aicore
36	static bool IsAiCoreSupport(const aclTensor* self)	45	static bool IsAiCoreSupport(const aclTensor* self)
37	{	46	{
38	- return CheckType(self->GetDataType(), AICORE910B_DTYPE_SUPPORT_LIST);	47	+ return CheckType(self->GetDataType(), GetAiCoreDtypeSupportList());
		48	+}
		49	+
		50	+static bool IsAiCpuTfSupport(const aclTensor* self)
		51	+{
		52	+ return CheckType(self->GetDataType(), GetAiCpuTfDtypeSupportList());
39	}	53	}
40		54
41	// AICORE算子kernel	55	// AICORE算子kernel
42	-static const aclTensor* EqualAiCore(	56	+static const aclTensor* EqualAiCore(const aclTensor* self, const aclTensor* other, aclTensor* eqOut,
43	- const aclTensor* self, const aclTensor* other, aclTensor* eqOut, aclOpExecutor* executor)	57	+ aclOpExecutor* executor)
44	{	58	{
45	L0_DFX(EqualAiCore, self, other, eqOut);	59	L0_DFX(EqualAiCore, self, other, eqOut);
46	// 使用框架宏 ADD_TO_LAUNCHER_LIST_AICORE，将Equal算子加入任务队列	60	// 使用框架宏 ADD_TO_LAUNCHER_LIST_AICORE，将Equal算子加入任务队列
47	auto ret = ADD_TO_LAUNCHER_LIST_AICORE(Equal, OP_INPUT(self, other), OP_OUTPUT(eqOut));	61	auto ret = ADD_TO_LAUNCHER_LIST_AICORE(Equal, OP_INPUT(self, other), OP_OUTPUT(eqOut));
48	- OP_CHECK(	62	+ OP_CHECK(ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "EqualAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),
49	- ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "EqualAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),	63	+ return nullptr);
50	- return nullptr);
51	return eqOut;	64	return eqOut;
52	}	65	}
53		66
54	// AICPU算子kernel	67	// AICPU算子kernel
55	-static const aclTensor* EqualAiCpu(	68	+static const aclTensor* EqualAiCpu(const aclTensor* self, const aclTensor* other, aclTensor* eqOut,
56	- const aclTensor* self, const aclTensor* other, aclTensor* eqOut, aclOpExecutor* executor)	69	+ aclOpExecutor* executor)
57	{	70	{
58	L0_DFX(EqualAiCpu, self, other);	71	L0_DFX(EqualAiCpu, self, other);
59		72
60	static internal::AicpuTaskSpace space("Equal");	73	static internal::AicpuTaskSpace space("Equal");
61	// 使用框架宏ADD_TO_LAUNCHER_LIST_AICPU，将Equal算子加入任务队列	74	// 使用框架宏ADD_TO_LAUNCHER_LIST_AICPU，将Equal算子加入任务队列
62	auto ret = ADD_TO_LAUNCHER_LIST_AICPU(Equal, OP_ATTR_NAMES(), OP_INPUT(self, other), OP_OUTPUT(eqOut));	75	auto ret = ADD_TO_LAUNCHER_LIST_AICPU(Equal, OP_ATTR_NAMES(), OP_INPUT(self, other), OP_OUTPUT(eqOut));
63	- OP_CHECK(	76	+ OP_CHECK(ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "EqualAiCpu ADD_TO_LAUNCHER_LIST_AICPU failed."),
64	- ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "EqualAiCpu ADD_TO_LAUNCHER_LIST_AICPU failed."),	77	+ return nullptr);
65	- return nullptr);
66	return eqOut;	78	return eqOut;
67	}	79	}
68		80
69	// AICPU算子 TF	81	// AICPU算子 TF
70	-static const aclTensor* EqualAiCpuTf(	82	+static const aclTensor* EqualAiCpuTf(const aclTensor* self, const aclTensor* other, aclTensor* eqOut,
71	- const aclTensor* self, const aclTensor* other, aclTensor* eqOut, aclOpExecutor* executor)	83	+ aclOpExecutor* executor)
72	{	84	{
73	L0_DFX(EqualAiCpuTf, self, other);	85	L0_DFX(EqualAiCpuTf, self, other);
74		86
	@@ -76,12 +88,10 @@ static const aclTensor* EqualAiCpuTf(
76	static internal::AicpuTaskSpace space("Equal", ge::DEPEND_IN_SHAPE, true);	88	static internal::AicpuTaskSpace space("Equal", ge::DEPEND_IN_SHAPE, true);
77		89
78	// 使用框架宏ADD_TO_LAUNCHER_LIST_AICPU，将Equal算子加入任务队列	90	// 使用框架宏ADD_TO_LAUNCHER_LIST_AICPU，将Equal算子加入任务队列
79	- auto ret = ADD_TO_LAUNCHER_LIST_AICPU(	91	+ auto ret = ADD_TO_LAUNCHER_LIST_AICPU(Equal, OP_ATTR_NAMES({"T", "incompatible_shape_error"}),
80	- Equal, OP_ATTR_NAMES({"T", "incompatible_shape_error"}), OP_INPUT(self, other), OP_OUTPUT(eqOut),	92	+ OP_INPUT(self, other), OP_OUTPUT(eqOut), OP_ATTR(self->GetDataType(), true));
81	- OP_ATTR(self->GetDataType(), true));	93	+ OP_CHECK(ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "EqualAiCpuTf ADD_TO_LAUNCHER_LIST_AICPU failed."),
82	- OP_CHECK(	94	+ return nullptr);
83	- ret == ACL_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "EqualAiCpuTf ADD_TO_LAUNCHER_LIST_AICPU failed."),
84	- return nullptr);
85	return eqOut;	95	return eqOut;
86	}	96	}
87		97
	@@ -89,22 +99,23 @@ const aclTensor* Equal(const aclTensor* self, const aclTensor* other, aclOpExecu
89	{	99	{
90	op::Shape outShape;	100	op::Shape outShape;
91	if (!BroadcastInferShape(self->GetViewShape(), other->GetViewShape(), outShape)) {	101	if (!BroadcastInferShape(self->GetViewShape(), other->GetViewShape(), outShape)) {
92	- OP_LOGE(	102	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Broadcast %s and %s failed.", op::ToString(self->GetViewShape()).GetString(),
93	- ACLNN_ERR_PARAM_INVALID, "Broadcast %s and %s failed.", op::ToString(self->GetViewShape()).GetString(),	103	+ op::ToString(other->GetViewShape()).GetString());
94	- op::ToString(other->GetViewShape()).GetString());
95	return nullptr;	104	return nullptr;
96	}	105	}
97		106
98	auto eqOut = executor->AllocTensor(outShape, op::DataType::DT_BOOL, op::Format::FORMAT_ND);	107	auto eqOut = executor->AllocTensor(outShape, op::DataType::DT_BOOL, op::Format::FORMAT_ND);
		108	+ OP_CHECK(eqOut != nullptr, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "Equal output tensor allocation failed."),
		109	+ return nullptr);
99		110
100	if (IsAiCoreSupport(self) && IsAiCoreSupport(other)) {	111	if (IsAiCoreSupport(self) && IsAiCoreSupport(other)) {
101	return EqualAiCore(self, other, eqOut, executor);	112	return EqualAiCore(self, other, eqOut, executor);
102	- } else {
103	- if (CheckType(self->GetDataType(), AICPU_TF_DTYPE_SUPPORT_LIST) \|\|
104	- CheckType(other->GetDataType(), AICPU_TF_DTYPE_SUPPORT_LIST)) {
105	- return EqualAiCpuTf(self, other, eqOut, executor);
106	- }
107	- return EqualAiCpu(self, other, eqOut, executor);
108	}	113	}
		114	+
		115	+ if (IsAiCpuTfSupport(self) \|\| IsAiCpuTfSupport(other)) {
		116	+ return EqualAiCpuTf(self, other, eqOut, executor);
		117	+ }
		118	+
		119	+ return EqualAiCpu(self, other, eqOut, executor);
109	}	120	}
110	} // namespace l0op	121	} // namespace l0op

Mexperimental/math/equal/op_api/equal.h+12-12

Mexperimental/math/equal/op_host/equal_def.cpp+29-32

	@@ -1,12 +1,12 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	/*!	11	/*!
12	* \file equal_def.cpp	12	* \file equal_def.cpp
	@@ -19,37 +19,34 @@ class Equal : public OpDef {
19	public:	19	public:
20	explicit Equal(const char* name) : OpDef(name)	20	explicit Equal(const char* name) : OpDef(name)
21	{	21	{
		22	+ // Direct Ascend C kernel dtypes. COMPLEX64/COMPLEX128 keep using the ACLNN AICPU-TF fallback.
		23	+ const std::initializer_list<ge::DataType> aicoreInputDtypes = {
		24	+ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8,
		25	+ ge::DT_INT16, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64, ge::DT_BOOL};
		26	+ const std::initializer_list<ge::Format> aicoreFormats = {
		27	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
		28	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
		29	+ const std::initializer_list<ge::DataType> aicoreOutputDtypes = {
		30	+ ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL,
		31	+ ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL};
		32	+
22	this->Input("x1")	33	this->Input("x1")
23	.ParamType(REQUIRED)	34	.ParamType(REQUIRED)
24	- .DataType(	35	+ .DataType(aicoreInputDtypes)
25	- {ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT32, ge::DT_UINT32})	36	+ .Format(aicoreFormats)
26	- .Format(	37	+ .UnknownShapeFormat(aicoreFormats);
27	- {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
28	- ge::FORMAT_ND})
29	- .UnknownShapeFormat(
30	- {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
31	- ge::FORMAT_ND});
32	this->Input("x2")	38	this->Input("x2")
33	.ParamType(REQUIRED)	39	.ParamType(REQUIRED)
34	- .DataType(	40	+ .DataType(aicoreInputDtypes)
35	- {ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT32, ge::DT_UINT32})	41	+ .Format(aicoreFormats)
36	- .Format(	42	+ .UnknownShapeFormat(aicoreFormats);
37	- {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
38	- ge::FORMAT_ND})
39	- .UnknownShapeFormat(
40	- {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
41	- ge::FORMAT_ND});
42	this->Output("y")	43	this->Output("y")
43	.ParamType(REQUIRED)	44	.ParamType(REQUIRED)
44	- .DataType({ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL, ge::DT_BOOL})	45	+ .DataType(aicoreOutputDtypes)
45	- .Format(	46	+ .Format(aicoreFormats)
46	- {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,	47	+ .UnknownShapeFormat(aicoreFormats);
47	- ge::FORMAT_ND})
48	- .UnknownShapeFormat(
49	- {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
50	- ge::FORMAT_ND});
51		48
52	- this->AICore().AddConfig("ascend910b");	49	+ this->AICore().AddConfig("ascend910b").AddConfig("ascend910_93");
53	}	50	}
54	};	51	};
55	OP_ADD(Equal); // 添加算子信息库	52	OP_ADD(Equal); // 添加算子信息库

Mexperimental/math/equal/op_host/equal_infershape.cpp+10-15

	@@ -1,19 +1,19 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	/*!	11	/*!
12	* \file equal_infershape.cpp	12	* \file equal_infershape.cpp
13	* \brief	13	* \brief
14	*/	14	*/
		15	+#include "op_host/infershape_broadcast_util.h"
15	#include "register/op_impl_registry.h"	16	#include "register/op_impl_registry.h"
16	-#include "log/log.h"
17		17
18	using namespace ge;	18	using namespace ge;
19		19
	@@ -21,12 +21,7 @@ namespace ops {
21		21
22	static ge::graphStatus InferShapeEqual(gert::InferShapeContext* context)	22	static ge::graphStatus InferShapeEqual(gert::InferShapeContext* context)
23	{	23	{
24	- const gert::Shape* x1_shape = context->GetInputShape(0);	24	+ return Ops::Base::InferShape4Broadcast(context);
25	- OP_CHECK_NULL_WITH_CONTEXT(context, x1_shape);
26	- gert::Shape* y_shape = context->GetOutputShape(0);
27	- OP_CHECK_NULL_WITH_CONTEXT(context, y_shape);
28	- y_shape = x1_shape;
29	- return GRAPH_SUCCESS;
30	}	25	}
31		26
32	IMPL_OP_INFERSHAPE(Equal).InferShape(InferShapeEqual);	27	IMPL_OP_INFERSHAPE(Equal).InferShape(InferShapeEqual);

Mexperimental/math/equal/op_host/equal_tiling.cpp+469-104

	@@ -1,142 +1,507 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	/*!	11	/*!
12	* \file equal_tiling.cpp	12	* \file equal_tiling.cpp
13	- * \brief	13	+ * \brief Host tiling for contiguous and broadcast Equal paths.
14	*/	14	*/
15		15
16	-#include "register/op_impl_registry.h"
17	-#include "tiling/platform/platform_ascendc.h"
18	#include "graph/utils/type_utils.h"	16	#include "graph/utils/type_utils.h"
19	#include "log/log.h"	17	#include "log/log.h"
20	-#include "util/math_util.h"	18	+#include "register/op_impl_registry.h"
		19	+#include "tiling/tiling_api.h"
		20	+#include "tiling/platform/platform_ascendc.h"
21	#include "../op_kernel/equal_tiling_data.h"	21	#include "../op_kernel/equal_tiling_data.h"
22	#include "../op_kernel/equal_tiling_key.h"	22	#include "../op_kernel/equal_tiling_key.h"
		23	+#include <algorithm>
		24	+#include <array>
		25	+#include <limits>
23		26
24	namespace optiling {	27	namespace optiling {
		28	+namespace {
25		29
26	-constexpr uint64_t BLOCK_SIZE = 32;	30	+struct EqualCompileInfo {};
27	-constexpr uint64_t ALIGN_BYTES = 256;
28	-constexpr uint64_t INT32_FP_UB_COMPUTE_RESERVE = 256 * 7; // 256B alignment × 7 slots reserved for int32/fp compute
29	-constexpr uint64_t INT32_FP_UB_SYS_RESERVE = 8 * 1024; // 8KB system reserve for int32/fp path
30	-constexpr uint64_t UB_DATA_NUMBER_DEFAULT = 6;
31	-constexpr uint64_t UB_DATA_NUMBER_INT8 = 12;
32	-constexpr uint64_t UB_DATA_NUMBER_INT32_FP = 5;
33		31
34	-// tiling 分发入口	32	+constexpr uint64_t REPEAT_BYTES = 256;
35	-static ge::graphStatus EqualTilingFunc(gert::TilingContext* context)	33	+constexpr uint64_t DEFAULT_BLOCK_ALIGN_ELEMENTS = 512;
		34	+constexpr uint64_t FP16_BLOCK_ALIGN_ELEMENTS = REPEAT_BYTES / sizeof(uint16_t);
		35	+constexpr uint64_t INT8_BLOCK_ALIGN_ELEMENTS = REPEAT_BYTES / sizeof(uint8_t);
		36	+constexpr uint64_t MIN_BYTES_PER_CORE = 4 * 1024;
		37	+constexpr uint64_t INT8_MIN_BYTES_PER_CORE = 2 * 1024;
		38	+constexpr uint64_t GENERAL_BROADCAST_MIN_BYTES_PER_CORE = 2 * 1024;
		39	+constexpr uint64_t CONTIGUOUS_QUEUE_BUFFER_NUM = 2;
		40	+constexpr uint64_t BROADCAST_QUEUE_BUFFER_NUM = 1;
		41	+constexpr uint64_t INPUT_QUEUE_NUM = 2;
		42	+constexpr uint64_t OUTPUT_QUEUE_NUM = 1;
		43	+constexpr uint64_t BOOL_BYTES = sizeof(uint8_t);
		44	+constexpr uint64_t DATA_BLOCK_BYTES = 32;
		45	+constexpr uint64_t SANDWICH_EXTRA_DATA_BLOCKS = 2;
		46	+constexpr uint64_t SELECT_UB_RESERVED_BYTES = 8 * 1024;
		47	+constexpr uint64_t DEFAULT_UB_RESERVED_BYTES = 2 * SELECT_UB_RESERVED_BYTES;
		48	+constexpr uint64_t FP16_UB_RESERVED_BYTES = 2 * SELECT_UB_RESERVED_BYTES;
		49	+constexpr uint64_t FP16_MAX_TILE_ELEMENTS = (30 * 1024) / sizeof(uint16_t);
		50	+constexpr uint64_t EXTENDED_TYPE_UB_RESERVED_BYTES = 4 * SELECT_UB_RESERVED_BYTES;
		51	+constexpr uint64_t INT64_MAX_TILE_ELEMENTS = 3072;
		52	+constexpr uint64_t INT64_TAIL_REUSE_MAX_TILE_ELEMENTS = 4608;
		53	+constexpr uint32_t SANDWICH_AXIS_COUNT = 3;
		54	+constexpr uint32_t SANDWICH_MIDDLE_AXIS_OFFSET = 1;
		55	+constexpr uint32_t SANDWICH_TAIL_AXIS_OFFSET = 2;
		56	+
		57	+struct BroadcastInfo {
		58	+ uint64_t totalLength = 1;
		59	+ uint64_t x1Length = 1;
		60	+ uint64_t x2Length = 1;
		61	+ uint32_t rank = 0;
		62	+ uint32_t mode = EQUAL_BROADCAST_CONTIGUOUS;
		63	+ uint32_t fastBroadcastInput = 0;
		64	+ std::array<uint64_t, EQUAL_MAX_BROADCAST_DIM> outShape{};
		65	+ std::array<uint64_t, EQUAL_MAX_BROADCAST_DIM> x1Stride{};
		66	+ std::array<uint64_t, EQUAL_MAX_BROADCAST_DIM> x2Stride{};
		67	+ uint64_t fastOuter = 0;
		68	+ uint64_t fastMiddle = 0;
		69	+ uint64_t fastTail = 0;
		70	+ uint64_t fastPlaneLength = 0;
		71	+ uint64_t fastSourceLength = 0;
		72	+ uint64_t fastTmpBytes = 0;
		73	+};
		74	+
		75	+uint64_t CeilDiv(uint64_t value, uint64_t divisor)
36	{	76	{
37	- EqualTilingData* tiling = context->GetTilingData<EqualTilingData>();	77	+ return value / divisor + static_cast<uint64_t>(value % divisor != 0);
38	- OP_CHECK_NULL_WITH_CONTEXT(context, tiling);	78	+}
39	- uint64_t ubSize;	79	+
40	- uint32_t bigprocessDataNum_computes = 0;	80	+bool CheckedMul(uint64_t lhs, uint64_t rhs, uint64_t& result)
41	- uint32_t tailbigprocessDataNum_computes = 0;	81	+{
42	- auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());	82	+ if (lhs != 0 && rhs > std::numeric_limits<uint64_t>::max() / lhs) {
43	- ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);	83	+ return false;
44	- auto coreNum = ascendcPlatform.GetCoreNumAiv();
45	- if (coreNum == 0) {
46	- coreNum = ascendcPlatform.GetCoreNum();
47	}	84	}
48	- auto socVersion = ascendcPlatform.GetSocVersion();	85	+ result = lhs * rhs;
49	- if (socVersion != platform_ascendc::SocVersion::ASCEND910B &&	86	+ return true;
50	- socVersion != platform_ascendc::SocVersion::ASCEND310B &&	87	+}
51	- context->GetInputDesc(0)->GetDataType() == ge::DT_BF16) {	88	+
52	- OP_LOGE(context, "bf16 is only supported on ASCEND910B and ASCEND310B");	89	+bool CheckedAdd(uint64_t lhs, uint64_t rhs, uint64_t& result)

Mexperimental/math/equal/op_kernel/equal.cpp+12-16

	@@ -1,12 +1,12 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	/*!	11	/*!
12	* \file equal.cpp	12	* \file equal.cpp
	@@ -18,18 +18,14 @@
18	using namespace NsEqual;	18	using namespace NsEqual;
19		19
20	template <uint32_t schMode>	20	template <uint32_t schMode>
		21	+// The lowercase entry name is fixed by the generated Equal kernel registration.
21	__global__ __aicore__ void equal(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)	22	__global__ __aicore__ void equal(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
22	{	23	{
23	REGISTER_TILING_DEFAULT(EqualTilingData);	24	REGISTER_TILING_DEFAULT(EqualTilingData);
24	- GET_TILING_DATA_WITH_STRUCT(EqualTilingData, tiling_data, tiling);	25	+ GET_TILING_DATA_WITH_STRUCT(EqualTilingData, tilingData, tiling);
25		26
26	- KernelEqual<DTYPE_X1> op;	27	+ KernelEqual<DTYPE_X1, schMode> op;
27		28
28	- op.Init(	29	+ op.Init(x1, x2, y, tilingData);
29	- x1, x2, y, tiling_data.smallCoreDataNum, tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
30	- tiling_data.finalSmallTileNum, tiling_data.tileDataNum, tiling_data.smallTailDataNum,
31	- tiling_data.bigTailDataNum, tiling_data.tailBlockNum, tiling_data.bigprocessDataNum_computes,
32	- tiling_data.smallprocessDataNum_computes, tiling_data.tailbigprocessDataNum_computes,
33	- tiling_data.tailsmallprocessDataNum_computes);
34	op.Process();	30	op.Process();
35	}	31	}

Mexperimental/math/equal/op_kernel/equal.h+558-198

Mexperimental/math/equal/op_kernel/equal_tiling_data.h+41-23

	@@ -1,33 +1,51 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	/*!	11	/*!
12	* \file equal_tiling_data.h	12	* \file equal_tiling_data.h
13	* \brief tiling data struct	13	* \brief tiling data struct
14	*/	14	*/
15		15
16	-#ifndef __EQUAL_TILLING_DATA_H__	16	+#ifndef EXPERIMENTAL_MATH_EQUAL_OP_KERNEL_EQUAL_TILING_DATA_H_
17	-#define __EQUAL_TILLING_DATA_H__	17	+#define EXPERIMENTAL_MATH_EQUAL_OP_KERNEL_EQUAL_TILING_DATA_H_
		18	+
		19	+constexpr uint32_t EQUAL_MAX_BROADCAST_DIM = 8;
		20	+
		21	+enum EqualBroadcastMode : uint32_t {
		22	+ EQUAL_BROADCAST_CONTIGUOUS = 0,
		23	+ EQUAL_BROADCAST_X1_SCALAR = 1,
		24	+ EQUAL_BROADCAST_X2_SCALAR = 2,
		25	+ EQUAL_BROADCAST_GENERAL = 3,
		26	+ EQUAL_BROADCAST_SANDWICH = 4,
		27	+ EQUAL_BROADCAST_TAIL_REUSE = 5,
		28	+};
18		29
19	struct EqualTilingData {	30	struct EqualTilingData {
20	- uint32_t smallCoreDataNum;	31	+ int64_t totalLength;
21	- uint32_t bigCoreDataNum;	32	+ int64_t blockLength;
22	- uint32_t finalBigTileNum;	33	+ int64_t tailBlockLength;
23	- uint32_t finalSmallTileNum;	34	+ int64_t tileLength;
24	- uint32_t tileDataNum;	35	+ uint32_t blockNum;
25	- uint32_t smallTailDataNum;	36	+ uint32_t broadcastMode;
26	- uint32_t bigTailDataNum;	37	+ uint32_t rank;
27	- uint32_t tailBlockNum;	38	+ uint32_t fastBroadcastInput;
28	- uint32_t bigprocessDataNum_computes;	39	+ int64_t x1Length;
29	- uint32_t smallprocessDataNum_computes;	40	+ int64_t x2Length;
30	- uint32_t tailbigprocessDataNum_computes;	41	+ int64_t outShape[EQUAL_MAX_BROADCAST_DIM];
31	- uint32_t tailsmallprocessDataNum_computes;	42	+ int64_t x1Stride[EQUAL_MAX_BROADCAST_DIM];
		43	+ int64_t x2Stride[EQUAL_MAX_BROADCAST_DIM];
		44	+ int64_t fastOuter;
		45	+ int64_t fastMiddle;
		46	+ int64_t fastTail;
		47	+ int64_t fastPlaneLength;
		48	+ int64_t fastSourceLength;
		49	+ uint64_t fastTmpBytes;
32	};	50	};
33	-#endif	51	+#endif // EXPERIMENTAL_MATH_EQUAL_OP_KERNEL_EQUAL_TILING_DATA_H_

Mexperimental/math/equal/op_kernel/equal_tiling_key.h+18-20

	@@ -1,35 +1,33 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	/*!	11	/*!
12	* \file equal_tiling_key.h	12	* \file equal_tiling_key.h
13	* \brief equal tiling key declare	13	* \brief equal tiling key declare
14	*/	14	*/
15		15
16	-#ifndef __EQUAL_TILING_KEY_H__	16	+#ifndef EXPERIMENTAL_MATH_EQUAL_OP_KERNEL_EQUAL_TILING_KEY_H_
17	-#define __EQUAL_TILING_KEY_H__	17	+#define EXPERIMENTAL_MATH_EQUAL_OP_KERNEL_EQUAL_TILING_KEY_H_
18		18
19	#include "ascendc/host_api/tiling/template_argument.h"	19	#include "ascendc/host_api/tiling/template_argument.h"
20		20
21	-/* Mode场景定义 */	21	+// Same-shape and broadcast scheduling modes.
22	#define ELEMENTWISE_TPL_SCH_MODE_0 0	22	#define ELEMENTWISE_TPL_SCH_MODE_0 0
23	#define ELEMENTWISE_TPL_SCH_MODE_1 1	23	#define ELEMENTWISE_TPL_SCH_MODE_1 1
24	-/* 继续定义其他Mode场景... */
25		24
26	-/* 模板参数 */	25	+// 模板参数
27	-ASCENDC_TPL_ARGS_DECL(	26	+ASCENDC_TPL_ARGS_DECL(Equal, ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0,
28	- Equal,	27	+ ELEMENTWISE_TPL_SCH_MODE_1));
29	- ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
30		28
31	-/* 模板参数组合 */	29	+// 模板参数组合
32	-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(	30	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0,
33	- ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));	31	+ ELEMENTWISE_TPL_SCH_MODE_1)));
34		32
35	-#endif	33	+#endif // EXPERIMENTAL_MATH_EQUAL_OP_KERNEL_EQUAL_TILING_KEY_H_

	@@ -1,21 +1,21 @@
1	/**	1	/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	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.	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.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	-#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_EQTENSOR_OP_H_	11	+#ifndef EXPERIMENTAL_MATH_EQUAL_OP_API_EQUAL_H_
12	-#define PTA_NPU_OP_API_INC_LEVEL0_OP_EQTENSOR_OP_H_	12	+#define EXPERIMENTAL_MATH_EQUAL_OP_API_EQUAL_H_
13		13
14	#include "opdev/op_executor.h"	14	#include "opdev/op_executor.h"
15		15
16	namespace l0op {	16	namespace l0op {
17	const aclTensor* Equal(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor);	17	const aclTensor* Equal(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor);
18		18
19	-}	19	+} // namespace l0op
20		20
21	-#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_EQTENSOR_OP_H_	21	+#endif // EXPERIMENTAL_MATH_EQUAL_OP_API_EQUAL_H_