Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[Py][CL]fix pool2d bugs and refactor python unittest #8591

Merged
merged 1 commit into from
Mar 10, 2022
Merged

[Py][CL]fix pool2d bugs and refactor python unittest #8591

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 5 additions & 5 deletions lite/backends/opencl/cl_kernel/image/pool_kernel.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,10 +37,10 @@ __kernel void pool(__read_only image2d_t input,
int start_h, start_w, end_h, end_w;
int pool_size = 1;
if (adaptive == 1) {
start_h = floor((out_h * in_height) / (float)out_height);
start_w = floor((out_w * in_width) / (float)out_width);
end_h = ceil(((out_h + 1) * in_height) / (float)out_height);
end_w = ceil(((out_w + 1) * in_width) / (float)out_width);
start_h = (out_h * in_height) / out_height;
start_w = (out_w * in_width) / out_width;
end_h = ((out_h + 1) * in_height + (out_height - 1)) / out_height;
end_w = ((out_w + 1) * in_width + (out_width - 1)) / out_width;
} else {
start_h = out_h * stride_h - pad_top;
start_w = out_w * stride_w - pad_left;
Expand Down Expand Up @@ -269,4 +269,4 @@ __kernel void pool_local(__read_only image2d_t input,
local_output[local_id]);
}
#endif // POOL_AVG
}
}
4 changes: 2 additions & 2 deletions lite/core/optimizer/mir/adaptive_1x1_pool2d_convert_global_pass.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,7 @@ void Adaptive1x1Pool2dConvertGlobalPass::Apply(
VLOG(1) << "check global_pooling:" << global_pooling;
VLOG(1) << "check ksize:" << ksize[0] << "," << ksize[1]
<< " | ksize_one:" << ksize_one;
if (adaptive && ksize_one) {
if (adaptive && ksize_one && !global_pooling) {
return true;
}
return false;
Expand Down Expand Up @@ -74,4 +74,4 @@ void Adaptive1x1Pool2dConvertGlobalPass::Apply(

REGISTER_MIR_PASS(adaptive_1x1_pool2d_convert_global_pass,
paddle::lite::mir::Adaptive1x1Pool2dConvertGlobalPass)
.BindTargets({TARGET(kOpenCL), TARGET(kARM)});
.BindTargets({TARGET(kARM)});
7 changes: 5 additions & 2 deletions lite/kernels/opencl/pool_image_compute.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -239,8 +239,11 @@ class PoolComputeImage2D : public KernelLite<TARGET(kOpenCL),
CL_CHECK_FATAL(status);
status = kernel_.setArg(arg_idx++, static_cast<int>(exclusive));
CL_CHECK_FATAL(status);
status = kernel_.setArg(arg_idx++, static_cast<int>(adaptive));
CL_CHECK_FATAL(status);
if (adaptive == true && pooling_type == "max") {
status = kernel_.setArg(arg_idx++, static_cast<int>(!adaptive));
} else {
status = kernel_.setArg(arg_idx++, static_cast<int>(adaptive));
}

#ifdef LITE_WITH_LOG
const std::vector<int>& paddings = *param.paddings;
Expand Down
47 changes: 22 additions & 25 deletions lite/tests/unittest_py/op/test_pool2d_op.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,9 @@
import hypothesis.strategies as st
import argparse

import numpy as np
from functools import partial


class TestPool2dOp(AutoScanTest):
def __init__(self, *args, **kwargs):
Expand Down Expand Up @@ -87,13 +90,15 @@ def sample_program_configs(self, draw):
ksize = draw(
st.lists(
st.integers(
min_value=1, max_value=32), min_size=2, max_size=2))
min_value=1, max_value=7), min_size=2, max_size=2))
strides = draw(
st.lists(
st.integers(
min_value=1, max_value=2), min_size=2, max_size=2))
min_value=1, max_value=16), min_size=2, max_size=2))
paddings = draw(
st.sampled_from([[0, 0], [0, 0, 0, 0], [1, 1], [1, 1, 1, 1]]))
st.lists(
st.integers(
min_value=0, max_value=16), min_size=2, max_size=2))
padding_algorithm = draw(
st.sampled_from(["EXPLICIT", "VALID", "SAME"]))
pooling_type = draw(st.sampled_from(["max", "avg"]))
Expand All @@ -104,21 +109,18 @@ def sample_program_configs(self, draw):
use_cudnn = False
use_mkldnn = False
use_quantizer = False
is_test = False
is_test = True
data_format = "NCHW"
assume(ksize[0] <= (in_shape[2] - strides[0] - 1))
assume(ksize[1] <= (in_shape[3] - strides[1] - 1))
if paddings[0] == 1:
assume((ksize[0] != 1 and ksize[1] != 1))

#This is the correct input when adaptive
if adaptive:
assume(in_shape[2] / ksize[0] == strides[0])
assume(in_shape[3] / ksize[1] == strides[1])
assume(ksize[0] <= in_shape[2])
assume(ksize[1] <= in_shape[3])
if adaptive == False:
assume(ksize[0] > paddings[0] and ksize[1] > paddings[1])
if adaptive == False and ceil_mode == True:
assume(strides[0] > 1 and strides[1] > 1)

#both paddle and lite have invalid output, so it is an invalid input.
if paddings == [0, 0] or paddings == [0, 0, 0, 0]:
assume(ceil_mode == False)
def generate_input(*args, **kwargs):
return np.random.normal(0.0, 1.0, in_shape).astype(np.float32)

build_ops = OpConfig(
type="pool2d",
Expand All @@ -143,7 +145,9 @@ def sample_program_configs(self, draw):
program_config = ProgramConfig(
ops=[build_ops],
weights={},
inputs={"input_data": TensorConfig(shape=in_shape)},
inputs={
"input_data": TensorConfig(data_gen=partial(generate_input))
},
outputs=["output_data"])
return program_config

Expand All @@ -167,9 +171,6 @@ def teller1(program_config, predictor_config):
if program_config.ops[0].attrs["ceil_mode"] == True \
and strides[0] != strides[1]:
return True
if predictor_config.target() == TargetType.OpenCL:
if program_config.ops[0].attrs["adaptive"] == True:
return True

self.add_ignore_check_case(
teller1, IgnoreReasons.ACCURACY_ERROR,
Expand All @@ -193,11 +194,7 @@ def teller2(program_config, predictor_config):

def teller3(program_config, predictor_config):
if predictor_config.target() == TargetType.ARM:
# This is an paddle error, when padding_algorithm == "Valid" with exclusive is False
if program_config.ops[0].attrs[
"padding_algorithm"] == "VALID" and program_config.ops[
0].attrs["exclusive"] == False:
return True
return True

self.add_ignore_check_case(
teller3, IgnoreReasons.PADDLE_NOT_SUPPORT,
Expand All @@ -209,7 +206,7 @@ def test(self, *args, **kwargs):
max_examples = 100
if target_str == "OpenCL":
# Make sure to generate enough valid cases for OpenCL
max_examples = 300
max_examples = 200
if target_str == "Metal":
# Make sure to generate enough valid cases for Metal
max_examples = 500
Expand Down