diff options
author | Daniil Efremov <daniil.efremov@xored.com> | 2017-11-22 00:26:51 +0700 |
---|---|---|
committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:41:04 +0000 |
commit | 02bf80d4554cfc824a76008905921cb564bee999 (patch) | |
tree | b86ebbed4d330af69c1107c10ce5e765705e88dd /src/core/CL | |
parent | 6194145681232bf59e0455434f15aba42956145b (diff) | |
download | ComputeLibrary-02bf80d4554cfc824a76008905921cb564bee999.tar.gz |
COMPMID-661: Fix scale border issue (#38)
Change-Id: If1dcca724e5e5f5ab363ffc16b0ef8c943e0b657
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/110105
Tested-by: BSG Visual Compute Jenkins server to access repositories on http://mpd-gerrit.cambridge.arm.com <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
Diffstat (limited to 'src/core/CL')
-rw-r--r-- | src/core/CL/cl_kernels/scale.cl | 16 | ||||
-rw-r--r-- | src/core/CL/kernels/CLScaleKernel.cpp | 3 |
2 files changed, 16 insertions, 3 deletions
diff --git a/src/core/CL/cl_kernels/scale.cl b/src/core/CL/cl_kernels/scale.cl index d533d970c7..a2ae8c4dd6 100644 --- a/src/core/CL/cl_kernels/scale.cl +++ b/src/core/CL/cl_kernels/scale.cl @@ -49,13 +49,23 @@ inline const float8 transform_nearest(const float2 coord, const float2 scale) inline const float8 transform_bilinear(const float2 coord, const float2 scale) { const float4 in_x_coords = (float4)(coord.s0, 1 + coord.s0, 2 + coord.s0, 3 + coord.s0); - const float4 new_x = (in_x_coords + ((float4)(0.5f))) * (float4)(scale.s0) - (float4)(0.5f); - const float4 new_y = (float4)((coord.s1 + 0.5f) * scale.s1 - 0.5f); +#ifdef SAMPLING_POLICY_TOP_LEFT + const float4 new_x = in_x_coords * (float4)(scale.s0); + const float4 new_y = (float4)(coord.s1 * scale.s1); return (float8)(new_x.s0, new_y.s0, new_x.s1, new_y.s1, new_x.s2, new_y.s2, new_x.s3, new_y.s3); +#elif SAMPLING_POLICY_CENTER + const float4 new_x = (in_x_coords + ((float4)(0.5f))) * (float4)(scale.s0) - (float4)(0.5f); + const float4 new_y = (float4)((coord.s1 + 0.5f) * scale.s1 - 0.5f); + return (float8)(new_x.s0, new_y.s0, new_x.s1, new_y.s1, new_x.s2, new_y.s2, new_x.s3, new_y.s3); +#else /* SAMPLING_POLICY */ +#error("Unsupported sampling policy"); +#endif /* SAMPLING_POLICY */ } /** Performs an affine transformation on an image interpolating with the NEAREAST NEIGHBOUR method. Input and output are single channel U8 or S16. * + * @note Sampling policy to used is passed as -DSAMPLING_POLICY_(TYPE) e.g. -DSAMPLING_POLICY_TOP_LEFT + * * @param[in] in_ptr Pointer to the source image. Supported data types: U8, S16. * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) * @param[in] in_step_x src_stride_x * number of elements along X processed per workitem(in bytes) @@ -90,6 +100,8 @@ __kernel void scale_nearest_neighbour( /** Performs an affine transformation on an image interpolating with the BILINEAR method. * + * @note Sampling policy to used is passed as -DSAMPLING_POLICY_(TYPE) e.g. -DSAMPLING_POLICY_TOP_LEFT + * * @param[in] in_ptr Pointer to the source image. Supported data types: U8, S16. * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) * @param[in] in_step_x src_stride_x * number of elements along X processed per workitem(in bytes) diff --git a/src/core/CL/kernels/CLScaleKernel.cpp b/src/core/CL/kernels/CLScaleKernel.cpp index 6a5d24c943..673304a271 100644 --- a/src/core/CL/kernels/CLScaleKernel.cpp +++ b/src/core/CL/kernels/CLScaleKernel.cpp @@ -44,7 +44,7 @@ BorderSize CLScaleKernel::border_size() const return BorderSize(1); } -void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, InterpolationPolicy policy, bool border_undefined) +void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, InterpolationPolicy policy, bool border_undefined, SamplingPolicy sampling_policy) { ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16, DataType::F16, DataType::F32); ARM_COMPUTE_ERROR_ON_NULLPTR(output); @@ -75,6 +75,7 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo CLBuildOptions build_opts; build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input->info()->data_type())); build_opts.add_option("-DBORDER_SIZE=" + support::cpp11::to_string(border.right)); + build_opts.add_option_if_else(sampling_policy == SamplingPolicy::CENTER, "-DSAMPLING_POLICY_CENTER", "-DSAMPLING_POLICY_TOP_LEFT"); std::string interpolation_name = string_from_interpolation_policy(policy); std::transform(interpolation_name.begin(), interpolation_name.end(), interpolation_name.begin(), ::tolower); |