feat(hypot): 新增Hypot算子Ascend C实现（A2/A3）-ops-math-AtomGit

Mdocs/zh/menu_aclnn_api.md+1-0

	@@ -122,6 +122,7 @@
122	- [aclnnHansEncode](../../math/hans_encode/docs/aclnnHansEncode.md)	122	- [aclnnHansEncode](../../math/hans_encode/docs/aclnnHansEncode.md)
123	- [aclnnHardtanh&aclnnInplaceHardtanh](../../conversion/clip_by_value_v2/docs/aclnnHardtanh&aclnnInplaceHardtanh.md)	123	- [aclnnHardtanh&aclnnInplaceHardtanh](../../conversion/clip_by_value_v2/docs/aclnnHardtanh&aclnnInplaceHardtanh.md)
124	- [aclnnHistc](../../math/histogram_v2/docs/aclnnHistc.md)	124	- [aclnnHistc](../../math/histogram_v2/docs/aclnnHistc.md)
		125	+- [aclnnHypot](../../math/hypot/docs/aclnnHypot.md)
125	- [aclnnIm2col](../../conversion/im2col/docs/aclnnIm2col.md)	126	- [aclnnIm2col](../../conversion/im2col/docs/aclnnIm2col.md)
126	- [aclnnInplaceBitwiseAndScalar](../../math/bitwise_and/docs/aclnnInplaceBitwiseAndScalar.md)	127	- [aclnnInplaceBitwiseAndScalar](../../math/bitwise_and/docs/aclnnInplaceBitwiseAndScalar.md)
127	- [aclnnInplaceBitwiseAndTensor](../../math/bitwise_and/docs/aclnnInplaceBitwiseAndTensor.md)	128	- [aclnnInplaceBitwiseAndTensor](../../math/bitwise_and/docs/aclnnInplaceBitwiseAndTensor.md)

Mdocs/zh/op_api_list.md+1-0

	@@ -149,6 +149,7 @@
149	\| [aclnnHansEncode](../../math/hans_encode/docs/aclnnHansEncode.md) \| 对张量的指数位所在字节实现PDF统计，按PDF分布统计进行无损压缩。 \| 默认确定性实现\| 默认确定性实现 \|	149	\| [aclnnHansEncode](../../math/hans_encode/docs/aclnnHansEncode.md) \| 对张量的指数位所在字节实现PDF统计，按PDF分布统计进行无损压缩。 \| 默认确定性实现\| 默认确定性实现 \|
150	\| [aclnnHardtanh&aclnnInplaceHardtanh](../../conversion/clip_by_value_v2/docs/aclnnHardtanh&aclnnInplaceHardtanh.md) \| 将输入的所有元素限制在[clipValueMin,clipValueMax]范围内，若元素大于clipValueMax则限制为clipValueMax，若元素小于clipValueMin则限制为clipValueMin，否则等于元素本身。 \| 默认确定性实现\| 默认确定性实现 \|	150	\| [aclnnHardtanh&aclnnInplaceHardtanh](../../conversion/clip_by_value_v2/docs/aclnnHardtanh&aclnnInplaceHardtanh.md) \| 将输入的所有元素限制在[clipValueMin,clipValueMax]范围内，若元素大于clipValueMax则限制为clipValueMax，若元素小于clipValueMin则限制为clipValueMin，否则等于元素本身。 \| 默认确定性实现\| 默认确定性实现 \|
151	\| [aclnnHistc](../../math/histogram_v2/docs/aclnnHistc.md) \| 以min和max作为统计上下限，在min和max之间划出等宽的数量为bins的区间，统计张量self中元素在各个区间的数量。如果min和max相等，则使用张量中所有元素的最小值和最大值作为统计的上下限。小于min和大于max的元素不会被统计。 \| 默认确定性实现\| 默认非确定性实现，支持配置开启 \|	151	\| [aclnnHistc](../../math/histogram_v2/docs/aclnnHistc.md) \| 以min和max作为统计上下限，在min和max之间划出等宽的数量为bins的区间，统计张量self中元素在各个区间的数量。如果min和max相等，则使用张量中所有元素的最小值和最大值作为统计的上下限。小于min和大于max的元素不会被统计。 \| 默认确定性实现\| 默认非确定性实现，支持配置开启 \|
		152	+\| [aclnnHypot](../../math/hypot/docs/aclnnHypot.md) \| 逐元素计算两个输入张量对应元素构成的直角三角形斜边长度。 \| 默认确定性实现\| 默认确定性实现\|
		C C caiwenwen5 days ago 补充mune_api
152	\| [aclnnIm2col](../../conversion/im2col/docs/aclnnIm2col.md) \| 图像到列，滑动局部窗口数据转为列向量，拼接为大张量。从批处理输入张量中提取滑动窗口。 \| 默认确定性实现\| 默认确定性实现\|	153	\| [aclnnIm2col](../../conversion/im2col/docs/aclnnIm2col.md) \| 图像到列，滑动局部窗口数据转为列向量，拼接为大张量。从批处理输入张量中提取滑动窗口。 \| 默认确定性实现\| 默认确定性实现\|
153	\| [aclnnInplaceBitwiseAndScalar](../../math/bitwise_and/docs/aclnnInplaceBitwiseAndScalar.md) \| 计算输入tensor中每个元素和输入标量other的按位与结果。输入的tensor与other必须是整型或者BOOL型变量。当输入为BOOL型时，计算逻辑与的结果。 \| 默认确定性实现\| 默认确定性实现 \|	154	\| [aclnnInplaceBitwiseAndScalar](../../math/bitwise_and/docs/aclnnInplaceBitwiseAndScalar.md) \| 计算输入tensor中每个元素和输入标量other的按位与结果。输入的tensor与other必须是整型或者BOOL型变量。当输入为BOOL型时，计算逻辑与的结果。 \| 默认确定性实现\| 默认确定性实现 \|
154	\| [aclnnInplaceBitwiseAndTensor](../../math/bitwise_and/docs/aclnnInplaceBitwiseAndTensor.md) \| 输入为BOOL型tensor时，进行逻辑与运算；输入为INT型时，进行位与运算。 \| 默认确定性实现\| 默认确定性实现 \|	155	\| [aclnnInplaceBitwiseAndTensor](../../math/bitwise_and/docs/aclnnInplaceBitwiseAndTensor.md) \| 输入为BOOL型tensor时，进行逻辑与运算；输入为INT型时，进行位与运算。 \| 默认确定性实现\| 默认确定性实现 \|

Mdocs/zh/op_list.md+10-0

	@@ -1056,6 +1056,16 @@
1056	<td>AI Core</td>	1056	<td>AI Core</td>
1057	<td>计算张量直方图。</td>	1057	<td>计算张量直方图。</td>
1058	</tr>	1058	</tr>
		1059	+ <tr>
		1060	+ <td>math</td>
		1061	+ <td><a href="../../math/hypot/README.md">hypot</a></td>
		1062	+ <td>√</td>
		1063	+ <td>√</td>
		1064	+ <td>√</td>
		1065	+ <td>√</td>
		1066	+ <td>AI Core</td>
		1067	+ <td>逐元素计算两个输入张量对应元素构成的直角三角形斜边长度。</td>
		1068	+ </tr>
1059	<tr>	1069	<tr>
1060	<td>math</td>	1070	<td>math</td>
1061	<td><a href="../../math/inplace_top_k_distance/README.md">inplace_top_k_distance</a></td>	1071	<td><a href="../../math/inplace_top_k_distance/README.md">inplace_top_k_distance</a></td>

Amath/hypot/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 "ascend910_93" "ascend910b")
	13	+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
	14	+set(SUPPORT_TILING_DIR "arch22" "arch22")
	15	+add_all_modules_sources(OPTYPE hypot ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR})

Amath/hypot/README.md+78-0

@@ -0,0 +1,78 @@
	1	+# Hypot
	2	+
	3	+## 产品支持情况
	4	+
	5	+\| 产品 \| 是否支持 \|
	6	+\| :----------------------------------------------------------- \| :------: \|
	7	+\| <term>Ascend 950PR&950DT系列产品</term> \| × \|
	8	+\| <term>Atlas A3系列产品</term> \| √ \|
	9	+\| <term>Atlas A2系列产品</term> \| √ \|
	10	+\| <term>Atlas 200I/500 A2推理产品</term> \| × \|
	11	+\| <term>Atlas推理系列产品</term> \| × \|
	12	+\| <term>Atlas训练系列产品</term> \| × \|
	13	+
	14	+## 功能说明
	15	+
	16	+- 算子功能：逐元素计算两个输入张量对应元素构成的直角三角形斜边长度（二维欧几里得范数），输出恒为非负值。
	17	+- 计算公式：
	18	+ $$
	19	+ out_i = \sqrt{self_i^{2} + other_i^{2}}
	20	+ $$
	21	+- 数值稳定性：采用以绝对值较大一方为基准缩放的实现，避免中间平方项上溢或下溢；双零输入输出 `+0`，`inf` 优先于 `NaN`。
	22	+
	23	+## 参数说明
	24	+
	25	+<table style="table-layout: fixed; width: 820px"><colgroup>
	26	+ <col style="width: 100px">
	27	+ <col style="width: 150px">
	28	+ <col style="width: 190px">
	29	+ <col style="width: 260px">
	30	+ <col style="width: 120px">
	31	+ </colgroup>
	32	+ <thead>
	33	+ <tr>
	34	+ <th>参数名</th>
	35	+ <th>输入/输出/属性</th>
	36	+ <th>描述</th>
	37	+ <th>数据类型</th>
	38	+ <th>数据格式</th>
	39	+ </tr></thead>
	40	+ <tbody>
	41	+ <tr>
	42	+ <td>self</td>
	43	+ <td>输入</td>
	44	+ <td>公式中的输入张量self。</td>
	45	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	46	+ <td>ND</td>
	47	+ </tr>
	48	+ <tr>
	49	+ <td>other</td>
	50	+ <td>输入</td>
	51	+ <td>公式中的输入张量other。</td>
	52	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	53	+ <td>ND</td>
	54	+ </tr>
	55	+ <tr>
	56	+ <td>out</td>
	57	+ <td>输出</td>
	58	+ <td>公式中的输出张量out。</td>
	59	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
	60	+ <td>ND</td>
	61	+ </tr>
	62	+ </tbody></table>
	63	+
	64	+- self 与 other 的 dtype 允许不一致，按互推导关系确定输出类型：任一侧为 FLOAT 或两侧 dtype 不同（FLOAT16 与 BFLOAT16）时输出 FLOAT，两侧同 dtype 时输出该 dtype。
	65	+
	66	+## 约束说明
	67	+
	68	+- 支持 0~8 维 ND Tensor，输入 shape 需满足 broadcast 关系。
	69	+- 支持空 Tensor 与非连续 Tensor。
	70	+- 不支持 DOUBLE、整型、BOOL 与复数类型。
	71	+- 本期仅提供 `aclnnHypot` 与图模式 `Hypot`，不提供原地（in-place）接口。
	72	+
	73	+## 调用说明
	74	+
	75	+\| 调用方式 \| 调用样例 \| 说明 \|
	76	+\|--------------\|------------------------------------------------------------------------\|----------------------------------------------------------------\|
	77	+\| aclnn调用 \| [test_aclnn_hypot](./examples/test_aclnn_hypot.cpp) \| 通过[aclnnHypot](./docs/aclnnHypot.md)接口方式调用Hypot算子。 \|
	78	+\| 图模式调用 \| [test_geir_hypot](./examples/test_geir_hypot.cpp) \| 通过[算子IR](./op_graph/hypot_proto.h)构图方式调用Hypot算子。 \|

Amath/hypot/README_en.md+78-0

@@ -0,0 +1,78 @@
	1	+# Hypot
	2	+
	3	+## Applicable Products
	4	+
	5	+\| Product \| Supported \|
	6	+\| :----------------------------------------------------------- \| :------: \|
	7	+\| <term>Ascend 950PR&950DT products</term> \| × \|
	8	+\| <term>Atlas A3 products</term> \| √ \|
	9	+\| <term>Atlas A2 products</term> \| √ \|
	10	+\| <term>Atlas 200I/500 A2 inference products</term> \| × \|
	11	+\| <term>Atlas inference products</term> \| × \|
	12	+\| <term>Atlas training products</term> \| × \|
	13	+
	14	+## Description
	15	+
	16	+- Operator function: The operator computes the hypotenuse length (two-dimensional Euclidean norm) elementwise from the corresponding elements of two input tensors. The output is always nonnegative.
	17	+- Formula:
	18	+ $$
	19	+ out_i = \sqrt{self_i^{2} + other_i^{2}}
	20	+ $$
	21	+- Numerical stability: The implementation scales inputs by the larger absolute value to avoid overflow or underflow in intermediate squared values. Two zero inputs produce `+0`; `inf` takes precedence over `NaN`.
	22	+
	23	+## Parameters
	24	+
	25	+<table style="table-layout: fixed; width: 820px"><colgroup>
	26	+ <col style="width: 100px">
	27	+ <col style="width: 150px">
	28	+ <col style="width: 190px">
	29	+ <col style="width: 260px">
	30	+ <col style="width: 120px">
	31	+ </colgroup>
	32	+ <thead>
	33	+ <tr>
	34	+ <th>Parameter</th>
	35	+ <th>Input/Output/Attribute</th>
	36	+ <th>Description</th>
	37	+ <th>Data Type</th>
	38	+ <th>Data Format</th>
	39	+ </tr></thead>
	40	+ <tbody>
	41	+ <tr>
	42	+ <td>self</td>
	43	+ <td>Input</td>
	44	+ <td>Input tensor self in the formula.</td>
	45	+ <td>FLOAT, FLOAT16, and BFLOAT16</td>
	46	+ <td>ND</td>
	47	+ </tr>
	48	+ <tr>
	49	+ <td>other</td>
	50	+ <td>Input</td>
	51	+ <td>Input tensor other in the formula.</td>
	52	+ <td>FLOAT, FLOAT16, and BFLOAT16</td>
	53	+ <td>ND</td>
	54	+ </tr>
	55	+ <tr>
	56	+ <td>out</td>
	57	+ <td>Output</td>
	58	+ <td>Output tensor out in the formula.</td>
	59	+ <td>FLOAT, FLOAT16, and BFLOAT16</td>
	60	+ <td>ND</td>
	61	+ </tr>
	62	+ </tbody></table>
	63	+
	64	+- self and other can have different dtypes. Type promotion determines the output dtype. The output dtype is FLOAT if either input is FLOAT or if one input is FLOAT16 and the other is BFLOAT16. If both inputs have the same dtype, the output has that dtype.
	65	+
	66	+## Constraints
	67	+
	68	+- The operator supports ND tensors with 0 to 8 dimensions. The input shapes must support broadcasting.
	69	+- The operator supports empty tensors and non-contiguous tensors.
	70	+- The operator does not support DOUBLE, integer types, BOOL, or complex types.
	71	+- This release provides only `aclnnHypot` and the graph-mode `Hypot` operator. It does not provide an in-place API.
	72	+
	73	+## Invocation
	74	+
	75	+\| Invocation Method \| Sample \| Description \|
	76	+\|--------------\|------------------------------------------------------------------------\|----------------------------------------------------------------\|
	77	+\| aclnn invocation \| [test_aclnn_hypot](./examples/test_aclnn_hypot.cpp) \| Invoke the Hypot operator using the [aclnnHypot](./docs/aclnnHypot.md) API. \|
	78	+\| Graph-mode invocation \| [test_geir_hypot](./examples/test_geir_hypot.cpp) \| Invoke the Hypot operator by constructing a graph using the [operator IR](./op_graph/hypot_proto.h). \|

Amath/hypot/docs/aclnnHypot.md+552-0

@@ -0,0 +1,552 @@
	1	+# aclnnHypot
	2	+
	3	+## 产品支持情况
	4	+
	5	+<!-- npu="950" id1 -->
	6	+- <term>Ascend 950PR&950DT系列产品</term>：不支持
	7	+<!-- end id1 -->
	8	+<!-- npu="A3" id2 -->
	9	+- <term>Atlas A3系列产品</term>：支持
	10	+<!-- end id2 -->
	11	+<!-- npu="910b" id3 -->
	12	+- <term>Atlas A2系列产品</term>：支持
	13	+<!-- end id3 -->
	14	+<!-- npu="310b" id4 -->
	15	+- <term>Atlas 200I/500 A2推理产品</term>：不支持
	16	+<!-- end id4 -->
	17	+<!-- npu="310p" id5 -->
	18	+- <term>Atlas推理系列产品</term>：不支持
	19	+<!-- end id5 -->
	20	+<!-- npu="910" id6 -->
	21	+- <term>Atlas训练系列产品</term>：不支持
	22	+<!-- end id6 -->
	23	+
	24	+## 功能说明
	25	+
	26	+- 接口功能：逐元素计算两个输入张量所对应的直角三角形斜边长度（二维欧几里得范数）。
	27	+- 计算公式：
	28	+
	29	+ $$
	30	+ out_i = \sqrt{self_i^{2} + other_i^{2}}
	31	+ $$
	32	+
	33	+- 计算原理：算子结果为$\sqrt{self_i^{2} + other_i^{2}}$，且不得因中间结果$self_i^{2}$、$other_i^{2}$的上溢或下溢而失真。本算子采用以绝对值较大一方为基准的稳定变形计算：
	34	+
	35	+ $$
	36	+ out_i =
	37	+ \begin{cases}
	38	+ \|self_i\| \cdot \sqrt{1 + \left( \dfrac{other_i}{self_i} \right)^{2}}, & \|self_i\| \ge \|other_i\| \\
	39	+ \|other_i\| \cdot \sqrt{1 + \left( \dfrac{self_i}{other_i} \right)^{2}}, & \|self_i\| \lt \|other_i\|
	40	+ \end{cases}
	41	+ $$
	42	+
	43	+ 其中$self_i$、$other_i$为输入在第$i$个元素位置的取值，$out_i$为输出在第$i$个元素位置的取值，当$self_i = other_i = 0$时按$out_i = 0$处理。
	44	+
	45	+## 函数原型
	46	+
	47	+每个算子分为[两段式接口](../../../docs/zh/context/two_phase_api.md)，必须先调用“aclnnHypotGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnHypot”接口执行计算。
	48	+
	49	+```Cpp
	50	+aclnnStatus aclnnHypotGetWorkspaceSize(
	51	+ const aclTensor* self,
	52	+ const aclTensor* other,
	53	+ aclTensor* out,
	54	+ uint64_t* workspaceSize,
	55	+ aclOpExecutor** executor)
	56	+```
	57	+
	58	+```Cpp
	59	+aclnnStatus aclnnHypot(
	60	+ void* workspace,
	61	+ uint64_t workspaceSize,
	62	+ aclOpExecutor* executor,
	63	+ aclrtStream stream)
	64	+```
	65	+
	66	+## aclnnHypotGetWorkspaceSize
	67	+
	68	+- 参数说明：
	69	+
	70	+ <table style="table-layout: fixed; width: 1500px"><colgroup>
	71	+ <col style="width: 180px">
	72	+ <col style="width: 120px">
	73	+ <col style="width: 300px">
	74	+ <col style="width: 350px">
	75	+ <col style="width: 250px">
	76	+ <col style="width: 100px">
	77	+ <col style="width: 100px">
	78	+ <col style="width: 100px">
	79	+ </colgroup>
	80	+ <thead>
	81	+ <tr>
	82	+ <th>参数名</th>
	83	+ <th>输入/输出</th>
	84	+ <th>描述</th>
	85	+ <th>使用说明</th>
	86	+ <th>数据类型</th>
	87	+ <th>数据格式</th>
	88	+ <th>维度(shape)</th>
	89	+ <th>非连续Tensor</th>
	90	+ </tr></thead>
	91	+ <tbody>
	92	+ <tr>
	93	+ <td>self（aclTensor*）</td>
	94	+ <td>输入</td>
	95	+ <td>公式中的self。</td>
	96	+ <td>
	97	+ <ul>
	98	+ <li>支持空Tensor。</li>
	99	+ <li>数据类型与other的数据类型需满足数据类型推导规则（参见<a href="../../../docs/zh/context/deduction_relationship.md" target="_blank">互推导关系</a>）。</li>
	100	+ <li>shape需要与other满足<a href="../../../docs/zh/context/broadcast_relationship.md" target="_blank">broadcast关系</a>。</li>

Amath/hypot/examples/test_aclnn_hypot.cpp+274-0

@@ -0,0 +1,274 @@
	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	+ * @file test_aclnn_hypot.cpp
	11	+ * aclnnHypot two-stage call example (eager mode).
	12	+ *
	13	+ * Covered scenarios (all FLOAT32):
	14	+ * 1) basic same-shape elementwise: out = sqrt(self^2 + other^2)
	15	+ * 2) broadcast: self (3,1) x other (1,4) -> out (3,4)
	16	+ * 3) special values: inf / NaN / signed zero / 1e30 stability demonstration
	17	+ *
	18	+ * Usage: ./test_aclnn_hypot [deviceId] (default deviceId = 0)
	19	+ */
	20	+
	21	+#include <cmath>
	22	+#include <cstdint>
	23	+#include <cstdio>
	24	+#include <cstdlib>
	25	+#include <limits>
	26	+#include <string>
	27	+#include <vector>
	28	+
	29	+#include "acl/acl.h"
	30	+#include "aclnn_hypot.h"
	31	+
	32	+#define CHECK_RET(cond, return_expr) \
	33	+ do { \
	34	+ if (!(cond)) { \
	35	+ return_expr; \
	36	+ } \
	37	+ } while (0)
	38	+
	39	+#define LOG_PRINT(message, ...) \
	40	+ do { \
	41	+ printf(message, ##__VA_ARGS__); \
	42	+ } while (0)
	43	+
	44	+static int64_t GetShapeSize(const std::vector<int64_t>& shape)
	45	+{
	46	+ int64_t shapeSize = 1;
	47	+ for (auto i : shape) {
	48	+ shapeSize *= i;
	49	+ }
	50	+ return shapeSize;
	51	+}
	52	+
	53	+static int Init(int32_t deviceId, aclrtStream* stream)
	54	+{
	55	+ // Fixed ACL resource initialization sequence
	56	+ auto ret = aclInit(nullptr);
	57	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	58	+ ret = aclrtSetDevice(deviceId);
	59	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); aclFinalize(); return ret);
	60	+ ret = aclrtCreateStream(stream);
	61	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); aclrtResetDevice(deviceId);
	62	+ aclFinalize(); return ret);
	63	+ return 0;
	64	+}
	65	+
	66	+template <typename T>
	67	+static int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
	68	+ aclDataType dataType, aclTensor** tensor)
	69	+{
	70	+ auto size = GetShapeSize(shape) * sizeof(T);
	71	+ // Allocate device memory
	72	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	73	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	74	+ // Copy host data to device
	75	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	76	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	77	+
	78	+ // Strides of a contiguous tensor
	79	+ std::vector<int64_t> strides(shape.size(), 1);
	80	+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; i--) {
	81	+ strides[i] = shape[i + 1] * strides[i + 1];
	82	+ }
	83	+
	84	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	85	+ shape.data(), shape.size(), *deviceAddr);
	86	+ return 0;
	87	+}
	88	+
	89	+// Golden reference: double-precision std::hypot.
	90	+static double GoldenHypot(float a, float b) { return std::hypot(static_cast<double>(a), static_cast<double>(b)); }
	91	+
	92	+// Map a linear index in the broadcast output back to the linear index of one input,
	93	+// following numpy broadcast rules (input dims of size 1 contribute offset 0).
	94	+static int64_t BroadcastLinearIndex(const std::vector<int64_t>& inShape, const std::vector<int64_t>& outShape,
	95	+ int64_t outLinear)
	96	+{
	97	+ size_t outRank = outShape.size();
	98	+ size_t inRank = inShape.size();
	99	+ std::vector<int64_t> outStrides(outRank, 1);
	100	+ for (int64_t d = static_cast<int64_t>(outRank) - 2; d >= 0; --d) {
	101	+ outStrides[d] = outStrides[d + 1] * outShape[d + 1];
	102	+ }
	103	+ std::vector<int64_t> inStrides(inRank, 1);
	104	+ for (int64_t d = static_cast<int64_t>(inRank) - 2; d >= 0; --d) {
	105	+ inStrides[d] = inStrides[d + 1] * inShape[d + 1];
	106	+ }
	107	+ size_t pad = outRank - inRank; // input is left-padded with 1s
	108	+ int64_t inLinear = 0;
	109	+ for (size_t od = 0; od < outRank; ++od) {
	110	+ int64_t coord = (outStrides[od] > 0) ? (outLinear / outStrides[od]) % outShape[od] : 0;
	111	+ if (od >= pad) {
	112	+ size_t id = od - pad;
	113	+ if (inShape[id] != 1) {
	114	+ inLinear += coord * inStrides[id];
	115	+ }
	116	+ }
	117	+ }
	118	+ return inLinear;
	119	+}
	120	+
	121	+// Verify one element against golden; fp32 output tolerance: \|a-g\| <= atol + rtol*\|g\|.
	122	+static bool CheckValue(float actual, double golden, double rtol, double atol, std::string* reason)
	123	+{
	124	+ if (std::isnan(golden)) {
	125	+ if (std::isnan(actual)) {
	126	+ return true;
	127	+ }
	128	+ *reason = "expect NaN";
	129	+ return false;
	130	+ }
	131	+ if (std::isnan(actual)) {
	132	+ *reason = "got NaN, expect finite/inf";
	133	+ return false;
	134	+ }
	135	+ if (std::isinf(golden)) {
	136	+ if (std::isinf(actual) && (std::signbit(actual) == std::signbit(golden))) {
	137	+ return true;
	138	+ }
	139	+ *reason = "expect same-sign inf";
	140	+ return false;
	141	+ }
	142	+ if (std::isinf(actual)) {
	143	+ *reason = "got inf, expect finite";
	144	+ return false;
	145	+ }
	146	+ double err = std::fabs(static_cast<double>(actual) - golden);
	147	+ if (err <= atol + rtol * std::fabs(golden)) {
	148	+ return true;
	149	+ }
	150	+ *reason = "abs err " + std::to_string(err) + " exceeds tolerance";
	151	+ return false;
	152	+}
	153	+
	154	+// Run one aclnnHypot scenario end to end (two-stage call) and verify against golden.
	155	+// Broadcast semantics: out[i] pairs self/other by numpy broadcast over selfShape/otherShape.
	156	+static bool RunScenario(const char* name, const std::vector<int64_t>& selfShape, const std::vector<int64_t>& otherShape,
	157	+ const std::vector<int64_t>& outShape, const std::vector<float>& selfData,
	158	+ const std::vector<float>& otherData, aclrtStream stream)
	159	+{
	160	+ LOG_PRINT("------------------------------------------------------------\n");
	161	+ LOG_PRINT("[scenario] %s\n", name);
	162	+
	163	+ void* selfDeviceAddr = nullptr;
	164	+ void* otherDeviceAddr = nullptr;
	165	+ void* outDeviceAddr = nullptr;
	166	+ aclTensor* self = nullptr;
	167	+ aclTensor* other = nullptr;
	168	+ aclTensor* out = nullptr;
	169	+ std::vector<float> outData(GetShapeSize(outShape), 0.0F);
	170	+
	171	+ int ret = CreateAclTensor(selfData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
	172	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("create self tensor failed\n"); return false);
	173	+ ret = CreateAclTensor(otherData, otherShape, &otherDeviceAddr, aclDataType::ACL_FLOAT, &other);
	174	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("create other tensor failed\n"); return false);
	175	+ ret = CreateAclTensor(outData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
	176	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("create out tensor failed\n"); return false);
	177	+
	178	+ // Stage 1: query workspace size and obtain the executor
	179	+ uint64_t workspaceSize = 0;
	180	+ aclOpExecutor* executor = nullptr;
	181	+ ret = aclnnHypotGetWorkspaceSize(self, other, out, &workspaceSize, &executor);
	182	+ if (ret != ACL_SUCCESS) {
	183	+ LOG_PRINT("aclnnHypotGetWorkspaceSize failed. ERROR: %d\n", ret);
	184	+ return false;
	185	+ }
	186	+ void* workspaceAddr = nullptr;
	187	+ if (workspaceSize > 0) {
	188	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	189	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return false);
	190	+ }
	191	+ // Stage 2: launch the computation (executor is consumed by this call)
	192	+ ret = aclnnHypot(workspaceAddr, workspaceSize, executor, stream);
	193	+ if (ret != ACL_SUCCESS) {
	194	+ LOG_PRINT("aclnnHypot failed. ERROR: %d\n", ret);
	195	+ return false;
	196	+ }
	197	+ ret = aclrtSynchronizeStream(stream);
	198	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return false);
	199	+
	200	+ // Copy result back to host
	201	+ auto outSize = GetShapeSize(outShape);
	202	+ std::vector<float> resultData(outSize, 0.0F);
	203	+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(float), outDeviceAddr, outSize * sizeof(float),
	204	+ ACL_MEMCPY_DEVICE_TO_HOST);
	205	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return false);
	206	+
	207	+ // Broadcast-aware golden comparison
	208	+ const double rtol = 1e-4;
	209	+ const double atol = 1e-5;
	210	+ bool pass = true;
	211	+ for (int64_t i = 0; i < outSize; i++) {
	212	+ int64_t selfIdx = BroadcastLinearIndex(selfShape, outShape, i);
	213	+ int64_t otherIdx = BroadcastLinearIndex(otherShape, outShape, i);
	214	+ double golden = GoldenHypot(selfData[selfIdx], otherData[otherIdx]);
	215	+ std::string reason;
	216	+ bool ok = CheckValue(resultData[i], golden, rtol, atol, &reason);
	217	+ if (!ok) {
	218	+ pass = false;
	219	+ }
	220	+ LOG_PRINT(" out[%ld] = %.8g (golden %.8g) %s%s\n", i, resultData[i], golden,
	221	+ ok ? "OK" : "MISMATCH: ", ok ? "" : reason.c_str());
	222	+ }
	223	+
	224	+ aclDestroyTensor(self);
	225	+ aclDestroyTensor(other);
	226	+ aclDestroyTensor(out);
	227	+ aclrtFree(selfDeviceAddr);
	228	+ aclrtFree(otherDeviceAddr);
	229	+ aclrtFree(outDeviceAddr);
	230	+ if (workspaceSize > 0) {
	231	+ aclrtFree(workspaceAddr);
	232	+ }
	233	+ LOG_PRINT("[scenario] %s : %s\n", name, pass ? "PASS" : "FAIL");
	234	+ return pass;
	235	+}
	236	+
	237	+int main(int argc, char* argv[])
	238	+{
	239	+ int32_t deviceId = 0;
	240	+ if (argc > 1) {
	241	+ deviceId = std::atoi(argv[1]);
	242	+ }
	243	+ LOG_PRINT("=== aclnnHypot example (two-stage aclnn call, FLOAT32) ===\n");
	244	+ LOG_PRINT("deviceId: %d\n", deviceId);
	245	+
	246	+ aclrtStream stream = nullptr;
	247	+ auto ret = Init(deviceId, &stream);
	248	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	249	+
	250	+ bool allPass = true;
	251	+
	252	+ // Scenario 1: basic same-shape elementwise (Pythagorean triples for easy eyeballing)
	253	+ allPass &= RunScenario("basic same-shape (4,2)", {4, 2}, {4, 2}, {4, 2}, {3, 4, 5, 12, 8, 15, 7, 24},
	254	+ {4, 3, 12, 5, 15, 8, 24, 7}, stream);
	255	+
	256	+ // Scenario 2: broadcast (3,1) x (1,4) -> (3,4)
	257	+ allPass &= RunScenario("broadcast (3,1)x(1,4)->(3,4)", {3, 1}, {1, 4}, {3, 4}, {3, 5, 9}, {4, 12, 40, 7}, stream);
	258	+
	259	+ // Scenario 3: special values: inf priority over NaN, NaN propagation, signed zero,
	260	+ // 1e30 numerical stability (no overflow of intermediate squares)
	261	+ const float kInf = std::numeric_limits<float>::infinity();
	262	+ const float kQnan = std::numeric_limits<float>::quiet_NaN();
	263	+ allPass &= RunScenario("special values inf/NaN/zero/1e30", {8}, {8}, {8},
	264	+ {kInf, kQnan, 0.0F, -0.0F, 1.0e30F, -3.0F, kInf, -kInf},
	265	+ {kQnan, 2.0F, -0.0F, 5.0F, 1.0e30F, -4.0F, -7.0F, kQnan}, stream);
	266	+
	267	+ aclrtDestroyStream(stream);
	268	+ aclrtResetDevice(deviceId);
	269	+ aclFinalize();
	270	+
	271	+ LOG_PRINT("------------------------------------------------------------\n");
	272	+ LOG_PRINT("aclnnHypot example overall: %s\n", allPass ? "PASS" : "FAIL");
	273	+ return allPass ? 0 : 1;
	274	+}

Amath/hypot/examples/test_geir_hypot.cpp+224-0

@@ -0,0 +1,224 @@
	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	+ * @file test_geir_hypot.cpp
	11	+ * Hypot graph-mode (GE IR) call example.
	12	+ *
	13	+ * Builds a single-op graph Data(self, other) -> Hypot -> out aligned with
	14	+ * op_graph/hypot_proto.h (REG_OP(Hypot)), then runs it via GE Session
	15	+ * AddGraph/RunGraph. Inputs are FLOAT32 with broadcast shapes
	16	+ * self (4,2) all 3.0 and other (1,2) all 4.0, so every output element must be 5.0.
	17	+ *
	18	+ * Usage: ./test_geir_hypot [deviceId] (default deviceId = 0)
	19	+ */
	20	+
	21	+#include <cmath>
	22	+#include <cstdint>
	23	+#include <cstdio>
	24	+#include <cstdlib>
	25	+#include <ctime>
	26	+#include <map>
	27	+#include <string>
	28	+#include <vector>
	29	+
	30	+#include "ge_api.h"
	31	+#include "ge_api_types.h"
	32	+#include "ge_error_codes.h"
	33	+#include "ge_ir_build.h"
	34	+#include "graph.h"
	35	+#include "tensor.h"
	36	+#include "types.h"
	37	+
	38	+// op::Data (graph input placeholder) comes from the GE built-in op proto header.
	39	+#include "array_ops.h"
	40	+// Custom proto of this operator (REG_OP(Hypot)), consistent with op_graph registration;
	41	+// provides ge::op::Hypot.
	42	+#include "../op_graph/hypot_proto.h"
	43	+
	44	+#define FAILED (-1)
	45	+#define SUCCESS 0
	46	+
	47	+using namespace ge;
	48	+
	49	+#define LOG_PRINT(message, ...) \
	50	+ do { \
	51	+ printf(message, ##__VA_ARGS__); \
	52	+ } while (0)
	53	+
	54	+static std::string GetTime()
	55	+{
	56	+ time_t timep;
	57	+ time(&timep);
	58	+ char tmp[64] = {0};
	59	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S", localtime(&timep));
	60	+ return std::string(tmp);
	61	+}
	62	+
	63	+// Fill a FLOAT32 tensor with a constant value.
	64	+static int32_t GenConstFloatData(const std::vector<int64_t>& shapes, Tensor& tensor, TensorDesc& desc, float value)
	65	+{
	66	+ desc.SetRealDimCnt(static_cast<uint32_t>(shapes.size()));
	67	+ size_t size = 1;
	68	+ for (auto d : shapes) {
	69	+ size *= static_cast<size_t>(d);
	70	+ }
	71	+ std::vector<float> buf(size, value);
	72	+ tensor = Tensor(desc, reinterpret_cast<uint8_t>(buf.data()), size sizeof(float));
	73	+ return SUCCESS;
	74	+}
	75	+
	76	+// Build the single-op graph: placeholder(self), placeholder(other) -> Hypot -> out (broadcast (4,2)x(1,2)).
	77	+static int BuildHypotGraph(std::vector<Tensor>& graphInputs, std::vector<Operator>& inputs,
	78	+ std::vector<Operator>& outputs, Graph& graph)
	79	+{
	80	+ Status ret = SUCCESS;
	81	+ const DataType dtype = DT_FLOAT;
	82	+
	83	+ std::vector<int64_t> selfShape = {4, 2};
	84	+ std::vector<int64_t> otherShape = {1, 2};
	85	+ TensorDesc selfDesc(ge::Shape(selfShape), FORMAT_ND, dtype);
	86	+ TensorDesc otherDesc(ge::Shape(otherShape), FORMAT_ND, dtype);
	87	+ selfDesc.SetPlacement(ge::kPlacementHost);
	88	+ otherDesc.SetPlacement(ge::kPlacementHost);
	89	+
	90	+ auto placeholderSelf = op::Data("self_placeholder").set_attr_index(0);
	91	+ placeholderSelf.update_input_desc_x(selfDesc);
	92	+ placeholderSelf.update_output_desc_y(selfDesc);
	93	+
	94	+ auto placeholderOther = op::Data("other_placeholder").set_attr_index(1);
	95	+ placeholderOther.update_input_desc_x(otherDesc);
	96	+ placeholderOther.update_output_desc_y(otherDesc);
	97	+
	98	+ // Runtime data: self = 3.0, other = 4.0 -> out = hypot(3,4) = 5.0 everywhere.
	99	+ Tensor selfTensor;
	100	+ Tensor otherTensor;
	101	+ ret = GenConstFloatData(selfShape, selfTensor, selfDesc, 3.0F);
	102	+ if (ret != SUCCESS) {
	103	+ LOG_PRINT("%s - ERROR - [XIR]: Generate self data failed\n", GetTime().c_str());
	104	+ return FAILED;
	105	+ }
	106	+ ret = GenConstFloatData(otherShape, otherTensor, otherDesc, 4.0F);
	107	+ if (ret != SUCCESS) {
	108	+ LOG_PRINT("%s - ERROR - [XIR]: Generate other data failed\n", GetTime().c_str());
	109	+ return FAILED;
	110	+ }
	111	+ graphInputs.push_back(selfTensor);
	112	+ graphInputs.push_back(otherTensor);
	113	+
	114	+ // Hypot node (OpType consistent with REG_OP(Hypot)).
	115	+ auto hypot = op::Hypot("hypot_0");
	116	+ hypot.set_input_self(placeholderSelf);
	117	+ hypot.set_input_other(placeholderOther);
	118	+ hypot.update_input_desc_self(selfDesc);
	119	+ hypot.update_input_desc_other(otherDesc);
	120	+ TensorDesc outDesc(ge::Shape(selfShape), FORMAT_ND, dtype); // broadcast result (4,2)
	121	+ hypot.update_output_desc_out(outDesc);
	122	+
	123	+ graph.AddOp(placeholderSelf);
	124	+ graph.AddOp(placeholderOther);
	125	+ graph.AddOp(hypot);
	126	+
	127	+ inputs.push_back(placeholderSelf);
	128	+ inputs.push_back(placeholderOther);
	129	+ outputs.push_back(hypot);
	130	+ return SUCCESS;
	131	+}
	132	+
	133	+int main(int argc, char* argv[])
	134	+{
	135	+ (void)argc;
	136	+ const char* deviceId = (argc > 1) ? argv[1] : "0";
	137	+ LOG_PRINT("=== Hypot GE IR (graph mode) example ===\n");
	138	+ LOG_PRINT("deviceId: %s\n", deviceId);
	139	+
	140	+ const char* graphName = "hypot_geir_test";
	141	+ Graph graph(graphName);
	142	+ std::vector<Tensor> graphInputs;
	143	+ std::vector<Operator> inputs;
	144	+ std::vector<Operator> outputs;
	145	+
	146	+ LOG_PRINT("%s - INFO - [XIR]: GEInitialize\n", GetTime().c_str());
	147	+ std::map<AscendString, AscendString> globalOptions = {
	148	+ {"ge.exec.deviceId", deviceId},
	149	+ {"ge.graphRunMode", "1"},
	150	+ };
	151	+ Status ret = ge::GEInitialize(globalOptions);
	152	+ if (ret != SUCCESS) {
	153	+ LOG_PRINT("%s - ERROR - [XIR]: GEInitialize failed\n", GetTime().c_str());
	154	+ return FAILED;
	155	+ }
	156	+
	157	+ ret = BuildHypotGraph(graphInputs, inputs, outputs, graph);
	158	+ if (ret != SUCCESS) {
	159	+ LOG_PRINT("%s - ERROR - [XIR]: Build graph failed\n", GetTime().c_str());
	160	+ ge::GEFinalize();
	161	+ return FAILED;
	162	+ }
	163	+ if (!inputs.empty() && !outputs.empty()) {
	164	+ graph.SetInputs(inputs).SetOutputs(outputs);
	165	+ }
	166	+
	167	+ LOG_PRINT("%s - INFO - [XIR]: create Session\n", GetTime().c_str());
	168	+ std::map<AscendString, AscendString> buildOptions;
	169	+ ge::Session* session = new (std::nothrow) Session(buildOptions);
	170	+ if (session == nullptr) {
	171	+ LOG_PRINT("%s - ERROR - [XIR]: create Session failed\n", GetTime().c_str());
	172	+ ge::GEFinalize();
	173	+ return FAILED;
	174	+ }
	175	+
	176	+ uint32_t graphId = 0;
	177	+ std::map<AscendString, AscendString> graphOptions;
	178	+ ret = session->AddGraph(graphId, graph, graphOptions);
	179	+ if (ret != SUCCESS) {
	180	+ LOG_PRINT("%s - ERROR - [XIR]: AddGraph failed\n", GetTime().c_str());
	181	+ delete session;
	182	+ ge::GEFinalize();
	183	+ return FAILED;
	184	+ }
	185	+ LOG_PRINT("%s - INFO - [XIR]: AddGraph success\n", GetTime().c_str());
	186	+
	187	+ LOG_PRINT("%s - INFO - [XIR]: RunGraph\n", GetTime().c_str());
	188	+ std::vector<Tensor> graphOutputs;
	189	+ ret = session->RunGraph(graphId, graphInputs, graphOutputs);
	190	+ if (ret != SUCCESS) {
	191	+ LOG_PRINT("%s - ERROR - [XIR]: RunGraph failed\n", GetTime().c_str());
	192	+ ge::AscendString errorMsg = ge::GEGetErrorMsgV2();
	193	+ LOG_PRINT("%s - ERROR - [XIR]: GE error message: %s\n", GetTime().c_str(), errorMsg.GetString());
	194	+ delete session;
	195	+ ge::GEFinalize();
	196	+ return FAILED;
	197	+ }
	198	+ LOG_PRINT("%s - INFO - [XIR]: RunGraph success\n", GetTime().c_str());
	199	+
	200	+ // Print outputs and verify against golden (hypot(3,4) = 5).
	201	+ bool pass = true;
	202	+ for (size_t i = 0; i < graphOutputs.size(); ++i) {
	203	+ const float* out = reinterpret_cast<const float*>(graphOutputs[i].GetData());
	204	+ int64_t numel = graphOutputs[i].GetTensorDesc().GetShape().GetShapeSize();
	205	+ for (int64_t j = 0; j < numel; ++j) {
	206	+ double golden = std::hypot(3.0, 4.0); // 5.0
	207	+ bool ok = std::fabs(static_cast<double>(out[j]) - golden) <= 1e-5;
	208	+ if (!ok) {
	209	+ pass = false;
	210	+ }
	211	+ LOG_PRINT(" out[%ld] = %f golden=%f %s\n", j, out[j], golden, ok ? "OK" : "MISMATCH");
	212	+ }
	213	+ }
	214	+
	215	+ delete session;
	216	+ ret = ge::GEFinalize();
	217	+ if (ret != SUCCESS) {
	218	+ LOG_PRINT("%s - ERROR - [XIR]: GEFinalize failed\n", GetTime().c_str());
	219	+ return FAILED;
	220	+ }
	221	+ LOG_PRINT("%s - INFO - [XIR]: done\n", GetTime().c_str());
	222	+ LOG_PRINT("Hypot GE IR example overall: %s\n", pass ? "PASS" : "FAIL");
	223	+ return pass ? SUCCESS : FAILED;
	224	+}

Amath/hypot/op_api/aclnn_hypot.cpp+147-0

@@ -0,0 +1,147 @@
	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 aclnn_hypot.cpp
	13	+ * \brief hypot L2 aclnn two-phase interface
	14	+ */
	15	+
	16	+#include "aclnn_hypot.h"
	17	+#include "hypot.h"
	18	+#include "aclnn_kernels/cast.h"
	19	+#include "aclnn_kernels/contiguous.h"
	20	+#include "aclnn_kernels/common/op_error_check.h"
	21	+#include "op_api/aclnn_check.h"
	22	+#include "opdev/common_types.h"
	23	+#include "opdev/data_type_utils.h"
	24	+#include "op_api/data_type_utils.h"
	25	+#include "opdev/format_utils.h"
	26	+#include "opdev/op_dfx.h"
	27	+#include "opdev/op_executor.h"
	28	+#include "opdev/op_log.h"
	29	+#include "opdev/shape_utils.h"
	30	+#include "opdev/tensor_view_utils.h"
	31	+#include "opdev/platform.h"
	32	+
	33	+using namespace op;
	34	+#ifdef __cplusplus
	35	+extern "C" {
	36	+#endif
	37	+
	38	+constexpr size_t MAX_DIM_LEN = 8;
	39	+
	40	+// FLOAT / FLOAT16 / BFLOAT16（异 dtype 组合由 kernel 内 DTYPE_SELF/DTYPE_OTHER 直读，
	41	+// 计算在 FP32 寄存器完成，不再插入 Cast 拼接算子）
	42	+static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16,
	43	+ op::DataType::DT_BF16};
	44	+
	45	+static bool CheckNotNull(const aclTensor* self, const aclTensor* other, const aclTensor* out)
	46	+{
	47	+ OP_CHECK_NULL(self, return false);
	48	+ OP_CHECK_NULL(other, return false);
	49	+ OP_CHECK_NULL(out, return false);
	50	+ return true;
	51	+}
	52	+
	53	+static bool CheckDtypeValid(const aclTensor* self, const aclTensor* other, const aclTensor* out)
	54	+{
	55	+ OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST, return false);
	56	+ OP_CHECK_DTYPE_NOT_SUPPORT(other, DTYPE_SUPPORT_LIST, return false);
	57	+ OP_CHECK_DTYPE_NOT_SUPPORT(out, DTYPE_SUPPORT_LIST, return false);
	58	+ return true;
	59	+}
	60	+
	61	+static bool CheckComputeDtype(const aclTensor* self, const aclTensor* other, const aclTensor* out)
	62	+{
	63	+ // dtype 互推导（np.promote_types 语义）
	64	+ op::DataType promoteType = op::PromoteType(self->GetDataType(), other->GetDataType());
	65	+ if (promoteType == DataType::DT_UNDEFINED) {
	66	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "self dtype %s and other dtype %s can not promote dtype.",
	67	+ op::ToString(self->GetDataType()).GetString(), op::ToString(other->GetDataType()).GetString());
	68	+ return false;
	69	+ }
	70	+ // 推导结果须为 FLOAT / FLOAT16 / BFLOAT16 之一
	71	+ if (!CheckType(promoteType, DTYPE_SUPPORT_LIST)) {
	72	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Hypot promote dtype %s is not supported.",
	73	+ op::ToString(promoteType).GetString());
	74	+ return false;
	75	+ }
	76	+ // 检查推导后的数据类型能否转换为输出的数据类型
	77	+ OP_CHECK_RESULT_DTYPE_CAST_FAILED(promoteType, out->GetDataType(), return false);
	78	+ return true;
	79	+}
	80	+
	81	+static bool CheckShape(const aclTensor* self, const aclTensor* other, const aclTensor* out)
	82	+{
	83	+ OP_CHECK_MAX_DIM(self, MAX_DIM_LEN, return false);
	84	+ OP_CHECK_MAX_DIM(other, MAX_DIM_LEN, return false);
	85	+
	86	+ op::Shape broadcastShape;
	87	+ OP_CHECK_BROADCAST_AND_INFER_SHAPE(self, other, broadcastShape, return false);
	88	+ OP_CHECK_SHAPE_NOT_EQUAL_WITH_EXPECTED_SIZE(out, broadcastShape, return false);
	89	+ return true;
	90	+}
	91	+
	92	+static aclnnStatus CheckParams(const aclTensor* self, const aclTensor* other, const aclTensor* out)
	93	+{
	94	+ // null pointer check failure reports ACLNN_ERR_PARAM_INVALID (161002)
	95	+ CHECK_RET(CheckNotNull(self, other, out), ACLNN_ERR_PARAM_INVALID);
	96	+ CHECK_RET(CheckDtypeValid(self, other, out), ACLNN_ERR_PARAM_INVALID);
	97	+ CHECK_RET(CheckComputeDtype(self, other, out), ACLNN_ERR_PARAM_INVALID);
	98	+ CHECK_RET(CheckShape(self, other, out), ACLNN_ERR_PARAM_INVALID);
	99	+ return ACLNN_SUCCESS;
	100	+}
	101	+
	102	+aclnnStatus aclnnHypotGetWorkspaceSize(const aclTensor* self, const aclTensor* other, aclTensor* out,
	103	+ uint64_t* workspaceSize, aclOpExecutor** executor)
	104	+{
	105	+ L2_DFX_PHASE_1(aclnnHypot, DFX_IN(self, other), DFX_OUT(out));
	106	+ auto uniqueExecutor = CREATE_EXECUTOR();
	107	+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
	108	+ auto ret = CheckParams(self, other, out);
	109	+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
	110	+ if (self->GetStorageFormat() != Format::FORMAT_ND) {
	111	+ OP_LOGW("Format only supports ND");
	112	+ }
	113	+ if (self->IsEmpty() \|\| other->IsEmpty()) {
	114	+ *workspaceSize = 0;
	115	+ uniqueExecutor.ReleaseTo(executor);
	116	+ return ACLNN_SUCCESS;
	117	+ }
	118	+ DataType computeType = op::PromoteType(self->GetDataType(), other->GetDataType());
	119	+ CHECK_RET(computeType != DataType::DT_UNDEFINED, ACLNN_ERR_INNER);
	120	+ bool needOutputCast = (computeType != out->GetDataType());
	121	+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
	122	+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	123	+ auto otherContiguous = l0op::Contiguous(other, uniqueExecutor.get());
	124	+ CHECK_RET(otherContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	125	+ const aclTensor* hypotOpOut = l0op::Hypot(selfContiguous, otherContiguous, uniqueExecutor.get());
	126	+ CHECK_RET(hypotOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
	127	+ const aclTensor* finalOut = hypotOpOut;
	128	+ if (needOutputCast) {
	129	+ finalOut = l0op::Cast(hypotOpOut, out->GetDataType(), uniqueExecutor.get());
	130	+ CHECK_RET(finalOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
	131	+ }
	132	+ auto viewCopyResult = l0op::ViewCopy(finalOut, out, uniqueExecutor.get());
	133	+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
	134	+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
	135	+ uniqueExecutor.ReleaseTo(executor);
	136	+ return ACLNN_SUCCESS;
	137	+}
	138	+
	139	+aclnnStatus aclnnHypot(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
	140	+{
	141	+ L2_DFX_PHASE_2(aclnnHypot);
	142	+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
	143	+}
	144	+
	145	+#ifdef __cplusplus
	146	+}
	147	+#endif

Amath/hypot/op_api/aclnn_hypot.h+55-0

@@ -0,0 +1,55 @@
	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	+#ifndef OP_API_INC_HYPOT_H_
	12	+#define OP_API_INC_HYPOT_H_
	13	+
	14	+#include "aclnn/aclnn_base.h"
	15	+#include "aclnn_util.h"
	16	+
	17	+#ifdef __cplusplus
	18	+extern "C" {
	19	+#endif
	20	+
	21	+/**
	22	+ * @brief aclnnHypot的第一段接口，根据具体的计算流程，计算workspace大小。
	23	+ * @domain aclnn_math
	24	+ *
	25	+ * 算子功能：计算两个张量对应元素直角三角形斜边 sqrt(self^2 + other^2)，数值稳定实现。
	26	+ *
	27	+ * @param [in] self: npu device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16，
	28	+ * 数据类型需要与other构成互相推导关系，shape需要与other满足broadcast关系。支持非连续的Tensor，数据格式支持ND。
	29	+ * @param [in] other: npu device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16，
	30	+ * 数据类型需要与self构成互相推导关系，shape需要与self满足broadcast关系。支持非连续的Tensor，数据格式支持ND。
	31	+ * @param [out] out: npu device侧的aclTensor，数据类型为self与other的推导结果，
	32	+ * shape需要是self与other broadcast之后的shape。支持非连续的Tensor，数据格式支持ND。
	33	+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
	34	+ * @param [out] executor: 返回op执行器，包含算子计算流程。
	35	+ * @return aclnnStatus: 返回状态码。
	36	+ */
	37	+ACLNN_API aclnnStatus aclnnHypotGetWorkspaceSize(const aclTensor* self, const aclTensor* other, aclTensor* out,
	38	+ uint64_t* workspaceSize, aclOpExecutor** executor);
	39	+
	40	+/**
	41	+ * @brief aclnnHypot的第二段接口，用于执行计算。
	42	+ *
	43	+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
	44	+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnHypotGetWorkspaceSize获取。
	45	+ * @param [in] executor: op执行器，包含了算子计算流程。
	46	+ * @param [in] stream: acl stream流。
	47	+ * @return aclnnStatus: 返回状态码。
	48	+ */
	49	+ACLNN_API aclnnStatus aclnnHypot(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
	50	+
	51	+#ifdef __cplusplus
	52	+}
	53	+#endif
	54	+
	55	+#endif // OP_API_INC_HYPOT_H_

Amath/hypot/op_api/hypot.cpp+71-0

@@ -0,0 +1,71 @@
	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 hypot.cpp
	13	+ * \brief hypot L0 implementation: broadcast infershape + AiCore launch
	14	+ */
	15	+
	16	+#include "hypot.h"
	17	+#include "opdev/make_op_executor.h"
	18	+#include "opdev/op_def.h"
	19	+#include "opdev/op_dfx.h"
	20	+#include "opdev/op_executor.h"
	21	+#include "opdev/op_log.h"
	22	+#include "opdev/shape_utils.h"
	23	+#include "opdev/data_type_utils.h"
	24	+#include "op_api/aclnn_check.h"
	25	+
	26	+using namespace op;
	27	+
	28	+namespace l0op {
	29	+
	30	+OP_TYPE_REGISTER(Hypot);
	31	+
	32	+// AiCore 支持 FLOAT / FLOAT16 / BFLOAT16
	33	+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
	34	+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
	35	+
	36	+// 根据dtype判断算子是否支持走AiCore
	37	+static bool IsAiCoreSupport(const aclTensor* self, const aclTensor* other)
	38	+{
	39	+ return CheckType(self->GetDataType(), AICORE_DTYPE_SUPPORT_LIST) &&
	40	+ CheckType(other->GetDataType(), AICORE_DTYPE_SUPPORT_LIST);
	41	+}
	42	+
	43	+// AICORE算子kernel
	44	+static const aclTensor* HypotAiCore(const aclTensor* self, const aclTensor* other, const aclTensor* hypotOut,
	45	+ aclOpExecutor* executor)
	46	+{
	47	+ L0_DFX(HypotAiCore, self, other, hypotOut)
	48	+ // 使用框架宏ADD_TO_LAUNCHER_LIST_AICORE，将Hypot算子加入任务队列
	49	+ ADD_TO_LAUNCHER_LIST_AICORE(Hypot, OP_INPUT(self, other), OP_OUTPUT(hypotOut));
	50	+ return hypotOut;
	51	+}
	52	+
	53	+const aclTensor* Hypot(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor)
	54	+{
	55	+ op::Shape broadcastShape;
	56	+ if (!BroadcastInferShape(self->GetViewShape(), other->GetViewShape(), broadcastShape)) {
	57	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Broadcast %s and %s failed.", op::ToString(self->GetViewShape()).GetString(),
	58	+ op::ToString(other->GetViewShape()).GetString());
	59	+ return nullptr;
	60	+ }
	61	+
	62	+ DataType outType = PromoteType(self->GetDataType(), other->GetDataType());
	63	+ auto hypotOut = executor->AllocTensor(broadcastShape, outType);
	64	+ if (IsAiCoreSupport(self, other)) {
	65	+ return HypotAiCore(self, other, hypotOut, executor);
	66	+ }
	67	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Hypot dtype %s is not supported on AiCore.",
	68	+ op::ToString(self->GetDataType()).GetString());
	69	+ return nullptr;
	70	+}
	71	+} // namespace l0op

Amath/hypot/op_api/hypot.h+27-0

@@ -0,0 +1,27 @@
	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 hypot.h
	13	+ * \brief hypot L0 interface
	14	+ */
	15	+
	16	+#ifndef OP_API_INC_LEVEL0_HYPOT_H_
	17	+#define OP_API_INC_LEVEL0_HYPOT_H_
	18	+
	19	+#include "opdev/op_executor.h"
	20	+
	21	+namespace l0op {
	22	+
	23	+const aclTensor* Hypot(const aclTensor* self, const aclTensor* other, aclOpExecutor* executor);
	24	+
	25	+} // namespace l0op
	26	+
	27	+#endif // OP_API_INC_LEVEL0_HYPOT_H_

Amath/hypot/op_graph/hypot_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 hypot_proto.h
	13	+ * \brief hypot graph-mode operator proto
	14	+ */
	15	+#ifndef OPS_OP_PROTO_INC_HYPOT_H_
	16	+#define OPS_OP_PROTO_INC_HYPOT_H_
	17	+
	18	+#include "graph/operator_reg.h"
	19	+#include "graph/types.h"
	20	+
	21	+namespace ge {
	22	+
	23	+/**
	24	+*@brief Computes the hypotenuse sqrt(self^2 + other^2) element-wise with a numerically
	25	+*stable implementation. Supports broadcasting. \n
	26	+
	27	+*@par Inputs:
	28	+*Two inputs, including:
	29	+* @li self: A ND Tensor. Must be float16, float32 or bfloat16. \n
	30	+* @li other: A ND Tensor. Must be float16, float32 or bfloat16, and must be broadcast-compatible
	31	+* with "self". Its dtype is promoted with "self" to determine the output dtype. \n
	32	+
	33	+*@par Outputs:
	34	+*out: A ND Tensor. Has the promoted dtype of "self" and "other". \n
	35	+
	36	+*@par Third-party framework compatibility
	37	+*Compatible with the PyTorch operator hypot.
	38	+*/
	39	+REG_OP(Hypot)
	40	+ .INPUT(self, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
	41	+ .INPUT(other, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
	42	+ .OUTPUT(out, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
	43	+ .OP_END_FACTORY_REG(Hypot)
	44	+
	45	+} // namespace ge
	46	+
	47	+#endif // OPS_OP_PROTO_INC_HYPOT_H_

Amath/hypot/op_host/arch22/hypot_tiling_arch22.cpp+984-0

Amath/hypot/op_host/arch22/hypot_tiling_arch22.h+27-0

@@ -0,0 +1,27 @@
	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 hypot_tiling_arch22.h
	13	+ * \brief hypot arch22 tiling: broadcast fusion, scene decision, UB/multi-core split
	14	+ */
	15	+
	16	+#ifndef HYPOT_TILING_ARCH22_H
	17	+#define HYPOT_TILING_ARCH22_H
	18	+
	19	+#include "../../op_kernel/hypot_tiling_data.h"
	20	+
	21	+namespace optiling {
	22	+struct HypotCompileInfoArch22 {
	23	+ int32_t totalCoreNum = 0;
	24	+ int64_t ubSize = 0;
	25	+};
	26	+} // namespace optiling
	27	+#endif // HYPOT_TILING_ARCH22_H

Amath/hypot/op_host/config/ascend910_93/hypot_binary.json+365-0

@@ -0,0 +1,365 @@
	1	+{
	2	+ "op_type": "Hypot",
	3	+ "op_list": [
	4	+ {
	5	+ "bin_filename": "Hypot_0000000000000000000000000000fp16",
	6	+ "inputs": [
	7	+ {
	8	+ "name": "self",
	9	+ "index": 0,
	10	+ "dtype": "float16",
	11	+ "format": "ND",
	12	+ "paramType": "required",
	13	+ "shape": [
	14	+ -2
	15	+ ],
	16	+ "format_match_mode": "FormatAgnostic"
	17	+ },
	18	+ {
	19	+ "name": "other",
	20	+ "index": 1,
	21	+ "dtype": "float16",
	22	+ "format": "ND",
	23	+ "paramType": "required",
	24	+ "shape": [
	25	+ -2
	26	+ ],
	27	+ "format_match_mode": "FormatAgnostic"
	28	+ }
	29	+ ],
	30	+ "outputs": [
	31	+ {
	32	+ "name": "out",
	33	+ "index": 0,
	34	+ "dtype": "float16",
	35	+ "format": "ND",
	36	+ "paramType": "required",
	37	+ "shape": [
	38	+ -2
	39	+ ],
	40	+ "format_match_mode": "FormatAgnostic"
	41	+ }
	42	+ ]
	43	+ },
	44	+ {
	45	+ "bin_filename": "Hypot_0000000000000000000000000000fp32",
	46	+ "inputs": [
	47	+ {
	48	+ "name": "self",
	49	+ "index": 0,
	50	+ "dtype": "float32",
	51	+ "format": "ND",
	52	+ "paramType": "required",
	53	+ "shape": [
	54	+ -2
	55	+ ],
	56	+ "format_match_mode": "FormatAgnostic"
	57	+ },
	58	+ {
	59	+ "name": "other",
	60	+ "index": 1,
	61	+ "dtype": "float32",
	62	+ "format": "ND",
	63	+ "paramType": "required",
	64	+ "shape": [
	65	+ -2
	66	+ ],
	67	+ "format_match_mode": "FormatAgnostic"
	68	+ }
	69	+ ],
	70	+ "outputs": [
	71	+ {
	72	+ "name": "out",
	73	+ "index": 0,
	74	+ "dtype": "float32",
	75	+ "format": "ND",
	76	+ "paramType": "required",
	77	+ "shape": [
	78	+ -2
	79	+ ],
	80	+ "format_match_mode": "FormatAgnostic"
	81	+ }
	82	+ ]
	83	+ },
	84	+ {
	85	+ "bin_filename": "Hypot_0000000000000000000000000000bf16",
	86	+ "inputs": [
	87	+ {
	88	+ "name": "self",
	89	+ "index": 0,
	90	+ "dtype": "bfloat16",
	91	+ "format": "ND",
	92	+ "paramType": "required",
	93	+ "shape": [
	94	+ -2
	95	+ ],
	96	+ "format_match_mode": "FormatAgnostic"
	97	+ },
	98	+ {
	99	+ "name": "other",
	100	+ "index": 1,
	101	+ "dtype": "bfloat16",
	102	+ "format": "ND",
	103	+ "paramType": "required",
	104	+ "shape": [
	105	+ -2
	106	+ ],
	107	+ "format_match_mode": "FormatAgnostic"
	108	+ }
	109	+ ],
	110	+ "outputs": [
	111	+ {
	112	+ "name": "out",
	113	+ "index": 0,
	114	+ "dtype": "bfloat16",
	115	+ "format": "ND",
	116	+ "paramType": "required",
	117	+ "shape": [
	118	+ -2
	119	+ ],
	120	+ "format_match_mode": "FormatAgnostic"
	121	+ }
	122	+ ]
	123	+ },
	124	+ {
	125	+ "bin_filename": "Hypot_0000000000000000000000000000mix_f16_bf16",
	126	+ "inputs": [
	127	+ {
	128	+ "name": "self",
	129	+ "index": 0,
	130	+ "dtype": "float16",
	131	+ "format": "ND",
	132	+ "paramType": "required",
	133	+ "shape": [
	134	+ -2
	135	+ ],
	136	+ "format_match_mode": "FormatAgnostic"
	137	+ },
	138	+ {
	139	+ "name": "other",
	140	+ "index": 1,
	141	+ "dtype": "bfloat16",
	142	+ "format": "ND",
	143	+ "paramType": "required",
	144	+ "shape": [
	145	+ -2
	146	+ ],
	147	+ "format_match_mode": "FormatAgnostic"
	148	+ }
	149	+ ],
	150	+ "outputs": [
	151	+ {
	152	+ "name": "out",
	153	+ "index": 0,
	154	+ "dtype": "float32",
	155	+ "format": "ND",
	156	+ "paramType": "required",
	157	+ "shape": [
	158	+ -2
	159	+ ],
	160	+ "format_match_mode": "FormatAgnostic"
	161	+ }
	162	+ ]
	163	+ },
	164	+ {
	165	+ "bin_filename": "Hypot_0000000000000000000000000000mix_bf16_f16",
	166	+ "inputs": [
	167	+ {
	168	+ "name": "self",
	169	+ "index": 0,
	170	+ "dtype": "bfloat16",
	171	+ "format": "ND",
	172	+ "paramType": "required",
	173	+ "shape": [
	174	+ -2
	175	+ ],
	176	+ "format_match_mode": "FormatAgnostic"
	177	+ },
	178	+ {
	179	+ "name": "other",
	180	+ "index": 1,
	181	+ "dtype": "float16",
	182	+ "format": "ND",
	183	+ "paramType": "required",
	184	+ "shape": [
	185	+ -2
	186	+ ],
	187	+ "format_match_mode": "FormatAgnostic"
	188	+ }
	189	+ ],
	190	+ "outputs": [
	191	+ {
	192	+ "name": "out",
	193	+ "index": 0,
	194	+ "dtype": "float32",
	195	+ "format": "ND",
	196	+ "paramType": "required",
	197	+ "shape": [
	198	+ -2
	199	+ ],
	200	+ "format_match_mode": "FormatAgnostic"
	201	+ }
	202	+ ]
	203	+ },
	204	+ {
	205	+ "bin_filename": "Hypot_0000000000000000000000000000mix_f16_f32",
	206	+ "inputs": [
	207	+ {
	208	+ "name": "self",
	209	+ "index": 0,
	210	+ "dtype": "float16",
	211	+ "format": "ND",
	212	+ "paramType": "required",
	213	+ "shape": [
	214	+ -2
	215	+ ],
	216	+ "format_match_mode": "FormatAgnostic"
	217	+ },
	218	+ {
	219	+ "name": "other",
	220	+ "index": 1,
	221	+ "dtype": "float32",
	222	+ "format": "ND",
	223	+ "paramType": "required",
	224	+ "shape": [
	225	+ -2
	226	+ ],
	227	+ "format_match_mode": "FormatAgnostic"
	228	+ }
	229	+ ],
	230	+ "outputs": [
	231	+ {
	232	+ "name": "out",
	233	+ "index": 0,
	234	+ "dtype": "float32",
	235	+ "format": "ND",
	236	+ "paramType": "required",
	237	+ "shape": [
	238	+ -2
	239	+ ],
	240	+ "format_match_mode": "FormatAgnostic"
	241	+ }
	242	+ ]
	243	+ },
	244	+ {
	245	+ "bin_filename": "Hypot_0000000000000000000000000000mix_f32_f16",
	246	+ "inputs": [
	247	+ {
	248	+ "name": "self",
	249	+ "index": 0,
	250	+ "dtype": "float32",
	251	+ "format": "ND",
	252	+ "paramType": "required",
	253	+ "shape": [
	254	+ -2
	255	+ ],
	256	+ "format_match_mode": "FormatAgnostic"
	257	+ },
	258	+ {
	259	+ "name": "other",
	260	+ "index": 1,
	261	+ "dtype": "float16",
	262	+ "format": "ND",
	263	+ "paramType": "required",
	264	+ "shape": [
	265	+ -2
	266	+ ],
	267	+ "format_match_mode": "FormatAgnostic"
	268	+ }
	269	+ ],
	270	+ "outputs": [
	271	+ {
	272	+ "name": "out",
	273	+ "index": 0,
	274	+ "dtype": "float32",
	275	+ "format": "ND",
	276	+ "paramType": "required",
	277	+ "shape": [
	278	+ -2
	279	+ ],
	280	+ "format_match_mode": "FormatAgnostic"
	281	+ }
	282	+ ]
	283	+ },
	284	+ {
	285	+ "bin_filename": "Hypot_0000000000000000000000000000mix_bf16_f32",
	286	+ "inputs": [
	287	+ {
	288	+ "name": "self",
	289	+ "index": 0,
	290	+ "dtype": "bfloat16",
	291	+ "format": "ND",
	292	+ "paramType": "required",
	293	+ "shape": [
	294	+ -2
	295	+ ],
	296	+ "format_match_mode": "FormatAgnostic"
	297	+ },
	298	+ {
	299	+ "name": "other",
	300	+ "index": 1,
	301	+ "dtype": "float32",
	302	+ "format": "ND",
	303	+ "paramType": "required",
	304	+ "shape": [
	305	+ -2
	306	+ ],
	307	+ "format_match_mode": "FormatAgnostic"
	308	+ }
	309	+ ],
	310	+ "outputs": [
	311	+ {
	312	+ "name": "out",
	313	+ "index": 0,
	314	+ "dtype": "float32",
	315	+ "format": "ND",
	316	+ "paramType": "required",
	317	+ "shape": [
	318	+ -2
	319	+ ],
	320	+ "format_match_mode": "FormatAgnostic"
	321	+ }
	322	+ ]
	323	+ },
	324	+ {
	325	+ "bin_filename": "Hypot_0000000000000000000000000000mix_f32_bf16",
	326	+ "inputs": [
	327	+ {
	328	+ "name": "self",
	329	+ "index": 0,
	330	+ "dtype": "float32",
	331	+ "format": "ND",
	332	+ "paramType": "required",
	333	+ "shape": [
	334	+ -2
	335	+ ],
	336	+ "format_match_mode": "FormatAgnostic"
	337	+ },
	338	+ {
	339	+ "name": "other",
	340	+ "index": 1,
	341	+ "dtype": "bfloat16",
	342	+ "format": "ND",
	343	+ "paramType": "required",
	344	+ "shape": [
	345	+ -2
	346	+ ],
	347	+ "format_match_mode": "FormatAgnostic"
	348	+ }
	349	+ ],
	350	+ "outputs": [
	351	+ {
	352	+ "name": "out",
	353	+ "index": 0,
	354	+ "dtype": "float32",
	355	+ "format": "ND",
	356	+ "paramType": "required",
	357	+ "shape": [
	358	+ -2
	359	+ ],
	360	+ "format_match_mode": "FormatAgnostic"
	361	+ }
	362	+ ]
	363	+ }
	364	+ ]
	365	+}

Amath/hypot/op_host/config/ascend910_93/hypot_simplified_key.ini+2-0

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

Amath/hypot/op_host/config/ascend910b/hypot_binary.json+365-0

@@ -0,0 +1,365 @@
	1	+{
	2	+ "op_type": "Hypot",
	3	+ "op_list": [
	4	+ {
	5	+ "bin_filename": "Hypot_0000000000000000000000000000fp16",
	6	+ "inputs": [
	7	+ {
	8	+ "name": "self",
	9	+ "index": 0,
	10	+ "dtype": "float16",
	11	+ "format": "ND",
	12	+ "paramType": "required",
	13	+ "shape": [
	14	+ -2
	15	+ ],
	16	+ "format_match_mode": "FormatAgnostic"
	17	+ },
	18	+ {
	19	+ "name": "other",
	20	+ "index": 1,
	21	+ "dtype": "float16",
	22	+ "format": "ND",
	23	+ "paramType": "required",
	24	+ "shape": [
	25	+ -2
	26	+ ],
	27	+ "format_match_mode": "FormatAgnostic"
	28	+ }
	29	+ ],
	30	+ "outputs": [
	31	+ {
	32	+ "name": "out",
	33	+ "index": 0,
	34	+ "dtype": "float16",
	35	+ "format": "ND",
	36	+ "paramType": "required",
	37	+ "shape": [
	38	+ -2
	39	+ ],
	40	+ "format_match_mode": "FormatAgnostic"
	41	+ }
	42	+ ]
	43	+ },
	44	+ {
	45	+ "bin_filename": "Hypot_0000000000000000000000000000fp32",
	46	+ "inputs": [
	47	+ {
	48	+ "name": "self",
	49	+ "index": 0,
	50	+ "dtype": "float32",
	51	+ "format": "ND",
	52	+ "paramType": "required",
	53	+ "shape": [
	54	+ -2
	55	+ ],
	56	+ "format_match_mode": "FormatAgnostic"
	57	+ },
	58	+ {
	59	+ "name": "other",
	60	+ "index": 1,
	61	+ "dtype": "float32",
	62	+ "format": "ND",
	63	+ "paramType": "required",
	64	+ "shape": [
	65	+ -2
	66	+ ],
	67	+ "format_match_mode": "FormatAgnostic"
	68	+ }
	69	+ ],
	70	+ "outputs": [
	71	+ {
	72	+ "name": "out",
	73	+ "index": 0,
	74	+ "dtype": "float32",
	75	+ "format": "ND",
	76	+ "paramType": "required",
	77	+ "shape": [
	78	+ -2
	79	+ ],
	80	+ "format_match_mode": "FormatAgnostic"
	81	+ }
	82	+ ]
	83	+ },
	84	+ {
	85	+ "bin_filename": "Hypot_0000000000000000000000000000bf16",
	86	+ "inputs": [
	87	+ {
	88	+ "name": "self",
	89	+ "index": 0,
	90	+ "dtype": "bfloat16",
	91	+ "format": "ND",
	92	+ "paramType": "required",
	93	+ "shape": [
	94	+ -2
	95	+ ],
	96	+ "format_match_mode": "FormatAgnostic"
	97	+ },
	98	+ {
	99	+ "name": "other",
	100	+ "index": 1,
	101	+ "dtype": "bfloat16",
	102	+ "format": "ND",
	103	+ "paramType": "required",
	104	+ "shape": [
	105	+ -2
	106	+ ],
	107	+ "format_match_mode": "FormatAgnostic"
	108	+ }
	109	+ ],
	110	+ "outputs": [
	111	+ {
	112	+ "name": "out",
	113	+ "index": 0,
	114	+ "dtype": "bfloat16",
	115	+ "format": "ND",
	116	+ "paramType": "required",
	117	+ "shape": [
	118	+ -2
	119	+ ],
	120	+ "format_match_mode": "FormatAgnostic"
	121	+ }
	122	+ ]
	123	+ },
	124	+ {
	125	+ "bin_filename": "Hypot_0000000000000000000000000000mix_f16_bf16",
	126	+ "inputs": [
	127	+ {
	128	+ "name": "self",
	129	+ "index": 0,
	130	+ "dtype": "float16",
	131	+ "format": "ND",
	132	+ "paramType": "required",
	133	+ "shape": [
	134	+ -2
	135	+ ],
	136	+ "format_match_mode": "FormatAgnostic"
	137	+ },
	138	+ {
	139	+ "name": "other",
	140	+ "index": 1,
	141	+ "dtype": "bfloat16",
	142	+ "format": "ND",
	143	+ "paramType": "required",
	144	+ "shape": [
	145	+ -2
	146	+ ],
	147	+ "format_match_mode": "FormatAgnostic"
	148	+ }
	149	+ ],
	150	+ "outputs": [
	151	+ {
	152	+ "name": "out",
	153	+ "index": 0,
	154	+ "dtype": "float32",
	155	+ "format": "ND",
	156	+ "paramType": "required",
	157	+ "shape": [
	158	+ -2
	159	+ ],
	160	+ "format_match_mode": "FormatAgnostic"
	161	+ }
	162	+ ]
	163	+ },
	164	+ {
	165	+ "bin_filename": "Hypot_0000000000000000000000000000mix_bf16_f16",
	166	+ "inputs": [
	167	+ {
	168	+ "name": "self",
	169	+ "index": 0,
	170	+ "dtype": "bfloat16",
	171	+ "format": "ND",
	172	+ "paramType": "required",
	173	+ "shape": [
	174	+ -2
	175	+ ],
	176	+ "format_match_mode": "FormatAgnostic"
	177	+ },
	178	+ {
	179	+ "name": "other",
	180	+ "index": 1,
	181	+ "dtype": "float16",
	182	+ "format": "ND",
	183	+ "paramType": "required",
	184	+ "shape": [
	185	+ -2
	186	+ ],
	187	+ "format_match_mode": "FormatAgnostic"
	188	+ }
	189	+ ],
	190	+ "outputs": [
	191	+ {
	192	+ "name": "out",
	193	+ "index": 0,
	194	+ "dtype": "float32",
	195	+ "format": "ND",
	196	+ "paramType": "required",
	197	+ "shape": [
	198	+ -2
	199	+ ],
	200	+ "format_match_mode": "FormatAgnostic"
	201	+ }
	202	+ ]
	203	+ },
	204	+ {
	205	+ "bin_filename": "Hypot_0000000000000000000000000000mix_f16_f32",
	206	+ "inputs": [
	207	+ {
	208	+ "name": "self",
	209	+ "index": 0,
	210	+ "dtype": "float16",
	211	+ "format": "ND",
	212	+ "paramType": "required",
	213	+ "shape": [
	214	+ -2
	215	+ ],
	216	+ "format_match_mode": "FormatAgnostic"
	217	+ },
	218	+ {
	219	+ "name": "other",
	220	+ "index": 1,
	221	+ "dtype": "float32",
	222	+ "format": "ND",
	223	+ "paramType": "required",
	224	+ "shape": [
	225	+ -2
	226	+ ],
	227	+ "format_match_mode": "FormatAgnostic"
	228	+ }
	229	+ ],
	230	+ "outputs": [
	231	+ {
	232	+ "name": "out",
	233	+ "index": 0,
	234	+ "dtype": "float32",
	235	+ "format": "ND",
	236	+ "paramType": "required",
	237	+ "shape": [
	238	+ -2
	239	+ ],
	240	+ "format_match_mode": "FormatAgnostic"
	241	+ }
	242	+ ]
	243	+ },
	244	+ {
	245	+ "bin_filename": "Hypot_0000000000000000000000000000mix_f32_f16",
	246	+ "inputs": [
	247	+ {
	248	+ "name": "self",
	249	+ "index": 0,
	250	+ "dtype": "float32",
	251	+ "format": "ND",
	252	+ "paramType": "required",
	253	+ "shape": [
	254	+ -2
	255	+ ],
	256	+ "format_match_mode": "FormatAgnostic"
	257	+ },
	258	+ {
	259	+ "name": "other",
	260	+ "index": 1,
	261	+ "dtype": "float16",
	262	+ "format": "ND",
	263	+ "paramType": "required",
	264	+ "shape": [
	265	+ -2
	266	+ ],
	267	+ "format_match_mode": "FormatAgnostic"
	268	+ }
	269	+ ],
	270	+ "outputs": [
	271	+ {
	272	+ "name": "out",
	273	+ "index": 0,
	274	+ "dtype": "float32",
	275	+ "format": "ND",
	276	+ "paramType": "required",
	277	+ "shape": [
	278	+ -2
	279	+ ],
	280	+ "format_match_mode": "FormatAgnostic"
	281	+ }
	282	+ ]
	283	+ },
	284	+ {
	285	+ "bin_filename": "Hypot_0000000000000000000000000000mix_bf16_f32",
	286	+ "inputs": [
	287	+ {
	288	+ "name": "self",
	289	+ "index": 0,
	290	+ "dtype": "bfloat16",
	291	+ "format": "ND",
	292	+ "paramType": "required",
	293	+ "shape": [
	294	+ -2
	295	+ ],
	296	+ "format_match_mode": "FormatAgnostic"
	297	+ },
	298	+ {
	299	+ "name": "other",
	300	+ "index": 1,
	301	+ "dtype": "float32",
	302	+ "format": "ND",
	303	+ "paramType": "required",
	304	+ "shape": [
	305	+ -2
	306	+ ],
	307	+ "format_match_mode": "FormatAgnostic"
	308	+ }
	309	+ ],
	310	+ "outputs": [
	311	+ {
	312	+ "name": "out",
	313	+ "index": 0,
	314	+ "dtype": "float32",
	315	+ "format": "ND",
	316	+ "paramType": "required",
	317	+ "shape": [
	318	+ -2
	319	+ ],
	320	+ "format_match_mode": "FormatAgnostic"
	321	+ }
	322	+ ]
	323	+ },
	324	+ {
	325	+ "bin_filename": "Hypot_0000000000000000000000000000mix_f32_bf16",
	326	+ "inputs": [
	327	+ {
	328	+ "name": "self",
	329	+ "index": 0,
	330	+ "dtype": "float32",
	331	+ "format": "ND",
	332	+ "paramType": "required",
	333	+ "shape": [
	334	+ -2
	335	+ ],
	336	+ "format_match_mode": "FormatAgnostic"
	337	+ },
	338	+ {
	339	+ "name": "other",
	340	+ "index": 1,
	341	+ "dtype": "bfloat16",
	342	+ "format": "ND",
	343	+ "paramType": "required",
	344	+ "shape": [
	345	+ -2
	346	+ ],
	347	+ "format_match_mode": "FormatAgnostic"
	348	+ }
	349	+ ],
	350	+ "outputs": [
	351	+ {
	352	+ "name": "out",
	353	+ "index": 0,
	354	+ "dtype": "float32",
	355	+ "format": "ND",
	356	+ "paramType": "required",
	357	+ "shape": [
	358	+ -2
	359	+ ],
	360	+ "format_match_mode": "FormatAgnostic"
	361	+ }
	362	+ ]
	363	+ }
	364	+ ]
	365	+}

Amath/hypot/op_host/config/ascend910b/hypot_simplified_key.ini+2-0

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

Amath/hypot/op_host/hypot_def.cpp+56-0

@@ -0,0 +1,56 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file hypot_def.cpp
	13	+ * \brief hypot operator definition
	14	+ */
	15	+#include <vector>
	16	+#include "register/op_def_registry.h"
	17	+
	18	+namespace ops {
	19	+
	20	+// dtype combinations: 3 same-dtype + 6 mixed (mixed always promotes to FP32 output; the
	21	+// kernel reads DTYPE_SELF/DTYPE_OTHER/DTYPE_OUT per combination, so no TilingKey dtype encoding)
	22	+static const std::vector<ge::DataType> selfDataTypes = {
	23	+ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16,
	24	+ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT,
	25	+};
	26	+
	27	+static const std::vector<ge::DataType> otherDataTypes = {
	28	+ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_BF16, ge::DT_FLOAT16,
	29	+ ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16,
	30	+};
	31	+
	32	+static const std::vector<ge::DataType> outDataTypes = {
	33	+ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT, ge::DT_FLOAT,
	34	+ ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
	35	+};
	36	+
	37	+static const std::vector<ge::Format> formats = {
	38	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
	39	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
	40	+};
	41	+
	42	+class Hypot : public OpDef {
	43	+public:
	44	+ explicit Hypot(const char* name) : OpDef(name)
	45	+ {
	46	+ this->Input("self").ParamType(REQUIRED).DataType(selfDataTypes).Format(formats).UnknownShapeFormat(formats);
	47	+ this->Input("other").ParamType(REQUIRED).DataType(otherDataTypes).Format(formats).UnknownShapeFormat(formats);
	48	+ this->Output("out").ParamType(REQUIRED).DataType(outDataTypes).Format(formats).UnknownShapeFormat(formats);
	49	+
	50	+ this->AICore().AddConfig("ascend910b");
	51	+ this->AICore().AddConfig("ascend910_93");
	52	+ }
	53	+};
	54	+
	55	+OP_ADD(Hypot);
	56	+} // namespace ops

Amath/hypot/op_host/hypot_infershape.cpp+45-0

@@ -0,0 +1,45 @@
	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 hypot_infershape.cpp
	13	+ * \brief hypot infershape: broadcast infershape
	14	+ */
	15	+
	16	+#include "infershape_broadcast_util.h"
	17	+#include "register/op_impl_registry.h"
	18	+#include "log/log.h"
	19	+
	20	+using namespace ge;
	21	+namespace ops {
	22	+
	23	+static ge::graphStatus InferShape4Hypot(gert::InferShapeContext* context)
	24	+{
	25	+ OP_LOGI(context, "Begin InferShape4Hypot");
	26	+ return Ops::Base::InferShape4Broadcast(context);
	27	+}
	28	+
	29	+static ge::graphStatus InferDataType4Hypot(gert::InferDataTypeContext* context)
	30	+{
	31	+ const ge::DataType x1Type = context->GetInputDataType(0);
	32	+ const ge::DataType x2Type = context->GetInputDataType(1);
	33	+ ge::DataType outType = x1Type;
	34	+ if (x1Type == ge::DT_FLOAT \|\| x2Type == ge::DT_FLOAT) {
	35	+ outType = ge::DT_FLOAT;
	36	+ } else if (x1Type != x2Type) {
	37	+ outType = ge::DT_FLOAT;
	38	+ }
	39	+ context->SetOutputDataType(0, outType);
	40	+ return ge::GRAPH_SUCCESS;
	41	+}
	42	+
	43	+IMPL_OP_INFERSHAPE(Hypot).InferShape(InferShape4Hypot).InferDataType(InferDataType4Hypot);
	44	+
	45	+} // namespace ops

Amath/hypot/op_kernel/hypot.cpp+34-0

@@ -0,0 +1,34 @@
	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 hypot.cpp
	13	+ * \brief hypot kernel entry (TilingKey template dispatch)
	14	+ */
	15	+
	16	+#include "hypot_struct.h"
	17	+#include "hypot_tiling_data.h"
	18	+
	19	+using namespace NsHypot;
	20	+
	21	+// Kernel function. The TilingKey carries only the broadcast scene; the concrete tensor dtypes are
	22	+// stamped per OpDef dtype combination by the opc framework and read through DTYPE_SELF/DTYPE_OTHER/DTYPE_OUT.
	23	+template <int BROADCAST_SCENE>
	24	+__global__ __aicore__ void hypot(GM_ADDR self, GM_ADDR other, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling)
	25	+{
	26	+ if (workspace == nullptr) {
	27	+ return;
	28	+ }
	29	+
	30	+ REGISTER_TILING_DEFAULT(HypotTilingData);
	31	+ GET_TILING_DATA_WITH_STRUCT(HypotTilingData, tilingData, tiling);
	32	+
	33	+ HypotKernelImpl<DTYPE_SELF, DTYPE_OTHER, DTYPE_OUT, BROADCAST_SCENE>(self, other, out, &tilingData);
	34	+}

Amath/hypot/op_kernel/hypot_struct.h+3564-0

Amath/hypot/op_kernel/hypot_tiling_data.h+106-0

@@ -0,0 +1,106 @@
	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 hypot_tiling_data.h
	13	+ * \brief hypot tiling template declaration and tiling data struct (shared by host tiling and kernel)
	14	+ */
	15	+
	16	+#ifndef HYPOT_TILING_DATA_H
	17	+#define HYPOT_TILING_DATA_H
	18	+
	19	+#include "ascendc/host_api/tiling/template_argument.h"
	20	+
	21	+#include <cstdint>
	22	+
	23	+namespace NsHypot {
	24	+
	25	+// TilingKey encodes only the broadcast scene; dtypes come from the per-binary
	26	+// DTYPE_SELF/DTYPE_OTHER/DTYPE_OUT aliases.
	27	+
	28	+// broadcast scene ids (the only TilingKey dimension)
	29	+#define HYPOT_SCENE_ONE_DIM 0U
	30	+#define HYPOT_SCENE_MULTI_DIM_BRC 1U
	31	+// MultiDimBroadcast plus structured small-shape extensions, isolated so the default BRC binary stays unchanged.
	32	+#define HYPOT_SCENE_MULTI_DIM_STRUCT 2U
	33	+// Large (>= 1M elems) same-dtype FP16/BF16 2D single-sided broadcast: the expansion engine buffers
	34	+// are not allocated, so the tile is larger. FP16/BF16 only.
	35	+#define HYPOT_SCENE_MULTI_DIM_STREAM 3U
	36	+// Small (<= 1024 elems) multi-dim broadcast: single core, whole output in UB, one Gather per input.
	37	+#define HYPOT_SCENE_MULTI_DIM_VECIDX 4U
	38	+
	39	+// per-input broadcast expansion mode (MultiDimBroadcast scene)
	40	+#define HYPOT_BRC_NONE 0U // no broadcast axis in tile, plain copy
	41	+#define HYPOT_BRC_ROW 1U // single broadcast group, trailing full product I > 1: [G,B,I] row expansion
	42	+#define HYPOT_BRC_COL 2U // single broadcast group, I == 1 (innermost broadcast): splat expansion
	43	+#define HYPOT_BRC_GENERIC 3U // multi-group broadcast: scalar index-mapping fallback
	44	+
	45	+// scalar input bitmap bits of HypotTilingData.scalarFlag: bit0 = x1 is scalar, bit1 = x2 is scalar
	46	+constexpr uint32_t HYPOT_SCALAR_FLAG_X1 = 0x1U;
	47	+constexpr uint32_t HYPOT_SCALAR_FLAG_X2 = 0x2U;
	48	+
	49	+ASCENDC_TPL_ARGS_DECL(Hypot,
	50	+ ASCENDC_TPL_UINT_DECL(BROADCAST_SCENE, ASCENDC_TPL_8_BW, ASCENDC_TPL_UI_LIST, HYPOT_SCENE_ONE_DIM,
	51	+ HYPOT_SCENE_MULTI_DIM_BRC, HYPOT_SCENE_MULTI_DIM_STRUCT,
	52	+ HYPOT_SCENE_MULTI_DIM_STREAM, HYPOT_SCENE_MULTI_DIM_VECIDX), );
	53	+
	54	+ASCENDC_TPL_SEL(
	55	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(BROADCAST_SCENE, ASCENDC_TPL_UI_RANGE, 1, HYPOT_SCENE_ONE_DIM,
	56	+ HYPOT_SCENE_ONE_DIM), ),
	57	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(BROADCAST_SCENE, ASCENDC_TPL_UI_RANGE, 1, HYPOT_SCENE_MULTI_DIM_BRC,
	58	+ HYPOT_SCENE_MULTI_DIM_BRC), ),
	59	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(BROADCAST_SCENE, ASCENDC_TPL_UI_RANGE, 1, HYPOT_SCENE_MULTI_DIM_STRUCT,
	60	+ HYPOT_SCENE_MULTI_DIM_STRUCT), ),
	61	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(BROADCAST_SCENE, ASCENDC_TPL_UI_RANGE, 1, HYPOT_SCENE_MULTI_DIM_STREAM,
	62	+ HYPOT_SCENE_MULTI_DIM_STREAM), ),
	63	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(BROADCAST_SCENE, ASCENDC_TPL_UI_RANGE, 1, HYPOT_SCENE_MULTI_DIM_VECIDX,
	64	+ HYPOT_SCENE_MULTI_DIM_VECIDX), ), );
	65	+
	66	+struct HypotTilingData {
	67	+ // ---- common: multi-core split (unit: tile) and UB split ----
	68	+ uint64_t totalLength = 0; // broadcast output total element num
	69	+ uint32_t usedCoreNum = 0; // actual used AIV core num (blockDim)
	70	+ uint32_t blockFormer = 0; // tiles per non-last core (unit: tile)
	71	+ uint32_t blockTail = 0; // tiles of last core (unit: tile)
	72	+ uint32_t ubFormer = 0; // OneDim: elements per tile; MultiDim: rows per tile on ubSplitAxis
	73	+ uint32_t ubTail = 0; // OneDim: tail tile elements; MultiDim: tail tile rows on ubSplitAxis
	74	+ uint32_t ubOuter = 0; // tile count on ubSplitAxis = ceil(dims[ubSplitAxis]/ubFormer)
	75	+ uint32_t tileNum = 0; // reserved (kernel derives per-core tile num from blockFormer/blockTail)
	76	+ uint32_t scalarFlag = 0; // scalar input bitmap: bit0=x1, bit1=x2 (fused input numel == 1)
	77	+ uint32_t x1PreloadLen = 0; // small x1 source staged in UB once at Init (0 = disabled)
	78	+ uint32_t x2PreloadLen = 0; // small x2 source staged in UB once at Init (0 = disabled)
	79	+ uint32_t nativeHalfFlag = 0; // FP16 native vector chain is allowed for this launch
	80	+ uint32_t x1Numel = 0; // compact input element count (mixed small-kernel path)
	81	+ uint32_t x2Numel = 0;
	82	+
	83	+ // ---- MultiDimBroadcast only (0 / ignored in OneDim) ----
	84	+ uint32_t dimNum = 0; // fused dim num (<= 8)
	85	+ uint32_t outShape[8] = {0}; // fused output dims
	86	+ uint32_t x1Strides[8] = {0}; // x1 strides in elements (broadcast axis stride = 0)
	87	+ uint32_t x2Strides[8] = {0}; // x2 strides in elements (broadcast axis stride = 0)
	88	+ uint32_t outStrides[8] = {0}; // output strides in elements
	89	+ uint32_t ubSplitAxis = 0; // UB split axis
	90	+ uint32_t dimProduct = 0; // product of dims[0..ubSplitAxis-1]
	91	+ uint32_t x1TileLen = 0; // x1 source elements of a full tile (unbroadcast)
	92	+ uint32_t x2TileLen = 0; // x2 source elements of a full tile (unbroadcast)
	93	+ uint32_t innerLen = 0; // C = product of dims[ubSplitAxis+1..dimNum-1]
	94	+ uint32_t rowPitch = 0; // ceil_align(innerLen, 8), padded row pitch in elements (32B)
	95	+ uint32_t x1BrcMode = 0; // HYPOT_BRC_*
	96	+ uint32_t x2BrcMode = 0; // HYPOT_BRC_*
	97	+ uint32_t x1LeadInner = 0; // L: product of full inner axes before the broadcast run (after ubSplitAxis)
	98	+ uint32_t x2LeadInner = 0;
	99	+ uint32_t x1BrcInner = 0; // B_inner: broadcast extent inside the inner region
	100	+ uint32_t x2BrcInner = 0;
	101	+ uint32_t x1SrcInner = 0; // I: trailing full product after the broadcast run
	102	+ uint32_t x2SrcInner = 0;
	103	+};
	104	+
	105	+} // namespace NsHypot
	106	+#endif // HYPOT_TILING_DATA_H

Amath/hypot/tests/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	12	+foreach(SUB_DIR ${CURRENT_DIRS})
	13	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	14	+ add_subdirectory(${SUB_DIR})
	15	+ endif()
	16	+endforeach()

Amath/hypot/tests/st/aclnnHypot/atk_aclnnHypot.json+10068-0

Amath/hypot/tests/st/aclnnHypot/executor_aclnnHypot.py+53-0

@@ -0,0 +1,53 @@
	1	+#!/usr/bin/env python3
	2	+# -- coding: utf-8 --
	3	+# ----------------------------------------------------------------------------
	4	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	5	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# ----------------------------------------------------------------------------
	12	+"""ATK adapter for the aclnnHypot operator (torch.hypot).
	13	+
	14	+Registers two entry points referenced by ``atk_aclnnHypot.json``:
	15	+
	16	+- ``function_npu_hypot`` (BaseApi): CPU golden. ``torch.hypot`` is evaluated in
	17	+ float64 (no intermediate fp16/bf16 rounding) and rounded once to the
	18	+ type-promoted dtype, which is what a correct NPU kernel produces as well.
	19	+- ``aclnn_hypot`` (AclnnBaseApi): dispatches the two tensor inputs to the
	20	+ ``aclnnHypot`` binary op (ND format).
	21	+"""
	22	+
	23	+import torch
	24	+from atk.configs.dataset_config import InputDataset
	25	+from atk.tasks.api_execute import register
	26	+from atk.tasks.api_execute.base_api import BaseApi
	27	+from atk.tasks.api_execute.aclnn_base_api import AclnnBaseApi
	28	+from atk.tasks.backends.lib_interface.acl_wrapper import AclFormat
	29	+
	30	+
	31	+@register("function_npu_hypot")
	32	+class FunctionHypotApi(BaseApi):
	33	+ """CPU golden (benchmark) implementation."""
	34	+
	35	+ def __call__(self, input_data: InputDataset, with_output: bool = False):
	36	+ output = None
	37	+ if self.device == "cpu":
	38	+ x = input_data.kwargs["x"]
	39	+ y = input_data.kwargs["y"]
	40	+ out_dtype = torch.promote_types(x.dtype, y.dtype)
	41	+ output = torch.hypot(
	42	+ x.detach().to("cpu").to(torch.float64),
	43	+ y.detach().to("cpu").to(torch.float64),
	44	+ ).to(out_dtype)
	45	+ return output
	46	+
	47	+
	48	+@register("aclnn_hypot")
	49	+class AclnnHypotApi(AclnnBaseApi):
	50	+ """aclnnHypot binary-op wrapper (x1, x2 -> y), ND format only."""
	51	+
	52	+ def get_format(self, input_data: InputDataset, index=None, name=None):
	53	+ return AclFormat.ACL_FORMAT_ND

Amath/hypot/tests/ut/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	12	+foreach(SUB_DIR ${CURRENT_DIRS})
	13	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	14	+ add_subdirectory(${SUB_DIR})
	15	+ endif()
	16	+endforeach()

Amath/hypot/tests/ut/op_api/test_aclnn_hypot.cpp+266-0

@@ -0,0 +1,266 @@
	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 <vector>
	12	+#include "gtest/gtest.h"
	13	+
	14	+#include "opdev/platform.h"
	15	+#include "../../../op_api/aclnn_hypot.h"
	16	+#include "op_api_ut_common/tensor_desc.h"
	17	+#include "op_api_ut_common/op_api_ut.h"
	18	+
	19	+using namespace std;
	20	+
	21	+class l2_hypot_test : public testing::Test {
	22	+protected:
	23	+ static void SetUpTestCase() { cout << "l2_hypot_test SetUp" << endl; }
	24	+
	25	+ static void TearDownTestCase() { cout << "l2_hypot_test TearDown" << endl; }
	26	+};
	27	+
	28	+TEST_F(l2_hypot_test, case_01_float)
	29	+{
	30	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	31	+
	32	+ auto x1Desc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND).Value(vector<float>{3.0, 5.0, 8.0, 7.0});
	33	+ auto x2Desc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND).Value(vector<float>{4.0, 12.0, 6.0, 24.0});
	34	+ auto outDesc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
	35	+
	36	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	37	+
	38	+ uint64_t workspaceSize = 0;
	39	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	40	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	41	+}
	42	+
	43	+TEST_F(l2_hypot_test, case_02_float16)
	44	+{
	45	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	46	+
	47	+ auto x1Desc = TensorDesc({2, 2}, ACL_FLOAT16, ACL_FORMAT_ND).Value(vector<float>{3.0, 5.0, 8.0, 7.0});
	48	+ auto x2Desc = TensorDesc({2, 2}, ACL_FLOAT16, ACL_FORMAT_ND).Value(vector<float>{4.0, 12.0, 6.0, 24.0});
	49	+ auto outDesc = TensorDesc({2, 2}, ACL_FLOAT16, ACL_FORMAT_ND);
	50	+
	51	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	52	+
	53	+ uint64_t workspaceSize = 0;
	54	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	55	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	56	+}
	57	+
	58	+TEST_F(l2_hypot_test, case_03_bfloat16)
	59	+{
	60	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	61	+
	62	+ auto x1Desc = TensorDesc({2, 2}, ACL_BF16, ACL_FORMAT_ND).Value(vector<float>{3.0, 5.0, 8.0, 7.0});
	63	+ auto x2Desc = TensorDesc({2, 2}, ACL_BF16, ACL_FORMAT_ND).Value(vector<float>{4.0, 12.0, 6.0, 24.0});
	64	+ auto outDesc = TensorDesc({2, 2}, ACL_BF16, ACL_FORMAT_ND);
	65	+
	66	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	67	+
	68	+ uint64_t workspaceSize = 0;
	69	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	70	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	71	+}
	72	+
	73	+TEST_F(l2_hypot_test, case_04_mixed_float_float16)
	74	+{
	75	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	76	+
	77	+ auto x1Desc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND).Value(vector<float>{3.0, 5.0, 8.0, 7.0});
	78	+ auto x2Desc = TensorDesc({2, 2}, ACL_FLOAT16, ACL_FORMAT_ND).Value(vector<float>{4.0, 12.0, 6.0, 24.0});
	79	+ auto outDesc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
	80	+
	81	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	82	+
	83	+ uint64_t workspaceSize = 0;
	84	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	85	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	86	+}
	87	+
	88	+TEST_F(l2_hypot_test, case_05_mixed_float16_bfloat16)
	89	+{
	90	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	91	+
	92	+ auto x1Desc = TensorDesc({2, 2}, ACL_FLOAT16, ACL_FORMAT_ND).Value(vector<float>{3.0, 5.0, 8.0, 7.0});
	93	+ auto x2Desc = TensorDesc({2, 2}, ACL_BF16, ACL_FORMAT_ND).Value(vector<float>{4.0, 12.0, 6.0, 24.0});
	94	+ auto outDesc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
	95	+
	96	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	97	+
	98	+ uint64_t workspaceSize = 0;
	99	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	100	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	101	+}
	102	+
	103	+TEST_F(l2_hypot_test, case_06_broadcast_right_pad)
	104	+{
	105	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	106	+
	107	+ auto x1Desc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-2.0, 2.0);
	108	+ auto x2Desc = TensorDesc({3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-2.0, 2.0);
	109	+ auto outDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	110	+
	111	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	112	+
	113	+ uint64_t workspaceSize = 0;
	114	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	115	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	116	+}
	117	+
	118	+TEST_F(l2_hypot_test, case_07_broadcast_cross)
	119	+{
	120	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	121	+
	122	+ auto x1Desc = TensorDesc({1, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-2.0, 2.0);
	123	+ auto x2Desc = TensorDesc({3, 1}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-2.0, 2.0);
	124	+ auto outDesc = TensorDesc({3, 4}, ACL_FLOAT, ACL_FORMAT_ND);
	125	+
	126	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	127	+
	128	+ uint64_t workspaceSize = 0;
	129	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	130	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	131	+}
	132	+
	133	+TEST_F(l2_hypot_test, case_08_empty_tensor)
	134	+{
	135	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	136	+
	137	+ auto x1Desc = TensorDesc({0}, ACL_FLOAT, ACL_FORMAT_ND);
	138	+ auto x2Desc = TensorDesc({0}, ACL_FLOAT, ACL_FORMAT_ND);
	139	+ auto outDesc = TensorDesc({0}, ACL_FLOAT, ACL_FORMAT_ND);
	140	+
	141	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	142	+
	143	+ uint64_t workspaceSize = 0;
	144	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	145	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	146	+}
	147	+
	148	+TEST_F(l2_hypot_test, case_09_not_contiguous)
	149	+{
	150	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	151	+
	152	+ auto x1Desc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND, {1, 2}, 0, {2, 2}).Value(vector<float>{1, 2, 3, 4});
	153	+ auto x2Desc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND, {1, 2}, 0, {2, 2}).Value(vector<float>{1, 2, 3, 4});
	154	+ auto outDesc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
	155	+
	156	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	157	+
	158	+ uint64_t workspaceSize = 0;
	159	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	160	+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
	161	+}
	162	+
	163	+TEST_F(l2_hypot_test, case_10_out_shape_mismatch)
	164	+{
	165	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	166	+
	167	+ auto x1Desc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	168	+ auto x2Desc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	169	+ auto outDesc = TensorDesc({3, 2}, ACL_FLOAT, ACL_FORMAT_ND);
	170	+
	171	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	172	+
	173	+ uint64_t workspaceSize = 0;
	174	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	175	+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
	176	+}
	177	+
	178	+TEST_F(l2_hypot_test, case_11_broadcast_incompatible)
	179	+{
	180	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	181	+
	182	+ auto x1Desc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	183	+ auto x2Desc = TensorDesc({4}, ACL_FLOAT, ACL_FORMAT_ND);
	184	+ auto outDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	185	+
	186	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	187	+
	188	+ uint64_t workspaceSize = 0;
	189	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	190	+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
	191	+}
	192	+
	193	+TEST_F(l2_hypot_test, case_12_x1_null)
	194	+{
	195	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	196	+
	197	+ auto x1Desc = nullptr;
	198	+ auto x2Desc = TensorDesc({3, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	199	+ auto outDesc = TensorDesc({3, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	200	+
	201	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	202	+
	203	+ uint64_t workspaceSize = 0;
	204	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	205	+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
	206	+}
	207	+
	208	+TEST_F(l2_hypot_test, case_13_x2_null)
	209	+{
	210	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	211	+
	212	+ auto x1Desc = TensorDesc({3, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	213	+ auto x2Desc = nullptr;
	214	+ auto outDesc = TensorDesc({3, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	215	+
	216	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	217	+
	218	+ uint64_t workspaceSize = 0;
	219	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	220	+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
	221	+}
	222	+
	223	+TEST_F(l2_hypot_test, case_14_out_null)
	224	+{
	225	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	226	+
	227	+ auto x1Desc = TensorDesc({3, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	228	+ auto x2Desc = TensorDesc({3, 3}, ACL_FLOAT, ACL_FORMAT_ND);
	229	+ auto outDesc = nullptr;
	230	+
	231	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	232	+
	233	+ uint64_t workspaceSize = 0;
	234	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	235	+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
	236	+}
	237	+
	238	+TEST_F(l2_hypot_test, case_15_dim_gt_8)
	239	+{
	240	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	241	+
	242	+ auto x1Desc = TensorDesc({3, 2, 1, 2, 1, 3, 2, 1, 2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
	243	+ auto x2Desc = TensorDesc({3, 2, 1, 2, 1, 3, 2, 1, 2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
	244	+ auto outDesc = TensorDesc({3, 2, 1, 2, 1, 3, 2, 1, 2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
	245	+
	246	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	247	+
	248	+ uint64_t workspaceSize = 0;
	249	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	250	+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
	251	+}
	252	+
	253	+TEST_F(l2_hypot_test, case_16_dtype_unsupported)
	254	+{
	255	+ op::SetPlatformNpuArch(NpuArch::DAV_2201);
	256	+
	257	+ auto x1Desc = TensorDesc({3, 3}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-10, 10);
	258	+ auto x2Desc = TensorDesc({3, 3}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-10, 10);
	259	+ auto outDesc = TensorDesc({3, 3}, ACL_INT32, ACL_FORMAT_ND);
	260	+
	261	+ auto ut = OP_API_UT(aclnnHypot, INPUT(x1Desc, x2Desc), OUTPUT(outDesc));
	262	+
	263	+ uint64_t workspaceSize = 0;
	264	+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
	265	+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
	266	+}

Amath/hypot/tests/ut/op_host/CMakeLists.txt+20-0

@@ -0,0 +1,20 @@
	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	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	12	+foreach(SUB_DIR ${CURRENT_DIRS})
	13	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	14	+ add_subdirectory(${SUB_DIR})
	15	+ endif()
	16	+endforeach()
	17	+
	18	+if(TILING_UT OR PROTO_UT OR OP_API_UT OR (UT_TEST_ALL AND NOT AICPU_ONLY))
	19	+ add_modules_llt_sources(HOSTNAME ${OPTEST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	20	+endif()

Amath/hypot/tests/ut/op_host/arch22/test_hypot_tiling_arch22.cpp+979-0

Amath/hypot/tests/ut/op_host/test_hypot_infershape.cpp+129-0

Amath/hypot/tests/ut/op_kernel/CMakeLists.txt+23-0

Amath/hypot/tests/ut/op_kernel/test_hypot.cpp+1212-0

Amath/hypot/tests/ut/op_kernel/test_hypot_launch.cpp+157-0

Mscripts/kernel/binary_config/ascendc_config.json+2-1