Star136
1532
代码
代码
Issues45
Pull Requests670
Actions
项目讨论
Wiki
分析
Star136
1532
已开启
【代码侦探Challenge 04】完成 DivCustomTemplate 工程化算子实现 #2156
iss_liqinghua创建于 6 天前
【代码侦探Challenge 04】完成 DivCustomTemplate 工程化算子实现 #2156
已开启
iss_liqinghua创建于 6 天前
从iss_liqinghua/cann-outreachchallenge04-divcustomtemplate合入到cann/cann-outreachmaster
共 18 个文件变更+965-0
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/.gitignore+5-0
@@ -0,0 +1,5 @@
1+custom_op/build_out/
2+execute_div_op
3+*.run
4+__pycache__/
5+*.pyc
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/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/iss_liqinghua/DivCustomTemplate/custom_op/CMakePresets.json+63-0
@@ -0,0 +1,63 @@
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_PYTHON_EXECUTABLE": {
41+ "type": "STRING",
42+ "value": "python3"
43+ },
44+ "CMAKE_INSTALL_PREFIX": {
45+ "type": "PATH",
46+ "value": "${sourceDir}/build_out"
47+ },
48+ "ENABLE_CROSS_COMPILE": {
49+ "type": "BOOL",
50+ "value": "False"
51+ },
52+ "CMAKE_CROSS_PLATFORM_COMPILER": {
53+ "type": "PATH",
54+ "value": "/usr/bin/aarch64-linux-gnu-g++"
55+ },
56+ "ASCEND_PACK_SHARED_LIBRARY": {
57+ "type": "BOOL",
58+ "value": "False"
59+ }
60+ }
61+ }
62+ ]
63+}
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/custom_op/build.sh+47-0
@@ -0,0 +1,47 @@
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+set -e
13+ 
14+if [ -z "$BASE_LIBS_PATH" ]; then
15+ if [ -z "$ASCEND_HOME_PATH" ]; then
16+ if [ -z "$ASCEND_AICPU_PATH" ]; then
17+ echo "please set env."
18+ exit 1
19+ else
20+ export ASCEND_HOME_PATH=$ASCEND_AICPU_PATH
21+ fi
22+ else
23+ export ASCEND_HOME_PATH=$ASCEND_HOME_PATH
24+ fi
25+else
26+ export ASCEND_HOME_PATH=$BASE_LIBS_PATH
27+fi
28+echo "using ASCEND_HOME_PATH: $ASCEND_HOME_PATH"
29+script_path=$(realpath "$(dirname "$0")")
30+ 
31+BUILD_DIR="build_out"
32+mkdir -p "$BUILD_DIR"
33+rm -rf "$BUILD_DIR"/*
34+opts=$(python3 "$ASCEND_HOME_PATH/tools/tikcpp/ascendc_kernel_cmake/fwk_modules/util/preset_parse.py" \
35+ "$script_path/CMakePresets.json")
36+cmake_version=$(cmake --version | grep "cmake version" | awk '{print $3}')
37+cmake_extra_args=(-DASCEND_CANN_PACKAGE_PATH="$ASCEND_HOME_PATH")
38+if [ -n "${ASCEND_COMPUTE_UNIT_OVERRIDE:-}" ]; then
39+ cmake_extra_args+=(-DASCEND_COMPUTE_UNIT="$ASCEND_COMPUTE_UNIT_OVERRIDE")
40+fi
41+ 
42+if [ "$(printf '%s\n' "$cmake_version" "3.19.0" | sort -V | head -n 1)" != "3.19.0" ]; then
43+ cmake -S . -B "$BUILD_DIR" $opts "${cmake_extra_args[@]}"
44+else
45+ cmake -S . -B "$BUILD_DIR" --preset=default "${cmake_extra_args[@]}"
46+fi
47+cmake --build "$BUILD_DIR" --target binary package -j"$(nproc)"
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/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/iss_liqinghua/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/iss_liqinghua/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/iss_liqinghua/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/iss_liqinghua/DivCustomTemplate/custom_op/op_host/div_custom_template.cpp+141-0
@@ -0,0 +1,141 @@
1+#include "../op_kernel/div_custom_template_tiling.h"
2+#include "register/op_def_registry.h"
3+ 
4+#include <algorithm>
5+#include <cstdint>
6+#include <limits>
7+ 
8+namespace optiling {
9+namespace {
10+constexpr uint32_t MAX_BLOCK_DIM = 8;
11+constexpr uint32_t MAX_TILE_LENGTH = 1024;
12+constexpr uint32_t DTYPE_FLOAT16 = 0;
13+constexpr uint32_t DTYPE_FLOAT32 = 1;
14+constexpr uint32_t DATA_COPY_ALIGNMENT_BYTES = 32;
15+ 
16+uint32_t SelectBlockDim(uint32_t totalLength, uint32_t alignmentElements)
17+{
18+ const uint32_t alignedUnits = totalLength / alignmentElements;
19+ uint32_t blockDim = std::min(MAX_BLOCK_DIM, alignedUnits);
20+ while (blockDim > 1 && alignedUnits % blockDim != 0) {
21+ --blockDim;
22+ }
23+ return blockDim;
24+}
25+} // namespace
26+ 
27+static ge::graphStatus TilingFunc(gert::TilingContext *context)
28+{
29+ const auto *xShape = context->GetInputShape(0);
30+ const auto *yShape = context->GetInputShape(1);
31+ if (xShape == nullptr || yShape == nullptr) {
32+ return ge::GRAPH_FAILED;
33+ }
34+ 
35+ const auto &xStorageShape = xShape->GetStorageShape();
36+ const auto &yStorageShape = yShape->GetStorageShape();
37+ if (xStorageShape.GetDimNum() != yStorageShape.GetDimNum()) {
38+ return ge::GRAPH_FAILED;
39+ }
40+ for (int32_t index = 0; index < xStorageShape.GetDimNum(); ++index) {
41+ if (xStorageShape.GetDim(index) != yStorageShape.GetDim(index)) {
42+ return ge::GRAPH_FAILED;
43+ }
44+ }
45+ 
46+ const int64_t xLength = xStorageShape.GetShapeSize();
47+ const int64_t yLength = yStorageShape.GetShapeSize();
48+ if (xLength <= 0 || xLength != yLength ||
49+ static_cast<uint64_t>(xLength) > std::numeric_limits<uint32_t>::max()) {
50+ return ge::GRAPH_FAILED;
51+ }
52+ 
53+ const auto *xDesc = context->GetInputDesc(0);
54+ const auto *yDesc = context->GetInputDesc(1);
55+ if (xDesc == nullptr || yDesc == nullptr) {
56+ return ge::GRAPH_FAILED;
57+ }
58+ const auto xDataType = xDesc->GetDataType();
59+ const auto yDataType = yDesc->GetDataType();
60+ if (xDataType != yDataType || (xDataType != ge::DT_FLOAT16 && xDataType != ge::DT_FLOAT)) {
61+ return ge::GRAPH_FAILED;
62+ }
63+ 
64+ const uint32_t totalLength = static_cast<uint32_t>(xLength);
65+ const uint32_t elementBytes = xDataType == ge::DT_FLOAT ? sizeof(float) : sizeof(uint16_t);
66+ const uint32_t alignmentElements = DATA_COPY_ALIGNMENT_BYTES / elementBytes;
67+ if (totalLength % alignmentElements != 0) {
68+ return ge::GRAPH_FAILED;
69+ }
70+ 
71+ const uint32_t blockDim = SelectBlockDim(totalLength, alignmentElements);
72+ const uint32_t blockLength = totalLength / blockDim;
73+ uint32_t tileLength = std::min(MAX_TILE_LENGTH, blockLength);
74+ tileLength -= tileLength % alignmentElements;
75+ if (tileLength == 0) {
76+ tileLength = alignmentElements;
77+ }
78+ 
79+ DivCustomTemplateTilingData tilingData;
80+ tilingData.set_totalLength(totalLength);
81+ tilingData.set_blockLength(blockLength);
82+ tilingData.set_tileLength(tileLength);
83+ tilingData.set_dataType(xDataType == ge::DT_FLOAT ? DTYPE_FLOAT32 : DTYPE_FLOAT16);
84+ 
85+ context->SetBlockDim(blockDim);
86+ tilingData.SaveToBuffer(context->GetRawTilingData()->GetData(),
87+ context->GetRawTilingData()->GetCapacity());
88+ context->GetRawTilingData()->SetDataSize(tilingData.GetDataSize());
89+ 
90+ size_t *workspaceSizes = context->GetWorkspaceSizes(1);
91+ workspaceSizes[0] = 0;
92+ return ge::GRAPH_SUCCESS;
93+}
94+} // namespace optiling
95+ 
96+namespace ge {
97+static graphStatus InferShape(gert::InferShapeContext *context)
98+{
99+ const gert::Shape *inputShape = context->GetInputShape(0);
100+ gert::Shape *outputShape = context->GetOutputShape(0);
101+ *outputShape = *inputShape;
102+ return GRAPH_SUCCESS;
103+}
104+ 
105+static graphStatus InferDataType(gert::InferDataTypeContext *context)
106+{
107+ context->SetOutputDataType(0, context->GetInputDataType(0));
108+ return GRAPH_SUCCESS;
109+}
110+} // namespace ge
111+ 
112+namespace ops {
113+class DivCustomTemplate : public OpDef {
114+public:
115+ explicit DivCustomTemplate(const char *name) : OpDef(name)
116+ {
117+ this->Input("x")
118+ .ParamType(REQUIRED)
119+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
120+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
121+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
122+ this->Input("y")
123+ .ParamType(REQUIRED)
124+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
125+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
126+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
127+ this->Output("z")
128+ .ParamType(REQUIRED)
129+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
130+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
131+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
132+ 
133+ this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);
134+ this->AICore().SetTiling(optiling::TilingFunc);
135+ this->AICore().AddConfig("ascend910b");
136+ this->AICore().AddConfig("ascend310p");
137+ }
138+};
139+ 
140+OP_ADD(DivCustomTemplate);
141+} // namespace ops
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/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/iss_liqinghua/DivCustomTemplate/custom_op/op_kernel/div_custom_template.cpp+99-0
@@ -0,0 +1,99 @@
1+#include "div_custom_template_tiling.h"
2+#include "kernel_operator.h"
3+ 
4+using namespace AscendC;
5+ 
6+namespace {
7+constexpr int32_t BUFFER_NUM = 2;
8+constexpr uint32_t DTYPE_FLOAT32 = 1;
9+}
10+ 
11+template <typename T>
12+class KernelDivCustomTemplate {
13+public:
14+ __aicore__ inline KernelDivCustomTemplate() {}
15+ 
16+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t blockLength, uint32_t tileLength)
17+ {
18+ this->blockLength = blockLength;
19+ this->tileLength = tileLength;
20+ 
21+ const uint32_t blockOffset = GetBlockIdx() * blockLength;
22+ xGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T *>(x) + blockOffset, blockLength);
23+ yGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T *>(y) + blockOffset, blockLength);
24+ zGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T *>(z) + blockOffset, blockLength);
25+ 
26+ pipe.InitBuffer(xQueue, BUFFER_NUM, tileLength * sizeof(T));
27+ pipe.InitBuffer(yQueue, BUFFER_NUM, tileLength * sizeof(T));
28+ pipe.InitBuffer(zQueue, BUFFER_NUM, tileLength * sizeof(T));
29+ }
30+ 
31+ __aicore__ inline void Process()
32+ {
33+ for (uint64_t offset = 0; offset < blockLength; offset += tileLength) {
34+ const uint32_t currentOffset = static_cast<uint32_t>(offset);
35+ const uint32_t length = blockLength - currentOffset < tileLength ? blockLength - currentOffset
36+ : tileLength;
37+ CopyIn(currentOffset, length);
38+ Compute(length);
39+ CopyOut(currentOffset, length);
40+ }
41+ }
42+ 
43+private:
44+ __aicore__ inline void CopyIn(uint32_t offset, uint32_t length)
45+ {
46+ LocalTensor<T> xLocal = xQueue.AllocTensor<T>();
47+ LocalTensor<T> yLocal = yQueue.AllocTensor<T>();
48+ DataCopy(xLocal, xGlobal[offset], length);
49+ DataCopy(yLocal, yGlobal[offset], length);
50+ xQueue.EnQue(xLocal);
51+ yQueue.EnQue(yLocal);
52+ }
53+ 
54+ __aicore__ inline void Compute(uint32_t length)
55+ {
56+ LocalTensor<T> xLocal = xQueue.DeQue<T>();
57+ LocalTensor<T> yLocal = yQueue.DeQue<T>();
58+ LocalTensor<T> zLocal = zQueue.AllocTensor<T>();
59+ Div(zLocal, xLocal, yLocal, length);
60+ zQueue.EnQue(zLocal);
61+ xQueue.FreeTensor(xLocal);
62+ yQueue.FreeTensor(yLocal);
63+ }
64+ 
65+ __aicore__ inline void CopyOut(uint32_t offset, uint32_t length)
66+ {
67+ LocalTensor<T> zLocal = zQueue.DeQue<T>();
68+ DataCopy(zGlobal[offset], zLocal, length);
69+ zQueue.FreeTensor(zLocal);
70+ }
71+ 
72+private:
73+ TPipe pipe;
74+ TQue<QuePosition::VECIN, BUFFER_NUM> xQueue;
75+ TQue<QuePosition::VECIN, BUFFER_NUM> yQueue;
76+ TQue<QuePosition::VECOUT, BUFFER_NUM> zQueue;
77+ GlobalTensor<T> xGlobal;
78+ GlobalTensor<T> yGlobal;
79+ GlobalTensor<T> zGlobal;
80+ uint32_t blockLength = 0;
81+ uint32_t tileLength = 0;
82+};
83+ 
84+extern "C" __global__ __aicore__ void div_custom_template(GM_ADDR x, GM_ADDR y, GM_ADDR z,
85+ GM_ADDR workspace, GM_ADDR tiling)
86+{
87+ (void)workspace;
88+ GET_TILING_DATA(tilingData, tiling);
89+ 
90+ if (tilingData.dataType == DTYPE_FLOAT32) {
91+ KernelDivCustomTemplate<float> op;
92+ op.Init(x, y, z, tilingData.blockLength, tilingData.tileLength);
93+ op.Process();
94+ } else {
95+ KernelDivCustomTemplate<half> op;
96+ op.Init(x, y, z, tilingData.blockLength, tilingData.tileLength);
97+ op.Process();
98+ }
99+}
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/custom_op/op_kernel/div_custom_template_tiling.h+19-0
@@ -0,0 +1,19 @@
1+#ifndef DIV_CUSTOM_TEMPLATE_TILING_H
2+#define DIV_CUSTOM_TEMPLATE_TILING_H
3+ 
4+#include "register/tilingdata_base.h"
5+ 
6+namespace optiling {
7+ 
8+BEGIN_TILING_DATA_DEF(DivCustomTemplateTilingData)
9+ TILING_DATA_FIELD_DEF(uint32_t, totalLength);
10+ TILING_DATA_FIELD_DEF(uint32_t, blockLength);
11+ TILING_DATA_FIELD_DEF(uint32_t, tileLength);
12+ TILING_DATA_FIELD_DEF(uint32_t, dataType);
13+END_TILING_DATA_DEF;
14+ 
15+REGISTER_TILING_DATA_CLASS(DivCustomTemplate, DivCustomTemplateTilingData)
16+ 
17+} // namespace optiling
18+ 
19+#endif // DIV_CUSTOM_TEMPLATE_TILING_H
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/div_custom_template.json+23-0
@@ -0,0 +1,23 @@
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+}]
23+ 
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/run.sh+87-0
@@ -0,0 +1,87 @@
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+ CANN_INSTALL_PATH="${ASCEND_TOOLKIT_HOME}"
14+elif [ -n "${ASCEND_HOME_PATH:-}" ] && [ -f "${ASCEND_HOME_PATH}/set_env.sh" ]; then
15+ CANN_INSTALL_PATH="${ASCEND_HOME_PATH}"
16+else
17+ echo "Error: Cannot find set_env.sh. Please check your CANN installation."
18+ exit 1
19+fi
20+source "${CANN_INSTALL_PATH}/set_env.sh"
21+export ASCEND_TOOLKIT_HOME="${ASCEND_TOOLKIT_HOME:-${CANN_INSTALL_PATH}}"
22+export ASCEND_HOME_PATH="${ASCEND_HOME_PATH:-${CANN_INSTALL_PATH}}"
23+echo "CANN Environment loaded successfully."
24+ 
25+echo "=========================================="
26+echo " 2. Checking Operator Project"
27+echo "=========================================="
28+if [ ! -d "custom_op" ] || [ ! -f "custom_op/build.sh" ]; then
29+ echo "Error: custom_op project is missing."
30+ exit 1
31+fi
32+echo ">>> custom_op project found."
33+ 
34+echo "=========================================="
35+echo " 3. Building Custom Operator"
36+echo "=========================================="
37+cd custom_op
38+# 清理旧的构建目录以确保干净编译
39+rm -rf build_out
40+echo ">>> Running build.sh..."
41+bash build.sh
42+ 
43+echo "=========================================="
44+echo " 4. Installing Custom Operator"
45+echo "=========================================="
46+RUN_FILE=$(ls build_out/custom_opp*.run 2>/dev/null | head -n 1)
47+ 
48+if [ -z "$RUN_FILE" ]; then
49+ echo "Error: .run file not found in build_out. Build might have failed."
50+ exit 1
51+fi
52+echo ">>> Found installer: $RUN_FILE"
53+# 安装算子到用户目录
54+"$RUN_FILE" --install-path="${HOME}/"
55+echo "Operator installed successfully."
56+ 
57+echo "=========================================="
58+echo " 5. Loading Custom Operator Environment"
59+echo "=========================================="
60+# 【关键】必须 source 自定义算子的环境变量，否则运行时找不到算子库
61+if [ -f "${HOME}/vendors/customize/bin/set_env.bash" ]; then
62+ source ${HOME}/vendors/customize/bin/set_env.bash
63+ echo "Custom operator environment loaded."
64+else
65+ echo "Warning: Custom operator env script not found at ${HOME}/vendors/customize/bin/set_env.bash"
66+fi
67+ 
68+echo "=========================================="
69+echo " 6. Building Test Case"
70+echo "=========================================="
71+cd "$SCRIPT_DIR"
72+# 编译测试代码
73+echo ">>> Compiling test/main.cpp..."
74+g++ -I"${ASCEND_TOOLKIT_HOME}/include" \
75+ -I"${HOME}/vendors/customize/op_api/include" \
76+ -L"${ASCEND_TOOLKIT_HOME}/lib64" \
77+ -L"${HOME}/vendors/customize/op_api/lib" \
78+ test/main.cpp \
79+ -lcust_opapi -lnnopbase -lacl_rt \
80+ -o execute_div_op
81+echo "Test case built successfully."
82+ 
83+echo "=========================================="
84+echo " 7. Running Test Case"
85+echo "=========================================="
86+echo ">>> Executing..."
87+./execute_div_op
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/test/CMakeLists.txt+61-0
@@ -0,0 +1,61 @@
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})
61+ 
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/test/main.cpp+200-0
@@ -0,0 +1,200 @@
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 <cstdio>
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 (int64_t dimension : shape) {
37+ shapeSize *= dimension;
38+ }
39+ return shapeSize;
40+}
41+ 
42+int Init(int32_t deviceId, aclrtStream *stream)
43+{
44+ auto ret = aclInit(nullptr);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return FAILED);
46+ 
47+ ret = aclrtSetDevice(deviceId);
48+ if (ret != ACL_SUCCESS) {
49+ LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret);
50+ aclFinalize();
51+ return FAILED;
52+ }
53+ 
54+ ret = aclrtCreateStream(stream);
55+ if (ret != ACL_SUCCESS) {
56+ LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret);
57+ aclrtResetDevice(deviceId);
58+ aclFinalize();
59+ return FAILED;
60+ }
61+ return SUCCESS;
62+}
63+ 
64+template <typename T>
65+int CreateAclTensor(const std::vector<T> &hostData, const std::vector<int64_t> &shape, void **deviceAddr,
66+ aclDataType dataType, aclTensor **tensor)
67+{
68+ const size_t byteSize = static_cast<size_t>(GetShapeSize(shape)) * sizeof(T);
69+ auto ret = aclrtMalloc(deviceAddr, byteSize, ACL_MEM_MALLOC_HUGE_FIRST);
70+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return FAILED);
71+ 
72+ ret = aclrtMemcpy(*deviceAddr, byteSize, hostData.data(), byteSize, ACL_MEMCPY_HOST_TO_DEVICE);
73+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return FAILED);
74+ 
75+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, nullptr, 0, aclFormat::ACL_FORMAT_ND,
76+ shape.data(), shape.size(), *deviceAddr);
77+ CHECK_RET(*tensor != nullptr, LOG_PRINT("aclCreateTensor failed.\n"); return FAILED);
78+ return SUCCESS;
79+}
80+ 
81+void DestroyCaseResources(aclTensor *inputX, aclTensor *inputY, aclTensor *outputZ, void *inputXDeviceAddr,
82+ void *inputYDeviceAddr, void *outputZDeviceAddr, void *workspaceAddr)
83+{
84+ if (inputX != nullptr) {
85+ aclDestroyTensor(inputX);
86+ }
87+ if (inputY != nullptr) {
88+ aclDestroyTensor(inputY);
89+ }
90+ if (outputZ != nullptr) {
91+ aclDestroyTensor(outputZ);
92+ }
93+ if (inputXDeviceAddr != nullptr) {
94+ aclrtFree(inputXDeviceAddr);
95+ }
96+ if (inputYDeviceAddr != nullptr) {
97+ aclrtFree(inputYDeviceAddr);
98+ }
99+ if (outputZDeviceAddr != nullptr) {
100+ aclrtFree(outputZDeviceAddr);
101+ }
102+ if (workspaceAddr != nullptr) {
103+ aclrtFree(workspaceAddr);
104+ }
105+}
106+ 
107+template <typename T, typename FloatToT, typename TToFloat>
108+int RunDivCase(aclrtStream stream, aclDataType dataType, const char *caseName, FloatToT toHostValue,
109+ TToFloat toFloatValue)
110+{
111+ const std::vector<int64_t> shape = {8, 2048};
112+ const int64_t elementCount = GetShapeSize(shape);
113+ std::vector<T> inputXHostData(elementCount, toHostValue(1.0F));
114+ std::vector<T> inputYHostData(elementCount, toHostValue(2.0F));
115+ std::vector<T> outputZHostData(elementCount, toHostValue(0.0F));
116+ 
117+ void *inputXDeviceAddr = nullptr;
118+ void *inputYDeviceAddr = nullptr;
119+ void *outputZDeviceAddr = nullptr;
120+ void *workspaceAddr = nullptr;
121+ aclTensor *inputX = nullptr;
122+ aclTensor *inputY = nullptr;
123+ aclTensor *outputZ = nullptr;
124+ 
125+ auto cleanup = [&]() {
126+ DestroyCaseResources(inputX, inputY, outputZ, inputXDeviceAddr, inputYDeviceAddr, outputZDeviceAddr,
127+ workspaceAddr);
128+ };
129+ 
130+ auto ret = CreateAclTensor(inputXHostData, shape, &inputXDeviceAddr, dataType, &inputX);
131+ CHECK_RET(ret == SUCCESS, cleanup(); return FAILED);
132+ ret = CreateAclTensor(inputYHostData, shape, &inputYDeviceAddr, dataType, &inputY);
133+ CHECK_RET(ret == SUCCESS, cleanup(); return FAILED);
134+ ret = CreateAclTensor(outputZHostData, shape, &outputZDeviceAddr, dataType, &outputZ);
135+ CHECK_RET(ret == SUCCESS, cleanup(); return FAILED);
136+ 
137+ uint64_t workspaceSize = 0;
138+ aclOpExecutor *executor = nullptr;
139+ ret = aclnnDivCustomTemplateGetWorkspaceSize(inputX, inputY, outputZ, &workspaceSize, &executor);
140+ CHECK_RET(ret == ACL_SUCCESS,
141+ LOG_PRINT("aclnnDivCustomTemplateGetWorkspaceSize failed. ERROR: %d\n", ret); cleanup();
142+ return FAILED);
143+ 
144+ if (workspaceSize > 0) {
145+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
146+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); cleanup();
147+ return FAILED);
148+ }
149+ 
150+ ret = aclnnDivCustomTemplate(workspaceAddr, workspaceSize, executor, stream);
151+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnDivCustomTemplate failed. ERROR: %d\n", ret); cleanup();
152+ return FAILED);
153+ ret = aclrtSynchronizeStream(stream);
154+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); cleanup();
155+ return FAILED);
156+ 
157+ std::vector<T> resultData(elementCount);
158+ const size_t resultBytes = resultData.size() * sizeof(T);
159+ ret = aclrtMemcpy(resultData.data(), resultBytes, outputZDeviceAddr, resultBytes, ACL_MEMCPY_DEVICE_TO_HOST);
160+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); cleanup();
161+ return FAILED);
162+ 
163+ cleanup();
164+ std::vector<T> goldenData(elementCount, toHostValue(0.5F));
165+ 
166+ LOG_PRINT("%s result is:\n", caseName);
167+ for (int64_t index = 0; index < 10; ++index) {
168+ LOG_PRINT("%.1f ", toFloatValue(resultData[index]));
169+ }
170+ LOG_PRINT("\n");
171+ if (!std::equal(resultData.begin(), resultData.end(), goldenData.begin())) {
172+ LOG_PRINT("%s test failed\n", caseName);
173+ return FAILED;
174+ }
175+ LOG_PRINT("%s test pass\n", caseName);
176+ return SUCCESS;
177+}
178+ 
179+int main(int argc, char **argv)
180+{
181+ (void)argc;
182+ (void)argv;
183+ const int32_t deviceId = 0;
184+ aclrtStream stream = nullptr;
185+ auto ret = Init(deviceId, &stream);
186+ CHECK_RET(ret == SUCCESS, return FAILED);
187+ 
188+ ret = RunDivCase<aclFloat16>(stream, aclDataType::ACL_FLOAT16, "float16",
189+ [](float value) { return aclFloatToFloat16(value); },
190+ [](aclFloat16 value) { return aclFloat16ToFloat(value); });
191+ if (ret == SUCCESS) {
192+ ret = RunDivCase<float>(stream, aclDataType::ACL_FLOAT, "float32", [](float value) { return value; },
193+ [](float value) { return value; });
194+ }
195+ 
196+ aclrtDestroyStream(stream);
197+ aclrtResetDevice(deviceId);
198+ aclFinalize();
199+ return ret;
200+}
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/verification.log+38-0
@@ -0,0 +1,38 @@
1+Script started on 2026-08-21 14:34:06+08:00 [TERM="dumb" TTY="/dev/pts/0" COLUMNS="120" LINES="40"]
2++ date '+%F %T %Z'
3+2026-08-21 14:34:06 CST
4++ sed -n 1,15p /usr/local/Ascend/ascend-toolkit/latest/compiler/version.info
5+Version=9.0.0
6+version_dir=cann
7+required_package_runtime_version="9.0"
8+required_package_ge-executor_version="9.0"
9+required_package_metadef_version="9.0"
10+required_package_opbase_version="9.0"
11+required_package_bisheng-compiler_version="9.0"
12+required_package_tbe-tik_version="9.0"
13+timestamp=20260428_134817545
14++ npu-smi info
15++--------------------------------------------------------------------------------------------------------+
16+| npu-smi 25.5.2 Version: 25.5.2 |
17++-------------------------------+-----------------+------------------------------------------------------+
18+| NPU Name | Health | Power(W) Temp(C) Hugepages-Usage(page) |
19+| Chip Device | Bus-Id | AICore(%) Memory-Usage(MB) |
20++===============================+=================+======================================================+
21+| 5 310P3 | OK | NA 52 0 / 0 |
22+| 0 0 | 0000:96:00.0 | 0 1901 / 44213 |
23++===============================+=================+======================================================+
24++-------------------------------+-----------------+------------------------------------------------------+
25+| NPU Chip | Process id | Process name | Process memory(MB) |
26++===============================+=================+======================================================+
27+| No running processes found in NPU 5 |
28++===============================+=================+======================================================+
29++ ./execute_div_op
30+float16 result is:
31+0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5
32+float16 test pass
33+float32 result is:
34+0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5 0.5
35+float32 test pass
36+ 
37+Script done on 2026-08-21 14:34:08+08:00 [COMMAND_EXIT_CODE="0"]
38+ 
A2026/CANN-Code-Detective/Challenge04-DivCustomTemplate/iss_liqinghua/DivCustomTemplate/verification.png+0-0