26 #if defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST) && defined(SCALE) && defined(OFFSET) 61 #if defined(LAST_ACCESSED_X) 64 const int xi = (int)(get_global_id(0) *
VEC_SIZE);
65 input.
ptr -= max(xi - (
int)LAST_ACCESSED_X, 0) * input_stride_x;
66 output.
ptr -= max(xi - (
int)LAST_ACCESSED_X, 0) * output_stride_x;
86 #else // !defined(LAST_ACCESSED_X) 87 *((__global DATA_TYPE_DST *)(output.
ptr)) = (DATA_TYPE_DST)((float)((
int)(*((__global DATA_TYPE_SRC *)(input.
ptr))) - (int)(OFFSET)) * (
float)(SCALE));
88 #endif // defined(LAST_ACCESSED_X) 90 #endif // defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST) && defined(SCALE) && defined(OFFSET) 92 #if defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST) 117 __kernel
void dequantization_layer_per_channel_nchw(
120 __global
float *
scale)
126 #if defined(LAST_ACCESSED_X) 129 const int xi = (int)(get_global_id(0) *
VEC_SIZE);
130 input.
ptr -= max(xi - (
int)LAST_ACCESSED_X, 0) * input_stride_x;
131 output.
ptr -= max(xi - (
int)LAST_ACCESSED_X, 0) * output_stride_x;
139 vscale = scale[get_global_id(2)];
148 #else // !defined(LAST_ACCESSED_X) 149 *((__global DATA_TYPE_DST *)(output.
ptr)) = (DATA_TYPE_DST)((float)((
int)(*((__global DATA_TYPE_SRC *)(input.
ptr)))) * scale[get_global_id(2)]);
150 #endif // defined(LAST_ACCESSED_X) 176 __kernel
void dequantization_layer_per_channel_nhwc(
179 __global
float *scale)
185 #if defined(LAST_ACCESSED_X) 188 const int xi = (int)(get_global_id(0) *
VEC_SIZE);
189 input.
ptr -= max(xi - (
int)LAST_ACCESSED_X, 0) * input_stride_x;
190 output.
ptr -= max(xi - (
int)LAST_ACCESSED_X, 0) * output_stride_x;
191 scale -= max(xi - (
int)LAST_ACCESSED_X, 0);
208 #else // !defined(LAST_ACCESSED_X) 209 *((__global DATA_TYPE_DST *)(output.
ptr)) = (DATA_TYPE_DST)((float)((
int)(*((__global DATA_TYPE_SRC *)(input.
ptr)))) * scale[get_global_id(0)]);
210 #endif // defined(LAST_ACCESSED_X) 212 #endif // defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST)
SimpleTensor< TOut > dequantization_layer(const SimpleTensor< TIn > &src)
Structure to hold 3D tensor information.
#define CONVERT_TO_TENSOR3D_STRUCT(name)
__global uchar * ptr
Pointer to the starting postion of the buffer.
#define TENSOR3D_DECLARATION(name)
#define VEC_DATA_TYPE(type, size)