26 #define VEC_FLOAT VEC_DATA_TYPE(float, VEC_SIZE)
28 #if defined(FLOAT_DOMAIN)
62 __kernel
void activation_layer_quant_f32(
70 uint x_offs = max((
int)(get_global_id(0) *
VEC_SIZE *
sizeof(DATA_TYPE) - (
VEC_SIZE - VEC_SIZE_LEFTOVER) %
VEC_SIZE *
sizeof(DATA_TYPE)), 0);
73 __global uchar *input_addr = input_ptr + input_offset_first_element_in_bytes + x_offs + get_global_id(1) * input_stride_y + get_global_id(2) * input_stride_z;
75 __global uchar *output_addr = input_addr;
77 __global uchar *output_addr = output_ptr + output_offset_first_element_in_bytes + x_offs + get_global_id(1) * output_stride_y + get_global_id(2) * output_stride_z;
85 data_flt =
round(data_flt - (
float)O1_VAL) * ((float)S1_VAL);
86 #else // defined(O1_VAL)
87 data_flt =
round(data_flt) * ((float)S1_VAL);
88 #endif // defined(O1_VAL)
93 #else // defined(O2_VAL)
95 #endif // defined(O2_VAL)
101 #else // defined(FLOAT_DOMAIN)
135 __kernel
void activation_layer_quant(
143 uint x_offs = max((
int)(get_global_id(0) *
VEC_SIZE *
sizeof(DATA_TYPE) - (
VEC_SIZE - VEC_SIZE_LEFTOVER) %
VEC_SIZE *
sizeof(DATA_TYPE)), 0);
146 __global uchar *input_addr = input_ptr + input_offset_first_element_in_bytes + x_offs + get_global_id(1) * input_stride_y + get_global_id(2) * input_stride_z;
148 __global uchar *output_addr = input_addr;
150 __global uchar *output_addr = output_ptr + output_offset_first_element_in_bytes + x_offs + get_global_id(1) * output_stride_y + get_global_id(2) * output_stride_z;
161 #endif // defined(ACT)
162 #endif // defined(FLOAT_DOMAIN)