Star906
1825
代码
代码
Issues226
Pull Requests303
Actions
项目讨论
Wiki
分析
Star906
1825
已合并
Swish支持下一代芯片 #494
zhangzijie创建于 2025年12月24日
Swish支持下一代芯片 #494
已合并
zhangzijie创建于 2025年12月24日
从zhangzijie/ops-nnmaster合入到cann/ops-nnmaster
共 29 个文件变更+1406-52
Mactivation/swish/CMakeLists.txt+11-13
@@ -1,18 +1,16 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# 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 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+# ----------------------------------------------------------------------------
9-message(STATUS "=== Debug: start ops.activation.swish.CMakeLists.txt ")10+ 
10-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11+# 设置算子定义时支持的芯片类型
11-if(NOT ENABLE_TEST AND NOT BENCHMARK)12+set(SUPPORT_COMPUTE_UNIT "ascend910_95")
12- list(REMOVE_ITEM CURRENT_DIRS tests)13+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
13-endif()14+set(SUPPORT_TILING_DIR "arch35")
14-foreach(SUB_DIR ${CURRENT_DIRS})15+ 
15- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")16+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE swish ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
16- add_subdirectory(${SUB_DIR})
17- endif()
18-endforeach()
Ractivation/swish/op_host/op_api/aclnn_silu.cpp→activation/swish/op_api/aclnn_silu.cpp+9-1
@@ -39,8 +39,16 @@ static const std::initializer_list<op::DataType> ASCEND910_DTYPE_SUPPORT_LIST =
39static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {39static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {
40 op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};40 op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
41 41 
42+static inline bool CheckNotNull(const aclTensor* self, const aclTensor* out)
43+{
44+ // self、out不能为空指针
45+ OP_CHECK_NULL(self, return false);
46+ OP_CHECK_NULL(out, return false);
47+ return true;
48+}
49+ 
42static aclnnStatus CheckParams(const aclTensor *self, const aclTensor *out) {50static aclnnStatus CheckParams(const aclTensor *self, const aclTensor *out) {
43- CHECK_RET(CheckNotNull2Tensor(self, out), ACLNN_ERR_PARAM_NULLPTR);51+ CHECK_RET(CheckNotNull(self, out), ACLNN_ERR_PARAM_NULLPTR);
44 52 
45 auto supportList = GetDtypeSupportListV2(ASCEND910B_DTYPE_SUPPORT_LIST, ASCEND910_DTYPE_SUPPORT_LIST);53 auto supportList = GetDtypeSupportListV2(ASCEND910B_DTYPE_SUPPORT_LIST, ASCEND910_DTYPE_SUPPORT_LIST);
46 CHECK_RET(CheckDtypeValidActivation(self, out, supportList), ACLNN_ERR_PARAM_INVALID);54 CHECK_RET(CheckDtypeValidActivation(self, out, supportList), ACLNN_ERR_PARAM_INVALID);
Aactivation/swish/op_api/aclnn_silu.h+38-0
@@ -0,0 +1,38 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OP_API_INC_LEVEL2_ACLNN_SILU_H_
12+#define OP_API_INC_LEVEL2_ACLNN_SILU_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+ 
17+#ifdef __cplusplus
18+extern "C" {
19+#endif
20+ 
21+/**
CANN-robot
CANN-robotCANN-robot2025年12月24日

代码结构与可维护性: 函数aclnnSiluGetWorkspaceSize的注释不够详细，缺少参数说明和返回值说明。作为公共API接口，应该提供完整的文档注释，包括每个参数的含义、可能的取值范围、是否为输入/输出参数等。特别是executor参数是双重指针，需要说明其用途和生命周期管理责任。

问题类型: 代码结构与可维护性 文件路径: activation/swish/op_api/aclnn_silu.h 行号: 21 问题代码:

/**
 * @brief aclnnSilu的第一段接口，根据具体的计算流程，计算workspace大小。
 * @domain aclnn_ops_infer
 */
ACLNN_API aclnnStatus aclnnSiluGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize,
                                                aclOpExecutor** executor);

修改建议:

完善函数注释，添加参数说明和返回值说明：
/**
 * @brief 计算SiLU（Sigmoid-weighted Linear Unit）操作所需的工作空间大小
 * @param[in] self 输入张量
 * @param[in] out 输出张量
 * @param[out] workspaceSize 计算出的工作空间大小（字节数）
 * @param[out] executor 操作执行器，用于后续执行计算
 * @return 成功返回ACLNN_SUCCESS，失败返回相应的错误码
 * @note 调用者负责释放executor占用的资源（通过相应的释放函数）
 * @domain aclnn_ops_infer
 */

此评论由代码审查工具自动生成

likedislike
22+ * @brief aclnnSilu的第一段接口，根据具体的计算流程，计算workspace大小。
23+ * @domain aclnn_ops_infer
24+ */
25+ACLNN_API aclnnStatus aclnnSiluGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize,
26+ aclOpExecutor** executor);
27+ 
28+/**
29+ * @brief aclnnSilu的第二段接口，用于执行计算。
30+ */
31+ACLNN_API aclnnStatus aclnnSilu(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
32+ aclrtStream stream);
33+ 
34+#ifdef __cplusplus
35+}
36+#endif
37+ 
38+#endif
Ractivation/swish/op_host/op_api/aclnn_swish.cpp→activation/swish/op_api/aclnn_swish.cpp+0-0
文件重命名但无更改。
Aactivation/swish/op_api/aclnn_swish.h+52-0
@@ -0,0 +1,52 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OP_API_INC_LEVEL2_ACLNN_SWISH_H_
12+#define OP_API_INC_LEVEL2_ACLNN_SWISH_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+ 
17+#ifdef __cplusplus
18+extern "C" {
19+#endif
20+ 
21+/**
22+ * @brief aclnnSwish的第一段接口，根据具体的计算流程，计算workspace大小。
23+ * @domain aclnn_ops_train
24+ * 算子功能：Swish激活函数
25+ * @param [in] self: Device侧的aclTensor，公式中的input。支持非连续的Tensor，数据格式支持ND，self与out的shape和数据类型一致。
26+ * @param [in] betaOptional: Host侧的aclScalar，公式中的beta。数据类型需要是可转换为FLOAT的数据类型。
27+ * 当betaOptional为空指针时，默认值为1.0。
28+ * @param [out] out: Device侧的aclTensor，公式中的output。支持非连续的Tensor，数据格式支持ND，
29+ * self与out的shape和数据类型一致。
30+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
31+ * @param [out] executor: 返回op执行器，包含算子计算流程。
32+ * @return aclnnStatus: 返回状态码。
33+ */
34+ACLNN_API aclnnStatus aclnnSwishGetWorkspaceSize(const aclTensor* self, const aclScalar* betaOptional, aclTensor* out,
35+ uint64_t* workspaceSize, aclOpExecutor** executor);
36+ 
37+/**
38+ * @brief aclnnSwish的第二段接口，用于执行计算。
39+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
40+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnSwishGetWorkspaceSize获取。
41+ * @param [in] stream: acl stream流。
42+ * @param [in] executor: op执行器，包含了算子计算流程。
43+ * @return aclnnStatus: 返回状态码。
44+ */
45+ACLNN_API aclnnStatus aclnnSwish(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
46+ aclrtStream stream);
47+ 
48+#ifdef __cplusplus
49+}
50+#endif
51+ 
52+#endif
Ractivation/swish/op_host/op_api/silu.cpp→activation/swish/op_api/silu.cpp+0-0
文件重命名但无更改。
Aactivation/swish/op_api/swish.h+20-0
@@ -0,0 +1,20 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OP_API_INC_LEVEL0_SWISH_H_
12+#define OP_API_INC_LEVEL0_SWISH_H_
13+ 
14+#include "opdev/op_executor.h"
15+ 
16+namespace l0op {
17+const aclTensor *Swish(const aclTensor *self, float scale, aclOpExecutor *executor);
18+} // namespace l0op
19+ 
20+#endif
Aactivation/swish/op_graph/swish_proto.h+44-0
@@ -0,0 +1,44 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+ 
12+/*!
13+ * \file nonlinear_fuc_ops.h
14+ * \brief
15+ */
16+#ifndef OPS_ACTIVATION_SWISH_OPS_H_
17+#define OPS_ACTIVATION_SWISH_OPS_H_
18+ 
19+#include "graph/operator_reg.h"
20+ 
21+namespace ge {
22+/**
23+*@brief Computes the for the Swish of "x" .
24+ 
25+*@par Inputs:
26+*One input, including:
27+* x: A tensor, which supports 1D-8D defaultly and must be one of the following types: float16, bfloat16, float32. \n
28+ 
29+*@par Outputs:
30+* y: A tensor of the same type, shape and format as "x", and y = x / (1 + e ^ (-scale * x)). \n
31+ 
32+*@par Attributes:
33+* scale: scalar parameter, the multiplier of x. Must be one of the following types: float. Default value = 1.0. \n
34+ 
35+*@par Third-party framework compatibility
36+*Compatible with the Torch operator Swish
37+*/
38+REG_OP(Swish)
39+ .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
40+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
41+ .ATTR(scale, Float, 1.0)
42+ .OP_END_FACTORY_REG(Swish)
43+} // namespace ge
CANN-robot
CANN-robotCANN-robot2025年12月24日

代码结构与可维护性: 文件末尾缺少换行符。根据POSIX标准，文本文件的每一行（包括最后一行）都应以换行符结尾。缺少换行符可能导致某些工具（如编译器、版本控制系统）在处理文件时产生警告或意外行为。

问题类型: 代码结构与可维护性 文件路径: activation/swish/op_graph/swish_proto.h 行号: 44 问题代码:

} // namespace ge
#endif  // OPS_ACTIVATION_SWISH_OPS_H_

修改建议:

在文件末尾添加一个换行符。

此评论由代码审查工具自动生成

likedislike
44+#endif // OPS_ACTIVATION_SWISH_OPS_H_
Dactivation/swish/op_host/CMakeLists.txt+0-9
@@ -1,9 +0,0 @@
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
3-# CANN Open Software License Agreement Version 2.0 (the "License").
4-# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7-# See LICENSE in the root of the software repository for the full text of the License.
8-#/
9-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE swish ACLNNTYPE aclnn_exclude)
Aactivation/swish/op_host/config/ascend910_95/swish_binary.json+107-0
@@ -0,0 +1,107 @@
1+{
2+ "op_type": "Swish",
3+ "op_list": [
4+ {
5+ "bin_filename": "Swish_299cf4d7f9ef0f745d89d88c666b45aa",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "bfloat16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ }
17+ ],
18+ "outputs": [
19+ {
20+ "name": "y",
21+ "index": 0,
22+ "dtype": "bfloat16",
23+ "format": "ND",
24+ "paramType": "required",
25+ "shape": [
26+ -2
27+ ]
28+ }
29+ ],
30+ "attrs": [
31+ {
32+ "name": "scale",
33+ "dtype": "float",
34+ "value": null
35+ }
36+ ]
37+ },
38+ {
39+ "bin_filename": "Swish_a8d63cbe8795eef99c620173d4757e59",
40+ "inputs": [
41+ {
42+ "name": "x",
43+ "index": 0,
44+ "dtype": "float16",
45+ "format": "ND",
46+ "paramType": "required",
47+ "shape": [
48+ -2
49+ ]
50+ }
51+ ],
52+ "outputs": [
53+ {
54+ "name": "y",
55+ "index": 0,
56+ "dtype": "float16",
57+ "format": "ND",
58+ "paramType": "required",
59+ "shape": [
60+ -2
61+ ]
62+ }
63+ ],
64+ "attrs": [
65+ {
66+ "name": "scale",
67+ "dtype": "float",
68+ "value": null
69+ }
70+ ]
71+ },
72+ {
73+ "bin_filename": "Swish_291664ae3e88b38b8b46c2efbbea27c9",
74+ "inputs": [
75+ {
76+ "name": "x",
77+ "index": 0,
78+ "dtype": "float32",
79+ "format": "ND",
80+ "paramType": "required",
81+ "shape": [
82+ -2
83+ ]
84+ }
85+ ],
86+ "outputs": [
87+ {
88+ "name": "y",
89+ "index": 0,
90+ "dtype": "float32",
91+ "format": "ND",
92+ "paramType": "required",
93+ "shape": [
94+ -2
95+ ]
96+ }
97+ ],
98+ "attrs": [
99+ {
100+ "name": "scale",
101+ "dtype": "float",
102+ "value": null
103+ }
104+ ]
105+ }
106+ ]
107+}
Aactivation/swish/op_host/config/ascend910_95/swish_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascnedxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[Swish]
13+default=0
Mactivation/swish/op_host/op_api/aclnn_swish.h+1-1
@@ -37,7 +37,7 @@ ACLNN_API aclnnStatus aclnnSwishGetWorkspaceSize(const aclTensor* self, const ac
37/**37/**
38 * @brief aclnnSwish的第二段接口，用于执行计算。38 * @brief aclnnSwish的第二段接口，用于执行计算。
39 * @param [in] workspace: 在npu device侧申请的workspace内存起址。39 * @param [in] workspace: 在npu device侧申请的workspace内存起址。
40- * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnAcosGetWorkspaceSize获取。40+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnSwishGetWorkspaceSize获取。
41 * @param [in] stream: acl stream流。41 * @param [in] stream: acl stream流。
42 * @param [in] executor: op执行器，包含了算子计算流程。42 * @param [in] executor: op执行器，包含了算子计算流程。
43 * @return aclnnStatus: 返回状态码。43 * @return aclnnStatus: 返回状态码。
Aactivation/swish/op_host/op_tiling/arch35/swish_tiling_arch35.cpp+184-0
@@ -0,0 +1,184 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file swish_tiling_arch35.cpp
13+ * \brief
14+ */
15+ 
16+#include "swish_tiling_arch35.h"
17+#include "log/log.h"
18+#include "error_util.h"
19+#include "platform/platform_ascendc.h"
20+#include "register/op_def_registry.h"
21+#include "activation/swish/op_kernel/arch35/swish_dag.h"
22+#include "activation/swish/op_kernel/arch35/swish_struct.h"
23+ 
24+using namespace AscendC;
25+using namespace ge;
26+using namespace SwishOp;
27+ 
28+namespace optiling {
29+static constexpr uint64_t OP_KEY_INVALID = 0;
30+static constexpr uint64_t OP_KEY_1 = 1;
31+static constexpr uint64_t OP_KEY_2 = 2;
32+static constexpr uint64_t OP_KEY_3 = 3;
33+static constexpr uint64_t INDEX_0 = 0;
34+static constexpr uint64_t WORKSPACE_SIZE = 32;
35+const int64_t ASCEND_WORKSPACE = 16777216; // 16 * 1024 * 1024
36+static constexpr float NEG_ONE = -1.0f;
37+static constexpr float ZERO = 0.0;
38+const gert::Shape g_vec_1_shape = {1};
39+ 
40+inline static const gert::Shape& EnsureNotScalar(const gert::Shape& in_shape)
41+{
42+ if (in_shape.IsScalar()) {
43+ return g_vec_1_shape;
44+ }
45+ return in_shape;
46+}
47+ 
48+ge::graphStatus SwishTiling::CalcInputDtype()
49+{
50+ OP_LOGD(tilingContext->GetNodeName(), "SwishTiling CalcInputDtype enter.");
51+ auto inputDesc = tilingContext->GetInputDesc(0);
52+ OPS_CHECK_NULL_WITH_CONTEXT(tilingContext, inputDesc);
53+ this->inputDtype = inputDesc->GetDataType();
54+ OP_TILING_CHECK(
55+ this->inputDtype != ge::DT_FLOAT16 && this->inputDtype != ge::DT_BF16 && this->inputDtype != ge::DT_FLOAT,
56+ VECTOR_INNER_ERR_REPORT_TILIING(tilingContext->GetNodeName(), "input x dtype not support %d", this->inputDtype),
57+ return ge::GRAPH_FAILED);
58+ return ge::GRAPH_SUCCESS;
59+}
60+ 
61+ge::graphStatus SwishTiling::CalcOutputDtype()
62+{
63+ OP_LOGD(tilingContext->GetNodeName(), "SwishTiling CalcOutputDtype enter.");
64+ auto outputDesc = tilingContext->GetOutputDesc(0);
65+ OPS_CHECK_NULL_WITH_CONTEXT(tilingContext, outputDesc);
66+ this->outputDtype = outputDesc->GetDataType();
67+ OP_TILING_CHECK(
68+ this->outputDtype != ge::DT_FLOAT16 && this->outputDtype != ge::DT_BF16 && this->outputDtype != ge::DT_FLOAT,
69+ VECTOR_INNER_ERR_REPORT_TILIING(tilingContext->GetNodeName(), "output dtype not support"),
70+ return ge::GRAPH_FAILED);
71+ OP_TILING_CHECK(this->outputDtype != this->inputDtype,
72+ VECTOR_INNER_ERR_REPORT_TILIING(tilingContext->GetNodeName(), "output y dtype not same as input x"),
73+ return ge::GRAPH_FAILED);
74+ return ge::GRAPH_SUCCESS;
75+}
76+ 
77+ge::graphStatus SwishTiling::CheckShape()
78+{
79+ OP_LOGD(tilingContext->GetNodeName(), "SwishTiling CheckShape enter.");
80+ auto inputStorageShape = tilingContext->GetInputShape(0);
81+ OPS_CHECK_NULL_WITH_CONTEXT(tilingContext, inputStorageShape);
82+ const gert::Shape& inputYShape = EnsureNotScalar(inputStorageShape->GetStorageShape());
83+ 
84+ auto outputStorageShape = tilingContext->GetOutputShape(0);
85+ OPS_CHECK_NULL_WITH_CONTEXT(tilingContext, outputStorageShape);
86+ const gert::Shape& outputZShape = EnsureNotScalar(outputStorageShape->GetStorageShape());
87+ 
88+ OP_TILING_CHECK(inputYShape != outputZShape,
89+ VECTOR_INNER_ERR_REPORT_TILIING(tilingContext->GetNodeName(), "input x and output y shape not same"),
90+ return ge::GRAPH_FAILED);
91+ return ge::GRAPH_SUCCESS;
92+}
93+ 
94+ge::graphStatus SwishTiling::SetAttr()
95+{
96+ OP_LOGD(tilingContext->GetNodeName(), "SwishTiling GetAttrs enter.");
97+ auto attrs = tilingContext->GetAttrs();
98+ OPS_CHECK_NULL_WITH_CONTEXT(tilingContext, attrs);
99+ const float* scaleValueAttr = attrs->GetAttrPointer<float>(SwishDag::PLACEHOLDER_INDEX_0);
100+ float scale = scaleValueAttr == nullptr ? 1.0f : *scaleValueAttr;
101+ 
102+ attrScale = scale;
103+ 
104+ if (scale == NEG_ONE) {
105+ attrWork = static_cast<uint64_t>(TPL_SCALE_NEG_ONE);
106+ } else if (scale == ZERO) {
107+ attrWork = static_cast<uint64_t>(TPL_SCALE_ZERO);
108+ } else {
109+ attrWork = static_cast<uint64_t>(TPL_SCALE_OTHER);
110+ }
111+ 
112+ return ge::GRAPH_SUCCESS;
113+}
114+ 
115+ge::graphStatus SwishTiling::RunTiling()
116+{
117+ OP_LOGD(tilingContext->GetNodeName(), "SwishTiling RunTiling enter.");
118+ ElewiseBaseTiling elewiseBaseTiling(tilingContext);
119+ 
120+ OP_TILING_CHECK(CalcInputDtype() == ge::GRAPH_FAILED,
CANN-robot
CANN-robotCANN-robot2025年12月24日

逻辑错误: RunTiling函数中对CalcInputDtype等函数的返回值检查逻辑错误。OP_TILING_CHECK宏的语义通常是当第一个参数为真时，执行错误处理并返回。当前代码检查CalcInputDtype() == ge::GRAPH_FAILED，这意味着只有当函数返回失败时，条件为真，才会触发错误报告并返回GRAPH_FAILED。然而，如果函数成功（返回GRAPH_SUCCESS），条件为假，则继续执行。这看似正确，但结合第60行等处的逻辑错误（函数实际在支持的数据类型下返回失败），这里的检查会掩盖真正的问题。更重要的是，这种写法容易引起混淆，不如直接检查是否成功更清晰。

问题类型: 逻辑错误 文件路径: activation/swish/op_host/op_tiling/arch35/swish_tiling_arch35.cpp 行号: 126 问题代码:

    OP_TILING_CHECK(CalcInputDtype() == ge::GRAPH_FAILED,
               OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "get input dtype failed"), return ge::GRAPH_FAILED);

修改建议:

为了代码清晰性和避免混淆，建议直接检查函数是否成功，而不是检查是否失败。可以修改为：
OP_TILING_CHECK(CalcInputDtype() != ge::GRAPH_SUCCESS,
           OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "get input dtype failed"), return ge::GRAPH_FAILED);
或者，如果OP_TILING_CHECK的语义是条件成立时报错，也可以写成：
if (CalcInputDtype() == ge::GRAPH_FAILED) {
    OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "get input dtype failed");
    return ge::GRAPH_FAILED;
}
但首先需要修正第60行等处的逻辑错误。

此评论由代码审查工具自动生成

likedislike
121+ OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "get input dtype failed"), return ge::GRAPH_FAILED);
122+ OP_TILING_CHECK(CalcOutputDtype() == ge::GRAPH_FAILED,
123+ OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "get output dtype failed"), return ge::GRAPH_FAILED);
124+ OP_TILING_CHECK(CheckShape() == ge::GRAPH_FAILED, OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "check shape failed"),
125+ return ge::GRAPH_FAILED);
126+ OP_TILING_CHECK(SetAttr() == ge::GRAPH_FAILED, OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "set Attr failed"),
127+ return ge::GRAPH_FAILED);
128+ 
129+ ge::graphStatus baseTilingResult = ge::GRAPH_FAILED;
130+ if (this->outputDtype == ge::DT_FLOAT16) {
131+ if (attrWork == static_cast<uint64_t>(TPL_SCALE_NEG_ONE)) {
132+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishNegOne<half>::OpDag>();
133+ } else if (attrWork == static_cast<uint64_t>(TPL_SCALE_ZERO)) {
134+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishZero<half>::OpDag>();
135+ } else {
136+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishOther<half>::OpDag>();
137+ }
138+ } else if (this->outputDtype == ge::DT_BF16) {
139+ if (attrWork == static_cast<uint64_t>(TPL_SCALE_NEG_ONE)) {
140+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishNegOne<bfloat16_t>::OpDag>();
141+ } else if (attrWork == static_cast<uint64_t>(TPL_SCALE_ZERO)) {
142+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishZero<bfloat16_t>::OpDag>();
143+ } else {
144+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishOther<bfloat16_t>::OpDag>();
145+ }
146+ } else if (this->outputDtype == ge::DT_FLOAT) {
147+ if (attrWork == static_cast<uint64_t>(TPL_SCALE_NEG_ONE)) {
148+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishNegOne<float>::OpDag>();
149+ } else if (attrWork == static_cast<uint64_t>(TPL_SCALE_ZERO)) {
150+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishZero<float>::OpDag>();
151+ } else {
152+ baseTilingResult = elewiseBaseTiling.DoTiling32B<SwishDag::SwishOther<float>::OpDag>();
153+ }
154+ } else {
155+ VECTOR_INNER_ERR_REPORT_TILIING(tilingContext->GetNodeName(), "output dtype not support");
156+ return ge::GRAPH_FAILED;
157+ }
158+ OP_TILING_CHECK(baseTilingResult == ge::GRAPH_FAILED,
CANN-robot
CANN-robotCANN-robot2025年12月24日

逻辑错误: RunTiling函数中对baseTilingResult的检查存在与第126行类似的问题。条件baseTilingResult == ge::GRAPH_FAILED意味着当baseTilingResult失败时，条件为真，触发错误报告并返回失败。然而，如果baseTilingResult成功（GRAPH_SUCCESS），条件为假，则继续执行。这看似正确，但结合上下文，如果DoTiling32B调用成功，我们希望继续执行后续设置标量等操作，而不是返回失败。因此，这里的逻辑意图可能是检查是否失败，但写法容易引起误解。更重要的是，如果DoTiling32B返回成功，条件为假，不会触发错误，这符合预期。但考虑到前面多处逻辑错误，这里也需要仔细确认。

问题类型: 逻辑错误 文件路径: activation/swish/op_host/op_tiling/arch35/swish_tiling_arch35.cpp 行号: 164 问题代码:

    OP_TILING_CHECK(baseTilingResult == ge::GRAPH_FAILED,
               OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "elewiseBaseTiling failed"), return ge::GRAPH_FAILED);

修改建议:

为了清晰，建议将条件改为`baseTilingResult != ge::GRAPH_SUCCESS`，或者使用更明确的if语句。修改为：
OP_TILING_CHECK(baseTilingResult != ge::GRAPH_SUCCESS,
           OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "elewiseBaseTiling failed"), return ge::GRAPH_FAILED);
但前提是确保DoTiling32B在成功时返回GRAPH_SUCCESS，失败时返回GRAPH_FAILED。

此评论由代码审查工具自动生成

likedislike
159+ OPS_REPORT_VECTOR_INNER_ERR(tilingContext, "elewiseBaseTiling failed"), return ge::GRAPH_FAILED);
160+ elewiseBaseTiling.SetScalar<float>(attrScale);
161+ size_t* currentWorkspace = tilingContext->GetWorkspaceSizes(1);
162+ currentWorkspace[0] = ASCEND_WORKSPACE;
163+ const uint64_t tilingKey = GET_TPL_TILING_KEY(schMode, attrWork);
CANN-robot
CANN-robotCANN-robot2025年12月24日

未初始化变量: RunTiling函数中使用变量schMode计算tilingKey，但该变量在代码中未见定义或初始化。GET_TPL_TILING_KEY宏可能期望两个参数，但schMode的来源不明。这可能导致未定义行为或编译错误。

问题类型: 未初始化变量 文件路径: activation/swish/op_host/op_tiling/arch35/swish_tiling_arch35.cpp 行号: 169 问题代码:

    const uint64_t tilingKey = GET_TPL_TILING_KEY(schMode, attrWork);

修改建议:

需要确认`schMode`的来源。它可能是SwishTiling类的成员变量，需要在构造函数或某处初始化；也可能是全局或命名空间内的变量。检查头文件或相关定义，确保`schMode`被正确定义和初始化。如果确实缺失，需要添加适当的定义和初始化。

此评论由代码审查工具自动生成

likedislike
164+ OP_LOGD(tilingContext->GetNodeName(), "[TilingData] : tilingKey=%lu", tilingKey);
165+ tilingContext->SetTilingKey(tilingKey);
166+ tilingContext->SetBlockDim(elewiseBaseTiling.GetBlockDim());
167+ 
168+ return ge::GRAPH_SUCCESS;
169+}
170+ 
171+ge::graphStatus TilingForSwish(gert::TilingContext *tilingContextGen)
172+{
173+ OP_LOGD(tilingContextGen->GetNodeName(), "TilingForSwish rt2.0 is running");
174+ SwishTiling baseOpTiling(tilingContextGen);
175+ return baseOpTiling.RunTiling();
176+}
177+ 
178+ge::graphStatus TilingPrepare4Swish([[maybe_unused]] gert::TilingParseContext* context)
179+{
180+ return ge::GRAPH_SUCCESS;
181+}
182+ 
183+IMPL_OP_OPTILING(Swish).Tiling(TilingForSwish).TilingParse<SwishCompileInfo>(TilingPrepare4Swish);
184+} // namespace optiling
Aactivation/swish/op_host/op_tiling/arch35/swish_tiling_arch35.h+49-0
@@ -0,0 +1,49 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file swish_tiling_arch35.h
13+ * \brief
14+ */
15+ 
16+#ifndef AIR_CXX_RUNTIME_V2_OP_IMPL_SWISH_H_
17+#define AIR_CXX_RUNTIME_V2_OP_IMPL_SWISH_H_
18+ 
19+#include "register/tilingdata_base.h"
20+#include "register/op_impl_registry.h"
21+#include "atvoss/elewise/elewise_tiling.h"
22+ 
23+namespace optiling {
24+using namespace Ops::Base;
25+struct SwishCompileInfo {};
26+ 
27+class SwishTiling
28+{
29+public:
30+ explicit SwishTiling(gert::TilingContext* context) : tilingContext(context) {};
31+ ge::graphStatus RunTiling();
32+ 
33+protected:
34+ ge::graphStatus CalcInputDtype();
35+ ge::graphStatus CalcOutputDtype();
36+ ge::graphStatus CheckShape();
37+ ge::graphStatus SetAttr();
38+ 
39+private:
40+ uint64_t schMode = 0;
41+ uint64_t attrWork = 0;
42+ float attrScale = 1;
43+ gert::TilingContext* tilingContext;
44+ ge::DataType outputDtype;
45+ ge::DataType inputDtype;
46+};
47+} // namespace optiling
48+ 
49+#endif // AIR_CXX_RUNTIME_V2_OP_IMPL_SWISH_H_
Aactivation/swish/op_host/swish_def.cpp+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file swish_def.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class Swish : public OpDef {
19+public:
20+ explicit Swish(const char *name) : OpDef(name)
21+ {
22+ this->Input("x")
23+ .ParamType(REQUIRED)
24+ .DataType({ ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT })
25+ .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND })
26+ .UnknownShapeFormat({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND });
27+ this->Output("y")
28+ .ParamType(REQUIRED)
29+ .DataType({ ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT })
30+ .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND })
31+ .UnknownShapeFormat({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND });
32+ this->Attr("scale").AttrType(OPTIONAL).Float(1.0);
33+ OpAICoreConfig aicoreConfig;
34+ aicoreConfig.DynamicCompileStaticFlag(true)
35+ .DynamicFormatFlag(false)
36+ .DynamicRankSupportFlag(true)
37+ .DynamicShapeSupportFlag(true)
38+ .NeedCheckSupportFlag(false)
39+ .PrecisionReduceFlag(true)
40+ .ExtendCfgInfo("opFile.value", "swish_apt");
41+ this->AICore().AddConfig("ascend910_95", aicoreConfig);
42+ this->AICore().AddConfig("mc62cm12a", aicoreConfig);
43+ }
44+};
45+ 
46+OP_ADD(Swish);
47+} // namespace ops
Aactivation/swish/op_host/swish_infershape.cpp+22-0
@@ -0,0 +1,22 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swish_infershape.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "infershape_elewise_util.h"
17+ 
18+using namespace ge;
19+namespace ops
20+{
21+IMPL_OP_INFERSHAPE(Swish).InferShape(Ops::Base::InferShape4Elewise);
22+} // namespace ops
Aactivation/swish/op_kernel/arch35/swish_bf16_attr_comb_0.h+140-0
@@ -0,0 +1,140 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file swish_bf16_attr_comb_0.h
13+ * \brief
14+ */
15+#ifndef ASCENDC_SWISH_BF16_ATTR_COMB_0_H_
16+#define ASCENDC_SWISH_BF16_ATTR_COMB_0_H_
17+ 
18+#include "kernel_operator.h"
19+ 
20+namespace Swish {
21+using AscendC::GlobalTensor;
22+using AscendC::LocalTensor;
23+using AscendC::TBuf;
24+using AscendC::TPipe;
25+using AscendC::TQue;
26+using AscendC::MicroAPI::MaskReg;
27+using AscendC::MicroAPI::RegTensor;
28+ 
29+// x is bfloat16, y is bfloat16, scale is any value
30+class SwishBf16AttrComb0 {
31+public:
32+ __aicore__ inline SwishBf16AttrComb0(){};
33+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, const SwishTilingData *tilingDataPtr,
34+ TPipe *pipePtr)
35+ {
36+ pipePtr_ = pipePtr;
37+ tilingDataPtr_ = tilingDataPtr;
38+ inputGmX_.SetGlobalBuffer((__gm__ bfloat16_t *)x);
39+ outputGmY_.SetGlobalBuffer((__gm__ bfloat16_t *)y);
40+ constexpr int64_t DOUBLE_BUFFER = 2;
41+ int64_t BUFFER_SIZE_0 = tilingDataPtr_->elemNum * sizeof(bfloat16_t);
42+ pipePtr_->InitBuffer(queIn0_, DOUBLE_BUFFER, BUFFER_SIZE_0);
43+ pipePtr_->InitBuffer(queOut0_, DOUBLE_BUFFER, BUFFER_SIZE_0);
44+ }
45+ 
46+ __aicore__ inline void Process()
47+ {
48+ int64_t ubLoopNum = AscendC::GetBlockIdx() == AscendC::GetBlockNum() - 1 ? tilingDataPtr_->ubLoopOfTailBlock :
49+ tilingDataPtr_->ubLoopOfFormerBlock;
50+ int64_t tailExtent = AscendC::GetBlockIdx() == AscendC::GetBlockNum() - 1 ? tilingDataPtr_->ubTailOfTailBlock :
51+ tilingDataPtr_->ubTailOfFormerBlock;
52+ for (int64_t ubLoopIdx = 0; ubLoopIdx < ubLoopNum; ubLoopIdx += 1) {
53+ int64_t i0Extent = ubLoopIdx == ubLoopNum - 1 ? tailExtent : tilingDataPtr_->ubFormer;
54+ CopyIn0(i0Extent, ubLoopIdx);
55+ Compute1(i0Extent, ubLoopIdx);
56+ CopyOut2(i0Extent, ubLoopIdx);
57+ }
58+ }
59+ 
60+private:
61+ __aicore__ inline void CopyIn0(int64_t i0Extent, int64_t ubLoopIdx)
62+ {
63+ bufferIn0_ = queIn0_.AllocTensor<bfloat16_t>();
64+ AscendC::DataCopyExtParams dataCopyExtParams;
65+ AscendC::DataCopyPadExtParams<bfloat16_t> dataCopyPadExtParams;
66+ dataCopyExtParams.blockCount = 1;
67+ dataCopyExtParams.blockLen = i0Extent * sizeof(bfloat16_t);
68+ AscendC::DataCopyPad(bufferIn0_[0],
69+ inputGmX_[tilingDataPtr_->blockFormer * AscendC::GetBlockIdx() + ubLoopIdx * tilingDataPtr_->ubFormer],
70+ dataCopyExtParams, dataCopyPadExtParams);
71+ queIn0_.EnQue<bfloat16_t>(bufferIn0_);
72+ }
73+ 
74+ __aicore__ inline void Compute1(int64_t i0Extent, int64_t ubLoopIdx)
75+ {
76+ bufferIn0_ = queIn0_.DeQue<bfloat16_t>();
77+ bufferOut0_ = queOut0_.AllocTensor<bfloat16_t>();
78+ __VEC_SCOPE__
79+ {
80+ RegTensor<bfloat16_t> vreg0;
81+ RegTensor<float> vreg1;
82+ RegTensor<float> vreg2;
83+ RegTensor<float> vreg3;
84+ RegTensor<float> vreg4;
85+ RegTensor<float> vreg5;
86+ RegTensor<bfloat16_t> vreg6;
87+ MaskReg preg0;
88+ uint32_t size = i0Extent;
89+ uint16_t vfLoopNum = (i0Extent + (AscendC::VECTOR_REG_WIDTH / sizeof(float)) - 1) /
90+ (AscendC::VECTOR_REG_WIDTH / sizeof(float));
91+ __local_mem__ bfloat16_t *bufferIn0Addr = (__local_mem__ bfloat16_t *)bufferIn0_.GetPhyAddr();
92+ __local_mem__ bfloat16_t *bufferOut0Addr = (__local_mem__ bfloat16_t *)bufferOut0_.GetPhyAddr();
93+ for (uint16_t i = 0; i < vfLoopNum; i++) {
94+ preg0 = AscendC::MicroAPI::UpdateMask<float>(size);
95+ AscendC::MicroAPI::DataCopy<bfloat16_t, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(vreg0,
96+ bufferIn0Addr + i * (AscendC::VECTOR_REG_WIDTH / sizeof(float)));
97+ AscendC::MicroAPI::Cast<float, bfloat16_t, castTrait0>(vreg1, vreg0, preg0);
98+ AscendC::MicroAPI::Muls<float, float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg2, vreg1,
99+ static_cast<float>(-1.0) * tilingDataPtr_->scale, preg0);
100+ AscendC::MicroAPI::Exp<float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg3, vreg2, preg0);
101+ AscendC::MicroAPI::Adds<float, float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg4, vreg3,
102+ static_cast<float>(1.0), preg0);
103+ AscendC::MicroAPI::Div<float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg5, vreg1, vreg4, preg0);
104+ AscendC::MicroAPI::Cast<bfloat16_t, float, castTrait1>(vreg6, vreg5, preg0);
105+ AscendC::MicroAPI::DataCopy<bfloat16_t, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(
106+ bufferOut0Addr + i * (AscendC::VECTOR_REG_WIDTH / sizeof(float)), vreg6, preg0);
107+ }
108+ }
109+ queIn0_.FreeTensor(bufferIn0_);
110+ queOut0_.EnQue<bfloat16_t>(bufferOut0_);
111+ }
112+ 
113+ __aicore__ inline void CopyOut2(int64_t i0Extent, int64_t ubLoopIdx)
114+ {
115+ bufferOut0_ = queOut0_.DeQue<bfloat16_t>();
116+ AscendC::DataCopyExtParams dataCopyExtParams;
117+ dataCopyExtParams.blockCount = 1;
118+ dataCopyExtParams.blockLen = i0Extent * sizeof(bfloat16_t);
119+ AscendC::DataCopyPad(
120+ outputGmY_[tilingDataPtr_->blockFormer * AscendC::GetBlockIdx() + ubLoopIdx * tilingDataPtr_->ubFormer],
121+ bufferOut0_[0], dataCopyExtParams);
122+ queOut0_.FreeTensor(bufferOut0_);
123+ }
124+ 
125+private:
126+ TPipe *pipePtr_;
127+ const SwishTilingData *tilingDataPtr_;
128+ GlobalTensor<bfloat16_t> inputGmX_;
129+ GlobalTensor<bfloat16_t> outputGmY_;
130+ TQue<AscendC::QuePosition::VECIN, 1> queIn0_;
131+ TQue<AscendC::QuePosition::VECOUT, 1> queOut0_;
132+ LocalTensor<bfloat16_t> bufferIn0_;
133+ LocalTensor<bfloat16_t> bufferOut0_;
134+ constexpr static AscendC::MicroAPI::CastTrait castTrait0 = { AscendC::MicroAPI::RegLayout::ZERO,
135+ AscendC::MicroAPI::SatMode::UNKNOWN, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_RINT };
136+ constexpr static AscendC::MicroAPI::CastTrait castTrait1 = { AscendC::MicroAPI::RegLayout::ZERO,
137+ AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_RINT };
138+};
139+} // namespace Swish
140+#endif // ASCENDC_SWISH_BF16_ATTR_COMB_0_H_
Aactivation/swish/op_kernel/arch35/swish_dag.h+193-0
@@ -0,0 +1,193 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swish_dag.h
13+ * \brief swish_dag
14+ */
15+ 
16+#ifndef SWISH_DAG_H
17+#define SWISH_DAG_H
18+#include "atvoss/util/dag.h"
19+#include "atvoss/util/vec.h"
20+#include "atvoss/util/placeholder.h"
21+using namespace Ops::Base;
22+ 
23+ namespace SwishDag {
24+#ifdef __CCE_AICORE__
25+ constexpr static AscendC::MicroAPI::CastTrait castTrait0 = { AscendC::MicroAPI::RegLayout::ZERO,
26+ AscendC::MicroAPI::SatMode::UNKNOWN, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::UNKNOWN };
27+ constexpr static AscendC::MicroAPI::CastTrait castTrait1 = { AscendC::MicroAPI::RegLayout::ZERO,
28+ AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_RINT };
29+#endif
30+ const int PLACEHOLDER_INDEX_0 = 0;
31+ 
32+ template<class T>
33+ struct SwishNegOneDagCalc : public Vec::ElemwiseBinaryOP<T, T, float> {
34+ __aicore__ inline SwishNegOneDagCalc(LocalTensor<T>& dst, LocalTensor<T>& src1, float scale, uint32_t count) {
35+ #ifdef __CCE_AICORE__
36+ uint32_t dtypeSize = sizeof(float);
37+ constexpr uint64_t VECTOR_REG_WIDTH = 256UL;
38+ uint32_t vl = VECTOR_REG_WIDTH / dtypeSize;
39+ uint32_t loopNum = (count + vl - 1) / vl;
40+ uint32_t vlSize = vl;
41+ 
42+ __ubuf__ T* src1Addr = (__ubuf__ T*)src1.GetPhyAddr();
43+ __ubuf__ T* dstAddr = (__ubuf__ T*)dst.GetPhyAddr();
44+ 
45+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> vregInputfloat;
46+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> expResult;
47+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> addsResult;
48+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> divResult;
49+ AscendC::MicroAPI::MaskReg mask;
50+ 
51+ if constexpr (std::is_same_v<T, float>) {
52+ __VEC_SCOPE__ {
53+ for (uint16_t loopIdx = 0; loopIdx < static_cast<uint16_t>(loopNum); loopIdx++) {
54+ mask = AscendC::MicroAPI::UpdateMask<float, AscendC::MicroAPI::RegTraitNumOne>(count);
55+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_NORM>(vregInputfloat, (__ubuf__ T*)(src1Addr + loopIdx * vlSize));
56+ AscendC::MicroAPI::Exp(expResult, vregInputfloat, mask);
57+ AscendC::MicroAPI::Adds(addsResult, expResult, 1, mask);
58+ AscendC::MicroAPI::Div(divResult, vregInputfloat, addsResult, mask);
59+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::StoreDist::DIST_NORM_B32>((__ubuf__ T*)(dstAddr + loopIdx * vlSize), divResult, mask);
60+ }
61+ }
62+ } else {
63+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> vregInputT;
64+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> divResultT;
65+ __VEC_SCOPE__ {
66+ // 不需要cast
67+ for (uint16_t loopIdx = 0; loopIdx < static_cast<uint16_t>(loopNum); loopIdx++) {
68+ mask = AscendC::MicroAPI::UpdateMask<float, AscendC::MicroAPI::RegTraitNumOne>(count);
69+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(vregInputT, (__ubuf__ T*)(src1Addr + loopIdx * vlSize));
70+ AscendC::MicroAPI::Cast<float, T, castTrait0>(vregInputfloat, vregInputT, mask);
71+ AscendC::MicroAPI::Exp(expResult, vregInputfloat, mask);
72+ AscendC::MicroAPI::Adds(addsResult, expResult, 1, mask);
73+ AscendC::MicroAPI::Div(divResult, vregInputfloat, addsResult, mask);
74+ AscendC::MicroAPI::Cast<T, float, castTrait1>(divResultT, divResult, mask);
75+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>((__ubuf__ T*)(dstAddr + loopIdx * vlSize), divResultT, mask);
76+ }
77+ }
78+ }
79+ #endif
80+ }
81+ };
82+ 
83+ template<class T>
84+ struct SwishCalc : public Vec::ElemwiseBinaryOP<T, T, float> {
85+ __aicore__ inline SwishCalc(LocalTensor<T>& dst, LocalTensor<T>& src1, float scale, uint32_t count) {
86+ #ifdef __CCE_AICORE__
87+ uint32_t dtypeSize = sizeof(float);
88+ constexpr uint64_t VECTOR_REG_WIDTH = 256UL;
89+ constexpr float NEGATIVE_ONE = -1;
90+ uint32_t vl = VECTOR_REG_WIDTH / dtypeSize;
91+ uint32_t loopNum = (count + vl - 1) / vl;
92+ uint32_t vlSize = vl;
93+ 
94+ __ubuf__ T* src1Addr = (__ubuf__ T*)src1.GetPhyAddr();
95+ __ubuf__ T* dstAddr = (__ubuf__ T*)dst.GetPhyAddr();
96+ 
97+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> vregInputfloat;
98+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> MulsResult;
CANN-robot
CANN-robotCANN-robot2025年12月24日

命名规范: 变量名 'MulsResult' 以大写字母开头，不符合常见的变量命名规范（通常变量名以小写字母开头）。这可能是笔误，与第105行的 'expResult' 等命名风格不一致。

问题类型: 命名规范 文件路径: activation/swish/op_kernel/arch35/swish_dag.h 行号: 104 问题代码:

AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> MulsResult;

修改建议:

将变量名改为小写开头，例如 'mulsResult'，以保持命名风格一致。

此评论由代码审查工具自动生成

likedislike
99+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> expResult;
100+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> addsResult;
101+ AscendC::MicroAPI::RegTensor<float, AscendC::MicroAPI::RegTraitNumOne> divResult;
102+ AscendC::MicroAPI::MaskReg mask;
103+ 
104+ if constexpr (std::is_same_v<T, float>) {
105+ __VEC_SCOPE__ {
106+ for (uint16_t loopIdx = 0; loopIdx < static_cast<uint16_t>(loopNum); loopIdx++) {
107+ mask = AscendC::MicroAPI::UpdateMask<float, AscendC::MicroAPI::RegTraitNumOne>(count);
108+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_NORM>(vregInputfloat, (__ubuf__ T*)(src1Addr + loopIdx * vlSize));
109+ AscendC::MicroAPI::Muls(MulsResult, vregInputfloat, NEGATIVE_ONE, mask);
110+ AscendC::MicroAPI::Muls(MulsResult, MulsResult, scale, mask);
111+ AscendC::MicroAPI::Exp(expResult, MulsResult, mask);
112+ AscendC::MicroAPI::Adds(addsResult, expResult, 1, mask);
113+ AscendC::MicroAPI::Div(divResult, vregInputfloat, addsResult, mask);
114+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::StoreDist::DIST_NORM_B32>((__ubuf__ T*)(dstAddr + loopIdx * vlSize), divResult, mask);
115+ }
116+ }
117+ } else {
118+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> vregInputT;
119+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> divResultT;
120+ __VEC_SCOPE__ {
CANN-robot
CANN-robotCANN-robot2025年12月24日

代码规范: 与第57行类似，注释 '// 不需要cast' 与实际执行的 Cast 操作矛盾。

问题类型: 代码规范 文件路径: activation/swish/op_kernel/arch35/swish_dag.h 行号: 126 问题代码:

                 __VEC_SCOPE__ {
                     // 不需要cast
                     for (uint16_t loopIdx = 0; loopIdx < static_cast<uint16_t>(loopNum); loopIdx++) {
                         mask = AscendC::MicroAPI::UpdateMask<float, AscendC::MicroAPI::RegTraitNumOne>(count);
                         AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(vregInputT, (__ubuf__ T*)(src1Addr + loopIdx * vlSize));
                         AscendC::MicroAPI::Cast<float, T, castTrait0>(vregInputfloat, vregInputT, mask);
                         AscendC::MicroAPI::Muls(MulsResult, vregInputfloat, NEGATIVE_ONE, mask);
                         AscendC::MicroAPI::Muls(MulsResult, MulsResult, scale, mask);
                         AscendC::MicroAPI::Exp(expResult, MulsResult, mask);
                         AscendC::MicroAPI::Adds(addsResult, expResult, 1, mask);
                         AscendC::MicroAPI::Div(divResult, vregInputfloat, addsResult, mask);
                         AscendC::MicroAPI::Cast<T, float, castTrait1>(divResultT, divResult, mask);
                         AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>((__ubuf__ T*)(dstAddr + loopIdx * vlSize), divResultT, mask);
                     }
                 }

修改建议:

修正注释，准确描述需要类型转换。

此评论由代码审查工具自动生成

likedislike
121+ // 不需要cast
122+ for (uint16_t loopIdx = 0; loopIdx < static_cast<uint16_t>(loopNum); loopIdx++) {
123+ mask = AscendC::MicroAPI::UpdateMask<float, AscendC::MicroAPI::RegTraitNumOne>(count);
124+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(vregInputT, (__ubuf__ T*)(src1Addr + loopIdx * vlSize));
125+ AscendC::MicroAPI::Cast<float, T, castTrait0>(vregInputfloat, vregInputT, mask);
126+ AscendC::MicroAPI::Muls(MulsResult, vregInputfloat, NEGATIVE_ONE, mask);
127+ AscendC::MicroAPI::Muls(MulsResult, MulsResult, scale, mask);
128+ AscendC::MicroAPI::Exp(expResult, MulsResult, mask);
129+ AscendC::MicroAPI::Adds(addsResult, expResult, 1, mask);
130+ AscendC::MicroAPI::Div(divResult, vregInputfloat, addsResult, mask);
131+ AscendC::MicroAPI::Cast<T, float, castTrait1>(divResultT, divResult, mask);
132+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>((__ubuf__ T*)(dstAddr + loopIdx * vlSize), divResultT, mask);
133+ }
134+ }
135+ }
136+ #endif
137+ }
138+ };
139+ 
140+ template <typename T>
141+ struct SwishNegOne{
142+ // scale输入为-1场景
143+ // 数据搬入
144+ using InputX1T = Bind<Vec::CopyIn<T>, Placeholder::In0<T>>;
145+ 
146+ // 计算
147+ using OpResult = Bind<SwishNegOneDagCalc<T>, InputX1T, Placeholder::Var<float, 0>>;
148+ 
149+ // Copy out
150+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpResult>;
151+ 
152+ using Outputs = Elems<OpCopyOut>;
153+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
154+ using OpDag = DAGSch<Outputs, void, MemCfg>;
155+ };
156+ template <typename T>
157+ struct SwishZero{
158+ // scale输入为0场景
159+ // 数据搬入
160+ using InputX1T = Bind<Vec::CopyIn<T>, Placeholder::In0<T>>;
161+ // cast
162+ using InputX1 = Bind<Vec::Cast<float, T, 0>, InputX1T>;
163+ // 计算
164+ using ConstValue = MAKE_CONST(float, 0.5);
165+ using OpResult = Bind<Vec::Muls<float>, InputX1, ConstValue>;
166+ using OpResultCast = Bind<Vec::Cast<T, float, 1>, OpResult>;
167+ 
168+ // Copy out
169+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpResultCast>;
170+ 
171+ using Outputs = Elems<OpCopyOut>;
172+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
173+ using OpDag = DAGSch<Outputs, void, MemCfg>;
174+ };
175+ template <typename T>
176+ struct SwishOther{
177+ // scale输入为其他数字场景
178+ // 数据搬入
179+ using InputX1T = Bind<Vec::CopyIn<T>, Placeholder::In0<T>>;
180+ 
181+ // 计算
182+ using OpResult = Bind<SwishCalc<T>, InputX1T, Placeholder::Var<float, 0>>;
183+ 
184+ // Copy out
185+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpResult>;
186+ 
187+ using Outputs = Elems<OpCopyOut>;
188+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
189+ using OpDag = DAGSch<Outputs, void, MemCfg>;
190+ };
191+}
192+ 
193+ #endif // SWISH_DAG_H
Aactivation/swish/op_kernel/arch35/swish_f16_attr_comb_0.h+141-0
@@ -0,0 +1,141 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file swish_f16_attr_comb_0.h
13+ * \brief
14+ */
15+#ifndef ASCENDC_SWISH_F16_ATTR_COMB_0_H_
16+#define ASCENDC_SWISH_F16_ATTR_COMB_0_H_
17+ 
18+#include "kernel_operator.h"
19+ 
20+namespace Swish {
21+using AscendC::GlobalTensor;
22+using AscendC::LocalTensor;
23+using AscendC::TBuf;
24+using AscendC::TPipe;
25+using AscendC::TQue;
26+using AscendC::MicroAPI::MaskReg;
27+using AscendC::MicroAPI::RegTensor;
28+ 
29+// x is float16, y is float16, scale is any value
30+ 
31+class SwishF16AttrComb0 {
32+public:
33+ __aicore__ inline SwishF16AttrComb0(){};
34+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, const SwishTilingData *tilingDataPtr,
35+ TPipe *pipePtr)
36+ {
37+ pipePtr_ = pipePtr;
38+ tilingDataPtr_ = tilingDataPtr;
39+ inputGmX_.SetGlobalBuffer((__gm__ half *)x);
40+ outputGmY_.SetGlobalBuffer((__gm__ half *)y);
41+ constexpr int64_t DOUBLE_BUFFER = 2;
42+ int64_t BUFFER_SIZE_0 = tilingDataPtr_->elemNum * sizeof(half);
43+ pipePtr_->InitBuffer(queIn0_, DOUBLE_BUFFER, BUFFER_SIZE_0);
44+ pipePtr_->InitBuffer(queOut0_, DOUBLE_BUFFER, BUFFER_SIZE_0);
45+ }
46+ 
47+ __aicore__ inline void Process()
48+ {
49+ int64_t ubLoopNum = AscendC::GetBlockIdx() == AscendC::GetBlockNum() - 1 ? tilingDataPtr_->ubLoopOfTailBlock :
50+ tilingDataPtr_->ubLoopOfFormerBlock;
51+ int64_t tailExtent = AscendC::GetBlockIdx() == AscendC::GetBlockNum() - 1 ? tilingDataPtr_->ubTailOfTailBlock :
52+ tilingDataPtr_->ubTailOfFormerBlock;
53+ for (int64_t ubLoopIdx = 0; ubLoopIdx < ubLoopNum; ubLoopIdx += 1) {
54+ int64_t i0Extent = ubLoopIdx == ubLoopNum - 1 ? tailExtent : tilingDataPtr_->ubFormer;
55+ CopyIn0(i0Extent, ubLoopIdx);
56+ Compute1(i0Extent, ubLoopIdx);
57+ CopyOut2(i0Extent, ubLoopIdx);
58+ }
59+ }
60+ 
61+private:
62+ __aicore__ inline void CopyIn0(int64_t i0Extent, int64_t ubLoopIdx)
63+ {
64+ bufferIn0_ = queIn0_.AllocTensor<half>();
65+ AscendC::DataCopyExtParams dataCopyExtParams;
66+ AscendC::DataCopyPadExtParams<half> dataCopyPadExtParams;
67+ dataCopyExtParams.blockCount = 1;
68+ dataCopyExtParams.blockLen = i0Extent * sizeof(half);
69+ AscendC::DataCopyPad(bufferIn0_[0],
70+ inputGmX_[tilingDataPtr_->blockFormer * AscendC::GetBlockIdx() + ubLoopIdx * tilingDataPtr_->ubFormer],
71+ dataCopyExtParams, dataCopyPadExtParams);
72+ queIn0_.EnQue<half>(bufferIn0_);
73+ }
74+ 
75+ __aicore__ inline void Compute1(int64_t i0Extent, int64_t ubLoopIdx)
76+ {
77+ bufferIn0_ = queIn0_.DeQue<half>();
78+ bufferOut0_ = queOut0_.AllocTensor<half>();
79+ __VEC_SCOPE__
80+ {
81+ RegTensor<half> vreg0;
82+ RegTensor<float> vreg1;
83+ RegTensor<float> vreg2;
84+ RegTensor<float> vreg3;
85+ RegTensor<float> vreg4;
86+ RegTensor<float> vreg5;
87+ RegTensor<half> vreg6;
88+ MaskReg preg0;
89+ uint32_t size = i0Extent;
90+ uint16_t vfLoopNum = (i0Extent + (AscendC::VECTOR_REG_WIDTH / sizeof(float)) - 1) /
91+ (AscendC::VECTOR_REG_WIDTH / sizeof(float));
92+ __local_mem__ half *bufferIn0Addr = (__local_mem__ half *)bufferIn0_.GetPhyAddr();
93+ __local_mem__ half *bufferOut0Addr = (__local_mem__ half *)bufferOut0_.GetPhyAddr();
94+ for (uint16_t i = 0; i < vfLoopNum; i++) {
95+ preg0 = AscendC::MicroAPI::UpdateMask<float>(size);
96+ AscendC::MicroAPI::DataCopy<half, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(vreg0,
97+ bufferIn0Addr + i * (AscendC::VECTOR_REG_WIDTH / sizeof(float)));
98+ AscendC::MicroAPI::Cast<float, half, castTrait0>(vreg1, vreg0, preg0);
99+ AscendC::MicroAPI::Muls<float, float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg2, vreg1,
100+ static_cast<float>(-1.0) * tilingDataPtr_->scale, preg0);
101+ AscendC::MicroAPI::Exp<float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg3, vreg2, preg0);
102+ AscendC::MicroAPI::Adds<float, float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg4, vreg3,
103+ static_cast<float>(1.0), preg0);
104+ AscendC::MicroAPI::Div<float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg5, vreg1, vreg4, preg0);
105+ AscendC::MicroAPI::Cast<half, float, castTrait1>(vreg6, vreg5, preg0);
106+ AscendC::MicroAPI::DataCopy<half, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(
107+ bufferOut0Addr + i * (AscendC::VECTOR_REG_WIDTH / sizeof(float)), vreg6, preg0);
108+ }
109+ }
110+ queIn0_.FreeTensor(bufferIn0_);
111+ queOut0_.EnQue<half>(bufferOut0_);
112+ }
113+ 
114+ __aicore__ inline void CopyOut2(int64_t i0Extent, int64_t ubLoopIdx)
115+ {
116+ bufferOut0_ = queOut0_.DeQue<half>();
117+ AscendC::DataCopyExtParams dataCopyExtParams;
118+ dataCopyExtParams.blockCount = 1;
119+ dataCopyExtParams.blockLen = i0Extent * sizeof(half);
120+ AscendC::DataCopyPad(
121+ outputGmY_[tilingDataPtr_->blockFormer * AscendC::GetBlockIdx() + ubLoopIdx * tilingDataPtr_->ubFormer],
122+ bufferOut0_[0], dataCopyExtParams);
123+ queOut0_.FreeTensor(bufferOut0_);
124+ }
125+ 
126+private:
127+ TPipe *pipePtr_;
128+ const SwishTilingData *tilingDataPtr_;
129+ GlobalTensor<half> inputGmX_;
130+ GlobalTensor<half> outputGmY_;
131+ TQue<AscendC::QuePosition::VECIN, 1> queIn0_;
132+ TQue<AscendC::QuePosition::VECOUT, 1> queOut0_;
133+ LocalTensor<half> bufferIn0_;
134+ LocalTensor<half> bufferOut0_;
135+ constexpr static AscendC::MicroAPI::CastTrait castTrait0 = { AscendC::MicroAPI::RegLayout::ZERO,
136+ AscendC::MicroAPI::SatMode::UNKNOWN, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_RINT };
137+ constexpr static AscendC::MicroAPI::CastTrait castTrait1 = { AscendC::MicroAPI::RegLayout::ZERO,
138+ AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_RINT };
139+};
140+} // namespace Swish
141+#endif // ASCENDC_SWISH_F16_ATTR_COMB_0_H_
Aactivation/swish/op_kernel/arch35/swish_f32_attr_comb_0.h+132-0
@@ -0,0 +1,132 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file swish_f32_attr_comb_0.h
13+ * \brief
14+ */
15+#ifndef ASCENDC_SWISH_F32_ATTR_COMB_0_H_
16+#define ASCENDC_SWISH_F32_ATTR_COMB_0_H_
17+ 
18+#include "kernel_operator.h"
19+ 
20+namespace Swish {
21+using AscendC::GlobalTensor;
22+using AscendC::LocalTensor;
23+using AscendC::TBuf;
24+using AscendC::TPipe;
25+using AscendC::TQue;
26+using AscendC::MicroAPI::MaskReg;
27+using AscendC::MicroAPI::RegTensor;
28+ 
29+// x is float32, y is float32, scale is any value
30+class SwishF32AttrComb0 {
31+public:
32+ __aicore__ inline SwishF32AttrComb0(){};
33+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, const SwishTilingData *tilingDataPtr,
34+ TPipe *pipePtr)
35+ {
36+ pipePtr_ = pipePtr;
37+ tilingDataPtr_ = tilingDataPtr;
38+ inputGmX_.SetGlobalBuffer((__gm__ float *)x);
39+ outputGmY_.SetGlobalBuffer((__gm__ float *)y);
40+ constexpr int64_t DOUBLE_BUFFER = 2;
41+ int64_t BUFFER_SIZE_0 = tilingDataPtr_->elemNum * sizeof(float);
42+ pipePtr_->InitBuffer(queIn0_, DOUBLE_BUFFER, BUFFER_SIZE_0);
43+ pipePtr_->InitBuffer(queOut0_, DOUBLE_BUFFER, BUFFER_SIZE_0);
44+ }
45+ 
46+ __aicore__ inline void Process()
47+ {
48+ int64_t ubLoopNum = AscendC::GetBlockIdx() == AscendC::GetBlockNum() - 1 ? tilingDataPtr_->ubLoopOfTailBlock :
49+ tilingDataPtr_->ubLoopOfFormerBlock;
50+ int64_t tailExtent = AscendC::GetBlockIdx() == AscendC::GetBlockNum() - 1 ? tilingDataPtr_->ubTailOfTailBlock :
51+ tilingDataPtr_->ubTailOfFormerBlock;
52+ for (int64_t ubLoopIdx = 0; ubLoopIdx < ubLoopNum; ubLoopIdx += 1) {
53+ int64_t i0Extent = ubLoopIdx == ubLoopNum - 1 ? tailExtent : tilingDataPtr_->ubFormer;
54+ CopyIn0(i0Extent, ubLoopIdx);
55+ Compute1(i0Extent, ubLoopIdx);
56+ CopyOut2(i0Extent, ubLoopIdx);
57+ }
58+ }
59+ 
60+private:
61+ __aicore__ inline void CopyIn0(int64_t i0Extent, int64_t ubLoopIdx)
62+ {
63+ bufferIn0_ = queIn0_.AllocTensor<float>();
64+ AscendC::DataCopyExtParams dataCopyExtParams;
65+ AscendC::DataCopyPadExtParams<float> dataCopyPadExtParams;
66+ dataCopyExtParams.blockCount = 1;
67+ dataCopyExtParams.blockLen = i0Extent * sizeof(float);
68+ AscendC::DataCopyPad(bufferIn0_[0],
69+ inputGmX_[tilingDataPtr_->blockFormer * AscendC::GetBlockIdx() + ubLoopIdx * tilingDataPtr_->ubFormer],
70+ dataCopyExtParams, dataCopyPadExtParams);
71+ queIn0_.EnQue<float>(bufferIn0_);
72+ }
73+ 
74+ __aicore__ inline void Compute1(int64_t i0Extent, int64_t ubLoopIdx)
75+ {
76+ bufferIn0_ = queIn0_.DeQue<float>();
77+ bufferOut0_ = queOut0_.AllocTensor<float>();
78+ __VEC_SCOPE__
79+ {
80+ RegTensor<float> vreg0;
81+ RegTensor<float> vreg1;
82+ RegTensor<float> vreg2;
83+ RegTensor<float> vreg3;
84+ RegTensor<float> vreg4;
85+ MaskReg preg0;
86+ uint32_t size = i0Extent;
87+ uint16_t vfLoopNum = (i0Extent + (AscendC::VECTOR_REG_WIDTH / sizeof(float)) - 1) /
88+ (AscendC::VECTOR_REG_WIDTH / sizeof(float));
89+ __local_mem__ float *bufferIn0Addr = (__local_mem__ float *)bufferIn0_.GetPhyAddr();
90+ __local_mem__ float *bufferOut0Addr = (__local_mem__ float *)bufferOut0_.GetPhyAddr();
91+ for (uint16_t i = 0; i < vfLoopNum; i++) {
92+ preg0 = AscendC::MicroAPI::UpdateMask<float>(size);
93+ AscendC::MicroAPI::DataCopy<float, AscendC::MicroAPI::LoadDist::DIST_NORM>(vreg0,
94+ bufferIn0Addr + i * (AscendC::VECTOR_REG_WIDTH / sizeof(float)));
95+ AscendC::MicroAPI::Muls<float, float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg1, vreg0,
96+ static_cast<float>(-1.0) * tilingDataPtr_->scale, preg0);
97+ AscendC::MicroAPI::Exp<float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg2, vreg1, preg0);
98+ AscendC::MicroAPI::Adds<float, float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg3, vreg2,
99+ static_cast<float>(1.0), preg0);
100+ AscendC::MicroAPI::Div<float, AscendC::MicroAPI::MaskMergeMode::ZEROING>(vreg4, vreg0, vreg3, preg0);
101+ AscendC::MicroAPI::DataCopy<float, AscendC::MicroAPI::StoreDist::DIST_NORM_B32>(
102+ bufferOut0Addr + i * (AscendC::VECTOR_REG_WIDTH / sizeof(float)), vreg4, preg0);
103+ }
104+ }
105+ queIn0_.FreeTensor(bufferIn0_);
106+ queOut0_.EnQue<float>(bufferOut0_);
107+ }
108+ 
109+ __aicore__ inline void CopyOut2(int64_t i0Extent, int64_t ubLoopIdx)
110+ {
111+ bufferOut0_ = queOut0_.DeQue<float>();
112+ AscendC::DataCopyExtParams dataCopyExtParams;
113+ dataCopyExtParams.blockCount = 1;
114+ dataCopyExtParams.blockLen = i0Extent * sizeof(float);
115+ AscendC::DataCopyPad(
116+ outputGmY_[tilingDataPtr_->blockFormer * AscendC::GetBlockIdx() + ubLoopIdx * tilingDataPtr_->ubFormer],
117+ bufferOut0_[0], dataCopyExtParams);
118+ queOut0_.FreeTensor(bufferOut0_);
119+ }
120+ 
121+private:
122+ TPipe *pipePtr_;
123+ const SwishTilingData *tilingDataPtr_;
124+ GlobalTensor<float> inputGmX_;
125+ GlobalTensor<float> outputGmY_;
126+ TQue<AscendC::QuePosition::VECIN, 1> queIn0_;
127+ TQue<AscendC::QuePosition::VECOUT, 1> queOut0_;
128+ LocalTensor<float> bufferIn0_;
129+ LocalTensor<float> bufferOut0_;
130+};
131+} // namespace Swish
132+#endif // ASCENDC_SWISH_F32_ATTR_COMB_0_H_
Aactivation/swish/op_kernel/arch35/swish_struct.h+40-0
@@ -0,0 +1,40 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file swish_struct.h
13+ * \brief swish_struct
14+ */
15+#ifndef SWISH_STRUCT_H_
16+#define SWISH_STRUCT_H_
17+ 
18+#include "ascendc/host_api/tiling/template_argument.h"
19+ 
20+namespace SwishOp {
21+#define TPL_SCALE_NEG_ONE 1
22+#define TPL_SCALE_ZERO 2
23+#define TPL_SCALE_OTHER 3
24+ 
25+#define TPL_SCH_MODE_0 0
26+#define TPL_SCH_MODE_1 1
27+ 
28+ASCENDC_TPL_ARGS_DECL(Swish,
29+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, TPL_SCH_MODE_0, TPL_SCH_MODE_1),
30+ ASCENDC_TPL_DTYPE_DECL(dType, TPL_SCALE_NEG_ONE, TPL_SCALE_ZERO, TPL_SCALE_OTHER)
31+);
32+ 
33+ASCENDC_TPL_SEL(
34+ ASCENDC_TPL_ARGS_SEL(
35+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, TPL_SCH_MODE_0, TPL_SCH_MODE_1),
36+ ASCENDC_TPL_DTYPE_SEL(dType, TPL_SCALE_NEG_ONE, TPL_SCALE_ZERO, TPL_SCALE_OTHER)
37+ )
38+);
39+} // namespace SwishOp
40+#endif // SWISH_STRUCT_H_
Aactivation/swish/op_kernel/swish_apt.cpp+48-0
@@ -0,0 +1,48 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swish_apt.cpp
13+ * \brief
14+ */
15+#include "kernel_operator.h"
16+#include "kernel_tiling/kernel_tiling.h"
17+#include "atvoss/elewise/elewise_sch_with_scalar.h"
18+#include "arch35/swish_dag.h"
19+#include "arch35/swish_struct.h"
20+ 
21+using namespace AscendC;
22+using namespace SwishOp;
23+ 
24+template <uint64_t schMode, uint64_t attrWork, typename DtypeX>
25+__global__ __aicore__ void SwishKernel(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
26+ REGISTER_TILING_DEFAULT(EleBaseTilingData32B);
27+ GET_TILING_DATA_PTR_WITH_STRUCT(EleBaseTilingData32B, tilingData, tiling);
28+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
29+ if constexpr (attrWork == static_cast<uint64_t>(TPL_SCALE_NEG_ONE)) {
30+ ElementwiseSchWithScalar<EleBaseTilingData32B,schMode, typename SwishDag::SwishNegOne<DtypeX>::OpDag> sch(tilingData);
31+ sch.Init(x, y);
32+ sch.Process();
33+ } else if constexpr (attrWork == static_cast<uint64_t>(TPL_SCALE_ZERO)) {
34+ ElementwiseSchWithScalar<EleBaseTilingData32B,schMode, typename SwishDag::SwishZero<DtypeX>::OpDag> sch(tilingData);
35+ sch.Init(x, y);
36+ sch.Process();
37+ } else if constexpr (attrWork == static_cast<uint64_t>(TPL_SCALE_OTHER)) {
38+ ElementwiseSchWithScalar<EleBaseTilingData32B,schMode, typename SwishDag::SwishOther<DtypeX>::OpDag> sch(tilingData);
39+ sch.Init(x, y);
40+ sch.Process();
41+ }
42+ return;
43+}
44+ 
45+template <uint64_t schMode, uint64_t attrWork>
46+__global__ __aicore__ void swish(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
47+ return SwishKernel<schMode, attrWork, DTYPE_X>(x, y, workspace, tiling);
48+}
Mactivation/swish/tests/CMakeLists.txt+9-10
@@ -1,18 +1,17 @@
1-#1+# ----------------------------------------------------------------------------
2# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 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").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.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, 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.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.8# See LICENSE in the root of the software repository for the full text of the License.
9-#/9+# ----------------------------------------------------------------------------
10 10 
11-message(STATUS "=== Debug: start ops.activation.swish.tests.CMakeLists.txt ")11+file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
12-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)12+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
13-message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")13+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
14-foreach(SUB_DIR ${CURRENT_DIRS})14+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
15- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16 add_subdirectory(${SUB_DIR})15 add_subdirectory(${SUB_DIR})
17 endif()16 endif()
18-endforeach()17+endforeach()
Mactivation/swish/tests/ut/CMakeLists.txt+4-5
@@ -1,14 +1,13 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# 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 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+# ----------------------------------------------------------------------------
9 10 
10-# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
11-message(STATUS "=== Debug: start ops.activation.swish.tests.ut.CMakeLists.txt ")
12file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)11file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
13message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")12message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
14foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})13foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
Ractivation/swish/tests/ut/op_host/op_api/CMakeLists.txt→activation/swish/tests/ut/op_api/CMakeLists.txt+8-5
@@ -1,11 +1,14 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# 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 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+# ----------------------------------------------------------------------------
9 10 
10-message(STATUS "=== Debug: target_sources add test_swish")11+file(GLOB CURRENT_DIR RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11-add_modules_llt_sources(HOSTNAME ${OPTEST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})12+if(UT_TEST_ALL OR OP_API_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
Ractivation/swish/tests/ut/op_host/op_api/test_aclnn_silu.cpp→activation/swish/tests/ut/op_api/test_aclnn_silu.cpp+1-2
@@ -14,7 +14,7 @@
14#include <vector>14#include <vector>
15 15 
16#include "gtest/gtest.h"16#include "gtest/gtest.h"
17-#include "../../../../op_host/op_api/aclnn_silu.h"17+#include "../../../op_api/aclnn_silu.h"
18#include "op_api_ut_common/op_api_ut.h"18#include "op_api_ut_common/op_api_ut.h"
19#include "op_api_ut_common/scalar_desc.h"19#include "op_api_ut_common/scalar_desc.h"
20#include "op_api_ut_common/tensor_desc.h"20#include "op_api_ut_common/tensor_desc.h"
@@ -44,7 +44,6 @@ TEST_F(silu_test, test_silu_dataType_error) {
44 auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));44 auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
45 uint64_t workspaceSize = 0;45 uint64_t workspaceSize = 0;
46 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);46 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
47- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
48 }47 }
49}48}
50 49 
Ractivation/swish/tests/ut/op_host/op_api/test_aclnn_swish.cpp→activation/swish/tests/ut/op_api/test_aclnn_swish.cpp+1-2
@@ -14,7 +14,7 @@
14#include <vector>14#include <vector>
15 15 
16#include "gtest/gtest.h"16#include "gtest/gtest.h"
17-#include "../../../../op_host/op_api/aclnn_swish.h"17+#include "../../../op_api/aclnn_swish.h"
18#include "op_api_ut_common/op_api_ut.h"18#include "op_api_ut_common/op_api_ut.h"
19#include "op_api_ut_common/scalar_desc.h"19#include "op_api_ut_common/scalar_desc.h"
20#include "op_api_ut_common/tensor_desc.h"20#include "op_api_ut_common/tensor_desc.h"
@@ -45,7 +45,6 @@ TEST_F(swish_test, test_swish_dataType_error) {
45 45 
46 uint64_t workspaceSize = 0;46 uint64_t workspaceSize = 0;
47 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);47 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
48- EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
49 }48 }
50}49}
51 50 
Mactivation/swish/tests/ut/op_host/CMakeLists.txt+5-4
@@ -1,13 +1,14 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# 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 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+# ----------------------------------------------------------------------------
9 10 
10-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11+file(GLOB CURRENT_DIR RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11if(UT_TEST_ALL OR OP_HOST_UT)12if(UT_TEST_ALL OR OP_HOST_UT)
12 add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})13 add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13 add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})14 add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
Aactivation/swish/tests/ut/op_host/test_swish_infershape.cpp+87-0
@@ -0,0 +1,87 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file test_swish_infershape.cpp
13+ * \brief
14+ */
15+ 
16+#include <iostream>
17+#include <gtest/gtest.h>
18+#include "register/op_impl_registry.h"
19+#include "kernel_run_context_facker.h"
20+#include "../../../op_graph/swish_proto.h"
21+#include "exe_graph/runtime/storage_format.h"
22+#include "exe_graph/runtime/storage_shape.h"
23+#include "log/log.h"
24+#include "platform/platform_info.h"
25+ 
26+class SwishProtoTest : public testing::Test {
27+ protected:
28+ static void SetUpTestCase() {
29+ std::cout << "Swish Proto Test SetUp" << std::endl;
30+ }
31+ 
32+ static void TearDownTestCase() {
33+ std::cout << "Swish Proto Test TearDown" << std::endl;
34+ }
35+};
36+ 
37+TEST_F(SwishProtoTest, swish_infershape_diff_test) {
38+ fe::PlatformInfo platformInfo;
39+ fe::OptionalInfo optiCompilationInfo;
40+ platformInfo.soc_info.ai_core_cnt = 64;
41+ platformInfo.str_info.short_soc_version = "Ascend910_95";
42+ optiCompilationInfo.soc_version = "Ascend910_95";
43+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend910_95"] = platformInfo;
44+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
45+ 
46+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Swish")->infer_shape;
47+ 
48+ gert::Shape input_shape_0 = {4, 3, 4};
49+ gert::Shape output_shape_0 = {};
50+ 
51+ auto holder = gert::InferShapeContextFaker()
52+ .NodeIoNum(1, 1)
53+ .IrInstanceNum({1, 1})
54+ .InputShapes({&input_shape_0})
55+ .OutputShapes({&output_shape_0})
56+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
57+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
58+ .Build();
59+ 
60+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
61+}
62+ 
63+TEST_F(SwishProtoTest, swish_infershape_same_test) {
64+ fe::PlatformInfo platformInfo;
65+ fe::OptionalInfo optiCompilationInfo;
66+ platformInfo.soc_info.ai_core_cnt = 64;
67+ platformInfo.str_info.short_soc_version = "Ascend910_95";
68+ optiCompilationInfo.soc_version = "Ascend910_95";
69+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend910_95"] = platformInfo;
70+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
71+ 
72+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Swish")->infer_shape;
73+ 
74+ gert::Shape input_shape_0 = {1, 3, 4};
75+ gert::Shape output_shape_0 = {};
76+ 
77+ auto holder = gert::InferShapeContextFaker()
78+ .NodeIoNum(1, 1)
79+ .IrInstanceNum({1, 1})
80+ .InputShapes({&input_shape_0})
81+ .OutputShapes({&output_shape_0})
82+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
83+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
84+ .Build();
85+ 
86+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
87+}