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
diff --git a/src/core/CL/cl_kernels/bounding_box_transform.cl b/src/core/CL/cl_kernels/bounding_box_transform.cl
index 14a0fad..0972355 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 @@
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 @@
// 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 09f3d33..bff28e3 100644
--- a/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp
+++ b/src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp
@@ -39,7 +39,7 @@
{
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 @@
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 @@
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 @@
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 @@
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{};
}