diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/CMakeLists.txt b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/CMakeLists.txt new file mode 100644 index 0000000000000000000000000000000000000000..a7251570af2fb9ceca52263aa06c727b24032256 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/CMakeLists.txt @@ -0,0 +1,21 @@ +cmake_minimum_required(VERSION 3.18.0) +project(opp) + +find_package(AscendCompile REQUIRED HINTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) + +npu_op_package(${vendor_name} + TYPE RUN + CONFIG + ENABLE_SOURCE_PACKAGE True + ENABLE_BINARY_PACKAGE True + INSTALL_PATH ${CMAKE_BINARY_DIR}/) + +if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/framework) + add_subdirectory(framework) +endif() +if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/op_host) + add_subdirectory(op_host) +endif() +if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/op_kernel) + add_subdirectory(op_kernel) +endif() diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/CMakePresets.json b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/CMakePresets.json new file mode 100644 index 0000000000000000000000000000000000000000..a6123671e378a3dd1f8fb7965c5d8054f606c9c6 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/CMakePresets.json @@ -0,0 +1,63 @@ +{ + "version": 1, + "cmakeMinimumRequired": { + "major": 3, + "minor": 19, + "patch": 0 + }, + "configurePresets": [ + { + "name": "default", + "displayName": "Default Config", + "description": "Default build using Unix Makefiles generator", + "generator": "Unix Makefiles", + "binaryDir": "${sourceDir}/build_out", + "cacheVariables": { + "CMAKE_BUILD_TYPE": { + "type": "STRING", + "value": "Release" + }, + "ENABLE_SOURCE_PACKAGE": { + "type": "BOOL", + "value": "True" + }, + "ENABLE_BINARY_PACKAGE": { + "type": "BOOL", + "value": "True" + }, + "ASCEND_COMPUTE_UNIT": { + "type": "STRING", + "value": "ascend910b" + }, + "ENABLE_TEST": { + "type": "BOOL", + "value": "True" + }, + "vendor_name": { + "type": "STRING", + "value": "customize" + }, + "ASCEND_CANN_PACKAGE_PATH": { + "type": "PATH", + "value": "/usr/local/Ascend/latest" + }, + "ASCEND_PYTHON_EXECUTABLE": { + "type": "STRING", + "value": "python3" + }, + "CMAKE_INSTALL_PREFIX": { + "type": "PATH", + "value": "${sourceDir}/build_out" + }, + "ENABLE_CROSS_COMPILE": { + "type": "BOOL", + "value": "False" + }, + "CMAKE_CROSS_PLATFORM_COMPILER": { + "type": "PATH", + "value": "/usr/bin/aarch64-linux-gnu-g++" + } + } + } + ] +} \ No newline at end of file diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/build.sh b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/build.sh new file mode 100755 index 0000000000000000000000000000000000000000..e3467ea9a7d9db476590cbfe090f6b85c44ead55 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/build.sh @@ -0,0 +1,33 @@ +#!/bin/bash +if [ -z "$BASE_LIBS_PATH" ]; then + if [ -z "$ASCEND_HOME_PATH" ]; then + if [ -z "$ASCEND_AICPU_PATH" ]; then + echo "please set env." + exit 1 + else + export ASCEND_HOME_PATH=$ASCEND_AICPU_PATH + fi + else + export ASCEND_HOME_PATH=$ASCEND_HOME_PATH + fi +else + export ASCEND_HOME_PATH=$BASE_LIBS_PATH +fi +echo "using ASCEND_HOME_PATH: $ASCEND_HOME_PATH" +script_path=$(realpath $(dirname $0)) + +BUILD_DIR="build_out" +HOST_NATIVE_DIR="host_native_tiling" +mkdir -p build_out +rm -rf build_out/* + +ENABLE_CROSS="-DENABLE_CROSS_COMPILE=True" +ENABLE_BINARY="-DENABLE_BINARY_PACKAGE=True" +ENABLE_LIBRARY="-DASCEND_PACK_SHARED_LIBRARY=True" +cmake_version=$(cmake --version | grep "cmake version" | awk '{print $3}') + +target=package +if [ "$1"x != ""x ]; then target=$1; fi + +cmake -S . -B "$BUILD_DIR" --preset=default +cmake --build "$BUILD_DIR" --target $target -j$(nproc) # for package diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/CMakeLists.txt b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/CMakeLists.txt new file mode 100644 index 0000000000000000000000000000000000000000..b6be9b492610f4d45b25bb7725648df9aac39a12 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/CMakeLists.txt @@ -0,0 +1,11 @@ +if(NOT EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/mindspore") + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/caffe_plugin") + add_subdirectory(caffe_plugin) + endif() + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/tf_plugin") + add_subdirectory(tf_plugin) + endif() + if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/onnx_plugin") + add_subdirectory(onnx_plugin) + endif() +endif() diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/tf_plugin/CMakeLists.txt b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/tf_plugin/CMakeLists.txt new file mode 100644 index 0000000000000000000000000000000000000000..ae714680da12efcb4caa0835bbbe18ad6cc9c492 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/tf_plugin/CMakeLists.txt @@ -0,0 +1,10 @@ +aux_source_directory(${CMAKE_CURRENT_SOURCE_DIR} plugin_srcs) + +npu_op_library(cust_tf_parsers TF_PLUGIN + ${plugin_srcs} +) + +npu_op_package_add(${vendor_name} + LIBRARY + cust_tf_parsers +) diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/tf_plugin/tensorflow_add_custom_plugin.cc b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/tf_plugin/tensorflow_add_custom_plugin.cc new file mode 100644 index 0000000000000000000000000000000000000000..bf6498770e1ef38f2b3e1c09de7b910c312257a7 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/framework/tf_plugin/tensorflow_add_custom_plugin.cc @@ -0,0 +1,23 @@ +/* Copyright (C) 2020-2021. Huawei Technologies Co., Ltd. All +rights reserved. + * + * This program is free software; you can redistribute it and/or modify + * it under the terms of the Apache License Version 2.0. + * You may not use this file except in compliance with the License. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * Apache License for more details at + * http://www.apache.org/licenses/LICENSE-2.0 + */ + +#include "register/register.h" + +namespace domi { +// register op info to GE +REGISTER_CUSTOM_OP("AddCustom") + .FrameworkType(TENSORFLOW) // type: CAFFE, TENSORFLOW + .OriginOpType("AddCustom") // name in tf module + .ParseParamsByOperatorFn(AutoMappingByOpFn); +} // namespace domi diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/CMakeLists.txt b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/CMakeLists.txt new file mode 100644 index 0000000000000000000000000000000000000000..fb57d336aa9e7bdc730e259a2922cfb35a17ccf6 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/CMakeLists.txt @@ -0,0 +1,52 @@ +aux_source_directory(${CMAKE_CURRENT_SOURCE_DIR} ops_srcs) + +npu_op_code_gen( + SRC ${ops_srcs} + PACKAGE ${vendor_name} + COMPILE_OPTIONS -g + OUT_DIR ${ASCEND_AUTOGEN_PATH} + OPTIONS + OPS_PRODUCT_NAME ${ASCEND_COMPUTE_UNIT} +) + +file(GLOB autogen_aclnn_src ${ASCEND_AUTOGEN_PATH}/aclnn_*.cpp) +set_source_files_properties(${autogen_aclnn_src} PROPERTIES GENERATED TRUE) +npu_op_library(cust_opapi ACLNN + ${autogen_aclnn_src} +) + +target_compile_options(cust_opapi PRIVATE + -fvisibility=hidden +) + +file(GLOB group_proto_src ${ASCEND_AUTOGEN_PATH}/group_proto/*.cc) +file(GLOB proto_src ${ASCEND_AUTOGEN_PATH}/op_proto.cc) +set_source_files_properties(${group_proto_src} PROPERTIES GENERATED TRUE) +set_source_files_properties(${proto_src} PROPERTIES GENERATED TRUE) +npu_op_library(cust_op_proto GRAPH + ${ops_srcs} + ${group_proto_src} + ${proto_src} +) +target_compile_options(cust_op_proto PRIVATE + -fvisibility=hidden +) + + +file(GLOB fallback_src ${ASCEND_AUTOGEN_PATH}/fallback_*.cpp) +set_source_files_properties(${fallback_src} PROPERTIES GENERATED TRUE) +npu_op_library(cust_optiling TILING + ${ops_srcs} + ${fallback_src} +) +target_compile_options(cust_optiling PRIVATE + -fvisibility=hidden +) + + +npu_op_package_add(${vendor_name} + LIBRARY + cust_opapi + cust_optiling + cust_op_proto +) diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/add_custom.cpp b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/add_custom.cpp new file mode 100644 index 0000000000000000000000000000000000000000..675b79b62643a25b8325ebe3cfdeb7255f5ca512 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/add_custom.cpp @@ -0,0 +1,84 @@ +/* + * Copyright (c) Huawei Technologies Co., Ltd. 2020-2020. All rights reserved. + */ +#include "add_custom_tiling.h" +#include "register/op_def_registry.h" + +namespace optiling { +const uint32_t BLOCK_DIM = 8; +const uint32_t TILE_NUM = 8; +static ge::graphStatus TilingFunc(gert::TilingContext* context) +{ + auto xDtype = context->GetInputDesc(0)->GetDataType(); + uint32_t totalLength = context->GetInputTensor(0)->GetShapeSize(); + context->SetBlockDim(BLOCK_DIM); + if (xDtype == ge::DT_FLOAT16 || xDtype == ge::DT_FLOAT) { + TilingData tiling; + tiling.D_in_tiling.C_in_D.B_in_C.A_in_B.set_totalLength(totalLength); + tiling.D_in_tiling.C_in_D.B_in_C.A_in_B.set_tileNum(TILE_NUM); + tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); + context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); + context->SetTilingKey(1); + } else if (xDtype == ge::DT_INT32) { + TilingDataShort tiling; + tiling.set_totalLength(totalLength); + tiling.set_tileNum(TILE_NUM); + tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); + context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); + context->SetTilingKey(2); + } + size_t *currentWorkspace = context->GetWorkspaceSizes(1); + currentWorkspace[0] = 0; + return ge::GRAPH_SUCCESS; +} +} // namespace optiling + +namespace ge { +static graphStatus InferShape(gert::InferShapeContext* context) +{ + const auto inputShape = context->GetInputShape(0); + auto outputShape = context->GetOutputShape(0); + *outputShape = *inputShape; + return GRAPH_SUCCESS; +} +static graphStatus InferDataType(gert::InferDataTypeContext* context) +{ + auto inputDataType = context->GetInputDataType(0); + context->SetOutputDataType(0, inputDataType); + return ge::GRAPH_SUCCESS; +} +} // namespace ge + +namespace ops { +class AddCustom : public OpDef { +public: + AddCustom(const char* name) : OpDef(name) + { + this->Input("x") + .ParamType(REQUIRED) + .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32}) + .FormatList({ ge::FORMAT_ND}); + this->Input("y") + .ParamType(REQUIRED) + .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32}) + .FormatList({ ge::FORMAT_ND}); + this->Output("z") + .ParamType(REQUIRED) + .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32}) + .FormatList({ ge::FORMAT_ND}); + + this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType); + + this->AICore() + .SetTiling(optiling::TilingFunc); + + this->AICore().AddConfig("ascend910"); + this->AICore().AddConfig("ascend310p"); + this->AICore().AddConfig("ascend910b"); + this->AICore().AddConfig("ascend910_93"); + this->AICore().AddConfig("ascend310b"); + } +}; + +OP_ADD(AddCustom); +} // namespace ops diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/add_custom_tiling.h b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/add_custom_tiling.h new file mode 100644 index 0000000000000000000000000000000000000000..a73f099d31e4a20d02be76cd4e1e827dbe43fb0f --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_host/add_custom_tiling.h @@ -0,0 +1,61 @@ +/* + * Copyright (c) Huawei Technologies Co., Ltd. 2022-2023. All rights reserved. + */ +#ifndef ADD_CUSTOM_TILING_H +#define ADD_CUSTOM_TILING_H +#include "register/tilingdata_base.h" + +namespace optiling { +BEGIN_TILING_DATA_DEF(TilingData_A) + TILING_DATA_FIELD_DEF(uint32_t, A_xxx1); + TILING_DATA_FIELD_DEF(uint32_t, A_xxx2); + TILING_DATA_FIELD_DEF(uint32_t, totalLength); + TILING_DATA_FIELD_DEF(uint32_t, tileNum); +END_TILING_DATA_DEF; +REGISTER_TILING_DATA_CLASS(TilingData_AOp, TilingData_A) + +BEGIN_TILING_DATA_DEF(TilingData_B) + TILING_DATA_FIELD_DEF(uint32_t, B_xxx1); + TILING_DATA_FIELD_DEF(uint32_t, B_xxx2); + TILING_DATA_FIELD_DEF_STRUCT(TilingData_A, A_in_B); +END_TILING_DATA_DEF; +REGISTER_TILING_DATA_CLASS(TilingData_BOp, TilingData_B) + +BEGIN_TILING_DATA_DEF(TilingData_C) + TILING_DATA_FIELD_DEF(uint32_t, C_xxx1); + TILING_DATA_FIELD_DEF(uint32_t, C_xxx2); + TILING_DATA_FIELD_DEF_STRUCT(TilingData_B, B_in_C); +END_TILING_DATA_DEF; +REGISTER_TILING_DATA_CLASS(TilingData_COp, TilingData_C) + +BEGIN_TILING_DATA_DEF(TilingData_E) + TILING_DATA_FIELD_DEF(uint32_t, E_xxx1); + TILING_DATA_FIELD_DEF(uint32_t, E_xxx2); + TILING_DATA_FIELD_DEF_STRUCT(TilingData_B, B_in_E); +END_TILING_DATA_DEF; +REGISTER_TILING_DATA_CLASS(TilingData_EOp, TilingData_E) + +BEGIN_TILING_DATA_DEF(TilingData_D) + TILING_DATA_FIELD_DEF(uint32_t, D_xxx1); + TILING_DATA_FIELD_DEF_STRUCT(TilingData_C, C_in_D); + TILING_DATA_FIELD_DEF(uint32_t, D_xxx2); + TILING_DATA_FIELD_DEF_STRUCT(TilingData_E, E_in_D); +END_TILING_DATA_DEF; +REGISTER_TILING_DATA_CLASS(TilingData_DOp, TilingData_D) + +BEGIN_TILING_DATA_DEF(TilingData) + TILING_DATA_FIELD_DEF(uint32_t, totalLength); + TILING_DATA_FIELD_DEF_STRUCT(TilingData_D, D_in_tiling); + TILING_DATA_FIELD_DEF(uint32_t, tileNum); +END_TILING_DATA_DEF; + +REGISTER_TILING_DATA_CLASS(AddCustom, TilingData) + +BEGIN_TILING_DATA_DEF(TilingDataShort) + TILING_DATA_FIELD_DEF(uint32_t, totalLength); + TILING_DATA_FIELD_DEF(uint32_t, tileNum); +END_TILING_DATA_DEF; + +REGISTER_TILING_DATA_CLASS(AddCustom_2, TilingDataShort) +} +#endif // ADD_CUSTOM_TILING_H diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_kernel/Add/add_custom.cpp b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_kernel/Add/add_custom.cpp new file mode 100644 index 0000000000000000000000000000000000000000..e8073b356dc415bdf0dbf43a1fedbf61c896178a --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_kernel/Add/add_custom.cpp @@ -0,0 +1,95 @@ +/* + * Copyright (c) Huawei Technologies Co., Ltd. 2022-2023. All rights reserved. + * + * Function : z = x + y + * This sample is a very basic sample that implements vector add on Ascend plaform. + */ +#include "kernel_operator.h" +using namespace AscendC; + +constexpr int32_t BUFFER_NUM = 2; + +class KernelAdd { +public: + __aicore__ inline KernelAdd() {} + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum) + { + ASSERT(GetBlockNum() != 0 && "block dim can not be zero!"); + this->blockLength = totalLength / GetBlockNum(); + this->tileNum = tileNum; + ASSERT(tileNum != 0 && "tile num can not be zero!"); + this->tileLength = this->blockLength / tileNum / BUFFER_NUM; + + xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x + this->blockLength * GetBlockIdx(), this->blockLength); + yGm.SetGlobalBuffer((__gm__ DTYPE_Y*)y + this->blockLength * GetBlockIdx(), this->blockLength); + zGm.SetGlobalBuffer((__gm__ DTYPE_Z*)z + this->blockLength * GetBlockIdx(), this->blockLength); + pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X)); + pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y)); + pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Z)); + } + __aicore__ inline void Process() + { + int32_t loopCount = this->tileNum * BUFFER_NUM; + for (int32_t i = 0; i < loopCount; i++) { + CopyIn(i); + Compute(i); + CopyOut(i); + } + } + +private: + __aicore__ inline void CopyIn(int32_t progress) + { + LocalTensor xLocal = inQueueX.AllocTensor(); + LocalTensor yLocal = inQueueY.AllocTensor(); + DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength); + DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength); + inQueueX.EnQue(xLocal); + inQueueY.EnQue(yLocal); + } + __aicore__ inline void Compute(int32_t progress) + { + LocalTensor xLocal = inQueueX.DeQue(); + LocalTensor yLocal = inQueueY.DeQue(); + LocalTensor zLocal = outQueueZ.AllocTensor(); + Add(zLocal, xLocal, yLocal, this->tileLength); + outQueueZ.EnQue(zLocal); + inQueueX.FreeTensor(xLocal); + inQueueY.FreeTensor(yLocal); + } + __aicore__ inline void CopyOut(int32_t progress) + { + LocalTensor zLocal = outQueueZ.DeQue(); + DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength); + outQueueZ.FreeTensor(zLocal); + } + +private: + TPipe pipe; + TQue inQueueX, inQueueY; + TQue outQueueZ; + GlobalTensor xGm; + GlobalTensor yGm; + GlobalTensor zGm; + uint32_t blockLength; + uint32_t tileNum; + uint32_t tileLength; +}; + +extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) +{ + KernelAdd op; + if (TILING_KEY_IS(1)) { + GET_TILING_DATA_WITH_STRUCT(TilingData, tilingData, tiling) + op.Init(x, y, z, tilingData.D_in_tiling.C_in_D.B_in_C.A_in_B.totalLength, tilingData.D_in_tiling.C_in_D.B_in_C.A_in_B.tileNum); + op.Process(); + } else if (TILING_KEY_IS(2)) { + GET_TILING_DATA_WITH_STRUCT(TilingDataShort, tilingData, tiling) + op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum); + op.Process(); + } +} + +void add_custom_do(uint32_t blockDim, void * stream, uint8_t* x, uint8_t *y, uint8_t* z) { + add_custom<<>>(x, y, z); +} diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_kernel/CMakeLists.txt b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_kernel/CMakeLists.txt new file mode 100644 index 0000000000000000000000000000000000000000..f9c5f89ba734a18b4b30e3e1e5c9b5021d748379 --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/op_kernel/CMakeLists.txt @@ -0,0 +1,24 @@ +if ("${CMAKE_BUILD_TYPE}x" STREQUAL "Debugx") + npu_op_kernel_options(ascendc_kernels ALL OPTIONS -gline-tables-only -O0) +endif() +npu_op_kernel_options(ascendc_kernels ALL OPTIONS --save-temp-files --cce-auto-sync=off --impl-mode="") + +npu_op_kernel_sources(ascendc_kernels + OP_NAME AddCustom + KERNEL_DIR ./Add + COMPUTE_UNIT Ascend910B1 Ascend910B2 + KERNEL_FILE add_custom.cpp +) + +npu_op_kernel_sources(ascendc_kernels + KERNEL_DIR ./ +) + +npu_op_kernel_library(ascendc_kernels + SRC_BASE ${CMAKE_SOURCE_DIR}/op_kernel/ + TILING_LIBRARY cust_optiling +) + +npu_op_package_add(${vendor_name} + LIBRARY ascendc_kernels +) \ No newline at end of file diff --git a/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/run.sh b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/run.sh new file mode 100644 index 0000000000000000000000000000000000000000..722fb835aeb16deab7b3ff26c627565e98c31c4b --- /dev/null +++ b/operator/ascendc/0_introduction/1_add_frameworklaunch/CustomOpNew/run.sh @@ -0,0 +1,42 @@ +#!/bin/bash +CURRENT_DIR=$( + cd $(dirname ${BASH_SOURCE:-$0}) + pwd +) +cd $CURRENT_DIR + +# 导出环境变量 +if [ ! $ASCEND_HOME_DIR ]; then + ASCEND_HOME_DIR=/usr/local/Ascend/latest + source $ASCEND_HOME_DIR/bin/setenv.bash +fi +export ASCEND_TENSOR_COMPILER_INCLUDE=$ASCEND_HOME_DIR/compiler/include + +function main { + is_library=$1 + # 1. 构建自定义算子包 + rm -rf build_out + bash build.sh + if [ $? -ne 0 ]; then + echo "ERROR: build custom op run package failed!" + return 1 + fi + + if [ "$is_library"x = ""x ]; then + # 2. 安装自定义算子包 + cd build_out + OS_ID=$(cat /etc/os-release | grep "^ID=" | awk -F= '{print $2}') + OS_ID=$(echo $OS_ID | sed -e 's/^"//' -e 's/"$//') + arch=$(uname -m) + ./custom_opp_${OS_ID}_${arch}.run --quiet + if [ $? -ne 0 ]; then + echo "ERROR: install custom op run package failed!" + return 1 + fi + echo "INFO: install custom op run package success!" + else + echo "INFO: build custom op library success!" + fi +} +IS_LIBRARY=$1 +main $IS_LIBRARY \ No newline at end of file