Star904
1821
代码
代码
Issues224
Pull Requests298
Actions
项目讨论
Wiki
分析
Star904
1821
已合并
selu_grad、softsign算子支持Ascend950 ascendc实现 #5767
yulianjie创建于 6月5日
selu_grad、softsign算子支持Ascend950 ascendc实现 #5767
已合并
yulianjie创建于 6月5日
从lianjieyu/ops-nnselu_grad合入到cann/ops-nnmaster
共 36 个文件变更+2752-83
Mactivation/selu_grad/CMakeLists.txt+13-14
@@ -1,18 +1,17 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.1+# ----------------------------------------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 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# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# 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.
5-# 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,
6# 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.
7# 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.
8-#/9+# ----------------------------------------------------------------------------------------------------------
9-message(STATUS "=== Debug: start ops.activation.selu_grad.CMakeLists.txt ")10+# Generated By CANNBot
T
Ttangweiwei26月27日

这个为什么要呈现是CANNBot？

likedislike
Y
yulianjie
6月27日 评论：
10-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11+ 
11-if(NOT ENABLE_TEST AND NOT BENCHMARK)12+# 设置算子定义时支持的芯片类型
12- list(REMOVE_ITEM CURRENT_DIRS tests)13+set(SUPPORT_COMPUTE_UNIT "ascend950")
13-endif()14+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14-foreach(SUB_DIR ${CURRENT_DIRS})15+set(SUPPORT_TILING_DIR "arch35")
15- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")16+ 
16- add_subdirectory(${SUB_DIR})17+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE selu_grad ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
17- endif()
18-endforeach()
Mactivation/selu_grad/README.md+87-1
@@ -1,3 +1,89 @@
1# SeluGrad1# SeluGrad
2 2 
3-本目录仅包含SeluGrad算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | √ |
12+| <term>Atlas 训练系列产品</term> | √ |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：对输入Tensor计算SELU（Scaled Exponential Linear Unit）激活函数的反向梯度。
17+- 计算公式：
18+ 
19+ $$
20+ y = \begin{cases} \text{scale} \times \text{gradients}, & \text{outputs} \ge 0 \\ \text{gradients} \times (\text{outputs} + \text{scale} \times \alpha), & \text{outputs} < 0 \end{cases}
21+ $$
22+ 
23+ 其中：
24+ 
25+ - $\alpha = 1.6732632423543772848170429916717$
26+ - $\text{scale} = 1.0507009873554804934193349852946$
27+ - $\text{scale} \times \alpha = 1.7580993408473768599402175208123$
C
Cchenjiao6月25日

1.6732632423543772848170429916717-----公式里这些魔鬼数字是什么？

likedislike
Y
yulianjie
6月25日 评论：
28+ 
C
Cchenqi3176月26日

md 建议写注明公式，方便社区阅读

likedislike
29+ 分段行为：
30+ 
31+ - 当 $\text{outputs} \ge 0$ 时：$y = \text{scale} \times \text{gradients}$（线性区梯度）
32+ - 当 $\text{outputs} < 0$ 时：$y = \text{gradients} \times (\text{outputs} + \text{scale} \times \alpha)$（指数饱和区梯度）
33+ 
34+## 参数说明
35+ 
36+<table style="table-layout: fixed; width: 1576px"><colgroup>
37+<col style="width: 170px">
38+<col style="width: 170px">
39+<col style="width: 200px">
40+<col style="width: 200px">
41+<col style="width: 170px">
42+<col style="width: 170px">
43+</colgroup>
44+<thead>
45+ <tr>
46+ <th>参数名</th>
47+ <th>输入/输出</th>
48+ <th>描述</th>
49+ <th>数据类型</th>
50+ <th>数据格式</th>
51+ <th>维度(shape)</th>
52+ </tr></thead>
53+<tbody>
54+ <tr>
55+ <td>gradients</td>
56+ <td>输入</td>
57+ <td>反向传播上游梯度。</td>
58+ <td>FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8</td>
59+ <td>ND</td>
60+ <td>1-8</td>
61+ </tr>
62+ <tr>
63+ <td>outputs</td>
64+ <td>输入</td>
65+ <td>SELU前向输出，shape与gradients支持numpy广播。</td>
66+ <td>FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8</td>
67+ <td>ND</td>
68+ <td>1-8</td>
69+ </tr>
70+ <tr>
71+ <td>y</td>
72+ <td>输出</td>
73+ <td>反向梯度结果，shape为gradients与outputs广播后的shape。</td>
74+ <td>FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8</td>
75+ <td>ND</td>
76+ <td>1-8</td>
77+ </tr>
78+</tbody></table>
79+ 
80+## 约束说明
81+ 
82+- 支持numpy广播：gradients和outputs的shape可以不同，输出y的shape为两者广播后的结果。
83+- 确定性计算：SeluGrad默认确定性实现。
84+ 
85+## 调用说明
86+ 
87+| 调用方式 | 调用样例 | 说明 |
88+|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
89+| aclnn调用 | [test_aclnn_selu_grad](./examples/arch35/test_aclnn_selu_grad.cpp) | 通过[aclnnSeluBackward](./docs/aclnnSeluBackward.md)接口方式调用SeluGrad算子。 |
Mactivation/selu_grad/docs/aclnnSeluBackward.md+1-1
@@ -6,7 +6,7 @@
6 6 
7|产品 | 是否支持 |7|产品 | 是否支持 |
8|:-------------------------|:----------:|8|:-------------------------|:----------:|
9-| <term>Ascend 950PR/Ascend 950DT</term> | × |9+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
11| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |11| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
12| <term>Atlas 200I/500 A2 推理产品</term> | × |12| <term>Atlas 200I/500 A2 推理产品</term> | × |
Ractivation/selu_grad/examples/test_aclnn_selu_grad.cpp→activation/selu_grad/examples/arch35/test_aclnn_selu_grad.cpp+62-56
@@ -9,19 +9,20 @@
9 */9 */
10#include <iostream>10#include <iostream>
11#include <vector>11#include <vector>
12+#include <cmath>
12#include "acl/acl.h"13#include "acl/acl.h"
13-#include "aclnnop/aclnn_selu_backward.h"14+#include "aclnn_selu_backward.h"
14 15 
15#define CHECK_RET(cond, return_expr) \16#define CHECK_RET(cond, return_expr) \
16 do { \17 do { \
17- if (!(cond)) { \18+ if (!(cond)) { \
18- return_expr; \19+ return_expr; \
19- } \20+ } \
20 } while(0)21 } while(0)
21 22 
22#define LOG_PRINT(message, ...) \23#define LOG_PRINT(message, ...) \
23 do { \24 do { \
24- printf(message, ##__VA_ARGS__); \25+ printf(message, ##__VA_ARGS__); \
25 } while(0)26 } while(0)
26 27 
27int64_t GetShapeSize(const std::vector<int64_t>& shape) {28int64_t GetShapeSize(const std::vector<int64_t>& shape) {
@@ -33,7 +34,6 @@ int64_t GetShapeSize(const std::vector<int64_t>& shape) {
33}34}
34 35 
35int Init(int32_t deviceId, aclrtStream* stream) {36int Init(int32_t deviceId, aclrtStream* stream) {
36- // 固定写法,资源初始化
37 auto ret = aclInit(nullptr);37 auto ret = aclInit(nullptr);
38 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);38 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
39 ret = aclrtSetDevice(deviceId);39 ret = aclrtSetDevice(deviceId);
@@ -47,96 +47,102 @@ template<typename T>
47int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,47int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
48 aclDataType dataType, aclTensor** tensor) {48 aclDataType dataType, aclTensor** tensor) {
49 auto size = GetShapeSize(shape) * sizeof(T);49 auto size = GetShapeSize(shape) * sizeof(T);
50- // 调用aclrtMalloc申请device侧引擎
51 auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);50 auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
52 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);51 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
53- 
54- // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
55 ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);52 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);53 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);54 std::vector<int64_t> strides(shape.size(), 1);
60 for (int64_t i = shape.size() - 2; i >= 0; i--) {55 for (int64_t i = shape.size() - 2; i >= 0; i--) {
61 strides[i] = shape[i + 1] * strides[i + 1];56 strides[i] = shape[i + 1] * strides[i + 1];
62 }57 }
63- 
64- // 调用aclCreateTensor接口创建aclTensor
65 *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,58 *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
66 shape.data(), shape.size(), *deviceAddr);59 shape.data(), shape.size(), *deviceAddr);
67 return 0;60 return 0;
68}61}
69 62 
70int main() {63int main() {
71- // 1. （固定写法）device/stream初始化, 参考acl API手册
72- // 根据自己的实际device填写deviceId
73 int32_t deviceId = 0;64 int32_t deviceId = 0;
74 aclrtStream stream;65 aclrtStream stream;
75 auto ret = Init(deviceId, &stream);66 auto ret = Init(deviceId, &stream);
76 CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);67 CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
77- // 2. 构造输入与输出，需要根据API的接口自定义构造
78- std::vector<int64_t> selfShape = {4, 2};
79- std::vector<int64_t> gradOutputShape = {4, 2};
80- std::vector<int64_t> gradInputShape = {4, 2};
81- void* selfDeviceAddr = nullptr;
82- void* gradOutputDeviceAddr = nullptr;
83- void* gradInputDeviceAddr = nullptr;
84- aclTensor* self = nullptr;
85- aclTensor* gradOutput = nullptr;
86- aclTensor* gradInput = nullptr;
87- std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
88- std::vector<float> gradOutputHostData = {1, 1, 1, 1, 1, 1, 1, 1};
89- std::vector<float> gradInputHostData = {0, 0, 0, 0, 0, 0, 0, 0};
90 68 
91- ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);69+ std::vector<int64_t> shape = {4, 2};
70+ void* gradDeviceAddr = nullptr;
71+ void* outDeviceAddr = nullptr;
72+ void* yDeviceAddr = nullptr;
73+ aclTensor* gradients = nullptr;
74+ aclTensor* outputs = nullptr;
75+ aclTensor* y = nullptr;
76+ 
77+ // SELU 常量
78+ const float SCALE = 1.0507009873554804f;
79+ const float ALPHA = 1.6732632423543772f;
80+ const float SCALE_ALPHA_PRODUCT = SCALE * ALPHA;
81+ 
82+ // 构造输入: gradients = 全1, outputs = [-2, -1, 0, 1, 2, 3, -0.5, 0.5]
83+ std::vector<float> gradHostData = {1, 1, 1, 1, 1, 1, 1, 1};
84+ std::vector<float> outHostData = {-2, -1, 0, 1, 2, 3, -0.5, 0.5};
85+ std::vector<float> yHostData(8, 0);
86+ 
87+ ret = CreateAclTensor(gradHostData, shape, &gradDeviceAddr, aclDataType::ACL_FLOAT, &gradients);
92 CHECK_RET(ret == ACL_SUCCESS, return ret);88 CHECK_RET(ret == ACL_SUCCESS, return ret);
93- ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT, &gradOutput);89+ ret = CreateAclTensor(outHostData, shape, &outDeviceAddr, aclDataType::ACL_FLOAT, &outputs);
94 CHECK_RET(ret == ACL_SUCCESS, return ret);90 CHECK_RET(ret == ACL_SUCCESS, return ret);
95- ret = CreateAclTensor(gradInputHostData, gradInputShape, &gradInputDeviceAddr, aclDataType::ACL_FLOAT, &gradInput);91+ ret = CreateAclTensor(yHostData, shape, &yDeviceAddr, aclDataType::ACL_FLOAT, &y);
96 CHECK_RET(ret == ACL_SUCCESS, return ret);92 CHECK_RET(ret == ACL_SUCCESS, return ret);
97 93 
98- // 3. 调用CANN算子库API，需要修改为具体的API
99 uint64_t workspaceSize = 0;94 uint64_t workspaceSize = 0;
100 aclOpExecutor* executor;95 aclOpExecutor* executor;
101- // 调用aclnnSeluBackward第一段接口96+ ret = aclnnSeluBackwardGetWorkspaceSize(gradients, outputs, y, &workspaceSize, &executor);
102- ret = aclnnSeluBackwardGetWorkspaceSize(gradOutput, self, gradInput, &workspaceSize, &executor);
103 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSeluBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);97 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSeluBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
104 98 
105- // 根据第一段接口计算出的workspaceSize申请device内存
106 void* workspaceAddr = nullptr;99 void* workspaceAddr = nullptr;
107 if (workspaceSize > 0) {100 if (workspaceSize > 0) {
108 ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);101 ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
109 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);102 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
110 }103 }
111- // 调用aclnnSeluBackward第二段接口
112 ret = aclnnSeluBackward(workspaceAddr, workspaceSize, executor, stream);104 ret = aclnnSeluBackward(workspaceAddr, workspaceSize, executor, stream);
113 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSeluBackward failed. ERROR: %d\n", ret); return ret);105 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSeluBackward failed. ERROR: %d\n", ret); return ret);
114- // 4. （固定写法）同步等待任务执行结束106+ 
115 ret = aclrtSynchronizeStream(stream);107 ret = aclrtSynchronizeStream(stream);
116 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);108 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
117- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
118- auto size = GetShapeSize(gradInputShape);
119- std::vector<float> resultData(size, 0);
120- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), gradInputDeviceAddr, size * sizeof(float),
121- ACL_MEMCPY_DEVICE_TO_HOST);
122- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
123- for (int64_t i = 0; i < size; i++) {
124- LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
125- }
126- // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
127- aclDestroyTensor(gradOutput);
128- aclDestroyTensor(self);
129- aclDestroyTensor(gradInput);
130 109 
131- // 7. 释放device资源，需要根据具体API的接口定义修改110+ auto size = GetShapeSize(shape);
132- aclrtFree(selfDeviceAddr);111+ std::vector<float> resultData(size, 0);
133- aclrtFree(gradOutputDeviceAddr);112+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(float), yDeviceAddr, size * sizeof(float),
134- aclrtFree(gradInputDeviceAddr);113+ ACL_MEMCPY_DEVICE_TO_HOST);
114+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result failed. ERROR: %d\n", ret); return ret);
115+ 
116+ LOG_PRINT("\n=== SeluGrad Results ===\n");
117+ LOG_PRINT("outputs >= 0: y = SCALE * gradients = %.6f * grad\n", SCALE);
118+ LOG_PRINT("outputs < 0: y = grad * (outputs + SCALE_ALPHA) = grad * (out + %.6f)\n\n", SCALE_ALPHA_PRODUCT);
119+ 
120+ int pass = 0, fail = 0;
121+ for (int64_t i = 0; i < size; i++) {
122+ float expected;
123+ if (outHostData[i] >= 0) {
124+ expected = SCALE * gradHostData[i];
125+ } else {
126+ expected = gradHostData[i] * (outHostData[i] + SCALE_ALPHA_PRODUCT);
127+ }
128+ bool ok = std::fabs(resultData[i] - expected) < 0.01f;
129+ if (ok) pass++; else fail++;
130+ LOG_PRINT(" [%ld] out=%.2f grad=%.2f => NPU=%.6f expected=%.6f %s\n",
131+ i, outHostData[i], gradHostData[i], resultData[i], expected, ok ? "PASS" : "FAIL");
132+ }
133+ LOG_PRINT("\nTotal: %d PASS, %d FAIL\n", pass, fail);
134+ 
135+ aclDestroyTensor(gradients);
136+ aclDestroyTensor(outputs);
137+ aclDestroyTensor(y);
138+ aclrtFree(gradDeviceAddr);
139+ aclrtFree(outDeviceAddr);
140+ aclrtFree(yDeviceAddr);
135 if (workspaceSize > 0) {141 if (workspaceSize > 0) {
136 aclrtFree(workspaceAddr);142 aclrtFree(workspaceAddr);
137 }143 }
138 aclrtDestroyStream(stream);144 aclrtDestroyStream(stream);
139 aclrtResetDevice(deviceId);145 aclrtResetDevice(deviceId);
140 aclFinalize();146 aclFinalize();
141- return 0;147+ return fail > 0 ? 1 : 0;
142-}148+}
Aactivation/selu_grad/examples/arch35/test_geir_selu_grad.cpp+225-0
@@ -0,0 +1,225 @@
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 <fstream>
13+#include <string.h>
14+#include <stdint.h>
15+#include <vector>
16+#include <string>
17+#include <map>
18+#include "assert.h"
19+ 
20+#include "graph.h"
21+#include "types.h"
22+#include "tensor.h"
23+#include "ge_error_codes.h"
24+#include "ge_api_types.h"
25+#include "ge_api.h"
26+#include "array_ops.h"
27+#include "ge_ir_build.h"
28+#include "../op_graph/selu_grad_proto.h"
29+ 
30+#define FAILED -1
31+#define SUCCESS 0
32+ 
33+using namespace ge;
34+using std::map;
35+using std::string;
36+using std::vector;
37+ 
38+#define ADD_INPUT_GRAD(intputIndex, intputName, intputDtype, inputShape) \
39+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
40+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
41+ TensorDesc placeholder##intputIndex##_desc = \
42+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
43+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
44+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
45+ Tensor tensor_placeholder##intputIndex; \
46+ ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, \
47+ tensor_placeholder##intputIndex, \
48+ placeholder##intputIndex##_desc, \
49+ 1); \
50+ if (ret != SUCCESS) { \
51+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
52+ return FAILED; \
53+ } \
54+ placeholder##intputIndex.update_input_desc_gradients(placeholder##intputIndex##_desc); \
55+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
56+ input.push_back(tensor_placeholder##intputIndex); \
57+ graph.AddOp(placeholder##intputIndex); \
58+ selu_grad_op.set_input_##intputName(placeholder##intputIndex); \
59+ inputs.push_back(placeholder##intputIndex);
60+ 
61+#define ADD_INPUT_OUT(intputIndex, intputName, intputDtype, inputShape) \
62+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
63+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
64+ TensorDesc placeholder##intputIndex##_desc = \
65+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
66+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
67+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
68+ Tensor tensor_placeholder##intputIndex; \
69+ ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, \
70+ tensor_placeholder##intputIndex, \
71+ placeholder##intputIndex##_desc, \
72+ 2); \
73+ if (ret != SUCCESS) { \
74+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
75+ return FAILED; \
76+ } \
77+ placeholder##intputIndex.update_input_desc_outputs(placeholder##intputIndex##_desc); \
78+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
79+ input.push_back(tensor_placeholder##intputIndex); \
80+ graph.AddOp(placeholder##intputIndex); \
81+ selu_grad_op.set_input_##intputName(placeholder##intputIndex); \
82+ inputs.push_back(placeholder##intputIndex);
83+ 
84+#define LOG_PRINT(message, ...) \
85+ do { \
86+ printf(message, ##__VA_ARGS__); \
87+ } while (0)
88+ 
89+string GetTime()
90+{
91+ time_t timep;
92+ time(&timep);
93+ char tmp[64];
94+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
95+ return tmp;
96+}
97+ 
98+uint32_t GetDataTypeSize(DataType dt)
99+{
100+ if (dt == ge::DT_FLOAT) return 4;
101+ if (dt == ge::DT_FLOAT16) return 2;
102+ if (dt == ge::DT_BF16) return 2;
103+ if (dt == ge::DT_INT32) return 4;
104+ if (dt == ge::DT_INT64) return 8;
105+ if (dt == ge::DT_INT8) return 1;
106+ return 4;
107+}
108+ 
109+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, float value)
110+{
111+ input_tensor_desc.SetRealDimCnt(shapes.size());
112+ size_t size = 1;
113+ for (uint32_t i = 0; i < shapes.size(); i++) {
114+ size *= shapes[i];
115+ }
116+ uint32_t data_len = size * 4;
117+ float *pData = new (std::nothrow) float[size];
118+ if (pData == nullptr) {
119+ return FAILED;
120+ }
121+ for (size_t i = 0; i < size; ++i) {
122+ *(pData + i) = value;
123+ }
124+ input_tensor = Tensor(input_tensor_desc, (uint8_t *)pData, data_len);
125+ delete[] pData;
126+ return SUCCESS;
127+}
128+ 
129+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t *inputData)
130+{
131+ FILE *fp = fopen(bin_file.c_str(), "w");
132+ if (fp == nullptr) return FAILED;
133+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
134+ fclose(fp);
135+ return SUCCESS;
136+}
137+ 
138+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
139+ std::vector<Operator> &outputs, Graph &graph)
140+{
141+ Status ret = SUCCESS;
142+ auto selu_grad_op = op::SeluGrad("selu_grad");
143+ std::vector<int64_t> shape = {4, 4};
144+ ADD_INPUT_GRAD(1, gradients, inDtype, shape);
145+ ADD_INPUT_OUT(2, outputs, inDtype, shape);
146+ 
147+ outputs.push_back(selu_grad_op);
148+ return SUCCESS;
149+}
150+ 
151+int main(int argc, char *argv[])
152+{
153+ const char *graph_name = "tc_ge_irrun_test";
154+ Graph graph(graph_name);
155+ std::vector<ge::Tensor> input;
156+ 
157+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
158+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
159+ Status ret = ge::GEInitialize(global_options);
160+ if (ret != SUCCESS) {
161+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
162+ return FAILED;
163+ }
164+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
165+ 
166+ std::vector<Operator> inputs{};
167+ std::vector<Operator> outputs{};
168+ 
169+ DataType inDtype = DT_FLOAT;
170+ 
171+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
172+ if (ret != SUCCESS) {
173+ printf("%s - ERROR - [XIR]: Create ir session failed\n", GetTime().c_str());
174+ return FAILED;
175+ }
176+ 
177+ if (!inputs.empty() && !outputs.empty()) {
178+ graph.SetInputs(inputs).SetOutputs(outputs);
179+ }
180+ 
181+ std::map<AscendString, AscendString> build_options = {};
182+ ge::Session *session = new Session(build_options);
183+ if (session == nullptr) {
184+ printf("%s - ERROR - [XIR]: Create session failed\n", GetTime().c_str());
185+ return FAILED;
186+ }
187+ 
188+ std::map<AscendString, AscendString> graph_options = {};
189+ uint32_t graph_id = 0;
190+ ret = session->AddGraph(graph_id, graph, graph_options);
191+ 
192+ std::vector<ge::Tensor> output;
193+ ret = session->RunGraph(graph_id, input, output);
194+ if (ret != SUCCESS) {
195+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
196+ delete session;
197+ GEFinalize();
198+ return FAILED;
199+ }
200+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
201+ 
202+ int output_num = output.size();
203+ for (int i = 0; i < output_num; i++) {
204+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
205+ string output_file = "./tc_ge_irrun_test_selu_grad_output_" + std::to_string(i) + ".bin";
206+ uint8_t *output_data_i = output[i].GetData();
207+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
208+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
209+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
210+ float *resultData = (float*)output_data_i;
211+ for (int64_t j = 0; j < output_shape; j++) {
212+ LOG_PRINT("result[%ld] is: %f\n", j, resultData[j]);
213+ }
214+ }
215+ 
216+ delete session;
217+ session = nullptr;
218+ ret = ge::GEFinalize();
219+ if (ret != SUCCESS) {
220+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
221+ return FAILED;
222+ }
223+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
224+ return SUCCESS;
225+}
Ractivation/selu_grad/op_host/op_api/aclnn_selu_backward.cpp→activation/selu_grad/op_api/aclnn_selu_backward.cpp+0-0
文件重命名但无更改。
Ractivation/selu_grad/op_host/op_api/aclnn_selu_backward.h→activation/selu_grad/op_api/aclnn_selu_backward.h+0-0
文件重命名但无更改。
Ractivation/selu_grad/op_host/op_api/selugrad.cpp→activation/selu_grad/op_api/selugrad.cpp+0-0
文件重命名但无更改。
Ractivation/selu_grad/op_host/op_api/selugrad.h→activation/selu_grad/op_api/selugrad.h+0-0
文件重命名但无更改。
Ractivation/selu_grad/op_host/CMakeLists.txt→activation/selu_grad/op_graph/CMakeLists.txt+6-4
@@ -1,10 +1,12 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.1+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").3# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.4# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.7# See LICENSE in the root of the software repository for the full text of the License.
8#/8#/
9+# Generated By CANNBot
9 10 
10-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE selu_grad ACLNNTYPE aclnn)11+message(STATUS "=== Debug: start ops.activation.selu_grad.op_graph.CMakeLists.txt ")
12+add_graph_plugin_sources()
Aactivation/selu_grad/op_graph/selu_grad_proto.h+47-0
@@ -0,0 +1,47 @@
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 selu_grad_proto.h
13+ * \brief
14+ */
15+#ifndef SELU_GRAD_PROTO_H_
16+#define SELU_GRAD_PROTO_H_
17+ 
18+#include "graph/operator_reg.h"
19+ 
20+namespace ge {
21+/**
22+* @brief Computes the gradient of SELU:
23+* y = scale * gradients if outputs >= 0
24+* y = gradients * (outputs + scale * alpha) if outputs < 0
25+*
26+* where alpha = 1.6732632423543772848170429916717
27+* scale = 1.0507009873554804934193349852946
28+*
29+* @par Inputs:
30+* Two inputs:
31+* gradients: A Tensor. Support 1D ~ 8D. Must be one of the following types: float16, float,
32+* bfloat16, int32, int8, uint8. format:ND.
T
Ttang-lei016月25日

RealNumberTypr定义的数据类型列表不一致，请参考CANNdev的方式

likedislike
Y
yulianjie
6月25日 评论：
33+* outputs: A Tensor. Has the same type, shape and format as "gradients".
34+*
35+* @par Outputs:
36+* y: A Tensor. Has the same type, shape and format as "gradients".
37+*
38+* @par Third-party framework compatibility
39+* @li Compatible with the Pytorch operator selu_backward.
40+*/
41+REG_OP(SeluGrad)
42+ .INPUT(gradients, TensorType::RealNumberType())
43+ .INPUT(outputs, TensorType::RealNumberType())
44+ .OUTPUT(y, TensorType::RealNumberType())
45+ .OP_END_FACTORY_REG(SeluGrad)
46+} // namespace ge
47+#endif // SELU_GRAD_PROTO_H_
Aactivation/selu_grad/op_host/arch35/selu_grad_tiling_arch35.cpp+405-0
@@ -0,0 +1,405 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file selu_grad_tiling.cpp
15+ * \brief SeluGrad 算子 Tiling 实现（arch35 架构）
16+ *
17+ * TilingKey_0 (OneDim) + TilingKey_1 (Broadcast)
18+ * 支持 float16, float32, bfloat16
19+ *
20+ * TilingKey 判定：
21+ * - 合轴后仅 1 维（shape 完全一致或标量广播）→ TilingKey_0 (OneDim)
22+ * - 合轴后 > 1 维（需要多维广播）→ TilingKey_1 (Broadcast)
23+ */
24+ 
25+#include "register/op_def_registry.h"
26+#include "op_common/log/log.h"
27+#include "op_common/op_host/util/math_util.h"
28+#include "op_common/op_host/util/platform_util.h"
29+#include "../../op_kernel/arch35/selu_grad_tiling_data.h"
30+#include "../../op_kernel/arch35/selu_grad_tiling_key.h"
31+ 
32+namespace optiling {
33+ 
34+using Ops::Base::CeilDiv;
35+using Ops::Base::CeilAlign;
36+using Ops::Base::FloorDiv;
37+using Ops::Base::FloorAlign;
38+using Ops::Base::GetUbBlockSize;
39+ 
40+constexpr uint32_t WS_SYS_SIZE = 0U;
41+constexpr size_t WORKSPACE_NUM = 1;
42+constexpr int32_t MAX_RANK = 8;
43+ 
44+// Buffer 数量常量
45+constexpr int64_t SELECT_UB_RESERVE = 8192; // Select 8K 预留
46+constexpr int64_t DEFAULT_BYTES_PER_ELEM = 40; // 默认每元素 UB 字节数（兜底）
47+ 
48+static const gert::Shape g_vec_1_shape = {1};
49+ 
50+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape)
51+{
52+ if (in_shape.GetDimNum() == 0) {
53+ return g_vec_1_shape;
54+ }
55+ return in_shape;
56+}
57+ 
58+// 获取平台信息
59+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)
60+{
61+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
62+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
63+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
64+ *coreNum = ascendcPlatform.GetCoreNumAiv();
65+ OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
66+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize);
67+ OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
68+ return ge::GRAPH_SUCCESS;
69+}
70+ 
71+// 获取 shape、dtype 信息（含广播推导）
72+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context,
73+ int64_t* totalElements,
74+ ge::DataType* dataType,
75+ gert::Shape* gradShapeOut,
76+ gert::Shape* outShapeOut,
77+ gert::Shape* yShapeOut)
78+{
79+ auto inputGrad = context->GetInputShape(0);
80+ OP_CHECK_NULL_WITH_CONTEXT(context, inputGrad);
81+ auto gradShape = inputGrad->GetStorageShape();
82+ 
83+ auto inputOut = context->GetInputShape(1);
84+ OP_CHECK_NULL_WITH_CONTEXT(context, inputOut);
85+ auto outShape = inputOut->GetStorageShape();
86+ 
87+ auto outputY = context->GetOutputShape(0);
88+ OP_CHECK_NULL_WITH_CONTEXT(context, outputY);
89+ auto yShape = outputY->GetStorageShape();
90+ 
91+ *gradShapeOut = gradShape;
92+ *outShapeOut = outShape;
93+ *yShapeOut = yShape;
94+ 
95+ // 计算输出总元素数
96+ if (yShape.GetDimNum() == 0) {
97+ *totalElements = 1; // rank=0 标量
98+ } else {
99+ *totalElements = yShape.GetShapeSize();
100+ }
101+ 
102+ auto inputDesc = context->GetInputDesc(0);
103+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
104+ *dataType = inputDesc->GetDataType();
105+ return ge::GRAPH_SUCCESS;
106+}
107+ 
108+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
109+{
110+ size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM);
111+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
112+ currentWorkspace[0] = WS_SYS_SIZE;
113+ return ge::GRAPH_SUCCESS;
114+}
115+ 
116+// 判断是否需要广播（shape 不完全一致）
117+static bool NeedsBroadcast(const gert::Shape& gradShape, const gert::Shape& outShape)
118+{
119+ if (gradShape.GetDimNum() != outShape.GetDimNum()) {
120+ return true;
121+ }
122+ for (size_t i = 0; i < gradShape.GetDimNum(); i++) {
123+ if (gradShape.GetDim(i) != outShape.GetDim(i)) {
124+ return true;
125+ }
126+ }
127+ return false;
128+}
129+ 
130+// 计算广播 stride（广播轴 stride=0）
131+static void ComputeBroadcastStrides(const gert::Shape& yShape,
132+ const gert::Shape& inputShape,
133+ int64_t* strides)
134+{
135+ int64_t yRank = static_cast<int64_t>(yShape.GetDimNum());
136+ int64_t inputRank = static_cast<int64_t>(inputShape.GetDimNum());
137+ 
138+ // 先计算 input 自身的 stride（从右到左累积）
139+ int64_t inputStrides[MAX_RANK] = {0};
140+ if (inputRank > 0) {
141+ inputStrides[inputRank - 1] = 1;
142+ for (int64_t d = inputRank - 2; d >= 0; d--) {
143+ inputStrides[d] = inputStrides[d + 1] * inputShape.GetDim(d + 1);
144+ }
145+ }
146+ 
147+ // 映射到 output 维度（右对齐）
148+ for (int64_t d = 0; d < yRank; d++) {
149+ int64_t inputDimIdx = d - (yRank - inputRank);
150+ if (inputDimIdx < 0) {
151+ // 该维度在 input 中不存在（补维），stride = 0
152+ strides[d] = 0;
153+ } else if (inputShape.GetDim(inputDimIdx) == 1 && yShape.GetDim(d) > 1) {
154+ // 广播轴：input dim=1, output dim>1 → stride=0
155+ strides[d] = 0;
156+ } else {
157+ // 非广播轴：使用 input 原始 stride
158+ strides[d] = inputStrides[inputDimIdx];
159+ }
160+ }
161+}
162+ 
163+// 计算连续内维大小（从最内层开始，连续非广播轴）
164+// 修复 B2/B3: 当任一输入在此维度有 stride=0（广播轴）时，内维到此结束
165+// 这确保了内维中的两个输入都是连续的，可以安全地用 DataCopyPad 搬入
166+static int64_t ComputeInnerSize(const gert::Shape& yShape,
167+ const int64_t* gradStrides,
168+ const int64_t* outStrides)
169+{
170+ int64_t rank = static_cast<int64_t>(yShape.GetDimNum());
171+ if (rank == 0) {
172+ return 1;
173+ }
174+ 
175+ // 从最内层开始，找到第一个在任一输入中是广播轴的维度
176+ // 内维 = 从最内层到第一个"任一广播轴"之间的所有维度
177+ int64_t innerSize = 1;
178+ for (int64_t d = rank - 1; d >= 0; d--) {
179+ // 如果该维度在任一输入中是广播轴（stride=0），则内维到此为止
180+ if (gradStrides[d] == 0 || outStrides[d] == 0) {
181+ break;
182+ }
183+ innerSize *= yShape.GetDim(d);
184+ }
185+ return innerSize;
186+}
187+ 
188+// 每元素 UB 字节数（OneDim 与 Broadcast 共用；buffer 数量随 dtype 的 Cast 路径而定）
189+static int64_t GetBytesPerElem(ge::DataType dataType)
190+{
191+ switch (dataType) {
192+ case ge::DT_FLOAT: return 7 * 4 + 1; // 29: 7 float 缓冲 + selMask(1B)
193+ case ge::DT_FLOAT16: return 3 * 2 + 7 * 4; // 34: 3 half 队列 + 7 float
194+ case ge::DT_BF16: return 3 * 2 + 7 * 4; // 34: 3 bf16 队列 + 7 float
195+ case ge::DT_INT32: return 3 * 4 + 2 * 2 + 7 * 4; // 44: 3 int32 队列 + 2 half + 7 float
196+ case ge::DT_INT8: return 3 * 1 + 2 * 2 + 7 * 4; // 35: 3 int8 队列 + 2 half + 7 float
197+ case ge::DT_UINT8: return 3 * 1 + 3 * 2 + 7 * 4; // 37: 3 uint8 队列 + 3 half + 7 float
198+ default: return DEFAULT_BYTES_PER_ELEM;
199+ }
200+}
201+ 
202+// TilingKey_0 (OneDim) 路径参数计算
203+static void ComputeOneDimTiling(SeluGradTilingData* tiling, int64_t totalElements,
204+ ge::DataType dataType, uint64_t ubSize,
205+ int64_t coreNum, int64_t ubBlockSize)
206+{
207+ tiling->totalElements = totalElements;
208+ 
209+ // 多核切分
210+ tiling->blockFormer = CeilAlign(CeilDiv(totalElements, coreNum), ubBlockSize);
211+ 
212+ // UB 切分（按 dtype 每元素字节数）
213+ int64_t availableUb = static_cast<int64_t>(ubSize) - SELECT_UB_RESERVE;
214+ int64_t bytesPerElem = GetBytesPerElem(dataType);
215+ tiling->ubFormer = FloorAlign(FloorDiv(availableUb, bytesPerElem), ubBlockSize);
216+ if (tiling->ubFormer < ubBlockSize) {
217+ tiling->ubFormer = ubBlockSize;
218+ }
219+}
220+ 
221+// Broadcast 内维分块 + UB/多核切分（已知 innerSize/totalRows/bytesPerElem 后）
222+static void ComputeBroadcastUbSplit(SeluGradTilingData* tiling, int64_t bytesPerElem,
223+ int64_t availableUb, int64_t coreNum, int64_t ubBlockSize)
224+{
225+ // 内维分块：如果 innerSize 太大无法放入 UB，则分块处理
226+ if (bytesPerElem <= 0) {
227+ return;
228+ }
229+ int64_t maxInnerSize = availableUb / bytesPerElem;
230+ if (maxInnerSize < 1) {
231+ return;
232+ }
233+ 
234+ int64_t innerChunkSize;
235+ int32_t numInnerChunks;
236+ if (tiling->innerSize <= maxInnerSize) {
237+ innerChunkSize = tiling->innerSize;
238+ numInnerChunks = 1;
239+ } else {
240+ innerChunkSize = FloorAlign(maxInnerSize, ubBlockSize);
241+ if (innerChunkSize < 1) {
242+ innerChunkSize = 1;
243+ }
244+ numInnerChunks = static_cast<int32_t>((tiling->innerSize + innerChunkSize - 1) / innerChunkSize);
245+ }
246+ 
247+ tiling->innerChunkSize = innerChunkSize;
248+ tiling->numInnerChunks = numInnerChunks;
249+ 
250+ // 每个 "work item" 处理一个 chunk（innerChunkSize 个元素）
251+ int64_t totalSubRows = tiling->totalRows * numInnerChunks;
252+ 
253+ // ubFormer = 每个 UB 能容纳的 work item 数
254+ int64_t bytesPerItem = innerChunkSize * bytesPerElem;
255+ tiling->ubFormer = availableUb / bytesPerItem;
256+ if (tiling->ubFormer < 1) {
257+ tiling->ubFormer = 1;
258+ }
259+ 
260+ // 多核切分（按 work item 数），blockFormer 不超过 ubFormer
261+ tiling->blockFormer = CeilDiv(totalSubRows, coreNum);
262+ if (tiling->blockFormer > tiling->ubFormer) {
263+ tiling->blockFormer = tiling->ubFormer;
264+ }
265+ if (tiling->blockFormer < 1) {
266+ tiling->blockFormer = 1;
267+ }
268+}
269+ 
270+// TilingKey_1 (Broadcast) 路径参数计算
271+static void ComputeBroadcastTiling(SeluGradTilingData* tiling,
272+ const gert::Shape& gradShape,
273+ const gert::Shape& outShape,
274+ const gert::Shape& yShape,
275+ int64_t totalElements,
276+ ge::DataType dataType,
277+ uint64_t ubSize,
278+ int64_t coreNum,
279+ int64_t ubBlockSize)
280+{
281+ tiling->totalElements = totalElements;
282+ tiling->needBroadcast = 1;
283+ tiling->shapeLen = static_cast<int32_t>(yShape.GetDimNum());
284+ 
285+ // 填充 outputDims
286+ for (int32_t d = 0; d < tiling->shapeLen; d++) {
287+ tiling->outputDims[d] = yShape.GetDim(d);
288+ }
289+ 
290+ // 计算广播 stride
291+ ComputeBroadcastStrides(yShape, gradShape, tiling->gradStrides);
292+ ComputeBroadcastStrides(yShape, outShape, tiling->outStrides);
293+ 
294+ // 计算连续内维大小
295+ tiling->innerSize = ComputeInnerSize(yShape, tiling->gradStrides, tiling->outStrides);
296+ tiling->totalRows = totalElements / tiling->innerSize;
297+ 
298+ // UB 切分：根据 dtype 每元素内存，支持内维分块
299+ int64_t availableUb = static_cast<int64_t>(ubSize) - SELECT_UB_RESERVE;
300+ int64_t bytesPerElem = GetBytesPerElem(dataType);
301+ ComputeBroadcastUbSplit(tiling, bytesPerElem, availableUb, coreNum, ubBlockSize);
302+}
303+ 
304+// 判定 schMode 并计算对应路径的 tiling 参数，返回 schMode；needBroadcast 经出参回传
305+static uint32_t DispatchTiling(SeluGradTilingData* tiling,
306+ const gert::Shape& gradShape, const gert::Shape& outShape,
307+ const gert::Shape& yShape, int64_t totalElements,
308+ ge::DataType dataType, uint64_t ubSize, int64_t coreNum,
309+ int64_t ubBlockSize, bool* needBroadcast)
310+{
311+ *needBroadcast = NeedsBroadcast(EnsureNotScalar(gradShape), EnsureNotScalar(outShape));
312+ if (*needBroadcast) {
313+ ComputeBroadcastTiling(tiling, EnsureNotScalar(gradShape), EnsureNotScalar(outShape),
314+ yShape, totalElements, dataType, ubSize, coreNum, ubBlockSize);
315+ return static_cast<uint32_t>(SELU_GRAD_BROADCAST);
316+ }
317+ ComputeOneDimTiling(tiling, totalElements, dataType, ubSize, coreNum, ubBlockSize);
318+ return static_cast<uint32_t>(SELU_GRAD_ONE_DIM);
319+}
320+ 
321+// 按 work item 数（Broadcast: 行×内维块；OneDim: 元素）设置使用核数
322+static void SetUsedCoreNum(gert::TilingContext* context, const SeluGradTilingData* tiling,
323+ int64_t totalElements, bool needBroadcast)
324+{
325+ int64_t workItems = needBroadcast ? (tiling->totalRows * tiling->numInnerChunks) : totalElements;
326+ int64_t usedCoreNum = CeilDiv(workItems, tiling->blockFormer);
327+ if (usedCoreNum < 1) {
328+ usedCoreNum = 1;
329+ }
330+ context->SetBlockDim(static_cast<uint32_t>(usedCoreNum));
331+}
332+ 
333+// Tiling 分发入口
334+static ge::graphStatus SeluGradTilingFunc(gert::TilingContext* context)
335+{
336+ // 1. 获取平台运行信息
337+ uint64_t ubSize;
338+ int64_t coreNum;
339+ OP_CHECK_IF(
340+ GetPlatformInfo(context, &ubSize, &coreNum) != ge::GRAPH_SUCCESS,
341+ OP_LOGE(context, "GetPlatformInfo error"),
342+ return ge::GRAPH_FAILED);
343+ 
344+ // 2. 获取 shape、属性信息
345+ int64_t totalElements;
346+ ge::DataType dataType;
347+ gert::Shape gradShape, outShape, yShape;
348+ OP_CHECK_IF(
349+ GetShapeAttrsInfo(context, &totalElements, &dataType, &gradShape, &outShape, &yShape) != ge::GRAPH_SUCCESS,
350+ OP_LOGE(context, "GetShapeAttrsInfo error"),
351+ return ge::GRAPH_FAILED);
352+ 
353+ // 3. 获取 WorkspaceSize
354+ OP_CHECK_IF(
355+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
356+ OP_LOGE(context, "GetWorkspaceSize error"),
357+ return ge::GRAPH_FAILED);
358+ 
359+ // 4. 设置 TilingData
360+ SeluGradTilingData* tiling = context->GetTilingData<SeluGradTilingData>();
361+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
362+ OP_CHECK_IF(
363+ memset_s(tiling, sizeof(SeluGradTilingData), 0, sizeof(SeluGradTilingData)) != EOK,
364+ OP_LOGE(context, "set tiling data error"),
365+ return ge::GRAPH_FAILED);
366+ 
367+ int64_t ubBlockSize = Ops::Base::GetUbBlockSize(context);
368+ 
369+ // 空 Tensor 快速返回
370+ if (totalElements == 0) {
371+ context->SetBlockDim(1);
372+ uint32_t dTypeXEmpty = static_cast<uint32_t>(dataType);
373+ ASCENDC_TPL_SEL_PARAM(context, dTypeXEmpty, static_cast<uint32_t>(SELU_GRAD_ONE_DIM));
374+ return ge::GRAPH_SUCCESS;
375+ }
376+ 
377+ // 5. 判定 TilingKey 并计算 tiling
378+ bool needBroadcast = false;
379+ uint32_t schMode = DispatchTiling(tiling, gradShape, outShape, yShape, totalElements,
380+ dataType, ubSize, coreNum, ubBlockSize, &needBroadcast);
381+ 
382+ // 6. 校验 tiling 参数有效性
383+ if (tiling->blockFormer < 1 || tiling->ubFormer < 1) {
384+ return ge::GRAPH_FAILED;
385+ }
386+ 
387+ // 7. 设置核数 + TilingKey（dtype + schMode 双参数）
388+ SetUsedCoreNum(context, tiling, totalElements, needBroadcast);
389+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
390+ ASCENDC_TPL_SEL_PARAM(context, dTypeX, schMode);
391+ 
392+ return ge::GRAPH_SUCCESS;
393+}
394+ 
395+static ge::graphStatus TilingParseForSeluGrad([[maybe_unused]] gert::TilingParseContext* context)
396+{
397+ return ge::GRAPH_SUCCESS;
398+}
399+ 
400+struct SeluGradCompileInfo {};
401+ 
402+// Tiling 注册入口
403+IMPL_OP_OPTILING(SeluGrad).Tiling(SeluGradTilingFunc).TilingParse<SeluGradCompileInfo>(TilingParseForSeluGrad);
404+ 
405+} // namespace optiling
Aactivation/selu_grad/op_host/selu_grad_def.cpp+66-0
@@ -0,0 +1,66 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file selu_grad_def.cpp
15+ * \brief SeluGrad 算子定义，声明输入输出和算子配置
16+ *
17+ * 2 输入（gradients, outputs）→ 1 输出（y），无属性
18+ * 迭代一：float16 + Ascend950
19+ */
20+#include "register/op_def_registry.h"
21+ 
22+namespace ops {
23+class SeluGrad : public OpDef {
24+public:
25+ explicit SeluGrad(const char* name) : OpDef(name)
26+ {
27+ this->Input("gradients") // 输入0：反向传播上游梯度
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16,
30+ ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
31+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
32+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
33+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
34+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
35+ .AutoContiguous();
36+ this->Input("outputs") // 输入1：SELU 前向输出
37+ .ParamType(REQUIRED)
38+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16,
39+ ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
40+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
41+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
42+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
43+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
44+ .AutoContiguous();
45+ this->Output("y") // 输出0：反向梯度结果
46+ .ParamType(REQUIRED)
47+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16,
48+ ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
49+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
50+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
51+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
52+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
53+ .AutoContiguous();
54+ 
55+ OpAICoreConfig aicoreConfig950;
56+ aicoreConfig950.DynamicCompileStaticFlag(true)
57+ .DynamicFormatFlag(false)
58+ .DynamicRankSupportFlag(true)
59+ .DynamicShapeSupportFlag(true)
60+ .NeedCheckSupportFlag(false)
61+ .PrecisionReduceFlag(true);
62+ this->AICore().AddConfig("ascend950", aicoreConfig950);
63+ }
64+};
65+OP_ADD(SeluGrad);
66+} // namespace ops
Aactivation/selu_grad/op_host/selu_grad_infershape.cpp+99-0
@@ -0,0 +1,99 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file selu_grad_infershape.cpp
15+ * \brief SeluGrad 算子形状推导实现
16+ *
17+ * 迭代二：实现 numpy broadcast 形状推导
18+ * 输出 shape = broadcast(gradients.shape, outputs.shape)
19+ */
20+ 
21+#include "register/op_impl_registry.h"
22+#include "exe_graph/runtime/infer_shape_context.h"
23+#include "op_common/log/log.h"
24+ 
25+using namespace ge;
26+ 
27+namespace ops {
28+ 
29+constexpr int64_t MAX_SUPPORTED_RANK = 8;
30+ 
31+static bool NumpyBroadcastShape(const gert::Shape& gradShape, const gert::Shape& outShape,
32+ gert::Shape& yShape)
33+{
34+ int64_t gradRank = static_cast<int64_t>(gradShape.GetDimNum());
35+ int64_t outRank = static_cast<int64_t>(outShape.GetDimNum());
36+ int64_t maxRank = (gradRank > outRank) ? gradRank : outRank;
37+ 
38+ if (maxRank > MAX_SUPPORTED_RANK) {
39+ return false;
40+ }
41+ 
42+ gert::Shape result;
43+ for (int64_t i = 0; i < maxRank; i++) {
44+ int64_t gradDim = (i < maxRank - gradRank) ? 1 : gradShape.GetDim(i - (maxRank - gradRank));
45+ int64_t outDim = (i < maxRank - outRank) ? 1 : outShape.GetDim(i - (maxRank - outRank));
46+ 
47+ if (gradDim == outDim) {
48+ result.AppendDim(gradDim);
49+ } else if (gradDim == 1) {
50+ result.AppendDim(outDim);
51+ } else if (outDim == 1) {
52+ result.AppendDim(gradDim);
53+ } else {
54+ return false; // 不可广播
55+ }
56+ }
57+ 
58+ yShape = result;
59+ return true;
60+}
61+ 
62+static ge::graphStatus InferShape4SeluGrad(gert::InferShapeContext* context)
63+{
64+ const gert::Shape* gradShape = context->GetInputShape(0);
65+ OP_CHECK_NULL_WITH_CONTEXT(context, gradShape);
66+ 
67+ const gert::Shape* outShape = context->GetInputShape(1);
68+ OP_CHECK_NULL_WITH_CONTEXT(context, outShape);
69+ 
70+ gert::Shape* yShape = context->GetOutputShape(0);
71+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
72+ 
73+ // 处理 rank=0 标量
74+ if (gradShape->GetDimNum() == 0 && outShape->GetDimNum() == 0) {
75+ // 两个标量输入，输出也是标量
76+ return ge::GRAPH_SUCCESS;
77+ }
78+ 
79+ // 处理 rank=0 标量 + tensor 广播
80+ if (gradShape->GetDimNum() == 0) {
81+ *yShape = *outShape;
82+ return ge::GRAPH_SUCCESS;
83+ }
84+ if (outShape->GetDimNum() == 0) {
85+ *yShape = *gradShape;
86+ return ge::GRAPH_SUCCESS;
87+ }
88+ 
89+ // numpy broadcast 推导
90+ OP_CHECK_IF(!NumpyBroadcastShape(*gradShape, *outShape, *yShape),
91+ OP_LOGE(context, "Shape broadcast failed"),
92+ return ge::GRAPH_FAILED);
93+ 
94+ return ge::GRAPH_SUCCESS;
95+}
96+ 
97+IMPL_OP_INFERSHAPE(SeluGrad).InferShape(InferShape4SeluGrad);
98+ 
99+} // namespace ops
Aactivation/selu_grad/op_kernel/arch35/selu_grad.h+474-0
@@ -0,0 +1,474 @@
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+/* Generated By CANNBot - Iter2 Fix */
12+ 
13+/*!
14+ * \file selu_grad.h
15+ * \brief SeluGrad 算子 Kernel 类定义（arch35 架构）
16+ *
17+ * 全 dtype × 2 SCH_MODE = 12 TilingKey。
18+ *
19+ * 设计：计算逻辑封装为两个 "Kit"，搬运/切分逻辑封装为两个驱动类，二者组合：
20+ * - SeluGradDirectKit<T> : float 直算（在 T==float 上原地计算）
21+ * - SeluGradTransitKit<T> : half/bfloat16/int32/int8/uint8 FP32 中转
22+ * (half/bf16 单步 Cast；int32/int8/uint8 经 half 两步 Cast)
23+ * - SeluGradOneDim<T, Kit> : 连续分块搬运 + 多核切分
24+ * - SeluGradBroadcast<T, Kit> : 多维 stride 偏移 + 内维分块逐行搬运
25+ *
26+ * 公式：
27+ * y = SCALE * gradients if outputs >= 0
28+ * y = gradients * (outputs + SCALE_ALPHA_PRODUCT) if outputs < 0
29+ */
30+ 
31+#ifndef SELU_GRAD_H
T
Ttangweiwei26月27日

这个头文件的宏没有引入ARCH，如何区别不同的平台算子？

likedislike
Y
yulianjie
6月27日 评论：
32+#define SELU_GRAD_H
33+ 
34+#include "kernel_operator.h"
35+#include "kernel_tiling/kernel_tiling.h"
36+#include "selu_grad_tiling_data.h"
37+#include "selu_grad_tiling_key.h"
38+ 
39+namespace NsSeluGrad {
40+ 
41+using namespace AscendC;
42+ 
43+// SELU 常量（FP32 精度）
44+constexpr float SCALE_F = 1.0507009873554804934193349852946f;
45+constexpr float SCALE_ALPHA_PRODUCT_F = 1.7580993408473768599402175208123f;
46+ 
47+// 需要经 half 两步 Cast（int → half → float）的 dtype（DAV_3510 不支持 int↔float 直接 Cast）
48+template <typename T>
49+constexpr bool kNeedsHalfTransit =
50+ std::is_same_v<T, int32_t> || std::is_same_v<T, int8_t> || std::is_same_v<T, uint8_t>;
51+ 
52+// ============================================================================
53+// 公共自由函数：搬运参数 / 计算核 / Cast / 广播偏移
54+// ============================================================================
55+ 
56+template <typename T>
57+__aicore__ inline DataCopyParams MakeCopyParams(int64_t count)
58+{
59+ DataCopyParams params;
60+ params.blockCount = 1;
61+ params.blockLen = count * sizeof(T);
62+ params.srcStride = 0;
63+ params.dstStride = 0;
64+ return params;
65+}
66+ 
67+// SELU 反向核（FP32）。direct 路径 T==float 时同样复用此核。
68+// Iter3 Fix #1: 使用 float 常量做 Muls/Adds，避免 half 精度截断。
69+__aicore__ inline void SeluGradSelectFp32(const LocalTensor<float>& yFp32,
70+ const LocalTensor<float>& gradFp32,
71+ const LocalTensor<float>& outFp32,
72+ const LocalTensor<float>& branchA,
73+ const LocalTensor<float>& branchB,
74+ const LocalTensor<float>& tmp,
75+ const LocalTensor<uint8_t>& selMask,
76+ int64_t n)
77+{
78+ CompareScalar(selMask, outFp32, (float)0.0f, CMPMODE::LT, n);
79+ Muls(branchA, gradFp32, SCALE_F, n);
80+ Adds(tmp, outFp32, SCALE_ALPHA_PRODUCT_F, n);
81+ Mul(branchB, gradFp32, tmp, n);
82+ Select(yFp32, selMask, branchB, branchA, SELMODE::VSEL_TENSOR_TENSOR_MODE, n);
83+}
84+ 
85+// 输入 Cast: T → float（整型经 half 两步，浮点单步）
86+template <typename T>
87+__aicore__ inline void SeluGradCastInToFp32(const LocalTensor<float>& gradFp32,
88+ const LocalTensor<float>& outFp32,
89+ const LocalTensor<T>& gradLocal,
90+ const LocalTensor<T>& outLocal,
91+ const LocalTensor<half>& gradHalf,
92+ const LocalTensor<half>& outHalf,
93+ int64_t n)
94+{
95+ if constexpr (kNeedsHalfTransit<T>) {
96+ Cast(gradHalf, gradLocal, RoundMode::CAST_NONE, n);
97+ Cast(outHalf, outLocal, RoundMode::CAST_NONE, n);
98+ Cast(gradFp32, gradHalf, RoundMode::CAST_NONE, n);
99+ Cast(outFp32, outHalf, RoundMode::CAST_NONE, n);
100+ } else {
101+ Cast(gradFp32, gradLocal, RoundMode::CAST_NONE, n);
102+ Cast(outFp32, outLocal, RoundMode::CAST_NONE, n);
103+ }
104+}
105+ 
106+// 输出 Cast: float → T（整型经 half 两步，浮点单步），统一 CAST_RINT
107+template <typename T>
108+__aicore__ inline void SeluGradCastFp32ToOut(const LocalTensor<T>& yLocal,
109+ const LocalTensor<float>& yFp32,
110+ const LocalTensor<half>& yHalf,
111+ int64_t n)
112+{
113+ if constexpr (kNeedsHalfTransit<T>) {
114+ Cast(yHalf, yFp32, RoundMode::CAST_RINT, n);
115+ Cast(yLocal, yHalf, RoundMode::CAST_RINT, n);
116+ } else {
117+ Cast(yLocal, yFp32, RoundMode::CAST_RINT, n);
118+ }
119+}
120+ 
121+// 广播路径：flatIdx → multiIdx 分解
122+__aicore__ inline void FlatIdxToMultiIdx(int64_t flatIdx, int32_t shapeLen,
123+ const int64_t* dims, int64_t* multiIdx)
124+{
125+ for (int32_t d = shapeLen - 1; d >= 0; d--) {
126+ if (dims[d] > 0) {
127+ multiIdx[d] = flatIdx % dims[d];
128+ flatIdx = flatIdx / dims[d];
129+ } else {
130+ multiIdx[d] = 0;
131+ }
132+ }
133+}
134+ 
135+__aicore__ inline int64_t ComputeStrideOffset(int32_t shapeLen, const int64_t* multiIdx,
136+ const int64_t* strides)
137+{
138+ int64_t offset = 0;
139+ for (int32_t d = 0; d < shapeLen; d++) {
140+ offset += multiIdx[d] * strides[d];
141+ }
142+ return offset;
143+}
144+ 
145+// 广播行的 GM 偏移与本块元素数（chunkOffset 已折入三个偏移）
146+struct RowGeom {
147+ int64_t gradOffset;
148+ int64_t outOffset;
149+ int64_t yOffset;
150+ int64_t count;
151+};
152+ 
153+__aicore__ inline RowGeom ComputeRowGeom(int64_t rowIdx, int32_t chunkIdx,
154+ int64_t innerSize, int32_t shapeLen,
155+ int64_t innerChunkSize,
156+ const int64_t* outputDims,
157+ const int64_t* gradStrides,
158+ const int64_t* outStrides)
159+{
160+ int64_t flatIdx = rowIdx * innerSize;
161+ int64_t multiIdx[SELU_GRAD_MAX_DIM];
162+ FlatIdxToMultiIdx(flatIdx, shapeLen, outputDims, multiIdx);
163+ 
164+ int64_t chunkOffset = (int64_t)chunkIdx * innerChunkSize;
165+ int64_t count = innerChunkSize;
166+ if (chunkOffset + count > innerSize) {
167+ count = innerSize - chunkOffset;
168+ }
169+ 
170+ RowGeom geom;
171+ geom.gradOffset = ComputeStrideOffset(shapeLen, multiIdx, gradStrides) + chunkOffset;
172+ geom.outOffset = ComputeStrideOffset(shapeLen, multiIdx, outStrides) + chunkOffset;
173+ geom.yOffset = rowIdx * innerSize + chunkOffset;
174+ geom.count = count;
175+ return geom;
176+}
177+ 
178+// ============================================================================
179+// 计算 Kit：封装计算所需缓冲与一次 (grad, out) -> y 的计算
180+// ============================================================================
181+ 
182+// 直算 Kit（half/float；实际仅 float 实例化）
183+template <typename T>
184+struct SeluGradDirectKit {
185+ TQue<QuePosition::VECCALC, 1> branchAQueue, branchBQueue, tmpQueue, selMaskQueue;
186+ TBuf<TPosition::VECCALC> nanMaskBuf;
187+ 
188+ __aicore__ inline void InitBufs(TPipe& pipe, int64_t n)
189+ {
190+ pipe.InitBuffer(branchAQueue, 1, n * sizeof(T));
191+ pipe.InitBuffer(branchBQueue, 1, n * sizeof(T));
192+ pipe.InitBuffer(tmpQueue, 1, n * sizeof(T));
193+ pipe.InitBuffer(selMaskQueue, 1, (n / 8) + 32);
194+ pipe.InitBuffer(nanMaskBuf, (n / 8) + 32);
195+ }
196+ 
197+ __aicore__ inline void Compute(const LocalTensor<T>& gradLocal, const LocalTensor<T>& outLocal,
198+ const LocalTensor<T>& yLocal, int64_t n)
199+ {
200+ LocalTensor<T> branchA = branchAQueue.template AllocTensor<T>();
201+ LocalTensor<T> branchB = branchBQueue.template AllocTensor<T>();
202+ LocalTensor<T> tmp = tmpQueue.template AllocTensor<T>();
203+ LocalTensor<uint8_t> selMask = selMaskQueue.template AllocTensor<uint8_t>();
204+ 
205+ SeluGradSelectFp32(yLocal, gradLocal, outLocal, branchA, branchB, tmp, selMask, n);
206+ 
207+ branchAQueue.FreeTensor(branchA);
208+ branchBQueue.FreeTensor(branchB);
209+ tmpQueue.FreeTensor(tmp);
210+ selMaskQueue.FreeTensor(selMask);
211+ }
212+};
213+ 
214+// FP32 中转 Kit（half/bfloat16/int32/int8/uint8）
215+template <typename T>
216+struct SeluGradTransitKit {
217+ TBuf<TPosition::VECCALC> gradHalfBuf, outHalfBuf, yHalfBuf;
218+ TBuf<TPosition::VECCALC> gradFp32Buf, outFp32Buf, branchABuf, branchBBuf, tmpBuf, yFp32Buf, maskBuf;
219+ 
220+ __aicore__ inline void InitBufs(TPipe& pipe, int64_t n)
221+ {
222+ if constexpr (kNeedsHalfTransit<T>) {
223+ pipe.InitBuffer(gradHalfBuf, n * sizeof(half));
224+ pipe.InitBuffer(outHalfBuf, n * sizeof(half));
225+ if constexpr (!std::is_same_v<T, int8_t>) {
226+ pipe.InitBuffer(yHalfBuf, n * sizeof(half)); // int8 复用 outHalfBuf
227+ }
228+ }
229+ pipe.InitBuffer(gradFp32Buf, n * sizeof(float));
230+ pipe.InitBuffer(outFp32Buf, n * sizeof(float));
231+ pipe.InitBuffer(branchABuf, n * sizeof(float));
232+ pipe.InitBuffer(branchBBuf, n * sizeof(float));
233+ pipe.InitBuffer(tmpBuf, n * sizeof(float));
234+ pipe.InitBuffer(yFp32Buf, n * sizeof(float));
235+ pipe.InitBuffer(maskBuf, (n / 8) + 32);
236+ }
237+ 
238+ __aicore__ inline void Compute(const LocalTensor<T>& gradLocal, const LocalTensor<T>& outLocal,
239+ const LocalTensor<T>& yLocal, int64_t n)
240+ {
241+ LocalTensor<float> gradFp32 = gradFp32Buf.template Get<float>();
242+ LocalTensor<float> outFp32 = outFp32Buf.template Get<float>();
243+ LocalTensor<float> branchA = branchABuf.template Get<float>();
244+ LocalTensor<float> branchB = branchBBuf.template Get<float>();
245+ LocalTensor<float> tmp = tmpBuf.template Get<float>();
246+ LocalTensor<uint8_t> selMask = maskBuf.template Get<uint8_t>();
247+ LocalTensor<float> yFp32 = yFp32Buf.template Get<float>();
248+ 
249+ LocalTensor<half> gradHalf;
250+ LocalTensor<half> outHalf;
251+ LocalTensor<half> yHalf;
252+ if constexpr (kNeedsHalfTransit<T>) {
253+ gradHalf = gradHalfBuf.template Get<half>();
254+ outHalf = outHalfBuf.template Get<half>();
255+ if constexpr (std::is_same_v<T, int8_t>) {
256+ yHalf = outHalfBuf.template Get<half>();
257+ } else {
258+ yHalf = yHalfBuf.template Get<half>();
259+ }
260+ }
261+ 
262+ SeluGradCastInToFp32<T>(gradFp32, outFp32, gradLocal, outLocal, gradHalf, outHalf, n);
263+ SeluGradSelectFp32(yFp32, gradFp32, outFp32, branchA, branchB, tmp, selMask, n);
264+ SeluGradCastFp32ToOut<T>(yLocal, yFp32, yHalf, n);
265+ }
266+};
267+ 
268+// ============================================================================
269+// 驱动类：连续分块搬运（OneDim） / 多维广播逐行搬运（Broadcast）
270+// ============================================================================
271+ 
272+template <typename T, typename Kit>
273+class SeluGradOneDim {
274+public:
275+ __aicore__ inline void Init(GM_ADDR gradients, GM_ADDR outputs, GM_ADDR y,
276+ const SeluGradTilingData* tilingData)
277+ {
278+ if (tilingData->totalElements == 0) {
279+ blockLen_ = 0;
280+ return;
281+ }
282+ int64_t blockOffset = tilingData->blockFormer * AscendC::GetBlockIdx();
283+ int64_t remainder = tilingData->totalElements - blockOffset;
284+ blockLen_ = (remainder > tilingData->blockFormer) ? tilingData->blockFormer : remainder;
285+ ubLen_ = tilingData->ubFormer;
286+ 
287+ gradGM.SetGlobalBuffer((__gm__ T*)gradients + blockOffset, blockLen_);
288+ outGM.SetGlobalBuffer((__gm__ T*)outputs + blockOffset, blockLen_);
289+ yGM.SetGlobalBuffer((__gm__ T*)y + blockOffset, blockLen_);
290+ 
291+ pipe.InitBuffer(gradQueue, 1, ubLen_ * sizeof(T));
292+ pipe.InitBuffer(outQueue, 1, ubLen_ * sizeof(T));
293+ pipe.InitBuffer(yQueue, 1, ubLen_ * sizeof(T));
294+ kit_.InitBufs(pipe, ubLen_);
295+ }
296+ 
297+ __aicore__ inline void Process()
298+ {
299+ if (blockLen_ == 0) {
300+ return;
301+ }
302+ int64_t loopCount = (blockLen_ + ubLen_ - 1) / ubLen_;
303+ for (int64_t i = 0; i < loopCount; i++) {
304+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLen_ - ubLen_ * i) : ubLen_;
305+ CopyIn(i, currentNum);
306+ Compute(currentNum);
307+ CopyOut(i, currentNum);
308+ }
309+ }
310+ 
311+private:
312+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum)
313+ {
314+ LocalTensor<T> gradLocal = gradQueue.template AllocTensor<T>();
315+ LocalTensor<T> outLocal = outQueue.template AllocTensor<T>();
316+ DataCopyParams copyParams = MakeCopyParams<T>(currentNum);
317+ DataCopyPad(gradLocal, gradGM[progress * ubLen_], copyParams, {false, 0, 0, 0});
318+ DataCopyPad(outLocal, outGM[progress * ubLen_], copyParams, {false, 0, 0, 0});
319+ gradQueue.EnQue(gradLocal);
320+ outQueue.EnQue(outLocal);
321+ }
322+ 
323+ __aicore__ inline void Compute(int64_t currentNum)
324+ {
325+ LocalTensor<T> gradLocal = gradQueue.template DeQue<T>();
326+ LocalTensor<T> outLocal = outQueue.template DeQue<T>();
327+ LocalTensor<T> yLocal = yQueue.template AllocTensor<T>();
328+ kit_.Compute(gradLocal, outLocal, yLocal, currentNum);
W
Wwang-xing0016月25日

为什么要用kit操作

likedislike
Y
yulianjie
6月25日 评论：
329+ yQueue.template EnQue<T>(yLocal);
330+ gradQueue.FreeTensor(gradLocal);
331+ outQueue.FreeTensor(outLocal);
332+ }
333+ 
334+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum)
335+ {
336+ LocalTensor<T> yLocal = yQueue.template DeQue<T>();
337+ DataCopyParams copyParams = MakeCopyParams<T>(currentNum);
338+ DataCopyPad(yGM[progress * ubLen_], yLocal, copyParams);
339+ yQueue.FreeTensor(yLocal);
340+ }
341+ 
342+ TPipe pipe;
T
Ttangweiwei26月27日

这边命名风格为什么不同，都是私有成员变量，一部分有_，一部分没有？

likedislike
Y
yulianjie
6月27日 评论：
343+ TQue<QuePosition::VECIN, 1> gradQueue, outQueue;
344+ TQue<QuePosition::VECOUT, 1> yQueue;
345+ GlobalTensor<T> gradGM, outGM, yGM;
346+ int64_t blockLen_ = 0;
347+ int64_t ubLen_ = 0;
348+ Kit kit_;
349+};
350+ 
351+template <typename T, typename Kit>
352+class SeluGradBroadcast {
353+public:
354+ __aicore__ inline void Init(GM_ADDR gradients, GM_ADDR outputs, GM_ADDR y,
355+ const SeluGradTilingData* tilingData)
356+ {
357+ totalElements_ = tilingData->totalElements;
358+ if (totalElements_ == 0) {
359+ return;
360+ }
361+ innerSize_ = tilingData->innerSize;
362+ totalRows_ = tilingData->totalRows;
363+ shapeLen_ = tilingData->shapeLen;
364+ blockRows_ = tilingData->blockFormer;
365+ innerChunkSize_ = tilingData->innerChunkSize;
366+ numInnerChunks_ = tilingData->numInnerChunks;
367+ 
368+ for (int32_t d = 0; d < shapeLen_; d++) {
369+ outputDims_[d] = tilingData->outputDims[d];
370+ gradStrides_[d] = tilingData->gradStrides[d];
371+ outStrides_[d] = tilingData->outStrides[d];
372+ }
373+ 
374+ int64_t totalItems = totalRows_ * numInnerChunks_;
375+ startItem_ = (int64_t)AscendC::GetBlockIdx() * blockRows_;
376+ endItem_ = startItem_ + blockRows_;
377+ if (endItem_ > totalItems) {
378+ endItem_ = totalItems;
379+ }
380+ 
381+ gradBase_ = (__gm__ T*)gradients;
382+ outBase_ = (__gm__ T*)outputs;
383+ yBase_ = (__gm__ T*)y;
384+ 
385+ pipe.InitBuffer(gradQueue, 1, innerChunkSize_ * sizeof(T));
386+ pipe.InitBuffer(outQueue, 1, innerChunkSize_ * sizeof(T));
387+ pipe.InitBuffer(yQueue, 1, innerChunkSize_ * sizeof(T));
388+ kit_.InitBufs(pipe, innerChunkSize_);
389+ }
390+ 
391+ __aicore__ inline void Process()
392+ {
393+ if (totalElements_ == 0) {
394+ return;
395+ }
396+ for (int64_t item = startItem_; item < endItem_; item++) {
397+ ProcessRow(item / numInnerChunks_, static_cast<int32_t>(item % numInnerChunks_));
398+ }
399+ }
400+ 
401+private:
402+ __aicore__ inline void ProcessRow(int64_t rowIdx, int32_t chunkIdx)
403+ {
404+ RowGeom geom = ComputeRowGeom(rowIdx, chunkIdx, innerSize_, shapeLen_, innerChunkSize_,
405+ outputDims_, gradStrides_, outStrides_);
406+ gradGM.SetGlobalBuffer(gradBase_ + geom.gradOffset, geom.count);
407+ outGM.SetGlobalBuffer(outBase_ + geom.outOffset, geom.count);
408+ yGM.SetGlobalBuffer(yBase_ + geom.yOffset, geom.count);
409+ 
410+ // CopyIn
411+ LocalTensor<T> gradLocal = gradQueue.template AllocTensor<T>();
412+ LocalTensor<T> outLocal = outQueue.template AllocTensor<T>();
413+ DataCopyParams copyParams = MakeCopyParams<T>(geom.count);
414+ DataCopyPad(gradLocal, gradGM, copyParams, {false, 0, 0, 0});
415+ DataCopyPad(outLocal, outGM, copyParams, {false, 0, 0, 0});
416+ gradQueue.EnQue(gradLocal);
417+ outQueue.EnQue(outLocal);
418+ 
419+ // Compute
420+ gradLocal = gradQueue.template DeQue<T>();
421+ outLocal = outQueue.template DeQue<T>();
422+ LocalTensor<T> yLocal = yQueue.template AllocTensor<T>();
423+ kit_.Compute(gradLocal, outLocal, yLocal, geom.count);
424+ yQueue.template EnQue<T>(yLocal);
425+ gradQueue.FreeTensor(gradLocal);
426+ outQueue.FreeTensor(outLocal);
427+ 
428+ // CopyOut
429+ yLocal = yQueue.template DeQue<T>();
430+ DataCopyPad(yGM, yLocal, MakeCopyParams<T>(geom.count));
431+ yQueue.FreeTensor(yLocal);
432+ }
433+ 
434+ TPipe pipe;
435+ TQue<QuePosition::VECIN, 1> gradQueue, outQueue;
436+ TQue<QuePosition::VECOUT, 1> yQueue;
437+ GlobalTensor<T> gradGM, outGM, yGM;
438+ 
439+ int64_t totalElements_ = 0;
440+ int64_t innerSize_ = 0;
441+ int64_t totalRows_ = 0;
442+ int64_t blockRows_ = 0;
443+ int32_t shapeLen_ = 0;
444+ int64_t innerChunkSize_ = 0;
445+ int32_t numInnerChunks_ = 0;
446+ 
447+ int64_t outputDims_[SELU_GRAD_MAX_DIM];
448+ int64_t gradStrides_[SELU_GRAD_MAX_DIM];
449+ int64_t outStrides_[SELU_GRAD_MAX_DIM];
450+ 
451+ int64_t startItem_ = 0;
452+ int64_t endItem_ = 0;
453+ 
454+ __gm__ T* gradBase_;
455+ __gm__ T* outBase_;
456+ __gm__ T* yBase_;
457+ 
458+ Kit kit_;
459+};
460+ 
461+// dtype 调度别名：float 走 Direct，其余走 Transit
462+template <typename T>
463+using SeluGradOneDimOp =
464+ SeluGradOneDim<T, std::conditional_t<std::is_same_v<T, float>,
465+ SeluGradDirectKit<T>, SeluGradTransitKit<T>>>;
466+ 
467+template <typename T>
468+using SeluGradBroadcastOp =
469+ SeluGradBroadcast<T, std::conditional_t<std::is_same_v<T, float>,
470+ SeluGradDirectKit<T>, SeluGradTransitKit<T>>>;
471+ 
472+} // namespace NsSeluGrad
473+ 
474+#endif // SELU_GRAD_H
Aactivation/selu_grad/op_kernel/arch35/selu_grad_tiling_data.h+46-0
@@ -0,0 +1,46 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file selu_grad_tiling_data.h
15+ * \brief TilingData 结构体定义（arch35 架构）
16+ *
17+ * 迭代二：TilingKey_0 (OneDim) + TilingKey_1 (Broadcast)
18+ * 支持 float16, float32, bfloat16 三种浮点 dtype
19+ */
20+ 
21+#ifndef _SELU_GRAD_TILING_DATA_H_
22+#define _SELU_GRAD_TILING_DATA_H_
23+ 
24+constexpr int32_t SELU_GRAD_MAX_DIM = 8;
25+ 
26+struct SeluGradTilingData {
27+ // === 基础信息（TilingKey_0 和 TilingKey_1 共用） ===
28+ int64_t totalElements = 0; // 输出总元素数（broadcast 后）
29+ int64_t blockFormer = 0; // OneDim: 每核元素数; Broadcast: 每核行数
30+ int64_t ubFormer = 0; // OneDim: 每次 UB 元素数; Broadcast: 每次 UB 行数
31+ 
32+ // === 广播信息（TilingKey_1 使用） ===
33+ int32_t needBroadcast = 0; // 0=无需广播, 1=需要广播
34+ int32_t shapeLen = 0; // 合轴后维度数
35+ int64_t innerSize = 0; // 连续内维大小（最内层连续非广播轴及其右侧元素积）
36+ int64_t totalRows = 0; // 外层行数 = totalElements / innerSize
37+ int64_t innerChunkSize = 0; // 内维分块大小（每块元素数，<= innerSize）
38+ int32_t numInnerChunks = 0; // 内维分块数 = CeilDiv(innerSize, innerChunkSize)
39+ 
40+ // === 多维广播参数（仅 TilingKey_1 使用） ===
41+ int64_t outputDims[SELU_GRAD_MAX_DIM]; // 合轴后输出 shape
42+ int64_t gradStrides[SELU_GRAD_MAX_DIM]; // gradients 合轴后 stride（0=广播轴）
43+ int64_t outStrides[SELU_GRAD_MAX_DIM]; // outputs 合轴后 stride（0=广播轴）
44+};
45+ 
46+#endif // _SELU_GRAD_TILING_DATA_H_
Aactivation/selu_grad/op_kernel/arch35/selu_grad_tiling_key.h+71-0
@@ -0,0 +1,71 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file selu_grad_tiling_key.h
15+ * \brief TilingKey 模板参数定义（arch35 架构）
16+ *
17+ * 迭代二修复：6 dtype × 2 SCH_MODE = 12 TilingKey
T
Ttangweiwei26月27日

这种agent内部迭代的信息不要暴露在外面。

likedislike
Y
yulianjie
6月27日 评论：
18+ *
19+ * dtype: float16, float32, bfloat16, int32, int8, uint8
20+ * SCH_MODE: OneDim(0), Broadcast(1)
21+ */
22+ 
23+#ifndef __SELU_GRAD_TILING_KEY_H__
24+#define __SELU_GRAD_TILING_KEY_H__
25+ 
26+#include "ascendc/host_api/tiling/template_argument.h"
27+ 
28+// 调度模式定义
29+#define SELU_GRAD_ONE_DIM 0 // OneDim 路径（合轴后仅 1 维，或 shape 完全一致）
30+#define SELU_GRAD_BROADCAST 1 // Broadcast 路径（合轴后 > 1 维，需要多维广播）
31+ 
32+ASCENDC_TPL_ARGS_DECL(SeluGrad,
33+ ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT16, C_DT_FLOAT, C_DT_BF16,
34+ C_DT_INT32, C_DT_INT8, C_DT_UINT8),
35+ ASCENDC_TPL_UINT_DECL(SCH_MODE, 8, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM, SELU_GRAD_BROADCAST)
36+);
37+ 
38+ASCENDC_TPL_SEL(
39+ // float16: OneDim + Broadcast
40+ ASCENDC_TPL_ARGS_SEL(
41+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16),
42+ ASCENDC_TPL_UINT_SEL(SCH_MODE, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM, SELU_GRAD_BROADCAST)
43+ ),
44+ // float32: OneDim + Broadcast
45+ ASCENDC_TPL_ARGS_SEL(
46+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT),
47+ ASCENDC_TPL_UINT_SEL(SCH_MODE, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM, SELU_GRAD_BROADCAST)
48+ ),
49+ // bfloat16: OneDim + Broadcast
50+ ASCENDC_TPL_ARGS_SEL(
51+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_BF16),
52+ ASCENDC_TPL_UINT_SEL(SCH_MODE, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM, SELU_GRAD_BROADCAST)
53+ ),
54+ // int32: OneDim + Broadcast
55+ ASCENDC_TPL_ARGS_SEL(
56+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT32),
57+ ASCENDC_TPL_UINT_SEL(SCH_MODE, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM, SELU_GRAD_BROADCAST)
58+ ),
59+ // int8: OneDim + Broadcast
60+ ASCENDC_TPL_ARGS_SEL(
61+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT8),
62+ ASCENDC_TPL_UINT_SEL(SCH_MODE, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM, SELU_GRAD_BROADCAST)
63+ ),
64+ // uint8: OneDim + Broadcast
65+ ASCENDC_TPL_ARGS_SEL(
66+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_UINT8),
67+ ASCENDC_TPL_UINT_SEL(SCH_MODE, ASCENDC_TPL_UI_LIST, SELU_GRAD_ONE_DIM, SELU_GRAD_BROADCAST)
68+ )
69+);
70+ 
71+#endif // __SELU_GRAD_TILING_KEY_H__
Aactivation/selu_grad/op_kernel/selu_grad.cpp+48-0
@@ -0,0 +1,48 @@
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+/* Generated By CANNBot - Iter2 Fix */
12+ 
13+/*!
14+ * \file selu_grad_arch35.cpp
15+ * \brief SeluGrad 算子 Kernel 入口（arch35 架构）
16+ *
17+ * 6 dtype × 2 SCH_MODE = 12 TilingKey
18+ *
19+ * 分发策略：
20+ * - float → Direct 直算路径（T 原地计算）
21+ * - half/bfloat16/int32/int8/uint8 → Transit FP32 中转路径
22+ * SCH_MODE=0 (OneDim) 与 SCH_MODE=1 (Broadcast) 分别选用对应的 OneDim/Broadcast 类。
23+ */
24+ 
25+#include "arch35/selu_grad.h"
26+ 
27+template <typename Op>
28+__aicore__ inline void RunSeluGrad(GM_ADDR gradients, GM_ADDR outputs, GM_ADDR y,
29+ const SeluGradTilingData* tilingData)
30+{
31+ Op op;
32+ op.Init(gradients, outputs, y, tilingData);
33+ op.Process();
34+}
35+ 
36+template <typename D_T_X, int SCH_MODE>
37+__global__ __aicore__ void selu_grad(GM_ADDR gradients, GM_ADDR outputs, GM_ADDR y,
38+ GM_ADDR workspace, GM_ADDR tiling)
39+{
40+ REGISTER_TILING_DEFAULT(SeluGradTilingData);
41+ GET_TILING_DATA_WITH_STRUCT(SeluGradTilingData, tilingData, tiling);
42+ 
43+ if constexpr (SCH_MODE == SELU_GRAD_ONE_DIM) {
44+ RunSeluGrad<NsSeluGrad::SeluGradOneDimOp<D_T_X>>(gradients, outputs, y, &tilingData);
45+ } else if constexpr (SCH_MODE == SELU_GRAD_BROADCAST) {
46+ RunSeluGrad<NsSeluGrad::SeluGradBroadcastOp<D_T_X>>(gradients, outputs, y, &tilingData);
47+ }
48+}
Mactivation/selu_grad/tests/ut/op_host/op_api/test_aclnn_selubackward.cpp+1-1
@@ -14,7 +14,7 @@
14#include <vector>14#include <vector>
15 15 
16#include "gtest/gtest.h"16#include "gtest/gtest.h"
17-#include "../../../../op_host/op_api/aclnn_selu_backward.h"17+#include "../../../../op_api/aclnn_selu_backward.h"
18#include "op_api_ut_common/op_api_ut.h"18#include "op_api_ut_common/op_api_ut.h"
19#include "op_api_ut_common/scalar_desc.h"19#include "op_api_ut_common/scalar_desc.h"
20#include "op_api_ut_common/tensor_desc.h"20#include "op_api_ut_common/tensor_desc.h"
Aactivation/softsign/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+# Generated By CANNBot
11+ 
12+# 设置算子定义时支持的芯片类型
13+set(SUPPORT_COMPUTE_UNIT "ascend950")
14+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
15+set(SUPPORT_TILING_DIR "arch35")
16+ 
17+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE softsign ACLNNTYPE aclnn COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Aactivation/softsign/README.md+81-0
@@ -0,0 +1,81 @@
1+# Softsign
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------- | :------:|
7+| Ascend 950PR/Ascend 950DT | √ |
8+| Atlas A3 训练系列产品/Atlas A3 推理系列产品 | √ |
9+| Atlas A2 训练系列产品/Atlas A2 推理系列产品 | √ |
10+| Atlas 200I/500 A2 推理产品 | × |
11+| Atlas 推理系列产品 | √ |
12+| Atlas 训练系列产品 | √ |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：完成 Softsign 激活函数计算，对输入张量的每个元素逐元素计算 softsign 值。
17+- 计算公式：
18+ 
19+ $$
20+ y = \frac{x}{1 + |x|}
21+ $$
22+ 
23+ 其中 $x$ 为输入张量，$y$ 为输出张量，输出值域为 $(-1, 1)$。
24+ 
25+- 数值特性：公式天然数值稳定，分母 $1 + |x| \geq 1$，无除零风险，无需 epsilon 保护。
26+ 
27+## 参数说明
28+ 
29+<table style="table-layout: fixed; width: 1576px"><colgroup>
30+<col style="width: 170px">
31+<col style="width: 170px">
32+<col style="width: 200px">
33+<col style="width: 200px">
34+<col style="width: 170px">
35+</colgroup>
36+<thead>
37+ <tr>
38+ <th>参数名</th>
39+ <th>输入/输出/属性</th>
40+ <th>描述</th>
41+ <th>数据类型</th>
42+ <th>数据格式</th>
43+ </tr></thead>
44+<tbody>
45+ <tr>
46+ <td>x</td>
47+ <td>输入</td>
48+ <td>公式中的输入 x，任意形状张量。</td>
49+ <td>FLOAT、FLOAT16、BFLOAT16</td>
50+ <td>ND</td>
51+ </tr>
52+ <tr>
53+ <td>y</td>
54+ <td>输出</td>
55+ <td>公式中的输出 y，shape 和 dtype 与输入 x 完全一致，值域 (-1, 1)。</td>
56+ <td>FLOAT、FLOAT16、BFLOAT16</td>
57+ <td>ND</td>
58+ </tr>
59+</tbody></table>
60+ 
61+## 约束说明
62+ 
63+- 输入维度限制：输入张量维度不超过 8 维。
64+- 数据类型限制：仅支持 FLOAT、FLOAT16、BFLOAT16，不支持 DOUBLE。
65+- FP16/BF16 精度说明：FP16 和 BF16 类型在内部通过 Cast 到 FP32 进行中间计算，再 Cast 回原始类型，以保证计算精度。
66+ 
67+## 调用说明
68+ 
69+<table><thead>
70+ <tr>
71+ <th>调用方式</th>
72+ <th>调用样例</th>
73+ <th>说明</th>
74+ </tr></thead>
75+<tbody>
76+ <tr>
77+ <td> 图模式调用 </td>
78+ <td><a href="./examples/arch35/test_geir_softsign.cpp">test_geir_softsign</a></td>
79+ <td> 通过[算子IR](./op_graph/softsign_proto.h)构图方式调用Softsign算子。 </td>
80+ </tr>
81+</tbody></table>
Aactivation/softsign/examples/arch35/test_geir_softsign.cpp+249-0
@@ -0,0 +1,249 @@
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 <fstream>
13+#include <string.h>
14+#include <stdint.h>
15+#include <vector>
16+#include <string>
17+#include <map>
18+#include "assert.h"
19+ 
20+#include "graph.h"
21+#include "types.h"
22+#include "tensor.h"
23+#include "ge_error_codes.h"
24+#include "ge_api_types.h"
25+#include "ge_api.h"
26+#include "array_ops.h"
27+#include "ge_ir_build.h"
28+#include "../../op_graph/softsign_proto.h"
29+ 
30+#define FAILED -1
31+#define SUCCESS 0
32+ 
33+using namespace ge;
34+using std::map;
35+using std::string;
36+using std::vector;
37+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
38+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
39+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
40+ TensorDesc placeholder##intputIndex##_desc = \
41+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
42+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
43+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
44+ Tensor tensor_placeholder##intputIndex; \
45+ ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, \
46+ tensor_placeholder##intputIndex, \
47+ placeholder##intputIndex##_desc, \
48+ 2); \
49+ if (ret != SUCCESS) { \
50+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
51+ return FAILED; \
52+ } \
53+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
54+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
55+ input.push_back(tensor_placeholder##intputIndex); \
56+ graph.AddOp(placeholder##intputIndex); \
57+ softsign_op.set_input_##intputName(placeholder##intputIndex); \
58+ inputs.push_back(placeholder##intputIndex);
59+ 
60+#define LOG_PRINT(message, ...) \
61+ do { \
62+ printf(message, ##__VA_ARGS__); \
63+ } while (0)
64+ 
65+string GetTime()
66+{
67+ time_t timep;
68+ time(&timep);
69+ char tmp[64];
70+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
71+ return tmp;
72+}
73+ 
74+uint32_t GetDataTypeSize(DataType dt)
75+{
76+ uint32_t oneByte = 1;
77+ uint32_t twoByte = 2;
78+ uint32_t fourByte = 4;
79+ uint32_t eightByte = 8;
80+ 
81+ if (dt == ge::DT_FLOAT) {
82+ return fourByte;
83+ } else if (dt == ge::DT_FLOAT16) {
84+ return twoByte;
85+ } else if (dt == ge::DT_BF16) {
86+ return twoByte;
87+ } else if (dt == ge::DT_INT32) {
88+ return fourByte;
89+ } else if (dt == ge::DT_INT64) {
90+ return eightByte;
91+ } else if (dt == ge::DT_INT8) {
92+ return oneByte;
93+ }
94+ return fourByte;
95+}
96+ 
97+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, float value)
98+{
99+ input_tensor_desc.SetRealDimCnt(shapes.size());
100+ size_t size = 1;
101+ for (uint32_t i = 0; i < shapes.size(); i++) {
102+ size *= shapes[i];
103+ }
104+ uint32_t byteSizeFloat32 = 4;
105+ uint32_t data_len = size * byteSizeFloat32;
106+ float *pData = new (std::nothrow) float[size];
107+ if (pData == nullptr) {
108+ return FAILED;
109+ }
110+ 
111+ for (size_t i = 0; i < size; ++i) {
112+ *(pData + i) = value;
113+ }
114+ input_tensor = Tensor(input_tensor_desc, (uint8_t *)pData, data_len);
115+ delete[] pData;
116+ return SUCCESS;
117+}
118+ 
119+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t *inputData)
120+{
121+ FILE *fp;
122+ fp = fopen(bin_file.c_str(), "w");
123+ if (fp == nullptr) {
124+ return FAILED;
125+ }
126+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
127+ fclose(fp);
128+ return SUCCESS;
129+}
130+ 
131+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
132+ std::vector<Operator> &outputs, Graph &graph)
133+{
134+ Status ret = SUCCESS;
135+ auto softsign_op = op::Softsign("softsign");
136+ std::vector<int64_t> xShape = {4, 4};
137+ ADD_INPUT(1, x, inDtype, xShape);
138+ 
139+ outputs.push_back(softsign_op);
140+ return SUCCESS;
141+}
142+ 
143+int main(int argc, char *argv[])
144+{
145+ const char *graph_name = "tc_ge_irrun_test";
146+ Graph graph(graph_name);
147+ std::vector<ge::Tensor> input;
148+ 
149+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
150+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
151+ Status ret = ge::GEInitialize(global_options);
152+ if (ret != SUCCESS) {
153+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
154+ return FAILED;
155+ }
156+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
157+ 
158+ std::vector<Operator> inputs{};
159+ std::vector<Operator> outputs{};
160+ 
161+ DataType inDtype = DT_FLOAT;
162+ if (argc > 1) {
163+ std::cout << argv[1] << std::endl;
164+ }
165+ std::cout << inDtype << std::endl;
166+ 
167+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
168+ if (ret != SUCCESS) {
169+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
170+ return FAILED;
171+ }
172+ 
173+ if (!inputs.empty() && !outputs.empty()) {
174+ graph.SetInputs(inputs).SetOutputs(outputs);
175+ }
176+ 
177+ std::map<AscendString, AscendString> build_options = {};
178+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
179+ ge::Session *session = new Session(build_options);
180+ 
181+ if (session == nullptr) {
182+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
183+ return FAILED;
184+ }
185+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
186+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
187+ 
188+ std::map<AscendString, AscendString> graph_options = {};
189+ uint32_t graph_id = 0;
190+ ret = session->AddGraph(graph_id, graph, graph_options);
191+ 
192+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
193+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
194+ std::string file_path = "./dump";
195+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
196+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
197+ std::vector<ge::Tensor> output;
198+ ret = session->RunGraph(graph_id, input, output);
199+ if (ret != SUCCESS) {
200+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
201+ delete session;
202+ GEFinalize();
203+ return FAILED;
204+ }
205+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
206+ 
207+ int input_num = input.size();
208+ for (int i = 0; i < input_num; i++) {
209+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
210+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
211+ uint8_t *input_data_i = input[i].GetData();
212+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
213+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
214+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
215+ WriteDataToFile((const char *)input_file.c_str(), data_size, input_data_i);
216+ }
217+ 
218+ int output_num = output.size();
219+ for (int i = 0; i < output_num; i++) {
220+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
221+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
222+ uint8_t *output_data_i = output[i].GetData();
223+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
224+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
225+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
226+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
227+ float *resultData = (float*)output_data_i;
228+ for (int64_t j = 0; j < output_shape; j++) {
229+ LOG_PRINT("result[%ld] is: %f\n", j, resultData[j]);
230+ }
231+ }
232+ 
233+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
234+ std::string error_str(error_msg.GetString());
235+ std::cout << "Error message: " << error_str << std::endl;
236+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
237+ std::string warning_str(warning_msg.GetString());
238+ std::cout << "Warning message: " << warning_str << std::endl;
239+ delete session;
240+ session = nullptr;
241+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
242+ ret = ge::GEFinalize();
243+ if (ret != SUCCESS) {
244+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
245+ return FAILED;
246+ }
247+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
248+ return SUCCESS;
249+}
Aactivation/softsign/op_graph/softsign_proto.h+49-0
@@ -0,0 +1,49 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file softsign_proto.h
15+ * \brief Softsign 算子 GE IR 图模式注册
16+ *
17+ * 算子功能：y = x / (1 + |x|)
18+ * 输入：x（float16, bfloat16, float32）
19+ * 输出：y（与输入 dtype/shape 一致）
20+ */
21+#ifndef OPS_OP_PROTO_INC_SOFTSIGN_H_
22+#define OPS_OP_PROTO_INC_SOFTSIGN_H_
23+ 
24+#include "graph/operator_reg.h"
25+#include "graph/types.h"
26+ 
27+namespace ge {
28+ 
29+/**
30+* @brief Computes softsign: x/(abs(x) + 1) .
31+ 
32+*@par Inputs:
33+* One input:
34+*x: A Tensor. Support 1D ~ 8D. Must be one of the following types: bfloat16, float16, float32 or double.
35+ 
36+*@par Outputs:
37+*y: The activations tensor. Has the same type and format as "x"
38+ 
39+*@par Third-party framework compatibility
40+* Compatible with the TensorFlow operator Softsign.
41+*/
42+REG_OP(Softsign)
43+ .INPUT(x, TensorType({FloatingDataType, DT_BF16}))
44+ .OUTPUT(y, TensorType({FloatingDataType, DT_BF16}))
45+ .OP_END_FACTORY_REG(Softsign)
46+ 
47+} // namespace ge
48+ 
49+#endif // OPS_OP_PROTO_INC_SOFTSIGN_H_
Aactivation/softsign/op_host/arch35/softsign_tiling_arch35.cpp+122-0
@@ -0,0 +1,122 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file softsign_tiling_arch35.cpp
15+ * \brief Softsign 算子 Host Tiling 实现（atvoss 框架 - Elewise 模式, arch35）
16+ *
17+ * Tiling 策略：
18+ * - FP32 (TilingKey=103): 直接计算，无 Cast
19+ * - FP16 (TilingKey=101): Cast→FP32→计算→Cast→FP16
20+ * - BF16 (TilingKey=102): Cast→FP32→计算→Cast→BF16
21+ */
22+ 
23+#include "register/op_def_registry.h"
24+#include "op_common/log/log.h"
25+#include "atvoss/elewise/elewise_tiling.h"
26+#include "../../op_kernel/arch35/softsign_dag.h"
27+#include "../../op_kernel/arch35/softsign_struct.h"
28+ 
29+namespace optiling {
30+ 
31+using namespace ge;
32+using Ops::Base::ElewiseBaseTiling;
33+using Ops::Base::half;
34+using Ops::Base::bfloat16_t;
35+ 
36+// TilingKey 定义（与 DESIGN.md §3.1 一致）
37+constexpr uint64_t TILING_KEY_FP16 = 101UL;
38+constexpr uint64_t TILING_KEY_BF16 = 102UL;
39+constexpr uint64_t TILING_KEY_FP32 = 103UL;
40+ 
41+static ge::graphStatus DoTilingByDtype(gert::TilingContext* context, ge::DataType dtype,
42+ SoftsignTilingData* tilingData, uint64_t& tilingKey)
43+{
44+ ElewiseBaseTiling eleTiling(context);
45+ ge::graphStatus ret;
46+ if (dtype == ge::DT_FLOAT) {
47+ using OpDag = SoftsignOp::GraphSoftsign<float, float>::OpDag;
48+ ret = eleTiling.DoTiling<OpDag>(tilingData->baseTiling);
49+ tilingKey = TILING_KEY_FP32;
50+ } else if (dtype == ge::DT_FLOAT16) {
51+ using OpDag = SoftsignOp::GraphSoftsign<half, float>::OpDag;
52+ ret = eleTiling.DoTiling<OpDag>(tilingData->baseTiling);
53+ tilingKey = TILING_KEY_FP16;
54+ } else if (dtype == ge::DT_BF16) {
55+ using OpDag = SoftsignOp::GraphSoftsign<bfloat16_t, float>::OpDag;
56+ ret = eleTiling.DoTiling<OpDag>(tilingData->baseTiling);
57+ tilingKey = TILING_KEY_BF16;
58+ } else {
59+ OP_LOGE(context, "Softsign: unsupported dtype=%d", static_cast<int>(dtype));
60+ return ge::GRAPH_FAILED;
61+ }
62+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS,
63+ OP_LOGE(context, "Softsign: DoTiling failed for dtype=%d", static_cast<int>(dtype)),
64+ return ret);
65+ return ge::GRAPH_SUCCESS;
66+}
67+ 
68+static ge::graphStatus PrepareWorkspace(gert::TilingContext* context)
69+{
70+ constexpr uint64_t WORKSPACE_RESERVE_BYTE = 16UL * 1024 * 1024;
71+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
72+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
73+ currentWorkspace[0] = WORKSPACE_RESERVE_BYTE;
74+ return ge::GRAPH_SUCCESS;
75+}
76+ 
77+static ge::graphStatus SoftsignTilingFunc(gert::TilingContext* context)
78+{
79+ auto inputDesc = context->GetInputDesc(0);
80+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
81+ ge::DataType dtype = inputDesc->GetDataType();
82+ 
83+ auto inputShape = context->GetInputShape(0);
84+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
85+ auto storageShape = inputShape->GetStorageShape();
86+ int64_t dim0 = (storageShape.GetDimNum() == 0) ? 1 : storageShape.GetShapeSize();
87+ 
88+ OP_CHECK_IF(PrepareWorkspace(context) != ge::GRAPH_SUCCESS,
89+ OP_LOGE(context, "Softsign: PrepareWorkspace failed"),
90+ return ge::GRAPH_FAILED);
91+ 
92+ if (dim0 == 0) {
93+ context->SetBlockDim(1);
94+ context->SetTilingKey(TILING_KEY_FP32);
95+ return ge::GRAPH_SUCCESS;
96+ }
97+ 
98+ uint64_t tilingKey = 0;
99+ auto tilingData = context->GetTilingData<SoftsignTilingData>();
100+ OP_CHECK_NULL_WITH_CONTEXT(context, tilingData);
101+ auto ret = DoTilingByDtype(context, dtype, tilingData, tilingKey);
102+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS,
103+ OP_LOGE(context, "Softsign: DoTilingByDtype failed"),
104+ return ret);
105+ 
106+ context->SetTilingKey(tilingKey);
107+ context->SetBlockDim(tilingData->baseTiling.blockNum);
108+ return ge::GRAPH_SUCCESS;
109+}
110+ 
111+static ge::graphStatus TilingParseForSoftsign([[maybe_unused]] gert::TilingParseContext* context)
112+{
113+ return ge::GRAPH_SUCCESS;
114+}
115+ 
116+struct SoftsignCompileInfo {};
117+ 
118+IMPL_OP_OPTILING(Softsign)
119+ .Tiling(SoftsignTilingFunc)
120+ .TilingParse<SoftsignCompileInfo>(TilingParseForSoftsign);
121+ 
122+} // namespace optiling
Aactivation/softsign/op_host/config/ascend950/softsign_binary.json+92-0
@@ -0,0 +1,92 @@
1+{
2+ "op_type": "Softsign",
3+ "op_list": [
4+ {
5+ "bin_filename": "Softsign_cf4dcbc13fa8539d71027ffc6922d575",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ }
18+ ],
19+ "outputs": [
20+ {
21+ "name": "y",
22+ "index": 0,
23+ "dtype": "float16",
24+ "format": "ND",
25+ "paramType": "required",
26+ "shape": [
27+ -2
28+ ],
29+ "format_match_mode": "FormatAgnostic"
30+ }
31+ ]
32+ },
33+ {
34+ "bin_filename": "Softsign_c386cb844329b3f1661c792e5d562d25",
35+ "inputs": [
36+ {
37+ "name": "x",
38+ "index": 0,
39+ "dtype": "bfloat16",
40+ "format": "ND",
41+ "paramType": "required",
42+ "shape": [
43+ -2
44+ ],
45+ "format_match_mode": "FormatAgnostic"
46+ }
47+ ],
48+ "outputs": [
49+ {
50+ "name": "y",
51+ "index": 0,
52+ "dtype": "bfloat16",
53+ "format": "ND",
54+ "paramType": "required",
55+ "shape": [
56+ -2
57+ ],
58+ "format_match_mode": "FormatAgnostic"
59+ }
60+ ]
61+ },
62+ {
63+ "bin_filename": "Softsign_4f0f237147fbd384bbd972faabe2557f",
64+ "inputs": [
65+ {
66+ "name": "x",
67+ "index": 0,
68+ "dtype": "float32",
69+ "format": "ND",
70+ "paramType": "required",
71+ "shape": [
72+ -2
73+ ],
74+ "format_match_mode": "FormatAgnostic"
75+ }
76+ ],
77+ "outputs": [
78+ {
79+ "name": "y",
80+ "index": 0,
81+ "dtype": "float32",
82+ "format": "ND",
83+ "paramType": "required",
84+ "shape": [
85+ -2
86+ ],
87+ "format_match_mode": "FormatAgnostic"
88+ }
89+ ]
90+ }
91+ ]
92+}
Aactivation/softsign/op_host/softsign_def.cpp+45-0
@@ -0,0 +1,45 @@
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+/* Generated By CANNBot */
12+ 
13+#include "register/op_def_registry.h"
14+ 
15+namespace ops {
16+class Softsign : public OpDef {
17+public:
18+ explicit Softsign(const char* name) : OpDef(name)
19+ {
20+ this->Input("x")
21+ .ParamType(REQUIRED)
22+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
23+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
24+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
25+ .AutoContiguous();
26+ this->Output("y")
27+ .ParamType(REQUIRED)
28+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
29+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
30+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
31+ .AutoContiguous();
32+ 
33+ OpAICoreConfig aicoreConfig950;
34+ aicoreConfig950.DynamicCompileStaticFlag(true)
35+ .DynamicFormatFlag(false)
36+ .DynamicRankSupportFlag(true)
37+ .DynamicShapeSupportFlag(true)
38+ .NeedCheckSupportFlag(false)
39+ .PrecisionReduceFlag(true)
40+ .ExtendCfgInfo("opFile.value", "softsign");
41+ this->AICore().AddConfig("ascend950", aicoreConfig950);
42+ }
43+};
44+OP_ADD(Softsign);
45+} // namespace ops
Aactivation/softsign/op_host/softsign_infershape.cpp+46-0
@@ -0,0 +1,46 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file softsign_infershape.cpp
15+ * \brief Softsign 算子形状推导实现
16+ *
17+ * 逐元素运算，输出 shape = 输入 shape
18+ */
19+ 
20+#include "register/op_impl_registry.h"
21+#include "exe_graph/runtime/infer_shape_context.h"
22+#include "op_common/log/log.h"
23+ 
24+using namespace ge;
25+ 
26+namespace ops {
27+ 
28+// 实现 element-wise 算子的通用形状推导
29+// 逻辑：输出形状 = 输入形状
30+static ge::graphStatus InferShape4Softsign(gert::InferShapeContext* context)
31+{
32+ const gert::Shape* input_shape = context->GetInputShape(0);
33+ OP_CHECK_NULL_WITH_CONTEXT(context, input_shape);
34+ 
35+ gert::Shape* output_shape = context->GetOutputShape(0);
36+ OP_CHECK_NULL_WITH_CONTEXT(context, output_shape);
37+ 
38+ // 设置输出形状 = 输入形状
39+ *output_shape = *input_shape;
40+ 
41+ return ge::GRAPH_SUCCESS;
42+}
43+ 
44+IMPL_OP_INFERSHAPE(Softsign).InferShape(InferShape4Softsign);
45+ 
46+} // namespace ops
Aactivation/softsign/op_kernel/arch35/softsign_dag.h+67-0
@@ -0,0 +1,67 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file softsign_dag.h
15+ * \brief Softsign 算子 DAG 计算图定义（atvoss 框架 - Elewise 模式）
16+ *
17+ * 计算公式: y = x / (1 + |x|)
18+ *
19+ * 数据流 (FP32):
20+ * x (GM) -> CopyIn -> SaveX(Vec::Copy) -----\
21+ * | -> Div -> CopyOut -> y (GM)
22+ * -> Abs -> Adds(+1) -----/
23+ *
24+ * 数据流 (FP16/BF16):
25+ * x (GM) -> CopyIn -> CastIn(→FP32) -> SaveX(Vec::Copy) -----\
26+ * | -> Div -> CastOut(→U) -> CopyOut -> y (GM)
27+ * -> Abs -> Adds(+1) ---------/
28+ */
29+ 
30+#ifndef SOFTSIGN_DAG_H
31+#define SOFTSIGN_DAG_H
32+ 
33+// Host 编译时 mock __aicore__（Kernel 编译器已内置定义）
34+#ifndef __CCE_AICORE__
35+#ifndef __aicore__
36+#define __aicore__
37+#endif
38+#endif
39+ 
40+#include "atvoss/util/dag.h"
41+#include "atvoss/util/vec.h"
42+#include "atvoss/util/placeholder.h"
43+ 
44+namespace SoftsignOp {
45+ 
46+using namespace Ops::Base;
47+ 
48+template <typename U, typename T = float>
49+struct GraphSoftsign {
50+ using ConstOne = MAKE_CONST(float, 1);
51+ 
52+ using OpCopyIn = Bind<Vec::CopyIn<U>, Placeholder::In0<U>>;
53+ using OpCastIn = Bind<Vec::Cast<T, U, 0>, OpCopyIn>;
54+ using OpSaveX = Bind<Vec::Copy<T>, OpCastIn>;
55+ using OpAbs = Bind<Vec::Abs<T>, OpCastIn>;
56+ using OpAdds = Bind<Vec::Adds<T>, OpAbs, ConstOne>;
57+ using OpDiv = Bind<Vec::Div<T>, OpSaveX, OpAdds>;
58+ using OpCastOut = Bind<Vec::Cast<U, T, 1>, OpDiv>;
59+ using OpCopyOut = Bind<Vec::CopyOut<U>, Placeholder::Out0<U>, OpCastOut>;
60+ 
61+ using Outputs = Elems<OpCopyOut>;
62+ using OpDag = DAGSch<Outputs>;
63+};
64+ 
65+} // namespace SoftsignOp
66+ 
67+#endif // SOFTSIGN_DAG_H
Aactivation/softsign/op_kernel/arch35/softsign_struct.h+30-0
@@ -0,0 +1,30 @@
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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file softsign_struct.h
15+ * \brief Softsign 算子 TilingData 和 TilingKey 定义（atvoss 框架 - Elewise 模式）
16+ *
17+ * TilingData: 使用 EleBaseTilingData16B（atvoss Elewise 标准结构体）
18+ * TilingKey: 按 dtype 区分（FP16=101, BF16=102, FP32=103）
19+ */
20+ 
21+#ifndef SOFTSIGN_STRUCT_H
22+#define SOFTSIGN_STRUCT_H
23+ 
24+#include "atvoss/elewise/elewise_base_struct.h"
25+ 
26+struct SoftsignTilingData {
27+ Ops::Base::EleBaseTilingData baseTiling;
28+};
29+ 
30+#endif // SOFTSIGN_STRUCT_H
Aactivation/softsign/op_kernel/softsign.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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file softsign.cpp
15+ * \brief Softsign 算子 Kernel 入口（atvoss 框架 - Elewise 模式）
16+ *
17+ * 按 TilingKey 分发到对应 dtype 的 DAG 模板：
18+ * - TilingKey 101: FP16 → Cast→FP32→计算→Cast→FP16
19+ * - TilingKey 102: BF16 → Cast→FP32→计算→Cast→BF16
20+ * - TilingKey 103: FP32 → 直接计算
21+ */
22+ 
23+#include "kernel_operator.h"
24+#include "kernel_tiling/kernel_tiling.h"
25+#include "atvoss/elewise/elewise_sch.h"
26+#include "arch35/softsign_dag.h"
27+#include "arch35/softsign_struct.h"
28+ 
29+using namespace AscendC;
30+using namespace SoftsignOp;
31+ 
32+extern "C" __global__ __aicore__ void softsign(GM_ADDR x, GM_ADDR y,
33+ GM_ADDR workspace, GM_ADDR tiling)
34+{
35+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
36+ REGISTER_TILING_DEFAULT(SoftsignTilingData);
37+ GET_TILING_DATA_WITH_STRUCT(SoftsignTilingData, tilingData, tiling);
38+ TPipe pipe;
39+ if (TILING_KEY_IS(103UL)) {
范其瑞
范其瑞范其瑞6月27日

不能把必选输入的dtype信息放到tilingKey里，信息冗余了。不同.o的dtype是不同的，.o内部就不需要再次区分dtype维度。另外新开发算子不能再使用TILING_KEY_IS

likedislike
40+ ElementwiseSch<0UL, GraphSoftsign<float, float>::OpDag> sch(&(tilingData.baseTiling), &pipe);
41+ sch.Init(x, y);
42+ sch.Process();
43+ } else if (TILING_KEY_IS(101UL)) {
44+ ElementwiseSch<0UL, GraphSoftsign<half, float>::OpDag> sch(&(tilingData.baseTiling), &pipe);
45+ sch.Init(x, y);
46+ sch.Process();
47+ } else if (TILING_KEY_IS(102UL)) {
48+ ElementwiseSch<0UL, GraphSoftsign<bfloat16_t, float>::OpDag> sch(&(tilingData.baseTiling), &pipe);
49+ sch.Init(x, y);
50+ sch.Process();
51+ }
52+}
Aactivation/softsign/tests/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# -----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 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+file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
13+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
14+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aactivation/softsign/tests/ut/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# -----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 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+file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
13+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
14+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aactivation/softsign/tests/ut/op_host/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+# -----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 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+file(GLOB CURRENT_DIR RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_HOST_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
15+endif()
Aactivation/softsign/tests/ut/op_host/test_softsign_infershape.cpp+186-0
@@ -0,0 +1,186 @@
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 <gtest/gtest.h>
13+#include "register/op_impl_registry.h"
14+#include "kernel_run_context_facker.h"
15+#include "../../../op_graph/softsign_proto.h"
16+#include "exe_graph/runtime/storage_format.h"
17+#include "exe_graph/runtime/storage_shape.h"
18+#include "log/log.h"
19+#include "platform/platform_info.h"
20+ 
21+class SoftsignProtoTest : public testing::Test {
22+ protected:
23+ static void SetUpTestCase() {
24+ std::cout << "Softsign Proto Test SetUp" << std::endl;
25+ }
26+ 
27+ static void TearDownTestCase() {
28+ std::cout << "Softsign Proto Test TearDown" << std::endl;
29+ }
30+};
31+ 
32+TEST_F(SoftsignProtoTest, softsign_infershape_1d_fp32_test) {
33+ fe::PlatformInfo platformInfo;
34+ fe::OptionalInfo optiCompilationInfo;
35+ platformInfo.soc_info.ai_core_cnt = 64;
36+ platformInfo.str_info.short_soc_version = "Ascend950";
37+ optiCompilationInfo.soc_version = "Ascend950";
38+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend950"] = platformInfo;
39+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
40+ 
41+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Softsign")->infer_shape;
42+ 
43+ gert::Shape input_shape_0 = {128};
44+ gert::Shape output_shape_0 = {};
45+ 
46+ auto holder = gert::InferShapeContextFaker()
47+ .NodeIoNum(1, 1)
48+ .IrInstanceNum({1, 1})
49+ .InputShapes({&input_shape_0})
50+ .OutputShapes({&output_shape_0})
51+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
52+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
53+ .Build();
54+ 
55+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
56+}
57+ 
58+TEST_F(SoftsignProtoTest, softsign_infershape_2d_fp16_test) {
59+ fe::PlatformInfo platformInfo;
60+ fe::OptionalInfo optiCompilationInfo;
61+ platformInfo.soc_info.ai_core_cnt = 64;
62+ platformInfo.str_info.short_soc_version = "Ascend950";
63+ optiCompilationInfo.soc_version = "Ascend950";
64+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend950"] = platformInfo;
65+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
66+ 
67+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Softsign")->infer_shape;
68+ 
69+ gert::Shape input_shape_0 = {8, 16};
70+ gert::Shape output_shape_0 = {};
71+ 
72+ auto holder = gert::InferShapeContextFaker()
73+ .NodeIoNum(1, 1)
74+ .IrInstanceNum({1, 1})
75+ .InputShapes({&input_shape_0})
76+ .OutputShapes({&output_shape_0})
77+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
78+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
79+ .Build();
80+ 
81+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
82+}
83+ 
84+TEST_F(SoftsignProtoTest, softsign_infershape_3d_bf16_test) {
85+ fe::PlatformInfo platformInfo;
86+ fe::OptionalInfo optiCompilationInfo;
87+ platformInfo.soc_info.ai_core_cnt = 64;
88+ platformInfo.str_info.short_soc_version = "Ascend950";
89+ optiCompilationInfo.soc_version = "Ascend950";
90+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend950"] = platformInfo;
91+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
92+ 
93+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Softsign")->infer_shape;
94+ 
95+ gert::Shape input_shape_0 = {4, 3, 4};
96+ gert::Shape output_shape_0 = {};
97+ 
98+ auto holder = gert::InferShapeContextFaker()
99+ .NodeIoNum(1, 1)
100+ .IrInstanceNum({1, 1})
101+ .InputShapes({&input_shape_0})
102+ .OutputShapes({&output_shape_0})
103+ .NodeInputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
104+ .NodeOutputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
105+ .Build();
106+ 
107+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
108+}
109+ 
110+TEST_F(SoftsignProtoTest, softsign_infershape_scalar_test) {
111+ fe::PlatformInfo platformInfo;
112+ fe::OptionalInfo optiCompilationInfo;
113+ platformInfo.soc_info.ai_core_cnt = 64;
114+ platformInfo.str_info.short_soc_version = "Ascend950";
115+ optiCompilationInfo.soc_version = "Ascend950";
116+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend950"] = platformInfo;
117+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
118+ 
119+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Softsign")->infer_shape;
120+ 
121+ gert::Shape input_shape_0 = {};
122+ gert::Shape output_shape_0 = {};
123+ 
124+ auto holder = gert::InferShapeContextFaker()
125+ .NodeIoNum(1, 1)
126+ .IrInstanceNum({1, 1})
127+ .InputShapes({&input_shape_0})
128+ .OutputShapes({&output_shape_0})
129+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
130+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
131+ .Build();
132+ 
133+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
134+}
135+ 
136+TEST_F(SoftsignProtoTest, softsign_infershape_empty_test) {
137+ fe::PlatformInfo platformInfo;
138+ fe::OptionalInfo optiCompilationInfo;
139+ platformInfo.soc_info.ai_core_cnt = 64;
140+ platformInfo.str_info.short_soc_version = "Ascend950";
141+ optiCompilationInfo.soc_version = "Ascend950";
142+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend950"] = platformInfo;
143+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
144+ 
145+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Softsign")->infer_shape;
146+ 
147+ gert::Shape input_shape_0 = {0};
148+ gert::Shape output_shape_0 = {};
149+ 
150+ auto holder = gert::InferShapeContextFaker()
151+ .NodeIoNum(1, 1)
152+ .IrInstanceNum({1, 1})
153+ .InputShapes({&input_shape_0})
154+ .OutputShapes({&output_shape_0})
155+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
156+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
157+ .Build();
158+ 
159+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
160+}
161+ 
162+TEST_F(SoftsignProtoTest, softsign_infershape_high_rank_test) {
163+ fe::PlatformInfo platformInfo;
164+ fe::OptionalInfo optiCompilationInfo;
165+ platformInfo.soc_info.ai_core_cnt = 64;
166+ platformInfo.str_info.short_soc_version = "Ascend950";
167+ optiCompilationInfo.soc_version = "Ascend950";
168+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend950"] = platformInfo;
169+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
170+ 
171+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Softsign")->infer_shape;
172+ 
173+ gert::Shape input_shape_0 = {2, 3, 4, 5, 6};
174+ gert::Shape output_shape_0 = {};
175+ 
176+ auto holder = gert::InferShapeContextFaker()
177+ .NodeIoNum(1, 1)
178+ .IrInstanceNum({1, 1})
179+ .InputShapes({&input_shape_0})
180+ .OutputShapes({&output_shape_0})
181+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
182+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
183+ .Build();
184+ 
185+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
186+}
Mdocs/zh/op_api_list.md+1-1
@@ -388,7 +388,7 @@
388| [aclnnScaledMaskedSoftmax](../../vfusion/scaled_masked_softmax_v2/docs/aclnnScaledMaskedSoftmax.md) | 将输入的数据x先进行scale缩放和mask，然后执行softmax的输出。 | 默认确定性实现 | - |388| [aclnnScaledMaskedSoftmax](../../vfusion/scaled_masked_softmax_v2/docs/aclnnScaledMaskedSoftmax.md) | 将输入的数据x先进行scale缩放和mask，然后执行softmax的输出。 | 默认确定性实现 | - |
389| [aclnnScaledMaskedSoftmaxBackward](../../vfusion/scaled_masked_softmax_grad_v2/docs/aclnnScaledMaskedSoftmaxBackward.md) | softmax的反向传播，并对结果进行缩放以及掩码。 | 默认非确定性实现，支持配置开启。 | - |389| [aclnnScaledMaskedSoftmaxBackward](../../vfusion/scaled_masked_softmax_grad_v2/docs/aclnnScaledMaskedSoftmaxBackward.md) | softmax的反向传播，并对结果进行缩放以及掩码。 | 默认非确定性实现，支持配置开启。 | - |
390| [aclnnSelu&aclnnInplaceSelu](../../activation/selu/docs/aclnnSelu&aclnnInplaceSelu.md) | 对输入Tensor逐元素计算SELU（Scaled Exponential Linear Unit）激活函数。 | 默认确定性实现 | - |390| [aclnnSelu&aclnnInplaceSelu](../../activation/selu/docs/aclnnSelu&aclnnInplaceSelu.md) | 对输入Tensor逐元素计算SELU（Scaled Exponential Linear Unit）激活函数。 | 默认确定性实现 | - |
391-| [aclnnSeluBackward](../../activation/selu_grad/docs/aclnnSeluBackward.md) | 完成aclnnSelu的反向。 | 默认确定性实现 | - |391+| [aclnnSeluBackward](../../activation/selu_grad/docs/aclnnSeluBackward.md) | 完成aclnnSelu的反向。 | 默认确定性实现 | 默认确定性实现 |
392| [aclnnShrink](../../activation/shrink/docs/aclnnShrink.md) | 对输入张量进行非线性变换，根据输入值self与阈值lambd的关系，对输入通过偏移量bias进行缩放和偏移处理。 | 默认确定性实现 | 默认确定性实现 |392| [aclnnShrink](../../activation/shrink/docs/aclnnShrink.md) | 对输入张量进行非线性变换，根据输入值self与阈值lambd的关系，对输入通过偏移量bias进行缩放和偏移处理。 | 默认确定性实现 | 默认确定性实现 |
393| [aclnnSigmoid&aclnnInplaceSigmoid](../../activation/sigmoid/docs/aclnnSigmoid&aclnnInplaceSigmoid.md) | 对输入Tensor完成sigmoid运算。 | 默认确定性实现 | 默认确定性实现 |393| [aclnnSigmoid&aclnnInplaceSigmoid](../../activation/sigmoid/docs/aclnnSigmoid&aclnnInplaceSigmoid.md) | 对输入Tensor完成sigmoid运算。 | 默认确定性实现 | 默认确定性实现 |
394| [aclnnSigmoidBackward](../../activation/sigmoid_grad/docs/aclnnSigmoidBackward.md) | 完成sigmoid的反向传播，根据sigmoid反向传播梯度与正向输出计算sigmoid的梯度输入。 | 默认确定性实现 | 默认确定性实现 |394| [aclnnSigmoidBackward](../../activation/sigmoid_grad/docs/aclnnSigmoidBackward.md) | 完成sigmoid的反向传播，根据sigmoid反向传播梯度与正向输出计算sigmoid的梯度输入。 | 默认确定性实现 | 默认确定性实现 |
Mdocs/zh/op_list.md+15-5
@@ -520,12 +520,12 @@
520 <tr>520 <tr>
521 <td>activation</td>521 <td>activation</td>
522 <td><a href="../../activation/selu_grad/README.md">selu_grad</a></td>522 <td><a href="../../activation/selu_grad/README.md">selu_grad</a></td>
523- <td>✗</td>
524- <td>✗</td>
525 <td>✓</td>523 <td>✓</td>
526- <td>✗</td>524+ <td>✓</td>
525+ <td>✓</td>
526+ <td>✓</td>
527 <td>AI Core</td>527 <td>AI Core</td>
528- <td>该算子暂无Ascend C代码实现，欢迎开发者补充贡献，贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>528+ <td>完成SELU激活函数的反向梯度计算。给定上游梯度gradients和前向输出outputs，按元素计算输出梯度。</td>
529 </tr>529 </tr>
530 <tr>530 <tr>
531 <td>activation</td>531 <td>activation</td>
@@ -629,7 +629,7 @@
629 </tr>629 </tr>
630 <tr>630 <tr>
631 <td>activation</td>631 <td>activation</td>
632- <td><a href="../../activation/softsign_grad/README.md">softshrink</a></td>632+ <td><a href="../../activation/softshrink/README.md">softshrink</a></td>
633 <td>✓</td>633 <td>✓</td>
634 <td>✓</td>634 <td>✓</td>
635 <td>✗</td>635 <td>✗</td>
@@ -647,6 +647,16 @@
647 <td>AI Core</td>647 <td>AI Core</td>
648 <td>该算子暂无Ascend C代码实现，欢迎开发者补充贡献，贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>648 <td>该算子暂无Ascend C代码实现，欢迎开发者补充贡献，贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>
649 </tr>649 </tr>
650+ <tr>
651+ <td>activation</td>
652+ <td><a href="../../activation/softsign/README.md">softsign</a></td>
653+ <td>✓</td>
654+ <td>✓</td>
655+ <td>✗</td>
656+ <td>✓</td>
657+ <td>AI Core</td>
658+ <td>对输入张量应用Softsign激活函数：y = x / (1 + |x|)，输出值域为(-1, 1)。</td>
659+ </tr>
650 <tr>660 <tr>
651 <td>activation</td>661 <td>activation</td>
652 <td><a href="../../activation/softsign_grad/README.md">softsign_grad</a></td>662 <td><a href="../../activation/softsign_grad/README.md">softsign_grad</a></td>