27 #if defined(FLOAT_DATA_TYPE) 28 #define ISGREATER(x, y) isgreater(x, y) 29 #define ISLESS(x, y) isless(x, y) 30 #else // !FLOAT_DATA_TYPE 32 #define ISGREATER(x, y) (x > y) ? 1 : 0 33 #define ISLESS(x, y) (x < y) ? 1 : 0 34 #else // !defined(WIDTH) 35 #define ISGREATER(x, y) select((int16)0, (int16)-1, x > y) 36 #define ISLESS(x, y) select((int16)0, (int16)-1, x < y) 37 #endif // defined(WIDTH) 38 #endif // defined(FLOAT_DATA_TYPE) 49 in = vload16(0, input);
53 in.s01234567 += in.s89ABCDEF;
57 return (in.s0 + in.s1);
69 in = vload16(0, input);
71 in.s01234567 += in.s89ABCDEF;
75 return (in.s0 + in.s1);
87 in = vload16(0, input);
89 in.s01234567 *= in.s89ABCDEF;
93 return (in.s0 * in.s1);
95 #if defined(OPERATION) 118 __kernel
void reduction_operation_x(
126 unsigned int lsize = get_local_size(0);
127 unsigned int lid = get_local_id(0);
129 for(
unsigned int y = 0; y < get_local_size(1); ++y)
131 local_results[lid] = OPERATION((__global
DATA_TYPE *)
offset(&src, 0, y));
132 barrier(CLK_LOCAL_MEM_FENCE);
135 for(
unsigned int i = lsize >> 1; i > 0; i >>= 1)
140 local_results[lid] *= local_results[lid + i];
141 #else // !defined(PROD) 142 local_results[lid] += local_results[lid + i];
143 #endif // defined(PROD) 145 barrier(CLK_LOCAL_MEM_FENCE);
150 #if defined(MEAN) && defined(WIDTH) 151 if(y == get_local_size(1) - 1)
153 local_results[0] /= WIDTH;
155 #endif // defined(MEAN) && defined(WIDTH) 156 ((__global
DATA_TYPE *)
offset(&partial_res, get_group_id(0), y))[0] = local_results[0];
160 #endif // defined(OPERATION) 179 __kernel
void reduction_operation_non_parallel_x(
189 #if defined(PROD) && defined(OFFSET) && defined(SCALE) 190 float res_f =
DEQUANTIZE(res, OFFSET, SCALE, DATA_TYPE_PROMOTED, 1);
191 #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) 193 for(
unsigned int x = 1; x < WIDTH; ++x)
201 #if defined(OFFSET) && defined(SCALE) 202 res_f *=
DEQUANTIZE(in, OFFSET, SCALE, DATA_TYPE_PROMOTED, 1);
203 #else // !(defined(OFFSET) && defined(SCALE)) 205 #endif // defined(OFFSET) && defined(SCALE) 206 #else // defined(SUM)) 208 #endif // defined(MAX) || defined(MIN) || defined(PROD) 214 #endif // defined(MEAN) 217 #if defined(SUM) && defined(OFFSET) && defined(SCALE) 218 res -= (WIDTH - 1) * OFFSET;
219 #endif // defined(OFFSET) && defined(OFFSET) && defined(SCALE) 222 #if defined(PROD) && defined(OFFSET) && defined(SCALE) 223 res =
QUANTIZE(res_f, OFFSET, SCALE, DATA_TYPE_PROMOTED, 1);
224 #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) 228 #endif // defined(WIDTH) 249 __kernel
void reduction_operation_y(
260 #if defined(PROD) && defined(OFFSET) && defined(SCALE) 261 float16 res_f =
DEQUANTIZE(res, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
262 #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) 264 #if defined(SUM_SQUARE) 266 #endif // defined(SUM_SQUARE) 268 for(
unsigned int y = 1; y < HEIGHT; ++y)
276 #else // !(defined(MAX) || defined(MIN)) 277 #if defined(SUM_SQUARE) 279 #endif // defined(SUM_SQUARE) 282 #if defined(OFFSET) && defined(SCALE) 283 res_f *=
DEQUANTIZE(in, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
284 #else // !(defined(OFFSET) && defined(SCALE)) 286 #endif // defined(OFFSET) && defined(SCALE) 288 #else // !defined(PROD) 290 #endif // defined(PROD) 291 #endif // defined(MAX) || defined(MIN) 296 #endif // defined(MEAN) 299 #if defined(SUM) && defined(OFFSET) && defined(SCALE) 300 res -= (HEIGHT - 1) * OFFSET;
301 #endif // defined(OFFSET) && defined(OFFSET) && defined(SCALE) 304 #if defined(PROD) && defined(OFFSET) && defined(SCALE) 305 res =
QUANTIZE(res_f, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
306 #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) 311 #endif // defined(HEIGHT) 336 __kernel
void reduction_operation_z(
347 #if defined(PROD) && defined(OFFSET) && defined(SCALE) 348 float16 res_f =
DEQUANTIZE(res, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
349 #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) 354 #endif // defined(COMPLEX) 355 #if defined(SUM_SQUARE) 357 #endif // defined(SUM_SQUARE) 359 for(
unsigned int z = 1; z < DEPTH; ++z)
367 #endif // defined(COMPLEX) 373 #else // !(defined(MAX) || defined(MIN)) 374 #if defined(SUM_SQUARE) 376 #endif // defined(SUM_SQUARE) 379 #if defined(OFFSET) && defined(SCALE) 380 res_f *=
DEQUANTIZE(in, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
381 #else // !(defined(OFFSET) && defined(SCALE)) 383 #endif // defined(OFFSET) && defined(SCALE) 385 #else // !defined(PROD) 389 #endif // defined(COMPLEX) 390 #endif // defined(PROD) 391 #endif // defined(MAX) || defined(MIN) 396 #endif // defined(MEAN) 399 #if defined(SUM) && defined(OFFSET) && defined(SCALE) 400 res -= (DEPTH - 1) * OFFSET;
401 #endif // defined(OFFSET) && defined(OFFSET) && defined(SCALE) 404 #if defined(PROD) && defined(OFFSET) && defined(SCALE) 405 res =
QUANTIZE(res_f, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
406 #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) 412 #endif // defined(COMPLEX) 416 #if defined(BATCH) && defined(DEPTH) 444 __kernel
void reduction_operation_w(
455 #if defined(PROD) && defined(OFFSET) && defined(SCALE) 456 float16 res_f =
DEQUANTIZE(res, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
457 #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) 459 #if defined(SUM_SQUARE) 461 #endif // defined(SUM_SQUARE) 463 for(
unsigned int w = 1;
w < BATCH; ++
w)
472 #else // !(defined(MAX) || defined(MIN)) 473 #if defined(SUM_SQUARE) 475 #endif // defined(SUM_SQUARE) 478 #if defined(OFFSET) && defined(SCALE) 479 res_f *=
DEQUANTIZE(in, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
480 #else // !(defined(OFFSET) && defined(SCALE)) 482 #endif // defined(OFFSET) && defined(SCALE) 484 #else // !defined(PROD) 486 #endif //defined(PROD) 487 #endif // defined(MAX) || defined(MIN) 492 #endif // defined(MEAN) 495 #if defined(SUM) && defined(OFFSET) && defined(SCALE) 496 res -= (BATCH - 1) * OFFSET;
497 #endif // defined(OFFSET) && defined(OFFSET) && defined(SCALE) 500 #if defined(PROD) && defined(OFFSET) && defined(SCALE) 501 res =
QUANTIZE(res_f, OFFSET, SCALE, DATA_TYPE_PROMOTED, 16);
502 #endif // defined(PROD) && defined(OFFSET) && defined(SCALE) Structure to hold Vector information.
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
DATA_TYPE product(__global const DATA_TYPE *input)
Calculate product of a vector.
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define CONVERT_TO_IMAGE_STRUCT(name)
DATA_TYPE sum(__global const DATA_TYPE *input)
Calculate sum of a vector.
#define IMAGE_DECLARATION(name)
#define CONVERT_SAT(a, b)
Structure to hold 3D tensor information.
SimpleTensor< float > src
Structure to hold 4D tensor information.
#define CONVERT_TO_VECTOR_STRUCT(name)
#define DEQUANTIZE(input, offset, scale, type, size)
#define VECTOR_DECLARATION(name)
__global const uchar * tensor4D_offset(const Tensor4D *tensor, int x, int y, int z, int w)
Get the pointer position of a Tensor4D.
#define CONVERT_TO_TENSOR4D_STRUCT(name, mod_size)
__global uchar * ptr
Pointer to the starting postion of the buffer.
Structure to hold Image information.
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define CONVERT_TO_TENSOR3D_STRUCT(name)
SimpleTensor< T > select(const SimpleTensor< uint8_t > &c, const SimpleTensor< T > &x, const SimpleTensor< T > &y)
__global const uchar * vector_offset(const Vector *vec, int x)
Get the pointer position of a Vector.
#define QUANTIZE(input, offset, scale, type, size)
#define TENSOR4D_DECLARATION(name)
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define TENSOR3D_DECLARATION(name)
DATA_TYPE square_sum(__global const DATA_TYPE *input)
Calculate square sum of a vector.
__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)