Star936
1856
代码
代码
Issues242
Pull Requests347
Actions
项目讨论
Wiki
分析
Star936
1856
已合并
【CANN训练营第二季社区任务】Mish算子开发-算子提交 #490
wenyidylan创建于 2025年12月24日
【CANN训练营第二季社区任务】Mish算子开发-算子提交 #490
已合并
wenyidylan创建于 2025年12月24日
从wendylan/ops-nn-mymishmaster合入到cann/ops-nnmaster
共 23 个文件变更+1774-0
Aexperimental/activation/mish/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
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+ 
10+message(STATUS "=== Debug: start ops.activation.mish.CMakeLists.txt ")
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST AND NOT BENCHMARK)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aexperimental/activation/mish/README.md+71-0
@@ -0,0 +1,71 @@
1+# Mish
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Atlas A2 训练系列产品/Atlas 800I A2推理产品</term> | √ |
8+ 
9+## 功能说明
10+ 
11+- 算子功能：逐元素计算张量的Mish激活函数。
12+- SUPER_PERFORMANCE模式的计算公式为：
13+$$y = x*(1-2/(1+(1+(1+x/64)^{64})^2))$$
14+其他模式的计算公式为：
15+$$y =
16+\begin{cases}
17+x*(2e^{-x} + 1) / (2e^{-2x} + 2e^{-x} + 1), & \text{if $x > 0$} \\
18+x*(2e^x + e^{2x}) / (2 + 2e^x + e^{2x}), & \text{if $x \leq 0$}
19+\end{cases}$$
20+ 
21+## 参数说明
22+ 
23+<table style="undefined;table-layout: fixed; width: 820px"><colgroup>
24+ <col style="width: 100px">
25+ <col style="width: 150px">
26+ <col style="width: 190px">
27+ <col style="width: 260px">
28+ <col style="width: 120px">
29+ </colgroup>
30+ <thead>
31+ <tr>
32+ <th>参数名</th>
33+ <th>输入/输出/属性</th>
34+ <th>描述</th>
35+ <th>数据类型</th>
36+ <th>数据格式</th>
37+ </tr></thead>
38+ <tbody>
39+ <tr>
40+ <td>x</td>
41+ <td>输入</td>
42+ <td>公式中的输入张量x</td>
43+ <td>FLOAT、FLOAT16、BFLOAT16</td>
44+ <td>ND</td>
45+ </tr>
46+ <tr>
47+ <td>y</td>
48+ <td>输出</td>
49+ <td>公式中的输出张量y</td>
50+ <td>FLOAT、FLOAT16、BFLOAT16</td>
51+ <td>ND</td>
52+ </tr>
53+ </tbody></table>
54+ 
55+## 约束说明
56+ 
57+ 
58+## 调用说明
59+ 
60+| 调用方式 | 调用样例 | 说明 |
61+|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|
62+| aclnn调用 | [test_aclnn_mish](./examples/test_aclnn_mish.cpp) | 通过aclnnMish接口方式调用Mish算子。 |
63+ 
condfuse_3
condfuse_3condfuse_32月9日

加一个贡献说明 参考：

贡献说明

 | 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
 | ---- | ---- | ---- | ---- | ---- |
 | ilovescrapy | 个人开发者 | ReluGrad | 2025/12/26 | ReluGrad算子适配开源仓 |
likedislike
64+ 
65+## 贡献说明
66+ 
67+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
68+| ---- | ---- | ---- | ---- | ---- |
69+| 闻毅 | 个人开发者 | Mish | 2025/12/23 | Mish算子适配开源仓 |
70+ 
71+ 
Aexperimental/activation/mish/examples/test_aclnn_mish.cpp+119-0
@@ -0,0 +1,119 @@
1+#include <iostream>
2+#include <vector>
3+#include "acl/acl.h"
4+#include "aclnn_mish.h"
5+ 
6+#define CHECK_RET(cond, return_expr) \
7+ do { \
8+ if (!(cond)) { \
9+ return_expr; \
10+ } \
11+ } while (0)
12+#define LOG_PRINT(message, ...) \
13+ do { \
14+ printf(message, ##__VA_ARGS__); \
15+ } while (0)
16+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
17+ int64_t shapeSize = 1;
18+ for (auto i : shape) {
19+ shapeSize *= i;
20+ }
21+ return shapeSize;
22+}
23+int Init(int32_t deviceId, aclrtStream* stream) {
24+ // 固定写法，资源初始化
25+ auto ret = aclInit(nullptr);
26+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
27+ ret = aclrtSetDevice(deviceId);
28+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
29+ ret = aclrtCreateStream(stream);
30+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
31+ return 0;
32+}
33+ 
34+template <typename T>
35+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
36+ aclDataType dataType, aclTensor** tensor) {
37+ auto size = GetShapeSize(shape) * sizeof(T);
38+ // 调用aclrtMalloc申请device侧内存
39+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
41+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
42+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
44+ // 计算连续tensor的strides
45+ std::vector<int64_t> strides(shape.size(), 1);
46+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
47+ strides[i] = shape[i + 1] * strides[i + 1];
48+ }
49+ // 调用aclCreateTensor接口创建aclTensor
50+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
51+ shape.data(), shape.size(), *deviceAddr);
52+ return 0;
53+}
54+int main() {
55+ // 1. （固定写法）device/stream初始化，参考acl API手册
56+ // 根据自己的实际device填写deviceId
57+ int32_t deviceId = 0;
58+ aclrtStream stream;
59+ auto ret = Init(deviceId, &stream);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
61+ // 2. 构造输入与输出，需要根据API的接口自定义构造
62+ std::vector<int64_t> selfShape = {2, 2};
63+ std::vector<int64_t> outShape = {2, 2};
64+ void* selfDeviceAddr = nullptr;
65+ void* outDeviceAddr = nullptr;
66+ aclTensor* self = nullptr;
67+ aclTensor* out = nullptr;
68+ std::vector<float> selfHostData = {0, 1, 2, 3};
69+ std::vector<float> outHostData = {0, 0, 0, 0};
70+ // 创建self aclTensor
71+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
72+ CHECK_RET(ret == ACL_SUCCESS, return ret);
73+ // 创建out aclTensor
74+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
75+ CHECK_RET(ret == ACL_SUCCESS, return ret);
76+ // 3. 调用CANN算子库API，需要修改为具体的API名称
77+ // aclnnMish接口调用示例
78+ uint64_t workspaceSize = 0;
79+ aclOpExecutor* executor;
80+ // 调用aclnnMish第一段接口
81+ ret = aclnnMishGetWorkspaceSize(self, out, &workspaceSize, &executor);
82+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMishGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
83+ // 根据第一段接口计算出的workspaceSize申请device内存
84+ void* workspaceAddr = nullptr;
85+ if (workspaceSize > 0) {
86+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
87+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
88+ }
89+ // 调用aclnnMish第二段接口
90+ ret = aclnnMish(workspaceAddr, workspaceSize, executor, stream);
91+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMish failed. ERROR: %d\n", ret); return ret);
92+ // 4. （固定写法）同步等待任务执行结束
93+ ret = aclrtSynchronizeStream(stream);
94+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
95+ 
96+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
97+ auto size = GetShapeSize(outShape);
98+ std::vector<float> resultData(size, 0);
99+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
100+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
101+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
102+ for (int64_t i = 0; i < size; i++) {
103+ LOG_PRINT("aclnnMish result[%ld] is: %f\n", i, resultData[i]);
104+ }
105+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
106+ aclDestroyTensor(self);
107+ aclDestroyTensor(out);
108+ 
109+ // 7. 释放device资源，需要根据具体API的接口定义修改
110+ aclrtFree(selfDeviceAddr);
111+ aclrtFree(outDeviceAddr);
112+ if (workspaceSize > 0) {
113+ aclrtFree(workspaceAddr);
114+ }
115+ aclrtDestroyStream(stream);
116+ aclrtResetDevice(deviceId);
117+ aclFinalize();
118+ return 0;
119+}
Aexperimental/activation/mish/op_host/CMakeLists.txt+10-0
@@ -0,0 +1,10 @@
1+# This program is free software, you can redistribute it and/or modify it.
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This file is a part of the CANN Open Software.
4+# Licensed under 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+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE mish ACLNNTYPE aclnn_exclude)
Aexperimental/activation/mish/op_host/mish_def.cpp+36-0
@@ -0,0 +1,36 @@
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, 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+ 
10+/*!
11+ * \file mish_def.cpp
12+ * \brief
13+*/
14+#include "register/op_def_registry.h"
15+ 
16+namespace ops {
17+class Mish : public OpDef {
18+public:
19+ explicit Mish(const char* name) : OpDef(name)
20+ {
21+ this->Input("x")
22+ .ParamType(REQUIRED)
23+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
24+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
25+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
26+ this->Output("y")
27+ .ParamType(REQUIRED)
28+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
29+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
30+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
31+
32+ this->AICore().AddConfig("ascend910b");
33+ }
34+};
35+OP_ADD(Mish); // 添加算子信息库
36+} // namespace ops
Aexperimental/activation/mish/op_host/mish_infershape.cpp+41-0
@@ -0,0 +1,41 @@
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, 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+ 
10+/*!
11+ * \file mish_infershape.cpp
12+ * \brief
13+*/
14+#include "register/op_impl_registry.h"
15+#include "log/log.h"
16+ 
17+using namespace ge;
18+ 
19+namespace ops {
20+static constexpr int64_t IDX_0 = 0;
21+ 
22+static ge::graphStatus InferShapeMish(gert::InferShapeContext* context)
23+{
24+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
25+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeMish");
26+ 
27+ // get input shapes
28+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
29+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
30+ 
31+ // get output shapes
32+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
33+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
34+ 
35+ *yShape = *xShape;
36+ OP_LOGD(context->GetNodeName(), "End to do InferShapeMish");
37+ return GRAPH_SUCCESS;
38+}
39+ 
40+IMPL_OP_INFERSHAPE(Mish).InferShape(InferShapeMish);
41+}
Aexperimental/activation/mish/op_host/mish_tiling.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, 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+ 
10+/*!
11+ * \file mish_tiling.cpp
12+ * \brief
13+ */
14+#include "log/log.h"
15+#include "util/math_util.h"
16+#include "op_host/tiling_util.h"
17+#include "tiling/platform/platform_ascendc.h"
18+#include "register/op_impl_registry.h"
19+#include "op_host/tiling_templates_registry.h"
20+#include "../op_kernel/mish_tiling_data.h"
21+#include "../op_kernel/mish_tiling_key.h"
22+#include "op_common/op_host/util/platform_util.h"
23+ 
24+namespace optiling {
25+ 
26+using namespace Ops::NN::OpTiling;
27+ 
28+#define BLOCK_SIZE Ops::Base::GetUbBlockSize(context)
29+#define OPT_CORE_SIZE 1024U
30+#define BLOCK_ALIGN_NUM 16U
31+#define UB_DATA_NUM_FLOAT 12U // 对应DT_FLOAT类型的ub分块数量
32+#define UB_DATA_NUM_OTHER 14U // 对应其他数据类型的ub分块数量
33+constexpr uint32_t BUFFER_NUM = 2;
34+constexpr uint32_t WS_SYS_SIZE = 0;
35+struct MishCompileInfo {};
36+ 
37+static ge::graphStatus TilingParseForMish([[maybe_unused]] gert::TilingParseContext* context)
38+{
39+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
40+ return ge::GRAPH_SUCCESS;
41+}
42+ 
43+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
44+{
45+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
46+ // 获取ubsize coreNum
47+ OP_CHECK_IF(context->GetPlatformInfo() == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
48+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
condfuse_3
condfuse_3condfuse_32025年12月24日

context->GetPlatformInfo()需要判空 OP_CHECK_IF(context->GetPlatformInfo() == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);

likedislike
49+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
50+ coreNum = ascendcPlatform.GetCoreNum();
51+ OP_CHECK_IF(coreNum <= 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
52+ OP_CHECK_IF(ubSize <= 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
53+ return ge::GRAPH_SUCCESS;
54+}
55+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
56+{
57+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
58+ size_t usrSize = 0;
59+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
60+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
61+ size_t* currentWorkspace = context->GetWorkspaceSizes(
62+ 1); // 通过框架获取workspace的指针，GetWorkspaceSizes入参为所需workspace的块数。当前限制使用一块。
63+ currentWorkspace[0] = usrSize + sysWorkspaceSize;
64+ return ge::GRAPH_SUCCESS;
65+}
66+ 
67+static ge::graphStatus GetShapeAttrsInfo(
68+ gert::TilingContext* context, uint64_t ubSize, uint64_t& inputNum, uint64_t& inputBytes, uint64_t& tileBlockNum,
69+ uint64_t& tileDataNum, uint64_t& inputLengthAlgin32)
70+{
71+ OP_CHECK_IF(
72+ context == nullptr || context->GetInputShape(0) == nullptr, OP_LOGE(context, "context is nullptr"),
73+ return ge::GRAPH_FAILED);
74+ inputNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
75+ uint32_t typeLength = 0;
76+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
77+ uint64_t inputLength = inputNum * typeLength;
78+ if (inputNum == 0) {
condfuse_3
condfuse_3condfuse_32025年12月24日

异常返回，添加打印

likedislike
79+ OP_LOGE(context, "inputNum is 0");
80+ return ge::GRAPH_FAILED;
81+ }
82+ inputBytes = inputLength / inputNum;
83+ uint64_t ubDataNumber =
84+ (context->GetInputDesc(0)->GetDataType() == ge::DT_FLOAT) ? UB_DATA_NUM_FLOAT : UB_DATA_NUM_OTHER;
85+ tileBlockNum = (ubSize / BLOCK_SIZE) / ubDataNumber;
86+ if (inputBytes == 0) {
condfuse_3
condfuse_3condfuse_32025年12月24日

异常返回，添加打印

likedislike
87+ OP_LOGE(context, "inputBytes is 0");
88+ return ge::GRAPH_FAILED;
89+ }
90+ tileBlockNum = tileBlockNum <= BLOCK_ALIGN_NUM ? tileBlockNum : tileBlockNum / BLOCK_ALIGN_NUM * BLOCK_ALIGN_NUM;
91+ tileDataNum = (tileBlockNum * BLOCK_SIZE) / inputBytes;
92+ inputLengthAlgin32 = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE);
93+ return ge::GRAPH_SUCCESS;
94+}
95+ 
96+static ge::graphStatus CalculateCoreBlockNums(
97+ gert::TilingContext* context,
98+ uint64_t inputLengthAlgin32, int64_t coreNum, uint64_t tileBlockNum, uint64_t inputBytes, uint64_t tileDataNum,
99+ uint64_t& smallCoreDataNum, uint64_t& bigCoreDataNum, uint64_t& smallTailDataNum, uint64_t& bigTailDataNum,
100+ uint64_t& finalSmallTileNum, uint64_t& finalBigTileNum, uint64_t& tailBlockNum)
101+{
102+ if (0 == BLOCK_SIZE || 0 == coreNum || 0 == tileBlockNum || 0 == inputBytes) {
condfuse_3
condfuse_3condfuse_32025年12月24日

异常返回，添加打印

likedislike
103+ OP_LOGE(context, "BLOCK_SIZE is 0 or coreNum is 0 or tileBlockNum is 0 or inputBytes is 0");
104+ return ge::GRAPH_FAILED;
105+ }
106+ uint64_t everyCoreInputBlockNum = inputLengthAlgin32 / BLOCK_SIZE / coreNum;
107+ tailBlockNum = (inputLengthAlgin32 / BLOCK_SIZE) % coreNum;
108+ smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
109+ uint64_t smallTileNum = everyCoreInputBlockNum / tileBlockNum;
110+ finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
111+ smallTailDataNum = smallCoreDataNum - (tileDataNum * smallTileNum);
112+ smallTailDataNum = smallTailDataNum == 0 ? tileDataNum : smallTailDataNum;
113+ 
114+ everyCoreInputBlockNum += 1;
115+ bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
116+ uint64_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
117+ finalBigTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
118+ bigTailDataNum = bigCoreDataNum - tileDataNum * bigTileNum;
119+ bigTailDataNum = bigTailDataNum == 0 ? tileDataNum : bigTailDataNum;
120+ 
121+ return ge::GRAPH_SUCCESS;
122+}
123+
124+static ge::graphStatus MishTilingFunc(gert::TilingContext* context)
125+{
126+ // MishTilingData tiling;
127+ MishTilingData* tiling = context->GetTilingData<MishTilingData>();
128+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
129+ OP_CHECK_IF(
130+ memset_s(tiling, sizeof(MishTilingData), 0, sizeof(MishTilingData)) != EOK,
131+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
132+ // 获取平台运行信息
133+ uint64_t ubSize;
134+ int64_t coreNum;
135+ int64_t usedcoreNum;
136+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, coreNum);
137+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
138+ // 获取输入数据信息
139+ uint64_t inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin32;
140+ ret = GetShapeAttrsInfo(context, ubSize, inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin32);
141+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
142+ 
143+ // 计算coreNum
144+ if (OPT_CORE_SIZE >= inputNum) {
145+ usedcoreNum = 1;
146+ } else {
147+ usedcoreNum = inputNum / OPT_CORE_SIZE > coreNum ? coreNum : inputNum / OPT_CORE_SIZE;
148+ if(inputNum % OPT_CORE_SIZE > 0) {
149+ if(usedcoreNum < coreNum) {
150+ usedcoreNum += 1;
151+ }
152+ }
153+ }
154+ // 计算每个core处理的数据块数
155+ uint64_t smallCoreDataNum, bigCoreDataNum, smallTailDataNum, bigTailDataNum;
156+ uint64_t finalSmallTileNum, finalBigTileNum, tailBlockNum;
157+ ret = CalculateCoreBlockNums(
158+ context,
159+ inputLengthAlgin32, usedcoreNum, tileBlockNum, inputBytes, tileDataNum, smallCoreDataNum, bigCoreDataNum,
160+ smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum);
161+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "CalculateCoreBlockNums error"), return ge::GRAPH_FAILED);
162+ // 设置tiling数据
163+ tiling->smallCoreDataNum = static_cast<uint64_t>(smallCoreDataNum);
164+ tiling->bigCoreDataNum = static_cast<uint64_t>(bigCoreDataNum);
165+ tiling->tileDataNum = static_cast<uint64_t>(tileDataNum);
166+ tiling->smallTailDataNum = static_cast<uint64_t>(smallTailDataNum);
167+ tiling->bigTailDataNum = static_cast<uint64_t>(bigTailDataNum);
168+ tiling->finalSmallTileNum = static_cast<uint64_t>(finalSmallTileNum);
169+ tiling->finalBigTileNum = static_cast<uint64_t>(finalBigTileNum);
170+ tiling->tailBlockNum = static_cast<uint64_t>(tailBlockNum);
171+ // 计算workspace大小
172+ OP_CHECK_IF(
173+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
174+ return ge::GRAPH_FAILED);
175+ uint64_t tilingKey = 0;
176+ tilingKey = GET_TPL_TILING_KEY(0);
177+ context->SetTilingKey(tilingKey);
178+ context->SetBlockDim(usedcoreNum);
179+ return ge::GRAPH_SUCCESS;
180+}
181+ 
182+// tiling注册入口.
183+IMPL_OP_OPTILING(Mish).Tiling(MishTilingFunc).TilingParse<MishCompileInfo>(TilingParseForMish);
184+} // namespace optiling
Aexperimental/activation/mish/op_host/op_api/aclnn_mish.cpp+123-0
@@ -0,0 +1,123 @@
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, 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+ 
10+#include "aclnn_kernels/cast.h"
11+#include "aclnn_kernels/contiguous.h"
12+#include "opdev/op_log.h"
13+#include "opdev/op_dfx.h"
14+#include "op_api/op_api_def.h"
15+#include "opdev/common_types.h"
16+#include "opdev/data_type_utils.h"
17+#include "opdev/make_op_executor.h"
18+#include "opdev/platform.h"
19+#include "aclnn_kernels/common/op_error_check.h"
20+#include "op_api/level2_base.h"
21+#include "mish.h"
22+#include "aclnn_mish.h"
23+ 
24+using namespace op;
25+ 
26+#ifdef __cplusplus
27+extern "C" {
28+#endif
29+ 
30+static const std::initializer_list<DataType> ASCEND910_DTYPE_SUPPORT_LIST = {
31+ DataType::DT_FLOAT, DataType::DT_FLOAT16};
32+ 
33+static const std::initializer_list<DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {
34+ DataType::DT_FLOAT, DataType::DT_FLOAT16, DataType::DT_BF16};
35+ 
36+static bool CheckDtypeValid(const aclTensor* self, const aclTensor* out) {
37+ // 检查self的数据类型是否在支持列表内
38+ auto supportList = GetDtypeSupportListV2(ASCEND910B_DTYPE_SUPPORT_LIST, ASCEND910_DTYPE_SUPPORT_LIST);
39+ OP_CHECK_DTYPE_NOT_SUPPORT(self, supportList, return false);
40+ OP_CHECK_DTYPE_NOT_SUPPORT(out, supportList, return false);
41+ // 检查self和out的数据类型是否相等
42+ OP_CHECK_DTYPE_NOT_SAME(self, out, return false);
43+ return true;
44+}
45+ 
46+static bool CheckShape(const aclTensor* self, const aclTensor* out) {
47+ OP_CHECK_MAX_DIM(self, MAX_SUPPORT_DIMS_NUMS, return false);
48+ OP_CHECK_MAX_DIM(out, MAX_SUPPORT_DIMS_NUMS, return false);
49+ // 输入和输出的shape必须一致
50+ OP_CHECK_SHAPE_NOT_EQUAL(out, self, return false);
51+ return true;
52+}
53+ 
54+static aclnnStatus CheckParams(const aclTensor* self, const aclTensor* out) {
55+ // 1. 检查参数是否为空指针
56+ CHECK_RET(CheckNotNull2Tensor(self, out), ACLNN_ERR_PARAM_NULLPTR);
57+ 
58+ // 2. 检查输入的数据类型是否在API支持的数据类型范围内，需要根据api定义校验
59+ CHECK_RET(CheckDtypeValid(self, out), ACLNN_ERR_PARAM_INVALID);
60+ 
61+ // 3. 检查shape是否满足约束
62+ CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID);
63+ 
64+ return ACLNN_SUCCESS;
65+}
66+ 
67+aclnnStatus aclnnMishGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize,
68+ aclOpExecutor** executor) {
69+ L2_DFX_PHASE_1(aclnnMish, DFX_IN(self), DFX_OUT(out));
70+ 
71+ // 固定写法，创建OpExecutor
72+ auto uniqueExecutor = CREATE_EXECUTOR();
73+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
74+ 
75+ // 固定写法，参数检查
76+ auto ret = CheckParams(self, out);
77+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
78+ 
79+ // 空Tensor处理
80+ if (self->IsEmpty()) {
81+ *workspaceSize = 0;
82+ uniqueExecutor.ReleaseTo(executor);
83+ return ACLNN_SUCCESS;
84+ }
85+ 
86+ // self如果非连续，需要转换
87+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
88+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
89+ 
90+ // 调用l0算子进行计算
91+ auto mishResult = l0op::Mish(selfContiguous, uniqueExecutor.get());
92+ CHECK_RET(mishResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
93+ 
94+ // 如果出参out是非连续Tensor，需要把计算完的连续Tensor转非连续
95+ auto viewCopyResult = l0op::ViewCopy(mishResult, out, uniqueExecutor.get());
96+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
97+ 
98+ // 固定写法，获取计算过程中需要使用的workspace大小
99+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
100+ uniqueExecutor.ReleaseTo(executor);
101+ return ACLNN_SUCCESS;
102+}
103+ 
104+aclnnStatus aclnnMish(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream) {
105+ L2_DFX_PHASE_2(aclnnMish);
106+ 
107+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
108+}
109+ 
110+aclnnStatus aclnnInplaceMishGetWorkspaceSize(aclTensor *selfRef, uint64_t *workspaceSize, aclOpExecutor **executor) {
111+ L2_DFX_PHASE_1(aclnnInplaceMish, DFX_IN(selfRef), DFX_OUT(selfRef));
112+ return aclnnMishGetWorkspaceSize(selfRef, selfRef, workspaceSize, executor);
113+}
114+ 
115+aclnnStatus aclnnInplaceMish(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream) {
116+ L2_DFX_PHASE_2(aclnnInplaceMish);
117+ // 调用框架能力，完成计算
118+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
119+}
120+ 
121+#ifdef __cplusplus
122+}
123+#endif
Aexperimental/activation/mish/op_host/op_api/aclnn_mish.h+81-0
@@ -0,0 +1,81 @@
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, 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+#ifndef OP_API_INC_MISH_H_
10+#define OP_API_INC_MISH_H_
11+ 
12+#include "aclnn/aclnn_base.h"
13+#include "aclnn_util.h"
14+ 
15+#ifdef __cplusplus
16+extern "C" {
17+#endif
18+ 
19+/**
20+ * @brief aclnnMish的第一段接口，根据具体的计算流程，计算workspace大小。
21+ * @domain aclnn_ops_infer
22+ *
23+ * 算子功能：一个自正则化的非单调神经网络激活函数。
24+ *
25+ * @param [in] self: npu device侧的aclTensor，数据类型支持FLOAT16、BFLOAT16、FLOAT。支持非连续的Tensor，数据格式支持ND。
26+ * @param [in] out: npu
27+ * device侧的aclTensor，数据类型支持FLOAT16、BFLOAT16、FLOAT。它的shape与self相同，且数据类型需要与self一致，数据格式支持ND。
28+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
29+ * @param [out] executor: 返回op执行器，包含算子计算流程。
30+ * @return aclnnStatus: 返回状态码。
31+ */
32+ACLNN_API aclnnStatus aclnnMishGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize,
33+ aclOpExecutor** executor);
34+ 
35+/**
36+ * @brief aclnnMish的第二段接口，用于执行计算。
37+ *
38+ * 算子功能：一个自正则化的非单调神经网络激活函数。
39+ *
40+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
41+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnMishGetWorkspaceSize获取。
42+ * @param [in] executor: op执行器，包含了算子计算流程。
43+ * @param [in] stream: acl stream流。
44+ * @return aclnnStatus: 返回状态码。
45+ */
46+ACLNN_API aclnnStatus aclnnMish(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
47+ 
48+/**
49+ * @brief aclnnInplaceMish的第一段接口，根据具体的计算流程，计算workspace大小。
50+ * @domain aclnn_ops_infer
51+ *
52+ * 算子功能：一个自正则化的非单调神经网络激活函数。
53+ *
54+ * @param [in] selfRef: npu
55+ * device侧的aclTensor，数据类型支持FLOAT16、BFLOAT16、FLOAT。支持非连续的Tensor，数据格式支持ND。
56+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
57+ * @param [out] executor: 返回op执行器，包含算子计算流程。
58+ * @return aclnnStatus: 返回状态码。
59+ */
60+ACLNN_API aclnnStatus aclnnInplaceMishGetWorkspaceSize(aclTensor* selfRef, uint64_t* workspaceSize,
61+ aclOpExecutor** executor);
62+ 
63+/**
64+ * @brief aclnnInplaceMish的第二段接口，用于执行计算。
65+ *
66+ * 算子功能：一个自正则化的非单调神经网络激活函数。
67+ *
68+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
69+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnInplaceMishGetWorkspaceSize获取。
70+ * @param [in] executor: op执行器，包含了算子计算流程。
71+ * @param [in] stream: acl stream流。
72+ * @return aclnnStatus: 返回状态码。
73+ */
74+ACLNN_API aclnnStatus aclnnInplaceMish(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
75+ aclrtStream stream);
76+ 
77+#ifdef __cplusplus
78+}
79+#endif
80+ 
81+#endif // OP_API_INC_MISH_H_
Aexperimental/activation/mish/op_host/op_api/mish.cpp+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, 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+#include "opdev/make_op_executor.h"
10+#include "opdev/op_def.h"
11+#include "opdev/op_dfx.h"
12+#include "opdev/op_executor.h"
13+#include "opdev/op_log.h"
14+#include "opdev/shape_utils.h"
15+#include "aclnn_kernels/common/op_error_check.h"
16+#include "mish.h"
17+ 
18+using namespace op;
19+ 
20+namespace l0op {
21+OP_TYPE_REGISTER(Mish);
22+ 
23+// AICORE算子kernel
24+static const aclTensor *MishAiCore(const aclTensor *self, aclTensor *out, aclOpExecutor *executor) {
25+ L0_DFX(MishAiCore, self, out);
26+ // 使用框架宏ADD_TO_LAUNCHER_LIST_AICORE，将Aicore Mish算子加入任务队列
27+ // Mish是算子的OpType，self是算子的输入，out是算子的输出
28+ auto retAicore = ADD_TO_LAUNCHER_LIST_AICORE(Mish, OP_INPUT(self), OP_OUTPUT(out));
29+ OP_CHECK_ADD_TO_LAUNCHER_LIST_AICORE(retAicore != ACLNN_SUCCESS, return nullptr,
30+ "Mish ADD_TO_LAUNCHER_LIST_AICORE failed.");
31+ return out;
32+}
33+ 
34+const aclTensor *Mish(const aclTensor *self, aclOpExecutor *executor) {
35+ auto mishOut = executor->AllocTensor(self->GetViewShape(), self->GetDataType());
36+ return MishAiCore(self, mishOut, executor);
37+}
38+} // namespace l0op
Aexperimental/activation/mish/op_host/op_api/mish.h+18-0
@@ -0,0 +1,18 @@
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, 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+#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_MISH_OP_H_
10+#define PTA_NPU_OP_API_INC_LEVEL0_OP_MISH_OP_H_
11+ 
12+#include "opdev/op_executor.h"
13+ 
14+namespace l0op {
15+const aclTensor* Mish(const aclTensor* self, aclOpExecutor* executor);
16+}
17+ 
18+#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_MISH_OP_H_
Aexperimental/activation/mish/op_kernel/mish.cpp+30-0
@@ -0,0 +1,30 @@
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, 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+ 
10+/*!
11+ * \file mish.cpp
12+ * \brief
13+*/
14+ 
15+#include "mish.h"
16+ 
17+enum class MishTilingKey : uint32_t
18+{
19+ TILING_KEY_EXAMPLE_FLOAT = 0,
20+ TILING_KEY_EXAMPLE_OTHER = 1,
21+};
22+template <uint32_t schMode>
23+__global__ __aicore__ void mish(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
24+{
25+ REGISTER_TILING_DEFAULT(MishTilingData);
26+ GET_TILING_DATA_WITH_STRUCT(MishTilingData, tilingData, tiling);
27+ MyMish::KernelMish<DTYPE_X,DTYPE_Y> op;
28+ op.Init(x, y, &tilingData);
29+ op.Process();
30+}
Aexperimental/activation/mish/op_kernel/mish.h+333-0
@@ -0,0 +1,333 @@
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, 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+ 
10+/*!
11+ * \file mish.h
12+ * \brief
13+ */
14+#ifndef MISH_H
15+#define MISH_H
16+ 
17+#include "kernel_operator.h"
18+#include "kernel_tiling/kernel_tiling.h"
19+#include "mish_tiling_data.h"
20+#include "mish_tiling_key.h"
21+ 
22+ 
23+namespace MyMish {
24+ 
25+using namespace AscendC;
26+ 
27+constexpr int32_t BUFFER_NUM = 2;
28+constexpr uint32_t COMPARE_ALIGN = 64;
29+ 
30+template <typename TYPE_X, typename TYPE_Y>
31+class KernelMish {
32+public:
33+ __aicore__ inline KernelMish(){};
34+ 
35+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const MishTilingData* tilingData);
36+ __aicore__ inline void Process();
37+ 
38+private:
39+ __aicore__ inline void CopyIn(int32_t progress);
40+ __aicore__ inline void CopyOut(int32_t progress);
41+ __aicore__ inline void Compute(int32_t progress);
42+ __aicore__ inline void ComputeSupPerfBf16(LocalTensor<TYPE_X> &xLocal, LocalTensor<TYPE_Y> &yLocal);
43+ __aicore__ inline void ComputeSupPerf(LocalTensor<TYPE_X> &xLocal, LocalTensor<TYPE_Y> &yLocal);
44+ __aicore__ inline void ComputeHighPerf16(LocalTensor<TYPE_X> &xLocal, LocalTensor<TYPE_Y> &yLocal);
45+ __aicore__ inline void ComputeHighPerf(LocalTensor<TYPE_X> &xLocal, LocalTensor<TYPE_Y> &yLocal);
46+ 
47+private:
48+ AscendC::TPipe pipe;
49+ AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX;
50+ AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueY;
51+ AscendC::TBuf<QuePosition::VECCALC> tmpBuffer1, tmpBuffer2, tmpBuffer3;
52+ AscendC::TBuf<QuePosition::VECCALC> QueueTmpX, QueueTmpY;
53+ AscendC::GlobalTensor<TYPE_X> xGm;
54+ AscendC::GlobalTensor<TYPE_Y> yGm;
55+ uint64_t coreDataNum;
56+ uint64_t tileNum;
57+ uint64_t tileDataNum;
58+ uint64_t tailDataNum;
59+ uint64_t processDataNum;
60+};
61+ 
62+template <typename TYPE_X, typename TYPE_Y>
63+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::Init(GM_ADDR x, GM_ADDR y, const MishTilingData* tilingData)
64+{
65+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
66+ uint64_t coreId = AscendC::GetBlockIdx();
67+ uint64_t globalBufferIndex = tilingData->bigCoreDataNum * coreId;
68+ this->tileDataNum = tilingData->tileDataNum;
69+ if (coreId < tilingData->tailBlockNum) {
70+ this->coreDataNum = tilingData->bigCoreDataNum;
71+ this->tileNum = tilingData->finalBigTileNum;
72+ this->tailDataNum = tilingData->bigTailDataNum;
73+ } else {
74+ this->coreDataNum = tilingData->smallCoreDataNum;
75+ this->tileNum = tilingData->finalSmallTileNum;
76+ this->tailDataNum = tilingData->smallTailDataNum;
77+ globalBufferIndex -= (tilingData->bigCoreDataNum - tilingData->smallCoreDataNum) * (coreId - tilingData->tailBlockNum);
78+ }
79+ xGm.SetGlobalBuffer((__gm__ TYPE_X*)x + globalBufferIndex, this->coreDataNum);
80+ yGm.SetGlobalBuffer((__gm__ TYPE_Y*)y + globalBufferIndex, this->coreDataNum);
81+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
82+ pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
83+
84+ pipe.InitBuffer(tmpBuffer1, this->tileDataNum * sizeof(float));
85+ pipe.InitBuffer(tmpBuffer2, this->tileDataNum * sizeof(float));
86+ pipe.InitBuffer(tmpBuffer3, this->tileDataNum * sizeof(uint8_t));
87+ pipe.InitBuffer(QueueTmpX, this->tileDataNum * sizeof(float));
88+ pipe.InitBuffer(QueueTmpY, this->tileDataNum * sizeof(float));
89+}
90+ 
91+template <typename TYPE_X, typename TYPE_Y>
92+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::CopyIn(int32_t progress)
93+{
94+ AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.AllocTensor<TYPE_X>();
95+ AscendC::DataCopy(xLocal, xGm[progress * this->tileDataNum], this->processDataNum);
96+ inQueueX.EnQue(xLocal);
97+}
98+ 
99+template <typename TYPE_X, typename TYPE_Y>
100+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::CopyOut(int32_t progress)
101+{
102+ AscendC::LocalTensor<TYPE_Y> yLocal = outQueueY.DeQue<TYPE_Y>();
103+ AscendC::DataCopy(yGm[progress * this->tileDataNum], yLocal, this->processDataNum);
104+ outQueueY.FreeTensor(yLocal);
105+}
106+ 
107+template <typename TYPE_X, typename TYPE_Y>
108+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::ComputeSupPerfBf16(LocalTensor<TYPE_X> &xLocal, LocalTensor<TYPE_Y> &yLocal)
109+{
110+ AscendC::LocalTensor<float> tmp1Local = tmpBuffer1.Get<float>();
111+ AscendC::LocalTensor<float> xLocalfp32 = QueueTmpX.Get<float>();
112+ AscendC::LocalTensor<float> yLocalfp32 = QueueTmpY.Get<float>();
113+ AscendC::Cast(xLocalfp32, xLocal, RoundMode::CAST_NONE, this->processDataNum);
114+ AscendC::PipeBarrier<PIPE_V>();
115+
116+ AscendC::Muls(yLocalfp32, xLocalfp32, float(0.015625), this->processDataNum);
117+ AscendC::PipeBarrier<PIPE_V>();
118+ AscendC::Adds(yLocalfp32, yLocalfp32, float(1), this->processDataNum);
119+ AscendC::PipeBarrier<PIPE_V>();
120+ AscendC::Mul(yLocalfp32, yLocalfp32, yLocalfp32, this->processDataNum);
121+ AscendC::PipeBarrier<PIPE_V>();
122+ AscendC::Mul(yLocalfp32, yLocalfp32, yLocalfp32, this->processDataNum);
123+ AscendC::PipeBarrier<PIPE_V>();
124+ AscendC::Mul(yLocalfp32, yLocalfp32, yLocalfp32, this->processDataNum);
125+ AscendC::Mul(yLocalfp32, yLocalfp32, yLocalfp32, this->processDataNum);
126+ AscendC::PipeBarrier<PIPE_V>();
127+ AscendC::Mul(yLocalfp32, yLocalfp32, yLocalfp32, this->processDataNum);
128+ AscendC::PipeBarrier<PIPE_V>();
129+ AscendC::Mul(yLocalfp32, yLocalfp32, yLocalfp32, this->processDataNum);
130+ AscendC::PipeBarrier<PIPE_V>();
131+ AscendC::Adds(yLocalfp32, yLocalfp32, float(1), this->processDataNum);
132+ AscendC::PipeBarrier<PIPE_V>();
133+ AscendC::Mul(yLocalfp32, yLocalfp32, yLocalfp32, this->processDataNum);
134+ AscendC::PipeBarrier<PIPE_V>();
135+ AscendC::Adds(yLocalfp32, yLocalfp32, float(1), this->processDataNum);
136+ AscendC::PipeBarrier<PIPE_V>();
137+ AscendC::Duplicate(tmp1Local, float(1), this->processDataNum);
138+ AscendC::PipeBarrier<PIPE_V>();
139+ AscendC::Div(yLocalfp32, tmp1Local, yLocalfp32, this->processDataNum);
140+ AscendC::PipeBarrier<PIPE_V>();
141+ AscendC::Muls(yLocalfp32, yLocalfp32, float(-2), this->processDataNum);
142+ AscendC::PipeBarrier<PIPE_V>();
143+ AscendC::Adds(yLocalfp32, yLocalfp32, float(1), this->processDataNum);
144+ AscendC::PipeBarrier<PIPE_V>();
145+ AscendC::Mul(yLocalfp32, xLocalfp32, yLocalfp32, this->processDataNum);
146+ AscendC::PipeBarrier<PIPE_V>();
147+ AscendC::Cast(yLocal, yLocalfp32, RoundMode::CAST_RINT, this->processDataNum);
148+ AscendC::PipeBarrier<PIPE_V>();
149+}
150+ 
151+template <typename TYPE_X, typename TYPE_Y>
152+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::ComputeSupPerf(LocalTensor<TYPE_X> &xLocal, LocalTensor<TYPE_Y> &yLocal)
153+{
154+ AscendC::LocalTensor<TYPE_X> tmp1Local = tmpBuffer1.Get<TYPE_X>();
155+ AscendC::Muls(yLocal, xLocal, TYPE_X(0.015625), this->processDataNum);
156+ AscendC::PipeBarrier<PIPE_V>();
157+ AscendC::Adds(yLocal, yLocal, TYPE_X(1), this->processDataNum);
158+ AscendC::PipeBarrier<PIPE_V>();
159+ AscendC::Mul(yLocal, yLocal, yLocal, this->processDataNum);
160+ AscendC::PipeBarrier<PIPE_V>();
161+ AscendC::Mul(yLocal, yLocal, yLocal, this->processDataNum);
162+ AscendC::PipeBarrier<PIPE_V>();
163+ AscendC::Mul(yLocal, yLocal, yLocal, this->processDataNum);
164+ AscendC::Mul(yLocal, yLocal, yLocal, this->processDataNum);
165+ AscendC::PipeBarrier<PIPE_V>();
166+ AscendC::Mul(yLocal, yLocal, yLocal, this->processDataNum);
167+ AscendC::PipeBarrier<PIPE_V>();
168+ AscendC::Mul(yLocal, yLocal, yLocal, this->processDataNum);
169+ AscendC::PipeBarrier<PIPE_V>();
170+ AscendC::Adds(yLocal, yLocal, TYPE_X(1), this->processDataNum);
171+ AscendC::PipeBarrier<PIPE_V>();
172+ AscendC::Mul(yLocal, yLocal, yLocal, this->processDataNum);
173+ AscendC::PipeBarrier<PIPE_V>();
174+ AscendC::Adds(yLocal, yLocal, TYPE_X(1), this->processDataNum);
175+ AscendC::PipeBarrier<PIPE_V>();
176+ AscendC::Duplicate(tmp1Local, TYPE_X(-2), this->processDataNum);
177+ AscendC::PipeBarrier<PIPE_V>();
178+ AscendC::Div(yLocal, tmp1Local, yLocal, this->processDataNum);
179+ AscendC::PipeBarrier<PIPE_V>();
180+ AscendC::Adds(yLocal, yLocal, TYPE_X(1), this->processDataNum);
181+ AscendC::PipeBarrier<PIPE_V>();
182+ AscendC::Mul(yLocal, xLocal, yLocal, this->processDataNum);
183+ AscendC::PipeBarrier<PIPE_V>();
184+}
185+ 
186+template <typename TYPE_X, typename TYPE_Y>
187+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::ComputeHighPerf16(LocalTensor<TYPE_X> &xLocal, LocalTensor<TYPE_Y> &yLocal)
188+{
189+ AscendC::LocalTensor<float> xLocalfp32 = QueueTmpX.Get<float>();
190+ AscendC::LocalTensor<float> yLocalfp32 = QueueTmpY.Get<float>();
191+ AscendC::LocalTensor<float> tmp1Local = tmpBuffer1.Get<float>();
192+ AscendC::LocalTensor<float> tmp2Local = tmpBuffer2.Get<float>();
193+ AscendC::LocalTensor<uint8_t> cmp1 = tmpBuffer3.Get<uint8_t>();
194+ 
195+ AscendC::Cast(xLocalfp32, xLocal, RoundMode::CAST_NONE, this->processDataNum);
196+ AscendC::PipeBarrier<PIPE_V>();
197+ uint32_t comparelength = (this->processDataNum + COMPARE_ALIGN - 1) / COMPARE_ALIGN * COMPARE_ALIGN;
198+ AscendC::CompareScalar(cmp1, xLocalfp32, float(0), CMPMODE::GT, comparelength);
199+ AscendC::PipeBarrier<PIPE_V>();
200+
201+ AscendC::Muls(tmp1Local , xLocalfp32, float(-1), this->processDataNum);
202+ AscendC::PipeBarrier<PIPE_V>();
203+ AscendC::Exp(tmp1Local, tmp1Local, this->processDataNum);
204+ AscendC::PipeBarrier<PIPE_V>();
205+ AscendC::Muls(tmp1Local, tmp1Local, float(2), this->processDataNum);
206+ AscendC::PipeBarrier<PIPE_V>();
207+ AscendC::Adds(tmp1Local, tmp1Local, float(1), this->processDataNum);
208+ AscendC::PipeBarrier<PIPE_V>();
209+ AscendC::Muls(tmp2Local , xLocalfp32, float(-2), this->processDataNum);
210+ AscendC::PipeBarrier<PIPE_V>();
211+ AscendC::Exp(tmp2Local, tmp2Local, this->processDataNum);
212+ AscendC::PipeBarrier<PIPE_V>();
213+ AscendC::Muls(tmp2Local, tmp2Local, float(2), this->processDataNum);
214+ AscendC::PipeBarrier<PIPE_V>();
215+ AscendC::Add(tmp2Local, tmp1Local, tmp2Local, this->processDataNum);
216+ AscendC::PipeBarrier<PIPE_V>();
217+ AscendC::Div(tmp1Local, tmp1Local, tmp2Local, this->processDataNum); // res for x > 0
218+ AscendC::PipeBarrier<PIPE_V>();
219+
220+ AscendC::Muls(tmp2Local , xLocalfp32, float(2), this->processDataNum);
221+ AscendC::PipeBarrier<PIPE_V>();
222+ AscendC::Exp(tmp2Local, tmp2Local, this->processDataNum);
223+ AscendC::PipeBarrier<PIPE_V>();
224+ AscendC::Exp(yLocalfp32, xLocalfp32, this->processDataNum);
225+ AscendC::PipeBarrier<PIPE_V>();
226+ AscendC::Muls(yLocalfp32, yLocalfp32, float(2), this->processDataNum);
227+ AscendC::PipeBarrier<PIPE_V>();
228+ AscendC::Add(yLocalfp32, yLocalfp32, tmp2Local, this->processDataNum);
229+ AscendC::PipeBarrier<PIPE_V>();
230+ AscendC::Adds(tmp2Local, yLocalfp32, float(2), this->processDataNum);
231+ AscendC::PipeBarrier<PIPE_V>();
232+ AscendC::Div(yLocalfp32, yLocalfp32, tmp2Local, this->processDataNum); //res for x <= 0
233+ AscendC::PipeBarrier<PIPE_V>();
234+
235+ AscendC::Select(yLocalfp32, cmp1, tmp1Local, yLocalfp32, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataNum);
236+ AscendC::Mul(yLocalfp32, yLocalfp32, xLocalfp32, this->processDataNum);
237+ AscendC::PipeBarrier<PIPE_V>();
238+ 
239+ AscendC::Cast(yLocal, yLocalfp32, RoundMode::CAST_RINT, this->processDataNum);
240+ AscendC::PipeBarrier<PIPE_V>();
241+}
242+ 
243+template <typename TYPE_X, typename TYPE_Y>
244+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::ComputeHighPerf(LocalTensor<TYPE_X> &xLocal, LocalTensor<TYPE_Y> &yLocal)
245+{
246+ AscendC::LocalTensor<TYPE_X> tmp1Local = tmpBuffer1.Get<TYPE_X>();
247+ AscendC::LocalTensor<TYPE_X> tmp2Local = tmpBuffer2.Get<TYPE_X>();
248+ AscendC::LocalTensor<uint8_t> cmp1 = tmpBuffer3.Get<uint8_t>();
249+ 
250+ uint32_t comparelength = (this->processDataNum + COMPARE_ALIGN - 1) / COMPARE_ALIGN * COMPARE_ALIGN;
251+ AscendC::CompareScalar(cmp1, xLocal, TYPE_X(0), CMPMODE::GT, comparelength);
252+ AscendC::PipeBarrier<PIPE_V>();
253+ AscendC::Muls(tmp1Local , xLocal, TYPE_X(-1), this->processDataNum);
254+ AscendC::PipeBarrier<PIPE_V>();
255+ AscendC::Exp(tmp1Local, tmp1Local, this->processDataNum);
256+ AscendC::PipeBarrier<PIPE_V>();
257+ AscendC::Muls(tmp1Local, tmp1Local, TYPE_X(2), this->processDataNum);
258+ AscendC::PipeBarrier<PIPE_V>();
259+ AscendC::Adds(tmp1Local, tmp1Local, TYPE_X(1), this->processDataNum);
260+ AscendC::PipeBarrier<PIPE_V>();
261+ AscendC::Muls(tmp2Local , xLocal, TYPE_X(-2), this->processDataNum);
262+ AscendC::PipeBarrier<PIPE_V>();
263+ AscendC::Exp(tmp2Local, tmp2Local, this->processDataNum);
264+ AscendC::PipeBarrier<PIPE_V>();
265+ AscendC::Muls(tmp2Local, tmp2Local, TYPE_X(2), this->processDataNum);
266+ AscendC::PipeBarrier<PIPE_V>();
267+ AscendC::Add(tmp2Local, tmp1Local, tmp2Local, this->processDataNum);
268+ AscendC::PipeBarrier<PIPE_V>();
269+ AscendC::Div(tmp1Local, tmp1Local, tmp2Local, this->processDataNum); // res for x > 0
270+ AscendC::PipeBarrier<PIPE_V>();
271+
272+ AscendC::Muls(tmp2Local , xLocal, TYPE_X(2), this->processDataNum);
273+ AscendC::PipeBarrier<PIPE_V>();
274+ AscendC::Exp(tmp2Local, tmp2Local, this->processDataNum);
275+ AscendC::PipeBarrier<PIPE_V>();
276+ AscendC::Exp(yLocal, xLocal, this->processDataNum);
277+ AscendC::PipeBarrier<PIPE_V>();
278+ AscendC::Muls(yLocal, yLocal, TYPE_X(2), this->processDataNum);
279+ AscendC::PipeBarrier<PIPE_V>();
280+ AscendC::Add(yLocal, yLocal, tmp2Local, this->processDataNum);
281+ AscendC::PipeBarrier<PIPE_V>();
282+ AscendC::Adds(tmp2Local, yLocal, TYPE_X(2), this->processDataNum);
283+ AscendC::PipeBarrier<PIPE_V>();
284+ AscendC::Div(yLocal, yLocal, tmp2Local, this->processDataNum); //res for x <= 0
285+ AscendC::PipeBarrier<PIPE_V>();
286+
287+ AscendC::Select(yLocal, cmp1, tmp1Local, yLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, this->processDataNum);
288+ AscendC::Mul(yLocal, yLocal, xLocal, this->processDataNum);
289+ AscendC::PipeBarrier<PIPE_V>();
290+}
291+ 
292+template <typename TYPE_X, typename TYPE_Y>
293+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::Compute(int32_t progress)
294+{
295+ AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.DeQue<TYPE_X>();
296+ AscendC::LocalTensor<TYPE_Y> yLocal = outQueueY.AllocTensor<TYPE_Y>();
297+ 
298+ #if defined(SUPER_PERFORMANCE) && SUPER_PERFORMANCE == 1
299+ if constexpr (std::is_same_v<TYPE_X, __bf16>) {
300+ ComputeSupPerfBf16(xLocal, yLocal);
301+ } else if constexpr (std::is_same_v<TYPE_X, float> || std::is_same_v<TYPE_X, half>){
302+ ComputeSupPerf(xLocal, yLocal);
303+ }
304+ #else
305+ if constexpr (std::is_same_v<TYPE_X, __bf16> || std::is_same_v<TYPE_X, half>) {
306+ ComputeHighPerf16(xLocal, yLocal);
307+ } else if constexpr (std::is_same_v<TYPE_X, float>) {
308+ ComputeHighPerf(xLocal, yLocal);
309+ }
310+ #endif
311+
312+ outQueueY.EnQue<TYPE_Y>(yLocal);
313+ inQueueX.FreeTensor(xLocal);
314+}
315+ 
316+template <typename TYPE_X, typename TYPE_Y>
317+__aicore__ inline void KernelMish<TYPE_X, TYPE_Y>::Process()
318+{
319+ int32_t loopCount = this->tileNum;
320+ this->processDataNum = this->tileDataNum;
321+ for (int32_t i = 0; i < loopCount - 1; i++) {
322+ CopyIn(i);
323+ Compute(i);
324+ CopyOut(i);
325+ }
326+ this->processDataNum = this->tailDataNum;
327+ CopyIn(loopCount - 1);
328+ Compute(loopCount - 1);
329+ CopyOut(loopCount - 1);
330+}
331+ 
332+} // namespace MyMish
333+#endif // MISH_H
Aexperimental/activation/mish/op_kernel/mish_tiling_data.h+26-0
@@ -0,0 +1,26 @@
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, 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+ 
10+/*!
11+ * \file mish_tiling_data.h
12+ * \brief tiling data struct
13+*/
14+#ifndef MISH_TILING_DATA_H_
15+#define MISH_TILING_DATA_H_
16+struct MishTilingData{
17+ uint64_t smallCoreDataNum;
18+ uint64_t bigCoreDataNum;
19+ uint64_t finalBigTileNum;
20+ uint64_t finalSmallTileNum;
21+ uint64_t tileDataNum;
22+ uint64_t smallTailDataNum;
23+ uint64_t bigTailDataNum;
24+ uint64_t tailBlockNum;
25+} ;
26+#endif
Aexperimental/activation/mish/op_kernel/mish_tiling_key.h+27-0
@@ -0,0 +1,27 @@
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, 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+ 
10+/*!
11+ * \file mish_tiling_key.h
12+ * \brief mish tiling key declare
13+*/
14+#include "ascendc/host_api/tiling/template_argument.h"
15+ 
16+#define ELEMENTWISE_TPL_SCH_MODE_0 0
17+#define ELEMENTWISE_TPL_SCH_MODE_1 1
18+ 
19+ASCENDC_TPL_ARGS_DECL(
20+ Mish,
21+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
22+ 
23+ASCENDC_TPL_SEL(
24+ ASCENDC_TPL_ARGS_SEL(
25+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,ELEMENTWISE_TPL_SCH_MODE_0,ELEMENTWISE_TPL_SCH_MODE_1)
26+ ),
27+);
Aexperimental/activation/mish/tests/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 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").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9+# the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aexperimental/activation/mish/tests/ut/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 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").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9+# the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aexperimental/activation/mish/tests/ut/op_host/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
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+ 
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_HOST_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_TILING_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})
15+endif()
16+if(UT_TEST_ALL OR OP_API_UT)
17+ add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
18+endif()
Aexperimental/activation/mish/tests/ut/op_host/test_mish_tiling.cpp+279-0
@@ -0,0 +1,279 @@
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+#include <iostream>
12+#include <fstream>
13+#include <sstream>
14+#include <string>
15+#include <map>
16+#include <vector>
17+#include <gtest/gtest.h>
18+#include "../../../op_kernel/mish_tiling_data.h"
19+#include "log/log.h"
20+#include "ut_op_common.h"
21+ 
22+#include "ut_op_util.h"
23+#include "platform/platform_infos_def.h"
24+#include "register/op_impl_registry.h"
25+#include "platform/platform_info.h"
26+#include "test_cube_util.h"
27+#include "exe_graph/runtime/storage_format.h"
28+#include "exe_graph/runtime/storage_shape.h"
29+#include "kernel_run_context_facker.h"
30+#include "array_ops.h"
31+ 
32+ 
33+using namespace ut_util;
34+using namespace std;
35+using namespace ge;
36+ 
37+class MishTilingTest : public testing::Test
38+{
39+protected:
40+ static void SetUpTestCase()
41+ {
42+ std::cout << "MishTilingTest Setup" << std::endl;
43+ }
44+ static void TearDownTestCase()
45+ {
46+ std::cout << "MishTilingTest TearDown" << std::endl;
47+ }
48+};
49+ 
50+TEST_F(MishTilingTest, mish_001)
51+{
52+ gert::StorageShape x_shape = {{2,2}, {2,2}};
53+ gert::StorageShape output_shapes = {{2,2}, {2,2}};
54+ 
55+ string compile_info_string = R"({
56+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
57+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
58+ "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
59+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
60+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
61+ "CORE_NUM": 40}
62+ })";
63+ map<string, string> soc_infos;
64+ map<string, string> aicore_spec;
65+ map<string, string> intrinsics;
66+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
67+ 
68+ fe::PlatFormInfos platform_info;
69+ platform_info.Init();
70+ 
71+ struct MishCompileInfo {
72+ };
73+ MishCompileInfo compile_info;
74+ 
75+ std::string op_type("Mish");
76+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
77+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
78+ auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
79+ 
80+ // tilingParseFunc simulate
81+ auto kernel_holder =
82+ gert::KernelRunContextFaker()
83+ .KernelIONum(1, 1)
84+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
85+ .Outputs({&compile_info})
86+ .Build();
87+ 
88+ ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
89+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
90+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
91+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
92+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
93+ "AICoreintrinsicDtypeMap", intrinsics);
94+ ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
95+ 
96+ // tilingFunc simulate
97+ 
98+ auto param = gert::TilingData::CreateCap(4096);
99+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
100+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
101+ ASSERT_NE(param, nullptr);
102+
103+ auto holder = gert::TilingContextFaker()
104+ .NodeIoNum(1, 1)
105+ .IrInstanceNum({1})
106+ .InputShapes({&x_shape})
107+ .OutputShapes({&output_shapes})
108+ .CompileInfo(&compile_info)
109+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
110+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
111+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
112+ .TilingData(param.get())
113+ .Workspace(ws_size)
114+ .Build();
115+ 
116+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
117+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
118+ 
119+ tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
120+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
121+ tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
122+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
123+ 
124+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
125+}
126+ 
127+TEST_F(MishTilingTest, mish_002)
128+{
129+ gert::StorageShape x_shape = {{267,54}, {267,54}};
130+ gert::StorageShape output_shapes = {{267,54}, {267,54}};
131+ 
132+ string compile_info_string = R"({
133+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
134+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
135+ "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
136+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
137+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
138+ "CORE_NUM": 40}
139+ })";
140+ map<string, string> soc_infos;
141+ map<string, string> aicore_spec;
142+ map<string, string> intrinsics;
143+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
144+ 
145+ fe::PlatFormInfos platform_info;
146+ platform_info.Init();
147+ 
148+ struct MishCompileInfo {
149+ };
150+ MishCompileInfo compile_info;
151+ 
152+ std::string op_type("Mish");
153+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
154+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
155+ auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
156+ 
157+ // tilingParseFunc simulate
158+ auto kernel_holder =
159+ gert::KernelRunContextFaker()
160+ .KernelIONum(1, 1)
161+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
162+ .Outputs({&compile_info})
163+ .Build();
164+ 
165+ ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
166+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
167+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
168+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
169+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
170+ "AICoreintrinsicDtypeMap", intrinsics);
171+ ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
172+ 
173+ // tilingFunc simulate
174+ 
175+ auto param = gert::TilingData::CreateCap(4096);
176+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
177+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
178+ ASSERT_NE(param, nullptr);
179+
180+ auto holder = gert::TilingContextFaker()
181+ .NodeIoNum(1, 1)
182+ .IrInstanceNum({1})
183+ .InputShapes({&x_shape})
184+ .OutputShapes({&output_shapes})
185+ .CompileInfo(&compile_info)
186+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
187+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
188+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
189+ .TilingData(param.get())
190+ .Workspace(ws_size)
191+ .Build();
192+ 
193+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
194+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
195+
196+ tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
197+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
198+ tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
199+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
200+ 
201+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
202+}
203+ 
204+TEST_F(MishTilingTest, mish_003)
205+{
206+ gert::StorageShape x_shape = {{267,354}, {267,354}};
207+ gert::StorageShape output_shapes = {{267,354}, {267,354}};
208+ 
209+ string compile_info_string = R"({
210+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
211+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
212+ "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
213+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
214+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
215+ "CORE_NUM": 40}
216+ })";
217+ map<string, string> soc_infos;
218+ map<string, string> aicore_spec;
219+ map<string, string> intrinsics;
220+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
221+ 
222+ fe::PlatFormInfos platform_info;
223+ platform_info.Init();
224+ 
225+ struct MishCompileInfo {
226+ };
227+ MishCompileInfo compile_info;
228+ 
229+ std::string op_type("Mish");
230+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
231+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
232+ auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
233+ 
234+ // tilingParseFunc simulate
235+ auto kernel_holder =
236+ gert::KernelRunContextFaker()
237+ .KernelIONum(1, 1)
238+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
239+ .Outputs({&compile_info})
240+ .Build();
241+ 
242+ ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
243+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
244+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
245+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
246+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
247+ "AICoreintrinsicDtypeMap", intrinsics);
248+ ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
249+ 
250+ // tilingFunc simulate
251+ 
252+ auto param = gert::TilingData::CreateCap(4096);
253+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
254+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
255+ ASSERT_NE(param, nullptr);
256+
257+ auto holder = gert::TilingContextFaker()
258+ .NodeIoNum(1, 1)
259+ .IrInstanceNum({1})
260+ .InputShapes({&x_shape})
261+ .OutputShapes({&output_shapes})
262+ .CompileInfo(&compile_info)
263+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
264+ .NodeInputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
265+ .NodeOutputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
266+ .TilingData(param.get())
267+ .Workspace(ws_size)
268+ .Build();
269+ 
270+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
271+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
272+
273+ tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
274+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
275+ tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
276+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
277+ 
278+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
279+}
Aexperimental/activation/mish/tests/ut/op_kernel/CMakeLists.txt+29-0
@@ -0,0 +1,29 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify it.
3+# Copyright (c) 2025 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").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
9+# the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+if (UT_TEST_ALL OR OP_KERNEL_UT)
13+ # 需要将Tiling依赖的文件添加到CMakeLists.txt中
14+ # set(elewise_common_tiling_files
15+ # ${CANN_ROOT}/ops/built-in/op_tiling/runtime/elewise_tiling.cc
16+ # )
17+ # 算子自己的tiling文件路径
18+ set(mish_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/mish_tiling.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/mish_infershape.cpp
21+ # ${elewise_common_tiling_files}
22+ )
23+ # 使用AddOpTestCase
24+ # param1：算子名称，以kernel方式命名
25+ # param2：soc版本，多个以分号分隔，例如："ascend910_9599;AscendB1"
26+ # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔，例如："-DDTYPE_X=float -DDTYPE_Y=float"
27+ # param4：该算子依赖的所有tiling源码文件
28+ # # AddOpTestCase(mish "ascend910B1" "" "${mish_tiling_files}")
condfuse_3
condfuse_3condfuse_32025年12月25日

把注释去掉

likedislike
29+endif()
Aexperimental/activation/mish/tests/ut/op_kernel/mish_data/compare_data.py+57-0
@@ -0,0 +1,57 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# This program is free software, you can redistribute it and/or modify it.
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6+# This file is a part of the CANN Open Software.
7+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
11+# the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------
13+ 
14+import sys
15+import numpy as np
16+import glob
17+import os
18+ 
19+curr_dir = os.path.dirname(os.path.realpath(__file__))
20+ 
21+def compare_data(golden_file_lists, output_file_lists, d_type):
22+ if d_type == "float16":
23+ np_dtype = np.float16
24+ elif d_type == "float32":
25+ np_dtype = np.float32
26+ else:
27+ raise ValueError("d_type must be float16 or float32")
28+
29+ data_same = True
30+ for gold, out in zip(golden_file_lists, output_file_lists):
31+ tmp_out = np.fromfile(out, np_dtype)
32+ tmp_gold = np.fromfile(gold, np_dtype)
33+ diff_res = np.isclose(tmp_out, tmp_gold, 0, 0, True)
34+ diff_idx = np.where(diff_res != True)[0]
35+ if len(diff_idx) == 0:
36+ print("PASSED!")
37+ else:
38+ print("FAILED!")
39+ for idx in diff_idx[:5]:
40+ print(f"index: {idx}, output: {tmp_out[idx]}, golden: {tmp_gold[idx]}")
41+ data_same = False
42+ return data_same
43+ 
44+def get_file_lists(dtype):
45+ golden_file_lists = sorted(glob.glob(curr_dir + "/*golden*.bin"))
46+ output_file_lists = sorted(glob.glob(curr_dir + "/*output*.bin"))
47+ return golden_file_lists, output_file_lists
48+ 
49+def process(d_type):
50+ golden_file_lists, output_file_lists = get_file_lists(d_type)
51+ result = compare_data(golden_file_lists, output_file_lists, d_type)
52+ print("compare result:", result)
53+ return result
54+ 
55+if __name__ == '__main__':
56+ ret = process(sys.argv[1])
57+ exit(0 if ret else 1)
Aexperimental/activation/mish/tests/ut/op_kernel/mish_data/gen_data.py+55-0
@@ -0,0 +1,55 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# This program is free software, you can redistribute it and/or modify it.
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6+# This file is a part of the CANN Open Software.
7+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
11+# the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------
13+ 
14+import sys
15+import os
16+import numpy as np
17+import re
18+import tensorflow as tf
19+ 
20+ 
21+def parse_str_to_shape_list(shape_str):
22+ shape_str = shape_str.strip('(').strip(')')
23+ shape_list = [int(x) for x in shape_str.split(",")]
24+ return np.array(shape_list)
25+ 
26+def mish(x):
27+ return x * tf.tanh(tf.math.softplus(x))
28+ 
29+def gen_data_and_golden(shape_str, d_type="float32"):
30+ d_type_dict = {
31+ "float32": np.float32,
32+ "float16": np.float16,
33+ "bfloat16": tf.bfloat16.as_numpy_dtype
34+ }
35+ np_type = d_type_dict[d_type]
36+ shape = parse_str_to_shape_list(shape_str)
37+ size = np.prod(shape)
38+ tmp_input = np.random.choice([0, 0.5, 1, 65504, np.nan, np.inf], size=size)
39+ tmp_input = tmp_input.reshape(shape).astype(np_type)
40+ tmp_golden = mish(tmp_input)
41+ 
42+ tmp_input.astype(np_type).tofile(f"{d_type}_input_t_mish.bin")
43+ tmp_golden.astype(np_type).tofile(f"{d_type}_golden_t_mish.bin")
44+ 
45+ 
46+if __name__ == "__main__":
47+ if len(sys.argv) != 3:
48+ print("Param num must be 3.")
49+ exit(1)
50+ # 清理bin文件
51+ os.system("rm -rf *.bin")
52+ gen_data_and_golden(sys.argv[1], sys.argv[2])
53+ 
54+ 
55+ 
Aexperimental/activation/mish/tests/ut/op_kernel/test_mish.cpp+146-0
@@ -0,0 +1,146 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify it.
3+ * Copyright (c) 2025 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").
6+ * Please refer to the License for details. You may not use this file except in compliance with the License.
7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file test_mish.cpp
14+ * \brief
15+ */
16+ 
17+#include <array>
18+#include <vector>
19+#include <iostream>
20+#include <string>
21+#include <cstdint>
22+#include "gtest/gtest.h"
23+#include "tikicpulib.h"
24+#include "data_utils.h"
25+#include "tiling_case_executor.h"
26+#include "../op_host/mish_tiling.h"
27+ 
28+using namespace std;
29+ 
30+extern "C" __global__ __aicore__ void mish(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
31+ 
32+class MishTest : public testing::Test {
33+protected:
34+ static void SetUpTestCase()
35+ {
36+ std::cout << "mish_test SetUp" << std::endl;
37+ const string cmd = "cp -rf " + dataPath + " ./";
38+ system(cmd.c_str());
39+ system("chmod -R 755 ./mish_data/");
40+ }
41+ static void TearDownTestCase()
42+ {
43+ std::cout << "mish_test TearDown" << std::endl;
44+ }
45+ 
46+private:
47+ const static std::string rootPath;
48+ const static std::string dataPath;
49+};
50+ 
51+const std::string MishTest::rootPath = "../../../../";
52+const std::string MishTest::dataPath = rootPath + "math/mish/tests/ut/op_kernel/mish_data";
53+ 
54+template <typename T1, typename T2>
55+inline T1 CeilAlign(T1 a, T2 b)
56+{
57+ return (a + b - 1) / b * b;
58+}
59+ 
60+TEST_F(MishTest, test_case_float16_1)
61+{
62+ optiling::MishCompileInfo compileInfo = {64, 262144, false};
63+ gert::TilingContextPara tilingContextPara(
64+ "Mish",
65+ {
66+ {{{128, 64}, {128, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
67+ },
68+ {
69+ {{{128, 64}, {128, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND},
70+ },
71+ &compileInfo);
72+ TilingInfo tilingInfo;
73+ auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
74+ EXPECT_EQ(tilingRet, true);
75+ 
76+ system("cd ./mish_data/ && python3 gen_data.py '(128, 64)' 'float16'");
77+ uint32_t dataCount = 128 * 64;
78+ size_t inputByteSize = dataCount * sizeof(half);
79+ std::string fileName = "./mish_data/float16_input_t_mish.bin";
80+ ;
81+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
82+ ReadFile(fileName, inputByteSize, x, inputByteSize);
83+ size_t outputByteSize = dataCount * sizeof(half);
84+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
85+ 
86+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(tilingInfo.workspaceSizes[0]);
87+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingInfo.tilingDataSize);
88+ std::memcpy(tiling, tilingInfo.tilingData.get(), tilingInfo.tilingDataSize);
89+ ICPU_SET_TILING_KEY(tilingInfo.tilingKey);
90+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
91+ ICPU_RUN_KF(mish, tilingInfo.blockNum, x, y, workspace, tiling);
92+ 
93+ fileName = "./mish_data/float16_output_t_mish.bin";
94+ WriteFile(fileName, y, outputByteSize);
95+ 
96+ AscendC::GmFree((void*)(x));
97+ AscendC::GmFree((void*)(y));
98+ AscendC::GmFree((void*)workspace);
99+ AscendC::GmFree((void*)tiling);
100+ 
101+ system("cd ./mish_data/ && python3 compare_data.py 'float16'");
102+}
103+ 
104+TEST_F(MishTest, test_case_float32_1)
105+{
106+ optiling::MishCompileInfo compileInfo = {64, 262144, false};
107+ gert::TilingContextPara tilingContextPara(
108+ "Mish",
109+ {
110+ {{{256, 33}, {256, 33}}, ge::DT_FLOAT, ge::FORMAT_ND},
111+ },
112+ {
113+ {{{256, 33}, {256, 33}}, ge::DT_FLOAT, ge::FORMAT_ND},
114+ },
115+ &compileInfo);
116+ TilingInfo tilingInfo;
117+ auto tilingRet = ExecuteTiling(tilingContextPara, tilingInfo);
118+ EXPECT_EQ(tilingRet, true);
119+ 
120+ system("cd ./mish_data/ && python3 gen_data.py '(256, 33)' 'float32'");
121+ uint32_t dataCount = 256 * 33;
122+ size_t inputByteSize = dataCount * sizeof(float);
123+ std::string fileName = "./mish_data/float32_input_t_mish.bin";
124+ ;
125+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
126+ ReadFile(fileName, inputByteSize, x, inputByteSize);
127+ size_t outputByteSize = dataCount * sizeof(float);
128+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
129+ 
130+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(tilingInfo.workspaceSizes[0]);
131+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingInfo.tilingDataSize);
132+ std::memcpy(tiling, tilingInfo.tilingData.get(), tilingInfo.tilingDataSize);
133+ ICPU_SET_TILING_KEY(tilingInfo.tilingKey);
134+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
135+ ICPU_RUN_KF(mish, tilingInfo.blockNum, x, y, workspace, tiling);
136+ 
137+ fileName = "./mish_data/float32_output_t_mish.bin";
138+ WriteFile(fileName, y, outputByteSize);
139+ 
140+ AscendC::GmFree((void*)(x));
141+ AscendC::GmFree((void*)(y));
142+ AscendC::GmFree((void*)workspace);
143+ AscendC::GmFree((void*)tiling);
144+ 
145+ system("cd ./mish_data/ && python3 compare_data.py 'float32'");
146+}