generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

added bitwise operations into gpu module

Browse Source
This commit is contained in:
Alexey Spizhevoy 2010-11-19 10:19:35 +00:00
parent 7df9aef99c
commit faf4d0bc74
4 changed files with 442 additions and 12 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

52
modules/gpu/include/opencv2/gpu/gpu.hpp
View File

@ -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:

128
modules/gpu/src/arithm.cpp
View File

@ -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) */

101
modules/gpu/src/cuda/mathfunc.cu
View File

@ -238,4 +238,105 @@ namespace cv { namespace gpu { namespace mathfunc
{
compare_ne<float, float>(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<<<grid, threads, 0, stream>>>(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<<<grid, threads, 0, stream>>>(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<<<grid, threads, 0, stream>>>(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<<<grid, threads, 0, stream>>>(cols * elemSize, rows, src1, src2, dst);
if (stream == 0)
cudaSafeCall(cudaThreadSynchronize());
}
}}}

173
tests/gpu/src/bitwise_oper.cpp Normal file
View File

@ -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 <iostream>
#include <limits>
#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;