已合并
refactor: flatten add_custom structure in parallel_ops_package #1603
jianwei10创建于 4月14日
refactor: flatten add_custom structure in parallel_ops_package #1603
已合并
共 12 个文件变更+79-217
| @@ -13,36 +13,37 @@ cmake_minimum_required(VERSION 3.16.0) | |||
| 13 | project(parallel_ops_package_proj) | 13 | project(parallel_ops_package_proj) |
| 14 | include(ExternalProject) | 14 | include(ExternalProject) |
| 15 | 15 | ||
| 16 | -set(ASCEND_COMPUTE_UNIT ascend910b ascend910_93 ascend950) | ||
| 17 | -set(PACKAGE_TYPE "RUN") | ||
| 18 | set(INSTALL_TARGET "package") | 16 | set(INSTALL_TARGET "package") |
| 19 | 17 | ||
| 20 | -ExternalProject_Add(add_custom | 18 | +set(ADD_CUSTOM_ASCEND_COMPUTE_UNIT ascend910 ascend310p ascend310b ascend910b ascend950) |
| 21 | - SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/add_custom | 19 | +set(LEAKY_RELU_ASCEND_COMPUTE_UNIT ascend910b) |
| 22 | - CMAKE_CACHE_ARGS | 20 | + |
| 23 | - -DASCEND_COMPUTE_UNIT:STRING=${ASCEND_COMPUTE_UNIT} | 21 | +function(add_parallel_vendor_project target_name source_dir vendor_name compute_units) |
| 24 | - CMAKE_ARGS | 22 | + ExternalProject_Add(${target_name} |
| 25 | - -Dvendor_name=add_custom | 23 | + SOURCE_DIR ${source_dir} |
| 26 | - -DCMAKE_INSTALL_PREFIX=${CMAKE_BINARY_DIR}/add_custom | 24 | + CMAKE_CACHE_ARGS |
| 27 | - -DPACKAGE_TYPE=${PACKAGE_TYPE} | 25 | + -DASCEND_COMPUTE_UNIT:STRING=${compute_units} |
| 28 | - BUILD_COMMAND | 26 | + CMAKE_ARGS |
| 29 | - ${CMAKE_COMMAND} --build <BINARY_DIR> --target binary | 27 | + -Dvendor_name=${vendor_name} |
| 30 | - COMMAND | 28 | + -DCMAKE_INSTALL_PREFIX=${CMAKE_BINARY_DIR}/${target_name} |
| 31 | - ${CMAKE_COMMAND} --build <BINARY_DIR> --target ${INSTALL_TARGET} | 29 | + BUILD_COMMAND |
| 32 | - INSTALL_COMMAND "" | 30 | + ${CMAKE_COMMAND} --build <BINARY_DIR> --target binary |
| 31 | + COMMAND | ||
| 32 | + ${CMAKE_COMMAND} --build <BINARY_DIR> --target ${INSTALL_TARGET} | ||
| 33 | + INSTALL_COMMAND "" | ||
| 34 | + ) | ||
| 35 | +endfunction() | ||
| 36 | + | ||
| 37 | +add_parallel_vendor_project( | ||
| 38 | + add_custom | ||
| 39 | + ${CMAKE_CURRENT_SOURCE_DIR}/add_custom | ||
| 40 | + add_custom | ||
| 41 | + "${ADD_CUSTOM_ASCEND_COMPUTE_UNIT}" | ||
| 33 | ) | 42 | ) |
| 34 | 43 | ||
| 35 | -ExternalProject_Add(leaky_relu_custom | 44 | +add_parallel_vendor_project( |
| 36 | - SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/leaky_relu_custom | 45 | + leaky_relu_custom |
| 37 | - CMAKE_CACHE_ARGS | 46 | + ${CMAKE_CURRENT_SOURCE_DIR}/leaky_relu_custom |
| 38 | - -DASCEND_COMPUTE_UNIT:STRING=${ASCEND_COMPUTE_UNIT} | 47 | + leaky_relu_custom |
| 39 | - CMAKE_ARGS | 48 | + "${LEAKY_RELU_ASCEND_COMPUTE_UNIT}" |
| 40 | - -Dvendor_name=leaky_relu_custom | ||
| 41 | - -DCMAKE_INSTALL_PREFIX=${CMAKE_BINARY_DIR}/leaky_relu_custom | ||
| 42 | - -DPACKAGE_TYPE=${PACKAGE_TYPE} | ||
| 43 | - BUILD_COMMAND | ||
| 44 | - ${CMAKE_COMMAND} --build <BINARY_DIR> --target binary | ||
| 45 | - COMMAND | ||
| 46 | - ${CMAKE_COMMAND} --build <BINARY_DIR> --target ${INSTALL_TARGET} | ||
| 47 | - INSTALL_COMMAND "" | ||
| 48 | ) | 49 | ) |
| @@ -6,6 +6,8 @@ | |||
| 6 | - `add_custom`（AddCustom） | 6 | - `add_custom`（AddCustom） |
| 7 | - `leaky_relu_custom`（LeakyReluCustom） | 7 | - `leaky_relu_custom`（LeakyReluCustom） |
| 8 | 8 | ||
| 9 | +其中 `add_custom` 采用扁平目录组织（host/kernel/tiling 源码位于同一层级），`leaky_relu_custom` 保持分层目录结构。 | ||
| 10 | + | ||
| 9 | 每个子工程会分别完成自定义算子的编译、打包，并生成独立的 `custom_opp_*.run` 安装包。 | 11 | 每个子工程会分别完成自定义算子的编译、打包，并生成独立的 `custom_opp_*.run` 安装包。 |
| 10 | 12 | ||
| 11 | ## 支持的产品 | 13 | ## 支持的产品 |
| @@ -27,20 +29,9 @@ parallel_ops_package | |||
| 27 | ├── README.md | 29 | ├── README.md |
| 28 | ├── add_custom | 30 | ├── add_custom |
| 29 | │ ├── CMakeLists.txt | 31 | │ ├── CMakeLists.txt |
| 30 | -│ ├── framework | 32 | +│ ├── add_custom_host.cpp |
| 31 | -│ │ ├── CMakeLists.txt | 33 | +│ ├── add_custom_kernel.cpp |
| 32 | -│ │ └── tf_plugin | 34 | +│ └── add_custom_tiling.h |
| 33 | -│ │ ├── CMakeLists.txt | ||
| 34 | -│ │ └── tensorflow_add_custom_plugin.cc | ||
| 35 | -│ ├── op_host | ||
| 36 | -│ │ ├── CMakeLists.txt | ||
| 37 | -│ │ └── add_custom | ||
| 38 | -│ │ └── add_custom_host.cpp | ||
| 39 | -│ └── op_kernel | ||
| 40 | -│ ├── CMakeLists.txt | ||
| 41 | -│ └── add_custom | ||
| 42 | -│ ├── add_custom_kernel.cpp | ||
| 43 | -│ └── add_custom_tiling.h | ||
| 44 | └── leaky_relu_custom | 35 | └── leaky_relu_custom |
| 45 | ├── CMakeLists.txt | 36 | ├── CMakeLists.txt |
| 46 | ├── framework | 37 | ├── framework |
| @@ -61,19 +52,19 @@ parallel_ops_package | |||
| 61 | 52 | ||
| 62 | ## 样例描述 | 53 | ## 样例描述 |
| 63 | 54 | ||
| 64 | -`parallel_ops_package` 与 `custom_op` 使用相同的 Add/LeakyRelu 样例描述，本文不重复维护，请参考： | 55 | +`parallel_ops_package` 与 `custom_op` 使用相同的 Add/LeakyRelu 样例描述，请参考： |
 | |||
| 65 | 56 | ||
| 66 | - [custom_op/README.md 的“样例描述”章节](../custom_op/README.md#样例描述) | 57 | - [custom_op/README.md 的“样例描述”章节](../custom_op/README.md#样例描述) |
| 67 | 58 | ||
| 68 | ## 样例规格描述 | 59 | ## 样例规格描述 |
| 69 | 60 | ||
| 70 | -`parallel_ops_package` 与 `custom_op` 使用相同的 Add/LeakyRelu 规格，本文不重复维护规格表，请参考： | 61 | +`parallel_ops_package` 与 `custom_op` 使用相同的 Add/LeakyRelu 规格描述，请参考： |
| 71 | 62 | ||
| 72 | - [custom_op/README.md 的“样例规格描述”章节](../custom_op/README.md#样例规格描述) | 63 | - [custom_op/README.md 的“样例规格描述”章节](../custom_op/README.md#样例规格描述) |
| 73 | 64 | ||
| 74 | ## 代码实现介绍 | 65 | ## 代码实现介绍 |
| 75 | 66 | ||
| 76 | -`parallel_ops_package` 的算子实现与 `custom_op` 保持一致，本文不重复维护实现细节，请参考： | 67 | +`parallel_ops_package` 的 Add/LeakyRelu 代码实现可参考 `custom_op` 文档说明。 |
| 77 | 68 | ||
| 78 | - [custom_op/README.md 的“代码实现介绍”章节](../custom_op/README.md#代码实现介绍) | 69 | - [custom_op/README.md 的“代码实现介绍”章节](../custom_op/README.md#代码实现介绍) |
| 79 | 70 | ||
Mexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/CMakeLists.txt+33-15
| @@ -13,23 +13,41 @@ cmake_minimum_required(VERSION 3.16.0) | |||
| 13 | project(opp) | 13 | project(opp) |
| 14 | 14 | ||
| 15 | find_package(ASC REQUIRED) | 15 | find_package(ASC REQUIRED) |
| 16 | -set(package_name ${vendor_name}) | ||
| 17 | -set(package_type ${PACKAGE_TYPE}) | ||
| 18 | 16 | ||
| 19 | -npu_op_package(${package_name} | 17 | +npu_op_package(${vendor_name} |
| 20 | - TYPE ${package_type} | 18 | + TYPE RUN |
| 21 | - CONFIG | ||
| 22 | - INSTALL_PATH ${CMAKE_INSTALL_PREFIX} | ||
| 23 | ) | 19 | ) |
| 24 | 20 | ||
| 25 | -if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/framework") | 21 | +file(GLOB host_ops_srcs ${CMAKE_CURRENT_SOURCE_DIR}/add_custom_host.cpp) |
| 26 | - add_subdirectory(framework) | 22 | +npu_op_code_gen( |
| 27 | -endif() | 23 | + SRC ${host_ops_srcs} |
| 24 | + PACKAGE ${vendor_name} | ||
| 25 | + OUT_DIR ${ASCEND_AUTOGEN_PATH} | ||
| 26 | +) | ||
| 28 | 27 | ||
| 29 | -if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_host") | 28 | +file(GLOB autogen_aclnn_srcs ${ASCEND_AUTOGEN_PATH}/aclnn_*.cpp) |
| 30 | - add_subdirectory(op_host) | 29 | +set_source_files_properties(${autogen_aclnn_srcs} PROPERTIES GENERATED TRUE) |
| 31 | -endif() | 30 | +npu_op_library(cust_opapi ACLNN |
| 31 | + ${autogen_aclnn_srcs} | ||
| 32 | +) | ||
| 32 | 33 | ||
| 33 | -if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_kernel") | 34 | +npu_op_library(cust_optiling TILING |
| 34 | - add_subdirectory(op_kernel) | 35 | + ${host_ops_srcs} |
| 35 | -endif() | 36 | +) |
| 37 | + | ||
| 38 | +npu_op_kernel_library(ascendc_kernels | ||
| 39 | + SRC_BASE ${CMAKE_CURRENT_SOURCE_DIR}/ | ||
| 40 | + TILING_LIBRARY cust_optiling | ||
| 41 | +) | ||
| 42 | + | ||
| 43 | +npu_op_kernel_sources(ascendc_kernels | ||
| 44 | + OP_TYPE AddCustom | ||
| 45 | + KERNEL_FILE add_custom_kernel.cpp | ||
| 46 | +) | ||
| 47 | + | ||
| 48 | +npu_op_package_add(${vendor_name} | ||
| 49 | + LIBRARY | ||
| 50 | + cust_opapi | ||
| 51 | + cust_optiling | ||
| 52 | + ascendc_kernels | ||
| 53 | +) | ||
Rexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/op_host/add_custom/add_custom_host.cpp→examples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/add_custom_host.cpp+4-23
| @@ -8,9 +8,9 @@ | |||
| 8 | * See LICENSE in the root of the software repository for the full text of the License. | 8 | * See LICENSE in the root of the software repository for the full text of the License. |
| 9 | */ | 9 | */ |
| 10 | 10 | ||
| 11 | - | 11 | +#include "add_custom_tiling.h" |
| 12 | - | ||
| 13 | 12 | ||
| 13 | + | ||
| 14 | 14 | ||
| 15 | namespace optiling { | 15 | namespace optiling { |
| 16 | const uint32_t NUM_BLOCKS = 8; | 16 | const uint32_t NUM_BLOCKS = 8; |
| @@ -22,28 +22,10 @@ static ge::graphStatus TilingFunc(gert::TilingContext *context) | |||
| 22 | context->SetBlockDim(NUM_BLOCKS); | 22 | context->SetBlockDim(NUM_BLOCKS); |
| 23 | tiling->totalLength = totalLength; | 23 | tiling->totalLength = totalLength; |
| 24 | tiling->tileNum = TILE_NUM; | 24 | tiling->tileNum = TILE_NUM; |
| 25 | - size_t *currentWorkspace = context->GetWorkspaceSizes(1); | ||
| 26 | - currentWorkspace[0] = 0; | ||
| 27 | return ge::GRAPH_SUCCESS; | 25 | return ge::GRAPH_SUCCESS; |
| 28 | } | 26 | } |
| 29 | } // namespace optiling | 27 | } // namespace optiling |
| 30 | 28 | ||
| 31 | -namespace ge { | ||
| 32 | -static graphStatus InferShape(gert::InferShapeContext *context) | ||
| 33 | -{ | ||
| 34 | - const gert::Shape *x1_shape = context->GetInputShape(0); | ||
| 35 | - gert::Shape *y_shape = context->GetOutputShape(0); | ||
| 36 | - *y_shape = *x1_shape; | ||
| 37 | - return GRAPH_SUCCESS; | ||
| 38 | -} | ||
| 39 | - | ||
| 40 | -static graphStatus InferDataType(gert::InferDataTypeContext *context) | ||
| 41 | -{ | ||
| 42 | - const auto inputDataType = context->GetInputDataType(0); | ||
| 43 | - context->SetOutputDataType(0, inputDataType); | ||
| 44 | - return ge::GRAPH_SUCCESS; | ||
| 45 | -} | ||
| 46 | -} // namespace ge | ||
| 47 | 29 | ||
| 48 | namespace ops { | 30 | namespace ops { |
| 49 | class AddCustom : public OpDef { | 31 | class AddCustom : public OpDef { |
| @@ -63,10 +45,9 @@ public: | |||
| 63 | .DataType({ge::DT_FLOAT16}) | 45 | .DataType({ge::DT_FLOAT16}) |
| 64 | .Format({ge::FORMAT_ND}); | 46 | .Format({ge::FORMAT_ND}); |
| 65 | 47 | ||
| 66 | - this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType); | ||
| 67 | this->AICore() | 48 | this->AICore() |
| 68 | .SetTiling(optiling::TilingFunc) | 49 | .SetTiling(optiling::TilingFunc) |
| 69 | - // at least one soc version must be configured. | 50 | + .AddConfig("ascend910") |
| 70 | .AddConfig("ascend310p") | 51 | .AddConfig("ascend310p") |
| 71 | .AddConfig("ascend310b") | 52 | .AddConfig("ascend310b") |
| 72 | .AddConfig("ascend910b") | 53 | .AddConfig("ascend910b") |
| @@ -75,4 +56,4 @@ public: | |||
| 75 | } | 56 | } |
| 76 | }; | 57 | }; |
| 77 | OP_ADD(AddCustom); | 58 | OP_ADD(AddCustom); |
| 78 | -} // namespace ops | 59 | +} // namespace ops |
Rexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/op_kernel/add_custom/add_custom_kernel.cpp→examples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/add_custom_kernel.cpp+3-4
| @@ -8,7 +8,6 @@ | |||
| 8 | * See LICENSE in the root of the software repository for the full text of the License. | 8 | * See LICENSE in the root of the software repository for the full text of the License. |
| 9 | */ | 9 | */ |
| 10 | 10 | ||
| 11 | - | ||
| 12 | 11 | ||
| 13 | 12 | ||
| 14 | constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue | 13 | constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue |
| @@ -16,9 +15,9 @@ constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue | |||
| 16 | class KernelAdd { | 15 | class KernelAdd { |
| 17 | public: | 16 | public: |
| 18 | __aicore__ inline KernelAdd() {} | 17 | __aicore__ inline KernelAdd() {} |
| 19 | - __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum, AscendC::TPipe* pipeIn) | 18 | + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum, |
| 19 | + AscendC::TPipe* pipeIn) | ||
| 20 | { | 20 | { |
| 21 | - ascendc_assert(tileNum != 0, "tileNum can not be zero.\n"); | ||
| 22 | this->pipe = pipeIn; | 21 | this->pipe = pipeIn; |
| 23 | this->blockLength = totalLength / AscendC::GetBlockNum(); | 22 | this->blockLength = totalLength / AscendC::GetBlockNum(); |
| 24 | this->tileNum = tileNum; | 23 | this->tileNum = tileNum; |
| @@ -88,4 +87,4 @@ extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z | |||
| 88 | KernelAdd op; | 87 | KernelAdd op; |
| 89 | op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum, &pipe); | 88 | op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum, &pipe); |
| 90 | op.Process(); | 89 | op.Process(); |
| 91 | -} | 90 | +} |
Rexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/op_kernel/add_custom/add_custom_tiling.h→examples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/add_custom_tiling.h+1-2文件内容审核中，请稍后刷新重试
| @@ -1,18 +0,0 @@ | |||
| 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 | -if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/tf_plugin") | ||
| 13 | - add_subdirectory(tf_plugin) | ||
| 14 | -endif() | ||
| 15 | - | ||
| 16 | -if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/onnx_plugin") | ||
| 17 | - add_subdirectory(onnx_plugin) | ||
| 18 | -endif() | ||
| @@ -1,21 +0,0 @@ | |||
| 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 | -aux_source_directory(${CMAKE_CURRENT_SOURCE_DIR} plugin_srcs) | ||
| 13 | - | ||
| 14 | -npu_op_library(cust_tf_parsers TF_PLUGIN | ||
| 15 | - ${plugin_srcs} | ||
| 16 | -) | ||
| 17 | - | ||
| 18 | -npu_op_package_add(${package_name} | ||
| 19 | - LIBRARY | ||
| 20 | - cust_tf_parsers | ||
| 21 | -) | ||
| @@ -1,19 +0,0 @@ | |||
| 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 | -namespace domi { | ||
| 14 | -// register op info to GE | ||
| 15 | -REGISTER_CUSTOM_OP("AddCustom") | ||
| 16 | - .FrameworkType(TENSORFLOW) // type: CAFFE, TENSORFLOW | ||
| 17 | - .OriginOpType("AddCustom") // name in tf module | ||
| 18 | - .ParseParamsByOperatorFn(AutoMappingByOpFn); | ||
| 19 | -} // namespace domi | ||
| @@ -1,42 +0,0 @@ | |||
| 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(GLOB host_ops_srcs | ||
| 13 | - add_custom/add_custom_host.cpp | ||
| 14 | -) | ||
| 15 | -npu_op_code_gen( | ||
| 16 | - SRC ${host_ops_srcs} | ||
| 17 | - PACKAGE ${package_name} | ||
| 18 | - OUT_DIR ${ASCEND_AUTOGEN_PATH} | ||
| 19 | -) | ||
| 20 | -npu_op_library(cust_optiling TILING | ||
| 21 | - ${host_ops_srcs} | ||
| 22 | -) | ||
| 23 | - | ||
| 24 | -file(GLOB autogen_aclnn_src ${ASCEND_AUTOGEN_PATH}/aclnn_*.cpp) | ||
| 25 | -set_source_files_properties(${autogen_aclnn_src} PROPERTIES GENERATED TRUE) | ||
| 26 | -npu_op_library(cust_opapi ACLNN | ||
| 27 | - ${autogen_aclnn_src} | ||
| 28 | -) | ||
| 29 | - | ||
| 30 | -file(GLOB proto_src ${ASCEND_AUTOGEN_PATH}/op_proto.cc) | ||
| 31 | -set_source_files_properties(${proto_src} PROPERTIES GENERATED TRUE) | ||
| 32 | -npu_op_library(cust_op_proto GRAPH | ||
| 33 | - ${host_ops_srcs} | ||
| 34 | - ${proto_src} | ||
| 35 | -) | ||
| 36 | - | ||
| 37 | -npu_op_package_add(${package_name} | ||
| 38 | - LIBRARY | ||
| 39 | - cust_optiling | ||
| 40 | - cust_opapi | ||
| 41 | - cust_op_proto | ||
| 42 | -) | ||
| @@ -1,26 +0,0 @@ | |||
| 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 | -npu_op_kernel_sources(ascendc_kernels | ||
| 13 | - OP_TYPE AddCustom | ||
| 14 | - KERNEL_DIR add_custom | ||
| 15 | - KERNEL_FILE add_custom_kernel.cpp | ||
| 16 | -) | ||
| 17 | - | ||
| 18 | -npu_op_kernel_library(ascendc_kernels | ||
| 19 | - SRC_BASE ${CMAKE_CURRENT_SOURCE_DIR} | ||
| 20 | - TILING_LIBRARY cust_optiling | ||
| 21 | -) | ||
| 22 | - | ||
| 23 | -npu_op_package_add(${package_name} | ||
| 24 | - LIBRARY | ||
| 25 | - ascendc_kernels | ||
| 26 | -) | ||
| @@ -14,10 +14,9 @@ project(opp) | |||
| 14 | 14 | ||
| 15 | find_package(ASC REQUIRED) | 15 | find_package(ASC REQUIRED) |
| 16 | set(package_name ${vendor_name}) | 16 | set(package_name ${vendor_name}) |
| 17 | -set(package_type ${PACKAGE_TYPE}) | ||
| 18 | 17 | ||
| 19 | npu_op_package(${package_name} | 18 | npu_op_package(${package_name} |
| 20 | - TYPE ${package_type} | 19 | + TYPE RUN |
| 21 | CONFIG | 20 | CONFIG |
| 22 | INSTALL_PATH ${CMAKE_INSTALL_PREFIX} | 21 | INSTALL_PATH ${CMAKE_INSTALL_PREFIX} |
| 23 | ) | 22 | ) |
这里写“Add/LeakyRelu 与 custom_op 使用相同样例描述”，但本 PR 已将 add_custom 改成扁平结构并删除 framework/tf_plugin、op_host/op_kernel 子工程，和 custom_op 的工程组织已经不一致。建议拆分成两段说明（add_custom 与 leaky_relu_custom 分开描述），避免读者按 custom_op 路径查找时产生误导。