Star570
391
代码
代码
Issues20
Pull Requests45
Actions
项目讨论
Wiki
分析
Star570
391
已合并
新增 SortedNMS 算子 Ascend 950 实现 #1236
黄晓彬创建于 8月5日
新增 SortedNMS 算子 Ascend 950 实现 #1236
已合并
黄晓彬创建于 8月5日
从ConanHuang/ops-cvmaster合入到cann/ops-cvmaster
共 16 个文件变更+1900-39
Mcommon/inc/op_graph/op_cv_proto_extend.h+0-39
@@ -1155,45 +1155,6 @@ REG_OP(RpnProposals)
1155 .OUTPUT(sorted_box, TensorType({DT_FLOAT16}))1155 .OUTPUT(sorted_box, TensorType({DT_FLOAT16}))
1156 .OP_END_FACTORY_REG(RpnProposals)1156 .OP_END_FACTORY_REG(RpnProposals)
1157 1157 
1158-/**
1159-* @brief Greedily selects a subset of bounding boxes in descending order of
1160-* score . \n
1161- 
1162-* @par Inputs:
1163-* @li boxes: A 2-D float tensor of shape [num_boxes, 4]. They are expected to be in (x1, y1, x2, y2) format
1164-* with 0 <= x1 < x2 and 0 <= y1 < y2. Supported type: float16, float32. Supported format: ND.
1165-* @li sorted_scores: A 1-D float tensor of shape [num_boxes] representing boxes' scores, which is sorted
1166-* by descending order. Supported type: float16, float32. Supported format: ND.
1167-* @li input_indices: A 1-D integer tensor of shape [num_boxes] representing the indices for each row of
1168-* boxes that would sort row of boxes by scores in descending order. Supported type: int32. Supported format: ND.
1169-* @li max_output_size: A scalar integer tensor representing the maximum number
1170-* of boxes to be selected by non max suppression. Supported type: int32. Supported format: ND.
1171-* @li iou_threshold: A 0-D float tensor representing the threshold for deciding
1172-* whether boxes overlap too much with respect to IOU. Supported type: float16, float32. Supported format: ND.
1173-* @li score_threshold: A 0-D float tensor representing the threshold for
1174-* deciding when to remove boxes based on score. Supported type: float16, float32. Supported format: ND . \n
1175- 
1176-* @par Attributes:
1177-* offset: An optional int. Defaults to "0". \n
1178- 
1179-* @par Outputs:
1180-* @li selected_indices: A 1-D integer tensor of shape [M] representing the selected
1181-* indices from the boxes tensor, where M <= max_output_size. Supported type: int32. Supported format: ND . \n
1182- 
1183-* @attention Constraints:
1184-* Input boxes and scores must be float type . \n
1185-*/
1186-REG_OP(SortedNMS)
1187- .INPUT(boxes, TensorType({DT_FLOAT16, DT_FLOAT}))
1188- .INPUT(sorted_scores, TensorType({DT_FLOAT16, DT_FLOAT}))
1189- .INPUT(input_indices, TensorType({DT_INT32}))
1190- .INPUT(max_output_size, TensorType({DT_INT32}))
1191- .INPUT(iou_threshold, TensorType({DT_FLOAT16, DT_FLOAT}))
1192- .INPUT(score_threshold, TensorType({DT_FLOAT16, DT_FLOAT}))
1193- .OUTPUT(selected_indices, TensorType({DT_INT32}))
1194- .ATTR(offset, Int, 0)
1195- .OP_END_FACTORY_REG(SortedNMS)
1196- 
1197/**1158/**
1198*@brief HDRNet and ISP direct data conversion1159*@brief HDRNet and ISP direct data conversion
1199returned tensor's dimension will correspond to input dimension [0, 3, 4, 2, 1],1160returned tensor's dimension will correspond to input dimension [0, 3, 4, 2, 1],
Aobjdetect/sorted_nms/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
2+ 
3+if(NOT ENABLE_TEST AND NOT BENCHMARK)
4+ list(REMOVE_ITEM CURRENT_DIRS tests)
5+endif()
6+ 
7+foreach(SUB_DIR ${CURRENT_DIRS})
8+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
9+ add_subdirectory(${SUB_DIR})
10+ endif()
11+endforeach()
12+ 
13+set(SUPPORT_COMPUTE_UNIT "ascend950")
14+set(SUPPORT_TILING_DIR "arch35")
15+add_all_modules_sources(OPTYPE sorted_nms ACLNNTYPE aclnn_exclude
16+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
17+ TILING_DIR ${SUPPORT_TILING_DIR})
Aobjdetect/sorted_nms/README.md+123-0
@@ -0,0 +1,123 @@
1+# SortedNMS
吴
吴吴瑶29 天前

新增算子需补充op_list.md

likedislike
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| Ascend 950PR/Ascend 950DT | √ |
吴
吴吴瑶29 天前

产品名加标签

likedislike
8+| Atlas A3 训练系列产品/Atlas A3 推理系列产品 | √ |
9+| Atlas A2 训练系列产品/Atlas A2 推理系列产品 | √ |
10+| Atlas 200I/500 A2 推理产品 | × |
11+| Atlas 推理系列产品 | × |
12+| Atlas 训练系列产品 | × |
13+| Kirin X90 处理器系列产品 | × |
14+| Kirin 9030 处理器系列产品 | × |
15+ 
16+## 功能说明
17+ 
18+- 算子功能：在已按分数降序排列的候选框序列上，按照交并比（IoU）阈值贪心选择非抑制框，输出被选中框在原始boxes中的索引。
19+ 
20+- 计算公式：
21+ 
22+ $$
23+ IoU = \frac {Area_{inter}} {Area_{current} + Area_{next} - Area_{inter}}
24+ $$
25+ 
26+ 其中，Area_current为当前选中框的面积，Area_next为候选框的面积，Area_inter为两个框的重叠面积，offset为坐标计算偏移量。
27+ 
28+ $$
29+ Area_i = max(X_{2i} - X_{1i} + offset, 0) * max(Y_{2i} - Y_{1i} + offset, 0) \\
30+ Area_{inter} = max(min(X_{2c}, X_{2n}) - max(X_{1c}, X_{1n}) + offset, 0) * max(min(Y_{2c}, Y_{2n}) - max(Y_{1c}, Y_{1n}) + offset, 0)
31+ $$
32+ 
33+ 算子按照sorted_scores的非递增顺序遍历候选框。当候选框的分数大于score_threshold且未被抑制时，将input_indices中对应的索引加入输出；当候选框与当前选中框的IoU大于iou_threshold时，抑制该候选框。
34+ 
35+## 参数说明
36+ 
37+<table style="undefined;table-layout: fixed; width: 1005px"><colgroup>
38+ <col style="width: 170px">
39+ <col style="width: 170px">
40+ <col style="width: 352px">
41+ <col style="width: 213px">
42+ <col style="width: 100px">
43+ </colgroup>
44+ <thead>
45+ <tr>
46+ <th>参数名</th>
47+ <th>输入/输出/属性</th>
48+ <th>描述</th>
49+ <th>数据类型</th>
50+ <th>数据格式</th>
51+ </tr></thead>
52+ <tbody>
53+ <tr>
54+ <td>boxes</td>
55+ <td>输入</td>
56+ <td>候选矩形框，shape为(N, 4)，坐标格式为(X1, Y1, X2, Y2)。</td>
57+ <td>FLOAT32、FLOAT16</td>
58+ <td>ND</td>
59+ </tr>
60+ <tr>
61+ <td>sorted_scores</td>
62+ <td>输入</td>
63+ <td>候选矩形框的分数，shape为(N,)，需要按非递增顺序排列。</td>
64+ <td>FLOAT32、FLOAT16</td>
65+ <td>ND</td>
66+ </tr>
67+ <tr>
68+ <td>input_indices</td>
69+ <td>输入</td>
70+ <td>sorted_scores对应的候选框索引，shape为(N)，取值范围为[0, N)。</td>
71+ <td>INT32</td>
72+ <td>ND</td>
73+ </tr>
74+ <tr>
75+ <td>max_output_size</td>
76+ <td>输入</td>
77+ <td>最多输出的候选框数量，输入为标量或shape为(1,)的张量。</td>
78+ <td>INT32</td>
79+ <td>ND</td>
80+ </tr>
81+ <tr>
82+ <td>iou_threshold</td>
83+ <td>输入</td>
84+ <td>判断候选框是否需要抑制的IoU阈值，输入为标量或shape为(1,)的张量。</td>
85+ <td>FLOAT32、FLOAT16</td>
86+ <td>ND</td>
87+ </tr>
88+ <tr>
89+ <td>score_threshold</td>
90+ <td>输入</td>
91+ <td>过滤候选框的分数阈值，输入为标量或shape为(1,)的张量。</td>
92+ <td>FLOAT32、FLOAT16</td>
93+ <td>ND</td>
94+ </tr>
95+ <tr>
96+ <td>offset</td>
97+ <td>属性</td>
98+ <td>计算坐标差值时使用的偏移量，取值为0或1，默认值为0。</td>
99+ <td>INT</td>
100+ <td>-</td>
101+ </tr>
102+ <tr>
103+ <td>selected_indices</td>
104+ <td>输出</td>
105+ <td>被选中候选框在原始boxes中的索引，shape为(M)，M为运行时计算结果且M不大于min(max_output_size, N)。</td>
106+ <td>INT32</td>
107+ <td>ND</td>
108+ </tr>
109+ </tbody></table>
110+ 
111+- boxes与iou_threshold的数据类型需要保持一致；sorted_scores与score_threshold的数据类型需要保持一致，两组数据类型可以不同。
112+ 
113+## 约束说明
114+ 
115+- 输入shape限制：boxes为(N, 4)的二维张量，sorted_scores和input_indices为(N,)的一维张量，max_output_size、iou_threshold和score_threshold为标量或shape为(1,)的张量。
116+ 
117+- sorted_scores需要按照非递增顺序排列，input_indices需要为合法的候选框索引。
118+ 
119+## 调用说明
120+ 
121+| 调用方式 | 调用样例 | 说明 |
122+|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
123+| 图模式调用 | [test_geir_sorted_nms](./examples/test_geir_sorted_nms.cpp) | 通过[算子IR](./op_graph/sorted_nms_proto.h)构图方式调用SortedNMS算子。 |
Aobjdetect/sorted_nms/examples/test_geir_sorted_nms.cpp+195-0
@@ -0,0 +1,195 @@
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 <array>
12+#include <cstdint>
13+#include <cstring>
14+#include <iostream>
15+#include <map>
16+#include <memory>
17+#include <new>
18+#include <vector>
19+ 
20+#include "ge_api.h"
21+#include "ge_api_types.h"
22+#include "array_ops.h"
23+#include "graph.h"
24+#include "tensor.h"
25+#include "types.h"
26+ 
27+#include "../op_graph/sorted_nms_proto.h"
28+ 
29+namespace {
30+constexpr int32_t kSuccess = 0;
31+constexpr int32_t kFailed = -1;
32+constexpr int64_t kBoxCount = 5;
33+constexpr size_t kExpectedCount = 3;
34+ 
35+ge::TensorDesc MakeTensorDesc(const std::vector<int64_t>& shape, ge::DataType dtype)
36+{
37+ ge::TensorDesc desc(ge::Shape(shape), ge::FORMAT_ND, dtype);
38+ desc.SetPlacement(ge::kPlacementHost);
39+ return desc;
40+}
41+ 
42+template <typename T>
43+ge::Tensor MakeTensor(const ge::TensorDesc& desc, std::vector<T>& data)
44+{
45+ return ge::Tensor(desc, reinterpret_cast<uint8_t*>(data.data()), data.size() * sizeof(T));
46+}
47+ 
48+bool CheckOutput(const std::vector<ge::Tensor>& outputs)
49+{
50+ const std::array<int32_t, kExpectedCount> expected = {0, 2, 3};
51+ const std::vector<int64_t> expectedShape = {static_cast<int64_t>(kExpectedCount)};
52+ if (outputs.size() != 1 || outputs[0].GetTensorDesc().GetDataType() != ge::DT_INT32 ||
53+ outputs[0].GetTensorDesc().GetShape().GetDims() != expectedShape) {
54+ std::cerr << "Unexpected SortedNMS output descriptor: output_count=" << outputs.size();
55+ if (!outputs.empty()) {
56+ const auto& outputDesc = outputs[0].GetTensorDesc();
57+ std::cerr << ", dtype=" << static_cast<int32_t>(outputDesc.GetDataType()) << ", shape=[";
58+ const auto& dims = outputDesc.GetShape().GetDims();
59+ for (size_t i = 0; i < dims.size(); ++i) {
60+ std::cerr << (i == 0 ? "" : ", ") << dims[i];
61+ }
62+ std::cerr << "], bytes=" << outputs[0].GetSize();
63+ const uint8_t* outputData = outputs[0].GetData();
64+ const size_t valueCount = outputs[0].GetSize() / sizeof(int32_t);
65+ if (outputData != nullptr && valueCount > 0) {
66+ const size_t printCount = valueCount < kExpectedCount ? valueCount : kExpectedCount;
67+ std::cerr << ", raw_values=[";
68+ for (size_t i = 0; i < printCount; ++i) {
69+ int32_t value = 0;
70+ std::memcpy(&value, outputData + i * sizeof(value), sizeof(value));
71+ std::cerr << (i == 0 ? "" : ", ") << value;
72+ }
73+ std::cerr << "]";
74+ }
75+ }
76+ std::cerr << std::endl;
77+ return false;
78+ }
79+ 
80+ const uint8_t* outputData = outputs[0].GetData();
81+ if (outputData == nullptr) {
82+ std::cerr << "SortedNMS output data is null" << std::endl;
83+ return false;
84+ }
85+ 
86+ std::array<int32_t, kExpectedCount> actual{};
87+ std::memcpy(actual.data(), outputData, sizeof(actual));
88+ for (size_t i = 0; i < actual.size(); ++i) {
89+ std::cout << "selected_indices[" << i << "] = " << actual[i] << std::endl;
90+ if (actual[i] != expected[i]) {
91+ std::cerr << "Unexpected SortedNMS result at " << i << ": expected " << expected[i] << ", got " << actual[i]
92+ << std::endl;
93+ return false;
94+ }
95+ }
96+ return true;
97+}
98+} // namespace
99+ 
100+int main()
101+{
102+ const std::map<ge::AscendString, ge::AscendString> globalOptions = {{"ge.exec.deviceId", "0"},
103+ {"ge.graphRunMode", "1"}};
104+ ge::Status ret = ge::GEInitialize(globalOptions);
105+ if (ret != ge::GRAPH_SUCCESS) {
106+ std::cerr << "GEInitialize failed: " << ret << std::endl;
107+ return kFailed;
108+ }
109+ 
110+ int32_t result = kFailed;
111+ {
112+ ge::Graph graph("sorted_nms_geir_example");
113+ auto boxes = ge::op::Data("boxes").set_attr_index(0);
114+ auto sortedScores = ge::op::Data("sorted_scores").set_attr_index(1);
115+ auto inputIndices = ge::op::Data("input_indices").set_attr_index(2);
116+ auto maxOutputSize = ge::op::Data("max_output_size").set_attr_index(3);
117+ auto iouThreshold = ge::op::Data("iou_threshold").set_attr_index(4);
118+ auto scoreThreshold = ge::op::Data("score_threshold").set_attr_index(5);
119+ auto sortedNms = ge::op::SortedNMS("sorted_nms");
120+ 
121+ const ge::TensorDesc boxesDesc = MakeTensorDesc({kBoxCount, 4}, ge::DT_FLOAT);
122+ const ge::TensorDesc scoresDesc = MakeTensorDesc({kBoxCount}, ge::DT_FLOAT);
123+ const ge::TensorDesc indicesDesc = MakeTensorDesc({kBoxCount}, ge::DT_INT32);
124+ // GEIR graph inputs use the supported single-element scalar representation.
125+ const ge::TensorDesc scalarFloatDesc = MakeTensorDesc({1}, ge::DT_FLOAT);
126+ const ge::TensorDesc scalarIntDesc = MakeTensorDesc({1}, ge::DT_INT32);
127+ const ge::TensorDesc outputDesc = MakeTensorDesc({ge::UNKNOWN_DIM}, ge::DT_INT32);
128+ 
129+ boxes.update_input_desc_x(boxesDesc);
130+ sortedScores.update_input_desc_x(scoresDesc);
131+ inputIndices.update_input_desc_x(indicesDesc);
132+ maxOutputSize.update_input_desc_x(scalarIntDesc);
133+ iouThreshold.update_input_desc_x(scalarFloatDesc);
134+ scoreThreshold.update_input_desc_x(scalarFloatDesc);
135+ sortedNms.set_input_boxes(boxes)
136+ .set_input_sorted_scores(sortedScores)
137+ .set_input_input_indices(inputIndices)
138+ .set_input_max_output_size(maxOutputSize)
139+ .set_input_iou_threshold(iouThreshold)
140+ .set_input_score_threshold(scoreThreshold)
141+ .set_attr_offset(0);
142+ sortedNms.update_output_desc_selected_indices(outputDesc);
143+ 
144+ graph.AddOp(boxes);
145+ graph.AddOp(sortedScores);
146+ graph.AddOp(inputIndices);
147+ graph.AddOp(maxOutputSize);
148+ graph.AddOp(iouThreshold);
149+ graph.AddOp(scoreThreshold);
150+ const std::vector<ge::Operator> graphInputs = {boxes, sortedScores, inputIndices,
151+ maxOutputSize, iouThreshold, scoreThreshold};
152+ const std::vector<ge::Operator> graphOutputs = {sortedNms};
153+ graph.SetInputs(graphInputs).SetOutputs(graphOutputs);
154+ 
155+ // Scores are descending. Box 1 overlaps box 0 and must be suppressed.
156+ std::vector<float> boxesData = {0.0F, 0.0F, 10.0F, 10.0F, 1.0F, 1.0F, 9.0F, 9.0F, 20.0F, 20.0F,
157+ 30.0F, 30.0F, 40.0F, 40.0F, 50.0F, 50.0F, 60.0F, 60.0F, 70.0F, 70.0F};
158+ std::vector<float> scoresData = {0.95F, 0.90F, 0.75F, 0.60F, 0.40F};
159+ std::vector<int32_t> indicesData = {0, 1, 2, 3, 4};
160+ std::vector<int32_t> maxOutputSizeData = {3};
161+ std::vector<float> iouThresholdData = {0.5F};
162+ std::vector<float> scoreThresholdData = {0.5F};
163+ std::vector<ge::Tensor> inputs = {MakeTensor(boxesDesc, boxesData),
164+ MakeTensor(scoresDesc, scoresData),
165+ MakeTensor(indicesDesc, indicesData),
166+ MakeTensor(scalarIntDesc, maxOutputSizeData),
167+ MakeTensor(scalarFloatDesc, iouThresholdData),
168+ MakeTensor(scalarFloatDesc, scoreThresholdData)};
169+ 
170+ const std::map<ge::AscendString, ge::AscendString> sessionOptions;
171+ const std::map<ge::AscendString, ge::AscendString> graphOptions;
172+ std::unique_ptr<ge::Session> session(new (std::nothrow) ge::Session(sessionOptions));
173+ if (session == nullptr) {
174+ std::cerr << "Failed to create GE session" << std::endl;
175+ } else if ((ret = session->AddGraph(0, graph, graphOptions)) != ge::GRAPH_SUCCESS) {
176+ std::cerr << "AddGraph failed: " << ret << std::endl;
177+ } else {
178+ std::vector<ge::Tensor> outputs;
179+ ret = session->RunGraph(0, inputs, outputs);
180+ if (ret != ge::GRAPH_SUCCESS) {
181+ std::cerr << "RunGraph failed: " << ret << std::endl;
182+ } else if (CheckOutput(outputs)) {
183+ std::cout << "SortedNMS GEIR example passed" << std::endl;
184+ result = kSuccess;
185+ }
186+ }
187+ }
188+ 
189+ ret = ge::GEFinalize();
190+ if (ret != ge::GRAPH_SUCCESS) {
191+ std::cerr << "GEFinalize failed: " << ret << std::endl;
192+ return kFailed;
193+ }
194+ return result;
195+}
Aobjdetect/sorted_nms/op_graph/sorted_nms_proto.h+47-0
@@ -0,0 +1,47 @@
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 OPS_CV_OBJDETECT_SORTED_NMS_PROTO_H_
12+#define OPS_CV_OBJDETECT_SORTED_NMS_PROTO_H_
13+ 
14+#include "graph/operator_reg.h"
15+ 
16+namespace ge {
17+/**
18+ * @brief Greedily selects non-suppressed boxes in the supplied descending score order.
19+ *
20+ * @par Inputs:
21+ * @li boxes: A 2-D tensor of shape [num_boxes, 4].
22+ * @li sorted_scores: A 1-D tensor of shape [num_boxes], sorted in descending order.
23+ * @li input_indices: A 1-D int32 tensor of shape [num_boxes], with values in [0, num_boxes).
24+ * @li max_output_size: A scalar int32 tensor giving the maximum number of selected boxes.
25+ * @li iou_threshold: A scalar float tensor in [0, 1].
26+ * @li score_threshold: A scalar float tensor.
27+ *
28+ * boxes and iou_threshold must have the same type (float16 or float32). sorted_scores and
29+ * score_threshold must have the same type (float16 or float32), independently of boxes.
30+ *
31+ * @par Outputs:
32+ * selected_indices: A 1-D int32 tensor of shape [M], where M is computed at runtime and
33+ * M <= min(max_output_size, num_boxes).
34+ */
35+REG_OP(SortedNMS)
36+ .INPUT(boxes, TensorType({DT_FLOAT16, DT_FLOAT}))
37+ .INPUT(sorted_scores, TensorType({DT_FLOAT16, DT_FLOAT}))
38+ .INPUT(input_indices, TensorType({DT_INT32}))
39+ .INPUT(max_output_size, TensorType({DT_INT32}))
40+ .INPUT(iou_threshold, TensorType({DT_FLOAT16, DT_FLOAT}))
41+ .INPUT(score_threshold, TensorType({DT_FLOAT16, DT_FLOAT}))
42+ .OUTPUT(selected_indices, TensorType({DT_INT32}))
43+ .ATTR(offset, Int, 0)
44+ .OP_END_FACTORY_REG(SortedNMS)
45+} // namespace ge
46+ 
47+#endif // OPS_CV_OBJDETECT_SORTED_NMS_PROTO_H_
Aobjdetect/sorted_nms/op_host/arch35/sorted_nms_tiling_arch35.cpp+232-0
@@ -0,0 +1,232 @@
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 <algorithm>
12+#include <limits>
13+#include <set>
14+#include "register/op_def_registry.h"
15+#include "op_common/log/log.h"
16+#include "tiling/platform/platform_ascendc.h"
17+#include "../../op_kernel/arch35/sorted_nms_tiling_data.h"
18+ 
19+namespace optiling {
20+namespace {
21+constexpr size_t WORKSPACE_NUM = 1;
22+constexpr int64_t INPUT_BOXES = 0;
23+constexpr int64_t INPUT_SORTED_SCORES = 1;
24+constexpr int64_t INPUT_INDICES = 2;
25+constexpr int64_t INPUT_MAX_OUTPUT_SIZE = 3;
26+constexpr int64_t INPUT_IOU_THRESHOLD = 4;
27+constexpr int64_t INPUT_SCORE_THRESHOLD = 5;
28+constexpr int64_t ATTR_OFFSET = 0;
29+constexpr int64_t BOX_RANK = 2;
30+constexpr int64_t BOX_COORDS = 4;
31+constexpr int64_t MASK_BITS = 32;
32+constexpr int64_t PAIR_MASK_WORDS_PER_CORE = 1024;
33+constexpr int64_t MULTI_CORE_MIN_BOXES = 1025;
34+constexpr int64_t MULTI_CORE_MAX_BOXES = 8192;
35+constexpr int32_t STRATEGY_SINGLE_CORE = 0;
36+constexpr int32_t STRATEGY_PAIRWISE_MASK = 1;
37+constexpr size_t WORK_CONTROL_NUM = 2U;
38+constexpr uint64_t UB_BLOCK_SIZE = 32U;
39+constexpr uint64_t SIMT_DATA_CACHE_RESERVE = 128U * 1024U;
40+constexpr int64_t MAX_BOXES_NUM = static_cast<int64_t>(std::numeric_limits<int32_t>::max());
41+} // namespace
42+ 
43+static uint64_t AlignUbBytes(uint64_t bytes) { return (bytes + UB_BLOCK_SIZE - 1U) / UB_BLOCK_SIZE * UB_BLOCK_SIZE; }
44+ 
45+static bool IsScalarOrSingleElement(const gert::Shape& shape)
46+{
47+ return shape.IsScalar() || (shape.GetDimNum() == 1U && shape.GetDim(0) == 1);
48+}
49+ 
50+static ge::graphStatus SetWorkspace(gert::TilingContext* context, int64_t boxesNum, int32_t strategy,
51+ int64_t maskWordNum)
52+{
53+ size_t* workspace = context->GetWorkspaceSizes(WORKSPACE_NUM);
54+ OP_CHECK_NULL_WITH_CONTEXT(context, workspace);
55+ OP_CHECK_IF(boxesNum < 0 || boxesNum > MAX_BOXES_NUM,
56+ OP_LOGE(context, "boxes_num must be in [0, %ld], got %ld", MAX_BOXES_NUM, boxesNum),
57+ return ge::GRAPH_FAILED);
58+ constexpr size_t MAX_WORKSPACE_INT32S = std::numeric_limits<size_t>::max() / sizeof(int32_t);
59+ const size_t boxesNumSize = static_cast<size_t>(boxesNum);
60+ size_t userWorkspaceInt32s = 0;
61+ if (strategy == STRATEGY_PAIRWISE_MASK) {
62+ const size_t maskWordNumSize = static_cast<size_t>(maskWordNum);
63+ OP_CHECK_IF(maskWordNumSize != 0 && boxesNumSize > MAX_WORKSPACE_INT32S / maskWordNumSize,
64+ OP_LOGE(context, "pairwise mask workspace overflows for boxes_num %ld", boxesNum),
65+ return ge::GRAPH_FAILED);
66+ const size_t pairMaskWords = boxesNumSize * maskWordNumSize;
67+ OP_CHECK_IF(maskWordNumSize > MAX_WORKSPACE_INT32S - WORK_CONTROL_NUM ||
68+ pairMaskWords > MAX_WORKSPACE_INT32S - WORK_CONTROL_NUM - maskWordNumSize,
69+ OP_LOGE(context, "pairwise mask workspace overflows for boxes_num %ld", boxesNum),
70+ return ge::GRAPH_FAILED);
71+ userWorkspaceInt32s = WORK_CONTROL_NUM + maskWordNumSize + pairMaskWords;
72+ } else {
73+ OP_CHECK_IF(boxesNumSize > MAX_WORKSPACE_INT32S - WORK_CONTROL_NUM,
74+ OP_LOGE(context, "workspace size overflows for boxes_num %ld", boxesNum), return ge::GRAPH_FAILED);
75+ userWorkspaceInt32s = WORK_CONTROL_NUM + boxesNumSize;
76+ }
77+ const size_t userWorkspaceSize = userWorkspaceInt32s * sizeof(int32_t);
78+ auto* platformInfo = context->GetPlatformInfo();
79+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
80+ const auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
81+ const size_t systemWorkspaceSize = static_cast<size_t>(ascendcPlatform.GetLibApiWorkSpaceSize());
82+ OP_CHECK_IF(userWorkspaceSize > std::numeric_limits<size_t>::max() - systemWorkspaceSize,
83+ OP_LOGE(context, "workspace size overflows for boxes_num %ld", boxesNum), return ge::GRAPH_FAILED);
84+ workspace[0] = systemWorkspaceSize + userWorkspaceSize;
85+ return ge::GRAPH_SUCCESS;
86+}
87+ 
88+static ge::graphStatus CheckInputDesc(gert::TilingContext* context)
89+{
90+ const std::set<ge::DataType> dataDtypes = {ge::DT_FLOAT16, ge::DT_FLOAT};
91+ const auto* boxesDesc = context->GetInputDesc(INPUT_BOXES);
92+ const auto* scoresDesc = context->GetInputDesc(INPUT_SORTED_SCORES);
93+ const auto* indicesDesc = context->GetInputDesc(INPUT_INDICES);
94+ const auto* maxOutputDesc = context->GetInputDesc(INPUT_MAX_OUTPUT_SIZE);
95+ const auto* iouDesc = context->GetInputDesc(INPUT_IOU_THRESHOLD);
96+ const auto* scoreThrDesc = context->GetInputDesc(INPUT_SCORE_THRESHOLD);
97+ OP_CHECK_IF(boxesDesc == nullptr || scoresDesc == nullptr || indicesDesc == nullptr || maxOutputDesc == nullptr ||
98+ iouDesc == nullptr || scoreThrDesc == nullptr,
99+ OP_LOGE(context, "input desc is nullptr"), return ge::GRAPH_FAILED);
100+ 
101+ OP_CHECK_IF(dataDtypes.count(boxesDesc->GetDataType()) == 0, OP_LOGE(context, "unsupported boxes dtype"),
102+ return ge::GRAPH_FAILED);
103+ OP_CHECK_IF(dataDtypes.count(scoresDesc->GetDataType()) == 0, OP_LOGE(context, "unsupported sorted_scores dtype"),
104+ return ge::GRAPH_FAILED);
105+ OP_CHECK_IF(dataDtypes.count(iouDesc->GetDataType()) == 0, OP_LOGE(context, "unsupported iou_threshold dtype"),
106+ return ge::GRAPH_FAILED);
107+ OP_CHECK_IF(dataDtypes.count(scoreThrDesc->GetDataType()) == 0,
108+ OP_LOGE(context, "unsupported score_threshold dtype"), return ge::GRAPH_FAILED);
109+ OP_CHECK_IF(indicesDesc->GetDataType() != ge::DT_INT32, OP_LOGE(context, "input_indices dtype must be int32"),
110+ return ge::GRAPH_FAILED);
111+ OP_CHECK_IF(maxOutputDesc->GetDataType() != ge::DT_INT32, OP_LOGE(context, "max_output_size dtype must be int32"),
112+ return ge::GRAPH_FAILED);
113+ OP_CHECK_IF(boxesDesc->GetDataType() != iouDesc->GetDataType(),
114+ OP_LOGE(context, "boxes and iou_threshold must share one dtype"), return ge::GRAPH_FAILED);
115+ OP_CHECK_IF(scoresDesc->GetDataType() != scoreThrDesc->GetDataType(),
116+ OP_LOGE(context, "sorted_scores and score_threshold must share one dtype"), return ge::GRAPH_FAILED);
117+ return ge::GRAPH_SUCCESS;
118+}
119+ 
120+static ge::graphStatus GetOffsetAttr(gert::TilingContext* context, int32_t* offset)
121+{
122+ auto attrs = context->GetAttrs();
123+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
124+ const auto* offsetPtr = attrs->GetAttrPointer<int64_t>(ATTR_OFFSET);
125+ OP_CHECK_NULL_WITH_CONTEXT(context, offsetPtr);
126+ OP_CHECK_IF(*offsetPtr != 0 && *offsetPtr != 1, OP_LOGE(context, "offset must be 0 or 1"), return ge::GRAPH_FAILED);
127+ *offset = static_cast<int32_t>(*offsetPtr);
128+ return ge::GRAPH_SUCCESS;
129+}
130+ 
131+static ge::graphStatus SortedNMSTilingFunc(gert::TilingContext* context)
132+{
133+ OP_LOGI(context->GetNodeName(), "Enter SortedNMSTilingFunc");
134+ OP_CHECK_IF(CheckInputDesc(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "check input desc failed"),
135+ return ge::GRAPH_FAILED);
136+ 
137+ auto boxesShapePtr = context->GetInputShape(INPUT_BOXES);
138+ auto scoresShapePtr = context->GetInputShape(INPUT_SORTED_SCORES);
139+ auto indicesShapePtr = context->GetInputShape(INPUT_INDICES);
140+ auto maxOutputShapePtr = context->GetInputShape(INPUT_MAX_OUTPUT_SIZE);
141+ auto iouShapePtr = context->GetInputShape(INPUT_IOU_THRESHOLD);
142+ auto scoreThrShapePtr = context->GetInputShape(INPUT_SCORE_THRESHOLD);
143+ OP_CHECK_IF(boxesShapePtr == nullptr || scoresShapePtr == nullptr || indicesShapePtr == nullptr ||
144+ maxOutputShapePtr == nullptr || iouShapePtr == nullptr || scoreThrShapePtr == nullptr,
145+ OP_LOGE(context, "input shape is nullptr"), return ge::GRAPH_FAILED);
146+ 
147+ auto boxesShape = boxesShapePtr->GetStorageShape();
148+ auto scoresShape = scoresShapePtr->GetStorageShape();
149+ auto indicesShape = indicesShapePtr->GetStorageShape();
150+ auto maxOutputShape = maxOutputShapePtr->GetStorageShape();
151+ auto iouShape = iouShapePtr->GetStorageShape();
152+ auto scoreThrShape = scoreThrShapePtr->GetStorageShape();
153+ OP_CHECK_IF(boxesShape.GetDimNum() != BOX_RANK, OP_LOGE(context, "boxes rank must be 2"), return ge::GRAPH_FAILED);
154+ OP_CHECK_IF(boxesShape.GetDim(1) != ge::UNKNOWN_DIM && boxesShape.GetDim(1) != BOX_COORDS,
155+ OP_LOGE(context, "boxes second dim must be 4"), return ge::GRAPH_FAILED);
156+ int64_t boxesNum = boxesShape.GetDim(0);
157+ OP_CHECK_IF(boxesNum < 0 || boxesNum > MAX_BOXES_NUM,
158+ OP_LOGE(context, "boxes first dim must be in [0, %ld] for tiling, got %ld", MAX_BOXES_NUM, boxesNum),
159+ return ge::GRAPH_FAILED);
160+ OP_CHECK_IF(scoresShape.GetDimNum() != 1 || scoresShape.GetDim(0) != boxesNum,
161+ OP_LOGE(context, "sorted_scores shape must be [boxes_num]"), return ge::GRAPH_FAILED);
162+ OP_CHECK_IF(indicesShape.GetDimNum() != 1 || indicesShape.GetDim(0) != boxesNum,
163+ OP_LOGE(context, "input_indices shape must be [boxes_num]"), return ge::GRAPH_FAILED);
164+ OP_CHECK_IF(!IsScalarOrSingleElement(maxOutputShape),
165+ OP_LOGE(context, "max_output_size shape must be scalar or [1]"), return ge::GRAPH_FAILED);
166+ OP_CHECK_IF(!IsScalarOrSingleElement(iouShape), OP_LOGE(context, "iou_threshold shape must be scalar or [1]"),
167+ return ge::GRAPH_FAILED);
168+ OP_CHECK_IF(!IsScalarOrSingleElement(scoreThrShape),
169+ OP_LOGE(context, "score_threshold shape must be scalar or [1]"), return ge::GRAPH_FAILED);
170+ const int32_t strategy = boxesNum >= MULTI_CORE_MIN_BOXES && boxesNum <= MULTI_CORE_MAX_BOXES ?
171+ STRATEGY_PAIRWISE_MASK :
172+ STRATEGY_SINGLE_CORE;
173+ const int64_t maskWordNum = strategy == STRATEGY_PAIRWISE_MASK ? (boxesNum + MASK_BITS - 1) / MASK_BITS : 0;
174+ OP_CHECK_IF(SetWorkspace(context, boxesNum, strategy, maskWordNum) != ge::GRAPH_SUCCESS,
175+ OP_LOGE(context, "set workspace failed"), return ge::GRAPH_FAILED);
176+ 
177+ SortedNMSTilingData* tiling = context->GetTilingData<SortedNMSTilingData>();
178+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
179+ OP_CHECK_IF(memset_s(tiling, sizeof(SortedNMSTilingData), 0, sizeof(SortedNMSTilingData)) != EOK,
180+ OP_LOGE(context, "set tiling data failed"), return ge::GRAPH_FAILED);
181+ 
182+ tiling->boxesNum = boxesNum;
183+ OP_CHECK_IF(GetOffsetAttr(context, &tiling->offset) != ge::GRAPH_SUCCESS, OP_LOGE(context, "get offset failed"),
184+ return ge::GRAPH_FAILED);
185+ auto* platformInfo = context->GetPlatformInfo();
186+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
187+ const auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
188+ const int64_t physicalCoreNum = static_cast<int64_t>(ascendcPlatform.GetCoreNumAiv());
189+ OP_CHECK_IF(physicalCoreNum <= 0, OP_LOGE(context, "AIV core num must be positive"), return ge::GRAPH_FAILED);
190+ if (strategy == STRATEGY_PAIRWISE_MASK) {
191+ const auto* boxesDesc = context->GetInputDesc(INPUT_BOXES);
192+ OP_CHECK_NULL_WITH_CONTEXT(context, boxesDesc);
193+ if (boxesDesc->GetDataType() == ge::DT_FLOAT16) {
194+ const uint64_t boxesBytes = AlignUbBytes(static_cast<uint64_t>(boxesNum) * BOX_COORDS * sizeof(uint16_t));
195+ const uint64_t areasBytes = AlignUbBytes(static_cast<uint64_t>(boxesNum) * sizeof(float));
196+ uint64_t ubSize = 0;
197+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
198+ OP_CHECK_IF(ubSize <= SIMT_DATA_CACHE_RESERVE,
199+ OP_LOGE(context, "%lu-byte UB cannot reserve 128KB SIMT Data Cache", ubSize),
200+ return ge::GRAPH_FAILED);
201+ const uint64_t localMemorySize = ubSize - SIMT_DATA_CACHE_RESERVE;
202+ OP_CHECK_IF(localMemorySize > std::numeric_limits<uint32_t>::max(),
203+ OP_LOGE(context, "local memory size %lu exceeds uint32 range", localMemorySize),
204+ return ge::GRAPH_FAILED);
205+ tiling->useLocalBoxes = boxesBytes <= localMemorySize && areasBytes <= localMemorySize - boxesBytes ? 1 : 0;
206+ OP_CHECK_IF(context->SetLocalMemorySize(static_cast<uint32_t>(localMemorySize)) != ge::GRAPH_SUCCESS,
207+ OP_LOGE(context, "SetLocalMemorySize failed for %lu bytes", localMemorySize),
208+ return ge::GRAPH_FAILED);
209+ }
210+ }
211+ int64_t requiredCoreNum = 1;
212+ if (strategy == STRATEGY_PAIRWISE_MASK) {
213+ const int64_t pairMaskWords = boxesNum * maskWordNum;
214+ requiredCoreNum = (pairMaskWords + PAIR_MASK_WORDS_PER_CORE - 1) / PAIR_MASK_WORDS_PER_CORE;
215+ }
216+ tiling->coreNum = static_cast<int32_t>(std::min(requiredCoreNum, physicalCoreNum));
217+ context->SetBlockDim(static_cast<uint32_t>(tiling->coreNum));
218+ if (strategy == STRATEGY_PAIRWISE_MASK) {
219+ context->SetScheduleMode(1);
220+ }
221+ return ge::GRAPH_SUCCESS;
222+}
223+ 
224+static ge::graphStatus TilingParseForSortedNMS([[maybe_unused]] gert::TilingParseContext* context)
225+{
226+ return ge::GRAPH_SUCCESS;
227+}
228+ 
229+struct SortedNMSCompileInfo {};
230+ 
231+IMPL_OP_OPTILING(SortedNMS).Tiling(SortedNMSTilingFunc).TilingParse<SortedNMSCompileInfo>(TilingParseForSortedNMS);
232+} // namespace optiling
Aobjdetect/sorted_nms/op_host/config/ascend950/sorted_nms_binary.json+313-0
@@ -0,0 +1,313 @@
1+{
2+ "op_type": "SortedNMS",
3+ "op_list": [
4+ {
5+ "bin_filename": "SortedNMS_fbc40177cc5f4223a624050998ede2f1",
6+ "inputs": [
7+ {
8+ "name": "boxes",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [-2],
14+ "format_match_mode": "FormatAgnostic"
15+ },
16+ {
17+ "name": "sorted_scores",
18+ "index": 1,
19+ "dtype": "float16",
20+ "format": "ND",
21+ "paramType": "required",
22+ "shape": [-2],
23+ "format_match_mode": "FormatAgnostic"
24+ },
25+ {
26+ "name": "input_indices",
27+ "index": 2,
28+ "dtype": "int32",
29+ "format": "ND",
30+ "paramType": "required",
31+ "shape": [-2],
32+ "format_match_mode": "FormatAgnostic"
33+ },
34+ {
35+ "name": "max_output_size",
36+ "index": 3,
37+ "dtype": "int32",
38+ "format": "ND",
39+ "paramType": "required",
40+ "shape": [-2],
41+ "format_match_mode": "FormatAgnostic"
42+ },
43+ {
44+ "name": "iou_threshold",
45+ "index": 4,
46+ "dtype": "float16",
47+ "format": "ND",
48+ "paramType": "required",
49+ "shape": [-2],
50+ "format_match_mode": "FormatAgnostic"
51+ },
52+ {
53+ "name": "score_threshold",
54+ "index": 5,
55+ "dtype": "float16",
56+ "format": "ND",
57+ "paramType": "required",
58+ "shape": [-2],
59+ "format_match_mode": "FormatAgnostic"
60+ }
61+ ],
62+ "outputs": [
63+ {
64+ "name": "selected_indices",
65+ "index": 0,
66+ "dtype": "int32",
67+ "format": "ND",
68+ "paramType": "required",
69+ "shape": [-2],
70+ "format_match_mode": "FormatAgnostic"
71+ }
72+ ],
73+ "attrs": [
74+ {
75+ "name": "offset",
76+ "dtype": "int",
77+ "value": 0
78+ }
79+ ]
80+ },
81+ {
82+ "bin_filename": "SortedNMS_d1e8cd8e75b0c427e4f4efcf546ca39b",
83+ "inputs": [
84+ {
85+ "name": "boxes",
86+ "index": 0,
87+ "dtype": "float16",
88+ "format": "ND",
89+ "paramType": "required",
90+ "shape": [-2],
91+ "format_match_mode": "FormatAgnostic"
92+ },
93+ {
94+ "name": "sorted_scores",
95+ "index": 1,
96+ "dtype": "float32",
97+ "format": "ND",
98+ "paramType": "required",
99+ "shape": [-2],
100+ "format_match_mode": "FormatAgnostic"
101+ },
102+ {
103+ "name": "input_indices",
104+ "index": 2,
105+ "dtype": "int32",
106+ "format": "ND",
107+ "paramType": "required",
108+ "shape": [-2],
109+ "format_match_mode": "FormatAgnostic"
110+ },
111+ {
112+ "name": "max_output_size",
113+ "index": 3,
114+ "dtype": "int32",
115+ "format": "ND",
116+ "paramType": "required",
117+ "shape": [-2],
118+ "format_match_mode": "FormatAgnostic"
119+ },
120+ {
121+ "name": "iou_threshold",
122+ "index": 4,
123+ "dtype": "float16",
124+ "format": "ND",
125+ "paramType": "required",
126+ "shape": [-2],
127+ "format_match_mode": "FormatAgnostic"
128+ },
129+ {
130+ "name": "score_threshold",
131+ "index": 5,
132+ "dtype": "float32",
133+ "format": "ND",
134+ "paramType": "required",
135+ "shape": [-2],
136+ "format_match_mode": "FormatAgnostic"
137+ }
138+ ],
139+ "outputs": [
140+ {
141+ "name": "selected_indices",
142+ "index": 0,
143+ "dtype": "int32",
144+ "format": "ND",
145+ "paramType": "required",
146+ "shape": [-2],
147+ "format_match_mode": "FormatAgnostic"
148+ }
149+ ],
150+ "attrs": [
151+ {
152+ "name": "offset",
153+ "dtype": "int",
154+ "value": 0
155+ }
156+ ]
157+ },
158+ {
159+ "bin_filename": "SortedNMS_b2480a92abd1aca211709bf6275fa350",
160+ "inputs": [
161+ {
162+ "name": "boxes",
163+ "index": 0,
164+ "dtype": "float32",
165+ "format": "ND",
166+ "paramType": "required",
167+ "shape": [-2],
168+ "format_match_mode": "FormatAgnostic"
169+ },
170+ {
171+ "name": "sorted_scores",
172+ "index": 1,
173+ "dtype": "float16",
174+ "format": "ND",
175+ "paramType": "required",
176+ "shape": [-2],
177+ "format_match_mode": "FormatAgnostic"
178+ },
179+ {
180+ "name": "input_indices",
181+ "index": 2,
182+ "dtype": "int32",
183+ "format": "ND",
184+ "paramType": "required",
185+ "shape": [-2],
186+ "format_match_mode": "FormatAgnostic"
187+ },
188+ {
189+ "name": "max_output_size",
190+ "index": 3,
191+ "dtype": "int32",
192+ "format": "ND",
193+ "paramType": "required",
194+ "shape": [-2],
195+ "format_match_mode": "FormatAgnostic"
196+ },
197+ {
198+ "name": "iou_threshold",
199+ "index": 4,
200+ "dtype": "float32",
201+ "format": "ND",
202+ "paramType": "required",
203+ "shape": [-2],
204+ "format_match_mode": "FormatAgnostic"
205+ },
206+ {
207+ "name": "score_threshold",
208+ "index": 5,
209+ "dtype": "float16",
210+ "format": "ND",
211+ "paramType": "required",
212+ "shape": [-2],
213+ "format_match_mode": "FormatAgnostic"
214+ }
215+ ],
216+ "outputs": [
217+ {
218+ "name": "selected_indices",
219+ "index": 0,
220+ "dtype": "int32",
221+ "format": "ND",
222+ "paramType": "required",
223+ "shape": [-2],
224+ "format_match_mode": "FormatAgnostic"
225+ }
226+ ],
227+ "attrs": [
228+ {
229+ "name": "offset",
230+ "dtype": "int",
231+ "value": 0
232+ }
233+ ]
234+ },
235+ {
236+ "bin_filename": "SortedNMS_55c1d7cf8ff616b8d9d92d67eed23d89",
237+ "inputs": [
238+ {
239+ "name": "boxes",
240+ "index": 0,
241+ "dtype": "float32",
242+ "format": "ND",
243+ "paramType": "required",
244+ "shape": [-2],
245+ "format_match_mode": "FormatAgnostic"
246+ },
247+ {
248+ "name": "sorted_scores",
249+ "index": 1,
250+ "dtype": "float32",
251+ "format": "ND",
252+ "paramType": "required",
253+ "shape": [-2],
254+ "format_match_mode": "FormatAgnostic"
255+ },
256+ {
257+ "name": "input_indices",
258+ "index": 2,
259+ "dtype": "int32",
260+ "format": "ND",
261+ "paramType": "required",
262+ "shape": [-2],
263+ "format_match_mode": "FormatAgnostic"
264+ },
265+ {
266+ "name": "max_output_size",
267+ "index": 3,
268+ "dtype": "int32",
269+ "format": "ND",
270+ "paramType": "required",
271+ "shape": [-2],
272+ "format_match_mode": "FormatAgnostic"
273+ },
274+ {
275+ "name": "iou_threshold",
276+ "index": 4,
277+ "dtype": "float32",
278+ "format": "ND",
279+ "paramType": "required",
280+ "shape": [-2],
281+ "format_match_mode": "FormatAgnostic"
282+ },
283+ {
284+ "name": "score_threshold",
285+ "index": 5,
286+ "dtype": "float32",
287+ "format": "ND",
288+ "paramType": "required",
289+ "shape": [-2],
290+ "format_match_mode": "FormatAgnostic"
291+ }
292+ ],
293+ "outputs": [
294+ {
295+ "name": "selected_indices",
296+ "index": 0,
297+ "dtype": "int32",
298+ "format": "ND",
299+ "paramType": "required",
300+ "shape": [-2],
301+ "format_match_mode": "FormatAgnostic"
302+ }
303+ ],
304+ "attrs": [
305+ {
306+ "name": "offset",
307+ "dtype": "int",
308+ "value": 0
309+ }
310+ ]
311+ }
312+ ]
313+}
Aobjdetect/sorted_nms/op_host/config/ascend950/sorted_nms_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascnedxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/built-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[SortedNMS]
13+default=0
Aobjdetect/sorted_nms/op_host/sorted_nms_def.cpp+112-0
@@ -0,0 +1,112 @@
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 SortedNMS : public OpDef {
15+public:
16+ explicit SortedNMS(const char* name) : OpDef(name)
17+ {
18+ this->Input("boxes")
19+ .ParamType(REQUIRED)
20+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT})
21+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
22+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
23+ .AutoContiguous();
24+ this->Input("sorted_scores")
25+ .ParamType(REQUIRED)
26+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT})
27+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
28+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
29+ .AutoContiguous();
30+ this->Input("input_indices")
31+ .ParamType(REQUIRED)
32+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
33+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
34+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
35+ .AutoContiguous();
36+ this->Input("max_output_size")
37+ .ParamType(REQUIRED)
38+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
39+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
40+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
41+ .AutoContiguous();
42+ this->Input("iou_threshold")
43+ .ParamType(REQUIRED)
44+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT})
45+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
46+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
47+ .AutoContiguous();
48+ this->Input("score_threshold")
49+ .ParamType(REQUIRED)
50+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT})
51+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
52+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
53+ .AutoContiguous();
54+ this->Output("selected_indices")
55+ .ParamType(REQUIRED)
56+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
57+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
58+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
59+ .AutoContiguous()
60+ .OutputShapeDependOnCompute();
61+ 
62+ this->Attr("offset").AttrType(OPTIONAL).Int(0);
63+ 
64+ OpAICoreConfig aiCoreConfig;
65+ aiCoreConfig.Input("boxes")
66+ .ParamType(REQUIRED)
67+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT})
68+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
69+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
70+ aiCoreConfig.Input("sorted_scores")
71+ .ParamType(REQUIRED)
72+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT})
73+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
74+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
75+ aiCoreConfig.Input("input_indices")
76+ .ParamType(REQUIRED)
77+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
78+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
79+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
80+ aiCoreConfig.Input("max_output_size")
81+ .ParamType(REQUIRED)
82+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
83+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
84+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
85+ aiCoreConfig.Input("iou_threshold")
86+ .ParamType(REQUIRED)
87+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT})
88+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
89+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
90+ aiCoreConfig.Input("score_threshold")
91+ .ParamType(REQUIRED)
92+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT})
93+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
94+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
95+ aiCoreConfig.Output("selected_indices")
96+ .ParamType(REQUIRED)
97+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
98+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
99+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
100+ .OutputShapeDependOnCompute();
101+ aiCoreConfig.DynamicCompileStaticFlag(true)
102+ .DynamicFormatFlag(false)
103+ .DynamicRankSupportFlag(true)
104+ .DynamicShapeSupportFlag(true)
105+ .NeedCheckSupportFlag(false)
106+ .PrecisionReduceFlag(false)
107+ .ExtendCfgInfo("opFile.value", "sorted_nms_apt");
108+ this->AICore().AddConfig("ascend950", aiCoreConfig);
109+ }
110+};
111+OP_ADD(SortedNMS);
112+} // namespace ops
Aobjdetect/sorted_nms/op_host/sorted_nms_infershape.cpp+82-0
@@ -0,0 +1,82 @@
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_impl_registry.h"
12+#include "log/log.h"
13+ 
14+namespace {
15+constexpr int64_t INPUT_BOXES = 0;
16+constexpr int64_t OUTPUT_SELECTED_INDICES = 0;
17+constexpr int64_t BOX_RANK = 2;
18+constexpr int64_t BOX_COORDS = 4;
19+constexpr int64_t OUTPUT_RANK = 1;
20+} // namespace
21+ 
22+namespace ops {
23+static ge::graphStatus InferShapeForSortedNMS(gert::InferShapeContext* context)
24+{
25+ auto boxesShape = context->GetInputShape(INPUT_BOXES);
26+ auto selectedShape = context->GetOutputShape(OUTPUT_SELECTED_INDICES);
27+ if (boxesShape == nullptr || selectedShape == nullptr) {
28+ OP_LOGE(context, "boxes shape or selected_indices shape is nullptr");
29+ return ge::GRAPH_FAILED;
30+ }
31+ 
32+ if (boxesShape->GetDimNum() != BOX_RANK) {
33+ OP_LOGE(context, "boxes rank must be 2");
34+ return ge::GRAPH_FAILED;
35+ }
36+ if (boxesShape->GetDim(1) != ge::UNKNOWN_DIM && boxesShape->GetDim(1) != BOX_COORDS) {
37+ OP_LOGE(context, "boxes second dim must be 4");
38+ return ge::GRAPH_FAILED;
39+ }
40+ 
41+ selectedShape->SetDimNum(OUTPUT_RANK);
42+ selectedShape->SetDim(0, ge::UNKNOWN_DIM);
43+ return ge::GRAPH_SUCCESS;
44+}
45+ 
46+static ge::graphStatus InferShapeRangeForSortedNMS(gert::InferShapeRangeContext* context)
47+{
48+ OP_CHECK_IF(context == nullptr, OP_LOGE("SortedNMS", "InferShapeRangeContext is nullptr"), return ge::GRAPH_FAILED);
49+ 
50+ auto boxesRange = context->GetInputShapeRange(INPUT_BOXES);
51+ auto selectedRange = context->GetOutputShapeRange(OUTPUT_SELECTED_INDICES);
52+ OP_CHECK_NULL_WITH_CONTEXT(context, boxesRange);
53+ OP_CHECK_NULL_WITH_CONTEXT(context, boxesRange->GetMin());
54+ OP_CHECK_NULL_WITH_CONTEXT(context, boxesRange->GetMax());
55+ OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange);
56+ OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange->GetMin());
57+ OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange->GetMax());
58+ 
59+ OP_CHECK_IF(boxesRange->GetMin()->GetDimNum() != BOX_RANK || boxesRange->GetMax()->GetDimNum() != BOX_RANK,
60+ OP_LOGE(context, "boxes shape range rank must be 2"), return ge::GRAPH_FAILED);
61+ const int64_t maxBoxesNum = boxesRange->GetMax()->GetDim(0);
62+ OP_CHECK_IF(maxBoxesNum < 0, OP_LOGE(context, "boxes shape range first dim must be known"),
63+ return ge::GRAPH_FAILED);
64+ 
65+ selectedRange->GetMin()->SetDimNum(OUTPUT_RANK);
66+ selectedRange->GetMin()->SetDim(0, 0);
67+ selectedRange->GetMax()->SetDimNum(OUTPUT_RANK);
68+ selectedRange->GetMax()->SetDim(0, maxBoxesNum);
69+ return ge::GRAPH_SUCCESS;
70+}
71+ 
72+static ge::graphStatus InferDataTypeForSortedNMS(gert::InferDataTypeContext* context)
73+{
74+ context->SetOutputDataType(OUTPUT_SELECTED_INDICES, ge::DT_INT32);
75+ return ge::GRAPH_SUCCESS;
76+}
77+ 
78+IMPL_OP_INFERSHAPE(SortedNMS)
79+ .InferShape(InferShapeForSortedNMS)
80+ .InferShapeRange(InferShapeRangeForSortedNMS)
81+ .InferDataType(InferDataTypeForSortedNMS);
82+} // namespace ops
Aobjdetect/sorted_nms/op_kernel/arch35/sorted_nms_simt.h+503-0
@@ -0,0 +1,503 @@
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 SORTED_NMS_SIMT_H_
12+#define SORTED_NMS_SIMT_H_
13+ 
14+#include "kernel_operator.h"
15+#include "kernel_tiling/kernel_tiling.h"
16+#include "simt_api/common_functions.h"
17+#include "simt_api/math_functions.h"
18+#include "simt_api/asc_fp16.h"
19+#include "sorted_nms_tiling_data.h"
20+ 
21+namespace NsSortedNMS {
22+using namespace AscendC;
23+ 
24+constexpr uint32_t THREAD_NUM = 1024;
25+constexpr uint32_t LOCAL_PAIRWISE_THREAD_NUM = 1024;
26+constexpr uint32_t MIN_THREAD_NUM = 32;
27+constexpr int32_t BOX_COORDS = 4;
28+constexpr int32_t INIT_INDEX = -1;
29+constexpr int32_t NOT_SUPPRESSED = 0;
30+constexpr int32_t SUPPRESSED = 1;
31+constexpr int32_t WORK_SELECTED_COUNT = 0;
32+constexpr int32_t WORK_CURRENT_INDEX = 1;
33+constexpr int32_t OUTPUT_SHAPE_INFO_SIZE = 9;
34+constexpr int32_t MASK_BITS = 32;
35+constexpr int64_t MULTI_CORE_MIN_BOXES = 1025;
36+constexpr int64_t MULTI_CORE_MAX_BOXES = 8192;
37+constexpr uint32_t UB_BLOCK_SIZE = 32;
38+// shape_out uses uint64_t entries. Bit 31 marks the rank field as uint64_t encoded.
39+constexpr uint64_t OUTPUT_SHAPE_RANK_ONE = 0x80000001ULL;
40+ 
41+__aicore__ inline uint32_t AlignUbBytes(uint32_t bytes)
42+{
43+ return (bytes + UB_BLOCK_SIZE - 1U) / UB_BLOCK_SIZE * UB_BLOCK_SIZE;
44+}
45+ 
46+__aicore__ inline uint32_t GetSimtThreadNum(int64_t workItems)
47+{
48+ uint32_t threadNum = MIN_THREAD_NUM;
49+ while (static_cast<int64_t>(threadNum) < workItems && threadNum < THREAD_NUM) {
50+ threadNum <<= 1;
51+ }
52+ return threadNum;
53+}
54+ 
55+__simt_callee__ __aicore__ inline void WriteOutputShape(__gm__ uint64_t* outputShape, int32_t selectedCount)
56+{
57+ for (int32_t index = 0; index < OUTPUT_SHAPE_INFO_SIZE; ++index) {
58+ outputShape[index] = 0;
59+ }
60+ outputShape[0] = OUTPUT_SHAPE_RANK_ONE;
61+ outputShape[1] = static_cast<uint64_t>(selectedCount);
62+}
63+ 
64+__simt_callee__ __aicore__ inline float ReadAsFloat(float val) { return val; }
65+ 
66+__simt_callee__ __aicore__ inline float ReadAsFloat(half val) { return __half2float(val); }
67+ 
68+template <typename TScore, typename TScoreThreshold>
69+__simt_vf__ __aicore__ inline void FindActiveBoxesNum(int64_t boxesNum, __gm__ TScore* sortedScores,
70+ __gm__ TScoreThreshold* scoreThreshold,
71+ __gm__ int32_t* activeBoxesNum)
72+{
73+ if (threadIdx.x != 0) {
74+ return;
75+ }
76+ const float scoreThr = ReadAsFloat(scoreThreshold[0]);
77+ int64_t left = 0;
78+ int64_t right = boxesNum;
79+ while (left < right) {
80+ const int64_t middle = left + (right - left) / 2;
81+ if (ReadAsFloat(sortedScores[middle]) > scoreThr) {
82+ left = middle + 1;
83+ } else {
84+ right = middle;
85+ }
86+ }
87+ activeBoxesNum[0] = static_cast<int32_t>(left);
88+}
89+ 
90+__simt_callee__ __aicore__ inline float ClampNonNegative(float val) { return val > 0.0f ? val : 0.0f; }
91+ 
92+template <typename T>
93+__simt_callee__ __aicore__ inline float BoxArea(__gm__ T* boxes, int32_t boxIdx, float offset)
94+{
95+ int64_t base = static_cast<int64_t>(boxIdx) * BOX_COORDS;
96+ float x1 = ReadAsFloat(boxes[base]);
97+ float y1 = ReadAsFloat(boxes[base + 1]);
98+ float x2 = ReadAsFloat(boxes[base + 2]);
99+ float y2 = ReadAsFloat(boxes[base + 3]);
100+ float width = ClampNonNegative(x2 - x1 + offset);
101+ float height = ClampNonNegative(y2 - y1 + offset);
102+ return width * height;
103+}
104+ 
105+template <typename T>
106+__simt_callee__ __aicore__ inline float Intersection(__gm__ T* boxes, int32_t lhs, int32_t rhs, float offset)
107+{
108+ int64_t lhsBase = static_cast<int64_t>(lhs) * BOX_COORDS;
109+ int64_t rhsBase = static_cast<int64_t>(rhs) * BOX_COORDS;
110+ float x1 = fmaxf(ReadAsFloat(boxes[lhsBase]), ReadAsFloat(boxes[rhsBase]));
111+ float y1 = fmaxf(ReadAsFloat(boxes[lhsBase + 1]), ReadAsFloat(boxes[rhsBase + 1]));
112+ float x2 = fminf(ReadAsFloat(boxes[lhsBase + 2]), ReadAsFloat(boxes[rhsBase + 2]));
113+ float y2 = fminf(ReadAsFloat(boxes[lhsBase + 3]), ReadAsFloat(boxes[rhsBase + 3]));
114+ float width = ClampNonNegative(x2 - x1 + offset);
115+ float height = ClampNonNegative(y2 - y1 + offset);
116+ return width * height;
117+}
118+ 
119+template <typename T>
120+__simt_callee__ __aicore__ inline float Iou(__gm__ T* boxes, int32_t lhs, int32_t rhs, float offset)
121+{
122+ float inter = Intersection(boxes, lhs, rhs, offset);
123+ float lhsArea = BoxArea(boxes, lhs, offset);
124+ float rhsArea = BoxArea(boxes, rhs, offset);
125+ float denom = lhsArea + rhsArea - inter;
126+ if (denom <= 0.0f) {
127+ return 0.0f;
128+ }
129+ return inter / denom;
130+}
131+ 
132+template <typename T>
133+__simt_callee__ __aicore__ inline float IntersectionLocal(__ubuf__ T* boxes, int32_t lhs, int32_t rhs, float offset)
134+{
135+ int64_t lhsBase = static_cast<int64_t>(lhs) * BOX_COORDS;
136+ int64_t rhsBase = static_cast<int64_t>(rhs) * BOX_COORDS;
137+ float x1 = fmaxf(ReadAsFloat(boxes[lhsBase]), ReadAsFloat(boxes[rhsBase]));
138+ float y1 = fmaxf(ReadAsFloat(boxes[lhsBase + 1]), ReadAsFloat(boxes[rhsBase + 1]));
139+ float x2 = fminf(ReadAsFloat(boxes[lhsBase + 2]), ReadAsFloat(boxes[rhsBase + 2]));
140+ float y2 = fminf(ReadAsFloat(boxes[lhsBase + 3]), ReadAsFloat(boxes[rhsBase + 3]));
141+ float width = ClampNonNegative(x2 - x1 + offset);
142+ float height = ClampNonNegative(y2 - y1 + offset);
143+ return width * height;
144+}
145+ 
146+template <typename TBox>
147+__simt_vf__ __aicore__ __launch_bounds__(THREAD_NUM) inline void BuildBoxAreas(int64_t boxesNum, int32_t offset,
148+ __ubuf__ TBox* boxes,
149+ __ubuf__ float* boxAreas)
150+{
151+ const float offsetVal = static_cast<float>(offset);
152+ for (int64_t boxIndex = threadIdx.x; boxIndex < boxesNum; boxIndex += blockDim.x) {
153+ const int64_t base = boxIndex * BOX_COORDS;
154+ const float width = ClampNonNegative(ReadAsFloat(boxes[base + 2]) - ReadAsFloat(boxes[base]) + offsetVal);
155+ const float height = ClampNonNegative(ReadAsFloat(boxes[base + 3]) - ReadAsFloat(boxes[base + 1]) + offsetVal);
156+ boxAreas[boxIndex] = width * height;
157+ }
158+}
159+ 
160+template <typename T>
161+__simt_callee__ __aicore__ inline float IouLocal(__ubuf__ T* boxes, __ubuf__ float* boxAreas, int32_t lhs, int32_t rhs,
162+ float offset)
163+{
164+ const float inter = IntersectionLocal(boxes, lhs, rhs, offset);
165+ const float denom = boxAreas[lhs] + boxAreas[rhs] - inter;
166+ if (denom <= 0.0f) {
167+ return 0.0f;
168+ }
169+ return inter / denom;
170+}
171+ 
172+template <typename TBox, typename TScore, typename TIouThreshold, typename TScoreThreshold>
173+__simt_vf__ __aicore__ __launch_bounds__(THREAD_NUM) inline void SortedNMSSingleCore(
174+ int64_t boxesNum, int32_t offset, __gm__ TBox* boxes, __gm__ TScore* sortedScores, __gm__ int32_t* inputIndices,
175+ __gm__ int32_t* maxOutputSize, __gm__ TIouThreshold* iouThreshold, __gm__ TScoreThreshold* scoreThreshold,
176+ __gm__ int32_t* selectedIndices, __gm__ uint64_t* outputShape, __gm__ int32_t* work)
177+{
178+ if (boxesNum <= 0) {
179+ if (threadIdx.x == 0) {
180+ WriteOutputShape(outputShape, 0);
181+ }
182+ return;
183+ }
184+ 
185+ __gm__ int32_t* control = work;
186+ __gm__ int32_t* suppressed = work + 2;
187+ 
188+ for (int64_t idx = threadIdx.x; idx < boxesNum; idx += blockDim.x) {
189+ suppressed[idx] = NOT_SUPPRESSED;
190+ }
191+ if (threadIdx.x == 0) {
192+ control[WORK_SELECTED_COUNT] = 0;
193+ control[WORK_CURRENT_INDEX] = INIT_INDEX;
194+ }
195+ asc_threadfence_block();
196+ asc_syncthreads();
197+ 
198+ int32_t maxOut = maxOutputSize[0];
199+ if (maxOut < 0) {
200+ maxOut = 0;
201+ }
202+ if (static_cast<int64_t>(maxOut) > boxesNum) {
203+ maxOut = static_cast<int32_t>(boxesNum);
204+ }
205+ float iouThr = ReadAsFloat(iouThreshold[0]);
206+ float scoreThr = ReadAsFloat(scoreThreshold[0]);
207+ float offsetVal = static_cast<float>(offset);
208+ 
209+ for (int64_t sortedPos = 0; sortedPos < boxesNum; ++sortedPos) {
210+ if (threadIdx.x == 0) {
211+ control[WORK_CURRENT_INDEX] = INIT_INDEX;
212+ int32_t selectedCount = control[WORK_SELECTED_COUNT];
213+ if (selectedCount < maxOut && suppressed[sortedPos] == NOT_SUPPRESSED) {
214+ float score = ReadAsFloat(sortedScores[sortedPos]);
215+ int32_t current = inputIndices[sortedPos];
216+ if (score > scoreThr && current >= 0 && static_cast<int64_t>(current) < boxesNum) {
217+ selectedIndices[selectedCount] = current;
218+ control[WORK_SELECTED_COUNT] = selectedCount + 1;
219+ control[WORK_CURRENT_INDEX] = current;
220+ }
221+ }
222+ }
223+ asc_threadfence_block();
224+ asc_syncthreads();
225+ 
226+ int32_t currentIndex = control[WORK_CURRENT_INDEX];
227+ if (currentIndex >= 0) {
228+ for (int64_t nextPos = sortedPos + 1 + threadIdx.x; nextPos < boxesNum; nextPos += blockDim.x) {
229+ if (suppressed[nextPos] == NOT_SUPPRESSED) {
230+ float score = ReadAsFloat(sortedScores[nextPos]);
231+ int32_t nextIndex = inputIndices[nextPos];
232+ if (score > scoreThr && nextIndex >= 0 && static_cast<int64_t>(nextIndex) < boxesNum) {
233+ float overlap = Iou(boxes, currentIndex, nextIndex, offsetVal);
234+ if (overlap > iouThr) {
235+ suppressed[nextPos] = SUPPRESSED;
236+ }
237+ }
238+ }
239+ }
240+ }
241+ asc_threadfence_block();
242+ asc_syncthreads();
243+ 
244+ if (control[WORK_SELECTED_COUNT] >= maxOut) {
245+ break;
246+ }
247+ }
248+ if (threadIdx.x == 0) {
249+ WriteOutputShape(outputShape, control[WORK_SELECTED_COUNT]);
250+ }
251+}
252+ 
253+template <typename TBox, typename TScore, typename TIouThreshold, typename TScoreThreshold>
254+__simt_vf__ __aicore__ __launch_bounds__(THREAD_NUM) inline void BuildPairwiseMasksLocal(
255+ int64_t boxesNum, int64_t activeBoxesNum, int64_t activeMaskWordNum, int32_t offset, int32_t coreIdx,
256+ int32_t coreNum, __ubuf__ TBox* boxes, __gm__ TScore* sortedScores, __gm__ int32_t* inputIndices,
257+ __ubuf__ float* boxAreas, __gm__ TIouThreshold* iouThreshold, __gm__ TScoreThreshold* scoreThreshold,
258+ __gm__ uint32_t* pairwiseMasks)
259+{
260+ float iouThr = ReadAsFloat(iouThreshold[0]);
261+ float offsetVal = static_cast<float>(offset);
262+ const int64_t globalThreadIdx = static_cast<int64_t>(coreIdx) * blockDim.x + threadIdx.x;
263+ const int64_t globalThreadNum = static_cast<int64_t>(coreNum) * blockDim.x;
264+ const int64_t pairwiseMaskWords = activeBoxesNum * activeMaskWordNum;
265+ 
266+ for (int64_t task = globalThreadIdx; task < pairwiseMaskWords; task += globalThreadNum) {
267+ const int64_t sortedPos = task / activeMaskWordNum;
268+ const int64_t wordIndex = task - sortedPos * activeMaskWordNum;
269+ const int64_t firstRelevantWord = (sortedPos + 1) / MASK_BITS;
270+ if (wordIndex < firstRelevantWord) {
271+ continue;
272+ }
273+ const int64_t firstNextPos = wordIndex * MASK_BITS;
274+ uint32_t mask = 0;
275+ const int32_t currentIndex = inputIndices[sortedPos];
276+ if (currentIndex >= 0 && static_cast<int64_t>(currentIndex) < boxesNum) {
277+ for (int32_t bit = 0; bit < MASK_BITS; ++bit) {
278+ const int64_t nextPos = firstNextPos + bit;
279+ if (nextPos <= sortedPos || nextPos >= activeBoxesNum) {
280+ continue;
281+ }
282+ const int32_t nextIndex = inputIndices[nextPos];
283+ if (nextIndex >= 0 && static_cast<int64_t>(nextIndex) < boxesNum &&
284+ IouLocal(boxes, boxAreas, currentIndex, nextIndex, offsetVal) > iouThr) {
285+ mask |= 1U << bit;
286+ }
287+ }
288+ }
289+ pairwiseMasks[task] = mask;
290+ }
291+}
292+ 
293+template <typename TBox, typename TScore, typename TIouThreshold, typename TScoreThreshold>
294+__simt_vf__ __aicore__ __launch_bounds__(THREAD_NUM) inline void BuildPairwiseMasks(
295+ int64_t boxesNum, int64_t activeBoxesNum, int64_t activeMaskWordNum, int32_t offset, int32_t coreIdx,
296+ int32_t coreNum, __gm__ TBox* boxes, __gm__ TScore* sortedScores, __gm__ int32_t* inputIndices,
297+ __gm__ TIouThreshold* iouThreshold, __gm__ TScoreThreshold* scoreThreshold, __gm__ uint32_t* pairwiseMasks)
298+{
299+ float iouThr = ReadAsFloat(iouThreshold[0]);
300+ float offsetVal = static_cast<float>(offset);
301+ const int64_t globalThreadIdx = static_cast<int64_t>(coreIdx) * blockDim.x + threadIdx.x;
302+ const int64_t globalThreadNum = static_cast<int64_t>(coreNum) * blockDim.x;
303+ const int64_t pairwiseMaskWords = activeBoxesNum * activeMaskWordNum;
304+ 
305+ for (int64_t task = globalThreadIdx; task < pairwiseMaskWords; task += globalThreadNum) {
306+ const int64_t sortedPos = task / activeMaskWordNum;
307+ const int64_t wordIndex = task - sortedPos * activeMaskWordNum;
308+ const int64_t firstRelevantWord = (sortedPos + 1) / MASK_BITS;
309+ if (wordIndex < firstRelevantWord) {
310+ continue;
311+ }
312+ const int64_t firstNextPos = wordIndex * MASK_BITS;
313+ uint32_t mask = 0;
314+ const int32_t currentIndex = inputIndices[sortedPos];
315+ if (currentIndex >= 0 && static_cast<int64_t>(currentIndex) < boxesNum) {
316+ for (int32_t bit = 0; bit < MASK_BITS; ++bit) {
317+ const int64_t nextPos = firstNextPos + bit;
318+ if (nextPos <= sortedPos || nextPos >= activeBoxesNum) {
319+ continue;
320+ }
321+ const int32_t nextIndex = inputIndices[nextPos];
322+ if (nextIndex >= 0 && static_cast<int64_t>(nextIndex) < boxesNum &&
323+ Iou(boxes, currentIndex, nextIndex, offsetVal) > iouThr) {
324+ mask |= 1U << bit;
325+ }
326+ }
327+ }
328+ pairwiseMasks[task] = mask;
329+ }
330+}
331+ 
332+template <typename TScore, typename TScoreThreshold>
333+__simt_vf__ __aicore__ __launch_bounds__(THREAD_NUM) inline void SelectFromPairwiseMasks(
334+ int64_t boxesNum, int64_t activeBoxesNum, int64_t activeMaskWordNum, __gm__ TScore* sortedScores,
335+ __gm__ int32_t* inputIndices, __gm__ int32_t* maxOutputSize, __gm__ TScoreThreshold* scoreThreshold,
336+ __gm__ int32_t* selectedIndices, __gm__ uint64_t* outputShape, __gm__ int32_t* work)
337+{
338+ __gm__ int32_t* control = work;
339+ __gm__ uint32_t* suppressedMasks = (__gm__ uint32_t*)(work + 2);
340+ __gm__ uint32_t* pairwiseMasks = suppressedMasks + activeMaskWordNum;
341+ 
342+ for (int64_t wordIndex = threadIdx.x; wordIndex < activeMaskWordNum; wordIndex += blockDim.x) {
343+ suppressedMasks[wordIndex] = 0;
344+ }
345+ if (threadIdx.x == 0) {
346+ control[WORK_SELECTED_COUNT] = 0;
347+ control[WORK_CURRENT_INDEX] = INIT_INDEX;
348+ }
349+ asc_threadfence_block();
350+ asc_syncthreads();
351+ 
352+ int32_t maxOut = maxOutputSize[0];
353+ maxOut = maxOut < 0 ? 0 : maxOut;
354+ if (static_cast<int64_t>(maxOut) > boxesNum) {
355+ maxOut = static_cast<int32_t>(boxesNum);
356+ }
357+ if (maxOut == 0 || activeBoxesNum == 0) {
358+ if (threadIdx.x == 0) {
359+ WriteOutputShape(outputShape, 0);
360+ }
361+ return;
362+ }
363+ 
364+ for (int64_t sortedPos = 0; sortedPos < activeBoxesNum; ++sortedPos) {
365+ if (threadIdx.x == 0) {
366+ control[WORK_CURRENT_INDEX] = INIT_INDEX;
367+ const int64_t wordIndex = sortedPos / MASK_BITS;
368+ const int32_t bitIndex = static_cast<int32_t>(sortedPos - wordIndex * MASK_BITS);
369+ const bool isSuppressed = (suppressedMasks[wordIndex] & (1U << bitIndex)) != 0;
370+ const int32_t currentIndex = inputIndices[sortedPos];
371+ if (!isSuppressed && currentIndex >= 0 && static_cast<int64_t>(currentIndex) < boxesNum) {
372+ const int32_t selectedCount = control[WORK_SELECTED_COUNT];
373+ selectedIndices[selectedCount] = currentIndex;
374+ control[WORK_SELECTED_COUNT] = selectedCount + 1;
375+ control[WORK_CURRENT_INDEX] = static_cast<int32_t>(sortedPos);
376+ }
377+ }
378+ asc_threadfence_block();
379+ asc_syncthreads();
380+ 
381+ const int32_t selectedRow = control[WORK_CURRENT_INDEX];
382+ if (selectedRow >= 0) {
383+ const int64_t rowOffset = static_cast<int64_t>(selectedRow) * activeMaskWordNum;
384+ const int64_t firstRelevantWord = (static_cast<int64_t>(selectedRow) + 1) / MASK_BITS;
385+ for (int64_t wordIndex = firstRelevantWord + threadIdx.x; wordIndex < activeMaskWordNum;
386+ wordIndex += blockDim.x) {
387+ suppressedMasks[wordIndex] |= pairwiseMasks[rowOffset + wordIndex];
388+ }
389+ }
390+ asc_threadfence_block();
391+ asc_syncthreads();
392+ 
393+ if (control[WORK_SELECTED_COUNT] >= maxOut) {
394+ break;
395+ }
396+ }
397+ if (threadIdx.x == 0) {
398+ WriteOutputShape(outputShape, control[WORK_SELECTED_COUNT]);
399+ }
400+}
401+ 
402+template <typename TBox, typename TScore, typename TIouThreshold, typename TScoreThreshold>
403+__aicore__ inline void Process(GM_ADDR boxes, GM_ADDR sortedScores, GM_ADDR inputIndices, GM_ADDR maxOutputSize,
404+ GM_ADDR iouThreshold, GM_ADDR scoreThreshold, GM_ADDR selectedIndices,
405+ GM_ADDR outputShape, GM_ADDR workspace, const SortedNMSTilingData* tilingData,
406+ TPipe* pipe)
407+{
408+ __gm__ TBox* boxesGm = (__gm__ TBox*)boxes;
409+ __gm__ TScore* sortedScoresGm = (__gm__ TScore*)sortedScores;
410+ __gm__ int32_t* inputIndicesGm = (__gm__ int32_t*)inputIndices;
411+ __gm__ int32_t* maxOutputSizeGm = (__gm__ int32_t*)maxOutputSize;
412+ __gm__ TIouThreshold* iouThresholdGm = (__gm__ TIouThreshold*)iouThreshold;
413+ __gm__ TScoreThreshold* scoreThresholdGm = (__gm__ TScoreThreshold*)scoreThreshold;
414+ __gm__ int32_t* selectedIndicesGm = (__gm__ int32_t*)selectedIndices;
415+ __gm__ uint64_t* outputShapeGm = (__gm__ uint64_t*)outputShape;
416+ __gm__ int32_t* workGm = (__gm__ int32_t*)workspace;
417+ const uint32_t singleCoreThreadNum = GetSimtThreadNum(tilingData->boxesNum);
418+ 
419+ if (tilingData->boxesNum >= MULTI_CORE_MIN_BOXES && tilingData->boxesNum <= MULTI_CORE_MAX_BOXES) {
420+ const int32_t coreIdx = static_cast<int32_t>(GetBlockIdx());
421+ const int64_t maskWordNum = (tilingData->boxesNum + MASK_BITS - 1) / MASK_BITS;
422+ int32_t maxOut = maxOutputSizeGm[0];
423+ maxOut = maxOut < 0 ? 0 : maxOut;
424+ if (static_cast<int64_t>(maxOut) > tilingData->boxesNum) {
425+ maxOut = static_cast<int32_t>(tilingData->boxesNum);
426+ }
427+ const int64_t pairwiseBreakEven = (tilingData->boxesNum + tilingData->coreNum - 1) / tilingData->coreNum;
428+ if (static_cast<int64_t>(maxOut) <= pairwiseBreakEven) {
429+ if (coreIdx == 0) {
430+ asc_vf_call<SortedNMSSingleCore<TBox, TScore, TIouThreshold, TScoreThreshold>>(
431+ dim3(singleCoreThreadNum), tilingData->boxesNum, tilingData->offset, boxesGm, sortedScoresGm,
432+ inputIndicesGm, maxOutputSizeGm, iouThresholdGm, scoreThresholdGm, selectedIndicesGm, outputShapeGm,
433+ workGm);
434+ }
435+ return;
436+ }
437+ 
438+ __gm__ uint32_t* suppressedMasks = (__gm__ uint32_t*)(workGm + 2);
439+ if (coreIdx == 0) {
440+ asc_vf_call<FindActiveBoxesNum<TScore, TScoreThreshold>>(dim3(1), tilingData->boxesNum, sortedScoresGm,
441+ scoreThresholdGm, workGm);
442+ }
443+ SyncAll();
444+ const int64_t activeBoxesNum = static_cast<int64_t>(workGm[0]);
445+ const int64_t activeMaskWordNum = (activeBoxesNum + MASK_BITS - 1) / MASK_BITS;
446+ __gm__ uint32_t* pairwiseMasks = suppressedMasks + activeMaskWordNum;
447+ if (tilingData->useLocalBoxes != 0) {
448+ const uint32_t boxesBytes = AlignUbBytes(
449+ static_cast<uint32_t>(tilingData->boxesNum * BOX_COORDS * sizeof(TBox)));
450+ const uint32_t areasBytes = AlignUbBytes(static_cast<uint32_t>(tilingData->boxesNum * sizeof(float)));
451+ TQue<QuePosition::VECIN, 1> boxesQueue;
452+ TBuf<TPosition::VECCALC> areasBuffer;
453+ pipe->InitBuffer(boxesQueue, 1, boxesBytes);
454+ pipe->InitBuffer(areasBuffer, areasBytes);
455+ 
456+ GlobalTensor<TBox> boxesTensor;
457+ boxesTensor.SetGlobalBuffer(boxesGm, tilingData->boxesNum * BOX_COORDS);
458+ LocalTensor<TBox> boxesLocal = boxesQueue.AllocTensor<TBox>();
459+ const DataCopyExtParams boxesCopyParams{
460+ 1, static_cast<uint32_t>(tilingData->boxesNum * BOX_COORDS * sizeof(TBox)), 0, 0, 0};
461+ DataCopyPad(boxesLocal, boxesTensor, boxesCopyParams, DataCopyPadExtParams<TBox>{false, 0, 0, 0});
462+ boxesQueue.EnQue(boxesLocal);
463+ boxesLocal = boxesQueue.DeQue<TBox>();
464+ LocalTensor<float> areasLocal = areasBuffer.Get<float>();
465+ __ubuf__ TBox* boxesUb = reinterpret_cast<__ubuf__ TBox*>(boxesLocal.GetPhyAddr());
466+ __ubuf__ float* areasUb = reinterpret_cast<__ubuf__ float*>(areasLocal.GetPhyAddr());
467+ 
468+ asc_vf_call<BuildBoxAreas<TBox>>(dim3(GetSimtThreadNum(tilingData->boxesNum)), tilingData->boxesNum,
469+ tilingData->offset, boxesUb, areasUb);
470+ event_t areasReady = static_cast<event_t>(pipe->FetchEventID(HardEvent::V_S));
471+ SetFlag<HardEvent::V_S>(areasReady);
472+ WaitFlag<HardEvent::V_S>(areasReady);
473+ asc_vf_call<BuildPairwiseMasksLocal<TBox, TScore, TIouThreshold, TScoreThreshold>>(
474+ dim3(LOCAL_PAIRWISE_THREAD_NUM), tilingData->boxesNum, activeBoxesNum, activeMaskWordNum,
475+ tilingData->offset, coreIdx, tilingData->coreNum, boxesUb, sortedScoresGm, inputIndicesGm, areasUb,
476+ iouThresholdGm, scoreThresholdGm, pairwiseMasks);
477+ event_t pairwiseDone = static_cast<event_t>(pipe->FetchEventID(HardEvent::V_S));
478+ SetFlag<HardEvent::V_S>(pairwiseDone);
479+ WaitFlag<HardEvent::V_S>(pairwiseDone);
480+ boxesQueue.FreeTensor(boxesLocal);
481+ } else {
482+ asc_vf_call<BuildPairwiseMasks<TBox, TScore, TIouThreshold, TScoreThreshold>>(
483+ dim3(THREAD_NUM), tilingData->boxesNum, activeBoxesNum, activeMaskWordNum, tilingData->offset, coreIdx,
484+ tilingData->coreNum, boxesGm, sortedScoresGm, inputIndicesGm, iouThresholdGm, scoreThresholdGm,
485+ pairwiseMasks);
486+ }
487+ SyncAll();
488+ if (coreIdx == 0) {
489+ const uint32_t selectThreadNum = GetSimtThreadNum(activeMaskWordNum);
490+ asc_vf_call<SelectFromPairwiseMasks<TScore, TScoreThreshold>>(
491+ dim3(selectThreadNum), tilingData->boxesNum, activeBoxesNum, activeMaskWordNum, sortedScoresGm,
492+ inputIndicesGm, maxOutputSizeGm, scoreThresholdGm, selectedIndicesGm, outputShapeGm, workGm);
493+ }
494+ return;
495+ }
496+ 
497+ asc_vf_call<SortedNMSSingleCore<TBox, TScore, TIouThreshold, TScoreThreshold>>(
498+ dim3(singleCoreThreadNum), tilingData->boxesNum, tilingData->offset, boxesGm, sortedScoresGm, inputIndicesGm,
499+ maxOutputSizeGm, iouThresholdGm, scoreThresholdGm, selectedIndicesGm, outputShapeGm, workGm);
500+}
501+} // namespace NsSortedNMS
502+ 
503+#endif // SORTED_NMS_SIMT_H_
Aobjdetect/sorted_nms/op_kernel/arch35/sorted_nms_tiling_data.h+23-0
@@ -0,0 +1,23 @@
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 SORTED_NMS_TILING_DATA_H_
12+#define SORTED_NMS_TILING_DATA_H_
13+ 
14+#include <cstdint>
15+ 
16+struct SortedNMSTilingData {
17+ int64_t boxesNum = 0;
18+ int32_t offset = 0;
19+ int32_t coreNum = 1;
20+ int32_t useLocalBoxes = 0;
21+};
22+ 
23+#endif // SORTED_NMS_TILING_DATA_H_
Aobjdetect/sorted_nms/op_kernel/sorted_nms_apt.cpp+35-0
@@ -0,0 +1,35 @@
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/sorted_nms_simt.h"
12+ 
13+extern "C" __global__ __aicore__ void sorted_nms(GM_ADDR boxes, GM_ADDR sorted_scores, GM_ADDR input_indices,
14+ GM_ADDR max_output_size, GM_ADDR iou_threshold,
15+ GM_ADDR score_threshold, GM_ADDR selected_indices, GM_ADDR shape_out,
16+ GM_ADDR workspace, GM_ADDR tiling)
17+{
18+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
19+ if (workspace == nullptr) {
20+ return;
21+ }
22+ AscendC::SetSysWorkspace(workspace);
23+ GM_ADDR userWorkspace = AscendC::GetUserWorkspace(workspace);
24+ if (userWorkspace == nullptr) {
25+ return;
26+ }
27+ 
28+ REGISTER_TILING_DEFAULT(SortedNMSTilingData);
29+ GET_TILING_DATA_WITH_STRUCT(SortedNMSTilingData, tilingData, tiling);
30+ 
31+ AscendC::TPipe pipe;
32+ NsSortedNMS::Process<DTYPE_BOXES, DTYPE_SORTED_SCORES, DTYPE_IOU_THRESHOLD, DTYPE_SCORE_THRESHOLD>(
33+ boxes, sorted_scores, input_indices, max_output_size, iou_threshold, score_threshold, selected_indices,
34+ shape_out, userWorkspace, &tilingData, &pipe);
35+}
Aobjdetect/sorted_nms/tests/assets/golden.py+153-0
@@ -0,0 +1,153 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+ 
11+import importlib.util
12+from pathlib import Path
13+ 
14+import numpy as np
15+ 
16+ 
17+def _load_customize_inputs():
18+ input_path = Path(__file__).with_name("input.py")
19+ spec = importlib.util.spec_from_file_location("sorted_nms_test_input", input_path)
20+ module = importlib.util.module_from_spec(spec)
21+ spec.loader.exec_module(module)
22+ return module.customize_inputs
23+ 
24+ 
25+__spec__ = {"sorted_nms": "SortedNmsTestSpec"}
26+ 
27+ 
28+class SortedNmsTestSpec:
29+ """Deterministic input generation and a sequential SortedNMS reference."""
30+ 
31+ tolerance = {"int32": {"standard": "binary_equal"}}
32+ customize_inputs = staticmethod(_load_customize_inputs())
33+ 
34+ @staticmethod
35+ def compare(*values, **kwargs):
36+ """Compare the logical output and the meaningful dynamic-shape metadata."""
37+ if len(values) != 4:
38+ return {
39+ "pass": False,
40+ "precision": "INVALID_OUTPUT_COUNT",
41+ "error_info": f"expected 2 outputs and 2 goldens, got {len(values)} values",
42+ }
43+ 
44+ selected_output, shape_output, selected_golden, shape_golden = values
45+ output_shape = np.asarray(shape_output, dtype=np.uint64).reshape(-1)
46+ golden_shape = np.asarray(shape_golden, dtype=np.uint64).reshape(-1)
47+ if output_shape.size < 2 or golden_shape.size < 2:
48+ return {
49+ "pass": False,
50+ "precision": "INVALID_SHAPE_METADATA",
51+ "error_info": "dynamic-shape metadata must contain rank and at least one dimension",
52+ }
53+ 
54+ # TTK marks the uint64 encoding in bit 31 of the first metadata word.
55+ rank_mask = np.uint64(0x7FFFFFFF)
56+ output_rank = int(output_shape[0] & rank_mask)
57+ golden_rank = int(golden_shape[0] & rank_mask)
58+ output_count = int(output_shape[1])
59+ golden_count = int(golden_shape[1])
60+ output_values = np.asarray(selected_output).reshape(-1)
61+ golden_values = np.asarray(selected_golden).reshape(-1)
62+ 
63+ errors = []
64+ if output_rank != 1 or golden_rank != 1:
65+ errors.append(f"rank mismatch: output={output_rank}, golden={golden_rank}")
66+ if output_count != golden_count or golden_count != golden_values.size:
67+ errors.append(
68+ f"selected count mismatch: output={output_count}, "
69+ f"metadata_golden={golden_count}, golden={golden_values.size}"
70+ )
71+ if output_count > output_values.size:
72+ errors.append(
73+ f"selected count {output_count} exceeds output buffer {output_values.size}"
74+ )
75+ elif not np.array_equal(output_values[:output_count], golden_values):
76+ errors.append("selected_indices differ from golden")
77+ 
78+ passed = not errors
79+ return {
80+ "pass": passed,
81+ "precision": "BINARY_EQUAL" if passed else "MISMATCH",
82+ "error_info": "; ".join(errors),
83+ }
84+ 
85+ @staticmethod
86+ def golden(
87+ boxes,
88+ sorted_scores,
89+ input_indices,
90+ max_output_size,
91+ iou_threshold,
92+ score_threshold,
93+ *,
94+ offset=0,
95+ **kwargs,
96+ ):
97+ boxes_num = boxes.shape[0]
98+ if sorted_scores.size > 1 and np.any(sorted_scores[:-1] < sorted_scores[1:]):
99+ raise ValueError("sorted_scores must be sorted in non-increasing order")
100+ max_out = min(max(int(max_output_size[0]), 0), boxes_num)
101+ iou_thr = float(iou_threshold[0])
102+ score_thr = float(score_threshold[0])
103+ boxes_f32 = boxes.astype(np.float32)
104+ suppressed = np.zeros(boxes_num, dtype=np.bool_)
105+ selected = []
106+ 
107+ for sorted_pos in range(boxes_num):
108+ if len(selected) >= max_out:
109+ break
110+ score = float(sorted_scores[sorted_pos])
111+ current = int(input_indices[sorted_pos])
112+ if (
113+ suppressed[sorted_pos]
114+ or score <= score_thr
115+ or current < 0
116+ or current >= boxes_num
117+ ):
118+ continue
119+ 
120+ selected.append(current)
121+ current_box = boxes_f32[current]
122+ for next_pos in range(sorted_pos + 1, boxes_num):
123+ if suppressed[next_pos] or float(sorted_scores[next_pos]) <= score_thr:
124+ continue
125+ next_index = int(input_indices[next_pos])
126+ if next_index < 0 or next_index >= boxes_num:
127+ continue
128+ next_box = boxes_f32[next_index]
129+ width = max(
130+ 0.0,
131+ min(current_box[2], next_box[2])
132+ - max(current_box[0], next_box[0])
133+ + offset,
134+ )
135+ height = max(
136+ 0.0,
137+ min(current_box[3], next_box[3])
138+ - max(current_box[1], next_box[1])
139+ + offset,
140+ )
141+ intersection = width * height
142+ current_area = max(0.0, current_box[2] - current_box[0] + offset) * max(
143+ 0.0, current_box[3] - current_box[1] + offset
144+ )
145+ next_area = max(0.0, next_box[2] - next_box[0] + offset) * max(
146+ 0.0, next_box[3] - next_box[1] + offset
147+ )
148+ union = current_area + next_area - intersection
149+ overlap = intersection / union if union > 0.0 else 0.0
150+ if overlap > iou_thr:
151+ suppressed[next_pos] = True
152+ 
153+ return [np.asarray(selected, dtype=np.int32)]
Aobjdetect/sorted_nms/tests/assets/input.py+51-0
@@ -0,0 +1,51 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+ 
11+import numpy as np
12+ 
13+ 
14+def customize_inputs(
15+ boxes,
16+ sorted_scores,
17+ input_indices,
18+ max_output_size,
19+ iou_threshold,
20+ score_threshold,
21+ **kwargs,
22+):
23+ """Generate deterministic SortedNMS inputs with retain and suppress paths."""
24+ boxes_num = boxes.shape[0]
25+ box_ids = np.arange(boxes_num, dtype=np.int64)
26+ group_ids = box_ids // 4
27+ 
28+ # Four boxes in each group are identical, while different groups are
29+ # spatially separated. This guarantees both suppress and retain paths.
30+ x1 = (group_ids % 64) * 8
31+ y1 = (group_ids // 64) * 8
32+ generated_boxes = np.stack((x1, y1, x1 + 4, y1 + 4), axis=1)
33+ boxes[...] = generated_boxes.astype(boxes.dtype)
34+ 
35+ # Scores are already sorted by position as required by SortedNMS.
36+ if boxes_num > 0:
37+ scores = np.linspace(0.99, 0.01, boxes_num, dtype=np.float32)
38+ sorted_scores[...] = scores.astype(sorted_scores.dtype)
39+ # A fixed coprime affine permutation exercises the input-index map.
40+ multiplier = boxes_num - 1 if boxes_num % 2 == 0 else boxes_num - 2
41+ multiplier = max(multiplier, 1)
42+ input_indices[...] = ((box_ids * multiplier + 1) % boxes_num).astype(np.int32)
43+ 
44+ return (
45+ boxes,
46+ sorted_scores,
47+ input_indices,
48+ max_output_size,
49+ iou_threshold,
50+ score_threshold,
51+ )
Mscripts/kernel/binary_config/ascendc_config.json+1-0
@@ -586,6 +586,7 @@
586 {"name":"Bincount", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},586 {"name":"Bincount", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
587 {"name":"Addr", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},587 {"name":"Addr", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
588 {"name":"NMSWithMask", "compute_units": ["ascend910b", "ascend950"], "auto_sync" : false, "impl_mode" : ""},588 {"name":"NMSWithMask", "compute_units": ["ascend910b", "ascend950"], "auto_sync" : false, "impl_mode" : ""},
589+ {"name": "SortedNMS", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": ""},
589 {"name":"Rasterizer", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false, "impl_mode" : ""},590 {"name":"Rasterizer", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false, "impl_mode" : ""},
590 {"name":"ExtractGlimpseV2", "compute_units": ["ascend950"], "auto_sync": false},591 {"name":"ExtractGlimpseV2", "compute_units": ["ascend950"], "auto_sync": false},
591 {"name":"RoiAlignV2", "compute_units": ["ascend910b"], "auto_sync":false, "impl_mode" : ""},592 {"name":"RoiAlignV2", "compute_units": ["ascend910b"], "auto_sync":false, "impl_mode" : ""},