[CANNBot]Experimental 新增Shrink子AscendC适配 Ascend950-ops-nn-AtomGit

Aexperimental/activation/shrink/CMakeLists.txt+21-0

@@ -0,0 +1,21 @@
	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	+# NOTE: Portions of this code were AI-generated and have been
	11	+# technically reviewed for functional accuracy and security
	12	+
	13	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	14	+if(NOT ENABLE_TEST)
	15	+ list(REMOVE_ITEM CURRENT_DIRS tests)
	16	+endif()
	17	+foreach(SUB_DIR ${CURRENT_DIRS})
	18	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	19	+ add_subdirectory(${SUB_DIR})
	20	+ endif()
	21	+endforeach()

Aexperimental/activation/shrink/README.md+48-0

@@ -0,0 +1,48 @@
	1	+# Shrink
	2	+
	3	+## 算子说明
	4	+
	5	+对输入张量进行非线性收缩变换。根据输入值与阈值 `lambd` 的大小关系，通过偏移量 `bias` 进行缩放和偏移处理。
	6	+
	7	+### 计算公式
	8	+
	9	+```
	10	+out_i = self_i - bias, if self_i > lambd
	11	+out_i = self_i + bias, if self_i < -lambd
	12	+out_i = 0, if -lambd <= self_i <= lambd
	13	+```
	14	+
	15	+## 支持平台
	16	+
	17	+\| 平台 \| 支持 \|
	18	+\|------\|------\|
	19	+\| Ascend950 \| Yes \|
	20	+
	21	+## 数据类型
	22	+
	23	+\| 输入 dtype \| 输出 dtype \|
	24	+\|-----------\|-----------\|
	25	+\| FLOAT16 \| FLOAT16 \|
	26	+\| FLOAT \| FLOAT \|
	27	+
	28	+## 目录结构
	29	+
	30	+```
	31	+shrink/
	32	+├── CMakeLists.txt
	33	+├── README.md
	34	+├── examples/
	35	+│ └── test_aclnn_shrink.cpp
	36	+├── op_host/
	37	+│ ├── CMakeLists.txt
	38	+│ ├── shrink_def.cpp
	39	+│ ├── shrink_infershape.cpp
	40	+│ └── shrink_tiling.cpp
	41	+├── op_kernel/
	42	+│ ├── shrink.cpp
	43	+│ ├── shrink.h
	44	+│ ├── shrink_tiling_data.h
	45	+│ └── shrink_tiling_key.h
	46	+└── tests/
	47	+ └── .gitkeep
	48	+```

Aexperimental/activation/shrink/examples/arch35/test_aclnn_shrink.cpp+144-0

@@ -0,0 +1,144 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+#include <iostream>
	16	+#include <vector>
	17	+#include "acl/acl.h"
	18	+#include "aclnn_shrink.h"
	19	+
	20	+
	21	+#define CHECK_RET(cond, return_expr) \
	22	+ do { \
	23	+ if (!(cond)) { \
	24	+ return_expr; \
	25	+ } \
	26	+} while (0)
	27	+
	28	+#define LOG_PRINT(message, ...) \
	29	+ do { \
	30	+ printf(message, ##__VA_ARGS__); \
	31	+} while (0)
	32	+
	33	+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
	34	+ int64_t shape_size = 1;
	35	+ for (auto i : shape) {
	36	+ shape_size *= i;
	37	+ }
	38	+ return shape_size;
	39	+}
	40	+
	41	+int Init(int32_t deviceId, aclrtStream* stream) {
	42	+ // 固定写法，资源初始化
	43	+ auto ret = aclInit(nullptr);
	44	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	45	+ ret = aclrtSetDevice(deviceId);
	46	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	47	+ ret = aclrtCreateStream(stream);
	48	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	49	+ return 0;
	50	+}
	51	+
	52	+template <typename T>
	53	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
	54	+ aclDataType dataType, aclTensor** tensor) {
	55	+ auto size = GetShapeSize(shape) * sizeof(T);
	56	+ // 调用aclrtMalloc申请Device侧内存
	57	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	58	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	59	+ // 调用aclrtMemcpy将Host侧数据拷贝到Device侧内存上
	60	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	61	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	62	+ // 计算连续tensor的strides
	63	+ std::vector<int64_t> strides(shape.size(), 1);
	64	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	65	+ strides[i] = shape[i + 1] * strides[i + 1];
	66	+ }
	67	+ // 调用aclCreateTensor接口创建aclTensor
	68	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	69	+ shape.data(), shape.size(), *deviceAddr);
	70	+ return 0;
	71	+}
	72	+
	73	+int main() {
	74	+ // 1. （固定写法）device/stream初始化, 参考acl对外接口列表
	75	+ // 根据自己的实际device填写deviceId
	76	+ int32_t deviceId = 0;
	77	+ aclrtStream stream;
	78	+ auto ret = Init(deviceId, &stream);
	79	+ // check根据自己的需要处理
	80	+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	81	+
	82	+ // 2. 构造输入与输出，需要根据API的接口自定义构造
	83	+ std::vector<int64_t> selfShape = {4, 2};
	84	+ std::vector<int64_t> outShape = {4, 2};
	85	+ void* selfDeviceAddr = nullptr;
	86	+ void* outDeviceAddr = nullptr;
	87	+ aclTensor* self = nullptr;
	88	+ aclTensor* out = nullptr;
	89	+ std::vector<float> selfHostData = {-3, -2, -1, 0, 1, 2, 3, -0.5};
	90	+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
	91	+ double lambd = 1.0;
	92	+ double bias = 1.0;
	93	+ // 创建self aclTensor
	94	+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
	95	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	96	+ // 创建out aclTensor
	97	+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
	98	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	99	+
	100	+ // 3. 调用CANN算子库API
	101	+ uint64_t workspaceSize = 0;
	102	+ aclOpExecutor* executor;
	103	+ // 调用aclnnShrink第一段接口
	104	+ ret = aclnnShrinkGetWorkspaceSize(self, lambd, bias, out, &workspaceSize, &executor);
	105	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnShrinkGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
	106	+ // 根据第一段接口计算出的workspaceSize申请device内存
	107	+ void* workspaceAddr = nullptr;
	108	+ if (workspaceSize > 0) {
	109	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	110	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret;);
	111	+ }
	112	+ // 调用aclnnShrink第二段接口
	113	+ ret = aclnnShrink(workspaceAddr, workspaceSize, executor, stream);
	114	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnShrink failed. ERROR: %d\n", ret); return ret);
	115	+
	116	+ // 4. （固定写法）同步等待任务执行结束
	117	+ ret = aclrtSynchronizeStream(stream);
	118	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	119	+
	120	+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
	121	+ auto size = GetShapeSize(outShape);
	122	+ std::vector<float> resultData(size, 0);
	123	+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(float),
	124	+ ACL_MEMCPY_DEVICE_TO_HOST);
	125	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
	126	+ for (int64_t i = 0; i < size; i++) {
	127	+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
	128	+ }
	129	+
	130	+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
	131	+ aclDestroyTensor(self);
	132	+ aclDestroyTensor(out);
	133	+
	134	+ // 7. 释放device资源，需要根据具体API的接口定义修改
	135	+ aclrtFree(selfDeviceAddr);
	136	+ aclrtFree(outDeviceAddr);
	137	+ if (workspaceSize > 0) {
	138	+ aclrtFree(workspaceAddr);
	139	+ }
	140	+ aclrtDestroyStream(stream);
	141	+ aclrtResetDevice(deviceId);
	142	+ aclFinalize();
	143	+ return 0;
	144	+}

Aexperimental/activation/shrink/op_host/CMakeLists.txt+13-0

@@ -0,0 +1,13 @@
	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	+# NOTE: Portions of this code were AI-generated and have been
	11	+# technically reviewed for functional accuracy and security
	12	+
	13	+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE shrink ACLNNTYPE aclnn)

Aexperimental/activation/shrink/op_host/shrink_def.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	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+
	16	+/*!
	17	+ * @file shrink_def.cpp
	18	+ * @brief Shrink operator definition, declares inputs, outputs, attributes and operator configuration
	19	+ */
	20	+#include "register/op_def_registry.h"
	21	+
	22	+namespace ops {
	23	+class Shrink : public OpDef {
	24	+public:
	25	+ explicit Shrink(const char* name) : OpDef(name)
	26	+ {
	27	+ this->Input("self")
	28	+ .ParamType(REQUIRED)
	29	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
	30	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	31	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	32	+ .AutoContiguous();
	33	+ this->Output("out")
	34	+ .ParamType(REQUIRED)
	35	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
	36	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	37	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	38	+ .AutoContiguous();
	39	+
	40	+ this->Attr("lambd").Float(1.0f);
	41	+ this->Attr("bias").Float(1.0f);
	42	+
	43	+ OpAICoreConfig aiCoreConfig;
	44	+ aiCoreConfig.DynamicCompileStaticFlag(true)
	45	+ .DynamicFormatFlag(false)
	46	+ .DynamicRankSupportFlag(true)
	47	+ .DynamicShapeSupportFlag(true)
	48	+ .NeedCheckSupportFlag(false)
	49	+ .PrecisionReduceFlag(true)
	50	+ .ExtendCfgInfo("opFile.value", "shrink");
	51	+ this->AICore().AddConfig("ascend950", aiCoreConfig);
	52	+ }
	53	+};
	54	+OP_ADD(Shrink);
	55	+} // namespace ops

Aexperimental/activation/shrink/op_host/shrink_infershape.cpp+56-0

@@ -0,0 +1,56 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+
	16	+/*!
	17	+ * @file shrink_infershape.cpp
	18	+ * @brief Shrink operator shape inference implementation
	19	+ */
	20	+#include "register/op_impl_registry.h"
	21	+#include "exe_graph/runtime/infer_shape_context.h"
	22	+
	23	+using namespace ge;
	24	+
	25	+namespace ops {
	26	+
	27	+static ge::graphStatus InferShape4Shrink(gert::InferShapeContext* context)
	28	+{
	29	+ const gert::Shape* inputShape = context->GetInputShape(0);
	30	+ if (inputShape == nullptr) {
	31	+ return ge::GRAPH_FAILED;
	32	+ }
	33	+
	34	+ // Propagate empty tensor: if input has 0 elements, output is also empty
	35	+ if (inputShape->GetShapeSize() == 0) {
	36	+ gert::Shape* outputShape = context->GetOutputShape(0);
	37	+ if (outputShape == nullptr) {
	38	+ return ge::GRAPH_FAILED;
	39	+ }
	40	+ outputShape = inputShape;
	41	+ return ge::GRAPH_SUCCESS;
	42	+ }
	43	+
	44	+ gert::Shape* outputShape = context->GetOutputShape(0);
	45	+ if (outputShape == nullptr) {
	46	+ return ge::GRAPH_FAILED;
	47	+ }
	48	+
	49	+ outputShape = inputShape;
	50	+
	51	+ return ge::GRAPH_SUCCESS;
	52	+}
	53	+
	54	+IMPL_OP_INFERSHAPE(Shrink).InferShape(InferShape4Shrink);
	55	+
	56	+} // namespace ops

Aexperimental/activation/shrink/op_host/shrink_tiling.cpp+206-0

@@ -0,0 +1,206 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * @file shrink_tiling.cpp
	13	+ * @brief Shrink Tiling implementation (arch35)
	14	+ */
	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/shrink_tiling_data.h"
	20	+#include "../op_kernel/shrink_tiling_key.h"
	21	+
	22	+#include <cmath>
	23	+
	24	+namespace optiling {
	25	+
	26	+using Ops::Base::CeilDiv;
	27	+using Ops::Base::FloorDiv;
	28	+using Ops::Base::FloorAlign;
	29	+using Ops::Base::GetUbBlockSize;
	30	+
	31	+constexpr uint32_t WS_SYS_SIZE = 0U;
	32	+// Double buffer switching threshold.
	33	+// 1024 elements is the minimum threshold; fp16 1024 elems = 2KB, fp32 1024 elems = 4KB.
	34	+// The actual optimal threshold should be determined by performance profiling;
	35	+// a higher threshold (e.g., 4096~8192) may yield better throughput.
	36	+constexpr int64_t MIN_SPLIT_THRESHOLD = 1024;
	37	+
	38	+static const gert::Shape g_vec_1_shape = {1};
	39	+
	40	+// 0-dim scalar tensors are treated as 1-dim with 1 element for Tiling.
	41	+// The inferShape preserves the original 0-dim shape, so output shape matches input.
	42	+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape) {
	43	+ if (in_shape.GetDimNum() == 0) {
	44	+ return g_vec_1_shape;
	45	+ }
	46	+ return in_shape;
	47	+}
	48	+
	49	+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
	50	+{
	51	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	52	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	53	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	54	+ coreNum = ascendcPlatform.GetCoreNumAiv();
	55	+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
	56	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
	57	+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
	58	+ return ge::GRAPH_SUCCESS;
	59	+}
	60	+
	61	+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context,
	62	+ int64_t& totalIdx,
	63	+ ge::DataType& dataType,
	64	+ float& lambd,
	65	+ float& bias)
	66	+{
	67	+ auto input = context->GetInputShape(0);
	68	+ OP_CHECK_NULL_WITH_CONTEXT(context, input);
	69	+ auto inputShape = EnsureNotScalar(input->GetStorageShape());
	70	+
	71	+ auto output = context->GetOutputShape(0);
	72	+ OP_CHECK_NULL_WITH_CONTEXT(context, output);
	73	+ auto outShape = EnsureNotScalar(output->GetStorageShape());
	74	+
	75	+ OP_CHECK_IF(
	76	+ inputShape.GetShapeSize() != outShape.GetShapeSize(),
	77	+ OP_LOGE(context, "Shrink: shape size mismatch: input=%ld, out=%ld",
	78	+ inputShape.GetShapeSize(), outShape.GetShapeSize()),
	79	+ return ge::GRAPH_FAILED);
	80	+
	81	+ totalIdx = inputShape.GetShapeSize();
	82	+
	83	+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT16, ge::DT_FLOAT};
	84	+ auto inputDesc = context->GetInputDesc(0);
	85	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	86	+ dataType = inputDesc->GetDataType();
	87	+ if (supportedDtype.count(dataType) == 0) {
	88	+ OP_LOGE(context, "Shrink: invalid dtype %d", static_cast<int>(dataType));
	89	+ return ge::GRAPH_FAILED;
	90	+ }
	91	+
	92	+ auto attrs = context->GetAttrs();
	93	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
	94	+ const float* lambdPtr = attrs->GetAttrPointer<float>(0);
	95	+ OP_CHECK_NULL_WITH_CONTEXT(context, lambdPtr);
	96	+ lambd = *lambdPtr;
	97	+
	98	+ const float* biasPtr = attrs->GetAttrPointer<float>(1);
	99	+ OP_CHECK_NULL_WITH_CONTEXT(context, biasPtr);
	100	+ bias = *biasPtr;
	101	+
	102	+ return ge::GRAPH_SUCCESS;
	103	+}
	104	+
	105	+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
	106	+{
	107	+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
	108	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	109	+ currentWorkspace[0] = WS_SYS_SIZE;
	110	+ return ge::GRAPH_SUCCESS;
	111	+}
	112	+
	113	+static int64_t CalcUbFactor(int64_t totalIdx, ge::DataType dataType, uint64_t ubSize,
	114	+ int64_t ubBlockSize, uint64_t useDoubleBuffer)
	115	+{
	116	+ int64_t typeSize = (dataType == ge::DT_FLOAT16) ? 2 : 4;
	117	+ int64_t computeAlignment = 256 / typeSize;
	118	+
	119	+ // Reserve for TPipe, queue management, Select mode 2 internal buffer, system overhead.
	120	+ // Ascend950 may need more headroom than Ascend910B.
	121	+ constexpr int64_t UB_RESERVED_BYTES = 48 * 1024;
	122	+ int64_t availableUb = static_cast<int64_t>(ubSize) - UB_RESERVED_BYTES;
	123	+ if (availableUb <= 0) {
	124	+ availableUb = static_cast<int64_t>(ubSize) / 2;
	125	+ }
	126	+
	127	+ // Buffer layout: 2 TQue (input+output) × BUFFER_NUM + 1 TBuf (mid) × 1
	128	+ int64_t numTQueBuffers = useDoubleBuffer ? 2 : 1;
	129	+ int64_t bufferBytes = (2 * numTQueBuffers + 1) * typeSize;
	130	+
	131	+ int64_t rawUbFactor = FloorDiv(availableUb, bufferBytes);
	132	+ int64_t alignFactor = (ubBlockSize > computeAlignment) ? ubBlockSize : computeAlignment;
	133	+ int64_t ubFactorAligned = FloorAlign(rawUbFactor, alignFactor);
	134	+
	135	+ if (totalIdx < ubFactorAligned) {
	136	+ ubFactorAligned = CeilDiv(totalIdx, alignFactor) * alignFactor;
	137	+ }
	138	+
	139	+ constexpr int64_t MAX_COPY_BYTES = 65535;
	140	+ int64_t maxElementsByType = MAX_COPY_BYTES / typeSize;
	141	+ if (ubFactorAligned > maxElementsByType) {
	142	+ ubFactorAligned = FloorAlign(maxElementsByType, alignFactor);
	143	+ }
	144	+ return ubFactorAligned;
	145	+}
	146	+
	147	+static ge::graphStatus ShrinkTilingFunc(gert::TilingContext* context)
	148	+{
	149	+ uint64_t ubSize = 0;
	150	+ int64_t coreNum = 0;
	151	+ OP_CHECK_IF(
	152	+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
	153	+ OP_LOGE(context, "GetPlatformInfo error"),
	154	+ return ge::GRAPH_FAILED);
	155	+
	156	+ int64_t totalIdx;
	157	+ ge::DataType dataType;
	158	+ float lambd;
	159	+ float bias;
	160	+ OP_CHECK_IF(
	161	+ GetShapeAttrsInfo(context, totalIdx, dataType, lambd, bias) != ge::GRAPH_SUCCESS,
	162	+ OP_LOGE(context, "GetShapeAttrsInfo error"),
	163	+ return ge::GRAPH_FAILED);
	164	+
	165	+ OP_CHECK_IF(
	166	+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
	167	+ OP_LOGE(context, "GetWorkspaceSize error"),
	168	+ return ge::GRAPH_FAILED);
	169	+
	170	+ ShrinkTilingData* tiling = context->GetTilingData<ShrinkTilingData>();
	171	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	172	+ OP_CHECK_IF(
	173	+ memset_s(tiling, sizeof(ShrinkTilingData), 0, sizeof(ShrinkTilingData)) != EOK,
	174	+ OP_LOGE(context, "set tiling data error"),
	175	+ return ge::GRAPH_FAILED);
	176	+
	177	+ tiling->totalNum = totalIdx;
	178	+ tiling->blockFactor = CeilDiv(totalIdx, coreNum);
	179	+ int64_t usedCoreNum = CeilDiv(totalIdx, tiling->blockFactor);
	180	+ tiling->lambd = (lambd < 0.0f) ? 0.0f : lambd;
	181	+ tiling->bias = bias;
	182	+
	183	+ int64_t ubBlockSize = GetUbBlockSize(context);
	184	+ uint64_t useDoubleBuffer = (totalIdx > MIN_SPLIT_THRESHOLD) ? 1 : 0;
	185	+ tiling->ubFactor = CalcUbFactor(totalIdx, dataType, ubSize, ubBlockSize, useDoubleBuffer);
	186	+
	187	+ context->SetBlockDim(usedCoreNum);
	188	+
	189	+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
	190	+ ASCENDC_TPL_SEL_PARAM(context, dTypeX, useDoubleBuffer);
	191	+
	192	+ return ge::GRAPH_SUCCESS;
	193	+}
	194	+
	195	+static ge::graphStatus TilingParseForShrink([[maybe_unused]] gert::TilingParseContext* context)
	196	+{
	197	+ return ge::GRAPH_SUCCESS;
	198	+}
	199	+
	200	+struct ShrinkCompileInfo {};
	201	+
	202	+IMPL_OP_OPTILING(Shrink)
	203	+ .Tiling(ShrinkTilingFunc)
	204	+ .TilingParse<ShrinkCompileInfo>(TilingParseForShrink);
	205	+
	206	+} // namespace optiling

Aexperimental/activation/shrink/op_kernel/shrink.cpp+32-0

@@ -0,0 +1,32 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+
	16	+/*!
	17	+ * @file shrink.cpp
	18	+ * @brief Shrink kernel entry
	19	+ */
	20	+#include "shrink.h"
	21	+
	22	+template <typename D_T_X, int BUFFER_MODE>
	23	+__global__ __aicore__ void shrink(
	24	+ GM_ADDR self, GM_ADDR out,
	25	+ [[maybe_unused]] GM_ADDR workspace, GM_ADDR tiling)
	26	+{
	27	+ REGISTER_TILING_DEFAULT(ShrinkTilingData);
	28	+ GET_TILING_DATA_WITH_STRUCT(ShrinkTilingData, tilingData, tiling);
	29	+ NsShrink::Shrink<D_T_X, BUFFER_MODE> op;
	30	+ op.Init(self, out, &tilingData);
	31	+ op.Process();
	32	+}

Aexperimental/activation/shrink/op_kernel/shrink.h+202-0

@@ -0,0 +1,202 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+
	16	+/*!
	17	+ * @file shrink.h
	18	+ * @brief Shrink Kernel class (arch35)
	19	+ *
	20	+ * Template parameters:
	21	+ * - T: data type (half / float)
	22	+ * - BUFFER_MODE: buffer mode (0=single, 1=double)
	23	+ *
	24	+ * Computation:
	25	+ * out = self - bias, if self > lambd
	26	+ * out = self + bias, if self < -lambd
	27	+ * out = 0, if -lambd <= self <= lambd
	28	+ *
	29	+ * Strategy (2 queues + 1 compute buffer: input, output, midBuf):
	30	+ * 1. Compares(self < -lambd) -> maskNeg
	31	+ * 2. Compares(self > lambd) -> maskPos
	32	+ * 3. outLocal = self + (-bias) using Adds
	33	+ * 4. Select using maskPos: where maskPos=1 keep self-bias, else pick 0
	34	+ * 5. midLocal = self + bias using Adds
	35	+ * 6. Select using maskNeg: where maskNeg=1 pick self+bias, else keep current
	36	+ *
	37	+ * Uses Compares for scalar comparison (outputs bitmask),
	38	+ * Select with VSEL_TENSOR_TENSOR_MODE, And for mask combination.
	39	+ */
	40	+#ifndef SHRINK_H
	41	+#define SHRINK_H
	42	+
	43	+#include "kernel_operator.h"
	44	+#include "kernel_tiling/kernel_tiling.h"
	45	+#include "shrink_tiling_data.h"
	46	+#include "shrink_tiling_key.h"
	47	+
	48	+namespace NsShrink {
	49	+
	50	+using namespace AscendC;
	51	+
	52	+template <typename T, int BUFFER_MODE>
	53	+class Shrink {
	54	+ static constexpr int32_t BUFFER_NUM = BUFFER_MODE ? 2 : 1;
	55	+
	56	+public:
	57	+ __aicore__ inline Shrink() {};
	58	+
	59	+ __aicore__ inline void Init(GM_ADDR self, GM_ADDR out,
	60	+ const ShrinkTilingData* tilingData);
	61	+ __aicore__ inline void Process();
	62	+
	63	+private:
	64	+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
	65	+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
	66	+ __aicore__ inline void Compute(int64_t currentNum);
	67	+
	68	+private:
	69	+ TPipe pipe_;
	70	+ // 2 queues (input, output) + 1 compute buffer (mid)
	71	+ TQue<QuePosition::VECIN, BUFFER_NUM> inputQueue_;
	72	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueue_;
	73	+
	74	+ // Mid buffer: pure compute, no MTE pipeline needed, saves double-buffer overhead
	75	+ TBuf<TPosition::VECCALC> midBuf_;
	76	+
	77	+ // Mask buffers for Compares output (bitmask: 1 bit per element)
	78	+ TBuf<TPosition::VECCALC> maskBuf1_;
	79	+ TBuf<TPosition::VECCALC> maskBuf2_;
	80	+
	81	+ GlobalTensor<T> gmIn_;
	82	+ GlobalTensor<T> gmOut_;
	83	+
	84	+ int64_t blockLength_ = 0;
	85	+ int64_t ubLength_ = 0;
	86	+ float lambd_ = 1.0f;
	87	+ float bias_ = 1.0f;
	88	+};
	89	+
	90	+template <typename T, int BUFFER_MODE>
	91	+__aicore__ inline void Shrink<T, BUFFER_MODE>::Init(
	92	+ GM_ADDR self, GM_ADDR out,
	93	+ const ShrinkTilingData* tilingData)
	94	+{
	95	+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx();
	96	+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
	97	+ ubLength_ = tilingData->ubFactor;
	98	+ lambd_ = tilingData->lambd;
	99	+ bias_ = tilingData->bias;
	100	+
	101	+ int64_t offset = tilingData->blockFactor * AscendC::GetBlockIdx();
	102	+ gmIn_.SetGlobalBuffer((__gm__ T*)self + offset, blockLength_);
	103	+ gmOut_.SetGlobalBuffer((__gm__ T*)out + offset, blockLength_);
	104	+
	105	+ // 2 data queues + 1 compute buffer
	106	+ pipe_.InitBuffer(inputQueue_, BUFFER_NUM, ubLength_ * sizeof(T));
	107	+ pipe_.InitBuffer(outputQueue_, BUFFER_NUM, ubLength_ * sizeof(T));
	108	+ pipe_.InitBuffer(midBuf_, ubLength_ * sizeof(T));
	109	+
	110	+ // Mask buffers: ubLength/8 bytes, 32-byte aligned, minimum 32 bytes
	111	+ int64_t maskBytes = (ubLength_ / 8 + 31) & ~31;
	112	+ if (maskBytes < 32) maskBytes = 32;
	113	+ pipe_.InitBuffer(maskBuf1_, maskBytes);
	114	+ pipe_.InitBuffer(maskBuf2_, maskBytes);
	115	+}
	116	+
	117	+template <typename T, int BUFFER_MODE>
	118	+__aicore__ inline void Shrink<T, BUFFER_MODE>::CopyIn(int64_t progress, int64_t currentNum)
	119	+{
	120	+ LocalTensor<T> inLocal = inputQueue_.template AllocTensor<T>();
	121	+
	122	+ // byteLen is uint16_t (max 65535); Tiling ensures currentNum * sizeof(T) <= 65535
	123	+ DataCopyPad(inLocal, gmIn_[progress * ubLength_],
	124	+ {1, static_cast<uint16_t>(currentNum * sizeof(T)), 0, 0},
	125	+ {false, 0, 0, 0});
	126	+
	127	+ inputQueue_.EnQue(inLocal);
	128	+}
	129	+
	130	+template <typename T, int BUFFER_MODE>
	131	+__aicore__ inline void Shrink<T, BUFFER_MODE>::CopyOut(int64_t progress, int64_t currentNum)
	132	+{
	133	+ LocalTensor<T> outLocal = outputQueue_.template DeQue<T>();
	134	+
	135	+ // byteLen is uint16_t (max 65535); Tiling ensures currentNum * sizeof(T) <= 65535
	136	+ DataCopyPad(gmOut_[progress * ubLength_], outLocal,
	137	+ {1, static_cast<uint16_t>(currentNum * sizeof(T)), 0, 0});
	138	+
	139	+ outputQueue_.FreeTensor(outLocal);
	140	+}
	141	+
	142	+template <typename T, int BUFFER_MODE>
	143	+__aicore__ inline void Shrink<T, BUFFER_MODE>::Compute(int64_t currentNum)
	144	+{
	145	+ LocalTensor<T> inLocal = inputQueue_.template DeQue<T>();
	146	+ LocalTensor<T> midLocal = midBuf_.template Get<T>();
	147	+ LocalTensor<T> outLocal = outputQueue_.template AllocTensor<T>();
	148	+ LocalTensor<uint8_t> maskNeg = maskBuf1_.template Get<uint8_t>();
	149	+ LocalTensor<uint8_t> maskPos = maskBuf2_.template Get<uint8_t>();
	150	+
	151	+ // computeNum: 256B-aligned count for Compares (Level 2 API requires count*sizeof(T) % 256 == 0)
	152	+ // Adds/Duplicate/Select use currentNum (actual valid element count) to avoid processing
	153	+ // invalid padding data in the last chunk.
	154	+ int64_t typeSize = sizeof(T);
	155	+ int64_t computeAlignment = 256 / typeSize;
	156	+ int64_t computeNum = ((currentNum + computeAlignment - 1) / computeAlignment) * computeAlignment;
	157	+
	158	+ // Phase 1: Compute masks from original self
	159	+ Compares(maskNeg, inLocal, static_cast<T>(-lambd_), CMPMODE::LT, computeNum);
	160	+ Compares(maskPos, inLocal, static_cast<T>(lambd_), CMPMODE::GT, computeNum);
	161	+
	162	+ // Phase 2: Compute result incrementally
	163	+ // Step 1: outLocal = self - bias (using Adds with negated bias)
	164	+ T negBias = static_cast<T>(-bias_);
	165	+ Adds(outLocal, inLocal, negBias, currentNum);
	166	+
	167	+ // Step 2: Duplicate midLocal as zero buffer for dead zone
	168	+ Duplicate(midLocal, static_cast<T>(0.0f), currentNum);
	169	+
	170	+ // Step 3: Where maskPos=0 (not positive region), pick 0; maskPos=1 keep self-bias
	171	+ Select(outLocal, maskPos, outLocal, midLocal,
	172	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, currentNum);
	173	+
	174	+ // Step 4: midLocal = self + bias (negative region value, inLocal still has original self)
	175	+ T posBias = static_cast<T>(bias_);
	176	+ Adds(midLocal, inLocal, posBias, currentNum);
	177	+
	178	+ // Step 5: Where maskNeg=1 (negative region), pick self+bias; else keep current
	179	+ Select(outLocal, maskNeg, midLocal, outLocal,
	180	+ SELMODE::VSEL_TENSOR_TENSOR_MODE, currentNum);
	181	+
	182	+ outputQueue_.template EnQue<T>(outLocal);
	183	+ inputQueue_.FreeTensor(inLocal);
	184	+}
	185	+
	186	+template <typename T, int BUFFER_MODE>
	187	+__aicore__ inline void Shrink<T, BUFFER_MODE>::Process()
	188	+{
	189	+ if (blockLength_ <= 0) return;
	190	+
	191	+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
	192	+ for (int64_t i = 0; i < loopCount; i++) {
	193	+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
	194	+ CopyIn(i, currentNum);
	195	+ Compute(currentNum);
	196	+ CopyOut(i, currentNum);
	197	+ }
	198	+}
	199	+
	200	+} // namespace NsShrink
	201	+
	202	+#endif // SHRINK_H

Aexperimental/activation/shrink/op_kernel/shrink_tiling_data.h+31-0

@@ -0,0 +1,31 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+
	16	+/*!
	17	+ * @file shrink_tiling_data.h
	18	+ * @brief Shrink TilingData structure definition
	19	+ */
	20	+#ifndef _SHRINK_TILING_DATA_H_
	21	+#define _SHRINK_TILING_DATA_H_
	22	+
	23	+struct ShrinkTilingData {
	24	+ int64_t totalNum = 0; // Total number of elements
	25	+ int64_t blockFactor = 0; // Elements processed per core
	26	+ int64_t ubFactor = 0; // Elements processed per UB loop (aligned to 256B)
	27	+ float lambd = 1.0f; // Threshold (preprocessed to non-negative)
	28	+ float bias = 1.0f; // Offset/bias parameter
	29	+};
	30	+
	31	+#endif // _SHRINK_TILING_DATA_H_

Aexperimental/activation/shrink/op_kernel/shrink_tiling_key.h+45-0

@@ -0,0 +1,45 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * NOTE: Portions of this code were AI-generated and have been
	13	+ * technically reviewed for functional accuracy and security
	14	+ */
	15	+
	16	+/*!
	17	+ * @file shrink_tiling_key.h
	18	+ * @brief Shrink Tiling template parameter definition
	19	+ *
	20	+ * Template parameters:
	21	+ * - D_T_X: data type (C_DT_FLOAT16, C_DT_FLOAT)
	22	+ * - BUFFER_MODE: buffer mode (0=single, 1=double)
	23	+ */
	24	+#ifndef __SHRINK_TILING_KEY_H__
	25	+#define __SHRINK_TILING_KEY_H__
	26	+
	27	+#include "ascendc/host_api/tiling/template_argument.h"
	28	+
	29	+ASCENDC_TPL_ARGS_DECL(Shrink,
	30	+ ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT16, C_DT_FLOAT, ASCENDC_TPL_INPUT(0)),
	31	+ ASCENDC_TPL_UINT_DECL(BUFFER_MODE, 8, ASCENDC_TPL_UI_LIST, 0, 1)
	32	+);
	33	+
	34	+ASCENDC_TPL_SEL(
	35	+ ASCENDC_TPL_ARGS_SEL(
	36	+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16),
	37	+ ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)
	38	+ ),
	39	+ ASCENDC_TPL_ARGS_SEL(
	40	+ ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT),
	41	+ ASCENDC_TPL_UINT_SEL(BUFFER_MODE, ASCENDC_TPL_UI_LIST, 0, 1)
	42	+ ),
	43	+);
	44	+
	45	+#endif // __SHRINK_TILING_KEY_H__

Aexperimental/activation/shrink/tests/.gitkeep+0-0

The file is empty

Mscripts/ci/check_example.sh+1-1

	@@ -124,7 +124,7 @@ do
124	#安装指定路径的自定义算子包	124	#安装指定路径的自定义算子包
125	echo "--------------------------------"	125	echo "--------------------------------"
126	echo "${name}"	126	echo "${name}"
127	- ./single/cann-ops-nn-${name}_linux*.run	127	+ ./single/cann-ops-nn-${name}_linux*.run --force
128	echo "[EXECUTE_COMMAND] bash build.sh --run_example $name eager cust --vendor_name=$name"	128	echo "[EXECUTE_COMMAND] bash build.sh --run_example $name eager cust --vendor_name=$name"
129	bash build.sh --run_example $name eager cust --vendor_name=$name	129	bash build.sh --run_example $name eager cust --vendor_name=$name
130	status=$?	130	status=$?