From 732bd621bbea99c9969ce2a511f0628e75e924ac Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Mon, 22 Nov 2010 09:39:34 +0000 Subject: [PATCH] added masks support for bitwise operations on GPU --- modules/gpu/include/opencv2/gpu/gpu.hpp | 56 +++++------ modules/gpu/src/arithm.cpp | 131 ++++++++++++++++++------- modules/gpu/src/cuda/mathfunc.cu | 166 +++++++++++++++++++++----------- tests/gpu/src/bitwise_oper.cpp | 56 ++++++++++- 4 files changed, 284 insertions(+), 125 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 29be204..42e30b1 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/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: diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index df3a3e7..0c2da06 100644 --- a/modules/gpu/src/arithm.cpp +++ b/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 + 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) { - ::bitwise_not_caller(src, dst, 0); + 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) { - ::bitwise_not_caller(src, dst, StreamAccessor::getStream(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) { - ::bitwise_or_caller(src1, src2, dst, 0); + 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) { - ::bitwise_or_caller(src1, src2, dst, StreamAccessor::getStream(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) { - ::bitwise_and_caller(src1, src2, dst, 0); + 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) { - ::bitwise_and_caller(src1, src2, dst, StreamAccessor::getStream(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) { - ::bitwise_xor_caller(src1, src2, dst, 0); + 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) { - ::bitwise_xor_caller(src1, src2, dst, StreamAccessor::getStream(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) diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu index ba5eb5d..73b76a8 100644 --- a/modules/gpu/src/cuda/mathfunc.cu +++ b/modules/gpu/src/cuda/mathfunc.cu @@ -243,100 +243,154 @@ namespace cv { namespace gpu { namespace mathfunc ////////////////////////////////////////////////////////////////////////////// // Per-element bit-wise logical matrix operations - - __global__ void bitwise_not_kernel(int cols, int rows, const PtrStep src, PtrStep dst) + struct Mask8U { - const int x = blockDim.x * blockIdx.x + threadIdx.x; - const int y = blockDim.y * blockIdx.y + threadIdx.y; - - if (x < cols && y < rows) - { - dst.ptr(y)[x] = ~src.ptr(y)[x]; - } - } - + 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; } }; - void bitwise_not_caller(const DevMem2D src, int elemSize, PtrStep dst, cudaStream_t stream) - { - dim3 threads(16, 16, 1); - dim3 grid(divUp(src.cols * elemSize, threads.x), divUp(src.rows, threads.y), 1); + // Unary operations - bitwise_not_kernel<<>>(src.cols * elemSize, src.rows, src, dst); + enum { UN_OP_NOT }; - if (stream == 0) - cudaSafeCall(cudaThreadSynchronize()); - } + template + struct UnOp { __device__ T operator()(T lhs, T rhs); }; + template + struct UnOp{ __device__ T operator()(T x) { return ~x; } }; - __global__ void bitwise_or_kernel(int cols, int rows, const PtrStep src1, const PtrStep src2, PtrStep dst) + template + __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] = src1.ptr(y)[x] | src2.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_or_caller(int cols, int rows, const PtrStep src1, const PtrStep src2, int elemSize, PtrStep dst, cudaStream_t stream) + template + 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(cols * elemSize, threads.x), divUp(rows, threads.y), 1); - - bitwise_or_kernel<<>>(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_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 2: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 3: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 4: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 6: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 8: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 12: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 16: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 24: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + case 32: bitwise_un_op<<>>(rows, cols, src, dst, UnOp(), mask); break; + } + if (stream == 0) cudaSafeCall(cudaThreadSynchronize()); } - - __global__ void bitwise_and_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(rows, cols, src, dst, elem_size, MaskTrue(), stream); + } - if (x < cols && y < rows) - { - dst.ptr(y)[x] = src1.ptr(y)[x] & src2.ptr(y)[x]; - } + void bitwise_not_caller(int rows, int cols,const PtrStep src, int elem_size, PtrStep dst, const PtrStep mask, cudaStream_t stream) + { + bitwise_un_op(rows, cols, src, dst, elem_size, Mask8U(mask), stream); } + // Binary operations - void bitwise_and_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_and_kernel<<>>(cols * elemSize, rows, src1, src2, dst); + template + struct BinOp { __device__ T operator()(T lhs, T rhs); }; - if (stream == 0) - cudaSafeCall(cudaThreadSynchronize()); - } + template + struct BinOp{ __device__ T operator()(T lhs, T rhs) { return lhs | rhs; } }; + template + struct BinOp{ __device__ T operator()(T lhs, T rhs) { return lhs & rhs; } }; + template + struct BinOp{ __device__ T operator()(T lhs, T rhs) { return lhs ^ rhs; } }; - __global__ void bitwise_xor_kernel(int cols, int rows, const PtrStep src1, const PtrStep src2, PtrStep dst) + template + __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]); } } + template + 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); + dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); + switch (elem_size) + { + case 1: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 2: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 3: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 4: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 6: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 8: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 12: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 16: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 24: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + case 32: bitwise_bin_op<<>>(rows, cols, src1, src2, dst, BinOp(), mask); break; + } + if (stream == 0) cudaSafeCall(cudaThreadSynchronize()); + } - void bitwise_xor_caller(int cols, int rows, 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 elem_size, PtrStep dst, cudaStream_t stream) { - dim3 threads(16, 16, 1); - dim3 grid(divUp(cols * elemSize, threads.x), divUp(rows, threads.y), 1); + bitwise_bin_op(rows, cols, src1, src2, dst, elem_size, MaskTrue(), stream); + } - bitwise_xor_kernel<<>>(cols * elemSize, rows, src1, src2, dst); + 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) + { + bitwise_bin_op(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream); + } - if (stream == 0) - cudaSafeCall(cudaThreadSynchronize()); + 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(rows, cols, src1, src2, dst, elem_size, MaskTrue(), 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) + { + bitwise_bin_op(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream); + } + + 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(rows, cols, src1, src2, dst, elem_size, MaskTrue(), stream); } + + 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(rows, cols, src1, src2, dst, elem_size, Mask8U(mask), stream); + } }}} diff --git a/tests/gpu/src/bitwise_oper.cpp b/tests/gpu/src/bitwise_oper.cpp index a3fcef2..7ad90ca 100644 --- a/tests/gpu/src/bitwise_oper.cpp +++ b/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; -- 2.7.4