26 #if defined(DATA_TYPE) && defined(OPERATION)
28 #define rsqrt_op(input) rsqrt(input)
30 #define VEC_FLOAT VEC_DATA_TYPE(float, VEC_SIZE)
31 #define VEC_INT VEC_DATA_TYPE(int, VEC_SIZE)
32 #define VEC_TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
33 #endif // defined(VEC_SIZE)
54 __kernel
void elementwise_unary_quantized(
63 const int xi = (int)(get_global_id(0) *
VEC_SIZE);
64 in.
ptr -= max(xi - (
int)LAST_ACCESSED_X, 0) * in_stride_x;
65 out.
ptr -= max(xi - (
int)LAST_ACCESSED_X, 0) * out_stride_x;
71 data_f32 = (data_f32 - (float)OFFSET_IN) * (float)SCALE_IN;
72 VEC_INT qres_int =
CONVERT_SAT((OPERATION(data_f32) / ((
VEC_FLOAT)(
float)SCALE_OUT)), VEC_INT) + ((VEC_INT)((
int)OFFSET_OUT));
73 const VEC_TYPE qres =
CONVERT_SAT(qres_int, VEC_TYPE);
75 (qres, 0, (__global DATA_TYPE *)out.
ptr);
77 #endif // defined(DATA_TYPE) && defined(OPERATION)