Skip to content

Commit 5161fc9

Browse files
shadedbluebirdshadedbluebird
authored
[LITE][OPENCL] add op bilinear_interp; fix output size error of op nearest_interp when scale == 0 (#3034)
* [LITE][OPENCL] add op bilinear_interp; fix output size error of op nearest_interp when scale == 0 * [LITE][OPENCL]fix codestyle
1 parent 35ea6f9 commit 5161fc9

File tree

5 files changed

+183
-1
lines changed

5 files changed

+183
-1
lines changed

‎mobile/src/operators/bilinear_interp_op.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,7 @@ limitations under the License. */
1515
#ifdef BILINEAR_INTERP_OP
1616

1717
#include "operators/bilinear_interp_op.h"
18+
#include <vector>
1819

1920
namespace paddle_mobile {
2021
namespace operators {
@@ -49,6 +50,10 @@ namespace ops = paddle_mobile::operators;
4950
REGISTER_OPERATOR_CPU(bilinear_interp, ops::BilinearOp);
5051
#endif
5152

53+
#if PADDLE_MOBILE_CL
54+
REGISTER_OPERATOR_CL(bilinear_interp, ops::BilinearOp)
55+
#endif
56+
5257
#ifdef PADDLE_MOBILE_FPGA
5358
#endif
5459

‎mobile/src/operators/kernel/cl/bilinear_interp_kernel.cpp

Lines changed: 85 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,85 @@
1+
/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
2+
3+
Licensed under the Apache License, Version 2.0 (the "License");
4+
you may not use this file except in compliance with the License.
5+
You may obtain a copy of the License at
6+
7+
http://www.apache.org/licenses/LICENSE-2.0
8+
9+
Unless required by applicable law or agreed to in writing, software
10+
distributed under the License is distributed on an "AS IS" BASIS,
11+
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+
See the License for the specific language governing permissions and
13+
limitations under the License. */
14+
15+
#ifdef BILINEAR_INTERP_OP
16+
17+
#include <operators/kernel/bilinear_interp_kernel.h>
18+
19+
namespace paddle_mobile {
20+
namespace operators {
21+
template <>
22+
bool BilinearInterpKernel<GPU_CL, float>::Init(
23+
paddle_mobile::operators::BilinearInterpParam<paddle_mobile::GPU_CL>
24+
*param) {
25+
this->cl_helper_.AddKernel("bilinear_interp", "bilinear_interp_kernel.cl");
26+
return true;
27+
}
28+
29+
template <>
30+
void BilinearInterpKernel<GPU_CL, float>::Compute(
31+
const paddle_mobile::operators::BilinearInterpParam<paddle_mobile::GPU_CL>
32+
&param) {
33+
auto kernel = this->cl_helper_.KernelAt(0);
34+
auto default_work_size = this->cl_helper_.DefaultWorkSize(*(param.Out()));
35+
auto input = param.InputX();
36+
cl_mem input_image = input->GetCLImage();
37+
auto output = param.Out();
38+
cl_mem output_image = output->GetCLImage();
39+
float scale_h, scale_w;
40+
if (param.AlignCorners()) {
41+
scale_h = (input->dims()[2] - 1.0f) / (output->dims()[2] - 1.0f);
42+
scale_w = (input->dims()[3] - 1.0f) / (output->dims()[3] - 1.0f);
43+
} else {
44+
scale_h = input->dims()[2] / static_cast<float> output->dims()[2];
45+
scale_w = input->dims()[3] / static_cast<float> output->dims()[3];
46+
}
47+
float align_delta = 0.0f;
48+
if (!param.AlignCorners() && param.AlignMode() == 0) {
49+
align_delta = 0.5f;
50+
}
51+
int in_dims_h = input->dims()[2];
52+
int out_dims_h = output->dims()[2];
53+
int in_dims_w = input->dims()[3];
54+
int out_dims_w = output->dims()[3];
55+
56+
cl_int status;
57+
58+
status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input_image);
59+
CL_CHECK_ERRORS(status)
60+
status = clSetKernelArg(kernel, 1, sizeof(cl_mem), &output_image);
61+
CL_CHECK_ERRORS(status)
62+
status = clSetKernelArg(kernel, 2, sizeof(float), &scale_h);
63+
CL_CHECK_ERRORS(status)
64+
status = clSetKernelArg(kernel, 3, sizeof(float), &scale_w);
65+
CL_CHECK_ERRORS(status)
66+
status = clSetKernelArg(kernel, 4, sizeof(int), &in_dims_h);
67+
CL_CHECK_ERRORS(status)
68+
status = clSetKernelArg(kernel, 5, sizeof(int), &out_dims_h);
69+
CL_CHECK_ERRORS(status)
70+
status = clSetKernelArg(kernel, 6, sizeof(int), &in_dims_w);
71+
CL_CHECK_ERRORS(status)
72+
status = clSetKernelArg(kernel, 7, sizeof(int), &out_dims_w);
73+
CL_CHECK_ERRORS(status)
74+
status = clSetKernelArg(kernel, 8, sizeof(float), &align_delta);
75+
CL_CHECK_ERRORS(status)
76+
status = clEnqueueNDRangeKernel(
77+
this->cl_helper_.CLCommandQueue(), kernel, default_work_size.size(), NULL,
78+
default_work_size.data(), NULL, 0, NULL, NULL);
79+
CL_CHECK_ERRORS(status)
80+
}
81+
template class BilinearInterpKernel<GPU_CL, float>;
82+
} // namespace operators
83+
} // namespace paddle_mobile
84+
85+
#endif

‎mobile/src/operators/kernel/cl/cl_kernel/bilinear_interp_kernel.cl

Lines changed: 82 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,82 @@
1+
/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
2+
3+
Licensed under the Apache License, Version 2.0 (the "License");
4+
you may not use this file except in compliance with the License.
5+
You may obtain a copy of the License at
6+
7+
http://www.apache.org/licenses/LICENSE-2.0
8+
9+
Unless required by applicable law or agreed to in writing, software
10+
distributed under the License is distributed on an "AS IS" BASIS,
11+
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+
See the License for the specific language governing permissions and
13+
limitations under the License. */
14+
15+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
16+
__kernel void bilinear_interp(__read_only image2d_t input, __write_only image2d_t output,
17+
__private const float scale_h, __private const float scale_w,
18+
__private const int in_dims_h, __private const int out_dims_h,
19+
__private const int in_dims_w, __private const int out_dims_w,
20+
__private const float align_delta) {
21+
const int c = get_global_id(0);
22+
const int w = get_global_id(1);
23+
const int nh = get_global_id(2);
24+
25+
int2 output_pos;
26+
output_pos.x = c * out_dims_w + w;
27+
output_pos.y = nh;
28+
29+
// calculate center pixel's pos
30+
int out_n = nh / out_dims_h;
31+
int out_h = nh % out_dims_h;
32+
float center_w = (w + align_delta) * scale_w - align_delta;
33+
float center_h = (out_h + align_delta) * scale_h - align_delta;
34+
35+
int floor_w = (int)center_w;
36+
int floor_h = (int)center_h;
37+
int ceil_w = floor_w + 1;
38+
int ceil_h = floor_h + 1;
39+
40+
if (ceil_w > in_dims_w) {
41+
ceil_w = floor_w;
42+
}
43+
if (ceil_h > in_dims_h) {
44+
ceil_h = floor_h;
45+
}
46+
float wight0_w = center_w - floor_w;
47+
float wight0_h = center_h - floor_h;
48+
float wight1_w = 1.0 - wight0_w;
49+
float wight1_h = 1.0 - wight0_h;
50+
51+
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
52+
53+
// get left up pixel data
54+
int2 left_up;
55+
left_up.x = c * in_dims_w + floor_w;
56+
left_up.y = out_n * in_dims_h + ceil_h;
57+
half4 left_up_data = read_imageh(input, sampler, left_up);
58+
59+
// get left down pixel data
60+
int2 left_down;
61+
left_down.x = c * in_dims_w + floor_w;
62+
left_down.y = out_n * in_dims_h + floor_h;
63+
half4 left_down_data = read_imageh(input, sampler, left_down);
64+
65+
// get right up pixel data
66+
int2 right_up;
67+
right_up.x = c * in_dims_w + ceil_w;
68+
right_up.y = out_n * in_dims_h + ceil_h;
69+
half4 right_up_data = read_imageh(input, sampler, right_up);
70+
71+
// get right down pixel's data
72+
int2 right_down;
73+
right_down.x = c * in_dims_w + ceil_w;
74+
right_down.y = out_n * in_dims_h + floor_h;
75+
half4 right_down_data = read_imageh(input, sampler, right_down);
76+
77+
// calculate output data
78+
half4 data = (left_down_data * wight1_w + right_down_data * wight0_w) * wight1_h
79+
+ (left_up_data * wight1_w + right_up_data * wight0_w) * wight0_h;
80+
81+
write_imageh(output, output_pos, data);
82+
}

‎mobile/src/operators/nearest_interp_op.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,8 +27,12 @@ void NearestInterpolationOp<DeviceType, T>::InferShape() const {
2727
auto dim_x = this->param_.InputX()->dims(); // NCHW format
2828
DLOG << "dim_x :" << dim_x;
2929

30+
bool ignore_scale = false;
3031
int out_h = this->param_.OutH();
3132
int out_w = this->param_.OutW();
33+
if (out_h > 0 && out_w > 0) {
34+
ignore_scale = true;
35+
}
3236
PADDLE_MOBILE_ENFORCE(dim_x.size() == 4, "X's dimension must be 4");
3337

3438
if (this->param_.InputOutPutSize() != nullptr) {
@@ -40,7 +44,7 @@ void NearestInterpolationOp<DeviceType, T>::InferShape() const {
4044
}
4145

4246
DLOG << "this->param_.HasScale(): " << this->param_.HasScale();
43-
if (this->param_.HasScale()) {
47+
if (this->param_.HasScale() && !ignore_scale) {
4448
const float scale = this->param_.Scale();
4549
DLOG << "scale_: " << scale;
4650
std::vector<int64_t> dim_out({dim_x[0], dim_x[1],

‎mobile/src/operators/op_param.h

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3081,19 +3081,25 @@ class BilinearInterpParam : public OpParam {
30813081
out_ = OutFrom<GType>(outputs, *scope);
30823082
out_h_ = GetAttr<int>("out_h", attrs);
30833083
out_w_ = GetAttr<int>("out_w", attrs);
3084+
align_corners = GetAttr<bool>("align_corners", attrs);
3085+
align_mode = GetAttr<int>("align_mode", attrs);
30843086
}
30853087
const GType *InputX() const { return input_x_; }
30863088
const GType *InputOutPutSize() const { return input_outsize_; }
30873089
GType *Out() const { return out_; }
30883090
int OutH() const { return out_h_; }
30893091
int OutW() const { return out_w_; }
3092+
bool AlignCorners() const { return align_corners; }
3093+
int AlignMode() const { return align_mode; }
30903094

30913095
private:
30923096
GType *input_x_;
30933097
GType *input_outsize_;
30943098
GType *out_;
30953099
int out_h_;
30963100
int out_w_;
3101+
bool align_corners;
3102+
int align_mode;
30973103
};
30983104
#endif
30993105

0 commit comments

Comments
 (0)