diff options
Diffstat (limited to 'src/core/CL')
-rw-r--r-- | src/core/CL/cl_kernels/pooling_layer.cl | 16 | ||||
-rw-r--r-- | src/core/CL/kernels/CLPoolingLayerKernel.cpp | 9 |
2 files changed, 19 insertions, 6 deletions
diff --git a/src/core/CL/cl_kernels/pooling_layer.cl b/src/core/CL/cl_kernels/pooling_layer.cl index 6379c288b3..635c44a849 100644 --- a/src/core/CL/cl_kernels/pooling_layer.cl +++ b/src/core/CL/cl_kernels/pooling_layer.cl @@ -186,10 +186,14 @@ DATA_TYPE calculate_avg_scale(const int pool_size, const int upper_bound_w, const int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y) { - const int start_x = get_global_id(0) * stride_x - pad_x; - const int start_y = get_global_id(1) * stride_y - pad_y; + int start_x = get_global_id(0) * stride_x - pad_x; + int start_y = get_global_id(1) * stride_y - pad_y; const int end_x = min(start_x + pool_size, upper_bound_w); const int end_y = min(start_y + pool_size, upper_bound_h); +#if defined(EXCLUDE_PADDING) + start_x = max(0, start_x); + start_y = max(0, start_y); +#endif /* defined(EXCLUDE_PADDING) */ return ((end_y - start_y) * (end_x - start_x)); } @@ -334,10 +338,14 @@ VEC_DATA_TYPE(DATA_TYPE, 4) calculate_avg_scale4(const int pool_size, const int upper_bound_w, const int upper_bound_h, const int pad_x, const int pad_y, const int stride_x, const int stride_y) { - const int4 start_x = ((int4)get_global_id(0) * 4 + (int4)(0, 1, 2, 3)) * (int4)stride_x - (int4)pad_x; - const int start_y = get_global_id(1) * stride_y - pad_y; + int4 start_x = ((int4)get_global_id(0) * 4 + (int4)(0, 1, 2, 3)) * (int4)stride_x - (int4)pad_x; + int start_y = get_global_id(1) * stride_y - pad_y; const int4 end_x = min(start_x + (int4)pool_size, (int4)upper_bound_w); const int end_y = min(start_y + pool_size, upper_bound_h); +#if defined(EXCLUDE_PADDING) + start_x = max((int4)0, start_x); + start_y = max(0, start_y); +#endif /* defined(EXCLUDE_PADDING) */ return (VEC_DATA_TYPE(DATA_TYPE, 4))(1.f) / CONVERT_VECTOR4(DATA_TYPE)(((int4)(end_y - start_y)) * (end_x - start_x)); } diff --git a/src/core/CL/kernels/CLPoolingLayerKernel.cpp b/src/core/CL/kernels/CLPoolingLayerKernel.cpp index 542d5dcf9f..8b8f61e621 100644 --- a/src/core/CL/kernels/CLPoolingLayerKernel.cpp +++ b/src/core/CL/kernels/CLPoolingLayerKernel.cpp @@ -61,6 +61,7 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, const PoolingType pool_type = pool_info.pool_type(); const int pool_size = pool_info.pool_size(); const PadStrideInfo pad_stride_info = pool_info.pad_stride_info(); + bool exclude_padding = pool_info.exclude_padding(); std::tie(pool_pad_x, pool_pad_y) = pad_stride_info.pad(); std::tie(pool_stride_x, pool_stride_y) = pad_stride_info.stride(); @@ -109,8 +110,12 @@ void CLPoolingLayerKernel::configure(const ICLTensor *input, ICLTensor *output, build_opts.emplace(("-DSTRIDE_X=" + support::cpp11::to_string(pool_stride_x))); if(pool_type != PoolingType::MAX) { - build_opts.emplace(("-DMAX_WIDTH=" + support::cpp11::to_string(input->info()->dimension(0) + pool_pad_x))); - build_opts.emplace(("-DMAX_HEIGHT=" + support::cpp11::to_string(input->info()->dimension(1) + pool_pad_y))); + if(exclude_padding) + { + build_opts.emplace("-DEXCLUDE_PADDING"); + } + build_opts.emplace(("-DMAX_WIDTH=" + support::cpp11::to_string(input->info()->dimension(0) + (exclude_padding ? 0 : pool_pad_x)))); + build_opts.emplace(("-DMAX_HEIGHT=" + support::cpp11::to_string(input->info()->dimension(1) + (exclude_padding ? 0 : pool_pad_y)))); build_opts.emplace(("-DSTRIDE_Y=" + support::cpp11::to_string(pool_stride_y))); build_opts.emplace(("-DPAD_X=" + support::cpp11::to_string(pool_pad_x))); build_opts.emplace(("-DPAD_Y=" + support::cpp11::to_string(pool_pad_y))); |