iou3d算子支持ascend950 ascendc实现-ops-cv-AtomGit

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

	@@ -798,29 +798,6 @@ REG_OP(ImageProjectiveTransformV2)
798	.OUTPUT(transformed_images, TensorType({DT_UINT8, DT_INT32, DT_INT64, DT_FLOAT16, DT_FLOAT, DT_DOUBLE}))	798	.OUTPUT(transformed_images, TensorType({DT_UINT8, DT_INT32, DT_INT64, DT_FLOAT16, DT_FLOAT, DT_DOUBLE}))
799	.OP_END_FACTORY_REG(ImageProjectiveTransformV2)	799	.OP_END_FACTORY_REG(ImageProjectiveTransformV2)
800		800
801	-/**
802	-* @brief Calculate the intersection ratio of two rotated cuboids . \n
803	-
804	-* @par Inputs:
805	-* @li bboxes : data of grad increment, a 3D Tensor of type float32 with
806	-* shape (B, 7, N). "N" indicates the number of boxes, and the value
807	-* "7" refers to [x, y, z, w, h, d, theta].
808	-* @li gtboxes: Bounding boxes, a 3D Tensor of type float32 with
809	-* shape (B, 7, K). "K" indcates the number of boxes, and the value
810	-* "7" refers to [x, y, z, w, h, d, theta].
811	-
812	-* @par Outputs:
813	-* iou: A 3D Tensor of float32 with shape [B, N, K].
814	-
815	-* @attention Constraints:
816	-* In each batch, the invalid box cannot appear before the valid box.
817	-*/
818	-REG_OP(Iou3D)
819	- .INPUT(bboxes, TensorType({DT_FLOAT}))
820	- .INPUT(gtboxes, TensorType({DT_FLOAT}))
821	- .OUTPUT(iou, TensorType({DT_FLOAT}))
822	- .OP_END_FACTORY_REG(Iou3D)
823	-
824	/**	801	/**
825	* @brief LUT3D	802	* @brief LUT3D
826	* Find the corresponding optimal pixel value for the pixel values in the input img. \n	803	* Find the corresponding optimal pixel value for the pixel values in the input img. \n

Mdocs/zh/op_list.md+10-0

	@@ -516,6 +516,16 @@
516	<td>AI Core</td>	516	<td>AI Core</td>
517	<td>计算Generalized-IoU (GIoU)损失函数的反向梯度。</td>	517	<td>计算Generalized-IoU (GIoU)损失函数的反向梯度。</td>
518	</tr>	518	</tr>
		519	+ <tr>
		520	+ <td>objdetect</td>
		521	+ <td><a href="../../objdetect/iou3d/README.md">iou3d</a></td>
		522	+ <td>&check;</td>
		523	+ <td>&check;</td>
		524	+ <td>&cross;</td>
		525	+ <td>&check;</td>
		526	+ <td>AI Core</td>
		527	+ <td>3D旋转框IoU算子：对两组7-DoF旋转框，先在BEV（鸟瞰）平面求旋转矩形交集面积，乘以Z轴重叠高度得到交集体积，再除以并集体积。</td>
		528	+ </tr>
519	<tr>	529	<tr>
520	<td>objdetect</td>	530	<td>objdetect</td>
521	<td><a href="../../objdetect/iou_v2/README.md">iou_v2</a></td>	531	<td><a href="../../objdetect/iou_v2/README.md">iou_v2</a></td>

Aobjdetect/iou3d/CMakeLists.txt+34-0

@@ -0,0 +1,34 @@
	1	+# -----------------------------------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# -----------------------------------------------------------------------------------------------------------
	10	+# Generated By CANNBot
	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	+# 设置算子定义时支持的芯片类型
	25	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	26	+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
	27	+set(SUPPORT_TILING_DIR "arch35")
	28	+# 目录名 iou3d（用户裁决保留，不改为 OpType 框架 snake_case 形式 iou3_d）。
	29	+# OPTYPE=iou3d 为本地文件命名域（仅用于 GLOB op_host/iou3d_def*.cpp + opFile.value=iou3d 的 src 路径解析），
	30	+# 与算子原型 OpType=Iou3D、aclnn 接口 aclnnIou3D、kernel 入口符号 iou3_d(=optype_snake(Iou3D)) 解耦：
	31	+# - 二进制配置走 checked-in op_host/config/ascend950/iou3d_binary.json（op_type:Iou3D），绕过 generate_bin_scripts
	32	+# 的 convert_to_camel('iou3d')='Iou3d'≠'Iou3D' round-trip 缺陷（该 gen 路径仅在无 checked-in json 时启用）。
	33	+# aclnn_exclude：使用自定义 aclnn 接口（op_api/aclnn_iou3d，对外接口名 aclnnIou3D），与本仓 ciou/iou_v2/roi_pooling_* 一致
	34	+add_all_modules_sources(OPTYPE iou3d ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
	A A andong_hw7月21日有了 add_all_modules_sources ，还需要前面add_subdirectory那一大段吗

Aobjdetect/iou3d/README.md+85-0

@@ -0,0 +1,85 @@
	1	+# Iou3D
	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	+- 算子功能：计算两组3D旋转框（7-DoF：`[x, y, z, w, h, d, theta]`）之间的3D IoU矩阵。
	17	+- 计算原理：
	18	+ - BEV投影：每个框在XY平面按`theta`旋转，得到4个顶点：`P = center ± 0.5w(cos, sin) ± 0.5h(-sin, cos)`。
	19	+ - 交集面积：顶点包含测试 + 边相交，收集交集多边形顶点；0/3顶点直接算三角形面积，>3顶点质心分解 + 极角排序求和。
	20	+ - Z轴重叠：`real_d = max(min(z1max, z2max) - max(z1min, z2min), 0)`，其中`z_min = z - 0.5d`，`z_max = z + 0.5d`。
	21	+- 计算公式：
	22	+
	23	+ $$
	24	+ \text{IoU}_{3D} = \frac{\text{bev\_area} \times \text{real\_d}}{V_A + V_B - V_{inter} + \varepsilon}, \quad \varepsilon = 10^{-6}
	25	+ $$
	26	+
	27	+ 其中`V = w * h * d`为框体积，`V_inter = bev_area * real_d`为交集体积。
	28	+- 数值稳定三守卫：分母加`epsilon = 1e-6`防除零；Z轴重叠clamp到非负；退化多边形（顶点 < 3）面积置0。
	29	+
	30	+## 参数说明
	31	+
	32	+<table style="table-layout: fixed; width: 1005px"><colgroup>
	33	+<col style="width: 170px">
	34	+<col style="width: 170px">
	35	+<col style="width: 352px">
	36	+<col style="width: 213px">
	37	+<col style="width: 100px">
	38	+</colgroup>
	39	+<thead>
	40	+ <tr>
	41	+ <th>参数名</th>
	42	+ <th>输入/输出/属性</th>
	43	+ <th>描述</th>
	44	+ <th>数据类型</th>
	45	+ <th>数据格式</th>
	46	+ </tr></thead>
	47	+<tbody>
	48	+ <tr>
	49	+ <td>bboxes</td>
	50	+ <td>输入</td>
	51	+ <td>预测框，shape <code>[B, 7, N]</code>。第二维为7-DoF <code>[x, y, z, w, h, d, theta]</code>。</td>
	52	+ <td>FLOAT</td>
	53	+ <td>ND</td>
	54	+ </tr>
	55	+ <tr>
	56	+ <td>gtboxes</td>
	57	+ <td>输入</td>
	58	+ <td>真值框，shape <code>[B, 7, K]</code>。第二维为7-DoF <code>[x, y, z, w, h, d, theta]</code>。</td>
	59	+ <td>FLOAT</td>
	60	+ <td>ND</td>
	61	+ </tr>
	62	+ <tr>
	63	+ <td>iou</td>
	64	+ <td>输出</td>
	65	+ <td>3D IoU矩阵，shape <code>[B, N, K]</code>，值域 <code>[0, 1]</code>。<code>iou[b, n, k]</code> = 第b批第n个预测框与第k个真值框的3D IoU。</td>
	66	+ <td>FLOAT</td>
	67	+ <td>ND</td>
	68	+ </tr>
	69	+</tbody></table>
	70	+
	71	+- 7-DoF语义：`x, y, z`为中心坐标，`w, h, d`为框在三个方向的尺寸，`theta`为绕Z轴的旋转角（弧度）。
	72	+- 三个Tensor的batch维`B`必须一致；`bboxes.shape[1] == gtboxes.shape[1] == 7`。
	73	+
	74	+## 约束说明
	75	+
	76	+- 数据类型：`bboxes`、`gtboxes`、`iou`均必须为float32。
	77	+- channel固定为7：`bboxes.shape[1]`与`gtboxes.shape[1]`必须等于7（7-DoF），否则拒绝。
	78	+- K 无上限：对标 mmcv，`gtboxes.shape[2]`（K）不设上限；逐对计算的UB与极角排序（固定32元素）缓冲与K无耦合，任意K成立。
	79	+- 无效框排序：同一batch内无效框（`w * h * d = 0`）不能排在有效框之前。
	80	+
	81	+## 调用说明
	82	+
	83	+\| 调用方式 \| 调用样例 \| 说明 \|
	84	+\| ------------ \| ------------ \| ------------ \|
	85	+\| GE图模式 \| [test_geir_iou3d](./examples/test_geir_iou3d.cpp) \| 通过[算子IR](./op_graph/iou3d_proto.h)构图方式调用Iou3D算子\|

Aobjdetect/iou3d/examples/arch35/test_geir_iou3d.cpp+204-0

@@ -0,0 +1,204 @@
	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	+ * Iou3D 算子 GE IR 图模式调用示例。
	11	+ *
	12	+ * 构图：Data(bboxes[B,7,N]) + Data(gtboxes[B,7,K]) -> Iou3D -> iou[B,N,K]
	13	+ * 通过 op::Iou3D（op_graph/iou3d_proto.h 注册的原型）建图，交给 ge::Session 编译执行。
	14	+ * 输入用 float32 常量数据（单位框 + 平移框），输出 dump 为 bin 供离线核对。
	15	+ */
	16	+
	17	+#include <cstdint>
	18	+#include <cstdio>
	19	+#include <cstring>
	20	+#include <ctime>
	21	+#include <map>
	22	+#include <string>
	23	+#include <vector>
	24	+
	25	+#include "ge_api.h"
	26	+#include "ge_api_types.h"
	27	+#include "ge_error_codes.h"
	28	+#include "ge_ir_build.h"
	29	+#include "graph.h"
	30	+#include "tensor.h"
	31	+#include "types.h"
	32	+#include "array_ops.h" // op::Data
	33	+
	34	+#include "../../op_graph/iou3d_proto.h"
	35	+
	36	+#define FAILED (-1)
	37	+#define SUCCESS 0
	38	+
	39	+using namespace ge;
	40	+using std::string;
	41	+using std::vector;
	42	+
	43	+static string GetTime()
	44	+{
	45	+ time_t timep;
	46	+ time(&timep);
	47	+ char tmp[64];
	48	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S", localtime(&timep));
	49	+ return tmp;
	50	+}
	51	+
	52	+// 构造 float32 host Tensor
	53	+static Tensor MakeFloatTensor(const vector<int64_t>& shape, const vector<float>& data)
	54	+{
	55	+ TensorDesc desc(ge::Shape(shape), FORMAT_ND, DT_FLOAT);
	56	+ desc.SetPlacement(ge::kPlacementHost);
	57	+ desc.SetRealDimCnt(shape.size());
	58	+ Tensor t(desc);
	59	+ t.SetData(reinterpret_cast<const uint8_t>(data.data()), data.size() sizeof(float));
	60	+ return t;
	61	+}
	62	+
	63	+static int32_t WriteBin(const string& path, const uint8_t* data, size_t size)
	64	+{
	65	+ FILE* fp = fopen(path.c_str(), "wb");
	66	+ if (fp == nullptr)
	67	+ return FAILED;
	68	+ fwrite(data, sizeof(uint8_t), size, fp);
	69	+ fclose(fp);
	70	+ return SUCCESS;
	71	+}
	72	+
	73	+int main()
	74	+{
	75	+ // B=1, N=2, K=3 —— 与 aclnn 示例同几何
	76	+ const int64_t B = 1, N = 2, K = 3;
	77	+ // bboxes[B,7,N]：DoF 顺序 x,y,z,w,h,d,theta，每 DoF 连续 N 个框
	78	+ vector<float> bboxesData = {
	79	+ 0.0f, 5.0f, // x
	80	+ 0.0f, 5.0f, // y
	81	+ 0.0f, 5.0f, // z
	82	+ 1.0f, 1.0f, // w
	83	+ 1.0f, 1.0f, // h
	84	+ 1.0f, 1.0f, // d
	85	+ 0.0f, 0.0f, // theta
	86	+ };
	87	+ // gtboxes[B,7,K]
	88	+ vector<float> gtboxesData = {
	89	+ 0.0f, 0.5f, 0.0f, // x
	90	+ 0.0f, 0.0f, 0.0f, // y
	91	+ 0.0f, 0.0f, 0.0f, // z
	92	+ 1.0f, 1.0f, 1.0f, // w
	93	+ 1.0f, 1.0f, 1.0f, // h
	94	+ 1.0f, 1.0f, 1.0f, // d
	95	+ 0.0f, 0.0f, 0.7853981634f, // theta (pi/4)
	96	+ };
	97	+ vector<int64_t> bboxesShape = {B, 7, N};
	98	+ vector<int64_t> gtboxesShape = {B, 7, K};
	99	+ vector<int64_t> iouShape = {B, N, K};
	100	+
	101	+ printf("%s - INFO - [GEIR]: GEInitialize\n", GetTime().c_str());
	102	+ std::map<AscendString, AscendString> globalOptions = {
	103	+ {"ge.exec.deviceId", "0"},
	104	+ {"ge.graphRunMode", "1"},
	105	+ };
	106	+ Status ret = ge::GEInitialize(globalOptions);
	107	+ if (ret != SUCCESS) {
	108	+ printf("%s - ERROR - [GEIR]: GEInitialize failed\n", GetTime().c_str());
	109	+ return FAILED;
	110	+ }
	111	+
	112	+ // 构图
	113	+ Graph graph("iou3d_geir_graph");
	114	+
	115	+ auto bboxesData0 = op::Data("bboxes").set_attr_index(0);
	116	+ TensorDesc bboxesDesc(ge::Shape(bboxesShape), FORMAT_ND, DT_FLOAT);
	117	+ bboxesData0.update_input_desc_x(bboxesDesc);
	118	+ bboxesData0.update_output_desc_y(bboxesDesc);
	119	+
	120	+ auto gtboxesData0 = op::Data("gtboxes").set_attr_index(1);
	121	+ TensorDesc gtboxesDesc(ge::Shape(gtboxesShape), FORMAT_ND, DT_FLOAT);
	122	+ gtboxesData0.update_input_desc_x(gtboxesDesc);
	123	+ gtboxesData0.update_output_desc_y(gtboxesDesc);
	124	+
	125	+ auto iou3d = op::Iou3D("iou3d_0");
	126	+ iou3d.set_input_bboxes(bboxesData0);
	127	+ iou3d.set_input_gtboxes(gtboxesData0);
	128	+ iou3d.update_input_desc_bboxes(bboxesDesc);
	129	+ iou3d.update_input_desc_gtboxes(gtboxesDesc);
	130	+ TensorDesc iouDesc(ge::Shape(iouShape), FORMAT_ND, DT_FLOAT);
	131	+ iou3d.update_output_desc_iou(iouDesc);
	132	+
	133	+ std::vector<Operator> inputs;
	134	+ inputs.push_back(bboxesData0);
	135	+ inputs.push_back(gtboxesData0);
	136	+ std::vector<Operator> outputs;
	137	+ outputs.push_back(iou3d);
	138	+ graph.SetInputs(inputs).SetOutputs(outputs);
	139	+
	140	+ // Session
	141	+ std::map<AscendString, AscendString> buildOptions;
	142	+ printf("%s - INFO - [GEIR]: create Session\n", GetTime().c_str());
	143	+ ge::Session* session = new Session(buildOptions);
	144	+ if (session == nullptr) {
	145	+ printf("%s - ERROR - [GEIR]: create Session failed\n", GetTime().c_str());
	146	+ ge::GEFinalize();
	147	+ return FAILED;
	148	+ }
	149	+
	150	+ std::map<AscendString, AscendString> graphOptions;
	151	+ uint32_t graphId = 0;
	152	+ ret = session->AddGraph(graphId, graph, graphOptions);
	153	+ if (ret != SUCCESS) {
	154	+ printf("%s - ERROR - [GEIR]: AddGraph failed\n", GetTime().c_str());
	155	+ delete session;
	156	+ ge::GEFinalize();
	157	+ return FAILED;
	158	+ }
	159	+ printf("%s - INFO - [GEIR]: AddGraph success\n", GetTime().c_str());
	160	+
	161	+ // 输入
	162	+ std::vector<Tensor> inputTensors;
	163	+ inputTensors.push_back(MakeFloatTensor(bboxesShape, bboxesData));
	164	+ inputTensors.push_back(MakeFloatTensor(gtboxesShape, gtboxesData));
	165	+
	166	+ printf("%s - INFO - [GEIR]: RunGraph\n", GetTime().c_str());
	167	+ std::vector<Tensor> outputTensors;
	168	+ ret = session->RunGraph(graphId, inputTensors, outputTensors);
	169	+ if (ret != SUCCESS) {
	170	+ printf("%s - ERROR - [GEIR]: RunGraph failed\n", GetTime().c_str());
	171	+ ge::AscendString errMsg = ge::GEGetErrorMsgV2();
	172	+ printf("%s - ERROR - [GEIR]: %s\n", GetTime().c_str(), errMsg.GetString());
	173	+ delete session;
	174	+ ge::GEFinalize();
	175	+ return FAILED;
	176	+ }
	177	+ printf("%s - INFO - [GEIR]: RunGraph success, output_num=%zu\n", GetTime().c_str(), outputTensors.size());
	178	+
	179	+ for (size_t i = 0; i < outputTensors.size(); ++i) {
	180	+ uint8_t* p = outputTensors[i].GetData();
	181	+ if (p == nullptr) {
	182	+ printf("%s - ERROR - [GEIR]: output[%zu] GetData returned null\n", GetTime().c_str(), i);
	183	+ continue;
	184	+ }
	185	+ size_t sz = outputTensors[i].GetSize();
	186	+ int64_t cnt = static_cast<int64_t>(sz / sizeof(float));
	187	+ printf("%s - INFO - [GEIR]: output[%zu] size=%zu (%ld floats)\n", GetTime().c_str(), i, sz, cnt);
	188	+ const float* fp = reinterpret_cast<const float*>(p);
	189	+ for (int64_t e = 0; e < cnt; ++e) {
	190	+ printf(" iou[%ld] = %.6f\n", e, fp[e]);
	191	+ }
	192	+ WriteBin("./iou3d_geir_output_" + std::to_string(i) + ".bin", p, sz);
	atomgit-bot7月13日 🟡 Medium Priority 与 `objdetect/iou3d/examples/test_geir_iou3d.cpp` 完全相同的逻辑缺陷：`examples/arch35/test_geir_iou3d.cpp` 第 179 行 `GetData()` 返回值未做空指针检查即解引用。建议：在 GetData() 后增加空指针检查：`if (p == nullptr) { printf("GetData null\n"); continue; }` 不准确？ Y yulianjie 7月21日评论： 7月21日评论：
	193	+ }
	194	+
	195	+ printf("%s - INFO - [GEIR]: GEFinalize\n", GetTime().c_str());
	196	+ delete session;
	197	+ ret = ge::GEFinalize();
	198	+ if (ret != SUCCESS) {
	199	+ printf("%s - ERROR - [GEIR]: GEFinalize failed\n", GetTime().c_str());
	200	+ return FAILED;
	201	+ }
	202	+ printf("%s - INFO - [GEIR]: done\n", GetTime().c_str());
	203	+ return SUCCESS;
	204	+}

Aobjdetect/iou3d/examples/test_geir_iou3d.cpp+204-0

@@ -0,0 +1,204 @@
	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	+ * Iou3D 算子 GE IR 图模式调用示例。
	11	+ *
	12	+ * 构图：Data(bboxes[B,7,N]) + Data(gtboxes[B,7,K]) -> Iou3D -> iou[B,N,K]
	13	+ * 通过 op::Iou3D（op_graph/iou3d_proto.h 注册的原型）建图，交给 ge::Session 编译执行。
	14	+ * 输入用 float32 常量数据（单位框 + 平移框），输出 dump 为 bin 供离线核对。
	15	+ */
	16	+
	17	+#include <cstdint>
	18	+#include <cstdio>
	19	+#include <cstring>
	20	+#include <ctime>
	21	+#include <map>
	22	+#include <string>
	23	+#include <vector>
	24	+
	25	+#include "ge_api.h"
	26	+#include "ge_api_types.h"
	27	+#include "ge_error_codes.h"
	28	+#include "ge_ir_build.h"
	29	+#include "graph.h"
	30	+#include "tensor.h"
	31	+#include "types.h"
	32	+#include "array_ops.h" // op::Data
	33	+
	34	+#include "../op_graph/iou3d_proto.h"
	35	+
	36	+#define FAILED (-1)
	37	+#define SUCCESS 0
	38	+
	39	+using namespace ge;
	40	+using std::string;
	41	+using std::vector;
	42	+
	43	+static string GetTime()
	44	+{
	45	+ time_t timep;
	46	+ time(&timep);
	47	+ char tmp[64];
	48	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S", localtime(&timep));
	49	+ return tmp;
	50	+}
	51	+
	52	+// 构造 float32 host Tensor
	53	+static Tensor MakeFloatTensor(const vector<int64_t>& shape, const vector<float>& data)
	54	+{
	55	+ TensorDesc desc(ge::Shape(shape), FORMAT_ND, DT_FLOAT);
	56	+ desc.SetPlacement(ge::kPlacementHost);
	57	+ desc.SetRealDimCnt(shape.size());
	58	+ Tensor t(desc);
	59	+ t.SetData(reinterpret_cast<const uint8_t>(data.data()), data.size() sizeof(float));
	60	+ return t;
	61	+}
	62	+
	63	+static int32_t WriteBin(const string& path, const uint8_t* data, size_t size)
	64	+{
	65	+ FILE* fp = fopen(path.c_str(), "wb");
	66	+ if (fp == nullptr)
	67	+ return FAILED;
	68	+ fwrite(data, sizeof(uint8_t), size, fp);
	69	+ fclose(fp);
	70	+ return SUCCESS;
	71	+}
	72	+
	73	+int main()
	74	+{
	75	+ // B=1, N=2, K=3 —— 与 aclnn 示例同几何
	76	+ const int64_t B = 1, N = 2, K = 3;
	77	+ // bboxes[B,7,N]：DoF 顺序 x,y,z,w,h,d,theta，每 DoF 连续 N 个框
	78	+ vector<float> bboxesData = {
	79	+ 0.0f, 5.0f, // x
	80	+ 0.0f, 5.0f, // y
	81	+ 0.0f, 5.0f, // z
	82	+ 1.0f, 1.0f, // w
	83	+ 1.0f, 1.0f, // h
	84	+ 1.0f, 1.0f, // d
	85	+ 0.0f, 0.0f, // theta
	86	+ };
	87	+ // gtboxes[B,7,K]
	88	+ vector<float> gtboxesData = {
	89	+ 0.0f, 0.5f, 0.0f, // x
	90	+ 0.0f, 0.0f, 0.0f, // y
	91	+ 0.0f, 0.0f, 0.0f, // z
	92	+ 1.0f, 1.0f, 1.0f, // w
	93	+ 1.0f, 1.0f, 1.0f, // h
	94	+ 1.0f, 1.0f, 1.0f, // d
	95	+ 0.0f, 0.0f, 0.7853981634f, // theta (pi/4)
	96	+ };
	97	+ vector<int64_t> bboxesShape = {B, 7, N};
	98	+ vector<int64_t> gtboxesShape = {B, 7, K};
	99	+ vector<int64_t> iouShape = {B, N, K};
	100	+
	101	+ printf("%s - INFO - [GEIR]: GEInitialize\n", GetTime().c_str());
	102	+ std::map<AscendString, AscendString> globalOptions = {
	103	+ {"ge.exec.deviceId", "0"},
	104	+ {"ge.graphRunMode", "1"},
	105	+ };
	106	+ Status ret = ge::GEInitialize(globalOptions);
	107	+ if (ret != SUCCESS) {
	108	+ printf("%s - ERROR - [GEIR]: GEInitialize failed\n", GetTime().c_str());
	109	+ return FAILED;
	110	+ }
	111	+
	112	+ // 构图
	113	+ Graph graph("iou3d_geir_graph");
	114	+
	115	+ auto bboxesData0 = op::Data("bboxes").set_attr_index(0);
	116	+ TensorDesc bboxesDesc(ge::Shape(bboxesShape), FORMAT_ND, DT_FLOAT);
	117	+ bboxesData0.update_input_desc_x(bboxesDesc);
	118	+ bboxesData0.update_output_desc_y(bboxesDesc);
	119	+
	120	+ auto gtboxesData0 = op::Data("gtboxes").set_attr_index(1);
	121	+ TensorDesc gtboxesDesc(ge::Shape(gtboxesShape), FORMAT_ND, DT_FLOAT);
	122	+ gtboxesData0.update_input_desc_x(gtboxesDesc);
	123	+ gtboxesData0.update_output_desc_y(gtboxesDesc);
	124	+
	125	+ auto iou3d = op::Iou3D("iou3d_0");
	126	+ iou3d.set_input_bboxes(bboxesData0);
	127	+ iou3d.set_input_gtboxes(gtboxesData0);
	128	+ iou3d.update_input_desc_bboxes(bboxesDesc);
	129	+ iou3d.update_input_desc_gtboxes(gtboxesDesc);
	130	+ TensorDesc iouDesc(ge::Shape(iouShape), FORMAT_ND, DT_FLOAT);
	131	+ iou3d.update_output_desc_iou(iouDesc);
	132	+
	133	+ std::vector<Operator> inputs;
	134	+ inputs.push_back(bboxesData0);
	135	+ inputs.push_back(gtboxesData0);
	136	+ std::vector<Operator> outputs;
	137	+ outputs.push_back(iou3d);
	138	+ graph.SetInputs(inputs).SetOutputs(outputs);
	139	+
	140	+ // Session
	141	+ std::map<AscendString, AscendString> buildOptions;
	142	+ printf("%s - INFO - [GEIR]: create Session\n", GetTime().c_str());
	143	+ ge::Session* session = new Session(buildOptions);
	144	+ if (session == nullptr) {
	145	+ printf("%s - ERROR - [GEIR]: create Session failed\n", GetTime().c_str());
	146	+ ge::GEFinalize();
	147	+ return FAILED;
	148	+ }
	149	+
	150	+ std::map<AscendString, AscendString> graphOptions;
	151	+ uint32_t graphId = 0;
	152	+ ret = session->AddGraph(graphId, graph, graphOptions);
	153	+ if (ret != SUCCESS) {
	154	+ printf("%s - ERROR - [GEIR]: AddGraph failed\n", GetTime().c_str());
	155	+ delete session;
	156	+ ge::GEFinalize();
	157	+ return FAILED;
	158	+ }
	159	+ printf("%s - INFO - [GEIR]: AddGraph success\n", GetTime().c_str());
	160	+
	161	+ // 输入
	162	+ std::vector<Tensor> inputTensors;
	163	+ inputTensors.push_back(MakeFloatTensor(bboxesShape, bboxesData));
	164	+ inputTensors.push_back(MakeFloatTensor(gtboxesShape, gtboxesData));
	165	+
	166	+ printf("%s - INFO - [GEIR]: RunGraph\n", GetTime().c_str());
	167	+ std::vector<Tensor> outputTensors;
	168	+ ret = session->RunGraph(graphId, inputTensors, outputTensors);
	169	+ if (ret != SUCCESS) {
	170	+ printf("%s - ERROR - [GEIR]: RunGraph failed\n", GetTime().c_str());
	171	+ ge::AscendString errMsg = ge::GEGetErrorMsgV2();
	172	+ printf("%s - ERROR - [GEIR]: %s\n", GetTime().c_str(), errMsg.GetString());
	173	+ delete session;
	174	+ ge::GEFinalize();
	175	+ return FAILED;
	176	+ }
	177	+ printf("%s - INFO - [GEIR]: RunGraph success, output_num=%zu\n", GetTime().c_str(), outputTensors.size());
	178	+
	179	+ for (size_t i = 0; i < outputTensors.size(); ++i) {
	180	+ uint8_t* p = outputTensors[i].GetData();
	181	+ if (p == nullptr) {
	182	+ printf("%s - ERROR - [GEIR]: output[%zu] GetData returned null\n", GetTime().c_str(), i);
	183	+ continue;
	184	+ }
	185	+ size_t sz = outputTensors[i].GetSize();
	186	+ int64_t cnt = static_cast<int64_t>(sz / sizeof(float));
	187	+ printf("%s - INFO - [GEIR]: output[%zu] size=%zu (%ld floats)\n", GetTime().c_str(), i, sz, cnt);
	188	+ const float* fp = reinterpret_cast<const float*>(p);
	189	+ for (int64_t e = 0; e < cnt; ++e) {
	190	+ printf(" iou[%ld] = %.6f\n", e, fp[e]);
	191	+ }
	192	+ WriteBin("./iou3d_geir_output_" + std::to_string(i) + ".bin", p, sz);
	atomgit-bot7月13日 🟡 Medium Priority 文件 `objdetect/iou3d/examples/test_geir_iou3d.cpp` 第 179 行（以及 `examples/arch35/test_geir_iou3d.cpp` 同一位置）：`uint8_t* p = outputTensors[i].GetData()` 返回的指针未做空指针检查，直接在第 183 行 `reinterpret_cast<const float*>(p)` 和第 185 行 `fp[e]` 解引用。虽然 RunGraph 成功后通常数据有效，但 GE Tensor 的 `GetData()` 在主机内存未正确分配时可能返回 nullptr。触发条件：输出 Tensor 数量 > 0 但数据未就绪（如 Session 异常状态），则空指针解引用导致测试段错误。作为测试代码，应对 API 返回值做防御性检查。建议：在 GetData() 后增加空指针检查：`if (p == nullptr) { printf("GetData null\n"); continue; }` 不准确？ Y yulianjie 7月21日评论： 7月21日评论：
	193	+ }
	194	+
	195	+ printf("%s - INFO - [GEIR]: GEFinalize\n", GetTime().c_str());
	196	+ delete session;
	197	+ ret = ge::GEFinalize();
	198	+ if (ret != SUCCESS) {
	199	+ printf("%s - ERROR - [GEIR]: GEFinalize failed\n", GetTime().c_str());
	200	+ return FAILED;
	201	+ }
	202	+ printf("%s - INFO - [GEIR]: done\n", GetTime().c_str());
	203	+ return SUCCESS;
	204	+}

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

The file is empty

Aobjdetect/iou3d/op_graph/iou3d_graph_infer.cpp+36-0

@@ -0,0 +1,36 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file iou3d_graph_infer.cpp
	13	+ * \brief Iou3D 算子图模式数据类型推导实现
	14	+ *
	15	+ * dtype_rule:
	16	+ * iou.dtype = bboxes.dtype == float32
	17	+ */
	18	+
	19	+#include "register/op_impl_registry.h"
	20	+#include "exe_graph/runtime/infer_datatype_context.h"
	21	+
	22	+using namespace ge;
	23	+
	24	+namespace ops {
	25	+
	26	+// dtype 推导：iou.dtype = bboxes.dtype
	27	+static ge::graphStatus InferDataType4Iou3D(gert::InferDataTypeContext* context)
	28	+{
	29	+ const auto inputDtype = context->GetInputDataType(0);
	30	+ context->SetOutputDataType(0, inputDtype);
	31	+ return ge::GRAPH_SUCCESS;
	32	+}
	33	+
	34	+IMPL_OP(Iou3D).InferDataType(InferDataType4Iou3D);
	35	+
	36	+} // namespace ops

Aobjdetect/iou3d/op_graph/iou3d_proto.h+48-0

@@ -0,0 +1,48 @@
	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 iou3d_proto.h
	13	+ * \brief Iou3D 图模式（GE IR）算子原型定义
	14	+ */
	15	+#ifndef OPS_OP_PROTO_INC_IOU3D_H_
	16	+#define OPS_OP_PROTO_INC_IOU3D_H_
	17	+
	18	+#include "graph/operator_reg.h"
	19	+#include "graph/types.h"
	20	+
	21	+namespace ge {
	22	+
	23	+/**
	24	+* @brief Calculate the intersection ratio of two rotated cuboids . \n
	25	+
	26	+* @par Inputs:
	27	+* @li bboxes : data of grad increment, a 3D Tensor of type float32 with
	28	+* shape (B, 7, N). "N" indicates the number of boxes, and the value
	29	+* "7" refers to [x, y, z, w, h, d, theta].
	30	+* @li gtboxes: Bounding boxes, a 3D Tensor of type float32 with
	31	+* shape (B, 7, K). "K" indcates the number of boxes, and the value
	32	+* "7" refers to [x, y, z, w, h, d, theta].
	33	+
	34	+* @par Outputs:
	35	+* iou: A 3D Tensor of float32 with shape [B, N, K].
	36	+
	37	+* @attention Constraints:
	38	+* In each batch, the invalid box cannot appear before the valid box.
	39	+*/
	40	+REG_OP(Iou3D)
	41	+ .INPUT(bboxes, TensorType({DT_FLOAT}))
	42	+ .INPUT(gtboxes, TensorType({DT_FLOAT}))
	43	+ .OUTPUT(iou, TensorType({DT_FLOAT}))
	44	+ .OP_END_FACTORY_REG(Iou3D)
	45	+
	46	+} // namespace ge
	47	+
	48	+#endif // OPS_OP_PROTO_INC_IOU3D_H_

Aobjdetect/iou3d/op_host/arch35/iou3d_tiling_arch35.cpp+179-0

@@ -0,0 +1,179 @@

1 +/**

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 iou3d_tiling_arch35.cpp

13 + * \brief Iou3D Host 侧 Tiling（arch35 / DAV_3510）

14 + *

15 + * def 驱动 dtype：dtype 由 _def.cpp 的 DataType({ge::DT_FLOAT}) 声明，构建系统按 dtype

16 + * 展开 kernel 变体；无 TilingKey 参数，空 Tensor 短路在 kernel 内部运行时判断。

17 + * - 多核切分：总 (b,i,j) 对数 = B*N*K，按 coreNum 均分为不相交子集

18 + * - UB 批处理：每核内按 tileLen 分批

19 + * - shape 校验：channel==7、同 batch（D5 对标 mmcv 已移除 K<=2000 上限）

20 + * - 空 Tensor：batch==0 || N==0 || K==0 → isEmpty=1 → kernel 侧运行时短路

21 + * - 极角排序（>3 顶点 Sort32）临时 buffer 由 kernel 侧按固定 32 元素分配，与逻辑规模无关，

22 + * Host 侧无需动态精算（历史遗留的 GetSortTmpSize 调用与 sortTmpSize 字段已移除）。

23 + */

24 +

25 +#include "register/op_def_registry.h"

26 +#include "op_common/log/log.h"

27 +#include "op_common/op_host/util/math_util.h"

28 +#include "op_common/op_host/util/platform_util.h"

29 +#include "../../op_kernel/arch35/iou3d_tiling_data.h"

30 +

31 +namespace optiling {

32 +

33 +using Ops::Base::CeilAlign;

34 +using Ops::Base::CeilDiv;

35 +using Ops::Base::FloorAlign;

36 +using Ops::Base::FloorDiv;

37 +

38 +constexpr uint32_t WS_SYS_SIZE = 0U; // 不使用 GM workspace（顶点/排序缓冲均 UB 内驻留）

39 +constexpr size_t WORKSPACE_NUM = 1;

40 +constexpr int64_t IOU3D_DOF = 7; // 7-DoF 通道

41 +// 单批处理的 (i,j) 对数（UB 批大小）。保守取 256（UB 预算 ~180KB < 248KB）。

42 +constexpr uint32_t IOU3D_TILE_LEN = 256U;

43 +

44 +// 获取平台信息（coreNum, ubSize）。

45 +static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)

46 +{

47 + fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();

48 + OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);

49 + auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);

50 + *coreNum = ascendcPlatform.GetCoreNumAiv();

51 + OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);

52 + ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize);

53 + OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);

54 + return ge::GRAPH_SUCCESS;

55 +}

56 +

57 +// 获取并校验 shape 信息，输出 B/N/K（dtype 校验内置：非 float32 直接报错，不编码进 tiling_key）

58 +static ge::graphStatus GetShapeInfo(gert::TilingContext* context, int64_t* batch, int64_t* numN, int64_t* numK)

59 +{

60 + auto bboxes = context->GetInputShape(0); // [B, 7, N]

61 + OP_CHECK_NULL_WITH_CONTEXT(context, bboxes);

62 + auto gtboxes = context->GetInputShape(1); // [B, 7, K]

63 + OP_CHECK_NULL_WITH_CONTEXT(context, gtboxes);

64 +

65 + const auto& bShape = bboxes->GetStorageShape();

66 + const auto& gShape = gtboxes->GetStorageShape();

67 +

68 + // rank-3 校验

69 + OP_CHECK_IF(

70 + bShape.GetDimNum() != 3 || gShape.GetDimNum() != 3,

71 + OP_LOGE(context, "Iou3D: bboxes/gtboxes must be rank-3, got %zu/%zu", bShape.GetDimNum(), gShape.GetDimNum()),

72 + return ge::GRAPH_FAILED);

73 +

74 + *batch = bShape.GetDim(0);

75 + *numN = bShape.GetDim(2);

76 + *numK = gShape.GetDim(2);

77 +

78 + // channel==7 校验

79 + OP_CHECK_IF(bShape.GetDim(1) != IOU3D_DOF || gShape.GetDim(1) != IOU3D_DOF,

80 + OP_LOGE(context, "Iou3D: channel dim must be 7 (7-DoF), got bboxes=%ld, gtboxes=%ld", bShape.GetDim(1),

A

A andong_hw7月21日

关于shape dtype的校验，建议使用新的日志接口

81 + gShape.GetDim(1)),

82 + return ge::GRAPH_FAILED);

83 + // 同 batch 校验

84 + OP_CHECK_IF(bShape.GetDim(0) != gShape.GetDim(0),

85 + OP_LOGE(context, "Iou3D: bboxes/gtboxes must share batch B, got %ld vs %ld", bShape.GetDim(0),

86 + gShape.GetDim(0)),

87 + return ge::GRAPH_FAILED);

88 + // D5 对标 mmcv：移除 K≤2000 上限（mmcv 无 K 限制）。多核切分按 totalPairs=B*N*K（int64）均分，

89 + // UB 批处理按固定 tileLen(256) 分批，Sort32 固定 32 元素（多边形顶点 <=8），均与 K 无耦合，任意 K 成立。

90 +

91 + // dtype 校验（float32）。dtype 由 def 文件驱动展开 kernel 变体，此处仅运行时友好报错，

92 + // 不再编码进 tiling_key（避免与 def 的 DataType 声明重复编码 dtype 维度）。

93 + auto inputDesc = context->GetInputDesc(0);

94 + OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);

95 + ge::DataType dataType = inputDesc->GetDataType();

96 + OP_CHECK_IF(dataType != ge::DT_FLOAT,

97 + OP_LOGE(context, "Iou3D: only float32 supported, got dtype=%d", static_cast<int>(dataType)),

98 + return ge::GRAPH_FAILED);

99 +

100 + return ge::GRAPH_SUCCESS;

101 +}

102 +

103 +// 声明 workspace slot（即使为 0 也必须声明，否则运行时不分配 workspace 指针）

104 +static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)

105 +{

106 + size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM);

107 + OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);

108 + currentWorkspace[0] = WS_SYS_SIZE;

109 + return ge::GRAPH_SUCCESS;

110 +}

111 +

112 +static ge::graphStatus Iou3DTilingFunc(gert::TilingContext* context)

113 +{

114 + // 1、平台信息（coreNum, ubSize）

115 + uint64_t ubSize;

116 + int64_t coreNum;

117 + OP_CHECK_IF(GetPlatformInfo(context, &ubSize, &coreNum) != ge::GRAPH_SUCCESS,

118 + OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);

119 +

120 + // 2、shape 信息 + 校验

121 + int64_t batch, numN, numK;

122 + OP_CHECK_IF(GetShapeInfo(context, &batch, &numN, &numK) != ge::GRAPH_SUCCESS,

123 + OP_LOGE(context, "GetShapeInfo error"), return ge::GRAPH_FAILED);

124 +

125 + // 3、workspace slot

126 + OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),

127 + return ge::GRAPH_FAILED);

128 +

129 + // 4、填 TilingData

130 + Iou3DTilingData* tiling = context->GetTilingData<Iou3DTilingData>();

131 + OP_CHECK_NULL_WITH_CONTEXT(context, tiling);

132 + OP_CHECK_IF(memset_s(tiling, sizeof(Iou3DTilingData), 0, sizeof(Iou3DTilingData)) != EOK,

133 + OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);

134 +

135 + tiling->batch = static_cast<uint32_t>(batch);

136 + tiling->numBboxes = static_cast<uint32_t>(numN);

137 + tiling->numGtboxes = static_cast<uint32_t>(numK);

138 +

139 + const uint32_t isEmpty = (batch == 0 || numN == 0 || numK == 0) ? 1U : 0U;

140 + tiling->isEmpty = isEmpty;

141 +

142 + // 空 Tensor（batch==0 || N==0 || K==0）→ 短路：输出为空矩阵，无 (b,i,j) 对，

143 + // 单核占位（block=1），kernel 侧运行时判断 isEmpty 直接返回。

144 + // 注：batch==0 时 totalPairs==0 → pairsPerCore==0 → usedCoreNum==0，若不短路会 SetBlockDim(0) 非法。

145 + if (isEmpty != 0U) {

146 + context->SetBlockDim(1);

147 + return ge::GRAPH_SUCCESS;

148 + }

149 +

150 + // 多核切分：总 (b,i,j) 对数按核均分，每核负责不相交子集

151 + const int64_t totalPairs = batch * numN * numK;

152 + const int64_t pairsPerCore = CeilDiv(totalPairs, coreNum);

153 + const int64_t usedCoreNum = CeilDiv(totalPairs, pairsPerCore);

atomgit-bot7月13日

🟡 Medium Priority

Tiling 函数中 isEmpty 仅检查 numN == 0 || numK == 0，未覆盖 batch == 0。

触发路径：当 batch=0 但 numN>0 且 numK>0 时（例如 shape [0,7,N]/[0,7,K]），isEmpty=0，进入多核切分：

即使 CeilDiv 内部有保护返回 0，后续 SetBlockDim(0) 也将传入非法核数，导致运行时错误。batch=0 的输入张量在动态批处理场景确有可能出现，应纳入空 Tensor 短路。

建议：在 isEmpty 判断中补充 batch==0 条件：const uint32_t isEmpty = (batch == 0 || numN == 0 || numK == 0) ? 1U : 0U;，或直接在获取 shape 后判断 totalPairs == 0 则短路。两种方式等价，选择后者更直观地保护后续除法。

改动建议

-    const int64_t usedCoreNum = CeilDiv(totalPairs, pairsPerCore);
+    const uint32_t isEmpty = (numN == 0 || numK == 0) ? 1U : 0U;
+    tiling->isEmpty = isEmpty;
+    // 空 Tensor（batch==0 || N==0 || K==0）→ TPL_EMPTY 短路
+    if (isEmpty != 0U || totalPairs == 0) {

应用建议

不准确？

Y

yulianjie

7月21日评论：

154 +

155 + tiling->coreNum = static_cast<uint32_t>(usedCoreNum);

156 + tiling->pairsPerCore = static_cast<uint32_t>(pairsPerCore);

157 +

158 + // UB 批处理粒度：每核内按 tileLen 分批（不超过本核 pairsPerCore）

159 + const uint32_t tileLen = (pairsPerCore < static_cast<int64_t>(IOU3D_TILE_LEN)) ?

160 + static_cast<uint32_t>(pairsPerCore) :

161 + IOU3D_TILE_LEN;

162 + tiling->tileLen = tileLen;

163 + tiling->tailLen = (tileLen == 0U) ? 0U : static_cast<uint32_t>(pairsPerCore % static_cast<int64_t>(tileLen));

164 +

165 + context->SetBlockDim(static_cast<uint32_t>(usedCoreNum));

166 +

167 + return ge::GRAPH_SUCCESS;

168 +}

169 +

170 +static ge::graphStatus TilingParseForIou3D([[maybe_unused]] gert::TilingParseContext* context)

171 +{

172 + return ge::GRAPH_SUCCESS;

173 +}

174 +

175 +struct Iou3DCompileInfo {}; // 必须定义，入图场景依赖

176 +

177 +IMPL_OP_OPTILING(Iou3D).Tiling(Iou3DTilingFunc).TilingParse<Iou3DCompileInfo>(TilingParseForIou3D);

178 +

179 +} // namespace optiling

Aobjdetect/iou3d/op_host/config/ascend950/iou3d_binary.json+40-0

@@ -0,0 +1,40 @@
	1	+{
	2	+ "op_type": "Iou3D",
	3	+ "op_list": [
	4	+ {
	5	+ "bin_filename": "Iou3D_float32",
	6	+ "inputs": [
	7	+ {
	8	+ "name": "bboxes",
	9	+ "index": 0,
	10	+ "dtype": "float32",
	11	+ "format": "ND",
	12	+ "paramType": "required",
	13	+ "shape": [-2],
	14	+ "format_match_mode": "FormatAgnostic"
	15	+ },
	16	+ {
	17	+ "name": "gtboxes",
	18	+ "index": 1,
	19	+ "dtype": "float32",
	20	+ "format": "ND",
	21	+ "paramType": "required",
	22	+ "shape": [-2],
	23	+ "format_match_mode": "FormatAgnostic"
	24	+ }
	25	+ ],
	26	+ "outputs": [
	27	+ {
	28	+ "name": "iou",
	29	+ "index": 0,
	30	+ "dtype": "float32",
	31	+ "format": "ND",
	32	+ "paramType": "required",
	33	+ "shape": [-2],
	34	+ "format_match_mode": "FormatAgnostic"
	35	+ }
	36	+ ],
	37	+ "attrs": []
	38	+ }
	39	+ ]
	40	+}

Aobjdetect/iou3d/op_host/config/ascend950/iou3d_simplified_key.ini+13-0

@@ -0,0 +1,13 @@
	1	+; 该文件主要影响 opc 工具编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
	2	+; [某算子]
	3	+; default=xx
	4	+; ascendxx=xx
	5	+; 其中，default为默认mode，ascnedxx为可选mode，如果不同芯片有差异化要求时，需要配置；
	6	+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
	7	+; 2)如果仅有default配置：各个版本按default配置
	8	+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
	9	+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
	10	+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
	11	+; 6)是否是AscendC算子，由 ops/built-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
	12	+[Iou3D]
	13	+default=0

Aobjdetect/iou3d/op_host/iou3d_def.cpp+60-0

@@ -0,0 +1,60 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file iou3d_def.cpp
	13	+ * \brief Iou3D 算子定义，声明输入输出和算子配置
	14	+ *
	15	+ * 3D 旋转框 IoU：BEV 旋转矩形交集面积 × Z 轴重叠高度 / 并集体积。
	16	+ * inputs : bboxes [B, 7, N] (float32) —— 预测框（7-DoF: x,y,z,w,h,d,theta）
	17	+ * gtboxes [B, 7, K] (float32) —— 真值框（K <= 2000）
	18	+ * output : iou [B, N, K] (float32)
	19	+ * 目标芯片：Ascend950PR / Ascend950DT（arch35 / DAV_3510）。
	20	+ */
	21	+#include "register/op_def_registry.h"
	22	+
	23	+namespace ops {
	24	+class Iou3D : public OpDef {
	25	+public:
	26	+ explicit Iou3D(const char* name) : OpDef(name)
	27	+ {
	28	+ this->Input("bboxes") // 预测框 [B, 7, N]
	29	+ .ParamType(REQUIRED)
	30	+ .DataType({ge::DT_FLOAT})
	31	+ .Format({ge::FORMAT_ND})
	32	+ .UnknownShapeFormat({ge::FORMAT_ND})
	33	+ .AutoContiguous();
	34	+ this->Input("gtboxes") // 真值框 [B, 7, K]
	35	+ .ParamType(REQUIRED)
	36	+ .DataType({ge::DT_FLOAT})
	37	+ .Format({ge::FORMAT_ND})
	38	+ .UnknownShapeFormat({ge::FORMAT_ND})
	39	+ .AutoContiguous();
	40	+ this->Output("iou") // 输出 [B, N, K]
	41	+ .ParamType(REQUIRED)
	42	+ .DataType({ge::DT_FLOAT})
	43	+ .Format({ge::FORMAT_ND})
	44	+ .UnknownShapeFormat({ge::FORMAT_ND})
	45	+ .AutoContiguous();
	46	+
	47	+ // 目标平台：仅 Ascend950（arch35 / DAV_3510）
	48	+ OpAICoreConfig aiCoreConfig;
	49	+ aiCoreConfig.DynamicCompileStaticFlag(true)
	50	+ .DynamicFormatFlag(false)
	51	+ .DynamicRankSupportFlag(true)
	52	+ .DynamicShapeSupportFlag(true)
	53	+ .NeedCheckSupportFlag(false)
	54	+ .PrecisionReduceFlag(true)
	55	+ .ExtendCfgInfo("opFile.value", "iou3d");
	56	+ this->AICore().AddConfig("ascend950", aiCoreConfig);
	57	+ }
	58	+};
	59	+OP_ADD(Iou3D); // 添加算子信息库
	60	+} // namespace ops

Aobjdetect/iou3d/op_host/iou3d_infershape.cpp+82-0

@@ -0,0 +1,82 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file iou3d_infershape.cpp
	13	+ * \brief Iou3D 算子形状推导实现
	14	+ *
	15	+ * shape_rule:
	16	+ * iou.shape = (bboxes.shape[0], bboxes.shape[2], gtboxes.shape[2]) == (B, N, K)
	17	+ */
	18	+
	19	+#include "register/op_impl_registry.h"
	20	+#include "exe_graph/runtime/infer_shape_context.h"
	21	+#include "op_common/log/log.h"
	22	+
	23	+using namespace ge;
	24	+
	25	+namespace ops {
	26	+
	27	+// iou3d 输出形状推导：iou[B, N, K]
	28	+static ge::graphStatus InferShape4Iou3D(gert::InferShapeContext* context)
	29	+{
	30	+ const gert::Shape* bboxesShape = context->GetInputShape(0); // [B, 7, N]
	31	+ OP_CHECK_NULL_WITH_CONTEXT(context, bboxesShape);
	32	+ const gert::Shape* gtboxesShape = context->GetInputShape(1); // [B, 7, K]
	33	+ OP_CHECK_NULL_WITH_CONTEXT(context, gtboxesShape);
	34	+
	35	+ gert::Shape* iouShape = context->GetOutputShape(0); // [B, N, K]
	36	+ OP_CHECK_NULL_WITH_CONTEXT(context, iouShape);
	A A andong_hw7月21日没有支持 unknownRank 场景 Y yulianjie 7月25日评论： 7月25日评论：
	37	+
	38	+ // 处理 unknownRank 场景：输入 rank 未知时，输出也设为 unknownRank
	39	+ if (bboxesShape->GetDimNum() == ge::UNKNOWN_RANK.size() \|\| gtboxesShape->GetDimNum() == ge::UNKNOWN_RANK.size()) {
	40	+ iouShape->SetDimNum(ge::UNKNOWN_RANK.size());
	41	+ return ge::GRAPH_SUCCESS;
	42	+ }
	43	+
	44	+ // 输入必须为 rank-3（[B, 7, N] / [B, 7, K]）
	45	+ OP_CHECK_IF(bboxesShape->GetDimNum() != 3 \|\| gtboxesShape->GetDimNum() != 3,
	46	+ OP_LOGE(context, "Iou3D: bboxes/gtboxes must be rank-3, got %zu/%zu", bboxesShape->GetDimNum(),
	47	+ gtboxesShape->GetDimNum()),
	48	+ return ge::GRAPH_FAILED);
	49	+
	50	+ // channel 维度必须为 7（7-DoF）
	51	+ const int64_t bboxesChannel = bboxesShape->GetDim(1);
	52	+ const int64_t gtboxesChannel = gtboxesShape->GetDim(1);
	53	+ OP_CHECK_IF(bboxesChannel != 7 \|\| gtboxesChannel != 7,
	54	+ OP_LOGE(context, "Iou3D: channel dimension must be 7, got bboxes=%ld, gtboxes=%ld", bboxesChannel,
	55	+ gtboxesChannel),
	56	+ return ge::GRAPH_FAILED);
	57	+
	58	+ const int64_t batch = bboxesShape->GetDim(0); // B
	59	+ const int64_t numN = bboxesShape->GetDim(2); // N
	60	+ const int64_t numK = gtboxesShape->GetDim(2); // K
	61	+
	62	+ // batch 一致性校验
	63	+ const int64_t gtboxesBatch = gtboxesShape->GetDim(0);
	64	+ OP_CHECK_IF(
	65	+ batch != gtboxesBatch,
	66	+ OP_LOGE(context, "Iou3D: batch dimension must be consistent, got bboxes=%ld, gtboxes=%ld", batch, gtboxesBatch),
	67	+ return ge::GRAPH_FAILED);
	68	+
	69	+ // D5 对标 mmcv：移除 K≤2000 上限（mmcv 无 K 限制）。逐对计算的 UB/Sort32 缓冲仅按
	70	+ // tileLen 与固定 32 元素多边形排序规模分配，与 K 无耦合，任意 K 成立。
	71	+
	72	+ iouShape->SetDimNum(3);
	73	+ iouShape->SetDim(0, batch);
	74	+ iouShape->SetDim(1, numN);
	75	+ iouShape->SetDim(2, numK);
	atomgit-bot7月13日 🟡 Medium Priority `iou3d_infershape.cpp` 的 `InferShape4Iou3D` 函数仅校验了输入的 rank 是否为 3（第 41 行），但未校验以下关键约束： channel 维度：`bboxesShape->GetDim(1)` 与 `gtboxesShape->GetDim(1)` 必须等于 7（7-DoF），否则后续计算会将非 7 通道的数据当 7-DoF 解析，产生错误结果。 batch 一致性：`bboxesShape->GetDim(0)` 必须等于 `gtboxesShape->GetDim(0)`（第 46 行直接取 bboxes 的 dim(0) 作为输出 batch，若两个输入的 batch 不一致，会静默使用错误的 batch 值推导输出 shape）。 K 上限：`gtboxesShape->GetDim(2)` 不得大于 2000（proto 声明与 README 均明确约束）。 Tiling 阶段（`iou3d_tiling_arch35.cpp` 第 88-104 行）虽然会完整校验以上所有约束，但 GE 图编译阶段 InferShape 先于 Tiling 执行：若 InferShape 不校验，框架会基于错误 shape 分配输出 Tensor 内存后才在 Tiling 阶段失败，造成内存浪费且错误信息延迟。同类算子（如 `anchor_response_flags_infershape.cpp` 第 66 行）均在 InferShape 中校验了特定维度值，本项目应遵循相同规范。触发条件：输入 bboxes/gtboxes 的 channel ≠ 7、batch 不一致、或 K > 2000 时触发。失败模式：InferShape 推导出错误输出 shape → 框架据此分配错误大小的输出 Tensor → Tiling 阶段检测失败（但浪费已分配内存）。建议：在第 46 行（获取 batch/numN/numK）之前，增加 channel==7、batch 一致性和 K≤2000 的校验，参考 tiling 侧 `GetShapeInfo` 函数（`iou3d_tiling_arch35.cpp` 第 88-104 行）的校验逻辑，在 InferShape 中提前拒绝非法 shape。不准确？ Y yulianjie 7月16日评论： 7月16日评论： Y yulianjie 7月25日评论： 7月25日评论：
	76	+
	77	+ return ge::GRAPH_SUCCESS;
	78	+}
	79	+
	80	+IMPL_OP_INFERSHAPE(Iou3D).InferShape(InferShape4Iou3D);
	81	+
	82	+} // namespace ops

Aobjdetect/iou3d/op_kernel/arch35/iou3d.h+647-0

@@ -0,0 +1,647 @@
	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 iou3d.h
	13	+ * \brief Iou3D 主模板 kernel（TPL_MAIN，arch35 / DAV_3510）
	14	+ *
	15	+ * def 驱动 dtype：dtype 由 _def.cpp 的 DataType({ge::DT_FLOAT}) 声明，构建系统通过
	16	+ * -DDTYPE_BBOXES 等编译宏注入；bboxes/gtboxes/iou 均 float32，类内统一用 DTYPE_BBOXES。
	17	+ *
	18	+ * 逐对 (b,i,j) 3D 旋转框 IoU 完整路径：
	19	+ * BEV 旋转顶点投影(向量 Sin/Cos) → 交集顶点收集(包含测试+边相交) →
	20	+ * 多边形面积(0/1/2 顶点置 0, 3 顶点直算三角形, >3 顶点质心分解 + 硬件 Sort32 极角排序 + 三角形叉积和) →
	21	+ * Z 轴重叠 clamp 非负 → epsilon 除法 → 普通 DataCopyPad 写回（(b,i,j) 单点不相交写，无需原子加）。
	22	+ *
	23	+ * 编程模型：RegBase(arch35) kernel-shell + Scalar 混合。
	24	+ * - 向量批算：AscendC::Sin / AscendC::Cos（4/pair 顶点 theta，adv_api math）。
	25	+ * - 硬件排序：>3 顶点走 AscendC::Sort32 + Extract（diamond-angle 极角键，替代标量选择排序）。
	26	+ * - 标量几何：分支密集（corners_num∈{0..8}），逐对用 LocalTensor::GetValue 读标量控制。
	27	+ * - 数值稳定三守卫：
	28	+ * epsilon_guard_division : iou = interVol / max(union, 1e-6)（对标 mmcv clamp，golden.py:377）
	29	+ * clamp_z_overlap_nonneg : real_d = max(min-max, 0)
	30	+ * degenerate_polygon_guard : corners_num < 3 → area = 0
	31	+ * - 近同旋转框 bowtie 半面积退化修复（fp32 特有）：
	32	+ * 幅值相对强去重 : CollectCorners 去重容差 tol=max(1e-6, 1e-4·max(1,\|x\|,\|y\|))，
	33	+ * 坍缩 fp32 向量 Sin/Cos 逐 lane 微差产生的近重复顶点（m=6/8→4），
	34	+ * 相对幅值以兼容大坐标场景（不误并真顶点）。
	35	+ * signed shoelace: SortPolygonArea 面积用有符号鞋带（不 per-triangle 取 abs，末尾一次 abs），
	36	+ * 对残留错序/自交鲁棒（错序正负抵消，不放大成半面积）。
	37	+ * 红线 : 面积层严禁任何 NaN 守卫；NaN/inf 按 IEEE754 自然传播。
	38	+ *
	39	+ * 布局：bboxes[B,7,N] 索引(b,c,i)=b7N + cN + i；gtboxes[B,7,K] 同理；iou[B,N,K] 索引(b,i,j)=bNK + iK + j。
	40	+ * 每核负责 flatten (b,i,j) 的一段连续区间（不相交），每元素只被一个核写一次，输出连续段用普通 DataCopyPad 写回
	41	+ * （单点不相交写 → 无需原子加，也不依赖外部对输出 GM 零初始化）。
	42	+ */
	43	+#ifndef IOU3D_ARCH35_H
	44	+#define IOU3D_ARCH35_H
	45	+
	46	+#include "kernel_operator.h"
	47	+#include "lib/math/sin.h"
	48	+#include "lib/math/cos.h"
	49	+#include "iou3d_tiling_data.h"
	50	+
	51	+namespace NsIou3D {
	52	+
	53	+using namespace AscendC;
	54	+
	55	+constexpr uint32_t IOU3D_DOF = 7; // 7-DoF 通道数
	56	+constexpr uint32_t IOU3D_CORNERS = 4; // 每框 BEV 顶点数
	57	+constexpr uint32_t IOU3D_MAX_INTER = 8; // 交集多边形顶点数上限（两凸四边形交集 <= 8）
	58	+constexpr float IOU3D_EPSILON = 1e-6f; // IoU 分母 clamp 下限（对标 mmcv golden.py:37 EPS_IOU=1e-6）
	59	+// 对标 mmcv：边相交退化判据阈值。\|s5-s1\|>EPS_KERNEL 用叉积公式，否则退化一般式直线方程
	60	+// （对标 mmcv golden.py:35/153 EPS_KERNEL=1e-8）。
	61	+constexpr float IOU3D_EPS_KERNEL = 1e-8f;
	62	+// 对标 mmcv：顶点包含测试绝对容差。反向旋转到框局部系后判 \|rot\|<half+MARGIN
	63	+// （对标 mmcv golden.py:36/194-195 MARGIN=1e-2，严格 <）。
	64	+constexpr float IOU3D_MARGIN = 1e-2f;
	65	+constexpr float IOU3D_DEDUP_TOL = 1e-6f; // 顶点去重绝对容差下限（小坐标场景，D6 保留现状）
	66	+// 幅值相对强去重容差（D6 保留现状）。fp32 向量 Sin/Cos 逐 lane 微差使「同一角点」在 bbox/gtbox 两通道间
	67	+// 偏差 ~1e-7~1e-5·幅值，叠加近共线边额外交点，使 identical box 去重后仍残留 m∈{6,8} 近重复
	68	+// 顶点 → Sort32 tie-break bowtie → 半面积 → IoU=1/3。用相对幅值容差 tol=max(ABS, REL·max(1,\|x\|,\|y\|))
	69	+// 把近重复坍缩回 4 真角点；相对而非绝对以兼容大坐标（1e20 场景不被误并——真顶点间距 >> REL·幅值）。
	70	+constexpr float IOU3D_DEDUP_REL = 1e-4f; // 顶点去重相对幅值容差（D6 保留现状）
	71	+// fp32 32B(one data block) = 8 元素。RegBase(arch35) 向量 adv_api Sin/Cos 要求 src/dst 32B 对齐
	72	+// （sin_3510_impl 用 Reg::StoreAlign<..,DIST_PACK_B32>），故所有参与向量计算的 UB 子段起址与长度按 8 对齐。
	73	+constexpr uint32_t IOU3D_ALIGN_ELEM = 8; // fp32 一个 datablock 的元素数（32B / 4B）
	74	+// 硬件极角排序（Sort32）常量。
	75	+constexpr uint32_t IOU3D_SORT32_LEN = 32U; // Sort32 一趟固定处理 32 元素（8 有效 + 24 padding 沉底）
	76	+constexpr float IOU3D_NEG_INF_KEY = -1.0e30f; // padding 槽键值（沉底到排序末尾）
	77	+constexpr uint32_t IOU3D_IDX_MASK = 0x07FFFFFFu; // Sort32 index 位宽 27bit，回读时按位与还原原始下标
	78	+
	79	+// 向上取整对齐到 IOU3D_ALIGN_ELEM 的倍数（32B 对齐）
	80	+__aicore__ inline uint32_t Iou3DCeilAlign(uint32_t n)
	81	+{
	82	+ return ((n + IOU3D_ALIGN_ELEM - 1U) / IOU3D_ALIGN_ELEM) * IOU3D_ALIGN_ELEM;
	83	+}
	84	+
	85	+class Iou3D {
	86	+public:
	87	+ __aicore__ inline Iou3D(){};
	88	+
	89	+ __aicore__ inline void Init(GM_ADDR bboxes, GM_ADDR gtboxes, GM_ADDR iou, const Iou3DTilingData* tilingData);
	90	+ __aicore__ inline void Process();
	91	+
	92	+private:
	93	+ // 逐对计算一个 (b,i,j) 的 IoU（标量几何，theta 的 sin/cos 由 tile 向量批算后传入）
	94	+ __aicore__ inline float ComputePairIou(int64_t b, int64_t i, int64_t j, float sin1, float cos1, float sin2,
	95	+ float cos2);
	96	+ // BEV 旋转 4 顶点（rect: x,y,z,w,h,d；sinT/cosT 为该框 theta 的 sin/cos）
	97	+ __aicore__ inline void BoxCorners(float x, float y, float w, float h, float sinT, float cosT,
	98	+ float cx[IOU3D_CORNERS], float cy[IOU3D_CORNERS]);
	99	+ // 顶点包含测试（对标 mmcv 反向旋转法 golden.py:171-195）：把点旋转 -heading 到框局部系，
	100	+ // 判 \|rot_x\|<halfDx+MARGIN && \|rot_y\|<halfDy+MARGIN（MARGIN=1e-2 绝对，严格 <）。
	101	+ __aicore__ inline bool PointInRect(float px, float py, float cx, float cy, float halfDx, float halfDy, float sinT,
	102	+ float cosT);
	103	+ // 线段相交（对标 mmcv check_rect_cross + 跨立实验 + 直线求交 golden.py:104-168），相交返回 true 并写 (ox, oy)
	104	+ __aicore__ inline bool SegIntersect(float a1x, float a1y, float a2x, float a2y, float b1x, float b1y, float b2x,
	105	+ float b2y, float& ox, float& oy);
	106	+ // 收集交集顶点（含去重），返回顶点数。box1/box2 各传 (cx,cy,halfDx,halfDy,sinT,cosT) 供反向旋转包含测试。
	107	+ __aicore__ inline uint32_t CollectCorners(const float c1x[IOU3D_CORNERS], const float c1y[IOU3D_CORNERS],
	108	+ const float c2x[IOU3D_CORNERS], const float c2y[IOU3D_CORNERS],
	109	+ float box1cx, float box1cy, float box1hx, float box1hy, float box1sin,
	110	+ float box1cos, float box2cx, float box2cy, float box2hx, float box2hy,
	111	+ float box2sin, float box2cos, float px[IOU3D_MAX_INTER],
	112	+ float py[IOU3D_MAX_INTER]);
	113	+ // 三角形面积（叉积法）
	114	+ __aicore__ inline float TriArea(float x1, float y1, float x2, float y2, float x3, float y3);
	115	+ // 多边形面积（0/1/2→0 守卫；3→三角形直算；>3→质心分解 + 硬件 Sort32 极角排序 + 三角形叉积和）
	116	+ __aicore__ inline float PolygonArea(float px[IOU3D_MAX_INTER], float py[IOU3D_MAX_INTER], uint32_t m);
	117	+ // >3 顶点：硬件 Sort32 极角排序（diamond-angle 键）+ 质心分解叉积面积。
	118	+ __aicore__ inline float SortPolygonArea(float px[IOU3D_MAX_INTER], float py[IOU3D_MAX_INTER], uint32_t m);
	119	+ // 读取一个框的 7-DoF 标量（layout [B,DOF,D]）
	120	+ __aicore__ inline void LoadBox(const GlobalTensor<DTYPE_BBOXES>& gm, int64_t b, int64_t idx, int64_t dimSize,
	121	+ float box[IOU3D_DOF]);
	122	+
	123	+ __aicore__ inline float ScalarAbs(float v) { return v < 0.0f ? -v : v; }
	124	+ __aicore__ inline float ScalarMax(float a, float b) { return a > b ? a : b; }
	125	+ __aicore__ inline float ScalarMin(float a, float b) { return a < b ? a : b; }
	126	+
	127	+private:
	128	+ TPipe pipe;
	129	+ // theta 向量批算 sin/cos 缓冲（bbox 与 gtbox 各一路）
	130	+ TBuf<TPosition::VECCALC> angleBuf; // 存 bbox_theta \| gtbox_theta \| sin/cos 中间（各段 32B 对齐）
	131	+ TBuf<TPosition::VECCALC> tmpBuf; // Sin/Cos adv_api 显式 sharedTmpBuffer（避免 PopStackBuffer 依赖）
	132	+ TBuf<TPosition::VECOUT> outBuf; // iou tile 输出缓冲（32B 对齐）
	133	+ // 硬件极角排序（Sort32）缓冲（>3 顶点分支复用，Init 一次分配）：
	134	+ TBuf<TPosition::VECCALC> sortKeyBuf; // 32 极角键（float）
	135	+ TBuf<TPosition::VECCALC> sortIdxBuf; // 32 原始下标（uint32）
	136	+ TBuf<TPosition::VECCALC> sortedBuf; // Sort32 输出 proposal 交织(value,index)：32*2 float
	137	+ TBuf<TPosition::VECCALC> sortOutValBuf; // Extract 输出排序 value（float）
	138	+ TBuf<TPosition::VECCALC> sortOutIdxBuf; // Extract 输出原始 index（uint32）
	139	+
	140	+ GlobalTensor<DTYPE_BBOXES> bboxesGM;
	141	+ GlobalTensor<DTYPE_BBOXES> gtboxesGM;
	142	+ GlobalTensor<DTYPE_BBOXES> iouGM;
	143	+
	144	+ uint32_t batch_ = 0;
	145	+ uint32_t numN_ = 0;
	146	+ uint32_t numK_ = 0;
	147	+ uint32_t tileLen_ = 0;
	148	+ uint32_t alignedTl_ = 0; // 32B 对齐后的单批粒度（各 UB 子段步长）
	149	+ uint32_t isEmpty_ = 0; // 空 Tensor 标志（从 tilingData 读取，用于运行时判断）
	150	+ int64_t pairStart_ = 0; // 本核 flatten (b,i,j) 起始
	151	+ int64_t pairCount_ = 0; // 本核 flatten 对数
	152	+ int64_t totalPairs_ = 0;
	153	+};
	154	+
	155	+// ---------------------------------------------------------------------------
	156	+// Init
	157	+// ---------------------------------------------------------------------------
	158	+
	159	+__aicore__ inline void Iou3D::Init(GM_ADDR bboxes, GM_ADDR gtboxes, GM_ADDR iou, const Iou3DTilingData* tilingData)
	160	+{
	161	+ batch_ = tilingData->batch;
	162	+ numN_ = tilingData->numBboxes;
	163	+ numK_ = tilingData->numGtboxes;
	164	+ tileLen_ = tilingData->tileLen;
	165	+ isEmpty_ = tilingData->isEmpty; // 保存空 Tensor 标志
	166	+
	167	+ totalPairs_ = static_cast<int64_t>(batch_) * numN_ * numK_;
	168	+ const int64_t pairsPerCore = tilingData->pairsPerCore;
	169	+ const int64_t blockIdx = GetBlockIdx();
	170	+ pairStart_ = pairsPerCore * blockIdx;
	atomgit-bot7月13日 🟡 Medium Priority `iou3d.h:181` 行：`pairStart_ = pairsPerCore * blockIdx;`。`pairsPerCore` 和 `blockIdx`（`GetBlockIdx()` 返回值）均为 `uint32_t`，乘法在 32 位域完成后再赋值给 `int64_t pairStart_`，形成典型的「32 位乘法链赋 64 位」溢出模式。触发条件：当 `totalPairs = B × N × K` 较大（如 > 4.3B）时，`pairsPerCore = CeilDiv(totalPairs, coreNum)` 可达数千万甚至上亿，`pairsPerCore × blockIdx` 可能超过 `UINT32_MAX(≈4.29B)`，在 32 位域溢出截断后得到一个错误的小值 → 本核起始偏移错乱 → 多个核写到重叠的 GM 区域或完全不写，输出数据错乱且难以排查。此模式命中 Ascend C 专项规则 1（乘法链转换滞后），属于跨 group 累加类高风险溢出（不适用 NPU 64GB 降级保护）。建议：将乘法链首操作数显式提升为 int64_t：`pairStart_ = static_cast<int64_t>(pairsPerCore) * static_cast<int64_t>(blockIdx);`。这确保乘法在 64 位域完成。不准确？ Y yulianjie 7月16日评论： 7月16日评论：
	171	+ int64_t remain = totalPairs_ - pairStart_;
	172	+ pairCount_ = (remain > pairsPerCore) ? pairsPerCore : (remain > 0 ? remain : 0);
	173	+
	174	+ bboxesGM.SetGlobalBuffer((__gm__ DTYPE_BBOXES*)bboxes);
	175	+ gtboxesGM.SetGlobalBuffer((__gm__ DTYPE_BBOXES*)gtboxes);
	176	+ iouGM.SetGlobalBuffer((__gm__ DTYPE_BBOXES*)iou);
	177	+
	178	+ // 空 Tensor 运行时短路：无 (b,i,j) 对，不分配任何 UB（Process 直接短路返回）
	179	+ if (isEmpty_ != 0U) {
	180	+ return;
	181	+ }
	182	+
	183	+ // theta / sin / cos 中间：6 段，每段步长 = 32B 对齐后的 alignedTl（保证每个子段起址 32B 对齐）
	184	+ // [0]=theta1 [1]=theta2 [2]=sin1 [3]=cos1 [4]=sin2 [5]=cos2
	185	+ // RegBase(arch35) Sin/Cos 要求 src/dst 32B 对齐；1 元素(4B) 非对齐子张量会触发 AIV 向量核异常(err 340)。
	186	+ uint32_t tl = (tileLen_ == 0U) ? 1U : tileLen_;
	187	+ alignedTl_ = Iou3DCeilAlign(tl);
	188	+ pipe.InitBuffer(angleBuf, 6U * alignedTl_ * sizeof(float));
	189	+ // Sin/Cos fp32 主路径不消耗 sharedTmpBuffer，但显式提供 32B 对齐的 tmp 以规避 PopStackBuffer 栈依赖。
	190	+ pipe.InitBuffer(tmpBuf, alignedTl_ * sizeof(float));
	191	+ pipe.InitBuffer(outBuf, alignedTl_ * sizeof(float));
	192	+
	193	+ // 硬件极角排序（Sort32）缓冲：一次分配、跨对复用（>3 顶点分支）。
	194	+ // proposal 交织格式 (value,index)，故 sortedBuf 需 32*2 个 float。固定 32 元素布局。
	195	+ pipe.InitBuffer(sortKeyBuf, IOU3D_SORT32_LEN * sizeof(float));
	196	+ pipe.InitBuffer(sortIdxBuf, IOU3D_SORT32_LEN * sizeof(uint32_t));
	197	+ pipe.InitBuffer(sortedBuf, IOU3D_SORT32_LEN * 2U * sizeof(float));
	198	+ pipe.InitBuffer(sortOutValBuf, IOU3D_SORT32_LEN * sizeof(float));
	199	+ pipe.InitBuffer(sortOutIdxBuf, IOU3D_SORT32_LEN * sizeof(uint32_t));
	200	+}
	201	+
	202	+// ---------------------------------------------------------------------------
	203	+// LoadBox：读取一个框的 7-DoF（layout [B, DOF, dimSize]，通道 stride = dimSize）
	204	+// ---------------------------------------------------------------------------
	205	+
	206	+__aicore__ inline void Iou3D::LoadBox(const GlobalTensor<DTYPE_BBOXES>& gm, int64_t b, int64_t idx, int64_t dimSize,
	207	+ float box[IOU3D_DOF])
	208	+{
	209	+ int64_t base = b * static_cast<int64_t>(IOU3D_DOF) * dimSize + idx;
	210	+ for (uint32_t c = 0; c < IOU3D_DOF; ++c) {
	211	+ box[c] = static_cast<float>(gm.GetValue(base + static_cast<int64_t>(c) * dimSize));
	212	+ }
	213	+}
	214	+
	215	+// ---------------------------------------------------------------------------
	216	+// BoxCorners（与 golden BoxCorners 顺序一致：P1 左上, P2 右上, P3 右下, P4 左下）
	217	+// ---------------------------------------------------------------------------
	218	+
	219	+__aicore__ inline void Iou3D::BoxCorners(float x, float y, float w, float h, float sinT, float cosT,
	220	+ float cx[IOU3D_CORNERS], float cy[IOU3D_CORNERS])
	221	+{
	222	+ float halfW = 0.5f * w;
	223	+ float halfH = 0.5f * h;
	224	+ float hwCos = halfW * cosT;
	225	+ float hwSin = halfW * sinT;
	226	+ float hhCos = halfH * cosT;
	227	+ float hhSin = halfH * sinT;
	228	+
	229	+ float xSubW = x - hwCos;
	230	+ float ySubW = y - hwSin;
	231	+ float xAddW = x + hwCos;
	232	+ float yAddW = y + hwSin;
	233	+
	234	+ cx[0] = xSubW - hhSin;
	235	+ cy[0] = ySubW + hhCos; // P1 左上
	236	+ cx[1] = xAddW - hhSin;
	237	+ cy[1] = yAddW + hhCos; // P2 右上
	238	+ cx[2] = xAddW + hhSin;
	239	+ cy[2] = yAddW - hhCos; // P3 右下
	240	+ cx[3] = xSubW + hhSin;
	241	+ cy[3] = ySubW - hhCos; // P4 左下
	242	+}
	243	+
	244	+// ---------------------------------------------------------------------------
	245	+// PointInRect（对标 mmcv 反向旋转法 golden.py:171-195 check_in_box2d）
	246	+// 把待测点 P 旋转 -heading 到框局部坐标系，检查是否在轴对齐矩形内（含 MARGIN=1e-2 绝对容差）。
	247	+// golden.py: angle_cos=cos(-h)=cosT, angle_sin=sin(-h)=-sinT
	248	+// rot_x = dpxangle_cos + dpy(-angle_sin) = dpxcosT + dpysinT
	249	+// rot_y = dpxangle_sin + dpyangle_cos = -dpxsinT + dpycosT
	250	+// return \|rot_x\| < dx/2 + MARGIN && \|rot_y\| < dy/2 + MARGIN （严格 <）
	251	+// 传入 sinT=sin(heading)、cosT=cos(heading)、halfDx=dx/2、halfDy=dy/2。
	252	+// ---------------------------------------------------------------------------
	253	+
	254	+__aicore__ inline bool Iou3D::PointInRect(float px, float py, float cx, float cy, float halfDx, float halfDy,
	255	+ float sinT, float cosT)
	256	+{
	257	+ float dpx = px - cx;
	258	+ float dpy = py - cy;
	259	+ float rotX = dpx * cosT + dpy * sinT;
	260	+ float rotY = -dpx * sinT + dpy * cosT;
	261	+ return (ScalarAbs(rotX) < halfDx + IOU3D_MARGIN) && (ScalarAbs(rotY) < halfDy + IOU3D_MARGIN);
	262	+}
	263	+
	264	+// ---------------------------------------------------------------------------
	265	+// SegIntersect（对标 mmcv intersection golden.py:104-168，逐分支复刻）
	266	+// ⚠️ 已知会重引入 θ=π fp32 伪交点风险（用户明确接受，为逐分支对标 mmcv 的权衡）。
	267	+// 变量映射（我方边 a=(a1→a2) 对应 mmcv p0→p1，边 b=(b1→b2) 对应 q0→q1）：
	268	+// p0=(a1x,a1y) p1=(a2x,a2y) q0=(b1x,b1y) q1=(b2x,b2y)
	269	+// ① check_rect_cross 快速排斥（golden.py:104-112 / 139）
	270	+// ② 跨立实验 s1s2>0 && s3s4>0（严格 >，共线/相切拒绝）（golden.py:143-149）
	271	+// cross_3pts(a,b,c)=(a.x-c.x)(b.y-c.y)-(b.x-c.x)(a.y-c.y)（golden.py:120）
	272	+// ③ \|s5-s1\|>EPS_KERNEL 用叉积公式，否则退化一般式直线方程（不判 D==0）（golden.py:152-166）
	273	+// ---------------------------------------------------------------------------
	274	+
	275	+__aicore__ inline bool Iou3D::SegIntersect(float a1x, float a1y, float a2x, float a2y, float b1x, float b1y, float b2x,
	276	+ float b2y, float& ox, float& oy)
	277	+{
	278	+ // p0=(a1x,a1y) p1=(a2x,a2y) q0=(b1x,b1y) q1=(b2x,b2y)
	279	+ // ① check_rect_cross 快速排斥（AABB 包围盒重叠判定），对标 golden.py:104-112
	280	+ if (!(ScalarMin(a1x, a2x) <= ScalarMax(b1x, b2x) && ScalarMin(b1x, b2x) <= ScalarMax(a1x, a2x) &&
	281	+ ScalarMin(a1y, a2y) <= ScalarMax(b1y, b2y) && ScalarMin(b1y, b2y) <= ScalarMax(a1y, a2y))) {
	282	+ return false;
	283	+ }
	284	+
	285	+ // cross_3pts(a,b,c) = (a.x-c.x)(b.y-c.y) - (b.x-c.x)(a.y-c.y)，对标 golden.py:115-120
	286	+ // ② 跨立实验（叉积判别），对标 golden.py:143-146
	287	+ float s1 = (b1x - a1x) * (a2y - a1y) - (a2x - a1x) * (b1y - a1y); // cross(q0, p1, p0)
	288	+ float s2 = (a2x - a1x) * (b2y - a1y) - (b2x - a1x) * (a2y - a1y); // cross(p1, q1, p0)
	289	+ float s3 = (a1x - b1x) * (b2y - b1y) - (b2x - b1x) * (a1y - b1y); // cross(p0, q1, q0)
	290	+ float s4 = (b2x - b1x) * (a2y - b1y) - (a2x - b1x) * (b2y - b1y); // cross(q1, p1, q0)
	291	+
	292	+ // 严格 >：共线/相切被拒（对标 golden.py:148）
	293	+ if (!(s1 * s2 > 0.0f && s3 * s4 > 0.0f)) {
	294	+ return false;
	295	+ }
	296	+
	297	+ // ③ 交点坐标计算（对标 golden.py:152-166）
	298	+ float s5 = (b2x - a1x) * (a2y - a1y) - (a2x - a1x) * (b2y - a1y); // cross(q1, p1, p0)
	299	+ if (ScalarAbs(s5 - s1) > IOU3D_EPS_KERNEL) {
	300	+ ox = (s5 * b1x - s1 * b2x) / (s5 - s1);
	301	+ oy = (s5 * b1y - s1 * b2y) / (s5 - s1);
	302	+ } else {
	303	+ // 退化情况：一般式直线方程（mmcv 不判 D==0，可产 inf/nan 并自然传播）
	304	+ float a0 = a1y - a2y;
	305	+ float b0 = a2x - a1x;
	306	+ float c0 = a1x * a2y - a2x * a1y;
	307	+ float a1c = b1y - b2y;
	308	+ float b1c = b2x - b1x;
	309	+ float c1c = b1x * b2y - b2x * b1y;
	310	+ float D = a0 * b1c - a1c * b0;
	311	+ ox = (b0 * c1c - b1c * c0) / D;
	312	+ oy = (a1c * c0 - a0 * c1c) / D;
	313	+ }
	314	+ return true;
	315	+}
	316	+
	317	+// ---------------------------------------------------------------------------
	318	+// CollectCorners（框1顶点∈框2 + 框2顶点∈框1 + 16 边对相交，含去重）
	319	+// ---------------------------------------------------------------------------
	320	+
	321	+__aicore__ inline uint32_t Iou3D::CollectCorners(const float c1x[IOU3D_CORNERS], const float c1y[IOU3D_CORNERS],
	322	+ const float c2x[IOU3D_CORNERS], const float c2y[IOU3D_CORNERS],
	323	+ float box1cx, float box1cy, float box1hx, float box1hy, float box1sin,
	324	+ float box1cos, float box2cx, float box2cy, float box2hx, float box2hy,
	325	+ float box2sin, float box2cos, float px[IOU3D_MAX_INTER],
	326	+ float py[IOU3D_MAX_INTER])
	327	+{
	328	+ float rawX[IOU3D_MAX_INTER + IOU3D_MAX_INTER + 16];
	329	+ float rawY[IOU3D_MAX_INTER + IOU3D_MAX_INTER + 16];
	330	+ uint32_t rawCnt = 0;
	331	+
	332	+ // 框1顶点 ∈ 框2（反向旋转到框2局部系测试），对标 mmcv golden.py:245 check_in_box2d(box_b, corners_a[k])
	333	+ for (uint32_t i = 0; i < IOU3D_CORNERS; ++i) {
	334	+ if (PointInRect(c1x[i], c1y[i], box2cx, box2cy, box2hx, box2hy, box2sin, box2cos)) {
	335	+ rawX[rawCnt] = c1x[i];
	336	+ rawY[rawCnt] = c1y[i];
	337	+ ++rawCnt;
	338	+ }
	339	+ }
	340	+ // 框2顶点 ∈ 框1，对标 mmcv golden.py:241 check_in_box2d(box_a, corners_b[k])
	341	+ for (uint32_t i = 0; i < IOU3D_CORNERS; ++i) {
	342	+ if (PointInRect(c2x[i], c2y[i], box1cx, box1cy, box1hx, box1hy, box1sin, box1cos)) {
	343	+ rawX[rawCnt] = c2x[i];
	344	+ rawY[rawCnt] = c2y[i];
	345	+ ++rawCnt;
	346	+ }
	347	+ }
	348	+ for (uint32_t i = 0; i < IOU3D_CORNERS; ++i) {
	349	+ float a1x = c1x[i], a1y = c1y[i];
	350	+ float a2x = c1x[(i + 1) % IOU3D_CORNERS], a2y = c1y[(i + 1) % IOU3D_CORNERS];
	351	+ for (uint32_t j = 0; j < IOU3D_CORNERS; ++j) {
	352	+ float b1x = c2x[j], b1y = c2y[j];
	353	+ float b2x = c2x[(j + 1) % IOU3D_CORNERS], b2y = c2y[(j + 1) % IOU3D_CORNERS];
	354	+ float ox, oy;
	355	+ if (SegIntersect(a1x, a1y, a2x, a2y, b1x, b1y, b2x, b2y, ox, oy)) {
	356	+ rawX[rawCnt] = ox;
	357	+ rawY[rawCnt] = oy;
	358	+ ++rawCnt;
	359	+ }
	360	+ }
	361	+ }
	362	+
	363	+ // 去重（幅值相对强去重）。tol = max(ABS_TOL, REL·max(1,\|x\|,\|y\|))，
	364	+ // 顶点自身幅值决定容差：小坐标退化为绝对 1e-6，大坐标（1e20）随幅值放大避免误并真顶点，
	365	+ // fp32 近重复顶点（identical box m=6/8）被坍缩回 4 真角点 → Sort32 键分离 → 稳定环序。
	366	+ // 最多保留 IOU3D_MAX_INTER 个。
	367	+ uint32_t m = 0;
	368	+ for (uint32_t k = 0; k < rawCnt; ++k) {
	369	+ float magK = ScalarMax(1.0f, ScalarMax(ScalarAbs(rawX[k]), ScalarAbs(rawY[k])));
	370	+ float tolK = ScalarMax(IOU3D_DEDUP_TOL, IOU3D_DEDUP_REL * magK);
	371	+ bool dup = false;
	372	+ for (uint32_t u = 0; u < m; ++u) {
	373	+ if (ScalarAbs(rawX[k] - px[u]) < tolK && ScalarAbs(rawY[k] - py[u]) < tolK) {
	374	+ dup = true;
	375	+ break;
	376	+ }
	377	+ }
	378	+ if (!dup && m < IOU3D_MAX_INTER) {
	379	+ px[m] = rawX[k];
	380	+ py[m] = rawY[k];
	381	+ ++m;
	382	+ }
	383	+ }
	384	+ return m;
	385	+}
	386	+
	387	+// ---------------------------------------------------------------------------
	388	+// TriArea（叉积法）
	389	+// ---------------------------------------------------------------------------
	390	+
	391	+__aicore__ inline float Iou3D::TriArea(float x1, float y1, float x2, float y2, float x3, float y3)
	392	+{
	393	+ float v = x1 * (y2 - y3) + x2 * (y3 - y1) + x3 * (y1 - y2);
	394	+ return ScalarAbs(v) * 0.5f;
	395	+}
	396	+
	397	+// ---------------------------------------------------------------------------
	398	+// PolygonArea：0/1/2 顶点 → 0（degenerate_polygon_guard）；3 → 三角形；
	399	+// >3 → 硬件 Sort32 极角排序（SortPolygonArea）
	400	+// ---------------------------------------------------------------------------
	401	+
	402	+__aicore__ inline float Iou3D::PolygonArea(float px[IOU3D_MAX_INTER], float py[IOU3D_MAX_INTER], uint32_t m)
	403	+{
	404	+ if (m < 3)
	405	+ return 0.0f; // degenerate_polygon_guard
	406	+ if (m == 3)
	407	+ return TriArea(px[0], py[0], px[1], py[1], px[2], py[2]);
	408	+ if (m > IOU3D_MAX_INTER)
	409	+ m = IOU3D_MAX_INTER; // 顶点数上限守卫（两凸四边形交集 <=8）
	410	+ return SortPolygonArea(px, py, m);
	411	+}
	412	+
	413	+// ---------------------------------------------------------------------------
	414	+// SortPolygonArea：>3 顶点，硬件 Sort32 极角排序（diamond-angle 键）+ 质心分解叉积面积。
	415	+// 1. 质心 (xc,yc)；相对质心坐标 (xr,yr)。
	416	+// 2. diamond-angle 键 key = (xr>=0)? t : (2-t)，t = yr/(\|xr\|+\|yr\|)，与 atan2(yr,xr) 单调同序；
	417	+// Sort32 为降序，故送 -key 得 angle 升序（逆时针）；padding 槽(idx>=m) 键置 -INF 沉底。
	418	+// 3. Sort32 + Extract 得排序后原始下标序 order[0..m-1]（逆时针环序）。
	419	+// 4. 质心分解：相邻 (C, v[order[i]], v[order[i+1]]) 三角形叉积绝对值累加。
	420	+// 面积对环序方向（顺/逆）不敏感（叉积取绝对值），与 atan2 升序等价。
	421	+// ---------------------------------------------------------------------------
	422	+
	423	+__aicore__ inline float Iou3D::SortPolygonArea(float px[IOU3D_MAX_INTER], float py[IOU3D_MAX_INTER], uint32_t m)
	424	+{
	425	+ LocalTensor<float> keyT = sortKeyBuf.Get<float>();
	426	+ LocalTensor<uint32_t> idxT = sortIdxBuf.Get<uint32_t>();
	427	+ LocalTensor<float> sortedT = sortedBuf.Get<float>();
	428	+ LocalTensor<float> outValT = sortOutValBuf.Get<float>();
	429	+ LocalTensor<uint32_t> outIdxT = sortOutIdxBuf.Get<uint32_t>();
	430	+
	431	+ // 质心（float32 全程，与 golden 一致）
	432	+ float xc = 0.0f, yc = 0.0f;
	433	+ for (uint32_t k = 0; k < m; ++k) {
	434	+ xc += px[k];
	435	+ yc += py[k];
	436	+ }
	437	+ float invm = 1.0f / static_cast<float>(m);
	438	+ xc *= invm;
	439	+ yc *= invm;
	440	+
	441	+ // diamond-angle 键 + 原始下标（padding 沉底）。32 槽全部填写，规避 Sort32 读到脏值。
	442	+ for (uint32_t i = 0; i < IOU3D_SORT32_LEN; ++i) {
	443	+ idxT.SetValue(i, i);
	444	+ float k;
	445	+ if (i < m) {
	446	+ float xr = px[i] - xc;
	447	+ float yr = py[i] - yc;
	448	+ float s = ScalarAbs(xr) + ScalarAbs(yr);
	449	+ float t = (s < 1.0e-20f) ? 0.0f : (yr / s);
	450	+ float keyRaw = (xr >= 0.0f) ? t : (2.0f - t);
	451	+ k = -keyRaw; // Sort32 降序 == angle 升序（逆时针）
	452	+ } else {
	453	+ k = IOU3D_NEG_INF_KEY; // padding 槽沉底
	454	+ }
	455	+ keyT.SetValue(i, k);
	456	+ }
	457	+ // SetValue(标量写 UB) → Sort32(向量读 UB) 之间同步
	458	+ PipeBarrier<PIPE_ALL>();
	459	+ Sort32<float>(sortedT, keyT, idxT, 1);
	460	+ PipeBarrier<PIPE_ALL>();
	461	+ Extract(outValT, outIdxT, sortedT, 1);
	462	+ // Extract(向量写 UB) → GetValue(标量读 UB) 之间同步
	463	+ PipeBarrier<PIPE_ALL>();
	464	+
	465	+ // 取前 m 个排序后原始下标（逆时针环序），越界守卫
	466	+ uint32_t order[IOU3D_MAX_INTER];
	467	+ for (uint32_t i = 0; i < m; ++i) {
	468	+ uint32_t rawIdx = outIdxT.GetValue(i) & IOU3D_IDX_MASK;
	469	+ if (rawIdx >= m)
	470	+ rawIdx = 0U; // 守卫：非法下标回退（不应发生，padding 已沉底）
	471	+ order[i] = rawIdx;
	472	+ }
	473	+
	474	+ // 面积（signed shoelace 收尾）。质心分解累加有符号叉积（不 per-triangle 取 abs），
	475	+ // 末尾对总和取一次 abs。去重已保证只剩真顶点、Sort32 环序为简单多边形，signed-shoelace
	476	+ // 与 per-triangle abs-fan 结果一致（对方向鲁棒）；但对「万一残留的近重复点/方向错序」多一层符号
	477	+ // 自洽保护——错序下正负三角形自然抵消，不会像 per-triangle abs-fan 那样把 bowtie 放大成半面积。
	478	+ // ⚠️ 严禁任何面积层 NaN 守卫；NaN/inf 必须按 IEEE754 自然传播到输出。
	479	+ float signedTwice = 0.0f; // 2·有符号面积（相对质心的鞋带和）
	480	+ for (uint32_t k = 0; k < m; ++k) {
	481	+ uint32_t a = order[k];
	482	+ uint32_t bnext = order[(k + 1U == m) ? 0U : (k + 1U)];
	483	+ float ax = px[a] - xc;
	484	+ float ay = py[a] - yc;
	485	+ float bx = px[bnext] - xc;
	486	+ float by = py[bnext] - yc;
	487	+ signedTwice += ax * by - ay * bx;
	488	+ }
	489	+ return ScalarAbs(signedTwice) * 0.5f;
	490	+}
	491	+
	492	+// ---------------------------------------------------------------------------
	493	+// ComputePairIou：单框对 3D IoU（对齐 golden Iou3dPair）
	494	+// ---------------------------------------------------------------------------
	495	+
	496	+__aicore__ inline float Iou3D::ComputePairIou(int64_t b, int64_t i, int64_t j, float sin1, float cos1, float sin2,
	497	+ float cos2)
	498	+{
	499	+ float rect1[IOU3D_DOF];
	500	+ float rect2[IOU3D_DOF];
	501	+ LoadBox(bboxesGM, b, i, static_cast<int64_t>(numN_), rect1); // [B,7,N]
	502	+ LoadBox(gtboxesGM, b, j, static_cast<int64_t>(numK_), rect2); // [B,7,K]
	503	+
	504	+ float v1 = rect1[3] * rect1[4] * rect1[5]; // whd
	505	+ float v2 = rect2[3] * rect2[4] * rect2[5];
	506	+ // D10 对标 mmcv：移除 v<=0 负体积短路（golden.py 无 v<=0 特判）。正尺寸零体积框自然得 iou=0
	507	+ // （interVol=0；D2 分母 clamp≥1e-6 保证不除零/nan）。
	508	+
	509	+ float c1x[IOU3D_CORNERS], c1y[IOU3D_CORNERS], c2x[IOU3D_CORNERS], c2y[IOU3D_CORNERS];
	510	+ BoxCorners(rect1[0], rect1[1], rect1[3], rect1[4], sin1, cos1, c1x, c1y);
	511	+ BoxCorners(rect2[0], rect2[1], rect2[3], rect2[4], sin2, cos2, c2x, c2y);
	512	+
	513	+ // Z 轴边界 + clamp 非负（clamp_z_overlap_nonneg）
	514	+ float z1Min = rect1[2] - 0.5f * rect1[5];
	515	+ float z1Max = rect1[2] + 0.5f * rect1[5];
	516	+ float z2Min = rect2[2] - 0.5f * rect2[5];
	517	+ float z2Max = rect2[2] + 0.5f * rect2[5];
	518	+ float realD = ScalarMax(ScalarMin(z1Max, z2Max) - ScalarMax(z1Min, z2Min), 0.0f);
	519	+
	520	+ float px[IOU3D_MAX_INTER], py[IOU3D_MAX_INTER];
	521	+ // 反向旋转包含测试需框中心/半尺寸/sin/cos：halfDx=0.5w, halfDy=0.5h。
	522	+ uint32_t m = CollectCorners(c1x, c1y, c2x, c2y, rect1[0], rect1[1], 0.5f * rect1[3], 0.5f * rect1[4], sin1, cos1,
	523	+ rect2[0], rect2[1], 0.5f * rect2[3], 0.5f * rect2[4], sin2, cos2, px, py);
	524	+ float area = PolygonArea(px, py, m);
	525	+
	526	+ float interVol = area * realD;
	527	+ float unionVol = v1 + v2 - interVol;
	528	+ // D2 对标 mmcv：分母 clamp（max(union, 1e-6)），对标 golden.py:377 clamp(min=1e-6)。
	529	+ float iou = interVol / ScalarMax(unionVol, IOU3D_EPSILON);
	530	+ return iou;
	531	+}
	532	+
	533	+// ---------------------------------------------------------------------------
	534	+// Process：逐 tile 向量批算 theta 的 sin/cos，逐对标量几何，普通 DataCopyPad 写回（单点不相交写）
	535	+// ---------------------------------------------------------------------------
	536	+
	537	+__aicore__ inline void Iou3D::Process()
	538	+{
	539	+ if (isEmpty_ != 0U) {
	540	+ return; // 空 Tensor：无输出（TPL_EMPTY 短路）
	541	+ }
	542	+ if (pairCount_ <= 0)
	543	+ return;
	544	+
	545	+ const uint32_t atl = alignedTl_; // 32B 对齐步长，保证每个子段起址 32B 对齐
	546	+ LocalTensor<float> angle = angleBuf.Get<float>(); // [6*atl]
	547	+ LocalTensor<float> theta1 = angle[0];
	548	+ LocalTensor<float> theta2 = angle[atl];
	549	+ LocalTensor<float> sin1Buf = angle[2U * atl];
	550	+ LocalTensor<float> cos1Buf = angle[3U * atl];
	551	+ LocalTensor<float> sin2Buf = angle[4U * atl];
	552	+ LocalTensor<float> cos2Buf = angle[5U * atl];
	553	+ LocalTensor<float> outTile = outBuf.Get<float>();
	554	+ LocalTensor<uint8_t> sinTmp = tmpBuf.Get<uint8_t>(); // 显式 sharedTmpBuffer（32B 对齐）
	555	+
	556	+ const int64_t nk = static_cast<int64_t>(numN_) * static_cast<int64_t>(numK_);
	557	+
	558	+ const uint32_t tl = (tileLen_ == 0U) ? 1U : tileLen_; // 逻辑批大小（未对齐）
	559	+ int64_t processed = 0;
	560	+ while (processed < pairCount_) {
	561	+ int64_t cur = pairCount_ - processed;
	562	+ if (cur > static_cast<int64_t>(tl))
	563	+ cur = static_cast<int64_t>(tl);
	564	+ uint32_t curNum = static_cast<uint32_t>(cur);
	565	+
	566	+ // 对齐后的向量计算长度（32B 倍数）。tail 元素(curNum..alignedCnt) 为 padding，
	567	+ // 预置为 0，使 Sin/Cos 在 padding 上产生合法值（sin0=0,cos0=1，无 NaN/Inf），结果不被读取。
	568	+ uint32_t alignedCnt = Iou3DCeilAlign(curNum);
	569	+ if (alignedCnt > atl)
	570	+ alignedCnt = atl; // 不超过子段容量
	571	+
	572	+ // 1) 收集本 tile 每对的两 theta（bbox_i.theta 通道=6，gtbox_j.theta 通道=6），
	573	+ // 并把对齐 padding 段清零。
	574	+ for (uint32_t t = 0; t < alignedCnt; ++t) {
	575	+ if (t >= curNum) {
	576	+ theta1.SetValue(t, 0.0f);
	577	+ theta2.SetValue(t, 0.0f);
	578	+ continue;
	579	+ }
	580	+ int64_t flat = pairStart_ + processed + static_cast<int64_t>(t);
	581	+ int64_t b = flat / nk;
	582	+ int64_t rem = flat % nk;
	583	+ int64_t ii = rem / static_cast<int64_t>(numK_);
	584	+ int64_t jj = rem % static_cast<int64_t>(numK_);
	585	+ float th1 = static_cast<float>(
	586	+ bboxesGM.GetValue(b * static_cast<int64_t>(IOU3D_DOF) * numN_ + 6 * static_cast<int64_t>(numN_) + ii));
	587	+ float th2 = static_cast<float>(
	588	+ gtboxesGM.GetValue(b * static_cast<int64_t>(IOU3D_DOF) * numK_ + 6 * static_cast<int64_t>(numK_) + jj));
	589	+ theta1.SetValue(t, th1);
	590	+ theta2.SetValue(t, th2);
	591	+ }
	592	+ // SetValue(标量写 UB) → Sin/Cos(向量读 UB) 之间插同步，确保标量写完成后向量再读。
	593	+ PipeBarrier<PIPE_ALL>();
	A A andong_hw7月21日不建议加 PIPE_ALL，插入需要的同步
	594	+
	595	+ // 2) 向量批算 sin/cos（32B 对齐子段 + 对齐计算长度 + 显式 sharedTmpBuffer）。
	596	+ // src/dst 起址均 32B 对齐，规避 RegBase Sin/Cos 的对齐约束（AIV err 340）。
	597	+ Sin(sin1Buf, theta1, sinTmp, alignedCnt);
	598	+ Cos(cos1Buf, theta1, sinTmp, alignedCnt);
	599	+ Sin(sin2Buf, theta2, sinTmp, alignedCnt);
	600	+ Cos(cos2Buf, theta2, sinTmp, alignedCnt);
	601	+ // Sin/Cos(向量写 UB) → GetValue(标量读 UB) 之间插同步，确保向量结果对标量可见。
	602	+ PipeBarrier<PIPE_ALL>();
	603	+
	604	+ // 3) 逐对标量几何计算 IoU
	605	+ for (uint32_t t = 0; t < curNum; ++t) {
	606	+ int64_t flat = pairStart_ + processed + static_cast<int64_t>(t);
	607	+ int64_t b = flat / nk;
	608	+ int64_t rem = flat % nk;
	609	+ int64_t ii = rem / static_cast<int64_t>(numK_);
	610	+ int64_t jj = rem % static_cast<int64_t>(numK_);
	611	+ float s1 = sin1Buf.GetValue(t);
	612	+ float c1 = cos1Buf.GetValue(t);
	613	+ float s2 = sin2Buf.GetValue(t);
	614	+ float c2 = cos2Buf.GetValue(t);
	615	+ float iou = ComputePairIou(b, ii, jj, s1, c1, s2, c2);
	616	+ outTile.SetValue(t, iou);
	617	+ }
	618	+
	619	+ // outTile 标量写 → DataCopyPad(MTE3 读 UB) 之间插同步，确保写完成后再搬出。
	620	+ PipeBarrier<PIPE_ALL>();
	621	+
	622	+ // 4) 写回：普通 DataCopyPad（本核连续段，(b,i,j) 单点写，无跨核冲突 → 无需原子加）。
	623	+ // 正确性论证（核实自 tiling 切分）：Host 侧 pairsPerCore=CeilDiv(total,coreNum)、
	624	+ // usedCoreNum=CeilDiv(total,pairsPerCore)、SetBlockDim(usedCoreNum)；kernel 侧核 k 负责
	625	+ // flatten 区间 [kpairsPerCore, kpairsPerCore + pairCount_)（pairCount_=min(pairsPerCore,remain)），
	626	+ // 相邻核区间连续且不相交，每个 (b,i,j) 输出元素恰被一个核写一次；核内 processed 单调递增，
	627	+ // 每元素只写一次。故写回为「不相交单点写」，普通 DataCopyPad 即完全正确。
	628	+ // 去掉 SetAtomicAdd/SetAtomicNone：原子累加语义要求目标 GM 预清零（累加到残留值），
	629	+ // 但算子自身不清零输出 GM，会对未零初始化的生产输出缓冲产生 garbage+value 错误；单点写
	630	+ // 无跨核累加需求，普通写既正确又消除「外部必须零初始化输出」的隐性契约、并省去原子开销。
	631	+ // outTile 源 UB 起址 32B 对齐；DataCopyPad 支持 sub-32B GM 长度（blockLen 按字节），
	632	+ // 尾段 padding(curNum..alignedCnt) 不搬出。
	633	+ DataCopyExtParams copyParams;
	634	+ copyParams.blockCount = 1;
	635	+ copyParams.blockLen = curNum * sizeof(float);
	636	+ copyParams.srcStride = 0;
	637	+ copyParams.dstStride = 0;
	638	+ copyParams.rsv = 0;
	639	+ DataCopyPad(iouGM[pairStart_ + processed], outTile, copyParams);
	640	+ PipeBarrier<PIPE_ALL>();
	641	+
	642	+ processed += cur;
	643	+ }
	644	+}
	645	+
	646	+} // namespace NsIou3D
	647	+#endif // IOU3D_ARCH35_H

Aobjdetect/iou3d/op_kernel/arch35/iou3d_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	+/*!
	12	+ * \file iou3d_tiling_data.h
	13	+ * \brief Iou3D TilingData 结构体定义（arch35）
	14	+ *
	15	+ * ✅ 使用标准 C++ struct 定义 TilingData
	16	+ * ❌ 禁止使用废弃的 BEGIN_TILING_DATA_DEF 宏
	17	+ *
	18	+ * 字段职责：
	19	+ * - batch/numBboxes/numGtboxes : 逻辑规模 B / N / K
	20	+ * - coreNum/pairsPerCore : 多核切分（总 (b,i,j) 对按核均分为不相交子集）
	21	+ * - tileLen/tailLen : 单核内 UB 批处理粒度（(i,j) 对数）
	22	+ * - isEmpty : 空 Tensor 标志（batch==0 \|\| N==0 \|\| K==0），用于 TPL_EMPTY 短路
	23	+ *
	24	+ * 注：极角排序（Sort32）临时 buffer 由 kernel 侧按固定 IOU3D_SORT32_LEN(32) 分配，与逻辑规模无关，
	25	+ * 故无需 Host 侧动态精算 sortTmpSize（历史遗留字段已移除）。
	26	+ */
	27	+
	28	+#ifndef _IOU3D_TILING_DATA_H_
	29	+#define _IOU3D_TILING_DATA_H_
	30	+
	31	+#include <cstdint>
	32	+
	33	+struct Iou3DTilingData {
	34	+ uint32_t batch = 0; // B
	35	+ uint32_t numBboxes = 0; // N（预测框数）
	36	+ uint32_t numGtboxes = 0; // K（真值框数，D5 对标 mmcv 已移除上限）
	37	+ uint32_t coreNum = 0; // 参与计算的核数
	38	+ uint32_t pairsPerCore = 0; // 每核负责的 (b,i,j) 对数（向上取整分配）
	39	+ uint32_t tileLen = 0; // 单批处理的 (i,j) 对数（UB 批大小）
	40	+ uint32_t tailLen = 0; // 尾批 (i,j) 对数
	41	+ uint32_t isEmpty = 0; // 空 Tensor 标志（batch==0 \|\| N==0 \|\| K==0）
	42	+};
	43	+#endif // _IOU3D_TILING_DATA_H_

Aobjdetect/iou3d/op_kernel/iou3d.cpp+33-0

@@ -0,0 +1,33 @@
	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 iou3d.cpp
	13	+ * \brief Iou3D 算子 kernel 入口（arch35 / DAV_3510）
	14	+ *
	15	+ * 一芯片一算子一入口：arch35 唯一 __global__ 入口。
	16	+ * 无 TilingKey 参数，单 dtype float32（由 def 文件驱动）。
	17	+ * - 空 Tensor 短路通过运行时判断（Init/Process 中判断 isEmpty_）
	18	+ *
	19	+ * def 驱动 dtype：dtype 由 _def.cpp 的 DataType({ge::DT_FLOAT}) 声明，构建系统通过
	20	+ * -DDTYPE_BBOXES 等编译宏注入，kernel 直接使用 DTYPE_BBOXES 宏获取实际类型。
	21	+ */
	22	+
	23	+#include "arch35/iou3d.h"
	24	+
	25	+// 核函数入口名须为 OpType(Iou3D) 的 snake_case 形式：iou3_d（框架 CamelCase→snake_case 约定）。
	26	+__global__ __aicore__ void iou3_d(GM_ADDR bboxes, GM_ADDR gtboxes, GM_ADDR iou, GM_ADDR workspace, GM_ADDR tiling)
	27	+{
	28	+ REGISTER_TILING_DEFAULT(Iou3DTilingData);
	29	+ GET_TILING_DATA_WITH_STRUCT(Iou3DTilingData, tilingData, tiling);
	30	+ NsIou3D::Iou3D op;
	31	+ op.Init(bboxes, gtboxes, iou, &tilingData);
	A A andong_hw7月21日是不是少了ut/st
	32	+ op.Process();
	33	+}

Aobjdetect/iou3d/tests/CMakeLists.txt+22-0

@@ -0,0 +1,22 @@
	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	+
	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	+# Generated By CANNBot
	12	+
	13	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	14	+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT OR OP_KERNEL_UT))
	15	+ list(REMOVE_ITEM CURRENT_DIRS ut)
	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()

Aobjdetect/iou3d/tests/assets/golden.py+312-0

@@ -0,0 +1,312 @@
	1	+#!/usr/bin/env python3
	2	+# -- coding: UTF-8 --
	3	+# ----------------------------------------------------------------------------
	4	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	5	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# ----------------------------------------------------------------------------
	12	+
	13	+# Generated By CANNBot
	14	+
	15	+"""
	16	+Iou3D Golden Reference —— 3D 旋转框 IoU 的独立 ground truth（对标 mmcv iou3d_cuda）。
	17	+
	18	+【独立性声明】
	19	+ 本 golden 直接从旋转框 3D IoU 的几何数学定义计算真值，**不镜像 AscendC kernel、
	20	+ 不引用 kernel 中间结果反推**。数学契约逐条对齐 mmcv iou3d_cuda_kernel.cuh + mmcv/iou3d.py：
	21	+ - 顶点顺序：P0=左下, P1=右下, P2=右上, P3=左上；rotate_around_center 先轴对齐后旋转
	22	+ - 边相交：check_rect_cross 快速排斥 + 跨立实验（严格 >）+ 条件叉积/一般式直线方程（不判 D==0）
	23	+ - 顶点包含：反向旋转法 + MARGIN=1e-2 绝对容差（严格 <）
	24	+ - 极角排序：atan2 冒泡排序；面积：扇形分解（以首顶点 p[0] 为基准），不去重
	25	+ - 常量：EPS_KERNEL=1e-8, MARGIN=1e-2, EPS_IOU=1e-6
	26	+ - IoU 分母：clamp(vol_a + vol_b - inter, min=1e-6)（clamp，非加法）
	27	+ - 无 v<=0 短路（正尺寸零体积框自然得 IoU=0）
	28	+ 与本仓 op_kernel/arch35/iou3d.h 的 mmcv 迁移决策一致（详见 tmp/iou3d_vs_mmcv_diff.md）。
	29	+
	30	+【接口】（参数名/序对齐 op_host/iou3d_def.cpp，供 ST/TTK 精度框架消费）
	31	+ iou3d_golden(bboxes, gtboxes, **kwargs) -> iou
	32	+ bboxes: numpy.ndarray [B, 7, N] float32 (7-DoF: x, y, z, w, h, d, theta)
	33	+ gtboxes: numpy.ndarray [B, 7, K] float32
	34	+ iou: numpy.ndarray [B, N, K] float32
	35	+ 几何计算全程 float64，末尾 cast 回 float32（高精度参考；与 kernel fp32 的容差由框架管理）。
	36	+"""
	37	+
	38	+import math
	39	+
	40	+import numpy as np
	41	+
	42	+__golden__ = {
	43	+ "kernel": {"iou3d": "iou3d_golden"},
	44	+}
	45	+
	46	+# 常量（与 mmcv iou3d_cuda_kernel.cuh / iou3d.py 一致）
	47	+EPS_KERNEL = 1e-8 # 边相交退化判据阈值 (iou3d_cuda_kernel.cuh)
	48	+MARGIN = 1e-2 # 顶点包含测试绝对容差 (iou3d_cuda_kernel.cuh)
	49	+EPS_IOU = 1e-6 # IoU 分母 clamp 下限 (mmcv iou3d.py)
	50	+
	51	+
	52	+def _rotate_around_center(cx, cy, angle_cos, angle_sin, px, py):
	53	+ """绕中心点旋转（mmcv rotate_around_center）。
	54	+ x' = (px-cx)cos - (py-cy)sin + cx ; y' = (px-cx)sin + (py-cy)cos + cy。"""
	55	+ dx = px - cx
	56	+ dy = py - cy
	57	+ new_x = dx * angle_cos - dy * angle_sin + cx
	58	+ new_y = dx * angle_sin + dy * angle_cos + cy
	59	+ return new_x, new_y
	60	+
	61	+
	62	+def _get_box_corners(box):
	63	+ """轴对齐顶点 + 整体旋转（mmcv box_overlap）。顶点顺序 P0 左下→P1 右下→P2 右上→P3 左上。
	64	+ box = [x, y, z, w(dx), h(dy), d(dz), theta]。"""
	65	+ cx, cy = float(box[0]), float(box[1])
	66	+ dx, dy = float(box[3]), float(box[4])
	67	+ heading = float(box[6])
	68	+ dx_half, dy_half = dx / 2.0, dy / 2.0
	69	+ x1, y1 = cx - dx_half, cy - dy_half
	70	+ x2, y2 = cx + dx_half, cy + dy_half
	71	+ pts = [(x1, y1), (x2, y1), (x2, y2), (x1, y2)] # 左下右下右上左上
	72	+ angle_cos, angle_sin = math.cos(heading), math.sin(heading)
	73	+ return [
	74	+ _rotate_around_center(cx, cy, angle_cos, angle_sin, px, py) for px, py in pts
	75	+ ]
	76	+
	77	+
	78	+def _check_rect_cross(p0, p1, q0, q1):
	79	+ """AABB 快速排斥实验（mmcv check_rect_cross）。"""
	80	+ return (
	81	+ min(p0[0], p1[0]) <= max(q0[0], q1[0])
	82	+ and min(q0[0], q1[0]) <= max(p0[0], p1[0])
	83	+ and min(p0[1], p1[1]) <= max(q0[1], q1[1])
	84	+ and min(q0[1], q1[1]) <= max(p0[1], p1[1])
	85	+ )
	86	+
	87	+
	88	+def _cross_3pts(a, b, c):
	89	+ """叉积 (a-c) × (b-c)（mmcv cross）。"""
	90	+ return (a[0] - c[0]) * (b[1] - c[1]) - (b[0] - c[0]) * (a[1] - c[1])
	91	+
	92	+
	93	+def _intersection(p1, p0, q1, q0):
	94	+ """线段 P0P1 × Q0Q1 求交（mmcv intersection，参数序 p1,p0,q1,q0）。
	95	+ 快速排斥 + 跨立实验（严格 >）+ \|s5-s1\|>EPS 叉积 / 否则一般式（不判 D==0）。返回 (x,y) 或 None。"""
	96	+ if not _check_rect_cross(p0, p1, q0, q1):
	97	+ return None
	98	+ s1 = _cross_3pts(q0, p1, p0)
	99	+ s2 = _cross_3pts(p1, q1, p0)
	100	+ s3 = _cross_3pts(p0, q1, q0)
	101	+ s4 = _cross_3pts(q1, p1, q0)
	102	+ if not (s1 * s2 > 0.0 and s3 * s4 > 0.0):
	103	+ return None
	104	+ s5 = _cross_3pts(q1, p1, p0)
	105	+ if abs(s5 - s1) > EPS_KERNEL:
	106	+ ans_x = (s5 * q0[0] - s1 * q1[0]) / (s5 - s1)
	107	+ ans_y = (s5 * q0[1] - s1 * q1[1]) / (s5 - s1)
	108	+ else:
	109	+ a0, b0 = p0[1] - p1[1], p1[0] - p0[0]
	110	+ c0 = p0[0] * p1[1] - p1[0] * p0[1]
	111	+ a1, b1 = q0[1] - q1[1], q1[0] - q0[0]
	112	+ c1 = q0[0] * q1[1] - q1[0] * q0[1]
	113	+ d = a0 * b1 - a1 * b0
	114	+ ans_x = (b0 * c1 - b1 * c0) / d
	115	+ ans_y = (a1 * c0 - a0 * c1) / d
	116	+ return (ans_x, ans_y)
	117	+
	118	+
	119	+def _check_in_box2d(box, p):
	120	+ """顶点包含测试 —— 反向旋转法 + MARGIN=1e-2（mmcv check_in_box2d）。
	121	+ 把点 P 旋转 -heading 到框局部系，判 \|rot_x\|<dx/2+MARGIN && \|rot_y\|<dy/2+MARGIN（严格 <）。"""
	122	+ cx, cy = float(box[0]), float(box[1])
	123	+ dx, dy = float(box[3]), float(box[4])
	124	+ heading = float(box[6])
	125	+ angle_cos, angle_sin = math.cos(-heading), math.sin(-heading)
	126	+ dpx, dpy = p[0] - cx, p[1] - cy
	127	+ rot_x = dpx * angle_cos + dpy * (-angle_sin)
	128	+ rot_y = dpx * angle_sin + dpy * angle_cos
	129	+ return abs(rot_x) < dx / 2.0 + MARGIN and abs(rot_y) < dy / 2.0 + MARGIN
	130	+
	131	+
	132	+def _point_cmp(a, b, center):
	133	+ """极角比较：atan2(a-center) > atan2(b-center)（mmcv point_cmp）。"""
	134	+ return math.atan2(a[1] - center[1], a[0] - center[0]) > math.atan2(
	135	+ b[1] - center[1], b[0] - center[0]
	136	+ )
	137	+
	138	+
	139	+def _box_overlap(box_a, box_b):
	140	+ """BEV 重叠面积（mmcv box_overlap，逐行复刻）。16 组边交点 + 8 次顶点包含，
	141	+ 质心=收集点均值，atan2 冒泡排序，扇形分解（以 p[0] 为基准），不去重。"""
	142	+ corners_a = _get_box_corners(box_a)
	143	+ corners_b = _get_box_corners(box_b)
	144	+ ca = corners_a + [corners_a[0]] # 闭合
	145	+ cb = corners_b + [corners_b[0]]
	146	+
	147	+ cross_points = []
	148	+ center = [0.0, 0.0]
	149	+
	150	+ # 边相交（4×4=16）
	151	+ for i in range(4):
	152	+ for j in range(4):
	153	+ pt = _intersection(ca[i + 1], ca[i], cb[j + 1], cb[j])
	154	+ if pt is not None:
	155	+ cross_points.append(pt)
	156	+ center[0] += pt[0]
	157	+ center[1] += pt[1]
	158	+ # 顶点包含（8 次）
	159	+ for k in range(4):
	160	+ if _check_in_box2d(box_a, corners_b[k]):
	161	+ cross_points.append(corners_b[k])
	162	+ center[0] += corners_b[k][0]
	163	+ center[1] += corners_b[k][1]
	164	+ if _check_in_box2d(box_b, corners_a[k]):
	165	+ cross_points.append(corners_a[k])
	166	+ center[0] += corners_a[k][0]
	167	+ center[1] += corners_a[k][1]
	168	+
	169	+ cnt = len(cross_points)
	170	+ if cnt == 0:
	171	+ return 0.0
	172	+ center[0] /= cnt
	173	+ center[1] /= cnt
	174	+
	175	+ # 极角冒泡排序
	176	+ for j in range(cnt - 1):
	177	+ for i in range(cnt - j - 1):
	178	+ if _point_cmp(cross_points[i], cross_points[i + 1], center):
	179	+ cross_points[i], cross_points[i + 1] = (
	180	+ cross_points[i + 1],
	181	+ cross_points[i],
	182	+ )
	183	+
	184	+ # 扇形分解（以 cross_points[0] 为基准）
	185	+ area = 0.0
	186	+ p0x, p0y = cross_points[0]
	187	+ for k in range(cnt - 1):
	188	+ ax, ay = cross_points[k][0] - p0x, cross_points[k][1] - p0y
	189	+ bx, by = cross_points[k + 1][0] - p0x, cross_points[k + 1][1] - p0y
	190	+ area += ax * by - ay * bx
	191	+ return abs(area) / 2.0
	192	+
	193	+
	194	+def _iou3d_pair(rect1, rect2):
	195	+ """单框对 3D IoU。rect = [x, y, z, w, h, d, theta]（float64 直算）。"""
	196	+ bev = _box_overlap(rect1, rect2)
	197	+ # Z 轴重叠高度：z ± d/2，clamp 非负
	198	+ z1_max, z1_min = rect1[2] + rect1[5] / 2.0, rect1[2] - rect1[5] / 2.0
	199	+ z2_max, z2_min = rect2[2] + rect2[5] / 2.0, rect2[2] - rect2[5] / 2.0
	200	+ overlaps_h = max(min(z1_max, z2_max) - max(z1_min, z2_min), 0.0)
	201	+ inter = bev * overlaps_h
	202	+ vol1 = rect1[3] * rect1[4] * rect1[5]
	203	+ vol2 = rect2[3] * rect2[4] * rect2[5]
	204	+ union = vol1 + vol2 - inter
	205	+ return inter / max(union, EPS_IOU) # clamp（对标 mmcv，非加法 epsilon）
	206	+
	207	+
	208	+def iou3d_golden(bboxes, gtboxes, **kwargs):
	209	+ """3D 旋转框 IoU golden。
	210	+
	211	+ bboxes [B,7,N], gtboxes [B,7,K] -> iou [B,N,K] (float32)。
	212	+ 参数名/序对齐 op_host/iou3d_def.cpp（无 outputs）；空 Tensor（N==0 或 K==0）返回对应空形状。
	213	+ """
	214	+ b_np = np.asarray(bboxes, dtype=np.float64)
	215	+ g_np = np.asarray(gtboxes, dtype=np.float64)
	216	+ batch, dof, num_n = b_np.shape
	217	+ num_k = g_np.shape[2]
	218	+ iou = np.zeros((batch, num_n, num_k), dtype=np.float32)
	219	+ if num_n == 0 or num_k == 0:
	220	+ return iou # 空 Tensor 短路
	221	+ for b in range(batch):
	222	+ for i in range(num_n):
	223	+ rect1 = [b_np[b, p, i] for p in range(7)]
	224	+ for j in range(num_k):
	225	+ rect2 = [g_np[b, p, j] for p in range(7)]
	226	+ iou[b, i, j] = np.float32(_iou3d_pair(rect1, rect2))
	227	+ return iou
	228	+
	229	+
	230	+# ============================================================================
	231	+# Golden 自测（解析真值锚点，交叉验证独立 golden 可靠性；无 NPU，CPU 即可跑）
	232	+# ============================================================================
	233	+def _self_test():
	234	+ # (name, rect1, rect2, expected_iou, tol)
	235	+ cases = [
	236	+ ("identical", [0, 0, 0, 2, 2, 2, 0], [0, 0, 0, 2, 2, 2, 0], 1.0, 1e-4),
	237	+ ("disjoint", [0, 0, 0, 1, 1, 1, 0], [10, 10, 10, 1, 1, 1, 0], 0.0, 1e-6),
	238	+ ("axis_partial", [0, 0, 0, 2, 2, 2, 0], [1, 0, 0, 2, 2, 2, 0], 4.0 / 12, 1e-3),
	239	+ ("z_disjoint", [0, 0, 0, 2, 2, 1, 0], [0, 0, 10, 2, 2, 1, 0], 0.0, 1e-6),
	240	+ (
	241	+ "theta_pi4",
	242	+ [0, 0, 0, 2, 2, 2, 0],
	243	+ [0, 0, 0, 2, 2, 2, math.pi / 4],
	244	+ 0.7071,
	245	+ 5e-2,
	246	+ ),
	247	+ ("contained", [0, 0, 0, 4, 4, 4, 0], [0, 0, 0, 2, 2, 2, 0], 8.0 / 64, 1e-3),
	248	+ (
	249	+ "theta_pi2",
	250	+ [0, 0, 0, 2, 4, 2, 0],
	251	+ [0, 0, 0, 2, 4, 2, math.pi / 2],
	252	+ 4.0 / 12,
	253	+ 1e-3,
	254	+ ),
	255	+ ("theta_pi", [0, 0, 0, 2, 3, 2, 0], [0, 0, 0, 2, 3, 2, math.pi], 1.0, 1e-3),
	256	+ ("zero_vol_iou0", [0, 0, 0, 0, 2, 2, 0], [0, 0, 0, 2, 2, 2, 0], 0.0, 1e-6),
	257	+ ]
	258	+ fail = 0
	259	+ for name, r1, r2, exp, tol in cases:
	260	+ got = float(_iou3d_pair([float(v) for v in r1], [float(v) for v in r2]))
	261	+ ok = abs(got - exp) < tol
	262	+ print(
	263	+ f" [{'PASS' if ok else 'FAIL'}] pair {name}: iou={got:.6f} (expect {exp:.6f}, tol={tol:.1e})"
	264	+ )
	265	+ fail += 0 if ok else 1
	266	+
	267	+ # tensor 接口自测：[1,7,2] × [1,7,2] → [1,2,2]，对角=1.0（identical），非对角=0（disjoint）
	268	+ b = np.array(
	269	+ [
	270	+ [
	271	+ [0.0, 100.0],
	272	+ [0.0, 100.0],
	273	+ [0.0, 0.0],
	274	+ [2.0, 2.0],
	275	+ [2.0, 2.0],
	276	+ [2.0, 2.0],
	277	+ [0.0, 0.0],
	278	+ ]
	279	+ ],
	280	+ dtype=np.float32,
	281	+ )
	282	+ iou = iou3d_golden(b, b)
	283	+ tensor_ok = (
	284	+ iou.shape == (1, 2, 2)
	285	+ and abs(iou[0, 0, 0] - 1.0) < 1e-4
	286	+ and abs(iou[0, 1, 1] - 1.0) < 1e-4
	287	+ and iou[0, 0, 1] < 1e-6
	288	+ and iou[0, 1, 0] < 1e-6
	289	+ )
	290	+ print(
	291	+ f" [{'PASS' if tensor_ok else 'FAIL'}] tensor [1,7,2]->[1,2,2] 对角=1.0 非对角=0.0"
	292	+ )
	293	+ fail += 0 if tensor_ok else 1
	294	+
	295	+ # 空 Tensor 自测：N==0 -> [1,0,2]
	296	+ empty = iou3d_golden(
	297	+ np.zeros((1, 7, 0), np.float32), np.zeros((1, 7, 2), np.float32)
	298	+ )
	299	+ empty_ok = empty.shape == (1, 0, 2)
	300	+ print(
	301	+ f" [{'PASS' if empty_ok else 'FAIL'}] empty tensor N=0 -> shape {empty.shape}"
	302	+ )
	303	+ fail += 0 if empty_ok else 1
	304	+
	305	+ print(f"golden self-check: {'ALL PASS' if fail == 0 else f'{fail} FAILED'}")
	306	+ return fail
	307	+
	308	+
	309	+if __name__ == "__main__":
	310	+ import sys
	311	+
	312	+ sys.exit(_self_test())

Aobjdetect/iou3d/tests/ut/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	+
	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	+# Generated By CANNBot
	12	+
	13	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	14	+
	15	+if(NOT (UT_TEST_ALL OR OP_HOST_UT OR OP_API_UT))
	16	+ list(REMOVE_ITEM CURRENT_DIRS op_host)
	17	+endif()
	18	+
	19	+if(NOT (UT_TEST_ALL OR OP_KERNEL_UT))
	20	+ list(REMOVE_ITEM CURRENT_DIRS op_kernel)
	21	+endif()
	22	+
	23	+foreach(SUB_DIR ${CURRENT_DIRS})
	24	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	25	+ add_subdirectory(${SUB_DIR})
	26	+ endif()
	27	+endforeach()

Aobjdetect/iou3d/tests/ut/op_host/CMakeLists.txt+20-0

@@ -0,0 +1,20 @@
	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	+
	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	+# Generated By CANNBot
	12	+
	13	+if(UT_TEST_ALL OR OP_HOST_UT)
	14	+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	15	+ add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	16	+
	17	+ if("${ASCEND_COMPUTE_UNIT}" STREQUAL "ascend950")
	18	+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}/arch35)
	19	+ endif()
	20	+endif()

Aobjdetect/iou3d/tests/ut/op_host/arch35/test_iou3d_tiling.cpp+175-0

@@ -0,0 +1,175 @@
	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 "any_value.h"
	18	+#include "../../../../op_kernel/arch35/iou3d_tiling_data.h"
	19	+
	20	+using namespace std;
	21	+using namespace ge;
	22	+
	23	+class Iou3DTiling : public testing::Test {
	24	+protected:
	25	+ static void SetUpTestCase() { std::cout << "Iou3DTiling SetUp" << std::endl; }
	26	+
	27	+ static void TearDownTestCase() { std::cout << "Iou3DTiling TearDown" << std::endl; }
	28	+};
	29	+
	30	+namespace {
	31	+constexpr uint64_t CORE_NUM = 64;
	32	+constexpr uint64_t UB_SIZE = 262144;
	33	+constexpr uint32_t IOU3D_TILE_LEN = 256U; // 与 tiling 源码 IOU3D_TILE_LEN 一致
	34	+
	35	+// CeilDiv（与 tiling 源码 Ops::Base::CeilDiv 语义一致）
	36	+int64_t CeilDiv(int64_t a, int64_t b) { return (b == 0) ? 0 : (a + b - 1) / b; }
	37	+
	38	+// Iou3D 的 CompileInfo（与 tiling 源码 struct Iou3DCompileInfo {} 对应，入图/tiling 场景需非空指针）。
	39	+struct Iou3DCompileInfo {};
	40	+static Iou3DCompileInfo g_iou3dCompileInfo;
	41	+
	42	+// 构造一个 Iou3D 的 TilingContextPara（无 attr，float32/ND，Ascend950）。
	43	+// bboxes[B,7,N] + gtboxes[B,7,K] -> iou[B,N,K]。gert::StorageShape 只接受 initializer_list，故显式传维。
	44	+gert::TilingContextPara MakePara(int64_t batch, int64_t numN, int64_t channelB, int64_t batchG, int64_t numK,
	45	+ int64_t channelG, ge::DataType dtype = ge::DT_FLOAT)
	46	+{
	47	+ return gert::TilingContextPara(
	48	+ "Iou3D",
	49	+ {
	50	+ {gert::StorageShape({batch, channelB, numN}, {batch, channelB, numN}), dtype, ge::FORMAT_ND},
	51	+ {gert::StorageShape({batchG, channelG, numK}, {batchG, channelG, numK}), dtype, ge::FORMAT_ND},
	52	+ },
	53	+ {
	54	+ {gert::StorageShape({batch, numN, numK}, {batch, numN, numK}), ge::DT_FLOAT, ge::FORMAT_ND},
	55	+ },
	56	+ {}, &g_iou3dCompileInfo, "Ascend950", CORE_NUM, UB_SIZE);
	57	+}
	58	+
	59	+// 常规同 batch / channel=7 的便捷构造：MakeParaBNK(B, N, K)
	60	+gert::TilingContextPara MakeParaBNK(int64_t batch, int64_t numN, int64_t numK, ge::DataType dtype = ge::DT_FLOAT)
	61	+{
	62	+ return MakePara(batch, numN, 7, batch, numK, 7, dtype);
	63	+}
	64	+
	65	+const Iou3DTilingData* AsTiling(const TilingInfo& info)
	66	+{
	67	+ EXPECT_EQ(info.tilingDataSize, sizeof(Iou3DTilingData));
	68	+ return reinterpret_cast<const Iou3DTilingData*>(info.tilingData.get());
	69	+}
	70	+} // namespace
	71	+
	72	+// 正例：小 shape [1,7,4] + [1,7,4]。校验 B/N/K、非空、tileLen>0、workspace==0。
	73	+TEST_F(Iou3DTiling, iou3d_tiling_small)
	74	+{
	75	+ auto para = MakeParaBNK(1, 4, 4);
	76	+ TilingInfo info;
	77	+ ASSERT_TRUE(ExecuteTiling(para, info));
	78	+ const Iou3DTilingData* t = AsTiling(info);
	79	+
	80	+ EXPECT_EQ(t->batch, 1U);
	81	+ EXPECT_EQ(t->numBboxes, 4U);
	82	+ EXPECT_EQ(t->numGtboxes, 4U);
	83	+ EXPECT_EQ(t->isEmpty, 0U);
	84	+ EXPECT_GT(t->tileLen, 0U);
	85	+ EXPECT_GE(t->coreNum, 1U);
	86	+ // totalPairs=16 < coreNum=64 -> pairsPerCore=1, tileLen=min(1,256)=1
	87	+ const int64_t totalPairs = 1 * 4 * 4;
	88	+ const int64_t pairsPerCore = CeilDiv(totalPairs, CORE_NUM);
	89	+ EXPECT_EQ(t->pairsPerCore, static_cast<uint32_t>(pairsPerCore));
	90	+ EXPECT_EQ(t->tileLen, static_cast<uint32_t>(std::min<int64_t>(pairsPerCore, IOU3D_TILE_LEN)));
	91	+ ASSERT_EQ(info.workspaceSizes.size(), 1U);
	92	+ EXPECT_EQ(info.workspaceSizes[0], 0U);
	93	+ EXPECT_EQ(info.blockNum, static_cast<size_t>(CeilDiv(totalPairs, pairsPerCore)));
	94	+}
	95	+
	96	+// 正例：大 shape [2,7,256] + [2,7,256]。校验多核切分 + tileLen 封顶 256。
	97	+TEST_F(Iou3DTiling, iou3d_tiling_large_multicore)
	98	+{
	99	+ auto para = MakeParaBNK(2, 256, 256);
	100	+ TilingInfo info;
	101	+ ASSERT_TRUE(ExecuteTiling(para, info));
	102	+ const Iou3DTilingData* t = AsTiling(info);
	103	+
	104	+ EXPECT_EQ(t->batch, 2U);
	105	+ EXPECT_EQ(t->numBboxes, 256U);
	106	+ EXPECT_EQ(t->numGtboxes, 256U);
	107	+ EXPECT_EQ(t->isEmpty, 0U);
	108	+ const int64_t totalPairs = 2 * 256 * 256; // 131072
	109	+ const int64_t pairsPerCore = CeilDiv(totalPairs, CORE_NUM);
	110	+ const int64_t usedCore = CeilDiv(totalPairs, pairsPerCore);
	111	+ EXPECT_EQ(t->pairsPerCore, static_cast<uint32_t>(pairsPerCore));
	112	+ EXPECT_EQ(t->coreNum, static_cast<uint32_t>(usedCore));
	113	+ // pairsPerCore(2048) > 256 -> tileLen 封顶 IOU3D_TILE_LEN
	114	+ EXPECT_EQ(t->tileLen, IOU3D_TILE_LEN);
	115	+ EXPECT_EQ(info.blockNum, static_cast<size_t>(usedCore));
	116	+}
	117	+
	118	+// 正例：大 K（D5 已移除 K<=2000 上限，K=4096 应成功）。
	119	+TEST_F(Iou3DTiling, iou3d_tiling_large_k)
	120	+{
	121	+ auto para = MakeParaBNK(1, 8, 4096);
	122	+ TilingInfo info;
	123	+ ASSERT_TRUE(ExecuteTiling(para, info));
	124	+ const Iou3DTilingData* t = AsTiling(info);
	125	+ EXPECT_EQ(t->numGtboxes, 4096U);
	126	+ EXPECT_EQ(t->isEmpty, 0U);
	127	+}
	128	+
	129	+// 正例：空 Tensor（N==0）-> isEmpty=1, blockNum=1。
	130	+TEST_F(Iou3DTiling, iou3d_tiling_empty_n)
	131	+{
	132	+ auto para = MakeParaBNK(1, 0, 4);
	133	+ TilingInfo info;
	134	+ ASSERT_TRUE(ExecuteTiling(para, info));
	135	+ const Iou3DTilingData* t = AsTiling(info);
	136	+ EXPECT_EQ(t->isEmpty, 1U);
	137	+ EXPECT_EQ(info.blockNum, 1U);
	138	+}
	139	+
	140	+// 正例：空 Tensor（batch==0）-> isEmpty=1, blockNum=1（batch=0 若不短路会 SetBlockDim(0) 非法）。
	141	+TEST_F(Iou3DTiling, iou3d_tiling_empty_batch)
	142	+{
	143	+ auto para = MakeParaBNK(0, 4, 4);
	144	+ TilingInfo info;
	145	+ ASSERT_TRUE(ExecuteTiling(para, info));
	146	+ const Iou3DTilingData* t = AsTiling(info);
	147	+ EXPECT_EQ(t->isEmpty, 1U);
	148	+ EXPECT_EQ(info.blockNum, 1U);
	149	+}
	150	+
	151	+// 反例：channel != 7 -> tiling FAILED
	152	+TEST_F(Iou3DTiling, iou3d_tiling_neg_channel_not_7)
	153	+{
	154	+ // channel=5（dim1≠7），bboxes[1,5,4] + gtboxes[1,5,4]
	155	+ auto para = MakePara(1, 4, 5, 1, 4, 5);
	156	+ TilingInfo info;
	157	+ EXPECT_FALSE(ExecuteTiling(para, info));
	158	+}
	159	+
	160	+// 反例：batch 不一致 -> tiling FAILED
	161	+TEST_F(Iou3DTiling, iou3d_tiling_neg_batch_mismatch)
	162	+{
	163	+ // bboxes[2,7,4] + gtboxes[3,7,4]
	164	+ auto para = MakePara(2, 4, 7, 3, 4, 7);
	165	+ TilingInfo info;
	166	+ EXPECT_FALSE(ExecuteTiling(para, info));
	167	+}
	168	+
	169	+// 反例：dtype != float32（float16）-> tiling FAILED
	170	+TEST_F(Iou3DTiling, iou3d_tiling_neg_dtype_fp16)
	171	+{
	172	+ auto para = MakeParaBNK(1, 4, 4, ge::DT_FLOAT16);
	173	+ TilingInfo info;
	174	+ EXPECT_FALSE(ExecuteTiling(para, info));
	175	+}

Aobjdetect/iou3d/tests/ut/op_host/test_iou3d_infershape.cpp+131-0

@@ -0,0 +1,131 @@
	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 Iou3DInfershape : public testing::Test {
	19	+protected:
	20	+ static void SetUpTestCase() { std::cout << "Iou3DInfershape SetUp" << std::endl; }
	21	+
	22	+ static void TearDownTestCase() { std::cout << "Iou3DInfershape TearDown" << std::endl; }
	23	+};
	24	+
	25	+// 正例：静态 shape，bboxes[2,7,8] + gtboxes[2,7,16] -> iou[2,8,16]
	26	+TEST_F(Iou3DInfershape, iou3d_infershape_static)
	27	+{
	28	+ gert::InfershapeContextPara infershapeContextPara("Iou3D",
	29	+ {
	30	+ {{{2, 7, 8}, {2, 7, 8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	31	+ {{{2, 7, 16}, {2, 7, 16}}, ge::DT_FLOAT, ge::FORMAT_ND},
	32	+ },
	33	+ {
	34	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	35	+ });
	36	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	37	+ {2, 8, 16},
	38	+ };
	39	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	40	+}
	41	+
	42	+// 正例：动态 shape，bboxes[-1,7,-1] + gtboxes[-1,7,-1] -> iou[-1,-1,-1]
	43	+TEST_F(Iou3DInfershape, iou3d_infershape_dynamic)
	44	+{
	45	+ gert::InfershapeContextPara infershapeContextPara("Iou3D",
	46	+ {
	47	+ {{{-1, 7, -1}, {-1, 7, -1}}, ge::DT_FLOAT, ge::FORMAT_ND},
	48	+ {{{-1, 7, -1}, {-1, 7, -1}}, ge::DT_FLOAT, ge::FORMAT_ND},
	49	+ },
	50	+ {
	51	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	52	+ });
	53	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	54	+ {-1, -1, -1},
	55	+ };
	56	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	57	+}
	58	+
	59	+// 正例：大 K（D5 对标 mmcv 已移除 K<=2000 上限，K=4096 应成功）
	60	+TEST_F(Iou3DInfershape, iou3d_infershape_large_k)
	61	+{
	62	+ gert::InfershapeContextPara infershapeContextPara("Iou3D",
	63	+ {
	64	+ {{{1, 7, 64}, {1, 7, 64}}, ge::DT_FLOAT, ge::FORMAT_ND},
	65	+ {{{1, 7, 4096}, {1, 7, 4096}}, ge::DT_FLOAT, ge::FORMAT_ND},
	66	+ },
	67	+ {
	68	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	69	+ });
	70	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	71	+ {1, 64, 4096},
	72	+ };
	73	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	74	+}
	75	+
	76	+// 正例：unknownRank（输入 rank 未知 -> 输出置 unknownRank）。仅校验返回 SUCCESS，
	77	+// 不断言输出维值（unknownRank 分支只 SetDimNum 不 SetDim，维值未定义）。
	78	+TEST_F(Iou3DInfershape, iou3d_infershape_unknown_rank)
	79	+{
	80	+ gert::InfershapeContextPara infershapeContextPara("Iou3D",
	81	+ {
	82	+ {{{-2}, {-2}}, ge::DT_FLOAT, ge::FORMAT_ND},
	83	+ {{{-2}, {-2}}, ge::DT_FLOAT, ge::FORMAT_ND},
	84	+ },
	85	+ {
	86	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	87	+ });
	88	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS);
	89	+}
	90	+
	91	+// 反例：channel != 7（bboxes dim1=5）-> FAILED
	92	+TEST_F(Iou3DInfershape, iou3d_infershape_neg_channel_not_7)
	93	+{
	94	+ gert::InfershapeContextPara infershapeContextPara("Iou3D",
	95	+ {
	96	+ {{{2, 5, 8}, {2, 5, 8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	97	+ {{{2, 7, 16}, {2, 7, 16}}, ge::DT_FLOAT, ge::FORMAT_ND},
	98	+ },
	99	+ {
	100	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	101	+ });
	102	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED);
	103	+}
	104	+
	105	+// 反例：rank != 3（bboxes 为 rank-2）-> FAILED
	106	+TEST_F(Iou3DInfershape, iou3d_infershape_neg_rank_not_3)
	107	+{
	108	+ gert::InfershapeContextPara infershapeContextPara("Iou3D",
	109	+ {
	110	+ {{{2, 7}, {2, 7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	111	+ {{{2, 7, 16}, {2, 7, 16}}, ge::DT_FLOAT, ge::FORMAT_ND},
	112	+ },
	113	+ {
	114	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	115	+ });
	116	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED);
	117	+}
	118	+
	119	+// 反例：batch 不一致（bboxes B=2, gtboxes B=3）-> FAILED
	120	+TEST_F(Iou3DInfershape, iou3d_infershape_neg_batch_mismatch)
	121	+{
	122	+ gert::InfershapeContextPara infershapeContextPara("Iou3D",
	123	+ {
	124	+ {{{2, 7, 8}, {2, 7, 8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	125	+ {{{3, 7, 16}, {3, 7, 16}}, ge::DT_FLOAT, ge::FORMAT_ND},
	126	+ },
	127	+ {
	128	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	129	+ });
	130	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED);
	131	+}

Aobjdetect/iou3d/tests/ut/op_kernel/CMakeLists.txt+21-0

@@ -0,0 +1,21 @@
	1	+# -----------------------------------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# -----------------------------------------------------------------------------------------------------------
	10	+
	11	+if (UT_TEST_ALL OR OP_KERNEL_UT)
	12	+ set(iou3d_tiling_files
	13	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/iou3d_tiling_arch35.cpp
	14	+ )
	15	+ # 使用AddOpTestCase
	16	+ # param1：算子名称，以kernel方式命名
	17	+ # param2：soc版本，多个以分号分隔，例如："ascend950PR_9599;AscendB1"
	18	+ # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔
	19	+ # param4：该算子依赖的所有tiling源码文件
	20	+ AddOpTestCase(iou3d "Ascend950PR_9599" "-DDTYPE_BBOXES=float" "${iou3d_tiling_files}")
	21	+endif()

Aobjdetect/iou3d/tests/ut/op_kernel/iou3d_tiling.h+62-0

@@ -0,0 +1,62 @@
	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 iou3d_tiling.h
	13	+ * \brief Iou3D kernel UT 的 tiling shim（CPU 仿真侧）。
	14	+ *
	15	+ * AddOpTestCase 宏会自动 `-include` 本文件（若存在），用于在 tikicpulib CPU 仿真下替换
	16	+ * kernel 里的 REGISTER_TILING_DEFAULT / GET_TILING_DATA_WITH_STRUCT 宏：真机侧这两个宏
	17	+ * 从 __gm__ 读 tiling 并做寄存器化，CPU 仿真侧改为 memcpy 到栈上结构体。
	18	+ * 结构体定义复用算子真实 op_kernel/arch35/iou3d_tiling_data.h（单一真值源，不重定义字段）。
	19	+ */
	20	+
	21	+#ifndef IOU3D_KERNEL_UT_TILING_H
	22	+#define IOU3D_KERNEL_UT_TILING_H
	23	+
	24	+#include <cstdint>
	25	+#include <cstring>
	26	+
	27	+#include "../../../op_kernel/arch35/iou3d_tiling_data.h"
	28	+#include "kernel_tiling/kernel_tiling.h"
	29	+
	30	+#ifdef __NPU_TILING__
	31	+inline[aicore] void InitTilingData(const __gm__ uint8_t* tiling, Iou3DTilingData* constData)
	32	+{
	33	+ const __gm__ uint32_t* src = (const __gm__ uint32_t*)tiling;
	34	+ uint32_t* dst = (uint32_t*)constData;
	35	+ for (auto i = 0; i < sizeof(Iou3DTilingData) / 4; i++)
	36	+ (dst + i) = (src + i);
	37	+}
	38	+#else
	39	+inline void InitTilingData(uint8_t* tiling, Iou3DTilingData* constData)
	40	+{
	41	+ memcpy(constData, tiling, sizeof(Iou3DTilingData));
	42	+}
	43	+#endif // __NPU_TILING__
	44	+
	45	+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
	46	+ __ubuf__ tilingStruct* tilingDataPointer = reinterpret_cast<__ubuf__ tilingStruct*>( \
	47	+ (__ubuf__ uint8_t*)(tilingPointer));
	48	+
	49	+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
	50	+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
	51	+
	52	+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingArg) \
	53	+ tilingStruct tilingData; \
	54	+ InitTilingData(tilingArg, &tilingData)
	55	+
	56	+#define GET_TILING_DATA(tilingData, tilingArg) \
	57	+ Iou3DTilingData tilingData; \
	58	+ InitTilingData(tilingArg, &tilingData)
	59	+
	60	+#define REGISTER_TILING_DEFAULT(T)
	61	+
	62	+#endif // IOU3D_KERNEL_UT_TILING_H

Aobjdetect/iou3d/tests/ut/op_kernel/test_iou3d.cpp+200-0

@@ -0,0 +1,200 @@
	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 test_iou3d.cpp
	13	+ * \brief Iou3D kernel 执行 UT（arch35 / tikicpulib CPU 仿真）。
	14	+ *
	15	+ * 定位：验证 kernel 主路径在 CPU 仿真下可执行不崩溃，并对典型锚点做宽松数值 sanity check
	16	+ * （非逐元素精度比对——逐元素精度由 assets/golden.py + 外部 ST/TTK 框架在真机上把关）。
	17	+ * 覆盖：
	18	+ * - identical box（IoU≈1）+ disjoint box（IoU≈0）：走 Sin/Cos 向量批算 + 顶点收集 + 面积。
	19	+ * - 旋转相交 box：触发 >3 交点的 Sort32 极角排序分支。
	20	+ * - 空 Tensor（isEmpty=1）：Process 运行时短路，无输出写。
	21	+ */
	22	+#include <array>
	23	+#include <vector>
	24	+#include <iostream>
	25	+#include <string>
	26	+#include <cstdint>
	27	+#include <cstring>
	28	+#include <cmath>
	29	+#include <cstdlib>
	30	+#include "data_utils.h"
	31	+#include "gtest/gtest.h"
	32	+#include "tikicpulib.h"
	33	+#include "kernel_tiling/kernel_tiling.h"
	34	+#include "../../../op_kernel/arch35/iou3d_tiling_data.h"
	35	+
	36	+using namespace std;
	37	+
	38	+// kernel 入口符号：OpType(Iou3D) 的 snake_case 形式 iou3_d（框架 CamelCase→snake_case 约定）。
	39	+// 仅前向声明，与 op_kernel/iou3d.cpp 的定义一致（C++ linkage，无 extern "C"）。
	40	+// kernel .cpp 由 AddOpTestCase 宏（cmake/ut.cmake）自动编入同一 case object，勿在此 #include 以免重复定义。
	41	+__global__ __aicore__ void iou3_d(GM_ADDR bboxes, GM_ADDR gtboxes, GM_ADDR iou, GM_ADDR workspace, GM_ADDR tiling);
	42	+
	43	+class iou3d_test : public testing::Test {
	44	+protected:
	45	+ static void SetUpTestCase() { cout << "iou3d SetUp\n" << endl; }
	46	+ static void TearDownTestCase() { cout << "iou3d_test TearDown\n" << endl; }
	47	+};
	48	+
	49	+namespace {
	50	+// 写入一个框的 7-DoF 到 [B,7,D] 布局：base=b7D + c*D + idx。
	51	+void SetBox(float* gm, int64_t b, int64_t idx, int64_t dimSize, const std::array<float, 7>& box)
	52	+{
	53	+ for (int c = 0; c < 7; ++c) {
	54	+ gm[b * 7 * dimSize + static_cast<int64_t>(c) * dimSize + idx] = box[c];
	55	+ }
	56	+}
	57	+
	58	+// 单核填 tiling：所有 (b,i,j) 对由 1 个核处理。
	59	+void FillTiling(Iou3DTilingData* t, uint32_t batch, uint32_t numN, uint32_t numK, uint32_t isEmpty)
	60	+{
	61	+ memset(t, 0, sizeof(Iou3DTilingData));
	62	+ t->batch = batch;
	63	+ t->numBboxes = numN;
	64	+ t->numGtboxes = numK;
	65	+ t->isEmpty = isEmpty;
	66	+ if (isEmpty != 0U) {
	67	+ t->coreNum = 1;
	68	+ return;
	69	+ }
	70	+ const uint32_t total = batch * numN * numK;
	71	+ t->coreNum = 1;
	72	+ t->pairsPerCore = total;
	73	+ t->tileLen = total < 256U ? total : 256U;
	74	+ t->tailLen = (t->tileLen == 0U) ? 0U : (total % t->tileLen);
	75	+}
	76	+} // namespace
	77	+
	78	+// 正常框对：[1,7,2] × [1,7,2]。box0 identical(IoU≈1)，box1 与 box0 不相交(IoU≈0)。
	79	+TEST_F(iou3d_test, test_iou3d_basic_fp32)
	80	+{
	81	+ const uint32_t B = 1, N = 2, K = 2;
	82	+ size_t bboxesByteSize = B * 7 * N * sizeof(float);
	83	+ size_t gtboxesByteSize = B * 7 * K * sizeof(float);
	84	+ size_t iouByteSize = B * N * K * sizeof(float);
	85	+ size_t tilingDataSize = sizeof(Iou3DTilingData);
	86	+
	87	+ uint8_t* bboxes = (uint8_t*)AscendC::GmAlloc(bboxesByteSize);
	88	+ uint8_t* gtboxes = (uint8_t*)AscendC::GmAlloc(gtboxesByteSize);
	89	+ uint8_t* iou = (uint8_t*)AscendC::GmAlloc(iouByteSize);
	90	+ uint8_t* workspace = (uint8_t)AscendC::GmAlloc(1024 16 * 1024);
	91	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);
	92	+ uint32_t numBlocks = 1;
	93	+
	94	+ // 输入布局 [B,7,N]：box0 在原点、box1 远离。gtbox 同 bbox。
	95	+ float* pb = reinterpret_cast<float*>(bboxes);
	96	+ float* pg = reinterpret_cast<float*>(gtboxes);
	97	+ memset(pb, 0, bboxesByteSize);
	98	+ memset(pg, 0, gtboxesByteSize);
	99	+ // [x,y,z,w,h,d,theta]
	100	+ SetBox(pb, 0, 0, N, {0.0f, 0.0f, 0.0f, 2.0f, 2.0f, 2.0f, 0.0f});
	101	+ SetBox(pb, 0, 1, N, {100.0f, 100.0f, 0.0f, 2.0f, 2.0f, 2.0f, 0.0f});
	102	+ SetBox(pg, 0, 0, K, {0.0f, 0.0f, 0.0f, 2.0f, 2.0f, 2.0f, 0.0f});
	103	+ SetBox(pg, 0, 1, K, {100.0f, 100.0f, 0.0f, 2.0f, 2.0f, 2.0f, 0.0f});
	104	+
	105	+ float* pi = reinterpret_cast<float*>(iou);
	106	+ memset(pi, 0xFF, iouByteSize); // 预置非零，验证 kernel 确实写回
	107	+
	108	+ Iou3DTilingData* tilingData = reinterpret_cast<Iou3DTilingData*>(tiling);
	109	+ FillTiling(tilingData, B, N, K, 0U);
	110	+
	111	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	112	+ ICPU_RUN_KF(iou3_d, numBlocks, bboxes, gtboxes, iou, workspace, tiling);
	113	+
	114	+ // 宽松 sanity check：对角（identical）≈1，非对角（disjoint）≈0。
	115	+ EXPECT_NEAR(pi[0 * K + 0], 1.0f, 1e-3f); // (i=0,j=0)
	116	+ EXPECT_NEAR(pi[1 * K + 1], 1.0f, 1e-3f); // (i=1,j=1)
	117	+ EXPECT_NEAR(pi[0 * K + 1], 0.0f, 1e-3f); // (i=0,j=1) disjoint
	118	+ EXPECT_NEAR(pi[1 * K + 0], 0.0f, 1e-3f); // (i=1,j=0) disjoint
	119	+
	120	+ AscendC::GmFree(bboxes);
	121	+ AscendC::GmFree(gtboxes);
	122	+ AscendC::GmFree(iou);
	123	+ AscendC::GmFree(workspace);
	124	+ AscendC::GmFree(tiling);
	125	+}
	126	+
	127	+// 旋转相交框对：[1,7,1] × [1,7,1]，theta 差 45° -> 交集为八边形（>3 顶点，触发 Sort32 分支）。
	128	+TEST_F(iou3d_test, test_iou3d_rotated_intersection_fp32)
	129	+{
	130	+ const uint32_t B = 1, N = 1, K = 1;
	131	+ size_t bboxesByteSize = B * 7 * N * sizeof(float);
	132	+ size_t gtboxesByteSize = B * 7 * K * sizeof(float);
	133	+ size_t iouByteSize = B * N * K * sizeof(float);
	134	+ size_t tilingDataSize = sizeof(Iou3DTilingData);
	135	+
	136	+ uint8_t* bboxes = (uint8_t*)AscendC::GmAlloc(bboxesByteSize);
	137	+ uint8_t* gtboxes = (uint8_t*)AscendC::GmAlloc(gtboxesByteSize);
	138	+ uint8_t* iou = (uint8_t*)AscendC::GmAlloc(iouByteSize);
	139	+ uint8_t* workspace = (uint8_t)AscendC::GmAlloc(1024 16 * 1024);
	140	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);
	141	+ uint32_t numBlocks = 1;
	142	+
	143	+ float* pb = reinterpret_cast<float*>(bboxes);
	144	+ float* pg = reinterpret_cast<float*>(gtboxes);
	145	+ memset(pb, 0, bboxesByteSize);
	146	+ memset(pg, 0, gtboxesByteSize);
	147	+ SetBox(pb, 0, 0, N, {0.0f, 0.0f, 0.0f, 2.0f, 2.0f, 2.0f, 0.0f});
	148	+ SetBox(pg, 0, 0, K, {0.0f, 0.0f, 0.0f, 2.0f, 2.0f, 2.0f, static_cast<float>(M_PI / 4.0)});
	149	+
	150	+ float* pi = reinterpret_cast<float*>(iou);
	151	+ memset(pi, 0, iouByteSize);
	152	+
	153	+ Iou3DTilingData* tilingData = reinterpret_cast<Iou3DTilingData*>(tiling);
	154	+ FillTiling(tilingData, B, N, K, 0U);
	155	+
	156	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	157	+ ICPU_RUN_KF(iou3_d, numBlocks, bboxes, gtboxes, iou, workspace, tiling);
	158	+
	159	+ // 45° 旋转的等尺寸同心框：BEV 交集为正八边形，3D IoU≈0.707（宽松容差，主要验证不崩溃）。
	160	+ EXPECT_GT(pi[0], 0.5f);
	161	+ EXPECT_LE(pi[0], 1.0f + 1e-3f);
	162	+
	163	+ AscendC::GmFree(bboxes);
	164	+ AscendC::GmFree(gtboxes);
	165	+ AscendC::GmFree(iou);
	166	+ AscendC::GmFree(workspace);
	167	+ AscendC::GmFree(tiling);
	168	+}
	169	+
	170	+// 空 Tensor：N=0 -> isEmpty=1，Process 运行时短路，不写输出（不崩溃即通过）。
	171	+TEST_F(iou3d_test, test_iou3d_empty_tensor)
	172	+{
	173	+ const uint32_t B = 1, N = 0, K = 4;
	174	+ size_t bboxesByteSize = 1 * 7 * 1 * sizeof(float); // N=0 无法分配 0 字节，占位 1
	175	+ size_t gtboxesByteSize = B * 7 * K * sizeof(float);
	176	+ size_t iouByteSize = 1 * sizeof(float); // 空输出占位
	177	+ size_t tilingDataSize = sizeof(Iou3DTilingData);
	178	+
	179	+ uint8_t* bboxes = (uint8_t*)AscendC::GmAlloc(bboxesByteSize);
	180	+ uint8_t* gtboxes = (uint8_t*)AscendC::GmAlloc(gtboxesByteSize);
	181	+ uint8_t* iou = (uint8_t*)AscendC::GmAlloc(iouByteSize);
	182	+ uint8_t* workspace = (uint8_t)AscendC::GmAlloc(1024 16 * 1024);
	183	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);
	184	+ uint32_t numBlocks = 1;
	185	+
	186	+ memset(bboxes, 0, bboxesByteSize);
	187	+ memset(gtboxes, 0, gtboxesByteSize);
	188	+
	189	+ Iou3DTilingData* tilingData = reinterpret_cast<Iou3DTilingData*>(tiling);
	190	+ FillTiling(tilingData, B, N, K, 1U);
	191	+
	192	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	193	+ ICPU_RUN_KF(iou3_d, numBlocks, bboxes, gtboxes, iou, workspace, tiling);
	194	+
	195	+ AscendC::GmFree(bboxes);
	196	+ AscendC::GmFree(gtboxes);
	197	+ AscendC::GmFree(iou);
	198	+ AscendC::GmFree(workspace);
	199	+ AscendC::GmFree(tiling);
	200	+}