Revert upsample_nearest2d_grad和upsample_nearest_exact2d

Dimage/upsample_nearest2d_grad/op_host/config/ascend950/upsample_nearest2d_grad_binary.json+0-140

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

Dimage/upsample_nearest2d_grad/op_host/config/ascend950/upsample_nearest2d_grad_simplified_key.ini+0-13

	@@ -1,13 +0,0 @@
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	-[UpsampleNearest2dGrad]
13	-default=0

Mimage/upsample_nearest2d_grad/op_host/upsample_nearest2d_grad_def.cpp+0-10

	@@ -39,16 +39,6 @@ public:
39		39
40	this->AICore().AddConfig("ascend910b");	40	this->AICore().AddConfig("ascend910b");
41	this->AICore().AddConfig("ascend910_93");	41	this->AICore().AddConfig("ascend910_93");
42	-
43	- OpAICoreConfig aicoreConfig;
44	- aicoreConfig.DynamicCompileStaticFlag(true)
45	- .DynamicFormatFlag(false)
46	- .DynamicRankSupportFlag(true)
47	- .DynamicShapeSupportFlag(true)
48	- .NeedCheckSupportFlag(false)
49	- .PrecisionReduceFlag(true)
50	- .ExtendCfgInfo("opFile.value", "upsample_nearest2d_grad_apt");
51	- this->AICore().AddConfig("ascend950", aicoreConfig);
52	}	42	}
53	};	43	};
54		44

Dimage/upsample_nearest2d_grad/op_host/upsample_nearest2d_grad_tiling_arch35.cpp+0-205

	@@ -1,205 +0,0 @@
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 upsample_nearest2d_grad_tiling.cpp
13	- * \brief Tiling implementation for upsample_nearest2d_grad operator
14	- */
15	-
16	-#include "log/log.h"
17	-#include "platform/platform_ascendc.h"
18	-#include "util/math_util.h"
19	-#include "util/platform_util.h"
20	-#include "op_host/tiling_util.h"
21	-#include "op_host/tiling_templates_registry.h"
22	-#include "../op_kernel/arch35/upsample_nearest2d_grad_tiling_data.h"
23	-#include "../op_kernel/arch35/upsample_nearest2d_grad_tiling_key.h"
24	-
25	-namespace optiling {
26	-
27	-constexpr int64_t PER_CORE_MIN_ELEMENTS = 1024;
28	-constexpr uint32_t DCACHE_SIZE = 128 * 1024;
29	-constexpr uint32_t STATIC_UB_ESTIMATE = 0;
30	-constexpr int32_t INPUT_SIZE_ATTR_LEN = 4;
31	-constexpr int32_t OUTPUT_SIZE_ATTR_LEN = 2;
32	-
33	-struct UpsampleNearest2dGradCompileInfo {};
34	-
35	-static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
36	-{
37	- fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
38	- OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
39	- auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
40	- coreNum = ascendcPlatform.GetCoreNumAiv();
41	- OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
42	- ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
43	- OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
44	- return ge::GRAPH_SUCCESS;
45	-}
46	-
47	-static ge::graphStatus GetTilingAttrs(gert::TilingContext* context, int32_t& dimN, int32_t& dimC, int32_t& dimHin,
48	- int32_t& dimWin, int32_t& dimHout, int32_t& dimWout, float& scaleH, float& scaleW)
49	-{
50	- auto attrs = context->GetAttrs();
51	- OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
52	-
53	- // Attr 0: output_size (ListInt) [H_out, W_out]
54	- const auto* outputSizeVec = attrs->GetListInt(0);
55	- OP_CHECK_NULL_WITH_CONTEXT(context, outputSizeVec);
56	- if (outputSizeVec->GetSize() != OUTPUT_SIZE_ATTR_LEN) {
57	- OP_LOGE_FOR_INVALID_LISTSIZE(context->GetNodeName(), "output_size",
58	- std::to_string(outputSizeVec->GetSize()).c_str(),
59	- std::to_string(OUTPUT_SIZE_ATTR_LEN).c_str());
60	- return ge::GRAPH_FAILED;
61	- }
62	- const int64_t* outputSizeData = outputSizeVec->GetData();
63	- dimHout = static_cast<int32_t>(outputSizeData[0]);
64	- dimWout = static_cast<int32_t>(outputSizeData[1]);
65	-
66	- // Attr 1: input_size (ListInt) [N, C, H_in, W_in]
67	- const auto* inputSizeVec = attrs->GetListInt(1);
68	- OP_CHECK_NULL_WITH_CONTEXT(context, inputSizeVec);
69	- if (inputSizeVec->GetSize() != INPUT_SIZE_ATTR_LEN) {
70	- OP_LOGE_FOR_INVALID_LISTSIZE(context->GetNodeName(), "input_size",
71	- std::to_string(inputSizeVec->GetSize()).c_str(),
72	- std::to_string(INPUT_SIZE_ATTR_LEN).c_str());
73	- return ge::GRAPH_FAILED;
74	- }
75	- const int64_t* inputSizeData = inputSizeVec->GetData();
76	- dimN = static_cast<int32_t>(inputSizeData[0]);
77	- dimC = static_cast<int32_t>(inputSizeData[1]);
78	- dimHin = static_cast<int32_t>(inputSizeData[2]);
79	- dimWin = static_cast<int32_t>(inputSizeData[3]);
80	-
81	- if (dimN <= 0) {
82	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "input_size", std::to_string(dimN).c_str(),
83	- "dimN must be greater than 0");
84	- return ge::GRAPH_FAILED;
85	- }
86	- if (dimC <= 0) {
87	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "input_size", std::to_string(dimC).c_str(),
88	- "dimC must be greater than 0");
89	- return ge::GRAPH_FAILED;
90	- }
91	- if (dimHin <= 0) {
92	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "input_size", std::to_string(dimHin).c_str(),
93	- "dimHin must be greater than 0");
94	- return ge::GRAPH_FAILED;
95	- }
96	- if (dimWin <= 0) {
97	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "input_size", std::to_string(dimWin).c_str(),
98	- "dimWin must be greater than 0");
99	- return ge::GRAPH_FAILED;
100	- }
101	- if (dimHout <= 0) {
102	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "output_size", std::to_string(dimHout).c_str(),
103	- "dimHout must be greater than 0");
104	- return ge::GRAPH_FAILED;
105	- }
106	- if (dimWout <= 0) {
107	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context->GetNodeName(), "output_size", std::to_string(dimWout).c_str(),
108	- "dimWout must be greater than 0");
109	- return ge::GRAPH_FAILED;
110	- }
111	-
112	- // Attr 2: scales_h (Float, default 0.0)
113	- float scalesHVal = 0.0f;
114	- const float* scalesHPtr = attrs->GetFloat(2);
115	- if (scalesHPtr != nullptr) {
116	- scalesHVal = *scalesHPtr;
117	- }
118	-
119	- // Attr 3: scales_w (Float, default 0.0)
120	- float scalesWVal = 0.0f;
121	- const float* scalesWPtr = attrs->GetFloat(3);
122	- if (scalesWPtr != nullptr) {
123	- scalesWVal = *scalesWPtr;
124	- }
125	-
126	- // Compute scale
127	- scaleH = (scalesHVal > 0.0f) ? scalesHVal : static_cast<float>(dimHout) / static_cast<float>(dimHin);
128	- scaleW = (scalesWVal > 0.0f) ? scalesWVal : static_cast<float>(dimWout) / static_cast<float>(dimWin);
129	-
130	- return ge::GRAPH_SUCCESS;
131	-}
132	-
133	-ge::graphStatus UpsampleNearest2dGradArch35TilingFunc(gert::TilingContext* context)
134	-{
135	- uint64_t ubSize = 0;
136	- int64_t coreNum = 0;
137	- OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
138	- OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
139	-
140	- int32_t dimN = 0, dimC = 0, dimHin = 0, dimWin = 0, dimHout = 0, dimWout = 0;
141	- float scaleH = 0.0f, scaleW = 0.0f;
142	- OP_CHECK_IF(
143	- GetTilingAttrs(context, dimN, dimC, dimHin, dimWin, dimHout, dimWout, scaleH, scaleW) != ge::GRAPH_SUCCESS,
144	- OP_LOGE(context, "GetTilingAttrs error"), return ge::GRAPH_FAILED);
145	-
146	- // Compute tiling parameters
147	- int64_t totalElements = (int64_t)dimC * dimHin * dimWin;
148	-
149	- int64_t perCoreElements = (totalElements > 0) ? Ops::Base::CeilDiv(totalElements, coreNum) : 0;
150	- if (perCoreElements > 0 && perCoreElements < PER_CORE_MIN_ELEMENTS) {
151	- perCoreElements = PER_CORE_MIN_ELEMENTS;
152	- }
153	- // Align to 32 (warp size)
154	- if (perCoreElements > 0) {
155	- perCoreElements = ((perCoreElements + 31) / 32) * 32;
156	- }
157	- int32_t needCoreNum = (totalElements > 0) ?
158	- static_cast<int32_t>(Ops::Base::CeilDiv(totalElements, perCoreElements)) :
159	- 1;
160	- if (needCoreNum < 1) {
161	- needCoreNum = 1;
162	- }
163	-
164	- // Fill tiling data
165	- UpsampleNearest2dGradTilingData* tiling = context->GetTilingData<UpsampleNearest2dGradTilingData>();
166	- OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
167	- OP_CHECK_IF(
168	- memset_s(tiling, sizeof(UpsampleNearest2dGradTilingData), 0, sizeof(UpsampleNearest2dGradTilingData)) != EOK,
169	- OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
170	-
171	- tiling->needCoreNum = needCoreNum;
172	- tiling->totalElements = totalElements;
173	- tiling->dimN = dimN;
174	- tiling->dimC = dimC;
175	- tiling->dimHin = dimHin;
176	- tiling->dimWin = dimWin;
177	- tiling->dimHout = dimHout;
178	- tiling->dimWout = dimWout;
179	- tiling->scaleH = scaleH;
180	- tiling->scaleW = scaleW;
181	-
182	- context->SetBlockDim(needCoreNum);
183	-
184	- // Set tiling key (single mode, dtype by DTYPE_ macro)
185	- uint64_t tilingKey = GET_TPL_TILING_KEY(UPSAMPLE_NEAREST2D_GRAD_MODE_DEFAULT);
186	- context->SetTilingKey(tilingKey);
187	-
188	- // Set workspace (system workspace only)
189	- auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
190	- uint64_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
191	- size_t* currentWorkspace = context->GetWorkspaceSizes(1);
192	- OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
193	- currentWorkspace[0] = static_cast<size_t>(sysWorkspaceSize);
194	-
195	- // Set local memory size
196	- OP_CHECK_IF((ubSize <= DCACHE_SIZE + STATIC_UB_ESTIMATE),
197	- OP_LOGE(context, "ubSize %lu <= DCACHE_SIZE + STATIC_UB_ESTIMATE", ubSize), return ge::GRAPH_FAILED);
198	- auto res = context->SetLocalMemorySize(static_cast<uint32_t>(ubSize - DCACHE_SIZE - STATIC_UB_ESTIMATE));
199	- OP_CHECK_IF((res != ge::GRAPH_SUCCESS), OP_LOGE(context, "SetLocalMemorySize failed, ubSize=%lu", ubSize),
200	- return ge::GRAPH_FAILED);
201	-
202	- return ge::GRAPH_SUCCESS;
203	-}
204	-
205	-} // namespace optiling

Dimage/upsample_nearest2d_grad/op_kernel/arch35/upsample_nearest2d_grad_simt.h+0-187

	@@ -1,187 +0,0 @@
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 upsample_nearest2d_grad_simt.h
13	- * \brief SIMT kernel implementation for upsample_nearest2d_grad operator
14	- *
15	- * Performance optimizations applied:
16	- * R002: UintDiv fast division for coordinate decomposition (3 divisions → 2 UintDiv + multiply-subtract)
17	- * R003: 32/64-bit index template (INDEX_SIZE_T) for grid-stride loop
18	- * R006: Thread count tuned by index width (1024 for uint32_t, 512 for uint64_t)
19	- */
20	-
21	-#ifndef UPSAMPLE_NEAREST2D_GRAD_SIMT_H_
22	-#define UPSAMPLE_NEAREST2D_GRAD_SIMT_H_
23	-
24	-#include "kernel_operator.h"
25	-#include "kernel_tiling/kernel_tiling.h"
26	-#include "simt_api/asc_simt.h"
27	-#include "simt_api/common_functions.h"
28	-#include "simt_api/math_functions.h"
29	-#include "simt_api/asc_fp16.h"
30	-#include "simt_api/asc_bf16.h"
31	-#include "upsample_nearest2d_grad_tiling_data.h"
32	-#include "upsample_nearest2d_grad_tiling_key.h"
33	-
34	-namespace NsUpsampleNearest2dGrad {
35	-
36	-using namespace AscendC;
37	-
38	-// R006: Thread count tuned by index width
39	-// uint32_t path: 1024 threads (lower register pressure, higher occupancy)
40	-// uint64_t path: 512 threads (higher register pressure from 64-bit ops)
41	-template <typename INDEX_SIZE_T>
42	-static constexpr uint32_t THREAD_NUM_OPT = (sizeof(INDEX_SIZE_T) == 4) ? 1024 : 512;
43	-
44	-// === Helper functions: type-safe read/write ===
45	-
46	-template <typename T>
47	-__simt_callee__ __aicore__ inline float ReadAsFloat(__gm__ T* ptr, int64_t idx);
48	-
49	-template <>
50	-__simt_callee__ __aicore__ inline float ReadAsFloat<half>(__gm__ half* ptr, int64_t idx)
51	-{
52	- return __half2float(ptr[idx]);
53	-}
54	-
55	-template <>
56	-__simt_callee__ __aicore__ inline float ReadAsFloat<bfloat16_t>(__gm__ bfloat16_t* ptr, int64_t idx)
57	-{
58	- return __bfloat162float(ptr[idx]);
59	-}
60	-
61	-template <>
62	-__simt_callee__ __aicore__ inline float ReadAsFloat<float>(__gm__ float* ptr, int64_t idx)
63	-{
64	- return ptr[idx];
65	-}
66	-
67	-template <typename T>
68	-__simt_callee__ __aicore__ inline void WriteFromFloat(__gm__ T* ptr, int64_t idx, float val);
69	-
70	-template <>
71	-__simt_callee__ __aicore__ inline void WriteFromFloat<half>(__gm__ half* ptr, int64_t idx, float val)
72	-{
73	- ptr[idx] = __float2half(val);
74	-}
75	-
76	-template <>
77	-__simt_callee__ __aicore__ inline void WriteFromFloat<bfloat16_t>(__gm__ bfloat16_t* ptr, int64_t idx, float val)
78	-{
79	- ptr[idx] = __float2bfloat16(val);
80	-}
81	-
82	-template <>
83	-__simt_callee__ __aicore__ inline void WriteFromFloat<float>(__gm__ float* ptr, int64_t idx, float val)
84	-{
85	- ptr[idx] = val;
86	-}
87	-
88	-// === Inner accumulation loop (split to keep function body <= 50 lines) ===
89	-
90	-template <typename T>
91	-__simt_callee__ __aicore__ inline float AccumulateGrad(__gm__ T* gradOutput, int32_t b, int64_t srcBase,
92	- int64_t srcNCStride, int32_t srcYStart, int32_t srcYEnd,
93	- int32_t srcXStart, int32_t srcXEnd, int32_t dimWout)
94	-{
95	- float grad = 0.0f;
96	- for (int32_t y = srcYStart; y < srcYEnd; y++) {
97	- for (int32_t x = srcXStart; x < srcXEnd; x++) {
98	- int64_t srcIdx = (int64_t)b * srcNCStride + srcBase + (int64_t)y * dimWout + (int64_t)x;
99	- grad += ReadAsFloat<T>(gradOutput, srcIdx);
100	- }
101	- }
102	- return grad;
103	-}
104	-
105	-// === SIMT VF Kernel (R002: UintDiv, R003: INDEX_SIZE_T template, R006: thread tuning) ===
106	-
107	-template <typename T, typename INDEX_SIZE_T>
108	-__simt_vf__ __aicore__ __launch_bounds__(THREAD_NUM_OPT<INDEX_SIZE_T>) inline void OpUpsampleNearest2dGradSimtKernel(
109	- INDEX_SIZE_T totalElements, int32_t dimN, int32_t dimC, int32_t dimHin, int32_t dimWin, int32_t dimHout,
110	- int32_t dimWout, float scaleH, float scaleW, INDEX_SIZE_T magic0, INDEX_SIZE_T shift0, INDEX_SIZE_T magic1,
111	- INDEX_SIZE_T shift1, __gm__ T* gradOutput, __gm__ T* gradInput)
112	-{
113	- const INDEX_SIZE_T dstCStride = (INDEX_SIZE_T)dimHin * (INDEX_SIZE_T)dimWin;
114	- const int32_t srcCStride = dimHout * dimWout;
115	- const int64_t dstNCStride = (int64_t)dimC * (int64_t)dstCStride;
116	- const int64_t srcNCStride = (int64_t)dimC * (int64_t)srcCStride;
117	-
118	- for (INDEX_SIZE_T index = static_cast<INDEX_SIZE_T>(blockIdx.x * blockDim.x + threadIdx.x); index < totalElements;
119	- index += static_cast<INDEX_SIZE_T>(blockDim.x * gridDim.x)) {
120	- // R002: UintDiv replaces 3 hardware divisions with 2 fast-div + multiply-subtract
121	- INDEX_SIZE_T cIdxFull = Simt::UintDiv<INDEX_SIZE_T>(index, magic0, shift0);
122	- int32_t cIdx = static_cast<int32_t>(cIdxFull % static_cast<INDEX_SIZE_T>(dimC));
123	- INDEX_SIZE_T rem0 = index - cIdxFull * dstCStride;
124	- INDEX_SIZE_T hInIdx = Simt::UintDiv<INDEX_SIZE_T>(rem0, magic1, shift1);
125	- int32_t hIn = static_cast<int32_t>(hInIdx);
126	- int32_t wIn = static_cast<int32_t>(rem0 - hInIdx * (INDEX_SIZE_T)dimWin);
127	-
128	- int32_t srcYStart = static_cast<int32_t>(min(
129	- static_cast<long long int>(ceilf(static_cast<float>(hIn) * scaleH)), static_cast<long long int>(dimHout)));
130	- int32_t srcYEnd = static_cast<int32_t>(
131	- min(static_cast<long long int>(ceilf(static_cast<float>(hIn + 1) * scaleH)),
132	- static_cast<long long int>(dimHout)));
133	- int32_t srcXStart = static_cast<int32_t>(min(
134	- static_cast<long long int>(ceilf(static_cast<float>(wIn) * scaleW)), static_cast<long long int>(dimWout)));
135	- int32_t srcXEnd = static_cast<int32_t>(
136	- min(static_cast<long long int>(ceilf(static_cast<float>(wIn + 1) * scaleW)),
137	- static_cast<long long int>(dimWout)));
138	-
139	- int64_t dstIdx = static_cast<int64_t>(index);
140	- int64_t srcBase = (int64_t)cIdx * srcCStride;
141	-
142	- for (int32_t b = 0; b < dimN; b++) {
143	- float grad = AccumulateGrad<T>(gradOutput, b, srcBase, srcNCStride, srcYStart, srcYEnd, srcXStart, srcXEnd,
144	- dimWout);
145	- WriteFromFloat<T>(gradInput, dstIdx, grad);
146	- dstIdx += dstNCStride;
147	- }
148	- }
149	-}
150	-
151	-// === Process entry (R002: precompute magic/shift, R003: dispatch by index width) ===
152	-
153	-template <typename T>
154	-__aicore__ inline void Process(GM_ADDR gradOutput, GM_ADDR gradInput, const UpsampleNearest2dGradTilingData* tilingData)
155	-{
156	- __gm__ T* gradOutputGm = (__gm__ T*)gradOutput;
157	- __gm__ T* gradInputGm = (__gm__ T*)gradInput;
158	-
159	- const int64_t totalElements = tilingData->totalElements;
160	- const bool use64Bit = (totalElements > static_cast<int64_t>(INT32_MAX));
161	-
162	- if (use64Bit) {
163	- using IDX_T = uint64_t;
164	- IDX_T magic0 = 0, shift0 = 0, magic1 = 0, shift1 = 0;
165	- IDX_T dstCStride = (IDX_T)tilingData->dimHin * (IDX_T)tilingData->dimWin;
166	- GetUintDivMagicAndShift<IDX_T>(magic0, shift0, dstCStride);
167	- GetUintDivMagicAndShift<IDX_T>(magic1, shift1, (IDX_T)tilingData->dimWin);
168	- asc_vf_call<OpUpsampleNearest2dGradSimtKernel<T, uint64_t>>(
169	- dim3(THREAD_NUM_OPT<uint64_t>), static_cast<IDX_T>(totalElements), tilingData->dimN, tilingData->dimC,
170	- tilingData->dimHin, tilingData->dimWin, tilingData->dimHout, tilingData->dimWout, tilingData->scaleH,
171	- tilingData->scaleW, magic0, shift0, magic1, shift1, gradOutputGm, gradInputGm);
172	- } else {
173	- using IDX_T = uint32_t;
174	- IDX_T magic0 = 0, shift0 = 0, magic1 = 0, shift1 = 0;
175	- IDX_T dstCStride = (IDX_T)tilingData->dimHin * (IDX_T)tilingData->dimWin;
176	- GetUintDivMagicAndShift<IDX_T>(magic0, shift0, dstCStride);
177	- GetUintDivMagicAndShift<IDX_T>(magic1, shift1, (IDX_T)tilingData->dimWin);
178	- asc_vf_call<OpUpsampleNearest2dGradSimtKernel<T, uint32_t>>(
179	- dim3(THREAD_NUM_OPT<uint32_t>), static_cast<IDX_T>(totalElements), tilingData->dimN, tilingData->dimC,
180	- tilingData->dimHin, tilingData->dimWin, tilingData->dimHout, tilingData->dimWout, tilingData->scaleH,
181	- tilingData->scaleW, magic0, shift0, magic1, shift1, gradOutputGm, gradInputGm);
182	- }
183	-}
184	-
185	-} // namespace NsUpsampleNearest2dGrad
186	-
187	-#endif // UPSAMPLE_NEAREST2D_GRAD_SIMT_H_

Dimage/upsample_nearest2d_grad/op_kernel/arch35/upsample_nearest2d_grad_tiling_data.h+0-32

	@@ -1,32 +0,0 @@
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 upsample_nearest2d_grad_tiling_data.h
13	- * \brief Tiling data struct for upsample_nearest2d_grad operator
14	- */
15	-
16	-#ifndef UPSAMPLE_NEAREST2D_GRAD_TILING_DATA_H_
17	-#define UPSAMPLE_NEAREST2D_GRAD_TILING_DATA_H_
18	-
19	-struct UpsampleNearest2dGradTilingData {
20	- int32_t needCoreNum = 0;
21	- int64_t totalElements = 0;
22	- int32_t dimN = 0;
23	- int32_t dimC = 0;
24	- int32_t dimHin = 0;
25	- int32_t dimWin = 0;
26	- int32_t dimHout = 0;
27	- int32_t dimWout = 0;
28	- float scaleH = 0.0f;
29	- float scaleW = 0.0f;
30	-};
31	-
32	-#endif // UPSAMPLE_NEAREST2D_GRAD_TILING_DATA_H_

Dimage/upsample_nearest2d_grad/op_kernel/arch35/upsample_nearest2d_grad_tiling_key.h+0-35

	@@ -1,35 +0,0 @@
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 upsample_nearest2d_grad_tiling_key.h
13	- * \brief Tiling key declaration for upsample_nearest2d_grad operator
14	- *
15	- * Single template parameter:
16	- * schMode (UINT 1-bit): scene mode (only MODE_DEFAULT = 0)
17	- * dtype is handled by DTYPE_GRAD_OUTPUT macro automatically
18	- */
19	-
20	-#ifndef UPSAMPLE_NEAREST2D_GRAD_TILING_KEY_H_
21	-#define UPSAMPLE_NEAREST2D_GRAD_TILING_KEY_H_
22	-
23	-#include "ascendc/host_api/tiling/template_argument.h"
24	-
25	-#define UPSAMPLE_NEAREST2D_GRAD_MODE_DEFAULT 0
26	-
27	-ASCENDC_TPL_ARGS_DECL(UpsampleNearest2dGrad,
28	- ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, UPSAMPLE_NEAREST2D_GRAD_MODE_DEFAULT));
29	-
30	-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),
31	- ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST,
32	- UPSAMPLE_NEAREST2D_GRAD_MODE_DEFAULT),
33	- ASCENDC_TPL_TILING_STRUCT_SEL(UpsampleNearest2dGradTilingData)));
34	-
35	-#endif // UPSAMPLE_NEAREST2D_GRAD_TILING_KEY_H_

Dimage/upsample_nearest2d_grad/op_kernel/upsample_nearest2d_grad_apt.cpp+0-30

	@@ -1,30 +0,0 @@
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 upsample_nearest2d_grad_apt.cpp
13	- * \brief Kernel entry for upsample_nearest2d_grad operator
14	- *
15	- * Single template parameter:
16	- * schMode (uint32_t): scene mode (only MODE_DEFAULT = 0)
17	- * dtype is handled by DTYPE_GRAD_OUTPUT macro automatically
18	- */
19	-
20	-#include "./arch35/upsample_nearest2d_grad_simt.h"
21	-
22	-template <uint32_t schMode>
23	-__global__ __aicore__ void upsample_nearest2d_grad(GM_ADDR gradOutput, GM_ADDR gradInput, GM_ADDR workspace,
24	- GM_ADDR tiling)
25	-{
26	- REGISTER_TILING_DEFAULT(UpsampleNearest2dGradTilingData);
27	- GET_TILING_DATA_WITH_STRUCT(UpsampleNearest2dGradTilingData, tilingData, tiling);
28	-
29	- NsUpsampleNearest2dGrad::Process<DTYPE_GRAD_OUTPUT>(gradOutput, gradInput, &tilingData);
30	-}

Mimage/upsample_nearest2d_grad/tests/ut/op_host/CMakeLists.txt+0-1

	@@ -10,7 +10,6 @@
10		10
11	file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)	11	file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12	if(UT_TEST_ALL OR OP_HOST_UT)	12	if(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})
14	endif()	13	endif()
15		14
16	if(NOT (UT_TEST_ALL OR OP_API_UT))	15	if(NOT (UT_TEST_ALL OR OP_API_UT))

Dimage/upsample_nearest2d_grad/tests/ut/op_host/test_upsample_nearest2d_grad_tiling_arch35.cpp+0-154

	@@ -1,154 +0,0 @@
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_nearest2d_grad_tiling_data.h"
18	-
19	-using namespace std;
20	-using namespace ge;
21	-
22	-class UpsampleNearest2dGradTiling : public testing::Test {
23	-protected:
24	- static void SetUpTestCase() { std::cout << "UpsampleNearest2dGradTiling SetUp" << std::endl; }
25	-
26	- static void TearDownTestCase() { std::cout << "UpsampleNearest2dGradTiling TearDown" << std::endl; }
27	-};
28	-
29	-TEST_F(UpsampleNearest2dGradTiling, upsample_nearest2d_grad_fp32)
30	-{
31	- struct UpsampleNearest2dGradCompileInfo {
32	- } compileInfo;
33	- // grad_output: [1, 1, 4, 4] float32
34	- // grad_input: [1, 1, 2, 2] float32
35	- // attrs: output_size=[4,4], input_size=[1,1,2,2], scales_h=0.0, scales_w=0.0
36	- std::vector<gert::TilingContextPara::OpAttr> attrs = {
37	- {"output_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{4, 4})},
38	- {"input_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{1, 1, 2, 2})},
39	- {"scales_h", Ops::Cv::AnyValue::CreateFrom(0.0f)},
40	- {"scales_w", Ops::Cv::AnyValue::CreateFrom(0.0f)}};
41	- gert::TilingContextPara tilingContextPara(
42	- "UpsampleNearest2dGrad",
43	- {
44	- {{{1, 1, 4, 4}, {1, 1, 4, 4}}, ge::DT_FLOAT, ge::FORMAT_ND}, // grad_output
45	- },
46	- {
47	- {{{1, 1, 2, 2}, {1, 1, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}, // grad_input
48	- },
49	- attrs, &compileInfo, "Ascend950",
50	- 64, // number of cores
51	- 262144, // ubsize
52	- 4096); // max tiling data size
53	- uint64_t expectTilingKey = 0;
54	- // Tiling data is read as int64_t values. Struct layout:
55	- // [int32 needCoreNum \| pad] [int64 totalElements] [int32 dimN \| int32 dimC] [int32 dimHin \| int32 dimWin]
56	- // [int32 dimHout \| int32 dimWout] [float scaleH \| float scaleW]
57	- string expectTilingData = "1 4 4294967297 8589934594 17179869188 4611686019501129728 ";
58	- std::vector<size_t> expectWorkspaces = {16777216};
59	- ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
60	-}
61	-
62	-TEST_F(UpsampleNearest2dGradTiling, upsample_nearest2d_grad_fp16)
63	-{
64	- struct UpsampleNearest2dGradCompileInfo {
65	- } compileInfo;
66	- // grad_output: [1, 64, 64, 64] float16
67	- // grad_input: [1, 64, 32, 32] float16
68	- // attrs: output_size=[64,64], input_size=[1,64,32,32], scales_h=0.0, scales_w=0.0
69	- std::vector<gert::TilingContextPara::OpAttr> attrs = {
70	- {"output_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{64, 64})},
71	- {"input_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{1, 64, 32, 32})},
72	- {"scales_h", Ops::Cv::AnyValue::CreateFrom(0.0f)},
73	- {"scales_w", Ops::Cv::AnyValue::CreateFrom(0.0f)}};
74	- gert::TilingContextPara tilingContextPara(
75	- "UpsampleNearest2dGrad",
76	- {
77	- {{{1, 64, 64, 64}, {1, 64, 64, 64}}, ge::DT_FLOAT16, ge::FORMAT_ND}, // grad_output
78	- },
79	- {
80	- {{{1, 64, 32, 32}, {1, 64, 32, 32}}, ge::DT_FLOAT16, ge::FORMAT_ND}, // grad_input
81	- },
82	- attrs, &compileInfo, "Ascend950", 64, 262144, 4096);
83	- uint64_t expectTilingKey = 0;
84	- // totalElements = 643232 = 65536, perCoreElements = ceil(65536/64)=1024, align32=1024, needCoreNum=64
85	- // scaleH = 64/32 = 2.0, scaleW = 64/32 = 2.0
86	- // Tiling data as int64_t: [needCoreNum\|pad] [totalElements] [dimN\|dimC] [dimHin\|dimWin] [dimHout\|dimWout]
87	- // [scaleH\|scaleW]
88	- string expectTilingData = "64 65536 274877906945 137438953504 274877907008 4611686019501129728 ";
89	- std::vector<size_t> expectWorkspaces = {16777216};
90	- ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
91	-}
92	-
93	-// Negative: dimHin=0 (input_size=[1,1,0,2])
94	-TEST_F(UpsampleNearest2dGradTiling, upsample_nearest2d_grad_neg_dimHin_zero)
95	-{
96	- struct UpsampleNearest2dGradCompileInfo {
97	- } compileInfo;
98	- std::vector<gert::TilingContextPara::OpAttr> attrs = {
99	- {"output_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{4, 4})},
100	- {"input_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{1, 1, 0, 2})},
101	- {"scales_h", Ops::Cv::AnyValue::CreateFrom(0.0f)},
102	- {"scales_w", Ops::Cv::AnyValue::CreateFrom(0.0f)}};
103	- gert::TilingContextPara tilingContextPara("UpsampleNearest2dGrad",
104	- {
105	- {{{1, 1, 4, 4}, {1, 1, 4, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
106	- },
107	- {
108	- {{{1, 1, 0, 2}, {1, 1, 0, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
109	- },
110	- attrs, &compileInfo, "Ascend950", 64, 262144, 4096);
111	- ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
112	-}
113	-
114	-// Negative: dimC=0 (input_size=[1,0,2,2])
115	-TEST_F(UpsampleNearest2dGradTiling, upsample_nearest2d_grad_neg_dimC_zero)
116	-{
117	- struct UpsampleNearest2dGradCompileInfo {
118	- } compileInfo;
119	- std::vector<gert::TilingContextPara::OpAttr> attrs = {
120	- {"output_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{4, 4})},
121	- {"input_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{1, 0, 2, 2})},
122	- {"scales_h", Ops::Cv::AnyValue::CreateFrom(0.0f)},
123	- {"scales_w", Ops::Cv::AnyValue::CreateFrom(0.0f)}};
124	- gert::TilingContextPara tilingContextPara("UpsampleNearest2dGrad",
125	- {
126	- {{{1, 1, 4, 4}, {1, 1, 4, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
127	- },
128	- {
129	- {{{1, 0, 2, 2}, {1, 0, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
130	- },
131	- attrs, &compileInfo, "Ascend950", 64, 262144, 4096);
132	- ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
133	-}
134	-
135	-// Negative: dimHout=0 (output_size=[0,4])
136	-TEST_F(UpsampleNearest2dGradTiling, upsample_nearest2d_grad_neg_dimHout_zero)
137	-{
138	- struct UpsampleNearest2dGradCompileInfo {
139	- } compileInfo;
140	- std::vector<gert::TilingContextPara::OpAttr> attrs = {
141	- {"output_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{0, 4})},
142	- {"input_size", Ops::Cv::AnyValue::CreateFrom(std::vector<int64_t>{1, 1, 2, 2})},
143	- {"scales_h", Ops::Cv::AnyValue::CreateFrom(0.0f)},
144	- {"scales_w", Ops::Cv::AnyValue::CreateFrom(0.0f)}};
145	- gert::TilingContextPara tilingContextPara("UpsampleNearest2dGrad",
146	- {
147	- {{{1, 1, 0, 4}, {1, 1, 0, 4}}, ge::DT_FLOAT, ge::FORMAT_ND},
148	- },
149	- {
150	- {{{1, 1, 2, 2}, {1, 1, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND},
151	- },
152	- attrs, &compileInfo, "Ascend950", 64, 262144, 4096);
153	- ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED);
154	-}

Mimage/upsample_nearest2d_grad/tests/ut/op_kernel/CMakeLists.txt+0-1

	@@ -12,7 +12,6 @@ set(upsample_nearest2d_grad_tiling_files
12	${CMAKE_CURRENT_SOURCE_DIR}/../../../../upsample_nearest_exact2d_grad/op_host/upsample_nearest_exact2d_grad_tiling.cpp	12	${CMAKE_CURRENT_SOURCE_DIR}/../../../../upsample_nearest_exact2d_grad/op_host/upsample_nearest_exact2d_grad_tiling.cpp
13	${CMAKE_CURRENT_SOURCE_DIR}/../../../../upsample_nearest_exact2d_grad/op_host/upsample_nearest_exact2d_grad_transpose_tiling.cpp	13	${CMAKE_CURRENT_SOURCE_DIR}/../../../../upsample_nearest_exact2d_grad/op_host/upsample_nearest_exact2d_grad_transpose_tiling.cpp
14	${CMAKE_CURRENT_SOURCE_DIR}/../../../../upsample_nearest_exact2d_grad/op_host/upsample_nearest_exact2d_grad_tiling_arch35.cpp	14	${CMAKE_CURRENT_SOURCE_DIR}/../../../../upsample_nearest_exact2d_grad/op_host/upsample_nearest_exact2d_grad_tiling_arch35.cpp
15	- ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_nearest2d_grad_tiling_arch35.cpp
16	)	15	)
17	# 使用AddOpTestCase	16	# 使用AddOpTestCase
18	# param1：算子名称，以kernel方式命名	17	# param1：算子名称，以kernel方式命名

Mimage/upsample_nearest_exact2d_grad/op_host/upsample_nearest_exact2d_grad_tiling.cpp+1-5

	@@ -49,7 +49,7 @@ constexpr int32_t RESERVED_VALUE = 4;
49		49
50	class UpsampleNearestExact2dGradTiling {	50	class UpsampleNearestExact2dGradTiling {
51	public:	51	public:
52	- explicit UpsampleNearestExact2dGradTiling(gert::TilingContext* context) : tilingContext(context){};	52	+ explicit UpsampleNearestExact2dGradTiling(gert::TilingContext* context) : tilingContext(context) {};
53	ge::graphStatus Init();	53	ge::graphStatus Init();
54	ge::graphStatus RunBigKernelTiling();	54	ge::graphStatus RunBigKernelTiling();
55		55
	@@ -162,10 +162,6 @@ ge::graphStatus UpsampleNearestExact2dGradTiling::RunBigKernelTiling()
162	OP_LOGI(tilingContext->GetNodeName(), "Enter Tiling4UpsampleNearestExact2dGradRegbase");	162	OP_LOGI(tilingContext->GetNodeName(), "Enter Tiling4UpsampleNearestExact2dGradRegbase");
163	return Tiling4UpsampleNearestExact2dGradRegbase(tilingContext);	163	return Tiling4UpsampleNearestExact2dGradRegbase(tilingContext);
164	}	164	}
165	- if (regBase && (opType == "UpsampleNearest2dGrad")) {
166	- OP_LOGI(tilingContext->GetNodeName(), "Enter UpsampleNearest2dGradArch35TilingFunc");
167	- return UpsampleNearest2dGradArch35TilingFunc(tilingContext);
168	- }
169	if (CheckTranspose()) {	165	if (CheckTranspose()) {
170	OP_LOGI(tilingContext->GetNodeName(), "Enter Tiling4UpsampleNearestExact2dGradTranspose");	166	OP_LOGI(tilingContext->GetNodeName(), "Enter Tiling4UpsampleNearestExact2dGradTranspose");
171	return tiling4UpsampleNearestExact2dGradTransposeTiling(tilingContext);	167	return tiling4UpsampleNearestExact2dGradTransposeTiling(tilingContext);

Mimage/upsample_nearest_exact2d_grad/op_host/upsample_nearest_exact2d_grad_tiling.h+0-1

	@@ -112,7 +112,6 @@ REGISTER_TILING_DATA_CLASS(UpsampleNearest2dGrad_110, UpsampleNearestExact2dGrad
112		112
113	ge::graphStatus tiling4UpsampleNearestExact2dGradTransposeTiling(gert::TilingContext* context);	113	ge::graphStatus tiling4UpsampleNearestExact2dGradTransposeTiling(gert::TilingContext* context);
114	ge::graphStatus Tiling4UpsampleNearestExact2dGradRegbase(gert::TilingContext* context);	114	ge::graphStatus Tiling4UpsampleNearestExact2dGradRegbase(gert::TilingContext* context);
115	-ge::graphStatus UpsampleNearest2dGradArch35TilingFunc(gert::TilingContext* context);
116		115
117	inline float compute_scale_value(int64_t in_size, int64_t out_size, const float* scale)	116	inline float compute_scale_value(int64_t in_size, int64_t out_size, const float* scale)
118	{	117	{

Mimage/upsample_nearest_exact2d_grad/tests/ut/op_kernel/CMakeLists.txt+0-1

	@@ -25,7 +25,6 @@ if (UT_TEST_ALL OR OP_KERNEL_UT)
25	${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_nearest_exact2d_grad_tiling.cpp	25	${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_nearest_exact2d_grad_tiling.cpp
26	${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_nearest_exact2d_grad_transpose_tiling.cpp	26	${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_nearest_exact2d_grad_transpose_tiling.cpp
27	${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_nearest_exact2d_grad_tiling_arch35.cpp	27	${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/upsample_nearest_exact2d_grad_tiling_arch35.cpp
28	- ${CMAKE_CURRENT_SOURCE_DIR}/../../../../upsample_nearest2d_grad/op_host/upsample_nearest2d_grad_tiling_arch35.cpp
29	)	28	)
30	# 使用AddOpTestCase	29	# 使用AddOpTestCase
31	# param1：算子名称，以kernel方式命名	30	# param1：算子名称，以kernel方式命名

Mscripts/kernel/binary_config/ascendc_config.json+1-1

	@@ -241,7 +241,7 @@
241	{"name":"UpsampleNearest3dGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},	241	{"name":"UpsampleNearest3dGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
242	{"name":"UpsampleNearestExact3dGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},	242	{"name":"UpsampleNearestExact3dGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
243	{"name":"UpsampleBicubic2dAAGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},	243	{"name":"UpsampleBicubic2dAAGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},
244	- {"name":"UpsampleNearest2dGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},	244	+ {"name":"UpsampleNearest2dGrad", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},
245	{"name":"UpsampleNearestExact2dGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},	245	{"name":"UpsampleNearestExact2dGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
246	{"name":"RmsNormGrad", "compute_units": ["ascend950", "ascend910_93", "ascend910b", "ascend910", "ascend310p"], "auto_sync" : false},	246	{"name":"RmsNormGrad", "compute_units": ["ascend950", "ascend910_93", "ascend910b", "ascend910", "ascend310p"], "auto_sync" : false},
247	{"name":"RotaryPositionEmbedding", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : true},	247	{"name":"RotaryPositionEmbedding", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync" : true},