Star1307
1424
代码
代码
Issues118
Pull Requests235
Actions
项目讨论
Wiki
分析
Star1307
1424
已合并
提交Ascend C实现的Neg算子 #648
skywang2创建于 1月4日
提交Ascend C实现的Neg算子 #648
已合并
skywang2创建于 1月4日
从skywang2/ops-mathmath-Neg合入到cann/ops-mathmaster
共 16 个文件变更+1537-0
Aexperimental/math/neg/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
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+add_all_modules_sources(OPTYPE neg ACLNNTYPE aclnn_exclude)
Aexperimental/math/neg/README.md+40-0
@@ -0,0 +1,40 @@
1+# Neg
2+## 贡献说明
3+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
F
Ffulltower3月27日

缺少docs目录 image.png

likedislike
4+|--------|------------------|-------|-----------|-----------|
5+| skywang2 | 个人开发者 | Neg | 2025/12/31 | 新增Neg算子 |
6+ 
7+### 算子描述
8+`Neg`算子对输入的数值型数据执行取负操作（y = -x）。
9+ 
10+### 算子规格描述
11+ 
12+<table>
13+<tr><th align="center">算子类型(OpType)</th><th colspan="4" align="center">Neg</th></tr>
14+<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> <tr><td rowspan="1" align="center">算子输入</td>
15+<td align="center">x</td><td align="center">tensor</td> <td align="center">int32, int8, float16, bfloat16, float32</td><td align="center">ND</td></tr>
16+ 
17+<tr><td rowspan="1" align="center">算子输出</td>
18+<td align="center">y</td><td align="center">tensor</td> <td align="center">与输入相同</td><td align="center">ND</td></tr>
19+<tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">neg</td></tr>
20+</table>
21+ 
22+### 支持的产品型号
23+本样例支持如下产品型号：
24+- Atlas A2训练系列产品
25+- Atlas 800I A2推理产品
26+ 
27+ 
28+### 环境要求
29+编译运行此样例前，请参考[《CANN软件安装指南》](https://hiascend.com/document/redirect/CannCommunityInstSoftware)完成开发运行环境的部署。
30+ 
31+ 
32+### 算子调用
33+测试命令调用方式：[build.sh](/docs/zh/invocation/quick_op_invocation.md)
34+<table>
35+ <th>目录</th><th>描述</th>
36+ <tr>
37+ <td><a href="./examples/test_aclnn_neg.cpp"> test_aclnn_neg.cpp</td><td>通过aclnn调用的方式调用Neg算子</td>
38+ </tr>
39+</table>
40+ 
Aexperimental/math/neg/docs/aclnnNeg&aclnnInplaceNeg.md+361-0
@@ -0,0 +1,361 @@
1+# aclnnNeg&aclnnInplaceNeg
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
8+ 
9+## 功能说明
10+ 
11+- 算子功能：对输入的每个元素完成相反数计算
12+- 计算公式：
13+ 
14+ $$
15+ output_i = (-1) * self_i
16+ $$
17+ 
18+## 函数原型
19+- aclnnNeg和aclnnInplaceNeg实现相同的功能，使用区别如下，请根据自身实际场景选择合适的算子。
20+ 
21+ - aclnnNeg：需新建一个输出张量对象存储计算结果。
22+ - aclnnInplaceNeg：无需新建输出张量对象，直接在输入张量的内存中存储计算结果。
23+ 
24+- 每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnNegGetWorkspaceSize”或者”aclnnInplaceNegGetWorkspaceSize“接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnNeg”或者”aclnnInplaceNeg“接口执行计算。
25+ - `aclnnStatus aclnnNegGetWorkspaceSize(const aclTensor *self, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor)`
26+ - `aclnnStatus aclnnNeg(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, const aclrtStream stream)`
27+ - `aclnnStatus aclnnInplaceNegGetWorkspaceSize(aclTensor *selfRef, uint64_t *workspaceSize, aclOpExecutor **executor)`
28+ - `aclnnStatus aclnnInplaceNeg(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)`
29+ 
30+## aclnnNegGetWorkspaceSize
31+ 
32+- **参数说明：**
33+ 
34+ - self(aclTensor*, 计算输入)：Device侧的aclTensor。支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。
35+ - out(aclTensor*, 计算输出)：Device侧的aclTensor。out的数据类型需要是self可转换的数据类型（[参考说明](../../../../docs/zh/context/互转换关系.md)）支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。
36+ - workspaceSize(uint64_t*, 计算输出)：返回用户需要在Device侧申请的workspace大小。
37+ - executor(aclOpExecutor**, 计算输出)：返回op执行器，包含了算子计算流程。
38+ 
39+- **返回值：**
40+ 
41+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
42+ 
43+ ```
44+ 第一段接口完成入参校验，出现以下场景时报错：
45+ 161001(ACLNN_ERR_PARAM_NULLPTR): 1. 传入的self、out是空指针。
46+ 161002(ACLNN_ERR_PARAM_INVALID): 1. self和out的数据类型不在支持的范围之内。
47+ 2. self与out的dtype不同。
48+ 3. self与out的shape不同。
49+ ```
50+ 
51+## aclnnNeg
52+ 
53+- **参数说明：**
54+ 
55+ - workspace(void*, 计算输入):在Device侧申请的workspace内存地址。
56+ - workspaceSize(uint64_t, 计算输入):在Device侧申请的workspace大小, 由第一段接口aclnnNegGetWorkspaceSize获取。
57+ - executor(aclOpExecutor*, 计算输入):op执行器, 包含了算子计算流程。
58+ - stream(aclrtStream, 计算输入):指定执行任务的Stream。
59+ 
60+- **返回值：**
61+ 
62+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)
63+ 
64+## aclnnInplaceNegGetWorkspaceSize
65+ 
66+- **参数说明：**
67+ 
68+ - selfRef(aclTensor*, 计算输入|计算输出)：Device侧的aclTensor。支持[非连续的Tensor](../../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../../docs/zh/context/数据格式.md)支持ND。
69+ - workspaceSize(uint64_t*, 计算输出)：返回用户需要在Device侧申请的workspace大小。
70+ - executor(aclOpExecutor**, 计算输出)：返回op执行器，包含了算子计算流程。
71+ 
72+- **返回值：**
73+ 
74+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
75+ 
76+ ```
77+ 第一段接口完成入参校验，出现以下场景时报错：
78+ 161001(ACLNN_ERR_PARAM_NULLPTR): 1. 传入的selfRef是空指针。
79+ 161002(ACLNN_ERR_PARAM_INVALID): 1. selfRef的数据类型不在支持的范围之内。
80+ 2. selfRef的维度范围超过8。
81+ ```
82+ 
83+## aclnnInplaceNeg
84+ 
85+- **参数说明：**
86+ 
87+ - workspace(void*, 计算输入):在Device侧申请的workspace内存地址。
88+ - workspaceSize(uint64_t, 计算输入):在Device侧申请的workspace大小, 由第一段接口aclnnInplaceNegGetWorkspaceSize获取。
89+ - executor(aclOpExecutor*, 计算输入):op执行器, 包含了算子计算流程。
90+ - stream(aclrtStream, 计算输入):指定执行任务的Stream。
91+ 
92+- **返回值：**
93+ 
94+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
95+ 
96+## 约束说明
97+ 
98+- 确定性计算：
99+ - aclnnNeg&aclnnInplaceNeg默认确定性实现。
100+ 
101+## 调用示例
102+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
103+**aclnnNeg示例代码：**
104+ 
105+```Cpp
106+#include <iostream>
107+#include <vector>
108+#include "acl/acl.h"
109+#include "aclnnop/aclnn_neg.h"
110+ 
111+#define CHECK_RET(cond, return_expr) \
112+ do { \
113+ if (!(cond)) { \
114+ return_expr; \
115+ } \
116+ } while (0)
117+ 
118+#define LOG_PRINT(message, ...) \
119+ do { \
120+ printf(message, ##__VA_ARGS__); \
121+ } while (0)
122+ 
123+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
124+ int64_t shapeSize = 1;
125+ for (auto i : shape) {
126+ shapeSize *= i;
127+ }
128+ return shapeSize;
129+}
130+ 
131+int Init(int32_t deviceId, aclrtStream* stream) {
132+ // 固定写法，资源初始化
133+ auto ret = aclInit(nullptr);
134+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
135+ ret = aclrtSetDevice(deviceId);
136+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
137+ ret = aclrtCreateStream(stream);
138+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
139+ return 0;
140+}
141+ 
142+template <typename T>
143+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
144+ aclDataType dataType, aclTensor** tensor) {
145+ auto size = GetShapeSize(shape) * sizeof(T);
146+ // 调用aclrtMalloc申请device侧内存
147+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
148+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
149+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
150+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
151+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
152+ 
153+ // 计算连续tensor的strides
154+ std::vector<int64_t> strides(shape.size(), 1);
155+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
156+ strides[i] = shape[i + 1] * strides[i + 1];
157+ }
158+ 
159+ // 调用aclCreateTensor接口创建aclTensor
160+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
161+ shape.data(), shape.size(), *deviceAddr);
162+ return 0;
163+}
164+ 
165+int main() {
166+ // 1. （固定写法）device/stream初始化，参考acl API手册
167+ // 根据自己的实际device填写deviceId
168+ int32_t deviceId = 0;
169+ aclrtStream stream;
170+ auto ret = Init(deviceId, &stream);
171+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
172+ 
173+ // 2. 构造输入与输出，需要根据API的接口自定义构造
174+ std::vector<int64_t> selfShape = {4, 2};
175+ std::vector<int64_t> outShape = {4, 2};
176+ void* selfDeviceAddr = nullptr;
177+ void* outDeviceAddr = nullptr;
178+ aclTensor* self = nullptr;
179+ aclTensor* out = nullptr;
180+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
181+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
182+ // 创建self aclTensor
183+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
184+ CHECK_RET(ret == ACL_SUCCESS, return ret);
185+ // 创建out aclTensor
186+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
187+ CHECK_RET(ret == ACL_SUCCESS, return ret);
188+ 
189+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
190+ uint64_t workspaceSize = 0;
191+ aclOpExecutor* executor;
192+ // 调用aclnnNeg第一段接口
193+ ret = aclnnNegGetWorkspaceSize(self, out, &workspaceSize, &executor);
194+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnNegGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
195+ // 根据第一段接口计算出的workspaceSize申请device内存
196+ void* workspaceAddr = nullptr;
197+ if (workspaceSize > 0) {
198+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
199+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
200+ }
201+ // 调用aclnnNeg第二段接口
202+ ret = aclnnNeg(workspaceAddr, workspaceSize, executor, stream);
203+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnNeg failed. ERROR: %d\n", ret); return ret);
204+ 
205+ // 4. （固定写法）同步等待任务执行结束
206+ ret = aclrtSynchronizeStream(stream);
207+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
208+ 
209+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
210+ auto size = GetShapeSize(outShape);
211+ std::vector<float> resultData(size, 0);
212+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
213+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
214+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
215+ for (int64_t i = 0; i < size; i++) {
216+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
217+ }
218+ 
219+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
220+ aclDestroyTensor(self);
221+ aclDestroyTensor(out);
222+ 
223+ // 7. 释放device资源，需要根据具体API的接口定义修改
224+ aclrtFree(selfDeviceAddr);
225+ aclrtFree(outDeviceAddr);
226+ if (workspaceSize > 0) {
227+ aclrtFree(workspaceAddr);
228+ }
229+ aclrtDestroyStream(stream);
230+ aclrtResetDevice(deviceId);
231+ aclFinalize();
232+ return 0;
233+}
234+```
235+**aclnnInplaceNeg示例代码：**
236+ 
237+```Cpp
238+#include <iostream>
239+#include <vector>
240+#include "acl/acl.h"
241+#include "aclnnop/aclnn_neg.h"
242+ 
243+#define CHECK_RET(cond, return_expr) \
244+ do { \
245+ if (!(cond)) { \
246+ return_expr; \
247+ } \
248+ } while (0)
249+ 
250+#define LOG_PRINT(message, ...) \
251+ do { \
252+ printf(message, ##__VA_ARGS__); \
253+ } while (0)
254+ 
255+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
256+ int64_t shapeSize = 1;
257+ for (auto i : shape) {
258+ shapeSize *= i;
259+ }
260+ return shapeSize;
261+}
262+ 
263+int Init(int32_t deviceId, aclrtStream* stream) {
264+ // 固定写法，资源初始化
265+ auto ret = aclInit(nullptr);
266+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
267+ ret = aclrtSetDevice(deviceId);
268+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
269+ ret = aclrtCreateStream(stream);
270+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
271+ return 0;
272+}
273+ 
274+template <typename T>
275+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
276+ aclDataType dataType, aclTensor** tensor) {
277+ auto size = GetShapeSize(shape) * sizeof(T);
278+ // 调用aclrtMalloc申请device侧内存
279+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
280+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
281+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
282+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
283+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
284+ 
285+ // 计算连续tensor的strides
286+ std::vector<int64_t> strides(shape.size(), 1);
287+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
288+ strides[i] = shape[i + 1] * strides[i + 1];
289+ }
290+ 
291+ // 调用aclCreateTensor接口创建aclTensor
292+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
293+ shape.data(), shape.size(), *deviceAddr);
294+ return 0;
295+}
296+ 
297+int main() {
298+ // 1. （固定写法）device/stream初始化，参考acl API手册
299+ // 根据自己的实际device填写deviceId
300+ int32_t deviceId = 0;
301+ aclrtStream stream;
302+ auto ret = Init(deviceId, &stream);
303+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
304+ 
305+ // 2. 构造输入与输出，需要根据API的接口自定义构造
306+ std::vector<int64_t> selfRefShape = {4, 2};
307+ void* selfRefDeviceAddr = nullptr;
308+ aclTensor* selfRef = nullptr;
309+ 
310+ std::vector<float> selfRefHostData = {0, 1, 2, 3, 4, 5, 6, 7};
311+ 
312+ // 创建selfRef aclTensor
313+ ret = CreateAclTensor(selfRefHostData, selfRefShape, &selfRefDeviceAddr, aclDataType::ACL_FLOAT, &selfRef);
314+ CHECK_RET(ret == ACL_SUCCESS, return ret);
315+ 
316+ // 3. 调用CANN算子库API
317+ uint64_t workspaceSize = 0;
318+ aclOpExecutor* executor;
319+ // 调用aclnnInplaceNeg第一段接口
320+ ret = aclnnInplaceNegGetWorkspaceSize(selfRef, &workspaceSize, &executor);
321+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceNegGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
322+ // 根据第一段接口计算出的workspaceSize申请device内存
323+ void* workspaceAddr = nullptr;
324+ if (workspaceSize > 0) {
325+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
326+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
327+ }
328+ // 调用aclnnInplaceNeg第二段接口
329+ ret = aclnnInplaceNeg(workspaceAddr, workspaceSize, executor, stream);
330+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceNeg failed. ERROR: %d\n", ret); return ret);
331+ 
332+ // 4. （固定写法）同步等待任务执行结束
333+ ret = aclrtSynchronizeStream(stream);
334+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
335+ 
336+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
337+ auto size = GetShapeSize(selfRefShape);
338+ std::vector<float> resultData(size, 0);
339+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]),
340+ selfRefDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
341+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy resultData from device to host failed. ERROR: %d\n", ret);
342+ return ret);
343+ for (int64_t i = 0; i < size; i++) {
344+ LOG_PRINT("resultData[%ld] is: %f\n", i, resultData[i]);
345+ }
346+ 
347+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
348+ aclDestroyTensor(selfRef);
349+ 
350+ // 7. 释放device资源，需要根据具体API的接口定义修改
351+ aclrtFree(selfRefDeviceAddr);
352+ if (workspaceSize > 0) {
353+ aclrtFree(workspaceAddr);
354+ }
355+ aclrtDestroyStream(stream);
356+ aclrtResetDevice(deviceId);
357+ aclFinalize();
358+ return 0;
359+}
360+```
361+ 
Aexperimental/math/neg/examples/test_aclnn_inplace_neg.cpp+139-0
@@ -0,0 +1,139 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <vector>
13+#include "acl/acl.h"
14+#include "aclnnop/aclnn_neg.h"
15+ 
16+#define CHECK_RET(cond, return_expr) \
17+ do { \
18+ if (!(cond)) { \
19+ return_expr; \
20+ } \
21+ } while (0)
22+ 
23+#define LOG_PRINT(message, ...) \
24+ do { \
25+ printf(message, ##__VA_ARGS__); \
26+ } while (0)
27+ 
28+int64_t GetShapeSize(const std::vector<int64_t>& shape)
29+{
30+ int64_t shapeSize = 1;
31+ for (auto i : shape) {
32+ shapeSize *= i;
33+ }
34+ return shapeSize;
35+}
36+ 
37+int Init(int32_t deviceId, aclrtStream* stream)
38+{
39+ // 固定写法，资源初始化
40+ auto ret = aclInit(nullptr);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtSetDevice(deviceId);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
44+ ret = aclrtCreateStream(stream);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
46+ return 0;
47+}
48+ 
49+template <typename T>
50+int CreateAclTensor(
51+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
52+ aclTensor** tensor)
53+{
54+ auto size = GetShapeSize(shape) * sizeof(T);
55+ // 调用aclrtMalloc申请device侧内存
56+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
57+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
58+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
59+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
61+ 
62+ // 计算连续tensor的strides
63+ std::vector<int64_t> strides(shape.size(), 1);
64+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
65+ strides[i] = shape[i + 1] * strides[i + 1];
66+ }
67+ 
68+ // 调用aclCreateTensor接口创建aclTensor
69+ *tensor = aclCreateTensor(
70+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
71+ *deviceAddr);
72+ return 0;
73+}
74+ 
75+int main()
76+{
77+ // 1. （固定写法）device/stream初始化，参考acl API手册
78+ // 根据自己的实际device填写deviceId
79+ int32_t deviceId = 0;
80+ aclrtStream stream;
81+ auto ret = Init(deviceId, &stream);
82+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
83+ 
84+ // 2. 构造输入与输出，需要根据API的接口自定义构造
85+ std::vector<int64_t> selfRefShape = {4, 2};
86+ void* selfRefDeviceAddr = nullptr;
87+ aclTensor* selfRef = nullptr;
88+ 
89+ std::vector<float> selfRefHostData = {0, 1, 2, 3, 4, 5, 6, 7};
90+ 
91+ // 创建selfRef aclTensor
92+ ret = CreateAclTensor(selfRefHostData, selfRefShape, &selfRefDeviceAddr, aclDataType::ACL_FLOAT, &selfRef);
93+ CHECK_RET(ret == ACL_SUCCESS, return ret);
94+ 
95+ // 3. 调用CANN算子库API
96+ uint64_t workspaceSize = 0;
97+ aclOpExecutor* executor;
98+ // 调用aclnnInplaceNeg第一段接口
99+ ret = aclnnInplaceNegGetWorkspaceSize(selfRef, &workspaceSize, &executor);
100+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceNegGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
101+ // 根据第一段接口计算出的workspaceSize申请device内存
102+ void* workspaceAddr = nullptr;
103+ if (workspaceSize > 0) {
104+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
105+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
106+ }
107+ // 调用aclnnInplaceNeg第二段接口
108+ ret = aclnnInplaceNeg(workspaceAddr, workspaceSize, executor, stream);
109+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceNeg failed. ERROR: %d\n", ret); return ret);
110+ 
111+ // 4. （固定写法）同步等待任务执行结束
112+ ret = aclrtSynchronizeStream(stream);
113+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
114+ 
115+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
116+ auto size = GetShapeSize(selfRefShape);
117+ std::vector<float> resultData(size, 0);
118+ ret = aclrtMemcpy(
119+ resultData.data(), resultData.size() * sizeof(resultData[0]), selfRefDeviceAddr, size * sizeof(resultData[0]),
120+ ACL_MEMCPY_DEVICE_TO_HOST);
121+ CHECK_RET(
122+ ret == ACL_SUCCESS, LOG_PRINT("copy resultData from device to host failed. ERROR: %d\n", ret); return ret);
123+ for (int64_t i = 0; i < size; i++) {
124+ LOG_PRINT("resultData[%ld] is: %f\n", i, resultData[i]);
125+ }
126+ 
127+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
128+ aclDestroyTensor(selfRef);
129+ 
130+ // 7. 释放device资源，需要根据具体API的接口定义修改
131+ aclrtFree(selfRefDeviceAddr);
132+ if (workspaceSize > 0) {
133+ aclrtFree(workspaceAddr);
134+ }
135+ aclrtDestroyStream(stream);
136+ aclrtResetDevice(deviceId);
137+ aclFinalize();
138+ return 0;
139+}
Aexperimental/math/neg/examples/test_aclnn_neg.cpp+138-0
@@ -0,0 +1,138 @@
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 <iostream>
12+#include <vector>
13+#include "acl/acl.h"
14+#include "aclnnop/aclnn_neg.h"
15+ 
16+#define CHECK_RET(cond, return_expr) \
17+ do { \
18+ if (!(cond)) { \
19+ return_expr; \
20+ } \
21+ } while (0)
22+ 
23+#define LOG_PRINT(message, ...) \
24+ do { \
25+ printf(message, ##__VA_ARGS__); \
26+ } while (0)
27+ 
28+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
29+ int64_t shapeSize = 1;
30+ for (auto i : shape) {
31+ shapeSize *= i;
32+ }
33+ return shapeSize;
34+}
35+ 
36+int Init(int32_t deviceId, aclrtStream* stream) {
37+ // 固定写法，资源初始化
38+ auto ret = aclInit(nullptr);
39+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
40+ ret = aclrtSetDevice(deviceId);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtCreateStream(stream);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
44+ return 0;
45+}
46+ 
47+template <typename T>
48+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
49+ aclDataType dataType, aclTensor** tensor) {
50+ auto size = GetShapeSize(shape) * sizeof(T);
51+ // 调用aclrtMalloc申请device侧内存
52+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
53+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
54+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
55+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
57+ 
58+ // 计算连续tensor的strides
59+ std::vector<int64_t> strides(shape.size(), 1);
60+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
61+ strides[i] = shape[i + 1] * strides[i + 1];
62+ }
63+ 
64+ // 调用aclCreateTensor接口创建aclTensor
65+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
66+ shape.data(), shape.size(), *deviceAddr);
67+ return 0;
68+}
69+ 
70+int main() {
71+ // 1. （固定写法）device/stream初始化，参考acl API手册
72+ // 根据自己的实际device填写deviceId
73+ int32_t deviceId = 0;
74+ aclrtStream stream;
75+ auto ret = Init(deviceId, &stream);
76+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
77+ 
78+ // 2. 构造输入与输出，需要根据API的接口自定义构造
79+ std::vector<int64_t> selfShape = {4, 2};
80+ std::vector<int64_t> outShape = {4, 2};
81+ void* selfDeviceAddr = nullptr;
82+ void* outDeviceAddr = nullptr;
83+ aclTensor* self = nullptr;
84+ aclTensor* out = nullptr;
85+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
86+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
87+ // 创建self aclTensor
88+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
89+ CHECK_RET(ret == ACL_SUCCESS, return ret);
90+ // 创建out aclTensor
91+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
92+ CHECK_RET(ret == ACL_SUCCESS, return ret);
93+ 
94+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
95+ uint64_t workspaceSize = 0;
96+ aclOpExecutor* executor;
97+ // 调用aclnnNeg第一段接口
98+ ret = aclnnNegGetWorkspaceSize(self, out, &workspaceSize, &executor);
99+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnNegGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
100+ // 根据第一段接口计算出的workspaceSize申请device内存
101+ void* workspaceAddr = nullptr;
102+ if (workspaceSize > 0) {
103+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
104+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
105+ }
106+ // 调用aclnnNeg第二段接口
107+ ret = aclnnNeg(workspaceAddr, workspaceSize, executor, stream);
108+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnNeg failed. ERROR: %d\n", ret); return ret);
109+ 
110+ // 4. （固定写法）同步等待任务执行结束
111+ ret = aclrtSynchronizeStream(stream);
112+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
113+ 
114+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
115+ auto size = GetShapeSize(outShape);
116+ std::vector<float> resultData(size, 0);
117+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
118+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
119+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
120+ for (int64_t i = 0; i < size; i++) {
121+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
122+ }
123+ 
124+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
125+ aclDestroyTensor(self);
126+ aclDestroyTensor(out);
127+ 
128+ // 7. 释放device资源，需要根据具体API的接口定义修改
129+ aclrtFree(selfDeviceAddr);
130+ aclrtFree(outDeviceAddr);
131+ if (workspaceSize > 0) {
132+ aclrtFree(workspaceAddr);
133+ }
134+ aclrtDestroyStream(stream);
135+ aclrtResetDevice(deviceId);
136+ aclFinalize();
137+ return 0;
138+}
Aexperimental/math/neg/op_api/aclnn_neg.cpp+154-0
@@ -0,0 +1,154 @@
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 "aclnn_neg.h"
12+#include "neg.h"
13+#include "aclnn_kernels/cast.h"
14+#include "aclnn_kernels/contiguous.h"
15+#include "aclnn_kernels/common/op_error_check.h"
16+#include "opdev/common_types.h"
17+#include "opdev/data_type_utils.h"
18+#include "opdev/format_utils.h"
19+#include "opdev/op_dfx.h"
20+#include "opdev/op_executor.h"
21+#include "opdev/op_log.h"
22+#include "opdev/shape_utils.h"
23+#include "opdev/tensor_view_utils.h"
24+#include "opdev/platform.h"
25+#include "platform/soc_spec.h"
26+ 
27+using namespace op;
28+#ifdef __cplusplus
29+extern "C" {
30+#endif
31+ 
32+static const std::initializer_list<op::DataType> dtype_support_list = {
33+ op::DataType::DT_INT8, op::DataType::DT_INT32, op::DataType::DT_INT64, op::DataType::DT_FLOAT16,
34+ op::DataType::DT_FLOAT, op::DataType::DT_DOUBLE, op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128};
35+ 
36+static const std::initializer_list<op::DataType> dtype_support_list_afterV200 = {
37+ op::DataType::DT_INT8, op::DataType::DT_INT32, op::DataType::DT_INT64,
38+ op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT, op::DataType::DT_DOUBLE,
39+ op::DataType::DT_COMPLEX64, op::DataType::DT_COMPLEX128, op::DataType::DT_BF16};
40+ 
41+static bool CheckNotNull(const aclTensor* self, const aclTensor* out)
42+{
43+ OP_CHECK_NULL(self, return false);
44+ OP_CHECK_NULL(out, return false);
45+ 
46+ return true;
47+}
48+ 
49+static bool CheckDtypeValid(const aclTensor* self)
50+{
51+ bool isAfterV200 =
52+ (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201 ||
53+ GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_3510);
54+ bool isSupport = isAfterV200 ? CheckType(self->GetDataType(), dtype_support_list_afterV200) :
55+ CheckType(self->GetDataType(), dtype_support_list);
56+ if (!isSupport) {
57+ OP_LOGE(
58+ ACLNN_ERR_PARAM_INVALID, "Self dtype %s should be in dtype support list %s.",
59+ ToString(self->GetDataType()).GetString(),
60+ isAfterV200 ? ToString(dtype_support_list_afterV200).GetString() :
61+ ToString(dtype_support_list).GetString());
62+ return false;
63+ }
64+ return true;
65+}
66+ 
67+static bool CheckPromoteType(const aclTensor* self, const aclTensor* out)
68+{
69+ OP_CHECK_RESULT_DTYPE_CAST_FAILED(self->GetDataType(), out->GetDataType(), return false);
70+ 
71+ return true;
72+}
73+ 
74+static bool CheckShape(const aclTensor* self, const aclTensor* out)
75+{
76+ const size_t MAX_DIM = 8;
77+ OP_CHECK_MAX_DIM(self, MAX_DIM, return false);
78+ OP_CHECK_SHAPE_NOT_EQUAL(self, out, return false);
CANN-robot
CANN-robotCANN-robot1月4日

逻辑运算与副作用: OP_CHECK_SHAPE_NOT_EQUAL 宏的名称 NOT_EQUAL 具有误导性。从上下文看，该宏应该是检查两个张量的形状是否相等（因为 Neg 算子要求输入和输出形状相同）。但宏名 NOT_EQUAL 暗示了检查“不相等”，这与实际功能相反，容易引起误解。

问题类型: 逻辑运算与副作用 文件路径: experimental/math/neg/op_host/op_api/aclnn_neg.cpp 行号: 79 问题代码:

    OP_CHECK_SHAPE_NOT_EQUAL(self, out, return false);

修改建议:

建议：1) 确认该宏的实际功能。如果确实是检查形状相等，应将宏名改为 `OP_CHECK_SHAPE_EQUAL` 或类似名称；2) 如果宏功能就是检查不相等（即允许形状不同），那么需要重新审视 `Neg` 算子的形状要求，因为通常取负操作要求输入输出形状一致。建议查阅框架文档或宏定义以确认其真实行为。

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

likedislike
79+ return true;
80+}
81+ 
82+static aclnnStatus CheckParams(const aclTensor* self, const aclTensor* out)
83+{
84+ // 1. 检查参数是否为空指针
85+ CHECK_RET(CheckNotNull(self, out), ACLNN_ERR_PARAM_NULLPTR);
86+ 
87+ // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
88+ CHECK_RET(CheckDtypeValid(self), ACLNN_ERR_PARAM_INVALID);
89+ 
90+ // 3. 检查self的数据类型能否转换为输出数据类型
91+ CHECK_RET(CheckPromoteType(self, out), ACLNN_ERR_PARAM_INVALID);
92+ 
93+ // 4. 检查shape
94+ CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID);
95+ 
96+ return ACLNN_SUCCESS;
97+}
98+ 
99+aclnnStatus aclnnNegGetWorkspaceSize(
100+ const aclTensor* self, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)
101+{
102+ L2_DFX_PHASE_1(aclnnNeg, DFX_IN(self), DFX_OUT(out));
103+ // 固定写法，参数检查
104+ auto ret = CheckParams(self, out);
105+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
106+ // 固定写法，创建OpExecutor
107+ auto uniqueExecutor = CREATE_EXECUTOR();
108+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
109+ // neg算子的空tensor在kernel中支持
110+ if (self->IsEmpty()) {
111+ *workspaceSize = 0;
112+ uniqueExecutor.ReleaseTo(executor);
113+ return ACLNN_SUCCESS;
114+ }
115+ // 固定写法，将输入self转换成连续的tensor
116+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
117+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
118+ // 调用Neg算子kernel
119+ auto negOpOut = l0op::Neg(selfContiguous, uniqueExecutor.get());
120+ CHECK_RET(negOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
121+ // 固定写法，将计算结果转换成输出out的数据类型
122+ auto castOut = l0op::Cast(negOpOut, out->GetDataType(), uniqueExecutor.get());
123+ CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
124+ // 固定写法，将计算结果拷贝到输出out上，out可能是非连续的tensor
125+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
126+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
127+ // 固定写法，获取计算过程中需要使用的workspace大小
128+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
129+ uniqueExecutor.ReleaseTo(executor);
130+ return ACLNN_SUCCESS;
131+}
132+ 
133+aclnnStatus aclnnInplaceNegGetWorkspaceSize(aclTensor* selfRef, uint64_t* workspaceSize, aclOpExecutor** executor)
134+{
135+ return aclnnNegGetWorkspaceSize(selfRef, selfRef, workspaceSize, executor);
136+}
137+ 
138+aclnnStatus aclnnNeg(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream)
139+{
140+ L2_DFX_PHASE_2(aclnnNeg);
141+ // 固定写法，调用框架能力，完成计算
142+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
143+}
144+ 
145+aclnnStatus aclnnInplaceNeg(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
146+{
147+ L2_DFX_PHASE_2(aclnnInplaceNeg);
148+ // 固定写法，调用框架能力，完成计算
149+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
150+}
151+ 
152+#ifdef __cplusplus
153+}
154+#endif
Aexperimental/math/neg/op_api/aclnn_neg.h+110-0
@@ -0,0 +1,110 @@
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+#ifndef OP_API_INC_NEG_H_
12+#define OP_API_INC_NEG_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+ 
17+#ifdef __cplusplus
18+extern "C" {
19+#endif
20+ 
21+/**
22+ * @brief aclnnNeg的第一段接口，根据具体的计算流程，计算workspace大小。
23+ * @domain aclnn_math
24+ *
25+ * 算子功能：对输入的每个元素完成相反数计算
26+ * 计算公式：
27+ * $$ output​=(-1) * self $$
28+ *
29+ * 实现说明：
30+ * api计算的基本路径：
31+ * ```mermaid
32+ * graph LR
33+ * A[(self)] -->B([Contiguous])
34+ * B -->D([Neg])
35+ * D-->E([Cast])
36+ * E-->F([ViewCopy])
37+ * F-->G[(out)]
38+ * ```
39+ *
40+ * @param [in] self: npu device侧的aclTensor，数据类型支持整型，浮点类型，支持非连续的Tensor，数据格式支持ND。
41+ * @param [in] out: npu device侧的aclTensor，数据类型支持整型，浮点类型，数据格式支持ND，且数据格式需要与self一致。
42+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
43+ * @param [out] executor: 返回op执行器，包含算子计算流程。
44+ * @return aclnnStatus: 返回状态码。
45+ */
46+ACLNN_API aclnnStatus
47+aclnnNegGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
48+ 
49+/**
50+ * @brief aclnnInplaceNeg的第一段接口，根据具体的计算流程，计算workspace大小。
51+ * @domain aclnn_math
52+ *
53+ * 算子功能：对输入的每个元素完成相反数计算
54+ * 计算公式：
55+ * $$ selfRef ​=(-1) * selfRef $$
56+ *
57+ * @param [in] selfRef: npu device侧的aclTensor，数据类型支持整型，浮点类型，支持非连续的Tensor，数据格式支持ND。
58+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
59+ * @param [out] executor: 返回op执行器，包含算子计算流程。
60+ * @return aclnnStatus: 返回状态码。
61+ */
62+ACLNN_API aclnnStatus
63+aclnnInplaceNegGetWorkspaceSize(aclTensor* selfRef, uint64_t* workspaceSize, aclOpExecutor** executor);
64+ 
65+/**
66+ * @brief aclnnNeg的第二段接口，用于执行计算。
67+ *
68+ * 算子功能：完成加法计算
69+ * 计算公式：
70+ * $$ output_i = self_i+alpha*other_i $$
71+ *
72+ * 实现说明：
73+ * api计算的基本路径：
74+ * ```mermaid
75+ * graph LR
76+ * A[(self)] -->B([Contiguous])
77+ * B -->D([Neg])
78+ * D-->E([Cast])
79+ * E-->F([ViewCopy])
80+ * F-->G[(out)]
81+ * ```
82+ *
83+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
84+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnNegGetWorkspaceSize获取。
85+ * @param [in] executor: op执行器，包含了算子计算流程。
86+ * @param [in] stream: acl stream流。
87+ * @return aclnnStatus: 返回状态码。
88+ */
89+ 
90+ACLNN_API aclnnStatus
91+aclnnNeg(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream);
92+ 
93+/**
94+ * @brief aclnnInplaceNeg的第二段接口，用于执行计算。
95+ *
96+ * 算子功能：对输入的每个元素完成相反数计算
97+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
98+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnInplaceNegGetWorkspaceSize获取。
99+ * @param [in] executor: op执行器，包含了算子计算流程。
100+ * @param [in] stream: acl stream流。
101+ * @return aclnnStatus: 返回状态码。
102+ */
103+ACLNN_API aclnnStatus
104+aclnnInplaceNeg(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
105+ 
106+#ifdef __cplusplus
107+}
108+#endif
109+ 
110+#endif /* OP_API_INC_NEG_H_ */
Aexperimental/math/neg/op_api/neg.cpp+80-0
@@ -0,0 +1,80 @@
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 "neg.h"
12+#include "opdev/aicpu/aicpu_task.h"
13+#include "opdev/op_log.h"
14+#include "opdev/op_executor.h"
15+#include "opdev/make_op_executor.h"
16+#include "opdev/op_def.h"
17+#include "opdev/op_dfx.h"
18+#include "aclnn_kernels/common/op_error_check.h"
19+ 
20+using namespace op;
21+ 
22+namespace l0op {
23+ 
24+OP_TYPE_REGISTER(Neg);
25+ 
26+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
27+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32, op::DataType::DT_INT8,
28+ op::DataType::DT_BF16};
29+ 
30+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST_AFTER_910B = {
31+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT32,
32+ op::DataType::DT_INT8, op::DataType::DT_BF16};
33+ 
34+// 根据芯片类型、dtype判断算子是否支持走aicore
35+static inline bool IsAiCoreSupport(DataType inputDtype)
36+{
37+ auto socVersion = GetCurrentPlatformInfo().GetCurNpuArch();
38+ if (socVersion == NpuArch::DAV_2201 || socVersion == NpuArch::DAV_3510)
39+ {
40+ return CheckType(inputDtype, AICORE_DTYPE_SUPPORT_LIST_AFTER_910B);
41+ }
42+ return CheckType(inputDtype, AICORE_DTYPE_SUPPORT_LIST);
43+}
44+ 
45+// AICORE算子kernel
46+static const aclTensor* NegAiCore(const aclTensor* self, aclTensor* negOut, aclOpExecutor* executor)
47+{
48+ L0_DFX(NegAiCore, self, negOut);
49+ // 使用框架宏ADD_TO_LAUNCHER_LIST_AICORE，将AiCore Neg算子加入任务队列
50+ // Neg是算子的OpType，self是算子的输入，negOut是算子的输出
51+ auto retAicore = ADD_TO_LAUNCHER_LIST_AICORE(Neg, OP_INPUT(self), OP_OUTPUT(negOut));
52+ OP_CHECK_ADD_TO_LAUNCHER_LIST_AICORE(
53+ retAicore != ACLNN_SUCCESS, return nullptr, "Neg ADD_TO_LAUNCHER_LIST_AICORE failed.");
54+ return negOut;
55+}
56+ 
57+// AICPU算子kernel
58+static const aclTensor* NegAiCpu(const aclTensor* self, aclTensor* negOut, aclOpExecutor* executor)
59+{
60+ // 使用框架宏ADD_TO_AICPU__LAUNCHER_LIST，将AiCpu Neg算子加入任务队列
61+ // Neg是算子的OpType，self是算子的输入，negOut是算子的输出
62+ L0_DFX(NegAiCpu, self, negOut);
63+ 
64+ static internal::AicpuTaskSpace space("Neg");
65+ auto ret = ADD_TO_LAUNCHER_LIST_AICPU(Neg, OP_ATTR_NAMES(), OP_INPUT(self), OP_OUTPUT(negOut));
66+ CHECK_RET(ret == ACLNN_SUCCESS, nullptr);
67+ return negOut;
68+}
69+ 
70+const aclTensor* Neg(const aclTensor* self, aclOpExecutor* executor)
71+{
72+ auto negOut = executor->AllocTensor(self->GetViewShape(), self->GetDataType());
73+ if (IsAiCoreSupport(self->GetDataType())) {
74+ return NegAiCore(self, negOut, executor);
75+ }
76+ 
77+ return NegAiCpu(self, negOut, executor);
78+} // namespace l0op
79+ 
80+} // namespace l0op
Aexperimental/math/neg/op_api/neg.h+20-0
@@ -0,0 +1,20 @@
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+#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_NEG_OP_H_
12+#define PTA_NPU_OP_API_INC_LEVEL0_OP_NEG_OP_H_
13+ 
14+#include "opdev/op_executor.h"
15+ 
16+namespace l0op {
17+const aclTensor* Neg(const aclTensor* self, aclOpExecutor* executor);
18+}
19+ 
20+#endif /* PTA_NPU_OP_API_INC_LEVEL0_OP_NEG_OP_H_ */
Aexperimental/math/neg/op_host/neg_def.cpp+37-0
@@ -0,0 +1,37 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file neg.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class Neg : public OpDef {
19+public:
20+ explicit Neg(const char* name) : OpDef(name)
CANN-robot
CANN-robotCANN-robot1月4日

代码格式/风格: 构造函数体内部缩进不一致。第23行开始的代码块使用了4个空格的缩进，但第21行的构造函数定义与第23行的代码块起始位置对齐，这不符合常见的C++编码规范（如Google C++ Style Guide建议使用2个空格缩进）。虽然这本身不是功能性问题，但会影响代码的可读性和维护一致性。

问题类型: 代码格式/风格 文件路径: experimental/math/neg/op_host/neg_def.cpp 行号: 21 问题代码:

    explicit Neg(const char* name) : OpDef(name)
    {
            this->Input("x")

修改建议:

统一代码块的缩进风格，建议使用2个或4个空格的统一缩进。例如将第23-32行的代码块调整为统一的缩进级别。

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

likedislike
21+ {
22+ this->Input("x")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
27+ this->Output("y")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
32+ 
33+ this->AICore().AddConfig("ascend910b");
34+ }
35+};
36+OP_ADD(Neg); // 添加算子信息库
37+} // namespace ops
Aexperimental/math/neg/op_host/neg_infershape.cpp+31-0
@@ -0,0 +1,31 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file neg_infer.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+ 
20+namespace ops {
21+ 
22+static ge::graphStatus InferShapeNeg(gert::InferShapeContext* context)
23+{
24+ const gert::Shape *x1_shape = context->GetInputShape(0);
25+ gert::Shape *y_shape = context->GetOutputShape(0);
26+ *y_shape = *x1_shape;
27+ return GRAPH_SUCCESS;
28+}
29+ 
30+IMPL_OP_INFERSHAPE(Neg).InferShape(InferShapeNeg);
31+} // namespace ops
Aexperimental/math/neg/op_host/neg_tiling.cpp+133-0
@@ -0,0 +1,133 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file neg_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "op_host/tiling_util.h"
19+#include "op_host/tiling_templates_registry.h"
20+#include "../op_kernel/neg_tiling_data.h"
21+#include "../op_kernel/neg_tiling_key.h"
22+ 
23+namespace optiling {
24+ 
25+const uint64_t BLOCK_SIZE = 32;
26+ 
27+struct NegCompileInfo {};
28+ 
29+// tiling 分发入口
30+static ge::graphStatus NegTilingFunc(gert::TilingContext* context)
31+{
32+ NegTilingData* tiling = context->GetTilingData<NegTilingData>();
33+ int64_t ubPartNum = 4;
34+ uint64_t dataTypeLength = 4;
35+ auto dt = context->GetInputDesc(0)->GetDataType();
36+ if (dt == ge::DT_INT8) {
37+ dataTypeLength = 1;
38+ ubPartNum += 4;
39+ }
40+ else if (dt == ge::DT_BF16)
41+ {
42+ dataTypeLength = 2;
43+ ubPartNum += 2;
44+ }
45+ else if (dt == ge::DT_FLOAT16)
46+ {
47+ dataTypeLength = 2;
48+ }
49+
50+ uint64_t ubLength = 0;
51+ uint64_t bigCoreDataNum = 0;
52+ uint64_t bigCoreLoopNum = 0;
53+ uint64_t bigCoreTailDataNum = 0;
54+
55+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
56+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubLength);
57+ auto coreNum = ascendcPlatform.GetCoreNum();
58+
59+ // Based on the input length and the number of inputs, the number of bytes of the input data type is obtained
60+ uint64_t inputDataNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
61+ uint64_t inputLength = inputDataNum * dataTypeLength;
62+ if (coreNum == 0 || BLOCK_SIZE == 0)
63+ {
F
Ffulltower3月27日

报错添加打印

likedislike
64+ OP_LOGE(context, "coreNum or BLOCK_SIZE is 0");
65+ return ge::GRAPH_FAILED;
66+ }
67+ uint64_t ubPartLength = ubLength / ubPartNum;
68+ // The number of 32B data blocks that can be used for each data. DOUBLE BUFFER is already counted here
69+ uint64_t ubPartBlockNum = ubPartLength / BLOCK_SIZE;
70+ uint64_t ubPartDataNum = (ubPartBlockNum * BLOCK_SIZE) / dataTypeLength;
71+ 
72+ // Input data for 32B alignment
73+ uint64_t inputLengthAlign32 = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE);
74+
75+ if(ubPartDataNum >= inputDataNum)
76+ {
77+ coreNum=1;
78+ }
79+ else
80+ {
81+ // There is at least 32B of data on each core, satisfying several settings for several cores. The maximum number of audits is the actual number of audits
82+ coreNum = (coreNum < inputLengthAlign32 / BLOCK_SIZE) ? coreNum : inputLengthAlign32 / BLOCK_SIZE;
CANN-robot
CANN-robotCANN-robot1月4日

逻辑运算与副作用: 在调整coreNum时，当inputLengthAlign32 / BLOCK_SIZE为0时，coreNum可能被设置为0。虽然第62行已经检查了coreNum == 0的情况，但这里重新赋值后可能再次变为0。

问题类型: 逻辑运算与副作用 文件路径: experimental/math/neg/op_host/neg_tiling.cpp 行号: 81 问题代码:

        coreNum = (coreNum <  inputLengthAlign32 / BLOCK_SIZE) ? coreNum : inputLengthAlign32 / BLOCK_SIZE;

修改建议:

建议在调整coreNum后再次检查其是否为0，或者确保inputLengthAlign32 / BLOCK_SIZE不为0。可以在调整后添加：if (coreNum == 0) coreNum = 1;

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

likedislike
83+ }
84+
85+ uint64_t everyCoreInputBlockNum = inputLengthAlign32 / BLOCK_SIZE / coreNum;
CANN-robot
CANN-robotCANN-robot1月4日

逻辑运算与副作用: 计算everyCoreInputBlockNum时存在除零风险。此时coreNum可能为0（虽然第62行检查过，但第81行可能将其重新赋值为0）。

问题类型: 逻辑运算与副作用 文件路径: experimental/math/neg/op_host/neg_tiling.cpp 行号: 84 问题代码:

    uint64_t everyCoreInputBlockNum = inputLengthAlign32 / BLOCK_SIZE / coreNum;

修改建议:

建议在计算前检查coreNum是否为0，或者确保在第81行调整后coreNum不为0。可以在计算前添加：if (coreNum == 0) return ge::GRAPH_FAILED;

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

likedislike
86+ uint64_t tailBlockNum = (inputLengthAlign32 / BLOCK_SIZE) % coreNum;
87+
88+ // Small chunks are calculated and sliced several times using the number of data on each core
89+ uint64_t smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / dataTypeLength;
90+ uint64_t smallCoreLoopNum = smallCoreDataNum / ubPartDataNum;
91+ smallCoreLoopNum = (everyCoreInputBlockNum % ubPartBlockNum) == 0 ? smallCoreLoopNum : smallCoreLoopNum + 1;
CANN-robot
CANN-robotCANN-robot1月4日

代码结构与可维护性: smallCoreLoopNum的计算逻辑复杂且重复。第89行计算了一次，第90行又根据条件调整。同样的逻辑在第99-100行又重复用于bigCoreLoopNum的计算。

问题类型: 代码结构与可维护性 文件路径: experimental/math/neg/op_host/neg_tiling.cpp 行号: 90 问题代码:

    smallCoreLoopNum = (everyCoreInputBlockNum % ubPartBlockNum) == 0 ? smallCoreLoopNum : smallCoreLoopNum + 1;

修改建议:

建议将循环次数的计算逻辑提取为一个独立的函数，提高代码可读性和可维护性，避免重复代码。

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

likedislike
92+ // Tail block calculation for small chunks of data
93+ uint64_t smallCoreTailDataNum = smallCoreDataNum - ubPartDataNum * (smallCoreLoopNum-1);
94+ smallCoreTailDataNum = smallCoreTailDataNum == 0 ? ubPartDataNum : smallCoreTailDataNum;
95+ 
96+ if(0 != tailBlockNum)
97+ {
98+ everyCoreInputBlockNum += 1;
99+ bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / dataTypeLength;
100+ bigCoreLoopNum = bigCoreDataNum / ubPartDataNum;
101+ bigCoreLoopNum = (everyCoreInputBlockNum % ubPartBlockNum) == 0 ? bigCoreLoopNum : bigCoreLoopNum + 1;
102+ bigCoreTailDataNum = bigCoreDataNum - ubPartDataNum * (bigCoreLoopNum-1);
103+ bigCoreTailDataNum = bigCoreTailDataNum == 0 ? ubPartDataNum : bigCoreTailDataNum;
104+ context->SetTilingKey(1);
105+ }
106+ else
107+ {
108+ context->SetTilingKey(0);
109+ }
110+
111+ tiling->smallCoreDataNum = smallCoreDataNum;
112+ tiling->bigCoreDataNum = bigCoreDataNum;
113+ tiling->ubPartDataNum = ubPartDataNum;
114+ tiling->smallCoreTailDataNum = smallCoreTailDataNum;
115+ tiling->bigCoreTailDataNum = bigCoreTailDataNum;
116+ tiling->smallCoreLoopNum = smallCoreLoopNum;
117+ tiling->bigCoreLoopNum = bigCoreLoopNum;
118+ tiling->tailBlockNum = tailBlockNum;
119+ 
120+ context->SetBlockDim(coreNum);
121+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
122+ currentWorkspace[0] = 0;
123+ return ge::GRAPH_SUCCESS;
124+}
125+ 
126+static ge::graphStatus TilingParseForNeg([[maybe_unused]] gert::TilingParseContext* context)
127+{
128+ return ge::GRAPH_SUCCESS;
129+}
130+ 
131+// tiling注册入口.
132+IMPL_OP_OPTILING(Neg).Tiling(NegTilingFunc).TilingParse<NegCompileInfo>(TilingParseForNeg);
133+} // namespace optiling
Aexperimental/math/neg/op_kernel/neg.cpp+52-0
@@ -0,0 +1,52 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file neg.cpp
13+ * \brief
14+ */
15+ 
16+#include "neg.h"
17+ 
18+using namespace NsNeg;
19+ 
20+template <uint32_t schMode>
21+__global__ __aicore__ void neg(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
22+{
23+ REGISTER_TILING_DEFAULT(NegTilingData);
24+ GET_TILING_DATA_WITH_STRUCT(NegTilingData, tiling_data, tiling);
25+
26+ if(TILING_KEY_IS(1))
27+ {
28+ TPipe pipe;
29+ KernelNeg<DTYPE_X, true> op;
30+ op.Init(x, y,
31+ tiling_data.smallCoreDataNum,
32+ tiling_data.bigCoreDataNum, tiling_data.bigCoreLoopNum,
33+ tiling_data.smallCoreLoopNum, tiling_data.ubPartDataNum,
34+ tiling_data.smallCoreTailDataNum, tiling_data.bigCoreTailDataNum,
35+ tiling_data.tailBlockNum,
36+ &pipe);
37+ op.Process();
38+ }
39+ else if(TILING_KEY_IS(0))
40+ {
41+ TPipe pipe;
42+ KernelNeg<DTYPE_X, false> op;
43+ op.Init(x, y,
44+ tiling_data.smallCoreDataNum,
45+ tiling_data.bigCoreDataNum, tiling_data.bigCoreLoopNum,
46+ tiling_data.smallCoreLoopNum, tiling_data.ubPartDataNum,
47+ tiling_data.smallCoreTailDataNum, tiling_data.bigCoreTailDataNum,
48+ tiling_data.tailBlockNum,
49+ &pipe);
50+ op.Process();
51+ }
52+}
Aexperimental/math/neg/op_kernel/neg.h+167-0
@@ -0,0 +1,167 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file neg.h
13+ * \brief
14+ */
15+#ifndef __NEG_H__
16+#define __NEG_H__
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "neg_tiling_data.h"
21+#include "neg_tiling_key.h"
22+ 
23+namespace NsNeg {
24+ 
25+using namespace AscendC;
26+ 
27+constexpr int32_t BUFFER_NUM = 2;
28+ 
29+template <typename T, bool IsExistBigCore>
30+class KernelNeg
31+{
32+public:
33+ __aicore__ inline KernelNeg() {}
34+ __aicore__ inline void Init(
35+ GM_ADDR src_gm,
36+ GM_ADDR dst_gm,
37+ uint32_t smallCoreDataNum,
38+ uint32_t bigCoreDataNum, uint32_t bigCoreLoopNum,
39+ uint32_t smallCoreLoopNum, uint32_t ubPartDataNum,
40+ uint32_t smallCoreTailDataNum, uint32_t bigCoreTailDataNum,
41+ uint32_t tailBlockNum,
42+ TPipe* pipeIn
43+ )
44+ {
45+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
46+ uint32_t coreNum = AscendC::GetBlockIdx();
47+ uint32_t globalBufferIndex = bigCoreDataNum * AscendC::GetBlockIdx();
48+ this->ubPartDataNum = ubPartDataNum;
49+ if constexpr (IsExistBigCore)
50+ {
51+ if (coreNum < tailBlockNum)
52+ {
53+ this->coreDataNum = bigCoreDataNum;
54+ this->tileNum = bigCoreLoopNum;
55+ this->tailDataNum = bigCoreTailDataNum;
56+ }
57+ else
58+ {
59+ this->coreDataNum = smallCoreDataNum;
60+ this->tileNum = smallCoreLoopNum;
61+ this->tailDataNum = smallCoreTailDataNum;
62+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (AscendC::GetBlockIdx() - tailBlockNum);
63+ }
64+ }
65+ else
66+ {
67+ this->coreDataNum = smallCoreDataNum;
68+ this->tileNum = smallCoreLoopNum;
69+ this->tailDataNum = smallCoreTailDataNum;
70+ globalBufferIndex = smallCoreDataNum * AscendC::GetBlockIdx();
71+ }
72+ src_global.SetGlobalBuffer((__gm__ T *)src_gm + globalBufferIndex, this->coreDataNum);
73+ dst_global.SetGlobalBuffer((__gm__ T *)dst_gm + globalBufferIndex, this->coreDataNum);
74+ pipe = pipeIn;
75+ pipe->InitBuffer(inQueueX, BUFFER_NUM, this->ubPartDataNum * sizeof(T));
76+ pipe->InitBuffer(outQueue, BUFFER_NUM, this->ubPartDataNum * sizeof(T));
77+ pipe->InitBuffer(QueueTmp, this->ubPartDataNum * sizeof(half));
78+ pipe->InitBuffer(QueueTmp2, this->ubPartDataNum * sizeof(half));
79+ pipe->InitBuffer(QueueTmp1, this->ubPartDataNum * sizeof(float));
80+ }
81+ 
82+ __aicore__ inline void Process()
83+ {
84+ int32_t loopCount = this->tileNum;
85+ this->processDataNum = this->ubPartDataNum;
86+ for (int32_t i = 0; i < loopCount-1; i++)
87+ {
88+ CopyIn(i);
89+ Compute(i);
90+ CopyOut(i);
91+ }
92+ this->processDataNum = this->tailDataNum;
93+ CopyIn(loopCount-1);
94+ Compute(loopCount-1);
95+ CopyOut(loopCount-1);
96+ }
97+ 
98+private:
99+ __aicore__ inline void CopyIn(uint32_t process)
100+ {
101+ LocalTensor<T> srcLocal = inQueueX.AllocTensor<T>();
102+ DataCopy(srcLocal, src_global[process * this->ubPartDataNum], this->processDataNum);
103+ inQueueX.EnQue(srcLocal);
104+ }
105+ __aicore__ inline void Compute(uint32_t process)
106+ {
107+ LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();
108+ LocalTensor<T> srcLocal = inQueueX.DeQue<T>();
109+ if constexpr (std::is_same_v<T, int32_t>){
110+ Duplicate(dstLocal, T(-1), this->processDataNum);
111+ Mul(dstLocal, srcLocal, dstLocal, this->processDataNum);
112+ }
113+ else if constexpr (std::is_same_v<T, int8_t>){
114+ LocalTensor<half> tmp = QueueTmp.Get<half>();
115+ Cast(tmp, srcLocal, RoundMode::CAST_NONE, this->processDataNum);
116+ Muls(tmp, tmp, half(-1), this->processDataNum);
117+ //移位操作实现溢出处理
118+ LocalTensor<int16_t> tmp2 = QueueTmp2.Get<int16_t>();
119+ Cast(tmp2, tmp, RoundMode::CAST_RINT, this->processDataNum); // float16 -> int16
120+ // 处理溢出 (模拟 int8 计算的行为)
121+ ShiftLeft(tmp2, tmp2, int16_t(8), this->processDataNum);
122+ ShiftRight(tmp2, tmp2, int16_t(8), this->processDataNum);
123+ // 转回 half
124+ Cast(tmp, tmp2, RoundMode::CAST_NONE, this->processDataNum);
125+ // 转回int8
126+ Cast(dstLocal, tmp, RoundMode::CAST_NONE, this->processDataNum);
127+ QueueTmp2.FreeTensor(tmp2);
128+ QueueTmp.FreeTensor(tmp);
129+ }
130+ else if constexpr (std::is_same_v<T, float> || std::is_same_v<T, half>) {
131+ Muls(dstLocal, srcLocal, T(-1), this->processDataNum);
132+ }
133+ //Muls不支持bfloat16类型
134+ else{
135+ LocalTensor<float> tmp1 = QueueTmp.Get<float>();
136+ Cast(tmp1, srcLocal, RoundMode::CAST_NONE, this->processDataNum);
137+ Muls(tmp1, tmp1, float(-1), this->processDataNum);
138+ Cast(dstLocal, tmp1, RoundMode::CAST_RINT, this->processDataNum);
139+ QueueTmp.FreeTensor(tmp1);
140+ }
141+ outQueue.EnQue<T>(dstLocal);
142+ inQueueX.FreeTensor(srcLocal);
143+ }
144+ 
145+ __aicore__ inline void CopyOut(uint32_t process)
146+ {
147+ LocalTensor<T> dstLocal = outQueue.DeQue<T>();
148+ DataCopy(dst_global[process * this->ubPartDataNum], dstLocal, this->processDataNum);
149+ outQueue.FreeTensor(dstLocal);
150+ }
151+ 
152+private:
153+ GlobalTensor<T> src_global;
154+ GlobalTensor<T> dst_global;
155+ TPipe* pipe;
156+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
157+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueue;
158+ TBuf<QuePosition::VECCALC> QueueTmp, QueueTmp2, QueueTmp1;
159+ uint32_t coreDataNum;
160+ uint32_t tileNum;
161+ uint32_t ubPartDataNum;
162+ uint32_t tailDataNum;
163+ uint32_t processDataNum;
164+};
165+ 
166+} // namespace NsNeg
167+#endif // NEG_H
Aexperimental/math/neg/op_kernel/neg_tiling_data.h+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file neg_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef __NEG_TILLING_DATA_H__
17+#define __NEG_TILLING_DATA_H__
18+ 
19+struct NegTilingData {
20+ uint64_t smallCoreDataNum;
21+ uint64_t bigCoreDataNum;
22+ uint64_t ubPartDataNum;
23+ uint64_t smallCoreTailDataNum;
24+ uint64_t bigCoreTailDataNum;
25+ uint64_t smallCoreLoopNum;
26+ uint64_t bigCoreLoopNum;
27+ uint64_t tailBlockNum;
28+};
29+#endif
Aexperimental/math/neg/op_kernel/neg_tiling_key.h+35-0
@@ -0,0 +1,35 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file neg_tiling_key.h
13+ * \brief neg tiling key declare
14+ */
15+ 
16+#ifndef __NEG_TILING_KEY_H__
17+#define __NEG_TILING_KEY_H__
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+/* Mode场景定义 */
22+#define ELEMENTWISE_TPL_SCH_MODE_0 0
23+#define ELEMENTWISE_TPL_SCH_MODE_1 1
24+/* 继续定义其他Mode场景... */
25+ 
26+/* 模板参数 */
27+ASCENDC_TPL_ARGS_DECL(
28+ Neg,
29+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
30+ 
31+/* 模板参数组合 */
32+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
33+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));
34+ 
35+#endif