diff options
| author | Giorgio Arena <giorgio.arena@arm.com> | 2021-01-07 14:01:47 +0000 |
|---|---|---|
| committer | Giorgio Arena <giorgio.arena@arm.com> | 2021-01-08 13:51:05 +0000 |
| commit | c216545a9ede448213e76bd50cba4ebd481b942c (patch) | |
| tree | fea258b7f8f44e8437fc8d02d8a84d19e1994090 | |
| parent | 33b103b8ad1b063ddf909a2c2027e25ef46c0d88 (diff) | |
| download | ComputeLibrary-c216545a9ede448213e76bd50cba4ebd481b942c.tar.gz | |
[Nightly Failure] Fix OpenCL kernel compilation for CLRange
- Change raw pointers in OpenCL kernel to __global uchar*
Resolves: COMPMID-4079
Change-Id: Ieeb99ced565bef59583216fd274958b29c7b2758
Signed-off-by: Giorgio Arena <giorgio.arena@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4774
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Pablo Marquez Tello <pablo.tello@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
| -rw-r--r-- | src/core/CL/cl_kernels/range.cl | 12 |
1 files changed, 6 insertions, 6 deletions
diff --git a/src/core/CL/cl_kernels/range.cl b/src/core/CL/cl_kernels/range.cl index 5569101615..467b962b0d 100644 --- a/src/core/CL/cl_kernels/range.cl +++ b/src/core/CL/cl_kernels/range.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2020 Arm Limited. + * Copyright (c) 2018-2021 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -55,8 +55,8 @@ __kernel void range( VECTOR_DECLARATION(out)) { - uint id = max((int)(get_global_id(0) * VECTOR_SIZE - (VECTOR_SIZE - VEC_SIZE_LEFTOVER) % VECTOR_SIZE), 0); - __global DATA_TYPE *dst_ptr = out_ptr + out_offset_first_element_in_bytes + id * sizeof(DATA_TYPE); + uint id = max((int)(get_global_id(0) * VECTOR_SIZE - (VECTOR_SIZE - VEC_SIZE_LEFTOVER) % VECTOR_SIZE), 0); + __global uchar *dst_ptr = out_ptr + out_offset_first_element_in_bytes + id * sizeof(DATA_TYPE); #if VECTOR_SIZE == 1 DATA_TYPE seq; seq = (DATA_TYPE)START + (DATA_TYPE)id * (DATA_TYPE)STEP; @@ -104,10 +104,10 @@ __kernel void range( __kernel void range_quantized( VECTOR_DECLARATION(out)) { - uint id = max((int)(get_global_id(0) * VECTOR_SIZE - (VECTOR_SIZE - VEC_SIZE_LEFTOVER) % VECTOR_SIZE), 0); - __global DATA_TYPE *dst_ptr = out_ptr + out_offset_first_element_in_bytes + id * sizeof(DATA_TYPE); + uint id = max((int)(get_global_id(0) * VECTOR_SIZE - (VECTOR_SIZE - VEC_SIZE_LEFTOVER) % VECTOR_SIZE), 0); + __global uchar *dst_ptr = out_ptr + out_offset_first_element_in_bytes + id * sizeof(DATA_TYPE); #if VECTOR_SIZE == 1 - float seq; + float seq; seq = (float)START + (float)id * (float)STEP; seq = (DATA_TYPE)(int)(seq / ((float)SCALE_OUT) + (float)OFFSET_OUT); seq = max(0.0f, min(seq, 255.0f)); |