Star906
1825
代码
代码
Issues226
Pull Requests303
Actions
项目讨论
Wiki
分析
Star906
1825
已合并
[CANNBOT]ThresholdGradV2D适配Ascend950 Ascend实现 #6179
liuyi2025创建于 6月17日
[CANNBOT]ThresholdGradV2D适配Ascend950 Ascend实现 #6179
已合并
liuyi2025创建于 6月17日
从liuyi2025/ops-nn0617合入到cann/ops-nnmaster
共 19 个文件变更+1737-29
Mactivation/threshold_grad_v2_d/CMakeLists.txt+11-13
@@ -1,18 +1,16 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.1+# ----------------------------------------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 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
3# CANN Open Software License Agreement Version 2.0 (the "License").4# 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# 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+# 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# 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# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------------------------------------
10+# Generated By CANNBot
9 11 
10-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)12+# 设置算子定义时支持的芯片类型
11-if(NOT ENABLE_TEST AND NOT BENCHMARK)13+set(SUPPORT_COMPUTE_UNIT "ascend950")
12- list(REMOVE_ITEM CURRENT_DIRS tests)14+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
13-endif()15+set(SUPPORT_TILING_DIR "arch35")
14-foreach(SUB_DIR ${CURRENT_DIRS})16+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE threshold_grad_v2_d ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES relu_grad)
15- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16- add_subdirectory(${SUB_DIR})
17- endif()
18-endforeach()
Mactivation/threshold_grad_v2_d/README.md+45-2
@@ -1,3 +1,46 @@
1-# ThresholdGradV21+# ThresholdGradV2D
2 2 
3-本目录仅包含ThresholdGradV2算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。3+## 产品支持情况
4+ 
5+| 产品　　　　　　　　　　　　　　　　　　　　　　　　　　 | 是否支持 |
6+| :---------------------------------------------------------| :--------:|
7+| <term>Ascend 950PR/Ascend 950DT</term>　　　　　　　　　 | √　　　　|
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √　　　　|
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √　　　　|
10+| <term>Atlas 200I/500 A2 推理产品</term>　　　　　　　　　| ×　　　　|
11+| <term>Atlas 推理系列产品</term>　　　　　　　　　　　　　| √　　　　|
12+| <term>Atlas 训练系列产品</term>　　　　　　　　　　　　　| √　　　　|
13+ 
14+## 功能说明
15+ 
16+- 算子功能：完成threshold正向操作的反向传播梯度计算。当threshold==0时等价于ReluGrad。
17+ 
18+- 计算公式：
19+ 
20+ $$
21+ out_i =
22+ \begin{cases}
23+ gradOutput_i, \quad self_i > threshold\\
24+ 0, \quad self_i \leq threshold
25+ \end{cases}
26+ $$
27+ 
28+## 参数说明
29+ 
30+| 参数名 | 输入/输出 | 描述 | 数据类型 | 数据格式 |
31+| :--------- | :-------- | :----------------------------------------------------------- | :--------------------------------------------- | :------- |
32+| gradOutput | 输入 | 上游梯度，与self广播兼容。 | FLOAT16、FLOAT、BF16、INT32、INT8、UINT8 | ND |
33+| self | 输入 | 正向输入，决定门控掩码。 | FLOAT16、FLOAT、BF16、INT32、INT8、UINT8 | ND |
34+| threshold | 属性 | 阈值标量，OPTIONAL，默认1.0。threshold==0时等价ReluGrad。 | FLOAT | - |
35+| out | 输出 | 反向梯度，dtype与self一致，shape为广播后形状。 | FLOAT16、FLOAT、BF16、INT32、INT8、UINT8 | ND |
36+ 
37+## 约束说明
38+ 
39+- gradOutput、self、out三者 dtype一致。
40+- threshold==0走ReluGrad路径（regbase上额外支持INT64）。
41+ 
42+## 调用说明
43+ 
44+| 调用方式 | 调用样例 | 说明 |
45+| :--------- | :------------------------------------------------------ | :------------------------------ |
46+| aclnn调用 | [test_aclnn_threshold_backward.cpp](examples/test_aclnn_threshold_backward.cpp) | 通过aclnnThresholdBackward接口调用 |
Aactivation/threshold_grad_v2_d/examples/arch35/test_aclnn_threshold_backward.cpp+152-0
@@ -0,0 +1,152 @@
1+/**
2+ * Copyright (c) 2025 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+#include <iostream>
11+#include <vector>
12+#include "acl/acl.h"
13+#include "aclnnop/aclnn_threshold_backward.h"
14+ 
15+#define CHECK_RET(cond, return_expr) \
16+ do { \
17+ if (!(cond)) { \
18+ return_expr; \
19+ } \
20+ } while (0)
21+ 
22+#define LOG_PRINT(message, ...) \
23+ do { \
24+ printf(message, ##__VA_ARGS__); \
25+ } while (0)
26+ 
27+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
28+ int64_t shapeSize = 1;
29+ for (auto i : shape) {
30+ shapeSize *= i;
31+ }
32+ return shapeSize;
33+}
34+ 
35+int Init(int32_t deviceId, aclrtStream* stream) {
36+ // 固定写法，资源初始化
37+ auto ret = aclInit(nullptr);
38+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
39+ ret = aclrtSetDevice(deviceId);
40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
41+ ret = aclrtCreateStream(stream);
42+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
43+ return 0;
44+}
45+ 
46+template <typename T>
47+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
48+ aclDataType dataType, aclTensor** tensor) {
49+ auto size = GetShapeSize(shape) * sizeof(T);
50+ // 调用aclrtMalloc申请device侧内存
51+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
52+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
53+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
54+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
55+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
56+ 
57+ // 计算连续tensor的strides
58+ std::vector<int64_t> strides(shape.size(), 1);
59+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
60+ strides[i] = shape[i + 1] * strides[i + 1];
61+ }
62+ 
63+ // 调用aclCreateTensor接口创建aclTensor
64+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
65+ shape.data(), shape.size(), *deviceAddr);
66+ return 0;
67+}
68+ 
69+int main() {
70+ // 1. （固定写法）device/stream初始化，参考acl API手册
71+ // 根据自己的实际device填写deviceId
72+ int32_t deviceId = 0;
73+ aclrtStream stream;
74+ auto ret = Init(deviceId, &stream);
75+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
76+ 
77+ // 2. 构造输入与输出，需要根据API的接口自定义构造
78+ std::vector<int64_t> selfShape = {2, 2};
79+ std::vector<int64_t> gradOutputShape = {2, 2};
80+ std::vector<int64_t> outShape = {2, 2};
81+ void* selfDeviceAddr = nullptr;
82+ void* gradOutputDeviceAddr = nullptr;
83+ void* outDeviceAddr = nullptr;
84+ aclTensor* self = nullptr;
85+ aclTensor* gradOutput = nullptr;
86+ aclScalar* threshold = nullptr;
87+ aclTensor* out = nullptr;
88+ std::vector<float> selfHostData = {0.2, 1.2, 2.2, 3.2};
89+ std::vector<float> gradOutputHostData = {4.5, 4.4, 4.3, 4.2};
90+ std::vector<float> outHostData = {0.0, 0.0, 0.0, 0.0};
91+ float thresholdValue = 1.0f;
92+ // 创建self aclTensor
93+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
94+ CHECK_RET(ret == ACL_SUCCESS, return ret);
95+ // 创建gradOutput aclTensor
96+ ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT, &gradOutput);
97+ CHECK_RET(ret == ACL_SUCCESS, return ret);
98+ // 创建threshold aclScalar
99+ threshold = aclCreateScalar(&thresholdValue, aclDataType::ACL_FLOAT);
100+ CHECK_RET(threshold != nullptr, return ret);
101+ // 创建out aclTensor
102+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
103+ CHECK_RET(ret == ACL_SUCCESS, return ret);
104+ 
105+ // 3. 调用CANN算子库API，需要修改为具体的API名称
106+ uint64_t workspaceSize = 0;
107+ aclOpExecutor* executor;
108+ // 调用aclnnThresholdBackward第一段接口
109+ ret = aclnnThresholdBackwardGetWorkspaceSize(gradOutput, self, threshold, out, &workspaceSize, &executor);
110+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnThresholdBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
111+ // 根据第一段接口计算出的workspaceSize申请device内存
112+ void* workspaceAddr = nullptr;
113+ if (workspaceSize > 0) {
114+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
115+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
116+ }
117+ // 调用aclnnThresholdBackward第二段接口
118+ ret = aclnnThresholdBackward(workspaceAddr, workspaceSize, executor, stream);
119+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnThresholdBackward failed. ERROR: %d\n", ret); return ret);
120+ 
121+ // 4. （固定写法）同步等待任务执行结束
122+ ret = aclrtSynchronizeStream(stream);
123+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
124+ 
125+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
126+ auto size = GetShapeSize(outShape);
127+ std::vector<float> resultData(size, 0);
128+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
129+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
130+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
131+ for (int64_t i = 0; i < size; i++) {
132+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
133+ }
134+ 
135+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
136+ aclDestroyTensor(self);
137+ aclDestroyTensor(gradOutput);
138+ aclDestroyScalar(threshold);
139+ aclDestroyTensor(out);
140+ 
141+ // 7. 释放device资源，需要根据具体API的接口定义修改
142+ aclrtFree(selfDeviceAddr);
143+ aclrtFree(gradOutputDeviceAddr);
144+ aclrtFree(outDeviceAddr);
145+ if (workspaceSize > 0) {
146+ aclrtFree(workspaceAddr);
147+ }
148+ aclrtDestroyStream(stream);
149+ aclrtResetDevice(deviceId);
150+ aclFinalize();
151+ return 0;
152+}
Ractivation/threshold_grad_v2_d/op_host/op_api/aclnn_threshold_backward.cpp→activation/threshold_grad_v2_d/op_api/aclnn_threshold_backward.cpp+1-1
@@ -8,7 +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_threshold_backward.h"10#include "aclnn_threshold_backward.h"
11-#include "../../../relu_grad/op_api/relu_grad.h"11+#include "../../relu_grad/op_api/relu_grad.h"
12#include "threshold_grad.h"12#include "threshold_grad.h"
13#include "aclnn_kernels/contiguous.h"13#include "aclnn_kernels/contiguous.h"
14#include "aclnn_kernels/common/op_error_check.h"14#include "aclnn_kernels/common/op_error_check.h"
Ractivation/threshold_grad_v2_d/op_host/op_api/aclnn_threshold_backward.h→activation/threshold_grad_v2_d/op_api/aclnn_threshold_backward.h+5-5
@@ -38,12 +38,12 @@ extern "C" {
38 * G --> H[(out)]38 * G --> H[(out)]
39 * ```39 * ```
40 *40 *
41- * @param [in] gradOutput: npu device侧的aclTensor，数据类型支持FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8，shape需要与self一致。41+ * @param [in] gradOutput: npu device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8，shape需要与self一致。
42 * 支持非连续的Tensor，数据格式支持ND，且数据格式需要与self一致。42 * 支持非连续的Tensor，数据格式支持ND，且数据格式需要与self一致。
43- * @param [in] self: npu device侧的aclTensor，数据类型支持FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8。43+ * @param [in] self: npu device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8。
44 * 支持非连续的Tensor，数据格式支持ND。44 * 支持非连续的Tensor，数据格式支持ND。
45 * @param [in] threshold: host侧的aclScalar，数据类型需要可转换成self与other推导后的数据类型。45 * @param [in] threshold: host侧的aclScalar，数据类型需要可转换成self与other推导后的数据类型。
46- * @param [in] out: npu device侧的aclTensor，数据类型支持FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8，shape需要与self一致。46+ * @param [in] out: npu device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8，shape需要与self一致。
47 * 支持非连续的Tensor，数据格式支持ND，且数据格式需要与self一致。47 * 支持非连续的Tensor，数据格式支持ND，且数据格式需要与self一致。
48 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。48 * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
49 * @param [out] executor: 返回op执行器，包含算子计算流程。49 * @param [out] executor: 返回op执行器，包含算子计算流程。
@@ -53,10 +53,10 @@ ACLNN_API aclnnStatus aclnnThresholdBackwardGetWorkspaceSize(const aclTensor *gr
53 const aclScalar *threshold, aclTensor *out,53 const aclScalar *threshold, aclTensor *out,
54 uint64_t *workspaceSize, aclOpExecutor **executor);54 uint64_t *workspaceSize, aclOpExecutor **executor);
55/**55/**
56- * @brief aclnnAdd的第二段接口，用于执行计算。56+ * @brief aclnnThresholdBackward的第二段接口，用于执行计算。
57 *57 *
58 * @param [in] workspace: 在npu device侧申请的workspace内存起址。58 * @param [in] workspace: 在npu device侧申请的workspace内存起址。
59- * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnAddGetWorkspaceSize获取。59+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnThresholdBackwardGetWorkspaceSize获取。
60 * @param [in] stream: acl stream流。60 * @param [in] stream: acl stream流。
61 * @param [in] executor: op执行器，包含了算子计算流程。61 * @param [in] executor: op执行器，包含了算子计算流程。
62 * @return aclnnStatus: 返回状态码。62 * @return aclnnStatus: 返回状态码。
Ractivation/threshold_grad_v2_d/op_host/op_api/threshold_grad.cpp→activation/threshold_grad_v2_d/op_api/threshold_grad.cpp+0-0
文件重命名但无更改。
Ractivation/threshold_grad_v2_d/op_host/op_api/threshold_grad.h→activation/threshold_grad_v2_d/op_api/threshold_grad.h+0-0
文件重命名但无更改。
Aactivation/threshold_grad_v2_d/op_host/arch35/threshold_grad_v2_d_tiling_arch35.cpp+301-0
@@ -0,0 +1,301 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file threshold_grad_v2_d_tiling_arch35.cpp
15+ * \brief ThresholdGradV2D Tiling — arch35 (RegBase Broadcast 范式)
16+ * PadAndSqueeze → FindSplitAxis → MultiCoreSplit, 无 workspace
17+ * 目标架构: DAV_3510 (ascend950)
18+ */
19+#include "register/op_def_registry.h"
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/threshold_grad_v2_d_tiling_struct.h"
24+#include "../../op_kernel/arch35/threshold_grad_v2_d_struct.h"
25+ 
26+#include <algorithm>
27+#include <sstream>
28+#include <vector>
29+ 
30+namespace optiling {
31+ 
32+namespace tgv2d {
33+ 
34+constexpr int64_t THRESHOLD_ATTR_IDX = 0;
35+constexpr float DEFAULT_THRESHOLD = 1.0f;
36+ 
37+// 5.0 CheckBroadcastShape — 补 1 后逐维检查 broadcast 兼容性：同维非 1 大小必须一致
38+static bool CheckBroadcastShape(
39+ const std::vector<std::vector<int64_t>>& padded_in,
40+ const std::vector<std::vector<int64_t>>& padded_out,
41+ int64_t max_rank)
42+{
43+ for (int64_t d = 0; d < max_rank; d++) {
44+ int64_t ref = -1;
45+ for (size_t i = 0; i < padded_in.size(); i++) {
46+ if (padded_in[i][d] != 1) {
47+ if (ref == -1) ref = padded_in[i][d];
48+ else if (padded_in[i][d] != ref) {
49+ OP_LOGE("PadAndSqueeze", "dim %d broadcast incompatible: input[%d] size %d != %d",
50+ (int)d, (int)i, (int)padded_in[i][d], (int)ref);
51+ return false;
52+ }
53+ }
54+ }
55+ for (size_t i = 0; i < padded_out.size(); i++) {
56+ if (padded_out[i][d] != 1) {
57+ if (ref == -1) ref = padded_out[i][d];
58+ else if (padded_out[i][d] != ref) {
59+ OP_LOGE("PadAndSqueeze", "dim %d broadcast incompatible: output[%d] size %d != %d",
60+ (int)d, (int)i, (int)padded_out[i][d], (int)ref);
61+ return false;
62+ }
63+ }
64+ }
65+ }
66+ return true;
67+}
68+ 
69+bool PadAndSqueeze(
70+ const std::vector<std::vector<int64_t>>& input_shapes,
71+ const std::vector<std::vector<int64_t>>& output_shapes,
72+ std::vector<int64_t>& maximum_bro_shape,
73+ std::vector<std::vector<int64_t>>& normal_input_shapes,
74+ std::vector<std::vector<int64_t>>& normal_output_shapes)
75+{
76+ int64_t num_inputs = (int64_t)input_shapes.size();
77+ int64_t num_outputs = (int64_t)output_shapes.size();
78+ int64_t max_rank = 0;
79+ for (auto& s : input_shapes) max_rank = std::max(max_rank, (int64_t)s.size());
80+ for (auto& s : output_shapes) max_rank = std::max(max_rank, (int64_t)s.size());
81+ auto pad = [&](const std::vector<int64_t>& s) {
82+ std::vector<int64_t> p;
83+ p.assign(max_rank - (int64_t)s.size(), 1);
84+ p.insert(p.end(), s.begin(), s.end());
85+ return p;
86+ };
87+ std::vector<std::vector<int64_t>> padded_in(num_inputs), padded_out(num_outputs);
88+ for (int64_t i = 0; i < num_inputs; i++) padded_in[i] = pad(input_shapes[i]);
89+ for (int64_t i = 0; i < num_outputs; i++) padded_out[i] = pad(output_shapes[i]);
90+ // 校验 — 补 1 后逐维 broadcast 兼容性，不兼容返回 false（调用方转 GRAPH_FAILED）
91+ if (!CheckBroadcastShape(padded_in, padded_out, max_rank)) return false;
92+ maximum_bro_shape.clear();
93+ normal_input_shapes.assign(num_inputs, std::vector<int64_t>());
94+ normal_output_shapes.assign(num_outputs, std::vector<int64_t>());
95+ for (int64_t d = 0; d < max_rank; d++) {
96+ bool all_one = true;
97+ int64_t max_dim = 0;
98+ for (int64_t i = 0; i < num_inputs; i++) { if (padded_in[i][d] != 1) all_one = false; max_dim = std::max(max_dim, padded_in[i][d]); }
99+ for (int64_t i = 0; i < num_outputs; i++) { if (padded_out[i][d] != 1) all_one = false; max_dim = std::max(max_dim, padded_out[i][d]); }
100+ if (!all_one) {
101+ maximum_bro_shape.push_back(max_dim);
102+ for (int64_t i = 0; i < num_inputs; i++) normal_input_shapes[i].push_back(padded_in[i][d]);
103+ for (int64_t i = 0; i < num_outputs; i++) normal_output_shapes[i].push_back(padded_out[i][d]);
104+ }
105+ }
106+ if (maximum_bro_shape.empty()) {
107+ maximum_bro_shape.push_back(1);
108+ for (int64_t i = 0; i < num_inputs; i++) normal_input_shapes[i].push_back(1);
109+ for (int64_t i = 0; i < num_outputs; i++) normal_output_shapes[i].push_back(1);
110+ }
111+ return true;
112+}
113+ 
114+bool FindSplitAxis(const std::vector<int64_t>& max_bro_shape,
115+ int64_t ub_per_core, int64_t phys_nodes, SplitResult& out)
116+{
117+ if (phys_nodes <= 0) { return false; }
118+ int64_t per_buf_bytes = (ub_per_core / phys_nodes) & ~31LL;
119+ int64_t per_buf_elems = per_buf_bytes / 4; // 统一 FP32
120+ int64_t rank = (int64_t)max_bro_shape.size();
121+ int64_t inner = 1;
122+ for (int64_t k = rank - 1; k >= 0; k--) {
123+ if (max_bro_shape[k] * inner > per_buf_elems) {
124+ out.a_i = per_buf_elems / inner;
125+ out.a_o = (max_bro_shape[k] + out.a_i - 1) / out.a_i;
126+ int64_t rem = max_bro_shape[k] % out.a_i;
127+ out.a_i_tail = (rem == 0) ? out.a_i : rem;
128+ out.axis = k;
129+ return true;
130+ }
131+ if (k == 0) { out.axis = 0; out.a_i = max_bro_shape[0]; out.a_o = 1; out.a_i_tail = max_bro_shape[0]; return true; }
132+ inner *= max_bro_shape[k];
133+ }
134+ return true;
135+}
136+ 
137+bool MultiCoreSplit(const std::vector<int64_t>& max_bro_shape,
138+ const SplitResult& ub_split, int64_t max_cores, MultiCoreResult& out)
139+{
140+ int64_t k = ub_split.axis, outer_prod = 1;
141+ for (int64_t j = 0; j < k; j++) outer_prod *= max_bro_shape[j];
142+ out.total_tiles = outer_prod * ub_split.a_o;
143+ out.num_cores = (out.total_tiles < max_cores) ? out.total_tiles : max_cores;
144+ out.tiles_main = out.total_tiles / out.num_cores;
145+ out.cores_tail = out.total_tiles % out.num_cores;
146+ return true;
147+}
148+ 
149+bool PrecomputeStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides) {
150+ int64_t rank = (int64_t)s.size();
151+ strides.assign(rank, 0);
152+ for (int64_t d = rank - 1; d >= 0; d--) {
153+ if (s[d] == 1) { strides[d] = 0; continue; }
154+ int64_t prod = 1;
155+ for (int64_t j = d + 1; j < rank; j++) prod *= s[j];
156+ strides[d] = prod;
157+ }
158+ return true;
159+}
160+ 
161+} // namespace tgv2d
162+ 
163+struct ThresholdGradV2DCompileInfo { uint64_t coreNum; uint64_t ubSize; };
164+ 
165+static std::string Arr2String(const int64_t* arr, int64_t n) {
166+ std::ostringstream oss; oss << "[";
167+ if (n > 0) { for (int64_t i = 0; i < n - 1; ++i) oss << arr[i] << ","; oss << arr[n - 1]; }
168+ oss << "]"; return oss.str();
169+}
170+ 
171+class ThresholdGradV2DTilingImpl {
172+public:
173+ explicit ThresholdGradV2DTilingImpl(gert::TilingContext* ctx) : ctx_(ctx) {}
174+ 
175+ ge::graphStatus RunTiling() {
176+ ge::graphStatus ret = GetShapeInfo();
177+ if (ret != ge::GRAPH_SUCCESS) return ret;
178+ int64_t total_out = 1;
179+ for (auto d : raw_output_shapes_[0]) total_out *= d;
180+ if (total_out == 0) { ctx_->SetBlockDim(1); ctx_->SetTilingKey(GET_TPL_TILING_KEY(THRESHOLD_GRAD_V2_D_RANK_4)); return ge::GRAPH_SUCCESS; }
181+ if (rank_ <= 4) { ret = DoTilingAndSet<4>(); ctx_->SetTilingKey(GET_TPL_TILING_KEY(THRESHOLD_GRAD_V2_D_RANK_4)); }
182+ else { ret = DoTilingAndSet<8>(); ctx_->SetTilingKey(GET_TPL_TILING_KEY(THRESHOLD_GRAD_V2_D_RANK_8)); }
183+ return ret;
184+ }
185+ 
186+private:
187+ ge::graphStatus GetShapeInfo() {
188+ fe::PlatFormInfos* platformInfo = ctx_->GetPlatformInfo();
189+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, platformInfo);
190+ auto ap = platform_ascendc::PlatformAscendC(platformInfo);
191+ coreNum_ = ap.GetCoreNumAiv();
192+ ap.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize_);
193+ 
194+ for (size_t i = 0; i < ctx_->GetComputeNodeInfo()->GetInputsNum(); ++i) {
195+ auto shape = ctx_->GetInputShape(i);
196+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, shape);
197+ std::vector<int64_t> dims; gert::Shape s = shape->GetStorageShape();
198+ for (size_t d = 0; d < s.GetDimNum(); ++d) dims.push_back(s.GetDim(d));
199+ if (dims.empty()) dims.push_back(1); // rank0 标量当 1
200+ raw_input_shapes_.push_back(dims);
201+ }
202+ for (size_t i = 0; i < ctx_->GetComputeNodeInfo()->GetOutputsNum(); ++i) {
203+ auto shape = ctx_->GetOutputShape(i);
204+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, shape);
205+ std::vector<int64_t> dims; gert::Shape s = shape->GetStorageShape();
206+ for (size_t d = 0; d < s.GetDimNum(); ++d) dims.push_back(s.GetDim(d));
207+ if (dims.empty()) dims.push_back(1);
208+ raw_output_shapes_.push_back(dims);
209+ }
210+ auto inputDesc = ctx_->GetInputDesc(0);
211+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, inputDesc);
212+ ge::DataType dtype = inputDesc->GetDataType();
213+ const std::set<ge::DataType> sup = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8};
214+ OP_CHECK_IF(sup.count(dtype) == 0, OP_LOGE(ctx_->GetNodeName(), "Unsupported dtype"), return ge::GRAPH_FAILED);
215+ 
216+ threshold_ = tgv2d::DEFAULT_THRESHOLD;
217+ const auto* attrs = ctx_->GetAttrs();
218+ if (attrs != nullptr) { const float* th = attrs->GetFloat(tgv2d::THRESHOLD_ATTR_IDX); if (th) threshold_ = *th; }
219+ 
220+ OP_CHECK_IF(!tgv2d::PadAndSqueeze(raw_input_shapes_, raw_output_shapes_, max_bro_shape_, normal_input_shapes_, normal_output_shapes_),
221+ OP_LOGE(ctx_->GetNodeName(), "broadcast shape incompatible"), return ge::GRAPH_FAILED);
222+ rank_ = (int64_t)max_bro_shape_.size();
223+ OP_CHECK_IF(rank_ > 8, OP_LOGE(ctx_->GetNodeName(), "rank>8 unsupported"), return ge::GRAPH_FAILED);
224+ return ge::GRAPH_SUCCESS;
225+ }
226+ 
227+ template<int64_t R>
228+ ge::graphStatus DoTilingAndSet() {
229+ auto* tiling = ctx_->GetTilingData<ThresholdGradV2DTilingData<R>>();
230+ OP_CHECK_NULL_WITH_CONTEXT(ctx_, tiling);
231+ int64_t ub_per_core = (int64_t)ubSize_;
232+ int64_t per_buf_bytes = (ub_per_core / kPhysNodes) & ~31LL;
233+ tgv2d::FindSplitAxis(max_bro_shape_, ub_per_core, kPhysNodes, tiling->split);
234+ tgv2d::MultiCoreSplit(max_bro_shape_, tiling->split, (int64_t)coreNum_, tiling->multicore);
235+ tiling->per_buf_bytes = per_buf_bytes;
236+ tiling->per_buf_elems = per_buf_bytes / 4;
237+ 
238+ int64_t num_in = (int64_t)normal_input_shapes_.size(), num_out = (int64_t)normal_output_shapes_.size();
239+ std::vector<std::vector<int64_t>> in_strides(num_in), out_strides(num_out);
240+ for (int64_t i = 0; i < num_in; i++) tgv2d::PrecomputeStrides(normal_input_shapes_[i], in_strides[i]);
241+ for (int64_t i = 0; i < num_out; i++) tgv2d::PrecomputeStrides(normal_output_shapes_[i], out_strides[i]);
242+ 
243+ tiling->rank = rank_;
244+ tiling->threshold = threshold_;
245+ int64_t delta = R - rank_;
246+ for (int64_t d = 0; d < delta; d++) tiling->max_bro_shape[d] = 1;
247+ for (int64_t d = 0; d < rank_; d++) tiling->max_bro_shape[d + delta] = max_bro_shape_[d];
248+ tiling->split.axis += delta;
249+ tiling->num_inputs = num_in; tiling->num_outputs = num_out;
250+ for (int64_t i = 0; i < num_in; i++) {
251+ for (int64_t d = 0; d < delta; d++) { tiling->input_shapes[i][d] = 1; tiling->input_strides[i][d] = 0; }
252+ for (int64_t d = 0; d < rank_; d++) { tiling->input_shapes[i][d + delta] = normal_input_shapes_[i][d]; tiling->input_strides[i][d + delta] = in_strides[i][d]; }
253+ }
254+ for (int64_t i = num_in; i < kMaxInputSlots; i++) for (int64_t d = 0; d < R; d++) { tiling->input_shapes[i][d] = 1; tiling->input_strides[i][d] = 0; }
255+ for (int64_t i = 0; i < num_out; i++) {
256+ for (int64_t d = 0; d < delta; d++) { tiling->output_shapes[i][d] = 1; tiling->output_strides[i][d] = 0; }
257+ for (int64_t d = 0; d < rank_; d++) { tiling->output_shapes[i][d + delta] = normal_output_shapes_[i][d]; tiling->output_strides[i][d + delta] = out_strides[i][d]; }
258+ }
259+ for (int64_t i = num_out; i < kMaxOutputSlots; i++) for (int64_t d = 0; d < R; d++) { tiling->output_shapes[i][d] = 1; tiling->output_strides[i][d] = 0; }
260+ 
261+ ctx_->SetBlockDim(tiling->multicore.num_cores);
262+ OP_LOGI(ctx_->GetNodeName(), "TGV2D per_buf=%ld rank=%ld->R=%d mb=%s split(ax=%ld ai=%ld ao=%ld tail=%ld) mc(c=%ld t=%ld m=%ld ct=%ld)",
263+ tiling->per_buf_bytes, rank_, (int)R, Arr2String(tiling->max_bro_shape, R).c_str(),
264+ tiling->split.axis, tiling->split.a_i, tiling->split.a_o, tiling->split.a_i_tail,
265+ tiling->multicore.num_cores, tiling->multicore.total_tiles, tiling->multicore.tiles_main, tiling->multicore.cores_tail);
266+ return ge::GRAPH_SUCCESS;
267+ }
268+ 
269+ gert::TilingContext* ctx_;
270+ std::vector<std::vector<int64_t>> raw_input_shapes_, raw_output_shapes_;
271+ std::vector<int64_t> max_bro_shape_;
272+ std::vector<std::vector<int64_t>> normal_input_shapes_, normal_output_shapes_;
273+ int64_t rank_ = 0; float threshold_ = 1.0f;
274+ uint64_t coreNum_ = 0, ubSize_ = 0;
275+};
276+ 
277+static ge::graphStatus ThresholdGradV2DTilingFunc(gert::TilingContext* context) {
278+ ThresholdGradV2DTilingImpl impl(context);
279+ auto ret = impl.RunTiling();
280+ if (ret != ge::GRAPH_SUCCESS) return ret;
281+ size_t* ws = context->GetWorkspaceSizes(1);
282+ ws[0] = 0;
283+ return ge::GRAPH_SUCCESS;
284+}
285+ 
286+static ge::graphStatus TilingPrepareForThresholdGradV2D(gert::TilingParseContext* context) {
287+ fe::PlatFormInfos* platformInfo = context->GetPlatformInfo();
288+ auto compileInfo = context->GetCompiledInfo<ThresholdGradV2DCompileInfo>();
289+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
290+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
291+ auto ap = platform_ascendc::PlatformAscendC(platformInfo);
292+ compileInfo->coreNum = ap.GetCoreNumAiv();
293+ ap.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfo->ubSize);
294+ return ge::GRAPH_SUCCESS;
295+}
296+ 
297+IMPL_OP_OPTILING(ThresholdGradV2D)
298+ .Tiling(ThresholdGradV2DTilingFunc)
299+ .TilingParse<ThresholdGradV2DCompileInfo>(TilingPrepareForThresholdGradV2D);
300+ 
301+} // namespace optiling
Aactivation/threshold_grad_v2_d/op_host/threshold_grad_v2_d_def.cpp+59-0
@@ -0,0 +1,59 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file threshold_grad_v2_d_def.cpp
15+ * \brief ThresholdGradV2D 算子定义：gradOutput/self -> out，fp16/fp32/bf16/int32/int8/uint8
16+ * 属性 threshold(Float, 默认 1.0)。out = self>threshold ? gradOutput : 0
17+ */
18+#include "register/op_def_registry.h"
19+ 
20+namespace ops {
21+class ThresholdGradV2D : public OpDef {
22+public:
23+ explicit ThresholdGradV2D(const char* name) : OpDef(name)
24+ {
25+ this->Input("gradOutput")
26+ .ParamType(REQUIRED)
27+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
28+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
29+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
30+ .AutoContiguous();
31+ this->Input("self")
32+ .ParamType(REQUIRED)
33+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
34+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
35+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
36+ .AutoContiguous();
37+ this->Output("out")
38+ .ParamType(REQUIRED)
39+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
40+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
41+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
42+ .AutoContiguous();
43+ 
44+ this->Attr("threshold").AttrType(OPTIONAL).Float(1.0);
45+ 
46+ // 目标芯片为 Ascend950PR/DT（arch35）。6 dtype（含 int8/uint8）依赖 arch35 矢量 ISA，arch22 不支持。
47+ OpAICoreConfig aiCoreConfig;
48+ aiCoreConfig.DynamicCompileStaticFlag(true)
49+ .DynamicFormatFlag(false)
50+ .DynamicRankSupportFlag(true)
51+ .DynamicShapeSupportFlag(true)
52+ .NeedCheckSupportFlag(false)
53+ .PrecisionReduceFlag(true)
54+ .ExtendCfgInfo("opFile.value", "threshold_grad_v2_d");
55+ this->AICore().AddConfig("ascend950", aiCoreConfig);
56+ }
57+};
58+OP_ADD(ThresholdGradV2D);
59+} // namespace ops
Aactivation/threshold_grad_v2_d/op_host/threshold_grad_v2_d_infershape.cpp+59-0
@@ -0,0 +1,59 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file threshold_grad_v2_d_infershape.cpp
15+ * \brief ThresholdGradV2D 形状/类型推导: out = broadcast(gradOutput, self), dtype=gradOutput
16+ */
17+ 
18+#include "register/op_impl_registry.h"
19+#include "exe_graph/runtime/infer_shape_context.h"
20+#include "exe_graph/runtime/infer_datatype_context.h"
21+#include "op_common/log/log.h"
22+ 
23+using namespace ge;
24+ 
25+namespace ops {
26+ 
27+// numpy 广播：取最大 rank，对齐到右，逐维取 max（1 维可广播）
28+static ge::graphStatus InferShape4ThresholdGradV2D(gert::InferShapeContext* context)
29+{
30+ const gert::Shape* g = context->GetInputShape(0);
31+ const gert::Shape* s = context->GetInputShape(1);
32+ OP_CHECK_NULL_WITH_CONTEXT(context, g);
33+ OP_CHECK_NULL_WITH_CONTEXT(context, s);
34+ gert::Shape* out = context->GetOutputShape(0);
35+ OP_CHECK_NULL_WITH_CONTEXT(context, out);
36+ 
37+ size_t gn = g->GetDimNum();
38+ size_t sn = s->GetDimNum();
39+ size_t rn = gn > sn ? gn : sn;
40+ out->SetDimNum(rn);
41+ for (size_t i = 0; i < rn; ++i) {
42+ int64_t gd = (i < rn - gn) ? 1 : g->GetDim(i - (rn - gn));
43+ int64_t sd = (i < rn - sn) ? 1 : s->GetDim(i - (rn - sn));
44+ out->SetDim(i, gd > sd ? gd : sd);
45+ }
46+ return ge::GRAPH_SUCCESS;
47+}
48+ 
49+static ge::graphStatus InferDataType4ThresholdGradV2D(gert::InferDataTypeContext* context)
50+{
51+ context->SetOutputDataType(0, context->GetInputDataType(0));
52+ return ge::GRAPH_SUCCESS;
53+}
54+ 
55+IMPL_OP_INFERSHAPE(ThresholdGradV2D)
56+ .InferShape(InferShape4ThresholdGradV2D)
57+ .InferDataType(InferDataType4ThresholdGradV2D);
58+ 
59+} // namespace ops
Aactivation/threshold_grad_v2_d/op_kernel/arch35/threshold_grad_v2_d.h+352-0
@@ -0,0 +1,352 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file threshold_grad_v2_d.h
15+ * \brief ThresholdGradV2D Kernel — ThresholdGradV2DKernel<T, RANK>
16+ * RegBase Broadcast 范式, 基于 DESIGN.md §3.4
17+ * 公式: out = (self > threshold) ? gradOutput : 0
18+ * - int32/int8/uint8/fp32: 原生零 cast, Compares<T> + Select<T> 直接在 T 寄存器
19+ * - fp16/bf16: 仅比较项升精度 — self 升 fp32 与 fp32 threshold 在 fp32 域比 self<=th 出 mask,
20+ * grad 仍原生 T (UNPACK_B16) 经 Select<T>(零,grad) 搬出, 不 cast grad, 不全量升精度。
21+ * 比较域与 golden(fp32) 对齐, 避免 threshold 量化成 fp16 造成边界回归。NaN<=th false 透传。
22+ * 目标架构: DAV_3510 (arch35)
23+ */
24+#pragma once
25+#include "kernel_operator.h"
26+#include "threshold_grad_v2_d_tiling_struct.h"
27+#include "threshold_grad_v2_d_struct.h"
28+ 
29+// ============================================================
30+// VF: dst = (self > th) ? grad : 0 (CompareScalar GT + Select, mask 入寄存器零 UB)
31+// int32/int8/uint8/fp32: T 原生比较 — 无量化, 比较域即 golden 域, 零 cast
32+// gateMode 1/2: 恒透传/恒置0, grad 原生搬, 零 cast
33+// ============================================================
34+template <typename T>
35+__simd_vf__ inline void ThresholdGradVF(
36+ __ubuf__ T* dstAddr, __ubuf__ T* selfAddr, __ubuf__ T* gradAddr,
37+ T threshold, int gateMode, uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes)
38+{
39+ AscendC::Reg::RegTensor<T> selfReg, gradReg, zeroReg, dstReg;
40+ AscendC::Reg::MaskReg cmpMask, mask;
41+ AscendC::Reg::AddrReg aReg;
42+ AscendC::Reg::Duplicate(zeroReg, static_cast<T>(0));
43+ for (uint16_t i = 0; i < repeatTimes; ++i) {
44+ aReg = AscendC::Reg::CreateAddrReg<T>(i, oneRepeatSize);
45+ uint32_t rem = count - (uint32_t)i * oneRepeatSize;
46+ mask = AscendC::Reg::UpdateMask<T>(rem);
47+ AscendC::Reg::LoadAlign(gradReg, gradAddr, aReg);
48+ if (gateMode == 1) { // 恒透传 grad
49+ dstReg = gradReg;
50+ } else if (gateMode == 2) { // 恒置 0
51+ dstReg = zeroReg;
52+ } else { // PyTorch: self<=th -> 0 ; 否则(含 NaN, NaN<=th 为 false) grad
53+ AscendC::Reg::LoadAlign(selfReg, selfAddr, aReg);
54+ AscendC::Reg::Compares<T, AscendC::CMPMODE::LE>(cmpMask, selfReg, threshold, mask);
55+ AscendC::Reg::Select<T>(dstReg, zeroReg, gradReg, cmpMask);
56+ }
57+ AscendC::Reg::StoreAlign(dstAddr, dstReg, aReg, mask);
58+ }
59+}
60+ 
61+// ============================================================
62+// VF(fp16/bf16): 仅比较项升 fp32 与 golden 比较域对齐 (self 升 fp32, 与 fp32 threshold 在 fp32 域比 self<=th)。
63+// grad 仍原生 T, Select<T>(零,grad) 原生搬 —— 不 cast grad, 不全量升精度。NaN<=th 为 false 透传。
64+// 按 fp32 VL 步进 (UNPACK_B16: 一拍 128 个 b16 -> 128 个 fp32, mask 与 b16 lane 1:1)。
65+// ============================================================
66+template <typename T>
67+__simd_vf__ inline void ThresholdGradVFb16(
68+ __ubuf__ T* dstAddr, __ubuf__ T* selfAddr, __ubuf__ T* gradAddr,
69+ float threshold, uint32_t count, uint32_t vlF, uint16_t repeatTimes)
70+{
71+ static constexpr AscendC::Reg::CastTrait kToFp32{
72+ AscendC::Reg::RegLayout::ZERO, AscendC::Reg::SatMode::UNKNOWN,
73+ AscendC::Reg::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_NONE };
74+ AscendC::Reg::RegTensor<T> selfReg, gradReg, zeroReg, dstReg;
75+ AscendC::Reg::RegTensor<float> selfF;
76+ AscendC::Reg::MaskReg cmpMask, mask;
77+ AscendC::Reg::Duplicate(zeroReg, static_cast<T>(0));
78+ for (uint16_t i = 0; i < repeatTimes; ++i) {
79+ int32_t off = static_cast<int32_t>(i) * static_cast<int32_t>(vlF);
80+ uint32_t rem = count - (uint32_t)i * vlF;
81+ mask = AscendC::Reg::UpdateMask<float>(rem);
82+ AscendC::Reg::LoadAlign<T, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(selfReg, selfAddr + off);
83+ AscendC::Reg::Cast<float, T, kToFp32>(selfF, selfReg, mask);
84+ AscendC::Reg::Compares<float, AscendC::CMPMODE::LE>(cmpMask, selfF, threshold, mask);
85+ AscendC::Reg::LoadAlign<T, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(gradReg, gradAddr + off);
86+ AscendC::Reg::Select<T>(dstReg, zeroReg, gradReg, cmpMask);
87+ AscendC::Reg::StoreAlign<T, AscendC::Reg::StoreDist::DIST_PACK_B32>(dstAddr + off, dstReg, mask);
88+ }
89+}
90+ 
91+// ============================================================
92+// Kernel 侧辅助函数 (int64_t* 版本)
93+// ============================================================
94+__aicore__ inline void GetCoreRange(int64_t core_id, int64_t tiles_main, int64_t cores_tail,
95+ int64_t& start, int64_t& end)
96+{
97+ if (core_id < cores_tail) {
98+ start = core_id * (tiles_main + 1);
99+ end = start + tiles_main + 1;
100+ } else {
101+ start = cores_tail * (tiles_main + 1) + (core_id - cores_tail) * tiles_main;
102+ end = start + tiles_main;
103+ }
104+}
105+ 
106+__aicore__ inline int64_t GetUBSplitRange(int64_t a_o_off, int64_t a_o, int64_t a_i, int64_t a_i_tail)
107+{
108+ return (a_o_off == a_o - 1) ? a_i_tail : a_i;
109+}
110+ 
111+__aicore__ inline bool FlatToEffectiveCoord(int64_t flat, const int64_t* max_bro_shape,
112+ int64_t rank, int64_t split_axis, int64_t a_i, int64_t a_o, int64_t* eff_coord)
113+{
114+ for (int64_t d = 0; d < rank; d++) eff_coord[d] = 0;
115+ if (a_o <= 0) { return false; }
116+ int64_t a_o_off = flat % a_o;
117+ int64_t outer = flat / a_o;
118+ for (int64_t d = split_axis - 1; d >= 0; d--) {
119+ eff_coord[d] = outer % max_bro_shape[d];
120+ outer /= max_bro_shape[d];
121+ }
122+ eff_coord[split_axis] = a_o_off * a_i;
123+ return true;
124+}
125+ 
126+__aicore__ inline int64_t CalcOffset(const int64_t* eff_coord, const int64_t* strides, int64_t rank)
127+{
128+ int64_t offset = 0;
129+ for (int64_t d = 0; d < rank; d++) offset += eff_coord[d] * strides[d];
130+ return offset;
131+}
132+ 
133+// ============================================================
134+// ThresholdGradV2DKernel<T, RANK>
135+// ============================================================
136+template <typename T, int64_t RANK>
137+class ThresholdGradV2DKernel {
138+ static constexpr int64_t ND = (RANK <= 5) ? RANK : 5;
139+ static constexpr uint32_t VL_T = AscendC::GetVecLen() / sizeof(T);
140+ 
141+ AscendC::TPipe pipe_;
142+ const ThresholdGradV2DTilingData<RANK>* td_;
143+ AscendC::GlobalTensor<T> gmIn_[kMaxInputSlots];
144+ AscendC::GlobalTensor<T> gmOut_[kMaxOutputSlots];
145+ AscendC::TBuf<AscendC::TPosition::VECCALC> buf_[kPhysNodes];
146+ AscendC::MultiCopyParams<T, ND> nddmaParams_[kMaxInputSlots];
147+ int64_t nddmaOuterIters_[kMaxInputSlots];
148+ int64_t nddma_dims_;
149+ float threshold_;
150+ 
151+public:
152+ static constexpr int IN_GRAD = 0;
153+ static constexpr int IN_SELF = 1;
154+ static constexpr int OUT_Y = 0;
155+ 
156+ __aicore__ inline float TypeMin() {
157+ if constexpr (std::is_same_v<T, int8_t>) return -128.0f;
158+ if constexpr (std::is_same_v<T, uint8_t>) return 0.0f;
159+ if constexpr (std::is_same_v<T, int32_t>) return -2147483648.0f;
160+ return -3.4e38f;
161+ }
162+ __aicore__ inline float TypeMax() {
163+ if constexpr (std::is_same_v<T, int8_t>) return 127.0f;
164+ if constexpr (std::is_same_v<T, uint8_t>) return 255.0f;
165+ if constexpr (std::is_same_v<T, int32_t>) return 2147483647.0f;
166+ return 3.4e38f;
167+ }
168+ 
169+ __aicore__ inline void Init(GM_ADDR inputs[kMaxInputSlots], GM_ADDR outputs[kMaxOutputSlots],
170+ const ThresholdGradV2DTilingData<RANK>* td)
171+ {
172+ td_ = td;
173+ threshold_ = td_->threshold;
174+ for (int i = 0; i < kMaxInputSlots; i++) gmIn_[i].SetGlobalBuffer((__gm__ T*)inputs[i]);
175+ for (int i = 0; i < kMaxOutputSlots; i++) gmOut_[i].SetGlobalBuffer((__gm__ T*)outputs[i]);
176+ for (int i = 0; i < kPhysNodes; i++) pipe_.InitBuffer(buf_[i], td_->per_buf_bytes);
177+ 
178+ const int64_t* dstShape = td_->max_bro_shape;
179+ int64_t k = td_->split.axis;
180+ nddma_dims_ = (RANK - k <= ND) ? (RANK - k) : ND;
181+ for (int inp = 0; inp < kMaxInputSlots; inp++) {
182+ int64_t inner = 1;
183+ int64_t nd = 0;
184+ for (int64_t d = RANK - 1; d >= k && nd < ND; d--) {
185+ nddmaParams_[inp].loopInfo.loopSize[nd] = (d == k) ? 0 : dstShape[d];
186+ nddmaParams_[inp].loopInfo.loopSrcStride[nd] = td_->input_strides[inp][d];
187+ nddmaParams_[inp].loopInfo.loopDstStride[nd] = inner;
188+ nddmaParams_[inp].loopInfo.loopLpSize[nd] = 0;
189+ nddmaParams_[inp].loopInfo.loopRpSize[nd] = 0;
190+ inner *= (d == k) ? td_->split.a_i : dstShape[d];
191+ nd++;
192+ }
193+ for (; nd < ND; nd++) {
194+ nddmaParams_[inp].loopInfo.loopSize[nd] = 1;
195+ nddmaParams_[inp].loopInfo.loopSrcStride[nd] = 0;
196+ nddmaParams_[inp].loopInfo.loopDstStride[nd] = inner;
197+ nddmaParams_[inp].loopInfo.loopLpSize[nd] = 0;
198+ nddmaParams_[inp].loopInfo.loopRpSize[nd] = 0;
199+ }
200+ nddmaOuterIters_[inp] = 1;
201+ for (int64_t d = k; d < RANK - nddma_dims_; d++)
202+ nddmaOuterIters_[inp] *= (d == k) ? td_->split.a_i : dstShape[d];
203+ }
204+ }
205+ 
206+ __aicore__ inline void Process()
207+ {
208+ ProcessNative();
209+ }
210+ 
211+private:
212+ // ============================================================
213+ // 全 dtype 原生域直算, 零 cast (fp32/fp16/bf16/int32/int8/uint8 同路)
214+ // ============================================================
215+ __aicore__ inline void ProcessNative()
216+ {
217+ int32_t evMTE2toV = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_V));
218+ int32_t evVtoMTE3 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_MTE3));
219+ int32_t evMTE3toMTE2 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE3_MTE2));
220+ 
221+ int64_t start, end;
222+ GetCoreRange(AscendC::GetBlockIdx(), td_->multicore.tiles_main, td_->multicore.cores_tail, start, end);
223+ 
224+ constexpr int B0 = 0, B1 = 1, B2 = 2;
225+ int64_t inner_count = 1;
226+ for (int64_t d = td_->split.axis + 1; d < RANK; d++) inner_count *= td_->max_bro_shape[d];
227+ 
228+ int64_t coord[8] = {};
229+ for (int64_t flat = start; flat < end; flat++) {
230+ int64_t a_i_seg = GetUBSplitRange(flat % td_->split.a_o, td_->split.a_o,
231+ td_->split.a_i, td_->split.a_i_tail);
232+ int64_t count = a_i_seg * inner_count;
233+ FlatToEffectiveCoord(flat, td_->max_bro_shape, RANK,
234+ td_->split.axis, td_->split.a_i, td_->split.a_o, coord);
235+ 
236+ if (flat != start) AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
237+ 
238+ CopyInBrc(coord, IN_SELF, B0, a_i_seg); // self → B0
239+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
240+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
241+ CopyInBrc(coord, IN_GRAD, B1, a_i_seg); // grad → B1
242+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
243+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
244+ 
245+ int gateMode = 0; // 0=比较, 1=全透传, 2=全置0
246+ T thT = ResolveThreshold(gateMode);
247+ ComputeThresholdGrad(B0, B1, B2, thT, gateMode, count);
248+ 
249+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
250+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
251+ CopyOutOne(coord, OUT_Y, B2, a_i_seg);
252+ if (flat != end - 1) AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
253+ }
254+ }
255+ 
256+ // 阈值解析: 整型阈值(float)可能超出 T 范围, self>th 整数语义。
257+ // th<type_min -> 恒透传(gateMode=1); th>=type_max -> 恒置0(gateMode=2); 范围内取 floor(th) 整型比较。
258+ __aicore__ inline T ResolveThreshold(int& gateMode)
259+ {
260+ if constexpr (std::is_same_v<T, float> || std::is_same_v<T, half> ||
261+ std::is_same_v<T, bfloat16_t>) {
262+ // 浮点: threshold(float) 直接转 T 原生比较 (含 NaN 透传: self<=th 置0)
263+ return static_cast<T>(threshold_);
264+ } else {
265+ // 整型: float 阈值可能超出 T 范围; 范围内取 floor (self>floor(th) <=> self>th)
266+ if (threshold_ < TypeMin()) { gateMode = 1; return static_cast<T>(0); }
267+ if (threshold_ >= TypeMax()) { gateMode = 2; return static_cast<T>(0); }
268+ float f = threshold_;
269+ int64_t fl = (int64_t)f;
270+ if (f < (float)fl) fl -= 1; // floor
271+ return static_cast<T>(fl);
272+ }
273+ }
274+ 
275+ __aicore__ inline void ComputeThresholdGrad(int b0, int b1, int b2, T thT, int gateMode, int64_t count)
276+ {
277+ uint16_t rep = AscendC::CeilDivision(count, VL_T);
278+ if constexpr (std::is_same_v<T, half> || std::is_same_v<T, bfloat16_t>) {
279+ if (gateMode == 0) {
280+ // fp16/bf16 比较项升 fp32 与 golden 比较域对齐; grad 原生 Select。fp32 VL 步进。
281+ constexpr uint32_t VL_F = AscendC::GetVecLen() / sizeof(float);
282+ uint16_t repF = AscendC::CeilDivision(count, VL_F);
283+ asc_vf_call<ThresholdGradVFb16<T>>(
284+ (__ubuf__ T*)buf_[b2].Get<T>().GetPhyAddr(),
285+ (__ubuf__ T*)buf_[b0].Get<T>().GetPhyAddr(),
286+ (__ubuf__ T*)buf_[b1].Get<T>().GetPhyAddr(),
287+ threshold_, count, VL_F, repF);
288+ return;
289+ }
290+ }
291+ asc_vf_call<ThresholdGradVF<T>>(
292+ (__ubuf__ T*)buf_[b2].Get<T>().GetPhyAddr(),
293+ (__ubuf__ T*)buf_[b0].Get<T>().GetPhyAddr(),
294+ (__ubuf__ T*)buf_[b1].Get<T>().GetPhyAddr(),
295+ thT, gateMode, count, VL_T, rep);
296+ }
297+ 
298+ // ============================================================
299+ // CopyInBrc — NDDMA 随路 broadcast (broadcast 轴 src stride=0)
300+ // ============================================================
301+ __aicore__ inline void CopyInBrc(const int64_t* coord, int inputIdx, int slot, int64_t a_i_seg)
302+ {
303+ int64_t k = td_->split.axis;
304+ int64_t off = CalcOffset(coord, td_->input_strides[inputIdx], RANK);
305+ const int64_t* dstShape = td_->max_bro_shape;
306+ 
307+ auto params = nddmaParams_[inputIdx];
308+ int64_t k_nd = RANK - 1 - k;
309+ int64_t inner = 1;
310+ for (int64_t nd = 0; nd < ND; nd++) {
311+ if (nd == k_nd) params.loopInfo.loopSize[nd] = a_i_seg;
312+ params.loopInfo.loopDstStride[nd] = inner;
313+ inner *= params.loopInfo.loopSize[nd];
314+ }
315+ 
316+ static constexpr AscendC::NdDmaConfig cfg = { false, AscendC::NdDmaConfig::unsetPad,
317+ AscendC::NdDmaConfig::unsetPad, false };
318+ if constexpr (RANK <= 5) {
319+ AscendC::DataCopy<T, ND, cfg>(buf_[slot].Get<T>(), gmIn_[inputIdx][off], params);
320+ } else {
321+ AscendC::LocalTensor<T> buf = buf_[slot].Get<T>();
322+ int64_t elem_base = off;
323+ for (int64_t oi = 0; oi < nddmaOuterIters_[inputIdx]; oi++) {
324+ int64_t elem_adj = 0, tmp = oi;
325+ for (int64_t d = RANK - nddma_dims_ - 1; d >= k; d--) {
326+ int64_t sz = (d == k) ? a_i_seg : dstShape[d];
327+ elem_adj += (tmp % sz) * td_->input_strides[inputIdx][d];
328+ tmp /= sz;
329+ }
330+ AscendC::DataCopy<T, ND, cfg>(buf[oi * inner], gmIn_[inputIdx][elem_base + elem_adj], params);
331+ }
332+ }
333+ }
334+ 
335+ // ============================================================
336+ // CopyOutOne — DataCopyPad 写回 (out.shape = max_bro, 无 broadcast)
337+ // ============================================================
338+ __aicore__ inline void CopyOutOne(const int64_t* coord, int outputIdx, int slot, int64_t a_i_seg)
339+ {
340+ int64_t off = CalcOffset(coord, td_->output_strides[outputIdx], RANK);
341+ int64_t inner_elems = 1;
342+ for (int64_t d = td_->split.axis + 1; d < RANK; d++) inner_elems *= td_->output_shapes[outputIdx][d];
343+ int64_t cnt = a_i_seg * inner_elems;
344+ 
345+ AscendC::DataCopyExtParams extParams;
346+ extParams.blockCount = 1;
347+ extParams.blockLen = cnt * sizeof(T);
348+ extParams.srcStride = 0;
349+ extParams.dstStride = 0;
350+ AscendC::DataCopyPad(gmOut_[outputIdx][off], buf_[slot].Get<T>(), extParams);
351+ }
352+};
Aactivation/threshold_grad_v2_d/op_kernel/arch35/threshold_grad_v2_d_struct.h+35-0
@@ -0,0 +1,35 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file threshold_grad_v2_d_struct.h
15+ * \brief ThresholdGradV2D TilingKey 模板参数 — 仅按 RANK 分叉, dtype 由模板 T 处理 (DESIGN §3.1)
16+ */
17+#ifndef THRESHOLD_GRAD_V2_D_STRUCT_H_
18+#define THRESHOLD_GRAD_V2_D_STRUCT_H_
19+ 
20+#include "ascendc/host_api/tiling/template_argument.h"
21+ 
22+#define THRESHOLD_GRAD_V2_D_RANK_4 4
23+#define THRESHOLD_GRAD_V2_D_RANK_8 8
24+ 
25+ASCENDC_TPL_ARGS_DECL(ThresholdGradV2D,
26+ ASCENDC_TPL_UINT_DECL(RANK, 8, ASCENDC_TPL_UI_LIST,
27+ THRESHOLD_GRAD_V2_D_RANK_4, THRESHOLD_GRAD_V2_D_RANK_8)
28+);
29+ 
30+ASCENDC_TPL_SEL(
31+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, THRESHOLD_GRAD_V2_D_RANK_4)),
32+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, THRESHOLD_GRAD_V2_D_RANK_8))
33+);
34+ 
35+#endif // THRESHOLD_GRAD_V2_D_STRUCT_H_
Aactivation/threshold_grad_v2_d/op_kernel/arch35/threshold_grad_v2_d_tiling_struct.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+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file threshold_grad_v2_d_tiling_struct.h
15+ * \brief ThresholdGradV2D TilingData — 按 rank 模板化，体积分两档 (RANK<=4 / RANK<=8)
16+ * RegBase Broadcast 范式: out = (self > threshold) ? gradOutput : 0
17+ */
18+#pragma once
19+#include <cstdint>
20+ 
21+// === 算子特定常量 ===
22+constexpr int64_t kMaxInputSlots = 2; // gradOutput, self
23+constexpr int64_t kMaxOutputSlots = 1; // out
24+constexpr int64_t kPhysNodes = 3; // self/grad/out 三槽 (fp16/bf16 比较升 fp32 仅在寄存器内, 无额外 UB 中转)
25+ 
26+struct SplitResult {
27+ int64_t axis;
28+ int64_t a_i;
29+ int64_t a_o;
30+ int64_t a_i_tail;
31+};
32+ 
33+struct MultiCoreResult {
34+ int64_t num_cores;
35+ int64_t total_tiles;
36+ int64_t tiles_main;
37+ int64_t cores_tail;
38+};
39+ 
40+template<int64_t kRank>
41+struct ThresholdGradV2DTilingData {
42+ SplitResult split;
43+ MultiCoreResult multicore;
44+ int64_t rank; // 实际 rank (1~8)
45+ int64_t per_buf_bytes; // (UB/P)&~31; Kernel 用此初始化 TBuf
46+ int64_t per_buf_elems; // per_buf_bytes / 4（统一按 FP32 计算）
47+ int64_t max_bro_shape[kRank];
48+ int64_t num_inputs;
49+ int64_t num_outputs;
50+ int64_t input_shapes [kMaxInputSlots][kRank];
51+ int64_t input_strides[kMaxInputSlots][kRank];
52+ int64_t output_shapes[kMaxOutputSlots][kRank];
53+ int64_t output_strides[kMaxOutputSlots][kRank];
54+ float threshold; // 属性标量, 默认 1.0
55+};
Aactivation/threshold_grad_v2_d/op_kernel/threshold_grad_v2_d.cpp+49-0
@@ -0,0 +1,49 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file threshold_grad_v2_d_apt.cpp
15+ * \brief ThresholdGradV2D Kernel 入口 — apt 模式
16+ * RANK 来自 TilingKey, DTYPE 来自 def 注册的 gradOutput dtype
17+ * 目标架构: DAV_3510 (arch35 / ascend950)
18+ * 参数顺序固定: gradOutput, self, out, workspace, tiling
19+ */
20+#include "kernel_operator.h"
21+#include "arch35/threshold_grad_v2_d.h"
22+#include "arch35/threshold_grad_v2_d_tiling_struct.h"
23+ 
24+using TilingData4 = ThresholdGradV2DTilingData<4>; // RANK<=4: 数组维度 [4]
25+using TilingData8 = ThresholdGradV2DTilingData<8>; // RANK>4: 数组维度 [8]
26+ 
27+template<int RANK>
28+__global__ __aicore__ void threshold_grad_v2_d(
29+ GM_ADDR gradOutput, GM_ADDR self, GM_ADDR out,
30+ GM_ADDR workspace, GM_ADDR tiling)
31+{
32+ GM_ADDR ins[2] = {gradOutput, self};
33+ GM_ADDR outs[1] = {out};
34+ 
35+ REGISTER_NONE_TILING;
36+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
37+ 
38+ if constexpr (RANK == 4) {
39+ GET_TILING_DATA_WITH_STRUCT(TilingData4, td, tiling);
40+ ThresholdGradV2DKernel<DTYPE_GRADOUTPUT, 4> kernel;
41+ kernel.Init(ins, outs, &td);
42+ kernel.Process();
43+ } else {
44+ GET_TILING_DATA_WITH_STRUCT(TilingData8, td, tiling);
45+ ThresholdGradV2DKernel<DTYPE_GRADOUTPUT, 8> kernel;
46+ kernel.Init(ins, outs, &td);
47+ kernel.Process();
48+ }
49+}
Ractivation/threshold_grad_v2_d/op_host/CMakeLists.txt→activation/threshold_grad_v2_d/tests/ut/op_host/arch35/CMakeLists.txt+13-5
@@ -1,10 +1,18 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.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 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-#/
9 8 
10-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE threshold_grad_v2_d ACLNNTYPE aclnn DEPENDENCIES relu_grad)9+if(UT_TEST_ALL OR OP_HOST_UT)
10+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
11+ if(TARGET ${OP_TILING_MODULE_NAME}_cases_obj AND NOT TARGET ${OPHOST_NAME}_tiling_obj)
12+ target_sources(${OP_TILING_MODULE_NAME}_cases_obj PRIVATE
13+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../../op_host/arch35/threshold_grad_v2_d_tiling_arch35.cpp)
14+ target_include_directories(${OP_TILING_MODULE_NAME}_cases_obj PRIVATE
15+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../../
16+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../../op_host)
17+ endif()
18+endif()
Aactivation/threshold_grad_v2_d/tests/ut/op_host/arch35/test_threshold_grad_v2_d_tiling.cpp+218-0
@@ -0,0 +1,218 @@
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+#include <iostream>
11+#include <vector>
12+#include <gtest/gtest.h>
13+#include "log/log.h"
14+#include "kernel_run_context_facker.h"
15+#include "test_cube_util.h"
16+#include "exe_graph/runtime/storage_format.h"
17+#include "exe_graph/runtime/storage_shape.h"
18+#include "platform/platform_infos_def.h"
19+#include "ut_op_util.h"
20+ 
21+using namespace std;
22+using namespace ge;
23+using namespace ut_util;
24+ 
25+struct ThresholdGradV2DCompileInfo { uint64_t coreNum; uint64_t ubSize; };
26+ 
27+class ThresholdGradV2DTilingTest : public testing::Test {
28+protected:
29+ static void SetUpTestCase() {
30+ std::cout << "ThresholdGradV2DTilingTest SetUp" << std::endl;
31+ }
32+ static void TearDownTestCase() {
33+ std::cout << "ThresholdGradV2DTilingTest TearDown" << std::endl;
34+ }
35+};
36+ 
37+static void InitPlatform(fe::PlatFormInfos& platFormInfo, map<string, string>& socInfos,
38+ map<string, string>& aicoreSpec, map<string, string>& intrinsics, map<string, string>& socVersion)
39+{
40+ string hardwareInfo = R"({
41+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
42+ "Intrinsic_fix_pipe_l0c2out": false,
43+ "Intrinsic_data_move_l12ub": true,
44+ "Intrinsic_data_move_l0c2ub": true,
45+ "Intrinsic_data_move_out2l1_nd2nz": false,
46+ "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
47+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
48+ "CORE_NUM": 64, "socVersion": "Ascend950"}})";
49+ GetPlatFormInfos(hardwareInfo.c_str(), socInfos, aicoreSpec, intrinsics, socVersion);
50+ platFormInfo.Init();
51+}
52+ 
53+static void DoThresholdGradV2DTilingCase(std::initializer_list<int64_t>& inputShape1,
54+ std::initializer_list<int64_t>& inputShape2, std::initializer_list<int64_t>& outputShape,
55+ ge::DataType inputDtype)
56+{
57+ fe::PlatFormInfos platFormInfo;
58+ map<string, string> socInfos;
59+ map<string, string> aicoreSpec;
60+ map<string, string> intrinsics;
61+ map<string, string> socVersion;
62+ InitPlatform(platFormInfo, socInfos, aicoreSpec, intrinsics, socVersion);
63+ 
64+ ThresholdGradV2DCompileInfo compileInfo;
65+ std::string opType("ThresholdGradV2D");
66+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
67+ 
68+ string compileInfoStr = R"({})";
69+ auto kernelHolder = gert::KernelRunContextFaker()
70+ .KernelIONum(2, 1)
71+ .Inputs({const_cast<char*>(compileInfoStr.c_str()), reinterpret_cast<void*>(&platFormInfo)})
72+ .Outputs({&compileInfo})
73+ .Build();
74+ 
75+ ASSERT_TRUE(kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
76+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("version", socVersion);
77+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
78+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
79+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
80+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
81+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
82+ ASSERT_NE(tilingParseFunc, nullptr) << "tiling_parse not registered for " << opType;
83+ ASSERT_EQ(tilingParseFunc(kernelHolder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
84+ 
85+ gert::StorageShape x1Shape = {inputShape1, inputShape1};
86+ gert::StorageShape x2Shape = {inputShape2, inputShape2};
87+ gert::StorageShape oShape = {outputShape, outputShape};
88+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
89+ ASSERT_NE(tilingFunc, nullptr) << "tiling not registered for " << opType;
90+ 
91+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(16 * 4096);
92+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
93+ auto param = gert::TilingData::CreateCap(4096);
94+ ASSERT_NE(param, nullptr);
95+ 
96+ auto holder = gert::TilingContextFaker()
97+ .NodeIoNum(2, 1)
98+ .IrInstanceNum({1, 1})
99+ .InputShapes({&x1Shape, &x2Shape})
100+ .OutputShapes({&oShape})
101+ .CompileInfo(&compileInfo)
102+ .PlatformInfo(reinterpret_cast<char*>(&platFormInfo))
103+ .NodeInputTd(0, inputDtype, ge::FORMAT_ND, ge::FORMAT_ND)
104+ .NodeInputTd(1, inputDtype, ge::FORMAT_ND, ge::FORMAT_ND)
105+ .NodeOutputTd(0, inputDtype, ge::FORMAT_ND, ge::FORMAT_ND)
106+ .TilingData(param.get())
107+ .Workspace(wsSize)
108+ .Build();
109+ 
110+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
111+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
112+ tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
113+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
114+ tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
115+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
116+ 
117+ EXPECT_EQ(tilingFunc(tiling_context), ge::GRAPH_SUCCESS);
118+}
119+ 
120+TEST_F(ThresholdGradV2DTilingTest, tiling_fp32_2d) {
121+ std::initializer_list<int64_t> inputShape1 = {2, 128};
122+ std::initializer_list<int64_t> inputShape2 = {2, 128};
123+ std::initializer_list<int64_t> outputShape = {2, 128};
124+ DoThresholdGradV2DTilingCase(inputShape1, inputShape2, outputShape, ge::DT_FLOAT);
125+}
126+ 
127+TEST_F(ThresholdGradV2DTilingTest, tiling_fp16_2d) {
128+ std::initializer_list<int64_t> inputShape1 = {2, 128};
129+ std::initializer_list<int64_t> inputShape2 = {2, 128};
130+ std::initializer_list<int64_t> outputShape = {2, 128};
131+ DoThresholdGradV2DTilingCase(inputShape1, inputShape2, outputShape, ge::DT_FLOAT16);
132+}
133+ 
134+TEST_F(ThresholdGradV2DTilingTest, tiling_bf16_2d) {
135+ std::initializer_list<int64_t> inputShape1 = {2, 128};
136+ std::initializer_list<int64_t> inputShape2 = {2, 128};
137+ std::initializer_list<int64_t> outputShape = {2, 128};
138+ DoThresholdGradV2DTilingCase(inputShape1, inputShape2, outputShape, ge::DT_BF16);
139+}
140+ 
141+TEST_F(ThresholdGradV2DTilingTest, tiling_int32_2d) {
142+ std::initializer_list<int64_t> inputShape1 = {2, 128};
143+ std::initializer_list<int64_t> inputShape2 = {2, 128};
144+ std::initializer_list<int64_t> outputShape = {2, 128};
145+ DoThresholdGradV2DTilingCase(inputShape1, inputShape2, outputShape, ge::DT_INT32);
146+}
147+ 
148+TEST_F(ThresholdGradV2DTilingTest, tiling_int8_2d) {
149+ std::initializer_list<int64_t> inputShape1 = {2, 128};
150+ std::initializer_list<int64_t> inputShape2 = {2, 128};
151+ std::initializer_list<int64_t> outputShape = {2, 128};
152+ DoThresholdGradV2DTilingCase(inputShape1, inputShape2, outputShape, ge::DT_INT8);
153+}
154+ 
155+TEST_F(ThresholdGradV2DTilingTest, tiling_uint8_2d) {
156+ std::initializer_list<int64_t> inputShape1 = {2, 128};
157+ std::initializer_list<int64_t> inputShape2 = {2, 128};
158+ std::initializer_list<int64_t> outputShape = {2, 128};
159+ DoThresholdGradV2DTilingCase(inputShape1, inputShape2, outputShape, ge::DT_UINT8);
160+}
161+ 
162+TEST_F(ThresholdGradV2DTilingTest, tiling_fp32_4d) {
163+ std::initializer_list<int64_t> inputShape1 = {2, 4, 8, 16};
164+ std::initializer_list<int64_t> inputShape2 = {2, 4, 8, 16};
165+ std::initializer_list<int64_t> outputShape = {2, 4, 8, 16};
166+ DoThresholdGradV2DTilingCase(inputShape1, inputShape2, outputShape, ge::DT_FLOAT);
167+}
168+ 
169+TEST_F(ThresholdGradV2DTilingTest, tiling_fp16_1d) {
170+ std::initializer_list<int64_t> inputShape1 = {1024};
171+ std::initializer_list<int64_t> inputShape2 = {1024};
172+ std::initializer_list<int64_t> outputShape = {1024};
173+ DoThresholdGradV2DTilingCase(inputShape1, inputShape2, outputShape, ge::DT_FLOAT16);
174+}
175+ 
176+TEST_F(ThresholdGradV2DTilingTest, tiling_failed_unsupported_dtype) {
177+ fe::PlatFormInfos platFormInfo;
178+ map<string, string> socInfos;
179+ map<string, string> aicoreSpec;
180+ map<string, string> intrinsics;
181+ map<string, string> socVersion;
182+ InitPlatform(platFormInfo, socInfos, aicoreSpec, intrinsics, socVersion);
183+ 
184+ std::string op_type("ThresholdGradV2D");
185+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
186+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
187+ ASSERT_NE(tiling_func, nullptr) << "tiling not registered for " << op_type;
188+ 
189+ auto param = gert::TilingData::CreateCap(4096);
190+ ASSERT_NE(param, nullptr);
191+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
192+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
193+ gert::StorageShape Shape1 = {{2, 128}, {2, 128}};
194+ 
195+ ThresholdGradV2DCompileInfo compile_info;
196+ compile_info.coreNum = 64;
197+ compile_info.ubSize = 262144;
198+ 
199+ auto holder = gert::TilingContextFaker()
200+ .NodeIoNum(2, 1)
201+ .IrInstanceNum({1, 1})
202+ .InputShapes({&Shape1, &Shape1})
203+ .OutputShapes({&Shape1})
204+ .CompileInfo(&compile_info)
205+ .PlatformInfo(reinterpret_cast<char*>(&platFormInfo))
206+ .NodeInputTd(0, ge::DT_BOOL, ge::FORMAT_ND, ge::FORMAT_ND)
207+ .NodeInputTd(1, ge::DT_BOOL, ge::FORMAT_ND, ge::FORMAT_ND)
208+ .NodeOutputTd(0, ge::DT_BOOL, ge::FORMAT_ND, ge::FORMAT_ND)
209+ .TilingData(param.get())
210+ .Workspace(ws_size)
211+ .Build();
212+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
213+ tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
214+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
215+ tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
216+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
217+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_FAILED);
218+}
Mactivation/threshold_grad_v2_d/tests/ut/op_host/op_api/test_aclnn_threshold_backward.cpp+137-0
@@ -207,4 +207,141 @@ TEST_F(l2_threshold_backward_test, ascend910B2_l2_test_relu_grad_check_max_dim)
207 uint64_t workspaceSize = 0;207 uint64_t workspaceSize = 0;
208 aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);208 aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
209 EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);209 EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
210+}
211+ 
212+TEST_F(l2_threshold_backward_test, l2_test_relu_grad_uint8_success) {
213+ auto gradOutputDesc = TensorDesc({10,}, ACL_UINT8, ACL_FORMAT_ND);
214+ auto selfDesc = TensorDesc({10,}, ACL_UINT8, ACL_FORMAT_ND);
215+ auto scalarDesc = ScalarDesc(0.0f);
216+ auto outDesc = TensorDesc(selfDesc);
217+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
218+ uint64_t workspaceSize = 0;
219+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
220+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
221+ ut.TestPrecision();
222+}
223+ 
224+TEST_F(l2_threshold_backward_test, l2_test_relu_grad_int64_unsupported_on_stub) {
225+ auto gradOutputDesc = TensorDesc({10,}, ACL_INT64, ACL_FORMAT_ND);
226+ auto selfDesc = TensorDesc({10,}, ACL_INT64, ACL_FORMAT_ND);
227+ auto scalarDesc = ScalarDesc(0.0f);
228+ auto outDesc = TensorDesc(selfDesc);
229+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
230+ uint64_t workspaceSize = 0;
231+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
232+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
233+}
234+ 
235+TEST_F(l2_threshold_backward_test, l2_test_int64_nonzero_threshold_fail) {
236+ auto gradOutputDesc = TensorDesc({10,}, ACL_INT64, ACL_FORMAT_ND);
237+ auto selfDesc = TensorDesc({10,}, ACL_INT64, ACL_FORMAT_ND);
238+ auto scalarDesc = ScalarDesc(1.0f);
239+ auto outDesc = TensorDesc(selfDesc);
240+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
241+ uint64_t workspaceSize = 0;
242+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
243+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
244+}
245+ 
246+TEST_F(l2_threshold_backward_test, l2_test_2d_shape_success) {
247+ auto gradOutputDesc = TensorDesc({4, 8}, ACL_FLOAT, ACL_FORMAT_ND);
248+ auto selfDesc = TensorDesc({4, 8}, ACL_FLOAT, ACL_FORMAT_ND);
249+ auto scalarDesc = ScalarDesc(0.5f);
250+ auto outDesc = TensorDesc(selfDesc);
251+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
252+ uint64_t workspaceSize = 0;
253+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
254+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
255+ ut.TestPrecision();
256+}
257+ 
258+TEST_F(l2_threshold_backward_test, l2_test_4d_shape_success) {
259+ auto gradOutputDesc = TensorDesc({2, 3, 4, 5}, ACL_FLOAT16, ACL_FORMAT_ND);
260+ auto selfDesc = TensorDesc({2, 3, 4, 5}, ACL_FLOAT16, ACL_FORMAT_ND);
261+ auto scalarDesc = ScalarDesc(0.0f);
262+ auto outDesc = TensorDesc(selfDesc);
263+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
264+ uint64_t workspaceSize = 0;
265+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
266+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
267+ ut.TestPrecision();
268+}
269+ 
270+TEST_F(l2_threshold_backward_test, l2_test_8d_shape_success) {
271+ auto gradOutputDesc = TensorDesc({2, 2, 2, 2, 2, 2, 2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
272+ auto selfDesc = TensorDesc({2, 2, 2, 2, 2, 2, 2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
273+ auto scalarDesc = ScalarDesc(0.0f);
274+ auto outDesc = TensorDesc(selfDesc);
275+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
276+ uint64_t workspaceSize = 0;
277+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
278+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
279+}
280+ 
281+TEST_F(l2_threshold_backward_test, l2_test_large_shape_success) {
282+ auto gradOutputDesc = TensorDesc({1024, 1024}, ACL_FLOAT16, ACL_FORMAT_ND);
283+ auto selfDesc = TensorDesc({1024, 1024}, ACL_FLOAT16, ACL_FORMAT_ND);
284+ auto scalarDesc = ScalarDesc(0.0f);
285+ auto outDesc = TensorDesc(selfDesc);
286+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
287+ uint64_t workspaceSize = 0;
288+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
289+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
290+}
291+ 
292+TEST_F(l2_threshold_backward_test, l2_test_bf16_threshold_zero_precision) {
293+ auto gradOutputDesc = TensorDesc({10,}, ACL_BF16, ACL_FORMAT_ND);
294+ auto selfDesc = TensorDesc({10,}, ACL_BF16, ACL_FORMAT_ND);
295+ auto scalarDesc = ScalarDesc(0.0f);
296+ auto outDesc = TensorDesc(selfDesc);
297+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
298+ uint64_t workspaceSize = 0;
299+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
300+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
301+ ut.TestPrecision();
302+}
303+ 
304+TEST_F(l2_threshold_backward_test, l2_test_uint8_nonzero_threshold_success) {
305+ auto gradOutputDesc = TensorDesc({10,}, ACL_UINT8, ACL_FORMAT_ND);
306+ auto selfDesc = TensorDesc({10,}, ACL_UINT8, ACL_FORMAT_ND);
307+ auto scalarDesc = ScalarDesc(1.0f);
308+ auto outDesc = TensorDesc(selfDesc);
309+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
310+ uint64_t workspaceSize = 0;
311+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
312+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
313+ ut.TestPrecision();
314+}
315+ 
316+TEST_F(l2_threshold_backward_test, l2_test_fp32_nonzero_threshold_precision) {
317+ auto gradOutputDesc = TensorDesc({10,}, ACL_FLOAT, ACL_FORMAT_ND);
318+ auto selfDesc = TensorDesc({10,}, ACL_FLOAT, ACL_FORMAT_ND);
319+ auto scalarDesc = ScalarDesc(2.0f);
320+ auto outDesc = TensorDesc(selfDesc);
321+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
322+ uint64_t workspaceSize = 0;
323+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
324+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
325+ ut.TestPrecision();
326+}
327+ 
328+TEST_F(l2_threshold_backward_test, l2_test_scalar_tensor_success) {
329+ auto gradOutputDesc = TensorDesc({}, ACL_FLOAT, ACL_FORMAT_ND);
330+ auto selfDesc = TensorDesc({}, ACL_FLOAT, ACL_FORMAT_ND);
331+ auto scalarDesc = ScalarDesc(0.0f);
332+ auto outDesc = TensorDesc(selfDesc);
333+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarDesc), OUTPUT(outDesc));
334+ uint64_t workspaceSize = 0;
335+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
336+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
337+}
338+ 
339+TEST_F(l2_threshold_backward_test, l2_test_threshold_nullptr) {
340+ auto gradOutputDesc = TensorDesc({10,}, ACL_FLOAT, ACL_FORMAT_ND);
341+ auto selfDesc = TensorDesc({10,}, ACL_FLOAT, ACL_FORMAT_ND);
342+ auto outDesc = TensorDesc(selfDesc);
343+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, (const aclScalar*)nullptr), OUTPUT(outDesc));
344+ uint64_t workspaceSize = 0;
345+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
346+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_NULLPTR);
210}347}
Aactivation/threshold_grad_v2_d/tests/ut/op_host/test_threshold_grad_v2_d_infershape.cpp+242-0
@@ -0,0 +1,242 @@
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+#include <iostream>
11+#include <gtest/gtest.h>
12+#include <vector>
13+#include "array_ops.h"
14+#include "register/op_impl_registry.h"
15+#include "kernel_run_context_facker.h"
16+#include "exe_graph/runtime/storage_format.h"
17+#include "exe_graph/runtime/storage_shape.h"
18+#include "log/log.h"
19+using namespace ge;
20+ 
21+class ThresholdGradV2DInferShapeTest : public testing::Test {
22+protected:
23+ static void SetUpTestCase() {
24+ std::cout << "ThresholdGradV2DInferShapeTest SetUp" << std::endl;
25+ }
26+ static void TearDownTestCase() {
27+ std::cout << "ThresholdGradV2DInferShapeTest TearDown" << std::endl;
28+ }
29+};
30+ 
31+TEST_F(ThresholdGradV2DInferShapeTest, test_same_shape_fp32) {
32+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
33+ gert::StorageShape gradOutput_shape = {{16, 16}, {16, 16}};
34+ gert::StorageShape self_shape = {{16, 16}, {16, 16}};
35+ gert::StorageShape yShape = {{}, {}};
36+ auto holder = gert::InferShapeContextFaker()
37+ .NodeIoNum(2, 1)
38+ .IrInstanceNum({1, 1})
39+ .InputShapes({&gradOutput_shape, &self_shape})
40+ .OutputShapes({&yShape})
41+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
42+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
43+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
44+ .Build();
45+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
46+ auto output_desc = holder.GetContext<gert::InferShapeContext>()->GetOutputShape(0);
47+ gert::Shape expected_output_shape = {16, 16};
48+ ASSERT_EQ(Ops::Base::ToString(*output_desc), Ops::Base::ToString(expected_output_shape));
49+}
50+ 
51+TEST_F(ThresholdGradV2DInferShapeTest, test_broadcast_shape) {
52+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
53+ gert::StorageShape gradOutput_shape = {{4, 1}, {4, 1}};
54+ gert::StorageShape self_shape = {{1, 8}, {1, 8}};
55+ gert::StorageShape yShape = {{}, {}};
56+ auto holder = gert::InferShapeContextFaker()
57+ .NodeIoNum(2, 1)
58+ .IrInstanceNum({1, 1})
59+ .InputShapes({&gradOutput_shape, &self_shape})
60+ .OutputShapes({&yShape})
61+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
62+ .NodeInputTd(1, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
63+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
64+ .Build();
65+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
66+ auto output_desc = holder.GetContext<gert::InferShapeContext>()->GetOutputShape(0);
67+ gert::Shape expected_output_shape = {4, 8};
68+ ASSERT_EQ(Ops::Base::ToString(*output_desc), Ops::Base::ToString(expected_output_shape));
69+}
70+ 
71+TEST_F(ThresholdGradV2DInferShapeTest, test_fp16_same_shape) {
72+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
73+ gert::StorageShape gradOutput_shape = {{32, 64, 128}, {32, 64, 128}};
74+ gert::StorageShape self_shape = {{32, 64, 128}, {32, 64, 128}};
75+ gert::StorageShape yShape = {{}, {}};
76+ auto holder = gert::InferShapeContextFaker()
77+ .NodeIoNum(2, 1)
78+ .IrInstanceNum({1, 1})
79+ .InputShapes({&gradOutput_shape, &self_shape})
80+ .OutputShapes({&yShape})
81+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
82+ .NodeInputTd(1, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
83+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
84+ .Build();
85+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
86+}
87+ 
88+TEST_F(ThresholdGradV2DInferShapeTest, test_bf16_same_shape) {
89+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
90+ gert::StorageShape gradOutput_shape = {{2, 64}, {2, 64}};
91+ gert::StorageShape self_shape = {{2, 64}, {2, 64}};
92+ gert::StorageShape yShape = {{}, {}};
93+ auto holder = gert::InferShapeContextFaker()
94+ .NodeIoNum(2, 1)
95+ .IrInstanceNum({1, 1})
96+ .InputShapes({&gradOutput_shape, &self_shape})
97+ .OutputShapes({&yShape})
98+ .NodeInputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
99+ .NodeInputTd(1, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
100+ .NodeOutputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
101+ .Build();
102+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
103+}
104+ 
105+TEST_F(ThresholdGradV2DInferShapeTest, test_int32_same_shape) {
106+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
107+ gert::StorageShape gradOutput_shape = {{4, 4}, {4, 4}};
108+ gert::StorageShape self_shape = {{4, 4}, {4, 4}};
109+ gert::StorageShape yShape = {{}, {}};
110+ auto holder = gert::InferShapeContextFaker()
111+ .NodeIoNum(2, 1)
112+ .IrInstanceNum({1, 1})
113+ .InputShapes({&gradOutput_shape, &self_shape})
114+ .OutputShapes({&yShape})
115+ .NodeInputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
116+ .NodeInputTd(1, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
117+ .NodeOutputTd(0, ge::DT_INT32, ge::FORMAT_ND, ge::FORMAT_ND)
118+ .Build();
119+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
120+}
121+ 
122+TEST_F(ThresholdGradV2DInferShapeTest, test_int8_same_shape) {
123+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
124+ gert::StorageShape gradOutput_shape = {{8, 8}, {8, 8}};
125+ gert::StorageShape self_shape = {{8, 8}, {8, 8}};
126+ gert::StorageShape yShape = {{}, {}};
127+ auto holder = gert::InferShapeContextFaker()
128+ .NodeIoNum(2, 1)
129+ .IrInstanceNum({1, 1})
130+ .InputShapes({&gradOutput_shape, &self_shape})
131+ .OutputShapes({&yShape})
132+ .NodeInputTd(0, ge::DT_INT8, ge::FORMAT_ND, ge::FORMAT_ND)
133+ .NodeInputTd(1, ge::DT_INT8, ge::FORMAT_ND, ge::FORMAT_ND)
134+ .NodeOutputTd(0, ge::DT_INT8, ge::FORMAT_ND, ge::FORMAT_ND)
135+ .Build();
136+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
137+}
138+ 
139+TEST_F(ThresholdGradV2DInferShapeTest, test_uint8_same_shape) {
140+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
141+ gert::StorageShape gradOutput_shape = {{16}, {16}};
142+ gert::StorageShape self_shape = {{16}, {16}};
143+ gert::StorageShape yShape = {{}, {}};
144+ auto holder = gert::InferShapeContextFaker()
145+ .NodeIoNum(2, 1)
146+ .IrInstanceNum({1, 1})
147+ .InputShapes({&gradOutput_shape, &self_shape})
148+ .OutputShapes({&yShape})
149+ .NodeInputTd(0, ge::DT_UINT8, ge::FORMAT_ND, ge::FORMAT_ND)
150+ .NodeInputTd(1, ge::DT_UINT8, ge::FORMAT_ND, ge::FORMAT_ND)
151+ .NodeOutputTd(0, ge::DT_UINT8, ge::FORMAT_ND, ge::FORMAT_ND)
152+ .Build();
153+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
154+}
155+ 
156+TEST_F(ThresholdGradV2DInferShapeTest, test_8d_shape) {
157+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
158+ gert::StorageShape gradOutput_shape = {{1, 2, 3, 4, 5, 6, 7, 8}, {1, 2, 3, 4, 5, 6, 7, 8}};
159+ gert::StorageShape self_shape = {{1, 2, 3, 4, 5, 6, 7, 8}, {1, 2, 3, 4, 5, 6, 7, 8}};
160+ gert::StorageShape yShape = {{}, {}};
161+ auto holder = gert::InferShapeContextFaker()
162+ .NodeIoNum(2, 1)
163+ .IrInstanceNum({1, 1})
164+ .InputShapes({&gradOutput_shape, &self_shape})
165+ .OutputShapes({&yShape})
166+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
167+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
168+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
169+ .Build();
170+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
171+}
172+ 
173+TEST_F(ThresholdGradV2DInferShapeTest, test_empty_tensor) {
174+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
175+ gert::StorageShape gradOutput_shape = {{0, 0}, {0, 0}};
176+ gert::StorageShape self_shape = {{0, 0}, {0, 0}};
177+ gert::StorageShape yShape = {{}, {}};
178+ auto holder = gert::InferShapeContextFaker()
179+ .NodeIoNum(2, 1)
180+ .IrInstanceNum({1, 1})
181+ .InputShapes({&gradOutput_shape, &self_shape})
182+ .OutputShapes({&yShape})
183+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
184+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
185+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
186+ .Build();
187+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
188+}
189+ 
190+TEST_F(ThresholdGradV2DInferShapeTest, test_scalar_input) {
191+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
192+ gert::StorageShape gradOutput_shape = {{}, {}};
193+ gert::StorageShape self_shape = {{}, {}};
194+ gert::StorageShape yShape = {{}, {}};
195+ auto holder = gert::InferShapeContextFaker()
196+ .NodeIoNum(2, 1)
197+ .IrInstanceNum({1, 1})
198+ .InputShapes({&gradOutput_shape, &self_shape})
199+ .OutputShapes({&yShape})
200+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
201+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
202+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
203+ .Build();
204+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
205+}
206+ 
207+TEST_F(ThresholdGradV2DInferShapeTest, test_large_shape) {
208+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
209+ gert::StorageShape gradOutput_shape = {{1024, 1024}, {1024, 1024}};
210+ gert::StorageShape self_shape = {{1024, 1024}, {1024, 1024}};
211+ gert::StorageShape yShape = {{}, {}};
212+ auto holder = gert::InferShapeContextFaker()
213+ .NodeIoNum(2, 1)
214+ .IrInstanceNum({1, 1})
215+ .InputShapes({&gradOutput_shape, &self_shape})
216+ .OutputShapes({&yShape})
217+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
218+ .NodeInputTd(1, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
219+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
220+ .Build();
221+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
222+}
223+ 
224+TEST_F(ThresholdGradV2DInferShapeTest, test_broadcast_different_rank) {
225+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("ThresholdGradV2D")->infer_shape;
226+ gert::StorageShape gradOutput_shape = {{2, 3, 4}, {2, 3, 4}};
227+ gert::StorageShape self_shape = {{4}, {4}};
228+ gert::StorageShape yShape = {{}, {}};
229+ auto holder = gert::InferShapeContextFaker()
230+ .NodeIoNum(2, 1)
231+ .IrInstanceNum({1, 1})
232+ .InputShapes({&gradOutput_shape, &self_shape})
233+ .OutputShapes({&yShape})
234+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
235+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
236+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
237+ .Build();
238+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
239+ auto output_desc = holder.GetContext<gert::InferShapeContext>()->GetOutputShape(0);
240+ gert::Shape expected_output_shape = {2, 3, 4};
241+ ASSERT_EQ(Ops::Base::ToString(*output_desc), Ops::Base::ToString(expected_output_shape));
242+}
Mdocs/zh/op_list.md+3-3
@@ -740,12 +740,12 @@
740 <tr>740 <tr>
741 <td>activation</td>741 <td>activation</td>
742 <td><a href="../../activation/threshold_grad_v2_d/README.md">threshold_grad_v2_d</a></td>742 <td><a href="../../activation/threshold_grad_v2_d/README.md">threshold_grad_v2_d</a></td>
743- <td>✗</td>743+ <td>✓</td>
744- <td>✗</td>744+ <td>✓</td>
745 <td>✓</td>745 <td>✓</td>
746 <td>✗</td>746 <td>✗</td>
747 <td>AI Core</td>747 <td>AI Core</td>
748- <td>该算子暂无Ascend C代码实现，欢迎开发者补充贡献，贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>748+ <td>完成threshold正向的反向计算：out = (self > threshold) ? gradOutput : 0；threshold==0时等价ReluGrad。</td>
749 </tr>749 </tr>
750 <tr>750 <tr>
751 <td>control</td>751 <td>control</td>