27 #if defined(WEIGHTS_OFFSET) && defined(INPUT_OFFSET) && defined(K_OFFSET) && defined(OUTPUT_OFFSET) && defined(OUTPUT_MULTIPLIER) && defined(OUTPUT_SHIFT) 29 #if defined(FUSED_ACTIVATION) 30 #define DATA_TYPE uchar 35 #define ACTIVATION_FUNC(x) PERFORM_ACTIVATION_QA8(FUSED_ACTIVATION, x) 37 #define ACTIVATION_FUNC(x) (x) 40 #if defined(CONV_STRIDE_Y) && defined(CONV_STRIDE_X) 43 #error "Stride X not supported" 46 #if CONV_STRIDE_X == 1 47 #define GET_VALUES(first_value, left, middle, right) \ 49 int8 temp0 = CONVERT(vload8(0, first_value), int8); \ 50 int2 temp1 = CONVERT(vload2(0, (first_value + 8 * sizeof(uchar))), int2); \ 52 left = CONVERT(temp0.s01234567, int8); \ 53 middle = CONVERT((int8)(temp0.s1234, temp0.s567, temp1.s0), int8); \ 54 right = CONVERT((int8)(temp0.s2345, temp0.s67, temp1.s01), int8); \ 56 #elif CONV_STRIDE_X == 2 57 #define GET_VALUES(first_value, left, middle, right) \ 59 int16 temp0 = CONVERT(vload16(0, first_value), int16); \ 60 int temp1 = CONVERT(*(first_value + 16 * sizeof(uchar)), int); \ 62 left = CONVERT(temp0.s02468ace, int8); \ 63 middle = CONVERT(temp0.s13579bdf, int8); \ 64 right = CONVERT((int8)(temp0.s2468, temp0.sace, temp1), int8); \ 67 #define GET_VALUES(first_value, left, middle, right) \ 69 int16 temp0 = CONVERT(vload16(0, first_value), int16); \ 70 int8 temp1 = CONVERT(vload8(0, (first_value + 16 * sizeof(uchar))), int8); \ 72 left = CONVERT((int8)(temp0.s0369, temp0.scf, temp1.s25), int8); \ 73 middle = CONVERT((int8)(temp0.s147a, temp0.sd, temp1.s036), int8); \ 74 right = CONVERT((int8)(temp0.s258b, temp0.se, temp1.s147), int8); \ 110 __kernel
void depthwise_convolution_3x3_quantized_nchw(
114 #
if defined(HAS_BIAS)
123 #if defined(HAS_BIAS) 126 int bias_value = *((__global
int *)(
vector_offset(&biases, get_global_id(2))));
127 #endif //defined(HAS_BIAS) 129 src.
ptr -= (get_global_id(2) - get_global_id(2) / DEPTH_MULTIPLIER) * src_step_z;
131 uchar3 w0 = vload3(0, weights.
ptr + 0 * weights_stride_y);
132 uchar3 w1 = vload3(0, weights.
ptr + 1 * weights_stride_y);
133 uchar3 w2 = vload3(0, weights.
ptr + 2 * weights_stride_y);
137 #if CONV_STRIDE_Y == 1 143 int8 left, middle, right;
144 GET_VALUES(src.
ptr + 0 * src_stride_y, left, middle, right);
145 values0 += left * (int8)(w0.s0);
146 values0 += middle * (int8)(w0.s1);
147 values0 += right * (int8)(w0.s2);
149 #if WEIGHTS_OFFSET != 0 150 sum0 += left + middle + right;
154 GET_VALUES(src.
ptr + 1 * src_stride_y, left, middle, right);
155 values0 += left * (int8)(w1.s0);
156 values0 += middle * (int8)(w1.s1);
157 values0 += right * (int8)(w1.s2);
158 #if CONV_STRIDE_Y == 1 159 values1 += left * (int8)(w0.s0);
160 values1 += middle * (int8)(w0.s1);
161 values1 += right * (int8)(w0.s2);
164 #if WEIGHTS_OFFSET != 0 165 int8 tmp = left + middle + right;
167 #if CONV_STRIDE_Y == 1 173 GET_VALUES(src.
ptr + 2 * src_stride_y, left, middle, right);
174 values0 += left * (int8)(w2.s0);
175 values0 += middle * (int8)(w2.s1);
176 values0 += right * (int8)(w2.s2);
177 #if CONV_STRIDE_Y == 1 178 values1 += left * (int8)(w1.s0);
179 values1 += middle * (int8)(w1.s1);
180 values1 += right * (int8)(w1.s2);
183 #if WEIGHTS_OFFSET != 0 184 tmp = left + middle + right;
186 #if CONV_STRIDE_Y == 1 191 #if CONV_STRIDE_Y == 1 193 GET_VALUES(src.
ptr + 3 * src_stride_y, left, middle, right);
194 values1 += left * (int8)(w2.s0);
195 values1 += middle * (int8)(w2.s1);
196 values1 += right * (int8)(w2.s2);
198 #if WEIGHTS_OFFSET != 0 199 sum1 += left + middle + right;
203 #if defined(HAS_BIAS) 204 values0 += (int8)(bias_value);
205 #if CONV_STRIDE_Y == 1 206 values1 += (int8)(bias_value);
208 #endif //defined(HAS_BIAS) 210 #if WEIGHTS_OFFSET != 0 211 values0 += sum0 * (int8)(WEIGHTS_OFFSET);
212 #if CONV_STRIDE_Y == 1 213 values1 += sum1 * (int8)(WEIGHTS_OFFSET);
217 #if INPUT_OFFSET != 0 218 ushort sum_weights = 0;
219 ushort3 tmp_we = convert_ushort3(w0) + convert_ushort3(w1) + convert_ushort3(w2);
220 sum_weights += tmp_we.s0 + tmp_we.s1 + tmp_we.s2;
221 values0 += sum_weights * (int8)(INPUT_OFFSET);
222 #if CONV_STRIDE_Y == 1 223 values1 += sum_weights * (int8)(INPUT_OFFSET);
228 values0 += (int8)(K_OFFSET);
229 #if CONV_STRIDE_Y == 1 230 values1 += (int8)(K_OFFSET);
235 values0 += (int8)OUTPUT_OFFSET;
236 uchar8 res0 = convert_uchar8_sat(values0);
237 res0 =
max(res0, (uchar8)0);
238 res0 =
min(res0, (uchar8)255);
240 vstore8(ACTIVATION_FUNC(res0), 0, dst.
ptr);
241 #if CONV_STRIDE_Y == 1 244 values1 += (int8)OUTPUT_OFFSET;
245 uchar8 res1 = convert_uchar8_sat(values1);
246 res1 =
max(res1, (uchar8)0);
247 res1 =
min(res1, (uchar8)255);
249 vstore8(ACTIVATION_FUNC(res1), 0, dst.
ptr + dst_stride_y);
255 #if defined(VEC_SIZE) && defined(SRC_DEPTH) && defined(CONV_PAD_TOP) && defined(ROWS_READ) 257 #define asymm_mult_by_quant_multiplier_less_than_one(x, y, z) ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(x, y, z, VEC_SIZE) 259 #define VEC_INT VEC_DATA_TYPE(int, VEC_SIZE) 260 #define VEC_UCHAR VEC_DATA_TYPE(uchar, VEC_SIZE) 262 #define BIFROST_MAD_4(acc, x, y) \ 264 acc.s0 += (ushort)x.s0 * (ushort)y.s0; \ 265 acc.s1 += (ushort)x.s1 * (ushort)y.s1; \ 266 acc.s2 += (ushort)x.s2 * (ushort)y.s2; \ 267 acc.s3 += (ushort)x.s3 * (ushort)y.s3; \ 270 #if WEIGHTS_OFFSET != 0 271 #define BIFROST_MAD_ACC_4(acc, sum, x, y) \ 273 sum += CONVERT(x, VEC_INT); \ 274 BIFROST_MAD_4(acc, x, y); \ 277 #define BIFROST_MAD_ACC_4(acc, sum, x, y) BIFROST_MAD_4(acc, x, y) 312 __kernel
void depthwise_convolution_3x3_quantized_nhwc_stride1(
316 #
if defined(HAS_BIAS)
323 #if defined(HAS_BIAS) 329 __global uchar *first_elem = src_ptr + src_offset_first_element_in_bytes;
331 const int z = get_global_id(2);
332 const int pad_offs = -ROWS_READ * src_stride_y;
333 const int src_offs0 = get_global_id(0) * src_step_x + get_global_id(1) * src_step_y + z * src_step_z - CONV_PAD_TOP * src_stride_z;
334 const int src_offs1 = src_offs0 + src_stride_z;
335 const int src_offs2 = src_offs1 + src_stride_z;
337 const int cond_top = z - CONV_PAD_TOP < 0;
338 const int cond_bottom = z * (src_step_z / src_stride_z) + 2 > SRC_DEPTH;
340 __global uchar *src_addr0 = first_elem + select(src_offs0, pad_offs, cond_top);
341 __global uchar *src_addr1 = first_elem + src_offs1;
342 __global uchar *src_addr2 = first_elem + select(src_offs2, pad_offs, cond_bottom);
345 VEC_INT acc0 = 0, acc1 = 0, acc2 = 0, acc3 = 0;
346 VEC_INT sum0 = 0, sum1 = 0, sum2 = 0, sum3 = 0;
350 w0 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 0 * weights_stride_y);
351 w1 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 1 * weights_stride_y);
352 w2 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 2 * weights_stride_y);
354 #if INPUT_OFFSET != 0 359 BIFROST_MAD_ACC_4(acc0, sum0, values, w0);
361 src_addr0 += src_stride_y;
362 values =
VLOAD(VEC_SIZE)(0, src_addr0);
363 BIFROST_MAD_ACC_4(acc0, sum0, values, w1);
364 BIFROST_MAD_ACC_4(acc1, sum1, values, w0);
366 src_addr0 += src_stride_y;
367 values =
VLOAD(VEC_SIZE)(0, src_addr0);
368 BIFROST_MAD_ACC_4(acc0, sum0, values, w2);
369 BIFROST_MAD_ACC_4(acc1, sum1, values, w1);
370 BIFROST_MAD_ACC_4(acc2, sum2, values, w0);
372 src_addr0 += src_stride_y;
373 values =
VLOAD(VEC_SIZE)(0, src_addr0);
374 BIFROST_MAD_ACC_4(acc1, sum1, values, w2);
375 BIFROST_MAD_ACC_4(acc2, sum2, values, w1);
376 BIFROST_MAD_ACC_4(acc3, sum3, values, w0);
378 src_addr0 += src_stride_y;
379 values =
VLOAD(VEC_SIZE)(0, src_addr0);
380 BIFROST_MAD_ACC_4(acc2, sum2, values, w2);
381 BIFROST_MAD_ACC_4(acc3, sum3, values, w1);
383 src_addr0 += src_stride_y;
384 values =
VLOAD(VEC_SIZE)(0, src_addr0);
385 BIFROST_MAD_ACC_4(acc3, sum3, values, w2);
387 weights.
ptr += weights_stride_z;
390 w0 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 0 * weights_stride_y);
391 w1 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 1 * weights_stride_y);
392 w2 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 2 * weights_stride_y);
394 #if INPUT_OFFSET != 0 398 values =
VLOAD(VEC_SIZE)(0, src_addr1);
399 BIFROST_MAD_ACC_4(acc0, sum0, values, w0);
401 src_addr1 += src_stride_y;
402 values =
VLOAD(VEC_SIZE)(0, src_addr1);
403 BIFROST_MAD_ACC_4(acc0, sum0, values, w1);
404 BIFROST_MAD_ACC_4(acc1, sum1, values, w0);
406 src_addr1 += src_stride_y;
407 values =
VLOAD(VEC_SIZE)(0, src_addr1);
408 BIFROST_MAD_ACC_4(acc0, sum0, values, w2);
409 BIFROST_MAD_ACC_4(acc1, sum1, values, w1);
410 BIFROST_MAD_ACC_4(acc2, sum2, values, w0);
412 src_addr1 += src_stride_y;
413 values =
VLOAD(VEC_SIZE)(0, src_addr1);
414 BIFROST_MAD_ACC_4(acc1, sum1, values, w2);
415 BIFROST_MAD_ACC_4(acc2, sum2, values, w1);
416 BIFROST_MAD_ACC_4(acc3, sum3, values, w0);
418 src_addr1 += src_stride_y;
419 values =
VLOAD(VEC_SIZE)(0, src_addr1);
420 BIFROST_MAD_ACC_4(acc2, sum2, values, w2);
421 BIFROST_MAD_ACC_4(acc3, sum3, values, w1);
423 src_addr1 += src_stride_y;
424 values =
VLOAD(VEC_SIZE)(0, src_addr1);
425 BIFROST_MAD_ACC_4(acc3, sum3, values, w2);
427 weights.
ptr += weights_stride_z;
430 w0 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 0 * weights_stride_y);
431 w1 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 1 * weights_stride_y);
432 w2 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 2 * weights_stride_y);
434 #if INPUT_OFFSET != 0 438 values =
VLOAD(VEC_SIZE)(0, src_addr2);
439 BIFROST_MAD_ACC_4(acc0, sum0, values, w0);
441 src_addr2 += src_stride_y;
442 values =
VLOAD(VEC_SIZE)(0, src_addr2);
443 BIFROST_MAD_ACC_4(acc0, sum0, values, w1);
444 BIFROST_MAD_ACC_4(acc1, sum1, values, w0);
446 src_addr2 += src_stride_y;
447 values =
VLOAD(VEC_SIZE)(0, src_addr2);
448 BIFROST_MAD_ACC_4(acc0, sum0, values, w2);
449 BIFROST_MAD_ACC_4(acc1, sum1, values, w1);
450 BIFROST_MAD_ACC_4(acc2, sum2, values, w0);
452 src_addr2 += src_stride_y;
453 values =
VLOAD(VEC_SIZE)(0, src_addr2);
454 BIFROST_MAD_ACC_4(acc1, sum1, values, w2);
455 BIFROST_MAD_ACC_4(acc2, sum2, values, w1);
456 BIFROST_MAD_ACC_4(acc3, sum3, values, w0);
458 src_addr2 += src_stride_y;
459 values =
VLOAD(VEC_SIZE)(0, src_addr2);
460 BIFROST_MAD_ACC_4(acc2, sum2, values, w2);
461 BIFROST_MAD_ACC_4(acc3, sum3, values, w1);
463 src_addr2 += src_stride_y;
464 values =
VLOAD(VEC_SIZE)(0, src_addr2);
465 BIFROST_MAD_ACC_4(acc3, sum3, values, w2);
467 #if defined(HAS_BIAS) 474 #if WEIGHTS_OFFSET != 0 475 acc0 += WEIGHTS_OFFSET * sum0;
476 acc1 += WEIGHTS_OFFSET * sum1;
477 acc2 += WEIGHTS_OFFSET * sum2;
478 acc3 += WEIGHTS_OFFSET * sum3;
481 #if INPUT_OFFSET != 0 482 VEC_INT offs = INPUT_OFFSET * sum_we;
497 acc0 = asymm_mult_by_quant_multiplier_less_than_one(acc0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT);
498 acc1 = asymm_mult_by_quant_multiplier_less_than_one(acc1, OUTPUT_MULTIPLIER, OUTPUT_SHIFT);
499 acc2 = asymm_mult_by_quant_multiplier_less_than_one(acc2, OUTPUT_MULTIPLIER, OUTPUT_SHIFT);
500 acc3 = asymm_mult_by_quant_multiplier_less_than_one(acc3, OUTPUT_MULTIPLIER, OUTPUT_SHIFT);
502 acc0 += (
VEC_INT)OUTPUT_OFFSET;
503 acc1 += (
VEC_INT)OUTPUT_OFFSET;
504 acc2 += (
VEC_INT)OUTPUT_OFFSET;
505 acc3 += (
VEC_INT)OUTPUT_OFFSET;
518 (res0, 0, dst.
ptr + 0 * dst_stride_y);
520 (res1, 0, dst.
ptr + 1 * dst_stride_y);
522 (res2, 0, dst.
ptr + 2 * dst_stride_y);
524 (res3, 0, dst.
ptr + 3 * dst_stride_y);
559 __kernel
void depthwise_convolution_3x3_quantized_nhwc_stride2(
563 #
if defined(HAS_BIAS)
570 #if defined(HAS_BIAS) 576 __global uchar *first_elem = src_ptr + src_offset_first_element_in_bytes;
578 const int z = get_global_id(2);
579 const int pad_offs = -ROWS_READ * src_stride_y;
580 const int src_offs0 = get_global_id(0) * src_step_x + get_global_id(1) * src_step_y + z * src_step_z - CONV_PAD_TOP * src_stride_z;
581 const int src_offs1 = src_offs0 + src_stride_z;
582 const int src_offs2 = src_offs1 + src_stride_z;
584 const int cond_top = z - CONV_PAD_TOP < 0;
585 const int cond_bottom = z * (src_step_z / src_stride_z) + 2 > SRC_DEPTH;
587 __global uchar *src_addr0 = first_elem + select(src_offs0, pad_offs, cond_top);
588 __global uchar *src_addr1 = first_elem + src_offs1;
589 __global uchar *src_addr2 = first_elem + select(src_offs2, pad_offs, cond_bottom);
597 w0 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 0 * weights_stride_y);
598 w1 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 1 * weights_stride_y);
599 w2 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 2 * weights_stride_y);
601 #if INPUT_OFFSET != 0 606 BIFROST_MAD_ACC_4(acc0, sum0, values, w0);
608 src_addr0 += src_stride_y;
609 values =
VLOAD(VEC_SIZE)(0, src_addr0);
610 BIFROST_MAD_ACC_4(acc0, sum0, values, w1);
612 src_addr0 += src_stride_y;
613 values =
VLOAD(VEC_SIZE)(0, src_addr0);
614 BIFROST_MAD_ACC_4(acc0, sum0, values, w2);
615 BIFROST_MAD_ACC_4(acc2, sum2, values, w0);
617 src_addr0 += src_stride_y;
618 values =
VLOAD(VEC_SIZE)(0, src_addr0);
619 BIFROST_MAD_ACC_4(acc2, sum2, values, w1);
621 src_addr0 += src_stride_y;
622 values =
VLOAD(VEC_SIZE)(0, src_addr0);
623 BIFROST_MAD_ACC_4(acc2, sum2, values, w2);
625 weights.
ptr += weights_stride_z;
628 w0 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 0 * weights_stride_y);
629 w1 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 1 * weights_stride_y);
630 w2 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 2 * weights_stride_y);
632 #if INPUT_OFFSET != 0 636 values =
VLOAD(VEC_SIZE)(0, src_addr1);
637 BIFROST_MAD_ACC_4(acc0, sum0, values, w0);
639 src_addr1 += src_stride_y;
640 values =
VLOAD(VEC_SIZE)(0, src_addr1);
641 BIFROST_MAD_ACC_4(acc0, sum0, values, w1);
643 src_addr1 += src_stride_y;
644 values =
VLOAD(VEC_SIZE)(0, src_addr1);
645 BIFROST_MAD_ACC_4(acc0, sum0, values, w2);
646 BIFROST_MAD_ACC_4(acc2, sum2, values, w0);
648 src_addr1 += src_stride_y;
649 values =
VLOAD(VEC_SIZE)(0, src_addr1);
650 BIFROST_MAD_ACC_4(acc2, sum2, values, w1);
652 src_addr1 += src_stride_y;
653 values =
VLOAD(VEC_SIZE)(0, src_addr1);
654 BIFROST_MAD_ACC_4(acc2, sum2, values, w2);
656 weights.
ptr += weights_stride_z;
659 w0 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 0 * weights_stride_y);
660 w1 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 1 * weights_stride_y);
661 w2 =
VLOAD(VEC_SIZE)(0, weights.
ptr + 2 * weights_stride_y);
663 #if INPUT_OFFSET != 0 667 values =
VLOAD(VEC_SIZE)(0, src_addr2);
668 BIFROST_MAD_ACC_4(acc0, sum0, values, w0);
670 src_addr2 += src_stride_y;
671 values =
VLOAD(VEC_SIZE)(0, src_addr2);
672 BIFROST_MAD_ACC_4(acc0, sum0, values, w1);
674 src_addr2 += src_stride_y;
675 values =
VLOAD(VEC_SIZE)(0, src_addr2);
676 BIFROST_MAD_ACC_4(acc0, sum0, values, w2);
677 BIFROST_MAD_ACC_4(acc2, sum2, values, w0);
679 src_addr2 += src_stride_y;
680 values =
VLOAD(VEC_SIZE)(0, src_addr2);
681 BIFROST_MAD_ACC_4(acc2, sum2, values, w1);
683 src_addr2 += src_stride_y;
684 values =
VLOAD(VEC_SIZE)(0, src_addr2);
685 BIFROST_MAD_ACC_4(acc2, sum2, values, w2);
687 #if defined(HAS_BIAS) 692 #if WEIGHTS_OFFSET != 0 693 acc0 += WEIGHTS_OFFSET * sum0;
694 acc2 += WEIGHTS_OFFSET * sum2;
697 #if INPUT_OFFSET != 0 698 VEC_INT offs = INPUT_OFFSET * sum_we;
709 acc0 = asymm_mult_by_quant_multiplier_less_than_one(acc0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT);
710 acc2 = asymm_mult_by_quant_multiplier_less_than_one(acc2, OUTPUT_MULTIPLIER, OUTPUT_SHIFT);
711 acc0 += (
VEC_INT)OUTPUT_OFFSET;
712 acc2 += (
VEC_INT)OUTPUT_OFFSET;
719 (res0, 0, dst.
ptr + 0 * dst_stride_y);
721 (res2, 0, dst.
ptr + 1 * dst_stride_y);
Structure to hold Vector information.
fixed_point< T > min(fixed_point< T > x, fixed_point< T > y)
#define CONVERT_TO_TENSOR3D_STRUCT(name)
#define CONVERT_TO_VECTOR_STRUCT(name)
#define CONVERT_SAT(a, b)
Structure to hold 3D tensor information.
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define CLAMP(x, min_val, max_val)
#define VECTOR_DECLARATION(name)
Structure to hold Image information.
#define TENSOR3D_DECLARATION(name)
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define CONVERT_TO_VECTOR_STRUCT_NO_STEP(name)
__global const uchar * vector_offset(const Vector *vec, int x)
Get the pointer position of a Vector.
#define ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(x, quantized_multiplier, right_shift, size)
#define CONVERT_TENSOR3D_TO_IMAGE_STRUCT(name)
__global uchar * ptr
Pointer to the starting postion of the buffer.
fixed_point< T > max(fixed_point< T > x, fixed_point< T > y)
convolution configure & src