Star570
391
代码
代码
Issues20
Pull Requests45
Actions
项目讨论
Wiki
分析
Star570
391
已合并
resize_nearest_neighbor_v2 nchw 格式性能优化 #822
chenfeng创建于 5月7日
resize_nearest_neighbor_v2 nchw 格式性能优化 #822
已合并
chenfeng创建于 5月7日
从chenfeng61/ops-cvnew_nearest合入到cann/ops-cvmaster
共 10 个文件变更+1076-22
Mimage/resize_nearest_neighbor_v2/CMakeLists.txt+16-1
@@ -1,5 +1,5 @@
1# -----------------------------------------------------------------------------------------------------------1# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.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 of3# 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").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.5# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -7,6 +7,21 @@
7# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.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.8# See LICENSE in the root of the software repository for the full text of the License.
9# -----------------------------------------------------------------------------------------------------------9# -----------------------------------------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+ 
13+if(NOT ENABLE_TEST AND NOT BENCHMARK)
14+ list(REMOVE_ITEM CURRENT_DIRS tests)
15+endif()
16+ 
17+foreach(SUB_DIR ${CURRENT_DIRS})
18+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
19+ add_subdirectory(${SUB_DIR})
20+ endif()
21+endforeach()
22+ 
10set(SUPPORT_COMPUTE_UNIT "ascend950" "mc62cm12a")23set(SUPPORT_COMPUTE_UNIT "ascend950" "mc62cm12a")
24+ 
11set(SUPPORT_TILING_DIR "arch35" "arch35")25set(SUPPORT_TILING_DIR "arch35" "arch35")
26+ 
12add_all_modules_sources(OPTYPE resize_nearest_neighbor_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)27add_all_modules_sources(OPTYPE resize_nearest_neighbor_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Mimage/resize_nearest_neighbor_v2/op_host/arch35/resize_nearest_neighbor_v2_tiling_base.cpp+153-19
@@ -18,6 +18,7 @@
18#include "log/log.h"18#include "log/log.h"
19#include "util/math_util.h"19#include "util/math_util.h"
20#include "op_host/tiling_util.h"20#include "op_host/tiling_util.h"
21+#include "util/platform_util.h"
21#include <cmath>22#include <cmath>
22 23 
23namespace optiling {24namespace optiling {
@@ -48,9 +49,10 @@ constexpr uint64_t SCHEDULE_ID_SIMT_INPUT_EQ_ONE = 5;
48constexpr uint64_t SCHEDULE_ID_NHWC_W_OUT_ALL = 6;49constexpr uint64_t SCHEDULE_ID_NHWC_W_OUT_ALL = 6;
49constexpr uint64_t SCHEDULE_ID_NHWC_NOT_W_OUT_ALL = 7;50constexpr uint64_t SCHEDULE_ID_NHWC_NOT_W_OUT_ALL = 7;
50constexpr uint64_t SCHEDULE_ID_NHWC_CORE_NH = 8;51constexpr uint64_t SCHEDULE_ID_NHWC_CORE_NH = 8;
52+constexpr uint64_t SCHEDULE_ID_GATHER_ALL_HW = 9;
53+constexpr uint64_t SCHEDULE_ID_GATHER_CUT_H = 10;
51constexpr float ENLARGE_SCALE_THRESHOLD = 4;54constexpr float ENLARGE_SCALE_THRESHOLD = 4;
52constexpr float REDUCE_SCALE_THRESHOLD = 0.25;55constexpr float REDUCE_SCALE_THRESHOLD = 0.25;
53-constexpr int64_t ONE_BLOCK_SIZE = 32;
54constexpr int64_t UNIT_PROC_BYTES = 256;56constexpr int64_t UNIT_PROC_BYTES = 256;
55constexpr int64_t TEMPLATE02_C_DIM_LOWER = 64;57constexpr int64_t TEMPLATE02_C_DIM_LOWER = 64;
56constexpr float HALF_PIXEL = 0.5;58constexpr float HALF_PIXEL = 0.5;
@@ -95,18 +97,22 @@ private:
95 int64_t CalTimes(int64_t a, int64_t b) const;97 int64_t CalTimes(int64_t a, int64_t b) const;
96 bool IsMatchTiling_NHWC();98 bool IsMatchTiling_NHWC();
97 bool IsMatchTiling_NHWC_UB2UB();99 bool IsMatchTiling_NHWC_UB2UB();
100+ bool IsMatchTiling_NCHW_Gather();
98 void DoTilingSmallC();101 void DoTilingSmallC();
99 void DoTilingBigC();102 void DoTilingBigC();
100 void DoTilingJHC();103 void DoTilingJHC();
101 void DoTilingUb2Ub();104 void DoTilingUb2Ub();
102 void ComputeHCut();105 void ComputeHCut();
103 void DoTilingUb2UbCutNH();106 void DoTilingUb2UbCutNH();
107+ void DoTilingGatherNchw();
108+ void ComputeGatherCutH(int64_t oneBlockNum);
104 109 
105private:110private:
106 uint64_t schId_ = 0;111 uint64_t schId_ = 0;
107 uint64_t idxUseInt32_ = 0;112 uint64_t idxUseInt32_ = 0;
108 int32_t dtypeSize_ = 0;113 int32_t dtypeSize_ = 0;
109 int64_t coreNum_ = 0;114 int64_t coreNum_ = 0;
115+ int64_t ubBlockSize_ = 0;
110 int64_t ubSize_ = 0;116 int64_t ubSize_ = 0;
111 int64_t realCoreNum_ = 0;117 int64_t realCoreNum_ = 0;
112 int64_t lenC_ = 0;118 int64_t lenC_ = 0;
@@ -117,6 +123,11 @@ private:
117 int64_t lenSrcW_ = 0;123 int64_t lenSrcW_ = 0;
118 int64_t lenDesH_ = 0;124 int64_t lenDesH_ = 0;
119 int64_t lenDesW_ = 0;125 int64_t lenDesW_ = 0;
126+ int64_t lenSrcHAlign_ = 0;
127+ int64_t lenSrcWAlign_ = 0;
128+ int64_t lenDesHAlign_ = 0;
129+ int64_t lenDesWAlign_ = 0;
130+ float bias_ = 0.0f;
120 int64_t splitBlockFactor_ = 0;131 int64_t splitBlockFactor_ = 0;
121 int64_t splitBlockTailFactor_ = 0;132 int64_t splitBlockTailFactor_ = 0;
122 int64_t alignCorners_ = 0;133 int64_t alignCorners_ = 0;
@@ -205,7 +216,11 @@ void ResizeNearestNeighborV2AscendCTilingImpl::SetDimsByFormat() {
205 lenSrcW_ = xShape_.GetDim(W_DIM_IDX_NHWC);216 lenSrcW_ = xShape_.GetDim(W_DIM_IDX_NHWC);
206 lenDesW_ = yShape_.GetDim(W_DIM_IDX_NHWC);217 lenDesW_ = yShape_.GetDim(W_DIM_IDX_NHWC);
207 }218 }
208- 219+ int64_t oneBlockNum = ubBlockSize_ / dtypeSize_;
L
Lloov5月9日

这里不需要做非零判断吗？

likedislike
chenfeng
chenfeng
5月9日 评论：
220+ lenSrcHAlign_ = Ops::Base::CeilDiv(lenSrcH_, oneBlockNum) * oneBlockNum;
221+ lenDesHAlign_ = Ops::Base::CeilDiv(lenDesH_, oneBlockNum) * oneBlockNum;
222+ lenSrcWAlign_ = Ops::Base::CeilDiv(lenSrcW_, oneBlockNum) * oneBlockNum;
223+ lenDesWAlign_ = Ops::Base::CeilDiv(lenDesW_, oneBlockNum) * oneBlockNum;
209 OP_LOGI(context_->GetNodeName(), "lenN_:%ld , lenC_: %ld, srcH:%ld, srcW:%ld, dstH:%ld, dstW:%ld",224 OP_LOGI(context_->GetNodeName(), "lenN_:%ld , lenC_: %ld, srcH:%ld, srcW:%ld, dstH:%ld, dstW:%ld",
210 lenN_, lenC_, lenSrcH_, lenSrcW_, lenDesH_, lenDesW_);225 lenN_, lenC_, lenSrcH_, lenSrcW_, lenDesH_, lenDesW_);
211 wcNum_ = lenSrcW_ * lenC_;226 wcNum_ = lenSrcW_ * lenC_;
@@ -213,10 +228,10 @@ void ResizeNearestNeighborV2AscendCTilingImpl::SetDimsByFormat() {
213 hwcNum_ = lenSrcH_ * wcNum_;228 hwcNum_ = lenSrcH_ * wcNum_;
214 dstHwcNum_ = lenDesH_ * dstWcNum_;229 dstHwcNum_ = lenDesH_ * dstWcNum_;
215 isAlign_ = false;230 isAlign_ = false;
216- if ((lenC_ * dtypeSize_) % ONE_BLOCK_SIZE == 0) {231+ if ((lenC_ * dtypeSize_) % ubBlockSize_ == 0) {
217 isAlign_ = true;232 isAlign_ = true;
218 }233 }
219- lenCAlign_ = ((lenC_ * dtypeSize_ + ONE_BLOCK_SIZE - 1) / ONE_BLOCK_SIZE * ONE_BLOCK_SIZE) / dtypeSize_;234+ lenCAlign_ = ((lenC_ * dtypeSize_ + ubBlockSize_ - 1) / ubBlockSize_ * ubBlockSize_) / dtypeSize_;
220}235}
221 236 
222void ResizeNearestNeighborV2AscendCTilingImpl::SetScales() {237void ResizeNearestNeighborV2AscendCTilingImpl::SetScales() {
@@ -282,19 +297,19 @@ int64_t ResizeNearestNeighborV2AscendCTilingImpl::CalcSrcLenByDesLen(int64_t des
282int64_t ResizeNearestNeighborV2AscendCTilingImpl::CalcUnitWCountPerUB(int64_t unitDesWBytes) {297int64_t ResizeNearestNeighborV2AscendCTilingImpl::CalcUnitWCountPerUB(int64_t unitDesWBytes) {
283 int64_t procUnitSrcWCount = CalcSrcLenByDesLen(unitDesWBytes / dtypeSize_, scaleW_);298 int64_t procUnitSrcWCount = CalcSrcLenByDesLen(unitDesWBytes / dtypeSize_, scaleW_);
284 int64_t srcWBytesAlignBlockSize = (procUnitSrcWCount > lenSrcW_) ?299 int64_t srcWBytesAlignBlockSize = (procUnitSrcWCount > lenSrcW_) ?
285- Ops::Base::CeilAlign(lenSrcW_ * dtypeSize_, ONE_BLOCK_SIZE) :300+ Ops::Base::CeilAlign(lenSrcW_ * dtypeSize_, ubBlockSize_) :
286- Ops::Base::CeilAlign(procUnitSrcWCount * dtypeSize_, ONE_BLOCK_SIZE);301+ Ops::Base::CeilAlign(procUnitSrcWCount * dtypeSize_, ubBlockSize_);
287 // 输出H轴先只切1302 // 输出H轴先只切1
288 int64_t unitSrcHCount = CalcSrcLenByDesLen(1, scaleH_);303 int64_t unitSrcHCount = CalcSrcLenByDesLen(1, scaleH_);
289 // UB内预留搬入块所需空间304 // UB内预留搬入块所需空间
290 int64_t unitSrcBytes = srcWBytesAlignBlockSize * unitSrcHCount * dtypeSize_;305 int64_t unitSrcBytes = srcWBytesAlignBlockSize * unitSrcHCount * dtypeSize_;
291 // UB内预留坐标计算所需空间306 // UB内预留坐标计算所需空间
292- int64_t unitLocationBytes = Ops::Base::CeilAlign(static_cast<int64_t>(1 * GetSizeByDataType(ge::DT_FLOAT)), ONE_BLOCK_SIZE);307+ int64_t unitLocationBytes = Ops::Base::CeilAlign(static_cast<int64_t>(1 * GetSizeByDataType(ge::DT_FLOAT)), ubBlockSize_);
293 int64_t unitWTotalBytes = unitDesWBytes + unitSrcBytes + unitLocationBytes;308 int64_t unitWTotalBytes = unitDesWBytes + unitSrcBytes + unitLocationBytes;
294 OP_LOGI(context_->GetNodeName(), "CalcUnitWCountPerUB: srcWFactor: %ld, srcHFactor: %ld, \309 OP_LOGI(context_->GetNodeName(), "CalcUnitWCountPerUB: srcWFactor: %ld, srcHFactor: %ld, \
295 unitDesWBytes: %ld, unitSrcBytes: %ld, unitLocationBytes: %ld",310 unitDesWBytes: %ld, unitSrcBytes: %ld, unitLocationBytes: %ld",
296 procUnitSrcWCount, unitSrcHCount, unitDesWBytes, unitSrcBytes, unitLocationBytes);311 procUnitSrcWCount, unitSrcHCount, unitDesWBytes, unitSrcBytes, unitLocationBytes);
297- return Ops::Base::FloorDiv((ubSize_ - ONE_BLOCK_SIZE) / NUM_2, unitWTotalBytes);312+ return Ops::Base::FloorDiv((ubSize_ - ubBlockSize_) / NUM_2, unitWTotalBytes);
298}313}
299 314 
300void ResizeNearestNeighborV2AscendCTilingImpl::DoTilingGatherWLessThanUnitProc(int64_t unitDesWBytes) {315void ResizeNearestNeighborV2AscendCTilingImpl::DoTilingGatherWLessThanUnitProc(int64_t unitDesWBytes) {
@@ -327,7 +342,7 @@ void ResizeNearestNeighborV2AscendCTilingImpl::DoTilingGatherWMoreThanUnitProc(i
327 unitDesWBytes += UNIT_PROC_BYTES;342 unitDesWBytes += UNIT_PROC_BYTES;
328 } else {343 } else {
329 isExceedWSize = true;344 isExceedWSize = true;
330- unitDesWBytes = Ops::Base::CeilAlign(lenDesW_ * dtypeSize_, ONE_BLOCK_SIZE);345+ unitDesWBytes = Ops::Base::CeilAlign(lenDesW_ * dtypeSize_, ubBlockSize_);
331 }346 }
332 unitWCountPerUB = CalcUnitWCountPerUB(unitDesWBytes);347 unitWCountPerUB = CalcUnitWCountPerUB(unitDesWBytes);
333 }348 }
@@ -358,7 +373,7 @@ void ResizeNearestNeighborV2AscendCTilingImpl::DoTilingStrategyGather() {
358 splitBlockTailFactor_ = numPerCore;373 splitBlockTailFactor_ = numPerCore;
359 splitBlockFactor_ = splitBlockFullCount_ == 0 ? 0 : numPerCore + 1;374 splitBlockFactor_ = splitBlockFullCount_ == 0 ? 0 : numPerCore + 1;
360 375 
361- int64_t unitDesWBytes = Ops::Base::CeilAlign(lenDesW_ * dtypeSize_, ONE_BLOCK_SIZE);376+ int64_t unitDesWBytes = Ops::Base::CeilAlign(lenDesW_ * dtypeSize_, ubBlockSize_);
362 if (unitDesWBytes <= UNIT_PROC_BYTES) {377 if (unitDesWBytes <= UNIT_PROC_BYTES) {
363 DoTilingGatherWLessThanUnitProc(unitDesWBytes);378 DoTilingGatherWLessThanUnitProc(unitDesWBytes);
364 } else {379 } else {
@@ -367,6 +382,84 @@ void ResizeNearestNeighborV2AscendCTilingImpl::DoTilingStrategyGather() {
367 }382 }
368}383}
369 384 
385+void ResizeNearestNeighborV2AscendCTilingImpl::ComputeGatherCutH(int64_t oneBlockNum)
386+{
387+ schId_ = SCHEDULE_ID_GATHER_CUT_H;
388+ int64_t dstWidxUb = lenDesWAlign_ * dtypeSize_;
389+ int64_t hFactor = static_cast<int64_t>(static_cast<float>(ubSize_ - ubBlockSize_ * NUM_3 - dstWidxUb) /
390+ (dstWidxUb + dtypeSize_ * NUM_2 + dstWidxUb * NUM_2 + scaleH_ * dtypeSize_ * NUM_2 * lenSrcWAlign_));
391+ int64_t hFactorAlgin = Ops::Base::CeilDiv(hFactor, oneBlockNum) * oneBlockNum;
392+ int64_t dstUb = hFactor * dstWidxUb * NUM_2;
393+ int64_t hInSize = static_cast<int64_t>((static_cast<float>(hFactor) + bias_) * scaleH_) + 1;
394+ if (alignCorners_ == 1) {
395+ hInSize = static_cast<int64_t>((static_cast<float>(hFactor) + bias_) * scaleH_ + HALF_PIXEL) + 1;
396+ }
397+ hInSize = hInSize + 1;
398+ if (hInSize > lenSrcH_) {
399+ hInSize = lenSrcH_;
400+ }
401+ int64_t srcUb = hInSize * lenSrcWAlign_ * dtypeSize_ * NUM_2;
402+ int64_t allUb = hFactor * dstWidxUb + hFactorAlgin * dtypeSize_ * NUM_2 + dstWidxUb + srcUb + dstUb;
403+ while (allUb > ubSize_) {
404+ hFactor = hFactor - 1;
405+ hFactorAlgin = Ops::Base::CeilDiv(hFactor, oneBlockNum) * oneBlockNum;
406+ dstUb = hFactor * dstWidxUb * NUM_2;
407+ hInSize = static_cast<int64_t>((static_cast<float>(hFactor) + bias_) * scaleH_) + 1;
408+ if (alignCorners_ == 1) {
409+ hInSize = static_cast<int64_t>((static_cast<float>(hFactor) + bias_) * scaleH_ + HALF_PIXEL) + 1;
410+ }
411+ hInSize = hInSize + 1;
412+ if (hInSize > lenSrcH_) {
413+ hInSize = lenSrcH_;
414+ }
415+ srcUb = hInSize * lenSrcWAlign_ * dtypeSize_ * NUM_2;
416+ allUb = hFactor * dstWidxUb + hFactorAlgin * dtypeSize_ * NUM_2 + dstWidxUb + srcUb + dstUb;
417+ }
418+ nLoop_ = splitBlockFactorTail_;
419+ splitFactorDesH_ = hFactor;
420+ splitCountDesH_ = Ops::Base::CeilDiv(lenDesH_, hFactor);
421+ nLoopTailLast_ = lenDesH_ - (splitCountDesH_ - 1) * hFactor;
422+ splitFactorTailDesW_ = hInSize * lenSrcWAlign_ * dtypeSize_;
423+ splitFactorDesW_ = hFactor * lenDesWAlign_ * dtypeSize_;
424+ ubSize_ = hFactor * dstWidxUb;
425+}
426+ 
427+void ResizeNearestNeighborV2AscendCTilingImpl::DoTilingGatherNchw()
428+{
429+ int64_t oneBlockNum = ubBlockSize_ / dtypeSize_;
430+ int64_t ncNum = lenN_ * lenC_;
431+ splitBlockFactor_ = Ops::Base::CeilDiv(ncNum, coreNum_);
432+ realCoreNum_ = Ops::Base::CeilDiv(ncNum, splitBlockFactor_);
433+ splitBlockFactorTail_ = ncNum - (realCoreNum_ - 1) * splitBlockFactor_;
434+ // 先按照hw可以全部放下计算
435+ int64_t dstHwIdxUb = lenDesH_ * lenDesWAlign_ * dtypeSize_;
436+ int64_t dstHIdxUb = lenDesHAlign_ * dtypeSize_;
437+ int64_t dstWIdxUb = lenDesWAlign_ * dtypeSize_;
438+ int64_t xUb = lenSrcH_ * lenSrcWAlign_ * dtypeSize_;
439+ int64_t srcUb = xUb * NUM_2;
440+ 
441+ int64_t yUb = lenDesH_ * lenDesWAlign_ * dtypeSize_;
442+ int64_t dstUb = yUb * NUM_2;
443+ 
444+ int64_t allUb = srcUb + dstUb + dstHwIdxUb + dstHIdxUb + dstWIdxUb;
445+ if (allUb <= ubSize_) {
446+ schId_ = SCHEDULE_ID_GATHER_ALL_HW;
447+ // 可以放多少个hw
448+ int64_t ncTimes = (ubSize_ - dstHwIdxUb - dstHIdxUb - dstWIdxUb) / (srcUb + dstUb);
449+ splitFactorDesH_ = ncTimes;
450+ nLoopTimesLast_ = Ops::Base::CeilDiv(splitBlockFactorTail_, ncTimes);
451+ nLoopTimesBefore_ = Ops::Base::CeilDiv(splitBlockFactor_, ncTimes);
452+ nLoopTailLast_ = splitBlockFactorTail_ - (nLoopTimesLast_ - 1) * ncTimes;
453+ nLoop_ = splitBlockFactor_ - (nLoopTimesBefore_ - 1) * ncTimes;
454+ ubSize_ = dstHwIdxUb;
455+ splitFactorTailDesW_ = ncTimes * xUb;
456+ splitFactorDesW_ = ncTimes * yUb;
457+ } else {
458+ // 需要切输出的h
459+ ComputeGatherCutH(oneBlockNum);
460+ }
461+}
462+ 
370bool ResizeNearestNeighborV2AscendCTilingImpl::IsMatchTiling_NHWC() {463bool ResizeNearestNeighborV2AscendCTilingImpl::IsMatchTiling_NHWC() {
371 OP_CHECK_IF((format_ != ge::FORMAT_NHWC),464 OP_CHECK_IF((format_ != ge::FORMAT_NHWC),
372 OP_LOGI(context_->GetNodeName(), "format is not eligible"),465 OP_LOGI(context_->GetNodeName(), "format is not eligible"),
@@ -375,7 +468,7 @@ bool ResizeNearestNeighborV2AscendCTilingImpl::IsMatchTiling_NHWC() {
375 OP_LOGI(context_->GetNodeName(), "c is small"),468 OP_LOGI(context_->GetNodeName(), "c is small"),
376 return false);469 return false);
377 470 
378- maxUbNum_ = (((ubSize_ / NUM_2) / ONE_BLOCK_SIZE) * ONE_BLOCK_SIZE) / dtypeSize_;471+ maxUbNum_ = (((ubSize_ / NUM_2) / ubBlockSize_) * ubBlockSize_) / dtypeSize_;
379 OP_LOGI(context_->GetNodeName(), "maxUbNum_ is %ld", maxUbNum_);472 OP_LOGI(context_->GetNodeName(), "maxUbNum_ is %ld", maxUbNum_);
380 splitBlockFactor_ = Ops::Base::CeilDiv(lenN_, coreNum_);473 splitBlockFactor_ = Ops::Base::CeilDiv(lenN_, coreNum_);
381 realCoreNum_ = Ops::Base::CeilDiv(lenN_, splitBlockFactor_);474 realCoreNum_ = Ops::Base::CeilDiv(lenN_, splitBlockFactor_);
@@ -390,6 +483,34 @@ bool ResizeNearestNeighborV2AscendCTilingImpl::IsMatchTiling_NHWC() {
390 return true;483 return true;
391}484}
392 485 
486+bool ResizeNearestNeighborV2AscendCTilingImpl::IsMatchTiling_NCHW_Gather()
487+{
488+ OP_CHECK_IF((format_ != ge::FORMAT_NCHW), OP_LOGI(context_->GetNodeName(), "format is not nchw"), return false);
489+ OP_CHECK_IF((lenDesW_ <= lenSrcW_) || (lenDesH_ < lenSrcH_),
490+ OP_LOGI(context_->GetNodeName(), "It must be an enlarged scene"), return false);
491+ OP_CHECK_IF((originalScaleH_ != 0.0f) || (originalScaleW_ != 0.0f),
492+ OP_LOGI(context_->GetNodeName(), "originalScaleH or originalScaleW is not 0.0f"), return false);
493+ OP_CHECK_IF((scaleH_ < REDUCE_SCALE_THRESHOLD) || (scaleW_ < REDUCE_SCALE_THRESHOLD),
494+ OP_LOGI(context_->GetNodeName(), "originalScaleH or originalScaleW samll than 0.25"), return false);
495+ int64_t xSize = xShape_.GetShapeSize();
496+ int64_t ySize = yShape_.GetShapeSize();
497+ OP_CHECK_IF(xSize >= UINT32_MAX || ySize >= UINT32_MAX,
498+ OP_LOGI(context_->GetNodeName(), "input or output size is too large"), return false);
499+ int64_t ncNum = lenN_ * lenC_;
500+ int64_t factorCore = Ops::Base::CeilDiv(ncNum, coreNum_);
501+ int64_t needCore = Ops::Base::CeilDiv(ncNum, factorCore);
502+ float scaleCore = 0.67; // 分核达不到2/3的总核数，不走此分支
503+ float rate = static_cast<float>(needCore) / static_cast<float>(coreNum_);
504+ OP_CHECK_IF(rate < scaleCore, OP_LOGI(context_->GetNodeName(), "nc is small"), return false);
505+ // 确保W不用切分
506+ int32_t idsUbSize = lenDesWAlign_ * dtypeSize_ * NUM_2;
507+ int32_t dstWUb = lenDesWAlign_ * dtypeSize_ * NUM_2;
508+ int32_t srcWUb = lenSrcWAlign_ * dtypeSize_ * NUM_2;
509+ int32_t leastUb = idsUbSize + dstWUb + srcWUb + ubBlockSize_ * NUM_2;
510+ OP_CHECK_IF(leastUb > ubSize_, OP_LOGI(context_->GetNodeName(), "ub is not enough"), return false);
511+ return true;
512+}
513+ 
393void ResizeNearestNeighborV2AscendCTilingImpl::DoTilingSmallC() {514void ResizeNearestNeighborV2AscendCTilingImpl::DoTilingSmallC() {
394 switchParams_ = cutNd_ ? 0 : NUM_2;515 switchParams_ = cutNd_ ? 0 : NUM_2;
395 nLoop_ = maxUbNum_ / lenCAlign_;516 nLoop_ = maxUbNum_ / lenCAlign_;
@@ -572,6 +693,10 @@ void ResizeNearestNeighborV2AscendCTilingImpl::MatchTilingStrategyAndSetTilingKe
572 alignCorners_ = static_cast<int64_t>(0);693 alignCorners_ = static_cast<int64_t>(0);
573 halfPixelCenters_ = static_cast<int64_t>(0);694 halfPixelCenters_ = static_cast<int64_t>(0);
574 idxUseInt32_ = static_cast<uint64_t>(0);695 idxUseInt32_ = static_cast<uint64_t>(0);
696+ } else if (IsMatchTiling_NCHW_Gather()) {
697+ schId_ = SCHEDULE_ID_GATHER_ALL_HW;
698+ idxUseInt32_ = static_cast<uint64_t>(1);
699+ format_ = ge::FORMAT_NCHW;
575 } else {700 } else {
576 idxUseInt32_ = xShape_.GetShapeSize() < UINT32_MAX && yShape_.GetShapeSize() < UINT32_MAX;701 idxUseInt32_ = xShape_.GetShapeSize() < UINT32_MAX && yShape_.GetShapeSize() < UINT32_MAX;
577 if (lenSrcH_ == lenDesH_ && lenSrcW_ == lenDesW_) {702 if (lenSrcH_ == lenDesH_ && lenSrcW_ == lenDesW_) {
@@ -709,6 +834,10 @@ void ResizeNearestNeighborV2AscendCTilingImpl::TilingStrategy()
709 DoTilingUb2Ub();834 DoTilingUb2Ub();
710 break;835 break;
711 }836 }
837+ case SCHEDULE_ID_GATHER_ALL_HW: {
838+ DoTilingGatherNchw();
839+ break;
840+ }
712 default: {841 default: {
713 realCoreNum_ = (yShape_.GetShapeSize() < coreNum_) ? yShape_.GetShapeSize() : coreNum_;842 realCoreNum_ = (yShape_.GetShapeSize() < coreNum_) ? yShape_.GetShapeSize() : coreNum_;
714 splitBlockFactor_ = Ops::Base::FloorDiv(yShape_.GetShapeSize(), realCoreNum_);843 splitBlockFactor_ = Ops::Base::FloorDiv(yShape_.GetShapeSize(), realCoreNum_);
@@ -744,7 +873,7 @@ void ResizeNearestNeighborV2AscendCTilingImpl::FillTilingData() {
744 tilingData_.set_nLoopTailLast(nLoopTailLast_);873 tilingData_.set_nLoopTailLast(nLoopTailLast_);
745 tilingData_.set_wcLoop(wcLoop_);874 tilingData_.set_wcLoop(wcLoop_);
746 tilingData_.set_wcLoopTimesBefore(wcLoopTimesBefore_);875 tilingData_.set_wcLoopTimesBefore(wcLoopTimesBefore_);
747- tilingData_.set_wcLoopTailBefore(wcLoopTailBefore_) ;876+ tilingData_.set_wcLoopTailBefore(wcLoopTailBefore_);
748 tilingData_.set_wcLoopTimesLast(wcLoopTimesLast_);877 tilingData_.set_wcLoopTimesLast(wcLoopTimesLast_);
749 tilingData_.set_wcLoopTailLast(wcLoopTailLast_);878 tilingData_.set_wcLoopTailLast(wcLoopTailLast_);
750 tilingData_.set_splitBlockFullCount(splitBlockFullCount_);879 tilingData_.set_splitBlockFullCount(splitBlockFullCount_);
@@ -770,7 +899,8 @@ void ResizeNearestNeighborV2AscendCTilingImpl::PrintTilingData() {
770 hwcNum: %ld, dstHwcNum:%ld, wcNum:%ld, dstWcNum:%ld, nLoop: %ld, nLoopTimesBefore: %ld, \899 hwcNum: %ld, dstHwcNum:%ld, wcNum:%ld, dstWcNum:%ld, nLoop: %ld, nLoopTimesBefore: %ld, \
771 nLoopTimesLast is %ld, nLoopTailLast: %ld, wcLoop: %ld, wcLoopTimesBefore: %ld, \900 nLoopTimesLast is %ld, nLoopTailLast: %ld, wcLoop: %ld, wcLoopTimesBefore: %ld, \
772 wcLoopTailBefore: %ld, wcLoopTimesLast: %ld, wcLoopTailLast: %ld, \901 wcLoopTailBefore: %ld, wcLoopTimesLast: %ld, wcLoopTailLast: %ld, \
773- splitFactorDesH %ld, splitFactorTailDesW %ld, splitFactorDesW %ld, scaleW: %f, scaleH：%f",902+ splitFactorDesH %ld, splitFactorTailDesW %ld, splitFactorDesW %ld, \
903+ splitCountDesH: %ld, scaleW: %f, scaleH：%f",
774 tilingData_.get_realCoreNum(),904 tilingData_.get_realCoreNum(),
775 tilingData_.get_ubSize(),905 tilingData_.get_ubSize(),
776 tilingData_.get_alignCorners(),906 tilingData_.get_alignCorners(),
@@ -802,6 +932,7 @@ void ResizeNearestNeighborV2AscendCTilingImpl::PrintTilingData() {
802 tilingData_.get_splitFactorDesH(),932 tilingData_.get_splitFactorDesH(),
803 tilingData_.get_splitFactorTailDesW(),933 tilingData_.get_splitFactorTailDesW(),
804 tilingData_.get_splitFactorDesW(),934 tilingData_.get_splitFactorDesW(),
935+ tilingData_.get_splitCountDesH(),
805 tilingData_.get_scaleW(),936 tilingData_.get_scaleW(),
806 tilingData_.get_scaleH());937 tilingData_.get_scaleH());
807}938}
@@ -899,7 +1030,7 @@ ge::graphStatus ResizeNearestNeighborV2AscendCTilingImpl::GetAttrInfo()
899 originalScaleW_ = scalesData[SCALE_W];1030 originalScaleW_ = scalesData[SCALE_W];
900 OP_LOGI(context_->GetNodeName(), "original scales(%f, %f)", originalScaleH_, originalScaleW_);1031 OP_LOGI(context_->GetNodeName(), "original scales(%f, %f)", originalScaleH_, originalScaleW_);
901 }1032 }
902- 1033+ bias_ = halfPixelCenters_ ? HALF_PIXEL : 0.0f;
903 return ge::GRAPH_SUCCESS;1034 return ge::GRAPH_SUCCESS;
904}1035}
905 1036 
@@ -951,12 +1082,13 @@ ge::graphStatus ResizeNearestNeighborV2AscendCTilingImpl::Init() {
951 return ge::GRAPH_FAILED);1082 return ge::GRAPH_FAILED);
952 uint64_t ubSize = 0;1083 uint64_t ubSize = 0;
953 ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);1084 ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
954- OP_CHECK_IF(1085+ ubBlockSize_ = Ops::Base::GetUbBlockSize(context_);
955- ubSize <= 0UL,1086+ OP_CHECK_IF(ubSize <= 0UL || ubBlockSize_ < 0,
956- OP_LOGE(context_, "ubSize must greater than zero, but is %lu", ubSize),1087+ OP_LOGE(context_, "ubSize and ubBlockSize must > 0, but ubSize is %lu, ubBlockSizeis %ld", ubSize,
1088+ ubBlockSize_),
957 return ge::GRAPH_FAILED);1089 return ge::GRAPH_FAILED);
958 ubSize_ = static_cast<int64_t>(ubSize);1090 ubSize_ = static_cast<int64_t>(ubSize);
959- OP_LOGI(context_->GetNodeName(), "coreNum_ is %ld, ubSize_ is %ld", coreNum_, ubSize_);1091+ OP_LOGI(context_->GetNodeName(), "coreNum_:%ld, ubSize_:%ld, ubBlockSize_:%ld", coreNum_, ubSize_, ubBlockSize_);
960 // Get attrs: alignCorners, halfPixelCenters, scales1092 // Get attrs: alignCorners, halfPixelCenters, scales
961 OP_CHECK_IF((GetAttrInfo() != ge::GRAPH_SUCCESS), OP_LOGE(context_->GetNodeName(), "GetAttrInfo failed."),1093 OP_CHECK_IF((GetAttrInfo() != ge::GRAPH_SUCCESS), OP_LOGE(context_->GetNodeName(), "GetAttrInfo failed."),
962 return ge::GRAPH_FAILED);1094 return ge::GRAPH_FAILED);
@@ -994,7 +1126,9 @@ ge::graphStatus ResizeNearestNeighborV2AscendCTilingImpl::DoTiling() {
994 TilingStrategy();1126 TilingStrategy();
995 FillTilingData();1127 FillTilingData();
996 PrintTilingData();1128 PrintTilingData();
997- 1129+ if (schId_ == SCHEDULE_ID_GATHER_ALL_HW || schId_ == SCHEDULE_ID_GATHER_CUT_H) {
1130+ halfPixelCenters_ = 0;
1131+ }
998 context_->SetBlockDim(tilingData_.get_realCoreNum());1132 context_->SetBlockDim(tilingData_.get_realCoreNum());
999 const uint64_t tilingKey = GET_TPL_TILING_KEY(schId_, (uint64_t)format_, (uint64_t)alignCorners_,1133 const uint64_t tilingKey = GET_TPL_TILING_KEY(schId_, (uint64_t)format_, (uint64_t)alignCorners_,
1000 (uint64_t)halfPixelCenters_, idxUseInt32_);1134 (uint64_t)halfPixelCenters_, idxUseInt32_);
Aimage/resize_nearest_neighbor_v2/op_kernel/arch35/resize_nearest_neighbor_v2_nchw_gather.h+696-0
@@ -0,0 +1,696 @@
1+/* *
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file resize_nearest_neighbor_v2_nchw_gather_hw.h
13+ * \brief resize_nearest_neighbor_v2_nchw_gather_hw.h
14+ */
15+#ifndef RESIZE_NEAREAST_NEIGHBOR_V2_NCHW_GATHER_H
16+#define RESIZE_NEAREAST_NEIGHBOR_V2_NCHW_GATHER_H
17+ 
18+#include "op_kernel/platform_util.h"
19+#include "kernel_operator.h"
20+#include "op_kernel/math_util.h"
21+#include "resize_nearest_neighbor_v2_tiling_key.h"
22+ 
23+namespace ResizeNearestNeighborV2 {
24+using namespace AscendC;
25+using AscendC::MicroAPI::AddrReg;
26+using AscendC::MicroAPI::CreateAddrReg;
27+using AscendC::MicroAPI::RegTensor;
28+using AscendC::MicroAPI::UpdateMask;
29+ 
30+template <typename T, typename T1, int schId, bool alignCorners>
31+class ResizeGather {
32+public:
33+ __aicore__ inline ResizeGather(){};
34+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR size, GM_ADDR y,
35+ const ResizeNearestNeighborV2TilingData *tilingData);
36+ __aicore__ inline void Process();
37+ 
38+private:
39+ __aicore__ inline int64_t ComputeOriIds(int64_t dstIds, int64_t maxIds, float scale, float bais);
40+ __aicore__ inline void GatherOutput(LocalTensor<T> &outUb, LocalTensor<T> &srcUb, int64_t ubFactor, int64_t srcLen,
41+ int64_t dstLen, int64_t srcHwStart);
42+ __aicore__ inline void DataCopyIn(LocalTensor<T> &xLocal, int64_t blockCount, int64_t blockLen, int64_t srcStride,
43+ int64_t offset);
44+ __aicore__ inline void DataCopyOut(LocalTensor<T> &yLocal, int64_t blockCount, int64_t blockLen, int64_t srcStride,
45+ int64_t offset);
46+ __aicore__ inline void ComputeDataCopyGather(LocalTensor<T> &srcUb, LocalTensor<T> &outUb, LocalTensor<T1> &idxHwUb, int64_t num);
47+ __aicore__ inline void ComputeAllHw();
48+ __aicore__ inline void ComputeIdsSpecial(int64_t dstWSizeAlgin, int64_t srcWSizeAlgin);
49+ __aicore__ inline void ComputeCutH();
50+ __aicore__ inline void ComputeOriHIdx(LocalTensor<T1> &idxHUb, LocalTensor<T1> &idxH1Ub, int64_t onceHsize,
51+ int64_t hoStart, int64_t hiStart);
52+ __aicore__ inline void ComputeOriHWidx(LocalTensor<T1> &idxWUb, LocalTensor<T1> &idxH1Ub, LocalTensor<T1> &idxHwUb,
53+ int64_t onceHsize);
54+ __aicore__ inline void ComputeHWids(LocalTensor<T1> &idxUb, LocalTensor<T1> &idxHUb, LocalTensor<T1> &idxWUb,
55+ int64_t dstWSizeAlgin);
56+ __aicore__ inline void ComputeHids(LocalTensor<T1> &idxHUb, int64_t hFactor);
57+ constexpr static int32_t bufferNum = 2;
58+ 
59+ constexpr static AscendC::MicroAPI::CastTrait castTraitRound = { AscendC::MicroAPI::RegLayout::UNKNOWN,
60+ AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_ROUND };
61+ constexpr static AscendC::MicroAPI::CastTrait castTraitFloor = { AscendC::MicroAPI::RegLayout::UNKNOWN,
62+ AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_FLOOR };
63+ constexpr static AscendC::MicroAPI::CastTrait castInt32ToF = { AscendC::MicroAPI::RegLayout::UNKNOWN,
64+ AscendC::MicroAPI::SatMode::UNKNOWN, AscendC::MicroAPI::MaskMergeMode::ZEROING,
65+ AscendC::RoundMode::CAST_FLOOR };
66+ 
67+private:
68+ TPipe pipe;
69+ TQue<QuePosition::VECIN, 1> inQue_;
70+ TQue<QuePosition::VECOUT, 1> outQue_;
71+ TBuf<QuePosition::VECCALC> idxBuf_;
72+ TBuf<QuePosition::VECCALC> idxDstBuf_;
73+ TBuf<QuePosition::VECCALC> idxHBuf_;
74+ TBuf<QuePosition::VECCALC> idxH1Buf_;
75+ TBuf<QuePosition::VECCALC> idxWBuf_;
76+ GlobalTensor<T> xGm_;
77+ GlobalTensor<T> yGm_;
78+ 
79+ int32_t blockIdx_ = 0;
80+ int32_t vlLen_ = Ops::Base::GetVRegSize() / sizeof(T);
81+ int32_t vlLenB32_ = Ops::Base::GetVRegSize() / sizeof(int32_t);
82+ float hScale_ = 0.0f;
83+ float wScale_ = 0.0f;
84+ int64_t srcHSize_ = 0;
85+ int64_t srcWSize_ = 0;
86+ int64_t srcWAlignSize_ = 0;
87+ int64_t dstHSize_ = 0;
88+ int64_t dstWSize_ = 0;
89+ int64_t dstWAlignSize_ = 0;
90+ int64_t ncFactor_ = 0;
91+ int64_t dstHwNum_ = 0;
92+ int64_t srcHwNum_ = 0;
93+ int64_t xUb_ = 0;
94+ int64_t yUb_ = 0;
95+ int64_t idsUb_ = 0;
96+ int64_t tailTimes_ = 0;
97+ int64_t hTimes_ = 0; // 切h时表示hLoopTimes
98+ int64_t beforeTimes_ = 0;
99+ int64_t tailTNum_ = 0;
100+ int64_t beforeTNum_ = 0;
101+ int64_t ubFactor_ = 0;
102+ int64_t blockFactor_ = 0;
103+ int64_t realCoreNum_ = 0;
104+ float bias_ = 0.0f;
105+ int32_t blockNumSize_ = Ops::Base::GetUbBlockSize() / sizeof(T);
106+ // tiling params
107+ const ResizeNearestNeighborV2TilingData *tiling_;
108+ DataCopyPadExtParams<T> padParams_{ false, 0, 0, 0 };
109+ DataCopyExtParams copyParams_{ 1, 1, 0, 0, 0 };
110+};
111+ 
112+template <typename T, typename T1, int schId, bool alignCorners>
113+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::Init(GM_ADDR x, GM_ADDR size, GM_ADDR y,
114+ const ResizeNearestNeighborV2TilingData *tilingData)
115+{
116+ blockIdx_ = GetBlockIdx();
117+ xGm_.SetGlobalBuffer((__gm__ T *)x);
118+ yGm_.SetGlobalBuffer((__gm__ T *)y);
119+ tiling_ = tilingData;
120+ 
121+ // 接收tilingdata信息
122+ 
123+ hScale_ = tiling_->scaleH;
124+ wScale_ = tiling_->scaleW;
125+ srcHSize_ = tiling_->lenSrcH;
126+ srcWSize_ = tiling_->lenSrcW;
127+ dstHSize_ = tiling_->lenDesH;
128+ dstWSize_ = tiling_->lenDesW;
129+ 
130+ dstWAlignSize_ = (dstWSize_ + blockNumSize_ - 1) / blockNumSize_ * blockNumSize_;
131+ srcWAlignSize_ = (srcWSize_ + blockNumSize_ - 1) / blockNumSize_ * blockNumSize_;
132+ 
133+ dstHwNum_ = dstHSize_ * dstWSize_;
134+ srcHwNum_ = srcHSize_ * srcWSize_;
135+ xUb_ = tilingData->splitFactorTailDesW; // 输入x的ub大小
136+ yUb_ = tilingData->splitFactorDesW; // 输出y的ub大小
137+ idsUb_ = tilingData->ubSize;
138+ realCoreNum_ = tilingData->realCoreNum;
139+ tailTimes_ = tilingData->nLoopTimesLast;
140+ beforeTimes_ = tilingData->nLoopTimesBefore;
141+ tailTNum_ = tilingData->nLoopTailLast;
142+ 
143+ beforeTNum_ = tilingData->nLoop;
144+ ubFactor_ = tilingData->splitFactorDesH;
145+ blockFactor_ = tilingData->splitBlockFactor;
146+ if (tilingData->halfPixelCenters == 1) {
147+ bias_ = 0.5f;
148+ }
149+ hTimes_ = tilingData->splitCountDesH; // 切h时,h反向循环次数
150+ 
151+ pipe.InitBuffer(inQue_, bufferNum, xUb_);
152+ pipe.InitBuffer(outQue_, bufferNum, yUb_);
153+ pipe.InitBuffer(idxBuf_, idsUb_); // 存放h*w坐标
154+ pipe.InitBuffer(idxWBuf_, dstWAlignSize_ * sizeof(T1)); // 放w的坐标
155+ if (schId == TPL_SCH_MODE_GATHER_ALL_HW) {
156+ int64_t dstHAlignSize = (dstHSize_ + blockNumSize_ - 1) / blockNumSize_ * blockNumSize_;
157+ pipe.InitBuffer(idxHBuf_, dstHAlignSize * sizeof(T1)); // 放h的坐标
158+ }
159+ if constexpr (schId == TPL_SCH_MODE_GATHER_CUT_H) {
160+ int64_t dstHAlignSize = (ubFactor_ + blockNumSize_ - 1) / blockNumSize_ * blockNumSize_;
161+ pipe.InitBuffer(idxH1Buf_, dstHAlignSize * sizeof(T1)); // 对应的输入h坐标
162+ pipe.InitBuffer(idxHBuf_, dstHAlignSize * sizeof(T1)); // 输出h坐标
163+ }
164+}
165+ 
166+template <typename T, typename T1, int schId, bool alignCorners>
167+__aicore__ inline int64_t ResizeGather<T, T1, schId, alignCorners>::ComputeOriIds(int64_t dstIds, int64_t maxIds,
168+ float scale, float bais)
169+{
170+ int64_t srcIds = 0;
171+ if constexpr (alignCorners) {
172+ // round
173+ srcIds = static_cast<int64_t>(static_cast<int32_t>(((static_cast<float>(dstIds) + bais) * scale) + 0.5f));
174+ } else {
175+ // floor
176+ srcIds = static_cast<int64_t>(static_cast<int32_t>((static_cast<float>(dstIds) + bais) * scale));
177+ }
178+ if (srcIds > maxIds) {
179+ srcIds = maxIds;
180+ }
181+ return srcIds;
182+}
183+ 
184+template <typename T, typename T1, int schId, bool alignCorners>
185+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::GatherOutput(LocalTensor<T> &outUb,
186+ LocalTensor<T> &srcUb, int64_t ubFactor, int64_t srcLen, int64_t dstLen, int64_t srcHwStart)
187+{
188+ uint16_t times = CeilDivision(dstLen, vlLen_);
189+ LocalTensor<T1> idxUb = idxBuf_.AllocTensor<T1>();
190+ auto idxUbAddr = (__ubuf__ T1 *)idxUb.GetPhyAddr();
191+ auto srcUbAddr = (__ubuf__ T *)srcUb.GetPhyAddr();
192+ auto dstUbAddr = (__ubuf__ T *)outUb.GetPhyAddr();
193+ T1 srcHwNum = srcHwStart;
194+ auto dstUbAddr1 = (__ubuf__ T *)outUb[dstLen].GetPhyAddr();
195+ 
196+ uint32_t vfLen = vlLen_;
197+ uint32_t dstHwAlign = dstLen;
198+ T1 srcLenNum = srcLen;
199+ uint16_t ubFactorTimes = ubFactor > 1 ? static_cast<uint16_t>(ubFactor) - 1 : 0;
200+ uint16_t timesNc = ubFactorTimes == 0 ? 0 : 1;
201+ uint32_t hwNum = dstLen;
202+ uint32_t hwNum1 = dstLen;
203+ __VEC_SCOPE__
204+ {
205+ AscendC::MicroAPI::RegTensor<T1> startReg;
206+ AscendC::MicroAPI::RegTensor<T1> idxRegT;
207+ AscendC::MicroAPI::RegTensor<T> dstReg;
208+ AscendC::MicroAPI::MaskReg preg;
209+ // 先处理第一行的hw
210+ Duplicate<T1>(startReg, srcHwNum);
211+ for (uint16_t j = 0; j < times; j++) {
212+ preg = AscendC::MicroAPI::UpdateMask<T>(hwNum);
213+ AscendC::MicroAPI::AddrReg srcIdxOffset = AscendC::MicroAPI::CreateAddrReg<T1>(j, vfLen);
214+ AscendC::MicroAPI::DataCopy(idxRegT, idxUbAddr, srcIdxOffset);
215+ AscendC::MicroAPI::Sub(idxRegT, idxRegT, startReg, preg);
216+ DataCopyGather(dstReg, srcUbAddr, idxRegT, preg);
217+ AscendC::MicroAPI::DataCopy(dstUbAddr, dstReg, srcIdxOffset, preg);
218+ }
219+ // 从第二行开始处理
220+ for (uint16_t nc = 0; nc < timesNc; nc++) {
221+ for (uint16_t jj = 0; jj < times; jj++) {
222+ preg = AscendC::MicroAPI::UpdateMask<T>(hwNum1);
223+ AscendC::MicroAPI::AddrReg srcIdxOffset = AscendC::MicroAPI::CreateAddrReg<T1>(jj, vfLen);
224+ AscendC::MicroAPI::DataCopy(idxRegT, idxUbAddr, srcIdxOffset);
225+ AscendC::MicroAPI::Sub(idxRegT, idxRegT, startReg, preg);
226+ for (uint16_t i = 0; i < ubFactorTimes; i++) {
227+ AscendC::MicroAPI::AddrReg outOffset =
228+ AscendC::MicroAPI::CreateAddrReg<T>(jj, vfLen, i, dstHwAlign);
229+ Adds(idxRegT, idxRegT, srcLenNum, preg);
230+ DataCopyGather(dstReg, srcUbAddr, idxRegT, preg);
231+ AscendC::MicroAPI::DataCopy(dstUbAddr1, dstReg, outOffset, preg);
232+ }
233+ }
234+ }
235+ }
236+}
237+ 
238+template <typename T, typename T1, int schId, bool alignCorners>
239+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::DataCopyIn(LocalTensor<T> &xLocal, int64_t blockCount,
240+ int64_t blockLen, int64_t srcStride, int64_t offset)
241+{
242+ copyParams_.blockCount = blockCount;
243+ copyParams_.blockLen = blockLen * sizeof(T);
244+ copyParams_.srcStride = srcStride * sizeof(T);
245+ copyParams_.dstStride = 0;
246+ AscendC::DataCopyPad(xLocal, xGm_[offset], copyParams_, padParams_);
247+}
248+ 
249+template <typename T, typename T1, int schId, bool alignCorners>
250+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::DataCopyOut(LocalTensor<T> &yLocal, int64_t blockCount,
251+ int64_t blockLen, int64_t srcStride, int64_t offset)
252+{
253+ copyParams_.blockCount = blockCount;
254+ copyParams_.blockLen = blockLen * sizeof(T);
255+ copyParams_.srcStride = srcStride * sizeof(T);
256+ copyParams_.dstStride = 0;
257+ AscendC::DataCopyPad(yGm_[offset], yLocal, copyParams_);
258+}
259+ 
260+template <typename T, typename T1, int schId, bool alignCorners>
261+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::ComputeDataCopyGather(LocalTensor<T> &srcUb,
262+ LocalTensor<T> &outUb, LocalTensor<T1> &idxHwUb, int64_t num)
263+{
264+ uint16_t times = CeilDivision(num, vlLen_);
265+ auto idxUbAddr = (__ubuf__ T1 *)idxHwUb.GetPhyAddr();
266+ auto srcUbAddr = (__ubuf__ T *)srcUb.GetPhyAddr();
267+ auto dstUbAddr = (__ubuf__ T *)outUb.GetPhyAddr();
268+ uint32_t vfLen = vlLen_;
269+ uint32_t onceSize = num;
270+ __VEC_SCOPE__
271+ {
272+ AscendC::MicroAPI::RegTensor<T1> idxRegT;
273+ AscendC::MicroAPI::RegTensor<T> dstReg;
274+ AscendC::MicroAPI::MaskReg preg;
275+ for (uint16_t j = 0; j < times; j++) {
276+ preg = AscendC::MicroAPI::UpdateMask<T>(onceSize);
277+ AscendC::MicroAPI::AddrReg srcIdxOffset = AscendC::MicroAPI::CreateAddrReg<T1>(j, vfLen);
278+ AscendC::MicroAPI::DataCopy(idxRegT, idxUbAddr, srcIdxOffset);
279+ DataCopyGather(dstReg, srcUbAddr, idxRegT, preg);
280+ AscendC::MicroAPI::DataCopy(dstUbAddr, dstReg, srcIdxOffset, preg);
281+ }
282+ }
283+}
284+ 
285+template <typename T, typename T1, int schId, bool alignCorners>
286+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::ComputeHids(LocalTensor<T1> &idxHUb, int64_t hFactor)
287+{
288+ auto idxUbAddr = (__ubuf__ T1 *)idxHUb.GetPhyAddr();
289+ auto idxUbRemainAddr = (__ubuf__ T1 *)idxHUb[vlLenB32_].GetPhyAddr();
290+ uint16_t times = 0;
291+ uint32_t remainNum = 0;
292+ uint32_t numH = vlLenB32_;
293+ if (hFactor > static_cast<int64_t>(vlLenB32_)) {
294+ remainNum = hFactor - static_cast<int64_t>(vlLenB32_);
295+ times = CeilDivision(remainNum, static_cast<uint32_t>(vlLenB32_));
296+ } else {
297+ numH = hFactor;
298+ }
299+ 
300+ __VEC_SCOPE__
301+ {
302+ AscendC::MicroAPI::RegTensor<int32_t> idxInt32Reg;
303+ AscendC::MicroAPI::MaskReg pregB32 = AscendC::MicroAPI::UpdateMask<uint32_t>(numH);
304+ AscendC::MicroAPI::MaskReg pregRemainB32;
305+ Arange(idxInt32Reg, 0);
306+ if constexpr (sizeof(T1) == sizeof(int32_t)) {
307+ DataCopy(idxUbAddr, (MicroAPI::RegTensor<T1> &)idxInt32Reg, pregB32);
308+ } else {
309+ DataCopy<T1, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(idxUbAddr, (MicroAPI::RegTensor<T1> &)idxInt32Reg,
310+ pregB32);
311+ }
312+ for (uint16_t i = 0; i < times; i++) {
313+ pregRemainB32 = AscendC::MicroAPI::UpdateMask<int32_t>(remainNum);
314+ Adds(idxInt32Reg, idxInt32Reg, 64, pregRemainB32);
315+ AscendC::MicroAPI::AddrReg dstOffset = AscendC::MicroAPI::CreateAddrReg<T1>(i, 64);
316+ if constexpr (sizeof(T1) == sizeof(int32_t)) {
317+ DataCopy(idxUbRemainAddr, (MicroAPI::RegTensor<T1> &)idxInt32Reg, dstOffset, pregRemainB32);
318+ } else {
319+ DataCopy<T1, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(idxUbRemainAddr,
320+ (MicroAPI::RegTensor<T1> &)idxInt32Reg, dstOffset, pregRemainB32);
321+ }
322+ }
323+ }
324+}
325+ 
326+template <typename T1, bool isH, bool alignCorners>
327+__aicore__ inline void ComputeHOrWids(LocalTensor<T1> &idxUb, float bias, float scale, int64_t srcSize, int64_t dstSize,
328+ int64_t srcWsize)
329+{
330+ uint32_t vfLenb32 = Ops::Base::GetVRegSize() / sizeof(float);
331+ auto idxUbAddr = (__ubuf__ T1 *)idxUb.GetPhyAddr();
332+ auto idxUbRemainAddr = (__ubuf__ T1 *)idxUb[vfLenb32].GetPhyAddr();
333+ uint32_t oneTimeNum = vfLenb32;
334+ uint32_t remainNum = 0;
335+ int32_t srcW = srcWsize;
336+ int32_t maxData = srcSize - 1;
337+ uint16_t times = 0;
338+ 
339+ if (dstSize <= static_cast<int64_t>(vfLenb32)) {
340+ oneTimeNum = dstSize;
341+ } else {
342+ remainNum = dstSize - vfLenb32;
343+ times = CeilDivision(remainNum, vfLenb32);
344+ }
345+ constexpr static AscendC::MicroAPI::CastTrait castTraitRound = { AscendC::MicroAPI::RegLayout::UNKNOWN,
346+ AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_ROUND };
347+ constexpr static AscendC::MicroAPI::CastTrait castTraitFloor = { AscendC::MicroAPI::RegLayout::UNKNOWN,
348+ AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING, AscendC::RoundMode::CAST_FLOOR };
349+ constexpr static AscendC::MicroAPI::CastTrait castInt32ToF = { AscendC::MicroAPI::RegLayout::UNKNOWN,
350+ AscendC::MicroAPI::SatMode::UNKNOWN, AscendC::MicroAPI::MaskMergeMode::ZEROING,
351+ AscendC::RoundMode::CAST_FLOOR };
352+ 
353+ __VEC_SCOPE__
354+ {
355+ AscendC::MicroAPI::RegTensor<float> idxInt32Reg;
356+ AscendC::MicroAPI::RegTensor<float> hIdxF;
357+ 
358+ AscendC::MicroAPI::RegTensor<int32_t> hIdxInt32;
359+ AscendC::MicroAPI::RegTensor<int32_t> hIdxInt32C;
360+ AscendC::MicroAPI::MaskReg pregB32 = AscendC::MicroAPI::UpdateMask<uint32_t>(oneTimeNum);
361+ AscendC::MicroAPI::MaskReg pregRemainB32;
362+ Arange(idxInt32Reg, 0.0f);
363+ 
364+ Adds(hIdxF, idxInt32Reg, bias, pregB32);
365+ 
366+ Muls(hIdxF, hIdxF, scale, pregB32);
367+ 
368+ if constexpr (alignCorners == 1) {
369+ Cast<int32_t, float, castTraitRound>(hIdxInt32, hIdxF, pregB32);
370+ } else {
371+ Cast<int32_t, float, castTraitFloor>(hIdxInt32, hIdxF, pregB32);
372+ }
373+ Mins(hIdxInt32C, hIdxInt32, maxData, pregB32);
374+ if constexpr (isH) {
375+ Muls(hIdxInt32C, hIdxInt32C, srcW, pregB32);
376+ }
377+ if constexpr (sizeof(T1) == sizeof(int32_t)) {
378+ DataCopy(idxUbAddr, (MicroAPI::RegTensor<T1> &)hIdxInt32C, pregB32);
379+ } else {
380+ DataCopy<T1, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(idxUbAddr, (MicroAPI::RegTensor<T1> &)hIdxInt32C,
381+ pregB32);
382+ }
383+ for (uint16_t i = 0; i < times; i++) {
384+ pregRemainB32 = AscendC::MicroAPI::UpdateMask<int32_t>(remainNum);
385+ Adds(idxInt32Reg, idxInt32Reg, 64.0f, pregRemainB32);
386+ Adds(hIdxF, idxInt32Reg, bias, pregRemainB32);
387+ Muls(hIdxF, hIdxF, scale, pregRemainB32);
388+ if constexpr (alignCorners == 1) {
389+ Cast<int32_t, float, castTraitRound>(hIdxInt32, hIdxF, pregRemainB32);
390+ } else {
391+ Cast<int32_t, float, castTraitFloor>(hIdxInt32, hIdxF, pregRemainB32);
392+ }
393+ Mins(hIdxInt32C, hIdxInt32, maxData, pregRemainB32);
394+ if constexpr (isH) {
395+ Muls(hIdxInt32C, hIdxInt32C, srcW, pregRemainB32);
396+ }
397+ AscendC::MicroAPI::AddrReg dstOffset = AscendC::MicroAPI::CreateAddrReg<T1>(i, 64);
398+ if constexpr (sizeof(T1) == sizeof(int32_t)) {
399+ DataCopy(idxUbRemainAddr, (MicroAPI::RegTensor<T1> &)hIdxInt32C, dstOffset, pregRemainB32);
400+ } else {
401+ DataCopy<T1, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(idxUbRemainAddr,
402+ (MicroAPI::RegTensor<T1> &)hIdxInt32C, dstOffset, pregRemainB32);
403+ }
404+ }
405+ }
406+}
407+ 
408+template <typename T, typename T1, int schId, bool alignCorners>
409+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::ComputeOriHIdx(LocalTensor<T1> &idxHUb,
410+ LocalTensor<T1> &idxH1Ub, int64_t onceHsize, int64_t hoStart, int64_t hiStart)
411+{
412+ // 先计算h的原始坐标
413+ auto idxHubAddr = (__ubuf__ T1 *)idxHUb.GetPhyAddr();
414+ auto idxH1UbAddr = (__ubuf__ T1 *)idxH1Ub.GetPhyAddr();
415+ uint32_t size = onceHsize;
416+ uint32_t vfLenB32 = vlLenB32_;
417+ uint16_t hTimes = CeilDivision(size, vfLenB32);
418+ float bias = bias_;
419+ float hScale = hScale_;
420+ int32_t maxData = srcHSize_ - 1;
421+ int32_t wSize = srcWAlignSize_;
422+ int32_t hoStartData = hoStart;
423+ int64_t hiStartData = hiStart;
424+ __VEC_SCOPE__
425+ {
426+ AscendC::MicroAPI::RegTensor<int32_t> hiStartReg;
427+ AscendC::MicroAPI::RegTensor<int32_t> idxInt32Reg;
428+ AscendC::MicroAPI::RegTensor<float> hIdxF;
429+ AscendC::MicroAPI::RegTensor<int32_t> idxInt32OriReg;
430+ AscendC::MicroAPI::RegTensor<int32_t> idxInt32OriWReg;
431+ AscendC::MicroAPI::MaskReg preg;
432+ Duplicate<int32_t>(hiStartReg, hiStartData);
433+ 
434+ for (uint16_t i = 0; i < hTimes; i++) {
435+ preg = AscendC::MicroAPI::UpdateMask<int32_t>(size);
436+ AscendC::MicroAPI::AddrReg srcIdxOffset = AscendC::MicroAPI::CreateAddrReg<T1>(i, vfLenB32);
437+ if constexpr (sizeof(T1) == sizeof(int32_t)) {
438+ DataCopy((MicroAPI::RegTensor<T1> &)idxInt32Reg, idxHubAddr, srcIdxOffset);
439+ } else {
440+ DataCopy<T1, MicroAPI::LoadDist::DIST_UNPACK_B16>((MicroAPI::RegTensor<T1> &)idxInt32Reg, idxHubAddr,
441+ srcIdxOffset);
442+ }
443+ //
444+ Adds(idxInt32Reg, idxInt32Reg, hoStartData, preg); // 输出位置
445+ Cast<float, int32_t, castInt32ToF>(hIdxF, idxInt32Reg, preg);
446+ Adds(hIdxF, hIdxF, bias, preg);
447+ Muls(hIdxF, hIdxF, hScale, preg);
448+ 
449+ if constexpr (alignCorners == 1) {
450+ Cast<int32_t, float, castTraitRound>(idxInt32OriReg, hIdxF, preg);
451+ } else {
452+ Cast<int32_t, float, castTraitFloor>(idxInt32OriReg, hIdxF, preg);
453+ }
454+ Mins(idxInt32OriReg, idxInt32OriReg, maxData, preg);
455+ Sub(idxInt32OriReg, idxInt32OriReg, hiStartReg, preg);
456+ Muls(idxInt32OriWReg, idxInt32OriReg, wSize, preg);
457+ if constexpr (sizeof(T1) == sizeof(int32_t)) {
458+ DataCopy(idxH1UbAddr, (MicroAPI::RegTensor<T1> &)idxInt32OriWReg, srcIdxOffset, preg);
459+ } else {
460+ DataCopy<T1, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(idxH1UbAddr,
461+ (MicroAPI::RegTensor<T1> &)idxInt32OriWReg, srcIdxOffset, preg);
462+ }
463+ }
464+ }
465+}
466+ 
467+template <typename T, typename T1, int schId, bool alignCorners>
468+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::ComputeOriHWidx(LocalTensor<T1> &idxWUb,
469+ LocalTensor<T1> &idxH1Ub, LocalTensor<T1> &idxHwUb, int64_t onceHsize)
470+{
471+ // 计算h*w的合轴的坐标
472+ auto idxHwUbAddr = (__ubuf__ T1 *)idxHwUb.GetPhyAddr();
473+ auto idxWubAddr = (__ubuf__ T1 *)idxWUb.GetPhyAddr();
474+ 
475+ uint16_t onceHTimes = onceHsize;
476+ uint32_t dstWSize = dstWAlignSize_;
477+ uint16_t wTimes = dstWAlignSize_ / vlLen_;
478+ uint16_t wTailTimes = 0;
479+ uint32_t tail = dstWAlignSize_ % vlLen_;
480+ if (tail != 0) {
481+ wTailTimes = 1;
482+ }
483+ uint32_t vfLen = vlLen_;
484+ auto idxWubAddr1 = (__ubuf__ T1 *)idxWUb[wTimes * vfLen].GetPhyAddr();
485+ auto idxHwUbAddr1 = (__ubuf__ T1 *)idxHwUb[wTimes * vfLen].GetPhyAddr();
486+ 
487+ __VEC_SCOPE__
488+ {
489+ AscendC::MicroAPI::RegTensor<T1> idxWReg;
490+ AscendC::MicroAPI::RegTensor<T1> addsReg;
491+ AscendC::MicroAPI::MaskReg preg = AscendC::MicroAPI::UpdateMask<T1>(tail);
492+ AscendC::MicroAPI::MaskReg pregB32 = AscendC::MicroAPI::CreateMask<T1, AscendC::MicroAPI::MaskPattern::ALL>();
493+ 
494+ for (uint16_t i = 0; i < onceHTimes; i++) {
495+ T1 hIdx = idxH1Ub.GetValue(i);
496+ AscendC::MicroAPI::AddrReg outIdxOffset1 = AscendC::MicroAPI::CreateAddrReg<T1>(i, dstWSize);
497+ for (uint16_t j = 0; j < static_cast<uint16_t>(wTimes); j++) {
498+ AscendC::MicroAPI::AddrReg srcIdxOffset = AscendC::MicroAPI::CreateAddrReg<T1>(j, vfLen);
499+ AscendC::MicroAPI::AddrReg outIdxOffset = AscendC::MicroAPI::CreateAddrReg<T1>(i, dstWSize, j, vfLen);
500+ DataCopy(idxWReg, idxWubAddr, srcIdxOffset);
501+ Adds(addsReg, idxWReg, hIdx, pregB32);
502+ DataCopy(idxHwUbAddr, addsReg, outIdxOffset, pregB32);
503+ }
504+ for (uint16_t jj = 0; jj < wTailTimes; jj++) {
505+ DataCopy(idxWReg, idxWubAddr1);
506+ Adds(addsReg, idxWReg, hIdx, preg);
507+ DataCopy(idxHwUbAddr1, addsReg, outIdxOffset1, preg);
508+ }
509+ }
510+ }
511+}
512+ 
513+template <typename T, typename T1, int schId, bool alignCorners>
514+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::ComputeCutH()
515+{
516+ LocalTensor<T1> idxWUb = idxWBuf_.Get<T1>();
517+ ComputeHOrWids<T1, false, alignCorners>(idxWUb, bias_, wScale_, srcWSize_, dstWAlignSize_, srcWAlignSize_);
518+ LocalTensor<T1> idxHUb = idxHBuf_.Get<T1>();
519+ ComputeHids(idxHUb, ubFactor_);
520+ LocalTensor<T1> idxH1Ub = idxH1Buf_.Get<T1>();
521+ LocalTensor<T1> idxHwUb = idxBuf_.Get<T1>();
522+ int64_t ncTimes = blockIdx_ == realCoreNum_ - 1 ? beforeTNum_ : blockFactor_; // 每个核处理的nc个数
523+ for (int64_t h = 0; h < hTimes_; h++) {
524+ int64_t onceHsize = h == hTimes_ - 1 ? tailTNum_ : ubFactor_; // 每次处理的h数量
525+ int64_t hoStart = h * ubFactor_;
526+ int64_t hoEnd = hoStart + onceHsize - 1;
527+ int64_t allSize = onceHsize * dstWAlignSize_;
528+ int64_t hiStart = static_cast<int32_t>(ComputeOriIds(hoStart, srcHSize_ - 1, hScale_, bias_));
529+ int64_t hiEnd = static_cast<int32_t>(ComputeOriIds(hoEnd, srcHSize_ - 1, hScale_, bias_));
530+ int64_t hiSize = hiEnd - hiStart + 1;
531+ ComputeOriHIdx(idxHUb, idxH1Ub, onceHsize, hoStart, hiStart);
532+ ComputeOriHWidx(idxWUb, idxH1Ub, idxHwUb, onceHsize);
533+ for (int64_t nc = 0; nc < ncTimes; nc++) {
534+ LocalTensor<T> xLocal = inQue_.AllocTensor<T>();
535+ int64_t inOffset = (blockIdx_ * blockFactor_ + nc) * srcHwNum_ + hiStart * srcWSize_;
536+ if (srcWAlignSize_ == srcWSize_) {
537+ DataCopyIn(xLocal, 1, hiSize * srcWSize_, 0, inOffset);
538+ } else {
539+ DataCopyIn(xLocal, hiSize, srcWSize_, 0, inOffset);
540+ }
541+ inQue_.EnQue<T>(xLocal);
542+ xLocal = inQue_.DeQue<T>();
543+ LocalTensor<T> yLocal = outQue_.AllocTensor<T>();
544+ ComputeDataCopyGather(xLocal, yLocal, idxHwUb, allSize);
545+ inQue_.FreeTensor(xLocal);
546+ int64_t outOffset = (blockIdx_ * blockFactor_ + nc) * dstHwNum_ + hoStart * dstWSize_;
547+ outQue_.EnQue<T>(yLocal);
548+ yLocal = outQue_.DeQue<T>();
549+ if (dstWAlignSize_ == dstWSize_) {
550+ DataCopyOut(yLocal, 1, allSize, 0, outOffset);
551+ } else {
552+ DataCopyOut(yLocal, onceHsize, dstWSize_, 0, outOffset);
553+ }
554+ outQue_.FreeTensor(yLocal);
555+ }
556+ }
557+}
558+ 
559+template <typename T, typename T1, int schId, bool alignCorners>
560+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::ComputeHWids(LocalTensor<T1> &idxUb,
561+ LocalTensor<T1> &idxHUb, LocalTensor<T1> &idxWUb, int64_t dstWSizeAlgin)
562+{
563+ T1 hSize = dstHSize_;
564+ uint32_t vfLen = vlLen_;
565+ T1 wTimes = dstWSizeAlgin / vlLen_;
566+ uint32_t tail = dstWSizeAlgin % vlLen_;
567+ uint16_t tailTimes = tail > 0 ? 1 : 0;
568+ auto idxUbAddr = (__ubuf__ T1 *)idxUb.GetPhyAddr();
569+ auto idxUbAddr1 = (__ubuf__ T1 *)idxUb[wTimes * vfLen].GetPhyAddr();
570+ 
571+ auto idxWUbAddr = (__ubuf__ T1 *)idxWUb.GetPhyAddr();
572+ auto idxWUbAddr1 = (__ubuf__ T1 *)idxWUb[wTimes * vfLen].GetPhyAddr();
573+ uint32_t wAlign = dstWSizeAlgin;
574+ __VEC_SCOPE__
575+ {
576+ AscendC::MicroAPI::RegTensor<T1> wIdxReg;
577+ AscendC::MicroAPI::RegTensor<T1> idxReg;
578+ 
579+ AscendC::MicroAPI::RegTensor<int32_t> hIdxInt32;
580+ AscendC::MicroAPI::RegTensor<int32_t> hIdxInt32C;
581+ AscendC::MicroAPI::MaskReg pregB32 = AscendC::MicroAPI::CreateMask<T1, AscendC::MicroAPI::MaskPattern::ALL>();
582+ AscendC::MicroAPI::MaskReg pregTail = AscendC::MicroAPI::UpdateMask<T1>(tail);
583+ for (uint16_t i = 0; i < static_cast<uint16_t>(hSize); i++) {
584+ T1 hIdx = idxHUb.GetValue(i);
585+ AscendC::MicroAPI::AddrReg outIdxOffset = AscendC::MicroAPI::CreateAddrReg<T1>(i, wAlign);
586+ for (uint16_t j = 0; j < static_cast<uint16_t>(wTimes); j++) {
587+ AscendC::MicroAPI::AddrReg srcIdxOffset = AscendC::MicroAPI::CreateAddrReg<T1>(j, vfLen);
588+ DataCopy(wIdxReg, idxWUbAddr, srcIdxOffset);
589+ Adds(idxReg, wIdxReg, hIdx, pregB32);
590+ AscendC::MicroAPI::AddrReg srcOutOffset = AscendC::MicroAPI::CreateAddrReg<T1>(i, wAlign, j, vfLen);
591+ DataCopy(idxUbAddr, idxReg, srcOutOffset, pregB32);
592+ }
593+ for (uint16_t jj = 0; jj < tailTimes; jj++) {
594+ DataCopy(wIdxReg, idxWUbAddr1);
595+ Adds(idxReg, wIdxReg, hIdx, pregTail);
596+ DataCopy(idxUbAddr1, idxReg, outIdxOffset, pregTail);
597+ }
598+ }
599+ }
600+}
601+ 
602+template <typename T, typename T1, int schId, bool alignCorners>
603+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::ComputeIdsSpecial(int64_t dstWSizeAlgin,
604+ int64_t srcWSizeAlgin)
605+{
606+ LocalTensor<T1> idxUb = idxBuf_.Get<T1>();
607+ LocalTensor<T1> idxHUb = idxHBuf_.Get<T1>();
608+ LocalTensor<T1> idxWUb = idxWBuf_.Get<T1>();
609+ ComputeHOrWids<T1, true, alignCorners>(idxHUb, bias_, hScale_, srcHSize_, dstHSize_, srcWSizeAlgin);
610+ ComputeHOrWids<T1, false, alignCorners>(idxWUb, bias_, wScale_, srcWSize_, dstWSizeAlgin, srcWSizeAlgin);
611+ ComputeHWids(idxUb, idxHUb, idxWUb, dstWSizeAlgin); // h和w坐标合成1维坐标
612+}
613+ 
614+template <typename T, typename T1, int schId, bool alignCorners>
615+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::ComputeAllHw()
616+{
617+ ComputeIdsSpecial(dstWAlignSize_, srcWAlignSize_); // h*w一次可以放下，对应的h*w坐标
618+ int64_t ncLoopTimes = blockIdx_ == realCoreNum_ - 1 ? tailTimes_ : beforeTimes_;
619+ int64_t ncLoopTailNum = blockIdx_ == realCoreNum_ - 1 ? tailTNum_ : beforeTNum_;
620+ int64_t hiStart = ComputeOriIds(0, srcHSize_ - 1, hScale_, bias_);
621+ int64_t hiEnd = ComputeOriIds(dstHSize_ - 1, srcHSize_ - 1, hScale_, bias_);
622+ int64_t srcHwOffsetStart = hiStart * srcWAlignSize_;
623+ int64_t hiSize = hiEnd - hiStart + 1;
624+ int64_t blockLen = hiSize * srcWAlignSize_;
625+ int64_t dstLen = dstHSize_ * dstWAlignSize_;
626+ for (int64_t i = 0; i < ncLoopTimes; i++) {
627+ int64_t onceNum = i == ncLoopTimes - 1 ? ncLoopTailNum : ubFactor_;
628+ int64_t srcOffset = (blockIdx_ * blockFactor_ + i * ubFactor_) * srcHwNum_ + hiStart * srcWSize_;
629+ LocalTensor<T> xLocal = inQue_.AllocTensor<T>();
630+ if (srcWAlignSize_ == srcWSize_) { // srcW对齐
631+ DataCopyIn(xLocal, onceNum, blockLen, srcHwNum_ - blockLen, srcOffset);
632+ } else {
633+ // srcW不对齐，需要用硬件4层for循环能力
634+ LoopModeParams loopParams;
635+ copyParams_.blockCount = hiSize;
636+ copyParams_.blockLen = srcWSize_ * sizeof(T);
637+ copyParams_.srcStride = 0;
638+ copyParams_.dstStride = 0;
639+ loopParams.loop2Size = 1;
640+ loopParams.loop1Size = onceNum;
641+ loopParams.loop2SrcStride = 0;
642+ loopParams.loop2DstStride = 0;
643+ loopParams.loop1SrcStride = srcHwNum_ * sizeof(T);
644+ loopParams.loop1DstStride = srcHSize_ * srcWAlignSize_ * sizeof(T);
645+ SetLoopModePara(loopParams, DataCopyMVType::OUT_TO_UB);
646+ DataCopyPad(xLocal, xGm_[srcOffset], copyParams_, padParams_);
647+ ResetLoopModePara(DataCopyMVType::OUT_TO_UB);
648+ }
649+ 
650+ inQue_.EnQue<T>(xLocal);
651+ xLocal = inQue_.DeQue<T>();
652+ LocalTensor<T> yLocal = outQue_.AllocTensor<T>();
653+ // 搬入后根据ids做gather选出输出数据
654+ GatherOutput(yLocal, xLocal, onceNum, blockLen, dstLen, srcHwOffsetStart);
655+ inQue_.FreeTensor(xLocal);
656+ outQue_.EnQue<T>(yLocal);
657+ yLocal = outQue_.DeQue<T>();
658+ int64_t outOffset = (blockIdx_ * blockFactor_ + i * ubFactor_) * dstHwNum_;
659+ if (dstWAlignSize_ == dstWSize_) { // dstW对齐
660+ DataCopyOut(yLocal, onceNum, dstHwNum_, 0, outOffset);
661+ } else {
662+ LoopModeParams loopParams;
663+ copyParams_.blockCount = dstHSize_;
664+ copyParams_.blockLen = dstWSize_ * sizeof(T);
665+ copyParams_.srcStride = 0;
666+ copyParams_.dstStride = 0;
667+ loopParams.loop2Size = 1;
668+ loopParams.loop1Size = onceNum;
669+ loopParams.loop2SrcStride = 0;
670+ loopParams.loop2DstStride = 0;
671+ loopParams.loop1SrcStride = dstHSize_ * dstWAlignSize_ * sizeof(T);
672+ loopParams.loop1DstStride = dstHwNum_ * sizeof(T);
673+ SetLoopModePara(loopParams, DataCopyMVType::UB_TO_OUT);
674+ DataCopyPad(yGm_[outOffset], yLocal, copyParams_);
675+ ResetLoopModePara(DataCopyMVType::UB_TO_OUT);
676+ }
677+ outQue_.FreeTensor(yLocal);
678+ }
679+}
680+ 
681+template <typename T, typename T1, int schId, bool alignCorners>
682+__aicore__ inline void ResizeGather<T, T1, schId, alignCorners>::Process()
683+{
684+ if (blockIdx_ >= realCoreNum_) {
685+ return;
686+ }
687+ 
688+ if constexpr (schId == TPL_SCH_MODE_GATHER_ALL_HW) {
689+ ComputeAllHw();
690+ }
691+ if constexpr (schId == TPL_SCH_MODE_GATHER_CUT_H) {
692+ ComputeCutH();
693+ }
694+}
695+} // namespace ResizeNearestNeighborV2
696+#endif // RESIZE_NEAREAST_NEIGHBOR_V2_NCHW_GATHER_H
Mimage/resize_nearest_neighbor_v2/op_kernel/arch35/resize_nearest_neighbor_v2_tiling_key.h+16-2
@@ -26,6 +26,8 @@
26#define TPL_SCH_MODE_DATA_COPY_ALL_W_OUT 626#define TPL_SCH_MODE_DATA_COPY_ALL_W_OUT 6
27#define TPL_SCH_MODE_DATA_COPY_NOT_ALL_W_OUT 727#define TPL_SCH_MODE_DATA_COPY_NOT_ALL_W_OUT 7
28#define TPL_SCH_MODE_DATA_COPY_CUT_NH 828#define TPL_SCH_MODE_DATA_COPY_CUT_NH 8
29+#define TPL_SCH_MODE_GATHER_ALL_HW 9
30+#define TPL_SCH_MODE_GATHER_CUT_H 10
29 31 
30#define TPL_FORMAT_NCHW 032#define TPL_FORMAT_NCHW 0
31#define TPL_FORMAT_NHWC 133#define TPL_FORMAT_NHWC 1
@@ -52,7 +54,9 @@
52 TPL_SCH_MODE_SIMT_INPUT_EQ_ONE, \54 TPL_SCH_MODE_SIMT_INPUT_EQ_ONE, \
53 TPL_SCH_MODE_DATA_COPY_ALL_W_OUT, \55 TPL_SCH_MODE_DATA_COPY_ALL_W_OUT, \
54 TPL_SCH_MODE_DATA_COPY_NOT_ALL_W_OUT, \56 TPL_SCH_MODE_DATA_COPY_NOT_ALL_W_OUT, \
55- TPL_SCH_MODE_DATA_COPY_CUT_NH), \57+ TPL_SCH_MODE_DATA_COPY_CUT_NH, \
58+ TPL_SCH_MODE_GATHER_ALL_HW, \
59+ TPL_SCH_MODE_GATHER_CUT_H), \
56 ASCENDC_TPL_UINT_DECL(format, ASCENDC_TPL_2_BW, ASCENDC_TPL_UI_LIST, TPL_FORMAT_NCHW, \60 ASCENDC_TPL_UINT_DECL(format, ASCENDC_TPL_2_BW, ASCENDC_TPL_UI_LIST, TPL_FORMAT_NCHW, \
57 TPL_FORMAT_NHWC, TPL_FORMAT_ND), \61 TPL_FORMAT_NHWC, TPL_FORMAT_ND), \
58 ASCENDC_TPL_UINT_DECL( \62 ASCENDC_TPL_UINT_DECL( \
@@ -91,8 +95,18 @@
91 halfPixelCenters, ASCENDC_TPL_UI_LIST, TPL_HALF_PIXEL_CENTERS_0), \95 halfPixelCenters, ASCENDC_TPL_UI_LIST, TPL_HALF_PIXEL_CENTERS_0), \
92 ASCENDC_TPL_UINT_SEL(idxInt32, ASCENDC_TPL_UI_LIST, TPL_IDX_INT32_1)96 ASCENDC_TPL_UINT_SEL(idxInt32, ASCENDC_TPL_UI_LIST, TPL_IDX_INT32_1)
93 97 
98+#define RESIZE_NEAREST_NEIGHBOR_V2_GATHER_TPL_KEY_SEL() \
99+ ASCENDC_TPL_UINT_SEL(schId, ASCENDC_TPL_UI_LIST, TPL_SCH_MODE_GATHER_ALL_HW, \
100+ TPL_SCH_MODE_GATHER_CUT_H), \
101+ ASCENDC_TPL_UINT_SEL(format, ASCENDC_TPL_UI_LIST, TPL_FORMAT_NCHW), \
102+ ASCENDC_TPL_UINT_SEL(alignCorners, ASCENDC_TPL_UI_LIST, TPL_ALIGN_CORNERS_0, TPL_ALIGN_CORNERS_1), \
103+ ASCENDC_TPL_UINT_SEL( \
104+ halfPixelCenters, ASCENDC_TPL_UI_LIST, TPL_HALF_PIXEL_CENTERS_0), \
105+ ASCENDC_TPL_UINT_SEL(idxInt32, ASCENDC_TPL_UI_LIST, TPL_IDX_INT32_1)
106+ 
94ASCENDC_TPL_ARGS_DECL(ResizeNearestNeighborV2, RESIZE_NEAREST_NEIGHBOR_V2_TPL_KEY_DECL());107ASCENDC_TPL_ARGS_DECL(ResizeNearestNeighborV2, RESIZE_NEAREST_NEIGHBOR_V2_TPL_KEY_DECL());
95ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(RESIZE_NEAREST_NEIGHBOR_V2_TPL_KEY_SEL()),108ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(RESIZE_NEAREST_NEIGHBOR_V2_TPL_KEY_SEL()),
96 ASCENDC_TPL_ARGS_SEL(RESIZE_NEAREST_NEIGHBOR_V2_SIMD_TPL_KEY_SEL()),109 ASCENDC_TPL_ARGS_SEL(RESIZE_NEAREST_NEIGHBOR_V2_SIMD_TPL_KEY_SEL()),
97- ASCENDC_TPL_ARGS_SEL(RESIZE_NEAREST_NEIGHBOR_V2_NHWC_TPL_KEY_SEL()));110+ ASCENDC_TPL_ARGS_SEL(RESIZE_NEAREST_NEIGHBOR_V2_NHWC_TPL_KEY_SEL()),
111+ ASCENDC_TPL_ARGS_SEL(RESIZE_NEAREST_NEIGHBOR_V2_GATHER_TPL_KEY_SEL()));
98#endif112#endif
Mimage/resize_nearest_neighbor_v2/op_kernel/resize_nearest_neighbor_v2_apt.cpp+13-0
@@ -18,6 +18,7 @@
18#include "./arch35/resize_nearest_neighbor_v2_data_copy_jh.h"18#include "./arch35/resize_nearest_neighbor_v2_data_copy_jh.h"
19#include "./arch35/resize_nearest_neighbor_v2_data_copy_small_c.h"19#include "./arch35/resize_nearest_neighbor_v2_data_copy_small_c.h"
20#include "./arch35/resize_nearest_neighbor_v2_data_copy_nhwc.h"20#include "./arch35/resize_nearest_neighbor_v2_data_copy_nhwc.h"
21+#include "./arch35/resize_nearest_neighbor_v2_nchw_gather.h"
21#include "./arch35/resize_nearest_neighbor_v2_tiling_key.h"22#include "./arch35/resize_nearest_neighbor_v2_tiling_key.h"
22 23 
23using namespace AscendC;24using namespace AscendC;
@@ -79,6 +80,18 @@ __global__ __aicore__ void resize_nearest_neighbor_v2(GM_ADDR x, GM_ADDR size, G
79 op.Process();80 op.Process();
80 return;81 return;
81 }82 }
83+ if constexpr (schId == TPL_SCH_MODE_GATHER_ALL_HW || schId == TPL_SCH_MODE_GATHER_CUT_H) {
84+ if constexpr(sizeof(DTYPE_X) == sizeof(float)) {
85+ ResizeNearestNeighborV2::ResizeGather<DTYPE_X, uint32_t ,schId, alignCorners> op;
86+ op.Init(x, size, y, &tilingData);
87+ op.Process();
88+ } else {
89+ ResizeNearestNeighborV2::ResizeGather<DTYPE_X, uint16_t ,schId, alignCorners> op;
90+ op.Init(x, size, y, &tilingData);
91+ op.Process();
92+ }
93+ return;
94+ }
82 if constexpr (idxInt32) {95 if constexpr (idxInt32) {
83 ResizeNearestNeighborV2::ResizeNearestNeighborV2Simt<DTYPE_X, uint32_t, format, schId, alignCorners,96 ResizeNearestNeighborV2::ResizeNearestNeighborV2Simt<DTYPE_X, uint32_t, format, schId, alignCorners,
84 halfPixelCenters> op;97 halfPixelCenters> op;
Aimage/resize_nearest_neighbor_v2/tests/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+ 
13+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aimage/resize_nearest_neighbor_v2/tests/ut/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+ 
13+foreach(SUB_DIR ${CURRENT_DIRS})
14+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
18+ 
Aimage/resize_nearest_neighbor_v2/tests/ut/op_host/CMakeLists.txt+24-0
@@ -0,0 +1,24 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_HOST_UT)
13+ add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
15+ 
16+if(NOT (UT_TEST_ALL OR OP_API_UT))
17+ list(REMOVE_ITEM CURRENT_DIRS op_api)
18+endif()
19+ 
20+foreach(SUB_DIR ${CURRENT_DIRS})
21+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
22+ add_subdirectory(${SUB_DIR})
23+ endif()
24+endforeach()
Aimage/resize_nearest_neighbor_v2/tests/ut/op_host/arch35/test_resize_nearest_neighbor_v2_tiling.cpp+78-0
@@ -0,0 +1,78 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <fstream>
13+#include <vector>
14+#include <gtest/gtest.h>
15+ 
16+#include "../../../../op_host/arch35/resize_nearest_neighbor_v2_tiling_base.h"
17+#include "tiling_context_faker.h"
18+#include "tiling_case_executor.h"
19+ 
20+using namespace ge;
21+using namespace optiling;
22+ 
23+class ResizeNearestNeighborV2TilingTest : public testing::Test {
24+ protected:
25+ static void SetUpTestCase() {
26+ std::cout << "ResizeNearestNeighborV2TilingTest SetUp" << std::endl;
27+ }
28+ 
29+ static void TearDownTestCase() {
30+ std::cout << "ResizeNearestNeighborV2TilingTest TearDown" << std::endl;
31+ }
32+};
33+ 
34+TEST_F(ResizeNearestNeighborV2TilingTest, resize_nearest_neighbor_v2_tiling_01)
35+{
36+ gert::StorageShape inputXShape = {{1, 64, 3, 32}, {1, 64, 3, 32}};
37+ gert::StorageShape inputSizeShape = {{2,}, {2,}};
38+ gert::StorageShape outputShape = {{1, 64, 6, 64}, {1, 64, 6, 64}};
39+ int size_value[2] = {6, 64};
40+ 
41+ ResizeNearestNeighborV2CompileInfo compileInfo = {64, 253952};
42+ 
43+ gert::TilingContextPara tilingContextPara(
44+ "ResizeNearestNeighborV2",
45+ {{inputXShape, ge::DT_FLOAT, ge::FORMAT_NCHW}, {inputSizeShape, ge::DT_INT32, ge::FORMAT_ND, true, size_value}},
46+ {{outputShape, ge::DT_FLOAT, ge::FORMAT_NCHW}},
47+ {gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
48+ gert::TilingContextPara::OpAttr("half_pixel_centers", Ops::Cv::AnyValue::CreateFrom<bool>(false))
49+ gert::TilingContextPara::OpAttr("scales", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0, 0.0}))},
50+ &compileInfo);
51+ uint64_t expectTilingKey = 4105;
52+ string expectTilingData = "64 1536 0 0 1 64 3 32 6 64 0 0 1 0 64 6144 24576 2048 4096 1 1 1 1 0 0 0 0 0 0 65 0 0 99840 24960 0 4539628425446424576 ";
53+ std::vector<size_t> expectWorkspaces = {16777216};
54+ 
55+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
56+}
57+TEST_F(ResizeNearestNeighborV2TilingTest, resize_nearest_neighbor_v2_tiling_02)
58+{
59+ gert::StorageShape inputXShape = {{1, 64, 64, 144}, {1, 64, 64, 144}};
60+ gert::StorageShape inputSizeShape = {{2,}, {2,}};
61+ gert::StorageShape outputShape = {{1, 64, 64, 128}, {1, 64, 128, 288}};
62+ int size_value[2] = {128, 288};
63+ 
64+ ResizeNearestNeighborV2CompileInfo compileInfo = {64, 253952};
65+ 
66+ gert::TilingContextPara tilingContextPara(
67+ "ResizeNearestNeighborV2",
68+ {{inputXShape, ge::DT_FLOAT, ge::FORMAT_NCHW}, {inputSizeShape, ge::DT_INT32, ge::FORMAT_ND, true, size_value}},
69+ {{outputShape, ge::DT_FLOAT, ge::FORMAT_NCHW}},
70+ {gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
71+ gert::TilingContextPara::OpAttr("half_pixel_centers", Ops::Cv::AnyValue::CreateFrom<bool>(false))
72+ gert::TilingContextPara::OpAttr("scales", Ops::Cv::AnyValue::CreateFrom<std::vector<float>>({0.0, 0.0}))},
73+ &compileInfo);
74+ uint64_t expectTilingKey = 4106;
75+ string expectTilingData = "64 71424 0 0 1 64 64 144 128 288 0 0 1 0 64 589824 2359296 9216 18432 1 0 0 4 0 0 0 0 0 0 62 0 3 71424 19008 0 4539628425446424576 ";
76+ std::vector<size_t> expectWorkspaces = {16777216};
77+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
78+}
Aimage/resize_nearest_neighbor_v2/tests/ut/op_host/test_resize_nearest_neighbor_v2_infershape.cpp+45-0
@@ -0,0 +1,45 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+#include <iostream>
13+#include <numeric>
14+#include "infershape_context_faker.h"
15+#include "infershape_case_executor.h"
16+#include "base/registry/op_impl_space_registry_v2.h"
17+ 
18+class ResizeNearestNeighborV2InfershapeTest : public testing::Test {
19+protected:
20+ static void SetUpTestCase()
21+ {
22+ std::cout << "ResizeNearestNeighborV2InfershapeTest SetUp" << std::endl;
23+ }
24+ 
25+ static void TearDownTestCase()
26+ {
27+ std::cout << "ResizeNearestNeighborV2InfershapeTest TearDown" << std::endl;
28+ }
29+};
30+ 
31+TEST_F(ResizeNearestNeighborV2InfershapeTest, resize_nearest_neighbor_v2_infershape_test_01)
32+{
33+ gert::StorageShape inputXShape = {{1, 2, 3, 32}, {1, 2, 3, 32}};
34+ gert::StorageShape inputSizeShape = {{2,}, {2,}};
35+ gert::StorageShape outputShape = {{1, 2, 6, 64}, {1, 2, 6, 64}};
36+ int size_value[2] = {6, 64};
37+ 
38+ gert::InfershapeContextPara infershapeContextPara(
39+ "ResizeNearestNeighborV2",
40+ {{inputXShape, ge::DT_FLOAT, ge::FORMAT_NCHW}, {inputSizeShape, ge::DT_INT32, ge::FORMAT_ND, true, size_value}},
41+ {{outputShape, ge::DT_FLOAT, ge::FORMAT_NCHW}});
42+ 
43+ std::vector<std::vector<int64_t>> expectOutputShape = {{1, 2, 6, 64},};
44+ ExecuteTestCase(infershapeContextPara, ge::GRAPH_SUCCESS, expectOutputShape);
45+}