Star869
1749
代码
代码
Issues217
Pull Requests310
Actions
项目讨论
Wiki
分析
Star869
1749
已合并
softplus算子贡献 #3225
松柏创建于 3月26日
softplus算子贡献 #3225
已合并
松柏创建于 3月26日
从ilovescrapy/softplusmaster合入到cann/ops-nnmaster
共 17 个文件变更+1085-0
Aexperimental/activation/softplus/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+ 
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+if(NOT ENABLE_TEST)
14+ list(REMOVE_ITEM CURRENT_DIRS tests)
15+endif()
16+foreach(SUB_DIR ${CURRENT_DIRS})
17+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18+ add_subdirectory(${SUB_DIR})
19+ endif()
20+endforeach()
Aexperimental/activation/softplus/README.md+68-0
@@ -0,0 +1,68 @@
1+# Softplus
2+ 
3+## 产品支持情况
4+ 
5+ 
6+| 产品 | 是否支持 |
7+| ------------------------------------------------------------------ | :------: |
8+| Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件 | √ |
9+ 
10+## 功能说明
11+ 
12+- 算子功能：求reluv2函数梯度。
13+- 计算公式：
14+ 
15+$$
16+y = log(1 + e^x)
17+$$
18+ 
19+## 参数说明
20+ 
21+<table style="undefined;table-layout: fixed; width: 820px"><colgroup>
22+ <col style="width: 100px">
23+ <col style="width: 150px">
24+ <col style="width: 190px">
25+ <col style="width: 260px">
26+ <col style="width: 120px">
27+ </colgroup>
28+ <thead>
29+ <tr>
30+ <th>参数名</th>
31+ <th>输入/输出/属性</th>
32+ <th>描述</th>
33+ <th>数据类型</th>
34+ <th>数据格式</th>
35+ </tr></thead>
36+ <tbody>
37+ <tr>
38+ <td>x</td>
39+ <td>输入</td>
40+ <td>反向传播梯度</td>
41+ <td>FLOAT、FLOAT16、BFLOAT16</td>
42+ <td>ND</td>
43+ </tr>
44+ <tr>
45+ <td>y</td>
46+ <td>输出</td>
47+ <td>公式中的输出张量</td>
48+ <td>FLOAT、FLOAT16、BFLOAT16</td>
49+ <td>ND</td>
50+ </tr>
51+ </tbody></table>
52+ 
53+## 调用说明
54+ 
55+| 调用方式 | 调用样例 | 说明 |
56+| --------- | ----------------------------------------------------------- | ---------------------------------------------------------------------- |
57+| aclnn调用 | [test_aclnn_softplus](./examples/test_aclnn_softplus.cpp) | 通过自动生成的AclnnSoftplus接口方式调用Softplus算子。 |
58+ 
59+## 约束说明
60+ 
61+无
62+ 
63+## 贡献说明
64+ 
65+ 
66+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
67+| ----------- | ---------- | ---------------- | ---------- | ------------------------------ |
68+| ilovescrapy | 个人开发者 | Softplus | 2026/3/21 | Softplus算子适配开源仓 |
Aexperimental/activation/softplus/examples/test_aclnn_softplus.cpp+161-0
@@ -0,0 +1,161 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <vector>
13+#include "acl/acl.h"
14+#include "aclnn_softplus.h"
15+// 修改测试数据类型
16+using DataType = float;
17+#define CHECK_RET(cond, return_expr) \
18+ do { \
19+ if (!(cond)) { \
20+ return_expr; \
21+ } \
22+ } while (0)
23+ 
24+#define LOG_PRINT(message, ...) \
25+ do { \
26+ printf(message, ##__VA_ARGS__); \
27+ } while (0)
28+ 
29+int64_t GetShapeSize(const std::vector<int64_t>& shape)
30+{
31+ int64_t shapeSize = 1;
32+ for (auto i : shape) {
33+ shapeSize *= i;
34+ }
35+ return shapeSize;
36+}
37+ 
38+void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr)
39+{
40+ auto size = GetShapeSize(shape);
41+ std::vector<DataType> resultData(size, 0);
42+ auto ret = aclrtMemcpy(
43+ resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]),
44+ ACL_MEMCPY_DEVICE_TO_HOST);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return);
46+ for (int64_t i = 0; i < size; i++) {
47+ LOG_PRINT("softplus result[%ld] is: %d\n", i, resultData[i]); // int
48+ }
49+}
50+ 
51+int Init(int32_t deviceId, aclrtStream* stream)
52+{
53+ // 固定写法，初始化
54+ auto ret = aclInit(nullptr);
55+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
56+ ret = aclrtSetDevice(deviceId);
57+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
58+ ret = aclrtCreateStream(stream);
59+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
60+ return 0;
61+}
62+ 
63+template <typename T>
64+int CreateAclTensor(
65+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
66+ aclTensor** tensor)
67+{
68+ auto size = GetShapeSize(shape) * sizeof(T);
69+ // 2. 申请device侧内存
70+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
71+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
72+ // 3. 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
73+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
74+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
75+ 
76+ // 计算连续tensor的strides
77+ std::vector<int64_t> strides(shape.size(), 1);
78+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
79+ strides[i] = shape[i + 1] * strides[i + 1];
80+ }
81+ 
82+ // 调用aclCreateTensor接口创建aclTensor
83+ *tensor = aclCreateTensor(
84+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
85+ *deviceAddr);
86+ return 0;
87+}
88+ 
89+int main()
90+{
91+ // 1. 调用acl进行device/stream初始化
92+ int32_t deviceId = 0;
93+ aclrtStream stream;
94+ auto ret = Init(deviceId, &stream);
95+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
96+ 
97+ // 2. 构造输入与输出，需要根据API的接口自定义构造
98+ aclTensor* x = nullptr;
99+ void* xDeviceAddr = nullptr;
100+ std::vector<int64_t> xShape = {8};
101+ std::vector<DataType> xHostData(8);
102+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT, &x);
103+ CHECK_RET(ret == ACL_SUCCESS, return ret);
104+ 
105+ aclTensor* out = nullptr;
106+ void* outDeviceAddr = nullptr;
107+ std::vector<int64_t> outShape = {8};
108+ std::vector<DataType> outHostData(8);
109+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
110+ CHECK_RET(ret == ACL_SUCCESS, return ret);
111+ 
112+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
113+ uint64_t workspaceSize = 0;
114+ aclOpExecutor* executor;
115+ 
116+ LOG_PRINT("Before GetWorkspaceSize: x=%p, out=%p\n", (void*)x, (void*)out);
117+ LOG_PRINT("Before GetWorkspaceSize: xDeviceAddr=%p, outDeviceAddr=%p\n",
118+ xDeviceAddr,outDeviceAddr);
119+ // 4. 调用aclnnAddExample第一段接口
120+ ret = aclnnSoftplusGetWorkspaceSize(x, out, &workspaceSize, &executor);
121+ LOG_PRINT("aclnnSoftplusGetWorkspaceSize returned %d, workspaceSize=%llu, executor=%p\n",
122+ ret, (unsigned long long)workspaceSize, (void*)executor);
123+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSoftplusExampleGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
124+ 
125+ // 根据第一段接口计算出的workspaceSize申请device内存
126+ void* workspaceAddr = nullptr;
127+ if (workspaceSize > static_cast<uint64_t>(0)) {
128+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
129+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
130+ }
131+ 
132+ // 5. 调用aclnnAddExample第二段接口
133+ ret = aclnnSoftplus(workspaceAddr, workspaceSize, executor, stream);
134+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAddExample failed. ERROR: %d\n", ret); return ret);
135+ 
136+ // 6. （固定写法）同步等待任务执行结束
137+ ret = aclrtSynchronizeStream(stream);
138+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
139+ 
140+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
141+ std::vector<int64_t> outShape1 = {8};
142+ PrintOutResult(outShape1, &outDeviceAddr);
143+ 
144+ // 7. 释放aclTensor，需要根据具体API的接口定义修改
145+ aclDestroyTensor(x);
146+ aclDestroyTensor(out);
147+ 
148+ // 8. 释放device资源
149+ aclrtFree(xDeviceAddr);
150+ aclrtFree(outDeviceAddr);
151+ if (workspaceSize > static_cast<uint64_t>(0)) {
152+ aclrtFree(workspaceAddr);
153+ }
154+ aclrtDestroyStream(stream);
155+ aclrtResetDevice(deviceId);
156+ 
157+ // 9. acl去初始化
158+ aclFinalize();
159+ 
160+ return 0;
161+}
Aexperimental/activation/softplus/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE softplus ACLNNTYPE aclnn)
Aexperimental/activation/softplus/op_host/softplus_def.cpp+44-0
@@ -0,0 +1,44 @@
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 softplus.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class Softplus : public OpDef {
19+public:
20+ explicit Softplus(const char* name) : OpDef(name)
21+ {
22+ this->Input("x")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
27+ this->Output("y")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
32+ OpAICoreConfig aicoreConfig;
33+ aicoreConfig.DynamicCompileStaticFlag(true)
34+ .DynamicFormatFlag(false)
35+ .DynamicRankSupportFlag(true)
36+ .DynamicShapeSupportFlag(true)
37+ .NeedCheckSupportFlag(false)
38+ .PrecisionReduceFlag(true)
39+ .ExtendCfgInfo("opFile.value", "softplus"); // 这里制定的值会对应到kernel入口文件名.cpp
40+ this->AICore().AddConfig("ascend910b", aicoreConfig); // 其他的soc版本补充部分配置项
41+ }
42+};
43+OP_ADD(Softplus); // 添加算子信息库
44+} // namespace ops
Aexperimental/activation/softplus/op_host/softplus_infershape.cpp+43-0
@@ -0,0 +1,43 @@
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 softplus_infershape.cpp
13+ * \brief
14+*/
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+ 
20+namespace ops {
21+static constexpr int64_t IDX_0 = 0;
22+ 
23+static ge::graphStatus InferShapeSoftplus(gert::InferShapeContext* context)
24+{
25+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
26+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeSoftplus");
27+ 
28+ // get input shapes
29+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
30+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
31+ 
32+ // get output shapes
33+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
34+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
35+ 
36+ // 填充输出shape大小
37+ *yShape = *xShape;
38+ OP_LOGD(context->GetNodeName(), "End to do InferShapeSoftplus");
39+ return GRAPH_SUCCESS;
40+}
41+ 
42+IMPL_OP_INFERSHAPE(Softplus).InferShape(InferShapeSoftplus);
43+}
Aexperimental/activation/softplus/op_host/softplus_tiling.cpp+204-0
@@ -0,0 +1,204 @@
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 softplus_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "util/platform_util.h"
19+#include "op_host/tiling_util.h"
20+#include "op_host/tiling_templates_registry.h"
21+#include "tiling/platform/platform_ascendc.h"
22+#include "register/op_impl_registry.h"
23+#include "../op_kernel/softplus_tiling_data.h"
24+#include "../op_kernel/softplus_tiling_key.h"
25+ 
26+namespace optiling {
27+ 
28+using namespace Ops::NN::OpTiling;
29+ 
30+// SINGLE BUFFER
31+#define UB_NUM_BF16_F16_ONE 8U
32+#define UB_NUM_F32_ONE 3U
33+// DOUBLE BUFFER
34+#define UB_NUM_BF16_F16_TWO 10U
35+#define UB_NUM_F32_TWO 5U
36+constexpr uint32_t BUFFER_NUM = 2;
37+constexpr uint32_t WS_SYS_SIZE = 0;
38+ 
39+struct SoftplusCompileInfo {};
40+ 
41+static ge::graphStatus TilingParseForSoftplus([[maybe_unused]] gert::TilingParseContext* context)
42+{
43+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
44+ return ge::GRAPH_SUCCESS;
45+}
46+ 
47+// 获取平台信息如ubSize, coreNum
48+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum, uint64_t& BLOCK_SIZE)
49+{
50+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
51+ // 获取ubsize coreNum
52+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
53+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
54+ coreNum = ascendcPlatform.GetCoreNumAiv();
55+ BLOCK_SIZE = Ops::Base::GetUbBlockSize(context);
56+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
57+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
58+ OP_CHECK_IF(BLOCK_SIZE == 0, OP_LOGE(context, "BLOCK_SIZE is 0"), return ge::GRAPH_FAILED);
59+ return ge::GRAPH_SUCCESS;
60+}
61+ 
62+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
63+{
64+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
65+ size_t usrSize = 0;
66+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
67+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
68+ size_t* currentWorkspace = context->GetWorkspaceSizes(
69+ 1); // 通过框架获取workspace的指针，GetWorkspaceSizes入参为所需workspace的块数。当前限制使用一块。
70+ currentWorkspace[0] = usrSize + sysWorkspaceSize;
71+ return ge::GRAPH_SUCCESS;
72+}
73+ 
74+static ge::graphStatus GetShapeAttrsInfo(
75+ gert::TilingContext* context, uint64_t ubSize, uint64_t coreNum, uint64_t BLOCK_SIZE, uint64_t& bufferOpen, uint64_t& inputNum, uint64_t& inputBytes,
76+ uint64_t& tileBlockNum, uint64_t& tileDataNum, uint64_t& inputLengthAlgin)
77+{
78+ OP_CHECK_IF(
79+ context == nullptr || context->GetInputShape(0) == nullptr, OP_LOGE(context, "context is nullptr"),
80+ return ge::GRAPH_FAILED);
81+ inputNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
82+ uint32_t typeLength = 0;
83+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
84+ uint64_t inputLength = inputNum * typeLength;
85+ if (inputNum == 0 || BLOCK_SIZE == 0 || UB_NUM_F32_ONE == 0 || UB_NUM_BF16_F16_ONE == 0) {
86+ OP_LOGE(context, "inputNum or BLOCK_SIZE or UB_NUM_F32_ONE or UB_NUM_BF16_F16_ONE is 0");
87+ return ge::GRAPH_FAILED;
88+ }
89+ inputBytes = inputLength / inputNum;
90+ inputLengthAlgin = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE);
91+ uint64_t ubDataNumber;
92+ // double buffer default open
93+ bufferOpen = 1;
94+ if (context->GetInputDesc(0)->GetDataType() == ge::DT_FLOAT && (inputLengthAlgin < coreNum * (((ubSize / BLOCK_SIZE) * BLOCK_SIZE) / UB_NUM_F32_ONE))) {
95+ bufferOpen = 0;
96+ ubDataNumber = UB_NUM_F32_ONE;
97+ } else {
98+ ubDataNumber = UB_NUM_F32_TWO;
99+ }
100+ if (context->GetInputDesc(0)->GetDataType() != ge::DT_FLOAT && (inputLengthAlgin < coreNum * (((ubSize / BLOCK_SIZE) * BLOCK_SIZE) / UB_NUM_BF16_F16_ONE))) {
101+ bufferOpen = 0;
102+ ubDataNumber = UB_NUM_BF16_F16_ONE;
103+ } else {
104+ ubDataNumber = UB_NUM_BF16_F16_TWO;
105+ }
106+ if (ubDataNumber == 0) {
107+ OP_LOGE(context, "ubDataNumber is 0");
108+ return ge::GRAPH_FAILED;
109+ }
110+ tileBlockNum = (ubSize / BLOCK_SIZE) / ubDataNumber;
111+ if (inputBytes == 0) {
112+ OP_LOGE(context, "inputBytes is 0");
113+ return ge::GRAPH_FAILED;
114+ }
115+ tileDataNum = (tileBlockNum * BLOCK_SIZE) / inputBytes;
116+ return ge::GRAPH_SUCCESS;
117+}
118+ 
119+static ge::graphStatus CalculateCoreBlockNums(
120+ gert::TilingContext* context, uint64_t inputLengthAlgin, int64_t coreNum, uint64_t BLOCK_SIZE, uint64_t tileBlockNum, uint64_t inputBytes, uint64_t tileDataNum,
121+ uint64_t& smallCoreDataNum, uint64_t& bigCoreDataNum, uint64_t& smallTailDataNum, uint64_t& bigTailDataNum, uint64_t& finalSmallTileNum, uint64_t& finalBigTileNum,
122+ uint64_t& tailBlockNum)
123+{
124+ if (0 == BLOCK_SIZE || 0 == coreNum || 0 == tileBlockNum || 0 == inputBytes) {
L
Lliujie123456784月2日

BLOCK_SIZE上面设置为常量256，是不是不需要进行判断是否等于0操作

likedislike
松柏
松柏
4月2日 评论：
125+ OP_LOGE(context, "BLOCK_SIZE or coreNum or tileBlockNum or inputBytes is 0");
126+ return ge::GRAPH_FAILED;
127+ }
128+ uint64_t everyCoreInputBlockNum = inputLengthAlgin / BLOCK_SIZE / coreNum;
129+ tailBlockNum = (inputLengthAlgin / BLOCK_SIZE) % coreNum;
130+ smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
131+ uint64_t smallTileNum = everyCoreInputBlockNum / tileBlockNum;
132+ finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
133+ smallTailDataNum = smallCoreDataNum - (tileDataNum * smallTileNum);
134+ smallTailDataNum = smallTailDataNum == 0 ? tileDataNum : smallTailDataNum;
135+ 
136+ everyCoreInputBlockNum += 1;
137+ bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
138+ uint64_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
139+ finalBigTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
140+ bigTailDataNum = bigCoreDataNum - tileDataNum * bigTileNum;
141+ bigTailDataNum = bigTailDataNum == 0 ? tileDataNum : bigTailDataNum;
142+ 
143+ return ge::GRAPH_SUCCESS;
144+}
145+ 
146+// tiling 分发入口
147+static ge::graphStatus SoftplusTilingFunc(gert::TilingContext* context)
148+{
149+ // 1、获取平台运行信息
150+ uint64_t ubSize;
151+ int64_t coreNum;
152+ // double buffer Open or Off
153+ uint64_t bufferOpen;
154+ uint64_t BLOCK_SIZE;
155+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, coreNum, BLOCK_SIZE);
156+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
157+ // 2、获取shape、属性信息
158+ uint64_t inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin;
159+ ret = GetShapeAttrsInfo(context, ubSize, coreNum, BLOCK_SIZE, bufferOpen, inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin);
160+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
161+ // 3、获取WorkspaceSize信息
162+ OP_CHECK_IF(
163+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
164+ return ge::GRAPH_FAILED);
165+ // 4、设置tiling信息
166+ SoftplusTilingData* tiling = context->GetTilingData<SoftplusTilingData>();
167+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
168+ OP_CHECK_IF(
169+ memset_s(tiling, sizeof(SoftplusTilingData), 0, sizeof(SoftplusTilingData)) != EOK,
170+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
171+ 
172+ if (tileDataNum >= inputNum) {
173+ coreNum = 1;
174+ } else {
175+ coreNum = (static_cast<uint64_t>(coreNum) < inputLengthAlgin / BLOCK_SIZE) ? coreNum : inputLengthAlgin / BLOCK_SIZE;
176+ }
177+ // 计算每个core处理的数据块数
178+ uint64_t smallCoreDataNum, bigCoreDataNum, smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum;
179+ ret = CalculateCoreBlockNums(
180+ context, inputLengthAlgin, coreNum, BLOCK_SIZE, tileBlockNum, inputBytes, tileDataNum, smallCoreDataNum, bigCoreDataNum,
181+ smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum);
182+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "CalculateCoreBlockNums error"), return ge::GRAPH_FAILED);
183+ // 设置tiling数据
184+ tiling->smallCoreDataNum = static_cast<uint64_t>(smallCoreDataNum);
185+ tiling->bigCoreDataNum = static_cast<uint64_t>(bigCoreDataNum);
186+ tiling->tileDataNum = static_cast<uint64_t>(tileDataNum);
187+ tiling->smallTailDataNum = static_cast<uint64_t>(smallTailDataNum);
188+ tiling->bigTailDataNum = static_cast<uint64_t>(bigTailDataNum);
189+ tiling->finalSmallTileNum = static_cast<uint64_t>(finalSmallTileNum);
190+ tiling->finalBigTileNum = static_cast<uint64_t>(finalBigTileNum);
191+ tiling->tailBlockNum = static_cast<uint64_t>(tailBlockNum);
192+ 
193+ tiling->bufferOpen = static_cast<uint64_t>(bufferOpen);
194+ 
195+ context->SetBlockDim(coreNum);
196+ uint64_t tilingKey = 0;
197+ tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0);
198+ context->SetTilingKey(tilingKey);
199+ return ge::GRAPH_SUCCESS;
200+}
201+ 
202+// tiling注册入口.
203+IMPL_OP_OPTILING(Softplus).Tiling(SoftplusTilingFunc).TilingParse<SoftplusCompileInfo>(TilingParseForSoftplus);
204+} // namespace optiling
Aexperimental/activation/softplus/op_kernel/softplus.cpp+28-0
@@ -0,0 +1,28 @@
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 softplus.cpp
13+ * \brief
14+ */
15+ 
16+#include "softplus.h"
17+ 
18+ 
19+template <uint32_t schMode>
20+__global__ __aicore__ void softplus(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
21+{
22+ REGISTER_TILING_DEFAULT(SoftplusTilingData);
23+ GET_TILING_DATA_WITH_STRUCT(SoftplusTilingData, tilingData, tiling);
24+ NsSoftplus::KernelSoftplus<DTYPE_X> op; // 算子kernel实例获取
25+ op.Init(x, y, tilingData.smallCoreDataNum, tilingData.bigCoreDataNum, tilingData.finalBigTileNum, tilingData.finalSmallTileNum, tilingData.tileDataNum,
26+ tilingData.smallTailDataNum, tilingData.bigTailDataNum, tilingData.tailBlockNum, tilingData.bufferOpen);
27+ op.Process();
28+}
Aexperimental/activation/softplus/op_kernel/softplus.h+179-0
@@ -0,0 +1,179 @@
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 softplus.h
13+ * \brief
14+ */
15+#ifndef SOFTPLUS_H
16+#define SOFTPLUS_H
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "softplus_tiling_data.h"
21+#include "softplus_tiling_key.h"
22+ 
23+#include "kernel_operator.h"
24+ 
25+namespace NsSoftplus {
26+ 
27+using namespace AscendC;
28+constexpr float MAGIC_NUM = -0.69314718055994530941723212145818f;
29+constexpr int32_t SINGLE_BUFFER_NUM = 1;
30+constexpr int32_t DOUBLE_BUFFER_NUM = 2;
31+ 
32+template <typename TYPE_X>
33+class KernelSoftplus {
34+public:
35+ __aicore__ inline KernelSoftplus(){};
36+ 
37+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
38+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum, uint64_t tailBlockNum, uint64_t bufferOpen);
39+ __aicore__ inline void Process();
40+ 
41+private:
42+ __aicore__ inline void CopyIn(int32_t progress);
43+ __aicore__ inline void CopyOut(int32_t progress);
44+ __aicore__ inline void Compute(int32_t progress);
45+ 
46+private:
47+ AscendC::TPipe pipe;
48+ AscendC::TQue<AscendC::TPosition::VECIN, DOUBLE_BUFFER_NUM> inQueueX;
49+ AscendC::TQue<AscendC::TPosition::VECOUT, DOUBLE_BUFFER_NUM> outQueueY;
50+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpQueue0, tmpQueue1, tmpQueue2;
51+ 
52+ AscendC::GlobalTensor<TYPE_X> xGm, yGm;
53+ uint64_t coreDataNum;
54+ uint64_t tileNum;
55+ uint64_t tileDataNum;
56+ uint64_t tailDataNum;
57+ uint64_t processDataNum;
58+};
59+ 
60+template <typename TYPE_X>
61+__aicore__ inline void KernelSoftplus<TYPE_X>::Init(GM_ADDR x, GM_ADDR y, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
62+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum, uint64_t tailBlockNum, uint64_t bufferOpen)
63+{
64+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
65+ uint64_t coreId = AscendC::GetBlockIdx();
66+ uint64_t globalBufferIndex = bigCoreDataNum * coreId;
67+ this->tileDataNum = tileDataNum;
68+ // default open DOUBLE BUFFER
69+ uint64_t BUFFER_NUM = DOUBLE_BUFFER_NUM;
70+ if (bufferOpen == 0) {
71+ BUFFER_NUM = SINGLE_BUFFER_NUM;
72+ }
73+ if (coreId < tailBlockNum) {
74+ this->coreDataNum = bigCoreDataNum;
75+ this->tileNum = finalBigTileNum;
76+ this->tailDataNum = bigTailDataNum;
77+ } else {
78+ this->coreDataNum = smallCoreDataNum;
79+ this->tileNum = finalSmallTileNum;
80+ this->tailDataNum = smallTailDataNum;
81+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (coreId - tailBlockNum);
82+ }
83+ xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
84+ yGm.SetGlobalBuffer((__gm__ TYPE_X *)y + globalBufferIndex, this->coreDataNum);
85+ 
86+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
87+ pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
88+ if (std::is_same_v<TYPE_X, float>) {
89+ pipe.InitBuffer(tmpQueue0, this->tileDataNum * sizeof(float));
90+ } else {
91+ pipe.InitBuffer(tmpQueue0, this->tileDataNum * sizeof(float));
92+ pipe.InitBuffer(tmpQueue1, this->tileDataNum * sizeof(float));
93+ pipe.InitBuffer(tmpQueue2, this->tileDataNum * sizeof(float));
94+ }
95+}
96+ 
97+template <typename TYPE_X>
98+__aicore__ inline void KernelSoftplus<TYPE_X>::CopyIn(int32_t progress)
99+{
100+ AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.AllocTensor<TYPE_X>();
101+ AscendC::DataCopy(xLocal, xGm[progress * this->tileDataNum], this->processDataNum);
102+ inQueueX.EnQue(xLocal);
103+}
104+ 
105+template <typename TYPE_X>
106+__aicore__ inline void KernelSoftplus<TYPE_X>::CopyOut(int32_t progress)
107+{
108+ AscendC::LocalTensor<TYPE_X> yLocal = outQueueY.DeQue<TYPE_X>();
109+ AscendC::DataCopy(yGm[progress * this->tileDataNum], yLocal, this->processDataNum);
110+ outQueueY.FreeTensor(yLocal);
111+}
112+ 
113+template <typename TYPE_X>
114+__aicore__ inline void KernelSoftplus<TYPE_X>::Compute(int32_t progress)
115+{
116+ if constexpr (std::is_same_v<TYPE_X, float>) {
117+ //float32
118+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
119+ AscendC::LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
120+ AscendC::LocalTensor<float> tmp0Local = tmpQueue0.AllocTensor<float>();
121+ AscendC::Maxs(tmp0Local, xLocal, static_cast<float>(0.0), this->processDataNum);
122+ AscendC::Muls(yLocal, tmp0Local, static_cast<float>(-1.0), this->processDataNum);
123+ AscendC::Exp(yLocal, yLocal, this->processDataNum);
124+ AscendC::Adds(yLocal, yLocal, static_cast<float>(1.0), this->processDataNum);
125+ AscendC::Ln(yLocal, yLocal, this->processDataNum);
126+ AscendC::Add(yLocal, yLocal, tmp0Local, this->processDataNum);
127+ AscendC::Mins(tmp0Local, xLocal, static_cast<float>(0.0), this->processDataNum);
128+ AscendC::Exp(tmp0Local, tmp0Local, this->processDataNum);
129+ AscendC::Adds(tmp0Local, tmp0Local, static_cast<float>(1.0), this->processDataNum);
130+ AscendC::Ln(tmp0Local, tmp0Local, this->processDataNum);
131+ AscendC::Add(yLocal, yLocal, tmp0Local, this->processDataNum);
132+ AscendC::Adds(yLocal, yLocal, static_cast<float>(MAGIC_NUM), this->processDataNum);
133+ outQueueY.EnQue<float>(yLocal);
134+ inQueueX.FreeTensor(xLocal);
135+ } else {
136+ // float16 or bfloat16
137+ // xLocal---tmp1Local tmp2Local---yLocal
138+ AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.DeQue<TYPE_X>();
139+ AscendC::LocalTensor<TYPE_X> yLocal = outQueueY.AllocTensor<TYPE_X>();
140+ AscendC::LocalTensor<float> tmp0Local = tmpQueue0.AllocTensor<float>();
141+ AscendC::LocalTensor<float> tmp1Local = tmpQueue1.AllocTensor<float>();
142+ AscendC::LocalTensor<float> tmp2Local = tmpQueue2.AllocTensor<float>();
143+ AscendC::Cast(tmp1Local, xLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
144+ AscendC::Maxs(tmp0Local, tmp1Local, static_cast<float>(0.0), this->processDataNum);
145+ AscendC::Muls(tmp2Local, tmp0Local, static_cast<float>(-1.0), this->processDataNum);
L
Lliujie123456784月2日

这里面应该缺少pipe_v同步指令，不会精度问题吗

likedislike
松柏
松柏
4月2日 评论：
F
fulltower
4月2日 评论：
146+ AscendC::Exp(tmp2Local, tmp2Local, this->processDataNum);
147+ AscendC::Adds(tmp2Local, tmp2Local, static_cast<float>(1.0), this->processDataNum);
148+ AscendC::Ln(tmp2Local, tmp2Local, this->processDataNum);
149+ AscendC::Add(tmp2Local, tmp2Local, tmp0Local, this->processDataNum);
150+ AscendC::Mins(tmp0Local, tmp1Local, static_cast<float>(0.0), this->processDataNum);
151+ AscendC::Exp(tmp0Local, tmp0Local, this->processDataNum);
152+ AscendC::Adds(tmp0Local, tmp0Local, static_cast<float>(1.0), this->processDataNum);
153+ AscendC::Ln(tmp0Local, tmp0Local, this->processDataNum);
154+ AscendC::Add(tmp2Local, tmp2Local, tmp0Local, this->processDataNum);
155+ AscendC::Adds(tmp2Local, tmp2Local, static_cast<float>(MAGIC_NUM), this->processDataNum);
156+ AscendC::Cast(yLocal, tmp2Local, AscendC::RoundMode::CAST_RINT, this->processDataNum);
157+ outQueueY.EnQue<TYPE_X>(yLocal);
158+ inQueueX.FreeTensor(xLocal);
159+ }
160+}
161+ 
162+template <typename TYPE_X>
163+__aicore__ inline void KernelSoftplus<TYPE_X>::Process()
164+{
165+ int32_t loopCount = this->tileNum;
166+ this->processDataNum = this->tileDataNum;
167+ for (int32_t i = 0; i < loopCount - 1; i++) {
168+ CopyIn(i);
169+ Compute(i);
170+ CopyOut(i);
171+ }
172+ this->processDataNum = this->tailDataNum;
173+ CopyIn(loopCount - 1);
174+ Compute(loopCount - 1);
175+ CopyOut(loopCount - 1);
176+}
177+ 
178+} // namespace NsSoftplus
179+#endif // SOFTPLUS_H
Aexperimental/activation/softplus/op_kernel/softplus_tiling_data.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+/*!
12+ * \file softplus_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef SOFTPLUS_TILING_DATA_H_
17+#define SOFTPLUS_TILING_DATA_H_
18+ 
19+struct SoftplusTilingData {
20+ uint64_t smallCoreDataNum;
21+ uint64_t bigCoreDataNum;
22+ uint64_t finalBigTileNum;
23+ uint64_t finalSmallTileNum;
24+ uint64_t tileDataNum;
25+ uint64_t smallTailDataNum;
26+ uint64_t bigTailDataNum;
27+ uint64_t tailBlockNum;
28+ uint64_t bufferOpen;
29+};
30+#endif
Aexperimental/activation/softplus/op_kernel/softplus_tiling_key.h+31-0
@@ -0,0 +1,31 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file softplus_tiling_key.h
13+ * \brief softplus tiling key declare
14+ */
15+ 
16+#ifndef __SOFTPLUS_TILING_KEY_H__
17+#define __SOFTPLUS_TILING_KEY_H__
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+#define ELEMENTWISE_TPL_SCH_MODE_0 0
22+#define ELEMENTWISE_TPL_SCH_MODE_1 1
23+ 
24+ASCENDC_TPL_ARGS_DECL(
25+ Softplus,
26+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
27+ 
28+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
29+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));
30+ 
31+#endif
Aexperimental/activation/softplus/tests/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
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+message(STATUS "=== Debug: start ops.activation.softplus.tests.CMakeLists.txt ")
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
14+foreach(SUB_DIR ${CURRENT_DIRS})
15+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16+ add_subdirectory(${SUB_DIR})
17+ endif()
18+endforeach()
Aexperimental/activation/softplus/tests/ut/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
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+# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
12+message(STATUS "=== Debug: start ops.activation.threshold_grad_v2_d.tests.ut.CMakeLists.txt ")
13+file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
14+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
15+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
16+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aexperimental/activation/softplus/tests/ut/op_host/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
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+file(GLOB CURRENT_DIRS 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()
Aexperimental/activation/softplus/tests/ut/op_host/test_threshold_grad_v2_d_tiling.cpp+103-0
@@ -0,0 +1,103 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <vector>
13+#include <gtest/gtest.h>
14+#include "log/log.h"
15+#include "kernel_run_context_facker.h"
16+#include "exe_graph/runtime/storage_format.h"
17+#include "exe_graph/runtime/storage_shape.h"
18+#include "test_cube_util.h"
19+#include "register/op_impl_registry.h"
20+#include "ut_op_util.h"
21+#include "ut_op_common.h"
22+#include "platform/platform_infos_def.h"
23+ 
24+using namespace ut_util;
25+using namespace std;
26+using namespace ge;
27+ 
28+class SoftplusTiling : public testing::Test {
29+protected:
30+ static void SetUpTestCase()
31+ {
32+ std::cout << "SoftplusTiling SetUp" << std::endl;
33+ }
34+ 
35+ static void TearDownTestCase()
36+ {
37+ std::cout << "SoftplusTiling TearDown" << std::endl;
38+ }
39+};
40+ 
41+TEST_F(SoftplusTiling, threshold_grad_v2_float32_success) {
42+ // input
43+ gert::StorageShape x1_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
44+ // output
45+ gert::StorageShape y_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
46+ string compile_info_string = R"({
47+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
48+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
49+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
50+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
51+ "CORE_NUM": 48}
52+ })";
53+ map<string, string> soc_infos;
54+ map<string, string> aicore_spec;
55+ map<string, string> intrinsics;
56+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
57+ 
58+ // platform info
59+ fe::PlatFormInfos platform_info;
60+ platform_info.Init();
61+ 
62+ // compile info
63+ struct SoftplusTilingCompileInfo {};
64+ SoftplusTilingCompileInfo compile_info;
65+ 
66+ std::string op_type("Softplus");
67+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
68+ 
69+ // tilingParseFunc simulate
70+ auto kernel_holder =
71+ gert::KernelRunContextFaker()
72+ .KernelIONum(1, 1)
73+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
74+ .Outputs({&compile_info})
75+ .Build();
76+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
77+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
78+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
79+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
80+ intrinsics);
81+ 
82+ // tilingFunc simulate
83+ auto param = gert::TilingData::CreateCap(4096);
84+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
85+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
86+ ASSERT_NE(param, nullptr);
87+ auto holder = gert::TilingContextFaker()
88+ .SetOpType("Softplus")
89+ .NodeIoNum(1, 1)
90+ .IrInstanceNum({0, 1})
91+ .InputShapes({&x1_shape})
92+ .OutputShapes({&y_shape})
93+ .CompileInfo(&compile_info)
94+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
95+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
96+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
97+ .TilingData(param.get())
98+ .Workspace(ws_size)
99+ .Build();
100+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
101+ ASSERT_NE(tiling_context, nullptr);
102+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
103+}
Aexperimental/activation/softplus/tests/ut/op_kernel/CMakeLists.txt+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+ 
12+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
13+ # 需要将Tiling依赖的文件添加到CMakeLists.txt中
14+ # set(elewise_common_tiling_files
15+ # ${CANN_ROOT}/ops/built-in/op_tiling/runtime/elewise_tiling.cc
16+ # )
17+ # 算子自己的tiling文件路径
18+ set(softplus_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/softplus_tiling.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/softplus_infershape.cpp
21+ # ${elewise_common_tiling_files}
22+ )
23+ # 使用AddOpTestCase
24+ # param1：算子名称，以kernel方式命名
25+ # param2：soc版本，多个以分号分隔，例如："ascend950pr_9599;AscendB1"
26+ # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔，例如："-DDTYPE_X=float -DDTYPE_Y=float"
27+ # param4：该算子依赖的所有tiling源码文件
28+ # AddOpTestCase(softplus "ascend910b" "${softplus_tiling_files}")
29+ AddOpTestCase(softplus "ascend910B1" "-DDTYPE_X=float" "${softplus_tiling_files}")
30+endif()
Aexperimental/activation/softplus/tests/ut/op_kernel/test_softplus.cpp+81-0
@@ -0,0 +1,81 @@
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 <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+ 
15+#ifdef __CCE_KT_TEST__
16+#include "tikicpulib.h"
17+#include "data_utils.h"
18+#include "string.h"
19+#include <iostream>
20+#include <string>
21+#endif
22+#include "../../../op_kernel/softplus.cpp"
23+#include "../../../op_kernel/softplus_tiling_data.h"
24+#include <cstdint>
25+ 
26+using namespace std;
27+ 
28+class softplus_test : public testing::Test {
29+protected:
30+ static void SetUpTestCase()
31+ {
32+ cout << "softplus_test SetUp\n" << endl;
33+ }
34+ static void TearDownTestCase()
35+ {
36+ cout << "softplus_test TearDown\n" << endl;
37+ }
38+};
39+ 
40+TEST_F(softplus_test, test_case_0)
41+{
42+ size_t xByteSize = 32 * 4 * 4 * 4 * sizeof(float);
43+ size_t yByteSize = 32 * 4 * 4 * 4 * sizeof(float);
44+ size_t tiling_data_size = sizeof(SoftplusTilingData);
45+ uint32_t blockDim = 1;
46+ 
47+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(xByteSize);
48+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(yByteSize);
49+ 
50+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(1024 * 1024 * 16);
51+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tiling_data_size);
52+ 
53+ char* path_ = get_current_dir_name();
54+ string path(path_);
55+ 
56+ SoftplusTilingData* tilingDatafromBin = reinterpret_cast<SoftplusTilingData*>(tiling);
57+ 
58+ tilingDatafromBin->smallCoreDataNum = 2048;
59+ tilingDatafromBin->bigCoreDataNum = 2112;
60+ tilingDatafromBin->tileDataNum = 4032;
61+ tilingDatafromBin->smallTailDataNum = 2048;
62+ tilingDatafromBin->bigTailDataNum = 2112;
63+ tilingDatafromBin->finalSmallTileNum = 1;
64+ tilingDatafromBin->finalBigTileNum = 1;
65+ tilingDatafromBin->tailBlockNum = 0;
66+ tilingDatafromBin->bufferOpen = 1;
67+ 
68+ auto SoftplusKernel = [](GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
69+ ::softplus<0>(x, y, workspace, tiling);
70+ };
71+ 
72+ ICPU_SET_TILING_KEY(0);
73+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
74+ ICPU_RUN_KF(SoftplusKernel, blockDim, x, y, workspace, (uint8_t *)(tilingDatafromBin));
75+ 
76+ AscendC::GmFree(x);
77+ AscendC::GmFree(y);
78+ AscendC::GmFree(workspace);
79+ AscendC::GmFree(tiling);
80+ free(path_);
81+}