Star1291
1406
代码
代码
Issues106
Pull Requests204
Actions
项目讨论
Wiki
分析
Star1291
1406
已合并
个人-AscendC实现Sinh算子贡献 #142
peihaobo创建于 2025年10月31日
个人-AscendC实现Sinh算子贡献 #142
已合并
peihaobo创建于 2025年10月31日
从xiaopei-1/ops-math_1062master合入到cann/ops-mathmaster
共 20 个文件变更+1361-1
Aexperimental/math/sinh/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 AND NOT BENCHMARK)
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/math/sinh/examples/test_aclnn_sinh.cpp+171-0
@@ -0,0 +1,171 @@
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+ * - Pei Haobo<@xiaopei-1>
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_sinh.h"
25+ 
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+ // 修改测试数据类型
51+ std::vector<float> resultData(size, 0);
52+ auto ret = aclrtMemcpy(
53+ resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]),
54+ ACL_MEMCPY_DEVICE_TO_HOST);
55+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return);
56+ for (int64_t i = 0; i < size; i++) {
57+ LOG_PRINT("mean result[%ld] is: %f\n", i, resultData[i]);
58+ }
59+}
60+ 
61+int Init(int32_t deviceId, aclrtStream* stream)
62+{
63+ // 固定写法，初始化
64+ auto ret = aclInit(nullptr);
65+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
66+ ret = aclrtSetDevice(deviceId);
67+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
68+ ret = aclrtCreateStream(stream);
69+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
70+ return 0;
71+}
72+ 
73+template <typename T>
74+int CreateAclTensor(
75+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
76+ aclTensor** tensor)
77+{
78+ auto size = GetShapeSize(shape) * sizeof(T);
79+ // 2. 申请device侧内存
80+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
81+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
82+ // 3. 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
83+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
84+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
85+ 
86+ // 计算连续tensor的strides
87+ std::vector<int64_t> strides(shape.size(), 1);
88+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
89+ strides[i] = shape[i + 1] * strides[i + 1];
90+ }
91+ 
92+ // 调用aclCreateTensor接口创建aclTensor
93+ *tensor = aclCreateTensor(
94+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
95+ *deviceAddr);
96+ return 0;
97+}
98+ 
99+int main()
100+{
101+ // 1. 调用acl进行device/stream初始化
102+ int32_t deviceId = 0;
103+ aclrtStream stream;
104+ auto ret = Init(deviceId, &stream);
105+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
106+ 
107+ // 2. 构造输入与输出，需要根据API的接口自定义构造
108+ aclTensor* selfX = nullptr;
109+ void* selfXDeviceAddr = nullptr;
110+ std::vector<int64_t> selfXShape = {4,8};
111+ std::vector<float> selfXHostData(32, 1);
112+ ret = CreateAclTensor(selfXHostData, selfXShape, &selfXDeviceAddr, aclDataType::ACL_FLOAT, &selfX);
113+ CHECK_RET(ret == ACL_SUCCESS, return ret);
114+ 
115+ aclTensor* out = nullptr;
116+ void* outDeviceAddr = nullptr;
117+ std::vector<int64_t> outShape = {4,8};
118+ std::vector<float> outHostData(32, 1);
119+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
120+ CHECK_RET(ret == ACL_SUCCESS, return ret);
121+ 
122+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
123+ uint64_t workspaceSize = 0;
124+ aclOpExecutor* executor;
125+ 
126+ // 4. 调用aclnnSinh第一段接口
127+ ret = aclnnSinhGetWorkspaceSize(selfX, out, &workspaceSize, &executor);
128+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSinhGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
129+ if (ret != ACL_SUCCESS) {
130+ const char* errMsg = aclGetRecentErrMsg(); // ← 关键一行
131+ LOG_PRINT("[ERROR] aclnnCeilV2GetWorkspaceSize failed, ret = %d\n", ret);
132+ LOG_PRINT(" recent AscendCL msg: %s\n",
133+ errMsg ? errMsg : "no additional message");
134+ return ret;
135+ }
136+ // 根据第一段接口计算出的workspaceSize申请device内存
137+ void* workspaceAddr = nullptr;
138+ if (workspaceSize > static_cast<uint64_t>(0)) {
139+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
140+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
141+ }
142+ 
143+ // 5. 调用aclnnSinh第二段接口
144+ ret = aclnnSinh(workspaceAddr, workspaceSize, executor, stream);
145+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSinh failed. ERROR: %d\n", ret); return ret);
146+ 
147+ // 6. （固定写法）同步等待任务执行结束
148+ ret = aclrtSynchronizeStream(stream);
149+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
150+ 
151+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
152+ PrintOutResult(outShape, &outDeviceAddr);
153+ 
154+ // 7. 释放aclTensor，需要根据具体API的接口定义修改
155+ aclDestroyTensor(selfX);
156+ aclDestroyTensor(out);
157+ 
158+ // 8. 释放device资源
159+ aclrtFree(selfXDeviceAddr);
160+ aclrtFree(outDeviceAddr);
161+ if (workspaceSize > static_cast<uint64_t>(0)) {
162+ aclrtFree(workspaceAddr);
163+ }
164+ aclrtDestroyStream(stream);
165+ aclrtResetDevice(deviceId);
166+ 
167+ // 9. acl去初始化
168+ aclFinalize();
169+ 
170+ return 0;
171+}
Aexperimental/math/sinh/examples/test_geir_sinh.cpp+323-0
@@ -0,0 +1,323 @@
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+ * - Pei Haobo<@xiaopei-1>
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/sinh_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+ } else if (dt == ge::DT_BF16) {
132+ dilation = twoByte;
133+ } else if (dt == ge::DT_INT16) {
134+ dilation = twoByte;
135+ } else if (dt == ge::DT_UINT16) {
136+ dilation = twoByte;
137+ } else if (dt == ge::DT_INT32) {
138+ dilation = fourByte;
139+ } else if (dt == ge::DT_UINT32) {
140+ dilation = fourByte;
141+ } else if (dt == ge::DT_INT64) {
142+ dilation = eightByte;
143+ } else if (dt == ge::DT_UINT64) {
144+ dilation = eightByte;
145+ } else if (dt == ge::DT_INT8) {
146+ dilation = oneByte;
147+ }
148+ return dilation;
149+}
150+ 
151+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value)
152+{
153+ input_tensor_desc.SetRealDimCnt(shapes.size());
154+ size_t size = 1;
155+ for (uint32_t i = 0; i < shapes.size(); i++) {
156+ size *= shapes[i];
157+ }
158+ uint32_t byteSizeFloat32 = 4;
159+ uint32_t data_len = size * byteSizeFloat32;
160+ float* pData = new (std::nothrow) float[size];
161+ 
162+ for (size_t i = 0; i < size; ++i) {
163+ *(pData + i) = value;
164+ }
165+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
166+ return SUCCESS;
167+}
168+ 
169+int32_t GenOnesData(
170+ vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type, int value)
171+{
172+ input_tensor_desc.SetRealDimCnt(shapes.size());
173+ size_t size = 1;
174+ for (uint32_t i = 0; i < shapes.size(); i++) {
175+ size *= shapes[i];
176+ }
177+ uint32_t data_len = size * GetDataTypeSize(data_type);
178+ int32_t* pData = new (std::nothrow) int32_t[size];
179+ for (uint32_t i = 0; i < size; ++i) {
180+ *(pData + i) = value;
181+ }
182+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
183+ return SUCCESS;
184+}
185+ 
186+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
187+{
188+ FILE* fp;
189+ fp = fopen(bin_file.c_str(), "w");
190+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
191+ fclose(fp);
192+ return SUCCESS;
193+}
194+ 
195+int CreateOppInGraph(
196+ DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, std::vector<Operator>& outputs,
197+ Graph& graph)
198+{
199+ Status ret = SUCCESS;
200+ // 自定义代码：添加单算子定义到图中
201+ auto add1 = op::Sinh("add1");
202+ std::vector<int64_t> xShape = {32, 4, 4, 4};
203+ ADD_INPUT(1, x, inDtype, xShape,1.0f);
204+ 
205+ ADD_OUTPUT(1, y, inDtype, xShape);
206+ 
207+ outputs.push_back(add1);
208+ // 添加完毕
209+ return SUCCESS;
210+}
211+ 
212+int main(int argc, char* argv[])
213+{
214+ const char* graph_name = "tc_ge_irrun_test";
215+ Graph graph(graph_name);
216+ std::vector<ge::Tensor> input;
217+ 
218+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
219+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
220+ Status ret = ge::GEInitialize(global_options);
221+ if (ret != SUCCESS) {
222+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
223+ return FAILED;
224+ }
225+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
226+ 
227+ std::vector<Operator> inputs{};
228+ std::vector<Operator> outputs{};
229+ 
230+ std::cout << argv[1] << std::endl;
231+ char* endptr;
232+ //修改类型
233+ DataType inDtype = DT_FLOAT;
234+ 
235+ std::cout << inDtype << std::endl;
236+ 
237+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
238+ if (ret != SUCCESS) {
239+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
240+ return FAILED;
241+ }
242+ 
243+ if (!inputs.empty() && !outputs.empty()) {
244+ graph.SetInputs(inputs).SetOutputs(outputs);
245+ }
246+ 
247+ std::map<AscendString, AscendString> build_options = {
248+ 
249+ };
250+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
251+ ge::Session* session = new Session(build_options);
252+ 
253+ if (session == nullptr) {
254+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
255+ return FAILED;
256+ }
257+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
258+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
259+ 
260+ std::map<AscendString, AscendString> graph_options = {
261+ 
262+ };
263+ uint32_t graph_id = 0;
264+ ret = session->AddGraph(graph_id, graph, graph_options);
265+ 
266+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
267+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
268+ std::string file_path = "./dump";
269+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
270+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
271+ std::vector<ge::Tensor> output;
272+ ret = session->RunGraph(graph_id, input, output);
273+ if (ret != SUCCESS) {
274+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
275+ delete session;
276+ GEFinalize();
277+ return FAILED;
278+ }
279+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
280+ 
281+ int input_num = input.size();
282+ for (int i = 0; i < input_num; i++) {
283+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
284+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
285+ uint8_t* input_data_i = input[i].GetData();
286+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
287+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
288+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
289+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
290+ }
291+ 
292+ int output_num = output.size();
293+ for (int i = 0; i < output_num; i++) {
294+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
295+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
296+ uint8_t* output_data_i = output[i].GetData();
297+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
298+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
299+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
300+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
301+ //新加3行打印
302+ float* result = (float*)output_data_i;
303+ for (int64_t j = 0; j < 8; j++) {
304+ LOG_PRINT("result[%ld] is: %f\n", j, result[j]);
305+ }
306+ }
307+ 
308+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
309+ std::string error_str(error_msg.GetString());
310+ std::cout << "Error message: " << error_str << std::endl;
311+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
312+ std::string warning_str(warning_msg.GetString());
313+ std::cout << "Warning message: " << warning_str << std::endl;
314+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
315+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
316+ ret = ge::GEFinalize();
317+ if (ret != SUCCESS) {
318+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
319+ return FAILED;
320+ }
321+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
322+ return SUCCESS;
323+}
Aexperimental/math/sinh/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/math/sinh/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/math/sinh/op_graph/sinh_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+ * - Pei Haobo<@xiaopei-1>
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 sinh_graph_infer.cpp
23+ * \brief sinh 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 InferDataTypeSinh(gert::InferDataTypeContext* context)
34+{
35+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypes=Sinh");
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 InferDataTypeSinh");
42+ return GRAPH_SUCCESS;
43+}
44+ 
45+IMPL_OP(Sinh).InferDataType(InferDataTypeSinh);
46+ 
47+}; // namespace ops
Aexperimental/math/sinh/op_graph/sinh_proto.h+50-0
@@ -0,0 +1,50 @@
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+ * - Pei Haobo<@xiaopei-1>
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 sinh_proto.h
23+ * \brief
24+*/
25+#ifndef OPS_OP_PROTO_INC_SINH_H_
26+#define OPS_OP_PROTO_INC_SINH_H_
27+ 
28+#include "graph/operator_reg.h"
29+#include "graph/types.h"
30+ 
31+namespace ge {
32+ 
33+/**
34+*@brief Returns element-wise hyperbolic sine of "x".
35+*@par Inputs:
36+*One input, including:
37+* @li x: A ND Tensor. Must be one of the following types: float32,float16.\n
38+*@par Outputs:
39+*y: A ND Tensor. Has the same dtype as "x".\n
40+*@par Third-party framework compatibility
41+*Compatible with the TensorFlow operator Sinh.
42+*/
43+REG_OP(Sinh)
44+ .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16}))
45+ .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16}))
46+ .OP_END_FACTORY_REG(Sinh)
47+ 
48+} // namespace ge
49+ 
50+#endif // OPS_OP_PROTO_INC_Sinh_H_
Aexperimental/math/sinh/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(OPTYPE sinh ACLNNTYPE aclnn)
Aexperimental/math/sinh/op_host/config/ascend910b/sinh_binary.json+67-0
@@ -0,0 +1,67 @@
1+{
2+ "op_type": "Sinh",
3+ "op_list": [
4+ {
5+ "bin_filename": "Sinh_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+ "outputs": [
22+ {
23+ "name": "y",
24+ "index": 0,
25+ "dtype": "float32",
26+ "format": "ND",
27+ "paramType": "required",
28+ "shape": [
29+ -2
30+ ],
31+ "format_match_mode": "FormatAgnostic"
32+ }
33+ ]
34+ },
35+ {
36+ "bin_filename": "Sinh_11132827238e1555db7b997c7bce2931",
37+ "inputs": [
38+ {
39+ "name": "x",
40+ "index": 0,
41+ "dtype": "float16",
42+ "format": "ND",
43+ "paramType": "required",
44+ "shape": [
45+ -2
46+ ],
47+ "format_match_mode": "FormatAgnostic"
48+ }
49+ ],
50+ "attrs": [
51+ ],
52+ "outputs": [
53+ {
54+ "name": "y",
55+ "index": 0,
56+ "dtype": "float16",
57+ "format": "ND",
58+ "paramType": "required",
59+ "shape": [
60+ -2
61+ ],
62+ "format_match_mode": "FormatAgnostic"
63+ }
64+ ]
65+ }
66+ ]
67+}
Aexperimental/math/sinh/op_host/config/ascend910b/sinh_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+[Sinh]
13+default=0
Aexperimental/math/sinh/op_host/sinh_def.cpp+56-0
@@ -0,0 +1,56 @@
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+ * - Pei Haobo<@xiaopei-1>
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 sinh.cpp
23+ * \brief
24+*/
25+#include "register/op_def_registry.h"
26+ 
27+namespace ops {
28+class Sinh : public OpDef {
29+public:
30+ explicit Sinh(const char* name) : OpDef(name)
31+ {
32+ this->Input("x") // 输入x1定义
33+ .ParamType(REQUIRED) // 必选输入
34+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16}) // 支持数据类型
35+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}) // 支持format格式
36+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND}) // 未确定大小shape对应format格式
37+ .AutoContiguous(); // 内存自动连续化
38+ this->Output("y") // 输出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+ OpAICoreConfig aicoreConfig;
45+ aicoreConfig.DynamicCompileStaticFlag(true)
46+ .DynamicFormatFlag(false)
47+ .DynamicRankSupportFlag(true)
48+ .DynamicShapeSupportFlag(true)
49+ .NeedCheckSupportFlag(false)
50+ .PrecisionReduceFlag(true)
51+ .ExtendCfgInfo("opFile.value", "sinh"); // 这里制定的值会对应到kernel入口文件名.cpp
52+ this->AICore().AddConfig("ascend910b", aicoreConfig); // 其他的soc版本补充部分配置项
53+ }
54+};
55+OP_ADD(Sinh); // 添加算子信息库
56+} // namespace ops
Aexperimental/math/sinh/op_host/sinh_infershape.cpp+53-0
@@ -0,0 +1,53 @@
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+ * - Pei Haobo<@xiaopei-1>
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 sinh_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 InferShapeSinh(gert::InferShapeContext* context)
34+{
35+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeSinh");
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+ *yShape = *xShape;
47+ 
48+ OP_LOGD(context->GetNodeName(), "End to do InferShapeSinh");
49+ return GRAPH_SUCCESS;
50+}
51+ 
52+IMPL_OP_INFERSHAPE(Sinh).InferShape(InferShapeSinh);
53+} // namespace ops
Aexperimental/math/sinh/op_host/sinh_tiling.cpp+229-0
@@ -0,0 +1,229 @@
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+ * - Pei Haobo<@xiaopei-1>
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 sinh_tiling.cpp
23+ * \brief
24+*/
25+#include "log/log.h"
26+#include "util/math_util.h"
27+#include "tiling_base/tiling_util.h"
28+#include "tiling_base/tiling_templates_registry.h"
29+#include "../op_kernel/sinh_tiling_data.h"
30+#include "../op_kernel/sinh_tiling_key.h"
31+ 
32+namespace optiling {
33+ 
34+using namespace Ops::Math::OpTiling;
35+constexpr uint32_t BLOCK_SIZE = 32;
36+constexpr uint32_t BUFFER_NUM = 2;
37+constexpr uint32_t WS_SYS_SIZE = 16U * 1024U * 1024U;
38+const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16};
39+ 
40+struct SinhCompileInfo {};
41+struct SinhShapeInfo {
42+ uint32_t smallCoreDataNum{0};
43+ uint32_t bigCoreDataNum{0};
44+ uint32_t finalSmallTileNum{0};
45+ uint32_t finalBigTileNum{0};
46+ uint32_t tileDataNum{0};
47+ uint32_t smallTailDataNum{0};
48+ uint32_t bigTailDataNum{0};
49+ uint32_t tailBlockNum{0};
50+ int64_t coreNum{0};
51+ };
52+// 获取平台信息如ubSize, coreNum
53+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
54+{
55+ // 获取ubsize coreNum
56+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
57+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
58+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
59+ coreNum = ascendcPlatform.GetCoreNumAiv();
60+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
61+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
62+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
63+ return ge::GRAPH_SUCCESS;
64+}
65+ 
66+// 获取属性，shape信息
67+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t& totalIdx, ge::DataType& dataType)
68+{
69+ // 获取输入shape信息
70+ auto inputX = context->GetInputShape(0);
71+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
72+ // 如果输入shape 是标量 转换为{1}，否则保持原 shape 不变
73+ auto inputShapeX = EnsureNotScalar(inputX->GetStorageShape());
74+
75+ auto outZ = context->GetOutputShape(0);
76+ OP_CHECK_NULL_WITH_CONTEXT(context, outZ);
77+ auto outShapeZ = EnsureNotScalar(outZ->GetStorageShape());
78+ 
79+ // shape校验
80+ // 校验维度数一致
81+ if (inputShapeX.GetDimNum() != outShapeZ.GetDimNum()) {
82+ OP_LOGE(
83+ context, "Sinh: inputx,outputz shape not match! dim num: x=%zu, z=%zu",
84+ inputShapeX.GetDimNum(), outShapeZ.GetDimNum());
85+ return ge::GRAPH_FAILED;
86+ } else {
87+ // 校验每个维度的大小一致
88+ size_t dimNum = inputShapeX.GetDimNum();
89+ for (size_t i = 0; i < dimNum; i++) {
90+ if (inputShapeX.GetDim(i) != outShapeZ.GetDim(i)) {
91+ OP_LOGE(
92+ context, "Sinh: inputx,outputz shape not match! dim num: x=%zu, z=%zu",
93+ inputShapeX.GetDimNum(), outShapeZ.GetDimNum());
94+ return ge::GRAPH_FAILED;
95+ }
96+ }
97+ }
98+
99+ totalIdx = inputX->GetOriginShape().GetShapeSize();
100+ // dtype校验
101+ auto inputDesc = context->GetInputDesc(0);
102+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
103+ dataType = inputDesc->GetDataType();
104+ if (supportedDtype.count(dataType) == 0) {
105+ OP_LOGE(context, "Sinh: invalid dtype! Current dtype: %d, supported dtypes: DT_FLOAT,DT_FLOAT16",
106+ static_cast<int>(dataType));
107+ return ge::GRAPH_FAILED;
108+ }
109+ return ge::GRAPH_SUCCESS;
110+}
111+ 
112+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
113+{
114+ auto ascendcPlatform = platform_ascendc:: PlatformAscendC(context->GetPlatformInfo());
115+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
116+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
117+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
118+ currentWorkspace[0] = WS_SYS_SIZE + sysWorkspaceSize;
119+ return ge::GRAPH_SUCCESS;
120+}
121+static ge::graphStatus CalculateCoreBlockNums(uint64_t ubSize, int64_t coreNum, uint32_t typeLength,int64_t totalIdx, SinhShapeInfo& info)
122+{
123+ if(0 == BLOCK_SIZE || 0 == coreNum ) {
124+ return ge::GRAPH_FAILED;
125+ }
126+ uint64_t inputBytes = static_cast<uint64_t>(typeLength);
127+ uint64_t inputLengthBytes = static_cast<uint64_t>(totalIdx) * inputBytes;
128+ 
129+ // ub-based tileBlockNum guard (避免为0)
130+ uint32_t ubDataNumber = (inputBytes == 1ULL) ? 4U : 2U;
131+ uint64_t tmp = (ubSize / BLOCK_SIZE / BUFFER_NUM);
132+ uint32_t tileBlockNum = 1U;
133+ if (tmp > 0) {
134+ uint64_t tb = tmp / ubDataNumber;
135+ tileBlockNum = (tb == 0) ? 1U : static_cast<uint32_t>(tb);
136+ }
137+ 
138+ // 每个 tile 包含的元素数（至少 1）
139+ info.tileDataNum = static_cast<uint32_t>((static_cast<uint64_t>(tileBlockNum) * BLOCK_SIZE) / inputBytes);
140+ if (info.tileDataNum == 0U) info.tileDataNum = 1U;
141+ 
142+ // 总 block 数（向上取整）
143+ uint64_t blocksTotal = (inputLengthBytes + BLOCK_SIZE - 1ULL) / BLOCK_SIZE;
144+ uint64_t coreNum64 = static_cast<uint64_t>(coreNum);
145+ if (coreNum64 > blocksTotal){
146+ coreNum64 = blocksTotal;
147+ }
148+ if (coreNum64 == 0ULL) coreNum64 = 1ULL; // 最少 1 core
149+ uint32_t finalCoreNum = static_cast<uint32_t>(coreNum64);
150+ 
151+ uint64_t everyCoreInputBlockNum = blocksTotal / coreNum64;
152+ info.tailBlockNum = static_cast<uint32_t>(blocksTotal % coreNum64);;
153+ 
154+ info.smallCoreDataNum = static_cast<uint32_t>(everyCoreInputBlockNum * BLOCK_SIZE / inputBytes);
155+ uint32_t smallTileNum = static_cast<uint32_t>(everyCoreInputBlockNum / static_cast<uint64_t>(tileBlockNum));
156+ info.finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
157+ int64_t smallTailDataNum_i = static_cast<int64_t>(info.smallCoreDataNum) - static_cast<int64_t>(info.tileDataNum) * static_cast<int64_t>(smallTileNum);
158+ info.smallTailDataNum = (smallTailDataNum_i == 0) ? info.tileDataNum : static_cast<uint32_t>(smallTailDataNum_i);
159+ 
160+ everyCoreInputBlockNum += 1ULL;
161+ info.bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
162+ uint32_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
163+ info.finalBigTileNum = ((everyCoreInputBlockNum % tileBlockNum) == 0)? bigTileNum : bigTileNum + 1;
164+ int64_t bigTailDataNum_i = static_cast<int64_t>(info.bigCoreDataNum) - static_cast<int64_t>(info.tileDataNum) * static_cast<int64_t>(bigTileNum);
165+ info.bigTailDataNum = (bigTailDataNum_i == 0) ? info.tileDataNum : static_cast<uint32_t>(bigTailDataNum_i);
166+ info.coreNum = finalCoreNum;
167+ return ge::GRAPH_SUCCESS;
168+}
169+// tiling 分发入口
170+// 可直接替换你的 SinhTilingFunc 内部实现（保留函数签名）
171+static ge::graphStatus SinhTilingFunc(gert::TilingContext* context)
172+{
173+ // 1. platform
174+ uint64_t ubSize = 0;
175+ int64_t coreNum = 0;
176+ OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
177+ OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
178+ 
179+ // 2. shapes & dtype
180+ int64_t totalIdx = 0;
181+ ge::DataType dataType;
182+ OP_CHECK_IF(GetShapeAttrsInfo(context, totalIdx, dataType) != ge::GRAPH_SUCCESS,
183+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
184+ 
185+ // 3. workspace
186+ OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
187+ OP_LOGE(context, "GetWorkspaceSize error"), return ge::GRAPH_FAILED);
188+ 
189+ SinhTilingData* tiling = context->GetTilingData<SinhTilingData>();
190+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
191+ OP_CHECK_IF(memset_s(tiling, sizeof(SinhTilingData), 0, sizeof(SinhTilingData)) != EOK,
192+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
193+ 
194+ // --- safer numeric types ---
195+ uint32_t typeLength = 0;
S
Ssu-yueming2025年12月16日

tiling部分建议提取成独立函数

likedislike
P
peihaobo
2025年12月17日 评论：
196+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
197+ if (typeLength == 0) {
198+ OP_LOGE(context, "typeLength is 0");
199+ return ge::GRAPH_FAILED;
200+ }
201+ SinhShapeInfo shapeInfo;
202+ ge::graphStatus ret = CalculateCoreBlockNums(ubSize,coreNum,typeLength, totalIdx,shapeInfo);
203+ if (ret != ge::GRAPH_SUCCESS) {
204+ return ret;
205+ }
206+ 
207+ // write back
208+ tiling->smallCoreDataNum = static_cast<int64_t>(shapeInfo.smallCoreDataNum);
209+ tiling->bigCoreDataNum = static_cast<int64_t>(shapeInfo.bigCoreDataNum);
210+ tiling->tileDataNum = static_cast<int64_t>(shapeInfo.tileDataNum);
211+ tiling->smallTailDataNum = static_cast<int64_t>(shapeInfo.smallTailDataNum);
212+ tiling->bigTailDataNum = static_cast<int64_t>(shapeInfo.bigTailDataNum);
213+ tiling->finalSmallTileNum = static_cast<int64_t>(shapeInfo.finalSmallTileNum);
214+ tiling->finalBigTileNum = static_cast<int64_t>(shapeInfo.finalBigTileNum);
215+ tiling->tailBlockNum = static_cast<int64_t>(shapeInfo.tailBlockNum);
216+ 
217+ context->SetBlockDim(shapeInfo.coreNum);
218+ 
219+ return ge::GRAPH_SUCCESS;
220+}
221+ 
222+static ge::graphStatus TilingParseForSinh([[maybe_unused]] gert::TilingParseContext* context)
223+{
224+ return ge::GRAPH_SUCCESS;
225+}
226+ 
227+// tiling注册入口.
228+IMPL_OP_OPTILING(Sinh).Tiling(SinhTilingFunc).TilingParse<SinhCompileInfo>(TilingParseForSinh);
229+} // namespace optiling
Aexperimental/math/sinh/op_kernel/sinh.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+ * - Pei Haobo<@xiaopei-1>
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 sinh.cpp
23+ * \brief
24+*/
25+ 
26+#include "sinh.h"
27+ 
28+template <uint32_t schMode>
29+__global__ __aicore__ void sinh(GM_ADDR x, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
30+{
31+ REGISTER_TILING_DEFAULT(SinhTilingData);
32+ GET_TILING_DATA_WITH_STRUCT(SinhTilingData, tilingData, tiling);
33+ NsSinh::Sinh<DTYPE_X> op;
34+ op.Init(x,z, &tilingData);
35+ op.Process();
36+}
Aexperimental/math/sinh/op_kernel/sinh.h+139-0
@@ -0,0 +1,139 @@
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+ * - Pei Haobo<@xiaopei-1>
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 sinh.h
23+ * \brief
24+*/
25+#ifndef SINH_H
26+#define SINH_H
27+ 
28+#include "kernel_operator.h"
29+#include "kernel_tiling/kernel_tiling.h"
30+#include "sinh_tiling_data.h"
31+#include "sinh_tiling_key.h"
32+ 
33+namespace NsSinh {
34+ 
35+using namespace AscendC;
36+ 
37+constexpr int32_t BUFFER_NUM = 2;
38+ 
39+template <typename T>
40+class Sinh {
41+public:
42+ __aicore__ inline Sinh(){};
43+ 
44+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, const SinhTilingData* 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+ TPipe pipe;
54+ TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX;
55+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueZ;
56+ GlobalTensor<T> inputGMX;
57+ GlobalTensor<T> outputGMZ;
58+ 
59+ uint32_t coreDataNum;
60+ uint32_t tileNum;
61+ uint32_t tileDataNum;
62+ uint32_t tailDataNum;
63+ uint32_t processDataNum;
64+};
65+ 
66+template <typename T>
67+__aicore__ inline void Sinh<T>::Init(GM_ADDR x, GM_ADDR z, const SinhTilingData* tilingData)
68+{
69+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
70+ uint32_t coreIdx = AscendC::GetBlockIdx();
71+ uint32_t globalBufferIndex = tilingData->bigCoreDataNum * AscendC::GetBlockIdx();
72+ this->tileDataNum = tilingData->tileDataNum;
73+ if (coreIdx < tilingData->tailBlockNum) {
74+ this->coreDataNum = tilingData->bigCoreDataNum;
75+ this->tileNum = tilingData->finalBigTileNum;
76+ this->tailDataNum = tilingData->bigTailDataNum;
77+ }
78+ else {
79+ this->coreDataNum = tilingData->smallCoreDataNum;
80+ this->tileNum = tilingData->finalSmallTileNum;
81+ this->tailDataNum = tilingData->smallTailDataNum;
82+ globalBufferIndex -= (tilingData->bigCoreDataNum - tilingData->smallCoreDataNum) * (AscendC::GetBlockIdx() - tilingData->tailBlockNum);
83+ }
84+ inputGMX.SetGlobalBuffer((__gm__ T*)x + globalBufferIndex, this->coreDataNum);
85+ outputGMZ.SetGlobalBuffer((__gm__ T*)z + globalBufferIndex, this->coreDataNum);
86+ pipe.InitBuffer(inputQueueX, BUFFER_NUM, this->tileDataNum * sizeof(T));
87+ pipe.InitBuffer(outputQueueZ, BUFFER_NUM, this->tileDataNum * sizeof(T));
88+ }
89+ 
90+template <typename T>
91+__aicore__ inline void Sinh<T>::CopyIn(int32_t progress)
92+{
93+ AscendC::LocalTensor<T> xLocal = inputQueueX.AllocTensor<T>();
94+ AscendC::DataCopy(xLocal, inputGMX[progress * this->tileDataNum], this->processDataNum);
95+ inputQueueX.EnQue(xLocal);
96+}
97+ 
98+template <typename T>
99+__aicore__ inline void Sinh<T>::CopyOut(int32_t progress)
100+{
101+ AscendC::LocalTensor<T> zLocal = outputQueueZ.DeQue<T>();
102+ AscendC::DataCopy(outputGMZ[progress * this->tileDataNum], zLocal, this->processDataNum);
103+ outputQueueZ.FreeTensor(zLocal);
104+}
105+ 
106+template <typename T>
107+__aicore__ inline void Sinh<T>::Compute(int32_t progress)
108+{
109+ AscendC::LocalTensor<T> xLocal = inputQueueX.DeQue<T>();
110+ AscendC::LocalTensor<T> zLocal = outputQueueZ.AllocTensor<T>();
111+ T scalar = 0.5;
112+ AscendC::Exp(xLocal, xLocal, this->processDataNum);
113+ PipeBarrier<PIPE_V>();
114+ AscendC::Reciprocal(zLocal, xLocal, this->processDataNum);
115+ PipeBarrier<PIPE_V>();
116+ AscendC::Sub(zLocal, xLocal, zLocal, this->processDataNum);
117+ PipeBarrier<PIPE_V>();
118+ AscendC::Muls(zLocal, zLocal, scalar, this->processDataNum);
119+ outputQueueZ.EnQue<T>(zLocal);
120+ inputQueueX.FreeTensor(xLocal);
121+}
122+ 
123+template <typename T>
124+__aicore__ inline void Sinh<T>::Process()
125+{
126+ int32_t loopCount = this->tileNum;
127+ this->processDataNum = this->tileDataNum;
128+ for (int32_t i = 0; i < loopCount; i++) {
129+ if (i == this->tileNum - 1) {
130+ this->processDataNum = this->tailDataNum;
131+ }
132+ CopyIn(i);
133+ Compute(i);
134+ CopyOut(i);
135+ }
136+}
137+ 
138+} // namespace NsSinh
139+#endif // Sinh_H
Aexperimental/math/sinh/op_kernel/sinh_tiling_data.h+38-0
@@ -0,0 +1,38 @@
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+ * - Pei Haobo<@xiaopei-1>
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 sinh_tiling_data.h
23+ * \brief tiling data struct
24+*/
25+#ifndef _SINH_TILING_DATA_H_
26+#define _SINH_TILING_DATA_H_
27+ 
28+struct SinhTilingData {
29+ int64_t smallCoreDataNum;
30+ int64_t bigCoreDataNum;
31+ int64_t finalBigTileNum;
32+ int64_t finalSmallTileNum;
33+ int64_t tileDataNum;
34+ int64_t smallTailDataNum;
35+ int64_t bigTailDataNum;
36+ int64_t tailBlockNum;
37+};
38+#endif
Aexperimental/math/sinh/op_kernel/sinh_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+ * - Pei Haobo<@xiaopei-1>
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 sinh_tiling_key.h
23+ * \brief sinh tiling key declare
24+*/
25+ 
26+#ifndef __SINH_TILING_KEY_H__
27+#define __SINH_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(Sinh,
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
Aexperimental/math/sinh/tests/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aexperimental/math/sinh/tests/ut/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Mscripts/kernel/binary_config/ascendc_config.json+4-1文件内容审核中，请稍后刷新重试