[CANNBot]950添加算子xlog1py-ops-math-AtomGit

Mdocs/zh/op_api_list.md+1-0

	@@ -317,6 +317,7 @@
317	\| [aclnnVar](../../math/reduce_var/docs/aclnnVar.md) \| 返回输入Tensor指定维度的值求得的方差。 \| 默认确定性实现\| 默认确定性实现\|	317	\| [aclnnVar](../../math/reduce_var/docs/aclnnVar.md) \| 返回输入Tensor指定维度的值求得的方差。 \| 默认确定性实现\| 默认确定性实现\|
318	\| [aclnnVarCorrection](../../math/reduce_var/docs/aclnnVarCorrection.md) \| 返回输入Tensor指定维度的值求得的方差。 \| 默认确定性实现\| 默认确定性实现 \|	318	\| [aclnnVarCorrection](../../math/reduce_var/docs/aclnnVarCorrection.md) \| 返回输入Tensor指定维度的值求得的方差。 \| 默认确定性实现\| 默认确定性实现 \|
319	\| [aclnnVarMean](../../math/reduce_var/docs/aclnnVarMean.md) \| 返回输入Tensor指定维度的值求得的均值及方差。 \| 默认确定性实现\| 默认确定性实现\|	319	\| [aclnnVarMean](../../math/reduce_var/docs/aclnnVarMean.md) \| 返回输入Tensor指定维度的值求得的均值及方差。 \| 默认确定性实现\| 默认确定性实现\|
		320	+\| [aclnnXlog1py](../../math/xlog1py/docs/aclnnXlog1py.md) \| 计算 x * log(1 + y)，当 x == 0 时结果为 0。支持 broadcast。 \| 默认确定性实现\| 默认确定性实现 \|
320	\| [aclnnXLogYScalarOther&aclnnInplaceXLogYScalarOther](../../math/x_log_y/docs/aclnnXLogYScalarOther&aclnnInplaceXLogYScalarOther.md) \| 计算self * log(other)的结果。 \| 默认确定性实现\| - \|	321	\| [aclnnXLogYScalarOther&aclnnInplaceXLogYScalarOther](../../math/x_log_y/docs/aclnnXLogYScalarOther&aclnnInplaceXLogYScalarOther.md) \| 计算self * log(other)的结果。 \| 默认确定性实现\| - \|
321	\| [aclnnXLogYScalarSelf](../../math/x_log_y/docs/aclnnXLogYScalarSelf.md) \| 计算self * log(other)的结果。 \| 默认确定性实现\| - \|	322	\| [aclnnXLogYScalarSelf](../../math/x_log_y/docs/aclnnXLogYScalarSelf.md) \| 计算self * log(other)的结果。 \| 默认确定性实现\| - \|
322	\| [aclnnXLogYTensor&aclnnInplaceXLogYTensor](../../math/x_log_y/docs/aclnnXLogYTensor&aclnnInplaceXLogYTensor.md) \| 计算self * log(other)的结果。 \| 默认确定性实现\| - \|	323	\| [aclnnXLogYTensor&aclnnInplaceXLogYTensor](../../math/x_log_y/docs/aclnnXLogYTensor&aclnnInplaceXLogYTensor.md) \| 计算self * log(other)的结果。 \| 默认确定性实现\| - \|

Mdocs/zh/op_list.md+10-0

	@@ -1826,6 +1826,16 @@
1826	<td>AI Core</td>	1826	<td>AI Core</td>
1827	<td>将输入的张量沿指定轴重复多次，但不进行广播。</td>	1827	<td>将输入的张量沿指定轴重复多次，但不进行广播。</td>
1828	</tr>	1828	</tr>
		1829	+ <tr>
		1830	+ <td>math</td>
		1831	+ <td><a href="../../math/xlog1py/README.md">xlog1py</a></td>
		1832	+ <td>√</td>
		1833	+ <td>√</td>
		1834	+ <td>√</td>
		1835	+ <td>√</td>
		1836	+ <td>AI Core</td>
		1837	+ <td>计算 x * log(1 + y)，当 x == 0 时结果为 0。支持 broadcast。</td>
		1838	+ </tr>
1829	<tr>	1839	<tr>
1830	<td>math</td>	1840	<td>math</td>
1831	<td><a href="../../math/zero_op/README.md">zeros_like</a></td>	1841	<td><a href="../../math/zero_op/README.md">zeros_like</a></td>

Amath/xlog1py/CMakeLists.txt+15-0

@@ -0,0 +1,15 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# 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	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	13	+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
	14	+set(SUPPORT_TILING_DIR "arch35")
	15	+add_all_modules_sources(OPTYPE xlog1py ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)

Amath/xlog1py/README.md+78-0

@@ -0,0 +1,78 @@
	1	+# Xlog1py
	C C chenjiao6月1日交付件必须配套修改： ①算子readme----同步检查op_list.md ②aclnn API----同步检查op_api_list.md
	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	+- 算子功能：计算 x * log(1 + y)，当 x == 0 时结果为 0。
	17	+- 计算公式：
	18	+
	19	+$$
	20	+z_i =
	21	+\begin{cases}
	22	+0, & x_i = 0 \\
	23	+x_i \cdot \log(1 + y_i), & x_i \neq 0
	24	+\end{cases}
	25	+$$
	26	+
	27	+## 参数说明
	28	+
	29	+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
	30	+ <col style="width: 100px">
	31	+ <col style="width: 150px">
	32	+ <col style="width: 280px">
	33	+ <col style="width: 330px">
	34	+ <col style="width: 120px">
	35	+ </colgroup>
	36	+ <thead>
	37	+ <tr>
	38	+ <th>参数名</th>
	39	+ <th>输入/输出/属性</th>
	40	+ <th>描述</th>
	41	+ <th>数据类型</th>
	42	+ <th>数据格式</th>
	43	+ </tr></thead>
	44	+ <tbody>
	45	+ <tr>
	46	+ <td>x</td>
	47	+ <td>输入</td>
	48	+ <td>公式中的 x，乘数因子。shape 需与 y 满足 broadcast 关系。</td>
	49	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	50	+ <td>ND</td>
	51	+ </tr>
	52	+ <tr>
	53	+ <td>y</td>
	54	+ <td>输入</td>
	55	+ <td>公式中的 y，log1p 的自变量。shape 需与 x 满足 broadcast 关系。</td>
	56	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	57	+ <td>ND</td>
	58	+ </tr>
	59	+ <tr>
	60	+ <td>z</td>
	61	+ <td>输出</td>
	62	+ <td>公式中的 z，计算结果。shape 为 x 与 y broadcast 后的最大值 shape。</td>
	63	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	64	+ <td>ND</td>
	65	+ </tr>
	66	+ </tbody></table>
	67	+
	68	+## 约束说明
	69	+
	70	+- 输入 shape 维度最大为 8。
	71	+- 输入 x、y、输出 z 的数据类型必须一致。
	72	+
	73	+## 调用说明
	C C chenjiao6月1日调用说明请参考https://gitcode.com/cann/ops-math/blob/621bc97babff15fe21d91bf7793da7ef2eb20128/math/abs/README.md#调用说明
	74	+
	75	+\| 调用方式 \| 调用样例 \| 说明 \|
	76	+\|--------------\|------------------------------------------------------------------------\|--------------------------------------------------------------\|
	77	+\| aclnn调用 \| [test_aclnn_xlog1py](./examples/arch35/test_aclnn_xlog1py.cpp) \| 通过[aclnnXlog1py](./docs/aclnnXlog1py.md)接口方式调用Xlog1py算子。 \|
	78	+\| 图模式调用 \| [test_geir_xlog1py](./examples/test_geir_xlog1py.cpp) \| 通过[算子IR](./op_graph/xlog1py_proto.h)构图方式调用Xlog1py算子。 \|

Amath/xlog1py/docs/aclnnXlog1py.md+451-0

@@ -0,0 +1,451 @@
	1	+# aclnnXlog1py
	2	+
	C C chenjiao6月1日增加查看源码跳转，绝对link
	3	+[📄 查看源码](https://gitcode.com/cann/ops-math/tree/master/math/xlog1py)
	4	+
	5	+## 产品支持情况
	6	+
	7	+\| 产品 \| 是否支持 \|
	8	+\|:------------------------------------------------\| :------: \|
	9	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
	10	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| × \|
	11	+\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| × \|
	12	+\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|
	13	+\| <term>Atlas 推理系列产品</term> \| × \|
	14	+\| <term>Atlas 训练系列产品</term> \| × \|
	15	+
	16	+## 功能说明
	17	+
	18	+- 接口功能：计算 x * log(1 + y)，当 x == 0 时结果为 0。
	19	+
	20	+- 计算公式：
	21	+
	22	+$$
	23	+z_i =
	24	+\begin{cases}
	25	+0, & x_i = 0 \\
	26	+x_i \cdot \log(1 + y_i), & x_i \neq 0
	27	+\end{cases}
	28	+$$
	29	+
	30	+- x 与 y 支持 broadcast，输出 shape 为广播后的最大值 shape。
	31	+- y 为 NaN 时输出 y 原值。
	32	+
	33	+## 函数原型
	34	+
	35	+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用"aclnnXlog1pyGetWorkspaceSize"接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用"aclnnXlog1py"接口执行计算。
	36	+
	37	+```Cpp
	38	+aclnnStatus aclnnXlog1pyGetWorkspaceSize(
	39	+ const aclTensor *x,
	40	+ const aclTensor *y,
	41	+ const aclTensor *z,
	42	+ uint64_t *workspaceSize,
	43	+ aclOpExecutor **executor)
	44	+```
	45	+
	46	+```Cpp
	47	+aclnnStatus aclnnXlog1py(
	48	+ void *workspace,
	49	+ uint64_t workspaceSize,
	50	+ aclOpExecutor *executor,
	51	+ const aclrtStream stream)
	52	+```
	53	+
	54	+## aclnnXlog1pyGetWorkspaceSize
	55	+
	56	+- 参数说明：
	57	+
	58	+ <table style="undefined;table-layout: fixed; width: 1494px"><colgroup>
	59	+ <col style="width: 146px">
	60	+ <col style="width: 110px">
	61	+ <col style="width: 301px">
	62	+ <col style="width: 219px">
	63	+ <col style="width: 328px">
	64	+ <col style="width: 101px">
	65	+ <col style="width: 143px">
	66	+ <col style="width: 146px">
	67	+ </colgroup>
	68	+ <thead>
	69	+ <tr>
	70	+ <th>参数名</th>
	71	+ <th>输入/输出</th>
	72	+ <th>描述</th>
	73	+ <th>使用说明</th>
	74	+ <th>数据类型</th>
	75	+ <th>数据格式</th>
	76	+ <th>维度(shape)</th>
	77	+ <th>非连续Tensor</th>
	78	+ </tr></thead>
	79	+ <tbody>
	80	+ <tr>
	81	+ <td>x</td>
	82	+ <td>输入</td>
	83	+ <td>表示乘数因子，对应公式中 x。</td>
	84	+ <td><ul><li>支持空Tensor。</li><li>数据类型需与 y、z 保持一致。</li><li>shape 需与 y 满足broadcast关系。</li></ul></td>
	85	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	86	+ <td>ND</td>
	87	+ <td>0-8</td>
	88	+ <td>√</td>
	89	+ </tr>
	90	+ <tr>
	91	+ <td>y</td>
	92	+ <td>输入</td>
	93	+ <td>表示 log1p 的自变量，对应公式中 y。</td>
	94	+ <td><ul><li>支持空Tensor。</li><li>数据类型需与 x 保持一致。</li><li>shape 需与 x 满足broadcast关系。</li></ul></td>
	95	+ <td>数据类型与 x 保持一致。</td>
	96	+ <td>ND</td>
	97	+ <td>0-8</td>
	98	+ <td>√</td>
	99	+ </tr>
	100	+ <tr>
	101	+ <td>z</td>
	102	+ <td>输出</td>
	103	+ <td>表示计算结果，对应公式中 z。</td>
	104	+ <td>shape 需为 x 与 y broadcast 后的 shape。</td>
	105	+ <td>数据类型与 x 保持一致。</td>
	106	+ <td>ND</td>
	107	+ <td>1-8</td>
	108	+ <td>√</td>
	109	+ </tr>
	110	+ <tr>
	111	+ <td>workspaceSize</td>
	112	+ <td>输出</td>
	113	+ <td>返回需要在Device侧申请的workspace大小。</td>
	114	+ <td>-</td>
	115	+ <td>-</td>
	116	+ <td>-</td>
	117	+ <td>-</td>
	118	+ <td>-</td>
	119	+ </tr>
	120	+ <tr>
	121	+ <td>executor</td>
	122	+ <td>输出</td>
	123	+ <td>返回op执行器，包含了算子计算流程。</td>
	124	+ <td>-</td>
	125	+ <td>-</td>
	126	+ <td>-</td>
	127	+ <td>-</td>
	128	+ <td>-</td>
	129	+ </tr>
	130	+ </tbody>
	131	+ </table>
	132	+
	133	+- 返回值：
	134	+
	135	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
	136	+
	137	+ 第一段接口会完成入参校验，出现以下场景时报错：
	138	+
	139	+ <table style="undefined;table-layout: fixed;width: 1155px"><colgroup>
	140	+ <col style="width: 319px">
	141	+ <col style="width: 144px">
	142	+ <col style="width: 671px">
	143	+ </colgroup>
	144	+ <thead>
	145	+ <tr>
	146	+ <th>返回码</th>
	147	+ <th>错误码</th>
	148	+ <th>描述</th>
	149	+ </tr>
	150	+ </thead>
	151	+ <tbody>
	152	+ <tr>
	153	+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
	154	+ <td>161001</td>
	155	+ <td>传入的tensor是空指针。</td>
	156	+ </tr>
	157	+ <tr>
	158	+ <td rowspan="4">ACLNN_ERR_PARAM_INVALID</td>
	159	+ <td rowspan="4">161002</td>
	160	+ <td>x、y 或 z 的数据类型和数据格式不在支持的范围之内。</td>
	161	+ </tr>
	162	+ <tr>
	163	+ <td>x、y、z 的数据类型不匹配。</td>
	164	+ </tr>
	165	+ <tr>
	166	+ <td>x、y 或 z 的 shape 维度不在支持的范围之内（最大 8 维）。</td>
	167	+ </tr>
	168	+ <tr>
	169	+ <td>x、y 或 z 使用了私有格式（private format）。</td>
	170	+ </tr>
	171	+ </tbody></table>
	172	+
	173	+## aclnnXlog1py
	174	+
	175	+- 参数说明：
	176	+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
	177	+ <col style="width: 173px">
	178	+ <col style="width: 112px">
	179	+ <col style="width: 668px">
	180	+ </colgroup>
	181	+ <thead>
	182	+ <tr>
	183	+ <th>参数名</th>
	184	+ <th>输入/输出</th>
	185	+ <th>描述</th>
	186	+ </tr></thead>
	187	+ <tbody>
	188	+ <tr>
	189	+ <td>workspace</td>
	190	+ <td>输入</td>
	191	+ <td>在Device侧申请的workspace内存地址。</td>
	192	+ </tr>
	193	+ <tr>
	194	+ <td>workspaceSize</td>
	195	+ <td>输入</td>
	196	+ <td>在Device侧申请的workspace大小，由第一段接口aclnnXlog1pyGetWorkspaceSize获取。</td>
	197	+ </tr>
	198	+ <tr>
	199	+ <td>executor</td>
	200	+ <td>输入</td>
	201	+ <td>op执行器，包含了算子计算流程。</td>
	202	+ </tr>
	203	+ <tr>
	204	+ <td>stream</td>
	205	+ <td>输入</td>
	206	+ <td>指定执行任务的Stream。</td>
	207	+ </tr>
	208	+ </tbody>
	209	+ </table>
	210	+
	211	+- 返回值：
	212	+
	213	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
	214	+
	215	+## 约束说明
	216	+
	217	+- 确定性计算：
	218	+ - aclnnXlog1py默认确定性实现。
	219	+- 输入数据类型必须为 FLOAT、FLOAT16 或 BFLOAT16。
	220	+- shape 维度范围为 0-8。
	221	+- 不支持私有格式（private format），仅支持 ND 格式。
	222	+- 输入 x、y，输出 z 的数据类型必须一致。
	223	+
	224	+## 调用示例
	225	+
	226	+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
	227	+
	228	+```Cpp
	229	+#include <iostream>
	230	+#include <vector>
	231	+#include <cmath>
	232	+#include "acl/acl.h"
	233	+#include "aclnnop/aclnn_xlog1py.h"
	234	+
	235	+#define CHECK_RET(cond, msg) \
	236	+ do { \
	237	+ if (!(cond)) { \
	238	+ printf("[FAIL] " msg "\n"); \
	239	+ return -1; \
	240	+ } \
	241	+ } while (0)
	242	+
	243	+#define LOG_PRINT(msg, ...) printf(msg "\n", ##__VA_ARGS__)
	244	+
	245	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
	246	+{
	247	+ int64_t size = 1;
	248	+ for (auto i : shape) size *= i;
	249	+ return size;
	250	+}
	251	+
	252	+// Broadcast index: map flat index in output to flat index in input
	253	+static int64_t BroadcastIdx(int64_t flat, const std::vector<int64_t>& inShape,
	254	+ const std::vector<int64_t>& outShape)
	255	+{
	256	+ int inRank = (int)inShape.size();
	257	+ int outRank = (int)outShape.size();
	258	+ int64_t outIdx = 0, outStride = 1;
	259	+ for (int d = 0; d < outRank; d++) {
	260	+ int dimIdx = outRank - 1 - d;
	261	+ int64_t dim = outShape[dimIdx];
	262	+ int64_t coord = (flat / outStride) % dim;
	263	+ int inDimIdx = dimIdx - (outRank - inRank);
	264	+ int64_t inDim = (inDimIdx >= 0) ? inShape[inDimIdx] : 1;
	265	+ int64_t inCoord = (inDim == 1) ? 0 : coord;
	266	+ int inStride = 1;
	267	+ for (int dd = inRank - 1; dd > inDimIdx; dd--) inStride *= inShape[dd];
	268	+ outIdx += inCoord * inStride;
	269	+ outStride *= dim;
	270	+ }
	271	+ return outIdx;
	272	+}
	273	+
	274	+std::vector<float> ComputeGolden(
	275	+ const std::vector<float>& x, const std::vector<int64_t>& shapeX,
	276	+ const std::vector<float>& y, const std::vector<int64_t>& shapeY,
	277	+ const std::vector<int64_t>& outShape)
	278	+{
	279	+ int64_t n = GetShapeSize(outShape);
	280	+ std::vector<float> result(n);
	281	+ for (int64_t i = 0; i < n; i++) {
	282	+ int64_t ix = BroadcastIdx(i, shapeX, outShape);
	283	+ int64_t iy = BroadcastIdx(i, shapeY, outShape);
	284	+ float fx = x[ix], fy = y[iy];
	285	+ if (fx == 0.0f) {
	286	+ result[i] = 0.0f;
	287	+ } else {
	288	+ result[i] = fx * std::log1p(fy);
	289	+ }
	290	+ }
	291	+ return result;
	292	+}
	293	+
	294	+int Init(int32_t deviceId, aclrtStream* stream)
	295	+{
	296	+ auto ret = aclInit(nullptr);
	297	+ CHECK_RET(ret == ACL_SUCCESS, "aclInit failed");
	298	+ ret = aclrtSetDevice(deviceId);
	299	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtSetDevice failed");
	300	+ ret = aclrtCreateStream(stream);
	301	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtCreateStream failed");
	302	+ return 0;
	303	+}
	304	+
	305	+template<typename T>
	306	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape,
	307	+ void deviceAddr, aclDataType dataType, aclTensor tensor)
	308	+{
	309	+ auto size = GetShapeSize(shape) * sizeof(T);
	310	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	311	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtMalloc failed");
	312	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	313	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtMemcpy H2D failed");
	314	+
	315	+ std::vector<int64_t> strides(shape.size(), 1);
	316	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	317	+ strides[i] = shape[i + 1] * strides[i + 1];
	318	+ }
	319	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType,
	320	+ strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	321	+ shape.data(), shape.size(), *deviceAddr);
	322	+ return 0;
	323	+}
	324	+
	325	+int RunXlog1py(const std::vector<int64_t>& shapeX, const std::vector<float>& dataX,
	326	+ const std::vector<int64_t>& shapeY, const std::vector<float>& dataY,
	327	+ const std::string& tag, aclrtStream stream)
	328	+{
	329	+ LOG_PRINT("--- Test %s ---", tag.c_str());
	330	+
	331	+ // Compute broadcast output shape
	332	+ int rank = std::max(shapeX.size(), shapeY.size());
	333	+ std::vector<int64_t> outShape(rank);
	334	+ for (int d = 0; d < rank; d++) {
	335	+ int dx = d - (rank - shapeX.size());
	336	+ int dy = d - (rank - shapeY.size());
	337	+ auto sx = (dx >= 0) ? shapeX[dx] : 1;
	338	+ auto sy = (dy >= 0) ? shapeY[dy] : 1;
	339	+ outShape[d] = std::max(sx, sy);
	340	+ }
	341	+ int64_t outSize = GetShapeSize(outShape);
	342	+
	343	+ // Compute golden
	344	+ auto golden = ComputeGolden(dataX, shapeX, dataY, shapeY, outShape);
	345	+
	346	+ // Allocate device tensors
	347	+ aclTensor* aclX = nullptr; void* devX = nullptr;
	348	+ auto ret = CreateAclTensor(dataX, shapeX, &devX, aclDataType::ACL_FLOAT, &aclX);
	349	+ CHECK_RET(ret == 0, "create tensor X failed");
	350	+
	351	+ aclTensor* aclY = nullptr; void* devY = nullptr;
	352	+ ret = CreateAclTensor(dataY, shapeY, &devY, aclDataType::ACL_FLOAT, &aclY);
	353	+ CHECK_RET(ret == 0, "create tensor Y failed");
	354	+
	355	+ std::vector<float> outHostData(outSize, 0);
	356	+ aclTensor* aclOut = nullptr; void* devOut = nullptr;
	357	+ ret = CreateAclTensor(outHostData, outShape, &devOut, aclDataType::ACL_FLOAT, &aclOut);
	358	+ CHECK_RET(ret == 0, "create tensor Out failed");
	359	+
	360	+ // Phase 1: GetWorkspaceSize
	361	+ uint64_t workspaceSize = 0;
	362	+ aclOpExecutor* executor = nullptr;
	363	+ ret = aclnnXlog1pyGetWorkspaceSize(aclX, aclY, aclOut, &workspaceSize, &executor);
	364	+ CHECK_RET(ret == ACL_SUCCESS, "aclnnXlog1pyGetWorkspaceSize failed");
	365	+
	366	+ void* workspaceAddr = nullptr;
	367	+ if (workspaceSize > 0) {
	368	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	369	+ CHECK_RET(ret == ACL_SUCCESS, "allocate workspace failed");
	370	+ }
	371	+
	372	+ // Phase 2: Execute
	373	+ ret = aclnnXlog1py(workspaceAddr, workspaceSize, executor, stream);
	374	+ CHECK_RET(ret == ACL_SUCCESS, "aclnnXlog1py execute failed");
	375	+
	376	+ ret = aclrtSynchronizeStream(stream);
	377	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtSynchronizeStream failed");
	378	+
	379	+ // Copy result back and compare
	380	+ std::vector<float> npuResult(outSize, 0);
	381	+ ret = aclrtMemcpy(npuResult.data(), outSize * sizeof(float), devOut,
	382	+ outSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
	383	+ CHECK_RET(ret == ACL_SUCCESS, "copy result D2H failed");
	384	+
	385	+ bool pass = true;
	386	+ for (int64_t i = 0; i < outSize; i++) {
	387	+ float g = golden[i];
	388	+ float r = npuResult[i];
	389	+ double mere = (std::fabs(g) > 1e-6) ?
	390	+ std::fabs(r - g) / std::fabs(g) : std::fabs(r - g);
	391	+ if (mere > 0.001) {
	392	+ LOG_PRINT(" [FAIL][%s][%ld] golden=%.6f npu=%.6f", tag.c_str(), i, g, r);
	393	+ pass = false;
	394	+ }
	395	+ }
	396	+ if (pass) LOG_PRINT(" [PASS][%s] all %ld elems OK", tag.c_str(), outSize);
	397	+
	398	+ // Cleanup
	399	+ aclDestroyTensor(aclX); aclDestroyTensor(aclY); aclDestroyTensor(aclOut);
	400	+ aclrtFree(devX); aclrtFree(devY); aclrtFree(devOut);
	401	+ if (workspaceSize > 0) aclrtFree(workspaceAddr);
	402	+
	403	+ return pass ? 0 : -1;
	404	+}
	405	+
	406	+int main()
	407	+{
	408	+ int32_t deviceId = 0;
	409	+ aclrtStream stream;
	410	+ auto ret = Init(deviceId, &stream);
	411	+ CHECK_RET(ret == 0, "Init failed");
	412	+
	413	+ int numPass = 0, numFail = 0;
	414	+
	415	+ // Test 1: same shape [1,2,4,4]
	416	+ {
	417	+ std::vector<int64_t> shape = {1, 2, 4, 4};
	418	+ std::vector<float> x(32), y(32);
	419	+ for (int i = 0; i < 32; i++) { x[i] = 2.0f; y[i] = 1.0f; }
	420	+ if (RunXlog1py(shape, x, shape, y, "same_shape", stream) == 0) numPass++; else numFail++;
	421	+ }
	422	+
	423	+ // Test 2: broadcast x=[1,2,1,4] y=[1,2,4,4]
	424	+ {
	425	+ std::vector<int64_t> shapeX = {1, 2, 1, 4};
	426	+ std::vector<int64_t> shapeY = {1, 2, 4, 4};
	427	+ std::vector<float> x(8), y(32);
	428	+ for (int i = 0; i < 8; i++) x[i] = 3.0f;
	429	+ for (int i = 0; i < 32; i++) y[i] = 2.0f;
	430	+ if (RunXlog1py(shapeX, x, shapeY, y, "broadcast", stream) == 0) numPass++; else numFail++;
	431	+ }
	432	+
	433	+ // Test 3: x == 0 boundary case
	434	+ {
	435	+ std::vector<int64_t> shape = {1, 1, 8, 8};
	436	+ std::vector<float> x(64, 0.0f);
	437	+ std::vector<float> y(64, 100.0f);
	438	+ if (RunXlog1py(shape, x, shape, y, "x_eq_0", stream) == 0) numPass++; else numFail++;
	439	+ }
	440	+
	441	+ LOG_PRINT("========================================");
	442	+ LOG_PRINT("ACLNN Xlog1py NPU results: PASS=%d FAIL=%d", numPass, numFail);
	443	+ LOG_PRINT("========================================");
	444	+
	445	+ aclrtDestroyStream(stream);
	446	+ aclrtResetDevice(deviceId);
	447	+ aclFinalize();
	448	+
	449	+ return (numFail == 0) ? 0 : -1;
	450	+}
	451	+```

Amath/xlog1py/examples/arch35/test_aclnn_xlog1py.cpp+273-0

@@ -0,0 +1,273 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+/**
	13	+ * @file test_aclnn_xlog1py.cpp
	14	+ * @brief aclnn xlog1py 算子 NPU 调用示例
	15	+ *
	16	+ * 算子: z = x * log1p(y), x==0 -> z=0
	17	+ * 支持 broadcast, FLOAT16/FLOAT/BF16
	18	+ */
	19	+
	20	+#include <iostream>
	21	+#include <vector>
	22	+#include <cmath>
	23	+#include "acl/acl.h"
	24	+#include "aclnnop/aclnn_xlog1py.h"
	25	+
	26	+#define CHECK_RET(cond, msg) \
	27	+ do { \
	28	+ if (!(cond)) { \
	29	+ printf("[FAIL] " msg "\n"); \
	30	+ return -1; \
	31	+ } \
	32	+ } while (0)
	33	+
	34	+#define LOG_PRINT(msg, ...) printf(msg "\n", ##__VA_ARGS__)
	35	+
	36	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
	37	+{
	38	+ int64_t size = 1;
	39	+ for (auto i : shape) size *= i;
	40	+ return size;
	41	+}
	42	+
	43	+// Broadcast index: map flat index in output to flat index in input
	44	+static int64_t BroadcastIdx(int64_t flat, const std::vector<int64_t>& inShape,
	45	+ const std::vector<int64_t>& outShape)
	46	+{
	47	+ int inRank = (int)inShape.size();
	48	+ int outRank = (int)outShape.size();
	49	+ int64_t outIdx = 0, outStride = 1;
	50	+ for (int d = 0; d < outRank; d++) {
	51	+ int dimIdx = outRank - 1 - d;
	52	+ int64_t dim = outShape[dimIdx];
	53	+ int64_t coord = (flat / outStride) % dim;
	54	+ int inDimIdx = dimIdx - (outRank - inRank);
	55	+ int64_t inDim = (inDimIdx >= 0) ? inShape[inDimIdx] : 1;
	56	+ int64_t inCoord = (inDim == 1) ? 0 : coord;
	57	+ int inStride = 1;
	58	+ for (int dd = inRank - 1; dd > inDimIdx; dd--) inStride *= inShape[dd];
	59	+ outIdx += inCoord * inStride;
	60	+ outStride *= dim;
	61	+ }
	62	+ return outIdx;
	63	+}
	64	+
	65	+std::vector<float> ComputeGolden(
	66	+ const std::vector<float>& x, const std::vector<int64_t>& shapeX,
	67	+ const std::vector<float>& y, const std::vector<int64_t>& shapeY,
	68	+ const std::vector<int64_t>& outShape)
	69	+{
	70	+ int64_t n = GetShapeSize(outShape);
	71	+ std::vector<float> result(n);
	72	+ for (int64_t i = 0; i < n; i++) {
	73	+ int64_t ix = BroadcastIdx(i, shapeX, outShape);
	74	+ int64_t iy = BroadcastIdx(i, shapeY, outShape);
	75	+ float fx = x[ix], fy = y[iy];
	76	+ if (fx == 0.0f) {
	77	+ result[i] = 0.0f;
	78	+ } else {
	79	+ result[i] = fx * std::log1p(fy);
	80	+ }
	81	+ }
	82	+ return result;
	83	+}
	84	+
	85	+template<typename T>
	86	+bool CompareResult(const std::vector<float>& golden, const std::vector<T>& npuResult,
	87	+ const std::vector<int64_t>& shape, const std::string& tag)
	88	+{
	89	+ int64_t n = GetShapeSize(shape);
	90	+ bool allPass = true;
	91	+ double maxMere = 0.0;
	92	+ for (int64_t i = 0; i < n; i++) {
	93	+ float g = golden[i];
	94	+ float r = static_cast<float>(npuResult[i]);
	95	+ double mere = 0.0;
	96	+ if (std::fabs(g) > 1e-6) {
	97	+ mere = std::fabs(static_cast<double>(r - g)) / std::fabs(static_cast<double>(g));
	98	+ } else {
	99	+ mere = std::fabs(static_cast<double>(r - g));
	100	+ }
	101	+ if (mere > maxMere) maxMere = mere;
	102	+ if (mere > 0.001) {
	103	+ LOG_PRINT(" [FAIL][%s][%ld] golden=%.6f npu=%.6f mere=%.6e", tag.c_str(), i, g, r, mere);
	104	+ allPass = false;
	105	+ }
	106	+ }
	107	+ if (allPass) {
	108	+ LOG_PRINT(" [PASS][%s] all %ld elems OK, max_mere=%.6e", tag.c_str(), n, maxMere);
	109	+ }
	110	+ return allPass;
	111	+}
	112	+
	113	+int Init(int32_t deviceId, aclrtStream* stream)
	114	+{
	115	+ auto ret = aclInit(nullptr);
	116	+ CHECK_RET(ret == ACL_SUCCESS, "aclInit failed");
	117	+ ret = aclrtSetDevice(deviceId);
	118	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtSetDevice failed");
	119	+ ret = aclrtCreateStream(stream);
	120	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtCreateStream failed");
	121	+ return 0;
	122	+}
	123	+
	124	+template<typename T>
	125	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape,
	126	+ void deviceAddr, aclDataType dataType, aclTensor tensor)
	127	+{
	128	+ auto size = GetShapeSize(shape) * sizeof(T);
	129	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	130	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtMalloc failed");
	131	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	132	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtMemcpy H2D failed");
	133	+
	134	+ std::vector<int64_t> strides(shape.size(), 1);
	135	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	136	+ strides[i] = shape[i + 1] * strides[i + 1];
	137	+ }
	138	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType,
	139	+ strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	140	+ shape.data(), shape.size(), *deviceAddr);
	141	+ return 0;
	142	+}
	143	+
	144	+int RunXlog1py(const std::vector<int64_t>& shapeX, const std::vector<float>& dataX,
	145	+ const std::vector<int64_t>& shapeY, const std::vector<float>& dataY,
	146	+ const std::string& tag, aclrtStream stream)
	147	+{
	148	+ LOG_PRINT("--- Test %s ---", tag.c_str());
	149	+
	150	+ // Compute broadcast output shape
	151	+ int rank = std::max(shapeX.size(), shapeY.size());
	152	+ std::vector<int64_t> outShape(rank);
	153	+ for (int d = 0; d < rank; d++) {
	154	+ int dx = d - (rank - shapeX.size());
	155	+ int dy = d - (rank - shapeY.size());
	156	+ auto sx = (dx >= 0) ? shapeX[dx] : 1;
	157	+ auto sy = (dy >= 0) ? shapeY[dy] : 1;
	158	+ outShape[d] = std::max(sx, sy);
	159	+ }
	160	+ int64_t outSize = GetShapeSize(outShape);
	161	+
	162	+ LOG_PRINT(" shapeX in=%ld outShape=[%ld,%ld,%ld,%ld]",
	163	+ dataX.size(), outShape[0], outShape[1], outShape[2], outShape[3]);
	164	+
	165	+ // Compute golden
	166	+ auto golden = ComputeGolden(dataX, shapeX, dataY, shapeY, outShape);
	167	+
	168	+ // Allocate device tensors
	169	+ aclTensor* aclX = nullptr; void* devX = nullptr;
	170	+ auto ret = CreateAclTensor(dataX, shapeX, &devX, aclDataType::ACL_FLOAT, &aclX);
	S S sunday5月30日已过期测试示例只覆盖了 FLOAT dtype，没有 FP16 和 BF16 的测试用例。作为声明支持 3 种 dtype 的新算子，测试应覆盖所有 dtype，尤其是 FP16/BF16 的 cast 逻辑和 NaN/零值边界场景。
	171	+ CHECK_RET(ret == 0, "create tensor X failed");
	172	+
	173	+ aclTensor* aclY = nullptr; void* devY = nullptr;
	174	+ ret = CreateAclTensor(dataY, shapeY, &devY, aclDataType::ACL_FLOAT, &aclY);
	175	+ CHECK_RET(ret == 0, "create tensor Y failed");
	176	+
	177	+ std::vector<float> outHostData(outSize, 0);
	178	+ aclTensor* aclOut = nullptr; void* devOut = nullptr;
	179	+ ret = CreateAclTensor(outHostData, outShape, &devOut, aclDataType::ACL_FLOAT, &aclOut);
	180	+ CHECK_RET(ret == 0, "create tensor Out failed");
	181	+
	182	+ // Phase 1: GetWorkspaceSize
	183	+ uint64_t workspaceSize = 0;
	184	+ aclOpExecutor* executor = nullptr;
	185	+ ret = aclnnXlog1pyGetWorkspaceSize(aclX, aclY, aclOut, &workspaceSize, &executor);
	186	+ CHECK_RET(ret == ACL_SUCCESS, "aclnnXlog1pyGetWorkspaceSize failed");
	187	+
	188	+ void* workspaceAddr = nullptr;
	189	+ if (workspaceSize > 0) {
	190	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	191	+ CHECK_RET(ret == ACL_SUCCESS, "allocate workspace failed");
	192	+ }
	193	+
	194	+ // Phase 2: Execute
	195	+ ret = aclnnXlog1py(workspaceAddr, workspaceSize, executor, stream);
	196	+ CHECK_RET(ret == ACL_SUCCESS, "aclnnXlog1py execute failed");
	197	+
	198	+ ret = aclrtSynchronizeStream(stream);
	199	+ CHECK_RET(ret == ACL_SUCCESS, "aclrtSynchronizeStream failed");
	200	+
	201	+ // Copy result back
	202	+ std::vector<float> npuResult(outSize, 0);
	203	+ ret = aclrtMemcpy(npuResult.data(), outSize * sizeof(float), devOut,
	204	+ outSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
	205	+ CHECK_RET(ret == ACL_SUCCESS, "copy result D2H failed");
	206	+
	207	+ // Compare
	208	+ bool pass = CompareResult(golden, npuResult, outShape, tag);
	209	+
	210	+ // Cleanup
	211	+ aclDestroyTensor(aclX); aclDestroyTensor(aclY); aclDestroyTensor(aclOut);
	212	+ aclrtFree(devX); aclrtFree(devY); aclrtFree(devOut);
	213	+ if (workspaceSize > 0) aclrtFree(workspaceAddr);
	214	+
	215	+ return pass ? 0 : -1;
	216	+}
	217	+
	218	+int main()
	219	+{
	220	+ int32_t deviceId = 0;
	221	+ aclrtStream stream;
	222	+ auto ret = Init(deviceId, &stream);
	223	+ CHECK_RET(ret == 0, "Init failed");
	224	+
	225	+ int numPass = 0, numFail = 0;
	226	+
	227	+ // Test 1: same shape [1,2,4,4]
	228	+ {
	229	+ std::vector<int64_t> shape = {1, 2, 4, 4};
	230	+ std::vector<float> x(32), y(32);
	231	+ for (int i = 0; i < 32; i++) { x[i] = 2.0f; y[i] = 1.0f; }
	232	+ if (RunXlog1py(shape, x, shape, y, "same_shape", stream) == 0) numPass++; else numFail++;
	233	+ }
	234	+
	235	+ // Test 2: broadcast x=[1,2,1,4] y=[1,2,4,4]
	236	+ {
	237	+ std::vector<int64_t> shapeX = {1, 2, 1, 4};
	238	+ std::vector<int64_t> shapeY = {1, 2, 4, 4};
	239	+ std::vector<float> x(8), y(32);
	240	+ for (int i = 0; i < 8; i++) x[i] = 3.0f;
	241	+ for (int i = 0; i < 32; i++) y[i] = 2.0f;
	242	+ if (RunXlog1py(shapeX, x, shapeY, y, "broadcast", stream) == 0) numPass++; else numFail++;
	243	+ }
	244	+
	245	+ // Test 3: x == 0 boundary case
	246	+ {
	247	+ std::vector<int64_t> shape = {1, 1, 8, 8};
	248	+ std::vector<float> x(64, 0.0f);
	249	+ std::vector<float> y(64, 100.0f);
	250	+ if (RunXlog1py(shape, x, shape, y, "x_eq_0", stream) == 0) numPass++; else numFail++;
	251	+ }
	252	+
	253	+ // Test 4: scalar broadcast x=scalar, y=[4,8,16,16]
	254	+ {
	255	+ std::vector<int64_t> shapeX = {1};
	256	+ std::vector<int64_t> shapeY = {4, 8, 16, 16};
	257	+ std::vector<float> x(1, 2.5f);
	258	+ int64_t n = 8192;
	259	+ std::vector<float> y(n);
	260	+ for (int64_t i = 0; i < n; i++) y[i] = 1.0f + 0.1f * (i % 5);
	261	+ if (RunXlog1py(shapeX, x, shapeY, y, "scalar_broadcast", stream) == 0) numPass++; else numFail++;
	262	+ }
	263	+
	264	+ LOG_PRINT("========================================");
	265	+ LOG_PRINT("ACLNN Xlog1py NPU results: PASS=%d FAIL=%d", numPass, numFail);
	266	+ LOG_PRINT("========================================");
	267	+
	268	+ aclrtDestroyStream(stream);
	269	+ aclrtResetDevice(deviceId);
	270	+ aclFinalize();
	271	+
	272	+ return (numFail == 0) ? 0 : -1;
	273	+}

Amath/xlog1py/examples/test_geir_xlog1py.cpp+314-0

@@ -0,0 +1,314 @@
	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	+/**
	13	+ * @file test_geir_xlog1py.cpp
	14	+ * @brief GE IR 图模式调用示例 - Xlog1py 算子
	15	+ *
	16	+ * z = x * log1p(y), x==0 -> z=0
	17	+ * 测试流程: GEInitialize -> 建图 -> AddGraph -> RunGraph -> 比对精度 -> GEFinalize
	18	+ */
	19	+
	20	+#include <iostream>
	21	+#include <fstream>
	22	+#include <string.h>
	23	+#include <stdint.h>
	24	+#include <vector>
	25	+#include <string>
	26	+#include <map>
	27	+#include <cmath>
	28	+#include "assert.h"
	29	+
	30	+#include "graph.h"
	31	+#include "types.h"
	32	+#include "tensor.h"
	33	+#include "ge_error_codes.h"
	34	+#include "ge_api_types.h"
	35	+#include "ge_api.h"
	36	+#include "ge_ir_build.h"
	37	+#include "array_ops.h"
	38	+
	39	+#include "../op_graph/xlog1py_proto.h"
	40	+
	41	+#define FAILED -1
	42	+#define SUCCESS 0
	43	+
	44	+using namespace ge;
	45	+using std::map;
	46	+using std::string;
	47	+using std::vector;
	48	+
	49	+static std::string GetTime()
	50	+{
	51	+ time_t timep;
	52	+ time(&timep);
	53	+ char tmp[64];
	54	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
	55	+ return tmp;
	56	+}
	57	+
	58	+static uint32_t GetDtypeSize(DataType dt)
	59	+{
	60	+ switch (dt) {
	61	+ case ge::DT_FLOAT: case ge::DT_INT32: case ge::DT_UINT32: return 4;
	62	+ case ge::DT_FLOAT16: case ge::DT_BF16: case ge::DT_INT16: case ge::DT_UINT16: return 2;
	63	+ case ge::DT_INT64: case ge::DT_UINT64: case ge::DT_DOUBLE: return 8;
	64	+ default: return 1;
	65	+ }
	66	+}
	67	+
	68	+static Status GenFloatData(const vector<int64_t>& shapes, Tensor& tensor,
	69	+ TensorDesc& desc, const vector<float>& values)
	70	+{
	71	+ desc.SetRealDimCnt(shapes.size());
	72	+ int64_t size = 1;
	73	+ for (auto s : shapes) size *= s;
	74	+ uint32_t dataLen = size * sizeof(float);
	75	+ float* pData = new (std::nothrow) float[size];
	76	+ for (int64_t i = 0; i < size; i++) pData[i] = values[i % values.size()];
	77	+ tensor = Tensor(desc, reinterpret_cast<uint8_t*>(pData), dataLen);
	78	+ return SUCCESS;
	79	+}
	80	+
	81	+static Status GenScalarData(const vector<int64_t>& shapes, Tensor& tensor,
	82	+ TensorDesc& desc, float value)
	83	+{
	84	+ desc.SetRealDimCnt(shapes.size());
	85	+ int64_t size = 1;
	86	+ for (auto s : shapes) size *= s;
	87	+ uint32_t dataLen = size * sizeof(float);
	88	+ float* pData = new (std::nothrow) float[size];
	89	+ for (int64_t i = 0; i < size; i++) pData[i] = value;
	90	+ tensor = Tensor(desc, reinterpret_cast<uint8_t*>(pData), dataLen);
	91	+ return SUCCESS;
	92	+}
	93	+
	94	+// CPU Golden: z = x * log1p(y), x==0 -> z=0
	95	+static std::vector<float> ComputeGolden(const vector<float>& x, const vector<int64_t>& shapeX,
	96	+ const vector<float>& y, const vector<int64_t>& shapeY, const vector<int64_t>& outShape)
	97	+{
	98	+ auto broadcastIdx = [](int64_t flat, const vector<int64_t>& inShape, const vector<int64_t>& outShape) {
	99	+ int inRank = (int)inShape.size(), outRank = (int)outShape.size();
	100	+ int64_t outIdx = 0, outStride = 1;
	101	+ for (int d = 0; d < outRank; d++) {
	102	+ int dimIdx = outRank - 1 - d;
	103	+ int64_t dim = outShape[dimIdx];
	104	+ int64_t coord = (flat / outStride) % dim;
	105	+ int inDimIdx = dimIdx - (outRank - inRank);
	106	+ int64_t inDim = (inDimIdx >= 0) ? inShape[inDimIdx] : 1;
	107	+ int64_t inCoord = (inDim == 1) ? 0 : coord;
	108	+ int inStride = 1;
	109	+ for (int dd = inRank - 1; dd > inDimIdx; dd--) inStride *= inShape[dd];
	110	+ outIdx += inCoord * inStride;
	111	+ outStride *= dim;
	112	+ }
	113	+ return outIdx;
	114	+ };
	115	+ int64_t n = 1;
	116	+ for (auto d : outShape) n *= d;
	117	+ vector<float> result(n);
	118	+ for (int64_t i = 0; i < n; i++) {
	119	+ int64_t ix = broadcastIdx(i, shapeX, outShape);
	120	+ int64_t iy = broadcastIdx(i, shapeY, outShape);
	121	+ float fx = x[ix], fy = y[iy];
	122	+ result[i] = (fx == 0.0f) ? 0.0f : fx * std::log1p(fy);
	123	+ }
	124	+ return result;
	125	+}
	126	+
	127	+static bool CompareOutput(const vector<float>& golden, const vector<float>& output,
	128	+ int64_t n, const string& tag)
	129	+{
	130	+ bool allPass = true;
	131	+ double maxMere = 0.0;
	132	+ for (int64_t i = 0; i < n; i++) {
	133	+ float g = golden[i], r = output[i];
	134	+ double mere = (std::fabs(g) > 1e-6) ? std::fabs(r - g) / std::fabs(g) : std::fabs(r - g);
	135	+ if (mere > maxMere) maxMere = mere;
	136	+ if (mere > 0.001) {
	137	+ printf(" [FAIL][%s][%ld] golden=%.6f geir=%.6f mere=%.6e\n",
	138	+ tag.c_str(), i, g, r, mere);
	139	+ allPass = false;
	140	+ }
	141	+ }
	142	+ if (allPass) {
	143	+ printf(" [PASS][%s] all %ld elems OK, max_mere=%.6e\n", tag.c_str(), n, maxMere);
	144	+ }
	145	+ return allPass;
	146	+}
	147	+
	148	+static int RunGEIRTest(const string& tag,
	149	+ const vector<int64_t>& shapeX, const vector<float>& dataX,
	150	+ const vector<int64_t>& shapeY, const vector<float>& dataY)
	151	+{
	152	+ printf("--- Test GEIR %s ---\n", tag.c_str());
	153	+
	154	+ // Compute output shape (broadcast)
	155	+ int rank = std::max((int)shapeX.size(), (int)shapeY.size());
	156	+ vector<int64_t> outShape(rank);
	157	+ for (int d = 0; d < rank; d++) {
	158	+ int dx = d - (rank - shapeX.size());
	159	+ int dy = d - (rank - shapeY.size());
	160	+ outShape[d] = std::max(dx >= 0 ? shapeX[dx] : 1, dy >= 0 ? shapeY[dy] : 1);
	161	+ }
	162	+ int64_t outSize = 1;
	163	+ for (auto d : outShape) outSize *= d;
	164	+ auto golden = ComputeGolden(dataX, shapeX, dataY, shapeY, outShape);
	165	+
	166	+ // Build graph
	167	+ Graph graph("tc_ge_irrun_xlog1py");
	168	+ vector<ge::Tensor> input;
	169	+ vector<Operator> inputs;
	170	+ vector<Operator> outputs;
	171	+ Status ret;
	172	+
	173	+ // Input X — Data node with index 0
	174	+ auto opX = op::Data("inputX").set_attr_index(0);
	175	+ TensorDesc descX(ge::Shape(shapeX), FORMAT_ND, DT_FLOAT);
	S S sunday5月30日图模式测试同样只用了 DT_FLOAT，缺少 FP16 和 BF16 的图模式测试覆盖。
	176	+ descX.SetPlacement(ge::kPlacementHost);
	177	+ descX.SetFormat(FORMAT_ND);
	178	+ Tensor tensorX;
	179	+ ret = GenFloatData(shapeX, tensorX, descX, dataX);
	180	+ if (ret != SUCCESS) { printf("Gen input X data failed\n"); return FAILED; }
	181	+ opX.update_input_desc_x(descX);
	182	+ input.push_back(tensorX);
	183	+ graph.AddOp(opX);
	184	+ inputs.push_back(opX);
	185	+
	186	+ // Input Y — Data node with index 1
	187	+ auto opY = op::Data("inputY").set_attr_index(1);
	188	+ TensorDesc descY(ge::Shape(shapeY), FORMAT_ND, DT_FLOAT);
	189	+ descY.SetPlacement(ge::kPlacementHost);
	190	+ descY.SetFormat(FORMAT_ND);
	191	+ Tensor tensorY;
	192	+ ret = GenFloatData(shapeY, tensorY, descY, dataY);
	193	+ if (ret != SUCCESS) { printf("Gen input Y data failed\n"); return FAILED; }
	194	+ opY.update_input_desc_x(descY);
	195	+ input.push_back(tensorY);
	196	+ graph.AddOp(opY);
	197	+ inputs.push_back(opY);
	198	+
	199	+ // Xlog1py operator
	200	+ auto xlog1py = op::Xlog1py("xlog1py_op");
	201	+ xlog1py.set_input_x(opX);
	202	+ xlog1py.set_input_y(opY);
	203	+
	204	+ TensorDesc descZ(ge::Shape(outShape), FORMAT_ND, DT_FLOAT);
	205	+ xlog1py.update_output_desc_z(descZ);
	206	+ outputs.push_back(xlog1py);
	207	+
	208	+ graph.SetInputs(inputs).SetOutputs(outputs);
	209	+
	210	+ // Run graph
	211	+ map<AscendString, AscendString> build_options;
	212	+ ge::Session* session = new Session(build_options);
	213	+ if (session == nullptr) { printf("Create session failed\n"); return FAILED; }
	214	+
	215	+ map<AscendString, AscendString> graph_options;
	216	+ uint32_t graphId = 0;
	217	+ ret = session->AddGraph(graphId, graph, graph_options);
	218	+ if (ret != SUCCESS) {
	219	+ printf("AddGraph failed, ret=%d\n", (int)ret);
	220	+ ge::AscendString errMsg = ge::GEGetErrorMsgV2();
	221	+ std::string errStr(errMsg.GetString());
	222	+ printf("GE Error: %s\n", errStr.c_str());
	223	+ delete session;
	224	+ return FAILED;
	225	+ }
	226	+
	227	+ vector<ge::Tensor> output;
	228	+ ret = session->RunGraph(graphId, input, output);
	229	+ if (ret != SUCCESS) {
	230	+ ge::AscendString errMsg = ge::GEGetErrorMsgV2();
	231	+ printf("RunGraph failed, ret=%d\n", (int)ret);
	232	+ printf("GE Error: %s\n", errMsg.GetString());
	233	+ delete session;
	234	+ return FAILED;
	235	+ }
	236	+ printf("RunGraph success\n");
	237	+
	238	+ // Read output
	239	+ if (output.empty()) { printf("No output\n"); delete session; return FAILED; }
	240	+ int64_t outputShapeSize = output[0].GetTensorDesc().GetShape().GetShapeSize();
	241	+ uint32_t dataSize = outputShapeSize * GetDtypeSize(output[0].GetTensorDesc().GetDataType());
	242	+ uint8_t* outputData = output[0].GetData();
	243	+ vector<float> npuResult(outputShapeSize);
	244	+ memcpy(npuResult.data(), outputData, dataSize);
	245	+
	246	+ // Compare
	247	+ bool pass = CompareOutput(golden, npuResult, outSize, tag);
	248	+
	249	+ delete session;
	250	+ return pass ? SUCCESS : FAILED;
	251	+}
	252	+
	253	+int main(int argc, char* argv[])
	254	+{
	255	+ printf("%s - INFO - [XIR]: Start to initialize GE\n", GetTime().c_str());
	256	+ map<AscendString, AscendString> globalOptions = {
	257	+ {"ge.exec.deviceId", "0"},
	258	+ {"ge.graphRunMode", "1"}
	259	+ };
	260	+ Status ret = ge::GEInitialize(globalOptions);
	261	+ if (ret != SUCCESS) {
	262	+ printf("%s - ERROR - [XIR]: GEInitialize failed\n", GetTime().c_str());
	263	+ return FAILED;
	264	+ }
	265	+ printf("%s - INFO - [XIR]: GEInitialize success\n", GetTime().c_str());
	266	+
	267	+ int numPass = 0, numFail = 0;
	268	+
	269	+ // Test 1: same shape
	270	+ {
	271	+ vector<int64_t> shape = {1, 2, 4, 4};
	272	+ vector<float> x(32, 2.0f), y(32, 1.0f);
	273	+ if (RunGEIRTest("same_shape", shape, x, shape, y) == SUCCESS) numPass++; else numFail++;
	274	+ }
	275	+
	276	+ // Test 2: broadcast
	277	+ {
	278	+ vector<int64_t> shapeX = {1, 2, 1, 4};
	279	+ vector<int64_t> shapeY = {1, 2, 4, 4};
	280	+ vector<float> x(8, 3.0f), y(32, 2.0f);
	281	+ if (RunGEIRTest("broadcast", shapeX, x, shapeY, y) == SUCCESS) numPass++; else numFail++;
	282	+ }
	283	+
	284	+ // Test 3: x == 0 boundary
	285	+ {
	286	+ vector<int64_t> shape = {1, 1, 8, 8};
	287	+ vector<float> x(64, 0.0f), y(64, 100.0f);
	288	+ if (RunGEIRTest("x_eq_0", shape, x, shape, y) == SUCCESS) numPass++; else numFail++;
	289	+ }
	290	+
	291	+ // Test 4: scalar broadcast
	292	+ {
	293	+ vector<int64_t> shapeX = {1};
	294	+ vector<int64_t> shapeY = {4, 8, 16, 16};
	295	+ vector<float> x(1, 2.5f);
	296	+ int64_t n = 8192;
	297	+ vector<float> y(n);
	298	+ for (int64_t i = 0; i < n; i++) y[i] = 1.0f + 0.1f * (i % 5);
	299	+ if (RunGEIRTest("scalar_broadcast", shapeX, x, shapeY, y) == SUCCESS) numPass++; else numFail++;
	300	+ }
	301	+
	302	+ printf("========================================\n");
	303	+ printf("GEIR Xlog1py NPU results: PASS=%d FAIL=%d\n", numPass, numFail);
	304	+ printf("========================================\n");
	305	+
	306	+ printf("%s - INFO - [XIR]: Finalize GE\n", GetTime().c_str());
	307	+ ret = ge::GEFinalize();
	308	+ if (ret != SUCCESS) {
	309	+ printf("%s - ERROR - [XIR]: GEFinalize failed\n", GetTime().c_str());
	310	+ return FAILED;
	311	+ }
	312	+ printf("%s - INFO - [XIR]: GEFinalize success\n", GetTime().c_str());
	313	+ return (numFail == 0) ? SUCCESS : FAILED;
	314	+}

Amath/xlog1py/op_api/aclnn_xlog1py.cpp+156-0

@@ -0,0 +1,156 @@
	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	+/**
	13	+ * @file aclnn_xlog1py.cpp
	14	+ * @brief ACLNN L2 API 实现 - xlog1py 算子
	15	+ *
	16	+ * 标准两段式流程:
	17	+ * 1. GetWorkspaceSize: 参数检查 → Contiguous → L0算子 → ViewCopy → 返回workspace
	18	+ * 2. Execute: 执行计算
	19	+ */
	20	+
	21	+#include "aclnn_xlog1py.h"
	22	+#include "xlog1py.h"
	23	+#include "aclnn_kernels/contiguous.h"
	24	+#include "aclnn_kernels/common/op_error_check.h"
	25	+#include "opdev/op_log.h"
	26	+#include "opdev/op_dfx.h"
	27	+#include "opdev/common_types.h"
	28	+#include "opdev/data_type_utils.h"
	29	+#include "opdev/make_op_executor.h"
	30	+
	31	+using namespace op;
	32	+
	33	+#define ACLNN_MAX_SHAPE_RANK 8
	34	+
	35	+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
	36	+ DataType::DT_FLOAT, DataType::DT_FLOAT16, DataType::DT_BF16
	37	+};
	38	+
	39	+static bool IsDtypeSupported(DataType dtype)
	40	+{
	41	+ return CheckType(dtype, AICORE_DTYPE_SUPPORT_LIST);
	42	+}
	43	+
	44	+static bool HasEmptyTensor(const aclTensor* x, const aclTensor* y)
	45	+{
	46	+ return x->IsEmpty() \|\| y->IsEmpty();
	47	+}
	48	+
	49	+static bool CheckNotNull(const aclTensor* x, const aclTensor* y, const aclTensor* z)
	50	+{
	51	+ OP_CHECK_NULL(x, return false);
	52	+ OP_CHECK_NULL(y, return false);
	53	+ OP_CHECK_NULL(z, return false);
	54	+ return true;
	55	+}
	56	+
	57	+static bool CheckDtypeValid(const aclTensor* x, const aclTensor* y, const aclTensor* z)
	58	+{
	59	+ OP_CHECK_DTYPE_NOT_MATCH(x, y->GetDataType(), return false);
	60	+ OP_CHECK_DTYPE_NOT_MATCH(z, x->GetDataType(), return false);
	61	+
	62	+ OP_CHECK(IsDtypeSupported(x->GetDataType()),
	63	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	64	+ "Dtype not supported: dtype=%d. Supported: FLOAT, FLOAT16, BF16.",
	65	+ static_cast<int>(x->GetDataType())),
	66	+ return false);
	67	+ return true;
	68	+}
	69	+
	70	+static bool CheckFormat(const aclTensor* x, const aclTensor* y, const aclTensor* z)
	71	+{
	72	+ auto fmtX = x->GetStorageFormat();
	73	+ auto fmtY = y->GetStorageFormat();
	74	+ auto fmtZ = z->GetStorageFormat();
	75	+
	76	+ OP_CHECK(!(IsPrivateFormat(fmtX) \|\| IsPrivateFormat(fmtY) \|\| IsPrivateFormat(fmtZ)),
	77	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	78	+ "Private format not supported: x=%d, y=%d, z=%d",
	79	+ static_cast<int>(fmtX), static_cast<int>(fmtY), static_cast<int>(fmtZ)),
	80	+ return false);
	81	+ return true;
	82	+}
	83	+
	84	+static bool CheckShape(const aclTensor* x, const aclTensor* y, const aclTensor* z)
	85	+{
	86	+ OP_CHECK_MAX_DIM(x, ACLNN_MAX_SHAPE_RANK, return false);
	87	+ OP_CHECK_MAX_DIM(y, ACLNN_MAX_SHAPE_RANK, return false);
	88	+ OP_CHECK_MAX_DIM(z, ACLNN_MAX_SHAPE_RANK, return false);
	89	+ return true;
	90	+}
	91	+
	92	+static aclnnStatus CheckParams(const aclTensor* x, const aclTensor* y, const aclTensor* z)
	93	+{
	94	+ CHECK_COND(CheckNotNull(x, y, z), ACLNN_ERR_PARAM_NULLPTR, "CheckNotNull failed");
	95	+ CHECK_COND(CheckDtypeValid(x, y, z), ACLNN_ERR_PARAM_INVALID,
	96	+ "CheckDtypeValid failed: x_dtype=%d, y_dtype=%d, z_dtype=%d",
	97	+ static_cast<int>(x->GetDataType()), static_cast<int>(y->GetDataType()),
	98	+ static_cast<int>(z->GetDataType()));
	99	+ CHECK_COND(CheckFormat(x, y, z), ACLNN_ERR_PARAM_INVALID,
	100	+ "CheckFormat failed: x_fmt=%d, y_fmt=%d, z_fmt=%d",
	101	+ static_cast<int>(x->GetStorageFormat()), static_cast<int>(y->GetStorageFormat()),
	102	+ static_cast<int>(z->GetStorageFormat()));
	103	+ CHECK_COND(CheckShape(x, y, z), ACLNN_ERR_PARAM_INVALID,
	104	+ "CheckShape failed: x_dim=%zu, y_dim=%zu, z_dim=%zu",
	105	+ x->GetViewShape().GetDimNum(), y->GetViewShape().GetDimNum(),
	106	+ z->GetViewShape().GetDimNum());
	107	+ return ACLNN_SUCCESS;
	108	+}
	109	+
	110	+extern "C" aclnnStatus aclnnXlog1pyGetWorkspaceSize(
	111	+ const aclTensor* x,
	112	+ const aclTensor* y,
	113	+ const aclTensor* z,
	114	+ uint64_t* workspaceSize,
	115	+ aclOpExecutor** executor)
	116	+{
	117	+ L2_DFX_PHASE_1(aclnnXlog1py, DFX_IN(x, y), DFX_OUT(z));
	118	+
	119	+ auto uniqueExecutor = CREATE_EXECUTOR();
	120	+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
	121	+
	122	+ auto ret = CheckParams(x, y, z);
	123	+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
	124	+
	125	+ if (HasEmptyTensor(x, y)) {
	126	+ *workspaceSize = 0;
	127	+ uniqueExecutor.ReleaseTo(executor);
	128	+ return ACLNN_SUCCESS;
	129	+ }
	130	+
	131	+ auto xContiguous = l0op::Contiguous(x, uniqueExecutor.get());
	132	+ CHECK_RET(xContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	133	+
	134	+ auto yContiguous = l0op::Contiguous(y, uniqueExecutor.get());
	135	+ CHECK_RET(yContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	136	+
	137	+ const aclTensor* opResult = l0op::Xlog1py(xContiguous, yContiguous, uniqueExecutor.get());
	138	+ CHECK_RET(opResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
	139	+
	140	+ auto viewCopyResult = l0op::ViewCopy(opResult, z, uniqueExecutor.get());
	141	+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
	142	+
	143	+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
	144	+ uniqueExecutor.ReleaseTo(executor);
	145	+ return ACLNN_SUCCESS;
	146	+}
	147	+
	148	+extern "C" aclnnStatus aclnnXlog1py(
	149	+ void* workspace,
	150	+ uint64_t workspaceSize,
	151	+ aclOpExecutor* executor,
	152	+ aclrtStream stream)
	153	+{
	154	+ L2_DFX_PHASE_2(aclnnXlog1py);
	155	+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
	156	+}

Amath/xlog1py/op_api/aclnn_xlog1py.h+80-0

@@ -0,0 +1,80 @@
	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	+/**
	13	+ * @file aclnn_xlog1py.h
	14	+ * @brief ACLNN L2 API - xlog1py: z = x * log1p(y), with x==0 → z=0
	15	+ *
	16	+ * 两段式接口:
	17	+ * - aclnnXlog1pyGetWorkspaceSize: 计算 workspace
	18	+ * - aclnnXlog1py: 执行计算
	19	+ */
	20	+
	21	+#ifndef ACLNN_XLOG1PY_H_
	22	+#define ACLNN_XLOG1PY_H_
	23	+
	24	+#include "aclnn/aclnn_base.h"
	25	+
	26	+#ifndef ACLNN_API
	27	+#define ACLNN_API __attribute__((visibility("default")))
	28	+#endif
	29	+
	30	+#ifdef __cplusplus
	31	+extern "C" {
	32	+#endif
	33	+
	34	+/**
	35	+ * @brief aclnnXlog1py的第一段接口，根据具体的计算流程，计算workspace大小。
	36	+ * @domain aclnn_math
	37	+ *
	38	+ * 算子功能：完成 x * log1p(y) 计算，x 为 0 时输出 0。
	39	+ * 参数描述：
	40	+ * @param [in] x: npu
	41	+ * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BF16，shape需要与y满足broadcast关系。
	42	+ * 支持非连续的Tensor，数据格式支持ND，维度不大于8。
	43	+ * @param [in] y: npu
	44	+ * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BF16，shape需要与x满足broadcast关系。
	45	+ * 支持非连续的Tensor，数据格式支持ND，维度不大于8。
	46	+ * @param [in] z: npu
	47	+ * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BF16，shape为x与y broadcast之后的shape。
	48	+ * 支持非连续的Tensor，数据格式支持ND，维度不大于8。
	49	+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
	50	+ * @param [out] executor: 返回op执行器，包含了算子计算流程。
	51	+ * @return aclnnStatus: 返回状态码。
	52	+ */
	53	+ACLNN_API aclnnStatus aclnnXlog1pyGetWorkspaceSize(
	W W wangyongguang6月5日补充注释
	54	+ const aclTensor *x,
	55	+ const aclTensor *y,
	56	+ const aclTensor *z,
	57	+ uint64_t *workspaceSize,
	58	+ aclOpExecutor **executor);
	59	+
	60	+/**
	61	+ * @brief aclnnXlog1py的第二段接口，用于执行计算。
	62	+ *
	63	+ * 算子功能：完成 x * log1p(y) 计算，x 为 0 时输出 0。
	64	+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
	65	+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnXlog1pyGetWorkspaceSize获取。
	66	+ * @param [in] executor: op执行器，包含了算子计算流程。
	67	+ * @param [in] stream: acl stream流。
	68	+ * @return aclnnStatus: 返回状态码。
	69	+ */
	70	+ACLNN_API aclnnStatus aclnnXlog1py(
	71	+ void *workspace,
	72	+ uint64_t workspaceSize,
	73	+ aclOpExecutor *executor,
	74	+ aclrtStream stream);
	75	+
	76	+#ifdef __cplusplus
	77	+}
	78	+#endif
	79	+
	80	+#endif // ACLNN_XLOG1PY_H_

Amath/xlog1py/op_api/xlog1py.cpp+94-0

@@ -0,0 +1,94 @@
	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	+/**
	13	+ * @file xlog1py.cpp
	14	+ * @brief ACLNN L0 API 实现 - xlog1py 算子
	15	+ *
	16	+ * z = x * log1p(y), 特殊规则: x == 0 → z = 0
	17	+ * 2入1出 broadcast，支持 FLOAT16, FLOAT, BF16
	18	+ */
	19	+
	20	+#include "xlog1py.h"
	21	+#include "opdev/op_log.h"
	22	+#include "opdev/op_dfx.h"
	23	+#include "opdev/shape_utils.h"
	24	+#include "opdev/make_op_executor.h"
	25	+#include "opdev/platform.h"
	26	+
	27	+using namespace op;
	28	+
	29	+namespace l0op {
	30	+
	31	+OP_TYPE_REGISTER(Xlog1py);
	32	+
	33	+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
	34	+ DataType::DT_FLOAT, DataType::DT_FLOAT16, DataType::DT_BF16
	35	+};
	36	+
	37	+static bool IsAiCoreSupport(const aclTensor* x, const aclTensor* y)
	38	+{
	39	+ auto npuArch = GetCurrentPlatformInfo().GetCurNpuArch();
	40	+ OP_CHECK(npuArch == NpuArch::DAV_3510,
	41	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	42	+ "Xlog1py only supported on ascend950 (DAV_3510). npuArch=%d.",
	43	+ static_cast<int>(npuArch)),
	44	+ return false);
	45	+ OP_CHECK(CheckType(x->GetDataType(), AICORE_DTYPE_SUPPORT_LIST) &&
	46	+ CheckType(y->GetDataType(), AICORE_DTYPE_SUPPORT_LIST),
	47	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	48	+ "Xlog1py dtype not supported: x_dtype=%d, y_dtype=%d.",
	49	+ static_cast<int>(x->GetDataType()),
	50	+ static_cast<int>(y->GetDataType())),
	51	+ return false);
	52	+ return true;
	53	+}
	54	+
	55	+static bool Xlog1pyInferShape(const op::Shape& xShape, const op::Shape& yShape, op::Shape& outShape)
	56	+{
	57	+ OP_CHECK(BroadcastInferShape(xShape, yShape, outShape),
	58	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Xlog1py shape broadcast failed."),
	59	+ return false);
	60	+ return true;
	61	+}
	62	+
	63	+static const aclTensor* Xlog1pyAiCore(const aclTensor* x, const aclTensor* y,
	64	+ const aclTensor* out, aclOpExecutor* executor)
	65	+{
	66	+ L0_DFX(Xlog1pyAiCore, x, y, out);
	67	+
	68	+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(Xlog1py,
	69	+ OP_INPUT(x, y), OP_OUTPUT(out));
	70	+ OP_CHECK(ret == ACLNN_SUCCESS,
	71	+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "Xlog1pyAiCore failed."),
	72	+ return nullptr);
	73	+ return out;
	74	+}
	75	+
	76	+const aclTensor* Xlog1py(const aclTensor* x, const aclTensor* y, aclOpExecutor* executor)
	77	+{
	78	+ Shape outShape;
	79	+ const aclTensor* out = nullptr;
	80	+
	81	+ OP_CHECK(Xlog1pyInferShape(x->GetViewShape(), y->GetViewShape(), outShape),
	82	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Xlog1py infer shape failed."),
	83	+ return nullptr);
	84	+
	85	+ out = executor->AllocTensor(outShape, x->GetDataType());
	86	+
	87	+ OP_CHECK(IsAiCoreSupport(x, y),
	88	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Xlog1py IsAiCoreSupport check failed."),
	89	+ return nullptr);
	90	+
	91	+ return Xlog1pyAiCore(x, y, out, executor);
	92	+}
	93	+
	94	+} // namespace l0op

Amath/xlog1py/op_api/xlog1py.h+28-0

@@ -0,0 +1,28 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+/**
	13	+ * @file xlog1py.h
	14	+ * @brief ACLNN L0 API - xlog1py: z = x * log1p(y), with x==0 → z=0
	15	+ */
	16	+
	17	+#ifndef OP_API_INC_LEVEL0_XLOG1PY_H_
	18	+#define OP_API_INC_LEVEL0_XLOG1PY_H_
	19	+
	20	+#include "opdev/op_executor.h"
	21	+
	22	+namespace l0op {
	23	+
	24	+const aclTensor* Xlog1py(const aclTensor* x, const aclTensor* y, aclOpExecutor* executor);
	25	+
	26	+} // namespace l0op
	27	+
	28	+#endif // OP_API_INC_LEVEL0_XLOG1PY_H_

Amath/xlog1py/op_graph/xlog1py_graph_infer.cpp+25-0

@@ -0,0 +1,25 @@
	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	+#include "register/op_impl_registry.h"
	12	+#include "log/log.h"
	13	+
	14	+using namespace ge;
	15	+namespace ops {
	16	+static ge::graphStatus InferDataTypeForXlog1py(gert::InferDataTypeContext *context)
	17	+{
	18	+ OP_LOGI("Begin InferDataTypeForXlog1py");
	19	+ const ge::DataType xDataType = context->GetInputDataType(0);
	20	+ context->SetOutputDataType(0, xDataType);
	21	+ return ge::GRAPH_SUCCESS;
	22	+}
	23	+
	24	+IMPL_OP(Xlog1py).InferDataType(InferDataTypeForXlog1py);
	25	+} // namespace ops

Amath/xlog1py/op_graph/xlog1py_proto.h+47-0

@@ -0,0 +1,47 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file xlog1py_proto.h
	13	+ * \brief
	14	+ */
	15	+#ifndef OPS_OP_PROTO_INC_XLOG1PY_H_
	16	+#define OPS_OP_PROTO_INC_XLOG1PY_H_
	17	+
	18	+#include "graph/operator_reg.h"
	19	+#include "graph/types.h"
	20	+
	21	+namespace ge {
	22	+/**
	23	+@brief Computes x log1p(y) element-wise. Returns 0 if x == 0.
	24	+
	25	+* @par Inputs:
	26	+* Two inputs, including:
	27	+* x: An ND Tensor. Must be one of: float16, float32, bfloat16, float64, complex64, complex128.
	28	+* y: An ND Tensor. Must be one of: float16, float32, bfloat16, float64, complex64, complex128.
	29	+
	30	+* @par Outputs:
	31	+* z: An ND Tensor with broadcast shape of x and y. Has the same dtype as x and y.
	32	+
	33	+* @par Third-party framework compatibility
	34	+* Compatible with the TensorFlow/PyTorch operator xlog1py.
	35	+*/
	36	+REG_OP(Xlog1py)
	37	+ .INPUT(x, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT, DT_DOUBLE, DT_COMPLEX64,
	38	+ DT_COMPLEX128}))
	39	+ .INPUT(y, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT, DT_DOUBLE, DT_COMPLEX64,
	40	+ DT_COMPLEX128}))
	41	+ .OUTPUT(z, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT, DT_DOUBLE, DT_COMPLEX64,
	42	+ DT_COMPLEX128}))
	43	+ .OP_END_FACTORY_REG(Xlog1py)
	44	+
	45	+} // namespace ge
	46	+
	47	+#endif // OPS_OP_PROTO_INC_XLOG1PY_H_

Amath/xlog1py/op_host/arch35/xlog1py_tiling_arch35.cpp+325-0

@@ -0,0 +1,325 @@
	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	+// Xlog1py Tiling — arch35 实现
	13	+#include "xlog1py_tiling_arch35.h"
	14	+#include "../../op_kernel/arch35/xlog1py_struct.h"
	15	+#include <algorithm>
	16	+#include <sstream>
	17	+#include <graph/utils/type_utils.h>
	18	+#include "register/op_impl_registry.h"
	19	+#include "log/log.h"
	20	+#include "tiling/platform/platform_ascendc.h"
	21	+#include "platform/soc_spec.h"
	22	+
	23	+using namespace ge;
	24	+
	25	+namespace xlog1py {
	26	+
	27	+constexpr int64_t kPhysNodesFp32 = 3;
	28	+constexpr int64_t kPhysNodesFp16 = 4;
	29	+constexpr int64_t kBufDtypeSize = 4;
	30	+
	31	+bool PadAndSqueeze(
	32	+ const std::vector<std::vector<int64_t>>& input_shapes,
	33	+ const std::vector<std::vector<int64_t>>& output_shapes,
	34	+ std::vector<int64_t>& maximum_bro_shape,
	35	+ std::vector<std::vector<int64_t>>& normal_input_shapes,
	36	+ std::vector<std::vector<int64_t>>& normal_output_shapes)
	37	+{
	38	+ int64_t num_inputs = (int64_t)input_shapes.size();
	39	+ int64_t num_outputs = (int64_t)output_shapes.size();
	40	+ int64_t max_rank = 0;
	41	+ for (auto& s : input_shapes) max_rank = std::max(max_rank, (int64_t)s.size());
	42	+ for (auto& s : output_shapes) max_rank = std::max(max_rank, (int64_t)s.size());
	43	+ auto pad = [&](const std::vector<int64_t>& s) {
	44	+ std::vector<int64_t> p;
	45	+ p.assign(max_rank - (int64_t)s.size(), 1);
	46	+ p.insert(p.end(), s.begin(), s.end());
	47	+ return p;
	48	+ };
	49	+ std::vector<std::vector<int64_t>> padded_in(num_inputs), padded_out(num_outputs);
	50	+ for (int64_t i = 0; i < num_inputs; i++) padded_in[i] = pad(input_shapes[i]);
	51	+ for (int64_t i = 0; i < num_outputs; i++) padded_out[i] = pad(output_shapes[i]);
	52	+ maximum_bro_shape.clear();
	53	+ normal_input_shapes.assign(num_inputs, std::vector<int64_t>());
	54	+ normal_output_shapes.assign(num_outputs, std::vector<int64_t>());
	55	+ for (int64_t d = 0; d < max_rank; d++) {
	56	+ bool all_one = true;
	57	+ int64_t max_dim = 0;
	58	+ for (int64_t i = 0; i < num_inputs; i++) {
	59	+ if (padded_in[i][d] != 1) all_one = false;
	60	+ max_dim = std::max(max_dim, padded_in[i][d]);
	61	+ }
	62	+ for (int64_t i = 0; i < num_outputs; i++) {
	63	+ if (padded_out[i][d] != 1) all_one = false;
	64	+ max_dim = std::max(max_dim, padded_out[i][d]);
	65	+ }
	66	+ if (!all_one) {
	67	+ maximum_bro_shape.push_back(max_dim);
	68	+ for (int64_t i = 0; i < num_inputs; i++) normal_input_shapes[i].push_back(padded_in[i][d]);
	69	+ for (int64_t i = 0; i < num_outputs; i++) normal_output_shapes[i].push_back(padded_out[i][d]);
	70	+ }
	71	+ }
	72	+ if (maximum_bro_shape.empty()) {
	73	+ maximum_bro_shape.push_back(1);
	74	+ for (int64_t i = 0; i < num_inputs; i++) normal_input_shapes[i].push_back(1);
	75	+ for (int64_t i = 0; i < num_outputs; i++) normal_output_shapes[i].push_back(1);
	76	+ }
	77	+ return true;
	78	+}
	79	+
	80	+bool FindSplitAxis(const std::vector<int64_t>& max_bro_shape,
	81	+ int64_t dtype_size, int64_t ub_per_core, int64_t phys_nodes, SplitResult& out)
	82	+{
	83	+ int64_t per_buf_bytes = (ub_per_core / phys_nodes) & ~31LL;
	84	+ int64_t per_buf_elems = per_buf_bytes / dtype_size;
	85	+ int64_t rank = (int64_t)max_bro_shape.size();
	86	+ int64_t inner = 1;
	87	+ for (int64_t k = rank - 1; k >= 0; k--) {
	88	+ if (max_bro_shape[k] * inner > per_buf_elems) {
	89	+ out.a_i = per_buf_elems / inner;
	90	+ out.a_o = (max_bro_shape[k] + out.a_i - 1) / out.a_i;
	91	+ int64_t rem = max_bro_shape[k] % out.a_i;
	92	+ out.a_i_tail = (rem == 0) ? out.a_i : rem;
	93	+ out.axis = k;
	94	+ return true;
	95	+ }
	96	+ if (k == 0) {
	97	+ out.axis = 0; out.a_i = max_bro_shape[0];
	98	+ out.a_o = 1; out.a_i_tail = max_bro_shape[0];
	99	+ return true;
	100	+ }
	101	+ inner *= max_bro_shape[k];
	102	+ }
	103	+ return true;
	104	+}
	105	+
	106	+bool MultiCoreSplit(const std::vector<int64_t>& max_bro_shape,
	107	+ const SplitResult& ub_split, int64_t max_cores, MultiCoreResult& out)
	108	+{
	109	+ int64_t k = ub_split.axis, outer_prod = 1;
	110	+ for (int64_t j = 0; j < k; j++) outer_prod *= max_bro_shape[j];
	111	+ out.total_tiles = outer_prod * ub_split.a_o;
	112	+ out.num_cores = (out.total_tiles < max_cores) ? out.total_tiles : max_cores;
	113	+ out.tiles_main = out.total_tiles / out.num_cores;
	114	+ out.cores_tail = out.total_tiles % out.num_cores;
	115	+ return true;
	116	+}
	117	+
	118	+bool PrecomputeInputStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides) {
	119	+ int64_t rank = (int64_t)s.size();
	120	+ strides.assign(rank, 0);
	121	+ for (int64_t d = rank - 1; d >= 0; d--) {
	122	+ if (s[d] == 1) { strides[d] = 0; continue; }
	123	+ int64_t prod = 1;
	124	+ for (int64_t j = d + 1; j < rank; j++) prod *= s[j];
	125	+ strides[d] = prod;
	126	+ }
	127	+ return true;
	128	+}
	129	+
	130	+bool PrecomputeOutputStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides) {
	131	+ int64_t rank = (int64_t)s.size();
	132	+ strides.assign(rank, 0);
	133	+ for (int64_t d = rank - 1; d >= 0; d--) {
	134	+ if (s[d] == 1) { strides[d] = 0; continue; }
	135	+ int64_t prod = 1;
	136	+ for (int64_t j = d + 1; j < rank; j++) prod *= s[j];
	137	+ strides[d] = prod;
	138	+ }
	139	+ return true;
	140	+}
	141	+
	142	+} // namespace xlog1py
	143	+
	144	+namespace optiling {
	145	+
	146	+using namespace xlog1py;
	147	+
	148	+static std::string Arr2String(const int64_t* arr, int64_t n)
	149	+{
	150	+ std::ostringstream oss;
	151	+ oss << "[";
	152	+ if (n > 0) {
	153	+ for (int64_t i = 0; i < n - 1; ++i) oss << arr[i] << ",";
	154	+ oss << arr[n - 1];
	155	+ }
	156	+ oss << "]";
	157	+ return oss.str();
	158	+}
	159	+
	160	+Xlog1pyTiling::Xlog1pyTiling(gert::TilingContext* ctx) : ctx_(ctx) {}
	161	+
	162	+ge::graphStatus Xlog1pyTiling::GetShapeInfo()
	163	+{
	164	+ auto compileInfo = reinterpret_cast<const Xlog1pyCompileInfo*>(ctx_->GetCompileInfo());
	165	+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, compileInfo);
	166	+
	167	+ for (size_t i = 0; i < ctx_->GetComputeNodeInfo()->GetInputsNum(); ++i) {
	168	+ auto shape = ctx_->GetInputShape(i); OP_CHECK_NULL_WITH_CONTEXT(ctx_, shape);
	169	+ std::vector<int64_t> dims;
	170	+ gert::Shape s = shape->GetStorageShape();
	171	+ for (size_t d = 0; d < s.GetDimNum(); ++d) dims.push_back(s.GetDim(d));
	172	+ raw_input_shapes_.push_back(dims);
	173	+ }
	174	+ for (size_t i = 0; i < ctx_->GetComputeNodeInfo()->GetOutputsNum(); ++i) {
	175	+ auto shape = ctx_->GetOutputShape(i); OP_CHECK_NULL_WITH_CONTEXT(ctx_, shape);
	176	+ std::vector<int64_t> dims;
	177	+ gert::Shape s = shape->GetStorageShape();
	178	+ for (size_t d = 0; d < s.GetDimNum(); ++d) dims.push_back(s.GetDim(d));
	179	+ raw_output_shapes_.push_back(dims);
	180	+ }
	181	+
	182	+ auto inputDesc = ctx_->GetInputDesc(0); OP_CHECK_NULL_WITH_CONTEXT(ctx_, inputDesc);
	183	+ ge::DataType dtype = inputDesc->GetDataType();
	184	+ if (dtype == ge::DT_FLOAT) {
	185	+ dtype_size_ = 4; phys_nodes_ = kPhysNodesFp32;
	186	+ } else if (dtype == ge::DT_FLOAT16 \|\| dtype == ge::DT_BF16) {
	187	+ dtype_size_ = 2; phys_nodes_ = kPhysNodesFp16;
	188	+ } else {
	189	+ OP_LOGE(ctx_->GetNodeName(), "Xlog1py unsupported dtype");
	190	+ return GRAPH_FAILED;
	191	+ }
	192	+
	193	+ PadAndSqueeze(raw_input_shapes_, raw_output_shapes_,
	194	+ max_bro_shape_, normal_input_shapes_, normal_output_shapes_);
	195	+ rank_ = (int64_t)max_bro_shape_.size();
	196	+
	197	+ OP_LOGI(ctx_->GetNodeName(), "Xlog1py GetShapeInfo done rank %ld phys_nodes %ld ub %lu core %lu",
	198	+ rank_, phys_nodes_, compileInfo->ubSize, compileInfo->coreNum);
	199	+
	200	+ return GRAPH_SUCCESS;
	201	+}
	202	+
	203	+template<int64_t R>
	204	+ge::graphStatus Xlog1pyTiling::DoTilingAndSet()
	205	+{
	206	+ auto* tiling = ctx_->GetTilingData<Xlog1pyTilingData<R>>();
	207	+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, tiling);
	208	+
	209	+ auto* compileInfo = reinterpret_cast<const Xlog1pyCompileInfo*>(ctx_->GetCompileInfo());
	210	+ int64_t ub_per_core = (int64_t)compileInfo->ubSize;
	211	+
	212	+ int64_t per_buf_bytes = (ub_per_core / phys_nodes_) & ~31LL;
	213	+ tiling->per_buf_bytes = per_buf_bytes;
	214	+
	215	+ FindSplitAxis(max_bro_shape_, kBufDtypeSize, ub_per_core, phys_nodes_, tiling->split);
	216	+ MultiCoreSplit(max_bro_shape_, tiling->split, (int64_t)compileInfo->coreNum, tiling->multicore);
	217	+
	218	+ int64_t num_in = (int64_t)normal_input_shapes_.size();
	219	+ int64_t num_out = (int64_t)normal_output_shapes_.size();
	220	+ std::vector<std::vector<int64_t>> in_strides(num_in), out_strides(num_out);
	221	+ for (int64_t i = 0; i < num_in; i++) PrecomputeInputStrides(normal_input_shapes_[i], in_strides[i]);
	222	+ for (int64_t i = 0; i < num_out; i++) PrecomputeOutputStrides(normal_output_shapes_[i], out_strides[i]);
	223	+
	224	+ tiling->rank = rank_;
	225	+ int64_t delta = R - rank_;
	226	+
	227	+ for (int64_t d = 0; d < delta; d++) tiling->max_bro_shape[d] = 1;
	228	+ for (int64_t d = 0; d < rank_; d++) tiling->max_bro_shape[d + delta] = max_bro_shape_[d];
	229	+
	230	+ tiling->split.axis += delta;
	231	+ tiling->num_inputs = num_in;
	232	+ tiling->num_outputs = num_out;
	233	+
	234	+ for (int64_t i = 0; i < num_in; i++) {
	235	+ for (int64_t d = 0; d < delta; d++) {
	236	+ tiling->input_shapes[i][d] = 1;
	237	+ tiling->input_strides[i][d] = 0;
	238	+ }
	239	+ for (int64_t d = 0; d < rank_; d++) {
	240	+ tiling->input_shapes[i][d + delta] = normal_input_shapes_[i][d];
	241	+ tiling->input_strides[i][d + delta] = in_strides[i][d];
	242	+ }
	243	+ }
	244	+ for (int64_t i = num_in; i < kMaxInputSlots; i++)
	245	+ for (int64_t d = 0; d < R; d++) {
	246	+ tiling->input_shapes[i][d] = 1;
	247	+ tiling->input_strides[i][d] = 0;
	248	+ }
	249	+
	250	+ for (int64_t i = 0; i < num_out; i++) {
	251	+ for (int64_t d = 0; d < delta; d++) {
	252	+ tiling->output_shapes[i][d] = 1;
	253	+ tiling->output_strides[i][d] = 0;
	254	+ }
	255	+ for (int64_t d = 0; d < rank_; d++) {
	256	+ tiling->output_shapes[i][d + delta] = normal_output_shapes_[i][d];
	257	+ tiling->output_strides[i][d + delta] = out_strides[i][d];
	258	+ }
	259	+ }
	260	+ for (int64_t i = num_out; i < kMaxOutputSlots; i++)
	261	+ for (int64_t d = 0; d < R; d++) {
	262	+ tiling->output_shapes[i][d] = 1;
	263	+ tiling->output_strides[i][d] = 0;
	264	+ }
	265	+
	266	+ ctx_->SetBlockDim(tiling->multicore.num_cores);
	267	+
	268	+ OP_LOGI(ctx_->GetNodeName(), "Xlog1py TilingData: per_buf_bytes=%ld rank=%ld->R=%d "
	269	+ "max_bro_shape=%s split(axis=%ld a_i=%ld a_o=%ld a_i_tail=%ld) "
	270	+ "multi(cores=%ld tiles=%ld main=%ld core_tail=%ld)",
	271	+ tiling->per_buf_bytes, rank_, (int)R,
	272	+ Arr2String(tiling->max_bro_shape, R).c_str(),
	273	+ tiling->split.axis, tiling->split.a_i, tiling->split.a_o, tiling->split.a_i_tail,
	274	+ tiling->multicore.num_cores, tiling->multicore.total_tiles,
	275	+ tiling->multicore.tiles_main, tiling->multicore.cores_tail);
	276	+
	277	+ for (int64_t i = 0; i < num_in; i++)
	278	+ OP_LOGI(ctx_->GetNodeName(), "Xlog1py TilingData input[%ld]: shape=%s stride=%s",
	279	+ i, Arr2String(tiling->input_shapes[i], R).c_str(),
	280	+ Arr2String(tiling->input_strides[i], R).c_str());
	281	+
	282	+ return GRAPH_SUCCESS;
	283	+}
	284	+
	285	+ge::graphStatus Xlog1pyTiling::RunTiling()
	286	+{
	287	+ ge::graphStatus ret = GetShapeInfo();
	288	+ if (ret != GRAPH_SUCCESS) return ret;
	289	+
	290	+ int64_t mapped = (rank_ <= 4) ? 4 : 8;
	291	+ if (mapped == 4) {
	292	+ ret = DoTilingAndSet<4>();
	293	+ ctx_->SetTilingKey(GET_TPL_TILING_KEY(XLOG1PY_RANK_4));
	294	+ } else {
	295	+ ret = DoTilingAndSet<8>();
	296	+ ctx_->SetTilingKey(GET_TPL_TILING_KEY(XLOG1PY_RANK_8));
	297	+ }
	298	+ return ret;
	299	+}
	300	+
	301	+static ge::graphStatus TilingFuncXlog1py(gert::TilingContext* context)
	302	+{
	303	+ Xlog1pyTiling xlog1pyTiling(context);
	304	+ auto ret = xlog1pyTiling.RunTiling();
	305	+ if (ret != GRAPH_SUCCESS) return ret;
	306	+ size_t* workspaces = context->GetWorkspaceSizes(1);
	307	+ workspaces[0] = 16 * 1024 * 1024;
	308	+ return GRAPH_SUCCESS;
	309	+}
	310	+
	311	+ge::graphStatus TilingPrepareForXlog1py(gert::TilingParseContext* context)
	312	+{
	313	+ fe::PlatFormInfos* platformInfo = context->GetPlatformInfo();
	314	+ auto compileInfo = context->GetCompiledInfo<Xlog1pyCompileInfo>();
	315	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
	316	+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
	317	+ auto ap = platform_ascendc::PlatformAscendC(platformInfo);
	318	+ compileInfo->coreNum = ap.GetCoreNumAiv();
	319	+ ap.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfo->ubSize);
	320	+ return GRAPH_SUCCESS;
	321	+}
	322	+
	323	+IMPL_OP_OPTILING(Xlog1py).Tiling(TilingFuncXlog1py).TilingParse<Xlog1pyCompileInfo>(TilingPrepareForXlog1py);
	324	+
	325	+} // namespace optiling

Amath/xlog1py/op_host/arch35/xlog1py_tiling_arch35.h+87-0

@@ -0,0 +1,87 @@
	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	+// Xlog1py Tiling — arch35 头文件
	13	+#pragma once
	14	+#include <cstdint>
	15	+#include <vector>
	16	+#include <exe_graph/runtime/tiling_context.h>
	17	+#include "../../op_kernel/arch35/xlog1py_tiling_struct.h"
	18	+
	19	+namespace xlog1py {
	20	+
	21	+bool PadAndSqueeze(
	22	+ const std::vector<std::vector<int64_t>>& input_shapes,
	23	+ const std::vector<std::vector<int64_t>>& output_shapes,
	24	+ std::vector<int64_t>& maximum_bro_shape,
	25	+ std::vector<std::vector<int64_t>>& normal_input_shapes,
	26	+ std::vector<std::vector<int64_t>>& normal_output_shapes);
	27	+
	28	+bool FindSplitAxis(
	29	+ const std::vector<int64_t>& max_bro_shape,
	30	+ int64_t dtype_size, int64_t ub_per_core, int64_t phys_nodes,
	31	+ SplitResult& out);
	32	+
	33	+bool MultiCoreSplit(
	34	+ const std::vector<int64_t>& max_bro_shape,
	35	+ const SplitResult& ub_split, int64_t max_cores,
	36	+ MultiCoreResult& out);
	37	+
	38	+inline void GetCoreRange(int64_t core_id, int64_t tiles_main, int64_t cores_tail,
	39	+ int64_t& start, int64_t& end)
	40	+{
	41	+ if (core_id < cores_tail) {
	42	+ start = core_id * (tiles_main + 1);
	43	+ end = start + tiles_main + 1;
	44	+ } else {
	45	+ start = cores_tail * (tiles_main + 1) + (core_id - cores_tail) * tiles_main;
	46	+ end = start + tiles_main;
	47	+ }
	48	+}
	49	+
	50	+inline int64_t GetUBSplitRange(int64_t a_o_off, int64_t a_o, int64_t a_i, int64_t a_i_tail)
	51	+{
	52	+ return (a_o_off == a_o - 1) ? a_i_tail : a_i;
	53	+}
	54	+
	55	+bool PrecomputeInputStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides);
	56	+bool PrecomputeOutputStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides);
	57	+
	58	+} // namespace xlog1py
	59	+
	60	+namespace optiling {
	61	+
	62	+struct Xlog1pyCompileInfo {
	63	+ uint64_t coreNum;
	64	+ uint64_t ubSize;
	65	+};
	66	+
	67	+class Xlog1pyTiling {
	68	+public:
	69	+ explicit Xlog1pyTiling(gert::TilingContext* ctx);
	70	+ ge::graphStatus RunTiling();
	71	+
	72	+private:
	73	+ ge::graphStatus GetShapeInfo();
	74	+ template<int64_t R> ge::graphStatus DoTilingAndSet();
	75	+
	76	+ gert::TilingContext* ctx_;
	77	+ std::vector<std::vector<int64_t>> raw_input_shapes_;
	78	+ std::vector<std::vector<int64_t>> raw_output_shapes_;
	79	+ std::vector<int64_t> max_bro_shape_;
	80	+ std::vector<std::vector<int64_t>> normal_input_shapes_;
	81	+ std::vector<std::vector<int64_t>> normal_output_shapes_;
	82	+ int64_t dtype_size_ = 0;
	83	+ int64_t phys_nodes_ = 3;
	84	+ int64_t rank_ = 0;
	85	+};
	86	+
	87	+} // namespace optiling

Amath/xlog1py/op_host/config/ascend950/xlog1py_binary.json+35-0

@@ -0,0 +1,35 @@
	1	+{
	2	+ "op_type": "Xlog1py",
	3	+ "op_list": [
	4	+ {
	5	+ "bin_filename": "Xlog1py_FLOAT16",
	6	+ "inputs": [
	7	+ {"name": "x","index": 0,"dtype": "float16","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"},
	8	+ {"name": "y","index": 1,"dtype": "float16","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"}
	9	+ ],
	10	+ "outputs": [
	11	+ {"name": "z","index": 0,"dtype": "float16","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"}
	12	+ ]
	13	+ },
	14	+ {
	15	+ "bin_filename": "Xlog1py_FLOAT32",
	16	+ "inputs": [
	17	+ {"name": "x","index": 0,"dtype": "float32","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"},
	18	+ {"name": "y","index": 1,"dtype": "float32","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"}
	19	+ ],
	20	+ "outputs": [
	21	+ {"name": "z","index": 0,"dtype": "float32","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"}
	22	+ ]
	23	+ },
	24	+ {
	25	+ "bin_filename": "Xlog1py_BF16",
	26	+ "inputs": [
	27	+ {"name": "x","index": 0,"dtype": "bfloat16","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"},
	28	+ {"name": "y","index": 1,"dtype": "bfloat16","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"}
	29	+ ],
	30	+ "outputs": [
	31	+ {"name": "z","index": 0,"dtype": "bfloat16","format": "ND","paramType": "required","shape": [-2],"format_match_mode": "FormatAgnostic"}
	32	+ ]
	33	+ }
	34	+ ]
	35	+}

Amath/xlog1py/op_host/config/ascend950/xlog1py_simplified_key.ini+2-0

@@ -0,0 +1,2 @@
	1	+[Xlog1py]
	2	+default=0

Amath/xlog1py/op_host/xlog1py_def.cpp+47-0

@@ -0,0 +1,47 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+// Xlog1py OpDef — 2入1出 broadcast, x * log1p(y), x==0 → 0
	13	+#include "register/op_def_registry.h"
	14	+
	15	+namespace ops {
	16	+class Xlog1py : public OpDef {
	17	+public:
	18	+ explicit Xlog1py(const char* name) : OpDef(name)
	19	+ {
	20	+ this->Input("x").ParamType(REQUIRED)
	21	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
	22	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	23	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
	24	+ this->Input("y").ParamType(REQUIRED)
	25	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
	26	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	27	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
	28	+
	29	+ this->Output("z").ParamType(REQUIRED)
	30	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
	31	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	32	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
	33	+
	34	+ OpAICoreConfig aicoreConfig;
	35	+ aicoreConfig.DynamicCompileStaticFlag(true)
	36	+ .DynamicFormatFlag(false)
	37	+ .DynamicRankSupportFlag(true)
	38	+ .DynamicShapeSupportFlag(true)
	39	+ .NeedCheckSupportFlag(false)
	40	+ .PrecisionReduceFlag(true)
	41	+ .ExtendCfgInfo("opFile.value", "xlog1py_apt");
	42	+ this->AICore().AddConfig("ascend950", aicoreConfig);
	43	+ }
	44	+};
	45	+
	46	+OP_ADD(Xlog1py);
	47	+} // namespace ops

Amath/xlog1py/op_host/xlog1py_infershape.cpp+25-0

@@ -0,0 +1,25 @@
	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	+// Xlog1py InferShape — output = broadcast_max(x, y)
	13	+#include "infershape_broadcast_util.h"
	14	+#include "log/log.h"
	15	+#include "register/op_impl_registry.h"
	16	+
	17	+using namespace ge;
	18	+namespace ops {
	19	+static ge::graphStatus InferShapeForXlog1py(gert::InferShapeContext* context)
	20	+{
	21	+ return Ops::Base::InferShape4Broadcast(context);
	22	+}
	23	+
	24	+IMPL_OP_INFERSHAPE(Xlog1py).InferShape(InferShapeForXlog1py);
	25	+} // namespace ops

Amath/xlog1py/op_kernel/arch35/xlog1py_kernel.h+330-0

@@ -0,0 +1,330 @@
	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	+// Xlog1py Kernel — 2入1出 broadcast, VF 全链融合 (Adds→Ln→Mul→Compare→Select→Compare→Select)
	13	+// 链长=7, P_FP32=3 P_FP16=4
	14	+#pragma once
	15	+#include "kernel_operator.h"
	16	+#include "xlog1py_tiling_struct.h"
	17	+#include "xlog1py_struct.h"
	18	+
	19	+__aicore__ inline void GetCoreRange(int64_t core_id, int64_t tiles_main, int64_t cores_tail,
	20	+ int64_t& start, int64_t& end)
	21	+{
	22	+ if (core_id < cores_tail) {
	23	+ start = core_id * (tiles_main + 1);
	24	+ end = start + tiles_main + 1;
	25	+ } else {
	26	+ start = cores_tail * (tiles_main + 1) + (core_id - cores_tail) * tiles_main;
	27	+ end = start + tiles_main;
	28	+ }
	29	+}
	30	+
	31	+__aicore__ inline int64_t GetUBSplitRange(
	32	+ int64_t a_o_off, int64_t a_o, int64_t a_i, int64_t a_i_tail)
	33	+{
	34	+ return (a_o_off == a_o - 1) ? a_i_tail : a_i;
	35	+}
	36	+
	37	+__aicore__ inline bool FlatToEffectiveCoord(int64_t flat, const int64_t* max_bro_shape,
	38	+ int64_t rank, int64_t split_axis, int64_t a_i, int64_t a_o, int64_t* eff_coord)
	39	+{
	40	+ for (int64_t d = 0; d < rank; d++) eff_coord[d] = 0;
	41	+ int64_t a_o_off = flat % a_o;
	42	+ int64_t outer = flat / a_o;
	43	+ for (int64_t d = split_axis - 1; d >= 0; d--) {
	44	+ eff_coord[d] = outer % max_bro_shape[d];
	45	+ outer /= max_bro_shape[d];
	46	+ }
	47	+ eff_coord[split_axis] = a_o_off * a_i;
	48	+ return true;
	49	+}
	50	+
	51	+__aicore__ inline int64_t CalcInputOffset(const int64_t* eff_coord, const int64_t* strides, int64_t rank)
	52	+{
	53	+ int64_t offset = 0;
	54	+ for (int64_t d = 0; d < rank; d++) offset += eff_coord[d] * strides[d];
	55	+ return offset;
	56	+}
	57	+
	58	+__aicore__ inline int64_t CalcOutputOffset(const int64_t* eff_coord, const int64_t* strides, int64_t rank)
	59	+{
	60	+ int64_t offset = 0;
	61	+ for (int64_t d = 0; d < rank; d++) offset += eff_coord[d] * strides[d];
	62	+ return offset;
	63	+}
	64	+
	65	+__aicore__ inline int64_t CalcInputTransferCount(
	66	+ const int64_t* normal_shape, int64_t rank, int64_t split_axis, int64_t a_i_seg)
	67	+{
	68	+ int64_t split_elems = (normal_shape[split_axis] == 1) ? 1 : a_i_seg;
	69	+ int64_t inner_elems = 1;
	70	+ for (int64_t d = split_axis + 1; d < rank; d++) inner_elems *= normal_shape[d];
	71	+ return split_elems * inner_elems;
	72	+}
	73	+
	74	+__aicore__ inline int64_t CalcOutputTransferCount(
	75	+ const int64_t* normal_shape, int64_t rank, int64_t split_axis, int64_t a_i_seg)
	76	+{
	77	+ int64_t split_elems = (normal_shape[split_axis] == 1) ? 1 : a_i_seg;
	78	+ int64_t inner_elems = 1;
	79	+ for (int64_t d = split_axis + 1; d < rank; d++) inner_elems *= normal_shape[d];
	80	+ return split_elems * inner_elems;
	81	+}
	82	+
	83	+template <typename T>
	84	+__simd_vf__ inline void Xlog1pyVF(
	85	+ __ubuf__ T* dstAddr, __ubuf__ T* srcXAddr, __ubuf__ T* srcYAddr,
	86	+ uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes,
	87	+ float scalar_one);
	88	+
	89	+// ============================================================
	90	+// Xlog1pyKernel
	91	+// ============================================================
	92	+
	93	+template <typename T, int64_t RANK>
	94	+class Xlog1pyKernel {
	95	+ static constexpr int64_t ND = (RANK <= 5) ? RANK : 5;
	96	+ static constexpr bool NEED_CAST = !std::is_same_v<T, float>;
	97	+ static constexpr int64_t NUM_BUF = NEED_CAST ? 4 : 3;
	98	+ static constexpr uint32_t VL = AscendC::GetVecLen() / sizeof(float);
	99	+
	100	+ AscendC::TPipe pipe_;
	101	+ const Xlog1pyTilingData<RANK>* td_;
	102	+ AscendC::GlobalTensor<T> gmIn_[kMaxInputSlots];
	103	+ AscendC::GlobalTensor<T> gmOut_[kMaxOutputSlots];
	104	+ AscendC::TBuf<AscendC::TPosition::VECCALC> buf_[NUM_BUF];
	105	+ AscendC::MultiCopyParams<T, ND> nddmaParams_[kMaxInputSlots];
	106	+ int64_t nddmaOuterIters_[kMaxInputSlots];
	107	+ int64_t nddma_dims_;
	108	+
	109	+public:
	110	+ __aicore__ inline void Init(GM_ADDR* inputs, GM_ADDR* outputs,
	111	+ const Xlog1pyTilingData<RANK>* td)
	112	+ {
	113	+ td_ = td;
	114	+ for (int i = 0; i < kMaxInputSlots; i++)
	115	+ gmIn_[i].SetGlobalBuffer((__gm__ T*)inputs[i]);
	116	+ for (int i = 0; i < kMaxOutputSlots; i++)
	117	+ gmOut_[i].SetGlobalBuffer((__gm__ T*)outputs[i]);
	118	+ for (int i = 0; i < NUM_BUF; i++)
	119	+ pipe_.InitBuffer(buf_[i], td_->per_buf_bytes);
	120	+
	121	+ const int64_t* dstShape = td_->max_bro_shape;
	122	+ int64_t k = td_->split.axis;
	123	+ nddma_dims_ = (RANK - k <= ND) ? (RANK - k) : ND;
	124	+ for (int inp = 0; inp < kMaxInputSlots; inp++) {
	125	+ int64_t inner = 1;
	126	+ int64_t nd = 0;
	127	+ for (int64_t d = RANK - 1; d >= k && nd < ND; d--) {
	128	+ nddmaParams_[inp].loopInfo.loopSize[nd] = (d == k) ? 0 : dstShape[d];
	129	+ nddmaParams_[inp].loopInfo.loopSrcStride[nd] = td_->input_strides[inp][d];
	130	+ nddmaParams_[inp].loopInfo.loopDstStride[nd] = inner;
	131	+ nddmaParams_[inp].loopInfo.loopLpSize[nd] = 0;
	132	+ nddmaParams_[inp].loopInfo.loopRpSize[nd] = 0;
	133	+ inner *= (d == k) ? td_->split.a_i : dstShape[d];
	134	+ nd++;
	135	+ }
	136	+ for (; nd < ND; nd++) {
	137	+ nddmaParams_[inp].loopInfo.loopSize[nd] = 1;
	138	+ nddmaParams_[inp].loopInfo.loopSrcStride[nd] = 0;
	139	+ nddmaParams_[inp].loopInfo.loopDstStride[nd] = inner;
	140	+ nddmaParams_[inp].loopInfo.loopLpSize[nd] = 0;
	141	+ nddmaParams_[inp].loopInfo.loopRpSize[nd] = 0;
	142	+ }
	143	+ nddmaOuterIters_[inp] = 1;
	144	+ for (int64_t d = k; d < RANK - nddma_dims_; d++)
	145	+ nddmaOuterIters_[inp] *= (d == k) ? td_->split.a_i : dstShape[d];
	146	+ }
	147	+ }
	148	+
	149	+ __aicore__ inline void Process()
	150	+ {
	151	+ int32_t evMTE2toV = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_V));
	152	+ int32_t evVtoMTE2 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_MTE2));
	153	+ int32_t evVtoMTE3 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_MTE3));
	154	+ int32_t evMTE3toMTE2 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE3_MTE2));
	155	+
	156	+ int64_t start, end;
	157	+ GetCoreRange(AscendC::GetBlockIdx(), td_->multicore.tiles_main,
	158	+ td_->multicore.cores_tail, start, end);
	159	+
	160	+ int64_t inner_count = 1;
	161	+ for (int64_t d = td_->split.axis + 1; d < RANK; d++)
	162	+ inner_count *= td_->max_bro_shape[d];
	163	+
	164	+ int64_t coord[8] = {};
	165	+ for (int64_t flat = start; flat < end; flat++) {
	166	+ int64_t a_i_seg = GetUBSplitRange(flat % td_->split.a_o, td_->split.a_o,
	167	+ td_->split.a_i, td_->split.a_i_tail);
	168	+ int64_t count = a_i_seg * inner_count;
	169	+ FlatToEffectiveCoord(flat, td_->max_bro_shape, RANK,
	170	+ td_->split.axis, td_->split.a_i, td_->split.a_o, coord);
	171	+
	172	+ if (flat != start) AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
	173	+
	174	+ constexpr int IN_X = 0, IN_Y = 1;
	175	+ constexpr int OUT_Z = 0;
	176	+
	177	+ if constexpr (NEED_CAST) {
	178	+ ProcessFP16(coord, IN_X, IN_Y, OUT_Z, a_i_seg, count,
	179	+ evMTE2toV, evVtoMTE2, evVtoMTE3);
	180	+ } else {
	181	+ ProcessFP32(coord, IN_X, IN_Y, OUT_Z, a_i_seg, count,
	182	+ evMTE2toV, evVtoMTE2, evVtoMTE3);
	183	+ }
	184	+
	185	+ if (flat != end - 1)
	186	+ AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
	187	+ }
	188	+ }
	189	+
	190	+private:
	191	+ __aicore__ inline void ProcessFP32(
	192	+ const int64_t* coord, int inX, int inY, int outZ,
	193	+ int64_t a_i_seg, int64_t count,
	194	+ int32_t evMTE2toV, int32_t evVtoMTE2, int32_t evVtoMTE3)
	195	+ {
	196	+ constexpr int B_X = 0, B_Y = 1, B_Z = 2;
	197	+ CopyInOne(coord, inX, B_X, a_i_seg);
	198	+ CopyInOne(coord, inY, B_Y, a_i_seg);
	199	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
	200	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
	201	+ CallXlog1pyVF(B_Z, B_X, B_Y, count);
	202	+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
	203	+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
	204	+ CopyOutOne(coord, outZ, B_Z, a_i_seg);
	205	+ }
	206	+
	207	+ __aicore__ inline void ProcessFP16(
	208	+ const int64_t* coord, int inX, int inY, int outZ,
	209	+ int64_t a_i_seg, int64_t count,
	210	+ int32_t evMTE2toV, int32_t evVtoMTE2, int32_t evVtoMTE3)
	211	+ {
	212	+ constexpr int B_TEMP_X = 0, B_X = 1, B_Y = 2, B_TEMP_Y = 3;
	213	+ CopyInOne(coord, inX, B_TEMP_X, a_i_seg);
	214	+ CopyInOne(coord, inY, B_TEMP_Y, a_i_seg);
	215	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
	216	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
	217	+ AscendC::Cast(buf_[B_X].template Get<float>(), buf_[B_TEMP_X].template Get<T>(),
	218	+ AscendC::RoundMode::CAST_NONE, count);
	219	+ AscendC::Cast(buf_[B_Y].template Get<float>(), buf_[B_TEMP_Y].template Get<T>(),
	220	+ AscendC::RoundMode::CAST_NONE, count);
	221	+ CallXlog1pyVF(B_X, B_X, B_Y, count);
	222	+ AscendC::Cast(buf_[B_TEMP_X].template Get<T>(), buf_[B_X].template Get<float>(),
	223	+ AscendC::RoundMode::CAST_RINT, count);
	224	+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
	225	+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
	226	+ CopyOutOne(coord, outZ, B_TEMP_X, a_i_seg);
	227	+ }
	228	+
	229	+ __aicore__ inline void CallXlog1pyVF(int dstSlot, int srcXSlot, int srcYSlot, int64_t count)
	230	+ {
	231	+ uint16_t repeatTimes = AscendC::CeilDivision(count, VL);
	232	+ __ubuf__ float* dstAddr = (__ubuf__ float*)buf_[dstSlot].template Get<float>().GetPhyAddr();
	233	+ __ubuf__ float* srcXAddr = (__ubuf__ float*)buf_[srcXSlot].template Get<float>().GetPhyAddr();
	234	+ __ubuf__ float* srcYAddr = (__ubuf__ float*)buf_[srcYSlot].template Get<float>().GetPhyAddr();
	235	+ asc_vf_call<Xlog1pyVF<float>>(dstAddr, srcXAddr, srcYAddr,
	236	+ (uint32_t)count, VL, repeatTimes, 1.0f);
	237	+ }
	238	+
	239	+ __aicore__ inline void CopyInOne(const int64_t* coord, int inputIdx, int slot, int64_t a_i_seg)
	240	+ {
	241	+ int64_t k = td_->split.axis;
	242	+ int64_t off = CalcInputOffset(coord, td_->input_strides[inputIdx], RANK);
	243	+ const int64_t* dstShape = td_->max_bro_shape;
	244	+ auto params = nddmaParams_[inputIdx];
	245	+ int64_t k_nd = RANK - 1 - k;
	246	+ int64_t inner = 1;
	247	+ for (int64_t nd = 0; nd < ND; nd++) {
	248	+ if (nd == k_nd) params.loopInfo.loopSize[nd] = a_i_seg;
	249	+ params.loopInfo.loopDstStride[nd] = inner;
	250	+ inner *= params.loopInfo.loopSize[nd];
	251	+ }
	252	+ static constexpr AscendC::NdDmaConfig cfg = { false, AscendC::NdDmaConfig::unsetPad,
	253	+ AscendC::NdDmaConfig::unsetPad, false };
	254	+ if constexpr (RANK <= 5) {
	255	+ AscendC::DataCopy<T, ND, cfg>(buf_[slot].template Get<T>(), gmIn_[inputIdx][off], params);
	256	+ } else {
	257	+ AscendC::LocalTensor<T> buf = buf_[slot].template Get<T>();
	258	+ int64_t elem_base = off;
	259	+ for (int64_t oi = 0; oi < nddmaOuterIters_[inputIdx]; oi++) {
	260	+ int64_t elem_adj = 0, tmp = oi;
	261	+ for (int64_t d = RANK - nddma_dims_ - 1; d >= k; d--) {
	262	+ int64_t sz = (d == k) ? a_i_seg : dstShape[d];
	263	+ elem_adj += (tmp % sz) * td_->input_strides[inputIdx][d];
	264	+ tmp /= sz;
	265	+ }
	266	+ AscendC::DataCopy<T, ND, cfg>(buf[oi * inner], gmIn_[inputIdx][elem_base + elem_adj], params);
	267	+ }
	268	+ }
	269	+ }
	270	+
	271	+ __aicore__ inline void CopyOutOne(const int64_t* coord, int outputIdx, int slot, int64_t a_i_seg)
	272	+ {
	273	+ int64_t off = CalcOutputOffset(coord, td_->output_strides[outputIdx], RANK);
	274	+ int64_t cnt = CalcOutputTransferCount(td_->output_shapes[outputIdx], RANK,
	275	+ td_->split.axis, a_i_seg);
	276	+ AscendC::DataCopyExtParams extParams;
	277	+ extParams.blockCount = 1;
	278	+ extParams.blockLen = cnt * sizeof(T);
	S S sunday5月30日 CopyOutOne 的 blockLen 在 FP16/BF16 场景下可能不满足 DMA 32 字节对齐要求。FP16 主 tile 的 blockLen = per_buf_bytes/2，而 per_buf_bytes 仅保证 32 对齐，除以 2 后可能只 16 对齐（如 UB=192KB 时 49120/2=24560，24560%32=16）。尾块更可能远小于 32 字节。建议对 FP16/BF16 场景将 blockLen 做 CeilAlign(cnt*sizeof(T), 32)，或使用带 padding 的 DataCopyPadExtParams 处理尾块。
	279	+ extParams.srcStride = 0;
	280	+ extParams.dstStride = 0;
	281	+ AscendC::DataCopyPad(gmOut_[outputIdx][off], buf_[slot].template Get<T>(), extParams);
	282	+ }
	283	+};
	284	+
	285	+// ============================================================
	286	+// Xlog1pyVF — Adds→Ln→Mul→Compare→Select→Compare→Select (链长=7)
	287	+// ============================================================
	288	+
	289	+template <typename T>
	290	+__simd_vf__ inline void Xlog1pyVF(
	291	+ __ubuf__ T* dstAddr, __ubuf__ T* srcXAddr, __ubuf__ T* srcYAddr,
	292	+ uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes,
	293	+ float scalar_one)
	294	+{
	295	+ AscendC::Reg::RegTensor<T> regX, regY, regYp1, regLogY, regMul, regRes0, regFinal, regZero;
	296	+ AscendC::Reg::MaskReg mask, maskEq, maskNaN;
	297	+ AscendC::Reg::AddrReg aReg;
	298	+
	299	+ for (uint16_t i = 0; i < repeatTimes; ++i) {
	300	+ aReg = AscendC::Reg::CreateAddrReg<T>(i, oneRepeatSize);
	301	+ mask = AscendC::Reg::UpdateMask<T>(count);
	302	+
	303	+ AscendC::Reg::LoadAlign(regX, srcXAddr, aReg);
	304	+ AscendC::Reg::LoadAlign(regY, srcYAddr, aReg);
	305	+
	306	+ // S1: Adds(y, 1.0) → y+1
	307	+ AscendC::Reg::Adds(regYp1, regY, scalar_one, mask);
	308	+
	309	+ // S2: Ln(y+1) → log(1+y)
	310	+ AscendC::Reg::Ln(regLogY, regYp1, mask);
	311	+
	312	+ // S3: Mul(x, log_y) → mul_res
	313	+ AscendC::Reg::Mul(regMul, regX, regLogY, mask);
	314	+
	315	+ // S4: Compare(x, 0, EQ) → mask_eq
	316	+ AscendC::Reg::Sub(regZero, regX, regX, mask);
	317	+ AscendC::Reg::Compare<float, AscendC::CMPMODE::EQ>(maskEq, regX, regZero, mask);
	318	+
	319	+ // S5: Select(mask_eq, 0, mul_res) → res0
	320	+ AscendC::Reg::Select<float>(regRes0, regZero, regMul, maskEq);
	321	+
	322	+ // S6: Compare(y, y, EQ) → mask_nan (NaN≠NaN → mask=0)
	323	+ AscendC::Reg::Compare<float, AscendC::CMPMODE::EQ>(maskNaN, regY, regY, mask);
	324	+
	325	+ // S7: Select(mask_nan, res0, y) → final
	S S sunday5月30日 VF链的NaN传播逻辑有问题：当 x==0 且 y==NaN 时，S5 算出 regRes0=0，S6 得到 maskNaN=0（NaN自比较不等），S7 的 Select(mask=0) 选择 regY=NaN，最终输出 NaN。但 PyTorch torch.xlog1py(0, NaN)=0，x==0 的零化规则应优先于 NaN 传播。建议将 S6-S7 改为：先用 maskEq 和 maskNaN 做 AND 得到 maskSafe（既非零又非NaN），或者让 x==0 的规则优先：当 maskEq=1 时直接选 0，跳过 NaN 传播判断。
	326	+ AscendC::Reg::Select<float>(regFinal, regRes0, regY, maskNaN);
	327	+
	328	+ AscendC::Reg::StoreAlign(dstAddr, regFinal, aReg, mask);
	329	+ }
	330	+}

Amath/xlog1py/op_kernel/arch35/xlog1py_struct.h+31-0

@@ -0,0 +1,31 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+// Xlog1py TilingKey
	13	+#ifndef XLOG1PY_STRUCT_H_
	14	+#define XLOG1PY_STRUCT_H_
	15	+
	16	+#include "ascendc/host_api/tiling/template_argument.h"
	17	+
	18	+#define XLOG1PY_RANK_4 4
	19	+#define XLOG1PY_RANK_8 8
	20	+
	21	+ASCENDC_TPL_ARGS_DECL(Xlog1py,
	22	+ ASCENDC_TPL_UINT_DECL(RANK, 8, ASCENDC_TPL_UI_LIST,
	23	+ XLOG1PY_RANK_4, XLOG1PY_RANK_8)
	24	+);
	25	+
	26	+ASCENDC_TPL_SEL(
	27	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, XLOG1PY_RANK_4)),
	28	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, XLOG1PY_RANK_8))
	29	+);
	30	+
	31	+#endif

Amath/xlog1py/op_kernel/arch35/xlog1py_tiling_struct.h+46-0

@@ -0,0 +1,46 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+// Xlog1py TilingData
	13	+#pragma once
	14	+#include <cstdint>
	15	+
	16	+constexpr int64_t kMaxInputSlots = 2;
	17	+constexpr int64_t kMaxOutputSlots = 1;
	18	+
	19	+struct SplitResult {
	20	+ int64_t axis;
	21	+ int64_t a_i;
	22	+ int64_t a_o;
	23	+ int64_t a_i_tail;
	24	+};
	25	+
	26	+struct MultiCoreResult {
	27	+ int64_t num_cores;
	28	+ int64_t total_tiles;
	29	+ int64_t tiles_main;
	30	+ int64_t cores_tail;
	31	+};
	32	+
	33	+template<int64_t kRank>
	34	+struct Xlog1pyTilingData {
	35	+ SplitResult split;
	36	+ MultiCoreResult multicore;
	37	+ int64_t rank;
	38	+ int64_t per_buf_bytes;
	39	+ int64_t max_bro_shape[kRank];
	40	+ int64_t num_inputs;
	41	+ int64_t num_outputs;
	42	+ int64_t input_shapes [kMaxInputSlots][kRank];
	43	+ int64_t input_strides[kMaxInputSlots][kRank];
	44	+ int64_t output_shapes[kMaxOutputSlots][kRank];
	45	+ int64_t output_strides[kMaxOutputSlots][kRank];
	46	+};

Amath/xlog1py/op_kernel/xlog1py_apt.cpp+41-0

@@ -0,0 +1,41 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+// Xlog1py Kernel 入口
	13	+#include "kernel_operator.h"
	14	+#include "arch35/xlog1py_kernel.h"
	15	+#include "arch35/xlog1py_tiling_struct.h"
	16	+
	17	+using TilingData4 = Xlog1pyTilingData<4>;
	18	+using TilingData8 = Xlog1pyTilingData<8>;
	19	+
	20	+template<int RANK>
	21	+__global__ __aicore__ void xlog1py(
	22	+ GM_ADDR x, GM_ADDR y,
	23	+ GM_ADDR z,
	24	+ GM_ADDR workspace, GM_ADDR tiling)
	25	+{
	26	+ GM_ADDR ins[2] = {x, y};
	27	+ GM_ADDR outs[1] = {z};
	28	+ REGISTER_NONE_TILING;
	29	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	30	+ if constexpr (RANK == 4) {
	31	+ GET_TILING_DATA_WITH_STRUCT(TilingData4, td, tiling);
	32	+ Xlog1pyKernel<DTYPE_X, 4> kernel;
	33	+ kernel.Init(ins, outs, &td);
	34	+ kernel.Process();
	35	+ } else {
	36	+ GET_TILING_DATA_WITH_STRUCT(TilingData8, td, tiling);
	37	+ Xlog1pyKernel<DTYPE_X, 8> kernel;
	38	+ kernel.Init(ins, outs, &td);
	39	+ kernel.Process();
	40	+ }
	41	+}