feat: support BoundingboxEncode ascend950-ops-cv-AtomGit

Mdocs/zh/op_list.md+10-0

	@@ -446,6 +446,16 @@
446	<td>AI Core</td>	446	<td>AI Core</td>
447	<td>将输入的新的背景图片与已有图片进行融合，通过掩码的方式将背景替换为新的背景。</td>	447	<td>将输入的新的背景图片与已有图片进行融合，通过掩码的方式将背景替换为新的背景。</td>
448	</tr>	448	</tr>
		449	+ <tr>
		450	+ <td>objdetect</td>
		451	+ <td><a href="../../objdetect/bounding_box_encode/README.md">bounding_box_encode</a></td>
		S S sunday7月9日该文件除了新增 `bounding_box_encode` 条目外，误删除了 6 个已有算子条目（extract_glimpse_v2, rgb2_yuv422, yuv4442_yuv422, d_io_u_grad, g_io_u_grad, anchor_response_flags），并修改了 ciou / roi_align_grad 的支持状态。这些变更是由于 PR 分支基于较早的 master snapshot 导致，与本次新增算子无关。建议 rebase 到最新 master 后重新生成 diff，确保只新增 bounding_box_encode 一行条目。 zhangyiyi 7月9日评论： 7月9日评论：
		452	+ <td>&check;</td>
		453	+ <td>&check;</td>
		454	+ <td>&check;</td>
		455	+ <td>&check;</td>
		456	+ <td>AI Core</td>
		457	+ <td>计算锚框与真实边界框之间的编码偏移量，生成目标检测回归目标。</td>
		458	+ </tr>
449	<tr>	459	<tr>
450	<td>objdetect</td>	460	<td>objdetect</td>
451	<td><a href="../../objdetect/ciou/README.md">ciou</a></td>	461	<td><a href="../../objdetect/ciou/README.md">ciou</a></td>

Aobjdetect/bounding_box_encode/CMakeLists.txt+28-0

@@ -0,0 +1,28 @@
	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	+# [CANNBot] Generated by CANNBot operator migration tool.
	11	+
	12	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	13	+
	14	+if(NOT ENABLE_TEST AND NOT BENCHMARK)
	15	+ list(REMOVE_ITEM CURRENT_DIRS tests)
	16	+endif()
	17	+
	18	+foreach(SUB_DIR ${CURRENT_DIRS})
	19	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	20	+ add_subdirectory(${SUB_DIR})
	21	+ endif()
	22	+endforeach()
	23	+
	24	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	25	+set(SUPPORT_TILING_DIR "arch35")
	26	+add_all_modules_sources(OPTYPE bounding_box_encode ACLNNTYPE aclnn_exclude
	27	+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
	28	+ TILING_DIR ${SUPPORT_TILING_DIR})

Aobjdetect/bounding_box_encode/README.md+108-0

@@ -0,0 +1,108 @@
	1	+# BoundingBoxEncode
	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	+- 算子功能：计算锚框（anchor box）与真实边界框（ground truth box）之间的编码偏移量，生成目标检测回归目标。
	17	+
	18	+- 计算公式：
	19	+
	20	+ 先将输入坐标 $(x_1, y_1, x_2, y_2)$ 转换为中心点+宽高格式：
	21	+
	22	+ $$
	23	+ cx = (x_1 + x_2) / 2, \quad cy = (y_1 + y_2) / 2, \quad w = x_2 - x_1 + 1, \quad h = y_2 - y_1 + 1
	24	+ $$
	25	+
	26	+ 再计算编码偏移量：
	27	+
	28	+ $$
	29	+ dx = \frac{g_{cx} - p_{cx}}{p_w}, \quad dy = \frac{g_{cy} - p_{cy}}{p_h}, \quad dw = \ln\left(\frac{g_w}{p_w}\right), \quad dh = \ln\left(\frac{g_h}{p_h}\right)
	30	+ $$
	31	+
	32	+ 最后做均值标准化：
	33	+
	34	+ $$
	35	+ \delta_i = \frac{raw_i - means_i}{stds_i}, \quad i \in \{0,1,2,3\}
	36	+ $$
	37	+
	38	+ 其中$p$为anchor_box对应值，$g$为ground_truth_box对应值。
	39	+
	40	+## 参数说明
	41	+
	42	+<table style="table-layout: fixed; width: 1100px"><colgroup>
	43	+ <col style="width: 170px">
	44	+ <col style="width: 170px">
	45	+ <col style="width: 350px">
	46	+ <col style="width: 200px">
	47	+ <col style="width: 100px">
	48	+ </colgroup>
	49	+ <thead>
	50	+ <tr>
	51	+ <th>参数名</th>
	52	+ <th>输入/输出/属性</th>
	53	+ <th>描述</th>
	54	+ <th>数据类型</th>
	55	+ <th>数据格式</th>
	56	+ </tr></thead>
	57	+ <tbody>
	58	+ <tr>
	59	+ <td>anchor_box</td>
	60	+ <td>输入</td>
	61	+ <td>锚框坐标张量，坐标格式为(x1, y1, x2, y2)。数据类型需与ground_truth_box一致。</td>
	62	+ <td>FLOAT16、FLOAT</td>
	63	+ <td>ND</td>
	64	+ </tr>
	65	+ <tr>
	66	+ <td>ground_truth_box</td>
	67	+ <td>输入</td>
	68	+ <td>真实边界框坐标张量，坐标格式为(x1, y1, x2, y2)。数据类型和shape需与anchor_box一致。</td>
	69	+ <td>FLOAT16、FLOAT</td>
	70	+ <td>ND</td>
	71	+ </tr>
	72	+ <tr>
	73	+ <td>means</td>
	74	+ <td>属性</td>
	75	+ <td>编码均值偏移量，长度为4。默认值为[0.0, 0.0, 0.0, 0.0]。</td>
	76	+ <td>ListFloat</td>
	77	+ <td>-</td>
	78	+ </tr>
	79	+ <tr>
	80	+ <td>stds</td>
	81	+ <td>属性</td>
	82	+ <td>编码标准差缩放量，长度为4，各元素不可为0。默认值为[1.0, 1.0, 1.0, 1.0]。</td>
	83	+ <td>ListFloat</td>
	84	+ <td>-</td>
	85	+ </tr>
	86	+ <tr>
	87	+ <td>delats</td>
	88	+ <td>输出</td>
	89	+ <td>编码偏移量输出张量。数据类型与anchor_box一致，shape与anchor_box相同。</td>
	90	+ <td>FLOAT16、FLOAT</td>
	91	+ <td>ND</td>
	92	+ </tr>
	93	+ </tbody></table>
	94	+
	95	+## 约束说明
	96	+
	97	+- anchor_box和ground_truth_box的数据类型必须相同，支持float16和float32。
	98	+- anchor_box和ground_truth_box的shape必须完全一致，均为(N, 4)。
	99	+- means和stds的长度必须为4，stds各元素不可为0。
	100	+- 坐标格式为标准(x1, y1, x2, y2)格式，即左上角和右下角坐标。
	101	+- 公式中宽高计算包含+1偏移（w = x2 - x1 + 1, h = y2 - y1 + 1），保证宽高至少为1，防止除零。
	102	+- 支持空Tensor（N=0时返回空输出）。
	103	+
	104	+## 调用说明
	105	+
	106	+\| 调用方式 \| 样例代码 \| 说明 \|
	107	+\| ---------------- \|-----------------------------------------------------------------------------\| --------------------------------------------------- \|
	108	+\| 图模式 \| [test_geir_bounding_box_encode](./examples/arch35/test_geir_bounding_box_encode.cpp) \| 通过[算子IR](./op_graph/bounding_box_encode_proto.h)构图方式调用BoundingBoxEncode算子。 \|

Aobjdetect/bounding_box_encode/examples/arch35/test_geir_bounding_box_encode.cpp+307-0

@@ -0,0 +1,307 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <iostream>
	12	+#include <fstream>
	13	+#include <string.h>
	14	+#include <stdint.h>
	15	+#include <vector>
	16	+#include <string>
	17	+#include <map>
	18	+#include "assert.h"
	19	+
	20	+#include "graph.h"
	21	+#include "types.h"
	22	+#include "tensor.h"
	23	+#include "ge_error_codes.h"
	24	+#include "ge_api_types.h"
	25	+#include "ge_api.h"
	26	+#include "array_ops.h"
	27	+#include "ge_ir_build.h"
	28	+
	29	+#include "../../op_graph/bounding_box_encode_proto.h"
	30	+
	31	+#define FAILED -1
	32	+#define SUCCESS 0
	33	+
	34	+using namespace ge;
	35	+using std::map;
	36	+using std::string;
	37	+using std::vector;
	38	+
	39	+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	40	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	41	+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
	42	+ TensorDesc placeholder##intputIndex##_desc = TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, \
	43	+ intputDtype); \
	44	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	45	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	46	+ Tensor tensor_placeholder##intputIndex; \
	47	+ ret = GenOnesData(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
	48	+ placeholder##intputIndex##_desc, intputDtype, 2); \
	49	+ if (ret != SUCCESS) { \
	50	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	51	+ return FAILED; \
	52	+ } \
	53	+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
	54	+ input.push_back(tensor_placeholder##intputIndex); \
	55	+ graph.AddOp(placeholder##intputIndex); \
	56	+ add1.set_input_##intputName(placeholder##intputIndex); \
	57	+ inputs.push_back(placeholder##intputIndex);
	58	+
	59	+#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	60	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	61	+ auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \
	62	+ TensorDesc placeholder##intputIndex##_desc = TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, \
	63	+ intputDtype); \
	64	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	65	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	66	+ Tensor tensor_placeholder##intputIndex; \
	67	+ ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
	68	+ placeholder##intputIndex##_desc, 1.0f); \
	69	+ if (ret != SUCCESS) { \
	70	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	71	+ return FAILED; \
	72	+ } \
	73	+ placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \
	74	+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
	75	+ graph.AddOp(placeholder##intputIndex); \
	76	+ add1.set_input_##intputName(placeholder##intputIndex); \
	77	+ add1.update_input_desc_##intputName(placeholder##intputIndex##_desc); \
	78	+ inputs.push_back(placeholder##intputIndex);
	79	+
	80	+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
	81	+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
	82	+ add1.update_output_desc_##outputName(outputName##outputIndex##_desc);
	83	+
	84	+string GetTime()
	85	+{
	86	+ time_t timep;
	87	+ time(&timep);
	88	+ char tmp[64];
	89	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
	90	+ return tmp;
	91	+}
	92	+
	93	+uint32_t GetDataTypeSize(DataType dt)
	94	+{
	95	+ uint32_t oneByte = 1;
	96	+ uint32_t twoByte = 2;
	97	+ uint32_t fourByte = 4;
	98	+ uint32_t eightByte = 8;
	99	+
	100	+ if (dt == ge::DT_FLOAT) {
	101	+ return fourByte;
	102	+ } else if (dt == ge::DT_FLOAT16) {
	103	+ return twoByte;
	104	+ } else if (dt == ge::DT_INT16) {
	105	+ return twoByte;
	106	+ } else if (dt == ge::DT_UINT16) {
	107	+ return twoByte;
	108	+ } else if (dt == ge::DT_INT32) {
	109	+ return fourByte;
	110	+ } else if (dt == ge::DT_UINT32) {
	111	+ return fourByte;
	112	+ } else if (dt == ge::DT_INT64) {
	113	+ return eightByte;
	114	+ } else if (dt == ge::DT_UINT64) {
	115	+ return eightByte;
	116	+ } else if (dt == ge::DT_INT8) {
	117	+ return oneByte;
	118	+ }
	119	+ return fourByte;
	120	+}
	121	+
	122	+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value)
	123	+{
	124	+ input_tensor_desc.SetRealDimCnt(shapes.size());
	125	+ size_t size = 1;
	126	+ for (uint32_t i = 0; i < shapes.size(); i++) {
	127	+ size *= shapes[i];
	128	+ }
	129	+ uint32_t byteSizeFloat32 = 4;
	130	+ uint32_t data_len = size * byteSizeFloat32;
	131	+ float* pData = new (std::nothrow) float[size];
	132	+ if (pData == nullptr) {
	133	+ return FAILED;
	134	+ }
	135	+ for (size_t i = 0; i < size; ++i) {
	136	+ *(pData + i) = value;
	atomgit-bot6月25日 🔵 Low Priority `GenOnesDataFloat32()` 行 135 和 `GenOnesData()` 行 153 使用 `new (std::nothrow)` 分配内存，但未检查返回值是否为 `nullptr`，随即在行 137-138 / 154-155 直接解引用 `pData`。当内存不足时 `new (std::nothrow)` 返回 `nullptr`，解引用将导致空指针崩溃。同时，分配的内存从未通过 `delete[]` 释放，每次调用都会泄漏内存。失败模式：在低内存环境下运行测试直接崩溃；长时间批量测试导致内存不断增长。
	137	+ }
	138	+ input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len);
	139	+ delete[] pData;
	140	+ return SUCCESS;
	141	+}
	142	+
	143	+int32_t GenOnesData(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type,
	144	+ int value)
	145	+{
	146	+ input_tensor_desc.SetRealDimCnt(shapes.size());
	147	+ size_t size = 1;
	148	+ for (uint32_t i = 0; i < shapes.size(); i++) {
	149	+ size *= shapes[i];
	150	+ }
	151	+ uint32_t data_len = size * GetDataTypeSize(data_type);
	152	+ int32_t* pData = new (std::nothrow) int32_t[data_len];
	153	+ if (pData == nullptr) {
	154	+ return FAILED;
	155	+ }
	156	+ for (uint32_t i = 0; i < size; ++i) {
	157	+ *(pData + i) = value;
	158	+ }
	159	+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
	160	+ delete[] pData;
	161	+ return SUCCESS;
	162	+}
	163	+
	164	+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
	165	+{
	166	+ FILE* fp = fopen(bin_file.c_str(), "wb");
	167	+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
	168	+ fclose(fp);
	169	+ return SUCCESS;
	170	+}
	171	+
	172	+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs,
	173	+ std::vector<Operator>& outputs, Graph& graph)
	174	+{
	175	+ Status ret = SUCCESS;
	176	+
	177	+ // 自定义代码：添加 BoundingBoxEncode 算子到图中
	178	+ auto add1 = op::BoundingBoxEncode("bounding_box_encode_1");
	179	+
	180	+ // anchor_box: (N, 4)
	181	+ std::vector<int64_t> boxShape = {2, 4};
	182	+ ADD_INPUT(1, anchor_box, inDtype, boxShape);
	183	+
	184	+ // ground_truth_box: (N, 4)
	185	+ ADD_INPUT(2, ground_truth_box, inDtype, boxShape);
	186	+
	187	+ // means: ListFloat attribute (length 4)
	188	+ std::vector<float> meansVec = {0.0f, 0.0f, 0.0f, 0.0f};
	189	+ add1.SetAttr("means", meansVec);
	190	+
	191	+ // stds: ListFloat attribute (length 4)
	192	+ std::vector<float> stdsVec = {1.0f, 1.0f, 1.0f, 1.0f};
	193	+ add1.SetAttr("stds", stdsVec);
	194	+
	195	+ // delats 输出: (N, 4)
	196	+ ADD_OUTPUT(1, delats, inDtype, boxShape);
	197	+
	198	+ outputs.push_back(add1);
	199	+
	200	+ return SUCCESS;
	201	+}
	202	+
	203	+int main(int argc, char* argv[])
	204	+{
	205	+ const char* graph_name = "tc_ge_irrun_test";
	206	+ Graph graph(graph_name);
	207	+ std::vector<ge::Tensor> input;
	208	+
	209	+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
	210	+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
	211	+ Status ret = ge::GEInitialize(global_options);
	212	+ if (ret != SUCCESS) {
	213	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
	214	+ return FAILED;
	215	+ }
	216	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
	217	+
	218	+ std::vector<Operator> inputs{};
	219	+ std::vector<Operator> outputs{};
	220	+
	221	+ if (argc > 1) {
	222	+ std::cout << argv[1] << std::endl;
	223	+ }
	224	+
	225	+ DataType inDtype = DT_FLOAT;
	226	+
	227	+ std::cout << "Input dtype: " << inDtype << std::endl;
	228	+
	229	+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
	230	+ if (ret != SUCCESS) {
	231	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	232	+ return FAILED;
	233	+ }
	234	+
	235	+ if (!inputs.empty() && !outputs.empty()) {
	236	+ graph.SetInputs(inputs).SetOutputs(outputs);
	237	+ }
	238	+
	239	+ std::map<AscendString, AscendString> build_options = {};
	240	+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
	241	+ ge::Session* session = new Session(build_options);
	242	+
	243	+ if (session == nullptr) {
	244	+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
	245	+ return FAILED;
	246	+ }
	247	+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
	248	+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
	249	+
	250	+ std::map<AscendString, AscendString> graph_options = {};
	251	+ uint32_t graph_id = 0;
	252	+ ret = session->AddGraph(graph_id, graph, graph_options);
	253	+
	254	+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
	255	+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
	256	+ std::string file_path = "./dump";
	257	+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
	258	+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
	259	+ std::vector<ge::Tensor> output;
	260	+ ret = session->RunGraph(graph_id, input, output);
	261	+ if (ret != SUCCESS) {
	262	+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
	263	+ delete session;
	264	+ GEFinalize();
	265	+ return FAILED;
	266	+ }
	267	+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
	268	+
	269	+ int input_num = input.size();
	270	+ for (int i = 0; i < input_num; i++) {
	271	+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
	272	+ string input_file = "./tc_ge_irrun_test_npu_input_" + std::to_string(i) + ".bin";
	273	+ uint8_t* input_data_i = input[i].GetData();
	274	+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
	275	+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
	276	+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
	277	+ WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i);
	278	+ }
	279	+
	280	+ int output_num = output.size();
	281	+ for (int i = 0; i < output_num; i++) {
	282	+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
	283	+ string output_file = "./tc_ge_irrun_test_npu_output_" + std::to_string(i) + ".bin";
	284	+ uint8_t* output_data_i = output[i].GetData();
	285	+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
	286	+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
	287	+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
	288	+ WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i);
	289	+ }
	290	+
	291	+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
	292	+ std::string error_str(error_msg.GetString());
	293	+ std::cout << "Error message: " << error_str << std::endl;
	294	+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
	295	+ std::string warning_str(warning_msg.GetString());
	296	+ std::cout << "Warning message: " << warning_str << std::endl;
	297	+ printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str());
	298	+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
	299	+ delete session;
	300	+ ret = ge::GEFinalize();
	301	+ if (ret != SUCCESS) {
	302	+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
	303	+ return FAILED;
	304	+ }
	305	+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
	306	+ return SUCCESS;
	atomgit-bot6月25日 🔵 Low Priority 在 `main()` 函数中，行 239 通过 `new Session(build_options)` 创建的 `session` 对象，仅在 `RunGraph` 失败的错误路径（行 261）执行 `delete session`；成功路径在行 303 `return SUCCESS` 之前未释放 `session`，造成内存泄漏。失败模式：每次测试运行泄漏一个 Session 对象。虽然这是测试/示例代码，单次泄漏影响有限，但若被集成到循环/批量测试中，会累积导致内存耗尽。建议：在 `return SUCCESS` 前添加 `delete session;` 释放 Session 对象。 zhangyiyi 6月29日评论： 6月29日评论：
	307	+}

Aobjdetect/bounding_box_encode/op_graph/bounding_box_encode_proto.h+67-0

@@ -0,0 +1,67 @@
	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	+/*!
	14	+ * \file bounding_box_encode_proto.h
	15	+ * \brief Encode the offset between anchor boxes and ground truth boxes.
	16	+ */
	17	+#ifndef OPS_OP_PROTO_INC_BOUNDING_BOX_ENCODE_H_
	18	+#define OPS_OP_PROTO_INC_BOUNDING_BOX_ENCODE_H_
	19	+
	20	+#include "graph/operator_reg.h"
	21	+#include "graph/types.h"
	22	+
	23	+namespace ge {
	24	+/**
	25	+* @brief Encodes the offset between anchor boxes and ground truth boxes to
	26	+* generate regression targets for object detection . \n
	27	+
	28	+* @par Inputs:
	29	+* Two inputs, including:
	30	+* @li anchor_box: Anchor boxes, a 2D Tensor of type float16 or float32
	31	+* with shape (N, 4). "N" indicates the number of anchor boxes,
	32	+* and the value "4" refers to "x1", "y1", "x2", and "y2".
	33	+* @li ground_truth_box: Ground-truth boxes, a 2D Tensor of type float16
	34	+* or float32 with shape (N, 4). It's dtype and shape should be
	35	+* same as anchor_box. The value "4" refers to "x1", "y1", "x2", and "y2" . \n
	36	+
	37	+* @par Attributes:
	38	+* @li means: A list of 4 floats, the mean values for encoding offset
	39	+* normalization, default value is [0.0, 0.0, 0.0, 0.0].
	40	+* @li stds: A list of 4 floats, the standard deviation values for encoding
	41	+* offset normalization, default value is [1.0, 1.0, 1.0, 1.0]. Each element
	42	+* must not be 0 . \n
	43	+
	44	+* @par Outputs:
	45	+* delats: Encoded offsets, a 2D Tensor of type float16 or float32
	46	+* with shape (N, 4), specifying the regression targets. It's dtype
	47	+* should be same as anchor_box . \n
	48	+
	49	+* @attention Constraints:
	50	+* The dtypes of "anchor_box" and "ground_truth_box" must be the same.
	51	+* The shapes of "anchor_box" and "ground_truth_box" must be exactly the same,
	52	+* both (N, 4). The length of "means" and "stds" must be 4, and each element of
	53	+* "stds" must not be 0. The coordinate format is standard (x1, y1, x2, y2).
	54	+* The width and height are calculated with a +1 offset (w = x2 - x1 + 1,
	55	+* h = y2 - y1 + 1) to ensure they are at least 1.
	56	+*/
	57	+REG_OP(BoundingBoxEncode)
	58	+ .INPUT(anchor_box, TensorType({DT_FLOAT16, DT_FLOAT}))
	59	+ .INPUT(ground_truth_box, TensorType({DT_FLOAT16, DT_FLOAT}))
	60	+ .OUTPUT(delats, TensorType({DT_FLOAT16, DT_FLOAT}))
	61	+ .ATTR(means, ListFloat, {0.0, 0.0, 0.0, 0.0})
	62	+ .ATTR(stds, ListFloat, {1.0, 1.0, 1.0, 1.0})
	63	+ .OP_END_FACTORY_REG(BoundingBoxEncode)
	64	+
	65	+} // namespace ge
	66	+
	67	+#endif

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

The file is empty

Aobjdetect/bounding_box_encode/op_host/arch35/bounding_box_encode_tiling_arch35.cpp+183-0

@@ -0,0 +1,183 @@
	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	+#include "op_common/log/log.h"
	15	+#include "op_common/op_host/util/math_util.h"
	16	+#include "op_common/op_host/util/platform_util.h"
	17	+#include "../../op_kernel/arch35/bounding_box_encode_tiling_data.h"
	18	+
	19	+namespace optiling {
	20	+
	21	+using Ops::Base::CeilAlign;
	22	+using Ops::Base::CeilDiv;
	23	+using Ops::Base::FloorAlign;
	24	+
	25	+constexpr uint32_t WS_SYS_SIZE = 0U;
	26	+constexpr size_t WORKSPACE_NUM = 1;
	27	+constexpr int64_t ELEM_ALIGN_FACTOR = 512;
	28	+constexpr int64_t BYTES_PER_FP16 = 2;
	29	+constexpr int64_t BYTES_PER_FP32 = 4;
	30	+constexpr int64_t BITS_PER_BYTE = 8;
	31	+constexpr int64_t ONE_BLOCK_SIZE = 32;
	32	+constexpr int64_t FP16_BITS = 16;
	33	+constexpr int64_t FP32_BITS = 32;
	34	+constexpr int64_t NUM_IO_BUFFERS = 3;
	35	+constexpr int64_t NUM_CALC_BUFFERS = 2;
	36	+constexpr int64_t NUM_TOTAL_BUFFERS = 5;
	37	+constexpr int64_t BLOCK_ELEM_THRESHOLD = 32768;
	38	+
	39	+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
	40	+{
	41	+ size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM);
	42	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	43	+ currentWorkspace[0] = WS_SYS_SIZE;
	44	+ return ge::GRAPH_SUCCESS;
	45	+}
	46	+
	47	+static ge::graphStatus ParseMeansStds(gert::TilingContext* context, BoundingBoxEncodeTilingData* tiling)
	48	+{
	49	+ auto attrs = context->GetAttrs();
	50	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
	51	+ const auto* meansList = attrs->GetListFloat(0);
	52	+ const auto* stdsList = attrs->GetListFloat(1);
	53	+ OP_CHECK_IF(meansList == nullptr \|\| stdsList == nullptr, OP_LOGE(context, "means or stds attr is null"),
	54	+ return ge::GRAPH_FAILED);
	55	+ const float* means = meansList->GetData();
	56	+ const float* stds = stdsList->GetData();
	57	+ OP_CHECK_IF(means == nullptr \|\| stds == nullptr, OP_LOGE(context, "means or stds data is null"),
	58	+ return ge::GRAPH_FAILED);
	59	+ OP_CHECK_IF(stds[0] == 0.0f \|\| stds[1] == 0.0f \|\| stds[2] == 0.0f \|\| stds[3] == 0.0f,
	60	+ OP_LOGE(context, "stds element is 0, division by zero"), return ge::GRAPH_FAILED);
	61	+
	62	+ tiling->means0 = means[0];
	63	+ tiling->means1 = means[1];
	64	+ tiling->means2 = means[2];
	65	+ tiling->means3 = means[3];
	66	+ tiling->invStds0 = 1.0f / stds[0];
	67	+ tiling->invStds1 = 1.0f / stds[1];
	68	+ tiling->invStds2 = 1.0f / stds[2];
	69	+ tiling->invStds3 = 1.0f / stds[3];
	70	+ return ge::GRAPH_SUCCESS;
	71	+}
	72	+
	73	+static ge::graphStatus CalcUbSizing(gert::TilingContext* context, BoundingBoxEncodeTilingData* tiling, int64_t dim0,
	74	+ ge::DataType dataType, uint64_t ubSize, int64_t coreNum, int64_t* outCoreNum)
	75	+{
	76	+ int64_t elemBytes = (dataType == ge::DT_FLOAT16) ? BYTES_PER_FP16 : BYTES_PER_FP32;
	77	+
	78	+ int64_t usedCoreNum = CeilDiv(dim0 * elemBytes * BITS_PER_BYTE, BLOCK_ELEM_THRESHOLD);
	79	+ if (usedCoreNum > coreNum)
	80	+ usedCoreNum = coreNum;
	81	+ if (usedCoreNum < 1)
	82	+ usedCoreNum = 1;
	83	+
	84	+ int64_t blockFactor = CeilAlign(CeilDiv(dim0, usedCoreNum), ELEM_ALIGN_FACTOR);
	85	+ usedCoreNum = CeilDiv(dim0, blockFactor);
	86	+
	87	+ // UB buffer sizing: anchor(T) + gt(T) + output(T) + calcBuf1(fp32) + calcBuf2(fp32)
	88	+ int64_t bufferDivisor;
	89	+ int64_t alignFactor;
	90	+ if (dataType == ge::DT_FLOAT16) {
	91	+ bufferDivisor = NUM_IO_BUFFERS * BYTES_PER_FP16 + NUM_CALC_BUFFERS * BYTES_PER_FP32;
	92	+ alignFactor = ONE_BLOCK_SIZE * BITS_PER_BYTE / FP16_BITS;
	93	+ } else {
	94	+ bufferDivisor = NUM_TOTAL_BUFFERS * BYTES_PER_FP32;
	95	+ alignFactor = ONE_BLOCK_SIZE * BITS_PER_BYTE / FP32_BITS;
	96	+ }
	97	+
	98	+ // Reserve 64 bytes for means+stds tensors (4 floats each, 32-byte aligned)
	99	+ constexpr int64_t RESERVED_BYTES = 64;
	100	+ int64_t maxElemNum = (static_cast<int64_t>(ubSize) - RESERVED_BYTES) / bufferDivisor;
	101	+ constexpr int64_t MAX_DATACOPY_BLOCKLEN = 65535;
	102	+ int64_t maxElemByBlockLen = MAX_DATACOPY_BLOCKLEN / elemBytes;
	103	+ if (maxElemNum > maxElemByBlockLen) {
	104	+ maxElemNum = maxElemByBlockLen;
	105	+ }
	106	+ int64_t ubFactor = FloorAlign(maxElemNum, alignFactor);
	107	+
	108	+ tiling->dim0 = dim0;
	109	+ tiling->blockFactor = blockFactor;
	110	+ tiling->ubFactor = ubFactor;
	111	+ *outCoreNum = usedCoreNum;
	112	+ return ge::GRAPH_SUCCESS;
	113	+}
	114	+
	115	+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)
	116	+{
	117	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	118	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	119	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	120	+ *coreNum = ascendcPlatform.GetCoreNumAiv();
	121	+ OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
	122	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize);
	123	+ OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
	124	+ return ge::GRAPH_SUCCESS;
	125	+}
	126	+
	127	+static ge::graphStatus BoundingBoxEncodeTilingFunc(gert::TilingContext* context)
	128	+{
	129	+ OP_LOGI(context->GetNodeName(), "Enter BoundingBoxEncodeTilingFunc");
	130	+
	131	+ uint64_t ubSize;
	132	+ int64_t coreNum;
	133	+ OP_CHECK_IF(GetPlatformInfo(context, &ubSize, &coreNum) != ge::GRAPH_SUCCESS,
	134	+ OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
	135	+
	136	+ OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
	137	+ return ge::GRAPH_FAILED);
	138	+
	139	+ auto inputAnchor = context->GetInputShape(0);
	140	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputAnchor);
	141	+ auto anchorShape = inputAnchor->GetStorageShape();
	142	+ int64_t dim0 = anchorShape.GetShapeSize();
	143	+
	144	+ auto inputDesc = context->GetInputDesc(0);
	145	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	146	+ ge::DataType dataType = inputDesc->GetDataType();
	147	+
	148	+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT16, ge::DT_FLOAT};
	149	+ OP_CHECK_IF(supportedDtype.count(dataType) == 0, OP_LOGE(context, "unsupported dtype"), return ge::GRAPH_FAILED);
	150	+
	151	+ BoundingBoxEncodeTilingData* tiling = context->GetTilingData<BoundingBoxEncodeTilingData>();
	152	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	153	+ OP_CHECK_IF(memset_s(tiling, sizeof(BoundingBoxEncodeTilingData), 0, sizeof(BoundingBoxEncodeTilingData)) != EOK,
	154	+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
	155	+
	156	+ if (dim0 == 0) {
	157	+ context->SetBlockDim(1);
	158	+ return ge::GRAPH_SUCCESS;
	159	+ }
	160	+
	161	+ OP_CHECK_IF(ParseMeansStds(context, tiling) != ge::GRAPH_SUCCESS, OP_LOGE(context, "ParseMeansStds error"),
	162	+ return ge::GRAPH_FAILED);
	163	+
	164	+ int64_t usedCoreNum = 0;
	165	+ OP_CHECK_IF(CalcUbSizing(context, tiling, dim0, dataType, ubSize, coreNum, &usedCoreNum) != ge::GRAPH_SUCCESS,
	166	+ OP_LOGE(context, "CalcUbSizing error"), return ge::GRAPH_FAILED);
	167	+
	168	+ context->SetBlockDim(usedCoreNum);
	169	+ return ge::GRAPH_SUCCESS;
	170	+}
	171	+
	172	+static ge::graphStatus TilingParseForBoundingBoxEncode([[maybe_unused]] gert::TilingParseContext* context)
	173	+{
	174	+ return ge::GRAPH_SUCCESS;
	175	+}
	176	+
	177	+struct BoundingBoxEncodeCompileInfo {};
	178	+
	179	+IMPL_OP_OPTILING(BoundingBoxEncode)
	180	+ .Tiling(BoundingBoxEncodeTilingFunc)
	181	+ .TilingParse<BoundingBoxEncodeCompileInfo>(TilingParseForBoundingBoxEncode);
	182	+
	183	+} // namespace optiling

Aobjdetect/bounding_box_encode/op_host/bounding_box_encode_def.cpp+54-0

@@ -0,0 +1,54 @@
	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 BoundingBoxEncode : public OpDef {
	17	+public:
	18	+ explicit BoundingBoxEncode(const char* name) : OpDef(name)
	19	+ {
	20	+ this->Input("anchor_box")
	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	+ this->Input("ground_truth_box")
	27	+ .ParamType(REQUIRED)
	28	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
	29	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	30	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	31	+ .AutoContiguous();
	32	+ this->Output("delats")
	33	+ .ParamType(REQUIRED)
	34	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
	35	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	36	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	37	+ .AutoContiguous();
	38	+
	39	+ this->Attr("means").AttrType(OPTIONAL).ListFloat({0.0f, 0.0f, 0.0f, 0.0f});
	40	+ this->Attr("stds").AttrType(OPTIONAL).ListFloat({1.0f, 1.0f, 1.0f, 1.0f});
	41	+
	42	+ OpAICoreConfig aiCoreConfig;
	43	+ aiCoreConfig.DynamicCompileStaticFlag(true)
	44	+ .DynamicFormatFlag(false)
	45	+ .DynamicRankSupportFlag(true)
	46	+ .DynamicShapeSupportFlag(true)
	47	+ .NeedCheckSupportFlag(false)
	48	+ .PrecisionReduceFlag(false)
	49	+ .ExtendCfgInfo("opFile.value", "bounding_box_encode_apt");
	50	+ this->AICore().AddConfig("ascend950", aiCoreConfig);
	51	+ }
	52	+};
	53	+OP_ADD(BoundingBoxEncode);
	54	+} // namespace ops

Aobjdetect/bounding_box_encode/op_host/bounding_box_encode_infershape.cpp+35-0

@@ -0,0 +1,35 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+#include "register/op_impl_registry.h"
	14	+#include "exe_graph/runtime/infer_shape_context.h"
	15	+#include "op_common/log/log.h"
	16	+
	17	+using namespace ge;
	18	+
	19	+namespace ops {
	20	+
	21	+static ge::graphStatus InferShape4BoundingBoxEncode(gert::InferShapeContext* context)
	22	+{
	23	+ const gert::Shape* inputShape = context->GetInputShape(0);
	24	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
	25	+
	26	+ gert::Shape* outputShape = context->GetOutputShape(0);
	27	+ OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);
	28	+
	29	+ outputShape = inputShape;
	30	+ return ge::GRAPH_SUCCESS;
	31	+}
	32	+
	33	+IMPL_OP_INFERSHAPE(BoundingBoxEncode).InferShape(InferShape4BoundingBoxEncode);
	34	+
	35	+} // namespace ops

Aobjdetect/bounding_box_encode/op_kernel/arch35/bounding_box_encode.h+326-0

@@ -0,0 +1,326 @@
	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	+#ifndef BOUNDING_BOX_ENCODE_H
	14	+#define BOUNDING_BOX_ENCODE_H
	15	+
	16	+#include "kernel_operator.h"
	17	+#include "kernel_tiling/kernel_tiling.h"
	18	+#include "bounding_box_encode_tiling_data.h"
	19	+
	20	+namespace NsBoundingBoxEncode {
	21	+
	22	+using namespace AscendC;
	23	+
	24	+// 公共坐标转换逻辑：消除 FP16/FP32 路径的重复代码（问题1）。
	25	+// 计算锚框/真实框的中心点、宽高及编码偏移量 dx/dy、宽高比 rw/rh。
	26	+// 注意：必须加 __aicore__ 修饰，否则在 device 侧引用参数行为异常。
	27	+struct BoxDeltaCalc {
	28	+ __aicore__ static inline void Run(float px1, float py1, float px2, float py2, float gx1, float gy1, float gx2,
	29	+ float gy2, float& dx, float& dy, float& rw, float& rh)
	30	+ {
	31	+ float pcx = (px1 + px2) * 0.5f;
	32	+ float pcy = (py1 + py2) * 0.5f;
	33	+ float pw = px2 - px1 + 1.0f;
	34	+ float ph = py2 - py1 + 1.0f;
	35	+ float gcx = (gx1 + gx2) * 0.5f;
	36	+ float gcy = (gy1 + gy2) * 0.5f;
	37	+ float gw = gx2 - gx1 + 1.0f;
	38	+ float gh = gy2 - gy1 + 1.0f;
	39	+ dx = (gcx - pcx) / pw;
	40	+ dy = (gcy - pcy) / ph;
	41	+ rw = gw / pw;
	42	+ rh = gh / ph;
	43	+ }
	44	+};
	45	+
	46	+template <typename T>
	47	+class BoundingBoxEncode {
	48	+public:
	49	+ __aicore__ inline BoundingBoxEncode(){};
	50	+ __aicore__ inline void Init(GM_ADDR anchorBox, GM_ADDR gtBox, GM_ADDR deltas, GM_ADDR workspace,
	51	+ const BoundingBoxEncodeTilingData* tilingData);
	52	+ __aicore__ inline void Process();
	53	+
	54	+private:
	55	+ __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum);
	56	+ __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum);
	57	+ __aicore__ inline void Compute(int64_t currentNum);
	58	+ // 拆分 Compute（问题3）：FP16/FP32 各自独立路径，主函数仅做分发。
	59	+ __aicore__ inline void ComputeFp16Path(AscendC::LocalTensor<T> anchorLocal, AscendC::LocalTensor<T> gtLocal,
	60	+ AscendC::LocalTensor<T> outLocal, int64_t currentNum, int64_t n);
	61	+ __aicore__ inline void ComputeFp32Path(AscendC::LocalTensor<T> anchorLocal, AscendC::LocalTensor<T> gtLocal,
	62	+ AscendC::LocalTensor<T> outLocal, int64_t currentNum, int64_t n);
	63	+
	64	+private:
	65	+ TPipe pipe;
	66	+ TQue<QuePosition::VECIN, 1> anchorQueue;
	67	+ TQue<QuePosition::VECIN, 1> gtQueue;
	68	+ TQue<QuePosition::VECOUT, 1> outputQueue;
	69	+ TBuf<QuePosition::VECCALC> calcBuf1;
	70	+ TBuf<QuePosition::VECCALC> calcBuf2;
	71	+
	72	+ GlobalTensor<T> anchorGM;
	73	+ GlobalTensor<T> gtGM;
	74	+ GlobalTensor<T> deltasGM;
	75	+
	76	+ int64_t blockLength_ = 0;
	77	+ int64_t ubLength_ = 0;
	78	+ float means0_ = 0.0f, means1_ = 0.0f, means2_ = 0.0f, means3_ = 0.0f;
	79	+ float invStds0_ = 0.0f, invStds1_ = 0.0f, invStds2_ = 0.0f, invStds3_ = 0.0f;
	80	+};
	81	+
	82	+template <typename T>
	83	+__aicore__ inline void BoundingBoxEncode<T>::Init(GM_ADDR anchorBox, GM_ADDR gtBox, GM_ADDR deltas, GM_ADDR workspace,
	84	+ const BoundingBoxEncodeTilingData* tilingData)
	85	+{
	86	+ int64_t remainderLength = tilingData->dim0 - tilingData->blockFactor * AscendC::GetBlockIdx();
	87	+ blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength;
	88	+ if (blockLength_ <= 0) {
	89	+ blockLength_ = 0;
	90	+ return;
	91	+ }
	92	+ ubLength_ = tilingData->ubFactor;
	93	+
	94	+ // Read means and stds from TilingData (passed as attributes)
	95	+ means0_ = tilingData->means0;
	96	+ means1_ = tilingData->means1;
	97	+ means2_ = tilingData->means2;
	98	+ means3_ = tilingData->means3;
	99	+ invStds0_ = tilingData->invStds0;
	100	+ invStds1_ = tilingData->invStds1;
	101	+ invStds2_ = tilingData->invStds2;
	102	+ invStds3_ = tilingData->invStds3;
	103	+
	104	+ int64_t offset = tilingData->blockFactor * AscendC::GetBlockIdx();
	105	+ anchorGM.SetGlobalBuffer((__gm__ T*)anchorBox + offset, blockLength_);
	106	+ gtGM.SetGlobalBuffer((__gm__ T*)gtBox + offset, blockLength_);
	107	+ deltasGM.SetGlobalBuffer((__gm__ T*)deltas + offset, blockLength_);
	108	+
	109	+ constexpr int64_t fpSize = static_cast<int64_t>(sizeof(float));
	110	+ if constexpr (std::is_same_v<T, half>) {
	111	+ pipe.InitBuffer(anchorQueue, 1, ubLength_ * static_cast<int64_t>(sizeof(half)));
	112	+ pipe.InitBuffer(gtQueue, 1, ubLength_ * static_cast<int64_t>(sizeof(half)));
	113	+ pipe.InitBuffer(outputQueue, 1, ubLength_ * static_cast<int64_t>(sizeof(half)));
	114	+ pipe.InitBuffer(calcBuf1, ubLength_ * fpSize);
	115	+ pipe.InitBuffer(calcBuf2, ubLength_ * fpSize);
	116	+ } else {
	117	+ pipe.InitBuffer(anchorQueue, 1, ubLength_ * fpSize);
	118	+ pipe.InitBuffer(gtQueue, 1, ubLength_ * fpSize);
	119	+ pipe.InitBuffer(outputQueue, 1, ubLength_ * fpSize);
	120	+ pipe.InitBuffer(calcBuf1, ubLength_ * fpSize);
	121	+ pipe.InitBuffer(calcBuf2, ubLength_ * fpSize);
	122	+ }
	123	+}
	124	+
	125	+template <typename T>
	126	+__aicore__ inline void BoundingBoxEncode<T>::CopyIn(int64_t progress, int64_t currentNum)
	127	+{
	128	+ AscendC::LocalTensor<T> anchorLocal = anchorQueue.template AllocTensor<T>();
	129	+ AscendC::LocalTensor<T> gtLocal = gtQueue.template AllocTensor<T>();
	130	+ // Use DataCopyExtParams (uint32_t blockLen) instead of DataCopyParams (uint16_t blockLen)
	131	+ // to avoid silent truncation when currentNum * sizeof(T) > 65535 (SEC-1.1 / SEC-2.2).
	132	+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(currentNum * static_cast<int64_t>(sizeof(T))), 0, 0,
	133	+ 0};
	134	+ AscendC::DataCopyPad(anchorLocal, anchorGM[progress * ubLength_], copyParams,
	135	+ AscendC::DataCopyPadExtParams<T>{false, 0, 0, 0});
	136	+ AscendC::DataCopyPad(gtLocal, gtGM[progress * ubLength_], copyParams,
	137	+ AscendC::DataCopyPadExtParams<T>{false, 0, 0, 0});
	138	+ anchorQueue.EnQue(anchorLocal);
	139	+ gtQueue.EnQue(gtLocal);
	140	+}
	141	+
	142	+template <typename T>
	143	+__aicore__ inline void BoundingBoxEncode<T>::CopyOut(int64_t progress, int64_t currentNum)
	144	+{
	145	+ AscendC::LocalTensor<T> outLocal = outputQueue.template DeQue<T>();
	146	+ // Use DataCopyExtParams (uint32_t blockLen) instead of DataCopyParams (uint16_t blockLen)
	147	+ // to avoid silent truncation when currentNum * sizeof(T) > 65535 (SEC-1.1 / SEC-2.2).
	148	+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(currentNum * static_cast<int64_t>(sizeof(T))), 0, 0,
	149	+ 0};
	150	+ AscendC::DataCopyPad(deltasGM[progress * ubLength_], outLocal, copyParams);
	151	+ outputQueue.FreeTensor(outLocal);
	152	+}
	153	+
	154	+template <typename T>
	155	+__aicore__ inline void BoundingBoxEncode<T>::Compute(int64_t currentNum)
	156	+{
	157	+ AscendC::LocalTensor<T> anchorLocal = anchorQueue.template DeQue<T>();
	158	+ AscendC::LocalTensor<T> gtLocal = gtQueue.template DeQue<T>();
	159	+ AscendC::LocalTensor<T> outLocal = outputQueue.template AllocTensor<T>();
	160	+ int64_t n = currentNum / 4;
	161	+
	162	+ if constexpr (std::is_same_v<T, half>) {
	163	+ ComputeFp16Path(anchorLocal, gtLocal, outLocal, currentNum, n);
	164	+ } else {
	165	+ ComputeFp32Path(anchorLocal, gtLocal, outLocal, currentNum, n);
	166	+ }
	167	+
	168	+ outputQueue.template EnQue<T>(outLocal);
	169	+ anchorQueue.FreeTensor(anchorLocal);
	170	+ gtQueue.FreeTensor(gtLocal);
	171	+}
	172	+
	173	+template <typename T>
	174	+__aicore__ inline void BoundingBoxEncode<T>::ComputeFp16Path(AscendC::LocalTensor<T> anchorLocal,
	175	+ AscendC::LocalTensor<T> gtLocal,
	176	+ AscendC::LocalTensor<T> outLocal, int64_t currentNum,
	177	+ int64_t n)
	178	+{
	179	+ AscendC::LocalTensor<float> buf1 = calcBuf1.Get<float>();
	180	+ AscendC::LocalTensor<float> buf2 = calcBuf2.Get<float>();
	181	+
	182	+ // Fetch cross-pipeline event IDs for Scalar<->Vector synchronization.
	183	+ int32_t evtSV = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::S_V));
	184	+ int32_t evtVS = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_S));
	185	+
	186	+ // --- Vector: Cast FP16 inputs to FP32 ---
	187	+ AscendC::Cast(buf1, anchorLocal, AscendC::RoundMode::CAST_NONE, currentNum);
	188	+ AscendC::Cast(buf2, gtLocal, AscendC::RoundMode::CAST_NONE, currentNum);
	189	+
	190	+ // Vector -> Scalar sync
	191	+ AscendC::SetFlag<AscendC::HardEvent::V_S>(evtVS);
	192	+ AscendC::WaitFlag<AscendC::HardEvent::V_S>(evtVS);
	193	+
	194	+ AscendC::LocalTensor<float> outFloat = outLocal.template ReinterpretCast<float>();
	195	+
	196	+ // --- Scalar loop: coordinate conversion + offset encoding (partial) ---
	197	+ for (int64_t b = 0; b < n; ++b) {
	198	+ int64_t base = b * 4;
	199	+ float dx, dy, rw, rh;
	200	+ BoxDeltaCalc::Run(buf1.GetValue(base + 0), buf1.GetValue(base + 1), buf1.GetValue(base + 2),
	201	+ buf1.GetValue(base + 3), buf2.GetValue(base + 0), buf2.GetValue(base + 1),
	202	+ buf2.GetValue(base + 2), buf2.GetValue(base + 3), dx, dy, rw, rh);
	203	+
	204	+ outFloat.SetValue(b * 2, dx);
	205	+ outFloat.SetValue(b * 2 + 1, dy);
	206	+
	207	+ buf1.SetValue(base + 0, 1.0f);
	208	+ buf1.SetValue(base + 1, 1.0f);
	209	+ buf1.SetValue(base + 2, rw);
	210	+ buf1.SetValue(base + 3, rh);
	211	+ }
	212	+
	213	+ // Scalar -> Vector sync
	214	+ AscendC::SetFlag<AscendC::HardEvent::S_V>(evtSV);
	215	+ AscendC::WaitFlag<AscendC::HardEvent::S_V>(evtSV);
	216	+
	217	+ // --- Vector: Ln ---
	218	+ AscendC::Ln(buf2, buf1, currentNum);
	219	+
	220	+ // Vector -> Scalar sync
	221	+ AscendC::SetFlag<AscendC::HardEvent::V_S>(evtVS);
	222	+ AscendC::WaitFlag<AscendC::HardEvent::V_S>(evtVS);
	223	+
	224	+ // --- Scalar loop: normalization ---
	225	+ for (int64_t b = 0; b < n; ++b) {
	226	+ int64_t base = b * 4;
	227	+ float v0 = (outFloat.GetValue(b * 2) - means0_) * invStds0_;
	228	+ float v1 = (outFloat.GetValue(b * 2 + 1) - means1_) * invStds1_;
	229	+ float v2 = (buf2.GetValue(base + 2) - means2_) * invStds2_;
	230	+ float v3 = (buf2.GetValue(base + 3) - means3_) * invStds3_;
	231	+ buf2.SetValue(base + 0, v0);
	232	+ buf2.SetValue(base + 1, v1);
	233	+ buf2.SetValue(base + 2, v2);
	234	+ buf2.SetValue(base + 3, v3);
	235	+ }
	236	+
	237	+ // Scalar -> Vector sync
	238	+ AscendC::SetFlag<AscendC::HardEvent::S_V>(evtSV);
	239	+ AscendC::WaitFlag<AscendC::HardEvent::S_V>(evtSV);
	240	+
	241	+ // --- Vector: Cast FP32 results back to FP16 ---
	242	+ AscendC::Cast(outLocal, buf2, AscendC::RoundMode::CAST_ROUND, currentNum);
	243	+}
	244	+
	245	+template <typename T>
	246	+__aicore__ inline void BoundingBoxEncode<T>::ComputeFp32Path(AscendC::LocalTensor<T> anchorLocal,
	247	+ AscendC::LocalTensor<T> gtLocal,
	248	+ AscendC::LocalTensor<T> outLocal, int64_t currentNum,
	249	+ int64_t n)
	250	+{
	251	+ AscendC::LocalTensor<float> buf1 = calcBuf1.Get<float>();
	252	+ AscendC::LocalTensor<float> buf2 = calcBuf2.Get<float>();
	253	+
	254	+ // Fetch cross-pipeline event IDs for Scalar<->Vector synchronization.
	255	+ int32_t evtSV = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::S_V));
	256	+ int32_t evtVS = static_cast<int32_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::V_S));
	257	+
	258	+ AscendC::LocalTensor<float> aFp32 = anchorLocal.template ReinterpretCast<float>();
	259	+ AscendC::LocalTensor<float> gFp32 = gtLocal.template ReinterpretCast<float>();
	260	+ AscendC::LocalTensor<float> oFp32 = outLocal.template ReinterpretCast<float>();
	261	+
	262	+ // --- Scalar loop: coordinate conversion + offset encoding (partial) ---
	263	+ for (int64_t b = 0; b < n; ++b) {
	264	+ int64_t base = b * 4;
	265	+ float dx, dy, rw, rh;
	266	+ BoxDeltaCalc::Run(aFp32.GetValue(base + 0), aFp32.GetValue(base + 1), aFp32.GetValue(base + 2),
	267	+ aFp32.GetValue(base + 3), gFp32.GetValue(base + 0), gFp32.GetValue(base + 1),
	268	+ gFp32.GetValue(base + 2), gFp32.GetValue(base + 3), dx, dy, rw, rh);
	269	+
	270	+ buf1.SetValue(b * 2, dx);
	271	+ buf1.SetValue(b * 2 + 1, dy);
	272	+
	273	+ aFp32.SetValue(base + 0, 1.0f);
	274	+ aFp32.SetValue(base + 1, 1.0f);
	275	+ aFp32.SetValue(base + 2, rw);
	276	+ aFp32.SetValue(base + 3, rh);
	277	+ }
	278	+
	279	+ // Scalar -> Vector sync
	280	+ AscendC::SetFlag<AscendC::HardEvent::S_V>(evtSV);
	281	+ AscendC::WaitFlag<AscendC::HardEvent::S_V>(evtSV);
	282	+
	283	+ // --- Vector: Ln ---
	284	+ AscendC::Ln(buf2, aFp32, currentNum);
	285	+
	286	+ // Vector -> Scalar sync
	287	+ AscendC::SetFlag<AscendC::HardEvent::V_S>(evtVS);
	288	+ AscendC::WaitFlag<AscendC::HardEvent::V_S>(evtVS);
	289	+
	290	+ // --- Scalar loop: normalization ---
	291	+ for (int64_t b = 0; b < n; ++b) {
	292	+ int64_t base = b * 4;
	293	+ float v0 = (buf1.GetValue(b * 2) - means0_) * invStds0_;
	294	+ float v1 = (buf1.GetValue(b * 2 + 1) - means1_) * invStds1_;
	295	+ float v2 = (buf2.GetValue(base + 2) - means2_) * invStds2_;
	296	+ float v3 = (buf2.GetValue(base + 3) - means3_) * invStds3_;
	297	+ aFp32.SetValue(base + 0, v0);
	298	+ aFp32.SetValue(base + 1, v1);
	299	+ aFp32.SetValue(base + 2, v2);
	300	+ aFp32.SetValue(base + 3, v3);
	301	+ }
	302	+
	303	+ // Scalar -> Vector sync
	304	+ AscendC::SetFlag<AscendC::HardEvent::S_V>(evtSV);
	305	+ AscendC::WaitFlag<AscendC::HardEvent::S_V>(evtSV);
	306	+
	307	+ // Vector copy to output buffer for correct EnQue -> MTE3 sync
	308	+ AscendC::Adds(oFp32, aFp32, 0.0f, currentNum);
	309	+}
	310	+
	311	+template <typename T>
	312	+__aicore__ inline void BoundingBoxEncode<T>::Process()
	313	+{
	314	+ if (blockLength_ <= 0)
	315	+ return;
	316	+ int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_;
	317	+ for (int64_t i = 0; i < loopCount; i++) {
	318	+ int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_;
	319	+ CopyIn(i, currentNum);
	320	+ Compute(currentNum);
	321	+ CopyOut(i, currentNum);
	322	+ }
	323	+}
	324	+
	325	+} // namespace NsBoundingBoxEncode
	326	+#endif

Aobjdetect/bounding_box_encode/op_kernel/arch35/bounding_box_encode_tiling_data.h+29-0

@@ -0,0 +1,29 @@
	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	+#ifndef _BOUNDING_BOX_ENCODE_TILING_DATA_H_
	14	+#define _BOUNDING_BOX_ENCODE_TILING_DATA_H_
	15	+
	16	+struct BoundingBoxEncodeTilingData {
	17	+ int64_t dim0 = 0;
	18	+ int64_t blockFactor = 0;
	19	+ int64_t ubFactor = 0;
	20	+ float means0 = 0.0f;
	21	+ float means1 = 0.0f;
	22	+ float means2 = 0.0f;
	23	+ float means3 = 0.0f;
	24	+ float invStds0 = 1.0f;
	25	+ float invStds1 = 1.0f;
	26	+ float invStds2 = 1.0f;
	27	+ float invStds3 = 1.0f;
	28	+};
	29	+#endif

Aobjdetect/bounding_box_encode/op_kernel/bounding_box_encode_apt.cpp+23-0

@@ -0,0 +1,23 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+#include "arch35/bounding_box_encode.h"
	14	+
	15	+extern "C" __global__ __aicore__ void bounding_box_encode(GM_ADDR anchorBox, GM_ADDR groundTruthBox, GM_ADDR delats,
	16	+ GM_ADDR workspace, GM_ADDR tiling)
	17	+{
	18	+ REGISTER_TILING_DEFAULT(BoundingBoxEncodeTilingData);
	19	+ GET_TILING_DATA_WITH_STRUCT(BoundingBoxEncodeTilingData, tilingData, tiling);
	20	+ NsBoundingBoxEncode::BoundingBoxEncode<DTYPE_ANCHOR_BOX> op;
	21	+ op.Init(anchorBox, groundTruthBox, delats, workspace, &tilingData);
	22	+ op.Process();
	23	+}