68 input.
ptr += (int)START_0 * input_stride_x;
69 #elif defined(START_0) && defined(STRIDE_0) && defined(VEC_SIZE) && defined(LAST_ACCESSED_X) 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) 84 #elif defined(START_1) && defined(STRIDE_1) 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) 91 #endif // defined(START_1) && defined(STRIDE_1) 96 #elif defined(START_2) && defined(STRIDE_2) 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) 107 #endif // defined(START_2) && defined(STRIDE_2) 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) 126 #if defined(VEC_SIZE) && defined(LAST_ACCESSED_X) 132 #else // defined(VEC_SIZE) && defined(LAST_ACCESSED_X) 134 #endif // defined(VEC_SIZE) && defined(LAST_ACCESSED_X) __global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
__global uchar * ptr
Pointer to the starting postion of the buffer.
const size_t input_stride_y
Structure to hold 4D tensor information.
#define CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(name, mod_size)
__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.
#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size)
#define TENSOR4D_DECLARATION(name)
const size_t input_stride_z
#define VEC_DATA_TYPE(type, size)