[CANNBOT]ascosh_grad适配Ascend950 Ascendc开发-ops-math-AtomGit

Mdocs/zh/op_list.md+10-0

	@@ -65,6 +65,16 @@
65	<td>×</td>	65	<td>×</td>
66	<td>AI CORE/AI CPU</td>	66	<td>AI CORE/AI CPU</td>
67	<td>该算子暂无Ascend C代码实现，欢迎开发者补充贡献，贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>	67	<td>该算子暂无Ascend C代码实现，欢迎开发者补充贡献，贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>
		68	+ </tr>
		69	+ <tr>
		70	+ <td>math</td>
		71	+ <td><a href="../../math/acosh_grad/README.md">acosh_grad</a></td>
		72	+ <td>√</td>
		73	+ <td>√</td>
		74	+ <td>×</td>
		陈陈陈佳良6月11日这里标记 op_api 为 ×，需要确认算子是否真的没有 aclnn 调用接口。如果有接口实现，应该标记为 √。 W wangweidong 6月11日评论： 6月11日评论：
		75	+ <td>√</td>
		76	+ <td>AI CORE</td>
		77	+ <td>计算Acosh算子的反向梯度，公式 z = dy / sqrt(y² - 1)。</td>
68	</tr>	78	</tr>
69	<tr>	79	<tr>
70	<td>math</td>	80	<td>math</td>

Amath/acosh_grad/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	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	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	14	+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
	15	+set(SUPPORT_TILING_DIR "arch35")
	16	+add_all_modules_sources(OPTYPE acosh_grad ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)

Amath/acosh_grad/README.md+78-0

@@ -0,0 +1,78 @@
	1	+# AcoshGrad
	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	+- 算子功能：计算Acosh（反双曲余弦）算子的反向梯度。
	17	+- 算子公式：$$
	18	+ z_i = dy_i \cdot \dfrac{1}{\sqrt{y_i^2 - 1}}
	19	+ $$其中：
	20	+ - $y_i$为前向Acosh算子的输入张量，值域期望$\geq 1$；
	21	+ - $dy_i$为上游传入的梯度；
	22	+ - $z_i$为对原始输入张量的梯度，等于上游梯度乘以$1/\sqrt{y_i^2 - 1}$。
	23	+
	24	+- 边界（$y_i \leq 1$）：$y_i = 1$时$\sqrt{0}=0$，除零结果为$+\text{Inf}$；$y_i < 1$时$\sqrt{负数}$结果为NaN，按IEEE 754自然传播。建议输入$y$的值域落在$[1, +\infty)$，以避免产生Inf或NaN结果。
	25	+
	26	+## 参数说明
	27	+
	28	+<table style="undefined;table-layout: fixed; width: 1576px"><colgroup>
	29	+ <col style="width: 170px">
	30	+ <col style="width: 170px">
	31	+ <col style="width: 310px">
	32	+ <col style="width: 212px">
	33	+ <col style="width: 100px">
	34	+ </colgroup>
	35	+ <thead>
	36	+ <tr>
	37	+ <th>参数名</th>
	38	+ <th>输入/输出/属性</th>
	39	+ <th>描述</th>
	40	+ <th>数据类型</th>
	41	+ <th>数据格式</th>
	42	+ </tr></thead>
	43	+ <tbody>
	44	+ <tr>
	45	+ <td>y</td>
	46	+ <td>输入</td>
	47	+ <td>前向Acosh算子的输入张量。值域期望落在[1, +∞)。</td>
	48	+ <td>FLOAT16, FLOAT32, BFLOAT16</td>
	49	+ <td>ND</td>
	50	+ </tr>
	51	+ <tr>
	52	+ <td>dy</td>
	53	+ <td>输入</td>
	54	+ <td>上游传入的梯度张量，shape与dtype与y一致。</td>
	55	+ <td>FLOAT16, FLOAT32, BFLOAT16</td>
	56	+ <td>ND</td>
	57	+ </tr>
	58	+ <tr>
	59	+ <td>z</td>
	60	+ <td>输出</td>
	61	+ <td>对原始输入张量的梯度，shape与dtype与y一致。</td>
	62	+ <td>FLOAT16, FLOAT32, BFLOAT16</td>
	63	+ <td>ND</td>
	64	+ </tr>
	65	+ </tbody></table>
	66	+
	67	+## 约束说明
	68	+
	69	+- y与dy的shape必须完全一致。
	70	+- y与dy的dtype必须完全一致。
	71	+- 仅支持ND格式。
	72	+- 支持[非连续的Tensor](../../docs/zh/context/非连续的Tensor.md)，非连续的Tensor维度不大于8。
	73	+
	74	+## 调用说明
	75	+
	76	+\| 调用方式 \| 调用样例 \| 说明 \|
	77	+\| -------- \| ----------------------------------------------------------------------------- \| ------------------------------------------------------------------------------------------ \|
	78	+\| 图模式调用 \| [test_geir_acosh_grad](./examples/test_geir_acosh_grad.cpp) \| 通过[算子IR](./op_graph/acosh_grad_proto.h) 构图方式调用AcoshGrad算子。 \|

Amath/acosh_grad/examples/arch35/test_geir_acosh_grad.cpp+327-0

@@ -0,0 +1,327 @@
	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 "ge_ir_build.h"
	29	+
	30	+#include "array_ops.h"
	31	+#include "nn_other.h"
	32	+#include "../../op_graph/acosh_grad_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	+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	42	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	43	+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
	44	+ TensorDesc placeholder##intputIndex##_desc = \
	45	+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
	46	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	47	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	48	+ Tensor tensor_placeholder##intputIndex; \
	49	+ ret = GenOnesData(placeholder##intputIndex##_shape, \
	50	+ tensor_placeholder##intputIndex, \
	51	+ placeholder##intputIndex##_desc, \
	52	+ intputDtype, \
	53	+ 2); \
	54	+ if (ret != SUCCESS) { \
	55	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	56	+ return FAILED; \
	57	+ } \
	58	+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
	59	+ input.push_back(tensor_placeholder##intputIndex); \
	60	+ graph.AddOp(placeholder##intputIndex); \
	61	+ acoshGrad.set_input_##intputName(placeholder##intputIndex); \
	62	+ inputs.push_back(placeholder##intputIndex)
	63	+
	64	+#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	65	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	66	+ auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \
	67	+ TensorDesc placeholder##intputIndex##_desc = \
	68	+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
	69	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	70	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	71	+ Tensor tensor_placeholder##intputIndex; \
	72	+ ret = GenOnesData(placeholder##intputIndex##_shape, \
	73	+ tensor_placeholder##intputIndex, \
	74	+ placeholder##intputIndex##_desc, \
	75	+ intputDtype, \
	76	+ 2); \
	77	+ if (ret != SUCCESS) { \
	78	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	79	+ return FAILED; \
	80	+ } \
	81	+ placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \
	82	+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
	83	+ graph.AddOp(placeholder##intputIndex); \
	84	+ acoshGrad.set_input_##intputName(placeholder##intputIndex); \
	85	+ acoshGrad.update_input_desc_##intputName(placeholder##intputIndex##_desc); \
	86	+ inputs.push_back(placeholder##intputIndex)
	87	+
	88	+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
	89	+ TensorDesc outputName##outputIndex##_desc = \
	90	+ TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
	91	+ acoshGrad.update_output_desc_##outputName(outputName##outputIndex##_desc)
	92	+
	93	+#define LOG_PRINT(message, ...) printf(message, ##__VA_ARGS__)
	94	+
	95	+string GetTime()
	96	+{
	97	+ time_t timep;
	98	+ time(&timep);
	99	+ char tmp[64];
	100	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
	101	+ return tmp;
	102	+}
	103	+
	104	+uint32_t GetDataTypeSize(DataType dt)
	105	+{
	106	+ uint32_t dilation = 1;
	107	+ uint32_t oneByte = 1;
	108	+ uint32_t twoByte = 2;
	109	+ uint32_t fourByte = 4;
	110	+ uint32_t eightByte = 8;
	111	+
	112	+ if (dt == ge::DT_FLOAT) {
	113	+ dilation = fourByte;
	114	+ } else if (dt == ge::DT_FLOAT16) {
	115	+ dilation = twoByte;
	116	+ } else if (dt == ge::DT_BF16) {
	117	+ dilation = twoByte;
	118	+ } else if (dt == ge::DT_INT16) {
	119	+ dilation = twoByte;
	120	+ } else if (dt == ge::DT_UINT16) {
	121	+ dilation = twoByte;
	122	+ } else if (dt == ge::DT_INT32) {
	123	+ dilation = fourByte;
	124	+ } else if (dt == ge::DT_UINT32) {
	125	+ dilation = fourByte;
	126	+ } else if (dt == ge::DT_INT64) {
	127	+ dilation = eightByte;
	128	+ } else if (dt == ge::DT_UINT64) {
	129	+ dilation = eightByte;
	130	+ } else if (dt == ge::DT_INT8) {
	131	+ dilation = oneByte;
	132	+ }
	133	+ return dilation;
	134	+}
	135	+
	136	+int32_t GenOnesData(
	137	+ vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, DataType data_type, int value)
	138	+{
	139	+ input_tensor_desc.SetRealDimCnt(shapes.size());
	140	+ size_t size = 1;
	141	+ for (uint32_t i = 0; i < shapes.size(); i++) {
	142	+ size *= shapes[i];
	143	+ }
	144	+ uint32_t data_len = size * GetDataTypeSize(data_type);
	145	+ if (data_type == ge::DT_FLOAT) {
	146	+ float *pData = new (std::nothrow) float[size];
	147	+ for (uint32_t i = 0; i < size; ++i) {
	148	+ *(pData + i) = static_cast<float>(value);
	149	+ }
	150	+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t *>(pData), data_len);
	151	+ return SUCCESS;
	152	+ }
	153	+ int32_t *pData = new (std::nothrow) int32_t[size];
	154	+ for (uint32_t i = 0; i < size; ++i) {
	155	+ *(pData + i) = value;
	156	+ }
	157	+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t *>(pData), data_len);
	158	+ return SUCCESS;
	159	+}
	160	+
	161	+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t *inputData)
	162	+{
	163	+ FILE *fp = fopen(bin_file.c_str(), "wb");
	164	+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
	165	+ fclose(fp);
	166	+ return SUCCESS;
	167	+}
	168	+
	169	+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
	170	+ std::vector<Operator> &outputs, Graph &graph)
	171	+{
	172	+ Status ret = SUCCESS;
	173	+ // 自定义代码：添加单算子定义到图中
	174	+ auto acoshGrad = op::AcoshGrad("acosh_grad");
	175	+
	176	+ std::vector<int64_t> yShape = {4, 2};
	177	+ ADD_INPUT(1, y, inDtype, yShape);
	178	+ ADD_INPUT(2, dy, inDtype, yShape);
	179	+
	180	+ ADD_OUTPUT(1, z, inDtype, yShape);
	181	+
	182	+ outputs.push_back(acoshGrad);
	183	+ // 添加完毕
	184	+ return SUCCESS;
	185	+}
	186	+
	187	+
	188	+int InitializeAndSetupGraph(Graph& graph, std::vector<ge::Tensor>& input, DataType inDtype) {
	189	+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
	190	+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
	191	+ Status ret = ge::GEInitialize(global_options);
	192	+ if (ret != SUCCESS) {
	193	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
	194	+ return FAILED;
	195	+ }
	196	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
	197	+
	198	+ std::vector<Operator> inputs{};
	199	+ std::vector<Operator> outputs{};
	200	+
	201	+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
	202	+ if (ret != SUCCESS) {
	203	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	204	+ return FAILED;
	205	+ }
	206	+
	207	+ if (!inputs.empty() && !outputs.empty()) {
	208	+ graph.SetInputs(inputs).SetOutputs(outputs);
	209	+ }
	210	+
	211	+ return SUCCESS;
	212	+}
	213	+
	214	+
	215	+int ExecuteGraph(Graph& graph, std::vector<ge::Tensor>& input, std::vector<ge::Tensor>& output) {
	216	+ std::map<AscendString, AscendString> build_options = {};
	217	+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
	218	+ ge::Session *session = new Session(build_options);
	219	+
	220	+ if (session == nullptr) {
	221	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	222	+ return FAILED;
	223	+ }
	224	+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
	225	+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
	226	+
	227	+ std::map<AscendString, AscendString> graph_options = {};
	228	+ uint32_t graph_id = 0;
	229	+ Status ret = session->AddGraph(graph_id, graph, graph_options);
	230	+
	231	+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
	232	+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
	233	+ std::string file_path = "./dump";
	234	+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
	235	+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
	236	+
	237	+ ret = session->RunGraph(graph_id, input, output);
	238	+ if (ret != SUCCESS) {
	239	+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
	240	+ delete session;
	241	+ GEFinalize();
	242	+ return FAILED;
	243	+ }
	244	+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
	245	+
	246	+ delete session;
	247	+ return SUCCESS;
	248	+}
	249	+
	250	+void ProcessIOData(std::vector<ge::Tensor>& input, std::vector<ge::Tensor>& output) {
	251	+ int input_num = input.size();
	252	+ for (int i = 0; i < input_num; i++) {
	253	+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
	254	+ string input_file = "./tc_ge_irrun_test_0009_npu_input_" + std::to_string(i) + ".bin";
	255	+ uint8_t *input_data_i = input[i].GetData();
	256	+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
	257	+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
	258	+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
	259	+ WriteDataToFile((const char *)input_file.c_str(), data_size, input_data_i);
	260	+ }
	261	+
	262	+ int output_num = output.size();
	263	+ for (int i = 0; i < output_num; i++) {
	264	+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
	265	+ string output_file = "./tc_ge_irrun_test_0009_npu_output_" + std::to_string(i) + ".bin";
	266	+ uint8_t *output_data_i = output[i].GetData();
	267	+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
	268	+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
	269	+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
	270	+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
	271	+ float* resultData = (float*)output_data_i;
	272	+ float expected = 2.0f / 1.7320508f; // dy/sqrt(y*y-1) = 2/sqrt(3)
	273	+ bool allok = true;
	274	+ for (int64_t j = 0; j < output_shape; j++) {
	275	+ LOG_PRINT("result[%ld] is: %f (expected %f)\n", j, resultData[j], expected);
	276	+ if (resultData[j] < expected - 0.01f \|\| resultData[j] > expected + 0.01f) {
	277	+ allok = false;
	278	+ }
	279	+ }
	280	+ LOG_PRINT("%s - INFO - [XIR]: AcoshGrad precision %s\n", GetTime().c_str(), allok ? "PASS" : "FAIL");
	281	+ }
	282	+}
	283	+
	284	+
	285	+int main(int argc, char *argv[])
	286	+{
	287	+ // 1、创建图对象
	288	+ const char *graph_name = "tc_ge_irrun_acosh_grad_test";
	289	+ Graph graph(graph_name);
	290	+ std::vector<ge::Tensor> input;
	291	+
	292	+ std::cout << argv[1] << std::endl;
	293	+
	294	+ DataType inDtype = DT_FLOAT;
	295	+
	296	+ std::cout << inDtype << std::endl;
	297	+
	298	+ // 初始化和设置图
	299	+ if (InitializeAndSetupGraph(graph, input, inDtype) != SUCCESS) {
	300	+ return FAILED;
	301	+ }
	302	+
	303	+ // 执行图计算
	304	+ std::vector<ge::Tensor> output;
	305	+ if (ExecuteGraph(graph, input, output) != SUCCESS) {
	306	+ return FAILED;
	307	+ }
	308	+
	309	+ // 处理输入输出数据
	310	+ ProcessIOData(input, output);
	311	+
	312	+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
	313	+ std::string error_str(error_msg.GetString());
	314	+ std::cout << "Error message: " << error_str << std::endl;
	315	+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
	316	+ std::string warning_str(warning_msg.GetString());
	317	+ std::cout << "Warning message: " << warning_str << std::endl;
	318	+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
	319	+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
	320	+ Status ret = ge::GEFinalize();
	321	+ if (ret != SUCCESS) {
	322	+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
	323	+ return FAILED;
	324	+ }
	325	+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
	326	+ return SUCCESS;
	327	+}

Amath/acosh_grad/examples/test_geir_acosh_grad.cpp+327-0

@@ -0,0 +1,327 @@
	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 "ge_ir_build.h"
	29	+
	30	+#include "array_ops.h"
	31	+#include "nn_other.h"
	32	+#include "../op_graph/acosh_grad_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	+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	42	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	43	+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
	44	+ TensorDesc placeholder##intputIndex##_desc = \
	45	+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
	46	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	47	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	48	+ Tensor tensor_placeholder##intputIndex; \
	49	+ ret = GenOnesData(placeholder##intputIndex##_shape, \
	50	+ tensor_placeholder##intputIndex, \
	51	+ placeholder##intputIndex##_desc, \
	52	+ intputDtype, \
	53	+ 2); \
	54	+ if (ret != SUCCESS) { \
	55	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	56	+ return FAILED; \
	57	+ } \
	58	+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
	59	+ input.push_back(tensor_placeholder##intputIndex); \
	60	+ graph.AddOp(placeholder##intputIndex); \
	61	+ acoshGrad.set_input_##intputName(placeholder##intputIndex); \
	62	+ inputs.push_back(placeholder##intputIndex)
	63	+
	64	+#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	65	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	66	+ auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \
	67	+ TensorDesc placeholder##intputIndex##_desc = \
	68	+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
	69	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	70	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	71	+ Tensor tensor_placeholder##intputIndex; \
	72	+ ret = GenOnesData(placeholder##intputIndex##_shape, \
	73	+ tensor_placeholder##intputIndex, \
	74	+ placeholder##intputIndex##_desc, \
	75	+ intputDtype, \
	76	+ 2); \
	77	+ if (ret != SUCCESS) { \
	78	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	79	+ return FAILED; \
	80	+ } \
	81	+ placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \
	82	+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
	83	+ graph.AddOp(placeholder##intputIndex); \
	84	+ acoshGrad.set_input_##intputName(placeholder##intputIndex); \
	85	+ acoshGrad.update_input_desc_##intputName(placeholder##intputIndex##_desc); \
	86	+ inputs.push_back(placeholder##intputIndex)
	87	+
	88	+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
	89	+ TensorDesc outputName##outputIndex##_desc = \
	90	+ TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
	91	+ acoshGrad.update_output_desc_##outputName(outputName##outputIndex##_desc)
	92	+
	93	+#define LOG_PRINT(message, ...) printf(message, ##__VA_ARGS__)
	94	+
	95	+string GetTime()
	96	+{
	97	+ time_t timep;
	98	+ time(&timep);
	99	+ char tmp[64];
	100	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
	101	+ return tmp;
	102	+}
	103	+
	104	+uint32_t GetDataTypeSize(DataType dt)
	105	+{
	106	+ uint32_t dilation = 1;
	107	+ uint32_t oneByte = 1;
	108	+ uint32_t twoByte = 2;
	109	+ uint32_t fourByte = 4;
	110	+ uint32_t eightByte = 8;
	111	+
	112	+ if (dt == ge::DT_FLOAT) {
	113	+ dilation = fourByte;
	114	+ } else if (dt == ge::DT_FLOAT16) {
	115	+ dilation = twoByte;
	116	+ } else if (dt == ge::DT_BF16) {
	117	+ dilation = twoByte;
	118	+ } else if (dt == ge::DT_INT16) {
	119	+ dilation = twoByte;
	120	+ } else if (dt == ge::DT_UINT16) {
	121	+ dilation = twoByte;
	122	+ } else if (dt == ge::DT_INT32) {
	123	+ dilation = fourByte;
	124	+ } else if (dt == ge::DT_UINT32) {
	125	+ dilation = fourByte;
	126	+ } else if (dt == ge::DT_INT64) {
	127	+ dilation = eightByte;
	128	+ } else if (dt == ge::DT_UINT64) {
	129	+ dilation = eightByte;
	130	+ } else if (dt == ge::DT_INT8) {
	131	+ dilation = oneByte;
	132	+ }
	133	+ return dilation;
	134	+}
	135	+
	136	+int32_t GenOnesData(
	137	+ vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, DataType data_type, int value)
	138	+{
	139	+ input_tensor_desc.SetRealDimCnt(shapes.size());
	140	+ size_t size = 1;
	141	+ for (uint32_t i = 0; i < shapes.size(); i++) {
	142	+ size *= shapes[i];
	143	+ }
	144	+ uint32_t data_len = size * GetDataTypeSize(data_type);
	145	+ if (data_type == ge::DT_FLOAT) {
	146	+ float *pData = new (std::nothrow) float[size];
	147	+ for (uint32_t i = 0; i < size; ++i) {
	148	+ *(pData + i) = static_cast<float>(value);
	149	+ }
	150	+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t *>(pData), data_len);
	151	+ return SUCCESS;
	152	+ }
	153	+ int32_t *pData = new (std::nothrow) int32_t[size];
	154	+ for (uint32_t i = 0; i < size; ++i) {
	155	+ *(pData + i) = value;
	156	+ }
	157	+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t *>(pData), data_len);
	158	+ return SUCCESS;
	159	+}
	160	+
	161	+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t *inputData)
	162	+{
	163	+ FILE *fp = fopen(bin_file.c_str(), "wb");
	164	+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
	165	+ fclose(fp);
	166	+ return SUCCESS;
	167	+}
	168	+
	169	+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
	170	+ std::vector<Operator> &outputs, Graph &graph)
	171	+{
	172	+ Status ret = SUCCESS;
	173	+ // 自定义代码：添加单算子定义到图中
	174	+ auto acoshGrad = op::AcoshGrad("acosh_grad");
	175	+
	176	+ std::vector<int64_t> yShape = {4, 2};
	177	+ ADD_INPUT(1, y, inDtype, yShape);
	178	+ ADD_INPUT(2, dy, inDtype, yShape);
	179	+
	180	+ ADD_OUTPUT(1, z, inDtype, yShape);
	181	+
	182	+ outputs.push_back(acoshGrad);
	183	+ // 添加完毕
	184	+ return SUCCESS;
	185	+}
	186	+
	187	+
	188	+int InitializeAndSetupGraph(Graph& graph, std::vector<ge::Tensor>& input, DataType inDtype) {
	189	+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
	190	+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
	191	+ Status ret = ge::GEInitialize(global_options);
	192	+ if (ret != SUCCESS) {
	193	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
	194	+ return FAILED;
	195	+ }
	196	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
	197	+
	198	+ std::vector<Operator> inputs{};
	199	+ std::vector<Operator> outputs{};
	200	+
	201	+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
	202	+ if (ret != SUCCESS) {
	203	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	204	+ return FAILED;
	205	+ }
	206	+
	207	+ if (!inputs.empty() && !outputs.empty()) {
	208	+ graph.SetInputs(inputs).SetOutputs(outputs);
	209	+ }
	210	+
	211	+ return SUCCESS;
	212	+}
	213	+
	214	+
	215	+int ExecuteGraph(Graph& graph, std::vector<ge::Tensor>& input, std::vector<ge::Tensor>& output) {
	216	+ std::map<AscendString, AscendString> build_options = {};
	217	+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
	218	+ ge::Session *session = new Session(build_options);
	219	+
	220	+ if (session == nullptr) {
	221	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	222	+ return FAILED;
	223	+ }
	224	+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
	225	+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
	226	+
	227	+ std::map<AscendString, AscendString> graph_options = {};
	228	+ uint32_t graph_id = 0;
	229	+ Status ret = session->AddGraph(graph_id, graph, graph_options);
	230	+
	231	+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
	232	+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
	233	+ std::string file_path = "./dump";
	234	+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
	235	+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
	236	+
	237	+ ret = session->RunGraph(graph_id, input, output);
	238	+ if (ret != SUCCESS) {
	239	+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
	240	+ delete session;
	241	+ GEFinalize();
	242	+ return FAILED;
	243	+ }
	244	+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
	245	+
	246	+ delete session;
	247	+ return SUCCESS;
	248	+}
	249	+
	250	+void ProcessIOData(std::vector<ge::Tensor>& input, std::vector<ge::Tensor>& output) {
	251	+ int input_num = input.size();
	252	+ for (int i = 0; i < input_num; i++) {
	253	+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
	254	+ string input_file = "./tc_ge_irrun_test_0009_npu_input_" + std::to_string(i) + ".bin";
	255	+ uint8_t *input_data_i = input[i].GetData();
	256	+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
	257	+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
	258	+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
	259	+ WriteDataToFile((const char *)input_file.c_str(), data_size, input_data_i);
	260	+ }
	261	+
	262	+ int output_num = output.size();
	263	+ for (int i = 0; i < output_num; i++) {
	264	+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
	265	+ string output_file = "./tc_ge_irrun_test_0009_npu_output_" + std::to_string(i) + ".bin";
	266	+ uint8_t *output_data_i = output[i].GetData();
	267	+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
	268	+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
	269	+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
	270	+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
	271	+ float* resultData = (float*)output_data_i;
	272	+ float expected = 2.0f / 1.7320508f; // dy/sqrt(y*y-1) = 2/sqrt(3)
	273	+ bool allok = true;
	274	+ for (int64_t j = 0; j < output_shape; j++) {
	275	+ LOG_PRINT("result[%ld] is: %f (expected %f)\n", j, resultData[j], expected);
	276	+ if (resultData[j] < expected - 0.01f \|\| resultData[j] > expected + 0.01f) {
	277	+ allok = false;
	278	+ }
	279	+ }
	280	+ LOG_PRINT("%s - INFO - [XIR]: AcoshGrad precision %s\n", GetTime().c_str(), allok ? "PASS" : "FAIL");
	281	+ }
	282	+}
	283	+
	284	+
	285	+int main(int argc, char *argv[])
	286	+{
	287	+ // 1、创建图对象
	288	+ const char *graph_name = "tc_ge_irrun_acosh_grad_test";
	289	+ Graph graph(graph_name);
	290	+ std::vector<ge::Tensor> input;
	291	+
	292	+ std::cout << argv[1] << std::endl;
	293	+
	294	+ DataType inDtype = DT_FLOAT;
	295	+
	296	+ std::cout << inDtype << std::endl;
	297	+
	298	+ // 初始化和设置图
	299	+ if (InitializeAndSetupGraph(graph, input, inDtype) != SUCCESS) {
	300	+ return FAILED;
	301	+ }
	302	+
	303	+ // 执行图计算
	304	+ std::vector<ge::Tensor> output;
	305	+ if (ExecuteGraph(graph, input, output) != SUCCESS) {
	306	+ return FAILED;
	307	+ }
	308	+
	309	+ // 处理输入输出数据
	310	+ ProcessIOData(input, output);
	311	+
	312	+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
	313	+ std::string error_str(error_msg.GetString());
	314	+ std::cout << "Error message: " << error_str << std::endl;
	315	+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
	316	+ std::string warning_str(warning_msg.GetString());
	317	+ std::cout << "Warning message: " << warning_str << std::endl;
	318	+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
	319	+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
	320	+ Status ret = ge::GEFinalize();
	321	+ if (ret != SUCCESS) {
	322	+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
	323	+ return FAILED;
	324	+ }
	325	+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
	326	+ return SUCCESS;
	327	+}

Amath/acosh_grad/op_graph/acosh_grad_proto.h+46-0

@@ -0,0 +1,46 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file acosh_grad_proto.h
	15	+ * \brief AcoshGrad GE IR 图模式注册
	16	+ */
	17	+#ifndef OPS_OP_PROTO_INC_ACOSHGRAD_H_
	18	+#define OPS_OP_PROTO_INC_ACOSHGRAD_H_
	19	+
	20	+#include "graph/operator_reg.h"
	21	+#include "graph/types.h"
	22	+
	23	+namespace ge {
	24	+
	25	+/**
	26	+ *@brief AcoshGrad: z = dy / sqrt(y² - 1), reverse gradient of acosh.
	27	+ *@par Inputs:
	28	+ *Two inputs:
	29	+ * @li y: A Tensor. Must be float16, bfloat16, or float32. Original input of forward acosh, domain ≥ 1.
	30	+ * @li dy: A Tensor. Must be float16, bfloat16, or float32. Upstream gradient, same shape/dtype as y.
	31	+ *
	32	+ *@par Outputs:
	33	+ *z: A Tensor. Must be float16, bfloat16, or float32. Reverse gradient, same shape/dtype as y.
	34	+ *
	35	+ *@par Third-party framework compatibility:
	36	+ *Compatible with TensorFlow AcoshGrad.
	37	+ */
	38	+REG_OP(AcoshGrad)
	39	+ .INPUT(y, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
	40	+ .INPUT(dy, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
	41	+ .OUTPUT(z, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
	42	+ .OP_END_FACTORY_REG(AcoshGrad)
	43	+
	44	+} // namespace ge
	45	+
	46	+#endif // OPS_OP_PROTO_INC_ACOSHGRAD_H_

Amath/acosh_grad/op_graph/fusion_pass/.gitkeep+0-0

The file is empty

Amath/acosh_grad/op_host/acosh_grad_def.cpp+57-0

@@ -0,0 +1,57 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file acosh_grad_def.cpp
	15	+ * \brief AcoshGrad 算子定义
	16	+ *
	17	+ * 迭代二：FP16 + BF16 + FP32 全 dtype 注册。
	18	+ */
	19	+#include "register/op_def_registry.h"
	20	+
	21	+namespace ops {
	22	+class AcoshGrad : public OpDef {
	23	+public:
	24	+ explicit AcoshGrad(const char* name) : OpDef(name)
	25	+ {
	26	+ this->Input("y")
	27	+ .ParamType(REQUIRED)
	28	+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
	29	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	30	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	31	+ .AutoContiguous();
	32	+ this->Input("dy")
	33	+ .ParamType(REQUIRED)
	34	+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
	35	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	36	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	37	+ .AutoContiguous();
	38	+ this->Output("z")
	39	+ .ParamType(REQUIRED)
	40	+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
	41	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	42	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	43	+ .AutoContiguous();
	44	+
	45	+ OpAICoreConfig aiCoreConfig;
	46	+ aiCoreConfig.DynamicCompileStaticFlag(true)
	47	+ .DynamicFormatFlag(false)
	48	+ .DynamicRankSupportFlag(true)
	49	+ .DynamicShapeSupportFlag(true)
	50	+ .NeedCheckSupportFlag(false)
	51	+ .PrecisionReduceFlag(true)
	52	+ .ExtendCfgInfo("opFile.value", "acosh_grad_apt");
	53	+ this->AICore().AddConfig("ascend950", aiCoreConfig);
	54	+ }
	55	+};
	56	+OP_ADD(AcoshGrad);
	57	+} // namespace ops

Amath/acosh_grad/op_host/acosh_grad_infershape.cpp+44-0

@@ -0,0 +1,44 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file acosh_grad_infershape.cpp
	15	+ * \brief AcoshGrad 算子形状推导
	16	+ *
	17	+ * z.shape = y.shape（element-wise，无广播）
	18	+ */
	19	+
	20	+#include "register/op_impl_registry.h"
	21	+#include "exe_graph/runtime/infer_shape_context.h"
	22	+#include "op_common/log/log.h"
	23	+
	24	+using namespace ge;
	25	+
	26	+namespace ops {
	27	+
	28	+static ge::graphStatus InferShape4AcoshGrad(gert::InferShapeContext* context)
	29	+{
	30	+ const gert::Shape* input_shape = context->GetInputShape(0);
	31	+ OP_CHECK_NULL_WITH_CONTEXT(context, input_shape);
	32	+
	33	+ gert::Shape* output_shape = context->GetOutputShape(0);
	34	+ OP_CHECK_NULL_WITH_CONTEXT(context, output_shape);
	35	+
	36	+ // z.shape = y.shape
	37	+ output_shape = input_shape;
	38	+
	39	+ return ge::GRAPH_SUCCESS;
	40	+}
	41	+
	42	+IMPL_OP_INFERSHAPE(AcoshGrad).InferShape(InferShape4AcoshGrad);
	43	+
	44	+} // namespace ops

Amath/acosh_grad/op_host/arch35/acosh_grad_tiling_arch35.cpp+182-0

@@ -0,0 +1,182 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file acosh_grad_tiling_arch35.cpp
	15	+ * \brief AcoshGrad Host 侧 Tiling 计算（arch35 / DAV_3510）
	16	+ *
	17	+ * 迭代二：FP16 + BF16 + FP32, Double Buffer, workspace = 0
	18	+ */
	19	+
	20	+#include "register/op_def_registry.h"
	21	+#include "op_common/log/log.h"
	22	+#include "op_common/op_host/util/math_util.h"
	23	+#include "op_common/op_host/util/platform_util.h"
	24	+#include "../../op_kernel/arch35/acosh_grad_tiling_data.h"
	25	+#include "../../op_kernel/arch35/acosh_grad_tiling_key.h"
	26	+
	27	+namespace optiling {
	28	+
	29	+using Ops::Base::CeilDiv;
	30	+using Ops::Base::CeilAlign;
	31	+using Ops::Base::FloorDiv;
	32	+using Ops::Base::FloorAlign;
	33	+using Ops::Base::GetUbBlockSize;
	34	+
	35	+constexpr uint32_t WS_SYS_SIZE = 0U;
	36	+constexpr size_t WORKSPACE_NUM = 1;
	37	+constexpr int64_t DOUBLE_BUF_TENSOR_COUNT = 6; // 3 queues × 2 slots
	38	+constexpr int64_t WORK_BUF_COUNT = 3; // 3 份 FP32 中间 buffer
	39	+constexpr int64_t MIN_SPLIT_THRESHOLD = 1024; // 双缓冲阈值
	40	+constexpr int64_t TYPE_SIZE_FP16 = 2;
	41	+constexpr int64_t TYPE_SIZE_FP32 = 4;
	42	+
	43	+static const gert::Shape g_vec_1_shape = {1};
	44	+
	45	+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape) {
	46	+ if (in_shape.GetDimNum() == 0) {
	47	+ return g_vec_1_shape;
	48	+ }
	49	+ return in_shape;
	50	+}
	51	+
	52	+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)
	53	+{
	54	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	55	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	56	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	57	+ *coreNum = ascendcPlatform.GetCoreNumAiv();
	58	+ OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
	59	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize);
	60	+ OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
	61	+ return ge::GRAPH_SUCCESS;
	62	+}
	63	+
	64	+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t* totalIdx, ge::DataType* dataType)
	65	+{
	66	+ auto inputY = context->GetInputShape(0);
	67	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputY);
	68	+ auto inputShapeY = EnsureNotScalar(inputY->GetStorageShape());
	69	+
	70	+ auto inputDy = context->GetInputShape(1);
	71	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDy);
	72	+ auto inputShapeDy = EnsureNotScalar(inputDy->GetStorageShape());
	73	+
	74	+ auto outZ = context->GetOutputShape(0);
	75	+ OP_CHECK_NULL_WITH_CONTEXT(context, outZ);
	76	+ auto outShapeZ = EnsureNotScalar(outZ->GetStorageShape());
	77	+
	78	+ // shape 校验：y.shape == dy.shape == z.shape
	79	+ OP_CHECK_IF(
	80	+ inputShapeY.GetShapeSize() != inputShapeDy.GetShapeSize() \|\|
	81	+ inputShapeY.GetShapeSize() != outShapeZ.GetShapeSize(),
	82	+ OP_LOGE(context, "AcoshGrad: shape mismatch: y=%ld, dy=%ld, z=%ld",
	83	+ inputShapeY.GetShapeSize(), inputShapeDy.GetShapeSize(), outShapeZ.GetShapeSize()),
	84	+ return ge::GRAPH_FAILED);
	85	+
	86	+ *totalIdx = inputShapeY.GetShapeSize();
	87	+
	88	+ auto inputDesc = context->GetInputDesc(0);
	89	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	90	+ *dataType = inputDesc->GetDataType();
	91	+
	92	+ // 迭代二：支持 FP16、BF16、FP32
	93	+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT};
	94	+ OP_CHECK_IF(supportedDtype.count(*dataType) == 0,
	95	+ OP_LOGE(context, "AcoshGrad: unsupported dtype %d", static_cast<int>(*dataType)),
	96	+ return ge::GRAPH_FAILED);
	97	+
	98	+ return ge::GRAPH_SUCCESS;
	99	+}
	100	+
	101	+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
	102	+{
	103	+ size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM);
	104	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	105	+ currentWorkspace[0] = WS_SYS_SIZE;
	106	+ return ge::GRAPH_SUCCESS;
	107	+}
	108	+
	109	+static ge::graphStatus AcoshGradTilingFunc(gert::TilingContext* context)
	110	+{
	111	+ // 1. 获取平台信息
	112	+ uint64_t ubSize;
	113	+ int64_t coreNum;
	114	+ OP_CHECK_IF(
	115	+ GetPlatformInfo(context, &ubSize, &coreNum) != ge::GRAPH_SUCCESS,
	116	+ OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
	117	+
	118	+ // 2. 获取 shape/属性
	119	+ int64_t totalIdx;
	120	+ ge::DataType dataType;
	121	+ OP_CHECK_IF(
	122	+ GetShapeAttrsInfo(context, &totalIdx, &dataType) != ge::GRAPH_SUCCESS,
	123	+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
	124	+
	125	+ // 3. Workspace
	126	+ OP_CHECK_IF(
	127	+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
	128	+ OP_LOGE(context, "GetWorkspaceSize error"), return ge::GRAPH_FAILED);
	129	+
	130	+ // 4. 设置 TilingData
	131	+ AcoshGradTilingData* tiling = context->GetTilingData<AcoshGradTilingData>();
	132	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	133	+ OP_CHECK_IF(
	134	+ memset_s(tiling, sizeof(AcoshGradTilingData), 0, sizeof(AcoshGradTilingData)) != EOK,
	135	+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
	136	+
	137	+ // 空 Tensor 检查
	138	+ if (totalIdx == 0) {
	139	+ context->SetBlockDim(1);
	140	+ context->SetTilingKey(GET_TPL_TILING_KEY(ACOSH_GRAD_TPL_SCH_MODE_0));
	141	+ return ge::GRAPH_SUCCESS;
	142	+ }
	143	+
	144	+ // 5. 多核切分
	145	+ int64_t ubBlockSize = Ops::Base::GetUbBlockSize(context);
	146	+ tiling->totalNum = totalIdx;
	147	+ tiling->blockFactor = CeilAlign(CeilDiv(totalIdx, coreNum), ubBlockSize);
	148	+ int64_t usedCoreNum = Ops::Base::CeilDiv(totalIdx, tiling->blockFactor);
	149	+
	150	+ // 6. UB 切分
	151	+ // FP16/BF16 + Double Buffer:
	152	+ // Queue 内存: 3 queues × 2 slots × ubFactor × 2 = 12 × ubFactor
	153	+ // Work 内存: 3 bufs × ubFactor × sizeof(float) = 12 × ubFactor
	154	+ // 总计: 24 × ubFactor bytes
	155	+ // FP32 + Double Buffer:
	156	+ // Queue 内存: 3 queues × 2 slots × ubFactor × 4 = 24 × ubFactor
	157	+ // Work 内存: 3 bufs × ubFactor × sizeof(float) = 12 × ubFactor
	158	+ // 总计: 36 × ubFactor bytes
	159	+ int64_t elemBytes = (dataType == ge::DT_FLOAT) ? TYPE_SIZE_FP32 : TYPE_SIZE_FP16;
	160	+ int64_t queueBufCount = DOUBLE_BUF_TENSOR_COUNT;
	161	+ int64_t totalBytesPerElem = queueBufCount * elemBytes + WORK_BUF_COUNT * static_cast<int64_t>(sizeof(float));
	162	+ tiling->ubFactor = FloorAlign(FloorDiv(static_cast<int64_t>(ubSize), totalBytesPerElem), ubBlockSize);
	163	+
	164	+ context->SetBlockDim(usedCoreNum);
	165	+
	166	+ // 7. 设置 TilingKey（使用 GET_TPL_TILING_KEY，与 relu6_grad 模式一致）
	167	+ context->SetTilingKey(GET_TPL_TILING_KEY(ACOSH_GRAD_TPL_SCH_MODE_0));
	168	+ return ge::GRAPH_SUCCESS;
	169	+}
	170	+
	171	+static ge::graphStatus TilingParseForAcoshGrad([[maybe_unused]] gert::TilingParseContext* context)
	172	+{
	173	+ return ge::GRAPH_SUCCESS;
	174	+}
	175	+
	176	+struct AcoshGradCompileInfo {};
	177	+
	178	+IMPL_OP_OPTILING(AcoshGrad)
	179	+ .Tiling(AcoshGradTilingFunc)
	180	+ .TilingParse<AcoshGradCompileInfo>(TilingParseForAcoshGrad);
	181	+
	182	+} // namespace optiling

Amath/acosh_grad/op_kernel/acosh_grad_apt.cpp+35-0

@@ -0,0 +1,35 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file acosh_grad_apt.cpp
	15	+ * \brief AcoshGrad 算子 Kernel 入口（arch35 / DAV_3510）
	16	+ *
	17	+ * 模板参数（对应 acosh_grad_tiling_key.h 中 ASCENDC_TPL_ARGS_DECL）：
	18	+ * - schMode: 调度模式（0 = 逐元素）
	19	+ *
	20	+ * dtype 通过 DTYPE_Y 宏指定（CANN 编译系统根据 op_def 输入名 "y" 自动生成）。
	21	+ * 每个 dtype 变体由编译系统独立编译，DTYPE_Y 自动替换为对应类型。
	22	+ */
	23	+
	24	+#include "arch35/acosh_grad_kernel.h"
	25	+
	26	+template <uint32_t schMode>
	27	+__global__ __aicore__ void acosh_grad(GM_ADDR y, GM_ADDR dy, GM_ADDR z,
	28	+ GM_ADDR workspace, GM_ADDR tiling)
	29	+{
	30	+ REGISTER_TILING_DEFAULT(AcoshGradTilingData);
	31	+ GET_TILING_DATA_WITH_STRUCT(AcoshGradTilingData, tilingData, tiling);
	32	+ NsAcoshGrad::AcoshGradKernel<DTYPE_Y> op;
	33	+ op.Init(y, dy, z, &tilingData);
	34	+ op.Process();
	35	+}

Amath/acosh_grad/op_kernel/arch35/acosh_grad_kernel.h+227-0

@@ -0,0 +1,227 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file acosh_grad_kernel.h
	15	+ * \brief AcoshGrad 算子 Kernel 类定义（arch35 / DAV_3510）
	16	+ *
	17	+ * 公式: z = dy / sqrt(y² - 1)
	18	+ * 计算路径 (FP16/BF16): Cast→FP32 → Mul → Adds(-1) → Sqrt → Div → Cast→原dtype
	19	+ * 计算路径 (FP32): Mul → Adds(-1) → Sqrt → Div
	20	+ * Double Buffer 流水线并行
	21	+ *
	22	+ * 关键设计约束（来自迭代二穿刺验证）：
	23	+ * 1. GlobalTensor 多 tile 不自动推进：每次 tile 循环前必须显式调用
	24	+ * SetGlobalBuffer(base + offset, remain) 更新 GM 指针偏移
	25	+ * 2. TBuf 必须用独立实例：3 个独立 TBuf<VECCALC>（bufSlot0_/1_/2_），
	26	+ * 禁止单 TBuf + GetWithOffset（totalNum=1 时仅 12 字节导致 +Inf）
	27	+ *
	28	+ * dtype 通过 DTYPE_Y 宏指定（由 CANN 编译系统根据 op_def 输入名 "y" 自动生成）
	29	+ */
	30	+#ifndef ACOSH_GRAD_KERNEL_H
	31	+#define ACOSH_GRAD_KERNEL_H
	32	+
	33	+#include "kernel_operator.h"
	34	+#include "kernel_tiling/kernel_tiling.h"
	35	+#include "acosh_grad_tiling_data.h"
	36	+#include "acosh_grad_tiling_key.h"
	37	+
	38	+namespace NsAcoshGrad {
	39	+
	40	+using namespace AscendC;
	41	+
	42	+template <typename T>
	43	+class AcoshGradKernel {
	44	+ static constexpr int32_t BUFFER_NUM = 2; // Double Buffer
	45	+
	46	+public:
	47	+ __aicore__ inline AcoshGradKernel() {};
	48	+
	49	+ __aicore__ inline void Init(GM_ADDR y, GM_ADDR dy, GM_ADDR z,
	50	+ const AcoshGradTilingData* tilingData);
	51	+ __aicore__ inline void Process();
	52	+
	53	+private:
	54	+ __aicore__ inline void CopyIn(int64_t currentNum);
	55	+ __aicore__ inline void Compute(int64_t currentNum);
	56	+ __aicore__ inline void CopyOut(int64_t currentNum);
	57	+
	58	+private:
	59	+ TPipe pipe;
	60	+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueY;
	61	+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueDy;
	62	+ TQue<TPosition::VECOUT, BUFFER_NUM> outQueueZ;
	63	+
	64	+ // 3 个独立 FP32 中间计算 buffer（禁止合并为单 TBuf + GetWithOffset）
	65	+ // 穿刺 probe_fp16_overflow 验证：单 TBuf 在 totalNum=1 时仅分配 12 字节，
	66	+ // 导致 slot 地址重叠，输出 +Inf
	67	+ TBuf<TPosition::VECCALC> bufSlot0_; // yFp32 → dyFp32 → result
	68	+ TBuf<TPosition::VECCALC> bufSlot1_; // ySq → ySqM1
	69	+ TBuf<TPosition::VECCALC> bufSlot2_; // sqrtVal
	70	+
	71	+ GlobalTensor<T> yGM;
	72	+ GlobalTensor<T> dyGM;
	73	+ GlobalTensor<T> zGM;
	74	+
	75	+ // Base GM 地址（用于 tile 偏移计算，穿刺 probe_large_shape 验证必须保存）
	76	+ __gm__ T* yBase_ = nullptr;
	77	+ __gm__ T* dyBase_ = nullptr;
	78	+ __gm__ T* zBase_ = nullptr;
	79	+
	80	+ int64_t blockLength_ = 0;
	81	+ int64_t ubLength_ = 0;
	82	+};
	83	+
	84	+template <typename T>
	85	+__aicore__ inline void AcoshGradKernel<T>::Init(
	86	+ GM_ADDR y, GM_ADDR dy, GM_ADDR z, const AcoshGradTilingData* tilingData)
	87	+{
	88	+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx();
	89	+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
	陈陈陈佳良6月11日 remainderLength 可能是负数（当 GetBlockIdx() 过大），但这里没有检查负数情况。Process() 中只检查了 blockLength_ == 0，负数会漏检导致后续计算出错。参考 ndtri_kernel.h:152-154 的做法，建议加上 if (blockLength_ < 0) blockLength_ = 0;
	90	+ // blockIdx 过大时 remainderLength 可能为负，钳为 0 避免后续计算出错（参考 ndtri_kernel.h:152-154）
	91	+ if (blockLength_ < 0) {
	92	+ blockLength_ = 0;
	93	+ }
	94	+ ubLength_ = tilingData->ubFactor;
	95	+
	96	+ // 保存 base GM 地址，Process 中按 tile 偏移调用 SetGlobalBuffer
	97	+ // （穿刺 probe_large_shape 验证：DataCopyPad 不自动推进 GlobalTensor 地址）
	98	+ yBase_ = (__gm__ T)y + tilingData->blockFactor AscendC::GetBlockIdx();
	99	+ dyBase_ = (__gm__ T)dy + tilingData->blockFactor AscendC::GetBlockIdx();
	100	+ zBase_ = (__gm__ T)z + tilingData->blockFactor AscendC::GetBlockIdx();
	101	+
	102	+ // GlobalTensor 初始设置：指向第一个 tile
	103	+ yGM.SetGlobalBuffer(yBase_, blockLength_);
	104	+ dyGM.SetGlobalBuffer(dyBase_, blockLength_);
	105	+ zGM.SetGlobalBuffer(zBase_, blockLength_);
	106	+
	107	+ pipe.InitBuffer(inQueueY, BUFFER_NUM, ubLength_ * sizeof(T));
	108	+ pipe.InitBuffer(inQueueDy, BUFFER_NUM, ubLength_ * sizeof(T));
	109	+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, ubLength_ * sizeof(T));
	110	+
	111	+ // FP32 中间 buffer：3 个独立 TBuf 实例（禁止合并为单 TBuf）
	112	+ // 穿刺 probe_fp16_overflow 验证：单 TBuf + GetWithOffset 在 totalNum=1 时
	113	+ // 仅分配 12 字节（3*sizeof(float)），slot 地址重叠导致输出 +Inf
	114	+ pipe.InitBuffer(bufSlot0_, ubLength_ * sizeof(float));
	115	+ pipe.InitBuffer(bufSlot1_, ubLength_ * sizeof(float));
	116	+ pipe.InitBuffer(bufSlot2_, ubLength_ * sizeof(float));
	117	+}
	118	+
	119	+template <typename T>
	120	+__aicore__ inline void AcoshGradKernel<T>::CopyIn(int64_t currentNum)
	121	+{
	122	+ LocalTensor<T> yLocal = inQueueY.template AllocTensor<T>();
	123	+ LocalTensor<T> dyLocal = inQueueDy.template AllocTensor<T>();
	124	+
	125	+ // DataCopyPad 直接从 GlobalTensor 当前地址读取（Process 中已推进到正确 tile）
	126	+ DataCopyParams copyParams;
	127	+ copyParams.blockCount = 1;
	128	+ copyParams.blockLen = currentNum * sizeof(T);
	129	+ copyParams.srcStride = 0;
	130	+ copyParams.dstStride = 0;
	131	+
	132	+ DataCopyPad(yLocal, yGM, copyParams, {false, 0, 0, 0});
	133	+ DataCopyPad(dyLocal, dyGM, copyParams, {false, 0, 0, 0});
	134	+
	135	+ inQueueY.EnQue(yLocal);
	136	+ inQueueDy.EnQue(dyLocal);
	137	+}
	138	+
	139	+template <typename T>
	140	+__aicore__ inline void AcoshGradKernel<T>::Compute(int64_t currentNum)
	141	+{
	142	+ LocalTensor<T> yLocal = inQueueY.template DeQue<T>();
	143	+ LocalTensor<T> dyLocal = inQueueDy.template DeQue<T>();
	144	+ LocalTensor<T> zLocal = outQueueZ.template AllocTensor<T>();
	145	+
	146	+ // 从 3 个独立 TBuf 获取 FP32 中间 Buffer（无地址冲突）
	147	+ // 穿刺 probe_fp16_overflow 验证：独立 TBuf 实例保证地址不重叠
	148	+ LocalTensor<float> slot0 = bufSlot0_.Get<float>(); // yFp32 → dyFp32 → result
	149	+ LocalTensor<float> slot1 = bufSlot1_.Get<float>(); // ySq → ySqM1
	150	+ LocalTensor<float> slot2 = bufSlot2_.Get<float>(); // sqrtVal
	151	+
	152	+ if constexpr (std::is_same_v<T, float>) {
	153	+ // FP32 路径：直接计算，无需 Cast
	154	+ // slot0 = y² (临时)
	155	+ AscendC::Mul<float>(slot0, yLocal, yLocal, currentNum);
	156	+ // slot1 = y² - 1
	157	+ AscendC::Adds<float>(slot1, slot0, -1.0f, currentNum);
	158	+ // slot2 = sqrt(y² - 1)
	159	+ AscendC::Sqrt<float>(slot2, slot1, currentNum);
	160	+ // zLocal = dy / sqrt(y² - 1)
	161	+ AscendC::Div<float>(zLocal, dyLocal, slot2, currentNum);
	162	+ } else {
	163	+ // FP16/BF16 路径：Cast→FP32→计算→Cast
	164	+ // Step 1: Cast y → FP32 (slot0)
	165	+ AscendC::Cast<float, T>(slot0, yLocal, AscendC::RoundMode::CAST_NONE, currentNum);
	166	+ // Step 2: y² (slot1)
	167	+ AscendC::Mul<float>(slot1, slot0, slot0, currentNum);
	168	+ // Step 3: y² - 1 (slot1, in-place)
	169	+ AscendC::Adds<float>(slot1, slot1, -1.0f, currentNum);
	170	+ // Step 4: sqrt(y² - 1) (slot2)
	171	+ AscendC::Sqrt<float>(slot2, slot1, currentNum);
	172	+ // Step 5: Cast dy → FP32 (slot0, 复用)
	173	+ AscendC::Cast<float, T>(slot0, dyLocal, AscendC::RoundMode::CAST_NONE, currentNum);
	174	+ // Step 6: dy / sqrt (slot0, 复用)
	175	+ AscendC::Div<float>(slot0, slot0, slot2, currentNum);
	176	+ // Step 7: Cast result → 原 dtype (zLocal)
	177	+ AscendC::Cast<T, float>(zLocal, slot0, AscendC::RoundMode::CAST_ROUND, currentNum);
	178	+ }
	179	+
	180	+ outQueueZ.template EnQue<T>(zLocal);
	181	+ inQueueY.FreeTensor(yLocal);
	182	+ inQueueDy.FreeTensor(dyLocal);
	183	+}
	184	+
	185	+template <typename T>
	186	+__aicore__ inline void AcoshGradKernel<T>::CopyOut(int64_t currentNum)
	187	+{
	188	+ LocalTensor<T> zLocal = outQueueZ.template DeQue<T>();
	189	+
	190	+ // DataCopyPad 直接写入 GlobalTensor 当前地址（Process 中已推进到正确 tile）
	191	+ DataCopyParams copyParams;
	192	+ copyParams.blockCount = 1;
	193	+ copyParams.blockLen = currentNum * sizeof(T);
	194	+ copyParams.srcStride = 0;
	195	+ copyParams.dstStride = 0;
	196	+
	197	+ DataCopyPad(zGM, zLocal, copyParams);
	198	+ outQueueZ.FreeTensor(zLocal);
	199	+}
	200	+
	201	+template <typename T>
	202	+__aicore__ inline void AcoshGradKernel<T>::Process()
	203	+{
	204	+ if (blockLength_ == 0) return;
	205	+
	206	+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
	207	+ for (int64_t i = 0; i < loopCount; i++) {
	208	+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
	209	+
	210	+ // 显式推进 GlobalTensor 到当前 tile 偏移
	211	+ // （穿刺 probe_large_shape 验证：DataCopyPad 不自动推进 GlobalTensor 地址，
	212	+ // 多 tile 循环必须在每次迭代前重新调用 SetGlobalBuffer）
	213	+ int64_t offset = i * ubLength_;
	214	+ int64_t remain = blockLength_ - offset;
	215	+ yGM.SetGlobalBuffer(yBase_ + offset, remain);
	216	+ dyGM.SetGlobalBuffer(dyBase_ + offset, remain);
	217	+ zGM.SetGlobalBuffer(zBase_ + offset, remain);
	218	+
	219	+ CopyIn(currentNum);
	220	+ Compute(currentNum);
	221	+ CopyOut(currentNum);
	222	+ }
	223	+}
	224	+
	225	+} // namespace NsAcoshGrad
	226	+
	227	+#endif // ACOSH_GRAD_KERNEL_H

Amath/acosh_grad/op_kernel/arch35/acosh_grad_tiling_data.h+31-0

@@ -0,0 +1,31 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file acosh_grad_tiling_data.h
	15	+ * \brief TilingData 结构体定义（acosh_grad, arch35）
	16	+ *
	17	+ * 使用标准 C++ struct，禁止 BEGIN_TILING_DATA_DEF 宏
	18	+ */
	19	+
	20	+#ifndef _ACOSH_GRAD_TILING_DATA_H_
	21	+#define _ACOSH_GRAD_TILING_DATA_H_
	22	+
	23	+#include <cstdint>
	24	+
	25	+struct AcoshGradTilingData {
	26	+ int64_t totalNum = 0; // 总元素数量
	27	+ int64_t blockFactor = 0; // 每个核处理的元素数量（对齐后）
	28	+ int64_t ubFactor = 0; // 每次 UB tile 处理的元素数量（对齐后）
	29	+};
	30	+
	31	+#endif // _ACOSH_GRAD_TILING_DATA_H_

Amath/acosh_grad/op_kernel/arch35/acosh_grad_tiling_key.h+39-0

@@ -0,0 +1,39 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file acosh_grad_tiling_key.h
	15	+ * \brief TilingKey 模板参数定义（acosh_grad, arch35）
	16	+ *
	17	+ * 迭代一：单 schMode（0 = 逐元素），dtype 由 DTYPE_Y 自动宏处理。
	18	+ * 使用 ASCENDC_TPL_ARGS_DECL 模板编程。
	19	+ *
	20	+ * 注意：不使用 ASCENDC_TPL_DATATYPE_DECL —— CANN 9.0 beta.2 的 compile_op
	21	+ * 生成的 wrapper 代码无法正确推导 typename 模板参数。dtype 通过 op_def 中
	22	+ * 输入名 "y" 自动生成的 DTYPE_Y 宏在 kernel entry 中指定。
	23	+ */
	24	+
	25	+#ifndef __ACOSH_GRAD_TILING_KEY_H__
	26	+#define __ACOSH_GRAD_TILING_KEY_H__
	27	+
	28	+#include "ascendc/host_api/tiling/template_argument.h"
	29	+
	30	+#define ACOSH_GRAD_TPL_SCH_MODE_0 0
	31	+
	32	+ASCENDC_TPL_ARGS_DECL(
	33	+ AcoshGrad,
	34	+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ACOSH_GRAD_TPL_SCH_MODE_0));
	35	+
	36	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
	37	+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ACOSH_GRAD_TPL_SCH_MODE_0)));
	38	+
	39	+#endif // __ACOSH_GRAD_TILING_KEY_H__

Amath/acosh_grad/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-2026 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()

Amath/acosh_grad/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-2026 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()

Amath/acosh_grad/tests/ut/op_host/CMakeLists.txt+21-0

@@ -0,0 +1,21 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify it.
	3	+# Copyright (c) 2025-2026 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	+if(UT_TEST_ALL OR OP_HOST_UT)
	13	+ add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	14	+endif()
	15	+
	16	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	17	+foreach(SUB_DIR ${CURRENT_DIRS})
	18	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	19	+ add_subdirectory(${SUB_DIR})
	20	+ endif()
	21	+endforeach()

Amath/acosh_grad/tests/ut/op_host/arch35/CMakeLists.txt+14-0

@@ -0,0 +1,14 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify it.
	3	+# Copyright (c) 2025-2026 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	+if(UT_TEST_ALL OR OP_HOST_UT)
	13	+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	14	+endif()

Amath/acosh_grad/tests/ut/op_host/arch35/test_acosh_grad_tiling_arch35.cpp+311-0

@@ -0,0 +1,311 @@
	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 test_acosh_grad_tiling_arch35.cpp
	15	+ * \brief AcoshGrad Tiling UT (arch35 / Ascend950)
	16	+ *
	17	+ * Covers paths in op_host/arch35/acosh_grad_tiling_arch35.cpp:
	18	+ * 1) Dtype paths: FP16 / BF16 / FP32 all map to the SAME TilingKey
	19	+ * (GET_TPL_TILING_KEY(ACOSH_GRAD_TPL_SCH_MODE_0)=0). dtype only changes
	20	+ * bytesPerElem in UB-split (fp32=36, fp16/bf16=24) → different ubFactor.
	21	+ * 2) Multi-core path (large shape) vs single-core path (small shape, tail).
	22	+ * 3) Non-aligned tail.
	23	+ * 4) Empty tensor (totalNum==0) → early-return, all fields zero.
	24	+ * 5) Failure paths:
	25	+ * - shape mismatch (different sizes) → GRAPH_FAILED
	26	+ * - unsupported dtype (DT_DOUBLE) → GRAPH_FAILED
	27	+ *
	28	+ * Platform constants (UT faker): coreNum=64, ubSize=262144, ubBlockSize=32.
	29	+ *
	30	+ * TilingData struct (acosh_grad_tiling_data.h):
	31	+ * int64_t totalNum; int64_t blockFactor; int64_t ubFactor;
	32	+ * Serialized as 3× int64 → "totalNum blockFactor ubFactor ".
	33	+ *
	34	+ * Formulas:
	35	+ * blockFactor = CeilAlign(CeilDiv(total, 64), 32)
	36	+ * perElem = 6elemBytes + 34 (fp32: 36, fp16/bf16: 24)
	37	+ * ubFactor = FloorAlign(262144 / perElem, 32) (fp32: 7264, fp16/bf16: 10912)
	38	+ */
	39	+
	40	+#include <cstdint>
	41	+#include <iostream>
	42	+#include <gtest/gtest.h>
	43	+#include "tiling_context_faker.h"
	44	+#include "tiling_case_executor.h"
	45	+
	46	+namespace optiling {
	47	+struct AcoshGradCompileInfo {};
	48	+} // namespace optiling
	49	+
	50	+using namespace std;
	51	+
	52	+class AcoshGradTilingTest : public testing::Test {
	53	+protected:
	54	+ static void SetUpTestCase()
	55	+ {
	56	+ std::cout << "AcoshGradTilingTest SetUp" << std::endl;
	57	+ }
	58	+
	59	+ static void TearDownTestCase()
	60	+ {
	61	+ std::cout << "AcoshGradTilingTest TearDown" << std::endl;
	62	+ }
	63	+};
	64	+
	65	+// ===========================================================================
	66	+// 1) FP32 multi-core aligned — 8192 elem {1,64,2,64}
	67	+// blockFactor=CeilAlign(CeilDiv(8192,64)=128,32)=128, ubFactor=7264
	68	+// TilingKey 0
	69	+// ===========================================================================
	70	+TEST_F(AcoshGradTilingTest, test_tiling_fp32_multi_core_001)
	71	+{
	72	+ optiling::AcoshGradCompileInfo compileInfo;
	73	+ gert::TilingContextPara tilingContextPara(
	74	+ "AcoshGrad",
	75	+ {
	76	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND}, // y
	77	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND}, // dy
	78	+ },
	79	+ {
	80	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND}, // z
	81	+ },
	82	+ &compileInfo);
	83	+ uint64_t expectTilingKey = 0;
	84	+ string expectTilingData = "8192 128 7264 ";
	85	+ std::vector<size_t> expectWorkspaces = {0};
	86	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	87	+}
	88	+
	89	+// ===========================================================================
	90	+// 2) FP16 multi-core aligned — same shape; ubFactor=10912
	91	+// ===========================================================================
	92	+TEST_F(AcoshGradTilingTest, test_tiling_fp16_multi_core_002)
	93	+{
	94	+ optiling::AcoshGradCompileInfo compileInfo;
	95	+ gert::TilingContextPara tilingContextPara(
	96	+ "AcoshGrad",
	97	+ {
	98	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	99	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	100	+ },
	101	+ {
	102	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	103	+ },
	104	+ &compileInfo);
	105	+ uint64_t expectTilingKey = 0;
	106	+ string expectTilingData = "8192 128 10912 ";
	107	+ std::vector<size_t> expectWorkspaces = {0};
	108	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	109	+}
	110	+
	111	+// ===========================================================================
	112	+// 3) BF16 multi-core aligned — same shape; ubFactor=10912
	113	+// ===========================================================================
	114	+TEST_F(AcoshGradTilingTest, test_tiling_bf16_multi_core_003)
	115	+{
	116	+ optiling::AcoshGradCompileInfo compileInfo;
	117	+ gert::TilingContextPara tilingContextPara(
	118	+ "AcoshGrad",
	119	+ {
	120	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_BF16, ge::FORMAT_ND},
	121	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_BF16, ge::FORMAT_ND},
	122	+ },
	123	+ {
	124	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_BF16, ge::FORMAT_ND},
	125	+ },
	126	+ &compileInfo);
	127	+ uint64_t expectTilingKey = 0;
	128	+ string expectTilingData = "8192 128 10912 ";
	129	+ std::vector<size_t> expectWorkspaces = {0};
	130	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	131	+}
	132	+
	133	+// ===========================================================================
	134	+// 4) FP32 small + single-core + tail — {7} elem
	135	+// blockFactor=CeilAlign(CeilDiv(7,64)=1,32)=32, usedCore=1, ubFactor=7264
	136	+// ===========================================================================
	137	+TEST_F(AcoshGradTilingTest, test_tiling_fp32_small_tail_004)
	138	+{
	139	+ optiling::AcoshGradCompileInfo compileInfo;
	140	+ gert::TilingContextPara tilingContextPara(
	141	+ "AcoshGrad",
	142	+ {
	143	+ {{{7}, {7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	144	+ {{{7}, {7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	145	+ },
	146	+ {
	147	+ {{{7}, {7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	148	+ },
	149	+ &compileInfo);
	150	+ uint64_t expectTilingKey = 0;
	151	+ string expectTilingData = "7 32 7264 ";
	152	+ std::vector<size_t> expectWorkspaces = {0};
	153	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	154	+}
	155	+
	156	+// ===========================================================================
	157	+// 5) FP16 unaligned single-core — {17} elem
	158	+// blockFactor=32, usedCore=1, ubFactor=10912
	159	+// ===========================================================================
	160	+TEST_F(AcoshGradTilingTest, test_tiling_fp16_unalign_005)
	161	+{
	162	+ optiling::AcoshGradCompileInfo compileInfo;
	163	+ gert::TilingContextPara tilingContextPara(
	164	+ "AcoshGrad",
	165	+ {
	166	+ {{{17}, {17}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	167	+ {{{17}, {17}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	168	+ },
	169	+ {
	170	+ {{{17}, {17}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	171	+ },
	172	+ &compileInfo);
	173	+ uint64_t expectTilingKey = 0;
	174	+ string expectTilingData = "17 32 10912 ";
	175	+ std::vector<size_t> expectWorkspaces = {0};
	176	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	177	+}
	178	+
	179	+// ===========================================================================
	180	+// 6) FP32 large multi-core — {416910} elem
	181	+// blockFactor=CeilAlign(CeilDiv(416910,64)=6515,32)=6528, ubFactor=7264
	182	+// ===========================================================================
	183	+TEST_F(AcoshGradTilingTest, test_tiling_fp32_large_multi_core_006)
	184	+{
	185	+ optiling::AcoshGradCompileInfo compileInfo;
	186	+ gert::TilingContextPara tilingContextPara(
	187	+ "AcoshGrad",
	188	+ {
	189	+ {{{416910}, {416910}}, ge::DT_FLOAT, ge::FORMAT_ND},
	190	+ {{{416910}, {416910}}, ge::DT_FLOAT, ge::FORMAT_ND},
	191	+ },
	192	+ {
	193	+ {{{416910}, {416910}}, ge::DT_FLOAT, ge::FORMAT_ND},
	194	+ },
	195	+ &compileInfo);
	196	+ uint64_t expectTilingKey = 0;
	197	+ string expectTilingData = "416910 6528 7264 ";
	198	+ std::vector<size_t> expectWorkspaces = {0};
	199	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	200	+}
	201	+
	202	+// ===========================================================================
	203	+// 7) FP16 2D large multi-core — {60882, 23} = 1400286 elem
	204	+// blockFactor=CeilAlign(CeilDiv(1400286,64)=21880,32)=21888, ubFactor=10912
	205	+// ===========================================================================
	206	+TEST_F(AcoshGradTilingTest, test_tiling_fp16_2d_large_multi_core_007)
	207	+{
	208	+ optiling::AcoshGradCompileInfo compileInfo;
	209	+ gert::TilingContextPara tilingContextPara(
	210	+ "AcoshGrad",
	211	+ {
	212	+ {{{60882, 23}, {60882, 23}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	213	+ {{{60882, 23}, {60882, 23}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	214	+ },
	215	+ {
	216	+ {{{60882, 23}, {60882, 23}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	217	+ },
	218	+ &compileInfo);
	219	+ uint64_t expectTilingKey = 0;
	220	+ string expectTilingData = "1400286 21888 10912 ";
	221	+ std::vector<size_t> expectWorkspaces = {0};
	222	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	223	+}
	224	+
	225	+// ===========================================================================
	226	+// 8) Empty tensor FP32 — early-return branch, all fields zero
	227	+// ===========================================================================
	228	+TEST_F(AcoshGradTilingTest, test_tiling_empty_fp32_008)
	229	+{
	230	+ optiling::AcoshGradCompileInfo compileInfo;
	231	+ gert::TilingContextPara tilingContextPara(
	232	+ "AcoshGrad",
	233	+ {
	234	+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
	235	+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
	236	+ },
	237	+ {
	238	+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
	239	+ },
	240	+ &compileInfo);
	241	+ uint64_t expectTilingKey = 0;
	242	+ string expectTilingData = "0 0 0 ";
	243	+ std::vector<size_t> expectWorkspaces = {0};
	244	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	245	+}
	246	+
	247	+// ===========================================================================
	248	+// 9) Empty tensor BF16 multi-dim — early-return
	249	+// ===========================================================================
	250	+TEST_F(AcoshGradTilingTest, test_tiling_empty_bf16_009)
	251	+{
	252	+ optiling::AcoshGradCompileInfo compileInfo;
	253	+ gert::TilingContextPara tilingContextPara(
	254	+ "AcoshGrad",
	255	+ {
	256	+ {{{0, 8}, {0, 8}}, ge::DT_BF16, ge::FORMAT_ND},
	257	+ {{{0, 8}, {0, 8}}, ge::DT_BF16, ge::FORMAT_ND},
	258	+ },
	259	+ {
	260	+ {{{0, 8}, {0, 8}}, ge::DT_BF16, ge::FORMAT_ND},
	261	+ },
	262	+ &compileInfo);
	263	+ uint64_t expectTilingKey = 0;
	264	+ string expectTilingData = "0 0 0 ";
	265	+ std::vector<size_t> expectWorkspaces = {0};
	266	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	267	+}
	268	+
	269	+// ===========================================================================
	270	+// 10) FAIL: shape mismatch (different sizes) → GRAPH_FAILED
	271	+// ===========================================================================
	272	+TEST_F(AcoshGradTilingTest, test_tiling_fail_shape_mismatch_010)
	273	+{
	274	+ optiling::AcoshGradCompileInfo compileInfo;
	275	+ gert::TilingContextPara tilingContextPara(
	276	+ "AcoshGrad",
	277	+ {
	278	+ {{{16}, {16}}, ge::DT_FLOAT, ge::FORMAT_ND},
	279	+ {{{8}, {8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	280	+ },
	281	+ {
	282	+ {{{16}, {16}}, ge::DT_FLOAT, ge::FORMAT_ND},
	283	+ },
	284	+ &compileInfo);
	285	+ uint64_t expectTilingKey = 0;
	286	+ string expectTilingData = "";
	287	+ std::vector<size_t> expectWorkspaces = {0};
	288	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);
	289	+}
	290	+
	291	+// ===========================================================================
	292	+// 11) FAIL: unsupported dtype (DT_DOUBLE) → GRAPH_FAILED
	293	+// ===========================================================================
	294	+TEST_F(AcoshGradTilingTest, test_tiling_fail_unsupported_dtype_011)
	295	+{
	296	+ optiling::AcoshGradCompileInfo compileInfo;
	297	+ gert::TilingContextPara tilingContextPara(
	298	+ "AcoshGrad",
	299	+ {
	300	+ {{{8}, {8}}, ge::DT_DOUBLE, ge::FORMAT_ND},
	301	+ {{{8}, {8}}, ge::DT_DOUBLE, ge::FORMAT_ND},
	302	+ },
	303	+ {
	304	+ {{{8}, {8}}, ge::DT_DOUBLE, ge::FORMAT_ND},
	305	+ },
	306	+ &compileInfo);
	307	+ uint64_t expectTilingKey = 0;
	308	+ string expectTilingData = "";
	309	+ std::vector<size_t> expectWorkspaces = {0};
	310	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);
	311	+}

Amath/acosh_grad/tests/ut/op_host/test_acosh_grad_infershape.cpp+168-0

@@ -0,0 +1,168 @@
	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 test_acosh_grad_infershape.cpp
	15	+ * \brief AcoshGrad InferShape UT — output shape copies input[0] shape.
	16	+ *
	17	+ * Op signature (aligned with proto.h / def.cpp):
	18	+ * Input 0: y (forward output of Acosh)
	19	+ * Input 1: dy (upstream gradient)
	20	+ * Output 0: z (gradient w.r.t. original input, same shape as y)
	21	+ *
	22	+ * Supported dtypes (per def.cpp): FLOAT16, BF16, FLOAT. Format: ND only.
	23	+ * Shape rule: z.shape = y.shape (acosh_grad_infershape.cpp: output_shape = input_shape).
	24	+ */
	25	+
	26	+#include <gtest/gtest.h>
	27	+#include <iostream>
	28	+#include "infershape_context_faker.h"
	29	+#include "infershape_case_executor.h"
	30	+
	31	+class AcoshGradInfershape : public testing::Test {
	32	+protected:
	33	+ static void SetUpTestCase()
	34	+ {
	35	+ std::cout << "AcoshGradInfershape SetUp" << std::endl;
	36	+ }
	37	+
	38	+ static void TearDownTestCase()
	39	+ {
	40	+ std::cout << "AcoshGradInfershape TearDown" << std::endl;
	41	+ }
	42	+};
	43	+
	44	+// ---------------------------------------------------------------------------
	45	+// 1D, fp32 — small shape
	46	+// ---------------------------------------------------------------------------
	47	+TEST_F(AcoshGradInfershape, acosh_grad_infershape_1d_fp32)
	48	+{
	49	+ gert::InfershapeContextPara infershapeContextPara(
	50	+ "AcoshGrad",
	51	+ {
	52	+ {{{16}, {16}}, ge::DT_FLOAT, ge::FORMAT_ND}, // y
	53	+ {{{16}, {16}}, ge::DT_FLOAT, ge::FORMAT_ND}, // dy
	54	+ },
	55	+ {
	56	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND}, // z (to infer)
	57	+ });
	58	+ std::vector<std::vector<int64_t>> expectOutputShape = {{16}};
	59	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	60	+}
	61	+
	62	+// ---------------------------------------------------------------------------
	63	+// 1D non-aligned, fp32
	64	+// ---------------------------------------------------------------------------
	65	+TEST_F(AcoshGradInfershape, acosh_grad_infershape_1d_tail_fp32)
	66	+{
	67	+ gert::InfershapeContextPara infershapeContextPara(
	68	+ "AcoshGrad",
	69	+ {
	70	+ {{{7}, {7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	71	+ {{{7}, {7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	72	+ },
	73	+ {
	74	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	75	+ });
	76	+ std::vector<std::vector<int64_t>> expectOutputShape = {{7}};
	77	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	78	+}
	79	+
	80	+// ---------------------------------------------------------------------------
	81	+// 2D, fp16
	82	+// ---------------------------------------------------------------------------
	83	+TEST_F(AcoshGradInfershape, acosh_grad_infershape_2d_fp16)
	84	+{
	85	+ gert::InfershapeContextPara infershapeContextPara(
	86	+ "AcoshGrad",
	87	+ {
	88	+ {{{4, 8}, {4, 8}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	89	+ {{{4, 8}, {4, 8}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	90	+ },
	91	+ {
	92	+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	93	+ });
	94	+ std::vector<std::vector<int64_t>> expectOutputShape = {{4, 8}};
	95	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	96	+}
	97	+
	98	+// ---------------------------------------------------------------------------
	99	+// 4D, bf16
	100	+// ---------------------------------------------------------------------------
	101	+TEST_F(AcoshGradInfershape, acosh_grad_infershape_4d_bf16)
	102	+{
	103	+ gert::InfershapeContextPara infershapeContextPara(
	104	+ "AcoshGrad",
	105	+ {
	106	+ {{{2, 3, 4, 5}, {2, 3, 4, 5}}, ge::DT_BF16, ge::FORMAT_ND},
	107	+ {{{2, 3, 4, 5}, {2, 3, 4, 5}}, ge::DT_BF16, ge::FORMAT_ND},
	108	+ },
	109	+ {
	110	+ {{{}, {}}, ge::DT_BF16, ge::FORMAT_ND},
	111	+ });
	112	+ std::vector<std::vector<int64_t>> expectOutputShape = {{2, 3, 4, 5}};
	113	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	114	+}
	115	+
	116	+// ---------------------------------------------------------------------------
	117	+// Large multi-core shape, fp32
	118	+// ---------------------------------------------------------------------------
	119	+TEST_F(AcoshGradInfershape, acosh_grad_infershape_large_fp32)
	120	+{
	121	+ gert::InfershapeContextPara infershapeContextPara(
	122	+ "AcoshGrad",
	123	+ {
	124	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
	125	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
	126	+ },
	127	+ {
	128	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	129	+ });
	130	+ std::vector<std::vector<int64_t>> expectOutputShape = {{1, 64, 2, 64}};
	131	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	132	+}
	133	+
	134	+// ---------------------------------------------------------------------------
	135	+// Large 1D, fp16 — stress test
	136	+// ---------------------------------------------------------------------------
	137	+TEST_F(AcoshGradInfershape, acosh_grad_infershape_large_1d_fp16)
	138	+{
	139	+ gert::InfershapeContextPara infershapeContextPara(
	140	+ "AcoshGrad",
	141	+ {
	142	+ {{{416910}, {416910}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	143	+ {{{416910}, {416910}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	144	+ },
	145	+ {
	146	+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	147	+ });
	148	+ std::vector<std::vector<int64_t>> expectOutputShape = {{416910}};
	149	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	150	+}
	151	+
	152	+// ---------------------------------------------------------------------------
	153	+// Empty tensor, fp32
	154	+// ---------------------------------------------------------------------------
	155	+TEST_F(AcoshGradInfershape, acosh_grad_infershape_empty_fp32)
	156	+{
	157	+ gert::InfershapeContextPara infershapeContextPara(
	158	+ "AcoshGrad",
	159	+ {
	160	+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
	161	+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
	162	+ },
	163	+ {
	164	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	165	+ });
	166	+ std::vector<std::vector<int64_t>> expectOutputShape = {{0}};
	167	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	168	+}