diff options
Diffstat (limited to 'src/core/CL/cl_kernels/common')
-rw-r--r-- | src/core/CL/cl_kernels/common/col2im.cl | 4 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/common/gather.cl | 9 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/common/instance_normalization.cl | 8 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/common/permute.cl | 4 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/common/reverse.cl | 4 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/common/slice_ops.cl | 6 | ||||
-rw-r--r-- | src/core/CL/cl_kernels/common/tile.cl | 6 |
7 files changed, 20 insertions, 21 deletions
diff --git a/src/core/CL/cl_kernels/common/col2im.cl b/src/core/CL/cl_kernels/common/col2im.cl index 89054dcb31..4dc005fd43 100644 --- a/src/core/CL/cl_kernels/common/col2im.cl +++ b/src/core/CL/cl_kernels/common/col2im.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2021 Arm Limited. + * Copyright (c) 2017-2021, 2024 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -67,7 +67,7 @@ __kernel void col2im( TENSOR4D_DECLARATION(dst)) { Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); - Tensor4D dst = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(dst, 0); + Tensor4D dst = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(dst); const uint xd = get_global_id(1) % WIDTH_OUTPUT; // x coordinate of the destination tensor const uint yd = get_global_id(1) / WIDTH_OUTPUT; // y coordinate of the destination tensor 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) diff --git a/src/core/CL/cl_kernels/common/instance_normalization.cl b/src/core/CL/cl_kernels/common/instance_normalization.cl index adfbebd67d..f9b3cd3620 100644 --- a/src/core/CL/cl_kernels/common/instance_normalization.cl +++ b/src/core/CL/cl_kernels/common/instance_normalization.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021 Arm Limited. + * Copyright (c) 2019-2021, 2024 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -53,7 +53,7 @@ __kernel void compute_mean_var( TENSOR4D_DECLARATION(input), 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_NO_STEP(output); #if defined(NHWC) @@ -176,10 +176,10 @@ __kernel void instance_normalization( #endif /* IN_PLACE */ ) { - Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input, 0); + Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input); Tensor3D mean_var = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(mean_var); #ifndef IN_PLACE - Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(output, 0); + Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(output); #endif /* IN_PLACE */ #if defined(NHWC) diff --git a/src/core/CL/cl_kernels/common/permute.cl b/src/core/CL/cl_kernels/common/permute.cl index a03eeb1a19..1a97ca7495 100644 --- a/src/core/CL/cl_kernels/common/permute.cl +++ b/src/core/CL/cl_kernels/common/permute.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021 Arm Limited. + * Copyright (c) 2018-2021, 2024 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -54,7 +54,7 @@ __kernel void permute(TENSOR4D_DECLARATION(input), { Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT(input, DEPTH_IN); - Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(output, 0); + Tensor4D out = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(output); int out_index[4] = { 0 }; int in_index[4] = { 0 }; diff --git a/src/core/CL/cl_kernels/common/reverse.cl b/src/core/CL/cl_kernels/common/reverse.cl index f94bfb6640..e6df3041c2 100644 --- a/src/core/CL/cl_kernels/common/reverse.cl +++ b/src/core/CL/cl_kernels/common/reverse.cl @@ -1,5 +1,5 @@ /* -* Copyright (c) 2018-2021, 2023 Arm Limited. + * Copyright (c) 2018-2021, 2023-2024 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -71,7 +71,7 @@ __kernel void reverse(TENSOR4D_DECLARATION(src), { Tensor4D src = CONVERT_TO_TENSOR4D_STRUCT(src, depth); Vector axis = CONVERT_TO_VECTOR_STRUCT_NO_STEP(axis); - Tensor4D dst = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(dst, depth); + Tensor4D dst = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(dst); const uint x_in = get_global_id(0); const uint y_in = get_global_id(1); diff --git a/src/core/CL/cl_kernels/common/slice_ops.cl b/src/core/CL/cl_kernels/common/slice_ops.cl index d12c60f5ea..189d414aba 100644 --- a/src/core/CL/cl_kernels/common/slice_ops.cl +++ b/src/core/CL/cl_kernels/common/slice_ops.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021 Arm Limited. + * Copyright (c) 2018-2021, 2024 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -28,7 +28,7 @@ * @attention Supported tensor rank: up to 4 * * @attention Data type can be passed using the -DDATA_TYPE compile flag, e.g. -DDATA_TYPE=float - * @attention Input and output tensor dephts should be given as a preprocessor arguments using -DSRC_DEPTH=size. and -DDST_DEPTH=size + * @attention Output tensor depht should be given as a preprocessor arguments using -DDST_DEPTH=size * @attention Absolute start coordinates for each dimension should be given as preprocessor -DSTART_index=value e.g. -DSTART_0=2 * @attention Strides for each dimension should be given as preprocessor -DSTRIDE_index=value e.g. -DSTRIDE_1=1 * @@ -58,7 +58,7 @@ __kernel void strided_slice( TENSOR4D_DECLARATION(output)) { // Get pixels pointer - Tensor4D input = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input, SRC_DEPTH); + Tensor4D input = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input); Tensor4D output = CONVERT_TO_TENSOR4D_STRUCT(output, DST_DEPTH); int offset = 0; diff --git a/src/core/CL/cl_kernels/common/tile.cl b/src/core/CL/cl_kernels/common/tile.cl index 971750b7b2..4d8f802ea1 100644 --- a/src/core/CL/cl_kernels/common/tile.cl +++ b/src/core/CL/cl_kernels/common/tile.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2021, 2023 Arm Limited. + * Copyright (c) 2018-2021, 2023-2024 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -50,8 +50,8 @@ __kernel void tile( TENSOR4D_DECLARATION(input), TENSOR4D_DECLARATION(output)) { - Tensor4D output = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(output, DST_DEPTH); - Tensor4D input = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input, SRC_DEPTH); + Tensor4D output = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(output); + Tensor4D input = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input); // For all coordinates but x, each tile copies from the input const int y = get_global_id(1); |