Star115
1505
代码
代码
Issues42
Pull Requests645
Actions
项目讨论
Wiki
分析
Star115
1505
已开启
【代码侦探Challenge 04 · DivCustomTemplate】完成 DivCustomTemplate 工程化算子开发（msopgen 工程 + 核函数/tiling 实现 + ACLNN 验证） #1939
A HOLLOWcsw创建于 5 天前
【代码侦探Challenge 04 · DivCustomTemplate】完成 DivCustomTemplate 工程化算子开发（msopgen 工程 + 核函数/tiling 实现 + ACLNN 验证） #1939
已开启
A HOLLOWcsw创建于 5 天前
从weixin_53562472/cann-outreachchallenge04-divcustomtemplate合入到cann/cann-outreachmaster
共 17 个文件变更+889-0
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/.gitignore+4-0
@@ -0,0 +1,4 @@
1+# 构建产物与本地环境文件，不纳入提交
2+custom_op/build_out/
3+execute_div_op
4+__pycache__/
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/CMakeLists.txt+31-0
@@ -0,0 +1,31 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 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+find_package(ASC REQUIRED)
15+set(package_name ${vendor_name})
16+ 
17+npu_op_package(${package_name}
18+ TYPE RUN
19+ CONFIG
20+ INSTALL_PATH ${CMAKE_BINARY_DIR}/
21+)
22+ 
23+if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/framework)
24+ add_subdirectory(framework)
25+endif()
26+if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/op_host)
27+ add_subdirectory(op_host)
28+endif()
29+if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/op_kernel)
30+ add_subdirectory(op_kernel)
31+endif()
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/CMakePresets.json+67-0
@@ -0,0 +1,67 @@
1+{
2+ "version": 1,
3+ "cmakeMinimumRequired": {
4+ "major": 3,
5+ "minor": 16,
6+ "patch": 0
7+ },
8+ "configurePresets": [
9+ {
10+ "name": "default",
11+ "displayName": "Default Config",
12+ "description": "Default build using Unix Makefiles generator",
13+ "generator": "Unix Makefiles",
14+ "binaryDir": "${sourceDir}/build_out",
15+ "cacheVariables": {
16+ "CMAKE_BUILD_TYPE": {
17+ "type": "STRING",
18+ "value": "Release"
19+ },
20+ "ENABLE_SOURCE_PACKAGE": {
21+ "type": "BOOL",
22+ "value": "True"
23+ },
24+ "ENABLE_BINARY_PACKAGE": {
25+ "type": "BOOL",
26+ "value": "True"
27+ },
28+ "ASCEND_COMPUTE_UNIT": {
29+ "type": "STRING",
30+ "value": "ascend910b"
31+ },
32+ "ENABLE_TEST": {
33+ "type": "BOOL",
34+ "value": "True"
35+ },
36+ "vendor_name": {
37+ "type": "STRING",
38+ "value": "customize"
39+ },
40+ "ASCEND_CANN_PACKAGE_PATH": {
41+ "type": "PATH",
42+ "value": "/usr/local/Ascend/cann-9.1.0-beta.1"
43+ },
44+ "ASCEND_PYTHON_EXECUTABLE": {
45+ "type": "STRING",
46+ "value": "python3"
47+ },
48+ "CMAKE_INSTALL_PREFIX": {
49+ "type": "PATH",
50+ "value": "${sourceDir}/build_out"
51+ },
52+ "ENABLE_CROSS_COMPILE": {
53+ "type": "BOOL",
54+ "value": "False"
55+ },
56+ "CMAKE_CROSS_PLATFORM_COMPILER": {
57+ "type": "PATH",
58+ "value": "/usr/bin/aarch64-linux-gnu-g++"
59+ },
60+ "ASCEND_PACK_SHARED_LIBRARY": {
61+ "type": "BOOL",
62+ "value": "False"
63+ }
64+ }
65+ }
66+ ]
67+}
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/build.sh+40-0
@@ -0,0 +1,40 @@
1+#!/bin/bash
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------------------------------------
11+ 
12+if [ -z "$BASE_LIBS_PATH" ]; then
13+ if [ -z "$ASCEND_HOME_PATH" ]; then
14+ if [ -z "$ASCEND_AICPU_PATH" ]; then
15+ echo "please set env."
16+ exit 1
17+ else
18+ export ASCEND_HOME_PATH=$ASCEND_AICPU_PATH
19+ fi
20+ else
21+ export ASCEND_HOME_PATH=$ASCEND_HOME_PATH
22+ fi
23+else
24+ export ASCEND_HOME_PATH=$BASE_LIBS_PATH
25+fi
26+echo "using ASCEND_HOME_PATH: $ASCEND_HOME_PATH"
27+script_path=$(realpath $(dirname $0))
28+ 
29+BUILD_DIR="build_out"
30+mkdir -p build_out
31+rm -rf build_out/*
32+opts=$(python3 $ASCEND_HOME_PATH/tools/tikcpp/ascendc_kernel_cmake/fwk_modules/util/preset_parse.py $script_path/CMakePresets.json)
33+cmake_version=$(cmake --version | grep "cmake version" | awk '{print $3}')
34+ 
35+if [ "$cmake_version" \< "3.19.0" ] ; then
36+ cmake -S . -B "$BUILD_DIR" $opts
37+else
38+ cmake -S . -B "$BUILD_DIR" --preset=default
39+fi
40+cmake --build "$BUILD_DIR" --target binary package -j$(nproc)
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/framework/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 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()
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/framework/tf_plugin/CMakeLists.txt+24-0
@@ -0,0 +1,24 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 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+if(NOT plugin_srcs)
14+ return()
15+endif()
16+ 
17+npu_op_library(cust_tf_parsers TF_PLUGIN
18+ ${plugin_srcs}
19+)
20+ 
21+npu_op_package_add(${package_name}
22+ LIBRARY
23+ cust_tf_parsers
24+)
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/framework/tf_plugin/tensorflow_div_custom_template_plugin.cc+25-0
@@ -0,0 +1,25 @@
1+/* -------------------------------------------------------------------------
2+ * This file is part of the MindStudio project.
3+ * Copyright (c) 2025 Huawei Technologies Co.,Ltd.
4+ *
5+ * MindStudio is licensed under Mulan PSL v2.
6+ * You can use this software according to the terms and conditions of the Mulan PSL v2.
7+ * You may obtain a copy of Mulan PSL v2 at:
8+ *
9+ * http://license.coscl.org.cn/MulanPSL2
10+ *
11+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND,
12+ * EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT,
13+ * MERCHANTABILITY OR FIT FOR A PARTICULAR PURPOSE.
14+ * See the Mulan PSL v2 for more details.
15+ * ------------------------------------------------------------------------- */
16+ 
17+#include "register/register.h"
18+ 
19+namespace domi {
20+// register op info to GE
21+REGISTER_CUSTOM_OP("DivCustomTemplate")
22+ .FrameworkType(TENSORFLOW) // type: CAFFE, TENSORFLOW
23+ .OriginOpType("DivCustomTemplate") // name in tf module
24+ .ParseParamsByOperatorFn(AutoMappingByOpFn);
25+} // namespace domi
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/op_host/CMakeLists.txt+58-0
@@ -0,0 +1,58 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 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} ops_srcs)
13+npu_op_code_gen(
14+ SRC ${ops_srcs}
15+ PACKAGE ${package_name}
16+ OUT_DIR ${ASCEND_AUTOGEN_PATH}
17+ JOIN_OP_DEF True
18+)
19+ 
20+file(GLOB autogen_aclnn_src ${ASCEND_AUTOGEN_PATH}/aclnn_*.cpp)
21+set_source_files_properties(${autogen_aclnn_src} PROPERTIES GENERATED TRUE)
22+npu_op_library(cust_opapi ACLNN
23+ ${autogen_aclnn_src}
24+)
25+ 
26+target_compile_options(cust_opapi PRIVATE
27+ -fvisibility=hidden
28+)
29+ 
30+file(GLOB group_proto_src ${ASCEND_AUTOGEN_PATH}/group_op_proto/*.cc)
31+file(GLOB proto_src ${ASCEND_AUTOGEN_PATH}/op_proto.cc)
32+set_source_files_properties(${group_proto_src} PROPERTIES GENERATED TRUE)
33+set_source_files_properties(${proto_src} PROPERTIES GENERATED TRUE)
34+npu_op_library(cust_op_proto GRAPH
35+ ${ops_srcs}
36+ ${group_proto_src}
37+ ${proto_src}
38+)
39+target_compile_options(cust_op_proto PRIVATE
40+ -fvisibility=hidden
41+)
42+ 
43+file(GLOB fallback_src ${ASCEND_AUTOGEN_PATH}/fallback_*.cpp)
44+set_source_files_properties(${fallback_src} PROPERTIES GENERATED TRUE)
45+npu_op_library(cust_optiling TILING
46+ ${ops_srcs}
47+ ${fallback_src}
48+)
49+target_compile_options(cust_optiling PRIVATE
50+ -fvisibility=hidden
51+)
52+ 
53+npu_op_package_add(${package_name}
54+ LIBRARY
55+ cust_optiling
56+ cust_opapi
57+ cust_op_proto
58+)
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/op_host/div_custom_template.cpp+97-0
@@ -0,0 +1,97 @@
1+/* div_custom_template.cpp (op_host)
2+ *
3+ * DivCustomTemplate 算子 host 侧实现：
4+ * - TilingFunc：计算输入总元素数、记录输入数据类型、设置核内 tile 切分数
5+ * 与参与计算的核数（BlockDim = 8，保证整除与 32B 对齐）；
6+ * - InferShape / InferDataType：输出 z 与输入 x 的形状、数据类型一致；
7+ * - OpDef 注册：输入 x/y、输出 z 均支持 float16 / float32，ND 格式，
8+ * 注册到 ascend910b 平台。
9+ */
10+ 
11+#include "../op_kernel/div_custom_template_tiling.h"
12+#include "register/op_def_registry.h"
13+ 
14+ 
15+namespace optiling {
16+// 参与计算的核数：本期任务数据量为 (8, 2048) = 16384 个元素，取 8 核均分，
17+// 每个核 2048 个元素，可继续被 tileNum * BUFFER_NUM = 16 整除（每片 128 个元素，
18+// float16 下 256B / float32 下 512B，均满足 32B 对齐要求）。
19+// 注：ascend910b1 实际有 48 个向量核，本算子数据量小，按 8 核切分即可保证
20+// 整除与对齐（与 Ascend C 官方 AddCustom 样例的切分方式一致）。
21+constexpr uint32_t BLOCK_DIM = 8;
22+ 
23+static ge::graphStatus TilingFunc(gert::TilingContext* context)
24+{
25+ DivCustomTemplateTilingData *tiling = context->GetTilingData<DivCustomTemplateTilingData>();
26+ 
27+ // 1. 计算输入张量的总元素数（各维度连乘）
28+ const gert::StorageShape* x1_shape = context->GetInputShape(0);
29+ uint32_t data_sz = 1;
30+ for (int i = 0; i < x1_shape->GetStorageShape().GetDimNum(); i++) {
31+ data_sz *= x1_shape->GetStorageShape().GetDim(i);
32+ }
33+ tiling->totalLength = data_sz;
34+ // 2. 核内 tile 切分数（配合双缓冲流水）
35+ tiling->tileNum = DIV_TILE_NUM;
36+ // 3. 记录输入数据类型，供 kernel 侧分发核函数模板实例
37+ const ge::DataType dtype = context->GetInputDesc(0)->GetDataType();
38+ tiling->dtype = (dtype == ge::DT_FLOAT16) ? DIV_DTYPE_FLOAT16 : DIV_DTYPE_FLOAT;
39+ // 4. 设置参与计算的核数
40+ context->SetBlockDim(BLOCK_DIM);
41+ // 5. 本算子不需要额外 workspace
42+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
43+ currentWorkspace[0] = 0;
44+ return ge::GRAPH_SUCCESS;
45+}
46+}
47+ 
48+ 
49+namespace ge {
50+static ge::graphStatus InferShape(gert::InferShapeContext* context)
51+{
52+ const gert::Shape* x1_shape = context->GetInputShape(0);
53+ gert::Shape* y_shape = context->GetOutputShape(0);
54+ *y_shape = *x1_shape;
55+ return GRAPH_SUCCESS;
56+}
57+static ge::graphStatus InferDataType(gert::InferDataTypeContext *context)
58+{
59+ const auto inputDataType = context->GetInputDataType(0);
60+ context->SetOutputDataType(0, inputDataType);
61+ return ge::GRAPH_SUCCESS;
62+}
63+}
64+ 
65+ 
66+namespace ops {
67+class DivCustomTemplate : public OpDef {
68+public:
69+ explicit DivCustomTemplate(const char* name) : OpDef(name)
70+ {
71+ this->Input("x")
72+ .ParamType(REQUIRED)
73+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
74+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
75+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
76+ this->Input("y")
77+ .ParamType(REQUIRED)
78+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
79+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
80+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
81+ this->Output("z")
82+ .ParamType(REQUIRED)
83+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
84+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
85+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
86+ 
87+ this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);
88+ 
89+ this->AICore()
90+ .SetTiling(optiling::TilingFunc);
91+ this->AICore().AddConfig("ascend910b");
92+ 
93+ }
94+};
95+ 
96+OP_ADD(DivCustomTemplate);
97+}
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/op_kernel/CMakeLists.txt+26-0
@@ -0,0 +1,26 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 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+if ("${CMAKE_BUILD_TYPE}x" STREQUAL "Debugx")
12+ npu_op_kernel_options(ascendc_kernels ALL OPTIONS -g -O0)
13+endif()
14+ 
15+npu_op_kernel_sources(ascendc_kernels
16+ KERNEL_DIR ./
17+)
18+ 
19+npu_op_kernel_library(ascendc_kernels
20+ SRC_BASE ${CMAKE_CURRENT_SOURCE_DIR}/
21+ TILING_LIBRARY cust_optiling
22+)
23+ 
24+npu_op_package_add(${package_name}
25+ LIBRARY ascendc_kernels
26+)
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/op_kernel/div_custom_template.cpp+117-0
@@ -0,0 +1,117 @@
1+/* div_custom_template.cpp (op_kernel)
2+ *
3+ * DivCustomTemplate 算子核函数实现：逐元素除法 z = x / y。
4+ *
5+ * 实现思路（Ascend C 矢量编程范式）：
6+ * 1. 将输入数据按 AI Core 数均分，每个核处理 blockLength 个元素；
7+ * 2. 核内再切分为 tileNum * BUFFER_NUM 个分片，循环执行
8+ * CopyIn(GM->UB) -> Compute(Div) -> CopyOut(UB->GM)；
9+ * 3. 队列深度 BUFFER_NUM = 2（双缓冲），相邻分片的搬运与计算流水重叠，
10+ * 隐藏数据搬运时延；
11+ * 4. 核函数实现为模板类，同时支持 float16 / float32，由 host 侧 tiling
12+ * 记录的 dtype 字段在入口处分发实例。
13+ */
14+ 
15+#include "kernel_operator.h"
16+#include "div_custom_template_tiling.h"
17+ 
18+constexpr uint32_t BUFFER_NUM = 2; // tensor num for each queue（队列缓冲级数：双缓冲）
19+ 
20+template <typename T>
21+class KernelDiv {
22+public:
23+ __aicore__ inline KernelDiv() {}
24+ 
25+ // 初始化：划分本核负责的数据段，绑定 GM，并在 UB 上申请队列缓冲
26+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
27+ {
28+ // 1. 总数据均分到每个 AI Core：本核负责的元素数 = 总长度 / 核数
29+ // （本期任务 shape 为 (8, 2048)，总长度 16384，可被 8 核整除）
30+ this->blockLength = totalLength / AscendC::GetBlockNum();
31+ this->tileNum = tileNum;
32+ // 2. 核内再按 tile 切分，每个 tile 均分给 BUFFER_NUM 级流水（双缓冲）
33+ this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
34+ 
35+ // 3. 绑定 GlobalTensor：按当前核号 GetBlockIdx() 偏移到本核负责的数据段
36+ this->xGm.SetGlobalBuffer((__gm__ T *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
37+ this->yGm.SetGlobalBuffer((__gm__ T *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
38+ this->zGm.SetGlobalBuffer((__gm__ T *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
39+ 
40+ // 4. 通过 TPipe 在 UB 上为队列分配缓冲，队列深度 = BUFFER_NUM（双缓冲）
41+ this->pipe.InitBuffer(this->inQueueX, BUFFER_NUM, this->tileLength * sizeof(T));
42+ this->pipe.InitBuffer(this->inQueueY, BUFFER_NUM, this->tileLength * sizeof(T));
43+ this->pipe.InitBuffer(this->outQueueZ, BUFFER_NUM, this->tileLength * sizeof(T));
44+ }
45+ 
46+ // 主流程：共 tileNum * BUFFER_NUM 个分片，逐片执行 CopyIn -> Compute -> CopyOut
47+ __aicore__ inline void Process()
48+ {
49+ const int32_t loopCount = this->tileNum * BUFFER_NUM;
50+ for (int32_t i = 0; i < loopCount; i++) {
51+ CopyIn(i);
52+ Compute(i);
53+ CopyOut(i);
54+ }
55+ }
56+ 
57+private:
58+ // 将本分片的 x/y 从 GM 搬入 UB，入队通知 Compute 输入已就绪
59+ __aicore__ inline void CopyIn(int32_t progress)
60+ {
61+ AscendC::LocalTensor<T> xLocal = this->inQueueX.template AllocTensor<T>();
62+ AscendC::LocalTensor<T> yLocal = this->inQueueY.template AllocTensor<T>();
63+ AscendC::DataCopy(xLocal, this->xGm[progress * this->tileLength], this->tileLength);
64+ AscendC::DataCopy(yLocal, this->yGm[progress * this->tileLength], this->tileLength);
65+ this->inQueueX.EnQue(xLocal);
66+ this->inQueueY.EnQue(yLocal);
67+ }
68+ 
69+ // 出队输入张量，调用 AscendC::Div 完成逐元素除法，结果入队等待搬出
70+ __aicore__ inline void Compute(int32_t progress)
71+ {
72+ AscendC::LocalTensor<T> xLocal = this->inQueueX.template DeQue<T>();
73+ AscendC::LocalTensor<T> yLocal = this->inQueueY.template DeQue<T>();
74+ AscendC::LocalTensor<T> zLocal = this->outQueueZ.template AllocTensor<T>();
75+ AscendC::Div(zLocal, xLocal, yLocal, this->tileLength);
76+ this->outQueueZ.template EnQue<T>(zLocal);
77+ this->inQueueX.FreeTensor(xLocal);
78+ this->inQueueY.FreeTensor(yLocal);
79+ }
80+ 
81+ // 出队计算结果，从 UB 搬回 GM 对应分片，并释放缓冲
82+ __aicore__ inline void CopyOut(int32_t progress)
83+ {
84+ AscendC::LocalTensor<T> zLocal = this->outQueueZ.template DeQue<T>();
85+ AscendC::DataCopy(this->zGm[progress * this->tileLength], zLocal, this->tileLength);
86+ this->outQueueZ.FreeTensor(zLocal);
87+ }
88+ 
89+private:
90+ AscendC::TPipe pipe;
91+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX;
92+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueY;
93+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
94+ AscendC::GlobalTensor<T> xGm;
95+ AscendC::GlobalTensor<T> yGm;
96+ AscendC::GlobalTensor<T> zGm;
97+ uint32_t blockLength; // 本核负责的元素数
98+ uint32_t tileNum; // 核内 tile 切分数
99+ uint32_t tileLength; // 每个分片的元素数
100+};
101+ 
102+extern "C" __global__ __aicore__ void div_custom_template(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace,
103+ GM_ADDR tiling)
104+{
105+ REGISTER_TILING_DEFAULT(DivCustomTemplateTilingData);
106+ GET_TILING_DATA(tilingData, tiling);
107+ // 根据 host 侧 tiling 记录的数据类型，分发对应的核函数模板实例
108+ if (tilingData.dtype == DIV_DTYPE_FLOAT16) {
109+ KernelDiv<half> op;
110+ op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);
111+ op.Process();
112+ } else {
113+ KernelDiv<float> op;
114+ op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);
115+ op.Process();
116+ }
117+}
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/custom_op/op_kernel/div_custom_template_tiling.h+24-0
@@ -0,0 +1,24 @@
1+/* div_custom_template_tiling.h
2+ *
3+ * DivCustomTemplate 算子的 tiling 数据结构定义（host 侧与 kernel 侧共享）。
4+ * host 侧 TilingFunc 在运行时填写本结构，随启动参数下发；
5+ * kernel 侧通过 GET_TILING_DATA 解析出同一份数据。
6+ */
7+ 
8+#ifndef DIV_CUSTOM_TEMPLATE_TILING_H
9+#define DIV_CUSTOM_TEMPLATE_TILING_H
10+#include <cstdint>
11+ 
12+// 输入数据类型编码：host 侧 tiling 写入，kernel 侧据此分发对应的核函数模板实例
13+constexpr uint32_t DIV_DTYPE_FLOAT16 = 0;
14+constexpr uint32_t DIV_DTYPE_FLOAT = 1;
15+// 每个 AI Core 内的 tile 切分数（配合队列双缓冲做流水）
16+constexpr uint32_t DIV_TILE_NUM = 8;
17+ 
18+struct DivCustomTemplateTilingData {
19+ uint32_t totalLength; // 输入张量的总元素数
20+ uint32_t tileNum; // 每个核内的 tile 切分数
21+ uint32_t dtype; // 输入数据类型编码（DIV_DTYPE_FLOAT16 / DIV_DTYPE_FLOAT）
22+};
23+ 
24+#endif // DIV_CUSTOM_TEMPLATE_TILING_H
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/div_custom_template.json+22-0
@@ -0,0 +1,22 @@
1+[{
2+ "op": "DivCustomTemplate",
3+ "input_desc": [{
4+ "name": "x",
5+ "param_type": "required",
6+ "format": ["ND", "ND"],
7+ "type": ["float16", "float"]
8+ },
9+ {
10+ "name": "y",
11+ "param_type": "required",
12+ "format": ["ND", "ND"],
13+ "type": ["float16", "float"]
14+ }
15+ ],
16+ "output_desc": [{
17+ "name": "z",
18+ "param_type": "required",
19+ "format": ["ND", "ND"],
20+ "type": ["float16", "float"]
21+ }]
22+}]
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/run.sh+90-0
@@ -0,0 +1,90 @@
1+#!/bin/bash
2+set -e
3+ 
4+# 获取当前脚本所在目录，确保在任何路径下执行都能找到正确文件
5+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )"
6+cd "$SCRIPT_DIR"
7+ 
8+echo "=========================================="
9+echo " 1. Loading CANN Environment"
10+echo "=========================================="
11+# 加载 CANN 环境变量
12+if [ -n "$ASCEND_TOOLKIT_HOME" ] && [ -f "$ASCEND_TOOLKIT_HOME/set_env.sh" ]; then
13+ source $ASCEND_TOOLKIT_HOME/set_env.sh
14+elif [ -n "$ASCEND_HOME_PATH" ] && [ -f "$ASCEND_HOME_PATH/set_env.sh" ]; then
15+ source $ASCEND_HOME_PATH/set_env.sh
16+else
17+ echo "Error: Cannot find set_env.sh. Please check your CANN installation."
18+ exit 1
19+fi
20+echo "CANN Environment loaded successfully."
21+ 
22+echo "=========================================="
23+echo " 2. Generating Operator Project (msopgen)"
24+echo "=========================================="
25+# 若算子工程不存在，则基于原型文件使用 msopgen 生成（工程化内容由参与者自行完成）
26+if [ -d "custom_op" ] && [ -f "custom_op/build.sh" ]; then
27+ echo ">>> custom_op already exists, skip msopgen generation."
28+else
29+ echo ">>> Generating operator project from div_custom_template.json..."
30+ msopgen gen -i div_custom_template.json -c ai_core-ascend910b1 -lan cpp -out ./custom_op
31+ echo ">>> Operator project generated at ./custom_op"
32+ echo ">>> 已基于910b1生成自定义算子工程 ./custom_op，如果芯片类型不一致或者有其他修改请手动生成相应工程"
33+ echo ">>> 请先在 custom_op/op_kernel 中完成算子核函数实现（参考 README.md），再重新执行 bash run.sh"
34+ exit 0
35+fi
36+ 
37+echo "=========================================="
38+echo " 3. Building Custom Operator"
39+echo "=========================================="
40+cd custom_op
41+# 清理旧的构建目录以确保干净编译
42+rm -rf build_out
43+echo ">>> Running build.sh..."
44+bash build.sh
45+ 
46+echo "=========================================="
47+echo " 4. Installing Custom Operator"
48+echo "=========================================="
49+RUN_FILE=$(ls build_out/custom_opp*.run 2>/dev/null | head -n 1)
50+ 
51+if [ -z "$RUN_FILE" ]; then
52+ echo "Error: .run file not found in build_out. Build might have failed."
53+ exit 1
54+fi
55+echo ">>> Found installer: $RUN_FILE"
56+# 安装算子到用户目录
57+$RUN_FILE --install-path=${HOME}/
58+echo "Operator installed successfully."
59+ 
60+echo "=========================================="
61+echo " 5. Loading Custom Operator Environment"
62+echo "=========================================="
63+# 【关键】必须 source 自定义算子的环境变量，否则运行时找不到算子库
64+if [ -f "${HOME}/vendors/customize/bin/set_env.bash" ]; then
65+ source ${HOME}/vendors/customize/bin/set_env.bash
66+ echo "Custom operator environment loaded."
67+else
68+ echo "Warning: Custom operator env script not found at ${HOME}/vendors/customize/bin/set_env.bash"
69+fi
70+ 
71+echo "=========================================="
72+echo " 6. Building Test Case"
73+echo "=========================================="
74+cd "$SCRIPT_DIR"
75+# 编译测试代码
76+echo ">>> Compiling test/main.cpp..."
77+g++ -I$ASCEND_TOOLKIT_HOME/include \
78+ -I${HOME}/vendors/customize/op_api/include \
79+ -L$ASCEND_TOOLKIT_HOME/lib64 \
80+ -L${HOME}/vendors/customize/op_api/lib \
81+ test/main.cpp \
82+ -lcust_opapi -lnnopbase -lacl_rt \
83+ -o execute_div_op
84+echo "Test case built successfully."
85+ 
86+echo "=========================================="
87+echo " 7. Running Test Case"
88+echo "=========================================="
89+echo ">>> Executing..."
90+./execute_div_op
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/test/CMakeLists.txt+60-0
@@ -0,0 +1,60 @@
1+# Copyright (c) Huawei Technologies Co., Ltd. 2020. All rights reserved.
2+ 
3+# CMake lowest version requirement
4+cmake_minimum_required(VERSION 3.5.1)
5+ 
6+# project information
7+project(acl_execute_div)
8+ 
9+# Compile options
10+add_compile_options(-std=c++11)
11+ 
12+set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "./")
13+ 
14+set(INC_PATH $ENV{DDK_PATH})
15+ 
16+if (NOT DEFINED ENV{DDK_PATH})
17+ set(INC_PATH "/usr/local/Ascend/ascend-toolkit/latest")
18+ message(STATUS "set default INC_PATH: ${INC_PATH}")
19+else ()
20+ message(STATUS "env INC_PATH: ${INC_PATH}")
21+endif()
22+ 
23+set(CUST_PKG_PATH "${INC_PATH}/opp/vendors/customize/op_api")
24+ 
25+set(LIB_PATH $ENV{NPU_HOST_LIB})
26+ 
27+# Dynamic libraries in the stub directory can only be used for compilation
28+if (NOT DEFINED ENV{NPU_HOST_LIB})
29+ string(TOLOWER "${CMAKE_SYSTEM_NAME}" SYSTEM_NAME_LOWER)
30+ set(LIB_PATH "/usr/local/Ascend/ascend-toolkit/latest/${CMAKE_SYSTEM_PROCESSOR}-${SYSTEM_NAME_LOWER}/devlib")
31+ message(STATUS "set default LIB_PATH: ${LIB_PATH}")
32+else ()
33+ message(STATUS "env LIB_PATH: ${LIB_PATH}")
34+endif()
35+ 
36+# Header path
37+include_directories(
38+ ${INC_PATH}/include
39+ ${CUST_PKG_PATH}/include
40+)
41+ 
42+# add host lib path
43+link_directories(
44+ ${LIB_PATH}
45+ ${CUST_PKG_PATH}/lib
46+)
47+ 
48+add_executable(execute_div_op
49+ main.cpp
50+)
51+ 
52+target_link_libraries(execute_div_op
53+ ascendcl
54+ cust_opapi
55+ acl_op_compiler
56+ nnopbase
57+ stdc++
58+)
59+ 
60+install(TARGETS execute_div_op DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/test/main.cpp+186-0
@@ -0,0 +1,186 @@
1+/**
2+ * @file main.cpp
3+ *
4+ * Copyright (C) 2024. Huawei Technologies Co., Ltd. All rights reserved.
5+ *
6+ * This program is distributed in the hope that it will be useful,
7+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
8+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
9+ */
10+#include <algorithm>
11+#include <cstdint>
12+#include <iostream>
13+#include <vector>
14+ 
15+#include "acl/acl.h"
16+#include "aclnn_div_custom_template.h"
17+ 
18+#define SUCCESS 0
19+#define FAILED 1
20+ 
21+#define CHECK_RET(cond, return_expr) \
22+ do { \
23+ if (!(cond)) { \
24+ return_expr; \
25+ } \
26+ } while (0)
27+ 
28+#define LOG_PRINT(message, ...) \
29+ do { \
30+ printf(message, ##__VA_ARGS__); \
31+ } while (0)
32+ 
33+int64_t GetShapeSize(const std::vector<int64_t> &shape)
34+{
35+ int64_t shapeSize = 1;
36+ for (auto i : shape) {
37+ shapeSize *= i;
38+ }
39+ return shapeSize;
40+}
41+ 
42+int Init(int32_t deviceId, aclrtStream *stream)
43+{
44+ // Fixed code, acl initialization
45+ auto ret = aclInit(nullptr);
46+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return FAILED);
47+ ret = aclrtSetDevice(deviceId);
48+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return FAILED);
49+ ret = aclrtCreateStream(stream);
50+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return FAILED);
51+ 
52+ return SUCCESS;
53+}
54+ 
55+template <typename T>
56+int CreateAclTensor(const std::vector<T> &hostData, const std::vector<int64_t> &shape, void **deviceAddr,
57+ aclDataType dataType, aclTensor **tensor)
58+{
59+ auto size = GetShapeSize(shape) * sizeof(T);
60+ // Call aclrtMalloc to allocate device memory
61+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
62+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return FAILED);
63+ 
64+ // Call aclrtMemcpy to copy host data to device memory
65+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
66+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return FAILED);
67+ 
68+ // Call aclCreateTensor to create a aclTensor object
69+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, nullptr, 0, aclFormat::ACL_FORMAT_ND, shape.data(),
70+ shape.size(), *deviceAddr);
atomgit-bot
atomgit-botatomgit-bot5 天前

🟡 Medium Priority

变更行：test/main.cpp 第 69–70 行，*tensor = aclCreateTensor(...) 后无条件 return SUCCESS。

受影响行为/契约：aclCreateTensor 是 C 接口，失败时返回 nullptr。该函数把 aclCreateTensor 的返回直接赋给输出参数却从不校验，而调用方 main 只检查 CreateAclTensor 的返回值（恒为 SUCCESS），nullptr 张量被当作合法输入继续向下传递。

失败模式：当 aclCreateTensor 因参数或内存问题失败时，inputX/inputY/outputZ 为 nullptr，随后 aclnnDivCustomTemplateGetWorkspaceSize(nullptr, ...) 会返回 ACL_ERROR_INVALID_PARAM 或直接触发异常/core dump。前者使测试打印误导性的"aclnnDivCustomTemplateGetWorkspaceSize failed"而掩盖真实根因，后者直接崩溃——错误被静默吞掉。规范明确要求"示例代码/单元测试/C 接口封装必须校验返回值"。

建议：在 aclCreateTensor 调用后立即校验返回值：CHECK_RET(*tensor != nullptr, LOG_PRINT("aclCreateTensor failed.\n"); return FAILED);，避免 nullptr 张量流入 aclnn 接口导致误导性报错或崩溃。

改动建议
70
+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, nullptr, 0, aclFormat::ACL_FORMAT_ND, shape.data(),
70
- shape.size(), *deviceAddr);
71
+ shape.size(), *deviceAddr);
72
+ CHECK_RET(*tensor != nullptr, LOG_PRINT("aclCreateTensor failed.\n"); return FAILED);
73
+ return SUCCESS;
应用建议
likedislike
71+ return SUCCESS;
72+}
73+ 
74+void DestroyResources(std::vector<void *> tensors, std::vector<void *> deviceAddrs, aclrtStream stream,
75+ int32_t deviceId, void *workspaceAddr = nullptr)
76+{
77+ // Release aclTensor and device
78+ for (uint32_t i = 0; i < tensors.size(); i++) {
79+ if (tensors[i] != nullptr) {
80+ aclDestroyTensor(reinterpret_cast<aclTensor *>(tensors[i]));
81+ }
82+ if (deviceAddrs[i] != nullptr) {
83+ aclrtFree(deviceAddrs[i]);
84+ }
85+ }
86+ if (workspaceAddr != nullptr) {
87+ aclrtFree(workspaceAddr);
88+ }
89+ // Destroy stream and reset device
90+ aclrtDestroyStream(stream);
91+ aclrtResetDevice(deviceId);
92+ aclFinalize();
93+}
94+ 
95+int main(int argc, char **argv)
96+{
97+ // 1. (Fixed code) Initialize device / stream, refer to the list of external interfaces of acl
98+ // Update deviceId to your own device id
99+ int32_t deviceId = 0;
100+ aclrtStream stream;
101+ auto ret = Init(deviceId, &stream);
102+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return FAILED);
103+ 
104+ // 2. Create input and output, need to customize according to the interface of the API
105+ std::vector<int64_t> inputXShape = {8, 2048};
106+ std::vector<int64_t> inputYShape = {8, 2048};
107+ std::vector<int64_t> outputZShape = {8, 2048};
108+ void *inputXDeviceAddr = nullptr;
109+ void *inputYDeviceAddr = nullptr;
110+ void *outputZDeviceAddr = nullptr;
111+ aclTensor *inputX = nullptr;
112+ aclTensor *inputY = nullptr;
113+ aclTensor *outputZ = nullptr;
114+ std::vector<aclFloat16> inputXHostData(inputXShape[0] * inputXShape[1]);
115+ std::vector<aclFloat16> inputYHostData(inputYShape[0] * inputYShape[1]);
116+ std::vector<aclFloat16> outputZHostData(outputZShape[0] * outputZShape[1]);
117+ for (int i = 0; i < inputXShape[0] * inputXShape[1]; ++i) {
118+ inputXHostData[i] = aclFloatToFloat16(1.0);
119+ inputYHostData[i] = aclFloatToFloat16(2.0);
120+ outputZHostData[i] = aclFloatToFloat16(0.0);
121+ }
122+ std::vector<void *> tensors = {inputX, inputY, outputZ};
123+ std::vector<void *> deviceAddrs = {inputXDeviceAddr, inputYDeviceAddr, outputZDeviceAddr};
124+ // Create inputX aclTensor
125+ ret = CreateAclTensor(inputXHostData, inputXShape, &inputXDeviceAddr, aclDataType::ACL_FLOAT16, &inputX);
126+ CHECK_RET(ret == ACL_SUCCESS, DestroyResources(tensors, deviceAddrs, stream, deviceId); return FAILED);
127+ // Create inputY aclTensor
128+ ret = CreateAclTensor(inputYHostData, inputYShape, &inputYDeviceAddr, aclDataType::ACL_FLOAT16, &inputY);
129+ CHECK_RET(ret == ACL_SUCCESS, DestroyResources(tensors, deviceAddrs, stream, deviceId); return FAILED);
130+ // Create outputZ aclTensor
131+ ret = CreateAclTensor(outputZHostData, outputZShape, &outputZDeviceAddr, aclDataType::ACL_FLOAT16, &outputZ);
132+ CHECK_RET(ret == ACL_SUCCESS, DestroyResources(tensors, deviceAddrs, stream, deviceId); return FAILED);
133+ 
134+ // 3. Call the API of the custom operator library
135+ uint64_t workspaceSize = 0;
136+ aclOpExecutor *executor;
137+ // Calculate the workspace size and allocate memory for it
138+ ret = aclnnDivCustomTemplateGetWorkspaceSize(inputX, inputY, outputZ, &workspaceSize, &executor);
139+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDivCustomTemplateGetWorkspaceSize failed. ERROR: %d\n", ret);
140+ DestroyResources(tensors, deviceAddrs, stream, deviceId); return FAILED);
141+ 
142+ void *workspaceAddr = nullptr;
143+ if (workspaceSize > 0) {
144+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
145+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret);
146+ DestroyResources(tensors, deviceAddrs, stream, deviceId, workspaceAddr); return FAILED);
147+ }
148+ // Execute the custom operator
149+ ret = aclnnDivCustomTemplate(workspaceAddr, workspaceSize, executor, stream);
150+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDivCustomTemplate failed. ERROR: %d\n", ret);
151+ DestroyResources(tensors, deviceAddrs, stream, deviceId, workspaceAddr); return FAILED);
152+ 
153+ // 4. (Fixed code) Synchronize and wait for the task to complete
154+ ret = aclrtSynchronizeStream(stream);
155+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret);
156+ DestroyResources(tensors, deviceAddrs, stream, deviceId, workspaceAddr); return FAILED);
157+ 
158+ // 5. Get the output value, copy the result from device memory to host memory, need to modify according to the
159+ // interface of the API
160+ auto size = GetShapeSize(outputZShape);
161+ std::vector<aclFloat16> resultData(size, 0);
162+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outputZDeviceAddr,
163+ size * sizeof(aclFloat16), ACL_MEMCPY_DEVICE_TO_HOST);
164+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret);
165+ DestroyResources(tensors, deviceAddrs, stream, deviceId, workspaceAddr); return FAILED);
166+ 
167+ // 6. Destroy resources, need to modify according to the interface of the API
168+ DestroyResources(tensors, deviceAddrs, stream, deviceId, workspaceAddr);
169+ 
170+ // print the output result
171+ // 期望结果：z = x / y = 1.0 / 2.0 = 0.5
172+ std::vector<aclFloat16> goldenData(size, aclFloatToFloat16(0.5));
173+ 
174+ LOG_PRINT("result is:\n");
175+ for (int64_t i = 0; i < 10; i++) {
176+ LOG_PRINT("%.1f ", aclFloat16ToFloat(resultData[i]));
177+ }
178+ LOG_PRINT("\n");
179+ if (std::equal(resultData.begin(), resultData.end(), goldenData.begin())) {
180+ LOG_PRINT("test pass\n");
181+ } else {
182+ LOG_PRINT("test failed\n");
183+ return FAILED;
184+ }
185+ return SUCCESS;
186+}
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/weixin_53562472/DivCustomTemplate/verification.png+0-0