From 4a237af814dd424c7409e654119951e4f38f68ce Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov 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 void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); + + template + 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 sourceTex(false, cudaFilterModePoint, cudaAddressModeClamp); + + template + __global__ void MHCdemosaic(PtrStepSz 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(PATTERN.y), saturate_cast(PATTERN.x), saturate_cast(C)) : + make_uchar3(saturate_cast(PATTERN.w), saturate_cast(C), saturate_cast(PATTERN.z))) : + ((alternate.x == 0) ? + make_uchar3(saturate_cast(PATTERN.z), saturate_cast(C), saturate_cast(PATTERN.w)) : + make_uchar3(saturate_cast(C), saturate_cast(PATTERN.x), saturate_cast(PATTERN.y))); + + dst(y, x) = toDst(pixelColor); + } + + template + void MHCdemosaic(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream) + { + typedef typename TypeVec::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<<>>((PtrStepSz)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 devInfo; + + virtual void SetUp() + { + devInfo = GetParam(); + + cv::gpu::setDevice(devInfo.deviceID()); + } + + static void mosaic(const cv::Mat_& src, cv::Mat_& 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_ 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_ 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_ 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_ 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_ 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_ 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_ 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_ 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