26 #if defined(FIXED_POINT_POSITION) 28 #endif // FIXED_POINT_POSITION 30 #if defined(DATA_TYPE) && defined(WIDTH_OUTPUT) && defined(ELEMENT_SIZE) && defined(WIDTH_INPUT) 31 #if !defined(FIXED_POINT_POSITION) 34 #define COND_DATA_TYPE char 35 #elif ELEMENT_SIZE == 2 36 #define COND_DATA_TYPE short 37 #elif ELEMENT_SIZE == 4 38 #define COND_DATA_TYPE int 40 #error "Element size not support" 41 #endif // ELEMENT_SIZE 76 data = vload8(0, (__global
DATA_TYPE *)src.ptr);
78 uint x = get_global_id(0) * 8;
79 uint8 x_clamped = x + (uint8)(0, 1, 2, 3, 4, 5, 6, 7);
85 x_clamped = select((uint8)x, x_clamped, convert_int8(cond0));
90 __global uchar *output_ptr = dst_ptr + dst_offset_first_element_in_bytes;
93 int idx = (get_global_id(1) / WIDTH_OUTPUT) * dst_stride_y + (get_global_id(1) % WIDTH_OUTPUT) * dst_stride_x + get_global_id(2) * dst_stride_w;
96 *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s0 * dst_stride_z)) = data.s0;
97 *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s1 * dst_stride_z)) = data.s1;
98 *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s2 * dst_stride_z)) = data.s2;
99 *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s3 * dst_stride_z)) = data.s3;
100 *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s4 * dst_stride_z)) = data.s4;
101 *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s5 * dst_stride_z)) = data.s5;
102 *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s6 * dst_stride_z)) = data.s6;
103 *((__global DATA_TYPE *)(output_ptr + idx + x_clamped.s7 * dst_stride_z)) = data.s7;
105 #else // !defined(FIXED_POINT_POSITION) 129 __kernel
void col2im(
138 int idx = get_global_id(0) * dst.
stride_z + (get_global_id(1) / WIDTH_OUTPUT) * dst_stride_y + (get_global_id(1) % WIDTH_OUTPUT) * dst_stride_x + get_global_id(2) * dst_stride_w;
141 *((__global DATA_TYPE *)(dst.
ptr + idx)) = *((__global DATA_TYPE *)(src.
ptr));
143 #endif // !defined(FIXED_POINT_POSITION) 144 #endif // defined(DATA_TYPE) && defined(WIDTH_OUTPUT) && defined(ELEMENT_SIZE) && defined(WIDTH_INPUT)
#define CONVERT_TO_TENSOR3D_STRUCT(name)
int stride_z
Stride of the image in Z dimension (in bytes)
Structure to hold 3D tensor information.
#define TENSOR3D_DECLARATION(name)
#define VEC_DATA_TYPE(type, size)
#define CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(name)
__global uchar * ptr
Pointer to the starting postion of the buffer.
convolution configure & src