26 #if defined(VECTOR_SIZE) && defined(START) && defined(STEP) && defined(DATA_TYPE) && defined(VEC_SIZE_LEFTOVER) 28 #if !defined(OFFSET_OUT) && !defined(SCALE_OUT) 31 #define STEP_VEC ((VEC_DATA_TYPE(DATA_TYPE, 2))(0, STEP)) 32 #elif VECTOR_SIZE == 3 33 #define STEP_VEC ((VEC_DATA_TYPE(DATA_TYPE, 3))(0, STEP, 2 * STEP)) 34 #elif VECTOR_SIZE == 4 35 #define STEP_VEC ((VEC_DATA_TYPE(DATA_TYPE, 4))(0, STEP, 2 * STEP, 3 * STEP)) 36 #elif VECTOR_SIZE == 8 37 #define STEP_VEC ((VEC_DATA_TYPE(DATA_TYPE, 8))(0, STEP, 2 * STEP, 3 * STEP, 4 * STEP, 5 * STEP, 6 * STEP, 7 * STEP)) 38 #elif VECTOR_SIZE == 16 39 #define STEP_VEC ((VEC_DATA_TYPE(DATA_TYPE, 16))(0, STEP, 2 * STEP, 3 * STEP, 4 * STEP, 5 * STEP, 6 * STEP, 7 * STEP, 8 * STEP, 9 * STEP, 10 * STEP, 11 * STEP, 12 * STEP, 13 * STEP, 14 * STEP, 15 * STEP)) 40 #endif // VECTOR_SIZE == 2 58 uint
id = max((
int)(get_global_id(0) * VECTOR_SIZE - (VECTOR_SIZE - VEC_SIZE_LEFTOVER) % VECTOR_SIZE), 0);
59 __global uchar *dst_ptr = out_ptr + out_offset_first_element_in_bytes +
id *
sizeof(
DATA_TYPE);
65 #else // VECTOR_SIZE == 1 68 seq0 = seq0 + STEP_VEC;
70 #endif //VECTOR_SIZE == 1 73 #else // !defined(OFFSET_OUT) && !defined(SCALE_OUT) 76 #define STEP_VEC ((VEC_DATA_TYPE(float, 2))(0, STEP)) 77 #elif VECTOR_SIZE == 3 78 #define STEP_VEC ((VEC_DATA_TYPE(float, 3))(0, STEP, 2 * STEP)) 79 #elif VECTOR_SIZE == 4 80 #define STEP_VEC ((VEC_DATA_TYPE(float, 4))(0, STEP, 2 * STEP, 3 * STEP)) 81 #elif VECTOR_SIZE == 8 82 #define STEP_VEC ((VEC_DATA_TYPE(float, 8))(0, STEP, 2 * STEP, 3 * STEP, 4 * STEP, 5 * STEP, 6 * STEP, 7 * STEP)) 83 #elif VECTOR_SIZE == 16 84 #define STEP_VEC ((VEC_DATA_TYPE(float, 16))(0, STEP, 2 * STEP, 3 * STEP, 4 * STEP, 5 * STEP, 6 * STEP, 7 * STEP, 8 * STEP, 9 * STEP, 10 * STEP, 11 * STEP, 12 * STEP, 13 * STEP, 14 * STEP, 15 * STEP)) 85 #endif // VECTOR_SIZE == 2 87 #define CONVERT_RTE(x, type) (convert_##type##_rte((x))) 88 #define CONVERT_DOWN(x, type) CONVERT_RTE(x, type) 104 __kernel
void range_quantized(
107 uint
id = max((
int)(get_global_id(0) * VECTOR_SIZE - (VECTOR_SIZE - VEC_SIZE_LEFTOVER) % VECTOR_SIZE), 0);
108 __global uchar *dst_ptr = out_ptr + out_offset_first_element_in_bytes +
id *
sizeof(
DATA_TYPE);
111 seq = (float)START + (
float)
id * (float)STEP;
112 seq = (
DATA_TYPE)(
int)(seq / ((float)SCALE_OUT) + (float)OFFSET_OUT);
113 seq = max(0.0f, min(seq, 255.0f));
115 #else // VECTOR_SIZE == 1 117 seq = (float)START +
id * (
float)STEP;
118 seq = seq + STEP_VEC;
119 seq = seq / ((
VEC_DATA_TYPE(
float, VECTOR_SIZE))((
float)SCALE_OUT)) + ((
VEC_DATA_TYPE(
float, VECTOR_SIZE))((
float)OFFSET_OUT));
124 #endif // VECTOR_SIZE == 1 126 #endif // !defined(OFFSET_OUT) && !defined(SCALE_OUT) 128 #endif // defined(VECTOR_SIZE) && defined(START) && defined(STEP) && defined(DATA_TYPE) && defined(VEC_SIZE_LEFTOVER) #define CONVERT_DOWN(x, type)
#define CONVERT_SAT(a, b)
#define VECTOR_DECLARATION(name)
SimpleTensor< T > range(SimpleTensor< T > &dst, float start, const size_t num_of_elements, float step)
#define STORE_VECTOR_SELECT(basename, data_type, ptr, vec_size, leftover, cond)
Store a vector that can only be partial in x.
#define VEC_DATA_TYPE(type, size)