Star906
1825
代码
代码
Issues226
Pull Requests303
Actions
项目讨论
Wiki
分析
Star906
1825
已合并
Add Mish for 950 #1529
ligen75创建于 2月4日
Add Mish for 950 #1529
已合并
ligen75创建于 2月4日
从已删除 :master合入到cann/ops-nnmaster
共 25 个文件变更+683-69
Mactivation/mish/CMakeLists.txt+11-14
@@ -1,18 +1,15 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.1+# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 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 "ascend950")
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+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE mish ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
15- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16- add_subdirectory(${SUB_DIR})
17- endif()
18-endforeach()
Mactivation/mish/examples/test_aclnn_mish.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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.
Mactivation/mish/framework/CMakeLists.txt+1-1
@@ -1,6 +1,6 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2# This program is free software, you can redistribute it and/or modify.2# This program is free software, you can redistribute it and/or modify.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4# This file is a part of the CANN Open Software.4# This file is a part of the CANN Open Software.
5# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").5# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6# Please refer to the License for details. You may not use this file except in compliance with the License.6# Please refer to the License for details. You may not use this file except in compliance with the License.
Mactivation/mish/framework/mish_onnx_plugin.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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.
Ractivation/mish/op_host/op_api/aclnn_mish.cpp→activation/mish/op_api/aclnn_mish.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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.
Ractivation/mish/op_host/op_api/aclnn_mish.h→activation/mish/op_api/aclnn_mish.h+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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.
Ractivation/mish/op_host/op_api/mish.cpp→activation/mish/op_api/mish.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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.
Ractivation/mish/op_host/op_api/mish.h→activation/mish/op_api/mish.h+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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.
Ractivation/mish/op_host/CMakeLists.txt→activation/mish/op_graph/CMakeLists.txt+8-6
@@ -1,10 +1,12 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.1+# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE mish ACLNNTYPE aclnn)11+message(STATUS "=== Debug: start ops.activation.mish.graph_plugin.CMakeLists.txt ")
12+add_graph_plugin_sources()
Aactivation/mish/op_graph/mish_proto.h+42-0
@@ -0,0 +1,42 @@
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 mish_proto.h
13+ * \brief
14+ */
15+#ifndef MISH_PROTO_H_
16+#define MISH_PROTO_H_
17+ 
18+#include "graph/operator_reg.h"
19+ 
20+namespace ge {
21+/**
22+*@brief Computes hyperbolic tangent of "x" element-wise .
CANN-robot
CANN-robotCANN-robot2月4日

代码结构与可维护性: 注释中的函数描述与Mish激活函数的实际数学定义不符。Mish函数的定义为：mish(x) = x * tanh(softplus(x)) = x * tanh(ln(1 + e^x))。而注释中写的是'Computes hyperbolic tangent of "x" element-wise'，即'逐元素计算x的双曲正切'，这描述的是tanh函数，不是Mish函数。这是一个严重的文档错误，会导致开发者误解该算子的功能。

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

*@brief Computes hyperbolic tangent of "x" element-wise .

修改建议:

将注释修改为准确的Mish函数描述。例如：'*@brief Computes the Mish activation of "x" element-wise. Mish(x) = x * tanh(softplus(x)) = x * tanh(ln(1 + e^x)).'

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

likedislike
23+
24+*@par Inputs:
25+* One input:
26+* x: An ND tensor. support 1D ~ 8D. Must be one of the following types:
27+* float16, float32, bfloat16.
28+*
29+*@par Outputs:
30+* y: A Tensor. Has the same type as "x" .
31+*
32+*@par Third-party framework compatibility
33+* Compatible with TensorFlow operator Mish.
34+*/
35+
36+REG_OP(Mish)
37+ .INPUT(x, TensorType({ DT_FLOAT, DT_FLOAT16, DT_BF16 }))
38+ .OUTPUT(y, TensorType({ DT_FLOAT, DT_FLOAT16, DT_BF16 }))
39+ .OP_END_FACTORY_REG(Mish)
40+
41+}
CANN-robot
CANN-robotCANN-robot2月4日

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

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

}
#endif

修改建议:

在文件最后一行（#endif之后）添加一个换行符。

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

likedislike
42+#endif
Aactivation/mish/op_host/arch35/mish_tiling_arch35.cpp+153-0文件内容审核中，请稍后刷新重试
Aactivation/mish/op_host/arch35/mish_tiling_arch35.h+38-0文件内容审核中，请稍后刷新重试
Aactivation/mish/op_host/config/ascend910_95/mish_binary.json+92-0
@@ -0,0 +1,92 @@
1+{
2+ "op_type": "Mish",
3+ "op_list": [
4+ {
5+ "bin_filename": "Mish_1c4543fdfe7149b6b0717bd88f061562",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "bfloat16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ }
18+ ],
19+ "outputs": [
20+ {
21+ "name": "y",
22+ "index": 0,
23+ "dtype": "bfloat16",
24+ "format": "ND",
25+ "paramType": "required",
26+ "shape": [
27+ -2
28+ ],
29+ "format_match_mode": "FormatAgnostic"
30+ }
31+ ]
32+ },
33+ {
34+ "bin_filename": "Mish_d2db1a80c523e7e59a032c95969880af",
35+ "inputs": [
36+ {
37+ "name": "x",
38+ "index": 0,
39+ "dtype": "float16",
40+ "format": "ND",
41+ "paramType": "required",
42+ "shape": [
43+ -2
44+ ],
45+ "format_match_mode": "FormatAgnostic"
46+ }
47+ ],
48+ "outputs": [
49+ {
50+ "name": "y",
51+ "index": 0,
52+ "dtype": "float16",
53+ "format": "ND",
54+ "paramType": "required",
55+ "shape": [
56+ -2
57+ ],
58+ "format_match_mode": "FormatAgnostic"
59+ }
60+ ]
61+ },
62+ {
63+ "bin_filename": "Mish_5c5e58e043224cccf51a06108010ae13",
64+ "inputs": [
65+ {
66+ "name": "x",
67+ "index": 0,
68+ "dtype": "float32",
69+ "format": "ND",
70+ "paramType": "required",
71+ "shape": [
72+ -2
73+ ],
74+ "format_match_mode": "FormatAgnostic"
75+ }
76+ ],
77+ "outputs": [
78+ {
79+ "name": "y",
80+ "index": 0,
81+ "dtype": "float32",
82+ "format": "ND",
83+ "paramType": "required",
84+ "shape": [
85+ -2
86+ ],
87+ "format_match_mode": "FormatAgnostic"
88+ }
89+ ]
90+ }
91+ ]
92+}
Aactivation/mish/op_host/config/ascend910_95/mish_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+[Mish]
13+default=0
Aactivation/mish/op_host/mish_def.cpp+46-0
@@ -0,0 +1,46 @@
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 mish_def.cpp
13+ * \brief mish def
14+ */
15+ 
16+#include "register/op_def_registry.h"
17+ 
18+namespace ops {
19+class Mish : public OpDef {
20+public:
21+ explicit Mish(const char* name) : OpDef(name)
22+ {
23+ this->Input("x")
24+ .ParamType(REQUIRED)
25+ .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT})
26+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
27+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
28+ this->Output("y")
29+ .ParamType(REQUIRED)
30+ .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT})
31+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
32+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
33+
34+ OpAICoreConfig aicoreConfig;
35+ aicoreConfig.DynamicCompileStaticFlag(true)
36+ .DynamicFormatFlag(false)
37+ .DynamicRankSupportFlag(true)
38+ .DynamicShapeSupportFlag(true)
39+ .NeedCheckSupportFlag(false)
40+ .PrecisionReduceFlag(true)
41+ .ExtendCfgInfo("opFile.value", "mish_apt");
42+ this->AICore().AddConfig("ascend950", aicoreConfig);
43+ }
44+};
45+OP_ADD(Mish);
46+}
Aactivation/mish/op_host/mish_infershape.cpp+29-0
@@ -0,0 +1,29 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file mish_infershape.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+#include "op_common/op_host/infershape_elewise_util.h"
18+using namespace ge;
19+namespace ops
20+{
21+ 
22+static ge::graphStatus InferShape4Mish(gert::InferShapeContext* context)
23+{
24+ return Ops::Base::InferShape4Elewise(context);
25+}
26+ 
27+IMPL_OP_INFERSHAPE(Mish).InferShape(InferShape4Mish);
28+ 
29+} // namespace ops
Aactivation/mish/op_kernel/arch35/mish_dag.h+123-0
@@ -0,0 +1,123 @@
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 mish_dag.h
13+ * \brief
14+ */
15+ 
16+#ifndef CANN_CUSTOM_OPS_MISH_DAG_H
17+#define CANN_CUSTOM_OPS_MISH_DAG_H
18+
19+#include "atvoss/util/dag.h"
20+#include "atvoss/util/vec.h"
21+#include "atvoss/util/placeholder.h"
22+ 
23+using namespace Ops::Base;
24+const int CAST_MODE_NONE = 0;
25+const int CAST_MODE_RINT = 1;
26+
27+const float FP32_ZERO = 0.0;
28+const float FP32_ONE = 1.0;
29+const float FP32_TWO = 2.0;
30+const float FP32_NEG_ONE = -1.0;
31+const float FP32_NEG_TWO = -2.0;
32+
33+#ifdef __CCE_AICORE__
34+constexpr static AscendC::MicroAPI::CastTrait castTrait0 = { AscendC::MicroAPI::RegLayout::ZERO,
35+AscendC::MicroAPI::SatMode::UNKNOWN, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::UNKNOWN };
36+constexpr static AscendC::MicroAPI::CastTrait castTrait1 = { AscendC::MicroAPI::RegLayout::ZERO,
37+AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_RINT };
38+#endif
39+namespace MishDag1 {
40+
41+template<class T>
42+struct MishCustom : public Vec::ElemwiseUnaryOP<T, T> {
43+ __aicore__ inline MishCustom(LocalTensor<T> &dst, LocalTensor<T> &src, uint32_t count) {
44+ #ifdef __CCE_AICORE__
45+ uint32_t dtypeSize = sizeof(float);
46+ uint32_t vl = VECTOR_REG_WIDTH / dtypeSize;
47+ uint16_t loopNum = CeilDivision(count, vl);
48+ uint32_t vlSize = vl;
49+ __ubuf__ T* srcAddr = (__ubuf__ T*)src.GetPhyAddr();
50+ __ubuf__ T* dstAddr = (__ubuf__ T*)dst.GetPhyAddr();
51+
52+ MicroAPI::RegTensor<float, MicroAPI::RegTraitNumOne> vregInput;
53+ MicroAPI::RegTensor<float, MicroAPI::RegTraitNumOne> vregInputNegNumerator;
54+ MicroAPI::RegTensor<float, MicroAPI::RegTraitNumOne> vregInputNegDenominator;
55+ MicroAPI::RegTensor<float, MicroAPI::RegTraitNumOne> vregInputNumerator;
56+ MicroAPI::RegTensor<float, MicroAPI::RegTraitNumOne> vregInputDenominator;
57+
58+ MicroAPI::RegTensor<float, MicroAPI::RegTraitNumOne> vregOutput;
59+ MicroAPI::MaskReg mask;
60+ MicroAPI::MaskReg cmpMaskReg;
61+ 
62+ MicroAPI::RegTensor<T, MicroAPI::RegTraitNumOne> vregInput16;
63+ MicroAPI::RegTensor<T, MicroAPI::RegTraitNumOne> vregOutput16;
64+ __VEC_SCOPE__ {
65+ for (uint16_t loopIdx = 0; loopIdx < loopNum; loopIdx++) {
66+ mask = MicroAPI::UpdateMask<float, MicroAPI::RegTraitNumOne>(count);
67+ // OpCopyIn
68+ if constexpr(std::is_same_v<T, float>) {
69+ MicroAPI::DataCopy<T, MicroAPI::LoadDist::DIST_NORM>(vregInput, (__ubuf__ T*)(srcAddr + loopIdx * vlSize));
70+ } else {
71+ MicroAPI::DataCopy<T, MicroAPI::LoadDist::DIST_UNPACK_B16>(vregInput16, (__ubuf__ T*)(srcAddr + loopIdx * vlSize));
72+ MicroAPI::Cast<float, T, castTrait0>(vregInput, vregInput16, mask);
73+ }
74+ MicroAPI::Muls(vregInputNegNumerator, vregInput, FP32_NEG_ONE, mask);// -x
75+ MicroAPI::Muls(vregInputNegDenominator, vregInput, FP32_NEG_TWO, mask);// -2x
76+ MicroAPI::Exp(vregInputNegNumerator, vregInputNegNumerator, mask);// e^-x
77+ MicroAPI::Exp(vregInputNegDenominator, vregInputNegDenominator, mask);// e^-2x
78+ MicroAPI::Muls(vregInputNegNumerator, vregInputNegNumerator, FP32_TWO, mask);// 2e^-x
79+ MicroAPI::Adds(vregInputNegNumerator, vregInputNegNumerator, FP32_ONE, mask);// 2e^-x + 1
80+ MicroAPI::Muls(vregInputNegDenominator, vregInputNegDenominator, FP32_TWO, mask);// 2e^-2x
81+ MicroAPI::Add(vregInputNegDenominator, vregInputNegNumerator, vregInputNegDenominator, mask);
82+ MicroAPI::Div(vregOutput, vregInputNegNumerator, vregInputNegDenominator, mask);
83+ 
84+ MicroAPI::Muls(vregInputNumerator, vregInput, FP32_TWO, mask);// 2x
85+ MicroAPI::Exp(vregInputNumerator, vregInputNumerator, mask);// e^2x
86+ MicroAPI::Exp(vregInputDenominator, vregInput, mask);// e^x
87+ MicroAPI::Axpy(vregInputNumerator, vregInputDenominator, FP32_TWO, mask);// e^2x + 2e^x
88+ MicroAPI::Adds(vregInputDenominator, vregInputNumerator, FP32_TWO, mask);// e^2x + 2e^x + 2
89+ MicroAPI::Div(vregInputNumerator, vregInputNumerator, vregInputDenominator, mask);
90+ 
91+ MicroAPI::CompareScalar<float, CMPMODE::LT>(cmpMaskReg, vregInput, FP32_ZERO, mask);
92+ MicroAPI::Select(vregOutput, vregInputNumerator, vregOutput, cmpMaskReg);
93+ MicroAPI::Mul(vregOutput, vregOutput, vregInput, mask);
94+ 
95+ // OpCopyOut
96+ if constexpr(std::is_same_v<T, float>) {
97+ MicroAPI::DataCopy<T, MicroAPI::StoreDist::DIST_NORM_B32>((__ubuf__ T*)(dstAddr + loopIdx * vlSize), vregOutput, mask);
98+ } else {
99+ MicroAPI::Cast<T, float, castTrait1>(vregOutput16, vregOutput, mask);
100+ MicroAPI::DataCopy<T, MicroAPI::StoreDist::DIST_PACK_B32>((__ubuf__ T*)(dstAddr + loopIdx * vlSize), vregOutput16, mask);
101+ }
102+ }
103+ }
104+ #endif
105+ }
106+};
107+}
108+ 
109+template <typename U, typename T=float>
110+struct MishDAG {
111+ using OpCopyIn0 = Bind<Vec::CopyIn<U>, Placeholder::In0<U>>;
112+
113+ using OpCopyIn0Cast = Bind<Vec::Cast<T, U, CAST_MODE_NONE>, OpCopyIn0>;
114+ using OpResult1 = Bind<MishDag1::MishCustom<T>, OpCopyIn0Cast>;
115+ using OpResultCast = Bind<Vec::Cast<U, T, CAST_MODE_RINT>, OpResult1>;
116+
117+ using OpCopyOut = Bind<Vec::CopyOut<U>, Placeholder::Out0<U>, OpResultCast>;
118+
119+ using Outputs = Elems<OpCopyOut>;
120+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
121+ using OpDag = DAGSch<Outputs, void, MemCfg>;
122+};
123+ #endif // CANN_CUSTOM_OPS_MISH_DAG_H
Aactivation/mish/op_kernel/arch35/mish_struct.h+39-0
@@ -0,0 +1,39 @@
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 mish_struct.h
13+ * \brief
14+ */
15+ 
16+#include "ascendc/host_api/tiling/template_argument.h"
17+ 
18+#ifndef CANN_CUSTOM_OPS_MISH_STRUCT_H_
19+#define CANN_CUSTOM_OPS_MISH_STRUCT_H_
20+ 
21+#define TPL_FP16 1
22+#define TPL_BF16 2
23+#define TPL_FP32 3
24+
25+#define TPL_SCH_MODE_0 0
26+#define TPL_SCH_MODE_1 1
27+
28+ASCENDC_TPL_ARGS_DECL(Mish,
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_FP16, TPL_BF16, TPL_FP32)
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_FP16, TPL_BF16, TPL_FP32)
37+ )
38+);
39+#endif // CANN_CUSTOM_OPS_GELU_STRUCT_H_
Aactivation/mish/op_kernel/mish_apt.cpp+43-0
@@ -0,0 +1,43 @@
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 mish_apt.cpp
13+ * \brief mish kernel
14+ */
15+#include "kernel_operator.h"
16+#include "kernel_tiling/kernel_tiling.h"
17+#include "arch35/mish_dag.h"
18+#include "arch35/mish_struct.h"
19+#include "atvoss/elewise/elewise_sch_16b.h"
20+#include "atvoss/util/dfx.h"
21+ 
22+using namespace AscendC;
23+ 
24+template <uint64_t schMode, uint64_t dType>
25+__global__ __aicore__ void mish(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
26+ REGISTER_TILING_DEFAULT(EleBaseTilingData16B);
27+ GET_TILING_DATA_PTR_WITH_STRUCT(EleBaseTilingData16B, tilingData, tiling);
28+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
29+ if constexpr (dType == TPL_FP16) {
30+ ElementwiseSch16B<schMode, MishDAG<half>::OpDag> sch(tilingData);
31+ sch.Init(x, y);
32+ sch.Process();
33+ } else if constexpr (dType == TPL_BF16) {
34+ ElementwiseSch16B<schMode, MishDAG<bfloat16_t>::OpDag> sch(tilingData);
35+ sch.Init(x, y);
36+ sch.Process();
37+ } else if constexpr (dType == TPL_FP32) {
38+ ElementwiseSch16B<schMode, MishDAG<float>::OpDag> sch(tilingData);
39+ sch.Init(x, y);
40+ sch.Process();
41+ }
42+ return;
43+ }
Mactivation/mish/tests/CMakeLists.txt+13-13
@@ -1,17 +1,17 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.1+# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-message(STATUS "=== Debug: start ops.activation.mish.tests.CMakeLists.txt ")11+ file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)12+ message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
12-message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")13+ foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
13-foreach(SUB_DIR ${CURRENT_DIRS})14+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
14- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")15+ add_subdirectory(${SUB_DIR})
15- add_subdirectory(${SUB_DIR})16+ endif()
16- endif()17+ endforeach()
17-endforeach()
Mactivation/mish/tests/ut/CMakeLists.txt+12-14
@@ -1,18 +1,16 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.1+# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~12411+ file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
11-message(STATUS "=== Debug: start ops.activation.mish.tests.ut.CMakeLists.txt ")12+ foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
12-file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)13+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
13-message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")14+ add_subdirectory(${SUB_DIR})
14-foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})15+ endif()
15- if(EXISTS "${SUB_DIR}/CMakeLists.txt")16+ endforeach()
16- add_subdirectory(${SUB_DIR})
17- endif()
18-endforeach()
Mactivation/mish/tests/ut/op_host/CMakeLists.txt+12-14
@@ -1,17 +1,15 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.1+ # -----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 2+# This program is free software, you can redistribute it and/or modify.
3-# CANN Open Software License Agreement Version 2.0 (the "License").3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under 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.6# 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, 7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
6-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.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}/*)
11-if(UT_TEST_ALL OR OP_HOST_UT)12+ if(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})
14-endif()15+ endif()
15-if(UT_TEST_ALL OR OP_API_UT)
16- add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
17-endif()
Mactivation/mish/tests/ut/op_host/test_aclnn_mish.cpp+1-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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.
Mscripts/kernel/binary_config/ascendc_config.json+1-0
@@ -35,6 +35,7 @@
35 {"name":"EluGradV2", "compute_units": ["ascend950"], "auto_sync": false, "compile_options": {"ascend950": ["--cce-simd-vf-fusion=true", "-mllvm -cce-aicore-dcci-before-kernel-end=false"]}, "impl_mode": ""},35 {"name":"EluGradV2", "compute_units": ["ascend950"], "auto_sync": false, "compile_options": {"ascend950": ["--cce-simd-vf-fusion=true", "-mllvm -cce-aicore-dcci-before-kernel-end=false"]}, "impl_mode": ""},
36 {"name":"LeakyRelu", "compute_units": ["ascend950"], "auto_sync" : false, "compile_options": {"ascend950":["--cce-simd-vf-fusion=true", "--cce-simd-at-voss=true"]}, "impl_mode" : ""},36 {"name":"LeakyRelu", "compute_units": ["ascend950"], "auto_sync" : false, "compile_options": {"ascend950":["--cce-simd-vf-fusion=true", "--cce-simd-at-voss=true"]}, "impl_mode" : ""},
37 {"name":"Softplus", "compute_units": ["ascend950"], "auto_sync" : false, "compile_options": {"ascend950":["--cce-simd-vf-fusion=true", "--cce-simd-at-voss=true"]}, "impl_mode" : ""},37 {"name":"Softplus", "compute_units": ["ascend950"], "auto_sync" : false, "compile_options": {"ascend950":["--cce-simd-vf-fusion=true", "--cce-simd-at-voss=true"]}, "impl_mode" : ""},
38+ {"name":"Mish", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
38 {"name":"Elu", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},39 {"name":"Elu", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
39 {"name":"RsqrtGrad", "compute_units": ["ascend950"], "auto_sync" : false, "compile_options": {"ascend950":["--cce-simd-vf-fusion=true", "--cce-simd-at-voss=true"]}, "impl_mode" : ""},40 {"name":"RsqrtGrad", "compute_units": ["ascend950"], "auto_sync" : false, "compile_options": {"ascend950":["--cce-simd-vf-fusion=true", "--cce-simd-at-voss=true"]}, "impl_mode" : ""},
40 {"name":"AddN", "compute_units": ["ascend950"], "auto_sync" : false, "compile_options": {"ascend950":["--cce-simd-vf-fusion=true", "--cce-simd-at-voss=true"]}, "impl_mode" : ""},41 {"name":"AddN", "compute_units": ["ascend950"], "auto_sync" : false, "compile_options": {"ascend950":["--cce-simd-vf-fusion=true", "--cce-simd-at-voss=true"]}, "impl_mode" : ""},