From cc96c30693c5c4ebc54f192eb6ab6450aa32179b Mon Sep 17 00:00:00 2001 From: wda1991 Date: Wed, 4 Sep 2024 19:35:25 +0800 Subject: [PATCH] remove broadcast --- examples/pad/broadcast/README.md | 67 ---- .../host_tiling/broadcast_custom_tiling.h | 53 --- .../broadcast/kernel_impl/broadcast_custom.h | 77 ----- .../CMakeLists.txt | 81 ----- .../kernel_launch_method_by_direct/README.md | 60 ---- .../broadcast_custom.cpp | 77 ----- .../broadcast_custom_tiling.cpp | 35 -- .../cmake/cpu_lib.cmake | 35 -- .../cmake/npu_lib.cmake | 28 -- .../kernel_launch_method_by_direct/main.cpp | 153 --------- .../kernel_launch_method_by_direct/run.sh | 63 ---- .../scripts/gen_data.py | 37 --- .../CMakeLists.txt | 78 ----- .../CMakePresets.json | 63 ---- .../README.md | 85 ----- .../build.sh | 85 ----- .../cmake/config.cmake | 33 -- .../cmake/func.cmake | 200 ----------- .../cmake/intf.cmake | 34 -- .../cmake/makeself.cmake | 26 -- .../op_host/CMakeLists.txt | 90 ----- .../op_host/broadcast_custom.cpp | 81 ----- .../op_kernel/CMakeLists.txt | 78 ----- .../op_kernel/broadcast_custom.cpp | 50 --- .../scripts/gen_data.py | 37 --- .../scripts/help.info | 1 - .../scripts/install.sh | 313 ------------------ .../scripts/upgrade.sh | 160 --------- .../testcases/CMakeLists.txt | 11 - .../testcases/cmake/fun.cmake | 61 ---- .../testcases/npu/CMakeLists.txt | 19 -- .../testcases/npu/broadcast_custom_main.cpp | 179 ---------- examples/readme.md | 5 - 33 files changed, 2455 deletions(-) delete mode 100644 examples/pad/broadcast/README.md delete mode 100644 examples/pad/broadcast/host_tiling/broadcast_custom_tiling.h delete mode 100644 examples/pad/broadcast/kernel_impl/broadcast_custom.h delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/CMakeLists.txt delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/README.md delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/broadcast_custom.cpp delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/broadcast_custom_tiling.cpp delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/cmake/cpu_lib.cmake delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/cmake/npu_lib.cmake delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/main.cpp delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/run.sh delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_direct/scripts/gen_data.py delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/CMakeLists.txt delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/CMakePresets.json delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/README.md delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/build.sh delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/cmake/config.cmake delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/cmake/func.cmake delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/cmake/intf.cmake delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/cmake/makeself.cmake delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/op_host/CMakeLists.txt delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/op_host/broadcast_custom.cpp delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/op_kernel/broadcast_custom.cpp delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/scripts/gen_data.py delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/scripts/help.info delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/scripts/install.sh delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/scripts/upgrade.sh delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/testcases/CMakeLists.txt delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/testcases/cmake/fun.cmake delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt delete mode 100644 examples/pad/broadcast/kernel_launch_method_by_framework/testcases/npu/broadcast_custom_main.cpp diff --git a/examples/pad/broadcast/README.md b/examples/pad/broadcast/README.md deleted file mode 100644 index 077a2587..00000000 --- a/examples/pad/broadcast/README.md +++ /dev/null @@ -1,67 +0,0 @@ - - -## 概述 - -本样例介绍了调用BroadCast高阶API实现broadcast单算子，并按照不同的算子调用方式分别给出了对应的端到端实现。 - -- Direct：使用核函数直调broadcast自定义算子。 - - 核函数的基础调用方式，开发者完成算子核函数的开发和Tiling实现后，即可通过AscendCL运行时接口，完成算子的调用。 - -- Framework：使用框架调用broadcast自定义算子。 - - 按照工程创建->算子实现->编译部署>算子调用的流程完成算子开发。整个过程都依赖于算子工程：基于工程代码框架完成算子核函数的开发和Tiling实现，通过工程编译脚本完成算子的编译部署，继而实现单算子调用或第三方框架中的算子调用。 - -本样例中包含如下调用方式： - -| 调用方式 | 目录 | **描述** | -| --------- | ------------------------------------------------------------ | ---------------------------------------------------------- | -| Direct | [kernel_launch_method_by_direct](./kernel_launch_method_by_direct) | host侧的核函数调用程序，包含CPU侧、NPU侧、仿真侧三种运行验证方法。 | -| Framework | [kernel_launch_method_by_framework](./kernel_launch_method_by_framework) | 通过aclnn调用的方式调用broadcast算子。 | - -## 样例支持的产品型号为： -- Atlas A2训练系列产品/Atlas 800I A2推理产品 -- Atlas推理系列产品（Ascend 310P处理器）AI Core - -## 目录结构 - -| 目录 | 描述 | -| ------------------------------------------------------------ | ------------------------------------------------------------ | -| [kernel_launch_method_by_direct](./kernel_launch_method_by_direct) | 通过kernel直调的方式调用自定义算子工程样例目录 | -| [kernel_launch_method_by_framework](./kernel_launch_method_by_framework) | 通过aclnn调用的方式调用自定义算子工程样例目录 | -| [host_tiling](./host_tiling) | 本样例tiling代码实现 | -| [kernel_impl](./kernel_impl) | 本样例kernel侧代码实现 | - -## 算子描述 - -broadcast单算子，对输入tensor做广播计算。 - -broadcast算子规格： - - - - - - - - - - - - - -
算子类型(OpType)BroadcastCustom
算子输入nameshapedata typeformat
x-floatND
算子输出y-floatND
核函数名broadcast_custom
- -## 算子实现介绍 - -本样例中实现了两种场景的broadcast算子，分别是[1, 48]到[96, 48]的广播和[96, 1]到[96, 96]的广播。 - -- kernel实现 - - 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用BroadCast高阶API接口完成broadcast计算，得到最终结果，再搬出到外部存储上。 - - broadcast算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor xGm搬运至Local Memory，存储在xLocal中，Compute任务负责对xLocal执行broadcast计算，然后存储在yLocal中，CopyOut任务负责将输出数据从yLocal搬运至Global Memory上的输出Tensor yGm中。 - -- tiling实现 - - broadcast算子的tiling实现流程如下：首先根据shape将广播的轴，input和output的二维shape填充到tiling中，由于样例的shape较小，所以这里只启动了一个核。 \ No newline at end of file diff --git a/examples/pad/broadcast/host_tiling/broadcast_custom_tiling.h b/examples/pad/broadcast/host_tiling/broadcast_custom_tiling.h deleted file mode 100644 index 884fa012..00000000 --- a/examples/pad/broadcast/host_tiling/broadcast_custom_tiling.h +++ /dev/null @@ -1,53 +0,0 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. - * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). - * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. - * See LICENSE in the root of the software repository for the full text of the License. - */ -#ifndef EXAMPLES_PAD_BROADCAST_CUSTOM_TILING_H -#define EXAMPLES_PAD_BROADCAST_CUSTOM_TILING_H -#include "graph/tensor.h" -#include "register/tilingdata_base.h" - -namespace optiling { -BEGIN_TILING_DATA_DEF(BroadcastTilingData) - TILING_DATA_FIELD_DEF(uint32_t, dim); - TILING_DATA_FIELD_DEF(uint32_t, axis); - TILING_DATA_FIELD_DEF(uint32_t, srcFirstDim); - TILING_DATA_FIELD_DEF(uint32_t, srcLastDim); - TILING_DATA_FIELD_DEF(uint32_t, dstFirstDim); - TILING_DATA_FIELD_DEF(uint32_t, dstLastDim); -END_TILING_DATA_DEF; - -REGISTER_TILING_DATA_CLASS(BroadcastCustom, BroadcastTilingData) -} - -void ComputeTiling(const ge::Shape &inputShape, const ge::Shape &outputShape, uint32_t dtypeSize, - optiling::BroadcastTilingData &tiling) -{ - int32_t axis = 0; - const uint32_t dim = inputShape.GetDimNum(); - if (dim == 1) { - tiling.set_srcFirstDim(inputShape.GetDim(0)); - tiling.set_srcLastDim(1); - tiling.set_dstFirstDim(outputShape.GetDim(0)); - tiling.set_dstLastDim(1); - } else { - tiling.set_srcFirstDim(inputShape.GetDim(0)); - tiling.set_srcLastDim(inputShape.GetDim(1)); - tiling.set_dstFirstDim(outputShape.GetDim(0)); - tiling.set_dstLastDim(outputShape.GetDim(1)); - if (inputShape.GetDim(1) == 1) { - axis = 1; - } - } - - tiling.set_axis(axis); - tiling.set_dim(dim); - return; -} - -#endif // EXAMPLES_PAD_BROADCAST_CUSTOM_TILING_H \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_impl/broadcast_custom.h b/examples/pad/broadcast/kernel_impl/broadcast_custom.h deleted file mode 100644 index 6df59bf8..00000000 --- a/examples/pad/broadcast/kernel_impl/broadcast_custom.h +++ /dev/null @@ -1,77 +0,0 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. - * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). - * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. - * See LICENSE in the root of the software repository for the full text of the License. - */ - -#ifndef EXAMPLES_PAD_BROADCAST_CUSTOM_H -#define EXAMPLES_PAD_BROADCAST_CUSTOM_H -#include "kernel_operator.h" - -constexpr int32_t BUFFER_NUM = 1; -template -class KernelBroadcastCustom { -public: - __aicore__ inline KernelBroadcastCustom() - {} - __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t srcLength, uint32_t dstLength, - const uint32_t srcShape[dim], const uint32_t dstShape[dim]) - { - AscendC::AscendCUtils::SetOverflow(1); - xGm.SetGlobalBuffer((__gm__ T *)x, srcLength); - yGm.SetGlobalBuffer((__gm__ T *)y, dstLength); - - pipe.InitBuffer(inQueueX, BUFFER_NUM, srcLength * sizeof(T)); - pipe.InitBuffer(outQueueY, BUFFER_NUM, dstLength * sizeof(T)); - - srcLength_ = srcLength; - dstLength_ = dstLength; - srcShape_ = srcShape; - dstShape_ = dstShape; - } - __aicore__ inline void Process() - { - CopyIn(); - Compute(); - CopyOut(); - } - -private: - __aicore__ inline void CopyIn() - { - AscendC::LocalTensor xLocal = inQueueX.AllocTensor(); - AscendC::DataCopy(xLocal, xGm, srcLength_); - inQueueX.EnQue(xLocal); - } - __aicore__ inline void Compute() - { - AscendC::LocalTensor xLocal = inQueueX.DeQue(); - AscendC::LocalTensor yLocal = outQueueY.AllocTensor(); - AscendC::BroadCast(yLocal, xLocal, dstShape_, srcShape_); - - outQueueY.EnQue(yLocal); - inQueueX.FreeTensor(xLocal); - } - __aicore__ inline void CopyOut() - { - AscendC::LocalTensor yLocal = outQueueY.DeQue(); - AscendC::DataCopy(yGm, yLocal, dstLength_); - outQueueY.FreeTensor(yLocal); - } - -private: - AscendC::TPipe pipe; - AscendC::TQue inQueueX; - AscendC::TQue outQueueY; - AscendC::GlobalTensor xGm; - AscendC::GlobalTensor yGm; - uint32_t srcLength_; - uint32_t dstLength_; - const uint32_t *srcShape_{nullptr}; - const uint32_t *dstShape_{nullptr}; -}; -#endif // EXAMPLES_PAD_BROADCAST_CUSTOM_H \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/CMakeLists.txt b/examples/pad/broadcast/kernel_launch_method_by_direct/CMakeLists.txt deleted file mode 100644 index af343dab..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/CMakeLists.txt +++ /dev/null @@ -1,81 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -cmake_minimum_required(VERSION 3.16) -project(Ascend_c) -if(${RUN_MODE}) - set(RUN_MODE "npu" CACHE STRING "cpu/sim/npu") -endif() -if (${SOC_VERSION}) - set(SOC_VERSION "Ascend910" CACHE STRING "system on chip type") -endif() - -set(ASCEND_CANN_PACKAGE_PATH "~/Ascend/ascend-toolkit/latest" CACHE STRING "ASCEND CANN package installation directory") -if(NOT CMAKE_BUILD_TYPE) - set(CMAKE_BUILD_TYPE "Debug" CACHE STRING "Build type Release/Debug (default Debug)" FORCE) -endif() - -if(CMAKE_INSTALL_PREFIX STREQUAL /usr/local) - set(CMAKE_INSTALL_PREFIX "${CMAKE_CURRENT_LIST_DIR}/out" CACHE STRING "path for install()" FORCE) -endif() - -file(GLOB KERNEL_FILES - ${CMAKE_CURRENT_SOURCE_DIR}/broadcast_custom.cpp -) -set(CUSTOM_ASCEND310P_LIST "Ascend310P1" "Ascend310P3") - -if("${RUN_MODE}" STREQUAL "cpu") - include(cmake/cpu_lib.cmake) -elseif("${RUN_MODE}" STREQUAL "sim" OR "${RUN_MODE}" STREQUAL "npu") - include(cmake/npu_lib.cmake) -else() - message("invalid RUN_MODE: ${RUN_MODE}") -endif() - -add_executable(broadcast_direct_kernel_op - ${CMAKE_CURRENT_SOURCE_DIR}/main.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/broadcast_custom_tiling.cpp -) - -target_compile_options(broadcast_direct_kernel_op PRIVATE - $:-g>> - -O2 - -std=c++17 - -D_GLIBCXX_USE_CXX11_ABI=0 -) - -target_compile_definitions(broadcast_direct_kernel_op PRIVATE - $<$>:CUSTOM_ASCEND310P> -) - -target_include_directories(broadcast_direct_kernel_op PRIVATE - ${CMAKE_CURRENT_SOURCE_DIR} - $:${ASCEND_CANN_PACKAGE_PATH}/include>> - $:${ASCEND_CANN_PACKAGE_PATH}/runtime/include>> -) - -target_link_libraries(broadcast_direct_kernel_op PRIVATE - $,$>:host_intf_pub>> - $:tikicpulib::${SOC_VERSION}>> - $:ascendcl>> - $:c_sec>> - ascendc_kernels_${RUN_MODE} - tiling_api - register - platform - ascendalog - dl - graph_base -) - -install(TARGETS broadcast_direct_kernel_op - LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} - ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} - RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} -) \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/README.md b/examples/pad/broadcast/kernel_launch_method_by_direct/README.md deleted file mode 100644 index e565510f..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/README.md +++ /dev/null @@ -1,60 +0,0 @@ - - -## 概述 - -本样例基于Kernel直调算子工程，介绍了调用BroadCast高阶API实现broadcast单算子，主要演示BroadCast高阶API在Kernel直调工程中的调用。 - - -## 目录结构介绍 -| 目录及文件 | 描述 | -|---------------------|----------------------| -| [cmake](./cmake) | 编译工程文件 | -| [scripts](./scripts) | 包含输入数据和真值数据生成脚本文件 | -| main.cpp | 主函数，调用算子的应用程序，含CPU域及NPU域调用 | -| broadcast_custom.cpp | 算子kernel实现 | -| broadcast_custom_tiling.cpp | 算子tiling实现 | -| CMakeLists.txt | 编译工程文件 | -| run.sh | 编译执行脚本 | - -## 编译运行样例 - - - 打开样例目录 - - ``` - cd examples/pad/broadcast/kernel_launch_method_by_direct - ``` - - 配置环境变量 - - 这里的\$ASCEND_CANN_PACKAGE_PATH需要替换为CANN包的存储路径。例如:/usr/local/Ascend/ascend-toolkit/latest - ``` - export ASCEND_HOME_DIR=$ASCEND_CANN_PACKAGE_PATH - ``` - 若执行sim仿真，可自行配置仿真日志文件目录，默认仿真日志会在build目录下生成。若需要详细了解sim仿真相关内容，请参考[《AscendC算子调测工具》](https://hiascend.com/document/redirect/CannCommunityToolAscendebug)中的 调测功能说明 > Simulator性能仿真功能 > CAModel性能仿真 章节。 - ``` - # 设置仿真模式日志生成目录（可选），需要自行确保设置的目录已存在。若设置为相对路径下的目录，则以程序执行时的目录作为当前目录。例如，执行如下设置时，需要确保./目录下存在xxx目录 - export CAMODEL_LOG_PATH=./xxx - ``` - - - 生成输入和真值 - ``` - python3 scripts/gen_data.py --testcase=[TEST_CASE] - ``` - 其中参数说明如下： - - TEST_CASE ：执行用例id，支持参数为[0/1]， 用例0实现了从[1, 48] -> [96, 48]的广播，用例1实现了从[96, 1] -> [96, 96]的广播 - - - 编译执行 - - ``` - bash run.sh -r [RUN_MODE] -v [SOC_VERSION] -t [TEST_CASE] - ``` - 其中cmake参数说明如下： - - RUN_MODE ：编译方式，可选择CPU调试，NPU仿真，NPU上板。支持参数为[cpu / sim/ npu] - - SOC_VERSION ：昇腾AI处理器型号，如果无法确定具体的[SOC_VERSION]，则在安装昇腾AI处理器的服务器执行npu-smi info命令进行查询，在查询到的“Name”前增加Ascend信息，例如“Name”对应取值为xxxyy，实际配置的[SOC_VERSION]值为Ascendxxxyy。支持以下参数取值（xxx请替换为具体取值）： - - Atlas 推理系列产品（Ascend 310P处理器）参数值：Ascend310P1、Ascend310P3 - - Atlas A2训练系列产品参数值：AscendxxxB1、AscendxxxB2、AscendxxxB3、AscendxxxB4 - - TEST_CASE ：执行用例id，支持参数为[0/1]， 用例0实现了从[1, 48] -> [96, 48]的广播，用例1实现了从[96, 1] -> [96, 96]的广播 - - 示例如下： - ``` - bash run.sh -r cpu -v Ascend310P1 -t 0 - ``` diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/broadcast_custom.cpp b/examples/pad/broadcast/kernel_launch_method_by_direct/broadcast_custom.cpp deleted file mode 100644 index 395ac034..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/broadcast_custom.cpp +++ /dev/null @@ -1,77 +0,0 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. - * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). - * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. - * See LICENSE in the root of the software repository for the full text of the License. - */ -#include "../kernel_impl/broadcast_custom.h" -struct BroadcastTilingData { - uint32_t dim{0}; - uint32_t axis{0}; - uint32_t srcFirstDim{0}; - uint32_t srcLastDim{0}; - uint32_t dstFirstDim{0}; - uint32_t dstLastDim{0}; -}; - -__aicore__ inline void CopyTiling(BroadcastTilingData *tiling, GM_ADDR tilingGM) -{ - uint32_t *ptr = reinterpret_cast(tiling); - auto tiling32 = reinterpret_cast<__gm__ uint32_t *>(tilingGM); - - for (int i = 0; i < sizeof(BroadcastTilingData) / sizeof(uint32_t); i++, ptr++) { - *ptr = *(tiling32 + i); - } - return; -} - -extern "C" __global__ __aicore__ void broadcast_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) -{ - BroadcastTilingData tilingData; - CopyTiling(&tilingData, tiling); - uint32_t axis = tilingData.axis; - uint32_t dim = tilingData.dim; - - if (dim == 1) { - const uint32_t srcShape[] = {tilingData.srcFirstDim}; - const uint32_t dstShape[] = {tilingData.dstFirstDim}; - KernelBroadcastCustom op; - op.Init(x, y, tilingData.srcFirstDim, tilingData.dstFirstDim, srcShape, dstShape); - op.Process(); - } else { - const uint32_t srcShape[] = {tilingData.srcFirstDim, tilingData.srcLastDim}; - const uint32_t dstShape[] = {tilingData.dstFirstDim, tilingData.dstLastDim}; - - if (axis == 0) { - KernelBroadcastCustom op; - op.Init(x, - y, - tilingData.srcFirstDim * tilingData.srcLastDim, - tilingData.dstFirstDim * tilingData.dstLastDim, - srcShape, - dstShape); - op.Process(); - } else { - KernelBroadcastCustom op; - op.Init(x, - y, - tilingData.srcFirstDim * tilingData.srcLastDim, - tilingData.dstFirstDim * tilingData.dstLastDim, - srcShape, - dstShape); - op.Process(); - } - } -} - -#ifndef ASCENDC_CPU_DEBUG -// call of kernel function -void broadcast_custom_do( - uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, uint8_t *workspace, uint8_t *tiling) -{ - broadcast_custom<<>>(x, y, workspace, tiling); -} -#endif \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/broadcast_custom_tiling.cpp b/examples/pad/broadcast/kernel_launch_method_by_direct/broadcast_custom_tiling.cpp deleted file mode 100644 index 0bb877a6..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/broadcast_custom_tiling.cpp +++ /dev/null @@ -1,35 +0,0 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. - * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). - * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. - * See LICENSE in the root of the software repository for the full text of the License. - */ - -#include -#include -#include -#include -#include - -#include "tiling/tiling_api.h" -#include "../host_tiling/broadcast_custom_tiling.h" - -using namespace std; - -uint8_t* GetTilingBuf(optiling::BroadcastTilingData* tilingData) -{ - uint32_t tilingSize = sizeof(optiling::BroadcastTilingData); - uint8_t* buf = (uint8_t*)malloc(tilingSize); - tilingData->SaveToBuffer(buf, tilingSize); - return buf; -} - -uint8_t* GenerateTiling(const ge::Shape &inputShape, const ge::Shape &outputShape, uint32_t dtypeSize) -{ - optiling::BroadcastTilingData tiling; - ComputeTiling(inputShape, outputShape, dtypeSize, tiling); - return GetTilingBuf(&tiling); -} diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/cmake/cpu_lib.cmake b/examples/pad/broadcast/kernel_launch_method_by_direct/cmake/cpu_lib.cmake deleted file mode 100644 index 7799cb22..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/cmake/cpu_lib.cmake +++ /dev/null @@ -1,35 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -if(NOT DEFINED ENV{CMAKE_PREFIX_PATH}) - set(CMAKE_PREFIX_PATH ${ASCEND_CANN_PACKAGE_PATH}/tools/tikicpulib/lib/cmake) -endif() -find_package(tikicpulib REQUIRED) - -add_library(ascendc_kernels_${RUN_MODE} SHARED - ${KERNEL_FILES} -) - -target_link_libraries(ascendc_kernels_${RUN_MODE} PRIVATE - tikicpulib::${SOC_VERSION} -) - -target_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE - $<$>:CUSTOM_ASCEND310P> -) - -target_compile_options(ascendc_kernels_${RUN_MODE} PRIVATE - -g - -O0 - -std=c++17 -) - -install(TARGETS ascendc_kernels_${RUN_MODE} -DESTINATION ${CMAKE_INSTALL_LIBDIR} -) \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/cmake/npu_lib.cmake b/examples/pad/broadcast/kernel_launch_method_by_direct/cmake/npu_lib.cmake deleted file mode 100644 index da74f37a..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/cmake/npu_lib.cmake +++ /dev/null @@ -1,28 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -if(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/compiler/tikcpp/ascendc_kernel_cmake) -elseif(EXISTS ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) - set(ASCENDC_CMAKE_DIR ${ASCEND_CANN_PACKAGE_PATH}/tools/tikcpp/ascendc_kernel_cmake) -else() - message(FATAL_ERROR "ascendc_kernel_cmake does not exist ,please check whether the cann package is installed") -endif() -include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) - -ascendc_library(ascendc_kernels_${RUN_MODE} STATIC - ${KERNEL_FILES} -) - -ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE - $<$>:CUSTOM_ASCEND310P> - -DASCENDC_DUMP - -DHAVE_WORKSPACE - -DHAVE_TILING - ) \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/main.cpp b/examples/pad/broadcast/kernel_launch_method_by_direct/main.cpp deleted file mode 100644 index 8c2c646f..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/main.cpp +++ /dev/null @@ -1,153 +0,0 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. - * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). - * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. - * See LICENSE in the root of the software repository for the full text of the License. - */ -#include "../../../common/data_utils.h" -#include "graph/tensor.h" -#include -#include -#ifndef ASCENDC_CPU_DEBUG -#include "acl/acl.h" -extern void broadcast_custom_do(uint32_t coreDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, uint8_t *workspace, uint8_t *tiling); -#else -#include "tikicpulib.h" -extern "C" __global__ __aicore__ void broadcast_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling); -#endif - -constexpr uint32_t TILINGDATA_SIZE = 6; - -int64_t CompareResult(void* outputData, int64_t outSize) -{ - void* goldenData; -#ifdef ASCENDC_CPU_DEBUG - goldenData = (uint8_t*)AscendC::GmAlloc(outSize); -#else - CHECK_ACL(aclrtMallocHost((void**)(&goldenData), outSize)); -#endif - size_t goldenSize = outSize; - bool ret = ReadFile("../output/golden.bin", goldenSize, goldenData, goldenSize); - if (ret) { - printf("ReadFile golden success!\n"); - } else { - return -1; - } - constexpr float EPS = 1e-5; - int64_t wrongNum = 0; - - for (int i = 0; i < outSize / sizeof(float); i++) { - float a = ((float*)outputData)[i]; - float b = ((float*)goldenData)[i]; - float ae = std::abs(a - b); - float re = ae / abs(b); - if(ae > EPS && re > EPS) { - printf("CompareResult failed output is %lf, golden is %lf\n",a,b); - wrongNum++; - } - } -#ifdef ASCENDC_CPU_DEBUG - AscendC::GmFree((void*)goldenData); -#else - CHECK_ACL(aclrtFreeHost(goldenData)); -#endif - return wrongNum; -} - -uint8_t* GenerateTiling(const ge::Shape &inputShape, const ge::Shape &outputShape, uint32_t dtypeSize); - -int32_t main(int32_t argc, char *argv[]) -{ - uint32_t blockDim = 1; - size_t inputSize = 48 * sizeof(float); - size_t outputSize = 96 * 48 * sizeof(float); - std::vector inputDims = {1, 48}; - std::vector outputDims = {96, 48}; - if(argc == 2) { - if(!strcmp(argv[1], "1")) { - inputSize = 96 * sizeof(float); - outputSize = 96 * 96 * sizeof(float); - inputDims = {96, 1}; - outputDims = {96, 96}; - } - } - ge::Shape inputShape(inputDims); - ge::Shape outputShape(outputDims); - size_t tilingSize = TILINGDATA_SIZE * sizeof(uint32_t); // tilingData size , defined in broadcast_custom_tiling.h - -#ifdef ASCENDC_CPU_DEBUG - uint8_t *x = (uint8_t *)AscendC::GmAlloc(inputSize); - uint8_t *y = (uint8_t *)AscendC::GmAlloc(outputSize); - uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingSize); - - ReadFile("../input/input.bin", inputSize, x, inputSize); - memcpy_s(tiling, tilingSize, GenerateTiling(inputShape, outputShape, sizeof(float)), tilingSize); - AscendC::SetKernelMode(KernelMode::AIV_MODE); - ICPU_RUN_KF(broadcast_custom, blockDim, x, y, nullptr, tiling); // use this macro for cpu debug - - WriteFile("../output/output.bin", y, outputSize); - - int64_t wrongNum = CompareResult(y, outputSize); - if (wrongNum != 0) { - printf("test failed!\n"); - } else { - printf("test pass!\n"); - } - - AscendC::GmFree((void *)x); - AscendC::GmFree((void *)tiling); - AscendC::GmFree((void *)y); -#else - CHECK_ACL(aclInit(nullptr)); - aclrtContext context; - int32_t deviceId = 0; - CHECK_ACL(aclrtSetDevice(deviceId)); - CHECK_ACL(aclrtCreateContext(&context, deviceId)); - aclrtStream stream = nullptr; - CHECK_ACL(aclrtCreateStream(&stream)); - - uint8_t *xHost, *yHost, *tilingHost; - uint8_t *xDevice, *yDevice, *tilingDevice; - - CHECK_ACL(aclrtMallocHost((void **)(&xHost), inputSize)); - CHECK_ACL(aclrtMallocHost((void **)(&yHost), outputSize)); - CHECK_ACL(aclrtMallocHost((void**)(&tilingHost), tilingSize)); - CHECK_ACL(aclrtMalloc((void **)&xDevice, inputSize, ACL_MEM_MALLOC_HUGE_FIRST)); - CHECK_ACL(aclrtMalloc((void **)&yDevice, outputSize, ACL_MEM_MALLOC_HUGE_FIRST)); - CHECK_ACL(aclrtMalloc((void**)&tilingDevice, tilingSize, ACL_MEM_MALLOC_HUGE_FIRST)); - - ReadFile("../input/input.bin", inputSize, xHost, inputSize); - - CHECK_ACL(aclrtMemcpy(xDevice, inputSize, xHost, inputSize, ACL_MEMCPY_HOST_TO_DEVICE)); - CHECK_ACL(aclrtMemcpy(tilingDevice, tilingSize, GenerateTiling(inputShape, outputShape, sizeof(float)), tilingSize, ACL_MEMCPY_HOST_TO_DEVICE)); - - broadcast_custom_do(blockDim, nullptr, stream, xDevice, yDevice, nullptr, tilingDevice); - CHECK_ACL(aclrtSynchronizeStream(stream)); - - CHECK_ACL(aclrtMemcpy(yHost, outputSize, yDevice, outputSize, ACL_MEMCPY_DEVICE_TO_HOST)); - WriteFile("../output/output.bin", yHost, outputSize); - - int64_t wrongNum = CompareResult(yHost, outputSize); - if (wrongNum != 0) { - printf("test failed!\n"); - } else { - printf("test pass!\n"); - } - - CHECK_ACL(aclrtFree(xDevice)); - CHECK_ACL(aclrtFree(yDevice)); - CHECK_ACL(aclrtFree(tilingDevice)); - CHECK_ACL(aclrtFreeHost(xHost)); - CHECK_ACL(aclrtFreeHost(yHost)); - CHECK_ACL(aclrtFreeHost(tilingHost)); - - CHECK_ACL(aclrtDestroyStream(stream)); - CHECK_ACL(aclrtDestroyContext(context)); - CHECK_ACL(aclrtResetDevice(deviceId)); - CHECK_ACL(aclFinalize()); -#endif - return 0; -} \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/run.sh b/examples/pad/broadcast/kernel_launch_method_by_direct/run.sh deleted file mode 100644 index 8e5d4503..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/run.sh +++ /dev/null @@ -1,63 +0,0 @@ -#!/bin/bash -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -SHORT=r:,v:, -LONG=run-mode:,soc-version:,testcase:, -OPTS=$(getopt -a --options $SHORT --longoptions $LONG -- "$@") -eval set -- "$OPTS" -while : -do - case "$1" in - (-r | --run-mode ) - RUN_MODE="$2" - shift 2;; - (-v | --soc-version ) - SOC_VERSION="$2" - shift 2;; - (-t | --testcase ) - TEST_CASE="$2" - shift 2;; - (--) - shift; - break;; - (*) - echo "[ERROR] Unexpected option: $1"; - break;; - esac -done - -rm -rf build -mkdir build -cd build - -# in case of running op in simulator, use stub so instead -if [ "${RUN_MODE}" = "sim" ]; then - export LD_LIBRARY_PATH=$(echo $LD_LIBRARY_PATH | sed 's/\/.*\/runtime\/lib64://g') - export LD_LIBRARY_PATH=$ASCEND_HOME_DIR/runtime/lib64/stub:$LD_LIBRARY_PATH - - if [ ! $CAMODEL_LOG_PATH ]; then - export CAMODEL_LOG_PATH=./ # default log save in build dir - else - export CAMODEL_LOG_PATH=../$CAMODEL_LOG_PATH - rm -rf $CAMODEL_LOG_PATH - mkdir -p $CAMODEL_LOG_PATH - fi -fi - -if [ "${RUN_MODE}" = "cpu" ]; then - export CAMODEL_LOG_PATH=./ # cpu run mode set fixed log path -fi - -source $ASCEND_HOME_DIR/bin/setenv.bash -export LD_LIBRARY_PATH=${ASCEND_HOME_DIR}/tools/simulator/${SOC_VERSION}/lib:$LD_LIBRARY_PATH - -cmake -DRUN_MODE=${RUN_MODE} -DSOC_VERSION=${SOC_VERSION} -DASCEND_CANN_PACKAGE_PATH=${ASCEND_HOME_DIR} .. -make -j16 -./broadcast_direct_kernel_op ${TEST_CASE} \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_direct/scripts/gen_data.py b/examples/pad/broadcast/kernel_launch_method_by_direct/scripts/gen_data.py deleted file mode 100644 index 9aba79bb..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_direct/scripts/gen_data.py +++ /dev/null @@ -1,37 +0,0 @@ -#!/usr/bin/python3 -# coding=utf-8 - -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -import numpy as np -import os -import argparse - -def gen_golden_data_simple(testcase): - x_shape = (1, 48) - y_shape = (96, 48) - if testcase == 1: - x_shape = (96, 1) - y_shape = (96, 96) - x = np.random.uniform(-10, 10, x_shape).astype(np.float32) - y = np.broadcast_to(x, y_shape).astype(np.float32) - os.system("mkdir -p input") - os.system("mkdir -p output") - x.tofile("./input/input.bin") - y.tofile("./output/golden.bin") - -if __name__ == "__main__": - parser = argparse.ArgumentParser() - parser.add_argument('--testcase', type=int, help='testcase,目前只有case 0 与 1') - args = parser.parse_args() - testcase = 0 - if args.testcase: - testcase = args.testcase - gen_golden_data_simple(testcase) \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/CMakeLists.txt b/examples/pad/broadcast/kernel_launch_method_by_framework/CMakeLists.txt deleted file mode 100644 index 4f0fdd44..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/CMakeLists.txt +++ /dev/null @@ -1,78 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -cmake_minimum_required(VERSION 3.16.0) -project(opp) -if(ENABLE_CROSS_COMPILE) - if(${CMAKE_SYSTEM_PROCESSOR} STREQUAL x86_64) - set(CROSS_COMPILE_PLATFORM aarch64) - else() - set(CROSS_COMPILE_PLATFORM x86_64) - endif() - set(PLATFORM ${CMAKE_SYSTEM_PROCESSOR}) - set(CMAKE_COMPILE_COMPILER_LIBRARY ${ASCEND_CANN_PACKAGE_PATH}/${PLATFORM}-linux/devlib/linux/${CROSS_COMPILE_PLATFORM}/) - set(CMAKE_COMPILE_RUNTIME_LIBRARY ${ASCEND_CANN_PACKAGE_PATH}/${PLATFORM}-linux/devlib/${CROSS_COMPILE_PLATFORM}/) - set(CMAKE_SYSTEM_PROCESSOR ${CROSS_COMPILE_PLATFORM}) - set(CMAKE_COMPILE ${CMAKE_CXX_COMPILER}) - set(CMAKE_CXX_COMPILER ${CMAKE_CROSS_PLATFORM_COMPILER}) -else() - set(CMAKE_COMPILE ${CMAKE_CXX_COMPILER}) -endif() - -include(cmake/config.cmake) -include(cmake/func.cmake) -include(cmake/intf.cmake) - -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() -if(ENABLE_TEST AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/testcases) - add_subdirectory(testcases) -endif() - -# modify vendor_name in install.sh and upgrade.sh -add_custom_command(OUTPUT ${CMAKE_BINARY_DIR}/scripts/install.sh ${CMAKE_BINARY_DIR}/scripts/upgrade.sh - COMMAND mkdir -p ${CMAKE_BINARY_DIR}/scripts - COMMAND cp -r ${CMAKE_SOURCE_DIR}/scripts/* ${CMAKE_BINARY_DIR}/scripts/ - COMMAND sed -i "s/vendor_name=customize/vendor_name=${vendor_name}/g" ${CMAKE_BINARY_DIR}/scripts/* -) -add_custom_target(modify_vendor ALL DEPENDS ${CMAKE_BINARY_DIR}/scripts/install.sh ${CMAKE_BINARY_DIR}/scripts/upgrade.sh) -install(DIRECTORY ${CMAKE_BINARY_DIR}/scripts/ DESTINATION . FILE_PERMISSIONS OWNER_EXECUTE OWNER_READ GROUP_READ) - -install(FILES ${CMAKE_SOURCE_DIR}/custom.proto DESTINATION packages OPTIONAL) - -get_system_info(SYSTEM_INFO) - -# gen version.info -add_custom_target(gen_version_info ALL - COMMAND bash ${CMAKE_CURRENT_SOURCE_DIR}/cmake/util/gen_version_info.sh ${ASCEND_CANN_PACKAGE_PATH} ${CMAKE_CURRENT_BINARY_DIR} -) - -install(FILES ${CMAKE_CURRENT_BINARY_DIR}/version.info - DESTINATION packages/vendors/${vendor_name}/) - -# CPack config -set(CPACK_PACKAGE_NAME ${CMAKE_PROJECT_NAME}) -set(CPACK_PACKAGE_VERSION ${CMAKE_PROJECT_VERSION}) -set(CPACK_PACKAGE_DESCRIPTION "CPack opp project") -set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "CPack opp project") -set(CPACK_PACKAGE_DIRECTORY ${CMAKE_INSTALL_PREFIX}) -set(CPACK_PACKAGE_FILE_NAME "custom_opp_${SYSTEM_INFO}.run") -set(CPACK_GENERATOR External) -set(CPACK_CMAKE_GENERATOR "Unix Makefiles") -set(CPACK_EXTERNAL_ENABLE_STAGING TRUE) -set(CPACK_EXTERNAL_PACKAGE_SCRIPT ${CMAKE_SOURCE_DIR}/cmake/makeself.cmake) -set(CPACK_EXTERNAL_BUILT_PACKAGES ${CPACK_PACKAGE_DIRECTORY}/_CPack_Packages/Linux/External/${CPACK_PACKAGE_FILE_NAME}/${CPACK_PACKAGE_FILE_NAME}) -include(CPack) diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/CMakePresets.json b/examples/pad/broadcast/kernel_launch_method_by_framework/CMakePresets.json deleted file mode 100644 index e56e9011..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/CMakePresets.json +++ /dev/null @@ -1,63 +0,0 @@ -{ - "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": "ascend310p;ascend910b" - }, - "ENABLE_TEST": { - "type": "BOOL", - "value": "True" - }, - "vendor_name": { - "type": "STRING", - "value": "customize" - }, - "ASCEND_CANN_PACKAGE_PATH": { - "type": "PATH", - "value": "~/Ascend/ascend-toolkit/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/examples/pad/broadcast/kernel_launch_method_by_framework/README.md b/examples/pad/broadcast/kernel_launch_method_by_framework/README.md deleted file mode 100644 index 561f6ebb..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/README.md +++ /dev/null @@ -1,85 +0,0 @@ - - -## 概述 - -本样例基于自定义算子工程，介绍了调用BroadCast高阶API实现broadcast单算子，主要演示BroadCast高阶API在自定义算子工程中的调用。 - - -## 目录结构 -| 目录 | 描述 | -|---------------------|----------------------| -| [cmake](./cmake) | 编译工程文件 | -| [op_host](./op_host) | 包含算子tiling实现 | -| [op_kernel](./op_kernel) | 包含算子kernel实现 | -| [scripts](./scripts) | 存放生成的输入数据脚本及安装算子包工程文件目录 | -| [testcases](./testcases) | aclnn调用实现目录 | -| build.sh | 编译算子的脚本 | -| CMakeLists.txt | 编译工程文件 | -| CMakePresets.json | 编译工程配置文件 | - -## 编译运行样例 - -### 1.配置环境变量 - - 这里的\$ASCEND_CANN_PACKAGE_PATH需要替换为CANN包的存储路径。eg:/usr/local/Ascend/ascend-toolkit/latest - ``` - export ASCEND_HOME_DIR=$ASCEND_CANN_PACKAGE_PATH - source $ASCEND_HOME_DIR/bin/setenv.bash - ``` -### 2.生成输入和真值 - ``` - python3 scripts/gen_data.py --testcase=[TEST_CASE] - ``` - 其中参数说明如下： - - TEST_CASE ：执行用例id，支持参数为[0/1]， 用例0实现了从[1, 48] -> [96, 48]的广播，用例1实现了从[96, 1] -> [96, 96]的广播 -### 3.编译算子工程 - - - 修改CMakePresets.json中ASCEND_CANN_PACKAGE_PATH为CANN软件包安装后的实际路径。 - - - ``` - { - …… - "configurePresets": [ - { - …… - "ASCEND_CANN_PACKAGE_PATH": { - "type": "PATH", - "value": "~/Ascend/ascend-toolkit/latest" //请替换为CANN软件包安装后的实际路径。eg:/home/HwHiAiUser/Ascend/ascend-toolkit/latest - }, - …… - } - ] - } - ``` - - 在当前算子工程目录下执行如下命令，进行算子工程编译。 - - ``` - bash build.sh - ``` - 编译成功后，会在当前目录下创建build_out目录，并在build_out目录下生成自定义算子安装包custom_opp_\_\.run，例如“custom_opp_ubuntu_x86_64.run”。 - - -### 4.部署算子包 - - - 执行如下命令，在自定义算子安装包所在路径下，安装自定义算子包。 - - ``` - cd build_out - ./custom_opp__.run - ``` - - 命令执行成功后，自定义算子包中的相关文件将部署至当前环境的OPP算子库的vendors/customize目录中。 - -### 5.执行样例 - - 在build_out目录下执行如下命令 - - ``` - ./broadcast_custom_npu [TEST_CASE] - ``` - - 其中参数说明如下： - - TEST_CASE ： ：执行用例id，支持参数为[0/1]， 用例0实现了从[1, 48] -> [96, 48]的广播，用例1实现了从[96, 1] -> [96, 96]的广播 - -### 注意事项 -本样例工程会自动识别执行的硬件平台，无需单独设置SOC_VERSION \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/build.sh b/examples/pad/broadcast/kernel_launch_method_by_framework/build.sh deleted file mode 100644 index e23a65b2..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/build.sh +++ /dev/null @@ -1,85 +0,0 @@ -#!/bin/bash -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -script_path=$(realpath $(dirname $0)) - -source $ASCEND_HOME_DIR/bin/setenv.bash -cp -rf ../host_tiling/* op_host/ -cp -rf $ASCEND_HOME_DIR/python/site-packages/op_gen/template/ascendc/cmake/ ./ -mkdir -p build_out -rm -rf build_out/* -cd build_out - -opts=$(python3 $script_path/cmake/util/preset_parse.py $script_path/CMakePresets.json) -ENABLE_CROSS="-DENABLE_CROSS_COMPILE=True" -ENABLE_BINARY="-DENABLE_BINARY_PACKAGE=True" -cmake_version=$(cmake --version | grep "cmake version" | awk '{print $3}') - -cmake_run_package() -{ - target=$1 - cmake --build . --target $target -j16 - if [ $? -ne 0 ]; then exit 1; fi - - if [ $target = "package" ]; then - if test -d ./op_kernel/binary ; then - ./cust*.run - if [ $? -ne 0 ]; then exit 1; fi - cmake --build . --target binary -j16 - if [ $? -ne 0 ]; then exit 1; fi - cmake --build . --target $target -j16 - fi - fi -} - -if [[ $opts =~ $ENABLE_CROSS ]] && [[ $opts =~ $ENABLE_BINARY ]] -then - target=package - if [ "$1"x != ""x ]; then target=$1; fi - if [ "$cmake_version" \< "3.19.0" ] ; then - cmake .. $opts -DENABLE_CROSS_COMPILE=0 - else - cmake .. --preset=default -DENABLE_CROSS_COMPILE=0 - fi - cmake_run_package $target - cp -r kernel ../ - rm -rf * - if [ "$cmake_version" \< "3.19.0" ] ; then - cmake .. $opts - else - cmake .. --preset=default - fi - - cmake --build . --target $target -j16 - if [ $? -ne 0 ]; then exit 1; fi - if [ $target = "package" ]; then - if test -d ./op_kernel/binary ; then - ./cust*.run - fi - fi - rm -rf ../kernel - -else - target=package - if [ "$1"x != ""x ]; then target=$1; fi - if [ "$cmake_version" \< "3.19.0" ] ; then - cmake .. $opts - else - cmake .. --preset=default - fi - cmake_run_package $target -fi - - -# for debug -# cd build_out -# make -# cpack -# verbose append -v \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/config.cmake b/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/config.cmake deleted file mode 100644 index d2195cf9..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/config.cmake +++ /dev/null @@ -1,33 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -set(CMAKE_CXX_FLAGS_DEBUG "") -set(CMAKE_CXX_FLAGS_RELEASE "") - -if (NOT DEFINED vendor_name) - set(vendor_name customize CACHE STRING "") -endif() -if (NOT DEFINED ASCEND_CANN_PACKAGE_PATH) - set(ASCEND_CANN_PACKAGE_PATH /usr/local/Ascend/latest CACHE PATH "") -endif() -if (NOT DEFINED ASCEND_PYTHON_EXECUTABLE) - set(ASCEND_PYTHON_EXECUTABLE python3 CACHE STRING "") -endif() -if (NOT DEFINED ASCEND_COMPUTE_UNIT) - message(FATAL_ERROR "ASCEND_COMPUTE_UNIT not set in CMakePreset.json ! -") -endif() -set(ASCEND_TENSOR_COMPILER_PATH ${ASCEND_CANN_PACKAGE_PATH}/compiler) -set(ASCEND_CCEC_COMPILER_PATH ${ASCEND_TENSOR_COMPILER_PATH}/ccec_compiler/bin) -set(ASCEND_AUTOGEN_PATH ${CMAKE_BINARY_DIR}/autogen) -set(ASCEND_FRAMEWORK_TYPE tensorflow) -file(MAKE_DIRECTORY ${ASCEND_AUTOGEN_PATH}) -set(CUSTOM_COMPILE_OPTIONS "custom_compile_options.ini") -execute_process(COMMAND rm -rf ${ASCEND_AUTOGEN_PATH}/${CUSTOM_COMPILE_OPTIONS} - COMMAND touch ${ASCEND_AUTOGEN_PATH}/${CUSTOM_COMPILE_OPTIONS}) diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/func.cmake b/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/func.cmake deleted file mode 100644 index 6301f39e..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/func.cmake +++ /dev/null @@ -1,200 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -function(get_system_info SYSTEM_INFO) - if (UNIX) - execute_process(COMMAND grep -i ^id= /etc/os-release OUTPUT_VARIABLE TEMP) - string(REGEX REPLACE "\n|id=|ID=|\"" "" SYSTEM_NAME ${TEMP}) - set(${SYSTEM_INFO} ${SYSTEM_NAME}_${CMAKE_SYSTEM_PROCESSOR} PARENT_SCOPE) - elseif (WIN32) - message(STATUS "System is Windows. Only for pre-build.") - else () - message(FATAL_ERROR "${CMAKE_SYSTEM_NAME} not support.") - endif () -endfunction() - -function(opbuild) - message(STATUS "Opbuild generating sources") - cmake_parse_arguments(OPBUILD "" "OUT_DIR;PROJECT_NAME;ACCESS_PREFIX" "OPS_SRC" ${ARGN}) - execute_process(COMMAND ${CMAKE_COMPILE} -g -fPIC -shared -std=c++11 ${OPBUILD_OPS_SRC} -D_GLIBCXX_USE_CXX11_ABI=0 - -I ${ASCEND_CANN_PACKAGE_PATH}/include -L ${ASCEND_CANN_PACKAGE_PATH}/lib64 -lexe_graph -lregister -ltiling_api - -o ${OPBUILD_OUT_DIR}/libascend_all_ops.so - RESULT_VARIABLE EXEC_RESULT - OUTPUT_VARIABLE EXEC_INFO - ERROR_VARIABLE EXEC_ERROR - ) - if (${EXEC_RESULT}) - message("build ops lib info: ${EXEC_INFO}") - message("build ops lib error: ${EXEC_ERROR}") - message(FATAL_ERROR "opbuild run failed!") - endif() - set(proj_env "") - set(prefix_env "") - if (NOT "${OPBUILD_PROJECT_NAME}x" STREQUAL "x") - set(proj_env "OPS_PROJECT_NAME=${OPBUILD_PROJECT_NAME}") - endif() - if (NOT "${OPBUILD_ACCESS_PREFIX}x" STREQUAL "x") - set(prefix_env "OPS_DIRECT_ACCESS_PREFIX=${OPBUILD_ACCESS_PREFIX}") - endif() - execute_process(COMMAND ${proj_env} ${prefix_env} ${ASCEND_CANN_PACKAGE_PATH}/toolkit/tools/opbuild/op_build - ${OPBUILD_OUT_DIR}/libascend_all_ops.so ${OPBUILD_OUT_DIR} - RESULT_VARIABLE EXEC_RESULT - OUTPUT_VARIABLE EXEC_INFO - ERROR_VARIABLE EXEC_ERROR - ) - if (${EXEC_RESULT}) - message("opbuild ops info: ${EXEC_INFO}") - message("opbuild ops error: ${EXEC_ERROR}") - endif() - message(STATUS "Opbuild generating sources - done") -endfunction() - -function(add_ops_info_target) - cmake_parse_arguments(OPINFO "" "TARGET;OPS_INFO;OUTPUT;INSTALL_DIR" "" ${ARGN}) - get_filename_component(opinfo_file_path "${OPINFO_OUTPUT}" DIRECTORY) - add_custom_command(OUTPUT ${OPINFO_OUTPUT} - COMMAND mkdir -p ${opinfo_file_path} - COMMAND ${ASCEND_PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/cmake/util/parse_ini_to_json.py - ${OPINFO_OPS_INFO} ${OPINFO_OUTPUT} - ) - add_custom_target(${OPINFO_TARGET} ALL - DEPENDS ${OPINFO_OUTPUT} - ) - install(FILES ${OPINFO_OUTPUT} - DESTINATION ${OPINFO_INSTALL_DIR} - ) -endfunction() - -function(add_ops_compile_options OP_TYPE) - cmake_parse_arguments(OP_COMPILE "" "OP_TYPE" "COMPUTE_UNIT;OPTIONS" ${ARGN}) - file(APPEND ${ASCEND_AUTOGEN_PATH}/${CUSTOM_COMPILE_OPTIONS} - "${OP_TYPE},${OP_COMPILE_COMPUTE_UNIT},${OP_COMPILE_OPTIONS}\n") -endfunction() - -function(add_ops_impl_target) - cmake_parse_arguments(OPIMPL "" "TARGET;OPS_INFO;IMPL_DIR;OUT_DIR;INSTALL_DIR" "OPS_BATCH;OPS_ITERATE" ${ARGN}) - add_custom_command(OUTPUT ${OPIMPL_OUT_DIR}/.impl_timestamp - COMMAND mkdir -m 700 -p ${OPIMPL_OUT_DIR}/dynamic - COMMAND ${ASCEND_PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/cmake/util/ascendc_impl_build.py - ${OPIMPL_OPS_INFO} - \"${OPIMPL_OPS_BATCH}\" \"${OPIMPL_OPS_ITERATE}\" - ${OPIMPL_IMPL_DIR} - ${OPIMPL_OUT_DIR}/dynamic - ${ASCEND_AUTOGEN_PATH} - - COMMAND rm -rf ${OPIMPL_OUT_DIR}/.impl_timestamp - COMMAND touch ${OPIMPL_OUT_DIR}/.impl_timestamp - DEPENDS ${OPIMPL_OPS_INFO} - ${CMAKE_SOURCE_DIR}/cmake/util/ascendc_impl_build.py - ) - add_custom_target(${OPIMPL_TARGET} ALL - DEPENDS ${OPIMPL_OUT_DIR}/.impl_timestamp) - if (${ENABLE_SOURCE_PACKAGE}) - install(DIRECTORY ${OPIMPL_OUT_DIR}/dynamic - DESTINATION ${OPIMPL_INSTALL_DIR} - ) - endif() -endfunction() - -function(add_npu_support_target) - cmake_parse_arguments(NPUSUP "" "TARGET;OPS_INFO_DIR;OUT_DIR;INSTALL_DIR" "" ${ARGN}) - get_filename_component(npu_sup_file_path "${NPUSUP_OUT_DIR}" DIRECTORY) - add_custom_command(OUTPUT ${NPUSUP_OUT_DIR}/npu_supported_ops.json - COMMAND mkdir -p ${NPUSUP_OUT_DIR} - COMMAND ${CMAKE_SOURCE_DIR}/cmake/util/gen_ops_filter.sh - ${NPUSUP_OPS_INFO_DIR} - ${NPUSUP_OUT_DIR} - ) - add_custom_target(npu_supported_ops ALL - DEPENDS ${NPUSUP_OUT_DIR}/npu_supported_ops.json - ) - install(FILES ${NPUSUP_OUT_DIR}/npu_supported_ops.json - DESTINATION ${NPUSUP_INSTALL_DIR} - ) -endfunction() - -function(add_bin_compile_target) - cmake_parse_arguments(BINCMP "" "TARGET;OPS_INFO;COMPUTE_UNIT;IMPL_DIR;ADP_DIR;OUT_DIR;INSTALL_DIR" "" ${ARGN}) - file(MAKE_DIRECTORY ${BINCMP_OUT_DIR}/src) - file(MAKE_DIRECTORY ${BINCMP_OUT_DIR}/bin) - file(MAKE_DIRECTORY ${BINCMP_OUT_DIR}/gen) - execute_process(COMMAND ${ASCEND_PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/cmake/util/ascendc_bin_param_build.py - ${BINCMP_OPS_INFO} ${BINCMP_OUT_DIR}/gen ${BINCMP_COMPUTE_UNIT} - RESULT_VARIABLE EXEC_RESULT - OUTPUT_VARIABLE EXEC_INFO - ERROR_VARIABLE EXEC_ERROR - ) - if (${EXEC_RESULT}) - message("ops binary compile scripts gen info: ${EXEC_INFO}") - message("ops binary compile scripts gen error: ${EXEC_ERROR}") - message(FATAL_ERROR "ops binary compile scripts gen failed!") - endif() - if (NOT TARGET binary) - add_custom_target(binary) - endif() - add_custom_target(${BINCMP_TARGET} - COMMAND cp -r ${BINCMP_IMPL_DIR}/*.* ${BINCMP_OUT_DIR}/src - ) - add_custom_target(${BINCMP_TARGET}_gen_ops_config - COMMAND ${ASCEND_PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/cmake/util/insert_simplified_keys.py -p ${BINCMP_OUT_DIR}/bin - COMMAND ${ASCEND_PYTHON_EXECUTABLE} ${CMAKE_SOURCE_DIR}/cmake/util/ascendc_ops_config.py -p ${BINCMP_OUT_DIR}/bin - -s ${BINCMP_COMPUTE_UNIT} - ) - add_dependencies(binary ${BINCMP_TARGET}_gen_ops_config) - file(GLOB bin_scripts ${BINCMP_OUT_DIR}/gen/*.sh) - foreach(bin_script ${bin_scripts}) - get_filename_component(bin_file ${bin_script} NAME_WE) - string(REPLACE "-" ";" bin_sep ${bin_file}) - list(GET bin_sep 0 op_type) - list(GET bin_sep 1 op_file) - list(GET bin_sep 2 op_index) - if (NOT TARGET ${BINCMP_TARGET}_${op_file}_copy) - file(MAKE_DIRECTORY ${BINCMP_OUT_DIR}/bin/${op_file}) - add_custom_target(${BINCMP_TARGET}_${op_file}_copy - COMMAND cp ${BINCMP_ADP_DIR}/${op_file}.py ${BINCMP_OUT_DIR}/src/${op_type}.py - ) - install(DIRECTORY ${BINCMP_OUT_DIR}/bin/${op_file} - DESTINATION ${BINCMP_INSTALL_DIR}/${BINCMP_COMPUTE_UNIT} OPTIONAL - ) - install(FILES ${BINCMP_OUT_DIR}/bin/${op_file}.json - DESTINATION ${BINCMP_INSTALL_DIR}/config/${BINCMP_COMPUTE_UNIT}/ OPTIONAL - ) - endif() - add_custom_target(${BINCMP_TARGET}_${op_file}_${op_index} - COMMAND export HI_PYTHON=${ASCEND_PYTHON_EXECUTABLE} && bash ${bin_script} ${BINCMP_OUT_DIR}/src/${op_type}.py ${BINCMP_OUT_DIR}/bin/${op_file} && echo $(MAKE) - WORKING_DIRECTORY ${BINCMP_OUT_DIR} - ) - add_dependencies(${BINCMP_TARGET}_${op_file}_${op_index} ${BINCMP_TARGET} ${BINCMP_TARGET}_${op_file}_copy) - add_dependencies(${BINCMP_TARGET}_gen_ops_config ${BINCMP_TARGET}_${op_file}_${op_index}) - endforeach() - install(FILES ${BINCMP_OUT_DIR}/bin/binary_info_config.json - DESTINATION ${BINCMP_INSTALL_DIR}/config/${BINCMP_COMPUTE_UNIT} OPTIONAL - ) - - install(DIRECTORY ${BINCMP_OUT_DIR}/bin/${op_file} - DESTINATION ${CMAKE_CURRENT_SOURCE_DIR}/../build_out/kernel/${BINCMP_COMPUTE_UNIT} OPTIONAL - ) - install(FILES ${BINCMP_OUT_DIR}/bin/binary_info_config.json - DESTINATION ${CMAKE_CURRENT_SOURCE_DIR}/../build_out/kernel/config/${BINCMP_COMPUTE_UNIT} OPTIONAL - ) - install(FILES ${BINCMP_OUT_DIR}/bin/${op_file}.json - DESTINATION ${CMAKE_CURRENT_SOURCE_DIR}/../build_out/kernel/config/${BINCMP_COMPUTE_UNIT} OPTIONAL - ) - -endfunction() - -function(add_cross_compile_target) - cmake_parse_arguments(CROSSMP "" "TARGET;OUT_DIR;INSTALL_DIR" "" ${ARGN}) - add_custom_target(${CROSSMP_TARGET} ALL - DEPENDS ${CROSSMP_OUT_DIR} - ) - install(DIRECTORY ${CROSSMP_OUT_DIR} - DESTINATION ${CROSSMP_INSTALL_DIR} - ) -endfunction() diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/intf.cmake b/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/intf.cmake deleted file mode 100644 index 35d1e555..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/intf.cmake +++ /dev/null @@ -1,34 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -add_library(intf_pub INTERFACE) -target_compile_options(intf_pub INTERFACE - -fPIC - -fvisibility=hidden - -fvisibility-inlines-hidden - $<$:-O2> - $<$:-O0 -g> - $<$:-std=c++11> - $<$,$>:-ftrapv -fstack-check> - $<$:-pthread -Wfloat-equal -Wshadow -Wformat=2 -Wno-deprecated -Wextra> - $,-fstack-protector-strong,-fstack-protector-all> -) -target_compile_definitions(intf_pub INTERFACE - _GLIBCXX_USE_CXX11_ABI=0 - $<$:_FORTIFY_SOURCE=2> -) -target_include_directories(intf_pub INTERFACE ${ASCEND_CANN_PACKAGE_PATH}/include) -target_link_options(intf_pub INTERFACE - $<$,EXECUTABLE>:-pie> - $<$:-s> - -Wl,-z,relro - -Wl,-z,now - -Wl,-z,noexecstack -) -target_link_directories(intf_pub INTERFACE ${ASCEND_CANN_PACKAGE_PATH}/lib64) diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/makeself.cmake b/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/makeself.cmake deleted file mode 100644 index fce82f3b..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/cmake/makeself.cmake +++ /dev/null @@ -1,26 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -execute_process(COMMAND chmod +x ${CMAKE_CURRENT_LIST_DIR}/util/makeself/makeself.sh) -execute_process(COMMAND ${CMAKE_CURRENT_LIST_DIR}/util/makeself/makeself.sh - --header ${CMAKE_CURRENT_LIST_DIR}/util/makeself/makeself-header.sh - --help-header ./help.info - --gzip --complevel 4 --nomd5 --sha256 - ./ ${CPACK_PACKAGE_FILE_NAME} "version:1.0" ./install.sh - WORKING_DIRECTORY ${CPACK_TEMPORARY_DIRECTORY} - RESULT_VARIABLE EXEC_RESULT - ERROR_VARIABLE EXEC_ERROR -) -if (NOT "${EXEC_RESULT}x" STREQUAL "0x") - message(FATAL_ERROR "CPack Command error: ${EXEC_RESULT}\n${EXEC_ERROR}") -endif() -execute_process(COMMAND cp ${CPACK_EXTERNAL_BUILT_PACKAGES} ${CPACK_PACKAGE_DIRECTORY}/ - COMMAND echo "Copy ${CPACK_EXTERNAL_BUILT_PACKAGES} to ${CPACK_PACKAGE_DIRECTORY}/" - WORKING_DIRECTORY ${CPACK_TEMPORARY_DIRECTORY} -) diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/op_host/CMakeLists.txt b/examples/pad/broadcast/kernel_launch_method_by_framework/op_host/CMakeLists.txt deleted file mode 100644 index 38aecf8e..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/op_host/CMakeLists.txt +++ /dev/null @@ -1,90 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -aux_source_directory(${CMAKE_CURRENT_SOURCE_DIR} ops_srcs) - -opbuild(OPS_SRC ${ops_srcs} - OUT_DIR ${ASCEND_AUTOGEN_PATH} -) - -add_library(cust_op_proto SHARED ${ops_srcs} ${ASCEND_AUTOGEN_PATH}/op_proto.cc) -target_compile_definitions(cust_op_proto PRIVATE OP_PROTO_LIB) -target_compile_options(cust_op_proto PRIVATE - -fvisibility=hidden -) -if(ENABLE_CROSS_COMPILE) - target_link_directories(cust_op_proto PRIVATE - ${CMAKE_COMPILE_COMPILER_LIBRARY} - ${CMAKE_COMPILE_RUNTIME_LIBRARY} - ) -endif() -target_link_libraries(cust_op_proto PRIVATE - intf_pub - exe_graph - register - tiling_api - -Wl,--whole-archive - rt2_registry - -Wl,--no-whole-archive -) -set_target_properties(cust_op_proto PROPERTIES OUTPUT_NAME - cust_opsproto_rt2.0 -) -add_library(cust_optiling SHARED ${ops_srcs}) -target_compile_definitions(cust_optiling PRIVATE OP_TILING_LIB) -target_compile_options(cust_optiling PRIVATE - -fvisibility=hidden -) -if(ENABLE_CROSS_COMPILE) - target_link_directories(cust_optiling PRIVATE - ${CMAKE_COMPILE_COMPILER_LIBRARY} - ${CMAKE_COMPILE_RUNTIME_LIBRARY} - ) -endif() -target_link_libraries(cust_optiling PRIVATE - intf_pub - exe_graph - register - tiling_api - -Wl,--whole-archive - rt2_registry - -Wl,--no-whole-archive -) -set_target_properties(cust_optiling PROPERTIES OUTPUT_NAME - cust_opmaster_rt2.0 -) - -file(GLOB aclnn_src ${ASCEND_AUTOGEN_PATH}/aclnn_*.cpp) -file(GLOB aclnn_inc ${ASCEND_AUTOGEN_PATH}/aclnn_*.h) -add_library(cust_opapi SHARED ${aclnn_src}) -if(ENABLE_CROSS_COMPILE) - target_link_directories(cust_opapi PRIVATE - ${CMAKE_COMPILE_COMPILER_LIBRARY} - ${CMAKE_COMPILE_RUNTIME_LIBRARY} - ) -endif() -target_link_libraries(cust_opapi PRIVATE intf_pub ascendcl nnopbase) - -add_custom_target(optiling_compat ALL - COMMAND ln -sf lib/linux/${CMAKE_SYSTEM_PROCESSOR}/$ - ${CMAKE_CURRENT_BINARY_DIR}/liboptiling.so -) - -install(TARGETS cust_op_proto - LIBRARY DESTINATION packages/vendors/${vendor_name}/op_proto/lib/linux/${CMAKE_SYSTEM_PROCESSOR}) -install(FILES ${ASCEND_AUTOGEN_PATH}/op_proto.h - DESTINATION packages/vendors/${vendor_name}/op_proto/inc) -install(TARGETS cust_optiling - LIBRARY DESTINATION packages/vendors/${vendor_name}/op_impl/ai_core/tbe/op_tiling/lib/linux/${CMAKE_SYSTEM_PROCESSOR}) -install(FILES ${CMAKE_CURRENT_BINARY_DIR}/liboptiling.so - DESTINATION packages/vendors/${vendor_name}/op_impl/ai_core/tbe/op_tiling) -install(TARGETS cust_opapi - LIBRARY DESTINATION packages/vendors/${vendor_name}/op_api/lib) -install(FILES ${aclnn_inc} - DESTINATION packages/vendors/${vendor_name}/op_api/include) diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/op_host/broadcast_custom.cpp b/examples/pad/broadcast/kernel_launch_method_by_framework/op_host/broadcast_custom.cpp deleted file mode 100644 index 504155f5..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/op_host/broadcast_custom.cpp +++ /dev/null @@ -1,81 +0,0 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. - * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). - * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. - * See LICENSE in the root of the software repository for the full text of the License. - */ -#include "broadcast_custom_tiling.h" -#include "register/op_def_registry.h" - -namespace optiling { -static ge::graphStatus TilingFunc(gert::TilingContext *context) -{ - auto dt = context->GetInputTensor(0)->GetDataType(); - uint32_t dtypeSize; - if (dt == ge::DT_FLOAT16) { - dtypeSize = 2; - } else if (dt == ge::DT_FLOAT) { - dtypeSize = 4; - } else { - dtypeSize = 1; - } - const auto &srcStorageShape = context->GetInputShape(0)->GetStorageShape(); - const auto &outputStorageShape = context->GetOutputShape(0)->GetStorageShape(); - ge::Shape inputShape; - ge::Shape outputShape; - BroadcastTilingData tiling; - const uint32_t dim = srcStorageShape.GetDimNum(); - if (dim == 1) { - inputShape = ge::Shape({srcStorageShape.GetDim(0)}); - outputShape = ge::Shape({outputStorageShape.GetDim(0)}); - } else { - inputShape = ge::Shape({srcStorageShape.GetDim(0), srcStorageShape.GetDim(1)}); - outputShape = ge::Shape({outputStorageShape.GetDim(0), outputStorageShape.GetDim(1)}); - } - - ComputeTiling(inputShape, outputShape, dtypeSize, tiling); - - context->SetBlockDim(1); - context->SetTilingKey(1); - tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); - context->GetRawTilingData()->SetDataSize(tiling.GetDataSize()); - size_t *currentWorkSpace = context->GetWorkspaceSizes(1); - currentWorkSpace[0] = 0; - return ge::GRAPH_SUCCESS; -} -} // namespace optiling - -namespace ge { -static graphStatus InferShape(gert::InferShapeContext *context) -{ - return GRAPH_SUCCESS; -} -} // namespace ge - -namespace ops { -class BroadcastCustom : public OpDef { -public: - BroadcastCustom(const char *name) : OpDef(name) - { - this->Input("x") - .ParamType(REQUIRED) - .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_UINT8}) - .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) - .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}); - this->Output("y") - .ParamType(REQUIRED) - .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT8, ge::DT_UINT8}) - .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) - .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}); - this->SetInferShape(ge::InferShape); - this->AICore().SetTiling(optiling::TilingFunc); - this->AICore().AddConfig("ascend910b"); - this->AICore().AddConfig("ascend310p"); - } -}; - -OP_ADD(BroadcastCustom); -} // namespace ops \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt b/examples/pad/broadcast/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt deleted file mode 100644 index 2102a7f3..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/op_kernel/CMakeLists.txt +++ /dev/null @@ -1,78 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -# set custom compile options -if ("${CMAKE_BUILD_TYPE}x" STREQUAL "Debugx") - add_ops_compile_options(ALL OPTIONS -g -O0) -endif() -add_ops_compile_options(ALL OPTIONS -mllvm -cce-aicore-jump-expand=true) - -foreach(compute_unit ${ASCEND_COMPUTE_UNIT}) - - # generate aic-${compute_unit}-ops-info.json - add_ops_info_target(TARGET ops_info_gen_${compute_unit} - OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/tbe/op_info_cfg/ai_core/${compute_unit}/aic-${compute_unit}-ops-info.json - OPS_INFO ${ASCEND_AUTOGEN_PATH}/aic-${compute_unit}-ops-info.ini - INSTALL_DIR packages/vendors/${vendor_name}/op_impl/ai_core/tbe/config/${compute_unit} - ) - - # generate ascendc impl py once - if (NOT TARGET ascendc_impl_gen) - add_ops_impl_target(TARGET ascendc_impl_gen - OPS_INFO ${ASCEND_AUTOGEN_PATH}/aic-${compute_unit}-ops-info.ini - IMPL_DIR ${CMAKE_CURRENT_SOURCE_DIR} - OUT_DIR ${CMAKE_CURRENT_BINARY_DIR}/tbe - INSTALL_DIR packages/vendors/${vendor_name}/op_impl/ai_core/tbe/${vendor_name}_impl - ) - endif() - - # dynamic shape binary compile - if (${ENABLE_BINARY_PACKAGE} AND NOT ${ENABLE_CROSS_COMPILE}) - add_bin_compile_target(TARGET ascendc_bin_${compute_unit} - OPS_INFO ${ASCEND_AUTOGEN_PATH}/aic-${compute_unit}-ops-info.ini - IMPL_DIR ${CMAKE_CURRENT_SOURCE_DIR} - ADP_DIR ${CMAKE_CURRENT_BINARY_DIR}/tbe/dynamic - OUT_DIR ${CMAKE_CURRENT_BINARY_DIR}/binary/${compute_unit} - INSTALL_DIR packages/vendors/${vendor_name}/op_impl/ai_core/tbe/kernel - COMPUTE_UNIT ${compute_unit} - ) - add_dependencies(ascendc_bin_${compute_unit} ascendc_impl_gen) - endif() - - if (${ENABLE_CROSS_COMPILE} AND ${ENABLE_BINARY_PACKAGE}) - add_cross_compile_target( - TARGET bin_${compute_unit} - OUT_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../kernel - INSTALL_DIR packages/vendors/${vendor_name}/op_impl/ai_core/tbe/ - ) - endif() -endforeach() - -# generate npu_supported_ops.json -add_npu_support_target(TARGET npu_supported_ops - OPS_INFO_DIR ${ASCEND_AUTOGEN_PATH} - OUT_DIR ${CMAKE_CURRENT_BINARY_DIR}/tbe/op_info_cfg/ai_core - INSTALL_DIR packages/vendors/${vendor_name}/framework/${ASCEND_FRAMEWORK_TYPE} -) - -if(ENABLE_TEST AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/testcases) - add_subdirectory(testcases) -endif() - -# install kernel file -if (${ENABLE_SOURCE_PACKAGE}) - file(GLOB KERNEL_FILES - ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/*.h - ${CMAKE_CURRENT_SOURCE_DIR}/*.py - ) - install(FILES ${KERNEL_FILES} - DESTINATION packages/vendors/${vendor_name}/op_impl/ai_core/tbe/${vendor_name}_impl/dynamic - ) -endif() diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/op_kernel/broadcast_custom.cpp b/examples/pad/broadcast/kernel_launch_method_by_framework/op_kernel/broadcast_custom.cpp deleted file mode 100644 index 230a5975..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/op_kernel/broadcast_custom.cpp +++ /dev/null @@ -1,50 +0,0 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. - * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). - * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. - * See LICENSE in the root of the software repository for the full text of the License. - */ -#include "../../../../../../kernel_impl/broadcast_custom.h" - -extern "C" __global__ __aicore__ void broadcast_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) -{ - KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); - GET_TILING_DATA(tilingData, tiling); - uint32_t axis = tilingData.axis; - uint32_t dim = tilingData.dim; - if (TILING_KEY_IS(1)) { - if (dim == 1) { - const uint32_t srcShape[] = {tilingData.srcFirstDim}; - const uint32_t dstShape[] = {tilingData.dstFirstDim}; - KernelBroadcastCustom op; - op.Init(x, y, tilingData.srcFirstDim, tilingData.dstFirstDim, srcShape, dstShape); - op.Process(); - } else { - const uint32_t srcShape[] = {tilingData.srcFirstDim, tilingData.srcLastDim}; - const uint32_t dstShape[] = {tilingData.dstFirstDim, tilingData.dstLastDim}; - - if (axis == 0) { - KernelBroadcastCustom op; - op.Init(x, - y, - tilingData.srcFirstDim * tilingData.srcLastDim, - tilingData.dstFirstDim * tilingData.dstLastDim, - srcShape, - dstShape); - op.Process(); - } else { - KernelBroadcastCustom op; - op.Init(x, - y, - tilingData.srcFirstDim * tilingData.srcLastDim, - tilingData.dstFirstDim * tilingData.dstLastDim, - srcShape, - dstShape); - op.Process(); - } - } - } -} \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/gen_data.py b/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/gen_data.py deleted file mode 100644 index 9aba79bb..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/gen_data.py +++ /dev/null @@ -1,37 +0,0 @@ -#!/usr/bin/python3 -# coding=utf-8 - -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -import numpy as np -import os -import argparse - -def gen_golden_data_simple(testcase): - x_shape = (1, 48) - y_shape = (96, 48) - if testcase == 1: - x_shape = (96, 1) - y_shape = (96, 96) - x = np.random.uniform(-10, 10, x_shape).astype(np.float32) - y = np.broadcast_to(x, y_shape).astype(np.float32) - os.system("mkdir -p input") - os.system("mkdir -p output") - x.tofile("./input/input.bin") - y.tofile("./output/golden.bin") - -if __name__ == "__main__": - parser = argparse.ArgumentParser() - parser.add_argument('--testcase', type=int, help='testcase,目前只有case 0 与 1') - args = parser.parse_args() - testcase = 0 - if args.testcase: - testcase = args.testcase - gen_golden_data_simple(testcase) \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/help.info b/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/help.info deleted file mode 100644 index f4b28d57..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/help.info +++ /dev/null @@ -1 +0,0 @@ - --install-path Install operator package to specific dir path \ No newline at end of file diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/install.sh b/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/install.sh deleted file mode 100644 index 17b4de60..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/install.sh +++ /dev/null @@ -1,313 +0,0 @@ -#!/bin/bash -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -vendor_name=customize -targetdir=/usr/local/Ascend/opp -target_custom=0 - -sourcedir=$PWD/packages -vendordir=vendors/$vendor_name - -QUIET="y" - -while true -do - case $1 in - --quiet) - QUIET="y" - shift - ;; - --install-path=*) - INSTALL_PATH=$(echo $1 | cut -d"=" -f2-) - INSTALL_PATH=${INSTALL_PATH%*/} - shift - ;; - --*) - shift - ;; - *) - break - ;; - esac -done - -log() { - cur_date=`date +"%Y-%m-%d %H:%M:%S"` - echo "[runtime] [$cur_date] "$1 -} - -if [ -n "${INSTALL_PATH}" ]; then - if [[ ! "${INSTALL_PATH}" = /* ]]; then - log "[ERROR] use absolute path for --install-path argument" - exit 1 - fi - if [ ! -d ${INSTALL_PATH} ]; then - mkdir ${INSTALL_PATH} >> /dev/null 2>&1 - if [ $? -ne 0 ]; then - log "[ERROR] create ${INSTALL_PATH} failed" - exit 1 - fi - fi - targetdir=${INSTALL_PATH} -elif [ -n "${ASCEND_CUSTOM_OPP_PATH}" ]; then - if [ ! -d ${ASCEND_CUSTOM_OPP_PATH} ]; then - mkdir -p ${ASCEND_CUSTOM_OPP_PATH} >> /dev/null 2>&1 - if [ $? -ne 0 ]; then - log "[ERROR] create ${ASCEND_CUSTOM_OPP_PATH} failed" - fi - fi - targetdir=${ASCEND_CUSTOM_OPP_PATH} -else - if [ "x${ASCEND_OPP_PATH}" == "x" ]; then - log "[ERROR] env ASCEND_OPP_PATH no exist" - exit 1 - fi - targetdir="${ASCEND_OPP_PATH}" -fi - -if [ ! -d $targetdir ];then - log "[ERROR] $targetdir no exist" - exit 1 -fi - -upgrade() -{ - if [ ! -d ${sourcedir}/$vendordir/$1 ]; then - log "[INFO] no need to upgrade ops $1 files" - return 0 - fi - - if [ ! -d ${targetdir}/$vendordir/$1 ];then - log "[INFO] create ${targetdir}/$vendordir/$1." - mkdir -p ${targetdir}/$vendordir/$1 - if [ $? -ne 0 ];then - log "[ERROR] create ${targetdir}/$vendordir/$1 failed" - return 1 - fi - else - has_same_file=-1 - for file_a in ${sourcedir}/$vendordir/$1/*; do - file_b=${file_a##*/}; - if [ "ls ${targetdir}/$vendordir/$1" = "" ]; then - log "[INFO] ${targetdir}/$vendordir/$1 is empty !!" - return 1 - fi - grep -q $file_b <<<`ls ${targetdir}/$vendordir/$1`; - if [[ $? -eq 0 ]]; then - echo -n "${file_b} " - has_same_file=0 - fi - done - if [ 0 -eq $has_same_file ]; then - if test $QUIET = "n"; then - echo "[INFO]: has old version in ${targetdir}/$vendordir/$1, \ - you want to Overlay Installation , please enter:[o]; \ - or replace directory installation , please enter: [r]; \ - or not install , please enter:[n]." - - while true - do - read orn - if [ "$orn" = n ]; then - return 0 - elif [ "$orn" = m ]; then - break; - elif [ "$0rn" = r ]; then - [ -n "${targetdir}/$vendordir/$1/" ] && rm -rf "${targetdir}/$vendordir/$1"/* - break; - else - echo "[ERROR] input error, please input again!" - fi - done - fi - fi - log "[INFO] replace or merge old ops $1 files .g....." - fi - - log "copy new ops $1 files ......" - if [ -d ${targetdir}/$vendordir/$1/ ]; then - chmod -R +w "$targetdir/$vendordir/$1/" >/dev/null 2>&1 - fi - cp -rf ${sourcedir}/$vendordir/$1/* $targetdir/$vendordir/$1/ - if [ $? -ne 0 ];then - log "[ERROR] copy new $1 files failed" - return 1 - fi - - return 0 -} -upgrade_proto() -{ - if [ ! -f ${sourcedir}/$vendordir/custom.proto ]; then - log "[INFO] no need to upgrade custom.proto files" - return 0 - fi - if [ ! -d ${targetdir}/$vendordir/framework/caffe ];then - log "[INFO] create ${targetdir}/$vendordir/framework/caffe." - mkdir -p ${targetdir}/$vendordir/framework/caffe - if [ $? -ne 0 ];then - log "[ERROR] create ${targetdir}/$vendordir/framework/caffe failed" - return 1 - fi - else - if [ -f ${targetdir}/$vendordir/framework/caffe/custom.proto ]; then - # 有老版本,判断是否要覆盖式安装 - if test $QUIET = "n"; then - echo "[INFO] ${targetdir}/$vendordir/framework/caffe has old version"\ - "custom.proto file. Do you want to replace? [y/n] " - - while true - do - read yn - if [ "$yn" = n ]; then - return 0 - elif [ "$yn" = y ]; then - break; - else - echo "[ERROR] input error, please input again!" - fi - done - fi - fi - log "[INFO] replace old caffe.proto files ......" - fi - chmod -R +w "$targetdir/$vendordir/framework/caffe/" >/dev/null 2>&1 - cp -rf ${sourcedir}/$vendordir/custom.proto ${targetdir}/$vendordir/framework/caffe/ - if [ $? -ne 0 ];then - log "[ERROR] copy new custom.proto failed" - return 1 - fi - log "[INFO] copy custom.proto success" - - return 0 -} - -upgrade_file() -{ - if [ ! -e ${sourcedir}/$vendordir/$1 ]; then - log "[INFO] no need to upgrade ops $1 file" - return 0 - fi - - log "copy new $1 files ......" - cp -f ${sourcedir}/$vendordir/$1 $targetdir/$vendordir/$1 - if [ $? -ne 0 ];then - log "[ERROR] copy new $1 file failed" - return 1 - fi - - return 0 -} - -log "[INFO] copy uninstall sh success" - -if [ ! -d ${targetdir}/vendors ];then - log "[INFO] create ${targetdir}/vendors." - mkdir -p ${targetdir}/vendors - if [ $? -ne 0 ];then - log "[ERROR] create ${targetdir}/vendors failed" - return 1 - fi -fi -chmod u+w ${targetdir}/vendors - -echo "[ops_custom]upgrade framework" -upgrade framework -if [ $? -ne 0 ];then - exit 1 -fi - -echo "[ops_custom]upgrade op proto" -upgrade op_proto -if [ $? -ne 0 ];then - exit 1 -fi - -echo "[ops_custom]upgrade version.info" -upgrade_file version.info -if [ $? -ne 0 ];then - exit 1 -fi - -echo "[ops_custom]upgrade op impl" -upgrade op_impl -if [ $? -ne 0 ];then - exit 1 -fi - -echo "[ops_custom]upgrade op api" -upgrade op_api -if [ $? -ne 0 ];then - exit 1 -fi - -upgrade_proto -if [ $? -ne 0 ];then - exit 1 -fi - -# set the set_env.bash -if [ -n "${INSTALL_PATH}" ] && [ -d ${INSTALL_PATH} ]; then - _ASCEND_CUSTOM_OPP_PATH=${targetdir}/${vendordir} - bin_path="${_ASCEND_CUSTOM_OPP_PATH}/bin" - set_env_variable="#!/bin/bash\nexport ASCEND_CUSTOM_OPP_PATH=${_ASCEND_CUSTOM_OPP_PATH}:\${ASCEND_CUSTOM_OPP_PATH}" - if [ ! -d ${bin_path} ]; then - mkdir -p ${bin_path} >> /dev/null 2>&1 - if [ $? -ne 0 ]; then - log "[ERROR] create ${bin_path} failed" - exit 1 - fi - fi - echo -e ${set_env_variable} > ${bin_path}/set_env.bash - if [ $? -ne 0 ]; then - log "[ERROR] write ASCEND_CUSTOM_OPP_PATH to set_env.bash failed" - exit 1 - else - log "[INFO] using requirements: when custom module install finished or before you run the custom module, \ - execute the command [ source ${bin_path}/set_env.bash ] to set the environment path" - fi -else - config_file=${targetdir}/vendors/config.ini - if [ ! -f ${config_file} ]; then - touch ${config_file} - chmod 640 ${config_file} - echo "load_priority=$vendor_name" > ${config_file} - if [ $? -ne 0 ];then - echo "echo load_priority failed" - exit 1 - fi - else - found_vendors="$(grep -w "load_priority" "$config_file" | cut --only-delimited -d"=" -f2-)" - found_vendor=$(echo $found_vendors | sed "s/$vendor_name//g" | tr ',' ' ') - vendor=$(echo $found_vendor | tr -s ' ' ',') - if [ "$vendor" != "" ]; then - sed -i "/load_priority=$found_vendors/s@load_priority=$found_vendors@load_priority=$vendor_name,$vendor@g" "$config_file" - fi - fi -fi - -chmod u-w ${targetdir}/vendors - -if [ -d ${targetdir}/$vendordir/op_impl/cpu/aicpu_kernel/impl/ ]; then - chmod -R 440 ${targetdir}/$vendordir/op_impl/cpu/aicpu_kernel/impl/* >/dev/null 2>&1 -fi -if [ -f ${targetdir}/ascend_install.info ]; then - chmod -R 440 ${targetdir}/ascend_install.info -fi -if [ -f ${targetdir}/scene.info ]; then - chmod -R 440 ${targetdir}/scene.info -fi -if [ -f ${targetdir}/version.info ]; then - chmod -R 440 ${targetdir}/version.info -fi - -echo "SUCCESS" -exit 0 - diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/upgrade.sh b/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/upgrade.sh deleted file mode 100644 index 8e7c7f80..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/scripts/upgrade.sh +++ /dev/null @@ -1,160 +0,0 @@ -#!/bin/bash -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -vendor_name=customize -targetdir=/usr/local/Ascend/opp -target_custom=0 - -sourcedir=$PWD/packages -vendordir=vendors/$vendor_name - -log() { - cur_date=`date +"%Y-%m-%d %H:%M:%S"` - echo "[runtime] [$cur_date] "$1 -} - -if [[ "x${ASCEND_OPP_PATH}" == "x" ]];then - log "[ERROR] env ASCEND_OPP_PATH no exist" - exit 1 -fi - -targetdir=${ASCEND_OPP_PATH} - -if [ ! -d $targetdir ];then - log "[ERROR] $targetdir no exist" - exit 1 -fi - -upgrade() -{ - if [ ! -d ${sourcedir}/$vendordir/$1 ]; then - log "[INFO] no need to upgrade ops $1 files" - return 0 - fi - - if [ ! -d ${targetdir}/$vendordir/$1 ];then - log "[INFO] create ${targetdir}/$vendordir/$1." - mkdir -p ${targetdir}/$vendordir/$1 - if [ $? -ne 0 ];then - log "[ERROR] create ${targetdir}/$vendordir/$1 failed" - return 1 - fi - else - vendor_installed_dir=$(ls "$targetdir/vendors" 2> /dev/null) - for i in $vendor_installed_dir;do - vendor_installed_file=$(ls "$vendor_installed_dir/$vendor_name/$i" 2> /dev/null) - if [ "$i" = "$vendor_name" ] && [ "$vendor_installed_file" != "" ]; then - echo "[INFO]: $vendor_name custom opp package has been installed on the path $vendor_installed_dir, \ - you want to Overlay Installation , please enter:[o]; \ - or replace directory installation , please enter: [r]; \ - or not install , please enter:[n]." - fi - while true - do - read mrn - if [ "$mrn" = m ]; then - break - elif [ "$mrn" = r ]; then - [ -n "$vendor_installed_file"] && rm -rf "$vendor_installed_file" - break - elif [ "$mrn" = n ]; then - return 0 - else - echo "[WARNING]: Input error, please input m or r or n to choose!" - fi - done - done - log "[INFO] replace old ops $1 files ......" - fi - - log "copy new ops $1 files ......" - cp -rf ${sourcedir}/$vendordir/$1/* $targetdir/$vendordir/$1/ - if [ $? -ne 0 ];then - log "[ERROR] copy new $1 files failed" - return 1 - fi - - return 0 -} - -upgrade_file() -{ - if [ ! -e ${sourcedir}/$vendordir/$1 ]; then - log "[INFO] no need to upgrade ops $1 file" - return 0 - fi - - log "copy new $1 files ......" - cp -f ${sourcedir}/$vendordir/$1 $targetdir/$vendordir/$1 - if [ $? -ne 0 ];then - log "[ERROR] copy new $1 file failed" - return 1 - fi - - return 0 -} - -log "[INFO] copy uninstall sh success" - -echo "[ops_custom]upgrade framework" -upgrade framework -if [ $? -ne 0 ];then - exit 1 -fi - -echo "[ops_custom]upgrade op proto" -upgrade op_proto -if [ $? -ne 0 ];then - exit 1 -fi - -echo "[ops_custom]upgrade op impl" -upgrade op_impl -if [ $? -ne 0 ];then - exit 1 -fi - -echo "[ops_custom]upgrade op api" -upgrade op_api -if [ $? -ne 0 ];then - exit 1 -fi - -echo "[ops_custom]upgrade version.info" -upgrade_file version.info -if [ $? -ne 0 ];then - exit 1 -fi - -config_file=${targetdir}/vendors/config.ini -found_vendors="$(grep -w "load_priority" "$config_file" | cut --only-delimited -d"=" -f2-)" -found_vendor=$(echo $found_vendors | sed "s/$vendor_name//g" | tr ',' ' ') -vendor=$(echo $found_vendor | tr -s ' ' ',') -if [ "$vendor" != "" ]; then - sed -i "/load_priority=$found_vendors/s@load_priority=$found_vendors@load_priority=$vendor_name,$vendor@g" "$config_file" -fi - -changemode() -{ - if [ -d ${targetdir} ];then - chmod -R 550 ${targetdir}>/dev/null 2>&1 - fi - - return 0 -} -echo "[ops_custom]changemode..." -#changemode -if [ $? -ne 0 ];then - exit 1 -fi - -echo "SUCCESS" -exit 0 - diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/CMakeLists.txt b/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/CMakeLists.txt deleted file mode 100644 index ef1f61f0..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/CMakeLists.txt +++ /dev/null @@ -1,11 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -include(cmake/fun.cmake) -add_subdirectory(npu) diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/cmake/fun.cmake b/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/cmake/fun.cmake deleted file mode 100644 index b7bdbf56..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/cmake/fun.cmake +++ /dev/null @@ -1,61 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -set(UPER_CHARS A B C D E F G H I J K L M N O P Q R S T U V W X Y Z) -function(string_to_snake str_in snake_out) - set(str_cam ${str_in}) - foreach(uper_char ${UPER_CHARS}) - string(TOLOWER "${uper_char}" lower_char) - string(REPLACE ${uper_char} "_${lower_char}" str_cam ${str_cam}) - endforeach() - string(SUBSTRING ${str_cam} 1 -1 str_cam) - set(${snake_out} "${str_cam}" PARENT_SCOPE) -endfunction() - -function(add_cpu_target) - cmake_parse_arguments(CPU_TEST "" "OP" "SRC" ${ARGN}) - string_to_snake("${CPU_TEST_OP}" op_snake) - add_custom_command(OUTPUT ${CMAKE_CURRENT_SOURCE_DIR}/${op_snake}_tiling.h - COMMAND python3 ${CMAKE_SOURCE_DIR}/cmake/util/tiling_data_def_build.py - ${CMAKE_SOURCE_DIR}/op_host/${op_snake}_tiling.h - ${CMAKE_CURRENT_SOURCE_DIR}/${op_snake}_tiling.h - DEPENDS ${CMAKE_SOURCE_DIR}/op_host/${op_snake}_tiling.h - ) - add_custom_target(gen_${op_snake}_tiling_header - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/${op_snake}_tiling.h - ) - - add_executable(${op_snake}_cpu ${CPU_TEST_SRC}) - add_dependencies(${op_snake}_cpu gen_${op_snake}_tiling_header) - target_compile_options(${op_snake}_cpu PRIVATE -g -include ${CMAKE_CURRENT_SOURCE_DIR}/${op_snake}_tiling.h) - target_link_libraries(${op_snake}_cpu PRIVATE tikicpulib::ascend910B1) - set_target_properties(${op_snake}_cpu PROPERTIES - RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR} - ) -endfunction() - -function(add_npu_target) - cmake_parse_arguments(NPU_TEST "" "OP" "SRC" ${ARGN}) - string_to_snake("${NPU_TEST_OP}" op_snake) - add_executable(${op_snake}_npu ${NPU_TEST_SRC}) - target_compile_options(${op_snake}_npu PRIVATE -g) - target_include_directories(${op_snake}_npu PRIVATE - ${ASCEND_CANN_PACKAGE_PATH}/include/acl - ${ASCEND_AUTOGEN_PATH} - ) - target_link_libraries(${op_snake}_npu PRIVATE - intf_pub - cust_opapi - ascendcl - nnopbase - ) - set_target_properties(${op_snake}_npu PROPERTIES - RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR} - ) -endfunction() diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt b/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt deleted file mode 100644 index b5a9c947..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/npu/CMakeLists.txt +++ /dev/null @@ -1,19 +0,0 @@ -# Copyright (c) 2024 Huawei Technologies Co., Ltd. -# This file is a part of the CANN Open Software. -# Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. -# ====================================================================================================================== - -add_npu_target(OP BroadcastCustom SRC broadcast_custom_main.cpp) - -add_custom_target(run_npu_test - COMMAND echo "===============================================================================" - COMMAND echo " Run NPU test at ${CMAKE_CURRENT_BINARY_DIR}" - COMMAND echo "===============================================================================" - COMMAND $ - COMMAND echo "===============================================================================" - ) -add_dependencies(run_npu_test broadcast_custom_npu) diff --git a/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/npu/broadcast_custom_main.cpp b/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/npu/broadcast_custom_main.cpp deleted file mode 100644 index b7376505..00000000 --- a/examples/pad/broadcast/kernel_launch_method_by_framework/testcases/npu/broadcast_custom_main.cpp +++ /dev/null @@ -1,179 +0,0 @@ -/** - * Copyright (c) 2024 Huawei Technologies Co., Ltd. - * This file is a part of the CANN Open Software. - * Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). - * Please refer to the License for details. You may not use this file except in compliance with the License. - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. - * See LICENSE in the root of the software repository for the full text of the License. - */ - -#include "aclnn_broadcast_custom.h" -#include "acl/acl_rt.h" -#include "acl/acl.h" -#include -#include -#include -#include "../../../../../common/data_utils.h" - -aclrtStream CreateStream(int device) -{ - if (aclInit(NULL) != ACL_SUCCESS) { - printf("acl init failed\n"); - return NULL; - } - if (aclrtSetDevice(device) != ACL_SUCCESS) { - printf("Set device failed\n"); - (void)aclFinalize(); - return NULL; - } - aclrtStream stream = nullptr; - if (aclrtCreateStream(&stream) != ACL_SUCCESS) { - printf("Create stream failed\n"); - return NULL; - } - return stream; -} - -void DestroyStream(aclrtStream stream, int device) -{ - (void)aclrtDestroyStream(stream); - if (aclrtResetDevice(device) != ACL_SUCCESS) { - printf("Reset device failed\n"); - } - if (aclFinalize() != ACL_SUCCESS) { - printf("Finalize acl failed\n"); - } -} - -struct tensorInfo { - int64_t *dims; - int64_t dimCnt; - aclDataType dtype; - aclFormat fmt; -}; - -int64_t GetDataSize(struct tensorInfo *desc) -{ - if (!desc->dims) - return 0; - int64_t size = 1; - for (auto i = 0; i < desc->dimCnt; i++) { - size *= desc->dims[i]; - } - return size * sizeof(float); -} - -int64_t CompareResult(void *outputData, int64_t outSize) -{ - void *goldenData; - CHECK_ACL(aclrtMallocHost((void **)(&goldenData), outSize)); - size_t goldenSize = outSize; - bool ret = ReadFile("../output/golden.bin", goldenSize, goldenData, goldenSize); - if (ret) { - printf("ReadFile golden success!\n"); - } else { - return -1; - } - constexpr float EPS = 1e-6; - int64_t wrongNum = 0; - - for (int i = 0; i < outSize / sizeof(float); i++) { - float a = ((float *)outputData)[i]; - float b = ((float *)goldenData)[i]; - float ae = std::abs(a - b); - float re = ae / abs(b); - if (ae > EPS && re > EPS) { - printf("CompareResult failed output is %lf, golden is %lf\n", a, b); - wrongNum++; - } - } - CHECK_ACL(aclrtFreeHost(goldenData)); - return wrongNum; -} - -int main(int32_t argc, char *argv[]) -{ - aclrtStream stream; - int64_t inputX[] = {1, 48}; - int64_t outputY[] = {96, 48}; - if (argc == 2) { - if (!strcmp(argv[1], "1")) { - inputX[0] = 96; - inputX[1] = 1; - outputY[0] = 96; - outputY[1] = 96; - } - } - struct tensorInfo tensorDesc[] = { - {inputX, 2, ACL_FLOAT, ACL_FORMAT_ND}, - {outputY, 2, ACL_FLOAT, ACL_FORMAT_ND}, - }; - stream = CreateStream(0); - aclTensor *tensors[sizeof(tensorDesc) / sizeof(struct tensorInfo)]; - void *devMem[sizeof(tensorDesc) / sizeof(struct tensorInfo)]; - for (auto i = 0; i < sizeof(tensorDesc) / sizeof(struct tensorInfo); i++) { - void *data; - struct tensorInfo *info = &(tensorDesc[i]); - int64_t size = GetDataSize(info); - if (size == 0) { - tensors[i] = NULL; - devMem[i] = NULL; - continue; - } - CHECK_ACL(aclrtMalloc(&data, size, ACL_MEM_MALLOC_HUGE_FIRST)); - // read input - if (i == 0) { - size_t inputSize = size; - void *dataHost; - CHECK_ACL(aclrtMallocHost((void **)(&dataHost), inputSize)); - ReadFile("../input/input.bin", inputSize, dataHost, inputSize); - CHECK_ACL(aclrtMemcpy(data, size, dataHost, size, ACL_MEMCPY_HOST_TO_DEVICE)); - CHECK_ACL(aclrtFreeHost(dataHost)); - } - devMem[i] = data; - tensors[i] = - aclCreateTensor(info->dims, info->dimCnt, info->dtype, NULL, 0, info->fmt, info->dims, info->dimCnt, data); - } - size_t workspaceSize = 0; - aclOpExecutor *handle; - int32_t ret; - ret = aclnnBroadcastCustomGetWorkspaceSize(tensors[0], tensors[1], &workspaceSize, &handle); - printf("aclnnBroadcastCustomGetWorkspaceSize ret %u workspace size %lu\n", ret, workspaceSize); - void *workspace = NULL; - if (workspaceSize != 0) { - CHECK_ACL(aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST)); - } - ret = aclnnBroadcastCustom(workspace, workspaceSize, handle, stream); - printf("aclnnBroadcastCustom ret %u\n", ret); - if (aclrtSynchronizeStreamWithTimeout(stream, 5000) != ACL_SUCCESS) { - printf("Synchronize stream failed\n"); - } - - uint8_t *yHost; - int64_t yHostSize = GetDataSize(&(tensorDesc[1])); - CHECK_ACL(aclrtMallocHost((void **)(&yHost), yHostSize)); - CHECK_ACL(aclrtMemcpy(yHost, yHostSize, devMem[1], yHostSize, ACL_MEMCPY_DEVICE_TO_HOST)); - WriteFile("../output/output.bin", yHost, yHostSize); - - int64_t wrongNum = CompareResult(yHost, yHostSize); - if (wrongNum != 0) { - printf("test failed!\n"); - } else { - printf("test pass!\n"); - } - printf("\n"); - - CHECK_ACL(aclrtFreeHost(yHost)); - - for (auto i = 0; i < sizeof(tensorDesc) / sizeof(struct tensorInfo); i++) { - if (!tensors[i]) - continue; - if (devMem[i]) { - CHECK_ACL(aclrtFree(devMem[i])); - } - aclDestroyTensor(tensors[i]); - } - DestroyStream(stream, 0); - return 0; -} diff --git a/examples/readme.md b/examples/readme.md index 071b5db8..2b64a770 100644 --- a/examples/readme.md +++ b/examples/readme.md @@ -41,11 +41,6 @@ layernorm 将输入数据收敛到[0, 1]之间，每个tensor的特征值减去该特征的均值然后除以该特征的标准差，实现数据归一化。 - - pad - broadcast - 对输入tensor的shape进行广播。 - -- Gitee