CUDA Device Layer:

header only library for CUDA programming

modules/cudev/CMakeLists.txt

+if(NOT HAVE_CUDA)
+  ocv_module_disable(cudev)
+endif()
+
+set(the_description "CUDA device layer")
+
+ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4189 /wd4505 -Wundef -Wmissing-declarations -Wunused-function -Wunused-variable)
+
+ocv_add_module(cudev)
+
+file(GLOB_RECURSE lib_hdrs "include/opencv2/*.hpp")
+
+add_custom_target(${the_module} SOURCES ${lib_hdrs})
+
+if(ENABLE_SOLUTION_FOLDERS)
+  set_target_properties(${the_module} PROPERTIES FOLDER "modules")
+endif()
+
+foreach(var CMAKE_CXX_FLAGS CMAKE_CXX_FLAGS_RELEASE CMAKE_CXX_FLAGS_DEBUG)
+  # we remove /EHa as it generates warnings under windows
+  string(REPLACE "/EHa" "" ${var} "${${var}}")
+
+  # we remove -ggdb3 flag as it leads to preprocessor errors when compiling CUDA files (CUDA 4.1)
+  string(REPLACE "-ggdb3" "" ${var} "${${var}}")
+
+  # we remove -Wsign-promo as it generates warnings under linux
+  string(REPLACE "-Wsign-promo" "" ${var} "${${var}}")
+
+  # we remove -fvisibility-inlines-hidden because it's used for C++ compiler
+  # but NVCC uses C compiler by default
+  string(REPLACE "-fvisibility-inlines-hidden" "" ${var} "${${var}}")
+endforeach()
+
+if(BUILD_TESTS)
+  add_subdirectory(test)
+endif()

modules/cudev/include/opencv2/cudev.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_HPP__
+#define __OPENCV_CUDEV_HPP__
+
+#include "cudev/common.hpp"
+
+#include "cudev/util/atomic.hpp"
+#include "cudev/util/limits.hpp"
+#include "cudev/util/saturate_cast.hpp"
+#include "cudev/util/simd_functions.hpp"
+#include "cudev/util/tuple.hpp"
+#include "cudev/util/type_traits.hpp"
+#include "cudev/util/vec_math.hpp"
+#include "cudev/util/vec_traits.hpp"
+
+#include "cudev/functional/color_cvt.hpp"
+#include "cudev/functional/functional.hpp"
+#include "cudev/functional/tuple_adapter.hpp"
+
+#include "cudev/warp/reduce.hpp"
+#include "cudev/warp/scan.hpp"
+#include "cudev/warp/shuffle.hpp"
+#include "cudev/warp/warp.hpp"
+
+#include "cudev/block/block.hpp"
+#include "cudev/block/dynamic_smem.hpp"
+#include "cudev/block/reduce.hpp"
+#include "cudev/block/scan.hpp"
+#include "cudev/block/vec_distance.hpp"
+
+#include "cudev/grid/copy.hpp"
+#include "cudev/grid/glob_reduce.hpp"
+#include "cudev/grid/histogram.hpp"
+#include "cudev/grid/integral.hpp"
+#include "cudev/grid/pyramids.hpp"
+#include "cudev/grid/reduce_to_vec.hpp"
+#include "cudev/grid/split_merge.hpp"
+#include "cudev/grid/transform.hpp"
+#include "cudev/grid/transpose.hpp"
+
+#include "cudev/ptr2d/constant.hpp"
+#include "cudev/ptr2d/deriv.hpp"
+#include "cudev/ptr2d/extrapolation.hpp"
+#include "cudev/ptr2d/glob.hpp"
+#include "cudev/ptr2d/gpumat.hpp"
+#include "cudev/ptr2d/interpolation.hpp"
+#include "cudev/ptr2d/lut.hpp"
+#include "cudev/ptr2d/mask.hpp"
+#include "cudev/ptr2d/remap.hpp"
+#include "cudev/ptr2d/resize.hpp"
+#include "cudev/ptr2d/texture.hpp"
+#include "cudev/ptr2d/traits.hpp"
+#include "cudev/ptr2d/transform.hpp"
+#include "cudev/ptr2d/warping.hpp"
+#include "cudev/ptr2d/zip.hpp"
+
+#include "cudev/expr/binary_func.hpp"
+#include "cudev/expr/binary_op.hpp"
+#include "cudev/expr/color.hpp"
+#include "cudev/expr/deriv.hpp"
+#include "cudev/expr/expr.hpp"
+#include "cudev/expr/per_element_func.hpp"
+#include "cudev/expr/reduction.hpp"
+#include "cudev/expr/unary_func.hpp"
+#include "cudev/expr/unary_op.hpp"
+#include "cudev/expr/warping.hpp"
+
+#endif

modules/cudev/include/opencv2/cudev/block/block.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_BLOCK_BLOCK_HPP__
+#define __OPENCV_CUDEV_BLOCK_BLOCK_HPP__
+
+#include "../common.hpp"
+
+namespace cv { namespace cudev {
+
+struct Block
+{
+    __device__ __forceinline__ static uint blockId()
+    {
+        return (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x;
+    }
+
+    __device__ __forceinline__ static uint blockSize()
+    {
+        return blockDim.x * blockDim.y * blockDim.z;
+    }
+
+    __device__ __forceinline__ static uint threadLineId()
+    {
+        return (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x;
+    }
+};
+
+template <class It, typename T>
+__device__ __forceinline__ static void blockFill(It beg, It end, const T& value)
+{
+    uint STRIDE = Block::blockSize();
+    It t = beg + Block::threadLineId();
+
+    for(; t < end; t += STRIDE)
+        *t = value;
+}
+
+template <class OutIt, typename T>
+__device__ __forceinline__ static void blockYota(OutIt beg, OutIt end, T value)
+{
+    uint STRIDE = Block::blockSize();
+    uint tid = Block::threadLineId();
+    value += tid;
+
+    for(OutIt t = beg + tid; t < end; t += STRIDE, value += STRIDE)
+        *t = value;
+}
+
+template <class InIt, class OutIt>
+__device__ __forceinline__ static void blockCopy(InIt beg, InIt end, OutIt out)
+{
+    uint STRIDE = Block::blockSize();
+    InIt  t = beg + Block::threadLineId();
+    OutIt o = out + (t - beg);
+
+    for(; t < end; t += STRIDE, o += STRIDE)
+        *o = *t;
+}
+
+template <class InIt, class OutIt, class UnOp>
+__device__ __forceinline__ static void blockTransfrom(InIt beg, InIt end, OutIt out, const UnOp& op)
+{
+    uint STRIDE = Block::blockSize();
+    InIt  t = beg + Block::threadLineId();
+    OutIt o = out + (t - beg);
+
+    for(; t < end; t += STRIDE, o += STRIDE)
+        *o = op(*t);
+}
+
+template <class InIt1, class InIt2, class OutIt, class BinOp>
+__device__ __forceinline__ static void blockTransfrom(InIt1 beg1, InIt1 end1, InIt2 beg2, OutIt out, const BinOp& op)
+{
+    uint STRIDE = Block::blockSize();
+    InIt1 t1 = beg1 + Block::threadLineId();
+    InIt2 t2 = beg2 + Block::threadLineId();
+    OutIt o  = out + (t1 - beg1);
+
+    for(; t1 < end1; t1 += STRIDE, t2 += STRIDE, o += STRIDE)
+        *o = op(*t1, *t2);
+}
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/block/dynamic_smem.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_BLOCK_DYNAMIC_SMEM_HPP__
+#define __OPENCV_CUDEV_BLOCK_DYNAMIC_SMEM_HPP__
+
+#include "../common.hpp"
+
+namespace cv { namespace cudev {
+
+template <class T> struct DynamicSharedMem
+{
+    __device__ __forceinline__ operator T*()
+    {
+        extern __shared__ int __smem[];
+        return (T*) __smem;
+    }
+
+    __device__ __forceinline__ operator const T*() const
+    {
+        extern __shared__ int __smem[];
+        return (T*) __smem;
+    }
+};
+
+// specialize for double to avoid unaligned memory access compile errors
+template <> struct DynamicSharedMem<double>
+{
+    __device__ __forceinline__ operator double*()
+    {
+        extern __shared__ double __smem_d[];
+        return (double*) __smem_d;
+    }
+
+    __device__ __forceinline__ operator const double*() const
+    {
+        extern __shared__ double __smem_d[];
+        return (double*) __smem_d;
+    }
+};
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/block/reduce.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_BLOCK_REDUCE_HPP__
+#define __OPENCV_CUDEV_BLOCK_REDUCE_HPP__
+
+#include "../common.hpp"
+#include "../util/tuple.hpp"
+#include "../warp/reduce.hpp"
+#include "detail/reduce.hpp"
+#include "detail/reduce_key_val.hpp"
+
+namespace cv { namespace cudev {
+
+// blockReduce
+
+template <int N, typename T, class Op>
+__device__ __forceinline__ void blockReduce(volatile T* smem, T& val, uint tid, const Op& op)
+{
+    block_reduce_detail::Dispatcher<N>::reductor::template reduce<volatile T*, T&, const Op&>(smem, val, tid, op);
+}
+
+template <int N,
+          typename P0, typename P1, typename P2, typename P3, typename P4, typename P5, typename P6, typename P7, typename P8, typename P9,
+          typename R0, typename R1, typename R2, typename R3, typename R4, typename R5, typename R6, typename R7, typename R8, typename R9,
+          class Op0, class Op1, class Op2, class Op3, class Op4, class Op5, class Op6, class Op7, class Op8, class Op9>
+__device__ __forceinline__ void blockReduce(const tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>& smem,
+                                            const tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9>& val,
+                                            uint tid,
+                                            const tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>& op)
+{
+    block_reduce_detail::Dispatcher<N>::reductor::template reduce<
+            const tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>&,
+            const tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9>&,
+            const tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>&>(smem, val, tid, op);
+}
+
+// blockReduceKeyVal
+
+template <int N, typename K, typename V, class Cmp>
+__device__ __forceinline__ void blockReduceKeyVal(volatile K* skeys, K& key, volatile V* svals, V& val, uint tid, const Cmp& cmp)
+{
+    block_reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, volatile V*, V&, const Cmp&>(skeys, key, svals, val, tid, cmp);
+}
+
+template <int N,
+          typename K,
+          typename VP0, typename VP1, typename VP2, typename VP3, typename VP4, typename VP5, typename VP6, typename VP7, typename VP8, typename VP9,
+          typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9,
+          class Cmp>
+__device__ __forceinline__ void blockReduceKeyVal(volatile K* skeys, K& key,
+                                                  const tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals,
+                                                  const tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val,
+                                                  uint tid, const Cmp& cmp)
+{
+    block_reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&,
+            const tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>&,
+            const tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&,
+            const Cmp&>(skeys, key, svals, val, tid, cmp);
+}
+
+template <int N,
+          typename KP0, typename KP1, typename KP2, typename KP3, typename KP4, typename KP5, typename KP6, typename KP7, typename KP8, typename KP9,
+          typename KR0, typename KR1, typename KR2, typename KR3, typename KR4, typename KR5, typename KR6, typename KR7, typename KR8, typename KR9,
+          typename VP0, typename VP1, typename VP2, typename VP3, typename VP4, typename VP5, typename VP6, typename VP7, typename VP8, typename VP9,
+          typename VR0, typename VR1, typename VR2, typename VR3, typename VR4, typename VR5, typename VR6, typename VR7, typename VR8, typename VR9,
+          class Cmp0, class Cmp1, class Cmp2, class Cmp3, class Cmp4, class Cmp5, class Cmp6, class Cmp7, class Cmp8, class Cmp9>
+__device__ __forceinline__ void blockReduceKeyVal(const tuple<KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9>& skeys,
+                                                  const tuple<KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9>& key,
+                                                  const tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals,
+                                                  const tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val,
+                                                  uint tid,
+                                                  const tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>& cmp)
+{
+    block_reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<
+            const tuple<KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9>&,
+            const tuple<KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9>&,
+            const tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>&,
+            const tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&,
+            const tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>&
+            >(skeys, key, svals, val, tid, cmp);
+}
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/block/scan.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_BLOCK_SCAN_HPP__
+#define __OPENCV_CUDEV_BLOCK_SCAN_HPP__
+
+#include "../common.hpp"
+#include "../warp/scan.hpp"
+
+namespace cv { namespace cudev {
+
+template <int THREADS_NUM, typename T>
+__device__ T blockScanInclusive(T data, volatile T* smem, uint tid)
+{
+    if (THREADS_NUM > WARP_SIZE)
+    {
+        // bottom-level inclusive warp scan
+        T warpResult = warpScanInclusive(data, smem, tid);
+
+        __syncthreads();
+
+        // save top elements of each warp for exclusive warp scan
+        // sync to wait for warp scans to complete (because s_Data is being overwritten)
+        if ((tid & (WARP_SIZE - 1)) == (WARP_SIZE - 1))
+        {
+            smem[tid >> LOG_WARP_SIZE] = warpResult;
+        }
+
+        __syncthreads();
+
+        if (tid < (THREADS_NUM / WARP_SIZE))
+        {
+            // grab top warp elements
+            T val = smem[tid];
+
+            // calculate exclusive scan and write back to shared memory
+            smem[tid] = warpScanExclusive(val, smem, tid);
+        }
+
+        __syncthreads();
+
+        // return updated warp scans with exclusive scan results
+        return warpResult + smem[tid >> LOG_WARP_SIZE];
+    }
+    else
+    {
+        return warpScanInclusive(data, smem, tid);
+    }
+}
+
+template <int THREADS_NUM, typename T>
+__device__ __forceinline__ T blockScanExclusive(T data, volatile T* smem, uint tid)
+{
+    return blockScanInclusive<THREADS_NUM>(data, smem, tid) - data;
+}
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/block/vec_distance.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_BLOCK_VEC_DISTANCE_HPP__
+#define __OPENCV_CUDEV_BLOCK_VEC_DISTANCE_HPP__
+
+#include "../common.hpp"
+#include "../functional/functional.hpp"
+#include "../warp/reduce.hpp"
+#include "reduce.hpp"
+
+namespace cv { namespace cudev {
+
+// NormL1
+
+template <typename T> struct NormL1
+{
+    typedef int value_type;
+    typedef uint result_type;
+
+    result_type mySum;
+
+    __device__ __forceinline__ NormL1() : mySum(0) {}
+
+    __device__ __forceinline__ void reduceThread(value_type val1, value_type val2)
+    {
+        mySum = __sad(val1, val2, mySum);
+    }
+
+    __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid)
+    {
+        warpReduce(smem, mySum, tid, plus<result_type>());
+    }
+
+    template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid)
+    {
+        blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>());
+    }
+
+    __device__ __forceinline__ operator result_type() const
+    {
+        return mySum;
+    }
+};
+template <> struct NormL1<float>
+{
+    typedef float value_type;
+    typedef float result_type;
+
+    result_type mySum;
+
+    __device__ __forceinline__ NormL1() : mySum(0.0f) {}
+
+    __device__ __forceinline__ void reduceThread(value_type val1, value_type val2)
+    {
+        mySum += ::fabsf(val1 - val2);
+    }
+
+    __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid)
+    {
+        warpReduce(smem, mySum, tid, plus<result_type>());
+    }
+
+    template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid)
+    {
+        blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>());
+    }
+
+    __device__ __forceinline__ operator result_type() const
+    {
+        return mySum;
+    }
+};
+
+// NormL2
+
+struct NormL2
+{
+    typedef float value_type;
+    typedef float result_type;
+
+    result_type mySum;
+
+    __device__ __forceinline__ NormL2() : mySum(0.0f) {}
+
+    __device__ __forceinline__ void reduceThread(value_type val1, value_type val2)
+    {
+        const float diff = val1 - val2;
+        mySum += diff * diff;
+    }
+
+    __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid)
+    {
+        warpReduce(smem, mySum, tid, plus<result_type>());
+    }
+
+    template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid)
+    {
+        blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>());
+    }
+
+    __device__ __forceinline__ operator result_type() const
+    {
+        return ::sqrtf(mySum);
+    }
+};
+
+// NormHamming
+
+struct NormHamming
+{
+    typedef int value_type;
+    typedef int result_type;
+
+    result_type mySum;
+
+    __device__ __forceinline__ NormHamming() : mySum(0) {}
+
+    __device__ __forceinline__ void reduceThread(value_type val1, value_type val2)
+    {
+        mySum += __popc(val1 ^ val2);
+    }
+
+    __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid)
+    {
+        warpReduce(smem, mySum, tid, plus<result_type>());
+    }
+
+    template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid)
+    {
+        blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>());
+    }
+
+    __device__ __forceinline__ operator result_type() const
+    {
+        return mySum;
+    }
+};
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/common.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_COMMON_HPP__
+#define __OPENCV_CUDEV_COMMON_HPP__
+
+#include <cuda_runtime.h>
+#include "opencv2/core/gpu.hpp"
+#include "opencv2/core/gpu_stream_accessor.hpp"
+
+namespace cv { namespace cudev {
+
+using namespace cv::gpu;
+
+// CV_CUDEV_ARCH
+
+#ifndef __CUDA_ARCH__
+#   define CV_CUDEV_ARCH 0
+#else
+#   define CV_CUDEV_ARCH __CUDA_ARCH__
+#endif
+
+// CV_CUDEV_SAFE_CALL
+
+__host__ __forceinline__ void checkCudaError(cudaError_t err, const char* file, const int line, const char* func)
+{
+    if (cudaSuccess != err)
+        cv::error(cv::Error::GpuApiCallError, cudaGetErrorString(err), func, file, line);
+}
+
+#ifdef __GNUC__
+#   define CV_CUDEV_SAFE_CALL(expr) cv::cudev::checkCudaError((expr), __FILE__, __LINE__, __func__)
+#else
+#   define CV_CUDEV_SAFE_CALL(expr) cv::cudev::checkCudaError((expr), __FILE__, __LINE__, "")
+#endif
+
+// divUp
+
+__host__ __device__ __forceinline__ int divUp(int total, int grain)
+{
+    return (total + grain - 1) / grain;
+}
+
+// math constants
+
+#define CV_PI_F   ((float)CV_PI)
+#define CV_LOG2_F ((float)CV_LOG2)
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/expr/binary_func.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_EXPR_BINARY_FUNC_HPP__
+#define __OPENCV_CUDEV_EXPR_BINARY_FUNC_HPP__
+
+#include "../common.hpp"
+#include "../util/type_traits.hpp"
+#include "../ptr2d/traits.hpp"
+#include "../ptr2d/transform.hpp"
+#include "../functional/functional.hpp"
+#include "expr.hpp"
+
+namespace cv { namespace cudev {
+
+#define CV_CUDEV_EXPR_BINARY_FUNC(name) \
+    template <class SrcPtr1, class SrcPtr2> \
+    __host__ Expr<BinaryTransformPtrSz<typename PtrTraits<SrcPtr1>::ptr_type, typename PtrTraits<SrcPtr2>::ptr_type, name ## _func<typename LargerType<typename PtrTraits<SrcPtr1>::value_type, typename PtrTraits<SrcPtr2>::value_type>::type> > > \
+    name ## _(const SrcPtr1& src1, const SrcPtr2& src2) \
+    { \
+        return makeExpr(transformPtr(src1, src2, name ## _func<typename LargerType<typename PtrTraits<SrcPtr1>::value_type, typename PtrTraits<SrcPtr2>::value_type>::type>())); \
+    }
+
+CV_CUDEV_EXPR_BINARY_FUNC(hypot)
+CV_CUDEV_EXPR_BINARY_FUNC(magnitude)
+CV_CUDEV_EXPR_BINARY_FUNC(atan2)
+CV_CUDEV_EXPR_BINARY_FUNC(absdiff)
+
+#undef CV_CUDEV_EXPR_BINARY_FUNC
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/expr/deriv.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_EXPR_DERIV_HPP__
+#define __OPENCV_CUDEV_EXPR_DERIV_HPP__
+
+#include "../common.hpp"
+#include "../ptr2d/traits.hpp"
+#include "../ptr2d/deriv.hpp"
+#include "expr.hpp"
+
+namespace cv { namespace cudev {
+
+// derivX
+
+template <class SrcPtr>
+__host__ Expr<DerivXPtrSz<typename PtrTraits<SrcPtr>::ptr_type> >
+derivX_(const SrcPtr& src)
+{
+    return makeExpr(derivXPtr(src));
+}
+
+// derivY
+
+template <class SrcPtr>
+__host__ Expr<DerivYPtrSz<typename PtrTraits<SrcPtr>::ptr_type> >
+derivY_(const SrcPtr& src)
+{
+    return makeExpr(derivYPtr(src));
+}
+
+// sobelX
+
+template <class SrcPtr>
+__host__ Expr<SobelXPtrSz<typename PtrTraits<SrcPtr>::ptr_type> >
+sobelX_(const SrcPtr& src)
+{
+    return makeExpr(sobelXPtr(src));
+}
+
+// sobelY
+
+template <class SrcPtr>
+__host__ Expr<SobelYPtrSz<typename PtrTraits<SrcPtr>::ptr_type> >
+sobelY_(const SrcPtr& src)
+{
+    return makeExpr(sobelYPtr(src));
+}
+
+// scharrX
+
+template <class SrcPtr>
+__host__ Expr<ScharrXPtrSz<typename PtrTraits<SrcPtr>::ptr_type> >
+scharrX_(const SrcPtr& src)
+{
+    return makeExpr(scharrXPtr(src));
+}
+
+// scharrY
+
+template <class SrcPtr>
+__host__ Expr<ScharrYPtrSz<typename PtrTraits<SrcPtr>::ptr_type> >
+scharrY_(const SrcPtr& src)
+{
+    return makeExpr(scharrYPtr(src));
+}
+
+// laplacian
+
+template <int ksize, class SrcPtr>
+__host__ Expr<LaplacianPtrSz<ksize, typename PtrTraits<SrcPtr>::ptr_type> >
+laplacian_(const SrcPtr& src)
+{
+    return makeExpr(laplacianPtr<ksize>(src));
+}
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/expr/expr.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_EXPR_EXPR_HPP__
+#define __OPENCV_CUDEV_EXPR_EXPR_HPP__
+
+#include "../common.hpp"
+#include "../ptr2d/traits.hpp"
+
+namespace cv { namespace cudev {
+
+template <class Body> struct Expr
+{
+    Body body;
+};
+
+template <class Body>
+__host__ Expr<Body> makeExpr(const Body& body)
+{
+    Expr<Body> e;
+    e.body = body;
+    return e;
+}
+
+template <class Body> struct PtrTraits< Expr<Body> >
+{
+    typedef Expr<Body>                         ptr_sz_type;
+    typedef typename PtrTraits<Body>::ptr_type ptr_type;
+
+    typedef typename ptr_type::value_type value_type;
+
+    __host__ static ptr_type shrinkPtr(const Expr<Body>& expr)
+    {
+        return PtrTraits<Body>::shrinkPtr(expr.body);
+    }
+
+    __host__ static int getRows(const Expr<Body>& expr)
+    {
+        return PtrTraits<Body>::getRows(expr.body);
+    }
+
+    __host__ static int getCols(const Expr<Body>& expr)
+    {
+        return PtrTraits<Body>::getCols(expr.body);
+    }
+};
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/expr/per_element_func.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_EXPR_PER_ELEMENT_FUNC_HPP__
+#define __OPENCV_CUDEV_EXPR_PER_ELEMENT_FUNC_HPP__
+
+#include "../common.hpp"
+#include "../util/type_traits.hpp"
+#include "../ptr2d/traits.hpp"
+#include "../ptr2d/transform.hpp"
+#include "../ptr2d/lut.hpp"
+#include "../functional/functional.hpp"
+#include "expr.hpp"
+
+namespace cv { namespace cudev {
+
+// min/max
+
+template <class SrcPtr1, class SrcPtr2>
+__host__ Expr<BinaryTransformPtrSz<typename PtrTraits<SrcPtr1>::ptr_type, typename PtrTraits<SrcPtr2>::ptr_type, minimum<typename LargerType<typename PtrTraits<SrcPtr1>::value_type, typename PtrTraits<SrcPtr2>::value_type>::type> > >
+min_(const SrcPtr1& src1, const SrcPtr2& src2)
+{
+    return makeExpr(transformPtr(src1, src2, minimum<typename LargerType<typename PtrTraits<SrcPtr1>::value_type, typename PtrTraits<SrcPtr2>::value_type>::type>()));
+}
+
+template <class SrcPtr1, class SrcPtr2>
+__host__ Expr<BinaryTransformPtrSz<typename PtrTraits<SrcPtr1>::ptr_type, typename PtrTraits<SrcPtr2>::ptr_type, maximum<typename LargerType<typename PtrTraits<SrcPtr1>::value_type, typename PtrTraits<SrcPtr2>::value_type>::type> > >
+max_(const SrcPtr1& src1, const SrcPtr2& src2)
+{
+    return makeExpr(transformPtr(src1, src2, maximum<typename LargerType<typename PtrTraits<SrcPtr1>::value_type, typename PtrTraits<SrcPtr2>::value_type>::type>()));
+}
+
+// threshold
+
+template <class SrcPtr>
+__host__ Expr<UnaryTransformPtrSz<typename PtrTraits<SrcPtr>::ptr_type, ThreshBinaryFunc<typename PtrTraits<SrcPtr>::value_type> > >
+threshBinary_(const SrcPtr& src, typename PtrTraits<SrcPtr>::value_type thresh, typename PtrTraits<SrcPtr>::value_type maxVal)
+{
+    return makeExpr(transformPtr(src, thresh_binary_func(thresh, maxVal)));
+}
+
+template <class SrcPtr>
+__host__ Expr<UnaryTransformPtrSz<typename PtrTraits<SrcPtr>::ptr_type, ThreshBinaryInvFunc<typename PtrTraits<SrcPtr>::value_type> > >
+threshBinaryInv_(const SrcPtr& src, typename PtrTraits<SrcPtr>::value_type thresh, typename PtrTraits<SrcPtr>::value_type maxVal)
+{
+    return makeExpr(transformPtr(src, thresh_binary_inv_func(thresh, maxVal)));
+}
+
+template <class SrcPtr>
+__host__ Expr<UnaryTransformPtrSz<typename PtrTraits<SrcPtr>::ptr_type, ThreshTruncFunc<typename PtrTraits<SrcPtr>::value_type> > >
+threshTrunc_(const SrcPtr& src, typename PtrTraits<SrcPtr>::value_type thresh)
+{
+    return makeExpr(transformPtr(src, thresh_trunc_func(thresh)));
+}
+
+template <class SrcPtr>
+__host__ Expr<UnaryTransformPtrSz<typename PtrTraits<SrcPtr>::ptr_type, ThreshToZeroFunc<typename PtrTraits<SrcPtr>::value_type> > >
+threshToZero_(const SrcPtr& src, typename PtrTraits<SrcPtr>::value_type thresh)
+{
+    return makeExpr(transformPtr(src, thresh_to_zero_func(thresh)));
+}
+
+template <class SrcPtr>
+__host__ Expr<UnaryTransformPtrSz<typename PtrTraits<SrcPtr>::ptr_type, ThreshToZeroInvFunc<typename PtrTraits<SrcPtr>::value_type> > >
+threshToZeroInv_(const SrcPtr& src, typename PtrTraits<SrcPtr>::value_type thresh)
+{
+    return makeExpr(transformPtr(src, thresh_to_zero_inv_func(thresh)));
+}
+
+// cvt
+
+template <typename D, class SrcPtr>
+__host__ Expr<UnaryTransformPtrSz<typename PtrTraits<SrcPtr>::ptr_type, saturate_cast_func<typename PtrTraits<SrcPtr>::value_type, D> > >
+cvt_(const SrcPtr& src)
+{
+    return makeExpr(transformPtr(src, saturate_cast_func<typename PtrTraits<SrcPtr>::value_type, D>()));
+}
+
+// lut
+
+template <class SrcPtr, class TablePtr>
+__host__ Expr<LutPtrSz<typename PtrTraits<SrcPtr>::ptr_type, typename PtrTraits<TablePtr>::ptr_type> >
+lut_(const SrcPtr& src, const TablePtr& tbl)
+{
+    return makeExpr(lutPtr(src, tbl));
+}
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/expr/reduction.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_EXPR_REDUCTION_HPP__
+#define __OPENCV_CUDEV_EXPR_REDUCTION_HPP__
+
+#include "../common.hpp"
+#include "../grid/glob_reduce.hpp"
+#include "../grid/histogram.hpp"
+#include "../grid/integral.hpp"
+#include "../grid/reduce_to_vec.hpp"
+#include "../ptr2d/traits.hpp"
+#include "expr.hpp"
+
+namespace cv { namespace cudev {
+
+// sum
+
+template <class SrcPtr> struct SumExprBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridCalcSum(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<SumExprBody<SrcPtr> >
+sum_(const SrcPtr& src)
+{
+    SumExprBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// minVal
+
+template <class SrcPtr> struct FindMinValExprBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridFindMinVal(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<FindMinValExprBody<SrcPtr> >
+minVal_(const SrcPtr& src)
+{
+    FindMinValExprBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// maxVal
+
+template <class SrcPtr> struct FindMaxValExprBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridFindMaxVal(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<FindMaxValExprBody<SrcPtr> >
+maxVal_(const SrcPtr& src)
+{
+    FindMaxValExprBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// minMaxVal
+
+template <class SrcPtr> struct FindMinMaxValExprBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridFindMinMaxVal(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<FindMinMaxValExprBody<SrcPtr> >
+minMaxVal_(const SrcPtr& src)
+{
+    FindMinMaxValExprBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// countNonZero
+
+template <class SrcPtr> struct CountNonZeroExprBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridCountNonZero(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<CountNonZeroExprBody<SrcPtr> >
+countNonZero_(const SrcPtr& src)
+{
+    CountNonZeroExprBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// reduceToRow
+
+template <class Reductor, class SrcPtr> struct ReduceToRowBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridReduceToRow<Reductor>(src, dst, stream);
+    }
+};
+
+template <class Reductor, class SrcPtr>
+__host__ Expr<ReduceToRowBody<Reductor, SrcPtr> >
+reduceToRow_(const SrcPtr& src)
+{
+    ReduceToRowBody<Reductor, SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// reduceToColumn
+
+template <class Reductor, class SrcPtr> struct ReduceToColumnBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridReduceToColumn<Reductor>(src, dst, stream);
+    }
+};
+
+template <class Reductor, class SrcPtr>
+__host__ Expr<ReduceToColumnBody<Reductor, SrcPtr> >
+reduceToColumn_(const SrcPtr& src)
+{
+    ReduceToColumnBody<Reductor, SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// histogram
+
+template <int BIN_COUNT, class SrcPtr> struct HistogramBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridHistogram<BIN_COUNT>(src, dst, stream);
+    }
+};
+
+template <int BIN_COUNT, class SrcPtr>
+__host__ Expr<HistogramBody<BIN_COUNT, SrcPtr> >
+histogram_(const SrcPtr& src)
+{
+    HistogramBody<BIN_COUNT, SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// integral
+
+template <class SrcPtr> struct IntegralBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridIntegral(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<IntegralBody<SrcPtr> >
+integral_(const SrcPtr& src)
+{
+    IntegralBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/expr/unary_func.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_EXPR_UNARY_FUNC_HPP__
+#define __OPENCV_CUDEV_EXPR_UNARY_FUNC_HPP__
+
+#include "../common.hpp"
+#include "../ptr2d/traits.hpp"
+#include "../ptr2d/transform.hpp"
+#include "../functional/functional.hpp"
+#include "expr.hpp"
+
+namespace cv { namespace cudev {
+
+#define CV_CUDEV_EXPR_UNARY_FUNC(name) \
+    template <class SrcPtr> \
+    __host__ Expr<UnaryTransformPtrSz<typename PtrTraits<SrcPtr>::ptr_type, name ## _func<typename PtrTraits<SrcPtr>::value_type> > > \
+    name ## _(const SrcPtr& src) \
+    { \
+        return makeExpr(transformPtr(src, name ## _func<typename PtrTraits<SrcPtr>::value_type>())); \
+    }
+
+CV_CUDEV_EXPR_UNARY_FUNC(abs)
+CV_CUDEV_EXPR_UNARY_FUNC(sqr)
+CV_CUDEV_EXPR_UNARY_FUNC(sqrt)
+CV_CUDEV_EXPR_UNARY_FUNC(exp)
+CV_CUDEV_EXPR_UNARY_FUNC(exp2)
+CV_CUDEV_EXPR_UNARY_FUNC(exp10)
+CV_CUDEV_EXPR_UNARY_FUNC(log)
+CV_CUDEV_EXPR_UNARY_FUNC(log2)
+CV_CUDEV_EXPR_UNARY_FUNC(log10)
+CV_CUDEV_EXPR_UNARY_FUNC(sin)
+CV_CUDEV_EXPR_UNARY_FUNC(cos)
+CV_CUDEV_EXPR_UNARY_FUNC(tan)
+CV_CUDEV_EXPR_UNARY_FUNC(asin)
+CV_CUDEV_EXPR_UNARY_FUNC(acos)
+CV_CUDEV_EXPR_UNARY_FUNC(atan)
+CV_CUDEV_EXPR_UNARY_FUNC(sinh)
+CV_CUDEV_EXPR_UNARY_FUNC(cosh)
+CV_CUDEV_EXPR_UNARY_FUNC(tanh)
+CV_CUDEV_EXPR_UNARY_FUNC(asinh)
+CV_CUDEV_EXPR_UNARY_FUNC(acosh)
+CV_CUDEV_EXPR_UNARY_FUNC(atanh)
+
+#undef CV_CUDEV_EXPR_UNARY_FUNC
+
+template <class SrcPtr>
+__host__ Expr<UnaryTransformPtrSz<typename PtrTraits<SrcPtr>::ptr_type, Binder2nd<pow_func<typename PtrTraits<SrcPtr>::value_type> > > >
+pow_(const SrcPtr& src, float power)
+{
+    return makeExpr(transformPtr(src, bind2nd(pow_func<typename PtrTraits<SrcPtr>::value_type>(), power)));
+}
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/expr/unary_op.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_EXPR_UNARY_OP_HPP__
+#define __OPENCV_CUDEV_EXPR_UNARY_OP_HPP__
+
+#include "../common.hpp"
+#include "../ptr2d/traits.hpp"
+#include "../ptr2d/transform.hpp"
+#include "../ptr2d/gpumat.hpp"
+#include "../ptr2d/texture.hpp"
+#include "../ptr2d/glob.hpp"
+#include "../functional/functional.hpp"
+#include "expr.hpp"
+
+namespace cv { namespace cudev {
+
+#define CV_CUDEV_EXPR_UNOP_INST(op, functor) \
+    template <typename T> \
+    __host__ Expr<UnaryTransformPtrSz<typename PtrTraits<GpuMat_<T> >::ptr_type, functor<T> > > \
+    operator op(const GpuMat_<T>& src) \
+    { \
+        return makeExpr(transformPtr(src, functor<T>())); \
+    } \
+    template <typename T> \
+    __host__ Expr<UnaryTransformPtrSz<typename PtrTraits<GlobPtrSz<T> >::ptr_type, functor<T> > > \
+    operator op(const GlobPtrSz<T>& src) \
+    { \
+        return makeExpr(transformPtr(src, functor<T>())); \
+    } \
+    template <typename T> \
+    __host__ Expr<UnaryTransformPtrSz<typename PtrTraits<Texture<T> >::ptr_type, functor<T> > > \
+    operator op(const Texture<T>& src) \
+    { \
+        return makeExpr(transformPtr(src, functor<T>())); \
+    } \
+    template <class Body> \
+    __host__ Expr<UnaryTransformPtrSz<typename PtrTraits<Body>::ptr_type, functor<typename Body::value_type> > > \
+    operator op(const Expr<Body>& src) \
+    { \
+        return makeExpr(transformPtr(src.body, functor<typename Body::value_type>())); \
+    }
+
+CV_CUDEV_EXPR_UNOP_INST(-, negate)
+CV_CUDEV_EXPR_UNOP_INST(!, logical_not)
+CV_CUDEV_EXPR_UNOP_INST(~, bit_not)
+
+#undef CV_CUDEV_EXPR_UNOP_INST
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/expr/warping.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_EXPR_WARPING_HPP__
+#define __OPENCV_CUDEV_EXPR_WARPING_HPP__
+
+#include "../common.hpp"
+#include "../ptr2d/traits.hpp"
+#include "../ptr2d/resize.hpp"
+#include "../ptr2d/remap.hpp"
+#include "../ptr2d/warping.hpp"
+#include "../grid/pyramids.hpp"
+#include "../grid/transpose.hpp"
+#include "expr.hpp"
+
+namespace cv { namespace cudev {
+
+// resize
+
+template <class SrcPtr>
+__host__ Expr<ResizePtrSz<typename PtrTraits<SrcPtr>::ptr_type> >
+resize_(const SrcPtr& src, float fx, float fy)
+{
+    return makeExpr(resizePtr(src, fx, fy));
+}
+
+// remap
+
+template <class SrcPtr, class MapPtr>
+__host__ Expr<RemapPtr1Sz<typename PtrTraits<SrcPtr>::ptr_type, typename PtrTraits<MapPtr>::ptr_type> >
+remap_(const SrcPtr& src, const MapPtr& map)
+{
+    return makeExpr(remapPtr(src, map));
+}
+
+template <class SrcPtr, class MapXPtr, class MapYPtr>
+__host__ Expr<RemapPtr2Sz<typename PtrTraits<SrcPtr>::ptr_type, typename PtrTraits<MapXPtr>::ptr_type, typename PtrTraits<MapYPtr>::ptr_type> >
+remap_(const SrcPtr& src, const MapXPtr& mapx, const MapYPtr& mapy)
+{
+    return makeExpr(remapPtr(src, mapx, mapy));
+}
+
+// warpAffine
+
+template <class SrcPtr>
+__host__ Expr<RemapPtr1Sz<typename PtrTraits<SrcPtr>::ptr_type, AffineMapPtr> >
+warpAffine_(const SrcPtr& src, Size dstSize, const GpuMat_<float>& warpMat)
+{
+    return makeExpr(warpAffinePtr(src, dstSize, warpMat));
+}
+
+// warpPerspective
+
+template <class SrcPtr>
+__host__ Expr<RemapPtr1Sz<typename PtrTraits<SrcPtr>::ptr_type, PerspectiveMapPtr> >
+warpPerspective_(const SrcPtr& src, Size dstSize, const GpuMat_<float>& warpMat)
+{
+    return makeExpr(warpPerspectivePtr(src, dstSize, warpMat));
+}
+
+// pyrDown
+
+template <class SrcPtr> struct PyrDownBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridPyrDown(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<PyrDownBody<SrcPtr> >
+pyrDown_(const SrcPtr& src)
+{
+    PyrDownBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// pyrUp
+
+template <class SrcPtr> struct PyrUpBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridPyrUp(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<PyrUpBody<SrcPtr> >
+pyrUp_(const SrcPtr& src)
+{
+    PyrUpBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+// transpose
+
+template <class SrcPtr> struct TransposeBody
+{
+    SrcPtr src;
+
+    template <typename T>
+    __host__ void assignTo(GpuMat_<T>& dst, Stream& stream = Stream::Null()) const
+    {
+        gridTranspose(src, dst, stream);
+    }
+};
+
+template <class SrcPtr>
+__host__ Expr<TransposeBody<SrcPtr> >
+transpose_(const SrcPtr& src)
+{
+    TransposeBody<SrcPtr> body;
+    body.src = src;
+    return makeExpr(body);
+}
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/functional/tuple_adapter.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_FUNCTIONAL_TUPLE_ADAPTER_HPP__
+#define __OPENCV_CUDEV_FUNCTIONAL_TUPLE_ADAPTER_HPP__
+
+#include "../common.hpp"
+#include "../util/tuple.hpp"
+
+namespace cv { namespace cudev {
+
+template <class Op, int n> struct UnaryTupleAdapter
+{
+    typedef typename Op::result_type result_type;
+
+    Op op;
+
+    template <class Tuple>
+    __device__ __forceinline__ typename Op::result_type operator ()(const Tuple& t) const
+    {
+        return op(get<n>(t));
+    }
+};
+
+template <int n, class Op>
+__host__ __device__ UnaryTupleAdapter<Op, n> unaryTupleAdapter(const Op& op)
+{
+    UnaryTupleAdapter<Op, n> a;
+    a.op = op;
+    return a;
+}
+
+template <class Op, int n0, int n1> struct BinaryTupleAdapter
+{
+    typedef typename Op::result_type result_type;
+
+    Op op;
+
+    template <class Tuple>
+    __device__ __forceinline__ typename Op::result_type operator ()(const Tuple& t) const
+    {
+        return op(get<n0>(t), get<n1>(t));
+    }
+};
+
+template <int n0, int n1, class Op>
+__host__ __device__ BinaryTupleAdapter<Op, n0, n1> binaryTupleAdapter(const Op& op)
+{
+    BinaryTupleAdapter<Op, n0, n1> a;
+    a.op = op;
+    return a;
+}
+
+}}
+
+#endif

modules/cudev/include/opencv2/cudev/grid/integral.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.
+// Copyright (C) 2013, OpenCV Foundation, 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*/
+
+#pragma once
+
+#ifndef __OPENCV_CUDEV_GRID_INTEGRAL_HPP__
+#define __OPENCV_CUDEV_GRID_INTEGRAL_HPP__
+
+#include "../common.hpp"
+#include "../ptr2d/traits.hpp"
+#include "../ptr2d/gpumat.hpp"
+#include "detail/integral.hpp"
+
+namespace cv { namespace cudev {
+
+template <class SrcPtr, typename DstType>
+__host__ void gridIntegral(const SrcPtr& src, GpuMat_<DstType>& dst, Stream& stream = Stream::Null())
+{
+    const int rows = getRows(src);
+    const int cols = getCols(src);
+
+    dst.create(rows, cols);
+
+    integral_detail::integral(shrinkPtr(src), shrinkPtr(dst), rows, cols, StreamAccessor::getStream(stream));
+}
+
+}}
+
+#endif