Star849
1706
代码
代码
Issues147
Pull Requests280
Actions
项目讨论
Wiki
分析
Star849
1706
已合并
[CANNBot]950新增softplusv2grad ascendc实现 #7046
Hana77创建于 7月6日
[CANNBot]950新增softplusv2grad ascendc实现 #7046
已合并
Hana77创建于 7月6日
从Hana77/ops-nnnew-softplus合入到cann/ops-nnmaster
共 28 个文件变更+2757-30
Mactivation/softplus_v2_grad/CMakeLists.txt+17-14
@@ -1,18 +1,21 @@
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# ----------------------------------------------------------------------------
9message(STATUS "=== Debug: start ops.activation.softplus_v2_grad.CMakeLists.txt ")10# Generated By CANNBot
10file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11 
11if(NOT ENABLE_TEST AND NOT BENCHMARK)12set(SUPPORT_COMPUTE_UNIT "ascend950")
12 list(REMOVE_ITEM CURRENT_DIRS tests)13set(SUPPORT_TILING_DIR "arch35")
13endif()14 
14foreach(SUB_DIR ${CURRENT_DIRS})15add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE
15 if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")16 DIR ${CMAKE_CURRENT_SOURCE_DIR}
16 add_subdirectory(${SUB_DIR})17 OPTYPE softplus_v2_grad
17 endif()18 ACLNNTYPE aclnn_exclude
18endforeach()19 COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
20 TILING_DIR ${SUPPORT_TILING_DIR}
21)
Aactivation/softplus_v2_grad/examples/test_geir_softplus_v2_grad.cpp+277-0
@@ -0,0 +1,277 @@
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#include <iostream>
14#include <fstream>
15#include <string.h>
16#include <stdint.h>
17#include <vector>
18#include <string>
19#include <map>
20#include "assert.h"
21 
22#include "graph.h"
23#include "types.h"
24#include "tensor.h"
25#include "ge_error_codes.h"
26#include "ge_api_types.h"
27#include "ge_api.h"
28#include "array_ops.h"
29#include "ge_ir_build.h"
30 
31#include "../op_graph/softplus_v2_grad_proto.h"
32 
33#define FAILED -1
34#define SUCCESS 0
35 
36using namespace ge;
37using std::map;
38using std::string;
39using std::vector;
40#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
41 vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
42 auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
43 TensorDesc placeholder##intputIndex##_desc = TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, \
44 intputDtype); \
45 placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
46 placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
47 Tensor tensor_placeholder##intputIndex; \
48 ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
49 placeholder##intputIndex##_desc, 2); \
50 if (ret != SUCCESS) { \
51 printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
52 return FAILED; \
53 } \
54 placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
55 placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
56 input.push_back(tensor_placeholder##intputIndex); \
57 graph.AddOp(placeholder##intputIndex); \
58 softplusV2Grad1.set_input_##intputName(placeholder##intputIndex); \
59 inputs.push_back(placeholder##intputIndex);
60 
61#define LOG_PRINT(message, ...) \
62 do { \
63 printf(message, ##__VA_ARGS__); \
64 } while (0)
65 
66string GetTime()
67{
68 time_t timep;
69 time(&timep);
70 char tmp[64];
71 strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
72 return tmp;
73}
74 
75uint32_t GetDataTypeSize(DataType dt)
76{
77 uint32_t dilation = 1;
78 uint32_t oneByte = 1;
79 uint32_t twoByte = 2;
80 uint32_t fourByte = 4;
81 uint32_t eightByte = 8;
82 
83 if (dt == ge::DT_FLOAT) {
84 dilation = fourByte;
85 } else if (dt == ge::DT_FLOAT16) {
86 dilation = twoByte;
87 } else if (dt == ge::DT_BF16) {
88 dilation = twoByte;
89 } else if (dt == ge::DT_INT16) {
90 dilation = twoByte;
91 } else if (dt == ge::DT_UINT16) {
92 dilation = twoByte;
93 } else if (dt == ge::DT_INT32) {
94 dilation = fourByte;
95 } else if (dt == ge::DT_UINT32) {
96 dilation = fourByte;
97 } else if (dt == ge::DT_INT64) {
98 dilation = eightByte;
99 } else if (dt == ge::DT_UINT64) {
100 dilation = eightByte;
101 } else if (dt == ge::DT_INT8) {
102 dilation = oneByte;
103 }
104 return dilation;
105}
106 
107int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value)
108{
109 input_tensor_desc.SetRealDimCnt(shapes.size());
110 size_t size = 1;
111 for (uint32_t i = 0; i < shapes.size(); i++) {
112 size *= shapes[i];
113 }
114 uint32_t byteSizeFloat32 = 4;
115 uint32_t data_len = size * byteSizeFloat32;
116 float* pData = new (std::nothrow) float[size];
117 
118 for (size_t i = 0; i < size; ++i) {
119 *(pData + i) = value;
120 }
121 input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
122 return SUCCESS;
123}
124 
125int32_t GenOnesData(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type,
126 int value)
127{
128 input_tensor_desc.SetRealDimCnt(shapes.size());
129 size_t size = 1;
130 for (uint32_t i = 0; i < shapes.size(); i++) {
131 size *= shapes[i];
132 }
133 uint32_t data_len = size * GetDataTypeSize(data_type);
134 int32_t* pData = new (std::nothrow) int32_t[data_len];
135 for (uint32_t i = 0; i < size; ++i) {
136 *(pData + i) = value;
137 }
138 input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
139 return SUCCESS;
140}
141 
142int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
143{
144 FILE* fp;
145 fp = fopen(bin_file.c_str(), "w");
146 fwrite(inputData, sizeof(uint8_t), data_size, fp);
147 fclose(fp);
148 return SUCCESS;
149}
150 
151int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs,
152 std::vector<Operator>& outputs, Graph& graph)
153{
154 Status ret = SUCCESS;
155 // 自定义代码：添加单算子定义到图中
156 auto softplusV2Grad1 = op::SoftplusV2Grad("softplusV2Grad1");
157 softplusV2Grad1.set_attr_beta(1.0f);
158 softplusV2Grad1.set_attr_threshold(20.0f);
159 std::vector<int64_t> xShape = {4, 2};
160 ADD_INPUT(1, input_gradients, inDtype, xShape);
161 ADD_INPUT(2, input_features, inDtype, xShape);
162 
163 outputs.push_back(softplusV2Grad1);
164 // 添加完毕
165 return SUCCESS;
166}
167 
168int main(int argc, char* argv[])
169{
170 const char* graph_name = "tc_ge_irrun_test";
171 Graph graph(graph_name);
172 std::vector<ge::Tensor> input;
173 
174 printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
175 std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
176 Status ret = ge::GEInitialize(global_options);
177 if (ret != SUCCESS) {
178 printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
179 return FAILED;
180 }
181 printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
182 
183 std::vector<Operator> inputs{};
184 std::vector<Operator> outputs{};
185 
186 std::cout << argv[1] << std::endl;
187 char* endptr;
188 
189 DataType inDtype = DT_FLOAT;
190 
191 std::cout << inDtype << std::endl;
192 
193 ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
194 if (ret != SUCCESS) {
195 printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
196 return FAILED;
197 }
198 
199 if (!inputs.empty() && !outputs.empty()) {
200 graph.SetInputs(inputs).SetOutputs(outputs);
201 }
202 
203 std::map<AscendString, AscendString> build_options = {
204 
205 };
206 printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
207 ge::Session* session = new Session(build_options);
208 
209 if (session == nullptr) {
210 printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
211 return FAILED;
212 }
213 printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
214 printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
215 
216 std::map<AscendString, AscendString> graph_options = {
217 
218 };
219 uint32_t graph_id = 0;
220 ret = session->AddGraph(graph_id, graph, graph_options);
221 
222 printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
223 printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
224 std::string file_path = "./dump";
225 aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
226 printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
227 std::vector<ge::Tensor> output;
228 ret = session->RunGraph(graph_id, input, output);
229 if (ret != SUCCESS) {
230 printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
231 delete session;
232 GEFinalize();
233 return FAILED;
234 }
235 printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
236 
237 int input_num = input.size();
238 for (int i = 0; i < input_num; i++) {
239 std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
240 string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
241 uint8_t* input_data_i = input[i].GetData();
242 int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
243 std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
244 uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
245 WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
246 }
247 
248 int output_num = output.size();
249 for (int i = 0; i < output_num; i++) {
250 std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
251 string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
252 uint8_t* output_data_i = output[i].GetData();
253 int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
254 std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
255 uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
256 WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
257 float* resultData = (float*)output_data_i;
258 for (int64_t j = 0; j < output_shape; j++) {
259 LOG_PRINT("result[%ld] is: %f\n", j, resultData[j]);
260 }
261 }
262 
263 ge::AscendString error_msg = ge::GEGetErrorMsgV2();
264 std::string error_str(error_msg.GetString());
265 std::cout << "Error message: " << error_str << std::endl;
266 ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
267 std::string warning_str(warning_msg.GetString());
268 std::cout << "Warning message: " << warning_str << std::endl;
269 printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
270 ret = ge::GEFinalize();
271 if (ret != SUCCESS) {
272 printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
273 return FAILED;
274 }
275 printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
276 return SUCCESS;
277}
Ractivation/softplus_v2_grad/op_host/op_api/aclnn_softplus_backward.cpp→activation/softplus_v2_grad/op_api/aclnn_softplus_backward.cpp+6-0
@@ -48,6 +48,9 @@ static const std::initializer_list<op::DataType> KERNEL_SUPPORT_LIST = {
48 48 
49static const std::initializer_list<DataType>& GetDtypeSupportList()49static const std::initializer_list<DataType>& GetDtypeSupportList()
50{50{
51 if (Ops::NN::AclnnUtil::IsRegbase()) {
52 return ASCEND910B_DTYPE_SUPPORT_LIST;
53 }
51 if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910B ||54 if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910B ||
52 GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_93) {55 GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_93) {
53 return ASCEND910B_DTYPE_SUPPORT_LIST;56 return ASCEND910B_DTYPE_SUPPORT_LIST;
@@ -58,6 +61,9 @@ static const std::initializer_list<DataType>& GetDtypeSupportList()
58 61 
59static const std::initializer_list<DataType>& GetSelfDtypeSupportList()62static const std::initializer_list<DataType>& GetSelfDtypeSupportList()
60{63{
64 if (Ops::NN::AclnnUtil::IsRegbase()) {
65 return SELF_ASCEND910B_DTYPE_SUPPORT_LIST;
66 }
61 if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910B ||67 if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910B ||
62 GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_93) {68 GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND910_93) {
63 return SELF_ASCEND910B_DTYPE_SUPPORT_LIST;69 return SELF_ASCEND910B_DTYPE_SUPPORT_LIST;
Ractivation/softplus_v2_grad/op_host/op_api/aclnn_softplus_backward.h→activation/softplus_v2_grad/op_api/aclnn_softplus_backward.h+0-0
文件重命名但无更改。
Ractivation/softplus_v2_grad/op_host/op_api/softplus_v2_grad.cpp→activation/softplus_v2_grad/op_api/softplus_v2_grad.cpp+0-0
文件重命名但无更改。
Ractivation/softplus_v2_grad/op_host/op_api/softplus_v2_grad.h→activation/softplus_v2_grad/op_api/softplus_v2_grad.h+0-0
文件重命名但无更改。
Aactivation/softplus_v2_grad/op_graph/softplus_v2_grad_proto.h+57-0
@@ -0,0 +1,57 @@
1/**
2 * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.
9 *
10 * Generated By CANNBot
11 */
12 
13/*!
14 * \file softplus_v2_grad_proto.h
15 * \brief SoftplusV2Grad 算子图模式 proto 定义
16 */
17 
18#ifndef SOFTPLUS_V2_GRAD_PROTO_H_
19#define SOFTPLUS_V2_GRAD_PROTO_H_
20 
21#include "graph/operator_reg.h"
22 
23namespace ge {
24 
25/**
26* @brief Calculates the reversed outputs of the function "softplus_v2".
27 
28* @par Inputs:
29* Two inputs, including:
30* @li input_gradients: A mutable tensor, which supports 1D-8D defaultly. Format support ND.
31* Must be one of the following types: float16, float32, bfloat16.
32* @li input_features: A mutable tensor of the same type, shape and format as "input_gradients".
33 
34* @par Attributes:
35* @li beta: An optional float. Defaults to "1.0".
36* Control the steepness of the beta function.
37 
38* @li threshold: An optional float. Defaults to "20.0".
39* Define a function to switch the threshold from nonlinear to linear.
40 
41* @par Outputs:
42* output_backprops: A mutable tensor of the same type, shape and format as "input_gradients".
43 
44* @par Third-party framework compatibility
45* Compatible with the Pytorch operator SoftplusGrad.
46*/
47 
48REG_OP(SoftplusV2Grad)
49 .INPUT(input_gradients, TensorType({DT_FLOAT, DT_FLOAT16, DT_BF16}))
50 .INPUT(input_features, TensorType({DT_FLOAT, DT_FLOAT16, DT_BF16}))
51 .OUTPUT(output_backprops, TensorType({DT_FLOAT, DT_FLOAT16, DT_BF16}))
52 .ATTR(beta, Float, 1.0)
53 .ATTR(threshold, Float, 20.0)
54 .OP_END_FACTORY_REG(SoftplusV2Grad)
55 
56} // namespace ge
57#endif // SOFTPLUS_V2_GRAD_PROTO_H_
Aactivation/softplus_v2_grad/op_host/arch35/softplus_v2_grad_tiling_arch35.cpp+479-0
@@ -0,0 +1,479 @@
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// SoftplusV2Grad Tiling — arch35 实现
14#include "softplus_v2_grad_tiling_arch35.h"
15#include "../../op_kernel/arch35/softplus_v2_grad_struct.h"
16#include <algorithm>
17#include <sstream>
18#include <graph/utils/type_utils.h>
19#include "register/op_impl_registry.h"
20#include "op_common/log/log.h"
21#include "tiling/platform/platform_ascendc.h"
22#include "platform/soc_spec.h"
23 
24using namespace ge;
25 
26// ============================================================
27// Tiling 函数模块 — namespace softplus_v2_grad
28// ============================================================
29 
30namespace softplus_v2_grad {
31 
32// CheckBroadcastShape — 逐维验证 broadcast 兼容性
33bool CheckBroadcastShape(const std::vector<std::vector<int64_t>>& padded_in,
34 const std::vector<std::vector<int64_t>>& padded_out, int64_t max_rank)
35{
36 for (int64_t d = 0; d < max_rank; d++) {
37 int64_t ref = -1;
38 for (size_t i = 0; i < padded_in.size(); i++) {
39 if (padded_in[i][d] != 1) {
40 if (ref == -1)
41 ref = padded_in[i][d];
42 else if (padded_in[i][d] != ref) {
43 OP_LOGE("CheckBroadcastShape", "dim %d broadcast incompatible: input[%d] size %d != %d", (int)d,
44 (int)i, (int)padded_in[i][d], (int)ref);
45 return false;
46 }
47 }
48 }
49 for (size_t i = 0; i < padded_out.size(); i++) {
50 if (padded_out[i][d] != 1) {
51 if (ref == -1)
52 ref = padded_out[i][d];
53 else if (padded_out[i][d] != ref) {
54 OP_LOGE("CheckBroadcastShape", "dim %d broadcast incompatible: output[%d] size %d != %d", (int)d,
55 (int)i, (int)padded_out[i][d], (int)ref);
56 return false;
57 }
58 }
59 }
60 }
61 return true;
62}
63 
64// ComputeBroadcastDims — 从等 rank padded shapes 中筛选非全 1 维，构建广播坐标和归一化 shape
65static void ComputeBroadcastDims(const std::vector<std::vector<int64_t>>& padded_in,
66 const std::vector<std::vector<int64_t>>& padded_out, int64_t max_rank,
67 std::vector<int64_t>& bro_shape, std::vector<std::vector<int64_t>>& norm_in,
68 std::vector<std::vector<int64_t>>& norm_out)
69{
70 for (int64_t d = 0; d < max_rank; d++) {
71 bool all_one = true;
72 int64_t max_dim = 0;
73 for (size_t i = 0; i < padded_in.size(); i++) {
74 if (padded_in[i][d] != 1)
75 all_one = false;
76 max_dim = std::max(max_dim, padded_in[i][d]);
77 }
78 for (size_t i = 0; i < padded_out.size(); i++) {
79 if (padded_out[i][d] != 1)
80 all_one = false;
81 max_dim = std::max(max_dim, padded_out[i][d]);
82 }
83 if (!all_one) {
84 bro_shape.push_back(max_dim);
85 for (size_t i = 0; i < padded_in.size(); i++)
86 norm_in[i].push_back(padded_in[i][d]);
87 for (size_t i = 0; i < padded_out.size(); i++)
88 norm_out[i].push_back(padded_out[i][d]);
89 }
90 }
91 if (bro_shape.empty()) {
92 bro_shape.push_back(1);
93 for (auto& v : norm_in)
94 v.push_back(1);
95 for (auto& v : norm_out)
96 v.push_back(1);
97 }
98}
99 
100// PadAndSqueeze — 将 input/output 的不同 shape 归一化到最大广播坐标系
101bool PadAndSqueeze(const std::vector<std::vector<int64_t>>& input_shapes,
102 const std::vector<std::vector<int64_t>>& output_shapes, std::vector<int64_t>& maximum_bro_shape,
103 std::vector<std::vector<int64_t>>& normal_input_shapes,
104 std::vector<std::vector<int64_t>>& normal_output_shapes)
105{
106 int64_t num_inputs = (int64_t)input_shapes.size();
107 int64_t num_outputs = (int64_t)output_shapes.size();
108 int64_t max_rank = 0;
109 for (auto& s : input_shapes)
110 max_rank = std::max(max_rank, (int64_t)s.size());
111 for (auto& s : output_shapes)
112 max_rank = std::max(max_rank, (int64_t)s.size());
113 auto pad = [&](const std::vector<int64_t>& s) {
114 std::vector<int64_t> p;
115 p.assign(max_rank - (int64_t)s.size(), 1);
116 p.insert(p.end(), s.begin(), s.end());
117 return p;
118 };
119 std::vector<std::vector<int64_t>> padded_in(num_inputs), padded_out(num_outputs);
120 for (int64_t i = 0; i < num_inputs; i++)
121 padded_in[i] = pad(input_shapes[i]);
122 for (int64_t i = 0; i < num_outputs; i++)
123 padded_out[i] = pad(output_shapes[i]);
124 maximum_bro_shape.clear();
125 normal_input_shapes.assign(num_inputs, std::vector<int64_t>());
126 normal_output_shapes.assign(num_outputs, std::vector<int64_t>());
127 ComputeBroadcastDims(padded_in, padded_out, max_rank, maximum_bro_shape, normal_input_shapes, normal_output_shapes);
128 return true;
129}
130 
131// FindSplitAxis — 确定 UB 切分轴和 tile 大小
132bool FindSplitAxis(const std::vector<int64_t>& max_bro_shape, int64_t dtype_size, int64_t ub_per_core,
133 int64_t phys_nodes, SplitResult& out)
134{
135 // 除零保护: phys_nodes / dtype_size 为除数，必须 > 0
136 if (phys_nodes <= 0 || dtype_size <= 0) {
137 OP_LOGE("FindSplitAxis", "invalid divisor: phys_nodes=%d dtype_size=%d", (int)phys_nodes, (int)dtype_size);
138 return false;
139 }
140 int64_t per_buf_bytes = (ub_per_core / phys_nodes) & ~(kUbAlignBytes - 1); // UB 对齐, TBuf 硬件要求
141 int64_t per_buf_elems = per_buf_bytes / dtype_size;
142 // 单 buffer 至少能容纳 1 个元素，否则后续 per_buf_elems/inner、%a_i 会除零
143 if (per_buf_elems <= 0) {
144 OP_LOGE("FindSplitAxis", "per_buf_elems=%d invalid (ub_per_core=%d phys_nodes=%d dtype_size=%d)",
145 (int)per_buf_elems, (int)ub_per_core, (int)phys_nodes, (int)dtype_size);
146 return false;
147 }
148 int64_t rank = (int64_t)max_bro_shape.size();
149 int64_t inner = 1;
150 for (int64_t k = rank - 1; k >= 0; k--) {
151 if (max_bro_shape[k] * inner > per_buf_elems) {
152 out.a_i = per_buf_elems / inner;
153 // 除零保护: 单行内层维度乘积超过单 buffer 容量时 a_i 会被整除成 0，
154 // 导致下方 /a_i、%a_i 除零。此 shape 当前不支持，返回失败。
155 if (out.a_i <= 0) {
156 OP_LOGE("FindSplitAxis",
157 "inner dims product exceeds per-buffer capacity (per_buf_elems=%d, inner=%d) at axis %d; shape "
158 "unsupported",
159 (int)per_buf_elems, (int)inner, (int)k);
160 return false;
161 }
162 out.a_o = (max_bro_shape[k] + out.a_i - 1) / out.a_i;
163 int64_t rem = max_bro_shape[k] % out.a_i;
164 out.a_i_tail = (rem == 0) ? out.a_i : rem;
165 out.axis = k;
166 return true;
167 }
168 if (k == 0) {
169 out.axis = 0;
170 out.a_i = max_bro_shape[0];
171 out.a_o = 1;
172 out.a_i_tail = max_bro_shape[0];
173 return true;
174 }
175 inner *= max_bro_shape[k];
176 }
177 return true;
178}
179 
180// MultiCoreSplit — 多核任务划分
181bool MultiCoreSplit(const std::vector<int64_t>& max_bro_shape, const SplitResult& ub_split, int64_t max_cores,
182 MultiCoreResult& out)
183{
184 int64_t k = ub_split.axis, outer_prod = 1;
185 for (int64_t j = 0; j < k; j++)
186 outer_prod *= max_bro_shape[j];
187 out.total_tiles = outer_prod * ub_split.a_o;
188 out.num_cores = (out.total_tiles < max_cores) ? out.total_tiles : max_cores;
189 // 除零保护: num_cores 为除数。total_tiles==0（空 tensor / 0 维）或 max_cores<=0
190 // 都会使 num_cores<=0，导致下方 /num_cores、%num_cores 除零。
191 if (out.num_cores <= 0) {
192 OP_LOGE("MultiCoreSplit", "num_cores=%d invalid (total_tiles=%d max_cores=%d)", (int)out.num_cores,
193 (int)out.total_tiles, (int)max_cores);
194 return false;
195 }
196 out.tiles_main = out.total_tiles / out.num_cores;
197 out.cores_tail = out.total_tiles % out.num_cores;
198 return true;
199}
200 
201// 地址偏移计算 — input/output 的 stride 计算逻辑一致，合并为一个函数
202bool PrecomputeStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides)
203{
204 int64_t rank = (int64_t)s.size();
205 strides.assign(rank, 0);
206 for (int64_t d = rank - 1; d >= 0; d--) {
207 if (s[d] == 1) {
208 strides[d] = 0;
209 continue;
210 }
211 int64_t prod = 1;
212 for (int64_t j = d + 1; j < rank; j++)
213 prod *= s[j];
214 strides[d] = prod;
215 }
216 return true;
217}
218 
219} // namespace softplus_v2_grad
220 
221// ============================================================
222// SoftplusV2GradTiling — CANN 主线
223// ============================================================
224 
225namespace optiling {
226 
227using namespace softplus_v2_grad;
228 
229static std::string Arr2String(const int64_t* arr, int64_t n)
230{
231 std::ostringstream oss;
232 oss << "[";
233 if (n > 0) {
S
Ssu-yueming7月6日

[High] GetComputeNodeInfo() 返回值未判空即解引用

文件位置: activation/softplus_v2_grad/op_host/arch35/softplus_v2_grad_tiling_arch35.cpp:233

问题说明

ctx_->GetComputeNodeInfo() 返回值在用于 ->GetInputsNum()/->GetOutputsNum() 前未判空，若返回 nullptr 将触发空指针解引用

假设检验分析

H0: GetComputeNodeInfo() 始终返回非空 → H1: 可能返回 nullptr。证据：代码对 compileInfo(第232行)和 shape(第234行)都做了判空，但 GetComputeNodeInfo() 返回值在 ->GetInputsNum()/->GetOutputsNum() 调用前未判空。对照 review_checklist.md §10.2(context指针nullptr检查，84条评审)和 §10.33(GetInputDesc每次调用均需判空)，多级指针需逐步检查。校验：若框架异常初始化返回 nullptr，->GetInputsNum() 触发空指针解引用。决策：确认存在指针保护缺失，自信值 75%

代码片段

229: ge::graphStatus SoftplusV2GradTiling::GetShapeInfo()
230: {
231:     auto compileInfo = reinterpret_cast<const SoftplusV2GradCompileInfo*>(ctx_->GetCompileInfo());
232:     OP_CHECK_NULL_WITH_CONTEXT(ctx_, compileInfo);
233:     for (size_t i = 0; i < ctx_->GetComputeNodeInfo()->GetInputsNum(); ++i) {   // <-- 未判空
234:         auto shape = ctx_->GetInputShape(i); OP_CHECK_NULL_WITH_CONTEXT(ctx_, shape);
...
240:     for (size_t i = 0; i < ctx_->GetComputeNodeInfo()->GetOutputsNum(); ++i) {  // <-- 未判空

引用规则

cpp-secure.md §3.5 指针操作使用前必须要判空; review_checklist.md §10.2 context指针nullptr检查

建议修复

auto* computeNodeInfo = ctx_->GetComputeNodeInfo();
OP_CHECK_NULL_WITH_CONTEXT(ctx_, computeNodeInfo);
for (size_t i = 0; i < computeNodeInfo->GetInputsNum(); ++i) { ... }
for (size_t i = 0; i < computeNodeInfo->GetOutputsNum(); ++i) { ... }
likedislike
234 for (int64_t i = 0; i < n - 1; ++i) {
235 oss << arr[i] << ",";
236 }
237 oss << arr[n - 1];
238 }
239 oss << "]";
240 return oss.str();
241}
242 
243SoftplusV2GradTiling::SoftplusV2GradTiling(gert::TilingContext* ctx) : ctx_(ctx) {}
244 
245ge::graphStatus SoftplusV2GradTiling::GetShapeInfo()
246{
247 auto compileInfo = reinterpret_cast<const SoftplusV2GradCompileInfo*>(ctx_->GetCompileInfo());
248 OP_CHECK_NULL_WITH_CONTEXT(ctx_, compileInfo);
249 auto* nodeInfo = ctx_->GetComputeNodeInfo();
250 OP_CHECK_NULL_WITH_CONTEXT(ctx_, nodeInfo);
251 for (size_t i = 0; i < nodeInfo->GetInputsNum(); ++i) {
252 auto shape = ctx_->GetInputShape(i);
253 OP_CHECK_NULL_WITH_CONTEXT(ctx_, shape);
254 std::vector<int64_t> dims;
255 gert::Shape s = shape->GetStorageShape();
256 for (size_t d = 0; d < s.GetDimNum(); ++d)
257 dims.push_back(s.GetDim(d));
258 raw_input_shapes_.push_back(dims);
259 }
260 for (size_t i = 0; i < nodeInfo->GetOutputsNum(); ++i) {
261 auto shape = ctx_->GetOutputShape(i);
262 OP_CHECK_NULL_WITH_CONTEXT(ctx_, shape);
263 std::vector<int64_t> dims;
264 gert::Shape s = shape->GetStorageShape();
265 for (size_t d = 0; d < s.GetDimNum(); ++d)
266 dims.push_back(s.GetDim(d));
267 raw_output_shapes_.push_back(dims);
268 }
269 auto inputDesc = ctx_->GetInputDesc(0);
270 OP_CHECK_NULL_WITH_CONTEXT(ctx_, inputDesc);
271 ge::DataType dtype = inputDesc->GetDataType();
272 if (dtype != ge::DT_FLOAT16 && dtype != ge::DT_BF16 && dtype != ge::DT_FLOAT) {
273 OP_LOGE(ctx_->GetNodeName(), "Unsupported dtype");
274 return GRAPH_FAILED;
275 }
276 dtype_size_ = ge::GetSizeByDataType(dtype);
277 
278 // 标量属性：默认值由成员初始化器（kDefaultBeta/kDefaultThreshold）设定，此处仅按 GetAttrs() 覆盖
279 {
280 auto runtimeAttrs = ctx_->GetAttrs();
281 if (runtimeAttrs != nullptr) {
282 const float* betaPtr = runtimeAttrs->GetAttrPointer<float>(kAttrIdxBeta);
283 const float* threshPtr = runtimeAttrs->GetAttrPointer<float>(kAttrIdxThreshold);
284 if (betaPtr != nullptr)
285 beta_ = *betaPtr;
286 if (threshPtr != nullptr)
287 threshold_ = *threshPtr;
288 }
289 }
290 
291 PadAndSqueeze(raw_input_shapes_, raw_output_shapes_, max_bro_shape_, normal_input_shapes_, normal_output_shapes_);
292 rank_ = (int64_t)max_bro_shape_.size();
293 
294 // 维测: 输入预处理结果
295 OP_LOGI(ctx_->GetNodeName(), "GetShapeInfo done rank %lld dtype %lld ub %llu core %llu", rank_, dtype_size_,
296 compileInfo->ubSize, compileInfo->coreNum);
297 
298 OP_CHECK_IF(!CheckBroadcastShape(normal_input_shapes_, normal_output_shapes_, rank_),
299 OP_LOGE(ctx_->GetNodeName(), "check broadcast shape failed"), return ge::GRAPH_FAILED);
300 
301 return GRAPH_SUCCESS;
302}
303 
304// 单行前补对齐: dst[0..delta) = padValue, dst[delta..delta+n) = src[0..n)
305static void PadRow(int64_t* dst, const int64_t* src, int64_t n, int64_t delta, int64_t padValue)
306{
307 for (int64_t d = 0; d < delta; d++)
308 dst[d] = padValue;
309 for (int64_t d = 0; d < n; d++)
310 dst[d + delta] = src[d];
311}
312 
313// 整行填充: dst[0..r) = value
314static void FillRow(int64_t* dst, int64_t r, int64_t value)
315{
316 for (int64_t d = 0; d < r; d++)
317 dst[d] = value;
318}
319 
320// 将 num 条 (shape, stride) 前补对齐写入槽位数组，剩余未用槽位填默认 (shape=1, stride=0)
321template <int64_t R>
322static void PadSlots(int64_t (*shapes)[R], int64_t (*strides)[R], int64_t maxSlots, int64_t num,
323 const std::vector<std::vector<int64_t>>& normShapes,
324 const std::vector<std::vector<int64_t>>& strideVecs, int64_t rank, int64_t delta)
325{
326 for (int64_t i = 0; i < num; i++) {
327 PadRow(shapes[i], normShapes[i].data(), rank, delta, 1);
328 PadRow(strides[i], strideVecs[i].data(), rank, delta, 0);
329 }
330 for (int64_t i = num; i < maxSlots; i++) {
331 FillRow(shapes[i], R, 1);
332 FillRow(strides[i], R, 0);
333 }
334}
335 
336// 归一化 shape/stride 补齐并写入 TilingData 各槽位（含 rank/delta、split.axis 右移、槽位数量）
337template <int64_t R>
338static void FillShapesAndStrides(SoftplusV2GradTilingData<R>* tiling,
339 const std::vector<std::vector<int64_t>>& normal_input_shapes,
340 const std::vector<std::vector<int64_t>>& normal_output_shapes,
341 const std::vector<int64_t>& max_bro_shape, int64_t rank)
342{
343 int64_t num_in = (int64_t)normal_input_shapes.size();
344 int64_t num_out = (int64_t)normal_output_shapes.size();
345 std::vector<std::vector<int64_t>> in_strides(num_in), out_strides(num_out);
346 for (int64_t i = 0; i < num_in; i++)
347 PrecomputeStrides(normal_input_shapes[i], in_strides[i]);
348 for (int64_t i = 0; i < num_out; i++)
349 PrecomputeStrides(normal_output_shapes[i], out_strides[i]);
350 
351 tiling->rank = rank;
352 int64_t delta = R - rank; // 前补维数
353 
354 // max_bro_shape: 前补 1，实际值右移
355 PadRow(tiling->max_bro_shape, max_bro_shape.data(), rank, delta, 1);
356 
357 // split axis 右平移
358 tiling->split.axis += delta;
359 
360 tiling->num_inputs = num_in;
361 tiling->num_outputs = num_out;
362 
363 // input/output 各槽位: 前补 shape=1 stride=0 后右移，未用槽位填默认
364 PadSlots<R>(tiling->input_shapes, tiling->input_strides, kMaxInputSlots, num_in, normal_input_shapes, in_strides,
365 rank, delta);
366 PadSlots<R>(tiling->output_shapes, tiling->output_strides, kMaxOutputSlots, num_out, normal_output_shapes,
367 out_strides, rank, delta);
368}
369 
370// 维测: 输出 TilingData 全部字段（含每个 input/output 槽位的 shape/stride）
371template <int64_t R>
372static void LogTilingData(const char* nodeName, int64_t rank, const SoftplusV2GradTilingData<R>* tiling, int64_t num_in,
373 int64_t num_out)
374{
375 OP_LOGI(nodeName,
376 "TilingData: per_buf_bytes=%lld rank=%lld->R=%d "
377 "max_bro_shape=%s "
378 "split(axis=%lld a_i=%lld a_o=%lld a_i_tail=%lld) "
379 "multi(cores=%lld tiles=%lld main=%lld core_tail=%lld) num_in=%lld num_out=%lld "
380 "beta=%f threshold=%f",
S
Ssu-yueming7月6日

[Low] DoTilingAndSet() 中 compileInfo 未重新判空（防御性编程建议）

文件位置: activation/softplus_v2_grad/op_host/arch35/softplus_v2_grad_tiling_arch35.cpp:378

问题说明

DoTilingAndSet() 通过 reinterpret_cast 重新获取 compileInfo 后直接解引用，未重新判空。虽然 RunTiling() 保证 GetShapeInfo() 先执行且已判空，但独立方法应防御性检查

假设检验分析

H0: compileInfo 在 GetShapeInfo() 已判空，此处安全 → H1: 独立调用路径应重新判空。证据：GetShapeInfo() 第232行已 OP_CHECK_NULL_WITH_CONTEXT 判空，但 DoTilingAndSet() 是独立 private 模板方法，重新获取 compileInfo 后直接解引用。对照 review_checklist.md §10.39(不可假设'前面已经检查过')。校验：RunTiling() 调用顺序保证安全，但防御性编程建议重检。决策：防御性建议（非确认 bug），自信值 55%

代码片段

372: template<int64_t R>
373: ge::graphStatus SoftplusV2GradTiling::DoTilingAndSet()
374: {
375:     auto* tiling = ctx_->GetTilingData<SoftplusV2GradTilingData<R>>();
376:     OP_CHECK_NULL_WITH_CONTEXT(ctx_, tiling);
377:
378:     auto* compileInfo = reinterpret_cast<const SoftplusV2GradCompileInfo*>(
379:         ctx_->GetCompileInfo());
380:     int64_t ub_per_core = (int64_t)compileInfo->ubSize;   // <-- compileInfo 未判空

引用规则

review_checklist.md §10.39 GetInputDesc每次调用均需判空（类比）; cpp-secure.md §3.5 指针操作使用前必须要判空

建议修复

auto* compileInfo = reinterpret_cast<const SoftplusV2GradCompileInfo*>(
    ctx_->GetCompileInfo());
OP_CHECK_NULL_WITH_CONTEXT(ctx_, compileInfo);
int64_t ub_per_core = (int64_t)compileInfo->ubSize;
likedislike
381 tiling->per_buf_bytes, rank, (int)R, Arr2String(tiling->max_bro_shape, R).c_str(), tiling->split.axis,
382 tiling->split.a_i, tiling->split.a_o, tiling->split.a_i_tail, tiling->multicore.num_cores,
383 tiling->multicore.total_tiles, tiling->multicore.tiles_main, tiling->multicore.cores_tail, num_in, num_out,
384 (double)tiling->beta, (double)tiling->threshold);
385 
386 for (int64_t i = 0; i < num_in; i++)
387 OP_LOGI(nodeName, "TilingData input[%lld]: shape=%s stride=%s", i,
388 Arr2String(tiling->input_shapes[i], R).c_str(), Arr2String(tiling->input_strides[i], R).c_str());
389 for (int64_t i = 0; i < num_out; i++)
390 OP_LOGI(nodeName, "TilingData output[%lld]: shape=%s stride=%s", i,
391 Arr2String(tiling->output_shapes[i], R).c_str(), Arr2String(tiling->output_strides[i], R).c_str());
392}
393 
394template <int64_t R>
395ge::graphStatus SoftplusV2GradTiling::DoTilingAndSet()
396{
397 auto* tiling = ctx_->GetTilingData<SoftplusV2GradTilingData<R>>();
398 OP_CHECK_NULL_WITH_CONTEXT(ctx_, tiling);
399 
400 auto* compileInfo = reinterpret_cast<const SoftplusV2GradCompileInfo*>(ctx_->GetCompileInfo());
401 OP_CHECK_NULL_WITH_CONTEXT(ctx_, compileInfo);
402 int64_t ub_per_core = (int64_t)compileInfo->ubSize;
403 int64_t per_buf_bytes = (ub_per_core / kPhysNodes) & ~(kUbAlignBytes - 1);
404 
405 // dtype_size 始终为 4 (sizeof(float))，per_buf_elems 以 fp32 元素数为准
406 if (!FindSplitAxis(max_bro_shape_, sizeof(float), ub_per_core, kPhysNodes, tiling->split)) {
407 OP_LOGE(ctx_->GetNodeName(), "FindSplitAxis failed");
408 return ge::GRAPH_FAILED;
409 }
410 if (!MultiCoreSplit(max_bro_shape_, tiling->split, (int64_t)compileInfo->coreNum, tiling->multicore)) {
411 OP_LOGE(ctx_->GetNodeName(), "MultiCoreSplit failed");
412 return ge::GRAPH_FAILED;
413 }
414 tiling->per_buf_bytes = per_buf_bytes;
415 tiling->beta = beta_;
416 tiling->threshold = threshold_;
417 
418 // 归一化 shape/stride 补齐并写入各槽位（含 rank/delta、split.axis 右移、槽位数量）
419 FillShapesAndStrides<R>(tiling, normal_input_shapes_, normal_output_shapes_, max_bro_shape_, rank_);
420 
421 ctx_->SetBlockDim(tiling->multicore.num_cores);
422 
423 // 维测: TilingData 全部字段
424 LogTilingData<R>(ctx_->GetNodeName(), rank_, tiling, tiling->num_inputs, tiling->num_outputs);
425 
426 return GRAPH_SUCCESS;
427}
428 
429ge::graphStatus SoftplusV2GradTiling::RunTiling()
430{
431 ge::graphStatus ret = GetShapeInfo();
432 if (ret != GRAPH_SUCCESS)
433 return ret;
434 
435 int64_t mapped = (rank_ <= SOFTPLUS_V2_GRAD_RANK_4) ? SOFTPLUS_V2_GRAD_RANK_4 : SOFTPLUS_V2_GRAD_RANK_8;
436 // TilingKey 仅在 DoTilingAndSet 成功后设置，避免失败时把有效 key 与未填充的 TilingData 绑定
437 if (mapped == SOFTPLUS_V2_GRAD_RANK_4) {
438 ret = DoTilingAndSet<SOFTPLUS_V2_GRAD_RANK_4>();
439 if (ret != GRAPH_SUCCESS)
440 return ret;
441 ctx_->SetTilingKey(GET_TPL_TILING_KEY(SOFTPLUS_V2_GRAD_RANK_4));
442 } else {
443 ret = DoTilingAndSet<SOFTPLUS_V2_GRAD_RANK_8>();
444 if (ret != GRAPH_SUCCESS)
445 return ret;
446 ctx_->SetTilingKey(GET_TPL_TILING_KEY(SOFTPLUS_V2_GRAD_RANK_8));
447 }
448 return ret;
449}
450 
451static ge::graphStatus TilingFuncSoftplusV2Grad(gert::TilingContext* context)
452{
453 SoftplusV2GradTiling softplusV2GradTiling(context);
454 auto ret = softplusV2GradTiling.RunTiling();
455 if (ret != GRAPH_SUCCESS)
456 return ret;
457 // Workspace: 当前不需要额外 workspace
458 size_t* workspaces = context->GetWorkspaceSizes(1);
459 workspaces[0] = 0;
460 return GRAPH_SUCCESS;
461}
462 
463ge::graphStatus TilingPrepareForSoftplusV2Grad(gert::TilingParseContext* context)
464{
465 fe::PlatFormInfos* platformInfo = context->GetPlatformInfo();
466 auto compileInfo = context->GetCompiledInfo<SoftplusV2GradCompileInfo>();
467 OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
468 OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
469 auto ap = platform_ascendc::PlatformAscendC(platformInfo);
470 compileInfo->coreNum = ap.GetCoreNumAiv();
471 ap.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfo->ubSize);
472 return GRAPH_SUCCESS;
473}
474 
475IMPL_OP_OPTILING(SoftplusV2Grad)
476 .Tiling(TilingFuncSoftplusV2Grad)
477 .TilingParse<SoftplusV2GradCompileInfo>(TilingPrepareForSoftplusV2Grad);
478 
479} // namespace optiling
Aactivation/softplus_v2_grad/op_host/arch35/softplus_v2_grad_tiling_arch35.h+77-0
@@ -0,0 +1,77 @@
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// SoftplusV2Grad Tiling — arch35 头文件
14#pragma once
15#include <cstdint>
16#include <vector>
17#include <exe_graph/runtime/tiling_context.h>
18#include "../../op_kernel/arch35/softplus_v2_grad_tiling_struct.h"
19 
20// ============================================================
21// Tiling 函数模块（无状态，与 CANN 框架零耦合）
22// 作用域：namespace softplus_v2_grad
23// ============================================================
24 
25namespace softplus_v2_grad {
26 
27bool CheckBroadcastShape(const std::vector<std::vector<int64_t>>& padded_in,
28 const std::vector<std::vector<int64_t>>& padded_out, int64_t max_rank);
29 
30bool PadAndSqueeze(const std::vector<std::vector<int64_t>>& input_shapes,
31 const std::vector<std::vector<int64_t>>& output_shapes, std::vector<int64_t>& maximum_bro_shape,
32 std::vector<std::vector<int64_t>>& normal_input_shapes,
33 std::vector<std::vector<int64_t>>& normal_output_shapes);
34 
35bool FindSplitAxis(const std::vector<int64_t>& max_bro_shape, int64_t dtype_size, int64_t ub_per_core,
36 int64_t phys_nodes, SplitResult& out);
37 
38bool MultiCoreSplit(const std::vector<int64_t>& max_bro_shape, const SplitResult& ub_split, int64_t max_cores,
39 MultiCoreResult& out);
40 
41bool PrecomputeStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides);
42 
43} // namespace softplus_v2_grad
44 
45// ============================================================
46// SoftplusV2GradTiling — CANN 主线：单次归一化 → 映射 rank → 模板填充
47// ============================================================
48namespace optiling {
49 
50struct SoftplusV2GradCompileInfo {
51 uint64_t coreNum;
52 uint64_t ubSize;
53};
54 
55class SoftplusV2GradTiling {
56public:
57 explicit SoftplusV2GradTiling(gert::TilingContext* ctx);
58 ge::graphStatus RunTiling();
59 
60private:
61 ge::graphStatus GetShapeInfo();
62 template <int64_t R>
63 ge::graphStatus DoTilingAndSet();
64 
65 gert::TilingContext* ctx_;
66 std::vector<std::vector<int64_t>> raw_input_shapes_;
67 std::vector<std::vector<int64_t>> raw_output_shapes_;
68 std::vector<int64_t> max_bro_shape_;
69 std::vector<std::vector<int64_t>> normal_input_shapes_;
70 std::vector<std::vector<int64_t>> normal_output_shapes_;
71 int64_t dtype_size_ = 0;
72 int64_t rank_ = 0;
73 float beta_ = kDefaultBeta;
74 float threshold_ = kDefaultThreshold;
75};
76 
77} // namespace optiling
Aactivation/softplus_v2_grad/op_host/softplus_v2_grad_def.cpp+58-0
@@ -0,0 +1,58 @@
1/**
2 * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.
9 *
10 * Generated By CANNBot
11 */
12 
13/**
14 * SoftplusV2Grad OpDef — 算子注册 (2入1出 + 2属性)
15 */
16#include "register/op_def_registry.h"
17 
18namespace ops {
19class SoftplusV2Grad : public OpDef {
20public:
21 explicit SoftplusV2Grad(const char* name) : OpDef(name)
22 {
23 this->Input("input_gradients")
24 .ParamType(REQUIRED)
25 .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
26 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
27 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
28 this->Input("input_features")
29 .ParamType(REQUIRED)
30 .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
31 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
32 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
33 this->Output("output_backprops")
34 .ParamType(REQUIRED)
35 .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
36 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
37 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
38 
39 // 标量属性（均为 OPTIONAL，有默认值）
40 // 注意：默认值须与 op_kernel/arch35/softplus_v2_grad_tiling_struct.h 中
41 // kDefaultBeta / kDefaultThreshold 保持一致（tiling 侧的运行期兜底默认）。
42 this->Attr("beta").AttrType(OPTIONAL).Float(1.0f);
43 this->Attr("threshold").AttrType(OPTIONAL).Float(20.0f);
44 
45 OpAICoreConfig aicoreConfig;
46 aicoreConfig.DynamicCompileStaticFlag(true)
47 .DynamicFormatFlag(false)
48 .DynamicRankSupportFlag(true)
49 .DynamicShapeSupportFlag(true)
50 .NeedCheckSupportFlag(false)
51 .PrecisionReduceFlag(true)
52 .ExtendCfgInfo("opFile.value", "softplus_v2_grad");
53 this->AICore().AddConfig("ascend950", aicoreConfig);
54 }
55};
56 
57OP_ADD(SoftplusV2Grad);
58} // namespace ops
Aactivation/softplus_v2_grad/op_host/softplus_v2_grad_infershape.cpp+28-0
@@ -0,0 +1,28 @@
1/**
2 * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.
9 *
10 * Generated By CANNBot
11 */
12 
13#include "register/op_impl_registry.h"
14#include "op_common/op_host/infershape_broadcast_util.h"
15 
16using namespace Ops::Base;
17namespace ops {
18 
19static ge::graphStatus InferShape4SoftplusV2Grad(gert::InferShapeContext* context)
20{
21 // Broadcast: input_gradients + input_features → output output_backprops
22 constexpr size_t kBroadcastInputNum = 2; // input_gradients, input_features
23 return InferShape4Broadcast(context, kBroadcastInputNum);
24}
25 
26IMPL_OP_INFERSHAPE(SoftplusV2Grad).InferShape(InferShape4SoftplusV2Grad);
27 
28} // namespace ops
Aactivation/softplus_v2_grad/op_kernel/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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 
13add_kernel_sources(
14 KERNEL_SRC arch35/softplus_v2_grad.cpp
15 COMPUTE_UNITS ascend950
16 AUTO_SYNC false
17)
Aactivation/softplus_v2_grad/op_kernel/arch35/softplus_v2_grad.cpp+44-0
@@ -0,0 +1,44 @@
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 * SoftplusV2Grad Kernel 入口
14 * RANK 来自 TilingKey, DTYPE 来自 CANN 框架 (def 注册的 Input("input_gradients") 类型)
15 */
16#include "kernel_operator.h"
17#include "softplus_v2_grad_kernel.h"
18#include "softplus_v2_grad_tiling_struct.h"
19 
20using TilingData4 = SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>;
21using TilingData8 = SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_8>;
22 
23template <int RANK>
24__global__ __aicore__ void softplus_v2_grad(GM_ADDR input_gradients, GM_ADDR input_features, GM_ADDR output_backprops,
25 GM_ADDR workspace, GM_ADDR tiling)
26{
27 GM_ADDR ins[kMaxInputSlots] = {input_gradients, input_features};
28 GM_ADDR outs[kMaxOutputSlots] = {output_backprops};
29 
30 REGISTER_NONE_TILING;
31 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
32 
33 if constexpr (RANK == SOFTPLUS_V2_GRAD_RANK_4) {
34 GET_TILING_DATA_WITH_STRUCT(TilingData4, td, tiling);
35 SoftplusV2GradKernel<DTYPE_INPUT_GRADIENTS, SOFTPLUS_V2_GRAD_RANK_4> kernel;
36 kernel.Init(ins, outs, &td);
37 kernel.Process();
38 } else {
39 GET_TILING_DATA_WITH_STRUCT(TilingData8, td, tiling);
40 SoftplusV2GradKernel<DTYPE_INPUT_GRADIENTS, SOFTPLUS_V2_GRAD_RANK_8> kernel;
41 kernel.Init(ins, outs, &td);
42 kernel.Process();
43 }
44}
Aactivation/softplus_v2_grad/op_kernel/arch35/softplus_v2_grad_kernel.h+383-0
@@ -0,0 +1,383 @@
1/**
2 * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.
9 *
10 * Generated By CANNBot
11 */
12 
13/**
14 * SoftplusV2Grad Kernel — SoftplusV2GradKernel<T, RANK>
15 */
16#pragma once
17#include "kernel_operator.h"
18#include "softplus_v2_grad_tiling_struct.h"
19#include "softplus_v2_grad_struct.h"
20 
21__aicore__ inline void GetCoreRange(int64_t core_id, int64_t tiles_main, int64_t cores_tail, int64_t& start,
22 int64_t& end)
23{
24 if (core_id < cores_tail) {
25 start = core_id * (tiles_main + 1);
26 end = start + tiles_main + 1;
27 } else {
28 start = cores_tail * (tiles_main + 1) + (core_id - cores_tail) * tiles_main;
29 end = start + tiles_main;
30 }
31}
32 
33__aicore__ inline int64_t GetUBSplitRange(int64_t a_o_off, int64_t a_o, int64_t a_i, int64_t a_i_tail)
34{
35 return (a_o_off == a_o - 1) ? a_i_tail : a_i;
36}
37 
38__aicore__ inline bool FlatToEffectiveCoord(int64_t flat, const int64_t* max_bro_shape, int64_t rank,
39 int64_t split_axis, int64_t a_i, int64_t a_o, int64_t* eff_coord)
40{
41 for (int64_t d = 0; d < rank; d++)
42 eff_coord[d] = 0;
43 a_o = (a_o == 0) ? 1 : a_o; // 除零兜底：a_o 由 host tiling 保证 >=1，此处防御非法 TilingData
44 int64_t a_o_off = flat % a_o;
45 int64_t outer = flat / a_o;
46 for (int64_t d = split_axis - 1; d >= 0; d--) {
47 eff_coord[d] = outer % max_bro_shape[d];
48 outer /= max_bro_shape[d];
49 }
50 eff_coord[split_axis] = a_o_off * a_i;
51 return true;
52}
53 
54// input/output 偏移计算逻辑一致（Σ eff_coord[d]*strides[d]）
55// 返回元素个数，用作 gmIn_[]/gmOut_[] 的 index
56__aicore__ inline int64_t CalcOffset(const int64_t* eff_coord, const int64_t* strides, int64_t rank)
57{
58 int64_t offset = 0;
59 for (int64_t d = 0; d < rank; d++)
60 offset += eff_coord[d] * strides[d];
61 return offset;
62}
63 
64// input/output 搬运元素数计算逻辑一致
65__aicore__ inline int64_t CalcTransferCount(const int64_t* normal_shape, int64_t rank, int64_t split_axis,
66 int64_t a_i_seg)
67{
68 int64_t split_elems = (normal_shape[split_axis] == 1) ? 1 : a_i_seg;
69 int64_t inner_elems = 1;
70 for (int64_t d = split_axis + 1; d < rank; d++)
71 inner_elems *= normal_shape[d];
72 return split_elems * inner_elems;
73}
74 
75template <typename T>
76__simd_vf__ inline void SoftplusV2GradVF(__ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint32_t count,
77 uint32_t VL, uint16_t rep, float beta, float threshold)
78{
79 // 寄存器声明
80 AscendC::Reg::RegTensor<T> src0Reg, src1Reg;
81 AscendC::Reg::RegTensor<T> regT1, regT2, regT3, regT4, regT5, dstReg;
82 AscendC::Reg::RegTensor<T> regThresh;
83 AscendC::Reg::MaskReg padMask;
84 AscendC::Reg::MaskReg cmpMask;
85 AscendC::Reg::AddrReg aReg;
86 
87 for (uint16_t i = 0; i < rep; i++) {
88 aReg = AscendC::Reg::CreateAddrReg<T>(i, VL);
89 padMask = AscendC::Reg::UpdateMask<T>(count);
90 
91 // --- Load: UB → Reg ---
92 AscendC::Reg::LoadAlign(src0Reg, src0, aReg); // 读 input_features
93 AscendC::Reg::LoadAlign(src1Reg, src1, aReg); // 读 input_gradients
94 
95 // S1: Mul(beta, input_features) → T1
96 // T1 = beta * input_features
97 AscendC::Reg::Muls<T>(regT1, src0Reg, beta, padMask);
98 
99 // S2: Neg(T1) → T2
100 AscendC::Reg::Neg<T>(regT2, regT1, padMask);
101 
102 // S3: Exp(T2) → T3
103 AscendC::Reg::Exp<T>(regT3, regT2, padMask);
104 
105 // S4: Adds(T3, 1.0f) → T4
106 AscendC::Reg::Adds<T>(regT4, regT3, static_cast<T>(1.0f), padMask);
107 
108 // S5: Div(input_gradients, T4) → T5
109 AscendC::Reg::Div<T>(regT5, src1Reg, regT4, padMask);
110 
111 // S6: Sub(T1, T1) → regZero（零寄存器）
112 AscendC::Reg::Sub<T>(regT2, regT1, regT1, padMask);
113 
114 // S7: Adds(regZero, threshold) → regThresh（阈值寄存器）
115 AscendC::Reg::Adds<T>(regThresh, regT2, threshold, padMask);
116 
117 // S8: Compare(T1, regThresh, GT) → cmpMask
118 AscendC::Reg::Compare<T, AscendC::CMPMODE::GT>(cmpMask, regT1, regThresh, padMask);
119 
120 // S9: Select(cmpMask, input_gradients, T5) → dstReg
121 AscendC::Reg::Select<T>(dstReg, src1Reg, regT5, cmpMask);
122 
123 // --- Store: Reg → UB ---
124 AscendC::Reg::StoreAlign(dst, dstReg, aReg, padMask);
125 }
126}
127 
128// ============================================================
129// SoftplusV2GradKernel — Kernel 类
130// 模板参数: T (dtype), RANK (有效 rank)
131// ============================================================
132template <typename T, int64_t RANK>
133class SoftplusV2GradKernel {
134 // NDDMA 维度数 (最大 kMaxNdDmaDims), RANK 超出时外层走 Flat loop
135 static constexpr int64_t ND = (RANK <= kMaxNdDmaDims) ? RANK : kMaxNdDmaDims;
136 static constexpr uint32_t VL_F = AscendC::GetVecLen() / sizeof(float);
137 
138 AscendC::TPipe pipe_;
139 const SoftplusV2GradTilingData<RANK>* td_;
140 AscendC::GlobalTensor<T> gmIn_[kMaxInputSlots];
141 AscendC::GlobalTensor<T> gmOut_[kMaxOutputSlots];
142 AscendC::TBuf<AscendC::TPosition::VECCALC> buf_[kPhysNodes];
143 AscendC::MultiCopyParams<T, ND> nddmaParams_[kMaxInputSlots];
144 int64_t nddmaOuterIters_[kMaxInputSlots];
145 int64_t nddma_dims_;
146 float beta_;
147 float threshold_;
148 
149public:
150 __aicore__ inline void Init(GM_ADDR inputs[kMaxInputSlots], GM_ADDR outputs[kMaxOutputSlots],
151 const SoftplusV2GradTilingData<RANK>* td)
152 {
153 td_ = td;
154 for (int i = 0; i < kMaxInputSlots; i++)
155 gmIn_[i].SetGlobalBuffer((__gm__ T*)inputs[i]);
156 for (int i = 0; i < kMaxOutputSlots; i++)
157 gmOut_[i].SetGlobalBuffer((__gm__ T*)outputs[i]);
158 for (int i = 0; i < kPhysNodes; i++)
159 pipe_.InitBuffer(buf_[i], td_->per_buf_bytes);
160 
161 // 标量参数从 TilingData 读取
162 beta_ = td_->beta;
163 threshold_ = td_->threshold;
164 
165 // NDDMA 参数预计算
166 const int64_t* dstShape = td_->max_bro_shape;
167 int64_t k = td_->split.axis;
168 nddma_dims_ = (RANK - k <= ND) ? (RANK - k) : ND;
169 for (int inp = 0; inp < kMaxInputSlots; inp++) {
170 int64_t inner = 1;
171 int64_t nd = 0;
172 for (int64_t d = RANK - 1; d >= k && nd < ND; d--) {
173 nddmaParams_[inp].loopInfo.loopSize[nd] = (d == k) ? 0 : dstShape[d];
174 nddmaParams_[inp].loopInfo.loopSrcStride[nd] = td_->input_strides[inp][d];
175 nddmaParams_[inp].loopInfo.loopDstStride[nd] = inner;
176 nddmaParams_[inp].loopInfo.loopLpSize[nd] = 0;
177 nddmaParams_[inp].loopInfo.loopRpSize[nd] = 0;
178 inner *= (d == k) ? td_->split.a_i : dstShape[d];
179 nd++;
180 }
181 for (; nd < ND; nd++) {
182 nddmaParams_[inp].loopInfo.loopSize[nd] = 1;
183 nddmaParams_[inp].loopInfo.loopSrcStride[nd] = 0;
184 nddmaParams_[inp].loopInfo.loopDstStride[nd] = inner;
185 nddmaParams_[inp].loopInfo.loopLpSize[nd] = 0;
186 nddmaParams_[inp].loopInfo.loopRpSize[nd] = 0;
187 }
188 nddmaOuterIters_[inp] = 1;
189 for (int64_t d = k; d < RANK - nddma_dims_; d++)
190 nddmaOuterIters_[inp] *= (d == k) ? td_->split.a_i : dstShape[d];
191 }
192 }
193 
194 __aicore__ inline void Process()
195 {
196 int32_t evMTE2toV = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_V));
197 int32_t evVtoMTE2 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_MTE2));
198 int32_t evVtoMTE3 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_MTE3));
199 int32_t evMTE3toMTE2 = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE3_MTE2));
200 
201 int64_t start, end;
202 GetCoreRange(AscendC::GetBlockIdx(), td_->multicore.tiles_main, td_->multicore.cores_tail, start, end);
203 
204 int64_t inner_count = 1;
205 for (int64_t d = td_->split.axis + 1; d < RANK; d++)
206 inner_count *= td_->max_bro_shape[d];
207 
208 int64_t coord[RANK] = {};
209 
210 if constexpr (std::is_same_v<T, float>) {
211 ProcessFP32(start, end, inner_count, coord, evMTE2toV, evVtoMTE3, evMTE3toMTE2);
212 } else {
213 ProcessFP16BF16(start, end, inner_count, coord, evMTE2toV, evVtoMTE2, evVtoMTE3, evMTE3toMTE2);
214 }
215 }
216 
217private:
218 // ============================================================
219 // fp32 Process: P=3, 3 buffer slot, 3 sync pairs
220 // ============================================================
221 __aicore__ inline void ProcessFP32(int64_t start, int64_t end, int64_t inner_count, int64_t* coord,
222 int32_t evMTE2toV, int32_t evVtoMTE3, int32_t evMTE3toMTE2)
223 {
224 constexpr int UB_GRAD_O = 0, UB_SELF = 1, UB_RESULT = 2;
225 
226 for (int64_t flat = start; flat < end; flat++) {
227 int64_t a_i_seg = GetUBSplitRange(flat % td_->split.a_o, td_->split.a_o, td_->split.a_i,
228 td_->split.a_i_tail);
229 int64_t count = a_i_seg * inner_count;
230 FlatToEffectiveCoord(flat, td_->max_bro_shape, RANK, td_->split.axis, td_->split.a_i, td_->split.a_o,
231 coord);
232 
233 if (flat != start)
234 AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
235 
236 // S1: CopyIn input_gradients → UB_GRAD_O
237 CopyInBrc(coord, 0, UB_GRAD_O, a_i_seg);
238 
239 // S2: CopyIn input_features → UB_SELF
240 CopyInBrc(coord, 1, UB_SELF, a_i_seg);
241 
242 // MTE2→V sync
243 AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
244 AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
245 
246 // S3: VF(UB_SELF, UB_GRAD_O → UB_RESULT)
247 uint16_t rep = AscendC::CeilDivision(count, VL_F);
248 asc_vf_call<SoftplusV2GradVF<float>>((__ubuf__ float*)buf_[UB_RESULT].Get<float>().GetPhyAddr(),
249 (__ubuf__ float*)buf_[UB_SELF].Get<float>().GetPhyAddr(),
250 (__ubuf__ float*)buf_[UB_GRAD_O].Get<float>().GetPhyAddr(), count,
251 VL_F, rep, beta_, threshold_);
252 
253 // V→MTE3 sync
254 AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
255 AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
256 
257 // S4: CopyOut UB_RESULT → GM
258 CopyOutOne(coord, 0, UB_RESULT, a_i_seg);
259 
260 if (flat != end - 1)
261 AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
262 }
263 }
264 
265 // ============================================================
266 // fp16/bf16 Process: P=4, 4 buffer slot, 4 sync pairs
267 // ============================================================
268 __aicore__ inline void ProcessFP16BF16(int64_t start, int64_t end, int64_t inner_count, int64_t* coord,
269 int32_t evMTE2toV, int32_t evVtoMTE2, int32_t evVtoMTE3,
270 int32_t evMTE3toMTE2)
271 {
272 constexpr int UB_TEMP = 0;
273 constexpr int UB_SELF_FP32 = 1;
274 constexpr int UB_GRAD_O_FP32 = 2;
275 constexpr int UB_RESULT = 3;
276 
277 for (int64_t flat = start; flat < end; flat++) {
278 int64_t a_i_seg = GetUBSplitRange(flat % td_->split.a_o, td_->split.a_o, td_->split.a_i,
279 td_->split.a_i_tail);
280 int64_t count = a_i_seg * inner_count;
281 FlatToEffectiveCoord(flat, td_->max_bro_shape, RANK, td_->split.axis, td_->split.a_i, td_->split.a_o,
282 coord);
283 
284 if (flat != start)
285 AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
286 
287 // S1a: CopyIn input_gradients(fp16) → UB_TEMP
288 CopyInBrc(coord, 0, UB_TEMP, a_i_seg);
289 AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
290 AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
291 
292 // S1b: Cast UB_TEMP(fp16)→UB_GRAD_O_FP32(fp32), CAST_NONE（扩精度）
293 AscendC::Cast(buf_[UB_GRAD_O_FP32].template Get<float>(), buf_[UB_TEMP].template Get<T>(),
294 AscendC::RoundMode::CAST_NONE, count);
295 AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(evVtoMTE2);
296 AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(evVtoMTE2);
297 
298 // S2a: CopyIn input_features(fp16)→UB_TEMP（复用）
299 CopyInBrc(coord, 1, UB_TEMP, a_i_seg);
300 AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
301 AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(evMTE2toV);
302 
303 // S2b: Cast UB_TEMP(fp16)→UB_SELF_FP32(fp32)
304 AscendC::Cast(buf_[UB_SELF_FP32].template Get<float>(), buf_[UB_TEMP].template Get<T>(),
305 AscendC::RoundMode::CAST_NONE, count);
306 AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(evVtoMTE2);
307 AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(evVtoMTE2);
308 
309 // S3: VF(UB_SELF_FP32, UB_GRAD_O_FP32 → UB_RESULT)
310 uint16_t rep = AscendC::CeilDivision(count, VL_F);
311 asc_vf_call<SoftplusV2GradVF<float>>((__ubuf__ float*)buf_[UB_RESULT].Get<float>().GetPhyAddr(),
312 (__ubuf__ float*)buf_[UB_SELF_FP32].Get<float>().GetPhyAddr(),
313 (__ubuf__ float*)buf_[UB_GRAD_O_FP32].Get<float>().GetPhyAddr(), count,
314 VL_F, rep, beta_, threshold_);
315 
316 // S4: Cast UB_RESULT(fp32)→UB_TEMP(fp16), CAST_RINT（缩精度）
317 AscendC::Cast(buf_[UB_TEMP].template Get<T>(), buf_[UB_RESULT].template Get<float>(),
318 AscendC::RoundMode::CAST_RINT, count);
319 AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
320 AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(evVtoMTE3);
321 
322 // S5: CopyOut UB_TEMP → GM
323 CopyOutOne(coord, 0, UB_TEMP, a_i_seg);
324 
325 if (flat != end - 1)
326 AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(evMTE3toMTE2);
327 }
328 }
329 
330 // ============================================================
331 // CopyInBrc — NDDMA 多维广播搬运
332 // ============================================================
333 __aicore__ inline void CopyInBrc(const int64_t* coord, int inputIdx, int slot, int64_t a_i_seg)
334 {
335 int64_t k = td_->split.axis;
336 int64_t off = CalcOffset(coord, td_->input_strides[inputIdx], RANK);
337 const int64_t* dstShape = td_->max_bro_shape;
338 
339 auto params = nddmaParams_[inputIdx];
340 int64_t k_nd = RANK - 1 - k;
341 int64_t inner = 1;
342 for (int64_t nd = 0; nd < ND; nd++) {
343 if (nd == k_nd)
344 params.loopInfo.loopSize[nd] = a_i_seg;
345 params.loopInfo.loopDstStride[nd] = inner;
346 inner *= params.loopInfo.loopSize[nd];
347 }
348 
349 static constexpr AscendC::NdDmaConfig cfg = {false, AscendC::NdDmaConfig::unsetPad,
350 AscendC::NdDmaConfig::unsetPad, false};
351 
352 if constexpr (RANK <= kMaxNdDmaDims) {
353 AscendC::DataCopy<T, ND, cfg>(buf_[slot].Get<T>(), gmIn_[inputIdx][off], params);
354 } else {
355 AscendC::LocalTensor<T> buf = buf_[slot].Get<T>();
356 int64_t elem_base = off;
357 for (int64_t oi = 0; oi < nddmaOuterIters_[inputIdx]; oi++) {
358 int64_t elem_adj = 0, tmp = oi;
359 for (int64_t d = RANK - nddma_dims_ - 1; d >= k; d--) {
360 int64_t sz = (d == k) ? a_i_seg : dstShape[d];
361 elem_adj += (tmp % sz) * td_->input_strides[inputIdx][d];
362 tmp /= sz;
363 }
364 AscendC::DataCopy<T, ND, cfg>(buf[oi * inner], gmIn_[inputIdx][elem_base + elem_adj], params);
365 }
366 }
367 }
368 
369 // ============================================================
370 // CopyOutOne — DataCopyPad 低维搬出
371 // ============================================================
372 __aicore__ inline void CopyOutOne(const int64_t* coord, int outputIdx, int slot, int64_t a_i_seg)
373 {
374 int64_t off = CalcOffset(coord, td_->output_strides[outputIdx], RANK);
375 int64_t cnt = CalcTransferCount(td_->output_shapes[outputIdx], RANK, td_->split.axis, a_i_seg);
376 AscendC::DataCopyExtParams extParams;
377 extParams.blockCount = 1;
378 extParams.blockLen = cnt * sizeof(T);
379 extParams.srcStride = 0;
380 extParams.dstStride = 0;
381 AscendC::DataCopyPad(gmOut_[outputIdx][off], buf_[slot].Get<T>(), extParams);
382 }
383};
Aactivation/softplus_v2_grad/op_kernel/arch35/softplus_v2_grad_struct.h+29-0
@@ -0,0 +1,29 @@
1/**
2 * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.
9 *
10 * Generated By CANNBot
11 */
12 
13// SoftplusV2Grad TilingKey 模板参数
14// 位置: operators/softplus_v2_grad/op_kernel/arch35/softplus_v2_grad_struct.h
15#ifndef SOFTPLUS_V2_GRAD_STRUCT_H_
16#define SOFTPLUS_V2_GRAD_STRUCT_H_
17 
18#include "ascendc/host_api/tiling/template_argument.h"
19 
20#define SOFTPLUS_V2_GRAD_RANK_4 4
21#define SOFTPLUS_V2_GRAD_RANK_8 8
22 
23ASCENDC_TPL_ARGS_DECL(SoftplusV2Grad, ASCENDC_TPL_UINT_DECL(RANK, 8, ASCENDC_TPL_UI_LIST, SOFTPLUS_V2_GRAD_RANK_4,
24 SOFTPLUS_V2_GRAD_RANK_8));
25 
26ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, SOFTPLUS_V2_GRAD_RANK_4)),
27 ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(RANK, ASCENDC_TPL_UI_LIST, SOFTPLUS_V2_GRAD_RANK_8)));
28 
29#endif // SOFTPLUS_V2_GRAD_STRUCT_H_
Aactivation/softplus_v2_grad/op_kernel/arch35/softplus_v2_grad_tiling_struct.h+64-0
@@ -0,0 +1,64 @@
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// SoftplusV2Grad TilingData — 按 rank 模板化，体积分两档
14#pragma once
15#include <cstdint>
16 
17// === 算子特定常量 ===
18// 2 输入 (input_gradients, input_features), 1 输出 (output_backprops)
19constexpr int64_t kMaxInputSlots = 2;
20constexpr int64_t kMaxOutputSlots = 1;
21// 物理存活节点 P = 4 (覆盖 fp16/bf16 含 Cast 路径)
22constexpr int64_t kPhysNodes = 4;
23 
24// UB 对齐粒度: TBuf 硬件要求 32B 对齐
25constexpr int64_t kUbAlignBytes = 32;
26// NDDMA 最大维数: 超出该维数的外层走 Flat loop
27constexpr int64_t kMaxNdDmaDims = 5;
28// 默认标量属性: 必须与 op_host/softplus_v2_grad_def.cpp 中 Attr 注册的默认值保持一致
29constexpr float kDefaultBeta = 1.0f;
30constexpr float kDefaultThreshold = 20.0f;
31// attr 索引: 与 def.cpp 中 Attr 注册顺序一致 (0=beta, 1=threshold)
32constexpr int64_t kAttrIdxBeta = 0;
33constexpr int64_t kAttrIdxThreshold = 1;
34 
35struct SplitResult {
36 int64_t axis;
37 int64_t a_i;
38 int64_t a_o;
39 int64_t a_i_tail;
40};
41 
42struct MultiCoreResult {
43 int64_t num_cores;
44 int64_t total_tiles;
45 int64_t tiles_main;
46 int64_t cores_tail;
47};
48 
49template <int64_t kRank>
50struct SoftplusV2GradTilingData {
51 SplitResult split;
52 MultiCoreResult multicore;
53 int64_t rank; // 实际 rank (0~8)，Kernel 运行期读取
54 int64_t per_buf_bytes; // UB/P 向下对齐 32B，Kernel 用此初始化 TBuf
55 int64_t max_bro_shape[kRank];
56 int64_t num_inputs;
57 int64_t num_outputs;
58 int64_t input_shapes[kMaxInputSlots][kRank];
59 int64_t input_strides[kMaxInputSlots][kRank];
60 int64_t output_shapes[kMaxOutputSlots][kRank];
61 int64_t output_strides[kMaxOutputSlots][kRank];
62 float beta; // 标量属性: softplus 陡峭度 (默认 1.0)
63 float threshold; // 标量属性: 数值稳定性阈值 (默认 20.0)
64};
Mactivation/softplus_v2_grad/tests/CMakeLists.txt+2-2
@@ -1,8 +1,8 @@
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#/
Mactivation/softplus_v2_grad/tests/ut/CMakeLists.txt+3-3
@@ -1,8 +1,8 @@
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#/
@@ -15,4 +15,4 @@ foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
15 if(EXISTS "${SUB_DIR}/CMakeLists.txt")15 if(EXISTS "${SUB_DIR}/CMakeLists.txt")
16 add_subdirectory(${SUB_DIR})16 add_subdirectory(${SUB_DIR})
17 endif()17 endif()
18endforeach()18endforeach()
Mactivation/softplus_v2_grad/tests/ut/op_host/CMakeLists.txt+10-3
@@ -1,13 +1,14 @@
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 
10message(STATUS "=== Debug: start ops.activation.softplus_v2_grad.tests.ut.CMakeLists.txt ")10message(STATUS "=== Debug: start ops.activation.softplus_v2_grad.tests.ut.op_host.CMakeLists.txt ")
11# 递归子目录（op_api UT 挂在 op_api/ 子目录下，此处保持不变）
11file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)12file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
12message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")13message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
13foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})14foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
@@ -15,3 +16,9 @@ foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
15 add_subdirectory(${SUB_DIR})16 add_subdirectory(${SUB_DIR})
16 endif()17 endif()
17endforeach()18endforeach()
19 
20# 本目录下的 tiling / infershape UT（DIR 为当前目录，仅收录本层 *.cpp，不递归子目录）
21if(UT_TEST_ALL OR OP_HOST_UT)
22 add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
23 add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
24endif()
Aactivation/softplus_v2_grad/tests/ut/op_host/arch35/test_softplus_v2_grad_tiling.cpp+490-0
@@ -0,0 +1,490 @@
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// SoftplusV2Grad Tiling UT — 纯函数 + 框架级全覆盖
14 
15#include <iostream>
16#include <vector>
17#include <sstream>
18 
19#include <gtest/gtest.h>
20#include "log/log.h"
21#include "kernel_run_context_facker.h"
22#include "test_cube_util.h"
23#include "exe_graph/runtime/storage_format.h"
24#include "exe_graph/runtime/storage_shape.h"
25#include "platform/platform_infos_def.h"
26#include "ut_op_util.h"
27#include "../../../../op_host/arch35/softplus_v2_grad_tiling_arch35.h"
28#include "../../../../op_kernel/arch35/softplus_v2_grad_struct.h"
29 
30using namespace std;
31using namespace ut_util;
32 
33class SoftplusV2GradTilingTest : public testing::Test {
34protected:
35 static void SetUpTestCase() { std::cout << "SoftplusV2GradTilingTest SetUp" << std::endl; }
36 static void TearDownTestCase() { std::cout << "SoftplusV2GradTilingTest TearDown" << std::endl; }
37};
38 
39// ============================================================
40// 1. CheckBroadcastShape
41// ============================================================
42TEST_F(SoftplusV2GradTilingTest, broadcast_same_shape_single_in_out)
43{
44 // input {2,3}, output {2,3} — 完全一致
45 std::vector<std::vector<int64_t>> padded_in = {{2, 3}};
46 std::vector<std::vector<int64_t>> padded_out = {{2, 3}};
47 EXPECT_TRUE(softplus_v2_grad::CheckBroadcastShape(padded_in, padded_out, 2));
48}
49 
50TEST_F(SoftplusV2GradTilingTest, broadcast_two_inputs_same)
51{
52 // 2 个 input 都是 {4,4}, 1 个 output {4,4}
53 std::vector<std::vector<int64_t>> padded_in = {{4, 4}, {4, 4}};
54 std::vector<std::vector<int64_t>> padded_out = {{4, 4}};
55 EXPECT_TRUE(softplus_v2_grad::CheckBroadcastShape(padded_in, padded_out, 2));
56}
57 
58TEST_F(SoftplusV2GradTilingTest, broadcast_input_dim_is_one)
59{
60 // input {1,4}, output {4,4} — input=1 的维度不冲突
61 std::vector<std::vector<int64_t>> padded_in = {{1, 4}};
62 std::vector<std::vector<int64_t>> padded_out = {{4, 4}};
63 EXPECT_TRUE(softplus_v2_grad::CheckBroadcastShape(padded_in, padded_out, 2));
64}
65 
66TEST_F(SoftplusV2GradTilingTest, broadcast_output_dim_is_one)
67{
68 // input {4,4}, output {1,4} — output=1 的维度不冲突
69 std::vector<std::vector<int64_t>> padded_in = {{4, 4}};
70 std::vector<std::vector<int64_t>> padded_out = {{1, 4}};
71 EXPECT_TRUE(softplus_v2_grad::CheckBroadcastShape(padded_in, padded_out, 2));
72}
73 
74TEST_F(SoftplusV2GradTilingTest, broadcast_two_inputs_different_non_one)
75{
76 // 两个 input 在同一维上取值不同且都不是 1 → 不兼容
77 std::vector<std::vector<int64_t>> padded_in = {{2, 3}, {4, 3}};
78 std::vector<std::vector<int64_t>> padded_out = {{4, 3}};
79 EXPECT_FALSE(softplus_v2_grad::CheckBroadcastShape(padded_in, padded_out, 2));
80}
81 
82TEST_F(SoftplusV2GradTilingTest, broadcast_input_output_diff_non_one)
83{
84 // input {2,3}, output {4,3} 且 input dim0=2 ≠ 1, output dim0=4 ≠ 1 → 不兼容
85 std::vector<std::vector<int64_t>> padded_in = {{2, 3}};
86 std::vector<std::vector<int64_t>> padded_out = {{4, 3}};
87 EXPECT_FALSE(softplus_v2_grad::CheckBroadcastShape(padded_in, padded_out, 2));
88}
89 
90// ============================================================
91// 2. PadAndSqueeze
92// ============================================================
93 
94TEST_F(SoftplusV2GradTilingTest, pad_squeeze_same_rank)
95{
96 // 两个 input 同 rank，output 同 rank → 无需补
97 std::vector<std::vector<int64_t>> in_shapes = {{2, 3}, {2, 3}};
98 std::vector<std::vector<int64_t>> out_shapes = {{2, 3}};
99 std::vector<int64_t> max_bro;
100 std::vector<std::vector<int64_t>> norm_in, norm_out;
101 EXPECT_TRUE(softplus_v2_grad::PadAndSqueeze(in_shapes, out_shapes, max_bro, norm_in, norm_out));
102 EXPECT_EQ(max_bro, std::vector<int64_t>({2, 3}));
103 EXPECT_EQ(norm_in.size(), 2u);
104 EXPECT_EQ(norm_out.size(), 1u);
105 EXPECT_EQ(norm_in[0], std::vector<int64_t>({2, 3}));
106 EXPECT_EQ(norm_out[0], std::vector<int64_t>({2, 3}));
107}
108 
109TEST_F(SoftplusV2GradTilingTest, pad_squeeze_different_rank)
110{
111 // input: {2,3} (rank 2), output: {2} (rank 1) → output 右对齐补 1
112 std::vector<std::vector<int64_t>> in_shapes = {{2, 3}};
113 std::vector<std::vector<int64_t>> out_shapes = {{3}};
114 std::vector<int64_t> max_bro;
115 std::vector<std::vector<int64_t>> norm_in, norm_out;
116 EXPECT_TRUE(softplus_v2_grad::PadAndSqueeze(in_shapes, out_shapes, max_bro, norm_in, norm_out));
117 // max_rank=2: input {2,3}, output padded to {1,3}
118 // broadcast result should be {2,3}
119 EXPECT_EQ(max_bro, std::vector<int64_t>({2, 3}));
120 EXPECT_EQ(norm_in[0], std::vector<int64_t>({2, 3}));
121 // output padded dim0=1 不挤压（因为有非1的输入维）→ 保持 {1,3}
122 EXPECT_EQ(norm_out[0], std::vector<int64_t>({1, 3}));
123}
124 
125TEST_F(SoftplusV2GradTilingTest, pad_squeeze_all_ones)
126{
127 // 全部是 1 → max_bro_shape 为 {1}
128 std::vector<std::vector<int64_t>> in_shapes = {{1}};
129 std::vector<std::vector<int64_t>> out_shapes = {{1}};
130 std::vector<int64_t> max_bro;
131 std::vector<std::vector<int64_t>> norm_in, norm_out;
132 EXPECT_TRUE(softplus_v2_grad::PadAndSqueeze(in_shapes, out_shapes, max_bro, norm_in, norm_out));
133 EXPECT_EQ(max_bro, std::vector<int64_t>({1}));
134 EXPECT_EQ(norm_in[0], std::vector<int64_t>({1}));
135 EXPECT_EQ(norm_out[0], std::vector<int64_t>({1}));
136}
137 
138TEST_F(SoftplusV2GradTilingTest, pad_squeeze_4d)
139{
140 // 4D tensor: {1,4,4,8}
141 std::vector<std::vector<int64_t>> in_shapes = {{1, 4, 4, 8}};
142 std::vector<std::vector<int64_t>> out_shapes = {{1, 4, 4, 8}};
143 std::vector<int64_t> max_bro;
144 std::vector<std::vector<int64_t>> norm_in, norm_out;
145 EXPECT_TRUE(softplus_v2_grad::PadAndSqueeze(in_shapes, out_shapes, max_bro, norm_in, norm_out));
146 // dim0=1 squeezed out, 剩下 {4,4,8}
147 EXPECT_EQ(max_bro, std::vector<int64_t>({4, 4, 8}));
148}
149 
150// ============================================================
151// 3. FindSplitAxis
152// ============================================================
153 
154TEST_F(SoftplusV2GradTilingTest, split_axis_small_shape_no_split)
155{
156 // shape {2,3}, dtype_size=4, ub=262144, phys_nodes=4
157 // per_buf_bytes = (262144/4) & ~31 = 65504
158 // per_buf_elems = 65504/4 = 16376
159 // inner: k=1: 3*1=3 ≤ 16376, inner*=3 → k=0: 2*3=6 ≤ 16376
160 // 不切: a_i=2, a_o=1, a_i_tail=2, axis=0
161 std::vector<int64_t> shape = {2, 3};
162 SplitResult out;
163 EXPECT_TRUE(softplus_v2_grad::FindSplitAxis(shape, /*dtype_size=*/4, /*ub=*/262144, /*phys_nodes=*/4, out));
164 EXPECT_EQ(out.axis, 0);
165 EXPECT_EQ(out.a_o, 1);
166 EXPECT_EQ(out.a_i, 2);
167 EXPECT_EQ(out.a_i_tail, 2);
168}
169 
170TEST_F(SoftplusV2GradTilingTest, split_axis_large_shape_needs_split)
171{
172 // shape {1, 65536}, dtype_size=4, ub=262144, phys_nodes=4
173 // per_buf_bytes = (262144/4) & ~31 = 65536; per_buf_elems = 65536/4 = 16384
174 // k=1: 65536 > 16384 → split on axis 1
175 // a_i = 16384/1 = 16384; a_o = ceil(65536/16384) = 4; rem = 0 → a_i_tail = 16384
176 std::vector<int64_t> shape = {1, 65536};
177 SplitResult out;
178 EXPECT_TRUE(softplus_v2_grad::FindSplitAxis(shape, /*dtype_size=*/4, /*ub=*/262144, /*phys_nodes=*/4, out));
179 EXPECT_EQ(out.axis, 1);
180 EXPECT_EQ(out.a_o, 4);
181 EXPECT_EQ(out.a_i, 16384);
182 EXPECT_EQ(out.a_i_tail, 16384);
183}
184 
185TEST_F(SoftplusV2GradTilingTest, split_axis_invalid_dtype_size)
186{
187 std::vector<int64_t> shape = {2, 3};
188 SplitResult out;
189 EXPECT_FALSE(softplus_v2_grad::FindSplitAxis(shape, /*dtype_size=*/0, /*ub=*/262144, /*phys_nodes=*/4, out));
190}
191 
192TEST_F(SoftplusV2GradTilingTest, split_axis_invalid_phys_nodes)
193{
194 std::vector<int64_t> shape = {2, 3};
195 SplitResult out;
196 EXPECT_FALSE(softplus_v2_grad::FindSplitAxis(shape, /*dtype_size=*/4, /*ub=*/262144, /*phys_nodes=*/0, out));
197}
198 
199TEST_F(SoftplusV2GradTilingTest, split_axis_very_small_ub)
200{
201 // ub=64: per_buf_bytes = (64/4) & ~31 = 0 → per_buf_elems = 0 → 直接失败
202 std::vector<int64_t> shape = {2, 3};
203 SplitResult out;
204 EXPECT_FALSE(softplus_v2_grad::FindSplitAxis(shape, /*dtype_size=*/4, /*ub=*/64, /*phys_nodes=*/4, out));
205}
206 
207// ============================================================
208// 4. MultiCoreSplit
209// ============================================================
210 
211TEST_F(SoftplusV2GradTilingTest, multicore_few_tiles)
212{
213 // split.axis=0, shape {2,3}: outer_prod=1, total_tiles = 1*1 = 1
214 // max_cores=64, num_cores = min(1,64) = 1
215 SplitResult split = {0, 2, 1, 2};
216 std::vector<int64_t> shape = {2, 3};
217 MultiCoreResult out;
218 EXPECT_TRUE(softplus_v2_grad::MultiCoreSplit(shape, split, 64, out));
219 EXPECT_EQ(out.num_cores, 1);
220 EXPECT_EQ(out.total_tiles, 1);
221 EXPECT_EQ(out.tiles_main, 1);
222 EXPECT_EQ(out.cores_tail, 0);
223}
224 
225TEST_F(SoftplusV2GradTilingTest, multicore_many_tiles)
226{
227 // split.axis=1, shape {1,65536}: outer_prod=1, a_o=5
228 // total_tiles = 1*5 = 5, num_cores=5
229 SplitResult split = {1, 16376, 5, 4};
230 std::vector<int64_t> shape = {1, 65536};
231 MultiCoreResult out;
232 EXPECT_TRUE(softplus_v2_grad::MultiCoreSplit(shape, split, 64, out));
233 EXPECT_EQ(out.num_cores, 5);
234 EXPECT_EQ(out.total_tiles, 5);
235 EXPECT_EQ(out.tiles_main, 1);
236 EXPECT_EQ(out.cores_tail, 0);
237}
238 
239TEST_F(SoftplusV2GradTilingTest, multicore_caps_at_max_cores)
240{
241 // shape {100, 8}: axis=0, a_o=5, outer_prod=1, total_tiles=5
242 // max_cores=3: num_cores=min(5,3)=3 → tiles_main=5/3=1, cores_tail=5%3=2
243 SplitResult split = {0, 2, 5, 2};
244 std::vector<int64_t> shape = {100, 8};
245 MultiCoreResult out;
246 EXPECT_TRUE(softplus_v2_grad::MultiCoreSplit(shape, split, 3, out));
247 EXPECT_EQ(out.num_cores, 3);
248 EXPECT_EQ(out.total_tiles, 5);
249 EXPECT_EQ(out.tiles_main, 1);
250 EXPECT_EQ(out.cores_tail, 2);
251}
252 
253TEST_F(SoftplusV2GradTilingTest, multicore_tail_cores)
254{
255 // shape {5}: split.axis=0, a_o=1, total_tiles=1 → outer_prod=1
256 // max_cores=3: num_cores=1 → tiles_main=1, cores_tail=0
257 SplitResult split = {0, 3, 2, 2}; // 5 元素: a_o=2, a_i=3, a_i_tail=2
258 std::vector<int64_t> shape = {5};
259 MultiCoreResult out;
260 EXPECT_TRUE(softplus_v2_grad::MultiCoreSplit(shape, split, 3, out));
261 EXPECT_EQ(out.total_tiles, 2); // outer_prod(1)*a_o(2)=2
262 EXPECT_EQ(out.num_cores, 2); // min(2,3)=2
263 EXPECT_EQ(out.tiles_main, 1); // 2/2=1
264 EXPECT_EQ(out.cores_tail, 0); // 2%2=0
265}
266 
267TEST_F(SoftplusV2GradTilingTest, multicore_zero_total_tiles)
268{
269 // shape {0} (空 tensor): outer_prod=1, a_o=0 → total_tiles=0, num_cores=0 → false
270 SplitResult split = {0, 1, 0, 1};
271 std::vector<int64_t> shape = {0};
272 MultiCoreResult out;
273 EXPECT_FALSE(softplus_v2_grad::MultiCoreSplit(shape, split, 64, out));
274}
275 
276// ============================================================
277// 5. PrecomputeStrides
278// ============================================================
279 
280TEST_F(SoftplusV2GradTilingTest, strides_1d)
281{
282 std::vector<int64_t> shape = {5};
283 std::vector<int64_t> strides;
284 EXPECT_TRUE(softplus_v2_grad::PrecomputeStrides(shape, strides));
285 // 1D: 最内维 stride=1
286 EXPECT_EQ(strides, std::vector<int64_t>({1}));
287}
288 
289TEST_F(SoftplusV2GradTilingTest, strides_2d)
290{
291 std::vector<int64_t> shape = {3, 4};
292 std::vector<int64_t> strides;
293 EXPECT_TRUE(softplus_v2_grad::PrecomputeStrides(shape, strides));
294 // d=1: s[1]=4 → prod=1 → stride[1]=1; d=0: prod=s[1]=4 → stride[0]=4
295 EXPECT_EQ(strides, std::vector<int64_t>({4, 1}));
296}
297 
298TEST_F(SoftplusV2GradTilingTest, strides_dim_one_yields_zero_stride)
299{
300 // dim=1 → stride=0 (broadcast dim)
301 std::vector<int64_t> shape = {1, 8, 1};
302 std::vector<int64_t> strides;
303 EXPECT_TRUE(softplus_v2_grad::PrecomputeStrides(shape, strides));
304 // d=2: s[2]=1 → 0; d=1: s[1]=8 → prod=1 → 1; d=0: s[0]=1 → 0
305 EXPECT_EQ(strides, std::vector<int64_t>({0, 1, 0}));
306}
307 
308TEST_F(SoftplusV2GradTilingTest, strides_3d)
309{
310 std::vector<int64_t> shape = {2, 3, 4};
311 std::vector<int64_t> strides;
312 EXPECT_TRUE(softplus_v2_grad::PrecomputeStrides(shape, strides));
313 // d=2: s[2]=4 → 1; d=1: prod=4 → 4; d=0: prod=3*4=12 → 12
314 EXPECT_EQ(strides, std::vector<int64_t>({12, 4, 1}));
315}
316 
317// ============================================================
318// 集成: PadAndSqueeze + CheckBroadcastShape 联合正确性
319// ============================================================
320 
321TEST_F(SoftplusV2GradTilingTest, integration_pad_and_check)
322{
323 // 确保 PadAndSqueeze 产出的归一化 shape 能通过 CheckBroadcastShape
324 std::vector<std::vector<int64_t>> in_shapes = {{4, 8}, {4, 8}};
325 std::vector<std::vector<int64_t>> out_shapes = {{4, 8}};
326 std::vector<int64_t> max_bro;
327 std::vector<std::vector<int64_t>> norm_in, norm_out;
328 ASSERT_TRUE(softplus_v2_grad::PadAndSqueeze(in_shapes, out_shapes, max_bro, norm_in, norm_out));
329 EXPECT_TRUE(softplus_v2_grad::CheckBroadcastShape(norm_in, norm_out, max_bro.size()));
330}
331 
332// ============================================================
333// 集成: FindSplitAxis + MultiCoreSplit 一致性
334// ============================================================
335 
336TEST_F(SoftplusV2GradTilingTest, integration_split_and_multicore)
337{
338 // 验证 split 和 multicore 对同一 shape 的结果一致性
339 std::vector<int64_t> shape = {32, 128};
340 SplitResult split;
341 ASSERT_TRUE(softplus_v2_grad::FindSplitAxis(shape, /*dtype_size=*/4, /*ub=*/262144, /*phys_nodes=*/4, split));
342 MultiCoreResult mc;
343 ASSERT_TRUE(softplus_v2_grad::MultiCoreSplit(shape, split, 64, mc));
344 // 确保多核划分 cover 所有 tile
345 EXPECT_EQ(mc.total_tiles, mc.num_cores * mc.tiles_main + mc.cores_tail);
346 // num_cores ≤ 64
347 EXPECT_LE(mc.num_cores, 64);
348}
349 
350// ============================================================
351// 框架级 Tiling UT — 使用 TilingContextFaker
352// 覆盖 GetShapeInfo / DoTilingAndSet / RunTiling / FillShapesAndStrides / LogTilingData
353// ============================================================
354 
355class SoftplusV2GradTilingIntegrationTest : public testing::Test {
356protected:
357 static void SetUpTestCase() { std::cout << "SoftplusV2GradTilingIntegration SetUp" << std::endl; }
358 static void TearDownTestCase() { std::cout << "SoftplusV2GradTilingIntegration TearDown" << std::endl; }
359};
360 
361// 辅助: 验证 TilingData 基本合法性
362// SoftplusV2GradTilingData 内存布局:
363// SplitResult(4xint64) + MultiCoreResult(4xint64) + rank(1) + per_buf_bytes(1) = 10 个 int64
364// raw[8] = rank, raw[9] = per_buf_bytes, raw[4] = multicore.num_cores
365static void VerifyTilingDataValid(const gert::TilingData* tiling_data)
366{
367 ASSERT_NE(tiling_data, nullptr);
368 ASSERT_GT(tiling_data->GetDataSize(), (size_t)0);
369 auto* raw = reinterpret_cast<const int64_t*>(tiling_data->GetData());
370 int64_t rank_val = raw[8];
371 int64_t per_buf = raw[9];
372 int64_t num_cores = raw[4];
373 EXPECT_GT(rank_val, 0);
374 EXPECT_GT(per_buf, 0);
375 EXPECT_GE(num_cores, 1);
376 EXPECT_LE(num_cores, 64);
377}
378 
379// 通用 tiling 全链路测试模板
380static void RunTilingIntegrationTest(const gert::StorageShape& in_a, const gert::StorageShape& in_b,
381 const gert::StorageShape& out, ge::DataType dtype, ge::DataType out_dtype,
382 ge::graphStatus expect_status)
383{
384 gert::StorageShape in_copy_a = in_a;
385 gert::StorageShape in_copy_b = in_b;
386 gert::StorageShape out_copy = out;
387 std::map<std::string, std::string> soc_infos, aicore_spec, intrinsics;
388 std::map<std::string, std::string> soc_version = {{"Short_SoC_version", "Ascend950"}, {"NpuArch", "3510"}};
389 std::string compile_str = R"({
390 "hardware_info": {
391 "BT_SIZE": 0, "load3d_constraints": "1",
392 "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
393 "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
394 "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
395 "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": 64
396 }
397 })";
398 GetPlatFormInfos(compile_str.c_str(), soc_infos, aicore_spec, intrinsics);
399 
400 fe::PlatFormInfos platform_info;
401 platform_info.Init();
402 optiling::SoftplusV2GradCompileInfo compile_info;
403 std::string op_type("SoftplusV2Grad");
404 
405 auto op_impl = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str());
406 ASSERT_NE(op_impl, nullptr);
407 auto tiling_func = op_impl->tiling;
408 auto tiling_parse_func = op_impl->tiling_parse;
409 
410 auto kernel_holder = gert::KernelRunContextFaker()
411 .KernelIONum(1, 1)
412 .Inputs({const_cast<char*>(compile_str.c_str()), reinterpret_cast<void*>(&platform_info)})
413 .Outputs({&compile_info})
414 .Build();
415 ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
416 kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
417 kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
418 kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
419 kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
420 intrinsics);
421 kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("version", soc_version);
422 ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
423 
424 auto param = gert::TilingData::CreateCap(8192);
425 ASSERT_NE(param, nullptr);
426 auto ws_holder = gert::ContinuousVector::Create<size_t>(4096);
427 auto ws_size = reinterpret_cast<gert::ContinuousVector*>(ws_holder.get());
428 
429 auto holder = gert::TilingContextFaker()
430 .SetOpType(op_type)
431 .NodeIoNum(2, 1)
432 .IrInstanceNum({1})
433 .InputShapes({&in_copy_a, &in_copy_b})
434 .OutputShapes({&out_copy})
435 .CompileInfo(&compile_info)
436 .PlatformInfo(reinterpret_cast<char*>(&platform_info))
437 .NodeInputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
438 .NodeInputTd(1, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
439 .NodeOutputTd(0, out_dtype, ge::FORMAT_ND, ge::FORMAT_ND)
440 .TilingData(param.get())
441 .Workspace(ws_size)
442 .Build();
443 auto* ctx = holder.GetContext<gert::TilingContext>();
444 ASSERT_NE(ctx->GetPlatformInfo(), nullptr);
445 ctx->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
446 ctx->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
447 ctx->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
448 ctx->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
449 
450 EXPECT_EQ(tiling_func(ctx), expect_status);
451 if (expect_status == ge::GRAPH_SUCCESS) {
452 VerifyTilingDataValid(ctx->GetRawTilingData());
453 }
454}
455 
456TEST_F(SoftplusV2GradTilingIntegrationTest, tiling_fp32_2d)
457{
458 gert::StorageShape shape = {{4, 8}, {4, 8}};
459 RunTilingIntegrationTest(shape, shape, shape, ge::DT_FLOAT, ge::DT_FLOAT, ge::GRAPH_SUCCESS);
460}
461 
462TEST_F(SoftplusV2GradTilingIntegrationTest, tiling_fp16_4d)
463{
464 gert::StorageShape shape = {{1, 4, 4, 8}, {1, 4, 4, 8}};
465 RunTilingIntegrationTest(shape, shape, shape, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::GRAPH_SUCCESS);
466}
467 
468TEST_F(SoftplusV2GradTilingIntegrationTest, tiling_bf16_5d_rank8)
469{
470 gert::StorageShape shape = {{1, 1, 2, 3, 4}, {1, 1, 2, 3, 4}};
471 RunTilingIntegrationTest(shape, shape, shape, ge::DT_BF16, ge::DT_BF16, ge::GRAPH_SUCCESS);
472}
473 
474TEST_F(SoftplusV2GradTilingIntegrationTest, tiling_fp32_1d)
475{
476 gert::StorageShape shape = {{256}, {256}};
477 RunTilingIntegrationTest(shape, shape, shape, ge::DT_FLOAT, ge::DT_FLOAT, ge::GRAPH_SUCCESS);
478}
479 
480TEST_F(SoftplusV2GradTilingIntegrationTest, tiling_fp32_large_2d)
481{
482 gert::StorageShape shape = {{1, 65536}, {1, 65536}};
483 RunTilingIntegrationTest(shape, shape, shape, ge::DT_FLOAT, ge::DT_FLOAT, ge::GRAPH_SUCCESS);
484}
485 
486TEST_F(SoftplusV2GradTilingIntegrationTest, tiling_unsupported_dtype_failed)
487{
488 gert::StorageShape shape = {{2, 3}, {2, 3}};
489 RunTilingIntegrationTest(shape, shape, shape, ge::DT_INT32, ge::DT_INT32, ge::GRAPH_FAILED);
490}
Mactivation/softplus_v2_grad/tests/ut/op_host/op_api/CMakeLists.txt+2-2
@@ -1,8 +1,8 @@
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#/
Mactivation/softplus_v2_grad/tests/ut/op_host/op_api/test_aclnn_softplus_backward.cpp+1-1
@@ -9,7 +9,7 @@
9 */9 */
10 10 
11#include "gtest/gtest.h"11#include "gtest/gtest.h"
12#include "../../../../op_host/op_api/aclnn_softplus_backward.h"12#include "../../../../op_api/aclnn_softplus_backward.h"
13#include "op_api_ut_common/tensor_desc.h"13#include "op_api_ut_common/tensor_desc.h"
14#include "op_api_ut_common/scalar_desc.h"14#include "op_api_ut_common/scalar_desc.h"
15#include "op_api_ut_common/op_api_ut.h"15#include "op_api_ut_common/op_api_ut.h"
Aactivation/softplus_v2_grad/tests/ut/op_host/test_softplus_v2_grad_infershape.cpp+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 * Generated By CANNBot
11 */
12 
13#include <gtest/gtest.h>
14#include <iostream>
15#include "infershape_test_util.h"
16#include "ut_op_common.h"
17#include "log/log.h"
18#include "../../../op_graph/softplus_v2_grad_proto.h"
19 
20class SoftplusV2Grad : public testing::Test {
21protected:
22 static void SetUpTestCase() { std::cout << "SoftplusV2Grad InferShape SetUp" << std::endl; }
23 
24 static void TearDownTestCase() { std::cout << "SoftplusV2Grad InferShape TearDown" << std::endl; }
25};
26 
27// 同 shape 广播: 2 输入相同 → 输出相同
28TEST_F(SoftplusV2Grad, SoftplusV2Grad_infershape_same_shape)
29{
30 ge::op::SoftplusV2Grad op;
31 op.UpdateInputDesc("input_gradients", create_desc({4, 1, 1280}, ge::DT_FLOAT16));
32 op.UpdateInputDesc("input_features", create_desc({4, 1, 1280}, ge::DT_FLOAT16));
33 
34 EXPECT_EQ(InferShapeTest(op), ge::GRAPH_SUCCESS);
35}
36 
37// 可广播 shape: {4,1,1280} + {1,3,1280} → {4,3,1280}
38TEST_F(SoftplusV2Grad, SoftplusV2Grad_infershape_broadcast)
39{
40 ge::op::SoftplusV2Grad op;
41 op.UpdateInputDesc("input_gradients", create_desc({4, 1, 1280}, ge::DT_FLOAT));
42 op.UpdateInputDesc("input_features", create_desc({1, 3, 1280}, ge::DT_FLOAT));
43 
44 EXPECT_EQ(InferShapeTest(op), ge::GRAPH_SUCCESS);
45}
46 
47// rank 不同 + 标量广播: {2,3,4} + {4}
48TEST_F(SoftplusV2Grad, SoftplusV2Grad_infershape_rank_diff)
49{
50 ge::op::SoftplusV2Grad op;
51 op.UpdateInputDesc("input_gradients", create_desc({2, 3, 4}, ge::DT_BF16));
52 op.UpdateInputDesc("input_features", create_desc({4}, ge::DT_BF16));
53 
54 EXPECT_EQ(InferShapeTest(op), ge::GRAPH_SUCCESS);
55}
Ractivation/softplus_v2_grad/op_host/CMakeLists.txt→activation/softplus_v2_grad/tests/ut/op_kernel/CMakeLists.txt+8-5
@@ -1,9 +1,12 @@
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#/8# Generated By CANNBot
9add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE softplus_v2_grad ACLNNTYPE aclnn)9 
10if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
11 AddOpTestCase(softplus_v2_grad "ascend950pr_9599" "-DDTYPE_INPUT_GRADIENTS=float")
12endif()
Aactivation/softplus_v2_grad/tests/ut/op_kernel/softplus_v2_grad_data/compare_data.py+76-0
@@ -0,0 +1,76 @@
1#!/usr/bin/env python3
2# -*- coding: utf-8 -*-
3# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5# CANN Open Software License Agreement Version 2.0 (the "License").
6# Please refer to the License for details. You may not use this file except in compliance with the License.
7# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9# See LICENSE in the root of the software repository for the full text of the License.
10 
11# Generated By CANNBot
12#
13# Compares <tag>_npuout.bin against <tag>_golden.bin.
14# Usage: python3 compare_data.py <tag> <dtype>
15# Exit code 0 on PASS, 1 on FAIL (so the C++ harness ExecuteCommand detects failures).
16 
17import sys
18import os
19import numpy as np
20 
21curr_dir = os.path.dirname(os.path.realpath(__file__))
22 
23 
24def bf16_u16_to_fp32(arr_u16):
25 u32 = arr_u16.astype(np.uint32) << 16
26 return u32.view(np.float32)
27 
28 
29def compare(tag, dtype):
30 golden_file = os.path.join(curr_dir, f"{tag}_golden.bin")
31 out_file = os.path.join(curr_dir, f"{tag}_npuout.bin")
32 
33 if not os.path.exists(golden_file) or not os.path.exists(out_file):
34 print(f"COMPARE DATA FAILED! missing file: {golden_file} or {out_file}")
35 return False
36 
37 if dtype == "float32":
38 rtol, atol = 1e-4, 1e-5
39 g = np.fromfile(golden_file, np.float32)
40 o = np.fromfile(out_file, np.float32)
41 elif dtype == "float16":
42 rtol, atol = 1e-2, 1e-3
43 g = np.fromfile(golden_file, np.float16).astype(np.float32)
44 o = np.fromfile(out_file, np.float16).astype(np.float32)
45 elif dtype == "bfloat16":
46 rtol, atol = 1e-2, 1e-2
47 g = bf16_u16_to_fp32(np.fromfile(golden_file, np.uint16))
48 o = bf16_u16_to_fp32(np.fromfile(out_file, np.uint16))
49 else:
50 print(f"COMPARE DATA FAILED! unsupported dtype {dtype}")
51 return False
52 
53 if g.shape != o.shape:
54 print(f"COMPARE DATA FAILED! shape mismatch golden={g.shape} out={o.shape}")
55 return False
56 
57 close = np.isclose(o, g, rtol=rtol, atol=atol, equal_nan=True)
58 bad = np.where(~close)[0]
59 if len(bad) == 0:
60 print("COMPARE DATA PASSED!")
61 return True
62 
63 print(
64 f"COMPARE DATA FAILED! {len(bad)}/{g.size} mismatches (rtol={rtol}, atol={atol})"
65 )
66 for idx in bad[:8]:
67 print(f" index {idx}: out={o[idx]} golden={g[idx]}")
68 return False
69 
70 
71if __name__ == "__main__":
72 if len(sys.argv) != 3:
73 print("Usage: python3 compare_data.py <tag> <dtype>")
74 sys.exit(1)
75 ok = compare(sys.argv[1], sys.argv[2])
76 sys.exit(0 if ok else 1)
Aactivation/softplus_v2_grad/tests/ut/op_kernel/softplus_v2_grad_data/gen_data.py+102-0
@@ -0,0 +1,102 @@
1#!/usr/bin/env python3
2# -*- coding: utf-8 -*-
3# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5# CANN Open Software License Agreement Version 2.0 (the "License").
6# Please refer to the License for details. You may not use this file except in compliance with the License.
7# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9# See LICENSE in the root of the software repository for the full text of the License.
10 
11# Generated By CANNBot
12#
13# SoftplusV2Grad golden data generator.
14# gradInput = gradOut * ( 1/(1+exp(-beta*self)) if beta*self <= threshold
15# else 1 )
16# Broadcasting of gradOut / features to out shape is supported (numpy semantics).
17#
18# Usage:
19# python3 gen_data.py <tag> <grad_shape> <feat_shape> <out_shape> <dtype> <beta> <threshold>
20# Example:
21# python3 gen_data.py case1 '(4,64)' '(4,64)' '(4,64)' float32 1.0 20.0
22#
23# Files written (in cwd):
24# <tag>_grad.bin, <tag>_features.bin, <tag>_golden.bin
25 
26import sys
27import numpy as np
28 
29 
30def parse_shape(shape_str):
31 s = shape_str.strip().strip("'").strip('"').strip("(").strip(")")
32 parts = [p for p in s.replace(" ", "").split(",") if p != ""]
33 return tuple(int(x) for x in parts)
34 
35 
36def fp32_to_bf16_u16(arr_fp32):
37 # Truncate fp32 to bf16 by taking the high 16 bits (little-endian: bytes 2..4).
38 arr_fp32 = np.ascontiguousarray(arr_fp32.astype(np.float32))
39 u32 = arr_fp32.view(np.uint32)
40 return (u32 >> 16).astype(np.uint16)
41 
42 
43def bf16_u16_to_fp32(arr_u16):
44 u32 = arr_u16.astype(np.uint32) << 16
45 return u32.view(np.float32)
46 
47 
48def softplus_v2_grad_golden(grad_f32, feat_f32, beta, threshold):
49 bx = beta * feat_f32
50 sig = 1.0 / (1.0 + np.exp(-bx))
51 # where beta*self <= threshold use sigmoid, else pass gradient through (factor 1).
52 factor = np.where(bx <= threshold, sig, np.ones_like(sig))
53 return grad_f32 * factor
54 
55 
56def gen(tag, grad_shape, feat_shape, out_shape, dtype, beta, threshold):
57 np.random.seed(1234 + (abs(hash(tag)) % 100000))
58 
59 # features spread wide enough (with beta) to straddle the threshold on both sides.
60 feat_fp32 = np.random.uniform(-6.0, 6.0, feat_shape).astype(np.float32)
61 grad_fp32 = np.random.uniform(-2.0, 2.0, grad_shape).astype(np.float32)
62 
63 if dtype == "bfloat16":
64 # Round inputs to bf16 first, then compute golden on the rounded values.
65 feat_u16 = fp32_to_bf16_u16(feat_fp32)
66 grad_u16 = fp32_to_bf16_u16(grad_fp32)
67 feat_r = bf16_u16_to_fp32(feat_u16)
68 grad_r = bf16_u16_to_fp32(grad_u16)
69 golden = softplus_v2_grad_golden(grad_r, feat_r, beta, threshold)
70 golden_b = np.broadcast_to(golden, out_shape)
71 grad_u16.reshape(grad_shape).tofile(f"./{tag}_grad.bin")
72 feat_u16.reshape(feat_shape).tofile(f"./{tag}_features.bin")
73 fp32_to_bf16_u16(golden_b).reshape(out_shape).tofile(f"./{tag}_golden.bin")
74 return
75 
76 np_dtype = {"float32": np.float32, "float16": np.float16}[dtype]
77 feat_c = feat_fp32.astype(np_dtype)
78 grad_c = grad_fp32.astype(np_dtype)
79 golden = softplus_v2_grad_golden(
80 grad_c.astype(np.float32), feat_c.astype(np.float32), beta, threshold
81 ).astype(np_dtype)
82 golden_b = np.broadcast_to(golden, out_shape)
83 grad_c.tofile(f"./{tag}_grad.bin")
84 feat_c.tofile(f"./{tag}_features.bin")
85 np.ascontiguousarray(golden_b).tofile(f"./{tag}_golden.bin")
86 
87 
88if __name__ == "__main__":
89 if len(sys.argv) != 8:
90 print(
91 "Usage: python3 gen_data.py <tag> <grad_shape> <feat_shape> "
92 "<out_shape> <dtype> <beta> <threshold>"
93 )
94 sys.exit(1)
95 tag = sys.argv[1]
96 gs = parse_shape(sys.argv[2])
97 fs = parse_shape(sys.argv[3])
98 os_ = parse_shape(sys.argv[4])
99 dt = sys.argv[5]
100 be = float(sys.argv[6])
101 th = float(sys.argv[7])
102 gen(tag, gs, fs, os_, dt, be, th)
Aactivation/softplus_v2_grad/tests/ut/op_kernel/softplus_v2_grad_tiling_def.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 * Generated By CANNBot
11 */
12 
13// UT tiling-definition shim for softplus_v2_grad.
14//
15// The op_kernel entry uses REGISTER_NONE_TILING, which the build-system tiling-header
16// generator does not recognize (it only handles REGISTER_TILING_DEFAULT / BroadcastSch),
17// so GET_TILING_DATA_WITH_STRUCT is not auto-generated for this op. When this
18// <op>_tiling_def.h exists, ut.cmake force-includes it (and skips generation), so we
19// define the macro here. The kernel simply needs a local TilingData struct populated
20// from the raw tiling GM buffer; a byte copy is exactly what the generated macro does
21// (minus the internal REGISTER_TILINGDATA_SIZE bookkeeping, which is unnecessary for UT).
22 
23#ifndef SOFTPLUS_V2_GRAD_TILING_DEF_H_
24#define SOFTPLUS_V2_GRAD_TILING_DEF_H_
25 
26#include <cstring>
27 
28#ifndef GET_TILING_DATA_WITH_STRUCT
29#define GET_TILING_DATA_WITH_STRUCT(tiling_struct, tiling_data, tiling_arg) \
30 tiling_struct tiling_data; \
31 std::memcpy(reinterpret_cast<void*>(&tiling_data), reinterpret_cast<const void*>(tiling_arg), sizeof(tiling_struct))
32#endif
33 
34#endif // SOFTPLUS_V2_GRAD_TILING_DEF_H_
Aactivation/softplus_v2_grad/tests/ut/op_kernel/test_softplus_v2_grad.cpp+438-0
@@ -0,0 +1,438 @@
1/**
2 * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.
9 *
10 * Generated By CANNBot
11 */
12 
13#include <array>
14#include <vector>
15#include <iostream>
16#include <string>
17#include <cstdint>
18#include <algorithm>
19#include "gtest/gtest.h"
20#include "tikicpulib.h"
21#include "data_utils.h"
22#include "kernel_ut_data_helper.h"
23#include "kernel_ut_data_executor.h"
24 
25// Bring in the real kernel entry (defines softplus_v2_grad<RANK> for DTYPE_INPUT_GRADIENTS=float)
26// as well as the kernel class / tiling struct headers it includes.
27#include "../../../op_kernel/arch35/softplus_v2_grad.cpp"
28 
29using namespace std;
30 
31// ============================================================
32// Host-side re-implementation of the tiling math (mirrors
33// op_host/arch35/softplus_v2_grad_tiling_arch35.cpp) so that the
34// TilingData handed to the kernel is byte-for-byte what the real
35// host tiling would produce. This is metadata construction only;
36// the kernel under test is the real implementation.
37// ============================================================
38namespace svg_ut {
39 
40constexpr int64_t kUB = 253952; // UB size (matches host compile info UB_SIZE)
41 
42// FindSplitAxis — verbatim algorithm from host tiling.
43static bool FindSplitAxis(const std::vector<int64_t>& max_bro_shape, int64_t dtype_size, int64_t ub_per_core,
44 int64_t phys_nodes, SplitResult& out)
45{
46 if (phys_nodes <= 0 || dtype_size <= 0)
47 return false;
48 int64_t per_buf_bytes = (ub_per_core / phys_nodes) & ~(kUbAlignBytes - 1);
49 int64_t per_buf_elems = per_buf_bytes / dtype_size;
50 if (per_buf_elems <= 0)
51 return false;
52 int64_t rank = (int64_t)max_bro_shape.size();
53 int64_t inner = 1;
54 for (int64_t k = rank - 1; k >= 0; k--) {
55 if (max_bro_shape[k] * inner > per_buf_elems) {
56 out.a_i = per_buf_elems / inner;
57 if (out.a_i <= 0)
58 return false;
59 out.a_o = (max_bro_shape[k] + out.a_i - 1) / out.a_i;
60 int64_t rem = max_bro_shape[k] % out.a_i;
61 out.a_i_tail = (rem == 0) ? out.a_i : rem;
62 out.axis = k;
63 return true;
64 }
65 if (k == 0) {
66 out.axis = 0;
67 out.a_i = max_bro_shape[0];
68 out.a_o = 1;
69 out.a_i_tail = max_bro_shape[0];
70 return true;
71 }
72 inner *= max_bro_shape[k];
73 }
74 return true;
75}
76 
77// MultiCoreSplit — verbatim algorithm from host tiling.
78static bool MultiCoreSplit(const std::vector<int64_t>& max_bro_shape, const SplitResult& ub_split, int64_t max_cores,
79 MultiCoreResult& out)
80{
81 int64_t k = ub_split.axis, outer_prod = 1;
82 for (int64_t j = 0; j < k; j++)
83 outer_prod *= max_bro_shape[j];
84 out.total_tiles = outer_prod * ub_split.a_o;
85 out.num_cores = (out.total_tiles < max_cores) ? out.total_tiles : max_cores;
86 if (out.num_cores <= 0)
87 return false;
88 out.tiles_main = out.total_tiles / out.num_cores;
89 out.cores_tail = out.total_tiles % out.num_cores;
90 return true;
91}
92 
93// PrecomputeStrides — verbatim from host tiling (size-1 dims -> stride 0).
94static void PrecomputeStrides(const std::vector<int64_t>& s, std::vector<int64_t>& strides)
95{
96 int64_t rank = (int64_t)s.size();
97 strides.assign(rank, 0);
98 for (int64_t d = rank - 1; d >= 0; d--) {
99 if (s[d] == 1) {
100 strides[d] = 0;
101 continue;
102 }
103 int64_t prod = 1;
104 for (int64_t j = d + 1; j < rank; j++)
105 prod *= s[j];
106 strides[d] = prod;
107 }
108}
109 
110// Front-pad src[0..n) into dst[delta..delta+n), dst[0..delta)=padValue.
111static void PadRow(int64_t* dst, const int64_t* src, int64_t n, int64_t delta, int64_t padValue)
112{
113 for (int64_t d = 0; d < delta; d++)
114 dst[d] = padValue;
115 for (int64_t d = 0; d < n; d++)
116 dst[d + delta] = src[d];
117}
118 
119// Build a SoftplusV2GradTilingData<R> from already-squeezed max_bro_shape + per-slot normal
120// shapes (all rank_ length, no dims squeezed for the shapes chosen in these tests).
121// Returns false if tiling is infeasible.
122template <int64_t R>
123static bool BuildTiling(uint8_t* buf, const std::vector<int64_t>& max_bro_shape,
124 const std::vector<std::vector<int64_t>>& in_shapes, const std::vector<int64_t>& out_shape,
125 float beta, float threshold, int64_t core_num)
126{
127 auto* t = reinterpret_cast<SoftplusV2GradTilingData<R>*>(buf);
128 int64_t rank_ = (int64_t)max_bro_shape.size();
129 int64_t delta = R - rank_;
130 if (delta < 0)
131 return false;
132 
133 int64_t per_buf_bytes = (kUB / kPhysNodes) & ~(kUbAlignBytes - 1);
134 if (!FindSplitAxis(max_bro_shape, sizeof(float), kUB, kPhysNodes, t->split))
135 return false;
136 if (!MultiCoreSplit(max_bro_shape, t->split, core_num, t->multicore))
137 return false;
138 
139 t->per_buf_bytes = per_buf_bytes;
140 t->beta = beta;
141 t->threshold = threshold;
142 t->rank = rank_;
143 t->num_inputs = (int64_t)in_shapes.size();
144 t->num_outputs = 1;
145 
146 // max_bro_shape: front-pad 1.
147 PadRow(t->max_bro_shape, max_bro_shape.data(), rank_, delta, 1);
148 // split axis shifted right by delta.
149 t->split.axis += delta;
150 
151 // input slots
152 for (int64_t i = 0; i < kMaxInputSlots; i++) {
153 if (i < (int64_t)in_shapes.size()) {
154 std::vector<int64_t> st;
155 PrecomputeStrides(in_shapes[i], st);
156 PadRow(t->input_shapes[i], in_shapes[i].data(), rank_, delta, 1);
157 PadRow(t->input_strides[i], st.data(), rank_, delta, 0);
158 } else {
159 for (int64_t d = 0; d < R; d++) {
160 t->input_shapes[i][d] = 1;
161 t->input_strides[i][d] = 0;
162 }
163 }
164 }
165 // output slot 0
166 {
167 std::vector<int64_t> st;
168 PrecomputeStrides(out_shape, st);
169 PadRow(t->output_shapes[0], out_shape.data(), rank_, delta, 1);
170 PadRow(t->output_strides[0], st.data(), rank_, delta, 0);
171 }
172 return true;
173}
174 
175} // namespace svg_ut
176 
177// ============================================================
178// Direct-invocation wrappers around the real kernel class.
179// These call the exact implementation in
180// op_kernel/arch35/softplus_v2_grad_kernel.h with the requested
181// dtype/RANK, letting one binary cover fp32 + fp16 + bf16.
182// ============================================================
183template <typename T, int64_t RANK>
184__global__ __aicore__ void RunSoftplusKernel(GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling)
185{
186 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
187 GM_ADDR ins[kMaxInputSlots] = {grad, feat};
188 GM_ADDR outs[kMaxOutputSlots] = {out};
189 const auto* td = reinterpret_cast<const SoftplusV2GradTilingData<RANK>*>(tiling);
190 SoftplusV2GradKernel<T, RANK> kernel;
191 kernel.Init(ins, outs, td);
192 kernel.Process();
193}
194 
195namespace {
196 
197const char* kDataDirRel = "activation/softplus_v2_grad/tests/ut/op_kernel/softplus_v2_grad_data";
198const char* kDataLocal = "softplus_v2_grad_data";
199 
200size_t Prod(const std::vector<int64_t>& s)
201{
202 size_t p = 1;
203 for (auto v : s)
204 p *= (size_t)v;
205 return p;
206}
207 
208// dtype tags used by gen_data.py / compare_data.py and their element byte sizes.
209struct DTypeInfo {
210 const char* pyName;
211 size_t bytes;
212};
213 
214class softplus_v2_grad_test : public testing::Test {
215protected:
216 static void SetUpTestCase() { cout << "softplus_v2_grad_test SetUp\n" << endl; }
217 static void TearDownTestCase()
218 {
219 cout << "softplus_v2_grad_test TearDown\n" << endl;
220 kernel_ut::CleanGeneratedBinFiles("./softplus_v2_grad_data");
221 }
222};
223 
224// Shared driver: generate data, build tiling, run kernel via ICPU_RUN_KF, compare golden.
225// KernelInvoker is a callable (lambda) matching the ICPU_RUN_KF signature.
226template <typename KernelInvoker>
227void RunAndVerify(const std::string& tag, const std::vector<int64_t>& grad_shape,
228 const std::vector<int64_t>& feat_shape, const std::vector<int64_t>& out_shape, const DTypeInfo& dt,
229 float beta, float threshold, uint32_t blockDim, uint8_t* tiling, KernelInvoker invoker)
230{
231 kernel_ut::SetupTestEnvironment(kDataDirRel, kDataLocal);
232 
233 auto shapeStr = [](const std::vector<int64_t>& s) {
234 std::string r = "'(";
235 for (size_t i = 0; i < s.size(); i++) {
236 r += std::to_string(s[i]);
237 if (i + 1 < s.size())
238 r += ",";
239 }
240 r += ")'";
241 return r;
242 };
243 
244 kernel_ut::RunGenData("./softplus_v2_grad_data",
245 {tag, shapeStr(grad_shape), shapeStr(feat_shape), shapeStr(out_shape), dt.pyName,
246 std::to_string(beta), std::to_string(threshold)});
247 
248 std::string path = kernel_ut::GetTestWorkDir();
249 size_t gradBytes = Prod(grad_shape) * dt.bytes;
250 size_t featBytes = Prod(feat_shape) * dt.bytes;
251 size_t outBytes = Prod(out_shape) * dt.bytes;
252 
253 uint8_t* grad = (uint8_t*)AscendC::GmAlloc(gradBytes);
254 uint8_t* feat = (uint8_t*)AscendC::GmAlloc(featBytes);
255 uint8_t* out = (uint8_t*)AscendC::GmAlloc(outBytes);
256 uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(16 * 1024 * 1024);
257 
258 ReadFile(path + "/softplus_v2_grad_data/" + tag + "_grad.bin", gradBytes, grad, gradBytes);
259 ReadFile(path + "/softplus_v2_grad_data/" + tag + "_features.bin", featBytes, feat, featBytes);
260 
261 AscendC::SetKernelMode(KernelMode::AIV_MODE);
262 ICPU_SET_TILING_KEY(0);
263 ICPU_RUN_KF(invoker, blockDim, grad, feat, out, workspace, tiling);
264 
265 WriteFile(path + "/softplus_v2_grad_data/" + tag + "_npuout.bin", out, outBytes);
266 
267 AscendC::GmFree(grad);
268 AscendC::GmFree(feat);
269 AscendC::GmFree(out);
270 AscendC::GmFree(workspace);
271 
272 // compare_data.py exits non-zero on mismatch; RunCompareData returns false in that case.
273 bool ok = kernel_ut::RunCompareData("./softplus_v2_grad_data", {tag, dt.pyName});
274 EXPECT_TRUE(ok) << "golden compare failed for tag=" << tag;
275}
276 
277const DTypeInfo kFp32{"float32", 4};
278const DTypeInfo kFp16{"float16", 2};
279const DTypeInfo kBf16{"bfloat16", 2};
280 
281// ---- Case 1: fp32, {4,64}, default attrs -> ProcessFP32 + VF formula branch, RANK_4 (NDDMA branch)
282TEST_F(softplus_v2_grad_test, case1_fp32_rank2_formula)
283{
284 std::vector<int64_t> g{4, 64}, f{4, 64}, o{4, 64};
285 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
286 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 20.0f, 1)));
287 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
288 RunSoftplusKernel<float, SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
289 };
290 RunAndVerify("case1", g, f, o, kFp32, 1.0f, 20.0f, 1, tiling, inv);
291 AscendC::GmFree(tiling);
292}
293 
294// ---- Case 2: fp32, small threshold -> exercises BOTH Select branches of the VF
295TEST_F(softplus_v2_grad_test, case2_fp32_threshold_both_branches)
296{
297 std::vector<int64_t> g{8, 64}, f{8, 64}, o{8, 64};
298 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
299 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 1.0f, 1)));
300 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
301 RunSoftplusKernel<float, SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
302 };
303 RunAndVerify("case2", g, f, o, kFp32, 1.0f, 1.0f, 1, tiling, inv);
304 AscendC::GmFree(tiling);
305}
306 
307// ---- Case 3: fp16 -> ProcessFP16BF16 + Cast NONE/RINT
308TEST_F(softplus_v2_grad_test, case3_fp16_rank2)
309{
310 std::vector<int64_t> g{8, 128}, f{8, 128}, o{8, 128};
311 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
312 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 20.0f, 1)));
313 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
314 RunSoftplusKernel<half, SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
315 };
316 RunAndVerify("case3", g, f, o, kFp16, 1.0f, 20.0f, 1, tiling, inv);
317 AscendC::GmFree(tiling);
318}
319 
320// ---- Case 4: bf16 -> ProcessFP16BF16 bf16
321TEST_F(softplus_v2_grad_test, case4_bf16_rank2)
322{
323 std::vector<int64_t> g{8, 128}, f{8, 128}, o{8, 128};
324 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
325 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 20.0f, 1)));
326 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
327 RunSoftplusKernel<bfloat16_t, SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
328 };
329 RunAndVerify("case4", g, f, o, kBf16, 1.0f, 20.0f, 1, tiling, inv);
330 AscendC::GmFree(tiling);
331}
332 
333// ---- Case 5: broadcast g{4,1,8} f{1,3,8} out{4,3,8} -> CopyInBrc broadcast strides, RANK_4
334TEST_F(softplus_v2_grad_test, case5_fp32_broadcast)
335{
336 std::vector<int64_t> g{4, 1, 8}, f{1, 3, 8}, o{4, 3, 8};
337 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
338 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 20.0f, 1)));
339 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
340 RunSoftplusKernel<float, SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
341 };
342 RunAndVerify("case5", g, f, o, kFp32, 1.0f, 20.0f, 1, tiling, inv);
343 AscendC::GmFree(tiling);
344}
345 
346// ---- Case 6: rank5 -> RANK_8 template, CopyInBrc flat-loop branch (R=8 > kMaxNdDmaDims)
347TEST_F(softplus_v2_grad_test, case6_fp32_rank5)
348{
349 std::vector<int64_t> s{2, 3, 2, 3, 4};
350 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_8>));
351 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_8>(tiling, s, {s, s}, s, 1.0f, 20.0f, 1)));
352 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
353 RunSoftplusKernel<float, SOFTPLUS_V2_GRAD_RANK_8>(grad, feat, out, ws, t);
354 };
355 RunAndVerify("case6", s, s, s, kFp32, 1.0f, 20.0f, 1, tiling, inv);
356 AscendC::GmFree(tiling);
357}
358 
359// ---- Case 7: rank6, explicit beta/threshold -> RANK_8 flat-loop, mixed Select branches
360TEST_F(softplus_v2_grad_test, case7_fp32_rank6_attrs)
361{
362 std::vector<int64_t> s{2, 2, 2, 2, 2, 4};
363 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_8>));
364 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_8>(tiling, s, {s, s}, s, 2.0f, 5.0f, 1)));
365 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
366 RunSoftplusKernel<float, SOFTPLUS_V2_GRAD_RANK_8>(grad, feat, out, ws, t);
367 };
368 RunAndVerify("case7", s, s, s, kFp32, 2.0f, 5.0f, 1, tiling, inv);
369 AscendC::GmFree(tiling);
370}
371 
372// ---- Case 8a/8b: apt.cpp entry (real softplus_v2_grad<RANK>) covers RANK dispatch + GET_TILING_DATA
373TEST_F(softplus_v2_grad_test, case8_apt_entry_rank4)
374{
375 std::vector<int64_t> g{4, 64}, f{4, 64}, o{4, 64};
376 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
377 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 20.0f, 1)));
378 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
379 ::softplus_v2_grad<SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
380 };
381 RunAndVerify("case8a", g, f, o, kFp32, 1.0f, 20.0f, 1, tiling, inv);
382 AscendC::GmFree(tiling);
383}
384 
385TEST_F(softplus_v2_grad_test, case8_apt_entry_rank8)
386{
387 std::vector<int64_t> s{2, 3, 2, 3, 4};
388 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_8>));
389 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_8>(tiling, s, {s, s}, s, 1.0f, 20.0f, 1)));
390 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
391 ::softplus_v2_grad<SOFTPLUS_V2_GRAD_RANK_8>(grad, feat, out, ws, t);
392 };
393 RunAndVerify("case8b", s, s, s, kFp32, 1.0f, 20.0f, 1, tiling, inv);
394 AscendC::GmFree(tiling);
395}
396 
397// ---- Case 9: multi-core (coreNum=4, not divisible) -> GetCoreRange tail-branch coverage
398TEST_F(softplus_v2_grad_test, case9_fp32_multicore)
399{
400 std::vector<int64_t> g{10, 64}, f{10, 64}, o{10, 64};
401 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
402 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 20.0f, 4)));
403 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
404 RunSoftplusKernel<float, SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
405 };
406 RunAndVerify("case9", g, f, o, kFp32, 1.0f, 20.0f, 4, tiling, inv);
407 AscendC::GmFree(tiling);
408}
409 
410// ---- Case 10: fp32 UB-split (a_o>1) + non-divisible cores -> ProcessFP32 multi-tile-per-core
411// (covers GetCoreRange cores_tail branch and the "flat != end-1" MTE3_MTE2 SetFlag)
412TEST_F(softplus_v2_grad_test, case10_fp32_ubsplit_multitile)
413{
414 std::vector<int64_t> g{2, 20000}, f{2, 20000}, o{2, 20000};
415 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
416 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 20.0f, 3)));
417 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
418 RunSoftplusKernel<float, SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
419 };
420 RunAndVerify("case10", g, f, o, kFp32, 1.0f, 20.0f, 3, tiling, inv);
421 AscendC::GmFree(tiling);
422}
423 
424// ---- Case 11: fp16 UB-split (a_o>1) + non-divisible cores -> ProcessFP16BF16 multi-tile-per-core
425// (covers the fp16/bf16 path "flat != end-1" MTE3_MTE2 SetFlag)
426TEST_F(softplus_v2_grad_test, case11_fp16_ubsplit_multitile)
427{
428 std::vector<int64_t> g{2, 20000}, f{2, 20000}, o{2, 20000};
429 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(SoftplusV2GradTilingData<SOFTPLUS_V2_GRAD_RANK_4>));
430 ASSERT_TRUE((svg_ut::BuildTiling<SOFTPLUS_V2_GRAD_RANK_4>(tiling, o, {g, f}, o, 1.0f, 20.0f, 3)));
431 auto inv = [](GM_ADDR grad, GM_ADDR feat, GM_ADDR out, GM_ADDR ws, GM_ADDR t) {
432 RunSoftplusKernel<half, SOFTPLUS_V2_GRAD_RANK_4>(grad, feat, out, ws, t);
433 };
434 RunAndVerify("case11", g, f, o, kFp16, 1.0f, 20.0f, 3, tiling, inv);
435 AscendC::GmFree(tiling);
436}
437 
438} // namespace