Compute Library
 22.08
slice_ops.cl
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2018-2021 Arm Limited.
3  *
4  * SPDX-License-Identifier: MIT
5  *
6  * Permission is hereby granted, free of charge, to any person obtaining a copy
7  * of this software and associated documentation files (the "Software"), to
8  * deal in the Software without restriction, including without limitation the
9  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10  * sell copies of the Software, and to permit persons to whom the Software is
11  * furnished to do so, subject to the following conditions:
12  *
13  * The above copyright notice and this permission notice shall be included in all
14  * copies or substantial portions of the Software.
15  *
16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22  * SOFTWARE.
23  */
24 #include "helpers.h"
25 
26 /** Perform a strided slice operation on a given input.
27  *
28  * @attention Supported tensor rank: up to 4
29  *
30  * @attention Data type can be passed using the -DDATA_TYPE compile flag, e.g. -DDATA_TYPE=float
31  * @attention Input and output tensor dephts should be given as a preprocessor arguments using -DSRC_DEPTH=size. and -DDST_DEPTH=size
32  * @attention Absolute start coordinates for each dimension should be given as preprocessor -DSTART_index=value e.g. -DSTART_0=2
33  * @attention Strides for each dimension should be given as preprocessor -DSTRIDE_index=value e.g. -DSTRIDE_1=1
34  *
35  * @param[in] input_ptr Pointer to the source tensor. Supported data types: All
36  * @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes)
37  * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
38  * @param[in] input_stride_y Stride of the source tensor in Y dimension (in bytes)
39  * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
40  * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
41  * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
42  * @param[in] input_stride_w Stride of the source tensor in W dimension (in bytes)
43  * @param[in] input_step_w input_stride_w * number of elements along W processed per workitem(in bytes)
44  * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source tensor
45  * @param[out] output_ptr Pointer to the destination tensor. Supported data types: same as @p input_ptr
46  * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes)
47  * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
48  * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes)
49  * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
50  * @param[in] output_stride_z Stride of the destination tensor in Z dimension (in bytes)
51  * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
52  * @param[in] output_stride_w Stride of the destination tensor in W dimension (in bytes)
53  * @param[in] output_step_w output_stride_w * number of elements along W processed per workitem(in bytes)
54  * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor
55  */
56 __kernel void strided_slice(
58  TENSOR4D_DECLARATION(output))
59 {
60  // Get pixels pointer
62  Tensor4D output = CONVERT_TO_TENSOR4D_STRUCT(output, DST_DEPTH);
63 
64  int offset = 0;
65 
66  // Offset X
67 #if defined(SHRINK_0)
68  input.ptr += (int)START_0 * input_stride_x;
69 #elif defined(START_0) && defined(STRIDE_0) && defined(VEC_SIZE) && defined(LAST_ACCESSED_X)
70  // Check if access on width gets out of bounds
71  // If it does shift access vector to access elements within bounds
72  const int xi = (int)(get_global_id(0) * VEC_SIZE);
73  offset = (int)START_0 + min(xi, (int)LAST_ACCESSED_X);
74  input.ptr += offset * input_stride_x;
75  output.ptr -= max(xi - (int)LAST_ACCESSED_X, 0) * output_stride_x;
76 #elif defined(START_0) && defined(STRIDE_0)
77  offset = (int)START_0 + (int)get_global_id(0) * (int)STRIDE_0;
78  input.ptr += offset * input_stride_x;
79 #endif // defined(START_0) && defined(STRIDE_0)
80 
81  // Offset Y
82 #if defined(SHRINK_1)
83  input.ptr += (int)START_1 * input_stride_y;
84 #elif defined(START_1) && defined(STRIDE_1)
85 #if defined(SHRINK_0)
86  offset = (int)START_1 + (int)get_global_id(0) * (int)STRIDE_1;
87 #else // defined(SHRINK_0)
88  offset = (int)START_1 + (int)get_global_id(1) * (int)STRIDE_1;
89 #endif // defined(SHRINK_0)
90  input.ptr += offset * input_stride_y;
91 #endif // defined(START_1) && defined(STRIDE_1)
92 
93  // Offset Z
94 #if defined(SHRINK_2)
95  input.ptr += (int)START_2 * input_stride_z;
96 #elif defined(START_2) && defined(STRIDE_2)
97 
98 #if defined(SHRINK_1) && defined(SHRINK_0)
99  offset = (int)START_2 + (int)get_global_id(0) * (int)STRIDE_2;
100 #elif defined(SHRINK_1) || defined(SHRINK_0)
101  offset = (int)START_2 + (int)get_global_id(1) * (int)STRIDE_2;
102 #else // defined(SHRINK_1) && defined(SHRINK_0)
103  offset = (int)START_2 + ((int)get_global_id(2) % (int)DST_DEPTH) * (int)STRIDE_2;
104 #endif // defined(SHRINK_1) && defined(SHRINK_0)
105 
106  input.ptr += offset * input_stride_z;
107 #endif // defined(START_2) && defined(STRIDE_2)
108 
109  // Offset depth
110 #if defined(SHRINK_3)
111  input.ptr += (int)START_3 * input_stride_w;
112 #elif defined(START_3) && defined(STRIDE_3)
113 #if defined(SHRINK_2) && defined(SHRINK_1) && defined(SHRINK_0)
114  offset = (int)START_3 + (int)get_global_id(0) * (int)STRIDE_3;
115 #elif !defined(SHRINK_2) && !defined(SHRINK_1) && !defined(SHRINK_0)
116  offset = (int)START_3 + ((int)get_global_id(2) / (int)DST_DEPTH) * (int)STRIDE_3;
117 #elif(defined(SHRINK_0) && defined(SHRINK_1)) || (defined(SHRINK_1) && defined(SHRINK_2)) || (defined(SHRINK_0) && defined(SHRINK_2))
118  offset = (int)START_3 + (int)get_global_id(1) * (int)STRIDE_3;
119 #else // defined(SHRINK_2) && defined(SHRINK_1) && defined(SHRINK_0)
120  offset = (int)START_3 + ((int)get_global_id(2) % (int)DST_DEPTH) * (int)STRIDE_3;
121 #endif // defined(SHRINK_2) && defined(SHRINK_1) && defined(SHRINK_0)
122  input.ptr += offset * input_stride_w;
123 #endif // defined(START_3) && defined(STRIDE_3)
124 
125  // Store result
126 #if defined(VEC_SIZE) && defined(LAST_ACCESSED_X)
127  VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
128  val = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(input.ptr));
129 
131  (val, 0, (__global DATA_TYPE *)(output.ptr));
132 #else // defined(VEC_SIZE) && defined(LAST_ACCESSED_X)
133  *((__global DATA_TYPE *)(output.ptr)) = *((__global DATA_TYPE *)(input.ptr));
134 #endif // defined(VEC_SIZE) && defined(LAST_ACCESSED_X)
135 }
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
Definition: helpers.h:1084
#define VEC_SIZE
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:918
Structure to hold 4D tensor information.
Definition: helpers.h:916
#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size)
Definition: helpers.h:881
__kernel void strided_slice(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_stride_z, uint input_step_z, uint input_stride_w, uint input_step_w, uint input_offset_first_element_in_bytes, __global uchar *output_ptr, uint output_stride_x, uint output_step_x, uint output_stride_y, uint output_step_y, uint output_stride_z, uint output_step_z, uint output_stride_w, uint output_step_w, uint output_offset_first_element_in_bytes)
Perform a strided slice operation on a given input.
Definition: slice_ops.cl:56
#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size)
Definition: helpers.h:877
const size_t input_stride_y
Definition: impl.cpp:50
#define VSTORE(size)
Definition: helpers.h:458
#define TENSOR4D_DECLARATION(name)
Definition: helpers.h:823
#define VLOAD(size)
Definition: helpers.h:204
const size_t input_stride_z
Definition: impl.cpp:51
#define VEC_DATA_TYPE(type, size)
Definition: helpers.h:728