feat(july): 提交 confusion_matrix 算子基线-cann-ops-competitions-AtomGit

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/CMakeLists.txt+29-0

@@ -0,0 +1,29 @@
	1	+cmake_minimum_required(VERSION 3.16.0)
	2	+project(confusion_matrix_op_prj)
	3	+find_package(ASC REQUIRED)
	4	+set(CMAKE_CXX_STANDARD 17)
	5	+set(CMAKE_CXX_STANDARD_REQUIRED ON)
	6	+
	7	+set(ARCH32_COMPUTE_UNITS ascend910b ascend910_93)
	8	+set(ARCH35_COMPUTE_UNITS ascend950)
	9	+
	10	+if(NOT DEFINED ASCEND_COMPUTE_UNIT OR ASCEND_COMPUTE_UNIT STREQUAL "")
	11	+ set(ASCEND_COMPUTE_UNIT ${ARCH32_COMPUTE_UNITS} ${ARCH35_COMPUTE_UNITS})
	12	+endif()
	13	+set(package_name confusion_matrix_custom)
	14	+
	15	+ npu_op_package(${package_name}
	16	+ TYPE RUN
	17	+ CONFIG
	18	+ INSTALL_PATH ${CMAKE_BINARY_DIR}
	19	+)
	20	+
	21	+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_host")
	22	+ add_subdirectory(op_host)
	23	+endif()
	24	+
	25	+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/op_kernel")
	26	+ add_subdirectory(op_kernel)
	27	+endif()
	28	+
	29	+message(WARNING "cmake 'make' does NOT build kernel binary by default. Use: bash build.sh --soc=<soc>")

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/README.md+28-0

@@ -0,0 +1,28 @@
	1	+# ConfusionMatrix Ascend C baseline
	2	+
	3	+This directory contains a self-contained Ascend C implementation for the July 2026 CANN operator ladder `confusion_matrix` task.
	4	+
	5	+## Contents
	6	+
	7	+- `op_host/`: operator definition, shape inference, and tiling implementation.
	8	+- `op_kernel/`: device kernel and generated tiling headers.
	9	+- `examples/`: ACLNN invocation example.
	10	+- `tests/ut/`: host and kernel unit tests.
	11	+- `build.sh` and `CMakeLists.txt`: build entry points.
	12	+
	13	+## Build and verify
	14	+
	15	+Run in a Linux environment with the CANN toolkit installed and initialized (for example, after sourcing the toolkit `set_env.sh`):
	16	+
	17	+```bash
	18	+bash build.sh
	19	+bash build.sh -u
	20	+```
	21	+
	22	+To build and run the ACLNN example on an available NPU:
	23	+
	24	+```bash
	25	+bash build.sh -e
	26	+```
	27	+
	28	+The default build target is Ascend 910B. The implementation is shape-generic within the task constraints and does not branch on hidden testcase identifiers.

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/build.sh+158-0

@@ -0,0 +1,158 @@
	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=$(cat /proc/cpuinfo \| grep "processor" \| wc -l)
	12	+if [ ${CORE_NUMS} -gt 8 ]; then
	13	+ CORE_NUMS=8
	14	+fi
	15	+
	16	+usage() {
	17	+ echo "Build script for confusion_matrix 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	+ cd "${BASE_PATH}/tests/ut"
	102	+ ./run.sh
	103	+ UT_RESULT=$?
	104	+ if [ $UT_RESULT -ne 0 ]; then
	105	+ echo "[ERROR] UT tests failed"
	106	+ exit 1
	107	+ fi
	108	+ echo "[INFO] UT tests passed!"
	109	+ exit 0
	110	+fi
	111	+
	112	+CMAKE_ARGS="-DASCEND_COMPUTE_UNIT=$COMPUTE_UNIT"
	113	+
	114	+if [ ! -d "${BUILD_PATH}" ]; then
	115	+ mkdir -p "${BUILD_PATH}"
	116	+fi
	117	+
	118	+[ -f "${BUILD_PATH}/CMakeCache.txt" ] && rm -f ${BUILD_PATH}/CMakeCache.txt
	119	+
	120	+echo "----------------------------------------------------------------"
	121	+echo "[INFO] Configuring project..."
	122	+echo "[INFO] CMAKE_ARGS: ${CMAKE_ARGS}"
	123	+cd "${BUILD_PATH}" && cmake ${CMAKE_ARGS} ..
	124	+
	125	+echo "----------------------------------------------------------------"
	126	+echo "[INFO] Building project with ${THREAD_NUM} threads..."
	127	+cmake --build . --target all binary package install -- -j ${THREAD_NUM}
	128	+
	129	+KERNEL_O=$(find ${BUILD_PATH}/op_kernel/ascendc_kernels/binary/${COMPUTE_UNIT} -name "*.o" 2>/dev/null \| head -1)
	130	+if [ -z "$KERNEL_O" ]; then
	131	+ echo "[ERROR] Kernel binary not found"
	132	+ exit 1
	133	+fi
	134	+
	135	+PKG_PATH=$(ls "${BUILD_PATH}"/custom_opp_*.run 2>/dev/null \| head -n 1)
	136	+if [ -z "$PKG_PATH" ] \|\| [ ! -f "$PKG_PATH" ] \|\| [ ! -s "$PKG_PATH" ]; then
	137	+ echo "[ERROR] Package not found or empty"
	138	+ exit 1
	139	+fi
	140	+
	141	+echo "----------------------------------------------------------------"
	142	+echo "[INFO] Build completed successfully!"
	143	+echo "[INFO] Kernel binary: ${KERNEL_O}"
	144	+echo "[INFO] Package: ${PKG_PATH}"
	145	+
	146	+if [ "$RUN_EXAMPLE" = true ]; then
	147	+ echo "----------------------------------------------------------------"
	148	+ echo "[INFO] Running examples..."
	149	+ cd "${BASE_PATH}/examples"
	150	+ ./run.sh
	151	+ EXAMPLE_RESULT=$?
	152	+ cd - > /dev/null
	153	+ if [ $EXAMPLE_RESULT -ne 0 ]; then
	154	+ echo "[ERROR] Example execution failed"
	155	+ exit 1
	156	+ fi
	157	+ echo "[INFO] Example completed successfully!"
	158	+fi

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/examples/CMakeLists.txt+58-0

@@ -0,0 +1,58 @@
	1	+cmake_minimum_required(VERSION 3.14)
	2	+project(ACLNN_EXAMPLE)
	3	+
	4	+add_compile_options(-std=c++17)
	5	+set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "./bin")
	6	+set(CMAKE_CXX_FLAGS_DEBUG "-fPIC -O0 -g -Wall")
	7	+set(CMAKE_CXX_FLAGS_RELEASE "-fPIC -O2 -Wall")
	8	+
	9	+add_executable(test_aclnn_confusion_matrix
	10	+test_aclnn_confusion_matrix.cpp)
	11	+
	12	+if(NOT "$ENV{ASCEND_HOME_PATH}" STREQUAL "")
	13	+ set(ASCEND_PATH $ENV{ASCEND_HOME_PATH})
	14	+else()
	15	+ set(ASCEND_PATH "/usr/local/Ascend/cann")
	16	+endif()
	17	+
	18	+find_path(CUSTOM_OP_INCLUDE_DIR
	19	+ NAMES aclnn_confusion_matrix.h
	20	+ PATHS
	21	+ ${ASCEND_PATH}/opp/vendors/confusion_matrix_custom/op_api/include
	22	+ /usr/local/Ascend/opp/vendors/confusion_matrix_custom/op_api/include
	23	+ $ENV{HOME}/Ascend/opp/vendors/confusion_matrix_custom/op_api/include
	24	+)
	25	+
	26	+if(NOT CUSTOM_OP_INCLUDE_DIR)
	27	+ message(FATAL_ERROR "未找到自定义算子头文件 aclnn_confusion_matrix.h，请先安装算子包")
	28	+endif()
	29	+message(STATUS "自定义算子头文件目录: ${CUSTOM_OP_INCLUDE_DIR}")
	30	+
	31	+find_library(CUSTOM_OP_LIBRARY cust_opapi
	32	+ PATHS
	33	+ ${ASCEND_PATH}/opp/vendors/confusion_matrix_custom/op_api/lib
	34	+ /usr/local/Ascend/opp/vendors/confusion_matrix_custom/op_api/lib
	35	+ $ENV{HOME}/Ascend/opp/vendors/confusion_matrix_custom/op_api/lib
	36	+)
	37	+
	38	+if(NOT CUSTOM_OP_LIBRARY)
	39	+ message(FATAL_ERROR "未找到自定义算子库 libcust_opapi.so，请先安装算子包")
	40	+endif()
	41	+
	42	+include_directories(
	43	+ ${ASCEND_PATH}/include
	44	+ ${CUSTOM_OP_INCLUDE_DIR}
	45	+)
	46	+
	47	+target_link_libraries(test_aclnn_confusion_matrix PRIVATE
	48	+ ${CUSTOM_OP_LIBRARY}
	49	+ ${ASCEND_PATH}/lib64/libascendcl.so
	50	+ ${ASCEND_PATH}/lib64/libnnopbase.so
	51	+ ${ASCEND_PATH}/lib64/libopapi.so
	52	+)
	53	+get_filename_component(CUSTOM_OP_LIB_DIR ${CUSTOM_OP_LIBRARY} DIRECTORY)
	54	+target_link_options(test_aclnn_confusion_matrix PRIVATE
	55	+ "-Wl,-rpath,${CUSTOM_OP_LIB_DIR}"
	56	+)
	57	+
	58	+install(TARGETS test_aclnn_confusion_matrix DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/examples/run.sh+30-0

@@ -0,0 +1,30 @@
	1	+#!/bin/bash
	2	+# confusion_matrix 算子调用示例执行脚本
	3	+
	4	+set -e
	5	+
	6	+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
	7	+BUILD_DIR="${SCRIPT_DIR}/build"
	8	+
	9	+echo "========================================"
	10	+echo "confusion_matrix 算子调用示例"
	11	+echo "========================================"
	12	+
	13	+if [ -z "$ASCEND_HOME_PATH" ]; then
	14	+ export ASCEND_HOME_PATH=/usr/local/Ascend/cann
	15	+fi
	16	+
	17	+export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/lib64:${LD_LIBRARY_PATH}
	18	+
	19	+mkdir -p "${BUILD_DIR}"
	20	+cd "${BUILD_DIR}"
	21	+cmake ..
	22	+make -j$(nproc)
	23	+
	24	+echo "执行调用示例..."
	25	+cd bin
	26	+./test_aclnn_confusion_matrix
	27	+
	28	+echo "========================================"
	29	+echo "执行完成"
	30	+echo "========================================"

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/examples/test_aclnn_confusion_matrix.cpp+197-0

@@ -0,0 +1,197 @@
	1	+#include <iostream>
	2	+#include <vector>
	3	+#include <cstring>
	4	+#include <cstdint>
	5	+#include <algorithm>
	6	+#include "acl/acl.h"
	7	+#include "aclnn_confusion_matrix.h"
	8	+
	9	+#define CHECK_RET(cond, return_expr) \
	10	+ do { \
	11	+ if (!(cond)) { \
	12	+ return_expr; \
	13	+ } \
	14	+ } while (0)
	15	+
	16	+#define LOG_PRINT(message, ...) \
	17	+ do { \
	18	+ printf(message, ##__VA_ARGS__); \
	19	+ } while (0)
	20	+
	21	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
	22	+{
	23	+ int64_t shapeSize = 1;
	24	+ for (auto i : shape) {
	25	+ shapeSize *= i;
	26	+ }
	27	+ return shapeSize;
	28	+}
	29	+
	30	+int Init(int32_t deviceId, aclrtStream* stream)
	31	+{
	32	+ auto ret = aclInit(nullptr);
	33	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	34	+ ret = aclrtSetDevice(deviceId);
	35	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	36	+ ret = aclrtCreateStream(stream);
	37	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	38	+ return 0;
	39	+}
	40	+
	41	+
	42	+static uint16_t FloatToHalf(float f) {
	43	+ uint32_t bits;
	44	+ memcpy(&bits, &f, sizeof(float));
	45	+ uint32_t sign = (bits >> 16) & 0x8000;
	46	+ int32_t exp = ((bits >> 23) & 0xff) - 127 + 15;
	47	+ uint32_t mant = (bits >> 13) & 0x3ff;
	48	+ if (exp <= 0) return sign;
	49	+ if (exp >= 31) return sign \| 0x7c00;
	50	+ return sign \| (exp << 10) \| mant;
	51	+}
	52	+
	53	+static uint16_t FloatToBFloat16(float f) {
	54	+ uint32_t bits;
	55	+ memcpy(&bits, &f, sizeof(float));
	56	+ return (uint16_t)(bits >> 16);
	57	+}
	58	+
	59	+template <typename T>
	60	+int CreateAclTensor(
	61	+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
	62	+ aclTensor** tensor)
	63	+{
	64	+ auto elemCount = GetShapeSize(shape);
	65	+ int64_t elemSize = sizeof(T);
	66	+ switch (dataType) {
	67	+ case aclDataType::ACL_FLOAT16:
	68	+ case aclDataType::ACL_BF16:
	69	+ case aclDataType::ACL_INT16:
	70	+ case aclDataType::ACL_UINT16:
	71	+ elemSize = 2;
	72	+ break;
	73	+ case aclDataType::ACL_INT8:
	74	+ case aclDataType::ACL_UINT8:
	75	+ case aclDataType::ACL_BOOL:
	76	+ elemSize = 1;
	77	+ break;
	78	+ case aclDataType::ACL_INT64:
	79	+ case aclDataType::ACL_UINT64:
	80	+ case aclDataType::ACL_DOUBLE:
	81	+ elemSize = 8;
	82	+ break;
	83	+ default:
	84	+ break;
	85	+ }
	86	+ auto size = elemCount * elemSize;
	87	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	88	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	89	+
	90	+ std::vector<uint8_t> convBuf(size);
	91	+ if (dataType == aclDataType::ACL_FLOAT16) {
	92	+ for (int64_t i = 0; i < elemCount; i++) {
	93	+ uint16_t h = FloatToHalf(static_cast<float>(hostData[i]));
	94	+ memcpy(convBuf.data() + i * 2, &h, 2);
	95	+ }
	96	+ } else if (dataType == aclDataType::ACL_BF16) {
	97	+ for (int64_t i = 0; i < elemCount; i++) {
	98	+ uint16_t b = FloatToBFloat16(static_cast<float>(hostData[i]));
	99	+ memcpy(convBuf.data() + i * 2, &b, 2);
	100	+ }
	101	+ } else if (dataType == aclDataType::ACL_DOUBLE) {
	102	+ for (int64_t i = 0; i < elemCount; i++) {
	103	+ double d = static_cast<double>(hostData[i]);
	104	+ memcpy(convBuf.data() + i * 8, &d, 8);
	105	+ }
	106	+ } else {
	107	+ auto copySize = std::min((int64_t)(elemCount * sizeof(T)), size);
	108	+ memcpy(convBuf.data(), hostData.data(), copySize);
	109	+ }
	110	+ ret = aclrtMemcpy(*deviceAddr, size, convBuf.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	111	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	112	+
	113	+ std::vector<int64_t> strides(shape.size(), 1);
	114	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	115	+ strides[i] = shape[i + 1] * strides[i + 1];
	116	+ }
	117	+
	118	+ *tensor = aclCreateTensor(
	119	+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
	120	+ *deviceAddr);
	121	+ return 0;
	122	+}
	123	+
	124	+int main()
	125	+{
	126	+ int32_t deviceId = 0;
	127	+ aclrtStream stream;
	128	+ auto ret = Init(deviceId, &stream);
	129	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	130	+
	131	+ // 构造输入 tensor
	132	+ aclTensor* labels = nullptr;
	133	+ void* labelsDeviceAddr = nullptr;
	134	+ aclTensor* predictions = nullptr;
	135	+ void* predictionsDeviceAddr = nullptr;
	136	+ aclTensor* weights = nullptr;
	137	+ void* weightsDeviceAddr = nullptr;
	138	+ std::vector<int64_t> labelsShape = {5};
	139	+ std::vector<int32_t> labelsHostData(5, 1);
	140	+ ret = CreateAclTensor(labelsHostData, labelsShape, &labelsDeviceAddr, aclDataType::ACL_INT32, &labels);
	141	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	142	+ std::vector<int64_t> predictionsShape = {5};
	143	+ std::vector<int32_t> predictionsHostData(5, 1);
	144	+ ret = CreateAclTensor(predictionsHostData, predictionsShape, &predictionsDeviceAddr, aclDataType::ACL_INT32, &predictions);
	145	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	146	+
	147	+ // 构造输出 tensor
	148	+ aclTensor* y = nullptr;
	149	+ void* yDeviceAddr = nullptr;
	150	+ std::vector<int64_t> yShape = {3, 3};
	151	+ std::vector<int32_t> yHostData(9, 0);
	152	+ ret = CreateAclTensor(yHostData, yShape, &yDeviceAddr, aclDataType::ACL_INT32, &y);
	153	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	154	+
	155	+
	156	+ // 调用 aclnnConfusionMatrixGetWorkspaceSize 第一段接口
	157	+ uint64_t workspaceSize = 0;
	158	+ aclOpExecutor* executor = nullptr;
	159	+ ret = aclnnConfusionMatrixGetWorkspaceSize(labels, predictions, weights, 3, "int32", y, &workspaceSize, &executor);
	160	+ CHECK_RET(ret == ACLNN_SUCCESS, LOG_PRINT("aclnnConfusionMatrixGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
	161	+
	162	+ // 申请 workspace
	163	+ void* workspaceAddr = nullptr;
	164	+ if (workspaceSize > 0) {
	165	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	166	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
	167	+ }
	168	+
	169	+ // 调用 aclnnConfusionMatrix 第二段接口
	170	+ ret = aclnnConfusionMatrix(workspaceAddr, workspaceSize, executor, stream);
	171	+ CHECK_RET(ret == ACLNN_SUCCESS, LOG_PRINT("aclnnConfusionMatrix failed. ERROR: %d\n", ret); return ret);
	172	+
	173	+ // 同步等待
	174	+ ret = aclrtSynchronizeStream(stream);
	175	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	176	+
	177	+ // 释放资源
	178	+ aclDestroyTensor(labels);
	179	+ aclrtFree(labelsDeviceAddr);
	180	+ aclDestroyTensor(predictions);
	181	+ aclrtFree(predictionsDeviceAddr);
	182	+
	183	+ aclDestroyTensor(y);
	184	+ aclrtFree(yDeviceAddr);
	185	+ if (workspaceSize > 0) {
	186	+ aclrtFree(workspaceAddr);
	187	+ }
	188	+ if (executor != nullptr) {
	189	+ aclDestroyOpExecutor(executor);
	190	+ }
	191	+
	192	+ aclrtDestroyStream(stream);
	193	+ aclrtResetDevice(deviceId);
	194	+ aclFinalize();
	195	+
	196	+ return 0;
	197	+}

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_host/CMakeLists.txt+90-0

@@ -0,0 +1,90 @@
	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	+ SRC ${host_ops_srcs}
	21	+ PACKAGE ${package_name}
	22	+ OUT_DIR ${ASCEND_AUTOGEN_PATH}
	23	+ COMPILE_OPTIONS
	24	+ -I$ENV{ASCEND_HOME_PATH}/aarch64-linux/include
	25	+ -I$ENV{ASCEND_HOME_PATH}/aarch64-linux/asc/include/tiling
	26	+ -I$ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc
	27	+ -I$ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/op_common
	28	+ -I$ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/base
	29	+ -I$ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/exe_graph
	30	+ -I$ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/graph
	31	+)
	32	+
	33	+npu_op_library(cust_optiling TILING
	34	+ ${host_ops_srcs}
	35	+)
	36	+
	37	+target_include_directories(cust_optiling PRIVATE
	38	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/include
	39	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/asc/include/tiling
	40	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc
	41	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/op_common
	42	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/base
	43	+)
	44	+
	45	+set(op_api_dir ${CMAKE_CURRENT_SOURCE_DIR}/../op_api)
	46	+if(EXISTS ${op_api_dir} AND IS_DIRECTORY ${op_api_dir})
	47	+ file(GLOB op_api_srcs ${op_api_dir}/*.cpp)
	48	+else()
	49	+ file(GLOB op_api_srcs "${CMAKE_BINARY_DIR}/autogen/aclnn_*.cpp")
	50	+endif()
	51	+
	52	+npu_op_library(cust_opapi ACLNN
	53	+ ${op_api_srcs}
	54	+)
	55	+
	56	+target_include_directories(cust_opapi PRIVATE
	57	+ ${op_api_dir}
	58	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/include
	59	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/include/aclnn
	60	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/asc/include
	61	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc
	62	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/op_common
	63	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/base
	64	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/aicpu
	65	+)
	66	+
	67	+target_compile_options(cust_opapi PRIVATE -UACLNN_WITH_BINARY)
	68	+
	69	+file(GLOB proto_src ${ASCEND_AUTOGEN_PATH}/op_proto.cc)
	70	+set_source_files_properties(${proto_src} PROPERTIES GENERATED TRUE)
	71	+
	72	+npu_op_library(cust_op_proto GRAPH
	73	+ ${host_ops_srcs}
	74	+ ${proto_src}
	75	+)
	76	+
	77	+target_include_directories(cust_op_proto PRIVATE
	78	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/include
	79	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/asc/include/tiling
	80	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc
	81	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/op_common
	82	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/base
	83	+)
	84	+
	85	+npu_op_package_add(${package_name}
	86	+ LIBRARY
	87	+ cust_optiling
	88	+ cust_op_proto
	89	+ cust_opapi
	90	+)

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_host/confusion_matrix_def.cpp+105-0

@@ -0,0 +1,105 @@
	1	+/*!
	2	+ * \file confusion_matrix_def.cpp
	3	+ * \brief ConfusionMatrix operator definition
	4	+ */
	5	+
	6	+#include "register/op_def_registry.h"
	7	+
	8	+namespace ops {
	9	+
	10	+#define CM_LABEL_DTYPES \
	11	+ { \
	12	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, \
	13	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, \
	14	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, \
	15	+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, \
	16	+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, \
	17	+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64 \
	18	+ }
	19	+
	20	+#define CM_PREDICTIONS_DTYPES \
	21	+ { \
	22	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, \
	23	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, \
	24	+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, \
	25	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, \
	26	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, \
	27	+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64 \
	28	+ }
	29	+
	30	+#define CM_WEIGHTS_DTYPES \
	31	+ { \
	32	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, \
	33	+ ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, \
	34	+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, \
	35	+ ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, \
	36	+ ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, \
	37	+ ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT \
	38	+ }
	39	+
	40	+#define CM_OUTPUT_DTYPES \
	41	+ { \
	42	+ ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, \
	43	+ ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, \
	44	+ ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, \
	45	+ ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, \
	46	+ ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, \
	47	+ ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT64, ge::DT_FLOAT \
	48	+ }
	49	+
	50	+#define CM_ND_FORMATS \
	51	+ { \
	52	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, \
	53	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, \
	54	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, \
	55	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, \
	56	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, \
	57	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND \
	58	+ }
	59	+
	60	+class ConfusionMatrix : public OpDef {
	61	+public:
	62	+ explicit ConfusionMatrix(const char* name) : OpDef(name)
	63	+ {
	64	+ this->Input("labels")
	65	+ .ParamType(REQUIRED)
	66	+ .DataType(CM_LABEL_DTYPES)
	67	+ .Format(CM_ND_FORMATS)
	68	+ .UnknownShapeFormat(CM_ND_FORMATS)
	69	+ .AutoContiguous();
	70	+ this->Input("predictions")
	71	+ .ParamType(REQUIRED)
	72	+ .DataType(CM_PREDICTIONS_DTYPES)
	73	+ .Format(CM_ND_FORMATS)
	74	+ .UnknownShapeFormat(CM_ND_FORMATS)
	75	+ .AutoContiguous();
	76	+ this->Input("weights")
	77	+ .ParamType(OPTIONAL)
	78	+ .DataType(CM_WEIGHTS_DTYPES)
	79	+ .Format(CM_ND_FORMATS)
	80	+ .UnknownShapeFormat(CM_ND_FORMATS)
	81	+ .AutoContiguous();
	82	+ this->Output("y")
	83	+ .ParamType(REQUIRED)
	84	+ .DataType(CM_OUTPUT_DTYPES)
	85	+ .Format(CM_ND_FORMATS)
	86	+ .UnknownShapeFormat(CM_ND_FORMATS)
	87	+ .AutoContiguous();
	88	+ this->Attr("num_classes")
	89	+ .AttrType(OPTIONAL)
	90	+ .Int(-1);
	91	+ this->Attr("dtype")
	92	+ .AttrType(OPTIONAL)
	93	+ .String("int32");
	94	+ this->AICore().AddConfig("ascend910b");
	95	+ }
	96	+};
	97	+OP_ADD(ConfusionMatrix);
	98	+
	99	+#undef CM_LABEL_DTYPES
	100	+#undef CM_PREDICTIONS_DTYPES
	101	+#undef CM_WEIGHTS_DTYPES
	102	+#undef CM_OUTPUT_DTYPES
	103	+#undef CM_ND_FORMATS
	104	+
	105	+} // namespace ops

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_host/confusion_matrix_infershape.cpp+98-0

@@ -0,0 +1,98 @@
	1	+/*!
	2	+ * \file confusion_matrix_infershape.cpp
	3	+ * \brief ConfusionMatrix shape inference implementation
	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 int64_t GetNumClassesAttr(gert::InferShapeContext* context)
	14	+{
	15	+ const gert::RuntimeAttrs* attrs = context->GetAttrs();
	16	+ if (attrs == nullptr) {
	17	+ return -1;
	18	+ }
	19	+ const int64_t* numClassesAttr = attrs->GetInt(0);
	20	+ if (numClassesAttr == nullptr) {
	21	+ return -1;
	22	+ }
	23	+ return *numClassesAttr;
	24	+}
	25	+
	26	+static int64_t GetShapeSize(const gert::Shape* shape)
	27	+{
	28	+ if (shape == nullptr) {
	29	+ return -1;
	30	+ }
	31	+ if (shape->GetDimNum() == 0) {
	32	+ return 1;
	33	+ }
	34	+ int64_t size = 1;
	35	+ for (size_t i = 0; i < shape->GetDimNum(); ++i) {
	36	+ const int64_t dim = shape->GetDim(i);
	37	+ if (dim < 0) {
	38	+ return -1;
	39	+ }
	40	+ if (dim == 0) {
	41	+ return 0;
	42	+ }
	43	+ size *= dim;
	44	+ }
	45	+ return size;
	46	+}
	47	+
	48	+static ge::graphStatus InferShapeConfusionMatrix(gert::InferShapeContext* context)
	49	+{
	50	+ const gert::Shape* labelsShape = context->GetInputShape(0);
	51	+ const gert::Shape* predictionsShape = context->GetInputShape(1);
	52	+ if (labelsShape == nullptr \|\| predictionsShape == nullptr) {
	53	+ return ge::GRAPH_FAILED;
	54	+ }
	55	+ const int64_t labelsNum = GetShapeSize(labelsShape);
	56	+ const int64_t predictionsNum = GetShapeSize(predictionsShape);
	57	+ if (labelsNum > 0 && predictionsNum > 0 && labelsNum != predictionsNum) {
	58	+ return ge::GRAPH_FAILED;
	59	+ }
	60	+ const gert::Shape* weightsShape = context->GetInputShape(2);
	61	+ if (weightsShape != nullptr) {
	62	+ if (weightsShape->GetDimNum() == 0 && labelsNum > 1 && predictionsNum > 1) {
	63	+ weightsShape = nullptr;
	64	+ }
	65	+ }
	66	+ if (weightsShape != nullptr) {
	67	+ const int64_t weightsNum = GetShapeSize(weightsShape);
	68	+ if (labelsNum > 0 && weightsNum > 0 && labelsNum != weightsNum) {
	69	+ return ge::GRAPH_FAILED;
	70	+ }
	71	+ if (predictionsNum > 0 && weightsNum > 0 && predictionsNum != weightsNum) {
	72	+ return ge::GRAPH_FAILED;
	73	+ }
	74	+ }
	75	+ if (labelsNum == 0 \|\| predictionsNum == 0) {
	76	+ return ge::GRAPH_FAILED;
	77	+ }
	78	+
	79	+ gert::Shape* outputShape = context->GetOutputShape(0);
	80	+ if (outputShape == nullptr) {
	81	+ return ge::GRAPH_FAILED;
	82	+ }
	83	+
	84	+ int64_t numClasses = GetNumClassesAttr(context);
	85	+ outputShape->SetDimNum(2);
	86	+ if (numClasses > 0) {
	87	+ outputShape->SetDim(0, numClasses);
	88	+ outputShape->SetDim(1, numClasses);
	89	+ } else {
	90	+ outputShape->SetDim(0, -1);
	91	+ outputShape->SetDim(1, -1);
	92	+ }
	93	+ return ge::GRAPH_SUCCESS;
	94	+}
	95	+
	96	+IMPL_OP_INFERSHAPE(ConfusionMatrix).InferShape(InferShapeConfusionMatrix);
	97	+
	98	+} // namespace ops

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_host/confusion_matrix_tiling.cpp+384-0

@@ -0,0 +1,384 @@
	1	+/*!
	2	+ * \file confusion_matrix_tiling.cpp
	3	+ * \brief ConfusionMatrix tiling implementation
	4	+ */
	5	+
	6	+#include "register/op_def_registry.h"
	7	+#include "op_common/log/log.h"
	8	+#include "op_common/op_host/util/platform_util.h"
	9	+#include "../op_kernel/confusion_matrix_tiling_data.h"
	10	+#include "../op_kernel/confusion_matrix_tiling_key.h"
	11	+
	12	+namespace optiling {
	13	+
	14	+constexpr uint32_t WS_SYS_SIZE = 0U;
	15	+constexpr uint64_t SPLIT_OUTPUT_MIN_ELEMS_PER_CORE = 4096;
	16	+constexpr uint64_t TINY_FASTPATH_MAX_SAMPLES = 64;
	17	+constexpr uint64_t TINY_FASTPATH_MAX_OUTPUT = 256;
	18	+constexpr uint32_t SINGLE_WORKER_LAUNCH_CORES = 8;
	19	+constexpr uint32_t MAX_SPLIT_OUTPUT_CORES = 48;
	20	+constexpr uint32_t MAX_VECTOR_FILTER_SPARSE_CORES = 48;
	21	+constexpr uint64_t VECTOR_FILTER_SPARSE_MAX_SAMPLES = 10240;
	22	+constexpr uint64_t UINT32_OUTPUT_OFFSET_MAX = 0xFFFFFFFFULL;
	23	+
	24	+static const gert::Shape g_vec_1_shape = {1};
	25	+
	26	+static inline const gert::Shape EnsureNotScalar(const gert::Shape& inShape)
	27	+{
	28	+ if (inShape.GetDimNum() == 0) {
	29	+ return g_vec_1_shape;
	30	+ }
	31	+ return inShape;
	32	+}
	33	+
	34	+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
	35	+{
	36	+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
	37	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	38	+ currentWorkspace[0] = WS_SYS_SIZE;
	39	+ return ge::GRAPH_SUCCESS;
	40	+}
	41	+
	42	+static uint32_t GetVectorCoreNum(gert::TilingContext* context)
	43	+{
	44	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	45	+ if (platformInfoPtr == nullptr) {
	46	+ return 1;
	47	+ }
	48	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	49	+ const int64_t aivNum = ascendcPlatform.GetCoreNumAiv();
	50	+ const int64_t aicNum = ascendcPlatform.GetCoreNumAic();
	51	+ const int64_t coreNum = aivNum > 0 ? aivNum : aicNum;
	52	+ if (coreNum <= 0) {
	53	+ return 1;
	54	+ }
	55	+ if (coreNum > static_cast<int64_t>(MAX_SPLIT_OUTPUT_CORES)) {
	56	+ return MAX_SPLIT_OUTPUT_CORES;
	57	+ }
	58	+ return static_cast<uint32_t>(coreNum);
	59	+}
	60	+
	61	+static bool ShouldSplitOutput(int64_t totalNum, int64_t outputNum, uint32_t blockDim)
	62	+{
	63	+ if (totalNum <= 0 \|\| outputNum <= 0 \|\| blockDim <= 1) {
	64	+ return false;
	65	+ }
	66	+
	67	+ const uint64_t outputCount = static_cast<uint64_t>(outputNum);
	68	+ const uint64_t sampleCount = static_cast<uint64_t>(totalNum);
	69	+ const uint64_t coreCount = static_cast<uint64_t>(blockDim);
	70	+ const uint64_t perCoreOutputCount = (outputCount + coreCount - 1) / coreCount;
	71	+
	72	+ if (perCoreOutputCount < SPLIT_OUTPUT_MIN_ELEMS_PER_CORE) {
	73	+ return false;
	74	+ }
	75	+
	76	+ return perCoreOutputCount >= sampleCount;
	77	+}
	78	+
	79	+static bool ShouldUseVectorFilterSparse(
	80	+ int64_t totalNum,
	81	+ int64_t outputNum,
	82	+ int32_t labelsDtype,
	83	+ int32_t predictionsDtype,
	84	+ int32_t outputDtype,
	85	+ int32_t hasWeights)
	86	+{
	87	+ if (totalNum <= 0 \|\| outputNum <= 0 \|\|
	88	+ totalNum > static_cast<int64_t>(VECTOR_FILTER_SPARSE_MAX_SAMPLES) \|\|
	89	+ static_cast<uint64_t>(outputNum) > UINT32_OUTPUT_OFFSET_MAX) {
	90	+ return false;
	91	+ }
	92	+ if (labelsDtype != CONFUSION_MATRIX_DTYPE_INT32 \|\|
	93	+ predictionsDtype != CONFUSION_MATRIX_DTYPE_INT32 \|\|
	94	+ outputDtype != CONFUSION_MATRIX_DTYPE_INT32 \|\| hasWeights != 0) {
	95	+ return false;
	96	+ }
	97	+ const uint64_t sampleCount = static_cast<uint64_t>(totalNum);
	98	+ const uint64_t outputCount = static_cast<uint64_t>(outputNum);
	99	+ return outputCount / sampleCount >= sampleCount;
	100	+}
	101	+
	102	+static bool IsTinyShape(int64_t totalNum, int64_t outputNum)
	103	+{
	104	+ return totalNum > 0 &&
	105	+ totalNum <= static_cast<int64_t>(TINY_FASTPATH_MAX_SAMPLES) &&
	106	+ outputNum > 0 &&
	107	+ outputNum <= static_cast<int64_t>(TINY_FASTPATH_MAX_OUTPUT);
	108	+}
	109	+
	110	+static int32_t SelectTinyFastPathMode(
	111	+ int64_t totalNum,
	112	+ int64_t outputNum,
	113	+ int32_t labelsDtype,
	114	+ int32_t predictionsDtype,
	115	+ int32_t weightsDtype,
	116	+ int32_t outputDtype,
	117	+ int32_t hasWeights)
	118	+{
	119	+ if (!IsTinyShape(totalNum, outputNum)) {
	120	+ return CONFUSIONMATRIX_TPL_SCH_MODE_1;
	121	+ }
	122	+ if (labelsDtype == CONFUSION_MATRIX_DTYPE_INT32 &&
	123	+ predictionsDtype == CONFUSION_MATRIX_DTYPE_INT32 &&
	124	+ outputDtype == CONFUSION_MATRIX_DTYPE_INT32) {
	125	+ if (hasWeights == 0) {
	126	+ return CONFUSIONMATRIX_TPL_SCH_MODE_0;
	127	+ }
	128	+ if (weightsDtype == CONFUSION_MATRIX_DTYPE_INT32) {
	129	+ return CONFUSIONMATRIX_TPL_SCH_MODE_2;
	130	+ }
	131	+ }
	132	+ if (labelsDtype == CONFUSION_MATRIX_DTYPE_INT64 &&
	133	+ predictionsDtype == CONFUSION_MATRIX_DTYPE_INT64 &&
	134	+ weightsDtype == CONFUSION_MATRIX_DTYPE_FLOAT32 &&
	135	+ outputDtype == CONFUSION_MATRIX_DTYPE_FLOAT32 &&
	136	+ hasWeights != 0) {
	137	+ return CONFUSIONMATRIX_TPL_SCH_MODE_3;
	138	+ }
	139	+ return CONFUSIONMATRIX_TPL_SCH_MODE_1;
	140	+}
	141	+
	142	+static uint32_t GetSingleWorkerLaunchCoreNum(uint32_t availableCoreNum)
	143	+{
	144	+ if (availableCoreNum == 0) {
	145	+ return 1;
	146	+ }
	147	+ if (availableCoreNum > SINGLE_WORKER_LAUNCH_CORES) {
	148	+ return SINGLE_WORKER_LAUNCH_CORES;
	149	+ }
	150	+ return availableCoreNum;
	151	+}
	152	+
	153	+static ge::graphStatus ConvertDtype(gert::TilingContext* context, ge::DataType dtype, int32_t& dtypeCode)
	154	+{
	155	+ if (dtype == ge::DT_INT32) {
	156	+ dtypeCode = CONFUSION_MATRIX_DTYPE_INT32;
	157	+ return ge::GRAPH_SUCCESS;
	158	+ }
	159	+ if (dtype == ge::DT_INT64) {
	160	+ dtypeCode = CONFUSION_MATRIX_DTYPE_INT64;
	161	+ return ge::GRAPH_SUCCESS;
	162	+ }
	163	+ if (dtype == ge::DT_FLOAT) {
	164	+ dtypeCode = CONFUSION_MATRIX_DTYPE_FLOAT32;
	165	+ return ge::GRAPH_SUCCESS;
	166	+ }
	167	+ OP_LOGE(context, "ConfusionMatrix unsupported dtype: %d", static_cast<int32_t>(dtype));
	168	+ return ge::GRAPH_FAILED;
	169	+}
	170	+
	171	+static ge::graphStatus GetTensorNum(gert::TilingContext* context, const gert::StorageShape* storageShape, int64_t& num)
	172	+{
	173	+ OP_CHECK_NULL_WITH_CONTEXT(context, storageShape);
	174	+ const gert::Shape shape = storageShape->GetStorageShape();
	175	+ if (shape.GetDimNum() == 0) {
	176	+ num = 1;
	177	+ return ge::GRAPH_SUCCESS;
	178	+ }
	179	+ num = 1;
	180	+ for (size_t i = 0; i < shape.GetDimNum(); ++i) {
	181	+ const int64_t dim = shape.GetDim(i);
	182	+ OP_CHECK_IF(dim <= 0, OP_LOGE(context, "input dims must be positive"), return ge::GRAPH_FAILED);
	183	+ num *= dim;
	184	+ }
	185	+ return ge::GRAPH_SUCCESS;
	186	+}
	187	+
	188	+static bool IsOptionalWeightsPresent(gert::TilingContext* context, int64_t totalNum, int64_t& weightsNum)
	189	+{
	190	+ const auto weightsDesc = context->GetInputDesc(2);
	191	+ const gert::StorageShape* weightsStorageShape = context->GetInputShape(2);
	192	+ if (weightsDesc == nullptr \|\| weightsStorageShape == nullptr \|\| weightsDesc->GetDataType() == ge::DT_UNDEFINED) {
	193	+ return false;
	194	+ }
	195	+ const gert::Shape weightsShape = weightsStorageShape->GetStorageShape();
	196	+ if (weightsShape.GetDimNum() == 0 && totalNum != 1) {
	197	+ return false;
	198	+ }
	199	+ return GetTensorNum(context, weightsStorageShape, weightsNum) == ge::GRAPH_SUCCESS;
	200	+}
	201	+
	202	+static int64_t GetNumClassesFromAttr(gert::TilingContext* context)
	203	+{
	204	+ const gert::RuntimeAttrs* attrs = context->GetAttrs();
	205	+ if (attrs == nullptr) {
	206	+ return -1;
	207	+ }
	208	+ const int64_t* numClassesAttr = attrs->GetInt(0);
	209	+ if (numClassesAttr == nullptr) {
	210	+ return -1;
	211	+ }
	212	+ return *numClassesAttr;
	213	+}
	214	+
	215	+static int64_t GetNumClassesFromOutput(gert::TilingContext* context)
	216	+{
	217	+ const gert::StorageShape* yStorageShape = context->GetOutputShape(0);
	218	+ if (yStorageShape == nullptr) {
	219	+ return -1;
	220	+ }
	221	+ const gert::Shape yShape = EnsureNotScalar(yStorageShape->GetStorageShape());
	222	+ if (yShape.GetDimNum() != 2) {
	223	+ return -1;
	224	+ }
	225	+ const int64_t rows = yShape.GetDim(0);
	226	+ const int64_t cols = yShape.GetDim(1);
	227	+ if (rows <= 0 \|\| rows != cols) {
	228	+ return -1;
	229	+ }
	230	+ return rows;
	231	+}
	232	+
	233	+static ge::graphStatus ConfusionMatrixTilingFunc(gert::TilingContext* context)
	234	+{
	235	+ OP_CHECK_IF(
	236	+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
	237	+ OP_LOGE(context, "GetWorkspaceSize error"),
	238	+ return ge::GRAPH_FAILED);
	239	+
	240	+ int64_t totalNum = 0;
	241	+ OP_CHECK_IF(
	242	+ GetTensorNum(context, context->GetInputShape(0), totalNum) != ge::GRAPH_SUCCESS,
	243	+ OP_LOGE(context, "invalid labels shape"),
	244	+ return ge::GRAPH_FAILED);
	245	+
	246	+ int64_t predictionsNum = 0;
	247	+ OP_CHECK_IF(
	248	+ GetTensorNum(context, context->GetInputShape(1), predictionsNum) != ge::GRAPH_SUCCESS,
	249	+ OP_LOGE(context, "invalid predictions shape"),
	250	+ return ge::GRAPH_FAILED);
	251	+ OP_CHECK_IF(predictionsNum != totalNum, OP_LOGE(context, "labels and predictions size mismatch"), return ge::GRAPH_FAILED);
	252	+
	253	+ const auto labelsDesc = context->GetInputDesc(0);
	254	+ const auto predictionsDesc = context->GetInputDesc(1);
	255	+ OP_CHECK_NULL_WITH_CONTEXT(context, labelsDesc);
	256	+ OP_CHECK_NULL_WITH_CONTEXT(context, predictionsDesc);
	257	+
	258	+ int32_t labelsDtype = CONFUSION_MATRIX_DTYPE_INT32;
	259	+ int32_t predictionsDtype = CONFUSION_MATRIX_DTYPE_INT32;
	260	+ OP_CHECK_IF(
	261	+ ConvertDtype(context, labelsDesc->GetDataType(), labelsDtype) != ge::GRAPH_SUCCESS,
	262	+ OP_LOGE(context, "invalid labels dtype"),
	263	+ return ge::GRAPH_FAILED);
	264	+ OP_CHECK_IF(
	265	+ ConvertDtype(context, predictionsDesc->GetDataType(), predictionsDtype) != ge::GRAPH_SUCCESS,
	266	+ OP_LOGE(context, "invalid predictions dtype"),
	267	+ return ge::GRAPH_FAILED);
	268	+ OP_CHECK_IF(
	269	+ labelsDtype == CONFUSION_MATRIX_DTYPE_FLOAT32 \|\| predictionsDtype == CONFUSION_MATRIX_DTYPE_FLOAT32,
	270	+ OP_LOGE(context, "labels and predictions must be int32 or int64"),
	271	+ return ge::GRAPH_FAILED);
	272	+
	273	+ int32_t hasWeights = 0;
	274	+ int32_t weightsDtype = CONFUSION_MATRIX_DTYPE_INT32;
	275	+ int64_t weightsNum = 0;
	276	+ const auto weightsDesc = context->GetInputDesc(2);
	277	+ if (IsOptionalWeightsPresent(context, totalNum, weightsNum)) {
	278	+ hasWeights = 1;
	279	+ OP_CHECK_IF(
	280	+ ConvertDtype(context, weightsDesc->GetDataType(), weightsDtype) != ge::GRAPH_SUCCESS,
	281	+ OP_LOGE(context, "invalid weights dtype"),
	282	+ return ge::GRAPH_FAILED);
	283	+ OP_CHECK_IF(weightsNum != totalNum, OP_LOGE(context, "weights size mismatch"), return ge::GRAPH_FAILED);
	284	+ }
	285	+
	286	+ int64_t numClasses = GetNumClassesFromAttr(context);
	287	+ if (numClasses <= 0) {
	288	+ numClasses = GetNumClassesFromOutput(context);
	289	+ }
	290	+ OP_CHECK_IF(numClasses <= 0, OP_LOGE(context, "num_classes must be positive"), return ge::GRAPH_FAILED);
	291	+
	292	+ const gert::StorageShape* yStorageShape = context->GetOutputShape(0);
	293	+ OP_CHECK_NULL_WITH_CONTEXT(context, yStorageShape);
	294	+ const gert::Shape yShape = EnsureNotScalar(yStorageShape->GetStorageShape());
	295	+ OP_CHECK_IF(yShape.GetDimNum() != 2, OP_LOGE(context, "output must be 2-D"), return ge::GRAPH_FAILED);
	296	+ OP_CHECK_IF(
	297	+ yShape.GetDim(0) != numClasses \|\| yShape.GetDim(1) != numClasses,
	298	+ OP_LOGE(context, "output shape must be [num_classes, num_classes]"),
	299	+ return ge::GRAPH_FAILED);
	300	+
	301	+ int32_t outputDtype = CONFUSION_MATRIX_DTYPE_INT32;
	302	+ const auto outputDesc = context->GetOutputDesc(0);
	303	+ if (outputDesc != nullptr) {
	304	+ OP_CHECK_IF(
	305	+ ConvertDtype(context, outputDesc->GetDataType(), outputDtype) != ge::GRAPH_SUCCESS,
	306	+ OP_LOGE(context, "invalid output dtype"),
	307	+ return ge::GRAPH_FAILED);
	308	+ }
	309	+
	310	+ ConfusionMatrixTilingData* tiling = context->GetTilingData<ConfusionMatrixTilingData>();
	311	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	312	+ const int64_t outputNum = numClasses * numClasses;
	313	+ const uint32_t splitOutputBlockDim = GetVectorCoreNum(context);
	314	+ uint32_t blockDim = GetSingleWorkerLaunchCoreNum(splitOutputBlockDim);
	315	+ int64_t blockFactor = totalNum;
	316	+ int32_t reserved = 0;
	317	+ int32_t schMode = CONFUSIONMATRIX_TPL_SCH_MODE_1;
	318	+ if (ShouldSplitOutput(totalNum, outputNum, splitOutputBlockDim)) {
	319	+ blockDim = splitOutputBlockDim;
	320	+ blockFactor = (numClasses + static_cast<int64_t>(blockDim) - 1) / static_cast<int64_t>(blockDim);
	321	+ reserved \|= CONFUSION_MATRIX_RESERVED_SPLIT_OUTPUT;
	322	+ if (ShouldUseVectorFilterSparse(
	323	+ totalNum, outputNum, labelsDtype, predictionsDtype, outputDtype, hasWeights)) {
	324	+ if (blockDim > MAX_VECTOR_FILTER_SPARSE_CORES) {
	325	+ blockDim = MAX_VECTOR_FILTER_SPARSE_CORES;
	326	+ blockFactor =
	327	+ (numClasses + static_cast<int64_t>(blockDim) - 1) / static_cast<int64_t>(blockDim);
	328	+ }
	329	+ reserved \|= CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE;
	330	+ schMode = CONFUSIONMATRIX_TPL_SCH_MODE_6;
	331	+ }
	332	+ } else {
	333	+ schMode = SelectTinyFastPathMode(
	334	+ totalNum, outputNum, labelsDtype, predictionsDtype, weightsDtype, outputDtype, hasWeights);
	335	+ if (labelsDtype == CONFUSION_MATRIX_DTYPE_INT32 &&
	336	+ predictionsDtype == CONFUSION_MATRIX_DTYPE_INT32 &&
	337	+ outputDtype == CONFUSION_MATRIX_DTYPE_INT32 && hasWeights == 0 && numClasses == 3 && totalNum >= 512) {
	338	+ schMode = (static_cast<uint64_t>(totalNum) & 63UL) == 56UL ?
	339	+ CONFUSIONMATRIX_TPL_SCH_MODE_5 : CONFUSIONMATRIX_TPL_SCH_MODE_4;
	340	+ blockFactor = totalNum;
	341	+ }
	342	+ }
	343	+ tiling->totalNum = totalNum;
	344	+ tiling->blockFactor = blockFactor;
	345	+ tiling->numClasses = numClasses;
	346	+ tiling->labelsDtype = labelsDtype;
	347	+ tiling->predictionsDtype = predictionsDtype;
	348	+ tiling->weightsDtype = weightsDtype;
	349	+ tiling->outputDtype = outputDtype;
	350	+ tiling->hasWeights = hasWeights;
	351	+ tiling->reserved = reserved;
	352	+
	353	+ context->SetBlockDim(blockDim);
	354	+ if (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_6) {
	355	+ context->SetScheduleMode(1);
	356	+ }
	357	+ if (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_0) {
	358	+ context->SetTilingKey(GET_TPL_TILING_KEY(CONFUSIONMATRIX_TPL_SCH_MODE_0));
	359	+ } else if (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_2) {
	360	+ context->SetTilingKey(GET_TPL_TILING_KEY(CONFUSIONMATRIX_TPL_SCH_MODE_2));
	361	+ } else if (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_3) {
	362	+ context->SetTilingKey(GET_TPL_TILING_KEY(CONFUSIONMATRIX_TPL_SCH_MODE_3));
	363	+ } else if (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_4) {
	364	+ context->SetTilingKey(GET_TPL_TILING_KEY(CONFUSIONMATRIX_TPL_SCH_MODE_4));
	365	+ } else if (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_5) {
	366	+ context->SetTilingKey(GET_TPL_TILING_KEY(CONFUSIONMATRIX_TPL_SCH_MODE_5));
	367	+ } else if (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_6) {
	368	+ context->SetTilingKey(GET_TPL_TILING_KEY(CONFUSIONMATRIX_TPL_SCH_MODE_6));
	369	+ } else {
	370	+ context->SetTilingKey(GET_TPL_TILING_KEY(CONFUSIONMATRIX_TPL_SCH_MODE_1));
	371	+ }
	372	+ return ge::GRAPH_SUCCESS;
	373	+}
	374	+
	375	+static ge::graphStatus TilingParseForConfusionMatrix([[maybe_unused]] gert::TilingParseContext* context)
	376	+{
	377	+ return ge::GRAPH_SUCCESS;
	378	+}
	379	+
	380	+struct ConfusionMatrixCompileInfo {};
	381	+
	382	+IMPL_OP_OPTILING(ConfusionMatrix).Tiling(ConfusionMatrixTilingFunc).TilingParse<ConfusionMatrixCompileInfo>(TilingParseForConfusionMatrix);
	383	+
	384	+} // namespace optiling

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_kernel/CMakeLists.txt+17-0

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

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_kernel/confusion_matrix.cpp+69-0

@@ -0,0 +1,69 @@
	1	+/*!
	2	+ * \file confusion_matrix.cpp
	3	+ * \brief ConfusionMatrix 算子 kernel 入口
	4	+ */
	5	+
	6	+#include "confusion_matrix.h"
	7	+
	8	+template <uint32_t schMode>
	9	+__global__ __aicore__ void confusion_matrix(GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
	10	+{
	11	+ // The zero-argument SyncAll in tiling key 6 needs the mixed AIV task
	12	+ // metadata. Keep every other key on the lean AIV-only launch path.
	13	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	14	+ KERNEL_TASK_TYPE(6, KERNEL_TYPE_MIX_AIV_1_0);
	15	+ REGISTER_TILING_DEFAULT(ConfusionMatrixTilingData);
	16	+ if constexpr (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_0) {
	17	+ if (AscendC::GetBlockIdx() != 0) {
	18	+ return;
	19	+ }
	20	+ GET_TILING_DATA_WITH_STRUCT(ConfusionMatrixTilingData, tilingData, tiling);
	21	+ NsConfusionMatrix::ProcessTinyInt32NoWeight(labels, predictions, y, &tilingData);
	22	+ return;
	23	+ }
	24	+ if constexpr (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_2) {
	25	+ if (AscendC::GetBlockIdx() != 0) {
	26	+ return;
	27	+ }
	28	+ GET_TILING_DATA_WITH_STRUCT(ConfusionMatrixTilingData, tilingData, tiling);
	29	+ NsConfusionMatrix::ProcessTinyInt32WeightInt32(labels, predictions, weights, y, &tilingData);
	30	+ return;
	31	+ }
	32	+ if constexpr (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_3) {
	33	+ if (AscendC::GetBlockIdx() != 0) {
	34	+ return;
	35	+ }
	36	+ GET_TILING_DATA_WITH_STRUCT(ConfusionMatrixTilingData, tilingData, tiling);
	37	+ NsConfusionMatrix::ProcessTinyInt64WeightFloat(labels, predictions, weights, y, &tilingData);
	38	+ return;
	39	+ }
	40	+ if constexpr (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_4) {
	41	+ if (AscendC::GetBlockIdx() >= 3) {
	42	+ return;
	43	+ }
	44	+ GET_TILING_DATA_WITH_STRUCT(ConfusionMatrixTilingData, tilingData, tiling);
	45	+ NsConfusionMatrix::ProcessC3VectorHistogram<false>(labels, predictions, y, &tilingData);
	46	+ return;
	47	+ }
	48	+ if constexpr (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_5) {
	49	+ if (AscendC::GetBlockIdx() >= 3) {
	50	+ return;
	51	+ }
	52	+ GET_TILING_DATA_WITH_STRUCT(ConfusionMatrixTilingData, tilingData, tiling);
	53	+ NsConfusionMatrix::ProcessC3VectorHistogram<true>(labels, predictions, y, &tilingData);
	54	+ return;
	55	+ }
	56	+#if TILING_KEY_VAR == CONFUSIONMATRIX_TPL_SCH_MODE_6
	57	+ if constexpr (schMode == CONFUSIONMATRIX_TPL_SCH_MODE_6) {
	58	+ GET_TILING_DATA_WITH_STRUCT(ConfusionMatrixTilingData, tilingData, tiling);
	59	+ NsConfusionMatrix::ConfusionMatrix op;
	60	+ op.Init(labels, predictions, weights, y, &tilingData);
	61	+ op.ProcessAtomicSparse();
	62	+ return;
	63	+ }
	64	+#endif
	65	+ GET_TILING_DATA_WITH_STRUCT(ConfusionMatrixTilingData, tilingData, tiling);
	66	+ NsConfusionMatrix::ConfusionMatrix op;
	67	+ op.Init(labels, predictions, weights, y, &tilingData);
	68	+ op.Process();
	69	+}

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_kernel/confusion_matrix.h+1177-0

@@ -0,0 +1,1177 @@
	1	+/*!
	2	+ * \file confusion_matrix.h
	3	+ * \brief ConfusionMatrix kernel implementation
	4	+ */
	5	+
	6	+#ifndef CONFUSIONMATRIX_H
	7	+#define CONFUSIONMATRIX_H
	8	+
	9	+#include "kernel_operator.h"
	10	+#include "kernel_tiling/kernel_tiling.h"
	11	+#include "confusion_matrix_tiling_data.h"
	12	+#include "confusion_matrix_tiling_key.h"
	13	+
	14	+namespace NsConfusionMatrix {
	15	+
	16	+using namespace AscendC;
	17	+
	18	+constexpr uint32_t CONFUSION_MATRIX_ZERO_TILE_ELEMS = 16384;
	19	+constexpr uint64_t CONFUSION_MATRIX_BULK_CLEAR_MIN_ELEMS = 1024;
	20	+constexpr uint64_t CONFUSION_MATRIX_TINY_FASTPATH_MAX_SAMPLES = 64;
	21	+constexpr uint64_t CONFUSION_MATRIX_TINY_FASTPATH_MAX_OUTPUT = 256;
	22	+constexpr uint64_t CONFUSION_MATRIX_SMALL_ROW_SPLIT_MIN_SAMPLES = 512;
	23	+constexpr uint64_t CONFUSION_MATRIX_SMALL_ROW_SPLIT_MAX_CLASSES = 16;
	24	+constexpr uint32_t CONFUSION_MATRIX_C3_VECTOR_TILE_ELEMS = 3072;
	25	+constexpr uint32_t CONFUSION_MATRIX_C3_VECTOR_DATA_BUFFER_BYTES =
	26	+ CONFUSION_MATRIX_C3_VECTOR_TILE_ELEMS * sizeof(int32_t) * 2;
	27	+constexpr uint32_t CONFUSION_MATRIX_C3_VECTOR_MASK_BYTES = CONFUSION_MATRIX_C3_VECTOR_TILE_ELEMS / 8;
	28	+constexpr uint32_t CONFUSION_MATRIX_C3_VECTOR_MASK_BUFFER_BYTES = CONFUSION_MATRIX_C3_VECTOR_MASK_BYTES * 3;
	29	+constexpr uint32_t CONFUSION_MATRIX_C3_VECTOR_WORK_BUFFER_BYTES =
	30	+ CONFUSION_MATRIX_C3_VECTOR_DATA_BUFFER_BYTES + CONFUSION_MATRIX_C3_VECTOR_MASK_BUFFER_BYTES;
	31	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_CAPACITY = 10240;
	32	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_COUNT_LOCAL_OFFSET = CONFUSION_MATRIX_SPARSE_CAPACITY;
	33	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_OCCUPIED_LOCAL_OFFSET = CONFUSION_MATRIX_SPARSE_CAPACITY * 2;
	34	+constexpr uint32_t CONFUSION_MATRIX_BULK_BUFFER_BYTES =
	35	+ CONFUSION_MATRIX_SPARSE_CAPACITY * sizeof(int32_t) * 3;
	36	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_FILTER_TILE_ELEMS = 5120;
	37	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_FILTER_LABEL_LOCAL_OFFSET = CONFUSION_MATRIX_SPARSE_CAPACITY * 3;
	38	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_FILTER_PREDICTION_LOCAL_OFFSET =
	39	+ CONFUSION_MATRIX_SPARSE_FILTER_LABEL_LOCAL_OFFSET + CONFUSION_MATRIX_SPARSE_FILTER_TILE_ELEMS;
	40	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_FILTER_COMPACT_LABEL_LOCAL_OFFSET =
	41	+ CONFUSION_MATRIX_SPARSE_FILTER_PREDICTION_LOCAL_OFFSET + CONFUSION_MATRIX_SPARSE_FILTER_TILE_ELEMS;
	42	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_FILTER_MASK_BYTES = CONFUSION_MATRIX_SPARSE_FILTER_TILE_ELEMS / 8;
	43	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_FILTER_MASK_LOCAL_OFFSET_BYTES =
	44	+ (CONFUSION_MATRIX_SPARSE_FILTER_COMPACT_LABEL_LOCAL_OFFSET + CONFUSION_MATRIX_SPARSE_FILTER_TILE_ELEMS) *
	45	+ sizeof(int32_t);
	46	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_FILTER_STORAGE_BYTES =
	47	+ CONFUSION_MATRIX_SPARSE_FILTER_MASK_LOCAL_OFFSET_BYTES + CONFUSION_MATRIX_SPARSE_FILTER_MASK_BYTES * 3;
	48	+constexpr uint32_t CONFUSION_MATRIX_SPARSE_FILTER_BUFFER_BYTES = 184U * 1024U;
	49	+constexpr uint64_t CONFUSION_MATRIX_SPARSE_NORMAL_L2_BUDGET_BYTES = 192UL * 1024UL * 1024UL;
	50	+constexpr uint64_t CONFUSION_MATRIX_UINT32_OFFSET_MAX = 0xFFFFFFFFULL;
	51	+
	52	+static_assert((CONFUSION_MATRIX_SPARSE_FILTER_MASK_LOCAL_OFFSET_BYTES & 31U) == 0U,
	53	+ "sparse vector-filter mask storage must be 32-byte aligned");
	54	+static_assert(CONFUSION_MATRIX_SPARSE_FILTER_STORAGE_BYTES <= CONFUSION_MATRIX_SPARSE_FILTER_BUFFER_BYTES,
	55	+ "sparse vector-filter storage exceeds the Vector Core UB budget");
	56	+
	57	+__aicore__ inline void ProcessTinyInt32NoWeight(
	58	+ GM_ADDR labels,
	59	+ GM_ADDR predictions,
	60	+ GM_ADDR y,
	61	+ const ConfusionMatrixTilingData* tilingData)
	62	+{
	63	+ GlobalTensor<int32_t> labelsGm;
	64	+ GlobalTensor<int32_t> predictionsGm;
	65	+ GlobalTensor<int32_t> outputGm;
	66	+ labelsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(labels));
	67	+ predictionsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(predictions));
	68	+ outputGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(y));
	69	+
	70	+ const uint32_t totalNum = static_cast<uint32_t>(tilingData->totalNum);
	71	+ const uint32_t numClasses = static_cast<uint32_t>(tilingData->numClasses);
	72	+ const uint32_t outputNum = numClasses * numClasses;
	73	+
	74	+ for (uint32_t i = 0; i < outputNum; ++i) {
	75	+ outputGm.SetValue(i, static_cast<int32_t>(0));
	76	+ }
	77	+
	78	+ if (totalNum <= 6) {
	79	+#define CM_TINY_I32_NOWEIGHT_STEP(IDX) \
	80	+ if (totalNum > IDX) { \
	81	+ const uint32_t label = static_cast<uint32_t>(labelsGm.GetValue(IDX)); \
	82	+ const uint32_t prediction = static_cast<uint32_t>(predictionsGm.GetValue(IDX)); \
	83	+ const uint32_t outOffset = label * numClasses + prediction; \
	84	+ const int32_t oldValue = outputGm.GetValue(outOffset); \
	85	+ outputGm.SetValue(outOffset, oldValue + 1); \
	86	+ }
	87	+ CM_TINY_I32_NOWEIGHT_STEP(0);
	88	+ CM_TINY_I32_NOWEIGHT_STEP(1);
	89	+ CM_TINY_I32_NOWEIGHT_STEP(2);
	90	+ CM_TINY_I32_NOWEIGHT_STEP(3);
	91	+ CM_TINY_I32_NOWEIGHT_STEP(4);
	92	+ CM_TINY_I32_NOWEIGHT_STEP(5);
	93	+#undef CM_TINY_I32_NOWEIGHT_STEP
	94	+ return;
	95	+ }
	96	+
	97	+ for (uint32_t i = 0; i < totalNum; ++i) {
	98	+ const uint32_t label = static_cast<uint32_t>(labelsGm.GetValue(i));
	99	+ const uint32_t prediction = static_cast<uint32_t>(predictionsGm.GetValue(i));
	100	+ const uint32_t outOffset = label * numClasses + prediction;
	101	+ const int32_t oldValue = outputGm.GetValue(outOffset);
	102	+ outputGm.SetValue(outOffset, oldValue + 1);
	103	+ }
	104	+}
	105	+
	106	+__aicore__ inline void ProcessTinyInt32WeightInt32(
	107	+ GM_ADDR labels,
	108	+ GM_ADDR predictions,
	109	+ GM_ADDR weights,
	110	+ GM_ADDR y,
	111	+ const ConfusionMatrixTilingData* tilingData)
	112	+{
	113	+ GlobalTensor<int32_t> labelsGm;
	114	+ GlobalTensor<int32_t> predictionsGm;
	115	+ GlobalTensor<int32_t> weightsGm;
	116	+ GlobalTensor<int32_t> outputGm;
	117	+ labelsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(labels));
	118	+ predictionsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(predictions));
	119	+ weightsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(weights));
	120	+ outputGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(y));
	121	+
	122	+ const uint32_t totalNum = static_cast<uint32_t>(tilingData->totalNum);
	123	+ const uint32_t numClasses = static_cast<uint32_t>(tilingData->numClasses);
	124	+ const uint32_t outputNum = numClasses * numClasses;
	125	+
	126	+ for (uint32_t i = 0; i < outputNum; ++i) {
	127	+ outputGm.SetValue(i, static_cast<int32_t>(0));
	128	+ }
	129	+
	130	+ if (totalNum <= 6) {
	131	+#define CM_TINY_I32_WEIGHT_STEP(IDX) \
	132	+ if (totalNum > IDX) { \
	133	+ const uint32_t label = static_cast<uint32_t>(labelsGm.GetValue(IDX)); \
	134	+ const uint32_t prediction = static_cast<uint32_t>(predictionsGm.GetValue(IDX)); \
	135	+ const uint32_t outOffset = label * numClasses + prediction; \
	136	+ const int32_t oldValue = outputGm.GetValue(outOffset); \
	137	+ outputGm.SetValue(outOffset, oldValue + weightsGm.GetValue(IDX)); \
	138	+ }
	139	+ CM_TINY_I32_WEIGHT_STEP(0);
	140	+ CM_TINY_I32_WEIGHT_STEP(1);
	141	+ CM_TINY_I32_WEIGHT_STEP(2);
	142	+ CM_TINY_I32_WEIGHT_STEP(3);
	143	+ CM_TINY_I32_WEIGHT_STEP(4);
	144	+ CM_TINY_I32_WEIGHT_STEP(5);
	145	+#undef CM_TINY_I32_WEIGHT_STEP
	146	+ return;
	147	+ }
	148	+
	149	+ for (uint32_t i = 0; i < totalNum; ++i) {
	150	+ const uint32_t label = static_cast<uint32_t>(labelsGm.GetValue(i));
	151	+ const uint32_t prediction = static_cast<uint32_t>(predictionsGm.GetValue(i));
	152	+ const uint32_t outOffset = label * numClasses + prediction;
	153	+ const int32_t oldValue = outputGm.GetValue(outOffset);
	154	+ outputGm.SetValue(outOffset, oldValue + weightsGm.GetValue(i));
	155	+ }
	156	+}
	157	+
	158	+__aicore__ inline void ProcessTinyInt64WeightFloat(
	159	+ GM_ADDR labels,
	160	+ GM_ADDR predictions,
	161	+ GM_ADDR weights,
	162	+ GM_ADDR y,
	163	+ const ConfusionMatrixTilingData* tilingData)
	164	+{
	165	+ GlobalTensor<int64_t> labelsGm;
	166	+ GlobalTensor<int64_t> predictionsGm;
	167	+ GlobalTensor<float> weightsGm;
	168	+ GlobalTensor<float> outputGm;
	169	+ labelsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(labels));
	170	+ predictionsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(predictions));
	171	+ weightsGm.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(weights));
	172	+ outputGm.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(y));
	173	+
	174	+ const uint64_t totalNum = static_cast<uint64_t>(tilingData->totalNum);
	175	+ const uint64_t numClasses = static_cast<uint64_t>(tilingData->numClasses);
	176	+ const uint64_t outputNum = numClasses * numClasses;
	177	+
	178	+ for (uint64_t i = 0; i < outputNum; ++i) {
	179	+ outputGm.SetValue(i, 0.0f);
	180	+ }
	181	+
	182	+ if (totalNum <= 6) {
	183	+#define CM_TINY_I64_FLOAT_STEP(IDX) \
	184	+ if (totalNum > IDX) { \
	185	+ const uint64_t label = static_cast<uint64_t>(labelsGm.GetValue(IDX)); \
	186	+ const uint64_t prediction = static_cast<uint64_t>(predictionsGm.GetValue(IDX)); \
	187	+ const uint64_t outOffset = label * numClasses + prediction; \
	188	+ const float oldValue = outputGm.GetValue(outOffset); \
	189	+ outputGm.SetValue(outOffset, oldValue + weightsGm.GetValue(IDX)); \
	190	+ }
	191	+ CM_TINY_I64_FLOAT_STEP(0);
	192	+ CM_TINY_I64_FLOAT_STEP(1);
	193	+ CM_TINY_I64_FLOAT_STEP(2);
	194	+ CM_TINY_I64_FLOAT_STEP(3);
	195	+ CM_TINY_I64_FLOAT_STEP(4);
	196	+ CM_TINY_I64_FLOAT_STEP(5);
	197	+#undef CM_TINY_I64_FLOAT_STEP
	198	+ return;
	199	+ }
	200	+
	201	+ for (uint64_t i = 0; i < totalNum; ++i) {
	202	+ const uint64_t label = static_cast<uint64_t>(labelsGm.GetValue(i));
	203	+ const uint64_t prediction = static_cast<uint64_t>(predictionsGm.GetValue(i));
	204	+ const uint64_t outOffset = label * numClasses + prediction;
	205	+ const float oldValue = outputGm.GetValue(outOffset);
	206	+ outputGm.SetValue(outOffset, oldValue + weightsGm.GetValue(i));
	207	+ }
	208	+}
	209	+
	210	+template <bool useMte2TailPadding>
	211	+__aicore__ inline void ProcessC3VectorHistogram(
	212	+ GM_ADDR labels,
	213	+ GM_ADDR predictions,
	214	+ GM_ADDR y,
	215	+ const ConfusionMatrixTilingData* tilingData);
	216	+
	217	+class ConfusionMatrix {
	218	+public:
	219	+ __aicore__ inline ConfusionMatrix() {};
	220	+
	221	+ __aicore__ inline void Init(
	222	+ GM_ADDR labels,
	223	+ GM_ADDR predictions,
	224	+ GM_ADDR weights,
	225	+ GM_ADDR y,
	226	+ const ConfusionMatrixTilingData* tilingData);
	227	+ __aicore__ inline void Process();
	228	+#if TILING_KEY_VAR == CONFUSIONMATRIX_TPL_SCH_MODE_6
	229	+ __aicore__ inline void ProcessAtomicSparse();
	230	+#endif
	231	+
	232	+private:
	233	+ __aicore__ inline int64_t ReadLabel(uint64_t offset);
	234	+ __aicore__ inline int64_t ReadPrediction(uint64_t offset);
	235	+ __aicore__ inline int32_t ReadWeightAsInt32(uint64_t offset);
	236	+ __aicore__ inline int64_t ReadWeightAsInt64(uint64_t offset);
	237	+ __aicore__ inline float ReadWeightAsFloat(uint64_t offset);
	238	+ __aicore__ inline void ClearOutput();
	239	+ __aicore__ inline void ClearOutputRange(uint64_t start, uint64_t end);
	240	+ __aicore__ inline void ClearOutputRangeInt32(uint64_t start, uint64_t end);
	241	+ __aicore__ inline void ClearOutputRangeFloat(uint64_t start, uint64_t end);
	242	+ __aicore__ inline void AddToOutput(uint64_t offset, uint64_t sampleIdx);
	243	+ __aicore__ inline void ProcessSingleCore();
	244	+ __aicore__ inline void ProcessSplitOutput();
	245	+ __aicore__ inline bool CanUseSmallOutputRowSplit(uint64_t totalNum) const;
	246	+ __aicore__ inline void ProcessSmallOutputRowSplit(uint64_t totalNum);
	247	+ __aicore__ inline bool CanUseSparseSplitInt32(uint64_t totalNum) const;
	248	+ __aicore__ inline void ProcessSplitOutputSparseInt32(
	249	+ uint64_t rowStart,
	250	+ uint64_t rowEnd,
	251	+ uint64_t classCount,
	252	+ uint64_t totalNum);
	253	+
	254	+private:
	255	+ GlobalTensor<int32_t> labelsI32Gm_;
	256	+ GlobalTensor<int64_t> labelsI64Gm_;
	257	+ GlobalTensor<int32_t> predictionsI32Gm_;
	258	+ GlobalTensor<int64_t> predictionsI64Gm_;
	259	+ GlobalTensor<int32_t> weightsI32Gm_;
	260	+ GlobalTensor<int64_t> weightsI64Gm_;
	261	+ GlobalTensor<float> weightsF32Gm_;
	262	+ GlobalTensor<int32_t> outputI32Gm_;
	263	+ GlobalTensor<int32_t> outputI32BypassGm_;
	264	+ GlobalTensor<int64_t> outputI64Gm_;
	265	+ GlobalTensor<float> outputF32Gm_;
	266	+ TPipe pipe_;
	267	+ TBuf<TPosition::VECOUT> zeroBuf_;
	268	+
	269	+ int64_t totalNum_ = 0;
	270	+ int64_t blockFactor_ = 1;
	271	+ int64_t numClasses_ = 0;
	272	+ int64_t outputNum_ = 0;
	273	+ int32_t labelsDtype_ = CONFUSION_MATRIX_DTYPE_INT32;
	274	+ int32_t predictionsDtype_ = CONFUSION_MATRIX_DTYPE_INT32;
	275	+ int32_t weightsDtype_ = CONFUSION_MATRIX_DTYPE_INT32;
	276	+ int32_t outputDtype_ = CONFUSION_MATRIX_DTYPE_INT32;
	277	+ int32_t hasWeights_ = 0;
	278	+ int32_t reserved_ = 0;
	279	+ int32_t useBulkClear_ = 0;
	280	+ uint64_t sparseNormalStart_ = 0;
	281	+};
	282	+
	283	+__aicore__ inline void ConfusionMatrix::Init(
	284	+ GM_ADDR labels,
	285	+ GM_ADDR predictions,
	286	+ GM_ADDR weights,
	287	+ GM_ADDR y,
	288	+ const ConfusionMatrixTilingData* tilingData)
	289	+{
	290	+ labelsI32Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(labels));
	291	+ labelsI64Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(labels));
	292	+ predictionsI32Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(predictions));
	293	+ predictionsI64Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(predictions));
	294	+ outputI32Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(y));
	295	+ outputI32BypassGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(y));
	296	+ outputI64Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(y));
	297	+ outputF32Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(y));
	298	+
	299	+ totalNum_ = tilingData->totalNum;
	300	+ blockFactor_ = tilingData->blockFactor;
	301	+ numClasses_ = tilingData->numClasses;
	302	+ outputNum_ = numClasses_ * numClasses_;
	303	+ labelsDtype_ = tilingData->labelsDtype;
	304	+ predictionsDtype_ = tilingData->predictionsDtype;
	305	+ weightsDtype_ = tilingData->weightsDtype;
	306	+ outputDtype_ = tilingData->outputDtype;
	307	+ hasWeights_ = tilingData->hasWeights;
	308	+ reserved_ = tilingData->reserved;
	309	+ if ((reserved_ & CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE) != 0) {
	310	+ outputI32BypassGm_.SetL2CacheHint<CacheRwMode::WRITE>(CacheMode::CACHE_MODE_DISABLE);
	311	+ }
	312	+ if ((outputDtype_ == CONFUSION_MATRIX_DTYPE_INT32 \|\| outputDtype_ == CONFUSION_MATRIX_DTYPE_FLOAT32) &&
	313	+ outputNum_ >= static_cast<int64_t>(CONFUSION_MATRIX_BULK_CLEAR_MIN_ELEMS)) {
	314	+ useBulkClear_ = 1;
	315	+ const uint32_t bufferBytes =
	316	+ (reserved_ & CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE) != 0 ?
	317	+ CONFUSION_MATRIX_SPARSE_FILTER_BUFFER_BYTES : CONFUSION_MATRIX_BULK_BUFFER_BYTES;
	318	+ pipe_.InitBuffer(zeroBuf_, bufferBytes);
	319	+ }
	320	+
	321	+ if (hasWeights_ != 0) {
	322	+ weightsI32Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(weights));
	323	+ weightsI64Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(weights));
	324	+ weightsF32Gm_.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(weights));
	325	+ }
	326	+}
	327	+
	328	+__aicore__ inline int64_t ConfusionMatrix::ReadLabel(uint64_t offset)
	329	+{
	330	+ if (labelsDtype_ == CONFUSION_MATRIX_DTYPE_INT64) {
	331	+ return labelsI64Gm_.GetValue(offset);
	332	+ }
	333	+ return static_cast<int64_t>(labelsI32Gm_.GetValue(offset));
	334	+}
	335	+
	336	+__aicore__ inline int64_t ConfusionMatrix::ReadPrediction(uint64_t offset)
	337	+{
	338	+ if (predictionsDtype_ == CONFUSION_MATRIX_DTYPE_INT64) {
	339	+ return predictionsI64Gm_.GetValue(offset);
	340	+ }
	341	+ return static_cast<int64_t>(predictionsI32Gm_.GetValue(offset));
	342	+}
	343	+
	344	+__aicore__ inline int32_t ConfusionMatrix::ReadWeightAsInt32(uint64_t offset)
	345	+{
	346	+ if (hasWeights_ == 0) {
	347	+ return 1;
	348	+ }
	349	+ if (weightsDtype_ == CONFUSION_MATRIX_DTYPE_INT64) {
	350	+ return static_cast<int32_t>(weightsI64Gm_.GetValue(offset));
	351	+ }
	352	+ if (weightsDtype_ == CONFUSION_MATRIX_DTYPE_FLOAT32) {
	353	+ return static_cast<int32_t>(weightsF32Gm_.GetValue(offset));
	354	+ }
	355	+ return weightsI32Gm_.GetValue(offset);
	356	+}
	357	+
	358	+__aicore__ inline int64_t ConfusionMatrix::ReadWeightAsInt64(uint64_t offset)
	359	+{
	360	+ if (hasWeights_ == 0) {
	361	+ return 1;
	362	+ }
	363	+ if (weightsDtype_ == CONFUSION_MATRIX_DTYPE_INT64) {
	364	+ return weightsI64Gm_.GetValue(offset);
	365	+ }
	366	+ if (weightsDtype_ == CONFUSION_MATRIX_DTYPE_FLOAT32) {
	367	+ return static_cast<int64_t>(weightsF32Gm_.GetValue(offset));
	368	+ }
	369	+ return static_cast<int64_t>(weightsI32Gm_.GetValue(offset));
	370	+}
	371	+
	372	+__aicore__ inline float ConfusionMatrix::ReadWeightAsFloat(uint64_t offset)
	373	+{
	374	+ if (hasWeights_ == 0) {
	375	+ return 1.0f;
	376	+ }
	377	+ if (weightsDtype_ == CONFUSION_MATRIX_DTYPE_INT64) {
	378	+ return static_cast<float>(weightsI64Gm_.GetValue(offset));
	379	+ }
	380	+ if (weightsDtype_ == CONFUSION_MATRIX_DTYPE_FLOAT32) {
	381	+ return weightsF32Gm_.GetValue(offset);
	382	+ }
	383	+ return static_cast<float>(weightsI32Gm_.GetValue(offset));
	384	+}
	385	+
	386	+__aicore__ inline void ConfusionMatrix::ClearOutput()
	387	+{
	388	+ ClearOutputRange(0, static_cast<uint64_t>(outputNum_));
	389	+}
	390	+
	391	+__aicore__ inline void ConfusionMatrix::ClearOutputRange(uint64_t start, uint64_t end)
	392	+{
	393	+ const uint64_t outputNum = static_cast<uint64_t>(outputNum_);
	394	+ if (start >= outputNum) {
	395	+ return;
	396	+ }
	397	+ if (end > outputNum) {
	398	+ end = outputNum;
	399	+ }
	400	+ if (outputDtype_ == CONFUSION_MATRIX_DTYPE_FLOAT32) {
	401	+ if (useBulkClear_ != 0 && end - start >= CONFUSION_MATRIX_BULK_CLEAR_MIN_ELEMS) {
	402	+ ClearOutputRangeFloat(start, end);
	403	+ return;
	404	+ }
	405	+ for (uint64_t i = start; i < end; ++i) {
	406	+ outputF32Gm_.SetValue(i, 0.0f);
	407	+ }
	408	+ } else if (outputDtype_ == CONFUSION_MATRIX_DTYPE_INT64) {
	409	+ for (uint64_t i = start; i < end; ++i) {
	410	+ outputI64Gm_.SetValue(i, static_cast<int64_t>(0));
	411	+ }
	412	+ } else {
	413	+ if (useBulkClear_ != 0 && end - start >= CONFUSION_MATRIX_BULK_CLEAR_MIN_ELEMS) {
	414	+ ClearOutputRangeInt32(start, end);
	415	+ return;
	416	+ }
	417	+ for (uint64_t i = start; i < end; ++i) {
	418	+ outputI32Gm_.SetValue(i, static_cast<int32_t>(0));
	419	+ }
	420	+ }
	421	+}
	422	+
	423	+__aicore__ inline void ConfusionMatrix::ClearOutputRangeInt32(uint64_t start, uint64_t end)
	424	+{
	425	+ const bool useMixedL2 =
	426	+ (reserved_ & CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE) != 0 &&
	427	+ sparseNormalStart_ > start && sparseNormalStart_ < end;
	428	+ uint64_t offset = start;
	429	+ while (offset < end && (offset & 7UL) != 0) {
	430	+ if (useMixedL2 && offset < sparseNormalStart_) {
	431	+ outputI32BypassGm_.SetValue(offset, static_cast<int32_t>(0));
	432	+ } else {
	433	+ outputI32Gm_.SetValue(offset, static_cast<int32_t>(0));
	434	+ }
	435	+ ++offset;
	436	+ }
	437	+
	438	+ const uint64_t alignedEnd = end & (~7UL);
	439	+ if (offset >= alignedEnd) {
	440	+ while (offset < end) {
	441	+ outputI32Gm_.SetValue(offset, static_cast<int32_t>(0));
	442	+ ++offset;
	443	+ }
	444	+ return;
	445	+ }
	446	+
	447	+ LocalTensor<int32_t> zeroLocal = zeroBuf_.Get<int32_t>();
	448	+ const uint32_t clearTileElems =
	449	+ (reserved_ & CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE) != 0 ?
	450	+ CONFUSION_MATRIX_SPARSE_FILTER_BUFFER_BYTES / sizeof(int32_t) :
	451	+ CONFUSION_MATRIX_ZERO_TILE_ELEMS;
	452	+ Duplicate(zeroLocal, static_cast<int32_t>(0), clearTileElems);
	453	+ TEventID eventIdVToMte3 = static_cast<TEventID>(0);
	454	+ SetFlag<HardEvent::V_MTE3>(eventIdVToMte3);
	455	+ WaitFlag<HardEvent::V_MTE3>(eventIdVToMte3);
	456	+
	457	+ // Keep the long case-8 clear copies on 512-byte GM boundaries. The row
	458	+ // partition only guarantees 32-byte alignment, so issue one short prefix
	459	+ // before the 184 KiB tiles. Other cases keep their original clear path.
	460	+ if ((reserved_ & CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE) != 0) {
	461	+ const uint32_t offsetIn512Bytes = static_cast<uint32_t>(offset & 127UL);
	462	+ if (offsetIn512Bytes != 0 && offset < alignedEnd) {
	463	+ uint64_t prefixElems = 128U - offsetIn512Bytes;
	464	+ const uint64_t remainingElems = alignedEnd - offset;
	465	+ if (prefixElems > remainingElems) {
	466	+ prefixElems = remainingElems;
	467	+ }
	468	+ if (useMixedL2 && offset < sparseNormalStart_) {
	469	+ const uint64_t bypassRemaining = sparseNormalStart_ - offset;
	470	+ if (prefixElems > bypassRemaining) {
	471	+ prefixElems = bypassRemaining;
	472	+ }
	473	+ DataCopy(outputI32BypassGm_[offset], zeroLocal, static_cast<uint32_t>(prefixElems));
	474	+ } else {
	475	+ DataCopy(outputI32Gm_[offset], zeroLocal, static_cast<uint32_t>(prefixElems));
	476	+ }
	477	+ offset += prefixElems;
	478	+ }
	479	+ }
	480	+
	481	+ while (offset < alignedEnd) {
	482	+ uint64_t segmentEnd = alignedEnd;
	483	+ if (useMixedL2 && offset < sparseNormalStart_) {
	484	+ segmentEnd = sparseNormalStart_;
	485	+ }
	486	+ uint64_t elems = segmentEnd - offset;
	487	+ if (elems > clearTileElems) {
	488	+ elems = clearTileElems;
	489	+ }
	490	+ if (useMixedL2 && offset < sparseNormalStart_) {
	491	+ DataCopy(outputI32BypassGm_[offset], zeroLocal, static_cast<uint32_t>(elems));
	492	+ } else {
	493	+ DataCopy(outputI32Gm_[offset], zeroLocal, static_cast<uint32_t>(elems));
	494	+ }
	495	+ offset += elems;
	496	+ }
	497	+ PipeBarrier<PIPE_ALL>();
	498	+
	499	+ while (offset < end) {
	500	+ if (useMixedL2 && offset < sparseNormalStart_) {
	501	+ outputI32BypassGm_.SetValue(offset, static_cast<int32_t>(0));
	502	+ } else {
	503	+ outputI32Gm_.SetValue(offset, static_cast<int32_t>(0));
	504	+ }
	505	+ ++offset;
	506	+ }
	507	+}
	508	+
	509	+__aicore__ inline void ConfusionMatrix::ClearOutputRangeFloat(uint64_t start, uint64_t end)
	510	+{
	511	+ uint64_t offset = start;
	512	+ while (offset < end && (offset & 7UL) != 0) {
	513	+ outputF32Gm_.SetValue(offset, 0.0f);
	514	+ ++offset;
	515	+ }
	516	+
	517	+ const uint64_t alignedEnd = end & (~7UL);
	518	+ if (offset >= alignedEnd) {
	519	+ while (offset < end) {
	520	+ outputF32Gm_.SetValue(offset, 0.0f);
	521	+ ++offset;
	522	+ }
	523	+ return;
	524	+ }
	525	+
	526	+ LocalTensor<float> zeroLocal = zeroBuf_.Get<float>();
	527	+ Duplicate(zeroLocal, 0.0f, CONFUSION_MATRIX_ZERO_TILE_ELEMS);
	528	+ TEventID eventIdVToMte3 = static_cast<TEventID>(0);
	529	+ SetFlag<HardEvent::V_MTE3>(eventIdVToMte3);
	530	+ WaitFlag<HardEvent::V_MTE3>(eventIdVToMte3);
	531	+
	532	+ while (offset < alignedEnd) {
	533	+ uint64_t elems = alignedEnd - offset;
	534	+ if (elems > CONFUSION_MATRIX_ZERO_TILE_ELEMS) {
	535	+ elems = CONFUSION_MATRIX_ZERO_TILE_ELEMS;
	536	+ }
	537	+ DataCopy(outputF32Gm_[offset], zeroLocal, static_cast<uint32_t>(elems));
	538	+ offset += elems;
	539	+ }
	540	+ PipeBarrier<PIPE_ALL>();
	541	+
	542	+ while (offset < end) {
	543	+ outputF32Gm_.SetValue(offset, 0.0f);
	544	+ ++offset;
	545	+ }
	546	+}
	547	+
	548	+__aicore__ inline void ConfusionMatrix::AddToOutput(uint64_t offset, uint64_t sampleIdx)
	549	+{
	550	+ if (outputDtype_ == CONFUSION_MATRIX_DTYPE_FLOAT32) {
	551	+ const float oldValue = outputF32Gm_.GetValue(offset);
	552	+ outputF32Gm_.SetValue(offset, oldValue + ReadWeightAsFloat(sampleIdx));
	553	+ } else if (outputDtype_ == CONFUSION_MATRIX_DTYPE_INT64) {
	554	+ const int64_t oldValue = outputI64Gm_.GetValue(offset);
	555	+ outputI64Gm_.SetValue(offset, oldValue + ReadWeightAsInt64(sampleIdx));
	556	+ } else {
	557	+ const int32_t oldValue = outputI32Gm_.GetValue(offset);
	558	+ outputI32Gm_.SetValue(offset, oldValue + ReadWeightAsInt32(sampleIdx));
	559	+ }
	560	+}
	561	+
	562	+__aicore__ inline void ConfusionMatrix::Process()
	563	+{
	564	+ if ((reserved_ & CONFUSION_MATRIX_RESERVED_SPLIT_OUTPUT) != 0) {
	565	+ ProcessSplitOutput();
	566	+ return;
	567	+ }
	568	+ ProcessSingleCore();
	569	+}
	570	+
	571	+#if TILING_KEY_VAR == CONFUSIONMATRIX_TPL_SCH_MODE_6
	572	+__aicore__ inline void ConfusionMatrix::ProcessAtomicSparse()
	573	+{
	574	+ const uint64_t blockIdx = static_cast<uint64_t>(GetBlockIdx());
	575	+ const uint64_t blockNum = static_cast<uint64_t>(GetBlockNum());
	576	+ const uint64_t classCount = static_cast<uint64_t>(numClasses_);
	577	+ const uint64_t rowsPerCore = static_cast<uint64_t>(blockFactor_);
	578	+ const uint64_t rowStart = blockIdx * rowsPerCore;
	579	+ uint64_t rowEnd = rowStart + rowsPerCore;
	580	+ if (rowEnd > classCount) {
	581	+ rowEnd = classCount;
	582	+ }
	583	+
	584	+ if (rowStart < classCount) {
	585	+ const uint64_t clearStart = rowStart * classCount;
	586	+ const uint64_t clearEnd = rowEnd * classCount;
	587	+ sparseNormalStart_ = clearStart;
	588	+ uint64_t normalTailElems =
	589	+ CONFUSION_MATRIX_SPARSE_NORMAL_L2_BUDGET_BYTES / sizeof(int32_t) / blockNum;
	590	+ normalTailElems &= ~127UL;
	591	+ if (normalTailElems != 0 && clearEnd - clearStart > normalTailElems) {
	592	+ sparseNormalStart_ = (clearEnd - normalTailElems + 127UL) & ~127UL;
	593	+ if (sparseNormalStart_ > clearEnd) {
	594	+ sparseNormalStart_ = clearEnd;
	595	+ }
	596	+ }
	597	+ ClearOutputRange(clearStart, clearEnd);
	598	+ }
	599	+
	600	+ SyncAll();
	601	+
	602	+ LocalTensor<int32_t> atomicLocal = zeroBuf_.Get<int32_t>();
	603	+ for (uint32_t lane = 0; lane < 8; ++lane) {
	604	+ atomicLocal.SetValue(lane * 8 + lane, static_cast<int32_t>(1));
	605	+ }
	606	+ const TEventID atomicDataReadyEvent = static_cast<TEventID>(0);
	607	+ SetFlag<HardEvent::S_MTE3>(atomicDataReadyEvent);
	608	+ WaitFlag<HardEvent::S_MTE3>(atomicDataReadyEvent);
	609	+
	610	+ const uint64_t totalNum = static_cast<uint64_t>(totalNum_);
	611	+ const uint64_t samplesPerCore = (totalNum + blockNum - 1) / blockNum;
	612	+ const uint64_t sampleStart = blockIdx * samplesPerCore;
	613	+ uint64_t sampleEnd = sampleStart + samplesPerCore;
	614	+ if (sampleEnd > totalNum) {
	615	+ sampleEnd = totalNum;
	616	+ }
	617	+
	618	+ SetAtomicAdd<int32_t>();
	619	+ for (uint64_t i = sampleStart; i < sampleEnd; ++i) {
	620	+ const int32_t labelValue = labelsI32Gm_.GetValue(i);
	621	+ const int32_t predictionValue = predictionsI32Gm_.GetValue(i);
	622	+ if (labelValue < 0 \|\| predictionValue < 0 \|\|
	623	+ labelValue >= numClasses_ \|\| predictionValue >= numClasses_) {
	624	+ continue;
	625	+ }
	626	+ const uint64_t outOffset = static_cast<uint64_t>(labelValue) * classCount +
	627	+ static_cast<uint64_t>(predictionValue);
	628	+ const uint32_t lane = static_cast<uint32_t>(outOffset & 7UL);
	629	+ DataCopy(outputI32Gm_[outOffset - lane], atomicLocal[lane * 8], 8);
	630	+ }
	631	+ SetAtomicNone();
	632	+}
	633	+#endif
	634	+
	635	+__aicore__ inline void ConfusionMatrix::ProcessSingleCore()
	636	+{
	637	+ const uint64_t totalNum = static_cast<uint64_t>(totalNum_);
	638	+ if (CanUseSmallOutputRowSplit(totalNum)) {
	639	+ ProcessSmallOutputRowSplit(totalNum);
	640	+ return;
	641	+ }
	642	+
	643	+ if (GetBlockIdx() != 0) {
	644	+ return;
	645	+ }
	646	+ ClearOutput();
	647	+
	648	+ const int64_t numClasses = numClasses_;
	649	+ if (labelsDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	650	+ predictionsDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	651	+ outputDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	652	+ hasWeights_ == 0) {
	653	+ if (numClasses == 3 && totalNum >= 512) {
	654	+ int32_t h0 = 0;
	655	+ int32_t h1 = 0;
	656	+ int32_t h2 = 0;
	657	+ int32_t h3 = 0;
	658	+ int32_t h4 = 0;
	659	+ int32_t h5 = 0;
	660	+ int32_t h6 = 0;
	661	+ int32_t h7 = 0;
	662	+ int32_t h8 = 0;
	663	+ uint64_t i = 0;
	664	+#define CM_C3_HIST9_STEP(IDX) \
	665	+ do { \
	666	+ const uint64_t label = static_cast<uint64_t>(labelsI32Gm_.GetValue(IDX)); \
	667	+ const uint64_t prediction = static_cast<uint64_t>(predictionsI32Gm_.GetValue(IDX)); \
	668	+ switch (label * 3 + prediction) { \
	669	+ case 0: ++h0; break; \
	670	+ case 1: ++h1; break; \
	671	+ case 2: ++h2; break; \
	672	+ case 3: ++h3; break; \
	673	+ case 4: ++h4; break; \
	674	+ case 5: ++h5; break; \
	675	+ case 6: ++h6; break; \
	676	+ case 7: ++h7; break; \
	677	+ case 8: ++h8; break; \
	678	+ default: break; \
	679	+ } \
	680	+ } while (false)
	681	+ for (; i + 5 < totalNum; i += 6) {
	682	+ CM_C3_HIST9_STEP(i);
	683	+ CM_C3_HIST9_STEP(i + 1);
	684	+ CM_C3_HIST9_STEP(i + 2);
	685	+ CM_C3_HIST9_STEP(i + 3);
	686	+ CM_C3_HIST9_STEP(i + 4);
	687	+ CM_C3_HIST9_STEP(i + 5);
	688	+ }
	689	+ for (; i < totalNum; ++i) {
	690	+ CM_C3_HIST9_STEP(i);
	691	+ }
	692	+#undef CM_C3_HIST9_STEP
	693	+ outputI32Gm_.SetValue(0, h0);
	694	+ outputI32Gm_.SetValue(1, h1);
	695	+ outputI32Gm_.SetValue(2, h2);
	696	+ outputI32Gm_.SetValue(3, h3);
	697	+ outputI32Gm_.SetValue(4, h4);
	698	+ outputI32Gm_.SetValue(5, h5);
	699	+ outputI32Gm_.SetValue(6, h6);
	700	+ outputI32Gm_.SetValue(7, h7);
	701	+ outputI32Gm_.SetValue(8, h8);
	702	+ return;
	703	+ }
	704	+ for (uint64_t i = 0; i < totalNum; ++i) {
	705	+ const uint64_t label = static_cast<uint64_t>(labelsI32Gm_.GetValue(i));
	706	+ const uint64_t prediction = static_cast<uint64_t>(predictionsI32Gm_.GetValue(i));
	707	+ const uint64_t outOffset = label * static_cast<uint64_t>(numClasses) + prediction;
	708	+ const int32_t oldValue = outputI32Gm_.GetValue(outOffset);
	709	+ outputI32Gm_.SetValue(outOffset, oldValue + 1);
	710	+ }
	711	+ return;
	712	+ }
	713	+
	714	+ for (uint64_t i = 0; i < totalNum; ++i) {
	715	+ const int64_t label = ReadLabel(i);
	716	+ const int64_t prediction = ReadPrediction(i);
	717	+ if (label < 0 \|\| prediction < 0 \|\| label >= numClasses \|\| prediction >= numClasses) {
	718	+ continue;
	719	+ }
	720	+ const uint64_t outOffset = static_cast<uint64_t>(label) * static_cast<uint64_t>(numClasses) +
	721	+ static_cast<uint64_t>(prediction);
	722	+ AddToOutput(outOffset, i);
	723	+ }
	724	+}
	725	+
	726	+template <bool useMte2TailPadding>
	727	+__aicore__ inline void ProcessC3VectorHistogram(
	728	+ GM_ADDR labels,
	729	+ GM_ADDR predictions,
	730	+ GM_ADDR y,
	731	+ const ConfusionMatrixTilingData* tilingData)
	732	+{
	733	+ GlobalTensor<int32_t> labelsGm;
	734	+ GlobalTensor<int32_t> predictionsGm;
	735	+ labelsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(labels));
	736	+ predictionsGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(predictions));
	737	+
	738	+ TPipe pipe;
	739	+ TBuf<TPosition::VECCALC> workBuf;
	740	+ pipe.InitBuffer(workBuf, CONFUSION_MATRIX_C3_VECTOR_WORK_BUFFER_BYTES);
	741	+
	742	+ const uint64_t totalNum = static_cast<uint64_t>(tilingData->totalNum);
	743	+ LocalTensor<int32_t> dataLocal = workBuf.Get<int32_t>();
	744	+ LocalTensor<int32_t> labelsLocal = dataLocal;
	745	+ LocalTensor<int32_t> predictionsLocal = dataLocal[CONFUSION_MATRIX_C3_VECTOR_TILE_ELEMS];
	746	+ LocalTensor<uint8_t> maskStorage =
	747	+ workBuf.Get<uint8_t>()[CONFUSION_MATRIX_C3_VECTOR_DATA_BUFFER_BYTES];
	748	+ LocalTensor<uint8_t> labelMask = maskStorage;
	749	+ LocalTensor<uint8_t> predictionMask0 = maskStorage[CONFUSION_MATRIX_C3_VECTOR_MASK_BYTES];
	750	+ LocalTensor<uint8_t> predictionMask1 = maskStorage[CONFUSION_MATRIX_C3_VECTOR_MASK_BYTES * 2];
	751	+
	752	+ int32_t h0 = 0;
	753	+ int32_t h1 = 0;
	754	+ int32_t h2 = 0;
	755	+ const int32_t row = static_cast<int32_t>(GetBlockIdx());
	756	+ uint64_t offset = 0;
	757	+ const TEventID copyEvent = static_cast<TEventID>(0);
	758	+ const TEventID countEvent = static_cast<TEventID>(1);
	759	+
	760	+ do {
	761	+ const uint64_t remaining = totalNum - offset;
	762	+ const bool usePaddedTail = remaining <= CONFUSION_MATRIX_C3_VECTOR_TILE_ELEMS && (remaining & 7UL) == 0;
	763	+ uint32_t copyCount = remaining > CONFUSION_MATRIX_C3_VECTOR_TILE_ELEMS ?
	764	+ CONFUSION_MATRIX_C3_VECTOR_TILE_ELEMS : static_cast<uint32_t>(remaining);
	765	+ uint32_t vectorCount = copyCount;
	766	+ const uint32_t paddingCount = (64U - (copyCount & 63U)) & 63U;
	767	+ if (usePaddedTail) {
	768	+ vectorCount = (copyCount + 63U) & ~63U;
	769	+ DataCopyParams inputParams{1, static_cast<uint16_t>(copyCount * sizeof(int32_t)), 0, 0};
	770	+ DataCopyPadParams inputPadParams{
	771	+ true, 0, static_cast<uint8_t>(useMte2TailPadding ? paddingCount : 0U),
	772	+ useMte2TailPadding ? static_cast<uint64_t>(static_cast<uint32_t>(-1)) : 0UL};
	773	+ DataCopyPad(labelsLocal, labelsGm[offset], inputParams, inputPadParams);
	774	+ DataCopyPad(predictionsLocal, predictionsGm[offset], inputParams, inputPadParams);
	775	+ } else {
	776	+ vectorCount &= ~63U;
	777	+ copyCount = vectorCount;
	778	+ DataCopy(labelsLocal, labelsGm[offset], vectorCount);
	779	+ DataCopy(predictionsLocal, predictionsGm[offset], vectorCount);
	780	+ }
	781	+ SetFlag<HardEvent::MTE2_V>(copyEvent);
	782	+ WaitFlag<HardEvent::MTE2_V>(copyEvent);
	783	+ if constexpr (!useMte2TailPadding) {
	784	+ if (vectorCount > copyCount) {
	785	+ Duplicate(labelsLocal[copyCount], static_cast<int32_t>(-1), vectorCount - copyCount);
	786	+ Duplicate(predictionsLocal[copyCount], static_cast<int32_t>(-1), vectorCount - copyCount);
	787	+ PipeBarrier<PIPE_V>();
	788	+ }
	789	+ }
	790	+
	791	+#define CM_C3_MASK_COUNT(MASK, COUNT) \
	792	+ do { \
	793	+ uint64_t selectedCount = 0; \
	794	+ GatherMask( \
	795	+ predictionsLocal, labelsLocal, MASK.ReinterpretCast<uint32_t>(), true, vectorCount, \
	796	+ {1, 1, 8, 0}, selectedCount); \
	797	+ PipeBarrier<PIPE_V>(); \
	798	+ SetFlag<HardEvent::V_S>(countEvent); \
	799	+ WaitFlag<HardEvent::V_S>(countEvent); \
	800	+ COUNT = static_cast<int32_t>(selectedCount); \
	801	+ } while (false)
	802	+ CompareScalar(labelMask, labelsLocal, row, CMPMODE::EQ, vectorCount);
	803	+ PipeBarrier<PIPE_V>();
	804	+ uint64_t selectedRowCount = 0;
	805	+ GatherMask(
	806	+ labelsLocal, predictionsLocal, labelMask.ReinterpretCast<uint32_t>(), true, vectorCount,
	807	+ {1, 1, 8, 0}, selectedRowCount);
	808	+ PipeBarrier<PIPE_V>();
	809	+ SetFlag<HardEvent::V_S>(countEvent);
	810	+ WaitFlag<HardEvent::V_S>(countEvent);
	811	+ const int32_t rowCount = static_cast<int32_t>(selectedRowCount);
	812	+ if (rowCount > 0) {
	813	+ const uint32_t compactVectorCount = (static_cast<uint32_t>(rowCount) + 63U) & ~63U;
	814	+ if (compactVectorCount > static_cast<uint32_t>(rowCount)) {
	815	+ Duplicate(
	816	+ labelsLocal[static_cast<uint32_t>(rowCount)], static_cast<int32_t>(-1),
	817	+ compactVectorCount - static_cast<uint32_t>(rowCount));
	818	+ PipeBarrier<PIPE_V>();
	819	+ }
	820	+ CompareScalar(predictionMask0, labelsLocal, static_cast<int32_t>(0), CMPMODE::EQ, compactVectorCount);
	821	+ CompareScalar(predictionMask1, labelsLocal, static_cast<int32_t>(1), CMPMODE::EQ, compactVectorCount);
	822	+ PipeBarrier<PIPE_V>();
	823	+ vectorCount = compactVectorCount;
	824	+ int32_t count0 = 0;
	825	+ int32_t count1 = 0;
	826	+ CM_C3_MASK_COUNT(predictionMask0, count0);
	827	+ CM_C3_MASK_COUNT(predictionMask1, count1);
	828	+ h0 += count0;
	829	+ h1 += count1;
	830	+ h2 += rowCount - count0 - count1;
	831	+ }
	832	+#undef CM_C3_MASK_COUNT
	833	+ offset += copyCount;
	834	+ } while (offset + 64 <= totalNum);
	835	+
	836	+ for (; offset < totalNum; ++offset) {
	837	+ const int32_t label = labelsGm.GetValue(offset);
	838	+ const int32_t prediction = predictionsGm.GetValue(offset);
	839	+ if (label != row \|\| prediction < 0 \|\| prediction >= 3) {
	840	+ continue;
	841	+ }
	842	+ switch (prediction) {
	843	+ case 0: ++h0; break;
	844	+ case 1: ++h1; break;
	845	+ case 2: ++h2; break;
	846	+ default: break;
	847	+ }
	848	+ }
	849	+
	850	+ GlobalTensor<int32_t> outputGm;
	851	+ outputGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(y));
	852	+ const uint64_t outputOffset = static_cast<uint64_t>(row) * 3;
	853	+ outputGm.SetValue(outputOffset, h0);
	854	+ outputGm.SetValue(outputOffset + 1, h1);
	855	+ outputGm.SetValue(outputOffset + 2, h2);
	856	+ PipeBarrier<PIPE_ALL>();
	857	+}
	858	+
	859	+__aicore__ inline bool ConfusionMatrix::CanUseSmallOutputRowSplit(uint64_t totalNum) const
	860	+{
	861	+ if (numClasses_ <= 1 \|\| numClasses_ > static_cast<int64_t>(CONFUSION_MATRIX_SMALL_ROW_SPLIT_MAX_CLASSES)) {
	862	+ return false;
	863	+ }
	864	+ const uint64_t classCount = static_cast<uint64_t>(numClasses_);
	865	+ const uint64_t rowBytes = classCount * sizeof(int32_t);
	866	+ return labelsDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	867	+ predictionsDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	868	+ outputDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	869	+ hasWeights_ == 0 &&
	870	+ totalNum >= CONFUSION_MATRIX_SMALL_ROW_SPLIT_MIN_SAMPLES &&
	871	+ rowBytes >= 32 &&
	872	+ (rowBytes & 31UL) == 0 &&
	873	+ GetBlockNum() > 1;
	874	+}
	875	+
	876	+__aicore__ inline void ConfusionMatrix::ProcessSmallOutputRowSplit(uint64_t totalNum)
	877	+{
	878	+ const uint64_t blockIdx = static_cast<uint64_t>(GetBlockIdx());
	879	+ const uint64_t blockNum = static_cast<uint64_t>(GetBlockNum());
	880	+ const uint64_t classCount = static_cast<uint64_t>(numClasses_);
	881	+ const uint64_t rowsPerCore = (classCount + blockNum - 1) / blockNum;
	882	+ const uint64_t rowStart = blockIdx * rowsPerCore;
	883	+ if (rowStart >= classCount) {
	884	+ return;
	885	+ }
	886	+ uint64_t rowEnd = rowStart + rowsPerCore;
	887	+ if (rowEnd > classCount) {
	888	+ rowEnd = classCount;
	889	+ }
	890	+
	891	+ ClearOutputRange(rowStart * classCount, rowEnd * classCount);
	892	+
	893	+ for (uint64_t i = 0; i < totalNum; ++i) {
	894	+ const int32_t labelValue = labelsI32Gm_.GetValue(i);
	895	+ if (labelValue < 0) {
	896	+ continue;
	897	+ }
	898	+ const uint64_t label = static_cast<uint64_t>(labelValue);
	899	+ if (label < rowStart \|\| label >= rowEnd) {
	900	+ continue;
	901	+ }
	902	+ const int32_t predictionValue = predictionsI32Gm_.GetValue(i);
	903	+ if (predictionValue < 0) {
	904	+ continue;
	905	+ }
	906	+ const uint64_t prediction = static_cast<uint64_t>(predictionValue);
	907	+ if (prediction >= classCount) {
	908	+ continue;
	909	+ }
	910	+ const uint64_t outOffset = label * classCount + prediction;
	911	+ const int32_t oldValue = outputI32Gm_.GetValue(outOffset);
	912	+ outputI32Gm_.SetValue(outOffset, oldValue + 1);
	913	+ }
	914	+}
	915	+
	916	+__aicore__ inline bool ConfusionMatrix::CanUseSparseSplitInt32(uint64_t totalNum) const
	917	+{
	918	+ if (useBulkClear_ == 0 \|\| totalNum == 0 \|\| totalNum > CONFUSION_MATRIX_SPARSE_CAPACITY \|\|
	919	+ outputNum_ > static_cast<int64_t>(CONFUSION_MATRIX_UINT32_OFFSET_MAX)) {
	920	+ return false;
	921	+ }
	922	+ const uint64_t outputNum = static_cast<uint64_t>(outputNum_);
	923	+ return outputNum / totalNum >= totalNum;
	924	+}
	925	+
	926	+__aicore__ inline void ConfusionMatrix::ProcessSplitOutputSparseInt32(
	927	+ uint64_t rowStart,
	928	+ uint64_t rowEnd,
	929	+ uint64_t classCount,
	930	+ uint64_t totalNum)
	931	+{
	932	+ const bool useVectorFilter =
	933	+ (reserved_ & CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE) != 0;
	934	+ LocalTensor<int32_t> offsetLowLocal = zeroBuf_.Get<int32_t>();
	935	+ LocalTensor<int32_t> countLocal = zeroBuf_.Get<int32_t>()[CONFUSION_MATRIX_SPARSE_COUNT_LOCAL_OFFSET];
	936	+ LocalTensor<int32_t> occupiedSlotLocal =
	937	+ zeroBuf_.Get<int32_t>()[CONFUSION_MATRIX_SPARSE_OCCUPIED_LOCAL_OFFSET];
	938	+ if (!useVectorFilter) {
	939	+ Duplicate(countLocal, static_cast<int32_t>(0), CONFUSION_MATRIX_SPARSE_CAPACITY);
	940	+ TEventID eventIdVToS = static_cast<TEventID>(1);
	941	+ SetFlag<HardEvent::V_S>(eventIdVToS);
	942	+ WaitFlag<HardEvent::V_S>(eventIdVToS);
	943	+ }
	944	+ uint32_t occupiedCount = 0;
	945	+ const int32_t rowStartI32 = static_cast<int32_t>(rowStart);
	946	+ const int32_t rowEndI32 = static_cast<int32_t>(rowEnd);
	947	+ const int32_t classCountI32 = static_cast<int32_t>(classCount);
	948	+#define CM_ACCUMULATE_SPARSE_OFFSET(OUT_OFFSET_LOW) \
	949	+ do { \
	950	+ const uint32_t currentOffsetLow = (OUT_OFFSET_LOW); \
	951	+ const uint32_t hashValue = currentOffsetLow ^ (currentOffsetLow >> 7); \
	952	+ uint32_t slot = hashValue % CONFUSION_MATRIX_SPARSE_CAPACITY; \
	953	+ for (uint32_t probe = 0; probe < CONFUSION_MATRIX_SPARSE_CAPACITY; ++probe) { \
	954	+ const int32_t oldCount = countLocal.GetValue(slot); \
	955	+ if (oldCount == 0) { \
	956	+ offsetLowLocal.SetValue(slot, static_cast<int32_t>(currentOffsetLow)); \
	957	+ countLocal.SetValue(slot, static_cast<int32_t>(1)); \
	958	+ occupiedSlotLocal.SetValue(occupiedCount, static_cast<int32_t>(slot)); \
	959	+ ++occupiedCount; \
	960	+ break; \
	961	+ } \
	962	+ if (static_cast<uint32_t>(offsetLowLocal.GetValue(slot)) == currentOffsetLow) { \
	963	+ countLocal.SetValue(slot, oldCount + 1); \
	964	+ break; \
	965	+ } \
	966	+ ++slot; \
	967	+ if (slot == CONFUSION_MATRIX_SPARSE_CAPACITY) { \
	968	+ slot = 0; \
	969	+ } \
	970	+ } \
	971	+ } while (false)
	972	+
	973	+ if (useVectorFilter) {
	974	+ LocalTensor<int32_t> labelsLocal =
	975	+ zeroBuf_.Get<int32_t>()[CONFUSION_MATRIX_SPARSE_FILTER_LABEL_LOCAL_OFFSET];
	976	+ LocalTensor<int32_t> predictionsLocal =
	977	+ zeroBuf_.Get<int32_t>()[CONFUSION_MATRIX_SPARSE_FILTER_PREDICTION_LOCAL_OFFSET];
	978	+ LocalTensor<int32_t> compactLabelsLocal =
	979	+ zeroBuf_.Get<int32_t>()[CONFUSION_MATRIX_SPARSE_FILTER_COMPACT_LABEL_LOCAL_OFFSET];
	980	+ LocalTensor<float> labelsFloatLocal = compactLabelsLocal.ReinterpretCast<float>();
	981	+ LocalTensor<uint8_t> maskStorage =
	982	+ zeroBuf_.Get<uint8_t>()[CONFUSION_MATRIX_SPARSE_FILTER_MASK_LOCAL_OFFSET_BYTES];
	983	+ LocalTensor<uint8_t> lowerMask = maskStorage;
	984	+ LocalTensor<uint8_t> upperMask = maskStorage[CONFUSION_MATRIX_SPARSE_FILTER_MASK_BYTES];
	985	+ LocalTensor<uint8_t> ownedMask = maskStorage[CONFUSION_MATRIX_SPARSE_FILTER_MASK_BYTES * 2];
	986	+
	987	+ const TEventID copyEvent = static_cast<TEventID>(0);
	988	+ const TEventID selectedCountEvent = static_cast<TEventID>(2);
	989	+ uint64_t offset = 0;
	990	+ while (offset + 64 <= totalNum) {
	991	+ const uint64_t remaining = totalNum - offset;
	992	+ const bool usePaddedTail = remaining <= CONFUSION_MATRIX_SPARSE_FILTER_TILE_ELEMS &&
	993	+ (remaining & 7UL) == 0;
	994	+ uint32_t copyCount = remaining > CONFUSION_MATRIX_SPARSE_FILTER_TILE_ELEMS ?
	995	+ CONFUSION_MATRIX_SPARSE_FILTER_TILE_ELEMS : static_cast<uint32_t>(remaining);
	996	+ uint32_t vectorCount = copyCount;
	997	+ if (usePaddedTail) {
	998	+ vectorCount = (copyCount + 63U) & ~63U;
	999	+ DataCopyParams inputParams{1, static_cast<uint16_t>(copyCount * sizeof(int32_t)), 0, 0};
	1000	+ DataCopyPadParams inputPadParams{true, 0, 0, 0};
	1001	+ DataCopyPad(labelsLocal, labelsI32Gm_[offset], inputParams, inputPadParams);
	1002	+ DataCopyPad(predictionsLocal, predictionsI32Gm_[offset], inputParams, inputPadParams);
	1003	+ } else {
	1004	+ vectorCount &= ~63U;
	1005	+ copyCount = vectorCount;
	1006	+ DataCopy(labelsLocal, labelsI32Gm_[offset], vectorCount);
	1007	+ DataCopy(predictionsLocal, predictionsI32Gm_[offset], vectorCount);
	1008	+ }
	1009	+ SetFlag<HardEvent::MTE2_V>(copyEvent);
	1010	+ WaitFlag<HardEvent::MTE2_V>(copyEvent);
	1011	+ if (vectorCount > copyCount) {
	1012	+ Duplicate(labelsLocal[copyCount], static_cast<int32_t>(-1), vectorCount - copyCount);
	1013	+ Duplicate(predictionsLocal[copyCount], static_cast<int32_t>(-1), vectorCount - copyCount);
	1014	+ PipeBarrier<PIPE_V>();
	1015	+ }
	1016	+
	1017	+ Cast(labelsFloatLocal, labelsLocal, RoundMode::CAST_NONE, vectorCount);
	1018	+ PipeBarrier<PIPE_V>();
	1019	+ CompareScalar(lowerMask, labelsFloatLocal, static_cast<float>(rowStartI32), CMPMODE::GE, vectorCount);
	1020	+ CompareScalar(upperMask, labelsFloatLocal, static_cast<float>(rowEndI32), CMPMODE::LT, vectorCount);
	1021	+ PipeBarrier<PIPE_V>();
	1022	+ And(
	1023	+ ownedMask.ReinterpretCast<uint16_t>(), lowerMask.ReinterpretCast<uint16_t>(),
	1024	+ upperMask.ReinterpretCast<uint16_t>(), vectorCount / 16);
	1025	+ PipeBarrier<PIPE_V>();
	1026	+
	1027	+ uint64_t selectedLabelCount = 0;
	1028	+ GatherMask(
	1029	+ compactLabelsLocal, labelsLocal, ownedMask.ReinterpretCast<uint32_t>(), true, vectorCount,
	1030	+ {1, 1, 8, 0}, selectedLabelCount);
	1031	+ PipeBarrier<PIPE_V>();
	1032	+ uint64_t selectedCount = 0;
	1033	+ GatherMask(
	1034	+ labelsLocal, predictionsLocal, ownedMask.ReinterpretCast<uint32_t>(), true, vectorCount,
	1035	+ {1, 1, 8, 0}, selectedCount);
	1036	+ PipeBarrier<PIPE_V>();
	1037	+ SetFlag<HardEvent::V_S>(selectedCountEvent);
	1038	+ WaitFlag<HardEvent::V_S>(selectedCountEvent);
	1039	+ for (uint32_t j = 0; j < static_cast<uint32_t>(selectedCount); ++j) {
	1040	+ const int32_t labelValue = compactLabelsLocal.GetValue(j);
	1041	+ const int32_t predictionValue = labelsLocal.GetValue(j);
	1042	+ if (predictionValue < 0 \|\| predictionValue >= classCountI32) {
	1043	+ continue;
	1044	+ }
	1045	+ const uint64_t outOffset = static_cast<uint64_t>(labelValue) * classCount +
	1046	+ static_cast<uint64_t>(predictionValue);
	1047	+ CM_ACCUMULATE_SPARSE_OFFSET(static_cast<uint32_t>(outOffset));
	1048	+ }
	1049	+ offset += copyCount;
	1050	+ }
	1051	+
	1052	+ for (; offset < totalNum; ++offset) {
	1053	+ const int32_t labelValue = labelsI32Gm_.GetValue(offset);
	1054	+ if (labelValue < rowStartI32 \|\| labelValue >= rowEndI32) {
	1055	+ continue;
	1056	+ }
	1057	+ const int32_t predictionValue = predictionsI32Gm_.GetValue(offset);
	1058	+ if (predictionValue < 0 \|\| predictionValue >= classCountI32) {
	1059	+ continue;
	1060	+ }
	1061	+ const uint64_t outOffset = static_cast<uint64_t>(labelValue) * classCount +
	1062	+ static_cast<uint64_t>(predictionValue);
	1063	+ CM_ACCUMULATE_SPARSE_OFFSET(static_cast<uint32_t>(outOffset));
	1064	+ }
	1065	+ } else {
	1066	+ for (uint64_t i = 0; i < totalNum; ++i) {
	1067	+ const int32_t labelValue = labelsI32Gm_.GetValue(i);
	1068	+ if (labelValue < rowStartI32 \|\| labelValue >= rowEndI32) {
	1069	+ continue;
	1070	+ }
	1071	+ const int32_t predictionValue = predictionsI32Gm_.GetValue(i);
	1072	+ if (predictionValue < 0 \|\| predictionValue >= classCountI32) {
	1073	+ continue;
	1074	+ }
	1075	+ const uint64_t outOffset = static_cast<uint64_t>(labelValue) * classCount +
	1076	+ static_cast<uint64_t>(predictionValue);
	1077	+ CM_ACCUMULATE_SPARSE_OFFSET(static_cast<uint32_t>(outOffset));
	1078	+ }
	1079	+ }
	1080	+#undef CM_ACCUMULATE_SPARSE_OFFSET
	1081	+
	1082	+ for (uint32_t i = 0; i < occupiedCount; ++i) {
	1083	+ const uint32_t k = static_cast<uint32_t>(occupiedSlotLocal.GetValue(i));
	1084	+ const int32_t count = countLocal.GetValue(k);
	1085	+ const uint64_t outOffset = static_cast<uint64_t>(static_cast<uint32_t>(offsetLowLocal.GetValue(k)));
	1086	+ if (useVectorFilter && outOffset < sparseNormalStart_) {
	1087	+ outputI32BypassGm_.SetValue(outOffset, count);
	1088	+ } else {
	1089	+ outputI32Gm_.SetValue(outOffset, count);
	1090	+ }
	1091	+ }
	1092	+}
	1093	+
	1094	+__aicore__ inline void ConfusionMatrix::ProcessSplitOutput()
	1095	+{
	1096	+ const uint64_t blockIdx = static_cast<uint64_t>(GetBlockIdx());
	1097	+ const uint64_t classCount = static_cast<uint64_t>(numClasses_);
	1098	+ const uint64_t rowsPerCore = static_cast<uint64_t>(blockFactor_);
	1099	+ const uint64_t rowStart = blockIdx * rowsPerCore;
	1100	+ if (rowStart >= classCount) {
	1101	+ return;
	1102	+ }
	1103	+ uint64_t rowEnd = rowStart + rowsPerCore;
	1104	+ if (rowEnd > classCount) {
	1105	+ rowEnd = classCount;
	1106	+ }
	1107	+
	1108	+ const uint64_t start = rowStart * classCount;
	1109	+ const uint64_t end = rowEnd * classCount;
	1110	+ sparseNormalStart_ = start;
	1111	+ if ((reserved_ & CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE) != 0) {
	1112	+ const uint64_t blockNum = static_cast<uint64_t>(GetBlockNum());
	1113	+ uint64_t normalTailElems =
	1114	+ CONFUSION_MATRIX_SPARSE_NORMAL_L2_BUDGET_BYTES / sizeof(int32_t) / blockNum;
	1115	+ normalTailElems &= ~127UL;
	1116	+ if (normalTailElems != 0 && end - start > normalTailElems) {
	1117	+ sparseNormalStart_ = (end - normalTailElems + 127UL) & ~127UL;
	1118	+ if (sparseNormalStart_ > end) {
	1119	+ sparseNormalStart_ = end;
	1120	+ }
	1121	+ }
	1122	+ }
	1123	+ ClearOutputRange(start, end);
	1124	+
	1125	+ const uint64_t totalNum = static_cast<uint64_t>(totalNum_);
	1126	+ if (labelsDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	1127	+ predictionsDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	1128	+ outputDtype_ == CONFUSION_MATRIX_DTYPE_INT32 &&
	1129	+ hasWeights_ == 0) {
	1130	+ if (CanUseSparseSplitInt32(totalNum)) {
	1131	+ ProcessSplitOutputSparseInt32(rowStart, rowEnd, classCount, totalNum);
	1132	+ return;
	1133	+ }
	1134	+ for (uint64_t i = 0; i < totalNum; ++i) {
	1135	+ const int32_t labelValue = labelsI32Gm_.GetValue(i);
	1136	+ if (labelValue < 0) {
	1137	+ continue;
	1138	+ }
	1139	+ const uint64_t label = static_cast<uint64_t>(labelValue);
	1140	+ if (label < rowStart \|\| label >= rowEnd) {
	1141	+ continue;
	1142	+ }
	1143	+ const int32_t predictionValue = predictionsI32Gm_.GetValue(i);
	1144	+ if (predictionValue < 0) {
	1145	+ continue;
	1146	+ }
	1147	+ const uint64_t prediction = static_cast<uint64_t>(predictionValue);
	1148	+ if (prediction >= classCount) {
	1149	+ continue;
	1150	+ }
	1151	+ const uint64_t outOffset = label * classCount + prediction;
	1152	+ const int32_t oldValue = outputI32Gm_.GetValue(outOffset);
	1153	+ outputI32Gm_.SetValue(outOffset, oldValue + 1);
	1154	+ }
	1155	+ return;
	1156	+ }
	1157	+
	1158	+ for (uint64_t i = 0; i < totalNum; ++i) {
	1159	+ const int64_t label = ReadLabel(i);
	1160	+ if (label < 0) {
	1161	+ continue;
	1162	+ }
	1163	+ const uint64_t labelOffset = static_cast<uint64_t>(label);
	1164	+ if (labelOffset < rowStart \|\| labelOffset >= rowEnd) {
	1165	+ continue;
	1166	+ }
	1167	+ const int64_t prediction = ReadPrediction(i);
	1168	+ if (prediction < 0 \|\| prediction >= numClasses_) {
	1169	+ continue;
	1170	+ }
	1171	+ const uint64_t outOffset = labelOffset * classCount + static_cast<uint64_t>(prediction);
	1172	+ AddToOutput(outOffset, i);
	1173	+ }
	1174	+}
	1175	+
	1176	+} // namespace NsConfusionMatrix
	1177	+#endif // CONFUSIONMATRIX_H

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_kernel/confusion_matrix_tiling_data.h+29-0

@@ -0,0 +1,29 @@
	1	+/*!
	2	+ * \file confusion_matrix_tiling_data.h
	3	+ * \brief tiling data struct
	4	+ */
	5	+
	6	+#ifndef _CONFUSIONMATRIX_TILING_DATA_H_
	7	+#define _CONFUSIONMATRIX_TILING_DATA_H_
	8	+
	9	+#include <cstdint>
	10	+
	11	+constexpr int32_t CONFUSION_MATRIX_DTYPE_INT32 = 0;
	12	+constexpr int32_t CONFUSION_MATRIX_DTYPE_INT64 = 1;
	13	+constexpr int32_t CONFUSION_MATRIX_DTYPE_FLOAT32 = 2;
	14	+constexpr int32_t CONFUSION_MATRIX_RESERVED_SPLIT_OUTPUT = 1;
	15	+constexpr int32_t CONFUSION_MATRIX_RESERVED_VECTOR_FILTER_SPARSE = 2;
	16	+
	17	+struct ConfusionMatrixTilingData {
	18	+ int64_t totalNum = 0;
	19	+ int64_t blockFactor = 1;
	20	+ int64_t numClasses = 0;
	21	+ int32_t labelsDtype = 0;
	22	+ int32_t predictionsDtype = 0;
	23	+ int32_t weightsDtype = 0;
	24	+ int32_t outputDtype = 0;
	25	+ int32_t hasWeights = 0;
	26	+ int32_t reserved = 0;
	27	+};
	28	+
	29	+#endif

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/op_kernel/confusion_matrix_tiling_key.h+45-0

@@ -0,0 +1,45 @@
	1	+/*!
	2	+ * \file confusion_matrix_tiling_key.h
	3	+ * \brief Tiling 模板参数定义
	4	+ */
	5	+
	6	+#ifndef __CONFUSIONMATRIX_TILING_KEY_H__
	7	+#define __CONFUSIONMATRIX_TILING_KEY_H__
	8	+
	9	+#include "ascendc/host_api/tiling/template_argument.h"
	10	+
	11	+#define CONFUSIONMATRIX_TPL_SCH_MODE_0 0
	12	+#define CONFUSIONMATRIX_TPL_SCH_MODE_1 1
	13	+#define CONFUSIONMATRIX_TPL_SCH_MODE_2 2
	14	+#define CONFUSIONMATRIX_TPL_SCH_MODE_3 3
	15	+#define CONFUSIONMATRIX_TPL_SCH_MODE_4 4
	16	+#define CONFUSIONMATRIX_TPL_SCH_MODE_5 5
	17	+#define CONFUSIONMATRIX_TPL_SCH_MODE_6 6
	18	+
	19	+ASCENDC_TPL_ARGS_DECL(
	20	+ ConfusionMatrix,
	21	+ ASCENDC_TPL_UINT_DECL(
	22	+ schMode,
	23	+ 3,
	24	+ ASCENDC_TPL_UI_LIST,
	25	+ CONFUSIONMATRIX_TPL_SCH_MODE_0,
	26	+ CONFUSIONMATRIX_TPL_SCH_MODE_1,
	27	+ CONFUSIONMATRIX_TPL_SCH_MODE_2,
	28	+ CONFUSIONMATRIX_TPL_SCH_MODE_3,
	29	+ CONFUSIONMATRIX_TPL_SCH_MODE_4,
	30	+ CONFUSIONMATRIX_TPL_SCH_MODE_5,
	31	+ CONFUSIONMATRIX_TPL_SCH_MODE_6));
	32	+
	33	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
	34	+ ASCENDC_TPL_UINT_SEL(
	35	+ schMode,
	36	+ ASCENDC_TPL_UI_LIST,
	37	+ CONFUSIONMATRIX_TPL_SCH_MODE_0,
	38	+ CONFUSIONMATRIX_TPL_SCH_MODE_1,
	39	+ CONFUSIONMATRIX_TPL_SCH_MODE_2,
	40	+ CONFUSIONMATRIX_TPL_SCH_MODE_3,
	41	+ CONFUSIONMATRIX_TPL_SCH_MODE_4,
	42	+ CONFUSIONMATRIX_TPL_SCH_MODE_5,
	43	+ CONFUSIONMATRIX_TPL_SCH_MODE_6)));
	44	+
	45	+#endif

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/CMakeLists.txt+69-0

@@ -0,0 +1,69 @@
	1	+cmake_minimum_required(VERSION 3.16.0)
	2	+project(confusion_matrix_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_LABELS=int32_t
	19	+ -DDTYPE_PREDICTIONS=int32_t
	20	+ -DDTYPE_Y=int32_t
	21	+)
	22	+
	23	+set(UT_COMMON_DIR ${CMAKE_CURRENT_SOURCE_DIR}/common)
	24	+set(UT_COMMON_SRCS
	25	+ ${UT_COMMON_DIR}/infershape_case_executor.cpp
	26	+ ${UT_COMMON_DIR}/infershape_context_faker.cpp
	27	+ ${UT_COMMON_DIR}/tiling_case_executor.cpp
	28	+ ${UT_COMMON_DIR}/tiling_context_faker.cpp
	29	+)
	30	+
	31	+set(UT_COMMON_INCLUDE_DIRS
	32	+ ${UT_COMMON_DIR}
	33	+ $ENV{ASCEND_HOME_PATH}/include
	34	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/include
	35	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/include/base
	36	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/include/base/context_builder
	37	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/asc/include
	38	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/asc/include/tiling
	39	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc
	40	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/op_common
	41	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/base
	42	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/exe_graph
	43	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/pkg_inc/graph
	44	+)
	45	+
	46	+set(UT_COMMON_LIB_DIRS
	47	+ $ENV{ASCEND_HOME_PATH}/lib64
	48	+ $ENV{ASCEND_HOME_PATH}/aarch64-linux/lib64
	49	+)
	50	+
	51	+add_compile_options(
	52	+ -Wall
	53	+ -Wno-deprecated-declarations
	54	+ -Wno-unused-variable
	55	+ -fno-access-control
	56	+)
	57	+
	58	+if(NOT CMAKE_BUILD_TYPE)
	59	+ set(CMAKE_BUILD_TYPE Debug)
	60	+endif()
	61	+
	62	+message(STATUS "Build Type: ${CMAKE_BUILD_TYPE}")
	63	+
	64	+add_subdirectory(op_host)
	65	+
	66	+if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/../../op_api)
	67	+ add_subdirectory(op_api)
	68	+endif()
	69	+add_subdirectory(op_kernel)

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/cmake/BuildGoogleTest.cmake+132-0

@@ -0,0 +1,132 @@
	1	+#
	2	+# BuildGoogleTest.cmake
	3	+#
	4	+# Purpose: Build Google Test from source with OLD ABI to match CANN libraries
	5	+#
	6	+# Reference: ops-math/cmake/third_party/gtest.cmake
	7	+# ----------------------------------------------------------------------------------------------------------
	8	+
	9	+include_guard(GLOBAL)
	10	+
	11	+# CRITICAL: Force build from source with OLD ABI to match CANN libraries
	12	+# We cannot use system Google Test because it uses new ABI by default
	13	+message(STATUS "")
	14	+message(STATUS "=== Building Google Test from source with OLD ABI ===")
	15	+message(STATUS " Reason: System gtest uses new ABI, incompatible with libplatform.so")
	16	+message(STATUS " Target ABI: _GLIBCXX_USE_CXX11_ABI=0 (old ABI)")
	17	+message(STATUS "")
	18	+
	19	+set(GTEST_VERSION "1.14.0")
	20	+set(GTEST_INSTALL_DIR ${CMAKE_BINARY_DIR}/3rd_party/gtest)
	21	+set(GTEST_SOURCE_DIR ${CMAKE_BINARY_DIR}/3rd_party/gtest-src)
	22	+
	23	+# Download URL (using gitcode mirror for China access)
	24	+set(GTEST_URL "https://gitcode.com/cann-src-third-party/googletest/releases/download/v${GTEST_VERSION}/googletest-${GTEST_VERSION}.tar.gz")
	25	+
	26	+# Compiler flags - CRITICAL: Use OLD ABI
	27	+set(GTEST_CXX_FLAGS "-D_GLIBCXX_USE_CXX11_ABI=0 -O2 -D_FORTIFY_SOURCE=2 -fPIC -fstack-protector-all -w")
	28	+set(GTEST_C_FLAGS "-D_GLIBCXX_USE_CXX11_ABI=0 -O2 -D_FORTIFY_SOURCE=2 -fPIC -fstack-protector-all -w")
	29	+
	30	+include(ExternalProject)
	31	+ExternalProject_Add(
	32	+ third_party_gtest
	33	+ URL ${GTEST_URL}
	34	+ TLS_VERIFY OFF
	35	+ DOWNLOAD_DIR ${CMAKE_BINARY_DIR}/downloads
	36	+ SOURCE_DIR ${GTEST_SOURCE_DIR}
	37	+ INSTALL_DIR ${GTEST_INSTALL_DIR}
	38	+
	39	+ CMAKE_ARGS
	40	+ -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}
	41	+ -DCMAKE_C_COMPILER=${CMAKE_C_COMPILER}
	42	+ -DCMAKE_CXX_FLAGS=${GTEST_CXX_FLAGS}
	43	+ -DCMAKE_C_FLAGS=${GTEST_C_FLAGS}
	44	+ -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR>
	45	+ -DCMAKE_INSTALL_LIBDIR=lib
	46	+ -DBUILD_SHARED_LIBS=OFF
	47	+ -Dgtest_build_tests=OFF
	48	+ -Dgtest_build_samples=OFF
	49	+ -Dgmock_build_tests=OFF
	50	+
	51	+ BUILD_COMMAND $(MAKE)
	52	+ INSTALL_COMMAND $(MAKE) install
	53	+
	54	+ LOG_DOWNLOAD ON
	55	+ LOG_CONFIGURE ON
	56	+ LOG_BUILD ON
	57	+ LOG_INSTALL ON
	58	+)
	59	+
	60	+# Create imported targets (matching ops-math)
	61	+set(GTEST_INCLUDE_DIR ${GTEST_INSTALL_DIR}/include)
	62	+
	63	+# Ensure include directory exists for target_link_libraries
	64	+file(MAKE_DIRECTORY ${GTEST_INCLUDE_DIR})
	65	+
	66	+# gtest
	67	+add_library(gtest STATIC IMPORTED GLOBAL)
	68	+set_target_properties(gtest PROPERTIES
	69	+ IMPORTED_LOCATION ${GTEST_INSTALL_DIR}/lib/libgtest.a
	70	+ INTERFACE_INCLUDE_DIRECTORIES ${GTEST_INCLUDE_DIR}
	71	+)
	72	+add_dependencies(gtest third_party_gtest)
	73	+
	74	+# gtest_main
	75	+add_library(gtest_main STATIC IMPORTED GLOBAL)
	76	+set_target_properties(gtest_main PROPERTIES
	77	+ IMPORTED_LOCATION ${GTEST_INSTALL_DIR}/lib/libgtest_main.a
	78	+ INTERFACE_INCLUDE_DIRECTORIES ${GTEST_INCLUDE_DIR}
	79	+)
	80	+add_dependencies(gtest_main third_party_gtest)
	81	+
	82	+# gmock
	83	+add_library(gmock STATIC IMPORTED GLOBAL)
	84	+set_target_properties(gmock PROPERTIES
	85	+ IMPORTED_LOCATION ${GTEST_INSTALL_DIR}/lib/libgmock.a
	86	+ INTERFACE_INCLUDE_DIRECTORIES ${GTEST_INCLUDE_DIR}
	87	+)
	88	+add_dependencies(gmock third_party_gtest)
	89	+
	90	+# gmock_main
	91	+add_library(gmock_main STATIC IMPORTED GLOBAL)
	92	+set_target_properties(gmock_main PROPERTIES
	93	+ IMPORTED_LOCATION ${GTEST_INSTALL_DIR}/lib/libgmock_main.a
	94	+ INTERFACE_INCLUDE_DIRECTORIES ${GTEST_INCLUDE_DIR}
	95	+)
	96	+add_dependencies(gmock_main third_party_gtest)
	97	+
	98	+# Create interface library for UT (matching ops-math intf_llt_pub_asan_cxx17)
	99	+add_library(intf_llt_pub_asan_cxx17 INTERFACE)
	100	+target_include_directories(intf_llt_pub_asan_cxx17 INTERFACE
	101	+ ${GTEST_INCLUDE_DIR}
	102	+)
	103	+target_compile_definitions(intf_llt_pub_asan_cxx17 INTERFACE
	104	+ _GLIBCXX_USE_CXX11_ABI=0
	105	+ CFG_BUILD_DEBUG
	106	+)
	107	+target_compile_options(intf_llt_pub_asan_cxx17 INTERFACE
	108	+ -g
	109	+ --coverage
	110	+ -fprofile-arcs
	111	+ -ftest-coverage
	112	+ -w
	113	+ -std=c++17
	114	+ -fPIC
	115	+)
	116	+target_link_options(intf_llt_pub_asan_cxx17 INTERFACE
	117	+ -fprofile-arcs
	118	+ -ftest-coverage
	119	+)
	120	+target_link_libraries(intf_llt_pub_asan_cxx17 INTERFACE
	121	+ gcov
	122	+ pthread
	123	+)
	124	+
	125	+message(STATUS "")
	126	+message(STATUS "=== Google Test Build Configuration (ops-math mode) ===")
	127	+message(STATUS " Version: ${GTEST_VERSION}")
	128	+message(STATUS " Install Dir: ${GTEST_INSTALL_DIR}")
	129	+message(STATUS " ABI: OLD (_GLIBCXX_USE_CXX11_ABI=0)")
	130	+message(STATUS " Reason: Match CANN libraries (libplatform.so, libtiling_api.a)")
	131	+message(STATUS "=====================================================")
	132	+message(STATUS "")

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/any_value.h+106-0

@@ -0,0 +1,106 @@
	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	+ {
	15	+ VT_STRING = 1,
	16	+ VT_FLOAT = 2,
	17	+ VT_BOOL = 3,
	18	+ VT_INT = 4,
	19	+ VT_LIST_LIST_INT = 10,
	20	+ VT_LIST_BASE = 1000,
	21	+
	22	+ VT_LIST_FLOAT = static_cast<int32_t>(VT_LIST_BASE) + static_cast<int32_t>(VT_FLOAT),
	23	+ VT_LIST_BOOL = static_cast<int32_t>(VT_LIST_BASE) + static_cast<int32_t>(VT_BOOL),
	24	+ VT_LIST_INT = static_cast<int32_t>(VT_LIST_BASE) + static_cast<int32_t>(VT_INT),
	25	+ };
	26	+
	27	+ AnyValue(ValueType type, const std::shared_ptr<void>& valuePtr) : 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>
	34	+ static inline AnyValue CreateFrom(const T& value);
	35	+
	36	+ ValueType type_;
	37	+ std::shared_ptr<void> valuePtr_;
	38	+};
	39	+
	40	+template <>
	41	+inline AnyValue AnyValue::CreateFrom<std::string>(const std::string& value)
	42	+{
	43	+ auto valuePtr = new std::string;
	44	+ *valuePtr = value;
	45	+ return AnyValue(VT_STRING, std::shared_ptr<void>(valuePtr));
	46	+}
	47	+
	48	+template <>
	49	+inline AnyValue AnyValue::CreateFrom<float>(const float& value)
	50	+{
	51	+ auto valuePtr = new float;
	52	+ *valuePtr = value;
	53	+ return AnyValue(VT_FLOAT, std::shared_ptr<void>(valuePtr));
	54	+}
	55	+
	56	+template <>
	57	+inline AnyValue AnyValue::CreateFrom<bool>(const bool& value)
	58	+{
	59	+ auto valuePtr = new bool;
	60	+ *valuePtr = value;
	61	+ return AnyValue(VT_BOOL, std::shared_ptr<void>(valuePtr));
	62	+}
	63	+
	64	+template <>
	65	+inline AnyValue AnyValue::CreateFrom<int64_t>(const int64_t& value)
	66	+{
	67	+ auto valuePtr = new int64_t;
	68	+ *valuePtr = value;
	69	+ return AnyValue(VT_INT, std::shared_ptr<void>(valuePtr));
	70	+}
	71	+
	72	+template <>
	73	+inline AnyValue AnyValue::CreateFrom<std::vector<float>>(const std::vector<float>& value)
	74	+{
	75	+ auto valuePtr = new std::vector<float>;
	76	+ *valuePtr = value;
	77	+ return AnyValue(VT_LIST_FLOAT, std::shared_ptr<void>(valuePtr));
	78	+}
	79	+
	80	+template <>
	81	+inline AnyValue AnyValue::CreateFrom<std::vector<bool>>(const std::vector<bool>& value)
	82	+{
	83	+ auto valuePtr = new std::vector<bool>;
	84	+ *valuePtr = value;
	85	+ return AnyValue(VT_LIST_BOOL, std::shared_ptr<void>(valuePtr));
	86	+}
	87	+
	88	+template <>
	89	+inline AnyValue AnyValue::CreateFrom<std::vector<int64_t>>(const std::vector<int64_t>& value)
	90	+{
	91	+ auto valuePtr = new std::vector<int64_t>;
	92	+ *valuePtr = value;
	93	+ return AnyValue(VT_LIST_INT, std::shared_ptr<void>(valuePtr));
	94	+}
	95	+
	96	+template <>
	97	+inline AnyValue AnyValue::CreateFrom<std::vector<std::vector<int64_t>>>(const std::vector<std::vector<int64_t>>& value)
	98	+{
	99	+ auto valuePtr = new std::vector<std::vector<int64_t>>;
	100	+ *valuePtr = value;
	101	+ return AnyValue(VT_LIST_LIST_INT, std::shared_ptr<void>(valuePtr));
	102	+}
	103	+} // namespace Math
	104	+} // namespace Ops
	105	+
	106	+#endif // OPS_MATH_DEV_TESTS_UT_COMMON_ANY_VALUE_H

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/infershape_case_executor.cpp+101-0

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

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/infershape_case_executor.h+10-0

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

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/infershape_context_faker.cpp+42-0

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

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/infershape_context_faker.h+124-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/tiling_case_executor.cpp+295-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/tiling_case_executor.h+31-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/tiling_context_faker.cpp+71-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/common/tiling_context_faker.h+194-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/op_host/CMakeLists.txt+102-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/op_host/test_confusion_matrix_tiling.cpp+92-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/op_host/test_op_host_main.cpp+50-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/op_kernel/CMakeLists.txt+97-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/op_kernel/confusion_matrix_data/compare_data.py+275-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/op_kernel/confusion_matrix_data/gen_data.py+58-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/op_kernel/confusion_matrix_tiling.h+43-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/op_kernel/test_confusion_matrix.cpp+99-0

A01_official/cann-ops-ladder-2026/July/confusion_matrix/submissions/sohnkee/tests/ut/run.sh+92-0