diff options
| -rw-r--r-- | arm_compute/core/Types.h | 29 | ||||
| -rw-r--r-- | src/core/CL/cl_kernels/bounding_box_transform.cl | 15 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp | 12 | ||||
| -rw-r--r-- | tests/validation/CL/BoundingBoxTransform.cpp | 15 | ||||
| -rw-r--r-- | tests/validation/reference/BoundingBoxTransform.cpp | 21 |
5 files changed, 59 insertions, 33 deletions
diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h index 4eb8129b62..38094ee56a 100644 --- a/arm_compute/core/Types.h +++ b/arm_compute/core/Types.h @@ -1209,21 +1209,24 @@ private: }; /** Bounding Box Transform information class */ -class BoundingBoxTransformInfo +class BoundingBoxTransformInfo final { public: /** Constructor * - * @param[in] img_width Width of the original image - * @param[in] img_height Height, of the original image - * @param[in] scale Scale of the original image - * @param[in] apply_scale (Optional)Re-apply scaling after transforming the boxes. Defaults to false - * @param[in] weights (Optional)Weights [wx, wy, ww, wh] for the deltas. Defaults to all ones - * @param[in] bbox_xform_clip (Optional)Minimum bounding box width and height after bounding box transformation in log-space. Defaults to log(1000/16) + * @param[in] img_width Width of the original image + * @param[in] img_height Height, of the original image + * @param[in] scale Scale of the original image + * @param[in] apply_scale (Optional)Re-apply scaling after transforming the boxes. Defaults to false + * @param[in] weights (Optional)Weights [wx, wy, ww, wh] for the deltas. Defaults to all ones + * @param[in] correct_transform_coords (Optional)Correct bounding box transform coordinates. Defaults to false + * @param[in] bbox_xform_clip (Optional)Minimum bounding box width and height after bounding box transformation in log-space. Defaults to log(1000/16) */ - BoundingBoxTransformInfo(float img_width, float img_height, float scale, bool apply_scale = false, const std::array<float, 4> weights = { { 1.f, 1.f, 1.f, 1.f } }, float bbox_xform_clip = - 4.135166556742356f) - : _img_width(img_width), _img_height(img_height), _scale(scale), _apply_scale(apply_scale), _weights(weights), _bbox_xform_clip(bbox_xform_clip) + BoundingBoxTransformInfo(float img_width, float img_height, float scale, bool apply_scale = false, const std::array<float, 4> weights = { { 1.f, 1.f, 1.f, 1.f } }, bool correct_transform_coords = + false, + float bbox_xform_clip = + 4.135166556742356f) + : _img_width(img_width), _img_height(img_height), _scale(scale), _apply_scale(apply_scale), _correct_transform_coords(correct_transform_coords), _weights(weights), _bbox_xform_clip(bbox_xform_clip) { } @@ -1257,11 +1260,17 @@ public: return _apply_scale; } + bool correct_transform_coords() const + { + return _correct_transform_coords; + } + private: float _img_width; float _img_height; float _scale; bool _apply_scale; + bool _correct_transform_coords; std::array<float, 4> _weights; float _bbox_xform_clip; }; 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{}; } diff --git a/tests/validation/CL/BoundingBoxTransform.cpp b/tests/validation/CL/BoundingBoxTransform.cpp index dba3a0eabe..c5856cae3d 100644 --- a/tests/validation/CL/BoundingBoxTransform.cpp +++ b/tests/validation/CL/BoundingBoxTransform.cpp @@ -50,7 +50,8 @@ const auto BboxInfoDataset = framework::dataset::make("BboxInfo", { BoundingBoxT 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 }), + BoundingBoxTransformInfo(800U, 600U, 4U, false, { 1.0, 0.5, 1.5, 2.0 }, true) }); const auto DeltaDataset = framework::dataset::make("DeltasShape", { TensorShape(36U, 1U), @@ -74,22 +75,26 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( 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), // 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, 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)})), - framework::dataset::make("Expected", { true, false, false, false, false })), + 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(CLBoundingBoxTransform::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); diff --git a/tests/validation/reference/BoundingBoxTransform.cpp b/tests/validation/reference/BoundingBoxTransform.cpp index 9918ff68c5..55dd165b51 100644 --- a/tests/validation/reference/BoundingBoxTransform.cpp +++ b/tests/validation/reference/BoundingBoxTransform.cpp @@ -50,7 +50,10 @@ SimpleTensor<T> bounding_box_transform(const SimpleTensor<T> &boxes, const Simpl const int img_h = floor(info.img_height() / info.scale() + 0.5f); const int img_w = floor(info.img_width() / info.scale() + 0.5f); - const T scale = (info.apply_scale() ? T(info.scale()) : T(1)); + const auto scale_after = (info.apply_scale() ? T(info.scale()) : T(1)); + const auto scale_before = T(info.scale()); + ARM_COMPUTE_ERROR_ON(scale_before <= 0); + const auto offset = (info.correct_transform_coords() ? T(1.f) : T(0.f)); const size_t box_fields = 4; const size_t class_fields = 4; @@ -59,10 +62,10 @@ SimpleTensor<T> bounding_box_transform(const SimpleTensor<T> &boxes, const Simpl { // Extract ROI information const size_t start_box = box_fields * i; - const T width = boxes[start_box + 2] - boxes[start_box] + T(1.0); - const T height = boxes[start_box + 3] - boxes[start_box + 1] + T(1.0); - const T ctr_x = boxes[start_box] + T(0.5) * width; - const T ctr_y = boxes[start_box + 1] + T(0.5) * height; + const T width = (boxes[start_box + 2] / scale_before) - (boxes[start_box] / scale_before) + T(1.f); + const T height = (boxes[start_box + 3] / scale_before) - (boxes[start_box + 1] / scale_before) + T(1.f); + const T ctr_x = (boxes[start_box] / scale_before) + T(0.5f) * width; + const T ctr_y = (boxes[start_box + 1] / scale_before) + T(0.5f) * height; for(size_t j = 0; j < num_classes; ++j) { @@ -84,10 +87,10 @@ SimpleTensor<T> bounding_box_transform(const SimpleTensor<T> &boxes, const Simpl const T pred_h = T(std::exp(dh)) * height; // Store the prediction into the output tensor - pred_boxes_ptr[start_delta] = scale * utility::clamp<T>(pred_ctr_x - T(0.5) * pred_w, T(0), T(img_w - 1)); - pred_boxes_ptr[start_delta + 1] = scale * utility::clamp<T>(pred_ctr_y - T(0.5) * pred_h, T(0), T(img_h - 1)); - pred_boxes_ptr[start_delta + 2] = scale * utility::clamp<T>(pred_ctr_x + T(0.5) * pred_w, T(0), T(img_w - 1)); - pred_boxes_ptr[start_delta + 3] = scale * utility::clamp<T>(pred_ctr_y + T(0.5) * pred_h, T(0), T(img_h - 1)); + pred_boxes_ptr[start_delta] = scale_after * utility::clamp<T>(pred_ctr_x - T(0.5f) * pred_w, T(0), T(img_w - 1)); + pred_boxes_ptr[start_delta + 1] = scale_after * utility::clamp<T>(pred_ctr_y - T(0.5f) * pred_h, T(0), T(img_h - 1)); + pred_boxes_ptr[start_delta + 2] = scale_after * utility::clamp<T>(pred_ctr_x + T(0.5f) * pred_w - offset, T(0), T(img_w - 1)); + pred_boxes_ptr[start_delta + 3] = scale_after * utility::clamp<T>(pred_ctr_y + T(0.5f) * pred_h - offset, T(0), T(img_h - 1)); } } return pred_boxes; |