Star1736
1647
Code
Code
Issues479
Pull Requests553
Actions
Discussion
Wiki
Analysis
Star1736
1647
Merged
feat: ConfusionMatrix 算子下一代实现 #5444
shengleAdd on Sep 8
feat: ConfusionMatrix 算子下一代实现 #5444
Merged
shengleAdd on Sep 8
Merged fromSeanDictionary/ops-mathfeat/confusion-matrixintocann/ops-mathmaster
Total: 25 files changed+3036-38
Mcommon/inc/op_graph/op_math_proto_extend.h+0-38
@@ -425,44 +425,6 @@ REG_OP(ExpandDims)
425 .ATTR(keep_dim, Bool, false)425 .ATTR(keep_dim, Bool, false)
426 .OP_END_FACTORY_REG(Bitcast)426 .OP_END_FACTORY_REG(Bitcast)
427 427 
428- /**
429- * @brief Computes the confusion matrix from predictions and labels .
430- 
431- * @par Inputs:
432- * Three inputs, including:
433- * @li labels: A Tensor. Must be one of the following types: float16, float32,
434- * int32, int8, uint8. 1D. Has format ND.
435- * @li predictions: A Tensor. Must be one of the following types: float16,
436- * float32, int32, int8, uint8. 1D. Has format ND.
437- * @li weights: A optional Tensor. Must be one of the following types: float16, float32,
438- * int32, int8, uint8. 1D. Has format ND. \n
439- 
440- * @par Attributes:
441- * @li num_classes: An integer for the shape of the output matrix.
442- * @li dtype: Data type of the confusion matrix. \n
443- 
444- * @par Outputs:
445- * y: A Tensor. 1D. Has format ND. Has the same type and format as input "labels" . \n
446- 
447- * @attention Constraints:
448- * @li "weights", "labels", and "predictions" are 1D tensors.
449- * @li The output is with shape (num_classes, num_classes),
450- * where, 1 <= num_classes <= 4096 . \n
451- 
452- * @see Region()
453- 
454- * @par Third-party framework compatibility
455- * Compatible with the TensorFlow operator ConfusionMatrix.
456- */
457- REG_OP(ConfusionMatrix)
458- .INPUT(labels, TensorType({DT_FLOAT, DT_INT32, DT_FLOAT16, DT_INT8, DT_UINT8}))
459- .INPUT(predictions, TensorType({DT_FLOAT, DT_INT32, DT_FLOAT16, DT_INT8, DT_UINT8}))
460- .OPTIONAL_INPUT(weights, TensorType({DT_FLOAT, DT_INT32, DT_FLOAT16, DT_INT8, DT_UINT8}))
461- .OUTPUT(y, TensorType({DT_FLOAT, DT_INT32, DT_FLOAT16, DT_INT8, DT_UINT8}))
462- .REQUIRED_ATTR(num_classes, Int)
463- .REQUIRED_ATTR(dtype, String)
464- .OP_END_FACTORY_REG(ConfusionMatrix)
465- 
466 /**428 /**
467 * @brief Computes the cumulative product of the tensor "x" along "axis" .429 * @brief Computes the cumulative product of the tensor "x" along "axis" .
468 430 
Mdocs/zh/op_list.md+10-0
@@ -605,6 +605,16 @@
605 <td>√</td>605 <td>√</td>
606 <td>AI Core/AI CPU</td>606 <td>AI Core/AI CPU</td>
607 <td>该算子对输入Tensor完成linear_cross运算。</td>607 <td>该算子对输入Tensor完成linear_cross运算。</td>
608+ </tr>
609+ <tr>
610+ <td>math</td>
611+ <td><a href="../../math/confusion_matrix/README.md">confusion_matrix</a></td>
612+ <td>√</td>
613+ <td>√</td>
614+ <td>×</td>
615+ <td>√</td>
616+ <td>AI Core</td>
617+ <td>计算分类任务的混淆矩阵，统计每个(标签, 预测)对的出现次数或加权累加和，输出num_classes×num_classes矩阵。</td>
608 </tr>618 </tr>
609 <tr>619 <tr>
610 <td>math</td>620 <td>math</td>
Amath/confusion_matrix/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+# 设置算子定义时支持的芯片类型
12+set(SUPPORT_COMPUTE_UNIT "ascend950")
13+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14+set(SUPPORT_TILING_DIR "arch35")
15+add_all_modules_sources(OPTYPE confusion_matrix ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Amath/confusion_matrix/README.md+106-0
@@ -0,0 +1,106 @@
1+# ConfusionMatrix
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR&950DT系列产品</term> | √ |
G
GgubaochengSep 29

[Major] README 产品支持表与实际注册范围不符（仅注册 ascend950）

README 产品支持表声明 6 大产品（950PR&950DT、A3、A2、200I/500 A2、推理系列、训练系列）全部支持，但 confusion_matrix_def.cpp 仅 AddConfig("ascend950")，tiling 实现仅位于 op_host/arch35/（DAV_3510 专属）且 binary json 仅在 op_host/config/ascend950/ 下。其余 5 类产品既无算子注册也无对应架构 tiling，用户按文档在 A2/A3 等平台调用会直接失败。建议将产品支持表收敛为实际注册的 Ascend 950 系列，或补齐其它产品的注册与 arch22 tiling 后再声明支持。

（来自自动检视代理，以 gubaocheng 身份提交）

likedislike
shengle
shengle
Oct 8 Comments：
8+| <term>Atlas A3系列产品</term> | √ |
9+| <term>Atlas A2系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2推理产品</term> | √ |
11+| <term>Atlas推理系列产品</term> | √ |
12+| <term>Atlas训练系列产品</term> | √ |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：计算分类任务的混淆矩阵。给定真实标签labels和预测标签predictions，统计每个(label, prediction)对的出现次数，生成num_classes×num_classes的矩阵。当weights不为空时，每次匹配的贡献为对应的权重值而非1。
17+ 
18+- 计算公式：
19+ 
20+ 如果指定了weights，则
21+ 
22+ $$
23+ y[labels_i, predictions_i] = y[labels_i, predictions_i] + weights_i
24+ $$
25+ 
26+ 否则：
27+ 
28+ $$
29+ y[labels_i, predictions_i] = y[labels_i, predictions_i] + 1
30+ $$
31+ 
32+## 参数说明
33+ 
34+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
35+ <col style="width: 100px">
36+ <col style="width: 150px">
37+ <col style="width: 280px">
38+ <col style="width: 330px">
39+ <col style="width: 120px">
40+ </colgroup>
41+ <thead>
42+ <tr>
43+ <th>参数名</th>
44+ <th>输入/输出/属性</th>
45+ <th>描述</th>
46+ <th>数据类型</th>
47+ <th>数据格式</th>
48+ </tr></thead>
49+ <tbody>
50+ <tr>
51+ <td>labels</td>
52+ <td>输入</td>
53+ <td>真实标签，1维tensor，取值范围[0, num_classes)。</td>
54+ <td>INT8、INT32、UINT8、FLOAT16、FLOAT</td>
55+ <td>ND</td>
56+ </tr>
57+ <tr>
58+ <td>predictions</td>
59+ <td>输入</td>
60+ <td>预测标签，1维tensor，shape和dtype与labels一致。</td>
61+ <td>INT8、INT32、UINT8、FLOAT16、FLOAT</td>
62+ <td>ND</td>
63+ </tr>
64+ <tr>
65+ <td>weights</td>
66+ <td>输入</td>
67+ <td>每个样本的权重，可为空。shape与labels一致。</td>
68+ <td>FLOAT、FLOAT16、INT32、INT8、UINT8</td>
69+ <td>ND</td>
70+ </tr>
71+ <tr>
72+ <td>num_classes</td>
73+ <td>属性</td>
74+ <td>类别数量，取值范围[1, 4096]。</td>
75+ <td>INT</td>
76+ <td>-</td>
77+ </tr>
78+ <tr>
79+ <td>dtype</td>
80+ <td>属性</td>
81+ <td>输出tensor的数据类型。</td>
82+ <td>STRING</td>
83+ <td>-</td>
84+ </tr>
85+ <tr>
86+ <td>y</td>
87+ <td>输出</td>
88+ <td>混淆矩阵，shape为[num_classes, num_classes]。</td>
89+ <td>FLOAT、FLOAT16、INT32、INT8、UINT8</td>
90+ <td>ND</td>
91+ </tr>
92+ </tbody></table>
93+ 
94+## 约束说明
95+ 
96+- labels和predictions的shape必须一致。
97+- labels和predictions的dtype必须一致。
98+- weights不为空时，其shape必须与labels一致。
99+- labels和predictions中的值必须在[0, num_classes)范围内。
100+- num_classes取值范围为[1, 4096]。
101+ 
102+## 调用说明
103+ 
104+| 调用方式 | 调用样例 | 说明 |
105+|--------------|--------|------|
106+| 图模式调用 | [test_geir_confusion_matrix](./examples/test_geir_confusion_matrix.cpp) | 通过[算子IR](./op_graph/confusion_matrix_proto.h)构图方式调用ConfusionMatrix算子。 |
Amath/confusion_matrix/examples/test_geir_confusion_matrix.cpp+296-0
@@ -0,0 +1,296 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <fstream>
13+#include <string.h>
14+#include <stdint.h>
15+#include <vector>
16+#include <string>
17+#include <map>
18+#include "assert.h"
19+ 
20+#include "graph.h"
21+#include "types.h"
22+#include "tensor.h"
23+#include "ge_error_codes.h"
24+#include "ge_api_types.h"
25+#include "ge_api.h"
26+#include "array_ops.h"
27+#include "ge_ir_build.h"
28+ 
29+#include "nn_other.h"
30+#include "../op_graph/confusion_matrix_proto.h"
31+ 
32+#define FAILED -1
33+#define SUCCESS 0
34+ 
35+using namespace ge;
36+using std::map;
37+using std::string;
38+using std::vector;
39+ 
40+#define ADD_INPUT(inputIndex, inputName, inputDtype, inputShape) \
41+ vector<int64_t> placeholder##inputIndex##_shape = inputShape; \
42+ auto placeholder##inputIndex = op::Data("placeholder" + std::to_string(inputIndex)); \
43+ placeholder##inputIndex.set_attr_index(0); \
44+ TensorDesc placeholder##inputIndex##_desc = TensorDesc(ge::Shape(placeholder##inputIndex##_shape), FORMAT_ND, \
45+ inputDtype); \
46+ placeholder##inputIndex##_desc.SetPlacement(ge::kPlacementHost); \
47+ placeholder##inputIndex##_desc.SetFormat(FORMAT_ND); \
48+ Tensor tensor_placeholder##inputIndex; \
49+ ret = GenOnesData(placeholder##inputIndex##_shape, tensor_placeholder##inputIndex, placeholder##inputIndex##_desc, \
50+ inputDtype, 1); \
51+ if (ret != SUCCESS) { \
52+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
53+ return FAILED; \
54+ } \
55+ placeholder##inputIndex.update_input_desc_x(placeholder##inputIndex##_desc); \
56+ input.push_back(tensor_placeholder##inputIndex); \
57+ graph.AddOp(placeholder##inputIndex); \
58+ op1.set_input_##inputName(placeholder##inputIndex); \
59+ inputs.push_back(placeholder##inputIndex);
60+ 
61+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
62+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
63+ op1.update_output_desc_##outputName(outputName##outputIndex##_desc);
64+ 
65+string GetTime()
66+{
67+ time_t timep;
68+ time(&timep);
69+ char tmp[64];
70+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
71+ return tmp;
72+}
73+ 
74+uint32_t GetDataTypeSize(DataType dt)
75+{
76+ uint32_t dilation = 1;
77+ if (dt == ge::DT_FLOAT) {
78+ dilation = 4;
79+ } else if (dt == ge::DT_FLOAT16) {
80+ dilation = 2;
81+ } else if (dt == ge::DT_BF16) {
82+ dilation = 2;
83+ } else if (dt == ge::DT_INT16) {
84+ dilation = 2;
85+ } else if (dt == ge::DT_UINT16) {
86+ dilation = 2;
87+ } else if (dt == ge::DT_INT32) {
88+ dilation = 4;
89+ } else if (dt == ge::DT_UINT32) {
90+ dilation = 4;
91+ } else if (dt == ge::DT_INT64) {
92+ dilation = 8;
93+ } else if (dt == ge::DT_UINT64) {
94+ dilation = 8;
95+ } else if (dt == ge::DT_INT8) {
96+ dilation = 1;
97+ }
98+ return dilation;
99+}
100+ 
101+int32_t GenOnesData(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type,
102+ int value)
103+{
104+ input_tensor_desc.SetRealDimCnt(shapes.size());
105+ size_t size = 1;
106+ for (uint32_t i = 0; i < shapes.size(); i++) {
107+ size *= shapes[i];
108+ }
109+ size_t data_len = size * GetDataTypeSize(data_type);
110+ uint8_t* pData = new (std::nothrow) uint8_t[data_len];
111+ if (pData == nullptr) {
112+ return FAILED;
113+ }
114+ // 按 dtype 实际宽度逐元素写入：固定按 int32 写会在 dtype < 4 字节时堆越界，
115+ // dtype > 4 字节时留下未初始化数据，且 float 会被写成 1.4e-45 而非 1.0
116+ for (size_t i = 0; i < size; ++i) {
117+ switch (data_type) {
118+ case ge::DT_FLOAT:
119+ *(reinterpret_cast<float*>(pData) + i) = static_cast<float>(value);
120+ break;
121+ case ge::DT_FLOAT16: // 1.0 的 FP16 位型
122+ *(reinterpret_cast<uint16_t*>(pData) + i) = 0x3C00;
123+ break;
124+ case ge::DT_BF16: // 1.0 的 BF16 位型
125+ *(reinterpret_cast<uint16_t*>(pData) + i) = 0x3F80;
126+ break;
127+ case ge::DT_INT8:
128+ *(reinterpret_cast<int8_t*>(pData) + i) = static_cast<int8_t>(value);
129+ break;
130+ case ge::DT_UINT8:
131+ *(reinterpret_cast<uint8_t*>(pData) + i) = static_cast<uint8_t>(value);
132+ break;
133+ case ge::DT_INT16:
134+ *(reinterpret_cast<int16_t*>(pData) + i) = static_cast<int16_t>(value);
135+ break;
136+ case ge::DT_UINT16:
137+ *(reinterpret_cast<uint16_t*>(pData) + i) = static_cast<uint16_t>(value);
138+ break;
139+ case ge::DT_INT32:
140+ *(reinterpret_cast<int32_t*>(pData) + i) = value;
141+ break;
142+ case ge::DT_UINT32:
143+ *(reinterpret_cast<uint32_t*>(pData) + i) = static_cast<uint32_t>(value);
144+ break;
145+ case ge::DT_INT64:
146+ *(reinterpret_cast<int64_t*>(pData) + i) = static_cast<int64_t>(value);
147+ break;
148+ case ge::DT_UINT64:
149+ *(reinterpret_cast<uint64_t*>(pData) + i) = static_cast<uint64_t>(value);
150+ break;
151+ default:
152+ printf("%s - ERROR - [XIR]: GenOnesData unsupported dtype %d\n", GetTime().c_str(),
153+ static_cast<int32_t>(data_type));
154+ return FAILED;
155+ }
156+ }
157+ input_tensor = Tensor(input_tensor_desc, pData, data_len);
158+ return SUCCESS;
159+}
160+ 
161+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
162+{
163+ FILE* fp = fopen(bin_file.c_str(), "w");
164+ if (fp == nullptr) {
165+ return FAILED;
166+ }
167+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
168+ fclose(fp);
169+ return SUCCESS;
170+}
171+ 
172+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs,
173+ std::vector<Operator>& outputs, Graph& graph)
174+{
175+ Status ret = SUCCESS;
176+ // ConfusionMatrix: 3 inputs (labels, predictions, weights), 1 output
177+ auto op1 = op::ConfusionMatrix("confusion_matrix1");
178+ 
179+ // labels: shape [8], dtype int32
180+ std::vector<int64_t> labelsShape = {8};
181+ ADD_INPUT(1, labels, ge::DT_INT32, labelsShape);
182+ 
183+ // predictions: shape [8], dtype int32
184+ std::vector<int64_t> predictionsShape = {8};
185+ ADD_INPUT(2, predictions, ge::DT_INT32, predictionsShape);
186+ 
187+ // weights: shape [8], dtype float32
188+ std::vector<int64_t> weightsShape = {8};
189+ ADD_INPUT(3, weights, ge::DT_FLOAT, weightsShape);
190+ 
191+ // output: shape [num_classes, num_classes], dtype float32
192+ std::vector<int64_t> outputShape = {4, 4};
193+ ADD_OUTPUT(1, y, ge::DT_FLOAT, outputShape);
194+ 
195+ // attrs: num_classes (required int), dtype (required string)
196+ op1.set_attr_num_classes(4);
197+ op1.set_attr_dtype("float32");
198+ 
199+ outputs.push_back(op1);
200+ return SUCCESS;
201+}
202+ 
203+int main(int argc, char* argv[])
204+{
205+ const char* graph_name = "tc_ge_irrun_test";
206+ Graph graph(graph_name);
207+ std::vector<ge::Tensor> input;
208+ 
209+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
210+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
211+ Status ret = ge::GEInitialize(global_options);
212+ if (ret != SUCCESS) {
213+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
214+ return FAILED;
215+ }
216+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
217+ 
218+ std::vector<Operator> inputs{};
219+ std::vector<Operator> outputs{};
220+ 
221+ DataType inDtype = DT_FLOAT;
222+ 
223+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
224+ if (ret != SUCCESS) {
225+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
226+ return FAILED;
227+ }
228+ 
229+ if (!inputs.empty() && !outputs.empty()) {
230+ graph.SetInputs(inputs).SetOutputs(outputs);
231+ }
232+ 
233+ std::map<AscendString, AscendString> build_options = {};
234+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
235+ ge::Session* session = new Session(build_options);
236+ 
237+ if (session == nullptr) {
238+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
239+ return FAILED;
240+ }
241+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
242+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
243+ 
244+ std::map<AscendString, AscendString> graph_options = {};
245+ uint32_t graph_id = 0;
246+ ret = session->AddGraph(graph_id, graph, graph_options);
247+ if (ret != SUCCESS) {
248+ printf("%s - ERROR - [XIR]: Add graph failed\n", GetTime().c_str());
249+ delete session;
250+ GEFinalize();
251+ return FAILED;
252+ }
253+ 
254+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
255+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
256+ std::vector<ge::Tensor> output;
257+ ret = session->RunGraph(graph_id, input, output);
258+ if (ret != SUCCESS) {
259+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
260+ delete session;
261+ GEFinalize();
262+ return FAILED;
263+ }
264+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
265+ 
266+ int input_num = input.size();
267+ for (int i = 0; i < input_num; i++) {
268+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
269+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
270+ uint8_t* input_data_i = input[i].GetData();
271+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
272+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
273+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
274+ }
275+ 
276+ int output_num = output.size();
277+ for (int i = 0; i < output_num; i++) {
278+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
279+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
280+ uint8_t* output_data_i = output[i].GetData();
281+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
282+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
283+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
284+ }
285+ 
286+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
287+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
288+ delete session;
289+ ret = ge::GEFinalize();
290+ if (ret != SUCCESS) {
291+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
292+ return FAILED;
293+ }
294+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
295+ return SUCCESS;
296+}
Amath/confusion_matrix/op_graph/confusion_matrix_graph_infer.cpp+60-0
@@ -0,0 +1,60 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_graph_infer.cpp
13+ * \brief Infer data type implementation for ConfusionMatrix operator
14+ */
15+ 
16+#include "register/op_impl_registry.h"
17+#include "log/log.h"
18+ 
19+#include <string>
20+ 
21+using namespace ge;
22+namespace ops {
23+static constexpr size_t kAttrIndexDtype = 1U;
24+static constexpr size_t kOutputIndexY = 0U;
25+ 
26+static ge::graphStatus InferDataTypeForConfusionMatrix(gert::InferDataTypeContext* context)
27+{
28+ OP_LOGI(context->GetNodeName(), "Begin to do InferDataTypeForConfusionMatrix");
29+ 
30+ auto attrs = context->GetAttrs();
31+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
32+ 
33+ const char* dtypeStr = attrs->GetAttrPointer<char>(kAttrIndexDtype);
34+ OP_CHECK_NULL_WITH_CONTEXT(context, dtypeStr);
35+ 
36+ ge::DataType outputDtype = ge::DT_FLOAT;
37+ std::string dtypeString(dtypeStr);
38+ if (dtypeString == "float32") {
39+ outputDtype = ge::DT_FLOAT;
40+ } else if (dtypeString == "int32") {
41+ outputDtype = ge::DT_INT32;
42+ } else if (dtypeString == "int8") {
43+ outputDtype = ge::DT_INT8;
44+ } else if (dtypeString == "float16") {
45+ outputDtype = ge::DT_FLOAT16;
46+ } else if (dtypeString == "uint8") {
47+ outputDtype = ge::DT_UINT8;
48+ } else {
49+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "dtype", dtypeStr,
50+ "dtype must be one of: float32, int32, int8, float16, uint8.");
51+ return ge::GRAPH_FAILED;
52+ }
53+ 
54+ context->SetOutputDataType(kOutputIndexY, outputDtype);
55+ OP_LOGI(context->GetNodeName(), "End to do InferDataTypeForConfusionMatrix");
56+ return ge::GRAPH_SUCCESS;
57+}
58+ 
59+IMPL_OP(ConfusionMatrix).InferDataType(InferDataTypeForConfusionMatrix);
60+} // namespace ops
Amath/confusion_matrix/op_graph/confusion_matrix_proto.h+64-0
@@ -0,0 +1,64 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_proto.h
13+ * \brief
14+ */
15+#ifndef OPS_OP_PROTO_INC_CONFUSION_MATRIX_H_
16+#define OPS_OP_PROTO_INC_CONFUSION_MATRIX_H_
17+ 
18+#include "graph/operator_reg.h"
19+#include "graph/types.h"
20+ 
21+namespace ge {
22+ 
23+/**
24+*@brief Computes the confusion matrix from predictions and labels .
25+ 
26+*@par Inputs:
27+*Three inputs, including:
28+*@li labels: A Tensor. Must be one of the following types: float16, float32,
29+*int32, int8, uint8. 1D. Has format ND.
30+*@li predictions: A Tensor. Must be one of the following types: float16,
31+*float32, int32, int8, uint8. 1D. Has format ND.
32+*@li weights: A optional Tensor. Must be one of the following types: float16, float32,
33+*int32, int8, uint8. 1D. Has format ND. \n
34+ 
35+*@par Attributes:
36+*@li num_classes: An integer for the shape of the output matrix.
37+*@li dtype: Data type of the confusion matrix. \n
38+ 
39+*@par Outputs:
40+*y: A Tensor. 1D. Has format ND. Has the same type and format as input "labels" . \n
41+ 
42+*@attention Constraints:
43+*@li "weights", "labels", and "predictions" are 1D tensors.
44+*@li The output is with shape (num_classes, num_classes),
45+*where, 1 <= num_classes <= 4096 . \n
46+ 
47+*@par Third-party framework compatibility
48+*Compatible with the TensorFlow operator ConfusionMatrix.
49+*/
50+#ifndef OPS_PROTO_DEF_CONFUSIONMATRIX
51+#define OPS_PROTO_DEF_CONFUSIONMATRIX
52+REG_OP(ConfusionMatrix)
53+ .INPUT(labels, TensorType({DT_FLOAT, DT_INT32, DT_FLOAT16, DT_INT8, DT_UINT8}))
54+ .INPUT(predictions, TensorType({DT_FLOAT, DT_INT32, DT_FLOAT16, DT_INT8, DT_UINT8}))
55+ .OPTIONAL_INPUT(weights, TensorType({DT_FLOAT, DT_INT32, DT_FLOAT16, DT_INT8, DT_UINT8}))
56+ .OUTPUT(y, TensorType({DT_FLOAT, DT_INT32, DT_FLOAT16, DT_INT8, DT_UINT8}))
57+ .REQUIRED_ATTR(num_classes, Int)
58+ .REQUIRED_ATTR(dtype, String)
59+ .OP_END_FACTORY_REG(ConfusionMatrix)
60+#endif // OPS_PROTO_DEF_CONFUSIONMATRIX
61+ 
62+} // namespace ge
63+ 
64+#endif // OPS_OP_PROTO_INC_CONFUSION_MATRIX_H_
Amath/confusion_matrix/op_host/arch35/confusion_matrix_tiling.cpp+358-0
@@ -0,0 +1,358 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software; you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_tiling.cpp
13+ * \brief confusion_matrix_tiling file
14+ */
15+ 
16+#include <iostream>
17+#include <cstring>
18+#include <string>
19+#include "register/op_impl_registry.h"
20+#include "platform/platform_info.h"
21+#include "op_host/tiling_base_util.h"
22+#include "log/log.h"
23+#include "confusion_matrix_tiling.h"
24+#include "util/math_util.h"
25+#include "util/const_util.h"
26+ 
27+namespace optiling {
28+ge::graphStatus ConfusionMatrixTiling::Init()
29+{
30+ OP_LOGD(context_->GetNodeName(), "ConfusionMatrixTiling init enter.");
31+ if (tilingData_ == nullptr) {
32+ tilingData_ = context_->GetTilingData<ConfusionMatrixTilingData>();
33+ OP_CHECK_IF(tilingData_ == nullptr, OP_LOGE(context_->GetNodeName(), "get tilingdata ptr failed"),
34+ return ge::GRAPH_FAILED);
35+ }
36+ OP_CHECK_IF((memset_s(tilingData_, sizeof(ConfusionMatrixTilingData), 0, sizeof(ConfusionMatrixTilingData)) != EOK),
37+ OP_LOGE(context_->GetNodeName(), "memset tilingdata failed"), return ge::GRAPH_FAILED);
38+ OP_LOGD(context_->GetNodeName(), "ConfusionMatrixTiling init exit.");
39+ return ge::GRAPH_SUCCESS;
40+}
41+ 
42+ge::graphStatus ConfusionMatrixTiling::ConfusionMatrixGetPlatformData(
43+ const AscendCConfusionMatrixCompileInfo* compileInfo)
44+{
45+ coreNum_ = compileInfo->totalCoreNum;
46+ ubSize_ = static_cast<int64_t>(compileInfo->totalUbSize - SIMD_SIMT_DCACHE_SIZE);
47+ isDetermine_ = context_->GetDeterministic() == 1 ? 1 : 0;
48+ OP_LOGI(context_->GetNodeName(),
49+ "ConfusionMatrixGetPlatformData ubSize is %ld, coreNum_ is %ld, isDetermine is %ld", ubSize_, coreNum_,
50+ isDetermine_);
51+ return ge::GRAPH_SUCCESS;
52+}
53+ 
54+ge::graphStatus ConfusionMatrixTiling::CheckShape()
55+{
56+ auto labelsShape = context_->GetInputShape(INPUT_IDX_LABELS);
57+ OP_CHECK_NULL_WITH_CONTEXT(context_, labelsShape);
58+ OP_CHECK_IF(labelsShape->GetStorageShape().GetDimNum() != DIM_1,
59+ OP_LOGE_FOR_INVALID_SHAPEDIM(context_->GetNodeName(), "labels",
60+ std::to_string(labelsShape->GetStorageShape().GetDimNum()).c_str(), "1D"),
61+ return ge::GRAPH_FAILED);
62+ 
63+ auto predictionsShape = context_->GetInputShape(INPUT_IDX_PREDICTIONS);
64+ OP_CHECK_NULL_WITH_CONTEXT(context_, predictionsShape);
65+ OP_CHECK_IF(
66+ predictionsShape->GetStorageShape().GetDimNum() != DIM_1,
67+ OP_LOGE_FOR_INVALID_SHAPEDIM(context_->GetNodeName(), "predictions",
68+ std::to_string(predictionsShape->GetStorageShape().GetDimNum()).c_str(), "1D"),
69+ return ge::GRAPH_FAILED);
70+ 
71+ auto labelsShapeSize = labelsShape->GetStorageShape().GetShapeSize();
72+ auto predictionsShapeSize = predictionsShape->GetStorageShape().GetShapeSize();
73+ OP_CHECK_IF(labelsShapeSize != predictionsShapeSize,
74+ OP_LOGE_FOR_INVALID_SHAPESIZE_WITH_REASON(
75+ context_->GetNodeName(), "labels and predictions",
76+ (std::to_string(labelsShapeSize) + " and " + std::to_string(predictionsShapeSize)).c_str(),
77+ "The shapes of labels and predictions must be the same"),
78+ return ge::GRAPH_FAILED);
79+ 
80+ auto weightsShape = context_->GetInputShape(INPUT_IDX_WEIGHTS);
81+ if (weightsShape != nullptr) {
82+ OP_CHECK_IF(
83+ weightsShape->GetStorageShape().GetDimNum() != DIM_1,
84+ OP_LOGE_FOR_INVALID_SHAPEDIM(context_->GetNodeName(), "weights",
85+ std::to_string(weightsShape->GetStorageShape().GetDimNum()).c_str(), "1D"),
86+ return ge::GRAPH_FAILED);
87+ auto weightsShapeSize = weightsShape->GetStorageShape().GetShapeSize();
88+ if (weightsShapeSize != 0) {
89+ OP_CHECK_IF(weightsShapeSize != labelsShapeSize,
90+ OP_LOGE_FOR_INVALID_SHAPESIZE_WITH_REASON(
91+ context_->GetNodeName(), "weights",
92+ (std::to_string(weightsShapeSize) + " and " + std::to_string(labelsShapeSize)).c_str(),
93+ "The shape of weights must be the same as labels when provided"),
94+ return ge::GRAPH_FAILED);
95+ }
96+ }
97+ 
98+ auto yShape = context_->GetOutputShape(OUTPUT_IDX_Y);
99+ OP_CHECK_NULL_WITH_CONTEXT(context_, yShape);
100+ OP_CHECK_IF(yShape->GetStorageShape().GetDimNum() != DIM_2,
101+ OP_LOGE_FOR_INVALID_SHAPEDIM(context_->GetNodeName(), "y",
102+ std::to_string(yShape->GetStorageShape().GetDimNum()).c_str(), "2D"),
103+ return ge::GRAPH_FAILED);
104+ 
105+ return ge::GRAPH_SUCCESS;
106+}
107+ 
108+ge::graphStatus ConfusionMatrixTiling::CheckDtype()
109+{
110+ auto attrs = context_->GetAttrs();
111+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
112+ 
113+ const char* dtypePtr = attrs->GetAttrPointer<char>(ATTR_IDX_DTYPE);
114+ OP_CHECK_NULL_WITH_CONTEXT(context_, dtypePtr);
115+ std::string dtypeStr(dtypePtr);
116+ 
117+ if (dtypeStr == "float32") {
118+ outputDtype_ = OUTPUT_DTYPE_FLOAT;
119+ sizeOfOutputDtype_ = static_cast<int64_t>(SIZE_DTYPE_FLOAT);
120+ } else if (dtypeStr == "int32") {
121+ outputDtype_ = OUTPUT_DTYPE_INT32;
122+ sizeOfOutputDtype_ = static_cast<int64_t>(SIZE_DTYPE_INT32);
123+ } else if (dtypeStr == "float16") {
124+ outputDtype_ = OUTPUT_DTYPE_FLOAT16;
125+ sizeOfOutputDtype_ = static_cast<int64_t>(SIZE_DTYPE_FLOAT16);
126+ } else if (dtypeStr == "int8") {
127+ outputDtype_ = OUTPUT_DTYPE_INT8;
128+ sizeOfOutputDtype_ = static_cast<int64_t>(SIZE_DTYPE_INT8);
129+ } else if (dtypeStr == "uint8") {
130+ outputDtype_ = OUTPUT_DTYPE_UINT8;
131+ sizeOfOutputDtype_ = static_cast<int64_t>(SIZE_DTYPE_UINT8);
132+ } else {
133+ OP_LOGE_FOR_INVALID_VALUE(context_->GetNodeName(), "dtype", dtypeStr.c_str(),
134+ "float32, int32, float16, int8 or uint8");
135+ return ge::GRAPH_FAILED;
136+ }
137+ 
138+ return ge::GRAPH_SUCCESS;
139+}
140+ 
141+ge::graphStatus ConfusionMatrixTiling::CheckInputParams()
142+{
143+ auto attrs = context_->GetAttrs();
144+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
145+ 
146+ const int64_t* numClassesPtr = attrs->GetAttrPointer<int64_t>(ATTR_IDX_NUM_CLASSES);
147+ OP_CHECK_NULL_WITH_CONTEXT(context_, numClassesPtr);
148+ numClasses_ = *numClassesPtr;
149+ OP_CHECK_IF((numClasses_ < NUM_CLASSES_MIN || numClasses_ > NUM_CLASSES_MAX),
150+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context_->GetNodeName(), "num_classes",
151+ std::to_string(numClasses_).c_str(),
152+ "num_classes must be in range [1, 4096]."),
153+ return ge::GRAPH_FAILED);
154+ 
155+ auto labelsShape = context_->GetInputShape(INPUT_IDX_LABELS);
156+ OP_CHECK_NULL_WITH_CONTEXT(context_, labelsShape);
157+ inputSize_ = labelsShape->GetStorageShape().GetShapeSize();
158+ 
159+ auto weightsShape = context_->GetInputShape(INPUT_IDX_WEIGHTS);
160+ if (weightsShape != nullptr) {
161+ int64_t weightsShapeSize = weightsShape->GetStorageShape().GetShapeSize();
162+ if (weightsShapeSize != 0) {
163+ isWeight_ = WEIGHT;
164+ }
165+ }
166+ OP_LOGI(context_->GetNodeName(), "inputSize is %ld, isWeight is %ld, numClasses is %ld", inputSize_, isWeight_,
167+ numClasses_);
168+ 
169+ outputSize_ = numClasses_ * numClasses_;
170+ OP_LOGI(context_->GetNodeName(), "outputSize is %ld", outputSize_);
171+ 
172+ return ge::GRAPH_SUCCESS;
173+}
174+ 
175+inline bool ConfusionMatrixTiling::IsMatchSimtBatchLoadMode()
176+{
177+ return (2 * inputSize_) > outputSize_ / GM_ATOMIC_ADD_FACTOR;
178+}
179+ 
180+ge::graphStatus ConfusionMatrixTiling::ComputeTilingStrategy()
181+{
182+ OP_LOGD(context_->GetNodeName(), "ComputeTilingStrategy enter.");
183+ ubNumCanUse_ = static_cast<int64_t>(ubSize_ / sizeOfOutputDtype_);
184+ 
185+ // For int8/uint8 output, asc_atomic_add does not support these types.
186+ // Force DETERMINE mode which uses direct += on disjoint output ranges
187+ // (no atomic operations needed).
188+ if (isDetermine_ || outputDtype_ == OUTPUT_DTYPE_INT8 || outputDtype_ == OUTPUT_DTYPE_UINT8) {
189+ schId_ = SCH_ID_SIMT_DETERMIN;
190+ return ComputeTilingSimtDetermine();
191+ }
192+ 
193+ if (outputSize_ < ubNumCanUse_) {
194+ schId_ = SCH_ID_SIMT_FULL_LOAD;
195+ } else if (IsMatchSimtBatchLoadMode()) {
196+ ubLoopNum_ = Ops::Base::CeilDiv(outputSize_, ubNumCanUse_);
197+ if (ubLoopNum_ <= MAX_UB_LOOP_FOR_BATCH) {
198+ schId_ = SCH_ID_SIMT_BATCH_LOAD;
199+ } else {
200+ schId_ = SCH_ID_SIMT_NOT_FULL_LOAD;
201+ }
202+ } else {
203+ schId_ = SCH_ID_SIMT_NOT_FULL_LOAD;
204+ }
205+ 
206+ return ComputeTilingSimtNotDetermine();
207+}
208+ 
209+ge::graphStatus ConfusionMatrixTiling::ComputeTilingSimtNotDetermine()
210+{
211+ OP_LOGD(context_->GetNodeName(), "ComputeTilingSimtNotDetermine enter.");
212+ ubLoopNum_ = Ops::Base::CeilDiv(outputSize_, ubNumCanUse_);
213+ if (inputSize_ > 0) {
214+ formerLength_ = Ops::Base::CeilDiv(inputSize_, coreNum_);
215+ OP_CHECK_IF(formerLength_ == 0, OP_LOGE(context_->GetNodeName(), "formerLength_ must not be 0."),
216+ return ge::GRAPH_FAILED);
217+ needXCoreNum_ = Ops::Base::CeilDiv(inputSize_, formerLength_);
218+ tailLength_ = inputSize_ - (needXCoreNum_ - 1) * formerLength_;
219+ } else {
220+ formerLength_ = 1;
221+ needXCoreNum_ = 1;
222+ tailLength_ = 0;
223+ }
224+ 
225+ clearYFactor_ = Ops::Base::CeilDiv(outputSize_, coreNum_);
226+ OP_CHECK_IF(clearYFactor_ == 0, OP_LOGE(context_->GetNodeName(), "clearYFactor_ must not be 0."),
227+ return ge::GRAPH_FAILED);
228+ clearYCoreNum_ = Ops::Base::CeilDiv(outputSize_, clearYFactor_);
229+ clearYTail_ = outputSize_ - (clearYCoreNum_ - 1) * clearYFactor_;
230+ needCoreNum_ = std::max(needXCoreNum_, clearYCoreNum_);
231+ OP_LOGD(context_->GetNodeName(), "ComputeTilingSimtNotDetermine end.");
232+ return ge::GRAPH_SUCCESS;
233+}
234+ 
235+ge::graphStatus ConfusionMatrixTiling::ComputeTilingSimtDetermine()
236+{
237+ OP_LOGD(context_->GetNodeName(), "ComputeTilingSimtDetermine enter.");
238+ binsFormerLength_ = Ops::Base::CeilDiv(outputSize_, coreNum_);
239+ OP_CHECK_IF(binsFormerLength_ == 0, OP_LOGE(context_->GetNodeName(), "binsFormerLength_ must not be 0."),
240+ return ge::GRAPH_FAILED);
241+ needBinsCoreNum_ = Ops::Base::CeilDiv(outputSize_, binsFormerLength_);
242+ binsTailLength_ = outputSize_ - (needBinsCoreNum_ - 1) * binsFormerLength_;
243+ needCoreNum_ = needBinsCoreNum_;
244+ OP_LOGD(context_->GetNodeName(), "ComputeTilingSimtDetermine end.");
245+ return ge::GRAPH_SUCCESS;
246+}
247+ 
248+void ConfusionMatrixTiling::PrintTilingData()
249+{
250+ OP_LOGI(context_->GetNodeName(),
251+ "ConfusionMatrix tilingData needCoreNum_ is %ld, numClasses is %ld,"
252+ "ubNumCanUse is %ld, ubLoopNum is %ld, needXCoreNum is %ld, formerLength is %ld, tailLength is %ld,"
253+ "clearYCoreNum is %ld, clearYFactor is %ld, clearYTail is %ld, binsFormerLength is %ld"
254+ "needBinsCoreNum is %ld, binsTailLength is %ld",
255+ needCoreNum_, tilingData_->numClasses, tilingData_->ubNumCanUse, tilingData_->ubLoopNum,
256+ tilingData_->needXCoreNum, tilingData_->formerLength, tilingData_->tailLength, tilingData_->clearYCoreNum,
257+ tilingData_->clearYFactor, tilingData_->clearYTail, tilingData_->binsFormerLength,
258+ tilingData_->needBinsCoreNum, tilingData_->binsTailLength);
259+ return;
260+}
261+ 
262+ge::graphStatus ConfusionMatrixTiling::SetTilingData()
263+{
264+ OP_LOGD(context_->GetNodeName(), "SetTilingData enter.");
265+ tilingData_->numClasses = numClasses_;
266+ tilingData_->inputSize = inputSize_;
267+ tilingData_->ubNumCanUse = ubNumCanUse_;
268+ tilingData_->ubLoopNum = ubLoopNum_;
269+ tilingData_->needXCoreNum = needXCoreNum_;
270+ tilingData_->formerLength = formerLength_;
271+ tilingData_->tailLength = tailLength_;
272+ tilingData_->clearYCoreNum = clearYCoreNum_;
273+ tilingData_->clearYFactor = clearYFactor_;
274+ tilingData_->clearYTail = clearYTail_;
275+ tilingData_->binsFormerLength = binsFormerLength_;
276+ tilingData_->needBinsCoreNum = needBinsCoreNum_;
277+ tilingData_->binsTailLength = binsTailLength_;
278+ // UB buffer 字节数由 host 一次算准，kernel 只透传（避免 host/内核两套算法）
279+ tilingData_->fullLoadBufSize = outputSize_ * sizeOfOutputDtype_;
280+ tilingData_->batchUbBufSize = ubNumCanUse_ * sizeOfOutputDtype_;
281+ OP_LOGI(context_->GetNodeName(), "schId is %ld, outputDtype is %ld, isWeight is %ld", schId_, outputDtype_,
282+ isWeight_);
283+ const uint64_t tilingKey = GET_TPL_TILING_KEY(schId_, outputDtype_, isWeight_);
284+ OP_LOGI(context_->GetNodeName(), "tilingKey is %ld", tilingKey);
285+ context_->SetTilingKey(tilingKey);
286+ context_->SetBlockDim(needCoreNum_);
287+ context_->SetLocalMemorySize(ubSize_);
288+ context_->SetScheduleMode(1); // SyncAll need set
289+ size_t* workspaces = context_->GetWorkspaceSizes(1);
290+ OP_CHECK_NULL_WITH_CONTEXT(context_, workspaces);
291+ workspaces[0] = WORK_SPACE_SIZE;
292+ PrintTilingData();
293+ return ge::GRAPH_SUCCESS;
294+}
295+ 
296+static ge::graphStatus Tiling4ConfusionMatrix(gert::TilingContext* context)
297+{
298+ auto compileInfo = reinterpret_cast<const AscendCConfusionMatrixCompileInfo*>(context->GetCompileInfo());
299+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
300+ ConfusionMatrixTiling tilingObject(context);
301+ if (tilingObject.Init() != ge::GRAPH_SUCCESS) {
302+ OP_LOGE(context->GetNodeName(), "Init failed.");
303+ return ge::GRAPH_FAILED;
304+ }
305+ if (tilingObject.ConfusionMatrixGetPlatformData(compileInfo) != ge::GRAPH_SUCCESS) {
306+ OP_LOGE(context->GetNodeName(), "ConfusionMatrixGetPlatformData return failed.");
307+ return ge::GRAPH_FAILED;
308+ }
309+ if (tilingObject.CheckShape() != ge::GRAPH_SUCCESS) {
310+ OP_LOGE(context->GetNodeName(), "CheckShape return failed.");
311+ return ge::GRAPH_FAILED;
312+ }
313+ if (tilingObject.CheckDtype() != ge::GRAPH_SUCCESS) {
314+ OP_LOGE(context->GetNodeName(), "CheckDtype return failed.");
315+ return ge::GRAPH_FAILED;
316+ }
317+ if (tilingObject.CheckInputParams() != ge::GRAPH_SUCCESS) {
318+ OP_LOGE(context->GetNodeName(), "CheckInputParams return failed.");
319+ return ge::GRAPH_FAILED;
320+ }
321+ if (tilingObject.ComputeTilingStrategy() != ge::GRAPH_SUCCESS) {
322+ OP_LOGE(context->GetNodeName(), "ComputeTilingStrategy return failed.");
323+ return ge::GRAPH_FAILED;
324+ }
325+ if (tilingObject.SetTilingData() != ge::GRAPH_SUCCESS) {
326+ OP_LOGE(context->GetNodeName(), "SetTilingData return failed.");
327+ return ge::GRAPH_FAILED;
328+ }
329+ return ge::GRAPH_SUCCESS;
330+}
331+ 
332+static ge::graphStatus TilingPrepare4ConfusionMatrix(gert::TilingParseContext* context)
333+{
334+ OP_LOGI(context->GetNodeName(), "TilingPrepare4ConfusionMatrix running.");
335+ auto compileInfo = context->GetCompiledInfo<AscendCConfusionMatrixCompileInfo>();
336+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
337+ auto platformInfo = context->GetPlatformInfo();
338+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
339+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
340+ compileInfo->totalCoreNum = static_cast<int32_t>(ascendcPlatform.GetCoreNumAiv());
341+ OP_CHECK_IF((compileInfo->totalCoreNum <= 0),
342+ OP_LOGE(context->GetNodeName(), "coreNum is invalid, must greater than zero"), return ge::GRAPH_FAILED);
343+ 
344+ uint64_t ubSizePlatForm = 0;
345+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm);
346+ compileInfo->totalUbSize = static_cast<int64_t>(ubSizePlatForm);
347+ OP_CHECK_IF((compileInfo->totalUbSize <= 0),
348+ OP_LOGE(context->GetNodeName(), "ubSize is invalid, must greater than zero"), return ge::GRAPH_FAILED);
349+ OP_LOGD(context->GetNodeName(), "totalUbSize is %lu.", compileInfo->totalUbSize);
350+ 
351+ return ge::GRAPH_SUCCESS;
352+}
353+ 
354+IMPL_OP_OPTILING(ConfusionMatrix)
355+ .Tiling(Tiling4ConfusionMatrix)
356+ .TilingParse<AscendCConfusionMatrixCompileInfo>(TilingPrepare4ConfusionMatrix);
357+ 
358+} // namespace optiling
Amath/confusion_matrix/op_host/arch35/confusion_matrix_tiling.h+121-0
@@ -0,0 +1,121 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_tiling.h
13+ * \brief confusion_matrix_tiling head file
14+ */
15+ 
16+#ifndef OPS_MATH_CONFUSION_MATRIX_TILING_H
17+#define OPS_MATH_CONFUSION_MATRIX_TILING_H
18+ 
19+#include "register/op_def_registry.h"
20+#include "op_host/tiling_base_class.h"
21+#include "../../op_kernel/arch35/confusion_matrix_tiling_key.h"
22+#include "../../op_kernel/arch35/confusion_matrix_tiling_data.h"
23+ 
24+namespace optiling {
25+ 
26+constexpr size_t WORK_SPACE_SIZE = static_cast<size_t>(16 * 1024 * 1024);
27+constexpr uint64_t DIM_0 = 0;
28+constexpr uint64_t DIM_1 = 1;
29+constexpr uint64_t DIM_2 = 2;
30+ 
31+constexpr int64_t INPUT_IDX_LABELS = 0;
32+constexpr int64_t INPUT_IDX_PREDICTIONS = 1;
33+constexpr int64_t INPUT_IDX_WEIGHTS = 2;
34+constexpr int64_t OUTPUT_IDX_Y = 0;
35+ 
36+constexpr int64_t ATTR_IDX_NUM_CLASSES = 0;
37+constexpr int64_t ATTR_IDX_DTYPE = 1;
38+ 
39+constexpr uint64_t SIZE_DTYPE_FLOAT = 4;
40+constexpr uint64_t SIZE_DTYPE_INT32 = 4;
41+constexpr uint64_t SIZE_DTYPE_FLOAT16 = 2;
42+constexpr uint64_t SIZE_DTYPE_INT8 = 1;
43+constexpr uint64_t SIZE_DTYPE_UINT8 = 1;
44+ 
45+constexpr uint64_t EMPTY_WEIGHT = 0;
46+constexpr uint64_t WEIGHT = 1;
47+ 
48+constexpr uint64_t OUTPUT_DTYPE_FLOAT = 0;
49+constexpr uint64_t OUTPUT_DTYPE_INT32 = 1;
50+constexpr uint64_t OUTPUT_DTYPE_FLOAT16 = 2;
51+constexpr uint64_t OUTPUT_DTYPE_INT8 = 3;
52+constexpr uint64_t OUTPUT_DTYPE_UINT8 = 4;
53+ 
54+constexpr uint64_t SCH_ID_SIMT_FULL_LOAD = 0; // full load to UB
55+constexpr uint64_t SCH_ID_SIMT_BATCH_LOAD = 1; // batch load to UB
56+constexpr uint64_t SCH_ID_SIMT_NOT_FULL_LOAD = 2; // not load to UB
57+constexpr uint64_t SCH_ID_SIMT_DETERMIN = 3;
58+ 
59+constexpr int64_t GM_ATOMIC_ADD_FACTOR = 100;
60+constexpr uint64_t SIMD_SIMT_DCACHE_SIZE = static_cast<uint64_t>(32 * 1024);
61+constexpr int64_t MAX_UB_LOOP_FOR_BATCH = 8;
62+ 
63+constexpr int64_t NUM_CLASSES_MIN = 1;
64+constexpr int64_t NUM_CLASSES_MAX = 4096;
65+ 
66+struct AscendCConfusionMatrixCompileInfo {
67+ int32_t totalCoreNum = 0;
68+ uint64_t totalUbSize = 0;
69+};
70+ 
71+class ConfusionMatrixTiling {
72+public:
73+ explicit ConfusionMatrixTiling(gert::TilingContext* context) : context_(context) {};
74+ ge::graphStatus Init();
75+ ge::graphStatus ConfusionMatrixGetPlatformData(const AscendCConfusionMatrixCompileInfo* compileInfo);
76+ ge::graphStatus CheckShape();
77+ ge::graphStatus CheckDtype();
78+ ge::graphStatus CheckInputParams();
79+ 
80+ inline bool IsMatchSimtBatchLoadMode();
81+ ge::graphStatus ComputeTilingStrategy();
82+ ge::graphStatus ComputeTilingSimtDetermine();
83+ ge::graphStatus ComputeTilingSimtNotDetermine();
84+ ge::graphStatus SetTilingData();
85+ void PrintTilingData();
86+ 
87+private:
88+ gert::TilingContext* context_ = nullptr;
89+ ConfusionMatrixTilingData* tilingData_ = nullptr;
90+ // tiling key template param
91+ uint64_t schId_ = 0;
92+ uint64_t isWeight_ = EMPTY_WEIGHT;
93+ uint64_t outputDtype_ = 0;
94+ // configed info
95+ int64_t coreNum_ = 0;
96+ int64_t ubSize_ = 0;
97+ int64_t isDetermine_ = 0;
98+ int64_t sizeOfOutputDtype_ = 0;
99+ 
100+ // tiling info
101+ int64_t inputSize_ = 0;
102+ int64_t numClasses_ = 0;
103+ int64_t outputSize_ = 0;
104+ int64_t needCoreNum_ = 0;
105+ 
106+ int64_t ubNumCanUse_ = 0;
107+ int64_t ubLoopNum_ = 0;
108+ int64_t needXCoreNum_ = 0;
109+ int64_t formerLength_ = 0;
110+ int64_t tailLength_ = 0;
111+ int64_t clearYCoreNum_ = 0;
112+ int64_t clearYFactor_ = 0;
113+ int64_t clearYTail_ = 0;
114+ int64_t binsFormerLength_ = 0;
115+ int64_t needBinsCoreNum_ = 0;
116+ int64_t binsTailLength_ = 0;
117+};
118+ 
119+} // namespace optiling
120+ 
121+#endif // OPS_MATH_CONFUSION_MATRIX_TILING_H
Amath/confusion_matrix/op_host/config/ascend950/confusion_matrix_binary.json+380-0
@@ -0,0 +1,380 @@
1+{
2+ "op_type": "ConfusionMatrix",
3+ "op_list": [
4+ {
5+ "bin_filename": "ConfusionMatrix_float16_float16_float16_float16",
6+ "inputs": [
7+ { "name": "labels", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
8+ { "name": "predictions", "index": 1, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
9+ { "name": "weights", "index": 2, "dtype": "float16", "format": "ND", "paramType": "optional", "shape": [-2] }
10+ ],
11+ "outputs": [
12+ { "name": "y", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] }
13+ ],
14+ "attrs": [
15+ { "name": "num_classes", "dtype": "int", "value": null },
16+ { "name": "dtype", "dtype": "string", "value": "float16" }
17+ ]
18+ },
19+ {
20+ "bin_filename": "ConfusionMatrix_float16_float16_float32_float32",
21+ "inputs": [
22+ { "name": "labels", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
23+ { "name": "predictions", "index": 1, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
24+ { "name": "weights", "index": 2, "dtype": "float32", "format": "ND", "paramType": "optional", "shape": [-2] }
25+ ],
26+ "outputs": [
27+ { "name": "y", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] }
28+ ],
29+ "attrs": [
30+ { "name": "num_classes", "dtype": "int", "value": null },
31+ { "name": "dtype", "dtype": "string", "value": "float32" }
32+ ]
33+ },
34+ {
35+ "bin_filename": "ConfusionMatrix_float16_float16_int32_int32",
36+ "inputs": [
37+ { "name": "labels", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
38+ { "name": "predictions", "index": 1, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
39+ { "name": "weights", "index": 2, "dtype": "int32", "format": "ND", "paramType": "optional", "shape": [-2] }
40+ ],
41+ "outputs": [
42+ { "name": "y", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] }
43+ ],
44+ "attrs": [
45+ { "name": "num_classes", "dtype": "int", "value": null },
46+ { "name": "dtype", "dtype": "string", "value": "int32" }
47+ ]
48+ },
49+ {
50+ "bin_filename": "ConfusionMatrix_float16_float16_int8_int8",
51+ "inputs": [
52+ { "name": "labels", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
53+ { "name": "predictions", "index": 1, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
54+ { "name": "weights", "index": 2, "dtype": "int8", "format": "ND", "paramType": "optional", "shape": [-2] }
55+ ],
56+ "outputs": [
57+ { "name": "y", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] }
58+ ],
59+ "attrs": [
60+ { "name": "num_classes", "dtype": "int", "value": null },
61+ { "name": "dtype", "dtype": "string", "value": "int8" }
62+ ]
63+ },
64+ {
65+ "bin_filename": "ConfusionMatrix_float16_float16_uint8_uint8",
66+ "inputs": [
67+ { "name": "labels", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
68+ { "name": "predictions", "index": 1, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] },
69+ { "name": "weights", "index": 2, "dtype": "uint8", "format": "ND", "paramType": "optional", "shape": [-2] }
70+ ],
71+ "outputs": [
72+ { "name": "y", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] }
73+ ],
74+ "attrs": [
75+ { "name": "num_classes", "dtype": "int", "value": null },
76+ { "name": "dtype", "dtype": "string", "value": "uint8" }
77+ ]
78+ },
79+ {
80+ "bin_filename": "ConfusionMatrix_float32_float32_float16_float16",
81+ "inputs": [
82+ { "name": "labels", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
83+ { "name": "predictions", "index": 1, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
84+ { "name": "weights", "index": 2, "dtype": "float16", "format": "ND", "paramType": "optional", "shape": [-2] }
85+ ],
86+ "outputs": [
87+ { "name": "y", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] }
88+ ],
89+ "attrs": [
90+ { "name": "num_classes", "dtype": "int", "value": null },
91+ { "name": "dtype", "dtype": "string", "value": "float16" }
92+ ]
93+ },
94+ {
95+ "bin_filename": "ConfusionMatrix_float32_float32_float32_float32",
96+ "inputs": [
97+ { "name": "labels", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
98+ { "name": "predictions", "index": 1, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
99+ { "name": "weights", "index": 2, "dtype": "float32", "format": "ND", "paramType": "optional", "shape": [-2] }
100+ ],
101+ "outputs": [
102+ { "name": "y", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] }
103+ ],
104+ "attrs": [
105+ { "name": "num_classes", "dtype": "int", "value": null },
106+ { "name": "dtype", "dtype": "string", "value": "float32" }
107+ ]
108+ },
109+ {
110+ "bin_filename": "ConfusionMatrix_float32_float32_int32_int32",
111+ "inputs": [
112+ { "name": "labels", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
113+ { "name": "predictions", "index": 1, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
114+ { "name": "weights", "index": 2, "dtype": "int32", "format": "ND", "paramType": "optional", "shape": [-2] }
115+ ],
116+ "outputs": [
117+ { "name": "y", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] }
118+ ],
119+ "attrs": [
120+ { "name": "num_classes", "dtype": "int", "value": null },
121+ { "name": "dtype", "dtype": "string", "value": "int32" }
122+ ]
123+ },
124+ {
125+ "bin_filename": "ConfusionMatrix_float32_float32_int8_int8",
126+ "inputs": [
127+ { "name": "labels", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
128+ { "name": "predictions", "index": 1, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
129+ { "name": "weights", "index": 2, "dtype": "int8", "format": "ND", "paramType": "optional", "shape": [-2] }
130+ ],
131+ "outputs": [
132+ { "name": "y", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] }
133+ ],
134+ "attrs": [
135+ { "name": "num_classes", "dtype": "int", "value": null },
136+ { "name": "dtype", "dtype": "string", "value": "int8" }
137+ ]
138+ },
139+ {
140+ "bin_filename": "ConfusionMatrix_float32_float32_uint8_uint8",
141+ "inputs": [
142+ { "name": "labels", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
143+ { "name": "predictions", "index": 1, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] },
144+ { "name": "weights", "index": 2, "dtype": "uint8", "format": "ND", "paramType": "optional", "shape": [-2] }
145+ ],
146+ "outputs": [
147+ { "name": "y", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] }
148+ ],
149+ "attrs": [
150+ { "name": "num_classes", "dtype": "int", "value": null },
151+ { "name": "dtype", "dtype": "string", "value": "uint8" }
152+ ]
153+ },
154+ {
155+ "bin_filename": "ConfusionMatrix_int32_int32_float16_float16",
156+ "inputs": [
157+ { "name": "labels", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
158+ { "name": "predictions", "index": 1, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
159+ { "name": "weights", "index": 2, "dtype": "float16", "format": "ND", "paramType": "optional", "shape": [-2] }
160+ ],
161+ "outputs": [
162+ { "name": "y", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] }
163+ ],
164+ "attrs": [
165+ { "name": "num_classes", "dtype": "int", "value": null },
166+ { "name": "dtype", "dtype": "string", "value": "float16" }
167+ ]
168+ },
169+ {
170+ "bin_filename": "ConfusionMatrix_int32_int32_float32_float32",
171+ "inputs": [
172+ { "name": "labels", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
173+ { "name": "predictions", "index": 1, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
174+ { "name": "weights", "index": 2, "dtype": "float32", "format": "ND", "paramType": "optional", "shape": [-2] }
175+ ],
176+ "outputs": [
177+ { "name": "y", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] }
178+ ],
179+ "attrs": [
180+ { "name": "num_classes", "dtype": "int", "value": null },
181+ { "name": "dtype", "dtype": "string", "value": "float32" }
182+ ]
183+ },
184+ {
185+ "bin_filename": "ConfusionMatrix_int32_int32_int32_int32",
186+ "inputs": [
187+ { "name": "labels", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
188+ { "name": "predictions", "index": 1, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
189+ { "name": "weights", "index": 2, "dtype": "int32", "format": "ND", "paramType": "optional", "shape": [-2] }
190+ ],
191+ "outputs": [
192+ { "name": "y", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] }
193+ ],
194+ "attrs": [
195+ { "name": "num_classes", "dtype": "int", "value": null },
196+ { "name": "dtype", "dtype": "string", "value": "int32" }
197+ ]
198+ },
199+ {
200+ "bin_filename": "ConfusionMatrix_int32_int32_int8_int8",
201+ "inputs": [
202+ { "name": "labels", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
203+ { "name": "predictions", "index": 1, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
204+ { "name": "weights", "index": 2, "dtype": "int8", "format": "ND", "paramType": "optional", "shape": [-2] }
205+ ],
206+ "outputs": [
207+ { "name": "y", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] }
208+ ],
209+ "attrs": [
210+ { "name": "num_classes", "dtype": "int", "value": null },
211+ { "name": "dtype", "dtype": "string", "value": "int8" }
212+ ]
213+ },
214+ {
215+ "bin_filename": "ConfusionMatrix_int32_int32_uint8_uint8",
216+ "inputs": [
217+ { "name": "labels", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
218+ { "name": "predictions", "index": 1, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] },
219+ { "name": "weights", "index": 2, "dtype": "uint8", "format": "ND", "paramType": "optional", "shape": [-2] }
220+ ],
221+ "outputs": [
222+ { "name": "y", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] }
223+ ],
224+ "attrs": [
225+ { "name": "num_classes", "dtype": "int", "value": null },
226+ { "name": "dtype", "dtype": "string", "value": "uint8" }
227+ ]
228+ },
229+ {
230+ "bin_filename": "ConfusionMatrix_int8_int8_float16_float16",
231+ "inputs": [
232+ { "name": "labels", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
233+ { "name": "predictions", "index": 1, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
234+ { "name": "weights", "index": 2, "dtype": "float16", "format": "ND", "paramType": "optional", "shape": [-2] }
235+ ],
236+ "outputs": [
237+ { "name": "y", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] }
238+ ],
239+ "attrs": [
240+ { "name": "num_classes", "dtype": "int", "value": null },
241+ { "name": "dtype", "dtype": "string", "value": "float16" }
242+ ]
243+ },
244+ {
245+ "bin_filename": "ConfusionMatrix_int8_int8_float32_float32",
246+ "inputs": [
247+ { "name": "labels", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
248+ { "name": "predictions", "index": 1, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
249+ { "name": "weights", "index": 2, "dtype": "float32", "format": "ND", "paramType": "optional", "shape": [-2] }
250+ ],
251+ "outputs": [
252+ { "name": "y", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] }
253+ ],
254+ "attrs": [
255+ { "name": "num_classes", "dtype": "int", "value": null },
256+ { "name": "dtype", "dtype": "string", "value": "float32" }
257+ ]
258+ },
259+ {
260+ "bin_filename": "ConfusionMatrix_int8_int8_int32_int32",
261+ "inputs": [
262+ { "name": "labels", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
263+ { "name": "predictions", "index": 1, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
264+ { "name": "weights", "index": 2, "dtype": "int32", "format": "ND", "paramType": "optional", "shape": [-2] }
265+ ],
266+ "outputs": [
267+ { "name": "y", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] }
268+ ],
269+ "attrs": [
270+ { "name": "num_classes", "dtype": "int", "value": null },
271+ { "name": "dtype", "dtype": "string", "value": "int32" }
272+ ]
273+ },
274+ {
275+ "bin_filename": "ConfusionMatrix_int8_int8_int8_int8",
276+ "inputs": [
277+ { "name": "labels", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
278+ { "name": "predictions", "index": 1, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
279+ { "name": "weights", "index": 2, "dtype": "int8", "format": "ND", "paramType": "optional", "shape": [-2] }
280+ ],
281+ "outputs": [
282+ { "name": "y", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] }
283+ ],
284+ "attrs": [
285+ { "name": "num_classes", "dtype": "int", "value": null },
286+ { "name": "dtype", "dtype": "string", "value": "int8" }
287+ ]
288+ },
289+ {
290+ "bin_filename": "ConfusionMatrix_int8_int8_uint8_uint8",
291+ "inputs": [
292+ { "name": "labels", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
293+ { "name": "predictions", "index": 1, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] },
294+ { "name": "weights", "index": 2, "dtype": "uint8", "format": "ND", "paramType": "optional", "shape": [-2] }
295+ ],
296+ "outputs": [
297+ { "name": "y", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] }
298+ ],
299+ "attrs": [
300+ { "name": "num_classes", "dtype": "int", "value": null },
301+ { "name": "dtype", "dtype": "string", "value": "uint8" }
302+ ]
303+ },
304+ {
305+ "bin_filename": "ConfusionMatrix_uint8_uint8_float16_float16",
306+ "inputs": [
307+ { "name": "labels", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
308+ { "name": "predictions", "index": 1, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
309+ { "name": "weights", "index": 2, "dtype": "float16", "format": "ND", "paramType": "optional", "shape": [-2] }
310+ ],
311+ "outputs": [
312+ { "name": "y", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2] }
313+ ],
314+ "attrs": [
315+ { "name": "num_classes", "dtype": "int", "value": null },
316+ { "name": "dtype", "dtype": "string", "value": "float16" }
317+ ]
318+ },
319+ {
320+ "bin_filename": "ConfusionMatrix_uint8_uint8_float32_float32",
321+ "inputs": [
322+ { "name": "labels", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
323+ { "name": "predictions", "index": 1, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
324+ { "name": "weights", "index": 2, "dtype": "float32", "format": "ND", "paramType": "optional", "shape": [-2] }
325+ ],
326+ "outputs": [
327+ { "name": "y", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2] }
328+ ],
329+ "attrs": [
330+ { "name": "num_classes", "dtype": "int", "value": null },
331+ { "name": "dtype", "dtype": "string", "value": "float32" }
332+ ]
333+ },
334+ {
335+ "bin_filename": "ConfusionMatrix_uint8_uint8_int32_int32",
336+ "inputs": [
337+ { "name": "labels", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
338+ { "name": "predictions", "index": 1, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
339+ { "name": "weights", "index": 2, "dtype": "int32", "format": "ND", "paramType": "optional", "shape": [-2] }
340+ ],
341+ "outputs": [
342+ { "name": "y", "index": 0, "dtype": "int32", "format": "ND", "paramType": "required", "shape": [-2] }
343+ ],
344+ "attrs": [
345+ { "name": "num_classes", "dtype": "int", "value": null },
346+ { "name": "dtype", "dtype": "string", "value": "int32" }
347+ ]
348+ },
349+ {
350+ "bin_filename": "ConfusionMatrix_uint8_uint8_int8_int8",
351+ "inputs": [
352+ { "name": "labels", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
353+ { "name": "predictions", "index": 1, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
354+ { "name": "weights", "index": 2, "dtype": "int8", "format": "ND", "paramType": "optional", "shape": [-2] }
355+ ],
356+ "outputs": [
357+ { "name": "y", "index": 0, "dtype": "int8", "format": "ND", "paramType": "required", "shape": [-2] }
358+ ],
359+ "attrs": [
360+ { "name": "num_classes", "dtype": "int", "value": null },
361+ { "name": "dtype", "dtype": "string", "value": "int8" }
362+ ]
363+ },
364+ {
365+ "bin_filename": "ConfusionMatrix_uint8_uint8_uint8_uint8",
366+ "inputs": [
367+ { "name": "labels", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
368+ { "name": "predictions", "index": 1, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] },
369+ { "name": "weights", "index": 2, "dtype": "uint8", "format": "ND", "paramType": "optional", "shape": [-2] }
370+ ],
371+ "outputs": [
372+ { "name": "y", "index": 0, "dtype": "uint8", "format": "ND", "paramType": "required", "shape": [-2] }
373+ ],
374+ "attrs": [
375+ { "name": "num_classes", "dtype": "int", "value": null },
376+ { "name": "dtype", "dtype": "string", "value": "uint8" }
377+ ]
378+ }
379+ ]
380+}
Amath/confusion_matrix/op_host/config/ascend950/confusion_matrix_simplified_key.ini+2-0
@@ -0,0 +1,2 @@
1+[ConfusionMatrix]
2+default=0
Amath/confusion_matrix/op_host/confusion_matrix_def.cpp+74-0
@@ -0,0 +1,74 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_def.cpp
13+ * \brief confusion_matrix def
14+ */
15+ 
16+#include "register/op_def_registry.h"
17+ 
18+namespace ops {
19+ 
20+static const std::vector<ge::DataType> labelsPredictionsDtypes = {
21+ ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT,
22+ ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
23+ ge::DT_INT32, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_UINT8,
24+ ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8, ge::DT_UINT8};
25+ 
26+static const std::vector<ge::DataType> weightsOutputDtypes = {
27+ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8, ge::DT_FLOAT16, ge::DT_FLOAT,
28+ ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT8,
29+ ge::DT_UINT8, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8, ge::DT_FLOAT16,
30+ ge::DT_FLOAT, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8};
31+ 
32+static const std::vector<ge::Format> kFormats = {
33+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
34+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
35+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
36+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
37+ 
38+class ConfusionMatrix : public OpDef {
39+public:
40+ explicit ConfusionMatrix(const char* name) : OpDef(name)
41+ {
42+ this->Input("labels")
43+ .ParamType(REQUIRED)
44+ .DataType(labelsPredictionsDtypes)
45+ .Format(kFormats)
46+ .UnknownShapeFormat(kFormats);
47+ this->Input("predictions")
48+ .ParamType(REQUIRED)
49+ .DataType(labelsPredictionsDtypes)
50+ .Format(kFormats)
51+ .UnknownShapeFormat(kFormats);
52+ this->Input("weights")
53+ .ParamType(OPTIONAL)
54+ .DataType(weightsOutputDtypes)
55+ .Format(kFormats)
56+ .UnknownShapeFormat(kFormats);
57+ this->Output("y")
58+ .ParamType(REQUIRED)
59+ .DataType(weightsOutputDtypes)
60+ .Format(kFormats)
61+ .UnknownShapeFormat(kFormats);
62+ this->Attr("num_classes").AttrType(REQUIRED).Int();
63+ this->Attr("dtype").AttrType(REQUIRED).String("float32");
64+ OpAICoreConfig aicoreConfig;
65+ aicoreConfig.DynamicCompileStaticFlag(true)
66+ .DynamicFormatFlag(false)
67+ .DynamicRankSupportFlag(true)
68+ .DynamicShapeSupportFlag(true)
69+ .ExtendCfgInfo("opFile.value", "confusion_matrix");
70+ this->AICore().AddConfig("ascend950", aicoreConfig);
71+ }
72+};
73+OP_ADD(ConfusionMatrix);
74+} // namespace ops
Amath/confusion_matrix/op_host/confusion_matrix_infershape.cpp+77-0
@@ -0,0 +1,77 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_infershape.cpp
13+ * \brief Infer shape implementation for ConfusionMatrix operator
14+ */
15+ 
16+#include "log/log.h"
17+#include "register/op_impl_registry.h"
18+#include "util/shape_util.h"
19+ 
20+#include <string>
21+ 
22+using namespace ge;
23+namespace ops {
24+static constexpr size_t kAttrIndexNumClasses = 0U;
25+static constexpr size_t kInputIndexLabels = 0U;
26+static constexpr size_t kInputIndexPredictions = 1U;
27+static constexpr size_t kOutputIndexY = 0U;
28+static constexpr int64_t kNumClassesMin = 1;
29+static constexpr int64_t kNumClassesMax = 4096;
30+static constexpr size_t kOutputDimNum = 2U;
31+static constexpr size_t kOutputDim0 = 0U;
32+static constexpr size_t kOutputDim1 = 1U;
33+ 
34+ge::graphStatus InferShapeForConfusionMatrix(gert::InferShapeContext* context)
35+{
36+ OP_LOGI(context->GetNodeName(), "Begin to do InferShapeForConfusionMatrix");
37+ 
38+ auto attrs = context->GetAttrs();
39+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
40+ 
41+ const int64_t* numClassesPtr = attrs->GetAttrPointer<int64_t>(kAttrIndexNumClasses);
42+ OP_CHECK_NULL_WITH_CONTEXT(context, numClassesPtr);
43+ 
44+ int64_t numClasses = *numClassesPtr;
45+ OP_CHECK_IF(
46+ (numClasses < kNumClassesMin || numClasses > kNumClassesMax),
47+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "num_classes", std::to_string(numClasses).c_str(),
48+ "num_classes must be in range [1, 4096]."),
49+ return ge::GRAPH_FAILED);
50+ 
51+ gert::Shape* yShape = context->GetOutputShape(kOutputIndexY);
52+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
53+ 
54+ const gert::Shape* labelsShape = context->GetInputShape(kInputIndexLabels);
55+ OP_CHECK_NULL_WITH_CONTEXT(context, labelsShape);
56+ const gert::Shape* predictionsShape = context->GetInputShape(kInputIndexPredictions);
57+ OP_CHECK_NULL_WITH_CONTEXT(context, predictionsShape);
58+ 
59+ // labels/predictions 为 unknown rank(-2) 时无法校验 1D 约束，输出透传 unknown rank；
60+ // unknown shape(-1，rank 已知)时输出仍由 num_classes 决定，走下方固定 [num_classes, num_classes] 分支
61+ if (Ops::Base::IsUnknownRank(*labelsShape) || Ops::Base::IsUnknownRank(*predictionsShape)) {
62+ Ops::Base::SetUnknownRank(*yShape);
63+ OP_LOGI(context->GetNodeName(), "labels/predictions is unknown rank, set output to unknown rank");
64+ return ge::GRAPH_SUCCESS;
65+ }
66+ 
67+ yShape->SetDimNum(kOutputDimNum);
68+ yShape->SetDim(kOutputDim0, numClasses);
69+ yShape->SetDim(kOutputDim1, numClasses);
70+ 
71+ OP_LOGI(context->GetNodeName(), "End to do InferShapeForConfusionMatrix");
72+ return ge::GRAPH_SUCCESS;
73+}
74+ 
75+IMPL_OP_INFERSHAPE(ConfusionMatrix).InferShape(InferShapeForConfusionMatrix);
76+ 
77+} // namespace ops
Amath/confusion_matrix/op_kernel/arch35/confusion_matrix_determine.h+126-0
@@ -0,0 +1,126 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software; you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_determine.h
13+ * \brief confusion_matrix determine realization
14+ */
15+ 
16+#ifndef CONFUSION_MATRIX_DETERMINE_H
17+#define CONFUSION_MATRIX_DETERMINE_H
18+ 
19+#include "confusion_matrix_tiling_data.h"
20+ 
21+namespace ConfusionMatrixSimt {
22+using namespace AscendC;
23+ 
24+template <typename LABEL_TYPE, typename Y_TYPE>
25+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void DetermineSimtCompute(
26+ __gm__ LABEL_TYPE* labelsGmAddr, __gm__ LABEL_TYPE* predictionsGmAddr, __gm__ Y_TYPE* yGmAddr,
27+ const int64_t binsHeadIndex, const int64_t binsDataLength, const int64_t inputSize, const int64_t numClasses)
28+{
29+ for (int64_t index = static_cast<int64_t>(Simt::GetThreadIdx()); index < binsDataLength;
30+ index += static_cast<int64_t>(Simt::GetThreadNum<0>())) {
31+ yGmAddr[binsHeadIndex + index] = static_cast<Y_TYPE>(0);
32+ for (int64_t inputIndex = 0; inputIndex < inputSize; inputIndex++) {
33+ int64_t label = static_cast<int64_t>(labelsGmAddr[inputIndex]);
34+ int64_t prediction = static_cast<int64_t>(predictionsGmAddr[inputIndex]);
35+ int64_t offset = label * numClasses + prediction;
36+ if (offset == binsHeadIndex + index) {
37+ yGmAddr[binsHeadIndex + index] += static_cast<Y_TYPE>(WEIGHT_ONE);
38+ }
39+ }
40+ }
41+}
42+ 
43+template <typename LABEL_TYPE, typename Y_TYPE>
44+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void DetermineSimtComputeWithWeight(
45+ __gm__ LABEL_TYPE* labelsGmAddr, __gm__ LABEL_TYPE* predictionsGmAddr, __gm__ Y_TYPE* weightsGmAddr,
46+ __gm__ Y_TYPE* yGmAddr, const int64_t binsHeadIndex, const int64_t binsDataLength, const int64_t inputSize,
47+ const int64_t numClasses)
48+{
49+ for (int64_t index = static_cast<int64_t>(Simt::GetThreadIdx()); index < binsDataLength;
50+ index += static_cast<int64_t>(Simt::GetThreadNum<0>())) {
51+ yGmAddr[binsHeadIndex + index] = static_cast<Y_TYPE>(0);
52+ for (int64_t inputIndex = 0; inputIndex < inputSize; inputIndex++) {
53+ int64_t label = static_cast<int64_t>(labelsGmAddr[inputIndex]);
54+ int64_t prediction = static_cast<int64_t>(predictionsGmAddr[inputIndex]);
55+ int64_t offset = label * numClasses + prediction;
56+ if (offset == binsHeadIndex + index) {
57+ yGmAddr[binsHeadIndex + index] += weightsGmAddr[inputIndex];
58+ }
59+ }
60+ }
61+}
62+ 
63+template <typename LABEL_TYPE, typename Y_TYPE>
64+class ConfusionMatrixDetermine : public ConfusionMatrixSimtBase<LABEL_TYPE, Y_TYPE> {
65+public:
66+ __aicore__ inline ConfusionMatrixDetermine(){};
67+ __aicore__ inline void Init(GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
68+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe,
69+ bool isWeightEmpty);
70+ __aicore__ inline void Process();
71+ 
72+private:
73+ __aicore__ inline void Compute();
74+ 
75+private:
76+ int64_t binsFormerLength_ = 0;
77+ int64_t needBinsCoreNum_ = 0;
78+ int64_t binsTailLength_ = 0;
79+};
80+ 
81+template <typename LABEL_TYPE, typename Y_TYPE>
82+__aicore__ inline void ConfusionMatrixDetermine<LABEL_TYPE, Y_TYPE>::Init(
83+ GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
84+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe, bool isWeightEmpty)
85+{
86+ this->BaseInit(labels, predictions, weights, y, tilingData, tPipe, isWeightEmpty);
87+ binsFormerLength_ = tilingData->binsFormerLength;
88+ needBinsCoreNum_ = tilingData->needBinsCoreNum;
89+ binsTailLength_ = tilingData->binsTailLength;
90+}
91+ 
92+template <typename LABEL_TYPE, typename Y_TYPE>
93+__aicore__ inline void ConfusionMatrixDetermine<LABEL_TYPE, Y_TYPE>::Process()
94+{
95+ if (this->blockIdx_ >= GetBlockNum()) {
96+ return;
97+ }
98+ 
99+ if (this->blockIdx_ < needBinsCoreNum_) {
100+ Compute();
101+ }
102+}
103+ 
104+template <typename LABEL_TYPE, typename Y_TYPE>
105+__aicore__ inline void ConfusionMatrixDetermine<LABEL_TYPE, Y_TYPE>::Compute()
106+{
107+ __gm__ LABEL_TYPE* labelsGmAddr = (__gm__ LABEL_TYPE*)this->labelsGm_.GetPhyAddr();
108+ __gm__ LABEL_TYPE* predictionsGmAddr = (__gm__ LABEL_TYPE*)this->predictionsGm_.GetPhyAddr();
109+ __gm__ Y_TYPE* yGmAddr = (__gm__ Y_TYPE*)this->yGm_.GetPhyAddr();
110+ int64_t binsHeadIndex = this->blockIdx_ * binsFormerLength_;
111+ int64_t binsDataLength = (this->blockIdx_ == needBinsCoreNum_ - 1) ? binsTailLength_ : binsFormerLength_;
112+ 
113+ if (this->isWeightEmpty_) {
114+ Simt::VF_CALL<DetermineSimtCompute<LABEL_TYPE, Y_TYPE>>(Simt::Dim3{THREAD_NUM}, labelsGmAddr, predictionsGmAddr,
115+ yGmAddr, binsHeadIndex, binsDataLength,
116+ this->inputSize_, this->numClasses_);
117+ } else {
118+ __gm__ Y_TYPE* weightsGmAddr = (__gm__ Y_TYPE*)this->weightsGm_.GetPhyAddr();
119+ Simt::VF_CALL<DetermineSimtComputeWithWeight<LABEL_TYPE, Y_TYPE>>(
120+ Simt::Dim3{THREAD_NUM}, labelsGmAddr, predictionsGmAddr, weightsGmAddr, yGmAddr, binsHeadIndex,
121+ binsDataLength, this->inputSize_, this->numClasses_);
122+ }
123+}
124+ 
125+} // namespace ConfusionMatrixSimt
126+#endif // CONFUSION_MATRIX_DETERMINE_H
Amath/confusion_matrix/op_kernel/arch35/confusion_matrix_simt_base.h+103-0
@@ -0,0 +1,103 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software; you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_simt_base.h
13+ * \brief The base class of confusion_matrix. The main function of this file is:
14+ */
15+ 
16+#ifndef CONFUSION_MATRIX_SIMT_BASE_H
17+#define CONFUSION_MATRIX_SIMT_BASE_H
18+ 
19+#include "confusion_matrix_tiling_data.h"
20+ 
21+namespace ConfusionMatrixSimt {
22+using namespace AscendC;
23+ 
24+constexpr uint32_t OUT_QUE_DEPTH = 1;
25+ 
26+constexpr uint32_t NUM_DOUBLE_BUFFER = 1;
27+ 
28+constexpr uint16_t DATA_COPY_PAD_BLOCK_COUNT = 1;
29+ 
30+constexpr int32_t WEIGHT_ONE = 1;
31+ 
32+constexpr uint32_t THREAD_NUM = 1024;
33+ 
34+template <typename LABEL_TYPE, typename Y_TYPE>
35+class ConfusionMatrixSimtBase {
36+public:
37+ __aicore__ inline ConfusionMatrixSimtBase(){};
38+ __aicore__ inline void BaseInit(GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
39+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe,
40+ bool isWeightEmpty);
41+ 
42+protected:
43+ __aicore__ inline void ParseTilingData(const ConfusionMatrixTilingData* __restrict tilingData);
44+ 
45+protected:
46+ GlobalTensor<LABEL_TYPE> labelsGm_;
47+ GlobalTensor<LABEL_TYPE> predictionsGm_;
48+ GlobalTensor<Y_TYPE> weightsGm_;
49+ GlobalTensor<Y_TYPE> yGm_;
50+ 
51+ int32_t blockIdx_ = 0;
52+ 
53+ bool isWeightEmpty_ = false;
54+ int64_t numClasses_ = 0;
55+ int64_t outputSize_ = 0;
56+ int64_t inputSize_ = 0;
57+ int64_t needCoreNum_ = 0;
58+ int64_t formerLength_ = 0;
59+ int64_t tailLength_ = 0;
60+ int64_t resetBinsCoreNum_ = 0;
61+ int64_t resetBinsLength_ = 0;
62+ int64_t resetBinsTailLength_ = 0;
63+ int64_t fullLoadBufSize_ = 0;
64+ int64_t batchUbBufSize_ = 0;
65+};
66+ 
67+template <typename LABEL_TYPE, typename Y_TYPE>
68+__aicore__ inline void ConfusionMatrixSimtBase<LABEL_TYPE, Y_TYPE>::BaseInit(
69+ GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
70+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe, bool isWeightEmpty)
71+{
72+ ParseTilingData(tilingData);
73+ isWeightEmpty_ = isWeightEmpty;
74+ labelsGm_.SetGlobalBuffer(reinterpret_cast<__gm__ LABEL_TYPE*>(labels));
75+ predictionsGm_.SetGlobalBuffer(reinterpret_cast<__gm__ LABEL_TYPE*>(predictions));
76+ if (!isWeightEmpty_) {
77+ weightsGm_.SetGlobalBuffer(reinterpret_cast<__gm__ Y_TYPE*>(weights));
78+ }
79+ yGm_.SetGlobalBuffer(reinterpret_cast<__gm__ Y_TYPE*>(y));
80+ 
81+ blockIdx_ = static_cast<int32_t>(GetBlockIdx());
82+}
83+ 
84+template <typename LABEL_TYPE, typename Y_TYPE>
85+__aicore__ inline void ConfusionMatrixSimtBase<LABEL_TYPE, Y_TYPE>::ParseTilingData(
86+ const ConfusionMatrixTilingData* __restrict tilingData)
87+{
88+ numClasses_ = tilingData->numClasses;
89+ outputSize_ = numClasses_ * numClasses_;
90+ inputSize_ = tilingData->inputSize;
91+ needCoreNum_ = tilingData->needXCoreNum;
92+ formerLength_ = tilingData->formerLength;
93+ tailLength_ = tilingData->tailLength;
94+ resetBinsCoreNum_ = tilingData->clearYCoreNum;
95+ resetBinsLength_ = tilingData->clearYFactor;
96+ resetBinsTailLength_ = tilingData->clearYTail;
97+ fullLoadBufSize_ = tilingData->fullLoadBufSize;
98+ batchUbBufSize_ = tilingData->batchUbBufSize;
99+}
100+ 
101+} // namespace ConfusionMatrixSimt
102+ 
103+#endif // CONFUSION_MATRIX_SIMT_BASE_H
Amath/confusion_matrix/op_kernel/arch35/confusion_matrix_simt_full_load.h+161-0
@@ -0,0 +1,161 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software; you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_simt_full_load.h
13+ * \brief The implementation of confusion_matrix by full load in ub.
14+ */
15+ 
16+#ifndef CONFUSION_MATRIX_SIMT_FULL_LOAD_H
17+#define CONFUSION_MATRIX_SIMT_FULL_LOAD_H
18+ 
19+#include "kernel_operator.h"
20+#include "confusion_matrix_simt_base.h"
21+#include "confusion_matrix_tiling_data.h"
22+#include "simt_api/asc_simt.h"
23+#include "simt_api/device_atomic_functions.h"
24+#include "simt_api/asc_fp16.h"
25+#include "simt_api/asc_bf16.h"
26+ 
27+namespace ConfusionMatrixSimt {
28+using namespace AscendC;
29+ 
30+template <typename LABEL_TYPE, typename Y_TYPE>
31+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void UbSimtCompute(
32+ __gm__ LABEL_TYPE* labelsGmAddr, __gm__ LABEL_TYPE* predictionsGmAddr, __ubuf__ Y_TYPE* yLocalAddr,
33+ const int64_t inputOffset, const int64_t inputDataLength, const int64_t numClasses)
34+{
35+ for (int64_t index = static_cast<int64_t>(threadIdx.x); index < inputDataLength;
36+ index += static_cast<int64_t>(blockDim.x)) {
37+ int64_t label = static_cast<int64_t>(labelsGmAddr[inputOffset + index]);
38+ int64_t prediction = static_cast<int64_t>(predictionsGmAddr[inputOffset + index]);
39+ int64_t offset = label * numClasses + prediction;
40+ asc_atomic_add(yLocalAddr + offset, static_cast<Y_TYPE>(WEIGHT_ONE));
41+ }
42+}
43+ 
44+template <typename LABEL_TYPE, typename Y_TYPE>
45+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void UbSimtComputeWithWeight(
46+ __gm__ LABEL_TYPE* labelsGmAddr, __gm__ LABEL_TYPE* predictionsGmAddr, __gm__ Y_TYPE* weightsGmAddr,
47+ __ubuf__ Y_TYPE* yLocalAddr, const int64_t inputOffset, const int64_t inputDataLength, const int64_t numClasses)
48+{
49+ for (int64_t index = static_cast<int64_t>(threadIdx.x); index < inputDataLength;
50+ index += static_cast<int64_t>(blockDim.x)) {
51+ int64_t label = static_cast<int64_t>(labelsGmAddr[inputOffset + index]);
52+ int64_t prediction = static_cast<int64_t>(predictionsGmAddr[inputOffset + index]);
53+ int64_t offset = label * numClasses + prediction;
54+ asc_atomic_add(yLocalAddr + offset, weightsGmAddr[inputOffset + index]);
55+ }
56+}
57+ 
58+template <typename LABEL_TYPE, typename Y_TYPE>
59+class ConfusionMatrixSimtFullLoad : public ConfusionMatrixSimtBase<LABEL_TYPE, Y_TYPE> {
60+public:
61+ __aicore__ inline ConfusionMatrixSimtFullLoad(){};
62+ __aicore__ inline void Init(GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
63+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe,
64+ bool isWeightEmpty);
65+ __aicore__ inline void Process();
66+ 
67+private:
68+ __aicore__ inline void ResetBins();
69+ __aicore__ inline void Compute();
70+ __aicore__ inline void CopyOut(int64_t offsetInGm, uint32_t stride);
71+ 
72+private:
73+ LocalTensor<Y_TYPE> yLocal_;
74+ TPipe* tPipe_;
75+ TQue<TPosition::VECOUT, OUT_QUE_DEPTH> yQue_;
76+};
77+ 
78+template <typename LABEL_TYPE, typename Y_TYPE>
79+__aicore__ inline void ConfusionMatrixSimtFullLoad<LABEL_TYPE, Y_TYPE>::Init(
80+ GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
81+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe, bool isWeightEmpty)
82+{
83+ this->BaseInit(labels, predictions, weights, y, tilingData, tPipe, isWeightEmpty);
84+ tPipe_ = tPipe;
85+ tPipe_->InitBuffer(yQue_, NUM_DOUBLE_BUFFER, this->fullLoadBufSize_);
86+}
87+ 
88+template <typename LABEL_TYPE, typename Y_TYPE>
89+__aicore__ inline void ConfusionMatrixSimtFullLoad<LABEL_TYPE, Y_TYPE>::Process()
90+{
91+ if (this->blockIdx_ >= GetBlockNum()) {
92+ return;
93+ }
94+ 
95+ ResetBins();
96+ SyncAll();
97+ Compute();
98+}
99+ 
100+template <typename LABEL_TYPE, typename Y_TYPE>
101+__aicore__ inline void ConfusionMatrixSimtFullLoad<LABEL_TYPE, Y_TYPE>::ResetBins()
102+{
103+ if (this->blockIdx_ >= this->resetBinsCoreNum_) {
104+ return;
105+ }
106+ 
107+ yLocal_ = yQue_.template AllocTensor<Y_TYPE>();
108+ Duplicate<Y_TYPE>(yLocal_, 0, this->outputSize_);
109+ yQue_.EnQue(yLocal_);
110+ yLocal_ = yQue_.template DeQue<Y_TYPE>();
111+ 
112+ int64_t resetBinsDataLength = this->blockIdx_ == this->resetBinsCoreNum_ - 1 ? this->resetBinsTailLength_ :
113+ this->resetBinsLength_;
114+ uint32_t resetBinsDataByteSize = static_cast<uint32_t>(resetBinsDataLength * sizeof(Y_TYPE));
115+ DataCopyExtParams dataCopyExtParams{DATA_COPY_PAD_BLOCK_COUNT, resetBinsDataByteSize, 0, 0, 0};
116+ DataCopyPad(this->yGm_[this->blockIdx_ * this->resetBinsLength_], yLocal_, dataCopyExtParams);
117+ yQue_.template FreeTensor<Y_TYPE>(yLocal_);
118+}
119+ 
120+template <typename LABEL_TYPE, typename Y_TYPE>
121+__aicore__ inline void ConfusionMatrixSimtFullLoad<LABEL_TYPE, Y_TYPE>::Compute()
122+{
123+ if (this->blockIdx_ >= this->needCoreNum_) {
124+ return;
125+ }
126+ 
127+ yLocal_ = yQue_.template AllocTensor<Y_TYPE>();
128+ Duplicate<Y_TYPE>(yLocal_, 0, this->outputSize_);
129+ int64_t inputOffset = this->blockIdx_ * this->formerLength_;
130+ int64_t inputDataLength = this->blockIdx_ == this->needCoreNum_ - 1 ? this->tailLength_ : this->formerLength_;
131+ __gm__ LABEL_TYPE* labelsGmAddr = (__gm__ LABEL_TYPE*)this->labelsGm_.GetPhyAddr();
132+ __gm__ LABEL_TYPE* predictionsGmAddr = (__gm__ LABEL_TYPE*)this->predictionsGm_.GetPhyAddr();
133+ __ubuf__ Y_TYPE* yLocalAddr = (__ubuf__ Y_TYPE*)yLocal_.GetPhyAddr();
134+ 
135+ if (this->isWeightEmpty_) {
136+ asc_vf_call<UbSimtCompute<LABEL_TYPE, Y_TYPE>>(dim3{THREAD_NUM}, labelsGmAddr, predictionsGmAddr, yLocalAddr,
137+ inputOffset, inputDataLength, this->numClasses_);
138+ } else {
139+ __gm__ Y_TYPE* weightsGmAddr = (__gm__ Y_TYPE*)this->weightsGm_.GetPhyAddr();
140+ asc_vf_call<UbSimtComputeWithWeight<LABEL_TYPE, Y_TYPE>>(dim3{THREAD_NUM}, labelsGmAddr, predictionsGmAddr,
141+ weightsGmAddr, yLocalAddr, inputOffset,
142+ inputDataLength, this->numClasses_);
143+ }
144+ 
145+ CopyOut(0, static_cast<uint32_t>(this->outputSize_ * sizeof(Y_TYPE)));
146+ yQue_.template FreeTensor<Y_TYPE>(yLocal_);
147+}
148+ 
149+template <typename LABEL_TYPE, typename Y_TYPE>
150+__aicore__ inline void ConfusionMatrixSimtFullLoad<LABEL_TYPE, Y_TYPE>::CopyOut(int64_t offsetInGm, uint32_t stride)
151+{
152+ yQue_.EnQue(yLocal_);
153+ yLocal_ = yQue_.template DeQue<Y_TYPE>();
154+ SetAtomicAdd<Y_TYPE>();
155+ DataCopyExtParams dataCopyExtParams{DATA_COPY_PAD_BLOCK_COUNT, stride, 0, 0, 0};
156+ DataCopyPad(this->yGm_[offsetInGm], yLocal_, dataCopyExtParams);
157+ SetAtomicNone();
158+}
159+ 
160+} // namespace ConfusionMatrixSimt
161+#endif // CONFUSION_MATRIX_SIMT_FULL_LOAD_H
Amath/confusion_matrix/op_kernel/arch35/confusion_matrix_simt_not_full_load_gm.h+129-0
@@ -0,0 +1,129 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software; you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_simt_not_full_load_gm.h
13+ * \brief The implementation of confusion_matrix by not full load in gm.
14+ */
15+ 
16+#ifndef CONFUSION_MATRIX_SIMT_NOT_FULL_LOAD_GM_H
17+#define CONFUSION_MATRIX_SIMT_NOT_FULL_LOAD_GM_H
18+ 
19+#include "confusion_matrix_tiling_data.h"
20+#include "simt_api/asc_simt.h"
21+#include "simt_api/device_atomic_functions.h"
22+#include "simt_api/asc_fp16.h"
23+#include "simt_api/asc_bf16.h"
24+ 
25+namespace ConfusionMatrixSimt {
26+using namespace AscendC;
27+ 
28+template <typename LABEL_TYPE, typename Y_TYPE>
29+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void GmSimtCompute(
30+ __gm__ LABEL_TYPE* labelsGmAddr, __gm__ LABEL_TYPE* predictionsGmAddr, __gm__ Y_TYPE* yGmAddr,
31+ const int64_t inputOffset, const int64_t inputDataLength, const int64_t numClasses)
32+{
33+ for (int64_t index = static_cast<int64_t>(threadIdx.x); index < inputDataLength;
34+ index += static_cast<int64_t>(blockDim.x)) {
35+ int64_t label = static_cast<int64_t>(labelsGmAddr[inputOffset + index]);
36+ int64_t prediction = static_cast<int64_t>(predictionsGmAddr[inputOffset + index]);
37+ int64_t offset = label * numClasses + prediction;
38+ asc_atomic_add(yGmAddr + offset, static_cast<Y_TYPE>(WEIGHT_ONE));
39+ }
40+}
41+ 
42+template <typename LABEL_TYPE, typename Y_TYPE>
43+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void GmSimtComputeWithWeight(
44+ __gm__ LABEL_TYPE* labelsGmAddr, __gm__ LABEL_TYPE* predictionsGmAddr, __gm__ Y_TYPE* weightsGmAddr,
45+ __gm__ Y_TYPE* yGmAddr, const int64_t inputOffset, const int64_t inputDataLength, const int64_t numClasses)
46+{
47+ for (int64_t index = static_cast<int64_t>(threadIdx.x); index < inputDataLength;
48+ index += static_cast<int64_t>(blockDim.x)) {
49+ int64_t label = static_cast<int64_t>(labelsGmAddr[inputOffset + index]);
50+ int64_t prediction = static_cast<int64_t>(predictionsGmAddr[inputOffset + index]);
51+ int64_t offset = label * numClasses + prediction;
52+ asc_atomic_add(yGmAddr + offset, weightsGmAddr[inputOffset + index]);
53+ }
54+}
55+ 
56+template <typename LABEL_TYPE, typename Y_TYPE>
57+class ConfusionMatrixSimtNotFullLoadGm : public ConfusionMatrixSimtBase<LABEL_TYPE, Y_TYPE> {
58+public:
59+ __aicore__ inline ConfusionMatrixSimtNotFullLoadGm(){};
60+ __aicore__ inline void Init(GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
61+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe,
62+ bool isWeightEmpty);
63+ __aicore__ inline void Process();
64+ 
65+private:
66+ __aicore__ inline void ResetBins();
67+ __aicore__ inline void Compute();
68+};
69+ 
70+template <typename LABEL_TYPE, typename Y_TYPE>
71+__aicore__ inline void ConfusionMatrixSimtNotFullLoadGm<LABEL_TYPE, Y_TYPE>::Init(
72+ GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
73+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe, bool isWeightEmpty)
74+{
75+ this->BaseInit(labels, predictions, weights, y, tilingData, tPipe, isWeightEmpty);
76+}
77+ 
78+template <typename LABEL_TYPE, typename Y_TYPE>
79+__aicore__ inline void ConfusionMatrixSimtNotFullLoadGm<LABEL_TYPE, Y_TYPE>::Process()
80+{
81+ if (this->blockIdx_ >= GetBlockNum()) {
82+ return;
83+ }
84+ 
85+ ResetBins();
86+ SyncAll();
87+ Compute();
88+}
89+ 
90+template <typename LABEL_TYPE, typename Y_TYPE>
91+__aicore__ inline void ConfusionMatrixSimtNotFullLoadGm<LABEL_TYPE, Y_TYPE>::ResetBins()
92+{
93+ if (this->blockIdx_ >= this->resetBinsCoreNum_) {
94+ return;
95+ }
96+ 
97+ int64_t resetBinsHeadIndex = this->blockIdx_ * this->resetBinsLength_;
98+ int64_t resetBinsDataLength = (this->blockIdx_ == this->resetBinsCoreNum_ - 1) ? this->resetBinsTailLength_ :
99+ this->resetBinsLength_;
100+ 
101+ InitOutput<Y_TYPE>(this->yGm_[resetBinsHeadIndex], resetBinsDataLength, static_cast<Y_TYPE>(0));
102+}
103+ 
104+template <typename LABEL_TYPE, typename Y_TYPE>
105+__aicore__ inline void ConfusionMatrixSimtNotFullLoadGm<LABEL_TYPE, Y_TYPE>::Compute()
106+{
107+ if (this->blockIdx_ >= this->needCoreNum_) {
108+ return;
109+ }
110+ 
111+ int64_t inputOffset = this->blockIdx_ * this->formerLength_;
112+ int64_t inputDataLength = (this->blockIdx_ == this->needCoreNum_ - 1) ? this->tailLength_ : this->formerLength_;
113+ __gm__ LABEL_TYPE* labelsGmAddr = (__gm__ LABEL_TYPE*)this->labelsGm_.GetPhyAddr();
114+ __gm__ LABEL_TYPE* predictionsGmAddr = (__gm__ LABEL_TYPE*)this->predictionsGm_.GetPhyAddr();
115+ __gm__ Y_TYPE* yGmAddr = (__gm__ Y_TYPE*)this->yGm_.GetPhyAddr();
116+ 
117+ if (this->isWeightEmpty_) {
118+ asc_vf_call<GmSimtCompute<LABEL_TYPE, Y_TYPE>>(dim3{THREAD_NUM}, labelsGmAddr, predictionsGmAddr, yGmAddr,
119+ inputOffset, inputDataLength, this->numClasses_);
120+ } else {
121+ __gm__ Y_TYPE* weightsGmAddr = (__gm__ Y_TYPE*)this->weightsGm_.GetPhyAddr();
122+ asc_vf_call<GmSimtComputeWithWeight<LABEL_TYPE, Y_TYPE>>(dim3{THREAD_NUM}, labelsGmAddr, predictionsGmAddr,
123+ weightsGmAddr, yGmAddr, inputOffset, inputDataLength,
124+ this->numClasses_);
125+ }
126+}
127+ 
128+} // namespace ConfusionMatrixSimt
129+#endif // CONFUSION_MATRIX_SIMT_NOT_FULL_LOAD_GM_H
Amath/confusion_matrix/op_kernel/arch35/confusion_matrix_simt_not_full_load_ub.h+177-0
@@ -0,0 +1,177 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software; you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_simt_not_full_load_ub.h
13+ * \brief The implementation of confusion_matrix by not full load in ub.
14+ */
15+ 
16+#ifndef CONFUSION_MATRIX_SIMT_NOT_FULL_LOAD_UB_H
17+#define CONFUSION_MATRIX_SIMT_NOT_FULL_LOAD_UB_H
18+ 
19+#include "kernel_operator.h"
20+#include "confusion_matrix_simt_base.h"
21+#include "confusion_matrix_tiling_data.h"
22+#include "simt_api/asc_simt.h"
23+#include "simt_api/device_atomic_functions.h"
24+#include "simt_api/asc_fp16.h"
25+#include "simt_api/asc_bf16.h"
26+ 
27+namespace ConfusionMatrixSimt {
28+using namespace AscendC;
29+ 
30+template <typename LABEL_TYPE, typename Y_TYPE>
31+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void UbSimtComputeLoop(
32+ __gm__ LABEL_TYPE* labelsGmAddr, __gm__ LABEL_TYPE* predictionsGmAddr, __ubuf__ Y_TYPE* yLocalAddr,
33+ const int64_t inputOffset, const int64_t inputDataLength, const int64_t ubLoop, const int64_t availableUbSize,
34+ const int64_t numClasses)
35+{
36+ for (int64_t index = static_cast<int64_t>(threadIdx.x); index < inputDataLength;
37+ index += static_cast<int64_t>(blockDim.x)) {
38+ int64_t label = static_cast<int64_t>(labelsGmAddr[inputOffset + index]);
39+ int64_t prediction = static_cast<int64_t>(predictionsGmAddr[inputOffset + index]);
40+ int64_t offset = label * numClasses + prediction;
41+ int64_t currentLoopMinAddr = ubLoop * availableUbSize;
42+ int64_t currentLoopMaxAddr = currentLoopMinAddr + availableUbSize;
43+ if (offset >= currentLoopMinAddr && offset < currentLoopMaxAddr) {
44+ int64_t arrayOffset = offset - currentLoopMinAddr;
45+ asc_atomic_add(yLocalAddr + arrayOffset, static_cast<Y_TYPE>(WEIGHT_ONE));
46+ }
47+ }
48+}
49+ 
50+template <typename LABEL_TYPE, typename Y_TYPE>
51+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void UbSimtComputeWithWeightLoop(
52+ __gm__ LABEL_TYPE* labelsGmAddr, __gm__ LABEL_TYPE* predictionsGmAddr, __gm__ Y_TYPE* weightsGmAddr,
53+ __ubuf__ Y_TYPE* yLocalAddr, const int64_t inputOffset, const int64_t inputDataLength, const int64_t ubLoop,
54+ const int64_t availableUbSize, const int64_t numClasses)
55+{
56+ for (int64_t index = static_cast<int64_t>(threadIdx.x); index < inputDataLength;
57+ index += static_cast<int64_t>(blockDim.x)) {
58+ int64_t label = static_cast<int64_t>(labelsGmAddr[inputOffset + index]);
59+ int64_t prediction = static_cast<int64_t>(predictionsGmAddr[inputOffset + index]);
60+ int64_t offset = label * numClasses + prediction;
61+ int64_t currentLoopMinAddr = ubLoop * availableUbSize;
62+ int64_t currentLoopMaxAddr = currentLoopMinAddr + availableUbSize;
63+ if (offset >= currentLoopMinAddr && offset < currentLoopMaxAddr) {
64+ int64_t arrayOffset = offset - currentLoopMinAddr;
65+ asc_atomic_add(yLocalAddr + arrayOffset, weightsGmAddr[inputOffset + index]);
66+ }
67+ }
68+}
69+ 
70+template <typename LABEL_TYPE, typename Y_TYPE>
71+class ConfusionMatrixSimtBatchLoad : public ConfusionMatrixSimtBase<LABEL_TYPE, Y_TYPE> {
72+public:
73+ __aicore__ inline ConfusionMatrixSimtBatchLoad(){};
74+ __aicore__ inline void Init(GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
75+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe,
76+ bool isWeightEmpty);
77+ __aicore__ inline void Process();
78+ 
79+private:
80+ __aicore__ inline void ResetBins();
81+ __aicore__ inline void Compute();
82+ __aicore__ inline void CopyOut(int64_t offsetInGm, uint32_t stride);
83+ 
84+private:
85+ LocalTensor<Y_TYPE> yLocal_;
86+ TPipe* tPipe_;
87+ TQue<TPosition::VECOUT, OUT_QUE_DEPTH> yQue_;
88+ 
89+ int64_t availableUbSize_ = 0;
90+ int64_t ubLoopCount_ = 0;
91+};
92+ 
93+template <typename LABEL_TYPE, typename Y_TYPE>
94+__aicore__ inline void ConfusionMatrixSimtBatchLoad<LABEL_TYPE, Y_TYPE>::Init(
95+ GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
96+ const ConfusionMatrixTilingData* __restrict tilingData, TPipe* tPipe, bool isWeightEmpty)
97+{
98+ this->BaseInit(labels, predictions, weights, y, tilingData, tPipe, isWeightEmpty);
99+ 
100+ availableUbSize_ = tilingData->ubNumCanUse;
101+ ubLoopCount_ = tilingData->ubLoopNum;
102+ 
103+ tPipe_ = tPipe;
104+ tPipe_->InitBuffer(yQue_, NUM_DOUBLE_BUFFER, this->batchUbBufSize_);
105+}
106+ 
107+template <typename LABEL_TYPE, typename Y_TYPE>
108+__aicore__ inline void ConfusionMatrixSimtBatchLoad<LABEL_TYPE, Y_TYPE>::Process()
109+{
110+ if (this->blockIdx_ >= GetBlockNum()) {
111+ return;
112+ }
113+ 
114+ ResetBins();
115+ SyncAll();
116+ Compute();
117+}
118+ 
119+template <typename LABEL_TYPE, typename Y_TYPE>
120+__aicore__ inline void ConfusionMatrixSimtBatchLoad<LABEL_TYPE, Y_TYPE>::ResetBins()
121+{
122+ int64_t resetBinsHeadIndex = this->blockIdx_ * this->resetBinsLength_;
123+ int64_t resetBinsDataLength = (this->blockIdx_ == this->resetBinsCoreNum_ - 1) ? this->resetBinsTailLength_ :
124+ this->resetBinsLength_;
125+ 
126+ InitOutput<Y_TYPE>(this->yGm_[resetBinsHeadIndex], resetBinsDataLength, static_cast<Y_TYPE>(0));
127+}
128+ 
129+template <typename LABEL_TYPE, typename Y_TYPE>
130+__aicore__ inline void ConfusionMatrixSimtBatchLoad<LABEL_TYPE, Y_TYPE>::Compute()
131+{
132+ if (this->blockIdx_ >= this->needCoreNum_) {
133+ return;
134+ }
135+ 
136+ int64_t inputOffset = this->blockIdx_ * this->formerLength_;
137+ int64_t inputDataLength = (this->blockIdx_ == this->needCoreNum_ - 1) ? this->tailLength_ : this->formerLength_;
138+ __gm__ LABEL_TYPE* labelsGmAddr = (__gm__ LABEL_TYPE*)this->labelsGm_.GetPhyAddr();
139+ __gm__ LABEL_TYPE* predictionsGmAddr = (__gm__ LABEL_TYPE*)this->predictionsGm_.GetPhyAddr();
140+ 
141+ for (int64_t index = 0; index < this->ubLoopCount_; index++) {
142+ int64_t outputLength = (index == this->ubLoopCount_ - 1) ?
143+ (this->outputSize_ - (this->ubLoopCount_ - 1) * this->availableUbSize_) :
144+ this->availableUbSize_;
145+ yLocal_ = yQue_.template AllocTensor<Y_TYPE>();
146+ Duplicate<Y_TYPE>(yLocal_, 0, outputLength);
147+ 
148+ __ubuf__ Y_TYPE* yLocalAddr = (__ubuf__ Y_TYPE*)yLocal_.GetPhyAddr();
149+ if (this->isWeightEmpty_) {
150+ asc_vf_call<UbSimtComputeLoop<LABEL_TYPE, Y_TYPE>>(dim3{THREAD_NUM}, labelsGmAddr, predictionsGmAddr,
151+ yLocalAddr, inputOffset, inputDataLength, index,
152+ this->availableUbSize_, this->numClasses_);
153+ } else {
154+ __gm__ Y_TYPE* weightsGmAddr = (__gm__ Y_TYPE*)this->weightsGm_.GetPhyAddr();
155+ asc_vf_call<UbSimtComputeWithWeightLoop<LABEL_TYPE, Y_TYPE>>(
156+ dim3{THREAD_NUM}, labelsGmAddr, predictionsGmAddr, weightsGmAddr, yLocalAddr, inputOffset,
157+ inputDataLength, index, this->availableUbSize_, this->numClasses_);
158+ }
159+ 
160+ CopyOut(index * this->availableUbSize_, static_cast<uint32_t>(outputLength * sizeof(Y_TYPE)));
161+ yQue_.template FreeTensor<Y_TYPE>(yLocal_);
162+ }
163+}
164+ 
165+template <typename LABEL_TYPE, typename Y_TYPE>
166+__aicore__ inline void ConfusionMatrixSimtBatchLoad<LABEL_TYPE, Y_TYPE>::CopyOut(int64_t offsetInGm, uint32_t stride)
167+{
168+ yQue_.EnQue(yLocal_);
169+ yLocal_ = yQue_.template DeQue<Y_TYPE>();
170+ SetAtomicAdd<Y_TYPE>();
171+ DataCopyExtParams dataCopyExtParams{DATA_COPY_PAD_BLOCK_COUNT, stride, 0, 0, 0};
172+ DataCopyPad(this->yGm_[offsetInGm], yLocal_, dataCopyExtParams);
173+ SetAtomicNone();
174+}
175+ 
176+} // namespace ConfusionMatrixSimt
177+#endif // CONFUSION_MATRIX_SIMT_NOT_FULL_LOAD_UB_H
Amath/confusion_matrix/op_kernel/arch35/confusion_matrix_tiling_data.h+36-0
@@ -0,0 +1,36 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_tiling_data.h
13+ * \brief confusion_matrix_tiling_data.h
14+ */
15+ 
16+#ifndef _CONFUSION_MATRIX_TILING_DATA_H_
17+#define _CONFUSION_MATRIX_TILING_DATA_H_
18+ 
19+struct ConfusionMatrixTilingData {
20+ int64_t numClasses = 0;
21+ int64_t inputSize = 0;
22+ int64_t ubNumCanUse = 0;
23+ int64_t ubLoopNum = 0;
24+ int64_t needXCoreNum = 0;
25+ int64_t formerLength = 0;
26+ int64_t tailLength = 0;
27+ int64_t clearYCoreNum = 0;
28+ int64_t clearYFactor = 0;
29+ int64_t clearYTail = 0;
30+ int64_t binsFormerLength = 0;
31+ int64_t needBinsCoreNum = 0;
32+ int64_t binsTailLength = 0;
33+ int64_t fullLoadBufSize = 0;
34+ int64_t batchUbBufSize = 0;
35+};
36+#endif
Amath/confusion_matrix/op_kernel/arch35/confusion_matrix_tiling_key.h+51-0
@@ -0,0 +1,51 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix_tiling_key.h
13+ * \brief confusion_matrix tiling key declare
14+ */
15+ 
16+#ifndef _CONFUSION_MATRIX_TILING_KEY_DECL_H_
17+#define _CONFUSION_MATRIX_TILING_KEY_DECL_H_
18+#include "ascendc/host_api/tiling/template_argument.h"
19+ 
20+#define TPL_SCH_ID_FULL_LOAD 0 // full load to UB
21+#define TPL_SCH_ID_BATCH_LOAD 1 // batch load to UB
22+#define TPL_SCH_ID_NOT_FULL_LOAD 2 // not load to UB
23+#define TPL_SCH_ID_DETERMINE 3 // determine
24+ 
25+#define TPL_OUTPUT_DTYPE_FLOAT 0
26+#define TPL_OUTPUT_DTYPE_INT32 1
27+#define TPL_OUTPUT_DTYPE_FLOAT16 2
28+#define TPL_OUTPUT_DTYPE_INT8 3
29+#define TPL_OUTPUT_DTYPE_UINT8 4
30+ 
31+#define TPL_EMPTY_WEIGHT 0
32+#define TPL_HAS_WEIGHT 1
33+ 
34+#define CONFUSION_MATRIX_TPL_KEY_DECL() \
35+ ASCENDC_TPL_UINT_DECL(schId, ASCENDC_TPL_8_BW, ASCENDC_TPL_UI_LIST, TPL_SCH_ID_FULL_LOAD, TPL_SCH_ID_BATCH_LOAD, \
36+ TPL_SCH_ID_NOT_FULL_LOAD, TPL_SCH_ID_DETERMINE), \
37+ ASCENDC_TPL_UINT_DECL(outputDtype, ASCENDC_TPL_8_BW, ASCENDC_TPL_UI_LIST, TPL_OUTPUT_DTYPE_FLOAT, \
38+ TPL_OUTPUT_DTYPE_INT32, TPL_OUTPUT_DTYPE_FLOAT16, TPL_OUTPUT_DTYPE_INT8, \
39+ TPL_OUTPUT_DTYPE_UINT8), \
40+ ASCENDC_TPL_UINT_DECL(isWeight, ASCENDC_TPL_8_BW, ASCENDC_TPL_UI_LIST, 0, 1),
41+ 
42+#define CONFUSION_MATRIX_TPL_KEY_SEL() \
43+ ASCENDC_TPL_UINT_SEL(schId, ASCENDC_TPL_UI_RANGE, 1, 0, TPL_SCH_ID_DETERMINE), \
44+ ASCENDC_TPL_UINT_SEL(outputDtype, ASCENDC_TPL_UI_LIST, TPL_OUTPUT_DTYPE_FLOAT, TPL_OUTPUT_DTYPE_INT32, \
45+ TPL_OUTPUT_DTYPE_FLOAT16, TPL_OUTPUT_DTYPE_INT8, TPL_OUTPUT_DTYPE_UINT8), \
46+ ASCENDC_TPL_UINT_SEL(isWeight, ASCENDC_TPL_UI_LIST, TPL_EMPTY_WEIGHT, TPL_HAS_WEIGHT),
47+ 
48+ASCENDC_TPL_ARGS_DECL(ConfusionMatrix, CONFUSION_MATRIX_TPL_KEY_DECL());
49+ 
50+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(CONFUSION_MATRIX_TPL_KEY_SEL()));
51+#endif
Amath/confusion_matrix/op_kernel/confusion_matrix.cpp+148-0
@@ -0,0 +1,148 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software; you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file confusion_matrix.cpp
13+ * \brief The kernel function of confusion_matrix and the entrance for using confusion_matrix.
14+ */
15+ 
16+#include "kernel_tiling/kernel_tiling.h"
17+#include "kernel_operator.h"
18+#include "arch35/confusion_matrix_tiling_data.h"
19+#include "arch35/confusion_matrix_tiling_key.h"
20+#include "arch35/confusion_matrix_simt_base.h"
21+#include "arch35/confusion_matrix_simt_full_load.h"
22+#include "arch35/confusion_matrix_simt_not_full_load_ub.h"
23+#include "arch35/confusion_matrix_simt_not_full_load_gm.h"
24+#include "arch35/confusion_matrix_determine.h"
25+ 
26+using namespace AscendC;
27+ 
28+#define TPL_SCH_ID_FULL_LOAD 0 // full load to UB
29+#define TPL_SCH_ID_BATCH_LOAD 1 // batch load to UB
30+#define TPL_SCH_ID_NOT_FULL_LOAD 2 // not load to UB
31+#define TPL_SCH_ID_DETERMINE 3 // determine
32+ 
33+#define TPL_OUTPUT_DTYPE_FLOAT 0
34+#define TPL_OUTPUT_DTYPE_INT32 1
35+#define TPL_OUTPUT_DTYPE_FLOAT16 2
36+#define TPL_OUTPUT_DTYPE_INT8 3
37+#define TPL_OUTPUT_DTYPE_UINT8 4
38+ 
39+#define TPL_EMPTY_WEIGHT 0
40+#define TPL_HAS_WEIGHT 1
41+ 
42+template <uint64_t schId, uint64_t outputDtype, uint64_t isWeight>
43+__global__ __aicore__ void confusion_matrix(GM_ADDR labels, GM_ADDR predictions, GM_ADDR weights, GM_ADDR y,
44+ GM_ADDR workspace, GM_ADDR tiling)
45+{
46+ if (g_coreType == AscendC::AIC) {
47+ return;
48+ }
49+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
50+ 
51+ using LABEL_TYPE = DTYPE_LABELS;
52+ 
53+ AscendC::TPipe tPipe;
54+ 
55+ REGISTER_TILING_DEFAULT(ConfusionMatrixTilingData);
56+ GET_TILING_DATA_WITH_STRUCT(ConfusionMatrixTilingData, tilingDataIn, tiling);
57+ const ConfusionMatrixTilingData* __restrict tilingData = &tilingDataIn;
58+ 
59+ bool isWeightEmpty = isWeight == TPL_EMPTY_WEIGHT ? true : false;
60+ 
61+ // For int8/uint8 output: only DETERMINE mode is supported.
62+ // asc_atomic_add does not support int8_t/uint8_t, so FULL_LOAD,
63+ // BATCH_LOAD, and NOT_FULL_LOAD (which rely on asc_atomic_add) cannot
64+ // be used. Tiling forces DETERMINE (schId=3) for these types.
65+ // DETERMINE mode uses direct += on disjoint output ranges per core,
66+ // so no atomic operations are needed and int8_t/uint8_t work correctly.
67+ if constexpr (outputDtype == TPL_OUTPUT_DTYPE_INT8 || outputDtype == TPL_OUTPUT_DTYPE_UINT8) {
68+ if (schId == TPL_SCH_ID_DETERMINE) {
69+ if (outputDtype == TPL_OUTPUT_DTYPE_INT8) {
70+ ConfusionMatrixSimt::ConfusionMatrixDetermine<LABEL_TYPE, int8_t> op;
71+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
72+ op.Process();
73+ } else {
74+ ConfusionMatrixSimt::ConfusionMatrixDetermine<LABEL_TYPE, uint8_t> op;
75+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
76+ op.Process();
77+ }
78+ }
79+ } else {
80+ // FULL_LOAD
81+ if (schId == TPL_SCH_ID_FULL_LOAD && outputDtype == TPL_OUTPUT_DTYPE_FLOAT) {
82+ ConfusionMatrixSimt::ConfusionMatrixSimtFullLoad<LABEL_TYPE, float> op;
83+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
84+ op.Process();
85+ }
86+ if (schId == TPL_SCH_ID_FULL_LOAD && outputDtype == TPL_OUTPUT_DTYPE_INT32) {
87+ ConfusionMatrixSimt::ConfusionMatrixSimtFullLoad<LABEL_TYPE, int32_t> op;
88+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
89+ op.Process();
90+ }
91+ if (schId == TPL_SCH_ID_FULL_LOAD && outputDtype == TPL_OUTPUT_DTYPE_FLOAT16) {
92+ ConfusionMatrixSimt::ConfusionMatrixSimtFullLoad<LABEL_TYPE, half> op;
93+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
94+ op.Process();
95+ }
96+ 
97+ // BATCH_LOAD
98+ if (schId == TPL_SCH_ID_BATCH_LOAD && outputDtype == TPL_OUTPUT_DTYPE_FLOAT) {
99+ ConfusionMatrixSimt::ConfusionMatrixSimtBatchLoad<LABEL_TYPE, float> op;
100+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
101+ op.Process();
102+ }
103+ if (schId == TPL_SCH_ID_BATCH_LOAD && outputDtype == TPL_OUTPUT_DTYPE_INT32) {
104+ ConfusionMatrixSimt::ConfusionMatrixSimtBatchLoad<LABEL_TYPE, int32_t> op;
105+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
106+ op.Process();
107+ }
108+ if (schId == TPL_SCH_ID_BATCH_LOAD && outputDtype == TPL_OUTPUT_DTYPE_FLOAT16) {
109+ ConfusionMatrixSimt::ConfusionMatrixSimtBatchLoad<LABEL_TYPE, half> op;
110+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
111+ op.Process();
112+ }
113+ 
114+ // NOT_FULL_LOAD (GM)
115+ if (schId == TPL_SCH_ID_NOT_FULL_LOAD && outputDtype == TPL_OUTPUT_DTYPE_FLOAT) {
116+ ConfusionMatrixSimt::ConfusionMatrixSimtNotFullLoadGm<LABEL_TYPE, float> op;
117+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
118+ op.Process();
119+ }
120+ if (schId == TPL_SCH_ID_NOT_FULL_LOAD && outputDtype == TPL_OUTPUT_DTYPE_INT32) {
121+ ConfusionMatrixSimt::ConfusionMatrixSimtNotFullLoadGm<LABEL_TYPE, int32_t> op;
122+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
123+ op.Process();
124+ }
125+ if (schId == TPL_SCH_ID_NOT_FULL_LOAD && outputDtype == TPL_OUTPUT_DTYPE_FLOAT16) {
126+ ConfusionMatrixSimt::ConfusionMatrixSimtNotFullLoadGm<LABEL_TYPE, half> op;
127+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
128+ op.Process();
129+ }
130+ 
131+ // DETERMINE
132+ if (schId == TPL_SCH_ID_DETERMINE && outputDtype == TPL_OUTPUT_DTYPE_FLOAT) {
133+ ConfusionMatrixSimt::ConfusionMatrixDetermine<LABEL_TYPE, float> op;
134+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
135+ op.Process();
136+ }
137+ if (schId == TPL_SCH_ID_DETERMINE && outputDtype == TPL_OUTPUT_DTYPE_INT32) {
138+ ConfusionMatrixSimt::ConfusionMatrixDetermine<LABEL_TYPE, int32_t> op;
139+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
140+ op.Process();
141+ }
142+ if (schId == TPL_SCH_ID_DETERMINE && outputDtype == TPL_OUTPUT_DTYPE_FLOAT16) {
143+ ConfusionMatrixSimt::ConfusionMatrixDetermine<LABEL_TYPE, half> op;
144+ op.Init(labels, predictions, weights, y, tilingData, &tPipe, isWeightEmpty);
145+ op.Process();
146+ }
147+ }
148+}
Amath/confusion_matrix/tests/assets/golden.py+184-0
@@ -0,0 +1,184 @@
1+#!/usr/bin/env python3
2+# -*- coding: UTF-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+"""ConfusionMatrix 多通路 golden（TestSpec 范式）。
13+ 
14+通路支持表（本算子交付面）：
15+| 通路 | 支持 | 依据 |
16+|--------|------|------|
17+| kernel | ✅ | op_kernel/arch35 有实现 |
18+| geir | ✅ | op_graph/confusion_matrix_proto.h 有 REG_OP(ConfusionMatrix) |
19+| aclnn | ❌ | Ascend910B/910C built-in 无 aclnn 接口，本算子对齐该支持面不交付 |
20+| e2e | ❌ | torch_npu 无绑定（无 aclnn 接口 → 无可调用入口） |
21+ 
22+int8/uint8 输出：整型计数语义、差 1 就是错 → binary_equal（非量化场景，输入含整型，
23+不适用 quant 的浮点输入前提）。int8/uint8 中间累加用 int32 防 scatter_add 回绕。
24+"""
25+ 
26+__spec__ = {
27+ "confusion_matrix": "ConfusionMatrixKernelSpec",
28+}
29+ 
30+import numpy as np
31+import torch
32+ 
33+_DTYPE_MAP = {
34+ "float32": torch.float32,
35+ "float64": torch.float64,
36+ "float16": torch.float16,
37+ "int32": torch.int32,
38+ "int8": torch.int8,
39+ "uint8": torch.uint8,
40+}
41+ 
42+_TOL = {
43+ "float32": {"standard": "cross_check", "level": "L1"},
44+ "float16": {"standard": "cross_check", "level": "L1"},
45+}
46+ 
47+ 
48+def _attr(kwargs, name, default):
49+ """CSV 里的 attributes 都是字符串，统一转成 default 的类型。"""
50+ v = kwargs.get(name, default)
51+ if isinstance(v, str):
52+ s = v.strip().lower()
53+ if s in ("true", "false", "yes", "no", "1", "0"):
54+ return s in ("true", "yes", "1")
55+ try:
56+ return type(default)(v)
57+ except Exception:
58+ return default
59+ return v
60+ 
61+ 
62+def _normalize(val):
63+ if isinstance(val, (np.ndarray, np.generic)):
64+ return val.item() if val.ndim == 0 else val.tolist()
65+ if isinstance(val, bytes):
66+ return val.decode()
67+ return val
68+ 
69+ 
70+def _resolve_compute_dtype(**kwargs):
71+ """累加精度基准：取 TTK 下发的 output_dtypes（Promote 后的值）。
72+ 
73+ 三方比对（cross_check）下 TTK 已按 Promote 抬高 output_dtypes（fp16->fp32、
74+ fp32->fp64），golden 据此在抬高后的精度上累加出高精度真值；两方比对未抬高，
75+ 则跟随算子声明精度。取不到 output_dtypes 时回落算子声明的 dtype。
76+ 返回 (compute_dtype, final_dtype)：int8/uint8 时 compute 为 int32 中转、
77+ final 为目标 dtype；其余两者相同。
78+ """
79+ dtype_str = _normalize(kwargs.get("dtype", "float32"))
80+ final_dtype = _DTYPE_MAP.get(dtype_str, torch.float32)
81+ od = kwargs.get("output_dtypes") or []
82+ if od:
83+ first = od[0]
84+ if isinstance(first, (list, tuple)):
85+ first = first[0] if first else None
86+ if first is not None:
87+ final_dtype = _DTYPE_MAP.get(str(first), final_dtype)
88+ compute_dtype = (
89+ torch.int32 if final_dtype in (torch.int8, torch.uint8) else final_dtype
90+ )
91+ return compute_dtype, final_dtype
92+ 
93+ 
94+def _compute(labels, predictions, weights, **kwargs):
95+ """计算核（torch 进 torch 出）。golden 用 torch.scatter_add 竞品接口拼接。
96+ 
97+ 累加精度按 _resolve_compute_dtype：三方比对下在 Promote 抬高后的精度上算
98+ 高精度真值；两方比对跟随算子声明精度。int8/uint8 输出用 int32 中间累加再
99+ cast 回（加法环同态下与窄整型累加等价，规避 torch 窄整型 scatter_add 的
100+ 中间量行为差异）。
101+ """
102+ num_classes = _attr(kwargs, "num_classes", 0)
103+ num_classes = _normalize(num_classes)
104+ 
105+ labels_int = labels.to(torch.int64)
106+ preds_int = predictions.to(torch.int64)
107+ flat_index = labels_int * num_classes + preds_int
108+ 
109+ compute_dtype, final_dtype = _resolve_compute_dtype(**kwargs)
110+ 
111+ y_flat = torch.zeros(num_classes * num_classes, dtype=compute_dtype)
112+ if weights is not None and weights.numel() > 0:
113+ y_flat.scatter_add_(0, flat_index, weights.to(compute_dtype))
114+ else:
115+ y_flat.scatter_add_(
116+ 0, flat_index, torch.ones(labels_int.shape[0], dtype=compute_dtype)
117+ )
118+ 
119+ if compute_dtype != final_dtype:
120+ y_flat = y_flat.to(final_dtype)
121+ return [y_flat.view(num_classes, num_classes)]
122+ 
123+ 
124+class _Compose:
125+ """三方标杆（远端 GPU server 执行）。属性喂 __init__、输入喂 __call__。
126+ 
127+ 参数名与 proto REG_OP 注册名逐字一致（labels/predictions/weights/num_classes/dtype）。
128+ 实现与 _compute 同为 torch.scatter_add 拼接（本算子即竞品拼接型，无更高层 API）。
129+ 浮点输出 cast 回 NPU 输出 dtype（防 ratio 失真）。
130+ """
131+ 
132+ def __init__(self, **kwargs):
133+ self.num_classes = _attr(kwargs, "num_classes", 0)
134+ self.num_classes = _normalize(self.num_classes)
135+ self.dtype = _normalize(kwargs.get("dtype", "float32"))
136+ self.out_dtype = _DTYPE_MAP.get(self.dtype, torch.float32)
137+ 
138+ def __call__(self, labels, predictions, weights, **kwargs):
139+ nc = self.num_classes
140+ flat_index = labels.to(torch.int64) * nc + predictions.to(torch.int64)
141+ compute_dtype = self.out_dtype
142+ if compute_dtype in (torch.int8, torch.uint8):
143+ compute_dtype = torch.int32
144+ y_flat = torch.zeros(nc * nc, dtype=compute_dtype, device=labels.device)
145+ if weights is not None and weights.numel() > 0:
146+ y_flat.scatter_add_(0, flat_index, weights.to(compute_dtype))
147+ else:
148+ y_flat.scatter_add_(
149+ 0,
150+ flat_index,
151+ torch.ones(
152+ flat_index.shape[0], dtype=compute_dtype, device=labels.device
153+ ),
154+ )
155+ if compute_dtype != self.out_dtype:
156+ y_flat = y_flat.to(self.out_dtype)
157+ return [y_flat.view(nc, nc)]
158+ 
159+ 
160+class ConfusionMatrixKernelSpec:
161+ """kernel + geir 共用。golden 收 numpy.ndarray，返 numpy.ndarray。
162+ 参数名照 op_host/confusion_matrix_def.cpp（labels/predictions/weights + num_classes/dtype）。
163+ """
164+ 
165+ def golden(*inputs, **kwargs):
166+ t = [
167+ torch.from_numpy(np.ascontiguousarray(a)) if a is not None else None
168+ for a in inputs
169+ ]
170+ outs = _compute(*t, **kwargs)
171+ od = kwargs.get("output_dtypes") or []
172+ od = [d[0] if isinstance(d, (list, tuple)) else str(d) for d in od]
173+ return [
174+ o.numpy().astype(od[i]) if i < len(od) else o.numpy()
175+ for i, o in enumerate(outs)
176+ ]
177+ 
178+ third_party = {"torch": _Compose}
179+ tolerance = _TOL
180+ 
181+ 
182+# 【不存在】aclnn 通路：Ascend910B/910C built-in 无 aclnn 源码 → 对齐该支持面，本算子不交付 aclnn。
183+# 【不存在】e2e 通路：torch_npu 无本算子绑定。无 aclnn 源码 → 无编译产物 → torch_npu 无可
184+# 调用入口；验证环境 torch_npu dispatcher 全表 grep "confusion" 无命中。
Amath/confusion_matrix/tests/st/arch35/ttk_kernel_confusion_matrix_st.csv+13-0
@@ -0,0 +1,13 @@
1+testcase_name,network_name,op_name,input_shapes,input_dtypes,input_formats,output_shapes,output_dtypes,output_formats,input_ori_shapes,input_ori_formats,output_ori_shapes,output_ori_formats,attributes,input_data_ranges,precision_tolerances,absolute_precision,output_inplace_indexes,output_shape_unknown_indexes,is_enabled,remark,soc_series,priority,dump_file_prefix,manual_input_binaries,manual_golden_binaries
2+confusion_matrix_000,,confusion_matrix,"[[100], [100], [100]]","['int32','int32','float32']","['ND','ND','ND']","[[5,5]]",['float32'],['ND'],,,,,"{'num_classes':5,'dtype':'float32'}","((0,4),(0,4),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
3+confusion_matrix_001,,confusion_matrix,"[[100], [100], [100]]","['int32','int32','int32']","['ND','ND','ND']","[[5,5]]",['int32'],['ND'],,,,,"{'num_classes':5,'dtype':'int32'}","((0,4),(0,4),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
4+confusion_matrix_002,,confusion_matrix,"[[100], [100], [100]]","['int32','int32','float16']","['ND','ND','ND']","[[5,5]]",['float16'],['ND'],,,,,"{'num_classes':5,'dtype':'float16'}","((0,4),(0,4),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
5+confusion_matrix_003,,confusion_matrix,"[[100], [100], [100]]","['int32','int32','int8']","['ND','ND','ND']","[[5,5]]",['int8'],['ND'],,,,,"{'num_classes':5,'dtype':'int8'}","((0,4),(0,4),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
6+confusion_matrix_004,,confusion_matrix,"[[100], [100], [100]]","['int32','int32','uint8']","['ND','ND','ND']","[[5,5]]",['uint8'],['ND'],,,,,"{'num_classes':5,'dtype':'uint8'}","((0,4),(0,4),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
7+confusion_matrix_005,,confusion_matrix,"[[100], [100], [0]]","['int32','int32','float32']","['ND','ND','ND']","[[10,10]]",['float32'],['ND'],,,,,"{'num_classes':10,'dtype':'float32'}","((0,9),(0,9),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
8+confusion_matrix_006,,confusion_matrix,"[[1000], [1000], [1000]]","['int32','int32','float32']","['ND','ND','ND']","[[10,10]]",['float32'],['ND'],,,,,"{'num_classes':10,'dtype':'float32'}","((0,9),(0,9),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
9+confusion_matrix_007,,confusion_matrix,"[[1], [1], [1]]","['int32','int32','float32']","['ND','ND','ND']","[[5,5]]",['float32'],['ND'],,,,,"{'num_classes':5,'dtype':'float32'}","((0,4),(0,4),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
10+confusion_matrix_008,,confusion_matrix,"[[10000], [10000], [10000]]","['int32','int32','float32']","['ND','ND','ND']","[[100,100]]",['float32'],['ND'],,,,,"{'num_classes':100,'dtype':'float32'}","((0,99),(0,99),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
11+confusion_matrix_009,,confusion_matrix,"[[5000], [5000], [5000]]","['float16','float16','float16']","['ND','ND','ND']","[[50,50]]",['float16'],['ND'],,,,,"{'num_classes':50,'dtype':'float16'}","((0,49),(0,49),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
12+confusion_matrix_010,,confusion_matrix,"[[5000], [5000], [5000]]","['int8','int8','int8']","['ND','ND','ND']","[[50,50]]",['int8'],['ND'],,,,,"{'num_classes':50,'dtype':'int8'}","((0,49),(0,49),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
13+confusion_matrix_011,,confusion_matrix,"[[5000], [5000], [5000]]","['uint8','uint8','uint8']","['ND','ND','ND']","[[50,50]]",['uint8'],['ND'],,,,,"{'num_classes':50,'dtype':'uint8'}","((0,49),(0,49),(0,1))",,1.00E-03,(),(),True,,,0,,(),()
Amath/confusion_matrix/tests/ut/op_host/arch35/test_confusion_matrix_tiling_arch35.cpp+211-0
@@ -0,0 +1,211 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <gtest/gtest.h>
13+#include "math/confusion_matrix/op_host/arch35/confusion_matrix_tiling.h"
14+#include "tiling_context_faker.h"
15+#include "tiling_case_executor.h"
16+#include "any_value.h"
17+ 
18+using namespace std;
19+ 
20+class ConfusionMatrixTilingTest : public testing::Test {
21+protected:
22+ static void SetUpTestCase() { std::cout << "ConfusionMatrixTilingTest SetUp" << std::endl; }
23+ 
24+ static void TearDownTestCase() { std::cout << "ConfusionMatrixTilingTest TearDown" << std::endl; }
25+};
26+ 
27+static gert::TilingContextPara::OpAttr MakeNumClassesAttr(int64_t numClasses)
28+{
29+ return gert::TilingContextPara::OpAttr("num_classes", Ops::Math::AnyValue::CreateFrom<int64_t>(numClasses));
30+}
31+ 
32+static gert::TilingContextPara::OpAttr MakeDtypeAttr(const std::string& dtype)
33+{
34+ return gert::TilingContextPara::OpAttr("dtype", Ops::Math::AnyValue::CreateFrom<std::string>(dtype));
35+}
36+ 
37+TEST_F(ConfusionMatrixTilingTest, confusion_matrix_tiling_float32_with_weights)
38+{
39+ optiling::AscendCConfusionMatrixCompileInfo compileInfo = {64, 262144};
40+ gert::TilingContextPara tilingContextPara("ConfusionMatrix",
41+ {
42+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
43+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
44+ {{{100}, {100}}, ge::DT_FLOAT, ge::FORMAT_ND},
45+ },
46+ {
47+ {{{5, 5}, {5, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
48+ },
49+ {
50+ MakeNumClassesAttr(5),
51+ MakeDtypeAttr("float32"),
52+ },
53+ &compileInfo);
54+ uint64_t expectTilingKey = 65536;
55+ std::vector<size_t> expectWorkspaces = {16777216};
56+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
57+}
58+ 
59+TEST_F(ConfusionMatrixTilingTest, confusion_matrix_tiling_float32_no_weights)
60+{
61+ optiling::AscendCConfusionMatrixCompileInfo compileInfo = {64, 262144};
62+ gert::TilingContextPara tilingContextPara("ConfusionMatrix",
63+ {
64+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
65+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
66+ {{{0}, {0}}, ge::DT_FLOAT, ge::FORMAT_ND},
67+ },
68+ {
69+ {{{5, 5}, {5, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
70+ },
71+ {
72+ MakeNumClassesAttr(5),
73+ MakeDtypeAttr("float32"),
74+ },
75+ &compileInfo);
76+ uint64_t expectTilingKey = 0;
77+ std::vector<size_t> expectWorkspaces = {16777216};
78+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
79+}
80+ 
81+TEST_F(ConfusionMatrixTilingTest, confusion_matrix_tiling_int32_with_weights)
82+{
83+ optiling::AscendCConfusionMatrixCompileInfo compileInfo = {64, 262144};
84+ gert::TilingContextPara tilingContextPara("ConfusionMatrix",
85+ {
86+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
87+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
88+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
89+ },
90+ {
91+ {{{10, 10}, {10, 10}}, ge::DT_INT32, ge::FORMAT_ND},
92+ },
93+ {
94+ MakeNumClassesAttr(10),
95+ MakeDtypeAttr("int32"),
96+ },
97+ &compileInfo);
98+ uint64_t expectTilingKey = 65792;
99+ std::vector<size_t> expectWorkspaces = {16777216};
100+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
101+}
102+ 
103+TEST_F(ConfusionMatrixTilingTest, confusion_matrix_tiling_int32_no_weights)
104+{
105+ optiling::AscendCConfusionMatrixCompileInfo compileInfo = {64, 262144};
106+ gert::TilingContextPara tilingContextPara("ConfusionMatrix",
107+ {
108+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
109+ {{{100}, {100}}, ge::DT_INT32, ge::FORMAT_ND},
110+ {{{0}, {0}}, ge::DT_INT32, ge::FORMAT_ND},
111+ },
112+ {
113+ {{{10, 10}, {10, 10}}, ge::DT_INT32, ge::FORMAT_ND},
114+ },
115+ {
116+ MakeNumClassesAttr(10),
117+ MakeDtypeAttr("int32"),
118+ },
119+ &compileInfo);
120+ uint64_t expectTilingKey = 256;
121+ std::vector<size_t> expectWorkspaces = {16777216};
122+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
123+}
124+ 
125+TEST_F(ConfusionMatrixTilingTest, confusion_matrix_tiling_float16_with_weights)
126+{
127+ optiling::AscendCConfusionMatrixCompileInfo compileInfo = {64, 262144};
128+ gert::TilingContextPara tilingContextPara("ConfusionMatrix",
129+ {
130+ {{{100}, {100}}, ge::DT_FLOAT16, ge::FORMAT_ND},
131+ {{{100}, {100}}, ge::DT_FLOAT16, ge::FORMAT_ND},
132+ {{{100}, {100}}, ge::DT_FLOAT16, ge::FORMAT_ND},
133+ },
134+ {
135+ {{{5, 5}, {5, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
136+ },
137+ {
138+ MakeNumClassesAttr(5),
139+ MakeDtypeAttr("float16"),
140+ },
141+ &compileInfo);
142+ uint64_t expectTilingKey = 66048;
143+ std::vector<size_t> expectWorkspaces = {16777216};
144+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
145+}
146+ 
147+TEST_F(ConfusionMatrixTilingTest, confusion_matrix_tiling_int8_no_weights)
148+{
149+ optiling::AscendCConfusionMatrixCompileInfo compileInfo = {64, 262144};
150+ gert::TilingContextPara tilingContextPara("ConfusionMatrix",
151+ {
152+ {{{100}, {100}}, ge::DT_INT8, ge::FORMAT_ND},
153+ {{{100}, {100}}, ge::DT_INT8, ge::FORMAT_ND},
154+ {{{0}, {0}}, ge::DT_INT8, ge::FORMAT_ND},
155+ },
156+ {
157+ {{{10, 10}, {10, 10}}, ge::DT_INT8, ge::FORMAT_ND},
158+ },
159+ {
160+ MakeNumClassesAttr(10),
161+ MakeDtypeAttr("int8"),
162+ },
163+ &compileInfo);
164+ uint64_t expectTilingKey = 771;
165+ std::vector<size_t> expectWorkspaces = {16777216};
166+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
167+}
168+ 
169+TEST_F(ConfusionMatrixTilingTest, confusion_matrix_tiling_uint8_with_weights)
170+{
171+ optiling::AscendCConfusionMatrixCompileInfo compileInfo = {64, 262144};
172+ gert::TilingContextPara tilingContextPara("ConfusionMatrix",
173+ {
174+ {{{100}, {100}}, ge::DT_UINT8, ge::FORMAT_ND},
175+ {{{100}, {100}}, ge::DT_UINT8, ge::FORMAT_ND},
176+ {{{100}, {100}}, ge::DT_UINT8, ge::FORMAT_ND},
177+ },
178+ {
179+ {{{10, 10}, {10, 10}}, ge::DT_UINT8, ge::FORMAT_ND},
180+ },
181+ {
182+ MakeNumClassesAttr(10),
183+ MakeDtypeAttr("uint8"),
184+ },
185+ &compileInfo);
186+ uint64_t expectTilingKey = 66563;
187+ std::vector<size_t> expectWorkspaces = {16777216};
188+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
189+}
190+ 
191+TEST_F(ConfusionMatrixTilingTest, confusion_matrix_tiling_large_num_classes)
192+{
193+ optiling::AscendCConfusionMatrixCompileInfo compileInfo = {64, 262144};
194+ gert::TilingContextPara tilingContextPara("ConfusionMatrix",
195+ {
196+ {{{10000}, {10000}}, ge::DT_INT32, ge::FORMAT_ND},
197+ {{{10000}, {10000}}, ge::DT_INT32, ge::FORMAT_ND},
198+ {{{10000}, {10000}}, ge::DT_FLOAT, ge::FORMAT_ND},
199+ },
200+ {
201+ {{{1000, 1000}, {1000, 1000}}, ge::DT_FLOAT, ge::FORMAT_ND},
202+ },
203+ {
204+ MakeNumClassesAttr(1000),
205+ MakeDtypeAttr("float32"),
206+ },
207+ &compileInfo);
208+ uint64_t expectTilingKey = 65538;
209+ std::vector<size_t> expectWorkspaces = {16777216};
210+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
211+}