Merge pull request #12585 from alalek:move_cuda_modules

cuda: move CUDA modules to opencv_contrib: https://github.com/opencv/opencv_contrib/pull/1781

modules/cudaarithm/CMakeLists.txt

-if(IOS OR WINRT OR (NOT HAVE_CUDA AND NOT BUILD_CUDA_STUBS))
-  ocv_module_disable(cudaarithm)
-endif()
-
-set(the_description "CUDA-accelerated Operations on Matrices")
-
-ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations -Wshadow)
-
-ocv_add_module(cudaarithm opencv_core OPTIONAL opencv_cudev WRAP python)
-
-ocv_module_include_directories()
-ocv_glob_module_sources()
-
-set(extra_libs "")
-
-if(HAVE_CUBLAS)
-  list(APPEND extra_libs ${CUDA_cublas_LIBRARY})
-endif()
-
-if(HAVE_CUFFT)
-  list(APPEND extra_libs ${CUDA_cufft_LIBRARY})
-endif()
-
-ocv_create_module(${extra_libs})
-
-ocv_add_accuracy_tests(DEPENDS_ON opencv_imgproc)
-ocv_add_perf_tests(DEPENDS_ON opencv_imgproc)

modules/cudaarithm/perf/perf_arithm.cpp

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-
-#include "perf_precomp.hpp"
-
-namespace opencv_test { namespace {
-
-//////////////////////////////////////////////////////////////////////
-// GEMM
-
-#ifdef HAVE_CUBLAS
-
-CV_FLAGS(GemmFlags, 0, cv::GEMM_1_T, cv::GEMM_2_T, cv::GEMM_3_T)
-#define ALL_GEMM_FLAGS Values(GemmFlags(0), GemmFlags(cv::GEMM_1_T), GemmFlags(cv::GEMM_2_T), GemmFlags(cv::GEMM_3_T), \
-                              GemmFlags(cv::GEMM_1_T | cv::GEMM_2_T), GemmFlags(cv::GEMM_1_T | cv::GEMM_3_T), GemmFlags(cv::GEMM_1_T | cv::GEMM_2_T | cv::GEMM_3_T))
-
-DEF_PARAM_TEST(Sz_Type_Flags, cv::Size, MatType, GemmFlags);
-
-PERF_TEST_P(Sz_Type_Flags, GEMM,
-            Combine(Values(cv::Size(512, 512), cv::Size(1024, 1024)),
-                    Values(CV_32FC1, CV_32FC2, CV_64FC1),
-                    ALL_GEMM_FLAGS))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int type = GET_PARAM(1);
-    const int flags = GET_PARAM(2);
-
-    cv::Mat src1(size, type);
-    declare.in(src1, WARMUP_RNG);
-
-    cv::Mat src2(size, type);
-    declare.in(src2, WARMUP_RNG);
-
-    cv::Mat src3(size, type);
-    declare.in(src3, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        declare.time(5.0);
-
-        const cv::cuda::GpuMat d_src1(src1);
-        const cv::cuda::GpuMat d_src2(src2);
-        const cv::cuda::GpuMat d_src3(src3);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() cv::cuda::gemm(d_src1, d_src2, 1.0, d_src3, 1.0, dst, flags);
-
-        CUDA_SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE);
-    }
-    else
-    {
-        declare.time(50.0);
-
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::gemm(src1, src2, 1.0, src3, 1.0, dst, flags);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-#endif
-
-//////////////////////////////////////////////////////////////////////
-// MulSpectrums
-
-CV_FLAGS(DftFlags, 0, cv::DFT_INVERSE, cv::DFT_SCALE, cv::DFT_ROWS, cv::DFT_COMPLEX_OUTPUT, cv::DFT_REAL_OUTPUT)
-
-DEF_PARAM_TEST(Sz_Flags, cv::Size, DftFlags);
-
-PERF_TEST_P(Sz_Flags, MulSpectrums,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    Values(0, DftFlags(cv::DFT_ROWS))))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int flag = GET_PARAM(1);
-
-    cv::Mat a(size, CV_32FC2);
-    cv::Mat b(size, CV_32FC2);
-    declare.in(a, b, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        const cv::cuda::GpuMat d_a(a);
-        const cv::cuda::GpuMat d_b(b);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() cv::cuda::mulSpectrums(d_a, d_b, dst, flag);
-
-        CUDA_SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::mulSpectrums(a, b, dst, flag);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// MulAndScaleSpectrums
-
-PERF_TEST_P(Sz, MulAndScaleSpectrums,
-            CUDA_TYPICAL_MAT_SIZES)
-{
-    const cv::Size size = GetParam();
-
-    const float scale = 1.f / size.area();
-
-    cv::Mat src1(size, CV_32FC2);
-    cv::Mat src2(size, CV_32FC2);
-    declare.in(src1,src2, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        const cv::cuda::GpuMat d_src1(src1);
-        const cv::cuda::GpuMat d_src2(src2);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() cv::cuda::mulAndScaleSpectrums(d_src1, d_src2, dst, cv::DFT_ROWS, scale, false);
-
-        CUDA_SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE);
-    }
-    else
-    {
-        FAIL_NO_CPU();
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// Dft
-
-PERF_TEST_P(Sz_Flags, Dft,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    Values(0, DftFlags(cv::DFT_ROWS), DftFlags(cv::DFT_INVERSE))))
-{
-    declare.time(10.0);
-
-    const cv::Size size = GET_PARAM(0);
-    const int flag = GET_PARAM(1);
-
-    cv::Mat src(size, CV_32FC2);
-    declare.in(src, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        const cv::cuda::GpuMat d_src(src);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() cv::cuda::dft(d_src, dst, size, flag);
-
-        CUDA_SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::dft(src, dst, flag);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// Convolve
-
-DEF_PARAM_TEST(Sz_KernelSz_Ccorr, cv::Size, int, bool);
-
-PERF_TEST_P(Sz_KernelSz_Ccorr, Convolve,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    Values(17, 27, 32, 64),
-                    Bool()))
-{
-    declare.time(10.0);
-
-    const cv::Size size = GET_PARAM(0);
-    const int templ_size = GET_PARAM(1);
-    const bool ccorr = GET_PARAM(2);
-
-    const cv::Mat image(size, CV_32FC1);
-    const cv::Mat templ(templ_size, templ_size, CV_32FC1);
-    declare.in(image, templ, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        cv::cuda::GpuMat d_image = cv::cuda::createContinuous(size, CV_32FC1);
-        d_image.upload(image);
-
-        cv::cuda::GpuMat d_templ = cv::cuda::createContinuous(templ_size, templ_size, CV_32FC1);
-        d_templ.upload(templ);
-
-        cv::Ptr<cv::cuda::Convolution> convolution = cv::cuda::createConvolution();
-
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() convolution->convolve(d_image, d_templ, dst, ccorr);
-
-        CUDA_SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE);
-    }
-    else
-    {
-        if (ccorr)
-            FAIL_NO_CPU();
-
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::filter2D(image, dst, image.depth(), templ);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-}} // namespace

modules/cudaarithm/perf/perf_core.cpp

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-
-#include "perf_precomp.hpp"
-
-namespace opencv_test { namespace {
-
-#define ARITHM_MAT_DEPTH Values(CV_8U, CV_16U, CV_32F, CV_64F)
-
-//////////////////////////////////////////////////////////////////////
-// Merge
-
-DEF_PARAM_TEST(Sz_Depth_Cn, cv::Size, MatDepth, MatCn);
-
-PERF_TEST_P(Sz_Depth_Cn, Merge,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    ARITHM_MAT_DEPTH,
-                    Values(2, 3, 4)))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int depth = GET_PARAM(1);
-    const int channels = GET_PARAM(2);
-
-    std::vector<cv::Mat> src(channels);
-    for (int i = 0; i < channels; ++i)
-    {
-        src[i].create(size, depth);
-        declare.in(src[i], WARMUP_RNG);
-    }
-
-    if (PERF_RUN_CUDA())
-    {
-        std::vector<cv::cuda::GpuMat> d_src(channels);
-        for (int i = 0; i < channels; ++i)
-            d_src[i].upload(src[i]);
-
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() cv::cuda::merge(d_src, dst);
-
-        CUDA_SANITY_CHECK(dst, 1e-10);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::merge(src, dst);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// Split
-
-PERF_TEST_P(Sz_Depth_Cn, Split,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    ARITHM_MAT_DEPTH,
-                    Values(2, 3, 4)))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int depth = GET_PARAM(1);
-    const int channels = GET_PARAM(2);
-
-    cv::Mat src(size, CV_MAKE_TYPE(depth, channels));
-    declare.in(src, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        const cv::cuda::GpuMat d_src(src);
-        std::vector<cv::cuda::GpuMat> dst;
-
-        TEST_CYCLE() cv::cuda::split(d_src, dst);
-
-        const cv::cuda::GpuMat& dst0 = dst[0];
-        const cv::cuda::GpuMat& dst1 = dst[1];
-
-        CUDA_SANITY_CHECK(dst0, 1e-10);
-        CUDA_SANITY_CHECK(dst1, 1e-10);
-    }
-    else
-    {
-        std::vector<cv::Mat> dst;
-
-        TEST_CYCLE() cv::split(src, dst);
-
-        const cv::Mat& dst0 = dst[0];
-        const cv::Mat& dst1 = dst[1];
-
-        CPU_SANITY_CHECK(dst0);
-        CPU_SANITY_CHECK(dst1);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// Transpose
-
-PERF_TEST_P(Sz_Type, Transpose,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    Values(CV_8UC1, CV_8UC4, CV_16UC2, CV_16SC2, CV_32SC1, CV_32SC2, CV_64FC1)))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int type = GET_PARAM(1);
-
-    cv::Mat src(size, type);
-    declare.in(src, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        const cv::cuda::GpuMat d_src(src);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() cv::cuda::transpose(d_src, dst);
-
-        CUDA_SANITY_CHECK(dst, 1e-10);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::transpose(src, dst);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// Flip
-
-enum {FLIP_BOTH = 0, FLIP_X = 1, FLIP_Y = -1};
-CV_ENUM(FlipCode, FLIP_BOTH, FLIP_X, FLIP_Y)
-
-DEF_PARAM_TEST(Sz_Depth_Cn_Code, cv::Size, MatDepth, MatCn, FlipCode);
-
-PERF_TEST_P(Sz_Depth_Cn_Code, Flip,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    Values(CV_8U, CV_16U, CV_32F),
-                    CUDA_CHANNELS_1_3_4,
-                    FlipCode::all()))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int depth = GET_PARAM(1);
-    const int channels = GET_PARAM(2);
-    const int flipCode = GET_PARAM(3);
-
-    const int type = CV_MAKE_TYPE(depth, channels);
-
-    cv::Mat src(size, type);
-    declare.in(src, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        const cv::cuda::GpuMat d_src(src);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() cv::cuda::flip(d_src, dst, flipCode);
-
-        CUDA_SANITY_CHECK(dst);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::flip(src, dst, flipCode);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// LutOneChannel
-
-PERF_TEST_P(Sz_Type, LutOneChannel,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    Values(CV_8UC1, CV_8UC3)))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int type = GET_PARAM(1);
-
-    cv::Mat src(size, type);
-    declare.in(src, WARMUP_RNG);
-
-    cv::Mat lut(1, 256, CV_8UC1);
-    declare.in(lut, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        cv::Ptr<cv::cuda::LookUpTable> lutAlg = cv::cuda::createLookUpTable(lut);
-
-        const cv::cuda::GpuMat d_src(src);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() lutAlg->transform(d_src, dst);
-
-        CUDA_SANITY_CHECK(dst);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::LUT(src, lut, dst);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// LutMultiChannel
-
-PERF_TEST_P(Sz_Type, LutMultiChannel,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    Values<MatType>(CV_8UC3)))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int type = GET_PARAM(1);
-
-    cv::Mat src(size, type);
-    declare.in(src, WARMUP_RNG);
-
-    cv::Mat lut(1, 256, CV_MAKE_TYPE(CV_8U, src.channels()));
-    declare.in(lut, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        cv::Ptr<cv::cuda::LookUpTable> lutAlg = cv::cuda::createLookUpTable(lut);
-
-        const cv::cuda::GpuMat d_src(src);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() lutAlg->transform(d_src, dst);
-
-        CUDA_SANITY_CHECK(dst);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::LUT(src, lut, dst);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// CopyMakeBorder
-
-DEF_PARAM_TEST(Sz_Depth_Cn_Border, cv::Size, MatDepth, MatCn, BorderMode);
-
-PERF_TEST_P(Sz_Depth_Cn_Border, CopyMakeBorder,
-            Combine(CUDA_TYPICAL_MAT_SIZES,
-                    Values(CV_8U, CV_16U, CV_32F),
-                    CUDA_CHANNELS_1_3_4,
-                    ALL_BORDER_MODES))
-{
-    const cv::Size size = GET_PARAM(0);
-    const int depth = GET_PARAM(1);
-    const int channels = GET_PARAM(2);
-    const int borderMode = GET_PARAM(3);
-
-    const int type = CV_MAKE_TYPE(depth, channels);
-
-    cv::Mat src(size, type);
-    declare.in(src, WARMUP_RNG);
-
-    if (PERF_RUN_CUDA())
-    {
-        const cv::cuda::GpuMat d_src(src);
-        cv::cuda::GpuMat dst;
-
-        TEST_CYCLE() cv::cuda::copyMakeBorder(d_src, dst, 5, 5, 5, 5, borderMode);
-
-        CUDA_SANITY_CHECK(dst);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::copyMakeBorder(src, dst, 5, 5, 5, 5, borderMode);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-}} // namespace

modules/cudaarithm/perf/perf_main.cpp

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-
-#include "perf_precomp.hpp"
-
-using namespace perf;
-
-CV_PERF_TEST_CUDA_MAIN(cudaarithm)

modules/cudaarithm/perf/perf_precomp.hpp

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-#ifndef __OPENCV_PERF_PRECOMP_HPP__
-#define __OPENCV_PERF_PRECOMP_HPP__
-
-#include "opencv2/ts.hpp"
-#include "opencv2/ts/cuda_perf.hpp"
-
-#include "opencv2/cudaarithm.hpp"
-
-namespace opencv_test {
-using namespace perf;
-using namespace testing;
-}
-
-#endif

modules/cudaarithm/src/core.cpp

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-
-#include "precomp.hpp"
-
-using namespace cv;
-using namespace cv::cuda;
-
-#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
-
-void cv::cuda::merge(const GpuMat*, size_t, OutputArray, Stream&) { throw_no_cuda(); }
-void cv::cuda::merge(const std::vector<GpuMat>&, OutputArray, Stream&) { throw_no_cuda(); }
-
-void cv::cuda::split(InputArray, GpuMat*, Stream&) { throw_no_cuda(); }
-void cv::cuda::split(InputArray, std::vector<GpuMat>&, Stream&) { throw_no_cuda(); }
-
-void cv::cuda::transpose(InputArray, OutputArray, Stream&) { throw_no_cuda(); }
-
-void cv::cuda::flip(InputArray, OutputArray, int, Stream&) { throw_no_cuda(); }
-
-Ptr<LookUpTable> cv::cuda::createLookUpTable(InputArray) { throw_no_cuda(); return Ptr<LookUpTable>(); }
-
-void cv::cuda::copyMakeBorder(InputArray, OutputArray, int, int, int, int, int, Scalar, Stream&) { throw_no_cuda(); }
-
-#else /* !defined (HAVE_CUDA) */
-
-////////////////////////////////////////////////////////////////////////
-// flip
-
-namespace
-{
-    template<int DEPTH> struct NppTypeTraits;
-    template<> struct NppTypeTraits<CV_8U>  { typedef Npp8u npp_t; };
-    template<> struct NppTypeTraits<CV_8S>  { typedef Npp8s npp_t; };
-    template<> struct NppTypeTraits<CV_16U> { typedef Npp16u npp_t; };
-    template<> struct NppTypeTraits<CV_16S> { typedef Npp16s npp_t; };
-    template<> struct NppTypeTraits<CV_32S> { typedef Npp32s npp_t; };
-    template<> struct NppTypeTraits<CV_32F> { typedef Npp32f npp_t; };
-    template<> struct NppTypeTraits<CV_64F> { typedef Npp64f npp_t; };
-
-    template <int DEPTH> struct NppMirrorFunc
-    {
-        typedef typename NppTypeTraits<DEPTH>::npp_t npp_t;
-
-        typedef NppStatus (*func_t)(const npp_t* pSrc, int nSrcStep, npp_t* pDst, int nDstStep, NppiSize oROI, NppiAxis flip);
-    };
-
-    template <int DEPTH, typename NppMirrorFunc<DEPTH>::func_t func> struct NppMirror
-    {
-        typedef typename NppMirrorFunc<DEPTH>::npp_t npp_t;
-
-        static void call(const GpuMat& src, GpuMat& dst, int flipCode, cudaStream_t stream)
-        {
-            NppStreamHandler h(stream);
-
-            NppiSize sz;
-            sz.width  = src.cols;
-            sz.height = src.rows;
-
-            nppSafeCall( func(src.ptr<npp_t>(), static_cast<int>(src.step),
-                dst.ptr<npp_t>(), static_cast<int>(dst.step), sz,
-                (flipCode == 0 ? NPP_HORIZONTAL_AXIS : (flipCode > 0 ? NPP_VERTICAL_AXIS : NPP_BOTH_AXIS))) );
-
-            if (stream == 0)
-                cudaSafeCall( cudaDeviceSynchronize() );
-        }
-    };
-}
-
-void cv::cuda::flip(InputArray _src, OutputArray _dst, int flipCode, Stream& stream)
-{
-    typedef void (*func_t)(const GpuMat& src, GpuMat& dst, int flipCode, cudaStream_t stream);
-    static const func_t funcs[6][4] =
-    {
-        {NppMirror<CV_8U, nppiMirror_8u_C1R>::call, 0, NppMirror<CV_8U, nppiMirror_8u_C3R>::call, NppMirror<CV_8U, nppiMirror_8u_C4R>::call},
-        {0,0,0,0},
-        {NppMirror<CV_16U, nppiMirror_16u_C1R>::call, 0, NppMirror<CV_16U, nppiMirror_16u_C3R>::call, NppMirror<CV_16U, nppiMirror_16u_C4R>::call},
-        {0,0,0,0},
-        {NppMirror<CV_32S, nppiMirror_32s_C1R>::call, 0, NppMirror<CV_32S, nppiMirror_32s_C3R>::call, NppMirror<CV_32S, nppiMirror_32s_C4R>::call},
-        {NppMirror<CV_32F, nppiMirror_32f_C1R>::call, 0, NppMirror<CV_32F, nppiMirror_32f_C3R>::call, NppMirror<CV_32F, nppiMirror_32f_C4R>::call}
-    };
-
-    GpuMat src = getInputMat(_src, stream);
-
-    CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32S || src.depth() == CV_32F);
-    CV_Assert(src.channels() == 1 || src.channels() == 3 || src.channels() == 4);
-
-    _dst.create(src.size(), src.type());
-    GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream);
-
-    funcs[src.depth()][src.channels() - 1](src, dst, flipCode, StreamAccessor::getStream(stream));
-
-    syncOutput(dst, _dst, stream);
-}
-
-#endif /* !defined (HAVE_CUDA) */

modules/cudaarithm/src/cuda/absdiff_mat.cu

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-
-#include "opencv2/opencv_modules.hpp"
-
-#ifndef HAVE_OPENCV_CUDEV
-
-#error "opencv_cudev is required"
-
-#else
-
-#include "opencv2/cudev.hpp"
-
-using namespace cv::cudev;
-
-void absDiffMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int);
-
-namespace
-{
-    __device__ __forceinline__ int _abs(int a)
-    {
-        return ::abs(a);
-    }
-    __device__ __forceinline__ float _abs(float a)
-    {
-        return ::fabsf(a);
-    }
-    __device__ __forceinline__ double _abs(double a)
-    {
-        return ::fabs(a);
-    }
-
-    template <typename T> struct AbsDiffOp1 : binary_function<T, T, T>
-    {
-        __device__ __forceinline__ T operator ()(T a, T b) const
-        {
-            return saturate_cast<T>(_abs(a - b));
-        }
-    };
-
-    template <typename ScalarDepth> struct TransformPolicy : DefaultTransformPolicy
-    {
-    };
-    template <> struct TransformPolicy<double> : DefaultTransformPolicy
-    {
-        enum {
-            shift = 1
-        };
-    };
-
-    template <typename T>
-    void absDiffMat_v1(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
-    {
-        gridTransformBinary_< TransformPolicy<T> >(globPtr<T>(src1), globPtr<T>(src2), globPtr<T>(dst), AbsDiffOp1<T>(), stream);
-    }
-
-    struct AbsDiffOp2 : binary_function<uint, uint, uint>
-    {
-        __device__ __forceinline__ uint operator ()(uint a, uint b) const
-        {
-            return vabsdiff2(a, b);
-        }
-    };
-
-    void absDiffMat_v2(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
-    {
-        const int vcols = src1.cols >> 1;
-
-        GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
-        GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
-        GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
-
-        gridTransformBinary(src1_, src2_, dst_, AbsDiffOp2(), stream);
-    }
-
-    struct AbsDiffOp4 : binary_function<uint, uint, uint>
-    {
-        __device__ __forceinline__ uint operator ()(uint a, uint b) const
-        {
-            return vabsdiff4(a, b);
-        }
-    };
-
-    void absDiffMat_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
-    {
-        const int vcols = src1.cols >> 2;
-
-        GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
-        GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
-        GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
-
-        gridTransformBinary(src1_, src2_, dst_, AbsDiffOp4(), stream);
-    }
-}
-
-void absDiffMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int)
-{
-    typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream);
-    static const func_t funcs[] =
-    {
-        absDiffMat_v1<uchar>,
-        absDiffMat_v1<schar>,
-        absDiffMat_v1<ushort>,
-        absDiffMat_v1<short>,
-        absDiffMat_v1<int>,
-        absDiffMat_v1<float>,
-        absDiffMat_v1<double>
-    };
-
-    const int depth = src1.depth();
-
-    CV_DbgAssert( depth <= CV_64F );
-
-    GpuMat src1_ = src1.reshape(1);
-    GpuMat src2_ = src2.reshape(1);
-    GpuMat dst_ = dst.reshape(1);
-
-    if (depth == CV_8U || depth == CV_16U)
-    {
-        const intptr_t src1ptr = reinterpret_cast<intptr_t>(src1_.data);
-        const intptr_t src2ptr = reinterpret_cast<intptr_t>(src2_.data);
-        const intptr_t dstptr = reinterpret_cast<intptr_t>(dst_.data);
-
-        const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0;
-
-        if (isAllAligned)
-        {
-            if (depth == CV_8U && (src1_.cols & 3) == 0)
-            {
-                absDiffMat_v4(src1_, src2_, dst_, stream);
-                return;
-            }
-            else if (depth == CV_16U && (src1_.cols & 1) == 0)
-            {
-                absDiffMat_v2(src1_, src2_, dst_, stream);
-                return;
-            }
-        }
-    }
-
-    const func_t func = funcs[depth];
-
-    if (!func)
-        CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types");
-
-    func(src1_, src2_, dst_, stream);
-}
-
-#endif

modules/cudaarithm/src/cuda/absdiff_scalar.cu

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-
-#include "opencv2/opencv_modules.hpp"
-
-#ifndef HAVE_OPENCV_CUDEV
-
-#error "opencv_cudev is required"
-
-#else
-
-#include "opencv2/cudev.hpp"
-
-using namespace cv::cudev;
-
-void absDiffScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int);
-
-namespace
-{
-    template <typename SrcType, typename ScalarType, typename DstType> struct AbsDiffScalarOp : unary_function<SrcType, DstType>
-    {
-        ScalarType val;
-
-        __device__ __forceinline__ DstType operator ()(SrcType a) const
-        {
-            abs_func<ScalarType> f;
-            return saturate_cast<DstType>(f(saturate_cast<ScalarType>(a) - val));
-        }
-    };
-
-    template <typename ScalarDepth> struct TransformPolicy : DefaultTransformPolicy
-    {
-    };
-    template <> struct TransformPolicy<double> : DefaultTransformPolicy
-    {
-        enum {
-            shift = 1
-        };
-    };
-
-    template <typename SrcType, typename ScalarDepth>
-    void absDiffScalarImpl(const GpuMat& src, cv::Scalar value, GpuMat& dst, Stream& stream)
-    {
-        typedef typename MakeVec<ScalarDepth, VecTraits<SrcType>::cn>::type ScalarType;
-
-        cv::Scalar_<ScalarDepth> value_ = value;
-
-        AbsDiffScalarOp<SrcType, ScalarType, SrcType> op;
-        op.val = VecTraits<ScalarType>::make(value_.val);
-        gridTransformUnary_< TransformPolicy<ScalarDepth> >(globPtr<SrcType>(src), globPtr<SrcType>(dst), op, stream);
-    }
-}
-
-void absDiffScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int)
-{
-    typedef void (*func_t)(const GpuMat& src, cv::Scalar val, GpuMat& dst, Stream& stream);
-    static const func_t funcs[7][4] =
-    {
-        {
-            absDiffScalarImpl<uchar, float>, absDiffScalarImpl<uchar2, float>, absDiffScalarImpl<uchar3, float>, absDiffScalarImpl<uchar4, float>
-        },
-        {
-            absDiffScalarImpl<schar, float>, absDiffScalarImpl<char2, float>, absDiffScalarImpl<char3, float>, absDiffScalarImpl<char4, float>
-        },
-        {
-            absDiffScalarImpl<ushort, float>, absDiffScalarImpl<ushort2, float>, absDiffScalarImpl<ushort3, float>, absDiffScalarImpl<ushort4, float>
-        },
-        {
-            absDiffScalarImpl<short, float>, absDiffScalarImpl<short2, float>, absDiffScalarImpl<short3, float>, absDiffScalarImpl<short4, float>
-        },
-        {
-            absDiffScalarImpl<int, float>, absDiffScalarImpl<int2, float>, absDiffScalarImpl<int3, float>, absDiffScalarImpl<int4, float>
-        },
-        {
-          absDiffScalarImpl<float, float>, absDiffScalarImpl<float2, float>, absDiffScalarImpl<float3, float>, absDiffScalarImpl<float4, float>
-        },
-        {
-          absDiffScalarImpl<double, double>, absDiffScalarImpl<double2, double>, absDiffScalarImpl<double3, double>, absDiffScalarImpl<double4, double>
-        }
-    };
-
-    const int sdepth = src.depth();
-    const int cn = src.channels();
-
-    CV_DbgAssert( sdepth <= CV_64F && cn <= 4 && src.type() == dst.type());
-
-    const func_t func = funcs[sdepth][cn - 1];
-    if (!func)
-        CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types");
-
-    func(src, val, dst, stream);
-}
-
-#endif

modules/cudaarithm/src/cuda/add_mat.cu

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-
-#include "opencv2/opencv_modules.hpp"
-
-#ifndef HAVE_OPENCV_CUDEV
-
-#error "opencv_cudev is required"
-
-#else
-
-#include "opencv2/cudev.hpp"
-
-using namespace cv::cudev;
-
-void addMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int);
-
-namespace
-{
-    template <typename T, typename D> struct AddOp1 : binary_function<T, T, D>
-    {
-        __device__ __forceinline__ D operator ()(T a, T b) const
-        {
-            return saturate_cast<D>(a + b);
-        }
-    };
-
-    template <typename T, typename D>
-    void addMat_v1(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, Stream& stream)
-    {
-        if (mask.data)
-            gridTransformBinary(globPtr<T>(src1), globPtr<T>(src2), globPtr<D>(dst), AddOp1<T, D>(), globPtr<uchar>(mask), stream);
-        else
-            gridTransformBinary(globPtr<T>(src1), globPtr<T>(src2), globPtr<D>(dst), AddOp1<T, D>(), stream);
-    }
-
-    struct AddOp2 : binary_function<uint, uint, uint>
-    {
-        __device__ __forceinline__ uint operator ()(uint a, uint b) const
-        {
-            return vadd2(a, b);
-        }
-    };
-
-    void addMat_v2(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
-    {
-        const int vcols = src1.cols >> 1;
-
-        GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
-        GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
-        GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
-
-        gridTransformBinary(src1_, src2_, dst_, AddOp2(), stream);
-    }
-
-    struct AddOp4 : binary_function<uint, uint, uint>
-    {
-        __device__ __forceinline__ uint operator ()(uint a, uint b) const
-        {
-            return vadd4(a, b);
-        }
-    };
-
-    void addMat_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)
-    {
-        const int vcols = src1.cols >> 2;
-
-        GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols);
-        GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols);
-        GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols);
-
-        gridTransformBinary(src1_, src2_, dst_, AddOp4(), stream);
-    }
-}
-
-void addMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int)
-{
-    typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, Stream& stream);
-    static const func_t funcs[7][7] =
-    {
-        {
-            addMat_v1<uchar, uchar>,
-            addMat_v1<uchar, schar>,
-            addMat_v1<uchar, ushort>,
-            addMat_v1<uchar, short>,
-            addMat_v1<uchar, int>,
-            addMat_v1<uchar, float>,
-            addMat_v1<uchar, double>
-        },
-        {
-            addMat_v1<schar, uchar>,
-            addMat_v1<schar, schar>,
-            addMat_v1<schar, ushort>,
-            addMat_v1<schar, short>,
-            addMat_v1<schar, int>,
-            addMat_v1<schar, float>,
-            addMat_v1<schar, double>
-        },
-        {
-            0 /*addMat_v1<ushort, uchar>*/,
-            0 /*addMat_v1<ushort, schar>*/,
-            addMat_v1<ushort, ushort>,
-            addMat_v1<ushort, short>,
-            addMat_v1<ushort, int>,
-            addMat_v1<ushort, float>,
-            addMat_v1<ushort, double>
-        },
-        {
-            0 /*addMat_v1<short, uchar>*/,
-            0 /*addMat_v1<short, schar>*/,
-            addMat_v1<short, ushort>,
-            addMat_v1<short, short>,
-            addMat_v1<short, int>,
-            addMat_v1<short, float>,
-            addMat_v1<short, double>
-        },
-        {
-            0 /*addMat_v1<int, uchar>*/,
-            0 /*addMat_v1<int, schar>*/,
-            0 /*addMat_v1<int, ushort>*/,
-            0 /*addMat_v1<int, short>*/,
-            addMat_v1<int, int>,
-            addMat_v1<int, float>,
-            addMat_v1<int, double>
-        },
-        {
-            0 /*addMat_v1<float, uchar>*/,
-            0 /*addMat_v1<float, schar>*/,
-            0 /*addMat_v1<float, ushort>*/,
-            0 /*addMat_v1<float, short>*/,
-            0 /*addMat_v1<float, int>*/,
-            addMat_v1<float, float>,
-            addMat_v1<float, double>
-        },
-        {
-            0 /*addMat_v1<double, uchar>*/,
-            0 /*addMat_v1<double, schar>*/,
-            0 /*addMat_v1<double, ushort>*/,
-            0 /*addMat_v1<double, short>*/,
-            0 /*addMat_v1<double, int>*/,
-            0 /*addMat_v1<double, float>*/,
-            addMat_v1<double, double>
-        }
-    };
-
-    const int sdepth = src1.depth();
-    const int ddepth = dst.depth();
-
-    CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F );
-
-    GpuMat src1_ = src1.reshape(1);
-    GpuMat src2_ = src2.reshape(1);
-    GpuMat dst_ = dst.reshape(1);
-
-    if (mask.empty() && (sdepth == CV_8U || sdepth == CV_16U) && ddepth == sdepth)
-    {
-        const intptr_t src1ptr = reinterpret_cast<intptr_t>(src1_.data);
-        const intptr_t src2ptr = reinterpret_cast<intptr_t>(src2_.data);
-        const intptr_t dstptr = reinterpret_cast<intptr_t>(dst_.data);
-
-        const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0;
-
-        if (isAllAligned)
-        {
-            if (sdepth == CV_8U && (src1_.cols & 3) == 0)
-            {
-                addMat_v4(src1_, src2_, dst_, stream);
-                return;
-            }
-            else if (sdepth == CV_16U && (src1_.cols & 1) == 0)
-            {
-                addMat_v2(src1_, src2_, dst_, stream);
-                return;
-            }
-        }
-    }
-
-    const func_t func = funcs[sdepth][ddepth];
-
-    if (!func)
-        CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types");
-
-    func(src1_, src2_, dst_, mask, stream);
-}
-
-#endif

modules/cudaarithm/src/cuda/countnonzero.cu

-/*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) 2000-2008, Intel Corporation, all rights reserved.
-// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// 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 materials 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 Intel 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 interruption) 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*/
-
-#include "opencv2/opencv_modules.hpp"
-
-#ifndef HAVE_OPENCV_CUDEV
-
-#error "opencv_cudev is required"
-
-#else
-
-#include "opencv2/cudaarithm.hpp"
-#include "opencv2/cudev.hpp"
-#include "opencv2/core/private.cuda.hpp"
-
-using namespace cv;
-using namespace cv::cuda;
-using namespace cv::cudev;
-
-namespace
-{
-    template <typename T, typename D>
-    void countNonZeroImpl(const GpuMat& _src, GpuMat& _dst, Stream& stream)
-    {
-        const GpuMat_<T>& src = (const GpuMat_<T>&) _src;
-        GpuMat_<D>& dst = (GpuMat_<D>&) _dst;
-
-        gridCountNonZero(src, dst, stream);
-    }
-}
-
-void cv::cuda::countNonZero(InputArray _src, OutputArray _dst, Stream& stream)
-{
-    typedef void (*func_t)(const GpuMat& src, GpuMat& dst, Stream& stream);
-    static const func_t funcs[] =
-    {
-        countNonZeroImpl<uchar, int>,
-        countNonZeroImpl<schar, int>,
-        countNonZeroImpl<ushort, int>,
-        countNonZeroImpl<short, int>,
-        countNonZeroImpl<int, int>,
-        countNonZeroImpl<float, int>,
-        countNonZeroImpl<double, int>,
-    };
-
-    GpuMat src = getInputMat(_src, stream);
-
-    CV_Assert( src.depth() <= CV_64F );
-    CV_Assert( src.channels() == 1 );
-
-    GpuMat dst = getOutputMat(_dst, 1, 1, CV_32SC1, stream);
-
-    const func_t func = funcs[src.depth()];
-    func(src, dst, stream);
-
-    syncOutput(dst, _dst, stream);
-}
-
-int cv::cuda::countNonZero(InputArray _src)
-{
-    Stream& stream = Stream::Null();
-
-    BufferPool pool(stream);
-    GpuMat buf = pool.getBuffer(1, 1, CV_32SC1);
-
-    countNonZero(_src, buf, stream);
-
-    int data;
-    buf.download(Mat(1, 1, CV_32SC1, &data));
-
-    return data;
-}
-
-#endif

modules/cudabgsegm/CMakeLists.txt

-if(IOS OR (NOT HAVE_CUDA AND NOT BUILD_CUDA_STUBS))
-  ocv_module_disable(cudabgsegm)
-endif()
-
-set(the_description "CUDA-accelerated Background Segmentation")
-
-ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations -Wshadow)
-
-ocv_define_module(cudabgsegm opencv_video WRAP python)

modules/cudafeatures2d/CMakeLists.txt

-if(IOS OR WINRT OR (NOT HAVE_CUDA AND NOT BUILD_CUDA_STUBS))
-  ocv_module_disable(cudafeatures2d)
-endif()
-
-set(the_description "CUDA-accelerated Feature Detection and Description")
-
-ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4100 /wd4324 /wd4512 /wd4515 -Wundef -Wmissing-declarations -Wshadow -Wunused-parameter -Wshadow)
-
-ocv_define_module(cudafeatures2d opencv_features2d opencv_cudafilters opencv_cudawarping WRAP python)