27 #define VEC_INT VEC_DATA_TYPE(int, VEC_SIZE) 29 #if defined(OFFSET_IN1) && defined(OFFSET_OUT) && defined(SCALE_IN1) && defined(SCALE_OUT) 30 #define VEC_FLOAT VEC_DATA_TYPE(float, VEC_SIZE) 31 #define VEC_QUANT VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) 32 #define CONVERT_RTE(x, type) (convert_##type##_rte((x))) 33 #define CONVERT_DOWN(x, type) CONVERT_RTE(x, type) 34 inline VEC_QUANT requantize(VEC_QUANT
input,
float in_offset,
float out_offset,
float in_scale,
float out_scale)
43 #if defined(DATA_TYPE) 44 #define VEC_TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) 46 #if defined(DEPTH) && defined(ELEMENT_SIZE) 47 #if defined(INPUT1_WIDTH) 49 #define SELECT_TYPE SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE) 50 #define SEQ VEC_OFFS(int, VEC_SIZE) 91 __kernel
void concatenate_width_x2(
98 const int y = get_global_id(1);
99 const int z = get_global_id(2) % (int)DEPTH;
100 const int w = get_global_id(2) / (int)DEPTH;
101 const int x1 = min(x, (
int)INPUT1_WIDTH - (
int)
VEC_SIZE);
102 const int x2 = max(x - (
int)INPUT1_WIDTH, 0);
105 const __global uchar *dst_addr = dst_ptr + (int)dst_offset_first_element_in_bytes + x *
sizeof(
DATA_TYPE) + y * (int)dst_stride_y + z * (
int)dst_stride_z + w * (int)dst_stride_w;
106 const __global uchar *src1_addr = src1_ptr + (int)src1_offset_first_element_in_bytes + x1 *
sizeof(
DATA_TYPE) + y * (int)src1_stride_y + z * (
int)src1_stride_z + w * (int)src1_stride_w;
107 const __global uchar *src2_addr = src2_ptr + (int)src2_offset_first_element_in_bytes + x2 *
sizeof(
DATA_TYPE) + y * (int)src2_stride_y + z * (
int)src2_stride_z + w * (int)src2_stride_w;
112 #if defined(OFFSET_IN1) && defined(OFFSET_IN2) && defined(OFFSET_OUT) && defined(SCALE_IN1) && defined(SCALE_IN2) && defined(SCALE_OUT) 113 src1_values = requantize(src1_values, OFFSET_IN1, OFFSET_OUT, SCALE_IN1, SCALE_OUT);
114 src2_values = requantize(src2_values, OFFSET_IN2, OFFSET_OUT, SCALE_IN2, SCALE_OUT);
124 const VEC_TYPE values0 =
select(src2_values, src1_values, cond);
129 #if defined(INPUT2_WIDTH) && defined(INPUT3_WIDTH) 191 __kernel
void concatenate_width_x4(
200 const int y = get_global_id(1);
201 const int z = get_global_id(2) % (int)DEPTH;
202 const int w = get_global_id(2) / (int)DEPTH;
204 const int x1 = min(x, (
int)INPUT1_WIDTH - (
int)
VEC_SIZE);
205 const int x2 = min(max(x - (
int)INPUT1_WIDTH, 0), (
int)INPUT2_WIDTH - (
int)
VEC_SIZE);
206 const int x3 = min(max(x - (
int)INPUT1_WIDTH - (
int)INPUT2_WIDTH, 0), (
int)INPUT3_WIDTH - (
int)
VEC_SIZE);
207 const int x4 = max(x - (
int)INPUT1_WIDTH - (
int)INPUT2_WIDTH - (
int)INPUT3_WIDTH, 0);
210 const __global uchar *dst_addr = dst_ptr + (int)dst_offset_first_element_in_bytes + x *
sizeof(
DATA_TYPE) + y * (int)dst_stride_y + z * (
int)dst_stride_z + w * (int)dst_stride_w;
211 const __global uchar *src1_addr = src1_ptr + (int)src1_offset_first_element_in_bytes + x1 *
sizeof(
DATA_TYPE) + y * (int)src1_stride_y + z * (
int)src1_stride_z + w * (int)src1_stride_w;
212 const __global uchar *src2_addr = src2_ptr + (int)src2_offset_first_element_in_bytes + x2 *
sizeof(
DATA_TYPE) + y * (int)src2_stride_y + z * (
int)src2_stride_z + w * (int)src2_stride_w;
213 const __global uchar *src3_addr = src3_ptr + (int)src3_offset_first_element_in_bytes + x3 *
sizeof(
DATA_TYPE) + y * (int)src3_stride_y + z * (
int)src3_stride_z + w * (int)src3_stride_w;
214 const __global uchar *src4_addr = src4_ptr + (int)src4_offset_first_element_in_bytes + x4 *
sizeof(
DATA_TYPE) + y * (int)src4_stride_y + z * (
int)src4_stride_z + w * (int)src4_stride_w;
221 #if defined(OFFSET_IN1) && defined(OFFSET_OUT) && defined(SCALE_IN1) && defined(SCALE_OUT) && defined(OFFSET_IN2) && defined(SCALE_IN2) && defined(OFFSET_IN3) && defined(SCALE_IN3) && defined(OFFSET_IN4) && defined(SCALE_IN4) 222 src1_values = requantize(src1_values, OFFSET_IN1, OFFSET_OUT, SCALE_IN1, SCALE_OUT);
223 src2_values = requantize(src2_values, OFFSET_IN2, OFFSET_OUT, SCALE_IN2, SCALE_OUT);
224 src3_values = requantize(src3_values, OFFSET_IN3, OFFSET_OUT, SCALE_IN3, SCALE_OUT);
225 src4_values = requantize(src4_values, OFFSET_IN4, OFFSET_OUT, SCALE_IN4, SCALE_OUT);
244 cond_in2 =
CONVERT(x_coords < (
VEC_INT)(INPUT1_WIDTH), SELECT_TYPE);
245 cond_in3 =
CONVERT(x_coords < (
VEC_INT)(INPUT1_WIDTH + INPUT2_WIDTH), SELECT_TYPE);
246 cond_in4 =
CONVERT(x_coords < (
VEC_INT)(INPUT1_WIDTH + INPUT2_WIDTH + INPUT3_WIDTH), SELECT_TYPE);
248 VEC_TYPE values0 =
select(src2_values, src1_values, cond_in2);
249 values0 =
select(src3_values, values0, cond_in3);
250 values0 =
select(src4_values, values0, cond_in4);
258 #if defined(WIDTH_OFFSET) && defined(DEPTH) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER) 289 __kernel
void concatenate_width(
295 const int y = get_global_id(1);
296 const int z = get_global_id(2) % (int)DEPTH;
297 const int w = get_global_id(2) / (int)DEPTH;
299 __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x *
sizeof(
DATA_TYPE) + y * src_stride_y + z * src_stride_z + w * src_stride_w;
300 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x *
sizeof(
DATA_TYPE) + y * dst_stride_y + z * dst_stride_z + w * dst_stride_w;
304 #if defined(OFFSET_IN1) && defined(OFFSET_OUT) && defined(SCALE_IN1) && defined(SCALE_OUT) 305 const VEC_QUANT out0 = requantize(source_values0, OFFSET_IN1, OFFSET_OUT, SCALE_IN1, SCALE_OUT);
314 #if defined(VEC_SIZE_LEFTOVER) 316 #if defined(HEIGHT_OFFSET) && defined(DEPTH) && defined(VEC_SIZE) 348 __kernel
void concatenate_height(
354 __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x_offs + get_global_id(1) * src_stride_y + (get_global_id(2) % DEPTH) * src_stride_z + (get_global_id(
355 2) / DEPTH) * src_stride_w;
356 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_offs + get_global_id(1) * dst_stride_y + (get_global_id(2) % DEPTH) * dst_stride_z + (get_global_id(
357 2) / DEPTH) * dst_stride_w;
361 #if defined(OFFSET_IN1) && defined(OFFSET_OUT) && defined(SCALE_IN1) && defined(SCALE_OUT) 362 const VEC_QUANT out0 = requantize(source_values0, OFFSET_IN1, OFFSET_OUT, SCALE_IN1, SCALE_OUT);
395 __kernel
void concatenate(
402 __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x_offs + get_global_id(1) * src_stride_y + get_global_id(2) * src_stride_z;
403 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + x_offs + get_global_id(1) * dst_stride_y + get_global_id(2) * dst_stride_z;
407 #if defined(OFFSET_IN1) && defined(OFFSET_OUT) && defined(SCALE_IN1) && defined(SCALE_OUT) 408 source_values0 = requantize(source_values0, OFFSET_IN1, OFFSET_OUT, SCALE_IN1, SCALE_OUT);
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
#define CONVERT_DOWN(x, type)
#define CONVERT_SAT(a, b)
SimpleTensor< float > src
#define STORE_VECTOR_SELECT(basename, data_type, ptr, vec_size, leftover, cond)
Store a vector that can only be partial in x.
SimpleTensor< T > select(const SimpleTensor< uint8_t > &c, const SimpleTensor< T > &x, const SimpleTensor< T > &y)
#define TENSOR4D_DECLARATION(name)
#define TENSOR3D_DECLARATION(name)