Merge pull request #2032 from ilya-lavrenov:tapi_mulSpectrums

modules/core/src/dxt.cpp

@@ -42,6 +42,7 @@
 #include "precomp.hpp"
 #include "opencv2/core/opencl/runtime/opencl_clamdfft.hpp"
 #include "opencv2/core/opencl/runtime/opencl_core.hpp"
+#include "opencl_kernels.hpp"
 
 namespace cv
 {
@@ -2134,9 +2135,46 @@ void cv::idft( InputArray src, OutputArray dst, int flags, int nonzero_rows )
     dft( src, dst, flags | DFT_INVERSE, nonzero_rows );
 }
 
+namespace cv {
+
+static bool ocl_mulSpectrums( InputArray _srcA, InputArray _srcB,
+                              OutputArray _dst, int flags, bool conjB )
+{
+    int atype = _srcA.type(), btype = _srcB.type();
+    Size asize = _srcA.size(), bsize = _srcB.size();
+    CV_Assert(asize == bsize);
+
+    if ( !(atype == CV_32FC2 && btype == CV_32FC2) || flags != 0 )
+        return false;
+
+    UMat A = _srcA.getUMat(), B = _srcB.getUMat();
+    CV_Assert(A.size() == B.size());
+
+    _dst.create(A.size(), atype);
+    UMat dst = _dst.getUMat();
+
+    ocl::Kernel k("mulAndScaleSpectrums",
+                  ocl::core::mulspectrums_oclsrc,
+                  format("%s", conjB ? "-D CONJ" : ""));
+    if (k.empty())
+        return false;
+
+    k.args(ocl::KernelArg::ReadOnlyNoSize(A), ocl::KernelArg::ReadOnlyNoSize(B),
+           ocl::KernelArg::WriteOnly(dst));
+
+    size_t globalsize[2] = { asize.width, asize.height };
+    return k.run(2, globalsize, NULL, false);
+}
+
+}
+
 void cv::mulSpectrums( InputArray _srcA, InputArray _srcB,
                        OutputArray _dst, int flags, bool conjB )
 {
+    if (ocl::useOpenCL() && _dst.isUMat() &&
+            ocl_mulSpectrums(_srcA, _srcB, _dst, flags, conjB))
+        return;
+
     Mat srcA = _srcA.getMat(), srcB = _srcB.getMat();
     int depth = srcA.depth(), cn = srcA.channels(), type = srcA.type();
     int rows = srcA.rows, cols = srcA.cols;

modules/core/src/opencl/mulspectrums.cl

@@ -0,0 +1,81 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Peng Xiao, pengxiao@multicorewareinc.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other oclMaterials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the uintel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business uinterruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+inline float2 cmulf(float2 a, float2 b)
+{
+    return (float2)(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x);
+}
+
+inline float2 conjf(float2 a)
+{
+    return (float2)(a.x, - a.y);
+}
+
+__kernel void mulAndScaleSpectrums(__global const uchar * src1ptr, int src1_step, int src1_offset,
+                                   __global const uchar * src2ptr, int src2_step, int src2_offset,
+                                   __global uchar * dstptr, int dst_step, int dst_offset,
+                                   int dst_rows, int dst_cols)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < dst_cols && y < dst_rows)
+    {
+        int src1_index = mad24(y, src1_step, x * (int)sizeof(float2) + src1_offset);
+        int src2_index = mad24(y, src2_step, x * (int)sizeof(float2) + src2_offset);
+        int dst_index = mad24(y, dst_step, x * (int)sizeof(float2) + dst_offset);
+
+        float2 src0 = *(__global const float2 *)(src1ptr + src1_index);
+        float2 src1 = *(__global const float2 *)(src2ptr + src2_index);
+        __global float2 * dst = (__global float2 *)(dstptr + dst_index);
+
+#ifdef CONJ
+        float2 v = cmulf(src0, conjf(src1));
+#else
+        float2 v = cmulf(src0, src1);
+#endif
+        dst[0] = v;
+    }
+}

modules/core/test/ocl/test_dft.cpp

@@ -99,6 +99,57 @@ OCL_TEST_P(Dft, C2C)
     EXPECT_MAT_NEAR(dst, udst, eps);
 }
 
+////////////////////////////////////////////////////////////////////////////
+// MulSpectrums
+
+PARAM_TEST_CASE(MulSpectrums, bool, bool)
+{
+    bool ccorr, useRoi;
+
+    TEST_DECLARE_INPUT_PARAMETER(src1)
+    TEST_DECLARE_INPUT_PARAMETER(src2)
+    TEST_DECLARE_OUTPUT_PARAMETER(dst)
+
+    virtual void SetUp()
+    {
+        ccorr = GET_PARAM(0);
+        useRoi = GET_PARAM(1);
+    }
+
+    void generateTestData()
+    {
+        Size srcRoiSize = randomSize(1, MAX_VALUE);
+        Border src1Border = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src1, src1_roi, srcRoiSize, src1Border, CV_32FC2, -11, 11);
+
+
+        Border src2Border = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(src2, src2_roi, srcRoiSize, src2Border, CV_32FC2, -11, 11);
+
+        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
+        randomSubMat(dst, dst_roi, srcRoiSize, dstBorder, CV_32FC2, 5, 16);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(src1)
+        UMAT_UPLOAD_INPUT_PARAMETER(src2)
+        UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
+    }
+};
+
+OCL_TEST_P(MulSpectrums, Mat)
+{
+    for (int i = 0; i < test_loop_times; ++i)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::mulSpectrums(src1_roi, src2_roi, dst_roi, 0, ccorr));
+        OCL_ON(cv::mulSpectrums(usrc1_roi, usrc2_roi, udst_roi, 0, ccorr));
+
+        OCL_EXPECT_MATS_NEAR_RELATIVE(dst, 1e-6);
+    }
+}
+
+OCL_INSTANTIATE_TEST_CASE_P(OCL_ImgProc, MulSpectrums, testing::Combine(Bool(), Bool()));
+
 OCL_INSTANTIATE_TEST_CASE_P(Core, Dft, Combine(Values(cv::Size(2, 3), cv::Size(5, 4), cv::Size(25, 20),
                                                        cv::Size(512, 1), cv::Size(1024, 768)),
                                                Values(CV_32F, CV_64F),