Star906
1826
代码
代码
Issues226
Pull Requests306
Actions
项目讨论
Wiki
分析
Star906
1826
已合并
Add SoftPlusV2 operator implementation (issue #263) #474
周珈民创建于 2025年12月23日
Add SoftPlusV2 operator implementation (issue #263) #474
已合并
周珈民创建于 2025年12月23日
从zhou-jiamin-666/ops-nn_prfeat/op_soft_plus_v2合入到cann/ops-nnmaster
共 18 个文件变更+1158-3
Mexperimental/activation/celu_v2/examples/test_aclnn_celu_v2.cpp+3-3
@@ -47,7 +47,7 @@ int64_t GetShapeSize(const std::vector<int64_t>& shape)
47void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr)47void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr)
48{48{
49 auto size = GetShapeSize(shape);49 auto size = GetShapeSize(shape);
50- std::vector<DataType> resultData(size, 0);50+ std::vector<float> resultData(size, 0);
51 auto ret = aclrtMemcpy(51 auto ret = aclrtMemcpy(
52 resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]),52 resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]),
53 ACL_MEMCPY_DEVICE_TO_HOST);53 ACL_MEMCPY_DEVICE_TO_HOST);
@@ -109,14 +109,14 @@ int main()
109 aclTensor* selfX = nullptr;109 aclTensor* selfX = nullptr;
110 void* selfXDeviceAddr = nullptr;110 void* selfXDeviceAddr = nullptr;
111 std::vector<int64_t> selfXShape = {128, 8, 8, 8};111 std::vector<int64_t> selfXShape = {128, 8, 8, 8};
112- std::vector<DataType> selfXHostData(2048, -2.0f);112+ std::vector<float> selfXHostData(2048, -2.0f);
113 ret = CreateAclTensor(selfXHostData, selfXShape, &selfXDeviceAddr, aclDataType::ACL_FLOAT, &selfX);113 ret = CreateAclTensor(selfXHostData, selfXShape, &selfXDeviceAddr, aclDataType::ACL_FLOAT, &selfX);
114 CHECK_RET(ret == ACL_SUCCESS, return ret);114 CHECK_RET(ret == ACL_SUCCESS, return ret);
115 115 
116 aclTensor* out = nullptr;116 aclTensor* out = nullptr;
117 void* outDeviceAddr = nullptr;117 void* outDeviceAddr = nullptr;
118 std::vector<int64_t> outShape = {128, 8, 8, 8};118 std::vector<int64_t> outShape = {128, 8, 8, 8};
119- std::vector<DataType> outHostData(2048, 0);119+ std::vector<float> outHostData(2048, 0);
120 ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);120 ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
121 CHECK_RET(ret == ACL_SUCCESS, return ret);121 CHECK_RET(ret == ACL_SUCCESS, return ret);
122 122 
Aexperimental/activation/soft_plus_v2/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 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
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
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/soft_plus_v2/examples/test_aclnn_soft_plus_v2.cpp+165-0
@@ -0,0 +1,165 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+#include <iostream>
22+#include <vector>
23+#include "acl/acl.h"
24+#include "aclnn_soft_plus_v2.h"
25+using DataType = float;
26+#define CHECK_RET(cond, return_expr) \
27+ do { \
28+ if (!(cond)) { \
29+ return_expr; \
30+ } \
31+ } while (0)
32+ 
33+#define LOG_PRINT(message, ...) \
34+ do { \
35+ printf(message, ##__VA_ARGS__); \
36+ } while (0)
37+ 
38+int64_t GetShapeSize(const std::vector<int64_t>& shape)
39+{
40+ int64_t shapeSize = 1;
41+ for (auto i : shape) {
42+ shapeSize *= i;
43+ }
44+ return shapeSize;
45+}
46+ 
47+void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr)
48+{
49+ auto size = GetShapeSize(shape);
50+ std::vector<float> resultData(size, 0);
51+ auto ret = aclrtMemcpy(
52+ resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]),
53+ ACL_MEMCPY_DEVICE_TO_HOST);
54+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return);
55+ for (int64_t i = 0; i < size; i++) {
56+ LOG_PRINT("mean result[%ld] is: %f\n", i, resultData[i]);
57+ }
58+}
59+ 
60+int Init(int32_t deviceId, aclrtStream* stream)
61+{
62+ // 固定写法，初始化
63+ auto ret = aclInit(nullptr);
64+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
65+ ret = aclrtSetDevice(deviceId);
66+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
67+ ret = aclrtCreateStream(stream);
68+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
69+ return 0;
70+}
71+ 
72+template <typename T>
73+int CreateAclTensor(
74+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
75+ aclTensor** tensor)
76+{
77+ auto size = GetShapeSize(shape) * sizeof(T);
78+ // 2. 申请device侧内存
79+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
80+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
81+ // 3. 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
82+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
83+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
84+ 
85+ // 计算连续tensor的strides
86+ std::vector<int64_t> strides(shape.size(), 1);
87+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
88+ strides[i] = shape[i + 1] * strides[i + 1];
89+ }
90+ 
91+ // 调用aclCreateTensor接口创建aclTensor
92+ *tensor = aclCreateTensor(
93+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
94+ *deviceAddr);
95+ return 0;
96+}
97+ 
98+int main()
99+{
100+ // 1. 调用acl进行device/stream初始化
101+ int32_t deviceId = 0;
102+ aclrtStream stream;
103+ auto ret = Init(deviceId, &stream);
104+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
105+ 
106+ // 2. 构造输入与输出，需要根据API的接口自定义构造
107+ aclTensor* selfX = nullptr;
108+ void* selfXDeviceAddr = nullptr;
109+ std::vector<int64_t> selfXShape = {32, 4, 4, 4};
110+ std::vector<float> selfXHostData(2048, 1.0f);
111+ ret = CreateAclTensor(selfXHostData, selfXShape, &selfXDeviceAddr, aclDataType::ACL_FLOAT, &selfX);
112+ CHECK_RET(ret == ACL_SUCCESS, return ret);
113+ 
114+ aclTensor* out = nullptr;
115+ void* outDeviceAddr = nullptr;
116+ std::vector<int64_t> outShape = {32, 4, 4, 4};
117+ std::vector<float> outHostData(2048, 0);
118+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
119+ CHECK_RET(ret == ACL_SUCCESS, return ret);
120+ 
121+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
122+ uint64_t workspaceSize = 0;
123+ aclOpExecutor* executor;
124+ 
125+ // 4. 调用aclnnSoftPlusV2Example第一段接口
126+ ret = aclnnSoftPlusV2GetWorkspaceSize(selfX, out, &workspaceSize, &executor);
127+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSoftPlusV2ExampleGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
128+ 
129+ // 根据第一段接口计算出的workspaceSize申请device内存
130+ void* workspaceAddr = nullptr;
131+ if (workspaceSize > static_cast<uint64_t>(0)) {
132+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
133+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
134+ }
135+ 
136+ // 5. 调用aclnnSoftPlusV2Example第二段接口
137+ ret = aclnnSoftPlusV2(workspaceAddr, workspaceSize, executor, stream);
138+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSoftPlusV2Example failed. ERROR: %d\n", ret); return ret);
139+ 
140+ // 6. （固定写法）同步等待任务执行结束
141+ ret = aclrtSynchronizeStream(stream);
142+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
143+ 
144+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
145+ std::vector<int64_t> outShape1 = {8};
146+ PrintOutResult(outShape1, &outDeviceAddr);
147+ 
148+ // 7. 释放aclTensor，需要根据具体API的接口定义修改
149+ aclDestroyTensor(selfX);
150+ aclDestroyTensor(out);
151+ 
152+ // 8. 释放device资源
153+ aclrtFree(selfXDeviceAddr);
154+ aclrtFree(outDeviceAddr);
155+ if (workspaceSize > static_cast<uint64_t>(0)) {
156+ aclrtFree(workspaceAddr);
157+ }
158+ aclrtDestroyStream(stream);
159+ aclrtResetDevice(deviceId);
160+ 
161+ // 9. acl去初始化
162+ aclFinalize();
163+ 
164+ return 0;
165+}
Aexperimental/activation/soft_plus_v2/examples/test_geir_soft_plus_v2.cpp+304-0
@@ -0,0 +1,304 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+#include <iostream>
22+#include <fstream>
23+#include <string.h>
24+#include <stdint.h>
25+#include <vector>
26+#include <string>
27+#include <map>
28+#include "assert.h"
29+ 
30+#include "graph.h"
31+#include "types.h"
32+#include "tensor.h"
33+#include "ge_error_codes.h"
34+#include "ge_api_types.h"
35+#include "ge_api.h"
36+#include "array_ops.h"
37+#include "ge_ir_build.h"
38+ 
39+#include "experiment_ops.h"
40+#include "nn_other.h"
41+#include "../op_graph/soft_plus_v2_proto.h"
42+ 
43+#define FAILED -1
44+#define SUCCESS 0
45+ 
46+using namespace ge;
47+using std::map;
48+using std::string;
49+using std::vector;
50+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
51+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
52+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
53+ TensorDesc placeholder##intputIndex##_desc = \
54+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
55+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
56+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
57+ Tensor tensor_placeholder##intputIndex; \
58+ ret = GenOnesData( \
59+ placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, placeholder##intputIndex##_desc, \
60+ intputDtype, 2); \
61+ if (ret != SUCCESS) { \
62+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
63+ return FAILED; \
64+ } \
65+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
66+ input.push_back(tensor_placeholder##intputIndex); \
67+ graph.AddOp(placeholder##intputIndex); \
68+ add1.set_input_##intputName(placeholder##intputIndex); \
69+ inputs.push_back(placeholder##intputIndex);
70+ 
71+#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape) \
72+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
73+ auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \
74+ TensorDesc placeholder##intputIndex##_desc = \
75+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
76+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
77+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
78+ Tensor tensor_placeholder##intputIndex; \
79+ ret = GenOnesData( \
80+ placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, placeholder##intputIndex##_desc, \
81+ intputDtype, 2); \
82+ if (ret != SUCCESS) { \
83+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
84+ return FAILED; \
85+ } \
86+ placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \
87+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
88+ graph.AddOp(placeholder##intputIndex); \
89+ add1.set_input_##intputName(placeholder##intputIndex); \
90+ add1.update_input_desc_##intputName(placeholder##intputIndex##_desc); \
91+ inputs.push_back(placeholder##intputIndex);
92+ 
93+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
94+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
95+ add1.update_output_desc_##outputName(outputName##outputIndex##_desc);
96+ 
97+string GetTime()
98+{
99+ time_t timep;
100+ time(&timep);
101+ char tmp[64];
102+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
103+ return tmp;
104+}
105+ 
106+uint32_t GetDataTypeSize(DataType dt)
107+{
108+ uint32_t dilation = 1;
109+ uint32_t oneByte = 1;
110+ uint32_t twoByte = 2;
111+ uint32_t fourByte = 4;
112+ uint32_t eightByte = 8;
113+ 
114+ if (dt == ge::DT_FLOAT) {
115+ dilation = fourByte;
116+ } else if (dt == ge::DT_FLOAT16) {
117+ dilation = twoByte;
118+ } else if (dt == ge::DT_BF16) {
119+ dilation = twoByte;
120+ } else if (dt == ge::DT_INT16) {
121+ dilation = twoByte;
122+ } else if (dt == ge::DT_UINT16) {
123+ dilation = twoByte;
124+ } else if (dt == ge::DT_INT32) {
125+ dilation = fourByte;
126+ } else if (dt == ge::DT_UINT32) {
127+ dilation = fourByte;
128+ } else if (dt == ge::DT_INT64) {
129+ dilation = eightByte;
130+ } else if (dt == ge::DT_UINT64) {
131+ dilation = eightByte;
132+ } else if (dt == ge::DT_INT8) {
133+ dilation = oneByte;
134+ }
135+ return dilation;
136+}
137+ 
138+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value)
139+{
140+ input_tensor_desc.SetRealDimCnt(shapes.size());
141+ size_t size = 1;
142+ for (uint32_t i = 0; i < shapes.size(); i++) {
143+ size *= shapes[i];
144+ }
145+ uint32_t byteSizeFloat32 = 4;
146+ uint32_t data_len = size * byteSizeFloat32;
147+ float* pData = new (std::nothrow) float[size];
148+ 
149+ for (size_t i = 0; i < size; ++i) {
150+ *(pData + i) = value;
151+ }
152+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
153+ return SUCCESS;
154+}
155+ 
156+int32_t GenOnesData(
157+ vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type, int value)
158+{
159+ input_tensor_desc.SetRealDimCnt(shapes.size());
160+ size_t size = 1;
161+ for (uint32_t i = 0; i < shapes.size(); i++) {
162+ size *= shapes[i];
163+ }
164+ uint32_t data_len = size * GetDataTypeSize(data_type);
165+ int32_t* pData = new (std::nothrow) int32_t[data_len];
166+ for (uint32_t i = 0; i < size; ++i) {
167+ *(pData + i) = value;
168+ }
169+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
170+ return SUCCESS;
171+}
172+ 
173+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
174+{
175+ FILE* fp;
176+ fp = fopen(bin_file.c_str(), "w");
177+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
178+ fclose(fp);
179+ return SUCCESS;
180+}
181+ 
182+int CreateOppInGraph(
183+ DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, std::vector<Operator>& outputs,
184+ Graph& graph)
185+{
186+ Status ret = SUCCESS;
187+ // 自定义代码：添加单算子定义到图中
188+ auto add1 = op::SoftPlusV2("add1");
189+ std::vector<int64_t> xShape = {32, 4, 4, 4};
190+ ADD_INPUT(1, x, inDtype, xShape);
191+ ADD_OUTPUT(1, y, inDtype, xShape);
192+ 
193+ outputs.push_back(add1);
194+ // 添加完毕
195+ return SUCCESS;
196+}
197+ 
198+int main(int argc, char* argv[])
199+{
200+ const char* graph_name = "tc_ge_irrun_test";
201+ Graph graph(graph_name);
202+ std::vector<ge::Tensor> input;
203+ 
204+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
205+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
206+ Status ret = ge::GEInitialize(global_options);
207+ if (ret != SUCCESS) {
208+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
209+ return FAILED;
210+ }
211+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
212+ 
213+ std::vector<Operator> inputs{};
214+ std::vector<Operator> outputs{};
215+ 
216+ std::cout << argv[1] << std::endl;
217+ char* endptr;
218+ 
219+ DataType inDtype = DT_FLOAT;
220+ 
221+ std::cout << inDtype << std::endl;
222+ 
223+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
224+ if (ret != SUCCESS) {
225+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
226+ return FAILED;
227+ }
228+ 
229+ if (!inputs.empty() && !outputs.empty()) {
230+ graph.SetInputs(inputs).SetOutputs(outputs);
231+ }
232+ 
233+ std::map<AscendString, AscendString> build_options = {
234+ 
235+ };
236+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
237+ ge::Session* session = new Session(build_options);
238+ 
239+ if (session == nullptr) {
240+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
241+ return FAILED;
242+ }
243+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
244+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
245+ 
246+ std::map<AscendString, AscendString> graph_options = {
247+ 
248+ };
249+ uint32_t graph_id = 0;
250+ ret = session->AddGraph(graph_id, graph, graph_options);
251+ 
252+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
253+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
254+ std::string file_path = "./dump";
255+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
256+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
257+ std::vector<ge::Tensor> output;
258+ ret = session->RunGraph(graph_id, input, output);
259+ if (ret != SUCCESS) {
260+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
261+ delete session;
262+ GEFinalize();
263+ return FAILED;
264+ }
265+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
266+ 
267+ int input_num = input.size();
268+ for (int i = 0; i < input_num; i++) {
269+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
270+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
271+ uint8_t* input_data_i = input[i].GetData();
272+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
273+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
274+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
275+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
276+ }
277+ 
278+ int output_num = output.size();
279+ for (int i = 0; i < output_num; i++) {
280+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
281+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
282+ uint8_t* output_data_i = output[i].GetData();
283+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
284+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
285+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
286+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
287+ }
288+ 
289+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
290+ std::string error_str(error_msg.GetString());
291+ std::cout << "Error message: " << error_str << std::endl;
292+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
293+ std::string warning_str(warning_msg.GetString());
294+ std::cout << "Warning message: " << warning_str << std::endl;
295+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
296+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
297+ ret = ge::GEFinalize();
298+ if (ret != SUCCESS) {
299+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
300+ return FAILED;
301+ }
302+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
303+ return SUCCESS;
304+}
Aexperimental/activation/soft_plus_v2/op_graph/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 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
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+add_graph_plugin_sources()
Aexperimental/activation/soft_plus_v2/op_graph/fusion_pass/CMakeLists.txt+10-0
@@ -0,0 +1,10 @@
1+# This program is free software, you can redistribute it and/or modify.
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This file is a part of the CANN Open Software.
4+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
7+# 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+ 
Aexperimental/activation/soft_plus_v2/op_graph/soft_plus_v2_graph_infer.cpp+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+/*!
22+ * \file soft_plus_v2_graph_infer.cpp
23+ * \brief soft_plus_v2 operater graph infer resource
24+ */
25+#include "register/op_impl_registry.h"
26+#include "log/log.h"
27+ 
28+namespace ops {
29+using namespace ge;
30+ 
31+static constexpr int64_t IDX_0 = 0;
32+ 
33+static ge::graphStatus InferDataTypeSoftPlusV2(gert::InferDataTypeContext* context)
34+{
35+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeSoftPlusV2");
36+ 
37+ // 设置输出的dtype
38+ ge::DataType sizeDtype = context->GetInputDataType(IDX_0);
39+ context->SetOutputDataType(IDX_0, sizeDtype);
40+ 
41+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeSoftPlusV2");
42+ return GRAPH_SUCCESS;
43+}
44+ 
45+IMPL_OP(SoftPlusV2).InferDataType(InferDataTypeSoftPlusV2);
46+ 
47+}; // namespace ops
Aexperimental/activation/soft_plus_v2/op_graph/soft_plus_v2_proto.h+72-0
@@ -0,0 +1,72 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+/*!
22+ * \file soft_plus_v2_proto.h
23+ * \brief Declaration of the SoftPlusV2 operator prototype, which defines the interface and attributes of the SoftPlusV2 operator.
24+ *
25+ * SoftPlusV2 is a smooth approximation of the ReLU activation function,
26+ * commonly used in neural networks to introduce non-linearity while avoiding
27+ * the dying ReLU problem.
28+ */
29+#ifndef OPS_OP_PROTO_INC_SOFTPLUSV2_H_
30+#define OPS_OP_PROTO_INC_SOFTPLUSV2_H_
31+ 
32+#include "graph/operator_reg.h"
33+#include "graph/types.h"
34+ 
35+namespace ge {
36+ 
37+/**
38+*@brief Computes the SoftPlus activation function: y = ln(1 + e^x)
39+*
40+* SoftPlusV2 applies the softplus function element-wise to the input tensor.
41+* The function is a smooth alternative to the rectified linear unit (ReLU),
42+* with the mathematical property that its derivative is the sigmoid function.
43+*
44+*@par Mathematical Formula:
45+* y = log(1 + exp(x))
46+* Where exp(x) is the exponential function, and log is the natural logarithm.
47+*
48+*@par Inputs:
49+*One required input:
50+* @li x: A tensor of type float32 or float16, with any valid shape (e.g., NCHW or NHWC format for image data).
51+* Represents the input features to be activated.
52+*
53+*@par Outputs:
54+*y: A tensor with the same shape and data type as input 'x'.
55+* Each element is the result of applying the softplus function to the corresponding element in 'x'.
56+*
57+*@par Third-party framework compatibility
58+*Compatible with TensorFlow's Softplus operator (consistent in mathematical behavior).
59+*
60+*@par Constraints:
61+* - Input and output tensors must have the same data type (float32 or float16).
62+* - For large positive values of x, the output approximates x (since exp(x) dominates 1, so log(exp(x)) ≈ x).
63+* - For large negative values of x, the output approximates 0 (since exp(x) becomes negligible, so log(1) ≈ 0).
64+*/
65+REG_OP(SoftPlusV2)
66+ .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16})) // Input tensor for softplus activation
67+ .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16})) // Output tensor after applying softplus
68+ .OP_END_FACTORY_REG(SoftPlusV2)
69+ 
70+} // namespace ge
71+ 
72+#endif // OPS_OP_PROTO_INC_SOFTPLUSV2_H_
Aexperimental/activation/soft_plus_v2/op_host/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 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
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE soft_plus_v2 ACLNNTYPE aclnn)
Aexperimental/activation/soft_plus_v2/op_host/config/ascend910b/soft_plus_v2_binary.json+63-0
@@ -0,0 +1,63 @@
1+{
2+ "op_type": "SoftPlusV2",
3+ "op_list": [
4+ {
5+ "bin_filename": "SoftPlusV2_801ec1428d40ccea7b2389990f06b489",
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": "SoftPlusV2_941e887efcf7e1780c383a682905aa0a",
35+ "inputs": [
36+ {
37+ "name": "x",
38+ "index": 0,
39+ "dtype": "float32",
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": "float32",
53+ "format": "ND",
54+ "paramType": "required",
55+ "shape": [
56+ -2
57+ ],
58+ "format_match_mode": "FormatAgnostic"
59+ }
60+ ]
61+ }
62+ ]
63+}
Aexperimental/activation/soft_plus_v2/op_host/config/ascend910b/soft_plus_v2_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[SoftPlusV2]
13+default=0
Aexperimental/activation/soft_plus_v2/op_host/soft_plus_v2_def.cpp+55-0
@@ -0,0 +1,55 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+/*!
22+ * \file soft_plus_v2.cpp
23+ * \brief
24+ */
25+#include "register/op_def_registry.h"
26+ 
27+namespace ops {
28+class SoftPlusV2 : public OpDef {
29+public:
30+ explicit SoftPlusV2(const char *name) : OpDef(name) {
31+ this->Input("x")
32+ .ParamType(REQUIRED)
33+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
34+ .Format({ge::FORMAT_ND, ge::FORMAT_ND});
35+ 
36+
37+ this->Output("z")
38+ .ParamType(REQUIRED)
39+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
40+ .Format({ge::FORMAT_ND, ge::FORMAT_ND});
41+ 
42+ 
43+ OpAICoreConfig aicoreConfig;
44+ aicoreConfig.DynamicCompileStaticFlag(true)
45+ .DynamicFormatFlag(false)
46+ .DynamicRankSupportFlag(true)
47+ .DynamicShapeSupportFlag(true)
48+ .NeedCheckSupportFlag(false)
49+ .PrecisionReduceFlag(true)
50+ .ExtendCfgInfo("opFile.value", "soft_plus_v2"); // 这里制定的值会对应到kernel入口文件名.cpp
51+ this->AICore().AddConfig("ascend910b", aicoreConfig);
52+ }
53+};
54+OP_ADD(SoftPlusV2);
55+} // namespace ops
Aexperimental/activation/soft_plus_v2/op_host/soft_plus_v2_infershape.cpp+36-0
@@ -0,0 +1,36 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+/*!
21+ * \file soft_plus_v2_infer.cpp
22+ * \brief/
23+ */
24+#include "register/op_impl_registry.h"
25+#include "log/log.h"
26+ 
27+namespace ops {
28+static ge::graphStatus InferShape(gert::InferShapeContext *context) {
29+ const gert::Shape *x_shape = context->GetInputShape(0);
30+ gert::Shape *z_shape = context->GetOutputShape(0);
31+ *z_shape = *x_shape;
32+ return ge::GRAPH_SUCCESS;
33+}
34+ 
35+IMPL_OP_INFERSHAPE(SoftPlusV2).InferShape(InferShape);
36+} // namespace ops
Aexperimental/activation/soft_plus_v2/op_host/soft_plus_v2_tiling.cpp+132-0
@@ -0,0 +1,132 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+/*!
22+ * \file soft_plus_v2_tiling.cpp
23+ * \brief
24+ */
25+ 
26+#include "log/log.h"
27+#include "util/math_util.h"
28+#include "op_host/tiling_util.h"
29+#include "op_host/tiling_templates_registry.h"
30+#include "../op_kernel/soft_plus_v2_tiling_data.h"
31+#include "../op_kernel/soft_plus_v2_tiling_key.h"
32+ 
33+namespace optiling {
34+ 
35+using namespace Ops::NN::OpTiling;
36+const uint32_t BLOCK_DIM = 8;
37+const uint32_t TILE_NUM = 8;
38+const uint32_t WS_SYS_SIZE = 16U * 1024U * 1024U;
39+ 
40+struct SoftPlusV2CompileInfo {};
41+ 
42+// 获取平台信息如ubSize, coreNum
43+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
44+{
45+ // 获取ubsize coreNum
46+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
47+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
48+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
49+ coreNum = ascendcPlatform.GetCoreNumAiv();
50+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
51+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
52+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
53+ return ge::GRAPH_SUCCESS;
54+}
55+ 
56+// 获取属性，shape信息
57+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t& totalIdx, ge::DataType& dataType)
58+{
59+ // 获取输入shape信息
60+ auto inputX = context->GetInputShape(0);
61+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
62+ totalIdx = inputX->GetStorageShape().GetShapeSize();
63+ // dtype校验
64+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16};
65+ auto inputDesc = context->GetInputDesc(0);
66+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
67+ dataType = inputDesc->GetDataType();
68+ if (supportedDtype.count(dataType) == 0) {
69+ OP_LOGE(context, "invalid dtype");
70+ return ge::GRAPH_FAILED;
71+ }
72+ return ge::GRAPH_SUCCESS;
73+}
74+ 
75+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
76+{
77+ auto ascendcPlatform = platform_ascendc:: PlatformAscendC(context->GetPlatformInfo());
78+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
79+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
80+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
81+ currentWorkspace[0] = WS_SYS_SIZE + sysWorkspaceSize;
82+ return ge::GRAPH_SUCCESS;
83+}
84+ 
85+static ge::graphStatus SoftPlusV2TilingFunc(gert::TilingContext *context) {
86+ // 获取平台运行信息
87+ uint64_t ubSize;
88+ int64_t coreNum;
89+ OP_CHECK_IF(
90+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"),
91+ return ge::GRAPH_FAILED);
92+ 
93+ // 获取WorkspaceSize信息
94+ OP_CHECK_IF(
95+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
96+ return ge::GRAPH_FAILED);
97+ 
98+ // 获取shape信息
99+ int64_t totalIdx = 0;
100+ ge::DataType dataType;
101+ OP_CHECK_IF(
102+ GetShapeAttrsInfo(context, totalIdx, dataType) != ge::GRAPH_SUCCESS,
103+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
104+ 
105+ // handle empty input
106+ if (totalIdx <= 0) {
107+ SoftPlusV2TilingData* tiling = context->GetTilingData<SoftPlusV2TilingData>();
108+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
109+ memset_s(tiling, sizeof(SoftPlusV2TilingData), 0, sizeof(SoftPlusV2TilingData));
110+ context->SetBlockDim(1);
111+ return ge::GRAPH_SUCCESS;
112+ }
113+ 
114+ // 设置tiling信息
115+ SoftPlusV2TilingData* tiling = context->GetTilingData<SoftPlusV2TilingData>();
116+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
117+ OP_CHECK_IF(
118+ memset_s(tiling, sizeof(SoftPlusV2TilingData), 0, sizeof(SoftPlusV2TilingData)) != EOK,
119+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
120+ uint32_t totalLength = static_cast<uint32_t>(totalIdx);
121+ // 确保totalLength能被BLOCK_DIM整除，向上取整
122+ uint32_t alignedLength = (totalLength + BLOCK_DIM - 1) / BLOCK_DIM * BLOCK_DIM;
123+ context->SetBlockDim(BLOCK_DIM);
124+ tiling->totalLength = alignedLength;
125+ tiling->tileNum = TILE_NUM;
126+ return ge::GRAPH_SUCCESS;
127+}
128+ 
129+// tiling注册入口.
130+IMPL_OP_OPTILING(SoftPlusV2).Tiling(SoftPlusV2TilingFunc);
131+}
132+// namespace optiling
Aexperimental/activation/soft_plus_v2/op_kernel/soft_plus_v2.cpp+33-0
@@ -0,0 +1,33 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+/*!
22+ * \file soft_plus_v2.cpp
23+ * \brief
24+ */
25+#include "soft_plus_v2.h"
26+ 
27+extern "C" __global__ __aicore__ void soft_plus_v2(GM_ADDR x, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) {
28+ REGISTER_TILING_DEFAULT(SoftPlusV2TilingData);
29+ GET_TILING_DATA_WITH_STRUCT(SoftPlusV2TilingData, tilingData, tiling);
30+ NsSoftPlusV2::SoftPlusV2 op;
31+ op.Init(x, z, &tilingData);
32+ op.Process();
33+}
Aexperimental/activation/soft_plus_v2/op_kernel/soft_plus_v2.h+118-0
@@ -0,0 +1,118 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+ #ifndef SOFTPLUSV2_H
22+#define SOFTPLUSV2_H
23+ 
24+#include "kernel_operator.h"
25+#include "kernel_tiling/kernel_tiling.h"
26+#include "soft_plus_v2_tiling_data.h"
27+#include "soft_plus_v2_tiling_key.h"
28+ 
29+namespace NsSoftPlusV2 {
30+
31+using namespace AscendC;
32+constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue
33+ 
34+class SoftPlusV2 {
35+public:
36+ __aicore__ inline SoftPlusV2() {}
37+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, const SoftPlusV2TilingData* tilingData);
38+ __aicore__ inline void Process();
39+ 
40+private:
41+ __aicore__ inline void CopyIn(int32_t progress);
42+ __aicore__ inline void Compute(int32_t progress);
43+ __aicore__ inline void CopyOut(int32_t progress);
44+ 
45+private:
46+ TPipe pipe;
47+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueX;
48+ TQue<TPosition::VECOUT, BUFFER_NUM> outQueueZ;
49+ GlobalTensor<DTYPE_X> xGm;
50+ GlobalTensor<DTYPE_Z> zGm;
51+ uint32_t blockLength;
52+ uint32_t tileNum;
53+ uint32_t tileLength;
54+};
55+ 
56+__aicore__ inline void SoftPlusV2::Init(GM_ADDR x, GM_ADDR z, const SoftPlusV2TilingData* tilingData) {
57+ this->blockLength = tilingData->totalLength / GetBlockNum();
58+ this->tileNum = tilingData->tileNum;
59+ this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
60+ 
61+ xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x +
62+ this->blockLength * GetBlockIdx(),
63+ this->blockLength);
64+ zGm.SetGlobalBuffer((__gm__ DTYPE_Z *)z +
65+ this->blockLength * GetBlockIdx(),
66+ this->blockLength);
67+ 
68+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
69+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Z));
70+}
71+ 
72+// 仅保留类成员Process函数
73+__aicore__ inline void SoftPlusV2::Process() {
74+ int32_t loopCount = this->tileNum * BUFFER_NUM;
75+ for (int32_t i = 0; i < loopCount; i++) {
76+ CopyIn(i);
77+ Compute(i);
78+ CopyOut(i);
79+ }
80+}
81+ 
82+__aicore__ inline void SoftPlusV2::CopyIn(int32_t progress) {
83+ LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
84+ DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
85+ inQueueX.EnQue(xLocal);
86+}
87+ 
88+__aicore__ inline void SoftPlusV2::Compute(int32_t progress) {
89+ LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>(); // 统一用DTYPE_X而非硬编码half
90+ LocalTensor<DTYPE_Z> yLocal = outQueueZ.AllocTensor<DTYPE_Z>();
91+ LocalTensor<DTYPE_X> tmpLocal = inQueueX.AllocTensor<DTYPE_X>();
92+ 
93+ // 修正为this->tileLength
94+ Exp(tmpLocal, xLocal, this->tileLength);
95+ 
96+ // yLocal = 1.0
97+ for (uint32_t i = 0; i < this->tileLength; ++i) {
98+ yLocal.SetValue(i, static_cast<DTYPE_Z>(1.0f));
99+ }
100+ 
101+ // 修正为this->tileLength
102+ Add(tmpLocal, tmpLocal, yLocal, this->tileLength);
103+ Ln(yLocal, tmpLocal, this->tileLength);
104+ 
105+ outQueueZ.EnQue<DTYPE_Z>(yLocal);
106+ 
107+ inQueueX.FreeTensor(xLocal);
108+ inQueueX.FreeTensor(tmpLocal);
109+}
110+ 
111+__aicore__ inline void SoftPlusV2::CopyOut(int32_t progress) {
112+ LocalTensor<DTYPE_Z> zLocal = outQueueZ.DeQue<DTYPE_Z>();
113+ DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
114+ outQueueZ.FreeTensor(zLocal);
115+}
116+ 
117+} // namespace NsSoftPlusV2
118+#endif // SOFTPLUSV2_H
Aexperimental/activation/soft_plus_v2/op_kernel/soft_plus_v2_tiling_data.h+32-0
@@ -0,0 +1,32 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+/*!
22+ * \file soft_plus_v2_tiling_data.h
23+ * \brief tiling data struct
24+ */
25+#ifndef _SOFT_PLUS_V2_TILING_DATA_H_
26+#define _SOFT_PLUS_V2_TILING_DATA_H_
27+ 
28+struct SoftPlusV2TilingData {
29+ uint32_t totalLength;
30+ uint32_t tileNum;
31+};
32+#endif // _SOFT_PLUS_V2_TILING_DATA_H_
Aexperimental/activation/soft_plus_v2/op_kernel/soft_plus_v2_tiling_key.h+31-0
@@ -0,0 +1,31 @@
1+/**
2+ * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT)
3+ * and is contributed to the CANN Open Software.
4+ *
5+ * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT).
6+ * All Rights Reserved.
7+ *
8+ * Authors (accounts):
9+ * - Zhou Jiamin <@zhou-jiamin-666>
10+ * - Su Tonghua <@sutonghua>
11+ *
12+ * This program is free software: you can redistribute it and/or modify it.
13+ * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License").
14+ * You may not use this file except in compliance with the License.
15+ * See the LICENSE file at the root of the repository for the full text of the License.
16+ *
17+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED,
18+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
19+ */
20+ 
21+/*!
22+ * \file soft_plus_v2_tiling_key.h
23+ * \brief soft_plus_v2 tiling key declare
24+ */
25+ 
26+#ifndef __SOFT_PLUS_V2_TILING_KEY_H__
27+#define __SOFT_PLUS_V2_TILING_KEY_H__
28+ 
29+/* 无tilingKey实现 */
30+ 
31+#endif // __SOFT_PLUS_V2_TILING_KEY_H__