从ops-cv迁移ROIPooling到ops-nn-ops-nn-AtomGit

Mcommon/inc/op_graph/op_nn_proto_extend.h+0-67

	@@ -2400,73 +2400,6 @@ Reserved.
2400	.ATTR(epsilon, Float, 1e-12f)	2400	.ATTR(epsilon, Float, 1e-12f)
2401	.OP_END_FACTORY_REG(LpNormUpdateV2)	2401	.OP_END_FACTORY_REG(LpNormUpdateV2)
2402		2402
2403	- /**
2404	- *@brief Performs Region of Interest (ROI) Pooling . \n
2405	-
2406	- *@par Inputs:
2407	- * Three inputs, including:
2408	- *@li x: A tensor of type float16 or float32, describing the feature
2409	- * map. The data of x must be greater than or equal to "0.0".
2410	- *@li rois: A tensor of type float16 or float32, with 3D shape
2411	- * [batch, 5, roi_max_num], describing the RIOs. Each ROI consists of five
2412	- * elements: "batch_id", "x1", "y1", "x2", and "y2", which "batch_id" indicates
2413	- * the index of the input feature map, "x1", "y1", "x2", or "y2" must be
2414	- * greater than or equal to "0.0".
2415	- * roi_max_num must be less than or equal to 6000 and must be divided by 16.
2416	- * The input data of the rois cannot exceed the width and height range of the x,
2417	- * otherwise, the accuracy of the output result may not be as expected.
2418	- *@li roi_actual_num: A optional tensor of type int32, with shape [batch, 8], specifying
2419	- * the number of ROIs per batch . \n
2420	-
2421	- *@par Attributes:
2422	- *@li pooled_h: A required int32, specifying the pooled H. Must be greater
2423	- * than 0.
2424	- *@li pooled_w: A required int32, specifying the pooled W. Must be greater
2425	- * than 0.
2426	- *@li spatial_scale_h: An required scaling factor for mapping the input
2427	- * coordinates of height to the ROI coordinates.
2428	- *@li spatial_scale_w: An required scaling factor for mapping the input
2429	- * coordinates of width to the ROI coordinates . \n
2430	-
2431	- *@par Outputs:
2432	- *y: A tensor of type float16 or float32, describing the result
2433	- * feature map . \n
2434	-
2435	- *@attention Constraints:
2436	- * For the feature map input:
2437	- *@li If pooled_h = pooled_w = 2, the feature map size must not exceed 50.
2438	- *@li If pooled_h = pooled_w = 3, the feature map size must not exceed 60.
2439	- *@li If pooled_h = pooled_w = 4, the feature map size must not exceed 70.
2440	- *@li If pooled_h = pooled_w = 5, the feature map size must not exceed 70.
2441	- *@li If pooled_h = pooled_w = 6, the feature map size must not exceed 80.
2442	- *@li If pooled_h = pooled_w = 7, the feature map size must not exceed 80.
2443	- *@li If pooled_h = pooled_w = 8, the feature map size must not exceed 80.
2444	- *@li If pooled_h = pooled_w = 9, the feature map size must not exceed 70.
2445	- *@li If pooled_h = pooled_w = 10, the feature map size must not exceed 70.
2446	- *@li If pooled_h = pooled_w = 11, the feature map size must not exceed 70.
2447	- *@li If pooled_h = pooled_w = 12, the feature map size must not exceed 70.
2448	- *@li If pooled_h = pooled_w = 13, the feature map size must not exceed 70.
2449	- *@li If pooled_h = pooled_w = 14, the feature map size must not exceed 70.
2450	- *@li If pooled_h = pooled_w = 15, the feature map size must not exceed 70.
2451	- *@li If pooled_h = pooled_w = 16, the feature map size must not exceed 70.
2452	- *@li If pooled_h = pooled_w = 17, the feature map size must not exceed 50.
2453	- *@li If pooled_h = pooled_w = 18, the feature map size must not exceed 40.
2454	- *@li If pooled_h = pooled_w = 19, the feature map size must not exceed 40.
2455	- *@li If pooled_h = pooled_w = 20, the feature map size must not exceed 40.
2456	- *@par Third-party framework compatibility
2457	- * It is a custom operator. It has no corresponding operator in Caffe.
2458	- */
2459	- REG_OP(ROIPooling)
2460	- .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16}))
2461	- .INPUT(rois, TensorType({DT_FLOAT, DT_FLOAT16}))
2462	- .OPTIONAL_INPUT(roi_actual_num, TensorType({DT_INT32}))
2463	- .REQUIRED_ATTR(pooled_h, Int)
2464	- .REQUIRED_ATTR(pooled_w, Int)
2465	- .REQUIRED_ATTR(spatial_scale_h, Float)
2466	- .REQUIRED_ATTR(spatial_scale_w, Float)
2467	- .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16}))
2468	- .OP_END_FACTORY_REG(ROIPooling)
2469	-
2470	/**	2403	/**
2471	* @brief Computes second-order gradients of the maxpooling function .	2404	* @brief Computes second-order gradients of the maxpooling function .
2472		2405

Mdocs/zh/op_list.md+10-0

	@@ -4738,6 +4738,16 @@
4738	<td>AI Core</td>	4738	<td>AI Core</td>
4739	<td>对输入特征图执行位置敏感RoI平均池化，适用于R-FCN等目标检测网络。</td>	4739	<td>对输入特征图执行位置敏感RoI平均池化，适用于R-FCN等目标检测网络。</td>
4740	</tr>	4740	</tr>
		4741	+ <tr>
		4742	+ <td>pooling</td>
		4743	+ <td><a href="../../pooling/roi_pooling/README.md">roi_pooling</a></td>
		4744	+ <td>✓</td>
		4745	+ <td>✓</td>
		4746	+ <td>✗</td>
		4747	+ <td>✓</td>
		4748	+ <td>AI Core</td>
		4749	+ <td>对输入特征图按ROI区域进行最大池化，输出固定尺寸的池化结果，用于目标检测等任务。</td>
		4750	+ </tr>
4741	<tr>	4751	<tr>
4742	<td>quant</td>	4752	<td>quant</td>
4743	<td><a href="../../quant/act_ulq_clamp_max_grad/README.md">act_ulq_clamp_max_grad</a></td>	4753	<td><a href="../../quant/act_ulq_clamp_max_grad/README.md">act_ulq_clamp_max_grad</a></td>

Apooling/roi_pooling/CMakeLists.txt+19-0

@@ -0,0 +1,19 @@
	1	+# ----------------------------------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the License).
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------------------------------------
	10	+# Generated By CANNBot
	11	+
	12	+# 设置算子定义时支持的芯片类型
	13	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	14	+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
	15	+set(SUPPORT_TILING_DIR "arch35")
	16	+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}
	17	+ OPTYPE roi_pooling ACLNNTYPE aclnn_exclude
	18	+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR}
	19	+ DISABLE_IN_OPP TRUE)

Apooling/roi_pooling/README.md+163-0

@@ -0,0 +1,163 @@
	1	+# ROIPooling
	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	+- 算子功能：对输入特征图按ROI（Region of Interest）区域进行最大池化，输出固定尺寸的池化结果，用于目标检测等任务。
	17	+
	18	+- 计算公式：
	19	+
	20	+对每个ROI n（rois[n] = [batch_idx, x1, y1, x2, y2]）、通道c、池化位置(ph,pw)：
	21	+
	22	+1. ROI坐标映射（使用分离的h/w缩放因子，roundf取整为int，终点坐标无+1偏移）：
	23	+
	24	+$$
	25	+\text{roi\_start}_w = \text{round}(x1 \cdot \text{spatial\_scale\_w}), \quad \text{roi\_start}_h = \text{round}(y1 \cdot \text{spatial\_scale\_h})
	26	+$$
	27	+
	28	+$$
	29	+\text{roi\_end}_w = \text{round}(x2 \cdot \text{spatial\_scale\_w}), \quad \text{roi\_end}_h = \text{round}(y2 \cdot \text{spatial\_scale\_h})
	30	+$$
	31	+
	32	+> 以上坐标均为整数（int类型），+1偏移在ROI尺寸上（见下），不在坐标上。
	33	+
	34	+2. ROI尺寸（int运算，+1在宽度/高度上，malformed ROI强制非空）：
	35	+
	36	+$$
	37	+\text{roi\_w} = \max(\text{roi\_end}_w - \text{roi\_start}_w + 1, 1), \quad \text{roi\_h} = \max(\text{roi\_end}_h - \text{roi\_start}_h + 1, 1)
	38	+$$
	39	+
	40	+3. Bin切分与边界裁剪（bin_size为float，基于int的roi_w/roi_h）：
	41	+
	42	+$$
	43	+\text{bin\_size}_w = \frac{\text{roi\_w}}{\text{pooled\_w}}, \quad \text{bin\_size}_h = \frac{\text{roi\_h}}{\text{pooled\_h}}
	44	+$$
	45	+
	46	+$$
	47	+\text{bin}_{x1} = \text{clamp}(\lfloor pw \cdot \text{bin\_size}_w \rfloor + \text{roi\_start}_w, 0, W)
	48	+$$
	49	+
	50	+$$
	51	+\text{bin}_{x2} = \text{clamp}(\lceil (pw+1) \cdot \text{bin\_size}_w \rceil + \text{roi\_start}_w, 0, W)
	52	+$$
	53	+
	54	+（y方向同理，clamp到[0, H]）
	55	+
	56	+> floor/ceil基于相对偏移（不含roi_start），再加整数roi_start。
	57	+
	58	+4. Max Pooling（空bin输出0）：
	59	+
	60	+$$
	61	+y[n, c, ph, pw] = \begin{cases} 0 & \text{if } \text{bin}_{x2} \le \text{bin}_{x1} \text{ or } \text{bin}_{y2} \le \text{bin}_{y1} \\ \max_{h, w \in [\text{bin}_{y1}, \text{bin}_{y2}) \times [\text{bin}_{x1}, \text{bin}_{x2})} x[\text{batch\_idx}, c, h, w] & \text{otherwise} \end{cases}
	62	+$$
	63	+
	64	+其中+1偏移在ROI尺寸上（`roi_width = roi_end - roi_start + 1`），是ROI Pooling通用标准（源自Fast R-CNN）
	65	+
	66	+## 参数说明
	67	+
	68	+<table><thead>
	69	+ <tr>
	70	+ <th>参数名</th>
	71	+ <th>输入/输出/属性</th>
	72	+ <th>描述</th>
	73	+ <th>数据类型</th>
	74	+ <th>数据格式</th>
	75	+ </tr></thead>
	76	+<tbody>
	77	+ <tr>
	78	+ <td>x</td>
	79	+ <td>输入</td>
	80	+ <td>特征图，shape为[N,C,H,W]，公式中的x。</td>
	81	+ <td>FLOAT、FLOAT16</td>
	82	+ <td>ND</td>
	83	+ </tr>
	84	+ <tr>
	85	+ <td>rois</td>
	86	+ <td>输入</td>
	87	+ <td>ROI框，shape为[K,5]，每行为[batch_idx,x1,y1,x2,y2]。</td>
	88	+ <td>FLOAT、FLOAT16</td>
	89	+ <td>ND</td>
	90	+ </tr>
	91	+ <tr>
	92	+ <td>roi_actual_num</td>
	93	+ <td>输入</td>
	94	+ <td>每个batch的实际ROI数量，可选输入，当前版本不参与计算。</td>
	95	+ <td>INT32</td>
	96	+ <td>ND</td>
	97	+ </tr>
	98	+ <tr>
	99	+ <td>y</td>
	100	+ <td>输出</td>
	101	+ <td>池化结果，shape为[K,C,pooled_h,pooled_w]，公式中的y。</td>
	102	+ <td>FLOAT、FLOAT16</td>
	103	+ <td>ND</td>
	104	+ </tr>
	105	+ <tr>
	106	+ <td>pooled_h</td>
	107	+ <td>属性</td>
	108	+ <td>池化输出高度。</td>
	109	+ <td>Int</td>
	110	+ <td>-</td>
	111	+ </tr>
	112	+ <tr>
	113	+ <td>pooled_w</td>
	114	+ <td>属性</td>
	115	+ <td>池化输出宽度。</td>
	116	+ <td>Int</td>
	117	+ <td>-</td>
	118	+ </tr>
	119	+ <tr>
	120	+ <td>spatial_scale_h</td>
	121	+ <td>属性</td>
	122	+ <td>高度方向缩放因子。</td>
	123	+ <td>Float</td>
	124	+ <td>-</td>
	125	+ </tr>
	126	+ <tr>
	127	+ <td>spatial_scale_w</td>
	128	+ <td>属性</td>
	129	+ <td>宽度方向缩放因子。</td>
	130	+ <td>Float</td>
	131	+ <td>-</td>
	132	+ </tr>
	133	+</tbody></table>
	134	+
	135	+## 约束说明
	136	+
	137	+- 输入x仅支持4维[N,C,H,W]。
	138	+- 输入rois仅支持2维[K,5]，每行为[batch_idx,x1,y1,x2,y2]。
	139	+- roi_actual_num为可选输入，当前版本不参与计算。
	140	+- pooled_h和pooled_w必须大于0。
	141	+- spatial_scale_h和spatial_scale_w必须大于0。
	142	+- x与rois的dtype必须一致。
	143	+- 输出y的shape为[K,C,pooled_h,pooled_w]。
	144	+- ROI坐标用roundf取整为int，+1偏移在ROI尺寸上（roi_width = roi_end - roi_start + 1）。
	145	+- malformed ROI强制非空，roi_size取max(size,1)（int运算）。
	146	+- 空bin输出0。
	147	+
	148	+## 调用说明
	149	+
	150	+<table><thead>
	151	+ <tr>
	152	+ <th>调用方式</th>
	153	+ <th>调用样例</th>
	154	+ <th>说明</th>
	155	+ </tr></thead>
	156	+<tbody>
	157	+ <tr>
	158	+ <td>图模式调用</td>
	159	+ <td><a href="./examples/test_geir_roi_pooling.cpp">test_geir_roi_pooling</a></td>
	160	+ <td>参见<a href="../../docs/zh/invocation/quick_op_invocation.md">算子调用</a>完成算子编译和验证。</td>
	161	+ </tr>
	162	+</tbody>
	163	+</table>

Apooling/roi_pooling/examples/test_geir_roi_pooling.cpp+385-0

@@ -0,0 +1,385 @@
	1	+
	2	+/**
	3	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	5	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	6	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	8	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+ * See LICENSE in the root of the software repository for the full text of the License.
	10	+ */
	11	+
	12	+/* Generated By CANNBot */
	13	+
	14	+#include <cstddef>
	15	+#include <cstdint>
	16	+#include <iostream>
	17	+#include <fstream>
	18	+#include <string.h>
	19	+#include <stdint.h>
	20	+#include <vector>
	21	+#include <string>
	22	+#include <map>
	23	+#include "assert.h"
	24	+
	25	+#include "graph.h"
	26	+#include "types.h"
	27	+#include "tensor.h"
	28	+#include "ge_error_codes.h"
	29	+#include "ge_api_types.h"
	30	+#include "ge_api.h"
	31	+#include "array_ops.h"
	32	+#include "ge_ir_build.h"
	33	+
	34	+#include "../op_graph/roi_pooling_proto.h"
	35	+
	36	+#define FAILED -1
	37	+#define SUCCESS 0
	38	+
	39	+using namespace ge;
	40	+using std::map;
	41	+using std::string;
	42	+using std::vector;
	43	+
	44	+enum RunMode { RUN_MODE_S = 0, RUN_MODE_D = 1 };
	45	+
	46	+struct CaseResult {
	47	+ std::string case_name;
	48	+ bool build_ok;
	49	+ bool run_ok;
	50	+ bool output_exists;
	51	+ int output_count;
	52	+ std::string err_msg;
	53	+};
	54	+
	55	+// ROI Pooling 输入 shape 各不相同且存在依赖关系：
	56	+// x : [N, C, H, W] 4D 特征图
	57	+// rois : [K, 5] 2D ROI 列表，每行 [batch_idx, x1, y1, x2, y2]
	58	+// y : [K, C, pooled_h, pooled_w] 4D 输出
	59	+// 因此不能共用单个 xShape，需用 ShapeCombo 结构分别为各输入构造合法 shape。
	60	+struct ShapeCombo {
	61	+ std::vector<int64_t> xShape;
	62	+ std::vector<int64_t> roisShape;
	63	+ std::string name;
	64	+};
	65	+
	66	+#define ADD_INPUT_MODE(intputIndex, intputName, intputDtype, inputShape, mode) \
	67	+ vector<int64_t> placeholder##intputIndex##_real_shape = inputShape; \
	68	+ vector<int64_t> placeholder##intputIndex##_graph_shape = ((mode) == RUN_MODE_D) ? \
	69	+ vector<int64_t>( \
	70	+ placeholder##intputIndex##_real_shape.size(), \
	71	+ -1) : \
	72	+ placeholder##intputIndex##_real_shape; \
	73	+ auto placeholder##intputIndex = op::Data("placeholder" #intputIndex).set_attr_index(0); \
	74	+ TensorDesc placeholder##intputIndex##_desc_graph = TensorDesc(ge::Shape(placeholder##intputIndex##_graph_shape), \
	75	+ FORMAT_ND, intputDtype); \
	76	+ placeholder##intputIndex##_desc_graph.SetPlacement(ge::kPlacementHost); \
	77	+ placeholder##intputIndex##_desc_graph.SetFormat(FORMAT_ND); \
	78	+ TensorDesc placeholder##intputIndex##_desc_real = TensorDesc(ge::Shape(placeholder##intputIndex##_real_shape), \
	79	+ FORMAT_ND, intputDtype); \
	80	+ placeholder##intputIndex##_desc_real.SetPlacement(ge::kPlacementHost); \
	81	+ placeholder##intputIndex##_desc_real.SetFormat(FORMAT_ND); \
	82	+ placeholder##intputIndex##_desc_real.SetRealDimCnt(placeholder##intputIndex##_real_shape.size()); \
	83	+ Tensor tensor_placeholder##intputIndex; \
	84	+ ret = GenOnesData(placeholder##intputIndex##_real_shape, tensor_placeholder##intputIndex, \
	85	+ placeholder##intputIndex##_desc_real, intputDtype, 2); \
	86	+ if (ret != SUCCESS) { \
	87	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	88	+ return FAILED; \
	89	+ } \
	90	+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc_graph); \
	91	+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc_graph); \
	92	+ input.push_back(tensor_placeholder##intputIndex); \
	93	+ graph.AddOp(placeholder##intputIndex); \
	94	+ add1.set_input_##intputName(placeholder##intputIndex); \
	95	+ inputs.push_back(placeholder##intputIndex);
	96	+
	97	+#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape) \
	98	+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
	99	+ auto placeholder##intputIndex = op::Const("placeholder" #intputIndex); \
	100	+ TensorDesc placeholder##intputIndex##_desc = TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, \
	101	+ intputDtype); \
	102	+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
	103	+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
	104	+ Tensor tensor_placeholder##intputIndex; \
	105	+ ret = GenOnesData(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \
	106	+ placeholder##intputIndex##_desc, intputDtype, 2); \
	107	+ if (ret != SUCCESS) { \
	108	+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
	109	+ return FAILED; \
	110	+ } \
	111	+ placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \
	112	+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
	113	+ graph.AddOp(placeholder##intputIndex); \
	114	+ add1.set_input_##intputName(placeholder##intputIndex); \
	115	+ add1.update_input_desc_##intputName(placeholder##intputIndex##_desc); \
	116	+ inputs.push_back(placeholder##intputIndex);
	117	+
	118	+// ADD_OUTPUT_MODE 内部封装了 S/D 模式 graph_shape 自动生成（D 模式含 -1）
	119	+#define ADD_OUTPUT_MODE(outputIndex, outputName, outputDtype, outputShape, mode) \
	120	+ vector<int64_t> output##outputIndex##_graph_shape = ((mode) == RUN_MODE_D) ? \
	121	+ vector<int64_t>(outputShape.size(), -1) : \
	122	+ outputShape; \
	123	+ TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(output##outputIndex##_graph_shape), FORMAT_ND, \
	124	+ outputDtype); \
	125	+ add1.update_output_desc_##outputName(outputName##outputIndex##_desc);
	126	+
	127	+string GetTime()
	128	+{
	129	+ time_t timep;
	130	+ time(&timep);
	131	+ char tmp[64];
	132	+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
	133	+ return tmp;
	134	+}
	135	+
	136	+uint32_t GetDataTypeSize(DataType dt)
	137	+{
	138	+ uint32_t dilation = 1;
	139	+ uint32_t oneByte = 1;
	140	+ uint32_t twoByte = 2;
	141	+ uint32_t fourByte = 4;
	142	+ uint32_t eightByte = 8;
	143	+
	144	+ if (dt == ge::DT_FLOAT) {
	145	+ dilation = fourByte;
	146	+ } else if (dt == ge::DT_FLOAT16) {
	147	+ dilation = twoByte;
	148	+ } else if (dt == ge::DT_BF16) {
	149	+ dilation = twoByte;
	150	+ } else if (dt == ge::DT_INT16) {
	151	+ dilation = twoByte;
	152	+ } else if (dt == ge::DT_UINT16) {
	153	+ dilation = twoByte;
	154	+ } else if (dt == ge::DT_INT32) {
	155	+ dilation = fourByte;
	156	+ } else if (dt == ge::DT_UINT32) {
	157	+ dilation = fourByte;
	158	+ } else if (dt == ge::DT_INT64) {
	159	+ dilation = eightByte;
	160	+ } else if (dt == ge::DT_UINT64) {
	161	+ dilation = eightByte;
	162	+ } else if (dt == ge::DT_INT8) {
	163	+ dilation = oneByte;
	164	+ }
	165	+ return dilation;
	166	+}
	167	+
	168	+int32_t GenOnesData(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type,
	169	+ int value)
	170	+{
	171	+ input_tensor_desc.SetRealDimCnt(shapes.size());
	172	+ size_t size = 1;
	173	+ for (uint32_t i = 0; i < shapes.size(); i++) {
	174	+ size *= shapes[i];
	175	+ }
	176	+ uint32_t data_len = size * GetDataTypeSize(data_type);
	177	+ // data_len 是字节数，不能直接用作 int32_t 元素个数（会多分配 3x 内存）
	178	+ // 元素个数用 size，int32_t(4字节) 可容纳 FLOAT/FLOAT16/INT32 的存储需求
	179	+ int32_t* pData = new (std::nothrow) int32_t[size];
	180	+ if (pData == nullptr) {
	181	+ printf("%s - ERROR - [XIR]: allocate int32 data failed\n", GetTime().c_str());
	182	+ return FAILED;
	183	+ }
	184	+ for (size_t i = 0; i < size; ++i) {
	185	+ *(pData + i) = value;
	186	+ }
	187	+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len);
	188	+ delete[] pData;
	189	+ return SUCCESS;
	190	+}
	191	+
	192	+int CreateOppInGraph(RunMode mode, DataType inDtype, const std::vector<int64_t>& xShape,
	193	+ const std::vector<int64_t>& roisShape, std::vector<ge::Tensor>& input,
	194	+ std::vector<Operator>& inputs, std::vector<Operator>& outputs, Graph& graph)
	195	+{
	196	+ Status ret = SUCCESS;
	197	+ // 自定义代码：添加单算子定义到图中
	198	+ auto add1 = op::ROIPooling("roi_pooling1");
	199	+ // x: 4D 特征图 [N, C, H, W]
	200	+ ADD_INPUT_MODE(1, x, inDtype, xShape, mode);
	201	+ // rois: 2D [K, 5]，每行 [batch_idx, x1, y1, x2, y2]
	202	+ ADD_INPUT_MODE(2, rois, inDtype, roisShape, mode);
	203	+ // roi_actual_num: 可选输入，冒烟测试不连接
	204	+
	205	+ // 输出 y: [K, C, pooled_h, pooled_w]
	206	+ // K = roisShape[0], C = xShape[1], pooled_h/pooled_w 取属性值
	207	+ int64_t pooledH = 2;
	208	+ int64_t pooledW = 2;
	209	+ std::vector<int64_t> yShape = {roisShape[0], xShape[1], pooledH, pooledW};
	210	+ ADD_OUTPUT_MODE(1, y, inDtype, yShape, mode);
	211	+
	212	+ // ====== 属性设置（REQUIRED_ATTR 必须设置）======
	213	+ // 强类型接口 set_attr_{name}(value)
	214	+ add1.set_attr_pooled_h(pooledH); // Int -> int64_t
	215	+ add1.set_attr_pooled_w(pooledW); // Int -> int64_t
	216	+ add1.set_attr_spatial_scale_h(1.0f); // Float -> float
	217	+ add1.set_attr_spatial_scale_w(1.0f); // Float -> float
	218	+
	219	+ outputs.push_back(add1);
	220	+ // 添加完毕
	221	+ return SUCCESS;
	222	+}
	223	+
	224	+CaseResult RunOneCase(ge::Session* session, uint32_t graph_id, RunMode mode, DataType dtype, const ShapeCombo& combo,
	225	+ const std::string& case_name)
	226	+{
	227	+ CaseResult r;
	228	+ r.case_name = case_name;
	229	+ r.build_ok = false;
	230	+ r.run_ok = false;
	231	+ r.output_exists = false;
	232	+ r.output_count = 0;
	233	+ r.err_msg = "";
	234	+
	235	+ std::string graph_name = "tc_ge_irrun_test_" + std::to_string(graph_id);
	236	+ Graph graph(graph_name.c_str());
	237	+ std::vector<ge::Tensor> input;
	238	+ std::vector<Operator> inputs{};
	239	+ std::vector<Operator> outputs{};
	240	+
	241	+ Status ret = CreateOppInGraph(mode, dtype, combo.xShape, combo.roisShape, input, inputs, outputs, graph);
	242	+ if (ret != SUCCESS) {
	243	+ r.err_msg = "CreateOppInGraph failed";
	244	+ return r;
	245	+ }
	246	+ if (!inputs.empty() && !outputs.empty()) {
	247	+ graph.SetInputs(inputs).SetOutputs(outputs);
	248	+ }
	249	+
	250	+ std::map<AscendString, AscendString> graph_options = {};
	251	+ ret = session->AddGraph(graph_id, graph, graph_options);
	252	+ if (ret != SUCCESS) {
	253	+ r.err_msg = "AddGraph failed, ret=" + std::to_string(ret);
	254	+ return r;
	255	+ }
	256	+ r.build_ok = true;
	257	+
	258	+ std::vector<ge::Tensor> output;
	259	+ ret = session->RunGraph(graph_id, input, output);
	260	+ session->RemoveGraph(graph_id);
	261	+ if (ret != SUCCESS) {
	262	+ r.err_msg = "RunGraph failed, ret=" + std::to_string(ret);
	263	+ return r;
	264	+ }
	265	+ r.run_ok = true;
	266	+ r.output_count = output.size();
	267	+ r.output_exists = (output.size() > 0);
	268	+
	269	+ for (size_t i = 0; i < output.size(); i++) {
	270	+ int64_t shape_size = output[i].GetTensorDesc().GetShape().GetShapeSize();
	271	+ printf(" [%s] output[%zu] dtype=%d shape_size=%lld\n", case_name.c_str(), i,
	272	+ output[i].GetTensorDesc().GetDataType(), (long long)shape_size);
	273	+ }
	274	+
	275	+ return r;
	276	+}
	277	+
	278	+void PrintReport(const std::vector<CaseResult>& results)
	279	+{
	280	+ printf("\n");
	281	+ printf("====================================================================================================\n");
	282	+ printf("\| %-22s \| %-8s \| %-9s \| %-12s \| %-7s \| %-20s\n", "Case", "Build", "RunGraph", "OutputExists", "OutCnt",
	283	+ "ErrMsg");
	284	+ printf("----------------------------------------------------------------------------------------------------\n");
	285	+ int pass_cnt = 0;
	286	+ int total = results.size();
	287	+ for (const auto& r : results) {
	288	+ bool pass = r.build_ok && r.run_ok && r.output_exists;
	289	+ if (pass)
	290	+ pass_cnt++;
	291	+ printf("\| %-22s \| %-8s \| %-9s \| %-12s \| %-7d \| %-20s\n", r.case_name.c_str(), r.build_ok ? "OK" : "FAIL",
	292	+ r.run_ok ? "OK" : "FAIL", r.output_exists ? "OK" : "FAIL", r.output_count,
	293	+ r.err_msg.empty() ? "-" : r.err_msg.c_str());
	294	+ }
	295	+ printf("====================================================================================================\n");
	296	+ printf("Summary: %d/%d passed\n", pass_cnt, total);
	297	+}
	298	+
	299	+int main(int argc, char* argv[])
	300	+{
	301	+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
	302	+ // 设置全局选项
	303	+ std::map<AscendString, AscendString> global_options = {
	304	+ {"ge.exec.deviceId", "0"}, // device_id=0，npu-smi info 查询 Health=OK
	305	+ {"ge.graphRunMode", "0"},
	306	+ {"ge.exec.precision_mode", "must_keep_origin_dtype"}};
	307	+ // 初始化图引擎
	308	+ Status ret = ge::GEInitialize(global_options);
	309	+ if (ret != SUCCESS) {
	310	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
	311	+ return FAILED;
	312	+ }
	313	+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
	314	+
	315	+ // dtype 矩阵（取自 reg_op dtype_set：DT_FLOAT, DT_FLOAT16）
	316	+ struct DtypeEntry {
	317	+ DataType dt;
	318	+ std::string name;
	319	+ };
	320	+ std::vector<DtypeEntry> dtype_list = {
	321	+ {DT_FLOAT, "FP32"},
	322	+ {DT_FLOAT16, "FP16"},
	323	+ };
	324	+
	325	+ // shape 场景矩阵
	326	+ // ROI Pooling 要求 x 为 4D [N,C,H,W]，rois 为 2D [K,5]。
	327	+ // scalar/1d/8d 场景对本算子无意义（维度不匹配），已裁剪。
	328	+ // 保留 regular/minimal/empty_rois/large 四类场景覆盖常规、最小、空、大尺寸边界。
	329	+ std::vector<ShapeCombo> shape_list = {
	330	+ {{2, 4, 8, 8}, {3, 5}, "regular"}, // 常规 4D 特征图 + 3 个 ROI
	331	+ {{1, 1, 2, 2}, {1, 5}, "minimal"}, // 最小非空场景：N=1,C=1,H=2,W=2，1 个 ROI
	332	+ {{1, 4, 8, 8}, {0, 5}, "empty_rois"}, // K=0 空输出边界
	333	+ {{4, 16, 16, 16}, {10, 5}, "large"}, // 较大尺寸场景
	334	+ };
	335	+
	336	+ // 单 session 复用
	337	+ std::map<AscendString, AscendString> build_options = {};
	338	+ ge::Session* session = new Session(build_options);
	339	+ if (session == nullptr) {
	340	+ printf("%s - ERROR - [XIR]: create session failed\n", GetTime().c_str());
	341	+ ge::GEFinalize();
	342	+ return FAILED;
	343	+ }
	344	+
	345	+ std::vector<CaseResult> results;
	346	+ uint32_t graph_id = 0;
	347	+
	348	+ // N_dtype × N_shape × 2 mode 全矩阵
	349	+ for (const auto& d : dtype_list) {
	350	+ for (const auto& s : shape_list) {
	351	+ for (auto mode : {RUN_MODE_S, RUN_MODE_D}) {
	352	+ std::string mode_name = (mode == RUN_MODE_S) ? "S" : "D";
	353	+ std::string case_name = d.name + "_" + s.name + "_" + mode_name;
	354	+ printf("\n%s - INFO - [XIR]: ===== %s =====\n", GetTime().c_str(), case_name.c_str());
	355	+ CaseResult r = RunOneCase(session, graph_id, mode, d.dt, s, case_name);
	356	+ results.push_back(r);
	357	+ graph_id++;
	358	+ }
	359	+ }
	360	+ }
	361	+
	362	+ PrintReport(results);
	363	+
	364	+ bool all_pass = true;
	365	+ for (const auto& r : results) {
	366	+ if (!r.build_ok \|\| !r.run_ok \|\| !r.output_exists) {
	367	+ all_pass = false;
	368	+ }
	369	+ }
	370	+ if (all_pass) {
	371	+ printf("\n%s - INFO - [XIR]: ALL CASES PASSED\n", GetTime().c_str());
	372	+ } else {
	373	+ printf("\n%s - ERROR - [XIR]: SOME CASES FAILED, see report above\n", GetTime().c_str());
	374	+ }
	375	+
	376	+ delete session;
	377	+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
	378	+ ret = ge::GEFinalize();
	379	+ if (ret != SUCCESS) {
	380	+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
	381	+ return FAILED;
	382	+ }
	383	+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
	384	+ return all_pass ? SUCCESS : FAILED;
	385	+}

Apooling/roi_pooling/op_graph/CMakeLists.txt+13-0

@@ -0,0 +1,13 @@
	1	+# ---------------------------------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the License).
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ---------------------------------------------------------------------------------------------------------
	10	+
	11	+# Generated By CANNBot
	12	+
	13	+add_graph_plugin_sources()

Apooling/roi_pooling/op_graph/roi_pooling_graph_infer.cpp+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	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file roi_pooling_graph_infer.cpp
	15	+ * \brief roi_pooling operater graph infer resource
	16	+ */
	17	+
	18	+#include "register/op_impl_registry.h"
	19	+#include "log/log.h"
	20	+
	21	+namespace ops {
	22	+using namespace ge;
	23	+
	24	+static constexpr int64_t IDX_0 = 0;
	25	+static constexpr int64_t IDX_1 = 1;
	26	+
	27	+static ge::graphStatus InferDataTypeRoiPooling(gert::InferDataTypeContext* context)
	28	+{
	29	+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeRoiPooling");
	30	+
	31	+ // 设置输出的dtype，与输入x的dtype一致
	32	+ ge::DataType xDtype = context->GetInputDataType(IDX_0);
	33	+ ge::DataType roisDtype = context->GetInputDataType(IDX_1);
	34	+ if (roisDtype != xDtype) {
	35	+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(context->GetNodeName(), "rois",
	36	+ std::to_string(static_cast<int32_t>(roisDtype)).c_str(),
	37	+ "rois dtype must match x dtype");
	38	+ return GRAPH_FAILED;
	39	+ }
	40	+ context->SetOutputDataType(IDX_0, xDtype);
	41	+
	42	+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeRoiPooling");
	43	+ return GRAPH_SUCCESS;
	44	+}
	45	+
	46	+IMPL_OP(ROIPooling).InferDataType(InferDataTypeRoiPooling);
	47	+
	48	+}; // namespace ops

Apooling/roi_pooling/op_graph/roi_pooling_proto.h+57-0

@@ -0,0 +1,57 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/* Generated By CANNBot */
	12	+
	13	+/*!
	14	+ * \file roi_pooling_proto.h
	15	+ * \brief
	16	+ */
	17	+#ifndef OPS_OP_PROTO_INC_ROI_POOLING_H_
	18	+#define OPS_OP_PROTO_INC_ROI_POOLING_H_
	19	+
	20	+#include "graph/operator_reg.h"
	21	+#include "graph/types.h"
	22	+
	23	+namespace ge {
	24	+
	25	+/**
	26	+ *@brief Performs ROI Pooling on the input feature map.
	27	+ *@par Inputs:
	28	+ *Three inputs, including:
	29	+ * @li x: A 4D Tensor (NCHW). Must be one of the following types: float32, float16.
	30	+ * @li rois: A 2D Tensor [num_rois, 5] describing ROI boxes. Must be one of the following types: float32, float16.
	31	+ * @li roi_actual_num: An optional 1D Tensor of type int32, specifying the actual number of valid ROIs per batch.
	32	+
	33	+ *@par Outputs:
	34	+ *y: A 4D Tensor [num_rois, channels, pooled_h, pooled_w]. Must be one of the following types: float32, float16.
	35	+
	36	+ *@par Attributes:
	37	+ * @li pooled_h: An required Int, the height of the pooled output.
	38	+ * @li pooled_w: An required Int, the width of the pooled output.
	39	+ * @li spatial_scale_h: An required Float, the vertical spatial scale factor mapping ROI coords to input feature map.
	40	+ * @li spatial_scale_w: An required Float, the horizontal spatial scale factor mapping ROI coords to input feature map.
	41	+ */
	42	+#ifndef OPS_PROTO_DEF_ROIPOOLING
	43	+#define OPS_PROTO_DEF_ROIPOOLING
	44	+REG_OP(ROIPooling)
	45	+ .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16}))
	46	+ .INPUT(rois, TensorType({DT_FLOAT, DT_FLOAT16}))
	47	+ .OPTIONAL_INPUT(roi_actual_num, TensorType({DT_INT32}))
	48	+ .REQUIRED_ATTR(pooled_h, Int)
	49	+ .REQUIRED_ATTR(pooled_w, Int)
	50	+ .REQUIRED_ATTR(spatial_scale_h, Float)
	51	+ .REQUIRED_ATTR(spatial_scale_w, Float)
	52	+ .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16}))
	53	+ .OP_END_FACTORY_REG(ROIPooling)
	54	+#endif
	55	+} // namespace ge
	56	+
	57	+#endif // OPS_OP_PROTO_INC_ROI_POOLING_H_

Apooling/roi_pooling/op_host/arch35/roi_pooling_tiling_arch35.cpp+294-0

@@ -0,0 +1,294 @@
	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 roi_pooling_tiling_arch35.cpp
	13	+ * \brief tiling: validate + grid-stride core split + workspace
	14	+ */
	15	+#include "log/log.h"
	16	+#include "platform/platform_ascendc.h"
	17	+#include "securec.h"
	18	+#include "util/math_util.h"
	19	+#include "util/platform_util.h"
	20	+#include "op_host/tiling_util.h"
	21	+#include "op_host/tiling_templates_registry.h"
	22	+#include "pooling/roi_pooling/op_kernel/arch35/roi_pooling_tiling_data.h"
	23	+#include "pooling/roi_pooling/op_kernel/arch35/roi_pooling_tiling_key.h"
	24	+
	25	+namespace optiling {
	26	+
	27	+using namespace Ops::NN::OpTiling;
	28	+
	29	+constexpr int64_t PER_CORE_MIN = 1024;
	30	+constexpr uint32_t DCACHE_SIZE = 32 * 1024;
	31	+constexpr uint32_t STATIC_UB_ESTIMATE = 0;
	32	+
	33	+static constexpr int32_t kXIdx = 0;
	34	+static constexpr int32_t kRoisIdx = 1;
	35	+static constexpr int32_t kRoiActualNumIdx = 2;
	36	+
	37	+constexpr int32_t ROI_COLS = 5; // rois 每行列数 [batch_idx, x1, y1, x2, y2]
	38	+constexpr int32_t X_DIM_NUM = 4; // x 维度数 [N, C, H, W]
	39	+constexpr int32_t ROIS_DIM_NUM = 2; // rois 维度数 [K, 5]
	40	+constexpr int32_t ROI_ACTUAL_NUM_DIM_NUM = 1; // roi_actual_num 维度数 [N]
	41	+
	42	+struct RoiPoolingCompileInfo {};
	43	+
	44	+// ══ GetPlatformInfo ══
	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 must be positive, got %ld", coreNum), 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	+static ge::graphStatus ValidateDtype(gert::TilingContext* context, ge::DataType& dataType)
	58	+{
	59	+ auto xDesc = context->GetInputDesc(kXIdx);
	60	+ OP_CHECK_NULL_WITH_CONTEXT(context, xDesc);
	61	+ dataType = xDesc->GetDataType();
	62	+ if (dataType != ge::DT_FLOAT && dataType != ge::DT_FLOAT16) {
	63	+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(context->GetNodeName(), "x",
	64	+ std::to_string(static_cast<int32_t>(dataType)).c_str(),
	65	+ "x dtype must be float16/float32");
	66	+ return ge::GRAPH_FAILED;
	67	+ }
	68	+
	69	+ auto roisDesc = context->GetInputDesc(kRoisIdx);
	70	+ OP_CHECK_NULL_WITH_CONTEXT(context, roisDesc);
	71	+ if (roisDesc->GetDataType() != dataType) {
	72	+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(context->GetNodeName(), "rois",
	73	+ std::to_string(static_cast<int32_t>(roisDesc->GetDataType())).c_str(),
	74	+ "rois dtype must match x dtype");
	75	+ return ge::GRAPH_FAILED;
	76	+ }
	77	+
	78	+ // roi_actual_num 为可选输入，传入时校验 dtype 为 INT32
	79	+ auto roiActualNumDesc = context->GetInputDesc(kRoiActualNumIdx);
	80	+ if (roiActualNumDesc != nullptr) {
	81	+ if (roiActualNumDesc->GetDataType() != ge::DT_INT32) {
	82	+ OP_LOGE_FOR_INVALID_DTYPE(context->GetNodeName(), "roi_actual_num",
	83	+ std::to_string(static_cast<int32_t>(roiActualNumDesc->GetDataType())).c_str(),
	84	+ "int32");
	85	+ return ge::GRAPH_FAILED;
	86	+ }
	87	+ }
	88	+ return ge::GRAPH_SUCCESS;
	89	+}
	90	+
	91	+static ge::graphStatus ValidateShape(gert::TilingContext* context, int64_t& N, int64_t& K, int64_t& C, int64_t& H,
	92	+ int64_t& W)
	93	+{
	94	+ // x: 4D [N, C, H, W]
	95	+ auto xInput = context->GetInputShape(kXIdx);
	96	+ OP_CHECK_NULL_WITH_CONTEXT(context, xInput);
	97	+ auto xShape = xInput->GetStorageShape();
	98	+ if (xShape.GetDimNum() != X_DIM_NUM) {
	99	+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "x", (std::to_string(xShape.GetDimNum()) + "D").c_str(),
	100	+ "4D");
	101	+ return ge::GRAPH_FAILED;
	102	+ }
	103	+ N = xShape.GetDim(0);
	104	+ C = xShape.GetDim(1);
	105	+ H = xShape.GetDim(2);
	106	+ W = xShape.GetDim(3);
	107	+ if (N <= 0 \|\| C <= 0 \|\| H <= 0 \|\| W <= 0) {
	108	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "x",
	109	+ ("N=" + std::to_string(N) + " C=" + std::to_string(C) +
	110	+ " H=" + std::to_string(H) + " W=" + std::to_string(W))
	111	+ .c_str(),
	112	+ "x dims must be positive");
	113	+ return ge::GRAPH_FAILED;
	114	+ }
	115	+
	116	+ // roiInput: 2D [K, 5]
	117	+ auto roisInput = context->GetInputShape(kRoisIdx);
	118	+ OP_CHECK_NULL_WITH_CONTEXT(context, roisInput);
	119	+ auto roisShape = roisInput->GetStorageShape();
	120	+ if (roisShape.GetDimNum() != ROIS_DIM_NUM) {
	121	+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "rois",
	122	+ (std::to_string(roisShape.GetDimNum()) + "D").c_str(), "2D");
	123	+ return ge::GRAPH_FAILED;
	124	+ }
	125	+ K = roisShape.GetDim(0);
	126	+ if (roisShape.GetDim(1) != ROI_COLS) {
	127	+ OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(context->GetNodeName(), "rois.shape[1]",
	128	+ std::to_string(roisShape.GetDim(1)).c_str(), "rois dim[1] must be 5");
	129	+ return ge::GRAPH_FAILED;
	130	+ }
	131	+
	132	+ // roi_actual_num 为可选输入，传入时校验为 1D
	133	+ auto roiActualNumInput = context->GetInputShape(kRoiActualNumIdx);
	134	+ if (roiActualNumInput != nullptr) {
	135	+ auto& roiActualNumShape = roiActualNumInput->GetStorageShape();
	136	+ if (roiActualNumShape.GetDimNum() != ROI_ACTUAL_NUM_DIM_NUM) {
	137	+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "roi_actual_num",
	138	+ (std::to_string(roiActualNumShape.GetDimNum()) + "D").c_str(), "1D");
	139	+ return ge::GRAPH_FAILED;
	140	+ }
	141	+ }
	142	+ return ge::GRAPH_SUCCESS;
	143	+}
	144	+
	145	+static ge::graphStatus ValidateAttr(gert::TilingContext* context, int64_t& pooledH, int64_t& pooledW,
	146	+ float& spatialScaleH, float& spatialScaleW)
	147	+{
	148	+ auto attrs = context->GetAttrs();
	149	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
	150	+ const auto pooledHPtr = attrs->GetAttrPointer<int64_t>(0);
	151	+ OP_CHECK_NULL_WITH_CONTEXT(context, pooledHPtr);
	152	+ pooledH = *pooledHPtr;
	153	+ const auto pooledWPtr = attrs->GetAttrPointer<int64_t>(1);
	154	+ OP_CHECK_NULL_WITH_CONTEXT(context, pooledWPtr);
	155	+ pooledW = *pooledWPtr;
	156	+ const auto spatialScaleHPtr = attrs->GetAttrPointer<float>(2);
	157	+ OP_CHECK_NULL_WITH_CONTEXT(context, spatialScaleHPtr);
	158	+ spatialScaleH = *spatialScaleHPtr;
	159	+ const auto spatialScaleWPtr = attrs->GetAttrPointer<float>(3);
	160	+ OP_CHECK_NULL_WITH_CONTEXT(context, spatialScaleWPtr);
	161	+ spatialScaleW = *spatialScaleWPtr;
	162	+ if (pooledH <= 0 \|\| pooledW <= 0) {
	163	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
	164	+ context->GetNodeName(), "pooled_h/pooled_w",
	165	+ ("pooledH=" + std::to_string(pooledH) + " pooledW=" + std::to_string(pooledW)).c_str(),
	166	+ "pooled_h/w must > 0");
	167	+ return ge::GRAPH_FAILED;
	168	+ }
	169	+ if (spatialScaleH <= 0.0f \|\| spatialScaleW <= 0.0f) {
	170	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
	171	+ context->GetNodeName(), "spatial_scale_h/spatial_scale_w",
	172	+ ("spatialScaleH=" + std::to_string(spatialScaleH) + " spatialScaleW=" + std::to_string(spatialScaleW))
	173	+ .c_str(),
	174	+ "spatial_scale_h/w must > 0");
	175	+ return ge::GRAPH_FAILED;
	176	+ }
	177	+ return ge::GRAPH_SUCCESS;
	178	+}
	179	+
	180	+// ══ ValidateInputs 调度器 ══
	181	+static ge::graphStatus ValidateInputs(gert::TilingContext* context, ge::DataType& dataType, int64_t& N, int64_t& K,
	182	+ int64_t& C, int64_t& H, int64_t& W, int64_t& pooledH, int64_t& pooledW,
	183	+ float& spatialScaleH, float& spatialScaleW)
	184	+{
	185	+ OP_CHECK_IF(ValidateDtype(context, dataType) != ge::GRAPH_SUCCESS, OP_LOGE(context, "ValidateDtype failed"),
	186	+ return ge::GRAPH_FAILED);
	187	+ OP_CHECK_IF(ValidateShape(context, N, K, C, H, W) != ge::GRAPH_SUCCESS, OP_LOGE(context, "ValidateShape failed"),
	188	+ return ge::GRAPH_FAILED);
	189	+ OP_CHECK_IF(ValidateAttr(context, pooledH, pooledW, spatialScaleH, spatialScaleW) != ge::GRAPH_SUCCESS,
	190	+ OP_LOGE(context, "ValidateAttr failed"), return ge::GRAPH_FAILED);
	191	+ return ge::GRAPH_SUCCESS;
	192	+}
	193	+
	194	+static ge::graphStatus ComputeTiling(RoiPoolingTilingData* tiling, int64_t totalElements, int64_t N, int64_t K,
	195	+ int64_t C, int64_t H, int64_t W, int64_t pooledH, int64_t pooledW,
	196	+ float spatialScaleH, float spatialScaleW, int64_t coreNum)
	197	+{
	198	+ tiling->totalElements = totalElements;
	199	+ tiling->N = N;
	200	+ tiling->K = K;
	201	+ tiling->C = C;
	202	+ tiling->H = H;
	203	+ tiling->W = W;
	204	+ tiling->pooledH = pooledH;
	205	+ tiling->pooledW = pooledW;
	206	+ tiling->spatialScaleH = spatialScaleH;
	207	+ tiling->spatialScaleW = spatialScaleW;
	208	+
	209	+ int64_t blockFactor = (totalElements + coreNum - 1) / coreNum;
	210	+ if (blockFactor < PER_CORE_MIN)
	211	+ blockFactor = PER_CORE_MIN;
	212	+ tiling->needCoreNum = (totalElements + blockFactor - 1) / blockFactor;
	213	+ if (tiling->needCoreNum > coreNum)
	214	+ tiling->needCoreNum = coreNum;
	215	+ if (tiling->needCoreNum <= 0)
	216	+ tiling->needCoreNum = 1;
	217	+ return ge::GRAPH_SUCCESS;
	218	+}
	219	+
	220	+static void DumpTilingData(gert::TilingContext* context, const RoiPoolingTilingData* tiling)
	221	+{
	222	+ OP_LOGD(context,
	223	+ "RoiPoolingTilingData: totalElements=%ld, needCoreNum=%ld, N=%ld, K=%ld, C=%ld, H=%ld, W=%ld, "
	224	+ "pooledH=%ld, pooledW=%ld, spatialScaleH=%.6f, spatialScaleW=%.6f",
	225	+ tiling->totalElements, tiling->needCoreNum, tiling->N, tiling->K, tiling->C, tiling->H, tiling->W,
	226	+ tiling->pooledH, tiling->pooledW, tiling->spatialScaleH, tiling->spatialScaleW);
	227	+}
	228	+
	229	+// ══ GetWorkspaceSize（系统 workspace，无用户 workspace）══
	230	+static ge::graphStatus SetupWorkspace(gert::TilingContext* context)
	231	+{
	232	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
	233	+ uint64_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
	234	+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
	235	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	236	+ currentWorkspace[0] = static_cast<size_t>(sysWorkspaceSize);
	237	+ return ge::GRAPH_SUCCESS;
	238	+}
	239	+
	240	+// ══ TilingFunc 主流程 ══
	241	+static ge::graphStatus RoiPoolingTilingFunc(gert::TilingContext* context)
	242	+{
	243	+ OP_LOGD(context, "RoiPoolingTilingFunc enter.");
	244	+ // 1. validate
	245	+ ge::DataType dataType;
	246	+ int64_t N = 0, K = 0, C = 0, H = 0, W = 0, pooledH = 0, pooledW = 0;
	247	+ float spatialScaleH = 0.0f, spatialScaleW = 0.0f;
	248	+ OP_CHECK_IF(ValidateInputs(context, dataType, N, K, C, H, W, pooledH, pooledW, spatialScaleH, spatialScaleW) !=
	249	+ ge::GRAPH_SUCCESS,
	250	+ OP_LOGE(context, "ValidateInputs failed"), return ge::GRAPH_FAILED);
	251	+
	252	+ // 2. platform
	253	+ uint64_t ubSize = 0;
	254	+ int64_t coreNum = 0;
	255	+ OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
	256	+ OP_LOGE(context, "GetPlatformInfo failed"), return ge::GRAPH_FAILED);
	257	+
	258	+ // 3. compute tiling
	259	+ int64_t totalElements = K * C * pooledH * pooledW;
	260	+ RoiPoolingTilingData* tiling = context->GetTilingData<RoiPoolingTilingData>();
	261	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	262	+ OP_CHECK_IF(memset_s(tiling, sizeof(RoiPoolingTilingData), 0, sizeof(RoiPoolingTilingData)) != EOK,
	263	+ OP_LOGE(context, "memset_s tiling failed"), return ge::GRAPH_FAILED);
	264	+ OP_CHECK_IF(ComputeTiling(tiling, totalElements, N, K, C, H, W, pooledH, pooledW, spatialScaleH, spatialScaleW,
	265	+ coreNum) != ge::GRAPH_SUCCESS,
	266	+ OP_LOGE(context, "ComputeTiling failed"), return ge::GRAPH_FAILED);
	267	+
	268	+ // 4. DFX log
	269	+ DumpTilingData(context, tiling);
	270	+
	271	+ // 5. workspace（仅系统 workspace）
	272	+ OP_CHECK_IF(SetupWorkspace(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "SetupWorkspace failed"),
	273	+ return ge::GRAPH_FAILED);
	274	+
	275	+ // 6. set block dim + local memory
	276	+ context->SetBlockDim(static_cast<uint32_t>(tiling->needCoreNum));
	277	+ OP_CHECK_IF(ubSize <= DCACHE_SIZE + STATIC_UB_ESTIMATE, OP_LOGE(context, "ubSize %lu <= DCache+Static", ubSize),
	278	+ return ge::GRAPH_FAILED);
	279	+ context->SetLocalMemorySize(static_cast<uint32_t>(ubSize - DCACHE_SIZE - STATIC_UB_ESTIMATE));
	280	+
	281	+ // 7. tiling key（单一场景模式，dtype 由 DTYPE_ 宏实例化）
	282	+ context->SetTilingKey(GET_TPL_TILING_KEY(ROI_POOLING_SCH_MODE_DEFAULT));
	283	+ return ge::GRAPH_SUCCESS;
	284	+}
	285	+
	286	+// ══ TilingParse ══
	287	+static ge::graphStatus TilingParseForRoiPooling([[maybe_unused]] gert::TilingParseContext* context)
	288	+{
	289	+ return ge::GRAPH_SUCCESS;
	290	+}
	291	+
	292	+IMPL_OP_OPTILING(ROIPooling).Tiling(RoiPoolingTilingFunc).TilingParse<RoiPoolingCompileInfo>(TilingParseForRoiPooling);
	293	+
	294	+} // namespace optiling

Apooling/roi_pooling/op_host/config/ascend950/roi_pooling_binary.json+151-0

@@ -0,0 +1,151 @@
	1	+{
	2	+ "op_type": "ROIPooling",
	3	+ "op_list": [
	4	+ {
	5	+ "bin_filename": "ROIPooling_26b375e9e74290c680050922753a7c15",
	6	+ "inputs": [
	7	+ {
	8	+ "name": "x",
	9	+ "index": 0,
	10	+ "dtype": "float32",
	11	+ "format": "ND",
	12	+ "paramType": "required",
	13	+ "shape": [
	14	+ -2
	15	+ ],
	16	+ "format_match_mode": "FormatAgnostic"
	17	+ },
	18	+ {
	19	+ "name": "rois",
	20	+ "index": 1,
	21	+ "dtype": "float32",
	22	+ "format": "ND",
	23	+ "paramType": "required",
	24	+ "shape": [
	25	+ -2
	26	+ ],
	27	+ "format_match_mode": "FormatAgnostic"
	28	+ },
	29	+ {
	30	+ "name": "roi_actual_num",
	31	+ "index": 2,
	32	+ "dtype": "int32",
	33	+ "format": "ND",
	34	+ "paramType": "optional",
	35	+ "shape": [
	36	+ -2
	37	+ ],
	38	+ "format_match_mode": "FormatAgnostic"
	39	+ }
	40	+ ],
	41	+ "outputs": [
	42	+ {
	43	+ "name": "y",
	44	+ "index": 0,
	45	+ "dtype": "float32",
	46	+ "format": "ND",
	47	+ "paramType": "required",
	48	+ "shape": [
	49	+ -2
	50	+ ],
	51	+ "format_match_mode": "FormatAgnostic"
	52	+ }
	53	+ ],
	54	+ "attrs": [
	55	+ {
	56	+ "name": "pooled_h",
	57	+ "dtype": "int",
	58	+ "value": 0
	59	+ },
	60	+ {
	61	+ "name": "pooled_w",
	62	+ "dtype": "int",
	63	+ "value": 0
	64	+ },
	65	+ {
	66	+ "name": "spatial_scale_h",
	67	+ "dtype": "float",
	68	+ "value": 0.0
	69	+ },
	70	+ {
	71	+ "name": "spatial_scale_w",
	72	+ "dtype": "float",
	73	+ "value": 0.0
	74	+ }
	75	+ ]
	76	+ },
	77	+ {
	78	+ "bin_filename": "ROIPooling_8c6aacae521127f659a1fcf5920e9de4",
	79	+ "inputs": [
	80	+ {
	81	+ "name": "x",
	82	+ "index": 0,
	83	+ "dtype": "float16",
	84	+ "format": "ND",
	85	+ "paramType": "required",
	86	+ "shape": [
	87	+ -2
	88	+ ],
	89	+ "format_match_mode": "FormatAgnostic"
	90	+ },
	91	+ {
	92	+ "name": "rois",
	93	+ "index": 1,
	94	+ "dtype": "float16",
	95	+ "format": "ND",
	96	+ "paramType": "required",
	97	+ "shape": [
	98	+ -2
	99	+ ],
	100	+ "format_match_mode": "FormatAgnostic"
	101	+ },
	102	+ {
	103	+ "name": "roi_actual_num",
	104	+ "index": 2,
	105	+ "dtype": "int32",
	106	+ "format": "ND",
	107	+ "paramType": "optional",
	108	+ "shape": [
	109	+ -2
	110	+ ],
	111	+ "format_match_mode": "FormatAgnostic"
	112	+ }
	113	+ ],
	114	+ "outputs": [
	115	+ {
	116	+ "name": "y",
	117	+ "index": 0,
	118	+ "dtype": "float16",
	119	+ "format": "ND",
	120	+ "paramType": "required",
	121	+ "shape": [
	122	+ -2
	123	+ ],
	124	+ "format_match_mode": "FormatAgnostic"
	125	+ }
	126	+ ],
	127	+ "attrs": [
	128	+ {
	129	+ "name": "pooled_h",
	130	+ "dtype": "int",
	131	+ "value": 0
	132	+ },
	133	+ {
	134	+ "name": "pooled_w",
	135	+ "dtype": "int",
	136	+ "value": 0
	137	+ },
	138	+ {
	139	+ "name": "spatial_scale_h",
	140	+ "dtype": "float",
	141	+ "value": 0.0
	142	+ },
	143	+ {
	144	+ "name": "spatial_scale_w",
	145	+ "dtype": "float",
	146	+ "value": 0.0
	147	+ }
	148	+ ]
	149	+ }
	150	+ ]
	151	+}

Apooling/roi_pooling/op_host/roi_pooling_def.cpp+64-0

@@ -0,0 +1,64 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file roi_pooling_def.cpp
	13	+ * \brief op definition for roi_pooling
	14	+ */
	15	+#include "register/op_def_registry.h"
	16	+
	17	+namespace ops {
	18	+class ROIPooling : public OpDef {
	19	+public:
	20	+ explicit ROIPooling(const char* name) : OpDef(name)
	21	+ {
	22	+ this->Input("x")
	23	+ .ParamType(REQUIRED)
	24	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
	25	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	26	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	27	+ .AutoContiguous();
	28	+ this->Input("rois")
	29	+ .ParamType(REQUIRED)
	30	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
	31	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	32	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	33	+ .AutoContiguous();
	34	+ this->Input("roi_actual_num")
	35	+ .ParamType(OPTIONAL)
	36	+ .DataType({ge::DT_INT32, ge::DT_INT32})
	37	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	38	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	39	+ .AutoContiguous();
	40	+ this->Output("y")
	41	+ .ParamType(REQUIRED)
	42	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16})
	43	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	44	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND})
	45	+ .AutoContiguous();
	46	+
	47	+ this->Attr("pooled_h").AttrType(REQUIRED).Int();
	48	+ this->Attr("pooled_w").AttrType(REQUIRED).Int();
	49	+ this->Attr("spatial_scale_h").AttrType(REQUIRED).Float();
	50	+ this->Attr("spatial_scale_w").AttrType(REQUIRED).Float();
	51	+
	52	+ OpAICoreConfig aicoreConfig;
	53	+ aicoreConfig.DynamicCompileStaticFlag(true)
	54	+ .DynamicFormatFlag(false)
	55	+ .DynamicRankSupportFlag(true)
	56	+ .DynamicShapeSupportFlag(true)
	57	+ .NeedCheckSupportFlag(false)
	58	+ .PrecisionReduceFlag(true)
	59	+ .ExtendCfgInfo("opFile.value", "roi_pooling_apt");
	60	+ this->AICore().AddConfig("ascend950", aicoreConfig);
	61	+ }
	62	+};
	63	+OP_ADD(ROIPooling);
	64	+} // namespace ops

Apooling/roi_pooling/op_host/roi_pooling_infershape.cpp+88-0

@@ -0,0 +1,88 @@
	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 roi_pooling_infershape.cpp
	13	+ * \brief infer shape for roi_pooling
	14	+ */
	15	+#include "register/op_impl_registry.h"
	16	+#include "log/log.h"
	17	+#include "util/shape_util.h"
	18	+
	19	+using namespace ge;
	20	+
	21	+namespace ops {
	22	+
	23	+constexpr int32_t ROI_COLS = 5; // rois 每行列数 [batch_idx, x1, y1, x2, y2]
	24	+constexpr int32_t X_DIM_NUM = 4; // x 维度数 [N, C, H, W]
	25	+constexpr int32_t ROIS_DIM_NUM = 2; // rois 维度数 [K, 5]
	26	+
	27	+// K = rois.shape[0], C = x.shape[1], pooled_h/w from attrs
	28	+// 动态 rank 保护：图模式下 shape 为 unknown rank {-2} 时 GetDimNum() 返回 1，
	29	+// 会被误判为"非 2D/4D"报错；unknown rank 时跳过维度校验，输出对应 dim 设为 UNKNOWN_DIM。
	30	+static ge::graphStatus InferShapeRoiPooling(gert::InferShapeContext* context)
	31	+{
	32	+ const gert::Shape* xShape = context->GetInputShape(0);
	33	+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
	34	+ const gert::Shape* roisShape = context->GetInputShape(1);
	35	+ OP_CHECK_NULL_WITH_CONTEXT(context, roisShape);
	36	+
	37	+ if (!Ops::Base::IsUnknownRank(*roisShape)) {
	38	+ OP_CHECK_IF(roisShape->GetDimNum() != ROIS_DIM_NUM,
	39	+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "rois",
	40	+ (std::to_string(roisShape->GetDimNum()) + "D").c_str(), "2D"),
	41	+ return ge::GRAPH_FAILED);
	42	+ if (roisShape->GetDim(1) != ROI_COLS && roisShape->GetDim(1) != ge::UNKNOWN_DIM) {
	43	+ OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(context->GetNodeName(), "rois.shape[1]",
	44	+ std::to_string(roisShape->GetDim(1)).c_str(),
	45	+ "rois dim[1] must be 5");
	46	+ return ge::GRAPH_FAILED;
	47	+ }
	48	+ }
	49	+ if (!Ops::Base::IsUnknownRank(*xShape)) {
	50	+ OP_CHECK_IF(xShape->GetDimNum() != X_DIM_NUM,
	51	+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "x",
	52	+ (std::to_string(xShape->GetDimNum()) + "D").c_str(), "4D"),
	53	+ return ge::GRAPH_FAILED);
	54	+ }
	55	+
	56	+ // unknown rank 时输出 dim 设为 UNKNOWN_DIM(-1)，正常时取实际值（含 -1 透传）
	57	+ int64_t K = Ops::Base::IsUnknownRank(*roisShape) ? ge::UNKNOWN_DIM : roisShape->GetDim(0);
	58	+ int64_t C = Ops::Base::IsUnknownRank(*xShape) ? ge::UNKNOWN_DIM : xShape->GetDim(1);
	59	+
	60	+ auto attrs = context->GetAttrs();
	61	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
	62	+ const auto pooledHPtr = attrs->GetAttrPointer<int64_t>(0);
	63	+ OP_CHECK_NULL_WITH_CONTEXT(context, pooledHPtr);
	64	+ int64_t pooledH = *pooledHPtr;
	65	+ const auto pooledWPtr = attrs->GetAttrPointer<int64_t>(1);
	66	+ OP_CHECK_NULL_WITH_CONTEXT(context, pooledWPtr);
	67	+ int64_t pooledW = *pooledWPtr;
	68	+ if (pooledH <= 0 \|\| pooledW <= 0) {
	69	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
	70	+ context->GetNodeName(), "pooled_h/pooled_w",
	71	+ ("pooledH=" + std::to_string(pooledH) + " pooledW=" + std::to_string(pooledW)).c_str(),
	72	+ "pooled_h/w must > 0");
	73	+ return ge::GRAPH_FAILED;
	74	+ }
	75	+ // 输出 [K, C, pooledH, pooledW]
	76	+ gert::Shape* yShape = context->GetOutputShape(0);
	77	+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
	78	+ yShape->SetDimNum(4);
	79	+ yShape->SetDim(0, K);
	80	+ yShape->SetDim(1, C);
	81	+ yShape->SetDim(2, pooledH);
	82	+ yShape->SetDim(3, pooledW);
	83	+ return ge::GRAPH_SUCCESS;
	84	+}
	85	+
	86	+IMPL_OP_INFERSHAPE(ROIPooling).InferShape(InferShapeRoiPooling);
	87	+
	88	+} // namespace ops

Apooling/roi_pooling/op_kernel/CMakeLists.txt+17-0

@@ -0,0 +1,17 @@
	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	+# arch35 (ascend950)
	13	+add_kernel_sources(
	14	+ KERNEL_SRC roi_pooling_apt.cpp
	15	+ COMPUTE_UNITS ascend950
	16	+ AUTO_SYNC false
	17	+)

Apooling/roi_pooling/op_kernel/arch35/roi_pooling_simt.h+198-0

@@ -0,0 +1,198 @@
	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 roi_pooling_simt.h
	13	+ * \brief SIMT kernel implementation for roi_pooling
	14	+ */
	15	+#ifndef ROI_POOLING_SIMT_H
	16	+#define ROI_POOLING_SIMT_H
	17	+
	18	+#include "kernel_operator.h"
	19	+#include "kernel_tiling/kernel_tiling.h"
	20	+#include "simt_api/common_functions.h"
	21	+#include "simt_api/asc_simt.h"
	22	+#include "simt_api/math_functions.h"
	23	+#include "simt_api/math_constants.h"
	24	+#include "simt_api/asc_fp16.h"
	25	+#include "roi_pooling_tiling_data.h"
	26	+#include "roi_pooling_tiling_key.h"
	27	+
	28	+namespace NsRoiPooling {
	29	+
	30	+using namespace AscendC;
	31	+
	32	+constexpr uint32_t THREAD_NUM = 512;
	33	+constexpr int32_t ROI_COLS = 5; // [batch_idx, x1, y1, x2, y2]
	34	+
	35	+template <typename T>
	36	+__simt_callee__ __aicore__ inline float ToFloat(T val)
	37	+{
	38	+ if constexpr (std::is_same_v<T, float>) {
	39	+ return val;
	40	+ } else {
	41	+ return __half2float(val);
	42	+ }
	43	+}
	44	+
	45	+template <typename T>
	46	+__simt_callee__ __aicore__ inline T FromFloat(float val)
	47	+{
	48	+ if constexpr (std::is_same_v<T, float>) {
	49	+ return val;
	50	+ } else {
	51	+ return __float2half(val);
	52	+ }
	53	+}
	54	+
	55	+// ========== 辅助函数：clip 到 [0, upper] ==========
	56	+// val 和 upper 均为 int64_t：与 int64_t 的 H/W 及 float→int64 转换结果对齐。
	57	+__simt_callee__ __aicore__ inline int64_t ClipInt(int64_t val, int64_t upper)
	58	+{
	59	+ if (val < 0)
	60	+ return 0;
	61	+ if (val > upper)
	62	+ return upper;
	63	+ return val;
	64	+}
	65	+
	66	+// ========== Bin 区域 max 扫描 ==========
	67	+// 在 float 域做 max 比较，避免 half -INF 和 half > 比较的潜在问题
	68	+// bin 坐标为 int64_t：与 ClipInt 返回类型对齐
	69	+template <typename T>
	70	+__simt_callee__ __aicore__ inline float ScanBinForMax(__gm__ T* x_gm, int64_t xBase, int64_t W, int64_t binY1,
	71	+ int64_t binY2, int64_t binX1, int64_t binX2)
	72	+{
	73	+ float maxVal = -ASCRT_INF_F; // float -INF
	74	+ for (int64_t h = binY1; h < binY2; h++) {
	75	+ int64_t rowBase = xBase + h * W;
	76	+ for (int64_t w = binX1; w < binX2; w++) {
	77	+ float val = ToFloat(x_gm[rowBase + w]);
	78	+ if (val > maxVal) {
	79	+ maxVal = val;
	80	+ }
	81	+ }
	82	+ }
	83	+ return maxVal;
	84	+}
	85	+
	86	+// ========== 处理单个输出元素 ==========
	87	+template <typename T>
	88	+__simt_callee__ __aicore__ inline void ProcessOneOutputElement(int64_t idx, int64_t N, int64_t C, int64_t H, int64_t W,
	89	+ int64_t pooledH, int64_t pooledW, float spatialScaleH,
	90	+ float spatialScaleW, int64_t strideC, int64_t strideN,
	91	+ int64_t xStrideC, int64_t xStrideN, __gm__ T* x_gm,
	92	+ __gm__ T* rois_gm, __gm__ T* y_gm)
	93	+{
	94	+ // ============ 1. 索引分解：idx → (n, c, ph, pw) ============
	95	+ // N/C/pooledH/pooledW 均为 int64_t，全程 int64 运算，无截断
	96	+ int64_t n = idx / strideN;
	97	+ int64_t rem1 = idx - n * strideN;
	98	+ int64_t c = rem1 / strideC;
	99	+ int64_t rem2 = rem1 - c * strideC;
	100	+ int64_t ph = rem2 / pooledW;
	101	+ int64_t pw = rem2 - ph * pooledW;
	102	+
	103	+ // ============ 2. 读取 rois 行，提升到 float ============
	104	+ float roiBatchF = ToFloat(rois_gm[n * ROI_COLS + 0]);
	105	+ float roiX1 = ToFloat(rois_gm[n * ROI_COLS + 1]);
	106	+ float roiY1 = ToFloat(rois_gm[n * ROI_COLS + 2]);
	107	+ float roiX2 = ToFloat(rois_gm[n * ROI_COLS + 3]);
	108	+ float roiY2 = ToFloat(rois_gm[n * ROI_COLS + 4]);
	109	+
	110	+ // ============ 3. batchIdx 越界双侧保护 ============
	111	+ int64_t batchIdx = static_cast<int64_t>(roiBatchF);
	112	+ if (batchIdx < 0 \|\| batchIdx >= N) {
	113	+ y_gm[idx] = FromFloat<T>(0.0f);
	114	+ return;
	115	+ }
	116	+
	117	+ // ============ 4. ROI 坐标映射（roundf 取整为 int，无 +1 偏移）============
	118	+ // roi_start = round(coord * spatial_scale) ← int64 类型
	119	+ // roi_end = round(coord * spatial_scale) ← int64 类型，无 +1 偏移
	120	+ // +1 偏移在 roi_width 上（见步骤5），不在坐标上
	121	+ // y 方向用 spatialScaleH，x 方向用 spatialScaleW
	122	+ int64_t roiStartW = static_cast<int64_t>(roundf(roiX1 * spatialScaleW));
	123	+ int64_t roiStartH = static_cast<int64_t>(roundf(roiY1 * spatialScaleH));
	124	+ int64_t roiEndW = static_cast<int64_t>(roundf(roiX2 * spatialScaleW)); // ← 无 +1 偏移
	125	+ int64_t roiEndH = static_cast<int64_t>(roundf(roiY2 * spatialScaleH)); // ← 无 +1 偏移
	126	+
	127	+ // ============ 5. ROI 尺寸（int64 运算，malformed 强制非空）============
	128	+ // +1 偏移在这里（Fast R-CNN 标准），不在坐标上
	129	+ // roiEndW/roiStartW 均为 int64_t，减法+1 不会溢出。
	130	+ int64_t roiWidth = roiEndW - roiStartW + 1;
	131	+ int64_t roiHeight = roiEndH - roiStartH + 1;
	132	+ if (roiWidth < 1)
	133	+ roiWidth = 1;
	134	+ if (roiHeight < 1)
	135	+ roiHeight = 1;
	136	+
	137	+ // ============ 6. Bin 大小（float，基于 int64 roi_width）============
	138	+ float binSizeW = static_cast<float>(roiWidth) / static_cast<float>(pooledW);
	139	+ float binSizeH = static_cast<float>(roiHeight) / static_cast<float>(pooledH);
	140	+
	141	+ // ============ 7. Bin 边界 floor/ceil → int64 → + roiStart → clip ============
	142	+ int64_t binX1 = ClipInt(static_cast<int64_t>(floorf(static_cast<float>(pw) * binSizeW)) + roiStartW, W);
	143	+ int64_t binY1 = ClipInt(static_cast<int64_t>(floorf(static_cast<float>(ph) * binSizeH)) + roiStartH, H);
	144	+ int64_t binX2 = ClipInt(static_cast<int64_t>(ceilf(static_cast<float>(pw + 1) * binSizeW)) + roiStartW, W);
	145	+ int64_t binY2 = ClipInt(static_cast<int64_t>(ceilf(static_cast<float>(ph + 1) * binSizeH)) + roiStartH, H);
	146	+
	147	+ // ============ 8. Max pooling ============
	148	+ if (binY2 <= binY1 \|\| binX2 <= binX1) {
	149	+ // 空 bin 输出 0
	150	+ y_gm[idx] = FromFloat<T>(0.0f);
	151	+ } else {
	152	+ int64_t xBase = batchIdx * xStrideN + c * xStrideC;
	153	+ float maxVal = ScanBinForMax<T>(x_gm, xBase, W, binY1, binY2, binX1, binX2);
	154	+ y_gm[idx] = FromFloat<T>(maxVal);
	155	+ }
	156	+}
	157	+
	158	+// ========== 主计算 VF ==========
	159	+template <typename T>
	160	+__simt_vf__ __aicore__ __launch_bounds__(THREAD_NUM) inline void OpRoiPoolingSimtKernel(
	161	+ int64_t totalElements, int64_t N, int64_t K, int64_t C, int64_t H, int64_t W, int64_t pooledH, int64_t pooledW,
	162	+ float spatialScaleH, float spatialScaleW, __gm__ T* x_gm, __gm__ T* rois_gm, __gm__ T* y_gm)
	163	+{
	164	+ // 预计算 stride（固定除数，VF 内计算，避免重复乘法）
	165	+ // N/C/H/W/pooledH/pooledW 均为 int64_t，乘法天然 int64，无需 static_cast
	166	+ const int64_t strideC = pooledH * pooledW;
	167	+ const int64_t strideN = C * strideC;
	168	+ const int64_t xStrideC = H * W;
	169	+ const int64_t xStrideN = C * xStrideC;
	170	+
	171	+ // Grid-Stride 循环：每个线程独占处理一个输出元素
	172	+ for (int64_t idx = static_cast<int64_t>(blockIdx.x * blockDim.x + threadIdx.x); idx < totalElements;
	173	+ idx += static_cast<int64_t>(blockDim.x * gridDim.x)) {
	174	+ ProcessOneOutputElement<T>(idx, N, C, H, W, pooledH, pooledW, spatialScaleH, spatialScaleW, strideC, strideN,
	175	+ xStrideC, xStrideN, x_gm, rois_gm, y_gm);
	176	+ }
	177	+}
	178	+
	179	+template <typename T>
	180	+__aicore__ inline void Process(GM_ADDR x, GM_ADDR rois, GM_ADDR roi_actual_num, GM_ADDR y, GM_ADDR workspace,
	181	+ GM_ADDR tiling, const RoiPoolingTilingData* tilingData)
	182	+{
	183	+ __gm__ T* x_gm = (__gm__ T*)x;
	184	+ __gm__ T* rois_gm = (__gm__ T*)rois;
	185	+ __gm__ T* y_gm = (__gm__ T*)y;
	186	+ // roi_actual_num 本算子未使用（保留接口对齐），不取地址
	187	+ (void)roi_actual_num;
	188	+ (void)workspace;
	189	+ (void)tiling;
	190	+ asc_vf_call<OpRoiPoolingSimtKernel<T>>(dim3(THREAD_NUM), tilingData->totalElements, tilingData->N, tilingData->K,
	191	+ tilingData->C, tilingData->H, tilingData->W, tilingData->pooledH,
	192	+ tilingData->pooledW, tilingData->spatialScaleH, tilingData->spatialScaleW,
	193	+ x_gm, rois_gm, y_gm);
	194	+}
	195	+
	196	+} // namespace NsRoiPooling
	197	+
	198	+#endif

Apooling/roi_pooling/op_kernel/arch35/roi_pooling_tiling_data.h+32-0

@@ -0,0 +1,32 @@
	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 roi_pooling_tiling_data.h
	13	+ * \brief tiling data struct for roi_pooling
	14	+ */
	15	+#ifndef ROI_POOLING_TILING_DATA_H
	16	+#define ROI_POOLING_TILING_DATA_H
	17	+
	18	+struct RoiPoolingTilingData {
	19	+ int64_t totalElements; // 输出总元素数 = K * C * pooled_h * pooled_w（Grid-Stride 遍历上界）
	20	+ int64_t needCoreNum; // 实际启动核数
	21	+ int64_t N; // 特征图 batch 数 = x.shape[0]（用于 kernel 内 batchIdx 越界检查）
	22	+ int64_t K; // ROI 数量 = rois.shape[0]
	23	+ int64_t C; // 通道数 = x.shape[1]
	24	+ int64_t H; // 特征图高 = x.shape[2]
	25	+ int64_t W; // 特征图宽 = x.shape[3]
	26	+ int64_t pooledH; // 池化输出高（来自属性 pooled_h）
	27	+ int64_t pooledW; // 池化输出宽（来自属性 pooled_w）
	28	+ float spatialScaleH; // y 方向缩放因子（来自属性 spatial_scale_h）
	29	+ float spatialScaleW; // x 方向缩放因子（来自属性 spatial_scale_w）
	30	+};
	31	+
	32	+#endif

Apooling/roi_pooling/op_kernel/arch35/roi_pooling_tiling_key.h+26-0

@@ -0,0 +1,26 @@
	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 roi_pooling_tiling_key.h
	13	+ * \brief tiling key declare for roi_pooling (scene-encoded, dtype via DTYPE_ macro)
	14	+ */
	15	+#ifndef ROI_POOLING_TILING_KEY_H
	16	+#define ROI_POOLING_TILING_KEY_H
	17	+
	18	+#include "ascendc/host_api/tiling/template_argument.h"
	19	+
	20	+#define ROI_POOLING_SCH_MODE_DEFAULT 0
	21	+
	22	+ASCENDC_TPL_ARGS_DECL(ROIPooling, ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ROI_POOLING_SCH_MODE_DEFAULT));
	23	+
	24	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ROI_POOLING_SCH_MODE_DEFAULT)));
	25	+
	26	+#endif

Apooling/roi_pooling/op_kernel/roi_pooling_apt.cpp+28-0

@@ -0,0 +1,28 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file roi_pooling_apt.cpp
	13	+ * \brief kernel entry for roi_pooling
	14	+ */
	15	+#include "arch35/roi_pooling_simt.h"
	16	+
	17	+template <uint32_t schMode>
	18	+__global__ __aicore__ void roi_pooling(GM_ADDR x, GM_ADDR rois, GM_ADDR roi_actual_num, GM_ADDR y, GM_ADDR workspace,
	19	+ GM_ADDR tiling)
	20	+{
	21	+ REGISTER_TILING_DEFAULT(RoiPoolingTilingData);
	22	+ GET_TILING_DATA_WITH_STRUCT(RoiPoolingTilingData, tilingData, tiling);
	23	+
	24	+ // 单一场景模式，dtype 由 DTYPE_ 宏自动实例化（DTYPE_X 对应 def.cpp Input("x")）
	25	+ if constexpr (schMode == static_cast<uint32_t>(ROI_POOLING_SCH_MODE_DEFAULT)) {
	26	+ NsRoiPooling::Process<DTYPE_X>(x, rois, roi_actual_num, y, workspace, tiling, &tilingData);
	27	+ }
	28	+}

Apooling/roi_pooling/tests/CMakeLists.txt+18-0

@@ -0,0 +1,18 @@
	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	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	14	+foreach(SUB_DIR ${CURRENT_DIRS})
	15	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	16	+ add_subdirectory(${SUB_DIR})
	17	+ endif()
	18	+endforeach()

Apooling/roi_pooling/tests/assets/golden.py+232-0

@@ -0,0 +1,232 @@
	1	+#!/usr/bin/env python3
	2	+# -- coding: utf-8 --
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	5	+# CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	8	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+# See LICENSE in the root of the software repository for the full text of the License.
	10	+# ----------------------------------------------------------------------------
	11	+# Generated By CANNBot
	12	+
	13	+"""roi_pooling golden (kernel/geir).
	14	+
	15	+Kernel/GEIR golden receives numpy.ndarray, converts to torch, returns numpy.
	16	+Equal scale (scale_h == scale_w) uses torchvision.ops.roi_pool;
	17	+Unequal scale uses torch implementation matching TBE (round + int +1).
	18	+Only supports 2D rois [K, 5].
	19	+"""
	20	+
	21	+__spec__ = {
	22	+ "roi_pooling": "RoiPoolingKernelSpec",
	23	+}
	24	+
	25	+import math
	26	+
	27	+import numpy as np
	28	+import torch
	29	+from torchvision.ops import roi_pool
	30	+
	31	+
	32	+def _fix_batch_idx(x, rois):
	33	+ """Fix rois[:, 0] (batch_idx) to [0, N-1], in-place, only changes batch_idx column."""
	34	+ N = x.shape[0]
	35	+ if isinstance(rois, np.ndarray):
	36	+ rois[:, 0] = (np.floor(rois[:, 0]).astype(np.int32) % N).astype(rois.dtype)
	37	+ else:
	38	+ rois[:, 0] = (torch.floor(rois[:, 0]).long() % N).to(rois.dtype)
	39	+ return rois
	40	+
	41	+
	42	+def _round_away_from_zero(t):
	43	+ """C++ round semantics: round half away from zero (vectorized, supports tensor and scalar)."""
	44	+ if isinstance(t, torch.Tensor):
	45	+ return torch.where(t >= 0, torch.floor(t + 0.5), -torch.floor(-t + 0.5)).long()
	46	+ return int(math.floor(t + 0.5)) if t >= 0 else -int(math.floor(-t + 0.5))
	47	+
	48	+
	49	+def _roi_pool_torch(x, rois, pooled_h, pooled_w, ssh, ssw):
	50	+ """ROI Pooling matching TBE/torchvision algorithm (vectorized, no per-element .item())."""
	51	+ N, C, H, W = x.shape
	52	+ K = rois.shape[0]
	53	+ if K == 0:
	54	+ return torch.zeros(0, C, pooled_h, pooled_w, dtype=x.dtype, device=x.device)
	55	+
	56	+ batch_idx = rois[:, 0].long()
	57	+ batch_idx = torch.clamp(batch_idx, 0, N - 1)
	58	+ x_gathered = x[batch_idx]
	59	+
	60	+ roi_start_w = _round_away_from_zero(rois[:, 1] * ssw)
	61	+ roi_start_h = _round_away_from_zero(rois[:, 2] * ssh)
	62	+ roi_end_w = _round_away_from_zero(rois[:, 3] * ssw)
	63	+ roi_end_h = _round_away_from_zero(rois[:, 4] * ssh)
	64	+
	65	+ roi_w = torch.clamp(roi_end_w - roi_start_w + 1, min=1)
	66	+ roi_h = torch.clamp(roi_end_h - roi_start_h + 1, min=1)
	67	+
	68	+ bin_size_w = roi_w.float() / pooled_w
	69	+ bin_size_h = roi_h.float() / pooled_h
	70	+
	71	+ ph_idx = torch.arange(pooled_h, device=x.device, dtype=torch.float32)
	72	+ pw_idx = torch.arange(pooled_w, device=x.device, dtype=torch.float32)
	73	+
	74	+ hstart = torch.clamp(
	75	+ torch.floor(ph_idx.unsqueeze(1) * bin_size_h.unsqueeze(0)).long()
	76	+ + roi_start_h.unsqueeze(0),
	77	+ 0,
	78	+ H,
	79	+ )
	80	+ hend = torch.clamp(
	81	+ torch.ceil((ph_idx.unsqueeze(1) + 1) * bin_size_h.unsqueeze(0)).long()
	82	+ + roi_start_h.unsqueeze(0),
	83	+ 0,
	84	+ H,
	85	+ )
	86	+ wstart = torch.clamp(
	87	+ torch.floor(pw_idx.unsqueeze(1) * bin_size_w.unsqueeze(0)).long()
	88	+ + roi_start_w.unsqueeze(0),
	89	+ 0,
	90	+ W,
	91	+ )
	92	+ wend = torch.clamp(
	93	+ torch.ceil((pw_idx.unsqueeze(1) + 1) * bin_size_w.unsqueeze(0)).long()
	94	+ + roi_start_w.unsqueeze(0),
	95	+ 0,
	96	+ W,
	97	+ )
	98	+
	99	+ max_bin_h = int((hend - hstart).clamp(min=1).max())
	100	+ max_bin_w = int((wend - wstart).clamp(min=1).max())
	101	+
	102	+ offsets_h = torch.arange(max_bin_h, device=x.device, dtype=torch.long)
	103	+ offsets_w = torch.arange(max_bin_w, device=x.device, dtype=torch.long)
	104	+
	105	+ result_h = torch.empty(pooled_h, K, C, W, dtype=x.dtype, device=x.device)
	106	+
	107	+ for ph in range(pooled_h):
	108	+ hs = hstart[ph]
	109	+ he = hend[ph]
	110	+ h_idx = hs.unsqueeze(1) + offsets_h
	111	+ h_valid = h_idx < he.unsqueeze(1)
	112	+ h_idx_clamped = h_idx.clamp(max=H - 1)
	113	+ h_idx_exp = h_idx_clamped.unsqueeze(1).unsqueeze(3).expand(K, C, max_bin_h, W)
	114	+ gathered = torch.gather(x_gathered, 2, h_idx_exp)
	115	+ invalid = ~h_valid.unsqueeze(1).unsqueeze(3) \| torch.isnan(gathered)
	116	+ gathered.masked_fill_(invalid, float("-inf"))
	117	+ result_h[ph] = gathered.amax(dim=2)
	118	+
	119	+ out = torch.empty(K, C, pooled_h, pooled_w, dtype=x.dtype, device=x.device)
	120	+
	121	+ for pw in range(pooled_w):
	122	+ ws = wstart[pw]
	123	+ we = wend[pw]
	124	+ w_idx = ws.unsqueeze(1) + offsets_w
	125	+ w_valid = w_idx < we.unsqueeze(1)
	126	+ w_idx_clamped = w_idx.clamp(max=W - 1)
	127	+ w_idx_exp = (
	128	+ w_idx_clamped.unsqueeze(0).unsqueeze(2).expand(pooled_h, K, C, max_bin_w)
	129	+ )
	130	+ gathered_w = torch.gather(result_h, 3, w_idx_exp)
	131	+ gathered_w.masked_fill_(~w_valid.unsqueeze(0).unsqueeze(2), float("-inf"))
	132	+ result_hw = gathered_w.amax(dim=3)
	133	+ out[:, :, :, pw] = result_hw.permute(1, 2, 0)
	134	+
	135	+ # 匹配 torchvision: 空 bin -> 0, 非空但无有效值(全 NaN) -> -FLT_MAX
	136	+ is_empty_h = hend <= hstart
	137	+ is_empty_w = wend <= wstart
	138	+ is_empty = (is_empty_h.unsqueeze(0) \| is_empty_w.unsqueeze(1)).permute(2, 1, 0)
	139	+ is_empty = is_empty.unsqueeze(1).expand_as(out)
	140	+ neg_min = torch.finfo(out.dtype).min
	141	+ out = torch.where(
	142	+ is_empty,
	143	+ torch.zeros_like(out),
	144	+ torch.where(out == float("-inf"), torch.full_like(out, neg_min), out),
	145	+ )
	146	+ return out
	147	+
	148	+
	149	+def _roi_pool_compute(x, rois, pooled_h, pooled_w, spatial_scale_h, spatial_scale_w):
	150	+ """ROI Pooling core computation (shared by golden).
	151	+
	152	+ Equal scale uses roi_pool; unequal scale uses _roi_pool_torch.
	153	+ """
	154	+ output_size = (int(pooled_h), int(pooled_w))
	155	+ ssh = float(spatial_scale_h)
	156	+ ssw = float(spatial_scale_w)
	157	+ if abs(ssh - ssw) < 1e-9:
	158	+ return roi_pool(x, rois, output_size=output_size, spatial_scale=ssh)
	159	+ return _roi_pool_torch(x, rois, int(pooled_h), int(pooled_w), ssh, ssw)
	160	+
	161	+
	162	+class RoiPoolingKernelSpec:
	163	+ """Kernel / GEIR golden"""
	164	+
	165	+ def golden(
	166	+ x,
	167	+ rois,
	168	+ roi_actual_num=None,
	169	+ *,
	170	+ pooled_h,
	171	+ pooled_w,
	172	+ spatial_scale_h,
	173	+ spatial_scale_w,
	174	+ **kwargs,
	175	+ ):
	176	+ x_t = torch.from_numpy(np.ascontiguousarray(x))
	177	+ rois_t = torch.from_numpy(np.ascontiguousarray(rois))
	178	+ orig_dtype = x_t.dtype
	179	+ if orig_dtype == torch.float16:
	180	+ x_t = x_t.float()
	181	+ rois_t = rois_t.float()
	182	+ out = _roi_pool_compute(
	183	+ x_t, rois_t, pooled_h, pooled_w, spatial_scale_h, spatial_scale_w
	184	+ )
	185	+ if orig_dtype == torch.float16:
	186	+ out = out.half()
	187	+ return [out.numpy()]
	188	+
	189	+ def customize_inputs(x, rois, roi_actual_num, **kwargs):
	190	+ rois = _fix_batch_idx(x, rois)
	191	+ return (x, rois, roi_actual_num)
	192	+
	193	+ class ThirdPartyImpl:
	194	+ def __init__(
	195	+ self,
	196	+ x,
	197	+ rois,
	198	+ *,
	199	+ pooled_h,
	200	+ pooled_w,
	201	+ spatial_scale_h,
	202	+ spatial_scale_w,
	203	+ **kwargs,
	204	+ ):
	205	+ self._output_size = (int(pooled_h), int(pooled_w))
	206	+ self._ssh = float(spatial_scale_h)
	207	+ self._ssw = float(spatial_scale_w)
	208	+ self._is_equal = abs(self._ssh - self._ssw) < 1e-9
	209	+ self._x = x
	210	+ self._rois = rois
	211	+
	212	+ def __call__(self, x, rois, **kwargs):
	213	+ if self._is_equal:
	214	+ return [
	215	+ roi_pool(
	216	+ self._x,
	217	+ self._rois,
	218	+ output_size=self._output_size,
	219	+ spatial_scale=self._ssh,
	220	+ )
	221	+ ]
	222	+ return [
	223	+ _roi_pool_torch(
	224	+ self._x, self._rois, *self._output_size, self._ssh, self._ssw
	225	+ )
	226	+ ]
	227	+
	228	+ third_party = {"torch": ThirdPartyImpl}
	229	+ tolerance = {
	230	+ "float16": {"standard": "cross_check", "level": "L1"},
	231	+ "float32": {"standard": "cross_check", "level": "L1"},
	232	+ }

Apooling/roi_pooling/tests/ut/CMakeLists.txt+18-0

@@ -0,0 +1,18 @@
	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	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	14	+foreach(SUB_DIR ${CURRENT_DIRS})
	15	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	16	+ add_subdirectory(${SUB_DIR})
	17	+ endif()
	18	+endforeach()

Apooling/roi_pooling/tests/ut/op_host/CMakeLists.txt+25-0

@@ -0,0 +1,25 @@
	1	+# -----------------------------------------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the License).
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# -----------------------------------------------------------------------------------------------------------
	10	+
	11	+# Generated By CANNBot
	12	+
	13	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	14	+if(UT_TEST_ALL OR OP_HOST_UT)
	15	+ if("${ASCEND_COMPUTE_UNIT}" STREQUAL "ascend950")
	16	+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}/arch35)
	17	+ endif()
	18	+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	19	+endif()
	20	+
	21	+foreach(SUB_DIR ${CURRENT_DIRS})
	22	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	23	+ add_subdirectory(${SUB_DIR})
	24	+ endif()
	25	+endforeach()

Apooling/roi_pooling/tests/ut/op_host/arch35/test_roi_pooling_tiling.cpp+152-0

@@ -0,0 +1,152 @@
	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	+#include <iostream>
	13	+#include <gtest/gtest.h>
	14	+#include "tiling_context_faker.h"
	15	+#include "tiling_case_executor.h"
	16	+#include "../../../../op_kernel/arch35/roi_pooling_tiling_data.h"
	17	+
	18	+using namespace std;
	19	+using namespace ge;
	20	+
	21	+class RoiPoolingTiling : public testing::Test {
	22	+protected:
	23	+ static void SetUpTestCase() { std::cout << "RoiPoolingTiling SetUp" << std::endl; }
	24	+
	25	+ static void TearDownTestCase() { std::cout << "RoiPoolingTiling TearDown" << std::endl; }
	26	+};
	27	+
	28	+TEST_F(RoiPoolingTiling, roi_pooling_float32_small_shape_success)
	29	+{
	30	+ struct RoiPoolingCompileInfo {
	31	+ } compileInfo;
	32	+ gert::TilingContextPara tilingContextPara(
	33	+ "ROIPooling",
	34	+ {
	35	+ {{{2, 4, 8, 8}, {2, 4, 8, 8}}, ge::DT_FLOAT, ge::FORMAT_ND}, // x [N, C, H, W]
	36	+ {{{10, 5}, {10, 5}}, ge::DT_FLOAT, ge::FORMAT_ND}, // rois [K, 5]
	37	+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND}, // roi_actual_num [N]
	38	+ },
	39	+ {
	40	+ {{{10, 4, 7, 7}, {10, 4, 7, 7}}, ge::DT_FLOAT, ge::FORMAT_ND}, // y [K, C, pooled_h, pooled_w]
	41	+ },
	42	+ {
	43	+ gert::TilingContextPara::OpAttr("pooled_h", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	44	+ gert::TilingContextPara::OpAttr("pooled_w", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	45	+ gert::TilingContextPara::OpAttr("spatial_scale_h", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	46	+ gert::TilingContextPara::OpAttr("spatial_scale_w", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	47	+ },
	48	+ &compileInfo, 32, 262144, 4096);
	49	+
	50	+ TilingInfo tilingInfo;
	51	+ EXPECT_TRUE(ExecuteTiling(tilingContextPara, tilingInfo));
	52	+
	53	+ EXPECT_EQ(tilingInfo.tilingKey, 0);
	54	+
	55	+ auto* tiling = reinterpret_cast<const RoiPoolingTilingData*>(tilingInfo.tilingData.get());
	56	+ int64_t expectedTotalElements = 10 * 4 * 7 * 7;
	57	+ EXPECT_EQ(tiling->totalElements, expectedTotalElements);
	58	+ EXPECT_EQ(tiling->needCoreNum, 2);
	59	+ EXPECT_EQ(tiling->N, 2);
	60	+ EXPECT_EQ(tiling->K, 10);
	61	+ EXPECT_EQ(tiling->C, 4);
	62	+ EXPECT_EQ(tiling->H, 8);
	63	+ EXPECT_EQ(tiling->W, 8);
	64	+ EXPECT_EQ(tiling->pooledH, 7);
	65	+ EXPECT_EQ(tiling->pooledW, 7);
	66	+ EXPECT_FLOAT_EQ(tiling->spatialScaleH, 1.0f);
	67	+ EXPECT_FLOAT_EQ(tiling->spatialScaleW, 1.0f);
	68	+}
	69	+
	70	+TEST_F(RoiPoolingTiling, roi_pooling_float16_medium_shape_success)
	71	+{
	72	+ struct RoiPoolingCompileInfo {
	73	+ } compileInfo;
	74	+ gert::TilingContextPara tilingContextPara(
	75	+ "ROIPooling",
	76	+ {
	77	+ {{{1, 3, 16, 16}, {1, 3, 16, 16}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	78	+ {{{100, 5}, {100, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	79	+ {{{1}, {1}}, ge::DT_INT32, ge::FORMAT_ND},
	80	+ },
	81	+ {
	82	+ {{{100, 3, 6, 6}, {100, 3, 6, 6}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	83	+ },
	84	+ {
	85	+ gert::TilingContextPara::OpAttr("pooled_h", Ops::NN::AnyValue::CreateFrom<int64_t>(6)),
	86	+ gert::TilingContextPara::OpAttr("pooled_w", Ops::NN::AnyValue::CreateFrom<int64_t>(6)),
	87	+ gert::TilingContextPara::OpAttr("spatial_scale_h", Ops::NN::AnyValue::CreateFrom<float>(0.5f)),
	88	+ gert::TilingContextPara::OpAttr("spatial_scale_w", Ops::NN::AnyValue::CreateFrom<float>(0.5f)),
	89	+ },
	90	+ &compileInfo, 32, 262144, 4096);
	91	+
	92	+ TilingInfo tilingInfo;
	93	+ EXPECT_TRUE(ExecuteTiling(tilingContextPara, tilingInfo));
	94	+
	95	+ EXPECT_EQ(tilingInfo.tilingKey, 0);
	96	+
	97	+ auto* tiling = reinterpret_cast<const RoiPoolingTilingData*>(tilingInfo.tilingData.get());
	98	+ int64_t expectedTotalElements = 100 * 3 * 6 * 6;
	99	+ EXPECT_EQ(tiling->totalElements, expectedTotalElements);
	100	+ EXPECT_EQ(tiling->needCoreNum, 11);
	101	+ EXPECT_EQ(tiling->N, 1);
	102	+ EXPECT_EQ(tiling->K, 100);
	103	+ EXPECT_EQ(tiling->C, 3);
	104	+ EXPECT_EQ(tiling->H, 16);
	105	+ EXPECT_EQ(tiling->W, 16);
	106	+ EXPECT_EQ(tiling->pooledH, 6);
	107	+ EXPECT_EQ(tiling->pooledW, 6);
	108	+ EXPECT_FLOAT_EQ(tiling->spatialScaleH, 0.5f);
	109	+ EXPECT_FLOAT_EQ(tiling->spatialScaleW, 0.5f);
	110	+}
	111	+
	112	+TEST_F(RoiPoolingTiling, roi_pooling_float32_large_shape_full_core_success)
	113	+{
	114	+ struct RoiPoolingCompileInfo {
	115	+ } compileInfo;
	116	+ gert::TilingContextPara tilingContextPara(
	117	+ "ROIPooling",
	118	+ {
	119	+ {{{4, 64, 32, 32}, {4, 64, 32, 32}}, ge::DT_FLOAT, ge::FORMAT_ND},
	120	+ {{{1000, 5}, {1000, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	121	+ {{{4}, {4}}, ge::DT_INT32, ge::FORMAT_ND},
	122	+ },
	123	+ {
	124	+ {{{1000, 64, 7, 7}, {1000, 64, 7, 7}}, ge::DT_FLOAT, ge::FORMAT_ND},
	125	+ },
	126	+ {
	127	+ gert::TilingContextPara::OpAttr("pooled_h", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	128	+ gert::TilingContextPara::OpAttr("pooled_w", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	129	+ gert::TilingContextPara::OpAttr("spatial_scale_h", Ops::NN::AnyValue::CreateFrom<float>(0.25f)),
	130	+ gert::TilingContextPara::OpAttr("spatial_scale_w", Ops::NN::AnyValue::CreateFrom<float>(0.25f)),
	131	+ },
	132	+ &compileInfo, 32, 262144, 4096);
	133	+
	134	+ TilingInfo tilingInfo;
	135	+ EXPECT_TRUE(ExecuteTiling(tilingContextPara, tilingInfo));
	136	+
	137	+ EXPECT_EQ(tilingInfo.tilingKey, 0);
	138	+
	139	+ auto* tiling = reinterpret_cast<const RoiPoolingTilingData*>(tilingInfo.tilingData.get());
	140	+ int64_t expectedTotalElements = 1000 * 64 * 7 * 7;
	141	+ EXPECT_EQ(tiling->totalElements, expectedTotalElements);
	142	+ EXPECT_EQ(tiling->needCoreNum, 32);
	143	+ EXPECT_EQ(tiling->N, 4);
	144	+ EXPECT_EQ(tiling->K, 1000);
	145	+ EXPECT_EQ(tiling->C, 64);
	146	+ EXPECT_EQ(tiling->H, 32);
	147	+ EXPECT_EQ(tiling->W, 32);
	148	+ EXPECT_EQ(tiling->pooledH, 7);
	149	+ EXPECT_EQ(tiling->pooledW, 7);
	150	+ EXPECT_FLOAT_EQ(tiling->spatialScaleH, 0.25f);
	151	+ EXPECT_FLOAT_EQ(tiling->spatialScaleW, 0.25f);
	152	+}

Apooling/roi_pooling/tests/ut/op_host/test_roi_pooling_infershape.cpp+144-0

@@ -0,0 +1,144 @@
	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	+#include <gtest/gtest.h>
	13	+#include <iostream>
	14	+#include "infer_shape_context_faker.h"
	15	+#include "infershape_case_executor.h"
	16	+
	17	+class RoiPoolingInfershape : public testing::Test {
	18	+protected:
	19	+ static void SetUpTestCase() { std::cout << "RoiPoolingInfershape SetUp" << std::endl; }
	20	+
	21	+ static void TearDownTestCase() { std::cout << "RoiPoolingInfershape TearDown" << std::endl; }
	22	+};
	23	+
	24	+TEST_F(RoiPoolingInfershape, roi_pooling_infershape_float32_success)
	25	+{
	26	+ gert::InfershapeContextPara infershapeContextPara(
	27	+ "ROIPooling",
	28	+ {
	29	+ {{{2, 4, 8, 8}, {2, 4, 8, 8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	30	+ {{{10, 5}, {10, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	31	+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND},
	32	+ },
	33	+ {
	34	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	35	+ },
	36	+ {
	37	+ gert::InfershapeContextPara::OpAttr("pooled_h", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	38	+ gert::InfershapeContextPara::OpAttr("pooled_w", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	39	+ gert::InfershapeContextPara::OpAttr("spatial_scale_h", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	40	+ gert::InfershapeContextPara::OpAttr("spatial_scale_w", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	41	+ });
	42	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	43	+ {10, 4, 7, 7},
	44	+ };
	45	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	46	+}
	47	+
	48	+TEST_F(RoiPoolingInfershape, roi_pooling_infershape_float16_success)
	49	+{
	50	+ gert::InfershapeContextPara infershapeContextPara(
	51	+ "ROIPooling",
	52	+ {
	53	+ {{{1, 3, 16, 16}, {1, 3, 16, 16}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	54	+ {{{100, 5}, {100, 5}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	55	+ {{{1}, {1}}, ge::DT_INT32, ge::FORMAT_ND},
	56	+ },
	57	+ {
	58	+ {{{}, {}}, ge::DT_FLOAT16, ge::FORMAT_ND},
	59	+ },
	60	+ {
	61	+ gert::InfershapeContextPara::OpAttr("pooled_h", Ops::NN::AnyValue::CreateFrom<int64_t>(6)),
	62	+ gert::InfershapeContextPara::OpAttr("pooled_w", Ops::NN::AnyValue::CreateFrom<int64_t>(6)),
	63	+ gert::InfershapeContextPara::OpAttr("spatial_scale_h", Ops::NN::AnyValue::CreateFrom<float>(0.5f)),
	64	+ gert::InfershapeContextPara::OpAttr("spatial_scale_w", Ops::NN::AnyValue::CreateFrom<float>(0.5f)),
	65	+ });
	66	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	67	+ {100, 3, 6, 6},
	68	+ };
	69	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	70	+}
	71	+
	72	+TEST_F(RoiPoolingInfershape, roi_pooling_infershape_3d_rois_failed)
	73	+{
	74	+ gert::InfershapeContextPara infershapeContextPara(
	75	+ "ROIPooling",
	76	+ {
	77	+ {{{2, 4, 8, 8}, {2, 4, 8, 8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	78	+ {{{1, 10, 5}, {1, 10, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	79	+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND},
	80	+ },
	81	+ {
	82	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	83	+ },
	84	+ {
	85	+ gert::InfershapeContextPara::OpAttr("pooled_h", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	86	+ gert::InfershapeContextPara::OpAttr("pooled_w", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	87	+ gert::InfershapeContextPara::OpAttr("spatial_scale_h", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	88	+ gert::InfershapeContextPara::OpAttr("spatial_scale_w", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	89	+ });
	90	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	91	+ {1, 10, 4, 7, 7},
	92	+ };
	93	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_FAILED, expectOutputShape);
	94	+}
	95	+
	96	+// Unknown shape：x 的 H 维 unknown dim(-1)，rois 已知，output 维取 rois[0]/x[1]/attr
	97	+TEST_F(RoiPoolingInfershape, roi_pooling_infershape_unknown_dim_x)
	98	+{
	99	+ gert::InfershapeContextPara infershapeContextPara(
	100	+ "ROIPooling",
	101	+ {
	102	+ {{{2, 4, -1, 8}, {2, 4, -1, 8}}, ge::DT_FLOAT, ge::FORMAT_ND},
	103	+ {{{10, 5}, {10, 5}}, ge::DT_FLOAT, ge::FORMAT_ND},
	104	+ {{{0}, {0}}, ge::DT_INT32, ge::FORMAT_ND},
	105	+ },
	106	+ {
	107	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	108	+ },
	109	+ {
	110	+ gert::InfershapeContextPara::OpAttr("pooled_h", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	111	+ gert::InfershapeContextPara::OpAttr("pooled_w", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	112	+ gert::InfershapeContextPara::OpAttr("spatial_scale_h", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	113	+ gert::InfershapeContextPara::OpAttr("spatial_scale_w", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	114	+ });
	115	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	116	+ {10, 4, 7, 7},
	117	+ };
	118	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	119	+}
	120	+
	121	+// Unknown rank：x 和 rois 均为 {-2}，输出 dim0/dim1 置 UNKNOWN_DIM(-1)，pooled_h/w 取 attr
	122	+TEST_F(RoiPoolingInfershape, roi_pooling_infershape_unknown_rank)
	123	+{
	124	+ gert::InfershapeContextPara infershapeContextPara(
	125	+ "ROIPooling",
	126	+ {
	127	+ {{{-2}, {-2}}, ge::DT_FLOAT, ge::FORMAT_ND},
	128	+ {{{-2}, {-2}}, ge::DT_FLOAT, ge::FORMAT_ND},
	129	+ {{{-2}, {-2}}, ge::DT_INT32, ge::FORMAT_ND},
	130	+ },
	131	+ {
	132	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
	133	+ },
	134	+ {
	135	+ gert::InfershapeContextPara::OpAttr("pooled_h", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	136	+ gert::InfershapeContextPara::OpAttr("pooled_w", Ops::NN::AnyValue::CreateFrom<int64_t>(7)),
	137	+ gert::InfershapeContextPara::OpAttr("spatial_scale_h", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	138	+ gert::InfershapeContextPara::OpAttr("spatial_scale_w", Ops::NN::AnyValue::CreateFrom<float>(1.0f)),
	139	+ });
	140	+ std::vector<std::vector<int64_t>> expectOutputShape = {
	141	+ {-1, -1, 7, 7},
	142	+ };
	143	+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
	144	+}