Open
[代码侦探(第十三期)] 个人开发者 - exercises #5638
[代码侦探(第十三期)] 个人开发者 - exercises #5638
Open
Code_demo111111Add on 1 day ago
Total: 39 files changed+2670-0
@@ -0,0 +1,11 @@
1+# 个人开发者
2+ 
3+## 团队信息
4+ 
5+- 提交者: Code_demo111111
6+- 身份: 企业员工
7+- 单位: 个人开发者
8+ 
9+## 成员
10+ 
11+- Code_demo111111 (Code_demo111111): 提交者
@@ -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()
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+ 
@@ -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
@@ -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")
@@ -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"
@@ -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");
atomgit-bot
atomgit-botatomgit-bot1 day ago

🟡 Medium Priority

变更行(sample_support.h:44-53):check(aclrtMalloc(&device, ...)) 成功分配 device 显存后,后续两步仍可能抛异常——check(aclrtMemcpy(...))(第 46 行)与 if (!tensor) throw(第 52-53 行)。 失败模式:C++ 规定构造函数抛出异常时该对象的析构函数不会执行(成员 host/file 会被析构,但 ~Tensor() 不会),因此已分配的 device GM 显存无人释放;同理 host 之外的 aclTensor 无需释放,但 device 一定会泄漏。 触发条件:aclrtMemcpy 返回非 0(如设备侧内存不足、非法拷贝参数)或 aclCreateTensor 返回 nullptr(strides/shape 非法、内存不足)。本 helper 被每个 Tensor 复用,一旦出现多次失败会累积泄漏。

Suggestions
53
+ check(aclrtMalloc(&device, std::max(size_t(1), size), ACL_MEM_MALLOC_HUGE_FIRST));
54
+ try {
55
+ if (size)
56
+ check(aclrtMemcpy(device, size, host.data(), size, ACL_MEMCPY_HOST_TO_DEVICE));
57
+ std::vector<int64_t> strides(shape.size(), 1);
58
+ for (size_t i = shape.size(); i > 1; --i)
59
+ strides[i - 2] = strides[i - 1] * shape[i - 1];
60
+ tensor = aclCreateTensor(shape.data(), shape.size(), type, strides.data(), 0, ACL_FORMAT_ND,
61
+ shape.data(), shape.size(), device);
62
+ if (!tensor)
53
- throw std::runtime_error("tensor creation failed");
63
+ throw std::runtime_error("tensor creation failed");
64
+ } catch (...) {
65
+ aclrtFree(device);
66
+ device = nullptr;
67
+ throw;
68
+ }
Apply suggestion
likedislike
Incorrect?
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+};
@@ -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+}
@@ -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()
@@ -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
@@ -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+ }
atomgit-bot
atomgit-botatomgit-bot1 day ago

🟡 Medium Priority

变更行:leakyrelu_infershape.cpp 第 25-32 行,if (in_shape != nullptr) { *output_shape = *in_shape; } 之后无条件 return ge::GRAPH_SUCCESS;。

契约/行为:Leakyrelu 只有 1 个必选输入,InferShape 的正常契约是「输出 shape = 输入 shape」;当 context->GetInputShape(0) 返回 nullptr(输入未找到、可选输入/子图切分边界、部分动态 shape 场景)时,本实现既不设置 *output_shape,也不报错,直接返回成功。

失效模式:框架把未被写入(rank/dims 处于初始/未知状态)的输出 Shape 当作合法结果继续下传,后续算子按错误的 rank/维度做 tiling 与内存分配,可能计算出错误的输出尺寸或触发越界读写;GetOutputShape(i) 的 nullptr 已经返回 GRAPH_FAILED,而更关键的输入侧缺失却被静默放过,属于错误处理不一致。

Suggestions
30
+ if (in_shape == nullptr) {
31
+ return ge::GRAPH_FAILED;
30
- }
32
+ }
33
+ *output_shape = *in_shape;
Apply suggestion
likedislike
Incorrect?
31+ }
32+ return ge::GRAPH_SUCCESS;
33+}
34+ 
35+IMPL_OP_INFERSHAPE(Leakyrelu).InferShape(InferShapeLeakyrelu);
36+ 
37+} // namespace ops
@@ -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
@@ -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+)
@@ -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+}
@@ -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;
atomgit-bot
atomgit-botatomgit-bot1 day ago

🟡 Medium Priority

变更行(leakyrelu.h:67-68):negativeSlope_ = (tilingData->slopeTag == LEAKYRELU_SLOPE_TAG) ? tilingData->negativeSlope : LEAKYRELU_DEFAULT_SLOPE;

  • 触发条件:host tiling 只要没有把 slopeTag 填为 LEAKYRELU_SLOPE_TAG(字段顺序/结构体版本不一致、属性解析遗漏、新增字段未同步初始化等),kernel 就会静默回退到硬编码 0.01。
likedislike
Incorrect?
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-bot
atomgit-botatomgit-bot1 day ago

🟡 Medium Priority

变更行(leakyrelu.h:117-118):const int64_t copyNum = AlignUp(currentNum, alignElems_); DataCopy(xLocal, inputGMX[offset], static_cast<int32_t>(copyNum));

  • 触发条件:totalNum 不是 alignElems_(float 8 / half 16)的整数倍时,最后一轮 tail 的 copyNum 会向上取整到 32B。
  • 失败模式:MTE2 会读到输入张量末尾之后最多 31B(half 为 30B)未分配的 GM,越界读;在开启越界检查的环境(仿真/UT/部分硬件校验)会报 AIC Error 或读到脏数据参与计算(虽然结果不写回,但可能污染生成 NaN/Inf 而影响 Add 的整块向量计算)。

建议:对非对齐尾块使用 DataCopyPad 按 currentNum*sizeof(T) 实际字节数搬入,或确保 host tiling 下发的 totalNum 已 32B 对齐,避免越界读。

likedislike
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));
atomgit-bot
atomgit-botatomgit-bot1 day ago

🟡 Medium Priority

变更行(leakyrelu.h:165):DataCopy(outputGMY[offset], yLocal, static_cast<int32_t>(currentNum));

  • 契约:基础 DataCopy(UB→GM,MTE3)要求搬运数据量为 32B 的整数倍;currentNum 在尾块由 blockLength_ - progress * ubLength_ 得到,当 totalNum 不是 alignElems_(float 8 个元素 / half 16 个元素)的整数倍时必然非对齐。
  • 失败模式:MTE3 按整块搬运时会把 32B 块尾部的冗余数据写到 outputGMY[offset+currentNum] 之后的 GM,踩坏输出张量尾部或相邻 tensor 数据;在开启越界检测时(kernel UT / 仿真)直接报 AIC Error。注释“只写回有效元素,避免越界污染输出张量尾部”的意图正确,但基础 DataCopy 无法保证,非对齐尾块必须用 DataCopyPad。

建议:尾块改 DataCopyPad 按实际字节数写回,或由 host tiling 保证 totalNum 32B 对齐后再用 DataCopy;同时 CopyOut 与 CopyIn 的对齐策略需保持一致。

likedislike
Incorrect?
166+ outputQueueY.FreeTensor(yLocal);
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
@@ -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
@@ -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+}
@@ -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+ 
@@ -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)
@@ -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)
@@ -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)
@@ -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
@@ -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+}
@@ -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
@@ -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
@@ -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
@@ -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+}
@@ -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
@@ -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
@@ -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
@@ -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()
atomgit-bot
atomgit-botatomgit-bot1 day ago

🟡 Medium Priority

变更行:op_host/CMakeLists.txt:4-6 设置 SOC_VERSION "Ascend910B",第 39-41 行以 BUILD_SOC_VERSION=${SOC_VERSION} 下发宏,tiling_case_executor.cpp:125 用 STR(BUILD_SOC_VERSION) 得到字符串 "Ascend910B"。

影响链路:tiling_case_executor.cpp:118-124 的 socToUpper 表键全部是小写("ascend910b"→"Ascend910B"、"ascend910_93" 等),第 126-127 行用 socToUpper[buildSocVersion] 查 "Ascend910B" 必然 miss;由于用的是 map::operator[],miss 时不是报错而是插入并返回空串,于是 buildSocVersion 变成 ""。

失败模式:run.sh 不传 -DSOC_VERSION,因此恒走该默认值。随后(1)第 144 行 socToArch[""] 再插入空值,socversions = {{"NpuArch",""},{"Short_SoC_version",""}} 被 SetPlatformRes("version", ...) 写入平台信息;(2)compileInfoString 结尾拼成 "socVersion":""。任何依赖 SoC 版本/NpuArch 做分支的 tiling(如按架构选择 UB 切分或特性开关)会取到空字符串走默认/错误分支,而测试仍会“通过”,把平台配置错误静默带过,属于 UT 侧算出的期望值不可信的可靠性问题。修复方向:默认值改为小写 "ascend910b"(或用 socToUpper.find + 大小写无关比较,miss 时直接报错)。

建议:将默认 SOC_VERSION 改为与 tiling_case_executor.cpp 中 socToUpper 键一致的小写形式("ascend910b"),或在该处统一转小写后再查表;同时确认 BUILD_SOC_VERSION 宏最终能命中 socToUpper 与 socToArch 两张表。

Suggestions
6
+ if(NOT DEFINED SOC_VERSION)
7
+ set(SOC_VERSION "ascend910b")
6
8
  endif()
Apply suggestion
likedislike
Incorrect?
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+)
@@ -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},
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+}
@@ -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+}
@@ -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}")
@@ -0,0 +1 @@
1+首点数据和比较入口为工程 samples/sample.py;opapi 内层 fixture 需自行指定。
@@ -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
@@ -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+}
@@ -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