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/nhwc/space_to_batch.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/nhwc/space_to_batch.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/nhwc/space_to_batch.cl | 8 |
1 files changed, 4 insertions, 4 deletions
diff --git a/src/core/CL/cl_kernels/nhwc/space_to_batch.cl b/src/core/CL/cl_kernels/nhwc/space_to_batch.cl index 785206e3b9..695bd4c217 100644 --- a/src/core/CL/cl_kernels/nhwc/space_to_batch.cl +++ b/src/core/CL/cl_kernels/nhwc/space_to_batch.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021 Arm Limited. + * Copyright (c) 2018-2021, 2024 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -64,7 +64,7 @@ __kernel void space_to_batch_nhwc( const int batch_id, TENSOR3D_DECLARATION(output)) { - Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input, 0); + Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input); Image pad = CONVERT_TO_IMAGE_STRUCT_NO_STEP(paddings); Vector block = CONVERT_TO_VECTOR_STRUCT_NO_STEP(block_shape); Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output); @@ -130,7 +130,7 @@ __kernel void space_to_batch_static_nhwc( const int batch_id, TENSOR3D_DECLARATION(output)) { - Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input, 0); + Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input); Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output); int block_x = BLOCK_SHAPE_X; @@ -152,4 +152,4 @@ __kernel void space_to_batch_static_nhwc( *((__global DATA_TYPE *)out.ptr) = *((__global DATA_TYPE *)tensor4D_offset(&in, z, in_x, in_y, w)); } } -#endif // defined(BATCH_SIZE) && defined(DATA_TYPE) && defined(BLOCK_SHAPE_X) && defined(BLOCK_SHAPE_Y) && defined(PAD_LEFT_X) && defined(PAD_RIGHT_X) && defined(PAD_LEFT_Y) && defined(PAD_RIGHT_Y) && defined(WIDTH_IN) && defined(HEIGHT_IN)
\ No newline at end of file +#endif // defined(BATCH_SIZE) && defined(DATA_TYPE) && defined(BLOCK_SHAPE_X) && defined(BLOCK_SHAPE_Y) && defined(PAD_LEFT_X) && defined(PAD_RIGHT_X) && defined(PAD_LEFT_Y) && defined(PAD_RIGHT_Y) && defined(WIDTH_IN) && defined(HEIGHT_IN) |