diff --git a/modules/core/src/dxt.cpp b/modules/core/src/dxt.cpp index f168478afa..c39f11d4fb 100644 --- a/modules/core/src/dxt.cpp +++ b/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; diff --git a/modules/core/src/opencl/mulspectrums.cl b/modules/core/src/opencl/mulspectrums.cl new file mode 100644 index 0000000000..65f0edf6a1 --- /dev/null +++ b/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; + } +} diff --git a/modules/core/test/ocl/test_dft.cpp b/modules/core/test/ocl/test_dft.cpp index 5a596f3823..cc9b06d38c 100644 --- a/modules/core/test/ocl/test_dft.cpp +++ b/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),