add hard_sigmoid for ascend950-ops-nn-AtomGit

MOAT.xml+1-0

	@@ -23,6 +23,7 @@
23	<!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->	23	<!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->
24	<!--filteritem type="filepath" name="abcdefg/.*.uvwxyz" desc="Describe the reason for filtering scan results"/-->	24	<!--filteritem type="filepath" name="abcdefg/.*.uvwxyz" desc="Describe the reason for filtering scan results"/-->
25	<!--filteritem type="filepath" name="projectroot/[a-zA-Z0-9]{20,}.sh" desc="Temp files"/-->	25	<!--filteritem type="filepath" name="projectroot/[a-zA-Z0-9]{20,}.sh" desc="Temp files"/-->
		26	+ <filteritem type="filename" name="*.csv" desc="skip csv files"/>
		Z Z zl_hw7月24日为何有OAT的修改？ T tianyu52 7月24日评论： 7月24日评论：
26	</filefilter>	27	</filefilter>
27	<filefilter name="defaultPolicyFilter" desc="Filters for compatibility，license header policies">	28	<filefilter name="defaultPolicyFilter" desc="Filters for compatibility，license header policies">
28	<!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->	29	<!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->

Mactivation/hard_sigmoid/README.md+10-9

	@@ -2,14 +2,14 @@
2		2
3	## 产品支持情况	3	## 产品支持情况
4		4
5	-\|产品 \| 是否支持 \|	5	+\| 产品 \| 是否支持 \|
6	-\|:-------------------------\|:----------:\|	6	+\| :----------------------------------------- \| :------:\|
7	-\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|	7	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
8	-\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|	8	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|
9	-\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|	9	+\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|
10	-\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|	10	+\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|
11	-\| <term>Atlas 推理系列产品</term> \| × \|	11	+\| <term>Atlas 推理系列产品</term> \| × \|
12	-\| <term>Atlas 训练系列产品</term> \| √ \|	12	+\| <term>Atlas 训练系列产品</term> \| √ \|
13		13
14	## 功能说明	14	## 功能说明
15		15
	@@ -26,7 +26,7 @@
26		26
27	## 参数说明	27	## 参数说明
28		28
29	- <table style="undefined;table-layout: fixed; width: 800px"><colgroup>	29	+ <table style="table-layout: fixed; width: 800px"><colgroup>
30	<col style="width: 110px">	30	<col style="width: 110px">
31	<col style="width: 130px">	31	<col style="width: 130px">
32	<col style="width: 300px">	32	<col style="width: 300px">
	@@ -71,3 +71,4 @@
71	\| ---------------- \|------------------------------------------------------------------------------\|------------------------------------------------------------------------------------\|	71	\| ---------------- \|------------------------------------------------------------------------------\|------------------------------------------------------------------------------------\|
72	\| aclnn接口 \| [test_aclnn_hard_sigmoid.cpp](examples/test_aclnn_hard_sigmoid.cpp) \| 通过[aclnnHardsigmoid](docs/aclnnHardsigmoid&aclnnInplaceHardsigmoid.md)接口方式调用HardSigmoid算子。 \|	72	\| aclnn接口 \| [test_aclnn_hard_sigmoid.cpp](examples/test_aclnn_hard_sigmoid.cpp) \| 通过[aclnnHardsigmoid](docs/aclnnHardsigmoid&aclnnInplaceHardsigmoid.md)接口方式调用HardSigmoid算子。 \|
73	\| aclnn接口 \| [test_aclnn_inplace_hard_sigmoid.cpp](examples/test_aclnn_inplace_hard_sigmoid.cpp) \| 通过[aclnnInplaceHardsigmoid](docs/aclnnHardsigmoid&aclnnInplaceHardsigmoid.md)接口方式调用HardSigmoid算子。 \|	73	\| aclnn接口 \| [test_aclnn_inplace_hard_sigmoid.cpp](examples/test_aclnn_inplace_hard_sigmoid.cpp) \| 通过[aclnnInplaceHardsigmoid](docs/aclnnHardsigmoid&aclnnInplaceHardsigmoid.md)接口方式调用HardSigmoid算子。 \|
		74	+\| 图模式(GE IR) \| [test_geir_hard_sigmoid.cpp](examples/arch35/test_geir_hard_sigmoid.cpp) \| 通过GE IR构图方式调用HardSigmoid算子。 \|

Mactivation/hard_sigmoid/docs/aclnnHardsigmoid&aclnnInplaceHardsigmoid.md+3-3

	@@ -5,7 +5,7 @@
5	## 产品支持情况	5	## 产品支持情况
6		6
7	<!-- npu="950" id1 -->	7	<!-- npu="950" id1 -->
8	-- <term>Ascend 950PR/Ascend 950DT</term>：不支持	8	+- <term>Ascend 950PR/Ascend 950DT</term>：支持
9	<!-- end id1 -->	9	<!-- end id1 -->
10	<!-- npu="A3" id2 -->	10	<!-- npu="A3" id2 -->
11	- <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：支持	11	- <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：支持
	@@ -102,7 +102,7 @@ aclnnStatus aclnnInplaceHardsigmoid(
102	<th>非连续Tensor</th>	102	<th>非连续Tensor</th>
103	</tr></thead>	103	</tr></thead>
104	<tbody>	104	<tbody>
105	- <tr>	105	+ <tr>
106	<td>self（aclTensor*）</td>	106	<td>self（aclTensor*）</td>
107	<td>输入</td>	107	<td>输入</td>
108	<td>表示激活函数的输入，公式中的输入self。</td>	108	<td>表示激活函数的输入，公式中的输入self。</td>
	@@ -258,7 +258,7 @@ aclnnStatus aclnnInplaceHardsigmoid(
258	<th>非连续Tensor</th>	258	<th>非连续Tensor</th>
259	</tr></thead>	259	</tr></thead>
260	<tbody>	260	<tbody>
261	- <tr>	261	+ <tr>
262	<td>self（aclTensor*）</td>	262	<td>self（aclTensor*）</td>
263	<td>输入</td>	263	<td>输入</td>
264	<td>公式中的输入self和输出Hardsigmoid(self)。</td>	264	<td>公式中的输入self和输出Hardsigmoid(self)。</td>

Aactivation/hard_sigmoid/examples/arch35/test_geir_hard_sigmoid.cpp+152-0

@@ -0,0 +1,152 @@
	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 test_geir_hard_sigmoid.cpp
	13	+ * \brief 通过GE IR构图方式调用HardSigmoid算子的样例，运行结果与CPU golden比对。
	14	+ */
	15	+
	16	+#include <cmath>
	17	+#include <cstdio>
	18	+#include <map>
	19	+#include <memory>
	20	+#include <new>
	21	+#include <string>
	22	+#include <vector>
	23	+
	24	+#include "graph.h"
	25	+#include "types.h"
	26	+#include "tensor.h"
	27	+#include "ge_error_codes.h"
	28	+#include "ge_api_types.h"
	29	+#include "ge_api.h"
	30	+#include "array_ops.h"
	31	+#include "../../op_graph/hard_sigmoid_proto.h"
	32	+
	33	+#define FAILED (-1)
	34	+#define SUCCESS 0
	35	+
	36	+using namespace ge;
	37	+using std::vector;
	38	+
	39	+namespace {
	40	+constexpr float kAlpha = 1.0f / 6.0f;
	41	+constexpr float kBeta = 0.5f;
	42	+constexpr float kTolerance = 1e-4f;
	43	+
	44	+// 取值覆盖y=0截断区(x<=-3)、线性区与y=1截断区(x>=3)
	45	+const vector<float> kInputData = {-4.0f, -3.0f, -2.0f, -1.0f, 0.0f, 1.0f, 2.0f, 3.0f};
	46	+const vector<int64_t> kInputShape = {4, 2};
	47	+
	48	+float HardSigmoidGolden(float x) { return std::fmax(0.0f, std::fmin(1.0f, kAlpha * x + kBeta)); }
	49	+
	50	+// 构造Data -> HardSigmoid单节点图
	51	+Status BuildGraph(Graph& graph, vector<ge::Tensor>& inputTensors)
	52	+{
	53	+ TensorDesc xDesc(ge::Shape(kInputShape), FORMAT_ND, DT_FLOAT);
	54	+ xDesc.SetPlacement(ge::kPlacementHost);
	55	+ xDesc.SetRealDimCnt(kInputShape.size());
	56	+
	57	+ auto data = op::Data("x").set_attr_index(0);
	58	+ data.update_input_desc_x(xDesc);
	59	+ data.update_output_desc_y(xDesc);
	60	+
	61	+ auto hardSigmoid = op::HardSigmoid("hard_sigmoid");
	62	+ hardSigmoid.set_attr_alpha(kAlpha);
	63	+ hardSigmoid.set_attr_beta(kBeta);
	64	+ hardSigmoid.set_input_input_x(data);
	65	+
	66	+ graph.AddOp(data);
	67	+ graph.SetInputs({data}).SetOutputs({hardSigmoid});
	68	+
	69	+ inputTensors.emplace_back(xDesc, reinterpret_cast<const uint8_t*>(kInputData.data()),
	70	+ kInputData.size() * sizeof(float));
	71	+ return SUCCESS;
	72	+}
	73	+
	74	+// 与CPU golden比对，返回不符合预期的个数
	75	+int VerifyOutput(const ge::Tensor& output)
	76	+{
	77	+ const auto* result = reinterpret_cast<const float*>(output.GetData());
	78	+ const int64_t count = output.GetTensorDesc().GetShape().GetShapeSize();
	79	+ const size_t expectedBytes = kInputData.size() * sizeof(float);
	80	+ if (result == nullptr \|\| count < 0 \|\| static_cast<size_t>(count) != kInputData.size() \|\|
	81	+ output.GetSize() < expectedBytes) {
	82	+ printf("[ERROR] Invalid output buffer: count=%ld, bytes=%zu, expected_count=%zu, expected_bytes=%zu\n", count,
	83	+ output.GetSize(), kInputData.size(), expectedBytes);
	84	+ return 1;
	85	+ }
	86	+ int failures = 0;
	87	+
	88	+ printf("%-10s %-12s %-12s %s\n", "x", "npu", "golden", "check");
	89	+ for (int64_t i = 0; i < count; i++) {
	90	+ float golden = HardSigmoidGolden(kInputData[i]);
	91	+ bool matched = std::fabs(result[i] - golden) <= kTolerance;
	92	+ printf("%-10.4f %-12.6f %-12.6f %s\n", kInputData[i], result[i], golden, matched ? "OK" : "MISMATCH");
	93	+ if (!matched) {
	94	+ failures++;
	95	+ }
	96	+ }
	97	+ return failures;
	98	+}
	99	+} // namespace
	100	+
	101	+int main()
	102	+{
	103	+ std::map<AscendString, AscendString> globalOptions = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
	104	+ if (ge::GEInitialize(globalOptions) != SUCCESS) {
	105	+ printf("[ERROR] GEInitialize failed: %s\n", ge::GEGetErrorMsgV2().GetString());
	106	+ return FAILED;
	107	+ }
	108	+
	109	+ Graph graph("hard_sigmoid_geir");
	110	+ vector<ge::Tensor> inputTensors;
	111	+ if (BuildGraph(graph, inputTensors) != SUCCESS) {
	112	+ printf("[ERROR] Build graph failed\n");
	113	+ (void)ge::GEFinalize();
	114	+ return FAILED;
	115	+ }
	116	+
	117	+ std::map<AscendString, AscendString> sessionOptions = {};
	118	+ std::unique_ptr<Session> session(new (std::nothrow) Session(sessionOptions));
	119	+ if (session == nullptr) {
	120	+ printf("[ERROR] Create GE session failed\n");
	121	+ (void)ge::GEFinalize();
	122	+ return FAILED;
	123	+ }
	124	+
	125	+ int failures = 0;
	126	+ std::map<AscendString, AscendString> graphOptions = {};
	127	+ const uint32_t graphId = 0;
	128	+ // AddGraph只注册图，算子编译在RunGraph阶段触发
	129	+ if (session->AddGraph(graphId, graph, graphOptions) != SUCCESS) {
	130	+ printf("[ERROR] AddGraph failed: %s\n", ge::GEGetErrorMsgV2().GetString());
	131	+ failures = 1;
	132	+ } else {
	133	+ vector<ge::Tensor> outputTensors;
	134	+ if (session->RunGraph(graphId, inputTensors, outputTensors) != SUCCESS) {
	135	+ printf("[ERROR] RunGraph failed: %s\n", ge::GEGetErrorMsgV2().GetString());
	136	+ failures = 1;
	137	+ } else if (outputTensors.empty()) {
	138	+ printf("[ERROR] RunGraph produced no output\n");
	139	+ failures = 1;
	140	+ } else {
	141	+ failures = VerifyOutput(outputTensors[0]);
	142	+ }
	143	+ }
	144	+
	145	+ session.reset();
	146	+ if (ge::GEFinalize() != SUCCESS) {
	147	+ printf("[WARN] GEFinalize failed: %s\n", ge::GEGetErrorMsgV2().GetString());
	148	+ }
	149	+
	150	+ printf("%s\n", failures == 0 ? "[INFO] test_geir_hard_sigmoid PASSED" : "[ERROR] test_geir_hard_sigmoid FAILED");
	151	+ return failures == 0 ? SUCCESS : FAILED;
	152	+}

Mactivation/hard_sigmoid/examples/test_aclnn_hard_sigmoid.cpp+1-1

	@@ -140,4 +140,4 @@ int main()
140	aclrtResetDevice(deviceId);	140	aclrtResetDevice(deviceId);
141	aclFinalize();	141	aclFinalize();
142	return 0;	142	return 0;
143	-}	143	+}

Mactivation/hard_sigmoid/examples/test_aclnn_inplace_hard_sigmoid.cpp+1-1

	@@ -132,4 +132,4 @@ int main()
132	aclrtResetDevice(deviceId);	132	aclrtResetDevice(deviceId);
133	aclFinalize();	133	aclFinalize();
134	return 0;	134	return 0;
135	-}	135	+}

Aactivation/hard_sigmoid/op_graph/CMakeLists.txt+9-0

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

Aactivation/hard_sigmoid/op_graph/hard_sigmoid_proto.h+46-0

@@ -0,0 +1,46 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file hard_sigmoid_proto.h
	13	+ * \brief HardSigmoid 图模式算子原型定义
	14	+ */
	15	+#ifndef OPS_BUILT_IN_OP_PROTO_INC_HARD_SIGMOID_H_
	16	+#define OPS_BUILT_IN_OP_PROTO_INC_HARD_SIGMOID_H_
	17	+
	18	+#include "graph/operator_reg.h"
	L L liubo757月25日这个算子补齐inferdatetype的注册 T tianyu52 7月25日评论： 7月25日评论：
	19	+
	20	+namespace ge {
	21	+/**
	22	+ * @brief Calculate the hard sigmoid function: output_y = max(0, min(1, alpha * input_x + beta)).
	23	+ *
	24	+ * @par Inputs:
	25	+ * input_x: An ND tensor. The shape should be within the range of 0D to 8D.
	26	+ * Must be one of the following types: float16, float32, int32, bfloat16.
	27	+ *
	28	+ * @par Attributes:
	29	+ * @li alpha: An optional float. Slope of the operator, defaults to 0.16666666.
	30	+ * @li beta: An optional float. Offset of the operator, defaults to 0.5.
	31	+ *
	32	+ * @par Outputs:
	33	+ * output_y: An ND tensor with the same dtype and shape as "input_x".
	34	+ *
	35	+ * @par Third-party framework compatibility
	36	+ * Compatible with the PyTorch operator torch.nn.Hardsigmoid.
	37	+ */
	38	+REG_OP(HardSigmoid)
	39	+ .INPUT(input_x, TensorType({DT_FLOAT, DT_FLOAT16, DT_BF16, DT_INT32}))
	40	+ .OUTPUT(output_y, TensorType({DT_FLOAT, DT_FLOAT16, DT_BF16, DT_INT32}))
	41	+ .ATTR(alpha, Float, 0.16666666)
	42	+ .ATTR(beta, Float, 0.5)
	43	+ .OP_END_FACTORY_REG(HardSigmoid)
	44	+} // namespace ge
	45	+
	46	+#endif // OPS_BUILT_IN_OP_PROTO_INC_HARD_SIGMOID_H_

Mactivation/hard_sigmoid/op_host/CMakeLists.txt+15-3

	@@ -1,10 +1,22 @@
1	# Copyright (c) 2025 Huawei Technologies Co., Ltd.	1	# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2	-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of	2	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3	# CANN Open Software License Agreement Version 2.0 (the "License").	3	# CANN Open Software License Agreement Version 2.0 (the "License").
4	# Please refer to the License for details. You may not use this file except in compliance with the License.	4	# Please refer to the License for details. You may not use this file except in compliance with the License.
5	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	5	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6	# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	6	# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7	# See LICENSE in the root of the software repository for the full text of the License.	7	# See LICENSE in the root of the software repository for the full text of the License.
8	#/	8	#/
9		9
10	-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE hard_sigmoid ACLNNTYPE aclnn)	10	+set(SUPPORT_COMPUTE_UNIT "ascend950")
		11	+set(SUPPORT_TILING_DIR "arch35")
		12	+
		13	+add_modules_sources(
		14	+ HOSTNAME ${OPHOST_NAME}
		15	+ MODE PRIVATE
		16	+ DIR ${CMAKE_CURRENT_SOURCE_DIR}
		17	+ OPTYPE hard_sigmoid
		18	+ ACLNNTYPE aclnn
		19	+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
		20	+ TILING_DIR ${SUPPORT_TILING_DIR}
		21	+ DISABLE_IN_OPP TRUE
		22	+)

Aactivation/hard_sigmoid/op_host/arch35/hard_sigmoid_tiling_arch35.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	+
	11	+/*!
	12	+ * \file hard_sigmoid_tiling_arch35.cpp
	13	+ * \brief HardSigmoid tiling（arch35 / DAV_3510）
	14	+ *
	15	+ * 多核切分和核内 UB 切分分别由 blockFactor、ubFactor 描述，并按 input0 dtype 选择模板实例。
	16	+ */
	17	+
	18	+#include "register/op_def_registry.h"
	19	+#include <string>
	20	+#include "op_common/log/log.h"
	21	+#include "op_common/op_host/util/math_util.h"
	22	+#include "op_common/op_host/util/platform_util.h"
	23	+#include "../../op_kernel/arch35/hard_sigmoid_tiling_data.h"
	24	+#include "../../op_kernel/arch35/hard_sigmoid_tiling_key.h"
	25	+
	26	+namespace optiling {
	27	+namespace {
	28	+using Ops::Base::CeilDiv;
	29	+using Ops::Base::FloorAlign;
	30	+using Ops::Base::FloorDiv;
	31	+using Ops::Base::GetUbBlockSize;
	32	+using Ops::Base::GetVRegSize;
	33	+
	34	+constexpr float DEFAULT_ALPHA = 1.0f / 6.0f;
	35	+constexpr float DEFAULT_BETA = 0.5f;
	36	+constexpr size_t WORKSPACE_NUM = 1;
	37	+constexpr size_t WS_SYS_SIZE = 0U;
	38	+
	39	+constexpr int64_t UB_RESERVE_BYTES = 8192; // 预留系统/对齐余量
	40	+constexpr int64_t MIN_COPY_BYTES = 16 * 1024;
	41	+constexpr int64_t F32_TEMP_BYTES = static_cast<int64_t>(sizeof(float)); // 非 fp32 路径的 fp32 中间缓冲
	42	+
	43	+struct HardSigmoidCompileInfo {
	44	+ uint64_t coreNum = 0;
	45	+ uint64_t ubSize = 0;
	46	+};
	47	+
	48	+int64_t StorageBytesPerElement(ge::DataType dtype) { return (dtype == ge::DT_FLOAT16 \|\| dtype == ge::DT_BF16) ? 2 : 4; }
	49	+
	50	+// 每元素 UB 占用，与 kernel 侧 buffer 布局严格对应（见 hard_sigmoid.cpp InitBuffer）：
	51	+// inQue/outQue 各 HARD_SIGMOID_BUFFER_NUM 份 sizeof(T)；非 fp32 另加 1 份 fp32 中间缓冲。
	52	+// 实测值：FLOAT=16、FLOAT16/BFLOAT16=12、INT32=20。
	53	+// BUFFER_NUM 取自 Host/Kernel 共享头，避免任一侧改动导致 UB 预算静默失配。
	54	+int64_t UbBytesPerElement(ge::DataType dtype)
	55	+{
	56	+ constexpr int64_t QUEUE_COUNT = 2; // inQue + outQue
	57	+ const int64_t queueBytes = HARD_SIGMOID_BUFFER_NUM * QUEUE_COUNT * StorageBytesPerElement(dtype);
	58	+ return queueBytes + ((dtype == ge::DT_FLOAT) ? 0 : F32_TEMP_BYTES);
	59	+}
	60	+
	61	+// vRegSize 为向量寄存器宽度（平台参数，由调用方经 GetVRegSize 获取），用于对齐 ubFactor。
	62	+int64_t ComputeUbFactor(ge::DataType dtype, int64_t ubSize, int64_t ubBlockSize, int64_t vRegSize)
	63	+{
	64	+ const int64_t usableUbSize = ubSize - UB_RESERVE_BYTES;
	65	+ if (usableUbSize <= 0 \|\| ubBlockSize <= 0 \|\| vRegSize <= 0) {
	66	+ return 0;
	67	+ }
	68	+
	69	+ const int64_t alignBytes = ubBlockSize > vRegSize ? ubBlockSize : vRegSize;
	70	+ const int64_t alignElements = alignBytes / StorageBytesPerElement(dtype);
	71	+ return FloorAlign(FloorDiv(usableUbSize, UbBytesPerElement(dtype)), alignElements);
	72	+}
	73	+
	74	+ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
	75	+{
	76	+ size_t* workspaces = context->GetWorkspaceSizes(WORKSPACE_NUM);
	77	+ OP_CHECK_NULL_WITH_CONTEXT(context, workspaces);
	78	+ workspaces[0] = WS_SYS_SIZE;
	79	+ return ge::GRAPH_SUCCESS;
	80	+}
	81	+} // namespace
	82	+
	83	+static ge::graphStatus HardSigmoidTilingFunc(gert::TilingContext* context)
	84	+{
	85	+ auto* inputShape = context->GetInputShape(0);
	86	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
	87	+
	88	+ auto* inputDesc = context->GetInputDesc(0);
	89	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	90	+ const ge::DataType dtype = inputDesc->GetDataType();
	91	+ OP_CHECK_IF(dtype != ge::DT_FLOAT && dtype != ge::DT_FLOAT16 && dtype != ge::DT_BF16 && dtype != ge::DT_INT32,
	92	+ OP_LOGE_FOR_INVALID_DTYPE(context->GetNodeName(), "input_x",
	93	+ std::to_string(static_cast<int32_t>(dtype)).c_str(),
	94	+ "DT_FLOAT, DT_FLOAT16, DT_BF16, DT_INT32"),
	95	+ return ge::GRAPH_FAILED);
	96	+
	97	+ auto* tilingData = context->GetTilingData<HardSigmoidTilingData>();
	98	+ OP_CHECK_NULL_WITH_CONTEXT(context, tilingData);
	99	+ *tilingData = HardSigmoidTilingData{};
	100	+
	101	+ // 激活系数：使用默认值，attrs 存在时按位置 0/1 覆盖。
	102	+ tilingData->alpha = DEFAULT_ALPHA;
	103	+ tilingData->beta = DEFAULT_BETA;
	104	+ const auto* attrs = context->GetAttrs();
	105	+ if (attrs != nullptr) {
	106	+ const float* alpha = attrs->GetFloat(0);
	107	+ const float* beta = attrs->GetFloat(1);
	108	+ if (alpha != nullptr) {
	109	+ tilingData->alpha = *alpha;
	110	+ }
	111	+ if (beta != nullptr) {
	112	+ tilingData->beta = *beta;
	113	+ }
	114	+ }
	115	+
	116	+ if (GetWorkspaceSize(context) != ge::GRAPH_SUCCESS) {
	117	+ OP_LOGE_WITHOUT_REPORT(context->GetNodeName(), "GetWorkspaceSize failed");
	118	+ return ge::GRAPH_FAILED;
	119	+ }
	120	+
	121	+ const auto& storageShape = inputShape->GetStorageShape();
	122	+ for (size_t i = 0; i < storageShape.GetDimNum(); ++i) {
	123	+ OP_CHECK_IF(storageShape.GetDim(i) < 0,
	124	+ OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(context->GetNodeName(), "input_x",
	125	+ std::to_string(storageShape.GetDim(i)).c_str(),
	126	+ "Storage shape dimensions must be non-negative"),
	127	+ return ge::GRAPH_FAILED);
	128	+ }
	129	+ const int64_t totalElements = storageShape.GetShapeSize();
	130	+ OP_CHECK_IF(totalElements < 0,
	131	+ OP_LOGE_FOR_INVALID_SHAPESIZE_WITH_REASON(context->GetNodeName(), "input_x",
	132	+ std::to_string(totalElements).c_str(),
	133	+ "The storage shape size must be non-negative"),
	134	+ return ge::GRAPH_FAILED);
	135	+ tilingData->totalElements = totalElements;
	136	+
	137	+ // tiling-key：按 input0 dtype 选择编译期计算路径。ge::DataType 与 C_DT_* 共享 c_types.h 枚举值。
	138	+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(dtype));
	139	+
	140	+ // 空张量：单核占位返回，kernel 不申请 UB。
	141	+ if (totalElements == 0) {
	142	+ context->SetBlockDim(1);
	143	+ return ge::GRAPH_SUCCESS;
	144	+ }
	145	+
	146	+ // 平台信息：UB 大小 + AIV 核数。binary tiling 场景下 platformInfo 可能为空，使用 TilingParse 写入的 compileInfo。
	147	+ auto platformInfoPtr = context->GetPlatformInfo();
	148	+ uint64_t ubSize = 0;
	149	+ int64_t coreNum = 0;
	150	+ if (platformInfoPtr == nullptr) {
	151	+ auto compileInfo = context->GetCompileInfo<HardSigmoidCompileInfo>();
	152	+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
	153	+ ubSize = compileInfo->ubSize;
	154	+ coreNum = static_cast<int64_t>(compileInfo->coreNum);
	155	+ } else {
	156	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	157	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
	158	+ coreNum = static_cast<int64_t>(ascendcPlatform.GetCoreNumAiv());
	159	+ }
	160	+ if (ubSize == 0) {
	161	+ OP_LOGE_WITHOUT_REPORT(context->GetNodeName(), "The platform UB size must be greater than zero, ubSize: %lu",
	162	+ ubSize);
	163	+ return ge::GRAPH_FAILED;
	164	+ }
	165	+ if (coreNum <= 0) {
	166	+ OP_LOGE_WITHOUT_REPORT(context->GetNodeName(),
	167	+ "The platform AIV core count must be greater than zero, coreNum: %ld", coreNum);
	168	+ return ge::GRAPH_FAILED;
	169	+ }
	170	+
	171	+ // ubBlockSize / vRegSize 的非法值统一由 ComputeUbFactor 返回 0，再由下方 ubFactor 守卫拦截并报错。
	172	+ const int64_t ubBlockSize = GetUbBlockSize(context);
	173	+ const int64_t vRegSize = static_cast<int64_t>(GetVRegSize(context));
	174	+ tilingData->ubFactor = ComputeUbFactor(dtype, static_cast<int64_t>(ubSize), ubBlockSize, vRegSize);
	175	+ const std::string platformValues = std::to_string(ubSize) + "," + std::to_string(ubBlockSize) + "," +
	176	+ std::to_string(vRegSize) + "," + std::to_string(static_cast<int32_t>(dtype));
	177	+ if (tilingData->ubFactor <= 0) {
	178	+ OP_LOGE_WITHOUT_REPORT(
	179	+ context->GetNodeName(),
	180	+ "The UB is too small or the platform parameters are invalid, ubSize,ubBlockSize,vRegSize,dtype: %s",
	181	+ platformValues.c_str());
	182	+ return ge::GRAPH_FAILED;
	183	+ }
	184	+
	185	+ // 限制常规核的 GM<->UB 搬运量不低于 16KB；总量不足时使用单核，最后一个尾核允许更小。
	186	+ const int64_t minCopyElements = CeilDiv(MIN_COPY_BYTES, StorageBytesPerElement(dtype));
	187	+ const int64_t maxCoreNumByCopy = totalElements / minCopyElements;
	188	+ const int64_t actualCoreNum = maxCoreNumByCopy <= 0 ? 1 : (maxCoreNumByCopy < coreNum ? maxCoreNumByCopy : coreNum);
	189	+ tilingData->blockFactor = CeilDiv(totalElements, actualCoreNum);
	190	+ // blockFactor 随即作为除数使用；此处就地守卫，避免其非零性仅依赖上游 totalElements>0 的跨行推导。
	191	+ if (tilingData->blockFactor <= 0) {
	192	+ OP_LOGE_WITHOUT_REPORT(context->GetNodeName(),
	193	+ "The per-core block factor must be greater than zero, blockFactor: %ld",
	194	+ tilingData->blockFactor);
	195	+ return ge::GRAPH_FAILED;
	196	+ }
	197	+ const int64_t usedCoreNum = CeilDiv(totalElements, tilingData->blockFactor);
	198	+ context->SetBlockDim(static_cast<uint32_t>(usedCoreNum));
	199	+ return ge::GRAPH_SUCCESS;
	200	+}
	201	+
	202	+static ge::graphStatus TilingParseForHardSigmoid(gert::TilingParseContext* context)
	203	+{
	204	+ auto compileInfo = context->GetCompiledInfo<HardSigmoidCompileInfo>();
	205	+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
	206	+ auto platformInfoPtr = context->GetPlatformInfo();
	207	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	208	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	209	+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
	210	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfo->ubSize);
	211	+ if (compileInfo->coreNum == 0 \|\| compileInfo->ubSize == 0) {
	212	+ OP_LOGE_WITHOUT_REPORT(context->GetNodeName(),
	213	+ "The platform AIV core count and UB size must be greater than zero, coreNum: %lu, "
	214	+ "ubSize: %lu",
	215	+ compileInfo->coreNum, compileInfo->ubSize);
	216	+ return ge::GRAPH_FAILED;
	217	+ }
	218	+ return ge::GRAPH_SUCCESS;
	219	+}
	220	+
	221	+IMPL_OP_OPTILING(HardSigmoid)
	222	+ .Tiling(HardSigmoidTilingFunc)
	223	+ .TilingParse<HardSigmoidCompileInfo>(TilingParseForHardSigmoid);
	224	+} // namespace optiling

Aactivation/hard_sigmoid/op_host/hard_sigmoid_def.cpp+47-0

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

Aactivation/hard_sigmoid/op_host/hard_sigmoid_infershape.cpp+35-0

@@ -0,0 +1,35 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include "register/op_impl_registry.h"
	12	+#include "exe_graph/runtime/infer_shape_context.h"
	13	+#include "op_common/log/log.h"
	14	+
	15	+namespace ops {
	16	+static ge::graphStatus InferShape4HardSigmoid(gert::InferShapeContext* context)
	17	+{
	18	+ const gert::Shape* inputShape = context->GetInputShape(0);
	19	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
	20	+
	21	+ gert::Shape* outputShape = context->GetOutputShape(0);
	22	+ OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);
	23	+
	24	+ outputShape = inputShape;
	25	+ return ge::GRAPH_SUCCESS;
	26	+}
	27	+
	28	+static ge::graphStatus InferDataType4HardSigmoid(gert::InferDataTypeContext* context)
	29	+{
	30	+ context->SetOutputDataType(0, context->GetInputDataType(0));
	31	+ return ge::GRAPH_SUCCESS;
	32	+}
	33	+
	34	+IMPL_OP_INFERSHAPE(HardSigmoid).InferShape(InferShape4HardSigmoid).InferDataType(InferDataType4HardSigmoid);
	35	+} // namespace ops

Mactivation/hard_sigmoid/op_host/op_api/aclnn_hardsigmoid.cpp+11-4

	@@ -8,6 +8,7 @@
8	* See LICENSE in the root of the software repository for the full text of the License.	8	* See LICENSE in the root of the software repository for the full text of the License.
9	*/	9	*/
10	#include "aclnn_hardsigmoid.h"	10	#include "aclnn_hardsigmoid.h"
		11	+#include <string>
		L L liubo757月25日这个op_api的目录咋还在op_host内部，要移动到外面同级 T tianyu52 7月25日评论： 7月25日评论：
11	#include "aclnn_kernels/contiguous.h"	12	#include "aclnn_kernels/contiguous.h"
12	#include "hardsigmoid.h"	13	#include "hardsigmoid.h"
13	#include "aclnn/aclnn_base.h"	14	#include "aclnn/aclnn_base.h"
	@@ -20,10 +21,12 @@
20	#include "opdev/op_log.h"	21	#include "opdev/op_log.h"
21	#include "opdev/tensor_view_utils.h"	22	#include "opdev/tensor_view_utils.h"
22	#include "aclnn_kernels/common/op_error_check.h"	23	#include "aclnn_kernels/common/op_error_check.h"
		24	+#include "log/log.h"
23	#include "op_api/level2_base.h"	25	#include "op_api/level2_base.h"
24	#include "op_api/op_api_def.h"	26	#include "op_api/op_api_def.h"
25		27
26	using namespace op;	28	using namespace op;
		29	+static constexpr const char* ACLNN_HARD_SIGMOID_NAME = "aclnnHardsigmoid";
27	#ifdef __cplusplus	30	#ifdef __cplusplus
28	extern "C" {	31	extern "C" {
29	#endif	32	#endif
	@@ -62,13 +65,17 @@ static aclnnStatus CheckParams(const aclTensor* self, const aclTensor* out)
62		65
63	// 4. 检查输入和输出的类型、数据格式是否一致	66	// 4. 检查输入和输出的类型、数据格式是否一致
64	if (self->GetDataType() != out->GetDataType()) {	67	if (self->GetDataType() != out->GetDataType()) {
65	- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Input tensor's dtype[%s] should be same with output's dtype[%s].",	68	+ const std::string dtypeStr = std::string(op::ToString(self->GetDataType()).GetString()) + ", " +
66	- op::ToString(self->GetDataType()).GetString(), op::ToString(out->GetDataType()).GetString());	69	+ op::ToString(out->GetDataType()).GetString();
		70	+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(ACLNN_HARD_SIGMOID_NAME, "self, out", dtypeStr.c_str(),
		71	+ "the dtype of self and out must be the same");
67	return ACLNN_ERR_PARAM_INVALID;	72	return ACLNN_ERR_PARAM_INVALID;
68	}	73	}
69	if (self->GetStorageFormat() != out->GetStorageFormat()) {	74	if (self->GetStorageFormat() != out->GetStorageFormat()) {
70	- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Input tensor's format[%s] and output's format[%s] should be same.",	75	+ const std::string formatStr = std::string(op::ToString(self->GetStorageFormat()).GetString()) + ", " +
71	- op::ToString(self->GetStorageFormat()).GetString(), op::ToString(out->GetStorageFormat()).GetString());	76	+ op::ToString(out->GetStorageFormat()).GetString();
		77	+ OP_LOGE_FOR_INVALID_FORMATS_WITH_REASON(ACLNN_HARD_SIGMOID_NAME, "self, out", formatStr.c_str(),
		78	+ "the format of self and out must be the same");
72	return ACLNN_ERR_PARAM_INVALID;	79	return ACLNN_ERR_PARAM_INVALID;
73	}	80	}
74	return ACLNN_SUCCESS;	81	return ACLNN_SUCCESS;

Mactivation/hard_sigmoid/op_host/op_api/aclnn_hardsigmoid.h+5-4

	@@ -40,10 +40,11 @@ extern "C" {
40	* ```	40	* ```
41	*	41	*
42	* @param [in] self: npu	42	* @param [in] self: npu
43	- * device侧的aclTensor，数据类型支持 FLOAT、FLOAT16、INT32，支持非连续的Tensor。	43	+ * device侧的aclTensor，数据类型支持 FLOAT、FLOAT16、INT32、BFLOAT16，支持非连续的Tensor。
		Z Z zl_hw7月24日 BFLOAT16是A5新增的么，如果是新增的，得区分开写 T tianyu52 7月24日评论： 7月24日评论：
44	- * 支持非连续的Tensor，数据格式支持ND	44	+ * Atlas 训练系列产品不支持BFLOAT16，数据格式支持ND。
45	* @param [in] out: npu	45	* @param [in] out: npu
46	- * device侧的aclTensor，数据类型支持 FLOAT、FLOAT16、INT32，支持非连续的Tensor，数据格式支持ND。	46	+ * device侧的aclTensor，数据类型支持 FLOAT、FLOAT16、INT32、BFLOAT16，支持非连续的Tensor。
		47	+ * Atlas 训练系列产品不支持BFLOAT16，数据格式支持ND。
47	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。	48	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
48	* @param [out] executor: 返回op执行器，包含算子计算流程。	49	* @param [out] executor: 返回op执行器，包含算子计算流程。
49	* @return aclnnStatus: 返回状态码。	50	* @return aclnnStatus: 返回状态码。
	@@ -102,7 +103,7 @@ ACLNN_API aclnnStatus aclnnHardsigmoid(void* workspace, uint64_t workspaceSize,
102	```	103	```
103	*	104	*
104	* @param [in] self: npu	105	* @param [in] self: npu
105	- * device侧的aclTensor，数据类型支持 FLOAT、FLOAT16、INT32，支持非连续的Tensor，数据格式支持ND	106	+ * device侧的aclTensor，数据类型支持 FLOAT、FLOAT16、INT32、BFLOAT16，支持非连续的Tensor，数据格式支持ND
106	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。	107	* @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
107	* @param [out] executor: 返回op执行器，包含算子计算流程。	108	* @param [out] executor: 返回op执行器，包含算子计算流程。
108	* @return aclnnStatus: 返回状态码。	109	* @return aclnnStatus: 返回状态码。

Mactivation/hard_sigmoid/op_host/op_api/hardsigmoid.cpp+2-2

	@@ -28,7 +28,7 @@ OP_TYPE_REGISTER(HardSigmoid);
28	const aclTensor* HardSigmoid(const aclTensor* self, aclOpExecutor* executor)	28	const aclTensor* HardSigmoid(const aclTensor* self, aclOpExecutor* executor)
29	{	29	{
30	L0_DFX(HardSigmoid, self);	30	L0_DFX(HardSigmoid, self);
31	- auto hardsigmoidOut = executor->AllocTensor(self->GetStorageShape(), self->GetDataType(), self->GetStorageFormat());	31	+ auto hardsigmoidOut = executor->AllocTensor(self->GetViewShape(), self->GetDataType());
32	CHECK_RET(hardsigmoidOut != nullptr, nullptr);	32	CHECK_RET(hardsigmoidOut != nullptr, nullptr);
33	auto ret = ADD_TO_LAUNCHER_LIST_AICORE(HardSigmoid, OP_INPUT(self), OP_ATTR(kHardSigmoidAlpha, kHardSigmoidBeta),	33	auto ret = ADD_TO_LAUNCHER_LIST_AICORE(HardSigmoid, OP_INPUT(self), OP_ATTR(kHardSigmoidAlpha, kHardSigmoidBeta),
34	OP_OUTPUT(hardsigmoidOut));	34	OP_OUTPUT(hardsigmoidOut));
	@@ -36,4 +36,4 @@ const aclTensor* HardSigmoid(const aclTensor* self, aclOpExecutor* executor)
36	OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "HardSigmoidAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."), return nullptr);	36	OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "HardSigmoidAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."), return nullptr);
37	return hardsigmoidOut;	37	return hardsigmoidOut;
38	}	38	}
39	-} // namespace l0op	39	+} // namespace l0op

Aactivation/hard_sigmoid/op_kernel/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	+add_kernel_sources(
	12	+ KERNEL_SRC arch35/hard_sigmoid.cpp
	13	+ COMPUTE_UNITS ascend950
	14	+ AUTO_SYNC false
	15	+)

Aactivation/hard_sigmoid/op_kernel/arch35/hard_sigmoid.cpp+176-0

@@ -0,0 +1,176 @@
	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 hard_sigmoid.cpp
	13	+ * \brief HardSigmoid kernel（arch35 / DAV_3510，非模板 regbase 手写实现）
	14	+ *
	15	+ * y = clamp(alpha * x + beta, 0, 1)
	16	+ *
	17	+ * 采用 EnQue/DeQue 三级流水（CopyIn / Compute / CopyOut），跨管道同步由 Queue 自动管理：
	18	+ * - FLOAT : 原生 fp32 计算
	19	+ * - FLOAT16/BFLOAT16: Cast 升 fp32 计算，再按目标类型舍入转回
	20	+ * - INT32 : CAST_RINT 升 fp32 计算，再 CAST_TRUNC 转回
	21	+ */
	22	+
	23	+#include <type_traits>
	24	+
	25	+#include "kernel_operator.h"
	26	+#include "kernel_tiling/kernel_tiling.h"
	27	+#include "hard_sigmoid_tiling_data.h"
	28	+#include "hard_sigmoid_tiling_key.h"
	29	+
	30	+using namespace AscendC;
	31	+
	32	+// 缓冲深度取自 Host/Kernel 共享头，确保与 tiling 的 UB 预算推导一致。
	33	+static constexpr int32_t BUFFER_NUM = static_cast<int32_t>(HARD_SIGMOID_BUFFER_NUM);
	34	+
	35	+template <typename T>
	36	+class HardSigmoidKernel {
	37	+public:
	38	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const HardSigmoidTilingData* tilingData)
	39	+ {
	40	+ totalElements_ = tilingData->totalElements;
	41	+ blockFactor_ = tilingData->blockFactor;
	42	+ ubFactor_ = tilingData->ubFactor;
	43	+ alpha_ = tilingData->alpha;
	44	+ beta_ = tilingData->beta;
	45	+
	46	+ const int64_t coreIdx = static_cast<int64_t>(GetBlockIdx());
	47	+ startIdx_ = coreIdx * blockFactor_;
	48	+ int64_t remainderLength = totalElements_ - startIdx_;
	49	+ blockLength_ = (remainderLength > blockFactor_) ? blockFactor_ : remainderLength;
	50	+ if (startIdx_ >= totalElements_ \|\| blockLength_ <= 0 \|\| ubFactor_ <= 0) {
	51	+ blockLength_ = 0;
	52	+ return;
	53	+ }
	54	+
	55	+ xGM_.SetGlobalBuffer((__gm__ T*)x + startIdx_, blockLength_);
	56	+ yGM_.SetGlobalBuffer((__gm__ T*)y + startIdx_, blockLength_);
	57	+
	58	+ pipe_.InitBuffer(inQue_, BUFFER_NUM, ubFactor_ * sizeof(T));
	59	+ pipe_.InitBuffer(outQue_, BUFFER_NUM, ubFactor_ * sizeof(T));
	60	+ if constexpr (!std::is_same_v<T, float>) {
	61	+ pipe_.InitBuffer(f32Buf_, ubFactor_ * sizeof(float));
	62	+ }
	63	+
	64	+ loopCount_ = (blockLength_ + ubFactor_ - 1) / ubFactor_;
	65	+ }
	66	+
	67	+ __aicore__ inline void Process()
	68	+ {
	69	+ if (blockLength_ <= 0) {
	70	+ return;
	71	+ }
	72	+ for (int64_t ci = 0; ci < loopCount_; ci++) {
	73	+ uint32_t currentChunk = static_cast<uint32_t>((ci == (loopCount_ - 1)) ? (blockLength_ - ubFactor_ * ci) :
	74	+ ubFactor_);
	75	+ CopyIn(ci, currentChunk);
	76	+ Compute(currentChunk);
	77	+ CopyOut(ci, currentChunk);
	78	+ }
	79	+ }
	80	+
	81	+private:
	82	+ __aicore__ inline void CopyIn(int64_t ci, uint32_t currentChunk)
	83	+ {
	84	+ LocalTensor<T> xLocal = inQue_.template AllocTensor<T>();
	85	+ DataCopyExtParams cpIn;
	86	+ cpIn.blockCount = 1;
	87	+ cpIn.blockLen = currentChunk * sizeof(T);
	88	+ cpIn.srcStride = 0;
	89	+ cpIn.dstStride = 0;
	90	+ DataCopyPad(xLocal, xGM_[ci * ubFactor_], cpIn, {false, 0, 0, 0});
	91	+ inQue_.EnQue(xLocal);
	92	+ }
	93	+
	94	+ __aicore__ inline void Compute(uint32_t currentChunk)
	95	+ {
	96	+ LocalTensor<T> xLocal = inQue_.template DeQue<T>();
	97	+ LocalTensor<T> yLocal = outQue_.template AllocTensor<T>();
	98	+
	99	+ if constexpr (std::is_same_v<T, float>) {
	100	+ // FLOAT: 原生 fp32，affine + clamp 全程 fp32
	101	+ Muls(yLocal, xLocal, alpha_, currentChunk);
	102	+ Adds(yLocal, yLocal, beta_, currentChunk);
	103	+ Mins(yLocal, yLocal, 1.0f, currentChunk);
	104	+ Maxs(yLocal, yLocal, 0.0f, currentChunk);
	105	+ } else if constexpr (std::is_same_v<T, int32_t>) {
	106	+ // INT32: RINT 升 fp32 -> affine+clamp(fp32) -> TRUNC 转回
	107	+ LocalTensor<float> f32 = f32Buf_.template Get<float>();
	108	+ Cast(f32, xLocal, RoundMode::CAST_RINT, currentChunk);
	109	+ Muls(f32, f32, alpha_, currentChunk);
	110	+ Adds(f32, f32, beta_, currentChunk);
	111	+ Mins(f32, f32, 1.0f, currentChunk);
	112	+ Maxs(f32, f32, 0.0f, currentChunk);
	113	+ Cast(yLocal, f32, RoundMode::CAST_TRUNC, currentChunk);
	114	+ } else {
	115	+ // FLOAT16/BFLOAT16: 升 fp32 -> affine(fp32) -> 按目标类型舍入降回 -> clamp 在原生 dtype
	116	+ // 0.0/1.0 在 fp16/bf16 均可精确表示，先降后 clamp 与 golden(fp32内clamp后降) 等价，
	117	+ // 且省去 2 个 fp32 域 clamp 算子（原生 dtype 向量吞吐更高）。
	118	+ LocalTensor<float> f32 = f32Buf_.template Get<float>();
	119	+ Cast(f32, xLocal, RoundMode::CAST_NONE, currentChunk);
	120	+ Muls(f32, f32, alpha_, currentChunk);
	121	+ Adds(f32, f32, beta_, currentChunk);
	122	+ if constexpr (std::is_same_v<T, bfloat16_t>) {
	123	+ Cast(yLocal, f32, RoundMode::CAST_ROUND, currentChunk);
	124	+ } else {
	125	+ Cast(yLocal, f32, RoundMode::CAST_RINT, currentChunk);
	126	+ }
	127	+ Mins(yLocal, yLocal, static_cast<T>(1.0f), currentChunk);
	128	+ Maxs(yLocal, yLocal, static_cast<T>(0.0f), currentChunk);
	129	+ }
	130	+
	131	+ outQue_.template EnQue<T>(yLocal);
	132	+ inQue_.FreeTensor(xLocal);
	133	+ }
	134	+
	135	+ __aicore__ inline void CopyOut(int64_t ci, uint32_t currentChunk)
	136	+ {
	137	+ LocalTensor<T> yLocal = outQue_.template DeQue<T>();
	138	+ DataCopyExtParams cpOut;
	139	+ cpOut.blockCount = 1;
	140	+ cpOut.blockLen = currentChunk * sizeof(T);
	141	+ cpOut.srcStride = 0;
	142	+ cpOut.dstStride = 0;
	143	+ DataCopyPad(yGM_[ci * ubFactor_], yLocal, cpOut);
	144	+ outQue_.FreeTensor(yLocal);
	145	+ }
	146	+
	147	+private:
	148	+ TPipe pipe_;
	149	+ TQue<QuePosition::VECIN, BUFFER_NUM> inQue_;
	150	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQue_;
	151	+ TBuf<TPosition::VECCALC> f32Buf_; // 非 fp32 dtype 的 fp32 中间计算缓冲
	152	+
	153	+ GlobalTensor<T> xGM_, yGM_;
	154	+
	155	+ int64_t totalElements_ = 0;
	156	+ int64_t blockFactor_ = 0;
	157	+ int64_t ubFactor_ = 0;
	158	+ int64_t startIdx_ = 0;
	159	+ int64_t blockLength_ = 0;
	160	+ int64_t loopCount_ = 0;
	161	+ float alpha_ = 0.0f;
	162	+ float beta_ = 0.0f;
	163	+};
	164	+
	165	+// D_T_X（input0 dtype tiling-key）分发：每个 dtype 计算路径独立实例化，无运行时 dtype 分支。
	166	+template <typename D_T_X>
	167	+__global__ __aicore__ void hard_sigmoid(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
	168	+{
	169	+ REGISTER_TILING_DEFAULT(HardSigmoidTilingData);
	170	+ GET_TILING_DATA_WITH_STRUCT(HardSigmoidTilingData, tilingData, tiling);
	171	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	172	+
	173	+ HardSigmoidKernel<D_T_X> kernel;
	174	+ kernel.Init(x, y, &tilingData);
	175	+ kernel.Process();
	176	+}

Aactivation/hard_sigmoid/op_kernel/arch35/hard_sigmoid_tiling_data.h+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 hard_sigmoid_tiling_data.h
	13	+ * \brief HardSigmoid TilingData（arch35 regbase 写法）。
	14	+ */
	15	+
	16	+#ifndef HARD_SIGMOID_TILING_DATA_H
	17	+#define HARD_SIGMOID_TILING_DATA_H
	18	+
	19	+#include <cstdint>
	20	+
	21	+// kernel 侧 in/out 队列的缓冲深度（Double Buffer）。
	22	+// Host tiling 依据它推导每元素 UB 占用（见 UbBytesPerElement），二者必须一致，
	23	+// 故置于 Host/Kernel 共享头，禁止任一侧另行定义。
	24	+constexpr int64_t HARD_SIGMOID_BUFFER_NUM = 2;
	25	+
	26	+struct __attribute__((aligned(8))) HardSigmoidTilingData {
	27	+ int64_t totalElements = 0; // 元素总数
	28	+ int64_t blockFactor = 0; // 单核处理的元素数
	29	+ int64_t ubFactor = 0; // 单次 UB 搬运/计算的元素数
	30	+ float alpha = 1.0f / 6.0f; // y = clamp(alpha*x + beta, 0, 1)
	31	+ float beta = 0.5f;
	32	+};
	33	+
	34	+#endif // HARD_SIGMOID_TILING_DATA_H

Aactivation/hard_sigmoid/op_kernel/arch35/hard_sigmoid_tiling_key.h+40-0

@@ -0,0 +1,40 @@
	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 hard_sigmoid_tiling_key.h
	13	+ * \brief HardSigmoid TilingKey 声明。
	14	+ *
	15	+ * D_T_X 作为单轴 tiling-key，绑定 input0 的 dtype 并枚举各 dtype 的计算路径。这里使用
	16	+ * DATATYPE 模板轴而不是自定义 uint 模式轴，保证 binary 构建产物和运行时 tiling 解析能从
	17	+ * supportInfo 的输入 dtype 直接选择模板实例。
	18	+ */
	19	+
	20	+#ifndef HARD_SIGMOID_TILING_KEY_H
	21	+#define HARD_SIGMOID_TILING_KEY_H
	22	+
	23	+#include "ascendc/host_api/tiling/template_argument.h"
	24	+
	25	+// Kernel UT host 编译时 ASCENDC_CPU_DEBUG 被定义，ASCENDC_TPL_DATATYPE_DECL
	26	+// 会展开为引用 C_DT_* 的 ParamStruct 构造；graph/c_types.h 提供这些枚举定义。
	27	+// 真机 kernel 编译路径不需要该头文件。
	28	+#ifdef ASCENDC_CPU_DEBUG
	29	+#include "graph/c_types.h"
	30	+#endif
	31	+
	32	+ASCENDC_TPL_ARGS_DECL(HardSigmoid, ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16, C_DT_INT32,
	W W wang-xing0017月25日 tilingKey不要使用D_T_X，可以直接在kernel入喉使用DTYPE_X变量
	33	+ ASCENDC_TPL_INPUT(0)));
	34	+
	35	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT)),
	36	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16)),
	37	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_BF16)),
	38	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT32)), );
	39	+
	40	+#endif // HARD_SIGMOID_TILING_KEY_H

Aactivation/hard_sigmoid/tests/assets/golden.py+101-0

@@ -0,0 +1,101 @@
	1	+#!/usr/bin/env python3
	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	+import numpy as np
	11	+import torch
	12	+
	13	+
	14	+__golden__ = {"kernel": {"hard_sigmoid": "hard_sigmoid_golden"}}
	15	+__input__ = {"kernel": {"hard_sigmoid": "hard_sigmoid_input"}}
	16	+
	17	+
	18	+def hard_sigmoid_golden(x, alpha=1.0 / 6.0, beta=0.5, **kwargs):
	19	+ """Compute HardSigmoid through PyTorch operator composition."""
	20	+ del kwargs
	21	+ input_dtype = x.dtype
	22	+ if input_dtype.name == "bfloat16":
	23	+ input_tensor = torch.tensor(
	24	+ np.asarray(x, dtype=np.float32), dtype=torch.bfloat16
	25	+ )
	26	+ else:
	27	+ input_tensor = torch.from_numpy(np.asarray(x))
	28	+ alpha32 = float(np.float32(alpha))
	29	+ beta32 = float(np.float32(beta))
	30	+ result = torch.clamp(
	31	+ input_tensor.to(torch.float32) * alpha32 + beta32, min=0.0, max=1.0
	32	+ )
	33	+ if not input_tensor.is_floating_point():
	34	+ result = torch.trunc(result)
	35	+ result = result.to(input_tensor.dtype)
	36	+ return (
	37	+ result.to(torch.float32).cpu().numpy()
	38	+ if input_dtype.name == "bfloat16"
	39	+ else result.cpu().numpy()
	40	+ )
	41	+
	42	+
	43	+def hard_sigmoid_input(x, alpha=1.0 / 6.0, beta=0.5, **kwargs):
	44	+ """Inject clamp boundaries and special values while retaining each case's requested dtype and shape."""
	45	+ if x.size == 0:
	46	+ return [x]
	47	+
	48	+ testcase_name = kwargs.get("testcase_name", "")
	49	+ if testcase_name == "hard_sigmoid_fp32_special":
	50	+ tiny = np.finfo(np.float32).tiny
	51	+ critical = np.array(
	52	+ [
	53	+ -np.inf,
	54	+ -8.0,
	55	+ -3.0001,
	56	+ -3.0,
	57	+ -2.9999,
	58	+ -1.0,
	59	+ -tiny,
	60	+ -0.0,
	61	+ 0.0,
	62	+ tiny,
	63	+ 1.0,
	64	+ 2.9999,
	65	+ 3.0,
	66	+ 3.0001,
	67	+ 8.0,
	68	+ np.inf,
	69	+ np.nan,
	70	+ ],
	71	+ dtype=np.float32,
	72	+ )
	73	+ else:
	74	+ alpha32 = np.float32(alpha)
	75	+ beta32 = np.float32(beta)
	76	+ if alpha32 == 0:
	77	+ critical = np.array([-8.0, -1.0, 0.0, 1.0, 8.0], dtype=np.float32)
	78	+ else:
	79	+ zero_boundary = -beta32 / alpha32
	80	+ one_boundary = (np.float32(1.0) - beta32) / alpha32
	81	+ epsilon = np.float32(1.0e-3)
	82	+ critical = np.array(
	83	+ [
	84	+ zero_boundary - epsilon,
	85	+ zero_boundary,
	86	+ zero_boundary + epsilon,
	87	+ -1.0,
	88	+ 0.0,
	89	+ 1.0,
	90	+ one_boundary - epsilon,
	91	+ one_boundary,
	92	+ one_boundary + epsilon,
	93	+ ],
	94	+ dtype=np.float32,
	95	+ )
	96	+
	97	+ result = np.array(x, copy=True)
	98	+ flat = result.reshape(-1)
	99	+ count = min(flat.size, critical.size)
	100	+ flat[:count] = critical[:count].astype(result.dtype, copy=False)
	101	+ return [flat.reshape(result.shape)]

Aactivation/hard_sigmoid/tests/st/arch35/ttk_kernel_hard_sigmoid.csv+7-0

@@ -0,0 +1,7 @@
	1	+testcase_name,network_name,op_name,input_shapes,input_dtypes,input_formats,output_shapes,output_dtypes,output_formats,input_ori_shapes,input_ori_formats,output_ori_shapes,output_ori_formats,attributes,input_data_ranges,precision_tolerances,absolute_precision,output_inplace_indexes,output_shape_unknown_indexes,is_enabled,remark,soc_series,priority,dump_file_prefix,manual_input_binaries,manual_golden_binaries
	2	+hard_sigmoid_fp32_special,UNKNOWN,hard_sigmoid,"((17,),)","('float32',)","('ND',)","((17,),)","('float32',)","('ND',)","((17,),)","('ND',)","((17,),)","('ND',)","{'alpha':0.16666666666666666,'beta':0.5}","((-8,8),)","((0.0001,0.0001),)",0.0001,(),(),True,boundary-inf-nan,"('Ascend950',)",0,,(),()
	3	+hard_sigmoid_fp32_custom,UNKNOWN,hard_sigmoid,"((2,3,5),)","('float32',)","('ND',)","((2,3,5),)","('float32',)","('ND',)","((2,3,5),)","('ND',)","((2,3,5),)","('ND',)","{'alpha':0.2,'beta':0.4}","((-8,8),)","((0.0001,0.0001),)",0.0001,(),(),True,custom-attributes,"('Ascend950',)",0,,(),()
	4	+hard_sigmoid_fp16_unaligned,UNKNOWN,hard_sigmoid,"((3,17),)","('float16',)","('ND',)","((3,17),)","('float16',)","('ND',)","((3,17),)","('ND',)","((3,17),)","('ND',)","{'alpha':0.16666666666666666,'beta':0.5}","((-8,8),)","((0.001,0.001),)",0.001,(),(),True,unaligned-fp16,"('Ascend950',)",0,,(),()
	5	+hard_sigmoid_bf16_unaligned,UNKNOWN,hard_sigmoid,"((5,13),)","('bfloat16',)","('ND',)","((5,13),)","('bfloat16',)","('ND',)","((5,13),)","('ND',)","((5,13),)","('ND',)","{'alpha':0.16666666666666666,'beta':0.5}","((-8,8),)","((0.004,0.004),)",0.004,(),(),True,unaligned-bf16,"('Ascend950',)",0,,(),()
	6	+hard_sigmoid_int32,UNKNOWN,hard_sigmoid,"((67,),)","('int32',)","('ND',)","((67,),)","('int32',)","('ND',)","((67,),)","('ND',)","((67,),)","('ND',)","{'alpha':0.2,'beta':0.4}","((-10,10),)","((0,0),)",0,(),(),True,int32,"('Ascend950',)",0,,(),()
	7	+hard_sigmoid_empty,UNKNOWN,hard_sigmoid,"((0,),)","('float32',)","('ND',)","((0,),)","('float32',)","('ND',)","((0,),)","('ND',)","((0,),)","('ND',)","{'alpha':0.16666666666666666,'beta':0.5}","((0,0),)","((0.0001,0.0001),)",0.0001,(),(),True,empty-tensor,"('Ascend950',)",0,,(),()

Aactivation/hard_sigmoid/tests/ut/op_host/arch35/CMakeLists.txt+19-0

@@ -0,0 +1,19 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+#
	10	+# Generated By CANNBot
	11	+# ----------------------------------------------------------------------------
	12	+
	13	+# arch35 tiling 用例依赖 arch35 tiling 实现(op_host/arch35/_tiling.cpp)被链接进 UT。
	14	+# 该实现仅在 ASCEND_COMPUTE_UNIT=ascend950 时才会被 add_tiling_sources 纳入编译；
	15	+# 未指定 soc 时 tiling 符号缺失，用例调用 tiling_parse 会踩空指针，故此处一并跳过。
	16	+# 运行方式：bash build.sh -u --ops=hard_sigmoid --soc=ascend950
	17	+if((UT_TEST_ALL OR OP_HOST_UT) AND ASCEND_COMPUTE_UNIT MATCHES "ascend950")
	18	+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	19	+endif()

Aactivation/hard_sigmoid/tests/ut/op_host/arch35/test_hard_sigmoid_tiling.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	+
	11	+#include <gtest/gtest.h>
	12	+#include <cstdint>
	13	+#include <cstring>
	14	+#include <map>
	15	+#include <string>
	16	+#include <vector>
	17	+#include "exe_graph/runtime/storage_shape.h"
	18	+#include "kernel_run_context_facker.h"
	19	+#include "platform/platform_infos_def.h"
	20	+#include "test_cube_util.h"
	21	+#include "ut_op_util.h"
	22	+#include "../../../../op_kernel/arch35/hard_sigmoid_tiling_data.h"
	23	+
	24	+using namespace ut_util;
	25	+
	26	+namespace optiling {
	27	+struct HardSigmoidCompileInfo {
	28	+ uint64_t coreNum = 0;
	29	+ uint64_t ubSize = 0;
	30	+};
	31	+} // namespace optiling
	32	+
	33	+namespace {
	34	+constexpr const char* COMPILE_INFO = R"({
	35	+ "hardware_info": {
	36	+ "BT_SIZE": 0,
	37	+ "load3d_constraints": "1",
	38	+ "Intrinsic_fix_pipe_l0c2out": false,
	39	+ "Intrinsic_data_move_l12ub": true,
	40	+ "Intrinsic_data_move_l0c2ub": true,
	41	+ "Intrinsic_data_move_out2l1_nd2nz": false,
	42	+ "UB_SIZE": 245760,
	43	+ "L2_SIZE": 33554432,
	44	+ "L1_SIZE": 524288,
	45	+ "L0A_SIZE": 65536,
	46	+ "L0B_SIZE": 65536,
	47	+ "L0C_SIZE": 131072,
	48	+ "CORE_NUM": 64
	49	+ }
	50	+})";
	51	+
	52	+struct TilingResult {
	53	+ ge::graphStatus status = ge::GRAPH_FAILED;
	54	+ uint32_t blockDim = 0;
	55	+ HardSigmoidTilingData tilingData{};
	56	+};
	57	+
	58	+TilingResult RunTilingCase(const std::vector<int64_t>& dims, ge::DataType dtype, float alpha = 1.0f / 6.0f,
	59	+ float beta = 0.5f)
	60	+{
	61	+ gert::StorageShape shape;
	62	+ for (const int64_t dim : dims) {
	63	+ shape.MutableOriginShape().AppendDim(dim);
	64	+ shape.MutableStorageShape().AppendDim(dim);
	65	+ }
	66	+ std::map<std::string, std::string> socInfos;
	67	+ std::map<std::string, std::string> aicoreSpec;
	68	+ std::map<std::string, std::string> intrinsics;
	69	+ std::map<std::string, std::string> version = {{"Short_SoC_version", "Ascend950"}, {"NpuArch", "3510"}};
	70	+ GetPlatFormInfos(COMPILE_INFO, socInfos, aicoreSpec, intrinsics);
	71	+
	72	+ fe::PlatFormInfos platformInfo;
	73	+ platformInfo.Init();
	74	+ optiling::HardSigmoidCompileInfo compileInfo;
	75	+ TilingResult result;
	76	+
	77	+ auto impl = gert::OpImplRegistry::GetInstance().GetOpImpl("HardSigmoid");
	78	+ if (impl == nullptr \|\| impl->tiling_parse == nullptr \|\| impl->tiling == nullptr) {
	79	+ ADD_FAILURE() << "HardSigmoid tiling callbacks are not registered";
	80	+ return result;
	81	+ }
	82	+ auto parseContextHolder = gert::KernelRunContextFaker()
	83	+ .KernelIONum(1, 1)
	84	+ .Inputs({const_cast<char>(COMPILE_INFO), reinterpret_cast<void>(&platformInfo)})
	85	+ .Outputs({&compileInfo})
	86	+ .Build();
	87	+ auto* parseContext = parseContextHolder.GetContext<gert::TilingParseContext>();
	88	+ if (parseContext == nullptr \|\| parseContext->GetPlatformInfo() == nullptr \|\|
	89	+ !parseContext->GetPlatformInfo()->Init()) {
	90	+ ADD_FAILURE() << "Failed to create the tiling parse context";
	91	+ return result;
	92	+ }
	93	+ auto* parsePlatformInfo = parseContext->GetPlatformInfo();
	94	+ parsePlatformInfo->SetPlatformRes("SoCInfo", socInfos);
	95	+ parsePlatformInfo->SetPlatformRes("AICoreSpec", aicoreSpec);
	96	+ parsePlatformInfo->SetCoreNumByCoreType("AICore");
	97	+ parsePlatformInfo->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
	98	+ parsePlatformInfo->SetPlatformRes("version", version);
	99	+ if (impl->tiling_parse(parseContextHolder.GetContext<gert::KernelContext>()) != ge::GRAPH_SUCCESS) {
	100	+ ADD_FAILURE() << "HardSigmoid tiling parse failed";
	101	+ return result;
	102	+ }
	103	+
	104	+ auto tilingData = gert::TilingData::CreateCap(4096);
	105	+ auto workspaceHolder = gert::ContinuousVector::Create<size_t>(4096);
	106	+ if (tilingData == nullptr \|\| workspaceHolder == nullptr) {
	107	+ ADD_FAILURE() << "Failed to allocate tiling test buffers";
	108	+ return result;
	109	+ }
	110	+ auto* workspace = reinterpret_cast<gert::ContinuousVector*>(workspaceHolder.get());
	111	+ auto contextHolder = gert::TilingContextFaker()
	112	+ .SetOpType("HardSigmoid")
	113	+ .NodeIoNum(1, 1)
	114	+ .IrInstanceNum({1})
	115	+ .InputShapes({&shape})
	116	+ .OutputShapes({&shape})
	117	+ .CompileInfo(&compileInfo)
	118	+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
	119	+ .NodeInputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	120	+ .NodeOutputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	121	+ .NodeAttrs({{"alpha", Ops::NN::AnyValue::CreateFrom<float>(alpha)},
	122	+ {"beta", Ops::NN::AnyValue::CreateFrom<float>(beta)}})
	123	+ .TilingData(tilingData.get())
	124	+ .Workspace(workspace)
	125	+ .Build();
	126	+ auto* context = contextHolder.GetContext<gert::TilingContext>();
	127	+ if (context == nullptr \|\| context->GetPlatformInfo() == nullptr) {
	128	+ ADD_FAILURE() << "Failed to create the tiling context";
	129	+ return result;
	130	+ }
	131	+ auto* tilingPlatformInfo = context->GetPlatformInfo();
	132	+ tilingPlatformInfo->SetPlatformRes("SoCInfo", socInfos);
	133	+ tilingPlatformInfo->SetPlatformRes("AICoreSpec", aicoreSpec);
	134	+ tilingPlatformInfo->SetCoreNumByCoreType("AICore");
	135	+ tilingPlatformInfo->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
	136	+ result.status = impl->tiling(context);
	137	+ result.blockDim = context->GetBlockDim();
	138	+ auto* rawTilingData = context->GetRawTilingData();
	139	+ if (result.status == ge::GRAPH_SUCCESS && rawTilingData != nullptr &&
	140	+ rawTilingData->GetDataSize() >= sizeof(HardSigmoidTilingData)) {
	141	+ std::memcpy(&result.tilingData, rawTilingData->GetData(), sizeof(HardSigmoidTilingData));
	142	+ }
	143	+ return result;
	144	+}
	145	+} // namespace
	146	+
	147	+TEST(HardSigmoidTilingTest, Float32UsesIndependentCoreAndUbSplits)
	148	+{
	149	+ const auto result = RunTilingCase({262144}, ge::DT_FLOAT);
	150	+ EXPECT_EQ(result.status, ge::GRAPH_SUCCESS);
	151	+ EXPECT_EQ(result.blockDim, 64U);
	152	+ EXPECT_EQ(result.tilingData.totalElements, 262144);
	153	+ EXPECT_EQ(result.tilingData.blockFactor, 4096);
	154	+ EXPECT_EQ(result.tilingData.ubFactor, 14848);
	155	+}
	156	+
	157	+TEST(HardSigmoidTilingTest, Float16)
	158	+{
	159	+ const auto result = RunTilingCase({3, 257}, ge::DT_FLOAT16);
	160	+ EXPECT_EQ(result.status, ge::GRAPH_SUCCESS);
	161	+ EXPECT_EQ(result.blockDim, 1U);
	162	+ EXPECT_EQ(result.tilingData.blockFactor, 771);
	163	+ EXPECT_EQ(result.tilingData.ubFactor, 19712);
	164	+}
	165	+
	166	+TEST(HardSigmoidTilingTest, Bfloat16)
	167	+{
	168	+ const auto result = RunTilingCase({2, 7, 65}, ge::DT_BF16);
	169	+ EXPECT_EQ(result.status, ge::GRAPH_SUCCESS);
	170	+ EXPECT_EQ(result.blockDim, 1U);
	171	+ EXPECT_EQ(result.tilingData.blockFactor, 910);
	172	+ EXPECT_EQ(result.tilingData.ubFactor, 19712);
	173	+}
	174	+
	175	+TEST(HardSigmoidTilingTest, Int32)
	176	+{
	177	+ const auto result = RunTilingCase({129}, ge::DT_INT32);
	178	+ EXPECT_EQ(result.status, ge::GRAPH_SUCCESS);
	179	+ EXPECT_EQ(result.blockDim, 1U);
	180	+ EXPECT_EQ(result.tilingData.blockFactor, 129);
	181	+ EXPECT_EQ(result.tilingData.ubFactor, 11840);
	182	+}
	183	+
	184	+TEST(HardSigmoidTilingTest, CustomAttributes)
	185	+{
	186	+ const auto result = RunTilingCase({1024}, ge::DT_FLOAT, 0.2f, 0.4f);
	187	+ EXPECT_EQ(result.status, ge::GRAPH_SUCCESS);
	188	+ EXPECT_EQ(result.blockDim, 1U);
	189	+ EXPECT_FLOAT_EQ(result.tilingData.alpha, 0.2f);
	190	+ EXPECT_FLOAT_EQ(result.tilingData.beta, 0.4f);
	191	+}
	192	+
	193	+TEST(HardSigmoidTilingTest, EmptyTensor)
	194	+{
	195	+ const auto result = RunTilingCase({0, 4}, ge::DT_FLOAT);
	196	+ EXPECT_EQ(result.status, ge::GRAPH_SUCCESS);
	197	+ EXPECT_EQ(result.blockDim, 1U);
	198	+ EXPECT_EQ(result.tilingData.totalElements, 0);
	199	+ EXPECT_EQ(result.tilingData.blockFactor, 0);
	200	+ EXPECT_EQ(result.tilingData.ubFactor, 0);
	201	+}
	202	+
	203	+TEST(HardSigmoidTilingTest, UnsupportedDtype) { EXPECT_EQ(RunTilingCase({128}, ge::DT_INT8).status, ge::GRAPH_FAILED); }
	204	+
	205	+TEST(HardSigmoidTilingTest, TwoMinimumCopyChunksUseTwoCores)
	206	+{
	207	+ const auto result = RunTilingCase({8192}, ge::DT_FLOAT);
	208	+ EXPECT_EQ(result.status, ge::GRAPH_SUCCESS);
	209	+ EXPECT_EQ(result.blockDim, 2U);
	210	+ EXPECT_EQ(result.tilingData.blockFactor, 4096);
	211	+}
	212	+
	213	+TEST(HardSigmoidTilingTest, JustBelowTwoMinimumCopyChunksUsesOneCore)
	214	+{
	215	+ const auto result = RunTilingCase({8191}, ge::DT_FLOAT);
	216	+ EXPECT_EQ(result.status, ge::GRAPH_SUCCESS);
	217	+ EXPECT_EQ(result.blockDim, 1U);
	218	+ EXPECT_EQ(result.tilingData.blockFactor, 8191);
	219	+}
	220	+
	221	+TEST(HardSigmoidTilingTest, RejectsNegativeDimensions)
	222	+{
	223	+ EXPECT_EQ(RunTilingCase({-1, -1}, ge::DT_FLOAT).status, ge::GRAPH_FAILED);
	224	+}

Aactivation/hard_sigmoid/tests/ut/op_host/test_hard_sigmoid_infershape.cpp+74-0

@@ -0,0 +1,74 @@
	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 <gtest/gtest.h>
	12	+#include "kernel_run_context_facker.h"
	13	+#include "register/op_impl_registry.h"
	14	+
	15	+namespace {
	16	+void RunInferShapeCase(gert::Shape inputShape)
	17	+{
	18	+ gert::Shape outputShape;
	19	+ auto holder = gert::InferShapeContextFaker()
	20	+ .NodeIoNum(1, 1)
	21	+ .IrInstanceNum({1, 1})
	22	+ .InputShapes({&inputShape})
	23	+ .OutputShapes({&outputShape})
	24	+ .Build();
	25	+
	26	+ auto* impl = gert::OpImplRegistry::GetInstance().GetOpImpl("HardSigmoid");
	27	+ ASSERT_NE(impl, nullptr);
	28	+ auto inferShape = impl->infer_shape;
	29	+ ASSERT_NE(inferShape, nullptr);
	30	+ auto* context = holder.GetContext<gert::InferShapeContext>();
	31	+ ASSERT_NE(context, nullptr);
	32	+ ASSERT_EQ(inferShape(context), ge::GRAPH_SUCCESS);
	33	+ const auto* inferredShape = context->GetOutputShape(0);
	34	+ ASSERT_NE(inferredShape, nullptr);
	35	+ EXPECT_EQ(*inferredShape, inputShape);
	36	+}
	37	+
	38	+void RunInferDataTypeCase(ge::DataType dtype)
	39	+{
	40	+ ge::DataType inputRef = dtype;
	41	+ ge::DataType outputRef = ge::DT_UNDEFINED;
	42	+ auto holder = gert::InferDataTypeContextFaker()
	43	+ .IrInputNum(1)
	44	+ .NodeIoNum(1, 1)
	45	+ .NodeInputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	46	+ .NodeOutputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	47	+ .InputDataTypes({&inputRef})
	48	+ .OutputDataTypes({&outputRef})
	49	+ .Build();
	50	+
	51	+ auto* impl = gert::OpImplRegistry::GetInstance().GetOpImpl("HardSigmoid");
	52	+ ASSERT_NE(impl, nullptr);
	53	+ auto inferDataType = impl->infer_datatype;
	54	+ ASSERT_NE(inferDataType, nullptr);
	55	+ auto* context = holder.GetContext<gert::InferDataTypeContext>();
	56	+ ASSERT_NE(context, nullptr);
	57	+ ASSERT_EQ(inferDataType(context), ge::GRAPH_SUCCESS);
	58	+ EXPECT_EQ(context->GetOutputDataType(0), dtype);
	59	+}
	60	+} // namespace
	61	+
	62	+TEST(HardSigmoidInferShapeTest, TensorShape) { RunInferShapeCase({2, 3, 5, 7}); }
	63	+
	64	+TEST(HardSigmoidInferShapeTest, ScalarShape) { RunInferShapeCase({}); }
	65	+
	66	+TEST(HardSigmoidInferShapeTest, EmptyShape) { RunInferShapeCase({0, 4}); }
	67	+
	68	+TEST(HardSigmoidInferShapeTest, InferDataType)
	69	+{
	70	+ RunInferDataTypeCase(ge::DT_FLOAT);
	71	+ RunInferDataTypeCase(ge::DT_FLOAT16);
	72	+ RunInferDataTypeCase(ge::DT_BF16);
	73	+ RunInferDataTypeCase(ge::DT_INT32);
	74	+}

Aactivation/hard_sigmoid/tests/ut/op_kernel/CMakeLists.txt+15-0

@@ -0,0 +1,15 @@
	1	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	2	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	3	+# CANN Open Software License Agreement Version 2.0 (the "License").
	4	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	5	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	6	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	7	+# See LICENSE in the root of the software repository for the full text of the License.
	8	+
	9	+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
	10	+ set(hard_sigmoid_tiling_files
	11	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/hard_sigmoid_tiling_arch35.cpp
	12	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/hard_sigmoid_infershape.cpp
	13	+ )
	14	+ AddOpTestCase(hard_sigmoid "ascend950pr_9599" "-DDTYPE_X=float" "${hard_sigmoid_tiling_files}")
	15	+endif()

Aactivation/hard_sigmoid/tests/ut/op_kernel/hard_sigmoid_data/compare_data.py+46-0

@@ -0,0 +1,46 @@
	1	+#!/usr/bin/env python3
	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	+import os
	11	+import sys
	12	+import numpy as np
	13	+
	14	+
	15	+def main():
	16	+ current_dir = os.path.dirname(os.path.realpath(__file__))
	17	+ diagnostic_path = os.path.join(current_dir, "compare_failure.txt")
	18	+ if os.path.exists(diagnostic_path):
	19	+ os.remove(diagnostic_path)
	20	+ golden = np.fromfile(os.path.join(current_dir, "golden.bin"), dtype=np.float32)
	21	+ output = np.fromfile(os.path.join(current_dir, "output.bin"), dtype=np.float32)
	22	+ if golden.shape != output.shape:
	23	+ message = f"shape mismatch: output={output.shape}, golden={golden.shape}"
	24	+ with open(diagnostic_path, "w", encoding="utf-8") as diagnostic:
	25	+ diagnostic.write(message + "\n")
	26	+ print(message)
	27	+ sys.exit(1)
	28	+ if not np.allclose(output, golden, rtol=1e-4, atol=1e-4, equal_nan=True):
	29	+ mismatch = np.flatnonzero(
	30	+ ~np.isclose(output, golden, rtol=1e-4, atol=1e-4, equal_nan=True)
	31	+ )
	32	+ if mismatch.size:
	33	+ index = int(mismatch[0])
	34	+ message = (
	35	+ f"mismatch_count={mismatch.size}; first mismatch at {index}: "
	36	+ f"output={output[index]}, golden={golden[index]}"
	37	+ )
	38	+ with open(diagnostic_path, "w", encoding="utf-8") as diagnostic:
	39	+ diagnostic.write(message + "\n")
	40	+ print(message)
	41	+ sys.exit(1)
	42	+ print("COMPARE DATA PASSED")
	43	+
	44	+
	45	+if __name__ == "__main__":
	46	+ main()

Aactivation/hard_sigmoid/tests/ut/op_kernel/hard_sigmoid_data/gen_data.py+45-0

@@ -0,0 +1,45 @@
	1	+#!/usr/bin/env python3
	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	+import sys
	11	+import numpy as np
	12	+
	13	+
	14	+def main():
	15	+ shape = tuple(int(value) for value in sys.argv[1].strip("()").split(",") if value)
	16	+ alpha = float(sys.argv[2])
	17	+ beta = float(sys.argv[3])
	18	+ values = np.array(
	19	+ [
	20	+ -np.inf,
	21	+ -8.0,
	22	+ -4.0,
	23	+ -3.0,
	24	+ -2.999,
	25	+ -1.0,
	26	+ 0.0,
	27	+ 1.0,
	28	+ 2.999,
	29	+ 3.0,
	30	+ 4.0,
	31	+ 8.0,
	32	+ np.inf,
	33	+ np.nan,
	34	+ ],
	35	+ dtype=np.float32,
	36	+ )
	37	+ count = int(np.prod(shape))
	38	+ x = np.resize(values, count).reshape(shape)
	39	+ golden = np.clip(alpha * x + beta, 0.0, 1.0).astype(np.float32)
	40	+ x.tofile("input_x.bin")
	41	+ golden.tofile("golden.bin")
	42	+
	43	+
	44	+if __name__ == "__main__":
	45	+ main()

Aactivation/hard_sigmoid/tests/ut/op_kernel/hard_sigmoid_tiling_def.h+19-0

@@ -0,0 +1,19 @@
	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 HARD_SIGMOID_TEST_TILING_DEF_H
	12	+#define HARD_SIGMOID_TEST_TILING_DEF_H
	13	+
	14	+#define REGISTER_TILING_DEFAULT(tilingStruct)
	15	+
	16	+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingPointer) \
	17	+ const tilingStruct& tilingData = reinterpret_cast<const tilingStruct>(tilingPointer)
	18	+
	19	+#endif // HARD_SIGMOID_TEST_TILING_DEF_H

Aactivation/hard_sigmoid/tests/ut/op_kernel/test_hard_sigmoid.cpp+110-0

@@ -0,0 +1,110 @@
	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 <gtest/gtest.h>
	12	+#include <memory>
	13	+#include "data_utils.h"
	14	+#include "kernel_ut_data_executor.h"
	15	+#include "kernel_ut_data_helper.h"
	16	+#include "tikicpulib.h"
	17	+
	18	+#include "../../../op_kernel/arch35/hard_sigmoid.cpp"
	19	+
	20	+namespace {
	21	+constexpr int64_t ELEMENT_COUNT = 256;
	22	+constexpr int64_t UB_ELEMENT_COUNT = 1024;
	23	+
	24	+struct GmDeleter {
	25	+ void operator()(uint8_t* ptr) const
	26	+ {
	27	+ if (ptr != nullptr) {
	28	+ AscendC::GmFree(ptr);
	29	+ }
	30	+ }
	31	+};
	32	+
	33	+using GmBuffer = std::unique_ptr<uint8_t, GmDeleter>;
	34	+
	35	+void RunHardSigmoidFloat(GM_ADDR input, GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling)
	36	+{
	37	+ (void)workspace;
	38	+ const auto* tilingData = reinterpret_cast<const HardSigmoidTilingData*>(tiling);
	39	+ HardSigmoidKernel<float> kernel;
	40	+ kernel.Init(input, output, tilingData);
	41	+ kernel.Process();
	42	+}
	43	+
	44	+void InitTiling(HardSigmoidTilingData* tilingData, float alpha, float beta)
	45	+{
	46	+ tilingData->totalElements = ELEMENT_COUNT;
	47	+ tilingData->blockFactor = ELEMENT_COUNT;
	48	+ tilingData->ubFactor = UB_ELEMENT_COUNT;
	49	+ tilingData->alpha = alpha;
	50	+ tilingData->beta = beta;
	51	+}
	52	+
	53	+void RunKernelCase(float alpha, float beta)
	54	+{
	55	+ size_t tensorBytes = ELEMENT_COUNT * sizeof(float);
	56	+ GmBuffer x(static_cast<uint8_t*>(AscendC::GmAlloc(tensorBytes)));
	57	+ GmBuffer y(static_cast<uint8_t*>(AscendC::GmAlloc(tensorBytes)));
	58	+ GmBuffer workspace(static_cast<uint8_t>(AscendC::GmAlloc(16 1024 * 1024)));
	59	+ GmBuffer tiling(static_cast<uint8_t*>(AscendC::GmAlloc(sizeof(HardSigmoidTilingData))));
	60	+ ASSERT_NE(x.get(), nullptr);
	61	+ ASSERT_NE(y.get(), nullptr);
	62	+ ASSERT_NE(workspace.get(), nullptr);
	63	+ ASSERT_NE(tiling.get(), nullptr);
	64	+
	65	+ kernel_ut::SetupTestEnvironment("activation/hard_sigmoid/tests/ut/op_kernel/hard_sigmoid_data",
	66	+ "hard_sigmoid_data");
	67	+ kernel_ut::RunGenData("./hard_sigmoid_data", {"'(256)'", std::to_string(alpha), std::to_string(beta)});
	68	+ const std::string path = kernel_ut::GetTestWorkDir();
	69	+ ReadFile(path + "/hard_sigmoid_data/input_x.bin", tensorBytes, x.get(), tensorBytes);
	70	+
	71	+ auto* tilingData = reinterpret_cast<HardSigmoidTilingData*>(tiling.get());
	72	+ InitTiling(tilingData, alpha, beta);
	73	+ ICPU_SET_TILING_KEY(0);
	74	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	75	+ ICPU_RUN_KF(RunHardSigmoidFloat, 1, x.get(), y.get(), workspace.get(), tiling.get());
	76	+ WriteFile(path + "/hard_sigmoid_data/output.bin", y.get(), tensorBytes);
	77	+
	78	+ ASSERT_TRUE(kernel_ut::RunCompareData("./hard_sigmoid_data", {}));
	79	+}
	80	+
	81	+void RunEmptyKernelCase()
	82	+{
	83	+ constexpr size_t bufferBytes = 32;
	84	+ GmBuffer x(static_cast<uint8_t*>(AscendC::GmAlloc(bufferBytes)));
	85	+ GmBuffer y(static_cast<uint8_t*>(AscendC::GmAlloc(bufferBytes)));
	86	+ GmBuffer workspace(static_cast<uint8_t*>(AscendC::GmAlloc(bufferBytes)));
	87	+ GmBuffer tiling(static_cast<uint8_t*>(AscendC::GmAlloc(sizeof(HardSigmoidTilingData))));
	88	+ ASSERT_NE(x.get(), nullptr);
	89	+ ASSERT_NE(y.get(), nullptr);
	90	+ ASSERT_NE(workspace.get(), nullptr);
	91	+ ASSERT_NE(tiling.get(), nullptr);
	92	+
	93	+ auto* tilingData = reinterpret_cast<HardSigmoidTilingData*>(tiling.get());
	94	+ *tilingData = HardSigmoidTilingData{};
	95	+ ICPU_SET_TILING_KEY(0);
	96	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	97	+ ICPU_RUN_KF(RunHardSigmoidFloat, 1, x.get(), y.get(), workspace.get(), tiling.get());
	98	+}
	99	+} // namespace
	100	+
	101	+class HardSigmoidKernelTest : public testing::Test {
	102	+protected:
	103	+ static void TearDownTestCase() { kernel_ut::CleanGeneratedBinFiles("./hard_sigmoid_data"); }
	104	+};
	105	+
	106	+TEST_F(HardSigmoidKernelTest, DefaultAttributes) { RunKernelCase(1.0f / 6.0f, 0.5f); }
	107	+
	108	+TEST_F(HardSigmoidKernelTest, CustomAttributes) { RunKernelCase(0.2f, 0.4f); }
	109	+
	110	+TEST_F(HardSigmoidKernelTest, EmptyTensor) { RunEmptyKernelCase(); }

Mdocs/zh/op_api_list.md+1-1

	@@ -268,7 +268,7 @@
268	\| [aclnnGroupNormSwishGrad](../../norm/group_norm_swish_grad/docs/aclnnGroupNormSwishGrad.md) \| [aclnnGroupNormSwish](../../norm/group_norm_swish/docs/aclnnGroupNormSwish.md)的反向操作。 \| 默认非确定性实现，支持配置开启 \| 默认确定性实现 \|	268	\| [aclnnGroupNormSwishGrad](../../norm/group_norm_swish_grad/docs/aclnnGroupNormSwishGrad.md) \| [aclnnGroupNormSwish](../../norm/group_norm_swish/docs/aclnnGroupNormSwish.md)的反向操作。 \| 默认非确定性实现，支持配置开启 \| 默认确定性实现 \|
269	\| [aclnnGroupQuant](../../quant/group_quant/docs/aclnnGroupQuant.md) \| 对输入x进行分组量化操作。 \| 默认确定性实现 \| - \|	269	\| [aclnnGroupQuant](../../quant/group_quant/docs/aclnnGroupQuant.md) \| 对输入x进行分组量化操作。 \| 默认确定性实现 \| - \|
270	\| [aclnnHardshrinkBackward](../../activation/hard_shrink_grad/docs/aclnnHardshrinkBackward.md) \| aclnnHardshrink计算反向传播的梯度gradInput。 \| 默认确定性实现 \| 默认确定性实现 \|	270	\| [aclnnHardshrinkBackward](../../activation/hard_shrink_grad/docs/aclnnHardshrinkBackward.md) \| aclnnHardshrink计算反向传播的梯度gradInput。 \| 默认确定性实现 \| 默认确定性实现 \|
271	-\| [aclnnHardsigmoid&aclnnInplaceHardsigmoid](../../activation/hard_sigmoid/docs/aclnnHardsigmoid&aclnnInplaceHardsigmoid.md) \| 激活函数变种，根据公式返回一个新的tensor。结果的形状与输入tensor相同。 \| 默认确定性实现 \| - \|	271	+\| [aclnnHardsigmoid&aclnnInplaceHardsigmoid](../../activation/hard_sigmoid/docs/aclnnHardsigmoid&aclnnInplaceHardsigmoid.md) \| 激活函数变种，根据公式返回一个新的tensor。结果的形状与输入tensor相同。 \| 默认确定性实现 \| 默认确定性实现 \|
272	\| [aclnnHardsigmoidBackward](../../activation/hard_sigmoid_grad/docs/aclnnHardsigmoidBackward.md) \| aclnnHardsigmoid的反向传播。 \| 默认确定性实现 \| 默认确定性实现 \|	272	\| [aclnnHardsigmoidBackward](../../activation/hard_sigmoid_grad/docs/aclnnHardsigmoidBackward.md) \| aclnnHardsigmoid的反向传播。 \| 默认确定性实现 \| 默认确定性实现 \|
273	\| [aclnnHardswishBackward](../../activation/hard_swish_grad/docs/aclnnHardswishBackward.md) \| aclnnHardswish的反向传播，完成张量self的梯度计算。 \| 默认确定性实现 \| 默认确定性实现 \|	273	\| [aclnnHardswishBackward](../../activation/hard_swish_grad/docs/aclnnHardswishBackward.md) \| aclnnHardswish的反向传播，完成张量self的梯度计算。 \| 默认确定性实现 \| 默认确定性实现 \|
274	\| [aclnnHardswishBackwardV2](../../activation/hard_swish_grad_v2/docs/aclnnHardswishBackwardV2.md) \| aclnnHardswish的反向传播，完成张量self的梯度计算。 \| 默认确定性实现 \| 默认确定性实现 \|	274	\| [aclnnHardswishBackwardV2](../../activation/hard_swish_grad_v2/docs/aclnnHardswishBackwardV2.md) \| aclnnHardswish的反向传播，完成张量self的梯度计算。 \| 默认确定性实现 \| 默认确定性实现 \|

Mdocs/zh/op_list.md+3-3

	@@ -290,10 +290,10 @@
290	<tr>	290	<tr>
291	<td>activation</td>	291	<td>activation</td>
292	<td><a href="../../activation/hard_sigmoid/README.md">hard_sigmoid</a></td>	292	<td><a href="../../activation/hard_sigmoid/README.md">hard_sigmoid</a></td>
293	- <td>✗</td>
294	- <td>✗</td>
295	<td>✓</td>	293	<td>✓</td>
296	- <td>✗</td>	294	+ <td>✓</td>
		295	+ <td>✓</td>
		296	+ <td>✓</td>
297	<td>AI Core</td>	297	<td>AI Core</td>
298	<td>激活函数，对输入张量self逐元素进行HardSigmoid变换，输出与输入shape相同的张量。</td>	298	<td>激活函数，对输入张量self逐元素进行HardSigmoid变换，输出与输入shape相同的张量。</td>
299	</tr>	299	</tr>