diff options
| author | Viet-Hoa Do <viet-hoa.do@arm.com> | 2024-01-16 16:23:24 +0000 |
|---|---|---|
| committer | Viet-Hoa Do <viet-hoa.do@arm.com> | 2024-01-18 11:32:09 +0000 |
| commit | 6829e0201e886cfd311e39f1d88f7452894bdfe5 (patch) | |
| tree | ac1b97df4b76d0d09652ef3dbdce2be3d222df3c /src/core/CL/cl_kernels/common/gather.cl | |
| parent | 8896cf7cb7df34c699e7453a0f0c683d1202ed15 (diff) | |
| download | ComputeLibrary-6829e0201e886cfd311e39f1d88f7452894bdfe5.tar.gz | |
Fix divide-by-zero compilation error
* CONVERT_TO_TENSOR4D_STRUCT_NO_STEP is implemented and used
in some CL kernels in the way that causes divide-by-zero issue.
- Since the steps are all zeros, the issue might have been
ignored by the compiler.
Resolves: COMPMID-6795
Signed-off-by: Viet-Hoa Do <viet-hoa.do@arm.com>
Change-Id: I0fb38fc62d63671b8abefa39b3d9b3ca6f49c7fe
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/10967
Reviewed-by: Gunes Bayir <gunes.bayir@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Benchmark: Arm Jenkins <bsgcomp@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/common/gather.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/common/gather.cl | 9 |
1 files changed, 4 insertions, 5 deletions
diff --git a/src/core/CL/cl_kernels/common/gather.cl b/src/core/CL/cl_kernels/common/gather.cl index 5d180f3781..e16f4bf315 100644 --- a/src/core/CL/cl_kernels/common/gather.cl +++ b/src/core/CL/cl_kernels/common/gather.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021, 2023 Arm Limited. + * Copyright (c) 2018-2021, 2023-2024 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -29,7 +29,6 @@ * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short * @note Axis should be given as a preprocessor argument using -DAXIS=axis. e.g. -DAXIS=1 * @attention Output tensor depth should be given as a preprocessor argument using -DOUTPUT_DIM_Z=size. e.g. -DOUTPUT_DIM_Z=16 - * @attention Input tensor depth should be given as a preprocessor argument using -DINPUT_DIM_Z=size. e.g. -DINPUT_DIM_Z=16 * * * @param[in] input_ptr Pointer to the source tensor. Supported data types: All @@ -67,8 +66,8 @@ __kernel void gather( const int pz = get_global_id(2) % OUTPUT_DIM_Z; const int pw = (get_global_id(2) / OUTPUT_DIM_Z ); - const Tensor4D input = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input, INPUT_DIM_Z); - const Tensor4D indices = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(indices, INDICES_DIM_Z); + const Tensor4D input = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input); + const Tensor4D indices = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(indices); Tensor4D output = CONVERT_TO_TENSOR4D_STRUCT(output, OUTPUT_DIM_Z); #if AXIS == 0 @@ -128,4 +127,4 @@ __kernel void gather( *(__global DATA_TYPE *)output.ptr = select((DATA_TYPE)0, *((__global const DATA_TYPE *)input_addr), (DATA_TYPE)(index < INDEX_LIMIT)); } -#endif //defined(DATA_TYPE) && defined(AXIS)
\ No newline at end of file +#endif //defined(DATA_TYPE) && defined(AXIS) |