24 #define PARTIAL_STORE_M0 VEC_SIZE_LEFTOVER_X 25 #define PARTIAL_STORE_N0 VEC_SIZE_LEFTOVER_Y 30 #if defined(DATA_TYPE_IN_BYTES) && defined(VEC_SIZE_X) && defined(VEC_SIZE_LEFTOVER_X) && defined(VEC_SIZE_Y) && defined(VEC_SIZE_LEFTOVER_Y) 34 #define TRANSPOSED_U(val) \ 39 #define TRANSPOSED_U(val) \ 44 #define TRANSPOSED_U(val) \ 49 #define TRANSPOSED_U(val) \ 54 #define TRANSPOSED_U(val) \ 56 u0, u1, u2, u3, u4, u5, u6, u7 \ 58 #elif VEC_SIZE_Y == 16 59 #define TRANSPOSED_U(val) \ 61 u0, u1, u2, u3, u4, u5, u6, u7, \ 62 u8, u9, u10, u11, u12, u13, u14, u15 \ 65 #else // VEC_SIZE_X == 1 67 #define TRANSPOSED_U(val) \ 72 #define TRANSPOSED_U(val) \ 77 #define TRANSPOSED_U(val) \ 79 u0.val, u1.val, u2.val \ 82 #define TRANSPOSED_U(val) \ 84 u0.val, u1.val, u2.val, u3.val \ 87 #define TRANSPOSED_U(val) \ 89 u0.val, u1.val, u2.val, u3.val, u4.val, u5.val, u6.val, u7.val \ 91 #elif VEC_SIZE_Y == 16 92 #define TRANSPOSED_U(val) \ 94 u0.val, u1.val, u2.val, u3.val, u4.val, u5.val, u6.val, u7.val, \ 95 u8.val, u9.val, u10.val, u11.val, u12.val, u13.val, u14.val, u15.val \ 98 #endif // VEC_SIZE_X == 1 100 #if DATA_TYPE_IN_BYTES == 4 101 #define DATA_TYPE uint 102 #elif DATA_TYPE_IN_BYTES == 2 103 #define DATA_TYPE ushort 104 #elif DATA_TYPE_IN_BYTES == 1 105 #define DATA_TYPE uchar 107 #error DATA_TYPE_IN_BYTES not supported for transpose 138 uint x_offs = max((
int)(get_global_id(0) * VEC_SIZE_X - (VEC_SIZE_X - VEC_SIZE_LEFTOVER_X) % VEC_SIZE_X), 0);
139 uint y_offs = max((
int)(get_global_id(1) * VEC_SIZE_Y - (VEC_SIZE_Y - VEC_SIZE_LEFTOVER_Y) % VEC_SIZE_Y), 0);
142 __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x_offs * DATA_TYPE_IN_BYTES + y_offs * src_stride_y;
143 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + y_offs * DATA_TYPE_IN_BYTES + x_offs * dst_stride_y;
150 u1 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + src_stride_y));
154 u2 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 2 * src_stride_y));
158 u3 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 3 * src_stride_y));
162 u4 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 4 * src_stride_y));
164 u5 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 5 * src_stride_y));
166 u6 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 6 * src_stride_y));
168 u7 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 7 * src_stride_y));
172 u8 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 8 * src_stride_y));
174 u9 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 9 * src_stride_y));
176 u10 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 10 * src_stride_y));
178 u11 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 11 * src_stride_y));
180 u12 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 12 * src_stride_y));
182 u13 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 13 * src_stride_y));
184 u14 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 14 * src_stride_y));
186 u15 =
VLOAD(VEC_SIZE_X)(0, (__global
DATA_TYPE *)(src_addr + 15 * src_stride_y));
191 t0 = TRANSPOSED_U(s0);
194 t1 = TRANSPOSED_U(s1);
198 t2 = TRANSPOSED_U(s2);
202 t3 = TRANSPOSED_U(s3);
206 t4 = TRANSPOSED_U(s4);
208 t5 = TRANSPOSED_U(s5);
210 t6 = TRANSPOSED_U(s6);
212 t7 = TRANSPOSED_U(s7);
216 t8 = TRANSPOSED_U(s8);
218 t9 = TRANSPOSED_U(s9);
220 tA = TRANSPOSED_U(sA);
222 tB = TRANSPOSED_U(sB);
224 tC = TRANSPOSED_U(sC);
226 tD = TRANSPOSED_U(sD);
228 tE = TRANSPOSED_U(sE);
230 tF = TRANSPOSED_U(sF);
235 STORE_BLOCK_BOUNDARY_AWARE(VEC_SIZE_X, VEC_SIZE_Y,
DATA_TYPE,
t, (__global uchar *)dst_addr, dst_stride_y, zout, VEC_SIZE_LEFTOVER_X, VEC_SIZE_LEFTOVER_Y, VEC_SIZE_LEFTOVER_X != 0
236 && get_global_id(0) == 0,
237 VEC_SIZE_LEFTOVER_Y != 0 && get_global_id(1) == 0);
240 #endif // defined(DATA_TYPE_IN_BYTES) && defined(VEC_SIZE_X) && defined(VEC_SIZE_LEFTOVER_X) && defined(VEC_SIZE_Y) && defined(VEC_SIZE_LEFTOVER_Y) #define REPEAT_VAR_INIT_TO_CONST(N, TYPE, VAR, VAL)
#define IMAGE_DECLARATION(name)
SimpleTensor< float > src
SimpleTensor< T > transpose(const SimpleTensor< T > &src)
#define VEC_DATA_TYPE(type, size)