Star1003
1932
代码
代码
Issues249
Pull Requests361
Actions
项目讨论
Wiki
分析
Star1003
1932
已合并
下一代activation，loss支持ctc_loss_v2_grad,fast_gelu,fast_gelu_grad #506
yuanqi_li创建于 2025年12月25日
下一代activation，loss支持ctc_loss_v2_grad,fast_gelu,fast_gelu_grad #506
已合并
yuanqi_li创建于 2025年12月25日
从yuanqi_li/ops-nnmigration1合入到cann/ops-nnmaster
共 77 个文件变更+5658-1768
Mactivation/fast_gelu/CMakeLists.txt+8-11
@@ -1,3 +1,4 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
@@ -5,15 +6,11 @@
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.fast_gelu.CMakeLists.txt ")11+# 设置算子定义时支持的芯片类型
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)12+set(SUPPORT_COMPUTE_UNIT "ascend910_95")
12-if(NOT ENABLE_TEST AND NOT BENCHMARK)13+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
13- list(REMOVE_ITEM CURRENT_DIRS tests)14+set(SUPPORT_TILING_DIR "arch35")
14-endif()15+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE fast_gelu ACLNNTYPE aclnn_exclude
15-foreach(SUB_DIR ${CURRENT_DIRS})16+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
16- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17- add_subdirectory(${SUB_DIR})
18- endif()
19-endforeach()
Mactivation/fast_gelu/README.md+56-1
@@ -1,3 +1,58 @@
1# FastGelu1# FastGelu
2 2 
3-本目录仅包含FastGelu算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT9</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+ 
10+## 功能说明
11+ 
12+- 算子功能：
13+ 输入fast_gelu的计算结果。
14+- 计算公式：
15+ $$ FastGelu(x_i) = \frac {x_i} {1 + \exp(-1.702 * \left| x_i \right|)} * \exp(0851 * (x_i - \left| x_i \right|)) $$
16+
17+## 参数说明
18+ 
19+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
20+ <col style="width: 100px">
21+ <col style="width: 150px">
22+ <col style="width: 280px">
23+ <col style="width: 330px">
24+ <col style="width: 120px">
25+ </colgroup>
26+ <thead>
27+ <tr>
28+ <th>参数名</th>
29+ <th>输入/输出/属性</th>
30+ <th>描述</th>
31+ <th>数据类型</th>
32+ <th>数据格式</th>
33+ </tr></thead>
34+ <tbody>
35+ <tr>
36+ <td>x</td>
37+ <td>输入</td>
38+ <td>输入一个张量。</td>
39+ <td>DT_BF16、FLOAT16、DT_FLOAT</td>
40+ <td>ND</td>
41+ </tr>
42+ <tr>
43+ <td>y</td>
44+ <td>输出</td>
45+ <td>输出一个张量。</td>
46+ <td>DT_BF16、FLOAT16、DT_FLOAT</td>
47+ <td>ND</td>
48+ </tr>
49+ </tbody></table>
50+ 
51+## 约束说明
52+ 
53+无
54+ 
55+## 调用说明
56+| 调用方式 | 调用样例 | 说明 |
57+|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
58+| aclnn调用 | [test_aclnn_fast_gelu](examples/arch35/test_aclnn_fast_gelu.cpp) | 通过[aclnnFastGeluGetWorkspaceSize](docs/aclnnFastGelu.md)接口方式调用fast_gelu算子。 |
Mactivation/fast_gelu/docs/aclnnFastGelu.md+136-326
@@ -1,326 +1,136 @@
1-# aclnnFastGelu1+# FastGelu
2- 2+ 
3-## 产品支持情况3+## 产品支持情况
4- 4+ 
5-|产品 | 是否支持 |5+| 产品 | 是否支持 |
6-|:-------------------------|:----------:|6+| :----------------------------------------------------------- | :------: |
7-| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |7+| <term>Ascend 950PR/Ascend 950DT9</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9- 9+ 
10-## 功能说明10+## 功能说明
11-快速高斯误差线性单元激活函数。11+ 
12-FastGelu定义如下：12+算子功能：返回fastgelu激活函数输出张量。
13- 13+ 
14-$$14+## 函数原型
15-FastGelu(x_i) = \frac {x_i} {1 + e^{-1.702 x_i}}15+ 
16-$$ 16+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnFastGeluGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnFastGelu”接口执行计算。
17- 17+ 
18-其中$x_{i}$是输入的元素。18+## aclnnFastGeluGetWorkspaceSize
19- 19+ 
20-## 函数原型20+- **参数说明：**
21- 21+ 
22-每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnFastGeluGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnFastGelu”接口执行计算。22+ <table style="undefined;table-layout: fixed; width: 1420px"><colgroup>
23- 23+ <col style="width: 173px">
24-```Cpp24+ <col style="width: 120px">
25-aclnnStatus aclnnFastGeluGetWorkspaceSize(25+ <col style="width: 222px">
26- const aclTensor *self,26+ <col style="width: 338px">
27- aclTensor *out,27+ <col style="width: 156px">
28- uint64_t *workspaceSize,28+ <col style="width: 104px">
29- aclOpExecutor **executor)29+ <col style="width: 162px">
30-```30+ <col style="width: 145px">
31- 31+ </colgroup>
32-```Cpp32+ <thead>
33-aclnnStatus aclnnFastGelu(33+ <tr>
34- void *workspace,34+ <th>参数名</th>
35- uint64_t workspaceSize,35+ <th>输入/输出</th>
36- aclOpExecutor *executor,36+ <th>描述</th>
37- aclrtStream stream)37+ <th>使用说明</th>
38-```38+ <th>数据类型</th>
39- 39+ <th>数据格式</th>
40-## aclnnFastGeluGetWorkspaceSize40+ <th>维度(shape)</th>
41- 41+ <th>非连续Tensor</th>
42-- **参数说明：**42+ </tr></thead>
43- 43+ <tbody>
44- <table style="undefined;table-layout: fixed; width: 1420px"><colgroup>44+ <tr>
45- <col style="width: 171px">45+ <td>self</td>
46- <col style="width: 115px">46+ <td>输入</td>
47- <col style="width: 220px">47+ <td>输入张量。</td>
48- <col style="width: 250px">48+ <td>数据类型支持FLOAT16、FLOAT32、BFLOAT16</td>
49- <col style="width: 177px">49+ <td>FLOAT16、FLOAT32、BFLOAT16</td>
50- <col style="width: 104px">50+ <td>ND</td>
51- <col style="width: 238px">51+ <td>(N)</td>
52- <col style="width: 145px">52+ <td>√</td>
53- </colgroup>53+ </tr>
54- <thead>54+ <tr>
55- <tr>55+ <td>out</td>
56- <th>参数名</th>56+ <td>输出</td>
57- <th>输入/输出</th>57+ <td>输出张量。</td>
58- <th>描述</th>58+ <td>数据类型支持FLOAT16、FLOAT32、BFLOAT16</td>
59- <th>使用说明</th>59+ <td>FLOAT16、FLOAT32、BFLOAT16</td>
60- <th>数据类型</th>60+ <td>ND</td>
61- <th>数据格式</th>61+ <td>(N)</td>
62- <th>维度(shape)</th>62+ <td>√</td>
63- <th>非连续Tensor</th>63+ </tr>
64- </tr></thead>64+ <tr>
65- <tbody>65+ <td>workspaceSize</td>
66- <tr>66+ <td>输出</td>
67- <td>self</td>67+ <td>返回需要在Device侧申请的workspace大小。</td>
68- <td>输入</td>68+ <td>-</td>
69- <td>表示输入张量，公式中的x_{i}。</td>69+ <td>-</td>
70- <td><ul><li>数据格式需要与out一致。</li><li>shape需要与out一致。</li><li>数据类型需要与out一致。</li></ul></td>70+ <td>-</td>
71- <td>FLOAT16、FLOAT32、BFLOAT16</td>71+ <td>-</td>
72- <td>ND</td>72+ <td>-</td>
73- <td>0-8</td>73+ </tr>
74- <td>√</td>74+ <tr>
75- </tr>75+ <td>executor</td>
76- <tr>76+ <td>输出</td>
77- <td>out</td>77+ <td>返回需要在Device侧申请的workspace大小。</td>
78- <td>输出</td>78+ <td>-</td>
79- <td>表示输出张量，公式中的FastGelu(x_i)。</td>79+ <td>-</td>
80- <td><ul><li>数据格式需要与self一致。</li><li>shape需要与self一致。</li><li>数据格式需要与self一致。</li></ul></td>80+ <td>-</td>
81- <td>FLOAT16、FLOAT32、BFLOAT16</td>81+ <td>-</td>
82- <td>ND</td>82+ <td>-</td>
83- <td>0-8</td>83+ </tr>
84- <td>√</td>84+ </tbody></table>
85- </tr>85+ 
86- <tr>86+- **返回值：**
87- <td>workspaceSize</td>87+ 
88- <td>输出</td>88+ aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
89- <td>返回需要在Device侧申请的workspace大小。</td>89+ 
90- <td>-</td>90+ 
91- <td>-</td>91+## aclnnFastGelu
92- <td>-</td>92+ 
93- <td>-</td>93+- **参数说明：**
94- <td>-</td>94+ 
95- </tr>95+ <table style="undefined;table-layout: fixed; width: 1244px"><colgroup>
96- <tr>96+ <col style="width: 200px">
97- <td>executor</td>97+ <col style="width: 162px">
98- <td>输出</td>98+ <col style="width: 882px">
99- <td>返回op执行器，包含了算子计算流程。</td>99+ </colgroup>
100- <td>-</td>100+ <thead>
101- <td>-</td>101+ <tr>
102- <td>-</td>102+ <th>参数名</th>
103- <td>-</td>103+ <th>输入/输出</th>
104- <td>-</td>104+ <th>描述</th>
105- </tr>105+ </tr></thead>
106- </tbody>106+ <tbody>
107- </table>107+ <tr>
108- 108+ <td>workspace</td>
109- 109+ <td>输入</td>
110-- **返回值：**110+ <td>在Device侧申请的workspace内存地址。</td>
111- 111+ </tr>
112- aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。112+ <tr>
113- 第一段接口会完成入参校验，出现以下场景时报错：113+ <td>workspaceSize</td>
114- <table style="undefined;table-layout: fixed;width: 979px"><colgroup>114+ <td>输入</td>
115- <col style="width: 272px">115+ <td>在Device侧申请的workspace大小，由第一段接口aclnnFastGeluGetWorkspaceSize获取。</td>
116- <col style="width: 103px">116+ </tr>
117- <col style="width: 604px">117+ <tr>
118- </colgroup>118+ <td>executor</td>
119- <thead>119+ <td>输入</td>
120- <tr>120+ <td>op执行器，包含了算子计算流程。</td>
121- <th>返回码</th>121+ </tr>
122- <th>错误码</th>122+ <tr>
123- <th>描述</th>123+ <td>stream</td>
124- </tr>124+ <td>输入</td>
125- </thead>125+ <td>指定执行任务的Stream。</td>
126- <tbody>126+ </tr>
127- <tr>127+ </tbody>
128- <td>ACLNN_ERR_PARAM_NULLPTR</td>128+ </table>
129- <td>161001</td>129+ 
130- <td>self或out是空指针。</td>130+- **返回值：**
131- </tr>131+ 
132- <tr>132+ aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
133- <td rowspan="8">ACLNN_ERR_PARAM_INVALID</td>133+ 
134- <td rowspan="8">161002</td>134+## 约束说明
135- <td>self与out的数据类型、数据格式不一致。</td>135+ 
136- </tr>136+无
137- <tr>
138- <td>self或out的数据类型不在支持的范围内。</td>
139- </tr>
140- <tr>
141- <td>self与out的shape不一致。</td>
142- </tr>
143- <tr>
144- <td>self的shape维度大于8维。</td>
145- </tr>
146- </tbody></table>
147- 
148-## aclnnFastGelu
149- 
150-- **参数说明：**
151- 
152- <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
153- <col style="width: 173px">
154- <col style="width: 112px">
155- <col style="width: 668px">
156- </colgroup>
157- <thead>
158- <tr>
159- <th>参数名</th>
160- <th>输入/输出</th>
161- <th>描述</th>
162- </tr></thead>
163- <tbody>
164- <tr>
165- <td>workspace</td>
166- <td>输入</td>
167- <td>在Device侧申请的workspace内存地址。</td>
168- </tr>
169- <tr>
170- <td>workspaceSize</td>
171- <td>输入</td>
172- <td>在Device侧申请的workspace大小，由第一段接口aclnnFastGeluGetWorkspaceSize获取。</td>
173- </tr>
174- <tr>
175- <td>executor</td>
176- <td>输入</td>
177- <td>op执行器，包含了算子计算流程。</td>
178- </tr>
179- <tr>
180- <td>stream</td>
181- <td>输入</td>
182- <td>指定执行任务的Stream。</td>
183- </tr>
184- </tbody>
185- </table>
186- 
187-- **返回值：**
188- 
189- aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
190- 
191-## 约束说明
192- 
193-- 确定性计算：
194- - aclnnFastGelu默认确定性实现。
195- 
196-## 调用示例
197- 
198-示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
199- 
200-```Cpp
201-#include <iostream>
202-#include <vector>
203-#include "acl/acl.h"
204-#include "aclnnop/aclnn_fast_gelu.h"
205- 
206-#define CHECK_RET(cond, return_expr) \
207- do { \
208- if (!(cond)) { \
209- return_expr; \
210- } \
211- } while (0)
212- 
213-#define LOG_PRINT(message, ...) \
214- do { \
215- printf(message, ##__VA_ARGS__); \
216- } while (0)
217- 
218-int64_t GetShapeSize(const std::vector<int64_t>& shape) {
219- int64_t shapeSize = 1;
220- for (auto i : shape) {
221- shapeSize *= i;
222- }
223- return shapeSize;
224-}
225- 
226-int Init(int32_t deviceId, aclrtStream* stream) {
227- // 固定写法，资源初始化
228- auto ret = aclInit(nullptr);
229- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
230- ret = aclrtSetDevice(deviceId);
231- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
232- ret = aclrtCreateStream(stream);
233- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
234- return 0;
235-}
236- 
237-template <typename T>
238-int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
239- aclDataType dataType, aclTensor** tensor) {
240- auto size = GetShapeSize(shape) * sizeof(T);
241- // 调用aclrtMalloc申请device侧内存
242- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
243- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
244- // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
245- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
246- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
247- 
248- // 计算连续tensor的strides
249- std::vector<int64_t> strides(shape.size(), 1);
250- for (int64_t i = shape.size() - 2; i >= 0; i--) {
251- strides[i] = shape[i + 1] * strides[i + 1];
252- }
253- 
254- // 调用aclCreateTensor接口创建aclTensor
255- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
256- shape.data(), shape.size(), *deviceAddr);
257- return 0;
258-}
259- 
260-int main() {
261- // 1. （固定写法）device/stream初始化，参考acl API手册
262- // 根据自己的实际device填写deviceId
263- int32_t deviceId = 0;
264- aclrtStream stream;
265- auto ret = Init(deviceId, &stream);
266- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
267- // 2. 构造输入与输出，需要根据API的接口自定义构造
268- std::vector<int64_t> selfShape = {3, 3};
269- std::vector<int64_t> outShape = {3, 3};
270- void* selfDeviceAddr = nullptr;
271- aclTensor* self = nullptr;
272- void* outDeviceAddr = nullptr;
273- aclTensor* out = nullptr;
274- std::vector<float> selfHostData(9, 0);
275- std::vector<float> outHostData(9, 0);
276- // 创建self和out aclTensor
277- ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
278- CHECK_RET(ret == ACL_SUCCESS, return ret);
279- ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
280- CHECK_RET(ret == ACL_SUCCESS, return ret);
281- 
282- // 3. 调用CANN算子库API，需要修改为具体的Api名称
283- uint64_t workspaceSize = 0;
284- aclOpExecutor* executor;
285- // 调用aclnnFastGelu第一段接口
286- ret = aclnnFastGeluGetWorkspaceSize(self, out, &workspaceSize, &executor);
287- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFastGeluGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
288- // 根据第一段接口计算出的workspaceSize申请device内存
289- void* workspaceAddr = nullptr;
290- if (workspaceSize > 0) {
291- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
292- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
293- }
294- // 调用aclnnFastGelu第二段接口
295- ret = aclnnFastGelu(workspaceAddr, workspaceSize, executor, stream);
296- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFastGelu failed. ERROR: %d\n", ret); return ret);
297- // 4. （固定写法）同步等待任务执行结束
298- ret = aclrtSynchronizeStream(stream);
299- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
300- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
301- auto size = GetShapeSize(outShape);
302- std::vector<float> resultData(size, 0);
303- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
304- size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
305- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
306- for (int64_t i = 0; i < size; i++) {
307- LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
308- }
309- 
310- // 6. 释放aclTensor，需要根据具体API的接口定义修改
311- aclDestroyTensor(self);
312- aclDestroyTensor(out);
313- 
314- // 7. 释放device资源，需要根据具体API的接口定义修改
315- aclrtFree(selfDeviceAddr);
316- aclrtFree(outDeviceAddr);
317- if (workspaceSize > 0) {
318- aclrtFree(workspaceAddr);
319- }
320- aclrtDestroyStream(stream);
321- aclrtResetDevice(deviceId);
322- aclFinalize();
323- return 0;
324-}
325-```
326- 
Aactivation/fast_gelu/examples/arch35/test_aclnn_fast_gelu.cpp+141-0
@@ -0,0 +1,141 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <vector>
13+#include "acl/acl.h"
14+#include "aclnnop/aclnn_fast_gelu.h"
15+ 
16+#define CHECK_RET(cond, return_expr) \
17+ do { \
18+ if (!(cond)) { \
19+ return_expr; \
20+ } \
21+ } while (0)
22+ 
23+#define LOG_PRINT(message, ...) \
24+ do { \
25+ printf(message, ##__VA_ARGS__); \
26+ } while (0)
27+ 
28+int64_t GetShapeSize(const std::vector<int64_t>& shape)
29+{
30+ int64_t shapeSize = 1;
31+ for (auto i : shape) {
32+ shapeSize *= i;
33+ }
34+ return shapeSize;
35+}
36+ 
37+int Init(int32_t deviceId, aclrtStream* stream)
38+{
39+ // 固定写法，资源初始化
40+ auto ret = aclInit(nullptr);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtSetDevice(deviceId);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
44+ ret = aclrtCreateStream(stream);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
46+ return 0;
47+}
48+ 
49+template <typename T>
50+int CreateAclTensor(
51+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
52+ aclTensor** tensor)
53+{
54+ auto size = GetShapeSize(shape) * sizeof(T);
55+ // 调用aclrtMalloc申请device侧内存
56+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
57+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
58+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
59+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
61+ 
62+ // 计算连续tensor的strides
63+ std::vector<int64_t> strides(shape.size(), 1);
64+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
65+ strides[i] = shape[i + 1] * strides[i + 1];
66+ }
67+ 
68+ // 调用aclCreateTensor接口创建aclTensor
69+ *tensor = aclCreateTensor(
70+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
71+ *deviceAddr);
72+ return 0;
73+}
74+ 
75+int main()
76+{
77+ // 1. （固定写法）device/stream初始化，参考acl API手册
78+ // 根据自己的实际device填写deviceId
79+ int32_t deviceId = 0;
80+ aclrtStream stream;
81+ auto ret = Init(deviceId, &stream);
82+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
83+ // 2. 构造输入与输出，需要根据API的接口自定义构造
84+ std::vector<int64_t> selfShape = {3, 3};
85+ std::vector<int64_t> outShape = {3, 3};
86+ void* selfDeviceAddr = nullptr;
87+ aclTensor* self = nullptr;
88+ void* outDeviceAddr = nullptr;
89+ aclTensor* out = nullptr;
90+ std::vector<float> selfHostData(9, 0);
91+ std::vector<float> outHostData(9, 0);
92+ // 创建self和out aclTensor
93+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
94+ CHECK_RET(ret == ACL_SUCCESS, return ret);
95+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
96+ CHECK_RET(ret == ACL_SUCCESS, return ret);
97+ 
98+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
99+ uint64_t workspaceSize = 0;
100+ aclOpExecutor* executor;
101+ // 调用aclnnFastGelu第一段接口
102+ ret = aclnnFastGeluGetWorkspaceSize(self, out, &workspaceSize, &executor);
103+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFastGeluGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
104+ // 根据第一段接口计算出的workspaceSize申请device内存
105+ void* workspaceAddr = nullptr;
106+ if (workspaceSize > 0) {
107+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
108+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
109+ }
110+ // 调用aclnnFastGelu第二段接口
111+ ret = aclnnFastGelu(workspaceAddr, workspaceSize, executor, stream);
112+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFastGelu failed. ERROR: %d\n", ret); return ret);
113+ // 4. （固定写法）同步等待任务执行结束
114+ ret = aclrtSynchronizeStream(stream);
115+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
116+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
117+ auto size = GetShapeSize(outShape);
118+ std::vector<float> resultData(size, 0);
119+ ret = aclrtMemcpy(
120+ resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]),
121+ ACL_MEMCPY_DEVICE_TO_HOST);
122+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
123+ for (int64_t i = 0; i < size; i++) {
124+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
125+ }
126+ 
127+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
128+ aclDestroyTensor(self);
129+ aclDestroyTensor(out);
130+ 
131+ // 7. 释放device资源，需要根据具体API的接口定义修改
132+ aclrtFree(selfDeviceAddr);
133+ aclrtFree(outDeviceAddr);
134+ if (workspaceSize > 0) {
135+ aclrtFree(workspaceAddr);
136+ }
137+ aclrtDestroyStream(stream);
138+ aclrtResetDevice(deviceId);
139+ aclFinalize();
140+ return 0;
141+}
Dactivation/fast_gelu/examples/test_aclnn_fast_gelu.cpp+0-134
@@ -1,134 +0,0 @@
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 <vector>
13-#include "acl/acl.h"
14-#include "aclnnop/aclnn_fast_gelu.h"
15- 
16-#define CHECK_RET(cond, return_expr) \
17- do { \
18- if (!(cond)) { \
19- return_expr; \
20- } \
21- } while (0)
22- 
23-#define LOG_PRINT(message, ...) \
24- do { \
25- printf(message, ##__VA_ARGS__); \
26- } while (0)
27- 
28-int64_t GetShapeSize(const std::vector<int64_t>& shape) {
29- int64_t shapeSize = 1;
30- for (auto i : shape) {
31- shapeSize *= i;
32- }
33- return shapeSize;
34-}
35- 
36-int Init(int32_t deviceId, aclrtStream* stream) {
37- // 固定写法，资源初始化
38- auto ret = aclInit(nullptr);
39- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
40- ret = aclrtSetDevice(deviceId);
41- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
42- ret = aclrtCreateStream(stream);
43- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
44- return 0;
45-}
46- 
47-template <typename T>
48-int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
49- aclDataType dataType, aclTensor** tensor) {
50- auto size = GetShapeSize(shape) * sizeof(T);
51- // 调用aclrtMalloc申请device侧内存
52- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
53- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
54- // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
55- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
56- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
57- 
58- // 计算连续tensor的strides
59- std::vector<int64_t> strides(shape.size(), 1);
60- for (int64_t i = shape.size() - 2; i >= 0; i--) {
61- strides[i] = shape[i + 1] * strides[i + 1];
62- }
63- 
64- // 调用aclCreateTensor接口创建aclTensor
65- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
66- shape.data(), shape.size(), *deviceAddr);
67- return 0;
68-}
69- 
70-int main() {
71- // 1. （固定写法）device/stream初始化，参考acl API手册
72- // 根据自己的实际device填写deviceId
73- int32_t deviceId = 0;
74- aclrtStream stream;
75- auto ret = Init(deviceId, &stream);
76- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
77- // 2. 构造输入与输出，需要根据API的接口自定义构造
78- std::vector<int64_t> selfShape = {3, 3};
79- std::vector<int64_t> outShape = {3, 3};
80- void* selfDeviceAddr = nullptr;
81- aclTensor* self = nullptr;
82- void* outDeviceAddr = nullptr;
83- aclTensor* out = nullptr;
84- std::vector<float> selfHostData(9, 0);
85- std::vector<float> outHostData(9, 0);
86- // 创建self和out aclTensor
87- ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
88- CHECK_RET(ret == ACL_SUCCESS, return ret);
89- ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
90- CHECK_RET(ret == ACL_SUCCESS, return ret);
91- 
92- // 3. 调用CANN算子库API，需要修改为具体的Api名称
93- uint64_t workspaceSize = 0;
94- aclOpExecutor* executor;
95- // 调用aclnnFastGelu第一段接口
96- ret = aclnnFastGeluGetWorkspaceSize(self, out, &workspaceSize, &executor);
97- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFastGeluGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
98- // 根据第一段接口计算出的workspaceSize申请device内存
99- void* workspaceAddr = nullptr;
100- if (workspaceSize > 0) {
101- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
102- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
103- }
104- // 调用aclnnFastGelu第二段接口
105- ret = aclnnFastGelu(workspaceAddr, workspaceSize, executor, stream);
106- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFastGelu failed. ERROR: %d\n", ret); return ret);
107- // 4. （固定写法）同步等待任务执行结束
108- ret = aclrtSynchronizeStream(stream);
109- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
110- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
111- auto size = GetShapeSize(outShape);
112- std::vector<float> resultData(size, 0);
113- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
114- size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
115- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
116- for (int64_t i = 0; i < size; i++) {
117- LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
118- }
119- 
120- // 6. 释放aclTensor，需要根据具体API的接口定义修改
121- aclDestroyTensor(self);
122- aclDestroyTensor(out);
123- 
124- // 7. 释放device资源，需要根据具体API的接口定义修改
125- aclrtFree(selfDeviceAddr);
126- aclrtFree(outDeviceAddr);
127- if (workspaceSize > 0) {
128- aclrtFree(workspaceAddr);
129- }
130- aclrtDestroyStream(stream);
131- aclrtResetDevice(deviceId);
132- aclFinalize();
133- return 0;
134-}
Ractivation/fast_gelu/op_host/op_api/aclnn_fast_gelu.cpp→activation/fast_gelu/op_api/aclnn_fast_gelu.cpp+0-0
文件重命名但无更改。
Ractivation/fast_gelu/op_host/op_api/aclnn_fast_gelu.h→activation/fast_gelu/op_api/aclnn_fast_gelu.h+6-5
@@ -1,11 +1,12 @@
1/**1/**
2+ * This program is free software, you can redistribute it and/or modify.
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.3 * 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 of4+ * This file is a part of the CANN Open Software.
4- * CANN Open Software License Agreement Version 2.0 (the "License").5+ * 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 * 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+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
8- * See LICENSE in the root of the software repository for the full text of the License.9+ * the software repository for the full text of the License.
9 */10 */
10 11 
11#ifndef OP_API_INC_LEVEL2_ACLNN_FAST_GELU_H_12#ifndef OP_API_INC_LEVEL2_ACLNN_FAST_GELU_H_
Ractivation/fast_gelu/op_host/op_api/fast_gelu.cpp→activation/fast_gelu/op_api/fast_gelu.cpp+6-5
@@ -1,11 +1,12 @@
1/**1/**
2+ * This program is free software, you can redistribute it and/or modify.
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.3 * 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 of4+ * This file is a part of the CANN Open Software.
4- * CANN Open Software License Agreement Version 2.0 (the "License").5+ * 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 * 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+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.8+ * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of
8- * See LICENSE in the root of the software repository for the full text of the License.9+ * the software repository for the full text of the License.
9 */10 */
10#include "fast_gelu.h"11#include "fast_gelu.h"
11#include "opdev/op_dfx.h"12#include "opdev/op_dfx.h"
Ractivation/fast_gelu/op_host/op_api/fast_gelu.h→activation/fast_gelu/op_api/fast_gelu.h+1-0
@@ -7,6 +7,7 @@
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10+ 
10#ifndef OP_API_INC_LEVEL0_FAST_GELU_H_11#ifndef OP_API_INC_LEVEL0_FAST_GELU_H_
11#define OP_API_INC_LEVEL0_FAST_GELU_H_12#define OP_API_INC_LEVEL0_FAST_GELU_H_
12#include "opdev/op_executor.h"13#include "opdev/op_executor.h"
Aactivation/fast_gelu/op_graph/fast_gelu_proto.h+40-0
@@ -0,0 +1,40 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OPS_BUILT_IN_OP_PROTO_INC_NONLINEAR_FUC_OPS_H_
12+#define OPS_BUILT_IN_OP_PROTO_INC_NONLINEAR_FUC_OPS_H_
13+ 
14+#include "graph/operator_reg.h"
15+ 
16+namespace ge {
17+ 
18+/**
19+* @brief
20+* For Atlas 200/300/500 Inference Product, Atlas Training Series Product,
21+ Atlas Inference Series Product, Ascend 610 AI Processor,
22+ the calculation formula is x*e^(0.851*x)*(x-|x|)/(1+e^(-1.702|x|)).
23+* For other chips, the calculation formula is x/(1+e^(-1.702*x)).
24+ 
25+* @par Inputs:
26+* One input, including:
27+* x: An ND or 5HD tensor. Support 1D~8D. Must be one of the following types:
28+* bfloat16, float16, float32
29+ 
30+* @par Outputs:
31+* y: A Tensor. Has the same type, format and shape as "x".
32+* @par Third-party framework compatibility
33+* Compatible with the TensorFlow operator FastGelu
34+*/
35+REG_OP(FastGelu)
36+ .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
37+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
38+ .OP_END_FACTORY_REG(FastGelu)
39+} // namespace ge
40+#endif // OPS_BUILT_IN_OP_PROTO_INC_NONLINEAR_FUC_OPS_H_
Dactivation/fast_gelu/op_host/CMakeLists.txt+0-10
@@ -1,10 +0,0 @@
1-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3-# CANN Open Software License Agreement Version 2.0 (the "License").
4-# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7-# See LICENSE in the root of the software repository for the full text of the License.
8-#/
9- 
10-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE fast_gelu ACLNNTYPE aclnn_exclude)
Aactivation/fast_gelu/op_host/arch35/fast_gelu_tiling_arch35.cpp+136-0
@@ -0,0 +1,136 @@
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 <graph/utils/type_utils.h>
13+#include "fast_gelu_tiling_arch35.h"
14+#include "log/log.h"
15+#include "platform/platform_ascendc.h"
16+#include "register/op_def_registry.h"
17+#include "register/tilingdata_base.h"
18+#include "activation/fast_gelu/op_kernel/arch35/fast_gelu_dag.h"
19+#include "activation/fast_gelu/op_kernel/arch35/fast_gelu_struct.h"
20+#include "atvoss/broadcast/broadcast_tiling.h"
21+#include "atvoss/elewise/elewise_tiling.h"
22+ 
23+using namespace FastGeluOp;
24+using namespace Ops::Base;
25+ 
26+namespace optiling {
27+const int64_t SYSWORKSPACE = 16777216; // 16 * 1024 * 1024
28+ 
29+ge::graphStatus FastGeluTiling::CalcInputDtype()
30+{
31+ auto inputDesc = tilingContext->GetInputDesc(0);
32+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, inputDesc);
33+ this->inputDtype = inputDesc->GetDataType();
34+ OP_CHECK_IF(
35+ this->inputDtype != ge::DT_FLOAT16 && this->inputDtype != ge::DT_BF16 && this->inputDtype != ge::DT_FLOAT,
36+ OP_LOGE(tilingContext->GetNodeName(), "x dtype not support"), return ge::GRAPH_FAILED);
37+ return ge::GRAPH_SUCCESS;
38+}
39+ 
40+ge::graphStatus FastGeluTiling::CalcOutputDtype()
41+{
42+ auto outputDesc = tilingContext->GetOutputDesc(0);
43+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outputDesc);
44+ this->outputDtype = outputDesc->GetDataType();
45+ OP_CHECK_IF(
46+ this->outputDtype != ge::DT_FLOAT16 && this->outputDtype != ge::DT_BF16 && this->outputDtype != ge::DT_FLOAT,
47+ OP_LOGE(tilingContext->GetNodeName(), "output y dtype not support"), return ge::GRAPH_FAILED);
48+ OP_CHECK_IF(
49+ this->outputDtype != this->inputDtype,
50+ OP_LOGE(tilingContext->GetNodeName(), "input x dtype not same as output y"), return ge::GRAPH_FAILED);
51+ return ge::GRAPH_SUCCESS;
52+}
53+ 
54+ge::graphStatus FastGeluTiling::CheckShape()
55+{
56+ auto selfStorageShape = tilingContext->GetInputShape(0);
57+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, selfStorageShape);
58+ const gert::Shape& inputShape = Ops::Base::EnsureNotScalar(selfStorageShape->GetStorageShape());
59+ 
60+ auto outStorageShape = tilingContext->GetOutputShape(0);
61+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outStorageShape);
62+ const gert::Shape& outputShape = Ops::Base::EnsureNotScalar(outStorageShape->GetStorageShape());
63+ 
64+ OP_CHECK_IF(
65+ inputShape != outputShape, OP_LOGE(tilingContext->GetNodeName(), "input x and output y shape not same"),
66+ return ge::GRAPH_FAILED);
67+ return ge::GRAPH_SUCCESS;
68+}
69+ 
70+std::string FastGeluTiling::DataTypeToSerialString(const ge::DataType type) const
71+{
72+ const auto it = DATATYPE_TO_STRING_MAP.find(type);
73+ if (it != DATATYPE_TO_STRING_MAP.end()) {
74+ return it->second;
75+ } else {
76+ OP_LOGE(tilingContext->GetNodeName(), "datatype %d not find", type);
77+ return "UNDEFINED";
78+ }
79+}
80+ 
81+ge::graphStatus FastGeluTiling::RunTiling()
82+{
83+ ElewiseBaseTiling elewiseBaseTiling(tilingContext);
84+ // 获取tiling计算所需的参数
85+ ge::graphStatus status = ge::GRAPH_FAILED;
86+ status = CalcInputDtype();
87+ OP_CHECK_IF(status == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "get input dtype failed"), return ge::GRAPH_FAILED);
88+ status = CalcOutputDtype();
89+ OP_CHECK_IF(status == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "get output dtype failed"), return ge::GRAPH_FAILED);
90+ status = CheckShape();
91+ OP_CHECK_IF(status == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "check shape failed"), return ge::GRAPH_FAILED);
92+ 
93+ auto tiling = tilingContext->GetTilingData<EleBaseTilingDataV2>();
94+ OP_CHECK_IF(
95+ (tiling == nullptr), OP_LOGE(tilingContext->GetNodeName(), "Get FastGeluTiling from GE context failed"),
96+ return ge::GRAPH_FAILED);
97+ if (this->outputDtype == ge::DT_FLOAT16) {
98+ dType = TPL_FP16;
99+ status = elewiseBaseTiling.DoTiling<FastGeluDag::FastGeluNeedCast<half>::OpDag>(*tiling);
100+ } else if (this->outputDtype == ge::DT_BF16) {
101+ dType = TPL_BF16;
102+ status = elewiseBaseTiling.DoTiling<FastGeluDag::FastGeluNeedCast<bfloat16_t>::OpDag>(*tiling);
103+ } else if (this->outputDtype == ge::DT_FLOAT) {
104+ dType = TPL_FP32;
105+ status = elewiseBaseTiling.DoTiling<FastGeluDag::FastGeluNoCast<float>::OpDag>(*tiling);
106+ } else {
107+ OP_LOGE(tilingContext->GetNodeName(), "elewiseBaseTiling DoTiling failed.");
108+ return ge::GRAPH_FAILED;
109+ }
110+ OP_CHECK_IF(
111+ status == ge::GRAPH_FAILED, OP_LOGE(tilingContext, "elewiseBaseTiling failed"), return ge::GRAPH_FAILED);
112+ const uint64_t tilingKey = GET_TPL_TILING_KEY(tiling->scheMode, dType);
113+ OP_LOGD(tilingContext->GetNodeName(), "[TilingData] : tilingKey=%ld.", tilingKey);
114+ tilingContext->SetTilingKey(tilingKey);
115+ tilingContext->SetBlockDim(tiling->blockNum);
116+ size_t usrWorkspaceSize = 0;
117+ size_t sysWorkspaceSize = SYSWORKSPACE;
118+ size_t* currentWorkspace = tilingContext->GetWorkspaceSizes(1);
119+ currentWorkspace[0] = sysWorkspaceSize + usrWorkspaceSize;
120+ return ge::GRAPH_SUCCESS;
121+}
122+ 
123+static ge::graphStatus TilingFuncFastGelu(gert::TilingContext* tilingContext)
124+{
125+ OP_LOGD(tilingContext->GetNodeName(), "START FastGelu AscendC Tiling \n");
126+ FastGeluTiling FastGeluTiling(tilingContext);
127+ return FastGeluTiling.RunTiling();
128+}
129+ 
130+ge::graphStatus TilingPrepareForFastGelu([[maybe_unused]] gert::TilingParseContext* context)
131+{
132+ return ge::GRAPH_SUCCESS;
133+}
134+ 
135+IMPL_OP_OPTILING(FastGelu).Tiling(TilingFuncFastGelu).TilingParse<FastGeluCompileInfo>(TilingPrepareForFastGelu);
136+} // namespace optiling
Aactivation/fast_gelu/op_host/arch35/fast_gelu_tiling_arch35.h+78-0
@@ -0,0 +1,78 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_FAST_GELU_TILING_H
12+#define OPS_BUILD_IN_OP_TILING_RUNTIME_FAST_GELU_TILING_H
13+ 
14+#include "register/tilingdata_base.h"
15+#include "register/op_impl_registry.h"
16+ 
17+namespace optiling {
18+ 
19+class FastGeluTiling {
20+public:
21+ explicit FastGeluTiling(gert::TilingContext* context) : tilingContext(context){};
22+ ge::graphStatus RunTiling();
23+ 
24+protected:
25+ ge::graphStatus CalcInputDtype();
26+ ge::graphStatus CalcOutputDtype();
27+ ge::graphStatus CheckShape();
28+ std::string DataTypeToSerialString(const ge::DataType type) const;
29+ 
30+private:
31+ uint64_t dType = 0;
32+ ge::DataType inputDtype;
33+ ge::DataType outputDtype;
34+ gert::TilingContext* tilingContext;
35+};
36+ 
37+const std::map<ge::DataType, std::string> DATATYPE_TO_STRING_MAP = {
38+ {ge::DT_UNDEFINED, "DT_UNDEFINED"}, // Used to indicate a DataType field has not been set.
39+ {ge::DT_FLOAT, "DT_FLOAT"}, // float type
40+ {ge::DT_FLOAT16, "DT_FLOAT16"}, // fp16 type
41+ {ge::DT_INT8, "DT_INT8"}, // int8 type
42+ {ge::DT_INT16, "DT_INT16"}, // int16 type
43+ {ge::DT_UINT16, "DT_UINT16"}, // uint16 type
44+ {ge::DT_UINT8, "DT_UINT8"}, // uint8 type
45+ {ge::DT_INT32, "DT_INT32"}, // uint32 type
46+ {ge::DT_INT64, "DT_INT64"}, // int64 type
47+ {ge::DT_UINT32, "DT_UINT32"}, // unsigned int32
48+ {ge::DT_UINT64, "DT_UINT64"}, // unsigned int64
49+ {ge::DT_BOOL, "DT_BOOL"}, // bool type
50+ {ge::DT_DOUBLE, "DT_DOUBLE"}, // double type
51+ {ge::DT_DUAL, "DT_DUAL"}, // dual output type
52+ {ge::DT_DUAL_SUB_INT8, "DT_DUAL_SUB_INT8"}, // dual output int8 type
53+ {ge::DT_DUAL_SUB_UINT8, "DT_DUAL_SUB_UINT8"}, // dual output uint8 type
54+ {ge::DT_COMPLEX32, "DT_COMPLEX32"}, // complex32 type
55+ {ge::DT_COMPLEX64, "DT_COMPLEX64"}, // complex64 type
56+ {ge::DT_COMPLEX128, "DT_COMPLEX128"}, // complex128 type
57+ {ge::DT_QINT8, "DT_QINT8"}, // qint8 type
58+ {ge::DT_QINT16, "DT_QINT16"}, // qint16 type
59+ {ge::DT_QINT32, "DT_QINT32"}, // qint32 type
60+ {ge::DT_QUINT8, "DT_QUINT8"}, // quint8 type
61+ {ge::DT_QUINT16, "DT_QUINT16"}, // quint16 type
62+ {ge::DT_RESOURCE, "DT_RESOURCE"}, // resource type
63+ {ge::DT_STRING_REF, "DT_STRING_REF"}, // string ref type
64+ {ge::DT_STRING, "DT_STRING"}, // string type
65+ {ge::DT_VARIANT, "DT_VARIANT"}, // dt_variant type
66+ {ge::DT_BF16, "DT_BFLOAT16"}, // dt_bfloat16 type
67+ {ge::DT_INT4, "DT_INT4"}, // dt_variant type
68+ {ge::DT_UINT1, "DT_UINT1"}, // dt_variant type
69+ {ge::DT_INT2, "DT_INT2"}, // dt_variant type
70+ {ge::DT_UINT2, "DT_UINT2"} // dt_variant type
71+};
72+ 
73+struct FastGeluCompileInfo {
74+ uint64_t coreNum = 0;
75+ uint64_t ubSize = 0;
76+};
77+} // namespace optiling
78+#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_FAST_GELU_TILING_H
Aactivation/fast_gelu/op_host/config/ascend910_95/fast_gelu_binary.json+92-0
@@ -0,0 +1,92 @@
1+{
2+ "op_type": "FastGelu",
3+ "op_list": [
4+ {
5+ "bin_filename": "FastGelu_bf16",
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": "FastGelu_fp16",
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": "FastGelu_fp32",
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/fast_gelu/op_host/config/ascend910_95/fast_gelu_simplified_key.ini+2-0
@@ -0,0 +1,2 @@
1+[FastGelu]
2+default=0
Aactivation/fast_gelu/op_host/fast_gelu_def.cpp+40-0
@@ -0,0 +1,40 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "register/op_def_registry.h"
12+ 
13+namespace ops {
14+class FastGelu : public OpDef {
15+public:
16+ explicit FastGelu(const char* name) : OpDef(name)
17+ {
18+ this->Input("x")
19+ .ParamType(REQUIRED)
20+ .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT})
21+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
22+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
23+ this->Output("y")
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+ OpAICoreConfig aicoreConfig;
29+ aicoreConfig.DynamicCompileStaticFlag(true)
30+ .DynamicFormatFlag(false)
31+ .DynamicRankSupportFlag(true)
32+ .DynamicShapeSupportFlag(true)
33+ .NeedCheckSupportFlag(false)
34+ .ExtendCfgInfo("opFile.value", "fast_gelu_apt");
35+ this->AICore().AddConfig("ascend910_95", aicoreConfig);
36+ }
37+};
38+ 
39+OP_ADD(FastGelu);
40+} // namespace ops
Ractivation/fast_gelu_grad/tests/ut/op_host/test_aclnn_fast_gelu_grad.cpp→activation/fast_gelu/op_host/fast_gelu_infershape.cpp+6-27
@@ -8,31 +8,10 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-#include "gtest/gtest.h"11+#include "infershape_elewise_util.h"
12-#include "../../../op_host/op_api/aclnn_fast_gelu_backward.h"12+#include "register/op_impl_registry.h"
13-#include "op_api_ut_common/tensor_desc.h"
14-#include "op_api_ut_common/op_api_ut.h"
15 13 
16-class l2_fast_gelu_grad_test : public testing::Test {14+using namespace ge;
17-protected:15+namespace ops {
18- static void SetUpTestCase() {16+IMPL_OP_INFERSHAPE(FastGelu).InferShape(Ops::Base::InferShape4Elewise);
19- std::cout << "l2_fast_gelu_grad_test SetUp" << std::endl;17+} // namespace ops
20- }
21- 
22- static void TearDownTestCase() { std::cout << "l2_fast_gelu_grad_test TearDown" << std::endl; }
23-};
24- 
25-// 正常路径，float32
26-TEST_F(l2_fast_gelu_grad_test, l2_fast_gelu_grad_test_005) {
27- auto gradOutDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
28- auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
29- auto outDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
30- 
31- auto ut = OP_API_UT(aclnnFastGeluBackward, INPUT(gradOutDesc,selfDesc), OUTPUT(outDesc));
32- 
33- uint64_t workspaceSize = 0;
34- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
35- EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
36- 
37- //ut.TestPrecision();
38-}
Aactivation/fast_gelu/op_kernel/arch35/fast_gelu_dag.h+87-0
@@ -0,0 +1,87 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef CANN_CUSTOM_OPS_FAST_GELU_DAG_H
12+#define CANN_CUSTOM_OPS_FAST_GELU_DAG_H
13+#include "atvoss/util/dag.h"
14+#include "atvoss/util/vec.h"
15+#include "atvoss/util/placeholder.h"
16+ 
17+namespace FastGeluDag {
18+using namespace Ops::Base;
19+ 
20+template <class T>
21+struct FastGeluCustom : public Vec::ElemwiseUnaryOP<T, T> {
22+ __aicore__ inline FastGeluCustom(LocalTensor<T>& dst, LocalTensor<T>& src, uint32_t count)
23+ {
24+#ifdef __CCE_AICORE__
25+ uint32_t dtypeSize = sizeof(T);
26+ uint32_t VL = AscendC::VECTOR_REG_WIDTH / dtypeSize;
27+ uint16_t loopNum = CeilDivision(count, VL);
28+ uint32_t vlSize = VL;
29+ T value1 = -1.702;
30+ T value2 = 1.0;
31+ 
32+ __VEC_SCOPE__
33+ {
34+ __ubuf__ T* srcAddr = (__ubuf__ T*)src.GetPhyAddr();
35+ __ubuf__ T* dstAddr = (__ubuf__ T*)dst.GetPhyAddr();
36+ 
37+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> x;
38+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> denominator;
39+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> result;
40+ static constexpr AscendC::MicroAPI::DivSpecificMode mode = {
41+ AscendC::MicroAPI::MaskMergeMode::ZEROING, true};
42+ AscendC::MicroAPI::MaskReg mask;
43+ 
44+ for (uint16_t loopIdx = 0; loopIdx < loopNum; loopIdx++) {
45+ mask = AscendC::MicroAPI::UpdateMask<T, AscendC::MicroAPI::RegTraitNumOne>(count);
46+ AscendC::MicroAPI::DataCopy(x, (__ubuf__ T*)(srcAddr + loopIdx * vlSize));
47+ AscendC::MicroAPI::Muls(denominator, x, value1, mask);
48+ AscendC::MicroAPI::Exp(denominator, denominator, mask);
49+ AscendC::MicroAPI::Adds(denominator, denominator, value2, mask);
50+ // result = x / (Exp(-1.702 * x) + 1)
51+ AscendC::MicroAPI::Div<T, &mode>(result, x, denominator, mask);
52+ // OpCopyOut
53+ AscendC::MicroAPI::DataCopy((__ubuf__ T*)(dstAddr + loopIdx * vlSize), result, mask);
54+ }
55+ }
56+#endif
57+ }
58+};
59+ 
60+template <typename T>
61+struct FastGeluNoCast {
62+ // 通过Compute构造计算图
63+ using OpCopyIn = Bind<Vec::CopyIn<T>, Placeholder::In0<T>>;
64+ using OpFastGelu = Bind<FastGeluCustom<float>, OpCopyIn>;
65+ 
66+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpFastGelu>;
67+ // 指定输出节点
68+ using Outputs = Elems<OpCopyOut>; // 设置输出
69+ using OpDag = DAGSch<Outputs>;
70+};
71+ 
72+template <typename T>
73+struct FastGeluNeedCast {
74+ // 通过Compute构造计算图
75+ using OpCopyIn = Bind<Vec::CopyIn<T>, Placeholder::In0<T>>; // x
76+ using CastIn = Bind<Vec::Cast<float, T, 0>, OpCopyIn>;
77+ using OpFastGelu = Bind<FastGeluCustom<float>, CastIn>;
78+ 
79+ using CastOut = Bind<Vec::Cast<T, float, 1>, OpFastGelu>;
80+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, CastOut>;
81+ // 指定输出节点
82+ using Outputs = Elems<OpCopyOut>; // 设置输出
83+ using OpDag = DAGSch<Outputs>;
84+};
85+ 
86+}; // namespace FastGeluDag
87+#endif // CANN_CUSTOM_OPS_FAST_GELU_DAG_H
Aactivation/fast_gelu/op_kernel/arch35/fast_gelu_struct.h+39-0
@@ -0,0 +1,39 @@
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 "ascendc/host_api/tiling/template_argument.h"
12+ 
13+#ifndef CANN_CUSTOM_OPS_FAST_GELU_STRUCT_H_
14+#define CANN_CUSTOM_OPS_FAST_GELU_STRUCT_H_
15+ 
16+namespace FastGeluOp {
17+#define TPL_FP16 1
18+#define TPL_BF16 2
19+#define TPL_FP32 3
20+#define FastGelu_TPL_SCH_MODE_0 0
21+#define FastGelu_TPL_SCH_MODE_1 1
22+ 
23+ASCENDC_TPL_ARGS_DECL(
24+ FastGelu, ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, FastGelu_TPL_SCH_MODE_0, FastGelu_TPL_SCH_MODE_1),
25+ ASCENDC_TPL_DTYPE_DECL(dType, TPL_FP16, TPL_BF16, TPL_FP32));
26+ 
27+ASCENDC_TPL_SEL(
28+ ASCENDC_TPL_ARGS_SEL(
29+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, FastGelu_TPL_SCH_MODE_0, FastGelu_TPL_SCH_MODE_1),
30+ ASCENDC_TPL_DTYPE_SEL(dType, TPL_FP16)),
31+ ASCENDC_TPL_ARGS_SEL(
32+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, FastGelu_TPL_SCH_MODE_0, FastGelu_TPL_SCH_MODE_1),
33+ ASCENDC_TPL_DTYPE_SEL(dType, TPL_BF16)),
34+ ASCENDC_TPL_ARGS_SEL(
35+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, FastGelu_TPL_SCH_MODE_0, FastGelu_TPL_SCH_MODE_1),
36+ ASCENDC_TPL_DTYPE_SEL(dType, TPL_FP32)));
37+} // namespace FastGeluOp
38+ 
39+#endif // CANN_CUSTOM_OPS_FAST_GELU_STRUCT_H_
Aactivation/fast_gelu/op_kernel/fast_gelu_apt.cpp+49-0
@@ -0,0 +1,49 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "kernel_operator.h"
12+#include "kernel_tiling/kernel_tiling.h"
13+#include "atvoss/elewise/elewise_sch.h"
14+#include "arch35/fast_gelu_dag.h"
15+#include "arch35/fast_gelu_struct.h"
16+ 
17+using namespace AscendC;
18+using namespace FastGeluOp;
19+using namespace Ops::Base;
20+ 
21+template <uint64_t schMode, uint64_t dType>
22+__global__ __aicore__ void fast_gelu(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
23+{
24+ if (workspace == nullptr) {
25+ return;
26+ }
27+ SetSysWorkspace(workspace);
28+ GM_ADDR userWS = GetUserWorkspace(workspace);
29+ if (userWS == nullptr) {
30+ return;
31+ }
32+ REGISTER_TILING_DEFAULT(EleBaseTilingDataV2);
33+ GET_TILING_DATA_WITH_STRUCT(EleBaseTilingDataV2, tilingData, tiling);
34+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
35+ TPipe pipe;
36+ if constexpr (dType == static_cast<uint64_t>(TPL_FP16)) {
37+ ElementwiseSch<schMode, FastGeluDag::FastGeluNeedCast<half>::OpDag> sch(&tilingData, &pipe);
38+ sch.Init(x, y);
39+ sch.Process();
40+ } else if constexpr (dType == static_cast<uint64_t>(TPL_BF16)) {
41+ ElementwiseSch<schMode, FastGeluDag::FastGeluNeedCast<bfloat16_t>::OpDag> sch(&tilingData, &pipe);
42+ sch.Init(x, y);
43+ sch.Process();
44+ } else if constexpr (dType == static_cast<uint64_t>(TPL_FP32)) {
45+ ElementwiseSch<schMode, FastGeluDag::FastGeluNoCast<float>::OpDag> sch(&tilingData, &pipe);
46+ sch.Init(x, y);
47+ sch.Process();
48+ }
49+}
Mactivation/fast_gelu/tests/CMakeLists.txt+8-8
@@ -1,17 +1,17 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-message(STATUS "=== Debug: start ops.activation.fast_gelu.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()
17endforeach()17endforeach()
Mactivation/fast_gelu/tests/ut/CMakeLists.txt+4-5
@@ -1,14 +1,13 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
11-message(STATUS "=== Debug: start ops.activation.fast_gelu.tests.ut.CMakeLists.txt ")
12file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)11file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
13message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")12message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
14foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})13foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
Rindex/index_put_with_sort_v2/op_host/CMakeLists.txt→activation/fast_gelu/tests/ut/op_api/CMakeLists.txt+6-5
@@ -1,13 +1,14 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
5# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11- 11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12-#optiling and opapi12+if(UT_TEST_ALL OR OP_API_UT)
13-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})13+ add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
Aactivation/fast_gelu/tests/ut/op_api/test_aclnn_fast_gelu.cpp+44-0
@@ -0,0 +1,44 @@
1+/**
2+ * This program is free software, you can redistribute it and/or modify.
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 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+#include <vector>
11+#include <array>
12+#include "gtest/gtest.h"
13+ 
14+#include "../../../op_api/aclnn_fast_gelu.h"
15+ 
16+#include "op_api_ut_common/tensor_desc.h"
17+#include "op_api_ut_common/scalar_desc.h"
18+#include "op_api_ut_common/op_api_ut.h"
19+ 
20+using namespace std;
21+ 
22+ 
23+class l2_fast_gelu_test : public testing::Test {
24+ protected:
25+ static void SetUpTestCase() {
26+ cout << "fast_gelu_test SetUp" << endl;
27+ }
28+ 
29+ static void TearDownTestCase() {
30+ cout << "fast_gelu_test TearDown" << endl;
31+ }
32+};
33+ 
34+TEST_F(l2_fast_gelu_test, test_fast_gelu_float32) {
35+ auto input_tensor = TensorDesc({1,16}, ACL_FLOAT, ACL_FORMAT_ND);
36+ auto output_tensor = TensorDesc({1,16}, ACL_FLOAT, ACL_FORMAT_ND);
37+ 
38+ auto ut = OP_API_UT(aclnnFastGelu, INPUT(input_tensor),
39+ OUTPUT(output_tensor));
40+ 
41+ uint64_t workspaceSize = 0;
42+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
43+ EXPECT_EQ(aclRet, ACL_SUCCESS);
44+}
Mactivation/fast_gelu/tests/ut/op_host/CMakeLists.txt+4-6
@@ -1,17 +1,15 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11if(UT_TEST_ALL OR OP_HOST_UT)12if(UT_TEST_ALL OR OP_HOST_UT)
12 add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})13 add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13 add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})14 add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14endif()15endif()
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()
Aactivation/fast_gelu/tests/ut/op_host/test_fast_gelu_infershape.cpp+57-0
@@ -0,0 +1,57 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
C
Cchenqi3172025年12月29日

arch35 UT 放到op_host arch35 目录下

likedislike
Y
yuanqi_li
2025年12月29日 评论：
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 <gtest/gtest.h>
13+#include "infershape_test_util.h"
14+#include "exe_graph/runtime/storage_format.h"
15+#include "exe_graph/runtime/storage_shape.h"
16+#include "kernel_run_context_facker.h"
17+#include "register/op_impl_registry.h"
18+#include "log/log.h"
19+#include "platform/platform_info.h"
20+#include "ut_op_common.h"
21+#include "../../../op_graph/fast_gelu_proto.h"
22+ 
23+using namespace ge;
24+ 
25+class FastGeluTest : public testing::Test {
26+protected:
27+ static void SetUpTestCase()
28+ {
29+ std::cout << "FastGeluTest SetUp" << std::endl;
30+ }
31+ 
32+ static void TearDownTestCase()
33+ {
34+ std::cout << "FastGeluTest TearDown" << std::endl;
35+ }
36+};
37+ 
38+TEST_F(FastGeluTest, fast_gelu_infershape_test)
39+{
40+ ge::op::FastGelu op;
41+ 
42+ std::vector<std::pair<int64_t, int64_t>> shape_range = {{1, 16}, {1, 16}};
43+ 
44+ auto input_tensor =
45+ create_desc_shape_range({-1, -1}, ge::DT_FLOAT16, ge::FORMAT_ND, {16, 16}, ge::FORMAT_ND, shape_range);
46+ op.UpdateInputDesc("x", input_tensor);
47+ 
48+ EXPECT_EQ(InferShapeTest(op), ge::GRAPH_SUCCESS);
49+ 
50+ auto output_desc = op.GetOutputDesc("y");
51+ auto output_shape = output_desc.GetShape();
52+ EXPECT_EQ(output_shape.GetDimNum(), 2);
53+ 
54+ std::vector<int64_t> expected_output_shape = {-1, -1};
55+ EXPECT_EQ(output_shape.GetDims(), expected_output_shape);
56+ 
57+}
Aactivation/fast_gelu/tests/ut/op_host/test_fast_gelu_tiling.cpp+121-0
@@ -0,0 +1,121 @@
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 <vector>
13+ 
14+#include <gtest/gtest.h>
15+#include "log/log.h"
16+#include "kernel_run_context_facker.h"
17+#include "test_cube_util.h"
18+#include "exe_graph/runtime/storage_format.h"
19+#include "exe_graph/runtime/storage_shape.h"
20+#include "platform/platform_infos_def.h"
21+#include "atvoss/elewise/elewise_tiling.h"
22+#include "../../../op_host/arch35/fast_gelu_tiling_arch35.h"
23+ 
24+using namespace std;
25+using namespace ge;
26+ 
27+class FastGeluTilingData : public testing::Test {
28+protected:
29+ static void SetUpTestCase()
30+ {
31+ std::cout << "FastGeluTilingData SetUp" << std::endl;
32+ }
33+ 
34+ static void TearDownTestCase()
35+ {
36+ std::cout << "FastGeluTilingData TearDown" << std::endl;
37+ }
38+};
39+ 
40+TEST_F(FastGeluTilingData, test_tiling_fp16_001)
41+{
42+ gert::StorageShape Shape = {{1, 64, 2, 64}, {1, 64, 2, 64}};
43+ 
44+ std::map<std::string, std::string> soc_infos;
45+ std::map<std::string, std::string> aicore_spec;
46+ std::map<std::string, std::string> intrinsics;
47+ std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend910_95"}};
48+ std::string compile_info_string = R"({
49+ "hardware_info": {
50+ "BT_SIZE": 0, "load3d_constraints": "1",
51+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
52+ "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
53+ "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
54+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": 64
55+ }
56+ })";
57+ std::string op_type("FastGelu");
58+ 
59+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
60+ 
61+ fe::PlatFormInfos platform_info;
62+ platform_info.Init();
63+ 
64+ Ops::Base::FastGeluCompileInfo compile_info;
65+ compile_info.coreNum = 64;
66+ compile_info.ubSize = 262144;
67+ 
68+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl("FastGelu")->tiling;
69+ auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl("FastGelu")->tiling_parse;
70+ auto gen_simplifiedkey_func = gert::OpImplRegistry::GetInstance().GetOpImpl("FastGelu")->gen_simplifiedkey;
71+ 
72+ auto kernel_holder =
73+ gert::KernelRunContextFaker()
74+ .KernelIONum(1, 1)
75+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
76+ .Outputs({&compile_info})
77+ .Build();
78+ 
79+ ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
80+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
81+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
82+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
83+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
84+ intrinsics);
85+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("version",
86+ soc_version_infos);
87+ ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
88+ 
89+ // // tilingFunc simulate
90+ auto param = gert::TilingData::CreateCap(4096);
91+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
92+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
93+ ASSERT_NE(param, nullptr);
94+ 
95+ auto holder = gert::TilingContextFaker()
96+ .SetOpType(op_type)
97+ .NodeIoNum(1, 1)
98+ .IrInstanceNum({1})
99+ .InputShapes({&Shape})
100+ .OutputShapes({&Shape})
101+ .CompileInfo(&compile_info)
102+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
103+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
104+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
105+ .TilingData(param.get())
106+ .Workspace(ws_size)
107+ .Build();
108+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
109+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
110+ 
111+ tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
112+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
113+ tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
114+ tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
115+ // workspaces nullptr return failed
116+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
117+ auto tiling_key = tiling_context->GetTilingKey();
118+ ASSERT_EQ(tiling_key, 3);
119+ auto block_dim = tiling_context->GetBlockDim();
120+ ASSERT_EQ(block_dim, 4);
121+}
Mactivation/fast_gelu_grad/CMakeLists.txt+10-13
@@ -1,19 +1,16 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-message(STATUS "=== Debug: start ops.activation.fast_gelu_grad.CMakeLists.txt ")11+# 设置算子定义时支持的芯片类型
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)12+set(SUPPORT_COMPUTE_UNIT "ascend910_95")
12-if(NOT ENABLE_TEST AND NOT BENCHMARK)13+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
13- list(REMOVE_ITEM CURRENT_DIRS tests)14+set(SUPPORT_TILING_DIR "arch35")
14-endif()15+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE fast_gelu_grad ACLNNTYPE aclnn_exclude
15-foreach(SUB_DIR ${CURRENT_DIRS})16+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
16- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17- add_subdirectory(${SUB_DIR})
18- endif()
19-endforeach()
Mactivation/fast_gelu_grad/README.md+61-1
@@ -1,3 +1,63 @@
1# FastGeluGrad1# FastGeluGrad
2 2 
3-本目录仅包含FastGeluGrad算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../CONTRIBUTING.md)。3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT9</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+ 
10+## 功能说明
11+ 
12+- 算子功能：
13+ fast_gelu反向计算。
14+ 
15+## 参数说明
16+ 
17+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
18+ <col style="width: 100px">
19+ <col style="width: 150px">
20+ <col style="width: 280px">
21+ <col style="width: 330px">
22+ <col style="width: 120px">
23+ </colgroup>
24+ <thead>
25+ <tr>
26+ <th>参数名</th>
27+ <th>输入/输出/属性</th>
28+ <th>描述</th>
29+ <th>数据类型</th>
30+ <th>数据格式</th>
31+ </tr></thead>
32+ <tbody>
33+ <tr>
34+ <td>x</td>
35+ <td>输入</td>
36+ <td>输入一个张量。</td>
37+ <td>DT_BF16、FLOAT16、DT_FLOAT</td>
38+ <td>ND</td>
39+ </tr>
40+ <tr>
41+ <td>dy</td>
42+ <td>输入</td>
43+ <td>梯度。</td>
44+ <td>DT_BF16、FLOAT16、DT_FLOAT</td>
45+ <td>ND</td>
46+ </tr>
47+ <tr>
48+ <td>y</td>
49+ <td>输出</td>
50+ <td>输出一个张量。</td>
51+ <td>DT_BF16、FLOAT16、DT_FLOAT</td>
52+ <td>ND</td>
53+ </tr>
54+ </tbody></table>
55+ 
56+## 约束说明
57+ 
58+无
59+ 
60+## 调用说明
61+| 调用方式 | 调用样例 | 说明 |
62+|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
63+| aclnn调用 | [test_aclnn_fast_gelu_grad](examples/arch35/test_aclnn_fast_gelu_grad.cpp) | 通过[aclnnFastGeluBackwardGetWorkspaceSize](docs/aclnnFastGeluGrad.md)接口方式调用fast_gelu算子。 |
Dactivation/fast_gelu_grad/docs/aclnnFastGeluBackward.md+0-342
@@ -1,342 +0,0 @@
1-# aclnnFastGeluBackward
2- 
3-## 产品支持情况
4- 
5-|产品 | 是否支持 |
6-|:-------------------------|:----------:|
7-| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
9- 
10-## 功能说明
11-[FastGelu](../../fast_gelu/docs/aclnnFastGelu.md)的反向计算。
12-FastGeluBackward定义如下：
13- 
14-$$
15-dx = dy * ((\frac{1}{e^{-1.702 x}+1} - 1) * -1.702x + 1) * \frac{1}{(e^{-1.702 x}+1)}
16-$$
17-
18-## 函数原型
19- 
20-每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnFastGeluBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnFastGeluBackward”接口执行计算。
21- 
22-```Cpp
23-aclnnStatus aclnnFastGeluBackwardGetWorkspaceSize(
24- const aclTensor *gradOutput,
25- const aclTensor *self,
26- aclTensor *gradInput,
27- uint64_t *workspaceSize,
28- aclOpExecutor **executor)
29-```
30- 
31-```Cpp
32-aclnnStatus aclnnFastGeluBackward(
33- void *workspace,
34- uint64_t workspaceSize,
35- aclOpExecutor *executor,
36- aclrtStream stream)
37-```
38- 
39-## aclnnFastGeluBackwardGetWorkspaceSize
40- 
41-- **参数说明：**
42- 
43- <table style="undefined;table-layout: fixed; width: 1450px"><colgroup>
44- <col style="width: 171px">
45- <col style="width: 115px">
46- <col style="width: 220px">
47- <col style="width: 280px">
48- <col style="width: 177px">
49- <col style="width: 104px">
50- <col style="width: 238px">
51- <col style="width: 145px">
52- </colgroup>
53- <thead>
54- <tr>
55- <th>参数名</th>
56- <th>输入/输出</th>
57- <th>描述</th>
58- <th>使用说明</th>
59- <th>数据类型</th>
60- <th>数据格式</th>
61- <th>维度(shape)</th>
62- <th>非连续Tensor</th>
63- </tr></thead>
64- <tbody>
65- <tr>
66- <td>gradOutput</td>
67- <td>输入</td>
68- <td>反向传播的梯度数据，公式中的dy。</td>
69- <td><ul><li>与self、gradInput的数据格式一致。</li><li>与self、gradInput的shape一致。</li><li>与self、gradInput的数据类型一致。</li></ul></td>
70- <td>FLOAT16、FLOAT32、BFLOAT16</td>
71- <td>ND</td>
72- <td>0-8</td>
73- <td>√</td>
74- </tr>
75- <tr>
76- <td>self</td>
77- <td>输入</td>
78- <td>表示输入张量，与正向的self相同，公式中的x。</td>
79- <td><ul><li>与gradOutput、gradInput的数据格式一致。</li><li>与gradOutput、gradInput的shape一致。</li><li>与gradOutput、gradInput的数据类型一致。</li></ul></td>
80- <td>FLOAT16、FLOAT32、BFLOAT16</td>
81- <td>ND</td>
82- <td>0-8</td>
83- <td>√</td>
84- </tr>
85- <tr>
86- <td>gradInput</td>
87- <td>输出</td>
88- <td>表示输出张量，公式中的dx。</td>
89- <td><ul><li>与gradOutput、self的数据格式一致。</li><li>与gradOutput、self的shape一致。</li><li>与gradOutput、self的数据类型一致。</li></ul></td>
90- <td>FLOAT16、FLOAT32、BFLOAT16</td>
91- <td>ND</td>
92- <td>0-8</td>
93- <td>√</td>
94- </tr>
95- <tr>
96- <td>workspaceSize</td>
97- <td>输出</td>
98- <td>返回需要在Device侧申请的workspace大小。</td>
99- <td>-</td>
100- <td>-</td>
101- <td>-</td>
102- <td>-</td>
103- <td>-</td>
104- </tr>
105- <tr>
106- <td>executor</td>
107- <td>输出</td>
108- <td>返回op执行器，包含了算子计算流程。</td>
109- <td>-</td>
110- <td>-</td>
111- <td>-</td>
112- <td>-</td>
113- <td>-</td>
114- </tr>
115- </tbody>
116- </table>
117- 
118-- **返回值：**
119- 
120- aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
121- 第一段接口会完成入参校验，出现以下场景时报错：
122- <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
123- <col style="width: 272px">
124- <col style="width: 103px">
125- <col style="width: 604px">
126- </colgroup>
127- <thead>
128- <tr>
129- <th>返回码</th>
130- <th>错误码</th>
131- <th>描述</th>
132- </tr>
133- </thead>
134- <tbody>
135- <tr>
136- <td>ACLNN_ERR_PARAM_NULLPTR</td>
137- <td>161001</td>
138- <td>gradOutput、self或gradInput是空指针。</td>
139- </tr>
140- <tr>
141- <td rowspan="8">ACLNN_ERR_PARAM_INVALID</td>
142- <td rowspan="8">161002</td>
143- <td>gradOutput、self和gradInput的数据类型、数据格式不一致。</td>
144- </tr>
145- <tr>
146- <td>gradInput的数据类型不在支持范围内。</td>
147- </tr>
148- <tr>
149- <td>gradOutput、self和gradInput的shape不一致。</td>
150- </tr>
151- <tr>
152- <td>self的shape维度大于8维。</td>
153- </tr>
154- </tbody></table>
155- 
156-## aclnnFastGeluBackward
157- 
158-- **参数说明：**
159- 
160- <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
161- <col style="width: 173px">
162- <col style="width: 112px">
163- <col style="width: 668px">
164- </colgroup>
165- <thead>
166- <tr>
167- <th>参数名</th>
168- <th>输入/输出</th>
169- <th>描述</th>
170- </tr></thead>
171- <tbody>
172- <tr>
173- <td>workspace</td>
174- <td>输入</td>
175- <td>在Device侧申请的workspace内存地址。</td>
176- </tr>
177- <tr>
178- <td>workspaceSize</td>
179- <td>输入</td>
180- <td>在Device侧申请的workspace大小，由第一段接口aclnnFastGeluBackwardGetWorkspaceSize获取。</td>
181- </tr>
182- <tr>
183- <td>executor</td>
184- <td>输入</td>
185- <td>op执行器，包含了算子计算流程。</td>
186- </tr>
187- <tr>
188- <td>stream</td>
189- <td>输入</td>
190- <td>指定执行任务的Stream。</td>
191- </tr>
192- </tbody>
193- </table>
194- 
195-- **返回值：**
196- 
197- aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
198- 
199-## 约束说明
200- 
201-- 确定性计算：
202- - aclnnFastGeluBackward默认确定性实现。
203- 
204-## 调用示例
205- 
206-示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
207- 
208-```Cpp
209-#include <iostream>
210-#include <vector>
211-#include "acl/acl.h"
212-#include "aclnnop/aclnn_fast_gelu_backward.h"
213- 
214-#define CHECK_RET(cond, return_expr) \
215- do { \
216- if (!(cond)) { \
217- return_expr; \
218- } \
219- } while (0)
220- 
221-#define LOG_PRINT(message, ...) \
222- do { \
223- printf(message, ##__VA_ARGS__); \
224- } while (0)
225- 
226-int64_t GetShapeSize(const std::vector<int64_t>& shape) {
227- int64_t shapeSize = 1;
228- for (auto i : shape) {
229- shapeSize *= i;
230- }
231- return shapeSize;
232-}
233- 
234-int Init(int32_t deviceId, aclrtStream* stream) {
235- // 固定写法，资源初始化
236- auto ret = aclInit(nullptr);
237- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
238- ret = aclrtSetDevice(deviceId);
239- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
240- ret = aclrtCreateStream(stream);
241- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
242- return 0;
243-}
244- 
245-template <typename T>
246-int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
247- aclDataType dataType, aclTensor** tensor) {
248- auto size = GetShapeSize(shape) * sizeof(T);
249- // 调用aclrtMalloc申请device侧内存
250- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
251- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
252- // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
253- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
254- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
255- 
256- // 计算连续tensor的strides
257- std::vector<int64_t> strides(shape.size(), 1);
258- for (int64_t i = shape.size() - 2; i >= 0; i--) {
259- strides[i] = shape[i + 1] * strides[i + 1];
260- }
261- 
262- // 调用aclCreateTensor接口创建aclTensor
263- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
264- shape.data(), shape.size(), *deviceAddr);
265- return 0;
266-}
267- 
268-int main() {
269- // 1. （固定写法）device/stream初始化，参考acl API手册
270- // 根据自己的实际device填写deviceId
271- int32_t deviceId = 0;
272- aclrtStream stream;
273- auto ret = Init(deviceId, &stream);
274- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
275- // 2. 构造输入与输出，需要根据API的接口自定义构造
276- std::vector<int64_t> gradOutputShape = {3, 3};
277- std::vector<int64_t> selfShape = {3, 3};
278- std::vector<int64_t> gradInputShape = {3, 3};
279- void* gradOutputDeviceAddr = nullptr;
280- aclTensor* gradOutput = nullptr;
281- void* selfDeviceAddr = nullptr;
282- aclTensor* self = nullptr;
283- void* gradInputDeviceAddr = nullptr;
284- aclTensor* gradInput = nullptr;
285- std::vector<float> gradOutputHostData(9, 0);
286- std::vector<float> selfHostData(9, 0);
287- std::vector<float> gradInputHostData(9, 0);
288- // 创建gradOutput、self和gradInput aclTensor
289- ret = CreateAclTensor(gradOutputHostData, selfShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT, &gradOutput);
290- CHECK_RET(ret == ACL_SUCCESS, return ret);
291- ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
292- CHECK_RET(ret == ACL_SUCCESS, return ret);
293- ret = CreateAclTensor(gradInputHostData, gradInputShape, &gradInputDeviceAddr, aclDataType::ACL_FLOAT, &gradInput);
294- CHECK_RET(ret == ACL_SUCCESS, return ret);
295- 
296- // 3. 调用CANN算子库API，需要修改为具体的Api名称
297- uint64_t workspaceSize = 0;
298- aclOpExecutor* executor;
299- // 调用aclnnFastGeluBackward第一段接口
300- ret = aclnnFastGeluBackwardGetWorkspaceSize(gradOutput, self, gradInput, &workspaceSize, &executor);
301- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFastGeluBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
302- // 根据第一段接口计算出的workspaceSize申请device内存
303- void* workspaceAddr = nullptr;
304- if (workspaceSize > 0) {
305- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
306- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
307- }
308- // 调用aclnnFastGeluBackward第二段接口
309- ret = aclnnFastGeluBackward(workspaceAddr, workspaceSize, executor, stream);
310- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnFastGeluBackward failed. ERROR: %d\n", ret); return ret);
311- // 4. （固定写法）同步等待任务执行结束
312- ret = aclrtSynchronizeStream(stream);
313- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
314- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
315- auto size = GetShapeSize(gradInputShape);
316- std::vector<float> resultData(size, 0);
317- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), gradInputDeviceAddr,
318- size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
319- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
320- for (int64_t i = 0; i < size; i++) {
321- LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
322- }
323- 
324- // 6. 释放aclTensor，需要根据具体API的接口定义修改
325- aclDestroyTensor(gradOutput);
326- aclDestroyTensor(self);
327- aclDestroyTensor(gradInput);
328- 
329- // 7. 释放device资源，需要根据具体API的接口定义修改
330- aclrtFree(gradOutputDeviceAddr);
331- aclrtFree(selfDeviceAddr);
332- aclrtFree(gradInputDeviceAddr);
333- if (workspaceSize > 0) {
334- aclrtFree(workspaceAddr);
335- }
336- aclrtDestroyStream(stream);
337- aclrtResetDevice(deviceId);
338- aclFinalize();
339- return 0;
340-}
341-```
342- 
Aactivation/fast_gelu_grad/docs/aclnnFastGeluGrad.md+136-0
@@ -0,0 +1,136 @@
1+# FastGeluGrad
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT9</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+ 
10+## 功能说明
11+ 
12+算子功能：aclnnFastGelu的反向计算。
13+ 
14+## 函数原型
15+ 
16+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnFastGeluBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnFastGeluBackward”接口执行计算。
17+ 
18+## aclnnFastGeluBackwardGetWorkspaceSize
19+ 
20+- **参数说明：**
21+ 
22+ <table style="undefined;table-layout: fixed; width: 1420px"><colgroup>
23+ <col style="width: 173px">
24+ <col style="width: 120px">
25+ <col style="width: 222px">
26+ <col style="width: 338px">
27+ <col style="width: 156px">
28+ <col style="width: 104px">
29+ <col style="width: 162px">
30+ <col style="width: 145px">
31+ </colgroup>
32+ <thead>
33+ <tr>
34+ <th>参数名</th>
35+ <th>输入/输出</th>
36+ <th>描述</th>
37+ <th>使用说明</th>
38+ <th>数据类型</th>
39+ <th>数据格式</th>
40+ <th>维度(shape)</th>
41+ <th>非连续Tensor</th>
42+ </tr></thead>
43+ <tbody>
44+ <tr>
45+ <td>self</td>
46+ <td>输入</td>
47+ <td>输入张量。</td>
48+ <td>数据类型支持FLOAT16、FLOAT32、BFLOAT16</td>
49+ <td>FLOAT16、FLOAT32、BFLOAT16</td>
50+ <td>ND</td>
51+ <td>(N)</td>
52+ <td>√</td>
53+ </tr>
54+ <tr>
55+ <td>out</td>
56+ <td>输出</td>
57+ <td>输出张量。</td>
58+ <td>数据类型支持FLOAT16、FLOAT32、BFLOAT16</td>
59+ <td>FLOAT16、FLOAT32、BFLOAT16</td>
60+ <td>ND</td>
61+ <td>(N)</td>
62+ <td>√</td>
63+ </tr>
64+ <tr>
65+ <td>workspaceSize</td>
66+ <td>输出</td>
67+ <td>返回需要在Device侧申请的workspace大小。</td>
68+ <td>-</td>
69+ <td>-</td>
70+ <td>-</td>
71+ <td>-</td>
72+ <td>-</td>
73+ </tr>
74+ <tr>
75+ <td>executor</td>
76+ <td>输出</td>
77+ <td>返回需要在Device侧申请的workspace大小。</td>
78+ <td>-</td>
79+ <td>-</td>
80+ <td>-</td>
81+ <td>-</td>
82+ <td>-</td>
83+ </tr>
84+ </tbody></table>
85+ 
86+- **返回值：**
87+ 
88+ aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
89+ 
90+ 
91+## aclnnFastGeluBackward
92+ 
93+- **参数说明：**
94+ 
95+ <table style="undefined;table-layout: fixed; width: 1244px"><colgroup>
96+ <col style="width: 200px">
97+ <col style="width: 162px">
98+ <col style="width: 882px">
99+ </colgroup>
100+ <thead>
101+ <tr>
102+ <th>参数名</th>
103+ <th>输入/输出</th>
104+ <th>描述</th>
105+ </tr></thead>
106+ <tbody>
107+ <tr>
108+ <td>workspace</td>
109+ <td>输入</td>
110+ <td>在Device侧申请的workspace内存地址。</td>
111+ </tr>
112+ <tr>
113+ <td>workspaceSize</td>
114+ <td>输入</td>
115+ <td>在Device侧申请的workspace大小，由第一段接口aclnnFastGeluGetWorkspaceSize获取。</td>
116+ </tr>
117+ <tr>
118+ <td>executor</td>
119+ <td>输入</td>
120+ <td>op执行器，包含了算子计算流程。</td>
121+ </tr>
122+ <tr>
123+ <td>stream</td>
124+ <td>输入</td>
125+ <td>指定执行任务的Stream。</td>
126+ </tr>
127+ </tbody>
128+ </table>
129+ 
130+- **返回值：**
131+ 
132+ aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
133+ 
134+## 约束说明
135+ 
136+无
Ractivation/fast_gelu_grad/examples/test_aclnn_fast_gelu_backward.cpp→activation/fast_gelu_grad/examples/arch35/test_aclnn_fast_gelu_backward.cpp+0-0
文件重命名但无更改。
Ractivation/fast_gelu_grad/op_host/op_api/aclnn_fast_gelu_backward.cpp→activation/fast_gelu_grad/op_api/aclnn_fast_gelu_backward.cpp+0-0
文件重命名但无更改。
Ractivation/fast_gelu_grad/op_host/op_api/aclnn_fast_gelu_backward.h→activation/fast_gelu_grad/op_api/aclnn_fast_gelu_backward.h+0-0
文件重命名但无更改。
Ractivation/fast_gelu_grad/op_host/op_api/fast_gelu_grad.cpp→activation/fast_gelu_grad/op_api/fast_gelu_grad.cpp+0-0
文件重命名但无更改。
Ractivation/fast_gelu_grad/op_host/op_api/fast_gelu_grad.h→activation/fast_gelu_grad/op_api/fast_gelu_grad.h+0-0
文件重命名但无更改。
Ractivation/fast_gelu/tests/ut/op_host/test_aclnn_fast_gelu.cpp→activation/fast_gelu_grad/op_graph/fast_gelu_grad_proto.h+22-23
@@ -7,31 +7,30 @@
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10+#ifndef OPS_BUILT_IN_OP_PROTO_INC_NONLINEAR_FUC_OPS_H_
11+#define OPS_BUILT_IN_OP_PROTO_INC_NONLINEAR_FUC_OPS_H_
10 12 
11-#include "gtest/gtest.h"13+#include "graph/operator_reg.h"
12-#include "../../../op_host/op_api/aclnn_fast_gelu.h"
13-#include "op_api_ut_common/tensor_desc.h"
14-#include "op_api_ut_common/op_api_ut.h"
15 14 
16-class l2_fast_gelu_test : public testing::Test {15+namespace ge {
17-protected:
18- static void SetUpTestCase() {
19- std::cout << "l2_fast_gelu_test SetUp" << std::endl;
20- }
21 16 
22- static void TearDownTestCase() { std::cout << "l2_fast_gelu_test TearDown" << std::endl; }17+/**
23-};18+*@brief Computes the gradient for the fast_gelu of "x" .
24 19 
25-// 正常路径，float3220+*@par Inputs:
26-TEST_F(l2_fast_gelu_test, l2_fast_gelu_test_005) {21+*Two inputs, including:
27- auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);22+* @li dy: A Tensor. Must be one of the following types: bfloat16, float16, float32
28- auto outDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);23+* @li x: A Tensor of the same type as "dy" . \n
29 24 
30- auto ut = OP_API_UT(aclnnFastGelu, INPUT(selfDesc), OUTPUT(outDesc));25+*@par Outputs:
31- 26+*z: A Tensor. Has the same type as "dy".
32- uint64_t workspaceSize = 0;27+*@par Third-party framework compatibility
33- aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);28+*Compatible with the TensorFlow operator FastGeluGrad
34- EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);29+*/
35- 30+REG_OP(FastGeluGrad)
36- //ut.TestPrecision();31+ .INPUT(dy, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
37-}32+ .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
33+ .OUTPUT(z, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
34+ .OP_END_FACTORY_REG(FastGeluGrad)
35+} // namespace ge
36+#endif // OPS_BUILT_IN_OP_PROTO_INC_NONLINEAR_FUC_OPS_H_
Aactivation/fast_gelu_grad/op_host/arch35/fast_gelu_grad_tiling_arch35.cpp+162-0
@@ -0,0 +1,162 @@
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 <graph/utils/type_utils.h>
13+#include "log/log.h"
14+#include "platform/platform_ascendc.h"
15+#include "register/op_def_registry.h"
16+#include "register/tilingdata_base.h"
17+#include "activation/fast_gelu_grad/op_kernel/arch35/fast_gelu_grad_dag.h"
18+#include "activation/fast_gelu_grad/op_kernel/arch35/fast_gelu_grad_struct.h"
19+#include "atvoss/elewise/elewise_tiling.h"
20+#include "atvoss/broadcast/broadcast_tiling.h"
21+#include "fast_gelu_grad_tiling_arch35.h"
22+ 
23+using namespace FastGeluGradOp;
24+ 
25+namespace optiling {
26+const int64_t SYSWORKSPACE = 16777216; // 16 * 1024 * 1024
27+ 
28+ge::graphStatus FastGeluGradTiling::CalcInputDtype()
29+{
30+ auto inputDesc = tilingContext->GetInputDesc(0);
31+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, inputDesc);
32+ this->inputDtype = inputDesc->GetDataType();
33+ OP_CHECK_IF(
34+ this->inputDtype != ge::DT_FLOAT16 && this->inputDtype != ge::DT_BF16 && this->inputDtype != ge::DT_FLOAT,
35+ OP_LOGE(tilingContext->GetNodeName(), "self dtype not support"), return ge::GRAPH_FAILED);
36+ 
37+ auto inputDesc1 = tilingContext->GetInputDesc(1);
38+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, inputDesc1);
39+ this->inputDtype1 = inputDesc1->GetDataType();
40+ OP_CHECK_IF(
41+ this->inputDtype1 != ge::DT_FLOAT16 && this->inputDtype1 != ge::DT_BF16 && this->inputDtype1 != ge::DT_FLOAT,
42+ OP_LOGE(tilingContext->GetNodeName(), "input dy dtype not support"), return ge::GRAPH_FAILED);
43+ 
44+ OP_CHECK_IF(
45+ this->inputDtype1 != this->inputDtype, OP_LOGE(tilingContext->GetNodeName(), "input dy dtype not support"),
46+ return ge::GRAPH_FAILED);
47+ 
48+ return ge::GRAPH_SUCCESS;
49+}
50+ 
51+ge::graphStatus FastGeluGradTiling::CalcOutputDtype()
52+{
53+ auto outputDesc = tilingContext->GetOutputDesc(0);
54+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outputDesc);
55+ this->outputDtype = outputDesc->GetDataType();
56+ OP_CHECK_IF(
57+ this->outputDtype != ge::DT_FLOAT16 && this->outputDtype != ge::DT_BF16 && this->outputDtype != ge::DT_FLOAT,
58+ OP_LOGE(tilingContext->GetNodeName(), "self dtype not support"), return ge::GRAPH_FAILED);
59+ OP_CHECK_IF(
60+ this->outputDtype != this->inputDtype, OP_LOGE(tilingContext->GetNodeName(), "out dtype not same as self"),
61+ return ge::GRAPH_FAILED);
62+ return ge::GRAPH_SUCCESS;
63+}
64+ 
65+ge::graphStatus FastGeluGradTiling::CheckShape()
66+{
67+ auto selfStorageShape = tilingContext->GetInputShape(0);
68+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, selfStorageShape);
69+ const gert::Shape& inputXShape = Ops::Base::EnsureNotScalar(selfStorageShape->GetStorageShape());
70+ 
71+ auto dyStorageShape = tilingContext->GetInputShape(1);
72+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, dyStorageShape);
73+ const gert::Shape& inputDyShape = Ops::Base::EnsureNotScalar(dyStorageShape->GetStorageShape());
74+ 
75+ auto outStorageShape = tilingContext->GetOutputShape(0);
76+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, outStorageShape);
77+ const gert::Shape& outputShape = Ops::Base::EnsureNotScalar(outStorageShape->GetStorageShape());
78+ 
79+ OP_CHECK_IF(
80+ inputXShape != inputDyShape, OP_LOGE(tilingContext->GetNodeName(), "input x and input dy shape not same"),
81+ return ge::GRAPH_FAILED);
82+ OP_CHECK_IF(
83+ inputXShape != outputShape, OP_LOGE(tilingContext->GetNodeName(), "input x and output z shape not same"),
84+ return ge::GRAPH_FAILED);
85+ return ge::GRAPH_SUCCESS;
86+}
87+ 
88+std::string FastGeluGradTiling::DataTypeToSerialString(const ge::DataType type) const
89+{
90+ const auto it = DATATYPE_TO_STRING_MAP.find(type);
91+ if (it != DATATYPE_TO_STRING_MAP.end()) {
92+ return it->second;
93+ } else {
94+ OP_LOGE("Neg", "datatype %d not find", type);
95+ return "UNDEFINED";
96+ }
97+}
98+ 
99+ge::graphStatus FastGeluGradTiling::RunTiling()
100+{
101+ ElewiseBaseTiling elewiseBaseTiling(tilingContext);
102+ // 获取tiling计算所需的参数
103+ ge::graphStatus status = ge::GRAPH_FAILED;
104+ status = CalcInputDtype();
105+ OP_CHECK_IF(
106+ status == ge::GRAPH_FAILED, OP_CHECK_NULL_WITH_CONTEXT(tilingContext, "get input dtype failed"),
107+ return ge::GRAPH_FAILED);
108+ status = CalcOutputDtype();
109+ OP_CHECK_IF(
110+ status == ge::GRAPH_FAILED, OP_CHECK_NULL_WITH_CONTEXT(tilingContext, "get output dtype failed"),
111+ return ge::GRAPH_FAILED);
112+ status = CheckShape();
113+ OP_CHECK_IF(
114+ status == ge::GRAPH_FAILED, OP_CHECK_NULL_WITH_CONTEXT(tilingContext, "check shape failed"),
115+ return ge::GRAPH_FAILED);
116+ auto tiling = tilingContext->GetTilingData<EleBaseTilingDataV2>();
117+ OP_CHECK_IF(
118+ (tiling == nullptr), OP_LOGE(tilingContext->GetNodeName(), "Get FastGeluGradTiling from GE context failed"),
119+ return ge::GRAPH_FAILED);
120+ if (this->outputDtype == ge::DT_FLOAT16) {
121+ dType = TPL_FP16;
122+ status = elewiseBaseTiling.DoTiling<FastGeluGradDag::FastGeluGradNeedCast<half>::OpDag>(*tiling);
123+ } else if (this->outputDtype == ge::DT_BF16) {
124+ dType = TPL_BF16;
125+ status = elewiseBaseTiling.DoTiling<FastGeluGradDag::FastGeluGradNeedCast<bfloat16_t>::OpDag>(*tiling);
126+ } else if (this->outputDtype == ge::DT_FLOAT) {
127+ dType = TPL_FP32;
128+ status = elewiseBaseTiling.DoTiling<FastGeluGradDag::FastGeluGradNoCast<float>::OpDag>(*tiling);
129+ } else {
130+ OP_LOGE(tilingContext->GetNodeName(), "elewiseBaseTiling DoTiling failed.");
131+ return ge::GRAPH_FAILED;
132+ }
133+ OP_CHECK_IF(
134+ status == ge::GRAPH_FAILED, OP_CHECK_NULL_WITH_CONTEXT(tilingContext, "elewiseBaseTiling failed"),
135+ return ge::GRAPH_FAILED);
136+ const uint64_t tilingKey = GET_TPL_TILING_KEY(tiling->scheMode, dType);
137+ OP_LOGD(tilingContext->GetNodeName(), "[TilingData] : tilingKey=%ld.", tilingKey);
138+ tilingContext->SetTilingKey(tilingKey);
139+ tilingContext->SetBlockDim(tiling->blockNum);
140+ size_t usr_workspaceSize = 0;
141+ size_t sys_workspaceSize = SYSWORKSPACE;
142+ size_t* current_workspace = tilingContext->GetWorkspaceSizes(1);
143+ current_workspace[0] = sys_workspaceSize + usr_workspaceSize;
144+ return ge::GRAPH_SUCCESS;
145+}
146+ 
147+static ge::graphStatus TilingFuncFastGeluGrad(gert::TilingContext* tilingContext)
148+{
149+ OP_LOGD(tilingContext->GetNodeName(), "START FastGeluGrad AscendC Tiling \n");
150+ FastGeluGradTiling FastGeluGradTiling(tilingContext);
151+ return FastGeluGradTiling.RunTiling();
152+}
153+ 
154+ge::graphStatus TilingPrepareForFastGeluGrad([[maybe_unused]] gert::TilingParseContext* context)
155+{
156+ return ge::GRAPH_SUCCESS;
157+}
158+ 
159+IMPL_OP_OPTILING(FastGeluGrad)
160+ .Tiling(TilingFuncFastGeluGrad)
161+ .TilingParse<FastGeluGradCompileInfo>(TilingPrepareForFastGeluGrad);
162+} // namespace optiling
Aactivation/fast_gelu_grad/op_host/arch35/fast_gelu_grad_tiling_arch35.h+79-0
@@ -0,0 +1,79 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_FAST_GELU_GRAD_TILING_H
12+#define OPS_BUILD_IN_OP_TILING_RUNTIME_FAST_GELU_GRAD_TILING_H
13+#include "register/tilingdata_base.h"
14+#include "register/op_impl_registry.h"
15+ 
16+namespace optiling {
17+using namespace Ops::Base;
18+ 
19+class FastGeluGradTiling {
20+public:
21+ explicit FastGeluGradTiling(gert::TilingContext* context) : tilingContext(context){};
22+ ge::graphStatus RunTiling();
23+ 
24+protected:
25+ ge::graphStatus CalcInputDtype();
26+ ge::graphStatus CalcOutputDtype();
27+ ge::graphStatus CheckShape();
28+ std::string DataTypeToSerialString(const ge::DataType type) const;
29+ 
30+private:
31+ uint64_t dType = 0;
32+ ge::DataType inputDtype;
33+ ge::DataType inputDtype1;
34+ ge::DataType outputDtype;
35+ gert::TilingContext* tilingContext;
36+};
37+ 
38+const std::map<ge::DataType, std::string> DATATYPE_TO_STRING_MAP = {
39+ {ge::DT_UNDEFINED, "DT_UNDEFINED"}, // Used to indicate a DataType field has not been set.
40+ {ge::DT_FLOAT, "DT_FLOAT"}, // float type
41+ {ge::DT_FLOAT16, "DT_FLOAT16"}, // fp16 type
42+ {ge::DT_INT8, "DT_INT8"}, // int8 type
43+ {ge::DT_INT16, "DT_INT16"}, // int16 type
44+ {ge::DT_UINT16, "DT_UINT16"}, // uint16 type
45+ {ge::DT_UINT8, "DT_UINT8"}, // uint8 type
46+ {ge::DT_INT32, "DT_INT32"}, // uint32 type
47+ {ge::DT_INT64, "DT_INT64"}, // int64 type
48+ {ge::DT_UINT32, "DT_UINT32"}, // unsigned int32
49+ {ge::DT_UINT64, "DT_UINT64"}, // unsigned int64
50+ {ge::DT_BOOL, "DT_BOOL"}, // bool type
51+ {ge::DT_DOUBLE, "DT_DOUBLE"}, // double type
52+ {ge::DT_DUAL, "DT_DUAL"}, // dual output type
53+ {ge::DT_DUAL_SUB_INT8, "DT_DUAL_SUB_INT8"}, // dual output int8 type
54+ {ge::DT_DUAL_SUB_UINT8, "DT_DUAL_SUB_UINT8"}, // dual output uint8 type
55+ {ge::DT_COMPLEX32, "DT_COMPLEX32"}, // complex32 type
56+ {ge::DT_COMPLEX64, "DT_COMPLEX64"}, // complex64 type
57+ {ge::DT_COMPLEX128, "DT_COMPLEX128"}, // complex128 type
58+ {ge::DT_QINT8, "DT_QINT8"}, // qint8 type
59+ {ge::DT_QINT16, "DT_QINT16"}, // qint16 type
60+ {ge::DT_QINT32, "DT_QINT32"}, // qint32 type
61+ {ge::DT_QUINT8, "DT_QUINT8"}, // quint8 type
62+ {ge::DT_QUINT16, "DT_QUINT16"}, // quint16 type
63+ {ge::DT_RESOURCE, "DT_RESOURCE"}, // resource type
64+ {ge::DT_STRING_REF, "DT_STRING_REF"}, // string ref type
65+ {ge::DT_STRING, "DT_STRING"}, // string type
66+ {ge::DT_VARIANT, "DT_VARIANT"}, // dt_variant type
67+ {ge::DT_BF16, "DT_BFLOAT16"}, // dt_bfloat16 type
68+ {ge::DT_INT4, "DT_INT4"}, // dt_variant type
69+ {ge::DT_UINT1, "DT_UINT1"}, // dt_variant type
70+ {ge::DT_INT2, "DT_INT2"}, // dt_variant type
71+ {ge::DT_UINT2, "DT_UINT2"} // dt_variant type
72+};
73+ 
74+struct FastGeluGradCompileInfo {
75+ uint64_t coreNum = 0;
76+ uint64_t ubSize = 0;
77+};
78+} // namespace optiling
79+#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_FAST_GELU_GRAD_TILING_H
Aactivation/fast_gelu_grad/op_host/config/ascend910_95/fast_gelu_grad_binary.json+125-0
@@ -0,0 +1,125 @@
1+{
2+ "op_type": "FastGeluGrad",
3+ "op_list": [
4+ {
5+ "bin_filename": "FastGeluGrad_bf16",
6+ "inputs": [
7+ {
8+ "name": "dy",
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+ "name": "x",
20+ "index": 1,
21+ "dtype": "bfloat16",
22+ "format": "ND",
23+ "paramType": "required",
24+ "shape": [
25+ -2
26+ ],
27+ "format_match_mode": "FormatAgnostic"
28+ }
29+ ],
30+ "outputs": [
31+ {
32+ "name": "z",
33+ "index": 0,
34+ "dtype": "bfloat16",
35+ "format": "ND",
36+ "paramType": "required",
37+ "shape": [
38+ -2
39+ ],
40+ "format_match_mode": "FormatAgnostic"
41+ }
42+ ]
43+ },
44+ {
45+ "bin_filename": "FastGeluGrad_fp16",
46+ "inputs": [
47+ {
48+ "name": "dy",
49+ "index": 0,
50+ "dtype": "float16",
51+ "format": "ND",
52+ "paramType": "required",
53+ "shape": [
54+ -2
55+ ],
56+ "format_match_mode": "FormatAgnostic"
57+ },
58+ {
59+ "name": "x",
60+ "index": 1,
61+ "dtype": "float16",
62+ "format": "ND",
63+ "paramType": "required",
64+ "shape": [
65+ -2
66+ ],
67+ "format_match_mode": "FormatAgnostic"
68+ }
69+ ],
70+ "outputs": [
71+ {
72+ "name": "z",
73+ "index": 0,
74+ "dtype": "float16",
75+ "format": "ND",
76+ "paramType": "required",
77+ "shape": [
78+ -2
79+ ],
80+ "format_match_mode": "FormatAgnostic"
81+ }
82+ ]
83+ },
84+ {
85+ "bin_filename": "FastGeluGrad_fp32",
86+ "inputs": [
87+ {
88+ "name": "dy",
89+ "index": 0,
90+ "dtype": "float32",
91+ "format": "ND",
92+ "paramType": "required",
93+ "shape": [
94+ -2
95+ ],
96+ "format_match_mode": "FormatAgnostic"
97+ },
98+ {
99+ "name": "x",
100+ "index": 1,
101+ "dtype": "float32",
102+ "format": "ND",
103+ "paramType": "required",
104+ "shape": [
105+ -2
106+ ],
107+ "format_match_mode": "FormatAgnostic"
108+ }
109+ ],
110+ "outputs": [
111+ {
112+ "name": "z",
113+ "index": 0,
114+ "dtype": "float32",
115+ "format": "ND",
116+ "paramType": "required",
117+ "shape": [
118+ -2
119+ ],
120+ "format_match_mode": "FormatAgnostic"
121+ }
122+ ]
123+ }
124+ ]
125+}
Aactivation/fast_gelu_grad/op_host/config/ascend910_95/fast_gelu_grad_simplified_key.ini+2-0
@@ -0,0 +1,2 @@
1+[FastGeluGrad]
2+default=0
Aactivation/fast_gelu_grad/op_host/fast_gelu_grad_def.cpp+45-0
@@ -0,0 +1,45 @@
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 "register/op_def_registry.h"
12+ 
13+namespace ops {
14+class FastGeluGrad : public OpDef {
15+public:
16+ explicit FastGeluGrad(const char* name) : OpDef(name)
17+ {
18+ this->Input("dy")
19+ .ParamType(REQUIRED)
20+ .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT})
21+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
22+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
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("z")
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+ OpAICoreConfig aicoreConfig;
34+ aicoreConfig.DynamicCompileStaticFlag(true)
35+ .DynamicFormatFlag(false)
36+ .DynamicRankSupportFlag(true)
37+ .DynamicShapeSupportFlag(true)
38+ .NeedCheckSupportFlag(false)
39+ .ExtendCfgInfo("opFile.value", "fast_gelu_grad_apt");
40+ this->AICore().AddConfig("ascend910_95", aicoreConfig);
41+ }
42+};
43+ 
44+OP_ADD(FastGeluGrad);
45+} // namespace ops
Aactivation/fast_gelu_grad/op_host/fast_gelu_grad_infershape.cpp+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,
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 "infershape_broadcast_util.h"
12+#include "register/op_impl_registry.h"
13+#include "log/log.h"
14+ 
15+using namespace ge;
16+namespace ops {
17+IMPL_OP_INFERSHAPE(FastGeluGrad).InferShape(Ops::Base::InferShape4Broadcast);
18+} // namespace ops
Aactivation/fast_gelu_grad/op_kernel/arch35/fast_gelu_grad_dag.h+103-0
@@ -0,0 +1,103 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef CANN_CUSTOM_OPS_FAST_GELU_GRAD_DAG_H
12+#define CANN_CUSTOM_OPS_FAST_GELU_GRAD_DAG_H
13+#include "atvoss/util/dag.h"
14+#include "atvoss/util/vec.h"
15+#include "atvoss/util/placeholder.h"
16+ 
17+namespace FastGeluGradDag {
18+using namespace Ops::Base;
19+ 
20+template <class T>
21+struct FastGeluGradCustom : public Vec::ElemwiseBinaryOP<T, T, T> {
22+ __aicore__ inline FastGeluGradCustom(
23+ LocalTensor<T>& dst, LocalTensor<T>& src1, LocalTensor<T>& src2, uint32_t count)
24+ {
25+#ifdef __CCE_AICORE__
26+ uint32_t dtypeSize = sizeof(T);
27+ uint32_t VL = AscendC::VECTOR_REG_WIDTH / dtypeSize;
28+ uint16_t loopNum = CeilDivision(count, VL);
29+ uint32_t vlSize = VL;
30+ T value1 = 1.702;
31+ T value2 = -1.702;
32+ T value3 = 1.0;
33+ T value4 = -1.0;
34+ __VEC_SCOPE__
35+ {
36+ __ubuf__ T* src1Addr = (__ubuf__ T*)src1.GetPhyAddr();
37+ __ubuf__ T* src2Addr = (__ubuf__ T*)src2.GetPhyAddr();
38+ __ubuf__ T* dstAddr = (__ubuf__ T*)dst.GetPhyAddr();
39+ 
40+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> dy;
41+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> x;
42+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> constantOne;
43+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> value1MulsX;
44+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> temp1Reg;
45+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> temp2Reg;
46+ AscendC::MicroAPI::RegTensor<T, AscendC::MicroAPI::RegTraitNumOne> divRes;
47+ static constexpr AscendC::MicroAPI::DivSpecificMode mode = {
48+ AscendC::MicroAPI::MaskMergeMode::ZEROING, true};
49+ AscendC::MicroAPI::MaskReg mask;
50+ AscendC::MicroAPI::Duplicate(constantOne, value3);
51+ for (uint16_t loopIdx = 0; loopIdx < loopNum; loopIdx++) {
52+ mask = AscendC::MicroAPI::UpdateMask<T, AscendC::MicroAPI::RegTraitNumOne>(count);
53+ // OpCopyIn0
54+ AscendC::MicroAPI::DataCopy(x, (__ubuf__ T*)(src2Addr + loopIdx * vlSize));
55+ // temp1Reg = e^(-1.702x) + 1
56+ AscendC::MicroAPI::Muls(value1MulsX, x, value2, mask);
57+ AscendC::MicroAPI::Exp(temp1Reg, value1MulsX, mask);
58+ AscendC::MicroAPI::Adds(temp1Reg, temp1Reg, value3, mask);
59+ // temp2Reg = (1/(e^(-1.702x) + 1)) - 1
60+ AscendC::MicroAPI::Div<T, &mode>(divRes, constantOne, temp1Reg, mask);
61+ AscendC::MicroAPI::Adds(temp2Reg, divRes, value4, mask);
62+ // divRes = (temp2Reg * -1.702x + 1) * 1/(e^(-1.702x) * dy
63+ AscendC::MicroAPI::Mul(temp2Reg, temp2Reg, value1MulsX, mask);
64+ AscendC::MicroAPI::Adds(temp2Reg, temp2Reg, value3, mask);
65+ AscendC::MicroAPI::Mul(divRes, temp2Reg, divRes, mask);
66+ AscendC::MicroAPI::DataCopy(dy, (__ubuf__ T*)(src1Addr + loopIdx * vlSize));
67+ AscendC::MicroAPI::Mul(divRes, dy, divRes, mask);
68+ 
69+ // OpCopyOut
70+ AscendC::MicroAPI::DataCopy((__ubuf__ T*)(dstAddr + loopIdx * vlSize), divRes, mask);
71+ }
72+ }
73+#endif
74+ }
75+};
76+ 
77+template <typename T>
78+struct FastGeluGradNeedCast {
79+ using OpCopyIn0 = Bind<Vec::CopyIn<T>, Placeholder::In0<T>>; // dy
80+ using OpCopyIn1 = Bind<Vec::CopyIn<T>, Placeholder::In1<T>>; // x
81+ using CastIn0 = Bind<Vec::Cast<float, T, 0>, OpCopyIn0>;
82+ using CastIn1 = Bind<Vec::Cast<float, T, 0>, OpCopyIn1>;
83+ using OpFastGeluGrad = Bind<FastGeluGradCustom<float>, CastIn0, CastIn1>;
84+ using CastOut1 = Bind<Vec::Cast<T, float, 1>, OpFastGeluGrad>;
85+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, CastOut1>; // dx
86+ // 指定输出节点
87+ using Outputs = Elems<OpCopyOut>; // 设置输出
88+ using OpDag = DAGSch<Outputs>;
89+};
90+ 
91+template <typename T>
92+struct FastGeluGradNoCast {
93+ using OpCopyIn0 = Bind<Vec::CopyIn<T>, Placeholder::In0<T>>; // dy
94+ using OpCopyIn1 = Bind<Vec::CopyIn<T>, Placeholder::In1<T>>; // x
95+ using OpFastGeluGrad = Bind<FastGeluGradCustom<float>, OpCopyIn0, OpCopyIn1>; // compute
96+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpFastGeluGrad>; // dx
97+ // 指定输出节点
98+ using Outputs = Elems<OpCopyOut>; // 设置输出
99+ using OpDag = DAGSch<Outputs>;
100+};
101+ 
102+}; // namespace FastGeluGradDag
103+#endif // CANN_CUSTOM_OPS_FAST_GELU_GRAD_DAG_H
Aactivation/fast_gelu_grad/op_kernel/arch35/fast_gelu_grad_struct.h+39-0
@@ -0,0 +1,39 @@
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 "ascendc/host_api/tiling/template_argument.h"
12+ 
13+#ifndef CANN_CUSTOM_OPS_FAST_GELU_GRAD_STRUCT_H_
14+#define CANN_CUSTOM_OPS_FAST_GELU_GRAD_STRUCT_H_
15+ 
16+namespace FastGeluGradOp {
17+#define TPL_FP16 1
18+#define TPL_BF16 2
19+#define TPL_FP32 3
20+#define TPL_SCH_MODE_0 0
21+#define TPL_SCH_MODE_1 1
22+ 
23+ASCENDC_TPL_ARGS_DECL(
24+ FastGeluGrad, ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, TPL_SCH_MODE_0, TPL_SCH_MODE_1),
25+ ASCENDC_TPL_DTYPE_DECL(dType, TPL_FP16, TPL_BF16, TPL_FP32));
26+ 
27+ASCENDC_TPL_SEL(
28+ ASCENDC_TPL_ARGS_SEL(
29+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, TPL_SCH_MODE_0, TPL_SCH_MODE_1),
30+ ASCENDC_TPL_DTYPE_SEL(dType, TPL_FP16)),
31+ ASCENDC_TPL_ARGS_SEL(
32+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, TPL_SCH_MODE_0, TPL_SCH_MODE_1),
33+ ASCENDC_TPL_DTYPE_SEL(dType, TPL_BF16)),
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_FP32)));
37+} // namespace FastGeluGradOp
38+ 
39+#endif // CANN_CUSTOM_OPS_FAST_GELU_GRAD_STRUCT_H_
Aactivation/fast_gelu_grad/op_kernel/fast_gelu_grad_apt.cpp+49-0
@@ -0,0 +1,49 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "kernel_operator.h"
12+#include "kernel_tiling/kernel_tiling.h"
13+#include "atvoss/elewise/elewise_sch.h"
14+#include "arch35/fast_gelu_grad_dag.h"
15+#include "arch35/fast_gelu_grad_struct.h"
16+ 
17+using namespace AscendC;
18+using namespace FastGeluGradOp;
19+using namespace Ops::Base;
20+ 
21+template <uint64_t schMode, uint64_t dType>
22+__global__ __aicore__ void fast_gelu_grad(GM_ADDR dy, GM_ADDR x, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
23+{
24+ if (workspace == nullptr) {
25+ return;
26+ }
27+ SetSysWorkspace(workspace);
28+ GM_ADDR userWS = GetUserWorkspace(workspace);
29+ if (userWS == nullptr) {
30+ return;
31+ }
32+ REGISTER_TILING_DEFAULT(EleBaseTilingDataV2);
33+ GET_TILING_DATA_WITH_STRUCT(EleBaseTilingDataV2, tilingData, tiling);
34+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
35+ TPipe pipe;
36+ if constexpr (dType == static_cast<uint64_t>(TPL_FP16)) {
37+ ElementwiseSch<schMode, FastGeluGradDag::FastGeluGradNeedCast<half>::OpDag> sch(&tilingData, &pipe);
38+ sch.Init(dy, x, z);
39+ sch.Process();
40+ } else if constexpr (dType == static_cast<uint64_t>(TPL_BF16)) {
41+ ElementwiseSch<schMode, FastGeluGradDag::FastGeluGradNeedCast<bfloat16_t>::OpDag> sch(&tilingData, &pipe);
42+ sch.Init(dy, x, z);
43+ sch.Process();
44+ } else if constexpr (dType == static_cast<uint64_t>(TPL_FP32)) {
45+ ElementwiseSch<schMode, FastGeluGradDag::FastGeluGradNoCast<float>::OpDag> sch(&tilingData, &pipe);
46+ sch.Init(dy, x, z);
47+ sch.Process();
48+ }
49+}
Mactivation/fast_gelu_grad/tests/CMakeLists.txt+8-8
@@ -1,17 +1,17 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-message(STATUS "=== Debug: start ops.activation.fast_gelu_grad.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()
17endforeach()17endforeach()
Mactivation/fast_gelu_grad/tests/ut/CMakeLists.txt+4-5
@@ -1,14 +1,13 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
11-message(STATUS "=== Debug: start ops.activation.fast_gelu_grad.tests.ut.CMakeLists.txt ")
12file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)11file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
13message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")12message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
14foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})13foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
Ractivation/fast_gelu_grad/op_host/CMakeLists.txt→activation/fast_gelu_grad/tests/ut/op_api/CMakeLists.txt+8-4
@@ -1,10 +1,14 @@
1+# ----------------------------------------------------------------------------
1# Copyright (c) 2025 Huawei Technologies Co., Ltd.2# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
4# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
5-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
8-#/9+# ----------------------------------------------------------------------------
9 10 
10-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE fast_gelu_grad ACLNNTYPE aclnn_exclude)11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_API_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
Rloss/ctc_loss_v2_grad/op_host/op_api/aclnn_ctc_loss_backward.cpp→loss/ctc_loss_v2_grad/op_api/aclnn_ctc_loss_backward.cpp+0-0
Rloss/ctc_loss_v2_grad/op_host/op_api/aclnn_ctc_loss_backward.h→loss/ctc_loss_v2_grad/op_api/aclnn_ctc_loss_backward.h+0-0
Rloss/ctc_loss_v2_grad/op_host/op_api/ctc_loss_v2_grad.cpp→loss/ctc_loss_v2_grad/op_api/ctc_loss_v2_grad.cpp+1-2
Rloss/ctc_loss_v2_grad/op_host/op_api/ctc_loss_v2_grad.h→loss/ctc_loss_v2_grad/op_api/ctc_loss_v2_grad.h+0-0
Rloss/ctc_loss_v2_grad/op_host/CMakeLists.txt→loss/ctc_loss_v2_grad/tests/ut/op_api/CMakeLists.txt+7-4