Added Elena's changes with implemented DFT_INVERSE C2C mode.

modules/core/perf/opencl/perf_dxt.cpp

@@ -67,7 +67,7 @@ typedef TestBaseWithParam<DftParams> DftFixture;
 
 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();

modules/core/src/dxt.cpp

@@ -2129,8 +2129,8 @@ struct OCL_FftPlan
 
                 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);
                 }
             }        
         }
@@ -2152,13 +2152,14 @@ struct OCL_FftPlan
         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";
         }
@@ -2166,7 +2167,7 @@ struct OCL_FftPlan
         {
             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";
         }
@@ -2269,15 +2270,12 @@ static bool ocl_dft(InputArray _src, OutputArray _dst, int flags, int nonzero_ro
 
     // if output format is not specified
     if (complex_output + real_output == 0)
-    {
-        if (!inv)
     {
         if (real_input)
             real_output = 1;
         else
             complex_output = 1;
     }
-    }
 
     // Forward Complex to CCS not supported
     if (complex_input && real_output && !inv)
@@ -2294,23 +2292,7 @@ static bool ocl_dft(InputArray _src, OutputArray _dst, int flags, int nonzero_ro
         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); 
@@ -2330,7 +2312,7 @@ static bool ocl_dft(InputArray _src, OutputArray _dst, int flags, int nonzero_ro
         }
     }
 
-    if (!ocl_dft_C2C_rows(input, output, nonzero_rows, flags))
+    if (!ocl_dft_C2C_rows(src, output, nonzero_rows, flags))
         return false;
 
     if (!is1d)

modules/core/src/opencl/fft.cl

@@ -425,3 +425,116 @@ __kernel void fft_multi_radix_cols(__global const uchar* src_ptr, int src_step,
 #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

modules/core/test/ocl/test_dft.cpp

@@ -62,7 +62,7 @@ namespace ocl {
 ////////////////////////////////////////////////////////////////////////////
 // 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;
@@ -91,9 +91,9 @@ PARAM_TEST_CASE(Dft, cv::Size, OCL_FFT_TYPE, bool, bool, bool)
             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;
@@ -190,9 +190,10 @@ OCL_INSTANTIATE_TEST_CASE_P(OCL_ImgProc, MulSpectrums, testing::Combine(Bool(),
 
 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
                                                )
                             );