COMPMID-1327: Add support for BBoxTransform operator in CL

Change-Id: I91865506166951b3bf7f06a0b2d4cde925cfefb6
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/153447
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Pablo Tello <pablo.tello@arm.com>

4 files changed, 293 insertions, 0 deletions

diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp
index 900cb04b1a..12a7c38dfd 100644
--- a/src/core/CL/CLKernelLibrary.cpp
+++ b/src/core/CL/CLKernelLibrary.cpp
@@ -164,6 +164,7 @@ const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map =
     { "bitwise_and", "bitwise_op.cl" },
     { "bitwise_xor", "bitwise_op.cl" },
     { "bitwise_not", "bitwise_op.cl" },
+    { "bounding_box_transform", "bounding_box_transform.cl" },
     { "channel_combine_NV", "channel_combine.cl" },
     { "channel_combine_RGB888", "channel_combine.cl" },
     { "channel_combine_RGBA8888", "channel_combine.cl" },
@@ -497,6 +498,10 @@ const std::map<std::string, std::string> CLKernelLibrary::_program_source_map =
 #include "./cl_kernels/bitwise_op.clembed"
     },
     {
+        "bounding_box_transform.cl",
+#include "./cl_kernels/bounding_box_transform.clembed"
+    },
+    {
         "canny.cl",
 #include "./cl_kernels/canny.clembed"
     },
diff --git a/src/core/CL/cl_kernels/bounding_box_transform.cl b/src/core/CL/cl_kernels/bounding_box_transform.cl
new file mode 100644
index 0000000000..a62635e052
--- /dev/null
+++ b/src/core/CL/cl_kernels/bounding_box_transform.cl
@@ -0,0 +1,118 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "helpers.h"
+
+#if defined(DATA_TYPE) && defined(WEIGHT_X) && defined(WEIGHT_Y) && defined(WEIGHT_W) && defined(WEIGHT_H) && defined(IMG_WIDTH) && defined(IMG_HEIGHT) && defined(BOX_FIELDS) // Check for compile time constants
+
+/** Perform a padded copy of input tensor to the output tensor. Padding values are defined at compile time
+ *
+ * @attention The following variables must be passed at compile time:
+ * -# -DDATA_TYPE = Tensor data type. Supported data types: F16/F32
+ * -# -DWEIGHT{X,Y,W,H}= Weights [wx, wy, ww, wh] for the deltas
+ * -# -DIMG_WIDTH= Original image width
+ * -# -DIMG_HEIGHT= Original image height
+ * -# -DBOX_FIELDS=Number of fields that are used to represent a box in boxes
+ *
+ * @param[in]  boxes_ptr                                Pointer to the boxes tensor. Supported data types: F16/F32
+ * @param[in]  boxes_stride_x                           Stride of the boxes tensor in X dimension (in bytes)
+ * @param[in]  boxes_step_x                             boxes_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  boxes_stride_y                           Stride of the boxes tensor in Y dimension (in bytes)
+ * @param[in]  boxes_step_y                             boxes_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  boxes_stride_z                           Stride of the source tensor in Z dimension (in bytes)
+ * @param[in]  boxes_step_z                             boxes_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in]  boxes_offset_first_element_in_bytes      The offset of the first element in the boxes tensor
+ * @param[out] pred_boxes_ptr                           Pointer to the predicted boxes. Supported data types: same as @p in_ptr
+ * @param[in]  pred_boxes_stride_x                      Stride of the predicted boxes in X dimension (in bytes)
+ * @param[in]  pred_boxes_step_x                        pred_boxes_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  pred_boxes_stride_y                      Stride of the predicted boxes in Y dimension (in bytes)
+ * @param[in]  pred_boxes_step_y                        pred_boxes_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  pred_boxes_stride_z                      Stride of the predicted boxes in Z dimension (in bytes)
+ * @param[in]  pred_boxes_step_z                        pred_boxes_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in]  pred_boxes_offset_first_element_in_bytes The offset of the first element in the predicted boxes
+ * @param[in]  deltas_ptr                               Pointer to the deltas tensor. Supported data types: same as @p in_ptr
+ * @param[in]  deltas_stride_x                          Stride of the deltas tensor in X dimension (in bytes)
+ * @param[in]  deltas_step_x                            deltas_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  deltas_stride_y                          Stride of the deltas tensor in Y dimension (in bytes)
+ * @param[in]  deltas_step_y                            deltas_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  deltas_stride_z                          Stride of the source tensor in Z dimension (in bytes)
+ * @param[in]  deltas_step_z                            deltas_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in]  deltas_offset_first_element_in_bytes     The offset of the first element in the deltas tensor
+ */
+__kernel void bounding_box_transform(
+    VECTOR_DECLARATION(boxes),
+    IMAGE_DECLARATION(pred_boxes),
+    IMAGE_DECLARATION(deltas))
+{
+    // Get pixels pointer
+    Vector boxes      = CONVERT_TO_VECTOR_STRUCT_NO_STEP(boxes);
+    Image  pred_boxes = CONVERT_TO_IMAGE_STRUCT(pred_boxes);
+    Image  deltas     = CONVERT_TO_IMAGE_STRUCT(deltas);
+
+    // Load delta and box values into registers
+    const DATA_TYPE one     = (DATA_TYPE)1.f;
+    const DATA_TYPE halfone = (DATA_TYPE)0.5f;
+
+    const int py = get_global_id(1); // box
+    VEC_DATA_TYPE(DATA_TYPE, 4)
+    delta = vload4(0, (__global DATA_TYPE *)deltas.ptr);
+    const VEC_DATA_TYPE(DATA_TYPE, 4)
+    box = vload4(0, (__global DATA_TYPE *)vector_offset(&boxes, BOX_FIELDS * py));
+
+    // Calculate width and centers of the old boxes
+    const VEC_DATA_TYPE(DATA_TYPE, 2)
+    dims = box.s23 - box.s01 + one;
+    const VEC_DATA_TYPE(DATA_TYPE, 2)
+    ctr = box.s01 + halfone * dims;
+    const VEC_DATA_TYPE(DATA_TYPE, 4)
+    weights = (VEC_DATA_TYPE(DATA_TYPE, 4))(WEIGHT_X, WEIGHT_Y, WEIGHT_W, WEIGHT_H);
+    delta /= weights;
+    delta.s23 = min(delta.s23, (DATA_TYPE)BBOX_XFORM_CLIP);
+
+    // Calculate widths and centers of the new boxes (translation + aspect ratio transformation)
+    const VEC_DATA_TYPE(DATA_TYPE, 2)
+    pred_ctr = delta.s01 * dims + ctr;
+    const VEC_DATA_TYPE(DATA_TYPE, 2)
+    pred_dims = exp(delta.s23) * dims;
+
+    // Useful vector constant definitions
+    const VEC_DATA_TYPE(DATA_TYPE, 4)
+    max_values = (VEC_DATA_TYPE(DATA_TYPE, 4))(IMG_WIDTH, IMG_HEIGHT, IMG_WIDTH, IMG_HEIGHT);
+    const VEC_DATA_TYPE(DATA_TYPE, 4)
+    sign = (VEC_DATA_TYPE(DATA_TYPE, 4))(-1, -1, 1, 1);
+    const VEC_DATA_TYPE(DATA_TYPE, 4)
+    min_values = 0;
+
+    // Calculate the coordinates of the new boxes
+    VEC_DATA_TYPE(DATA_TYPE, 4)
+    pred_box = pred_ctr.s0101 + sign * halfone * pred_dims.s0101;
+    pred_box = CLAMP(pred_box, min_values, max_values);
+#ifdef SCALE // Possibly scale the predicted boxes
+    pred_box *= SCALE;
+#endif // Possibly scale the predicted boxes
+
+    // Store them into the output
+    vstore4(pred_box, 0, (__global DATA_TYPE *)pred_boxes.ptr);
+}
+
+#endif // defined(DATA_TYPE) && defined(WEIGHT_X) && defined(WEIGHT_Y) && defined(WEIGHT_W) && defined(WEIGHT_H) && defined(IMG_WIDTH) && defined(IMG_HEIGHT) && defined(BOX_FIELDS)
diff --git a/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp b/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp
new file mode 100644
index 0000000000..09f3d33f5b
--- /dev/null
+++ b/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp
@@ -0,0 +1,127 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/core/CL/kernels/CLBoundingBoxTransformKernel.h"
+
+#include "arm_compute/core/AccessWindowStatic.h"
+#include "arm_compute/core/CL/CLHelpers.h"
+#include "arm_compute/core/CL/CLKernelLibrary.h"
+#include "arm_compute/core/CL/CLValidate.h"
+#include "arm_compute/core/CL/ICLArray.h"
+#include "arm_compute/core/CL/ICLTensor.h"
+#include "arm_compute/core/CL/OpenCL.h"
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/TensorInfo.h"
+#include "arm_compute/core/Utils.h"
+#include "arm_compute/core/Window.h"
+
+namespace arm_compute
+{
+namespace
+{
+Status validate_arguments(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas)
+{
+    ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(boxes, pred_boxes, deltas);
+    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_NOT_IN(boxes, DataType::F32, DataType::F16);
+    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_NOT_IN(deltas, DataType::F32, DataType::F16);
+    ARM_COMPUTE_RETURN_ERROR_ON(deltas->tensor_shape()[1] != boxes->tensor_shape()[1]);
+    ARM_COMPUTE_RETURN_ERROR_ON(deltas->tensor_shape()[0] % 4 != 0);
+    ARM_COMPUTE_RETURN_ERROR_ON(boxes->tensor_shape()[0] != 4);
+    ARM_COMPUTE_RETURN_ERROR_ON(deltas->num_dimensions() > 2);
+    ARM_COMPUTE_RETURN_ERROR_ON(boxes->num_dimensions() > 2);
+
+    if(pred_boxes->total_size() > 0)
+    {
+        ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DIMENSIONS(pred_boxes->tensor_shape(), deltas->tensor_shape());
+        ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(pred_boxes, deltas);
+        ARM_COMPUTE_RETURN_ERROR_ON(pred_boxes->num_dimensions() > 2);
+    }
+    return Status{};
+}
+} // namespace
+
+CLBoundingBoxTransformKernel::CLBoundingBoxTransformKernel()
+    : _boxes(nullptr), _pred_boxes(nullptr), _deltas(nullptr)
+{
+}
+
+void CLBoundingBoxTransformKernel::configure(const ICLTensor *boxes, ICLTensor *pred_boxes, const ICLTensor *deltas, const BoundingBoxTransformInfo &info)
+{
+    ARM_COMPUTE_ERROR_ON_NULLPTR(boxes, pred_boxes, deltas);
+    auto_init_if_empty(*pred_boxes->info(), *deltas->info());
+
+    // Set instance variables
+    _boxes      = boxes;
+    _pred_boxes = pred_boxes;
+    _deltas     = deltas;
+
+    // Get image height and widht (rescaled)
+    const int img_h = floor(info.img_height() / info.scale() + 0.5f);
+    const int img_w = floor(info.img_width() / info.scale() + 0.5f);
+
+    // Set build options
+    CLBuildOptions build_opts;
+    build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(deltas->info()->data_type()));
+    build_opts.add_option("-DWEIGHT_X=" + float_to_string_with_full_precision(info.weights()[0]));
+    build_opts.add_option("-DWEIGHT_Y=" + float_to_string_with_full_precision(info.weights()[1]));
+    build_opts.add_option("-DWEIGHT_W=" + float_to_string_with_full_precision(info.weights()[2]));
+    build_opts.add_option("-DWEIGHT_H=" + float_to_string_with_full_precision(info.weights()[3]));
+    build_opts.add_option("-DBBOX_XFORM_CLIP=" + float_to_string_with_full_precision(info.bbox_xform_clip()));
+    build_opts.add_option("-DIMG_WIDTH=" + support::cpp11::to_string(img_w));
+    build_opts.add_option("-DIMG_HEIGHT=" + support::cpp11::to_string(img_h));
+    build_opts.add_option("-DBOX_FIELDS=" + support::cpp11::to_string(4));
+    build_opts.add_option_if(info.apply_scale(), "-DSCALE=" + float_to_string_with_full_precision(info.scale()));
+
+    // Create kernel
+    _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("bounding_box_transform", build_opts.options()));
+
+    // Since the number of columns is a multiple of 4 by definition, we don't need to pad the tensor
+    const unsigned int num_elems_processed_per_iteration = 4;
+    Window             win                               = calculate_max_window(*deltas->info(), Steps(num_elems_processed_per_iteration));
+    ICLKernel::configure_internal(win);
+}
+
+Status CLBoundingBoxTransformKernel::validate(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas, const BoundingBoxTransformInfo &info)
+{
+    ARM_COMPUTE_UNUSED(info);
+    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(boxes, pred_boxes, deltas));
+    return Status{};
+}
+
+void CLBoundingBoxTransformKernel::run(const Window &window, cl::CommandQueue &queue)
+{
+    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
+    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window);
+
+    Window slice = window.first_slice_window_2D();
+
+    // Set arguments
+    unsigned int idx = 0;
+    add_1D_tensor_argument(idx, _boxes, slice);
+    add_2D_tensor_argument(idx, _pred_boxes, slice);
+    add_2D_tensor_argument(idx, _deltas, slice);
+
+    // Note that we don't need to loop over the slices, as we are sure that we are dealing with all 2D tensors
+    enqueue(queue, *this, slice);
+}
+} // namespace arm_compute
diff --git a/src/runtime/CL/functions/CLBoundingBoxTransform.cpp b/src/runtime/CL/functions/CLBoundingBoxTransform.cpp
new file mode 100644
index 0000000000..46a6b8ea96
--- /dev/null
+++ b/src/runtime/CL/functions/CLBoundingBoxTransform.cpp
@@ -0,0 +1,43 @@
+/*
+ * Copyright (c) 2018 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/runtime/CL/functions/CLBoundingBoxTransform.h"
+
+#include "arm_compute/core/CL/kernels/CLBoundingBoxTransformKernel.h"
+#include "support/ToolchainSupport.h"
+
+namespace arm_compute
+{
+void CLBoundingBoxTransform::configure(const ICLTensor *boxes, ICLTensor *pred_boxes, const ICLTensor *deltas, const BoundingBoxTransformInfo &info)
+{
+    // Configure Bounding Box kernel
+    auto k = arm_compute::support::cpp14::make_unique<CLBoundingBoxTransformKernel>();
+    k->configure(boxes, pred_boxes, deltas, info);
+    _kernel = std::move(k);
+}
+
+Status CLBoundingBoxTransform::validate(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas, const BoundingBoxTransformInfo &info)
+{
+    return CLBoundingBoxTransformKernel::validate(boxes, pred_boxes, deltas, info);
+}
+} // namespace arm_compute