From 26769ad792f0ead12059dce2d92807cd2e0081be Mon Sep 17 00:00:00 2001 From: hjchen2 Date: Sun, 9 Dec 2018 22:46:00 +0800 Subject: [PATCH] Fix pooling3x3 bug --- src/operators/math/pooling.cpp | 8 +- src/operators/math/pooling.h | 2 +- src/operators/math/pooling3x3.cpp | 59 +++++++---- test/operators/test_pool_op.cpp | 157 ++++++++++++++++++------------ 4 files changed, 145 insertions(+), 81 deletions(-) diff --git a/src/operators/math/pooling.cpp b/src/operators/math/pooling.cpp index 3bb6c27c5a..b4aba52b9b 100644 --- a/src/operators/math/pooling.cpp +++ b/src/operators/math/pooling.cpp @@ -60,12 +60,18 @@ void Pooling

::operator()(const framework::Tensor &input, wstart = std::max(wstart, 0); PoolingVal

val; + // std::cout << "output[" << ph * output_width + pw << "]:" + // << std::endl; for (int h = hstart; h < hend; ++h) { for (int w = wstart; w < wend; ++w) { val += input_ptr[h * input_width + w]; + // std::cout << "input[" << h << "][" << w << "] = " + // << input_ptr[h * input_width + w] << std::endl; } } - output_data[ph * output_width + pw] = val.Value(); + output_ptr[ph * output_width + pw] = val.Value(); + // std::cout << "output[" << ph * output_width + pw << "] = " + // << val.Value() << std::endl; } } } diff --git a/src/operators/math/pooling.h b/src/operators/math/pooling.h index 9407270a47..9fcfdf811f 100644 --- a/src/operators/math/pooling.h +++ b/src/operators/math/pooling.h @@ -35,7 +35,7 @@ struct PoolingVal { float val; int count; PoolingVal() { - val = std::numeric_limits::min(); + val = -std::numeric_limits::max(); count = 0; } inline PoolingVal

&operator+=(const float &x) { diff --git a/src/operators/math/pooling3x3.cpp b/src/operators/math/pooling3x3.cpp index 3918001d76..8f1defa6cb 100644 --- a/src/operators/math/pooling3x3.cpp +++ b/src/operators/math/pooling3x3.cpp @@ -161,9 +161,9 @@ struct Pooling3x3 { x0.val[0] = vld1q_f32(input_ptr0); x0.val[1] = vld1q_f32(input_ptr0 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); y0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); @@ -172,9 +172,9 @@ struct Pooling3x3 { x0.val[0] = vld1q_f32(input_ptr1); x0.val[1] = vld1q_f32(input_ptr1 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); y1.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); @@ -185,9 +185,9 @@ struct Pooling3x3 { x0.val[0] = vld1q_f32(input_ptr2); x0.val[1] = vld1q_f32(input_ptr2 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); y2.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); @@ -204,9 +204,9 @@ struct Pooling3x3 { x0.val[0] = vld1q_f32(input_ptr3); x0.val[1] = vld1q_f32(input_ptr3 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); y0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); @@ -223,9 +223,9 @@ struct Pooling3x3 { x0.val[0] = vld1q_f32(input_ptr4); x0.val[1] = vld1q_f32(input_ptr4 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); @@ -242,9 +242,9 @@ struct Pooling3x3 { x0.val[0] = vld1q_f32(input_ptr5); x0.val[1] = vld1q_f32(input_ptr5 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); @@ -270,12 +270,14 @@ struct Pooling3x3 { // remain w if (remain >= 4) { x0.val[0] = vld1q_f32(input_ptr0); + x0.val[1] = vld1q_f32(input_ptr0 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); y0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); x0.val[0] = vld1q_f32(input_ptr1); + x0.val[1] = vld1q_f32(input_ptr1 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); @@ -283,6 +285,7 @@ struct Pooling3x3 { y0.val[0] = vPoolPreq_f32

(y1.val[0], y0.val[0]); x0.val[0] = vld1q_f32(input_ptr2); + x0.val[1] = vld1q_f32(input_ptr2 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); @@ -293,6 +296,7 @@ struct Pooling3x3 { vst1q_f32(output_ptr0, y0.val[0]); x0.val[0] = vld1q_f32(input_ptr3); + x0.val[1] = vld1q_f32(input_ptr3 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); @@ -303,6 +307,7 @@ struct Pooling3x3 { vst1q_f32(output_ptr1, y1.val[0]); x0.val[0] = vld1q_f32(input_ptr4); + x0.val[1] = vld1q_f32(input_ptr4 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); @@ -313,6 +318,7 @@ struct Pooling3x3 { vst1q_f32(output_ptr2, y2.val[0]); x0.val[0] = vld1q_f32(input_ptr5); + x0.val[1] = vld1q_f32(input_ptr5 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); @@ -372,9 +378,9 @@ struct Pooling3x3 { x0.val[0] = vld1q_f32(input_ptr0); x0.val[1] = vld1q_f32(input_ptr0 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); y0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); @@ -383,22 +389,26 @@ struct Pooling3x3 { x0.val[0] = vld1q_f32(input_ptr1); x0.val[1] = vld1q_f32(input_ptr1 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); + x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); + x0.val[1] = vPoolPreq_f32

(x0.val[1], x2.val[1]); y0.val[0] = vPoolPreq_f32

(x0.val[0], y0.val[0]); y0.val[1] = vPoolPreq_f32

(x0.val[1], y0.val[1]); x0.val[0] = vld1q_f32(input_ptr2); x0.val[1] = vld1q_f32(input_ptr2 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); - x1.val[1] = vextq_f32(x0.val[0], x0.val[1], 1); + x1.val[1] = vextq_f32(x0.val[1], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); - x2.val[1] = vextq_f32(x0.val[0], x0.val[1], 2); + x2.val[1] = vextq_f32(x0.val[1], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); x0.val[1] = vPoolPreq_f32

(x0.val[1], x1.val[1]); + x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); + x0.val[1] = vPoolPreq_f32

(x0.val[1], x2.val[1]); y0.val[0] = vPoolPreq_f32

(x0.val[0], y0.val[0]); y0.val[1] = vPoolPreq_f32

(x0.val[1], y0.val[1]); y0.val[0] = vPoolPostq_f32

(y0.val[0]); @@ -414,21 +424,26 @@ struct Pooling3x3 { // remain w if (remain >= 4) { x0.val[0] = vld1q_f32(input_ptr0); + x0.val[1] = vld1q_f32(input_ptr0 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); y0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); x0.val[0] = vld1q_f32(input_ptr1); + x0.val[1] = vld1q_f32(input_ptr1 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); + x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); y0.val[0] = vPoolPreq_f32

(x0.val[0], y0.val[0]); x0.val[0] = vld1q_f32(input_ptr2); + x0.val[1] = vld1q_f32(input_ptr2 + 4); x1.val[0] = vextq_f32(x0.val[0], x0.val[1], 1); x2.val[0] = vextq_f32(x0.val[0], x0.val[1], 2); x0.val[0] = vPoolPreq_f32

(x0.val[0], x1.val[0]); + x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); y0.val[0] = vPoolPreq_f32

(x0.val[0], y0.val[0]); y0.val[0] = vPoolPostq_f32

(y0.val[0]); vst1q_f32(output_ptr0, y0.val[0]); @@ -540,6 +555,8 @@ struct Pooling3x3 { x2.val[1] = vextq_f32(x1.val[0], x1.val[0], 1); x0.val[0] = vPoolPreq_f32

(x0.val[0], x0.val[1]); x0.val[1] = vPoolPreq_f32

(x1.val[0], x1.val[1]); + x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); + x0.val[1] = vPoolPreq_f32

(x0.val[1], x2.val[1]); y0.val[0] = vPoolPreq_f32

(x0.val[0], y0.val[0]); y0.val[1] = vPoolPreq_f32

(x0.val[1], y0.val[1]); @@ -616,6 +633,7 @@ struct Pooling3x3 { input_ptr3 += 12; input_ptr4 += 12; input_ptr5 += 12; + input_ptr6 += 12; output_ptr0 += 6; output_ptr1 += 6; output_ptr2 += 6; @@ -632,6 +650,7 @@ struct Pooling3x3 { x1.val[0] = vdupq_n_f32(input_ptr1[8]); x2.val[0] = vextq_f32(x0.val[0], x1.val[0], 1); x0.val[0] = vPoolPreq_f32

(x0.val[0], x0.val[1]); + x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); y0.val[0] = vPoolPreq_f32

(x0.val[0], y0.val[0]); x0 = vld2q_f32(input_ptr2); @@ -681,6 +700,7 @@ struct Pooling3x3 { input_ptr3 += 8; input_ptr4 += 8; input_ptr5 += 8; + input_ptr6 += 8; output_ptr0 += 4; output_ptr1 += 4; output_ptr2 += 4; @@ -738,6 +758,8 @@ struct Pooling3x3 { x2.val[1] = vextq_f32(x1.val[0], x1.val[0], 1); x0.val[0] = vPoolPreq_f32

(x0.val[0], x0.val[1]); x0.val[1] = vPoolPreq_f32

(x1.val[0], x1.val[1]); + x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); + x0.val[1] = vPoolPreq_f32

(x0.val[1], x2.val[1]); y0.val[0] = vPoolPreq_f32

(x0.val[0], y0.val[0]); y0.val[1] = vPoolPreq_f32

(x0.val[1], y0.val[1]); @@ -773,6 +795,7 @@ struct Pooling3x3 { x1.val[0] = vdupq_n_f32(input_ptr1[8]); x2.val[0] = vextq_f32(x0.val[0], x1.val[0], 1); x0.val[0] = vPoolPreq_f32

(x0.val[0], x0.val[1]); + x0.val[0] = vPoolPreq_f32

(x0.val[0], x2.val[0]); y0.val[0] = vPoolPreq_f32

(x0.val[0], y0.val[0]); x0 = vld2q_f32(input_ptr2); diff --git a/test/operators/test_pool_op.cpp b/test/operators/test_pool_op.cpp index ae5ff9d3f7..b38123a7d6 100644 --- a/test/operators/test_pool_op.cpp +++ b/test/operators/test_pool_op.cpp @@ -21,20 +21,7 @@ namespace paddle_mobile { namespace math = operators::math; -static int PoolOutputSize(int input_size, int filter_size, int padding, - int stride, bool ceil_mode) { - int output_size; - if (!ceil_mode) { - output_size = (input_size - filter_size + 2 * padding) / stride + 1; - } else { - output_size = - (input_size - filter_size + 2 * padding + stride - 1) / stride + 1; - } - return output_size; -} - -template +template int TestPoolOp(int in_channels, int in_height, int in_width) { int kernel_h = Kernel; int kernel_w = Kernel; @@ -42,7 +29,6 @@ int TestPoolOp(int in_channels, int in_height, int in_width) { int pad_w = Pad; int stride_h = Stride; int stride_w = Stride; - bool ceil_mode = CeilMode != 0; std::string pooling_type = (PoolType == 0 ? "max" : "avg"); int batch_size = 1; @@ -53,14 +39,6 @@ int TestPoolOp(int in_channels, int in_height, int in_width) { framework::DDim input_shape = framework::make_ddim({batch_size, input_c, input_h, input_w}); - std::vector output_shape_v({batch_size, input_c}); - output_shape_v.push_back( - PoolOutputSize(input_h, kernel_h, pad_h, stride_h, ceil_mode)); - output_shape_v.push_back( - PoolOutputSize(input_w, kernel_w, pad_w, stride_w, ceil_mode)); - - framework::DDim output_shape = framework::make_ddim(output_shape_v); - VariableNameMap inputs; VariableNameMap outputs; auto scope = std::make_shared(); @@ -69,7 +47,11 @@ int TestPoolOp(int in_channels, int in_height, int in_width) { auto input_var = scope.get()->Var("input"); auto input = input_var->template GetMutable(); - SetupTensor(input, input_shape, -127, 127); + SetupTensor(input, input_shape, -127, 127); + + // for (int i = 0; i < input->numel(); ++i) { + // DLOG << "input[" << i << "] = " << input->data()[i]; + // } auto output_var = scope.get()->Var("output"); framework::AttributeMap attrs; @@ -86,8 +68,9 @@ int TestPoolOp(int in_channels, int in_height, int in_width) { op->Init(); op->Run(); + auto output = output_var->template Get(); framework::Tensor output_cmp; - output_cmp.mutable_data(output_shape); + output_cmp.mutable_data(output->dims()); if (pooling_type == "avg") { math::Pooling()(*input, std::vector{kernel_h, kernel_w}, @@ -100,13 +83,19 @@ int TestPoolOp(int in_channels, int in_height, int in_width) { } // compare results - auto output = output_var->template Get(); - const T *output_data = output->data(); - T *output_cmp_data = output_cmp.data(); + const float *output_data = output->data(); + float *output_cmp_data = output_cmp.data(); for (int i = 0; i < output->numel(); ++i) { - PADDLE_MOBILE_ENFORCE(output_data[i] == output_cmp_data[i], - "output[%d] = %d, output_cmp[%d] = %d", i, - output_data[i], i, output_cmp_data[i]); + float gap = output_data[i] - output_cmp_data[i]; + // PADDLE_MOBILE_ENFORCE(output_data[i] == output_cmp_data[i], + // "output[%d] = %d, output_cmp[%d] = %d", i, + // output_data[i], i, output_cmp_data[i]); + if (gap > 1e-5 && std::abs(gap / (output_data[i] + 1e-5)) > 1e-3) { + LOG(kLOG_INFO) << "output_data[" << i << "] = " << output_data[i] + << ", output_cmp_data[" << i + << "] = " << output_cmp_data[i]; + exit(1); + } } delete op; return 0; @@ -127,34 +116,80 @@ int main(int argc, char *argv[]) { int in_channels = atoi(argv[1]); int in_height = atoi(argv[2]); int in_width = atoi(argv[3]); - // kernel = 3, pad = 1, stride = 1 - LOG(paddle_mobile::kLOG_INFO) - << "float, ceil_mode=false, pooling_type=max, kernel=3, pad=1, stride=1"; - paddle_mobile::TestPoolOp(in_channels, in_height, - in_width); - // kernel = 3, pad = 0, stride = 2 - LOG(paddle_mobile::kLOG_INFO) - << "float, ceil_mode=false, pooling_type=max, kernel=3, pad=0, stride=2"; - paddle_mobile::TestPoolOp(in_channels, in_height, - in_width); - // kernel = 5, pad = 0, stride = 1 - LOG(paddle_mobile::kLOG_INFO) - << "float, ceil_mode=false, pooling_type=avg, kernel=5, pad=0, stride=1"; - paddle_mobile::TestPoolOp(in_channels, in_height, - in_width); - // kernel = 5, pad = 0, stride = 2 - LOG(paddle_mobile::kLOG_INFO) - << "float, ceil_mode=false, pooling_type=avg, kernel=5, pad=0, stride=1"; - paddle_mobile::TestPoolOp(in_channels, in_height, - in_width); - // kernel = 7, pad = 0, stride = 1 - LOG(paddle_mobile::kLOG_INFO) - << "float, ceil_mode=false, pooling_type=avg, kernel=7, pad=0, stride=1"; - paddle_mobile::TestPoolOp(in_channels, in_height, - in_width); - // kernel = 7, pad = 0, stride = 4 - LOG(paddle_mobile::kLOG_INFO) - << "float, ceil_mode=false, pooling_type=avg, kernel=7, pad=0, stride=4"; - paddle_mobile::TestPoolOp(in_channels, in_height, - in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=max, kernel=3, pad=0, stride=1"; + paddle_mobile::TestPoolOp<0, 3, 0, 1>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=max, kernel=3, pad=1, stride=1"; + paddle_mobile::TestPoolOp<0, 3, 1, 1>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=max, kernel=3, pad=2, stride=1"; + paddle_mobile::TestPoolOp<0, 3, 2, 1>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=max, kernel=3, pad=5, stride=1"; + paddle_mobile::TestPoolOp<0, 3, 5, 1>(in_channels, in_height, in_width); + + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=avg, kernel=3, pad=0, stride=1"; + paddle_mobile::TestPoolOp<1, 3, 0, 1>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=avg, kernel=3, pad=1, stride=1"; + paddle_mobile::TestPoolOp<1, 3, 1, 1>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=avg, kernel=3, pad=2, stride=1"; + paddle_mobile::TestPoolOp<1, 3, 2, 1>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=avg, kernel=3, pad=5, stride=1"; + paddle_mobile::TestPoolOp<1, 3, 5, 1>(in_channels, in_height, in_width); + + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=max, kernel=3, pad=0, stride=2"; + paddle_mobile::TestPoolOp<0, 3, 0, 2>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=max, kernel=3, pad=1, stride=2"; + paddle_mobile::TestPoolOp<0, 3, 1, 2>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=max, kernel=3, pad=2, stride=2"; + paddle_mobile::TestPoolOp<0, 3, 2, 2>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=max, kernel=3, pad=5, stride=2"; + paddle_mobile::TestPoolOp<0, 3, 5, 2>(in_channels, in_height, in_width); + + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=avg, kernel=3, pad=0, stride=2"; + paddle_mobile::TestPoolOp<1, 3, 0, 2>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=avg, kernel=3, pad=1, stride=2"; + paddle_mobile::TestPoolOp<1, 3, 1, 2>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=avg, kernel=3, pad=2, stride=2"; + paddle_mobile::TestPoolOp<1, 3, 2, 2>(in_channels, in_height, in_width); + LOG(paddle_mobile::kLOG_INFO) + << "float, pooling_type=avg, kernel=3, pad=5, stride=2"; + paddle_mobile::TestPoolOp<1, 3, 5, 2>(in_channels, in_height, in_width); + + // // kernel = 5, pad = 0, stride = 1 + // LOG(paddle_mobile::kLOG_INFO) + // << "float, ceil_mode=false, pooling_type=avg, kernel=5, pad=0, + // stride=1"; + // paddle_mobile::TestPoolOp(in_channels, in_height, + // in_width); + // // kernel = 5, pad = 0, stride = 2 + // LOG(paddle_mobile::kLOG_INFO) + // << "float, ceil_mode=false, pooling_type=avg, kernel=5, pad=0, + // stride=1"; + // paddle_mobile::TestPoolOp(in_channels, in_height, + // in_width); + // // kernel = 7, pad = 0, stride = 1 + // LOG(paddle_mobile::kLOG_INFO) + // << "float, ceil_mode=false, pooling_type=avg, kernel=7, pad=0, + // stride=1"; + // paddle_mobile::TestPoolOp(in_channels, in_height, + // in_width); + // // kernel = 7, pad = 0, stride = 4 + // LOG(paddle_mobile::kLOG_INFO) + // << "float, ceil_mode=false, pooling_type=avg, kernel=7, pad=0, + // stride=4"; + // paddle_mobile::TestPoolOp(in_channels, in_height, + // in_width); } -- GitLab