diff options
Diffstat (limited to 'src/core')
-rw-r--r-- | src/core/CL/cl_kernels/scale.cl | 16 | ||||
-rw-r--r-- | src/core/CL/kernels/CLScaleKernel.cpp | 3 | ||||
-rw-r--r-- | src/core/NEON/kernels/NEScaleKernel.cpp | 5 |
3 files changed, 20 insertions, 4 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); diff --git a/src/core/NEON/kernels/NEScaleKernel.cpp b/src/core/NEON/kernels/NEScaleKernel.cpp index b1ced7e38d..1918a77300 100644 --- a/src/core/NEON/kernels/NEScaleKernel.cpp +++ b/src/core/NEON/kernels/NEScaleKernel.cpp @@ -48,12 +48,15 @@ BorderSize NEScaleKernel::border_size() const return BorderSize(1); } -void NEScaleKernel::configure(const ITensor *input, const ITensor *dx, const ITensor *dy, const ITensor *offsets, ITensor *output, InterpolationPolicy policy, bool border_undefined) +void NEScaleKernel::configure(const ITensor *input, const ITensor *dx, const ITensor *dy, const ITensor *offsets, ITensor *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::F32); ARM_COMPUTE_ERROR_ON_NULLPTR(output); ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); ARM_COMPUTE_ERROR_ON(output == input); + ARM_COMPUTE_ERROR_ON(sampling_policy != SamplingPolicy::CENTER); + ARM_COMPUTE_UNUSED(sampling_policy); if(policy == InterpolationPolicy::NEAREST_NEIGHBOR) { |