added masks support for bitwise operations on GPU

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -474,53 +474,53 @@ namespace cv
         //! computes magnitude of each (x(i), y(i)) vector
         //! supports only floating-point source
         CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude);
-        //! Async version
+        //! async version
         CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, const Stream& stream);
 
         //! computes squared magnitude of each (x(i), y(i)) vector
         //! supports only floating-point source
         CV_EXPORTS void magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& magnitude);
-        //! Async version
+        //! async version
         CV_EXPORTS void magnitudeSqr(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, const Stream& stream);
 
         //! computes angle (angle(i)) of each (x(i), y(i)) vector
         //! supports only floating-point source
         CV_EXPORTS void phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees = false);
-        //! Async version
+        //! async version
         CV_EXPORTS void phase(const GpuMat& x, const GpuMat& y, GpuMat& angle, bool angleInDegrees, const Stream& stream);
 
         //! converts Cartesian coordinates to polar
         //! supports only floating-point source
         CV_EXPORTS void cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, GpuMat& angle, bool angleInDegrees = false);
-        //! Async version
+        //! async version
         CV_EXPORTS void cartToPolar(const GpuMat& x, const GpuMat& y, GpuMat& magnitude, GpuMat& angle, bool angleInDegrees, const Stream& stream);
 
         //! converts polar coordinates to Cartesian
         //! supports only floating-point source
         CV_EXPORTS void polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees = false);
-        //! Async version
+        //! async version
         CV_EXPORTS void polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& x, GpuMat& y, bool angleInDegrees, const Stream& stream);
 
 
-        //! Perfroms per-elements bit-wise inversion 
-        CV_EXPORTS void bitwise_not(const GpuMat& src, GpuMat& dst);
-        //! Async version
-        CV_EXPORTS void bitwise_not(const GpuMat& src, GpuMat& dst, const Stream& stream);
+        //! perfroms per-elements bit-wise inversion 
+        CV_EXPORTS void bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask=GpuMat());
+        //! async version
+        CV_EXPORTS void bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, const Stream& stream);
 
-        //! Calculates per-element bit-wise disjunction of two arrays
-        CV_EXPORTS void bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst);
-        //! Async version
-        CV_EXPORTS void bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream);
+        //! calculates per-element bit-wise disjunction of two arrays
+        CV_EXPORTS void bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask=GpuMat());
+        //! async version
+        CV_EXPORTS void bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, const Stream& stream);
 
-        //! Calculates per-element bit-wise conjunction of two arrays
-        CV_EXPORTS void bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst);
-        //! Async version
-        CV_EXPORTS void bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream);
+        //! calculates per-element bit-wise conjunction of two arrays
+        CV_EXPORTS void bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask=GpuMat());
+        //! async version
+        CV_EXPORTS void bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, const Stream& stream);
 
-        //! Calculates per-element bit-wise "exclusive or" operation
-        CV_EXPORTS void bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst);
-        //! Async version
-        CV_EXPORTS void bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream);
+        //! calculates per-element bit-wise "exclusive or" operation
+        CV_EXPORTS void bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask=GpuMat());
+        //! async version
+        CV_EXPORTS void bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, const Stream& stream);
 
         //! Logical operators
         CV_EXPORTS GpuMat operator ~ (const GpuMat& src);
@@ -551,7 +551,7 @@ namespace cv
         //! Supported types of input disparity: CV_8U, CV_16S.
         //! Output disparity has CV_8UC4 type in BGRA format (alpha = 255).
         CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndisp);
-        //! Async version
+        //! async version
         CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndisp, const Stream& stream);
 
         //! Reprojects disparity image to 3D space. 
@@ -560,12 +560,12 @@ namespace cv
         //! Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map.
         //! Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify.
         CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q);
-        //! Async version
+        //! async version
         CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const Stream& stream);
 
         //! converts image from one color space to another
         CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0);
-        //! Async version
+        //! async version
         CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream);
 
         //! applies fixed threshold to the image. 
@@ -821,7 +821,7 @@ namespace cv
             //! Output disparity has CV_8U type.
             void operator() ( const GpuMat& left, const GpuMat& right, GpuMat& disparity);
 
-            //! Async version
+            //! async version
             void operator() ( const GpuMat& left, const GpuMat& right, GpuMat& disparity, const Stream & stream);
 
             //! Some heuristics that tries to estmate
@@ -876,7 +876,7 @@ namespace cv
             //! if disparity is empty output type will be CV_16S else output type will be disparity.type().
             void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity);
 
-            //! Async version
+            //! async version
             void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream);
 
 
@@ -935,7 +935,7 @@ namespace cv
             //! if disparity is empty output type will be CV_16S else output type will be disparity.type().
             void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity);
 
-            //! Async version
+            //! async version
             void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream);
 
             int ndisp;
@@ -991,7 +991,7 @@ namespace cv
             //! disparity must have CV_8U or CV_16S type, image must have CV_8UC1 or CV_8UC3 type.
             void operator()(const GpuMat& disparity, const GpuMat& image, GpuMat& dst);
 
-            //! Async version
+            //! async version
             void operator()(const GpuMat& disparity, const GpuMat& image, GpuMat& dst, Stream& stream);
 
         private:

modules/gpu/src/arithm.cpp

@@ -81,14 +81,14 @@ void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool)
 void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, const Stream&) { throw_nogpu(); }
 void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool) { throw_nogpu(); }
 void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, const Stream&) { throw_nogpu(); }
-void cv::gpu::bitwise_not(const GpuMat&, GpuMat&) { throw_nogpu(); }
-void cv::gpu::bitwise_not(const GpuMat&, GpuMat&, const Stream& stream) { throw_nogpu(); }
-void cv::gpu::bitwise_or(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }
-void cv::gpu::bitwise_or(const GpuMat&, const GpuMat&, GpuMat&, const Stream& stream) { throw_nogpu(); }
-void cv::gpu::bitwise_and(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }
-void cv::gpu::bitwise_and(const GpuMat&, const GpuMat&, GpuMat&, const Stream& stream) { throw_nogpu(); }
-void cv::gpu::bitwise_xor(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }
-void cv::gpu::bitwise_xor(const GpuMat&, const GpuMat&, GpuMat&, const Stream& stream) { throw_nogpu(); }
+void cv::gpu::bitwise_not(const GpuMat&, GpuMat&, const GpuMat&) { throw_nogpu(); }
+void cv::gpu::bitwise_not(const GpuMat&, GpuMat&, const GpuMat&, const Stream&) { throw_nogpu(); }
+void cv::gpu::bitwise_or(const GpuMat&, const GpuMat&, GpuMat&, const GpuMat&) { throw_nogpu(); }
+void cv::gpu::bitwise_or(const GpuMat&, const GpuMat&, GpuMat&, const GpuMat&, const Stream&) { throw_nogpu(); }
+void cv::gpu::bitwise_and(const GpuMat&, const GpuMat&, GpuMat&, const GpuMat&) { throw_nogpu(); }
+void cv::gpu::bitwise_and(const GpuMat&, const GpuMat&, GpuMat&, const GpuMat&, const Stream&) { throw_nogpu(); }
+void cv::gpu::bitwise_xor(const GpuMat&, const GpuMat&, GpuMat&, const GpuMat&) { throw_nogpu(); }
+void cv::gpu::bitwise_xor(const GpuMat&, const GpuMat&, GpuMat&, const GpuMat&, const Stream&) { throw_nogpu(); }
 cv::gpu::GpuMat cv::gpu::operator ~ (const GpuMat&) { throw_nogpu(); return GpuMat(); }
 cv::gpu::GpuMat cv::gpu::operator | (const GpuMat&, const GpuMat&) { throw_nogpu(); return GpuMat(); }
 cv::gpu::GpuMat cv::gpu::operator & (const GpuMat&, const GpuMat&) { throw_nogpu(); return GpuMat(); }
@@ -873,10 +873,18 @@ void cv::gpu::polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat&
 
 namespace cv { namespace gpu { namespace mathfunc 
 {
-    void bitwise_not_caller(const DevMem2D src, int elemSize, PtrStep dst, cudaStream_t stream);
-    void bitwise_or_caller(int cols, int rows, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream);
-    void bitwise_and_caller(int cols, int rows, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream);
-    void bitwise_xor_caller(int cols, int rows, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream);
+    void bitwise_not_caller(int rows, int cols, const PtrStep src, int elemSize, PtrStep dst, cudaStream_t stream);
+    void bitwise_not_caller(int rows, int cols, const PtrStep src, int elemSize, PtrStep dst, const PtrStep mask, cudaStream_t stream);
+    void bitwise_or_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream);
+    void bitwise_or_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, const PtrStep mask, cudaStream_t stream);
+    void bitwise_and_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream);
+    void bitwise_and_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, const PtrStep mask, cudaStream_t stream);
+    void bitwise_xor_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream);
+    void bitwise_xor_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, const PtrStep mask, cudaStream_t stream);
+
+
+    template <int opid, typename Mask>
+    void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, int elem_size, Mask mask, cudaStream_t stream);
 }}}
 
 namespace 
@@ -884,75 +892,126 @@ namespace
     void bitwise_not_caller(const GpuMat& src, GpuMat& dst, cudaStream_t stream)
     {
         dst.create(src.size(), src.type());
-        mathfunc::bitwise_not_caller(src, src.elemSize(), dst, stream);
+        mathfunc::bitwise_not_caller(src.rows, src.cols, src, src.elemSize(), dst, stream);
+    }
+
+    void bitwise_not_caller(const GpuMat& src, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)
+    {
+        CV_Assert(mask.type() == CV_8U && mask.size() == src.size());
+        dst.create(src.size(), src.type());
+        mathfunc::bitwise_not_caller(src.rows, src.cols, src, src.elemSize(), dst, mask, stream);
     }
 
     void bitwise_or_caller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream)
     {
-        CV_Assert(src1.size() == src2.size());
-        CV_Assert(src1.type() == src2.type());
+        CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());
+        dst.create(src1.size(), src1.type());
+        mathfunc::bitwise_or_caller(dst.rows, dst.cols, src1, src2, dst.elemSize(), dst, stream);
+    }
 
+    void bitwise_or_caller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)
+    {
+        CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());
+        CV_Assert(mask.type() == CV_8U && mask.size() == src1.size());
         dst.create(src1.size(), src1.type());
-        mathfunc::bitwise_or_caller(dst.cols, dst.rows, src1, src2, dst.elemSize(), dst, stream);
+        mathfunc::bitwise_or_caller(dst.rows, dst.cols, src1, src2, dst.elemSize(), dst, mask, stream);
     }
 
     void bitwise_and_caller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream)
     {
-        CV_Assert(src1.size() == src2.size());
-        CV_Assert(src1.type() == src2.type());
+        CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());
+        dst.create(src1.size(), src1.type());
+        mathfunc::bitwise_and_caller(dst.rows, dst.cols, src1, src2, dst.elemSize(), dst, stream);
+    }
 
+    void bitwise_and_caller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)
+    {
+        CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());
+        CV_Assert(mask.type() == CV_8U && mask.size() == src1.size());
         dst.create(src1.size(), src1.type());
-        mathfunc::bitwise_and_caller(dst.cols, dst.rows, src1, src2, dst.elemSize(), dst, stream);
+        mathfunc::bitwise_and_caller(dst.rows, dst.cols, src1, src2, dst.elemSize(), dst, mask, stream);
     }
 
     void bitwise_xor_caller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream)
     {
         CV_Assert(src1.size() == src2.size());
         CV_Assert(src1.type() == src2.type());
+        dst.create(src1.size(), src1.type());
+        mathfunc::bitwise_xor_caller(dst.rows, dst.cols, src1, src2, dst.elemSize(), dst, stream);
+    }
 
+    void bitwise_xor_caller(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, cudaStream_t stream)
+    {
+        CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());
+        CV_Assert(mask.type() == CV_8U && mask.size() == src1.size());
         dst.create(src1.size(), src1.type());
-        mathfunc::bitwise_xor_caller(dst.cols, dst.rows, src1, src2, dst.elemSize(), dst, stream);
+        mathfunc::bitwise_xor_caller(dst.rows, dst.cols, src1, src2, dst.elemSize(), dst, mask, stream);
     }
 }
 
-void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst)
+void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask)
 {
+    if (mask.empty()) 
         ::bitwise_not_caller(src, dst, 0);
+    else
+        ::bitwise_not_caller(src, dst, mask, 0);
 }
 
-void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const Stream& stream)
+void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const GpuMat& mask, const Stream& stream)
 {
+    if (mask.empty())
         ::bitwise_not_caller(src, dst, StreamAccessor::getStream(stream));
+    else
+        ::bitwise_not_caller(src, dst, mask, StreamAccessor::getStream(stream));
 }
 
-void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst)
+void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask)
 {
+    if (mask.empty())
         ::bitwise_or_caller(src1, src2, dst, 0);
+    else
+        ::bitwise_or_caller(src1, src2, dst, mask, 0);
 }
 
-void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream)
+void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, const Stream& stream)
 {
+    if (mask.empty())
         ::bitwise_or_caller(src1, src2, dst, StreamAccessor::getStream(stream));
+    else
+        ::bitwise_or_caller(src1, src2, dst, mask, StreamAccessor::getStream(stream));
 }
 
-void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst)
+void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask)
 {
+    if (mask.empty())
         ::bitwise_and_caller(src1, src2, dst, 0);
+    else
+        ::bitwise_and_caller(src1, src2, dst, mask, 0);
 }
 
-void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream)
+void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, const Stream& stream)
 {
+    if (mask.empty())
         ::bitwise_and_caller(src1, src2, dst, StreamAccessor::getStream(stream));
+    else
+        ::bitwise_and_caller(src1, src2, dst, mask, StreamAccessor::getStream(stream));
 }
 
-void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst)
+void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask)
 {
+    if (mask.empty())
         ::bitwise_xor_caller(src1, src2, dst, 0);
+    else
+        ::bitwise_xor_caller(src1, src2, dst, mask, 0);
 }
 
-void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream)
+void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, const Stream& stream)
 {
+    if (mask.empty())
         ::bitwise_xor_caller(src1, src2, dst, StreamAccessor::getStream(stream));
+    else
+        ::bitwise_xor_caller(src1, src2, dst, mask, StreamAccessor::getStream(stream));
+
 }
 
 cv::gpu::GpuMat cv::gpu::operator ~ (const GpuMat& src)

modules/gpu/src/cuda/mathfunc.cu

@@ -243,100 +243,154 @@ namespace cv { namespace gpu { namespace mathfunc
 //////////////////////////////////////////////////////////////////////////////
 // Per-element bit-wise logical matrix operations
 
+    struct Mask8U
+    {
+        explicit Mask8U(PtrStep mask): mask(mask) {}
+        __device__ bool operator()(int y, int x) { return mask.ptr(y)[x]; }
+        PtrStep mask;
+    };
+    struct MaskTrue { __device__ bool operator()(int y, int x) { return true; } };
+
+    // Unary operations
+
+    enum { UN_OP_NOT };
 
-    __global__ void bitwise_not_kernel(int cols, int rows, const PtrStep src, PtrStep dst)
+    template <typename T, int opid>
+    struct UnOp { __device__ T operator()(T lhs, T rhs); };
+
+    template <typename T>
+    struct UnOp<T, UN_OP_NOT>{ __device__ T operator()(T x) { return ~x; } };
+
+    template <typename T, int cn, typename UnOp, typename Mask>
+    __global__ void bitwise_un_op(int rows, int cols, const PtrStep src, PtrStep dst, UnOp op, Mask mask)
     {
         const int x = blockDim.x * blockIdx.x + threadIdx.x;
         const int y = blockDim.y * blockIdx.y + threadIdx.y;
 
-        if (x < cols && y < rows)
+        if (x < cols && y < rows && mask(y, x)) 
         {
-            dst.ptr(y)[x] = ~src.ptr(y)[x];
+            T* dsty = (T*)dst.ptr(y);
+            const T* srcy = (const T*)src.ptr(y);
+
+            #pragma unroll
+            for (int i = 0; i < cn; ++i)
+                dsty[cn * x + i] = op(srcy[cn * x + i]);
         }
     }
 
-
-    void bitwise_not_caller(const DevMem2D src, int elemSize, PtrStep dst, cudaStream_t stream)
+    template <int opid, typename Mask>
+    void bitwise_un_op(int rows, int cols, const PtrStep src, PtrStep dst, int elem_size, Mask mask, cudaStream_t stream)
     {
-        dim3 threads(16, 16, 1);
-        dim3 grid(divUp(src.cols * elemSize, threads.x), divUp(src.rows, threads.y), 1);
-
-        bitwise_not_kernel<<<grid, threads, 0, stream>>>(src.cols * elemSize, src.rows, src, dst);
-
-        if (stream == 0)
-            cudaSafeCall(cudaThreadSynchronize());
+        dim3 threads(16, 16);
+        dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));
+        switch (elem_size)
+        {
+        case 1: bitwise_un_op<unsigned char, 1><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned char, opid>(), mask); break;
+        case 2: bitwise_un_op<unsigned short, 1><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned short, opid>(), mask); break;
+        case 3: bitwise_un_op<unsigned char, 3><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned char, opid>(), mask); break;
+        case 4: bitwise_un_op<unsigned int, 1><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned int, opid>(), mask); break;
+        case 6: bitwise_un_op<unsigned short, 3><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned short, opid>(), mask); break;
+        case 8: bitwise_un_op<unsigned int, 2><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned int, opid>(), mask); break;       
+        case 12: bitwise_un_op<unsigned int, 3><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned int, opid>(), mask); break;
+        case 16: bitwise_un_op<unsigned int, 4><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned int, opid>(), mask); break;
+        case 24: bitwise_un_op<unsigned int, 6><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned int, opid>(), mask); break;
+        case 32: bitwise_un_op<unsigned int, 8><<<grid, threads>>>(rows, cols, src, dst, UnOp<unsigned int, opid>(), mask); break;
+        }
+        if (stream == 0) cudaSafeCall(cudaThreadSynchronize());        
     }
 
-
-    __global__ void bitwise_or_kernel(int cols, int rows, const PtrStep src1, const PtrStep src2, PtrStep dst)
+    void bitwise_not_caller(int rows, int cols,const PtrStep src, int elem_size, PtrStep dst, cudaStream_t stream)
     {
-        const int x = blockDim.x * blockIdx.x + threadIdx.x;
-        const int y = blockDim.y * blockIdx.y + threadIdx.y;
+        bitwise_un_op<UN_OP_NOT>(rows, cols, src, dst, elem_size, MaskTrue(), stream);
+    }
 
-        if (x < cols && y < rows)
+    void bitwise_not_caller(int rows, int cols,const PtrStep src, int elem_size, PtrStep dst, const PtrStep mask, cudaStream_t stream)
     {
-            dst.ptr(y)[x] = src1.ptr(y)[x] | src2.ptr(y)[x];
-        }
+        bitwise_un_op<UN_OP_NOT>(rows, cols, src, dst, elem_size, Mask8U(mask), stream);
     }
 
+    // Binary operations
 
-    void bitwise_or_caller(int cols, int rows, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream)
-    {
-        dim3 threads(16, 16, 1);
-        dim3 grid(divUp(cols * elemSize, threads.x), divUp(rows, threads.y), 1);
+    enum { BIN_OP_OR, BIN_OP_AND, BIN_OP_XOR };
 
-        bitwise_or_kernel<<<grid, threads, 0, stream>>>(cols * elemSize, rows, src1, src2, dst);
+    template <typename T, int opid>
+    struct BinOp { __device__ T operator()(T lhs, T rhs); };
 
-        if (stream == 0)
-            cudaSafeCall(cudaThreadSynchronize());
-    }
+    template <typename T>
+    struct BinOp<T, BIN_OP_OR>{ __device__ T operator()(T lhs, T rhs) { return lhs | rhs; } };
 
+    template <typename T>
+    struct BinOp<T, BIN_OP_AND>{ __device__ T operator()(T lhs, T rhs) { return lhs & rhs; } };
 
-    __global__ void bitwise_and_kernel(int cols, int rows, const PtrStep src1, const PtrStep src2, PtrStep dst)
+    template <typename T>
+    struct BinOp<T, BIN_OP_XOR>{ __device__ T operator()(T lhs, T rhs) { return lhs ^ rhs; } };
+
+    template <typename T, int cn, typename BinOp, typename Mask>
+    __global__ void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, BinOp op, Mask mask)
     {
         const int x = blockDim.x * blockIdx.x + threadIdx.x;
         const int y = blockDim.y * blockIdx.y + threadIdx.y;
 
-        if (x < cols && y < rows)
+        if (x < cols && y < rows && mask(y, x)) 
         {
-            dst.ptr(y)[x] = src1.ptr(y)[x] & src2.ptr(y)[x];
+            T* dsty = (T*)dst.ptr(y);
+            const T* src1y = (const T*)src1.ptr(y);
+            const T* src2y = (const T*)src2.ptr(y);
+
+            #pragma unroll
+            for (int i = 0; i < cn; ++i)
+                dsty[cn * x + i] = op(src1y[cn * x + i], src2y[cn * x + i]);
         }
     }
 
-
-    void bitwise_and_caller(int cols, int rows, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream)
+    template <int opid, typename Mask>
+    void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, int elem_size, Mask mask, cudaStream_t stream)
     {
-        dim3 threads(16, 16, 1);
-        dim3 grid(divUp(cols * elemSize, threads.x), divUp(rows, threads.y), 1);
-
-        bitwise_and_kernel<<<grid, threads, 0, stream>>>(cols * elemSize, rows, src1, src2, dst);
-
-        if (stream == 0)
-            cudaSafeCall(cudaThreadSynchronize());
+        dim3 threads(16, 16);
+        dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));
+        switch (elem_size)
+        {
+        case 1: bitwise_bin_op<unsigned char, 1><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned char, opid>(), mask); break;
+        case 2: bitwise_bin_op<unsigned short, 1><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned short, opid>(), mask); break;
+        case 3: bitwise_bin_op<unsigned char, 3><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned char, opid>(), mask); break;
+        case 4: bitwise_bin_op<unsigned int, 1><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned int, opid>(), mask); break;
+        case 6: bitwise_bin_op<unsigned short, 3><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned short, opid>(), mask); break;
+        case 8: bitwise_bin_op<unsigned int, 2><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned int, opid>(), mask); break;       
+        case 12: bitwise_bin_op<unsigned int, 3><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned int, opid>(), mask); break;
+        case 16: bitwise_bin_op<unsigned int, 4><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned int, opid>(), mask); break;
+        case 24: bitwise_bin_op<unsigned int, 6><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned int, opid>(), mask); break;
+        case 32: bitwise_bin_op<unsigned int, 8><<<grid, threads>>>(rows, cols, src1, src2, dst, BinOp<unsigned int, opid>(), mask); break;
+        }
+        if (stream == 0) cudaSafeCall(cudaThreadSynchronize());        
     }
 
-
-
-    __global__ void bitwise_xor_kernel(int cols, int rows, const PtrStep src1, const PtrStep src2, PtrStep dst)
+    void bitwise_or_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, cudaStream_t stream)
     {
-        const int x = blockDim.x * blockIdx.x + threadIdx.x;
-        const int y = blockDim.y * blockIdx.y + threadIdx.y;
+        bitwise_bin_op<BIN_OP_OR>(rows, cols, src1, src2, dst, elem_size, MaskTrue(), stream);
+    }
 
-        if (x < cols && y < rows)
+    void bitwise_or_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, const PtrStep mask, cudaStream_t stream)
     {
-            dst.ptr(y)[x] = src1.ptr(y)[x] ^ src2.ptr(y)[x];
-        }
+        bitwise_bin_op<BIN_OP_OR>(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream);
     }
 
+    void bitwise_and_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, cudaStream_t stream)
+    {
+        bitwise_bin_op<BIN_OP_AND>(rows, cols, src1, src2, dst, elem_size, MaskTrue(), stream);
+    }
 
-    void bitwise_xor_caller(int cols, int rows, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream)
+    void bitwise_and_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, const PtrStep mask, cudaStream_t stream)
     {
-        dim3 threads(16, 16, 1);
-        dim3 grid(divUp(cols * elemSize, threads.x), divUp(rows, threads.y), 1);
+        bitwise_bin_op<BIN_OP_AND>(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream);
+    }
 
-        bitwise_xor_kernel<<<grid, threads, 0, stream>>>(cols * elemSize, rows, src1, src2, dst);
+    void bitwise_xor_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, cudaStream_t stream)
+    {
+        bitwise_bin_op<BIN_OP_XOR>(rows, cols, src1, src2, dst, elem_size, MaskTrue(), stream);
+    }
 
-        if (stream == 0)
-            cudaSafeCall(cudaThreadSynchronize());
+    void bitwise_xor_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, const PtrStep mask, cudaStream_t stream)
+    {
+        bitwise_bin_op<BIN_OP_XOR>(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream);
     }  
 }}}

tests/gpu/src/bitwise_oper.cpp

@@ -60,7 +60,7 @@ struct CV_GpuBitwiseTest: public CvTest
         int rows, cols;
         for (int depth = CV_8U; depth <= CV_64F; ++depth)
             for (int cn = 1; cn <= 4; ++cn)
-                for (int attempt = 0; attempt < 5; ++attempt)
+                for (int attempt = 0; attempt < 3; ++attempt)
                 {
                     rows = 1 + rand() % 100;
                     cols = 1 + rand() % 100;
@@ -83,7 +83,12 @@ struct CV_GpuBitwiseTest: public CvTest
         }
 
         Mat dst_gold = ~src;
-        gpu::GpuMat dst = ~gpu::GpuMat(src);
+
+        gpu::GpuMat mask(src.size(), CV_8U);
+        mask.setTo(Scalar(1));
+
+        gpu::GpuMat dst;
+        gpu::bitwise_not(gpu::GpuMat(src), dst, mask);
 
         CHECK(dst_gold.size() == dst.size(), CvTS::FAIL_INVALID_OUTPUT);
         CHECK(dst_gold.type() == dst.type(), CvTS::FAIL_INVALID_OUTPUT);        
@@ -112,10 +117,23 @@ struct CV_GpuBitwiseTest: public CvTest
 
         CHECK(dst_gold.size() == dst.size(), CvTS::FAIL_INVALID_OUTPUT);
         CHECK(dst_gold.type() == dst.type(), CvTS::FAIL_INVALID_OUTPUT);        
-
         Mat dsth(dst);
         for (int i = 0; i < dst_gold.rows; ++i)       
             CHECK(memcmp(dst_gold.ptr(i), dsth.ptr(i), dst_gold.cols * dst_gold.elemSize()) == 0, CvTS::FAIL_INVALID_OUTPUT)
+
+        Mat mask(src1.size(), CV_8U);
+        randu(mask, Scalar(0), Scalar(255));
+
+        Mat dst_gold2(dst_gold.size(), dst_gold.type()); dst_gold2.setTo(Scalar::all(0));
+        gpu::GpuMat dst2(dst.size(), dst.type()); dst2.setTo(Scalar::all(0));
+        bitwise_or(src1, src2, dst_gold2, mask);
+        gpu::bitwise_or(gpu::GpuMat(src1), gpu::GpuMat(src2), dst2, gpu::GpuMat(mask));
+
+        CHECK(dst_gold2.size() == dst2.size(), CvTS::FAIL_INVALID_OUTPUT);
+        CHECK(dst_gold2.type() == dst2.type(), CvTS::FAIL_INVALID_OUTPUT);        
+        dsth = dst2;
+        for (int i = 0; i < dst_gold.rows; ++i)       
+            CHECK(memcmp(dst_gold2.ptr(i), dsth.ptr(i), dst_gold2.cols * dst_gold2.elemSize()) == 0, CvTS::FAIL_INVALID_OUTPUT)
     }
 
     void test_bitwise_and(int rows, int cols, int type)
@@ -138,10 +156,24 @@ struct CV_GpuBitwiseTest: public CvTest
 
         CHECK(dst_gold.size() == dst.size(), CvTS::FAIL_INVALID_OUTPUT);
         CHECK(dst_gold.type() == dst.type(), CvTS::FAIL_INVALID_OUTPUT);        
-
         Mat dsth(dst);
         for (int i = 0; i < dst_gold.rows; ++i)       
             CHECK(memcmp(dst_gold.ptr(i), dsth.ptr(i), dst_gold.cols * dst_gold.elemSize()) == 0, CvTS::FAIL_INVALID_OUTPUT)
+
+
+        Mat mask(src1.size(), CV_8U);
+        randu(mask, Scalar(0), Scalar(255));
+
+        Mat dst_gold2(dst_gold.size(), dst_gold.type()); dst_gold2.setTo(Scalar::all(0));
+        gpu::GpuMat dst2(dst.size(), dst.type()); dst2.setTo(Scalar::all(0));
+        bitwise_and(src1, src2, dst_gold2, mask);
+        gpu::bitwise_and(gpu::GpuMat(src1), gpu::GpuMat(src2), dst2, gpu::GpuMat(mask));
+
+        CHECK(dst_gold2.size() == dst2.size(), CvTS::FAIL_INVALID_OUTPUT);
+        CHECK(dst_gold2.type() == dst2.type(), CvTS::FAIL_INVALID_OUTPUT);        
+        dsth = dst2;
+        for (int i = 0; i < dst_gold.rows; ++i)       
+            CHECK(memcmp(dst_gold2.ptr(i), dsth.ptr(i), dst_gold2.cols * dst_gold2.elemSize()) == 0, CvTS::FAIL_INVALID_OUTPUT)
     }
 
     void test_bitwise_xor(int rows, int cols, int type)
@@ -164,10 +196,24 @@ struct CV_GpuBitwiseTest: public CvTest
 
         CHECK(dst_gold.size() == dst.size(), CvTS::FAIL_INVALID_OUTPUT);
         CHECK(dst_gold.type() == dst.type(), CvTS::FAIL_INVALID_OUTPUT);        
-
         Mat dsth(dst);
         for (int i = 0; i < dst_gold.rows; ++i)       
             CHECK(memcmp(dst_gold.ptr(i), dsth.ptr(i), dst_gold.cols * dst_gold.elemSize()) == 0, CvTS::FAIL_INVALID_OUTPUT)
+
+
+        Mat mask(src1.size(), CV_8U);
+        randu(mask, Scalar(0), Scalar(255));
+
+        Mat dst_gold2(dst_gold.size(), dst_gold.type()); dst_gold2.setTo(Scalar::all(0));
+        gpu::GpuMat dst2(dst.size(), dst.type()); dst2.setTo(Scalar::all(0));
+        bitwise_xor(src1, src2, dst_gold2, mask);
+        gpu::bitwise_xor(gpu::GpuMat(src1), gpu::GpuMat(src2), dst2, gpu::GpuMat(mask));
+
+        CHECK(dst_gold2.size() == dst2.size(), CvTS::FAIL_INVALID_OUTPUT);
+        CHECK(dst_gold2.type() == dst2.type(), CvTS::FAIL_INVALID_OUTPUT);        
+        dsth = dst2;
+        for (int i = 0; i < dst_gold.rows; ++i)       
+            CHECK(memcmp(dst_gold2.ptr(i), dsth.ptr(i), dst_gold2.cols * dst_gold2.elemSize()) == 0, CvTS::FAIL_INVALID_OUTPUT)
     }
 } gpu_bitwise_test;