COMPMID-1451: Fixes for BoundingBoxTransform

- Fixing a bug for which we did not scale the boxes before transforming them
- Adding the correct_transform_coords option to BoundingBoxTransformInfo

Change-Id: I40281254bcf87e7c8583c119e99562414fe59822

2 files changed, 18 insertions, 9 deletions

diff --git a/src/core/CL/cl_kernels/bounding_box_transform.cl b/src/core/CL/cl_kernels/bounding_box_transform.cl
index 14a0fadc2f..097235549b 100644
--- a/src/core/CL/cl_kernels/bounding_box_transform.cl
+++ b/src/core/CL/cl_kernels/bounding_box_transform.cl
@@ -23,7 +23,7 @@
  */
 #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
+#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) && defined(SCALE_BEFORE) // Check for compile time constants
 
 /** Perform a padded copy of input tensor to the output tensor. Padding values are defined at compile time
  *
@@ -74,10 +74,12 @@ __kernel void bounding_box_transform(
     const DATA_TYPE halfone = (DATA_TYPE)0.5f;
 
     const int py = get_global_id(1); // box
+    const VEC_DATA_TYPE(DATA_TYPE, 4)
+    scale_before = (VEC_DATA_TYPE(DATA_TYPE, 4))SCALE_BEFORE;
     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));
+    box = vload4(0, (__global DATA_TYPE *)vector_offset(&boxes, BOX_FIELDS * py)) / scale_before;
 
     // Calculate width and centers of the old boxes
     const VEC_DATA_TYPE(DATA_TYPE, 2)
@@ -106,13 +108,16 @@ __kernel void bounding_box_transform(
     // Calculate the coordinates of the new boxes
     VEC_DATA_TYPE(DATA_TYPE, 4)
     pred_box = pred_ctr.s0101 + sign * halfone * pred_dims.s0101;
+#ifdef OFFSET // Possibly adjust the predicted boxes
+    pred_box.s23 -= one;
+#endif // Possibly adjust the predicted boxes
     pred_box = CLAMP(pred_box, min_values, max_values);
-#ifdef SCALE // Possibly scale the predicted boxes
-    pred_box *= (VEC_DATA_TYPE(DATA_TYPE, 4))SCALE;
+#ifdef SCALE_AFTER // Possibly scale the predicted boxes
+    pred_box *= (VEC_DATA_TYPE(DATA_TYPE, 4))SCALE_AFTER;
 #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)
+#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) && defined(SCALE_BEFORE)
diff --git a/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp b/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp
index 09f3d33f5b..bff28e3ed9 100644
--- a/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp
+++ b/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp
@@ -39,7 +39,7 @@ namespace arm_compute
 {
 namespace
 {
-Status validate_arguments(const ITensorInfo *boxes, const ITensorInfo *pred_boxes, const ITensorInfo *deltas)
+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_DATA_TYPE_NOT_IN(boxes, DataType::F32, DataType::F16);
@@ -56,6 +56,7 @@ Status validate_arguments(const ITensorInfo *boxes, const ITensorInfo *pred_boxe
         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{};
 }
 } // namespace
@@ -70,6 +71,8 @@ void CLBoundingBoxTransformKernel::configure(const ICLTensor *boxes, ICLTensor *
     ARM_COMPUTE_ERROR_ON_NULLPTR(boxes, pred_boxes, deltas);
     auto_init_if_empty(*pred_boxes->info(), *deltas->info());
 
+    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(boxes->info(), pred_boxes->info(), deltas->info(), info));
+
     // Set instance variables
     _boxes      = boxes;
     _pred_boxes = pred_boxes;
@@ -90,7 +93,9 @@ void CLBoundingBoxTransformKernel::configure(const ICLTensor *boxes, ICLTensor *
     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()));
+    build_opts.add_option("-DSCALE_BEFORE=" + float_to_string_with_full_precision(info.scale()));
+    build_opts.add_option_if(info.apply_scale(), "-DSCALE_AFTER=" + float_to_string_with_full_precision(info.scale()));
+    build_opts.add_option_if(info.correct_transform_coords(), "-DOFFSET=1");
 
     // Create kernel
     _kernel = static_cast<cl::Kernel>(CLKernelLibrary::get().create_kernel("bounding_box_transform", build_opts.options()));
@@ -103,8 +108,7 @@ void CLBoundingBoxTransformKernel::configure(const ICLTensor *boxes, ICLTensor *
 
 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));
+    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(boxes, pred_boxes, deltas, info));
     return Status{};
 }