diff options
| author | Manuel Bottini <manuel.bottini@arm.com> | 2020-01-28 17:25:48 +0000 |
|---|---|---|
| committer | Manuel Bottini <manuel.bottini@arm.com> | 2020-02-03 16:55:07 +0000 |
| commit | 5c829cab49b0dcffd893065af05f0bf028e7a548 (patch) | |
| tree | e534e4377985499d6af7ab2c3ffd8b96af2c9ea0 /src/core/CL/cl_kernels/arg_min_max.cl | |
| parent | 94672fb2af6535adc6ea7fe8b8498580ad8cf3f4 (diff) | |
| download | ComputeLibrary-5c829cab49b0dcffd893065af05f0bf028e7a548.tar.gz | |
COMPMID-3045: CTS failures in ARGMAX/MIN
Change-Id: I35276a3d95dc99a7f4dea00e89c8ed206a5f13f1
Signed-off-by: Manuel Bottini <manuel.bottini@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2669
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michalis Spyrou <michalis.spyrou@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/arg_min_max.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/arg_min_max.cl | 9 |
1 files changed, 7 insertions, 2 deletions
diff --git a/src/core/CL/cl_kernels/arg_min_max.cl b/src/core/CL/cl_kernels/arg_min_max.cl index 06dcc8ddde..104d30d8f3 100644 --- a/src/core/CL/cl_kernels/arg_min_max.cl +++ b/src/core/CL/cl_kernels/arg_min_max.cl @@ -254,10 +254,15 @@ __kernel void arg_min_max_x( barrier(CLK_LOCAL_MEM_FENCE); + // Looking for the next highest power of 2 (maximum value of lsize is 8) + unsigned int middle = lsize - 1; + middle |= middle >> 1; + middle |= middle >> 2; + middle += 1; // Perform parallel reduction - for(unsigned int i = lsize >> 1; i > 0; i >>= 1) + for(unsigned int i = middle; i > 0; i >>= 1) { - if(lid < i) + if( lid < i && lid + i < lsize) { DATA_TYPE tmp0 = *(src_in_row + local_results[lid]); DATA_TYPE tmp1 = *(src_in_row + local_results[lid + i]); |