OCL_PERF_TEST_P(DftFixture, Dft, ::testing::Combine(Values(C2C, R2R, C2R, R2C),
Values(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3, Size(1024, 1024), Size(512, 512), Size(2048, 2048)),
- Values((int)DFT_ROWS, (int) 0/*, (int)DFT_SCALE, (int)DFT_INVERSE,
+ Values((int)DFT_ROWS, (int) 0, (int)DFT_SCALE/*, (int)DFT_INVERSE,
(int)DFT_INVERSE | DFT_SCALE, (int)DFT_ROWS | DFT_INVERSE*/)))
{
const DftParams params = GetParam();
for (int k=0; k<(n/radix); k++)
{
- ptr[ptr_index++] = cos(k*theta);
- ptr[ptr_index++] = sin(k*theta);
+ ptr[ptr_index++] = (float) cos(k*theta);
+ ptr[ptr_index++] = (float) sin(k*theta);
}
}
}
String kernel_name;
bool is1d = (flags & DFT_ROWS) != 0 || dft_size == 1;
+ bool inv = (flags & DFT_INVERSE) != 0;
String options = buildOptions;
if (rows)
{
globalsize[0] = thread_count; globalsize[1] = dft_size;
localsize[0] = thread_count; localsize[1] = 1;
- kernel_name = "fft_multi_radix_rows";
+ kernel_name = !inv ? "fft_multi_radix_rows" : "ifft_multi_radix_rows";
if (is1d && (flags & DFT_SCALE))
options += " -D DFT_SCALE";
}
{
globalsize[0] = dft_size; globalsize[1] = thread_count;
localsize[0] = 1; localsize[1] = thread_count;
- kernel_name = "fft_multi_radix_cols";
+ kernel_name = !inv ? "fft_multi_radix_cols" : "ifft_multi_radix_cols";
if (flags & DFT_SCALE)
options += " -D DFT_SCALE";
}
// if output format is not specified
if (complex_output + real_output == 0)
{
- if (!inv)
- {
- if (real_input)
- real_output = 1;
- else
- complex_output = 1;
- }
+ if (real_input)
+ real_output = 1;
+ else
+ complex_output = 1;
}
// Forward Complex to CCS not supported
real_output = 1;
}
- UMat input, output;
- if (complex_input)
- {
- input = src;
- }
- else
- {
- if (!inv)
- {
- input = src;
- }
- else
- {
- // TODO: unpack from CCS format
- }
- }
-
+ UMat output;
if (complex_output)
{
_dst.create(src.size(), CV_32FC2);
}
}
- if (!ocl_dft_C2C_rows(input, output, nonzero_rows, flags))
+ if (!ocl_dft_C2C_rows(src, output, nonzero_rows, flags))
return false;
if (!is1d)
}
#endif
}
+}
+
+__kernel void ifft_multi_radix_rows(__global const uchar* src_ptr, int src_step, int src_offset, int src_rows, int src_cols,
+ __global uchar* dst_ptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ __constant float2 * twiddles_ptr, const int t, const int nz)
+{
+ const int x = get_global_id(0);
+ const int y = get_group_id(1);
+
+ if (y < nz)
+ {
+ __local float2 smem[LOCAL_SIZE];
+ __constant const float2* twiddles = (__constant float2*) twiddles_ptr;
+ const int ind = x;
+ const int block_size = LOCAL_SIZE/kercn;
+#ifdef IS_1D
+ float scale = 1.f/dst_cols;
+#else
+ float scale = 1.f/(dst_cols*dst_rows);
+#endif
+
+#ifndef REAL_INPUT
+ __global const float2* src = (__global const float2*)(src_ptr + mad24(y, src_step, mad24(x, (int)(sizeof(float)*2), src_offset)));
+ #pragma unroll
+ for (int i=0; i<kercn; i++)
+ {
+ smem[x+i*block_size].x = src[i*block_size].x;
+ smem[x+i*block_size].y = -src[i*block_size].y;
+ }
+#else
+ __global const float2* src = (__global const float2*)(src_ptr + mad24(y, src_step, mad24(1, (int)sizeof(float), src_offset)));
+ #pragma unroll
+ for (int i=x; i<(LOCAL_SIZE-1)/2; i+=block_size)
+ {
+ smem[i+1].x = src[i].x;
+ smem[i+1].y = -src[i].y;
+ smem[LOCAL_SIZE-i-1] = src[i];
+ }
+ if (x==0)
+ {
+ smem[0].x = *(__global const float*)(src_ptr + mad24(y, src_step, src_offset));
+ smem[0].y = 0.f;
+
+ if(LOCAL_SIZE % 2 ==0)
+ {
+ smem[LOCAL_SIZE/2].x = src[LOCAL_SIZE/2-1].x;
+ smem[LOCAL_SIZE/2].y = 0.f;
+ }
+ }
+#endif
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ RADIX_PROCESS;
+
+ // copy data to dst
+#ifndef REAL_INPUT
+ __global float2* dst = (__global float*)(dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)*2), dst_offset)));
+ #pragma unroll
+ for (int i=0; i<kercn; i++)
+ {
+ dst[i*block_size].x = VAL(smem[x + i*block_size].x, scale);
+ dst[i*block_size].y = VAL(-smem[x + i*block_size].y, scale);
+ }
+#else
+ __global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)), dst_offset)));
+ #pragma unroll
+ for (int i=0; i<kercn; i++)
+ {
+ dst[i*block_size] = smem[x + i*block_size].x;
+ }
+#endif
+ }
+}
+
+__kernel void ifft_multi_radix_cols(__global const uchar* src_ptr, int src_step, int src_offset, int src_rows, int src_cols,
+ __global uchar* dst_ptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ __constant float2 * twiddles_ptr, const int t, const int nz)
+{
+ const int x = get_group_id(0);
+ const int y = get_global_id(1);
+
+ if (x < nz)
+ {
+ __local float2 smem[LOCAL_SIZE];
+ __global const uchar* src = src_ptr + mad24(y, src_step, mad24(x, (int)(sizeof(float)*2), src_offset));
+ __global uchar* dst = dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)*2), dst_offset));
+ __constant const float2* twiddles = (__constant float2*) twiddles_ptr;
+ const int ind = y;
+ const int block_size = LOCAL_SIZE/kercn;
+ float scale = 1.f/(dst_rows*dst_cols);
+
+ #pragma unroll
+ for (int i=0; i<kercn; i++)
+ {
+ float2 temp = *((__global const float2*)(src + i*block_size*src_step));
+ smem[y+i*block_size].x = temp.x;
+ smem[y+i*block_size].y = -temp.y;
+ }
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ RADIX_PROCESS;
+
+ // copy data to dst
+ #pragma unroll
+ for (int i=0; i<kercn; i++)
+ {
+ __global float2* rez = (__global float2*)(dst + i*block_size*src_step);
+ rez[0].x = VAL(smem[y + i*block_size].x, scale);
+ rez[0].y = VAL(-smem[y + i*block_size].y, scale);
+ }
+ }
}
\ No newline at end of file
////////////////////////////////////////////////////////////////////////////
// Dft
-PARAM_TEST_CASE(Dft, cv::Size, OCL_FFT_TYPE, bool, bool, bool)
+PARAM_TEST_CASE(Dft, cv::Size, OCL_FFT_TYPE, bool, bool, bool, bool)
{
cv::Size dft_size;
int dft_flags, depth, cn, dft_type;
dft_flags |= cv::DFT_ROWS;
if (GET_PARAM(3))
dft_flags |= cv::DFT_SCALE;
- //if (GET_PARAM(4))
- // dft_flags |= cv::DFT_INVERSE;
- inplace = GET_PARAM(4);
+ if (GET_PARAM(4))
+ dft_flags |= cv::DFT_INVERSE;
+ inplace = GET_PARAM(5);
is1d = (dft_flags & DFT_ROWS) != 0 || dft_size.height == 1;
OCL_INSTANTIATE_TEST_CASE_P(Core, Dft, Combine(Values(cv::Size(6, 4), cv::Size(5, 8), cv::Size(6, 6),
cv::Size(512, 1), cv::Size(1280, 768)),
- Values((OCL_FFT_TYPE) R2C, (OCL_FFT_TYPE) C2C, (OCL_FFT_TYPE) R2R, (OCL_FFT_TYPE) C2R),
+ Values(/*(OCL_FFT_TYPE) R2C, */(OCL_FFT_TYPE) C2C/*, (OCL_FFT_TYPE) R2R, (OCL_FFT_TYPE) C2R*/),
Bool(), // DFT_ROWS
Bool(), // DFT_SCALE
+ Bool(), // DFT_INVERSE
Bool() // inplace
)
);