Star879
1780
代码
代码
Issues229
Pull Requests325
Actions
项目讨论
Wiki
分析
Star879
1780
已合并
个人-AscendC实现Celu算子贡献 #378
zhoujianhua创建于 2025年12月3日
个人-AscendC实现Celu算子贡献 #378
已合并
zhoujianhua创建于 2025年12月3日
从LePenseur/ops-nn_zjhfeat/celu_v2合入到cann/ops-nnmaster
共 19 个文件变更+1356-0
Aexperimental/activation/celu_v2/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
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/celu_v2/README.md+32-0
@@ -0,0 +1,32 @@
1+# Celu
2+## 产品支持情况
3+ 
4+| 产品 | 是否支持 |
5+| :----------------------------------------------------------- | :------- |
6+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
7+## 功能说明
8+ 
9+- 算子功能：激活函数Celu（Continuously Differentiable Exponential Linear Units）。
10+- 计算公式：
11+ 
12+ $$ \text{CELU}(x) = \max(0,x) + \min(0, \alpha * (\exp(x/\alpha) - 1)) $$
13+ 
14+## 参数说明
15+ 
16+| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |
17+| :----- | :------------- | :----------------------------------------------------------- | :------- | :------- |
18+| x | 输入 | 输入张量。 | 见下方 | ND |
19+| alpha | 属性 | CELU公式中的$\alpha$值，默认值为1.0。 | FLOAT | - |
20+| z | 输出 | 输出张量，shape和dtype与输入x一致。 | 同self | ND |
21+ 
22+- <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term>：数据类型支持FLOAT、FLOAT16。
23+ 
24+## 约束说明
25+ 
26+无
27+ 
28+## 调用说明
29+ 
30+| 调用方式 | 样例代码 | 说明 |
31+| :-------- | :----------------------------------------------------------- | :----------------------------------------------------------- |
32+| aclnn接口 | [test_celu](ops-nn/activation/celu_v2/examples/test_aclnn_celu_v2.cpp) | 通过[aclnnCelu]接口方式调用Celu算子。 |
Aexperimental/activation/celu_v2/examples/test_aclnn_celu_v2.cpp+172-0文件内容审核中，请稍后刷新重试
Aexperimental/activation/celu_v2/examples/test_geir_celu_v2.cpp+307-0
@@ -0,0 +1,307 @@
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 Jianhua <@LePenseur>
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/celu_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,value) \
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+ if (intputDtype == DT_FLOAT || intputDtype == DT_FLOAT16) { \
59+ ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
60+ placeholder##intputIndex##_desc, static_cast<float>(value)); \
61+ } else if(intputDtype == DT_INT32 || intputDtype == DT_INT16){ \
62+ ret = GenOnesData(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
63+ placeholder##intputIndex##_desc, intputDtype, static_cast<int>(value)); \
64+ } \
65+ if (ret != SUCCESS) { \
66+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
67+ return FAILED; \
68+ } \
69+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
70+ input.push_back(tensor_placeholder##intputIndex); \
71+ graph.AddOp(placeholder##intputIndex); \
72+ add1.set_input_##intputName(placeholder##intputIndex); \
73+ inputs.push_back(placeholder##intputIndex);
74+ 
75+#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape,value) \
76+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
77+ auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \
78+ TensorDesc placeholder##intputIndex##_desc = \
79+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
80+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
81+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
82+ Tensor tensor_placeholder##intputIndex; \
83+ if (intputDtype == DT_FLOAT || intputDtype == DT_FLOAT16) { \
84+ ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
85+ placeholder##intputIndex##_desc, static_cast<float>(value)); \
86+ } else if(intputDtype == DT_INT32 || intputDtype == DT_INT16){ \
87+ ret = GenOnesData(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
88+ placeholder##intputIndex##_desc, intputDtype, static_cast<int>(value)); \
89+ } \
90+ if (ret != SUCCESS) { \
91+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
92+ return FAILED; \
93+ } \
94+ placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \
95+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
96+ graph.AddOp(placeholder##intputIndex); \
97+ add1.set_input_##intputName(placeholder##intputIndex); \
98+ add1.update_input_desc_##intputName(placeholder##intputIndex##_desc); \
99+ inputs.push_back(placeholder##intputIndex);
100+ 
101+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
102+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
103+ add1.update_output_desc_##outputName(outputName##outputIndex##_desc);
104+ //新加
105+ #define LOG_PRINT(message, ...) \
106+ do { \
107+ printf(message, ##__VA_ARGS__); \
108+ } while (0)
109+ 
110+string GetTime()
111+{
112+ time_t timep;
113+ time(&timep);
114+ char tmp[64];
115+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
116+ return tmp;
117+}
118+ 
119+uint32_t GetDataTypeSize(DataType dt)
120+{
121+ uint32_t dilation = 1;
122+ uint32_t oneByte = 1;
123+ uint32_t twoByte = 2;
124+ uint32_t fourByte = 4;
125+ uint32_t eightByte = 8;
126+ 
127+ if (dt == ge::DT_FLOAT) {
128+ dilation = fourByte;
129+ } else if (dt == ge::DT_FLOAT16) {
130+ dilation = twoByte;
131+ }
132+ return dilation;
133+}
134+ 
135+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value)
136+{
137+ input_tensor_desc.SetRealDimCnt(shapes.size());
138+ size_t size = 1;
139+ for (uint32_t i = 0; i < shapes.size(); i++) {
140+ size *= shapes[i];
141+ }
142+ uint32_t byteSizeFloat32 = 4;
143+ uint32_t data_len = size * byteSizeFloat32;
144+ float* pData = new (std::nothrow) float[size];
145+ 
146+ for (size_t i = 0; i < size; ++i) {
147+ *(pData + i) = value;
148+ }
149+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
150+ return SUCCESS;
151+}
152+ 
153+int32_t GenOnesData(
154+ vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type, int value)
155+{
156+ input_tensor_desc.SetRealDimCnt(shapes.size());
157+ size_t size = 1;
158+ for (uint32_t i = 0; i < shapes.size(); i++) {
159+ size *= shapes[i];
160+ }
161+ uint32_t data_len = size * GetDataTypeSize(data_type);
162+ int32_t* pData = new (std::nothrow) int32_t[size];
163+ for (uint32_t i = 0; i < size; ++i) {
164+ *(pData + i) = value;
165+ }
166+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
167+ return SUCCESS;
168+}
169+ 
170+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
171+{
172+ FILE* fp;
173+ fp = fopen(bin_file.c_str(), "w");
174+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
175+ fclose(fp);
176+ return SUCCESS;
177+}
178+ 
179+int CreateOppInGraph(
180+ DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, std::vector<Operator>& outputs,
181+ Graph& graph)
182+{
183+ Status ret = SUCCESS;
184+ auto add1 = op::CeluV2("add1");
185+ std::vector<int64_t> xShape = {32, 4, 4, 4};
186+ ADD_INPUT(1, x1, inDtype, xShape,4);
187+ 
188+ ADD_OUTPUT(1, y, inDtype, xShape);
189+ 
190+ outputs.push_back(add1);
191+ // 添加完毕
192+ return SUCCESS;
193+}
194+ 
195+int main(int argc, char* argv[])
196+{
197+ const char* graph_name = "tc_ge_irrun_test";
198+ Graph graph(graph_name);
199+ std::vector<ge::Tensor> input;
200+ 
201+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
202+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
203+ Status ret = ge::GEInitialize(global_options);
204+ if (ret != SUCCESS) {
205+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
206+ return FAILED;
207+ }
208+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
209+ 
210+ std::vector<Operator> inputs{};
211+ std::vector<Operator> outputs{};
212+ std::cout << argv[1] << std::endl;
213+ char* endptr;
214+ //修改类型
215+ DataType inDtype = DT_FLOAT;
216+ 
217+ std::cout << inDtype << std::endl;
218+ 
219+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
220+ if (ret != SUCCESS) {
221+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
222+ return FAILED;
223+ }
224+ 
225+ if (!inputs.empty() && !outputs.empty()) {
226+ graph.SetInputs(inputs).SetOutputs(outputs);
227+ }
228+ 
229+ std::map<AscendString, AscendString> build_options = {
230+ 
231+ };
232+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
233+ ge::Session* session = new Session(build_options);
234+ 
235+ if (session == nullptr) {
236+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
237+ return FAILED;
238+ }
239+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
240+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
241+ 
242+ std::map<AscendString, AscendString> graph_options = {
243+ 
244+ };
245+ uint32_t graph_id = 0;
246+ ret = session->AddGraph(graph_id, graph, graph_options);
247+ 
248+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
249+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
250+ std::string file_path = "./dump";
251+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
252+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
253+ std::vector<ge::Tensor> output;
254+ ret = session->RunGraph(graph_id, input, output);
255+ if (ret != SUCCESS) {
256+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
257+ delete session;
258+ GEFinalize();
259+ return FAILED;
260+ }
261+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
262+ 
263+ int input_num = input.size();
264+ for (int i = 0; i < input_num; i++) {
265+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
266+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
267+ uint8_t* input_data_i = input[i].GetData();
268+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
269+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
270+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
271+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
272+ }
273+ 
274+ int output_num = output.size();
275+ for (int i = 0; i < output_num; i++) {
276+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
277+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
278+ uint8_t* output_data_i = output[i].GetData();
279+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
280+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
281+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
282+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
283+ //新加3行打印
284+ float* result = (float*)output_data_i;
285+ // int32_t* result = (int32_t*)output_data_i;
286+ for (int64_t j = 0; j < 8; j++) {
287+ LOG_PRINT("result[%ld] is: %f\n", j, result[j]);
288+ // LOG_PRINT("result[%ld] is: %d\n", j, result[j]);
289+ }
290+ }
291+ 
292+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
293+ std::string error_str(error_msg.GetString());
294+ std::cout << "Error message: " << error_str << std::endl;
295+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
296+ std::string warning_str(warning_msg.GetString());
297+ std::cout << "Warning message: " << warning_str << std::endl;
298+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
299+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
300+ ret = ge::GEFinalize();
301+ if (ret != SUCCESS) {
302+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
303+ return FAILED;
304+ }
305+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
306+ return SUCCESS;
307+}
Aexperimental/activation/celu_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 it.
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/celu_v2/op_graph/celu_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 Jianhua <@LePenseur>
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 celu_v2_graph_infer.cpp
23+ * \brief celu_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 InferDataTypeCeluV2(gert::InferDataTypeContext* context)
34+{
35+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeCeluV2");
36+ 
37+ // 设置输出的dtype CeluV2算子逻辑是两个数相除，因此输出dataType与输入dataType一致
38+ ge::DataType sizeDtype = context->GetInputDataType(IDX_0);
39+ context->SetOutputDataType(IDX_0, sizeDtype);
40+ 
41+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeCeluV2");
42+ return GRAPH_SUCCESS;
43+}
44+ 
45+IMPL_OP(CeluV2).InferDataType(InferDataTypeCeluV2);
46+ 
47+}; // namespace ops
Aexperimental/activation/celu_v2/op_graph/celu_v2_proto.h+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 Jianhua <@LePenseur>
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 celu_v2_proto.h
23+ * \brief
24+ */
25+#ifndef OPS_OP_PROTO_INC_CELU_V2_H_
26+#define OPS_OP_PROTO_INC_CELU_V2_H_
27+ 
28+#include "graph/operator_reg.h"
29+#include "graph/types.h"
30+ 
31+namespace ge {
32+ 
33+/**
34+*@brief Computes the Continuously Differentiable Exponential Linear Unit (CELU) of input tensors.
35+*@par Inputs:
36+*One input, including:
37+* @li x: A Tensor. Must be one of the following types: float16, float32. \n
38+ 
39+*@par Attributes:
40+* @li a: A Float. The alpha value for the CELU formulation. Defaults to 1.0. \n
41+ 
42+*@par Outputs:
43+*z: A Tensor. Has the same type and format as input "x".
44+*@par Third-party framework compatibility
45+*Compatible with the PyTorch operator CELU.
46+*/
47+REG_OP(CeluV2)
48+ .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16}))
49+ .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16}))
50+ .ATTR(alpha, Float, 1.0)
51+ .OP_END_FACTORY_REG(CeluV2)
52+ 
53+} // namespace ge
54+ 
55+#endif // OPS_OP_PROTO_INC_CELU_V2_H_
Aexperimental/activation/celu_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 it.
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/celu_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 it.
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 celu_v2 ACLNNTYPE aclnn)
Aexperimental/activation/celu_v2/op_host/celu_v2_def.cpp+59-0
@@ -0,0 +1,59 @@
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 Jianhua <@LePenseur>
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 celu_v2.cpp
23+ * \brief
24+ */
25+#include "register/op_def_registry.h"
26+ 
27+namespace ops {
28+class CeluV2 : public OpDef {
29+public:
30+ explicit CeluV2(const char* name) : OpDef(name)
31+ {
32+ this->Input("x")
33+ .ParamType(REQUIRED)
34+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
35+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
36+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
37+ .AutoContiguous();
38+ this->Output("y")
39+ .ParamType(REQUIRED)
40+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
41+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
42+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
43+ .AutoContiguous();
44+ this->Attr("alpha")
45+ .AttrType(OPTIONAL)
46+ .Float();
47+ OpAICoreConfig aicoreConfig;
48+ aicoreConfig.DynamicCompileStaticFlag(true)
49+ .DynamicFormatFlag(false)
50+ .DynamicRankSupportFlag(true)
51+ .DynamicShapeSupportFlag(true)
52+ .NeedCheckSupportFlag(false)
53+ .PrecisionReduceFlag(true)
54+ .ExtendCfgInfo("opFile.value", "celu_v2");
55+ this->AICore().AddConfig("ascend910b", aicoreConfig);
56+ }
57+};
58+OP_ADD(CeluV2);
59+} // namespace ops
Aexperimental/activation/celu_v2/op_host/celu_v2_infershape.cpp+58-0
@@ -0,0 +1,58 @@
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 Jianhua <@LePenseur>
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 celu_v2_infer.cpp
23+ * \brief
24+ */
25+#include "register/op_impl_registry.h"
26+#include "log/log.h"
27+ 
28+using namespace ge;
29+ 
30+namespace ops {
31+static constexpr int64_t IDX_0 = 0;
32+ 
33+static ge::graphStatus InferShapeCeluV2(gert::InferShapeContext* context)
34+{
35+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeCeluV2");
36+ 
37+ // get input shapes
38+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
39+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
40+ 
41+ // get output shapes
42+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
43+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
44+ 
45+ // 填充输出shape大小
46+ auto xShapeSize = xShape->GetDimNum();
47+ yShape->SetDimNum(xShapeSize);
48+ for (size_t i = 0; i < xShapeSize; i++) {
49+ int64_t dim = xShape->GetDim(i);
50+ yShape->SetDim(i, dim);
51+ }
52+ 
53+ OP_LOGD(context->GetNodeName(), "End to do InferShapeCeluV2");
54+ return GRAPH_SUCCESS;
55+}
56+ 
57+IMPL_OP_INFERSHAPE(CeluV2).InferShape(InferShapeCeluV2);
58+} // namespace ops
Aexperimental/activation/celu_v2/op_host/celu_v2_tiling.cpp+206-0
@@ -0,0 +1,206 @@
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 Jianhua <@LePenseur>
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 celu_v2_tiling.cpp
23+ * \brief
24+ */
25+ 
26+#include "log/log.h"
27+#include "util/math_util.h"
28+#include "tiling_base/tiling_util.h"
29+#include "tiling_base/tiling_templates_registry.h"
30+#include "../op_kernel/celu_v2_tiling_data.h"
31+#include "../op_kernel/celu_v2_tiling_key.h"
32+ 
33+namespace optiling {
34+ 
35+using namespace Ops::NN::OpTiling;
36+constexpr uint32_t BLOCK_SIZE = 32;
37+constexpr uint32_t BUFFER_NUM = 2;
38+constexpr uint32_t WS_SYS_SIZE = 512U;
39+ 
40+struct CeluV2CompileInfo {};
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+// tiling 分发入口
86+static ge::graphStatus CeluV2TilingFunc(gert::TilingContext* context)
87+{
88+ // 1、获取平台运行信息
89+ uint64_t ubSize = 0;
90+ int64_t coreNum = 0;
91+ OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
92+ OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
93+ 
94+ // 2、获取shape、属性信息
95+ int64_t totalIdx = 0;
96+ ge::DataType dataType;
97+ OP_CHECK_IF(GetShapeAttrsInfo(context, totalIdx, dataType) != ge::GRAPH_SUCCESS,
98+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
99+ 
100+ // 处理空输入
101+ if (totalIdx <= 0) {
102+ CeluV2TilingData* tiling = context->GetTilingData<CeluV2TilingData>();
103+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
104+ memset_s(tiling, sizeof(CeluV2TilingData), 0, sizeof(CeluV2TilingData));
105+ context->SetBlockDim(1);
106+ context->SetTilingKey(GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0));
107+ return ge::GRAPH_SUCCESS;
108+ }
109+ 
110+ // 3、获取WorkspaceSize信息
111+ OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
112+ OP_LOGE(context, "GetWorkspaceSize error"), return ge::GRAPH_FAILED);
113+ 
114+ // 4、设置tiling信息
115+ CeluV2TilingData* tiling = context->GetTilingData<CeluV2TilingData>();
116+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
117+ OP_CHECK_IF(memset_s(tiling, sizeof(CeluV2TilingData), 0, sizeof(CeluV2TilingData)) != EOK,
118+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
119+ 
120+ auto attrs = context->GetAttrs();
121+ // 默认取1，实际alpha的值由aclnn测试文件传入
122+ float alpha = 1.0f;
123+ if (attrs) {
124+ const float* attrA = attrs->GetFloat(0);
125+ if (attrA != nullptr) {
126+ alpha = *attrA;
127+ }
128+ }
129+ OP_LOGI(context, "CeluV2 tiling: attr a final value = %f", alpha);
130+ float invAlpha = 1.0f / alpha;
131+ 
132+ uint32_t typeLength = 0;
133+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
134+ if (typeLength == 0) {
135+ OP_LOGE(context, "typeLength is 0");
136+ return ge::GRAPH_FAILED;
137+ }
138+ uint64_t inputBytes = static_cast<uint64_t>(typeLength);
139+ uint64_t inputLengthBytes = static_cast<uint64_t>(totalIdx) * inputBytes;
140+ 
141+ // ub-based tileBlockNum guard (避免为0)
142+ uint32_t ubDataNumber = (inputBytes == 1ULL) ? 5U : 3U;
143+ uint64_t tmp = (ubSize / BLOCK_SIZE / BUFFER_NUM);
144+ uint32_t tileBlockNum = 1U;
145+ if (tmp > 0) {
146+ uint64_t tb = tmp / ubDataNumber;
147+ tileBlockNum = (tb == 0) ? 1U : static_cast<uint32_t>(tb);
148+ }
149+ 
150+ // 每个 tile 包含的元素数（至少 1）
151+ uint32_t tileDataNum = static_cast<uint32_t>((static_cast<uint64_t>(tileBlockNum) * BLOCK_SIZE) / inputBytes);
152+ if (tileDataNum == 0U) tileDataNum = 1U;
153+ 
154+ // 总 block 数（向上取整）
155+ uint64_t blocksTotal = (inputLengthBytes + BLOCK_SIZE - 1ULL) / BLOCK_SIZE;
156+ uint64_t coreNum64 = static_cast<uint64_t>(coreNum);
157+ if (coreNum64 > blocksTotal) coreNum64 = blocksTotal;
158+ if (coreNum64 == 0ULL) coreNum64 = 1ULL; // 最少 1 core
159+ uint32_t finalCoreNum = static_cast<uint32_t>(coreNum64);
160+ 
161+ uint64_t everyCoreInputBlockNum = blocksTotal / coreNum64; // 基本块数
162+ uint32_t tailBlockNum = static_cast<uint32_t>(blocksTotal % coreNum64); // 前 tailBlockNum 个核是 big-core
163+ 
164+ // small-core 数量（元素）
165+ uint64_t smallCoreDataNum_u = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
166+ uint32_t smallCoreDataNum = static_cast<uint32_t>(smallCoreDataNum_u);
167+ 
168+ uint32_t smallTileNum = static_cast<uint32_t>(everyCoreInputBlockNum / static_cast<uint64_t>(tileBlockNum));
169+ uint32_t finalSmallTileNum = ((everyCoreInputBlockNum % tileBlockNum) == 0) ? smallTileNum : (smallTileNum + 1);
170+ int64_t smallTailDataNum_i = static_cast<int64_t>(smallCoreDataNum) - static_cast<int64_t>(tileDataNum) * static_cast<int64_t>(smallTileNum);
171+ uint32_t smallTailDataNum = (smallTailDataNum_i <= 0) ? tileDataNum : static_cast<uint32_t>(smallTailDataNum_i);
172+ 
173+ // big-core（每个多一个 block）
174+ uint64_t bigEveryCoreBlockNum = everyCoreInputBlockNum + 1ULL;
175+ uint64_t bigCoreDataNum_u = bigEveryCoreBlockNum * BLOCK_SIZE / inputBytes;
176+ uint32_t bigCoreDataNum = static_cast<uint32_t>(bigCoreDataNum_u);
177+ uint32_t bigTileNum = static_cast<uint32_t>(bigEveryCoreBlockNum / static_cast<uint64_t>(tileBlockNum));
178+ uint32_t finalBigTileNum = ((bigEveryCoreBlockNum % tileBlockNum) == 0) ? bigTileNum : (bigTileNum + 1);
179+ int64_t bigTailDataNum_i = static_cast<int64_t>(bigCoreDataNum) - static_cast<int64_t>(tileDataNum) * static_cast<int64_t>(bigTileNum);
180+ uint32_t bigTailDataNum = (bigTailDataNum_i <= 0) ? tileDataNum : static_cast<uint32_t>(bigTailDataNum_i);
181+ 
182+ // write back
183+ tiling->smallCoreDataNum = static_cast<int64_t>(smallCoreDataNum);
184+ tiling->bigCoreDataNum = static_cast<int64_t>(bigCoreDataNum);
185+ tiling->tileDataNum = static_cast<int64_t>(tileDataNum);
186+ tiling->smallTailDataNum = static_cast<int64_t>(smallTailDataNum);
187+ tiling->bigTailDataNum = static_cast<int64_t>(bigTailDataNum);
188+ tiling->finalSmallTileNum = static_cast<int64_t>(finalSmallTileNum);
189+ tiling->finalBigTileNum = static_cast<int64_t>(finalBigTileNum);
190+ tiling->tailBlockNum = static_cast<int64_t>(tailBlockNum);
191+ tiling->alpha = alpha;
192+ tiling->invAlpha = invAlpha;
193+ 
194+ context->SetBlockDim(finalCoreNum);
195+ return ge::GRAPH_SUCCESS;
196+}
197+ 
198+ 
199+static ge::graphStatus TilingParseForCeluV2([[maybe_unused]] gert::TilingParseContext* context)
200+{
201+ return ge::GRAPH_SUCCESS;
202+}
203+ 
204+// tiling注册入口.
205+IMPL_OP_OPTILING(CeluV2).Tiling(CeluV2TilingFunc).TilingParse<CeluV2CompileInfo>(TilingParseForCeluV2);
206+} // namespace optiling
Aexperimental/activation/celu_v2/op_host/config/ascend910b/celu_v2_binary.json+77-0
@@ -0,0 +1,77 @@
1+{
2+ "op_type": "CeluV2",
3+ "op_list": [
4+ {
5+ "bin_filename": "CeluV2_a1532827238e1555db7b997c7bce2928",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float32",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ }
18+ ],
19+ "attrs": [
20+ {
21+ "name": "alpha",
22+ "dtype": "float",
23+ "value": null
24+ }
25+ ],
26+ "outputs": [
27+ {
28+ "name": "y",
29+ "index": 0,
30+ "dtype": "float32",
31+ "format": "ND",
32+ "paramType": "required",
33+ "shape": [
34+ -2
35+ ],
36+ "format_match_mode": "FormatAgnostic"
37+ }
38+ ]
39+ },
40+ {
41+ "bin_filename": "CeluV2_11132827238e1555db7b997c7bce2928",
42+ "inputs": [
43+ {
44+ "name": "x",
45+ "index": 0,
46+ "dtype": "float16",
47+ "format": "ND",
48+ "paramType": "required",
49+ "shape": [
50+ -2
51+ ],
52+ "format_match_mode": "FormatAgnostic"
53+ }
54+ ],
55+ "attrs": [
56+ {
57+ "name": "alpha",
58+ "dtype": "float",
59+ "value": null
60+ }
61+ ],
62+ "outputs": [
63+ {
64+ "name": "y",
65+ "index": 0,
66+ "dtype": "float16",
67+ "format": "ND",
68+ "paramType": "required",
69+ "shape": [
70+ -2
71+ ],
72+ "format_match_mode": "FormatAgnostic"
73+ }
74+ ]
75+ }
76+ ]
77+}
Aexperimental/activation/celu_v2/op_host/config/ascend910b/celu_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+[CeluV2]
13+default=0
Aexperimental/activation/celu_v2/op_kernel/celu_v2.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 Jianhua <@LePenseur>
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 celu_v2.cpp
23+ * \brief
24+ */
25+ 
26+#include "celu_v2.h"
27+ 
28+template <uint32_t schMode>
29+__global__ __aicore__ void celu_v2(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
30+{
31+ REGISTER_TILING_DEFAULT(CeluV2TilingData);
32+ GET_TILING_DATA_WITH_STRUCT(CeluV2TilingData, tilingData, tiling);
33+ NsCeluV2::CeluV2<DTYPE_X> op;
34+ op.Init(x, y, &tilingData);
35+ op.Process();
36+}
Aexperimental/activation/celu_v2/op_kernel/celu_v2.h+151-0
@@ -0,0 +1,151 @@
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 Jianhua <@LePenseur>
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 celu_v2.h
23+ * \brief
24+ */
25+#ifndef CELU_V2_H
26+#define CELU_V2_H
27+ 
28+#include "kernel_operator.h"
29+#include "kernel_tiling/kernel_tiling.h"
30+#include "celu_v2_tiling_data.h"
31+#include "celu_v2_tiling_key.h"
32+ 
33+namespace NsCeluV2 {
34+ 
35+using namespace AscendC;
36+ 
37+constexpr int32_t BUFFER_NUM = 2;
38+ 
39+template <typename T>
40+class CeluV2 {
41+public:
42+ __aicore__ inline CeluV2(){};
43+ 
44+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const CeluV2TilingData* tilingData);
45+ __aicore__ inline void Process();
46+ 
47+private:
48+ __aicore__ inline void CopyIn(int32_t progress);
49+ __aicore__ inline void CopyOut(int32_t progress);
50+ __aicore__ inline void Compute(int32_t progress);
51+ 
52+private:
53+ AscendC::TPipe pipe;
54+ AscendC::TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX;
55+ AscendC::TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueY;
56+ AscendC::TBuf<QuePosition::VECCALC> tmpBuffer1, tmpBuffer2;
57+ AscendC::GlobalTensor<T> inputGMX;
58+ AscendC::GlobalTensor<T> outputGMY;
59+ 
60+ uint32_t coreDataNum;
61+ uint32_t tileNum;
62+ uint32_t tileDataNum;
63+ uint32_t tailDataNum;
64+ uint32_t processDataNum;
65+ float alpha;
66+ float invAlpha;
67+};
68+ 
69+template <typename T>
70+__aicore__ inline void CeluV2<T>::Init(GM_ADDR x, GM_ADDR y, const CeluV2TilingData* tilingData)
71+{
72+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
73+ uint32_t coreNum = AscendC::GetBlockIdx();
74+ uint32_t globalBufferIndex = tilingData->bigCoreDataNum * AscendC::GetBlockIdx();
75+ this->tileDataNum = tilingData->tileDataNum;
76+ this->alpha = tilingData->alpha;
77+ this->invAlpha = tilingData->invAlpha;
78+ if (coreNum < tilingData->tailBlockNum) {
79+ this->coreDataNum = tilingData->bigCoreDataNum;
80+ this->tileNum = tilingData->finalBigTileNum;
81+ this->tailDataNum = tilingData->bigTailDataNum;
82+ }
83+ else {
84+ this->coreDataNum = tilingData->smallCoreDataNum;
85+ this->tileNum = tilingData->finalSmallTileNum;
86+ this->tailDataNum = tilingData->smallTailDataNum;
87+ globalBufferIndex -= (tilingData->bigCoreDataNum - tilingData->smallCoreDataNum) * (AscendC::GetBlockIdx() - tilingData->tailBlockNum);
88+ }
89+ inputGMX.SetGlobalBuffer((__gm__ T*)x + globalBufferIndex, this->coreDataNum);
90+ outputGMY.SetGlobalBuffer((__gm__ T*)y + globalBufferIndex, this->coreDataNum);
91+ pipe.InitBuffer(inputQueueX, BUFFER_NUM, this->tileDataNum * sizeof(T));
92+ pipe.InitBuffer(outputQueueY, BUFFER_NUM, this->tileDataNum * sizeof(T));
93+ pipe.InitBuffer(tmpBuffer1, this->tileDataNum * sizeof(T));
94+ pipe.InitBuffer(tmpBuffer2, this->tileDataNum * sizeof(T));
95+}
96+ 
97+template <typename T>
98+__aicore__ inline void CeluV2<T>::CopyIn(int32_t progress)
99+{
100+ AscendC::LocalTensor<T> xLocal = inputQueueX.AllocTensor<T>();
101+ AscendC::DataCopy(xLocal, inputGMX[progress * this->tileDataNum], this->processDataNum);
102+ inputQueueX.EnQue(xLocal);
103+}
104+ 
105+template <typename T>
106+__aicore__ inline void CeluV2<T>::CopyOut(int32_t progress)
107+{
108+ AscendC::LocalTensor<T> yLocal = outputQueueY.DeQue<T>();
109+ AscendC::DataCopy(outputGMY[progress * this->tileDataNum], yLocal, this->processDataNum);
110+ outputQueueY.FreeTensor(yLocal);
111+}
112+ 
113+template <typename T>
114+__aicore__ inline void CeluV2<T>::Compute(int32_t progress)
115+{
116+ AscendC::LocalTensor<T> xLocal = inputQueueX.DeQue<T>();
117+ AscendC::LocalTensor<T> yLocal = outputQueueY.AllocTensor<T>();
118+ LocalTensor<T> tmpTensor1 = tmpBuffer1.Get<T>();
119+ LocalTensor<T> tmpTensor2 = tmpBuffer2.Get<T>();
120+ AscendC::Muls(tmpTensor1, xLocal, static_cast<T>(invAlpha), this->tileDataNum);
121+ AscendC::PipeBarrier<PIPE_V>();
122+ AscendC::Exp(tmpTensor1, tmpTensor1, this->tileDataNum);
123+ AscendC::PipeBarrier<PIPE_V>();
124+ AscendC::Adds(tmpTensor1, tmpTensor1, static_cast<T>(-1), this->tileDataNum);
125+ AscendC::Muls(tmpTensor1, tmpTensor1, static_cast<T>(alpha), this->tileDataNum);
Z
Zzhajianqing1232025年12月21日

请检查相关计算逻辑，计算过程有计算前后依赖的地方，建议手动加同步。示例：PipeBarrier<PIPE_V>();

likedislike
126+ AscendC::PipeBarrier<PIPE_V>();
127+ AscendC::Mins(tmpTensor1, tmpTensor1, static_cast<T>(0), this->tileDataNum);
128+ AscendC::Maxs(tmpTensor2, xLocal, static_cast<T>(0), this->tileDataNum);
129+ AscendC::PipeBarrier<PIPE_V>();
130+ AscendC::Add(yLocal, tmpTensor2, tmpTensor1, this->tileDataNum);
131+ outputQueueY.EnQue<T>(yLocal);
132+ inputQueueX.FreeTensor(xLocal);
133+}
134+ 
135+template <typename T>
136+__aicore__ inline void CeluV2<T>::Process()
137+{
138+ int32_t loopCount = this->tileNum;
139+ this->processDataNum = this->tileDataNum;
140+ for (int32_t i = 0; i < loopCount; i++) {
141+ if (i == this->tileNum - 1) {
142+ this->processDataNum = this->tailDataNum;
143+ }
144+ CopyIn(i);
145+ Compute(i);
146+ CopyOut(i);
147+ }
148+}
149+ 
150+} // namespace NsCeluV2
151+#endif // LEAKYRELU_V2_H
Aexperimental/activation/celu_v2/op_kernel/celu_v2_tiling_data.h+41-0
@@ -0,0 +1,41 @@
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 Jianhua <@LePenseur>
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 celu_v2_tiling_data.h
23+ * \brief tiling data struct
24+ */
25+ 
26+#ifndef _ROTARY_POSITION_EMBEDDING_GRAD_TILING_DATA_H_
27+#define _ROTARY_POSITION_EMBEDDING_GRAD_TILING_DATA_H_
28+ 
29+struct CeluV2TilingData {
30+ int64_t smallCoreDataNum;
31+ int64_t bigCoreDataNum;
32+ int64_t finalBigTileNum;
33+ int64_t finalSmallTileNum;
34+ int64_t tileDataNum;
35+ int64_t smallTailDataNum;
36+ int64_t bigTailDataNum;
37+ int64_t tailBlockNum;
38+ float alpha;
39+ float invAlpha;
40+};
41+#endif
Aexperimental/activation/celu_v2/op_kernel/celu_v2_tiling_key.h+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 Jianhua <@LePenseur>
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 celu_v2_tiling_key.h
23+ * \brief celu_v2 tiling key declare
24+ */
25+ 
26+#ifndef __CELU_V2_TILING_KEY_H__
27+#define __CELU_V2_TILING_KEY_H__
28+ 
29+#include "ascendc/host_api/tiling/template_argument.h"
30+ 
31+#define ELEMENTWISE_TPL_SCH_MODE_0 0
32+#define ELEMENTWISE_TPL_SCH_MODE_1 1
33+ 
34+ASCENDC_TPL_ARGS_DECL(CeluV2,
35+ ASCENDC_TPL_UINT_DECL(schMode, 1,
36+ ASCENDC_TPL_UI_LIST,
37+ ELEMENTWISE_TPL_SCH_MODE_0,
38+ ELEMENTWISE_TPL_SCH_MODE_1),);
39+ 
40+ASCENDC_TPL_SEL(
41+ ASCENDC_TPL_ARGS_SEL(
42+ ASCENDC_TPL_UINT_SEL(schMode,
43+ ASCENDC_TPL_UI_LIST,
44+ ELEMENTWISE_TPL_SCH_MODE_0,
45+ ELEMENTWISE_TPL_SCH_MODE_1)),);
46+ 
47+#endif
Mscripts/kernel/binary_config/ascendc_config.json+1-0
@@ -591,6 +591,7 @@
591 {"name":"QuantizedBatchNorm", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},591 {"name":"QuantizedBatchNorm", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},
592 {"name":"UnfoldGrad", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync": true},592 {"name":"UnfoldGrad", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync": true},
593 {"name":"Bincount", "compute_units": ["ascend910_95"], "auto_sync" : false, "impl_mode" : ""},593 {"name":"Bincount", "compute_units": ["ascend910_95"], "auto_sync" : false, "impl_mode" : ""},
594+ {"name":"CeluV2", "compute_units": ["ascend910b"], "auto_sync" : true, "impl_mode" : "high_performance"},
594 {"name":"Addr", "compute_units": ["ascend910_95"], "auto_sync" : false, "impl_mode" : ""},595 {"name":"Addr", "compute_units": ["ascend910_95"], "auto_sync" : false, "impl_mode" : ""},
595 {"name":"SparseTensorDenseMatMul", "compute_units": ["ascend910_95"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend910_95": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}}596 {"name":"SparseTensorDenseMatMul", "compute_units": ["ascend910_95"], "auto_sync" : false, "impl_mode" : "", "compile_options": {"ascend910_95": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}}
596]597]