diff --git a/examples/pad/broadcast/README.md b/examples/pad/broadcast/README.md
deleted file mode 100644
index 077a25872bb8a0ee721c7a9c5a56291af3c14356..0000000000000000000000000000000000000000
--- 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 |
-
-算子输入 | name | shape | data type | format |
-x | - | float | ND |
-
-
-
-
-算子输出 | y | - | float | ND |
-
-核函数名 | 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 884fa012b7d5aef4f032c4fe481080198c611ff4..0000000000000000000000000000000000000000
--- 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 6df59bf8caa2d5ea9e845a5a7e1fc4ac42c48523..0000000000000000000000000000000000000000
--- 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 af343dab3af6e080e96f1056e6bb03ea1441fb88..0000000000000000000000000000000000000000
--- 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 e565510f5cfc49381884261fec4a27266b5b8b0f..0000000000000000000000000000000000000000
--- 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 395ac0349a140c5a2cee5bec5ef3cf23d45723fd..0000000000000000000000000000000000000000
--- 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 0bb877a6bdf0d8a38e7f832ce2c4ecb905bf41e9..0000000000000000000000000000000000000000
--- 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