Star1411
1496
代码
代码
Issues100
Pull Requests327
Actions
项目讨论
Wiki
分析
Star1411
1496
已合并
feat:add dawse operate #3432
Almost_Miao创建于 6月18日
feat:add dawse operate #3432
已合并
Almost_Miao创建于 6月18日
从Almost_CANN/ops-mathmhf_ascendc_dawsn合入到cann/ops-mathmaster
共 15 个文件变更+1854-0
Mdocs/zh/op_list.md+10-0
@@ -576,6 +576,16 @@
576 <td>AI Core</td>576 <td>AI Core</td>
577 <td>逐元素比较两个相同shape和dtype的输入张量，统计差异超出容差范围的元素总个数。</td>577 <td>逐元素比较两个相同shape和dtype的输入张量，统计差异超出容差范围的元素总个数。</td>
578 </tr>578 </tr>
579+ <tr>
580+ <td>math</td>
581+ <td><a href="../../math/dawsn/README.md">dawsn</a></td>
582+ <td>√</td>
583+ <td>√</td>
584+ <td>×</td>
585+ <td>√</td>
586+ <td>AI Core</td>
587+ <td>计算Dawson积分 F(x) = exp(-x^2) * integral_0^x exp(t^2) dt。</td>
588+ </tr>
579 <tr>589 <tr>
580 <td>math</td>590 <td>math</td>
581 <td><a href="../../math/diag_part/README.md">diag_part</a></td>591 <td><a href="../../math/diag_part/README.md">diag_part</a></td>
Amath/dawsn/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+ 
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,
8+# INCLUDING 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+# Generated By CANNBot
12+ 
13+set(SUPPORT_COMPUTE_UNIT "ascend950")
14+set(SUPPORT_TILING_DIR "arch35")
15+add_all_modules_sources(OPTYPE dawsn ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Amath/dawsn/README.md+64-0
@@ -0,0 +1,64 @@
1+# Dawsn
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+|:------------------------------------------------| :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | √ |
12+| <term>Atlas 训练系列产品</term> | √ |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：计算Dawson积分 F(x) = exp(-x^2) * integral_0^x exp(t^2) dt。
17+ 
18+- 计算公式：
19+ 
20+$$F(x) = e^{-x^{2}} \int_{0}^{x} e^{t^{2}} \mathrm{d}t$$
21+ 
22+## 参数说明
23+ 
24+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
25+ <col style="width: 100px">
26+ <col style="width: 150px">
27+ <col style="width: 280px">
28+ <col style="width: 330px">
29+ <col style="width: 120px">
30+ </colgroup>
31+ <thead>
32+ <tr>
33+ <th>参数名</th>
34+ <th>输入/输出/属性</th>
35+ <th>描述</th>
36+ <th>数据类型</th>
37+ <th>数据格式</th>
38+ </tr></thead>
39+ <tbody>
40+ <tr>
41+ <td>x</td>
42+ <td>输入</td>
43+ <td>待进行Dawson积分计算的入参，公式中的x。</td>
44+ <td>BFLOAT16、FLOAT16、FLOAT</td>
45+ <td>ND</td>
46+ </tr>
47+ <tr>
48+ <td>y</td>
49+ <td>输出</td>
50+ <td>Dawson积分计算的出参，公式中的F(x)。</td>
51+ <td>BFLOAT16、FLOAT16、FLOAT</td>
52+ <td>ND</td>
53+ </tr>
54+ </tbody></table>
55+ 
56+## 约束说明
57+ 
58+- 不支持DOUBLE（FP64）数据类型。
59+ 
60+## 调用说明
61+ 
62+| 调用方式 | 调用样例 | 说明 |
63+|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
64+| 图模式调用 | [test_geir_dawsn](./examples/test_geir_dawsn.cpp) | 通过图模式调用Dawsn算子。 |
Amath/dawsn/examples/arch35/test_geir_dawsn.cpp+273-0
@@ -0,0 +1,273 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+#include <iostream>
14+#include <fstream>
15+#include <string.h>
16+#include <stdint.h>
17+#include <vector>
18+#include <string>
19+#include <map>
20+#include "assert.h"
21+ 
22+#include "graph.h"
23+#include "types.h"
24+#include "tensor.h"
25+#include "ge_error_codes.h"
26+#include "ge_api_types.h"
27+#include "ge_api.h"
28+#include "array_ops.h"
29+#include "ge_ir_build.h"
30+ 
31+#include "nn_other.h"
32+#include "../op_graph/dawsn_proto.h"
33+ 
34+#define FAILED -1
35+#define SUCCESS 0
36+ 
37+using namespace ge;
38+using std::map;
39+using std::string;
40+using std::vector;
41+ 
42+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
43+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
44+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
45+ TensorDesc placeholder##intputIndex##_desc = \
46+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
47+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
48+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
49+ Tensor tensor_placeholder##intputIndex; \
50+ ret = GenOnesData( \
51+ placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, placeholder##intputIndex##_desc, \
52+ intputDtype, 2); \
A
AAlmost_Miao6月25日

魔鬼数字需要修改。

likedislike
53+ if (ret != SUCCESS) { \
54+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
55+ return FAILED; \
56+ } \
57+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
58+ input.push_back(tensor_placeholder##intputIndex); \
59+ graph.AddOp(placeholder##intputIndex); \
60+ dawsn1.set_input_##intputName(placeholder##intputIndex); \
61+ inputs.push_back(placeholder##intputIndex);
62+ 
63+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
64+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
65+ dawsn1.update_output_desc_##outputName(outputName##outputIndex##_desc);
66+ 
67+string GetTime()
68+{
69+ time_t timep;
70+ time(&timep);
71+ char tmp[64];
72+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
73+ return tmp;
74+}
75+ 
76+uint32_t GetDataTypeSize(DataType dt)
77+{
78+ uint32_t dilation = 1;
79+ uint32_t twoByte = 2;
80+ uint32_t fourByte = 4;
81+ uint32_t eightByte = 8;
82+ 
83+ if (dt == ge::DT_FLOAT) {
84+ dilation = fourByte;
85+ } else if (dt == ge::DT_FLOAT16) {
86+ dilation = twoByte;
87+ } else if (dt == ge::DT_BF16) {
88+ dilation = twoByte;
89+ } else if (dt == ge::DT_INT16) {
90+ dilation = twoByte;
91+ } else if (dt == ge::DT_UINT16) {
92+ dilation = twoByte;
93+ } else if (dt == ge::DT_INT32) {
94+ dilation = fourByte;
95+ } else if (dt == ge::DT_UINT32) {
96+ dilation = fourByte;
97+ } else if (dt == ge::DT_INT64) {
98+ dilation = eightByte;
99+ } else if (dt == ge::DT_UINT64) {
100+ dilation = eightByte;
101+ } else if (dt == ge::DT_INT8) {
102+ dilation = 1;
103+ }
104+ return dilation;
105+}
106+ 
107+int32_t GenOnesData(
108+ vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type, int value)
109+{
110+ input_tensor_desc.SetRealDimCnt(shapes.size());
111+ size_t size = 1;
112+ for (uint32_t i = 0; i < shapes.size(); i++) {
113+ size *= shapes[i];
114+ }
115+ uint32_t data_len = size * GetDataTypeSize(data_type);
116+ uint8_t* pData = new (std::nothrow) uint8_t[data_len];
117+ if (pData == nullptr) {
118+ return FAILED;
119+ }
120+ FILE* fp = fopen("./test_input.bin", "rb");
121+ if (fp != nullptr) {
122+ size_t read = fread(pData, 1, data_len, fp);
123+ fclose(fp);
124+ if (read != data_len) {
125+ printf("[WARN] test_input.bin: read %zu bytes, expected %u\n", read, data_len);
126+ }
127+ } else {
128+ memset(pData, 0, data_len);
129+ printf("[WARN] test_input.bin not found, using zeros\n");
130+ }
131+ input_tensor = Tensor(input_tensor_desc, pData, data_len);
132+ delete[] pData;
133+ pData = nullptr;
134+ return SUCCESS;
135+}
136+ 
137+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
138+{
139+ FILE* fp = fopen(bin_file.c_str(), "w");
140+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
141+ fclose(fp);
142+ return SUCCESS;
143+}
144+ 
145+int CreateOppInGraph(
146+ DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, std::vector<Operator>& outputs,
147+ Graph& graph)
148+{
149+ Status ret = SUCCESS;
150+ auto dawsn1 = op::Dawsn("dawsn1");
151+ std::vector<int64_t> xShape = {2, 4};
152+ ADD_INPUT(1, x, inDtype, xShape);
153+ 
154+ ADD_OUTPUT(1, y, inDtype, xShape);
155+ 
156+ outputs.push_back(dawsn1);
157+ return SUCCESS;
158+}
159+ 
160+int main(int argc, char* argv[])
161+{
162+ const char* graph_name = "tc_ge_irrun_test";
163+ Graph graph(graph_name);
164+ std::vector<ge::Tensor> input;
165+ 
166+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
167+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
168+ Status ret = ge::GEInitialize(global_options);
169+ if (ret != SUCCESS) {
170+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
171+ return FAILED;
172+ }
173+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
174+ 
175+ std::vector<Operator> inputs{};
176+ std::vector<Operator> outputs{};
177+ 
178+ if (argc < 2) {
179+ printf("%s - ERROR - [XIR]: Missing required argument. Usage: %s <dtype_index>\n", GetTime().c_str(), argv[0]);
180+ ge::GEFinalize();
181+ return FAILED;
182+ }
183+ std::cout << argv[1] << std::endl;
184+ char* endptr;
185+ 
186+ DataType inDtype = DT_FLOAT;
187+ 
188+ std::cout << inDtype << std::endl;
189+ 
190+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
191+ if (ret != SUCCESS) {
192+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
193+ return FAILED;
194+ }
195+ 
196+ if (!inputs.empty() && !outputs.empty()) {
197+ graph.SetInputs(inputs).SetOutputs(outputs);
198+ }
199+ 
200+ std::map<AscendString, AscendString> build_options = {
201+ 
202+ };
203+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
204+ ge::Session* session = new Session(build_options);
205+ 
206+ if (session == nullptr) {
207+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
208+ return FAILED;
209+ }
210+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
211+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
212+ 
213+ std::map<AscendString, AscendString> graph_options = {
214+ 
215+ };
216+ uint32_t graph_id = 0;
217+ ret = session->AddGraph(graph_id, graph, graph_options);
218+ 
219+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
220+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
221+ std::string file_path = "./dump";
222+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
223+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
224+ std::vector<ge::Tensor> output;
225+ ret = session->RunGraph(graph_id, input, output);
226+ if (ret != SUCCESS) {
227+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
228+ delete session;
229+ GEFinalize();
230+ return FAILED;
231+ }
232+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
233+ 
234+ int input_num = input.size();
235+ for (int i = 0; i < input_num; i++) {
236+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
237+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
238+ uint8_t* input_data_i = input[i].GetData();
239+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
240+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
241+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
242+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
243+ }
244+ 
245+ int output_num = output.size();
246+ for (int i = 0; i < output_num; i++) {
247+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
248+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
249+ uint8_t* output_data_i = output[i].GetData();
250+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
251+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
252+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
253+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
254+ }
255+ 
256+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
257+ std::string error_str(error_msg.GetString());
258+ std::cout << "Error message: " << error_str << std::endl;
259+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
260+ std::string warning_str(warning_msg.GetString());
261+ std::cout << "Warning message: " << warning_str << std::endl;
262+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
263+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
264+ delete session;
265+ session = nullptr;
266+ ret = ge::GEFinalize();
267+ if (ret != SUCCESS) {
268+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
269+ return FAILED;
270+ }
271+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
272+ return SUCCESS;
273+}
Amath/dawsn/examples/test_geir_dawsn.cpp+273-0
@@ -0,0 +1,273 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+#include <iostream>
14+#include <fstream>
15+#include <string.h>
16+#include <stdint.h>
17+#include <vector>
18+#include <string>
19+#include <map>
20+#include "assert.h"
21+ 
22+#include "graph.h"
23+#include "types.h"
24+#include "tensor.h"
25+#include "ge_error_codes.h"
26+#include "ge_api_types.h"
27+#include "ge_api.h"
28+#include "array_ops.h"
29+#include "ge_ir_build.h"
30+ 
31+#include "nn_other.h"
32+#include "../op_graph/dawsn_proto.h"
33+ 
34+#define FAILED -1
35+#define SUCCESS 0
36+ 
37+using namespace ge;
38+using std::map;
39+using std::string;
40+using std::vector;
41+ 
42+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
43+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
44+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
45+ TensorDesc placeholder##intputIndex##_desc = \
46+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
47+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
48+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
49+ Tensor tensor_placeholder##intputIndex; \
50+ ret = GenOnesData( \
51+ placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, placeholder##intputIndex##_desc, \
52+ intputDtype, 2); \
53+ if (ret != SUCCESS) { \
54+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
55+ return FAILED; \
56+ } \
57+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
58+ input.push_back(tensor_placeholder##intputIndex); \
59+ graph.AddOp(placeholder##intputIndex); \
60+ dawsn1.set_input_##intputName(placeholder##intputIndex); \
61+ inputs.push_back(placeholder##intputIndex);
62+ 
63+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
64+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
65+ dawsn1.update_output_desc_##outputName(outputName##outputIndex##_desc);
66+ 
67+string GetTime()
68+{
69+ time_t timep;
70+ time(&timep);
71+ char tmp[64];
72+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
73+ return tmp;
74+}
75+ 
76+uint32_t GetDataTypeSize(DataType dt)
77+{
78+ uint32_t dilation = 1;
79+ uint32_t twoByte = 2;
80+ uint32_t fourByte = 4;
81+ uint32_t eightByte = 8;
82+ 
83+ if (dt == ge::DT_FLOAT) {
84+ dilation = fourByte;
85+ } else if (dt == ge::DT_FLOAT16) {
86+ dilation = twoByte;
87+ } else if (dt == ge::DT_BF16) {
88+ dilation = twoByte;
89+ } else if (dt == ge::DT_INT16) {
90+ dilation = twoByte;
91+ } else if (dt == ge::DT_UINT16) {
92+ dilation = twoByte;
93+ } else if (dt == ge::DT_INT32) {
94+ dilation = fourByte;
95+ } else if (dt == ge::DT_UINT32) {
96+ dilation = fourByte;
97+ } else if (dt == ge::DT_INT64) {
98+ dilation = eightByte;
99+ } else if (dt == ge::DT_UINT64) {
100+ dilation = eightByte;
101+ } else if (dt == ge::DT_INT8) {
102+ dilation = 1;
103+ }
104+ return dilation;
105+}
106+ 
107+int32_t GenOnesData(
108+ vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type, int value)
109+{
110+ input_tensor_desc.SetRealDimCnt(shapes.size());
111+ size_t size = 1;
112+ for (uint32_t i = 0; i < shapes.size(); i++) {
113+ size *= shapes[i];
114+ }
115+ uint32_t data_len = size * GetDataTypeSize(data_type);
116+ uint8_t* pData = new (std::nothrow) uint8_t[data_len];
117+ if (pData == nullptr) {
118+ return FAILED;
119+ }
120+ FILE* fp = fopen("./test_input.bin", "rb");
121+ if (fp != nullptr) {
122+ size_t read = fread(pData, 1, data_len, fp);
123+ fclose(fp);
124+ if (read != data_len) {
125+ printf("[WARN] test_input.bin: read %zu bytes, expected %u\n", read, data_len);
126+ }
127+ } else {
128+ memset(pData, 0, data_len);
129+ printf("[WARN] test_input.bin not found, using zeros\n");
130+ }
131+ input_tensor = Tensor(input_tensor_desc, pData, data_len);
132+ delete[] pData;
133+ pData = nullptr;
134+ return SUCCESS;
135+}
136+ 
137+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
138+{
139+ FILE* fp = fopen(bin_file.c_str(), "w");
140+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
141+ fclose(fp);
142+ return SUCCESS;
143+}
144+ 
145+int CreateOppInGraph(
146+ DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, std::vector<Operator>& outputs,
147+ Graph& graph)
148+{
149+ Status ret = SUCCESS;
150+ auto dawsn1 = op::Dawsn("dawsn1");
151+ std::vector<int64_t> xShape = {2, 4};
152+ ADD_INPUT(1, x, inDtype, xShape);
153+ 
154+ ADD_OUTPUT(1, y, inDtype, xShape);
155+ 
156+ outputs.push_back(dawsn1);
157+ return SUCCESS;
158+}
159+ 
160+int main(int argc, char* argv[])
161+{
162+ const char* graph_name = "tc_ge_irrun_test";
163+ Graph graph(graph_name);
164+ std::vector<ge::Tensor> input;
165+ 
166+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
167+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
168+ Status ret = ge::GEInitialize(global_options);
169+ if (ret != SUCCESS) {
170+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
171+ return FAILED;
172+ }
173+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
174+ 
175+ std::vector<Operator> inputs{};
176+ std::vector<Operator> outputs{};
177+ 
178+ if (argc < 2) {
179+ printf("%s - ERROR - [XIR]: Missing required argument. Usage: %s <dtype_index>\n", GetTime().c_str(), argv[0]);
180+ ge::GEFinalize();
181+ return FAILED;
182+ }
183+ std::cout << argv[1] << std::endl;
184+ char* endptr;
185+ 
186+ DataType inDtype = DT_FLOAT;
187+ 
188+ std::cout << inDtype << std::endl;
189+ 
190+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
191+ if (ret != SUCCESS) {
192+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
193+ return FAILED;
194+ }
195+ 
196+ if (!inputs.empty() && !outputs.empty()) {
197+ graph.SetInputs(inputs).SetOutputs(outputs);
198+ }
199+ 
200+ std::map<AscendString, AscendString> build_options = {
201+ 
202+ };
203+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
204+ ge::Session* session = new Session(build_options);
205+ 
206+ if (session == nullptr) {
207+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
208+ return FAILED;
209+ }
210+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
211+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
212+ 
213+ std::map<AscendString, AscendString> graph_options = {
214+ 
215+ };
216+ uint32_t graph_id = 0;
217+ ret = session->AddGraph(graph_id, graph, graph_options);
218+ 
219+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
220+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
221+ std::string file_path = "./dump";
222+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
223+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
224+ std::vector<ge::Tensor> output;
225+ ret = session->RunGraph(graph_id, input, output);
226+ if (ret != SUCCESS) {
227+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
228+ delete session;
229+ GEFinalize();
230+ return FAILED;
231+ }
232+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
233+ 
234+ int input_num = input.size();
235+ for (int i = 0; i < input_num; i++) {
236+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
237+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
238+ uint8_t* input_data_i = input[i].GetData();
239+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
240+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
241+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
242+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
243+ }
244+ 
245+ int output_num = output.size();
246+ for (int i = 0; i < output_num; i++) {
247+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
248+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
249+ uint8_t* output_data_i = output[i].GetData();
250+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
251+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
252+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
253+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
254+ }
255+ 
256+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
257+ std::string error_str(error_msg.GetString());
258+ std::cout << "Error message: " << error_str << std::endl;
259+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
260+ std::string warning_str(warning_msg.GetString());
261+ std::cout << "Warning message: " << warning_str << std::endl;
262+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
263+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
264+ delete session;
265+ session = nullptr;
266+ ret = ge::GEFinalize();
267+ if (ret != SUCCESS) {
268+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
269+ return FAILED;
270+ }
271+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
272+ return SUCCESS;
273+}
Amath/dawsn/op_graph/dawsn_proto.h+36-0
@@ -0,0 +1,36 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+#ifndef OPS_OP_PROTO_INC_DAWSN_H_
14+#define OPS_OP_PROTO_INC_DAWSN_H_
15+ 
16+#include "graph/operator_reg.h"
17+#include "graph/types.h"
18+ 
19+namespace ge {
20+ 
21+/**
22+ * @brief Computes Dawson's integral: F(x) = exp(-x^2) * integral_0^x exp(t^2) dt
23+ * @par Inputs:
24+ * @li x: A ND Tensor. Must be one of: float16, float32, bfloat16, double.
25+ *
26+ * @par Outputs:
27+ * @li y: A ND Tensor. Same type and shape as x.
28+ */
29+REG_OP(Dawsn)
30+ .INPUT(x, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT, DT_DOUBLE}))
31+ .OUTPUT(y, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT, DT_DOUBLE}))
32+ .OP_END_FACTORY_REG(Dawsn)
33+ 
34+} // namespace ge
35+ 
36+#endif // OPS_OP_PROTO_INC_DAWSN_H_
Amath/dawsn/op_host/arch35/dawsn_tiling_arch35.cpp+218-0
@@ -0,0 +1,218 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+/*!
14+ * \file dawsn_tiling_arch35.cpp
15+ * \brief Dawsn 算子 Host 侧 Tiling 计算（arch35 / DAV_3510）
16+ */
17+ 
18+#include "register/op_def_registry.h"
19+#include "op_common/log/log.h"
20+#include "op_common/op_host/util/math_util.h"
21+#include "op_common/op_host/util/platform_util.h"
22+#include "../../op_kernel/arch35/dawsn_tiling_data.h"
23+#include "../../op_kernel/arch35/dawsn_tiling_key.h"
24+ 
25+namespace optiling {
26+ 
27+using Ops::Base::CeilDiv;
28+using Ops::Base::CeilAlign;
29+using Ops::Base::FloorDiv;
30+using Ops::Base::FloorAlign;
31+using Ops::Base::GetUbBlockSize;
32+ 
33+constexpr uint32_t WS_SYS_SIZE = 0U;
34+constexpr size_t WORKSPACE_NUM = 1;
35+constexpr int64_t MIN_TILING_BITS = 32768; // 4KB per core minimum, in bits
36+constexpr int64_t ELEM_ALIGN_FACTOR = 512; // multi-core element alignment
37+constexpr int64_t ALIGN_256_BYTES = 256; // UB alignment in bytes
38+constexpr int64_t BITS_PER_BYTE = 8; // bits per byte
39+ 
40+// Buffer planning constants for GetBufferDivisor
41+constexpr int64_t FP32_REG_BUF_COUNT = 7; // FP32 path: 6 RegTensor buffers + x/y buffers
42+constexpr int64_t FP32_MASK_BYTES = 1; // mask buffer size per element
43+constexpr int64_t FP16_ELEM_BYTES = 2; // half/bf16 element size in bytes
44+constexpr int64_t FP16_ORIG_BUF_COUNT = 2; // FP16/BF16: 2 original dtype buffers
45+constexpr int64_t FP16_FP32_BUF_COUNT = 5; // FP16/BF16: 5 FP32 intermediate buffers
46+ 
47+static const gert::Shape g_vec_1_shape = {1};
48+ 
49+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape) {
50+ if (in_shape.GetDimNum() == 0) {
51+ return g_vec_1_shape;
52+ }
53+ return in_shape;
54+}
55+ 
56+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)
57+{
58+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
59+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
60+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
61+ *coreNum = ascendcPlatform.GetCoreNumAiv();
62+ OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
63+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize);
64+ OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
65+ return ge::GRAPH_SUCCESS;
66+}
67+ 
68+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t* totalIdx, ge::DataType* dataType)
69+{
70+ auto inputX = context->GetInputShape(0);
71+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
72+ auto inputShapeX = EnsureNotScalar(inputX->GetStorageShape());
73+ 
74+ *totalIdx = inputShapeX.GetShapeSize();
75+ 
76+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
77+ auto inputDesc = context->GetInputDesc(0);
78+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
79+ *dataType = inputDesc->GetDataType();
80+ OP_CHECK_IF(supportedDtype.count(*dataType) == 0, OP_LOGE(context, "invalid dtype"), return ge::GRAPH_FAILED);
81+ return ge::GRAPH_SUCCESS;
82+}
83+ 
84+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
85+{
86+ size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM);
87+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
88+ currentWorkspace[0] = WS_SYS_SIZE;
89+ return ge::GRAPH_SUCCESS;
90+}
91+ 
92+static int64_t GetElemBytes(ge::DataType dtype)
93+{
94+ switch (dtype) {
95+ case ge::DT_FLOAT16: return 2;
96+ case ge::DT_BF16: return 2;
97+ case ge::DT_FLOAT: return 4;
98+ default: return 4;
99+ }
100+}
101+ 
102+static int64_t GetBufferDivisor(ge::DataType dtype)
103+{
104+ // bufferDivisor accounts for all UB buffers needed per element
105+ if (dtype == ge::DT_FLOAT) {
106+ // FP32 path: FP32_REG_BUF_COUNT FP32 buffers + FP32_MASK_BYTES mask
107+ // Plus x_buf and y_buf
108+ return FP32_REG_BUF_COUNT * sizeof(float) + FP32_MASK_BYTES;
109+ } else if (dtype == ge::DT_FLOAT16 || dtype == ge::DT_BF16) {
110+ // FP16/BF16 path: FP16_ORIG_BUF_COUNT original dtype buffers + FP16_FP32_BUF_COUNT FP32 buffers
111+ return FP16_ORIG_BUF_COUNT * FP16_ELEM_BYTES + FP16_FP32_BUF_COUNT * sizeof(float);
112+ } else {
113+ // Fallback: use FP32-like buffer planning
114+ return FP32_REG_BUF_COUNT * sizeof(float) + FP32_MASK_BYTES;
115+ }
116+}
117+ 
118+static void ComputeTilingParams(int64_t totalIdx, int64_t availableCoreNum, uint64_t ubSize,
119+ ge::DataType dataType, DawsnTilingData* tiling)
120+{
121+ int64_t elemBytes = GetElemBytes(dataType);
122+ int64_t minDtypeBits = elemBytes * BITS_PER_BYTE;
123+ 
124+ // Multi-core splitting
125+ int64_t coreNum = (totalIdx * minDtypeBits + MIN_TILING_BITS - 1) / MIN_TILING_BITS;
126+ coreNum = std::min(coreNum, availableCoreNum);
127+ int64_t blockFormer = ((totalIdx + coreNum - 1) / coreNum + ELEM_ALIGN_FACTOR - 1) /
128+ ELEM_ALIGN_FACTOR * ELEM_ALIGN_FACTOR;
129+ int64_t blockNum = (totalIdx + blockFormer - 1) / blockFormer;
130+ 
131+ // UB splitting
132+ int64_t bufferDivisor = GetBufferDivisor(dataType);
133+ int64_t maxElemNum = static_cast<int64_t>(ubSize) / bufferDivisor;
134+ int64_t alignFactor = ALIGN_256_BYTES / elemBytes;
135+ int64_t ubFormer = (maxElemNum / alignFactor) * alignFactor;
136+ if (ubFormer <= 0) {
137+ ubFormer = alignFactor;
138+ }
139+ 
140+ // Loop counts
141+ int64_t ubLoopOfFormerBlock = (blockFormer + ubFormer - 1) / ubFormer;
142+ int64_t ubTailOfFormerBlock = blockFormer - (ubLoopOfFormerBlock - 1) * ubFormer;
143+ int64_t blockTail = totalIdx - (blockNum - 1) * blockFormer;
144+ int64_t ubLoopOfTailBlock = (blockTail + ubFormer - 1) / ubFormer;
145+ int64_t ubTailOfTailBlock = blockTail - (ubLoopOfTailBlock - 1) * ubFormer;
146+ 
147+ // Fill tiling data
148+ tiling->dim0 = totalIdx;
149+ tiling->coreNum = coreNum;
150+ tiling->blockFormer = blockFormer;
151+ tiling->blockNum = blockNum;
152+ tiling->ubFormer = ubFormer;
153+ tiling->ubLoopOfFormerBlock = ubLoopOfFormerBlock;
154+ tiling->ubTailOfFormerBlock = ubTailOfFormerBlock;
155+ tiling->ubLoopOfTailBlock = ubLoopOfTailBlock;
156+ tiling->ubTailOfTailBlock = ubTailOfTailBlock;
157+}
158+ 
159+static ge::graphStatus DawsnTilingFunc(gert::TilingContext* context)
160+{
161+ OP_LOGI(context->GetNodeName(), "Enter DawsnTilingFunc");
162+ // 1. Get platform info
163+ uint64_t ubSize;
164+ int64_t availableCoreNum;
165+ OP_CHECK_IF(
166+ GetPlatformInfo(context, &ubSize, &availableCoreNum) != ge::GRAPH_SUCCESS,
167+ OP_LOGE(context, "GetPlatformInfo error"),
168+ return ge::GRAPH_FAILED);
169+ 
170+ // 2. Get shape and dtype info
171+ int64_t totalIdx;
172+ ge::DataType dataType = ge::DT_FLOAT;
173+ OP_CHECK_IF(
174+ GetShapeAttrsInfo(context, &totalIdx, &dataType) != ge::GRAPH_SUCCESS,
175+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
176+ 
177+ // 3. Set workspace
178+ OP_CHECK_IF(
179+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
180+ OP_LOGE(context, "GetWorkspaceSize error"),
181+ return ge::GRAPH_FAILED);
182+ 
183+ // 4. Get tiling data pointer
184+ DawsnTilingData* tiling = context->GetTilingData<DawsnTilingData>();
185+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
186+ OP_CHECK_IF(
187+ memset_s(tiling, sizeof(DawsnTilingData), 0, sizeof(DawsnTilingData)) != EOK,
188+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
189+ 
190+ // Empty tensor check
191+ if (totalIdx == 0) {
192+ context->SetBlockDim(1);
193+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(dataType));
194+ return ge::GRAPH_SUCCESS;
195+ }
196+ 
197+ // 5-8. Compute tiling parameters
198+ ComputeTilingParams(totalIdx, availableCoreNum, ubSize, dataType, tiling);
199+ 
200+ context->SetBlockDim(tiling->blockNum);
201+ 
202+ // 9. Set TilingKey (template parameter selection)
203+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
204+ ASCENDC_TPL_SEL_PARAM(context, dTypeX);
205+ 
206+ return ge::GRAPH_SUCCESS;
207+}
208+ 
209+static ge::graphStatus TilingParseForDawsn([[maybe_unused]] gert::TilingParseContext* context)
210+{
211+ return ge::GRAPH_SUCCESS;
212+}
213+ 
214+struct DawsnCompileInfo {};
215+ 
216+IMPL_OP_OPTILING(Dawsn).Tiling(DawsnTilingFunc).TilingParse<DawsnCompileInfo>(TilingParseForDawsn);
217+ 
218+} // namespace optiling
Amath/dawsn/op_host/dawsn_def.cpp+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+/*!
14+ * \file dawsn_def.cpp
15+ * \brief dawsn def
16+ */
17+#include "register/op_def_registry.h"
18+ 
19+namespace ops {
20+class Dawsn : public OpDef {
21+ public:
22+ explicit Dawsn(const char* name) : OpDef(name)
23+ {
24+ this->Input("x")
25+ .ParamType(REQUIRED)
26+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
27+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
28+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
29+ this->Output("y")
30+ .ParamType(REQUIRED)
31+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
32+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
33+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
34+ OpAICoreConfig aiCoreConfig;
35+ aiCoreConfig.DynamicCompileStaticFlag(true)
36+ .DynamicFormatFlag(false)
37+ .DynamicRankSupportFlag(true)
38+ .DynamicShapeSupportFlag(true)
39+ .NeedCheckSupportFlag(false)
40+ .PrecisionReduceFlag(true)
41+ .ExtendCfgInfo("opFile.value", "dawsn_apt");
42+ this->AICore().AddConfig("ascend950", aiCoreConfig);
43+ }
44+};
45+ 
46+OP_ADD(Dawsn);
47+} // namespace ops
Amath/dawsn/op_host/dawsn_infershape.cpp+41-0
@@ -0,0 +1,41 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+/*!
14+ * \file dawsn_infershape.cpp
15+ * \brief Dawsn 算子形状推导（输出 shape = 输入 shape）
16+ */
17+ 
18+#include "register/op_impl_registry.h"
19+#include "exe_graph/runtime/infer_shape_context.h"
20+#include "op_common/log/log.h"
21+ 
22+using namespace ge;
23+ 
24+namespace ops {
25+ 
26+static ge::graphStatus InferShape4Dawsn(gert::InferShapeContext* context)
27+{
28+ const gert::Shape* input_shape = context->GetInputShape(0);
29+ OP_CHECK_NULL_WITH_CONTEXT(context, input_shape);
30+ 
31+ gert::Shape* output_shape = context->GetOutputShape(0);
32+ OP_CHECK_NULL_WITH_CONTEXT(context, output_shape);
33+ 
34+ *output_shape = *input_shape;
35+ 
36+ return ge::GRAPH_SUCCESS;
37+}
38+ 
39+IMPL_OP_INFERSHAPE(Dawsn).InferShape(InferShape4Dawsn);
40+ 
41+} // namespace ops
Amath/dawsn/op_kernel/arch35/dawsn_kernel.h+498-0
@@ -0,0 +1,498 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+#ifndef DAWSN_KERNEL_H
14+#define DAWSN_KERNEL_H
15+ 
16+#include "kernel_operator.h"
17+#include "kernel_tiling/kernel_tiling.h"
18+#include "dawsn_tiling_data.h"
19+#include "dawsn_tiling_key.h"
20+ 
21+namespace NsDawsn {
22+ 
23+using namespace AscendC;
24+ 
25+// ============================================================================
26+// Cephes dawsn.c rational polynomial coefficients for Dawson's integral
27+// Source: scipy/special/cephes/dawsn.c (Cephes Math Library Release 2.1)
28+// ============================================================================
29+ 
30+// Segment 1: F(x) = x * polevl(x^2, AN, 9) / polevl(x^2, AD, 10) for |x| < 3.25
31+static constexpr float S1_AN[] = {
32+ 1.13681498971755967716E-11f, 8.49262267667473820670E-10f,
33+ 1.94434204175553062245E-08f, 9.53151741254484376186E-07f,
34+ 3.07828309874913215953E-06f, 3.52513368520288765742E-04f,
35+ -8.50149846724410876604E-04f, 4.22618223005546619908E-02f,
36+ -9.17480371773452402229E-02f, 1.00000000000000000000E+00f
37+};
38+static constexpr float S1_AD[] = {
39+ 2.40372073066762607321E-11f, 1.48864681368493399468E-09f,
40+ 5.21265281010541632356E-08f, 1.27258478273186967251E-06f,
41+ 2.32490249820789501151E-05f, 3.25524741826057934131E-04f,
42+ 3.48805814657162597398E-03f, 2.79448531198828974131E-02f,
43+ 1.58874241960120565764E-01f, 5.74918629489320376358E-01f,
44+ 1.00000000000000000000E+00f
45+};
46+ 
47+// Segment 2: 0.5*(1/x + (1/x^2)*polevl(1/x^2,BN,10)/(p1evl(1/x^2,BD,10)*x))
48+// for 3.25 <= |x| < 6.25
49+static constexpr float S2_BN[] = {
50+ 5.08955156417900900223E-01f, -2.44754418142697854899E-01f,
51+ 9.41512335303534425579E-02f, -2.18711255142039033561E-02f,
52+ 3.66207612329569201851E-03f, -4.23209114460388778275E-04f,
53+ 3.59641304793896659001E-05f, -2.14640351719968955485E-06f,
54+ 9.10010780076391391280E-08f, -2.40274520828250976051E-09f,
55+ 3.59233385440928385474E-11f
56+};
57+static constexpr float S2_BD[] = {
58+ -6.31839869873368153996E-01f, 2.36706788228248699024E-01f,
59+ -5.31806367003223243173E-02f, 8.48041718586295401039E-03f,
60+ -9.47996768486665357389E-04f, 7.81025592944552368886E-05f,
61+ -4.55875153252442628812E-06f, 1.89100358111421858017E-07f,
62+ -4.91324691331920598557E-09f, 7.18466403235734505828E-11f
63+};
64+ 
65+// Segment 3: 0.5*(1/x + (1/x^2)*polevl(1/x^2,CN,4)/(p1evl(1/x^2,CD,5)*x))
66+// for 6.25 <= |x| <= 1e9
67+static constexpr float S3_CN[] = {
68+ -5.90592860534773222980E-01f, 6.29235242724368815104E-01f,
69+ -1.72858975380388135701E-01f, 1.64837047825189646211E-02f,
70+ -4.86827613020462685988E-04f
71+};
72+static constexpr float S3_CD[] = {
73+ -2.69820057197544915795E+00f, 1.73270799045947843098E+00f,
74+ -3.93708582281939500902E-01f, 3.44278924041233405640E-02f,
75+ -9.73655226040941201328E-04f
76+};
77+ 
78+// NaN/Inf handling constants
79+// F32_NAN: IEEE 754 quiet NaN bit pattern for float (used with int32 Duplicate + MERGING)
80+static constexpr int32_t F32_NAN = 0x7fc00000;
81+// FLT_MAX: largest finite float; abs(x) > FLT_MAX means abs(x) is Inf
82+static constexpr float FLT_MAX_VAL = 3.4028235e+38f;
83+ 
84+// Segment boundary constants (from Cephes dawsn.c)
85+static constexpr float SEG1_UPPER_BOUND = 3.25f; // |x| < 3.25 → Seg1
86+static constexpr float SEG2_UPPER_BOUND = 6.25f; // 3.25 <= |x| < 6.25 → Seg2
87+static constexpr float SEG3_UPPER_BOUND = 1.0e9f; // 6.25 <= |x| <= 1e9 → Seg3, |x| > 1e9 → Seg4
88+static constexpr float DAWSON_HALF = 0.5f; // coefficient in Seg2/Seg3/Seg4 formula
89+ 
90+// VF register size in bytes
91+static constexpr uint32_t VF_REG_SIZE_BYTES = 256;
92+ 
93+// CastTrait constants (per cast-api.md §4.2 type conversion matrix)
94+// half/bf16 -> float (widening): RegLayout=ZERO, RoundMode=UNKNOWN
95+static constexpr AscendC::Reg::CastTrait WIDEN_F16_TO_F32 = {
96+ AscendC::Reg::RegLayout::ZERO, AscendC::Reg::SatMode::UNKNOWN,
97+ AscendC::Reg::MaskMergeMode::ZEROING, AscendC::RoundMode::UNKNOWN};
98+static constexpr AscendC::Reg::CastTrait WIDEN_BF16_TO_F32 = {
99+ AscendC::Reg::RegLayout::ZERO, AscendC::Reg::SatMode::UNKNOWN,
100+ AscendC::Reg::MaskMergeMode::ZEROING, AscendC::RoundMode::UNKNOWN};
101+// float -> half/bf16 (narrowing): RegLayout=ZERO, SatMode=NO_SAT, RoundMode=CAST_ROUND
102+static constexpr AscendC::Reg::CastTrait NARROW_F32_TO_F16 = {
103+ AscendC::Reg::RegLayout::ZERO, AscendC::Reg::SatMode::NO_SAT,
104+ AscendC::Reg::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_ROUND};
105+static constexpr AscendC::Reg::CastTrait NARROW_F32_TO_BF16 = {
106+ AscendC::Reg::RegLayout::ZERO, AscendC::Reg::SatMode::NO_SAT,
107+ AscendC::Reg::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_ROUND};
108+ 
109+// ============================================================================
110+// VF function declaration
111+// ============================================================================
112+template <typename T>
113+__simd_vf__ inline void DawsnVF(
114+ __ubuf__ T* dstAddr, __ubuf__ T* srcAddr,
115+ uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes);
116+ 
117+// ============================================================================
118+// Kernel class
119+// ============================================================================
120+template <typename T>
121+class DawsnKernel {
122+public:
123+ __aicore__ inline DawsnKernel() {}
124+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const DawsnTilingData* tilingData);
125+ __aicore__ inline void Process();
126+private:
127+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
128+ __aicore__ inline void Compute(int64_t currentNum);
129+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
130+ 
131+ TPipe pipe;
132+ TQue<QuePosition::VECIN, 2> inputQueueX;
133+ TQue<QuePosition::VECOUT, 2> outputQueueY;
134+ GlobalTensor<T> inputGMX;
135+ GlobalTensor<T> outputGMY;
136+ int64_t blockLength_ = 0;
137+ int64_t ubLength_ = 0;
138+};
139+ 
140+template <typename T>
141+__aicore__ inline void DawsnKernel<T>::Init(GM_ADDR x, GM_ADDR y, const DawsnTilingData* tilingData)
142+{
143+ int64_t blockIdx = AscendC::GetBlockIdx();
144+ int64_t remainderLength = tilingData->dim0 - tilingData->blockFormer * blockIdx;
145+ blockLength_ = (remainderLength > tilingData->blockFormer) ? tilingData->blockFormer : remainderLength;
146+ ubLength_ = tilingData->ubFormer;
147+ inputGMX.SetGlobalBuffer((__gm__ T*)x + tilingData->blockFormer * blockIdx, blockLength_);
148+ outputGMY.SetGlobalBuffer((__gm__ T*)y + tilingData->blockFormer * blockIdx, blockLength_);
149+ pipe.InitBuffer(inputQueueX, 2, ubLength_ * sizeof(T));
150+ pipe.InitBuffer(outputQueueY, 2, ubLength_ * sizeof(T));
151+}
152+ 
153+template <typename T>
154+__aicore__ inline void DawsnKernel<T>::CopyIn(int64_t progress, int64_t currentNum)
155+{
156+ AscendC::LocalTensor<T> xLocal = inputQueueX.template AllocTensor<T>();
157+ AscendC::DataCopyParams copyParams;
158+ copyParams.blockCount = 1;
159+ copyParams.blockLen = currentNum * sizeof(T);
160+ copyParams.srcStride = 0;
161+ copyParams.dstStride = 0;
162+ AscendC::DataCopyPad(xLocal, inputGMX[progress * ubLength_], copyParams, {false, 0, 0, 0});
163+ inputQueueX.EnQue(xLocal);
164+}
165+ 
166+template <typename T>
167+__aicore__ inline void DawsnKernel<T>::CopyOut(int64_t progress, int64_t currentNum)
168+{
169+ AscendC::LocalTensor<T> yLocal = outputQueueY.template DeQue<T>();
170+ AscendC::DataCopyParams copyParams;
171+ copyParams.blockCount = 1;
172+ copyParams.blockLen = currentNum * sizeof(T);
173+ copyParams.srcStride = 0;
174+ copyParams.dstStride = 0;
175+ AscendC::DataCopyPad(outputGMY[progress * ubLength_], yLocal, copyParams);
176+ outputQueueY.FreeTensor(yLocal);
177+}
178+ 
179+template <typename T>
180+__aicore__ inline void DawsnKernel<T>::Compute(int64_t currentNum)
181+{
182+ AscendC::LocalTensor<T> xLocal = inputQueueX.template DeQue<T>();
183+ AscendC::LocalTensor<T> yLocal = outputQueueY.template AllocTensor<T>();
184+ constexpr uint32_t oneRepeatSize = AscendC::GetVecLen() / sizeof(T);
185+ uint16_t repeatTimes = AscendC::CeilDivision(static_cast<uint32_t>(currentNum), oneRepeatSize);
186+ asc_vf_call<DawsnVF<T>>(
187+ (__ubuf__ T*)yLocal.GetPhyAddr(),
188+ (__ubuf__ T*)xLocal.GetPhyAddr(),
189+ static_cast<uint32_t>(currentNum), oneRepeatSize, repeatTimes);
190+ outputQueueY.template EnQue<T>(yLocal);
191+ inputQueueX.FreeTensor(xLocal);
192+}
193+ 
194+template <typename T>
195+__aicore__ inline void DawsnKernel<T>::Process()
196+{
197+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
198+ if (loopCount == 1) {
199+ CopyIn(0, blockLength_);
200+ Compute(blockLength_);
201+ CopyOut(0, blockLength_);
202+ return;
203+ }
204+ CopyIn(0, ubLength_);
205+ int64_t currentNum = ubLength_;
206+ for (int64_t i = 0; i < loopCount - 1; i++) {
207+ int64_t nextNum = (i + 1 == loopCount - 1) ? (blockLength_ - ubLength_ * (i + 1)) : ubLength_;
208+ Compute(currentNum);
209+ CopyIn(i + 1, nextNum);
210+ CopyOut(i, currentNum);
211+ currentNum = nextNum;
212+ }
213+ Compute(currentNum);
214+ CopyOut(loopCount - 1, currentNum);
215+}
216+ 
217+// ============================================================================
218+// Helper: Horner evaluation of S1_AN (degree 9, 10 coeffs) with variable u
219+// ============================================================================
220+#define DAWSN_HORNER_AN(u, pReg, tmpReg, mask) \
221+ AscendC::Reg::Duplicate(pReg, S1_AN[0], mask); \
222+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[1], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
223+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[2], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
224+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[3], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
225+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[4], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
226+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[5], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
227+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[6], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
228+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[7], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
229+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[8], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
230+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AN[9], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask);
231+ 
232+// Helper: Horner evaluation of S1_AD (degree 10, 11 coeffs) with variable u
233+#define DAWSN_HORNER_AD(u, qReg, tmpReg, mask) \
234+ AscendC::Reg::Duplicate(qReg, S1_AD[0], mask); \
235+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[1], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
236+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[2], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
237+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[3], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
238+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[4], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
239+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[5], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
240+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[6], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
241+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[7], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
242+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[8], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
243+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[9], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
244+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S1_AD[10], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask);
245+ 
246+// Helper: polevl(1/x^2, BN, 10) - degree 10, 11 coeffs
247+#define DAWSN_HORNER_BN(u, pReg, tmpReg, mask) \
248+ AscendC::Reg::Duplicate(pReg, S2_BN[0], mask); \
249+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[1], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
250+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[2], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
251+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[3], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
252+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[4], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
253+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[5], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
254+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[6], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
255+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[7], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
256+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[8], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
257+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[9], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
258+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BN[10], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask);
259+ 
260+// Helper: p1evl(1/x^2, BD, 10) - degree 10, leading 1 implicit, 10 stored coeffs
261+#define DAWSN_HORNER_BD(u, qReg, tmpReg, mask) \
262+ AscendC::Reg::Duplicate(qReg, S2_BD[0], mask); AscendC::Reg::Add(qReg, qReg, u, mask); \
263+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[1], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
264+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[2], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
265+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[3], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
266+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[4], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
267+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[5], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
268+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[6], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
269+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[7], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
270+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[8], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
271+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S2_BD[9], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask);
272+ 
273+// Helper: polevl(1/x^2, CN, 4) - degree 4, 5 coeffs
274+#define DAWSN_HORNER_CN(u, pReg, tmpReg, mask) \
275+ AscendC::Reg::Duplicate(pReg, S3_CN[0], mask); \
276+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S3_CN[1], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
277+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S3_CN[2], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
278+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S3_CN[3], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask); \
279+ AscendC::Reg::Mul(pReg, pReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S3_CN[4], mask); AscendC::Reg::Add(pReg, pReg, tmpReg, mask);
280+ 
281+// Helper: p1evl(1/x^2, CD, 5) - degree 5, leading 1 implicit, 5 stored coeffs
282+#define DAWSN_HORNER_CD(u, qReg, tmpReg, mask) \
283+ AscendC::Reg::Duplicate(qReg, S3_CD[0], mask); AscendC::Reg::Add(qReg, qReg, u, mask); \
284+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S3_CD[1], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
285+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S3_CD[2], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
286+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S3_CD[3], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask); \
287+ AscendC::Reg::Mul(qReg, qReg, u, mask); AscendC::Reg::Duplicate(tmpReg, S3_CD[4], mask); AscendC::Reg::Add(qReg, qReg, tmpReg, mask);
288+ 
289+// Helper: Segment 2/3 formula: 0.5*(1/|x| + inv_x2 * P / (Q * |x|))
290+// Inputs: invXReg=1/|x|, invX2Reg=1/x^2, pReg=P(inv_x2), qReg=Q(inv_x2), absReg=|x|
291+// Output: segResult
292+#define DAWSN_SEG23_FORMULA(invXReg, invX2Reg, pReg, qReg, absReg, segResult, tmpReg, tmp2Reg, mask) \
293+ AscendC::Reg::Mul(tmpReg, qReg, absReg, mask); \
294+ AscendC::Reg::Div(tmpReg, pReg, tmpReg, mask); \
295+ AscendC::Reg::Mul(segResult, invX2Reg, tmpReg, mask); \
296+ AscendC::Reg::Add(segResult, invXReg, segResult, mask); \
297+ AscendC::Reg::Duplicate(tmpReg, DAWSON_HALF, mask); \
298+ AscendC::Reg::Mul(segResult, segResult, tmpReg, mask);
299+ 
300+// ============================================================================
301+// Common compute logic for all dtypes (shared by float/half/bf16 VF functions)
302+// Parameters: xReg=input, absReg, xSqReg, pReg, qReg, tmpReg, resultReg,
303+// invXReg, segResult, invX2Reg, mask, maskSeg1-4, maskZero,
304+// maskNeg, maskNan, maskInf, maskSpecial
305+// ============================================================================
306+#define DAWSN_COMPUTE(xReg, absReg, xSqReg, pReg, qReg, tmpReg, resultReg, \
307+ invXReg, segResult, invX2Reg, mask, \
308+ maskSeg1, maskSeg2, maskSeg3, maskSeg4, \
309+ maskZero, maskNeg, maskNan, maskInf, maskSpecial) \
310+ AscendC::Reg::Abs(absReg, xReg, mask); \
311+ AscendC::Reg::Compare<float, AscendC::CMPMODE::NE>(maskNan, absReg, absReg, mask); \
312+ AscendC::Reg::Duplicate(tmpReg, FLT_MAX_VAL, mask); \
313+ AscendC::Reg::Compare<float, AscendC::CMPMODE::GT>(maskInf, absReg, tmpReg, mask); \
314+ AscendC::Reg::MaskOr(maskSpecial, maskNan, maskInf, mask); \
315+ AscendC::Reg::Duplicate(tmpReg, 0.0f, mask); \
316+ AscendC::Reg::Compare<float, AscendC::CMPMODE::EQ>(maskZero, absReg, tmpReg, mask); \
317+ AscendC::Reg::Duplicate(tmpReg, 1.0f, mask); \
318+ AscendC::Reg::Select<float>(absReg, tmpReg, absReg, maskZero); \
319+ AscendC::Reg::Select<float>(absReg, tmpReg, absReg, maskSpecial); \
320+ AscendC::Reg::Mul(xSqReg, absReg, absReg, mask); \
321+ AscendC::Reg::Duplicate(tmpReg, SEG1_UPPER_BOUND, mask); \
322+ AscendC::Reg::Compare<float, AscendC::CMPMODE::LT>(maskSeg1, absReg, tmpReg, mask); \
323+ AscendC::Reg::Duplicate(tmpReg, SEG2_UPPER_BOUND, mask); \
324+ AscendC::Reg::Compare<float, AscendC::CMPMODE::LT>(maskSeg2, absReg, tmpReg, mask); \
325+ AscendC::Reg::Duplicate(tmpReg, SEG3_UPPER_BOUND, mask); \
326+ AscendC::Reg::Compare<float, AscendC::CMPMODE::LE>(maskSeg3, absReg, tmpReg, mask); \
327+ AscendC::Reg::Compare<float, AscendC::CMPMODE::GT>(maskSeg4, absReg, tmpReg, mask); \
328+ AscendC::Reg::Duplicate(tmpReg, 1.0f, mask); \
329+ AscendC::Reg::Div(invXReg, tmpReg, absReg, mask); \
330+ AscendC::Reg::Mul(invX2Reg, invXReg, invXReg, mask); \
331+ AscendC::Reg::Duplicate(tmpReg, 0.0f, mask); \
332+ AscendC::Reg::Select<float>(invX2Reg, tmpReg, invX2Reg, maskSeg1); \
333+ DAWSN_HORNER_AN(xSqReg, pReg, tmpReg, mask) \
334+ DAWSN_HORNER_AD(xSqReg, qReg, tmpReg, mask) \
335+ AscendC::Reg::Mul(segResult, absReg, pReg, mask); \
336+ AscendC::Reg::Div(segResult, segResult, qReg, mask); \
337+ AscendC::Reg::Duplicate(resultReg, 0.0f, mask); \
338+ AscendC::Reg::Select<float>(resultReg, segResult, resultReg, maskSeg1); \
339+ DAWSN_HORNER_BN(invX2Reg, pReg, tmpReg, mask) \
340+ DAWSN_HORNER_BD(invX2Reg, qReg, tmpReg, mask) \
341+ DAWSN_SEG23_FORMULA(invXReg, invX2Reg, pReg, qReg, absReg, segResult, tmpReg, resultReg, mask) \
342+ AscendC::Reg::Select<float>(resultReg, segResult, resultReg, maskSeg2); \
343+ DAWSN_HORNER_CN(invX2Reg, pReg, tmpReg, mask) \
344+ DAWSN_HORNER_CD(invX2Reg, qReg, tmpReg, mask) \
345+ DAWSN_SEG23_FORMULA(invXReg, invX2Reg, pReg, qReg, absReg, segResult, tmpReg, resultReg, mask) \
346+ AscendC::Reg::Select<float>(resultReg, segResult, resultReg, maskSeg3); \
347+ AscendC::Reg::Duplicate(tmpReg, DAWSON_HALF, mask); \
348+ AscendC::Reg::Div(segResult, tmpReg, absReg, mask); \
349+ AscendC::Reg::Select<float>(resultReg, segResult, resultReg, maskSeg4); \
350+ DAWSN_HORNER_AN(xSqReg, pReg, tmpReg, mask) \
351+ DAWSN_HORNER_AD(xSqReg, qReg, tmpReg, mask) \
352+ AscendC::Reg::Mul(segResult, absReg, pReg, mask); \
353+ AscendC::Reg::Div(segResult, segResult, qReg, mask); \
354+ AscendC::Reg::Select<float>(resultReg, segResult, resultReg, maskSeg1); \
355+ AscendC::Reg::Duplicate(tmpReg, 0.0f, mask); \
356+ AscendC::Reg::Select<float>(resultReg, tmpReg, resultReg, maskZero); \
357+ AscendC::Reg::Duplicate(tmpReg, 0.0f, mask); \
358+ AscendC::Reg::Compare<float, AscendC::CMPMODE::LT>(maskNeg, xReg, tmpReg, mask); \
359+ AscendC::Reg::Neg(segResult, resultReg, mask); \
360+ AscendC::Reg::Select<float>(resultReg, segResult, resultReg, maskNeg); \
361+ AscendC::Reg::Duplicate(tmpReg, 0.0f, mask); \
362+ AscendC::Reg::Select<float>(resultReg, tmpReg, resultReg, maskInf); \
363+ AscendC::Reg::Duplicate<int32_t, AscendC::Reg::MaskMergeMode::MERGING>( \
364+ (AscendC::Reg::RegTensor<int32_t>&)resultReg, F32_NAN, maskNan);
365+ 
366+// ============================================================================
367+// VF function: float - 4-segment Dawson integral
368+// ============================================================================
369+template <>
370+__simd_vf__ inline void DawsnVF<float>(
371+ __ubuf__ float* dstAddr, __ubuf__ float* srcAddr,
372+ uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes)
373+{
374+ AscendC::Reg::RegTensor<float> xReg, absReg, xSqReg, pReg, qReg, tmpReg;
375+ AscendC::Reg::RegTensor<float> resultReg, invXReg, segResult, invX2Reg;
376+ AscendC::Reg::MaskReg mask, maskSeg1, maskSeg2, maskSeg3, maskSeg4, maskZero, maskNeg;
377+ AscendC::Reg::MaskReg maskNan, maskInf, maskSpecial;
378+ 
379+ for (uint16_t i = 0; i < repeatTimes; ++i) {
380+ mask = AscendC::Reg::UpdateMask<float>(count);
381+ AscendC::Reg::LoadAlign(xReg, srcAddr + i * oneRepeatSize);
382+ DAWSN_COMPUTE(xReg, absReg, xSqReg, pReg, qReg, tmpReg, resultReg,
383+ invXReg, segResult, invX2Reg, mask,
384+ maskSeg1, maskSeg2, maskSeg3, maskSeg4,
385+ maskZero, maskNeg, maskNan, maskInf, maskSpecial)
386+ AscendC::Reg::StoreAlign(dstAddr + i * oneRepeatSize, resultReg, mask);
387+ }
388+}
389+ 
390+// ============================================================================
391+// VF function: half (FP16 -> Cast -> FP32 -> compute -> Cast -> FP16)
392+//
393+// Cast chain widest type = float (4B) → 64 elements per VF register.
394+// Per cast-vf.md §2:
395+// - LoadAlign must use DIST_UNPACK_B16 (half 2B → float 4B slot)
396+// - StoreAlign must use DIST_PACK_B32 (float 4B slot → half 2B)
397+// - Mask must use UpdateMask<float> (widest type determines mask format)
398+// - oneRepeatSize must be based on float (64 elements, not 128)
399+// ============================================================================
400+template <>
401+__simd_vf__ inline void DawsnVF<half>(
402+ __ubuf__ half* dstAddr, __ubuf__ half* srcAddr,
403+ uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes)
404+{
405+ AscendC::Reg::RegTensor<half> xHalfReg, yHalfReg;
406+ AscendC::Reg::RegTensor<float> xFloatReg;
407+ AscendC::Reg::RegTensor<float> absReg, xSqReg, pReg, qReg, tmpReg;
408+ AscendC::Reg::RegTensor<float> resultReg, invXReg, segResult, invX2Reg;
409+ AscendC::Reg::MaskReg mask;
410+ AscendC::Reg::MaskReg maskSeg1, maskSeg2, maskSeg3, maskSeg4, maskZero, maskNeg;
411+ AscendC::Reg::MaskReg maskNan, maskInf, maskSpecial;
412+ 
413+ // Cast chain widest type is float → 64 elements per register
414+ constexpr uint32_t fOneRepeatSize = VF_REG_SIZE_BYTES / sizeof(float); // 64
415+ uint16_t fRepeatTimes = AscendC::CeilDivision(count, fOneRepeatSize);
416+ 
417+ for (uint16_t i = 0; i < fRepeatTimes; ++i) {
418+ // Use float mask (widest type in Cast chain determines mask format)
419+ mask = AscendC::Reg::UpdateMask<float>(count);
420+ 
421+ // Load half with UNPACK_B16: each 16-bit half → 32-bit slot
422+ AscendC::Reg::LoadAlign<half, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(
423+ xHalfReg, srcAddr + i * fOneRepeatSize);
424+ 
425+ // Widen: half (in 32-bit slots) → float
426+ AscendC::Reg::Cast<float, half, WIDEN_F16_TO_F32>(xFloatReg, xHalfReg, mask);
427+ 
428+ // Compute in FP32
429+ DAWSN_COMPUTE(xFloatReg, absReg, xSqReg, pReg, qReg, tmpReg, resultReg,
430+ invXReg, segResult, invX2Reg, mask,
431+ maskSeg1, maskSeg2, maskSeg3, maskSeg4,
432+ maskZero, maskNeg, maskNan, maskInf, maskSpecial)
433+ 
434+ // Narrow: float → half (in 32-bit slots)
435+ AscendC::Reg::Cast<half, float, NARROW_F32_TO_F16>(yHalfReg, resultReg, mask);
436+ 
437+ // Store half with PACK_B32: each 32-bit slot → 16-bit half
438+ AscendC::Reg::StoreAlign<half, AscendC::Reg::StoreDist::DIST_PACK_B32>(
439+ dstAddr + i * fOneRepeatSize, yHalfReg, mask);
440+ }
441+}
442+ 
443+// ============================================================================
444+// VF function: bfloat16 (BF16 -> Cast -> FP32 -> compute -> Cast -> BF16)
445+//
446+// Cast chain widest type = float (4B) → 64 elements per VF register.
447+// Per cast-vf.md §2:
448+// - LoadAlign must use DIST_UNPACK_B16 (bf16 2B → float 4B slot)
449+// - StoreAlign must use DIST_PACK_B32 (float 4B slot → bf16 2B)
450+// - Mask must use UpdateMask<float> (widest type determines mask format)
451+// - oneRepeatSize must be based on float (64 elements, not 128)
452+// ============================================================================
453+template <>
454+__simd_vf__ inline void DawsnVF<bfloat16_t>(
455+ __ubuf__ bfloat16_t* dstAddr, __ubuf__ bfloat16_t* srcAddr,
456+ uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes)
457+{
458+ AscendC::Reg::RegTensor<bfloat16_t> xBfReg, yBfReg;
459+ AscendC::Reg::RegTensor<float> xFloatReg;
460+ AscendC::Reg::RegTensor<float> absReg, xSqReg, pReg, qReg, tmpReg;
461+ AscendC::Reg::RegTensor<float> resultReg, invXReg, segResult, invX2Reg;
462+ AscendC::Reg::MaskReg mask;
463+ AscendC::Reg::MaskReg maskSeg1, maskSeg2, maskSeg3, maskSeg4, maskZero, maskNeg;
464+ AscendC::Reg::MaskReg maskNan, maskInf, maskSpecial;
465+ 
466+ // Cast chain widest type is float → 64 elements per register
467+ constexpr uint32_t fOneRepeatSize = VF_REG_SIZE_BYTES / sizeof(float); // 64
468+ uint16_t fRepeatTimes = AscendC::CeilDivision(count, fOneRepeatSize);
469+ 
470+ for (uint16_t i = 0; i < fRepeatTimes; ++i) {
471+ // Use float mask (widest type in Cast chain determines mask format)
472+ mask = AscendC::Reg::UpdateMask<float>(count);
473+ 
474+ // Load bf16 with UNPACK_B16: each 16-bit bf16 → 32-bit slot
475+ AscendC::Reg::LoadAlign<bfloat16_t, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(
476+ xBfReg, srcAddr + i * fOneRepeatSize);
477+ 
478+ // Widen: bf16 (in 32-bit slots) → float
479+ AscendC::Reg::Cast<float, bfloat16_t, WIDEN_BF16_TO_F32>(xFloatReg, xBfReg, mask);
480+ 
481+ // Compute in FP32
482+ DAWSN_COMPUTE(xFloatReg, absReg, xSqReg, pReg, qReg, tmpReg, resultReg,
483+ invXReg, segResult, invX2Reg, mask,
484+ maskSeg1, maskSeg2, maskSeg3, maskSeg4,
485+ maskZero, maskNeg, maskNan, maskInf, maskSpecial)
486+ 
487+ // Narrow: float → bf16 (in 32-bit slots)
488+ AscendC::Reg::Cast<bfloat16_t, float, NARROW_F32_TO_BF16>(yBfReg, resultReg, mask);
489+ 
490+ // Store bf16 with PACK_B32: each 32-bit slot → 16-bit bf16
491+ AscendC::Reg::StoreAlign<bfloat16_t, AscendC::Reg::StoreDist::DIST_PACK_B32>(
492+ dstAddr + i * fOneRepeatSize, yBfReg, mask);
493+ }
494+}
495+ 
496+} // namespace NsDawsn
497+ 
498+#endif // DAWSN_KERNEL_H
Amath/dawsn/op_kernel/arch35/dawsn_tiling_data.h+28-0
@@ -0,0 +1,28 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+#ifndef DAWSN_TILING_DATA_H
A
AAlmost_Miao6月18日

CANNBot生成的代码需要增加对应的说明

likedislike
A
Almost_Miao
6月25日 评论：
14+#define DAWSN_TILING_DATA_H
15+ 
16+struct DawsnTilingData {
17+ int64_t dim0; // 元素总数量（所有维度展平）
18+ int64_t coreNum; // 实际使用核数
19+ int64_t blockFormer; // 每个核的基础元素数（512 对齐）
20+ int64_t blockNum; // 虚拟 block 数量
21+ int64_t ubFormer; // 每个 UB 块的基础元素数（256B 对齐）
22+ int64_t ubLoopOfFormerBlock; // 首块 UB 循环次数
23+ int64_t ubTailOfFormerBlock; // 首块尾部元素数
24+ int64_t ubLoopOfTailBlock; // 尾块 UB 循环次数
25+ int64_t ubTailOfTailBlock; // 尾块尾部元素数
26+};
27+ 
28+#endif // DAWSN_TILING_DATA_H
Amath/dawsn/op_kernel/arch35/dawsn_tiling_key.h+37-0
@@ -0,0 +1,37 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+#ifndef DAWSN_TILING_KEY_H
14+#define DAWSN_TILING_KEY_H
15+ 
16+#include "ascendc/host_api/tiling/template_argument.h"
17+ 
18+// Dawsn 仅有 1 个 TilingKey，dtype 差异通过模板参数在 VF 内部处理
19+// 模板参数：D_T_X = 输入数据类型
20+ASCENDC_TPL_ARGS_DECL(Dawsn,
21+ ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16,
22+ ASCENDC_TPL_INPUT(0))
23+);
24+ 
25+ASCENDC_TPL_SEL(
26+ ASCENDC_TPL_ARGS_SEL(
27+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT)
28+ ),
29+ ASCENDC_TPL_ARGS_SEL(
30+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16)
31+ ),
32+ ASCENDC_TPL_ARGS_SEL(
33+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_BF16)
34+ ),
35+);
36+ 
37+#endif // DAWSN_TILING_KEY_H
Amath/dawsn/op_kernel/dawsn_apt.cpp+28-0
@@ -0,0 +1,28 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+/*!
14+ * \file dawsn_apt.cpp
15+ * \brief Dawsn 算子 kernel 入口（arch35 架构，DAV_3510）
16+ */
17+ 
18+#include "arch35/dawsn_kernel.h"
19+ 
20+template <typename D_T_X>
21+__global__ __aicore__ void dawsn(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
22+{
23+ REGISTER_TILING_DEFAULT(DawsnTilingData);
24+ GET_TILING_DATA_WITH_STRUCT(DawsnTilingData, tilingData, tiling);
25+ NsDawsn::DawsnKernel<D_T_X> op;
26+ op.Init(x, y, &tilingData);
27+ op.Process();
28+}
Amath/dawsn/tests/ut/op_host/arch35/test_dawsn_tiling_arch35.cpp+155-0
@@ -0,0 +1,155 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+#include <iostream>
14+#include <gtest/gtest.h>
15+#include "tiling_context_faker.h"
16+#include "tiling_case_executor.h"
17+#include "../../../../op_kernel/arch35/dawsn_tiling_data.h"
18+ 
19+namespace DawsnUT {
20+using namespace std;
21+using namespace ge;
22+using namespace gert;
23+static const std::string OP_NAME = "Dawsn";
24+ 
25+struct DawsnCompileInfo {};
26+ 
27+struct DawsnTestParam {
28+ std::string caseName;
29+ std::initializer_list<int64_t> xShape;
30+ ge::DataType xDtype;
31+ ge::Format xFormat;
32+ std::initializer_list<int64_t> yShape;
33+ ge::DataType yDtype;
34+ ge::Format yFormat;
35+ ge::graphStatus status;
36+ std::vector<size_t> expectWorkspaces;
37+ uint64_t maxAIVNum;
38+ uint64_t ubSize;
39+ uint64_t tilingDataMaxSize;
40+};
41+ 
42+// Dawsn tiling test cases
43+// Tiling key mapping (from dawsn_tiling_key.h ASCENDC_TPL_SEL):
44+// FP32 -> key = DT_FLOAT(0), FP16 -> key = DT_FLOAT16(1), BF16 -> key = DT_BF16(27)
45+// UB size for DAV_3510 (Ascend950) = 253952 (248KB)
46+// Core num for Ascend950 = 48
47+static DawsnTestParam testCases[] = {
48+ // === 正常场景 - 不同数据类型 ===
49+ {"dawsn_tiling_0_small_2d_fp32",
50+ {2, 3}, ge::DT_FLOAT, ge::FORMAT_ND,
51+ {2, 3}, ge::DT_FLOAT, ge::FORMAT_ND,
52+ ge::GRAPH_SUCCESS, {0}, 48, 253952, 4096},
53+ 
54+ {"dawsn_tiling_1_1d_fp16",
55+ {1024}, ge::DT_FLOAT16, ge::FORMAT_ND,
56+ {1024}, ge::DT_FLOAT16, ge::FORMAT_ND,
57+ ge::GRAPH_SUCCESS, {0}, 48, 253952, 4096},
58+ 
59+ {"dawsn_tiling_2_4d_bf16",
60+ {2, 3, 4, 5}, ge::DT_BF16, ge::FORMAT_ND,
61+ {2, 3, 4, 5}, ge::DT_BF16, ge::FORMAT_ND,
62+ ge::GRAPH_SUCCESS, {0}, 48, 253952, 4096},
63+ 
64+ // === 正常场景 - 不同形状 ===
65+ {"dawsn_tiling_3_large_1d_fp32",
66+ {1000000}, ge::DT_FLOAT, ge::FORMAT_ND,
67+ {1000000}, ge::DT_FLOAT, ge::FORMAT_ND,
68+ ge::GRAPH_SUCCESS, {0}, 48, 253952, 4096},
69+ 
70+ {"dawsn_tiling_4_scalar_fp32",
71+ {1}, ge::DT_FLOAT, ge::FORMAT_ND,
72+ {1}, ge::DT_FLOAT, ge::FORMAT_ND,
73+ ge::GRAPH_SUCCESS, {0}, 48, 253952, 4096},
74+ 
75+ {"dawsn_tiling_5_8d_fp16",
76+ {1, 1, 2, 2, 2, 2, 2, 2}, ge::DT_FLOAT16, ge::FORMAT_ND,
77+ {1, 1, 2, 2, 2, 2, 2, 2}, ge::DT_FLOAT16, ge::FORMAT_ND,
78+ ge::GRAPH_SUCCESS, {0}, 48, 253952, 4096},
79+ 
80+ {"dawsn_tiling_6_3d_fp32_medium",
81+ {32, 32, 32}, ge::DT_FLOAT, ge::FORMAT_ND,
82+ {32, 32, 32}, ge::DT_FLOAT, ge::FORMAT_ND,
83+ ge::GRAPH_SUCCESS, {0}, 48, 253952, 4096},
84+ 
85+ // === 边界场景 - 空Tensor ===
86+ {"dawsn_tiling_7_empty_tensor",
87+ {0, 3}, ge::DT_FLOAT, ge::FORMAT_ND,
88+ {0, 3}, ge::DT_FLOAT, ge::FORMAT_ND,
89+ ge::GRAPH_SUCCESS, {0}, 48, 253952, 4096},
90+ 
91+ // === 异常场景 - 不支持的数据类型 ===
92+ {"dawsn_tiling_8_unsupport_double",
93+ {2, 3}, ge::DT_DOUBLE, ge::FORMAT_ND,
94+ {2, 3}, ge::DT_DOUBLE, ge::FORMAT_ND,
95+ ge::GRAPH_FAILED, {0}, 48, 253952, 4096},
96+ 
97+ {"dawsn_tiling_9_unsupport_int32",
98+ {2, 3}, ge::DT_INT32, ge::FORMAT_ND,
99+ {2, 3}, ge::DT_INT32, ge::FORMAT_ND,
100+ ge::GRAPH_FAILED, {0}, 48, 253952, 4096},
101+};
102+ 
103+class DawsnTilingTest : public testing::TestWithParam<DawsnTestParam> {
104+protected:
105+ static void SetUpTestCase()
106+ {
107+ std::cout << "DawsnTilingTest SetUp." << std::endl;
108+ }
109+ 
110+ static void TearDownTestCase()
111+ {
112+ std::cout << "DawsnTilingTest TearDown." << std::endl;
113+ }
114+};
115+ 
116+static DawsnCompileInfo compileInfo;
117+ 
118+static uint64_t GetExpectedTilingKey(ge::DataType dtype)
119+{
120+ return static_cast<uint64_t>(dtype);
121+}
122+ 
123+static void TestOneParamCase(const DawsnTestParam &param)
124+{
125+ std::cout << "[TEST_CASE] " << param.caseName << std::endl;
126+ gert::StorageShape xShape = {param.xShape, param.xShape};
127+ gert::StorageShape yShape = {param.yShape, param.yShape};
128+ std::vector<gert::TilingContextPara::TensorDescription> inputTensorDesc_(
129+ {{xShape, param.xDtype, param.xFormat}});
130+ std::vector<gert::TilingContextPara::TensorDescription> outputTensorDesc_(
131+ {{yShape, param.yDtype, param.yFormat}});
132+ gert::TilingContextPara tilingContextPara(
133+ OP_NAME,
134+ inputTensorDesc_,
135+ outputTensorDesc_,
136+ &compileInfo,
137+ param.maxAIVNum,
138+ param.ubSize,
139+ param.tilingDataMaxSize);
140+ uint64_t expectedKey = GetExpectedTilingKey(param.xDtype);
141+ ExecuteTestCase(tilingContextPara, param.status, expectedKey, EMPTY_EXPECT_TILING_DATA, param.expectWorkspaces);
142+}
143+ 
144+TEST_P(DawsnTilingTest, tiling_test)
145+{
146+ const DawsnTestParam &param = GetParam();
147+ TestOneParamCase(param);
148+}
149+ 
150+INSTANTIATE_TEST_SUITE_P(
151+ DawsnTilingTests,
152+ DawsnTilingTest,
153+ testing::ValuesIn(testCases));
154+ 
155+} // namespace DawsnUT
Amath/dawsn/tests/ut/op_host/test_dawsn_infershape.cpp+131-0
@@ -0,0 +1,131 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ *
10+ * Generated By CANNBot
11+ */
12+ 
13+#include <gtest/gtest.h>
14+#include <iostream>
15+#include "infershape_context_faker.h"
16+#include "infershape_case_executor.h"
17+ 
18+class DawsnInfershape : public testing::Test {
19+protected:
20+ static void SetUpTestCase()
21+ {
22+ std::cout << "DawsnInfershape SetUp" << std::endl;
23+ }
24+ 
25+ static void TearDownTestCase()
26+ {
27+ std::cout << "DawsnInfershape TearDown" << std::endl;
28+ }
29+};
30+ 
31+// ========== 正常场景测试 - 不同数据类型 ==========
32+ 
33+TEST_F(DawsnInfershape, infershape_2d_fp32)
34+{
35+ gert::StorageShape shape = {{2, 3}, {2, 3}};
36+ gert::InfershapeContextPara::TensorDescription x(shape, ge::DT_FLOAT, ge::FORMAT_ND);
37+ gert::InfershapeContextPara::TensorDescription y(shape, ge::DT_FLOAT, ge::FORMAT_ND);
38+ gert::InfershapeContextPara infershapeContextPara("Dawsn", {x}, {y});
39+ std::vector<std::vector<int64_t>> expectOutputShape = {
40+ {2, 3},
41+ };
42+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
43+}
44+ 
45+TEST_F(DawsnInfershape, infershape_1d_fp16)
46+{
47+ gert::StorageShape shape = {{128}, {128}};
48+ gert::InfershapeContextPara::TensorDescription x(shape, ge::DT_FLOAT16, ge::FORMAT_ND);
49+ gert::InfershapeContextPara::TensorDescription y(shape, ge::DT_FLOAT16, ge::FORMAT_ND);
50+ gert::InfershapeContextPara infershapeContextPara("Dawsn", {x}, {y});
51+ std::vector<std::vector<int64_t>> expectOutputShape = {
52+ {128},
53+ };
54+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
55+}
56+ 
57+TEST_F(DawsnInfershape, infershape_4d_bf16)
58+{
59+ gert::StorageShape shape = {{2, 3, 4, 5}, {2, 3, 4, 5}};
60+ gert::InfershapeContextPara::TensorDescription x(shape, ge::DT_BF16, ge::FORMAT_ND);
61+ gert::InfershapeContextPara::TensorDescription y(shape, ge::DT_BF16, ge::FORMAT_ND);
62+ gert::InfershapeContextPara infershapeContextPara("Dawsn", {x}, {y});
63+ std::vector<std::vector<int64_t>> expectOutputShape = {
64+ {2, 3, 4, 5},
65+ };
66+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
67+}
68+ 
69+// ========== 正常场景测试 - 不同维度 ==========
70+ 
71+TEST_F(DawsnInfershape, infershape_scalar_fp32)
72+{
73+ gert::StorageShape shape = {{}, {}};
74+ gert::InfershapeContextPara::TensorDescription x(shape, ge::DT_FLOAT, ge::FORMAT_ND);
75+ gert::InfershapeContextPara::TensorDescription y(shape, ge::DT_FLOAT, ge::FORMAT_ND);
76+ gert::InfershapeContextPara infershapeContextPara("Dawsn", {x}, {y});
77+ std::vector<std::vector<int64_t>> expectOutputShape = {
78+ {},
79+ };
80+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
81+}
82+ 
83+TEST_F(DawsnInfershape, infershape_3d_fp32)
84+{
85+ gert::StorageShape shape = {{4, 8, 16}, {4, 8, 16}};
86+ gert::InfershapeContextPara::TensorDescription x(shape, ge::DT_FLOAT, ge::FORMAT_ND);
87+ gert::InfershapeContextPara::TensorDescription y(shape, ge::DT_FLOAT, ge::FORMAT_ND);
88+ gert::InfershapeContextPara infershapeContextPara("Dawsn", {x}, {y});
89+ std::vector<std::vector<int64_t>> expectOutputShape = {
90+ {4, 8, 16},
91+ };
92+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
93+}
94+ 
95+TEST_F(DawsnInfershape, infershape_8d_fp32)
96+{
97+ gert::StorageShape shape = {{1, 1, 2, 2, 2, 2, 2, 2}, {1, 1, 2, 2, 2, 2, 2, 2}};
98+ gert::InfershapeContextPara::TensorDescription x(shape, ge::DT_FLOAT, ge::FORMAT_ND);
99+ gert::InfershapeContextPara::TensorDescription y(shape, ge::DT_FLOAT, ge::FORMAT_ND);
100+ gert::InfershapeContextPara infershapeContextPara("Dawsn", {x}, {y});
101+ std::vector<std::vector<int64_t>> expectOutputShape = {
102+ {1, 1, 2, 2, 2, 2, 2, 2},
103+ };
104+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
105+}
106+ 
107+// ========== 边界场景测试 - 空Tensor ==========
108+ 
109+TEST_F(DawsnInfershape, infershape_empty_tensor)
110+{
111+ gert::StorageShape shape = {{0, 3}, {0, 3}};
112+ gert::InfershapeContextPara::TensorDescription x(shape, ge::DT_FLOAT, ge::FORMAT_ND);
113+ gert::InfershapeContextPara::TensorDescription y(shape, ge::DT_FLOAT, ge::FORMAT_ND);
114+ gert::InfershapeContextPara infershapeContextPara("Dawsn", {x}, {y});
115+ std::vector<std::vector<int64_t>> expectOutputShape = {
116+ {0, 3},
117+ };
118+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
119+}
120+ 
121+TEST_F(DawsnInfershape, infershape_single_element)
122+{
123+ gert::StorageShape shape = {{1}, {1}};
124+ gert::InfershapeContextPara::TensorDescription x(shape, ge::DT_FLOAT, ge::FORMAT_ND);
125+ gert::InfershapeContextPara::TensorDescription y(shape, ge::DT_FLOAT, ge::FORMAT_ND);
126+ gert::InfershapeContextPara infershapeContextPara("Dawsn", {x}, {y});
127+ std::vector<std::vector<int64_t>> expectOutputShape = {
128+ {1},
129+ };
130+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
131+}