From faf4d0bc74160bfa246cca5d93e270cc7a9b528f Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Fri, 19 Nov 2010 10:19:35 +0000 Subject: [PATCH] added bitwise operations into gpu module --- modules/gpu/include/opencv2/gpu/gpu.hpp | 52 +++++++--- modules/gpu/src/arithm.cpp | 128 +++++++++++++++++++++++ modules/gpu/src/cuda/mathfunc.cu | 101 +++++++++++++++++++ tests/gpu/src/bitwise_oper.cpp | 173 ++++++++++++++++++++++++++++++++ 4 files changed, 442 insertions(+), 12 deletions(-) create mode 100644 tests/gpu/src/bitwise_oper.cpp diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index cc0b438..2af921f 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -474,33 +474,61 @@ 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); - //! Acync 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); - //! Acync 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); - //! Acync 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); - //! Acync 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); - //! Acync 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); + + //! 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 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 "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); + + //! Logical operators + CV_EXPORTS GpuMat operator ~ (const GpuMat& src); + CV_EXPORTS GpuMat operator | (const GpuMat& src1, const GpuMat& src2); + CV_EXPORTS GpuMat operator & (const GpuMat& src1, const GpuMat& src2); + CV_EXPORTS GpuMat operator ^ (const GpuMat& src1, const GpuMat& src2); + + ////////////////////////////// Image processing ////////////////////////////// //! DST[x,y] = SRC[xmap[x,y],ymap[x,y]] with bilinear interpolation. @@ -523,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); - //! Acync 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. @@ -532,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); - //! Acync 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); - //! Acync 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. @@ -793,7 +821,7 @@ namespace cv //! Output disparity has CV_8U type. void operator() ( const GpuMat& left, const GpuMat& right, GpuMat& disparity); - //! Acync version + //! Async version void operator() ( const GpuMat& left, const GpuMat& right, GpuMat& disparity, const Stream & stream); //! Some heuristics that tries to estmate @@ -848,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); - //! Acync version + //! Async version void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream); @@ -907,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); - //! Acync version + //! Async version void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream); int ndisp; @@ -963,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); - //! Acync 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 510ec11..df3a3e7 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -81,6 +81,18 @@ 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(); } +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(); } +cv::gpu::GpuMat cv::gpu::operator ^ (const GpuMat&, const GpuMat&) { throw_nogpu(); return GpuMat(); } #else /* !defined (HAVE_CUDA) */ @@ -856,4 +868,120 @@ void cv::gpu::polarToCart(const GpuMat& magnitude, const GpuMat& angle, GpuMat& ::polarToCart_caller(magnitude, angle, x, y, angleInDegrees, StreamAccessor::getStream(stream)); } +////////////////////////////////////////////////////////////////////////////// +// Per-element bit-wise logical matrix operations + +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); +}}} + +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); + } + + 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()); + + dst.create(src1.size(), src1.type()); + mathfunc::bitwise_or_caller(dst.cols, dst.rows, src1, src2, dst.elemSize(), dst, 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()); + + dst.create(src1.size(), src1.type()); + mathfunc::bitwise_and_caller(dst.cols, dst.rows, src1, src2, dst.elemSize(), dst, 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.cols, dst.rows, src1, src2, dst.elemSize(), dst, stream); + } +} + +void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst) +{ + ::bitwise_not_caller(src, dst, 0); +} + +void cv::gpu::bitwise_not(const GpuMat& src, GpuMat& dst, const Stream& stream) +{ + ::bitwise_not_caller(src, dst, StreamAccessor::getStream(stream)); +} + +void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) +{ + ::bitwise_or_caller(src1, src2, dst, 0); +} + +void cv::gpu::bitwise_or(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream) +{ + ::bitwise_or_caller(src1, src2, dst, StreamAccessor::getStream(stream)); +} + +void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) +{ + ::bitwise_and_caller(src1, src2, dst, 0); +} + +void cv::gpu::bitwise_and(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream) +{ + ::bitwise_and_caller(src1, src2, dst, StreamAccessor::getStream(stream)); +} + +void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) +{ + ::bitwise_xor_caller(src1, src2, dst, 0); +} + +void cv::gpu::bitwise_xor(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const Stream& stream) +{ + ::bitwise_xor_caller(src1, src2, dst, StreamAccessor::getStream(stream)); +} + +cv::gpu::GpuMat cv::gpu::operator ~ (const GpuMat& src) +{ + GpuMat dst; + bitwise_not(src, dst); + return dst; +} + +cv::gpu::GpuMat cv::gpu::operator | (const GpuMat& src1, const GpuMat& src2) +{ + GpuMat dst; + bitwise_or(src1, src2, dst); + return dst; +} + +cv::gpu::GpuMat cv::gpu::operator & (const GpuMat& src1, const GpuMat& src2) +{ + GpuMat dst; + bitwise_and(src1, src2, dst); + return dst; +} + +cv::gpu::GpuMat cv::gpu::operator ^ (const GpuMat& src1, const GpuMat& src2) +{ + GpuMat dst; + bitwise_xor(src1, src2, dst); + return dst; +} + + #endif /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu index 7fe87bf..ba5eb5d 100644 --- a/modules/gpu/src/cuda/mathfunc.cu +++ b/modules/gpu/src/cuda/mathfunc.cu @@ -238,4 +238,105 @@ namespace cv { namespace gpu { namespace mathfunc { compare_ne(src1, src2, dst); } + + +////////////////////////////////////////////////////////////////////////////// +// Per-element bit-wise logical matrix operations + + + __global__ void bitwise_not_kernel(int cols, int rows, const PtrStep src, PtrStep dst) + { + 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]; + } + } + + + 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); + + bitwise_not_kernel<<>>(src.cols * elemSize, src.rows, src, dst); + + if (stream == 0) + cudaSafeCall(cudaThreadSynchronize()); + } + + + __global__ void bitwise_or_kernel(int cols, int rows, const PtrStep src1, const PtrStep src2, PtrStep dst) + { + 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] = src1.ptr(y)[x] | src2.ptr(y)[x]; + } + } + + + 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); + + bitwise_or_kernel<<>>(cols * elemSize, rows, src1, src2, dst); + + if (stream == 0) + cudaSafeCall(cudaThreadSynchronize()); + } + + + __global__ void bitwise_and_kernel(int cols, int rows, const PtrStep src1, const PtrStep src2, PtrStep dst) + { + 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] = src1.ptr(y)[x] & src2.ptr(y)[x]; + } + } + + + 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); + + bitwise_and_kernel<<>>(cols * elemSize, rows, src1, src2, dst); + + if (stream == 0) + cudaSafeCall(cudaThreadSynchronize()); + } + + + + __global__ void bitwise_xor_kernel(int cols, int rows, const PtrStep src1, const PtrStep src2, PtrStep dst) + { + 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] = src1.ptr(y)[x] ^ src2.ptr(y)[x]; + } + } + + + void bitwise_xor_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); + + bitwise_xor_kernel<<>>(cols * elemSize, rows, src1, src2, dst); + + if (stream == 0) + cudaSafeCall(cudaThreadSynchronize()); + } }}} diff --git a/tests/gpu/src/bitwise_oper.cpp b/tests/gpu/src/bitwise_oper.cpp new file mode 100644 index 0000000..a3fcef2 --- /dev/null +++ b/tests/gpu/src/bitwise_oper.cpp @@ -0,0 +1,173 @@ +/*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. +// +// +// Intel License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000, Intel Corporation, 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 Intel Corporation may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include +#include +#include "gputest.hpp" + +#define CHECK(pred, err) if (!(pred)) { \ + ts->printf(CvTS::LOG, "Fail: \"%s\" at line: %d\n", #pred, __LINE__); \ + ts->set_failed_test_info(err); \ + return; } + +using namespace cv; +using namespace std; + +struct CV_GpuBitwiseTest: public CvTest +{ + CV_GpuBitwiseTest(): CvTest("GPU-BitwiseOpers", "bitwiseMatOperators") {} + + void run(int) + { + 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) + { + rows = 1 + rand() % 100; + cols = 1 + rand() % 100; + test_bitwise_not(rows, cols, CV_MAKETYPE(depth, cn)); + test_bitwise_or(rows, cols, CV_MAKETYPE(depth, cn)); + test_bitwise_and(rows, cols, CV_MAKETYPE(depth, cn)); + test_bitwise_xor(rows, cols, CV_MAKETYPE(depth, cn)); + } + } + + void test_bitwise_not(int rows, int cols, int type) + { + Mat src(rows, cols, type); + + RNG rng; + for (int i = 0; i < src.rows; ++i) + { + Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i)); + rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255)); + } + + Mat dst_gold = ~src; + gpu::GpuMat dst = ~gpu::GpuMat(src); + + 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) + } + + void test_bitwise_or(int rows, int cols, int type) + { + Mat src1(rows, cols, type); + Mat src2(rows, cols, type); + + RNG rng; + for (int i = 0; i < src1.rows; ++i) + { + Mat row1(1, src1.cols * src1.elemSize(), CV_8U, src1.ptr(i)); + rng.fill(row1, RNG::UNIFORM, Scalar(0), Scalar(255)); + Mat row2(1, src2.cols * src2.elemSize(), CV_8U, src2.ptr(i)); + rng.fill(row2, RNG::UNIFORM, Scalar(0), Scalar(255)); + } + + Mat dst_gold = src1 | src2; + gpu::GpuMat dst = gpu::GpuMat(src1) | gpu::GpuMat(src2); + + 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) + } + + void test_bitwise_and(int rows, int cols, int type) + { + Mat src1(rows, cols, type); + Mat src2(rows, cols, type); + + RNG rng; + for (int i = 0; i < src1.rows; ++i) + { + Mat row1(1, src1.cols * src1.elemSize(), CV_8U, src1.ptr(i)); + rng.fill(row1, RNG::UNIFORM, Scalar(0), Scalar(255)); + Mat row2(1, src2.cols * src2.elemSize(), CV_8U, src2.ptr(i)); + rng.fill(row2, RNG::UNIFORM, Scalar(0), Scalar(255)); + } + + Mat dst_gold = src1 & src2; + + gpu::GpuMat dst = gpu::GpuMat(src1) & gpu::GpuMat(src2); + + 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) + } + + void test_bitwise_xor(int rows, int cols, int type) + { + Mat src1(rows, cols, type); + Mat src2(rows, cols, type); + + RNG rng; + for (int i = 0; i < src1.rows; ++i) + { + Mat row1(1, src1.cols * src1.elemSize(), CV_8U, src1.ptr(i)); + rng.fill(row1, RNG::UNIFORM, Scalar(0), Scalar(255)); + Mat row2(1, src2.cols * src2.elemSize(), CV_8U, src2.ptr(i)); + rng.fill(row2, RNG::UNIFORM, Scalar(0), Scalar(255)); + } + + Mat dst_gold = src1 ^ src2; + + gpu::GpuMat dst = gpu::GpuMat(src1) ^ gpu::GpuMat(src2); + + 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) + } +} gpu_bitwise_test; + -- 2.7.4