Star542
368
代码
代码
Issues17
Pull Requests35
Actions
项目讨论
Wiki
分析
Star542
368
已合并
feat: 新增 ascend950 BatchMultiClassNonMaxSuppression 算子 #1176
Huang-Peng创建于 25 天前
feat: 新增 ascend950 BatchMultiClassNonMaxSuppression 算子 #1176
已合并
Huang-Peng创建于 25 天前
从Huang-Peng/ops-cv-publicfeat/ascend950-batch-multiclass-nms合入到cann/ops-cvmaster
共 23 个文件变更+3188-60
Mcommon/inc/op_graph/op_cv_proto_extend.h+0-60
@@ -225,66 +225,6 @@ REG_OP(NonMaxSuppressionV4)
225 .OUTPUT(valid_outputs, TensorType({DT_INT32}))225 .OUTPUT(valid_outputs, TensorType({DT_INT32}))
226 .ATTR(pad_to_max_output_size, Bool, false)226 .ATTR(pad_to_max_output_size, Bool, false)
227 .OP_END_FACTORY_REG(NonMaxSuppressionV4)227 .OP_END_FACTORY_REG(NonMaxSuppressionV4)
228-/* ===== Migrated from canndev ops_proto_legacy.h (macro guards stripped) ===== */
229- 
230-/**
231-*@brief Computes nms for input boxes and score, support multiple batch and classes.
232-* will do clip to window, score filter, top_k, and nms
233- 
234-*@par Inputs:
235-* Four inputs, including:
236-*@li boxes: boxes, a 4D Tensor of type float16 or float32 with
237-* shape (batch, num_anchors, num_classes, 4). "batch" indicates the batch size of image,
238-* and "num_anchors" indicates num of boxes, and "num_classes" indicates classes of detect.
239-* and the value "4" refers to "x0", "x1", "y0", and "y1".
240-*@li scores: boxes, a 4D Tensor of type float16 or float32 with
241-* shape (batch, num_anchors, num_classes).
242-*@li clip_window: window size, a 2D Tensor of type float16 or float32 with
243-* shape (batch, 4). 4" refers to "anchor_x0", "anchor_x1", "anchor_y0", and "anchor_y1".
244-*@li num_valid_boxes: valid boxes number for each batch, a 1D Tensor of type int32 with
245-* shape (batch,) . \n
246- 
247-*@par Attributes:
248-*@li score_threshold: A required attribute of type float32, specifying the score filter iou iou_threshold.
249-*@li iou_threshold: A required attribute of type float32, specifying the nms iou iou_threshold.
250-*@li max_size_per_class: A required attribute of type int, specifying the nms output num per class.
251-*@li max_total_size: A required attribute of type int, specifying the the nms output num per batch.
252-*@li change_coordinate_frame: A optional attribute of type bool, whether to normalize coordinates after clipping.
253-* @li transpose_box: A optional attribute of type bool, whether inserted transpose before this op. must be "false".
254-* @li image_size: A optional attribute of type ListInt, the size of the image. \n
255- 
256-*@par Outputs:
257-*@li nmsed_boxes: A 3D Tensor of type float16 or float32 with shape (batch, max_total_size, 4),
258-* specifying the output nms boxes per batch.
259-*@li nmsed_scores: A 2D Tensor of type float16 or float32 with shape (batch, max_total_size),
260-* specifying the output nms score per batch.
261-*@li nmsed_classes: A 2D Tensor of type float16 or float32 with shape (batch, max_total_size),
262-* specifying the output nms class per batch.
263-*@li nmsed_num: A 1D Tensor of type int32 with shape (batch), specifying the valid num of nmsed_boxes . \n
264- 
265-*@attention Constraints:
266-* Only computation of float16 or float32 data is supported.
267-* Note: when the class num per image * max_size_per_class is too big, will compile fail with ERROR-insufficient
268-memory
269-*/
270-REG_OP(BatchMultiClassNonMaxSuppression)
271- .INPUT(boxes, TensorType({DT_FLOAT16, DT_FLOAT}))
272- .INPUT(scores, TensorType({DT_FLOAT16, DT_FLOAT}))
273- .OPTIONAL_INPUT(clip_window, TensorType({DT_FLOAT16, DT_FLOAT}))
274- .OPTIONAL_INPUT(num_valid_boxes, TensorType({DT_INT32}))
275- .OUTPUT(nmsed_boxes, TensorType({DT_FLOAT16, DT_FLOAT}))
276- .OUTPUT(nmsed_scores, TensorType({DT_FLOAT16, DT_FLOAT}))
277- .OUTPUT(nmsed_classes, TensorType({DT_FLOAT16, DT_FLOAT}))
278- .OUTPUT(nmsed_num, TensorType({DT_INT32}))
279- .REQUIRED_ATTR(score_threshold, Float)
280- .REQUIRED_ATTR(iou_threshold, Float)
281- .REQUIRED_ATTR(max_size_per_class, Int)
282- .REQUIRED_ATTR(max_total_size, Int)
283- .ATTR(change_coordinate_frame, Bool, false)
284- .ATTR(transpose_box, Bool, false)
285- .ATTR(image_size, ListInt, {})
286- .OP_END_FACTORY_REG(BatchMultiClassNonMaxSuppression)
287- 
288/**228/**
289*@brief Generates bounding boxes based on "rois" and "deltas".229*@brief Generates bounding boxes based on "rois" and "deltas".
290* It is a customized FasterRcnn operator . \n230* It is a customized FasterRcnn operator . \n
Aobjdetect/batch_multi_class_non_max_suppression/CMakeLists.txt+25-0
@@ -0,0 +1,25 @@
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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST AND NOT BENCHMARK)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+ 
16+foreach(SUB_DIR ${CURRENT_DIRS})
17+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18+ add_subdirectory(${SUB_DIR})
19+ endif()
20+endforeach()
21+ 
22+set(SUPPORT_COMPUTE_UNIT "ascend950")
23+set(SUPPORT_TILING_DIR "arch35")
24+add_all_modules_sources(OPTYPE batch_multi_class_non_max_suppression ACLNNTYPE aclnn_exclude
25+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR})
Aobjdetect/batch_multi_class_non_max_suppression/README.md+159-0
@@ -0,0 +1,159 @@
1+# BatchMultiClassNonMaxSuppression
黄
黄黄晓彬16 天前

按照规范写文档

likedislike
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :--- | :---: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | × |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：对每个 batch、每个类别的候选框执行贪心非极大值抑制（NMS），再从所有类别的保留结果中按分数选择最多 `max_total_size` 个检测框。
17+- boxes 坐标格式为 `[y_min, x_min, y_max, x_max]`。当 `q=1` 时各类别共享 boxes；当 `q=C` 时每个类别使用自己的 boxes。
18+- 当提供 `clip_window` 时，算子先执行裁剪；`change_coordinate_frame=true` 时，再以窗口左上角为原点、窗口宽高为尺度进行归一化。
19+ 
20+对候选框 `a` 和已选框 `b`，IoU 的计算为：
21+ 
22+$$
23+IoU(a,b)=\frac{Area(a\cap b)}{max(Area(a)+Area(b)-Area(a\cap b), 1e-12)}
24+$$
25+ 
26+分数仅在 `score > score_threshold` 时保留；IoU 严格大于 `iou_threshold` 时抑制。每个类别最多保留 `max_size_per_class` 个框，最终结果不足 `max_total_size` 时以 0 填充。
27+ 
28+## 参数说明
29+ 
30+<table style="undefined;table-layout: fixed; width: 1005px"><colgroup>
31+ <col style="width: 170px">
32+ <col style="width: 170px">
33+ <col style="width: 352px">
34+ <col style="width: 213px">
35+ <col style="width: 100px">
36+ </colgroup>
37+ <thead>
38+ <tr>
39+ <th>参数名</th>
40+ <th>输入/输出/属性</th>
41+ <th>描述</th>
42+ <th>数据类型</th>
43+ <th>数据格式</th>
44+ </tr></thead>
45+ <tbody>
46+ <tr>
47+ <td>boxes</td>
48+ <td>输入</td>
49+ <td>候选框，坐标格式为[y_min, x_min, y_max, x_max]。transpose_box为false时shape为[B,N,q,4]，为true时shape为[B,q,4,N]；q为1或C。</td>
50+ <td>FLOAT16、FLOAT</td>
51+ <td>ND</td>
52+ </tr>
53+ <tr>
54+ <td>scores</td>
55+ <td>输入</td>
56+ <td>每个候选框、每个类别的分数，shape为[B,N,C]，数据类型必须与boxes一致。</td>
57+ <td>FLOAT16、FLOAT</td>
58+ <td>ND</td>
59+ </tr>
60+ <tr>
61+ <td>clip_window</td>
62+ <td>可选输入</td>
63+ <td>裁剪窗口，坐标格式为yxyx，shape为[B,4]。可传空指针，数据类型必须与boxes一致。</td>
64+ <td>FLOAT16、FLOAT</td>
65+ <td>ND</td>
66+ </tr>
67+ <tr>
68+ <td>num_valid_boxes</td>
69+ <td>可选输入</td>
70+ <td>每个batch的有效候选框数，shape为[B]。可传空指针。</td>
71+ <td>INT32</td>
72+ <td>ND</td>
73+ </tr>
74+ <tr>
75+ <td>score_threshold</td>
76+ <td>属性</td>
77+ <td>分数阈值，仅保留严格大于该值的候选框；必须为有限值。</td>
78+ <td>FLOAT</td>
79+ <td>-</td>
80+ </tr>
81+ <tr>
82+ <td>iou_threshold</td>
83+ <td>属性</td>
84+ <td>IoU抑制阈值；必须为[0,1]内的有限值。</td>
85+ <td>FLOAT</td>
86+ <td>-</td>
87+ </tr>
88+ <tr>
89+ <td>max_size_per_class</td>
90+ <td>属性</td>
91+ <td>每个类别最多保留的候选框数量，取值范围为[1,1000]。</td>
92+ <td>INT64</td>
93+ <td>-</td>
94+ </tr>
95+ <tr>
96+ <td>max_total_size</td>
97+ <td>属性</td>
98+ <td>每个batch最多保留的候选框数量，取值范围为[1,1000]。</td>
99+ <td>INT64</td>
100+ <td>-</td>
101+ </tr>
102+ <tr>
103+ <td>change_coordinate_frame</td>
104+ <td>属性</td>
105+ <td>是否将裁剪后的坐标归一化到窗口坐标系，默认为false；为true时必须提供clip_window。</td>
106+ <td>BOOL</td>
107+ <td>-</td>
108+ </tr>
109+ <tr>
110+ <td>transpose_box</td>
111+ <td>属性</td>
112+ <td>是否按[B,q,4,N]解释boxes，默认为false。</td>
113+ <td>BOOL</td>
114+ <td>-</td>
115+ </tr>
116+ <tr>
117+ <td>nmsed_boxes</td>
118+ <td>输出</td>
119+ <td>NMS后的检测框，shape为[B,M,4]，其中M为max_total_size。</td>
120+ <td>FLOAT16、FLOAT</td>
121+ <td>ND</td>
122+ </tr>
123+ <tr>
124+ <td>nmsed_scores</td>
125+ <td>输出</td>
126+ <td>NMS后的检测分数，shape为[B,M]。</td>
127+ <td>FLOAT16、FLOAT</td>
128+ <td>ND</td>
129+ </tr>
130+ <tr>
131+ <td>nmsed_classes</td>
132+ <td>输出</td>
133+ <td>NMS后的类别编号，shape为[B,M]，以浮点类型表示。</td>
134+ <td>FLOAT16、FLOAT</td>
135+ <td>ND</td>
136+ </tr>
137+ <tr>
138+ <td>nmsed_num</td>
139+ <td>输出</td>
140+ <td>每个batch的有效输出数量，shape为[B]。</td>
141+ <td>INT32</td>
142+ <td>ND</td>
143+ </tr>
144+ </tbody></table>
145+ 
146+## 约束说明
147+ 
148+- 仅支持 ND format、float16/float32 的 boxes 和 scores；三个浮点输出与 boxes dtype 保持一致。
149+- boxes 为 4 维、scores 为 3 维，B、N、C、q 必须为正；q 必须为 1 或 C。
150+- `clip_window` 的形状必须为 `[B,4]`，`num_valid_boxes` 的形状必须为 `[B]`。
151+- 当前 Ascend950 tiling 只接受具体的正 shape，不支持动态 rank 和未知维度。
152+- `image_size` 为图模式原型兼容属性，当前通用 NMS 路径不读取该属性；旧平台 `norm_class` 专用语义不在本实现范围内。
153+- 相同分数的相对顺序不是公共接口承诺；调用方应以 `nmsed_num` 确定有效输出范围。
154+ 
155+## 调用说明
156+ 
157+| 调用方式 | 样例代码 | 说明 |
158+| :--- | :--- | :--- |
159+| 图模式 | [test_geir_batch_multi_class_non_max_suppression](examples/arch35/test_geir_batch_multi_class_non_max_suppression.cpp) | 通过[算子IR](./op_graph/batch_multi_class_non_max_suppression_proto.h)构图方式调用BatchMultiClassNonMaxSuppression算子。 |
Aobjdetect/batch_multi_class_non_max_suppression/examples/arch35/test_geir_batch_multi_class_non_max_suppression.cpp+187-0
@@ -0,0 +1,187 @@
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 <cmath>
12+#include <cstddef>
13+#include <cstdint>
14+#include <iostream>
15+#include <map>
16+#include <string>
17+#include <vector>
18+ 
19+#include "ge_api.h"
20+#include "ge_api_types.h"
21+#include "ge_error_codes.h"
22+#include "ge_ir_build.h"
23+#include "graph.h"
24+#include "tensor.h"
25+#include "types.h"
26+ 
27+#include "../../op_graph/batch_multi_class_non_max_suppression_proto.h"
28+ 
29+#define FAILED -1
30+#define SUCCESS 0
31+ 
32+using namespace ge;
33+using std::map;
34+using std::string;
35+using std::vector;
36+ 
37+namespace {
38+constexpr float kTolerance = 1e-5F;
39+ 
40+TensorDesc MakeTensorDesc(const vector<int64_t>& shape, DataType dataType)
41+{
42+ TensorDesc desc(Shape(shape), FORMAT_ND, dataType);
43+ desc.SetPlacement(kPlacementHost);
44+ desc.SetFormat(FORMAT_ND);
45+ desc.SetRealDimCnt(shape.size());
46+ return desc;
47+}
48+ 
49+template <typename T>
50+Tensor MakeTensor(const vector<int64_t>& shape, vector<T>& data, DataType dataType)
51+{
52+ TensorDesc desc = MakeTensorDesc(shape, dataType);
53+ return Tensor(desc, reinterpret_cast<uint8_t*>(data.data()), data.size() * sizeof(T));
54+}
55+ 
56+template <typename T>
57+op::Data AddInput(const string& name, int32_t index, const vector<int64_t>& shape, vector<T>& data, DataType dataType,
58+ vector<Tensor>& input, Graph& graph)
59+{
60+ auto placeholder = op::Data(name.c_str()).set_attr_index(index);
61+ TensorDesc desc = MakeTensorDesc(shape, dataType);
62+ placeholder.update_input_desc_x(desc);
63+ placeholder.update_output_desc_y(desc);
64+ input.push_back(MakeTensor(shape, data, dataType));
65+ graph.AddOp(placeholder);
66+ return placeholder;
67+}
68+ 
69+bool VerifyOutput(const vector<Tensor>& output)
70+{
71+ const vector<int64_t> expectedElements = {16, 4, 4, 1};
72+ const vector<DataType> expectedDataTypes = {DT_FLOAT, DT_FLOAT, DT_FLOAT, DT_INT32};
73+ if (output.size() != expectedElements.size()) {
74+ std::cerr << "Unexpected output count: " << output.size() << std::endl;
75+ return false;
76+ }
77+ for (size_t index = 0; index < output.size(); ++index) {
78+ if (output[index].GetTensorDesc().GetDataType() != expectedDataTypes[index] ||
79+ output[index].GetTensorDesc().GetShape().GetShapeSize() != expectedElements[index] ||
80+ output[index].GetData() == nullptr) {
81+ std::cerr << "Output " << index << " has an unexpected descriptor" << std::endl;
82+ return false;
83+ }
84+ }
85+ 
86+ const vector<float> expectedBoxes = {0.00F, 0.00F, 0.50F, 0.50F, 0.05F, 0.05F, 0.55F, 0.55F,
87+ 0.55F, 0.55F, 1.00F, 1.00F, 0.00F, 0.00F, 0.00F, 0.00F};
88+ const vector<float> expectedScores = {0.95F, 0.85F, 0.80F, 0.00F};
89+ const vector<float> expectedClasses = {0.00F, 1.00F, 0.00F, 0.00F};
90+ const auto* boxes = reinterpret_cast<const float*>(output[0].GetData());
91+ const auto* scores = reinterpret_cast<const float*>(output[1].GetData());
92+ const auto* classes = reinterpret_cast<const float*>(output[2].GetData());
93+ const auto* num = reinterpret_cast<const int32_t*>(output[3].GetData());
94+ for (size_t index = 0; index < expectedBoxes.size(); ++index) {
95+ if (std::fabs(boxes[index] - expectedBoxes[index]) > kTolerance) {
96+ std::cerr << "nmsed_boxes mismatch at " << index << std::endl;
97+ return false;
98+ }
99+ }
100+ for (size_t index = 0; index < expectedScores.size(); ++index) {
101+ if (std::fabs(scores[index] - expectedScores[index]) > kTolerance ||
102+ std::fabs(classes[index] - expectedClasses[index]) > kTolerance) {
103+ std::cerr << "nmsed_scores or nmsed_classes mismatch at " << index << std::endl;
104+ return false;
105+ }
106+ }
107+ if (num[0] != 3) {
108+ std::cerr << "nmsed_num mismatch: " << num[0] << std::endl;
109+ return false;
110+ }
111+ return true;
112+}
113+} // namespace
114+ 
115+int CreateOppInGraph(vector<Tensor>& input, vector<Operator>& inputs, vector<Operator>& outputs, Graph& graph)
116+{
117+ // boxes: [B, N, q, 4], scores: [B, N, C]. q=1 means boxes are shared by all classes.
118+ static vector<float> boxesData = {0.00F, 0.00F, 0.50F, 0.50F, 0.05F, 0.05F, 0.55F, 0.55F,
119+ 0.55F, 0.55F, 1.00F, 1.00F, 0.10F, 0.55F, 0.45F, 1.00F};
120+ static vector<float> scoresData = {0.95F, 0.20F, 0.90F, 0.85F, 0.80F, 0.10F, 0.99F, 0.98F};
121+ static vector<float> clipWindowData = {0.00F, 0.00F, 1.00F, 1.00F};
122+ static vector<int32_t> numValidBoxesData = {3};
123+ 
124+ auto boxes = AddInput("boxes", 0, {1, 4, 1, 4}, boxesData, DT_FLOAT, input, graph);
125+ auto scores = AddInput("scores", 1, {1, 4, 2}, scoresData, DT_FLOAT, input, graph);
126+ auto clipWindow = AddInput("clip_window", 2, {1, 4}, clipWindowData, DT_FLOAT, input, graph);
127+ auto numValidBoxes = AddInput("num_valid_boxes", 3, {1}, numValidBoxesData, DT_INT32, input, graph);
128+ 
129+ auto nms = op::BatchMultiClassNonMaxSuppression("batch_multi_class_non_max_suppression");
130+ nms.set_input_boxes(boxes);
131+ nms.set_input_scores(scores);
132+ nms.set_input_clip_window(clipWindow);
133+ nms.set_input_num_valid_boxes(numValidBoxes);
134+ nms.set_attr_score_threshold(0.20F);
135+ nms.set_attr_iou_threshold(0.50F);
136+ nms.set_attr_max_size_per_class(2);
137+ nms.set_attr_max_total_size(4);
138+ nms.set_attr_change_coordinate_frame(false);
139+ nms.set_attr_transpose_box(false);
140+ nms.update_output_desc_nmsed_boxes(MakeTensorDesc({1, 4, 4}, DT_FLOAT));
141+ nms.update_output_desc_nmsed_scores(MakeTensorDesc({1, 4}, DT_FLOAT));
142+ nms.update_output_desc_nmsed_classes(MakeTensorDesc({1, 4}, DT_FLOAT));
143+ nms.update_output_desc_nmsed_num(MakeTensorDesc({1}, DT_INT32));
144+ 
145+ inputs = {boxes, scores, clipWindow, numValidBoxes};
146+ outputs = {nms};
147+ return SUCCESS;
148+}
149+ 
150+int main(int argc, char* argv[])
151+{
152+ (void)argc;
153+ (void)argv;
154+ map<AscendString, AscendString> globalOptions = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
155+ if (GEInitialize(globalOptions) != SUCCESS) {
156+ std::cerr << "GEInitialize failed" << std::endl;
157+ return FAILED;
158+ }
159+ 
160+ Graph graph("batch_multi_class_non_max_suppression_graph");
161+ vector<Tensor> input;
162+ vector<Operator> inputs;
163+ vector<Operator> outputs;
164+ int ret = CreateOppInGraph(input, inputs, outputs, graph);
165+ if (ret == SUCCESS) {
166+ graph.SetInputs(inputs).SetOutputs(outputs);
167+ map<AscendString, AscendString> buildOptions = {};
168+ map<AscendString, AscendString> graphOptions = {};
169+ Session session(buildOptions);
170+ ret = session.AddGraph(0, graph, graphOptions);
171+ vector<Tensor> output;
172+ if (ret == SUCCESS) {
173+ ret = session.RunGraph(0, input, output);
174+ }
175+ if (ret == SUCCESS && !VerifyOutput(output)) {
176+ ret = FAILED;
177+ }
178+ }
179+ 
180+ const Status finalizeRet = GEFinalize();
181+ if (ret != SUCCESS || finalizeRet != SUCCESS) {
182+ std::cerr << "BatchMultiClassNonMaxSuppression GE IR example failed" << std::endl;
183+ return FAILED;
184+ }
185+ std::cout << "BatchMultiClassNonMaxSuppression GE IR example passed" << std::endl;
186+ return SUCCESS;
187+}
Aobjdetect/batch_multi_class_non_max_suppression/op_graph/batch_multi_class_non_max_suppression_proto.h+66-0
@@ -0,0 +1,66 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
黄
黄黄晓彬16 天前

怎么是一个空文件

likedislike
H
Huang-Peng
16 天前 评论：
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 batch_multi_class_non_max_suppression_proto.h
13+ * \brief Operator prototype definition for BatchMultiClassNonMaxSuppression.
14+ */
15+#ifndef OPS_OP_PROTO_INC_BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_OPS_H_
16+#define OPS_OP_PROTO_INC_BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_OPS_H_
17+ 
18+#include "graph/operator_reg.h"
19+ 
20+namespace ge {
21+/**
22+ * @brief Selects high-scoring bounding boxes independently for each class and suppresses boxes with excessive IoU.
23+ *
24+ * @par Inputs:
25+ * @li boxes: A float16 or float32 tensor. When transpose_box is false, its shape is [B, N, q, 4]; when true, its
26+ * shape is [B, q, 4, N]. B is the batch size, N is the number of boxes, and q is 1 or the number of classes.
27+ * @li scores: A float16 or float32 tensor of shape [B, N, C], where C is the number of classes. Its type must equal
28+ * the type of boxes.
29+ * @li clip_window: An optional float16 or float32 tensor of shape [B, 4]. Its type must equal the type of boxes.
30+ * @li num_valid_boxes: An optional int32 tensor of shape [B].
31+ *
32+ * @par Attributes:
33+ * @li score_threshold: A required float, the score filtering threshold.
34+ * @li iou_threshold: A required float in [0, 1], the IoU suppression threshold.
35+ * @li max_size_per_class: A required int, the maximum selected boxes per class.
36+ * @li max_total_size: A required int, the maximum selected boxes per batch.
37+ * @li change_coordinate_frame: An optional bool. When true, clip_window must be provided. Defaults to false.
38+ * @li transpose_box: An optional bool indicating the layout of boxes. Defaults to false.
39+ * @li image_size: An optional list of int values reserved for graph compatibility.
40+ *
41+ * @par Outputs:
42+ * @li nmsed_boxes: A float16 or float32 tensor of shape [B, max_total_size, 4].
43+ * @li nmsed_scores: A float16 or float32 tensor of shape [B, max_total_size].
44+ * @li nmsed_classes: A float16 or float32 tensor of shape [B, max_total_size].
45+ * @li nmsed_num: An int32 tensor of shape [B].
46+ */
47+REG_OP(BatchMultiClassNonMaxSuppression)
48+ .INPUT(boxes, TensorType({DT_FLOAT16, DT_FLOAT}))
49+ .INPUT(scores, TensorType({DT_FLOAT16, DT_FLOAT}))
50+ .OPTIONAL_INPUT(clip_window, TensorType({DT_FLOAT16, DT_FLOAT}))
51+ .OPTIONAL_INPUT(num_valid_boxes, TensorType({DT_INT32}))
52+ .OUTPUT(nmsed_boxes, TensorType({DT_FLOAT16, DT_FLOAT}))
53+ .OUTPUT(nmsed_scores, TensorType({DT_FLOAT16, DT_FLOAT}))
54+ .OUTPUT(nmsed_classes, TensorType({DT_FLOAT16, DT_FLOAT}))
55+ .OUTPUT(nmsed_num, TensorType({DT_INT32}))
56+ .REQUIRED_ATTR(score_threshold, Float)
57+ .REQUIRED_ATTR(iou_threshold, Float)
58+ .REQUIRED_ATTR(max_size_per_class, Int)
59+ .REQUIRED_ATTR(max_total_size, Int)
60+ .ATTR(change_coordinate_frame, Bool, false)
61+ .ATTR(transpose_box, Bool, false)
62+ .ATTR(image_size, ListInt, {})
63+ .OP_END_FACTORY_REG(BatchMultiClassNonMaxSuppression)
64+} // namespace ge
65+ 
66+#endif // OPS_OP_PROTO_INC_BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_OPS_H_
Aobjdetect/batch_multi_class_non_max_suppression/op_host/arch35/batch_multi_class_non_max_suppression_tiling_arch35.cpp+496-0
@@ -0,0 +1,496 @@
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 "batch_multi_class_non_max_suppression_tiling_arch35.h"
12+#include "objdetect/batch_multi_class_non_max_suppression/op_kernel/arch35/batch_multi_class_non_max_suppression_tiling_key.h"
13+ 
14+#include <algorithm>
15+#include <cmath>
16+#include <cstdint>
17+#include <limits>
18+#include <securec.h>
19+ 
20+#include "graph/utils/type_utils.h"
21+#include "log/log.h"
22+#include "register/op_impl_registry.h"
23+#include "tiling/platform/platform_ascendc.h"
24+ 
25+namespace {
26+constexpr int64_t kBoxesIndex = 0;
27+constexpr int64_t kScoresIndex = 1;
28+constexpr int64_t kClipWindowIndex = 2;
29+constexpr int64_t kNumValidBoxesIndex = 3;
30+constexpr int64_t kNmsedBoxesIndex = 0;
31+constexpr int64_t kNmsedScoresIndex = 1;
32+constexpr int64_t kNmsedClassesIndex = 2;
33+constexpr int64_t kNmsedNumIndex = 3;
34+constexpr int64_t kScoreThresholdAttrIndex = 0;
35+constexpr int64_t kIouThresholdAttrIndex = 1;
36+constexpr int64_t kMaxSizePerClassAttrIndex = 2;
37+constexpr int64_t kMaxTotalSizeAttrIndex = 3;
38+constexpr int64_t kChangeCoordinateFrameAttrIndex = 4;
39+constexpr int64_t kTransposeBoxAttrIndex = 5;
40+constexpr int64_t kMaxOutputSize = 1000;
41+constexpr uint32_t kScheduleMode = 1;
42+constexpr int64_t kScratchFloatFieldCount = 5;
43+constexpr int64_t kVectorFloatFieldCount = 10;
44+constexpr int64_t kTileAlignment = 64;
45+constexpr int64_t kMaxTileSize = 4096;
46+// Leave 32 KiB of the MIX_AIV UB budget outside AIV tile and TopK planning
47+// for the SIMT side of this kernel.
48+constexpr uint64_t kSimtUbReserveBytes = 32 * 1024;
49+constexpr uint64_t kWorkspaceAlignment = 32;
50+constexpr int64_t kTopKAlignment = 32;
51+// Keep the cross-class TopK input comfortably below the Ascend950 UB budget.
52+// Larger candidate lists are merged incrementally in the kernel.
53+constexpr int64_t kMergeTopKInputCapacity = 2048;
54+ 
55+bool AddAligned(uint64_t offset, uint64_t size, uint64_t& nextOffset)
56+{
57+ if (offset > std::numeric_limits<uint64_t>::max() - size) {
58+ return false;
59+ }
60+ const uint64_t endOffset = offset + size;
61+ if (endOffset > std::numeric_limits<uint64_t>::max() - (kWorkspaceAlignment - 1)) {
62+ return false;
63+ }
64+ nextOffset = (endOffset + kWorkspaceAlignment - 1) / kWorkspaceAlignment * kWorkspaceAlignment;
65+ return true;
66+}
67+ 
68+bool ProductFitsUint32(uint64_t dim0, uint64_t dim1, uint64_t dim2, uint64_t dim3)
69+{
70+ constexpr uint64_t kUint32Max = std::numeric_limits<uint32_t>::max();
71+ uint64_t product = dim0;
72+ if (product > kUint32Max / dim1) {
73+ return false;
74+ }
75+ product *= dim1;
76+ if (product > kUint32Max / dim2) {
77+ return false;
78+ }
79+ product *= dim2;
80+ return product <= kUint32Max / dim3;
81+}
82+ 
83+bool HasPositiveDims(const gert::Shape& shape)
84+{
85+ for (size_t i = 0; i < shape.GetDimNum(); ++i) {
86+ if (shape.GetDim(i) <= 0) {
87+ return false;
88+ }
89+ }
90+ return true;
91+}
92+ 
93+struct ParsedTilingParams {
94+ int64_t batch{0};
95+ int64_t boxesNum{0};
96+ int64_t classesNum{0};
97+ int64_t boxClassesNum{0};
98+ int64_t maxSizePerClass{0};
99+ int64_t maxTotalSize{0};
100+ float scoreThreshold{0.0F};
101+ float iouThreshold{0.0F};
102+ ge::DataType boxesType{ge::DT_UNDEFINED};
103+ bool changeCoordinateFrame{false};
104+ bool transposeBox{false};
105+ bool hasClipWindow{false};
106+ bool hasNumValidBoxes{false};
107+};
108+ 
109+ge::graphStatus ParseRequiredInputsAndAttrs(gert::TilingContext* context, ParsedTilingParams& params)
110+{
111+ const gert::StorageShape* boxesInput = context->GetInputShape(kBoxesIndex);
112+ const gert::StorageShape* scoresInput = context->GetInputShape(kScoresIndex);
113+ OP_CHECK_NULL_WITH_CONTEXT(context, boxesInput);
114+ OP_CHECK_NULL_WITH_CONTEXT(context, scoresInput);
115+ const gert::Shape& boxesShape = boxesInput->GetStorageShape();
116+ const gert::Shape& scoresShape = scoresInput->GetStorageShape();
117+ OP_CHECK_IF(boxesShape.GetDimNum() != 4 || scoresShape.GetDimNum() != 3,
118+ OP_LOGE(context, "boxes must be rank 4 and scores must be rank 3."), return ge::GRAPH_FAILED);
119+ OP_CHECK_IF(!HasPositiveDims(boxesShape) || !HasPositiveDims(scoresShape),
120+ OP_LOGE(context, "Dynamic or non-positive dimensions are not supported by this tiling."),
121+ return ge::GRAPH_FAILED);
122+ 
123+ const gert::RuntimeAttrs* attrs = context->GetAttrs();
124+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
125+ const float* scoreThreshold = attrs->GetAttrPointer<float>(kScoreThresholdAttrIndex);
126+ const float* iouThreshold = attrs->GetAttrPointer<float>(kIouThresholdAttrIndex);
127+ const int64_t* maxSizePerClass = attrs->GetAttrPointer<int64_t>(kMaxSizePerClassAttrIndex);
128+ const int64_t* maxTotalSize = attrs->GetAttrPointer<int64_t>(kMaxTotalSizeAttrIndex);
129+ const bool* changeCoordinateFrame = attrs->GetAttrPointer<bool>(kChangeCoordinateFrameAttrIndex);
130+ const bool* transposeBox = attrs->GetAttrPointer<bool>(kTransposeBoxAttrIndex);
131+ OP_CHECK_NULL_WITH_CONTEXT(context, scoreThreshold);
132+ OP_CHECK_NULL_WITH_CONTEXT(context, iouThreshold);
133+ OP_CHECK_NULL_WITH_CONTEXT(context, maxSizePerClass);
134+ OP_CHECK_NULL_WITH_CONTEXT(context, maxTotalSize);
135+ OP_CHECK_NULL_WITH_CONTEXT(context, changeCoordinateFrame);
136+ OP_CHECK_NULL_WITH_CONTEXT(context, transposeBox);
137+ OP_CHECK_IF(!std::isfinite(*scoreThreshold) || !std::isfinite(*iouThreshold) || *iouThreshold < 0.0F ||
138+ *iouThreshold > 1.0F,
139+ OP_LOGE(context, "score_threshold must be finite and iou_threshold must be finite in [0, 1]."),
140+ return ge::GRAPH_FAILED);
141+ OP_CHECK_IF(*maxSizePerClass <= 0 || *maxSizePerClass > kMaxOutputSize || *maxTotalSize <= 0 ||
142+ *maxTotalSize > kMaxOutputSize,
143+ OP_LOGE(context, "max_size_per_class and max_total_size must be in [1, %ld].", kMaxOutputSize),
144+ return ge::GRAPH_FAILED);
atomgit-bot
atomgit-botatomgit-bot25 天前

🟡 Medium Priority

变更位置：CheckAndParse() 第 69 行获取 scoreThreshold，第 81 行只对 iouThreshold 做了 [0, 1] 范围校验，但 scoreThreshold 完全没有校验。

kernel 中（batch_multi_class_non_max_suppression_kernel.h 第 241 行）使用 !(score >= scoreThreshold) 做过滤。如果用户传入 scoreThreshold = NaN，根据 IEEE 754 规则 score >= NaN 对任意有限 score 恒为 false，!(false) = true 会使全部候选框通过过滤，导致输出结果异常膨胀且用户毫无感知。类似地，scoreThreshold = -Inf 也会让所有框通过。该属性为 REQUIRED，应在 tiling 阶段拦截非法值。

likedislike
145+ 
146+ params.batch = boxesShape.GetDim(0);
147+ params.transposeBox = *transposeBox;
148+ if (params.transposeBox) {
149+ params.boxClassesNum = boxesShape.GetDim(1);
150+ OP_CHECK_IF(boxesShape.GetDim(2) != 4, OP_LOGE(context, "boxes dim 2 must be 4 when transpose_box is true."),
151+ return ge::GRAPH_FAILED);
152+ params.boxesNum = boxesShape.GetDim(3);
153+ } else {
154+ params.boxesNum = boxesShape.GetDim(1);
155+ params.boxClassesNum = boxesShape.GetDim(2);
156+ OP_CHECK_IF(boxesShape.GetDim(3) != 4, OP_LOGE(context, "boxes last dimension must be 4."),
157+ return ge::GRAPH_FAILED);
158+ }
159+ params.classesNum = scoresShape.GetDim(2);
160+ OP_CHECK_IF(scoresShape.GetDim(0) != params.batch || scoresShape.GetDim(1) != params.boxesNum ||
161+ (params.boxClassesNum != 1 && params.boxClassesNum != params.classesNum),
162+ OP_LOGE(context, "boxes and scores shapes are incompatible."), return ge::GRAPH_FAILED);
163+ OP_CHECK_IF(params.batch > std::numeric_limits<int64_t>::max() / params.classesNum,
164+ OP_LOGE(context, "batch-by-class task count overflow."), return ge::GRAPH_FAILED);
165+ 
166+ const auto* boxesDesc = context->GetInputDesc(kBoxesIndex);
167+ const auto* scoresDesc = context->GetInputDesc(kScoresIndex);
168+ OP_CHECK_NULL_WITH_CONTEXT(context, boxesDesc);
169+ OP_CHECK_NULL_WITH_CONTEXT(context, scoresDesc);
170+ params.boxesType = boxesDesc->GetDataType();
171+ OP_CHECK_IF((params.boxesType != ge::DT_FLOAT16 && params.boxesType != ge::DT_FLOAT) ||
172+ scoresDesc->GetDataType() != params.boxesType,
173+ OP_LOGE(context, "boxes and scores must have the same dtype in {float16, float}."),
174+ return ge::GRAPH_FAILED);
175+ 
176+ params.scoreThreshold = *scoreThreshold;
177+ params.iouThreshold = *iouThreshold;
178+ params.maxSizePerClass = *maxSizePerClass;
179+ params.maxTotalSize = *maxTotalSize;
180+ params.changeCoordinateFrame = *changeCoordinateFrame;
181+ return ge::GRAPH_SUCCESS;
182+}
183+ 
184+ge::graphStatus ValidateOptionalInputs(gert::TilingContext* context, ParsedTilingParams& params)
185+{
186+ const gert::StorageShape* clipWindowInput = context->GetInputShape(kClipWindowIndex);
187+ const gert::StorageShape* numValidBoxesInput = context->GetInputShape(kNumValidBoxesIndex);
188+ if (clipWindowInput != nullptr) {
189+ const auto* clipWindowDesc = context->GetInputDesc(kClipWindowIndex);
190+ OP_CHECK_NULL_WITH_CONTEXT(context, clipWindowDesc);
191+ // The tiling context compacts a later optional tensor into slot 2 when
192+ // clip_window is absent. Distinguish that valid [B]/int32 NVB input
193+ // from clip_window here; the kernel ABI itself retains the original
194+ // fourth argument position for num_valid_boxes.
195+ if (clipWindowDesc->GetDataType() == ge::DT_INT32) {
196+ const gert::Shape& numValidBoxesShape = clipWindowInput->GetStorageShape();
197+ OP_CHECK_IF(numValidBoxesShape.GetDimNum() != 1 || numValidBoxesShape.GetDim(0) != params.batch,
198+ OP_LOGE(context, "num_valid_boxes must have shape [B]."), return ge::GRAPH_FAILED);
199+ OP_CHECK_IF(numValidBoxesInput != nullptr,
200+ OP_LOGE(context, "num_valid_boxes is present in both optional input slots."),
201+ return ge::GRAPH_FAILED);
202+ params.hasNumValidBoxes = true;
203+ } else {
204+ params.hasClipWindow = true;
205+ const gert::Shape& clipWindowShape = clipWindowInput->GetStorageShape();
206+ OP_CHECK_IF(clipWindowShape.GetDimNum() != 2 || clipWindowShape.GetDim(0) != params.batch ||
207+ clipWindowShape.GetDim(1) != 4,
208+ OP_LOGE(context, "clip_window must have shape [B, 4]."), return ge::GRAPH_FAILED);
209+ OP_CHECK_IF(clipWindowDesc->GetDataType() != params.boxesType,
210+ OP_LOGE(context, "clip_window dtype must equal boxes dtype."), return ge::GRAPH_FAILED);
211+ }
212+ }
213+ if (numValidBoxesInput != nullptr) {
214+ const gert::Shape& numValidBoxesShape = numValidBoxesInput->GetStorageShape();
215+ OP_CHECK_IF(numValidBoxesShape.GetDimNum() != 1 || numValidBoxesShape.GetDim(0) != params.batch,
216+ OP_LOGE(context, "num_valid_boxes must have shape [B]."), return ge::GRAPH_FAILED);
217+ const auto* numValidBoxesDesc = context->GetInputDesc(kNumValidBoxesIndex);
218+ OP_CHECK_NULL_WITH_CONTEXT(context, numValidBoxesDesc);
219+ OP_CHECK_IF(numValidBoxesDesc->GetDataType() != ge::DT_INT32,
220+ OP_LOGE(context, "num_valid_boxes must be int32."), return ge::GRAPH_FAILED);
221+ params.hasNumValidBoxes = true;
222+ }
223+ OP_CHECK_IF(params.changeCoordinateFrame && !params.hasClipWindow,
224+ OP_LOGE(context, "change_coordinate_frame requires clip_window."), return ge::GRAPH_FAILED);
225+ return ge::GRAPH_SUCCESS;
226+}
227+ 
228+ge::graphStatus ValidateOutputContract(gert::TilingContext* context, const ParsedTilingParams& params)
229+{
230+ const gert::StorageShape* outputBoxes = context->GetOutputShape(kNmsedBoxesIndex);
231+ const gert::StorageShape* outputScores = context->GetOutputShape(kNmsedScoresIndex);
232+ const gert::StorageShape* outputClasses = context->GetOutputShape(kNmsedClassesIndex);
233+ const gert::StorageShape* outputNum = context->GetOutputShape(kNmsedNumIndex);
234+ OP_CHECK_NULL_WITH_CONTEXT(context, outputBoxes);
235+ OP_CHECK_NULL_WITH_CONTEXT(context, outputScores);
236+ OP_CHECK_NULL_WITH_CONTEXT(context, outputClasses);
237+ OP_CHECK_NULL_WITH_CONTEXT(context, outputNum);
238+ const gert::Shape& outputBoxesShape = outputBoxes->GetStorageShape();
239+ const gert::Shape& outputScoresShape = outputScores->GetStorageShape();
240+ const gert::Shape& outputClassesShape = outputClasses->GetStorageShape();
241+ const gert::Shape& outputNumShape = outputNum->GetStorageShape();
242+ OP_CHECK_IF(outputBoxesShape.GetDimNum() != 3 || outputBoxesShape.GetDim(0) != params.batch ||
243+ outputBoxesShape.GetDim(1) != params.maxTotalSize || outputBoxesShape.GetDim(2) != 4 ||
244+ outputScoresShape.GetDimNum() != 2 || outputScoresShape.GetDim(0) != params.batch ||
245+ outputScoresShape.GetDim(1) != params.maxTotalSize || outputClassesShape.GetDimNum() != 2 ||
246+ outputClassesShape.GetDim(0) != params.batch ||
247+ outputClassesShape.GetDim(1) != params.maxTotalSize || outputNumShape.GetDimNum() != 1 ||
248+ outputNumShape.GetDim(0) != params.batch,
249+ OP_LOGE(context, "Output shapes do not match BatchMultiClassNonMaxSuppression contract."),
250+ return ge::GRAPH_FAILED);
251+ const auto* outputBoxesDesc = context->GetOutputDesc(kNmsedBoxesIndex);
252+ const auto* outputScoresDesc = context->GetOutputDesc(kNmsedScoresIndex);
253+ const auto* outputClassesDesc = context->GetOutputDesc(kNmsedClassesIndex);
254+ const auto* outputNumDesc = context->GetOutputDesc(kNmsedNumIndex);
255+ OP_CHECK_NULL_WITH_CONTEXT(context, outputBoxesDesc);
256+ OP_CHECK_NULL_WITH_CONTEXT(context, outputScoresDesc);
257+ OP_CHECK_NULL_WITH_CONTEXT(context, outputClassesDesc);
258+ OP_CHECK_NULL_WITH_CONTEXT(context, outputNumDesc);
259+ OP_CHECK_IF(
260+ outputBoxesDesc->GetDataType() != params.boxesType || outputScoresDesc->GetDataType() != params.boxesType ||
261+ outputClassesDesc->GetDataType() != params.boxesType || outputNumDesc->GetDataType() != ge::DT_INT32,
262+ OP_LOGE(context, "Output dtypes do not match BatchMultiClassNonMaxSuppression contract."),
263+ return ge::GRAPH_FAILED);
264+ return ge::GRAPH_SUCCESS;
265+}
266+} // namespace
267+ 
268+namespace optiling {
269+ge::graphStatus BatchMultiClassNonMaxSuppressionTiling::CheckAndParse()
270+{
271+ ParsedTilingParams params;
272+ if (ParseRequiredInputsAndAttrs(context_, params) != ge::GRAPH_SUCCESS) {
273+ return ge::GRAPH_FAILED;
274+ }
275+ if (ValidateOptionalInputs(context_, params) != ge::GRAPH_SUCCESS) {
276+ return ge::GRAPH_FAILED;
277+ }
278+ if (ValidateOutputContract(context_, params) != ge::GRAPH_SUCCESS) {
279+ return ge::GRAPH_FAILED;
280+ }
281+ 
282+ batch_ = params.batch;
283+ boxesNum_ = params.boxesNum;
284+ classesNum_ = params.classesNum;
285+ boxClassesNum_ = params.boxClassesNum;
286+ maxSizePerClass_ = params.maxSizePerClass;
287+ maxTotalSize_ = params.maxTotalSize;
288+ scoreThreshold_ = params.scoreThreshold;
289+ iouThreshold_ = params.iouThreshold;
290+ hasClipWindow_ = params.hasClipWindow;
291+ hasNumValidBoxes_ = params.hasNumValidBoxes;
292+ changeCoordinateFrame_ = params.changeCoordinateFrame;
293+ transposeBox_ = params.transposeBox;
294+ return ge::GRAPH_SUCCESS;
295+}
296+ 
297+ge::graphStatus BatchMultiClassNonMaxSuppressionTiling::SetTilingData()
298+{
299+ const auto* compileInfo = static_cast<const BatchMultiClassNonMaxSuppressionCompileInfo*>(
300+ context_->GetCompileInfo());
301+ OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo);
302+ OP_CHECK_IF(compileInfo->coreNum == 0 || compileInfo->ubSize == 0,
303+ OP_LOGE(context_, "Ascend950 platform resource query failed."), return ge::GRAPH_FAILED);
304+ 
305+ tilingData_ = context_->GetTilingData<BatchMultiClassNonMaxSuppressionTilingData>();
306+ OP_CHECK_NULL_WITH_CONTEXT(context_, tilingData_);
307+ OP_CHECK_IF(memset_s(tilingData_, sizeof(BatchMultiClassNonMaxSuppressionTilingData), 0,
308+ sizeof(BatchMultiClassNonMaxSuppressionTilingData)) != EOK,
309+ OP_LOGE(context_, "Failed to initialize tiling data."), return ge::GRAPH_FAILED);
310+ const int64_t taskCount = batch_ * classesNum_;
311+ OP_CHECK_IF(taskCount <= 0, OP_LOGE(context_, "Invalid batch-by-class task count."), return ge::GRAPH_FAILED);
312+ const int64_t usedCoreNum = std::min<int64_t>(taskCount, static_cast<int64_t>(compileInfo->coreNum));
313+ 
314+ OP_CHECK_IF(classesNum_ > std::numeric_limits<int64_t>::max() / maxSizePerClass_,
315+ OP_LOGE(context_, "Cross-class TopK input size overflow."), return ge::GRAPH_FAILED);
316+ const int64_t mergeInputCount = classesNum_ * maxSizePerClass_;
317+ OP_CHECK_IF(mergeInputCount <= 0 || mergeInputCount > std::numeric_limits<int32_t>::max() - kTopKAlignment,
318+ OP_LOGE(context_, "Cross-class TopK input size is not supported."), return ge::GRAPH_FAILED);
319+ const int64_t mergeInputSize = std::min<int64_t>(
320+ kMergeTopKInputCapacity, (mergeInputCount + kTopKAlignment - 1) / kTopKAlignment * kTopKAlignment);
321+ const int64_t mergeOutputCount = std::min<int64_t>(maxTotalSize_, mergeInputCount);
322+ const int64_t mergeOutputSize = (maxTotalSize_ + kTopKAlignment - 1) / kTopKAlignment * kTopKAlignment;
323+ OP_CHECK_IF(mergeInputSize < mergeOutputSize,
324+ OP_LOGE(context_, "Cross-class TopK merge input capacity is smaller than its output size."),
325+ return ge::GRAPH_FAILED);
326+ 
327+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context_->GetPlatformInfo());
328+ AscendC::tiling::TopkTiling mergeTopKTiling{};
329+ OP_CHECK_IF(!AscendC::TopKTilingFunc(ascendcPlatform, static_cast<uint32_t>(mergeInputSize), 1,
330+ static_cast<uint32_t>(mergeOutputCount), sizeof(float), true,
331+ AscendC::TopKMode::TOPK_NORMAL, true, mergeTopKTiling),
332+ OP_LOGE(context_, "Ascend950 TopK tiling failed for cross-class merge."), return ge::GRAPH_FAILED);
333+ uint32_t topKMaxTmpBytes = 0;
334+ uint32_t topKMinTmpBytes = 0;
335+ OP_CHECK_IF(!AscendC::GetTopKMaxMinTmpSize(ascendcPlatform, static_cast<uint32_t>(mergeInputSize), 1, false, true,
336+ AscendC::TopKMode::TOPK_NORMAL, true, sizeof(float), topKMaxTmpBytes,
337+ topKMinTmpBytes),
338+ OP_LOGE(context_, "Ascend950 TopK workspace query failed for cross-class merge."),
339+ return ge::GRAPH_FAILED);
340+ const uint64_t topKTempBytes = std::max<uint64_t>(topKMaxTmpBytes, kWorkspaceAlignment);
341+ const uint64_t mergeFixedUbBytes = static_cast<uint64_t>(mergeInputSize) * (sizeof(float) + sizeof(int32_t)) +
342+ static_cast<uint64_t>(mergeOutputSize) * (sizeof(float) + sizeof(int32_t)) +
343+ topKTempBytes;
344+ OP_CHECK_IF(
345+ compileInfo->ubSize <= kSimtUbReserveBytes || mergeFixedUbBytes >= compileInfo->ubSize - kSimtUbReserveBytes,
346+ OP_LOGE(context_, "Insufficient UB for cross-class TopK merge."), return ge::GRAPH_FAILED);
347+ const uint64_t availableUbBytes = compileInfo->ubSize - kSimtUbReserveBytes - mergeFixedUbBytes;
348+ const uint64_t bytesPerTileElement = kVectorFloatFieldCount * sizeof(float) + sizeof(uint8_t);
349+ int64_t tileSize = std::min<int64_t>(boxesNum_,
350+ std::min<int64_t>(kMaxTileSize, availableUbBytes / bytesPerTileElement));
351+ if (tileSize < mergeOutputSize) {
352+ const uint64_t minimumReduceBytes = static_cast<uint64_t>(mergeOutputSize) * sizeof(float);
353+ OP_CHECK_IF(minimumReduceBytes > availableUbBytes,
354+ OP_LOGE(context_, "Insufficient UB for cross-class TopK reduction."), return ge::GRAPH_FAILED);
355+ tileSize = std::min<int64_t>(tileSize, static_cast<int64_t>((availableUbBytes - minimumReduceBytes) /
356+ (bytesPerTileElement - sizeof(float))));
357+ }
358+ if (tileSize >= kTileAlignment) {
359+ tileSize = tileSize / kTileAlignment * kTileAlignment;
360+ }
361+ OP_CHECK_IF(tileSize <= 0, OP_LOGE(context_, "Insufficient UB for BatchMultiClassNonMaxSuppression tile."),
362+ return ge::GRAPH_FAILED);
363+ const int64_t reduceBufferSize = std::max<int64_t>(tileSize, mergeOutputSize);
364+ 
365+ const uint64_t boxesNum = static_cast<uint64_t>(boxesNum_);
366+ const uint64_t taskCountU64 = static_cast<uint64_t>(taskCount);
367+ const uint64_t maxSizePerClass = static_cast<uint64_t>(maxSizePerClass_);
368+ OP_CHECK_IF(boxesNum > std::numeric_limits<uint64_t>::max() / sizeof(float),
369+ OP_LOGE(context_, "Scratch workspace size overflow."), return ge::GRAPH_FAILED);
370+ const uint64_t scratchFieldBytes = boxesNum * sizeof(float);
371+ OP_CHECK_IF(scratchFieldBytes > std::numeric_limits<uint64_t>::max() - (kWorkspaceAlignment - 1),
372+ OP_LOGE(context_, "Scratch workspace size overflow."), return ge::GRAPH_FAILED);
373+ const uint64_t scratchFieldStride = (scratchFieldBytes + kWorkspaceAlignment - 1) / kWorkspaceAlignment *
374+ kWorkspaceAlignment;
375+ OP_CHECK_IF(scratchFieldStride > std::numeric_limits<uint64_t>::max() / kScratchFloatFieldCount,
376+ OP_LOGE(context_, "Scratch workspace size overflow."), return ge::GRAPH_FAILED);
377+ const uint64_t scratchBytesPerCore = scratchFieldStride * kScratchFloatFieldCount;
378+ // Non-transpose inputs are gathered by one persistent SIMT VF per AIV
379+ // core. Each batch/class task therefore needs its own staging region so
380+ // the AIV NMS stage can consume it without relaunching the VF task.
381+ OP_CHECK_IF(taskCountU64 > std::numeric_limits<uint64_t>::max() / scratchBytesPerCore,
382+ OP_LOGE(context_, "Scratch workspace size overflow."), return ge::GRAPH_FAILED);
383+ const uint64_t scratchBytes = taskCountU64 * scratchBytesPerCore;
384+ OP_CHECK_IF(taskCountU64 > std::numeric_limits<uint64_t>::max() / maxSizePerClass,
385+ OP_LOGE(context_, "Class result workspace size overflow."), return ge::GRAPH_FAILED);
386+ const uint64_t classResultCount = taskCountU64 * maxSizePerClass;
387+ OP_CHECK_IF(classResultCount > std::numeric_limits<uint64_t>::max() / (4 * sizeof(float)) ||
388+ taskCountU64 > std::numeric_limits<uint64_t>::max() / sizeof(int32_t),
389+ OP_LOGE(context_, "Class result workspace size overflow."), return ge::GRAPH_FAILED);
390+ const uint64_t classBoxesBytes = classResultCount * 4 * sizeof(float);
391+ const uint64_t classScoresBytes = classResultCount * sizeof(float);
392+ const uint64_t classCountsBytes = taskCountU64 * sizeof(float);
393+ OP_CHECK_IF(
394+ static_cast<uint64_t>(batch_) > std::numeric_limits<uint64_t>::max() / static_cast<uint64_t>(maxTotalSize_),
395+ OP_LOGE(context_, "Cross-class TopK workspace size overflow."), return ge::GRAPH_FAILED);
396+ const uint64_t mergeResultCount = static_cast<uint64_t>(batch_) * static_cast<uint64_t>(maxTotalSize_);
397+ OP_CHECK_IF(mergeResultCount > std::numeric_limits<uint64_t>::max() / sizeof(float),
398+ OP_LOGE(context_, "Cross-class TopK workspace size overflow."), return ge::GRAPH_FAILED);
399+ const uint64_t mergeScoresBytes = mergeResultCount * sizeof(float);
400+ const uint64_t mergeIndicesBytes = mergeResultCount * sizeof(int32_t);
401+ uint64_t classBoxesOffset = scratchBytes;
402+ uint64_t classScoresOffset = 0;
403+ uint64_t classCountsOffset = 0;
404+ uint64_t mergeScoresOffset = 0;
405+ uint64_t mergeIndicesOffset = 0;
406+ uint64_t userWorkspaceBytes = 0;
407+ OP_CHECK_IF(!AddAligned(classBoxesOffset, classBoxesBytes, classScoresOffset) ||
408+ !AddAligned(classScoresOffset, classScoresBytes, classCountsOffset) ||
409+ !AddAligned(classCountsOffset, classCountsBytes, mergeScoresOffset) ||
410+ !AddAligned(mergeScoresOffset, mergeScoresBytes, mergeIndicesOffset) ||
411+ !AddAligned(mergeIndicesOffset, mergeIndicesBytes, userWorkspaceBytes),
412+ OP_LOGE(context_, "Class result workspace size overflow."), return ge::GRAPH_FAILED);
413+ tilingData_->batch = batch_;
414+ tilingData_->boxesNum = boxesNum_;
415+ tilingData_->classesNum = classesNum_;
416+ tilingData_->boxClassesNum = boxClassesNum_;
417+ tilingData_->maxSizePerClass = maxSizePerClass_;
418+ tilingData_->maxTotalSize = maxTotalSize_;
419+ tilingData_->usedCoreNum = usedCoreNum;
420+ tilingData_->tileSize = tileSize;
421+ tilingData_->reduceBufferSize = reduceBufferSize;
422+ tilingData_->mergeInputCount = mergeInputCount;
423+ tilingData_->mergeInputSize = mergeInputSize;
424+ tilingData_->mergeOutputCount = mergeOutputCount;
425+ tilingData_->mergeOutputSize = mergeOutputSize;
426+ tilingData_->scratchFieldStride = scratchFieldStride;
427+ tilingData_->scratchBytesPerCore = scratchBytesPerCore;
428+ tilingData_->classBoxesOffset = classBoxesOffset;
429+ tilingData_->classScoresOffset = classScoresOffset;
430+ tilingData_->classCountsOffset = classCountsOffset;
431+ tilingData_->mergeScoresOffset = mergeScoresOffset;
432+ tilingData_->mergeIndicesOffset = mergeIndicesOffset;
433+ tilingData_->topKTempBytes = topKTempBytes;
434+ tilingData_->mergeTopKTiling = mergeTopKTiling;
435+ tilingData_->scoreThreshold = scoreThreshold_;
436+ tilingData_->iouThreshold = iouThreshold_;
437+ tilingData_->hasClipWindow = static_cast<uint8_t>(hasClipWindow_);
438+ tilingData_->hasNumValidBoxes = static_cast<uint8_t>(hasNumValidBoxes_);
439+ tilingData_->changeCoordinateFrame = static_cast<uint8_t>(changeCoordinateFrame_);
440+ tilingData_->transposeBox = static_cast<uint8_t>(transposeBox_);
441+ tilingData_->use32Index = static_cast<uint8_t>(
442+ ProductFitsUint32(static_cast<uint64_t>(batch_), boxesNum, static_cast<uint64_t>(classesNum_), 1) &&
443+ ProductFitsUint32(static_cast<uint64_t>(batch_), boxesNum, static_cast<uint64_t>(boxClassesNum_), 4));
444+ 
445+ size_t* workspace = context_->GetWorkspaceSizes(1);
446+ OP_CHECK_NULL_WITH_CONTEXT(context_, workspace);
447+ OP_CHECK_IF(userWorkspaceBytes > std::numeric_limits<size_t>::max() - ascendcPlatform.GetLibApiWorkSpaceSize(),
448+ OP_LOGE(context_, "Workspace size exceeds platform limit."), return ge::GRAPH_FAILED);
449+ workspace[0] = ascendcPlatform.GetLibApiWorkSpaceSize() + static_cast<size_t>(userWorkspaceBytes);
450+ context_->SetBlockDim(usedCoreNum);
451+ context_->SetScheduleMode(kScheduleMode);
452+ context_->SetTilingKey(BATCH_MULTI_CLASS_NMS_TILING_KEY);
453+ 
454+ auto rawTilingData = context_->GetRawTilingData();
455+ OP_CHECK_NULL_WITH_CONTEXT(context_, rawTilingData);
456+ const size_t tilingDataSize = sizeof(BatchMultiClassNonMaxSuppressionTilingData);
457+ OP_CHECK_IF(memcpy_s(rawTilingData->GetData(), rawTilingData->GetCapacity(), tilingData_, tilingDataSize) != EOK,
458+ OP_LOGE(context_, "Failed to serialize tiling data."), return ge::GRAPH_FAILED);
459+ rawTilingData->SetDataSize(tilingDataSize);
460+ return ge::GRAPH_SUCCESS;
461+}
462+ 
463+ge::graphStatus BatchMultiClassNonMaxSuppressionTiling::RunTiling()
464+{
465+ const ge::graphStatus status = CheckAndParse();
466+ if (status != ge::GRAPH_SUCCESS) {
467+ return status;
468+ }
469+ return SetTilingData();
470+}
471+ 
472+static ge::graphStatus TilingPrepareForBatchMultiClassNonMaxSuppression(gert::TilingParseContext* context)
473+{
474+ auto* compileInfo = context->GetCompiledInfo<BatchMultiClassNonMaxSuppressionCompileInfo>();
475+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
476+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
477+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
478+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfo->ubSize);
479+ OP_CHECK_IF(compileInfo->coreNum == 0 || compileInfo->ubSize == 0,
480+ OP_LOGE(context, "Ascend950 platform resource query failed."), return ge::GRAPH_FAILED);
481+ return ge::GRAPH_SUCCESS;
482+}
483+ 
484+static ge::graphStatus TilingForBatchMultiClassNonMaxSuppression(gert::TilingContext* context)
485+{
486+ if (context == nullptr) {
487+ return ge::GRAPH_FAILED;
488+ }
489+ BatchMultiClassNonMaxSuppressionTiling tiling(context);
490+ return tiling.RunTiling();
491+}
492+ 
493+IMPL_OP_OPTILING(BatchMultiClassNonMaxSuppression)
494+ .Tiling(TilingForBatchMultiClassNonMaxSuppression)
495+ .TilingParse<BatchMultiClassNonMaxSuppressionCompileInfo>(TilingPrepareForBatchMultiClassNonMaxSuppression);
496+} // namespace optiling
Aobjdetect/batch_multi_class_non_max_suppression/op_host/arch35/batch_multi_class_non_max_suppression_tiling_arch35.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+#ifndef BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_ARCH35_H_
12+#define BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_ARCH35_H_
13+ 
14+#include "register/op_impl_registry.h"
15+#include "register/tilingdata_base.h"
16+#include "tiling/tiling_api.h"
17+#include "../../op_kernel/arch35/batch_multi_class_non_max_suppression_tiling_data.h"
18+ 
19+namespace optiling {
20+struct BatchMultiClassNonMaxSuppressionCompileInfo {
21+ uint32_t coreNum{0};
22+ uint64_t ubSize{0};
23+};
24+ 
25+class BatchMultiClassNonMaxSuppressionTiling {
26+public:
27+ explicit BatchMultiClassNonMaxSuppressionTiling(gert::TilingContext* context) : context_(context) {}
28+ ge::graphStatus RunTiling();
29+ 
30+private:
31+ ge::graphStatus CheckAndParse();
32+ ge::graphStatus SetTilingData();
33+ 
34+ gert::TilingContext* context_;
35+ BatchMultiClassNonMaxSuppressionTilingData* tilingData_{nullptr};
36+ int64_t batch_{0};
37+ int64_t boxesNum_{0};
38+ int64_t classesNum_{0};
39+ int64_t boxClassesNum_{0};
40+ int64_t maxSizePerClass_{0};
41+ int64_t maxTotalSize_{0};
42+ float scoreThreshold_{0.0F};
43+ float iouThreshold_{0.0F};
44+ bool hasClipWindow_{false};
45+ bool hasNumValidBoxes_{false};
46+ bool changeCoordinateFrame_{false};
47+ bool transposeBox_{false};
48+};
49+} // namespace optiling
50+ 
51+#endif // BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_ARCH35_H_
Aobjdetect/batch_multi_class_non_max_suppression/op_host/batch_multi_class_non_max_suppression_def.cpp+84-0
@@ -0,0 +1,84 @@
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 "register/op_def_registry.h"
12+ 
13+namespace ops {
14+class BatchMultiClassNonMaxSuppression : public OpDef {
15+public:
16+ explicit BatchMultiClassNonMaxSuppression(const char* name) : OpDef(name)
17+ {
18+ this->Input("boxes")
19+ .ParamType(REQUIRED)
20+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
21+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
22+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
23+ .AutoContiguous();
24+ this->Input("scores")
25+ .ParamType(REQUIRED)
26+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
27+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
28+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
29+ .AutoContiguous();
30+ this->Input("clip_window")
31+ .ParamType(OPTIONAL)
32+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
33+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
34+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
35+ .AutoContiguous();
36+ this->Input("num_valid_boxes")
37+ .ParamType(OPTIONAL)
38+ .DataType({ge::DT_INT32, ge::DT_INT32})
39+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
40+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
41+ .AutoContiguous();
42+ this->Output("nmsed_boxes")
43+ .ParamType(REQUIRED)
44+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
45+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
46+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
47+ this->Output("nmsed_scores")
48+ .ParamType(REQUIRED)
49+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
50+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
51+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
52+ this->Output("nmsed_classes")
53+ .ParamType(REQUIRED)
54+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
55+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
56+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
57+ this->Output("nmsed_num")
58+ .ParamType(REQUIRED)
59+ .DataType({ge::DT_INT32, ge::DT_INT32})
60+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
61+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
62+ 
63+ this->Attr("score_threshold").AttrType(REQUIRED).Float();
64+ this->Attr("iou_threshold").AttrType(REQUIRED).Float();
65+ this->Attr("max_size_per_class").AttrType(REQUIRED).Int();
66+ this->Attr("max_total_size").AttrType(REQUIRED).Int();
67+ this->Attr("change_coordinate_frame").AttrType(OPTIONAL).Bool(false);
68+ this->Attr("transpose_box").AttrType(OPTIONAL).Bool(false);
69+ this->Attr("image_size").AttrType(OPTIONAL).ListInt();
70+ 
71+ OpAICoreConfig aicoreConfig;
72+ aicoreConfig.DynamicCompileStaticFlag(true)
73+ .DynamicFormatFlag(false)
74+ .DynamicRankSupportFlag(false)
75+ .DynamicShapeSupportFlag(true)
76+ .NeedCheckSupportFlag(false)
77+ .PrecisionReduceFlag(false)
78+ .ExtendCfgInfo("opFile.value", "batch_multi_class_non_max_suppression_apt");
79+ this->AICore().AddConfig("ascend950", aicoreConfig);
80+ }
81+};
82+ 
83+OP_ADD(BatchMultiClassNonMaxSuppression);
84+} // namespace ops
Aobjdetect/batch_multi_class_non_max_suppression/op_host/batch_multi_class_non_max_suppression_infershape.cpp+79-0
@@ -0,0 +1,79 @@
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 "log/log.h"
12+#include "register/op_impl_registry.h"
13+ 
14+namespace {
15+constexpr int64_t kBoxesIndex = 0;
16+constexpr int64_t kScoresIndex = 1;
17+constexpr int64_t kNmsedBoxesIndex = 0;
18+constexpr int64_t kNmsedScoresIndex = 1;
19+constexpr int64_t kNmsedClassesIndex = 2;
20+constexpr int64_t kNmsedNumIndex = 3;
21+constexpr int64_t kMaxTotalSizeAttrIndex = 3;
22+} // namespace
23+ 
24+namespace ops {
25+static ge::graphStatus InferShapeForBatchMultiClassNonMaxSuppression(gert::InferShapeContext* context)
26+{
27+ const gert::Shape* boxesShape = context->GetInputShape(kBoxesIndex);
28+ const gert::Shape* scoresShape = context->GetInputShape(kScoresIndex);
29+ const gert::RuntimeAttrs* attrs = context->GetAttrs();
30+ OP_CHECK_NULL_WITH_CONTEXT(context, boxesShape);
31+ OP_CHECK_NULL_WITH_CONTEXT(context, scoresShape);
32+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
33+ 
34+ OP_CHECK_IF(boxesShape->GetDimNum() != 4 || scoresShape->GetDimNum() != 3,
35+ OP_LOGE(context, "boxes must be rank 4 and scores must be rank 3."), return ge::GRAPH_FAILED);
36+ const int64_t batch = boxesShape->GetDim(0);
37+ const int64_t scoresBatch = scoresShape->GetDim(0);
38+ const int64_t* maxTotalSize = attrs->GetAttrPointer<int64_t>(kMaxTotalSizeAttrIndex);
39+ OP_CHECK_NULL_WITH_CONTEXT(context, maxTotalSize);
40+ OP_CHECK_IF(batch <= 0 || scoresBatch != batch || *maxTotalSize <= 0,
41+ OP_LOGE(context, "Invalid batch dimension or max_total_size."), return ge::GRAPH_FAILED);
atomgit-bot
atomgit-botatomgit-bot25 天前

🔴 Critical

在 infershape.cpp 第 41 行的 OP_LOGE(context_, ...) 中，使用了未定义的变量名 context_（带下划线）。该函数的参数名为 context（无下划线），在函数作用域内不存在 context_。同一函数中其他 OP_LOGE 调用（如第 35 行）均正确使用 context，唯独此处误写为 context_。

证据链：

失效模式：编译器报 use of undeclared identifier 'context_'，导致算子模块无法编译通过。

建议：将 context_ 改为 context，与函数参数名一致。

likedislike
42+ 
43+ gert::Shape* nmsedBoxesShape = context->GetOutputShape(kNmsedBoxesIndex);
44+ gert::Shape* nmsedScoresShape = context->GetOutputShape(kNmsedScoresIndex);
45+ gert::Shape* nmsedClassesShape = context->GetOutputShape(kNmsedClassesIndex);
46+ gert::Shape* nmsedNumShape = context->GetOutputShape(kNmsedNumIndex);
47+ OP_CHECK_NULL_WITH_CONTEXT(context, nmsedBoxesShape);
48+ OP_CHECK_NULL_WITH_CONTEXT(context, nmsedScoresShape);
49+ OP_CHECK_NULL_WITH_CONTEXT(context, nmsedClassesShape);
50+ OP_CHECK_NULL_WITH_CONTEXT(context, nmsedNumShape);
51+ 
52+ nmsedBoxesShape->SetDimNum(3);
53+ nmsedBoxesShape->SetDim(0, batch);
54+ nmsedBoxesShape->SetDim(1, *maxTotalSize);
55+ nmsedBoxesShape->SetDim(2, 4);
56+ for (gert::Shape* output : {nmsedScoresShape, nmsedClassesShape}) {
57+ output->SetDimNum(2);
58+ output->SetDim(0, batch);
59+ output->SetDim(1, *maxTotalSize);
60+ }
61+ nmsedNumShape->SetDimNum(1);
62+ nmsedNumShape->SetDim(0, batch);
63+ return ge::GRAPH_SUCCESS;
64+}
65+ 
66+static ge::graphStatus InferDataTypeForBatchMultiClassNonMaxSuppression(gert::InferDataTypeContext* context)
67+{
68+ const ge::DataType boxesType = context->GetInputDataType(kBoxesIndex);
69+ context->SetOutputDataType(kNmsedBoxesIndex, boxesType);
70+ context->SetOutputDataType(kNmsedScoresIndex, boxesType);
71+ context->SetOutputDataType(kNmsedClassesIndex, boxesType);
72+ context->SetOutputDataType(kNmsedNumIndex, ge::DT_INT32);
73+ return ge::GRAPH_SUCCESS;
74+}
75+ 
76+IMPL_OP_INFERSHAPE(BatchMultiClassNonMaxSuppression)
77+ .InferShape(InferShapeForBatchMultiClassNonMaxSuppression)
78+ .InferDataType(InferDataTypeForBatchMultiClassNonMaxSuppression);
79+} // namespace ops
Aobjdetect/batch_multi_class_non_max_suppression/op_host/config/ascend950/batch_multi_class_non_max_suppression_binary.json+269-0
@@ -0,0 +1,269 @@
1+{
2+ "op_type": "BatchMultiClassNonMaxSuppression",
3+ "op_list": [
4+ {
5+ "bin_filename": "BatchMultiClassNonMaxSuppression_float16",
6+ "inputs": [
7+ {
8+ "name": "boxes",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ },
18+ {
19+ "name": "scores",
20+ "index": 1,
21+ "dtype": "float16",
22+ "format": "ND",
23+ "paramType": "required",
24+ "shape": [
25+ -2
26+ ],
27+ "format_match_mode": "FormatAgnostic"
28+ },
29+ {
30+ "name": "clip_window",
31+ "index": 2,
32+ "dtype": "float16",
33+ "format": "ND",
34+ "paramType": "optional",
35+ "shape": [
36+ -2
37+ ],
38+ "format_match_mode": "FormatAgnostic"
39+ },
40+ {
41+ "name": "num_valid_boxes",
42+ "index": 3,
43+ "dtype": "int32",
44+ "format": "ND",
45+ "paramType": "optional",
46+ "shape": [
47+ -2
48+ ],
49+ "format_match_mode": "FormatAgnostic"
50+ }
51+ ],
52+ "outputs": [
53+ {
54+ "name": "nmsed_boxes",
55+ "index": 0,
56+ "dtype": "float16",
57+ "format": "ND",
58+ "paramType": "required",
59+ "shape": [
60+ -2
61+ ],
62+ "format_match_mode": "FormatAgnostic"
63+ },
64+ {
65+ "name": "nmsed_scores",
66+ "index": 1,
67+ "dtype": "float16",
68+ "format": "ND",
69+ "paramType": "required",
70+ "shape": [
71+ -2
72+ ],
73+ "format_match_mode": "FormatAgnostic"
74+ },
75+ {
76+ "name": "nmsed_classes",
77+ "index": 2,
78+ "dtype": "float16",
79+ "format": "ND",
80+ "paramType": "required",
81+ "shape": [
82+ -2
83+ ],
84+ "format_match_mode": "FormatAgnostic"
85+ },
86+ {
87+ "name": "nmsed_num",
88+ "index": 3,
89+ "dtype": "int32",
90+ "format": "ND",
91+ "paramType": "required",
92+ "shape": [
93+ -2
94+ ],
95+ "format_match_mode": "FormatAgnostic"
96+ }
97+ ],
98+ "attrs": [
99+ {
100+ "name": "score_threshold",
101+ "dtype": "float",
102+ "value": 0.0
103+ },
104+ {
105+ "name": "iou_threshold",
106+ "dtype": "float",
107+ "value": 0.0
108+ },
109+ {
110+ "name": "max_size_per_class",
111+ "dtype": "int",
112+ "value": 0
113+ },
114+ {
115+ "name": "max_total_size",
116+ "dtype": "int",
117+ "value": 0
118+ },
119+ {
120+ "name": "change_coordinate_frame",
121+ "dtype": "bool",
122+ "value": false
123+ },
124+ {
125+ "name": "transpose_box",
126+ "dtype": "bool",
127+ "value": false
128+ },
129+ {
130+ "name": "image_size",
131+ "dtype": "list_int",
132+ "value": []
133+ }
134+ ]
135+ },
136+ {
137+ "bin_filename": "BatchMultiClassNonMaxSuppression_float32",
138+ "inputs": [
139+ {
140+ "name": "boxes",
141+ "index": 0,
142+ "dtype": "float32",
143+ "format": "ND",
144+ "paramType": "required",
145+ "shape": [
146+ -2
147+ ],
148+ "format_match_mode": "FormatAgnostic"
149+ },
150+ {
151+ "name": "scores",
152+ "index": 1,
153+ "dtype": "float32",
154+ "format": "ND",
155+ "paramType": "required",
156+ "shape": [
157+ -2
158+ ],
159+ "format_match_mode": "FormatAgnostic"
160+ },
161+ {
162+ "name": "clip_window",
163+ "index": 2,
164+ "dtype": "float32",
165+ "format": "ND",
166+ "paramType": "optional",
167+ "shape": [
168+ -2
169+ ],
170+ "format_match_mode": "FormatAgnostic"
171+ },
172+ {
173+ "name": "num_valid_boxes",
174+ "index": 3,
175+ "dtype": "int32",
176+ "format": "ND",
177+ "paramType": "optional",
178+ "shape": [
179+ -2
180+ ],
181+ "format_match_mode": "FormatAgnostic"
182+ }
183+ ],
184+ "outputs": [
185+ {
186+ "name": "nmsed_boxes",
187+ "index": 0,
188+ "dtype": "float32",
189+ "format": "ND",
190+ "paramType": "required",
191+ "shape": [
192+ -2
193+ ],
194+ "format_match_mode": "FormatAgnostic"
195+ },
196+ {
197+ "name": "nmsed_scores",
198+ "index": 1,
199+ "dtype": "float32",
200+ "format": "ND",
201+ "paramType": "required",
202+ "shape": [
203+ -2
204+ ],
205+ "format_match_mode": "FormatAgnostic"
206+ },
207+ {
208+ "name": "nmsed_classes",
209+ "index": 2,
210+ "dtype": "float32",
211+ "format": "ND",
212+ "paramType": "required",
213+ "shape": [
214+ -2
215+ ],
216+ "format_match_mode": "FormatAgnostic"
217+ },
218+ {
219+ "name": "nmsed_num",
220+ "index": 3,
221+ "dtype": "int32",
222+ "format": "ND",
223+ "paramType": "required",
224+ "shape": [
225+ -2
226+ ],
227+ "format_match_mode": "FormatAgnostic"
228+ }
229+ ],
230+ "attrs": [
231+ {
232+ "name": "score_threshold",
233+ "dtype": "float",
234+ "value": 0.0
235+ },
236+ {
237+ "name": "iou_threshold",
238+ "dtype": "float",
239+ "value": 0.0
240+ },
241+ {
242+ "name": "max_size_per_class",
243+ "dtype": "int",
244+ "value": 0
245+ },
246+ {
247+ "name": "max_total_size",
248+ "dtype": "int",
249+ "value": 0
250+ },
251+ {
252+ "name": "change_coordinate_frame",
253+ "dtype": "bool",
254+ "value": false
255+ },
256+ {
257+ "name": "transpose_box",
258+ "dtype": "bool",
259+ "value": false
260+ },
261+ {
262+ "name": "image_size",
263+ "dtype": "list_int",
264+ "value": []
265+ }
266+ ]
267+ }
268+ ]
269+}
Aobjdetect/batch_multi_class_non_max_suppression/op_host/config/ascend950/batch_multi_class_non_max_suppression_simplified_key.ini+3-0
@@ -0,0 +1,3 @@
1+; The initial Ascend950 implementation uses the default Ascend C simplified-key mode.
2+[BatchMultiClassNonMaxSuppression]
3+default=0
Aobjdetect/batch_multi_class_non_max_suppression/op_kernel/arch35/batch_multi_class_non_max_suppression_kernel.h+830-0
@@ -0,0 +1,830 @@
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+#ifndef BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_KERNEL_H_
12+#define BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_KERNEL_H_
13+ 
14+#include <cstdint>
15+ 
16+#include "kernel_operator.h"
17+#include "simt_api/asc_simt.h"
18+#include "batch_multi_class_non_max_suppression_tiling_data.h"
19+ 
20+namespace BatchMultiClassNonMaxSuppressionOp {
21+using namespace AscendC;
22+ 
23+constexpr uint32_t kGatherThreadNum32 = 1024;
24+constexpr uint32_t kGatherThreadNum64 = 512;
25+constexpr uint32_t kMergeThreadNum = 256;
26+constexpr float kNoCandidate = -(__builtin_inff());
27+constexpr float kMinPositive = 1.0e-12F;
28+constexpr TopKConfig kMergeTopKConfig{TopKAlgo::RADIX_SELECT, TopKOrder::UNSET, true};
29+ 
30+template <bool Use32Bit>
31+struct GatherIndexType {
32+ using type = uint64_t;
33+};
34+ 
35+template <>
36+struct GatherIndexType<true> {
37+ using type = uint32_t;
38+};
39+ 
40+template <bool Use32Bit>
41+using GatherIndex = typename GatherIndexType<Use32Bit>::type;
42+ 
43+// Scores are strided by class and boxes can be either [B, N, q, 4] or
44+// [B, q, 4, N]. This is an irregular GM access pattern, so use the 950 SIMT
45+// unit to compact one (batch, class) task into five contiguous FP32 arrays.
46+// Subsequent score reduction and IoU work is entirely vectorized on UB tiles.
47+template <typename T, bool Use32Bit>
48+__simt_vf__ __aicore__
49+__launch_bounds__(Use32Bit ? kGatherThreadNum32 : kGatherThreadNum64) inline void GatherClassInput(
50+ const __gm__ T* boxesYMin, const __gm__ T* boxesXMin, const __gm__ T* boxesYMax, const __gm__ T* boxesXMax,
51+ const __gm__ T* scores, const __gm__ int32_t* numValidBoxes, __gm__ float* stageScores, __gm__ float* stageYMin,
52+ __gm__ float* stageXMin, __gm__ float* stageYMax, __gm__ float* stageXMax, GatherIndex<Use32Bit> batchIndex,
53+ GatherIndex<Use32Bit> classIndex, GatherIndex<Use32Bit> boxesNum, GatherIndex<Use32Bit> classesNum,
54+ bool hasNumValidBoxes)
55+{
56+ using IndexT = GatherIndex<Use32Bit>;
57+ IndexT validBoxes = boxesNum;
58+ if (hasNumValidBoxes) {
59+ const int32_t validBoxesRaw = numValidBoxes[batchIndex];
60+ validBoxes = validBoxesRaw <= 0 ?
61+ 0 :
62+ (static_cast<IndexT>(validBoxesRaw) > boxesNum ? boxesNum :
63+ static_cast<IndexT>(validBoxesRaw));
64+ }
65+ for (IndexT boxIndex = static_cast<IndexT>(threadIdx.x); boxIndex < boxesNum;
66+ boxIndex += static_cast<IndexT>(blockDim.x)) {
67+ const IndexT scoreOffset = (batchIndex * boxesNum + boxIndex) * classesNum + classIndex;
68+ stageScores[boxIndex] = boxIndex < validBoxes ? static_cast<float>(scores[scoreOffset]) : kNoCandidate;
69+ // Each [B, q, 4, N] field is contiguous. Their bases are computed
70+ // by AIV before the SIMT launch, so this path only needs a per-box
71+ // offset. Non-transpose inputs use the persistent gather below.
72+ stageYMin[boxIndex] = static_cast<float>(boxesYMin[boxIndex]);
73+ stageXMin[boxIndex] = static_cast<float>(boxesXMin[boxIndex]);
74+ stageYMax[boxIndex] = static_cast<float>(boxesYMax[boxIndex]);
75+ stageXMax[boxIndex] = static_cast<float>(boxesXMax[boxIndex]);
76+ }
77+}
78+ 
79+// Keep one asynchronous VF task alive per AIV core for the complete class
80+// wave. Re-launching a short VF task for every class exhausts the 950 VF
81+// task queue after several waves. Each task writes to an independent staging
82+// slice, which the AIV NMS stage subsequently consumes.
83+template <typename T, bool Use32Bit>
84+__simt_vf__ __aicore__
85+__launch_bounds__(Use32Bit ? kGatherThreadNum32 : kGatherThreadNum64) inline void GatherAllClassInputs(
86+ const __gm__ T* boxes, const __gm__ T* scores, const __gm__ int32_t* numValidBoxes, __gm__ float* scratch,
87+ GatherIndex<Use32Bit> taskStart, GatherIndex<Use32Bit> taskStride, GatherIndex<Use32Bit> taskCount,
88+ GatherIndex<Use32Bit> boxesNum, GatherIndex<Use32Bit> classesNum, GatherIndex<Use32Bit> boxClassesNum,
89+ uint64_t scratchFieldStride, bool hasNumValidBoxes)
90+{
91+ using IndexT = GatherIndex<Use32Bit>;
92+ const uint64_t scratchFieldElements = scratchFieldStride / sizeof(float);
93+ for (IndexT taskIndex = taskStart; taskIndex < taskCount; taskIndex += taskStride) {
94+ const IndexT batchIndex = taskIndex / classesNum;
95+ const IndexT classIndex = taskIndex % classesNum;
96+ IndexT validBoxes = boxesNum;
97+ if (hasNumValidBoxes) {
98+ const int32_t validBoxesRaw = numValidBoxes[batchIndex];
99+ validBoxes = validBoxesRaw <= 0 ?
100+ 0 :
101+ (static_cast<IndexT>(validBoxesRaw) > boxesNum ? boxesNum :
102+ static_cast<IndexT>(validBoxesRaw));
103+ }
104+ const uint64_t taskScratchOffset = static_cast<uint64_t>(taskIndex) * scratchFieldElements * 5;
105+ __gm__ float* const stageScores = scratch + taskScratchOffset;
106+ __gm__ float* const stageYMin = stageScores + scratchFieldElements;
107+ __gm__ float* const stageXMin = stageYMin + scratchFieldElements;
108+ __gm__ float* const stageYMax = stageXMin + scratchFieldElements;
109+ __gm__ float* const stageXMax = stageYMax + scratchFieldElements;
110+ const IndexT boxClass = boxClassesNum == 1 ? 0 : classIndex;
111+ for (IndexT boxIndex = static_cast<IndexT>(threadIdx.x); boxIndex < boxesNum;
112+ boxIndex += static_cast<IndexT>(blockDim.x)) {
113+ const IndexT scoreOffset = (batchIndex * boxesNum + boxIndex) * classesNum + classIndex;
114+ const IndexT boxOffset = ((batchIndex * boxesNum + boxIndex) * boxClassesNum + boxClass) * 4;
115+ stageScores[boxIndex] = boxIndex < validBoxes ? static_cast<float>(scores[scoreOffset]) : kNoCandidate;
116+ stageYMin[boxIndex] = static_cast<float>(boxes[boxOffset]);
117+ stageXMin[boxIndex] = static_cast<float>(boxes[boxOffset + 1]);
118+ stageYMax[boxIndex] = static_cast<float>(boxes[boxOffset + 2]);
119+ stageXMax[boxIndex] = static_cast<float>(boxes[boxOffset + 3]);
120+ }
121+ }
122+}
123+ 
124+// The cross-class candidates are already compacted in class-major order. AIV
125+// TopK produces their flattened positions; this SIMT epilogue turns those
126+// positions into the irregular box/class output writes without scalar GM
127+// gathers in the AIV pipeline.
128+template <typename T>
129+__simt_vf__ __aicore__ __launch_bounds__(kMergeThreadNum) inline void GatherMergedOutput(
130+ const __gm__ float* classBoxes, const __gm__ float* mergeScores, const __gm__ int32_t* mergeIndices,
131+ __gm__ T* nmsedBoxes, __gm__ T* nmsedScores, __gm__ T* nmsedClasses, __gm__ int32_t* nmsedNum, uint64_t batchIndex,
132+ uint64_t classesNum, uint64_t maxSizePerClass, uint64_t maxTotalSize, uint64_t validOutputCount)
133+{
134+ // The TopK input is already compacted by class counts, so thread zero can
135+ // publish the exact output size while this existing SIMT epilogue starts.
136+ // This replaces the former serial scan of mergeScores without adding a
137+ // separate AIV-to-GM transfer.
138+ if (threadIdx.x == 0) {
139+ nmsedNum[batchIndex] = static_cast<int32_t>(validOutputCount);
140+ }
141+ for (uint64_t outputIndex = static_cast<uint64_t>(threadIdx.x); outputIndex < maxTotalSize;
142+ outputIndex += static_cast<uint64_t>(blockDim.x)) {
143+ const uint64_t outputOffset = batchIndex * maxTotalSize + outputIndex;
144+ if (outputIndex < validOutputCount) {
145+ const float score = mergeScores[outputOffset];
146+ const uint64_t flatIndex = static_cast<uint64_t>(mergeIndices[outputOffset]);
147+ const uint64_t classIndex = flatIndex / maxSizePerClass;
148+ const uint64_t classPosition = flatIndex % maxSizePerClass;
149+ const uint64_t candidateOffset = (batchIndex * classesNum + classIndex) * maxSizePerClass + classPosition;
150+ const uint64_t boxOffset = candidateOffset * 4;
151+ const uint64_t outputBoxOffset = outputOffset * 4;
152+ nmsedBoxes[outputBoxOffset] = static_cast<T>(classBoxes[boxOffset]);
153+ nmsedBoxes[outputBoxOffset + 1] = static_cast<T>(classBoxes[boxOffset + 1]);
154+ nmsedBoxes[outputBoxOffset + 2] = static_cast<T>(classBoxes[boxOffset + 2]);
155+ nmsedBoxes[outputBoxOffset + 3] = static_cast<T>(classBoxes[boxOffset + 3]);
156+ nmsedScores[outputOffset] = static_cast<T>(score);
157+ nmsedClasses[outputOffset] = static_cast<T>(classIndex);
158+ } else {
159+ const uint64_t outputBoxOffset = outputOffset * 4;
160+ nmsedBoxes[outputBoxOffset] = static_cast<T>(0);
161+ nmsedBoxes[outputBoxOffset + 1] = static_cast<T>(0);
162+ nmsedBoxes[outputBoxOffset + 2] = static_cast<T>(0);
163+ nmsedBoxes[outputBoxOffset + 3] = static_cast<T>(0);
164+ nmsedScores[outputOffset] = static_cast<T>(0);
165+ nmsedClasses[outputOffset] = static_cast<T>(0);
166+ }
167+ }
168+}
169+ 
170+template <typename T>
171+class BatchMultiClassNonMaxSuppressionKernel {
172+public:
173+ __aicore__ inline void Init(GM_ADDR boxes, GM_ADDR scores, GM_ADDR clipWindow, GM_ADDR numValidBoxes,
174+ GM_ADDR nmsedBoxes, GM_ADDR nmsedScores, GM_ADDR nmsedClasses, GM_ADDR nmsedNum,
175+ GM_ADDR workspace, const BatchMultiClassNonMaxSuppressionTilingData* tilingData);
176+ __aicore__ inline void Process();
177+ 
178+private:
179+ __aicore__ inline void InitWorkspace();
180+ __aicore__ inline void InitTileBuffers();
181+ __aicore__ inline void GatherClass(int64_t taskIndex);
182+ __aicore__ inline bool FilterStageBoxes(int64_t batchIndex);
183+ __aicore__ inline bool FindBestCandidate(int64_t& bestIndex, float& bestScore);
184+ __aicore__ inline void SuppressBySelectedBox(float yMin, float xMin, float yMax, float xMax);
185+ __aicore__ inline void ProcessClass(int64_t taskIndex);
186+ __aicore__ inline void PadClassScores(int64_t taskIndex, int64_t selectedCount);
187+ __aicore__ inline void StoreClassCount(int64_t taskIndex, int64_t selectedCount);
188+ __aicore__ inline void MergeBatch(int64_t batchIndex);
189+ __aicore__ inline void LoadTile(int64_t offset, int64_t count);
190+ __aicore__ inline void StoreTile(int64_t offset, int64_t count);
191+ __aicore__ inline void LoadScores(int64_t offset, int64_t count);
192+ __aicore__ inline void StoreScores(int64_t offset, int64_t count);
193+ __aicore__ inline void CopyIn(LocalTensor<float>& dst, const GlobalTensor<float>& src, int64_t offset,
194+ int64_t count);
195+ __aicore__ inline void CopyOut(GlobalTensor<float>& dst, int64_t offset, LocalTensor<float>& src, int64_t count);
196+ TPipe pipe_;
197+ TBuf<QuePosition::VECCALC> scoreBuffer_;
198+ TBuf<QuePosition::VECCALC> yMinBuffer_;
199+ TBuf<QuePosition::VECCALC> xMinBuffer_;
200+ TBuf<QuePosition::VECCALC> yMaxBuffer_;
201+ TBuf<QuePosition::VECCALC> xMaxBuffer_;
202+ TBuf<QuePosition::VECCALC> temp0Buffer_;
203+ TBuf<QuePosition::VECCALC> temp1Buffer_;
204+ TBuf<QuePosition::VECCALC> temp2Buffer_;
205+ TBuf<QuePosition::VECCALC> temp3Buffer_;
206+ TBuf<QuePosition::VECCALC> reduceWorkBuffer_;
207+ TBuf<QuePosition::VECCALC> reduceOutputBuffer_;
208+ TBuf<QuePosition::VECCALC> compareMaskBuffer_;
209+ TBuf<QuePosition::VECCALC> mergeInputScoresBuffer_;
210+ TBuf<QuePosition::VECCALC> mergeInputIndicesBuffer_;
211+ TBuf<QuePosition::VECCALC> mergeOutputScoresBuffer_;
212+ TBuf<QuePosition::VECCALC> mergeOutputIndicesBuffer_;
213+ TBuf<QuePosition::VECCALC> mergeTopKTempBuffer_;
214+ 
215+ LocalTensor<float> scoreLocal_;
216+ LocalTensor<float> yMinLocal_;
217+ LocalTensor<float> xMinLocal_;
218+ LocalTensor<float> yMaxLocal_;
219+ LocalTensor<float> xMaxLocal_;
220+ LocalTensor<float> temp0Local_;
221+ LocalTensor<float> temp1Local_;
222+ LocalTensor<float> temp2Local_;
223+ LocalTensor<float> temp3Local_;
224+ LocalTensor<float> reduceWorkLocal_;
225+ LocalTensor<float> reduceOutputLocal_;
226+ LocalTensor<uint8_t> compareMaskLocal_;
227+ LocalTensor<float> mergeInputScoresLocal_;
228+ LocalTensor<int32_t> mergeInputIndicesLocal_;
229+ LocalTensor<float> mergeOutputScoresLocal_;
230+ LocalTensor<int32_t> mergeOutputIndicesLocal_;
231+ LocalTensor<uint8_t> mergeTopKTempLocal_;
232+ 
233+ GlobalTensor<T> clipWindowGm_;
234+ GlobalTensor<T> nmsedBoxesGm_;
235+ GlobalTensor<T> nmsedScoresGm_;
236+ GlobalTensor<T> nmsedClassesGm_;
237+ GlobalTensor<int32_t> nmsedNumGm_;
238+ GlobalTensor<float> stageScoresGm_;
239+ GlobalTensor<float> stageYMinGm_;
240+ GlobalTensor<float> stageXMinGm_;
241+ GlobalTensor<float> stageYMaxGm_;
242+ GlobalTensor<float> stageXMaxGm_;
243+ GlobalTensor<float> classBoxesGm_;
244+ GlobalTensor<float> classScoresGm_;
245+ GlobalTensor<float> classCountsGm_;
246+ GlobalTensor<float> mergeScoresGm_;
247+ GlobalTensor<int32_t> mergeIndicesGm_;
248+ 
249+ GM_ADDR boxesAddr_{nullptr};
250+ GM_ADDR scoresAddr_{nullptr};
251+ GM_ADDR numValidBoxesAddr_{nullptr};
252+ GM_ADDR userWorkspace_{nullptr};
253+ const BatchMultiClassNonMaxSuppressionTilingData* tilingData_{nullptr};
254+ int64_t cachedClipBatch_{-1};
255+ float cachedClipYMin_{0.0F};
256+ float cachedClipXMin_{0.0F};
257+ float cachedClipYMax_{0.0F};
258+ float cachedClipXMax_{0.0F};
259+ bool workspaceReady_{false};
260+};
261+ 
262+template <typename T>
263+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::Init(
264+ GM_ADDR boxes, GM_ADDR scores, GM_ADDR clipWindow, GM_ADDR numValidBoxes, GM_ADDR nmsedBoxes, GM_ADDR nmsedScores,
265+ GM_ADDR nmsedClasses, GM_ADDR nmsedNum, GM_ADDR workspace,
266+ const BatchMultiClassNonMaxSuppressionTilingData* tilingData)
267+{
268+ tilingData_ = tilingData;
269+ boxesAddr_ = boxes;
270+ scoresAddr_ = scores;
271+ numValidBoxesAddr_ = numValidBoxes;
272+ workspaceReady_ = workspace != nullptr;
273+ if (!workspaceReady_) {
274+ return;
275+ }
276+ userWorkspace_ = GetUserWorkspace(workspace);
277+ workspaceReady_ = userWorkspace_ != nullptr;
278+ if (!workspaceReady_) {
279+ return;
280+ }
281+ 
282+ if (tilingData_->hasClipWindow != 0U) {
283+ clipWindowGm_.SetGlobalBuffer((__gm__ T*)clipWindow, tilingData_->batch * 4);
284+ }
285+ const int64_t resultElements = tilingData_->batch * tilingData_->maxTotalSize;
286+ nmsedBoxesGm_.SetGlobalBuffer((__gm__ T*)nmsedBoxes, resultElements * 4);
287+ nmsedScoresGm_.SetGlobalBuffer((__gm__ T*)nmsedScores, resultElements);
288+ nmsedClassesGm_.SetGlobalBuffer((__gm__ T*)nmsedClasses, resultElements);
289+ nmsedNumGm_.SetGlobalBuffer((__gm__ int32_t*)nmsedNum, tilingData_->batch);
290+ InitWorkspace();
291+ InitTileBuffers();
292+}
293+ 
294+template <typename T>
295+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::InitWorkspace()
296+{
297+ const int64_t taskCount = tilingData_->batch * tilingData_->classesNum;
298+ const int64_t classResultCount = taskCount * tilingData_->maxSizePerClass;
299+ const int64_t classBoxesElements = classResultCount * 4;
300+ classBoxesGm_.SetGlobalBuffer((__gm__ float*)(userWorkspace_ + tilingData_->classBoxesOffset), classBoxesElements);
301+ classScoresGm_.SetGlobalBuffer((__gm__ float*)(userWorkspace_ + tilingData_->classScoresOffset), classResultCount);
302+ classCountsGm_.SetGlobalBuffer((__gm__ float*)(userWorkspace_ + tilingData_->classCountsOffset), taskCount);
303+ const int64_t mergeResultCount = tilingData_->batch * tilingData_->maxTotalSize;
304+ mergeScoresGm_.SetGlobalBuffer((__gm__ float*)(userWorkspace_ + tilingData_->mergeScoresOffset), mergeResultCount);
305+ mergeIndicesGm_.SetGlobalBuffer((__gm__ int32_t*)(userWorkspace_ + tilingData_->mergeIndicesOffset),
306+ mergeResultCount);
307+}
308+ 
309+template <typename T>
310+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::InitTileBuffers()
311+{
312+ const int64_t tileSize = tilingData_->tileSize;
313+ const int64_t floatBytes = tileSize * static_cast<int64_t>(sizeof(float));
314+ pipe_.InitBuffer(scoreBuffer_, floatBytes);
315+ pipe_.InitBuffer(yMinBuffer_, floatBytes);
316+ pipe_.InitBuffer(xMinBuffer_, floatBytes);
317+ pipe_.InitBuffer(yMaxBuffer_, floatBytes);
318+ pipe_.InitBuffer(xMaxBuffer_, floatBytes);
319+ pipe_.InitBuffer(temp0Buffer_, floatBytes);
320+ pipe_.InitBuffer(temp1Buffer_, floatBytes);
321+ pipe_.InitBuffer(temp2Buffer_, floatBytes);
322+ pipe_.InitBuffer(temp3Buffer_, floatBytes);
323+ pipe_.InitBuffer(reduceWorkBuffer_, tilingData_->reduceBufferSize * static_cast<int64_t>(sizeof(float)));
324+ pipe_.InitBuffer(reduceOutputBuffer_, 64);
325+ pipe_.InitBuffer(compareMaskBuffer_, tileSize * static_cast<int64_t>(sizeof(uint8_t)));
326+ pipe_.InitBuffer(mergeInputScoresBuffer_, tilingData_->mergeInputSize * static_cast<int64_t>(sizeof(float)));
327+ pipe_.InitBuffer(mergeInputIndicesBuffer_, tilingData_->mergeInputSize * static_cast<int64_t>(sizeof(int32_t)));
328+ pipe_.InitBuffer(mergeOutputScoresBuffer_, tilingData_->mergeOutputSize * static_cast<int64_t>(sizeof(float)));
329+ pipe_.InitBuffer(mergeOutputIndicesBuffer_, tilingData_->mergeOutputSize * static_cast<int64_t>(sizeof(int32_t)));
330+ pipe_.InitBuffer(mergeTopKTempBuffer_, static_cast<int64_t>(tilingData_->topKTempBytes));
331+ 
332+ scoreLocal_ = scoreBuffer_.Get<float>();
333+ yMinLocal_ = yMinBuffer_.Get<float>();
334+ xMinLocal_ = xMinBuffer_.Get<float>();
335+ yMaxLocal_ = yMaxBuffer_.Get<float>();
336+ xMaxLocal_ = xMaxBuffer_.Get<float>();
337+ temp0Local_ = temp0Buffer_.Get<float>();
338+ temp1Local_ = temp1Buffer_.Get<float>();
339+ temp2Local_ = temp2Buffer_.Get<float>();
340+ temp3Local_ = temp3Buffer_.Get<float>();
341+ reduceWorkLocal_ = reduceWorkBuffer_.Get<float>();
342+ reduceOutputLocal_ = reduceOutputBuffer_.Get<float>();
343+ compareMaskLocal_ = compareMaskBuffer_.Get<uint8_t>();
344+ mergeInputScoresLocal_ = mergeInputScoresBuffer_.Get<float>();
345+ mergeInputIndicesLocal_ = mergeInputIndicesBuffer_.Get<int32_t>();
346+ mergeOutputScoresLocal_ = mergeOutputScoresBuffer_.Get<float>();
347+ mergeOutputIndicesLocal_ = mergeOutputIndicesBuffer_.Get<int32_t>();
348+ mergeTopKTempLocal_ = mergeTopKTempBuffer_.Get<uint8_t>();
349+}
350+ 
351+template <typename T>
352+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::CopyIn(LocalTensor<float>& dst,
353+ const GlobalTensor<float>& src, int64_t offset,
354+ int64_t count)
355+{
356+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(count * static_cast<int64_t>(sizeof(float))), 0, 0, 0};
357+ DataCopyPadExtParams<float> padParams{false, 0, 0, 0};
358+ DataCopyPad(dst, src[offset], copyParams, padParams);
359+}
360+ 
361+template <typename T>
362+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::CopyOut(GlobalTensor<float>& dst, int64_t offset,
363+ LocalTensor<float>& src, int64_t count)
364+{
365+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(count * static_cast<int64_t>(sizeof(float))), 0, 0, 0};
366+ DataCopyPad(dst[offset], src, copyParams);
367+}
368+ 
369+template <typename T>
370+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::LoadTile(int64_t offset, int64_t count)
371+{
372+ CopyIn(scoreLocal_, stageScoresGm_, offset, count);
373+ CopyIn(yMinLocal_, stageYMinGm_, offset, count);
374+ CopyIn(xMinLocal_, stageXMinGm_, offset, count);
375+ CopyIn(yMaxLocal_, stageYMaxGm_, offset, count);
376+ CopyIn(xMaxLocal_, stageXMaxGm_, offset, count);
377+ const event_t eventMte2ToV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));
378+ SetFlag<HardEvent::MTE2_V>(eventMte2ToV);
379+ WaitFlag<HardEvent::MTE2_V>(eventMte2ToV);
380+}
381+ 
382+template <typename T>
383+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::StoreTile(int64_t offset, int64_t count)
384+{
385+ PipeBarrier<PIPE_ALL>();
386+ CopyOut(stageScoresGm_, offset, scoreLocal_, count);
387+ CopyOut(stageYMinGm_, offset, yMinLocal_, count);
388+ CopyOut(stageXMinGm_, offset, xMinLocal_, count);
389+ CopyOut(stageYMaxGm_, offset, yMaxLocal_, count);
390+ CopyOut(stageXMaxGm_, offset, xMaxLocal_, count);
391+ PipeBarrier<PIPE_ALL>();
392+}
393+ 
394+template <typename T>
395+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::LoadScores(int64_t offset, int64_t count)
396+{
397+ CopyIn(scoreLocal_, stageScoresGm_, offset, count);
398+ const event_t eventMte2ToV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));
399+ SetFlag<HardEvent::MTE2_V>(eventMte2ToV);
400+ WaitFlag<HardEvent::MTE2_V>(eventMte2ToV);
401+}
402+ 
403+template <typename T>
404+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::StoreScores(int64_t offset, int64_t count)
405+{
406+ PipeBarrier<PIPE_ALL>();
407+ CopyOut(stageScoresGm_, offset, scoreLocal_, count);
408+ PipeBarrier<PIPE_ALL>();
409+}
410+ 
411+template <typename T>
412+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::GatherClass(int64_t taskIndex)
413+{
414+ const int64_t batchIndex = taskIndex / tilingData_->classesNum;
415+ const int64_t classIndex = taskIndex % tilingData_->classesNum;
416+ const int64_t scratchOffset = taskIndex * tilingData_->scratchBytesPerCore;
417+ GM_ADDR scratchBase = userWorkspace_ + scratchOffset;
418+ const int64_t scratchElements = tilingData_->boxesNum;
419+ stageScoresGm_.SetGlobalBuffer((__gm__ float*)scratchBase, scratchElements);
420+ scratchBase += tilingData_->scratchFieldStride;
421+ stageYMinGm_.SetGlobalBuffer((__gm__ float*)scratchBase, scratchElements);
422+ scratchBase += tilingData_->scratchFieldStride;
423+ stageXMinGm_.SetGlobalBuffer((__gm__ float*)scratchBase, scratchElements);
424+ scratchBase += tilingData_->scratchFieldStride;
425+ stageYMaxGm_.SetGlobalBuffer((__gm__ float*)scratchBase, scratchElements);
426+ scratchBase += tilingData_->scratchFieldStride;
427+ stageXMaxGm_.SetGlobalBuffer((__gm__ float*)scratchBase, scratchElements);
428+ 
429+ if (tilingData_->transposeBox == 0U) {
430+ return;
431+ }
432+ 
433+ const int64_t boxClass = tilingData_->boxClassesNum == 1 ? 0 : classIndex;
434+ const int64_t transposeBoxBase = ((batchIndex * tilingData_->boxClassesNum + boxClass) * 4) * tilingData_->boxesNum;
435+ const __gm__ T* transposeYMin = (__gm__ T*)boxesAddr_ + transposeBoxBase;
436+ const __gm__ T* transposeXMin = transposeYMin + tilingData_->boxesNum;
437+ const __gm__ T* transposeYMax = transposeXMin + tilingData_->boxesNum;
438+ const __gm__ T* transposeXMax = transposeYMax + tilingData_->boxesNum;
439+ 
440+ if (tilingData_->use32Index != 0U) {
441+ asc_vf_call<GatherClassInput<T, true>>(
442+ dim3{kGatherThreadNum32}, transposeYMin, transposeXMin, transposeYMax, transposeXMax,
443+ (__gm__ T*)scoresAddr_, (__gm__ int32_t*)numValidBoxesAddr_, (__gm__ float*)stageScoresGm_.GetPhyAddr(),
444+ (__gm__ float*)stageYMinGm_.GetPhyAddr(), (__gm__ float*)stageXMinGm_.GetPhyAddr(),
445+ (__gm__ float*)stageYMaxGm_.GetPhyAddr(), (__gm__ float*)stageXMaxGm_.GetPhyAddr(),
446+ static_cast<uint32_t>(batchIndex), static_cast<uint32_t>(classIndex),
447+ static_cast<uint32_t>(tilingData_->boxesNum), static_cast<uint32_t>(tilingData_->classesNum),
448+ tilingData_->hasNumValidBoxes != 0U);
449+ } else {
450+ asc_vf_call<GatherClassInput<T, false>>(
451+ dim3{kGatherThreadNum64}, transposeYMin, transposeXMin, transposeYMax, transposeXMax,
452+ (__gm__ T*)scoresAddr_, (__gm__ int32_t*)numValidBoxesAddr_, (__gm__ float*)stageScoresGm_.GetPhyAddr(),
453+ (__gm__ float*)stageYMinGm_.GetPhyAddr(), (__gm__ float*)stageXMinGm_.GetPhyAddr(),
454+ (__gm__ float*)stageYMaxGm_.GetPhyAddr(), (__gm__ float*)stageXMaxGm_.GetPhyAddr(),
455+ static_cast<uint64_t>(batchIndex), static_cast<uint64_t>(classIndex),
456+ static_cast<uint64_t>(tilingData_->boxesNum), static_cast<uint64_t>(tilingData_->classesNum),
457+ tilingData_->hasNumValidBoxes != 0U);
458+ }
459+}
460+ 
461+template <typename T>
462+__aicore__ inline bool BatchMultiClassNonMaxSuppressionKernel<T>::FilterStageBoxes(int64_t batchIndex)
463+{
464+ float clipYMin = 0.0F;
465+ float clipXMin = 0.0F;
466+ float clipYMax = 0.0F;
467+ float clipXMax = 0.0F;
468+ if (tilingData_->hasClipWindow != 0U) {
469+ // A core commonly handles several classes from the same batch in
470+ // adjacent task waves. Keep these four scalar parameters in core
471+ // state so only the first class performs scalar GM reads.
472+ if (cachedClipBatch_ != batchIndex) {
473+ const int64_t clipOffset = batchIndex * 4;
474+ cachedClipYMin_ = static_cast<float>(clipWindowGm_.GetValue(clipOffset));
475+ cachedClipXMin_ = static_cast<float>(clipWindowGm_.GetValue(clipOffset + 1));
476+ cachedClipYMax_ = static_cast<float>(clipWindowGm_.GetValue(clipOffset + 2));
477+ cachedClipXMax_ = static_cast<float>(clipWindowGm_.GetValue(clipOffset + 3));
478+ cachedClipBatch_ = batchIndex;
479+ }
480+ clipYMin = cachedClipYMin_;
481+ clipXMin = cachedClipXMin_;
482+ clipYMax = cachedClipYMax_;
483+ clipXMax = cachedClipXMax_;
484+ if (tilingData_->changeCoordinateFrame != 0U && (clipYMax <= clipYMin || clipXMax <= clipXMin)) {
485+ return false;
486+ }
487+ }
488+ 
489+ for (int64_t offset = 0; offset < tilingData_->boxesNum; offset += tilingData_->tileSize) {
490+ const int64_t count = (tilingData_->boxesNum - offset) < tilingData_->tileSize ?
491+ (tilingData_->boxesNum - offset) :
492+ tilingData_->tileSize;
493+ LoadTile(offset, count);
494+ if (tilingData_->hasClipWindow != 0U) {
495+ Maxs(yMinLocal_, yMinLocal_, clipYMin, count);
496+ Maxs(xMinLocal_, xMinLocal_, clipXMin, count);
497+ Mins(yMaxLocal_, yMaxLocal_, clipYMax, count);
498+ Mins(xMaxLocal_, xMaxLocal_, clipXMax, count);
499+ if (tilingData_->changeCoordinateFrame != 0U) {
500+ Adds(yMinLocal_, yMinLocal_, -clipYMin, count);
501+ Adds(yMaxLocal_, yMaxLocal_, -clipYMin, count);
502+ Adds(xMinLocal_, xMinLocal_, -clipXMin, count);
503+ Adds(xMaxLocal_, xMaxLocal_, -clipXMin, count);
504+ Muls(yMinLocal_, yMinLocal_, 1.0F / (clipYMax - clipYMin), count);
505+ Muls(yMaxLocal_, yMaxLocal_, 1.0F / (clipYMax - clipYMin), count);
506+ Muls(xMinLocal_, xMinLocal_, 1.0F / (clipXMax - clipXMin), count);
507+ Muls(xMaxLocal_, xMaxLocal_, 1.0F / (clipXMax - clipXMin), count);
508+ }
509+ }
510+ CompareScalar(compareMaskLocal_, scoreLocal_, tilingData_->scoreThreshold, CMPMODE::GT, count);
511+ Select(scoreLocal_, compareMaskLocal_, scoreLocal_, kNoCandidate, SELMODE::VSEL_TENSOR_SCALAR_MODE, count);
512+ Compare(compareMaskLocal_, yMaxLocal_, yMinLocal_, CMPMODE::GT, count);
513+ Select(scoreLocal_, compareMaskLocal_, scoreLocal_, kNoCandidate, SELMODE::VSEL_TENSOR_SCALAR_MODE, count);
514+ Compare(compareMaskLocal_, xMaxLocal_, xMinLocal_, CMPMODE::GT, count);
515+ Select(scoreLocal_, compareMaskLocal_, scoreLocal_, kNoCandidate, SELMODE::VSEL_TENSOR_SCALAR_MODE, count);
516+ StoreTile(offset, count);
517+ }
518+ return true;
519+}
520+ 
521+template <typename T>
522+__aicore__ inline bool BatchMultiClassNonMaxSuppressionKernel<T>::FindBestCandidate(int64_t& bestIndex,
523+ float& bestScore)
524+{
525+ bestIndex = -1;
526+ bestScore = kNoCandidate;
527+ for (int64_t offset = 0; offset < tilingData_->boxesNum; offset += tilingData_->tileSize) {
528+ const int64_t count = (tilingData_->boxesNum - offset) < tilingData_->tileSize ?
529+ (tilingData_->boxesNum - offset) :
530+ tilingData_->tileSize;
531+ LoadScores(offset, count);
532+ ReduceMax<float>(reduceOutputLocal_, scoreLocal_, reduceWorkLocal_, static_cast<int32_t>(count), true);
533+ PipeBarrier<PIPE_V>();
534+ const event_t eventVS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_S));
535+ SetFlag<HardEvent::V_S>(eventVS);
536+ WaitFlag<HardEvent::V_S>(eventVS);
537+ const float tileScore = reduceOutputLocal_.GetValue(0);
538+ // ReduceMax's index result is invalid for this Ascend 950 kernel.
539+ // Its value result is reliable: reduce matching negative indices with
540+ // it, so the maximum gives the first matching index on a tied score.
541+ CompareScalar(compareMaskLocal_, scoreLocal_, tileScore, CMPMODE::EQ, count);
542+ ArithProgression<float>(temp0Local_, 0.0F, -1.0F, static_cast<int32_t>(count));
543+ Select(temp0Local_, compareMaskLocal_, temp0Local_, -static_cast<float>(count),
544+ SELMODE::VSEL_TENSOR_SCALAR_MODE, count);
545+ // The mask/select chain and the following reduction share the V
546+ // pipeline. Make the data dependency explicit on 950 before
547+ // consuming the selected negative indices.
548+ PipeBarrier<PIPE_V>();
549+ ReduceMax<float>(reduceOutputLocal_, temp0Local_, reduceWorkLocal_, static_cast<int32_t>(count), false);
550+ PipeBarrier<PIPE_V>();
551+ const event_t eventIndexVS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_S));
552+ SetFlag<HardEvent::V_S>(eventIndexVS);
553+ WaitFlag<HardEvent::V_S>(eventIndexVS);
554+ const int64_t tileIndex = -static_cast<int64_t>(reduceOutputLocal_.GetValue(0));
555+ const int64_t candidateIndex = offset + tileIndex;
556+ if (tileScore > bestScore || (tileScore == bestScore && (bestIndex < 0 || candidateIndex < bestIndex))) {
557+ bestScore = tileScore;
558+ bestIndex = candidateIndex;
559+ }
560+ }
561+ return bestIndex >= 0 && bestScore > tilingData_->scoreThreshold;
562+}
563+ 
564+template <typename T>
565+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::SuppressBySelectedBox(float selectedYMin,
566+ float selectedXMin,
567+ float selectedYMax,
568+ float selectedXMax)
569+{
570+ const float selectedArea = (selectedYMax - selectedYMin) * (selectedXMax - selectedXMin);
571+ for (int64_t offset = 0; offset < tilingData_->boxesNum; offset += tilingData_->tileSize) {
572+ const int64_t count = (tilingData_->boxesNum - offset) < tilingData_->tileSize ?
573+ (tilingData_->boxesNum - offset) :
574+ tilingData_->tileSize;
575+ LoadTile(offset, count);
576+ Maxs(temp0Local_, yMinLocal_, selectedYMin, count);
577+ Mins(temp1Local_, yMaxLocal_, selectedYMax, count);
578+ Sub(temp0Local_, temp1Local_, temp0Local_, count);
579+ Maxs(temp0Local_, temp0Local_, 0.0F, count);
580+ Maxs(temp1Local_, xMinLocal_, selectedXMin, count);
581+ Mins(temp2Local_, xMaxLocal_, selectedXMax, count);
582+ Sub(temp1Local_, temp2Local_, temp1Local_, count);
583+ Maxs(temp1Local_, temp1Local_, 0.0F, count);
584+ Mul(temp2Local_, temp0Local_, temp1Local_, count);
585+ Sub(temp0Local_, yMaxLocal_, yMinLocal_, count);
586+ Sub(temp1Local_, xMaxLocal_, xMinLocal_, count);
587+ Mul(temp3Local_, temp0Local_, temp1Local_, count);
588+ Adds(temp3Local_, temp3Local_, selectedArea, count);
589+ Sub(temp3Local_, temp3Local_, temp2Local_, count);
590+ Maxs(temp3Local_, temp3Local_, kMinPositive, count);
591+ Div(temp2Local_, temp2Local_, temp3Local_, count);
592+ CompareScalar(compareMaskLocal_, temp2Local_, tilingData_->iouThreshold, CMPMODE::GT, count);
593+ Duplicate(temp3Local_, kNoCandidate, count);
594+ Select(scoreLocal_, compareMaskLocal_, temp3Local_, scoreLocal_, SELMODE::VSEL_TENSOR_TENSOR_MODE, count);
595+ StoreScores(offset, count);
596+ }
597+}
598+ 
599+template <typename T>
600+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::ProcessClass(int64_t taskIndex)
601+{
602+ const int64_t classResultBase = taskIndex * tilingData_->maxSizePerClass;
603+ const int64_t batchIndex = taskIndex / tilingData_->classesNum;
604+ if (!FilterStageBoxes(batchIndex)) {
605+ PadClassScores(taskIndex, 0);
606+ StoreClassCount(taskIndex, 0);
607+ return;
608+ }
609+ int64_t selectedCount = 0;
610+ while (selectedCount < tilingData_->maxSizePerClass) {
611+ int64_t bestIndex = -1;
612+ float bestScore = kNoCandidate;
613+ if (!FindBestCandidate(bestIndex, bestScore)) {
614+ break;
615+ }
616+ const float yMin = stageYMinGm_.GetValue(bestIndex);
617+ const float xMin = stageXMinGm_.GetValue(bestIndex);
618+ const float yMax = stageYMaxGm_.GetValue(bestIndex);
619+ const float xMax = stageXMaxGm_.GetValue(bestIndex);
620+ const int64_t resultOffset = classResultBase + selectedCount;
621+ temp0Local_.SetValue(0, yMin);
622+ temp0Local_.SetValue(1, xMin);
623+ temp0Local_.SetValue(2, yMax);
624+ temp0Local_.SetValue(3, xMax);
625+ PipeBarrier<PIPE_ALL>();
626+ CopyOut(classBoxesGm_, resultOffset * 4, temp0Local_, 4);
627+ PipeBarrier<PIPE_ALL>();
628+ // Scalar GM writes are not visible to the MTE2 class-score merge on
629+ // Ascend 950. Materialize the selected score through the vector
630+ // pipeline and publish it with MTE3, as PadClassScores does for the
631+ // sentinel tail of this class.
632+ Duplicate(temp0Local_, bestScore, 1);
633+ const event_t eventVToMte3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));
634+ SetFlag<HardEvent::V_MTE3>(eventVToMte3);
635+ WaitFlag<HardEvent::V_MTE3>(eventVToMte3);
636+ CopyOut(classScoresGm_, resultOffset, temp0Local_, 1);
637+ PipeBarrier<PIPE_ALL>();
638+ // Publish removal of the selected candidate through MTE3. A scalar
639+ // GM store is invisible to the subsequent MTE2 score reload on 950,
640+ // which otherwise selects this same candidate repeatedly.
641+ Duplicate(temp0Local_, kNoCandidate, 1);
642+ const event_t invalidateVToMte3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));
643+ SetFlag<HardEvent::V_MTE3>(invalidateVToMte3);
644+ WaitFlag<HardEvent::V_MTE3>(invalidateVToMte3);
645+ CopyOut(stageScoresGm_, bestIndex, temp0Local_, 1);
646+ PipeBarrier<PIPE_ALL>();
647+ SuppressBySelectedBox(yMin, xMin, yMax, xMax);
648+ ++selectedCount;
649+ }
650+ PadClassScores(taskIndex, selectedCount);
651+ StoreClassCount(taskIndex, selectedCount);
652+}
653+ 
654+template <typename T>
655+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::PadClassScores(int64_t taskIndex,
656+ int64_t selectedCount)
657+{
658+ const int64_t classResultBase = taskIndex * tilingData_->maxSizePerClass;
659+ for (int64_t offset = selectedCount; offset < tilingData_->maxSizePerClass; offset += tilingData_->tileSize) {
660+ const int64_t count = (tilingData_->maxSizePerClass - offset) < tilingData_->tileSize ?
661+ (tilingData_->maxSizePerClass - offset) :
662+ tilingData_->tileSize;
663+ Duplicate(scoreLocal_, kNoCandidate, count);
664+ PipeBarrier<PIPE_ALL>();
665+ CopyOut(classScoresGm_, classResultBase + offset, scoreLocal_, count);
666+ PipeBarrier<PIPE_ALL>();
667+ }
668+}
669+ 
670+template <typename T>
671+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::StoreClassCount(int64_t taskIndex,
672+ int64_t selectedCount)
673+{
674+ // Publish the count through the vector/MTE3 path used for class scores.
675+ // A scalar LocalTensor write is not visible to a later DMA consumer on
676+ // Ascend 950, which would make the merge use stale workspace content.
677+ Duplicate(scoreLocal_, static_cast<float>(selectedCount), 1);
678+ const event_t eventVToMte3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));
679+ SetFlag<HardEvent::V_MTE3>(eventVToMte3);
680+ WaitFlag<HardEvent::V_MTE3>(eventVToMte3);
681+ CopyOut(classCountsGm_, taskIndex, scoreLocal_, 1);
682+ PipeBarrier<PIPE_ALL>();
683+}
684+ 
685+template <typename T>
686+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::MergeBatch(int64_t batchIndex)
687+{
688+ const int64_t classResultBase = batchIndex * tilingData_->classesNum * tilingData_->maxSizePerClass;
689+ const int64_t mergeInputCapacity = tilingData_->mergeInputSize;
690+ int64_t selectedAcrossClasses = 0;
691+ const int64_t batchTaskBase = batchIndex * tilingData_->classesNum;
692+ for (int64_t classIndex = 0; classIndex < tilingData_->classesNum; ++classIndex) {
693+ selectedAcrossClasses += static_cast<int64_t>(classCountsGm_.GetValue(batchTaskBase + classIndex));
694+ }
695+ const int64_t mergeOutputCount = selectedAcrossClasses < tilingData_->maxTotalSize ? selectedAcrossClasses :
696+ tilingData_->maxTotalSize;
697+ const int64_t mergeActiveSize = tilingData_->mergeOutputSize;
698+ int64_t activeCount = 0;
699+ int64_t sourceOffset = 0;
700+ while (mergeOutputCount > 0 && sourceOffset < tilingData_->mergeInputCount) {
701+ // Start every round with a fully aligned sentinel fill. The physical
702+ // carry length is mergeActiveSize, while only mergeOutputCount entries
703+ // are real TopK results. Filling from offset zero avoids a vector
704+ // write at an unaligned k boundary (for example k = 100).
705+ Duplicate(mergeInputScoresLocal_, kNoCandidate, mergeInputCapacity);
706+ if (activeCount > 0) {
707+ // Existing candidates are from earlier class-major positions. Put
708+ // them before the next chunk so TopK's stable tie handling keeps
709+ // the original lower-index-first ordering.
710+ Copy(mergeInputScoresLocal_, mergeOutputScoresLocal_, static_cast<uint32_t>(mergeOutputCount));
711+ Copy(mergeInputIndicesLocal_, mergeOutputIndicesLocal_, static_cast<uint32_t>(mergeOutputCount));
712+ }
713+ const int64_t inputCapacityLeft = mergeInputCapacity - activeCount;
714+ const int64_t sourceCountLeft = tilingData_->mergeInputCount - sourceOffset;
715+ const int64_t chunkCount = inputCapacityLeft < sourceCountLeft ? inputCapacityLeft : sourceCountLeft;
716+ const int64_t inputCount = activeCount + chunkCount;
717+ // The DMA writes only the current chunk, after any vector copy/fill
718+ // that touches the same local input buffers has completed.
719+ const event_t eventVToMte2 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE2));
720+ SetFlag<HardEvent::V_MTE2>(eventVToMte2);
721+ WaitFlag<HardEvent::V_MTE2>(eventVToMte2);
722+ DataCopyExtParams copyInParams{1, static_cast<uint32_t>(chunkCount * static_cast<int64_t>(sizeof(float))), 0, 0,
723+ 0};
724+ DataCopyPadExtParams<float> copyInPadParams{false, 0, 0, 0};
725+ DataCopyPad(mergeInputScoresLocal_[activeCount], classScoresGm_[classResultBase + sourceOffset], copyInParams,
726+ copyInPadParams);
727+ const event_t eventMte2V = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));
728+ SetFlag<HardEvent::MTE2_V>(eventMte2V);
729+ WaitFlag<HardEvent::MTE2_V>(eventMte2V);
730+ ArithProgression<int32_t>(mergeInputIndicesLocal_[activeCount], static_cast<int32_t>(sourceOffset), 1,
731+ static_cast<int32_t>(chunkCount));
732+ PipeBarrier<PIPE_V>();
733+ 
734+ TopKInfo topKInfo{1, static_cast<int32_t>(mergeInputCapacity), static_cast<int32_t>(inputCount)};
735+ LocalTensor<bool> emptyFinishLocal;
736+ TopK<float, true, false, false, TopKMode::TOPK_NORMAL, kMergeTopKConfig>(
737+ mergeOutputScoresLocal_, mergeOutputIndicesLocal_, mergeInputScoresLocal_, mergeInputIndicesLocal_,
738+ emptyFinishLocal, mergeTopKTempLocal_, static_cast<int32_t>(mergeOutputCount), tilingData_->mergeTopKTiling,
739+ topKInfo, true);
740+ PipeBarrier<PIPE_V>();
741+ activeCount = mergeActiveSize;
742+ sourceOffset += chunkCount;
743+ }
744+ const int64_t outputOffset = batchIndex * tilingData_->maxTotalSize;
745+ PipeBarrier<PIPE_ALL>();
746+ if (mergeOutputCount > 0) {
747+ DataCopyExtParams copyOutParams{
748+ 1, static_cast<uint32_t>(mergeOutputCount * static_cast<int64_t>(sizeof(float))), 0, 0, 0};
749+ DataCopyPad(mergeScoresGm_[outputOffset], mergeOutputScoresLocal_, copyOutParams);
750+ DataCopyPad(mergeIndicesGm_[outputOffset], mergeOutputIndicesLocal_, copyOutParams);
751+ }
752+ // Never ask TopK to rank padded scores. Materialize the unwritten tail
753+ // directly in GM so the SIMT gather neither counts it nor reads its index.
754+ for (int64_t tailOffset = mergeOutputCount; tailOffset < tilingData_->maxTotalSize;
755+ tailOffset += tilingData_->tileSize) {
756+ const int64_t tailCount = (tilingData_->maxTotalSize - tailOffset) < tilingData_->tileSize ?
757+ (tilingData_->maxTotalSize - tailOffset) :
758+ tilingData_->tileSize;
759+ Duplicate(scoreLocal_, kNoCandidate, tailCount);
760+ PipeBarrier<PIPE_ALL>();
761+ CopyOut(mergeScoresGm_, outputOffset + tailOffset, scoreLocal_, tailCount);
762+ }
763+ PipeBarrier<PIPE_ALL>();
764+ 
765+ asc_vf_call<GatherMergedOutput<T>>(
766+ dim3{kMergeThreadNum}, (__gm__ float*)classBoxesGm_.GetPhyAddr(), (__gm__ float*)mergeScoresGm_.GetPhyAddr(),
767+ (__gm__ int32_t*)mergeIndicesGm_.GetPhyAddr(), (__gm__ T*)nmsedBoxesGm_.GetPhyAddr(),
768+ (__gm__ T*)nmsedScoresGm_.GetPhyAddr(), (__gm__ T*)nmsedClassesGm_.GetPhyAddr(),
769+ (__gm__ int32_t*)nmsedNumGm_.GetPhyAddr(), static_cast<uint64_t>(batchIndex),
770+ static_cast<uint64_t>(tilingData_->classesNum), static_cast<uint64_t>(tilingData_->maxSizePerClass),
771+ static_cast<uint64_t>(tilingData_->maxTotalSize), static_cast<uint64_t>(mergeOutputCount));
772+}
773+ 
774+template <typename T>
775+__aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::Process()
776+{
777+ if (!workspaceReady_ || static_cast<int64_t>(GetBlockIdx()) >= tilingData_->usedCoreNum) {
778+ return;
779+ }
780+ const int64_t taskCount = tilingData_->batch * tilingData_->classesNum;
781+ if (tilingData_->transposeBox == 0U) {
782+ if (tilingData_->use32Index != 0U) {
783+ asc_vf_call<GatherAllClassInputs<T, true>>(
784+ dim3{kGatherThreadNum32}, (__gm__ T*)boxesAddr_, (__gm__ T*)scoresAddr_,
785+ (__gm__ int32_t*)numValidBoxesAddr_, (__gm__ float*)userWorkspace_,
786+ static_cast<uint32_t>(GetBlockIdx()), static_cast<uint32_t>(tilingData_->usedCoreNum),
787+ static_cast<uint32_t>(taskCount), static_cast<uint32_t>(tilingData_->boxesNum),
788+ static_cast<uint32_t>(tilingData_->classesNum), static_cast<uint32_t>(tilingData_->boxClassesNum),
789+ tilingData_->scratchFieldStride, tilingData_->hasNumValidBoxes != 0U);
790+ } else {
791+ asc_vf_call<GatherAllClassInputs<T, false>>(
792+ dim3{kGatherThreadNum64}, (__gm__ T*)boxesAddr_, (__gm__ T*)scoresAddr_,
793+ (__gm__ int32_t*)numValidBoxesAddr_, (__gm__ float*)userWorkspace_,
794+ static_cast<uint64_t>(GetBlockIdx()), static_cast<uint64_t>(tilingData_->usedCoreNum),
795+ static_cast<uint64_t>(taskCount), static_cast<uint64_t>(tilingData_->boxesNum),
796+ static_cast<uint64_t>(tilingData_->classesNum), static_cast<uint64_t>(tilingData_->boxClassesNum),
797+ tilingData_->scratchFieldStride, tilingData_->hasNumValidBoxes != 0U);
798+ }
799+ // All cores have launched their persistent gather task before any
800+ // class enters the AIV NMS stage.
801+ SyncAll();
802+ }
803+ for (int64_t taskBase = 0; taskBase < taskCount; taskBase += tilingData_->usedCoreNum) {
804+ const int64_t taskIndex = taskBase + static_cast<int64_t>(GetBlockIdx());
805+ const bool hasTask = taskIndex < taskCount;
806+ if (hasTask) {
807+ GatherClass(taskIndex);
808+ }
809+ // asc_vf_call is synchronized with the AIV pipeline at the task-wave
810+ // boundary. All blocks enter the barrier, including the last partial
811+ // wave, so this remains valid when batch * class is not core-aligned.
812+ SyncAll();
813+ if (hasTask) {
814+ ProcessClass(taskIndex);
815+ }
816+ SyncAll();
817+ }
818+ for (int64_t batchBase = 0; batchBase < tilingData_->batch; batchBase += tilingData_->usedCoreNum) {
819+ const int64_t batchIndex = batchBase + static_cast<int64_t>(GetBlockIdx());
820+ if (batchIndex < tilingData_->batch) {
821+ MergeBatch(batchIndex);
822+ }
823+ // Keep the AIV->SIMT hand-off collective. All blocks enter this
824+ // barrier, including inactive tail blocks, before the next batch wave.
825+ SyncAll();
826+ }
827+}
828+} // namespace BatchMultiClassNonMaxSuppressionOp
829+ 
830+#endif // BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_KERNEL_H_
Aobjdetect/batch_multi_class_non_max_suppression/op_kernel/arch35/batch_multi_class_non_max_suppression_tiling_data.h+49-0
@@ -0,0 +1,49 @@
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+#ifndef BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_DATA_H_
12+#define BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_DATA_H_
13+ 
14+#include <cstdint>
15+#include "kernel_tiling/kernel_tiling.h"
16+ 
17+struct BatchMultiClassNonMaxSuppressionTilingData {
18+ int64_t batch{0};
19+ int64_t boxesNum{0};
20+ int64_t classesNum{0};
21+ int64_t boxClassesNum{0};
22+ int64_t maxSizePerClass{0};
23+ int64_t maxTotalSize{0};
24+ int64_t usedCoreNum{0};
25+ int64_t tileSize{0};
26+ int64_t reduceBufferSize{0};
27+ int64_t mergeInputCount{0};
28+ int64_t mergeInputSize{0};
29+ int64_t mergeOutputCount{0};
30+ int64_t mergeOutputSize{0};
31+ uint64_t scratchFieldStride{0};
32+ uint64_t scratchBytesPerCore{0};
33+ uint64_t classBoxesOffset{0};
34+ uint64_t classScoresOffset{0};
35+ uint64_t classCountsOffset{0};
36+ uint64_t mergeScoresOffset{0};
37+ uint64_t mergeIndicesOffset{0};
38+ uint64_t topKTempBytes{0};
39+ AscendC::tiling::TopkTiling mergeTopKTiling{};
40+ float scoreThreshold{0.0F};
41+ float iouThreshold{0.0F};
42+ uint8_t hasClipWindow{0};
43+ uint8_t hasNumValidBoxes{0};
44+ uint8_t changeCoordinateFrame{0};
45+ uint8_t transposeBox{0};
46+ uint8_t use32Index{0};
47+};
48+ 
49+#endif // BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_DATA_H_
Aobjdetect/batch_multi_class_non_max_suppression/op_kernel/arch35/batch_multi_class_non_max_suppression_tiling_key.h+19-0
@@ -0,0 +1,19 @@
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+#ifndef BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_KEY_H_
12+#define BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_KEY_H_
13+ 
14+// TILING_KEY_IS consumes this value in the preprocessor, so it must remain a
15+// numeric macro rather than a C++ constexpr. The host tiler includes this
16+// header too, keeping the generated binary and runtime tiling key aligned.
17+#define BATCH_MULTI_CLASS_NMS_TILING_KEY 10000UL
18+ 
19+#endif // BATCH_MULTI_CLASS_NON_MAX_SUPPRESSION_TILING_KEY_H_
Aobjdetect/batch_multi_class_non_max_suppression/op_kernel/batch_multi_class_non_max_suppression_apt.cpp+34-0
@@ -0,0 +1,34 @@
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 "arch35/batch_multi_class_non_max_suppression_kernel.h"
12+#include "arch35/batch_multi_class_non_max_suppression_tiling_key.h"
13+ 
14+extern "C" __global__ __aicore__ void batch_multi_class_non_max_suppression(
15+ GM_ADDR boxes, GM_ADDR scores, GM_ADDR clip_window, GM_ADDR num_valid_boxes, GM_ADDR nmsed_boxes,
16+ GM_ADDR nmsed_scores, GM_ADDR nmsed_classes, GM_ADDR nmsed_num, GM_ADDR workspace, GM_ADDR tiling)
17+{
18+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0);
19+ if (workspace == nullptr) {
20+ return;
21+ }
22+ AscendC::SetSysWorkspace(workspace);
23+ if (AscendC::GetUserWorkspace(workspace) == nullptr) {
24+ return;
25+ }
26+ REGISTER_TILING_DEFAULT(BatchMultiClassNonMaxSuppressionTilingData);
27+ GET_TILING_DATA_WITH_STRUCT(BatchMultiClassNonMaxSuppressionTilingData, tilingData, tiling);
28+ if (TILING_KEY_IS(BATCH_MULTI_CLASS_NMS_TILING_KEY)) {
29+ BatchMultiClassNonMaxSuppressionOp::BatchMultiClassNonMaxSuppressionKernel<DTYPE_BOXES> op;
30+ op.Init(boxes, scores, clip_window, num_valid_boxes, nmsed_boxes, nmsed_scores, nmsed_classes, nmsed_num,
31+ workspace, &tilingData);
32+ op.Process();
33+ }
34+}
Aobjdetect/batch_multi_class_non_max_suppression/tests/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
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+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT (UT_TEST_ALL OR OP_HOST_UT))
13+ list(REMOVE_ITEM CURRENT_DIRS ut)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aobjdetect/batch_multi_class_non_max_suppression/tests/assets/golden.py+398-0
@@ -0,0 +1,398 @@
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+ 
13+import torch
14+import numpy
15+ 
16+__spec__ = {
17+ "batch_multi_class_non_max_suppression": "BatchMultiClassNonMaxSuppressionTestSpec",
18+}
19+ 
20+ 
21+def _compute_iou_vectorized(boxes_a, box_b):
22+ iy_min = torch.max(boxes_a[:, 0], box_b[0])
23+ ix_min = torch.max(boxes_a[:, 1], box_b[1])
24+ iy_max = torch.min(boxes_a[:, 2], box_b[2])
25+ ix_max = torch.min(boxes_a[:, 3], box_b[3])
26+ inter_h = torch.clamp(iy_max - iy_min, min=0.0)
27+ inter_w = torch.clamp(ix_max - ix_min, min=0.0)
28+ inter_area = inter_h * inter_w
29+ area_a = (boxes_a[:, 2] - boxes_a[:, 0]) * (boxes_a[:, 3] - boxes_a[:, 1])
30+ area_b = (box_b[2] - box_b[0]) * (box_b[3] - box_b[1])
31+ union_area = torch.clamp(area_a + area_b - inter_area, min=1e-12)
32+ return inter_area / union_area
33+ 
34+ 
35+def _nms_single_class(
36+ class_boxes, class_scores, score_threshold, iou_threshold, max_size
37+):
38+ valid_mask = class_scores > score_threshold
39+ valid_mask &= class_boxes[:, 2] > class_boxes[:, 0]
40+ valid_mask &= class_boxes[:, 3] > class_boxes[:, 1]
41+ valid_indices = torch.where(valid_mask)[0]
42+ if len(valid_indices) == 0:
43+ return []
44+ valid_scores = class_scores[valid_indices]
45+ valid_boxes = class_boxes[valid_indices]
46+ sorted_order = torch.argsort(valid_scores, descending=True, stable=True)
47+ sorted_scores = valid_scores[sorted_order]
48+ sorted_boxes = valid_boxes[sorted_order]
49+ num_candidates = len(sorted_scores)
50+ suppressed = torch.zeros(num_candidates, dtype=torch.bool)
51+ kept = []
52+ for i in range(num_candidates):
53+ if suppressed[i]:
54+ continue
55+ kept.append((sorted_scores[i].item(), sorted_boxes[i].clone()))
56+ if len(kept) >= max_size:
57+ break
58+ if i < num_candidates - 1:
59+ rem_indices = torch.where(~suppressed)[0]
60+ rem_indices = rem_indices[rem_indices > i]
61+ if len(rem_indices) == 0:
62+ continue
63+ iou = _compute_iou_vectorized(sorted_boxes[rem_indices], sorted_boxes[i])
64+ suppressed[rem_indices[iou > iou_threshold]] = True
65+ return kept
66+ 
67+ 
68+def _golden_impl(
69+ boxes,
70+ scores,
71+ clip_window,
72+ num_valid_boxes,
73+ score_threshold,
74+ iou_threshold,
75+ max_size_per_class,
76+ max_total_size,
77+ change_coordinate_frame,
78+ transpose_box,
79+):
80+ out_dtype = boxes.dtype
81+ batch_size = scores.shape[0]
82+ num_classes = scores.shape[2]
83+ out_boxes = torch.zeros(batch_size, max_total_size, 4, dtype=torch.float32)
84+ out_scores = torch.zeros(batch_size, max_total_size, dtype=torch.float32)
85+ out_classes = torch.zeros(batch_size, max_total_size, dtype=torch.float32)
86+ out_num = torch.zeros(batch_size, dtype=torch.int32)
87+ if transpose_box:
88+ num_boxes = boxes.shape[3]
89+ box_classes = boxes.shape[1]
90+ else:
91+ num_boxes = boxes.shape[1]
92+ box_classes = boxes.shape[2]
93+ for b in range(batch_size):
94+ valid_count = num_boxes
95+ if num_valid_boxes is not None:
96+ nvb = int(num_valid_boxes[b].item())
97+ valid_count = max(0, min(nvb, num_boxes))
98+ has_clip = clip_window is not None
99+ if has_clip:
100+ cw = clip_window[b].float()
101+ cy_min, cx_min, cy_max, cx_max = (
102+ cw[0].item(),
103+ cw[1].item(),
104+ cw[2].item(),
105+ cw[3].item(),
106+ )
107+ if change_coordinate_frame and (cy_max <= cy_min or cx_max <= cx_min):
108+ continue
109+ candidates = []
110+ for c in range(num_classes):
111+ class_scores = scores[b, :, c].float().clone()
112+ if valid_count < num_boxes:
113+ class_scores[valid_count:] = float("-inf")
114+ box_class_idx = 0 if box_classes == 1 else c
115+ if transpose_box:
116+ class_boxes = (
117+ boxes[b, box_class_idx, :, :].float().T.contiguous().clone()
118+ )
119+ else:
120+ class_boxes = boxes[b, :, box_class_idx, :].float().clone()
121+ if has_clip:
122+ class_boxes[:, 0] = torch.clamp(class_boxes[:, 0], min=cy_min)
123+ class_boxes[:, 1] = torch.clamp(class_boxes[:, 1], min=cx_min)
124+ class_boxes[:, 2] = torch.clamp(class_boxes[:, 2], max=cy_max)
125+ class_boxes[:, 3] = torch.clamp(class_boxes[:, 3], max=cx_max)
126+ if change_coordinate_frame:
127+ win_h = cy_max - cy_min
128+ win_w = cx_max - cx_min
129+ class_boxes[:, 0] = (class_boxes[:, 0] - cy_min) / win_h
130+ class_boxes[:, 2] = (class_boxes[:, 2] - cy_min) / win_h
131+ class_boxes[:, 1] = (class_boxes[:, 1] - cx_min) / win_w
132+ class_boxes[:, 3] = (class_boxes[:, 3] - cx_min) / win_w
133+ kept = _nms_single_class(
134+ class_boxes,
135+ class_scores,
136+ score_threshold,
137+ iou_threshold,
138+ max_size_per_class,
139+ )
140+ for score_val, box_val in kept:
141+ candidates.append((score_val, c, box_val))
142+ candidates.sort(key=lambda x: -x[0])
143+ num_valid = min(len(candidates), max_total_size)
144+ out_num[b] = num_valid
145+ for i in range(num_valid):
146+ score_val, class_id, box_val = candidates[i]
147+ out_scores[b, i] = score_val
148+ out_classes[b, i] = float(class_id)
149+ out_boxes[b, i] = box_val
150+ return [
151+ out_boxes.to(out_dtype),
152+ out_scores.to(out_dtype),
153+ out_classes.to(out_dtype),
154+ out_num,
155+ ]
156+ 
157+ 
158+def _to_numpy(arr):
159+ """Convert torch tensor or numpy array to numpy."""
160+ if isinstance(arr, torch.Tensor):
161+ return arr.detach().cpu().numpy()
162+ return numpy.asarray(arr)
163+ 
164+ 
165+def _to_torch(arr):
166+ """Adapt kernel-direct numpy inputs to torch tensors."""
167+ if arr is None or isinstance(arr, torch.Tensor):
168+ return arr
169+ return torch.from_numpy(numpy.asarray(arr))
170+ 
171+ 
172+def _numpy_iou_matrix(boxes_a, boxes_b):
173+ """Compute IoU matrix between two sets of boxes [N,4] and [M,4] in yxyx format."""
174+ n = boxes_a.shape[0]
175+ m = boxes_b.shape[0]
176+ if n == 0 or m == 0:
177+ return numpy.zeros((n, m), dtype=numpy.float64)
178+ iy_min = numpy.maximum(boxes_a[:, 0:1], boxes_b[:, 0:1].T)
179+ ix_min = numpy.maximum(boxes_a[:, 1:2], boxes_b[:, 1:2].T)
180+ iy_max = numpy.minimum(boxes_a[:, 2:3], boxes_b[:, 2:3].T)
181+ ix_max = numpy.minimum(boxes_a[:, 3:4], boxes_b[:, 3:4].T)
182+ inter_h = numpy.clip(iy_max - iy_min, 0, None)
183+ inter_w = numpy.clip(ix_max - ix_min, 0, None)
184+ inter_area = inter_h * inter_w
185+ area_a = (boxes_a[:, 2] - boxes_a[:, 0]) * (boxes_a[:, 3] - boxes_a[:, 1])
186+ area_b = (boxes_b[:, 2] - boxes_b[:, 0]) * (boxes_b[:, 3] - boxes_b[:, 1])
187+ union_area = numpy.clip(area_a[:, None] + area_b[None, :] - inter_area, 1e-12, None)
188+ return inter_area / union_area
189+ 
190+ 
191+class BatchMultiClassNonMaxSuppressionTestSpec:
192+ def golden(
193+ boxes,
194+ scores,
195+ clip_window=None,
196+ num_valid_boxes=None,
197+ score_threshold=0.5,
198+ iou_threshold=0.5,
199+ max_size_per_class=100,
200+ max_total_size=100,
201+ change_coordinate_frame=False,
202+ transpose_box=False,
203+ **kwargs,
204+ ):
205+ is_torch_input = isinstance(boxes, torch.Tensor)
206+ results = _golden_impl(
207+ _to_torch(boxes),
208+ _to_torch(scores),
209+ _to_torch(clip_window),
210+ _to_torch(num_valid_boxes),
211+ float(score_threshold),
212+ float(iou_threshold),
213+ int(max_size_per_class),
214+ int(max_total_size),
215+ bool(change_coordinate_frame),
216+ bool(transpose_box),
217+ )
218+ if is_torch_input:
219+ return results
220+ return [result.cpu().numpy() for result in results]
221+ 
222+ def pre_compare(*outputs, **kwargs):
223+ """Sort detections by score descending in both NPU and golden outputs.
224+ 
225+ Layout: func(*npu_outputs, *golden_outputs)
226+ outputs[0]=npu_boxes, outputs[1]=npu_scores, outputs[2]=npu_classes, outputs[3]=npu_num
227+ outputs[4]=golden_boxes, outputs[5]=golden_scores, outputs[6]=golden_classes, outputs[7]=golden_num
228+ In-place mode: modify via [:], return None.
229+ """
230+ npu_boxes = _to_numpy(outputs[0])
231+ npu_scores = _to_numpy(outputs[1])
232+ npu_classes = _to_numpy(outputs[2])
233+ npu_num = _to_numpy(outputs[3])
234+ golden_boxes = _to_numpy(outputs[4])
235+ golden_scores = _to_numpy(outputs[5])
236+ golden_classes = _to_numpy(outputs[6])
237+ golden_num = _to_numpy(outputs[7])
238+ 
239+ batch_size = npu_boxes.shape[0]
240+ for b in range(batch_size):
241+ n_cnt = int(npu_num.flatten()[b])
242+ g_cnt = int(golden_num.flatten()[b])
243+ 
244+ if n_cnt > 1:
245+ n_order = numpy.argsort(
246+ -npu_scores[b, :n_cnt].astype(numpy.float64), kind="stable"
247+ )
248+ npu_boxes[b, :n_cnt] = npu_boxes[b, n_order]
249+ npu_scores[b, :n_cnt] = npu_scores[b, n_order]
250+ npu_classes[b, :n_cnt] = npu_classes[b, n_order]
251+ 
252+ if g_cnt > 1:
253+ g_order = numpy.argsort(
254+ -golden_scores[b, :g_cnt].astype(numpy.float64), kind="stable"
255+ )
256+ golden_boxes[b, :g_cnt] = golden_boxes[b, g_order]
257+ golden_scores[b, :g_cnt] = golden_scores[b, g_order]
258+ golden_classes[b, :g_cnt] = golden_classes[b, g_order]
259+ 
260+ def compare(*outputs, **kwargs):
261+ """IoU-based matching comparison for NMS outputs.
262+ 
263+ Layout: func(*npu_outputs, *golden_outputs)
264+ Returns list[dict], one per output (boxes, scores, classes, num).
265+ """
266+ npu_boxes = _to_numpy(outputs[0]).astype(numpy.float64)
267+ npu_scores = _to_numpy(outputs[1]).astype(numpy.float64)
268+ npu_classes = _to_numpy(outputs[2]).astype(numpy.float64)
269+ npu_num = _to_numpy(outputs[3]).flatten()
270+ golden_boxes = _to_numpy(outputs[4]).astype(numpy.float64)
271+ golden_scores = _to_numpy(outputs[5]).astype(numpy.float64)
272+ golden_classes = _to_numpy(outputs[6]).astype(numpy.float64)
273+ golden_num = _to_numpy(outputs[7]).flatten()
274+ 
275+ batch_size = npu_boxes.shape[0]
276+ iou_match_threshold = 0.5
277+ score_rtol = 0.05
278+ 
279+ total_npu_det = 0
280+ total_golden_det = 0
281+ total_matched = 0
282+ total_score_close = 0
283+ total_class_match = 0
284+ total_box_iou_sum = 0.0
285+ num_diff_sum = 0
286+ 
287+ for b in range(batch_size):
288+ n_cnt = int(npu_num[b])
289+ g_cnt = int(golden_num[b])
290+ total_npu_det += n_cnt
291+ total_golden_det += g_cnt
292+ num_diff_sum += abs(n_cnt - g_cnt)
293+ 
294+ if n_cnt == 0 and g_cnt == 0:
295+ continue
296+ if n_cnt == 0 or g_cnt == 0:
297+ continue
298+ 
299+ n_boxes = npu_boxes[b, :n_cnt]
300+ g_boxes = golden_boxes[b, :g_cnt]
301+ n_scores_b = npu_scores[b, :n_cnt]
302+ g_scores_b = golden_scores[b, :g_cnt]
303+ n_classes_b = npu_classes[b, :n_cnt]
304+ g_classes_b = golden_classes[b, :g_cnt]
305+ 
306+ iou_mat = _numpy_iou_matrix(n_boxes, g_boxes)
307+ # A single geometric box may legitimately be retained by several
308+ # classes. Pure IoU ordering then pairs identical boxes
309+ # arbitrarily and can report an otherwise exact result as having
310+ # wrong scores/classes. Prefer semantic agreement before IoU.
311+ class_equal = n_classes_b[:, None] == g_classes_b[None, :]
312+ score_close = numpy.isclose(
313+ n_scores_b[:, None], g_scores_b[None, :], rtol=score_rtol, atol=1e-3
314+ )
315+ matched_n = set()
316+ matched_g = set()
317+ pairs = []
318+ flat_order = numpy.lexsort(
319+ (
320+ numpy.arange(iou_mat.size),
321+ -iou_mat.flatten(),
322+ -score_close.astype(numpy.int8).flatten(),
323+ -class_equal.astype(numpy.int8).flatten(),
324+ )
325+ )
326+ for idx in flat_order:
327+ ni = int(idx // g_cnt)
328+ gi = int(idx % g_cnt)
329+ if iou_mat[ni, gi] < iou_match_threshold:
330+ continue
331+ if ni in matched_n or gi in matched_g:
332+ continue
333+ matched_n.add(ni)
334+ matched_g.add(gi)
335+ pairs.append((ni, gi, iou_mat[ni, gi]))
336+ 
337+ total_matched += len(pairs)
338+ for ni, gi, iou_val in pairs:
339+ total_box_iou_sum += iou_val
340+ if numpy.isclose(
341+ n_scores_b[ni], g_scores_b[gi], rtol=score_rtol, atol=1e-3
342+ ):
343+ total_score_close += 1
344+ if n_classes_b[ni] == g_classes_b[gi]:
345+ total_class_match += 1
346+ 
347+ max_det = max(total_npu_det, total_golden_det, 1)
348+ if total_npu_det == 0 and total_golden_det == 0:
349+ # Empty detection sets are equivalent; treating the absence of
350+ # pairs as a 0% match creates a false negative for valid inputs.
351+ match_rate = 100.0
352+ avg_iou = 1.0
353+ score_rate = 100.0
354+ class_rate = 100.0
355+ else:
356+ match_rate = total_matched / max_det * 100.0
357+ avg_iou = total_box_iou_sum / max(total_matched, 1)
358+ score_rate = total_score_close / max(total_matched, 1) * 100.0
359+ class_rate = total_class_match / max(total_matched, 1) * 100.0
360+ num_accuracy = max(0.0, 100.0 - num_diff_sum / max(batch_size, 1) * 10)
361+ 
362+ boxes_result = {
363+ "pass": match_rate >= 70.0 and avg_iou >= 0.8,
364+ "precision": round(match_rate, 2),
365+ "error_info": f"match_rate={match_rate:.1f}% avg_iou={avg_iou:.3f} "
366+ f"npu_det={total_npu_det} golden_det={total_golden_det} matched={total_matched}",
367+ "metrics": {
368+ "match_rate": match_rate,
369+ "avg_iou": avg_iou,
370+ "total_npu_det": total_npu_det,
371+ "total_golden_det": total_golden_det,
372+ "total_matched": total_matched,
373+ },
374+ }
375+ scores_result = {
376+ "pass": score_rate >= 80.0,
377+ "precision": round(score_rate, 2),
378+ "error_info": f"score_close={total_score_close}/{total_matched}",
379+ "metrics": {"score_match_rate": score_rate},
380+ }
381+ classes_result = {
382+ "pass": class_rate >= 80.0,
383+ "precision": round(class_rate, 2),
384+ "error_info": f"class_match={total_class_match}/{total_matched}",
385+ "metrics": {"class_match_rate": class_rate},
386+ }
387+ num_result = {
388+ "pass": num_accuracy >= 50.0,
389+ "precision": round(num_accuracy, 2),
390+ "error_info": f"num_diff_sum={num_diff_sum} over {batch_size} batches",
391+ "metrics": {"num_diff_sum": int(num_diff_sum), "batch_size": batch_size},
392+ }
393+ return [boxes_result, scores_result, classes_result, num_result]
394+ 
395+ tolerance = {
396+ "float32": {"standard": "stat_rel_err"},
397+ "float16": {"standard": "stat_rel_err"},
398+ }
Aobjdetect/batch_multi_class_non_max_suppression/tests/assets/inputs.py+139-0
@@ -0,0 +1,139 @@
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+"""TTK input generator for BatchMultiClassNonMaxSuppression.
13+ 
14+Random coordinates do not necessarily form valid ``[y_min, x_min, y_max,
15+x_max]`` boxes. NMS would then filter most candidates before the suppression
16+path is exercised. This generator normalizes the generated data into valid
17+boxes and, when optional inputs are present, makes their values valid too.
18+"""
19+ 
20+import numpy as np
21+ 
22+ 
23+__input__ = {
24+ "kernel": {
25+ "batch_multi_class_non_max_suppression": (
26+ "batch_multi_class_non_max_suppression_inputs"
27+ ),
28+ },
29+}
30+ 
31+ 
32+def _to_numpy(value):
33+ if hasattr(value, "detach"):
34+ return value.detach().cpu().numpy()
35+ return np.asarray(value)
36+ 
37+ 
38+def _copy_to(destination, value):
39+ if hasattr(destination, "copy_"):
40+ destination.copy_(destination.new_tensor(value))
41+ else:
42+ destination[...] = value.astype(destination.dtype, copy=False)
43+ 
44+ 
45+def _transpose_box_from_args(extra_args, kwargs):
46+ del extra_args
47+ return bool(kwargs["transpose_box"]) if "transpose_box" in kwargs else None
48+ 
49+ 
50+def _valid_boxes(boxes, transpose_box):
51+ values = np.nan_to_num(_to_numpy(boxes).astype(np.float32, copy=True))
52+ if values.ndim != 4:
53+ raise ValueError("boxes must be rank 4")
54+ 
55+ if transpose_box is None:
56+ if values.shape[-1] == 4 and values.shape[2] != 4:
57+ transpose_box = False
58+ elif values.shape[2] == 4 and values.shape[-1] != 4:
59+ transpose_box = True
60+ elif values.shape[-1] == 4:
61+ transpose_box = False
62+ else:
63+ raise ValueError("boxes must contain a coordinate dimension of size 4")
64+ 
65+ canonical = values.transpose(0, 3, 1, 2) if transpose_box else values
66+ if canonical.shape[-1] != 4:
67+ raise ValueError("boxes must use [B, N, q, 4] or [B, q, 4, N] layout")
68+ 
69+ mins = np.minimum(canonical[..., :2], canonical[..., 2:])
70+ maxs = np.maximum(canonical[..., :2], canonical[..., 2:])
71+ # A sizable minimum extent survives float16 conversion and guarantees that
72+ # every generated candidate passes the geometric validity check.
73+ canonical[..., :2] = mins
74+ canonical[..., 2:] = np.maximum(maxs, mins + 0.125)
75+ normalized = canonical.transpose(0, 2, 3, 1) if transpose_box else canonical
76+ return normalized, canonical
77+ 
78+ 
79+def _valid_scores(scores):
80+ values = np.nan_to_num(_to_numpy(scores).astype(np.float32, copy=True))
81+ if values.ndim != 3:
82+ raise ValueError("scores must be rank 3")
83+ # CSV cases commonly request [0, 1]. For broader default ranges, sigmoid
84+ # gives score_threshold meaningful coverage without changing valid ranges.
85+ if np.any(values < 0.0) or np.any(values > 1.0):
86+ values = 1.0 / (1.0 + np.exp(-np.clip(values, -20.0, 20.0)))
87+ return values
88+ 
89+ 
90+def _valid_clip_window(clip_window, canonical_boxes):
91+ if clip_window is None:
92+ return
93+ values = _to_numpy(clip_window)
94+ if values.ndim != 2 or values.shape[1] != 4:
95+ raise ValueError("clip_window must have shape [B, 4]")
96+ if values.shape[0] != canonical_boxes.shape[0]:
97+ raise ValueError("clip_window batch size must match boxes")
98+ 
99+ y_min = canonical_boxes[..., 0].min(axis=(1, 2)) - 1.0
100+ x_min = canonical_boxes[..., 1].min(axis=(1, 2)) - 1.0
101+ y_max = canonical_boxes[..., 2].max(axis=(1, 2)) + 1.0
102+ x_max = canonical_boxes[..., 3].max(axis=(1, 2)) + 1.0
103+ _copy_to(clip_window, np.stack((y_min, x_min, y_max, x_max), axis=1))
104+ 
105+ 
106+def _valid_num_boxes(num_valid_boxes, batch_size, boxes_num):
107+ if num_valid_boxes is None:
108+ return
109+ values = _to_numpy(num_valid_boxes)
110+ if values.ndim != 1 or values.shape[0] != batch_size:
111+ raise ValueError("num_valid_boxes must have shape [B]")
112+ # Cover both the full-input and truncated-input paths when B > 1.
113+ counts = boxes_num - (np.arange(batch_size) % min(boxes_num, 2))
114+ _copy_to(num_valid_boxes, counts)
115+ 
116+ 
117+def batch_multi_class_non_max_suppression_inputs(
118+ boxes, scores, clip_window=None, num_valid_boxes=None, *extra_args, **kwargs
119+):
120+ """Make generated NMS inputs mutually valid in place for TTK."""
121+ transpose_box = _transpose_box_from_args(extra_args, kwargs)
122+ valid_boxes, canonical_boxes = _valid_boxes(boxes, transpose_box)
123+ score_values = _valid_scores(scores)
124+ 
125+ if canonical_boxes.shape[0] != score_values.shape[0]:
126+ raise ValueError("boxes and scores batch sizes must match")
127+ if canonical_boxes.shape[1] != score_values.shape[1]:
128+ raise ValueError("boxes and scores candidate counts must match")
129+ if canonical_boxes.shape[2] not in (1, score_values.shape[2]):
130+ raise ValueError("boxes class dimension must be 1 or match scores")
131+ 
132+ _copy_to(boxes, valid_boxes)
133+ _copy_to(scores, score_values)
134+ _valid_clip_window(clip_window, canonical_boxes)
135+ _valid_num_boxes(
136+ num_valid_boxes, canonical_boxes.shape[0], canonical_boxes.shape[1]
137+ )
138+ # Kernel direct-call mode consumes the normalized input arrays.
139+ return [boxes, scores, clip_window, num_valid_boxes]
Aobjdetect/batch_multi_class_non_max_suppression/tests/golden/cpu_golden.py+201-0
@@ -0,0 +1,201 @@
1+#!/usr/bin/env python3
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+"""Torch CPU golden plugin for kernel-direct TTK Promote validation."""
10+ 
11+from __future__ import annotations
12+ 
13+import numpy as np
14+import pickle
15+import subprocess
16+import sys
17+import tempfile
18+from pathlib import Path
19+ 
20+__spec__ = {
21+ "batch_multi_class_non_max_suppression": "BatchMultiClassNonMaxSuppressionTestSpec"
22+}
23+ 
24+ 
25+def _iou(boxes: torch.Tensor, box: torch.Tensor) -> torch.Tensor:
26+ mins = torch.maximum(boxes[:, :2], box[:2])
27+ maxes = torch.minimum(boxes[:, 2:], box[2:])
28+ inter = torch.clamp(maxes - mins, min=0).prod(dim=1)
29+ areas = (boxes[:, 2:] - boxes[:, :2]).prod(dim=1)
30+ box_area = (box[2:] - box[:2]).prod()
31+ return inter / torch.clamp(areas + box_area - inter, min=1e-12)
32+ 
33+ 
34+def _nms(boxes, scores, score_threshold, iou_threshold, max_size):
35+ valid = scores > score_threshold
36+ valid &= boxes[:, 2] > boxes[:, 0]
37+ valid &= boxes[:, 3] > boxes[:, 1]
38+ indices = torch.where(valid)[0]
39+ if indices.numel() == 0:
40+ return []
41+ boxes = boxes[indices]
42+ scores = scores[indices]
43+ order = torch.argsort(scores, descending=True, stable=True)
44+ boxes = boxes[order]
45+ scores = scores[order]
46+ suppressed = torch.zeros(scores.numel(), dtype=torch.bool)
47+ kept = []
48+ for index in range(scores.numel()):
49+ if bool(suppressed[index]):
50+ continue
51+ kept.append((scores[index], boxes[index].clone()))
52+ if len(kept) >= max_size:
53+ break
54+ remaining = torch.where(~suppressed)[0]
55+ remaining = remaining[remaining > index]
56+ if remaining.numel():
57+ suppressed[
58+ remaining[_iou(boxes[remaining], boxes[index]) > iou_threshold]
59+ ] = True
60+ return kept
61+ 
62+ 
63+def _torch_golden(
64+ boxes,
65+ scores,
66+ clip_window=None,
67+ num_valid_boxes=None,
68+ score_threshold=0.5,
69+ iou_threshold=0.5,
70+ max_size_per_class=100,
71+ max_total_size=100,
72+ change_coordinate_frame=False,
73+ transpose_box=False,
74+ **kwargs,
75+):
76+ boxes = torch.as_tensor(boxes, device="cpu")
77+ scores = torch.as_tensor(scores, device="cpu")
78+ clip_window = (
79+ None if clip_window is None else torch.as_tensor(clip_window, device="cpu")
80+ )
81+ num_valid_boxes = (
82+ None
83+ if num_valid_boxes is None
84+ else torch.as_tensor(num_valid_boxes, device="cpu")
85+ )
86+ dtype = boxes.dtype
87+ batch_size, num_classes = scores.shape[0], scores.shape[2]
88+ num_boxes = boxes.shape[3] if transpose_box else boxes.shape[1]
89+ box_classes = boxes.shape[1] if transpose_box else boxes.shape[2]
90+ out_boxes = torch.zeros((batch_size, max_total_size, 4), dtype=dtype)
91+ out_scores = torch.zeros((batch_size, max_total_size), dtype=dtype)
92+ out_classes = torch.zeros((batch_size, max_total_size), dtype=dtype)
93+ out_num = torch.zeros(batch_size, dtype=torch.int32)
94+ for batch in range(batch_size):
95+ valid_count = (
96+ num_boxes
97+ if num_valid_boxes is None
98+ else max(0, min(int(num_valid_boxes[batch]), num_boxes))
99+ )
100+ candidates = []
101+ for class_id in range(num_classes):
102+ class_scores = scores[batch, :, class_id].clone()
103+ class_scores[valid_count:] = -torch.inf
104+ box_class = 0 if box_classes == 1 else class_id
105+ class_boxes = (
106+ boxes[batch, box_class].T.contiguous().clone()
107+ if transpose_box
108+ else boxes[batch, :, box_class].clone()
109+ )
110+ if clip_window is not None:
111+ y_min, x_min, y_max, x_max = clip_window[batch]
112+ class_boxes[:, 0].clamp_(min=y_min)
113+ class_boxes[:, 1].clamp_(min=x_min)
114+ class_boxes[:, 2].clamp_(max=y_max)
115+ class_boxes[:, 3].clamp_(max=x_max)
116+ if change_coordinate_frame:
117+ if bool(y_max <= y_min) or bool(x_max <= x_min):
118+ continue
119+ class_boxes[:, (0, 2)] = (class_boxes[:, (0, 2)] - y_min) / (
120+ y_max - y_min
121+ )
122+ class_boxes[:, (1, 3)] = (class_boxes[:, (1, 3)] - x_min) / (
123+ x_max - x_min
124+ )
125+ for score, box in _nms(
126+ class_boxes,
127+ class_scores,
128+ float(score_threshold),
129+ float(iou_threshold),
130+ int(max_size_per_class),
131+ ):
132+ candidates.append((float(score), class_id, box))
133+ candidates.sort(key=lambda item: -item[0])
134+ count = min(len(candidates), int(max_total_size))
135+ out_num[batch] = count
136+ for index, (score, class_id, box) in enumerate(candidates[:count]):
137+ out_boxes[batch, index] = box
138+ out_scores[batch, index] = score
139+ out_classes[batch, index] = class_id
140+ return [value.numpy() for value in (out_boxes, out_scores, out_classes, out_num)]
141+ 
142+ 
143+def _golden(*inputs, **attrs):
144+ """Run Torch in a clean process, outside TTK's CANN-initialized worker."""
145+ with tempfile.TemporaryDirectory(prefix="batch_nms_cpu_golden_") as directory:
146+ request = Path(directory) / "request.pkl"
147+ response = Path(directory) / "response.pkl"
148+ request.write_bytes(
149+ pickle.dumps((inputs, attrs), protocol=pickle.HIGHEST_PROTOCOL)
150+ )
151+ subprocess.run(
152+ [
153+ sys.executable,
154+ str(Path(__file__).resolve()),
155+ "--torch-worker",
156+ str(request),
157+ str(response),
158+ ],
159+ check=True,
160+ )
161+ return pickle.loads(response.read_bytes())
162+ 
163+ 
164+class BatchMultiClassNonMaxSuppressionTestSpec:
165+ golden = staticmethod(_golden)
166+ tolerance = {
167+ "float16": {"standard": "cross_check", "level": "L1"},
168+ "float32": {"standard": "cross_check", "level": "L1"},
169+ }
170+ 
171+ @staticmethod
172+ def pre_compare(*outputs, **kwargs):
173+ half = len(outputs) // 2
174+ for offset in (0, half):
175+ boxes, scores, classes, count = outputs[offset : offset + 4]
176+ for batch in range(boxes.shape[0]):
177+ valid = int(np.asarray(count).reshape(-1)[batch])
178+ if valid <= 1:
179+ continue
180+ order = np.argsort(
181+ -np.asarray(scores)[batch, :valid].astype(np.float64), kind="stable"
182+ )
183+ boxes[batch, :valid] = boxes[batch, order]
184+ scores[batch, :valid] = scores[batch, order]
185+ classes[batch, :valid] = classes[batch, order]
186+ 
187+ 
188+def _worker(request: str, response: str) -> int:
189+ global torch
190+ import torch
191+ 
192+ inputs, attrs = pickle.loads(Path(request).read_bytes())
193+ outputs = _torch_golden(*inputs, **attrs)
194+ Path(response).write_bytes(pickle.dumps(outputs, protocol=pickle.HIGHEST_PROTOCOL))
195+ return 0
196+ 
197+ 
198+if __name__ == "__main__":
199+ if len(sys.argv) != 4 or sys.argv[1] != "--torch-worker":
200+ raise SystemExit("expected --torch-worker REQUEST RESPONSE")
201+ raise SystemExit(_worker(sys.argv[2], sys.argv[3]))
Aobjdetect/batch_multi_class_non_max_suppression/tests/ut/CMakeLists.txt+13-0
@@ -0,0 +1,13 @@
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+if(UT_TEST_ALL OR OP_HOST_UT)
12+ add_subdirectory(op_host)
13+endif()
Aobjdetect/batch_multi_class_non_max_suppression/tests/ut/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
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+add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
Aobjdetect/batch_multi_class_non_max_suppression/tests/ut/op_host/test_batch_multi_class_non_max_suppression_infershape.cpp+55-0
@@ -0,0 +1,55 @@
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 <gtest/gtest.h>
12+#include <vector>
13+ 
14+#include "infershape_case_executor.h"
15+#include "infershape_context_faker.h"
16+ 
17+namespace {
18+std::vector<gert::InfershapeContextPara::OpAttr> GetAttrs(int64_t maxTotalSize)
19+{
20+ return {
21+ gert::InfershapeContextPara::OpAttr("score_threshold", Ops::Cv::AnyValue::CreateFrom<float>(0.3F)),
22+ gert::InfershapeContextPara::OpAttr("iou_threshold", Ops::Cv::AnyValue::CreateFrom<float>(0.5F)),
23+ gert::InfershapeContextPara::OpAttr("max_size_per_class", Ops::Cv::AnyValue::CreateFrom<int64_t>(3)),
24+ gert::InfershapeContextPara::OpAttr("max_total_size", Ops::Cv::AnyValue::CreateFrom<int64_t>(maxTotalSize)),
25+ gert::InfershapeContextPara::OpAttr("change_coordinate_frame", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
26+ gert::InfershapeContextPara::OpAttr("transpose_box", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
27+ };
28+}
29+} // namespace
30+ 
31+TEST(BatchMultiClassNonMaxSuppressionInfershape, derives_public_output_layout)
32+{
33+ gert::InfershapeContextPara context("BatchMultiClassNonMaxSuppression",
34+ {{{{2, 8, 1, 4}, {2, 8, 1, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
35+ {{{2, 8, 3}, {2, 8, 3}}, ge::DT_FLOAT16, ge::FORMAT_ND}},
36+ {{{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
37+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
38+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
39+ {{{}, {}}, ge::DT_INT32, ge::FORMAT_ND}},
40+ GetAttrs(5));
41+ ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{2, 5, 4}, {2, 5}, {2, 5}, {2}});
42+}
43+ 
44+TEST(BatchMultiClassNonMaxSuppressionInfershape, rejects_non_rank_four_boxes)
45+{
46+ gert::InfershapeContextPara context(
47+ "BatchMultiClassNonMaxSuppression",
48+ {{{{2, 8, 4}, {2, 8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND}, {{{2, 8, 3}, {2, 8, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}},
49+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
50+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
51+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
52+ {{{}, {}}, ge::DT_INT32, ge::FORMAT_ND}},
53+ GetAttrs(5));
54+ ExecuteTestCase(context, ge::GRAPH_FAILED);
55+}
Mscripts/kernel/binary_config/ascendc_config.json+1-0
@@ -304,6 +304,7 @@
304 {"name":"AddRmsNormDynamicQuant", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : false},304 {"name":"AddRmsNormDynamicQuant", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : false},
305 {"name":"AddRmsNormDynamicQuantV2", "compute_units": ["ascend910b"], "auto_sync" : false},305 {"name":"AddRmsNormDynamicQuantV2", "compute_units": ["ascend910b"], "auto_sync" : false},
306 {"name":"BatchNormGradV3", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},306 {"name":"BatchNormGradV3", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
307+ {"name":"BatchMultiClassNonMaxSuppression", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
307 {"name":"GroupNormV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},308 {"name":"GroupNormV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
308 {"name":"SwiGlu", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync": true},309 {"name":"SwiGlu", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync": true},
309 {"name":"Axpy", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},310 {"name":"Axpy", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},