diff options
| author | giuros01 <giuseppe.rossini@arm.com> | 2018-11-16 10:39:59 +0000 |
|---|---|---|
| committer | Giuseppe Rossini <giuseppe.rossini@arm.com> | 2018-11-16 17:57:46 +0000 |
| commit | d696cb6d18c2fe66f1abce88bbd14faf2137ef89 (patch) | |
| tree | 9428fea2613651e2d7ffb53b15c62c80594eb8e3 /src/core/CL/cl_kernels | |
| parent | c8df89f477c3dc63f396ad37bee8ed5d50dee4ac (diff) | |
| download | ComputeLibrary-d696cb6d18c2fe66f1abce88bbd14faf2137ef89.tar.gz | |
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
Diffstat (limited to 'src/core/CL/cl_kernels')
| -rw-r--r-- | src/core/CL/cl_kernels/bounding_box_transform.cl | 15 |
1 files changed, 10 insertions, 5 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) |