aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/common
diff options
context:
space:
mode:
Diffstat (limited to 'src/core/CL/cl_kernels/common')
-rw-r--r--src/core/CL/cl_kernels/common/col2im.cl4
-rw-r--r--src/core/CL/cl_kernels/common/gather.cl9
-rw-r--r--src/core/CL/cl_kernels/common/instance_normalization.cl8
-rw-r--r--src/core/CL/cl_kernels/common/permute.cl4
-rw-r--r--src/core/CL/cl_kernels/common/reverse.cl4
-rw-r--r--src/core/CL/cl_kernels/common/slice_ops.cl6
-rw-r--r--src/core/CL/cl_kernels/common/tile.cl6
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);