[CANNBot]新增ndtri算子AscendC适配Ascend950-ops-math-AtomGit

Mdocs/zh/op_list.md+10-0

	@@ -1305,6 +1305,16 @@
1305	<td>AI Core</td>	1305	<td>AI Core</td>
1306	<td>将输入中的NaN、正无穷大和负无穷大值分别替换为nan、posinf、neginf指定的值。</td>	1306	<td>将输入中的NaN、正无穷大和负无穷大值分别替换为nan、posinf、neginf指定的值。</td>
1307	</tr>	1307	</tr>
		1308	+ <tr>
		1309	+ <td>math</td>
		1310	+ <td><a href="../../math/ndtri/README.md">ndtri</a></td>
		1311	+ <td>√</td>
		1312	+ <td>√</td>
		1313	+ <td>×</td>
		1314	+ <td>√</td>
		1315	+ <td>AI Core</td>
		1316	+ <td>逆标准正态累积分布函数（probit），对输入概率张量逐元素计算标准正态分位点。</td>
		1317	+ </tr>
1308	<tr>	1318	<tr>
1309	<td>math</td>	1319	<td>math</td>
1310	<td><a href="../../math/neg/README.md">neg</a></td>	1320	<td><a href="../../math/neg/README.md">neg</a></td>

Amath/ndtri/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 文件目录
	15	+set(SUPPORT_TILING_DIR "arch35")
	16	+add_all_modules_sources(OPTYPE ndtri ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)

Amath/ndtri/README.md+72-0

@@ -0,0 +1,72 @@
	1	+# Ndtri
	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	+- 算子功能：逆标准正态累积分布函数（probit / inverse normal CDF）。对输入概率张量逐元素计算标准正态分位点，对齐 PyTorch `torch.special.ndtri` 与 SciPy `scipy.special.ndtri` 语义。底层采用 Cephes 数学库的分区间有理逼近算法。
	17	+
	18	+- 计算公式：
	19	+
	20	+$$y_i = \mathrm{ndtri}(x_i) = \Phi^{-1}(x_i) = \sqrt{2} \cdot \mathrm{erf}^{-1}(2 x_i - 1), \quad x_i \in (0, 1)$$
	21	+
	22	+- 特殊值处理（与 PyTorch / SciPy 一致）：
	23	+
	24	+\| 输入 $x_i$ \| 输出 $y_i$ \|
	25	+\| :---: \| :---: \|
	26	+\| `0` \| `-inf` \|
	27	+\| `1` \| `+inf` \|
	28	+\| $x_i < 0$ 或 $x_i > 1$ \| `NaN` \|
	29	+\| `NaN` / `+inf` / `-inf` \| `NaN` \|
	30	+
	31	+## 参数说明
	32	+
	33	+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
	34	+ <col style="width: 100px">
	35	+ <col style="width: 150px">
	36	+ <col style="width: 280px">
	37	+ <col style="width: 330px">
	38	+ <col style="width: 120px">
	39	+ </colgroup>
	40	+ <thead>
	41	+ <tr>
	42	+ <th>参数名</th>
	43	+ <th>输入/输出/属性</th>
	44	+ <th>描述</th>
	45	+ <th>数据类型</th>
	46	+ <th>数据格式</th>
	47	+ </tr></thead>
	48	+ <tbody>
	49	+ <tr>
	50	+ <td>x</td>
	51	+ <td>输入</td>
	52	+ <td>待进行 ndtri 计算的概率张量，公式中的 x_i，推荐值域 (0, 1)。</td>
	53	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	54	+ <td>ND</td>
	55	+ </tr>
	56	+ <tr>
	57	+ <td>y</td>
	58	+ <td>输出</td>
	59	+ <td>逆标准正态 CDF 的计算结果，公式中的 y_i。</td>
	60	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	61	+ <td>ND</td>
	62	+ </tr>
	63	+ </tbody></table>
	64	+
	65	+## 约束说明
	66	+无
	67	+
	68	+## 调用说明
	69	+
	70	+\| 调用方式 \| 调用样例 \| 说明 \|
	71	+\|--------------\|--------------------------------------------\|-----------------------------------------------------------------\|
	72	+\| 图模式调用 \| [test_geir_ndtri](./examples/test_geir_ndtri.cpp) \| 通过[算子IR](./op_graph/ndtri_proto.h)构图方式调用 Ndtri 算子。 \|

Amath/ndtri/examples/arch35/test_geir_ndtri.cpp+256-0

@@ -0,0 +1,256 @@
	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	+// 注：本文件为 examples/test_geir_ndtri.cpp 的 arch35 副本（include 深度 ../../），
	14	+// 由 ops-math 构建流水线扫描 examples/arch35/ 而存在。两份内容需保持同步，
	15	+// 任何修改必须同时应用到两份。建议通过工具脚本 diff 校验。
	16	+
	17	+#include <iostream>
	18	+#include <fstream>
	19	+#include <string.h>
	20	+#include <stdint.h>
	21	+#include <vector>
	22	+#include <string>
	23	+#include <map>
	24	+#include "assert.h"
	25	+
	26	+#include "graph.h"
	27	+#include "types.h"
	28	+#include "tensor.h"
	29	+#include "ge_error_codes.h"
	30	+#include "ge_api_types.h"
	31	+#include "ge_api.h"
	32	+#include "array_ops.h"
	33	+#include "ge_ir_build.h"
	34	+
	35	+#include "nn_other.h"
	36	+#include "../../op_graph/ndtri_proto.h"
	37	+
	38	+#define FAILED -1
	39	+#define SUCCESS 0
	40	+
	41	+using namespace ge;
	42	+using std::map;
	43	+using std::string;
	44	+using std::vector;
	45	+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	46	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	47	+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
	48	+ TensorDesc placeholder##intputIndex##_desc = \
	49	+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
	50	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	51	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	52	+ Tensor tensor_placeholder##intputIndex; \
	53	+ ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, \
	54	+ tensor_placeholder##intputIndex, \
	55	+ placeholder##intputIndex##_desc, \
	56	+ 0.5f); \
	57	+ if (ret != SUCCESS) { \
	58	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	59	+ return FAILED; \
	60	+ } \
	61	+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
	62	+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
	63	+ input.push_back(tensor_placeholder##intputIndex); \
	64	+ graph.AddOp(placeholder##intputIndex); \
	65	+ ndtri_op.set_input_##intputName(placeholder##intputIndex); \
	66	+ inputs.push_back(placeholder##intputIndex);
	67	+
	68	+#define LOG_PRINT(message, ...) \
	69	+ do { \
	70	+ printf(message, ##__VA_ARGS__); \
	71	+ } while (0)
	72	+
	73	+string GetTime()
	74	+{
	75	+ time_t timep;
	76	+ time(&timep);
	77	+ char tmp[64];
	78	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
	79	+ return tmp;
	80	+}
	81	+
	82	+uint32_t GetDataTypeSize(DataType dt)
	83	+{
	84	+ static const std::map<DataType, uint32_t> kSize = {
	85	+ {ge::DT_FLOAT, 4}, {ge::DT_FLOAT16, 2}, {ge::DT_BF16, 2},
	86	+ {ge::DT_INT16, 2}, {ge::DT_UINT16, 2},
	87	+ {ge::DT_INT32, 4}, {ge::DT_UINT32, 4},
	88	+ {ge::DT_INT64, 8}, {ge::DT_UINT64, 8},
	89	+ {ge::DT_INT8, 1},
	90	+ };
	91	+ auto it = kSize.find(dt);
	92	+ return it != kSize.end() ? it->second : 1;
	93	+}
	94	+
	95	+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, float value)
	96	+{
	97	+ input_tensor_desc.SetRealDimCnt(shapes.size());
	98	+ size_t size = 1;
	99	+ for (uint32_t i = 0; i < shapes.size(); i++) {
	100	+ size *= shapes[i];
	101	+ }
	102	+ uint32_t byteSizeFloat32 = 4;
	103	+ uint32_t data_len = size * byteSizeFloat32;
	104	+ float *pData = new (std::nothrow) float[size];
	S S sunday5月25日同上，arch35 副本也存在相同内存泄漏问题。两份文件需同步修复。
	105	+ if (pData == nullptr) {
	106	+ printf("alloc failed\n");
	107	+ return FAILED;
	108	+ }
	109	+
	110	+ for (size_t i = 0; i < size; ++i) {
	111	+ *(pData + i) = value;
	112	+ }
	113	+ input_tensor = Tensor(input_tensor_desc, (uint8_t *)pData, data_len);
	114	+ return SUCCESS;
	115	+}
	116	+
	117	+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t *inputData)
	118	+{
	119	+ FILE *fp = fopen(bin_file.c_str(), "w");
	120	+ if (fp == nullptr) {
	121	+ printf("fopen %s failed\n", bin_file.c_str());
	122	+ return FAILED;
	123	+ }
	124	+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
	125	+ fclose(fp);
	126	+ return SUCCESS;
	127	+}
	128	+
	129	+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
	130	+ std::vector<Operator> &outputs, Graph &graph)
	131	+{
	132	+ Status ret = SUCCESS;
	133	+ // 添加 Ndtri 算子到图：输入 x（概率 ∈ (0,1)），输出 y（正态分位点）
	134	+ auto ndtri_op = op::Ndtri("ndtri_op");
	135	+ std::vector<int64_t> xShape = {4, 2};
	136	+ ADD_INPUT(1, x, inDtype, xShape);
	137	+
	138	+ outputs.push_back(ndtri_op);
	139	+ return SUCCESS;
	140	+}
	141	+
	142	+int main(int argc, char *argv[])
	143	+{
	144	+ if (argc < 2) {
	145	+ printf("INFO: %s invoked without dtype arg, defaulting to DT_FLOAT\n",
	146	+ argv[0] ? argv[0] : "test_geir_ndtri");
	147	+ }
	148	+ const char *graph_name = "tc_ge_irrun_test";
	149	+ Graph graph(graph_name);
	150	+ std::vector<ge::Tensor> input;
	151	+
	152	+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
	153	+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
	154	+ Status ret = ge::GEInitialize(global_options);
	155	+ if (ret != SUCCESS) {
	156	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
	157	+ return FAILED;
	158	+ }
	159	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
	160	+
	161	+ std::vector<Operator> inputs{};
	162	+ std::vector<Operator> outputs{};
	163	+
	164	+ if (argc >= 2) {
	165	+ std::cout << argv[1] << std::endl;
	166	+ }
	167	+
	168	+ DataType inDtype = DT_FLOAT;
	169	+
	170	+ std::cout << inDtype << std::endl;
	171	+
	172	+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
	173	+ if (ret != SUCCESS) {
	174	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	175	+ return FAILED;
	176	+ }
	177	+
	178	+ if (!inputs.empty() && !outputs.empty()) {
	179	+ graph.SetInputs(inputs).SetOutputs(outputs);
	180	+ }
	181	+
	182	+ std::map<AscendString, AscendString> build_options = {
	183	+
	184	+ };
	185	+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
	186	+ ge::Session *session = new Session(build_options);
	187	+
	188	+ if (session == nullptr) {
	189	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	190	+ return FAILED;
	191	+ }
	192	+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
	193	+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
	194	+
	195	+ std::map<AscendString, AscendString> graph_options = {
	196	+
	197	+ };
	198	+ uint32_t graph_id = 0;
	199	+ ret = session->AddGraph(graph_id, graph, graph_options);
	200	+
	201	+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
	202	+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
	203	+ std::string file_path = "./dump";
	204	+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
	205	+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
	206	+ std::vector<ge::Tensor> output;
	207	+ ret = session->RunGraph(graph_id, input, output);
	208	+ if (ret != SUCCESS) {
	209	+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
	210	+ delete session;
	211	+ GEFinalize();
	212	+ return FAILED;
	213	+ }
	214	+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
	215	+
	216	+ int input_num = input.size();
	217	+ for (int i = 0; i < input_num; i++) {
	218	+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
	219	+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
	220	+ uint8_t *input_data_i = input[i].GetData();
	221	+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
	222	+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
	223	+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
	224	+ WriteDataToFile((const char *)input_file.c_str(), data_size, input_data_i);
	225	+ }
	226	+
	227	+ int output_num = output.size();
	228	+ for (int i = 0; i < output_num; i++) {
	229	+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
	230	+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
	231	+ uint8_t *output_data_i = output[i].GetData();
	232	+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
	233	+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
	234	+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
	235	+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
	236	+ float resultData = (float)output_data_i;
	237	+ for (int64_t j = 0; j < output_shape; j++) {
	238	+ LOG_PRINT("result[%ld] is: %f\n", j, resultData[j]);
	239	+ }
	240	+ }
	241	+
	242	+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
	243	+ std::string error_str(error_msg.GetString());
	244	+ std::cout << "Error message: " << error_str << std::endl;
	245	+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
	246	+ std::string warning_str(warning_msg.GetString());
	247	+ std::cout << "Warning message: " << warning_str << std::endl;
	248	+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
	249	+ ret = ge::GEFinalize();
	250	+ if (ret != SUCCESS) {
	251	+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
	252	+ return FAILED;
	253	+ }
	254	+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
	255	+ return SUCCESS;
	256	+}

Amath/ndtri/examples/test_geir_ndtri.cpp+252-0

@@ -0,0 +1,252 @@
	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/ndtri_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 = GenOnesDataFloat32(placeholder##intputIndex##_shape, \
	50	+ tensor_placeholder##intputIndex, \
	51	+ placeholder##intputIndex##_desc, \
	52	+ 0.5f); \
	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	+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
	59	+ input.push_back(tensor_placeholder##intputIndex); \
	60	+ graph.AddOp(placeholder##intputIndex); \
	61	+ ndtri_op.set_input_##intputName(placeholder##intputIndex); \
	62	+ inputs.push_back(placeholder##intputIndex);
	63	+
	64	+#define LOG_PRINT(message, ...) \
	65	+ do { \
	66	+ printf(message, ##__VA_ARGS__); \
	67	+ } while (0)
	68	+
	69	+string GetTime()
	70	+{
	71	+ time_t timep;
	72	+ time(&timep);
	73	+ char tmp[64];
	74	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
	75	+ return tmp;
	76	+}
	77	+
	78	+uint32_t GetDataTypeSize(DataType dt)
	79	+{
	80	+ static const std::map<DataType, uint32_t> kSize = {
	81	+ {ge::DT_FLOAT, 4}, {ge::DT_FLOAT16, 2}, {ge::DT_BF16, 2},
	82	+ {ge::DT_INT16, 2}, {ge::DT_UINT16, 2},
	83	+ {ge::DT_INT32, 4}, {ge::DT_UINT32, 4},
	84	+ {ge::DT_INT64, 8}, {ge::DT_UINT64, 8},
	85	+ {ge::DT_INT8, 1},
	86	+ };
	87	+ auto it = kSize.find(dt);
	88	+ return it != kSize.end() ? it->second : 1;
	89	+}
	90	+
	91	+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, float value)
	92	+{
	93	+ input_tensor_desc.SetRealDimCnt(shapes.size());
	94	+ size_t size = 1;
	95	+ for (uint32_t i = 0; i < shapes.size(); i++) {
	96	+ size *= shapes[i];
	97	+ }
	98	+ uint32_t byteSizeFloat32 = 4;
	99	+ uint32_t data_len = size * byteSizeFloat32;
	100	+ float *pData = new (std::nothrow) float[size];
	S S sunday5月25日这里 `pData = new float[size]` 分配的内存没有释放。Tensor 构造函数不会接管内存 ownership，pData 需要在函数结束后或 input_tensor 销毁时 delete[]，否则存在内存泄漏。建议使用 std::vector 管理内存，或者添加显式释放逻辑。 Nerddddddddddd 5月26日评论： 5月26日评论：
	101	+ if (pData == nullptr) {
	102	+ printf("alloc failed\n");
	103	+ return FAILED;
	104	+ }
	105	+
	106	+ for (size_t i = 0; i < size; ++i) {
	107	+ *(pData + i) = value;
	108	+ }
	109	+ input_tensor = Tensor(input_tensor_desc, (uint8_t *)pData, data_len);
	110	+ return SUCCESS;
	111	+}
	112	+
	113	+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t *inputData)
	114	+{
	115	+ FILE *fp = fopen(bin_file.c_str(), "w");
	116	+ if (fp == nullptr) {
	117	+ printf("fopen %s failed\n", bin_file.c_str());
	118	+ return FAILED;
	119	+ }
	120	+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
	121	+ fclose(fp);
	122	+ return SUCCESS;
	123	+}
	124	+
	125	+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
	126	+ std::vector<Operator> &outputs, Graph &graph)
	127	+{
	128	+ Status ret = SUCCESS;
	129	+ // 添加 Ndtri 算子到图：输入 x（概率 ∈ (0,1)），输出 y（正态分位点）
	130	+ auto ndtri_op = op::Ndtri("ndtri_op");
	131	+ std::vector<int64_t> xShape = {4, 2};
	132	+ ADD_INPUT(1, x, inDtype, xShape);
	133	+
	134	+ outputs.push_back(ndtri_op);
	135	+ return SUCCESS;
	136	+}
	137	+
	138	+int main(int argc, char *argv[])
	139	+{
	140	+ if (argc < 2) {
	141	+ printf("INFO: %s invoked without dtype arg, defaulting to DT_FLOAT\n",
	142	+ argv[0] ? argv[0] : "test_geir_ndtri");
	143	+ }
	144	+ const char *graph_name = "tc_ge_irrun_test";
	145	+ Graph graph(graph_name);
	146	+ std::vector<ge::Tensor> input;
	147	+
	148	+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
	149	+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
	150	+ Status ret = ge::GEInitialize(global_options);
	151	+ if (ret != SUCCESS) {
	152	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
	153	+ return FAILED;
	154	+ }
	155	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
	156	+
	157	+ std::vector<Operator> inputs{};
	158	+ std::vector<Operator> outputs{};
	159	+
	160	+ if (argc >= 2) {
	161	+ std::cout << argv[1] << std::endl;
	162	+ }
	163	+
	164	+ DataType inDtype = DT_FLOAT;
	165	+
	166	+ std::cout << inDtype << std::endl;
	167	+
	168	+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
	169	+ if (ret != SUCCESS) {
	170	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	171	+ return FAILED;
	172	+ }
	173	+
	174	+ if (!inputs.empty() && !outputs.empty()) {
	175	+ graph.SetInputs(inputs).SetOutputs(outputs);
	176	+ }
	177	+
	178	+ std::map<AscendString, AscendString> build_options = {
	179	+
	180	+ };
	181	+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
	182	+ ge::Session *session = new Session(build_options);
	183	+
	184	+ if (session == nullptr) {
	185	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	186	+ return FAILED;
	187	+ }
	188	+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
	189	+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
	190	+
	191	+ std::map<AscendString, AscendString> graph_options = {
	192	+
	193	+ };
	194	+ uint32_t graph_id = 0;
	195	+ ret = session->AddGraph(graph_id, graph, graph_options);
	196	+
	197	+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
	198	+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
	199	+ std::string file_path = "./dump";
	200	+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
	201	+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
	202	+ std::vector<ge::Tensor> output;
	203	+ ret = session->RunGraph(graph_id, input, output);
	204	+ if (ret != SUCCESS) {
	205	+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
	206	+ delete session;
	207	+ GEFinalize();
	208	+ return FAILED;
	209	+ }
	210	+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
	211	+
	212	+ int input_num = input.size();
	213	+ for (int i = 0; i < input_num; i++) {
	214	+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
	215	+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
	216	+ uint8_t *input_data_i = input[i].GetData();
	217	+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
	218	+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
	219	+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
	220	+ WriteDataToFile((const char *)input_file.c_str(), data_size, input_data_i);
	221	+ }
	222	+
	223	+ int output_num = output.size();
	224	+ for (int i = 0; i < output_num; i++) {
	225	+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
	226	+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
	227	+ uint8_t *output_data_i = output[i].GetData();
	228	+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
	229	+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
	230	+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
	231	+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
	232	+ float resultData = (float)output_data_i;
	233	+ for (int64_t j = 0; j < output_shape; j++) {
	234	+ LOG_PRINT("result[%ld] is: %f\n", j, resultData[j]);
	235	+ }
	236	+ }
	237	+
	238	+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
	239	+ std::string error_str(error_msg.GetString());
	240	+ std::cout << "Error message: " << error_str << std::endl;
	241	+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
	242	+ std::string warning_str(warning_msg.GetString());
	243	+ std::cout << "Warning message: " << warning_str << std::endl;
	244	+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
	245	+ ret = ge::GEFinalize();
	246	+ if (ret != SUCCESS) {
	247	+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
	248	+ return FAILED;
	249	+ }
	250	+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
	251	+ return SUCCESS;
	252	+}

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

The file is empty

Amath/ndtri/op_graph/ndtri_proto.h+50-0

@@ -0,0 +1,50 @@
	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 ndtri_proto.h
	15	+ * \brief
	16	+ */
	17	+#ifndef OP_PROTO_NDTRI_PROTO_H_
	18	+#define OP_PROTO_NDTRI_PROTO_H_
	19	+
	20	+#include "graph/operator_reg.h"
	21	+
	22	+namespace ge
	23	+{
	24	+/**
	25	+@brief Computes ndtri element-wise (y = sqrt(2) erfinv(2 * x - 1)).
	26	+
	27	+* @par Inputs:
	28	+* One input, including:
	29	+* x: An ND Tensor. Support 1D~8D. Must be one of the following types:
	30	+* bfloat16, float16, float32, double.
	31	+
	32	+* @par Outputs:
	33	+* y: An ND Tensor. Has the same dtype as "x".
	34	+
	35	+* @par Implementation notes
	36	+* On Ascend 950 (arch35) the aicore kernel currently implements
	37	+* {bfloat16, float16, float32} only. DT_DOUBLE is declared in the
	38	+* proto for future host fallback / aicpu compatibility.
	39	+
	40	+* @par Third-party framework compatibility:
	41	+* Compatible with the TensorFlow operator Ndtri and torch.special.ndtri.
	42	+*/
	43	+REG_OP(Ndtri)
	44	+ .INPUT(x, "T")
	45	+ .OUTPUT(y, "T")
	46	+ .DATATYPE(T, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT, DT_DOUBLE}))
	47	+ .OP_END_FACTORY_REG(Ndtri)
	48	+
	49	+} // namespace ge
	50	+#endif // OP_PROTO_NDTRI_PROTO_H_

Amath/ndtri/op_host/arch35/ndtri_tiling_arch35.cpp+215-0

@@ -0,0 +1,215 @@
	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 ndtri_tiling_arch35.cpp
	15	+ * \brief Ndtri Tiling 实现（arch35 / Ascend950）
	16	+ */
	17	+
	18	+#include "register/op_def_registry.h"
	19	+#include "log/log.h"
	20	+#include "util/math_util.h"
	21	+#include "util/platform_util.h"
	22	+#include "../../op_kernel/arch35/ndtri_tiling_data.h"
	23	+#include "../../op_kernel/arch35/ndtri_tiling_key.h"
	24	+
	25	+namespace optiling {
	26	+
	27	+constexpr uint32_t WS_USER_SIZE = 0U; // 算子自身不需要额外 workspace
	28	+static constexpr size_t IDX_SELF = 0;
	29	+constexpr int64_t TYPE_SIZE_FP32 = 4;
	30	+constexpr int64_t TYPE_SIZE_FP16_BF16 = 2;
	31	+constexpr int64_t RESERVED_UB = 48 * 1024;
	32	+// 按 Kernel 实际 UB 占用估算：
	33	+// in/out (DB): 2 * 2 * sizeof(T) → FP32 16, FP16/BF16 8
	34	+// pBuf + yBuf: 2 * 4 = 8
	35	+// tmpBuf0..10: 11 * 4 = 44
	36	+// maskBuf0..3: ~0.5 (按 bit)
	37	+// FP32 合计 ≈ 68.5 → 取 80（保留余量）
	38	+// FP16/BF16 合计 ≈ 60.5 → 取 64
	39	+constexpr int64_t BYTE_PER_ELEM_FP32 = 80;
	40	+constexpr int64_t BYTE_PER_ELEM_FP16_BF16 = 64;
	41	+constexpr int64_t TILE_ALIGN = 256;
	42	+
	43	+static const gert::Shape K_VEC_1_SHAPE = {1};
	44	+
	45	+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape)
	46	+{
	47	+ if (in_shape.GetDimNum() == 0) {
	48	+ return K_VEC_1_SHAPE;
	49	+ }
	50	+ return in_shape;
	51	+}
	52	+
	53	+// 平台信息
	54	+static ge::graphStatus GetPlatformInfo(
	55	+ gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum,
	56	+ uint32_t& sysWorkspaceSize)
	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"),
	63	+ return ge::GRAPH_FAILED);
	64	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
	65	+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"),
	66	+ return ge::GRAPH_FAILED);
	67	+ sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
	68	+ return ge::GRAPH_SUCCESS;
	69	+}
	70	+
	71	+// dtype 校验
	72	+static ge::graphStatus CheckDtype(gert::TilingContext* context, ge::DataType& dtype)
	73	+{
	74	+ auto selfDesc = context->GetInputDesc(IDX_SELF);
	75	+ OP_CHECK_NULL_WITH_CONTEXT(context, selfDesc);
	76	+ dtype = selfDesc->GetDataType();
	77	+ const std::set<ge::DataType> supported = {
	78	+ ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
	79	+ OP_CHECK_IF(supported.count(dtype) == 0,
	80	+ OP_LOGE(context, "Ndtri: unsupported dtype %d",
	81	+ static_cast<int>(dtype)),
	82	+ return ge::GRAPH_FAILED);
	83	+ return ge::GRAPH_SUCCESS;
	84	+}
	85	+
	86	+// 提取 totalNum
	87	+static ge::graphStatus GetTotalNum(gert::TilingContext* context, int64_t& totalNum)
	88	+{
	89	+ auto selfShapePtr = context->GetInputShape(IDX_SELF);
	90	+ OP_CHECK_NULL_WITH_CONTEXT(context, selfShapePtr);
	91	+ auto selfShape = EnsureNotScalar(selfShapePtr->GetStorageShape());
	92	+ totalNum = selfShape.GetShapeSize();
	93	+ OP_CHECK_IF(totalNum <= 0,
	94	+ OP_LOGE(context, "Ndtri: totalNum must > 0, got %ld", totalNum),
	95	+ return ge::GRAPH_FAILED);
	96	+ return ge::GRAPH_SUCCESS;
	97	+}
	98	+
	99	+// 多核 + UB 切分
	100	+static ge::graphStatus DoTiling(
	101	+ gert::TilingContext* context, ge::DataType dtype, int64_t totalNum,
	102	+ uint64_t ubSize, int64_t coreNum,
	103	+ NdtriTilingData* tiling, int64_t& usedCoreNum, int64_t& alignElem)
	104	+{
	105	+ int64_t ubBlockSize = Ops::Base::GetUbBlockSize(context);
	106	+ int64_t typeSize = (dtype == ge::DT_FLOAT) ? TYPE_SIZE_FP32 : TYPE_SIZE_FP16_BF16;
	107	+ OP_CHECK_IF(typeSize <= 0, OP_LOGE(context, "typeSize<=0"),
	108	+ return ge::GRAPH_FAILED);
	109	+ alignElem = ubBlockSize / typeSize;
	110	+ OP_CHECK_IF(alignElem <= 0, OP_LOGE(context, "alignElem<=0"),
	111	+ return ge::GRAPH_FAILED);
	112	+
	113	+ if (totalNum < alignElem) {
	114	+ tiling->blockFactor = totalNum;
	115	+ usedCoreNum = 1;
	116	+ } else {
	117	+ int64_t perCoreRaw = Ops::Base::CeilDiv(totalNum, coreNum);
	118	+ tiling->blockFactor = Ops::Base::CeilAlign(perCoreRaw, alignElem);
	119	+ usedCoreNum = Ops::Base::CeilDiv(totalNum, tiling->blockFactor);
	120	+ }
	121	+ OP_CHECK_IF(usedCoreNum == 0, OP_LOGE(context, "usedCoreNum is 0"),
	122	+ return ge::GRAPH_FAILED);
	123	+
	124	+ int64_t availableUb = static_cast<int64_t>(ubSize) - RESERVED_UB;
	125	+ OP_CHECK_IF(availableUb <= 0, OP_LOGE(context, "availableUb<=0"),
	126	+ return ge::GRAPH_FAILED);
	127	+ int64_t bytePerElem = (dtype == ge::DT_FLOAT) ? BYTE_PER_ELEM_FP32 : BYTE_PER_ELEM_FP16_BF16;
	128	+ int64_t tileElem = availableUb / bytePerElem;
	129	+ tileElem = Ops::Base::FloorAlign(tileElem, TILE_ALIGN);
	130	+ if (tileElem < alignElem) {
	131	+ tileElem = alignElem;
	132	+ }
	133	+ tiling->ubFactor = tileElem;
	134	+ OP_LOGD(context, "Ndtri UB plan: ubFactor=%ld, bytePerElem=%ld, estUbUsage=%ldB, reserved=%ldB, available=%ldB",
	135	+ static_cast<long>(tileElem),
	136	+ static_cast<long>(bytePerElem),
	137	+ static_cast<long>(tileElem * bytePerElem),
	138	+ static_cast<long>(RESERVED_UB),
	139	+ static_cast<long>(availableUb));
	140	+ return ge::GRAPH_SUCCESS;
	141	+}
	142	+
	143	+// TilingKey 派发
	144	+static void DispatchTilingKey(
	145	+ gert::TilingContext* context, ge::DataType dtype, int64_t totalNum, int64_t alignElem)
	146	+{
	147	+ uint32_t dtypeKey;
	148	+ if (dtype == ge::DT_FLOAT) {
	149	+ dtypeKey = static_cast<uint32_t>(C_DT_FLOAT);
	150	+ } else if (dtype == ge::DT_FLOAT16) {
	151	+ dtypeKey = static_cast<uint32_t>(C_DT_FLOAT16);
	152	+ } else {
	153	+ dtypeKey = static_cast<uint32_t>(C_DT_BF16);
	154	+ }
	155	+ uint32_t isAlign = (alignElem > 0 && totalNum > 0 && totalNum % alignElem == 0) ? 1U : 0U;
	156	+ ASCENDC_TPL_SEL_PARAM(context, dtypeKey, isAlign);
	157	+}
	158	+
	159	+// Tiling 入口
	160	+static ge::graphStatus NdtriTilingFunc(gert::TilingContext* context)
	161	+{
	162	+ OP_LOGD(context->GetNodeName(), "Enter NdtriTilingFunc");
	163	+ uint64_t ubSize = 0;
	164	+ int64_t coreNum = 0;
	165	+ uint32_t sysWorkspaceSize = 0;
	166	+ OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum, sysWorkspaceSize) != ge::GRAPH_SUCCESS,
	167	+ OP_LOGE(context, "GetPlatformInfo error"),
	168	+ return ge::GRAPH_FAILED);
	169	+
	170	+ ge::DataType dtype;
	171	+ OP_CHECK_IF(CheckDtype(context, dtype) != ge::GRAPH_SUCCESS,
	172	+ OP_LOGE(context, "CheckDtype error"),
	173	+ return ge::GRAPH_FAILED);
	174	+
	175	+ int64_t totalNum = 0;
	176	+ OP_CHECK_IF(GetTotalNum(context, totalNum) != ge::GRAPH_SUCCESS,
	177	+ OP_LOGE(context, "GetTotalNum error"),
	178	+ return ge::GRAPH_FAILED);
	179	+
	180	+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
	181	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	182	+ currentWorkspace[0] = WS_USER_SIZE + sysWorkspaceSize;
	183	+
	184	+ NdtriTilingData* tiling = context->GetTilingData<NdtriTilingData>();
	185	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	186	+ OP_CHECK_IF(memset_s(tiling, sizeof(NdtriTilingData), 0, sizeof(NdtriTilingData)) != EOK,
	187	+ OP_LOGE(context, "set tiling data error"),
	188	+ return ge::GRAPH_FAILED);
	189	+ tiling->totalNum = totalNum;
	190	+
	191	+ int64_t usedCoreNum = 0;
	192	+ int64_t alignElem = 0;
	193	+ OP_CHECK_IF(DoTiling(context, dtype, totalNum, ubSize, coreNum,
	194	+ tiling, usedCoreNum, alignElem) != ge::GRAPH_SUCCESS,
	195	+ OP_LOGE(context, "DoTiling error"),
	196	+ return ge::GRAPH_FAILED);
	197	+
	198	+ context->SetBlockDim(usedCoreNum);
	199	+ DispatchTilingKey(context, dtype, totalNum, alignElem);
	200	+ return ge::GRAPH_SUCCESS;
	201	+}
	202	+
	203	+static ge::graphStatus TilingParseForNdtri(
	204	+ [[maybe_unused]] gert::TilingParseContext* context)
	205	+{
	206	+ return ge::GRAPH_SUCCESS;
	207	+}
	208	+
	209	+struct NdtriCompileInfo {};
	210	+
	211	+IMPL_OP_OPTILING(Ndtri)
	212	+ .Tiling(NdtriTilingFunc)
	213	+ .TilingParse<NdtriCompileInfo>(TilingParseForNdtri);
	214	+
	215	+} // namespace optiling

Amath/ndtri/op_host/ndtri_def.cpp+50-0

@@ -0,0 +1,50 @@
	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 ndtri_def.cpp
	15	+ * \brief Ndtri 算子定义
	16	+ */
	17	+#include "register/op_def_registry.h"
	18	+
	19	+namespace ops {
	20	+class Ndtri : public OpDef {
	21	+public:
	22	+ explicit Ndtri(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	+ .AutoContiguous();
	30	+
	31	+ this->Output("y")
	32	+ .ParamType(REQUIRED)
	33	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
	34	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	35	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	36	+ .AutoContiguous();
	37	+
	38	+ OpAICoreConfig aiCoreConfig;
	39	+ aiCoreConfig.DynamicCompileStaticFlag(true)
	40	+ .DynamicFormatFlag(false)
	41	+ .DynamicRankSupportFlag(true)
	42	+ .DynamicShapeSupportFlag(true)
	43	+ .NeedCheckSupportFlag(false)
	44	+ .PrecisionReduceFlag(true)
	45	+ .ExtendCfgInfo("opFile.value", "ndtri_apt");
	46	+ this->AICore().AddConfig("ascend950", aiCoreConfig);
	47	+ }
	48	+};
	49	+OP_ADD(Ndtri);
	50	+} // namespace ops

Amath/ndtri/op_host/ndtri_infershape.cpp+51-0

@@ -0,0 +1,51 @@
	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 ndtri_infershape.cpp
	15	+ * \brief Ndtri 算子形状/类型推导实现
	16	+ *
	17	+ * out shape == self shape；out dtype == self dtype（逐元素算子）。
	18	+ */
	19	+
	20	+#include "register/op_impl_registry.h"
	21	+#include "exe_graph/runtime/infer_shape_context.h"
	22	+#include "exe_graph/runtime/infer_datatype_context.h"
	23	+#include "log/log.h"
	24	+
	25	+using namespace ge;
	26	+
	27	+namespace ops {
	28	+
	29	+static constexpr size_t IDX_SELF = 0;
	30	+static constexpr size_t OUT_OUT = 0;
	31	+
	32	+static ge::graphStatus InferShape4Ndtri(gert::InferShapeContext* context)
	33	+{
	34	+ const gert::Shape* selfShape = context->GetInputShape(IDX_SELF);
	35	+ OP_CHECK_NULL_WITH_CONTEXT(context, selfShape);
	36	+ gert::Shape* outShape = context->GetOutputShape(OUT_OUT);
	37	+ OP_CHECK_NULL_WITH_CONTEXT(context, outShape);
	38	+ outShape = selfShape;
	39	+ return ge::GRAPH_SUCCESS;
	40	+}
	41	+
	42	+static ge::graphStatus InferDataType4Ndtri(gert::InferDataTypeContext* context)
	43	+{
	44	+ const ge::DataType selfDtype = context->GetInputDataType(IDX_SELF);
	45	+ context->SetOutputDataType(OUT_OUT, selfDtype);
	46	+ return ge::GRAPH_SUCCESS;
	47	+}
	48	+
	49	+IMPL_OP_INFERSHAPE(Ndtri).InferShape(InferShape4Ndtri).InferDataType(InferDataType4Ndtri);
	50	+
	51	+} // namespace ops

Amath/ndtri/op_kernel/arch35/ndtri_coeffs.h+134-0

@@ -0,0 +1,134 @@
	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	+ * Cephes Math Library, Stephen L. Moshier, https://netlib.org/cephes/
	12	+ * Licensed under BSD-like terms. Redistribution with attribution.
	13	+ * Source: cephes/cprob/ndtri.c
	14	+ *
	15	+ * The rational approximation coefficients below are derived from the
	16	+ * Cephes Math Library (double precision). They are converted to single
	17	+ * precision (float) for use on NPU arch35 vector pipeline. The loss of
	18	+ * precision compared to double (~7-8 significant decimal digits) is well
	19	+ * within the FP32 threshold (2^-13 ≈ 1.22e-4).
	20	+ * ------------------------------------------------------------------------
	21	+ */
	22	+
	23	+/**
	24	+ * NOTE: Portions of this code were AI-generated and have been
	25	+ * technically reviewed for functional accuracy and security
	26	+ */
	27	+/*!
	28	+ * \file ndtri_coeffs.h
	29	+ * \brief Cephes Ndtri 分区间有理逼近系数（FP32）。
	30	+ */
	31	+
	32	+#ifndef NDTRI_COEFFS_H_
	33	+#define NDTRI_COEFFS_H_
	34	+
	35	+#include "kernel_operator.h"
	36	+
	37	+namespace NsNdtri {
	38	+
	39	+// ---------------------------------------------------------------
	40	+// 中心区 \|p - 0.5\| <= 0.5 - exp(-2)，z = (p - 0.5)^2
	41	+// Rational approximation: ndtri(p) ≈ sqrt(2π) * ((p-0.5) + (p-0.5)^3 * P0(z)/Q0(z))
	42	+// ---------------------------------------------------------------
	43	+
	44	+// LIST_P0: 5 项（Cephes 约定 polevl：coefs[0] 为最高次 x^(n-1) 的系数）
	45	+__aicore__ constexpr float LIST_P0[5] = {
	46	+ -5.99633501014107895267e1f,
	47	+ 9.80010754185999661536e1f,
	48	+ -5.66762857469070293439e1f,
	49	+ 1.39312609387279679503e1f,
	50	+ -1.23916583867381258016e0f,
	51	+};
	52	+
	53	+// LIST_Q0: 8 项（Cephes 约定 p1evl：x^n 项系数恒为 1，不显式存入；
	54	+// coefs[0] 为 x^(n-1) 的系数，依次降至 coefs[n-1] 为常数项）
	55	+__aicore__ constexpr float LIST_Q0[8] = {
	56	+ 1.95448858338141759834e0f,
	57	+ 4.67627912898881538453e0f,
	58	+ 8.63602421390890590575e1f,
	59	+ -2.25462687854119370527e2f,
	60	+ 2.00260212380060660359e2f,
	61	+ -8.20372256168333339912e1f,
	62	+ 1.59056225126211695515e1f,
	63	+ -1.18331621121330003142e0f,
	64	+};
	65	+
	66	+// ---------------------------------------------------------------
	67	+// 尾部区 P1/Q1：x = sqrt(-2 ln q) ∈ [2, 8]
	68	+// ---------------------------------------------------------------
	69	+
	70	+// LIST_P1: 9 项
	71	+__aicore__ constexpr float LIST_P1[9] = {
	72	+ 4.05544892305962419923e0f,
	73	+ 3.15251094599893866154e1f,
	74	+ 5.71628192246421288162e1f,
	75	+ 4.40805073893200834700e1f,
	76	+ 1.46849561928858024014e1f,
	77	+ 2.18663306850790267539e0f,
	78	+ -1.40256079171354495875e-1f,
	79	+ -3.50424626827848203418e-2f,
	80	+ -8.57456785154685413611e-4f,
	81	+};
	82	+
	83	+// LIST_Q1: 8 项（首项系数为 1，不显式存入）
	84	+__aicore__ constexpr float LIST_Q1[8] = {
	85	+ 1.57799883256466749731e1f,
	86	+ 4.53907635128879210584e1f,
	87	+ 4.13172038254672030440e1f,
	88	+ 1.50425385692907503408e1f,
	89	+ 2.50464946208309415979e0f,
	90	+ -1.42182922854787788574e-1f,
	91	+ -3.80806407691578277194e-2f,
	92	+ -9.33259480895457427372e-4f,
	93	+};
	94	+
	95	+// ---------------------------------------------------------------
	96	+// 尾部区 P2/Q2：x = sqrt(-2 ln q) > 8（极端小概率，p < exp(-32)）
	97	+// ---------------------------------------------------------------
	98	+
	99	+// LIST_P2: 9 项
	100	+__aicore__ constexpr float LIST_P2[9] = {
	101	+ 3.23774891776946035970e0f,
	102	+ 6.91522889068984211695e0f,
	103	+ 3.93881025292474443415e0f,
	104	+ 1.33303460815807542389e0f,
	105	+ 2.01485389549179081538e-1f,
	106	+ 1.23716634817820021358e-2f,
	107	+ 3.01581553508235416007e-4f,
	108	+ 2.65806974686737550832e-6f,
	109	+ 6.23974539184983293730e-9f,
	110	+};
	111	+
	112	+// LIST_Q2: 8 项（首项系数为 1，不显式存入）
	113	+__aicore__ constexpr float LIST_Q2[8] = {
	114	+ 6.02427039364742014255e0f,
	115	+ 3.67983563856160859403e0f,
	116	+ 1.37702099489081330271e0f,
	117	+ 2.16236993594496635890e-1f,
	118	+ 1.34204006088543189037e-2f,
	119	+ 3.28014464682127739104e-4f,
	120	+ 2.89247864745380683936e-6f,
	121	+ 6.79019408009981274425e-9f,
	122	+};
	123	+
	124	+// 边界常量
	125	+__aicore__ constexpr float NDTRI_VAL_SUB = 0.1353352832366127f; // e^-2
	126	+__aicore__ constexpr float NDTRI_RES_EXP = 0.8646647167633873f; // 1 - e^-2
	127	+__aicore__ constexpr float NDTRI_SQRT_2PI = 2.50662827463100050242f; // sqrt(2π)
	128	+__aicore__ constexpr float NDTRI_X_BOUNDARY = 8.0f; // P1/Q1 vs P2/Q2 分界
	129	+__aicore__ constexpr float NDTRI_SAFE_LO = 1.1754944e-38f; // FLT_MIN，pSafe 下限
	130	+__aicore__ constexpr float NDTRI_SAFE_HI = 0.99999994f; // nextafter(1.0f, 0.0f) = 1 - 2^-24
	131	+
	132	+} // namespace NsNdtri
	133	+
	134	+#endif // NDTRI_COEFFS_H_

Amath/ndtri/op_kernel/arch35/ndtri_compute.h+316-0

@@ -0,0 +1,316 @@
	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	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file ndtri_compute.h
	17	+ * \brief Cephes Ndtri 分区间有理逼近的 Tensor 化实现（FP32 域）。
	18	+ *
	19	+ * 模块划分（与详细设计 §4.4 对齐）：
	20	+ * - _polevl(x) : P(x) Horner 多项式
	21	+ * - _plevl(x) : Q(x) = 1 + ... 首项为 1 的变体
	22	+ * - polevl_plevl(x) : P(x) / Q(x) 有理函数
	23	+ * - cal_p0(p) : 中心区 y = sqrt(2π) * pm * (1 + pm^2 * P0(z)/Q0(z))
	24	+ * - cal_sub(q) : 尾部 x = sqrt(-2 ln q)，x0 = x - ln(x)/x
	25	+ * - cal_p12(x) : 尾部修正 1/x * P12(1/x)/Q12(1/x)（按 x<8 / x>=8 掩码合并）
	26	+ * - cal_tail(pSafe) : 尾部 y_tail = sign * (x0 - cal_p12(x))
	27	+ *
	28	+ * 所有函数在 FP32 域工作。输入/输出 LocalTensor 均由调用者（Kernel 主体）管理 UB 分配，
	29	+ * 本文件仅做计算逻辑组合。
	30	+ */
	31	+
	32	+#ifndef NDTRI_COMPUTE_H_
	33	+#define NDTRI_COMPUTE_H_
	34	+
	35	+#include "kernel_operator.h"
	36	+#include "ndtri_coeffs.h"
	37	+
	38	+namespace NsNdtri {
	39	+
	40	+using namespace AscendC;
	41	+
	42	+// ---------------------------------------------------------------
	43	+// Cephes polevl 约定（coefs[0] 为最高次项系数）：
	44	+// P(x) = coefs[0]x^(n-1) + coefs[1]x^(n-2) + ... + coefs[n-1]
	45	+// Horner: ans = coefs[0]; for i=1..n-1: ans = ans*x + coefs[i]
	46	+//
	47	+// 设计说明：
	48	+// 保留语义最清晰的原始 Mul+Adds 写法，交由编译器自动融合。
	49	+// ---------------------------------------------------------------
	50	+__aicore__ inline void PolEvl(
	51	+ const LocalTensor<float>& dst,
	52	+ const LocalTensor<float>& x,
	53	+ const float* coefs, int n,
	54	+ const LocalTensor<float>& scratch,
	55	+ int32_t len)
	56	+{
	57	+ // dst = coefs[0]（最高次）
	58	+ Duplicate(dst, coefs[0], len);
	59	+ for (int i = 1; i < n; ++i) {
	60	+ // scratch = dst * x
	61	+ Mul(scratch, dst, x, len);
	62	+ // dst = scratch + coefs[i]
	63	+ Adds(dst, scratch, coefs[i], len);
	64	+ }
	65	+}
	66	+
	67	+// ---------------------------------------------------------------
	68	+// Cephes p1evl 约定（首项系数为 1，不显式存入 coefs）：
	69	+// Q(x) = x^n + coefs[0]x^(n-1) + coefs[1]x^(n-2) + ... + coefs[n-1]
	70	+// 等价 Horner: ans = 1; for i=0..n-1: ans = ans*x + coefs[i]
	71	+// （FMA 自动融合同 PolEvl。）
	72	+// ---------------------------------------------------------------
	73	+__aicore__ inline void PlEvl(
	74	+ const LocalTensor<float>& dst,
	75	+ const LocalTensor<float>& x,
	76	+ const float* coefs, int n,
	77	+ const LocalTensor<float>& scratch,
	78	+ int32_t len)
	79	+{
	80	+ // dst = 1.0（隐式 x^n 项系数）
	81	+ Duplicate(dst, 1.0f, len);
	82	+ for (int i = 0; i < n; ++i) {
	83	+ Mul(scratch, dst, x, len);
	84	+ Adds(dst, scratch, coefs[i], len);
	85	+ }
	86	+}
	87	+
	88	+// ---------------------------------------------------------------
	89	+// 有理函数 R(x) = P(x) / Q(x)（P 无首项约束，Q 首项为 1）
	90	+// - tmpP / tmpQ: 存 P(x) / Q(x) 中间结果
	91	+// - scratch : Horner 内部 scratch
	92	+// ---------------------------------------------------------------
	93	+__aicore__ inline void PolEvlPlEvl(
	94	+ const LocalTensor<float>& dst,
	95	+ const LocalTensor<float>& x,
	96	+ const float* coefsP, int nP,
	97	+ const float* coefsQ, int nQ,
	98	+ const LocalTensor<float>& tmpP,
	99	+ const LocalTensor<float>& tmpQ,
	100	+ const LocalTensor<float>& scratch,
	101	+ int32_t len)
	102	+{
	103	+ PolEvl(tmpP, x, coefsP, nP, scratch, len);
	104	+ PlEvl (tmpQ, x, coefsQ, nQ, scratch, len);
	105	+ Div(dst, tmpP, tmpQ, len);
	106	+}
	107	+
	108	+// ---------------------------------------------------------------
	109	+// cal_p0: 中心区
	110	+// y = sqrt(2π) * (pm + pm^3 * R(z))
	111	+// = sqrt(2π) * pm * (1 + z * R(z))
	112	+// 其中 pm = p - 0.5, z = pm^2
	113	+//
	114	+// Buffer 约定（由调用者传入，大小 = len * sizeof(float)）：
	115	+// - y : 输出
	116	+// - p : 输入
	117	+// - tmpPm : pm 中间（可以复用 y 做输入→输出 inplace，不推荐；保持独立更清晰）
	118	+// - tmpZ : z 中间
	119	+// - tmpP : P0(z) 结果
	120	+// - tmpQ : Q0(z) 结果
	121	+// - scratch : Horner scratch
	122	+// ---------------------------------------------------------------
	123	+__aicore__ inline void CalP0(
	124	+ const LocalTensor<float>& y,
	125	+ const LocalTensor<float>& p,
	126	+ const LocalTensor<float>& tmpPm,
	127	+ const LocalTensor<float>& tmpZ,
	128	+ const LocalTensor<float>& tmpP,
	129	+ const LocalTensor<float>& tmpQ,
	130	+ const LocalTensor<float>& scratch,
	131	+ int32_t len)
	132	+{
	133	+ // pm = p - 0.5
	134	+ Adds(tmpPm, p, -0.5f, len);
	135	+
	136	+ // z = pm * pm
	137	+ Mul(tmpZ, tmpPm, tmpPm, len);
	138	+
	139	+ // R = P0(z) / Q0(z)
	140	+ PolEvlPlEvl(y, tmpZ, LIST_P0, 5, LIST_Q0, 8, tmpP, tmpQ, scratch, len);
	141	+ // y 临时存 R(z)
	142	+
	143	+ // y = z * R
	144	+ Mul(y, y, tmpZ, len);
	145	+ // y = 1 + z * R
	146	+ Adds(y, y, 1.0f, len);
	147	+ // y = pm * (1 + z * R)
	148	+ Mul(y, y, tmpPm, len);
	149	+ // y = sqrt(2π) * y
	150	+ Muls(y, y, NDTRI_SQRT_2PI, len);
	151	+}
	152	+
	153	+// ---------------------------------------------------------------
	154	+// cal_sub: 尾部基础
	155	+// x = sqrt(-2 ln q)
	156	+// x0 = x - ln(x) / x
	157	+//
	158	+// 输入 q ∈ (0, e^-2]（由调用者在 cal_tail 中通过 q = select(mask_neg, 1 - pSafe, pSafe) 保证），
	159	+// pSafe 钳制已确保 q > 0。
	160	+//
	161	+// Buffer 约定：
	162	+// - x0 : 输出 x0
	163	+// - xOut : 输出 x（供 cal_p12 使用）
	164	+// - q : 输入
	165	+// - tmp : 工作 buffer
	166	+// ---------------------------------------------------------------
	167	+__aicore__ inline void CalSub(
	168	+ const LocalTensor<float>& x0,
	169	+ const LocalTensor<float>& xOut,
	170	+ const LocalTensor<float>& q,
	171	+ const LocalTensor<float>& tmp,
	172	+ int32_t len)
	173	+{
	174	+ // tmp = ln(q)
	175	+ Ln(tmp, q, len);
	176	+ // tmp = -2 * ln(q)
	177	+ Muls(tmp, tmp, -2.0f, len);
	178	+ // xOut = sqrt(-2 ln q)
	179	+ Sqrt(xOut, tmp, len);
	180	+ // tmp = ln(xOut)
	181	+ Ln(tmp, xOut, len);
	182	+ // tmp = ln(x) / x
	183	+ Div(tmp, tmp, xOut, len);
	184	+ // x0 = x - ln(x) / x
	185	+ Sub(x0, xOut, tmp, len);
	186	+}
	187	+
	188	+// ---------------------------------------------------------------
	189	+// cal_p12: 尾部修正
	190	+// z = 1 / x
	191	+// r1 = z * P1(z) / Q1(z) （对 x < 8 使用）
	192	+// r2 = z * P2(z) / Q2(z) （对 x >= 8 使用）
	193	+// corr = select(x < 8, r1, r2)
	194	+//
	195	+// Buffer 约定：
	196	+// - corr : 输出
	197	+// - x : 输入 x = sqrt(-2 ln q)
	198	+// - tmpZ : z = 1/x
	199	+// - tmpR1 : r1 = P1(z)/Q1(z)
	200	+// - tmpR2 : r2 = P2(z)/Q2(z)
	201	+// - tmpP : Horner 多项式 P(x) 结果
	202	+// - tmpQ : Horner 多项式 Q(x) 结果
	203	+// - scratch : Horner scratch
	204	+// - maskX : uint8 mask buffer
	205	+//
	206	+// ISSUE-001：调用者传入的 len 必须是 64 倍数（FP32 下 256B 对齐），
	207	+// 由 Kernel 层的 lenAligned 保证。
	208	+// ---------------------------------------------------------------
	209	+__aicore__ inline void CalP12(
	210	+ const LocalTensor<float>& corr,
	211	+ const LocalTensor<float>& x,
	212	+ const LocalTensor<float>& tmpZ,
	213	+ const LocalTensor<float>& tmpR1,
	214	+ const LocalTensor<float>& tmpR2,
	215	+ const LocalTensor<float>& tmpP,
	216	+ const LocalTensor<float>& tmpQ,
	217	+ const LocalTensor<float>& scratch,
	218	+ const LocalTensor<uint8_t>& maskX,
	219	+ int32_t len)
	220	+{
	221	+ // z = 1 / x => tmpZ = 1.0, tmpZ /= x
	222	+ Duplicate(tmpZ, 1.0f, len);
	223	+ Div(tmpZ, tmpZ, x, len);
	224	+
	225	+ // r1 = P1(z) / Q1(z)
	226	+ PolEvlPlEvl(tmpR1, tmpZ, LIST_P1, 9, LIST_Q1, 8, tmpP, tmpQ, scratch, len);
	227	+
	228	+ // r2 = P2(z) / Q2(z)
	229	+ PolEvlPlEvl(tmpR2, tmpZ, LIST_P2, 9, LIST_Q2, 8, tmpP, tmpQ, scratch, len);
	230	+
	231	+ // mask: x < 8 -> 选 r1，否则 r2
	232	+ CompareScalar(maskX, x, NDTRI_X_BOUNDARY, CMPMODE::LT, len);
	233	+
	234	+ // corr_raw = select(mask, r1, r2)
	235	+ Select(corr, maskX, tmpR1, tmpR2,
	236	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, len);
	237	+
	238	+ // corr = z * corr_raw
	239	+ Mul(corr, corr, tmpZ, len);
	240	+}
	241	+
	242	+// ---------------------------------------------------------------
	243	+// cal_tail: 尾部完整流程
	244	+// q = select(maskNeg, 1 - pSafe, pSafe)
	245	+// x = sqrt(-2 ln q)
	246	+// x0 = x - ln(x)/x
	247	+// corr = cal_p12(x)
	248	+// base = x0 - corr （Cephes 源码 x0 -= ...）
	249	+// y_tail = select(maskNeg, +base, -base)
	250	+//
	251	+// Buffer 约定：
	252	+// - yTail : 输出
	253	+// - pSafe : 输入（已经 clamp 到 [FLT_MIN, 1-FLT_MIN]）
	254	+// - maskNeg : p >= 0.5 的掩码
	255	+// - tmpQ : q 中间（复用为 "1 - pSafe"）
	256	+// - tmpX : x
	257	+// - tmpX0 : x0
	258	+// - tmpCorr : corr
	259	+// - tmp1..5 : 5 个 fp32 scratch buffer（供 cal_sub / cal_p12 使用）
	260	+// - maskX : uint8 scratch mask
	261	+// ---------------------------------------------------------------
	262	+__aicore__ inline void CalTail(
	263	+ const LocalTensor<float>& yTail,
	264	+ const LocalTensor<float>& pSafe,
	265	+ const LocalTensor<uint8_t>& maskNeg,
	266	+ const LocalTensor<float>& tmpQ,
	267	+ const LocalTensor<float>& tmpX,
	268	+ const LocalTensor<float>& tmpX0,
	269	+ const LocalTensor<float>& tmpCorr,
	270	+ const LocalTensor<float>& tmp1, // cal_sub 的 tmp / cal_p12 的 tmpZ
	271	+ const LocalTensor<float>& tmp2, // cal_p12 的 tmpR1
	272	+ const LocalTensor<float>& tmp3, // cal_p12 的 tmpR2
	273	+ const LocalTensor<float>& tmp4, // cal_p12 的 tmpP
	274	+ const LocalTensor<float>& tmp5, // cal_p12 的 tmpQ / cal_p12 的 scratch
	275	+ const LocalTensor<uint8_t>& maskX,
	276	+ int32_t len)
	277	+{
	278	+ // Step 1: q = select(maskNeg, 1 - pSafe, pSafe)
	279	+ // oneMinusP = 1 - pSafe
	280	+ Muls(tmpQ, pSafe, -1.0f, len);
	281	+ Adds(tmpQ, tmpQ, 1.0f, len);
	282	+ // Select: maskNeg=1 -> tmpQ (1-pSafe), maskNeg=0 -> pSafe
	283	+ Select(tmpQ, maskNeg, tmpQ, pSafe,
	284	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, len);
	285	+
	286	+ // Step 2: x = sqrt(-2 ln q)，x0 = x - ln(x)/x
	287	+ // CalSub 使用 tmp1 作为工作 buffer
	288	+ CalSub(tmpX0, tmpX, tmpQ, tmp1, len);
	289	+
	290	+ // Step 3: corr = cal_p12(x)
	291	+ // CalP12 内部 Horner 需要 tmpP / tmpQ / scratch：复用 tmp4 / tmp5 / tmpQ
	292	+ // 注意：tmpQ 在此时已经不再需要（q 在 Step 2 中已经消费）
	293	+ CalP12(tmpCorr, tmpX,
	294	+ /tmpZ /tmp1,
	295	+ /tmpR1 /tmp2,
	296	+ /tmpR2 /tmp3,
	297	+ /tmpP /tmp4,
	298	+ /tmpQ /tmp5,
	299	+ /scratch/tmpQ,
	300	+ maskX, len);
	301	+
	302	+ // Step 4: base = x0 - corr
	303	+ Sub(tmpX0, tmpX0, tmpCorr, len);
	304	+
	305	+ // Step 5: sign：
	306	+ // - p < 0.5 (maskNeg=0) -> y_tail = -base
	307	+ // - p >= 0.5 (maskNeg=1) -> y_tail = +base
	308	+ Muls(tmpCorr, tmpX0, -1.0f, len); // -base 存 tmpCorr
	309	+ // Select: maskNeg=1 -> +base, maskNeg=0 -> -base
	310	+ Select(yTail, maskNeg, tmpX0, tmpCorr,
	311	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, len);
	312	+}
	313	+
	314	+} // namespace NsNdtri
	315	+
	316	+#endif // NDTRI_COMPUTE_H_

Amath/ndtri/op_kernel/arch35/ndtri_kernel.h+495-0

@@ -0,0 +1,495 @@
	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	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file ndtri_kernel.h
	17	+ * \brief Ndtri Kernel 实现（arch35 / Ascend950）
	18	+ *
	19	+ * 公式：
	20	+ * y = ndtri(p) = sqrt(2) * erf^{-1}(2p - 1)
	21	+ * 对齐 PyTorch torch.special.ndtri / SciPy scipy.special.ndtri（Cephes 算法）。
	22	+ *
	23	+ * 计算流（单 tile）：
	24	+ * Step 0: 输入 Cast→fp32 (fp16/bf16: CAST_NONE, fp32: ReinterpretCast)
	25	+ * Step 1: 构造 maskTail / maskNeg / maskSpecial
	26	+ * Step 2: pSafe = clamp(p, FLT_MIN, 1 - 2^-24)
	27	+ * Step 3: yTail = cal_tail(pSafe, maskNeg)
	28	+ * Step 4: yCenter = cal_p0(pSafe)
	29	+ * Step 5: y = select(maskTail, yTail, yCenter)
	30	+ * Step 6: y = select(maskSpecial, ySpecial, y)
	31	+ * Step 7: 输出 Cast→T (fp16/bf16: CAST_RINT, fp32: ReinterpretCast)
	32	+ *
	33	+ * 迭代二范围（本次整合）：
	34	+ * - FP32 / FP16 / BF16 × 对齐/非对齐共 6 个 TilingKey 真实实现
	35	+ * - FP16 路径：Cast fp16→fp32 → 统一算法 → Cast fp32→fp16 (穿刺 P-2 已验证 bit-exact)
	36	+ * - BF16 路径：Cast bf16→fp32 → 统一算法 → Cast fp32→bf16 (穿刺 P-3 已验证 bit-exact)
	37	+ * - 非对齐路径：DataCopyPad 处理尾块（本来就用 DataCopyPad，天然兼容）
	38	+ *
	39	+ * TilingKey 矩阵：{fp32, fp16, bf16} × {对齐, 非对齐} = 6 个
	40	+ */
	41	+
	42	+#ifndef NDTRI_KERNEL_H
	43	+#define NDTRI_KERNEL_H
	44	+
	45	+#include "kernel_operator.h"
	46	+#include "kernel_tiling/kernel_tiling.h"
	47	+#include "ndtri_tiling_data.h"
	48	+#include "ndtri_tiling_key.h"
	49	+#include "ndtri_coeffs.h"
	50	+#include "ndtri_compute.h"
	51	+
	52	+namespace NsNdtri {
	53	+
	54	+using namespace AscendC;
	55	+
	56	+template <typename T, int K_ALIGN>
	57	+class Ndtri {
	58	+ static constexpr int32_t BUFFER_NUM = 2;
	59	+ static constexpr bool IS_FP32 = AscendC::IsSameType<T, float>::value;
	60	+
	61	+ // Compare/CompareScalar API 对 `count` 所占空间有 256B 对齐硬约束（见
	62	+ // ascendc-api-best-practices §2.1）。FP32 下等价 count 为 64 元素倍数。
	63	+ // 稳态 tile（ubFactor）已由 Host Tiling 的 FloorAlign(_, 256) 对齐；
	64	+ // 尾块 currentNum 可能任意非 64 倍数，需在 Kernel 层向上取 64 对齐并在
	65	+ // padding 区域填入中性值 0.5f（中心区值，既不触发 tail 分支也不触发
	66	+ // special 分支），由 DataCopyPad 的 blockLen=currentNum*sizeof(T) 保证
	67	+ // padding 位置最终不会被写回 GM。
	68	+ static constexpr int32_t CMP_ALIGN_ELEM = 64;
	69	+ __aicore__ inline static int32_t AlignCmpLen(int32_t len)
	70	+ {
	71	+ return (len + CMP_ALIGN_ELEM - 1) / CMP_ALIGN_ELEM * CMP_ALIGN_ELEM;
	72	+ }
	73	+
	74	+public:
	75	+ __aicore__ inline Ndtri() = default;
	76	+
	77	+ __aicore__ inline void Init(
	78	+ GM_ADDR self, GM_ADDR out,
	79	+ const NdtriTilingData* tilingData);
	80	+
	81	+ __aicore__ inline void Process();
	82	+
	83	+private:
	84	+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
	85	+ __aicore__ inline void Compute(int64_t currentNum);
	86	+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
	87	+
	88	+ // 构造三个 mask：tail / neg / special
	89	+ __aicore__ inline void BuildMasks(
	90	+ const LocalTensor<float>& p,
	91	+ const LocalTensor<uint8_t>& maskTail,
	92	+ const LocalTensor<uint8_t>& maskNeg,
	93	+ const LocalTensor<uint8_t>& maskSpecial,
	94	+ const LocalTensor<float>& scratch,
	95	+ int32_t len);
	96	+
	97	+ // 构造 y_special：p==0 -> -inf, p==1 -> +inf, otherwise -> NaN
	98	+ __aicore__ inline void BuildSpecialY(
	99	+ const LocalTensor<float>& ySpecial,
	100	+ const LocalTensor<float>& p,
	101	+ const LocalTensor<float>& scratch,
	102	+ int32_t len);
	103	+
	104	+private:
	105	+ TPipe pipe;
	106	+
	107	+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueSelf;
	108	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueY;
	109	+
	110	+ // fp32 域：p / y（所有 dtype 统一使用，FP32 路径通过 Cast(CAST_NONE) 等价搬运）
	111	+ TBuf<TPosition::VECCALC> pBuf;
	112	+ TBuf<TPosition::VECCALC> yBuf;
	113	+
	114	+ // fp32 计算 Buffer
	115	+ TBuf<TPosition::VECCALC> tmpBuf0;
	116	+ TBuf<TPosition::VECCALC> tmpBuf1;
	117	+ TBuf<TPosition::VECCALC> tmpBuf2;
	118	+ TBuf<TPosition::VECCALC> tmpBuf3;
	119	+ TBuf<TPosition::VECCALC> tmpBuf4;
	120	+ TBuf<TPosition::VECCALC> tmpBuf5;
	121	+ TBuf<TPosition::VECCALC> tmpBuf6;
	122	+ TBuf<TPosition::VECCALC> tmpBuf7;
	123	+ TBuf<TPosition::VECCALC> tmpBuf8;
	124	+ TBuf<TPosition::VECCALC> tmpBuf9; // CalTail tmp4（避免与 q 别名）
	125	+ TBuf<TPosition::VECCALC> tmpBuf10; // CalTail tmp5（避免与 x 别名）
	126	+
	127	+ // mask buffer（uint8）
	128	+ TBuf<TPosition::VECCALC> maskBuf0;
	129	+ TBuf<TPosition::VECCALC> maskBuf1;
	130	+ TBuf<TPosition::VECCALC> maskBuf2;
	131	+ TBuf<TPosition::VECCALC> maskBuf3; // scratch mask for cal_p12 / BuildSpecialY
	132	+
	133	+ GlobalTensor<T> selfGm;
	134	+ GlobalTensor<T> outGm;
	135	+
	136	+ int64_t blockLength_ = 0;
	137	+ int64_t ubLength_ = 0;
	138	+};
	139	+
	140	+// ---------------------------------------------------------------
	141	+// Init
	142	+// ---------------------------------------------------------------
	143	+template <typename T, int K_ALIGN>
	144	+__aicore__ inline void Ndtri<T, K_ALIGN>::Init(
	145	+ GM_ADDR self, GM_ADDR out,
	146	+ const NdtriTilingData* tilingData)
	147	+{
	148	+ int64_t blockIdx = AscendC::GetBlockIdx();
	149	+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * blockIdx;
	150	+ blockLength_ = (remainderLength > tilingData->blockFactor) ?
	151	+ tilingData->blockFactor : remainderLength;
	152	+ if (blockLength_ < 0) {
	153	+ blockLength_ = 0;
	154	+ }
	155	+ ubLength_ = tilingData->ubFactor;
	156	+ if (ubLength_ <= 0) {
	157	+ ubLength_ = 1;
	158	+ }
	159	+
	160	+ int64_t offset = tilingData->blockFactor * blockIdx;
	161	+ selfGm.SetGlobalBuffer((__gm__ T*)self + offset, blockLength_);
	162	+ outGm.SetGlobalBuffer((__gm__ T*)out + offset, blockLength_);
	163	+
	164	+ // InQue / OutQue（DB）
	165	+ pipe.InitBuffer(inQueSelf, BUFFER_NUM, ubLength_ * sizeof(T));
	166	+ pipe.InitBuffer(outQueY, BUFFER_NUM, ubLength_ * sizeof(T));
	167	+
	168	+ // fp32 域 p / y buffer（独立分配，保证 fp16/bf16 Cast 链路有足够空间）
	169	+ pipe.InitBuffer(pBuf, ubLength_ * sizeof(float));
	170	+ pipe.InitBuffer(yBuf, ubLength_ * sizeof(float));
	171	+
	172	+ // fp32 工作 Buffer
	173	+ pipe.InitBuffer(tmpBuf0, ubLength_ * sizeof(float));
	174	+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(float));
	175	+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(float));
	176	+ pipe.InitBuffer(tmpBuf3, ubLength_ * sizeof(float));
	177	+ pipe.InitBuffer(tmpBuf4, ubLength_ * sizeof(float));
	178	+ pipe.InitBuffer(tmpBuf5, ubLength_ * sizeof(float));
	179	+ pipe.InitBuffer(tmpBuf6, ubLength_ * sizeof(float));
	180	+ pipe.InitBuffer(tmpBuf7, ubLength_ * sizeof(float));
	181	+ pipe.InitBuffer(tmpBuf8, ubLength_ * sizeof(float));
	182	+ pipe.InitBuffer(tmpBuf9, ubLength_ * sizeof(float));
	183	+ pipe.InitBuffer(tmpBuf10, ubLength_ * sizeof(float));
	184	+
	185	+ // uint8 mask：按 bit 存储，大小 = ceil(len/8)；保守分配 len/8 + 32 字节冗余
	186	+ // uint8 mask 按 bit 存储；+32B 余量保证 mask 写操作（按 256bit/cycle 批量写）
	187	+ // 末尾即使越过 logical 边界也不踩到下一段 buffer。32B = 一条 V 写入的最小粒度。
	188	+ int64_t maskBytes = (ubLength_ + 7) / 8 + 32;
	189	+ pipe.InitBuffer(maskBuf0, maskBytes);
	190	+ pipe.InitBuffer(maskBuf1, maskBytes);
	191	+ pipe.InitBuffer(maskBuf2, maskBytes);
	192	+ pipe.InitBuffer(maskBuf3, maskBytes);
	193	+}
	194	+
	195	+// ---------------------------------------------------------------
	196	+// Process
	197	+// ---------------------------------------------------------------
	198	+template <typename T, int K_ALIGN>
	199	+__aicore__ inline void Ndtri<T, K_ALIGN>::Process()
	200	+{
	201	+ if (blockLength_ <= 0) {
	202	+ return;
	203	+ }
	204	+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
	205	+ for (int64_t i = 0; i < loopCount; ++i) {
	206	+ int64_t currentNum = (i == loopCount - 1) ?
	207	+ (blockLength_ - ubLength_ * i) : ubLength_;
	208	+ CopyIn(i, currentNum);
	209	+ Compute(currentNum);
	210	+ CopyOut(i, currentNum);
	211	+ }
	212	+}
	213	+
	214	+// ---------------------------------------------------------------
	215	+// CopyIn
	216	+// DataCopyPad 天然支持对齐/非对齐两种路径，K_ALIGN 仅影响 Host 的 TilingKey 派发。
	217	+// ---------------------------------------------------------------
	218	+template <typename T, int K_ALIGN>
	219	+__aicore__ inline void Ndtri<T, K_ALIGN>::CopyIn(int64_t progress, int64_t currentNum)
	220	+{
	221	+ LocalTensor<T> inLocal = inQueSelf.template AllocTensor<T>();
	222	+ DataCopyExtParams copyParams{
	223	+ 1, static_cast<uint32_t>(currentNum * sizeof(T)), 0, 0, 0};
	224	+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
	225	+ int64_t gmOffset = progress * ubLength_;
	226	+ DataCopyPad(inLocal, selfGm[gmOffset], copyParams, padParams);
	227	+ inQueSelf.EnQue(inLocal);
	228	+}
	229	+
	230	+// ---------------------------------------------------------------
	231	+// BuildMasks
	232	+// - maskSpecial = (p <= 0) \| (p >= 1) \| (p != p) [NaN 也视为特殊值]
	233	+// - maskTail = \|p - 0.5\| >= (0.5 - VAL_SUB)
	234	+// <=> p <= VAL_SUB or p >= RES_EXP
	235	+// - maskNeg = p >= 0.5
	236	+//
	237	+// NaN 处理：CompareScalar(NaN, x, any) 返回 false；因此单纯 (p<=0)\|(p>=1)
	238	+// 不会把 NaN 放入 maskSpecial。使用 CompareScalar(p, p, NE) 等价 isnan：
	239	+// 浮点 NaN 有 NaN != NaN，正常数有 x == x。将 isnan 与 (p<=0)\|(p>=1) 合并。
	240	+//
	241	+// ISSUE-001：调用者传入的 len 必须是 64 倍数（FP32 下 256B 对齐），
	242	+// 由 Compute 层的 lenAligned 保证。
	243	+// ---------------------------------------------------------------
	244	+template <typename T, int K_ALIGN>
	245	+__aicore__ inline void Ndtri<T, K_ALIGN>::BuildMasks(
	246	+ const LocalTensor<float>& p,
	247	+ const LocalTensor<uint8_t>& maskTail,
	248	+ const LocalTensor<uint8_t>& maskNeg,
	249	+ const LocalTensor<uint8_t>& maskSpecial,
	250	+ const LocalTensor<float>& scratch,
	251	+ int32_t len)
	252	+{
	253	+ // maskTail = (\|p - 0.5\| >= 0.5 - VAL_SUB)
	254	+ Adds(scratch, p, -0.5f, len);
	255	+ Abs(scratch, scratch, len);
	256	+ CompareScalar(maskTail, scratch,
	257	+ 0.5f - NDTRI_VAL_SUB, CMPMODE::GE, len);
	258	+
	259	+ // maskSpecial：由 3 个条件组合：
	260	+ // c1 = (p <= 0)
	261	+ // c2 = (p >= 1)
	262	+ // c3 = isnan(p) → 通过 (p != p) 判断（等价：Compare(p, p, NE)）
	263	+ // 但 AscendC 的 CompareScalar 不支持 src1 == src2 同一张量的比较（要求不同 tensor）。
	264	+ //
	265	+ // NaN 检测策略：
	266	+ // 通过 scratch = Mins(Maxs(p, -INF), +INF) 把所有有限/无穷值映射到自身，
	267	+ // 并把 NaN 映射为 -INF（arch35 上 Maxs/Mins 对 NaN 采用 IEEE 754 "minNumber/maxNumber" 语义：
	268	+ // 当任一操作数为 NaN 时返回非 NaN 操作数）。然后 Compare(p, scratch, NE) 即可捕获 NaN。
	269	+ // 该行为在 arch35 IEEE 754 兼容模式下硬件保证。已通过 docs/precision-report.md 的
	270	+ // NaN 用例（boundary_nan_*）验证。
	271	+ //
	272	+ // 本实现策略：先在 maskSpecial 存 (p <= 0)，然后 OR 进 (p >= 1)（借 maskNeg 中转），
	273	+ // 最后 OR 进 isnan（借 maskNeg 中转），最后重建 maskNeg。
	274	+
	275	+ constexpr uint32_t NEG_INF_BITS_U = 0xFF800000U;
	276	+ constexpr uint32_t POS_INF_BITS_U = 0x7F800000U;
	277	+ float negInf, posInf;
	278	+ {
	279	+ union { uint32_t u; float f; } cvt;
	280	+ cvt.u = NEG_INF_BITS_U; negInf = cvt.f;
	281	+ cvt.u = POS_INF_BITS_U; posInf = cvt.f;
	282	+ }
	283	+
	284	+ // maskSpecial = (p <= 0)
	285	+ CompareScalar(maskSpecial, p, 0.0f, CMPMODE::LE, len);
	286	+
	287	+ // maskNeg_tmp = (p >= 1)
	288	+ CompareScalar(maskNeg, p, 1.0f, CMPMODE::GE, len);
	289	+ Or(maskSpecial, maskSpecial, maskNeg, len);
	290	+
	291	+ // isnan(p)：scratch = Mins(Maxs(p, -INF), +INF) （NaN 经此回填为非 NaN）
	292	+ // 然后 maskNeg_tmp = Compare(p, scratch, NE) → 对 NaN 位置置 1
	293	+ Maxs(scratch, p, negInf, len);
	294	+ Mins(scratch, scratch, posInf, len);
	295	+ Compare(maskNeg, p, scratch, CMPMODE::NE, len);
	296	+ Or(maskSpecial, maskSpecial, maskNeg, len);
	297	+
	298	+ // 重建 maskNeg = (p >= 0.5)
	299	+ CompareScalar(maskNeg, p, 0.5f, CMPMODE::GE, len);
	300	+}
	301	+
	302	+// ---------------------------------------------------------------
	303	+// BuildSpecialY
	304	+// - p == 0 → -inf
	305	+// - p == 1 → +inf
	306	+// - 其他（包括 NaN / p < 0 / p > 1） → NaN
	307	+// ---------------------------------------------------------------
	308	+template <typename T, int K_ALIGN>
	309	+__aicore__ inline void Ndtri<T, K_ALIGN>::BuildSpecialY(
	310	+ const LocalTensor<float>& ySpecial,
	311	+ const LocalTensor<float>& p,
	312	+ const LocalTensor<float>& scratch,
	313	+ int32_t len)
	314	+{
	315	+ constexpr uint32_t NAN_BITS = 0x7FC00000U;
	316	+ constexpr uint32_t POS_INF_BITS = 0x7F800000U;
	317	+ constexpr uint32_t NEG_INF_BITS = 0xFF800000U;
	318	+
	319	+ float nanVal, posInf, negInf;
	320	+ {
	321	+ union { uint32_t u; float f; } cvt;
	322	+ cvt.u = NAN_BITS; nanVal = cvt.f;
	323	+ cvt.u = POS_INF_BITS; posInf = cvt.f;
	324	+ cvt.u = NEG_INF_BITS; negInf = cvt.f;
	325	+ }
	326	+
	327	+ // ISSUE-001：调用者传入的 len 由 Compute 层已对齐到 256B（64 元素）。
	328	+
	329	+ // 默认 NaN
	330	+ Duplicate(ySpecial, nanVal, len);
	331	+
	332	+ // 临时复用 maskBuf3 scratch 存 mask_eq0 / mask_eq1
	333	+ LocalTensor<uint8_t> maskEq = maskBuf3.Get<uint8_t>();
	334	+
	335	+ // p == 0 -> -inf
	336	+ CompareScalar(maskEq, p, 0.0f, CMPMODE::EQ, len);
	337	+ Duplicate(scratch, negInf, len);
	338	+ Select(ySpecial, maskEq, scratch, ySpecial,
	339	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, len);
	340	+
	341	+ // p == 1 -> +inf
	342	+ CompareScalar(maskEq, p, 1.0f, CMPMODE::EQ, len);
	343	+ Duplicate(scratch, posInf, len);
	344	+ Select(ySpecial, maskEq, scratch, ySpecial,
	345	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, len);
	346	+}
	347	+
	348	+// ---------------------------------------------------------------
	349	+// Compute
	350	+// ---------------------------------------------------------------
	351	+template <typename T, int K_ALIGN>
	352	+__aicore__ inline void Ndtri<T, K_ALIGN>::Compute(int64_t currentNum)
	353	+{
	354	+ LocalTensor<T> inLocal = inQueSelf.template DeQue<T>();
	355	+ LocalTensor<T> outLocal = outQueY.template AllocTensor<T>();
	356	+ int32_t len = static_cast<int32_t>(currentNum);
	357	+ // ISSUE-001：Compare/CompareScalar API 对 count 所占空间要求 256B 对齐
	358	+ // （FP32: 64 元素倍数）。尾块 currentNum 可能任意非 64 倍数，此处向上
	359	+ // 取 64 对齐；padding 区域 [len, lenAligned) 需要在 Step 0 填入中性值
	360	+ // 0.5f，以保证 BuildMasks / BuildSpecialY 在 padding 位置计算结果不会
	361	+ // 触发越界异常或产生 NaN/Inf 干扰流水线。padding 位置的输出不会被
	362	+ // CopyOut 写回 GM（DataCopyPad.blockLen = len * sizeof(T)）。
	363	+ int32_t lenAligned = AlignCmpLen(len);
	364	+
	365	+ // fp32 域别名
	366	+ LocalTensor<float> p = pBuf.Get<float>();
	367	+ LocalTensor<float> y = yBuf.Get<float>();
	368	+
	369	+ // ISSUE-001：先对 p 的整个对齐区间 [0, lenAligned) 填入中性值 0.5f，
	370	+ // 然后再用有效数据覆盖前 len 个位置。这样避免 LocalTensor 切片
	371	+ // `p[len]` 可能产生的非 32B 对齐起始地址问题（Duplicate 对起始地址
	372	+ // 有 32B 对齐要求）。中性值 0.5f 在中心区，既不触发 tail 分支也不
	373	+ // 触发 special 分支。
	374	+ if (lenAligned > len) {
	375	+ Duplicate(p, 0.5f, lenAligned);
	376	+ }
	377	+
	378	+ // Step 0: 输入 Cast → fp32
	379	+ // - FP32: ReinterpretCast 零拷贝（inLocal 本身就是 fp32 bit pattern）
	380	+ // 但为了保持与 fp16/bf16 路径的语义一致（p 是独立 fp32 buffer），
	381	+ // 此处统一使用 Cast(CAST_NONE)。对于 float→float，Cast(CAST_NONE) 等价于 Copy。
	382	+ // - FP16: Cast(fp16→fp32, CAST_NONE) 无损
	383	+ // - BF16: Cast(bf16→fp32, CAST_NONE) 无损
	384	+ if constexpr (IS_FP32) {
	385	+ // float → float：用 Muls(x, 1.0f) 拷贝到独立 buffer（等价 Copy）
	386	+ // AscendC 没有裸 Copy API，使用 DataCopy 或 Adds(x, 0.0f) 均可；
	387	+ // 此处 Adds(p, inLocal_fp32, 0.0f) 最直观。
	388	+ LocalTensor<float> inFp32 = inLocal.template ReinterpretCast<float>();
	389	+ Adds(p, inFp32, 0.0f, len);
	390	+ } else {
	391	+ // half / bf16 → fp32
	392	+ Cast(p, inLocal, RoundMode::CAST_NONE, len);
	393	+ }
	394	+
	395	+ // Buffer 别名
	396	+ LocalTensor<float> tmpPm = tmpBuf0.Get<float>();
	397	+ LocalTensor<float> tmpZ = tmpBuf1.Get<float>();
	398	+ LocalTensor<float> tmpP = tmpBuf2.Get<float>();
	399	+ LocalTensor<float> tmpQ = tmpBuf3.Get<float>();
	400	+ LocalTensor<float> scratch = tmpBuf4.Get<float>();
	401	+ LocalTensor<float> pSafe = tmpBuf5.Get<float>();
	402	+ LocalTensor<float> yCenter = tmpBuf6.Get<float>();
	403	+ LocalTensor<float> yTail = tmpBuf7.Get<float>();
	404	+ LocalTensor<float> ySpecial = tmpBuf8.Get<float>();
	405	+
	406	+ LocalTensor<uint8_t> maskTail = maskBuf0.Get<uint8_t>();
	407	+ LocalTensor<uint8_t> maskNeg = maskBuf1.Get<uint8_t>();
	408	+ LocalTensor<uint8_t> maskSpecial = maskBuf2.Get<uint8_t>();
	409	+
	410	+ // ISSUE-001：后续所有 Vector 计算统一使用 lenAligned 长度运行。
	411	+ // 理由：
	412	+ // 1) Compare/CompareScalar 要求 count 所占空间 256B 对齐（64 元素）
	413	+ // 2) p 在 [len, lenAligned) 已填 0.5f 中性值，计算中不会产生 NaN/Inf
	414	+ // 3) 所有 UB buffer 均已分配 ubLength_ * sizeof(float)，且 ubLength_
	415	+ // 已由 Host Tiling FloorAlign(_, 256) 保证 ≥ lenAligned
	416	+ // 4) CopyOut 仍按 len 字节写回 GM，padding 位置的输出被自然丢弃
	417	+
	418	+ // Step 1: 构造 mask_tail / mask_neg / mask_special
	419	+ BuildMasks(p, maskTail, maskNeg, maskSpecial, scratch, lenAligned);
	420	+
	421	+ // Step 2: pSafe = clamp(p, FLT_MIN, 1 - 2^-24)
	422	+ Maxs(pSafe, p, NDTRI_SAFE_LO, lenAligned);
	423	+ Mins(pSafe, pSafe, NDTRI_SAFE_HI, lenAligned);
	424	+
	425	+ // Step 3: 计算 yTail（cal_tail 内部调用 cal_sub + cal_p12）
	426	+ // 采用"先 yTail 再 yCenter"策略，以使 yCenter (tmpBuf6) 可作为 CalTail 的 tmp3 使用。
	427	+ // buffer 映射（CalTail 参数名 → tmpBuf 来源）：
	428	+ // tmpQ(q) = tmpPm (tmpBuf0)
	429	+ // tmpX = tmpZ (tmpBuf1)
	430	+ // tmpX0 = tmpP (tmpBuf2)
	431	+ // tmpCorr = tmpQ (tmpBuf3)
	432	+ // tmp1 = scratch (tmpBuf4)
	433	+ // tmp2 = ySpecial (tmpBuf8) [Step 6 前 ySpecial 可借用]
	434	+ // tmp3 = yCenter (tmpBuf6) [随后会被 Step 4 覆盖]
	435	+ // tmp4 = tmpBuf9
	436	+ // tmp5 = tmpBuf10
	437	+ CalTail(yTail, pSafe, maskNeg,
	438	+ /tmpQ /tmpPm,
	439	+ /tmpX /tmpZ,
	440	+ /tmpX0 /tmpP,
	441	+ /tmpCorr/tmpQ,
	442	+ /tmp1 /scratch,
	443	+ /tmp2 /ySpecial,
	444	+ /tmp3 /yCenter,
	445	+ /tmp4 /tmpBuf9.Get<float>(),
	446	+ /tmp5 /tmpBuf10.Get<float>(),
	447	+ /maskX /maskBuf3.Get<uint8_t>(),
	448	+ lenAligned);
	449	+
	450	+ // Step 4: 计算 yCenter（覆盖 yCenter 暂借值）
	451	+ CalP0(yCenter, pSafe, tmpPm, tmpZ, tmpP, tmpQ, scratch, lenAligned);
	452	+
	453	+ // Step 5: y = select(maskTail, yTail, yCenter)
	454	+ Select(y, maskTail, yTail, yCenter,
	455	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, lenAligned);
	456	+
	457	+ // Step 6: y = select(maskSpecial, ySpecial, y)
	458	+ BuildSpecialY(ySpecial, p, scratch, lenAligned);
	459	+ Select(y, maskSpecial, ySpecial, y,
	460	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, lenAligned);
	461	+
	462	+ // Step 7: 输出 Cast → T
	463	+ // - FP32: Adds(outLocal_fp32, y, 0.0f) 等价 Copy（只处理前 len 位置，
	464	+ // padding 位置不需要写回 GM）
	465	+ // - FP16: Cast(fp32 → fp16, CAST_RINT)
	466	+ // - BF16: Cast(fp32 → bf16, CAST_RINT)
	467	+ if constexpr (IS_FP32) {
	468	+ LocalTensor<float> outFp32 = outLocal.template ReinterpretCast<float>();
	469	+ Adds(outFp32, y, 0.0f, len);
	470	+ } else {
	471	+ Cast(outLocal, y, RoundMode::CAST_RINT, len);
	472	+ }
	473	+
	474	+ outQueY.template EnQue<T>(outLocal);
	475	+ inQueSelf.FreeTensor(inLocal);
	476	+}
	477	+
	478	+// ---------------------------------------------------------------
	479	+// CopyOut
	480	+// DataCopyPad 天然支持对齐/非对齐两种路径。
	481	+// ---------------------------------------------------------------
	482	+template <typename T, int K_ALIGN>
	483	+__aicore__ inline void Ndtri<T, K_ALIGN>::CopyOut(int64_t progress, int64_t currentNum)
	484	+{
	485	+ LocalTensor<T> outLocal = outQueY.template DeQue<T>();
	486	+ DataCopyExtParams copyParams{
	487	+ 1, static_cast<uint32_t>(currentNum * sizeof(T)), 0, 0, 0};
	488	+ int64_t gmOffset = progress * ubLength_;
	489	+ DataCopyPad(outGm[gmOffset], outLocal, copyParams);
	490	+ outQueY.FreeTensor(outLocal);
	491	+}
	492	+
	493	+} // namespace NsNdtri
	494	+
	495	+#endif // NDTRI_KERNEL_H

Amath/ndtri/op_kernel/arch35/ndtri_tiling_data.h+29-0

@@ -0,0 +1,29 @@
	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	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file ndtri_tiling_data.h
	17	+ * \brief Ndtri Tiling 数据结构定义
	18	+ */
	19	+
	20	+#ifndef NDTRI_TILING_DATA_H_
	21	+#define NDTRI_TILING_DATA_H_
	22	+
	23	+struct NdtriTilingData {
	24	+ int64_t totalNum = 0; // self 元素总数
	25	+ int64_t blockFactor = 0; // 每核主体处理元素数（按 alignElem 向上对齐）
	26	+ int64_t ubFactor = 0; // 单次 UB 循环处理元素数（对齐 256）
	27	+};
	28	+
	29	+#endif // NDTRI_TILING_DATA_H_

Amath/ndtri/op_kernel/arch35/ndtri_tiling_key.h+56-0

@@ -0,0 +1,56 @@
	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	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+/*!
	16	+ * \file ndtri_tiling_key.h
	17	+ * \brief Ndtri Tiling 模板参数定义
	18	+ *
	19	+ * 模板参数维度：
	20	+ * - D_T: 输入/输出 Tensor 的数据类型（C_DT_FLOAT / C_DT_FLOAT16 / C_DT_BF16）
	21	+ * - K_ALIGN: 32B 对齐标记（1=对齐, 0=非对齐）
	22	+ *
	23	+ * 6 个 TilingKey：{fp32, fp16, bf16} × {对齐, 非对齐}
	24	+ * 全部通过 Cast(fp16↔fp32) / Cast(bf16↔fp32) 链路实现；fp32 走 Adds(*, 0.0f) 等价 Copy。
	25	+ */
	26	+
	27	+#ifndef NDTRI_TILING_KEY_H_
	28	+#define NDTRI_TILING_KEY_H_
	29	+
	30	+#include "ascendc/host_api/tiling/template_argument.h"
	31	+
	32	+ASCENDC_TPL_ARGS_DECL(Ndtri,
	33	+ ASCENDC_TPL_DATATYPE_DECL(D_T,
	34	+ C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16,
	35	+ ASCENDC_TPL_INPUT(0)),
	36	+ ASCENDC_TPL_UINT_DECL(K_ALIGN, 8, ASCENDC_TPL_UI_LIST, 0, 1)
	37	+);
	38	+
	39	+// 6 个 TilingKey：{fp32, fp16, bf16} × {对齐, 非对齐}
	40	+// 全部通过 Cast(fp16↔fp32) / Cast(bf16↔fp32) 链路实现；fp32 走 Adds(*, 0.0f) 等价 Copy。
	41	+ASCENDC_TPL_SEL(
	42	+ ASCENDC_TPL_ARGS_SEL(
	43	+ ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_FLOAT),
	44	+ ASCENDC_TPL_UINT_SEL(K_ALIGN, ASCENDC_TPL_UI_LIST, 0, 1)
	45	+ ),
	46	+ ASCENDC_TPL_ARGS_SEL(
	47	+ ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_FLOAT16),
	48	+ ASCENDC_TPL_UINT_SEL(K_ALIGN, ASCENDC_TPL_UI_LIST, 0, 1)
	49	+ ),
	50	+ ASCENDC_TPL_ARGS_SEL(
	51	+ ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_BF16),
	52	+ ASCENDC_TPL_UINT_SEL(K_ALIGN, ASCENDC_TPL_UI_LIST, 0, 1)
	53	+ ),
	54	+);
	55	+
	56	+#endif // NDTRI_TILING_KEY_H_

Amath/ndtri/op_kernel/ndtri_apt.cpp+38-0

@@ -0,0 +1,38 @@
	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 ndtri_apt.cpp
	15	+ * \brief Ndtri Kernel 入口（arch35 / Ascend950）
	16	+ *
	17	+ * 模板参数：
	18	+ * - D_T: 输入 Tensor 数据类型（float / half / bfloat16_t）
	19	+ * - K_ALIGN: 32B 对齐标记（0=非对齐, 1=对齐）
	20	+ *
	21	+ * 核函数参数顺序：1 输入 + 1 输出 + workspace + tiling
	22	+ */
	23	+
	24	+#include "arch35/ndtri_kernel.h"
	25	+
	26	+template <typename D_T, int K_ALIGN>
	27	+__global__ __aicore__ void ndtri(
	28	+ GM_ADDR self,
	29	+ GM_ADDR out,
	30	+ GM_ADDR workspace,
	31	+ GM_ADDR tiling)
	32	+{
	33	+ REGISTER_TILING_DEFAULT(NdtriTilingData);
	34	+ GET_TILING_DATA_WITH_STRUCT(NdtriTilingData, tilingData, tiling);
	35	+ NsNdtri::Ndtri<D_T, K_ALIGN> op;
	36	+ op.Init(self, out, &tilingData);
	37	+ op.Process();
	38	+}

Amath/ndtri/tests/.gitkeep+0-0

The file is empty

Amath/ndtri/tests/ut/op_host/arch35/test_ndtri_tiling_arch35.cpp+305-0

@@ -0,0 +1,305 @@
	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_ndtri_tiling_arch35.cpp
	15	+ * \brief Ndtri Tiling UT（arch35 / Ascend950）
	16	+ *
	17	+ * 覆盖路径（参见 op_host/arch35/ndtri_tiling_arch35.cpp）:
	18	+ * 1) dtype 路径（DispatchTilingKey）：FP32 / FP16 / BF16 → C_DT_FLOAT/C_DT_FLOAT16/C_DT_BF16
	19	+ * 2) isAlign 路径（DispatchTilingKey）：
	20	+ * - totalNum % alignElem == 0 → isAlign=1
	21	+ * - totalNum % alignElem != 0 → isAlign=0
	22	+ * 3) 多核 vs 单核（DoTiling）：
	23	+ * - totalNum < alignElem → 单核 usedCoreNum=1
	24	+ * - totalNum >= alignElem → 多核
	25	+ * 4) 失败路径：
	26	+ * - 不支持 dtype（INT32/DOUBLE 等）→ CheckDtype 返回 GRAPH_FAILED
	27	+ * - totalNum=0（zero-shape）→ GetTotalNum 返回 GRAPH_FAILED
	28	+ *
	29	+ * 期望 tilingData 字段（NdtriTilingData，定义在 op_kernel/arch35/ndtri_tiling_data.h）：
	30	+ * "totalNum blockFactor ubFactor "
	31	+ *
	32	+ * 期望 TilingKey: ASCENDC_TPL_SEL_PARAM(context, dtypeKey, isAlign)
	33	+ * dtypeKey: 0=C_DT_FLOAT, 1=C_DT_FLOAT16, 2=C_DT_BF16（具体取决于框架编码）
	34	+ * 实际 TilingKey 由框架根据 (dtypeKey, isAlign) 计算，本 UT 通过 expectTilingKey 固化运行时实际值。
	35	+ */
	36	+
	37	+#include <cstdint>
	38	+#include <iostream>
	39	+#include <gtest/gtest.h>
	40	+#include "../../../../op_kernel/arch35/ndtri_tiling_data.h"
	41	+#include "tiling_context_faker.h"
	42	+#include "tiling_case_executor.h"
	43	+
	44	+namespace optiling {
	45	+// Ndtri 的 TilingParse 不读取任何字段（NdtriCompileInfo 是空 struct）。
	46	+// 这里定义一个本地占位类型，方便 UT 提供 compileInfo 指针。
	47	+struct NdtriCompileInfoStub {
	48	+ uint64_t reserved = 0;
	49	+};
	50	+} // namespace optiling
	51	+
	52	+using namespace std;
	53	+
	54	+class NdtriTilingTest : public testing::Test {
	55	+protected:
	56	+ static void SetUpTestCase()
	57	+ {
	58	+ std::cout << "NdtriTilingTest SetUp" << std::endl;
	59	+ }
	60	+
	61	+ static void TearDownTestCase()
	62	+ {
	63	+ std::cout << "NdtriTilingTest TearDown" << std::endl;
	64	+ }
	65	+};
	66	+
	67	+// ==========================================================================
	68	+// 主路径：FP32 对齐 + 多核
	69	+// shape {1, 64, 2, 64} = 8192 elem
	70	+// FP32: alignElem = 32/4 = 8
	71	+// 8192 % 8 == 0 → isAlign=1
	72	+// 8192 >= alignElem → 多核
	73	+// 期望 tilingData: "8192 <blockFactor> <ubFactor> "
	74	+// coreNum=64: perCoreRaw=ceil(8192/64)=128，blockFactor=ceil_align(128,8)=128
	75	+// ubFactor: availableUb=(262144-49152)/80=2662 → floor_align(2662,256)=2560
	76	+// ==========================================================================
	77	+TEST_F(NdtriTilingTest, test_tiling_fp32_align_001)
	78	+{
	79	+ optiling::NdtriCompileInfoStub compileInfo;
	80	+ gert::TilingContextPara tilingContextPara(
	81	+ "Ndtri",
	82	+ {
	83	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
	84	+ },
	85	+ {
	86	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
	87	+ },
	88	+ &compileInfo);
	89	+ // TilingKey = 256：fp32 + isAlign=1（由 ASCENDC_TPL_SEL_PARAM 编码）
	90	+ uint64_t expectTilingKey = 256;
	91	+ string expectTilingData = "8192 128 2560 ";
	92	+ std::vector<size_t> expectWorkspaces = {16777216};
	93	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	94	+}
	95	+
	96	+// ==========================================================================
	97	+// FP32 非对齐 + 单核（totalNum < alignElem）
	98	+// shape {1, 7} = 7 elem
	99	+// FP32: alignElem=8 → 7 < 8 → 单核 usedCoreNum=1
	100	+// 7 % 8 != 0 → isAlign=0
	101	+// blockFactor=7
	102	+// ==========================================================================
	103	+TEST_F(NdtriTilingTest, test_tiling_fp32_unalign_small_single_core_002)
	104	+{
	105	+ optiling::NdtriCompileInfoStub compileInfo;
	106	+ gert::TilingContextPara tilingContextPara(
	107	+ "Ndtri",
	108	+ {
	109	+ {{{1, 7}, {1, 7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	110	+ },
	111	+ {
	112	+ {{{1, 7}, {1, 7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	113	+ },
	114	+ &compileInfo);
	115	+ // TilingKey = 0：fp32 + isAlign=0
	116	+ uint64_t expectTilingKey = 0;
	117	+ string expectTilingData = "7 7 2560 ";
	118	+ std::vector<size_t> expectWorkspaces = {16777216};
	119	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	120	+}
	121	+
	122	+// ==========================================================================
	123	+// FP16 对齐 + 多核
	124	+// shape {1, 64, 2, 64} = 8192 elem
	125	+// FP16: alignElem=32/2=16；8192 % 16 == 0 → isAlign=1
	126	+// perCoreRaw=ceil(8192/64)=128，blockFactor=ceil_align(128,16)=128
	127	+// ubFactor: availableUb/64=(262144-49152)/64=3328 → floor_align(3328,256)=3328
	128	+// ==========================================================================
	129	+TEST_F(NdtriTilingTest, test_tiling_fp16_align_003)
	130	+{
	131	+ optiling::NdtriCompileInfoStub compileInfo;
	132	+ gert::TilingContextPara tilingContextPara(
	133	+ "Ndtri",
	134	+ {
	135	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	136	+ },
	137	+ {
	138	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	139	+ },
	140	+ &compileInfo);
	141	+ // TilingKey = 257：fp16 + isAlign=1
	142	+ uint64_t expectTilingKey = 257;
	143	+ string expectTilingData = "8192 128 3328 ";
	144	+ std::vector<size_t> expectWorkspaces = {16777216};
	145	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	146	+}
	147	+
	148	+// ==========================================================================
	149	+// FP16 小规模单核（< alignElem）
	150	+// shape {3} = 3 elem，FP16 alignElem=16
	151	+// 3 < 16 → 单核
	152	+// 3 % 16 != 0 → isAlign=0
	153	+// ==========================================================================
	154	+TEST_F(NdtriTilingTest, test_tiling_fp16_small_single_core_004)
	155	+{
	156	+ optiling::NdtriCompileInfoStub compileInfo;
	157	+ gert::TilingContextPara tilingContextPara(
	158	+ "Ndtri",
	159	+ {
	160	+ {{{3}, {3}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	161	+ },
	162	+ {
	163	+ {{{3}, {3}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	164	+ },
	165	+ &compileInfo);
	166	+ // TilingKey = 1：fp16 + isAlign=0
	167	+ uint64_t expectTilingKey = 1;
	168	+ string expectTilingData = "3 3 3328 ";
	169	+ std::vector<size_t> expectWorkspaces = {16777216};
	170	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	171	+}
	172	+
	173	+// ==========================================================================
	174	+// BF16 对齐 + 多核（与 FP16 共享对齐路径）
	175	+// ==========================================================================
	176	+TEST_F(NdtriTilingTest, test_tiling_bf16_align_005)
	177	+{
	178	+ optiling::NdtriCompileInfoStub compileInfo;
	179	+ gert::TilingContextPara tilingContextPara(
	180	+ "Ndtri",
	181	+ {
	182	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_BF16, ge::FORMAT_ND},
	183	+ },
	184	+ {
	185	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_BF16, ge::FORMAT_ND},
	186	+ },
	187	+ &compileInfo);
	188	+ // TilingKey = 283：bf16 + isAlign=1（与 fp32/fp16 的连续编码不一致，由 ASCENDC_TPL_SEL_PARAM 内部映射决定）
	189	+ uint64_t expectTilingKey = 283;
	190	+ string expectTilingData = "8192 128 3328 ";
	191	+ std::vector<size_t> expectWorkspaces = {16777216};
	192	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	193	+}
	194	+
	195	+// ==========================================================================
	196	+// FP32 非对齐 + 多核（大 shape）
	197	+// shape {2, 1000} = 2000 elem
	198	+// FP32 alignElem=8；2000 % 8 == 0（实际对齐）→ 改为 {1, 997} 制造非对齐
	199	+// 997 not divisible by 8 → isAlign=0
	200	+// 997 >= 8 → 多核
	201	+// perCoreRaw=ceil(997/64)=16，blockFactor=ceil_align(16,8)=16
	202	+// ==========================================================================
	203	+TEST_F(NdtriTilingTest, test_tiling_fp32_unalign_multi_core_006)
	204	+{
	205	+ optiling::NdtriCompileInfoStub compileInfo;
	206	+ gert::TilingContextPara tilingContextPara(
	207	+ "Ndtri",
	208	+ {
	209	+ {{{1, 997}, {1, 997}}, ge::DT_FLOAT, ge::FORMAT_ND},
	210	+ },
	211	+ {
	212	+ {{{1, 997}, {1, 997}}, ge::DT_FLOAT, ge::FORMAT_ND},
	213	+ },
	214	+ &compileInfo);
	215	+ // TilingKey = 0：fp32 + isAlign=0
	216	+ uint64_t expectTilingKey = 0;
	217	+ string expectTilingData = "997 16 2560 ";
	218	+ std::vector<size_t> expectWorkspaces = {16777216};
	219	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
	220	+}
	221	+
	222	+// ==========================================================================
	223	+// 失败路径：不支持的 dtype（INT32）
	224	+// CheckDtype 应返回 GRAPH_FAILED
	225	+// ==========================================================================
	226	+TEST_F(NdtriTilingTest, test_tiling_fail_unsupported_dtype_int32_007)
	227	+{
	228	+ optiling::NdtriCompileInfoStub compileInfo;
	229	+ gert::TilingContextPara tilingContextPara(
	230	+ "Ndtri",
	231	+ {
	232	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_INT32, ge::FORMAT_ND},
	233	+ },
	234	+ {
	235	+ {{{1, 64, 2, 64}, {1, 64, 2, 64}}, ge::DT_INT32, ge::FORMAT_ND},
	236	+ },
	237	+ &compileInfo);
	238	+ uint64_t expectTilingKey = 0;
	239	+ string expectTilingData = "";
	240	+ std::vector<size_t> expectWorkspaces = {0};
	241	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);
	242	+}
	243	+
	244	+// ==========================================================================
	245	+// 失败路径：不支持的 dtype（DOUBLE）
	246	+// ==========================================================================
	247	+TEST_F(NdtriTilingTest, test_tiling_fail_unsupported_dtype_double_008)
	248	+{
	249	+ optiling::NdtriCompileInfoStub compileInfo;
	250	+ gert::TilingContextPara tilingContextPara(
	251	+ "Ndtri",
	252	+ {
	253	+ {{{1, 1, 2, 64}, {1, 1, 2, 64}}, ge::DT_DOUBLE, ge::FORMAT_ND},
	254	+ },
	255	+ {
	256	+ {{{1, 1, 2, 64}, {1, 1, 2, 64}}, ge::DT_DOUBLE, ge::FORMAT_ND},
	257	+ },
	258	+ &compileInfo);
	259	+ uint64_t expectTilingKey = 0;
	260	+ string expectTilingData = "";
	261	+ std::vector<size_t> expectWorkspaces = {0};
	262	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);
	263	+}
	264	+
	265	+// ==========================================================================
	266	+// 失败路径：zero-shape 1D —— shape {0} → totalNum=0 → GetTotalNum 返回 GRAPH_FAILED
	267	+// ==========================================================================
	268	+TEST_F(NdtriTilingTest, test_tiling_fail_zero_shape_1d_009)
	269	+{
	270	+ optiling::NdtriCompileInfoStub compileInfo;
	271	+ gert::TilingContextPara tilingContextPara(
	272	+ "Ndtri",
	273	+ {
	274	+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
	275	+ },
	276	+ {
	277	+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
	278	+ },
	279	+ &compileInfo);
	280	+ uint64_t expectTilingKey = 0;
	281	+ string expectTilingData = "";
	282	+ std::vector<size_t> expectWorkspaces = {0};
	283	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);
	284	+}
	285	+
	286	+// ==========================================================================
	287	+// 失败路径：zero-shape 3D —— shape {4, 0, 3} → totalNum=0 → GetTotalNum 返回 GRAPH_FAILED
	288	+// ==========================================================================
	289	+TEST_F(NdtriTilingTest, test_tiling_fail_zero_shape_3d_010)
	290	+{
	291	+ optiling::NdtriCompileInfoStub compileInfo;
	292	+ gert::TilingContextPara tilingContextPara(
	293	+ "Ndtri",
	294	+ {
	295	+ {{{4, 0, 3}, {4, 0, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
	296	+ },
	297	+ {
	298	+ {{{4, 0, 3}, {4, 0, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
	299	+ },
	300	+ &compileInfo);
	301	+ uint64_t expectTilingKey = 0;
	302	+ string expectTilingData = "";
	303	+ std::vector<size_t> expectWorkspaces = {0};
	304	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);
	305	+}

Amath/ndtri/tests/ut/op_host/test_ndtri_infershape.cpp+183-0

@@ -0,0 +1,183 @@
	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_ndtri_infershape.cpp
	15	+ * \brief Ndtri InferShape / InferDataType UT
	16	+ *
	17	+ * 覆盖：
	18	+ * - 多维 shape：1D / 2D / 3D / 4D / 5D
	19	+ * - 多 dtype：FP32 / FP16 / BF16
	20	+ * - 空 tensor：含 0 维度（InferShape 仅做形状/类型推导，不校验 0 维）
	21	+ * - 标量
	22	+ * - 大 shape
	23	+ *
	24	+ * 关键真值：
	25	+ * - OpType: "Ndtri"
	26	+ * - 输出 shape == 输入 shape（逐元素）
	27	+ * - 输出 dtype == 输入 dtype（透传）
	28	+ * - format == FORMAT_ND
	29	+ */
	30	+
	31	+#include <gtest/gtest.h>
	32	+#include <iostream>
	33	+#include "infershape_context_faker.h"
	34	+#include "infershape_case_executor.h"
	35	+
	36	+class NdtriInfershape : public testing::Test {
	37	+protected:
	38	+ static void SetUpTestCase()
	39	+ {
	40	+ std::cout << "NdtriInfershape SetUp" << std::endl;
	41	+ }
	42	+
	43	+ static void TearDownTestCase()
	44	+ {
	45	+ std::cout << "NdtriInfershape TearDown" << std::endl;
	46	+ }
	47	+};
	48	+
	49	+// FP32 4D —— 与 tiling 主路径完全一致的 shape，作为基线
	50	+TEST_F(NdtriInfershape, ndtri_infershape_float_test)
	51	+{
	52	+ gert::InfershapeContextPara infershapeContextPara(
	53	+ "Ndtri",
	54	+ {
	55	+ {{{2, 3, 4}, {2, 3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	56	+ },
	57	+ {
	58	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	59	+ });
	60	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	61	+ {2, 3, 4},
	62	+ };
	63	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	64	+}
	65	+
	66	+// FP16 多维
	67	+TEST_F(NdtriInfershape, ndtri_infershape_float16_test)
	68	+{
	69	+ gert::InfershapeContextPara infershapeContextPara(
	70	+ "Ndtri",
	71	+ {
	72	+ {{{4, 3, 4}, {4, 3, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	73	+ },
	74	+ {
	75	+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	76	+ });
	77	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	78	+ {4, 3, 4},
	79	+ };
	80	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	81	+}
	82	+
	83	+// BF16 多维
	84	+TEST_F(NdtriInfershape, ndtri_infershape_bf16_test)
	85	+{
	86	+ gert::InfershapeContextPara infershapeContextPara(
	87	+ "Ndtri",
	88	+ {
	89	+ {{{2, 3, 4}, {2, 3, 4}}, ge::DT_BF16, ge::FORMAT_ND},
	90	+ },
	91	+ {
	92	+ {{{}, {}}, ge::DT_BF16, ge::FORMAT_ND},
	93	+ });
	94	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	95	+ {2, 3, 4},
	96	+ };
	97	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	98	+}
	99	+
	100	+// 1D
	101	+TEST_F(NdtriInfershape, ndtri_infershape_1d_tensor_test)
	102	+{
	103	+ gert::InfershapeContextPara infershapeContextPara(
	104	+ "Ndtri",
	105	+ {
	106	+ {{{10}, {10}}, ge::DT_FLOAT, ge::FORMAT_ND},
	107	+ },
	108	+ {
	109	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	110	+ });
	111	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	112	+ {10},
	113	+ };
	114	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	115	+}
	116	+
	117	+// 2D
	118	+TEST_F(NdtriInfershape, ndtri_infershape_2d_tensor_test)
	119	+{
	120	+ gert::InfershapeContextPara infershapeContextPara(
	121	+ "Ndtri",
	122	+ {
	123	+ {{{8, 16}, {8, 16}}, ge::DT_FLOAT, ge::FORMAT_ND},
	124	+ },
	125	+ {
	126	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	127	+ });
	128	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	129	+ {8, 16},
	130	+ };
	131	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	132	+}
	133	+
	134	+// 5D 高维
	135	+TEST_F(NdtriInfershape, ndtri_infershape_5d_tensor_test)
	136	+{
	137	+ gert::InfershapeContextPara infershapeContextPara(
	138	+ "Ndtri",
	139	+ {
	140	+ {{{1, 2, 3, 4, 5}, {1, 2, 3, 4, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	141	+ },
	142	+ {
	143	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	144	+ });
	145	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	146	+ {1, 2, 3, 4, 5},
	147	+ };
	148	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	149	+}
	150	+
	151	+// 含 0 维度的 empty tensor —— InferShape 阶段不校验 0 维，仍 SUCCESS
	152	+TEST_F(NdtriInfershape, ndtri_infershape_empty_tensor_test)
	153	+{
	154	+ gert::InfershapeContextPara infershapeContextPara(
	155	+ "Ndtri",
	156	+ {
	157	+ {{{0, 3, 4}, {0, 3, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	158	+ },
	159	+ {
	160	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	161	+ });
	162	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	163	+ {0, 3, 4},
	164	+ };
	165	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	166	+}
	167	+
	168	+// 大 shape
	169	+TEST_F(NdtriInfershape, ndtri_infershape_large_shape_test)
	170	+{
	171	+ gert::InfershapeContextPara infershapeContextPara(
	172	+ "Ndtri",
	173	+ {
	174	+ {{{1024, 1024}, {1024, 1024}}, ge::DT_FLOAT, ge::FORMAT_ND},
	175	+ },
	176	+ {
	177	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	178	+ });
	179	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	180	+ {1024, 1024},
	181	+ };
	182	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	183	+}