Star536
364
代码
代码
Issues16
Pull Requests35
Actions
项目讨论
Wiki
分析
Star536
364
已合并
aclnnUpsampleTrilinear3d算子支持950 #927
tommyCR7创建于 5月30日
aclnnUpsampleTrilinear3d算子支持950 #927
已合并
tommyCR7创建于 5月30日
从已删除 :master合入到cann/ops-cvmaster
共 15 个文件变更+1596-25
Aimage/resize_upsample_trilinear/op_host/config/ascend950/resize_upsample_trilinear_binary.json+167-0
@@ -0,0 +1,167 @@
1{
2 "op_type": "ResizeUpsampleTrilinear",
3 "op_list": [
4 {
5 "bin_filename": "ResizeUpsampleTrilinear_bin_float32",
6 "inputs": [
7 {
8 "name": "input",
9 "index": 0,
10 "dtype": "float32",
11 "format": "ND",
12 "paramType": "required",
13 "shape": [
14 -2
15 ]
16 }
17 ],
18 "outputs": [
19 {
20 "name": "output",
21 "index": 0,
22 "dtype": "float32",
23 "format": "ND",
24 "paramType": "required",
25 "shape": [
26 -2
27 ]
28 }
29 ],
30 "attrs": [
31 {
32 "name": "output_size",
33 "dtype": "list_int",
34 "value": []
35 },
36 {
37 "name": "align_corners",
38 "dtype": "bool",
39 "value": false
40 },
41 {
42 "name": "scales_d",
43 "dtype": "float",
44 "value": 0.0
45 },
46 {
47 "name": "scales_h",
48 "dtype": "float",
49 "value": 0.0
50 },
51 {
52 "name": "scales_w",
53 "dtype": "float",
54 "value": 0.0
55 }
56 ]
57 },
58 {
59 "bin_filename": "ResizeUpsampleTrilinear_bin_float16",
60 "inputs": [
61 {
62 "name": "input",
63 "index": 0,
64 "dtype": "float16",
65 "format": "ND",
66 "paramType": "required",
67 "shape": [
68 -2
69 ]
70 }
71 ],
72 "outputs": [
73 {
74 "name": "output",
75 "index": 0,
76 "dtype": "float16",
77 "format": "ND",
78 "paramType": "required",
79 "shape": [
80 -2
81 ]
82 }
83 ],
84 "attrs": [
85 {
86 "name": "output_size",
87 "dtype": "list_int",
88 "value": []
89 },
90 {
91 "name": "align_corners",
92 "dtype": "bool",
93 "value": false
94 },
95 {
96 "name": "scales_d",
97 "dtype": "float",
98 "value": 0.0
99 },
100 {
101 "name": "scales_h",
102 "dtype": "float",
103 "value": 0.0
104 },
105 {
106 "name": "scales_w",
107 "dtype": "float",
108 "value": 0.0
109 }
110 ]
111 },
112 {
113 "bin_filename": "ResizeUpsampleTrilinear_bin_bfloat16",
114 "inputs": [
115 {
116 "name": "input",
117 "index": 0,
118 "dtype": "bfloat16",
119 "format": "ND",
120 "paramType": "required",
121 "shape": [
122 -2
123 ]
124 }
125 ],
126 "outputs": [
127 {
128 "name": "output",
129 "index": 0,
130 "dtype": "bfloat16",
131 "format": "ND",
132 "paramType": "required",
133 "shape": [
134 -2
135 ]
136 }
137 ],
138 "attrs": [
139 {
140 "name": "output_size",
141 "dtype": "list_int",
142 "value": []
143 },
144 {
145 "name": "align_corners",
146 "dtype": "bool",
147 "value": false
148 },
149 {
150 "name": "scales_d",
151 "dtype": "float",
152 "value": 0.0
153 },
154 {
155 "name": "scales_h",
156 "dtype": "float",
157 "value": 0.0
158 },
159 {
160 "name": "scales_w",
161 "dtype": "float",
162 "value": 0.0
163 }
164 ]
165 }
166 ]
167 }
Aimage/resize_upsample_trilinear/op_host/config/ascend950/resize_upsample_trilinear_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2; [某算子]
3; default=xx
4; ascendxx=xx
5; 其中，default为默认mode，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[ResizeUpsampleTrilinear]
13default=0
Mimage/resize_upsample_trilinear/op_host/op_api/aclnn_upsample_trilinear_3d.cpp+19-24
@@ -37,6 +37,8 @@ static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {
37 op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_DOUBLE, op::DataType::DT_BF16};37 op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_DOUBLE, op::DataType::DT_BF16};
38static const std::initializer_list<op::DataType> ASCEND310P_DTYPE_SUPPORT_LIST = {38static const std::initializer_list<op::DataType> ASCEND310P_DTYPE_SUPPORT_LIST = {
39 op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT};39 op::DataType::DT_FLOAT16, op::DataType::DT_FLOAT};
40static const std::initializer_list<op::DataType> A5_DTYPE_SUPPORT_LIST = {
41 op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
40static constexpr size_t EXPECT_SIZE = 3;42static constexpr size_t EXPECT_SIZE = 3;
41static constexpr float MAX_SUPPORT_SCALE = 50.0;43static constexpr float MAX_SUPPORT_SCALE = 50.0;
42 44 
@@ -53,6 +55,8 @@ static bool CheckDtypeValid(const aclTensor* self, const aclTensor* out)
53 auto curArch = GetCurrentPlatformInfo().GetCurNpuArch();55 auto curArch = GetCurrentPlatformInfo().GetCurNpuArch();
54 if (curArch == NpuArch::DAV_2002 || curArch == NpuArch::DAV_3002) {56 if (curArch == NpuArch::DAV_2002 || curArch == NpuArch::DAV_3002) {
55 OP_CHECK_DTYPE_NOT_SUPPORT(self, ASCEND310P_DTYPE_SUPPORT_LIST, return false);57 OP_CHECK_DTYPE_NOT_SUPPORT(self, ASCEND310P_DTYPE_SUPPORT_LIST, return false);
58 } else if (curArch == NpuArch::DAV_3510) {
59 OP_CHECK_DTYPE_NOT_SUPPORT(self, A5_DTYPE_SUPPORT_LIST, return false);
56 } else {60 } else {
57 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST, return false);61 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST, return false);
58 }62 }
@@ -167,37 +171,28 @@ static float AsComputeScale(bool alignCorners, int64_t inputSize, int64_t output
167 }171 }
168}172}
169 173 
170static bool CheckUplimit(const aclTensor* self, const aclTensor* out)174static bool CheckTensorShapeUplimit(const aclTensor* tensor, const char* name)
171{175{
172 if (IsRegBase()) {176 int64_t dim0 = tensor->GetViewShape().GetDim(DIM_ZERO);
173 return true;177 int64_t dim1 = tensor->GetViewShape().GetDim(DIM_ONE);
174 }178 int64_t dim2 = tensor->GetViewShape().GetDim(DIM_TWO);
175 int64_t inN = self->GetViewShape().GetDim(DIM_ZERO);179 int64_t dim3 = tensor->GetViewShape().GetDim(DIM_THREE);
176 int64_t inC = self->GetViewShape().GetDim(DIM_ONE);180 int64_t dim4 = tensor->GetViewShape().GetDim(DIM_FOUR);
177 int64_t inD = self->GetViewShape().GetDim(DIM_TWO);
178 int64_t inH = self->GetViewShape().GetDim(DIM_THREE);
179 int64_t inW = self->GetViewShape().GetDim(DIM_FOUR);
180 int64_t outN = out->GetViewShape().GetDim(DIM_ZERO);
181 int64_t outC = out->GetViewShape().GetDim(DIM_ONE);
182 int64_t outD = out->GetViewShape().GetDim(DIM_TWO);
183 int64_t outH = out->GetViewShape().GetDim(DIM_THREE);
184 int64_t outW = out->GetViewShape().GetDim(DIM_THREE);
185 
186 OP_CHECK(181 OP_CHECK(
187 inN <= INT32_MAX && inC <= INT32_MAX && inD <= INT32_MAX && inH <= INT32_MAX && inW <= INT32_MAX,182 dim0 <= INT32_MAX && dim1 <= INT32_MAX && dim2 <= INT32_MAX && dim3 <= INT32_MAX && dim4 <= INT32_MAX,
188 OP_LOGE(183 OP_LOGE(
189 ACLNN_ERR_PARAM_INVALID, "Self sizes should not be greater than %d, bug got self(%ld, %ld, %ld, %ld, %ld)",184 ACLNN_ERR_PARAM_INVALID,
190 INT32_MAX, inN, inC, inD, inH, inW),185 "%s sizes should not be greater than %d, but got %s(%ld, %ld, %ld, %ld, %ld)",
191 return false);186 name, INT32_MAX, name, dim0, dim1, dim2, dim3, dim4),
192 OP_CHECK(
193 outN <= INT32_MAX && outC <= INT32_MAX && outD <= INT32_MAX && outH <= INT32_MAX && outW <= INT32_MAX,
194 OP_LOGE(
195 ACLNN_ERR_PARAM_INVALID, "Out sizes should not be greater than %d, bug got out(%ld, %ld, %ld, %ld, %ld)",
196 INT32_MAX, outN, outC, outD, outH, outW),
197 return false);187 return false);
198 return true;188 return true;
199}189}
200 190 
191static bool CheckUplimit(const aclTensor* self, const aclTensor* out)
192{
193 return CheckTensorShapeUplimit(self, "Self") && CheckTensorShapeUplimit(out, "Out");
194}
195 
201static aclnnStatus CheckParams(const aclTensor* self, const aclIntArray* outputSize, const aclTensor* out)196static aclnnStatus CheckParams(const aclTensor* self, const aclIntArray* outputSize, const aclTensor* out)
202{197{
203 // 1. 检查参数是否为空指针198 // 1. 检查参数是否为空指针
Mimage/resize_upsample_trilinear/op_host/op_api/upsample_trilinear_3d.cpp+17-0
@@ -20,6 +20,7 @@
20#include "opdev/op_log.h"20#include "opdev/op_log.h"
21#include "opdev/platform.h"21#include "opdev/platform.h"
22#include "aclnn_kernels/cast.h"22#include "aclnn_kernels/cast.h"
23#include "op_api/aclnn_check.h"
23 24 
24using namespace op;25using namespace op;
25 26 
@@ -107,6 +108,22 @@ const aclTensor* UpsampleTrilinear3dNcdhw(
107 out = l0op::Cast(out, op::DataType::DT_BF16, executor);108 out = l0op::Cast(out, op::DataType::DT_BF16, executor);
108 }109 }
109 return out;110 return out;
111 } else if (curArch == NpuArch::DAV_3510 &&
112 CheckType(self->GetDataType(), AICORE_DTYPE_SUPPORT_LIST) &&
113 CheckScales(checkScaleW, checkScaleH, checkScaleD)) {
114 // A5 SIMT Kernel handles FP16/BF16->FP32 Cast internally via if constexpr,
115 // no external Cast needed (unlike DAV_2201 path which Casts externally).
116 // NDHWC format is handled by L2 layer upsampleTrilinear3dCompute
117 // (transpose NDHWC->NCDHW before calling this L0 function).
118 const aclTensor* out = executor->AllocTensor(outShape, self->GetDataType(), self->GetStorageFormat());
119 ret = ADD_TO_LAUNCHER_LIST_AICORE(
120 ResizeUpsampleTrilinear, OP_INPUT(self), OP_OUTPUT(out),
121 OP_ATTR(outputSize, alignCorners, scalesD, scalesH, scalesW));
122 OP_CHECK(
123 ret == ACLNN_SUCCESS,
124 OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "ResizeUpsampleTrilinearA5AiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),
125 return nullptr);
126 return out;
110 } else if (127 } else if (
111 (curArch == NpuArch::DAV_2002) && CheckType(self->GetDataType(), AICORE_310P_SUPPORT_LIST) &&128 (curArch == NpuArch::DAV_2002) && CheckType(self->GetDataType(), AICORE_310P_SUPPORT_LIST) &&
112 CheckScales(checkScaleW, checkScaleH, checkScaleD)) {129 CheckScales(checkScaleW, checkScaleH, checkScaleD)) {
Mimage/resize_upsample_trilinear/op_host/resize_upsample_trilinear_def.cpp+19-0
@@ -59,6 +59,25 @@ public:
59 .Format({ge::FORMAT_ND, ge::FORMAT_ND})59 .Format({ge::FORMAT_ND, ge::FORMAT_ND})
60 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});60 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
61 this->AICore().AddConfig("ascend310p", config310p);61 this->AICore().AddConfig("ascend310p", config310p);
62 
63 OpAICoreConfig config950;
64 config950.DynamicCompileStaticFlag(true)
65 .DynamicFormatFlag(true)
66 .DynamicRankSupportFlag(true)
67 .DynamicShapeSupportFlag(true)
68 .NeedCheckSupportFlag(false)
69 .PrecisionReduceFlag(true);
70 config950.Input("input")
71 .ParamType(REQUIRED)
72 .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
73 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
74 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
75 config950.Output("output")
76 .ParamType(REQUIRED)
77 .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
78 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
79 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
80 this->AICore().AddConfig("ascend950", config950);
62 }81 }
63};82};
64 83 
Mimage/resize_upsample_trilinear/op_host/resize_upsample_trilinear_tiling.cpp+5-0
@@ -555,6 +555,11 @@ static uint64_t GetWorkSpaceSize(
555 555 
556static ge::graphStatus Tiling4UpsampleTrilinear(gert::TilingContext* context)556static ge::graphStatus Tiling4UpsampleTrilinear(gert::TilingContext* context)
557{557{
558 auto templateStatus = Ops::Cv::OpTiling::TilingRegistry::GetInstance().DoTilingImpl(context);
559 if (templateStatus == ge::GRAPH_SUCCESS) {
560 return ge::GRAPH_SUCCESS;
561 }
562 
558 int batchIdx = OUTPUT_DEPTH_IDX;563 int batchIdx = OUTPUT_DEPTH_IDX;
559 int channelIdx = OUTPUT_HEIGHT_IDX;564 int channelIdx = OUTPUT_HEIGHT_IDX;
560 int weightIdx = SCALE_W_IDX;565 int weightIdx = SCALE_W_IDX;
Aimage/resize_upsample_trilinear/op_host/resize_upsample_trilinear_tiling_arch35.cpp+393-0
@@ -0,0 +1,393 @@
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/*!
12 * \file resize_upsample_trilinear_tiling_arch35.cpp
13 * \brief resize_upsample_trilinear tiling implementation for A5 architecture (SIMT paradigm)
14 */
15#include "resize_upsample_trilinear_tiling_arch35.h"
16#include "image/resize_upsample_trilinear/op_kernel/arch35/resize_upsample_trilinear_tiling_key.h"
17#include "image/resize_upsample_trilinear/op_kernel/arch35/resize_upsample_trilinear_tiling_data.h"
18#include "log/log.h"
19#include "securec.h"
20#include <cmath>
21#include <algorithm>
22 
23namespace optiling {
24static constexpr size_t DIM_0 = 0;
25static constexpr size_t DIM_1 = 1;
26static constexpr size_t DIM_2 = 2;
27static constexpr size_t DIM_3 = 3;
28static constexpr size_t DIM_4 = 4;
29static constexpr size_t DIM_5 = 5;
30static constexpr uint32_t MIN_THREADS_PER_BLOCK = 32;
31static constexpr uint32_t MAX_THREADS_PER_BLOCK = 512;
32static constexpr uint32_t MIN_ELEMENTS_PER_THREAD = 1;
33static constexpr uint32_t MAX_BLOCKS = 1024;
34static constexpr uint32_t MAX_ELEMENTS_PER_THREAD = 1024;
35static constexpr float MAX_SUPPORT_SCALE = 50.0f;
36 
37static float ComputeSourceCoordScale(bool alignCorners, int64_t inputSize, int64_t outputSize, float scale)
38{
39 if (outputSize == inputSize) {
40 return 1.0f;
41 }
42 if (alignCorners) {
43 if (outputSize > 1) {
44 return static_cast<float>(inputSize - 1) / static_cast<float>(outputSize - 1);
45 } else {
46 return 0.0f;
47 }
48 } else {
49 if (scale > 0.0f) {
50 return 1.0f / scale;
51 } else {
52 return static_cast<float>(inputSize) / static_cast<float>(outputSize);
53 }
54 }
55}
56 
57bool ResizeUpsampleTrilinearArch35Tiling::IsCapable()
58{
59 auto platformInfo = context_->GetPlatformInfo();
60 if (platformInfo == nullptr) {
61 OP_LOGE(context_->GetNodeName(), "platformInfo is nullptr, cannot get platform info for arch35 tiling.");
R
Rrenruhai6月2日

return false这里建议加上日志，方便后面定位问题

likedislike
T
tommyCR7
6月2日 评论：
likedislike
62 return false;
63 }
64 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
65 auto npuArch = ascendcPlatform.GetCurNpuArch();
66 if (npuArch == NpuArch::DAV_3510) {
67 coreNum_ = ascendcPlatform.GetCoreNumAiv();
68 platformInfoCached_ = true;
69 return true;
70 }
71 return false;
72}
73 
74ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::GetPlatformInfo()
75{
76 if (platformInfoCached_) {
77 OP_CHECK_IF(
78 coreNum_ <= 0,
79 OP_LOGE(context_->GetNodeName(), "coreNum is error: %d", coreNum_),
80 return ge::GRAPH_FAILED);
81 OP_LOGI(context_->GetNodeName(), "A5 coreNum(AIV)=%d", coreNum_);
82 return ge::GRAPH_SUCCESS;
83 }
84 auto platformInfo = context_->GetPlatformInfo();
85 OP_CHECK_NULL_WITH_CONTEXT(context_, platformInfo);
86 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
87 coreNum_ = ascendcPlatform.GetCoreNumAiv();
88 OP_CHECK_IF(
89 coreNum_ <= 0,
90 OP_LOGE(context_->GetNodeName(), "coreNum is error: %d", coreNum_),
91 return ge::GRAPH_FAILED);
92 OP_LOGI(context_->GetNodeName(), "A5 coreNum(AIV)=%d", coreNum_);
93 return ge::GRAPH_SUCCESS;
94}
95 
96ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::ValidateAndGetInputShape()
97{
98 auto inputShape = context_->GetInputShape(DIM_0);
99 OP_CHECK_NULL_WITH_CONTEXT(context_, inputShape);
100 gert::Shape inShape = inputShape->GetStorageShape();
101 int32_t inDims = inShape.GetDimNum();
102 OP_CHECK_IF(
103 inDims != DIM_5,
104 OP_LOGE(context_->GetNodeName(), "input dims must be 5, but got %d", inDims),
105 return ge::GRAPH_FAILED);
106 inN_ = inShape.GetDim(DIM_0);
107 inC_ = inShape.GetDim(DIM_1);
108 inputD_ = inShape.GetDim(DIM_2);
109 inputH_ = inShape.GetDim(DIM_3);
110 inputW_ = inShape.GetDim(DIM_4);
111 return ge::GRAPH_SUCCESS;
112}
113 
114ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::ValidateAndGetOutputShape()
115{
116 auto outputShape = context_->GetOutputShape(DIM_0);
117 OP_CHECK_NULL_WITH_CONTEXT(context_, outputShape);
118 gert::Shape outShape = outputShape->GetStorageShape();
119 int32_t outDims = outShape.GetDimNum();
120 OP_CHECK_IF(
121 outDims != DIM_5,
122 OP_LOGE(context_->GetNodeName(), "output dims must be 5, but got %d", outDims),
123 return ge::GRAPH_FAILED);
124 outN_ = outShape.GetDim(DIM_0);
125 outC_ = outShape.GetDim(DIM_1);
126 outputD_ = outShape.GetDim(DIM_2);
127 outputH_ = outShape.GetDim(DIM_3);
128 outputW_ = outShape.GetDim(DIM_4);
129 return ge::GRAPH_SUCCESS;
130}
131 
132ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::ValidateShapeValues()
133{
134 OP_CHECK_IF(
135 inputD_ <= 0 || inputH_ <= 0 || inputW_ <= 0,
136 OP_LOGE(context_->GetNodeName(), "input D/H/W must be positive, got D=%ld H=%ld W=%ld",
137 inputD_, inputH_, inputW_),
138 return ge::GRAPH_FAILED);
139 OP_CHECK_IF(
140 outputD_ <= 0 || outputH_ <= 0 || outputW_ <= 0,
141 OP_LOGE(context_->GetNodeName(), "output D/H/W must be positive, got D=%ld H=%ld W=%ld",
142 outputD_, outputH_, outputW_),
143 return ge::GRAPH_FAILED);
144 OP_CHECK_IF(
145 inN_ <= 0 || inC_ <= 0,
146 OP_LOGE(context_->GetNodeName(), "input N/C must be positive, got N=%ld C=%ld", inN_, inC_),
147 return ge::GRAPH_FAILED);
148 batchCount_ = inN_ * inC_;
149 return ge::GRAPH_SUCCESS;
150}
151 
152ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::ExtractAttrsAndComputeScales()
153{
154 auto attrs = context_->GetAttrs();
155 OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
156 const bool* alignCornersPtr = attrs->GetAttrPointer<bool>(DIM_1);
157 alignCorners_ = (alignCornersPtr != nullptr && *alignCornersPtr) ? 1 : 0;
158 const float* scalesDPtr = attrs->GetAttrPointer<float>(DIM_2);
159 const float* scalesHPtr = attrs->GetAttrPointer<float>(DIM_3);
160 const float* scalesWPtr = attrs->GetAttrPointer<float>(DIM_4);
161 float attrScaleD = (scalesDPtr != nullptr) ? *scalesDPtr : 0.0f;
162 float attrScaleH = (scalesHPtr != nullptr) ? *scalesHPtr : 0.0f;
163 float attrScaleW = (scalesWPtr != nullptr) ? *scalesWPtr : 0.0f;
164 scaleD_ = ComputeSourceCoordScale(alignCorners_ != 0, inputD_, outputD_, attrScaleD);
165 scaleH_ = ComputeSourceCoordScale(alignCorners_ != 0, inputH_, outputH_, attrScaleH);
166 scaleW_ = ComputeSourceCoordScale(alignCorners_ != 0, inputW_, outputW_, attrScaleW);
167 return ge::GRAPH_SUCCESS;
168}
169 
170ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::DetermineDtypeKey()
171{
172 auto dataType = context_->GetInputDesc(DIM_0)->GetDataType();
173 if (dataType == ge::DT_FLOAT) {
174 dtypeKey_ = TPL_DTYPE_FP32;
175 } else if (dataType == ge::DT_FLOAT16) {
176 dtypeKey_ = TPL_DTYPE_FP16;
177 } else if (dataType == ge::DT_BF16) {
178 dtypeKey_ = TPL_DTYPE_BF16;
179 } else {
180 OP_LOGE(context_->GetNodeName(), "unsupported data type: %d", static_cast<int>(dataType));
181 return ge::GRAPH_FAILED;
182 }
183 return ge::GRAPH_SUCCESS;
184}
185 
186ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::GetShapeAttrsInfo()
187{
188 OP_CHECK_IF(ValidateAndGetInputShape() != ge::GRAPH_SUCCESS, , return ge::GRAPH_FAILED);
189 OP_CHECK_IF(ValidateAndGetOutputShape() != ge::GRAPH_SUCCESS, , return ge::GRAPH_FAILED);
190 OP_CHECK_IF(ValidateShapeValues() != ge::GRAPH_SUCCESS, , return ge::GRAPH_FAILED);
191 OP_CHECK_IF(ExtractAttrsAndComputeScales() != ge::GRAPH_SUCCESS, , return ge::GRAPH_FAILED);
192 OP_CHECK_IF(DetermineDtypeKey() != ge::GRAPH_SUCCESS, , return ge::GRAPH_FAILED);
193 return ge::GRAPH_SUCCESS;
194}
195 
196ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::ValidateDimensionsAndComputeTotal()
197{
198 OP_CHECK_IF(
199 inN_ != outN_ || inC_ != outC_,
200 OP_LOGE(context_->GetNodeName(), "input and output N/C dimensions must match"),
201 return ge::GRAPH_FAILED);
202 int64_t dhw = outputD_ * outputH_;
203 OP_CHECK_IF(
204 outputH_ > 0 && dhw > INT64_MAX / outputW_,
205 OP_LOGE(context_->GetNodeName(), "output D*H*W overflow, outputD=%ld outputH=%ld outputW=%ld",
206 outputD_, outputH_, outputW_),
207 return ge::GRAPH_FAILED);
208 dhw *= outputW_;
209 OP_CHECK_IF(
210 batchCount_ > INT64_MAX / dhw,
211 OP_LOGE(context_->GetNodeName(), "totalElements overflow detected, batchCount=%ld outputD=%ld outputH=%ld outputW=%ld",
212 batchCount_, outputD_, outputH_, outputW_),
213 return ge::GRAPH_FAILED);
214 totalElements_ = static_cast<uint64_t>(batchCount_) * static_cast<uint64_t>(dhw);
215 float checkScaleD = (scaleD_ > 0.0f) ? 1.0f / scaleD_ : 0.0f;
216 float checkScaleH = (scaleH_ > 0.0f) ? 1.0f / scaleH_ : 0.0f;
217 float checkScaleW = (scaleW_ > 0.0f) ? 1.0f / scaleW_ : 0.0f;
218 OP_CHECK_IF(
219 checkScaleD > MAX_SUPPORT_SCALE || checkScaleH > MAX_SUPPORT_SCALE || checkScaleW > MAX_SUPPORT_SCALE,
220 OP_LOGE(context_->GetNodeName(),
221 "scales exceed max support scale %f, got scaleD=%f scaleH=%f scaleW=%f",
222 MAX_SUPPORT_SCALE, checkScaleD, checkScaleH, checkScaleW),
223 return ge::GRAPH_FAILED);
224 return ge::GRAPH_SUCCESS;
225}
226 
227ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::HandleZeroElements()
228{
229 if (totalElements_ == 0) {
230 usedCoreNum_ = 1;
231 blockCount_ = 1;
232 baseElementsPerBlock_ = 0;
233 tailElements_ = 0;
234 threadsPerBlock_ = MIN_THREADS_PER_BLOCK;
235 elementsPerThread_ = MIN_ELEMENTS_PER_THREAD;
236 useInt32_ = 1;
237 return ge::GRAPH_SUCCESS;
238 }
239 return ge::GRAPH_FAILED;
240}
241 
242ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::ComputeThreadBlockConfig()
243{
244 OP_CHECK_IF(
245 totalElements_ > static_cast<uint64_t>(MAX_ELEMENTS_PER_THREAD) * MAX_THREADS_PER_BLOCK * MAX_BLOCKS,
246 OP_LOGE(context_->GetNodeName(), "totalElements too large for SIMT processing: %lu", totalElements_),
247 return ge::GRAPH_FAILED);
248 threadsPerBlock_ = MAX_THREADS_PER_BLOCK;
249 elementsPerThread_ = MIN_ELEMENTS_PER_THREAD;
250 uint64_t totalThreads = (totalElements_ + elementsPerThread_ - 1) / elementsPerThread_;
251 blockCount_ = static_cast<uint32_t>((totalThreads + threadsPerBlock_ - 1) / threadsPerBlock_);
252 if (blockCount_ == 0) {
253 blockCount_ = 1;
254 }
255 if (blockCount_ > MAX_BLOCKS) {
256 uint32_t minEPT = static_cast<uint32_t>(
257 (totalElements_ + static_cast<uint64_t>(MAX_BLOCKS) * threadsPerBlock_ - 1) /
258 (static_cast<uint64_t>(MAX_BLOCKS) * threadsPerBlock_));
259 elementsPerThread_ = std::max(minEPT, elementsPerThread_);
260 OP_CHECK_IF(
261 elementsPerThread_ > MAX_ELEMENTS_PER_THREAD,
262 OP_LOGE(context_->GetNodeName(), "cannot reduce blockCount below MAX_BLOCKS after max elementsPerThread"),
263 return ge::GRAPH_FAILED);
264 if(elementsPerThread_ == 0){
265 return ge::GRAPH_FAILED;}
266 totalThreads = (totalElements_ + elementsPerThread_ - 1) / elementsPerThread_;
267 blockCount_ = static_cast<uint32_t>((totalThreads + threadsPerBlock_ - 1) / threadsPerBlock_);
268 if (blockCount_ == 0) {
269 blockCount_ = 1;
270 }
271 } else if (totalThreads < threadsPerBlock_) {
272 threadsPerBlock_ = static_cast<uint32_t>(totalThreads);
273 if (threadsPerBlock_ < MIN_THREADS_PER_BLOCK) {
274 threadsPerBlock_ = MIN_THREADS_PER_BLOCK;
275 }
276 blockCount_ = 1;
277 }
278 threadsPerBlock_ = ((threadsPerBlock_ + 31) / 32) * 32;
279 return ge::GRAPH_SUCCESS;
280}
281 
282ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::ComputeFinalTilingConfig()
283{
284 usedCoreNum_ = (blockCount_ < static_cast<uint32_t>(coreNum_)) ? blockCount_ : static_cast<uint32_t>(coreNum_);
285 baseElementsPerBlock_ = threadsPerBlock_ * elementsPerThread_;
286 if (blockCount_ == 1) {
287 tailElements_ = static_cast<uint32_t>(totalElements_);
288 } else {
289 tailElements_ = static_cast<uint32_t>(
290 totalElements_ - static_cast<uint64_t>(blockCount_ - 1) * baseElementsPerBlock_);
291 }
292 
293 static constexpr int64_t INT32_MAX_VAL = 2147483647LL;
294 int64_t strideBcInput = inputD_ * inputH_ * inputW_;
295 int64_t strideBcOutput = outputD_ * outputH_ * outputW_;
296 int64_t strideDInput = inputH_ * inputW_;
297 int64_t strideDOutput = outputH_ * outputW_;
298 useInt32_ = (totalElements_ <= static_cast<uint64_t>(INT32_MAX_VAL) &&
299 strideBcInput <= INT32_MAX_VAL &&
300 strideBcOutput <= INT32_MAX_VAL &&
301 strideDInput <= INT32_MAX_VAL &&
302 strideDOutput <= INT32_MAX_VAL &&
303 inputW_ <= INT32_MAX_VAL &&
304 outputW_ <= INT32_MAX_VAL &&
305 inputD_ <= INT32_MAX_VAL &&
306 inputH_ <= INT32_MAX_VAL &&
307 outputD_ <= INT32_MAX_VAL &&
308 outputH_ <= INT32_MAX_VAL) ? 1 : 0;
309 
310 return ge::GRAPH_SUCCESS;
311}
312 
313ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::DoOpTiling()
314{
315 OP_CHECK_IF(ValidateDimensionsAndComputeTotal() != ge::GRAPH_SUCCESS, , return ge::GRAPH_FAILED);
316 if (HandleZeroElements() == ge::GRAPH_SUCCESS) {
317 return ge::GRAPH_SUCCESS;
318 }
319 OP_CHECK_IF(ComputeThreadBlockConfig() != ge::GRAPH_SUCCESS, , return ge::GRAPH_FAILED);
320 OP_CHECK_IF(ComputeFinalTilingConfig() != ge::GRAPH_SUCCESS, , return ge::GRAPH_FAILED);
321 return ge::GRAPH_SUCCESS;
322}
323 
324ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::DoLibApiTiling()
325{
326 return ge::GRAPH_SUCCESS;
327}
328 
329uint64_t ResizeUpsampleTrilinearArch35Tiling::GetTilingKey() const
330{
331 return GET_TPL_TILING_KEY(dtypeKey_);
332}
333 
334ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::GetWorkspaceSize()
335{
336 workspaceSize_ = 0;
337 return ge::GRAPH_SUCCESS;
338}
339 
340ge::graphStatus ResizeUpsampleTrilinearArch35Tiling::PostTiling()
341{
342 ResizeUpsampleTrilinearArch35TilingData tilingData;
343 tilingData.elements_per_thread = elementsPerThread_;
344 tilingData.block_count = blockCount_;
345 tilingData.used_core_num = usedCoreNum_;
346 tilingData.base_elements_per_block = baseElementsPerBlock_;
347 tilingData.tail_elements = tailElements_;
348 tilingData.total_elements = totalElements_;
349 tilingData.batch_count = batchCount_;
350 tilingData.input_d = inputD_;
351 tilingData.input_h = inputH_;
352 tilingData.input_w = inputW_;
353 tilingData.output_d = outputD_;
354 tilingData.output_h = outputH_;
355 tilingData.output_w = outputW_;
356 tilingData.scale_d = scaleD_;
357 tilingData.scale_h = scaleH_;
358 tilingData.scale_w = scaleW_;
359 tilingData.align_corners = alignCorners_;
360 tilingData.use_int32 = useInt32_;
361 
362 auto* rawTilingData = context_->GetRawTilingData();
363 OP_CHECK_NULL_WITH_CONTEXT(context_, rawTilingData);
364 uint32_t tilingSize = sizeof(tilingData);
365 OP_CHECK_IF(
366 rawTilingData->GetCapacity() < tilingSize,
367 OP_LOGE(context_->GetNodeName(), "tiling data capacity %zu is less than required %u",
368 rawTilingData->GetCapacity(), tilingSize),
369 return ge::GRAPH_FAILED);
370 errno_t cpyRet = memcpy_s(rawTilingData->GetData(), rawTilingData->GetCapacity(), &tilingData, tilingSize);
371 if (cpyRet != EOK) {
372 OP_LOGE(context_->GetNodeName(), "memcpy_s tiling data failed, ret=%d.", cpyRet);
373 return ge::GRAPH_FAILED;
374 }
375 rawTilingData->SetDataSize(tilingSize);
376 
377 context_->SetBlockDim(usedCoreNum_);
378 
379 OP_LOGI(context_->GetNodeName(),
380 "ResizeUpsampleTrilinear A5 tiling: dtypeKey=%lu, elementsPerThread=%u, "
381 "blockCount=%u, usedCoreNum=%u, baseElementsPerBlock=%u, tailElements=%u, totalElements=%lu",
382 dtypeKey_, elementsPerThread_, blockCount_, usedCoreNum_,
383 baseElementsPerBlock_, tailElements_, totalElements_);
384 OP_LOGI(context_->GetNodeName(),
385 "ResizeUpsampleTrilinear A5 tiling: batchCount=%ld, inputD=%ld, inputH=%ld, inputW=%ld, "
386 "outputD=%ld, outputH=%ld, outputW=%ld, scaleD=%f, scaleH=%f, scaleW=%f, alignCorners=%d, useInt32=%d",
387 batchCount_, inputD_, inputH_, inputW_, outputD_, outputH_, outputW_,
388 scaleD_, scaleH_, scaleW_, alignCorners_, useInt32_);
389 return ge::GRAPH_SUCCESS;
390}
391 
392REGISTER_OPS_TILING_TEMPLATE(ResizeUpsampleTrilinear, ResizeUpsampleTrilinearArch35Tiling, 2000);
393} // namespace optiling
Aimage/resize_upsample_trilinear/op_host/resize_upsample_trilinear_tiling_arch35.h+80-0
@@ -0,0 +1,80 @@
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/*!
12 * \file resize_upsample_trilinear_tiling_arch35.h
13 * \brief resize_upsample_trilinear tiling for A5 architecture
14 */
15#ifndef RESIZE_UPSAMPLE_TRILINEAR_TILING_ARCH35_H_
16#define RESIZE_UPSAMPLE_TRILINEAR_TILING_ARCH35_H_
17 
18#include "op_host/tiling_base.h"
19#include "op_host/tiling_templates_registry.h"
20#include "register/op_impl_registry.h"
21#include "tiling/platform/platform_ascendc.h"
22 
23namespace optiling {
24 
25class ResizeUpsampleTrilinearArch35Tiling : public Ops::Cv::OpTiling::TilingBaseClass {
26public:
27 explicit ResizeUpsampleTrilinearArch35Tiling(gert::TilingContext* context) : TilingBaseClass(context) {}
28 
29protected:
30 bool IsCapable() override;
31 ge::graphStatus GetPlatformInfo() override;
32 ge::graphStatus GetShapeAttrsInfo() override;
33 ge::graphStatus DoOpTiling() override;
34 ge::graphStatus DoLibApiTiling() override;
35 uint64_t GetTilingKey() const override;
36 ge::graphStatus GetWorkspaceSize() override;
37 ge::graphStatus PostTiling() override;
38 
39private:
40 int32_t coreNum_ = 0;
41 bool platformInfoCached_ = false;
42 int64_t inN_ = 0;
43 int64_t inC_ = 0;
44 int64_t outN_ = 0;
45 int64_t outC_ = 0;
46 int64_t batchCount_ = 0;
47 int64_t inputD_ = 0;
48 int64_t inputH_ = 0;
49 int64_t inputW_ = 0;
50 int64_t outputD_ = 0;
51 int64_t outputH_ = 0;
52 int64_t outputW_ = 0;
53 float scaleD_ = 0.0f;
54 float scaleH_ = 0.0f;
55 float scaleW_ = 0.0f;
56 int32_t alignCorners_ = 0;
57 uint64_t totalElements_ = 0;
58 uint32_t threadsPerBlock_ = 512;
59 uint32_t elementsPerThread_ = 1;
60 uint32_t blockCount_ = 1;
61 uint32_t usedCoreNum_ = 1;
62 uint32_t baseElementsPerBlock_ = 0;
63 uint32_t tailElements_ = 0;
64 uint64_t dtypeKey_ = 0;
65 int32_t useInt32_ = 0;
66 
67 ge::graphStatus ValidateAndGetInputShape();
68 ge::graphStatus ValidateAndGetOutputShape();
69 ge::graphStatus ValidateShapeValues();
70 ge::graphStatus ExtractAttrsAndComputeScales();
71 ge::graphStatus DetermineDtypeKey();
72 ge::graphStatus ValidateDimensionsAndComputeTotal();
73 ge::graphStatus HandleZeroElements();
74 ge::graphStatus ComputeThreadBlockConfig();
75 ge::graphStatus ComputeFinalTilingConfig();
76};
77 
78} // namespace optiling
79 
80#endif // RESIZE_UPSAMPLE_TRILINEAR_TILING_ARCH35_H_
Aimage/resize_upsample_trilinear/op_kernel/arch35/resize_upsample_trilinear_simt.h+98-0
@@ -0,0 +1,98 @@
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/*!
12 * \file resize_upsample_trilinear_simt.h
13 * \brief ResizeUpsampleTrilinear SIMT kernel implementation for A5
14 */
15 
16#ifndef RESIZE_UPSAMPLE_TRILINEAR_SIMT_H
17#define RESIZE_UPSAMPLE_TRILINEAR_SIMT_H
18 
19#include "kernel_operator.h"
20#include "kernel_tiling/kernel_tiling.h"
21#include "./resize_upsample_trilinear_simt_base.h"
22#include "./resize_upsample_trilinear_tiling_data.h"
23 
24namespace ResizeUpsampleTrilinear {
25using namespace AscendC;
26 
27template <typename T>
28class ResizeUpsampleTrilinearSimt {
29public:
30 __aicore__ inline ResizeUpsampleTrilinearSimt() = default;
31 __aicore__ inline void Init(GM_ADDR output, GM_ADDR input,
32 const ResizeUpsampleTrilinearArch35TilingData* __restrict tiling);
33 __aicore__ inline void Process();
34 
35private:
36 const ResizeUpsampleTrilinearArch35TilingData* tilingData;
37 uint32_t bid = 0;
38 GlobalTensor<T> outputGm;
39 GlobalTensor<T> inputGm;
40};
41 
42template <typename T>
43__aicore__ inline void ResizeUpsampleTrilinearSimt<T>::Init(
44 GM_ADDR output, GM_ADDR input, const ResizeUpsampleTrilinearArch35TilingData* __restrict tiling)
45{
46 outputGm.SetGlobalBuffer((__gm__ T*)output);
47 inputGm.SetGlobalBuffer((__gm__ T*)input);
48 tilingData = tiling;
49 bid = static_cast<uint32_t>(GetBlockIdx());
50}
51 
52template <typename T>
53__aicore__ inline void ResizeUpsampleTrilinearSimt<T>::Process()
54{
55 uint32_t baseElementsPerBlock = tilingData->base_elements_per_block;
56 uint32_t blockCount = tilingData->block_count;
57 uint32_t tailElements = tilingData->tail_elements;
58 
59 if (bid >= blockCount) {
60 return;
61 }
62 
63 __gm__ T* outputPtr = (__gm__ T*)outputGm.GetAddr();
64 __gm__ T* inputPtr = (__gm__ T*)inputGm.GetAddr();
65 
66 if (tilingData->use_int32 != 0) {
67 int32_t blkStartOffset = static_cast<int32_t>(bid) * static_cast<int32_t>(baseElementsPerBlock);
68 int32_t blkProcessNum;
69 if (bid == blockCount - 1) {
70 blkProcessNum = static_cast<int32_t>(tailElements);
71 } else {
72 blkProcessNum = static_cast<int32_t>(baseElementsPerBlock);
73 }
74 
75 if (blkProcessNum <= 0) {
76 return;
77 }
78 
79 calleeInt32<T, int32_t>(outputPtr, inputPtr, blkStartOffset, blkProcessNum, tilingData);
80 } else {
81 int64_t blkStartOffset = static_cast<int64_t>(bid) * static_cast<int64_t>(baseElementsPerBlock);
82 int64_t blkProcessNum;
83 if (bid == blockCount - 1) {
84 blkProcessNum = static_cast<int64_t>(tailElements);
85 } else {
86 blkProcessNum = static_cast<int64_t>(baseElementsPerBlock);
87 }
88 
89 if (blkProcessNum <= 0) {
90 return;
91 }
92 
93 calleeInt64<T, int64_t>(outputPtr, inputPtr, blkStartOffset, blkProcessNum, tilingData);
94 }
95}
96} // namespace ResizeUpsampleTrilinear
97 
98#endif // RESIZE_UPSAMPLE_TRILINEAR_SIMT_H
Aimage/resize_upsample_trilinear/op_kernel/arch35/resize_upsample_trilinear_simt_base.h+385-0
@@ -0,0 +1,385 @@
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/*!
12 * \file resize_upsample_trilinear_simt_base.h
13 * \brief SIMT compute base functions for ResizeUpsampleTrilinear
14 */
15#ifndef RESIZE_UPSAMPLE_TRILINEAR_SIMT_BASE_H
16#define RESIZE_UPSAMPLE_TRILINEAR_SIMT_BASE_H
17 
18#include "kernel_operator.h"
19#include "kernel_tiling/kernel_tiling.h"
20#include "./resize_upsample_trilinear_tiling_data.h"
21 
22namespace ResizeUpsampleTrilinear {
23using namespace AscendC;
24 
25const int32_t THREAD_NUM_B32 = 512;
26const int32_t THREAD_NUM_B64 = 512;
27 
28static __simt_callee__ __aicore__ __attribute__((always_inline)) inline float CalcSourceCoord(
29 int outIdx, float scale, bool alignCorners)
30{
31 if (alignCorners) {
32 return static_cast<float>(outIdx) * scale;
33 }
34 return fmaxf((static_cast<float>(outIdx) + 0.5f) * scale - 0.5f, 0.0f);
35}
36 
37static __simt_callee__ __aicore__ __attribute__((always_inline)) inline int ClampIndex(
38 int idx, int maxVal)
39{
40 return (idx < 0) ? 0 : (idx > maxVal) ? maxVal : idx;
41}
42 
43template <typename T2>
44struct DepthInterpCoeffs {
45 int x0;
46 int x1;
47 float lambda0;
48 float lambda1;
49 T2 addr_d0;
50 T2 addr_d1;
51};
52 
53template <typename T2>
54static __simt_callee__ __aicore__ __attribute__((always_inline)) inline DepthInterpCoeffs<T2> CalcDepthInterp(
55 float src_d, int maxD, T2 input_base, T2 stride_d_input)
56{
57 DepthInterpCoeffs<T2> c;
58 c.x0 = ClampIndex(static_cast<int>(floorf(src_d)), maxD);
59 c.x1 = (c.x0 + 1 > maxD) ? maxD : c.x0 + 1;
60 c.lambda0 = static_cast<float>(c.x1) - src_d;
61 c.lambda1 = 1.0f - c.lambda0;
62 c.addr_d0 = input_base + static_cast<T2>(c.x0) * stride_d_input;
63 c.addr_d1 = input_base + static_cast<T2>(c.x1) * stride_d_input;
64 return c;
65}
66 
67template <typename T2>
68struct HeightInterpCoeffs {
69 int y0;
70 int y1;
71 float mu0;
72 float mu1;
73 T2 addr_d0_h0;
74 T2 addr_d0_h1;
75 T2 addr_d1_h0;
76 T2 addr_d1_h1;
77};
78 
79template <typename T2>
80static __simt_callee__ __aicore__ __attribute__((always_inline)) inline HeightInterpCoeffs<T2> CalcHeightInterp(
81 float src_h, int maxH, T2 addr_d0, T2 addr_d1, T2 stride_h_input)
82{
83 HeightInterpCoeffs<T2> c;
84 c.y0 = ClampIndex(static_cast<int>(floorf(src_h)), maxH);
85 c.y1 = (c.y0 + 1 > maxH) ? maxH : c.y0 + 1;
86 c.mu0 = static_cast<float>(c.y1) - src_h;
87 c.mu1 = 1.0f - c.mu0;
88 c.addr_d0_h0 = addr_d0 + static_cast<T2>(c.y0) * stride_h_input;
89 c.addr_d0_h1 = addr_d0 + static_cast<T2>(c.y1) * stride_h_input;
90 c.addr_d1_h0 = addr_d1 + static_cast<T2>(c.y0) * stride_h_input;
91 c.addr_d1_h1 = addr_d1 + static_cast<T2>(c.y1) * stride_h_input;
92 return c;
93}
94 
95struct CombinedCoeffs {
96 float lambda0_mu0;
97 float lambda0_mu1;
98 float lambda1_mu0;
99 float lambda1_mu1;
100};
101 
102static __simt_callee__ __aicore__ __attribute__((always_inline)) inline CombinedCoeffs CalcCombinedCoeffs(
103 float lambda0, float lambda1, float mu0, float mu1)
104{
105 CombinedCoeffs c;
106 c.lambda0_mu0 = lambda0 * mu0;
107 c.lambda0_mu1 = lambda0 * mu1;
108 c.lambda1_mu0 = lambda1 * mu0;
109 c.lambda1_mu1 = lambda1 * mu1;
110 return c;
111}
112 
113template <typename T2>
114struct SimtComputeParams {
115 T2 input_d;
116 T2 input_h;
117 T2 input_w;
118 T2 output_d;
119 T2 output_h;
120 T2 output_w;
121 float scale_d;
122 float scale_h;
123 float scale_w;
124 bool alignCorners;
125 float half_scale_w;
126 int maxD;
127 int maxH;
128 int maxW;
129 T2 stride_bc_input;
130 T2 stride_d_input;
131 T2 stride_h_input;
132 T2 stride_bc_output;
133 T2 stride_d_output;
134 T2 stride_h_output;
135 T2 plane_size;
136 T2 row_size;
137};
138 
139template <typename T2>
140static __simt_callee__ __aicore__ __attribute__((always_inline)) inline SimtComputeParams<T2> InitSimtParams(
141 const ResizeUpsampleTrilinearArch35TilingData* __restrict tilingData)
142{
143 SimtComputeParams<T2> p;
144 p.input_d = static_cast<T2>(tilingData->input_d);
145 p.input_h = static_cast<T2>(tilingData->input_h);
146 p.input_w = static_cast<T2>(tilingData->input_w);
147 p.output_d = static_cast<T2>(tilingData->output_d);
148 p.output_h = static_cast<T2>(tilingData->output_h);
149 p.output_w = static_cast<T2>(tilingData->output_w);
150 p.scale_d = tilingData->scale_d;
151 p.scale_h = tilingData->scale_h;
152 p.scale_w = tilingData->scale_w;
153 p.alignCorners = (tilingData->align_corners != 0);
154 p.half_scale_w = 0.5f * p.scale_w;
155 p.maxD = static_cast<int>(p.input_d - 1);
156 p.maxH = static_cast<int>(p.input_h - 1);
157 p.maxW = static_cast<int>(p.input_w - 1);
158 p.stride_bc_input = p.input_d * p.input_h * p.input_w;
159 p.stride_d_input = p.input_h * p.input_w;
160 p.stride_h_input = p.input_w;
161 p.stride_bc_output = p.output_d * p.output_h * p.output_w;
162 p.stride_d_output = p.output_h * p.output_w;
163 p.stride_h_output = p.output_w;
164 p.plane_size = p.output_d * p.output_h * p.output_w;
165 p.row_size = p.output_h * p.output_w;
166 return p;
167}
168 
169template <typename T2>
170struct SimtThreadState {
171 T2 bc;
172 int od;
173 int oh;
174 int ow;
175 T2 flat_idx;
176 T2 input_base;
177 T2 output_base;
178 float src_d;
179 float src_h;
180 DepthInterpCoeffs<T2> depthCoeffs;
181 HeightInterpCoeffs<T2> heightCoeffs;
182 CombinedCoeffs combinedCoeffs;
183};
184 
185template <typename T2>
186static __simt_callee__ __aicore__ __attribute__((always_inline)) inline SimtThreadState<T2> InitThreadState(
187 T2 blkStartOffset, T2 threadOffset, const SimtComputeParams<T2>& p,
188 const ResizeUpsampleTrilinearArch35TilingData* __restrict tilingData)
189{
190 SimtThreadState<T2> s;
191 s.flat_idx = blkStartOffset + threadOffset;
192 s.bc = s.flat_idx / p.plane_size;
193 T2 spatial_idx = s.flat_idx - s.bc * p.plane_size;
194 s.od = static_cast<int>(spatial_idx / p.row_size);
195 T2 hw_idx = spatial_idx - static_cast<T2>(s.od) * p.row_size;
196 s.oh = static_cast<int>(hw_idx / p.output_w);
197 s.ow = static_cast<int>(hw_idx - static_cast<T2>(s.oh) * p.output_w);
198 s.input_base = s.bc * p.stride_bc_input;
199 s.output_base = s.bc * p.stride_bc_output + static_cast<T2>(s.od) * p.stride_d_output +
200 static_cast<T2>(s.oh) * p.stride_h_output;
201 s.src_d = CalcSourceCoord(s.od, p.scale_d, p.alignCorners);
202 s.src_h = CalcSourceCoord(s.oh, p.scale_h, p.alignCorners);
203 s.depthCoeffs = CalcDepthInterp<T2>(s.src_d, p.maxD, s.input_base, p.stride_d_input);
204 s.heightCoeffs = CalcHeightInterp<T2>(s.src_h, p.maxH, s.depthCoeffs.addr_d0, s.depthCoeffs.addr_d1, p.stride_h_input);
205 s.combinedCoeffs = CalcCombinedCoeffs(s.depthCoeffs.lambda0, s.depthCoeffs.lambda1,
206 s.heightCoeffs.mu0, s.heightCoeffs.mu1);
207 return s;
208}
209 
210template <typename T1, typename T2>
211static __simt_callee__ __aicore__ __attribute__((always_inline)) inline float ComputeTrilinearValue(
212 __gm__ T1* input, const HeightInterpCoeffs<T2>& hc, int z0, int z1,
213 float nu0, float nu1, const CombinedCoeffs& cc)
214{
215 T1 raw_v000 = asc_ldcg(&input[hc.addr_d0_h0 + z0]);
216 T1 raw_v001 = asc_ldcg(&input[hc.addr_d0_h0 + z1]);
217 T1 raw_v010 = asc_ldcg(&input[hc.addr_d0_h1 + z0]);
218 T1 raw_v011 = asc_ldcg(&input[hc.addr_d0_h1 + z1]);
219 T1 raw_v100 = asc_ldcg(&input[hc.addr_d1_h0 + z0]);
220 T1 raw_v101 = asc_ldcg(&input[hc.addr_d1_h0 + z1]);
221 T1 raw_v110 = asc_ldcg(&input[hc.addr_d1_h1 + z0]);
222 T1 raw_v111 = asc_ldcg(&input[hc.addr_d1_h1 + z1]);
223 
224 float v000, v001, v010, v011, v100, v101, v110, v111;
225 if constexpr (std::is_same<T1, half>::value || std::is_same<T1, bfloat16_t>::value) {
226 v000 = static_cast<float>(raw_v000); v001 = static_cast<float>(raw_v001);
227 v010 = static_cast<float>(raw_v010); v011 = static_cast<float>(raw_v011);
228 v100 = static_cast<float>(raw_v100); v101 = static_cast<float>(raw_v101);
229 v110 = static_cast<float>(raw_v110); v111 = static_cast<float>(raw_v111);
230 } else {
231 v000 = raw_v000; v001 = raw_v001;
232 v010 = raw_v010; v011 = raw_v011;
233 v100 = raw_v100; v101 = raw_v101;
234 v110 = raw_v110; v111 = raw_v111;
235 }
236 
237 return v000 * cc.lambda0_mu0 * nu0 + v001 * cc.lambda0_mu0 * nu1
238 + v010 * cc.lambda0_mu1 * nu0 + v011 * cc.lambda0_mu1 * nu1
239 + v100 * cc.lambda1_mu0 * nu0 + v101 * cc.lambda1_mu0 * nu1
240 + v110 * cc.lambda1_mu1 * nu0 + v111 * cc.lambda1_mu1 * nu1;
241}
242 
243template <typename T1, typename T2>
244static __simt_callee__ __aicore__ __attribute__((always_inline)) inline void StoreOutputValue(
245 __gm__ T1* output, T2 output_addr, float result)
246{
247 if constexpr (std::is_same<T1, half>::value || std::is_same<T1, bfloat16_t>::value) {
248 T1 output_val = static_cast<T1>(result);
249 asc_stcg(&output[output_addr], output_val);
250 } else {
251 asc_stcg(&output[output_addr], result);
252 }
253}
254 
255template <typename T2>
256static __simt_callee__ __aicore__ __attribute__((always_inline)) inline void UpdateCoordsOnNewBatch(
257 SimtThreadState<T2>& s, const SimtComputeParams<T2>& p)
258{
259 s.input_base = s.bc * p.stride_bc_input;
260 s.output_base = s.bc * p.stride_bc_output;
261 s.src_d = CalcSourceCoord(0, p.scale_d, p.alignCorners);
262 s.src_h = CalcSourceCoord(0, p.scale_h, p.alignCorners);
263 s.depthCoeffs = CalcDepthInterp<T2>(s.src_d, p.maxD, s.input_base, p.stride_d_input);
264 s.heightCoeffs = CalcHeightInterp<T2>(s.src_h, p.maxH, s.depthCoeffs.addr_d0, s.depthCoeffs.addr_d1, p.stride_h_input);
265 s.combinedCoeffs = CalcCombinedCoeffs(s.depthCoeffs.lambda0, s.depthCoeffs.lambda1,
266 s.heightCoeffs.mu0, s.heightCoeffs.mu1);
267}
268 
269template <typename T2>
270static __simt_callee__ __aicore__ __attribute__((always_inline)) inline void UpdateCoordsOnNewDepth(
271 SimtThreadState<T2>& s, const SimtComputeParams<T2>& p)
272{
273 s.src_d = CalcSourceCoord(s.od, p.scale_d, p.alignCorners);
274 s.depthCoeffs = CalcDepthInterp<T2>(s.src_d, p.maxD, s.input_base, p.stride_d_input);
275 s.heightCoeffs = CalcHeightInterp<T2>(s.src_h, p.maxH, s.depthCoeffs.addr_d0, s.depthCoeffs.addr_d1, p.stride_h_input);
276 s.combinedCoeffs = CalcCombinedCoeffs(s.depthCoeffs.lambda0, s.depthCoeffs.lambda1,
277 s.heightCoeffs.mu0, s.heightCoeffs.mu1);
278 s.output_base = s.bc * p.stride_bc_output + static_cast<T2>(s.od) * p.stride_d_output;
279}
280 
281template <typename T2>
282static __simt_callee__ __aicore__ __attribute__((always_inline)) inline void UpdateCoordsOnNewHeight(
283 SimtThreadState<T2>& s, const SimtComputeParams<T2>& p)
284{
285 s.src_h = CalcSourceCoord(s.oh, p.scale_h, p.alignCorners);
286 s.heightCoeffs = CalcHeightInterp<T2>(s.src_h, p.maxH, s.depthCoeffs.addr_d0, s.depthCoeffs.addr_d1, p.stride_h_input);
287 s.combinedCoeffs = CalcCombinedCoeffs(s.depthCoeffs.lambda0, s.depthCoeffs.lambda1,
288 s.heightCoeffs.mu0, s.heightCoeffs.mu1);
289 s.output_base = s.bc * p.stride_bc_output + static_cast<T2>(s.od) * p.stride_d_output +
290 static_cast<T2>(s.oh) * p.stride_h_output;
291}
292 
293template <typename T2>
294static __simt_callee__ __aicore__ __attribute__((always_inline)) inline void AdvanceOutputCoords(
295 SimtThreadState<T2>& s, const SimtComputeParams<T2>& p)
296{
297 s.ow++;
298 s.flat_idx++;
299 if (s.ow >= p.output_w) {
300 s.ow = 0;
301 s.oh++;
302 if (s.oh >= p.output_h) {
303 s.oh = 0;
304 s.od++;
305 if (s.od >= p.output_d) {
306 s.od = 0;
307 s.bc++;
308 UpdateCoordsOnNewBatch(s, p);
309 } else {
310 UpdateCoordsOnNewDepth(s, p);
311 }
312 } else {
313 UpdateCoordsOnNewHeight(s, p);
314 }
315 }
316}
317 
318template <typename T1, typename T2>
319__simt_callee__ __aicore__ __attribute__((always_inline)) inline void SimtCompute(
320 __gm__ T1* output, __gm__ T1* input, T2 blkStartOffset, T2 blkProcessNum,
321 const ResizeUpsampleTrilinearArch35TilingData* __restrict tilingData)
322{
323 SimtComputeParams<T2> p = InitSimtParams<T2>(tilingData);
324 uint32_t elements_per_thread = tilingData->elements_per_thread;
325 
326 T2 tid = static_cast<T2>(threadIdx.x);
327 T2 threadOffset = tid * static_cast<T2>(elements_per_thread);
328 
329 if (threadOffset >= blkProcessNum) {
330 return;
331 }
332 
333 T2 elementsToProcess = static_cast<T2>(elements_per_thread);
334 T2 remaining = blkProcessNum - threadOffset;
335 if (remaining < elementsToProcess) {
336 elementsToProcess = remaining;
337 }
338 
339 SimtThreadState<T2> s = InitThreadState<T2>(blkStartOffset, threadOffset, p, tilingData);
340 if (s.flat_idx >= static_cast<T2>(tilingData->total_elements)) {
341 return;
342 }
343 
344 for (T2 e = 0; e < elementsToProcess; e++) {
345 if (s.flat_idx >= static_cast<T2>(tilingData->total_elements)) {
346 break;
347 }
348 
349 float src_w;
350 if (p.alignCorners) {
351 src_w = static_cast<float>(s.ow) * p.scale_w;
352 } else {
353 src_w = fmaxf(static_cast<float>(s.ow) * p.scale_w + p.half_scale_w - 0.5f, 0.0f);
354 }
355 
356 int z0 = ClampIndex(static_cast<int>(floorf(src_w)), p.maxW);
357 int z1 = (z0 + 1 > p.maxW) ? p.maxW : z0 + 1;
358 float nu0 = static_cast<float>(z1) - src_w;
359 float nu1 = 1.0f - nu0;
360 
361 float result = ComputeTrilinearValue<T1, T2>(input, s.heightCoeffs, z0, z1, nu0, nu1, s.combinedCoeffs);
362 T2 output_addr = s.output_base + static_cast<T2>(s.ow);
363 StoreOutputValue<T1, T2>(output, output_addr, result);
364 
365 AdvanceOutputCoords(s, p);
366 }
367}
368 
369template <typename T1, typename T2>
370__simt_vf__ LAUNCH_BOUND(THREAD_NUM_B32) __aicore__ void calleeInt32(
371 __gm__ T1* output, __gm__ T1* input, T2 blkStartOffset, T2 blkProcessNum,
372 const ResizeUpsampleTrilinearArch35TilingData* __restrict tilingData)
373{
374 SimtCompute<T1, T2>(output, input, blkStartOffset, blkProcessNum, tilingData);
375}
376 
377template <typename T1, typename T2>
378__simt_vf__ LAUNCH_BOUND(THREAD_NUM_B64) __aicore__ void calleeInt64(
379 __gm__ T1* output, __gm__ T1* input, T2 blkStartOffset, T2 blkProcessNum,
380 const ResizeUpsampleTrilinearArch35TilingData* __restrict tilingData)
381{
382 SimtCompute<T1, T2>(output, input, blkStartOffset, blkProcessNum, tilingData);
383}
384} // namespace ResizeUpsampleTrilinear
385#endif // RESIZE_UPSAMPLE_TRILINEAR_SIMT_BASE_H
Aimage/resize_upsample_trilinear/op_kernel/arch35/resize_upsample_trilinear_tiling_data.h+39-0
@@ -0,0 +1,39 @@
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/*!
12 * \file resize_upsample_trilinear_tiling_data.h
13 * \brief tiling data struct for arch35
14 */
15 
16#ifndef RESIZE_UPSAMPLE_TRILINEAR_ARCH35_TILING_DATA_H
17#define RESIZE_UPSAMPLE_TRILINEAR_ARCH35_TILING_DATA_H
18 
19struct ResizeUpsampleTrilinearArch35TilingData {
20 uint32_t elements_per_thread;
21 uint32_t block_count;
22 uint32_t used_core_num;
23 uint32_t base_elements_per_block;
24 uint32_t tail_elements;
25 uint64_t total_elements;
26 int64_t batch_count;
27 int64_t input_d;
28 int64_t input_h;
29 int64_t input_w;
30 int64_t output_d;
31 int64_t output_h;
32 int64_t output_w;
33 float scale_d;
34 float scale_h;
35 float scale_w;
36 int32_t align_corners;
37 int32_t use_int32;
38};
39#endif
Aimage/resize_upsample_trilinear/op_kernel/arch35/resize_upsample_trilinear_tiling_key.h+33-0
@@ -0,0 +1,33 @@
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/*!
12 * \file resize_upsample_trilinear_tiling_key.h
13 * \brief resize_upsample_trilinear tiling key declare for arch35
14 */
15 
16#ifndef RESIZE_UPSAMPLE_TRILINEAR_ARCH35_TILING_KEY_H
17#define RESIZE_UPSAMPLE_TRILINEAR_ARCH35_TILING_KEY_H
18 
19#include "ascendc/host_api/tiling/template_argument.h"
20 
21#define TPL_DTYPE_FP32 0
22#define TPL_DTYPE_FP16 1
23#define TPL_DTYPE_BF16 2
24 
25ASCENDC_TPL_ARGS_DECL(ResizeUpsampleTrilinear,
26 ASCENDC_TPL_UINT_DECL(dtypeKey, ASCENDC_TPL_8_BW, ASCENDC_TPL_UI_LIST, TPL_DTYPE_FP32, TPL_DTYPE_FP16, TPL_DTYPE_BF16)
27);
28 
29ASCENDC_TPL_SEL(
30 ASCENDC_TPL_ARGS_SEL(
31 ASCENDC_TPL_UINT_SEL(dtypeKey, ASCENDC_TPL_UI_RANGE, 1, TPL_DTYPE_FP32, TPL_DTYPE_BF16)));
32 
33#endif // RESIZE_UPSAMPLE_TRILINEAR_ARCH35_TILING_KEY_H
Aimage/resize_upsample_trilinear/op_kernel/resize_upsample_trilinear_apt.cpp+43-0
@@ -0,0 +1,43 @@
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/*!
12 * \file resize_upsample_trilinear_apt.cpp
13 * \brief ResizeUpsampleTrilinear APT kernel entry function for A5
14 */
15#include "./arch35/resize_upsample_trilinear_tiling_key.h"
16#include "./arch35/resize_upsample_trilinear_tiling_data.h"
17#include "./arch35/resize_upsample_trilinear_simt.h"
18#include "./arch35/resize_upsample_trilinear_simt_base.h"
19#include "kernel_operator.h"
20#include "kernel_tiling/kernel_tiling.h"
21 
22template <uint32_t dtypeKey>
23__global__ __aicore__ void resize_upsample_trilinear(GM_ADDR output, GM_ADDR input, GM_ADDR workspace, GM_ADDR tiling)
24{
25 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
26 REGISTER_TILING_DEFAULT(ResizeUpsampleTrilinearArch35TilingData);
27 GET_TILING_DATA_WITH_STRUCT(ResizeUpsampleTrilinearArch35TilingData, tilingData, tiling);
28 AscendC::InitSocState();
29 
30 if (TILING_KEY_IS(TPL_DTYPE_FP32)) {
31 ResizeUpsampleTrilinear::ResizeUpsampleTrilinearSimt<float> op;
32 op.Init(output, input, &tilingData);
33 op.Process();
34 } else if (TILING_KEY_IS(TPL_DTYPE_FP16)) {
35 ResizeUpsampleTrilinear::ResizeUpsampleTrilinearSimt<half> op;
36 op.Init(output, input, &tilingData);
37 op.Process();
38 } else if (TILING_KEY_IS(TPL_DTYPE_BF16)) {
39 ResizeUpsampleTrilinear::ResizeUpsampleTrilinearSimt<bfloat16_t> op;
40 op.Init(output, input, &tilingData);
41 op.Process();
42 }
43}
Aimage/resize_upsample_trilinear/tests/ut/op_host/arch35/test_resize_upsample_trilinear_tiling_arch35.cpp+284-0
@@ -0,0 +1,284 @@
1/**
2 * Copyright (c) 2025-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 <vector>
14#include "../../../../op_host/resize_upsample_trilinear_tiling_arch35.h"
15#include "tiling_context_faker.h"
16#include "tiling_case_executor.h"
17 
18class ResizeUpsampleTrilinearArch35TilingTest : public testing::Test {
19protected:
20 static void SetUpTestCase()
21 {
22 std::cout << "ResizeUpsampleTrilinearArch35TilingTest SetUp" << std::endl;
23 }
24 
25 static void TearDownTestCase()
26 {
27 std::cout << "ResizeUpsampleTrilinearArch35TilingTest TearDown" << std::endl;
28 }
29};
30 
31struct ResizeUpsampleTrilinearArch35CompileInfo {
32 int32_t coreNum = 0;
33 int32_t socVersionType = 0;
34};
35 
36TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_fp32_output_size_case01)
37{
38 gert::StorageShape inputShape = {{1, 1, 4, 128, 128}, {1, 1, 4, 128, 128}};
39 gert::StorageShape outShape = {{1, 1, 12, 256, 256}, {1, 1, 12, 256, 256}};
40 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
41 std::vector<int64_t> output_size = {12, 256, 256};
42 gert::TilingContextPara tilingContextPara(
43 "ResizeUpsampleTrilinear",
44 {{inputShape, ge::DT_FLOAT, ge::FORMAT_ND}},
45 {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}},
46 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
47 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
48 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
49 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
50 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
51 &compileInfo,
52 "Ascend950",
53 64,
54 262144);
55 uint64_t expectTilingKey = 0;
56 string expectTilingData = "3298534883330 4398046511168 1024 786432 1 4 128 128 12 256 256 4539628425440832171 1056964608 ";
57 std::vector<size_t> expectWorkspaces = {16777216};
58 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
59}
60 
61TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_fp16_output_size_case02)
62{
63 gert::StorageShape inputShape = {{1, 2, 8, 64, 64}, {1, 2, 8, 64, 64}};
64 gert::StorageShape outShape = {{1, 2, 16, 128, 128}, {1, 2, 16, 128, 128}};
65 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
66 std::vector<int64_t> output_size = {16, 128, 128};
67 gert::TilingContextPara tilingContextPara(
68 "ResizeUpsampleTrilinear",
69 {{inputShape, ge::DT_FLOAT16, ge::FORMAT_ND}},
70 {{outShape, ge::DT_FLOAT16, ge::FORMAT_ND}},
71 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
72 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
73 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
74 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
75 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
76 &compileInfo,
77 "Ascend950",
78 64,
79 262144);
80 uint64_t expectTilingKey = 1;
81 string expectTilingData = "4398046511105 2199023255616 512 524288 2 8 64 64 16 128 128 4539628425446424576 1056964608 ";
82 std::vector<size_t> expectWorkspaces = {16777216};
83 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
84}
85 
86TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_bf16_output_size_case03)
87{
88 gert::StorageShape inputShape = {{2, 3, 10, 48, 48}, {2, 3, 10, 48, 48}};
89 gert::StorageShape outShape = {{2, 3, 20, 96, 96}, {2, 3, 20, 96, 96}};
90 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
91 std::vector<int64_t> output_size = {20, 96, 96};
92 gert::TilingContextPara tilingContextPara(
93 "ResizeUpsampleTrilinear",
94 {{inputShape, ge::DT_BF16, ge::FORMAT_ND}},
95 {{outShape, ge::DT_BF16, ge::FORMAT_ND}},
96 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
97 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
98 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
99 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
100 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
101 &compileInfo,
102 "Ascend950",
103 64,
104 262144);
105 uint64_t expectTilingKey = 2;
106 string expectTilingData = "3092376453123 6597069766720 1536 1105920 6 10 48 48 20 96 96 4539628425446424576 1056964608 ";
107 std::vector<size_t> expectWorkspaces = {16777216};
108 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
109}
110 
111TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_fp32_align_corners_case04)
112{
113 gert::StorageShape inputShape = {{1, 1, 4, 4, 4}, {1, 1, 4, 4, 4}};
114 gert::StorageShape outShape = {{1, 1, 8, 8, 8}, {1, 1, 8, 8, 8}};
115 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
116 std::vector<int64_t> output_size = {8, 8, 8};
117 gert::TilingContextPara tilingContextPara(
118 "ResizeUpsampleTrilinear",
119 {{inputShape, ge::DT_FLOAT, ge::FORMAT_ND}},
120 {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}},
121 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
122 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(true)),
123 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
124 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
125 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
126 &compileInfo,
127 "Ascend950",
128 64,
129 262144);
130 uint64_t expectTilingKey = 0;
131 string expectTilingData = "4294967297 2199023255553 512 512 1 4 4 4 8 8 8 4529334484052176311 5349535159 ";
132 std::vector<size_t> expectWorkspaces = {16777216};
133 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
134}
135 
136TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_fp32_scales_case05)
137{
138 gert::StorageShape inputShape = {{1, 1, 4, 4, 4}, {1, 1, 4, 4, 4}};
139 gert::StorageShape outShape = {{1, 1, 8, 8, 8}, {1, 1, 8, 8, 8}};
140 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
141 std::vector<int64_t> output_size = {0, 0, 0};
142 gert::TilingContextPara tilingContextPara(
143 "ResizeUpsampleTrilinear",
144 {{inputShape, ge::DT_FLOAT, ge::FORMAT_ND}},
145 {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}},
146 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
147 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
148 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(2.0)),
149 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(2.0)),
150 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(2.0))},
151 &compileInfo,
152 "Ascend950",
153 64,
154 262144);
155 uint64_t expectTilingKey = 0;
156 string expectTilingData = "4294967297 2199023255553 512 512 1 4 4 4 8 8 8 4539628425446424576 1056964608 ";
157 std::vector<size_t> expectWorkspaces = {16777216};
158 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
159}
160 
161TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_fp32_small_shape_case06)
162{
163 gert::StorageShape inputShape = {{1, 1, 2, 2, 2}, {1, 1, 2, 2, 2}};
164 gert::StorageShape outShape = {{1, 1, 3, 3, 3}, {1, 1, 3, 3, 3}};
165 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
166 std::vector<int64_t> output_size = {3, 3, 3};
167 gert::TilingContextPara tilingContextPara(
168 "ResizeUpsampleTrilinear",
169 {{inputShape, ge::DT_FLOAT, ge::FORMAT_ND}},
170 {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}},
171 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
172 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
173 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
174 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
175 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
176 &compileInfo,
177 "Ascend950",
178 64,
179 262144);
180 uint64_t expectTilingKey = 0;
181 string expectTilingData = "4294967297 137438953473 27 27 1 2 2 2 3 3 3 4551638025887197867 1059760811 ";
182 std::vector<size_t> expectWorkspaces = {16777216};
183 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
184}
185 
186TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_fp32_downsample_case07)
187{
188 gert::StorageShape inputShape = {{1, 2, 60, 256, 512}, {1, 2, 60, 256, 512}};
189 gert::StorageShape outShape = {{1, 2, 30, 128, 256}, {1, 2, 30, 128, 256}};
190 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
191 std::vector<int64_t> output_size = {30, 128, 256};
192 gert::TilingContextPara tilingContextPara(
193 "ResizeUpsampleTrilinear",
194 {{inputShape, ge::DT_FLOAT, ge::FORMAT_ND}},
195 {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}},
196 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
197 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
198 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
199 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
200 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
201 &compileInfo,
202 "Ascend950",
203 64,
204 262144);
205 uint64_t expectTilingKey = 0;
206 string expectTilingData = "4123168604164 8796093022272 2048 1966080 2 60 256 512 30 128 256 4611686019501129728 1073741824 ";
207 std::vector<size_t> expectWorkspaces = {16777216};
208 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
209}
210 
211TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_fp32_large_shape_case08)
212{
213 gert::StorageShape inputShape = {{4, 8, 32, 64, 64}, {4, 8, 32, 64, 64}};
214 gert::StorageShape outShape = {{4, 8, 64, 128, 128}, {4, 8, 64, 128, 128}};
215 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
216 std::vector<int64_t> output_size = {64, 128, 128};
217 gert::TilingContextPara tilingContextPara(
218 "ResizeUpsampleTrilinear",
219 {{inputShape, ge::DT_FLOAT, ge::FORMAT_ND}},
220 {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}},
221 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
222 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
223 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
224 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
225 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
226 &compileInfo,
227 "Ascend950",
228 64,
229 262144);
230 uint64_t expectTilingKey = 0;
231 string expectTilingData = "4398046511168 140737488355392 32768 33554432 32 32 64 64 64 128 128 4539628425446424576 1056964608 ";
232 std::vector<size_t> expectWorkspaces = {16777216};
233 ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
234}
235 
236TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_fp32_scale_exceed_max_case09)
237{
238 gert::StorageShape inputShape = {{1, 1, 2, 2, 2}, {1, 1, 2, 2, 2}};
239 gert::StorageShape outShape = {{1, 1, 101, 101, 101}, {1, 1, 101, 101, 101}};
240 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
241 std::vector<int64_t> output_size = {101, 101, 101};
242 gert::TilingContextPara tilingContextPara(
243 "ResizeUpsampleTrilinear",
244 {{inputShape, ge::DT_FLOAT, ge::FORMAT_ND}},
245 {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}},
246 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
247 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
248 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
249 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
250 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
251 &compileInfo,
252 "Ascend950",
253 64,
254 262144);
255 uint64_t expectTilingKey = 0;
256 string expectTilingData = "";
257 std::vector<size_t> expectWorkspaces = {16777216};
258 ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);
259}
260 
261TEST_F(ResizeUpsampleTrilinearArch35TilingTest, upsample_trilinear3d_nc_mismatch_case10)
262{
263 gert::StorageShape inputShape = {{1, 2, 4, 4, 4}, {1, 2, 4, 4, 4}};
264 gert::StorageShape outShape = {{2, 2, 8, 8, 8}, {2, 2, 8, 8, 8}};
265 ResizeUpsampleTrilinearArch35CompileInfo compileInfo = {64, 3510};
266 std::vector<int64_t> output_size = {8, 8, 8};
267 gert::TilingContextPara tilingContextPara(
268 "ResizeUpsampleTrilinear",
269 {{inputShape, ge::DT_FLOAT, ge::FORMAT_ND}},
270 {{outShape, ge::DT_FLOAT, ge::FORMAT_ND}},
271 {gert::TilingContextPara::OpAttr("output_size", Ops::Cv::AnyValue::CreateFrom<std::vector<int64_t>>(output_size)),
272 gert::TilingContextPara::OpAttr("align_corners", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
273 gert::TilingContextPara::OpAttr("scales_d", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
274 gert::TilingContextPara::OpAttr("scales_h", Ops::Cv::AnyValue::CreateFrom<float>(0.0)),
275 gert::TilingContextPara::OpAttr("scales_w", Ops::Cv::AnyValue::CreateFrom<float>(0.0))},
276 &compileInfo,
277 "Ascend950",
278 64,
279 262144);
280 uint64_t expectTilingKey = 0;
281 string expectTilingData = "";
282 std::vector<size_t> expectWorkspaces = {16777216};
283 ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);
284}
Mscripts/kernel/binary_config/ascendc_config.json+1-1
@@ -233,7 +233,7 @@
233 {"name":"UpsampleBicubic2d", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend310b", "kirinx90", "kirin9030"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},233 {"name":"UpsampleBicubic2d", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend310b", "kirinx90", "kirin9030"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},
234 {"name":"UpsampleBicubic2dAA", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "kirinx90", "kirin9030"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},234 {"name":"UpsampleBicubic2dAA", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "kirinx90", "kirin9030"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
235 {"name":"UpsampleBicubic2dGrad", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},235 {"name":"UpsampleBicubic2dGrad", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},
236 {"name":"ResizeUpsampleTrilinear", "compute_units": ["ascend910_93", "ascend910b", "ascend310p"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},236 {"name":"ResizeUpsampleTrilinear", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend950"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
237 {"name":"CTCLossV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},237 {"name":"CTCLossV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
238 {"name":"UpsampleTrilinear3dBackward", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},238 {"name":"UpsampleTrilinear3dBackward", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync" : true, "compile_options": {"ascend950": ["-DENABLE_CV_COMM_VIA_SSBUF=true"]}},
239 {"name":"UpsampleNearest3d", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend950", "kirinx90", "kirin9030"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},239 {"name":"UpsampleNearest3d", "compute_units": ["ascend910_93", "ascend910b", "ascend310p", "ascend950", "kirinx90", "kirin9030"], "auto_sync" : {"ascend910_93": true, "ascend910b": true, "ascend310p": true, "ascend950": false}, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},