26 #if defined(VEC_SIZE) && defined(DATA_TYPE) 28 #if defined(FIXED_POINT_POSITION) 31 #define ADD_OP(a, b) ADD_SAT_OP_EXPAND((a), (b), DATA_TYPE, VEC_SIZE) 32 #define SUB_OP(a, b) SUB_SAT_OP_EXPAND((a), (b), DATA_TYPE, VEC_SIZE) 33 #define MUL_OP(a, b) MUL_SAT_OP_EXPAND((a), (b), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION) 34 #define INVSQRT_OP(a) INVSQRT_OP_EXPAND((a), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION) 35 #define SQCVT_SAT(a) SQCVT_SAT_OP_EXPAND((a), DATA_TYPE, FIXED_POINT_POSITION) 39 #define ADD_OP(a, b) ((a) + (b)) 40 #define SUB_OP(a, b) ((a) - (b)) 41 #define MUL_OP(a, b) ((a) * (b)) 42 #define INVSQRT_OP(a) rsqrt((a)) 43 #define SQCVT_SAT(a) (a) 48 #define ACTIVATION_FUNC(x) CLAMP(x, (DATA_TYPE)B_VAL, (DATA_TYPE)A_VAL) 50 #define ACTIVATION_FUNC(x) CLAMP(x, (DATA_TYPE)0, (DATA_TYPE)A_VAL) 52 #define ACTIVATION_FUNC(x) max(x, (DATA_TYPE)0) 54 #define ACTIVATION_FUNC(x) (x) 127 const
int current_slice = get_global_id(2);
129 data =
VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in.ptr);
130 denominator = *((__global DATA_TYPE *)(var.ptr + current_slice * var.stride_x));
134 numerator = *((__global DATA_TYPE *)(mean.ptr + current_slice * mean.stride_x));
135 numerator =
SUB_OP(data, numerator);
136 x_bar =
MUL_OP(numerator, denominator);
138 gamma_vec = *((__global DATA_TYPE *)(gamma.ptr + current_slice * gamma.stride_x));
139 beta_vec = *((__global DATA_TYPE *)(beta.ptr + current_slice * beta.stride_x));
144 res = ACTIVATION_FUNC(res);
147 (res, 0, (__global DATA_TYPE *)out.ptr);
Structure to hold Vector information.
#define CONVERT_TO_TENSOR3D_STRUCT(name)
#define CONVERT_TO_VECTOR_STRUCT(name)
Structure to hold 3D tensor information.
#define VECTOR_DECLARATION(name)
#define TENSOR3D_DECLARATION(name)
#define VEC_DATA_TYPE(type, size)