aboutsummaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
Diffstat (limited to 'src')
-rw-r--r--src/core/CL/cl_kernels/remap.cl16
-rw-r--r--src/core/CL/kernels/CLRemapKernel.cpp32
-rw-r--r--src/core/CL/kernels/CLRemapKernel.h18
-rw-r--r--src/runtime/CL/functions/CLRemap.cpp32
-rw-r--r--src/runtime/NEON/functions/NERemap.cpp6
5 files changed, 67 insertions, 37 deletions
diff --git a/src/core/CL/cl_kernels/remap.cl b/src/core/CL/cl_kernels/remap.cl
index 8ea4e84e96..cb67c2df1e 100644
--- a/src/core/CL/cl_kernels/remap.cl
+++ b/src/core/CL/cl_kernels/remap.cl
@@ -24,6 +24,7 @@
#include "helpers.h"
#include "warp_helpers.h"
+#ifndef DEPTH_OUT
/** Performs a remapping of an input image to an output given two remapping image using nearest neighbor as interpolation.
*
* This kernel performs remapping with this method of pixel coordinate translation:
@@ -129,20 +130,20 @@ __kernel void remap_bilinear_nchw(
vstore4(bilinear_interpolate(&in, clamp_to_border(map_coords, width, height), width, height), 0, out.ptr);
}
-
+#else // DEPTH_OUT
/** Performs a remapping of an input image to an output given two remapping image using nearest neighbor as interpolation.
* Also applies constant border value, "border_val", if "CONSTANT_BORDER" is set.
*
* This kernel performs remapping with this method of pixel coordinate translation:
* out(x,y) = in(mapx(x,y), mapy(x,y));
*
- * @param[in] in_ptr Pointer to the source image. Supported data types: U8.
+ * @param[in] in_ptr Pointer to the source image. Supported data types: U8,F16.
* @param[in] in_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] in_step_x in_stride_x * number of elements along X processed per work item (in bytes)
* @param[in] in_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] in_step_y in_stride_y * number of elements along Y processed per work item (in bytes)
* @param[in] in_offset_first_element_in_bytes Offset of the first element in the source image
- * @param[out] out_ptr Pointer to the destination image. Supported data types: U8.
+ * @param[out] out_ptr Pointer to the destination image. Supported data types: U8,F16.
* @param[in] out_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] out_step_x out_stride_x * number of elements along X processed per work item (in bytes)
* @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes)
@@ -162,10 +163,8 @@ __kernel void remap_bilinear_nchw(
* @param[in] mapy_offset_first_element_in_bytes Offset of the first element in the remapping image
* @param[in] width Width of the input image
* @param[in] height Height of the input image
+ * @param[in] border_val Value to use for border around input tensor when in CONSTANT border is selected
*/
-
-#if defined(DEPTH_OUT)
-
__kernel void remap_nearest_neighbour_nhwc(
TENSOR4D_DECLARATION(in),
TENSOR4D_DECLARATION(out),
@@ -206,13 +205,13 @@ __kernel void remap_nearest_neighbour_nhwc(
* This kernel performs remapping with this method of pixel coordinate translation:
* out(x,y) = in(mapx(x,y), mapy(x,y));
*
- * @param[in] in_ptr Pointer to the source image. Supported data types: U8.
+ * @param[in] in_ptr Pointer to the source image. Supported data types: U8,F16.
* @param[in] in_stride_x Stride of the source image in X dimension (in bytes)
* @param[in] in_step_x in_stride_x * number of elements along X processed per work item (in bytes)
* @param[in] in_stride_y Stride of the source image in Y dimension (in bytes)
* @param[in] in_step_y in_stride_y * number of elements along Y processed per work item (in bytes)
* @param[in] in_offset_first_element_in_bytes Offset of the first element in the source image
- * @param[out] out_ptr Pointer to the destination image. Supported data types: U8.
+ * @param[out] out_ptr Pointer to the destination image. Supported data types: U8,F16.
* @param[in] out_stride_x Stride of the destination image in X dimension (in bytes)
* @param[in] out_step_x out_stride_x * number of elements along X processed per work item (in bytes)
* @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes)
@@ -232,6 +231,7 @@ __kernel void remap_nearest_neighbour_nhwc(
* @param[in] mapy_offset_first_element_in_bytes Offset of the first element in the remapping image
* @param[in] width Width of the input image
* @param[in] height Height of the input image
+ * @param[in] border_val Value to use for border around input tensor when in CONSTANT border is selected
*/
__kernel void remap_bilinear_nhwc(
TENSOR4D_DECLARATION(in),
diff --git a/src/core/CL/kernels/CLRemapKernel.cpp b/src/core/CL/kernels/CLRemapKernel.cpp
index 6edd744db7..7e3157c99d 100644
--- a/src/core/CL/kernels/CLRemapKernel.cpp
+++ b/src/core/CL/kernels/CLRemapKernel.cpp
@@ -32,8 +32,6 @@
#include "src/core/AccessWindowStatic.h"
#include "src/core/helpers/WindowHelpers.h"
-#include <algorithm>
-
namespace arm_compute
{
CLRemapKernel::CLRemapKernel()
@@ -54,11 +52,18 @@ void CLRemapKernel::set_constant_border(unsigned int idx, const PixelValue &cons
ICLKernel::add_argument<T>(idx, static_cast<T>(value));
}
-Status CLRemapKernel::validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, ITensorInfo *output, RemapInfo info)
+Status CLRemapKernel::validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, const ITensorInfo *output, RemapInfo info)
{
ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, map_x, map_y, output);
- ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8);
- ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8);
+ if(input->data_layout() == DataLayout::NCHW)
+ {
+ ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8);
+ }
+ else
+ {
+ ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::F16);
+ }
+ ARM_COMPUTE_RETURN_ERROR_ON_MSG(input->data_type() != output->data_type(), "Input/output have different data types");
ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(map_x, 1, DataType::F32);
ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(map_y, 1, DataType::F32);
ARM_COMPUTE_RETURN_ERROR_ON_MSG(info.policy == InterpolationPolicy::AREA, "Area interpolation is not supported!");
@@ -68,7 +73,8 @@ Status CLRemapKernel::validate(const ITensorInfo *input, const ITensorInfo *map_
void CLRemapKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, RemapInfo info)
{
- CLRemapKernel::validate(input->info(), map_x->info(), map_y->info(), output->info(), info);
+ ARM_COMPUTE_ERROR_ON_NULLPTR(input, map_x, map_y, output);
+ ARM_COMPUTE_ERROR_THROW_ON(CLRemapKernel::validate(input->info(), map_x->info(), map_y->info(), output->info(), info));
_input = input;
_output = output;
@@ -118,7 +124,19 @@ void CLRemapKernel::configure(const CLCompileContext &compile_context, const ICL
_kernel.setArg<cl_float>(idx++, input_height);
if(is_nhwc && is_constant_border)
{
- set_constant_border<uint8_t>(idx, info.constant_border_value);
+ switch(input->info()->data_type())
+ {
+ case DataType::U8:
+ set_constant_border<uint8_t>(idx, info.constant_border_value);
+ break;
+ case DataType::F16:
+ static_assert(sizeof(cl_half) == sizeof(half), "Half must be same size as cl_half");
+ static_assert(sizeof(cl_half) == 2, "Half must be 16 bit");
+ set_constant_border<half>(idx, info.constant_border_value);
+ break;
+ default:
+ ARM_COMPUTE_ERROR("Data Type not handled");
+ }
}
}
diff --git a/src/core/CL/kernels/CLRemapKernel.h b/src/core/CL/kernels/CLRemapKernel.h
index 1e3a4ad13f..93b0b4e660 100644
--- a/src/core/CL/kernels/CLRemapKernel.h
+++ b/src/core/CL/kernels/CLRemapKernel.h
@@ -49,28 +49,22 @@ public:
/** Initialize the kernel's input, output and border mode.
*
* @param[in] compile_context The compile context to be used.
- * @param[in] input Source tensor. Data types supported: U8.
+ * @param[in] input Source tensor. Data types supported: U8 (or F16 when layout is NHWC).
* @param[in] map_x Map for X coordinates. Data types supported: F32.
* @param[in] map_y Map for Y coordinates. Data types supported: F32.
- * @param[out] output Destination tensor. Data types supported: U8. All but the lowest two dimensions must be the same size as in the input tensor, i.e. remapping is only performed within the XY-plane.
+ * @param[out] output Destination tensor. Data types supported: Same as @p input. All but the lowest two dimensions must be the same size as in the input tensor, i.e. remapping is only performed within the XY-plane.
* @param[in] info RemapInfo struct:
* - policy Interpolation policy to use. Only NEAREST and BILINEAR are supported.
* - border_mode Border mode to use on the input tensor. Only CONSTANT and UNDEFINED are supported.
* - constant_border_value Constant value to use for borders if border_mode is set to CONSTANT.
*/
void configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, RemapInfo info);
- /** Validate the kernel's input, output and border mode.
+ /** Checks if the kernel's input, output and border mode will lead to a valid configuration of @ref CLRemapKernel
+ *
+ * Similar to @ref CLRemapKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, RemapInfo info)
*
- * @param[in] input Source tensor. Data types supported: U8.
- * @param[in] map_x Map for X coordinates. Data types supported: F32.
- * @param[in] map_y Map for Y coordinates. Data types supported: F32.
- * @param[out] output Destination tensor. Data types supported: U8. All but the lowest two dimensions must be the same size as in the input tensor, i.e. remapping is only performed within the XY-plane.
- * @param[in] info RemapInfo struct:
- * - policy Interpolation policy to use. Only NEAREST and BILINEAR are supported.
- * - border_mode Border mode to use on the input tensor. Only CONSTANT and UNDEFINED are supported.
- * - constant_border_value Constant value to use for borders if border_mode is set to CONSTANT.
*/
- static Status validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, ITensorInfo *output, RemapInfo info);
+ static Status validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, const ITensorInfo *output, RemapInfo info);
/** Function to set the constant value on fill border kernel depending on type.
*
* @param[in] idx Index of the kernel argument to set.
diff --git a/src/runtime/CL/functions/CLRemap.cpp b/src/runtime/CL/functions/CLRemap.cpp
index 0a1f864543..de9f0a0148 100644
--- a/src/runtime/CL/functions/CLRemap.cpp
+++ b/src/runtime/CL/functions/CLRemap.cpp
@@ -21,31 +21,45 @@
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*/
-#include "arm_compute/runtime/CL/functions/CLRemap.h"
+#include "arm_compute/runtime/CL/functions/CLRemap.h"
#include "arm_compute/core/CL/ICLTensor.h"
#include "arm_compute/core/Error.h"
-#include "arm_compute/core/PixelValue.h"
#include "arm_compute/core/TensorInfo.h"
#include "arm_compute/core/Validate.h"
#include "src/core/CL/kernels/CLFillBorderKernel.h"
#include "src/core/CL/kernels/CLRemapKernel.h"
-#include <utility>
-
-using namespace arm_compute;
+namespace arm_compute
+{
+void CLRemap::configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, InterpolationPolicy policy,
+ BorderMode border_mode, uint8_t constant_border_value)
+{
+ configure(compile_context, input, map_x, map_y, output, policy, border_mode, PixelValue{ constant_border_value });
+}
void CLRemap::configure(ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
{
+ configure(CLKernelLibrary::get().get_compile_context(), input, map_x, map_y, output, policy, border_mode, PixelValue{ constant_border_value });
+}
+
+void CLRemap::configure(ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value)
+{
configure(CLKernelLibrary::get().get_compile_context(), input, map_x, map_y, output, policy, border_mode, constant_border_value);
}
void CLRemap::configure(const CLCompileContext &compile_context, ICLTensor *input, const ICLTensor *map_x, const ICLTensor *map_y, ICLTensor *output, InterpolationPolicy policy,
- BorderMode border_mode,
- uint8_t constant_border_value)
+ BorderMode border_mode, PixelValue constant_border_value)
{
auto k = std::make_unique<CLRemapKernel>();
- k->configure(compile_context, input, map_x, map_y, output, RemapInfo{ policy, border_mode, PixelValue(constant_border_value) });
+ k->configure(compile_context, input, map_x, map_y, output, RemapInfo{ policy, border_mode, constant_border_value });
_kernel = std::move(k);
- _border_handler->configure(compile_context, input, _kernel->border_size(), border_mode, PixelValue(constant_border_value));
+ _border_handler->configure(compile_context, input, _kernel->border_size(), border_mode, constant_border_value);
+}
+
+Status CLRemap::validate(const ITensorInfo *input, const ITensorInfo *map_x, const ITensorInfo *map_y, const ITensorInfo *output, const InterpolationPolicy policy, const BorderMode border_mode,
+ PixelValue constant_border_value)
+{
+ return CLRemapKernel::validate(input, map_x, map_y, output, RemapInfo{ policy, border_mode, constant_border_value });
}
+} // namespace arm_compute
diff --git a/src/runtime/NEON/functions/NERemap.cpp b/src/runtime/NEON/functions/NERemap.cpp
index d9fd987480..dc2c4b4f91 100644
--- a/src/runtime/NEON/functions/NERemap.cpp
+++ b/src/runtime/NEON/functions/NERemap.cpp
@@ -36,6 +36,10 @@ namespace arm_compute
{
void NERemap::configure(ITensor *input, const ITensor *map_x, const ITensor *map_y, ITensor *output, InterpolationPolicy policy, BorderMode border_mode, uint8_t constant_border_value)
{
+ configure(input, map_x, map_y, output, policy, border_mode, PixelValue{ constant_border_value });
+}
+void NERemap::configure(ITensor *input, const ITensor *map_x, const ITensor *map_y, ITensor *output, InterpolationPolicy policy, BorderMode border_mode, PixelValue constant_border_value)
+{
ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8);
ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::U8);
ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(map_x, 1, DataType::F32);
@@ -43,7 +47,7 @@ void NERemap::configure(ITensor *input, const ITensor *map_x, const ITensor *map
ARM_COMPUTE_ERROR_ON_MSG(policy == InterpolationPolicy::AREA, "Area interpolation is not supported");
auto k = std::make_unique<NERemapKernel>();
- k->configure(input, map_x, map_y, output, policy, border_mode, constant_border_value);
+ k->configure(input, map_x, map_y, output, policy, border_mode, constant_border_value.get<uint8_t>());
_kernel = std::move(k);
}
} // namespace arm_compute
generated by cgit v1.2.1 (git 2.23.0) at 2024-07-19 01:21:10 +0000