已合并
[CANNBot]Relu6算子支持Ascend950 AscendC实现 #3794
wangweidong创建于 4月14日
[CANNBot]Relu6算子支持Ascend950 AscendC实现 #3794
已合并
共 11 个文件变更+900-0
| @@ -0,0 +1,21 @@ | |||
| 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 | +# NOTE: Portions of this code were AI-generated and have been | ||
| 11 | +# technically reviewed for functional accuracy and security | ||
| 12 | + | ||
| 13 | +file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) | ||
| 14 | +if(NOT ENABLE_TEST) | ||
| 15 | + list(REMOVE_ITEM CURRENT_DIRS tests) | ||
| 16 | +endif() | ||
| 17 | +foreach(SUB_DIR ${CURRENT_DIRS}) | ||
| 18 | + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt") | ||
| 19 | + add_subdirectory(${SUB_DIR}) | ||
| 20 | + endif() | ||
| 21 | +endforeach() | ||
| @@ -0,0 +1,104 @@ | |||
| 1 | +# Relu6 算子 | ||
| 2 | + | ||
| 3 | +Relu6 激活函数算子的 Ascend C 自定义算子实现，面向 Ascend 950 芯片。 | ||
| 4 | + | ||
| 5 | +## 功能简介 | ||
| 6 | + | ||
| 7 | +Relu6 算子按元素对输入张量执行 Relu6 激活函数计算，将输入张量中的每个元素值限制在 `[0, 6]` 区间内。 | ||
| 8 | + | ||
| 9 | +### 数学公式 | ||
| 10 | + | ||
| 11 | +``` | ||
| 12 | +y = min(max(x, 0), 6) | ||
| 13 | +``` | ||
| 14 | + | ||
| 15 | +等价于分段函数： | ||
| 16 | + | ||
| 17 | +``` | ||
| 18 | +y = { 0, if x < 0 | ||
| 19 | + { x, if 0 <= x <= 6 | ||
| 20 | + { 6, if x > 6 | ||
| 21 | +``` | ||
| 22 | + | ||
| 23 | +## 支持特性 | ||
| 24 | + | ||
| 25 | +### 数据类型 | ||
| 26 | + | ||
| 27 | +| 数据类型 | C 类型 | | ||
| 28 | +|---------|--------| | ||
| 29 | +| float16 | half | | ||
| 30 | +| float | float | | ||
| 31 | +| int32 | int32_t | | ||
| 32 | +| bfloat16 | bfloat16_t | | ||
| 33 | + | ||
| 34 | +### Shape 约束 | ||
| 35 | + | ||
| 36 | +- 输入 x 和输出 y 的 shape 必须相同 | ||
| 37 | +- 最高支持 8 维张量 | ||
| 38 | +- 支持任意维度（1D ~ 8D），使用 FORMAT_ND（任意维度连续排布） | ||
| 39 | + | ||
| 40 | +### 芯片支持 | ||
| 41 | + | ||
| 42 | +| 芯片 | 状态 | | ||
| 43 | +|-----|------| | ||
| 44 | +| Ascend 950PR | 支持 | | ||
| 45 | +| Ascend 950DT | 支持 | | ||
| 46 | + | ||
| 47 | +## 工程结构 | ||
| 48 | + | ||
| 49 | +``` | ||
| 50 | +relu6/ | ||
| 51 | +├── examples/ # 调用示例 | ||
| 52 | +│ └── test_aclnn_relu6.cpp # C++ 调用示例 | ||
| 53 | +├── op_host/ # 主机端算子注册代码 | ||
| 54 | +│ ├── relu6_def.cpp # 算子定义 | ||
| 55 | +│ └── relu6_tiling.cpp # Tiling 策略 | ||
| 56 | +├── op_kernel/ # Kernel 实现代码 | ||
| 57 | +│ ├── relu6.cpp # Kernel 入口 | ||
| 58 | +│ ├── relu6.h # Kernel 类实现 | ||
| 59 | +│ ├── relu6_tiling_data.h # TilingData 结构体 | ||
| 60 | +│ └── relu6_tiling_key.h # TilingKey 模板参数 | ||
| 61 | +├── tests/ # 测试代码 | ||
| 62 | +├── CMakeLists.txt # 顶层 CMake 配置 | ||
| 63 | +└── README.md # 本文件 | ||
| 64 | +``` | ||
| 65 | + | ||
| 66 | +## API 接口 | ||
| 67 | + | ||
| 68 | +算子提供两段式 ACLNN 接口： | ||
| 69 | + | ||
| 70 | +### aclnnRelu6GetWorkspaceSize | ||
| 71 | + | ||
| 72 | +计算执行 Relu6 算子所需的 workspace 大小。 | ||
| 73 | + | ||
| 74 | +```cpp | ||
| 75 | +extern "C" aclnnStatus aclnnRelu6GetWorkspaceSize( | ||
| 76 | + const aclTensor* x, // 输入张量 | ||
| 77 | + const aclTensor* y, // 输出张量 | ||
| 78 | + uint64_t* workspaceSize, // 输出：workspace 大小 | ||
| 79 | + aclOpExecutor** executor); // 输出：执行器指针 | ||
| 80 | +``` | ||
| 81 | + | ||
| 82 | +### aclnnRelu6 | ||
| 83 | + | ||
| 84 | +执行 Relu6 算子计算。 | ||
| 85 | + | ||
| 86 | +```cpp | ||
| 87 | +extern "C" aclnnStatus aclnnRelu6( | ||
| 88 | + void* workspace, // workspace 内存地址 | ||
| 89 | + uint64_t workspaceSize, // workspace 大小 | ||
| 90 | + aclOpExecutor* executor, // 执行器指针 | ||
| 91 | + aclrtStream stream); // 运行流 | ||
| 92 | +``` | ||
| 93 | + | ||
| 94 | +## 实现原理 | ||
| 95 | + | ||
| 96 | +Relu6 算子在 AI Core 上通过两次矢量计算实现： | ||
| 97 | + | ||
| 98 | +1. **Maxs 计算**：`tmp = max(x, 0)` -- 将小于 0 的值截断为 0 | ||
| 99 | +2. **Mins 计算**：`y = min(tmp, 6)` -- 将大于 6 的值截断为 6 | ||
| 100 | + | ||
| 101 | +内部采用多核并行 + UB 缓冲优化策略： | ||
| 102 | +- 根据数据量和 AI Core 数量自动切分数据，实现多核并行处理 | ||
| 103 | +- UB 中使用 3 个 LocalTensor（input、tmp、output），根据 UB 大小动态计算单次循环处理量 | ||
| 104 | +- LocalTensor 起始地址 32 字节对齐，尾部数据通过 DataCopyPad 自动处理 | ||
| @@ -0,0 +1,209 @@ | |||
| 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 | + * NOTE: Portions of this code were AI-generated and have been | ||
| 13 | + * technically reviewed for functional accuracy and security | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +/** | ||
| 17 | + * @file test_aclnn_relu6.cpp | ||
| 18 | + * @brief Relu6 算子调用示例 | ||
| 19 | + * | ||
| 20 | + * Relu6 计算公式：y = min(max(x, 0), 6) | ||
| 21 | + */ | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | +using DataType = float; | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + do { \ | ||
| 32 | + if (!(cond)) { \ | ||
| 33 | + return_expr; \ | ||
| 34 | + } \ | ||
| 35 | + } while (0) | ||
| 36 | + | ||
| 37 | + | ||
| 38 | + do { \ | ||
| 39 | + printf(message, ##__VA_ARGS__); \ | ||
| 40 | + } while (0) | ||
| 41 | + | ||
| 42 | +int64_t GetShapeSize(const std::vector<int64_t>& shape) | ||
| 43 | +{ | ||
| 44 | + int64_t shapeSize = 1; | ||
| 45 | + for (auto i : shape) { | ||
| 46 | + shapeSize *= i; | ||
| 47 | + } | ||
| 48 | + return shapeSize; | ||
| 49 | +} | ||
| 50 | + | ||
| 51 | +void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr) | ||
| 52 | +{ | ||
| 53 | + auto size = GetShapeSize(shape); | ||
| 54 | + std::vector<DataType> resultData(size, 0); | ||
| 55 | + auto ret = aclrtMemcpy( | ||
| 56 | + resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]), | ||
| 57 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 58 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return); | ||
| 59 | + for (int64_t i = 0; i < size; i++) { | ||
| 60 | + LOG_PRINT("relu6 result[%ld] is: %f\n", i, resultData[i]); | ||
| 61 | + } | ||
| 62 | +} | ||
| 63 | + | ||
| 64 | +/** | ||
| 65 | + * @brief 在 CPU 上计算 Relu6 参考结果 | ||
| 66 | + */ | ||
| 67 | +void Relu6CpuReference(const DataType* input, DataType* output, size_t size) | ||
| 68 | +{ | ||
| 69 | + for (size_t i = 0; i < size; ++i) { | ||
| 70 | + float val = input[i]; | ||
| 71 | + if (val < 0.0f) { | ||
| 72 | + output[i] = 0.0f; | ||
| 73 | + } else if (val > 6.0f) { | ||
| 74 | + output[i] = 6.0f; | ||
| 75 | + } else { | ||
| 76 | + output[i] = val; | ||
| 77 | + } | ||
| 78 | + } | ||
| 79 | +} | ||
| 80 | + | ||
| 81 | +int Init(int32_t deviceId, aclrtStream* stream) | ||
| 82 | +{ | ||
| 83 | + auto ret = aclInit(nullptr); | ||
| 84 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 85 | + ret = aclrtSetDevice(deviceId); | ||
| 86 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 87 | + ret = aclrtCreateStream(stream); | ||
| 88 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 89 | + return 0; | ||
| 90 | +} | ||
| 91 | + | ||
| 92 | +template <typename T> | ||
| 93 | +int CreateAclTensor( | ||
| 94 | + const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, | ||
| 95 | + aclDataType dataType, aclTensor** tensor) | ||
| 96 | +{ | ||
| 97 | + auto size = GetShapeSize(shape) * sizeof(T); | ||
| 98 | + auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 99 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 100 | + ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 101 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); | ||
| 102 | + | ||
| 103 | + std::vector<int64_t> strides(shape.size(), 1); | ||
| 104 | + for (int64_t i = shape.size() - 2; i >= 0; i--) { | ||
| 105 | + strides[i] = shape[i + 1] * strides[i + 1]; | ||
| 106 | + } | ||
| 107 | + | ||
| 108 | + *tensor = aclCreateTensor( | ||
| 109 | + shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), | ||
| 110 | + *deviceAddr); | ||
| 111 | + return 0; | ||
| 112 | +} | ||
| 113 | + | ||
| 114 | +int main() | ||
| 115 | +{ | ||
| 116 | + // 1. 调用acl进行device/stream初始化 | ||
| 117 | + int32_t deviceId = 0; | ||
| 118 | + aclrtStream stream; | ||
| 119 | + auto ret = Init(deviceId, &stream); | ||
| 120 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); | ||
| 121 | + | ||
| 122 | + // 2. 构造输入与输出 | ||
| 123 | + // 输入 shape: [2, 4]，包含典型 Relu6 测试值：负数、0、正数、大于 6 的数 | ||
| 124 | + aclTensor* selfX = nullptr; | ||
| 125 | + void* selfXDeviceAddr = nullptr; | ||
| 126 | + std::vector<int64_t> selfXShape = {2, 4}; | ||
| 127 | + std::vector<DataType> selfXHostData = {-2.0f, 0.0f, 3.0f, 7.0f, 1.5f, -0.5f, 6.0f, 10.0f}; | ||
| 128 | + ret = CreateAclTensor(selfXHostData, selfXShape, &selfXDeviceAddr, aclDataType::ACL_FLOAT, &selfX); | ||
| 129 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 130 | + | ||
| 131 | + aclTensor* out = nullptr; | ||
| 132 | + void* outDeviceAddr = nullptr; | ||
| 133 | + std::vector<int64_t> outShape = {2, 4}; | ||
| 134 | + std::vector<DataType> outHostData(8, 0); | ||
| 135 | + ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out); | ||
| 136 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 137 | + | ||
| 138 | + // 计算 CPU 参考结果用于对比 | ||
| 139 | + std::vector<DataType> cpuRefOut(8); | ||
| 140 | + Relu6CpuReference(selfXHostData.data(), cpuRefOut.data(), 8); | ||
| 141 | + | ||
| 142 | + LOG_PRINT("Before GetWorkspaceSize: selfX=%p, out=%p\n", (void*)selfX, (void*)out); | ||
| 143 | + LOG_PRINT("Before GetWorkspaceSize: selfXDeviceAddr=%p, outDeviceAddr=%p\n", | ||
| 144 | + selfXDeviceAddr, outDeviceAddr); | ||
| 145 | + | ||
| 146 | + // 3. 调用 aclnnRelu6GetWorkspaceSize（第一段接口） | ||
| 147 | + uint64_t workspaceSize = 0; | ||
| 148 | + aclOpExecutor* executor; | ||
| 149 | + ret = aclnnRelu6GetWorkspaceSize(selfX, out, &workspaceSize, &executor); | ||
| 150 | + LOG_PRINT("aclnnRelu6GetWorkspaceSize returned %d, workspaceSize=%llu, executor=%p\n", | ||
| 151 | + ret, (unsigned long long)workspaceSize, (void*)executor); | ||
| 152 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnRelu6GetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 153 | + | ||
| 154 | + // 根据 workspaceSize 申请 device 内存 | ||
| 155 | + void* workspaceAddr = nullptr; | ||
| 156 | + if (workspaceSize > static_cast<uint64_t>(0)) { | ||
| 157 | + ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 158 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret); | ||
| 159 | + } | ||
| 160 | + | ||
| 161 | + // 4. 调用 aclnnRelu6（第二段接口） | ||
| 162 | + ret = aclnnRelu6(workspaceAddr, workspaceSize, executor, stream); | ||
| 163 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnRelu6 failed. ERROR: %d\n", ret); return ret); | ||
| 164 | + | ||
| 165 | + // 5. 同步等待任务执行结束 | ||
| 166 | + ret = aclrtSynchronizeStream(stream); | ||
| 167 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 168 | + | ||
| 169 | + // 6. 获取输出的值 | ||
| 170 | + PrintOutResult(outShape, &outDeviceAddr); | ||
| 171 | + | ||
| 172 | + // 7. 简单验证 | ||
| 173 | + std::vector<DataType> npuOut(8); | ||
| 174 | + ret = aclrtMemcpy(npuOut.data(), npuOut.size() * sizeof(DataType), outDeviceAddr, | ||
| 175 | + 8 * sizeof(DataType), ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 176 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result for verification failed. ERROR: %d\n", ret); return ret); | ||
| 177 | + | ||
| 178 | + bool pass = true; | ||
| 179 | + for (size_t i = 0; i < 8; i++) { | ||
| 180 | + float diff = std::abs(npuOut[i] - cpuRefOut[i]); | ||
| 181 | + if (diff > 1e-2f) { | ||
| 182 | + LOG_PRINT("Mismatch at index %ld: NPU=%f, CPU=%f, diff=%f\n", i, npuOut[i], cpuRefOut[i], diff); | ||
| 183 | + pass = false; | ||
| 184 | + } | ||
| 185 | + } | ||
| 186 | + if (pass) { | ||
| 187 | + LOG_PRINT("Verification passed! NPU output matches CPU reference.\n"); | ||
| 188 | + } else { | ||
| 189 | + LOG_PRINT("Verification FAILED!\n"); | ||
| 190 | + } | ||
| 191 | + | ||
| 192 | + // 8. 释放 aclTensor | ||
| 193 | + aclDestroyTensor(selfX); | ||
| 194 | + aclDestroyTensor(out); | ||
| 195 | + | ||
| 196 | + // 9. 释放 device 资源 | ||
| 197 | + aclrtFree(selfXDeviceAddr); | ||
| 198 | + aclrtFree(outDeviceAddr); | ||
| 199 | + if (workspaceSize > static_cast<uint64_t>(0)) { | ||
| 200 | + aclrtFree(workspaceAddr); | ||
| 201 | + } | ||
| 202 | + aclrtDestroyStream(stream); | ||
| 203 | + aclrtResetDevice(deviceId); | ||
| 204 | + | ||
| 205 | + // 10. acl 去初始化 | ||
| 206 | + aclFinalize(); | ||
| 207 | + | ||
| 208 | + return pass ? 0 : 1; | ||
| 209 | +} | ||
| @@ -0,0 +1,13 @@ | |||
| 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 | +# NOTE: Portions of this code were AI-generated and have been | ||
| 11 | +# technically reviewed for functional accuracy and security | ||
| 12 | + | ||
| 13 | +add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE relu6 ACLNNTYPE aclnn) | ||
陈 | |||
| @@ -0,0 +1,58 @@ | |||
| 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 | + * NOTE: Portions of this code were AI-generated and have been | ||
| 13 | + * technically reviewed for functional accuracy and security | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +/*! | ||
| 17 | + * \file relu6_def.cpp | ||
| 18 | + * \brief Relu6 算子定义，声明输入输出和算子配置 | ||
| 19 | + * | ||
| 20 | + * 算子公式: y = min(max(x, 0), 6) | ||
| 21 | + * 支持数据类型: float16, float, int32, bfloat16 | ||
| 22 | + * 目标芯片: Ascend950 (arch35) | ||
| 23 | + */ | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | +namespace ops { | ||
| 28 | +class Relu6 : public OpDef { | ||
| 29 | +public: | ||
| 30 | + explicit Relu6(const char* name) : OpDef(name) | ||
| 31 | + { | ||
| 32 | + this->Input("x") // 输入 x 定义 | ||
| 33 | + .ParamType(REQUIRED) // 必选输入 | ||
| 34 | + .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_BF16}) // 支持数据类型 | ||
| 35 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) // 支持 format 格式 | ||
| 36 | + .UnknownShapeFormat({ge::FORMAT_ND}) // 未确定大小 shape 对应 format | ||
| 37 | + .AutoContiguous(); // 内存自动连续化 | ||
| 38 | + this->Output("y") // 输出 y 定义 | ||
| 39 | + .ParamType(REQUIRED) | ||
| 40 | + .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_BF16}) | ||
| 41 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 42 | + .UnknownShapeFormat({ge::FORMAT_ND}) | ||
| 43 | + .AutoContiguous(); | ||
| 44 | + | ||
| 45 | + // Ascend950 AI Core 配置 | ||
| 46 | + OpAICoreConfig aiCoreConfig; | ||
| 47 | + aiCoreConfig.DynamicCompileStaticFlag(true) | ||
| 48 | + .DynamicFormatFlag(false) | ||
| 49 | + .DynamicRankSupportFlag(true) | ||
| 50 | + .DynamicShapeSupportFlag(true) | ||
| 51 | + .NeedCheckSupportFlag(false) | ||
| 52 | + .PrecisionReduceFlag(true) | ||
| 53 | + .ExtendCfgInfo("opFile.value", "relu6"); // 指定 Kernel 入口文件名 | ||
| 54 | + this->AICore().AddConfig("ascend950", aiCoreConfig); | ||
| 55 | + } | ||
| 56 | +}; | ||
| 57 | +OP_ADD(Relu6); // 注册算子到算子信息库 | ||
| 58 | +} // namespace ops | ||
| @@ -0,0 +1,175 @@ | |||
| 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 | + * NOTE: Portions of this code were AI-generated and have been | ||
| 13 | + * technically reviewed for functional accuracy and security | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +/*! | ||
| 17 | + * \file relu6_tiling.cpp | ||
| 18 | + * \brief Relu6 Tiling 实现（arch35 = Ascend950） | ||
| 19 | + * | ||
| 20 | + * Tiling 策略： | ||
| 21 | + * 1. 获取平台信息（UB 大小、AI Core 数量） | ||
| 22 | + * 2. 获取 shape/dtype 信息 | ||
| 23 | + * 3. 多核切分：blockFactor = ceil(totalNum / coreNum) | ||
| 24 | + * 4. UB 切分：ubFactor = floor_align(floor_div(ubSize / typeSize / 3), ubBlockSize) | ||
| 25 | + * Relu6 需要 3 个 LocalTensor: inputLocal, tmpLocal, outputLocal | ||
| 26 | + * 5. 设置 TilingKey（由模板参数选择 dtype） | ||
| 27 | + */ | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | +namespace optiling { | ||
| 37 | + | ||
| 38 | +using Ops::Base::CeilDiv; | ||
| 39 | +using Ops::Base::FloorAlign; | ||
| 40 | +using Ops::Base::FloorDiv; | ||
| 41 | +using Ops::Base::GetUbBlockSize; | ||
| 42 | + | ||
| 43 | +constexpr uint32_t WS_SYS_SIZE = 0U; | ||
| 44 | +// Relu6 需要 3 个 LocalTensor：inputLocal + tmpLocal（Maxs输出）+ outputLocal | ||
| 45 | +constexpr int64_t BUFFER_NUM = 3; | ||
| 46 | +// 双缓冲阈值：数据量大于此值时启用双缓冲（迭代三扩展） | ||
| 47 | +constexpr int64_t MIN_SPLIT_THRESHOLD = 1024; | ||
| 48 | + | ||
| 49 | +static const gert::Shape g_vec_1_shape = {1}; | ||
| 50 | + | ||
| 51 | +/// \brief 确保 shape 不为标量（0 维），若为标量则转换为 {1} | ||
| 52 | +static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape) | ||
| 53 | +{ | ||
| 54 | + if (in_shape.GetDimNum() == 0) { | ||
| 55 | + return g_vec_1_shape; | ||
| 56 | + } | ||
| 57 | + return in_shape; | ||
| 58 | +} | ||
| 59 | + | ||
| 60 | +/// \brief 获取平台信息：UB 大小、AI Core 数量 | ||
| 61 | +static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum) | ||
| 62 | +{ | ||
| 63 | + fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo(); | ||
| 64 | + OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr); | ||
| 65 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr); | ||
| 66 | + coreNum = ascendcPlatform.GetCoreNumAiv(); | ||
| 67 | + OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED); | ||
| 68 | + ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize); | ||
| 69 | + OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED); | ||
| 70 | + return ge::GRAPH_SUCCESS; | ||
| 71 | +} | ||
| 72 | + | ||
| 73 | +/// \brief 获取 shape 和 dtype 信息 | ||
| 74 | +static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t& totalNum, ge::DataType& dataType) | ||
| 75 | +{ | ||
| 76 | + auto inputX = context->GetInputShape(0); | ||
| 77 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputX); | ||
| 78 | + auto inputShapeX = EnsureNotScalar(inputX->GetStorageShape()); | ||
| 79 | + auto outY = context->GetOutputShape(0); | ||
| 80 | + OP_CHECK_NULL_WITH_CONTEXT(context, outY); | ||
| 81 | + auto outShapeY = EnsureNotScalar(outY->GetStorageShape()); | ||
| 82 | + | ||
| 83 | + // Shape 校验：输入输出元素数一致 | ||
| 84 | + OP_CHECK_IF( | ||
| 85 | + inputShapeX.GetShapeSize() != outShapeY.GetShapeSize(), | ||
| 86 | + OP_LOGE(context, "Relu6: input and output shape size mismatch"), return ge::GRAPH_FAILED); | ||
| 87 | + | ||
| 88 | + totalNum = inputShapeX.GetShapeSize(); | ||
| 89 | + | ||
| 90 | + // dtype 校验 | ||
| 91 | + const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_BF16}; | ||
| 92 | + auto inputDesc = context->GetInputDesc(0); | ||
| 93 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc); | ||
| 94 | + dataType = inputDesc->GetDataType(); | ||
| 95 | + if (supportedDtype.count(dataType) == 0) { | ||
| 96 | + OP_LOGE(context, "invalid dtype"); | ||
| 97 | + return ge::GRAPH_FAILED; | ||
| 98 | + } | ||
| 99 | + return ge::GRAPH_SUCCESS; | ||
| 100 | +} | ||
| 101 | + | ||
| 102 | +/// \brief 设置 workspace 大小（Relu6 不需要额外 workspace） | ||
| 103 | +static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context) | ||
| 104 | +{ | ||
| 105 | + size_t* currentWorkspace = context->GetWorkspaceSizes(1); | ||
| 106 | + OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace); | ||
| 107 | + currentWorkspace[0] = WS_SYS_SIZE; | ||
| 108 | + return ge::GRAPH_SUCCESS; | ||
| 109 | +} | ||
| 110 | + | ||
| 111 | +/// \brief Tiling 分发入口 | ||
| 112 | +static ge::graphStatus Relu6TilingFunc(gert::TilingContext* context) | ||
| 113 | +{ | ||
| 114 | + // 1. 获取平台运行信息 | ||
| 115 | + uint64_t ubSize; | ||
| 116 | + int64_t coreNum; | ||
| 117 | + OP_CHECK_IF( | ||
| 118 | + GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), | ||
| 119 | + return ge::GRAPH_FAILED); | ||
| 120 | + | ||
| 121 | + // 2. 获取 shape、dtype 信息 | ||
| 122 | + int64_t totalNum; | ||
| 123 | + ge::DataType dataType; | ||
| 124 | + OP_CHECK_IF( | ||
| 125 | + GetShapeAttrsInfo(context, totalNum, dataType) != ge::GRAPH_SUCCESS, | ||
| 126 | + OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED); | ||
| 127 | + | ||
| 128 | + // 3. 获取 WorkspaceSize | ||
| 129 | + OP_CHECK_IF( | ||
| 130 | + GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"), | ||
| 131 | + return ge::GRAPH_FAILED); | ||
| 132 | + | ||
| 133 | + // 4. 设置 TilingData | ||
| 134 | + Relu6TilingData* tiling = context->GetTilingData<Relu6TilingData>(); | ||
| 135 | + OP_CHECK_NULL_WITH_CONTEXT(context, tiling); | ||
| 136 | + OP_CHECK_IF( | ||
| 137 | + memset_s(tiling, sizeof(Relu6TilingData), 0, sizeof(Relu6TilingData)) != EOK, | ||
| 138 | + OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED); | ||
| 139 | + | ||
| 140 | + // 多核切分：将总元素按核数均分 | ||
| 141 | + tiling->totalNum = totalNum; | ||
| 142 | + tiling->blockFactor = CeilDiv(totalNum, coreNum); | ||
| 143 | + int64_t usedCoreNum = CeilDiv(totalNum, tiling->blockFactor); | ||
| 144 | + | ||
| 145 | + // UB 切分： | ||
| 146 | + // 公式：(UB总大小 / 类型大小) / buffer数量，然后按 UB 块大小对齐 | ||
| 147 | + // bufferNum = 3（inputLocal + tmpLocal + outputLocal） | ||
| 148 | + int64_t typeSize = ge::GetSizeByDataType(dataType); | ||
| 149 | + OP_CHECK_IF(typeSize == 0, OP_LOGE(context, "typeSize is 0, invalid dataType"), return ge::GRAPH_FAILED); | ||
| 150 | + int64_t ubBlockSize = GetUbBlockSize(context); | ||
| 151 | + tiling->ubFactor = FloorAlign(FloorDiv(static_cast<int64_t>(ubSize) / typeSize, BUFFER_NUM), ubBlockSize); | ||
| 152 | + | ||
| 153 | + // 设置 dataType 字段：保存原始 ge::DataType 枚举值，供 Kernel 侧模板参数实例化使用 | ||
| 154 | + tiling->dataType = static_cast<int32_t>(dataType); | ||
| 155 | + | ||
| 156 | + context->SetBlockDim(usedCoreNum); | ||
| 157 | + | ||
| 158 | + // 5. 设置 TilingKey（dtype 模板参数选择） | ||
| 159 | + uint32_t dType = static_cast<uint32_t>(dataType); | ||
| 160 | + ASCENDC_TPL_SEL_PARAM(context, dType); | ||
| 161 | + return ge::GRAPH_SUCCESS; | ||
| 162 | +} | ||
| 163 | + | ||
| 164 | +/// \brief Tiling 解析（Relu6 不需要额外解析） | ||
| 165 | +static ge::graphStatus TilingParseForRelu6([[maybe_unused]] gert::TilingParseContext* context) | ||
| 166 | +{ | ||
| 167 | + return ge::GRAPH_SUCCESS; | ||
| 168 | +} | ||
| 169 | + | ||
| 170 | +struct Relu6CompileInfo {}; // 必须定义，入图场景依赖 | ||
| 171 | + | ||
| 172 | +// Tiling 注册入口 | ||
| 173 | +IMPL_OP_OPTILING(Relu6).Tiling(Relu6TilingFunc).TilingParse<Relu6CompileInfo>(TilingParseForRelu6); | ||
| 174 | + | ||
| 175 | +} // namespace optiling | ||
| @@ -0,0 +1,38 @@ | |||
| 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 | + * NOTE: Portions of this code were AI-generated and have been | ||
| 13 | + * technically reviewed for functional accuracy and security | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +/*! | ||
| 17 | + * \file relu6.cpp | ||
| 18 | + * \brief Relu6 Kernel 入口（arch35 = Ascend950） | ||
| 19 | + * | ||
| 20 | + * 模板参数说明（与 relu6_tiling_key.h 中 ASCENDC_TPL_ARGS_DECL 定义对应）： | ||
| 21 | + * - D_T: 数据类型，由 ASCENDC_TPL_DATATYPE_DECL 定义 | ||
| 22 | + * | ||
| 23 | + * 迭代一：单核 float16 骨架，D_T 模板自动推导实例化 | ||
| 24 | + * 核函数参数顺序（固定）： | ||
| 25 | + * GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling | ||
| 26 | + */ | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | +template <typename D_T> | ||
| 31 | +__global__ __aicore__ void relu6(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) | ||
| 32 | +{ | ||
| 33 | + REGISTER_TILING_DEFAULT(Relu6TilingData); | ||
| 34 | + GET_TILING_DATA_WITH_STRUCT(Relu6TilingData, tilingData, tiling); | ||
| 35 | + NsRelu6::Relu6<D_T> op; | ||
| 36 | + op.Init(x, y, &tilingData); | ||
| 37 | + op.Process(); | ||
| 38 | +} | ||
| @@ -0,0 +1,190 @@ | |||
| 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 | + * NOTE: Portions of this code were AI-generated and have been | ||
| 13 | + * technically reviewed for functional accuracy and security | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +/*! | ||
| 17 | + * \file relu6.h | ||
| 18 | + * \brief Relu6 Kernel 类实现（arch35 = Ascend950） | ||
| 19 | + * | ||
| 20 | + * 公式: y = min(max(x, 0), 6) | ||
| 21 | + * 迭代一：单核 + 单 dtype（float16）骨架 | ||
| 22 | + * 数据流: GM -> UB(inputLocal) -> UB(tmpLocal=Maxs) -> UB(outputLocal=Mins) -> GM | ||
| 23 | + * UB 使用: 3 个 LocalTensor（inputLocal + tmpLocal + outputLocal） | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | +namespace NsRelu6 { | ||
| 35 | + | ||
| 36 | +using namespace AscendC; | ||
| 37 | + | ||
| 38 | +/// \brief Relu6 Kernel 类 | ||
| 39 | +/// \tparam T 数据类型（迭代一使用 half，迭代二扩展为模板参数） | ||
| 40 | +template <typename T> | ||
| 41 | +class Relu6 { | ||
| 42 | +public: | ||
| 43 | + __aicore__ inline Relu6() {} | ||
| 44 | + | ||
| 45 | + /// \brief 初始化 GlobalBuffer 和 UB Buffer | ||
| 46 | + /// \param x 输入 GM 地址 | ||
| 47 | + /// \param y 输出 GM 地址 | ||
| 48 | + /// \param tilingData Tiling 参数 | ||
| 49 | + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const Relu6TilingData* tilingData); | ||
| 50 | + | ||
| 51 | + /// \brief 执行主循环：CopyIn -> Compute -> CopyOut | ||
| 52 | + __aicore__ inline void Process(); | ||
| 53 | + | ||
| 54 | +private: | ||
| 55 | + /// \brief 从 GM 搬运数据到 UB | ||
| 56 | + /// \param progress 当前循环进度 | ||
| 57 | + /// \param currentNum 本次搬运元素数 | ||
| 58 | + __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum); | ||
| 59 | + | ||
| 60 | + /// \brief 从 UB 搬运结果到 GM | ||
| 61 | + /// \param progress 当前循环进度 | ||
| 62 | + /// \param currentNum 本次搬运元素数 | ||
| 63 | + __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum); | ||
| 64 | + | ||
| 65 | + /// \brief UB 内计算: max(x, 0) -> min(result, 6) | ||
| 66 | + /// \param currentNum 本次计算元素数 | ||
| 67 | + __aicore__ inline void Compute(int64_t currentNum); | ||
| 68 | + | ||
| 69 | +private: | ||
| 70 | + TPipe pipe; | ||
| 71 | + TQue<QuePosition::VECIN, 2> inputQueue; // 双缓冲：2 个 input buffer 支持乒乓调度 | ||
| 72 | + TQue<QuePosition::VECOUT, 1> outputQueue; | ||
| 73 | + | ||
| 74 | + GlobalTensor<T> inputGM; | ||
| 75 | + GlobalTensor<T> outputGM; | ||
| 76 | + | ||
| 77 | + int64_t blockLength_ = 0; // 当前 Core 需处理的元素数 | ||
| 78 | + int64_t ubLength_ = 0; // 单次 UB 循环处理的元素数 | ||
| 79 | +}; | ||
| 80 | + | ||
| 81 | +template <typename T> | ||
| 82 | +__aicore__ inline void Relu6<T>::Init(GM_ADDR x, GM_ADDR y, const Relu6TilingData* tilingData) | ||
| 83 | +{ | ||
| 84 | + // 计算当前 AI Core 需要处理的元素数量 | ||
| 85 | + int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx(); | ||
| 86 | + blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength; | ||
| 87 | + ubLength_ = tilingData->ubFactor; | ||
| 88 | + | ||
| 89 | + // 设置 Global Buffer，根据 blockIdx 偏移 | ||
| 90 | + inputGM.SetGlobalBuffer((__gm__ T*)x + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_); | ||
| 91 | + outputGM.SetGlobalBuffer((__gm__ T*)y + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_); | ||
| 92 | + | ||
| 93 | + // 初始化 UB Buffer：每个 LocalTensor 大小为 ubLength_ * sizeof(T) | ||
| 94 | + // 32 字节对齐由 AscendC AllocTensor 自动保证 | ||
| 95 | + // inputQueue: 2 个 buffer（双缓冲），支持 CopyIn 与 Compute/CopyOut 流水并行 | ||
| 96 | + // outputQueue: 动态管理 2 个 LocalTensor（tmpLocal + outputLocal） | ||
| 97 | + pipe.InitBuffer(inputQueue, 2, ubLength_ * sizeof(T)); | ||
| 98 | + pipe.InitBuffer(outputQueue, 2, ubLength_ * sizeof(T)); | ||
| 99 | +} | ||
| 100 | + | ||
| 101 | +template <typename T> | ||
| 102 | +__aicore__ inline void Relu6<T>::CopyIn(int64_t progress, int64_t currentNum) | ||
| 103 | +{ | ||
| 104 | + AscendC::LocalTensor<T> inputLocal = inputQueue.template AllocTensor<T>(); | ||
| 105 | + AscendC::DataCopyParams copyParams; | ||
| 106 | + copyParams.blockCount = 1; | ||
| 107 | + copyParams.blockLen = currentNum * sizeof(T); | ||
| 108 | + copyParams.srcStride = 0; | ||
| 109 | + copyParams.dstStride = 0; | ||
| 110 | + // DataCopyPad 自动处理末尾不足 32 字节对齐的尾部数据 | ||
| 111 | + AscendC::DataCopyPad(inputLocal, inputGM[progress * ubLength_], copyParams, {false, 0, 0, 0}); | ||
| 112 | + inputQueue.EnQue(inputLocal); | ||
| 113 | +} | ||
| 114 | + | ||
| 115 | +template <typename T> | ||
| 116 | +__aicore__ inline void Relu6<T>::CopyOut(int64_t progress, int64_t currentNum) | ||
| 117 | +{ | ||
| 118 | + AscendC::LocalTensor<T> outputLocal = outputQueue.template DeQue<T>(); | ||
| 119 | + AscendC::DataCopyParams copyParams; | ||
| 120 | + copyParams.blockCount = 1; | ||
| 121 | + copyParams.blockLen = currentNum * sizeof(T); | ||
| 122 | + copyParams.srcStride = 0; | ||
| 123 | + copyParams.dstStride = 0; | ||
| 124 | + AscendC::DataCopyPad(outputGM[progress * ubLength_], outputLocal, copyParams); | ||
| 125 | + outputQueue.FreeTensor(outputLocal); | ||
| 126 | +} | ||
| 127 | + | ||
| 128 | +template <typename T> | ||
| 129 | +__aicore__ inline void Relu6<T>::Compute(int64_t currentNum) | ||
| 130 | +{ | ||
| 131 | + AscendC::LocalTensor<T> inputLocal = inputQueue.template DeQue<T>(); | ||
| 132 | + AscendC::LocalTensor<T> outputLocal = outputQueue.template AllocTensor<T>(); | ||
| 133 | + | ||
| 134 | + // Step 1: tmpLocal = max(inputLocal, 0) | ||
| 135 | + // 使用 Maxs API 实现 Relu 的下界 | ||
| 136 | + AscendC::LocalTensor<T> tmpLocal = outputQueue.template AllocTensor<T>(); | ||
| 137 | + AscendC::Maxs(tmpLocal, inputLocal, static_cast<T>(0), currentNum); | ||
| 138 | + | ||
| 139 | + // Step 2: outputLocal = min(tmpLocal, 6) | ||
| 140 | + // 使用 Mins API 实现 Relu6 的上界 | ||
| 141 | + AscendC::Mins(outputLocal, tmpLocal, static_cast<T>(6), currentNum); | ||
| 142 | + | ||
| 143 | + // 释放 tmpLocal | ||
| 144 | + outputQueue.FreeTensor(tmpLocal); | ||
| 145 | + | ||
| 146 | + outputQueue.EnQue(outputLocal); | ||
| 147 | + inputQueue.FreeTensor(inputLocal); | ||
| 148 | +} | ||
| 149 | + | ||
| 150 | +template <typename T> | ||
| 151 | +__aicore__ inline void Relu6<T>::Process() | ||
| 152 | +{ | ||
| 153 | + int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_; | ||
| 154 | + if (loopCount == 0) { | ||
| 155 | + return; | ||
| 156 | + } | ||
| 157 | + | ||
| 158 | + // 单循环场景：直接串行执行，无需流水 | ||
| 159 | + if (loopCount == 1) { | ||
| 160 | + int64_t currentNum = blockLength_; | ||
| 161 | + CopyIn(0, currentNum); | ||
| 162 | + Compute(currentNum); | ||
| 163 | + CopyOut(0, currentNum); | ||
| 164 | + return; | ||
| 165 | + } | ||
| 166 | + | ||
| 167 | + // 多循环场景：双缓冲流水调度 | ||
| 168 | + // 第 0 轮：仅 CopyIn，启动输入搬运 | ||
| 169 | + int64_t currentNum = ubLength_; | ||
| 170 | + CopyIn(0, currentNum); | ||
| 171 | + | ||
| 172 | + for (int64_t i = 0; i < loopCount; i++) { | ||
| 173 | + // 计算当前轮的元素数 | ||
| 174 | + currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_; | ||
| 175 | + | ||
| 176 | + // Compute + CopyOut：处理当前轮数据 | ||
| 177 | + Compute(currentNum); | ||
| 178 | + CopyOut(i, currentNum); | ||
| 179 | + | ||
| 180 | + // CopyIn：预取下一轮数据（与当前轮的 Compute/CopyOut 流水并行） | ||
| 181 | + if (i < loopCount - 1) { | ||
| 182 | + int64_t nextNum = ubLength_; | ||
| 183 | + CopyIn(i + 1, nextNum); | ||
| 184 | + } | ||
| 185 | + } | ||
| 186 | +} | ||
| 187 | + | ||
| 188 | +} // namespace NsRelu6 | ||
| 189 | + | ||
| 190 | + | ||
| @@ -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 | + * NOTE: Portions of this code were AI-generated and have been | ||
| 13 | + * technically reviewed for functional accuracy and security | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +/*! | ||
| 17 | + * \file relu6_tiling_data.h | ||
| 18 | + * \brief Relu6 TilingData 结构体定义（arch35 = Ascend950） | ||
| 19 | + * | ||
| 20 | + * 迭代一：仅 float16 单 dtype 骨架 | ||
| 21 | + * 迭代二：新增 dataType 字段，支持多 dtype 分发 | ||
| 22 | + * 字段说明： | ||
| 23 | + * - totalNum: 输入张量展平后的总元素数量 | ||
| 24 | + * - blockFactor: 每个 AI Core 处理的元素数量 | ||
| 25 | + * - ubFactor: UB 单次循环处理的元素数量 | ||
| 26 | + * - dataType: 数据类型标识（0=float16, 1=float, 2=int32, 3=bfloat16） | ||
| 27 | + */ | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | +struct Relu6TilingData { | ||
| 35 | + int64_t totalNum = 0; // 总元素数量 | ||
| 36 | + int64_t blockFactor = 0; // 每个核处理的元素数量 | ||
| 37 | + int64_t ubFactor = 0; // 每次 UB 循环处理的元素数量 | ||
| 38 | + int32_t dataType = 0; // 数据类型：0=float16, 1=float, 2=int32, 3=bfloat16 | ||
| 39 | +}; | ||
| 40 | + | ||
| 41 | + | ||
| @@ -0,0 +1,51 @@ | |||
| 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 | + * NOTE: Portions of this code were AI-generated and have been | ||
| 13 | + * technically reviewed for functional accuracy and security | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | +/*! | ||
| 17 | + * \file relu6_tiling_key.h | ||
| 18 | + * \brief Relu6 TilingKey 模板参数定义（arch35 = Ascend950） | ||
| 19 | + * | ||
| 20 | + * 迭代一：仅 float16 单 dtype 骨架，预留 float/int32/bfloat16 扩展位置 | ||
| 21 | + * 模板参数类型参考： | ||
| 22 | + * - DATATYPE: 原生数据类型（C_DT_FLOAT16, C_DT_FLOAT, C_DT_INT32, C_DT_BF16） | ||
| 23 | + * 参考：ascendc/host_api/tiling/template_argument.h | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | +// 迭代一：仅 float16；迭代二扩展全部 4 种 dtype | ||
| 32 | +ASCENDC_TPL_ARGS_DECL(Relu6, | ||
| 33 | + ASCENDC_TPL_DATATYPE_DECL(D_T, C_DT_FLOAT16, C_DT_FLOAT, C_DT_INT32, C_DT_BF16, ASCENDC_TPL_INPUT(0)) | ||
| 34 | +); | ||
| 35 | + | ||
| 36 | +ASCENDC_TPL_SEL( | ||
| 37 | + ASCENDC_TPL_ARGS_SEL( | ||
| 38 | + ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_FLOAT16) | ||
| 39 | + ), | ||
| 40 | + ASCENDC_TPL_ARGS_SEL( | ||
| 41 | + ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_FLOAT) | ||
| 42 | + ), | ||
| 43 | + ASCENDC_TPL_ARGS_SEL( | ||
| 44 | + ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_INT32) | ||
| 45 | + ), | ||
| 46 | + ASCENDC_TPL_ARGS_SEL( | ||
| 47 | + ASCENDC_TPL_DATATYPE_SEL(D_T, C_DT_BF16) | ||
| 48 | + ), | ||
| 49 | +); | ||
| 50 | + | ||
| 51 | + | ||
The file is empty
CMakeLists 缺少架构和 tiling 目录参数。虽然只支持 ascend950，但 examples 在 arch35 目录下，建议添加
COMPUTE_UNIT和TILING_DIR参数以保持配置一致性。