diff options
Diffstat (limited to 'src/core/CL')
-rw-r--r-- | src/core/CL/cl_kernels/bounding_box_transform.cl | 15 | ||||
-rw-r--r-- | src/core/CL/kernels/CLBoundingBoxTransformKernel.cpp | 12 |
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{}; } |