26 #if defined(VEC_SIZE) && defined(DATA_TYPE)
50 __kernel
void fft_digit_reverse_axis_0(
60 const unsigned int iidx = *((__global uint *)(idx.
ptr));
64 DATA_TYPE data = *((__global DATA_TYPE *)
tensor3D_offset(&
src, iidx, get_global_id(1), get_global_id(2)));
67 data = vload2(0, (__global DATA_TYPE *)
tensor3D_offset(&
src, iidx, get_global_id(1), get_global_id(2)));
68 #else // VEC_SIZE == 1
69 #error "vec_size of 1 and 2 are supported"
70 #endif // VEC_SIZE == 1
79 #else // VEC_SIZE == 1
80 #error "vec_size of 1 and 2 are supported"
81 #endif // VEC_SIZE == 1
85 vstore2((
VEC_DATA_TYPE(DATA_TYPE, 2))(res.s0, -res.s1), 0, (__global DATA_TYPE *)
dst.ptr);
86 #else // defined(CONJ)
87 vstore2(res, 0, (__global DATA_TYPE *)
dst.ptr);
88 #endif // defined(CONJ)
114 __kernel
void fft_digit_reverse_axis_1(
124 const unsigned int iidx = *((__global uint *)
vector_offset(&idx, (
int)(get_global_id(1))));
128 DATA_TYPE data = *((__global DATA_TYPE *)
tensor3D_offset(&
src, get_global_id(0), iidx, get_global_id(2)));
131 data = vload2(0, (__global DATA_TYPE *)
tensor3D_offset(&
src, get_global_id(0), iidx, get_global_id(2)));
132 #else // VEC_SIZE == 1
133 #error "vec_size of 1 and 2 are supported"
134 #endif // VEC_SIZE == 1
143 #else // VEC_SIZE == 1
144 #error "vec_size of 1 and 2 are supported"
145 #endif // VEC_SIZE == 1
149 vstore2((
VEC_DATA_TYPE(DATA_TYPE, 2))(res.s0, -res.s1), 0, (__global DATA_TYPE *)
dst.ptr);
150 #else // defined(CONJ)
151 vstore2(res, 0, (__global DATA_TYPE *)
dst.ptr);
152 #endif // defined(CONJ)
154 #endif // defined(VEC_SIZE) && defined(DATA_TYPE)