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
142 uint x_offs = max((
int)(get_global_id(0) * VEC_SIZE_X - (VEC_SIZE_X - VEC_SIZE_LEFTOVER_X) % VEC_SIZE_X), 0);
143 uint y_offs = max((
int)(get_global_id(1) * VEC_SIZE_Y - (VEC_SIZE_Y - VEC_SIZE_LEFTOVER_Y) % VEC_SIZE_Y), 0);
144 uint z_offs = get_global_id(2);
147 __global uchar *src_addr = src_ptr + src_offset_first_element_in_bytes + x_offs * DATA_TYPE_IN_BYTES + y_offs * src_stride_y + z_offs * src_stride_z;
148 __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + y_offs * DATA_TYPE_IN_BYTES + x_offs * dst_stride_y + z_offs * dst_stride_z;
152 u0 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)src_addr);
155 u1 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + src_stride_y));
159 u2 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y));
163 u3 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y));
167 u4 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y));
169 u5 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 5 * src_stride_y));
171 u6 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 6 * src_stride_y));
173 u7 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 7 * src_stride_y));
177 u8 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 8 * src_stride_y));
179 u9 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 9 * src_stride_y));
181 u10 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 10 * src_stride_y));
183 u11 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 11 * src_stride_y));
185 u12 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 12 * src_stride_y));
187 u13 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 13 * src_stride_y));
189 u14 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 14 * src_stride_y));
191 u15 =
VLOAD(VEC_SIZE_X)(0, (__global DATA_TYPE *)(src_addr + 15 * src_stride_y));
196 t0 = TRANSPOSED_U(s0);
199 t1 = TRANSPOSED_U(s1);
203 t2 = TRANSPOSED_U(s2);
207 t3 = TRANSPOSED_U(s3);
211 t4 = TRANSPOSED_U(s4);
213 t5 = TRANSPOSED_U(s5);
215 t6 = TRANSPOSED_U(s6);
217 t7 = TRANSPOSED_U(s7);
221 t8 = TRANSPOSED_U(s8);
223 t9 = TRANSPOSED_U(s9);
225 tA = TRANSPOSED_U(sA);
227 tB = TRANSPOSED_U(sB);
229 tC = TRANSPOSED_U(sC);
231 tD = TRANSPOSED_U(sD);
233 tE = TRANSPOSED_U(sE);
235 tF = TRANSPOSED_U(sF);
240 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
241 && get_global_id(0) == 0,
242 VEC_SIZE_LEFTOVER_Y != 0 && get_global_id(1) == 0);
245 #endif // defined(DATA_TYPE_IN_BYTES) && defined(VEC_SIZE_X) && defined(VEC_SIZE_LEFTOVER_X) && defined(VEC_SIZE_Y) && defined(VEC_SIZE_LEFTOVER_Y)