已合并
[CANNBot] 增加 ExtractGlimpseV2,UpsampleBicubic2d 支持 Ascend950 实现 SIMT #1039
[CANNBot] 增加 ExtractGlimpseV2,UpsampleBicubic2d 支持 Ascend950 实现 SIMT #1039
已合并
ideal_创建于 6月25日
30 个文件变更+1994-4
@@ -66,6 +66,16 @@
66 <td>AI CPU</td>66 <td>AI CPU</td>
67 <td>从输入图像中提取多个裁剪区域,并将它们统一调整为指定大小,支持双线性插值和最近邻插值。</td>67 <td>从输入图像中提取多个裁剪区域,并将它们统一调整为指定大小,支持双线性插值和最近邻插值。</td>
68 </tr>68 </tr>
69+ <tr>
70+ <td>image</td>
71+ <td><a href="../../image/extract_glimpse_v2/README.md">extract_glimpse_v2</a></td>
72+ <td>&check;</td>
73+ <td>&check;</td>
74+ <td>&cross;</td>
75+ <td>&check;</td>
76+ <td>AI Core</td>
77+ <td>从批量输入图像中提取指定位置和大小的子图像(glimpse)。</td>
78+ </tr>
69 <tr>79 <tr>
70 <td>image</td>80 <td>image</td>
71 <td><a href="../../image/grid_sample/README.md">grid_sample</a></td>81 <td><a href="../../image/grid_sample/README.md">grid_sample</a></td>
@@ -0,0 +1,27 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+ 
11+# Generated By CANNBot
12+ 
13+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
14+ 
15+if(NOT ENABLE_TEST AND NOT BENCHMARK)
16+ list(REMOVE_ITEM CURRENT_DIRS tests)
17+endif()
18+ 
19+foreach(SUB_DIR ${CURRENT_DIRS})
20+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
21+ add_subdirectory(${SUB_DIR})
22+ endif()
23+endforeach()
24+ 
25+set(SUPPORT_COMPUTE_UNIT "ascend950")
26+set(SUPPORT_TILING_DIR "arch35")
27+add_all_modules_sources(OPTYPE extract_glimpse_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
@@ -0,0 +1,125 @@
1+# ExtractGlimpseV2
C
Cchenjiao6月25日

新增算子需同步刷新op_list.md

likedislike
ideal_
ideal_
6月25日 评论:
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+- 算子功能:从批量输入图像中提取指定位置和大小的子图像(glimpse)。根据给定的偏移坐标和裁剪尺寸,从每张输入图像中裁剪出一个子区域,输出为一批裁剪后的图像。
17+ 
18+- 计算公式:
19+ 
20+$$
21+\text{对于每个 batch item } i:
22+$$
23+$$
24+\text{1. 坐标变换: } offset\_y, offset\_x \text{ (根据 normalized/centered 属性)}
25+$$
26+$$
27+\text{2. 裁剪区域: } [start\_y, end\_y) \times [start\_x, end\_x) \text{ (clamp 到图像边界)}
28+$$
29+$$
30+\text{3. 数据拷贝: } glimpse[i][base\_y:base\_y+copy\_h][base\_x:base\_x+copy\_w][:] = input[i][start\_y:end\_y][start\_x:end\_x][:]
31+$$
32+$$
33+\text{4. 越界填充: } glimpse \text{ 中未拷贝区域填零 (noise="zero")}
34+$$
35+ 
36+## 参数说明
37+ 
38+<table style="undefined;table-layout: fixed; width: 980px"><colgroup>
39+ <col style="width: 100px">
40+ <col style="width: 150px">
41+ <col style="width: 280px">
42+ <col style="width: 330px">
43+ <col style="width: 120px">
44+ </colgroup>
45+ <thead>
46+ <tr>
47+ <th>参数名</th>
48+ <th>输入/输出/属性</th>
49+ <th>描述</th>
50+ <th>数据类型</th>
51+ <th>数据格式</th>
52+ </tr></thead>
53+ <tbody>
54+ <tr>
55+ <td>input</td>
56+ <td>输入</td>
57+ <td>输入图像批次,4D tensor,shape为[batch, height, width, channels]。</td>
58+ <td>FLOAT</td>
59+ <td>ND</td>
60+ </tr>
61+ <tr>
62+ <td>size</td>
63+ <td>输入</td>
64+ <td>裁剪尺寸[glimpse_h, glimpse_w],1D const tensor,shape为[2]。</td>
65+ <td>INT32</td>
66+ <td>ND</td>
67+ </tr>
68+ <tr>
69+ <td>offsets</td>
70+ <td>输入</td>
71+ <td>每个batch item的偏移坐标(y, x),2D tensor,shape为[batch, 2]。</td>
72+ <td>FLOAT</td>
73+ <td>ND</td>
74+ </tr>
75+ <tr>
76+ <td>glimpse</td>
77+ <td>输出</td>
78+ <td>裁剪后的图像批次,4D tensor,shape为[batch, size_h, size_w, channels]。</td>
79+ <td>FLOAT</td>
80+ <td>ND</td>
81+ </tr>
82+ <tr>
83+ <td>centered</td>
84+ <td>属性</td>
85+ <td>偏移坐标是否相对于图像居中。默认值:true。</td>
86+ <td>Bool</td>
87+ <td>-</td>
88+ </tr>
89+ <tr>
90+ <td>normalized</td>
91+ <td>属性</td>
92+ <td>偏移坐标是否归一化到[-1, 1]。默认值:true。</td>
93+ <td>Bool</td>
94+ <td>-</td>
95+ </tr>
96+ <tr>
97+ <td>uniform_noise</td>
98+ <td>属性</td>
99+ <td>是否使用均匀分布噪声(当前实现仅支持false)。默认值:true。</td>
100+ <td>Bool</td>
101+ <td>-</td>
102+ </tr>
103+ <tr>
104+ <td>noise</td>
105+ <td>属性</td>
106+ <td>噪声类型(当前实现仅支持"zero")。默认值:"uniform"。</td>
107+ <td>String</td>
108+ <td>-</td>
109+ </tr>
110+ </tbody></table>
111+ 
112+## 约束说明
113+ 
114+- input 必须为 4D tensor。
115+- offsets 必须为 2D tensor,且第二维必须为 2。
116+- size 必须为 1D const tensor,且长度为 2。
117+- input 和 offsets 的 batch 维度必须一致。
118+- 当前仅支持 float32 数据类型。
119+- 当前仅支持 noise="zero" 的越界填充模式。
C
Cchenjiao6月25日

缺少调用说明

likedislike
ideal_
ideal_
6月25日 评论:
120+ 
121+## 调用说明
122+ 
123+| 调用方式 | 样例代码 | 说明 |
124+| ---------------- | --------------------------- | --------------------------------------------------- |
125+| 图模式 | - | 通过[算子IR](op_graph/extract_glimpse_v2_proto.h)构图方式调用ExtractGlimpseV2算子。 |
@@ -0,0 +1,40 @@
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 extract_glimpse_v2_graph_infer.cpp
15+ * \brief extract_glimpse_v2 operator 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+ 
26+static ge::graphStatus InferDataTypeExtractGlimpseV2(gert::InferDataTypeContext* context)
27+{
28+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeExtractGlimpseV2");
29+ 
30+ // Output dtype is same as input dtype
31+ ge::DataType inputDtype = context->GetInputDataType(IDX_0);
32+ context->SetOutputDataType(IDX_0, inputDtype);
33+ 
34+ OP_LOGD(context->GetNodeName(), "End to do InferDataTypeExtractGlimpseV2");
35+ return GRAPH_SUCCESS;
36+}
37+ 
38+IMPL_OP(ExtractGlimpseV2).InferDataType(InferDataTypeExtractGlimpseV2);
39+ 
40+}; // namespace ops
@@ -0,0 +1,56 @@
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 extract_glimpse_v2_proto.h
15+ * \brief Operator proto declaration for extract_glimpse_v2
16+ */
17+#ifndef OPS_OP_PROTO_INC_EXTRACT_GLIMPSE_V2_H_
18+#define OPS_OP_PROTO_INC_EXTRACT_GLIMPSE_V2_H_
19+ 
20+#include "graph/operator_reg.h"
21+#include "graph/types.h"
22+ 
23+namespace ge {
24+ 
25+/**
26+ *@brief Extracts a glimpse of specified size from input images at given offsets.
27+ *@par Inputs:
28+ *Three inputs, including:
29+ * @li input: A 4D NHWC Tensor. Must be one of the following types: float32.
30+ * @li size: A 1D Tensor of shape [2]. Must be one of the following types: int32.
31+ * @li offsets: A 2D Tensor of shape [batch, 2]. Must be one of the following types: float32. \n
32+ 
33+ *@par Outputs:
34+ *glimpse: A 4D NHWC Tensor of shape [batch, size_h, size_w, channels]. Must be one of the following types: float32.
35+ *@par Attributes:
36+ * @li centered: Bool. Whether offsets are centered. Default: true.
37+ * @li normalized: Bool. Whether offsets are normalized to [-1, 1]. Default: true.
38+ * @li uniform_noise: Bool. Whether to use uniform noise. Default: true.
39+ * @li noise: String. Noise type. Default: "uniform".
40+ *@par Third-party framework compatibility
41+ *Compatible with the TensorFlow operator ExtractGlimpseV2.
42+ */
43+REG_OP(ExtractGlimpseV2)
44+ .INPUT(input, TensorType({DT_FLOAT}))
45+ .INPUT(size, TensorType({DT_INT32}))
46+ .INPUT(offsets, TensorType({DT_FLOAT}))
47+ .OUTPUT(glimpse, TensorType({DT_FLOAT}))
48+ .ATTR(centered, Bool, true)
49+ .ATTR(normalized, Bool, true)
50+ .ATTR(uniform_noise, Bool, true)
51+ .ATTR(noise, String, "uniform")
52+ .OP_END_FACTORY_REG(ExtractGlimpseV2)
53+ 
54+} // namespace ge
55+ 
56+#endif // OPS_OP_PROTO_INC_EXTRACT_GLIMPSE_V2_H_
@@ -0,0 +1,210 @@
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 extract_glimpse_v2_tiling.cpp
15+ * \brief Tiling implementation for extract_glimpse_v2
16+ */
17+ 
18+#include "log/log.h"
19+#include "platform/platform_ascendc.h"
20+#include "util/math_util.h"
21+#include "util/platform_util.h"
22+#include "op_host/tiling_util.h"
23+#include "op_host/tiling_templates_registry.h"
24+#include "../../op_kernel/arch35/extract_glimpse_v2_tiling_data.h"
25+#include "../../op_kernel/arch35/extract_glimpse_v2_tiling_key.h"
26+ 
27+namespace 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+constexpr uint32_t ATTR_IDX_CENTERED = 0;
33+constexpr uint32_t ATTR_IDX_NORMALIZED = 1;
34+constexpr uint32_t ATTR_IDX_UNIFORM_NOISE = 2;
35+constexpr uint32_t ATTR_IDX_NOISE = 3;
36+constexpr uint32_t SCENE_CENTERED_BIT = 2;
37+constexpr uint32_t SCENE_NORMALIZED_BIT = 1;
38+constexpr uint32_t WORKSPACE_COUNT = 1;
39+ 
40+struct ExtractGlimpseV2CompileInfo {};
41+ 
42+struct ShapeInfo {
43+ int32_t batchSize;
44+ int32_t inputH;
45+ int32_t inputW;
46+ int32_t channels;
47+ int32_t glimpseH;
48+ int32_t glimpseW;
49+ int64_t maxCoreNum;
50+ uint64_t ubSize;
51+};
52+ 
53+static ge::graphStatus GetPlatformAndShapeInfo(gert::TilingContext* context, ShapeInfo& info)
54+{
55+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
56+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
57+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
58+ info.maxCoreNum = ascendcPlatform.GetCoreNumAiv();
59+ OP_CHECK_IF(info.maxCoreNum == 0, OP_LOGE(context, "maxCoreNum is 0"), return ge::GRAPH_FAILED);
60+ info.ubSize = 0;
61+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, info.ubSize);
62+ OP_CHECK_IF(info.ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
63+ 
64+ auto inputShape = context->GetInputShape(0);
65+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
66+ int64_t batchSize = inputShape->GetShape().GetDim(0);
67+ int64_t inputH = inputShape->GetShape().GetDim(1);
68+ int64_t inputW = inputShape->GetShape().GetDim(2);
69+ int64_t channels = inputShape->GetShape().GetDim(3);
70+ 
71+ auto outputShape = context->GetOutputShape(0);
72+ OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);
73+ int64_t glimpseH = outputShape->GetShape().GetDim(1);
74+ int64_t glimpseW = outputShape->GetShape().GetDim(2);
75+ 
liu-wei
liu-weiliu-wei6月27日

🟡 Medium | extract_glimpse_v2 缺少 glimpseH/W > 0 校验

int64_t glimpseH = outputShape->GetShape().GetDim(1);
int64_t glimpseW = outputShape->GetShape().GetDim(2);

OP_CHECK_IF(batchSize > INT32_MAX || inputH > INT32_MAX || inputW > INT32_MAX || channels > INT32_MAX ||
                glimpseH > INT32_MAX || glimpseW > INT32_MAX,
            OP_LOGE(context, "dimension exceeds int32_t range"), return ge::GRAPH_FAILED);

只校验 glimpseH > INT32_MAX没有校验 glimpseH > 0

问题链路

  1. infer_shapeReadSizeData 接受 int32_t 输入,符号扩展到 int64_t 后赋给 glimpseH/W
  2. 如果用户传入 size[0] = -1(或任意负数),glimpseH = -1
  3. infer_shape 设输出 shape 为 (N, -1, -1, C),非法 shape
  4. tilingoutputTotalElements = batchSize * (-1) * (-1) * channels = 正数(巧合)
  5. perCoreElements = (outputTotalElements + maxCoreNum - 1) / maxCoreNum,可能产生奇怪结果
  6. SetBlockDim(needCoreNum)needCoreNum 可能异常

对比 upsample_bicubic2d_tiling_arch35.cpp:102

if (outputHeight <= 0 || outputWidth <= 0) {
    OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(...);
    return ge::GRAPH_FAILED;
}

upsample_bicubic2d 有显式校验,extract_glimpse_v2 缺。

修复建议

OP_CHECK_IF(glimpseH <= 0 || glimpseW <= 0,
            OP_LOGE(context, "glimpseH/W must be > 0"), return ge::GRAPH_FAILED);
likedislike
76+ if (glimpseH <= 0 || glimpseW <= 0) {
77+ std::string valMsg = "[" + std::to_string(glimpseH) + ", " + std::to_string(glimpseW) + "]";
78+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "output_size", valMsg.c_str(),
79+ "output_size elements must be greater than 0");
80+ return ge::GRAPH_FAILED;
81+ }
82+ 
83+ OP_CHECK_IF(batchSize > INT32_MAX || inputH > INT32_MAX || inputW > INT32_MAX || channels > INT32_MAX ||
84+ glimpseH > INT32_MAX || glimpseW > INT32_MAX,
85+ OP_LOGE(context, "dimension exceeds int32_t range"), return ge::GRAPH_FAILED);
86+ 
87+ info.batchSize = static_cast<int32_t>(batchSize);
88+ info.inputH = static_cast<int32_t>(inputH);
89+ info.inputW = static_cast<int32_t>(inputW);
90+ info.channels = static_cast<int32_t>(channels);
91+ info.glimpseH = static_cast<int32_t>(glimpseH);
92+ info.glimpseW = static_cast<int32_t>(glimpseW);
93+ return ge::GRAPH_SUCCESS;
94+}
95+ 
96+static ge::graphStatus ValidateAttributes(gert::TilingContext* context, bool& centered, bool& normalized)
97+{
98+ const gert::RuntimeAttrs* attrs = context->GetAttrs();
99+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
100+ const bool* centeredPtr = attrs->GetBool(ATTR_IDX_CENTERED);
101+ const bool* normalizedPtr = attrs->GetBool(ATTR_IDX_NORMALIZED);
102+ const bool* uniformNoisePtr = attrs->GetBool(ATTR_IDX_UNIFORM_NOISE);
103+ const char* noiseStr = attrs->GetStr(ATTR_IDX_NOISE);
104+ OP_CHECK_IF(centeredPtr == nullptr || normalizedPtr == nullptr || uniformNoisePtr == nullptr || noiseStr == nullptr,
105+ OP_LOGE(context, "attr pointer is null"), return ge::GRAPH_FAILED);
106+ if (*uniformNoisePtr) {
107+ OP_LOGE_FOR_INVALID_VALUE(context->GetNodeName(), "uniform_noise", "true", "false");
108+ return ge::GRAPH_FAILED;
109+ }
110+ if (strcmp(noiseStr, "zero") != 0) {
111+ OP_LOGE_FOR_INVALID_VALUE(context->GetNodeName(), "noise", noiseStr, "zero");
112+ return ge::GRAPH_FAILED;
113+ }
114+ centered = *centeredPtr;
115+ normalized = *normalizedPtr;
116+ return ge::GRAPH_SUCCESS;
117+}
118+ 
119+static ge::graphStatus ComputeAndSetTiling(gert::TilingContext* context, const ShapeInfo& info, bool centered,
120+ bool normalized)
121+{
122+ int64_t outputTotalElements = static_cast<int64_t>(info.batchSize) * info.glimpseH * info.glimpseW * info.channels;
123+ 
124+ int64_t perCoreElements = (outputTotalElements + info.maxCoreNum - 1) / info.maxCoreNum;
125+ if (perCoreElements < PER_CORE_MIN) {
126+ perCoreElements = PER_CORE_MIN;
127+ }
128+ int32_t needCoreNum = static_cast<int32_t>((outputTotalElements + perCoreElements - 1) / perCoreElements);
129+ if (needCoreNum > static_cast<int32_t>(info.maxCoreNum)) {
130+ needCoreNum = static_cast<int32_t>(info.maxCoreNum);
131+ }
132+ if (needCoreNum < 1) {
133+ needCoreNum = 1;
134+ }
135+ 
136+ auto* tiling = context->GetTilingData<ExtractGlimpseV2TilingData>();
137+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
138+ OP_CHECK_IF(memset_s(tiling, sizeof(ExtractGlimpseV2TilingData), 0, sizeof(ExtractGlimpseV2TilingData)) != EOK,
139+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
140+ tiling->needCoreNum = needCoreNum;
141+ tiling->batchSize = info.batchSize;
142+ tiling->inputH = info.inputH;
143+ tiling->inputW = info.inputW;
144+ tiling->channels = info.channels;
145+ tiling->glimpseH = info.glimpseH;
146+ tiling->glimpseW = info.glimpseW;
147+ tiling->centered = centered ? 1 : 0;
148+ tiling->normalized = normalized ? 1 : 0;
149+ tiling->outputTotalElements = outputTotalElements;
150+ tiling->perCoreElements = perCoreElements;
151+ 
152+ context->SetBlockDim(needCoreNum);
153+ uint32_t sceneMode = (centered ? SCENE_CENTERED_BIT : 0) + (normalized ? SCENE_NORMALIZED_BIT : 0);
154+ context->SetTilingKey(GET_TPL_TILING_KEY(sceneMode));
155+ return ge::GRAPH_SUCCESS;
156+}
157+ 
158+static ge::graphStatus SetupMemoryAndWorkspace(gert::TilingContext* context, uint64_t ubSize)
159+{
160+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
161+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
162+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
163+ 
164+ OP_CHECK_IF((ubSize <= DCACHE_SIZE + STATIC_UB_ESTIMATE),
165+ OP_LOGE(context, "ubSize %lu <= DCACHE_SIZE + STATIC_UB_ESTIMATE", ubSize), return ge::GRAPH_FAILED);
166+ auto res = context->SetLocalMemorySize(static_cast<uint32_t>(ubSize - DCACHE_SIZE - STATIC_UB_ESTIMATE));
167+ OP_CHECK_IF((res != ge::GRAPH_SUCCESS), OP_LOGE(context, "SetLocalMemorySize failed"), return ge::GRAPH_FAILED);
168+ 
169+ uint64_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
170+ size_t* ws = context->GetWorkspaceSizes(WORKSPACE_COUNT);
171+ OP_CHECK_NULL_WITH_CONTEXT(context, ws);
172+ ws[0] = static_cast<size_t>(sysWorkspaceSize);
173+ return ge::GRAPH_SUCCESS;
174+}
175+ 
176+static ge::graphStatus ExtractGlimpseV2TilingFunc(gert::TilingContext* context)
177+{
178+ ShapeInfo shapeInfo{};
179+ if (GetPlatformAndShapeInfo(context, shapeInfo) != ge::GRAPH_SUCCESS) {
180+ return ge::GRAPH_FAILED;
181+ }
182+ 
183+ bool centered = false;
184+ bool normalized = false;
185+ if (ValidateAttributes(context, centered, normalized) != ge::GRAPH_SUCCESS) {
186+ return ge::GRAPH_FAILED;
187+ }
188+ 
189+ if (ComputeAndSetTiling(context, shapeInfo, centered, normalized) != ge::GRAPH_SUCCESS) {
190+ return ge::GRAPH_FAILED;
191+ }
192+ 
193+ if (SetupMemoryAndWorkspace(context, shapeInfo.ubSize) != ge::GRAPH_SUCCESS) {
194+ return ge::GRAPH_FAILED;
195+ }
196+ 
197+ return ge::GRAPH_SUCCESS;
198+}
199+ 
200+static ge::graphStatus TilingParseForExtractGlimpseV2([[maybe_unused]] gert::TilingParseContext* context)
201+{
202+ return ge::GRAPH_SUCCESS;
203+}
204+ 
205+IMPL_OP_OPTILING(ExtractGlimpseV2)
206+ .InputsDataDependency({1})
207+ .Tiling(ExtractGlimpseV2TilingFunc)
208+ .TilingParse<ExtractGlimpseV2CompileInfo>(TilingParseForExtractGlimpseV2);
209+ 
210+} // namespace optiling
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时, --simplified_key_mode 选项中填写的值,格式如下所示:
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中,default为默认mode,ascendxx为可选mode,如果不同芯片有差异化要求时,需要配置;
6+; 1)如果没有配置:非ascendC算子继续按空处理,即opc编译命令中不添加 --simplified_key_mode 选项,AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置:各个版本按default配置
8+; 3)如果仅有某些平台的配置,没有default配置:对应平台的按照配置的值传递,非对应平台的:非AscendC算子继续按空处理,AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有:对应平台的使用平台的配置,非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况,需要在binary_simplified_key_mode.ini 文件中显式配置为None,不传入 --simplified_key_mode 选项,由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子,由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[ExtractGlimpseV2]
13+default=0
@@ -0,0 +1,70 @@
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 extract_glimpse_v2_def.cpp
15+ * \brief Operator definition for extract_glimpse_v2
16+ */
17+ 
18+#include "register/op_def_registry.h"
19+ 
20+namespace ops {
21+ 
22+class ExtractGlimpseV2 : public OpDef {
23+public:
24+ explicit ExtractGlimpseV2(const char* name) : OpDef(name)
25+ {
26+ this->Input("input")
27+ .ParamType(REQUIRED)
28+ .DataType({ge::DT_FLOAT})
29+ .Format({ge::FORMAT_ND})
30+ .UnknownShapeFormat({ge::FORMAT_ND})
31+ .AutoContiguous();
32+ this->Input("size")
33+ .ParamType(REQUIRED)
34+ .ValueDepend(OPTIONAL)
35+ .DataType({ge::DT_INT32})
36+ .Format({ge::FORMAT_ND})
37+ .UnknownShapeFormat({ge::FORMAT_ND});
38+ this->Input("offsets")
39+ .ParamType(REQUIRED)
40+ .DataType({ge::DT_FLOAT})
41+ .Format({ge::FORMAT_ND})
42+ .UnknownShapeFormat({ge::FORMAT_ND})
43+ .AutoContiguous();
44+ this->Output("glimpse")
45+ .ParamType(REQUIRED)
46+ .DataType({ge::DT_FLOAT})
47+ .Format({ge::FORMAT_ND})
48+ .UnknownShapeFormat({ge::FORMAT_ND})
49+ .AutoContiguous();
50+ 
51+ this->Attr("centered").AttrType(OPTIONAL).Bool(true);
52+ this->Attr("normalized").AttrType(OPTIONAL).Bool(true);
53+ this->Attr("uniform_noise").AttrType(OPTIONAL).Bool(true);
54+ this->Attr("noise").AttrType(OPTIONAL).String("uniform");
55+ 
56+ OpAICoreConfig aicoreConfig;
57+ aicoreConfig.DynamicCompileStaticFlag(true)
58+ .DynamicFormatFlag(false)
59+ .DynamicRankSupportFlag(true)
60+ .DynamicShapeSupportFlag(true)
61+ .NeedCheckSupportFlag(false)
62+ .PrecisionReduceFlag(false)
63+ .ExtendCfgInfo("opFile.value", "extract_glimpse_v2");
64+ this->AICore().AddConfig("ascend950", aicoreConfig);
65+ }
66+};
67+ 
68+OP_ADD(ExtractGlimpseV2);
69+ 
70+} // namespace ops
@@ -0,0 +1,122 @@
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 extract_glimpse_v2_infershape.cpp
15+ * \brief InferShape implementation for extract_glimpse_v2
16+ */
17+ 
18+#include "register/op_impl_registry.h"
19+#include "log/log.h"
20+ 
21+using namespace ge;
22+ 
23+namespace ops {
24+ 
25+static constexpr int64_t IDX_0 = 0;
26+static constexpr int64_t IDX_1 = 1;
27+static constexpr int64_t IDX_2 = 2;
28+static constexpr int64_t IDX_3 = 3;
29+static constexpr size_t INPUT_REQUIRED_RANK = 4;
30+static constexpr size_t OFFSETS_REQUIRED_RANK = 2;
31+static constexpr int64_t OFFSETS_COORD_DIM = 2;
32+static constexpr size_t SIZE_REQUIRED_RANK = 1;
33+static constexpr int64_t SIZE_ELEMENT_COUNT = 2;
34+static constexpr size_t OUTPUT_RANK = 4;
35+ 
36+static ge::graphStatus ValidateInputShapes(gert::InferShapeContext* context, const gert::Shape* inputShape,
37+ const gert::Shape* sizeShape, const gert::Shape* offsetsShape)
38+{
39+ if (inputShape->GetDimNum() != INPUT_REQUIRED_RANK) {
40+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "input",
41+ (std::to_string(inputShape->GetDimNum()) + "D").c_str(),
42+ (std::to_string(INPUT_REQUIRED_RANK) + "D").c_str());
43+ return ge::GRAPH_FAILED;
44+ }
45+ if (offsetsShape->GetDimNum() != OFFSETS_REQUIRED_RANK) {
46+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "offsets",
47+ (std::to_string(offsetsShape->GetDimNum()) + "D").c_str(),
48+ (std::to_string(OFFSETS_REQUIRED_RANK) + "D").c_str());
49+ return ge::GRAPH_FAILED;
50+ }
51+ if (offsetsShape->GetDim(1) != OFFSETS_COORD_DIM) {
52+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "offsets", std::to_string(offsetsShape->GetDim(1)).c_str(),
53+ std::to_string(OFFSETS_COORD_DIM).c_str());
54+ return ge::GRAPH_FAILED;
55+ }
56+ if (inputShape->GetDim(0) != offsetsShape->GetDim(0)) {
57+ std::string dimMsg = std::to_string(inputShape->GetDim(0)) + " and " + std::to_string(offsetsShape->GetDim(0));
58+ OP_LOGE_FOR_INVALID_SHAPEDIMS_WITH_REASON(context->GetNodeName(), "input and offsets", dimMsg.c_str(),
59+ "batch size of input and offsets must be equal");
60+ return ge::GRAPH_FAILED;
61+ }
62+ if (sizeShape->GetDimNum() != SIZE_REQUIRED_RANK || sizeShape->GetDim(0) != SIZE_ELEMENT_COUNT) {
63+ OP_LOGE_FOR_INVALID_SHAPESIZE(
64+ context->GetNodeName(), "size", (std::to_string(sizeShape->GetDimNum()) + "D").c_str(),
65+ (std::to_string(SIZE_REQUIRED_RANK) + "D with " + std::to_string(SIZE_ELEMENT_COUNT) + " elements")
66+ .c_str());
67+ return ge::GRAPH_FAILED;
68+ }
69+ return ge::GRAPH_SUCCESS;
70+}
71+ 
72+static ge::graphStatus ReadSizeData(gert::InferShapeContext* context, int64_t& glimpseH, int64_t& glimpseW)
73+{
74+ const gert::Tensor* sizeTensor = context->GetInputTensor(IDX_1);
75+ OP_CHECK_NULL_WITH_CONTEXT(context, sizeTensor);
76+ const int32_t* sizeData = static_cast<const int32_t*>(sizeTensor->GetAddr());
77+ OP_CHECK_IF(sizeData == nullptr, OP_LOGE(context, "size data is null"), return ge::GRAPH_FAILED);
78+ glimpseH = static_cast<int64_t>(sizeData[0]);
79+ glimpseW = static_cast<int64_t>(sizeData[1]);
80+ return ge::GRAPH_SUCCESS;
81+}
82+ 
83+static ge::graphStatus SetOutputShapeAndDtype(gert::InferShapeContext* context, const gert::Shape* inputShape,
84+ int64_t glimpseH, int64_t glimpseW)
85+{
86+ gert::Shape* outShape = context->GetOutputShape(IDX_0);
87+ OP_CHECK_NULL_WITH_CONTEXT(context, outShape);
88+ outShape->SetDimNum(OUTPUT_RANK);
89+ outShape->SetDim(0, inputShape->GetDim(0));
90+ outShape->SetDim(IDX_1, glimpseH);
91+ outShape->SetDim(IDX_2, glimpseW);
92+ outShape->SetDim(IDX_3, inputShape->GetDim(IDX_3));
93+ auto* mutableOutDesc = context->GetComputeNodeInfo()->MutableOutputTdInfo(0);
94+ OP_CHECK_NULL_WITH_CONTEXT(context, mutableOutDesc);
95+ mutableOutDesc->SetDataType(context->GetInputDesc(0)->GetDataType());
96+ return ge::GRAPH_SUCCESS;
97+}
98+ 
99+static ge::graphStatus InferShapeExtractGlimpseV2(gert::InferShapeContext* context)
100+{
101+ const gert::Shape* inputShape = context->GetInputShape(IDX_0);
102+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
103+ const gert::Shape* sizeShape = context->GetInputShape(IDX_1);
104+ OP_CHECK_NULL_WITH_CONTEXT(context, sizeShape);
105+ const gert::Shape* offsetsShape = context->GetInputShape(IDX_2);
106+ OP_CHECK_NULL_WITH_CONTEXT(context, offsetsShape);
107+ if (ValidateInputShapes(context, inputShape, sizeShape, offsetsShape) != ge::GRAPH_SUCCESS) {
108+ return ge::GRAPH_FAILED;
109+ }
110+ int64_t glimpseH = 0;
111+ int64_t glimpseW = 0;
112+ if (ReadSizeData(context, glimpseH, glimpseW) != ge::GRAPH_SUCCESS) {
113+ return ge::GRAPH_FAILED;
114+ }
115+ return SetOutputShapeAndDtype(context, inputShape, glimpseH, glimpseW);
116+}
117+ 
118+IMPL_OP_INFERSHAPE(ExtractGlimpseV2)
119+ .InferShape(InferShapeExtractGlimpseV2)
120+ .InputsDataDependency({1});
121+ 
122+} // namespace ops
@@ -0,0 +1,170 @@
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 extract_glimpse_v2_simt.h
15+ * \brief SIMT kernel implementation for extract_glimpse_v2
16+ */
17+ 
18+#ifndef __EXTRACT_GLIMPSE_V2_SIMT_H__
19+#define __EXTRACT_GLIMPSE_V2_SIMT_H__
20+ 
21+#include "kernel_operator.h"
22+#include "kernel_tiling/kernel_tiling.h"
23+#include "simt_api/common_functions.h"
24+#include "simt_api/math_functions.h"
25+#include "extract_glimpse_v2_tiling_data.h"
26+#include "extract_glimpse_v2_tiling_key.h"
27+ 
28+namespace NsExtractGlimpseV2 {
29+ 
30+using namespace AscendC;
31+ 
32+constexpr uint32_t THREADS_UINT32 = 2048;
33+constexpr uint32_t THREADS_UINT64 = 1024;
34+constexpr size_t UINT32_BYTE_WIDTH = 4;
35+constexpr int32_t OFFSETS_STRIDE = 2;
36+constexpr float CENTERED_SCALE = 0.5f;
37+ 
38+template <typename IDX_T>
39+static constexpr uint32_t THREADS = (sizeof(IDX_T) == UINT32_BYTE_WIDTH) ? THREADS_UINT32 : THREADS_UINT64;
40+ 
41+template <typename T, typename IDX_T>
42+__simt_callee__ inline void ProcessPixel(IDX_T idx, int32_t b, int32_t oh, int32_t ow, int32_t c, int32_t inputW,
43+ int32_t channels, int32_t startY, int32_t startX, int32_t baseY, int32_t baseX,
44+ int32_t copyH, int32_t copyW, int64_t inputSliceSize, __gm__ T* input,
45+ __gm__ T* output)
46+{
47+ int32_t localY = oh - baseY;
48+ int32_t localX = ow - baseX;
atomgit-bot
atomgit-botatomgit-bot6月25日

🔵 Low Priority

ProcessPixel 在每像素调用中都会重新计算 startY, startX, endY, endX, baseY, baseX, copyH, copyW 共 8 个值(第 50-61 行),这些值完全由 offsetYoffsetXinputHinputWglimpseHglimpseW 决定,与具体的 ohowc(输出像素坐标和通道)无关。对于同一个 batch item 的所有像素(如 64×64×3 = 12288 像素),这 8 个值完全相同却被重复计算了上万次。

建议将 startY, startX, endY, endX, baseY, baseX, copyH, copyW 的计算提升到 OpExtractGlimpseV2SimtKernel 中(在计算完 offsetY/offsetX 之后),并作为参数传入 ProcessPixel,消除循环内的冗余浮点运算。

建议:将 startY/startX/endY/endX/baseY/baseX/copyH/copyW 的计算从 ProcessPixel 提升到 OpExtractGlimpseV2SimtKernel 中(在调用 ProcessPixel 之前),作为额外参数传入,避免每个输出像素重复计算相同的边界值。

likedislike
ideal_
ideal_
6月25日 评论:
49+ if (copyH > 0 && copyW > 0 && localY >= 0 && localY < copyH && localX >= 0 && localX < copyW) {
50+ int32_t srcY = startY + localY;
51+ int32_t srcX = startX + localX;
52+ int64_t inputIdx = static_cast<int64_t>(b) * inputSliceSize + static_cast<int64_t>(srcY) * inputW * channels +
53+ static_cast<int64_t>(srcX) * channels + c;
54+ output[idx] = input[inputIdx];
55+ } else {
56+ output[idx] = static_cast<T>(0);
57+ }
58+}
59+ 
60+template <typename T, bool Centered, bool Normalized, typename IDX_T>
61+__simt_vf__ __aicore__ __launch_bounds__(THREADS<IDX_T>) inline void OpExtractGlimpseV2SimtKernel(
62+ IDX_T outputTotalElements, int32_t inputH, int32_t inputW, int32_t channels, int32_t glimpseH, int32_t glimpseW,
63+ IDX_T magicChannels, IDX_T shiftChannels, IDX_T magicGlimpseW, IDX_T shiftGlimpseW, IDX_T magicGlimpseH,
64+ IDX_T shiftGlimpseH, __gm__ T* input, __gm__ float* offsets, __gm__ T* output)
65+{
66+ const int64_t inputSliceSize = static_cast<int64_t>(inputH) * inputW * channels;
67+ 
68+ for (IDX_T idx = static_cast<IDX_T>(blockIdx.x * blockDim.x + threadIdx.x); idx < outputTotalElements;
69+ idx += static_cast<IDX_T>(blockDim.x * gridDim.x)) {
70+ IDX_T remaining = idx;
71+ 
liu-wei
liu-weiliu-wei6月27日

🟢 Low | Process 函数接收 sizeworkspace 参数但未使用

template <typename T, bool Centered, bool Normalized>
__aicore__ inline void Process(GM_ADDR input, GM_ADDR size, GM_ADDR offsets, GM_ADDR output, GM_ADDR workspace,
                               const ExtractGlimpseV2TilingData* tilingData)
{
    int64_t outputTotalElements = tilingData->outputTotalElements;
    if (outputTotalElements == 0) {
        return;
    }
    // size 和 workspace 未使用
    ...
}

size 是算子输入(inputSize tensor,存放在 GM),workspace 是 workspace 内存。ExtractGlimpseV2 算子的 kernel 不需要这些:

  • size 信息已在 tiling 时通过 outputShape 提取(glimpseH/W),传给 tilingData
  • workspace kernel 不需要

这是设计选择,因为算子签名与 TensorFlow 算子签名一致(input, size, offsets, output)。kernel 只用 input/offsets/output

修复建议

  • 添加注释说明 sizeworkspace 在 SIMT 实现中不需要
  • 或用 [[maybe_unused]] 标记参数(但 GM_ADDR 是 typedef,不容易用)
likedislike
ideal_
ideal_
6月27日 评论:
72+ IDX_T q1 = Simt::UintDiv<IDX_T>(remaining, magicChannels, shiftChannels);
73+ int32_t c = static_cast<int32_t>(remaining - q1 * static_cast<IDX_T>(channels));
74+ remaining = q1;
75+ 
76+ IDX_T q2 = Simt::UintDiv<IDX_T>(remaining, magicGlimpseW, shiftGlimpseW);
77+ int32_t ow = static_cast<int32_t>(remaining - q2 * static_cast<IDX_T>(glimpseW));
78+ remaining = q2;
79+ 
80+ IDX_T q3 = Simt::UintDiv<IDX_T>(remaining, magicGlimpseH, shiftGlimpseH);
81+ int32_t oh = static_cast<int32_t>(remaining - q3 * static_cast<IDX_T>(glimpseH));
82+ int32_t b = static_cast<int32_t>(q3);
83+ 
84+ float offsetY = offsets[b * OFFSETS_STRIDE];
85+ float offsetX = offsets[b * OFFSETS_STRIDE + 1];
86+ 
87+ if (offsetY != offsetY || offsetX != offsetX) {
88+ output[idx] = static_cast<T>(0);
89+ continue;
90+ }
91+ 
92+ if constexpr (Normalized) {
93+ offsetY = offsetY * static_cast<float>(inputH);
94+ offsetX = offsetX * static_cast<float>(inputW);
95+ }
96+ if constexpr (Centered) {
97+ if constexpr (Normalized) {
98+ offsetY = offsetY * CENTERED_SCALE + static_cast<float>(inputH) * CENTERED_SCALE -
99+ static_cast<float>(glimpseH) * CENTERED_SCALE;
100+ offsetX = offsetX * CENTERED_SCALE + static_cast<float>(inputW) * CENTERED_SCALE -
101+ static_cast<float>(glimpseW) * CENTERED_SCALE;
102+ } else {
103+ offsetY = offsetY + static_cast<float>(inputH) * CENTERED_SCALE;
104+ offsetX = offsetX + static_cast<float>(inputW) * CENTERED_SCALE;
105+ }
106+ }
107+ 
108+ int32_t startY = static_cast<int32_t>(floorf(fmaxf(0.0f, offsetY)));
109+ int32_t startX = static_cast<int32_t>(floorf(fmaxf(0.0f, offsetX)));
110+ int32_t endY = static_cast<int32_t>(
111+ ceilf(fminf(static_cast<float>(inputH), offsetY + static_cast<float>(glimpseH))));
112+ int32_t endX = static_cast<int32_t>(
113+ ceilf(fminf(static_cast<float>(inputW), offsetX + static_cast<float>(glimpseW))));
114+ int32_t baseY = static_cast<int32_t>(floorf(fmaxf(0.0f, -offsetY)));
115+ int32_t baseX = static_cast<int32_t>(floorf(fmaxf(0.0f, -offsetX)));
116+ int32_t copyH = endY - startY;
117+ int32_t copyW = endX - startX;
118+ 
119+ ProcessPixel<T, IDX_T>(idx, b, oh, ow, c, inputW, channels, startY, startX, baseY, baseX, copyH, copyW,
120+ inputSliceSize, input, output);
121+ }
122+}
123+ 
124+template <typename T, bool Centered, bool Normalized>
125+__aicore__ inline void Process(GM_ADDR input, GM_ADDR size, GM_ADDR offsets, GM_ADDR output, GM_ADDR workspace,
126+ const ExtractGlimpseV2TilingData* tilingData)
127+{
128+ int64_t outputTotalElements = tilingData->outputTotalElements;
129+ if (outputTotalElements == 0) {
130+ return;
131+ }
132+ 
133+ __gm__ T* inputGm = (__gm__ T*)input;
134+ __gm__ float* offsetsGm = (__gm__ float*)offsets;
135+ __gm__ T* outputGm = (__gm__ T*)output;
136+ 
137+ int32_t inputH = tilingData->inputH;
138+ int32_t inputW = tilingData->inputW;
139+ int32_t channels = tilingData->channels;
140+ int32_t glimpseH = tilingData->glimpseH;
141+ int32_t glimpseW = tilingData->glimpseW;
142+ 
143+ if (outputTotalElements <= static_cast<int64_t>(INT32_MAX)) {
144+ uint32_t mCh = 0, sCh = 0;
145+ uint32_t mGW = 0, sGW = 0;
146+ uint32_t mGH = 0, sGH = 0;
147+ GetUintDivMagicAndShift<uint32_t>(mCh, sCh, static_cast<uint32_t>(channels));
148+ GetUintDivMagicAndShift<uint32_t>(mGW, sGW, static_cast<uint32_t>(glimpseW));
149+ GetUintDivMagicAndShift<uint32_t>(mGH, sGH, static_cast<uint32_t>(glimpseH));
150+ 
151+ asc_vf_call<OpExtractGlimpseV2SimtKernel<T, Centered, Normalized, uint32_t>>(
152+ dim3(THREADS<uint32_t>), static_cast<uint32_t>(outputTotalElements), inputH, inputW, channels, glimpseH,
153+ glimpseW, mCh, sCh, mGW, sGW, mGH, sGH, inputGm, offsetsGm, outputGm);
154+ } else {
155+ uint64_t mCh = 0, sCh = 0;
156+ uint64_t mGW = 0, sGW = 0;
157+ uint64_t mGH = 0, sGH = 0;
158+ GetUintDivMagicAndShift<uint64_t>(mCh, sCh, static_cast<uint64_t>(channels));
159+ GetUintDivMagicAndShift<uint64_t>(mGW, sGW, static_cast<uint64_t>(glimpseW));
160+ GetUintDivMagicAndShift<uint64_t>(mGH, sGH, static_cast<uint64_t>(glimpseH));
161+ 
162+ asc_vf_call<OpExtractGlimpseV2SimtKernel<T, Centered, Normalized, uint64_t>>(
163+ dim3(THREADS<uint64_t>), static_cast<uint64_t>(outputTotalElements), inputH, inputW, channels, glimpseH,
164+ glimpseW, mCh, sCh, mGW, sGW, mGH, sGH, inputGm, offsetsGm, outputGm);
165+ }
166+}
167+ 
168+} // namespace NsExtractGlimpseV2
169+ 
170+#endif // __EXTRACT_GLIMPSE_V2_SIMT_H__
@@ -0,0 +1,35 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file extract_glimpse_v2_tiling_data.h
15+ * \brief tiling data struct for extract_glimpse_v2
16+ */
17+ 
18+#ifndef __EXTRACT_GLIMPSE_V2_TILING_DATA_H__
19+#define __EXTRACT_GLIMPSE_V2_TILING_DATA_H__
20+ 
21+struct ExtractGlimpseV2TilingData {
22+ int32_t needCoreNum;
23+ int32_t batchSize;
24+ int32_t inputH;
25+ int32_t inputW;
26+ int32_t channels;
27+ int32_t glimpseH;
28+ int32_t glimpseW;
29+ int32_t centered;
30+ int32_t normalized;
31+ int64_t outputTotalElements;
32+ int64_t perCoreElements;
33+};
34+ 
35+#endif // __EXTRACT_GLIMPSE_V2_TILING_DATA_H__
@@ -0,0 +1,37 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file extract_glimpse_v2_tiling_key.h
15+ * \brief tiling key declaration for extract_glimpse_v2
16+ */
17+ 
18+#ifndef __EXTRACT_GLIMPSE_V2_TILING_KEY_H__
19+#define __EXTRACT_GLIMPSE_V2_TILING_KEY_H__
20+ 
21+#include "ascendc/host_api/tiling/template_argument.h"
22+ 
23+#define SCENE_MODE_0 0
24+#define SCENE_MODE_1 1
25+#define SCENE_MODE_2 2
26+#define SCENE_MODE_3 3
27+ 
28+ASCENDC_TPL_ARGS_DECL(
29+ ExtractGlimpseV2,
30+ ASCENDC_TPL_UINT_DECL(schMode, 2, ASCENDC_TPL_UI_LIST,
31+ SCENE_MODE_0, SCENE_MODE_1, SCENE_MODE_2, SCENE_MODE_3));
32+ 
33+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
34+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,
35+ SCENE_MODE_0, SCENE_MODE_1, SCENE_MODE_2, SCENE_MODE_3)));
36+ 
37+#endif // __EXTRACT_GLIMPSE_V2_TILING_KEY_H__
@@ -0,0 +1,41 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file extract_glimpse_v2.cpp
15+ * \brief kernel entry for extract_glimpse_v2
16+ */
17+ 
18+#include "arch35/extract_glimpse_v2_simt.h"
19+ 
20+template <uint32_t schMode>
21+__global__ __aicore__ void extract_glimpse_v2(
22+ GM_ADDR input, GM_ADDR size, GM_ADDR offsets,
23+ GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling)
24+{
25+ REGISTER_TILING_DEFAULT(ExtractGlimpseV2TilingData);
26+ GET_TILING_DATA_WITH_STRUCT(ExtractGlimpseV2TilingData, tilingData, tiling);
27+ 
28+ if constexpr (schMode == SCENE_MODE_0) {
29+ NsExtractGlimpseV2::Process<DTYPE_INPUT, false, false>(
30+ input, size, offsets, output, workspace, &tilingData);
31+ } else if constexpr (schMode == SCENE_MODE_1) {
32+ NsExtractGlimpseV2::Process<DTYPE_INPUT, false, true>(
33+ input, size, offsets, output, workspace, &tilingData);
34+ } else if constexpr (schMode == SCENE_MODE_2) {
35+ NsExtractGlimpseV2::Process<DTYPE_INPUT, true, false>(
36+ input, size, offsets, output, workspace, &tilingData);
37+ } else {
38+ NsExtractGlimpseV2::Process<DTYPE_INPUT, true, true>(
39+ input, size, offsets, output, workspace, &tilingData);
40+ }
41+}
@@ -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()
@@ -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()
@@ -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(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}/arch35)
17+ endif()
18+ add_modules_ut_sources(UT_NAME ${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()
@@ -0,0 +1,181 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+#include <iostream>
14+#include <gtest/gtest.h>
15+#include "tiling_context_faker.h"
16+#include "tiling_case_executor.h"
17+#include "../../../../op_kernel/arch35/extract_glimpse_v2_tiling_data.h"
18+ 
19+using namespace std;
20+using namespace ge;
21+ 
22+class ExtractGlimpseV2Tiling : public testing::Test {
23+protected:
24+ static void SetUpTestCase()
25+ {
26+ std::cout << "ExtractGlimpseV2Tiling SetUp" << std::endl;
27+ }
28+ 
29+ static void TearDownTestCase()
30+ {
31+ std::cout << "ExtractGlimpseV2Tiling TearDown" << std::endl;
32+ }
33+};
34+ 
35+static std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend950"}};
36+ 
37+TEST_F(ExtractGlimpseV2Tiling, extract_glimpse_v2_0)
38+{
39+ struct ExtractGlimpseV2CompileInfo {
40+ } compileInfo;
41+ gert::TilingContextPara tilingContextPara(
42+ "ExtractGlimpseV2",
43+ {
44+ {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
45+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND},
46+ {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
47+ },
48+ {
49+ {{{1, 2, 2, 3}, {1, 2, 2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
50+ },
51+ {
52+ gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
53+ gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
54+ gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
55+ gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("zero")),
56+ },
57+ &compileInfo,
58+ "Ascend950",
59+ 64,
60+ 262144,
61+ 4096);
62+ uint64_t expectTilingKey = 3;
63+ string expectTilingData = "4294967297 12884901891 8589934595 4294967298 1 12 1024 ";
64+ std::vector<size_t> expectWorkspaces = {16777216};
65+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
66+}
67+ 
68+TEST_F(ExtractGlimpseV2Tiling, extract_glimpse_v2_1)
69+{
70+ struct ExtractGlimpseV2CompileInfo {
71+ } compileInfo;
72+ gert::TilingContextPara tilingContextPara(
73+ "ExtractGlimpseV2",
74+ {
75+ {{{15, 13, 16, 16}, {15, 13, 16, 16}}, ge::DT_FLOAT, ge::FORMAT_ND},
76+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND},
77+ {{{15, 2}, {15, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
78+ },
79+ {
80+ {{{15, 2, 9, 16}, {15, 2, 9, 16}}, ge::DT_FLOAT, ge::FORMAT_ND},
81+ },
82+ {
83+ gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
84+ gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
85+ gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
86+ gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("zero")),
87+ },
88+ &compileInfo,
89+ "Ascend950",
90+ 64,
91+ 262144,
92+ 4096);
93+ uint64_t expectTilingKey = 0;
94+ string expectTilingData = "64424509445 68719476749 8589934608 9 0 4320 1024 ";
95+ std::vector<size_t> expectWorkspaces = {16777216};
96+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
97+}
98+ 
99+TEST_F(ExtractGlimpseV2Tiling, fail_uniform_noise_true)
100+{
101+ struct ExtractGlimpseV2CompileInfo {
102+ } compileInfo;
103+ gert::TilingContextPara tilingContextPara(
104+ "ExtractGlimpseV2",
105+ {
106+ {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
107+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND},
108+ {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
109+ },
110+ {
111+ {{{1, 2, 2, 3}, {1, 2, 2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
112+ },
113+ {
114+ gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
115+ gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
116+ gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
117+ gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("zero")),
118+ },
119+ &compileInfo,
120+ "Ascend950",
121+ 64,
122+ 262144,
123+ 4096);
124+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, 0, "", {});
125+}
126+ 
127+TEST_F(ExtractGlimpseV2Tiling, fail_noise_not_zero)
128+{
129+ struct ExtractGlimpseV2CompileInfo {
130+ } compileInfo;
131+ gert::TilingContextPara tilingContextPara(
132+ "ExtractGlimpseV2",
133+ {
134+ {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
135+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND},
136+ {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
137+ },
138+ {
139+ {{{1, 2, 2, 3}, {1, 2, 2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
140+ },
141+ {
142+ gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
143+ gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
144+ gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
145+ gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("uniform")),
146+ },
147+ &compileInfo,
148+ "Ascend950",
149+ 64,
150+ 262144,
151+ 4096);
152+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, 0, "", {});
153+}
154+ 
155+TEST_F(ExtractGlimpseV2Tiling, fail_noise_gaussian)
156+{
157+ struct ExtractGlimpseV2CompileInfo {
158+ } compileInfo;
159+ gert::TilingContextPara tilingContextPara(
160+ "ExtractGlimpseV2",
161+ {
162+ {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
163+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND},
164+ {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
165+ },
166+ {
167+ {{{1, 2, 2, 3}, {1, 2, 2, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
168+ },
169+ {
170+ gert::TilingContextPara::OpAttr("centered", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
171+ gert::TilingContextPara::OpAttr("normalized", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
172+ gert::TilingContextPara::OpAttr("uniform_noise", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
173+ gert::TilingContextPara::OpAttr("noise", Ops::Cv::AnyValue::CreateFrom<std::string>("gaussian")),
174+ },
175+ &compileInfo,
176+ "Ascend950",
177+ 64,
178+ 262144,
179+ 4096);
180+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, 0, "", {});
181+}
@@ -0,0 +1,68 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+#include <gtest/gtest.h>
14+#include <iostream>
15+#include "infershape_context_faker.h"
16+#include "infershape_case_executor.h"
17+ 
18+class ExtractGlimpseV2Infershape : public testing::Test
19+{
20+protected:
21+ static void SetUpTestCase()
22+ {
23+ std::cout << "ExtractGlimpseV2Infershape SetUp" << std::endl;
24+ }
25+ 
26+ static void TearDownTestCase()
27+ {
28+ std::cout << "ExtractGlimpseV2Infershape TearDown" << std::endl;
29+ }
30+};
31+ 
32+TEST_F(ExtractGlimpseV2Infershape, extract_glimpse_v2_infershape_test1)
33+{
34+ int32_t sizeData1[2] = {2, 2};
35+ gert::InfershapeContextPara infershapeContextPara(
36+ "ExtractGlimpseV2",
37+ {
38+ {{{1, 3, 3, 3}, {1, 3, 3, 3}}, ge::DT_FLOAT, ge::FORMAT_ND},
39+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND, true, sizeData1},
40+ {{{1, 2}, {1, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
41+ },
42+ {
43+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
44+ });
45+ std::vector<std::vector<int64_t>> expectOutputShape = {
46+ {1, 2, 2, 3},
47+ };
48+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
49+}
50+ 
51+TEST_F(ExtractGlimpseV2Infershape, extract_glimpse_v2_infershape_test2)
52+{
53+ int32_t sizeData2[2] = {2, 9};
54+ gert::InfershapeContextPara infershapeContextPara(
55+ "ExtractGlimpseV2",
56+ {
57+ {{{15, 13, 16, 16}, {15, 13, 16, 16}}, ge::DT_FLOAT, ge::FORMAT_ND},
58+ {{{2}, {2}}, ge::DT_INT32, ge::FORMAT_ND, true, sizeData2},
59+ {{{15, 2}, {15, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
60+ },
61+ {
62+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
63+ });
64+ std::vector<std::vector<int64_t>> expectOutputShape = {
65+ {15, 2, 9, 16},
66+ };
67+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
68+}
@@ -4,7 +4,7 @@
4 4 
5|产品 | 是否支持 |5|产品 | 是否支持 |
6|:-------------------------|:----------:|6|:-------------------------|:----------:|
7-| <term>Ascend 950PR/Ascend 950DT</term> | × |7+| <term>Ascend 950PR/Ascend 950DT</term> | |
8| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |8| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |9| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10| <term>Atlas 200I/500 A2 推理产品</term> | √ |10| <term>Atlas 200I/500 A2 推理产品</term> | √ |
@@ -8,6 +8,8 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11+/* Generated By CANNBot */
12+ 
11/*!13/*!
12 * \file upsample_bicubic2d_def.cpp14 * \file upsample_bicubic2d_def.cpp
13 * \brief15 * \brief
@@ -38,6 +40,15 @@ public:
38 this->AICore().AddConfig("ascend910b");40 this->AICore().AddConfig("ascend910b");
39 this->AICore().AddConfig("ascend910_93");41 this->AICore().AddConfig("ascend910_93");
40 42 
43+ OpAICoreConfig aicoreConfig;
44+ aicoreConfig.DynamicCompileStaticFlag(true)
45+ .DynamicRankSupportFlag(true)
46+ .DynamicShapeSupportFlag(true)
47+ .DynamicFormatFlag(true)
48+ .NeedCheckSupportFlag(false)
49+ .ExtendCfgInfo("opFile.value", "upsample_bicubic2d_apt");
50+ this->AICore().AddConfig("ascend950", aicoreConfig);
51+ 
41 OpAICoreConfig config310p;52 OpAICoreConfig config310p;
42 config310p.DynamicCompileStaticFlag(true)53 config310p.DynamicCompileStaticFlag(true)
43 .DynamicFormatFlag(true)54 .DynamicFormatFlag(true)
@@ -8,6 +8,8 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11+/* Generated By CANNBot */
12+ 
11/*!13/*!
12 * \file upsample_bicubic2d_tiling.cpp14 * \file upsample_bicubic2d_tiling.cpp
13 * \brief15 * \brief
@@ -18,7 +20,11 @@
18#include "log/log.h"20#include "log/log.h"
19#include "tiling/platform/platform_ascendc.h"21#include "tiling/platform/platform_ascendc.h"
20#include "upsample_bicubic2d_tiling.h"22#include "upsample_bicubic2d_tiling.h"
21-#include "op_host/tiling_base.h"23+#include "op_host/tiling_util.h"
24+ 
25+namespace optiling {
26+ge::graphStatus UpsampleBicubic2dTilingFunc(gert::TilingContext* context);
27+}
22 28 
23namespace optiling {29namespace optiling {
24constexpr uint32_t BEST_PERFORMANCE_SIZE_1 = 16;30constexpr uint32_t BEST_PERFORMANCE_SIZE_1 = 16;
@@ -583,6 +589,10 @@ void UpsampleBicubic2dTiling::FillTilingData()
583 589 
584static ge::graphStatus tiling4UpsampleBicubic2dTiling(gert::TilingContext* context)590static ge::graphStatus tiling4UpsampleBicubic2dTiling(gert::TilingContext* context)
585{591{
592+ bool regBase = Ops::Cv::OpTiling::IsRegbaseSocVersion(context);
593+ if (regBase) {
594+ return optiling::UpsampleBicubic2dTilingFunc(context);
595+ }
586 UpsampleBicubic2dTiling tilingObject(context);596 UpsampleBicubic2dTiling tilingObject(context);
587 context->SetScheduleMode(SCHEDULE_MODE);597 context->SetScheduleMode(SCHEDULE_MODE);
588 return tilingObject.RunBigKernelTiling();598 return tilingObject.RunBigKernelTiling();
@@ -0,0 +1,228 @@
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 upsample_bicubic2d_tiling.cpp
15+ * \brief tiling implementation for upsample_bicubic2d
16+ */
17+ 
18+#include "log/log.h"
19+#include "tiling/platform/platform_ascendc.h"
20+#include "util/math_util.h"
21+#include "util/platform_util.h"
22+#include "op_host/tiling_util.h"
23+#include "op_host/tiling_templates_registry.h"
24+#include "../op_kernel/arch35/upsample_bicubic2d_tiling_data.h"
25+#include "../op_kernel/arch35/upsample_bicubic2d_tiling_key.h"
26+ 
27+namespace optiling {
28+ 
29+using namespace Ops::Cv::OpTiling;
30+ 
31+constexpr uint32_t DCACHE_SIZE = 128 * 1024;
32+constexpr uint32_t STATIC_UB_ESTIMATE = 0;
33+constexpr int64_t PER_CORE_MIN = 1024;
34+constexpr int32_t INPUT_DIM_NUM = 4;
35+constexpr int64_t PER_CORE_ALIGN = 32;
36+constexpr int64_t ALIGN_MASK = PER_CORE_ALIGN - 1;
37+constexpr int32_t OUTPUT_SIZE_ATTR_COUNT = 2;
38+constexpr uint32_t ATTR_IDX_OUTPUT_SIZE = 0;
39+constexpr uint32_t ATTR_IDX_ALIGN_CORNERS = 1;
40+constexpr uint32_t ATTR_IDX_SCALES_H = 2;
41+constexpr uint32_t ATTR_IDX_SCALES_W = 3;
42+ 
43+struct UpsampleBicubic2dCompileInfo {};
44+ 
45+// 获取平台信息
46+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
47+{
48+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
49+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
50+ 
51+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
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+ 
55+ coreNum = ascendcPlatform.GetCoreNumAiv();
56+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
57+
58+ return ge::GRAPH_SUCCESS;
59+}
60+ 
61+// 获取shape和属性信息
62+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int32_t& batchSize, int32_t& channels,
63+ int32_t& inputHeight, int32_t& inputWidth, int32_t& outputHeight,
64+ int32_t& outputWidth, bool& alignCorners, float& scalesH, float& scalesW)
65+{
66+ // 获取输入shape
67+ auto inputX = context->GetInputShape(0);
68+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
69+ auto inputShapeX = inputX->GetStorageShape();
70+ 
71+ if (inputShapeX.GetDimNum() != INPUT_DIM_NUM) {
72+ OP_LOGE_FOR_INVALID_SHAPEDIM(context->GetNodeName(), "x",
73+ (std::to_string(inputShapeX.GetDimNum()) + "D").c_str(),
74+ (std::to_string(INPUT_DIM_NUM) + "D").c_str());
75+ return ge::GRAPH_FAILED;
76+ }
77+ 
78+ batchSize = static_cast<int32_t>(inputShapeX.GetDim(0));
79+ channels = static_cast<int32_t>(inputShapeX.GetDim(1));
80+ inputHeight = static_cast<int32_t>(inputShapeX.GetDim(2));
81+ inputWidth = static_cast<int32_t>(inputShapeX.GetDim(3));
82+ 
83+ // 获取属性
84+ const gert::RuntimeAttrs* attrs = context->GetAttrs();
85+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
86+ 
87+ // output_size (REQUIRED_ATTR, ListInt, index 0)
88+ const auto* outputSizeAttr = attrs->GetListInt(ATTR_IDX_OUTPUT_SIZE);
89+ OP_CHECK_NULL_WITH_CONTEXT(context, outputSizeAttr);
90+ if (outputSizeAttr->GetSize() != OUTPUT_SIZE_ATTR_COUNT) {
91+ OP_LOGE_FOR_INVALID_LISTSIZE(context->GetNodeName(), "output_size",
92+ std::to_string(outputSizeAttr->GetSize()).c_str(),
93+ std::to_string(OUTPUT_SIZE_ATTR_COUNT).c_str());
94+ return ge::GRAPH_FAILED;
95+ }
96+ const int64_t* outputSizeData = outputSizeAttr->GetData();
97+ outputHeight = static_cast<int32_t>(outputSizeData[0]);
98+ outputWidth = static_cast<int32_t>(outputSizeData[1]);
99+ if (outputHeight <= 0 || outputWidth <= 0) {
100+ std::string valMsg = "[" + std::to_string(outputHeight) + ", " + std::to_string(outputWidth) + "]";
101+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "output_size", valMsg.c_str(),
102+ "output_size elements must be greater than 0");
103+ return ge::GRAPH_FAILED;
104+ }
105+ 
106+ // align_corners (ATTR, Bool, index 1)
107+ const bool* alignCornersPtr = attrs->GetBool(ATTR_IDX_ALIGN_CORNERS);
108+ OP_CHECK_NULL_WITH_CONTEXT(context, alignCornersPtr);
109+ alignCorners = *alignCornersPtr;
110+ 
111+ // scales_h (ATTR, Float, index 2)
112+ const float* scalesHPtr = attrs->GetFloat(ATTR_IDX_SCALES_H);
113+ OP_CHECK_NULL_WITH_CONTEXT(context, scalesHPtr);
114+ scalesH = *scalesHPtr;
115+ 
116+ // scales_w (ATTR, Float, index 3)
117+ const float* scalesWPtr = attrs->GetFloat(ATTR_IDX_SCALES_W);
118+ OP_CHECK_NULL_WITH_CONTEXT(context, scalesWPtr);
119+ scalesW = *scalesWPtr;
120+ 
121+ return ge::GRAPH_SUCCESS;
122+}
123+ 
124+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
125+{
126+ int64_t userWorkspaceSize = 0;
127+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
128+ uint64_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
129+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
130+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
131+ currentWorkspace[0] = static_cast<size_t>(userWorkspaceSize + static_cast<int64_t>(sysWorkspaceSize));
132+ return ge::GRAPH_SUCCESS;
133+}
134+ 
135+static void ComputeScaleFactors(int32_t inputHeight, int32_t inputWidth, int32_t outputHeight, int32_t outputWidth,
136+ bool alignCorners, float& actualScaleH, float& actualScaleW)
137+{
138+ if (alignCorners) {
139+ actualScaleH = (outputHeight > 1) ? static_cast<float>(inputHeight - 1) / static_cast<float>(outputHeight - 1) :
140+ 0.0f;
141+ actualScaleW = (outputWidth > 1) ? static_cast<float>(inputWidth - 1) / static_cast<float>(outputWidth - 1) :
142+ 0.0f;
143+ } else {
144+ actualScaleH = static_cast<float>(inputHeight) / static_cast<float>(outputHeight);
145+ actualScaleW = static_cast<float>(inputWidth) / static_cast<float>(outputWidth);
146+ }
147+}
148+ 
149+static ge::graphStatus SetTilingAndCoreSplit(gert::TilingContext* context, int32_t batchSize, int32_t channels,
150+ int32_t inputHeight, int32_t inputWidth, int32_t outputHeight,
151+ int32_t outputWidth, float actualScaleH, float actualScaleW,
152+ int64_t coreNum)
153+{
154+ UpsampleBicubic2dRegBaseTilingData* tiling = context->GetTilingData<UpsampleBicubic2dRegBaseTilingData>();
155+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
156+ OP_CHECK_IF(memset_s(tiling, sizeof(UpsampleBicubic2dRegBaseTilingData), 0,
157+ sizeof(UpsampleBicubic2dRegBaseTilingData)) != EOK,
158+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
159+ 
160+ int64_t totalElements = static_cast<int64_t>(batchSize) * channels * outputHeight * outputWidth;
161+ tiling->totalElements = totalElements;
162+ 
163+ int64_t perCoreElements = Ops::Base::CeilDiv(totalElements, coreNum);
164+ if (perCoreElements < PER_CORE_MIN) {
165+ perCoreElements = PER_CORE_MIN;
166+ }
167+ perCoreElements = ((perCoreElements + ALIGN_MASK) / PER_CORE_ALIGN) * PER_CORE_ALIGN;
168+ int64_t needCoreNum = Ops::Base::CeilDiv(totalElements, perCoreElements);
169+ if (needCoreNum < 1) {
170+ needCoreNum = 1;
171+ }
172+ tiling->needCoreNum = static_cast<int32_t>(needCoreNum);
173+ context->SetBlockDim(static_cast<uint32_t>(needCoreNum));
174+ tiling->ncSize = batchSize * channels;
175+ tiling->inputHeight = inputHeight;
176+ tiling->inputWidth = inputWidth;
177+ tiling->outputHeight = outputHeight;
178+ tiling->outputWidth = outputWidth;
179+ tiling->scaleH = actualScaleH;
180+ tiling->scaleW = actualScaleW;
181+ return ge::GRAPH_SUCCESS;
182+}
183+ 
184+static ge::graphStatus SetMemoryAndTilingKey(gert::TilingContext* context, uint64_t ubSize, bool alignCorners)
185+{
186+ OP_CHECK_IF((ubSize <= DCACHE_SIZE + STATIC_UB_ESTIMATE),
187+ OP_LOGE(context, "ubSize %lu <= DCACHE_SIZE + STATIC_UB_ESTIMATE", ubSize), return ge::GRAPH_FAILED);
188+ auto res = context->SetLocalMemorySize(static_cast<uint32_t>(ubSize - DCACHE_SIZE - STATIC_UB_ESTIMATE));
189+ OP_CHECK_IF((res != ge::GRAPH_SUCCESS), OP_LOGE(context, "SetLocalMemorySize failed"), return ge::GRAPH_FAILED);
190+ 
191+ uint64_t tilingKey = GET_TPL_TILING_KEY(static_cast<uint64_t>(
192+ alignCorners ? UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE : UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE));
193+ context->SetTilingKey(tilingKey);
194+ return ge::GRAPH_SUCCESS;
liu-wei
liu-weiliu-wei6月27日

🟢 Low | static_cast<uint64_t> 三目运算结果是冗余 cast

uint64_t tilingKey = GET_TPL_TILING_KEY(static_cast<uint64_t>(
    alignCorners ? UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE : UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE));

UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE/FALSE#define 宏(值为 0/1),三目运算结果类型是 intstatic_cast<uint64_t>(int) 是隐式转换,explicit cast 是冗余的。

实际影响:无功能影响,但代码可读性差。

修复建议

uint64_t schMode = alignCorners ? UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE
                                : UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE;
uint64_t tilingKey = GET_TPL_TILING_KEY(schMode);
likedislike
ideal_
ideal_
6月27日 评论:
195+}
196+ 
197+ge::graphStatus UpsampleBicubic2dTilingFunc(gert::TilingContext* context)
198+{
199+ uint64_t ubSize;
200+ int64_t coreNum;
201+ OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
202+ OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
203+ 
204+ int32_t batchSize, channels, inputHeight, inputWidth;
205+ int32_t outputHeight, outputWidth;
206+ bool alignCorners;
207+ float scalesH, scalesW;
208+ OP_CHECK_IF(GetShapeAttrsInfo(context, batchSize, channels, inputHeight, inputWidth, outputHeight, outputWidth,
209+ alignCorners, scalesH, scalesW) != ge::GRAPH_SUCCESS,
210+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
211+ 
212+ OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
213+ return ge::GRAPH_FAILED);
214+ 
215+ float actualScaleH, actualScaleW;
216+ ComputeScaleFactors(inputHeight, inputWidth, outputHeight, outputWidth, alignCorners, actualScaleH, actualScaleW);
217+ 
218+ OP_CHECK_IF(SetTilingAndCoreSplit(context, batchSize, channels, inputHeight, inputWidth, outputHeight, outputWidth,
219+ actualScaleH, actualScaleW, coreNum) != ge::GRAPH_SUCCESS,
220+ OP_LOGE(context, "SetTilingAndCoreSplit error"), return ge::GRAPH_FAILED);
221+ 
222+ OP_CHECK_IF(SetMemoryAndTilingKey(context, ubSize, alignCorners) != ge::GRAPH_SUCCESS,
223+ OP_LOGE(context, "SetMemoryAndTilingKey error"), return ge::GRAPH_FAILED);
224+ 
225+ return ge::GRAPH_SUCCESS;
226+}
227+ 
228+} // namespace optiling
@@ -0,0 +1,272 @@
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 upsample_bicubic2d_simt.h
15+ * \brief SIMT kernel implementation for upsample_bicubic2d
16+ *
17+ * 性能优化(Stage 1):
18+ * R001/R007: AlignCorners 模板参数化,消除循环内运行期分支
19+ * R002: 3D 坐标拆解使用 Simt::UintDiv 快除替代原生除法
20+ * R003: 索引类型模板化(uint32/uint64),32 位路径吞吐翻倍
21+ * R005: N*C 合轴降为 3D 拆解,减少 1 次 UintDiv
22+ */
23+ 
24+#ifndef __UPSAMPLE_BICUBIC2D_SIMT_H__
25+#define __UPSAMPLE_BICUBIC2D_SIMT_H__
26+ 
27+#include "kernel_operator.h"
28+#include "kernel_tiling/kernel_tiling.h"
29+#include "simt_api/common_functions.h"
30+#include "simt_api/asc_simt.h"
31+#include "simt_api/math_functions.h"
32+#include "simt_api/asc_fp16.h"
33+#include "simt_api/asc_bf16.h"
34+#include "upsample_bicubic2d_tiling_data.h"
35+#include "upsample_bicubic2d_tiling_key.h"
36+ 
37+#include <type_traits>
38+ 
39+namespace NsUpsampleBicubic2d {
40+ 
41+using namespace AscendC;
42+ 
43+constexpr uint32_t BICUBIC_THREADS_32 = 512;
44+constexpr uint32_t BICUBIC_THREADS_64 = 256;
45+constexpr size_t UINT32_BYTE_SIZE = 4;
46+constexpr int32_t BICUBIC_WINDOW_SIZE = 4;
47+constexpr float CUBIC_A = -0.75f;
48+constexpr float CUBIC_C2_COEFF_5A = 5.0f;
49+constexpr float CUBIC_C2_COEFF_8A = 8.0f;
50+constexpr float CUBIC_C2_COEFF_4A = 4.0f;
51+constexpr float CUBIC_C1_COEFF_2 = 2.0f;
52+constexpr float CUBIC_C1_COEFF_3 = 3.0f;
53+constexpr float CUBIC_OFFSET_1 = 1.0f;
54+constexpr float CUBIC_OFFSET_2 = 2.0f;
55+constexpr float HALF_PIXEL_OFFSET = 0.5f;
56+ 
57+// ===== 线程数按索引位宽分级(R003/R006)=====
58+// bicubic 插值寄存器压力中等偏高(16 像素读取 + 5 次三次插值)
59+// uint32 路径: 512 线程;uint64 路径: 256 线程
60+template <typename UintT>
61+static constexpr uint32_t BICUBIC_THREADS = (sizeof(UintT) == UINT32_BYTE_SIZE) ? BICUBIC_THREADS_32 : BICUBIC_THREADS_64;
62+ 
63+// ===== 类型转换辅助函数(__simt_callee__)=====
64+ 
65+template <typename T>
66+__simt_callee__ __aicore__ inline float ToFloat(T v)
67+{
68+ if constexpr (std::is_same_v<T, float>) {
69+ return v;
70+ } else if constexpr (std::is_same_v<T, half>) {
71+ return __half2float(v);
72+ } else {
73+ return __bfloat162float(v);
74+ }
75+}
76+ 
77+template <typename T>
78+__simt_callee__ __aicore__ inline T FromFloat(float v)
79+{
80+ if constexpr (std::is_same_v<T, float>) {
81+ return v;
82+ } else if constexpr (std::is_same_v<T, half>) {
83+ return __float2half(v);
84+ } else {
85+ return __float2bfloat16(v);
86+ }
87+}
88+ 
89+// ===== 三次插值核心函数(__simt_callee__)=====
90+ 
91+__simt_callee__ __aicore__ inline float CubicInterp1d(
92+ float x0, float x1, float x2, float x3, float t)
93+{
94+ const float A = CUBIC_A;
95+ // coeffs[0]: cubic_convolution2(t + 1, A)
96+ float x_0 = t + CUBIC_OFFSET_1;
97+ float c0 = ((A * x_0 - CUBIC_C2_COEFF_5A * A) * x_0 + CUBIC_C2_COEFF_8A * A) * x_0 - CUBIC_C2_COEFF_4A * A;
98+ // coeffs[1]: cubic_convolution1(t, A)
99+ float c1 = ((A + CUBIC_C1_COEFF_2) * t - (A + CUBIC_C1_COEFF_3)) * t * t + CUBIC_OFFSET_1;
100+ // coeffs[2]: cubic_convolution1(1 - t, A)
101+ float x2v = CUBIC_OFFSET_1 - t;
102+ float c2 = ((A + CUBIC_C1_COEFF_2) * x2v - (A + CUBIC_C1_COEFF_3)) * x2v * x2v + CUBIC_OFFSET_1;
103+ // coeffs[3]: cubic_convolution2(2 - t, A)
104+ float x_3 = CUBIC_OFFSET_2 - t;
105+ float c3 = ((A * x_3 - CUBIC_C2_COEFF_5A * A) * x_3 + CUBIC_C2_COEFF_8A * A) * x_3 - CUBIC_C2_COEFF_4A * A;
106+ return x0 * c0 + x1 * c1 + x2 * c2 + x3 * c3;
107+}
108+ 
109+// ===== 边界安全像素读取(__simt_callee__)=====
110+ 
111+template <typename T>
112+__simt_callee__ __aicore__ inline float GetPixelBounded(
113+ __gm__ T* data, int64_t ncBase,
114+ int32_t height, int32_t width,
115+ int32_t y, int32_t x)
116+{
117+ int32_t accessY = (y < 0) ? 0 : ((y >= height) ? (height - 1) : y);
118+ int32_t accessX = (x < 0) ? 0 : ((x >= width) ? (width - 1) : x);
119+ int64_t pixelIdx = ncBase + static_cast<int64_t>(accessY) * width + accessX;
120+ return ToFloat<T>(data[pixelIdx]);
121+}
122+ 
123+// ===== 双三次插值子计算(__simt_callee__,拆分以满足50行限制)=====
124+ 
125+template <typename T>
126+__simt_callee__ __aicore__ inline float ComputeBicubicPixel(
127+ __gm__ T* input, int64_t ncBase,
128+ int32_t inputHeight, int32_t inputWidth,
129+ int32_t inY, int32_t inX,
130+ float tX, float tY)
131+{
132+ float coefficients[BICUBIC_WINDOW_SIZE];
133+ for (int32_t k = 0; k < BICUBIC_WINDOW_SIZE; k++) {
134+ int32_t iy = inY - 1 + k;
135+ float v0 = GetPixelBounded<T>(input, ncBase, inputHeight, inputWidth, iy, inX - 1);
136+ float v1 = GetPixelBounded<T>(input, ncBase, inputHeight, inputWidth, iy, inX + 0);
137+ float v2 = GetPixelBounded<T>(input, ncBase, inputHeight, inputWidth, iy, inX + 1);
138+ float v3 = GetPixelBounded<T>(input, ncBase, inputHeight, inputWidth, iy, inX + 2);
139+ coefficients[k] = CubicInterp1d(v0, v1, v2, v3, tX);
140+ }
141+ return CubicInterp1d(coefficients[0], coefficients[1],
142+ coefficients[2], coefficients[3], tY);
143+}
144+ 
145+// ===== SIMT VF Kernel(R001/R002/R003/R005 全部应用)=====
146+//
147+// 模板参数:
148+// T - 数据类型(float / half / bfloat16_t)
149+// AlignCorners - 编译期常量,消除循环内 if-else 分支(R001/R007)
150+// UintT - 索引/快除统一类型(uint32_t 或 uint64_t)(R003)
151+//
152+// 坐标拆解策略(R005 + R002):
153+// N*C 合轴为 nc → 3D 拆解 (nc, oy, ox)
154+// 仅 2 次 Simt::UintDiv(outputWidth, outputHeight),比原始 4D 少 1 次
155+// ncBase = nc * inputHeight * inputWidth,无需再分解 n 和 c
156+ 
157+template <typename T, bool AlignCorners, typename UintT>
158+__simt_vf__ __aicore__ __launch_bounds__(BICUBIC_THREADS<UintT>)
159+inline void OpUpsampleBicubic2dSimtKernel(
160+ UintT totalElements,
161+ int32_t ncSize,
162+ int32_t inputHeight,
163+ int32_t inputWidth,
164+ int32_t outputHeight,
165+ int32_t outputWidth,
166+ float scaleH,
167+ float scaleW,
168+ UintT owMagic,
169+ UintT owShift,
170+ UintT ohMagic,
171+ UintT ohShift,
172+ __gm__ T* input,
173+ __gm__ T* output)
174+{
175+ for (UintT idx = static_cast<UintT>(blockIdx.x) * static_cast<UintT>(blockDim.x) + static_cast<UintT>(threadIdx.x);
176+ idx < totalElements;
177+ idx += static_cast<UintT>(blockDim.x) * static_cast<UintT>(gridDim.x)) {
178+ 
179+ // 3D 坐标拆解 (nc, oy, ox) — R005 合轴 + R002 快除
180+ // idx = nc * (outputHeight * outputWidth) + oy * outputWidth + ox
181+ UintT ox = idx - Simt::UintDiv<UintT>(idx, owMagic, owShift) * static_cast<UintT>(outputWidth);
182+ UintT tmp = Simt::UintDiv<UintT>(idx, owMagic, owShift);
183+ UintT oy = tmp - Simt::UintDiv<UintT>(tmp, ohMagic, ohShift) * static_cast<UintT>(outputHeight);
184+ UintT nc = Simt::UintDiv<UintT>(tmp, ohMagic, ohShift);
185+ 
186+ // 短路优化:输入输出尺寸相同时直接拷贝
187+ if (inputHeight == outputHeight && inputWidth == outputWidth) {
188+ int64_t inIdx = static_cast<int64_t>(nc) * inputHeight * inputWidth +
189+ static_cast<int64_t>(oy) * inputWidth + static_cast<int64_t>(ox);
190+ output[idx] = input[inIdx];
191+ continue;
192+ }
193+ 
194+ // 计算源坐标 — AlignCorners 为编译期常量,if constexpr 零运行期开销(R001/R007)
195+ float realX, realY;
196+ if constexpr (AlignCorners) {
197+ realX = scaleW * static_cast<float>(ox);
198+ realY = scaleH * static_cast<float>(oy);
199+ } else {
200+ realX = scaleW * (static_cast<float>(ox) + HALF_PIXEL_OFFSET) - HALF_PIXEL_OFFSET;
201+ realY = scaleH * (static_cast<float>(oy) + HALF_PIXEL_OFFSET) - HALF_PIXEL_OFFSET;
202+ }
203+ 
204+ int32_t inX = static_cast<int32_t>(floorf(realX));
205+ float tX = realX - static_cast<float>(inX);
206+ int32_t inY = static_cast<int32_t>(floorf(realY));
207+ float tY = realY - static_cast<float>(inY);
208+ 
209+ // 双三次插值 — ncBase 直接用合轴后的 nc 计算,无需分解 n 和 c(R005)
210+ int64_t ncBase = static_cast<int64_t>(nc) * inputHeight * inputWidth;
211+ 
212+ float result = ComputeBicubicPixel<T>(
213+ input, ncBase, inputHeight, inputWidth,
214+ inY, inX, tX, tY);
215+ 
216+ output[idx] = FromFloat<T>(result);
217+ }
218+}
219+ 
220+// ===== Process 入口(R003: 按数据量分派 32/64 位路径)=====
221+ 
222+template <typename T, bool AlignCorners>
223+__aicore__ inline void Process(GM_ADDR x, GM_ADDR y,
224+ const UpsampleBicubic2dRegBaseTilingData* tilingData)
225+{
226+ __gm__ T* inputGm = (__gm__ T*)x;
liu-wei
liu-weiliu-wei6月27日

🟢 Low | upsample_bicubic2d Process 函数缺少 totalElements == 0 拦截

template <typename T, bool AlignCorners>
__aicore__ inline void Process(GM_ADDR x, GM_ADDR y,
                                const UpsampleBicubic2dRegBaseTilingData* tilingData)
{
    __gm__ T* inputGm = (__gm__ T*)x;
    __gm__ T* outputGm = (__gm__ T*)y;

    if (tilingData->totalElements <= static_cast<int64_t>(INT32_MAX)) {
        // 直接调用 asc_vf_call,没有 totalElements == 0 拦截
        ...
    }
}

对比 extract_glimpse_v2_simt.h

int64_t outputTotalElements = tilingData->outputTotalElements;
if (outputTotalElements == 0) {
    return;
}

upsample_bicubic2d 的 Process 缺少 totalElements == 0 早期返回。

实际影响

  • totalElements = 0 时,asc_vf_call 启动 BICUBIC_THREADS=512/256 个 thread,但每个 thread 的 for 循环条件 idx < 0 立即失败
  • 实际不会执行计算,但仍然启动 kernel launch
  • 算子调用方通常不会传 0 size,但增加显式保护更安全

修复建议

if (tilingData->totalElements == 0) {
    return;
}
likedislike
ideal_
ideal_
6月27日 评论:
227+ __gm__ T* outputGm = (__gm__ T*)y;
228+ 
229+ if (tilingData->totalElements <= static_cast<int64_t>(INT32_MAX)) {
230+ // 32 位路径:索引运算吞吐翻倍(R003)
231+ uint32_t owMagic = 0, owShift = 0;
232+ GetUintDivMagicAndShift<uint32_t>(owMagic, owShift, static_cast<uint32_t>(tilingData->outputWidth));
233+ uint32_t ohMagic = 0, ohShift = 0;
234+ GetUintDivMagicAndShift<uint32_t>(ohMagic, ohShift, static_cast<uint32_t>(tilingData->outputHeight));
235+ 
236+ asc_vf_call<OpUpsampleBicubic2dSimtKernel<T, AlignCorners, uint32_t>>(
237+ dim3(BICUBIC_THREADS<uint32_t>),
238+ static_cast<uint32_t>(tilingData->totalElements),
239+ tilingData->ncSize,
240+ tilingData->inputHeight,
241+ tilingData->inputWidth,
242+ tilingData->outputHeight,
243+ tilingData->outputWidth,
244+ tilingData->scaleH,
245+ tilingData->scaleW,
246+ owMagic, owShift, ohMagic, ohShift,
247+ inputGm, outputGm);
248+ } else {
249+ // 64 位路径:大数据量回退
250+ uint64_t owMagic = 0, owShift = 0;
251+ GetUintDivMagicAndShift<uint64_t>(owMagic, owShift, static_cast<uint64_t>(tilingData->outputWidth));
252+ uint64_t ohMagic = 0, ohShift = 0;
253+ GetUintDivMagicAndShift<uint64_t>(ohMagic, ohShift, static_cast<uint64_t>(tilingData->outputHeight));
254+ 
255+ asc_vf_call<OpUpsampleBicubic2dSimtKernel<T, AlignCorners, uint64_t>>(
256+ dim3(BICUBIC_THREADS<uint64_t>),
257+ static_cast<uint64_t>(tilingData->totalElements),
258+ static_cast<uint64_t>(tilingData->ncSize),
259+ static_cast<uint64_t>(tilingData->inputHeight),
260+ static_cast<uint64_t>(tilingData->inputWidth),
261+ static_cast<uint64_t>(tilingData->outputHeight),
262+ static_cast<uint64_t>(tilingData->outputWidth),
263+ tilingData->scaleH,
264+ tilingData->scaleW,
265+ owMagic, owShift, ohMagic, ohShift,
266+ inputGm, outputGm);
267+ }
268+}
269+ 
270+} // namespace NsUpsampleBicubic2d
271+ 
272+#endif // __UPSAMPLE_BICUBIC2D_SIMT_H__
@@ -0,0 +1,33 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+/*!
14+ * \file upsample_bicubic2d_tiling_data.h
15+ * \brief tiling data struct for upsample_bicubic2d
16+ */
17+ 
18+#ifndef __UPSAMPLE_BICUBIC2D_REGBASE_TILING_DATA_H__
19+#define __UPSAMPLE_BICUBIC2D_REGBASE_TILING_DATA_H__
20+ 
21+struct UpsampleBicubic2dRegBaseTilingData {
22+ int32_t needCoreNum; // 实际启动核数
23+ int64_t totalElements; // 总输出元素数: N * C * H_out * W_out
24+ int32_t ncSize; // N * C(合轴后的批次-通道合并维度)
25+ int32_t inputHeight; // H_in
26+ int32_t inputWidth; // W_in
27+ int32_t outputHeight; // H_out
28+ int32_t outputWidth; // W_out
29+ float scaleH; // 高度方向缩放因子(已计算好的实际值)
30+ float scaleW; // 宽度方向缩放因子(已计算好的实际值)
31+};
32+ 
33+#endif // __UPSAMPLE_BICUBIC2D_REGBASE_TILING_DATA_H__
@@ -0,0 +1,39 @@
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 upsample_bicubic2d_tiling_key.h
15+ * \brief tiling key declare for upsample_bicubic2d
16+ */
17+ 
18+#ifndef __UPSAMPLE_BICUBIC2D_TILING_KEY_H__
19+#define __UPSAMPLE_BICUBIC2D_TILING_KEY_H__
20+ 
21+#include "ascendc/host_api/tiling/template_argument.h"
22+ 
23+// 编码 alignCorners 场景模式,dtype 由 DTYPE_ 宏自动实例化
24+#define UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE 0
25+#define UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE 1
26+ 
27+ASCENDC_TPL_ARGS_DECL(
28+ UpsampleBicubic2d,
29+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST,
30+ UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE,
31+ UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE));
32+ 
33+ASCENDC_TPL_SEL(
34+ ASCENDC_TPL_ARGS_SEL(
35+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,
36+ UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_FALSE,
37+ UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE)));
38+ 
39+#endif // __UPSAMPLE_BICUBIC2D_TILING_KEY_H__
@@ -0,0 +1,38 @@
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 upsample_bicubic2d_apt.cpp
15+ * \brief arch35 SIMT kernel entry for upsample_bicubic2d
16+ */
17+ 
18+#include "kernel_tiling/kernel_tiling.h"
19+#include "kernel_operator.h"
20+#include "./arch35/upsample_bicubic2d_simt.h"
21+#include "./arch35/upsample_bicubic2d_tiling_key.h"
22+#include "./arch35/upsample_bicubic2d_tiling_data.h"
23+ 
24+using namespace AscendC;
25+ 
26+template <uint32_t schMode>
27+__global__ __aicore__ void upsample_bicubic2d(GM_ADDR x, GM_ADDR y,
28+ GM_ADDR workspace, GM_ADDR tiling)
29+{
30+ REGISTER_TILING_DEFAULT(UpsampleBicubic2dRegBaseTilingData);
31+ GET_TILING_DATA_WITH_STRUCT(UpsampleBicubic2dRegBaseTilingData, tilingData, tiling);
32+ 
33+ if constexpr (schMode == UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE) {
34+ NsUpsampleBicubic2d::Process<DTYPE_INPUT, true>(x, y, &tilingData);
35+ } else {
36+ NsUpsampleBicubic2d::Process<DTYPE_INPUT, false>(x, y, &tilingData);
liu-wei
liu-weiliu-wei6月27日

🟢 Low | APT 中 if constexpr ... else 完整覆盖 2 个 mode(最佳实践)

if constexpr (schMode == UPSAMPLE_BICUBIC2D_MODE_ALIGN_CORNERS_TRUE) {
    NsUpsampleBicubic2d::Process<DTYPE_INPUT, true>(x, y, &tilingData);
} else {
    NsUpsampleBicubic2d::Process<DTYPE_INPUT, false>(x, y, &tilingData);
}

这是 if constexpr 链的最佳实践写法:用 else 兜底,覆盖所有未匹配 mode。相比 PR 1035 和 PR 1036 的 6-mode 链写法更简洁。

对比

  • PR 1035: 6 个 if constexpr 链无 else
  • PR 1036: 1 个 if constexpr 无 else
  • PR 1039 (upsample_bicubic2d): 1 个 if constexpr + 1 个 else(完整覆盖)
  • PR 1039 (extract_glimpse_v2): 3 个 if constexpr + 1 个 else(完整覆盖 4 个 mode)

正向建议:可以把这个最佳实践(用 else 兜底)作为模板,让 PR 1035/1036 的 APT 后续修复时参考。

likedislike
ideal_
ideal_
6月27日 评论:
37+ }
38+}
@@ -10,7 +10,11 @@
10 10 
11file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12if(UT_TEST_ALL OR OP_HOST_UT)12if(UT_TEST_ALL OR OP_HOST_UT)
13- add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})13+ if("${ASCEND_COMPUTE_UNIT}" STREQUAL "ascend950")
14+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}/arch35)
15+ else()
16+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
17+ endif()
14endif()18endif()
15 19 
16if(NOT (UT_TEST_ALL OR OP_API_UT))20if(NOT (UT_TEST_ALL OR OP_API_UT))
@@ -0,0 +1,87 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* Generated By CANNBot */
12+ 
13+#include <iostream>
14+#include <gtest/gtest.h>
15+#include "tiling_context_faker.h"
16+#include "tiling_case_executor.h"
17+#include "../../../../op_kernel/arch35/upsample_bicubic2d_tiling_data.h"
18+ 
19+using namespace std;
20+using namespace ge;
21+ 
22+class UpsampleBicubic2dTiling : public testing::Test {
23+protected:
24+ static void SetUpTestCase() { std::cout << "UpsampleBicubic2dTiling SetUp" << std::endl; }
25+ 
26+ static void TearDownTestCase() { std::cout << "UpsampleBicubic2dTiling TearDown" << std::endl; }
27+};
28+ 
29+std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend950"}};
30+ 
31+TEST_F(UpsampleBicubic2dTiling, upsample_bicubic2d_float32_align_false)
32+{
33+ struct UpsampleBicubic2dCompileInfo {
34+ } compileInfo;
35+ gert::TilingContextPara tilingContextPara(
36+ "UpsampleBicubic2d",
37+ {
38+ {{{1, 1, 4, 4}, {1, 1, 4, 4}}, ge::DT_FLOAT, ge::FORMAT_ND}, // input x
39+ },
40+ {
41+ {{{1, 1, 8, 8}, {1, 1, 8, 8}}, ge::DT_FLOAT, ge::FORMAT_ND}, // output y
42+ },
43+ {
44+ gert::TilingContextPara::OpAttr(
45+ "output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(std::vector<int64_t>{8, 8})),
46+ gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
47+ gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0f)),
48+ gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0f)),
49+ },
50+ &compileInfo, "Ascend950",
51+ 64, // number of cores
52+ 262144, // ub size
53+ 4096); // tiling data max size
54+ uint64_t expectTilingKey = 0; // ALIGN_CORNERS_FALSE
55+ string expectTilingData = "1 64 17179869185 34359738372 4539628424389459976 1056964608 ";
56+ std::vector<size_t> expectWorkspaces = {16777216};
57+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
58+}
59+ 
60+TEST_F(UpsampleBicubic2dTiling, upsample_bicubic2d_float16_align_true)
61+{
62+ struct UpsampleBicubic2dCompileInfo {
63+ } compileInfo;
64+ gert::TilingContextPara tilingContextPara(
65+ "UpsampleBicubic2d",
66+ {
67+ {{{2, 3, 4, 4}, {2, 3, 4, 4}}, ge::DT_FLOAT16, ge::FORMAT_ND}, // input x
68+ },
69+ {
70+ {{{2, 3, 8, 8}, {2, 3, 8, 8}}, ge::DT_FLOAT16, ge::FORMAT_ND}, // output y
71+ },
72+ {
73+ gert::TilingContextPara::OpAttr(
74+ "output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(std::vector<int64_t>{8, 8})),
75+ gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
76+ gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0f)),
77+ gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0f)),
78+ },
79+ &compileInfo, "Ascend950",
80+ 64, // number of cores
81+ 262144, // ub size
82+ 4096); // tiling data max size
83+ uint64_t expectTilingKey = 1; // ALIGN_CORNERS_TRUE
84+ string expectTilingData = "1 384 17179869190 34359738372 4529334482997608456 1054567863 ";
85+ std::vector<size_t> expectWorkspaces = {16777216};
86+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
87+}
@@ -16,6 +16,7 @@ if (UT_TEST_ALL OR OP_KERNEL_UT)
16 # 算子自己的tiling文件路径16 # 算子自己的tiling文件路径
17 set(upsample_bicubic2d_tiling_files17 set(upsample_bicubic2d_tiling_files
18 ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_bicubic2d_tiling.cpp18 ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_bicubic2d_tiling.cpp
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_bicubic2d_tiling_arch35.cpp
19 # ${elewise_common_tiling_files}20 # ${elewise_common_tiling_files}
20 )21 )
21 # 使用AddOpTestCase22 # 使用AddOpTestCase
@@ -230,7 +230,7 @@
230 {"name":"ReluGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},230 {"name":"ReluGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
231 {"name":"LeakyReluGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},231 {"name":"LeakyReluGrad", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
232 {"name":"UpsampleNearest", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "kirinx90", "kirin9030"], "auto_sync" : true},232 {"name":"UpsampleNearest", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "kirinx90", "kirin9030"], "auto_sync" : true},
233- {"name":"UpsampleBicubic2d", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend310b", "kirinx90", "kirin9030"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},233+ {"name":"UpsampleBicubic2d", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend310b", "kirinx90", "kirin9030", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend310b": true, "kirinx90": true, "kirin9030": true, "ascend950": false}},
234 {"name":"UpsampleBicubic2dAA", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "kirinx90", "kirin9030"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},234 {"name":"UpsampleBicubic2dAA", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "kirinx90", "kirin9030"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
235 {"name":"UpsampleBicubic2dGrad", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},235 {"name":"UpsampleBicubic2dGrad", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},
236 {"name":"ResizeUpsampleTrilinear", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},236 {"name":"ResizeUpsampleTrilinear", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
@@ -587,6 +587,7 @@
587 {"name":"Addr", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},587 {"name":"Addr", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
588 {"name":"NMSWithMask", "compute_units": ["ascend910b", "ascend950"], "auto_sync" : false, "impl_mode" : ""},588 {"name":"NMSWithMask", "compute_units": ["ascend910b", "ascend950"], "auto_sync" : false, "impl_mode" : ""},
589 {"name":"Rasterizer", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false, "impl_mode" : ""},589 {"name":"Rasterizer", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false, "impl_mode" : ""},
590+ {"name":"ExtractGlimpseV2", "compute_units": ["ascend950"], "auto_sync": false},
590 {"name":"RoiAlignV2", "compute_units": ["ascend910b"], "auto_sync":false, "impl_mode" : ""},591 {"name":"RoiAlignV2", "compute_units": ["ascend910b"], "auto_sync":false, "impl_mode" : ""},
591 {"name":"DIoUGrad", "compute_units": ["ascend950"], "auto_sync": false},592 {"name":"DIoUGrad", "compute_units": ["ascend950"], "auto_sync": false},
592 {"name":"GIoUGrad", "compute_units": ["ascend950"], "auto_sync": false}593 {"name":"GIoUGrad", "compute_units": ["ascend950"], "auto_sync": false}