From 4a237af814dd424c7409e654119951e4f38f68ce Mon Sep 17 00:00:00 2001
From: Vladislav Vinogradov <vlad.vinogradov@itseez.com>
Date: Wed, 13 Mar 2013 14:54:06 +0400
Subject: [PATCH] implemented Malvar, He, and Cutler Bayer Demosaicing on gpu
---
modules/gpu/include/opencv2/gpu/gpu.hpp | 20 +++
modules/gpu/perf/perf_imgproc.cpp | 44 ++++++
modules/gpu/src/color.cpp | 81 ++++++++++--
modules/gpu/src/cuda/debayer.cu | 160 ++++++++++++++++++++++
modules/gpu/test/test_color.cpp | 169 ++++++++++++++++++++++++
5 files changed, 466 insertions(+), 8 deletions(-)
diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index a9481b281..802954c71 100644
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -627,6 +627,26 @@ CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat&
//! converts image from one color space to another
CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0, Stream& stream = Stream::Null());
+enum
+{
+ // Bayer Demosaicing (Malvar, He, and Cutler)
+ COLOR_BayerBG2BGR_MHT = 256,
+ COLOR_BayerGB2BGR_MHT = 257,
+ COLOR_BayerRG2BGR_MHT = 258,
+ COLOR_BayerGR2BGR_MHT = 259,
+
+ COLOR_BayerBG2RGB_MHT = COLOR_BayerRG2BGR_MHT,
+ COLOR_BayerGB2RGB_MHT = COLOR_BayerGR2BGR_MHT,
+ COLOR_BayerRG2RGB_MHT = COLOR_BayerBG2BGR_MHT,
+ COLOR_BayerGR2RGB_MHT = COLOR_BayerGB2BGR_MHT,
+
+ COLOR_BayerBG2GRAY_MHT = 260,
+ COLOR_BayerGB2GRAY_MHT = 261,
+ COLOR_BayerRG2GRAY_MHT = 262,
+ COLOR_BayerGR2GRAY_MHT = 263
+};
+CV_EXPORTS void demosaicing(const GpuMat& src, GpuMat& dst, int code, int dcn = -1, Stream& stream = Stream::Null());
+
//! swap channels
//! dstOrder - Integer array describing how channel values are permutated. The n-th entry
//! of the array contains the number of the channel that is stored in the n-th channel of
diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp
index ab445dc87..d1a87968d 100644
--- a/modules/gpu/perf/perf_imgproc.cpp
+++ b/modules/gpu/perf/perf_imgproc.cpp
@@ -1374,6 +1374,50 @@ PERF_TEST_P(Sz_Depth_Code, ImgProc_CvtColorBayer,
}
}
+CV_ENUM(DemosaicingCode,
+ cv::COLOR_BayerBG2BGR, cv::COLOR_BayerGB2BGR, cv::COLOR_BayerRG2BGR, cv::COLOR_BayerGR2BGR,
+ cv::COLOR_BayerBG2GRAY, cv::COLOR_BayerGB2GRAY, cv::COLOR_BayerRG2GRAY, cv::COLOR_BayerGR2GRAY,
+ cv::gpu::COLOR_BayerBG2BGR_MHT, cv::gpu::COLOR_BayerGB2BGR_MHT, cv::gpu::COLOR_BayerRG2BGR_MHT, cv::gpu::COLOR_BayerGR2BGR_MHT,
+ cv::gpu::COLOR_BayerBG2GRAY_MHT, cv::gpu::COLOR_BayerGB2GRAY_MHT, cv::gpu::COLOR_BayerRG2GRAY_MHT, cv::gpu::COLOR_BayerGR2GRAY_MHT)
+
+DEF_PARAM_TEST(Sz_Code, cv::Size, DemosaicingCode);
+
+PERF_TEST_P(Sz_Code, ImgProc_Demosaicing,
+ Combine(GPU_TYPICAL_MAT_SIZES,
+ ValuesIn(DemosaicingCode::all())))
+{
+ const cv::Size size = GET_PARAM(0);
+ const int code = GET_PARAM(1);
+
+ cv::Mat src(size, CV_8UC1);
+ declare.in(src, WARMUP_RNG);
+
+ if (PERF_RUN_GPU())
+ {
+ const cv::gpu::GpuMat d_src(src);
+ cv::gpu::GpuMat dst;
+
+ TEST_CYCLE() cv::gpu::demosaicing(d_src, dst, code);
+
+ GPU_SANITY_CHECK(dst);
+ }
+ else
+ {
+ if (code >= cv::COLOR_COLORCVT_MAX)
+ {
+ FAIL_NO_CPU();
+ }
+ else
+ {
+ cv::Mat dst;
+
+ TEST_CYCLE() cv::cvtColor(src, dst, code);
+
+ CPU_SANITY_CHECK(dst);
+ }
+ }
+}
+
//////////////////////////////////////////////////////////////////////
// SwapChannels
diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp
index 09986e8c3..76793d520 100644
--- a/modules/gpu/src/color.cpp
+++ b/modules/gpu/src/color.cpp
@@ -48,6 +48,7 @@ using namespace cv::gpu;
#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, Stream&) { throw_nogpu(); }
+void cv::gpu::demosaicing(const GpuMat&, GpuMat&, int, int, Stream&) { throw_nogpu(); }
void cv::gpu::swapChannels(GpuMat&, const int[], Stream&) { throw_nogpu(); }
void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_nogpu(); }
@@ -62,6 +63,9 @@ namespace cv { namespace gpu {
void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
template <int cn>
void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
+
+ template <int cn>
+ void MHCdemosaic(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
}
}}
@@ -1620,26 +1624,23 @@ namespace
funcs[src.depth()][dcn - 1](src, dst, blue_last, start_with_green, StreamAccessor::getStream(stream));
}
-
void bayerBG_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
{
bayer_to_bgr(src, dst, dcn, false, false, stream);
}
-
void bayerGB_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
{
bayer_to_bgr(src, dst, dcn, false, true, stream);
}
-
void bayerRG_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
{
bayer_to_bgr(src, dst, dcn, true, false, stream);
}
-
void bayerGR_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)
{
bayer_to_bgr(src, dst, dcn, true, true, stream);
}
+
void bayer_to_gray(const GpuMat& src, GpuMat& dst, bool blue_last, bool start_with_green, Stream& stream)
{
typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
@@ -1657,22 +1658,18 @@ namespace
funcs[src.depth()](src, dst, blue_last, start_with_green, StreamAccessor::getStream(stream));
}
-
void bayerBG_to_gray(const GpuMat& src, GpuMat& dst, int /*dcn*/, Stream& stream)
{
bayer_to_gray(src, dst, false, false, stream);
}
-
void bayerGB_to_gray(const GpuMat& src, GpuMat& dst, int /*dcn*/, Stream& stream)
{
bayer_to_gray(src, dst, false, true, stream);
}
-
void bayerRG_to_gray(const GpuMat& src, GpuMat& dst, int /*dcn*/, Stream& stream)
{
bayer_to_gray(src, dst, true, false, stream);
}
-
void bayerGR_to_gray(const GpuMat& src, GpuMat& dst, int /*dcn*/, Stream& stream)
{
bayer_to_gray(src, dst, true, true, stream);
@@ -1862,6 +1859,74 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
func(src, dst, dcn, stream);
}
+void cv::gpu::demosaicing(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream)
+{
+ const int depth = src.depth();
+
+ CV_Assert( src.channels() == 1 );
+
+ switch (code)
+ {
+ case CV_BayerBG2GRAY: case CV_BayerGB2GRAY: case CV_BayerRG2GRAY: case CV_BayerGR2GRAY:
+ bayer_to_gray(src, dst, code == CV_BayerBG2GRAY || code == CV_BayerGB2GRAY, code == CV_BayerGB2GRAY || code == CV_BayerGR2GRAY, stream);
+ break;
+
+ case CV_BayerBG2BGR: case CV_BayerGB2BGR: case CV_BayerRG2BGR: case CV_BayerGR2BGR:
+ bayer_to_bgr(src, dst, dcn, code == CV_BayerBG2BGR || code == CV_BayerGB2BGR, code == CV_BayerGB2BGR || code == CV_BayerGR2BGR, stream);
+ break;
+
+ case COLOR_BayerBG2BGR_MHT: case COLOR_BayerGB2BGR_MHT: case COLOR_BayerRG2BGR_MHT: case COLOR_BayerGR2BGR_MHT:
+ {
+ if (dcn <= 0)
+ dcn = 3;
+
+ CV_Assert( depth == CV_8U );
+ CV_Assert( dcn == 3 || dcn == 4 );
+
+ dst.create(src.size(), CV_MAKETYPE(depth, dcn));
+ dst.setTo(Scalar::all(0));
+
+ Size wholeSize;
+ Point ofs;
+ src.locateROI(wholeSize, ofs);
+ PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
+
+ const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1,
+ code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);
+
+ if (dcn == 3)
+ device::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+ else
+ device::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+
+ break;
+ }
+
+ case COLOR_BayerBG2GRAY_MHT: case COLOR_BayerGB2GRAY_MHT: case COLOR_BayerRG2GRAY_MHT: case COLOR_BayerGR2GRAY_MHT:
+ {
+ CV_Assert( depth == CV_8U );
+
+ dst.create(src.size(), CV_MAKETYPE(depth, 1));
+ dst.setTo(Scalar::all(0));
+
+ Size wholeSize;
+ Point ofs;
+ src.locateROI(wholeSize, ofs);
+ PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);
+
+ const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1,
+ code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);
+
+ device::MHCdemosaic<1>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
+
+ break;
+ }
+
+ default:
+ CV_Error( CV_StsBadFlag, "Unknown / unsupported color conversion code" );
+ }
+}
+
void cv::gpu::swapChannels(GpuMat& image, const int dstOrder[4], Stream& s)
{
CV_Assert(image.type() == CV_8UC4);
diff --git a/modules/gpu/src/cuda/debayer.cu b/modules/gpu/src/cuda/debayer.cu
index fc4372629..1d2f18e7a 100644
--- a/modules/gpu/src/cuda/debayer.cu
+++ b/modules/gpu/src/cuda/debayer.cu
@@ -47,6 +47,7 @@
#include "opencv2/gpu/device/vec_math.hpp"
#include "opencv2/gpu/device/limits.hpp"
#include "opencv2/gpu/device/color.hpp"
+#include "opencv2/gpu/device/saturate_cast.hpp"
namespace cv { namespace gpu { namespace device
{
@@ -379,6 +380,165 @@ namespace cv { namespace gpu { namespace device
template void Bayer2BGR_16u_gpu<1>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
template void Bayer2BGR_16u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
template void Bayer2BGR_16u_gpu<4>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
+
+ //////////////////////////////////////////////////////////////
+ // Bayer Demosaicing (Malvar, He, and Cutler)
+ //
+ // by Morgan McGuire, Williams College
+ // http://graphics.cs.williams.edu/papers/BayerJGT09/#shaders
+ //
+ // ported to CUDA
+
+ texture<uchar, cudaTextureType2D, cudaReadModeElementType> sourceTex(false, cudaFilterModePoint, cudaAddressModeClamp);
+
+ template <typename DstType>
+ __global__ void MHCdemosaic(PtrStepSz<DstType> dst, const int2 sourceOffset, const int2 firstRed)
+ {
+ const float kAx = -1.0f / 8.0f, kAy = -1.5f / 8.0f, kAz = 0.5f / 8.0f /*kAw = -1.0f / 8.0f*/;
+ const float kBx = 2.0f / 8.0f, /*kBy = 0.0f / 8.0f,*/ /*kBz = 0.0f / 8.0f,*/ kBw = 4.0f / 8.0f ;
+ const float kCx = 4.0f / 8.0f, kCy = 6.0f / 8.0f, kCz = 5.0f / 8.0f /*kCw = 5.0f / 8.0f*/;
+ const float /*kDx = 0.0f / 8.0f,*/ kDy = 2.0f / 8.0f, kDz = -1.0f / 8.0f /*kDw = -1.0f / 8.0f*/;
+ const float kEx = -1.0f / 8.0f, kEy = -1.5f / 8.0f, /*kEz = -1.0f / 8.0f,*/ kEw = 0.5f / 8.0f ;
+ const float kFx = 2.0f / 8.0f, /*kFy = 0.0f / 8.0f,*/ kFz = 4.0f / 8.0f /*kFw = 0.0f / 8.0f*/;
+
+ const int x = blockIdx.x * blockDim.x + threadIdx.x;
+ const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (x == 0 || x >= dst.cols - 1 || y == 0 || y >= dst.rows - 1)
+ return;
+
+ int2 center;
+ center.x = x + sourceOffset.x;
+ center.y = y + sourceOffset.y;
+
+ int4 xCoord;
+ xCoord.x = center.x - 2;
+ xCoord.y = center.x - 1;
+ xCoord.z = center.x + 1;
+ xCoord.w = center.x + 2;
+
+ int4 yCoord;
+ yCoord.x = center.y - 2;
+ yCoord.y = center.y - 1;
+ yCoord.z = center.y + 1;
+ yCoord.w = center.y + 2;
+
+ float C = tex2D(sourceTex, center.x, center.y); // ( 0, 0)
+
+ float4 Dvec;
+ Dvec.x = tex2D(sourceTex, xCoord.y, yCoord.y); // (-1,-1)
+ Dvec.y = tex2D(sourceTex, xCoord.y, yCoord.z); // (-1, 1)
+ Dvec.z = tex2D(sourceTex, xCoord.z, yCoord.y); // ( 1,-1)
+ Dvec.w = tex2D(sourceTex, xCoord.z, yCoord.z); // ( 1, 1)
+
+ float4 value;
+ value.x = tex2D(sourceTex, center.x, yCoord.x); // ( 0,-2) A0
+ value.y = tex2D(sourceTex, center.x, yCoord.y); // ( 0,-1) B0
+ value.z = tex2D(sourceTex, xCoord.x, center.y); // (-2, 0) E0
+ value.w = tex2D(sourceTex, xCoord.y, center.y); // (-1, 0) F0
+
+ // (A0 + A1), (B0 + B1), (E0 + E1), (F0 + F1)
+ value.x += tex2D(sourceTex, center.x, yCoord.w); // ( 0, 2) A1
+ value.y += tex2D(sourceTex, center.x, yCoord.z); // ( 0, 1) B1
+ value.z += tex2D(sourceTex, xCoord.w, center.y); // ( 2, 0) E1
+ value.w += tex2D(sourceTex, xCoord.z, center.y); // ( 1, 0) F1
+
+ float4 PATTERN;
+ PATTERN.x = kCx * C;
+ PATTERN.y = kCy * C;
+ PATTERN.z = kCz * C;
+ PATTERN.w = PATTERN.z;
+
+ float D = Dvec.x + Dvec.y + Dvec.z + Dvec.w;
+
+ // There are five filter patterns (identity, cross, checker,
+ // theta, phi). Precompute the terms from all of them and then
+ // use swizzles to assign to color channels.
+ //
+ // Channel Matches
+ // x cross (e.g., EE G)
+ // y checker (e.g., EE B)
+ // z theta (e.g., EO R)
+ // w phi (e.g., EO B)
+
+ #define A value.x // A0 + A1
+ #define B value.y // B0 + B1
+ #define E value.z // E0 + E1
+ #define F value.w // F0 + F1
+
+ float3 temp;
+
+ // PATTERN.yzw += (kD.yz * D).xyy;
+ temp.x = kDy * D;
+ temp.y = kDz * D;
+ PATTERN.y += temp.x;
+ PATTERN.z += temp.y;
+ PATTERN.w += temp.y;
+
+ // PATTERN += (kA.xyz * A).xyzx;
+ temp.x = kAx * A;
+ temp.y = kAy * A;
+ temp.z = kAz * A;
+ PATTERN.x += temp.x;
+ PATTERN.y += temp.y;
+ PATTERN.z += temp.z;
+ PATTERN.w += temp.x;
+
+ // PATTERN += (kE.xyw * E).xyxz;
+ temp.x = kEx * E;
+ temp.y = kEy * E;
+ temp.z = kEw * E;
+ PATTERN.x += temp.x;
+ PATTERN.y += temp.y;
+ PATTERN.z += temp.x;
+ PATTERN.w += temp.z;
+
+ // PATTERN.xw += kB.xw * B;
+ PATTERN.x += kBx * B;
+ PATTERN.w += kBw * B;
+
+ // PATTERN.xz += kF.xz * F;
+ PATTERN.x += kFx * F;
+ PATTERN.z += kFz * F;
+
+ // Determine which of four types of pixels we are on.
+ int2 alternate;
+ alternate.x = (x + firstRed.x) % 2;
+ alternate.y = (y + firstRed.y) % 2;
+
+ // in BGR sequence;
+ uchar3 pixelColor =
+ (alternate.y == 0) ?
+ ((alternate.x == 0) ?
+ make_uchar3(saturate_cast<uchar>(PATTERN.y), saturate_cast<uchar>(PATTERN.x), saturate_cast<uchar>(C)) :
+ make_uchar3(saturate_cast<uchar>(PATTERN.w), saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.z))) :
+ ((alternate.x == 0) ?
+ make_uchar3(saturate_cast<uchar>(PATTERN.z), saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.w)) :
+ make_uchar3(saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.x), saturate_cast<uchar>(PATTERN.y)));
+
+ dst(y, x) = toDst<DstType>(pixelColor);
+ }
+
+ template <int cn>
+ void MHCdemosaic(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream)
+ {
+ typedef typename TypeVec<uchar, cn>::vec_type dst_t;
+
+ const dim3 block(32, 8);
+ const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y));
+
+ bindTexture(&sourceTex, src);
+
+ MHCdemosaic<dst_t><<<grid, block, 0, stream>>>((PtrStepSz<dst_t>)dst, sourceOffset, firstRed);
+ cudaSafeCall( cudaGetLastError() );
+
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+
+ template void MHCdemosaic<1>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
+ template void MHCdemosaic<3>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
+ template void MHCdemosaic<4>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
}}}
#endif /* CUDA_DISABLER */
diff --git a/modules/gpu/test/test_color.cpp b/modules/gpu/test/test_color.cpp
index 81831af8c..3657107d9 100644
--- a/modules/gpu/test/test_color.cpp
+++ b/modules/gpu/test/test_color.cpp
@@ -2288,6 +2288,175 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CvtColor, testing::Combine(
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),
WHOLE_SUBMAT));
+///////////////////////////////////////////////////////////////////////////////////////////////////////
+// Demosaicing
+
+struct Demosaicing : testing::TestWithParam<cv::gpu::DeviceInfo>
+{
+ cv::gpu::DeviceInfo devInfo;
+
+ virtual void SetUp()
+ {
+ devInfo = GetParam();
+
+ cv::gpu::setDevice(devInfo.deviceID());
+ }
+
+ static void mosaic(const cv::Mat_<cv::Vec3b>& src, cv::Mat_<uchar>& dst, cv::Point firstRed)
+ {
+ dst.create(src.size());
+
+ for (int y = 0; y < src.rows; ++y)
+ {
+ for (int x = 0; x < src.cols; ++x)
+ {
+ cv::Vec3b pix = src(y, x);
+
+ cv::Point alternate;
+ alternate.x = (x + firstRed.x) % 2;
+ alternate.y = (y + firstRed.y) % 2;
+
+ if (alternate.y == 0)
+ {
+ if (alternate.x == 0)
+ {
+ // RG
+ // GB
+ dst(y, x) = pix[2];
+ }
+ else
+ {
+ // GR
+ // BG
+ dst(y, x) = pix[1];
+ }
+ }
+ else
+ {
+ if (alternate.x == 0)
+ {
+ // GB
+ // RG
+ dst(y, x) = pix[1];
+ }
+ else
+ {
+ // BG
+ // GR
+ dst(y, x) = pix[0];
+ }
+ }
+ }
+ }
+ }
+};
+
+GPU_TEST_P(Demosaicing, BayerBG2BGR)
+{
+ cv::Mat img = readImage("stereobm/aloe-L.png");
+
+ cv::Mat_<uchar> src;
+ mosaic(img, src, cv::Point(1, 1));
+
+ cv::gpu::GpuMat dst;
+ cv::gpu::demosaicing(loadMat(src), dst, cv::COLOR_BayerBG2BGR);
+
+ EXPECT_MAT_SIMILAR(img, dst, 2e-2);
+}
+
+GPU_TEST_P(Demosaicing, BayerGB2BGR)
+{
+ cv::Mat img = readImage("stereobm/aloe-L.png");
+
+ cv::Mat_<uchar> src;
+ mosaic(img, src, cv::Point(0, 1));
+
+ cv::gpu::GpuMat dst;
+ cv::gpu::demosaicing(loadMat(src), dst, cv::COLOR_BayerGB2BGR);
+
+ EXPECT_MAT_SIMILAR(img, dst, 2e-2);
+}
+
+GPU_TEST_P(Demosaicing, BayerRG2BGR)
+{
+ cv::Mat img = readImage("stereobm/aloe-L.png");
+
+ cv::Mat_<uchar> src;
+ mosaic(img, src, cv::Point(0, 0));
+
+ cv::gpu::GpuMat dst;
+ cv::gpu::demosaicing(loadMat(src), dst, cv::COLOR_BayerRG2BGR);
+
+ EXPECT_MAT_SIMILAR(img, dst, 2e-2);
+}
+
+GPU_TEST_P(Demosaicing, BayerGR2BGR)
+{
+ cv::Mat img = readImage("stereobm/aloe-L.png");
+
+ cv::Mat_<uchar> src;
+ mosaic(img, src, cv::Point(1, 0));
+
+ cv::gpu::GpuMat dst;
+ cv::gpu::demosaicing(loadMat(src), dst, cv::COLOR_BayerGR2BGR);
+
+ EXPECT_MAT_SIMILAR(img, dst, 2e-2);
+}
+
+GPU_TEST_P(Demosaicing, BayerBG2BGR_MHT)
+{
+ cv::Mat img = readImage("stereobm/aloe-L.png");
+
+ cv::Mat_<uchar> src;
+ mosaic(img, src, cv::Point(1, 1));
+
+ cv::gpu::GpuMat dst;
+ cv::gpu::demosaicing(loadMat(src), dst, cv::gpu::COLOR_BayerBG2BGR_MHT);
+
+ EXPECT_MAT_SIMILAR(img, dst, 5e-3);
+}
+
+GPU_TEST_P(Demosaicing, BayerGB2BGR_MHT)
+{
+ cv::Mat img = readImage("stereobm/aloe-L.png");
+
+ cv::Mat_<uchar> src;
+ mosaic(img, src, cv::Point(0, 1));
+
+ cv::gpu::GpuMat dst;
+ cv::gpu::demosaicing(loadMat(src), dst, cv::gpu::COLOR_BayerGB2BGR_MHT);
+
+ EXPECT_MAT_SIMILAR(img, dst, 5e-3);
+}
+
+GPU_TEST_P(Demosaicing, BayerRG2BGR_MHT)
+{
+ cv::Mat img = readImage("stereobm/aloe-L.png");
+
+ cv::Mat_<uchar> src;
+ mosaic(img, src, cv::Point(0, 0));
+
+ cv::gpu::GpuMat dst;
+ cv::gpu::demosaicing(loadMat(src), dst, cv::gpu::COLOR_BayerRG2BGR_MHT);
+
+ EXPECT_MAT_SIMILAR(img, dst, 5e-3);
+}
+
+GPU_TEST_P(Demosaicing, BayerGR2BGR_MHT)
+{
+ cv::Mat img = readImage("stereobm/aloe-L.png");
+
+ cv::Mat_<uchar> src;
+ mosaic(img, src, cv::Point(1, 0));
+
+ cv::gpu::GpuMat dst;
+ cv::gpu::demosaicing(loadMat(src), dst, cv::gpu::COLOR_BayerGR2BGR_MHT);
+
+ EXPECT_MAT_SIMILAR(img, dst, 5e-3);
+}
+
+INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Demosaicing, ALL_DEVICES);
+
///////////////////////////////////////////////////////////////////////////////////////////////////////
// swapChannels