Star912
1827
代码
代码
Issues224
Pull Requests308
Actions
项目讨论
Wiki
分析
Star912
1827
已合并
[CANNBot] 增加 experimental 下 Softshrink 实现 #3298
陈展熹创建于 3月28日
[CANNBot] 增加 experimental 下 Softshrink 实现 #3298
已合并
陈展熹创建于 3月28日
从yourealize/ops-nnfeature/softshrink合入到cann/ops-nnmaster
共 12 个文件变更+801-0
Aexperimental/activation/softshrink/CMakeLists.txt+9-0
@@ -0,0 +1,9 @@
1+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
2+if(NOT ENABLE_TEST)
3+ list(REMOVE_ITEM CURRENT_DIRS tests)
4+endif()
5+foreach(SUB_DIR ${CURRENT_DIRS})
6+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
7+ add_subdirectory(${SUB_DIR})
8+ endif()
9+endforeach()
Aexperimental/activation/softshrink/README.md+74-0
@@ -0,0 +1,74 @@
1+# Softshrink
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| ---- | :----:|
7+|Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件|√|
8+ 
9+## 功能说明
10+ 
11+- 算子功能：对输入Tensor逐元素执行Softshrink激活函数。
12+ 
13+- 计算公式：
14+ 
15+$$
16+y_i = \begin{cases} x_i - \lambda, & \text{if } x_i > \lambda \\ x_i + \lambda, & \text{if } x_i < -\lambda \\ 0, & \text{otherwise} \end{cases}
17+$$
18+ 
19+## 参数说明
20+ 
21+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
22+ <col style="width: 100px">
23+ <col style="width: 150px">
24+ <col style="width: 280px">
25+ <col style="width: 330px">
26+ <col style="width: 120px">
27+ </colgroup>
28+ <thead>
29+ <tr>
30+ <th>参数名</th>
31+ <th>输入/输出/属性</th>
32+ <th>描述</th>
33+ <th>数据类型</th>
34+ <th>数据格式</th>
35+ </tr></thead>
36+ <tbody>
37+ <tr>
38+ <td>input_x</td>
39+ <td>输入</td>
40+ <td>待进行Softshrink计算的输入Tensor。</td>
41+ <td>fp16、fp32、bf16</td>
42+ <td>ND</td>
43+ </tr>
44+ <tr>
45+ <td>lambd</td>
46+ <td>属性</td>
47+ <td>Softshrink公式中的lambda值，默认0.5，需 >= 0。</td>
48+ <td>float</td>
49+ <td>标量</td>
50+ </tr>
51+ <tr>
52+ <td>output_y</td>
53+ <td>输出</td>
54+ <td>Softshrink计算结果，shape与input_x相同。</td>
55+ <td>fp16、fp32、bf16</td>
56+ <td>ND</td>
57+ </tr>
58+ </tbody></table>
59+ 
60+## 约束说明
61+ 
62+- lambd属性值必须 >= 0。
63+ 
64+## 调用说明
65+ 
66+| 调用方式 | 调用样例 | 说明 |
67+|----------|---------|------|
68+| aclnn调用 | [test_aclnn_softshrink.cpp](./examples/test_aclnn_softshrink.cpp) | 通过aclnn接口方式调用Softshrink算子。 |
69+ 
70+## 贡献说明
71+ 
72+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
73+| ---- | ---- | ---- | ---- | ---- |
74+| yourealize | 个人开发者 | Softshrink | 2025/03 | Softshrink算子适配开源仓 |
Aexperimental/activation/softshrink/examples/test_aclnn_softshrink.cpp+133-0
@@ -0,0 +1,133 @@
1+/**
2+ * Copyright (c) 2025-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+ * NOTE: Portions of this code were AI-generated and have been
11+ * technically reviewed for functional accuracy and security
12+ */
13+ 
14+#include <iostream>
15+#include <vector>
16+#include "acl/acl.h"
17+#include "aclnn_softshrink.h"
18+ 
19+#define CHECK_RET(cond, return_expr) \
20+ do { \
21+ if (!(cond)) { \
22+ return_expr; \
23+ } \
24+ } while (0)
25+ 
26+#define LOG_PRINT(message, ...) \
27+ do { \
28+ printf(message, ##__VA_ARGS__); \
29+ } while (0)
30+ 
31+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
32+ int64_t shapeSize = 1;
33+ for (auto i : shape) {
34+ shapeSize *= i;
35+ }
36+ return shapeSize;
37+}
38+ 
39+int Init(int32_t deviceId, aclrtStream* stream) {
40+ auto ret = aclInit(nullptr);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtSetDevice(deviceId);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
44+ ret = aclrtCreateStream(stream);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
46+ return 0;
47+}
48+ 
49+template <typename T>
50+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
51+ aclDataType dataType, aclTensor** tensor) {
52+ auto size = GetShapeSize(shape) * sizeof(T);
53+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
54+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
55+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
57+ 
58+ std::vector<int64_t> strides(shape.size(), 1);
59+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
60+ strides[i] = shape[i + 1] * strides[i + 1];
61+ }
62+ 
63+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
64+ shape.data(), shape.size(), *deviceAddr);
65+ return 0;
66+}
67+ 
68+int main() {
69+ // 1. device/stream初始化
70+ int32_t deviceId = 0;
71+ aclrtStream stream;
72+ auto ret = Init(deviceId, &stream);
73+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
74+ 
75+ // 2. 构造输入与输出
76+ std::vector<int64_t> inputShape = {2, 3};
77+ std::vector<int64_t> outputShape = {2, 3};
78+ void* inputDeviceAddr = nullptr;
79+ void* outputDeviceAddr = nullptr;
80+ aclTensor* inputX = nullptr;
81+ aclTensor* outputY = nullptr;
82+ std::vector<float> inputHostData = {-1.0, -0.3, 0.0, 0.3, 1.0, 2.0};
83+ std::vector<float> outputHostData(6, 0);
84+ float lambdValue = 0.5f;
85+ 
86+ ret = CreateAclTensor(inputHostData, inputShape, &inputDeviceAddr, aclDataType::ACL_FLOAT, &inputX);
87+ CHECK_RET(ret == ACL_SUCCESS, return ret);
88+ 
89+ ret = CreateAclTensor(outputHostData, outputShape, &outputDeviceAddr, aclDataType::ACL_FLOAT, &outputY);
90+ CHECK_RET(ret == ACL_SUCCESS, return ret);
91+ 
92+ // 3. 调用aclnnSoftshrink算子
93+ uint64_t workspaceSize = 0;
94+ aclOpExecutor* executor;
95+ ret = aclnnSoftshrinkGetWorkspaceSize(inputX, lambdValue, outputY, &workspaceSize, &executor);
96+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSoftshrinkGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
97+ 
98+ void* workspaceAddr = nullptr;
99+ if (workspaceSize > 0) {
100+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
101+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
102+ }
103+ 
104+ ret = aclnnSoftshrink(workspaceAddr, workspaceSize, executor, stream);
105+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSoftshrink failed. ERROR: %d\n", ret); return ret);
106+ 
107+ // 4. 同步等待
108+ ret = aclrtSynchronizeStream(stream);
109+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
110+ 
111+ // 5. 获取输出
112+ auto size = GetShapeSize(outputShape);
113+ std::vector<float> outData(size, 0);
114+ ret = aclrtMemcpy(outData.data(), outData.size() * sizeof(outData[0]), outputDeviceAddr,
115+ size * sizeof(outData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
116+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
117+ for (int64_t i = 0; i < size; i++) {
118+ LOG_PRINT("out result[%ld] is: %f\n", i, outData[i]);
119+ }
120+ 
121+ // 6. 释放资源
122+ aclDestroyTensor(inputX);
123+ aclDestroyTensor(outputY);
124+ aclrtFree(inputDeviceAddr);
125+ aclrtFree(outputDeviceAddr);
126+ if (workspaceSize > 0) {
127+ aclrtFree(workspaceAddr);
128+ }
129+ aclrtDestroyStream(stream);
130+ aclrtResetDevice(deviceId);
131+ aclFinalize();
132+ return 0;
133+}
Aexperimental/activation/softshrink/op_host/CMakeLists.txt+1-0
@@ -0,0 +1 @@
1+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE softshrink ACLNNTYPE aclnn)
Aexperimental/activation/softshrink/op_host/softshrink_def.cpp+49-0
@@ -0,0 +1,49 @@
1+/**
2+ * Copyright (c) 2025-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+ * NOTE: Portions of this code were AI-generated and have been
11+ * technically reviewed for functional accuracy and security
12+*/
13+ 
14+#include "register/op_def_registry.h"
15+ 
16+namespace ops {
17+class Softshrink : public OpDef {
18+public:
19+ explicit Softshrink(const char* name) : OpDef(name)
20+ {
21+ this->Input("input_x")
22+ .ParamType(REQUIRED)
23+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
24+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
25+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
26+ .AutoContiguous();
27+ this->Attr("lambd")
28+ .AttrType(OPTIONAL)
29+ .Float(0.5);
30+ this->Output("output_y")
31+ .ParamType(REQUIRED)
32+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
33+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
34+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
35+ .AutoContiguous();
36+ 
37+ OpAICoreConfig aicoreConfig;
38+ aicoreConfig.DynamicCompileStaticFlag(true)
39+ .DynamicFormatFlag(false)
40+ .DynamicRankSupportFlag(true)
41+ .DynamicShapeSupportFlag(true)
42+ .NeedCheckSupportFlag(false)
43+ .PrecisionReduceFlag(true)
44+ .ExtendCfgInfo("opFile.value", "softshrink");
45+ this->AICore().AddConfig("ascend910b", aicoreConfig);
46+ }
47+};
48+OP_ADD(Softshrink);
49+} // namespace ops
Aexperimental/activation/softshrink/op_host/softshrink_infershape.cpp+43-0
@@ -0,0 +1,43 @@
1+/**
2+ * Copyright (c) 2025-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+ * NOTE: Portions of this code were AI-generated and have been
11+ * technically reviewed for functional accuracy and security
12+ */
13+ 
14+#include "register/op_impl_registry.h"
15+#include "exe_graph/runtime/infer_shape_context.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+ 
20+namespace ops {
21+ 
22+static ge::graphStatus InferShape4Softshrink(gert::InferShapeContext* context)
23+{
24+ OP_LOGI(context, "Softshrink InferShape start");
25+
26+ const gert::Shape* input_shape = context->GetInputShape(0);
27+ OP_CHECK_NULL_WITH_CONTEXT(context, input_shape);
28+ OP_LOGI(context, "Get input shape success, shape size: %ld", input_shape->GetShapeSize());
29+
30+ gert::Shape* output_shape = context->GetOutputShape(0);
31+ OP_CHECK_NULL_WITH_CONTEXT(context, output_shape);
32+ OP_LOGI(context, "Get output shape success");
33+
34+ *output_shape = *input_shape;
35+ OP_LOGI(context, "Set output shape success, output shape size: %ld", output_shape->GetShapeSize());
36+
37+ OP_LOGI(context, "Softshrink InferShape success");
38+ return ge::GRAPH_SUCCESS;
39+}
40+ 
41+IMPL_OP_INFERSHAPE(Softshrink).InferShape(InferShape4Softshrink);
42+ 
43+} // namespace ops
Aexperimental/activation/softshrink/op_host/softshrink_tiling.cpp+160-0
@@ -0,0 +1,160 @@
1+/**
2+ * Copyright (c) 2025-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+ * NOTE: Portions of this code were AI-generated and have been
11+ * technically reviewed for functional accuracy and security
12+ */
13+ 
14+#include <algorithm>
15+#include "register/op_def_registry.h"
16+#include "op_common/log/log.h"
17+#include "op_common/op_host/util/math_util.h"
18+#include "op_common/op_host/util/platform_util.h"
19+#include "../op_kernel/softshrink_tiling_data.h"
20+#include "../op_kernel/softshrink_tiling_key.h"
21+ 
22+namespace optiling {
23+ 
24+using Ops::Base::CeilDiv;
25+using Ops::Base::FloorDiv;
26+using Ops::Base::FloorAlign;
27+using Ops::Base::GetUbBlockSize;
28+ 
29+constexpr uint32_t WS_SYS_SIZE = 0U;
30+constexpr int64_t MIN_SPLIT_THRESHOLD = 1024;
31+ 
32+static const gert::Shape g_vec_1_shape = {1};
33+ 
34+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape) {
35+ if (in_shape.GetDimNum() == 0) {
36+ return g_vec_1_shape;
37+ }
38+ return in_shape;
39+}
40+ 
41+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context,
42+ uint64_t& ubSize, int64_t& coreNum)
43+{
44+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
45+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
46+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
47+ coreNum = ascendcPlatform.GetCoreNumAiv();
48+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"),
49+ return ge::GRAPH_FAILED);
50+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
51+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"),
52+ return ge::GRAPH_FAILED);
53+ return ge::GRAPH_SUCCESS;
54+}
55+ 
56+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context,
57+ int64_t& totalIdx, ge::DataType& dataType, float& lambd)
58+{
59+ auto inputX = context->GetInputShape(0);
60+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
61+ auto shapeX = EnsureNotScalar(inputX->GetStorageShape());
62+ 
63+ totalIdx = shapeX.GetShapeSize();
64+ 
65+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
66+ auto inputDesc = context->GetInputDesc(0);
67+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
68+ dataType = inputDesc->GetDataType();
69+ if (supportedDtype.count(dataType) == 0) {
70+ OP_LOGE(context, "Softshrink: unsupported dtype");
71+ return ge::GRAPH_FAILED;
72+ }
73+ 
74+ const float* lambdPtr = context->GetAttrs()->GetAttrPointer<float>(0);
75+ lambd = (lambdPtr != nullptr) ? *lambdPtr : 0.5f;
76+ OP_CHECK_IF(lambd < 0.0f,
77+ OP_LOGE(context, "Softshrink: lambd must be >= 0, got %f", lambd),
78+ return ge::GRAPH_FAILED);
79+ 
80+ return ge::GRAPH_SUCCESS;
81+}
82+ 
83+static ge::graphStatus SoftshrinkTilingFunc(gert::TilingContext* context)
84+{
85+ uint64_t ubSize;
86+ int64_t coreNum;
87+ OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
88+ OP_LOGE(context, "GetPlatformInfo error"),return ge::GRAPH_FAILED);
89+ 
90+ int64_t totalIdx;
91+ ge::DataType dataType;
92+ float lambd;
93+ OP_CHECK_IF(GetShapeAttrsInfo(context, totalIdx, dataType, lambd) != ge::GRAPH_SUCCESS,
94+ OP_LOGE(context, "GetShapeAttrsInfo error"),return ge::GRAPH_FAILED);
95+ 
96+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
97+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
98+ currentWorkspace[0] = WS_SYS_SIZE;
99+ 
100+ SoftshrinkTilingData* tiling =context->GetTilingData<SoftshrinkTilingData>();
101+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
102+ OP_CHECK_IF(memset_s(tiling, sizeof(SoftshrinkTilingData), 0,
103+ sizeof(SoftshrinkTilingData)) != EOK,OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
104+ 
105+ if (totalIdx == 0) {
106+ tiling->totalNum = 0;
107+ tiling->blockFactor = 0;
108+ tiling->ubFactor = 0;
109+ tiling->lambd = lambd;
110+ context->SetBlockDim(1);
111+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
112+ ASCENDC_TPL_SEL_PARAM(context, dTypeX, 0ULL);
113+ return ge::GRAPH_SUCCESS;
114+ }
115+ 
116+ int64_t typeSize = (dataType == ge::DT_FLOAT) ? 4 : 2;
117+ tiling->totalNum = totalIdx;tiling->lambd = lambd;
118+ 
119+ constexpr int64_t MIN_PER_CORE_ELEMS = 1024;
120+ int64_t maxCores = std::max(1L, totalIdx / MIN_PER_CORE_ELEMS);
121+ int64_t effectiveCoreNum = std::min(coreNum, maxCores);
122+ 
123+ tiling->blockFactor = CeilDiv(totalIdx, effectiveCoreNum);
124+ int64_t usedCoreNum = CeilDiv(totalIdx, tiling->blockFactor);
125+ 
126+ uint64_t useDoubleBuffer = (tiling->blockFactor > MIN_SPLIT_THRESHOLD) ? 1 : 0;
127+ 
128+ int64_t bufferNum;
129+ if (dataType == ge::DT_FLOAT) {
130+ bufferNum = useDoubleBuffer ? 6 : 4;
131+ } else {
132+ bufferNum = useDoubleBuffer ? 9 : 7;
133+ }
134+ 
135+ constexpr int64_t VECTOR_ALIGN_ELEM = 256 / static_cast<int64_t>(sizeof(float));
136+ int64_t ubBlockSize = GetUbBlockSize(context);
137+ int64_t alignUnit = std::max(ubBlockSize, VECTOR_ALIGN_ELEM);
138+ tiling->ubFactor = FloorAlign(FloorDiv((static_cast<int64_t>(ubSize) / typeSize), bufferNum),alignUnit);
139+ 
140+ context->SetBlockDim(usedCoreNum);
141+ 
142+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
143+ ASCENDC_TPL_SEL_PARAM(context, dTypeX, useDoubleBuffer);
144+ 
145+ return ge::GRAPH_SUCCESS;
146+}
147+ 
148+static ge::graphStatus TilingParseForSoftshrink(
149+ [[maybe_unused]] gert::TilingParseContext* context)
150+{
151+ return ge::GRAPH_SUCCESS;
152+}
153+ 
154+struct SoftshrinkCompileInfo {};
155+ 
156+IMPL_OP_OPTILING(Softshrink)
157+ .Tiling(SoftshrinkTilingFunc)
158+ .TilingParse<SoftshrinkCompileInfo>(TilingParseForSoftshrink);
159+ 
160+} // namespace optiling
Aexperimental/activation/softshrink/op_kernel/softshrink.cpp+28-0
@@ -0,0 +1,28 @@
1+/**
2+ * Copyright (c) 2025-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+ * NOTE: Portions of this code were AI-generated and have been
11+ * technically reviewed for functional accuracy and security
12+ */
13+ 
14+#include "softshrink.h"
15+ 
16+using namespace AscendC;
17+ 
18+template <typename D_T_X, int BUFFER_MODE>
19+__global__ __aicore__ void softshrink(
20+ GM_ADDR input_x, GM_ADDR output_y,
21+ GM_ADDR workspace, GM_ADDR tiling)
22+{
23+ REGISTER_TILING_DEFAULT(SoftshrinkTilingData);
24+ GET_TILING_DATA_WITH_STRUCT(SoftshrinkTilingData, tilingData, tiling);
25+ NsSoftshrink::Softshrink<D_T_X, BUFFER_MODE> op;
26+ op.Init(input_x, output_y, &tilingData);
27+ op.Process();
28+}
Aexperimental/activation/softshrink/op_kernel/softshrink.h+241-0
@@ -0,0 +1,241 @@
1+/**
2+ * Copyright (c) 2025-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+ * NOTE: Portions of this code were AI-generated and have been
11+ * technically reviewed for functional accuracy and security
12+ */
13+ 
14+#ifndef SOFTSHRINK_H
15+#define SOFTSHRINK_H
16+ 
17+#include "kernel_operator.h"
18+#include "kernel_tiling/kernel_tiling.h"
19+#include "softshrink_tiling_data.h"
20+#include "softshrink_tiling_key.h"
21+ 
22+namespace NsSoftshrink {
23+ 
24+using namespace AscendC;
25+ 
26+// SoftShrink forward:
27+// y[i] = x[i] - lambd, if x[i] > lambd
28+// y[i] = x[i] + lambd, if x[i] < -lambd
29+// y[i] = 0, otherwise
30+ 
31+template <typename T, int BUFFER_MODE>
32+class Softshrink {
33+ static constexpr int32_t BUFFER_NUM = BUFFER_MODE ? 2 : 1;
34+ static constexpr bool NEED_CAST = std::is_same_v<T, half> || std::is_same_v<T, bfloat16_t>;
35+ using ComputeType = std::conditional_t<NEED_CAST, float, T>;
36+ 
37+public:
38+ __aicore__ inline Softshrink() {};
39+ __aicore__ inline void Init(GM_ADDR inputX, GM_ADDR outputY,
40+ const SoftshrinkTilingData* tilingData);
41+ __aicore__ inline void Process();
42+ 
43+private:
44+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
45+ __aicore__ inline void Compute(int64_t currentNum);
46+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
47+ 
48+private:
49+ TPipe pipe;
50+ TQue<QuePosition::VECIN, BUFFER_NUM> inputQueue;
51+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueue;
52+ TBuf<QuePosition::VECCALC> tmpBuf;
53+ TBuf<QuePosition::VECCALC> maskBuf;
54+ 
55+ GlobalTensor<T> inputGM;
56+ GlobalTensor<T> outputGM;
57+ 
58+ int64_t blockLength_ = 0;
59+ int64_t ubLength_ = 0;
60+ float lambd_;
61+};
62+ 
63+template <typename T, int BUFFER_MODE>
64+__aicore__ inline void Softshrink<T, BUFFER_MODE>::Init(
65+ GM_ADDR inputX, GM_ADDR outputY,
66+ const SoftshrinkTilingData* tilingData)
67+{
68+ int64_t remainderLength =
69+ tilingData->totalNum - tilingData->blockFactor * GetBlockIdx();
70+ blockLength_ = (remainderLength > tilingData->blockFactor)
71+ ? tilingData->blockFactor : remainderLength;
72+ ubLength_ = tilingData->ubFactor;
73+ lambd_ = tilingData->lambd;
74+ 
75+ int64_t offset = tilingData->blockFactor * GetBlockIdx();
76+ inputGM.SetGlobalBuffer((__gm__ T*)inputX + offset, blockLength_);
77+ outputGM.SetGlobalBuffer((__gm__ T*)outputY + offset, blockLength_);
78+ 
79+ pipe.InitBuffer(inputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
80+ pipe.InitBuffer(outputQueue, BUFFER_NUM, ubLength_ * sizeof(T));
81+ 
82+ if constexpr (NEED_CAST) {
83+ // fp16/bf16: need fp32Buf for computation (reused from tmpBuf allocation)
84+ // tmpBuf serves as both fp32 workspace and tmp buffer
85+ // Allocate enough for 2 fp32 buffers (fp32Buf + tmp) back-to-back
86+ pipe.InitBuffer(tmpBuf, 2 * ubLength_ * sizeof(float));
87+ } else {
88+ pipe.InitBuffer(tmpBuf, ubLength_ * sizeof(T));
89+ }
90+ 
91+ int64_t maskBytes = (ubLength_ + 7) / 8;
92+ maskBytes = (maskBytes + 255) / 256 * 256;
93+ pipe.InitBuffer(maskBuf, maskBytes);
94+}
95+ 
96+template <typename T, int BUFFER_MODE>
97+__aicore__ inline void Softshrink<T, BUFFER_MODE>::CopyIn(
98+ int64_t progress, int64_t currentNum)
99+{
100+ LocalTensor<T> inputLocal = inputQueue.template AllocTensor<T>();
101+ 
102+ DataCopyParams copyParams;
103+ copyParams.blockCount = 1;
104+ copyParams.blockLen = currentNum * sizeof(T);
105+ copyParams.srcStride = 0;
106+ copyParams.dstStride = 0;
107+ 
108+ DataCopyPad(inputLocal, inputGM[progress * ubLength_], copyParams,
109+ {false, 0, 0, 0});
110+ 
111+ inputQueue.EnQue(inputLocal);
112+}
113+ 
114+template <typename T, int BUFFER_MODE>
115+__aicore__ inline void Softshrink<T, BUFFER_MODE>::Compute(
116+ int64_t currentNum)
117+{
118+ LocalTensor<T> inputLocal = inputQueue.template DeQue<T>();
119+ LocalTensor<T> resultLocal = outputQueue.template AllocTensor<T>();
120+ 
121+ constexpr int64_t ALIGN_ELEM = 256 / static_cast<int64_t>(sizeof(ComputeType));
122+ int64_t alignedNum = (currentNum + ALIGN_ELEM - 1) / ALIGN_ELEM * ALIGN_ELEM;
123+ 
124+ LocalTensor<uint8_t> maskLocal = maskBuf.Get<uint8_t>();
125+ 
126+ if constexpr (NEED_CAST) {
127+ // fp16/bf16: cast to fp32, compute, cast back
128+ // tmpBuf is allocated as 2 * ubLength_ * sizeof(float)
129+ // Split into fp32Buf (first half) and fp32Tmp (second half)
130+ LocalTensor<float> fp32Buf = tmpBuf.Get<float>();
131+ LocalTensor<float> fp32Tmp = fp32Buf[ubLength_];
132+ 
133+ // Cast input to fp32
134+ Cast(fp32Buf, inputLocal, RoundMode::CAST_NONE, alignedNum);
135+ PipeBarrier<PIPE_V>();
136+ 
137+ // Step 1: Positive branch — mask = x > lambd; tmp = x - lambd; select
138+ Compares(maskLocal, fp32Buf, lambd_, CMPMODE::GT, alignedNum);
139+ Adds(fp32Tmp, fp32Buf, -lambd_, alignedNum);
140+ PipeBarrier<PIPE_V>();
141+ Select(fp32Tmp, maskLocal, fp32Tmp, 0.0f,
142+ SELMODE::VSEL_TENSOR_SCALAR_MODE, alignedNum);
143+ PipeBarrier<PIPE_V>();
144+ 
145+ // Step 2: Negative branch — mask = x < -lambd; fp32Buf = x + lambd; select
146+ Compares(maskLocal, fp32Buf, -lambd_, CMPMODE::LT, alignedNum);
147+ Adds(fp32Buf, fp32Buf, lambd_, alignedNum);
148+ PipeBarrier<PIPE_V>();
149+ Select(fp32Buf, maskLocal, fp32Buf, 0.0f,
150+ SELMODE::VSEL_TENSOR_SCALAR_MODE, alignedNum);
151+ PipeBarrier<PIPE_V>();
152+ 
153+ // Step 3: Merge — result = positive_branch + negative_branch
154+ Add(fp32Buf, fp32Buf, fp32Tmp, alignedNum);
155+ PipeBarrier<PIPE_V>();
156+ 
157+ // Cast back to original type
158+ Cast(resultLocal, fp32Buf, RoundMode::CAST_ROUND, alignedNum);
159+ PipeBarrier<PIPE_V>();
160+ } else {
161+ // fp32: compute directly
162+ LocalTensor<T> tmpLocal = tmpBuf.Get<T>();
163+ 
164+ // Step 1: Positive branch — mask = x > lambd; tmp = x - lambd; select
165+ Compares(maskLocal, inputLocal, lambd_, CMPMODE::GT, alignedNum);
166+ Adds(tmpLocal, inputLocal, (T)(-lambd_), alignedNum);
167+ PipeBarrier<PIPE_V>();
168+ Select(tmpLocal, maskLocal, tmpLocal, (T)0.0,
169+ SELMODE::VSEL_TENSOR_SCALAR_MODE, alignedNum);
170+ PipeBarrier<PIPE_V>();
171+ 
172+ // Step 2: Negative branch — mask = x < -lambd; result = x + lambd; select
173+ Compares(maskLocal, inputLocal, (T)(-lambd_), CMPMODE::LT, alignedNum);
174+ Adds(resultLocal, inputLocal, (T)(lambd_), alignedNum);
175+ PipeBarrier<PIPE_V>();
176+ Select(resultLocal, maskLocal, resultLocal, (T)0.0,
177+ SELMODE::VSEL_TENSOR_SCALAR_MODE, alignedNum);
178+ PipeBarrier<PIPE_V>();
179+ 
180+ // Step 3: Merge — result = positive_branch + negative_branch
181+ Add(resultLocal, resultLocal, tmpLocal, alignedNum);
182+ PipeBarrier<PIPE_V>();
183+ }
184+ 
185+ outputQueue.template EnQue<T>(resultLocal);
186+ inputQueue.FreeTensor(inputLocal);
187+}
188+ 
189+template <typename T, int BUFFER_MODE>
190+__aicore__ inline void Softshrink<T, BUFFER_MODE>::CopyOut(
191+ int64_t progress, int64_t currentNum)
192+{
193+ LocalTensor<T> resultLocal = outputQueue.template DeQue<T>();
194+ DataCopyParams copyParams;
195+ copyParams.blockCount = 1;
196+ copyParams.blockLen = currentNum * sizeof(T);
197+ copyParams.srcStride = 0;
198+ copyParams.dstStride = 0;
199+ DataCopyPad(outputGM[progress * ubLength_], resultLocal, copyParams);
200+ outputQueue.FreeTensor(resultLocal);
201+}
202+ 
203+template <typename T, int BUFFER_MODE>
204+__aicore__ inline void Softshrink<T, BUFFER_MODE>::Process()
205+{
206+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
207+ 
208+ if constexpr (BUFFER_NUM == 2) {
209+ int64_t curNum0 = (loopCount == 1) ? blockLength_ : ubLength_;
210+ if (loopCount >= 2) {
211+ CopyIn(0, curNum0);
212+ for (int64_t i = 1; i < loopCount; i++) {
213+ int64_t prevNum = (i == 1) ? curNum0 :
214+ ((i - 1 == loopCount - 1) ? (blockLength_ - ubLength_ * (i - 1)) : ubLength_);
215+ int64_t curNum = (i == loopCount - 1)
216+ ? (blockLength_ - ubLength_ * i) : ubLength_;
217+ CopyIn(i, curNum);
218+ Compute(prevNum);
219+ CopyOut(i - 1, prevNum);
220+ }
221+ int64_t lastNum = blockLength_ - ubLength_ * (loopCount - 1);
222+ Compute(lastNum);
223+ CopyOut(loopCount - 1, lastNum);
224+ } else if (loopCount == 1) {
225+ CopyIn(0, curNum0);
226+ Compute(curNum0);
227+ CopyOut(0, curNum0);
228+ }
229+ } else {
230+ for (int64_t i = 0; i < loopCount; i++) {
231+ int64_t currentNum = (i == (loopCount - 1))
232+ ? (blockLength_ - ubLength_ * i) : ubLength_;
233+ CopyIn(i, currentNum);
234+ Compute(currentNum);
235+ CopyOut(i, currentNum);
236+ }
237+ }
238+}
239+ 
240+} // namespace NsSoftshrink
241+#endif // SOFTSHRINK_H
Aexperimental/activation/softshrink/op_kernel/softshrink_tiling_data.h+24-0
@@ -0,0 +1,24 @@
1+/**
2+ * Copyright (c) 2025-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+ * NOTE: Portions of this code were AI-generated and have been
11+ * technically reviewed for functional accuracy and security
12+ */
13+ 
14+#ifndef _SOFTSHRINK_TILING_DATA_H_
15+#define _SOFTSHRINK_TILING_DATA_H_
16+ 
17+struct SoftshrinkTilingData {
18+ int64_t totalNum = 0;
CANN-robot
CANN-robotCANN-robot3月28日
代码规范/可移植性: 在C语言结构体定义中使用了C++风格的成员初始化（int64_t totalNum = 0;）。C语言标准（C89/C99/C11）不支持在结构体声明中直接初始化成员变量。这会导致在C编译器下编译失败，或在某些C编译器中产生语法错误。虽然代码可能在某些支持C++扩展的C编译器下工作，但这会破坏代码的可移植性，特别是在需要严格C语言编译的环境中。
问题类型: 代码规范/可移植性
文件路径: experimental/activation/softshrink/op_kernel/softshrink_tiling_data.h
行号: 24
问题代码:
    int64_t totalNum = 0;
修改建议:
移除结构体成员声明中的初始化值，改为在结构体变量定义时或使用前进行初始化。例如：
1. 在定义结构体变量时初始化：struct SoftshrinkTilingData data = {0, 0, 0, 0.5f};
2. 使用memset或逐个成员赋值进行初始化。
修改后的结构体定义应为：
struct SoftshrinkTilingData {
    int64_t totalNum;
    int64_t blockFactor;
    int64_t ubFactor;
    float lambd;
};
---
此评论由代码审查工具自动生成
likedislike
19+ int64_t blockFactor = 0;
20+ int64_t ubFactor = 0;
21+ float lambd = 0.5f;
CANN-robot
CANN-robotCANN-robot3月28日
代码规范/可移植性: 在C语言结构体定义中使用了C++风格的成员初始化（float lambd = 0.5f;）。与上一个问题相同，C语言标准不支持这种语法。这会导致代码在纯C编译环境中编译失败，影响代码的可移植性和跨平台兼容性。
问题类型: 代码规范/可移植性
文件路径: experimental/activation/softshrink/op_kernel/softshrink_tiling_data.h
行号: 27
问题代码:
    float lambd = 0.5f;
修改建议:
移除成员初始化，将默认值0.5f的使用移到结构体初始化的地方。例如在创建结构体实例时指定该值：struct SoftshrinkTilingData data = {0, 0, 0, 0.5f};。或者，如果0.5f是硬编码的默认值，可以在使用该结构体的代码中显式设置。
---
此评论由代码审查工具自动生成
likedislike
22+};
23+ 
24+#endif
Aexperimental/activation/softshrink/op_kernel/softshrink_tiling_key.h+39-0
@@ -0,0 +1,39 @@
1+/**
2+ * Copyright (c) 2025-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+ * NOTE: Portions of this code were AI-generated and have been
11+ * technically reviewed for functional accuracy and security
12+ */
13+ 
14+#ifndef __SOFTSHRINK_TILING_KEY_H__
CANN-robot
CANN-robotCANN-robot3月28日
代码结构与可维护性: 头文件保护宏使用了双下划线开头（__SOFTSHRINK_TILING_KEY_H__）。根据C/C++标准，以双下划线开头或包含双下划线的标识符是保留给实现使用的。在用户代码中使用这样的标识符可能导致与系统头文件或编译器内部定义冲突，是未定义行为。
问题类型: 代码结构与可维护性
文件路径: experimental/activation/softshrink/op_kernel/softshrink_tiling_key.h
行号: 20
问题代码:
#ifndef __SOFTSHRINK_TILING_KEY_H__
#define __SOFTSHRINK_TILING_KEY_H__
修改建议:
将头文件保护宏改为不以双下划线开头，例如：SOFTSHRINK_TILING_KEY_H_ 或 SOFT_SHRINK_TILING_KEY_H。同时确保项目中的所有头文件都遵循相同的命名约定。
---
此评论由代码审查工具自动生成
likedislike
15+#define __SOFTSHRINK_TILING_KEY_H__
16+ 
17+#include "ascendc/host_api/tiling/template_argument.h"
18+ 
19+ASCENDC_TPL_ARGS_DECL(Softshrink,
20+ ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16, ASCENDC_TPL_INPUT(0)),
21+ ASCENDC_TPL_UINT_DECL(BUFFER_MODE, 8, ASCENDC_TPL_UI_LIST, 0, 1)
22+);
23+ 
24+ASCENDC_TPL_SEL(
25+ ASCENDC_TPL_ARGS_SEL(
26+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT),
27+ ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)
28+ ),
29+ ASCENDC_TPL_ARGS_SEL(
30+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16),
31+ ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)
32+ ),
33+ ASCENDC_TPL_ARGS_SEL(
34+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_BF16),
35+ ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)
36+ ),
37+);
38+ 
39+#endif