Star725
1026
代码
代码
Issues61
Pull Requests201
Actions
项目讨论
Wiki
分析
Star725
1026
已合并
feat: add parallel_ops_package multi-vendor build example and update sample index #1466
jianwei10创建于 4月7日
feat: add parallel_ops_package multi-vendor build example and update sample index #1466
已合并
jianwei10创建于 4月7日
从Geminiver/asc-devkitsamples合入到cann/asc-devkitmaster
共 21 个文件变更+956-1
Mexamples/01_simd_cpp_api/02_features/00_compilation/README.md+2-1
@@ -9,4 +9,5 @@
9| 目录名称 | 功能描述 |9| 目录名称 | 功能描述 |
10| --------- | --------- |10| --------- | --------- |
11| [custom_op](./custom_op/) | 本样例介绍了自定义算子编译工程的实现方法，包括Add和LeakyRelu算子 |11| [custom_op](./custom_op/) | 本样例介绍了自定义算子编译工程的实现方法，包括Add和LeakyRelu算子 |
12-| [custom_op_static_lib](./custom_op_static_lib/) | 本样例介绍了静态Aclnn调用的实现方法 |12+| [custom_op_static_lib](./custom_op_static_lib/) | 本样例介绍了静态Aclnn调用的实现方法 |
13+| [parallel_ops_package](./parallel_ops_package/) | 本样例介绍了自定义算子多vender并行编译工程的实现方法，包括Add和LeakyRelu算子 |
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/CMakeLists.txt+48-0
@@ -0,0 +1,48 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+cmake_minimum_required(VERSION 3.16.0)
13+project(parallel_ops_package_proj)
14+include(ExternalProject)
15+ 
16+set(ASCEND_COMPUTE_UNIT ascend910b ascend910_93 ascend950)
17+set(PACKAGE_TYPE "RUN")
18+set(INSTALL_TARGET "package")
19+ 
20+ExternalProject_Add(add_custom
21+ SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/add_custom
22+ CMAKE_CACHE_ARGS
23+ -DASCEND_COMPUTE_UNIT:STRING=${ASCEND_COMPUTE_UNIT}
24+ CMAKE_ARGS
25+ -Dvendor_name=add_custom
26+ -DCMAKE_INSTALL_PREFIX=${CMAKE_BINARY_DIR}/add_custom
27+ -DPACKAGE_TYPE=${PACKAGE_TYPE}
28+ BUILD_COMMAND
29+ ${CMAKE_COMMAND} --build <BINARY_DIR> --target binary
30+ COMMAND
31+ ${CMAKE_COMMAND} --build <BINARY_DIR> --target ${INSTALL_TARGET}
32+ INSTALL_COMMAND ""
33+)
34+ 
35+ExternalProject_Add(leaky_relu_custom
36+ SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/leaky_relu_custom
37+ CMAKE_CACHE_ARGS
38+ -DASCEND_COMPUTE_UNIT:STRING=${ASCEND_COMPUTE_UNIT}
39+ CMAKE_ARGS
40+ -Dvendor_name=leaky_relu_custom
41+ -DCMAKE_INSTALL_PREFIX=${CMAKE_BINARY_DIR}/leaky_relu_custom
42+ -DPACKAGE_TYPE=${PACKAGE_TYPE}
43+ BUILD_COMMAND
44+ ${CMAKE_COMMAND} --build <BINARY_DIR> --target binary
45+ COMMAND
46+ ${CMAKE_COMMAND} --build <BINARY_DIR> --target ${INSTALL_TARGET}
47+ INSTALL_COMMAND ""
48+)
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/README.md+129-0
@@ -0,0 +1,129 @@
1+# 自定义算子工程多 Vendor 并行编译、打包和部署样例
2+ 
3+## 概述
4+ 
5+本样例展示如何在一个顶层 CMake 工程中，使用 `ExternalProject_Add` 并行编译两个独立自定义算子工程：
6+- `add_custom`（AddCustom）
7+- `leaky_relu_custom`（LeakyReluCustom）
8+ 
9+每个子工程会分别完成自定义算子的编译、打包，并生成独立的 `custom_opp_*.run` 安装包。
10+ 
11+## 支持的产品
12+ 
13+本样例支持如下产品型号：
14+- Ascend 950PR/Ascend 950DT
15+- Atlas A3 训练系列产品/Atlas A3 推理系列产品
16+- Atlas A2 训练系列产品/Atlas A2 推理系列产品
17+- Atlas 200I/500 A2 推理产品
18+- Atlas 推理系列产品
19+ 
20+> 注意: 本样例中涉及多个算子示例，请以各个算子示例实际支持的产品型号为准。
21+ 
22+## 目录结构
23+ 
24+```text
25+parallel_ops_package
26+├── CMakeLists.txt
27+├── README.md
28+├── add_custom
29+│ ├── CMakeLists.txt
30+│ ├── framework
31+│ │ ├── CMakeLists.txt
32+│ │ └── tf_plugin
33+│ │ ├── CMakeLists.txt
34+│ │ └── tensorflow_add_custom_plugin.cc
35+│ ├── op_host
36+│ │ ├── CMakeLists.txt
37+│ │ └── add_custom
38+│ │ └── add_custom_host.cpp
39+│ └── op_kernel
40+│ ├── CMakeLists.txt
41+│ └── add_custom
42+│ ├── add_custom_kernel.cpp
43+│ └── add_custom_tiling.h
44+└── leaky_relu_custom
45+ ├── CMakeLists.txt
46+ ├── framework
47+ │ ├── CMakeLists.txt
48+ │ └── onnx_plugin
49+ │ ├── CMakeLists.txt
50+ │ └── leaky_relu_custom_plugin.cc
51+ ├── op_host
52+ │ ├── CMakeLists.txt
53+ │ └── leaky_relu_custom
54+ │ └── leaky_relu_custom_host.cpp
55+ └── op_kernel
56+ ├── CMakeLists.txt
57+ └── leaky_relu_custom
58+ ├── leaky_relu_custom_kernel.cpp
59+ └── leaky_relu_custom_tiling.h
60+```
61+ 
62+## 样例描述
63+ 
64+`parallel_ops_package` 与 `custom_op` 使用相同的 Add/LeakyRelu 样例描述，本文不重复维护，请参考：
65+ 
66+- [custom_op/README.md 的“样例描述”章节](../custom_op/README.md#样例描述)
67+ 
68+## 样例规格描述
69+ 
70+`parallel_ops_package` 与 `custom_op` 使用相同的 Add/LeakyRelu 规格，本文不重复维护规格表，请参考：
71+ 
72+- [custom_op/README.md 的“样例规格描述”章节](../custom_op/README.md#样例规格描述)
73+ 
74+## 代码实现介绍
75+ 
76+`parallel_ops_package` 的算子实现与 `custom_op` 保持一致，本文不重复维护实现细节，请参考：
77+ 
78+- [custom_op/README.md 的“代码实现介绍”章节](../custom_op/README.md#代码实现介绍)
79+ 
80+## 编译运行
81+ 
82+在本样例根目录下执行如下步骤，编译、打包并部署自定义样例包。
83+ 
84+- 配置环境变量
85+ 
86+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
87+ - 默认路径，root用户安装CANN软件包
88+ 
89+ ```bash
90+ source /usr/local/Ascend/cann/set_env.sh
91+ ```
92+ 
93+ - 默认路径，非root用户安装CANN软件包
94+ 
95+ ```bash
96+ source $HOME/Ascend/cann/set_env.sh
97+ ```
98+ 
99+ - 指定路径install_path，安装CANN软件包
100+ 
101+ ```bash
102+ source ${install_path}/cann/set_env.sh
103+ ```
104+ 
105+- 编译、打包样例并部署两个算子包
106+ 
107+ ```bash
108+ cmake -S . -B build
109+ cmake --build build -j
110+ # add_custom 包
111+ ./add_custom/custom_opp_*.run
112+ 
113+ # leaky_relu_custom 包
114+ ./leaky_relu_custom/custom_opp_*.run
115+ ```
116+ 
117+ 执行结果如下，说明执行成功。
118+ 
119+ ```log
120+ SUCCESS
121+ ```
122+ 
123+## 构建结果说明
124+ 
125+顶层工程会在 `build/` 下生成两个子目录：
126+- `build/add_custom/`：AddCustom 的中间产物与安装包
127+- `build/leaky_relu_custom/`：LeakyReluCustom 的中间产物与安装包
128+ 
129+这两个目录互相独立，便于多 Vendor 场景下并行开发与发布。
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/CMakeLists.txt+35-0
@@ -0,0 +1,35 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+cmake_minimum_required(VERSION 3.16.0)
13+project(opp)
14+ 
15+find_package(ASC REQUIRED)
16+set(package_name ${vendor_name})
17+set(package_type ${PACKAGE_TYPE})
18+ 
19+npu_op_package(${package_name}
20+ TYPE ${package_type}
21+ CONFIG
22+ INSTALL_PATH ${CMAKE_INSTALL_PREFIX}
23+)
24+ 
25+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/framework")
26+ add_subdirectory(framework)
27+endif()
28+ 
29+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_host")
30+ add_subdirectory(op_host)
31+endif()
32+ 
33+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_kernel")
34+ add_subdirectory(op_kernel)
35+endif()
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/framework/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/tf_plugin")
13+ add_subdirectory(tf_plugin)
14+endif()
15+ 
16+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/onnx_plugin")
17+ add_subdirectory(onnx_plugin)
18+endif()
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/framework/tf_plugin/CMakeLists.txt+21-0
@@ -0,0 +1,21 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+aux_source_directory(${CMAKE_CURRENT_SOURCE_DIR} plugin_srcs)
13+ 
14+npu_op_library(cust_tf_parsers TF_PLUGIN
15+ ${plugin_srcs}
16+)
17+ 
18+npu_op_package_add(${package_name}
19+ LIBRARY
20+ cust_tf_parsers
21+)
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/framework/tf_plugin/tensorflow_add_custom_plugin.cc+19-0
@@ -0,0 +1,19 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+#include "register/register.h"
12+ 
13+namespace domi {
14+// register op info to GE
15+REGISTER_CUSTOM_OP("AddCustom")
16+ .FrameworkType(TENSORFLOW) // type: CAFFE, TENSORFLOW
17+ .OriginOpType("AddCustom") // name in tf module
18+ .ParseParamsByOperatorFn(AutoMappingByOpFn);
19+} // namespace domi
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/op_host/CMakeLists.txt+42-0
@@ -0,0 +1,42 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+file(GLOB host_ops_srcs
13+ add_custom/add_custom_host.cpp
14+)
15+npu_op_code_gen(
16+ SRC ${host_ops_srcs}
17+ PACKAGE ${package_name}
18+ OUT_DIR ${ASCEND_AUTOGEN_PATH}
19+)
20+npu_op_library(cust_optiling TILING
21+ ${host_ops_srcs}
22+)
23+ 
24+file(GLOB autogen_aclnn_src ${ASCEND_AUTOGEN_PATH}/aclnn_*.cpp)
25+set_source_files_properties(${autogen_aclnn_src} PROPERTIES GENERATED TRUE)
26+npu_op_library(cust_opapi ACLNN
27+ ${autogen_aclnn_src}
28+)
29+ 
30+file(GLOB proto_src ${ASCEND_AUTOGEN_PATH}/op_proto.cc)
31+set_source_files_properties(${proto_src} PROPERTIES GENERATED TRUE)
32+npu_op_library(cust_op_proto GRAPH
33+ ${host_ops_srcs}
34+ ${proto_src}
35+)
36+ 
37+npu_op_package_add(${package_name}
38+ LIBRARY
39+ cust_optiling
40+ cust_opapi
41+ cust_op_proto
42+)
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/op_host/add_custom/add_custom_host.cpp+78-0
@@ -0,0 +1,78 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+#include "../../op_kernel/add_custom/add_custom_tiling.h"
13+#include "register/op_def_registry.h"
14+ 
15+namespace optiling {
16+const uint32_t NUM_BLOCKS = 8;
17+const uint32_t TILE_NUM = 8;
18+static ge::graphStatus TilingFunc(gert::TilingContext *context)
19+{
20+ AddCustomTilingData *tiling = context->GetTilingData<AddCustomTilingData>();
21+ uint32_t totalLength = context->GetInputShape(0)->GetOriginShape().GetShapeSize();
22+ context->SetBlockDim(NUM_BLOCKS);
23+ tiling->totalLength = totalLength;
24+ tiling->tileNum = TILE_NUM;
25+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
26+ currentWorkspace[0] = 0;
27+ return ge::GRAPH_SUCCESS;
28+}
29+} // namespace optiling
30+ 
31+namespace ge {
32+static graphStatus InferShape(gert::InferShapeContext *context)
33+{
34+ const gert::Shape *x1_shape = context->GetInputShape(0);
35+ gert::Shape *y_shape = context->GetOutputShape(0);
36+ *y_shape = *x1_shape;
37+ return GRAPH_SUCCESS;
38+}
39+ 
40+static graphStatus InferDataType(gert::InferDataTypeContext *context)
41+{
42+ const auto inputDataType = context->GetInputDataType(0);
43+ context->SetOutputDataType(0, inputDataType);
44+ return ge::GRAPH_SUCCESS;
45+}
46+} // namespace ge
47+ 
48+namespace ops {
49+class AddCustom : public OpDef {
50+public:
51+ explicit AddCustom(const char *name) : OpDef(name)
52+ {
53+ this->Input("x")
54+ .ParamType(REQUIRED)
55+ .DataType({ge::DT_FLOAT16})
56+ .Format({ge::FORMAT_ND});
57+ this->Input("y")
58+ .ParamType(REQUIRED)
59+ .DataType({ge::DT_FLOAT16})
60+ .Format({ge::FORMAT_ND});
61+ this->Output("z")
62+ .ParamType(REQUIRED)
63+ .DataType({ge::DT_FLOAT16})
64+ .Format({ge::FORMAT_ND});
65+ 
66+ this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);
67+ this->AICore()
68+ .SetTiling(optiling::TilingFunc)
69+ // at least one soc version must be configured.
70+ .AddConfig("ascend310p")
71+ .AddConfig("ascend310b")
72+ .AddConfig("ascend910b")
73+ .AddConfig("ascend910_93")
74+ .AddConfig("ascend950");
75+ }
76+};
77+OP_ADD(AddCustom);
78+} // namespace ops
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/op_kernel/CMakeLists.txt+26-0
@@ -0,0 +1,26 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+npu_op_kernel_sources(ascendc_kernels
13+ OP_TYPE AddCustom
14+ KERNEL_DIR add_custom
15+ KERNEL_FILE add_custom_kernel.cpp
16+)
17+ 
18+npu_op_kernel_library(ascendc_kernels
19+ SRC_BASE ${CMAKE_CURRENT_SOURCE_DIR}
20+ TILING_LIBRARY cust_optiling
21+)
22+ 
23+npu_op_package_add(${package_name}
24+ LIBRARY
25+ ascendc_kernels
26+)
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/op_kernel/add_custom/add_custom_kernel.cpp+91-0
@@ -0,0 +1,91 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+#include "kernel_operator.h"
13+#include "add_custom_tiling.h"
14+constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue
15+ 
16+class KernelAdd {
17+public:
18+ __aicore__ inline KernelAdd() {}
19+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum, AscendC::TPipe* pipeIn)
20+ {
21+ ascendc_assert(tileNum != 0, "tileNum can not be zero.\n");
22+ this->pipe = pipeIn;
23+ this->blockLength = totalLength / AscendC::GetBlockNum();
24+ this->tileNum = tileNum;
25+ this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
26+ 
27+ xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
28+ yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
29+ zGm.SetGlobalBuffer((__gm__ DTYPE_Z *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
30+ this->pipe->InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
31+ this->pipe->InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));
32+ this->pipe->InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Z));
33+ }
34+ __aicore__ inline void Process()
35+ {
36+ int32_t loopCount = this->tileNum * BUFFER_NUM;
37+ for (int32_t i = 0; i < loopCount; i++) {
38+ CopyIn(i);
39+ Compute(i);
40+ CopyOut(i);
41+ }
42+ }
43+ 
44+private:
45+ __aicore__ inline void CopyIn(int32_t progress)
46+ {
47+ AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
48+ AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.AllocTensor<DTYPE_Y>();
49+ AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
50+ AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);
51+ inQueueX.EnQue(xLocal);
52+ inQueueY.EnQue(yLocal);
53+ }
54+ __aicore__ inline void Compute(int32_t progress)
55+ {
56+ AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
57+ AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.DeQue<DTYPE_Y>();
58+ AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.AllocTensor<DTYPE_Z>();
59+ AscendC::Add(zLocal, xLocal, yLocal, this->tileLength);
60+ outQueueZ.EnQue<DTYPE_Z>(zLocal);
61+ inQueueX.FreeTensor(xLocal);
62+ inQueueY.FreeTensor(yLocal);
63+ }
64+ __aicore__ inline void CopyOut(int32_t progress)
65+ {
66+ AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.DeQue<DTYPE_Z>();
67+ AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
68+ outQueueZ.FreeTensor(zLocal);
69+ }
70+ 
71+private:
72+ AscendC::TPipe* pipe = nullptr;
73+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
74+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
75+ AscendC::GlobalTensor<DTYPE_X> xGm;
76+ AscendC::GlobalTensor<DTYPE_Y> yGm;
77+ AscendC::GlobalTensor<DTYPE_Z> zGm;
78+ uint32_t blockLength;
79+ uint32_t tileNum;
80+ uint32_t tileLength;
81+};
82+ 
83+extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
84+{
85+ REGISTER_TILING_DEFAULT(AddCustomTilingData);
86+ GET_TILING_DATA(tilingData, tiling);
87+ AscendC::TPipe pipe;
88+ KernelAdd op;
89+ op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum, &pipe);
90+ op.Process();
91+}
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/add_custom/op_kernel/add_custom/add_custom_tiling.h+21-0
@@ -0,0 +1,21 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+#ifndef ADD_CUSTOM_TILING_H
13+#define ADD_CUSTOM_TILING_H
14+#include <cstdint>
15+ 
16+struct AddCustomTilingData {
17+ uint32_t totalLength;
18+ uint32_t tileNum;
19+};
20+ 
21+#endif // ADD_CUSTOM_TILING_H
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/CMakeLists.txt+35-0
@@ -0,0 +1,35 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+cmake_minimum_required(VERSION 3.16.0)
13+project(opp)
14+ 
15+find_package(ASC REQUIRED)
16+set(package_name ${vendor_name})
17+set(package_type ${PACKAGE_TYPE})
18+ 
19+npu_op_package(${package_name}
20+ TYPE ${package_type}
21+ CONFIG
22+ INSTALL_PATH ${CMAKE_INSTALL_PREFIX}
23+)
24+ 
25+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/framework")
26+ add_subdirectory(framework)
27+endif()
28+ 
29+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_host")
30+ add_subdirectory(op_host)
31+endif()
32+ 
33+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_kernel")
34+ add_subdirectory(op_kernel)
35+endif()
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/framework/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/tf_plugin")
13+ add_subdirectory(tf_plugin)
14+endif()
15+ 
16+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/onnx_plugin")
17+ add_subdirectory(onnx_plugin)
18+endif()
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/framework/onnx_plugin/CMakeLists.txt+50-0
@@ -0,0 +1,50 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+aux_source_directory(${CMAKE_CURRENT_SOURCE_DIR} plugin_srcs)
13+ 
14+set(NLOHMANN_JSON_URL "https://gitcode.com/cann-src-third-party/json/releases/download/v3.11.3/include.zip")
15+set(NLOHMANN_JSON_DIR "${CMAKE_CURRENT_BINARY_DIR}/third_party/json")
16+set(NLOHMANN_JSON_ARCHIVE "${NLOHMANN_JSON_DIR}/include.zip")
17+set(NLOHMANN_JSON_HEADER "${NLOHMANN_JSON_DIR}/include/nlohmann/json.hpp")
18+ 
19+file(MAKE_DIRECTORY "${NLOHMANN_JSON_DIR}")
20+if(NOT EXISTS "${NLOHMANN_JSON_HEADER}")
21+ file(DOWNLOAD "${NLOHMANN_JSON_URL}" "${NLOHMANN_JSON_ARCHIVE}"
22+ STATUS download_status
23+ LOG download_log
24+ )
25+ list(GET download_status 0 download_code)
26+ if(NOT download_code EQUAL 0)
27+ message(FATAL_ERROR "Failed to download nlohmann/json include.zip: ${download_status}\n${download_log}")
28+ endif()
29+ execute_process(
30+ COMMAND ${CMAKE_COMMAND} -E tar xf "${NLOHMANN_JSON_ARCHIVE}" --format=zip
31+ WORKING_DIRECTORY "${NLOHMANN_JSON_DIR}"
32+ RESULT_VARIABLE extract_code
33+ OUTPUT_VARIABLE extract_out
34+ ERROR_VARIABLE extract_err
35+ )
36+ if(NOT extract_code EQUAL 0)
37+ message(FATAL_ERROR "Failed to extract nlohmann/json include.zip: ${extract_code}\n${extract_out}\n${extract_err}")
38+ endif()
39+endif()
40+ 
41+npu_op_library(cust_onnx_parsers ONNX_PLUGIN
42+ ${plugin_srcs}
43+)
44+ 
45+target_include_directories(cust_onnx_parsers PRIVATE "${NLOHMANN_JSON_DIR}/include")
46+ 
47+npu_op_package_add(${package_name}
48+ LIBRARY
49+ cust_onnx_parsers
50+)
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/framework/onnx_plugin/leaky_relu_custom_plugin.cc+53-0
@@ -0,0 +1,53 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+#include "graph/operator.h"
12+#include "register/register.h"
13+#include <nlohmann/json.hpp>
14+ 
15+using namespace ge;
16+using json = nlohmann::json;
17+ 
18+namespace domi {
19+namespace {
20+const int kTypeFloat = 1;
21+}
22+Status ParseOnnxParamsLeakyReluCustom(const ge::Operator& op_src, ge::Operator& op_dest) {
23+ // trans op_src to op_dest
24+ // if op_src get required attr failed, need to return Failed
25+ // if op_src get optional attr failed, need to return Failed or set a default value
26+ float negative_slope = 0.01f;
27+ string negative_slope_str;
28+ AscendString attrs_string;
29+ if (ge::GRAPH_SUCCESS == op_src.GetAttr("attribute", attrs_string)) {
30+ json attrs = json::parse(attrs_string.GetString());
31+ for (json attr : attrs["attribute"]) {
32+ if (attr["name"] == "alpha" && attr["type"] == kTypeFloat) {
33+ negative_slope_str = attr["f"]; // float type in json has accuracy loss, so we use string type to store it
34+ negative_slope = atof(negative_slope_str.c_str());
35+ }
36+ }
37+ }
38+ 
39+ op_dest.SetAttr("negative_slope", negative_slope);
40+ return SUCCESS;
41+}
42+ 
43+REGISTER_CUSTOM_OP("LeakyReluCustom")
44+ .FrameworkType(ONNX)
45+ .OriginOpType({ge::AscendString("ai.onnx::8::LeakyRelu"),
46+ ge::AscendString("ai.onnx::9::LeakyRelu"),
47+ ge::AscendString("ai.onnx::10::LeakyRelu"),
48+ ge::AscendString("ai.onnx::11::LeakyRelu"),
49+ ge::AscendString("ai.onnx::12::LeakyRelu"),
50+ ge::AscendString("ai.onnx::13::LeakyRelu")})
51+ .ParseParamsByOperatorFn(ParseOnnxParamsLeakyReluCustom)
52+ .ImplyType(ImplyType::TVM);
53+} // namespace domi
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/op_host/CMakeLists.txt+42-0
@@ -0,0 +1,42 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+file(GLOB host_ops_srcs
13+ leaky_relu_custom/leaky_relu_custom_host.cpp
14+)
15+npu_op_code_gen(
16+ SRC ${host_ops_srcs}
17+ PACKAGE ${package_name}
18+ OUT_DIR ${ASCEND_AUTOGEN_PATH}
19+)
20+npu_op_library(cust_optiling TILING
21+ ${host_ops_srcs}
22+)
23+ 
24+file(GLOB autogen_aclnn_src ${ASCEND_AUTOGEN_PATH}/aclnn_*.cpp)
25+set_source_files_properties(${autogen_aclnn_src} PROPERTIES GENERATED TRUE)
26+npu_op_library(cust_opapi ACLNN
27+ ${autogen_aclnn_src}
28+)
29+ 
30+file(GLOB proto_src ${ASCEND_AUTOGEN_PATH}/op_proto.cc)
31+set_source_files_properties(${proto_src} PROPERTIES GENERATED TRUE)
32+npu_op_library(cust_op_proto GRAPH
33+ ${host_ops_srcs}
34+ ${proto_src}
35+)
36+ 
37+npu_op_package_add(${package_name}
38+ LIBRARY
39+ cust_optiling
40+ cust_opapi
41+ cust_op_proto
42+)
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/op_host/leaky_relu_custom/leaky_relu_custom_host.cpp+73-0
@@ -0,0 +1,73 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+#include "../../op_kernel/leaky_relu_custom/leaky_relu_custom_tiling.h"
13+#include "register/op_def_registry.h"
14+ 
15+namespace optiling {
16+const uint32_t NUM_BLOCKS = 8;
17+const uint32_t TILE_NUM = 16;
18+ 
19+static ge::graphStatus TilingFunc(gert::TilingContext *context)
20+{
21+ LeakyReluCustomTilingData *tiling = context->GetTilingData<LeakyReluCustomTilingData>();
22+ uint32_t totalLength = context->GetInputShape(0)->GetOriginShape().GetShapeSize();
23+ const gert::RuntimeAttrs *attrs = context->GetAttrs();
24+ const float *negativeSlope = attrs->GetAttrPointer<float>(0);
25+ 
26+ context->SetBlockDim(NUM_BLOCKS);
27+ tiling->totalLength = totalLength;
28+ tiling->tileNum = TILE_NUM;
29+ tiling->negativeSlope = *negativeSlope;
30+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
31+ currentWorkspace[0] = 0;
32+ return ge::GRAPH_SUCCESS;
33+}
34+} // namespace optiling
35+ 
36+namespace ge {
37+static ge::graphStatus InferShape(gert::InferShapeContext *context)
38+{
39+ const gert::Shape *xShape = context->GetInputShape(0);
40+ gert::Shape *yShape = context->GetOutputShape(0);
41+ *yShape = *xShape;
42+ return GRAPH_SUCCESS;
43+}
44+static ge::graphStatus InferDataType(gert::InferDataTypeContext *context)
45+{
46+ const ge::DataType xDtype = context->GetInputDataType(0);
47+ context->SetOutputDataType(0, xDtype);
48+ return GRAPH_SUCCESS;
49+}
50+} // namespace ge
51+ 
52+namespace ops {
53+class LeakyReluCustom : public OpDef {
54+public:
55+ LeakyReluCustom(const char *name) : OpDef(name)
56+ {
57+ this->Input("x")
58+ .ParamType(REQUIRED)
59+ .DataType({ge::DT_FLOAT})
60+ .Format({ge::FORMAT_ND});
61+ this->Output("y")
62+ .ParamType(REQUIRED)
63+ .DataType({ge::DT_FLOAT})
64+ .Format({ge::FORMAT_ND});
65+ this->Attr("negative_slope").AttrType(OPTIONAL).Float(0.0f);
66+ this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);
67+ this->AICore()
68+ .SetTiling(optiling::TilingFunc)
69+ .AddConfig("ascend910b");
70+ }
71+};
72+OP_ADD(LeakyReluCustom);
73+} // namespace ops
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/op_kernel/CMakeLists.txt+26-0
@@ -0,0 +1,26 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+npu_op_kernel_sources(ascendc_kernels
13+ OP_TYPE LeakyReluCustom
14+ KERNEL_DIR leaky_relu_custom
15+ KERNEL_FILE leaky_relu_custom_kernel.cpp
16+)
17+ 
18+npu_op_kernel_library(ascendc_kernels
19+ SRC_BASE ${CMAKE_CURRENT_SOURCE_DIR}
20+ TILING_LIBRARY cust_optiling
21+)
22+ 
23+npu_op_package_add(${package_name}
24+ LIBRARY
25+ ascendc_kernels
26+)
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/op_kernel/leaky_relu_custom/leaky_relu_custom_kernel.cpp+108-0
@@ -0,0 +1,108 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+#include "kernel_operator.h"
13+#include "leaky_relu_custom_tiling.h"
14+constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue
15+ 
16+class KernelLeakyRelu {
17+public:
18+ __aicore__ inline KernelLeakyRelu() {}
19+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength, uint32_t tileNum, float negativeSlope, AscendC::TPipe* pipeIn)
20+ {
21+ ascendc_assert(tileNum != 0, "tileNum can not be zero.\n");
22+ this->pipe = pipeIn;
23+ this->blockLength = totalLength / AscendC::GetBlockNum();
24+ this->tileNum = tileNum;
25+ this->negativeSlope = static_cast<float>(negativeSlope);
26+ this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
27+ 
28+ // get start index for current core, core parallel
29+ xGm.SetGlobalBuffer((__gm__ float *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
30+ yGm.SetGlobalBuffer((__gm__ float *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
31+ // pipe alloc memory to queue, the unit is Bytes
32+ this->pipe->InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
33+ this->pipe->InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
34+ this->pipe->InitBuffer(tmpBuffer1, this->tileLength * sizeof(float));
35+ this->pipe->InitBuffer(tmpBuffer2, this->tileLength * sizeof(float));
36+ }
37+ __aicore__ inline void Process()
38+ {
39+ // loop count need to be doubled, due to double buffer
40+ int32_t loopCount = this->tileNum * BUFFER_NUM;
41+ // tiling strategy, pipeline parallel
42+ for (int32_t i = 0; i < loopCount; i++) {
43+ CopyIn(i);
44+ Compute(i);
45+ CopyOut(i);
46+ }
47+ }
48+ 
49+private:
50+ __aicore__ inline void CopyIn(int32_t progress)
51+ {
52+ // alloc tensor from queue memory
53+ AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
54+ // copy progress_th tile from global tensor to local tensor
55+ AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
56+ // enque input tensors to VECIN queue
57+ inQueueX.EnQue(xLocal);
58+ }
59+ __aicore__ inline void Compute(int32_t progress)
60+ {
61+ // deque input tensors from VECIN queue
62+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
63+ AscendC::LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
64+ AscendC::LocalTensor<float> tmpTensor1 = tmpBuffer1.Get<float>();
65+ AscendC::LocalTensor<float> tmpTensor2 = tmpBuffer2.Get<float>();
66+ float inputVal = 0.0f;
67+ AscendC::Maxs(tmpTensor1, xLocal, inputVal, this->tileLength);
68+ AscendC::Mins(tmpTensor2, xLocal, inputVal, this->tileLength);
69+ AscendC::Muls(tmpTensor2, tmpTensor2, this->negativeSlope, this->tileLength);
70+ AscendC::Add(yLocal, tmpTensor1, tmpTensor2, this->tileLength);
71+ // enque the output tensor to VECOUT queue
72+ outQueueY.EnQue<float>(yLocal);
73+ // free input tensors for reuse
74+ inQueueX.FreeTensor(xLocal);
75+ }
76+ __aicore__ inline void CopyOut(int32_t progress)
77+ {
78+ // deque output tensor from VECOUT queue
79+ AscendC::LocalTensor<float> yLocal = outQueueY.DeQue<float>();
80+ // copy progress_th tile from local tensor to global tensor
81+ AscendC::DataCopy(yGm[progress * this->tileLength], yLocal, this->tileLength);
82+ // free output tensor for reuse
83+ outQueueY.FreeTensor(yLocal);
84+ }
85+ 
86+private:
87+ AscendC::TPipe* pipe = nullptr;
88+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpBuffer1, tmpBuffer2;
89+ // create queues for input, in this case depth is equal to buffer num
90+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX;
91+ // create queue for output, in this case depth is equal to buffer num
92+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueY;
93+ AscendC::GlobalTensor<float> xGm, yGm;
94+ uint32_t blockLength;
95+ uint32_t tileNum;
96+ uint32_t tileLength;
97+ float negativeSlope;
98+};
99+ 
100+extern "C" __global__ __aicore__ void leaky_relu_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
101+{
102+ REGISTER_TILING_DEFAULT(LeakyReluCustomTilingData);
103+ GET_TILING_DATA(tiling_data, tiling);
104+ AscendC::TPipe pipe;
105+ KernelLeakyRelu op;
106+ op.Init(x, y, tiling_data.totalLength, tiling_data.tileNum, tiling_data.negativeSlope, &pipe);
107+ op.Process();
108+}
Aexamples/01_simd_cpp_api/02_features/00_compilation/parallel_ops_package/leaky_relu_custom/op_kernel/leaky_relu_custom/leaky_relu_custom_tiling.h+21-0
@@ -0,0 +1,21 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+#ifndef LEAKY_RELU_CUSTOM_TILING_H
13+#define LEAKY_RELU_CUSTOM_TILING_H
14+#include <cstdint>
15+ 
16+struct LeakyReluCustomTilingData {
17+ uint32_t totalLength;
18+ uint32_t tileNum;
19+ float negativeSlope;
20+};
21+#endif // LEAKY_RELU_CUSTOM_TILING_H