59 #if defined(VEC_SIZE_IN) && defined(VEC_SIZE_OUT) && defined(LAST_ACCESSED_X_IN) && defined(LAST_ACCESSED_X_OUT) 62 const int xi_in = (int)(get_global_id(0) * VEC_SIZE_IN);
63 const int xi_out = (int)(get_global_id(0) * VEC_SIZE_OUT);
64 src.
ptr -= max(xi_in - (
int)LAST_ACCESSED_X_IN, 0) * src_stride_x;
65 dst.
ptr -= max(xi_out - (
int)LAST_ACCESSED_X_OUT, 0) * dst_stride_x;
68 data = vload8(0, (__global DATA_TYPE *)src.
ptr);
71 data_out = (
VEC_DATA_TYPE(DATA_TYPE, 16))(data.s0, data.s0, data.s1, data.s1, data.s2, data.s2, data.s3, data.s3, data.s4, data.s4, data.s5, data.s5, data.s6, data.s6, data.s7, data.s7);
73 vstore16(data_out, 0, (__global DATA_TYPE *)dst.
ptr);
74 vstore16(data_out, 0, (__global DATA_TYPE *)
tensor3D_offset(&dst, 0, 1, 0));
75 #else // !defined(VEC_SIZE_IN) && defined(VEC_SIZE_OUT) && defined(LAST_ACCESSED_X_IN) && defined(LAST_ACCESSED_X_OUT) 76 *((__global DATA_TYPE *)
tensor3D_offset(&dst, 0, 0, 0)) = *((__global DATA_TYPE *)src.
ptr);
77 *((__global DATA_TYPE *)
tensor3D_offset(&dst, 0, 1, 0)) = *((__global DATA_TYPE *)src.
ptr);
78 #endif // defined(VEC_SIZE_IN) && defined(VEC_SIZE_OUT) && defined(LAST_ACCESSED_X_IN) && defined(LAST_ACCESSED_X_OUT) Structure to hold 3D tensor information.
SimpleTensor< float > src
__kernel void upsample_layer_nchw(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_stride_z, uint src_step_z, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_stride_z, uint dst_step_z, uint dst_offset_first_element_in_bytes)
This function applies upsample on an input image.
#define CONVERT_TO_TENSOR3D_STRUCT(name)
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define TENSOR3D_DECLARATION(name)
__global const uchar * tensor3D_offset(const Tensor3D *tensor, int x, int y, int z)
Get the pointer position of a Tensor3D.
#define VEC_DATA_TYPE(type, size)