bounding_box_decode-ops-cv-AtomGit

Mcommon/inc/op_graph/op_cv_proto_extend.h+0-38

	@@ -225,44 +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	-/**
229	-*@brief Generates bounding boxes based on "rois" and "deltas".
230	-* It is a customized FasterRcnn operator . \n
231	-
232	-*@par Inputs:
233	-* Two inputs, including:
234	-*@li rois: Region of interests (ROIs) generated by the region proposal
235	-* network (RPN). A 2D Tensor of type float32 or float16 with shape (N, 4).
236	-* "N" indicates the number of ROIs, and the value "4" refers to "x0", "x1",
237	-* "y0", and "y1".
238	-*@li deltas: Absolute variation between the ROIs generated by the RPN and
239	-* ground truth boxes. A 2D Tensor of type float32 or float16 with shape (N, 4).
240	-* "N" indicates the number of errors, and 4 indicates "dx", "dy", "dw", and "dh" . \n
241	-
242	-*@par Attributes:
243	-*@li means: An index of type float32. Defaults to [0.0, 0.0, 0.0, 0.0].
244	-* "deltas" = "deltas" x "stds" + "means".
245	-*@li stds: An index of type float32. Defaults to [1.0, 1.0, 1.0, 1.0].
246	-* "deltas" = "deltas" x "stds" + "means".
247	-*@li max_shape: Shape [h, w], specifying the size of the image transferred to
248	-* the network. Used to ensure that the bbox shape after conversion does not
249	-* exceed "max_shape".
250	-*@li wh_ratio_clip: Defaults to "16/1000". The values of "dw" and "dh" fall
251	-* within (-wh_ratio_clip, wh_ratio_clip) . \n
252	-
253	-*@par Outputs:
254	-*bboxes: Bboxes generated based on "rois" and "deltas". Have the same format
255	-* and type as "rois".
256	-*/
257	-REG_OP(BoundingBoxDecode)
258	- .INPUT(rois, TensorType({DT_FLOAT16, DT_FLOAT}))
259	- .INPUT(deltas, TensorType({DT_FLOAT16, DT_FLOAT}))
260	- .OUTPUT(bboxes, TensorType({DT_FLOAT16, DT_FLOAT}))
261	- .ATTR(means, ListFloat, {0.0, 0.0, 0.0, 0.0})
262	- .ATTR(stds, ListFloat, {1.0, 1.0, 1.0, 1.0})
263	- .REQUIRED_ATTR(max_shape, ListInt)
264	- .ATTR(wh_ratio_clip, Float, 0.016f)
265	- .OP_END_FACTORY_REG(BoundingBoxDecode)
266		228
267	/**	229	/**
268	*@brief Judges whether the bounding box is valid. It is a customized	230	*@brief Judges whether the bounding box is valid. It is a customized

Mdocs/zh/op_list.md+10-0

	@@ -576,6 +576,16 @@
576	<td>AI Core</td>	576	<td>AI Core</td>
577	<td>计算锚框与真实边界框之间的编码偏移量，生成目标检测回归目标。</td>	577	<td>计算锚框与真实边界框之间的编码偏移量，生成目标检测回归目标。</td>
578	</tr>	578	</tr>
		579	+ <tr>
		580	+ <td>objdetect</td>
		581	+ <td><a href="../../objdetect/bounding_box_decode/README.md">bounding_box_decode</a></td>
		582	+ <td>&check;</td>
		583	+ <td>&check;</td>
		584	+ <td>&check;</td>
		585	+ <td>&check;</td>
		586	+ <td>AI Core</td>
		587	+ <td>将模型输出的相对于先验框（或锚点）的偏移量与缩放参数，转换为原始图像中真实的绝对边界框坐标。</td>
		588	+ </tr>
579	<tr>	589	<tr>
580	<td>objdetect</td>	590	<td>objdetect</td>
581	<td><a href="../../objdetect/batch_multi_class_non_max_suppression/README.md">batch_multi_class_non_max_suppression</a></td>	591	<td><a href="../../objdetect/batch_multi_class_non_max_suppression/README.md">batch_multi_class_non_max_suppression</a></td>

Aobjdetect/bounding_box_decode/CMakeLists.txt+27-0

@@ -0,0 +1,27 @@
	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	+
	13	+if(NOT ENABLE_TEST AND NOT BENCHMARK)
	14	+ list(REMOVE_ITEM CURRENT_DIRS tests)
	15	+endif()
	16	+
	17	+foreach(SUB_DIR ${CURRENT_DIRS})
	18	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	19	+ add_subdirectory(${SUB_DIR})
	20	+ endif()
	21	+endforeach()
	22	+
	23	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	24	+set(SUPPORT_TILING_DIR "arch35")
	25	+add_all_modules_sources(OPTYPE bounding_box_decode ACLNNTYPE aclnn_exclude
	26	+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
	27	+ TILING_DIR ${SUPPORT_TILING_DIR})

Aobjdetect/bounding_box_decode/README.md+141-0

@@ -0,0 +1,141 @@
	1	+# BoundingBoxDecode
	C C chenjiao10 天前新增算子需同步刷新op_list.md
	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	+将目标检测回归分支预测的偏移量`deltas`结合锚框`anchor_box`解码为图像坐标空间的绝对框`boxes`，用于后续NMS、画框或COCO mAP评估。该算子是`bounding_box_encode`的逆运算，两者共享同一套`(means, stds)`标准化参数与`(x1, y1, x2, y2)`锚框坐标约定，配套使用才能保证训练与推理编码一致。
	17	+
	18	+设锚框`anchor_box=(x1, y1, x2, y2)`、偏移量`deltas=(dx', dy', dw', dh')`、`means=(m0,m1,m2,m3)`、`stds=(s0,s1,s2,s3)`、`max_shape=(H, W)`，解码步骤为：
	19	+
	20	+$$
	21	+pw = x2 - x1 + 1,\quad pcx = (x1 + x2) \times 0.5
	22	+$$
	23	+
	24	+$$
	25	+dx = dx' \times s0 + m0,\quad dw = dw' \times s2 + m2
	26	+$$
	27	+
	28	+$$
	29	+gw = pw \times \exp(dw),\quad gx = pcx + pw \times dx
	30	+$$
	31	+
	32	+$$
	33	+x1_{out} = \mathrm{clip}(gx - gw \times 0.5 + 0.5,\ 0,\ W)
	34	+$$
	35	+
	36	+y维度（y1、y2）同理，裁剪上界为H。fp16输入时中间`exp`及乘加计算在fp32域进行后回cast到fp16。
	37	+
	38	+## 参数说明
	39	+
	40	+<table style="table-layout: fixed; width: 1576px">
	41	+<colgroup>
	42	+<col style="width: 170px">
	43	+<col style="width: 170px">
	44	+<col style="width: 200px">
	45	+<col style="width: 200px">
	46	+<col style="width: 170px">
	47	+</colgroup>
	48	+<thead>
	49	+<tr>
	50	+<th>参数名</th>
	51	+<th>输入/输出/属性</th>
	52	+<th>描述</th>
	53	+<th>数据类型</th>
	54	+<th>数据格式</th>
	55	+</tr>
	56	+</thead>
	57	+<tbody>
	58	+<tr>
	59	+<td>anchor_box</td>
	60	+<td>输入</td>
	61	+<td>锚框坐标(x1,y1,x2,y2)，对应公式中anchor_box。</td>
	62	+<td>FLOAT16、FLOAT</td>
	63	+<td>ND</td>
	64	+</tr>
	65	+<tr>
	66	+<td>deltas</td>
	67	+<td>输入</td>
	68	+<td>回归偏移量(dx',dy',dw',dh')，对应公式中deltas。</td>
	69	+<td>FLOAT16、FLOAT</td>
	70	+<td>ND</td>
	71	+</tr>
	72	+<tr>
	73	+<td>boxes</td>
	74	+<td>输出</td>
	75	+<td>解码后的绝对框(x1,y1,x2,y2)，对应公式中boxes。</td>
	76	+<td>FLOAT16、FLOAT</td>
	77	+<td>ND</td>
	78	+</tr>
	79	+<tr>
	80	+<td>means</td>
	81	+<td>可选属性</td>
	82	+<td>deltas反标准化均值，长度4，默认[0.0, 0.0, 0.0, 0.0]。</td>
	83	+<td>FLOAT</td>
	84	+<td>-</td>
	85	+</tr>
	86	+<tr>
	87	+<td>stds</td>
	88	+<td>可选属性</td>
	89	+<td>deltas反标准化标准差，长度4且各元素非0，默认[1.0, 1.0, 1.0, 1.0]。</td>
	90	+<td>FLOAT</td>
	91	+<td>-</td>
	92	+</tr>
	93	+<tr>
	94	+<td>max_shape</td>
	95	+<td>属性</td>
	96	+<td>解码框裁剪上限(H, W)，长度2。</td>
	97	+<td>INT64</td>
	98	+<td>-</td>
	99	+</tr>
	100	+<tr>
	101	+<td>wh_ratio_clip</td>
	102	+<td>可选属性</td>
	103	+<td>宽高比裁剪阈值，默认0.016。</td>
	104	+<td>FLOAT</td>
	105	+<td>-</td>
	106	+</tr>
	107	+</tbody>
	108	+</table>
	109	+
	110	+## 约束说明
	111	+
	112	+- anchor_box与deltas的shape必须完全一致(N, 4)，无广播。
	113	+- anchor_box与deltas的dtype必须一致。
	114	+- means与stds长度必须为4，stds各元素不能为0。
	115	+- max_shape长度必须为2。
	116	+- wh_ratio_clip必须大于0，当前不参与核心解码公式，仅做入参校验。
	117	+- 输入仅支持ND数据格式。
	118	+
	119	+## 调用说明
	120	+
	121	+<table style="table-layout: fixed; width: 1000px">
	122	+<colgroup>
	123	+<col style="width: 180px">
	124	+<col style="width: 200px">
	125	+<col style="width: 620px">
	126	+</colgroup>
	127	+<thead>
	128	+<tr>
	129	+<th>调用方式</th>
	130	+<th>样例代码</th>
	131	+<th>说明</th>
	132	+</tr>
	133	+</thead>
	134	+<tbody>
	135	+<tr>
	136	+<td>GE图模式</td>
	137	+<td>-</td>
	138	+<td><a href="examples/arch35/test_geir_bounding_box_decode.cpp">test_geir_bounding_box_decode.cpp</a></td>
	139	+</tr>
	140	+</tbody>
	141	+</table>

Aobjdetect/bounding_box_decode/examples/arch35/test_geir_bounding_box_decode.cpp+269-0

@@ -0,0 +1,269 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <iostream>
	12	+#include <fstream>
	13	+#include <string.h>
	14	+#include <stdint.h>
	15	+#include <vector>
	16	+#include <string>
	17	+#include <map>
	18	+#include "assert.h"
	19	+
	20	+#include "graph.h"
	21	+#include "types.h"
	22	+#include "tensor.h"
	23	+#include "ge_error_codes.h"
	24	+#include "ge_api_types.h"
	25	+#include "ge_api.h"
	26	+#include "array_ops.h"
	27	+#include "ge_ir_build.h"
	28	+#include "../../op_graph/bounding_box_decode_proto.h"
	29	+
	30	+#define FAILED -1
	31	+#define SUCCESS 0
	32	+
	33	+using namespace ge;
	34	+using std::map;
	35	+using std::string;
	36	+using std::vector;
	37	+
	38	+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	39	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	40	+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
	41	+ TensorDesc placeholder##intputIndex##_desc = TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, \
	42	+ intputDtype); \
	43	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	44	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	45	+ Tensor tensor_placeholder##intputIndex; \
	46	+ ret = GenOnesData(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
	47	+ placeholder##intputIndex##_desc, intputDtype, 2); \
	48	+ if (ret != SUCCESS) { \
	49	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	50	+ return FAILED; \
	51	+ } \
	52	+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
	53	+ input.push_back(tensor_placeholder##intputIndex); \
	54	+ graph.AddOp(placeholder##intputIndex); \
	55	+ add1.set_input_##intputName(placeholder##intputIndex); \
	56	+ inputs.push_back(placeholder##intputIndex);
	57	+
	58	+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
	59	+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
	60	+ add1.update_output_desc_##outputName(outputName##outputIndex##_desc);
	61	+
	62	+string GetTime()
	63	+{
	64	+ time_t timep;
	65	+ time(&timep);
	66	+ char tmp[64];
	67	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
	68	+ return string(tmp);
	69	+}
	70	+
	71	+uint32_t GetDataTypeSize(DataType dt)
	72	+{
	73	+ uint32_t oneByte = 1;
	74	+ uint32_t twoByte = 2;
	75	+ uint32_t fourByte = 4;
	76	+ uint32_t eightByte = 8;
	77	+
	78	+ if (dt == ge::DT_FLOAT) {
	79	+ return fourByte;
	80	+ } else if (dt == ge::DT_FLOAT16) {
	81	+ return twoByte;
	82	+ } else if (dt == ge::DT_INT16) {
	83	+ return twoByte;
	84	+ } else if (dt == ge::DT_UINT16) {
	85	+ return twoByte;
	86	+ } else if (dt == ge::DT_INT32) {
	87	+ return fourByte;
	88	+ } else if (dt == ge::DT_UINT32) {
	89	+ return fourByte;
	90	+ } else if (dt == ge::DT_INT64) {
	91	+ return eightByte;
	92	+ } else if (dt == ge::DT_UINT64) {
	93	+ return eightByte;
	94	+ } else if (dt == ge::DT_INT8) {
	95	+ return oneByte;
	96	+ }
	97	+ return fourByte;
	98	+}
	99	+
	100	+int32_t GenOnesData(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type,
	101	+ int value)
	102	+{
	103	+ input_tensor_desc.SetRealDimCnt(shapes.size());
	104	+ size_t size = 1;
	105	+ for (uint32_t i = 0; i < shapes.size(); i++) {
	106	+ size *= shapes[i];
	107	+ }
	108	+ uint32_t data_len = size * GetDataTypeSize(data_type);
	109	+ int32_t* pData = new (std::nothrow) int32_t[data_len];
	110	+ if (pData == nullptr) {
	111	+ return FAILED;
	112	+ }
	113	+ for (size_t i = 0; i < size; ++i) {
	114	+ *(pData + i) = value;
	115	+ }
	116	+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
	117	+ delete[] pData;
	118	+ return SUCCESS;
	119	+}
	120	+
	121	+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
	122	+{
	123	+ FILE* fp = fopen(bin_file.c_str(), "wb");
	124	+ if (fp == nullptr) {
	125	+ return FAILED;
	126	+ }
	127	+ size_t written = fwrite(inputData, sizeof(uint8_t), data_size, fp);
	128	+ fclose(fp);
	129	+ if (written != data_size) {
	130	+ return FAILED;
	131	+ }
	132	+ return SUCCESS;
	133	+}
	134	+
	135	+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs,
	136	+ std::vector<Operator>& outputs, Graph& graph)
	137	+{
	138	+ Status ret = SUCCESS;
	139	+
	140	+ auto add1 = op::BoundingBoxDecode("bounding_box_decode_1");
	141	+
	142	+ std::vector<int64_t> boxShape = {2, 4};
	143	+ ADD_INPUT(1, rois, inDtype, boxShape);
	144	+
	145	+ ADD_INPUT(2, deltas, inDtype, boxShape);
	146	+
	147	+ std::vector<float> meansVec = {0.0f, 0.0f, 0.0f, 0.0f};
	148	+ add1.SetAttr("means", meansVec);
	149	+
	150	+ std::vector<float> stdsVec = {1.0f, 1.0f, 1.0f, 1.0f};
	151	+ add1.SetAttr("stds", stdsVec);
	152	+
	153	+ std::vector<int64_t> maxShapeVec = {768, 1280};
	154	+ add1.SetAttr("max_shape", maxShapeVec);
	155	+
	156	+ add1.SetAttr("wh_ratio_clip", 0.016f);
	157	+
	158	+ ADD_OUTPUT(1, bboxes, inDtype, boxShape);
	159	+
	160	+ outputs.push_back(add1);
	161	+
	162	+ return SUCCESS;
	163	+}
	164	+
	165	+int main(int argc, char* argv[])
	166	+{
	167	+ const char* graph_name = "tc_ge_irrun_test";
	168	+ Graph graph(graph_name);
	169	+ std::vector<ge::Tensor> input;
	170	+
	171	+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
	172	+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
	173	+ Status ret = ge::GEInitialize(global_options);
	174	+ if (ret != SUCCESS) {
	175	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
	176	+ return FAILED;
	177	+ }
	178	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
	179	+
	180	+ std::vector<Operator> inputs{};
	181	+ std::vector<Operator> outputs{};
	182	+
	183	+ if (argc > 1) {
	184	+ std::cout << argv[1] << std::endl;
	185	+ }
	186	+
	187	+ DataType inDtype = DT_FLOAT;
	188	+
	189	+ std::cout << "Input dtype: " << inDtype << std::endl;
	190	+
	191	+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
	192	+ if (ret != SUCCESS) {
	193	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	194	+ return FAILED;
	195	+ }
	196	+
	197	+ if (!inputs.empty() && !outputs.empty()) {
	198	+ graph.SetInputs(inputs).SetOutputs(outputs);
	199	+ }
	200	+
	201	+ std::map<AscendString, AscendString> build_options = {};
	202	+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
	203	+ ge::Session* session = new Session(build_options);
	204	+
	205	+ if (session == nullptr) {
	206	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	207	+ return FAILED;
	208	+ }
	209	+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
	210	+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
	211	+
	212	+ std::map<AscendString, AscendString> graph_options = {};
	213	+ uint32_t graph_id = 0;
	214	+ ret = session->AddGraph(graph_id, graph, graph_options);
	215	+
	216	+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
	217	+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
	218	+ std::string file_path = "./dump";
	219	+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
	220	+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
	221	+ std::vector<ge::Tensor> output;
	222	+ ret = session->RunGraph(graph_id, input, output);
	223	+ if (ret != SUCCESS) {
	224	+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
	225	+ delete session;
	226	+ GEFinalize();
	227	+ return FAILED;
	228	+ }
	229	+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
	230	+
	231	+ int input_num = input.size();
	232	+ for (int i = 0; i < input_num; i++) {
	233	+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
	234	+ string input_file = "./tc_ge_irrun_test_npu_input_" + std::to_string(i) + ".bin";
	235	+ uint8_t* input_data_i = input[i].GetData();
	236	+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
	237	+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
	238	+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
	239	+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
	240	+ }
	241	+
	242	+ int output_num = output.size();
	243	+ for (int i = 0; i < output_num; i++) {
	244	+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
	245	+ string output_file = "./tc_ge_irrun_test_npu_output_" + std::to_string(i) + ".bin";
	246	+ uint8_t* output_data_i = output[i].GetData();
	247	+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
	248	+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
	249	+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
	250	+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
	251	+ }
	252	+
	253	+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
	254	+ std::string error_str(error_msg.GetString());
	255	+ std::cout << "Error message: " << error_str << std::endl;
	256	+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
	257	+ std::string warning_str(warning_msg.GetString());
	258	+ std::cout << "Warning message: " << warning_str << std::endl;
	259	+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
	260	+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
	261	+ delete session;
	262	+ ret = ge::GEFinalize();
	263	+ if (ret != SUCCESS) {
	264	+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
	265	+ return FAILED;
	266	+ }
	267	+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
	268	+ return SUCCESS;
	269	+}

Aobjdetect/bounding_box_decode/op_graph/bounding_box_decode_graph_infer.cpp+46-0

@@ -0,0 +1,46 @@
	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	+
	13	+using namespace ge;
	14	+
	15	+namespace ops {
	16	+
	17	+/**
	18	+ * InferDataTypeForBoundingBoxDecode: GE data-type inference callback.
	19	+ *
	20	+ * boxes.dtype = anchor_box.dtype (same_as_first_input). Validates that both
	21	+ * inputs have the same dtype (DESIGN §3.2, no type promotion).
	22	+ */
	23	+static ge::graphStatus InferDataTypeForBoundingBoxDecode(gert::InferDataTypeContext* context)
	24	+{
	25	+ const ge::DataType anchorDtype = context->GetInputDataType(0);
	26	+ const ge::DataType deltasDtype = context->GetInputDataType(1);
	27	+
	28	+ if (anchorDtype != deltasDtype) {
	29	+ return ge::GRAPH_FAILED;
	30	+ }
	31	+
	32	+ context->SetOutputDataType(0, anchorDtype);
	33	+ return ge::GRAPH_SUCCESS;
	34	+}
	35	+
	36	+/**
	37	+ * IMPL_OP(BoundingBoxDecode).InferDataType(...): registers the dtype inference
	38	+ * function for the operator named "BoundingBoxDecode" at static init time.
	39	+ * When GE encounters a BoundingBoxDecode node during graph compilation, it
	40	+ * calls InferDataTypeForBoundingBoxDecode to determine the output type.
	41	+ * This static registration is required by the GE runtime to locate the op
	42	+ * implementation via the op_impl_registry.
	43	+ */
	44	+IMPL_OP(BoundingBoxDecode).InferDataType(InferDataTypeForBoundingBoxDecode);
	45	+
	46	+} // namespace ops

Aobjdetect/bounding_box_decode/op_graph/bounding_box_decode_proto.h+64-0

@@ -0,0 +1,64 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file bounding_box_decode_proto.h
	13	+ * \brief decode the offset between anchor boxes and ground truth boxes.
	14	+ */
	15	+#ifndef OPS_OP_PROTO_INC_BOUNDING_BOX_DECODE_H_
	16	+#define OPS_OP_PROTO_INC_BOUNDING_BOX_DECODE_H_
	17	+
	18	+#include "graph/operator_reg.h"
	19	+#include "graph/types.h"
	20	+
	21	+namespace ge {
	22	+/**
	23	+*@brief Generates bounding boxes based on "rois" and "deltas".
	24	+* It is a customized FasterRcnn operator . \n
	25	+
	26	+*@par Inputs:
	27	+* Two inputs, including:
	28	+*@li rois: Region of interests (ROIs) generated by the region proposal
	29	+* network (RPN). A 2D Tensor of type float32 or float16 with shape (N, 4).
	30	+* "N" indicates the number of ROIs, and the value "4" refers to "x0", "x1",
	31	+* "y0", and "y1".
	32	+*@li deltas: Absolute variation between the ROIs generated by the RPN and
	33	+* ground truth boxes. A 2D Tensor of type float32 or float16 with shape (N, 4).
	34	+* "N" indicates the number of errors, and 4 indicates "dx", "dy", "dw", and "dh" . \n
	35	+
	36	+*@par Attributes:
	37	+*@li means: An index of type float32. Defaults to [0.0, 0.0, 0.0, 0.0].
	38	+* "deltas" = "deltas" x "stds" + "means".
	39	+*@li stds: An index of type float32. Defaults to [1.0, 1.0, 1.0, 1.0].
	40	+* "deltas" = "deltas" x "stds" + "means".
	41	+*@li max_shape: Shape [h, w], specifying the size of the image transferred to
	42	+* the network. Used to ensure that the bbox shape after conversion does not
	43	+* exceed "max_shape".
	44	+*@li wh_ratio_clip: Defaults to "16/1000". The values of "dw" and "dh" fall
	45	+* within (-wh_ratio_clip, wh_ratio_clip) . \n
	46	+
	47	+*@par Outputs:
	48	+*bboxes: Bboxes generated based on "rois" and "deltas". Have the same format
	49	+* and type as "rois".
	50	+*/
	51	+
	52	+REG_OP(BoundingBoxDecode)
	W W wangyongguang10 天前添加去重宏 W wuao 10 天前评论： 10 天前评论：
	53	+ .INPUT(rois, TensorType({DT_FLOAT16, DT_FLOAT}))
	54	+ .INPUT(deltas, TensorType({DT_FLOAT16, DT_FLOAT}))
	55	+ .OUTPUT(bboxes, TensorType({DT_FLOAT16, DT_FLOAT}))
	56	+ .ATTR(means, ListFloat, {0.0, 0.0, 0.0, 0.0})
	57	+ .ATTR(stds, ListFloat, {1.0, 1.0, 1.0, 1.0})
	58	+ .REQUIRED_ATTR(max_shape, ListInt)
	59	+ .ATTR(wh_ratio_clip, Float, 0.016f)
	60	+ .OP_END_FACTORY_REG(BoundingBoxDecode)
	61	+
	62	+} // namespace ge
	63	+
	64	+#endif

Aobjdetect/bounding_box_decode/op_graph/fusion_pass/.gitkeep+0-0

The file is empty

Aobjdetect/bounding_box_decode/op_host/arch35/bounding_box_decode_tiling_arch35.cpp+342-0

@@ -0,0 +1,342 @@
	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 <cstdint>
	13	+#include <cstring>
	14	+
	15	+// BoundingBoxDecodeTilingData struct (§7, 13 fields) — shared host/kernel.
	16	+#include "../../op_kernel/arch35/bounding_box_decode_tiling_struct.h"
	17	+// gert::TilingContext — GetInputShape / GetInputDesc / GetAttrs / GetTilingData /
	18	+// SetBlockDim / SetTilingKey / GetWorkspaceSizes / GetPlatformInfo.
	19	+#include "exe_graph/runtime/tiling_context.h"
	20	+// gert::Tensor — GetDataType / GetStorageFormat (input tensor metadata).
	21	+#include "exe_graph/runtime/tensor.h"
	22	+// ge::DataType / ge::Format — DT_FLOAT16, DT_FLOAT, FORMAT_ND, etc.
	23	+#include "graph/types.h"
	24	+// platform_ascendc::PlatformAscendC — GetCoreNumAiv / GetCoreMemSize(UB).
	25	+#include "tiling/platform/platform_ascendc.h"
	26	+// op_def_registry.h — IMPL_OP_OPTILING macro for tiling registration.
	27	+#include "register/op_def_registry.h"
	28	+// Own header — declares optiling::TilingFunc.
	29	+#include "bounding_box_decode_tiling_arch35.h"
	30	+
	31	+namespace optiling {
	32	+
	33	+// =========================================================================
	34	+// Constants (DESIGN §9.3 / §9.4 / §9.5 / §9.7 / §6)
	35	+// =========================================================================
	36	+
	37	+// §6 TilingKey values — TPL key = T_value (datatype only, no BOOL param)
	38	+// kIsEmpty TPL parameter removed (Task 41 fix): the BOOL TPL parameter
	39	+// caused the framework's NnopbaseExecutorDoTiling to fail allocating the
	40	+// tiling data buffer in aclnn e2e mode (chicken-and-egg: tilingKey encodes
	41	+// the BOOL value but is set by TilingFunc, which needs the buffer first).
	42	+// Now the kernel handles empty tensors via runtime check (td.dim0 == 0).
	43	+// (T=FP32=0) → 0 (fp32, handles both normal and empty)
	44	+// (T=FP16=1) → 1 (fp16, handles both normal and empty)
	45	+constexpr uint64_t BOUNDING_BOX_DECODE_FP32 = 0;
	46	+constexpr uint64_t BOUNDING_BOX_DECODE_FP16 = 1;
	47	+
	48	+// §9.3 validation constants
	49	+constexpr int64_t MAX_RANK = 2; // spec.yaml rank_range:[2,2], fixed rank=2
	50	+constexpr int64_t ELEMS_PER_BOX = 4; // C=4 fixed (kElemsPerBox)
	51	+
	52	+// §9.5 multi-core split constants
	53	+constexpr int64_t MIN_TILING_BITS = 32768; // 4 KB, unit: bits
	54	+constexpr int64_t ELEM_ALIGN_FACTOR = 512; // multi-core box alignment factor
	55	+
	56	+// §9.4 UB split constants
	57	+constexpr int64_t ALIGN_256 = 256; // UB alignment, unit: bytes
	58	+constexpr int64_t RESERVED_BYTES = 8 * 1024; // UB reserved for sync/TBuf
	59	+constexpr int64_t NUM_IO_BUFS = 3; // B_anchor, B_deltas, B_boxes (§9.1)
	60	+
	61	+// =========================================================================
	62	+// Inline math helpers (mirror DESIGN §9.4 / §9.5 formulas)
	63	+// =========================================================================
	64	+static inline int64_t CeilDiv(int64_t a, int64_t b) { return (a + b - 1) / b; }
	65	+static inline int64_t AlignUp(int64_t v, int64_t f) { return CeilDiv(v, f) * f; }
	66	+static inline int64_t AlignDown(int64_t v, int64_t f) { return (v / f) * f; }
	67	+
	68	+// =========================================================================
	69	+// §9.5 ComputeMultiCoreSplit — by box count, 512-box alignment
	70	+//
	71	+// Fills: dim0, coreNum, blockFormer, blockNum.
	72	+// coreNum = min(CeilDiv(dim0 × minDtypeBits, MIN_TILING_BITS), availableCoreNum); ≥1
	73	+// blockFormer = AlignUp(CeilDiv(dim0, coreNum), ELEM_ALIGN_FACTOR)
	74	+// blockNum = CeilDiv(dim0, blockFormer)
	75	+// =========================================================================
	76	+static void ComputeMultiCoreSplit(int64_t dim0, int64_t sizeofT, int64_t availableCoreNum,
	77	+ BoundingBoxDecodeTilingData* td)
	78	+{
	79	+ int64_t minDtypeBits = ELEMS_PER_BOX * sizeofT * 8; // bits per box
	80	+ int64_t coreNum = CeilDiv(dim0 * minDtypeBits, MIN_TILING_BITS);
	81	+ coreNum = std::min(coreNum, availableCoreNum);
	82	+ coreNum = std::max(coreNum, static_cast<int64_t>(1)); // at least 1 core
	83	+
	84	+ int64_t blockFormer = AlignUp(CeilDiv(dim0, coreNum), ELEM_ALIGN_FACTOR);
	85	+ int64_t blockNum = CeilDiv(dim0, blockFormer);
	86	+ // 512-box alignment can enlarge blockFormer so that blockNum < coreNum;
	87	+ // cap coreNum to blockNum to avoid launching idle cores (SetBlockDim).
	88	+ coreNum = std::min(coreNum, blockNum);
	89	+
	90	+ td->dim0 = dim0;
	91	+ td->coreNum = static_cast<int32_t>(coreNum);
	92	+ td->blockFormer = blockFormer;
	93	+ td->blockNum = blockNum;
	94	+}
	95	+
	96	+// =========================================================================
	97	+// §9.4 ComputeUbSplit — by box count, 256B alignment
	98	+//
	99	+// Fills: ubFormer, ubLoopOfFormerBlock, ubTailOfFormerBlock,
	100	+// ubLoopOfTailBlock, ubTailOfTailBlock.
	101	+// perBoxBytes = NUM_IO_BUFS×4×sizeof(T) + K×4×sizeof(float) (§9.1)
	102	+// alignFactor = ALIGN_256 / (4 × sizeof(T)) [box count]
	103	+// ubFormer = max(AlignDown((ubSize-RESERVED)/perBoxBytes, alignFactor), alignFactor)
	104	+// blockTail = dim0 - (blockNum-1)×blockFormer
	105	+// =========================================================================
	106	+static void ComputeUbSplit(int64_t ubSize, int64_t sizeofT, int64_t K, int64_t blockFormer, int64_t dim0,
	107	+ int64_t blockNum, BoundingBoxDecodeTilingData* td)
	108	+{
	109	+ int64_t perBoxBytes = NUM_IO_BUFS * ELEMS_PER_BOX * sizeofT + K * ELEMS_PER_BOX * sizeof(float);
	110	+ int64_t alignFactor = ALIGN_256 / (ELEMS_PER_BOX * sizeofT);
	111	+ int64_t maxBoxNum = (ubSize - RESERVED_BYTES) / perBoxBytes;
	112	+ int64_t ubFormer = AlignDown(maxBoxNum, alignFactor);
	113	+ ubFormer = std::max(ubFormer, alignFactor); // floor at 1 alignment block
	114	+
	115	+ int64_t ubLoopOfFormerBlock = CeilDiv(blockFormer, ubFormer);
	116	+ int64_t ubTailOfFormerBlock = blockFormer - (ubLoopOfFormerBlock - 1) * ubFormer;
	117	+
	118	+ int64_t blockTail = dim0 - (blockNum - 1) * blockFormer;
	119	+ int64_t ubLoopOfTailBlock = CeilDiv(blockTail, ubFormer);
	120	+ int64_t ubTailOfTailBlock = blockTail - (ubLoopOfTailBlock - 1) * ubFormer;
	121	+
	122	+ td->ubFormer = ubFormer;
	123	+ td->ubLoopOfFormerBlock = ubLoopOfFormerBlock;
	124	+ td->ubTailOfFormerBlock = ubTailOfFormerBlock;
	125	+ td->ubLoopOfTailBlock = ubLoopOfTailBlock;
	126	+ td->ubTailOfTailBlock = ubTailOfTailBlock;
	127	+}
	128	+
	129	+// =========================================================================
	130	+// §9.3 BoundingBoxDecodePreCheck — 异常值校验
	131	+//
	132	+// Fixed order: null → dtype → format → rank(rank==2) → attr → shape.
	133	+// On success returns true and fills N (box count) + isEmpty (N==0).
	134	+// On any failure returns false (caller returns GRAPH_FAILED, no kernel launch).
	135	+//
	136	+// Attr layout (OpDef order, indices match AppendAttr order):
	137	+// index 0: means (ListFloat→ GetListFloat(0))
	138	+// index 1: stds (ListFloat→ GetListFloat(1))
	139	+// index 2: max_shape (ListInt → GetListInt(2))
	140	+// index 3: wh_ratio_clip(Float → GetFloat(3))
	141	+// =========================================================================
	142	+static bool BoundingBoxDecodePreCheck(gert::TilingContext* ctx, int64_t& N, bool& isEmpty)
	143	+{
	144	+ // 1. null check — GetInputShape bounds-checks (returns nullptr for missing input).
	145	+ const gert::StorageShape* aShape = ctx->GetInputShape(0);
	146	+ const gert::StorageShape* dShape = ctx->GetInputShape(1);
	147	+ if (aShape == nullptr \|\| dShape == nullptr) {
	148	+ return false; // → null_input
	149	+ }
	150	+
	151	+ // 2. dtype check: anchor_box/deltas ∈ {fp16,fp32} and equal.
	152	+ const gert::CompileTimeTensorDesc* aDesc = ctx->GetInputDesc(0);
	153	+ const gert::CompileTimeTensorDesc* dDesc = ctx->GetInputDesc(1);
	154	+ if (aDesc == nullptr \|\| dDesc == nullptr) {
	155	+ return false;
	156	+ }
	157	+ ge::DataType aDt = aDesc->GetDataType();
	158	+ ge::DataType dDt = dDesc->GetDataType();
	159	+ if (aDt != ge::DT_FLOAT16 && aDt != ge::DT_FLOAT) {
	160	+ return false; // → dtype_not_supported
	161	+ }
	162	+ if (aDt != dDt) {
	163	+ return false; // → dtype_not_supported (mix)
	164	+ }
	165	+
	166	+ // 3. format check: inputs must be ND.
	167	+ ge::Format aFmt = aDesc->GetStorageFormat();
	168	+ ge::Format dFmt = dDesc->GetStorageFormat();
	169	+ if (aFmt != ge::FORMAT_ND \|\| dFmt != ge::FORMAT_ND) {
	170	+ return false; // → shape_mismatch (non-ND format rejected)
	171	+ }
	172	+
	173	+ // 4. rank check: both inputs rank == 2 (MAX_RANK).
	174	+ const gert::Shape& aShp = aShape->GetStorageShape();
	175	+ const gert::Shape& dShp = dShape->GetStorageShape();
	176	+ if (aShp.GetDimNum() != static_cast<size_t>(MAX_RANK) \|\| dShp.GetDimNum() != static_cast<size_t>(MAX_RANK)) {
	177	+ return false; // → shape_mismatch (rank != 2)
	178	+ }
	179	+
	180	+ // 5. attr value-range check.
	181	+ const gert::RuntimeAttrs* attrs = ctx->GetAttrs();
	182	+ if (attrs == nullptr) {
	183	+ return false; // → attribute_value_out_of_range (attrs missing)
	184	+ }
	185	+ const auto* meansVec = attrs->GetListFloat(0); // means ListFloat index 0
	186	+ const auto* stdsVec = attrs->GetListFloat(1); // stds ListFloat index 1
	187	+ const auto* maxShapeVec = attrs->GetListInt(2); // max_shape ListInt index 2
	188	+ const float* clipPtr = attrs->GetFloat(3); // wh_ratio_clip Float index 3
	189	+ if (maxShapeVec == nullptr \|\| meansVec == nullptr \|\| stdsVec == nullptr \|\| clipPtr == nullptr) {
	190	+ return false; // → attribute_value_out_of_range (required attr missing)
	191	+ }
	192	+ // means/stds length == 4
	193	+ if (meansVec->GetSize() != 4 \|\| stdsVec->GetSize() != 4) {
	194	+ return false; // → attribute_value_out_of_range (length != 4)
	195	+ }
	196	+ const float* means = meansVec->GetData();
	197	+ const float* stds = stdsVec->GetData();
	198	+ // stds elements must be non-zero
	199	+ if (stds[0] == 0.0f \|\| stds[1] == 0.0f \|\| stds[2] == 0.0f \|\| stds[3] == 0.0f) {
	200	+ return false; // → attribute_value_out_of_range (stds has zero)
	201	+ }
	202	+ // wh_ratio_clip > 0
	203	+ if (*clipPtr <= 0.0f) {
	204	+ return false; // → attribute_value_out_of_range (wh_ratio_clip <= 0)
	205	+ }
	206	+ // max_shape length == 2 (H, W)
	207	+ if (maxShapeVec->GetSize() != 2) {
	208	+ return false; // → attribute_value_out_of_range (max_shape length != 2)
	209	+ }
	210	+
	211	+ // 6. shape check: last dim == 4 and N (dim 0) matches across inputs.
	212	+ if (aShp.GetDim(1) != ELEMS_PER_BOX \|\| dShp.GetDim(1) != ELEMS_PER_BOX) {
	213	+ return false; // → shape_mismatch (last dim != 4)
	214	+ }
	215	+ if (aShp.GetDim(0) != dShp.GetDim(0)) {
	216	+ return false; // → shape_mismatch (N mismatch)
	217	+ }
	218	+
	219	+ N = aShp.GetDim(0);
	220	+ isEmpty = (N == 0);
	221	+ return true;
	222	+}
	223	+
	224	+// =========================================================================
	225	+// §9.7 FillTilingData — empty/normal fork + attr passthrough + SetTilingKey/SetBlockDim
	226	+//
	227	+// empty branch: coreNum=1, compute fields=0, SetTilingKey(dtype), SetBlockDim(1).
	228	+// normal branch: ComputeMultiCoreSplit → ComputeUbSplit,
	229	+// SetTilingKey(dtype), SetBlockDim(coreNum).
	230	+// tilingKey is now dtype-only (no BOOL component) — see struct.h comment.
	231	+// Attr scalars (means/stds/maxShapeH/maxShapeW) are passed through in both branches.
	232	+// =========================================================================
	233	+static void FillTilingData(gert::TilingContext* ctx, BoundingBoxDecodeTilingData* td, int64_t N, bool isEmpty,
	234	+ int64_t sizeofT, int64_t ubSize, int64_t availableCoreNum)
	235	+{
	236	+ td->dim0 = N;
	237	+
	238	+ // Attr passthrough (§7 field table; aclnn-supplied, kernel Compute consumes).
	239	+ const gert::RuntimeAttrs* attrs = ctx->GetAttrs();
	240	+ const auto* meansVec = attrs->GetListFloat(0);
	241	+ const auto* stdsVec = attrs->GetListFloat(1);
	242	+ const auto* maxShapeVec = attrs->GetListInt(2);
	243	+ const float* means = meansVec->GetData();
	244	+ const float* stds = stdsVec->GetData();
	245	+ const int64_t* ms = maxShapeVec->GetData();
	246	+ for (int i = 0; i < 4; i++) {
	247	+ td->means[i] = means[i];
	248	+ td->stds[i] = stds[i];
	249	+ }
	250	+ td->maxShapeH = ms[0];
	251	+ td->maxShapeW = ms[1];
	252	+ // wh_ratio_clip is not referenced by the core formula (§1.3) — not in TilingData.
	253	+
	254	+ if (isEmpty) {
	255	+ // empty branch: coreNum=1, compute fields=0 (kernel short-circuits via td.dim0==0).
	256	+ td->coreNum = 1;
	257	+ td->blockFormer = 0;
	258	+ td->blockNum = 1; // 1 (not 0) to avoid downstream div-by-zero
	259	+ td->ubFormer = 0;
	260	+ td->ubLoopOfFormerBlock = 0;
	261	+ td->ubTailOfFormerBlock = 0;
	262	+ td->ubLoopOfTailBlock = 0;
	263	+ td->ubTailOfTailBlock = 0;
	264	+ ctx->SetBlockDim(1);
	265	+ ctx->SetTilingKey((sizeofT == 2) ? BOUNDING_BOX_DECODE_FP16 : BOUNDING_BOX_DECODE_FP32);
	266	+ } else {
	267	+ // normal branch: K by dtype — fp32→0, fp16→2 (§9.1 NUM_CALC_BUFS).
	268	+ int64_t K = (sizeofT == static_cast<int64_t>(sizeof(float))) ? 0 : 2;
	269	+ ComputeMultiCoreSplit(N, sizeofT, availableCoreNum, td);
	270	+ ComputeUbSplit(ubSize, sizeofT, K, td->blockFormer, td->dim0, td->blockNum, td);
	271	+ ctx->SetBlockDim(static_cast<uint32_t>(td->coreNum));
	272	+ ctx->SetTilingKey((sizeofT == 2) ? BOUNDING_BOX_DECODE_FP16 : BOUNDING_BOX_DECODE_FP32);
	273	+ }
	274	+}
	275	+
	276	+// =========================================================================
	277	+// §9.9 TilingFunc — entry point
	278	+//
	279	+// Workflow:
	280	+// 1. Get TilingData buffer (GetTilingData<T> sets data size = sizeof(T)).
	281	+// 2. Query platform info (ubSize / availableCoreNum) via PlatformAscendC.
	282	+// 3. BoundingBoxDecodePreCheck (§9.3) — fail → GRAPH_FAILED, no kernel.
	283	+// 4. FillTilingData (§9.7) — empty/normal fork + attr passthrough.
	284	+// 5. Set workspace = 0 (no cross-core partial merge, §9.6).
	285	+// =========================================================================
	286	+ge::graphStatus TilingFunc(gert::TilingContext* context)
	287	+{
	288	+ // 1. TilingData buffer
	289	+ BoundingBoxDecodeTilingData* td = context->GetTilingData<BoundingBoxDecodeTilingData>();
	290	+ if (td == nullptr) {
	291	+ return ge::GRAPH_FAILED;
	292	+ }
	293	+ std::memset(td, 0, sizeof(BoundingBoxDecodeTilingData));
	294	+
	295	+ // 2. Platform info — ubSize / availableCoreNum (never hard-coded)
	296	+ int64_t ubSize = 0;
	297	+ int64_t availableCoreNum = 0;
	298	+ fe::PlatFormInfos* platformInfo = context->GetPlatformInfo();
	299	+ if (platformInfo != nullptr) {
	300	+ platform_ascendc::PlatformAscendC ascendcPlatform(platformInfo);
	301	+ uint64_t ub = 0;
	302	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ub);
	303	+ ubSize = static_cast<int64_t>(ub);
	304	+ availableCoreNum = static_cast<int64_t>(ascendcPlatform.GetCoreNumAiv());
	305	+ }
	306	+ if (ubSize <= 0) {
	307	+ ubSize = 196608; // 192 KB fallback (ascend950 default)
	308	+ }
	309	+ if (availableCoreNum <= 0) {
	310	+ availableCoreNum = 1;
	311	+ }
	312	+
	313	+ // 3. PreCheck (§9.3): null → dtype → format → rank → attr → shape
	314	+ int64_t N = 0;
	315	+ bool isEmpty = false;
	316	+ if (!BoundingBoxDecodePreCheck(context, N, isEmpty)) {
	317	+ return ge::GRAPH_FAILED; // validation failure → no kernel launch
	318	+ }
	319	+
	320	+ // 4. FillTilingData (§9.7): empty/normal fork + attr passthrough
	321	+ ge::DataType aDt = context->GetInputDesc(0)->GetDataType();
	322	+ int64_t sizeofT = (aDt == ge::DT_FLOAT) ? 4 : 2;
	323	+ FillTilingData(context, td, N, isEmpty, sizeofT, ubSize, availableCoreNum);
	324	+
	325	+ // 5. Workspace = 0 (Elementwise, no cross-core partial merge, §9.6)
	326	+ size_t* ws = context->GetWorkspaceSizes(1);
	327	+ if (ws != nullptr) {
	328	+ ws[0] = 0;
	329	+ }
	330	+
	331	+ return ge::GRAPH_SUCCESS;
	332	+}
	333	+
	334	+// =============================================================================
	335	+// §9.8 Host-side registration — IMPL_OP_OPTILING
	336	+// .Tiling(TilingFunc): registers the runtime tiling callback.
	337	+// Platform info is queried live via GetPlatformInfo() + PlatformAscendC
	338	+// (broadcast_tiling pattern), so no TilingParse/CompileInfo cache is required.
	339	+// =============================================================================
	340	+IMPL_OP_OPTILING(BoundingBoxDecode).Tiling(TilingFunc);
	341	+
	342	+} // namespace optiling

Aobjdetect/bounding_box_decode/op_host/arch35/bounding_box_decode_tiling_arch35.h+37-0

@@ -0,0 +1,37 @@
	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 BOUNDING_BOX_DECODE_TILING_ARCH35_H
	12	+#define BOUNDING_BOX_DECODE_TILING_ARCH35_H
	13	+
	14	+namespace optiling {
	15	+
	16	+/**
	17	+ * TilingFunc: the tiling callback invoked by the CANN framework before kernel launch.
	18	+ *
	19	+ * Parameters:
	20	+ * context — [in/out] tiling context providing input shapes, dtypes, and
	21	+ * accepting the computed BoundingBoxDecodeTilingData.
	22	+ *
	23	+ * Returns:
	24	+ * ge::GRAPH_SUCCESS on success.
	25	+ *
	26	+ * Side effects:
	27	+ * - Reads input tensor metadata from context.
	28	+ * - Computes totalLength, blockLength, tileLength, numBlocks.
	29	+ * - Writes these values into BoundingBoxDecodeTilingData via context->GetTilingData.
	30	+ * - Sets the block dimension via context->SetBlockDim.
	31	+ * - Optionally sets workspace sizes via context->GetWorkspaceSizes.
	32	+ */
	33	+ge::graphStatus TilingFunc(gert::TilingContext* context);
	34	+
	35	+} // namespace optiling
	36	+
	37	+#endif

Aobjdetect/bounding_box_decode/op_host/bounding_box_decode_def.cpp+59-0

@@ -0,0 +1,59 @@
	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	+/* Generated By CANNBot */
	12	+
	13	+#include "register/op_def_registry.h"
	14	+
	15	+namespace ops {
	16	+class BoundingBoxDecode : public OpDef {
	17	+public:
	18	+ explicit BoundingBoxDecode(const char* name) : OpDef(name)
	19	+ {
	20	+ this->Input("rois")
	21	+ .ParamType(REQUIRED)
	22	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
	23	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	24	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	25	+ .AutoContiguous();
	26	+
	27	+ this->Input("deltas")
	28	+ .ParamType(REQUIRED)
	29	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
	30	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	31	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	32	+ .AutoContiguous();
	33	+
	34	+ this->Output("bboxes")
	35	+ .ParamType(REQUIRED)
	36	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
	37	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	38	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	39	+ .AutoContiguous();
	40	+
	41	+ this->Attr("means").AttrType(OPTIONAL).ListFloat({0.0f, 0.0f, 0.0f, 0.0f});
	42	+ this->Attr("stds").AttrType(OPTIONAL).ListFloat({1.0f, 1.0f, 1.0f, 1.0f});
	43	+ this->Attr("max_shape").AttrType(REQUIRED).ListInt();
	44	+ this->Attr("wh_ratio_clip").AttrType(OPTIONAL).Float(0.016f);
	45	+
	46	+ OpAICoreConfig aicoreConfig;
	47	+ aicoreConfig.DynamicCompileStaticFlag(true)
	48	+ .DynamicFormatFlag(false)
	49	+ .DynamicRankSupportFlag(true)
	50	+ .DynamicShapeSupportFlag(true)
	51	+ .NeedCheckSupportFlag(false)
	52	+ .PrecisionReduceFlag(false)
	53	+ .ExtendCfgInfo("opFile.value", "bounding_box_decode_apt");
	54	+
	55	+ this->AICore().AddConfig("ascend950", aicoreConfig);
	56	+ }
	57	+};
	58	+OP_ADD(BoundingBoxDecode);
	59	+} // namespace ops

Aobjdetect/bounding_box_decode/op_host/bounding_box_decode_infershape.cpp+87-0

@@ -0,0 +1,87 @@
	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	+
	13	+#include <cstdint>
	14	+#include <vector>
	15	+
	16	+#include "graph/types.h"
	17	+
	18	+using namespace ge;
	19	+
	20	+namespace ops {
	21	+
	22	+namespace {
	23	+
	24	+bool IsUnknownRank(const gert::Shape* shape)
	25	+{
	26	+ return shape != nullptr && shape->GetDimNum() == 1 && shape->GetDim(0) == ge::UNKNOWN_DIM_NUM;
	27	+}
	28	+
	29	+} // namespace
	30	+
	31	+/**
	32	+ * InferShapeForBoundingBoxDecode: GE shape inference callback.
	33	+ *
	34	+ * boxes.shape = anchor_box.shape (identity). Validates that both inputs are
	35	+ * non-null, rank == 2, last dim == 4, and the two input shapes are equal
	36	+ * (DESIGN §3.1). Empty tensor (N==0) is supported — the identity copy
	37	+ * propagates (0, 4) to the output.
	38	+ */
	39	+ge::graphStatus InferShapeForBoundingBoxDecode(gert::InferShapeContext* context)
	40	+{
	41	+ const gert::Shape* anchorShape = context->GetInputShape(0);
	42	+ const gert::Shape* deltasShape = context->GetInputShape(1);
	43	+ gert::Shape* boxesShape = context->GetOutputShape(0);
	44	+ if (anchorShape == nullptr \|\| deltasShape == nullptr \|\| boxesShape == nullptr) {
	45	+ return GRAPH_FAILED;
	46	+ }
	47	+
	48	+ if (IsUnknownRank(anchorShape) \|\| IsUnknownRank(deltasShape)) {
	49	+ boxesShape->SetDimNum(1);
	50	+ boxesShape->SetDim(0, ge::UNKNOWN_DIM_NUM);
	51	+ return GRAPH_SUCCESS;
	52	+ }
	53	+
	54	+ const size_t rank = anchorShape->GetDimNum();
	55	+ if (rank != 2 \|\| deltasShape->GetDimNum() != 2) {
	56	+ return GRAPH_FAILED;
	57	+ }
	58	+ if (anchorShape->GetDim(1) != 4 \|\| deltasShape->GetDim(1) != 4) {
	59	+ return GRAPH_FAILED;
	60	+ }
	61	+
	62	+ // Validate anchor_box.shape == deltas.shape (no broadcast, DESIGN §3.1)
	63	+ const int64_t anchorN = anchorShape->GetDim(0);
	64	+ const int64_t deltasN = deltasShape->GetDim(0);
	65	+ if (anchorN != deltasN && anchorN != ge::UNKNOWN_DIM && deltasN != ge::UNKNOWN_DIM) {
	66	+ return GRAPH_FAILED;
	67	+ }
	68	+
	69	+ // boxes.shape = anchor_box.shape (identity copy)
	70	+ boxesShape->SetDimNum(rank);
	71	+ for (size_t i = 0; i < rank; ++i) {
	72	+ boxesShape->SetDim(i, anchorShape->GetDim(i));
	73	+ }
	74	+ return GRAPH_SUCCESS;
	75	+}
	76	+
	77	+/**
	78	+ * IMPL_OP_INFERSHAPE(BoundingBoxDecode).InferShape(...):
	79	+ * Registers the shape inference function at static init time. When the
	80	+ * framework needs to determine the output shape of a BoundingBoxDecode node,
	81	+ * it calls InferShapeForBoundingBoxDecode. This static registration is
	82	+ * required by the GE runtime (NnopbaseExecutorDoTiling) to locate the op
	83	+ * implementation via the op_impl_registry.
	84	+ */
	85	+IMPL_OP_INFERSHAPE(BoundingBoxDecode).InferShape(InferShapeForBoundingBoxDecode);
	86	+
	87	+} // namespace ops

Aobjdetect/bounding_box_decode/op_kernel/arch35/bounding_box_decode_kernel.h+462-0

@@ -0,0 +1,462 @@
	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 BOUNDING_BOX_DECODE_KERNEL_H
	12	+#define BOUNDING_BOX_DECODE_KERNEL_H
	13	+
	14	+#include "kernel_operator.h"
	15	+#include "bounding_box_decode_tiling_data.h"
	16	+
	17	+// Cast traits for fp16<->fp32 conversion inside VF functions
	18	+constexpr AscendC::MicroAPI::CastTrait kCastB162B32 = {
	19	+ AscendC::MicroAPI::RegLayout::ZERO,
	20	+ AscendC::MicroAPI::SatMode::UNKNOWN,
	21	+ AscendC::MicroAPI::MaskMergeMode::ZEROING,
	22	+ AscendC::RoundMode::UNKNOWN,
	23	+};
	24	+constexpr AscendC::MicroAPI::CastTrait kCastB322B16 = {
	25	+ AscendC::MicroAPI::RegLayout::ZERO,
	26	+ AscendC::MicroAPI::SatMode::NO_SAT,
	27	+ AscendC::MicroAPI::MaskMergeMode::ZEROING,
	28	+ AscendC::RoundMode::CAST_RINT,
	29	+};
	30	+
	31	+constexpr int64_t kPhysNodes = 5;
	32	+constexpr int64_t kMaxInputSlots = 2;
	33	+constexpr int64_t kMaxOutputSlots = 1;
	34	+static constexpr uint32_t VL_F32 = 256U / sizeof(float); // 64
	35	+
	36	+// High-precision exp constants (Cody-Waite range reduction for fp32 path).
	37	+// exp(dw) = 2^n * exp(r), where dw = n*ln2 + r, \|r\| <= ln2/2 ≈ 0.347.
	38	+// The 2^n scaling is exact (power-of-2 multiply = IEEE exponent shift, zero
	39	+// mantissa rounding). NPU Reg::Exp on the reduced range [-0.347, 0.347] is
	40	+// more accurate than on the full input range, matching libm expf more closely.
	41	+constexpr float kInvLn2 = 1.44269502162933349609375f; // 0x3fb8aa3b, 1/ln(2)
	42	+constexpr float kLn2Hi = 0.693145751953125f; // 0x3f317000, ln(2) high (12 sig bits → n*ln2Hi exact for \|n\|<2048)
	43	+constexpr float kLn2Lo = 1.428606765330187045e-06f; // ln(2) - kLn2Hi, low residual
	44	+constexpr int32_t kExpBias = 127;
	45	+constexpr int16_t kExpShift = 23;
	46	+constexpr uint32_t kF32PInf = 0x7F800000u;
	47	+constexpr uint32_t kF32NInf = 0xFF800000u;
	48	+
	49	+constexpr AscendC::MicroAPI::CastTrait kCastF2IRound = {
	50	+ AscendC::MicroAPI::RegLayout::UNKNOWN, AscendC::MicroAPI::SatMode::SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING,
	51	+ AscendC::RoundMode::CAST_ROUND};
	52	+constexpr AscendC::MicroAPI::CastTrait kCastI2FRound = {
	53	+ AscendC::MicroAPI::RegLayout::UNKNOWN, AscendC::MicroAPI::SatMode::NO_SAT,
	54	+ AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_ROUND};
	55	+
	56	+// =========================================================================
	57	+// VF 1: AnchorPreVF — pw=aHi-aLo+1, pcx=(aLo+aHi)*0.5
	58	+// Loads T from channel-contiguous aChLo/aChHi (fp16 upcast to fp32),
	59	+// stores fp32 pw/pcx to calc slots. DESIGN §5.2 S3a/S3f.
	60	+// =========================================================================
	61	+template <typename T>
	62	+__simd_vf__ inline void AnchorPreVF(__ubuf__ float* pwOut, __ubuf__ float* pcxOut, __ubuf__ T* aChLo, __ubuf__ T* aChHi,
	63	+ uint32_t count, uint16_t repeatTime)
	64	+{
	65	+ uint32_t remaining = count;
	66	+ for (uint16_t i = 0; i < repeatTime; ++i) {
	67	+ uint32_t off = static_cast<uint32_t>(i) * VL_F32;
	68	+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remaining);
	69	+ AscendC::Reg::RegTensor<float> aLo, aHi, pw, pcx;
	70	+ if constexpr (std::is_same_v<T, half>) {
	71	+ AscendC::Reg::RegTensor<half> hLo, hHi;
	72	+ AscendC::Reg::LoadAlign<half, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(hLo, aChLo + off);
	73	+ AscendC::Reg::LoadAlign<half, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(hHi, aChHi + off);
	74	+ AscendC::Reg::Cast<float, half, kCastB162B32>(aLo, hLo, mask);
	75	+ AscendC::Reg::Cast<float, half, kCastB162B32>(aHi, hHi, mask);
	76	+ } else {
	77	+ AscendC::Reg::LoadAlign(aLo, aChLo + off);
	78	+ AscendC::Reg::LoadAlign(aHi, aChHi + off);
	79	+ }
	80	+ AscendC::Reg::Sub<float>(pw, aHi, aLo, mask);
	81	+ AscendC::Reg::Adds<float>(pw, pw, 1.0f, mask);
	82	+ AscendC::Reg::Add<float>(pcx, aLo, aHi, mask);
	83	+ AscendC::Reg::Muls<float>(pcx, pcx, 0.5f, mask);
	84	+ AscendC::Reg::StoreAlign(pwOut + off, pw, mask);
	85	+ AscendC::Reg::StoreAlign(pcxOut + off, pcx, mask);
	86	+ }
	87	+}
	88	+
	89	+// =========================================================================
	90	+// VF 2: DeltaDeStdVF — dx=dLostdsLo+meansLo, dw=dHistdsHi+meansHi
	91	+// Loads T from channel-contiguous dChLo/dChHi (fp16 upcast), stores fp32
	92	+// dx/dw. DESIGN §5.2 S3b/S3g.
	93	+// =========================================================================
	94	+template <typename T>
	95	+__simd_vf__ inline void DeltaDeStdVF(__ubuf__ float* dxOut, __ubuf__ float* dwOut, __ubuf__ T* dChLo, __ubuf__ T* dChHi,
	96	+ float stdsLo, float stdsHi, float meansLo, float meansHi, uint32_t count,
	97	+ uint16_t repeatTime)
	98	+{
	99	+ uint32_t remaining = count;
	100	+ for (uint16_t i = 0; i < repeatTime; ++i) {
	101	+ uint32_t off = static_cast<uint32_t>(i) * VL_F32;
	102	+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remaining);
	103	+ AscendC::Reg::RegTensor<float> dLo, dHi, dx, dw;
	104	+ if constexpr (std::is_same_v<T, half>) {
	105	+ AscendC::Reg::RegTensor<half> hLo, hHi;
	106	+ AscendC::Reg::LoadAlign<half, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(hLo, dChLo + off);
	107	+ AscendC::Reg::LoadAlign<half, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(hHi, dChHi + off);
	108	+ AscendC::Reg::Cast<float, half, kCastB162B32>(dLo, hLo, mask);
	109	+ AscendC::Reg::Cast<float, half, kCastB162B32>(dHi, hHi, mask);
	110	+ } else {
	111	+ AscendC::Reg::LoadAlign(dLo, dChLo + off);
	112	+ AscendC::Reg::LoadAlign(dHi, dChHi + off);
	113	+ }
	114	+ AscendC::Reg::Muls<float>(dx, dLo, stdsLo, mask);
	115	+ AscendC::Reg::Adds<float>(dx, dx, meansLo, mask);
	116	+ AscendC::Reg::Muls<float>(dw, dHi, stdsHi, mask);
	117	+ AscendC::Reg::Adds<float>(dw, dw, meansHi, mask);
	118	+ AscendC::Reg::StoreAlign(dxOut + off, dx, mask);
	119	+ AscendC::Reg::StoreAlign(dwOut + off, dw, mask);
	120	+ }
	121	+}
	122	+
	123	+// =========================================================================
	124	+// VF 3: DecodeGwGxVF<T> — gw=pwexp(dw), gx=pcx+pwdx (pure fp32)
	125	+// Reads pw/pcx/dx/dw from calc slots, writes gw/gx. Slot reuse safe:
	126	+// all inputs loaded to registers before any store. DESIGN §5.2 S3c/S3h.
	127	+//
	128	+// Template parameter T selects the exp implementation:
	129	+// - T=float (fp32 path): custom Cody-Waite+Taylor FMA exp (~0.5 ULP/step,
	130	+// matching libm expf to meet fp32 rtol=atol=1e-5, DESIGN §4.3/§7.2).
	131	+// - T=half (fp16 path): NPU Reg::Exp hardware instruction. The fp16
	132	+// tolerance (rtol=atol=1e-3) is 100× more lenient than fp32, so the
	133	+// hardware exp's ~1 ULP error is sufficient. Using Reg::Exp for fp16
	134	+// also aligns the kernel's exp with np.exp (which uses the same vexp
	135	+// instruction family on SIMD platforms), reducing catastrophic-
	136	+// cancellation-induced mare spikes at near-zero outputs.
	137	+// =========================================================================
	138	+template <typename T>
	139	+__simd_vf__ inline void DecodeGwGxVF(__ubuf__ float* gwOut, __ubuf__ float* gxOut, __ubuf__ float* pwIn,
	140	+ __ubuf__ float* pcxIn, __ubuf__ float* dxIn, __ubuf__ float* dwIn, uint32_t count,
	141	+ uint16_t repeatTime)
	142	+{
	143	+ uint32_t remaining = count;
	144	+ for (uint16_t i = 0; i < repeatTime; ++i) {
	145	+ uint32_t off = static_cast<uint32_t>(i) * VL_F32;
	146	+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remaining);
	147	+ AscendC::Reg::RegTensor<float> pw, pcx, dx, dw, gw, gx, expDw, tmp;
	148	+ AscendC::Reg::LoadAlign(pw, pwIn + off);
	149	+ AscendC::Reg::LoadAlign(pcx, pcxIn + off);
	150	+ AscendC::Reg::LoadAlign(dx, dxIn + off);
	151	+ AscendC::Reg::LoadAlign(dw, dwIn + off);
	152	+ // Common ±Inf/overflow correction masks and constants (shared by
	153	+ // both exp paths below).
	154	+ AscendC::Reg::MaskReg infMask, ninfMask, ovfMask, udfMask;
	155	+ AscendC::Reg::RegTensor<float> zeroReg, infReg;
	156	+ AscendC::Reg::Duplicate<float>(zeroReg, 0.0f);
	157	+ AscendC::Reg::Duplicate((AscendC::Reg::RegTensor<uint32_t>&)infReg, kF32PInf);
	158	+ constexpr float kExpOvfThreshold = 89.0f;
	159	+ constexpr float kExpUdfThreshold = -88.0f;
	160	+ AscendC::Reg::CompareScalar<float, AscendC::CMPMODE::GE>(ovfMask, dw, kExpOvfThreshold, mask);
	161	+ AscendC::Reg::CompareScalar<float, AscendC::CMPMODE::LE>(udfMask, dw, kExpUdfThreshold, mask);
	162	+ AscendC::Reg::CompareScalar<uint32_t, AscendC::CMPMODE::EQ>(infMask, (AscendC::Reg::RegTensor<uint32_t>&)dw,
	163	+ kF32PInf, mask);
	164	+ AscendC::Reg::CompareScalar<uint32_t, AscendC::CMPMODE::EQ>(ninfMask, (AscendC::Reg::RegTensor<uint32_t>&)dw,
	165	+ kF32NInf, mask);
	166	+
	167	+ if constexpr (std::is_same_v<T, float>) {
	168	+ // ---- fp32 path: high-precision Cody-Waite+Taylor FMA exp ----
	169	+ // dw = nln2 + r (\|r\| <= ln2/2), exp(dw) = 2^n exp(r)
	170	+ // 2^n is exact (IEEE exponent shift). FMA Horner gives 0.5 ULP
	171	+ // per step, matching libm expf for fp32 rtol=atol=1e-5.
	172	+ // Range guards fix 2^n bit-pattern wrap for \|dw\| > ~88.7.
	173	+ AscendC::Reg::RegTensor<float> t, nF, r, poly, coef;
	174	+ AscendC::Reg::RegTensor<int32_t> nI, biased;
	175	+ AscendC::Reg::Muls<float>(t, dw, kInvLn2, mask);
	176	+ AscendC::Reg::Cast<int32_t, float, kCastF2IRound>(nI, t, mask);
	177	+ AscendC::Reg::Cast<float, int32_t, kCastI2FRound>(nF, nI, mask);
	178	+ AscendC::Reg::Muls<float>(r, nF, kLn2Hi, mask);
	179	+ AscendC::Reg::Sub<float>(r, dw, r, mask);
	180	+ AscendC::Reg::Muls<float>(t, nF, kLn2Lo, mask);
	181	+ AscendC::Reg::Sub<float>(r, r, t, mask);
	182	+ // exp(r) via degree-7 Taylor, FMA Horner
	183	+ AscendC::Reg::Duplicate<float>(poly, 1.9841269841269841e-04f);
	184	+ AscendC::Reg::Duplicate<float>(coef, 1.3888888888888889e-03f);
	185	+ AscendC::Reg::FusedMulDstAdd<float>(poly, r, coef, mask);
	186	+ AscendC::Reg::Duplicate<float>(coef, 8.3333333333333333e-03f);
	187	+ AscendC::Reg::FusedMulDstAdd<float>(poly, r, coef, mask);
	188	+ AscendC::Reg::Duplicate<float>(coef, 4.1666666666666667e-02f);
	189	+ AscendC::Reg::FusedMulDstAdd<float>(poly, r, coef, mask);
	190	+ AscendC::Reg::Duplicate<float>(coef, 1.6666666666666667e-01f);
	191	+ AscendC::Reg::FusedMulDstAdd<float>(poly, r, coef, mask);
	192	+ AscendC::Reg::Duplicate<float>(coef, 5.0000000000000000e-01f);
	193	+ AscendC::Reg::FusedMulDstAdd<float>(poly, r, coef, mask);
	194	+ AscendC::Reg::Duplicate<float>(coef, 1.0000000000000000e+00f);
	195	+ AscendC::Reg::FusedMulDstAdd<float>(poly, r, coef, mask);
	196	+ AscendC::Reg::Duplicate<float>(coef, 1.0000000000000000e+00f);
	197	+ AscendC::Reg::FusedMulDstAdd<float>(poly, r, coef, mask);
	198	+ AscendC::Reg::Adds<int32_t>(biased, nI, kExpBias, mask);
	199	+ AscendC::Reg::ShiftLefts<int32_t, int16_t>(biased, biased, kExpShift, mask);
	200	+ AscendC::Reg::Mul<float>(expDw, poly, (AscendC::Reg::RegTensor<float>&)biased, mask);
	201	+ } else {
	202	+ // ---- fp16 path: NPU hardware Reg::Exp ----
	203	+ // fp16 tolerance (1e-3) is 100× more lenient than fp32 (1e-5),
	204	+ // so the hardware exp's ~1 ULP error is well within tolerance.
	205	+ // Using the same vexp instruction family as np.exp reduces
	206	+ // ULP divergence at near-zero outputs (catastrophic cancellation).
	207	+ AscendC::Reg::Exp<float>(expDw, dw, mask);
	208	+ }
	209	+ // Override for ±Inf and \|dw\| beyond valid exp range (both paths).
	210	+ AscendC::Reg::Select(expDw, infReg, expDw, ovfMask); // dw >= 89 → +Inf
	211	+ AscendC::Reg::Select(expDw, zeroReg, expDw, udfMask); // dw <= -88 → 0
	212	+ AscendC::Reg::Select(expDw, infReg, expDw, infMask); // dw == +Inf → +Inf
	213	+ AscendC::Reg::Select(expDw, zeroReg, expDw, ninfMask); // dw == -Inf → 0
	214	+ AscendC::Reg::Mul<float>(gw, pw, expDw, mask);
	215	+ AscendC::Reg::Mul<float>(tmp, pw, dx, mask);
	216	+ AscendC::Reg::Add<float>(gx, pcx, tmp, mask);
	217	+ AscendC::Reg::StoreAlign(gwOut + off, gw, mask);
	218	+ AscendC::Reg::StoreAlign(gxOut + off, gx, mask);
	219	+ }
	220	+}
	221	+
	222	+// =========================================================================
	223	+// VF 4: BoxClipVF<T> — clip(gx+halfSign*gw+offset, clipLo, clipHi) → outCh
	224	+// Loads fp32 gw/gx from calc, clips, casts to T on store (fp16 downcast).
	225	+// DESIGN §5.2 S3d/e/S3i/j. clipLo=0.0f (clipped_non_negative invariant).
	226	+// =========================================================================
	227	+template <typename T>
	228	+__simd_vf__ inline void BoxClipVF(__ubuf__ T* outCh, __ubuf__ float* gwIn, __ubuf__ float* gxIn, float halfSign,
	229	+ float offset, float clipLo, float clipHi, uint32_t count, uint16_t repeatTime)
	230	+{
	231	+ uint32_t remaining = count;
	232	+ for (uint16_t i = 0; i < repeatTime; ++i) {
	233	+ uint32_t off = static_cast<uint32_t>(i) * VL_F32;
	234	+ AscendC::Reg::MaskReg mask = AscendC::Reg::UpdateMask<float>(remaining);
	235	+ AscendC::Reg::RegTensor<float> gw, gx, tmp, loReg, hiReg, out;
	236	+ AscendC::Reg::MaskReg nanMask;
	237	+ AscendC::Reg::LoadAlign(gw, gwIn + off);
	238	+ AscendC::Reg::LoadAlign(gx, gxIn + off);
	239	+ AscendC::Reg::Duplicate(loReg, clipLo);
	240	+ AscendC::Reg::Duplicate(hiReg, clipHi);
	241	+ AscendC::Reg::Muls<float>(tmp, gw, halfSign, mask);
	242	+ AscendC::Reg::Add<float>(tmp, gx, tmp, mask);
	243	+ AscendC::Reg::Adds<float>(tmp, tmp, offset, mask);
	244	+ AscendC::Reg::Max<float>(tmp, tmp, loReg, mask);
	245	+ AscendC::Reg::Min<float>(out, tmp, hiReg, mask);
	246	+ // NaN→0: extreme anchor_box (±3.4e38) causes Inf-Inf=NaN in coordinate
	247	+ // computation. Max/Min propagate NaN on NPU. Replace NaN with clipLo (0)
	248	+ // to satisfy clipped_non_negative invariant (DESIGN §4.3).
	249	+ // IEEE 754: NaN==NaN is false → EQ mask is false for NaN → Select picks loReg.
	250	+ AscendC::Reg::Compare<float, AscendC::CMPMODE::EQ>(nanMask, out, out, mask);
	251	+ AscendC::Reg::Select(out, out, loReg, nanMask);
	252	+ if constexpr (std::is_same_v<T, half>) {
	253	+ AscendC::Reg::RegTensor<half> hOut;
	254	+ AscendC::Reg::Cast<half, float, kCastB322B16>(hOut, out, mask);
	255	+ AscendC::Reg::StoreAlign<half, AscendC::Reg::StoreDist::DIST_PACK_B32>(outCh + off, hOut, mask);
	256	+ } else {
	257	+ AscendC::Reg::StoreAlign(outCh + off, out, mask);
	258	+ }
	259	+ }
	260	+}
	261	+
	262	+// =========================================================================
	263	+// Kernel Class
	264	+// buf_[0]: B_anchor [4, ubFormer] channel-contiguous (T)
	265	+// buf_[1]: B_deltas [4, ubFormer] channel-contiguous (T)
	266	+// buf_[2]: B_calc0 [4, ubFormer] fp32 (fp16 only; slots pw\|pcx\|gw\|gx)
	267	+// buf_[3]: B_calc1 [4, ubFormer] fp32 (fp16 only; slots dx\|dw\|..)
	268	+// buf_[4]: B_boxes [4, ubFormer] channel-contiguous (T)
	269	+// fp32 path: buf_[2]/buf_[3] not allocated; calc aliases IO views (K=0).
	270	+// =========================================================================
	271	+template <typename T>
	272	+class BoundingBoxDecodeKernel {
	273	+public:
	274	+ __aicore__ inline void Init(GM_ADDR anchor_box, GM_ADDR deltas, GM_ADDR boxes,
	275	+ const BoundingBoxDecodeTilingData* td)
	276	+ {
	277	+ td_ = td;
	278	+ // Empty-tensor short-circuit (DESIGN-BRANCH-1 §3/§5, DESIGN §10.4):
	279	+ // dim0==0 (N==0) returns early, skipping all buffer allocation and
	280	+ // GM setup — no CopyIn/Compute/CopyOut, no UB occupancy.
	281	+ // (was compile-time kIsEmpty; changed to runtime check to fix aclnn
	282	+ // e2e tiling buffer allocation failure — see struct.h comment.)
	283	+ if (td_->dim0 == 0) {
	284	+ return;
	285	+ }
	286	+ gmIn_[0].SetGlobalBuffer((__gm__ T*)anchor_box);
	287	+ gmIn_[1].SetGlobalBuffer((__gm__ T*)deltas);
	288	+ gmOut_[0].SetGlobalBuffer((__gm__ T*)boxes);
	289	+
	290	+ const int64_t cap = td_->ubFormer; // full tile capacity (box count)
	291	+ const int64_t ioBytes = cap * kElemsPerBox * sizeof(T);
	292	+ const int64_t calcBytes = cap * kElemsPerBox * sizeof(float);
	293	+ pipe_.InitBuffer(buf_[0], static_cast<uint32_t>(ioBytes));
	294	+ pipe_.InitBuffer(buf_[1], static_cast<uint32_t>(ioBytes));
	295	+ if constexpr (std::is_same_v<T, half>) {
	296	+ pipe_.InitBuffer(buf_[2], static_cast<uint32_t>(calcBytes));
	297	+ pipe_.InitBuffer(buf_[3], static_cast<uint32_t>(calcBytes));
	298	+ }
	299	+ pipe_.InitBuffer(buf_[4], static_cast<uint32_t>(ioBytes));
	300	+ }
	301	+
	302	+ __aicore__ inline void Process()
	303	+ {
	304	+ // Empty-tensor short-circuit (DESIGN-BRANCH-1 §3/§5, DESIGN §10.3):
	305	+ // dim0==0 (N==0) returns early — no CopyIn/Compute/CopyOut pipeline.
	306	+ if (td_->dim0 == 0) {
	307	+ return;
	308	+ }
	309	+ const int64_t blockIdx = AscendC::GetBlockIdx();
	310	+ // 512-box alignment in ComputeMultiCoreSplit can make blockNum < coreNum;
	311	+ // idle cores (blockIdx >= blockNum) have no valid data — return early to
	312	+ // avoid GM OOB read/write beyond tensor bounds.
	313	+ if (blockIdx >= td_->blockNum) {
	314	+ return;
	315	+ }
	316	+ const int64_t isLastBlock = (blockIdx == td_->blockNum - 1);
	317	+ const int64_t boxBase = blockIdx * td_->blockFormer;
	318	+ const int64_t loopNum = isLastBlock ? td_->ubLoopOfTailBlock : td_->ubLoopOfFormerBlock;
	319	+ const int64_t tailNum = isLastBlock ? td_->ubTailOfTailBlock : td_->ubTailOfFormerBlock;
	320	+
	321	+ int64_t offset = 0;
	322	+ for (int64_t i = 0; i < loopNum - 1; ++i) {
	323	+ ProcessTile(boxBase + offset, td_->ubFormer);
	324	+ offset += td_->ubFormer;
	325	+ }
	326	+ if (tailNum > 0) {
	327	+ ProcessTile(boxBase + offset, tailNum);
	328	+ }
	329	+ }
	330	+
	331	+private:
	332	+ __aicore__ inline void ProcessTile(int64_t boxOffset, int64_t boxCount)
	333	+ {
	334	+ // ===== CopyIn: NDDMA 2D stride gather, 4 channels × 2 inputs =====
	335	+ // DESIGN §3.2 / §10.7. GM [N,4] interleaved → UB [4,N] ch-contiguous.
	336	+ // blockCount=boxCount, blockLen=sizeof(T), srcStride=(4-1)*sizeof(T) bytes,
	337	+ // dstStride=0 (UB contiguous packing). Channel view stride = ubFormer.
	338	+ CopyInChannels(buf_[0], gmIn_[0], boxOffset, boxCount);
	339	+ CopyInChannels(buf_[1], gmIn_[1], boxOffset, boxCount);
	340	+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(0);
	341	+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(0);
	342	+
	343	+ // ===== Compute: x-chain (ch0,ch2) then y-chain (ch1,ch3) =====
	344	+ ProcessAxis(boxCount, 0, 2, td_->stds[0], td_->stds[2], td_->means[0], td_->means[2],
	345	+ static_cast<float>(td_->maxShapeW));
	346	+ ProcessAxis(boxCount, 1, 3, td_->stds[1], td_->stds[3], td_->means[1], td_->means[3],
	347	+ static_cast<float>(td_->maxShapeH));
	348	+
	349	+ // ===== CopyOut: 2D stride scatter, 4 channels → GM [N,4] =====
	350	+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(0);
	351	+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(0);
	352	+ CopyOutChannels(buf_[4], gmOut_[0], boxOffset, boxCount);
	353	+ AscendC::SetFlag<AscendC::HardEvent::MTE3_MTE2>(0);
	354	+ AscendC::WaitFlag<AscendC::HardEvent::MTE3_MTE2>(0);
	355	+ }
	356	+
	357	+ // One axis (x or y): chLo/chHi are the two channels for this axis.
	358	+ __aicore__ inline void ProcessAxis(int64_t boxCount, int32_t chLo, int32_t chHi, float stdsLo, float stdsHi,
	359	+ float meansLo, float meansHi, float clipHi)
	360	+ {
	361	+ const uint32_t cap = static_cast<uint32_t>(td_->ubFormer);
	362	+ const uint32_t count = static_cast<uint32_t>(boxCount);
	363	+ const uint16_t rep = static_cast<uint16_t>((count + VL_F32 - 1) / VL_F32);
	364	+
	365	+ __ubuf__ T* aBase = (__ubuf__ T*)buf_[0].Get<T>().GetPhyAddr();
	366	+ __ubuf__ T* dBase = (__ubuf__ T*)buf_[1].Get<T>().GetPhyAddr();
	367	+ __ubuf__ T* oBase = (__ubuf__ T*)buf_[4].Get<T>().GetPhyAddr();
	368	+
	369	+ // Channel views (stride = ubFormer for 32B alignment).
	370	+ __ubuf__ T* aChLo = aBase + static_cast<uint64_t>(chLo) * cap;
	371	+ __ubuf__ T* aChHi = aBase + static_cast<uint64_t>(chHi) * cap;
	372	+ __ubuf__ T* dChLo = dBase + static_cast<uint64_t>(chLo) * cap;
	373	+ __ubuf__ T* dChHi = dBase + static_cast<uint64_t>(chHi) * cap;
	374	+
	375	+ // Calc slot pointers.
	376	+ __ubuf__ float* pwOut;
	377	+ __ubuf__ float* pcxOut;
	378	+ __ubuf__ float* dxOut;
	379	+ __ubuf__ float* dwOut;
	380	+ __ubuf__ float* gwIn;
	381	+ __ubuf__ float* gxIn;
	382	+ if constexpr (std::is_same_v<T, half>) {
	383	+ // fp16: separate fp32 calc buffers. 4 slots in c0 (pw\|pcx\|gw\|gx), 2 in c1 (dx\|dw).
	384	+ __ubuf__ float* c0 = (__ubuf__ float*)buf_[2].Get<float>().GetPhyAddr();
	385	+ __ubuf__ float* c1 = (__ubuf__ float*)buf_[3].Get<float>().GetPhyAddr();
	386	+ pwOut = c0 + 0 * cap;
	387	+ pcxOut = c0 + 1 * cap;
	388	+ dxOut = c1 + 0 * cap;
	389	+ dwOut = c1 + 1 * cap;
	390	+ gwIn = c0 + 2 * cap;
	391	+ gxIn = c0 + 3 * cap;
	392	+ } else {
	393	+ // fp32: calc aliases IO buffer views (K=0, no separate calc TBuf).
	394	+ // pw→a_chLo, pcx→a_chHi, dx→d_chLo, dw→d_chHi; gw reuses pw slot, gx reuses pcx slot.
	395	+ pwOut = (__ubuf__ float*)aChLo;
	396	+ pcxOut = (__ubuf__ float*)aChHi;
	397	+ dxOut = (__ubuf__ float*)dChLo;
	398	+ dwOut = (__ubuf__ float*)dChHi;
	399	+ gwIn = (__ubuf__ float*)aChLo;
	400	+ gxIn = (__ubuf__ float*)aChHi;
	401	+ }
	402	+
	403	+ // S3a/S3f: anchor preprocess → pw, pcx
	404	+ asc_vf_call<AnchorPreVF<T>>(pwOut, pcxOut, aChLo, aChHi, count, rep);
	405	+ // S3b/S3g: deltas de-standardize → dx, dw
	406	+ asc_vf_call<DeltaDeStdVF<T>>(dxOut, dwOut, dChLo, dChHi, stdsLo, stdsHi, meansLo, meansHi, count, rep);
	407	+ // S3c/S3h: decode gw, gx (reads pw,pcx,dx,dw; writes gw,gx)
	408	+ asc_vf_call<DecodeGwGxVF<T>>(gwIn, gxIn, pwOut, pcxOut, dxOut, dwOut, count, rep);
	409	+ // S3d: ox1/oy1 = clip(gx - gw*0.5 + 0.5, 0, clipHi) → B_boxes chLo
	410	+ asc_vf_call<BoxClipVF<T>>(oBase + static_cast<uint64_t>(chLo) * cap, gwIn, gxIn, -0.5f, 0.5f, 0.0f, clipHi,
	411	+ count, rep);
	412	+ // S3e: ox2/oy2 = clip(gx + gw*0.5 - 0.5, 0, clipHi) → B_boxes chHi
	413	+ asc_vf_call<BoxClipVF<T>>(oBase + static_cast<uint64_t>(chHi) * cap, gwIn, gxIn, 0.5f, -0.5f, 0.0f, clipHi,
	414	+ count, rep);
	415	+ }
	416	+
	417	+ // NDDMA 2D stride gather: GM [N,4] interleaved → UB [4,N] ch-contiguous.
	418	+ // 4 DataCopyPad calls, one per channel. dst channel view stride = ubFormer.
	419	+ // PaddingMode::Compact packs sub-32B blocks contiguously in UB (Normal mode
	420	+ // would 32B-align each block, breaking channel-contiguous layout).
	421	+ __aicore__ inline void CopyInChannels(AscendC::TBuf<AscendC::TPosition::VECCALC>& buf, AscendC::GlobalTensor<T>& gm,
	422	+ int64_t boxOffset, int64_t boxCount)
	423	+ {
	424	+ AscendC::DataCopyExtParams params(static_cast<uint16_t>(boxCount), static_cast<uint32_t>(sizeof(T)),
	425	+ static_cast<int64_t>((kElemsPerBox - 1) * sizeof(T)), // srcStride bytes (GM)
	426	+ 0, // dstStride 0 (UB contiguous)
	427	+ 0);
	428	+ AscendC::DataCopyPadExtParams<T> pad(false, 0, 0, 0);
	429	+ const uint32_t cap = static_cast<uint32_t>(td_->ubFormer);
	430	+ for (int32_t c = 0; c < kElemsPerBox; ++c) {
	431	+ AscendC::DataCopyPad<T, AscendC::PaddingMode::Compact>(
	432	+ buf.Get<T>()[static_cast<uint32_t>(c) * cap], gm[static_cast<uint64_t>(boxOffset * kElemsPerBox + c)],
	433	+ params, pad);
	434	+ }
	435	+ }
	436	+
	437	+ // MTE3 2D stride scatter: UB [4,N] ch-contiguous → GM [N,4] interleaved.
	438	+ // 4 DataCopyPad calls, one per channel. src channel view stride = ubFormer.
	439	+ // PaddingMode::Compact reads sub-32B blocks contiguously from UB.
	440	+ __aicore__ inline void CopyOutChannels(AscendC::TBuf<AscendC::TPosition::VECCALC>& buf,
	441	+ AscendC::GlobalTensor<T>& gm, int64_t boxOffset, int64_t boxCount)
	442	+ {
	443	+ AscendC::DataCopyExtParams params(static_cast<uint16_t>(boxCount), static_cast<uint32_t>(sizeof(T)),
	444	+ 0, // srcStride 0 (UB contiguous)
	445	+ static_cast<int64_t>((kElemsPerBox - 1) * sizeof(T)), // dstStride bytes (GM)
	446	+ 0);
	447	+ const uint32_t cap = static_cast<uint32_t>(td_->ubFormer);
	448	+ for (int32_t c = 0; c < kElemsPerBox; ++c) {
	449	+ AscendC::DataCopyPad<T, AscendC::PaddingMode::Compact>(
	450	+ gm[static_cast<uint64_t>(boxOffset * kElemsPerBox + c)], buf.Get<T>()[static_cast<uint32_t>(c) * cap],
	451	+ params);
	452	+ }
	453	+ }
	454	+
	455	+ AscendC::TPipe pipe_;
	456	+ const BoundingBoxDecodeTilingData* td_;
	457	+ AscendC::GlobalTensor<T> gmIn_[kMaxInputSlots];
	458	+ AscendC::GlobalTensor<T> gmOut_[kMaxOutputSlots];
	459	+ AscendC::TBuf<AscendC::TPosition::VECCALC> buf_[kPhysNodes];
	460	+};
	461	+
	462	+#endif // BOUNDING_BOX_DECODE_KERNEL_H

Aobjdetect/bounding_box_decode/op_kernel/arch35/bounding_box_decode_struct.h+41-0

@@ -0,0 +1,41 @@
	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 BOUNDING_BOX_DECODE_STRUCT_H
	12	+#define BOUNDING_BOX_DECODE_STRUCT_H
	13	+
	14	+#include "ascendc/host_api/tiling/template_argument.h" // ASCENDC_TPL macros
	15	+
	16	+// DataType numeric constants (match ge::DataType enum values)
	17	+#define BOUNDING_BOX_DECODE_TPL_FP32 0 // ge::DT_FLOAT
	18	+#define BOUNDING_BOX_DECODE_TPL_FP16 1 // ge::DT_FLOAT16
	19	+
	20	+// ---------------------------------------------------------------------------
	21	+// TPL parameter declarations
	22	+// T : DATATYPE, selectable from {FP16, FP32}
	23	+//
	24	+// NOTE: ASCENDC_TPL_DATATYPE_* (not ASCENDC_TPL_DTYPE_*) is used so that
	25	+// the codegen emits TypeFromId<N>::type as the template argument, which
	26	+// resolves to the C++ type (half / float) and matches the
	27	+// `template <typename T>` kernel signature in proto.md.
	28	+// ASCENDC_TPL_DTYPE_* would emit raw integers, only matching `int T`.
	29	+// ---------------------------------------------------------------------------
	30	+ASCENDC_TPL_ARGS_DECL(BoundingBoxDecode,
	31	+ ASCENDC_TPL_DATATYPE_DECL(T, BOUNDING_BOX_DECODE_TPL_FP16, BOUNDING_BOX_DECODE_TPL_FP32));
	32	+
	33	+// ---------------------------------------------------------------------------
	34	+// TPL specialisations — tilingKey = T_value
	35	+// (T=FP32) → tilingKey = 0 (fp32, handles both normal and empty at runtime)
	36	+// (T=FP16) → tilingKey = 1 (fp16, handles both normal and empty at runtime)
	37	+// ---------------------------------------------------------------------------
	38	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(T, BOUNDING_BOX_DECODE_TPL_FP16)),
	39	+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(T, BOUNDING_BOX_DECODE_TPL_FP32)));
	40	+
	41	+#endif

Aobjdetect/bounding_box_decode/op_kernel/arch35/bounding_box_decode_tiling_data.h+43-0

@@ -0,0 +1,43 @@
	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 BOUNDING_BOX_DECODE_TILING_DATA_H_
	12	+#define BOUNDING_BOX_DECODE_TILING_DATA_H_
	13	+
	14	+#include <cstdint>
	15	+
	16	+// bounding_box_decode TilingData — Elementwise 按 box 数切分（tiling.md §四/§七）+ 属性标量
	17	+// 非模板：固定 rank=2，无 broadcast，无 Group；IS_EMPTY 属 TilingKey（§6）不进本结构
	18	+constexpr int64_t kElemsPerBox = 4; // C=4，每 box 4 元素 (x1,y1,x2,y2) / (dx',dy',dw',dh')
	19	+
	20	+struct BoundingBoxDecodeTilingData {
	21	+ // —— 多核切分（按 box 数，§9 多核切分；tiling.md §一/§七）—— 单位均为 box 数（除 coreNum 为核数）
	22	+ int64_t dim0; // box 总数 N（= anchor_box.shape[0]；非元素数，元素数 = dim0 × kElemsPerBox）
	23	+ int32_t coreNum; // 实际参与核数 = min(CeilDiv(dim0×minDtypeBits, MIN_TILING_BITS), availableCoreNum)
	24	+ int64_t blockFormer; // 每核基础 box 数（512 box 对齐，ELEM_ALIGN_FACTOR=512）
	25	+ int64_t blockNum; // 虚拟 block 数 = CeilDiv(dim0, blockFormer)
	26	+ // —— UB 切分（按 box 数，§9 UB 切分；tiling.md §二/§七）—— 单位均为 box 数
	27	+ int64_t ubFormer; // 每 UB 块基础 box 数（256B 对齐；alignFactor = 256 / (kElemsPerBox × sizeof(T))，fp16=32 box /
	28	+ // fp32=16 box）
	29	+ int64_t ubLoopOfFormerBlock; // 首 block 的 UB 循环次数 = CeilDiv(blockFormer, ubFormer)
	30	+ int64_t ubTailOfFormerBlock; // 首 block 尾部 box 数 = blockFormer - (ubLoopOfFormerBlock-1)×ubFormer
	31	+ int64_t ubLoopOfTailBlock; // 尾 block 的 UB 循环次数 = CeilDiv(blockTail, ubFormer)
	32	+ int64_t ubTailOfTailBlock; // 尾 block 尾部 box 数 = blockTail - (ubLoopOfTailBlock-1)×ubFormer
	33	+ // —— 属性标量（aclnn 传入，kernel Compute 消费）——
	34	+ float means[4]; // deltas 反标准化均值 m0..m3（spec.yaml attributes.means，默认 0）
	35	+ float stds[4]; // deltas 反标准化标准差 s0..s3（各元素 ≠ 0，Host 校验保证；默认 1）
	36	+ int64_t maxShapeH; // max_shape[0] = H，y 维度 (y1_out/y2_out) 裁剪上界
	37	+ int64_t maxShapeW; // max_shape[1] = W，x 维度 (x1_out/x2_out) 裁剪上界
	38	+ // 注：wh_ratio_clip 不参与核心公式（§1.3），不进 TilingData
	39	+ // 注：IS_EMPTY 不进 TilingData（属 TilingKey 模板参数，§6）
	40	+ // 注：empty 路径（IS_EMPTY=true）kernel 短路，上述计算字段填 0 不被消费
	41	+};
	42	+
	43	+#endif // BOUNDING_BOX_DECODE_TILING_DATA_H_

Aobjdetect/bounding_box_decode/op_kernel/arch35/bounding_box_decode_tiling_struct.h+16-0

@@ -0,0 +1,16 @@
	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 BOUNDING_BOX_DECODE_TILING_STRUCT_H
	12	+#define BOUNDING_BOX_DECODE_TILING_STRUCT_H
	13	+
	14	+#include "bounding_box_decode_tiling_data.h"
	15	+
	16	+#endif

Aobjdetect/bounding_box_decode/op_kernel/bounding_box_decode_apt.cpp+50-0

@@ -0,0 +1,50 @@
	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 "kernel_operator.h"
	12	+#include "arch35/bounding_box_decode_tiling_data.h"
	13	+#include "arch35/bounding_box_decode_struct.h"
	14	+#include "arch35/bounding_box_decode_kernel.h"
	15	+
	16	+template <typename T>
	17	+__global__ __aicore__ void bounding_box_decode(GM_ADDR anchor_box, GM_ADDR deltas, GM_ADDR boxes, GM_ADDR workspace,
	18	+ GM_ADDR tiling)
	19	+{
	20	+ // §10.1: TilingData registration + task type (pure Vector AIV_ONLY)
	21	+ REGISTER_TILING_DEFAULT(BoundingBoxDecodeTilingData);
	22	+ GET_TILING_DATA_WITH_STRUCT(BoundingBoxDecodeTilingData, td, tiling);
	23	+
	24	+ // §10.1: instantiate kernel class and run
	25	+ BoundingBoxDecodeKernel<T> kernel;
	26	+ kernel.Init(anchor_box, deltas, boxes, &td);
	27	+ kernel.Process();
	28	+
	29	+ (void)workspace; // workspaceSize=0 (DESIGN §9.6, no cross-core partial merge)
	30	+}
	31	+
	32	+// [REF_SAMPLE] --- original sample kernel logic (element-wise add) ---
	33	+// The block below is the reference sample implementation from the AddCustom
	34	+// template, preserved for traceability.
	35	+#if 0
	36	+ AscendC::InitSocState();
	37	+
	38	+ REGISTER_TILING_DEFAULT(BoundingBoxDecodeTilingData);
	39	+
	40	+ GET_TILING_DATA(tilingData, tiling);
	41	+
	42	+ KernelBoundingBoxDecode<DTYPE_X> op;
	43	+
	44	+ op.Init(x, y, z, tilingData.totalLength, tilingData.blockLength, tilingData.tileLength);
	45	+
	46	+ op.Process();
	47	+
	48	+ AscendC::PipeBarrier<PIPE_ALL>();
	49	+#endif
	50	+// [REF_SAMPLE] --- end ---

Aobjdetect/bounding_box_decode/tests/CMakeLists.txt+9-0

@@ -0,0 +1,9 @@
	1	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	2	+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT OR OP_KERNEL_UT))
	3	+ list(REMOVE_ITEM CURRENT_DIRS ut)
	4	+endif()
	5	+foreach(SUB_DIR ${CURRENT_DIRS})
	6	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	7	+ add_subdirectory(${SUB_DIR})
	8	+ endif()
	9	+endforeach()

Aobjdetect/bounding_box_decode/tests/assets/golden.py+142-0

@@ -0,0 +1,142 @@
	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	+
	12	+import numpy as np
	13	+import torch
	14	+
	15	+__spec__ = {"bounding_box_decode": "BoundingBoxDecodeTestSpec"}
	16	+
	17	+
	18	+class BoundingBoxDecodeTestSpec:
	19	+ """bounding_box_decode 算子测试规范（kernel 流程）
	20	+
	21	+ golden 和 third_party 均用 torch 算子拼接实现，供 cross_check 交叉比对
	22	+ """
	23	+
	24	+ def golden(
	25	+ anchor_box,
	26	+ deltas,
	27	+ *,
	28	+ max_shape,
	29	+ means=(0.0, 0.0, 0.0, 0.0),
	30	+ stds=(1.0, 1.0, 1.0, 1.0),
	31	+ wh_ratio_clip=0.016,
	32	+ **kwargs,
	33	+ ):
	34	+ x_dtype = anchor_box.dtype
	35	+
	36	+ means = [float(v) for v in means]
	37	+ stds = [float(v) for v in stds]
	38	+ max_h = float(max_shape[0])
	39	+ max_w = float(max_shape[1])
	40	+
	41	+ a = torch.from_numpy(anchor_box.astype(np.float32))
	42	+ d = torch.from_numpy(deltas.astype(np.float32))
	43	+
	44	+ a_x1, a_y1, a_x2, a_y2 = a[:, 0], a[:, 1], a[:, 2], a[:, 3]
	45	+ dx_in, dy_in, dw_in, dh_in = d[:, 0], d[:, 1], d[:, 2], d[:, 3]
	46	+
	47	+ pw = a_x2 - a_x1 + 1.0
	48	+ ph = a_y2 - a_y1 + 1.0
	49	+ pcx = (a_x1 + a_x2) * 0.5
	50	+ pcy = (a_y1 + a_y2) * 0.5
	51	+
	52	+ dx = dx_in * stds[0] + means[0]
	53	+ dy = dy_in * stds[1] + means[1]
	54	+ dw = dw_in * stds[2] + means[2]
	55	+ dh = dh_in * stds[3] + means[3]
	56	+
	57	+ gw = pw * torch.exp(dw)
	58	+ gh = ph * torch.exp(dh)
	59	+ gx = pcx + pw * dx
	60	+ gy = pcy + ph * dy
	61	+
	62	+ ox1 = gx - gw * 0.5 + 0.5
	63	+ oy1 = gy - gh * 0.5 + 0.5
	64	+ ox2 = gx + gw * 0.5 - 0.5
	65	+ oy2 = gy + gh * 0.5 - 0.5
	66	+
	67	+ ox1 = torch.nan_to_num(ox1, nan=0.0, posinf=max_w, neginf=0.0)
	68	+ oy1 = torch.nan_to_num(oy1, nan=0.0, posinf=max_h, neginf=0.0)
	69	+ ox2 = torch.nan_to_num(ox2, nan=0.0, posinf=max_w, neginf=0.0)
	70	+ oy2 = torch.nan_to_num(oy2, nan=0.0, posinf=max_h, neginf=0.0)
	71	+ ox1 = torch.clamp(ox1, 0.0, max_w)
	72	+ oy1 = torch.clamp(oy1, 0.0, max_h)
	73	+ ox2 = torch.clamp(ox2, 0.0, max_w)
	74	+ oy2 = torch.clamp(oy2, 0.0, max_h)
	75	+
	76	+ boxes = torch.stack([ox1, oy1, ox2, oy2], dim=1)
	77	+ return [boxes.numpy().astype(x_dtype)]
	78	+
	79	+ class TorchImpl:
	80	+ """torch 算子拼接实现：用 torch.exp + torch.clamp + torch.nan_to_num 拼接解码公式。"""
	81	+
	82	+ def __init__(
	83	+ self,
	84	+ *,
	85	+ max_shape,
	86	+ means=(0.0, 0.0, 0.0, 0.0),
	87	+ stds=(1.0, 1.0, 1.0, 1.0),
	88	+ wh_ratio_clip=0.016,
	89	+ **kwargs,
	90	+ ):
	91	+ self.max_h = float(max_shape[0])
	92	+ self.max_w = float(max_shape[1])
	93	+ self.means = [float(v) for v in means]
	94	+ self.stds = [float(v) for v in stds]
	95	+
	96	+ def __call__(self, anchor_box, deltas, **kwargs):
	97	+ a = anchor_box.to(torch.float32)
	98	+ d = deltas.to(torch.float32)
	99	+
	100	+ a_x1, a_y1, a_x2, a_y2 = a[:, 0], a[:, 1], a[:, 2], a[:, 3]
	101	+ dx_in, dy_in, dw_in, dh_in = d[:, 0], d[:, 1], d[:, 2], d[:, 3]
	102	+
	103	+ pw = a_x2 - a_x1 + 1.0
	104	+ ph = a_y2 - a_y1 + 1.0
	105	+ pcx = (a_x1 + a_x2) * 0.5
	106	+ pcy = (a_y1 + a_y2) * 0.5
	107	+
	108	+ dx = dx_in * self.stds[0] + self.means[0]
	109	+ dy = dy_in * self.stds[1] + self.means[1]
	110	+ dw = dw_in * self.stds[2] + self.means[2]
	111	+ dh = dh_in * self.stds[3] + self.means[3]
	112	+
	113	+ gw = pw * torch.exp(dw)
	114	+ gh = ph * torch.exp(dh)
	115	+ gx = pcx + pw * dx
	116	+ gy = pcy + ph * dy
	117	+
	118	+ ox1 = gx - gw * 0.5 + 0.5
	119	+ oy1 = gy - gh * 0.5 + 0.5
	120	+ ox2 = gx + gw * 0.5 - 0.5
	121	+ oy2 = gy + gh * 0.5 - 0.5
	122	+
	123	+ ox1 = torch.nan_to_num(ox1, nan=0.0, posinf=self.max_w, neginf=0.0)
	124	+ oy1 = torch.nan_to_num(oy1, nan=0.0, posinf=self.max_h, neginf=0.0)
	125	+ ox2 = torch.nan_to_num(ox2, nan=0.0, posinf=self.max_w, neginf=0.0)
	126	+ oy2 = torch.nan_to_num(oy2, nan=0.0, posinf=self.max_h, neginf=0.0)
	127	+ ox1 = torch.clamp(ox1, 0.0, self.max_w)
	128	+ oy1 = torch.clamp(oy1, 0.0, self.max_h)
	129	+ ox2 = torch.clamp(ox2, 0.0, self.max_w)
	130	+ oy2 = torch.clamp(oy2, 0.0, self.max_h)
	131	+
	132	+ boxes = torch.stack([ox1, oy1, ox2, oy2], dim=1)
	133	+ return [boxes.to(anchor_box.dtype)]
	134	+
	135	+ third_party = {
	136	+ "torch": TorchImpl,
	137	+ }
	138	+
	139	+ tolerance = {
	140	+ "float32": {"standard": "stat_rel_err"},
	141	+ "float16": {"standard": "stat_rel_err"},
	142	+ }

Aobjdetect/bounding_box_decode/tests/ut/CMakeLists.txt+12-0

@@ -0,0 +1,12 @@
	1	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	2	+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT))
	3	+ list(REMOVE_ITEM CURRENT_DIRS op_host)
	4	+endif()
	5	+if(NOT (UT_TEST_ALL OR OP_KERNEL_UT))
	6	+ list(REMOVE_ITEM CURRENT_DIRS op_kernel)
	7	+endif()
	8	+foreach(SUB_DIR ${CURRENT_DIRS})
	9	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	10	+ add_subdirectory(${SUB_DIR})
	11	+ endif()
	12	+endforeach()

Aobjdetect/bounding_box_decode/tests/ut/op_host/CMakeLists.txt+6-0

@@ -0,0 +1,6 @@
	1	+if(UT_TEST_ALL OR OP_HOST_UT)
	2	+ if("${ASCEND_COMPUTE_UNIT}" STREQUAL "ascend950")
	3	+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}/arch35)
	4	+ endif()
	5	+ add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	6	+endif()

Aobjdetect/bounding_box_decode/tests/ut/op_host/arch35/test_bounding_box_decode_tiling.cpp+359-0

@@ -0,0 +1,359 @@
	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	+/* Generated By CANNBot */
	12	+
	13	+#include <iostream>
	14	+#include <gtest/gtest.h>
	15	+#include "tiling_context_faker.h"
	16	+#include "tiling_case_executor.h"
	17	+#include "../../../../op_kernel/arch35/bounding_box_decode_tiling_data.h"
	18	+
	19	+using namespace std;
	20	+using namespace ge;
	21	+
	22	+class BoundingBoxDecodeTiling : public testing::Test {
	23	+protected:
	24	+ static void SetUpTestCase() { std::cout << "BoundingBoxDecodeTiling SetUp" << std::endl; }
	25	+
	26	+ static void TearDownTestCase() { std::cout << "BoundingBoxDecodeTiling TearDown" << std::endl; }
	27	+};
	28	+
	29	+std::map<std::string, std::string> soc_versions_infos = {{"Short_SoC_version", "Ascend950"}};
	30	+
	31	+constexpr size_t BBD_SYS_WORKSPACE_SIZE = 16777216;
	32	+
	33	+TEST_F(BoundingBoxDecodeTiling, bbd_fp32_normal)
	34	+{
	35	+ struct BoundingBoxDecodeCompileInfo {
	36	+ } compileInfo;
	37	+ gert::TilingContextPara tilingContextPara(
	38	+ "BoundingBoxDecode",
	39	+ {
	40	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	41	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	42	+ },
	43	+ {
	44	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	45	+ },
	46	+ {
	47	+ gert::TilingContextPara::OpAttr(
	48	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	49	+ gert::TilingContextPara::OpAttr(
	50	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	51	+ gert::TilingContextPara::OpAttr("max_shape",
	52	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	53	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	54	+ },
	55	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	56	+ uint64_t expectTilingKey = 0;
	57	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey,
	58	+ "8 1 512 1 5280 1 512 1 8 0 0 4575657222473777152 4575657222473777152 768 1280 ", {0});
	59	+}
	60	+
	61	+TEST_F(BoundingBoxDecodeTiling, bbd_fp16_normal)
	62	+{
	63	+ struct BoundingBoxDecodeCompileInfo {
	64	+ } compileInfo;
	65	+ gert::TilingContextPara tilingContextPara(
	66	+ "BoundingBoxDecode",
	67	+ {
	68	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	69	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	70	+ },
	71	+ {
	72	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	73	+ },
	74	+ {
	75	+ gert::TilingContextPara::OpAttr(
	76	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	77	+ gert::TilingContextPara::OpAttr(
	78	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	79	+ gert::TilingContextPara::OpAttr("max_shape",
	80	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	81	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	82	+ },
	83	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	84	+ uint64_t expectTilingKey = 1;
	85	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey,
	86	+ "8 1 512 1 4512 1 512 1 8 0 0 4575657222473777152 4575657222473777152 768 1280 ", {0});
	87	+}
	88	+
	89	+TEST_F(BoundingBoxDecodeTiling, bbd_fp32_empty)
	90	+{
	91	+ struct BoundingBoxDecodeCompileInfo {
	92	+ } compileInfo;
	93	+ gert::TilingContextPara tilingContextPara(
	94	+ "BoundingBoxDecode",
	95	+ {
	96	+ {{{0, 4}, {0, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	97	+ {{{0, 4}, {0, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	98	+ },
	99	+ {
	100	+ {{{0, 4}, {0, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	101	+ },
	102	+ {
	103	+ gert::TilingContextPara::OpAttr(
	104	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	105	+ gert::TilingContextPara::OpAttr(
	106	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	107	+ gert::TilingContextPara::OpAttr("max_shape",
	108	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	109	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	110	+ },
	111	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	112	+ uint64_t expectTilingKey = 0;
	113	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey,
	114	+ "0 1 0 1 0 0 0 0 0 0 0 4575657222473777152 4575657222473777152 768 1280 ", {0});
	115	+}
	116	+
	117	+TEST_F(BoundingBoxDecodeTiling, bbd_fp16_empty)
	118	+{
	119	+ struct BoundingBoxDecodeCompileInfo {
	120	+ } compileInfo;
	121	+ gert::TilingContextPara tilingContextPara(
	122	+ "BoundingBoxDecode",
	123	+ {
	124	+ {{{0, 4}, {0, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	125	+ {{{0, 4}, {0, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	126	+ },
	127	+ {
	128	+ {{{0, 4}, {0, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	129	+ },
	130	+ {
	131	+ gert::TilingContextPara::OpAttr(
	132	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	133	+ gert::TilingContextPara::OpAttr(
	134	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	135	+ gert::TilingContextPara::OpAttr("max_shape",
	136	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	137	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	138	+ },
	139	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	140	+ uint64_t expectTilingKey = 1;
	141	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey,
	142	+ "0 1 0 1 0 0 0 0 0 0 0 4575657222473777152 4575657222473777152 768 1280 ", {0});
	143	+}
	144	+
	145	+TEST_F(BoundingBoxDecodeTiling, bbd_fp32_non_default_attr)
	146	+{
	147	+ struct BoundingBoxDecodeCompileInfo {
	148	+ } compileInfo;
	149	+ gert::TilingContextPara tilingContextPara(
	150	+ "BoundingBoxDecode",
	151	+ {
	152	+ {{{16, 4}, {16, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	153	+ {{{16, 4}, {16, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	154	+ },
	155	+ {
	156	+ {{{16, 4}, {16, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	157	+ },
	158	+ {
	159	+ gert::TilingContextPara::OpAttr(
	160	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({-4.0f, -4.0f, -1.0f, -1.0f})),
	161	+ gert::TilingContextPara::OpAttr(
	162	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.1f, 0.1f, 0.2f, 0.2f})),
	163	+ gert::TilingContextPara::OpAttr("max_shape",
	164	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({1080, 1920})),
	165	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.967f)),
	166	+ },
	167	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	168	+ uint64_t expectTilingKey = 0;
	169	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey,
	170	+ "16 1 512 1 5280 1 512 1 16 -4575657218178809856 -4647714812233515008 4453159313439771853 "
	171	+ "4489188110467124429 1080 1920 ",
	172	+ {0});
	173	+}
	174	+
	175	+TEST_F(BoundingBoxDecodeTiling, bbd_fp32_large_multicore)
	176	+{
	177	+ struct BoundingBoxDecodeCompileInfo {
	178	+ } compileInfo;
	179	+ gert::TilingContextPara tilingContextPara(
	180	+ "BoundingBoxDecode",
	181	+ {
	182	+ {{{100000, 4}, {100000, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	183	+ {{{100000, 4}, {100000, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	184	+ },
	185	+ {
	186	+ {{{100000, 4}, {100000, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	187	+ },
	188	+ {
	189	+ gert::TilingContextPara::OpAttr(
	190	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	191	+ gert::TilingContextPara::OpAttr(
	192	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	193	+ gert::TilingContextPara::OpAttr("max_shape",
	194	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({1080, 1920})),
	195	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	196	+ },
	197	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	198	+ uint64_t expectTilingKey = 0;
	199	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey,
	200	+ "100000 49 2048 49 5280 1 2048 1 1696 0 0 4575657222473777152 4575657222473777152 1080 1920 ", {0});
	201	+}
	202	+
	203	+TEST_F(BoundingBoxDecodeTiling, bbd_fp16_nonaligned)
	204	+{
	205	+ struct BoundingBoxDecodeCompileInfo {
	206	+ } compileInfo;
	207	+ gert::TilingContextPara tilingContextPara(
	208	+ "BoundingBoxDecode",
	209	+ {
	210	+ {{{33, 4}, {33, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	211	+ {{{33, 4}, {33, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	212	+ },
	213	+ {
	214	+ {{{33, 4}, {33, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	215	+ },
	216	+ {
	217	+ gert::TilingContextPara::OpAttr(
	218	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	219	+ gert::TilingContextPara::OpAttr(
	220	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	221	+ gert::TilingContextPara::OpAttr("max_shape",
	222	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	223	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	224	+ },
	225	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	226	+ uint64_t expectTilingKey = 1;
	227	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey,
	228	+ "33 1 512 1 4512 1 512 1 33 0 0 4575657222473777152 4575657222473777152 768 1280 ", {0});
	229	+}
	230	+
	231	+TEST_F(BoundingBoxDecodeTiling, bbd_invalid_rank1)
	232	+{
	233	+ struct BoundingBoxDecodeCompileInfo {
	234	+ } compileInfo;
	235	+ gert::TilingContextPara tilingContextPara(
	236	+ "BoundingBoxDecode",
	237	+ {
	238	+ {{{8}, {8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	239	+ {{{8}, {8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	240	+ },
	241	+ {
	242	+ {{{8}, {8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	243	+ },
	244	+ {
	245	+ gert::TilingContextPara::OpAttr(
	246	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	247	+ gert::TilingContextPara::OpAttr(
	248	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	249	+ gert::TilingContextPara::OpAttr("max_shape",
	250	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	251	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	252	+ },
	253	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	254	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
	255	+}
	256	+
	257	+TEST_F(BoundingBoxDecodeTiling, bbd_invalid_last_dim)
	258	+{
	259	+ struct BoundingBoxDecodeCompileInfo {
	260	+ } compileInfo;
	261	+ gert::TilingContextPara tilingContextPara(
	262	+ "BoundingBoxDecode",
	263	+ {
	264	+ {{{8, 3}, {8, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
	265	+ {{{8, 3}, {8, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
	266	+ },
	267	+ {
	268	+ {{{8, 3}, {8, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
	269	+ },
	270	+ {
	271	+ gert::TilingContextPara::OpAttr(
	272	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	273	+ gert::TilingContextPara::OpAttr(
	274	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	275	+ gert::TilingContextPara::OpAttr("max_shape",
	276	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	277	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	278	+ },
	279	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	280	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
	281	+}
	282	+
	283	+TEST_F(BoundingBoxDecodeTiling, bbd_invalid_shape_mismatch)
	284	+{
	285	+ struct BoundingBoxDecodeCompileInfo {
	286	+ } compileInfo;
	287	+ gert::TilingContextPara tilingContextPara(
	288	+ "BoundingBoxDecode",
	289	+ {
	290	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	291	+ {{{16, 4}, {16, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	292	+ },
	293	+ {
	294	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	295	+ },
	296	+ {
	297	+ gert::TilingContextPara::OpAttr(
	298	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	299	+ gert::TilingContextPara::OpAttr(
	300	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	301	+ gert::TilingContextPara::OpAttr("max_shape",
	302	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	303	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	304	+ },
	305	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	306	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
	307	+}
	308	+
	309	+TEST_F(BoundingBoxDecodeTiling, bbd_invalid_dtype_int64)
	310	+{
	311	+ struct BoundingBoxDecodeCompileInfo {
	312	+ } compileInfo;
	313	+ gert::TilingContextPara tilingContextPara(
	314	+ "BoundingBoxDecode",
	315	+ {
	316	+ {{{8, 4}, {8, 4}}, ge::DT_INT64, ge::FORMAT_ND},
	317	+ {{{8, 4}, {8, 4}}, ge::DT_INT64, ge::FORMAT_ND},
	318	+ },
	319	+ {
	320	+ {{{8, 4}, {8, 4}}, ge::DT_INT64, ge::FORMAT_ND},
	321	+ },
	322	+ {
	323	+ gert::TilingContextPara::OpAttr(
	324	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	325	+ gert::TilingContextPara::OpAttr(
	326	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	327	+ gert::TilingContextPara::OpAttr("max_shape",
	328	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	329	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	330	+ },
	331	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	332	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
	333	+}
	334	+
	335	+TEST_F(BoundingBoxDecodeTiling, bbd_invalid_dtype_mismatch)
	336	+{
	337	+ struct BoundingBoxDecodeCompileInfo {
	338	+ } compileInfo;
	339	+ gert::TilingContextPara tilingContextPara(
	340	+ "BoundingBoxDecode",
	341	+ {
	342	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	343	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	344	+ },
	345	+ {
	346	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	347	+ },
	348	+ {
	349	+ gert::TilingContextPara::OpAttr(
	350	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	351	+ gert::TilingContextPara::OpAttr(
	352	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	353	+ gert::TilingContextPara::OpAttr("max_shape",
	354	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	355	+ gert::TilingContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	356	+ },
	357	+ &compileInfo, "Ascend950", 64, 262144, 4096);
	358	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
	359	+}

Aobjdetect/bounding_box_decode/tests/ut/op_host/test_bounding_box_decode_infershape.cpp+259-0

@@ -0,0 +1,259 @@
	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	+/* Generated By CANNBot */
	12	+
	13	+#include <gtest/gtest.h>
	14	+#include <iostream>
	15	+#include "infershape_context_faker.h"
	16	+#include "infershape_case_executor.h"
	17	+
	18	+class BoundingBoxDecodeInfershape : public testing::Test {
	19	+protected:
	20	+ static void SetUpTestCase() { std::cout << "BoundingBoxDecodeInfershape SetUp" << std::endl; }
	21	+
	22	+ static void TearDownTestCase() { std::cout << "BoundingBoxDecodeInfershape TearDown" << std::endl; }
	23	+};
	24	+
	25	+TEST_F(BoundingBoxDecodeInfershape, bbd_fp32_normal)
	26	+{
	27	+ gert::InfershapeContextPara infershapeContextPara(
	28	+ "BoundingBoxDecode",
	29	+ {
	30	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	31	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	32	+ },
	33	+ {
	34	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	35	+ },
	36	+ {
	37	+ gert::InfershapeContextPara::OpAttr(
	38	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	39	+ gert::InfershapeContextPara::OpAttr(
	40	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	41	+ gert::InfershapeContextPara::OpAttr("max_shape",
	42	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	43	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	44	+ });
	45	+ std::vector<std::vector<int64_t>> expectOutputShape = {{8, 4}};
	46	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	47	+}
	48	+
	49	+TEST_F(BoundingBoxDecodeInfershape, bbd_fp16_normal)
	50	+{
	51	+ gert::InfershapeContextPara infershapeContextPara(
	52	+ "BoundingBoxDecode",
	53	+ {
	54	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	55	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	56	+ },
	57	+ {
	58	+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	59	+ },
	60	+ {
	61	+ gert::InfershapeContextPara::OpAttr(
	62	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	63	+ gert::InfershapeContextPara::OpAttr(
	64	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	65	+ gert::InfershapeContextPara::OpAttr("max_shape",
	66	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	67	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	68	+ });
	69	+ std::vector<std::vector<int64_t>> expectOutputShape = {{8, 4}};
	70	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	71	+}
	72	+
	73	+TEST_F(BoundingBoxDecodeInfershape, bbd_fp32_empty)
	74	+{
	75	+ gert::InfershapeContextPara infershapeContextPara(
	76	+ "BoundingBoxDecode",
	77	+ {
	78	+ {{{0, 4}, {0, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	79	+ {{{0, 4}, {0, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	80	+ },
	81	+ {
	82	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	83	+ },
	84	+ {
	85	+ gert::InfershapeContextPara::OpAttr(
	86	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	87	+ gert::InfershapeContextPara::OpAttr(
	88	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	89	+ gert::InfershapeContextPara::OpAttr("max_shape",
	90	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	91	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	92	+ });
	93	+ std::vector<std::vector<int64_t>> expectOutputShape = {{0, 4}};
	94	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	95	+}
	96	+
	97	+TEST_F(BoundingBoxDecodeInfershape, bbd_fp32_non_default_attr)
	98	+{
	99	+ gert::InfershapeContextPara infershapeContextPara(
	100	+ "BoundingBoxDecode",
	101	+ {
	102	+ {{{16, 4}, {16, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	103	+ {{{16, 4}, {16, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	104	+ },
	105	+ {
	106	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	107	+ },
	108	+ {
	109	+ gert::InfershapeContextPara::OpAttr(
	110	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({-4.0f, -4.0f, -1.0f, -1.0f})),
	111	+ gert::InfershapeContextPara::OpAttr(
	112	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.1f, 0.1f, 0.2f, 0.2f})),
	113	+ gert::InfershapeContextPara::OpAttr("max_shape",
	114	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({1080, 1920})),
	115	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.967f)),
	116	+ });
	117	+ std::vector<std::vector<int64_t>> expectOutputShape = {{16, 4}};
	118	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	119	+}
	120	+
	121	+TEST_F(BoundingBoxDecodeInfershape, bbd_fp32_large)
	122	+{
	123	+ gert::InfershapeContextPara infershapeContextPara(
	124	+ "BoundingBoxDecode",
	125	+ {
	126	+ {{{100000, 4}, {100000, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	127	+ {{{100000, 4}, {100000, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	128	+ },
	129	+ {
	130	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	131	+ },
	132	+ {
	133	+ gert::InfershapeContextPara::OpAttr(
	134	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	135	+ gert::InfershapeContextPara::OpAttr(
	136	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	137	+ gert::InfershapeContextPara::OpAttr("max_shape",
	138	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({1080, 1920})),
	139	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	140	+ });
	141	+ std::vector<std::vector<int64_t>> expectOutputShape = {{100000, 4}};
	142	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	143	+}
	144	+
	145	+TEST_F(BoundingBoxDecodeInfershape, bbd_invalid_rank1)
	146	+{
	147	+ gert::InfershapeContextPara infershapeContextPara(
	148	+ "BoundingBoxDecode",
	149	+ {
	150	+ {{{8}, {8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	151	+ {{{8}, {8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	152	+ },
	153	+ {
	154	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	155	+ },
	156	+ {
	157	+ gert::InfershapeContextPara::OpAttr(
	158	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	159	+ gert::InfershapeContextPara::OpAttr(
	160	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	161	+ gert::InfershapeContextPara::OpAttr("max_shape",
	162	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	163	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	164	+ });
	165	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED);
	166	+}
	167	+
	168	+TEST_F(BoundingBoxDecodeInfershape, bbd_invalid_rank3)
	169	+{
	170	+ gert::InfershapeContextPara infershapeContextPara(
	171	+ "BoundingBoxDecode",
	172	+ {
	173	+ {{{8, 4, 1}, {8, 4, 1}}, ge::DT_FLOAT, ge::FORMAT_ND},
	174	+ {{{8, 4, 1}, {8, 4, 1}}, ge::DT_FLOAT, ge::FORMAT_ND},
	175	+ },
	176	+ {
	177	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	178	+ },
	179	+ {
	180	+ gert::InfershapeContextPara::OpAttr(
	181	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	182	+ gert::InfershapeContextPara::OpAttr(
	183	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	184	+ gert::InfershapeContextPara::OpAttr("max_shape",
	185	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	186	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	187	+ });
	188	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED);
	189	+}
	190	+
	191	+TEST_F(BoundingBoxDecodeInfershape, bbd_invalid_last_dim)
	192	+{
	193	+ gert::InfershapeContextPara infershapeContextPara(
	194	+ "BoundingBoxDecode",
	195	+ {
	196	+ {{{8, 3}, {8, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
	197	+ {{{8, 3}, {8, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
	198	+ },
	199	+ {
	200	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	201	+ },
	202	+ {
	203	+ gert::InfershapeContextPara::OpAttr(
	204	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	205	+ gert::InfershapeContextPara::OpAttr(
	206	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	207	+ gert::InfershapeContextPara::OpAttr("max_shape",
	208	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	209	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	210	+ });
	211	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED);
	212	+}
	213	+
	214	+TEST_F(BoundingBoxDecodeInfershape, bbd_invalid_shape_mismatch)
	215	+{
	216	+ gert::InfershapeContextPara infershapeContextPara(
	217	+ "BoundingBoxDecode",
	218	+ {
	219	+ {{{8, 4}, {8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	220	+ {{{16, 4}, {16, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
	221	+ },
	222	+ {
	223	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	224	+ },
	225	+ {
	226	+ gert::InfershapeContextPara::OpAttr(
	227	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	228	+ gert::InfershapeContextPara::OpAttr(
	229	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	230	+ gert::InfershapeContextPara::OpAttr("max_shape",
	231	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	232	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	233	+ });
	234	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED);
	235	+}
	236	+
	237	+TEST_F(BoundingBoxDecodeInfershape, bbd_unknown_rank)
	238	+{
	239	+ gert::InfershapeContextPara infershapeContextPara(
	240	+ "BoundingBoxDecode",
	241	+ {
	242	+ {{{-2}, {-2}}, ge::DT_FLOAT, ge::FORMAT_ND},
	243	+ {{{-2}, {-2}}, ge::DT_FLOAT, ge::FORMAT_ND},
	244	+ },
	245	+ {
	246	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	247	+ },
	248	+ {
	249	+ gert::InfershapeContextPara::OpAttr(
	250	+ "means", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0f, 0.0f, 0.0f, 0.0f})),
	251	+ gert::InfershapeContextPara::OpAttr(
	252	+ "stds", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({1.0f, 1.0f, 1.0f, 1.0f})),
	253	+ gert::InfershapeContextPara::OpAttr("max_shape",
	254	+ Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>({768, 1280})),
	255	+ gert::InfershapeContextPara::OpAttr("wh_ratio_clip", Ops::Cv::AnyValue::CreateFrom<float>(0.016f)),
	256	+ });
	257	+ std::vector<std::vector<int64_t>> expectOutputShape = {{-2}};
	258	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	259	+}