Fix validation error in CL generate proposals kernel

This fix modifies some of the conversions done in the generate proposals kernel that causes DDK issues while compiling the kernel.

The issues are mostly related to conversion from i64 to fp16, and it doesn't affect fp32. Firstly, type identifier size_t is converted into unsigned int. But, this alone was compiling but causing mismatches, even in older devices, where it was passing before. Therefore, the fp16 conversion delayed until vector construction where the integers are now converted to fp32, and then fp16. This, although may not be ideal, seems like the best solution.

Resolves: COMPMID-6756
Signed-off-by: Gunes Bayir <gunes.bayir@arm.com>
Change-Id: Iee61216c908fe51431985b80c3653fc32add4741
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/10879
Benchmark: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Jakub Sujak <jakub.sujak@arm.com>
Reviewed-by: Pablo Marquez Tello <pablo.tello@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>

1 files changed, 6 insertions, 8 deletions

diff --git a/src/core/CL/cl_kernels/common/generate_proposals.cl b/src/core/CL/cl_kernels/common/generate_proposals.cl
index 5b8502072a..bfe1922ac2 100644
--- a/src/core/CL/cl_kernels/common/generate_proposals.cl
+++ b/src/core/CL/cl_kernels/common/generate_proposals.cl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2021 Arm Limited.
+ * Copyright (c) 2019-2021, 2023 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -59,18 +59,16 @@ __kernel void generate_proposals_compute_all_anchors(
     Vector anchors = CONVERT_TO_VECTOR_STRUCT_NO_STEP(anchors);
     Vector rois    = CONVERT_TO_VECTOR_STRUCT(rois);
 
-    const size_t idx = get_global_id(0);
+    const unsigned int idx = get_global_id(0);
     // Find the index of the anchor
-    const size_t anchor_idx = idx % NUM_ANCHORS;
+    const unsigned int anchor_idx = idx % NUM_ANCHORS;
 
     // Find which shift is this thread using
-    const size_t shift_idx = idx / NUM_ANCHORS;
+    const unsigned int shift_idx = idx / NUM_ANCHORS;
 
     // Compute the shift on the X and Y direction (the shift depends exclusively by the index thread id)
-    const DATA_TYPE
-    shift_x = (DATA_TYPE)(shift_idx % WIDTH) * STRIDE;
-    const DATA_TYPE
-    shift_y = (DATA_TYPE)(shift_idx / WIDTH) * STRIDE;
+    const float shift_x = (float)(shift_idx % WIDTH) * STRIDE;
+    const float shift_y = (float)(shift_idx / WIDTH) * STRIDE;
 
     const VEC_DATA_TYPE(DATA_TYPE, NUM_ROI_FIELDS)
     shift = (VEC_DATA_TYPE(DATA_TYPE, NUM_ROI_FIELDS))(shift_x, shift_y, shift_x, shift_y);