【CANNbot】增加Fast_gelu_v2算子-ops-nn-AtomGit

Aexperimental/activation/fast_gelu_v2/CMakeLists.txt+13-0

@@ -0,0 +1,13 @@
	1	+# Copyright (c) 2026 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	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	11	+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
	12	+set(SUPPORT_TILING_DIR "arch35")
	13	+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE fast_gelu_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)

Aexperimental/activation/fast_gelu_v2/README.md+3-0

@@ -0,0 +1,3 @@
	1	+# FastGeluV2
	2	+
	3	+本目录仅包含FastGeluV2算子对应的aclnn接口；如您想要贡献该算子的AscendC实现，请参考[贡献流程](../../../CONTRIBUTING.md)。

Aexperimental/activation/fast_gelu_v2/docs/aclnnFastGeluV2.md+175-0

@@ -0,0 +1,175 @@
	1	+# aclnnFastGeluV2
	2	+
	3	+## 支持的产品型号
	4	+
	5	+\| 产品型号 \| 芯片号 \| 支持的数据类型 \|
	6	+\|---------\|--------\|---------------\|
	7	+\| Atlas A5 训练/推理系列 \| Ascend950PR \| float32, float16, bfloat16 \|
	8	+
	9	+## 功能描述
	10	+
	11	+计算 FastGeluV2 激活函数。FastGeluV2 是 GELU 激活函数的一种高效近似实现变体，通过分段多项式近似替代标准 GELU 中的 erf/tanh 计算。
	12	+
	13	+计算公式如下：
	14	+
	15	+```
	16	+FastGeluV2(x) = x * (sgn(x) * [-0.1444 * (clip(\|0.7071 * x\|, max=1.769) - 1.769)^2 + 0.5] + 0.5)
	17	+```
	18	+
	19	+其中：
	20	+
	21	+```
	22	+sgn(x) = (x + 1e-12) / \|x + 1e-12\|
	23	+```
	24	+
	25	+## aclnnFastGeluV2GetWorkspaceSize
	26	+
	27	+### 函数原型
	28	+
	29	+```cpp
	30	+aclnnStatus aclnnFastGeluV2GetWorkspaceSize(
	31	+ const aclTensor *x,
	32	+ const aclTensor *out,
	33	+ uint64_t *workspaceSize,
	34	+ aclOpExecutor **executor);
	35	+```
	36	+
	37	+### 参数说明
	38	+
	39	+\| 参数名 \| 输入/输出 \| 说明 \|
	40	+\|--------\|---------\|------\|
	41	+\| x \| 输入 \| 公式中的输入 `x`，Device 侧的 aclTensor，数据类型支持 float32、float16、bfloat16。支持非连续 Tensor，数据格式支持 ND，最大支持 8 维。 \|
	42	+\| out \| 输入 \| 公式中的输出 `y`，Device 侧的 aclTensor，数据类型需要与 `x` 一致，shape 需要与 `x` 一致。支持非连续 Tensor，数据格式支持 ND。 \|
	43	+\| workspaceSize \| 输出 \| 返回用户需要在 Device 侧申请的 workspace 大小。 \|
	44	+\| executor \| 输出 \| 返回 op 执行器，包含了算子计算流程。 \|
	45	+
	46	+### 返回值
	47	+
	48	+\| 返回值 \| 说明 \|
	49	+\|--------\|------\|
	50	+\| ACLNN_SUCCESS (0) \| 成功 \|
	51	+\| ACLNN_ERR_PARAM_NULLPTR (161001) \| 必选参数为空指针 \|
	52	+\| ACLNN_ERR_PARAM_INVALID (161002) \| 参数校验失败（dtype 不支持、shape 不匹配等） \|
	53	+\| 其他值 \| 失败 \|
	54	+
	55	+## aclnnFastGeluV2
	56	+
	57	+### 函数原型
	58	+
	59	+```cpp
	60	+aclnnStatus aclnnFastGeluV2(
	61	+ void *workspace,
	62	+ uint64_t workspaceSize,
	63	+ aclOpExecutor *executor,
	64	+ aclrtStream stream);
	65	+```
	66	+
	67	+### 参数说明
	68	+
	69	+\| 参数名 \| 输入/输出 \| 说明 \|
	70	+\|--------\|---------\|------\|
	71	+\| workspace \| 输入 \| 在 Device 侧申请的 workspace 内存地址。 \|
	72	+\| workspaceSize \| 输入 \| 在 Device 侧申请的 workspace 大小，由第一段接口 aclnnFastGeluV2GetWorkspaceSize 获取。 \|
	73	+\| executor \| 输入 \| op 执行器，包含了算子计算流程。 \|
	74	+\| stream \| 输入 \| 指定执行任务的 AscendCL Stream 流。 \|
	75	+
	76	+### 返回值
	77	+
	78	+\| 返回值 \| 说明 \|
	79	+\|--------\|------\|
	80	+\| ACLNN_SUCCESS (0) \| 成功 \|
	81	+\| 其他值 \| 失败 \|
	82	+
	83	+## 约束说明
	84	+
	85	+1. 输入 `x` 和输出 `out` 的数据类型必须一致。
	86	+2. 输出 `out` 的 shape 必须与输入 `x` 的 shape 一致。
	87	+3. 输入张量维度不超过 8 维。
	88	+4. 不支持私有数据格式（Private Format）。
	89	+5. 输入 `x` 为空 Tensor 时，`workspaceSize` 返回 0，直接返回成功。
	90	+
	91	+## 调用示例
	92	+
	93	+以下代码摘自 [examples/test_aclnn_fast_gelu_v2.cpp](../examples/test_aclnn_fast_gelu_v2.cpp)，演示两段式 ACLNN 接口的完整调用流程。
	94	+
	95	+### 1. 初始化 ACL 运行时
	96	+
	97	+```cpp
	98	+#include "acl/acl.h"
	99	+#include "aclnn_fast_gelu_v2.h"
	100	+
	101	+// 初始化 ACL 并设置设备
	102	+aclInit(nullptr);
	103	+int32_t deviceId = 0;
	104	+aclrtSetDevice(deviceId);
	105	+aclrtStream stream = nullptr;
	106	+aclrtCreateStream(&stream);
	107	+```
	108	+
	109	+### 2. 准备输入输出张量
	110	+
	111	+```cpp
	112	+// 定义 shape 和数据
	113	+std::vector<int64_t> shape = {2, 8};
	114	+int64_t totalElements = 16;
	115	+size_t dataBytes = totalElements * sizeof(float);
	116	+
	117	+// 分配 Device 内存并拷贝输入数据
	118	+void* xDevAddr = nullptr;
	119	+aclrtMalloc(&xDevAddr, dataBytes, ACL_MEM_MALLOC_HUGE_FIRST);
	120	+aclrtMemcpy(xDevAddr, dataBytes, hostInput.data(), dataBytes, ACL_MEMCPY_HOST_TO_DEVICE);
	121	+
	122	+// 计算 strides 并创建 aclTensor
	123	+auto strides = ComputeStrides(shape); // row-major contiguous
	124	+aclTensor* xTensor = aclCreateTensor(
	125	+ shape.data(), shape.size(), ACL_FLOAT,
	126	+ strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	127	+ shape.data(), shape.size(), xDevAddr);
	128	+
	129	+// 同样创建输出 tensor（outDevAddr, outTensor）
	130	+```
	131	+
	132	+### 3. Phase 1: GetWorkspaceSize
	133	+
	134	+```cpp
	135	+uint64_t workspaceSize = 0;
	136	+aclOpExecutor* executor = nullptr;
	137	+aclnnFastGeluV2GetWorkspaceSize(xTensor, outTensor, &workspaceSize, &executor);
	138	+
	139	+// 按需分配 workspace
	140	+void* workspace = nullptr;
	141	+if (workspaceSize > 0) {
	142	+ aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	143	+}
	144	+```
	145	+
	146	+### 4. Phase 2: 执行算子
	147	+
	148	+```cpp
	149	+aclnnFastGeluV2(workspace, workspaceSize, executor, stream);
	150	+aclrtSynchronizeStream(stream);
	151	+```
	152	+
	153	+### 5. 获取结果并释放资源
	154	+
	155	+```cpp
	156	+// D2H 拷贝结果
	157	+aclrtMemcpy(hostOutput.data(), dataBytes, outDevAddr, dataBytes, ACL_MEMCPY_DEVICE_TO_HOST);
	158	+
	159	+// 释放资源
	160	+if (workspace) aclrtFree(workspace);
	161	+aclDestroyTensor(xTensor);
	162	+aclDestroyTensor(outTensor);
	163	+aclrtFree(xDevAddr);
	164	+aclrtFree(outDevAddr);
	165	+aclrtDestroyStream(stream);
	166	+aclrtResetDevice(deviceId);
	167	+aclFinalize();
	168	+```
	169	+
	170	+### 编译运行
	171	+
	172	+```bash
	173	+# 编译并运行示例
	174	+cd examples && bash run.sh --eager
	175	+```

Aexperimental/activation/fast_gelu_v2/examples/arch35/test_aclnn_fast_gelu_v2.cpp+289-0

@@ -0,0 +1,289 @@
	1	+/**
	2	+ * @file test_aclnn_fast_gelu_v2.cpp
	3	+ * @brief ACLNN two-phase invocation example for FastGeluV2 custom operator.
	4	+ *
	5	+ * This example demonstrates the standard ACLNN two-phase calling pattern:
	6	+ * Phase 1: aclnnFastGeluV2GetWorkspaceSize -- compute required workspace
	7	+ * Phase 2: aclnnFastGeluV2 -- execute the operator
	8	+ *
	9	+ * Prerequisites:
	10	+ * - CANN Toolkit installed (source set_env.sh)
	11	+ * - FastGeluV2 custom operator package installed (build.sh --soc=ascend950, then run .run file)
	12	+ * - NPU device available
	13	+ *
	14	+ * Build & run:
	15	+ * cd examples && bash run.sh --eager
	16	+ */
	17	+
	18	+#include <iostream>
	19	+#include <vector>
	20	+#include <cmath>
	21	+#include <cstring>
	22	+#include <algorithm>
	23	+
	24	+#include "acl/acl.h"
	25	+#include "aclnn_fast_gelu_v2.h"
	26	+
	27	+// ============================================================================
	28	+// Macros
	29	+// ============================================================================
	30	+#define CHECK_ACL(expr) \
	31	+ do { \
	32	+ auto _ret = (expr); \
	33	+ if (_ret != ACL_SUCCESS) { \
	34	+ std::cerr << "[ERROR] " << #expr << " failed: " << _ret \
	35	+ << " at " << __FILE__ << ":" << __LINE__ \
	36	+ << std::endl; \
	37	+ return 1; \
	38	+ } \
	39	+ } while (0)
	40	+
	41	+// ============================================================================
	42	+// Helper: compute strides from shape (row-major contiguous)
	43	+// ============================================================================
	44	+static std::vector<int64_t> ComputeStrides(const std::vector<int64_t>& shape) {
	45	+ std::vector<int64_t> strides(shape.size(), 1);
	46	+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) {
	47	+ strides[i] = shape[i + 1] * strides[i + 1];
	48	+ }
	49	+ return strides;
	50	+}
	51	+
	52	+// ============================================================================
	53	+// Helper: create aclTensor from host data
	54	+// ============================================================================
	55	+static int CreateAclTensor(const void* hostData, size_t dataBytes,
	56	+ const std::vector<int64_t>& shape,
	57	+ void** deviceAddr,
	58	+ aclDataType dataType,
	59	+ aclTensor** tensor) {
	60	+ auto ret = aclrtMalloc(deviceAddr, dataBytes, ACL_MEM_MALLOC_HUGE_FIRST);
	61	+ if (ret != ACL_SUCCESS) {
	62	+ std::cerr << "[ERROR] aclrtMalloc failed: " << ret << std::endl;
	63	+ return ret;
	64	+ }
	65	+
	66	+ ret = aclrtMemcpy(*deviceAddr, dataBytes, hostData, dataBytes,
	67	+ ACL_MEMCPY_HOST_TO_DEVICE);
	68	+ if (ret != ACL_SUCCESS) {
	69	+ std::cerr << "[ERROR] aclrtMemcpy H2D failed: " << ret << std::endl;
	70	+ aclrtFree(*deviceAddr);
	71	+ return ret;
	72	+ }
	73	+
	74	+ auto strides = ComputeStrides(shape);
	75	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType,
	76	+ strides.data(), 0, aclFormat::ACL_FORMAT_ND,
	77	+ shape.data(), shape.size(), *deviceAddr);
	78	+ if (*tensor == nullptr) {
	79	+ std::cerr << "[ERROR] aclCreateTensor returned nullptr" << std::endl;
	80	+ aclrtFree(*deviceAddr);
	81	+ return -1;
	82	+ }
	83	+ return ACL_SUCCESS;
	84	+}
	85	+
	86	+// ============================================================================
	87	+// CPU golden: FastGeluV2(x) using float32 precision
	88	+// ============================================================================
	89	+static void ComputeGolden(const float* x, float* output, size_t size) {
	90	+ for (size_t i = 0; i < size; ++i) {
	91	+ float xi = x[i];
	92	+ float eps = 1e-12f;
	93	+ float x_eps = xi + eps;
	94	+ float abs_x_eps = std::fabs(x_eps);
	95	+ float sgn_x = x_eps / abs_x_eps;
	96	+ float abs_x = std::fabs(xi);
	97	+ float scaled_x = 0.7071f * abs_x;
	98	+ float clipped_x = std::min(scaled_x, 1.769f);
	99	+ float diff = clipped_x - 1.769f;
	100	+ float sq = diff * diff;
	101	+ float inner = -0.1444f * sq + 0.5f;
	102	+ float bracket = sgn_x * inner + 0.5f;
	103	+ output[i] = xi * bracket;
	104	+ }
	105	+}
	106	+
	107	+// ============================================================================
	108	+// Main
	109	+// ============================================================================
	110	+int main() {
	111	+ std::cout << "========================================" << std::endl;
	112	+ std::cout << "FastGeluV2 ACLNN Example" << std::endl;
	113	+ std::cout << "========================================" << std::endl;
	114	+
	115	+ // -----------------------------------------------------------------------
	116	+ // Step 1: Initialize ACL runtime
	117	+ // -----------------------------------------------------------------------
	118	+ std::cout << "\n[Step 1] Initializing ACL..." << std::endl;
	119	+ CHECK_ACL(aclInit(nullptr));
	120	+
	121	+ int32_t deviceId = 0;
	122	+ CHECK_ACL(aclrtSetDevice(deviceId));
	123	+
	124	+ aclrtStream stream = nullptr;
	125	+ CHECK_ACL(aclrtCreateStream(&stream));
	126	+ std::cout << " ACL initialized, device=" << deviceId << std::endl;
	127	+
	128	+ // -----------------------------------------------------------------------
	129	+ // Step 2: Prepare input data on host
	130	+ // -----------------------------------------------------------------------
	131	+ std::cout << "\n[Step 2] Preparing input data..." << std::endl;
	132	+
	133	+ std::vector<int64_t> shape = {2, 8}; // 2D tensor, 16 elements total
	134	+ int64_t totalElements = 1;
	135	+ for (auto d : shape) totalElements *= d;
	136	+
	137	+ // Generate test data: linearly spaced from -5.0 to 5.0
	138	+ std::vector<float> hostInput(totalElements);
	139	+ for (int64_t i = 0; i < totalElements; ++i) {
	140	+ hostInput[i] = -5.0f + 10.0f * static_cast<float>(i) /
	141	+ static_cast<float>(totalElements - 1);
	142	+ }
	143	+
	144	+ size_t dataBytes = totalElements * sizeof(float);
	145	+
	146	+ std::cout << " Shape: [" << shape[0] << ", " << shape[1] << "]" << std::endl;
	147	+ std::cout << " Elements: " << totalElements << std::endl;
	148	+ std::cout << " Dtype: float32" << std::endl;
	149	+ std::cout << " Input range: [" << hostInput.front() << ", "
	150	+ << hostInput.back() << "]" << std::endl;
	151	+
	152	+ // -----------------------------------------------------------------------
	153	+ // Step 3: Create aclTensors (allocate device memory, copy H2D)
	154	+ // -----------------------------------------------------------------------
	155	+ std::cout << "\n[Step 3] Creating aclTensors..." << std::endl;
	156	+
	157	+ void* xDevAddr = nullptr;
	158	+ aclTensor* xTensor = nullptr;
	159	+ if (CreateAclTensor(hostInput.data(), dataBytes, shape,
	160	+ &xDevAddr, ACL_FLOAT, &xTensor) != ACL_SUCCESS) {
	161	+ std::cerr << " Failed to create input tensor" << std::endl;
	162	+ return 1;
	163	+ }
	164	+
	165	+ // Allocate output tensor (initialized to zero)
	166	+ std::vector<float> hostOutput(totalElements, 0.0f);
	167	+ void* outDevAddr = nullptr;
	168	+ aclTensor* outTensor = nullptr;
	169	+ if (CreateAclTensor(hostOutput.data(), dataBytes, shape,
	170	+ &outDevAddr, ACL_FLOAT, &outTensor) != ACL_SUCCESS) {
	171	+ std::cerr << " Failed to create output tensor" << std::endl;
	172	+ aclDestroyTensor(xTensor);
	173	+ aclrtFree(xDevAddr);
	174	+ return 1;
	175	+ }
	176	+ std::cout << " Input and output tensors created on device" << std::endl;
	177	+
	178	+ // -----------------------------------------------------------------------
	179	+ // Step 4: Phase 1 -- GetWorkspaceSize
	180	+ // -----------------------------------------------------------------------
	181	+ std::cout << "\n[Step 4] Calling aclnnFastGeluV2GetWorkspaceSize..." << std::endl;
	182	+
	183	+ uint64_t workspaceSize = 0;
	184	+ aclOpExecutor* executor = nullptr;
	185	+
	186	+ CHECK_ACL(aclnnFastGeluV2GetWorkspaceSize(xTensor, outTensor,
	187	+ &workspaceSize, &executor));
	188	+ std::cout << " workspaceSize = " << workspaceSize << " bytes" << std::endl;
	189	+
	190	+ // -----------------------------------------------------------------------
	191	+ // Step 5: Allocate workspace on device
	192	+ // -----------------------------------------------------------------------
	193	+ void* workspace = nullptr;
	194	+ if (workspaceSize > 0) {
	195	+ std::cout << "\n[Step 5] Allocating workspace..." << std::endl;
	196	+ CHECK_ACL(aclrtMalloc(&workspace, workspaceSize,
	197	+ ACL_MEM_MALLOC_HUGE_FIRST));
	198	+ std::cout << " Workspace allocated: " << workspaceSize
	199	+ << " bytes" << std::endl;
	200	+ } else {
	201	+ std::cout << "\n[Step 5] No workspace needed (size=0)" << std::endl;
	202	+ }
	203	+
	204	+ // -----------------------------------------------------------------------
	205	+ // Step 6: Phase 2 -- Execute the operator
	206	+ // -----------------------------------------------------------------------
	207	+ std::cout << "\n[Step 6] Calling aclnnFastGeluV2..." << std::endl;
	208	+ CHECK_ACL(aclnnFastGeluV2(workspace, workspaceSize, executor, stream));
	209	+
	210	+ // Synchronize to ensure execution completes
	211	+ CHECK_ACL(aclrtSynchronizeStream(stream));
	212	+ std::cout << " Operator execution completed" << std::endl;
	213	+
	214	+ // -----------------------------------------------------------------------
	215	+ // Step 7: Copy result from device to host
	216	+ // -----------------------------------------------------------------------
	217	+ std::cout << "\n[Step 7] Copying results D2H..." << std::endl;
	218	+ CHECK_ACL(aclrtMemcpy(hostOutput.data(), dataBytes, outDevAddr, dataBytes,
	219	+ ACL_MEMCPY_DEVICE_TO_HOST));
	220	+
	221	+ // -----------------------------------------------------------------------
	222	+ // Step 8: Compare with CPU golden
	223	+ // -----------------------------------------------------------------------
	224	+ std::cout << "\n[Step 8] Verifying results against CPU golden..." << std::endl;
	225	+
	226	+ std::vector<float> golden(totalElements);
	227	+ ComputeGolden(hostInput.data(), golden.data(), totalElements);
	228	+
	229	+ // Print first few results
	230	+ int printCount = std::min(static_cast<int64_t>(8), totalElements);
	231	+ std::cout << "\n Index \| Input \| Golden \| NPU Output \| Match" << std::endl;
	232	+ std::cout << " ------+-----------+-----------+------------+------" << std::endl;
	233	+
	234	+ double maxRelErr = 0.0;
	235	+ double sumRelErr = 0.0;
	236	+ const double eps = 1e-7;
	237	+
	238	+ for (int64_t i = 0; i < totalElements; ++i) {
	239	+ double g = static_cast<double>(golden[i]);
	240	+ double a = static_cast<double>(hostOutput[i]);
	241	+ double relErr = std::abs(a - g) / (std::abs(g) + eps);
	242	+ if (relErr > maxRelErr) maxRelErr = relErr;
	243	+ sumRelErr += relErr;
	244	+
	245	+ if (i < printCount) {
	246	+ const char* matchStr = (relErr < 1.22e-4) ? "OK" : "MISMATCH";
	247	+ printf(" %5ld \| %9.4f \| %9.6f \| %10.6f \| %s\n",
	248	+ (long)i, hostInput[i], golden[i], hostOutput[i], matchStr);
	249	+ }
	250	+ }
	251	+
	252	+ double meanRelErr = sumRelErr / static_cast<double>(totalElements);
	253	+ double threshold = 1.22e-4; // 2^-13 for float32
	254	+
	255	+ std::cout << "\n MERE (mean relative error) = " << meanRelErr << std::endl;
	256	+ std::cout << " MARE (max relative error) = " << maxRelErr << std::endl;
	257	+ std::cout << " Threshold = " << threshold << std::endl;
	258	+
	259	+ bool pass = (meanRelErr < threshold) && (maxRelErr < 10.0 * threshold);
	260	+
	261	+ // -----------------------------------------------------------------------
	262	+ // Step 9: Cleanup
	263	+ // -----------------------------------------------------------------------
	264	+ std::cout << "\n[Step 9] Cleaning up..." << std::endl;
	265	+
	266	+ if (workspace) aclrtFree(workspace);
	267	+ aclDestroyTensor(xTensor);
	268	+ aclDestroyTensor(outTensor);
	269	+ aclrtFree(xDevAddr);
	270	+ aclrtFree(outDevAddr);
	271	+ aclrtDestroyStream(stream);
	272	+ aclrtResetDevice(deviceId);
	273	+ aclFinalize();
	274	+
	275	+ std::cout << " Resources released" << std::endl;
	276	+
	277	+ // -----------------------------------------------------------------------
	278	+ // Final result
	279	+ // -----------------------------------------------------------------------
	280	+ std::cout << "\n========================================" << std::endl;
	281	+ if (pass) {
	282	+ std::cout << "Result: PASS" << std::endl;
	283	+ } else {
	284	+ std::cout << "Result: FAIL" << std::endl;
	285	+ }
	286	+ std::cout << "========================================" << std::endl;
	287	+
	288	+ return pass ? 0 : 1;
	289	+}

Aexperimental/activation/fast_gelu_v2/examples/arch35/test_geir_fast_gelu_v2.cpp+72-0

@@ -0,0 +1,72 @@
	1	+/**
	2	+ * @file test_geir_fast_gelu_v2.cpp
	3	+ * @brief GE IR (Graph Engine) mode example for FastGeluV2 custom operator.
	4	+ *
	5	+ * NOTE: FastGeluV2 currently does NOT support GE IR graph mode.
	6	+ * Supported calling method: ACLNN single-operator mode only.
	7	+ *
	8	+ * This file serves as a placeholder to explain why graph mode is not
	9	+ * available and provides a reference for future graph-mode integration
	10	+ * if needed.
	11	+ *
	12	+ * For the working ACLNN single-operator example, see:
	13	+ * test_aclnn_fast_gelu_v2.cpp
	14	+ *
	15	+ * Build & run:
	16	+ * cd examples && bash run.sh --geir
	17	+ */
	18	+
	19	+#include <iostream>
	20	+
	21	+// ============================================================================
	22	+// Graph mode is not supported for FastGeluV2
	23	+// ============================================================================
	24	+//
	25	+// Reason:
	26	+// FastGeluV2 is a user-defined custom activation function intended for
	27	+// single-operator (ACLNN) invocation only. The following graph-mode
	28	+// calling methods are NOT supported:
	29	+//
	30	+// - torch_npu single-operator dispatch
	31	+// - torch.compile graph capture
	32	+// - GE static-shape graph mode
	33	+// - GE dynamic-shape graph mode
	34	+//
	35	+// To use this operator, call the ACLNN two-phase interface directly:
	36	+//
	37	+// 1. aclnnFastGeluV2GetWorkspaceSize(x, out, &workspaceSize, &executor)
	38	+// 2. aclnnFastGeluV2(workspace, workspaceSize, executor, stream)
	39	+//
	40	+// See test_aclnn_fast_gelu_v2.cpp for a complete working example.
	41	+//
	42	+// If graph-mode support is required in the future, the following steps
	43	+// would be needed:
	44	+// 1. Define GE IR proto for FastGeluV2 (op type, inputs, outputs, attrs)
	45	+// 2. Register the operator with GE graph engine via op_proto
	46	+// 3. Implement GE graph-mode plugin (framework adapter)
	47	+// 4. Validate through graph compilation and execution pipeline
	48	+//
	49	+// ============================================================================
	50	+
	51	+int main() {
	52	+ std::cout << "========================================" << std::endl;
	53	+ std::cout << "FastGeluV2 GE IR Graph Mode Example" << std::endl;
	54	+ std::cout << "========================================" << std::endl;
	55	+ std::cout << std::endl;
	56	+ std::cout << "FastGeluV2 does NOT support GE IR graph mode." << std::endl;
	57	+ std::cout << std::endl;
	58	+ std::cout << "Supported calling method:" << std::endl;
	59	+ std::cout << " - ACLNN single-operator mode (two-phase interface)" << std::endl;
	60	+ std::cout << std::endl;
	61	+ std::cout << "Please refer to test_aclnn_fast_gelu_v2.cpp for a" << std::endl;
	62	+ std::cout << "complete working example of ACLNN invocation." << std::endl;
	63	+ std::cout << std::endl;
	64	+ std::cout << "Run the ACLNN example:" << std::endl;
	65	+ std::cout << " cd examples && bash run.sh --eager" << std::endl;
	66	+ std::cout << std::endl;
	67	+ std::cout << "========================================" << std::endl;
	68	+ std::cout << "Result: PASS (informational only)" << std::endl;
	69	+ std::cout << "========================================" << std::endl;
	70	+
	71	+ return 0;
	72	+}

Aexperimental/activation/fast_gelu_v2/op_api/aclnn_fast_gelu_v2.cpp+174-0

@@ -0,0 +1,174 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * @file aclnn_fast_gelu_v2.cpp
	13	+ * @brief ACLNN L2 API implementation for FastGeluV2
	14	+ */
	15	+
	16	+#include "aclnn_fast_gelu_v2.h"
	17	+#include "fast_gelu_v2.h"
	18	+#include "aclnn_kernels/contiguous.h"
	19	+#include "aclnn_kernels/common/op_error_check.h"
	20	+#include "opdev/op_log.h"
	21	+#include "opdev/op_dfx.h"
	22	+#include "opdev/common_types.h"
	23	+#include "opdev/data_type_utils.h"
	24	+#include "opdev/make_op_executor.h"
	25	+#include "opdev/platform.h"
	26	+
	27	+using namespace op;
	28	+
	29	+#define ACLNN_MAX_SHAPE_RANK 8
	30	+
	31	+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
	32	+ DataType::DT_FLOAT, DataType::DT_FLOAT16, DataType::DT_BF16
	33	+};
	34	+
	35	+static bool HasEmptyTensor(const aclTensor* x)
	36	+{
	37	+ return x->IsEmpty();
	38	+}
	39	+
	40	+static bool CheckNotNull(const aclTensor* x, const aclTensor* out)
	41	+{
	42	+ OP_CHECK_NULL(x, return false);
	43	+ OP_CHECK_NULL(out, return false);
	44	+ return true;
	45	+}
	46	+
	47	+static bool CheckDtypeValid(const aclTensor* x, const aclTensor* out)
	48	+{
	49	+ OP_CHECK_DTYPE_NOT_MATCH(out, x->GetDataType(), return false);
	50	+
	51	+ if (!CheckType(x->GetDataType(), AICORE_DTYPE_SUPPORT_LIST)) {
	52	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	53	+ "Dtype not supported: dtype=%d. Supported: FLOAT, FLOAT16, BF16.",
	54	+ static_cast<int>(x->GetDataType()));
	55	+ return false;
	56	+ }
	57	+ return true;
	58	+}
	59	+
	60	+static bool CheckFormat(const aclTensor* x, const aclTensor* out)
	61	+{
	62	+ auto formatX = x->GetStorageFormat();
	63	+ auto formatOut = out->GetStorageFormat();
	64	+
	65	+ if (IsPrivateFormat(formatX) \|\| IsPrivateFormat(formatOut)) {
	66	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	67	+ "Private format not supported: x=%d, out=%d",
	68	+ static_cast<int>(formatX), static_cast<int>(formatOut));
	69	+ return false;
	70	+ }
	71	+ return true;
	72	+}
	73	+
	74	+static bool CheckShape(const aclTensor* x, const aclTensor* out)
	75	+{
	76	+ OP_CHECK_MAX_DIM(x, ACLNN_MAX_SHAPE_RANK, return false);
	77	+ OP_CHECK_MAX_DIM(out, ACLNN_MAX_SHAPE_RANK, return false);
	78	+
	79	+ // Verify output shape matches input shape (elementwise: y.shape == x.shape)
	80	+ auto xShape = x->GetViewShape();
	81	+ auto outShape = out->GetViewShape();
	82	+ if (xShape.GetDimNum() != outShape.GetDimNum()) {
	83	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	84	+ "Shape rank mismatch: x has %zu dims, out has %zu dims",
	85	+ xShape.GetDimNum(), outShape.GetDimNum());
	86	+ return false;
	87	+ }
	88	+ for (size_t i = 0; i < xShape.GetDimNum(); ++i) {
	89	+ if (xShape.GetDim(i) != outShape.GetDim(i)) {
	90	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	91	+ "Shape mismatch at dim %zu: x=%ld, out=%ld",
	92	+ i, xShape.GetDim(i), outShape.GetDim(i));
	93	+ return false;
	94	+ }
	95	+ }
	96	+ return true;
	97	+}
	98	+
	99	+static aclnnStatus CheckParams(const aclTensor* x, const aclTensor* out)
	100	+{
	101	+ if (!CheckNotNull(x, out)) {
	102	+ OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "CheckNotNull failed");
	103	+ return ACLNN_ERR_PARAM_NULLPTR;
	104	+ }
	105	+ if (!CheckDtypeValid(x, out)) {
	106	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "CheckDtypeValid failed: x_dtype=%d, out_dtype=%d",
	107	+ static_cast<int>(x->GetDataType()), static_cast<int>(out->GetDataType()));
	108	+ return ACLNN_ERR_PARAM_INVALID;
	109	+ }
	110	+ if (!CheckFormat(x, out)) {
	111	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "CheckFormat failed");
	112	+ return ACLNN_ERR_PARAM_INVALID;
	113	+ }
	114	+ if (!CheckShape(x, out)) {
	115	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "CheckShape failed: x_dim=%zu, out_dim=%zu",
	116	+ x->GetViewShape().GetDimNum(), out->GetViewShape().GetDimNum());
	117	+ return ACLNN_ERR_PARAM_INVALID;
	118	+ }
	119	+ return ACLNN_SUCCESS;
	120	+}
	121	+
	122	+extern "C" aclnnStatus aclnnFastGeluV2GetWorkspaceSize(
	123	+ const aclTensor* x,
	124	+ const aclTensor* out,
	125	+ uint64_t* workspaceSize,
	126	+ aclOpExecutor** executor)
	127	+{
	128	+ L2_DFX_PHASE_1(aclnnFastGeluV2, DFX_IN(x), DFX_OUT(out));
	129	+
	130	+ // Validate output parameters before any other work
	131	+ if (workspaceSize == nullptr) {
	132	+ OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "workspaceSize is nullptr");
	133	+ return ACLNN_ERR_PARAM_NULLPTR;
	134	+ }
	135	+ if (executor == nullptr) {
	136	+ OP_LOGE(ACLNN_ERR_PARAM_NULLPTR, "executor is nullptr");
	137	+ return ACLNN_ERR_PARAM_NULLPTR;
	138	+ }
	139	+
	140	+ auto uniqueExecutor = CREATE_EXECUTOR();
	141	+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
	142	+
	143	+ auto ret = CheckParams(x, out);
	144	+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
	145	+
	146	+ if (HasEmptyTensor(x)) {
	147	+ *workspaceSize = 0;
	148	+ uniqueExecutor.ReleaseTo(executor);
	149	+ return ACLNN_SUCCESS;
	150	+ }
	151	+
	152	+ auto xContiguous = l0op::Contiguous(x, uniqueExecutor.get());
	153	+ CHECK_RET(xContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
	154	+
	155	+ const aclTensor* opResult = l0op::FastGeluV2(xContiguous, uniqueExecutor.get());
	156	+ CHECK_RET(opResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
	157	+
	158	+ auto viewCopyResult = l0op::ViewCopy(opResult, out, uniqueExecutor.get());
	159	+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
	160	+
	161	+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
	162	+ uniqueExecutor.ReleaseTo(executor);
	163	+ return ACLNN_SUCCESS;
	164	+}
	165	+
	166	+extern "C" aclnnStatus aclnnFastGeluV2(
	167	+ void* workspace,
	168	+ uint64_t workspaceSize,
	169	+ aclOpExecutor* executor,
	170	+ aclrtStream stream)
	171	+{
	172	+ L2_DFX_PHASE_2(aclnnFastGeluV2);
	173	+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
	174	+}

Aexperimental/activation/fast_gelu_v2/op_api/aclnn_fast_gelu_v2.h+61-0

@@ -0,0 +1,61 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * @file aclnn_fast_gelu_v2.h
	13	+ * @brief ACLNN L2 API declaration for FastGeluV2
	14	+ */
	15	+
	16	+#ifndef ACLNN_FAST_GELU_V2_H_
	17	+#define ACLNN_FAST_GELU_V2_H_
	18	+
	19	+#include "aclnn/aclnn_base.h"
	20	+
	21	+#ifndef ACLNN_API
	22	+#define ACLNN_API __attribute__((visibility("default")))
	23	+#endif
	24	+
	25	+#ifdef __cplusplus
	26	+extern "C" {
	27	+#endif
	28	+
	29	+/**
	30	+ * @brief Compute workspace size for aclnnFastGeluV2
	31	+ * @param x [in] Input tensor
	32	+ * @param out [in] Output tensor
	33	+ * @param workspaceSize [out] Required workspace size in bytes
	34	+ * @param executor [out] Executor handle
	35	+ * @return aclnnStatus
	36	+ */
	37	+ACLNN_API aclnnStatus aclnnFastGeluV2GetWorkspaceSize(
	38	+ const aclTensor *x,
	39	+ const aclTensor *out,
	40	+ uint64_t *workspaceSize,
	41	+ aclOpExecutor **executor);
	42	+
	43	+/**
	44	+ * @brief Execute FastGeluV2 computation
	45	+ * @param workspace [in] Workspace memory address
	46	+ * @param workspaceSize [in] Workspace size in bytes
	47	+ * @param executor [in] Executor handle
	48	+ * @param stream [in] ACL stream
	49	+ * @return aclnnStatus
	50	+ */
	51	+ACLNN_API aclnnStatus aclnnFastGeluV2(
	52	+ void *workspace,
	53	+ uint64_t workspaceSize,
	54	+ aclOpExecutor *executor,
	55	+ aclrtStream stream);
	56	+
	57	+#ifdef __cplusplus
	58	+}
	59	+#endif
	60	+
	61	+#endif // ACLNN_FAST_GELU_V2_H_

Aexperimental/activation/fast_gelu_v2/op_api/fast_gelu_v2.cpp+79-0

@@ -0,0 +1,79 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * @file fast_gelu_v2.cpp
	13	+ * @brief ACLNN L0 API implementation for FastGeluV2
	14	+ */
	15	+
	16	+#include "fast_gelu_v2.h"
	17	+#include "opdev/op_log.h"
	18	+#include "opdev/op_dfx.h"
	19	+#include "opdev/shape_utils.h"
	20	+#include "opdev/make_op_executor.h"
	21	+
	22	+using namespace op;
	23	+
	24	+namespace l0op {
	25	+
	26	+OP_TYPE_REGISTER(FastGeluV2);
	27	+
	28	+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
	29	+ DataType::DT_FLOAT, DataType::DT_FLOAT16, DataType::DT_BF16
	30	+};
	31	+
	32	+static bool IsAiCoreSupport(const aclTensor* x)
	33	+{
	34	+ return CheckType(x->GetDataType(), AICORE_DTYPE_SUPPORT_LIST);
	35	+}
	36	+
	37	+static bool FastGeluV2InferShape(const op::Shape& xShape, op::Shape& outShape)
	38	+{
	39	+ outShape = xShape;
	40	+ return true;
	41	+}
	42	+
	43	+static const aclTensor* FastGeluV2AiCore(const aclTensor* x, const aclTensor* out,
	44	+ aclOpExecutor* executor)
	45	+{
	46	+ L0_DFX(FastGeluV2AiCore, x, out);
	47	+
	48	+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(FastGeluV2,
	49	+ OP_INPUT(x), OP_OUTPUT(out));
	50	+ OP_CHECK(
	51	+ ret == ACLNN_SUCCESS,
	52	+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "FastGeluV2AiCore failed."),
	53	+ return nullptr);
	54	+ return out;
	55	+}
	56	+
	57	+const aclTensor* FastGeluV2(const aclTensor* x, aclOpExecutor* executor)
	58	+{
	59	+ Shape outShape;
	60	+ const aclTensor* out = nullptr;
	61	+
	62	+ if (!FastGeluV2InferShape(x->GetViewShape(), outShape)) {
	63	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Infer shape failed.");
	64	+ return nullptr;
	65	+ }
	66	+
	67	+ if (!IsAiCoreSupport(x)) {
	68	+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
	69	+ "FastGeluV2 not supported: dtype x=%d. Supported: FLOAT, FLOAT16, BF16.",
	70	+ static_cast<int>(x->GetDataType()));
	71	+ return nullptr;
	72	+ }
	73	+
	74	+ out = executor->AllocTensor(outShape, x->GetDataType());
	75	+
	76	+ return FastGeluV2AiCore(x, out, executor);
	77	+}
	78	+
	79	+} // namespace l0op

Aexperimental/activation/fast_gelu_v2/op_api/fast_gelu_v2.h+27-0

@@ -0,0 +1,27 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/**
	12	+ * @file fast_gelu_v2.h
	13	+ * @brief ACLNN L0 API declaration for FastGeluV2
	14	+ */
	15	+
	16	+#ifndef OP_API_INC_LEVEL0_FAST_GELU_V2_H_
	17	+#define OP_API_INC_LEVEL0_FAST_GELU_V2_H_
	18	+
	19	+#include "opdev/op_executor.h"
	20	+
	21	+namespace l0op {
	22	+
	23	+const aclTensor* FastGeluV2(const aclTensor* x, aclOpExecutor* executor);
	24	+
	25	+} // namespace l0op
	26	+
	27	+#endif // OP_API_INC_LEVEL0_FAST_GELU_V2_H_

Aexperimental/activation/fast_gelu_v2/op_graph/fast_gelu_v2_proto.h+41-0

@@ -0,0 +1,41 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file fast_gelu_v2_proto.h
	13	+ * \brief
	14	+ */
	15	+#ifndef OPS_BUILT_IN_OP_PROTO_INC_FAST_GELU_V2_OPS_H_
	16	+#define OPS_BUILT_IN_OP_PROTO_INC_FAST_GELU_V2_OPS_H_
	17	+
	18	+#include "graph/operator_reg.h"
	19	+
	20	+namespace ge {
	21	+/**
	22	+* @brief The FastGeluV2 activation function is
	23	+* FastGeluV2(x) = x * (sgn(x) * [-0.1444 * (clip(\|0.7071 * x\|, max=1.769) - 1.769)^2 + 0.5] + 0.5),
	24	+* where sgn(x) function is (x+0.000000000001)/\|(x+0.000000000001)\|.
	25	+
	26	+* @par Inputs:
	27	+* One input, including:
	28	+* x: An ND or 5HD tensor. Support 1D~8D. Must be one of the following types: bfloat16, float16, float32
	29	+
	30	+* @par Outputs:
	31	+* y: A Tensor. Has the same type as "x".
	32	+* @par Third-party framework compatibility
	33	+* Compatible with the TensorFlow operator FastGeluV2
	34	+*/
	35	+REG_OP(FastGeluV2)
	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(FastGeluV2)
	39	+} // namespace ge
	40	+
	41	+#endif // OPS_BUILT_IN_OP_PROTO_INC_FAST_GELU_V2_OPS_H_

Aexperimental/activation/fast_gelu_v2/op_host/fast_gelu_v2_def.cpp+53-0

@@ -0,0 +1,53 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file fast_gelu_v2_def.cpp
	13	+ * \brief FastGeluV2 operator definition
	14	+ *
	15	+ * Defines the operator prototype: single input "x", single output "y",
	16	+ * both supporting float32/float16/bfloat16 in ND format.
	17	+ * Currently registered for ascend950 (arch35) only.
	18	+ */
	19	+#include "register/op_def_registry.h"
	20	+
	21	+namespace ops {
	22	+class FastGeluV2 : public OpDef {
	23	+public:
	24	+ explicit FastGeluV2(const char* name) : OpDef(name)
	25	+ {
	26	+ // Input tensor: required, supports float32/float16/bfloat16, ND format
	27	+ this->Input("x")
	28	+ .ParamType(REQUIRED)
	29	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
	30	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	31	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	32	+ .AutoContiguous();
	33	+ // Output tensor: same dtype/format as input (elementwise operation)
	34	+ this->Output("y")
	35	+ .ParamType(REQUIRED)
	36	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
	37	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	38	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	39	+ .AutoContiguous();
	40	+
	41	+ OpAICoreConfig aicoreConfig950;
	42	+ aicoreConfig950.DynamicCompileStaticFlag(true)
	43	+ .DynamicFormatFlag(false)
	44	+ .DynamicRankSupportFlag(true)
	45	+ .DynamicShapeSupportFlag(true)
	46	+ .NeedCheckSupportFlag(false)
	47	+ .PrecisionReduceFlag(true)
	48	+ .ExtendCfgInfo("opFile.value", "fast_gelu_v2_apt");
	49	+ this->AICore().AddConfig("ascend950", aicoreConfig950);
	50	+ }
	51	+};
	52	+OP_ADD(FastGeluV2);
	53	+} // namespace ops

Aexperimental/activation/fast_gelu_v2/op_host/fast_gelu_v2_infershape.cpp+46-0

@@ -0,0 +1,46 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file fast_gelu_v2_infershape.cpp
	13	+ * \brief FastGeluV2 shape inference: output shape = input shape
	14	+ *
	15	+ * This is an elementwise operator with no broadcasting, so the output shape
	16	+ * is always identical to the input shape. Scalar inputs (dim=0) are handled
	17	+ * transparently by the framework and the Tiling layer (EnsureNotScalar).
	18	+ */
	19	+
	20	+#include "register/op_impl_registry.h"
	21	+#include "exe_graph/runtime/infer_shape_context.h"
	22	+
	23	+using namespace ge;
	24	+
	25	+namespace ops {
	26	+
	27	+static ge::graphStatus InferShape4FastGeluV2(gert::InferShapeContext* context)
	28	+{
	29	+ const gert::Shape* input_shape = context->GetInputShape(0);
	30	+ if (input_shape == nullptr) {
	31	+ return ge::GRAPH_FAILED;
	32	+ }
	33	+
	34	+ gert::Shape* output_shape = context->GetOutputShape(0);
	35	+ if (output_shape == nullptr) {
	36	+ return ge::GRAPH_FAILED;
	37	+ }
	38	+
	39	+ output_shape = input_shape;
	40	+
	41	+ return ge::GRAPH_SUCCESS;
	42	+}
	43	+
	44	+IMPL_OP_INFERSHAPE(FastGeluV2).InferShape(InferShape4FastGeluV2);
	45	+
	46	+} // namespace ops

Aexperimental/activation/fast_gelu_v2/op_host/fast_gelu_v2_tiling.cpp+221-0

@@ -0,0 +1,221 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file fast_gelu_v2_tiling.cpp
	13	+ * \brief FastGeluV2 tiling implementation (arch35)
	14	+ */
	15	+
	16	+#include "register/op_def_registry.h"
	17	+#include "op_common/log/log.h"
	18	+#include "op_common/op_host/util/math_util.h"
	19	+#include "op_common/op_host/util/platform_util.h"
	20	+#include "../op_kernel/fast_gelu_v2_tiling_data.h"
	21	+#include "../op_kernel/fast_gelu_v2_tiling_key.h"
	22	+
	23	+namespace optiling {
	24	+
	25	+using Ops::Base::CeilDiv;
	26	+using Ops::Base::CeilAlign;
	27	+using Ops::Base::FloorDiv;
	28	+using Ops::Base::FloorAlign;
	29	+using Ops::Base::GetUbBlockSize;
	30	+
	31	+constexpr uint32_t WS_SYS_SIZE = 0U;
	32	+// Double-buffer threshold: enable double buffer when total > 1024
	33	+constexpr int64_t MIN_SPLIT_THRESHOLD = 1024;
	34	+
	35	+static const gert::Shape g_vec_1_shape = {1};
	36	+
	37	+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape) {
	38	+ if (in_shape.GetDimNum() == 0) {
	39	+ return g_vec_1_shape;
	40	+ }
	41	+ return in_shape;
	42	+}
	43	+
	44	+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
	45	+{
	46	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	47	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	48	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	49	+ coreNum = ascendcPlatform.GetCoreNumAiv();
	50	+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
	51	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
	52	+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
	53	+ return ge::GRAPH_SUCCESS;
	54	+}
	55	+
	56	+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t& totalIdx, ge::DataType& dataType)
	57	+{
	58	+ auto inputX = context->GetInputShape(0);
	59	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
	60	+ auto inputShapeX = EnsureNotScalar(inputX->GetStorageShape());
	61	+
	62	+ auto outY = context->GetOutputShape(0);
	63	+ OP_CHECK_NULL_WITH_CONTEXT(context, outY);
	64	+ auto outShapeY = EnsureNotScalar(outY->GetStorageShape());
	65	+
	66	+ OP_CHECK_IF(
	67	+ inputShapeX.GetShapeSize() != outShapeY.GetShapeSize(),
	68	+ OP_LOGE(context, "FastGeluV2: input and output shape size mismatch: x=%ld, y=%ld",
	69	+ inputShapeX.GetShapeSize(), outShapeY.GetShapeSize()),
	70	+ return ge::GRAPH_FAILED);
	71	+
	72	+ totalIdx = inputShapeX.GetShapeSize();
	73	+
	74	+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
	75	+ auto inputDesc = context->GetInputDesc(0);
	76	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	77	+ dataType = inputDesc->GetDataType();
	78	+ if (supportedDtype.count(dataType) == 0) {
	79	+ OP_LOGE(context, "invalid dtype");
	80	+ return ge::GRAPH_FAILED;
	81	+ }
	82	+ return ge::GRAPH_SUCCESS;
	83	+}
	84	+
	85	+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
	86	+{
	87	+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
	88	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	89	+ currentWorkspace[0] = WS_SYS_SIZE;
	90	+ return ge::GRAPH_SUCCESS;
	91	+}
	92	+
	93	+static ge::graphStatus FastGeluV2TilingFunc(gert::TilingContext* context)
	94	+{
	95	+ // 1. Get platform info
	96	+ uint64_t ubSize;
	97	+ int64_t coreNum;
	98	+ OP_CHECK_IF(
	99	+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
	100	+ OP_LOGE(context, "GetPlatformInfo error"),
	101	+ return ge::GRAPH_FAILED);
	102	+
	103	+ // 2. Get shape and attribute info
	104	+ int64_t totalIdx;
	105	+ ge::DataType dataType;
	106	+ OP_CHECK_IF(
	107	+ GetShapeAttrsInfo(context, totalIdx, dataType) != ge::GRAPH_SUCCESS,
	108	+ OP_LOGE(context, "GetShapeAttrsInfo error"),
	109	+ return ge::GRAPH_FAILED);
	110	+
	111	+ // 3. Get workspace size
	112	+ OP_CHECK_IF(
	113	+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
	114	+ OP_LOGE(context, "GetWorkspaceSize error"),
	115	+ return ge::GRAPH_FAILED);
	116	+
	117	+ // Handle empty tensor: set safe tiling defaults so the kernel's Process()
	118	+ // computes loopCount = 0 and never enters the loop body.
	119	+ if (totalIdx == 0) {
	120	+ FastGeluV2TilingData* emptyTiling = context->GetTilingData<FastGeluV2TilingData>();
	121	+ OP_CHECK_NULL_WITH_CONTEXT(context, emptyTiling);
	122	+ OP_CHECK_IF(
	123	+ memset_s(emptyTiling, sizeof(FastGeluV2TilingData), 0, sizeof(FastGeluV2TilingData)) != EOK,
	124	+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
	125	+ emptyTiling->totalNum = 0;
	126	+ emptyTiling->blockFactor = 1;
	127	+ emptyTiling->ubFactor = 1;
	128	+ context->SetBlockDim(1);
	129	+ return ge::GRAPH_SUCCESS;
	130	+ }
	131	+
	132	+ // 4. Set tiling data
	133	+ FastGeluV2TilingData* tiling = context->GetTilingData<FastGeluV2TilingData>();
	134	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	135	+ OP_CHECK_IF(
	136	+ memset_s(tiling, sizeof(FastGeluV2TilingData), 0, sizeof(FastGeluV2TilingData)) != EOK,
	137	+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
	138	+
	139	+ // Determine type size based on dtype
	140	+ int64_t typeSize = 4; // default float32
	141	+ if (dataType == ge::DT_FLOAT16 \|\| dataType == ge::DT_BF16) {
	142	+ typeSize = 2;
	143	+ }
	144	+
	145	+ int64_t ubBlockSize = Ops::Base::GetUbBlockSize(context);
	146	+
	147	+ // Multi-core split: distribute elements evenly across available cores.
	148	+ // blockFactor is aligned to ubBlockSize to satisfy DMA alignment constraints.
	149	+ tiling->totalNum = totalIdx;
	150	+ tiling->blockFactor = CeilAlign(CeilDiv(totalIdx, coreNum), ubBlockSize);
	151	+ int64_t usedCoreNum = CeilDiv(totalIdx, tiling->blockFactor);
	152	+
	153	+ // UB split: calculate how many elements fit in one UB iteration.
	154	+ //
	155	+ // Buffer layout varies by dtype:
	156	+ // FP16: I/O buffers are FP16 (2B), compute buffers are FP32 (4B) for precision.
	157	+ // Single-buffer: 12(in) + 12(out) + 34(tmp) + 24(cast) = 24 bytes/elem
	158	+ // Double-buffer: 22(in) + 22(out) + 34(tmp) + 24(cast) = 28 bytes/elem
	159	+ // BF16: All buffers use bfloat16 (2B), no FP32 promotion.
	160	+ // Single: 5 * 2 = 10 bytes/elem; Double: 7 * 2 = 14 bytes/elem
	161	+ // FP32: All buffers use float (4B).
	162	+ // Single: 5 * 4 = 20 bytes/elem; Double: 7 * 4 = 28 bytes/elem
	163	+ //
	164	+ // Double-buffer is enabled when totalNum > MIN_SPLIT_THRESHOLD (1024) to allow
	165	+ // overlap of CopyIn/Compute/CopyOut via pipeline parallelism.
	166	+ uint64_t useDoubleBuffer = (totalIdx > MIN_SPLIT_THRESHOLD) ? 1 : 0;
	167	+
	168	+ int64_t bytesPerElement;
	169	+ if (dataType == ge::DT_FLOAT16) {
	170	+ // FP16: I/O buffers are FP16 (2 bytes), compute buffers are FP32 (4 bytes)
	171	+ // Single: 12(in) + 12(out) + 34(tmp) + 24(cast) = 24
	172	+ // Double: 22(in) + 22(out) + 34(tmp) + 24(cast) = 28
	173	+ int64_t ioBufCount = useDoubleBuffer ? 4 : 2; // input + output, doubled if DB
	174	+ bytesPerElement = ioBufCount * 2 + 5 * 4; // 5 FP32 buffers (3 tmp + 2 cast)
	175	+ } else if (dataType == ge::DT_BF16) {
	176	+ // BF16: all buffers use bfloat16 (2 bytes), no cast to FP32
	177	+ // Single: 1 input + 3 tmp + 1 output = 5 buffers * 2 bytes = 10
	178	+ // Double: (1 input + 1 output)2 + 3 tmp = 7 buffers 2 bytes = 14
	179	+ int64_t bufferNum = useDoubleBuffer ? 7 : 5;
	180	+ bytesPerElement = bufferNum * typeSize;
	181	+ } else {
	182	+ // FP32: all buffers same size (4 bytes)
	183	+ // Single: 1 input + 3 tmp + 1 output = 5 buffers * 4 bytes = 20
	184	+ // Double: (1 input + 1 output)2 + 3 tmp = 7 buffers 4 bytes = 28
	185	+ int64_t bufferNum = useDoubleBuffer ? 7 : 5;
	186	+ bytesPerElement = bufferNum * typeSize;
	187	+ }
	188	+ tiling->ubFactor = FloorAlign(
	189	+ FloorDiv(static_cast<int64_t>(ubSize), bytesPerElement), ubBlockSize);
	190	+
	191	+ // Safeguard: ubFactor must be at least 1 ubBlockSize to make progress
	192	+ if (tiling->ubFactor <= 0) {
	193	+ OP_LOGE(context, "FastGeluV2: ubFactor is 0; UB too small for even one block "
	194	+ "(ubSize=%lu, bytesPerElement=%ld)", ubSize, bytesPerElement);
	195	+ return ge::GRAPH_FAILED;
	196	+ }
	197	+
	198	+ // Safeguard: blockFactor must be at least ubBlockSize
	199	+ if (tiling->blockFactor <= 0) {
	200	+ OP_LOGE(context, "FastGeluV2: blockFactor is 0; unexpected tiling calculation error");
	201	+ return ge::GRAPH_FAILED;
	202	+ }
	203	+
	204	+ context->SetBlockDim(usedCoreNum);
	205	+
	206	+ // 5. Set TilingKey
	207	+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
	208	+ ASCENDC_TPL_SEL_PARAM(context, dTypeX, static_cast<uint32_t>(useDoubleBuffer));
	209	+ return ge::GRAPH_SUCCESS;
	210	+}
	211	+
	212	+static ge::graphStatus TilingParseForFastGeluV2([[maybe_unused]] gert::TilingParseContext* context)
	213	+{
	214	+ return ge::GRAPH_SUCCESS;
	215	+}
	216	+
	217	+struct FastGeluV2CompileInfo {};
	218	+
	219	+IMPL_OP_OPTILING(FastGeluV2).Tiling(FastGeluV2TilingFunc).TilingParse<FastGeluV2CompileInfo>(TilingParseForFastGeluV2);
	220	+
	221	+} // namespace optiling

Aexperimental/activation/fast_gelu_v2/op_kernel/fast_gelu_v2.h+264-0

@@ -0,0 +1,264 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file fast_gelu_v2.h
	13	+ * \brief FastGeluV2 kernel class definition (arch35)
	14	+ *
	15	+ * float32/float16 path (original formula):
	16	+ * y = x * (sgn(x) * [-0.1444 * (clip(\|0.7071*x\|, max=1.769) - 1.769)^2 + 0.5] + 0.5)
	17	+ * where sgn(x) = (x + 1e-12) / \|x + 1e-12\|
	18	+ *
	19	+ * bfloat16 path (equivalent formula, avoids Abs/Div):
	20	+ * y = \|x\| * inner + 0.5 * x
	21	+ * where inner = -0.1444 * (clip(0.7071 * \|x\|, max=1.769) - 1.769)^2 + 0.5
	22	+ * \|x\| via Muls(-1) + Max
	23	+ *
	24	+ * Template parameters:
	25	+ * - T: data type (half, float, bfloat16_t)
	26	+ * - BUFFER_MODE: 0=single buffer, 1=double buffer
	27	+ */
	28	+#ifndef FAST_GELU_V2_H
	29	+#define FAST_GELU_V2_H
	30	+
	31	+#include "kernel_operator.h"
	32	+#include "kernel_tiling/kernel_tiling.h"
	33	+#include "fast_gelu_v2_tiling_data.h"
	34	+#include "fast_gelu_v2_tiling_key.h"
	35	+
	36	+namespace NsFastGeluV2 {
	37	+
	38	+using namespace AscendC;
	39	+
	40	+// Helper trait: compute type is float for half (FP16->FP32 promotion), T otherwise
	41	+template <typename T>
	42	+struct ComputeType { using type = T; };
	43	+template <>
	44	+struct ComputeType<half> { using type = float; };
	45	+
	46	+template <typename T, int BUFFER_MODE>
	47	+class FastGeluV2 {
	48	+ static constexpr int32_t BUFFER_NUM = BUFFER_MODE ? 2 : 1;
	49	+ using CT = typename ComputeType<T>::type; // Compute type: float for half, T otherwise
	50	+ static constexpr bool NEED_CAST = !std::is_same<T, CT>::value;
	51	+
	52	+public:
	53	+ __aicore__ inline FastGeluV2() {};
	54	+
	55	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const FastGeluV2TilingData* tilingData);
	56	+ __aicore__ inline void Process();
	57	+
	58	+private:
	59	+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
	60	+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
	61	+ __aicore__ inline void Compute(int64_t currentNum);
	62	+
	63	+private:
	64	+ TPipe pipe;
	65	+ TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX;
	66	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueY;
	67	+ TBuf<TPosition::VECCALC> tmpBuf1;
	68	+ TBuf<TPosition::VECCALC> tmpBuf2;
	69	+ TBuf<TPosition::VECCALC> tmpBuf3;
	70	+ // Additional buffers for FP16->FP32 cast (input cast and output cast)
	71	+ TBuf<TPosition::VECCALC> castInBuf;
	72	+ TBuf<TPosition::VECCALC> castOutBuf;
	73	+
	74	+ GlobalTensor<T> inputGMX;
	75	+ GlobalTensor<T> outputGMY;
	76	+
	77	+ int64_t blockLength_ = 0;
	78	+ int64_t ubLength_ = 0;
	79	+};
	80	+
	81	+template <typename T, int BUFFER_MODE>
	82	+__aicore__ inline void FastGeluV2<T, BUFFER_MODE>::Init(GM_ADDR x, GM_ADDR y,
	83	+ const FastGeluV2TilingData* tilingData)
	84	+{
	85	+ int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx();
	86	+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
	87	+ ubLength_ = tilingData->ubFactor;
	88	+
	89	+ inputGMX.SetGlobalBuffer((__gm__ T)x + tilingData->blockFactor AscendC::GetBlockIdx(), blockLength_);
	90	+ outputGMY.SetGlobalBuffer((__gm__ T)y + tilingData->blockFactor AscendC::GetBlockIdx(), blockLength_);
	91	+
	92	+ pipe.InitBuffer(inputQueueX, BUFFER_NUM, ubLength_ * sizeof(T));
	93	+ pipe.InitBuffer(outputQueueY, BUFFER_NUM, ubLength_ * sizeof(T));
	94	+ // Temporary buffers use CT (compute type) size: float for FP16, T for others
	95	+ pipe.InitBuffer(tmpBuf1, ubLength_ * sizeof(CT));
	96	+ pipe.InitBuffer(tmpBuf2, ubLength_ * sizeof(CT));
	97	+ pipe.InitBuffer(tmpBuf3, ubLength_ * sizeof(CT));
	98	+ if constexpr (NEED_CAST) {
	99	+ // Cast buffers: castInBuf holds FP32 input, castOutBuf holds FP32 output before cast back
	100	+ pipe.InitBuffer(castInBuf, ubLength_ * sizeof(CT));
	101	+ pipe.InitBuffer(castOutBuf, ubLength_ * sizeof(CT));
	102	+ }
	103	+}
	104	+
	105	+template <typename T, int BUFFER_MODE>
	106	+__aicore__ inline void FastGeluV2<T, BUFFER_MODE>::CopyIn(int64_t progress, int64_t currentNum)
	107	+{
	108	+ AscendC::LocalTensor<T> xLocal = inputQueueX.template AllocTensor<T>();
	109	+ AscendC::DataCopyParams copyParams;
	110	+ copyParams.blockCount = 1;
	111	+ copyParams.blockLen = currentNum * sizeof(T);
	112	+ copyParams.srcStride = 0;
	113	+ copyParams.dstStride = 0;
	114	+ AscendC::DataCopyPad(xLocal, inputGMX[progress * ubLength_], copyParams, {false, 0, 0, 0});
	115	+ inputQueueX.EnQue(xLocal);
	116	+}
	117	+
	118	+template <typename T, int BUFFER_MODE>
	119	+__aicore__ inline void FastGeluV2<T, BUFFER_MODE>::CopyOut(int64_t progress, int64_t currentNum)
	120	+{
	121	+ AscendC::LocalTensor<T> yLocal = outputQueueY.template DeQue<T>();
	122	+ AscendC::DataCopyParams copyParams;
	123	+ copyParams.blockCount = 1;
	124	+ copyParams.blockLen = currentNum * sizeof(T);
	125	+ copyParams.srcStride = 0;
	126	+ copyParams.dstStride = 0;
	127	+ AscendC::DataCopyPad(outputGMY[progress * ubLength_], yLocal, copyParams);
	128	+ outputQueueY.FreeTensor(yLocal);
	129	+}
	130	+
	131	+template <typename T, int BUFFER_MODE>
	132	+__aicore__ inline void FastGeluV2<T, BUFFER_MODE>::Compute(int64_t currentNum)
	133	+{
	134	+ AscendC::LocalTensor<T> xLocal = inputQueueX.template DeQue<T>();
	135	+ AscendC::LocalTensor<T> yLocal = outputQueueY.template AllocTensor<T>();
	136	+
	137	+ // Get compute-type tensors (CT = float for FP16, T for FP32)
	138	+ AscendC::LocalTensor<CT> tmp1 = tmpBuf1.Get<CT>();
	139	+ AscendC::LocalTensor<CT> tmp2 = tmpBuf2.Get<CT>();
	140	+ AscendC::LocalTensor<CT> tmp3 = tmpBuf3.Get<CT>();
	141	+
	142	+ // For FP16: Cast input to FP32 for computation, use FP32 throughout, cast back at end
	143	+ AscendC::LocalTensor<CT> xCompute;
	144	+ if constexpr (NEED_CAST) {
	145	+ xCompute = castInBuf.Get<CT>();
	146	+ AscendC::Cast(xCompute, xLocal, AscendC::RoundMode::CAST_NONE, currentNum);
	147	+ } else {
	148	+ // T == CT, xLocal is already the compute type
	149	+ xCompute = xLocal;
	150	+ }
	151	+
	152	+ if constexpr (std::is_same<T, bfloat16_t>::value) {
	153	+ // bfloat16 equivalent formula: y = \|x\| * inner + 0.5 * x
	154	+ // Avoids Abs and Div which do not support bfloat16
	155	+
	156	+ // Step 1: neg_x = -x
	157	+ AscendC::Muls(tmp1, xCompute, (CT)(-1.0f), currentNum);
	158	+
	159	+ // Step 2: abs_x = max(x, -x)
	160	+ AscendC::Max(tmp2, xCompute, tmp1, currentNum);
	161	+
	162	+ // Step 3: scaled_x = 0.7071 * abs_x
	163	+ AscendC::Muls(tmp1, tmp2, (CT)(0.7071f), currentNum);
	164	+
	165	+ // Step 4: clipped_x = min(scaled_x, 1.769)
	166	+ AscendC::Mins(tmp1, tmp1, (CT)(1.769f), currentNum);
	167	+
	168	+ // Step 5: diff = clipped_x - 1.769
	169	+ AscendC::Adds(tmp1, tmp1, (CT)(-1.769f), currentNum);
	170	+
	171	+ // Step 6: sq = diff^2
	172	+ AscendC::Mul(tmp1, tmp1, tmp1, currentNum);
	173	+
	174	+ // Step 7a: inner = -0.1444 * sq
	175	+ AscendC::Muls(tmp1, tmp1, (CT)(-0.1444f), currentNum);
	176	+
	177	+ // Step 7b: inner = inner + 0.5
	178	+ AscendC::Adds(tmp1, tmp1, (CT)(0.5f), currentNum);
	179	+
	180	+ // Step 8: term1 = abs_x * inner
	181	+ AscendC::Mul(tmp3, tmp2, tmp1, currentNum);
	182	+
	183	+ // Step 9: term2 = 0.5 * x
	184	+ AscendC::Muls(tmp1, xCompute, (CT)(0.5f), currentNum);
	185	+
	186	+ // Step 10: y = term1 + term2
	187	+ AscendC::Add(yLocal, tmp3, tmp1, currentNum);
	188	+ } else {
	189	+ // float32 / float16 path: original formula
	190	+ // FastGeluV2(x) = x * (sgn(x) * inner + 0.5)
	191	+ // where inner = -0.1444 * (clip(0.7071 * \|x\|, max=1.769) - 1.769)^2 + 0.5
	192	+ // and sgn(x) = (x + eps) / \|x + eps\|
	193	+
	194	+ // Step 1: x_eps = x + eps
	195	+ AscendC::Adds(tmp1, xCompute, (CT)1e-12, currentNum);
	196	+
	197	+ // Step 2a: abs_xeps = \|x_eps\|
	198	+ AscendC::Abs(tmp2, tmp1, currentNum);
	199	+
	200	+ // Step 2b: sgn_x = x_eps / abs_xeps
	201	+ AscendC::Div(tmp1, tmp1, tmp2, currentNum);
	S S sunday4月16日这里 `sgn(x)` 的计算在极小概率点会出现除零：当 `x = -1e-12` 时，`x_eps = 0` 且 `abs_xeps = 0`，随后 `x_eps / abs_xeps` 变成 0/0，结果可能是 NaN。建议在分母上加一个最小保护值（例如 `max(abs_xeps, eps)`），或者改成无除法的符号实现，避免数值不稳定。陈陈佳良 4月16日评论： 4月16日评论：
	202	+
	203	+ // Step 3: abs_x = \|x\|
	204	+ AscendC::Abs(tmp2, xCompute, currentNum);
	205	+
	206	+ // Step 4: scaled_x = 0.7071 * abs_x
	207	+ AscendC::Muls(tmp2, tmp2, (CT)0.7071, currentNum);
	208	+
	209	+ // Step 5: clipped_x = min(scaled_x, 1.769)
	210	+ AscendC::Mins(tmp2, tmp2, (CT)1.769, currentNum);
	211	+
	212	+ // Step 6: diff = clipped_x - 1.769
	213	+ AscendC::Adds(tmp2, tmp2, (CT)(-1.769), currentNum);
	214	+
	215	+ // Step 7: sq = diff^2
	216	+ AscendC::Mul(tmp2, tmp2, tmp2, currentNum);
	217	+
	218	+ // Step 8a: inner = -0.1444 * sq
	219	+ AscendC::Muls(tmp2, tmp2, (CT)(-0.1444), currentNum);
	220	+
	221	+ // Step 8b: inner = inner + 0.5
	222	+ AscendC::Adds(tmp2, tmp2, (CT)0.5, currentNum);
	223	+
	224	+ // Step 9a: bracket = sgn_x * inner
	225	+ AscendC::Mul(tmp3, tmp1, tmp2, currentNum);
	226	+
	227	+ // Step 9b: bracket = bracket + 0.5
	228	+ AscendC::Adds(tmp3, tmp3, (CT)0.5, currentNum);
	229	+
	230	+ // Step 10: y = x * bracket
	231	+ if constexpr (NEED_CAST) {
	232	+ // Compute result in FP32, then cast back to FP16
	233	+ AscendC::LocalTensor<CT> resultFloat = castOutBuf.Get<CT>();
	234	+ AscendC::Mul(resultFloat, xCompute, tmp3, currentNum);
	235	+ AscendC::Cast(yLocal, resultFloat, AscendC::RoundMode::CAST_ROUND, currentNum);
	236	+ } else {
	237	+ AscendC::Mul(yLocal, xCompute, tmp3, currentNum);
	238	+ }
	239	+ }
	240	+
	241	+ outputQueueY.template EnQue<T>(yLocal);
	242	+ inputQueueX.FreeTensor(xLocal);
	243	+}
	244	+
	245	+template <typename T, int BUFFER_MODE>
	246	+__aicore__ inline void FastGeluV2<T, BUFFER_MODE>::Process()
	247	+{
	248	+ // Guard: skip computation if this core has no work assigned
	249	+ if (blockLength_ <= 0) {
	250	+ return;
	251	+ }
	252	+ // Loop over UB-sized tiles within this core's block.
	253	+ // The last tile may process fewer than ubLength_ elements (tail handling).
	254	+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
	255	+ for (int64_t i = 0; i < loopCount; i++) {
	256	+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
	257	+ CopyIn(i, currentNum);
	258	+ Compute(currentNum);
	259	+ CopyOut(i, currentNum);
	260	+ }
	261	+}
	262	+
	263	+} // namespace NsFastGeluV2
	264	+#endif // FAST_GELU_V2_H

Aexperimental/activation/fast_gelu_v2/op_kernel/fast_gelu_v2_apt.cpp+36-0

@@ -0,0 +1,36 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file fast_gelu_v2_arch35.cpp
	13	+ * \brief FastGeluV2 kernel entry (arch35)
	14	+ *
	15	+ * This is the global kernel function registered via ASCENDC_TPL_ARGS_DECL.
	16	+ * The template parameters D_T_X and BUFFER_MODE are resolved at compile time
	17	+ * based on the TilingKey set by the host-side Tiling function.
	18	+ *
	19	+ * Parameters (in order, fixed by CANN convention):
	20	+ * x - Input tensor in Global Memory
	21	+ * y - Output tensor in Global Memory
	22	+ * workspace - System workspace (unused by this operator)
	23	+ * tiling - Serialized TilingData structure from host
	24	+ */
	25	+
	26	+#include "fast_gelu_v2.h"
	27	+
	28	+template <typename D_T_X, int BUFFER_MODE>
	29	+__global__ __aicore__ void fast_gelu_v2(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
	30	+{
	31	+ REGISTER_TILING_DEFAULT(FastGeluV2TilingData);
	32	+ GET_TILING_DATA_WITH_STRUCT(FastGeluV2TilingData, tilingData, tiling);
	33	+ NsFastGeluV2::FastGeluV2<D_T_X, BUFFER_MODE> op;
	34	+ op.Init(x, y, &tilingData);
	35	+ op.Process();
	36	+}

Aexperimental/activation/fast_gelu_v2/op_kernel/fast_gelu_v2_tiling_data.h+28-0

@@ -0,0 +1,28 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file fast_gelu_v2_tiling_data.h
	13	+ * \brief FastGeluV2 TilingData structure definition
	14	+ *
	15	+ * This structure is filled by the host-side Tiling function and read by the
	16	+ * device-side Kernel via GET_TILING_DATA_WITH_STRUCT. All fields must be
	17	+ * int64_t for cross-platform compatibility.
	18	+ */
	19	+
	20	+#ifndef _FAST_GELU_V2_TILING_DATA_H_
	21	+#define _FAST_GELU_V2_TILING_DATA_H_
	22	+
	23	+struct FastGeluV2TilingData {
	24	+ int64_t totalNum = 0; // Total number of input elements across all dimensions
	25	+ int64_t blockFactor = 0; // Number of elements assigned to each AI Core (aligned to ubBlockSize)
	26	+ int64_t ubFactor = 0; // Number of elements processed per UB iteration (aligned to ubBlockSize)
	27	+};
	28	+#endif

Aexperimental/activation/fast_gelu_v2/op_kernel/fast_gelu_v2_tiling_key.h+47-0

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