已合并
【社区任务】LogicalNot算子开发--算子提交 #1993
lyzin创建于 3月30日
【社区任务】LogicalNot算子开发--算子提交 #1993
已合并
共 15 个文件变更+1131-0
| @@ -0,0 +1,11 @@ | |||
| 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 | +add_all_modules_sources(OPTYPE logical_not ACLNNTYPE aclnn_exclude) | ||
| @@ -0,0 +1,97 @@ | |||
| 1 | +# LogicalNot | ||
| 2 | + | ||
| 3 | +## 贡献说明 | ||
| 4 | + | ||
| 5 | +| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 | | ||
| 6 | +|----|----|----|------|------| | ||
| 7 | +| 周世星 | 浙江工业大学-智能计算研究所 | LogicalNot | 2025/06/24 | 新增LogicalNot算子，实现了逻辑否功能。 | | ||
| 8 | + | ||
| 9 | +## 支持的产品型号 | ||
| 10 | + | ||
| 11 | +- Atlas A2训练系列产品 | ||
| 12 | + | ||
| 13 | +产品形态详细说明请参见[昇腾产品形态说明](https://www.hiascend.com/document/redirect/CannCommunityProductForm)。 | ||
| 14 | + | ||
| 15 | +## 算子描述 | ||
| 16 | + | ||
| 17 | +- 功能描述 | ||
| 18 | + | ||
| 19 | + `LogicalNot` 算子对输入的布尔值进行逻辑非运算（取反操作）。 | ||
| 20 | + | ||
| 21 | +- 计算公式： | ||
| 22 | + | ||
| 23 | + $$ | ||
| 24 | + y = \neg x | ||
| 25 | + $$ | ||
| 26 | + | ||
| 27 | +- 原型信息 | ||
| 28 | + | ||
| 29 | + <table> | ||
| 30 | + <tr> | ||
| 31 | + <th align="center">算子类型(OpType)</th> | ||
| 32 | + <th colspan="4" align="center">LogicalNot</th> | ||
| 33 | + </tr> | ||
| 34 | + <tr> | ||
| 35 | + <td align="center"></td> | ||
| 36 | + <td align="center">name</td> | ||
| 37 | + <td align="center">Type</td> | ||
| 38 | + <td align="center">data type</td> | ||
| 39 | + <td align="center">format</td> | ||
| 40 | + </tr> | ||
| 41 | + <tr> | ||
| 42 | + <td rowspan="2" align="center">算子输入</td> | ||
| 43 | + </tr> | ||
| 44 | + <tr> | ||
| 45 | + <td align="center">x</td> | ||
| 46 | + <td align="center">tensor</td> | ||
| 47 | + <td align="center">bool</td> | ||
| 48 | + <td align="center">ND</td> | ||
| 49 | + </tr> | ||
| 50 | + <tr> | ||
| 51 | + <td rowspan="1" align="center">算子输出</td> | ||
| 52 | + <td align="center">y</td> | ||
| 53 | + <td align="center">tensor</td> | ||
| 54 | + <td align="center">bool</td> | ||
| 55 | + <td align="center">ND</td> | ||
| 56 | + </tr> | ||
| 57 | + <tr> | ||
| 58 | + <td rowspan="1" align="center">核函数名</td> | ||
| 59 | + <td colspan="4" align="center">logical_not</td> | ||
| 60 | + </tr> | ||
| 61 | + </table> | ||
| 62 | + | ||
| 63 | + | ||
| 64 | +## 约束与限制 | ||
| 65 | + | ||
| 66 | +- 无 | ||
| 67 | + | ||
| 68 | +## 算子使用 | ||
F | |||
| 69 | + | ||
| 70 | +编译运行此样例前，请参考[《CANN软件安装指南》](https://hiascend.com/document/redirect/CannCommunityInstSoftware)完成开发运行环境的部署。 | ||
| 71 | + | ||
| 72 | +### 编译部署 | ||
| 73 | + | ||
| 74 | + - 进入到仓库目录 | ||
| 75 | + | ||
| 76 | + ```bash | ||
| 77 | + cd ${git_clone_path}/ops-math | ||
| 78 | + ``` | ||
| 79 | + | ||
| 80 | + - 执行编译 | ||
| 81 | + | ||
| 82 | + ```bash | ||
| 83 | + bash build.sh --pkg --experimental --soc=ascend910b --ops=logical_not | ||
| 84 | + ``` | ||
| 85 | + | ||
| 86 | + - 部署算子包 | ||
| 87 | + | ||
| 88 | + ```bash | ||
| 89 | + ./build_out/cann-ops-<vendor_name>-linux.<arch>.run | ||
| 90 | + ``` | ||
| 91 | + | ||
| 92 | +### 算子调用 | ||
| 93 | + | ||
| 94 | + - 执行调用 | ||
| 95 | + | ||
| 96 | + ```bash | ||
| 97 | + bash build.sh --run_example --experimental logical_not eager cust --vendor_name=custom | ||
| @@ -0,0 +1,109 @@ | |||
| 1 | +# aclnnLogicalNot | ||
| 2 | + | ||
| 3 | +## 支持的产品型号 | ||
| 4 | + | ||
| 5 | +| 产品 | 是否支持 | | ||
| 6 | +| :------------------------------------------------------------------------------ | :------: | | ||
| 7 | +| <term>Atlas A2 训练系列产品</term> | √ | | ||
| 8 | + | ||
| 9 | +## 功能描述 | ||
| 10 | + | ||
| 11 | +- 算子功能：LogicalNot算子提供逻辑非运算功能，对输入的布尔值进行取反操作。 | ||
| 12 | +- 计算公式： | ||
| 13 | + | ||
| 14 | + $$ | ||
| 15 | + y = \neg x | ||
| 16 | + $$ | ||
| 17 | + | ||
| 18 | +## 实现原理 | ||
| 19 | + | ||
| 20 | +输入的`bool`类型数据在kernel侧以`int8`进行处理，通过调用`Ascend C`的`Cast`函数将输入的`int8`数据转换为`float16`后进行计算，最后通过`Cast`函数将`float16`数据转换回`int8`，实现对输入的布尔值进行取反操作。 | ||
| 21 | + | ||
| 22 | +## 算子执行接口 | ||
| 23 | + | ||
| 24 | +每个算子分为两段式接口，必须先调用“aclnnLogicalNotGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnLogicalNot”接口执行计算。 | ||
| 25 | + | ||
| 26 | +* `aclnnStatus aclnnLogicalNotGetWorkspaceSize(const aclTensor* x, const aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)` | ||
| 27 | +* `aclnnStatus aclnnLogicalNot(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)` | ||
| 28 | + | ||
| 29 | +**说明**： | ||
| 30 | + | ||
| 31 | +- 算子执行接口对外屏蔽了算子内部实现逻辑以及不同代际NPU的差异，且开发者无需编译算子，实现了算子的精简调用。 | ||
| 32 | +- 若开发者不使用算子执行接口的调用算子，也可以定义基于Ascend IR的算子描述文件，通过ATC工具编译获得算子om文件，然后加载模型文件执行算子。 | ||
| 33 | + | ||
| 34 | +### aclnnLogicalNotGetWorkspaceSize | ||
| 35 | + | ||
| 36 | +- **参数说明：** | ||
| 37 | + | ||
| 38 | + - x（aclTensor\*，计算输入）：必选参数，Device侧的aclTensor，公式中的输入x，数据类型支持BOOL，数据格式支持ND。 | ||
| 39 | + - out（aclTensor\*，计算输出）：Device侧的aclTensor，公式中的输出y，数据类型支持BOOL，数据格式支持ND，输出维度与x一致。 | ||
| 40 | + - workspaceSize（uint64\_t\*，出参）：返回用户需要在Device侧申请的workspace大小。 | ||
| 41 | + - executor（aclOpExecutor\*\*，出参）：返回op执行器，包含了算子计算流程。 | ||
| 42 | + | ||
| 43 | +- **返回值：** | ||
| 44 | + | ||
| 45 | + 返回aclnnStatus状态码。 | ||
| 46 | + | ||
| 47 | + ``` | ||
| 48 | + 第一段接口完成入参校验，若出现以下错误码，则对应原因为： | ||
| 49 | + - 返回161001（ACLNN_ERR_PARAM_NULLPTR）：如果传入参数是必选输入，输出或者必选属性，且是空指针，则返回161001。 | ||
| 50 | + - 返回161002（ACLNN_ERR_PARAM_INVALID）：x、out的数据类型和数据格式不在支持的范围内。 | ||
| 51 | + ``` | ||
| 52 | + | ||
| 53 | +### aclnnLogicalNot | ||
| 54 | + | ||
| 55 | +- **参数说明：** | ||
| 56 | + | ||
| 57 | + - workspace（void\*，入参）：在Device侧申请的workspace内存起址。 | ||
| 58 | + - workspaceSize（uint64\_t，入参）：在Device侧申请的workspace大小，由第一段接口aclnnLogicalNotGetWorkspaceSize获取。 | ||
| 59 | + - executor（aclOpExecutor\*，入参）：op执行器，包含了算子计算流程。 | ||
| 60 | + - stream（aclrtStream，入参）：指定执行任务的AscendCL stream流。 | ||
| 61 | +- **返回值：** | ||
| 62 | + | ||
| 63 | + 返回aclnnStatus状态码。 | ||
| 64 | + | ||
| 65 | + | ||
| 66 | +## 约束与限制 | ||
| 67 | + | ||
| 68 | +- x，out的数据类型只支持BOOL，数据格式只支持ND | ||
| 69 | + | ||
| 70 | +## 算子原型 | ||
| 71 | + | ||
| 72 | +<table> | ||
| 73 | + <tr> | ||
| 74 | + <th align="center">算子类型(OpType)</th> | ||
| 75 | + <th colspan="4" align="center">LogicalNot</th> | ||
| 76 | + </tr> | ||
| 77 | + <tr> | ||
| 78 | + <td align="center"></td> | ||
| 79 | + <td align="center">name</td> | ||
| 80 | + <td align="center">type</td> | ||
| 81 | + <td align="center">data type</td> | ||
| 82 | + <td align="center">format</td> | ||
| 83 | + </tr> | ||
| 84 | + <tr> | ||
| 85 | + <td rowspan="2" align="center">算子输入</td> | ||
| 86 | + </tr> | ||
| 87 | + <tr> | ||
| 88 | + <td align="center">x</td> | ||
| 89 | + <td align="center">tensor</td> | ||
| 90 | + <td align="center">bool</td> | ||
| 91 | + <td align="center">ND</td> | ||
| 92 | + </tr> | ||
| 93 | + <tr> | ||
| 94 | + <td rowspan="1" align="center">算子输出</td> | ||
| 95 | + <td align="center">y</td> | ||
| 96 | + <td align="center">tensor</td> | ||
| 97 | + <td align="center">bool</td> | ||
| 98 | + <td align="center">ND</td> | ||
| 99 | + </tr> | ||
| 100 | + <tr> | ||
| 101 | + <td rowspan="1" align="center">核函数名</td> | ||
| 102 | + <td colspan="4" align="center">logical_not</td> | ||
| 103 | + </tr> | ||
| 104 | +</table> | ||
| 105 | + | ||
| 106 | + | ||
| 107 | +## 调用示例 | ||
| 108 | + | ||
| 109 | +详见[test_aclnn_logical_not.cpp](../examples/test_aclnn_logical_not.cpp) | ||
| @@ -0,0 +1,171 @@ | |||
| 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 test_aclnn_logical_not.cpp | ||
| 13 | + * \brief | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + do { \ | ||
| 22 | + if (!(cond)) { \ | ||
| 23 | + return_expr; \ | ||
| 24 | + } \ | ||
| 25 | + } while (0) | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + do { \ | ||
| 29 | + printf(message, ##__VA_ARGS__); \ | ||
| 30 | + } while (0) | ||
| 31 | + | ||
| 32 | +int64_t GetShapeSize(const std::vector<int64_t>& shape) { | ||
| 33 | + int64_t shape_size = 1; | ||
| 34 | + for (auto i : shape) { | ||
| 35 | + shape_size *= i; | ||
| 36 | + } | ||
| 37 | + return shape_size; | ||
| 38 | +} | ||
| 39 | + | ||
| 40 | +int Init(int32_t deviceId, aclrtStream* stream) { | ||
| 41 | + // 固定写法，资源初始化 | ||
| 42 | + auto ret = aclInit(nullptr); | ||
| 43 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 44 | + ret = aclrtSetDevice(deviceId); | ||
| 45 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 46 | + ret = aclrtCreateStream(stream); | ||
| 47 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 48 | + return 0; | ||
| 49 | +} | ||
| 50 | + | ||
| 51 | +template <typename T> | ||
| 52 | +int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, | ||
| 53 | + aclDataType dataType, aclTensor** tensor) { | ||
| 54 | + auto size = GetShapeSize(shape) * sizeof(T); | ||
| 55 | + // 调用aclrtMalloc申请device侧内存 | ||
| 56 | + auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 57 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 58 | + | ||
| 59 | + // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 | ||
| 60 | + ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 61 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); | ||
| 62 | + | ||
| 63 | + // 计算连续tensor的strides | ||
| 64 | + std::vector<int64_t> strides(shape.size(), 1); | ||
| 65 | + for (int64_t i = shape.size() - 2; i >= 0; i--) { | ||
| 66 | + strides[i] = shape[i + 1] * strides[i + 1]; | ||
| 67 | + } | ||
| 68 | + | ||
| 69 | + // 调用aclCreateTensor接口创建aclTensor | ||
| 70 | + *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, | ||
| 71 | + shape.data(), shape.size(), *deviceAddr); | ||
| 72 | + return 0; | ||
| 73 | +} | ||
| 74 | + | ||
| 75 | +int main() { | ||
| 76 | + // 1. 固定写法，device/stream初始化, 参考acl API手册 | ||
| 77 | + // 根据自己的实际device填写deviceId | ||
| 78 | + int32_t deviceId = 0; | ||
| 79 | + aclrtStream stream; | ||
| 80 | + auto ret = Init(deviceId, &stream); | ||
| 81 | + // check根据自己的需要处理 | ||
| 82 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); | ||
| 83 | + // 2. 构造输入与输出，需要根据API的接口定义构造 | ||
| 84 | + std::vector<int64_t> selfShape = {3, 3}; | ||
| 85 | + std::vector<int64_t> outShape = {3, 3}; | ||
| 86 | + | ||
| 87 | + void* selfDeviceAddr = nullptr; | ||
| 88 | + void* outDeviceAddr = nullptr; | ||
| 89 | + aclTensor* self = nullptr; | ||
| 90 | + aclTensor* out = nullptr; | ||
| 91 | + | ||
| 92 | + std::vector<int> selfHostData = {0, 1, 0, 2, 0, 3, 4, 0, 0}; | ||
| 93 | + std::vector<int> outHostData ={0, 0, 0, 0, 0, 0, 0, 0, 0}; | ||
| 94 | + | ||
| 95 | + // 创建self aclTensor | ||
| 96 | + ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_INT32, &self); | ||
| 97 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 98 | + // 创建out aclTensor | ||
| 99 | + ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_INT32, &out); | ||
| 100 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 101 | + | ||
| 102 | + uint64_t workspaceSize = 0; | ||
| 103 | + aclOpExecutor* executor; | ||
| 104 | + | ||
| 105 | + // aclnnLogicalNot接口调用示例 | ||
| 106 | + // 3. 调用aclnnLogicalNot第一段接口 | ||
| 107 | + ret = aclnnLogicalNotGetWorkspaceSize(self, out, &workspaceSize, &executor); | ||
| 108 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnLogicalNotGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 109 | + // 根据第一段接口计算出的workspaceSize申请device内存 | ||
| 110 | + void* workspaceAddr = nullptr; | ||
| 111 | + if (workspaceSize > 0) { | ||
| 112 | + ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 113 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret;); | ||
| 114 | + } | ||
| 115 | + // 调用aclnnLogicalNot第二段接口 | ||
| 116 | + ret = aclnnLogicalNot(workspaceAddr, workspaceSize, executor, stream); | ||
| 117 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnLogicalNot failed. ERROR: %d\n", ret); return ret); | ||
| 118 | + // 4. 固定写法，同步等待任务执行结束 | ||
| 119 | + ret = aclrtSynchronizeStream(stream); | ||
| 120 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 121 | + // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改 | ||
| 122 | + auto size = GetShapeSize(outShape); | ||
| 123 | + std::vector<int> resultData(size, 0); | ||
| 124 | + ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, | ||
| 125 | + size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 126 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); | ||
| 127 | + | ||
| 128 | + for (int64_t i = 0; i < size; i++) { | ||
| 129 | + LOG_PRINT("result[%ld] is: %d\n", i, resultData[i]); | ||
| 130 | + } | ||
| 131 | + | ||
| 132 | + // aclnnInplaceLogicalNot接口调用示例 | ||
| 133 | + // step3. 调用aclnnInplaceLogicalNot第一段接口 | ||
| 134 | + ret = aclnnInplaceLogicalNotGetWorkspaceSize(self, &workspaceSize, &executor); | ||
| 135 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceLogicalNotGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 136 | + // 根据第一段接口计算出的workspaceSize申请device内存 | ||
| 137 | + if (workspaceSize > 0) { | ||
| 138 | + ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 139 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret;); | ||
| 140 | + } | ||
| 141 | + // 调用aclnnInplaceLogicalNot第二段接口 | ||
| 142 | + ret = aclnnInplaceLogicalNot(workspaceAddr, workspaceSize, executor, stream); | ||
| 143 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnInplaceLogicalNot failed. ERROR: %d\n", ret); return ret); | ||
| 144 | + // step4. 固定写法，同步等待任务执行结束 | ||
| 145 | + ret = aclrtSynchronizeStream(stream); | ||
| 146 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 147 | + // step5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改 | ||
| 148 | + ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), selfDeviceAddr, | ||
| 149 | + size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 150 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); | ||
| 151 | + | ||
| 152 | + for (int64_t i = 0; i < size; i++) { | ||
| 153 | + LOG_PRINT("result[%ld] is: %d\n", i, resultData[i]); | ||
| 154 | + } | ||
| 155 | + | ||
| 156 | + // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改 | ||
| 157 | + aclDestroyTensor(self); | ||
| 158 | + aclDestroyTensor(out); | ||
| 159 | + | ||
| 160 | + // 7. 释放device 资源 | ||
| 161 | + aclrtFree(selfDeviceAddr); | ||
| 162 | + aclrtFree(outDeviceAddr); | ||
| 163 | + if (workspaceSize > 0) { | ||
| 164 | + aclrtFree(workspaceAddr); | ||
| 165 | + } | ||
| 166 | + aclrtDestroyStream(stream); | ||
| 167 | + aclrtResetDevice(deviceId); | ||
| 168 | + aclFinalize(); | ||
| 169 | + | ||
| 170 | + return 0; | ||
| 171 | +} | ||
| @@ -0,0 +1,157 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | +using namespace op; | ||
| 28 | + | ||
| 29 | +static const std::initializer_list<DataType> dtype_support_list = { | ||
| 30 | + op::DataType::DT_UINT8, op::DataType::DT_INT8, op::DataType::DT_INT16, | ||
| 31 | + op::DataType::DT_INT32, op::DataType::DT_INT64, op::DataType::DT_FLOAT16, | ||
| 32 | + op::DataType::DT_FLOAT, op::DataType::DT_DOUBLE, op::DataType::DT_BOOL, op::DataType::DT_BF16}; | ||
| 33 | + | ||
| 34 | + | ||
| 35 | +static bool CheckNotNull(const aclTensor* self, const aclTensor* out) { | ||
| 36 | + OP_CHECK_NULL(self, return false); | ||
| 37 | + OP_CHECK_NULL(out, return false); | ||
| 38 | + | ||
| 39 | + return true; | ||
| 40 | +} | ||
| 41 | + | ||
| 42 | +static bool CheckDtypeValid(const aclTensor* self, const aclTensor* out) { | ||
| 43 | + // 检查self的数据类型是否在支持列表内 | ||
| 44 | + OP_CHECK_DTYPE_NOT_SUPPORT(self, dtype_support_list, return false); | ||
| 45 | + | ||
| 46 | + // 检查out的数据类型是否在支持列表内 | ||
| 47 | + OP_CHECK_DTYPE_NOT_SUPPORT(out, dtype_support_list, return false); | ||
| 48 | + | ||
| 49 | + return true; | ||
| 50 | +} | ||
| 51 | + | ||
| 52 | +static bool CheckShape(const aclTensor* self, const aclTensor* out) { | ||
| 53 | + // self和out的shape必须一致 | ||
| 54 | + OP_CHECK_SHAPE_NOT_EQUAL(self, out, return false); | ||
| 55 | + | ||
| 56 | + // 输入输出连续，不限制维度数 | ||
| 57 | + if (IsContiguous(self) && IsContiguous(out)) { | ||
| 58 | + return true; | ||
| 59 | + } | ||
| 60 | + | ||
| 61 | + OP_CHECK_MAX_DIM(self, MAX_SUPPORT_DIMS_NUMS, return false); | ||
| 62 | + OP_CHECK_MAX_DIM(out, MAX_SUPPORT_DIMS_NUMS, return false); | ||
| 63 | + | ||
| 64 | + return true; | ||
| 65 | +} | ||
| 66 | + | ||
| 67 | +static aclnnStatus CheckParams(const aclTensor* self, const aclTensor* out) { | ||
| 68 | + // 检查参数是否为空指针 | ||
| 69 | + CHECK_RET(CheckNotNull(self, out), ACLNN_ERR_PARAM_NULLPTR); | ||
| 70 | + | ||
| 71 | + // 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验 | ||
| 72 | + CHECK_RET(CheckDtypeValid(self, out), ACLNN_ERR_PARAM_INVALID); | ||
| 73 | + | ||
| 74 | + // 检查shape是否满足约束 | ||
| 75 | + CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID); | ||
| 76 | + return ACLNN_SUCCESS; | ||
| 77 | +} | ||
| 78 | + | ||
| 79 | +static aclnnStatus CalculateResult(const aclTensor* self, aclTensor* out, aclOpExecutor* executor) { | ||
| 80 | + // 固定写法，参数检查 | ||
| 81 | + auto ret = CheckParams(self, out); | ||
| 82 | + CHECK_RET(ret == ACLNN_SUCCESS, ret); | ||
| 83 | + | ||
| 84 | + // 空Tensor处理 | ||
| 85 | + if (self->IsEmpty()) { | ||
| 86 | + return ACLNN_SUCCESS; | ||
| 87 | + } | ||
| 88 | + | ||
| 89 | + // self如果非连续，需要转连续 | ||
| 90 | + auto selfContiguous = l0op::Contiguous(self, executor); | ||
| 91 | + CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR); | ||
| 92 | + | ||
| 93 | + // 将输入self的数据类型转换成隐式数据类型，根据具体算子语义按需调用 | ||
| 94 | + auto selfCasted = l0op::Cast(selfContiguous, DataType::DT_BOOL, executor); | ||
| 95 | + CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR); | ||
| 96 | + | ||
| 97 | + // 调用l0算子LogicalNot行计算 | ||
| 98 | + auto result = l0op::LogicalNot(selfCasted, executor); | ||
| 99 | + CHECK_RET(result != nullptr, ACLNN_ERR_INNER_NULLPTR); | ||
| 100 | + | ||
| 101 | + auto resultCasted = l0op::Cast(result, out->GetDataType(), executor); | ||
| 102 | + CHECK_RET(resultCasted != nullptr, ACLNN_ERR_INNER_NULLPTR); | ||
| 103 | + | ||
| 104 | + // 如果出参out是非连续Tensor，需要把计算完的连续Tensor转非连续 | ||
| 105 | + auto viewCopyResult = l0op::ViewCopy(resultCasted, out, executor); | ||
| 106 | + CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR); | ||
| 107 | + | ||
| 108 | + return ACLNN_SUCCESS; | ||
| 109 | +} | ||
| 110 | + | ||
| 111 | +aclnnStatus aclnnLogicalNotGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize, | ||
| 112 | + aclOpExecutor** executor) { | ||
| 113 | + L2_DFX_PHASE_1(aclnnLogicalNot, DFX_IN(self), DFX_OUT(out)); | ||
| 114 | + | ||
| 115 | + // 固定写法，创建OpExecutor | ||
| 116 | + auto uniqueExecutor = CREATE_EXECUTOR(); | ||
| 117 | + CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR); | ||
| 118 | + | ||
| 119 | + auto ret = CalculateResult(self, out, uniqueExecutor.get()); | ||
| 120 | + CHECK_RET(ret == ACLNN_SUCCESS, ret); | ||
| 121 | + | ||
| 122 | + // 固定写法，获取计算过程中需要使用的workspace大小 | ||
| 123 | + *workspaceSize = uniqueExecutor->GetWorkspaceSize(); | ||
| 124 | + uniqueExecutor.ReleaseTo(executor); | ||
| 125 | + return ACLNN_SUCCESS; | ||
| 126 | +} | ||
| 127 | + | ||
| 128 | +aclnnStatus aclnnInplaceLogicalNotGetWorkspaceSize(aclTensor* selfRef, uint64_t* workspaceSize, | ||
| 129 | + aclOpExecutor** executor) { | ||
| 130 | + L2_DFX_PHASE_1(aclnnInplaceLogicalNot, DFX_IN(selfRef), DFX_OUT(selfRef)); | ||
| 131 | + | ||
| 132 | + // 固定写法，创建OpExecutor | ||
| 133 | + auto uniqueExecutor = CREATE_EXECUTOR(); | ||
| 134 | + CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR); | ||
| 135 | + | ||
| 136 | + auto ret = CalculateResult(selfRef, selfRef, uniqueExecutor.get()); | ||
| 137 | + CHECK_RET(ret == ACLNN_SUCCESS, ret); | ||
| 138 | + | ||
| 139 | + // 固定写法，获取计算过程中需要使用的workspace大小 | ||
| 140 | + *workspaceSize = uniqueExecutor->GetWorkspaceSize(); | ||
| 141 | + uniqueExecutor.ReleaseTo(executor); | ||
| 142 | + return ACLNN_SUCCESS; | ||
| 143 | +} | ||
| 144 | + | ||
| 145 | +aclnnStatus aclnnLogicalNot(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, | ||
| 146 | + aclrtStream stream) { | ||
| 147 | + L2_DFX_PHASE_2(aclnnLogicalNot); | ||
| 148 | + // 固定写法，调用框架能力，完成计算 | ||
| 149 | + return CommonOpExecutorRun(workspace, workspaceSize, executor, stream); | ||
| 150 | +} | ||
| 151 | + | ||
| 152 | +aclnnStatus aclnnInplaceLogicalNot(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, | ||
| 153 | + aclrtStream stream) { | ||
| 154 | + L2_DFX_PHASE_2(aclnnInplaceLogicalNot); | ||
| 155 | + // 固定写法，调用框架能力，完成计算 | ||
| 156 | + return CommonOpExecutorRun(workspace, workspaceSize, executor, stream); | ||
| 157 | +} | ||
| @@ -0,0 +1,94 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | +extern "C" { | ||
| 19 | + | ||
| 20 | + | ||
| 21 | +/** | ||
| 22 | + * @brief aclnnLogicalNot的第一段接口，根据具体的计算流程，计算workspace大小。 | ||
| 23 | + * @domain aclnn_math | ||
| 24 | + * | ||
| 25 | + * 算子功能：完成逻辑取反计算 | ||
| 26 | + * | ||
| 27 | + * 实现说明： | ||
| 28 | + * api计算的基本路径： | ||
| 29 | + * ```mermaid | ||
| 30 | + * graph LR | ||
| 31 | + * A[(self)] -->B([l0op::Contiguous]) | ||
| 32 | + * B --> L([l0op::Cast]) | ||
| 33 | + * L --> E([l0op::LogicalNot]) | ||
| 34 | + * E --> E1([l0op::Cast]) | ||
| 35 | + * E1 --> G([l0op::ViewCopy]) | ||
| 36 | + * G --> H[(out)] | ||
| 37 | + * ``` | ||
| 38 | + * | ||
| 39 | + * @param [in] self: npu | ||
| 40 | + * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL， | ||
| 41 | + * 支持非连续的Tensor，数据格式支持ND。 | ||
| 42 | + * @param [in] out: npu | ||
| 43 | + * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL， | ||
| 44 | + * shape与self一致，数据格式支持ND，且数据格式需要与self一致。 | ||
| 45 | + * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。 | ||
| 46 | + * @param [out] executor: 返回op执行器，包含算子计算流程。 | ||
| 47 | + * @return aclnnStatus: 返回状态码。 | ||
| 48 | + */ | ||
| 49 | +ACLNN_API aclnnStatus aclnnLogicalNotGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize, | ||
| 50 | + aclOpExecutor** executor); | ||
| 51 | + | ||
| 52 | +/** | ||
| 53 | + * @brief aclnnLogicalNot的第二段接口，用于执行计算。 | ||
| 54 | + * @param [in] workspace: 在npu device侧申请的workspace内存起址。 | ||
| 55 | + * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnLogicalNotGetWorkspaceSize获取。 | ||
| 56 | + * @param [in] executor: op执行器，包含了算子计算流程。 | ||
| 57 | + * @param [in] stream: acl stream流。 | ||
| 58 | + * @return aclnnStatus: 返回状态码。 | ||
| 59 | + */ | ||
| 60 | +ACLNN_API aclnnStatus aclnnLogicalNot(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, | ||
| 61 | + aclrtStream stream); | ||
| 62 | + | ||
| 63 | +/** | ||
| 64 | + * @brief aclnnInplaceLogicalNot的第一段接口，根据具体的计算流程，计算workspace大小。 | ||
| 65 | + * @domain aclnn_math | ||
| 66 | + * | ||
| 67 | + * 算子功能：完成逻辑取反计算 | ||
| 68 | + * | ||
| 69 | + * @param [in] selfRef: npu | ||
| 70 | + * device侧的aclTensor，数据类型支持FLOAT、FLOAT16、BFLOAT16、DOUBLE、INT32、INT64、INT16、INT8、UINT8、BOOL， | ||
| 71 | + * s持非连续的Tensor，数据格式支持ND。 | ||
| 72 | + * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。 | ||
| 73 | + * @param [out] executor: 返回op执行器，包含算子计算流程。 | ||
| 74 | + * @return aclnnStatus: 返回状态码。 | ||
| 75 | + */ | ||
| 76 | +ACLNN_API aclnnStatus aclnnInplaceLogicalNotGetWorkspaceSize(aclTensor* selfRef, uint64_t* workspaceSize, | ||
| 77 | + aclOpExecutor** executor); | ||
| 78 | + | ||
| 79 | +/** | ||
| 80 | + * @brief aclnnInplaceLogicalNot的第二段接口，用于执行计算。 | ||
| 81 | + * @param [in] workspace: 在npu device侧申请的workspace内存起址。 | ||
| 82 | + * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnLogicalNotGetWorkspaceSize获取。 | ||
| 83 | + * @param [in] executor: op执行器，包含了算子计算流程。 | ||
| 84 | + * @param [in] stream: acl stream流。 | ||
| 85 | + * @return aclnnStatus: 返回状态码。 | ||
| 86 | + */ | ||
| 87 | +ACLNN_API aclnnStatus aclnnInplaceLogicalNot(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, | ||
| 88 | + aclrtStream stream); | ||
| 89 | + | ||
| 90 | + | ||
| 91 | +} | ||
| 92 | + | ||
| 93 | + | ||
| 94 | + | ||
| @@ -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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | +using namespace op; | ||
| 22 | + | ||
| 23 | +namespace l0op { | ||
| 24 | +OP_TYPE_REGISTER(LogicalNot); | ||
| 25 | + | ||
| 26 | +// AICORE算子kernel | ||
| 27 | +static const aclTensor *LogicalNotAiCore(const aclTensor *self, aclTensor *out, aclOpExecutor *executor) { | ||
| 28 | + L0_DFX(LogicalNotAiCore, self, out); | ||
| 29 | + // 使用框架宏ADD_TO_LAUNCHER_LIST_AICORE，将Aicore LogicalNot算子加入任务队列 | ||
| 30 | + // LogicalNot是算子的OpType，self是算子的输入，out是算子的输出 | ||
| 31 | + auto retAicore = ADD_TO_LAUNCHER_LIST_AICORE(LogicalNot, OP_INPUT(self), OP_OUTPUT(out)); | ||
| 32 | + OP_CHECK_ADD_TO_LAUNCHER_LIST_AICORE(retAicore != ACLNN_SUCCESS, return nullptr, | ||
| 33 | + "LogicalNot ADD_TO_LAUNCHER_LIST_AICORE failed."); | ||
| 34 | + return out; | ||
| 35 | +} | ||
| 36 | + | ||
| 37 | +const aclTensor *LogicalNot(const aclTensor *self, aclOpExecutor *executor) { | ||
| 38 | + auto logicalNotOut = executor->AllocTensor(self->GetViewShape(), self->GetDataType()); | ||
| 39 | + return LogicalNotAiCore(self, logicalNotOut, executor); | ||
| 40 | +} | ||
| 41 | +} // namespace l0op | ||
| @@ -0,0 +1,20 @@ | |||
| 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 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | +namespace l0op { | ||
| 17 | + const aclTensor *LogicalNot(const aclTensor *self, aclOpExecutor *executor); | ||
| 18 | +} | ||
| 19 | + | ||
| 20 | + | ||
| @@ -0,0 +1,37 @@ | |||
| 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 logical_and.cpp | ||
| 13 | + * \brief | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | +namespace ops { | ||
| 18 | +class LogicalNot : public OpDef { | ||
| 19 | +public: | ||
| 20 | + explicit LogicalNot(const char* name) : OpDef(name) | ||
| 21 | + { | ||
| 22 | + this->Input("x") | ||
| 23 | + .ParamType(REQUIRED) | ||
| 24 | + .DataType({ge::DT_BOOL}) | ||
| 25 | + .Format({ge::FORMAT_ND}) | ||
| 26 | + .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 27 | + this->Output("y") | ||
| 28 | + .ParamType(REQUIRED) | ||
| 29 | + .DataType({ge::DT_BOOL}) | ||
| 30 | + .Format({ge::FORMAT_ND}) | ||
| 31 | + .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 32 | + | ||
| 33 | + this->AICore().AddConfig("ascend910b").AddConfig("ascend310p"); | ||
| 34 | + } | ||
| 35 | +}; | ||
| 36 | +OP_ADD(LogicalNot); | ||
| 37 | +} | ||
| @@ -0,0 +1,35 @@ | |||
| 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 logical_not_infer.cpp | ||
| 13 | + * \brief | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | +using namespace ge; | ||
| 19 | + | ||
| 20 | +namespace ops { | ||
| 21 | + | ||
| 22 | +static ge::graphStatus InferShapeLogicalNot(gert::InferShapeContext* context) | ||
| 23 | +{ | ||
| 24 | + OP_LOGD(context->GetNodeName(), "Begin to do InferShapeLogicalNot"); | ||
| 25 | + | ||
| 26 | + const gert::Shape *x1_shape = context->GetInputShape(0); | ||
| 27 | + gert::Shape *y_shape = context->GetOutputShape(0); | ||
| 28 | + *y_shape = *x1_shape; | ||
| 29 | + | ||
| 30 | + OP_LOGD(context->GetNodeName(), "End to do InferShapeLogicalNot"); | ||
| 31 | + return GRAPH_SUCCESS; | ||
| 32 | +} | ||
| 33 | + | ||
| 34 | +IMPL_OP_INFERSHAPE(LogicalNot).InferShape(InferShapeLogicalNot); | ||
| 35 | +} | ||
| @@ -0,0 +1,116 @@ | |||
| 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 logical_not_tiling.cpp | ||
| 13 | + * \brief | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | +namespace optiling { | ||
| 24 | + | ||
| 25 | +const uint32_t BLOCK_SIZE = 32; | ||
| 26 | +const uint32_t BUFFER_NUM = 2; | ||
| 27 | + | ||
| 28 | +struct LogicalNotCompileInfo {}; | ||
| 29 | + | ||
| 30 | +static ge::graphStatus LogicalNotTilingFunc(gert::TilingContext* context) | ||
| 31 | +{ | ||
| 32 | + LogicalNotTilingData* tiling = context->GetTilingData<LogicalNotTilingData>(); | ||
| 33 | + OP_CHECK_NULL_WITH_CONTEXT(context, tiling); | ||
| 34 | + OP_CHECK_IF(memset_s(tiling, sizeof(LogicalNotTilingData), 0, sizeof(LogicalNotTilingData)) != EOK, | ||
| 35 | + OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED); | ||
| 36 | + uint64_t ubLength = 0; | ||
| 37 | + uint32_t bigCoreDataNum = 0; | ||
| 38 | + uint32_t bigCoreLoopNum = 0; | ||
| 39 | + uint32_t bigCoreTailDataNum = 0; | ||
| 40 | + | ||
| 41 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo()); | ||
| 42 | + ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubLength); | ||
| 43 | + auto coreNum = ascendcPlatform.GetCoreNum(); | ||
| 44 | + | ||
| 45 | + // Based on the input length and the number of inputs, the number of bytes of the input data type is obtained | ||
| 46 | + uint32_t inputDataNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize(); | ||
| 47 | + uint32_t dataTypeLength = 0; | ||
| 48 | + ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), dataTypeLength); | ||
| 49 | + uint32_t inputLength = inputDataNum * dataTypeLength; | ||
| 50 | + if (coreNum == 0 || BLOCK_SIZE == 0) { | ||
| 51 | + OP_LOGE(context, "coreNum or BLOCK_SIZE is 0"); | ||
| 52 | + return ge::GRAPH_FAILED; | ||
| 53 | + } | ||
| 54 | + | ||
| 55 | + // There are a total of 3 shared UB spaces in the input and output. If it's int8, there are 2 more TBUFs | ||
| 56 | + uint32_t ubPartNum = 3; | ||
| 57 | + uint32_t ubPartLength = ubLength / ubPartNum / BUFFER_NUM; | ||
| 58 | + // The number of 32B data blocks that can be used for each data. DOUBLE BUFFER is already counted here | ||
| 59 | + uint32_t ubPartBlockNum = ubPartLength / BLOCK_SIZE; | ||
| 60 | + uint32_t ubPartDataNum = (ubPartBlockNum * BLOCK_SIZE) / dataTypeLength; | ||
| 61 | + | ||
| 62 | + // Input data for 32B alignment | ||
| 63 | + uint32_t inputLengthAlign32 = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE); | ||
| 64 | + | ||
| 65 | + if (ubPartDataNum >= inputDataNum) { | ||
| 66 | + coreNum = 1; | ||
| 67 | + } else { | ||
| 68 | + // There is at least 32B of data on each core, satisfying several settings for several cores. The maximum number of audits is the actual number of audits | ||
| 69 | + coreNum = (coreNum < inputLengthAlign32 / BLOCK_SIZE) ? coreNum : inputLengthAlign32 / BLOCK_SIZE; | ||
| 70 | + } | ||
| 71 | + | ||
| 72 | + uint32_t everyCoreInputBlockNum = inputLengthAlign32 / BLOCK_SIZE / coreNum; | ||
| 73 | + uint32_t tailBlockNum = (inputLengthAlign32 / BLOCK_SIZE) % coreNum; | ||
| 74 | + | ||
| 75 | + // Small chunks are calculated and sliced several times using the number of data on each core | ||
| 76 | + uint32_t smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / dataTypeLength; | ||
| 77 | + uint32_t smallCoreLoopNum = smallCoreDataNum / ubPartDataNum; | ||
| 78 | + smallCoreLoopNum = (everyCoreInputBlockNum % ubPartBlockNum) == 0 ? smallCoreLoopNum : smallCoreLoopNum + 1; | ||
| 79 | + // Tail block calculation for small chunks of data | ||
| 80 | + uint32_t smallCoreTailDataNum = smallCoreDataNum - ubPartDataNum * (smallCoreLoopNum - 1); | ||
| 81 | + smallCoreTailDataNum = smallCoreTailDataNum == 0 ? ubPartDataNum : smallCoreTailDataNum; | ||
| 82 | + | ||
| 83 | + if (0 != tailBlockNum) { | ||
| 84 | + everyCoreInputBlockNum += 1; | ||
| 85 | + bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / dataTypeLength; | ||
| 86 | + bigCoreLoopNum = bigCoreDataNum / ubPartDataNum; | ||
| 87 | + bigCoreLoopNum = (everyCoreInputBlockNum % ubPartBlockNum) == 0 ? bigCoreLoopNum : bigCoreLoopNum + 1; | ||
| 88 | + bigCoreTailDataNum = bigCoreDataNum - ubPartDataNum * (bigCoreLoopNum - 1); | ||
| 89 | + bigCoreTailDataNum = bigCoreTailDataNum == 0 ? ubPartDataNum : bigCoreTailDataNum; | ||
| 90 | + context->SetTilingKey(1); | ||
| 91 | + } else { | ||
| 92 | + context->SetTilingKey(0); | ||
| 93 | + } | ||
| 94 | + | ||
| 95 | + tiling->smallCoreDataNum = smallCoreDataNum; | ||
| 96 | + tiling->bigCoreDataNum = bigCoreDataNum; | ||
| 97 | + tiling->ubPartDataNum = ubPartDataNum; | ||
| 98 | + tiling->smallCoreTailDataNum = smallCoreTailDataNum; | ||
| 99 | + tiling->bigCoreTailDataNum = bigCoreTailDataNum; | ||
| 100 | + tiling->smallCoreLoopNum = smallCoreLoopNum; | ||
| 101 | + tiling->bigCoreLoopNum = bigCoreLoopNum; | ||
| 102 | + tiling->tailBlockNum = tailBlockNum; | ||
| 103 | + context->SetBlockDim(coreNum); | ||
| 104 | + | ||
| 105 | + size_t *currentWorkspace = context->GetWorkspaceSizes(1); | ||
| 106 | + currentWorkspace[0] = 0; | ||
| 107 | + return ge::GRAPH_SUCCESS; | ||
| 108 | +} | ||
| 109 | + | ||
| 110 | +static ge::graphStatus TilingParseForLogicalNot([[maybe_unused]] gert::TilingParseContext* context) | ||
| 111 | +{ | ||
| 112 | + return ge::GRAPH_SUCCESS; | ||
| 113 | +} | ||
| 114 | + | ||
| 115 | +IMPL_OP_OPTILING(LogicalNot).Tiling(LogicalNotTilingFunc).TilingParse<LogicalNotCompileInfo>(TilingParseForLogicalNot); | ||
| 116 | +} | ||
| @@ -0,0 +1,50 @@ | |||
| 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 logical_not.cpp | ||
| 13 | + * \brief | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | +enum class LogicalNotTilingKey : uint32_t | ||
| 19 | +{ | ||
| 20 | + TILING_KEY_EXAMPLE_FLOAT = 0, | ||
| 21 | + TILING_KEY_EXAMPLE_INT32 = 1, | ||
| 22 | +}; | ||
| 23 | + | ||
| 24 | +template <uint32_t schMode> | ||
| 25 | +__global__ __aicore__ void logical_not(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) | ||
| 26 | +{ | ||
| 27 | + REGISTER_TILING_DEFAULT(LogicalNotTilingData); | ||
| 28 | + GET_TILING_DATA_WITH_STRUCT(LogicalNotTilingData, tilingData, tiling); | ||
| 29 | + | ||
| 30 | + if (TILING_KEY_IS(1)) | ||
| 31 | + { | ||
| 32 | + NsLogicalNot::LogicalNot<int8_t, int8_t, true> op; | ||
| 33 | + op.Init(x, y, tilingData.smallCoreDataNum, | ||
| 34 | + tilingData.bigCoreDataNum, tilingData.bigCoreLoopNum, | ||
| 35 | + tilingData.smallCoreLoopNum, tilingData.ubPartDataNum, | ||
| 36 | + tilingData.smallCoreTailDataNum, tilingData.bigCoreTailDataNum, | ||
| 37 | + tilingData.tailBlockNum); | ||
| 38 | + op.Process(); | ||
| 39 | + } | ||
| 40 | + else if (TILING_KEY_IS(0)) | ||
| 41 | + { | ||
| 42 | + NsLogicalNot::LogicalNot<int8_t, int8_t, false> op; | ||
| 43 | + op.Init(x, y, tilingData.smallCoreDataNum, | ||
| 44 | + tilingData.bigCoreDataNum, tilingData.bigCoreLoopNum, | ||
| 45 | + tilingData.smallCoreLoopNum, tilingData.ubPartDataNum, | ||
| 46 | + tilingData.smallCoreTailDataNum, tilingData.bigCoreTailDataNum, | ||
| 47 | + tilingData.tailBlockNum); | ||
| 48 | + op.Process(); | ||
| 49 | + } | ||
| 50 | +} | ||
| @@ -0,0 +1,129 @@ | |||
| 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 logical_not.h | ||
| 13 | + * \brief | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | +namespace NsLogicalNot { | ||
| 24 | + | ||
| 25 | +using namespace AscendC; | ||
| 26 | + | ||
| 27 | +constexpr int32_t BUFFER_NUM = 2; | ||
| 28 | +constexpr half ONE = 1.0f; | ||
| 29 | +constexpr half NEGATIVE_ONE = -1.0f; | ||
| 30 | + | ||
| 31 | +template <typename TYPE_X, typename TYPE_Y, bool IsExistBigCore> | ||
| 32 | +class LogicalNot { | ||
| 33 | +public: | ||
| 34 | + __aicore__ inline LogicalNot() {} | ||
| 35 | + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t smallCoreDataNum, | ||
| 36 | + uint32_t bigCoreDataNum, uint32_t bigCoreLoopNum, | ||
| 37 | + uint32_t smallCoreLoopNum, uint32_t ubPartDataNum, | ||
| 38 | + uint32_t smallCoreTailDataNum, uint32_t bigCoreTailDataNum, | ||
| 39 | + uint32_t tailBlockNum) | ||
| 40 | + { | ||
| 41 | + ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!"); | ||
| 42 | + uint32_t blockIdx = AscendC::GetBlockIdx(); | ||
| 43 | + uint32_t globalBufferIndex = bigCoreDataNum * AscendC::GetBlockIdx(); | ||
| 44 | + this->ubPartDataNum = ubPartDataNum; | ||
| 45 | + if constexpr (IsExistBigCore) { | ||
| 46 | + if (blockIdx < tailBlockNum) { | ||
| 47 | + this->coreDataNum = bigCoreDataNum; | ||
| 48 | + this->tileNum = bigCoreLoopNum; | ||
| 49 | + this->tailDataNum = bigCoreTailDataNum; | ||
| 50 | + } else { | ||
| 51 | + this->coreDataNum = smallCoreDataNum; | ||
| 52 | + this->tileNum = smallCoreLoopNum; | ||
| 53 | + this->tailDataNum = smallCoreTailDataNum; | ||
| 54 | + globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (AscendC::GetBlockIdx() - tailBlockNum); | ||
| 55 | + } | ||
| 56 | + } else { | ||
| 57 | + this->coreDataNum = smallCoreDataNum; | ||
| 58 | + this->tileNum = smallCoreLoopNum; | ||
| 59 | + this->tailDataNum = smallCoreTailDataNum; | ||
| 60 | + globalBufferIndex = smallCoreDataNum * AscendC::GetBlockIdx(); | ||
| 61 | + } | ||
| 62 | + | ||
| 63 | + xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum); | ||
| 64 | + yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum); | ||
| 65 | + pipe.InitBuffer(inQueueX, BUFFER_NUM, this->ubPartDataNum * sizeof(TYPE_X)); | ||
| 66 | + pipe.InitBuffer(tmp1, this->ubPartDataNum * sizeof(half)); | ||
| 67 | + pipe.InitBuffer(outQueueY, BUFFER_NUM, this->ubPartDataNum * sizeof(TYPE_Y)); | ||
| 68 | + } | ||
| 69 | + | ||
| 70 | + __aicore__ inline void Process() | ||
| 71 | + { | ||
| 72 | + int32_t loopCount = this->tileNum; | ||
| 73 | + this->processDataNum = this->ubPartDataNum; | ||
| 74 | + for (int32_t i = 0; i < loopCount - 1; i++) { | ||
| 75 | + CopyIn(i); | ||
| 76 | + Compute(i); | ||
| 77 | + CopyOut(i); | ||
| 78 | + } | ||
| 79 | + this->processDataNum = this->tailDataNum; | ||
| 80 | + CopyIn(loopCount - 1); | ||
| 81 | + Compute(loopCount - 1); | ||
| 82 | + CopyOut(loopCount - 1); | ||
| 83 | + } | ||
| 84 | + | ||
| 85 | +private: | ||
| 86 | + __aicore__ inline void CopyIn(int32_t progress) | ||
| 87 | + { | ||
| 88 | + AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.AllocTensor<TYPE_X>(); | ||
| 89 | + AscendC::DataCopy(xLocal, xGm[progress * this->ubPartDataNum], this->processDataNum); | ||
| 90 | + inQueueX.EnQue(xLocal); | ||
| 91 | + } | ||
| 92 | + | ||
| 93 | + __aicore__ inline void Compute(int32_t progress) | ||
| 94 | + { | ||
| 95 | + AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.DeQue<TYPE_X>(); | ||
| 96 | + AscendC::LocalTensor<TYPE_Y> yLocal = outQueueY.AllocTensor<TYPE_Y>(); | ||
| 97 | + AscendC::LocalTensor<half> tmp1Local = tmp1.Get<half>(); | ||
| 98 | + AscendC::Cast(tmp1Local, xLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum); | ||
| 99 | + AscendC::Abs(tmp1Local, tmp1Local, this->processDataNum); | ||
| 100 | + AscendC::Mins(tmp1Local, tmp1Local, ONE, this->processDataNum); | ||
| 101 | + AscendC::Adds(tmp1Local, tmp1Local, NEGATIVE_ONE, this->processDataNum); | ||
| 102 | + AscendC::Abs(tmp1Local, tmp1Local, this->processDataNum); | ||
| 103 | + AscendC::Cast(yLocal, tmp1Local, AscendC::RoundMode::CAST_NONE, this->processDataNum); | ||
| 104 | + outQueueY.EnQue<TYPE_Y>(yLocal); | ||
| 105 | + inQueueX.FreeTensor(xLocal); | ||
| 106 | + } | ||
| 107 | + | ||
| 108 | + __aicore__ inline void CopyOut(int32_t progress) | ||
| 109 | + { | ||
| 110 | + AscendC::LocalTensor<TYPE_Y> yLocal = outQueueY.DeQue<TYPE_Y>(); | ||
| 111 | + AscendC::DataCopy(yGm[progress * this->ubPartDataNum], yLocal, this->processDataNum); | ||
| 112 | + outQueueY.FreeTensor(yLocal); | ||
| 113 | + } | ||
| 114 | + | ||
| 115 | +private: | ||
| 116 | + AscendC::TPipe pipe; | ||
| 117 | + AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX; | ||
| 118 | + AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueY; | ||
| 119 | + AscendC::TBuf<AscendC::QuePosition::VECCALC> tmp1; | ||
| 120 | + AscendC::GlobalTensor<TYPE_X> xGm; | ||
| 121 | + AscendC::GlobalTensor<TYPE_Y> yGm; | ||
| 122 | + uint32_t coreDataNum = 0; | ||
| 123 | + uint32_t tileNum = 0; | ||
| 124 | + uint32_t ubPartDataNum = 0; | ||
| 125 | + uint32_t tailDataNum = 0; | ||
| 126 | + uint32_t processDataNum = 0; | ||
| 127 | +}; | ||
| 128 | +} | ||
| 129 | + | ||
| @@ -0,0 +1,29 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + */ | ||
| 10 | + | ||
| 11 | +/*! | ||
| 12 | + * \file logical_not_tiling_data.h | ||
| 13 | + * \brief tiling data struct | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +struct LogicalNotTilingData { | ||
| 20 | + uint32_t smallCoreDataNum; | ||
| 21 | + uint32_t bigCoreDataNum; | ||
| 22 | + uint32_t bigCoreLoopNum; | ||
| 23 | + uint32_t smallCoreLoopNum; | ||
| 24 | + uint32_t ubPartDataNum; | ||
| 25 | + uint32_t smallCoreTailDataNum; | ||
| 26 | + uint32_t bigCoreTailDataNum; | ||
| 27 | + uint32_t tailBlockNum; | ||
| 28 | +}; | ||
| 29 | + | ||
| @@ -0,0 +1,35 @@ | |||
| 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 logical_not_tiling_key.h | ||
| 13 | + * \brief logical_not tiling key declare | ||
| 14 | + */ | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | +/* Mode场景定义 */ | ||
| 22 | + | ||
| 23 | + | ||
| 24 | +/* 继续定义其他Mode场景... */ | ||
| 25 | + | ||
| 26 | +/* 模板参数 */ | ||
| 27 | +ASCENDC_TPL_ARGS_DECL( | ||
| 28 | + LogicalNot, | ||
| 29 | + ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)); | ||
| 30 | + | ||
| 31 | +/* 模板参数组合 */ | ||
| 32 | +ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL( | ||
| 33 | + ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1))); | ||
| 34 | + | ||
| 35 | + | ||
后面内容根据实际情况填写