updated other gpu's bitwise operations

diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu
index 5fbb73a..2de34ec 100644 (file)
--- a/modules/gpu/src/cuda/mathfunc.cu
+++ b/modules/gpu/src/cuda/mathfunc.cu
@@ -291,16 +291,11 @@ namespace cv { namespace gpu { namespace mathfunc
     //------------------------------------------------------------------------\r
     // Unary operations\r
 \r
-    enum \r
-    { \r
-        UN_OP_NOT \r
-    };\r
-\r
+    enum { UN_OP_NOT };\r
 \r
     template <typename T, int opid>\r
     struct UnOp;\r
 \r
-\r
     template <typename T>\r
     struct UnOp<T, UN_OP_NOT>\r
     { \r
@@ -380,7 +375,8 @@ namespace cv { namespace gpu { namespace mathfunc
         dim3 grid(divUp(cols * elem_size, threads.x * sizeof(uint)), \r
                   divUp(rows, threads.y));\r
         bitwise_un_op<opid><<<grid, threads>>>(rows, cols * elem_size, src, dst);\r
-        if (stream == 0) cudaSafeCall(cudaThreadSynchronize());\r
+        if (stream == 0) \r
+            cudaSafeCall(cudaThreadSynchronize());\r
     }\r
 \r
 \r
@@ -422,7 +418,8 @@ namespace cv { namespace gpu { namespace mathfunc
             bitwise_un_op_two_loads<uint, 4, opid><<<grid, threads>>>(rows, cols, src, dst, mask); \r
             break;\r
         }\r
-        if (stream == 0) cudaSafeCall(cudaThreadSynchronize());\r
+        if (stream == 0) \r
+            cudaSafeCall(cudaThreadSynchronize());\r
     }\r
 \r
 \r
@@ -442,134 +439,201 @@ namespace cv { namespace gpu { namespace mathfunc
 \r
     enum { BIN_OP_OR, BIN_OP_AND, BIN_OP_XOR };\r
 \r
-\r
     template <typename T, int opid>\r
     struct BinOp;\r
 \r
-\r
     template <typename T>\r
     struct BinOp<T, BIN_OP_OR>\r
     { \r
-        static __device__ T call(T lhs, T rhs) \r
-        { \r
-            return lhs | rhs; \r
-        } \r
+        typedef typename TypeVec<T, 2>::vec_t Vec2;\r
+        typedef typename TypeVec<T, 3>::vec_t Vec3;\r
+        typedef typename TypeVec<T, 4>::vec_t Vec4;\r
+        static __device__ T call(T a, T b) { return a | b; } \r
+        static __device__ Vec2 call(Vec2 a, Vec2 b) { return VecTraits<Vec2>::make(a.x | b.x, a.y | b.y); } \r
+        static __device__ Vec3 call(Vec3 a, Vec3 b) { return VecTraits<Vec3>::make(a.x | b.x, a.y | b.y, a.z | b.z); } \r
+        static __device__ Vec4 call(Vec4 a, Vec4 b) { return VecTraits<Vec4>::make(a.x | b.x, a.y | b.y, a.z | b.z, a.w | b.w); } \r
     };\r
 \r
 \r
     template <typename T>\r
     struct BinOp<T, BIN_OP_AND>\r
     { \r
-        static __device__ T call(T lhs, T rhs) \r
-        { \r
-            return lhs & rhs; \r
-        } \r
+        typedef typename TypeVec<T, 2>::vec_t Vec2;\r
+        typedef typename TypeVec<T, 3>::vec_t Vec3;\r
+        typedef typename TypeVec<T, 4>::vec_t Vec4;\r
+        static __device__ T call(T a, T b) { return a & b; } \r
+        static __device__ Vec2 call(Vec2 a, Vec2 b) { return VecTraits<Vec2>::make(a.x & b.x, a.y & b.y); } \r
+        static __device__ Vec3 call(Vec3 a, Vec3 b) { return VecTraits<Vec3>::make(a.x & b.x, a.y & b.y, a.z & b.z); } \r
+        static __device__ Vec4 call(Vec4 a, Vec4 b) { return VecTraits<Vec4>::make(a.x & b.x, a.y & b.y, a.z & b.z, a.w & b.w); } \r
     };\r
 \r
     template <typename T>\r
     struct BinOp<T, BIN_OP_XOR>\r
     { \r
-        static __device__ T call(T lhs, T rhs) \r
-        { \r
-            return lhs ^ rhs; \r
-        } \r
+        typedef typename TypeVec<T, 2>::vec_t Vec2;\r
+        typedef typename TypeVec<T, 3>::vec_t Vec3;\r
+        typedef typename TypeVec<T, 4>::vec_t Vec4;\r
+        static __device__ T call(T a, T b) { return a ^ b; } \r
+        static __device__ Vec2 call(Vec2 a, Vec2 b) { return VecTraits<Vec2>::make(a.x ^ b.x, a.y ^ b.y); } \r
+        static __device__ Vec3 call(Vec3 a, Vec3 b) { return VecTraits<Vec3>::make(a.x ^ b.x, a.y ^ b.y, a.z ^ b.z); } \r
+        static __device__ Vec4 call(Vec4 a, Vec4 b) { return VecTraits<Vec4>::make(a.x ^ b.x, a.y ^ b.y, a.z ^ b.z, a.w ^ b.w); } \r
     };\r
 \r
 \r
-    template <typename T, int cn, typename BinOp, typename Mask>\r
-    __global__ void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, Mask mask)\r
+    template <int opid>\r
+    __global__ void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst)\r
     {\r
+        const int x = (blockDim.x * blockIdx.x + threadIdx.x) * 4;\r
+        const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+\r
+        if (y < rows) \r
+        {\r
+            uchar* dst_ptr = dst.ptr(y) + x;\r
+            const uchar* src1_ptr = src1.ptr(y) + x;\r
+            const uchar* src2_ptr = src2.ptr(y) + x;\r
+            if (x + sizeof(uint) - 1 < cols)\r
+            {\r
+                *(uint*)dst_ptr = BinOp<uint, opid>::call(*(uint*)src1_ptr, *(uint*)src2_ptr);\r
+            }\r
+            else\r
+            {\r
+                const uchar* src1_end = src1.ptr(y) + cols;\r
+                while (src1_ptr < src1_end)\r
+                {\r
+                    *dst_ptr++ = BinOp<uchar, opid>::call(*src1_ptr++, *src2_ptr++);\r
+                }\r
+            }\r
+        }\r
+    }\r
+\r
+\r
+    template <typename T, int cn, int opid>\r
+    __global__ void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, \r
+                                   PtrStep dst, const PtrStep mask)\r
+    {\r
+        typedef typename TypeVec<T, cn>::vec_t Type;\r
         const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
         const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
 \r
-        if (x < cols && y < rows && mask(y, x)) \r
+        if (x < cols && y < rows && mask.ptr(y)[x]) \r
         {\r
-            T* dsty = (T*)dst.ptr(y);\r
-            const T* src1y = (const T*)src1.ptr(y);\r
-            const T* src2y = (const T*)src2.ptr(y);\r
+            Type* dst_row = (Type*)dst.ptr(y);\r
+            const Type* src1_row = (const Type*)src1.ptr(y);\r
+            const Type* src2_row = (const Type*)src2.ptr(y);\r
+            dst_row[x] = BinOp<T, opid>::call(src1_row[x], src2_row[x]);\r
+        }\r
+    }\r
 \r
-            #pragma unroll\r
-            for (int i = 0; i < cn; ++i)\r
-                dsty[cn * x + i] = BinOp::call(src1y[cn * x + i], src2y[cn * x + i]);\r
+\r
+    template <typename T, int cn, int opid>\r
+    __global__ void bitwise_bin_op_two_loads(int rows, int cols, const PtrStep src1, const PtrStep src2, \r
+                                             PtrStep dst, const PtrStep mask)\r
+    {\r
+        typedef typename TypeVec<T, cn>::vec_t Type;\r
+        const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+        const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+\r
+        if (x < cols && y < rows && mask.ptr(y)[x]) \r
+        {\r
+            Type* dst_row = (Type*)dst.ptr(y);\r
+            const Type* src1_row = (const Type*)src1.ptr(y);\r
+            const Type* src2_row = (const Type*)src2.ptr(y);\r
+            dst_row[2 * x] = BinOp<T, opid>::call(src1_row[2 * x], src2_row[2 * x]);\r
+            dst_row[2 * x + 1] = BinOp<T, opid>::call(src1_row[2 * x + 1], src2_row[2 * x + 1]);\r
         }\r
     }\r
 \r
 \r
-    template <int opid, typename Mask>\r
-    void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, int elem_size, Mask mask, cudaStream_t stream)\r
+    template <int opid>\r
+    void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, \r
+                        int elem_size, cudaStream_t stream)\r
+    {\r
+        dim3 threads(16, 16);\r
+        dim3 grid(divUp(cols * elem_size, threads.x * sizeof(uint)), \r
+                  divUp(rows, threads.y));\r
+        bitwise_bin_op<opid><<<grid, threads>>>(rows, cols * elem_size, src1, src2, dst);\r
+        if (stream == 0) \r
+            cudaSafeCall(cudaThreadSynchronize());\r
+    }\r
+\r
+\r
+    template <int opid>\r
+    void bitwise_bin_op(int rows, int cols, const PtrStep src1, const PtrStep src2, PtrStep dst, \r
+                        int elem_size, const PtrStep mask, cudaStream_t stream)\r
     {\r
         dim3 threads(16, 16);\r
         dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
         switch (elem_size)\r
         {\r
         case 1: \r
-            bitwise_bin_op<uchar, 1, BinOp<uchar, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op<uchar, 1, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 2: \r
-            bitwise_bin_op<ushort, 1, BinOp<ushort, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op<ushort, 1, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 3: \r
-            bitwise_bin_op<uchar, 3, BinOp<uchar, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op<uchar, 3, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 4: \r
-            bitwise_bin_op<uint, 1, BinOp<uint, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op<uint, 1, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 6: \r
-            bitwise_bin_op<ushort, 3, BinOp<ushort, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op<ushort, 3, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 8: \r
-            bitwise_bin_op<uint, 2, BinOp<uint, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op<uint, 2, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 12: \r
-            bitwise_bin_op<uint, 3, BinOp<uint, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op<uint, 3, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 16: \r
-            bitwise_bin_op<uint, 4, BinOp<uint, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op<uint, 4, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 24: \r
-            bitwise_bin_op<uint, 6, BinOp<uint, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op_two_loads<uint, 3, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         case 32: \r
-            bitwise_bin_op<uint, 8, BinOp<uint, opid> ><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
+            bitwise_bin_op_two_loads<uint, 4, opid><<<grid, threads>>>(rows, cols, src1, src2, dst, mask); \r
             break;\r
         }\r
-        if (stream == 0) cudaSafeCall(cudaThreadSynchronize());        \r
+        if (stream == 0) \r
+            cudaSafeCall(cudaThreadSynchronize());\r
     }\r
 \r
 \r
     void bitwise_or_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, cudaStream_t stream)\r
     {\r
-        bitwise_bin_op<BIN_OP_OR>(rows, cols, src1, src2, dst, elem_size, MaskTrue(), stream);\r
+        bitwise_bin_op<BIN_OP_OR>(rows, cols, src1, src2, dst, elem_size, stream);\r
     }\r
 \r
 \r
     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)\r
     {\r
-        bitwise_bin_op<BIN_OP_OR>(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream);\r
+        bitwise_bin_op<BIN_OP_OR>(rows, cols, src1, src2, dst, elem_size, mask, stream);\r
     }\r
 \r
 \r
     void bitwise_and_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, cudaStream_t stream)\r
     {\r
-        bitwise_bin_op<BIN_OP_AND>(rows, cols, src1, src2, dst, elem_size, MaskTrue(), stream);\r
+        bitwise_bin_op<BIN_OP_AND>(rows, cols, src1, src2, dst, elem_size, stream);\r
     }\r
 \r
 \r
     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)\r
     {\r
-        bitwise_bin_op<BIN_OP_AND>(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream);\r
+        bitwise_bin_op<BIN_OP_AND>(rows, cols, src1, src2, dst, elem_size, mask, stream);\r
     }\r
 \r
 \r
     void bitwise_xor_caller(int rows, int cols, const PtrStep src1, const PtrStep src2, int elem_size, PtrStep dst, cudaStream_t stream)\r
     {\r
-        bitwise_bin_op<BIN_OP_XOR>(rows, cols, src1, src2, dst, elem_size, MaskTrue(), stream);\r
+        bitwise_bin_op<BIN_OP_XOR>(rows, cols, src1, src2, dst, elem_size, stream);\r
     }\r
 \r
 \r
     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)\r
     {\r
-        bitwise_bin_op<BIN_OP_XOR>(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream);\r
+        bitwise_bin_op<BIN_OP_XOR>(rows, cols, src1, src2, dst, elem_size, mask, stream);\r
     }  \r
 \r
 \r
@@ -2247,3 +2311,4 @@ namespace cv { namespace gpu { namespace mathfunc
 }}}\r
 \r
 \r
+\r