COMPMID-2246: Implement NEBoundingBoxTransform.

Change-Id: I69faa21ccdc295f5369450a12e6e217d873349e3
Signed-off-by: Pablo Tello <pablo.tello@arm.com>
Reviewed-on: https://review.mlplatform.org/c/1854
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>

7 files changed, 537 insertions, 0 deletions

diff --git a/arm_compute/core/NEON/NEKernels.h b/arm_compute/core/NEON/NEKernels.h
index 5e3389744f..f86a1d0507 100644
--- a/arm_compute/core/NEON/NEKernels.h
+++ b/arm_compute/core/NEON/NEKernels.h
@@ -37,6 +37,7 @@
 #include "arm_compute/core/NEON/kernels/NEBitwiseNotKernel.h"
 #include "arm_compute/core/NEON/kernels/NEBitwiseOrKernel.h"
 #include "arm_compute/core/NEON/kernels/NEBitwiseXorKernel.h"
+#include "arm_compute/core/NEON/kernels/NEBoundingBoxTransformKernel.h"
 #include "arm_compute/core/NEON/kernels/NEBox3x3Kernel.h"
 #include "arm_compute/core/NEON/kernels/NECannyEdgeKernel.h"
 #include "arm_compute/core/NEON/kernels/NEChannelCombineKernel.h"
diff --git a/arm_compute/core/NEON/kernels/NEBoundingBoxTransformKernel.h b/arm_compute/core/NEON/kernels/NEBoundingBoxTransformKernel.h
new file mode 100644
index 0000000000..0f0a1f6a2d
--- /dev/null
+++ b/arm_compute/core/NEON/kernels/NEBoundingBoxTransformKernel.h
@@ -0,0 +1,93 @@
+/*
+ * Copyright (c) 2019 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.
+ */
+#ifndef __ARM_COMPUTE_NEBOUNDINGBOXTRANSFORMKERNEL_H__
+#define __ARM_COMPUTE_NEBOUNDINGBOXTRANSFORMKERNEL_H__
+
+#include "arm_compute/core/NEON/INEKernel.h"
+
+namespace arm_compute
+{
+class ITensor;
+
+/** Interface for the bounding box kernel */
+class NEBoundingBoxTransformKernel : public INEKernel
+{
+public:
+    const char *name() const override
+    {
+        return "NEBoundingBoxTransformKernel";
+    }
+
+    /** Default constructor */
+    NEBoundingBoxTransformKernel();
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    NEBoundingBoxTransformKernel(const NEBoundingBoxTransformKernel &) = delete;
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    NEBoundingBoxTransformKernel &operator=(const NEBoundingBoxTransformKernel &) = delete;
+    /** Allow instances of this class to be moved */
+    NEBoundingBoxTransformKernel(NEBoundingBoxTransformKernel &&) = default;
+    /** Allow instances of this class to be moved */
+    NEBoundingBoxTransformKernel &operator=(NEBoundingBoxTransformKernel &&) = default;
+    /** Default destructor */
+    ~NEBoundingBoxTransformKernel() = default;
+
+    /** Set the input and output tensors.
+     *
+     * @param[in]  boxes      Source tensor. Bounding box proposals in pixel coordinates. Size(M, 4), format [x1, y1, x2, y2]. Data types supported: F16/F32.
+     * @param[out] pred_boxes Destination tensor. Pixel coordinates of the transformed bounding boxes. Size (M, 4*K), format [x1, y1, x2, y2]. Data types supported: Same as @p input
+     * @param[in]  deltas     Bounding box translations and scales. Size (M, 4*K), format [dx, dy, dw, dh], K  is the number of classes. Data types supported: Same as @p input
+     * @param[in]  info       Contains BoundingBox operation information described in @ref BoundingBoxTransformInfo.
+     *
+     * @note Only single image prediction is supported. Height and Width (and scale) of the image will be contained in the BoundingBoxTransformInfo struct.
+     *
+     */
+    void configure(const ITensor *boxes, ITensor *pred_boxes, const ITensor *deltas, const BoundingBoxTransformInfo &info);
+
+    /** Static function to check if given info will lead to a valid configuration of @ref CLBoundingBoxTransform
+     *
+     * @param[in] boxes      Source tensor info. Bounding box proposals in pixel coordinates. Size(M, 4), format [x1, y1, x2, y2]. Data types supported: F16/F32.
+     * @param[in] pred_boxes Destination tensor info. Pixel coordinates of the transformed bounding boxes. Size (M, 4*K), format [x1, y1, x2, y2]. Data types supported: Same as @p input
+     * @param[in] deltas     Bounding box translations and scales. Size (M, 4*K), format [dx, dy, dw, dh], K  is the number of classes. Data types supported: Same as @p input
+     * @param[in] info       Contains BoundingBox operation information described in @ref BoundingBoxTransformInfo.
+     *
+     * @note Only single image prediction is supported. Height and Width (and scale) of the image will be contained in the BoundingBoxTransformInfo struct.
+     *
+     * @return a Status
+     */
+    static Status validate(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas, const BoundingBoxTransformInfo &info);
+
+    // Inherited methods overridden:
+    void run(const Window &window, const ThreadInfo &info) override;
+
+private:
+    template <typename T>
+    void internal_run(const Window &window, const ThreadInfo &info);
+
+    const ITensor           *_boxes;
+    ITensor                 *_pred_boxes;
+    const ITensor           *_deltas;
+    BoundingBoxTransformInfo _bbinfo;
+};
+} // namespace arm_compute
+#endif /*__ARM_COMPUTE_NEBOUNDINGBOXTRANSFORMKERNEL_H__ */
diff --git a/arm_compute/runtime/NEON/NEFunctions.h b/arm_compute/runtime/NEON/NEFunctions.h
index 515e4070f4..9dd7e5e5e7 100644
--- a/arm_compute/runtime/NEON/NEFunctions.h
+++ b/arm_compute/runtime/NEON/NEFunctions.h
@@ -37,6 +37,7 @@
 #include "arm_compute/runtime/NEON/functions/NEBitwiseNot.h"
 #include "arm_compute/runtime/NEON/functions/NEBitwiseOr.h"
 #include "arm_compute/runtime/NEON/functions/NEBitwiseXor.h"
+#include "arm_compute/runtime/NEON/functions/NEBoundingBoxTransform.h"
 #include "arm_compute/runtime/NEON/functions/NEBox3x3.h"
 #include "arm_compute/runtime/NEON/functions/NECannyEdge.h"
 #include "arm_compute/runtime/NEON/functions/NECast.h"
diff --git a/arm_compute/runtime/NEON/functions/NEBoundingBoxTransform.h b/arm_compute/runtime/NEON/functions/NEBoundingBoxTransform.h
new file mode 100644
index 0000000000..1f512b938b
--- /dev/null
+++ b/arm_compute/runtime/NEON/functions/NEBoundingBoxTransform.h
@@ -0,0 +1,67 @@
+/*
+ * Copyright (c) 2019 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.
+ */
+#ifndef __ARM_COMPUTE_NEBOUNDINGBOXTRANSOFORM_H__
+#define __ARM_COMPUTE_NEBOUNDINGBOXTRANSOFORM_H__
+
+#include "arm_compute/core/NEON/kernels/NEBoundingBoxTransformKernel.h"
+#include "arm_compute/runtime/NEON/INESimpleFunction.h"
+
+namespace arm_compute
+{
+class ITensor;
+
+/** Basic function to run @ref NEBoundingBoxTransformKernel.
+ *
+ * This function calls the following Neon kernels:
+ * -# @ref NEBoundingBoxTransformKernel
+ */
+class NEBoundingBoxTransform : public INESimpleFunction
+{
+public:
+    /** Set the input and output tensors.
+     *
+     * @param[in]  boxes      Source tensor. Bounding box proposals in pixel coordinates. Size(M, 4), format [x1, y1, x2, y2]. Data types supported: F16/F32.
+     * @param[out] pred_boxes Destination tensor. Pixel coordinates of the transformed bounding boxes. Size (M, 4*K), format [x1, y1, x2, y2]. Data types supported: Same as @p input
+     * @param[in]  deltas     Bounding box translations and scales. Size (M, 4*K), format [dx, dy, dw, dh], K  is the number of classes. Data types supported: Same as @p input
+     * @param[in]  info       Contains BoundingBox operation information described in @ref BoundingBoxTransformInfo.
+     *
+     * @note Only single image prediction is supported. Height and Width (and scale) of the image will be contained in the BoundingBoxTransformInfo struct.
+     */
+    void configure(const ITensor *boxes, ITensor *pred_boxes, const ITensor *deltas, const BoundingBoxTransformInfo &info);
+
+    /** Static function to check if given info will lead to a valid configuration of @ref NEBoundingBoxTransform
+     *
+     * @param[in] boxes      Source tensor info. Bounding box proposals in pixel coordinates. Size(M, 4), format [x1, y1, x2, y2]. Data types supported: F16/F32.
+     * @param[in] pred_boxes Destination tensor info. Pixel coordinates of the transformed bounding boxes. Size (M, 4*K), format [x1, y1, x2, y2]. Data types supported: Same as @p input
+     * @param[in] deltas     Bounding box translations and scales. Size (M, 4*K), format [dx, dy, dw, dh], K  is the number of classes. Data types supported: Same as @p input
+     * @param[in] info       Contains BoundingBox operation information described in @ref BoundingBoxTransformInfo.
+     *
+     * @note Only single image prediction is supported. Height and Width (and scale) of the image will be contained in the BoundingBoxTransformInfo struct.
+     *
+     * @return a Status
+     */
+    static Status validate(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas, const BoundingBoxTransformInfo &info);
+};
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_CLBOUNDINGBOXTRANSFORM_H__ */
diff --git a/src/core/NEON/kernels/NEBoundingBoxTransformKernel.cpp b/src/core/NEON/kernels/NEBoundingBoxTransformKernel.cpp
new file mode 100644
index 0000000000..cfd3e708e5
--- /dev/null
+++ b/src/core/NEON/kernels/NEBoundingBoxTransformKernel.cpp
@@ -0,0 +1,191 @@
+/*
+ * Copyright (c) 2019 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/NEON/kernels/NEBoundingBoxTransformKernel.h"
+
+#include "arm_compute/core/AccessWindowStatic.h"
+#include "arm_compute/core/CPP/Validate.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"
+
+#include <arm_neon.h>
+
+namespace arm_compute
+{
+namespace
+{
+Status validate_arguments(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas, const BoundingBoxTransformInfo &info)
+{
+    ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(boxes, pred_boxes, deltas);
+    ARM_COMPUTE_RETURN_ERROR_ON_CPU_F16_UNSUPPORTED(boxes);
+    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);
+    }
+    ARM_COMPUTE_RETURN_ERROR_ON(info.scale() <= 0);
+    return Status{};
+}
+
+std::pair<Status, Window> validate_and_configure_window(ITensorInfo *boxes, ITensorInfo *pred_boxes, ITensorInfo *deltas, const BoundingBoxTransformInfo &bb_info)
+{
+    ARM_COMPUTE_ERROR_ON_NULLPTR(boxes, pred_boxes);
+
+    auto_init_if_empty(*pred_boxes, *deltas);
+
+    const unsigned int num_boxes = boxes->dimension(1);
+
+    Window window;
+    window.set(Window::DimX, Window::Dimension(0, 1u));
+    window.set(Window::DimY, Window::Dimension(0, num_boxes));
+
+    AccessWindowHorizontal input_access(boxes, 0, 4u);
+    AccessWindowHorizontal output_access(pred_boxes, 0, 4u);
+
+    const bool window_changed = update_window_and_padding(window, input_access, output_access);
+    output_access.set_valid_region(window, ValidRegion(Coordinates(), pred_boxes->tensor_shape()));
+
+    Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
+    return std::make_pair(err, window);
+}
+
+} // namespace
+
+NEBoundingBoxTransformKernel::NEBoundingBoxTransformKernel()
+    : _boxes(nullptr), _pred_boxes(nullptr), _deltas(nullptr), _bbinfo(0, 0, 0)
+{
+}
+
+void NEBoundingBoxTransformKernel::configure(const ITensor *boxes, ITensor *pred_boxes, const ITensor *deltas, const BoundingBoxTransformInfo &info)
+{
+    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(boxes->info(), pred_boxes->info(), deltas->info(), info));
+
+    // Configure kernel window
+    auto win_config = validate_and_configure_window(boxes->info(), pred_boxes->info(), deltas->info(), info);
+    ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
+
+    // Set instance variables
+    _boxes      = boxes;
+    _pred_boxes = pred_boxes;
+    _deltas     = deltas;
+    _bbinfo     = info;
+
+    INEKernel::configure(win_config.second);
+}
+
+Status NEBoundingBoxTransformKernel::validate(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas, const BoundingBoxTransformInfo &info)
+{
+    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(boxes, pred_boxes, deltas, info));
+    return Status{};
+}
+
+template <typename T>
+void NEBoundingBoxTransformKernel::internal_run(const Window &window, const ThreadInfo &info)
+{
+    const size_t num_classes  = _deltas->info()->tensor_shape()[0] >> 2;
+    const size_t deltas_width = _deltas->info()->tensor_shape()[0];
+    const int    img_h        = floor(_bbinfo.img_height() / _bbinfo.scale() + 0.5f);
+    const int    img_w        = floor(_bbinfo.img_width() / _bbinfo.scale() + 0.5f);
+
+    const auto scale_after  = (_bbinfo.apply_scale() ? T(_bbinfo.scale()) : T(1));
+    const auto scale_before = T(_bbinfo.scale());
+    ARM_COMPUTE_ERROR_ON(scale_before <= 0);
+    const auto offset = (_bbinfo.correct_transform_coords() ? T(1.f) : T(0.f));
+
+    auto pred_ptr  = reinterpret_cast<T *>(_pred_boxes->buffer() + _pred_boxes->info()->offset_first_element_in_bytes());
+    auto delta_ptr = reinterpret_cast<T *>(_deltas->buffer() + _deltas->info()->offset_first_element_in_bytes());
+
+    Iterator box_it(_boxes, window);
+    execute_window_loop(window, [&](const Coordinates & id)
+    {
+        const auto ptr    = reinterpret_cast<T *>(box_it.ptr());
+        const auto b0     = *ptr;
+        const auto b1     = *(ptr + 1);
+        const auto b2     = *(ptr + 2);
+        const auto b3     = *(ptr + 3);
+        const T    width  = (b2 / scale_before) - (b0 / scale_before) + T(1.f);
+        const T    height = (b3 / scale_before) - (b1 / scale_before) + T(1.f);
+        const T    ctr_x  = (b0 / scale_before) + T(0.5f) * width;
+        const T    ctr_y  = (b1 / scale_before) + T(0.5f) * height;
+        for(size_t j = 0; j < num_classes; ++j)
+        {
+            // Extract deltas
+            const size_t delta_id = id.y() * deltas_width + 4u * j;
+            const T      dx       = delta_ptr[delta_id] / T(_bbinfo.weights()[0]);
+            const T      dy       = delta_ptr[delta_id + 1] / T(_bbinfo.weights()[1]);
+            T            dw       = delta_ptr[delta_id + 2] / T(_bbinfo.weights()[2]);
+            T            dh       = delta_ptr[delta_id + 3] / T(_bbinfo.weights()[3]);
+            // Clip dw and dh
+            dw = std::min(dw, T(_bbinfo.bbox_xform_clip()));
+            dh = std::min(dh, T(_bbinfo.bbox_xform_clip()));
+            // Determine the predictions
+            const T pred_ctr_x = dx * width + ctr_x;
+            const T pred_ctr_y = dy * height + ctr_y;
+            const T pred_w     = T(std::exp(dw)) * width;
+            const T pred_h     = T(std::exp(dh)) * height;
+            // Store the prediction into the output tensor
+            pred_ptr[delta_id]     = scale_after * utility::clamp<T>(pred_ctr_x - T(0.5f) * pred_w, T(0), T(img_w - 1));
+            pred_ptr[delta_id + 1] = scale_after * utility::clamp<T>(pred_ctr_y - T(0.5f) * pred_h, T(0), T(img_h - 1));
+            pred_ptr[delta_id + 2] = scale_after * utility::clamp<T>(pred_ctr_x + T(0.5f) * pred_w - offset, T(0), T(img_w - 1));
+            pred_ptr[delta_id + 3] = scale_after * utility::clamp<T>(pred_ctr_y + T(0.5f) * pred_h - offset, T(0), T(img_h - 1));
+        }
+    },
+    box_it);
+}
+
+void NEBoundingBoxTransformKernel::run(const Window &window, const ThreadInfo &info)
+{
+    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
+    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window);
+    switch(_boxes->info()->data_type())
+    {
+        case DataType::F32:
+        {
+            internal_run<float>(window, info);
+            break;
+        }
+#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
+        case DataType::F16:
+        {
+            internal_run<float16_t>(window, info);
+            break;
+        }
+#endif // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
+        default:
+        {
+            ARM_COMPUTE_ERROR("Data type not supported");
+        }
+    }
+}
+} // namespace arm_compute
diff --git a/src/runtime/NEON/functions/NEBoundingBoxTransform.cpp b/src/runtime/NEON/functions/NEBoundingBoxTransform.cpp
new file mode 100644
index 0000000000..818e228d52
--- /dev/null
+++ b/src/runtime/NEON/functions/NEBoundingBoxTransform.cpp
@@ -0,0 +1,42 @@
+/*
+ * Copyright (c) 2019 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/NEON/functions/NEBoundingBoxTransform.h"
+
+#include "support/ToolchainSupport.h"
+
+namespace arm_compute
+{
+void NEBoundingBoxTransform::configure(const ITensor *boxes, ITensor *pred_boxes, const ITensor *deltas, const BoundingBoxTransformInfo &info)
+{
+    // Configure Bounding Box kernel
+    auto k = arm_compute::support::cpp14::make_unique<NEBoundingBoxTransformKernel>();
+    k->configure(boxes, pred_boxes, deltas, info);
+    _kernel = std::move(k);
+}
+
+Status NEBoundingBoxTransform::validate(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas, const BoundingBoxTransformInfo &info)
+{
+    return NEBoundingBoxTransformKernel::validate(boxes, pred_boxes, deltas, info);
+}
+} // namespace arm_compute
diff --git a/tests/validation/NEON/BoundingBoxTransform.cpp b/tests/validation/NEON/BoundingBoxTransform.cpp
new file mode 100644
index 0000000000..070ea8b3ba
--- /dev/null
+++ b/tests/validation/NEON/BoundingBoxTransform.cpp
@@ -0,0 +1,142 @@
+/*
+ * Copyright (c) 2019 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/Types.h"
+
+#include "arm_compute/runtime/NEON/functions/NEBoundingBoxTransform.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "tests/Globals.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/datasets/ShapeDatasets.h"
+#include "tests/framework/Macros.h"
+#include "tests/framework/datasets/Datasets.h"
+#include "tests/validation/Validation.h"
+#include "tests/validation/fixtures/BoundingBoxTransformFixture.h"
+#include "utils/TypePrinter.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+RelativeTolerance<float> relative_tolerance_f32(0.01f);
+AbsoluteTolerance<float> absolute_tolerance_f32(0.001f);
+#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
+RelativeTolerance<half>  relative_tolerance_f16(half(0.2));
+AbsoluteTolerance<float> absolute_tolerance_f16(half(0.02f));
+#endif // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
+
+// *INDENT-OFF*
+// clang-format off
+const auto BboxInfoDataset = framework::dataset::make("BboxInfo", { BoundingBoxTransformInfo(20U, 20U, 2U, true),
+                                                                    BoundingBoxTransformInfo(128U, 128U, 4U, true),
+                                                                    BoundingBoxTransformInfo(800U, 600U, 1U, false),
+                                                                    BoundingBoxTransformInfo(800U, 600U, 2U, true, { { 1.0, 0.5, 1.5, 2.0 } }),
+                                                                    BoundingBoxTransformInfo(800U, 600U, 4U, false, { { 1.0, 0.5, 1.5, 2.0 } }),
+                                                                    BoundingBoxTransformInfo(800U, 600U, 4U, false, { { 1.0, 0.5, 1.5, 2.0 } }, true)
+                                                                  });
+
+const auto DeltaDataset = framework::dataset::make("DeltasShape", { TensorShape(36U, 1U),
+                                                                    TensorShape(36U, 2U),
+                                                                    TensorShape(36U, 2U),
+                                                                    TensorShape(40U, 1U),
+                                                                    TensorShape(40U, 20U),
+                                                                    TensorShape(40U, 100U),
+                                                                    TensorShape(40U, 200U)
+                                                                  });
+// clang-format on
+// *INDENT-ON*
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(BBoxTransform)
+
+// *INDENT-OFF*
+// clang-format off
+DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip(
+               framework::dataset::make("BoxesInfo", { TensorInfo(TensorShape(4U, 128U), 1, DataType::F32),
+                                                       TensorInfo(TensorShape(5U, 128U), 1, DataType::F32), // Wrong number of box fields
+                                                       TensorInfo(TensorShape(4U, 128U), 1, DataType::F16), // Wrong data type
+                                                       TensorInfo(TensorShape(4U, 128U), 1, DataType::F32), // Wrong number of classes
+                                                       TensorInfo(TensorShape(4U, 128U), 1, DataType::F32),  // Deltas and predicted boxes have different dimensions
+                                                       TensorInfo(TensorShape(4U, 128U), 1, DataType::F32)}),  // Scaling is zero
+               framework::dataset::make("PredBoxesInfo",{ TensorInfo(TensorShape(128U, 128U), 1, DataType::F32),
+                                                          TensorInfo(TensorShape(128U, 128U), 1, DataType::F32),
+                                                          TensorInfo(TensorShape(127U, 128U), 1, DataType::F32),
+                                                          TensorInfo(TensorShape(128U, 100U), 1, DataType::F32),
+                                                          TensorInfo(TensorShape(128U, 100U), 1, DataType::F32),
+                                                          TensorInfo(TensorShape(128U, 128U), 1, DataType::F32)})),
+               framework::dataset::make("DeltasInfo", { TensorInfo(TensorShape(128U, 128U), 1, DataType::F32),
+                                                        TensorInfo(TensorShape(128U, 128U), 1, DataType::F32),
+                                                        TensorInfo(TensorShape(127U, 128U), 1, DataType::F32),
+                                                        TensorInfo(TensorShape(128U, 100U), 1, DataType::F32),
+                                                        TensorInfo(TensorShape(128U, 128U), 1, DataType::F32),
+                                                        TensorInfo(TensorShape(128U, 128U), 1, DataType::F32)})),
+               framework::dataset::make("BoundingBoxTransofmInfo", { BoundingBoxTransformInfo(800.f, 600.f, 1.f),
+                                                                     BoundingBoxTransformInfo(800.f, 600.f, 1.f),
+                                                                     BoundingBoxTransformInfo(800.f, 600.f, 1.f),
+                                                                     BoundingBoxTransformInfo(800.f, 600.f, 1.f),
+                                                                     BoundingBoxTransformInfo(800.f, 600.f, 0.f)})),
+               framework::dataset::make("Expected", { true, false, false, false, false, false})),
+               boxes_info, pred_boxes_info, deltas_info, bbox_info, expected)
+{
+    ARM_COMPUTE_EXPECT(bool(NEBoundingBoxTransform::validate(&boxes_info.clone()->set_is_resizable(true), &pred_boxes_info.clone()->set_is_resizable(true), &deltas_info.clone()->set_is_resizable(true), bbox_info)) == expected, framework::LogLevel::ERRORS);
+}
+// clang-format on
+// *INDENT-ON*
+
+template <typename T>
+using NEBoundingBoxTransformFixture = BoundingBoxTransformFixture<Tensor, Accessor, NEBoundingBoxTransform, T>;
+
+TEST_SUITE(Float)
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(BoundingBox, NEBoundingBoxTransformFixture<float>, framework::DatasetMode::ALL,
+                       combine(combine(DeltaDataset, BboxInfoDataset), framework::dataset::make("DataType", { DataType::F32 })))
+{
+    // Validate output
+    validate(Accessor(_target), _reference, relative_tolerance_f32, 0.f, absolute_tolerance_f32);
+}
+TEST_SUITE_END() // FP32
+
+#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
+TEST_SUITE(FP16)
+FIXTURE_DATA_TEST_CASE(BoundingBox, NEBoundingBoxTransformFixture<half>, framework::DatasetMode::ALL,
+                       combine(combine(DeltaDataset, BboxInfoDataset), framework::dataset::make("DataType", { DataType::F16 })))
+{
+    // Validate output
+    validate(Accessor(_target), _reference, relative_tolerance_f16, 0.03f, absolute_tolerance_f16);
+}
+TEST_SUITE_END() // FP16
+#endif           // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
+
+TEST_SUITE_END() // Float
+
+TEST_SUITE_END() // BBoxTransform
+TEST_SUITE_END() // NEON
+} // namespace validation
+} // namespace test
+} // namespace arm_compute