【代码侦探Challenge13】个人开发者 - exercises-cann-outreach-AtomGit

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/README.md+11-0

@@ -0,0 +1,11 @@
	1	+# 个人开发者
	2	+
	3	+## 团队信息
	4	+
	5	+- 提交者: weixin_48190404
	6	+- 身份: 企业员工
	7	+- 单位: 个人开发者
	8	+
	9	+## 成员
	10	+
	11	+- weixin_48190404 (weixin_48190404): 提交者

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/CMakeLists.txt+33-0

@@ -0,0 +1,33 @@
	1	+cmake_minimum_required(VERSION 3.16.0)
	2	+project(leakyrelu_op_prj)
	3	+find_package(ASC REQUIRED)
	4	+if(NOT DEFINED CANNJUDGE_COMPILE_JOBS)
	5	+ set(CANNJUDGE_COMPILE_JOBS 4)
	6	+endif()
	7	+set(CANN_HOST_ARCH "${CMAKE_SYSTEM_PROCESSOR}-linux")
	8	+set(CMAKE_CXX_STANDARD 17)
	9	+set(CMAKE_CXX_STANDARD_REQUIRED ON)
	10	+
	11	+set(ARCH32_COMPUTE_UNITS ascend910b ascend910_93)
	12	+set(ARCH35_COMPUTE_UNITS ascend950)
	13	+
	14	+if(NOT DEFINED ASCEND_COMPUTE_UNIT OR ASCEND_COMPUTE_UNIT STREQUAL "")
	15	+ set(ASCEND_COMPUTE_UNIT ${ARCH32_COMPUTE_UNITS} ${ARCH35_COMPUTE_UNITS})
	16	+endif()
	atomgit-bot22 hours ago 🟡 Medium Priority 变更点：本文件第 14–16 行在未传入 -DASCEND_COMPUTE_UNIT 时，把 ASCEND_COMPUTE_UNIT 默认为 `ascend910b ascend910_93 ascend950`（5–16 行的 ARCH32/ARCH35 列表）。受影响契约：npu_op_package 会按 ASCEND_COMPUTE_UNIT 逐 SoC 打包，而 op_host/leakyrelu_def.cpp:26 只写了 `this->AICore().AddConfig("ascend910b")`，即算子仅在 ascend910b 上注册。失效模式：一旦不显式传参（例如直接 `cmake -S . -B build && cmake --build build`）构建，产出的 vendor 包里 ascend910_93 / ascend950 目录没有该算子（或 arch35 下 kernel 编译/调优路径无对应 op def），在这些 SoC 上运行 examples 会报“算子未注册/找不到算子”，而不是构建期就暴露。当前 build.sh:50/114 显式传了 `-DASCEND_COMPUTE_UNIT=ascend910b`，所以默认路径才触发，实测概率中等。建议：将默认值收敛为算子实际注册的 SoC（ascend910b），或在 op_host/leakyrelu_def.cpp 中为 ARCH32/ARCH35 列表中的每个 SoC 追加 AICore().AddConfig(...)，避免出现“包里有 SoC 目录、但没有该算子”的不一致。 Incorrect?
	17	+set(package_name leakyrelu_custom)
	18	+
	19	+ npu_op_package(${package_name}
	20	+ TYPE RUN
	21	+ CONFIG
	22	+ INSTALL_PATH ${CMAKE_BINARY_DIR}
	23	+)
	24	+
	25	+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_host")
	26	+ add_subdirectory(op_host)
	27	+endif()
	28	+
	29	+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_kernel")
	30	+ add_subdirectory(op_kernel)
	31	+endif()
	32	+
	33	+

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/README.md+5-0

@@ -0,0 +1,5 @@
	1	+# exercises
	2	+
	3	+提交团队: 个人开发者
	4	+提交者: weixin_48190404
	5	+

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/build.sh+177-0

@@ -0,0 +1,177 @@
	1	+#!/bin/bash
	2	+set -e
	3	+
	4	+export BASE_PATH=$(
	5	+ cd "$(dirname $0)"
	6	+ pwd
	7	+)
	8	+export BUILD_PATH="${BASE_PATH}/build"
	9	+export BUILD_OUT_PATH="${BASE_PATH}/build_out"
	10	+
	11	+CORE_NUMS=4
	12	+if [ ${CORE_NUMS} -gt 8 ]; then
	13	+ CORE_NUMS=8
	14	+fi
	15	+
	16	+usage() {
	17	+ echo "Build script for leakyrelu operator"
	18	+ echo "Usage: bash build.sh [OPTIONS]"
	19	+ echo ""
	20	+ echo "Options:"
	21	+ echo " -h, --help Print this help message"
	22	+ echo " -j[n] Compile thread nums, default is ${CORE_NUMS}, eg: -j8"
	23	+ echo " --make_clean Clean build artifacts"
	24	+ echo " -u, --ut Run UT (Unit Tests)"
	25	+ echo " -e, --example Run examples (requires NPU)"
	26	+ echo ""
	27	+ echo "Examples:"
	28	+ echo " bash build.sh # Build with default soc (ascend910b)"
	29	+ echo " bash build.sh -j8 # Build with 8 threads"
	30	+ echo " bash build.sh --make_clean"
	31	+ echo " bash build.sh -u # Run UT tests"
	32	+ echo " bash build.sh -e # Run aclnn example (requires NPU)"
	33	+}
	34	+
	35	+clean_build() {
	36	+ if [ -d "${BUILD_PATH}" ]; then
	37	+ echo "Cleaning build directory..."
	38	+ rm -rf ${BUILD_PATH}/*
	39	+ fi
	40	+}
	41	+
	42	+clean_build_out() {
	43	+ if [ -d "${BUILD_OUT_PATH}" ]; then
	44	+ echo "Cleaning build_out directory..."
	45	+ rm -rf ${BUILD_OUT_PATH}/*
	46	+ fi
	47	+}
	48	+
	49	+THREAD_NUM=${CORE_NUMS}
	50	+COMPUTE_UNIT="ascend910b"
	51	+ENABLE_CLEAN=FALSE
	52	+RUN_UT=FALSE
	53	+RUN_EXAMPLE=FALSE
	54	+
	55	+while [[ $# -gt 0 ]]; do
	56	+ case "$1" in
	57	+ -h\|--help)
	58	+ usage
	59	+ exit 0
	60	+ ;;
	61	+ -j*)
	62	+ THREAD_NUM="${1:2}"
	63	+ if [ -z "$THREAD_NUM" ]; then
	64	+ THREAD_NUM=${CORE_NUMS}
	65	+ fi
	66	+ shift
	67	+ ;;
	68	+ -u\|--ut)
	69	+ RUN_UT=true
	70	+ shift
	71	+ ;;
	72	+ -e\|--example)
	73	+ RUN_EXAMPLE=true
	74	+ shift
	75	+ ;;
	76	+ --make_clean)
	77	+ ENABLE_CLEAN=TRUE
	78	+ shift
	79	+ ;;
	80	+ -*)
	81	+ echo "[ERROR] Invalid option: $1"
	82	+ usage
	83	+ exit 1
	84	+ ;;
	85	+ *)
	86	+ echo "[ERROR] Unexpected argument: $1"
	87	+ usage
	88	+ exit 1
	89	+ ;;
	90	+ esac
	91	+done
	92	+
	93	+if [ "$ENABLE_CLEAN" = "TRUE" ]; then
	94	+ clean_build
	95	+ clean_build_out
	96	+ exit 0
	97	+fi
	98	+
	99	+if [ "$RUN_UT" = true ]; then
	100	+ echo "[INFO] Running UT tests..."
	101	+ export CANNJUDGE_COMPILE_JOBS="$THREAD_NUM"
	102	+ cd "${BASE_PATH}/tests/ut"
	103	+ bash ./run.sh
	104	+ UT_RESULT=$?
	105	+ if [ $UT_RESULT -ne 0 ]; then
	106	+ echo "[ERROR] UT tests failed"
	107	+ exit 1
	108	+ fi
	109	+ echo "[INFO] UT command completed; see PASS/SKIPPED results above."
	110	+ exit 0
	111	+fi
	112	+
	113	+export CANNJUDGE_COMPILE_JOBS="$THREAD_NUM"
	114	+CMAKE_ARGS="-DASCEND_COMPUTE_UNIT=$COMPUTE_UNIT -DCANNJUDGE_COMPILE_JOBS=$THREAD_NUM"
	115	+
	116	+if [ ! -d "${BUILD_PATH}" ]; then
	117	+ mkdir -p "${BUILD_PATH}"
	118	+fi
	119	+
	120	+[ -f "${BUILD_PATH}/CMakeCache.txt" ] && rm -f ${BUILD_PATH}/CMakeCache.txt
	121	+
	122	+echo "----------------------------------------------------------------"
	123	+SOURCE_FINGERPRINT=$(
	124	+ { printf '%s\n' "$ASCEND_HOME_PATH";
	125	+ for d in op_host op_kernel op_api; do
	126	+ if [ -d "$BASE_PATH/$d" ]; then find "$BASE_PATH/$d" -type f -print0; fi
	127	+ done \| sort -z \| xargs -0 sha256sum;
	128	+ sha256sum "$BASE_PATH/CMakeLists.txt";
	129	+ } \| sha256sum
	130	+)
	131	+if [ -L "$BUILD_PATH" ]; then echo "build must not be a symlink"; exit 2; fi
	132	+if [ ! -f "$BUILD_PATH/.source-fingerprint" ] \|\| [ "$(cat "$BUILD_PATH/.source-fingerprint")" != "$SOURCE_FINGERPRINT" ]; then
	133	+ rm -rf "$BUILD_PATH"
	134	+ mkdir -p "$BUILD_PATH"
	135	+ printf '%s' "$SOURCE_FINGERPRINT" > "$BUILD_PATH/.source-fingerprint"
	136	+fi
	137	+
	138	+echo "[INFO] Configuring project..."
	139	+echo "[INFO] CMAKE_ARGS: ${CMAKE_ARGS}"
	140	+cd "${BUILD_PATH}" && cmake ${CMAKE_ARGS} ..
	141	+
	142	+echo "----------------------------------------------------------------"
	143	+echo "[INFO] Building project with ${THREAD_NUM} threads..."
	144	+for target in all binary package install; do
	145	+ cmake --build . --target "$target" --parallel "$THREAD_NUM"
	146	+done
	147	+
	148	+KERNEL_O=$(find ${BUILD_PATH}/op_kernel/ascendc_kernels/binary/${COMPUTE_UNIT} -name "*.o" 2>/dev/null \| head -1)
	149	+if [ -z "$KERNEL_O" ]; then
	150	+ echo "[ERROR] Kernel binary not found"
	151	+ exit 1
	152	+fi
	153	+
	154	+PKG_PATH=$(ls "${BUILD_PATH}"/custom_opp_*.run 2>/dev/null \| head -n 1)
	155	+if [ -z "$PKG_PATH" ] \|\| [ ! -f "$PKG_PATH" ] \|\| [ ! -s "$PKG_PATH" ]; then
	156	+ echo "[ERROR] Package not found or empty"
	157	+ exit 1
	158	+fi
	159	+
	160	+echo "----------------------------------------------------------------"
	161	+echo "[INFO] Build completed successfully!"
	162	+echo "[INFO] Kernel binary: ${KERNEL_O}"
	163	+echo "[INFO] Package: ${PKG_PATH}"
	164	+
	165	+if [ "$RUN_EXAMPLE" = true ]; then
	166	+ echo "----------------------------------------------------------------"
	167	+ echo "[INFO] Running examples..."
	168	+ cd "${BASE_PATH}/examples"
	169	+ bash ./run.sh
	170	+ EXAMPLE_RESULT=$?
	171	+ cd - > /dev/null
	172	+ if [ $EXAMPLE_RESULT -ne 0 ]; then
	173	+ echo "[ERROR] Example execution failed"
	174	+ exit 1
	175	+ fi
	176	+ echo "[INFO] Example completed successfully!"
	177	+fi

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/examples/CMakeLists.txt+9-0

@@ -0,0 +1,9 @@
	1	+cmake_minimum_required(VERSION 3.16)
	2	+project(local_example CXX)
	3	+set(CMAKE_CXX_STANDARD 17)
	4	+set(OPP_ROOT "${CMAKE_CURRENT_SOURCE_DIR}/../build/packages/vendors/leakyrelu_custom" CACHE PATH "Built operator package")
	5	+find_path(OP_INCLUDE aclnn_leakyrelu.h PATHS "${OPP_ROOT}/op_api/include" NO_DEFAULT_PATH REQUIRED)
	6	+find_library(OP_LIBRARY cust_opapi PATHS "${OPP_ROOT}/op_api/lib" NO_DEFAULT_PATH REQUIRED)
	7	+add_executable(sample test_aclnn_leakyrelu.cpp)
	8	+target_include_directories(sample PRIVATE "${OP_INCLUDE}" "$ENV{ASCEND_HOME_PATH}/include")
	9	+target_link_libraries(sample PRIVATE "${OP_LIBRARY}" "$ENV{ASCEND_HOME_PATH}/lib64/libascendcl.so" "$ENV{ASCEND_HOME_PATH}/lib64/libnnopbase.so")

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/examples/run.sh+11-0

@@ -0,0 +1,11 @@
	1	+#!/bin/bash
	2	+set -euo pipefail
	3	+ROOT="$(cd "$(dirname "$0")/.." && pwd)"
	4	+: "${ASCEND_HOME_PATH:?source the CANN set_env.sh first}"
	5	+export ASCEND_CUSTOM_OPP_PATH="$ROOT/build/packages/vendors/leakyrelu_custom${ASCEND_CUSTOM_OPP_PATH:+:$ASCEND_CUSTOM_OPP_PATH}"
	6	+export LD_LIBRARY_PATH="$ROOT/build/packages/vendors/leakyrelu_custom/op_api/lib:$ASCEND_HOME_PATH/lib64:${LD_LIBRARY_PATH:-}"
	7	+cmake -S "$ROOT/examples" -B "$ROOT/examples/build"
	8	+cmake --build "$ROOT/examples/build" --parallel "${CANNJUDGE_COMPILE_JOBS:-4}"
	9	+python3 "$ROOT/samples/sample.py" generate --directory "$ROOT/samples/build"
	10	+"$ROOT/examples/build/sample" "$ROOT/samples/build" "${DEVICE_ID:-0}"
	11	+python3 "$ROOT/samples/sample.py" verify --directory "$ROOT/samples/build"

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/examples/sample_support.h+82-0

@@ -0,0 +1,82 @@
	1	+#pragma once
	2	+#include <acl/acl.h>
	3	+#include <aclnn/aclnn_base.h>
	4	+#include <array>
	5	+#include <vector>
	6	+#include <string>
	7	+#include <fstream>
	8	+#include <iostream>
	9	+#include <stdexcept>
	10	+#include <algorithm>
	11	+#include <cstdlib>
	12	+
	13	+inline void check(int rc) {
	14	+ if (rc)
	15	+ throw std::runtime_error("ACL error " + std::to_string(rc));
	16	+}
	17	+inline std::vector<char> read_file(const std::string& name, size_t size) {
	18	+ std::ifstream file(name, std::ios::binary \| std::ios::ate);
	19	+ if (!file \|\| file.tellg() != static_cast<std::streamoff>(size))
	20	+ throw std::runtime_error("input size: " + name);
	21	+ std::vector<char> data(size);
	22	+ file.seekg(0);
	23	+ if (size && !file.read(data.data(), size))
	24	+ throw std::runtime_error("read: " + name);
	25	+ return data;
	26	+}
	27	+struct Tensor {
	28	+ void* device = nullptr;
	29	+ aclTensor* tensor = nullptr;
	30	+ aclScalar* scalar = nullptr;
	31	+ size_t bytes;
	32	+ std::string file;
	33	+ std::vector<char> host;
	34	+ Tensor(std::vector<int64_t> shape, aclDataType type, size_t size, const char* name, bool input,
	35	+ bool is_scalar, const std::string& root)
	36	+ : bytes(size), file(name) {
	37	+ host = input ? read_file(root + "/input/" + file, size) : std::vector<char>(size);
	38	+ if (is_scalar) {
	39	+ scalar = aclCreateScalar(host.data(), type);
	40	+ if (!scalar)
	41	+ throw std::runtime_error("scalar creation failed");
	42	+ return;
	43	+ }
	44	+ check(aclrtMalloc(&device, std::max(size_t(1), size), ACL_MEM_MALLOC_HUGE_FIRST));
	45	+ if (size)
	46	+ check(aclrtMemcpy(device, size, host.data(), size, ACL_MEMCPY_HOST_TO_DEVICE));
	47	+ std::vector<int64_t> strides(shape.size(), 1);
	48	+ for (size_t i = shape.size(); i > 1; --i)
	49	+ strides[i - 2] = strides[i - 1] * shape[i - 1];
	50	+ tensor = aclCreateTensor(shape.data(), shape.size(), type, strides.data(), 0, ACL_FORMAT_ND,
	51	+ shape.data(), shape.size(), device);
	52	+ if (!tensor)
	53	+ throw std::runtime_error("tensor creation failed");
	54	+ }
	55	+ void save(const std::string& root) {
	56	+ if (bytes)
	57	+ check(aclrtMemcpy(host.data(), bytes, device, bytes, ACL_MEMCPY_DEVICE_TO_HOST));
	58	+ std::ofstream out(root + "/output/" + file, std::ios::binary);
	59	+ if (!out \|\| (bytes && !out.write(host.data(), bytes)))
	60	+ throw std::runtime_error("write: " + file);
	61	+ }
	62	+ ~Tensor() {
	63	+ if (scalar)
	64	+ aclDestroyScalar(scalar);
	65	+ if (tensor)
	66	+ aclDestroyTensor(tensor);
	67	+ if (device)
	68	+ aclrtFree(device);
	69	+ }
	70	+};
	71	+struct TensorList {
	72	+ aclTensorList* list;
	73	+ TensorList(std::initializer_list<aclTensor*> values) {
	74	+ std::vector<aclTensor*> items(values);
	75	+ list = aclCreateTensorList(items.data(), items.size());
	76	+ if (!list)
	77	+ throw std::runtime_error("tensor list creation failed");
	78	+ }
	79	+ ~TensorList() {
	80	+ aclDestroyTensorList(list);
	81	+ }
	82	+};

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/examples/test_aclnn_leakyrelu.cpp+41-0

@@ -0,0 +1,41 @@
	1	+#include "sample_support.h"
	2	+#include "aclnn_leakyrelu.h"
	3	+
	4	+int main(int argc, char** argv) {
	5	+ try {
	6	+ std::string root = argc > 1 ? argv[1] : "../samples/build";
	7	+ int device = argc > 2 ? std::stoi(argv[2]) : 0;
	8	+ check(aclInit(nullptr));
	9	+ check(aclrtSetDevice(device));
	10	+ aclrtStream stream;
	11	+ check(aclrtCreateStream(&stream));
	12	+ {
	13	+ Tensor input_0_0({8, 2048}, ACL_FLOAT, 65536, "x.bin", true, false, root);
	14	+ Tensor output_0_0({8, 2048}, ACL_FLOAT, 65536, "y.bin", false, false, root);
	15	+ double attr_0 = static_cast<double>(0.01);
	16	+ uint64_t size = 0;
	17	+ aclOpExecutor* executor = nullptr;
	18	+ check(aclnnLeakyreluGetWorkspaceSize(
	19	+ input_0_0.tensor,
	20	+ attr_0,
	21	+ output_0_0.tensor,
	22	+ &size,
	23	+ &executor));
	24	+ void* workspace = nullptr;
	25	+ if (size)
	26	+ check(aclrtMalloc(&workspace, size, ACL_MEM_MALLOC_HUGE_FIRST));
	27	+ check(aclnnLeakyrelu(workspace, size, executor, stream));
	28	+ check(aclrtSynchronizeStream(stream));
	29	+ output_0_0.save(root);
	30	+ if (workspace)
	31	+ check(aclrtFree(workspace));
	32	+ }
	33	+ check(aclrtDestroyStream(stream));
	34	+ check(aclrtResetDevice(device));
	35	+ check(aclFinalize());
	36	+ return 0;
	37	+ } catch (const std::exception& e) {
	38	+ std::cerr << e.what() << std::endl;
	39	+ return 2;
	40	+ }
	41	+}

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_host/CMakeLists.txt+102-0

@@ -0,0 +1,102 @@
	1	+file(GLOB host_ops_def_srcs
	2	+ ${CMAKE_CURRENT_SOURCE_DIR}/*def.cpp
	3	+)
	4	+
	5	+file(GLOB host_ops_infershape_srcs
	6	+ ${CMAKE_CURRENT_SOURCE_DIR}/*_infershape.cpp
	7	+)
	8	+
	9	+set(host_ops_tiling_srcs)
	10	+file(GLOB TILING_FILES ${CMAKE_CURRENT_SOURCE_DIR}/*tiling.cpp)
	11	+list(APPEND host_ops_tiling_srcs ${TILING_FILES})
	12	+
	13	+set(host_ops_srcs
	14	+ ${host_ops_def_srcs}
	15	+ ${host_ops_infershape_srcs}
	16	+ ${host_ops_tiling_srcs}
	17	+)
	18	+
	19	+npu_op_code_gen(
	20	+ JOBS ${CANNJUDGE_COMPILE_JOBS}
	21	+ SRC ${host_ops_srcs}
	22	+ PACKAGE ${package_name}
	23	+ OUT_DIR ${ASCEND_AUTOGEN_PATH}
	24	+ COMPILE_OPTIONS
	25	+ -I$ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/include
	26	+ -I$ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include/tiling
	27	+ -I$ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc
	28	+ -I$ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/op_common
	29	+ -I$ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/base
	30	+ -I$ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/exe_graph
	31	+ -I$ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/graph
	32	+)
	33	+
	34	+npu_op_library(cust_optiling TILING
	35	+ ${host_ops_srcs}
	36	+)
	37	+
	38	+target_include_directories(cust_optiling PRIVATE
	39	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/include
	40	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include/tiling
	41	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc
	42	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/op_common
	43	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/base
	44	+)
	45	+
	46	+set(op_api_dir ${CMAKE_CURRENT_SOURCE_DIR}/../op_api)
	47	+if(EXISTS ${op_api_dir} AND IS_DIRECTORY ${op_api_dir})
	48	+ file(GLOB op_api_srcs ${op_api_dir}/*.cpp)
	49	+else()
	50	+ file(GLOB op_api_srcs "${CMAKE_BINARY_DIR}/autogen/aclnn_*.cpp")
	51	+endif()
	52	+
	53	+npu_op_library(cust_opapi ACLNN
	54	+ ${op_api_srcs}
	55	+)
	56	+
	57	+target_include_directories(cust_opapi PRIVATE
	58	+ ${op_api_dir}
	59	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/include
	60	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/include/aclnn
	61	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include
	62	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc
	63	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/op_common
	64	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/base
	65	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/aicpu
	66	+)
	67	+
	68	+target_compile_options(cust_opapi PRIVATE -UACLNN_WITH_BINARY)
	69	+
	70	+file(GLOB proto_src ${ASCEND_AUTOGEN_PATH}/op_proto.cc)
	71	+set_source_files_properties(${proto_src} PROPERTIES GENERATED TRUE)
	72	+
	73	+npu_op_library(cust_op_proto GRAPH
	74	+ ${host_ops_srcs}
	75	+ ${proto_src}
	76	+)
	77	+
	78	+target_include_directories(cust_op_proto PRIVATE
	79	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/include
	80	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include/tiling
	81	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc
	82	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/op_common
	83	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/base
	84	+)
	85	+
	86	+npu_op_package_add(${package_name}
	87	+ LIBRARY
	88	+ cust_optiling
	89	+ cust_op_proto
	90	+ cust_opapi
	91	+)
	92	+
	93	+# Install hand-written op_api headers into RUN package
	94	+if(EXISTS ${op_api_dir} AND IS_DIRECTORY ${op_api_dir})
	95	+ file(GLOB op_api_aclnn_headers ${op_api_dir}/aclnn_.h ${op_api_dir}/acl_.h)
	96	+ if(op_api_aclnn_headers)
	97	+ npu_op_package_add(${package_name}
	98	+ FILES ${op_api_aclnn_headers}
	99	+ TYPE ACLNN
	100	+ )
	101	+ endif()
	102	+endif()

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_host/leakyrelu_def.cpp+30-0

@@ -0,0 +1,30 @@
	1	+/*!
	2	+ * \file leakyrelu_def.cpp
	3	+ * \brief Leakyrelu 算子定义
	4	+ */
	5	+#include "register/op_def_registry.h"
	6	+
	7	+namespace ops {
	8	+class Leakyrelu : public OpDef {
	9	+public:
	10	+ explicit Leakyrelu(const char* name) : OpDef(name)
	11	+ {
	12	+ this->Input("x")
	13	+ .ParamType(REQUIRED)
	14	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
	15	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	16	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	17	+ .AutoContiguous();
	18	+ this->Output("y")
	19	+ .ParamType(REQUIRED)
	20	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
	21	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	22	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	23	+ .AutoContiguous();
	24	+ this->Attr("negative_slope")
	25	+ .Float();
	26	+ this->AICore().AddConfig("ascend910b");
	27	+ }
	28	+};
	29	+OP_ADD(Leakyrelu);
	30	+} // namespace ops

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_host/leakyrelu_infershape.cpp+37-0

@@ -0,0 +1,37 @@
	1	+/*!
	2	+ * \file leakyrelu_infershape.cpp
	3	+ * \brief Leakyrelu 算子形状推导实现
	4	+ */
	5	+
	6	+#include "register/op_impl_registry.h"
	7	+#include "exe_graph/runtime/infer_shape_context.h"
	8	+
	9	+using namespace ge;
	10	+
	11	+namespace ops {
	12	+
	13	+static ge::graphStatus InferShapeLeakyrelu(gert::InferShapeContext* context)
	14	+{
	15	+ // TODO: 实现形状推导逻辑
	16	+ const gert::Shape* input_shape = (1 > 0) ? context->GetInputShape(0) : nullptr;
	17	+ // 注意：无输入算子时 input_shape 为 nullptr，需在此处手动设置输出 shape
	18	+
	19	+ for (size_t i = 0; i < 1; i++) {
	20	+ gert::Shape* output_shape = context->GetOutputShape(i);
	21	+ if (output_shape == nullptr) {
	22	+ return ge::GRAPH_FAILED;
	23	+ }
	24	+ const gert::Shape* in_shape = (i < 1) ? context->GetInputShape(i) : input_shape;
	25	+ if (in_shape == nullptr) {
	26	+ in_shape = input_shape;
	27	+ }
	28	+ if (in_shape != nullptr) {
	29	+ output_shape = in_shape;
	30	+ }
	31	+ }
	32	+ return ge::GRAPH_SUCCESS;
	33	+}
	34	+
	35	+IMPL_OP_INFERSHAPE(Leakyrelu).InferShape(InferShapeLeakyrelu);
	36	+
	37	+} // namespace ops

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_host/leakyrelu_tiling.cpp+156-0

@@ -0,0 +1,156 @@
	1	+/*!
	2	+ * \file leakyrelu_tiling.cpp
	3	+ * \brief Leakyrelu 算子 Tiling 实现
	4	+ */
	5	+
	6	+#include <algorithm>
	7	+
	8	+#include "register/op_def_registry.h"
	9	+#include "op_common/log/log.h"
	10	+#include "op_common/op_host/util/math_util.h"
	11	+#include "op_common/op_host/util/platform_util.h"
	12	+#include "../op_kernel/leakyrelu_tiling_data.h"
	13	+#include "../op_kernel/leakyrelu_tiling_key.h"
	14	+
	15	+namespace optiling {
	16	+
	17	+using Ops::Base::CeilDiv;
	18	+using Ops::Base::CeilAlign;
	19	+using Ops::Base::FloorDiv;
	20	+using Ops::Base::FloorAlign;
	21	+using Ops::Base::GetUbBlockSize;
	22	+
	23	+constexpr uint32_t WS_SYS_SIZE = 0U;
	24	+constexpr int64_t TYPE_SIZE = 4;
	25	+constexpr int64_t MIN_SPLIT_THRESHOLD = 1024;
	26	+constexpr int64_t BLOCK_BYTES = 32;
	27	+constexpr int64_t BUFFER_NUM = 2;
	28	+
	29	+static const gert::Shape g_vec_1_shape = {1};
	30	+
	31	+static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape) {
	32	+ if (in_shape.GetDimNum() == 0) {
	33	+ return g_vec_1_shape;
	34	+ }
	35	+ return in_shape;
	36	+}
	37	+
	38	+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
	39	+{
	40	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	41	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	42	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	43	+ coreNum = ascendcPlatform.GetCoreNumAiv();
	44	+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
	45	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
	46	+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
	47	+ return ge::GRAPH_SUCCESS;
	48	+}
	49	+
	50	+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
	51	+{
	52	+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
	53	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	54	+ currentWorkspace[0] = WS_SYS_SIZE;
	55	+ return ge::GRAPH_SUCCESS;
	56	+}
	57	+
	58	+// 读取 negative_slope 属性；属性缺失时退回默认值
	59	+static float GetNegativeSlope(gert::TilingContext* context)
	60	+{
	61	+ const auto* attrs = context->GetAttrs();
	62	+ if (attrs == nullptr) {
	63	+ return LEAKYRELU_DEFAULT_SLOPE;
	64	+ }
	65	+ const float* slopePtr = attrs->GetFloat(0);
	66	+ if (slopePtr == nullptr) {
	67	+ slopePtr = attrs->GetAttrPointer<float>(0);
	68	+ }
	69	+ return (slopePtr == nullptr) ? LEAKYRELU_DEFAULT_SLOPE : *slopePtr;
	70	+}
	71	+
	72	+static ge::graphStatus LeakyreluTilingFunc(gert::TilingContext* context)
	73	+{
	74	+ uint64_t ubSize = 0;
	75	+ int64_t coreNum = 0;
	76	+ OP_CHECK_IF(
	77	+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
	78	+ OP_LOGE(context, "GetPlatformInfo error"),
	79	+ return ge::GRAPH_FAILED);
	80	+
	81	+ OP_CHECK_IF(
	82	+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
	83	+ OP_LOGE(context, "GetWorkspaceSize error"),
	84	+ return ge::GRAPH_FAILED);
	85	+
	86	+ const auto* inputDesc = context->GetInputDesc(0);
	87	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	88	+ const bool isHalf = (inputDesc->GetDataType() == ge::DT_FLOAT16 \|\|
	89	+ inputDesc->GetDataType() == ge::DT_BF16);
	90	+ const int64_t typeSize = isHalf ? 2 : TYPE_SIZE;
	91	+ const int64_t alignElems = BLOCK_BYTES / typeSize; // float:8, half:16
	92	+
	93	+ const gert::StorageShape* storageShape = context->GetInputShape(0);
	94	+ OP_CHECK_NULL_WITH_CONTEXT(context, storageShape);
	95	+ const gert::Shape shape = EnsureNotScalar(storageShape->GetStorageShape());
	96	+ int64_t totalNum = 1;
	97	+ for (size_t i = 0; i < shape.GetDimNum(); ++i) {
	98	+ totalNum *= shape.GetDim(i);
	99	+ }
	100	+ OP_CHECK_IF(totalNum < 0, OP_LOGE(context, "invalid input shape"), return ge::GRAPH_FAILED);
	101	+
	102	+ // 1) 多核切分：每核至少 MIN_SPLIT_THRESHOLD 个元素，块长 32B 对齐
	103	+ const int64_t splitNum = (totalNum > 0) ? totalNum : 1;
	104	+ int64_t usedCoreNum = std::min(coreNum, std::max<int64_t>(1, CeilDiv(splitNum, MIN_SPLIT_THRESHOLD)));
	105	+ int64_t blockFactor = CeilAlign(CeilDiv(splitNum, usedCoreNum), alignElems);
	106	+ if (blockFactor < alignElems) {
	107	+ blockFactor = alignElems;
	108	+ }
	109	+ usedCoreNum = std::min(coreNum, std::max<int64_t>(1, CeilDiv(splitNum, blockFactor)));
	110	+
	111	+ // 2) UB 切分：输入/输出各 2 块双缓冲，half 分支另需 2 块 float 中间缓冲
	112	+ // 按最坏情况估算，保证所有 dtype 都不会超出 UB 容量
	113	+ const int64_t bytesPerElem = 2 * BUFFER_NUM * typeSize + 2 * static_cast<int64_t>(sizeof(float));
	114	+ int64_t ubCapacity = FloorAlign(static_cast<int64_t>(ubSize) / bytesPerElem, alignElems);
	115	+ if (ubCapacity < alignElems) {
	116	+ ubCapacity = alignElems;
	117	+ }
	118	+ const int64_t ubFactor = std::min(blockFactor, ubCapacity);
	119	+
	120	+ // 3) 填充 tiling 数据，并把属性透传给 kernel
	121	+ LeakyreluTilingData* tiling = context->GetTilingData<LeakyreluTilingData>();
	122	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	123	+ tiling->totalNum = totalNum;
	124	+ tiling->blockFactor = blockFactor;
	125	+ tiling->ubFactor = ubFactor;
	126	+ tiling->negativeSlope = GetNegativeSlope(context);
	127	+ tiling->slopeTag = LEAKYRELU_SLOPE_TAG;
	128	+ tiling->reserved = 0.0f;
	129	+
	130	+ auto* rawTilingData = context->GetRawTilingData();
	131	+ OP_CHECK_NULL_WITH_CONTEXT(context, rawTilingData);
	132	+ rawTilingData->SetDataSize(sizeof(LeakyreluTilingData));
	133	+
	134	+ context->SetBlockDim(static_cast<uint32_t>(usedCoreNum));
	135	+
	136	+ // 根据输入 dtype 选择 tilingKey：half -> mode 0，float -> mode 1
	137	+ uint64_t tilingKey;
	138	+ if (isHalf) {
	139	+ tilingKey = GET_TPL_TILING_KEY(LEAKYRELU_TPL_SCH_MODE_0);
	140	+ } else {
	141	+ tilingKey = GET_TPL_TILING_KEY(LEAKYRELU_TPL_SCH_MODE_1);
	142	+ }
	143	+ context->SetTilingKey(tilingKey);
	144	+ return ge::GRAPH_SUCCESS;
	145	+}
	146	+
	147	+static ge::graphStatus TilingParseForLeakyrelu([[maybe_unused]] gert::TilingParseContext* context)
	148	+{
	149	+ return ge::GRAPH_SUCCESS;
	150	+}
	151	+
	152	+struct LeakyreluCompileInfo {};
	153	+
	154	+IMPL_OP_OPTILING(Leakyrelu).Tiling(LeakyreluTilingFunc).TilingParse<LeakyreluCompileInfo>(TilingParseForLeakyrelu);
	155	+
	156	+} // namespace optiling

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_kernel/CMakeLists.txt+17-0

@@ -0,0 +1,17 @@
	1	+file(GLOB_RECURSE ALL_KERNEL_FILES RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} *.cpp)
	2	+
	3	+npu_op_kernel_sources(ascendc_kernels
	4	+ OP_TYPE OP
	5	+ KERNEL_DIR .
	6	+ KERNEL_FILE ${ALL_KERNEL_FILES}
	7	+)
	8	+
	9	+npu_op_kernel_library(ascendc_kernels
	10	+ SRC_BASE ${CMAKE_CURRENT_SOURCE_DIR}
	11	+ TILING_LIBRARY cust_optiling
	12	+)
	13	+
	14	+npu_op_package_add(${package_name}
	15	+ LIBRARY
	16	+ ascendc_kernels
	17	+)

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_kernel/leakyrelu.cpp+29-0

@@ -0,0 +1,29 @@
	1	+/*!
	2	+ * \file leakyrelu.cpp
	3	+ * \brief Leakyrelu 算子 kernel 入口
	4	+ */
	5	+
	6	+#include "leakyrelu.h"
	7	+
	8	+enum class LeakyreluTilingKey : uint32_t
	9	+{
	10	+ TILING_KEY_LEAKYRELU_MODE_0 = 0,
	11	+ TILING_KEY_LEAKYRELU_MODE_1 = 1,
	12	+};
	13	+
	14	+template <uint32_t schMode>
	15	+__global__ __aicore__ void leakyrelu(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
	16	+{
	17	+ REGISTER_TILING_DEFAULT(LeakyreluTilingData);
	18	+ GET_TILING_DATA_WITH_STRUCT(LeakyreluTilingData, tilingData, tiling);
	19	+ if constexpr (schMode == static_cast<uint32_t>(LeakyreluTilingKey::TILING_KEY_LEAKYRELU_MODE_0)) {
	20	+ NsLeakyrelu::Leakyrelu<half> op;
	21	+ op.Init(x, y, &tilingData);
	22	+ op.Process();
	23	+ }
	24	+ if constexpr (schMode == static_cast<uint32_t>(LeakyreluTilingKey::TILING_KEY_LEAKYRELU_MODE_1)) {
	25	+ NsLeakyrelu::Leakyrelu<float> op;
	26	+ op.Init(x, y, &tilingData);
	27	+ op.Process();
	28	+ }
	29	+}

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_kernel/leakyrelu.h+186-0

@@ -0,0 +1,186 @@
	1	+/*!
	2	+ * \file leakyrelu.h
	3	+ * \brief Leakyrelu 算子 kernel 类定义
	4	+ */
	5	+
	6	+#ifndef LEAKYRELU_H
	7	+#define LEAKYRELU_H
	8	+
	9	+#include "kernel_operator.h"
	10	+#include "kernel_tiling/kernel_tiling.h"
	11	+#include "leakyrelu_tiling_data.h"
	12	+#include "leakyrelu_tiling_key.h"
	13	+
	14	+namespace NsLeakyrelu {
	15	+
	16	+using namespace AscendC;
	17	+
	18	+constexpr int32_t BUFFER_NUM = 2;
	19	+constexpr int64_t BLOCK_BYTES = 32; // GM/UB 搬运的对齐粒度
	20	+
	21	+// 用 sizeof 判断 dtype（T 只会是 half=2 字节或 float=4 字节），
	22	+// 不依赖 type_traits，也不依赖 half 的具体类型定义
	23	+constexpr int64_t HALF_TYPE_SIZE = 2;
	24	+
	25	+template <typename T>
	26	+class Leakyrelu {
	27	+public:
	28	+ __aicore__ inline Leakyrelu(){};
	29	+
	30	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const LeakyreluTilingData* tilingData);
	31	+ __aicore__ inline void Process();
	32	+
	33	+private:
	34	+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
	35	+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
	36	+ __aicore__ inline void Compute(int64_t currentNum);
	37	+
	38	+ // y = max(x, 0) + negative_slope * min(x, 0)，在 float 域上完成
	39	+ __aicore__ inline void ComputeFloat(const LocalTensor<float>& src, const LocalTensor<float>& dst,
	40	+ int64_t count);
	41	+ __aicore__ inline int64_t AlignUp(int64_t value, int64_t align)
	42	+ {
	43	+ return (value + align - 1) / align * align;
	44	+ }
	45	+
	46	+private:
	47	+ TPipe pipe;
	48	+ TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX;
	49	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueY;
	50	+ TBuf<TPosition::VECCALC> floatBufX; // 仅 half 分支使用
	51	+ TBuf<TPosition::VECCALC> floatBufY; // 仅 half 分支使用
	52	+
	53	+ GlobalTensor<T> inputGMX;
	54	+ GlobalTensor<T> outputGMY;
	55	+
	56	+ int64_t blockLength_ = 0;
	57	+ int64_t ubLength_ = 0;
	58	+ int64_t alignElems_ = 1;
	59	+ float negativeSlope_ = LEAKYRELU_DEFAULT_SLOPE;
	60	+};
	61	+
	62	+template <typename T>
	63	+__aicore__ inline void Leakyrelu<T>::Init(GM_ADDR x, GM_ADDR y, const LeakyreluTilingData* tilingData)
	64	+{
	65	+ alignElems_ = BLOCK_BYTES / static_cast<int64_t>(sizeof(T));
	66	+ // 兼容只填写了前三个字段的 CPU kernel UT fixture（未下发属性时使用默认 slope）
	67	+ negativeSlope_ = (tilingData->slopeTag == LEAKYRELU_SLOPE_TAG) ? tilingData->negativeSlope
	68	+ : LEAKYRELU_DEFAULT_SLOPE;
	69	+ blockLength_ = 0;
	70	+ ubLength_ = (tilingData->ubFactor > 0) ? tilingData->ubFactor : tilingData->blockFactor;
	71	+ if (ubLength_ <= 0) {
	72	+ ubLength_ = alignElems_;
	73	+ }
	74	+
	75	+ const int64_t blockIdx = static_cast<int64_t>(GetBlockIdx());
	76	+ const int64_t offset = blockIdx * tilingData->blockFactor;
	77	+ if (offset >= tilingData->totalNum) {
	78	+ return; // 多余的核空跑
	79	+ }
	80	+ const int64_t remain = tilingData->totalNum - offset;
	81	+ blockLength_ = (remain < tilingData->blockFactor) ? remain : tilingData->blockFactor;
	82	+ if (blockLength_ <= 0) {
	83	+ return;
	84	+ }
	85	+ if (ubLength_ > blockLength_) {
	86	+ ubLength_ = AlignUp(blockLength_, alignElems_);
	87	+ }
	88	+ if (ubLength_ < alignElems_) {
	89	+ ubLength_ = alignElems_;
	90	+ }
	91	+ ubLength_ = AlignUp(ubLength_, alignElems_);
	92	+
	93	+ // 每个核只看自己那一段，起始地址天然 32B 对齐（blockFactor 已对齐）；
	94	+ // 声明长度按对齐长度取，兼容尾块按 copyNum 对齐搬入
	95	+ const int64_t gmLength = AlignUp(blockLength_, alignElems_);
	96	+ inputGMX.SetGlobalBuffer((__gm__ T*)x + offset, gmLength);
	97	+ outputGMY.SetGlobalBuffer((__gm__ T*)y + offset, gmLength);
	98	+
	99	+ // 输入/输出双缓冲；尾块按对齐长度计算时仍不会超过该容量
	100	+ const uint32_t bufferBytes = static_cast<uint32_t>(ubLength_ * static_cast<int64_t>(sizeof(T)));
	101	+ pipe.InitBuffer(inputQueueX, BUFFER_NUM, bufferBytes);
	102	+ pipe.InitBuffer(outputQueueY, BUFFER_NUM, bufferBytes);
	103	+ if constexpr (static_cast<int64_t>(sizeof(T)) == HALF_TYPE_SIZE) {
	104	+ // half 直接乘 slope 会引入半精度舍入误差，这里升到 float 计算
	105	+ const uint32_t floatBytes = static_cast<uint32_t>(ubLength_ * static_cast<int64_t>(sizeof(float)));
	106	+ pipe.InitBuffer(floatBufX, floatBytes);
	107	+ pipe.InitBuffer(floatBufY, floatBytes);
	108	+ }
	109	+}
	110	+
	111	+template <typename T>
	112	+__aicore__ inline void Leakyrelu<T>::CopyIn(int64_t progress, int64_t currentNum)
	113	+{
	114	+ LocalTensor<T> xLocal = inputQueueX.AllocTensor<T>();
	115	+ const int64_t offset = progress * ubLength_;
	116	+ // 按 32B 对齐长度搬入；尾部多读的元素只参与计算，不会被写回 GM
	117	+ const int64_t copyNum = AlignUp(currentNum, alignElems_);
	118	+ DataCopy(xLocal, inputGMX[offset], static_cast<int32_t>(copyNum));
	atomgit-bot22 hours ago 🟡 Medium Priority 变更行：leakyrelu.h:95-97 用 `gmLength = AlignUp(blockLength_, alignElems_)` 声明 GlobalTensor 长度，leakyrelu.h:117-118 再用 `copyNum = AlignUp(currentNum, alignElems_)` 搬入。证据链：host 侧 tiling（op_host/leakyrelu_tiling.cpp:96-107）计算 `totalNum` 为各维乘积，未做 32B 对齐，只有 `blockFactor = CeilAlign(CeilDiv(splitNum, usedCoreNum), alignElems)` 对齐；因此最后一个核的 `blockLength_ = totalNum - offset` 可以是任意值（例如 float 输入 9 个元素、blockFactor=8 时，末核 blockLength_=1）。此时 `gmLength = 8`、`copyNum = 8`，但该核实际可用元素只有 1 个。失败模式：DataCopy 会读取 `x + totalNum` 之后最多 alignElems_-1 个元素（float 为 28 字节、half 为 30 字节）的越界 GM 数据；SetGlobalBuffer 声明的长度本身也已超过真实张量长度，越界读到的内容依赖相邻内存，虽只用于计算不回写，但越过输入张量末尾属于典型 GM 越界访问，在尾部恰好压到内存池边界时可能触发 AIC Error/精度校验异常。建议：把尾块搬入改为 DataCopyPad（blockLen 取 currentNum*sizeof(T)，越界部分由 pad 参数补齐），并与 CopyOut 的对齐策略保持一致，避免按 AlignUp 后的长度读取输入张量之外的内存。 Incorrect?
	119	+ inputQueueX.EnQue(xLocal);
	120	+}
	121	+
	122	+template <typename T>
	123	+__aicore__ inline void Leakyrelu<T>::ComputeFloat(const LocalTensor<float>& src,
	124	+ const LocalTensor<float>& dst, int64_t count)
	125	+{
	126	+ const int32_t n = static_cast<int32_t>(count);
	127	+ // dst = max(x, 0)，即正半轴
	128	+ Relu(dst, src, n);
	129	+ // src = x - max(x, 0) = min(x, 0)（原地复用输入缓冲）
	130	+ Sub(src, src, dst, n);
	131	+ // src = negative_slope * min(x, 0)
	132	+ Muls(src, src, negativeSlope_, n);
	133	+ // dst = max(x, 0) + negative_slope * min(x, 0)
	134	+ Add(dst, dst, src, n);
	135	+}
	136	+
	137	+template <typename T>
	138	+__aicore__ inline void Leakyrelu<T>::Compute(int64_t currentNum)
	139	+{
	140	+ const int64_t count = AlignUp(currentNum, alignElems_);
	141	+ LocalTensor<T> xLocal = inputQueueX.DeQue<T>();
	142	+ LocalTensor<T> yLocal = outputQueueY.AllocTensor<T>();
	143	+
	144	+ if constexpr (static_cast<int64_t>(sizeof(T)) == HALF_TYPE_SIZE) {
	145	+ LocalTensor<float> xFloat = floatBufX.Get<float>();
	146	+ LocalTensor<float> yFloat = floatBufY.Get<float>();
	147	+ Cast(xFloat, xLocal, AscendC::RoundMode::CAST_NONE, static_cast<int32_t>(count));
	148	+ ComputeFloat(xFloat, yFloat, count);
	149	+ Cast(yLocal, yFloat, AscendC::RoundMode::CAST_RINT, static_cast<int32_t>(count));
	150	+ } else {
	151	+ // T == float：直接复用输入/输出缓冲，无需额外 UB 空间
	152	+ ComputeFloat(xLocal, yLocal, count);
	153	+ }
	154	+
	155	+ outputQueueY.EnQue(yLocal);
	156	+ inputQueueX.FreeTensor(xLocal);
	157	+}
	158	+
	159	+template <typename T>
	160	+__aicore__ inline void Leakyrelu<T>::CopyOut(int64_t progress, int64_t currentNum)
	161	+{
	162	+ LocalTensor<T> yLocal = outputQueueY.DeQue<T>();
	163	+ const int64_t offset = progress * ubLength_;
	164	+ // 只写回有效元素，避免越界污染输出张量尾部
	165	+ DataCopy(outputGMY[offset], yLocal, static_cast<int32_t>(currentNum));
	166	+ outputQueueY.FreeTensor(yLocal);
	atomgit-bot22 hours ago 🟡 Medium Priority 变更行：leakyrelu.h:165 `DataCopy(outputGMY[offset], yLocal, static_cast<int32_t>(currentNum));`。证据链：同一文件的 CopyIn（leakyrelu.h:117）刻意把搬入长度 AlignUp 到 alignElems_ 以满足 AscendC DataCopy 的 32B（datablock）长度对齐要求，而 CopyOut 却直接用未对齐的 `currentNum`。`currentNum` 来自 Process（第 177-178 行）`remain < ubLength_ ? remain : ubLength_`；host tiling 未对齐 totalNum（op_host/leakyrelu_tiling.cpp:96-105），所以末核最后一块的 `currentNum` 常常不是 alignElems_（float 8 / half 16）的整数倍。失败模式：DataCopy 的搬运长度按 datablock 粒度处理，长度非 32B 整数倍时行为不符合接口约定——按 32B 向上取整会把 UB 中尾部未初始化/无意义数据写到 GM，对最后一个核就是写到输出张量末尾之外（破坏相邻输出或内存池数据，验证脚本按精确字节数比对时表现为越界污染）；若硬件严格校验则直接报 AIC Error。建议：尾块写回使用 DataCopyPad（按 currentNum*sizeof(T) 的真实字节长度），或统一按对齐长度写回并在 host 侧保证输出张量有对齐余量，避免向 GM 写入非 32B 对齐长度。 Incorrect?
	167	+}
	168	+
	169	+template <typename T>
	170	+__aicore__ inline void Leakyrelu<T>::Process()
	171	+{
	172	+ if (blockLength_ <= 0) {
	173	+ return;
	174	+ }
	175	+ const int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
	176	+ for (int64_t progress = 0; progress < loopCount; ++progress) {
	177	+ const int64_t remain = blockLength_ - progress * ubLength_;
	178	+ const int64_t currentNum = (remain < ubLength_) ? remain : ubLength_;
	179	+ CopyIn(progress, currentNum);
	180	+ Compute(currentNum);
	181	+ CopyOut(progress, currentNum);
	182	+ }
	183	+}
	184	+
	185	+} // namespace NsLeakyrelu
	186	+#endif // LEAKYRELU_H

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_kernel/leakyrelu_tiling_data.h+27-0

@@ -0,0 +1,27 @@
	1	+/*!
	2	+ * \file leakyrelu_tiling_data.h
	3	+ * \brief tiling data struct
	4	+ */
	5	+
	6	+#ifndef _LEAKYRELU_TILING_DATA_H_
	7	+#define _LEAKYRELU_TILING_DATA_H_
	8	+
	9	+#include <cstdint>
	10	+
	11	+// negative_slope 的默认值（与算子属性缺省值、样例取值保持一致）。
	12	+#define LEAKYRELU_DEFAULT_SLOPE 0.01f
	13	+
	14	+// Host Tiling 写入的标记：表示 negativeSlope 字段是真实属性值。
	15	+// 直接构造 TilingData 的 CPU kernel UT fixture 不会写该标记，
	16	+// 此时 kernel 会退回 LEAKYRELU_DEFAULT_SLOPE，保证旧 fixture 仍可正确计算。
	17	+#define LEAKYRELU_SLOPE_TAG 0x4C524C554C524C55LL // "LRLULRLU"
	18	+
	19	+struct LeakyreluTilingData {
	20	+ int64_t totalNum = 0; // 总元素数量
	21	+ int64_t blockFactor = 1; // 每个核处理的元素数量（32B 对齐）
	22	+ int64_t ubFactor = 0; // 每次 UB 循环处理的元素数量（32B 对齐）
	23	+ int64_t slopeTag = 0; // 属性有效标记，见 LEAKYRELU_SLOPE_TAG
	24	+ float negativeSlope = LEAKYRELU_DEFAULT_SLOPE; // negative_slope 属性，随 tiling 下发到 kernel
	25	+ float reserved = 0.0f; // 显式占位，保证整个结构体按 8 字节整齐对齐
	26	+};
	27	+#endif

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/op_kernel/leakyrelu_tiling_key.h+21-0

@@ -0,0 +1,21 @@
	1	+/*!
	2	+ * \file leakyrelu_tiling_key.h
	3	+ * \brief Tiling 模板参数定义
	4	+ */
	5	+
	6	+#ifndef __LEAKYRELU_TILING_KEY_H__
	7	+#define __LEAKYRELU_TILING_KEY_H__
	8	+
	9	+#include "ascendc/host_api/tiling/template_argument.h"
	10	+
	11	+#define LEAKYRELU_TPL_SCH_MODE_0 0
	12	+#define LEAKYRELU_TPL_SCH_MODE_1 1
	13	+
	14	+ASCENDC_TPL_ARGS_DECL(
	15	+ Leakyrelu,
	16	+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, LEAKYRELU_TPL_SCH_MODE_0, LEAKYRELU_TPL_SCH_MODE_1));
	17	+
	18	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
	19	+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, LEAKYRELU_TPL_SCH_MODE_0, LEAKYRELU_TPL_SCH_MODE_1)));
	20	+
	21	+#endif

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/samples/case.json+53-0

@@ -0,0 +1,53 @@
	1	+{
	2	+ "source_case": "npu_cases[0]",
	3	+ "function": "impl",
	4	+ "attributes": {
	5	+ "negative_slope": 0.01
	6	+ },
	7	+ "inputs": [
	8	+ {
	9	+ "name": "x",
	10	+ "kind": "required",
	11	+ "tensors": [
	12	+ {
	13	+ "file": "x.bin",
	14	+ "dtype": "float32",
	15	+ "shape": [
	16	+ 8,
	17	+ 2048
	18	+ ],
	19	+ "bytes": 65536,
	20	+ "elements": 16384,
	21	+ "range": [
	22	+ -10,
	23	+ 10
	24	+ ],
	25	+ "value": null,
	26	+ "value_url": null,
	27	+ "value_sha256": null
	28	+ }
	29	+ ]
	30	+ }
	31	+ ],
	32	+ "outputs": [
	33	+ {
	34	+ "name": "y",
	35	+ "kind": "required",
	36	+ "tensors": [
	37	+ {
	38	+ "file": "y.bin",
	39	+ "dtype": "float32",
	40	+ "shape": [
	41	+ 8,
	42	+ 2048
	43	+ ],
	44	+ "bytes": 65536,
	45	+ "elements": 16384,
	46	+ "rtol": 0.0001,
	47	+ "atol": 0.0001,
	48	+ "tol": 0.0001
	49	+ }
	50	+ ]
	51	+ }
	52	+ ]
	53	+}

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/samples/reference.py+20-0

@@ -0,0 +1,20 @@
	1	+import numpy as np
	2	+
	3	+def impl(x, negative_slope=0.01):
	4	+ # 验算函数：逐元素 leaky relu，y = max(x, 0) + negative_slope * min(x, 0)
	5	+ # float64 中间计算，最后一次性转回输入 dtype
	6	+ dtype = x.dtype
	7	+ x64 = x.astype(np.float64)
	8	+ y = np.where(x64 > 0, x64, np.float64(negative_slope) * x64)
	9	+ return y.astype(dtype)
	10	+
	11	+if __name__ == "__main__":
	12	+ for dtype in (np.float16, np.float32):
	13	+ x = np.random.uniform(-10, 10, (8, 2)).astype(dtype)
	14	+ y = impl(x, 0.01)
	15	+ x64 = x.astype(np.float64)
	16	+ ref = np.where(x64 > 0, x64, 0.01 * x64).astype(dtype)
	17	+ ok = np.array_equal(y, ref)
	18	+ print("dtype={}, shape={}, output dtype={}, match={}".format(
	19	+ np.dtype(dtype).name, x.shape, y.dtype, ok))
	20	+

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/samples/sample.py+87-0

@@ -0,0 +1,87 @@
	1	+import argparse
	2	+from pathlib import Path
	3	+
	4	+import numpy as np
	5	+import reference
	6	+
	7	+
	8	+ROOT = Path(__file__).resolve().parent
	9	+
	10	+OUTPUTS = [
	11	+ ('y.bin', [8, 2048], np.dtype('float32'), 65536, 0.0001, 0.0001, 0.0001),
	12	+]
	13	+
	14	+
	15	+def generate(directory):
	16	+ rng = np.random.default_rng(42)
	17	+ for name in ('input', 'output', 'golden'):
	18	+ (directory / name).mkdir(parents=True, exist_ok=True)
	19	+ for item in OUTPUTS:
	20	+ (directory / 'output' / item[0]).unlink(missing_ok=True)
	21	+ (directory / 'golden' / item[0]).unlink(missing_ok=True)
	22	+ x0 = rng.uniform(-10, 10, [8, 2048]).astype(np.dtype('float32'))
	23	+ (directory / 'input' / 'x.bin').write_bytes(x0.tobytes())
	24	+ result = reference.impl(x0, **{'negative_slope': 0.01})
	25	+ result = (result,)
	26	+ values = [result[0]]
	27	+ for item, array in zip(OUTPUTS, values):
	28	+ name, shape, dtype = item[:3]
	29	+ if np.isscalar(array):
	30	+ array = np.asarray(array).reshape(shape)
	31	+ array = np.asarray(array)
	32	+ if array.dtype != dtype or list(array.shape) != shape:
	33	+ raise ValueError('reference dtype/shape mismatch: ' + name)
	34	+ (directory / 'golden' / name).write_bytes(array.tobytes())
	35	+ print('Sample data generated.')
	36	+
	37	+
	38	+def verify(directory):
	39	+ passed = True
	40	+ expected = {item[0] for item in OUTPUTS}
	41	+ for folder in ('output', 'golden'):
	42	+ actual = {p.name for p in (directory / folder).iterdir()}
	43	+ if actual != expected:
	44	+ print('VERIFY ERROR: unexpected or missing files in ' + folder)
	45	+ return 1
	46	+ for item in OUTPUTS:
	47	+ name, shape, dtype, size, rtol, atol, tol = item[:7]
	48	+ try:
	49	+ raw = (directory / 'output' / name).read_bytes()
	50	+ golden = (directory / 'golden' / name).read_bytes()
	51	+ if len(raw) != size or len(golden) != size:
	52	+ raise ValueError('exact byte size mismatch')
	53	+ a, b = [np.frombuffer(v, dtype=dtype) for v in (raw, golden)]
	54	+ close = np.isclose(
	55	+ a.astype(np.float64),
	56	+ b.astype(np.float64),
	57	+ rtol=rtol,
	58	+ atol=atol,
	59	+ equal_nan=True
	60	+ )
	61	+ bad = int(np.count_nonzero(~close))
	62	+ ok = bad <= tol * a.size
	63	+ print(
	64	+ ('PASS' if ok else 'Wrong Answer')
	65	+ + ': ' + name + f', mismatched={bad}/{a.size}'
	66	+ )
	67	+ passed &= ok
	68	+ except (OSError, ValueError) as error:
	69	+ print('VERIFY ERROR: ' + name + ': ' + str(error))
	70	+ passed = False
	71	+ print('PASSED' if passed else 'FAILED')
	72	+ return 0 if passed else 1
	73	+
	74	+
	75	+if __name__ == '__main__':
	76	+ parser = argparse.ArgumentParser()
	77	+ parser.add_argument('action', choices=('generate', 'verify'))
	78	+ parser.add_argument('--directory', type=Path, default=ROOT / 'build-npu')
	79	+ args = parser.parse_args()
	80	+ try:
	81	+ if args.action == 'generate':
	82	+ generate(args.directory)
	83	+ else:
	84	+ raise SystemExit(verify(args.directory))
	85	+ except Exception as error:
	86	+ print('SAMPLE ERROR: ' + str(error))
	87	+ raise SystemExit(2)

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/CMakeLists.txt+68-0

@@ -0,0 +1,68 @@
	1	+cmake_minimum_required(VERSION 3.16.0)
	2	+project(leakyrelu_ut CXX C)
	3	+
	4	+set(CMAKE_CXX_STANDARD 17)
	5	+set(CMAKE_CXX_STANDARD_REQUIRED ON)
	6	+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
	7	+
	8	+if(NOT DEFINED ENV{ASCEND_HOME_PATH})
	9	+ message(FATAL_ERROR "ASCEND_HOME_PATH environment variable is not set!")
	10	+endif()
	11	+
	12	+message(STATUS "ASCEND_HOME_PATH: $ENV{ASCEND_HOME_PATH}")
	13	+
	14	+include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/BuildGoogleTest.cmake)
	15	+
	16	+# DTYPE 宏定义（根据算子原型和测试用例生成）
	17	+add_definitions(
	18	+ -DDTYPE_X=float
	19	+ -DDTYPE_Y=float
	20	+)
	21	+
	22	+set(UT_COMMON_DIR ${CMAKE_CURRENT_SOURCE_DIR}/common)
	23	+set(UT_COMMON_SRCS
	24	+ ${UT_COMMON_DIR}/infershape_case_executor.cpp
	25	+ ${UT_COMMON_DIR}/infershape_context_faker.cpp
	26	+ ${UT_COMMON_DIR}/tiling_case_executor.cpp
	27	+ ${UT_COMMON_DIR}/tiling_context_faker.cpp
	28	+)
	29	+
	30	+set(UT_COMMON_INCLUDE_DIRS
	31	+ ${UT_COMMON_DIR}
	32	+ $ENV{ASCEND_HOME_PATH}/include
	33	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/include
	34	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/include/base
	35	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/include/base/context_builder
	36	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include
	37	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include/tiling
	38	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc
	39	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/op_common
	40	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/base
	41	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/exe_graph
	42	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/pkg_inc/graph
	43	+)
	44	+
	45	+set(UT_COMMON_LIB_DIRS
	46	+ $ENV{ASCEND_HOME_PATH}/lib64
	47	+ $ENV{ASCEND_HOME_PATH}/${CMAKE_SYSTEM_PROCESSOR}-linux/lib64
	48	+)
	49	+
	50	+add_compile_options(
	51	+ -Wall
	52	+ -Wno-deprecated-declarations
	53	+ -Wno-unused-variable
	54	+ -fno-access-control
	55	+)
	56	+
	57	+if(NOT CMAKE_BUILD_TYPE)
	58	+ set(CMAKE_BUILD_TYPE Debug)
	59	+endif()
	60	+
	61	+message(STATUS "Build Type: ${CMAKE_BUILD_TYPE}")
	62	+
	63	+add_subdirectory(op_host)
	64	+
	65	+if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/../../op_api)
	66	+ add_subdirectory(op_api)
	67	+endif()
	68	+add_subdirectory(op_kernel)

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/cmake/BuildGoogleTest.cmake+11-0

@@ -0,0 +1,11 @@
	1	+include_guard(GLOBAL)
	2	+find_path(NLOHMANN_JSON_INCLUDE_DIR nlohmann/json.hpp HINTS "$ENV{NLOHMANN_JSON_INCLUDE_DIR}" REQUIRED)
	3	+include_directories(${NLOHMANN_JSON_INCLUDE_DIR})
	4	+set(GTEST_SOURCE_DIR "$ENV{GTEST_SOURCE_DIR}" CACHE PATH "GoogleTest source directory")
	5	+if(NOT EXISTS "${GTEST_SOURCE_DIR}/CMakeLists.txt")
	6	+ message(FATAL_ERROR "Set GTEST_SOURCE_DIR to local GoogleTest sources (1.14.0 recommended)")
	7	+endif()
	8	+add_compile_definitions(_GLIBCXX_USE_CXX11_ABI=0)
	9	+add_subdirectory(${GTEST_SOURCE_DIR} ${CMAKE_BINARY_DIR}/gtest EXCLUDE_FROM_ALL)
	10	+add_library(intf_llt_pub_asan_cxx17 INTERFACE)
	11	+target_link_libraries(intf_llt_pub_asan_cxx17 INTERFACE gtest)

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/any_value.h+94-0

@@ -0,0 +1,94 @@
	1	+#ifndef OPS_MATH_DEV_TESTS_UT_COMMON_ANY_VALUE_H
	2	+#define OPS_MATH_DEV_TESTS_UT_COMMON_ANY_VALUE_H
	3	+
	4	+#include <memory>
	5	+#include <cstdint>
	6	+#include <string>
	7	+#include <vector>
	8	+
	9	+namespace Ops {
	10	+namespace Math {
	11	+class AnyValue {
	12	+ public:
	13	+ enum ValueType {
	14	+ VT_STRING = 1,
	15	+ VT_FLOAT = 2,
	16	+ VT_BOOL = 3,
	17	+ VT_INT = 4,
	18	+ VT_LIST_LIST_INT = 10,
	19	+ VT_LIST_BASE = 1000,
	20	+
	21	+ VT_LIST_FLOAT = static_cast<int32_t>(VT_LIST_BASE) + static_cast<int32_t>(VT_FLOAT),
	22	+ VT_LIST_BOOL = static_cast<int32_t>(VT_LIST_BASE) + static_cast<int32_t>(VT_BOOL),
	23	+ VT_LIST_INT = static_cast<int32_t>(VT_LIST_BASE) + static_cast<int32_t>(VT_INT),
	24	+ };
	25	+
	26	+ AnyValue(ValueType type, const std::shared_ptr<void>& valuePtr)
	27	+ : type_(type), valuePtr_(valuePtr) {
	28	+ }
	29	+ ~AnyValue() = default;
	30	+ AnyValue(const AnyValue& anyValue) : type_(anyValue.type_), valuePtr_(anyValue.valuePtr_) {
	31	+ }
	32	+
	33	+ template <typename T> static inline AnyValue CreateFrom(const T& value);
	34	+
	35	+ ValueType type_;
	36	+ std::shared_ptr<void> valuePtr_;
	37	+};
	38	+
	39	+template <> inline AnyValue AnyValue::CreateFrom<std::string>(const std::string& value) {
	40	+ auto valuePtr = new std::string;
	41	+ *valuePtr = value;
	42	+ return AnyValue(VT_STRING, std::shared_ptr<void>(valuePtr));
	43	+}
	44	+
	45	+template <> inline AnyValue AnyValue::CreateFrom<float>(const float& value) {
	46	+ auto valuePtr = new float;
	47	+ *valuePtr = value;
	48	+ return AnyValue(VT_FLOAT, std::shared_ptr<void>(valuePtr));
	49	+}
	50	+
	51	+template <> inline AnyValue AnyValue::CreateFrom<bool>(const bool& value) {
	52	+ auto valuePtr = new bool;
	53	+ *valuePtr = value;
	54	+ return AnyValue(VT_BOOL, std::shared_ptr<void>(valuePtr));
	55	+}
	56	+
	57	+template <> inline AnyValue AnyValue::CreateFrom<int64_t>(const int64_t& value) {
	58	+ auto valuePtr = new int64_t;
	59	+ *valuePtr = value;
	60	+ return AnyValue(VT_INT, std::shared_ptr<void>(valuePtr));
	61	+}
	62	+
	63	+template <>
	64	+inline AnyValue AnyValue::CreateFrom<std::vector<float>>(const std::vector<float>& value) {
	65	+ auto valuePtr = new std::vector<float>;
	66	+ *valuePtr = value;
	67	+ return AnyValue(VT_LIST_FLOAT, std::shared_ptr<void>(valuePtr));
	68	+}
	69	+
	70	+template <>
	71	+inline AnyValue AnyValue::CreateFrom<std::vector<bool>>(const std::vector<bool>& value) {
	72	+ auto valuePtr = new std::vector<bool>;
	73	+ *valuePtr = value;
	74	+ return AnyValue(VT_LIST_BOOL, std::shared_ptr<void>(valuePtr));
	75	+}
	76	+
	77	+template <>
	78	+inline AnyValue AnyValue::CreateFrom<std::vector<int64_t>>(const std::vector<int64_t>& value) {
	79	+ auto valuePtr = new std::vector<int64_t>;
	80	+ *valuePtr = value;
	81	+ return AnyValue(VT_LIST_INT, std::shared_ptr<void>(valuePtr));
	82	+}
	83	+
	84	+template <>
	85	+inline AnyValue AnyValue::CreateFrom<std::vector<std::vector<int64_t>>>(
	86	+ const std::vector<std::vector<int64_t>>& value) {
	87	+ auto valuePtr = new std::vector<std::vector<int64_t>>;
	88	+ *valuePtr = value;
	89	+ return AnyValue(VT_LIST_LIST_INT, std::shared_ptr<void>(valuePtr));
	90	+}
	91	+} // namespace Math
	92	+} // namespace Ops
	93	+
	94	+#endif // OPS_MATH_DEV_TESTS_UT_COMMON_ANY_VALUE_H

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/infershape_case_executor.cpp+121-0

@@ -0,0 +1,121 @@
	1	+#include "infershape_case_executor.h"
	2	+#include <gtest/gtest.h>
	3	+#include "base/registry/op_impl_space_registry_v2.h"
	4	+
	5	+#define DO_INFERSHAPE(infershapeContextPara) \
	6	+ auto contextFaker = gert::InferShapeContextFaker(); \
	7	+ /* 1. input/output information */ \
	8	+ size_t inputNum = infershapeContextPara.inputTensorDesc_.size(); \
	9	+ size_t outputNum = infershapeContextPara.outputTensorDesc_.size(); \
	10	+ if (infershapeContextPara.inputInstanceNum_.size() != 0 \|\| \
	11	+ infershapeContextPara.outputInstanceNum_.size() != 0) { \
	12	+ contextFaker.IrInstanceNum(infershapeContextPara.inputInstanceNum_, \
	13	+ infershapeContextPara.outputInstanceNum_); \
	14	+ } else { \
	15	+ contextFaker.NodeIoNum(inputNum, outputNum); \
	16	+ } \
	17	+ std::vector<gert::Tensor*> inputTensors = {}; \
	18	+ std::vector<std::unique_ptr<gert::Tensor>> inputTensorsKeepAlive = {}; \
	19	+ for (size_t index = 0; index < inputNum; index++) { \
	20	+ std::unique_ptr<gert::Tensor> curTensor = std::make_unique<gert::Tensor>( \
	21	+ infershapeContextPara.inputTensorDesc_[index].shape_, \
	22	+ gert::StorageFormat(infershapeContextPara.inputTensorDesc_[index].format_, \
	23	+ infershapeContextPara.inputTensorDesc_[index].format_, \
	24	+ gert::ExpandDimsType()), \
	25	+ gert::TensorPlacement::kOnHost, infershapeContextPara.inputTensorDesc_[index].dtype_, \
	26	+ infershapeContextPara.inputTensorDesc_[index].isConst_ \
	27	+ ? infershapeContextPara.inputTensorDesc_[index].constValue_ \
	28	+ : nullptr); \
	29	+ inputTensors.push_back(curTensor.get()); \
	30	+ inputTensorsKeepAlive.push_back(std::move(curTensor)); \
	31	+ } \
	32	+ for (size_t index = 0; index < outputNum; index++) { \
	33	+ contextFaker.NodeOutputTd(index, infershapeContextPara.outputTensorDesc_[index].dtype_, \
	34	+ infershapeContextPara.outputTensorDesc_[index].format_, \
	35	+ infershapeContextPara.outputTensorDesc_[index].format_); \
	36	+ } \
	37	+ contextFaker.InputTensors(inputTensors); \
	38	+ for (auto& attrInfo : infershapeContextPara.attrs_) { \
	39	+ switch (attrInfo.attr_.type_) { \
	40	+ case Ops::Math::AnyValue::ValueType::VT_BOOL: { \
	41	+ contextFaker.Attr(attrInfo.attrName_, \
	42	+ reinterpret_cast<bool>(attrInfo.attr_.valuePtr_.get())); \
	43	+ break; \
	44	+ } \
	45	+ case Ops::Math::AnyValue::ValueType::VT_INT: { \
	46	+ contextFaker.Attr(attrInfo.attrName_, \
	47	+ reinterpret_cast<int64_t>(attrInfo.attr_.valuePtr_.get())); \
	48	+ break; \
	49	+ } \
	50	+ case Ops::Math::AnyValue::ValueType::VT_FLOAT: { \
	51	+ contextFaker.Attr(attrInfo.attrName_, \
	52	+ reinterpret_cast<float>(attrInfo.attr_.valuePtr_.get())); \
	53	+ break; \
	54	+ } \
	55	+ case Ops::Math::AnyValue::ValueType::VT_STRING: { \
	56	+ contextFaker.Attr( \
	57	+ attrInfo.attrName_, \
	58	+ ge::AscendString( \
	59	+ reinterpret_cast<std::string*>(attrInfo.attr_.valuePtr_.get())->c_str())); \
	60	+ break; \
	61	+ } \
	62	+ case Ops::Math::AnyValue::ValueType::VT_LIST_BOOL: { \
	63	+ contextFaker.Attr(attrInfo.attrName_, reinterpret_cast<std::vector<bool>>( \
	64	+ attrInfo.attr_.valuePtr_.get())); \
	65	+ break; \
	66	+ } \
	67	+ case Ops::Math::AnyValue::ValueType::VT_LIST_INT: { \
	68	+ contextFaker.Attr(attrInfo.attrName_, reinterpret_cast<std::vector<int64_t>>( \
	69	+ attrInfo.attr_.valuePtr_.get())); \
	70	+ break; \
	71	+ } \
	72	+ case Ops::Math::AnyValue::ValueType::VT_LIST_LIST_INT: { \
	73	+ contextFaker.Attr(attrInfo.attrName_, \
	74	+ reinterpret_cast<std::vector<std::vector<int64_t>>>( \
	75	+ attrInfo.attr_.valuePtr_.get())); \
	76	+ break; \
	77	+ } \
	78	+ case Ops::Math::AnyValue::ValueType::VT_LIST_FLOAT: { \
	79	+ contextFaker.Attr(attrInfo.attrName_, reinterpret_cast<std::vector<float>>( \
	80	+ attrInfo.attr_.valuePtr_.get())); \
	81	+ break; \
	82	+ } \
	83	+ default: \
	84	+ std::cout << "[ERROR]" << __FILE__ << ":" << __LINE__ << "The ValueType " \
	85	+ << attrInfo.attr_.type_ << "is not supported!" << std::endl; \
	86	+ } \
	87	+ } \
	88	+ auto contextHolder = contextFaker.SetOpType(infershapeContextPara.opName_.c_str()).Build(); \
	89	+ /* 2. get infershape func */ \
	90	+ auto spaceRegistry = gert::DefaultOpImplSpaceRegistryV2::GetInstance().GetSpaceRegistry(); \
	91	+ auto infershapeFunc = \
	92	+ spaceRegistry->GetOpImpl(infershapeContextPara.opName_.c_str())->infer_shape; \
	93	+ /* 3. check infershape func */ \
	94	+ auto infershapeRet = infershapeFunc(contextHolder.GetContext());
	95	+
	96	+static std::vector<int64_t> ToVector(const gert::Shape& shape) {
	97	+ size_t shapeSize = shape.GetDimNum();
	98	+ std::vector<int64_t> shapeVec(shapeSize, 0);
	99	+
	100	+ for (size_t i = 0; i < shapeSize; i++) {
	101	+ shapeVec[i] = shape.GetDim(i);
	102	+ }
	103	+ return shapeVec;
	104	+}
	105	+
	106	+void ExecuteTestCase(gert::InfershapeContextPara& infershapeContextPara,
	107	+ ge::graphStatus expectResult,
	108	+ const std::vector<std::vector<int64_t>>& expectOutputShape) {
	109	+ DO_INFERSHAPE(infershapeContextPara);
	110	+
	111	+ // check infershape func
	112	+ EXPECT_EQ(infershapeRet, expectResult);
	113	+ if (expectResult == ge::GRAPH_FAILED) {
	114	+ return;
	115	+ }
	116	+
	117	+ // check output shape
	118	+ for (size_t i = 0; i < expectOutputShape.size(); i++) {
	119	+ EXPECT_EQ(ToVector(*contextHolder.GetContext()->GetOutputShape(i)), expectOutputShape[i]);
	120	+ }
	121	+}

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/infershape_case_executor.h+10-0

@@ -0,0 +1,10 @@
	1	+#ifndef OPS_MATH_DEV_TESTS_UT_COMMON_INFERSHAPE_CASE_EXECUTOR_H
	2	+#define OPS_MATH_DEV_TESTS_UT_COMMON_INFERSHAPE_CASE_EXECUTOR_H
	3	+
	4	+#include "infershape_context_faker.h"
	5	+
	6	+void ExecuteTestCase(gert::InfershapeContextPara& infershapeContextPara,
	7	+ ge::graphStatus expectResult = ge::GRAPH_FAILED,
	8	+ const std::vector<std::vector<int64_t>>& expectOutputShape = {});
	9	+
	10	+#endif // OPS_MATH_DEV_TESTS_UT_COMMON_INFERSHAPE_CASE_EXECUTOR_H

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/infershape_context_faker.cpp+39-0

@@ -0,0 +1,39 @@
	1	+#include "infershape_context_faker.h"
	2	+
	3	+namespace gert {
	4	+
	5	+InferShapeContextFaker& InferShapeContextFaker::SetOpType(const std::string opType) {
	6	+ OpInferShapeContextBuilder::OpType(opType.c_str()).OpName(opType.c_str());
	7	+ return *this;
	8	+}
	9	+
	10	+InferShapeContextFaker& InferShapeContextFaker::NodeIoNum(size_t inputNum, size_t outputNum) {
	11	+ OpInferShapeContextBuilder::IONum(inputNum, outputNum);
	12	+ return *this;
	13	+}
	14	+
	15	+InferShapeContextFaker&
	16	+InferShapeContextFaker::IrInstanceNum(const std::vector<uint32_t>& inputInstanceNum,
	17	+ const std::vector<uint32_t>& outputInstanceNum) {
	18	+ OpInferShapeContextBuilder::IOInstanceNum(inputInstanceNum, outputInstanceNum);
	19	+ return *this;
	20	+}
	21	+
	22	+InferShapeContextFaker& InferShapeContextFaker::NodeOutputTd(int32_t index, ge::DataType dtype,
	23	+ ge::Format originFormat,
	24	+ ge::Format storageFormat) {
	25	+ OpInferShapeContextBuilder::OutputTensorDesc(index, dtype, originFormat, storageFormat);
	26	+ return *this;
	27	+}
	28	+
	29	+InferShapeContextFaker&
	30	+InferShapeContextFaker::InputTensors(const std::vector<Tensor*>& inputTensors) {
	31	+ OpInferShapeContextBuilder::InputTensors(inputTensors);
	32	+ return *this;
	33	+}
	34	+
	35	+ContextHolder<InferShapeContext> InferShapeContextFaker::Build() {
	36	+ return OpInferShapeContextBuilder::Build();
	37	+}
	38	+
	39	+} // namespace gert

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/infershape_context_faker.h+129-0

@@ -0,0 +1,129 @@
	1	+#ifndef OPS_MATH_DEV_TESTS_UT_COMMON_INFERSHAPE_CONTEXT_FAKER_H
	2	+#define OPS_MATH_DEV_TESTS_UT_COMMON_INFERSHAPE_CONTEXT_FAKER_H
	3	+
	4	+#include "op_infer_shape_context_builder.h"
	5	+#include "any_value.h"
	6	+
	7	+namespace gert {
	8	+
	9	+class InfershapeContextPara {
	10	+ public:
	11	+ class TensorDescription {
	12	+ public:
	13	+ TensorDescription(const gert::StorageShape& shape, ge::DataType dtype, ge::Format format,
	14	+ bool isConst = false, void* constValue = nullptr)
	15	+ : shape_(shape), dtype_(dtype), format_(format), isConst_(isConst),
	16	+ constValue_(constValue) {
	17	+ }
	18	+
	19	+ public:
	20	+ gert::StorageShape shape_;
	21	+ ge::DataType dtype_ = ge::DT_FLOAT;
	22	+ ge::Format format_ = ge::FORMAT_ND;
	23	+ bool isConst_ = false;
	24	+ void* constValue_ = nullptr;
	25	+ };
	26	+
	27	+ class OpAttr {
	28	+ public:
	29	+ OpAttr(const std::string& attrName, const Ops::Math::AnyValue& attr)
	30	+ : attrName_(attrName), attr_(attr) {
	31	+ }
	32	+
	33	+ public:
	34	+ std::string attrName_;
	35	+ Ops::Math::AnyValue attr_;
	36	+ };
	37	+
	38	+ public:
	39	+ InfershapeContextPara(const std::string& opName,
	40	+ const std::vector<TensorDescription>& inputTensorDesc,
	41	+ const std::vector<TensorDescription>& outputTensorDesc,
	42	+ const std::vector<OpAttr>& attrs,
	43	+ const std::vector<uint32_t>& inputInstanceNum = {},
	44	+ const std::vector<uint32_t>& outputInstanceNum = {})
	45	+ : opName_(opName), inputInstanceNum_(inputInstanceNum),
	46	+ outputInstanceNum_(outputInstanceNum), inputTensorDesc_(inputTensorDesc),
	47	+ outputTensorDesc_(outputTensorDesc), attrs_(attrs) {
	48	+ }
	49	+
	50	+ InfershapeContextPara(const std::string& opName,
	51	+ const std::vector<TensorDescription>& inputTensorDesc,
	52	+ const std::vector<TensorDescription>& outputTensorDesc,
	53	+ const std::vector<uint32_t>& inputInstanceNum = {},
	54	+ const std::vector<uint32_t>& outputInstanceNum = {})
	55	+ : opName_(opName), inputInstanceNum_(inputInstanceNum),
	56	+ outputInstanceNum_(outputInstanceNum), inputTensorDesc_(inputTensorDesc),
	57	+ outputTensorDesc_(outputTensorDesc), attrs_() {
	58	+ }
	59	+
	60	+ public:
	61	+ std::string opName_;
	62	+ std::vector<uint32_t> inputInstanceNum_;
	63	+ std::vector<uint32_t> outputInstanceNum_;
	64	+ std::vector<TensorDescription> inputTensorDesc_;
	65	+ std::vector<TensorDescription> outputTensorDesc_;
	66	+ std::vector<OpAttr> attrs_;
	67	+};
	68	+
	69	+class InferShapeContextFaker : public OpInferShapeContextBuilder {
	70	+ public:
	71	+ InferShapeContextFaker& SetOpType(const std::string opType);
	72	+
	73	+ /* only one can be choosed from IrInstanceNum */
	74	+ InferShapeContextFaker& NodeIoNum(size_t inputNum, size_t outputNum);
	75	+
	76	+ /* can be used for dynamic inputs/outputs
	77	+ * only one can be choosed from NodeIoNum */
	78	+ InferShapeContextFaker& IrInstanceNum(const std::vector<uint32_t>& inputInstanceNum,
	79	+ const std::vector<uint32_t>& outputInstanceNum);
	80	+
	81	+ InferShapeContextFaker& NodeOutputTd(int32_t index, ge::DataType dtype, ge::Format originFormat,
	82	+ ge::Format storageFormat);
	83	+
	84	+ InferShapeContextFaker& Attr(const std::string& attrName, bool attr) {
	85	+ OpInferShapeContextBuilder::AppendAttr(attr);
	86	+ return *this;
	87	+ }
	88	+ InferShapeContextFaker& Attr(const std::string& attrName, int64_t attr) {
	89	+ OpInferShapeContextBuilder::AppendAttr(attr);
	90	+ return *this;
	91	+ }
	92	+ InferShapeContextFaker& Attr(const std::string& attrName, float attr) {
	93	+ OpInferShapeContextBuilder::AppendAttr(attr);
	94	+ return *this;
	95	+ }
	96	+ InferShapeContextFaker& Attr(const std::string& attrName, const ge::AscendString& attr) {
	97	+ OpInferShapeContextBuilder::AppendAttr(attr);
	98	+ return *this;
	99	+ }
	100	+ InferShapeContextFaker& Attr(const std::string& attrName, const std::vector<bool>& attr) {
	101	+ OpInferShapeContextBuilder::AppendAttr(attr);
	102	+ return *this;
	103	+ }
	104	+ InferShapeContextFaker& Attr(const std::string& attrName, const std::vector<int64_t>& attr) {
	105	+ OpInferShapeContextBuilder::AppendAttr(attr);
	106	+ return *this;
	107	+ }
	108	+ InferShapeContextFaker& Attr(const std::string& attrName, const std::vector<float>& attr) {
	109	+ OpInferShapeContextBuilder::AppendAttr(attr);
	110	+ return *this;
	111	+ }
	112	+ InferShapeContextFaker& Attr(const std::string& attrName,
	113	+ const std::vector<ge::AscendString>& attr) {
	114	+ OpInferShapeContextBuilder::AppendAttr(attr);
	115	+ return *this;
	116	+ }
	117	+ InferShapeContextFaker& Attr(const std::string& attrName,
	118	+ const std::vector<std::vector<int64_t>>& attr) {
	119	+ OpInferShapeContextBuilder::AppendAttr(attr);
	120	+ return *this;
	121	+ }
	122	+
	123	+ InferShapeContextFaker& InputTensors(const std::vector<Tensor*>& inputTensors);
	124	+
	125	+ ContextHolder<InferShapeContext> Build();
	126	+};
	127	+
	128	+} // namespace gert
	129	+#endif // OPS_MATH_DEV_TESTS_UT_COMMON_INFERSHAPE_CONTEXT_FAKER_H

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/tiling_case_executor.cpp+310-0

@@ -0,0 +1,310 @@
	1	+#include "tiling_case_executor.h"
	2	+#include <gtest/gtest.h>
	3	+#include <nlohmann/json.hpp>
	4	+#include "platform/platform_infos_def.h"
	5	+#include "base/registry/op_impl_space_registry_v2.h"
	6	+
	7	+#define STR_IMPL(x) #x
	8	+#define STR(x) STR_IMPL(x)
	9	+#define DO_TILING(tilingContextPara) \
	10	+ auto contextFaker = gert::TilingContextFaker(); \
	11	+ /* 1. input/output information */ \
	12	+ size_t inputNum = tilingContextPara.inputTensorDesc_.size(); \
	13	+ size_t outputNum = tilingContextPara.outputTensorDesc_.size(); \
	14	+ if (tilingContextPara.inputInstanceNum_.size() != 0 \|\| \
	15	+ tilingContextPara.outputInstanceNum_.size() != 0) { \
	16	+ contextFaker.IrInstanceNum(tilingContextPara.inputInstanceNum_, \
	17	+ tilingContextPara.outputInstanceNum_); \
	18	+ } else { \
	19	+ contextFaker.NodeIoNum(inputNum, outputNum); \
	20	+ } \
	21	+ std::vector<gert::Tensor*> inputTensors = {}; \
	22	+ std::vector<gert::Tensor*> outputTensors = {}; \
	23	+ std::vector<std::unique_ptr<gert::Tensor>> inputTensorsKeepAlive = {}; \
	24	+ std::vector<std::unique_ptr<gert::Tensor>> outputTensorsKeepAlive = {}; \
	25	+ for (size_t index = 0; index < inputNum; index++) { \
	26	+ std::unique_ptr<gert::Tensor> curTensor = std::make_unique<gert::Tensor>( \
	27	+ tilingContextPara.inputTensorDesc_[index].shape_, \
	28	+ gert::StorageFormat(tilingContextPara.inputTensorDesc_[index].format_, \
	29	+ tilingContextPara.inputTensorDesc_[index].format_, \
	30	+ gert::ExpandDimsType()), \
	31	+ gert::TensorPlacement::kOnHost, tilingContextPara.inputTensorDesc_[index].dtype_, \
	32	+ tilingContextPara.inputTensorDesc_[index].isConst_ \
	33	+ ? tilingContextPara.inputTensorDesc_[index].constValue_ \
	34	+ : nullptr); \
	35	+ inputTensors.push_back(curTensor.get()); \
	36	+ inputTensorsKeepAlive.push_back(std::move(curTensor)); \
	37	+ } \
	38	+ for (size_t index = 0; index < outputNum; index++) { \
	39	+ std::unique_ptr<gert::Tensor> curTensor = std::make_unique<gert::Tensor>( \
	40	+ tilingContextPara.outputTensorDesc_[index].shape_, \
	41	+ gert::StorageFormat(tilingContextPara.outputTensorDesc_[index].format_, \
	42	+ tilingContextPara.outputTensorDesc_[index].format_, \
	43	+ gert::ExpandDimsType()), \
	44	+ gert::TensorPlacement::kOnHost, tilingContextPara.outputTensorDesc_[index].dtype_, \
	45	+ tilingContextPara.outputTensorDesc_[index].isConst_ \
	46	+ ? tilingContextPara.outputTensorDesc_[index].constValue_ \
	47	+ : nullptr); \
	48	+ outputTensors.push_back(curTensor.get()); \
	49	+ outputTensorsKeepAlive.push_back(std::move(curTensor)); \
	50	+ } \
	51	+ contextFaker.InputTensors(inputTensors).OutputTensors(outputTensors); \
	52	+ for (auto& attrInfo : tilingContextPara.attrs_) { \
	53	+ switch (attrInfo.attr_.type_) { \
	54	+ case Ops::Math::AnyValue::ValueType::VT_BOOL: { \
	55	+ contextFaker.Attr(attrInfo.attrName_, \
	56	+ reinterpret_cast<bool>(attrInfo.attr_.valuePtr_.get())); \
	57	+ break; \
	58	+ } \
	59	+ case Ops::Math::AnyValue::ValueType::VT_INT: { \
	60	+ contextFaker.Attr(attrInfo.attrName_, \
	61	+ reinterpret_cast<int64_t>(attrInfo.attr_.valuePtr_.get())); \
	62	+ break; \
	63	+ } \
	64	+ case Ops::Math::AnyValue::ValueType::VT_FLOAT: { \
	65	+ contextFaker.Attr(attrInfo.attrName_, \
	66	+ reinterpret_cast<float>(attrInfo.attr_.valuePtr_.get())); \
	67	+ break; \
	68	+ } \
	69	+ case Ops::Math::AnyValue::ValueType::VT_STRING: { \
	70	+ contextFaker.Attr( \
	71	+ attrInfo.attrName_, \
	72	+ ge::AscendString( \
	73	+ reinterpret_cast<std::string*>(attrInfo.attr_.valuePtr_.get())->c_str())); \
	74	+ break; \
	75	+ } \
	76	+ case Ops::Math::AnyValue::ValueType::VT_LIST_BOOL: { \
	77	+ contextFaker.Attr(attrInfo.attrName_, reinterpret_cast<std::vector<bool>>( \
	78	+ attrInfo.attr_.valuePtr_.get())); \
	79	+ break; \
	80	+ } \
	81	+ case Ops::Math::AnyValue::ValueType::VT_LIST_INT: { \
	82	+ contextFaker.Attr(attrInfo.attrName_, reinterpret_cast<std::vector<int64_t>>( \
	83	+ attrInfo.attr_.valuePtr_.get())); \
	84	+ break; \
	85	+ } \
	86	+ case Ops::Math::AnyValue::ValueType::VT_LIST_LIST_INT: { \
	87	+ contextFaker.Attr(attrInfo.attrName_, \
	88	+ reinterpret_cast<std::vector<std::vector<int64_t>>>( \
	89	+ attrInfo.attr_.valuePtr_.get())); \
	90	+ break; \
	91	+ } \
	92	+ case Ops::Math::AnyValue::ValueType::VT_LIST_FLOAT: { \
	93	+ contextFaker.Attr(attrInfo.attrName_, reinterpret_cast<std::vector<float>>( \
	94	+ attrInfo.attr_.valuePtr_.get())); \
	95	+ break; \
	96	+ } \
	97	+ default: \
	98	+ std::cout << "[ERROR]" << __FILE__ << ":" << __LINE__ << "The ValueType " \
	99	+ << attrInfo.attr_.type_ << "is not supported!" << std::endl; \
	100	+ } \
	101	+ } \
	102	+ /* 2. base information */ \
	103	+ fe::PlatFormInfos platformInfo; \
	104	+ platformInfo.Init(); \
	105	+ auto tilingData = gert::TilingData::CreateCap(tilingContextPara.tilingDataSize_); \
	106	+ auto workspace = gert::ContinuousVector::Create<size_t>(4096); \
	107	+ auto contextHolder = \
	108	+ contextFaker.SetOpType(tilingContextPara.opName_.c_str()) \
	109	+ .CompileInfo(tilingContextPara.compileInfo_) \
	110	+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo)) \
	111	+ .TilingData(tilingData.get()) \
	112	+ .Workspace(reinterpret_cast<gert::ContinuousVector*>(workspace.get())) \
	113	+ .Build(); \
	114	+ string compileInfoStringPrefix = \
	115	+ R"({"hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1", "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false, "UB_SIZE": )"; \
	116	+ string compileInfoStringMiddle = \
	117	+ R"(, "L2_SIZE": 33554432, "L1_SIZE": 524288, "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": )"; \
	118	+ map<string, string> socToUpper = { \
	119	+ {"ascend910b", "Ascend910B"}, {"ascend910_93", "Ascend910_93"}, \
	120	+ {"ascend950", "Ascend950"}, {"ascend310p", "Ascend310P"}, \
	121	+ {"ascend910", "Ascend910"}, {"ascend310b", "Ascend310B"}, \
	122	+ {"ascend610lite", "Ascend610Lite"}, {"ascend031", "Ascend031"}, \
	123	+ {"ascend035", "Ascend035"}, {"kirinx90", "KrinX90"}, \
	124	+ {"kirin9030", "Kirin9030"}, {"mc62cm12a", "MC62CM12A"}}; \
	125	+ std::string buildSocVersion = STR(BUILD_SOC_VERSION); \
	126	+ if (!buildSocVersion.empty()) { \
	127	+ buildSocVersion = socToUpper[buildSocVersion]; \
	128	+ } \
	129	+ string compileInfoStringSuffix = R"(, "socVersion":)" \
	130	+ R"(")" + \
	131	+ buildSocVersion + R"("} })"; \
	132	+ string compileInfoString = \
	133	+ compileInfoStringPrefix + std::to_string(tilingContextPara.ubSize_) + \
	134	+ compileInfoStringMiddle + std::to_string(tilingContextPara.coreNum_) + \
	135	+ compileInfoStringSuffix; \
	136	+ map<string, string> socToArch = {{"Ascend310P", "2002"}, \
	137	+ {"Ascend910B", "2201"}, \
	138	+ {"Ascend910_93", "2201"}, \
	139	+ {"Ascend950", "3510"}, \
	140	+ {"Ascend910", "1001"}}; \
	141	+ map<string, string> socInfos; \
	142	+ map<string, string> aicoreSpec; \
	143	+ map<string, string> intrinsics; \
	144	+ map<string, string> socversions = {{"NpuArch", socToArch[buildSocVersion]}, \
	145	+ {"Short_SoC_version", buildSocVersion}}; \
	146	+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics); \
	147	+ auto tilingContext = contextHolder.GetContext(); \
	148	+ tilingContext->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos); \
	149	+ tilingContext->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec); \
	150	+ tilingContext->GetPlatformInfo()->SetCoreNumByCoreType("AICore"); \
	151	+ tilingContext->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics); \
	152	+ tilingContext->GetPlatformInfo()->SetPlatformRes("version", socversions); \
	153	+ /* 3. get tiling func */ \
	154	+ auto spaceRegistry = gert::DefaultOpImplSpaceRegistryV2::GetInstance().GetSpaceRegistry(); \
	155	+ if (spaceRegistry == nullptr) { \
	156	+ throw std::invalid_argument("not found spaceRegistry"); \
	157	+ } \
	158	+ auto functionStruct = spaceRegistry->GetOpImpl(tilingContextPara.opName_.c_str()); \
	159	+ if (functionStruct == nullptr) { \
	160	+ throw std::invalid_argument("not found " + tilingContextPara.opName_); \
	161	+ } \
	162	+ auto tilingFunc = functionStruct->tiling; /* 4. check tiling func */ \
	163	+ /* 4. check tiling func */ \
	164	+ auto tilingRet = tilingFunc(tilingContext);
	165	+
	166	+template <typename T> static string to_string(void* buf, size_t size) {
	167	+ string result;
	168	+ const T* data = reinterpret_cast<const T*>(buf);
	169	+ size_t len = size / sizeof(T);
	170	+ for (size_t i = 0; i < len; i++) {
	171	+ result += std::to_string(data[i]);
	172	+ result += " ";
	173	+ }
	174	+ return result;
	175	+}
	176	+
	177	+static void GetPlatFormInfos(const char* compileInfoStr, map<string, string>& socInfos,
	178	+ map<string, string>& aicoreSpec, map<string, string>& intrinsics) {
	179	+ string default_hardward_info = R"({
	180	+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1", "Intrinsic_fix_pipe_l0c2out": false,
	181	+ "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true,
	182	+ "Intrinsic_data_move_out2l1_nd2nz": false, "UB_SIZE": 262144, "L2_SIZE": 33554432,
	183	+ "L1_SIZE": 1048576, "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 262144,
	184	+ "CORE_NUM": 32}})";
	185	+ nlohmann::json compileInfoJson = nlohmann::json::parse(compileInfoStr);
	186	+ if (compileInfoJson.type() != nlohmann::json::value_t::object) {
	187	+ compileInfoJson = nlohmann::json::parse(default_hardward_info.c_str());
	188	+ }
	189	+
	190	+ map<string, string> socInfoKeys = {{"ai_core_cnt", "CORE_NUM"},
	191	+ {"l2_size", "L2_SIZE"},
	192	+ {"cube_core_cnt", "cube_core_cnt"},
	193	+ {"vector_core_cnt", "vector_core_cnt"},
	194	+ {"core_type_list", "core_type_list"}};
	195	+ socInfos["core_type_list"] = "AICore";
	196	+
	197	+ for (auto& t : socInfoKeys) {
	198	+ if (compileInfoJson.contains("hardware_info") &&
	199	+ compileInfoJson["hardware_info"].contains(t.second)) {
	200	+ auto& objJson = compileInfoJson["hardware_info"][t.second];
	201	+ if (objJson.is_number_integer()) {
	202	+ socInfos[t.first] =
	203	+ to_string(compileInfoJson["hardware_info"][t.second].get<uint32_t>());
	204	+ } else if (objJson.is_string()) {
	205	+ socInfos[t.first] = objJson;
	206	+ }
	207	+ }
	208	+ }
	209	+ map<string, string> aicoreSpecKeys = {{"ub_size", "UB_SIZE"},
	210	+ {"l0_a_size", "L0A_SIZE"},
	211	+ {"l0_b_size", "L0B_SIZE"},
	212	+ {"l0_c_size", "L0C_SIZE"},
	213	+ {"l1_size", "L1_SIZE"},
	214	+ {"bt_size", "BT_SIZE"},
	215	+ {"load3d_constraints", "load3d_constraints"}};
	216	+ aicoreSpec["cube_freq"] = "cube_freq";
	217	+ for (auto& t : aicoreSpecKeys) {
	218	+ if (compileInfoJson.contains("hardware_info") &&
	219	+ compileInfoJson["hardware_info"].contains(t.second)) {
	220	+ if (t.second == "load3d_constraints") {
	221	+ aicoreSpec[t.first] = compileInfoJson["hardware_info"][t.second].get<string>();
	222	+ } else {
	223	+ aicoreSpec[t.first] =
	224	+ to_string(compileInfoJson["hardware_info"][t.second].get<uint32_t>());
	225	+ }
	226	+ }
	227	+ }
	228	+
	229	+ std::string intrinsicsKeys[] = {
	230	+ "Intrinsic_data_move_l12ub", "Intrinsic_data_move_l0c2ub",
	231	+ "Intrinsic_fix_pipe_l0c2out", "Intrinsic_data_move_out2l1_nd2nz",
	232	+ "Intrinsic_matmul_ub_to_ub", "Intrinsic_conv_ub_to_ub",
	233	+ "Intrinsic_data_move_l12bt"};
	234	+ for (string key : intrinsicsKeys) {
	235	+ if (compileInfoJson.contains("hardware_info") &&
	236	+ compileInfoJson["hardware_info"].contains(key) &&
	237	+ compileInfoJson["hardware_info"][key].get<bool>()) {
	238	+ intrinsics[key] = "float16";
	239	+ if (key.find("Intrinsic_data_move_l12bt") != string::npos) {
	240	+ intrinsics[key] = "bf16";
	241	+ }
	242	+ }
	243	+ }
	244	+}
	245	+
	246	+void ExecuteTestCase(const gert::TilingContextPara& tilingContextPara, ge::graphStatus expectResult,
	247	+ uint64_t expectTilingKey, const string& expectTilingData,
	248	+ const std::vector<size_t>& expectWorkspaces) {
	249	+ DO_TILING(tilingContextPara);
	250	+
	251	+ // check tiling func
	252	+ EXPECT_EQ(tilingRet, expectResult);
	253	+ if (expectResult == ge::GRAPH_FAILED) {
	254	+ return;
	255	+ }
	256	+
	257	+ // check workspace
	258	+ size_t workspaceCount = tilingContext->GetWorkspaceNum();
	259	+ if (workspaceCount > 0) {
	260	+ ASSERT_EQ(workspaceCount, expectWorkspaces.size());
	261	+ auto workspaceSizes = tilingContext->GetWorkspaceSizes(workspaceCount);
	262	+ for (size_t i = 0; i < workspaceCount; i++) {
	263	+ ASSERT_EQ(workspaceSizes[i], expectWorkspaces[i]);
	264	+ }
	265	+ }
	266	+
	267	+ // check tiling key
	268	+ auto tilingKeyResult = tilingContext->GetTilingKey();
	269	+ ASSERT_EQ(tilingKeyResult, expectTilingKey);
	270	+
	271	+ // check tiling data
	272	+ if (expectTilingData == EMPTY_EXPECT_TILING_DATA) {
	273	+ return;
	274	+ }
	275	+ auto rawTilingData = tilingContext->GetRawTilingData();
	276	+ auto tilingDataResult =
	277	+ to_string<int64_t>(rawTilingData->GetData(), rawTilingData->GetDataSize());
	278	+ EXPECT_EQ(tilingDataResult, expectTilingData);
	279	+}
	280	+
	281	+void ExecuteTestCase(const gert::TilingContextPara& tilingContextPara, ge::graphStatus expectResult,
	282	+ uint64_t expectTilingKey, const std::vector<size_t>& expectWorkspaces) {
	283	+ ExecuteTestCase(tilingContextPara, expectResult, expectTilingKey, EMPTY_EXPECT_TILING_DATA,
	284	+ expectWorkspaces);
	285	+}
	286	+
	287	+bool ExecuteTiling(const gert::TilingContextPara& tilingContextPara, TilingInfo& tilingInfo) {
	288	+ DO_TILING(tilingContextPara);
	289	+
	290	+ if (tilingRet != ge::GRAPH_SUCCESS) {
	291	+ return false;
	292	+ }
	293	+
	294	+ tilingInfo.tilingKey = tilingContext->GetTilingKey();
	295	+ tilingInfo.blockNum = tilingContext->GetBlockDim();
	296	+ size_t workspaceCount = tilingContext->GetWorkspaceNum();
	297	+ if (workspaceCount > 0) {
	298	+ auto workSpaceSizes = tilingContext->GetWorkspaceSizes(workspaceCount);
	299	+ for (size_t i = 0; i < workspaceCount; i++) {
	300	+ tilingInfo.workspaceSizes.push_back(workSpaceSizes[i]);
	301	+ }
	302	+ }
	303	+ auto rawTilingData = tilingContext->GetRawTilingData();
	304	+ tilingInfo.tilingData = std::make_unique<uint8_t[]>(rawTilingData->GetDataSize());
	305	+ tilingInfo.tilingDataSize = rawTilingData->GetDataSize();
	306	+ std::memcpy(tilingInfo.tilingData.get(), rawTilingData->GetData(),
	307	+ rawTilingData->GetDataSize());
	308	+
	309	+ return true;
	310	+}

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/tiling_case_executor.h+28-0

@@ -0,0 +1,28 @@
	1	+#ifndef OPS_MATH_DEV_TESTS_UT_COMMON_TILING_CASE_EXECUTOR_H
	2	+#define OPS_MATH_DEV_TESTS_UT_COMMON_TILING_CASE_EXECUTOR_H
	3	+
	4	+#include "tiling_context_faker.h"
	5	+
	6	+using namespace std;
	7	+
	8	+const string EMPTY_EXPECT_TILING_DATA = "EMPTY_EXPECT_TILING_DATA";
	9	+
	10	+struct TilingInfo {
	11	+ int64_t tilingKey = -1;
	12	+ std::vector<int64_t> workspaceSizes;
	13	+ std::unique_ptr<uint8_t[]> tilingData;
	14	+ size_t tilingDataSize = 0;
	15	+ size_t blockNum = 0;
	16	+};
	17	+
	18	+void ExecuteTestCase(const gert::TilingContextPara& tilingContextPara,
	19	+ ge::graphStatus expectResult = ge::GRAPH_FAILED, uint64_t expectTilingKey = 0,
	20	+ const string& expectTilingData = "",
	21	+ const std::vector<size_t>& expectWorkspaces = {});
	22	+
	23	+void ExecuteTestCase(const gert::TilingContextPara& tilingContextPara, ge::graphStatus expectResult,
	24	+ uint64_t expectTilingKey, const std::vector<size_t>& expectWorkspaces);
	25	+
	26	+bool ExecuteTiling(const gert::TilingContextPara& tilingContextPara, TilingInfo& tilingInfo);
	27	+
	28	+#endif // OPS_MATH_DEV_TESTS_UT_COMMON_TILING_CASE_EXECUTOR_H

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/tiling_context_faker.cpp+61-0

@@ -0,0 +1,61 @@
	1	+#include "tiling_context_faker.h"
	2	+
	3	+namespace gert {
	4	+
	5	+TilingContextFaker& TilingContextFaker::SetOpType(const std::string opType) {
	6	+ OpTilingContextBuilder::OpType(opType.c_str()).OpName(opType.c_str());
	7	+ return *this;
	8	+}
	9	+
	10	+TilingContextFaker& TilingContextFaker::NodeIoNum(size_t inputNum, size_t outputNum) {
	11	+ OpTilingContextBuilder::IONum(inputNum, outputNum);
	12	+ return *this;
	13	+}
	14	+
	15	+TilingContextFaker&
	16	+TilingContextFaker::IrInstanceNum(const std::vector<uint32_t>& inputInstanceNum,
	17	+ const std::vector<uint32_t>& outputInstanceNum) {
	18	+ OpTilingContextBuilder::IOInstanceNum(inputInstanceNum, outputInstanceNum);
	19	+ return *this;
	20	+}
	21	+
	22	+TilingContextFaker& TilingContextFaker::InputTensors(const std::vector<Tensor*>& inputTensors) {
	23	+ OpTilingContextBuilder::InputTensors(inputTensors);
	24	+ return *this;
	25	+}
	26	+
	27	+TilingContextFaker& TilingContextFaker::OutputTensors(const std::vector<Tensor*>& outputTensors) {
	28	+ OpTilingContextBuilder::OutputTensors(outputTensors);
	29	+ return *this;
	30	+}
	31	+
	32	+TilingContextFaker& TilingContextFaker::CompileInfo(const void* compileInfo) {
	33	+ OpTilingContextBuilder::CompileInfo(compileInfo);
	34	+ return *this;
	35	+}
	36	+
	37	+TilingContextFaker& TilingContextFaker::PlatformInfo(const void* platformInfo) {
	38	+ OpTilingContextBuilder::PlatformInfo(platformInfo);
	39	+ return *this;
	40	+}
	41	+
	42	+TilingContextFaker& TilingContextFaker::DeterministicInfo(int32_t* deterministicInfo) {
	43	+ OpTilingContextBuilder::Deterministic(*deterministicInfo);
	44	+ return *this;
	45	+}
	46	+
	47	+TilingContextFaker& TilingContextFaker::TilingData(const void* tilingData) {
	48	+ OpTilingContextBuilder::TilingData(static_cast<const gert::TilingData*>(tilingData));
	49	+ return *this;
	50	+}
	51	+
	52	+TilingContextFaker& TilingContextFaker::Workspace(const ContinuousVector* workspace) {
	53	+ OpTilingContextBuilder::Workspace(workspace);
	54	+ return *this;
	55	+}
	56	+
	57	+ContextHolder<TilingContext> TilingContextFaker::Build() {
	58	+ return OpTilingContextBuilder::Build();
	59	+}
	60	+
	61	+} // namespace gert

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/common/tiling_context_faker.h+171-0

@@ -0,0 +1,171 @@
	1	+#ifndef OPS_MATH_DEV_TESTS_UT_COMMON_TILING_CONTEXT_FAKER_H
	2	+#define OPS_MATH_DEV_TESTS_UT_COMMON_TILING_CONTEXT_FAKER_H
	3	+
	4	+#include "op_tiling_context_builder.h"
	5	+#include "any_value.h"
	6	+
	7	+namespace gert {
	8	+
	9	+class TilingContextPara {
	10	+ public:
	11	+ class TensorDescription {
	12	+ public:
	13	+ TensorDescription(const gert::StorageShape& shape, ge::DataType dtype, ge::Format format,
	14	+ bool isConst = false, void* constValue = nullptr)
	15	+ : shape_(shape), dtype_(dtype), format_(format), isConst_(isConst),
	16	+ constValue_(constValue) {
	17	+ }
	18	+
	19	+ public:
	20	+ gert::StorageShape shape_;
	21	+ ge::DataType dtype_ = ge::DT_FLOAT;
	22	+ ge::Format format_ = ge::FORMAT_ND;
	23	+ bool isConst_ = false;
	24	+ void* constValue_ = nullptr;
	25	+ };
	26	+
	27	+ class OpAttr {
	28	+ public:
	29	+ OpAttr(const std::string& attrName, const Ops::Math::AnyValue& attr)
	30	+ : attrName_(attrName), attr_(attr) {
	31	+ }
	32	+
	33	+ public:
	34	+ std::string attrName_;
	35	+ Ops::Math::AnyValue attr_;
	36	+ };
	37	+
	38	+ public:
	39	+ TilingContextPara(const std::string& opName,
	40	+ const std::vector<TensorDescription>& inputTensorDesc,
	41	+ const std::vector<TensorDescription>& outputTensorDesc,
	42	+ const std::vector<OpAttr>& attrs, void* compileInfo = nullptr,
	43	+ uint64_t coreNum = 64, uint64_t ubSize = 262144,
	44	+ uint64_t tilingDataSize = 4096)
	45	+ : opName_(opName), inputInstanceNum_(), outputInstanceNum_(),
	46	+ inputTensorDesc_(inputTensorDesc), outputTensorDesc_(outputTensorDesc), attrs_(attrs),
	47	+ coreNum_(coreNum), ubSize_(ubSize), tilingDataSize_(tilingDataSize),
	48	+ compileInfo_(compileInfo) {
	49	+ }
	50	+
	51	+ TilingContextPara(const std::string& opName,
	52	+ const std::vector<TensorDescription>& inputTensorDesc,
	53	+ const std::vector<TensorDescription>& outputTensorDesc,
	54	+ void* compileInfo = nullptr, uint64_t coreNum = 64, uint64_t ubSize = 262144,
	55	+ uint64_t tilingDataSize = 4096)
	56	+ : opName_(opName), inputInstanceNum_(), outputInstanceNum_(),
	57	+ inputTensorDesc_(inputTensorDesc), outputTensorDesc_(outputTensorDesc), attrs_(),
	58	+ coreNum_(coreNum), ubSize_(ubSize), tilingDataSize_(tilingDataSize),
	59	+ compileInfo_(compileInfo) {
	60	+ }
	61	+
	62	+ TilingContextPara(const std::string& opName,
	63	+ const std::vector<TensorDescription>& inputTensorDesc,
	64	+ const std::vector<TensorDescription>& outputTensorDesc,
	65	+ const std::vector<OpAttr>& attrs,
	66	+ const std::vector<uint32_t>& inputInstanceNum,
	67	+ const std::vector<uint32_t>& outputInstanceNum, void* compileInfo = nullptr,
	68	+ uint64_t coreNum = 64, uint64_t ubSize = 262144,
	69	+ uint64_t tilingDataSize = 4096)
	70	+ : opName_(opName), inputInstanceNum_(inputInstanceNum),
	71	+ outputInstanceNum_(outputInstanceNum), inputTensorDesc_(inputTensorDesc),
	72	+ outputTensorDesc_(outputTensorDesc), attrs_(attrs), coreNum_(coreNum), ubSize_(ubSize),
	73	+ tilingDataSize_(tilingDataSize), compileInfo_(compileInfo) {
	74	+ }
	75	+
	76	+ TilingContextPara(const std::string& opName,
	77	+ const std::vector<TensorDescription>& inputTensorDesc,
	78	+ const std::vector<TensorDescription>& outputTensorDesc,
	79	+ const std::vector<uint32_t>& inputInstanceNum,
	80	+ const std::vector<uint32_t>& outputInstanceNum, void* compileInfo = nullptr,
	81	+ uint64_t coreNum = 64, uint64_t ubSize = 262144,
	82	+ uint64_t tilingDataSize = 4096)
	83	+ : opName_(opName), inputInstanceNum_(inputInstanceNum),
	84	+ outputInstanceNum_(outputInstanceNum), inputTensorDesc_(inputTensorDesc),
	85	+ outputTensorDesc_(outputTensorDesc), attrs_(), coreNum_(coreNum), ubSize_(ubSize),
	86	+ tilingDataSize_(tilingDataSize), compileInfo_(compileInfo) {
	87	+ }
	88	+
	89	+ public:
	90	+ std::string opName_;
	91	+ std::vector<uint32_t> inputInstanceNum_;
	92	+ std::vector<uint32_t> outputInstanceNum_;
	93	+ std::vector<TensorDescription> inputTensorDesc_;
	94	+ std::vector<TensorDescription> outputTensorDesc_;
	95	+ std::vector<OpAttr> attrs_;
	96	+ uint64_t coreNum_ = 64;
	97	+ uint64_t ubSize_ = 262144;
	98	+ uint64_t tilingDataSize_ = 4096;
	99	+ void* compileInfo_ = nullptr;
	100	+};
	101	+
	102	+class TilingContextFaker : public OpTilingContextBuilder {
	103	+ public:
	104	+ TilingContextFaker& SetOpType(const std::string opType);
	105	+
	106	+ /* only one can be choosed from IrInstanceNum */
	107	+ TilingContextFaker& NodeIoNum(size_t inputNum, size_t outputNum);
	108	+
	109	+ /* can be used for dynamic inputs/outputs
	110	+ * only one can be choosed from NodeIoNum */
	111	+ TilingContextFaker& IrInstanceNum(const std::vector<uint32_t>& inputInstanceNum,
	112	+ const std::vector<uint32_t>& outputInstanceNum);
	113	+
	114	+ TilingContextFaker& Attr(const std::string& attrName, bool attr) {
	115	+ OpTilingContextBuilder::AppendAttr(attr);
	116	+ return *this;
	117	+ }
	118	+ TilingContextFaker& Attr(const std::string& attrName, int64_t attr) {
	119	+ OpTilingContextBuilder::AppendAttr(attr);
	120	+ return *this;
	121	+ }
	122	+ TilingContextFaker& Attr(const std::string& attrName, float attr) {
	123	+ OpTilingContextBuilder::AppendAttr(attr);
	124	+ return *this;
	125	+ }
	126	+ TilingContextFaker& Attr(const std::string& attrName, const ge::AscendString& attr) {
	127	+ OpTilingContextBuilder::AppendAttr(attr);
	128	+ return *this;
	129	+ }
	130	+ TilingContextFaker& Attr(const std::string& attrName, const std::vector<bool>& attr) {
	131	+ OpTilingContextBuilder::AppendAttr(attr);
	132	+ return *this;
	133	+ }
	134	+ TilingContextFaker& Attr(const std::string& attrName, const std::vector<int64_t>& attr) {
	135	+ OpTilingContextBuilder::AppendAttr(attr);
	136	+ return *this;
	137	+ }
	138	+ TilingContextFaker& Attr(const std::string& attrName, const std::vector<float>& attr) {
	139	+ OpTilingContextBuilder::AppendAttr(attr);
	140	+ return *this;
	141	+ }
	142	+ TilingContextFaker& Attr(const std::string& attrName,
	143	+ const std::vector<ge::AscendString>& attr) {
	144	+ OpTilingContextBuilder::AppendAttr(attr);
	145	+ return *this;
	146	+ }
	147	+ TilingContextFaker& Attr(const std::string& attrName,
	148	+ const std::vector<std::vector<int64_t>>& attr) {
	149	+ OpTilingContextBuilder::AppendAttr(attr);
	150	+ return *this;
	151	+ }
	152	+
	153	+ TilingContextFaker& InputTensors(const std::vector<Tensor*>& inputTensors);
	154	+
	155	+ TilingContextFaker& OutputTensors(const std::vector<Tensor*>& outputTensors);
	156	+
	157	+ TilingContextFaker& CompileInfo(const void* compileInfo);
	158	+
	159	+ TilingContextFaker& PlatformInfo(const void* platformInfo);
	160	+
	161	+ TilingContextFaker& DeterministicInfo(int32_t* deterministicInfo);
	162	+
	163	+ TilingContextFaker& TilingData(const void* tilingData);
	164	+
	165	+ TilingContextFaker& Workspace(const ContinuousVector* workspace);
	166	+
	167	+ ContextHolder<TilingContext> Build();
	168	+};
	169	+
	170	+} // namespace gert
	171	+#endif // OPS_MATH_DEV_TESTS_UT_COMMON_INFERSHAPE_CONTEXT_FAKER_H

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/op_host/CMakeLists.txt+102-0

@@ -0,0 +1,102 @@
	1	+set(OP_NAME leakyrelu)
	2	+set(UT_EXE ${OP_NAME}_op_host_ut)
	3	+
	4	+if(NOT DEFINED SOC_VERSION)
	5	+ set(SOC_VERSION "Ascend910B")
	6	+endif()
	7	+
	8	+message(STATUS "Building UT for SOC_VERSION: ${SOC_VERSION}")
	9	+
	10	+set(OP_HOST_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host)
	11	+
	12	+set(TILING_SRC ${OP_HOST_DIR}/${OP_NAME}_tiling.cpp)
	13	+set(TILING_DATA_H ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_kernel/${OP_NAME}_tiling_data.h)
	14	+
	15	+set(OP_HOST_SRCS
	16	+ ${OP_HOST_DIR}/${OP_NAME}_def.cpp
	17	+ ${OP_HOST_DIR}/${OP_NAME}_infershape.cpp
	18	+ ${TILING_SRC}
	19	+)
	20	+
	21	+set(UT_TEST_SRCS
	22	+ test_${OP_NAME}_tiling.cpp
	23	+)
	24	+
	25	+set(UT_MAIN_SRC ${CMAKE_CURRENT_SOURCE_DIR}/test_op_host_main.cpp)
	26	+
	27	+set(OP_HOST_SO ${OP_NAME}_op_host_ut_lib)
	28	+
	29	+add_library(${OP_HOST_SO} SHARED
	30	+ ${OP_HOST_SRCS}
	31	+)
	32	+
	33	+target_include_directories(${OP_HOST_SO} PRIVATE
	34	+ ${UT_COMMON_INCLUDE_DIRS}
	35	+ ${OP_HOST_DIR}
	36	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_kernel
	37	+)
	38	+
	39	+target_compile_definitions(${OP_HOST_SO} PRIVATE
	40	+ BUILD_SOC_VERSION=${SOC_VERSION}
	41	+ _GLIBCXX_USE_CXX11_ABI=0
	42	+)
	43	+
	44	+target_link_libraries(${OP_HOST_SO} PRIVATE
	45	+ -Wl,--no-as-needed
	46	+ -Wl,--whole-archive
	47	+ $ENV{ASCEND_HOME_PATH}/lib64/librt2_registry.a
	48	+ $ENV{ASCEND_HOME_PATH}/lib64/libtiling_api.a
	49	+ -Wl,--no-whole-archive
	50	+ $ENV{ASCEND_HOME_PATH}/lib64/libopp_registry.so
	51	+ $ENV{ASCEND_HOME_PATH}/lib64/libregister.so
	52	+ $ENV{ASCEND_HOME_PATH}/lib64/libmmpa.so
	53	+)
	54	+
	55	+add_executable(${UT_EXE}
	56	+ ${UT_MAIN_SRC}
	57	+ ${UT_TEST_SRCS}
	58	+ ${UT_COMMON_SRCS}
	59	+)
	60	+
	61	+target_include_directories(${UT_EXE} PRIVATE
	62	+ ${UT_COMMON_INCLUDE_DIRS}
	63	+ ${OP_HOST_DIR}
	64	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_kernel
	65	+)
	66	+
	67	+target_compile_definitions(${UT_EXE} PRIVATE
	68	+ BUILD_SOC_VERSION=${SOC_VERSION}
	69	+ _GLIBCXX_USE_CXX11_ABI=0
	70	+)
	71	+
	72	+target_compile_options(${UT_EXE} PUBLIC
	73	+ -fPIE
	74	+ -fno-access-control
	75	+)
	76	+
	77	+target_link_directories(${UT_EXE} PRIVATE
	78	+ ${UT_COMMON_LIB_DIRS}
	79	+ ${CMAKE_CURRENT_BINARY_DIR}
	80	+)
	81	+
	82	+target_link_libraries(${UT_EXE} PRIVATE
	83	+ gtest
	84	+ gtest_main
	85	+ -Wl,--no-as-needed
	86	+ $ENV{ASCEND_HOME_PATH}/lib64/libmetadef.so
	87	+ $ENV{ASCEND_HOME_PATH}/lib64/libunified_dlog.so
	88	+ $ENV{ASCEND_HOME_PATH}/lib64/libopp_registry.so
	89	+ $ENV{ASCEND_HOME_PATH}/lib64/libregister.so
	90	+ $ENV{ASCEND_HOME_PATH}/lib64/libgraph.so
	91	+ $ENV{ASCEND_HOME_PATH}/lib64/libgraph_base.so
	92	+ $ENV{ASCEND_HOME_PATH}/lib64/libplatform.so
	93	+ $ENV{ASCEND_HOME_PATH}/lib64/libc_sec.so
	94	+ $ENV{ASCEND_HOME_PATH}/lib64/libmmpa.so
	95	+ dl
	96	+ pthread
	97	+)
	98	+
	99	+set_target_properties(${UT_EXE} PROPERTIES
	100	+ INSTALL_RPATH "$ORIGIN;$ENV{ASCEND_HOME_PATH}/lib64"
	101	+ BUILD_WITH_INSTALL_RPATH TRUE
	102	+)

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/op_host/test_leakyrelu_tiling.cpp+85-0

@@ -0,0 +1,85 @@
	1	+#include <iostream>
	2	+#include <gtest/gtest.h>
	3	+#include "tiling_context_faker.h"
	4	+#include "tiling_case_executor.h"
	5	+#include "leakyrelu_tiling_data.h"
	6	+
	7	+namespace LeakyreluUT {
	8	+using namespace std;
	9	+using namespace ge;
	10	+using namespace gert;
	11	+static const std::string OP_NAME = "Leakyrelu";
	12	+
	13	+struct LeakyreluTestParam {
	14	+ std::string caseName;
	15	+ std::initializer_list<int64_t> xShape;
	16	+ ge::DataType xDtype;
	17	+ ge::Format xFormat;
	18	+ std::initializer_list<int64_t> yShape;
	19	+ ge::DataType yDtype;
	20	+ ge::Format yFormat;
	21	+ std::string socVersion;
	22	+ ge::graphStatus status;
	23	+ uint64_t expectTilingKey;
	24	+ std::string expectTilingData;
	25	+ std::vector<size_t> expectWorkspaces;
	26	+ uint64_t maxAIVNum;
	27	+ uint64_t ubSize;
	28	+ uint64_t tilingDataMaxSize;
	29	+};
	30	+
	31	+// TODO: 以下期望值基于初始模板实现，修改 tiling 逻辑后请更新：
	32	+// expectTilingKey: 参考 op_kernel/leakyrelu_tiling_key.h 和 op_host/leakyrelu_tiling.cpp 中 tilingKey 的逻辑
	33	+// expectTilingData: 参考 op_host/leakyrelu_tiling.cpp 中 TilingData 各字段的赋值
	34	+// expectWorkspaces: 参考 op_host/leakyrelu_tiling.cpp 中 GetWorkspaceSize 的逻辑
	35	+static LeakyreluTestParam testCases[] = {
	36	+ {"leakyrelu_0", {8, 2048}, ge::DT_FLOAT, ge::FORMAT_ND, {8, 2048}, ge::DT_FLOAT, ge::FORMAT_ND, "Ascend910B", ge::GRAPH_SUCCESS, 1UL, "0 1 0 ", {0}, 64, 262144, 4096},
	atomgit-bot22 hours ago 🟠 High Priority 建议：先执行一次 tiling（可用 ExecuteTiling 取 TilingInfo），把 totalNum/blockFactor/ubFactor/slopeTag/negativeSlope+reserved 这 5 个 int64 槽位的真实值填入 expectTilingData；或改用字段级断言替代整段字符串比较。 Incorrect?
	37	+};
	38	+
	39	+class LeakyreluTilingTest : public testing::TestWithParam<LeakyreluTestParam> {
	40	+protected:
	41	+ static void SetUpTestCase() {
	42	+ std::cout << "LeakyreluTilingTest SetUp." << std::endl;
	43	+ }
	44	+ static void TearDownTestCase() {
	45	+ std::cout << "LeakyreluTilingTest TearDown." << std::endl;
	46	+ }
	47	+};
	48	+
	49	+struct LeakyreluCompileInfo {} compileInfo;
	50	+
	51	+static void TestOneParamCase(const LeakyreluTestParam &param)
	52	+{
	53	+ gert::StorageShape xShape = {param.xShape, param.xShape};
	54	+ gert::StorageShape yShape = {param.yShape, param.yShape};
	55	+ std::vector<gert::TilingContextPara::TensorDescription> inputTensorDesc_(
	56	+ {{xShape, param.xDtype, param.xFormat}});
	57	+ std::vector<gert::TilingContextPara::TensorDescription> outputTensorDesc_(
	58	+ {{yShape, param.yDtype, param.yFormat}});
	59	+ std::vector<gert::TilingContextPara::OpAttr> attrs_;
	60	+ attrs_.push_back(gert::TilingContextPara::OpAttr("negative_slope", Ops::Math::AnyValue::CreateFrom<float>(0.01)));
	61	+ gert::TilingContextPara tilingContextPara(
	62	+ OP_NAME,
	63	+ inputTensorDesc_,
	64	+ outputTensorDesc_,
	65	+ attrs_,
	66	+ &compileInfo,
	67	+ param.maxAIVNum,
	68	+ param.ubSize,
	69	+ param.tilingDataMaxSize);
	70	+ ExecuteTestCase(tilingContextPara, param.status, param.expectTilingKey,
	71	+ param.expectTilingData, param.expectWorkspaces);
	72	+}
	73	+
	74	+TEST_P(LeakyreluTilingTest, tiling_test)
	75	+{
	76	+ const LeakyreluTestParam &param = GetParam();
	77	+ TestOneParamCase(param);
	78	+}
	79	+
	80	+INSTANTIATE_TEST_SUITE_P(
	81	+ LeakyreluTilingTests,
	82	+ LeakyreluTilingTest,
	83	+ testing::ValuesIn(testCases));
	84	+
	85	+}

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/op_host/test_op_host_main.cpp+50-0

@@ -0,0 +1,50 @@
	1	+#include <filesystem>
	2	+#include <iostream>
	3	+#include <gtest/gtest.h>
	4	+#include "base/registry/op_impl_space_registry_v2.h"
	5	+
	6	+using namespace std;
	7	+
	8	+class OpHostUtEnvironment : public testing::Environment {
	9	+public:
	10	+ OpHostUtEnvironment(char** argv) : argv_(argv)
	11	+ {}
	12	+
	13	+ virtual void SetUp() {
	14	+ cout << "Global Environment SetUp." << endl;
	15	+
	16	+ std::filesystem::path currDir = argv_[0];
	17	+ if (currDir.is_relative()) {
	18	+ currDir = std::filesystem::weakly_canonical(std::filesystem::current_path() / currDir);
	19	+ } else {
	20	+ currDir = std::filesystem::canonical(currDir);
	21	+ }
	22	+
	23	+ string opHostSoPath = currDir.parent_path().string() + string("/libleakyrelu_op_host_ut_lib.so");
	24	+ cout << "Loading op_host .so from: " << opHostSoPath << endl;
	25	+
	26	+ gert::OppSoDesc oppSoDesc({ge::AscendString(opHostSoPath.c_str())}, "op_host_so");
	27	+
	28	+ shared_ptr<gert::OpImplSpaceRegistryV2> opImplSpaceRegistryV2 = make_shared<gert::OpImplSpaceRegistryV2>();
	29	+ if (opImplSpaceRegistryV2->AddSoToRegistry(oppSoDesc) == ge::GRAPH_FAILED) {
	30	+ cerr << "Failed to add .so to registry." << endl;
	31	+ return;
	32	+ }
	33	+
	34	+ gert::DefaultOpImplSpaceRegistryV2::GetInstance().SetSpaceRegistry(opImplSpaceRegistryV2);
	35	+ cout << "OpImplSpaceRegistryV2 initialized successfully" << endl;
	36	+ }
	37	+
	38	+ virtual void TearDown() {
	39	+ cout << "Global Environment TearDown" << endl;
	40	+ }
	41	+
	42	+private:
	43	+ char** argv_;
	44	+};
	45	+
	46	+int main(int argc, char** argv) {
	47	+ testing::InitGoogleTest(&argc, argv);
	48	+ testing::AddGlobalTestEnvironment(new OpHostUtEnvironment(argv));
	49	+ return RUN_ALL_TESTS();
	50	+}

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/op_kernel/CMakeLists.txt+97-0

@@ -0,0 +1,97 @@
	1	+set(OP_NAME leakyrelu)
	2	+set(UT_EXE ${OP_NAME}_op_kernel_ut)
	3	+
	4	+message(STATUS "Building op_kernel UT for: ${OP_NAME}")
	5	+
	6	+if(DEFINED ENV{ASCEND_HOME_PATH} AND EXISTS "$ENV{ASCEND_HOME_PATH}")
	7	+ set(CANN_HOME "$ENV{ASCEND_HOME_PATH}")
	8	+elseif(EXISTS "/usr/local/Ascend/cann-8.5.0")
	9	+ set(CANN_HOME "/usr/local/Ascend/cann-8.5.0")
	10	+else()
	11	+ message(FATAL_ERROR "Cannot find CANN installation")
	12	+endif()
	13	+
	14	+message(STATUS "Using CANN_HOME: ${CANN_HOME}")
	15	+
	16	+set(SOC_VERSION "Ascend910B1")
	17	+
	18	+set(tikicpulib_DIR ${CANN_HOME}/tools/tikicpulib/lib/cmake)
	19	+find_package(tikicpulib REQUIRED)
	20	+
	21	+set(OP_KERNEL_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_kernel)
	22	+
	23	+set(UT_TEST_SRCS
	24	+ test_${OP_NAME}.cpp
	25	+)
	26	+
	27	+set(UT_INCLUDE_DIRS
	28	+ ${CMAKE_CURRENT_SOURCE_DIR}
	29	+ ${OP_KERNEL_DIR}
	30	+ ${UT_COMMON_INCLUDE_DIRS}
	31	+ ${CANN_HOME}/tools/tikicpulib/lib/include
	32	+ ${CANN_HOME}/compiler/ascendc/include/highlevel_api
	33	+ ${CANN_HOME}/compiler/ascendc/include/basic_api/impl
	34	+ ${CANN_HOME}/compiler/ascendc/include/basic_api
	35	+ ${CANN_HOME}/compiler/ascendc/include
	36	+ ${CANN_HOME}/compiler/ascendc/include/ascendc/host_api/tiling
	37	+ ${CANN_HOME}/compiler/ascendc/include/ascendc/host_api
	38	+ ${CANN_HOME}/compiler/ascendc/include/ascendc
	39	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/include
	40	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include/basic_api
	41	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include
	42	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/include/interface
	43	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/impl/basic_api
	44	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc
	45	+ ${CANN_HOME}/compiler/tikcpp/tikcfw
	46	+ ${CANN_HOME}/compiler/tikcpp/tikcfw/impl
	47	+ ${CANN_HOME}/compiler/tikcpp/tikcfw/interface
	48	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/include/graph
	49	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/ascendc/include/highlevel_api
	50	+ ${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/asc/impl/basic_api/utils
	51	+)
	52	+
	53	+add_executable(${UT_EXE}
	54	+ ${UT_TEST_SRCS}
	55	+)
	56	+
	57	+target_include_directories(${UT_EXE} PRIVATE
	58	+ ${UT_INCLUDE_DIRS}
	59	+)
	60	+
	61	+target_compile_definitions(${UT_EXE} PRIVATE
	62	+ _GLIBCXX_USE_CXX11_ABI=0
	63	+ __aicore__=
	64	+ __NPU_TILING__=1
	65	+)
	66	+
	67	+target_compile_options(${UT_EXE} PRIVATE
	68	+ -Wall
	69	+ -Wno-deprecated-declarations
	70	+ -Wno-unused-variable
	71	+ -Wno-unknown-pragmas
	72	+ -fno-access-control
	73	+ -std=c++17
	74	+)
	75	+
	76	+target_link_directories(${UT_EXE} PRIVATE
	77	+ ${UT_COMMON_LIB_DIRS}
	78	+ ${CANN_HOME}/lib64
	79	+)
	80	+
	81	+target_link_libraries(${UT_EXE} PRIVATE
	82	+ gtest_main
	83	+ gtest
	84	+ ${CANN_HOME}/lib64/libmmpa.so
	85	+ tikicpulib::${SOC_VERSION}
	86	+ dl
	87	+ pthread
	88	+)
	89	+
	90	+set_target_properties(${UT_EXE} PROPERTIES
	91	+ INSTALL_RPATH "$ORIGIN;${CANN_HOME}/lib64;${CANN_HOME}/tools/tikicpulib/lib;${CANN_HOME}/tools/tikicpulib/lib/Ascend910B1;${CANN_HOME}/${CMAKE_SYSTEM_PROCESSOR}-linux/simulator/Ascend910B1/lib"
	92	+ BUILD_WITH_INSTALL_RPATH TRUE
	93	+)
	94	+
	95	+message(STATUS " Kernel UT executable: ${UT_EXE}")
	96	+message(STATUS " CANN Home: ${CANN_HOME}")
	97	+message(STATUS " SOC Version: ${SOC_VERSION}")

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/op_kernel/leakyrelu_data/README.md+1-0

	@@ -0,0 +1 @@
		1	+首点数据和比较入口为工程 samples/sample.py；opapi 内层 fixture 需自行指定。

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/op_kernel/leakyrelu_tiling.h+43-0

@@ -0,0 +1,43 @@
	1	+/*!
	2	+ * \file leakyrelu_tiling.h
	3	+ * \brief Leakyrelu tiling 数据定义
	4	+ */
	5	+
	6	+#ifndef _I_LEAKYRELU_TILING_H_
	7	+#define _I_LEAKYRELU_TILING_H_
	8	+
	9	+#include <cstdint>
	10	+#include <cstring>
	11	+#include "../../../op_kernel/leakyrelu_tiling_data.h"
	12	+#include "tikicpulib.h"
	13	+#include "kernel_operator.h"
	14	+#include "kernel_tiling/kernel_tiling.h"
	15	+#include "graph/c_types.h"
	16	+#include "ascendc/host_api/tiling/template_argument.h"
	17	+
	18	+#ifndef __aicore__
	19	+#define __aicore__
	20	+#endif
	21	+
	22	+#ifndef __gm__
	23	+#define __gm__
	24	+#endif
	25	+
	26	+#ifndef __ubuf__
	27	+#define __ubuf__
	28	+#endif
	29	+
	30	+inline void InitTilingData(uint8_t tiling, LeakyreluTilingData constData)
	31	+{
	32	+ memcpy(constData, tiling, sizeof(LeakyreluTilingData));
	33	+}
	34	+
	35	+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingArg) \
	36	+ tilingStruct tilingData; \
	37	+ InitTilingData(tilingArg, &tilingData)
	38	+
	39	+#define GET_TILING_DATA(tilingData, tilingArg) \
	40	+ LeakyreluTilingData tilingData; \
	41	+ InitTilingData(tilingArg, &tilingData)
	42	+
	43	+#endif

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/op_kernel/test_leakyrelu.cpp+99-0

@@ -0,0 +1,99 @@
	1	+/*!
	2	+ * \file test_leakyrelu.cpp
	3	+ * \brief Leakyrelu 算子 kernel UT 测试
	4	+ *
	5	+ * 独立运行，直接构造 tilingData，不依赖 op_host UT
	6	+ */
	7	+
	8	+#include "leakyrelu_tiling.h"
	9	+#include "../../../op_kernel/leakyrelu.cpp"
	10	+
	11	+#include <array>
	12	+#include <vector>
	13	+#include <iostream>
	14	+#include <cstdint>
	15	+#include <cstdlib>
	16	+#include <cstring>
	17	+#include <fstream>
	18	+#include "gtest/gtest.h"
	19	+#include "tikicpulib.h"
	20	+
	21	+using namespace std;
	22	+
	23	+static uint16_t FloatToHalf(float f) {
	24	+ uint32_t bits;
	25	+ memcpy(&bits, &f, sizeof(float));
	26	+ uint32_t sign = (bits >> 16) & 0x8000;
	27	+ int32_t exp = ((bits >> 23) & 0xff) - 127 + 15;
	28	+ uint32_t mant = (bits >> 13) & 0x3ff;
	29	+ if (exp <= 0) return sign;
	30	+ if (exp >= 31) return sign \| 0x7c00;
	31	+ return sign \| (exp << 10) \| mant;
	32	+}
	33	+
	34	+static uint16_t FloatToBFloat16(float f) {
	35	+ uint32_t bits;
	36	+ memcpy(&bits, &f, sizeof(float));
	37	+ return (uint16_t)(bits >> 16);
	38	+}
	39	+
	40	+class LeakyreluKernelTest : public testing::Test {
	41	+protected:
	42	+ static void SetUpTestCase()
	43	+ {
	44	+ cout << "LeakyreluKernelTest SetUp" << endl;
	45	+ }
	46	+ static void TearDownTestCase()
	47	+ {
	48	+ cout << "LeakyreluKernelTest TearDown" << endl;
	49	+ }
	50	+};
	51	+
	52	+TEST_F(LeakyreluKernelTest, test_kernel_run)
	53	+{
	54	+ constexpr size_t size = 16384;
	55	+ constexpr size_t tilingDataSize = sizeof(LeakyreluTilingData);
	56	+ constexpr uint32_t numBlocks = 1;
	57	+
	58	+ constexpr size_t xByteSize = 16384 * 4;
	59	+ constexpr size_t yByteSize = 16384 * 4;
	60	+ std::vector<uint8_t> xHost(xByteSize);
	61	+ std::vector<uint8_t> yHost(yByteSize);
	62	+
	63	+
	64	+ ASSERT_NE(std::getenv("SAMPLE_DIR"), nullptr);
	65	+ const std::string root = std::getenv("SAMPLE_DIR");
	66	+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(xByteSize);
	67	+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(yByteSize);
	68	+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(32);
	69	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);
	70	+
	71	+ { std::ifstream f(root + "/input/x.bin", std::ios::binary); ASSERT_TRUE(f.good()); f.read(reinterpret_cast<char*>(x), 65536); ASSERT_EQ(f.gcount(), 65536); }
	72	+
	73	+ // TODO: 以下 tilingData 字段基于初始模板的 TilingData 结构。
	74	+ // 修改 TilingData 结构后请更新字段名和赋值：
	75	+ // 参考 op_kernel/leakyrelu_tiling_data.h 中的字段定义
	76	+ // 参考 op_host/leakyrelu_tiling.cpp 中的 tiling 计算逻辑
	77	+ LeakyreluTilingData* tilingData = reinterpret_cast<LeakyreluTilingData*>(tiling);
	78	+ tilingData->totalNum = size;
	79	+ tilingData->blockFactor = size;
	80	+ tilingData->ubFactor = size;
	81	+
	82	+ // TODO: tilingKey 应与 op_host/leakyrelu_tiling.cpp 中 SetTilingKey 设置的值一致
	83	+ ICPU_SET_TILING_KEY(1);
	84	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	85	+
	86	+ ICPU_RUN_KF((leakyrelu<1>), numBlocks, x, y, workspace, tiling);
	87	+
	88	+ // 将动态输出的 packed buffer 拆回 individual buffers
	89	+
	90	+
	91	+ // 将 output 数据保存到 bin 文件供 compare_data.py 比对
	92	+ memcpy(yHost.data(), y, yByteSize);
	93	+ { std::ofstream _ofs((root + "/output/y.bin"), std::ios::binary); _ofs.write(reinterpret_cast<const char*>(yHost.data()), yByteSize); }
	94	+
	95	+ AscendC::GmFree(x);
	96	+ AscendC::GmFree(y);
	97	+ AscendC::GmFree(workspace);
	98	+ AscendC::GmFree(tiling);
	99	+}

A2026/CANN-Code-Detective/Challenge13-LeakyReluCustom/PR_code/ZY_weixin_48190404_13/code/exercises/tests/ut/run.sh+15-0

@@ -0,0 +1,15 @@
	1	+#!/bin/bash
	2	+set -euo pipefail
	3	+ROOT="$(cd "$(dirname "$0")/../.." && pwd)"
	4	+BUILD="$ROOT/tests/ut/build"
	5	+cmake -S "$ROOT/tests/ut" -B "$BUILD" -DGTEST_SOURCE_DIR="${GTEST_SOURCE_DIR:-}" -DENABLE_INNER_KERNEL_UT="${ENABLE_INNER_KERNEL_UT:-OFF}"
	6	+cmake --build "$BUILD" --parallel "${CANNJUDGE_COMPILE_JOBS:-4}"
	7	+"$BUILD/op_host/leakyrelu_op_host_ut"
	8	+python3 "$ROOT/samples/sample.py" generate --directory "$ROOT/samples/build-cpu"
	9	+export SAMPLE_DIR="$ROOT/samples/build-cpu"
	10	+"$BUILD/op_kernel/leakyrelu_op_kernel_ut"
	11	+python3 "$ROOT/samples/sample.py" verify --directory "$SAMPLE_DIR"
	12	+if [ -x "$BUILD/op_api/leakyrelu_op_api_ut" ]; then
	13	+ export ASCEND_CUSTOM_OPP_PATH="$ROOT/build/packages/vendors/leakyrelu_custom${ASCEND_CUSTOM_OPP_PATH:+:$ASCEND_CUSTOM_OPP_PATH}"
	14	+ "$BUILD/op_api/leakyrelu_op_api_ut"
	15	+fi