From 8624d18ca5694be59d8a689c28c5f572751f9f71 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 6 Aug 2012 10:58:11 +0400 Subject: [PATCH] new implementation of gpu debayer * without border extrapolation * with aligned write * process 4 pixels per thread in 8u case --- modules/gpu/perf/perf_imgproc.cpp | 3 + modules/gpu/src/color.cpp | 10 +- modules/gpu/src/cuda/debayer.cu | 351 ++++++++++++++++++++---------- modules/gpu/test/test_color.cpp | 53 +++-- 4 files changed, 273 insertions(+), 144 deletions(-) diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp index c1e348f794..b5c986d220 100644 --- a/modules/gpu/perf/perf_imgproc.cpp +++ b/modules/gpu/perf/perf_imgproc.cpp @@ -1148,6 +1148,9 @@ GPU_PERF_TEST(CvtColor, cv::gpu::DeviceInfo, cv::Size, MatDepth, CvtColorInfo) cv::gpu::GpuMat src(src_host); cv::gpu::GpuMat dst; + if (info.code >= cv::COLOR_BayerBG2BGR && info.code <= cv::COLOR_BayerGR2BGR) + info.dcn = 4; + cv::gpu::cvtColor(src, dst, info.code, info.dcn); TEST_CYCLE() diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp index d5615aba4f..403e2dfae3 100644 --- a/modules/gpu/src/color.cpp +++ b/modules/gpu/src/color.cpp @@ -58,8 +58,10 @@ void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_nog namespace cv { namespace gpu { namespace device { - template - void Bayer2BGR_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); + template + void Bayer2BGR_8u_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); + template + void Bayer2BGR_16u_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); } }} @@ -1337,9 +1339,9 @@ namespace typedef void (*func_t)(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); static const func_t funcs[3][4] = { - {0,0,Bayer2BGR_gpu, Bayer2BGR_gpu}, + {0,0,Bayer2BGR_8u_gpu<3>, Bayer2BGR_8u_gpu<4>}, {0,0,0,0}, - {0,0,Bayer2BGR_gpu, Bayer2BGR_gpu} + {0,0,Bayer2BGR_16u_gpu<3>, Bayer2BGR_16u_gpu<4>} }; if (dcn <= 0) dcn = 3; diff --git a/modules/gpu/src/cuda/debayer.cu b/modules/gpu/src/cuda/debayer.cu index 61adfd64b1..fc7533b5b3 100644 --- a/modules/gpu/src/cuda/debayer.cu +++ b/modules/gpu/src/cuda/debayer.cu @@ -42,167 +42,286 @@ #include #include +#include +#include namespace cv { namespace gpu { namespace device { - template - __global__ void Bayer2BGR(const SrcPtr src, PtrStep_ dst, const int width, const int height, const bool glob_blue_last, const bool glob_start_with_green) + template + __global__ void Bayer2BGR_8u(const PtrStepb src, DevMem2D_ dst, const bool blue_last, const bool start_with_green) { - const int tx = blockIdx.x * blockDim.x + threadIdx.x; - const int y = blockIdx.y * blockDim.y + threadIdx.y; + const int s_x = blockIdx.x * blockDim.x + threadIdx.x; + int s_y = blockIdx.y * blockDim.y + threadIdx.y; - if (y >= height) + if (s_y >= dst.rows || (s_x << 2) >= dst.cols) return; - const bool blue_last = (y & 1) ? !glob_blue_last : glob_blue_last; - const bool start_with_green = (y & 1) ? !glob_start_with_green : glob_start_with_green; + s_y = ::min(::max(s_y, 1), dst.rows - 2); - int x = tx * 2; + uchar4 patch[3][3]; + patch[0][1] = ((const uchar4*) src.ptr(s_y - 1))[s_x]; + patch[0][0] = ((const uchar4*) src.ptr(s_y - 1))[::max(s_x - 1, 0)]; + patch[0][2] = ((const uchar4*) src.ptr(s_y - 1))[::min(s_x + 1, ((dst.cols + 3) >> 2) - 1)]; - if (start_with_green) + patch[1][1] = ((const uchar4*) src.ptr(s_y))[s_x]; + patch[1][0] = ((const uchar4*) src.ptr(s_y))[::max(s_x - 1, 0)]; + patch[1][2] = ((const uchar4*) src.ptr(s_y))[::min(s_x + 1, ((dst.cols + 3) >> 2) - 1)]; + + patch[2][1] = ((const uchar4*) src.ptr(s_y + 1))[s_x]; + patch[2][0] = ((const uchar4*) src.ptr(s_y + 1))[::max(s_x - 1, 0)]; + patch[2][2] = ((const uchar4*) src.ptr(s_y + 1))[::min(s_x + 1, ((dst.cols + 3) >> 2) - 1)]; + + D res0 = VecTraits::all(numeric_limits::max()); + D res1 = VecTraits::all(numeric_limits::max()); + D res2 = VecTraits::all(numeric_limits::max()); + D res3 = VecTraits::all(numeric_limits::max()); + + if ((s_y & 1) ^ start_with_green) { - --x; + const int t0 = (patch[0][1].x + patch[2][1].x + 1) >> 1; + const int t1 = (patch[1][0].w + patch[1][1].y + 1) >> 1; - if (tx == 0) + const int t2 = (patch[0][1].x + patch[0][1].z + patch[2][1].x + patch[2][1].z + 2) >> 2; + const int t3 = (patch[0][1].y + patch[1][1].x + patch[1][1].z + patch[2][1].y + 2) >> 2; + + const int t4 = (patch[0][1].z + patch[2][1].z + 1) >> 1; + const int t5 = (patch[1][1].y + patch[1][1].w + 1) >> 1; + + const int t6 = (patch[0][1].z + patch[0][2].x + patch[2][1].z + patch[2][2].x + 2) >> 2; + const int t7 = (patch[0][1].w + patch[1][1].z + patch[1][2].x + patch[2][1].w + 2) >> 2; + + if ((s_y & 1) ^ blue_last) { - const int t0 = (src(y, 1) + src(y + 2, 1) + 1) >> 1; - const int t1 = (src(y + 1, 0) + src(y + 1, 2) + 1) >> 1; + res0.x = t1; + res0.y = patch[1][1].x; + res0.z = t0; - T res; - res.x = blue_last ? t0 : t1; - res.y = src(y + 1, 1); - res.z = blue_last ? t1 : t0; + res1.x = patch[1][1].y; + res1.y = t3; + res1.z = t2; - dst(y + 1, 0) = dst(y + 1, 1) = res; - if (y == 0) - { - dst(0, 0) = dst(0, 1) = res; - } - else if (y == height - 1) - { - dst(height + 1, 0) = dst(height + 1, 1) = res; - } - } - } + res2.x = t5; + res2.y = patch[1][1].z; + res2.z = t4; - if (x >= 0 && x <= width - 2) - { - const int t0 = (src(y, x) + src(y, x + 2) + src(y + 2, x) + src(y + 2, x + 2) + 2) >> 2; - const int t1 = (src(y, x + 1) + src(y + 1, x) + src(y + 1, x + 2) + src(y + 2, x + 1) + 2) >> 2; - - const int t2 = (src(y, x + 2) + src(y + 2, x + 2) + 1) >> 1; - const int t3 = (src(y + 1, x + 1) + src(y + 1, x + 3) + 1) >> 1; - - T res1, res2; - - if (blue_last) - { - res1.x = t0; - res1.y = t1; - res1.z = src(y + 1, x + 1); - - res2.x = t2; - res2.y = src(y + 1, x + 2); - res2.z = t3; + res3.x = patch[1][1].w; + res3.y = t7; + res3.z = t6; } else { - res1.x = src(y + 1, x + 1); - res1.y = t1; - res1.z = t0; + res0.x = t0; + res0.y = patch[1][1].x; + res0.z = t1; - res2.x = t3; - res2.y = src(y + 1, x + 2); - res2.z = t2; - } + res1.x = t2; + res1.y = t3; + res1.z = patch[1][1].y; - dst(y + 1, x + 1) = res1; - dst(y + 1, x + 2) = res2; + res2.x = t4; + res2.y = patch[1][1].z; + res2.z = t5; - if (y == 0) - { - dst(0, x + 1) = res1; - dst(0, x + 2) = res2; - - if (x == 0) - { - dst(0, 0) = res1; - } - else if (x == width - 2) - { - dst(0, width + 1) = res2; - } - } - else if (y == height - 1) - { - dst(height + 1, x + 1) = res1; - dst(height + 1, x + 2) = res2; - - if (x == 0) - { - dst(height + 1, 0) = res1; - } - else if (x == width - 2) - { - dst(height + 1, width + 1) = res2; - } - } - - if (x == 0) - { - dst(y + 1, 0) = res1; - } - else if (x == width - 2) - { - dst(y + 1, width + 1) = res2; + res3.x = t6; + res3.y = t7; + res3.z = patch[1][1].w; } } - else if (x == width - 1) + else { - const int t0 = (src(y, x) + src(y, x + 2) + src(y + 2, x) + src(y + 2, x + 2) + 2) >> 2; - const int t1 = (src(y, x + 1) + src(y + 1, x) + src(y + 1, x + 2) + src(y + 2, x + 1) + 2) >> 2; + const int t0 = (patch[0][0].w + patch[0][1].y + patch[2][0].w + patch[2][1].y + 2) >> 2; + const int t1 = (patch[0][1].x + patch[1][0].w + patch[1][1].y + patch[2][1].x + 2) >> 2; - T res; - res.x = blue_last ? t0 : src(y + 1, x + 1); - res.y = t1; - res.z = blue_last ? src(y + 1, x + 1) : t0; + const int t2 = (patch[0][1].y + patch[2][1].y + 1) >> 1; + const int t3 = (patch[1][1].x + patch[1][1].z + 1) >> 1; - dst(y + 1, x + 1) = dst(y + 1, x + 2) = res; - if (y == 0) + const int t4 = (patch[0][1].y + patch[0][1].w + patch[2][1].y + patch[2][1].w + 2) >> 2; + const int t5 = (patch[0][1].z + patch[1][1].y + patch[1][1].w + patch[2][1].z + 2) >> 2; + + const int t6 = (patch[0][1].w + patch[2][1].w + 1) >> 1; + const int t7 = (patch[1][1].z + patch[1][2].x + 1) >> 1; + + if ((s_y & 1) ^ blue_last) { - dst(0, x + 1) = dst(0, x + 2) = res; + res0.x = patch[1][1].x; + res0.y = t1; + res0.z = t0; + + res1.x = t3; + res1.y = patch[1][1].y; + res1.z = t2; + + res2.x = patch[1][1].z; + res2.y = t5; + res2.z = t4; + + res3.x = t7; + res3.y = patch[1][1].w; + res3.z = t6; } - else if (y == height - 1) + else { - dst(height + 1, x + 1) = dst(height + 1, x + 2) = res; + res0.x = t0; + res0.y = t1; + res0.z = patch[1][1].x; + + res1.x = t2; + res1.y = patch[1][1].y; + res1.z = t3; + + res2.x = t4; + res2.y = t5; + res2.z = patch[1][1].z; + + res3.x = t6; + res3.y = patch[1][1].w; + res3.z = t7; } } + + const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 2; + const int d_y = blockIdx.y * blockDim.y + threadIdx.y; + + dst(d_y, d_x) = res0; + if (d_x + 1 < dst.cols) + dst(d_y, d_x + 1) = res1; + if (d_x + 2 < dst.cols) + dst(d_y, d_x + 2) = res2; + if (d_x + 3 < dst.cols) + dst(d_y, d_x + 3) = res3; } - template - void Bayer2BGR_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream) + template + __global__ void Bayer2BGR_16u(const PtrStepb src, DevMem2D_ dst, const bool blue_last, const bool start_with_green) { - typedef typename TypeVec::vec_type dst_t; + const int s_x = blockIdx.x * blockDim.x + threadIdx.x; + int s_y = blockIdx.y * blockDim.y + threadIdx.y; - const int width = src.cols - 2; - const int height = src.rows - 2; + if (s_y >= dst.rows || (s_x << 1) >= dst.cols) + return; - const dim3 total(divUp(width, 2), height); + s_y = ::min(::max(s_y, 1), dst.rows - 2); + + ushort2 patch[3][3]; + patch[0][1] = ((const ushort2*) src.ptr(s_y - 1))[s_x]; + patch[0][0] = ((const ushort2*) src.ptr(s_y - 1))[::max(s_x - 1, 0)]; + patch[0][2] = ((const ushort2*) src.ptr(s_y - 1))[::min(s_x + 1, ((dst.cols + 1) >> 1) - 1)]; + + patch[1][1] = ((const ushort2*) src.ptr(s_y))[s_x]; + patch[1][0] = ((const ushort2*) src.ptr(s_y))[::max(s_x - 1, 0)]; + patch[1][2] = ((const ushort2*) src.ptr(s_y))[::min(s_x + 1, ((dst.cols + 1) >> 1) - 1)]; + + patch[2][1] = ((const ushort2*) src.ptr(s_y + 1))[s_x]; + patch[2][0] = ((const ushort2*) src.ptr(s_y + 1))[::max(s_x - 1, 0)]; + patch[2][2] = ((const ushort2*) src.ptr(s_y + 1))[::min(s_x + 1, ((dst.cols + 1) >> 1) - 1)]; + + D res0 = VecTraits::all(numeric_limits::max()); + D res1 = VecTraits::all(numeric_limits::max()); + + if ((s_y & 1) ^ start_with_green) + { + const int t0 = (patch[0][1].x + patch[2][1].x + 1) >> 1; + const int t1 = (patch[1][0].y + patch[1][1].y + 1) >> 1; + + const int t2 = (patch[0][1].x + patch[0][2].x + patch[2][1].x + patch[2][2].x + 2) >> 2; + const int t3 = (patch[0][1].y + patch[1][1].x + patch[1][2].x + patch[2][1].y + 2) >> 2; + + if ((s_y & 1) ^ blue_last) + { + res0.x = t1; + res0.y = patch[1][1].x; + res0.z = t0; + + res1.x = patch[1][1].y; + res1.y = t3; + res1.z = t2; + } + else + { + res0.x = t0; + res0.y = patch[1][1].x; + res0.z = t1; + + res1.x = t2; + res1.y = t3; + res1.z = patch[1][1].y; + } + } + else + { + const int t0 = (patch[0][0].y + patch[0][1].y + patch[2][0].y + patch[2][1].y + 2) >> 2; + const int t1 = (patch[0][1].x + patch[1][0].y + patch[1][1].y + patch[2][1].x + 2) >> 2; + + const int t2 = (patch[0][1].y + patch[2][1].y + 1) >> 1; + const int t3 = (patch[1][1].x + patch[1][2].x + 1) >> 1; + + if ((s_y & 1) ^ blue_last) + { + res0.x = patch[1][1].x; + res0.y = t1; + res0.z = t0; + + res1.x = t3; + res1.y = patch[1][1].y; + res1.z = t2; + } + else + { + res0.x = t0; + res0.y = t1; + res0.z = patch[1][1].x; + + res1.x = t2; + res1.y = patch[1][1].y; + res1.z = t3; + } + } + + const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 1; + const int d_y = blockIdx.y * blockDim.y + threadIdx.y; + + dst(d_y, d_x) = res0; + if (d_x + 1 < dst.cols) + dst(d_y, d_x + 1) = res1; + } + + template + void Bayer2BGR_8u_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream) + { + typedef typename TypeVec::vec_type dst_t; const dim3 block(32, 8); - const dim3 grid(divUp(total.x, block.x), divUp(total.y, block.y)); + const dim3 grid(divUp(dst.cols, 4 * block.x), divUp(dst.rows, block.y)); - Bayer2BGR, dst_t><<>>((DevMem2D_)src, (DevMem2D_)dst, width, height, blue_last, start_with_green); + cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u, cudaFuncCachePreferL1) ); + + Bayer2BGR_8u<<>>(src, (DevMem2D_)dst, blue_last, start_with_green); + cudaSafeCall( cudaGetLastError() ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } + template + void Bayer2BGR_16u_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream) + { + typedef typename TypeVec::vec_type dst_t; + + const dim3 block(32, 8); + const dim3 grid(divUp(dst.cols, 2 * block.x), divUp(dst.rows, block.y)); + + cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u, cudaFuncCachePreferL1) ); + + Bayer2BGR_16u<<>>(src, (DevMem2D_)dst, blue_last, start_with_green); cudaSafeCall( cudaGetLastError() ); if (stream == 0) cudaSafeCall( cudaDeviceSynchronize() ); } - template void Bayer2BGR_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); - template void Bayer2BGR_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); - template void Bayer2BGR_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); - template void Bayer2BGR_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); + template void Bayer2BGR_8u_gpu<3>(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); + template void Bayer2BGR_8u_gpu<4>(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); + template void Bayer2BGR_16u_gpu<3>(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); + template void Bayer2BGR_16u_gpu<4>(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); } }} diff --git a/modules/gpu/test/test_color.cpp b/modules/gpu/test/test_color.cpp index bb522eb140..645967ef27 100644 --- a/modules/gpu/test/test_color.cpp +++ b/modules/gpu/test/test_color.cpp @@ -41,6 +41,8 @@ #include "precomp.hpp" +#ifdef HAVE_CUDA + namespace { /////////////////////////////////////////////////////////////////////////////////////////////////////// @@ -1746,10 +1748,10 @@ TEST_P(CvtColor, RGBA2mRGBA) TEST_P(CvtColor, BayerBG2BGR) { - if (depth != CV_8U && depth != CV_16U) + if ((depth != CV_8U && depth != CV_16U) || useRoi) return; - cv::Mat src = randomMat(size, CV_8UC1); + cv::Mat src = randomMat(size, depth); cv::gpu::GpuMat dst; cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerBG2BGR); @@ -1757,15 +1759,15 @@ TEST_P(CvtColor, BayerBG2BGR) cv::Mat dst_gold; cv::cvtColor(src, dst_gold, cv::COLOR_BayerBG2BGR); - EXPECT_MAT_NEAR(dst_gold, dst, 0); + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0); } TEST_P(CvtColor, BayerBG2BGR4) { - if (depth != CV_8U && depth != CV_16U) + if ((depth != CV_8U && depth != CV_16U) || useRoi) return; - cv::Mat src = randomMat(size, CV_8UC1); + cv::Mat src = randomMat(size, depth); cv::gpu::GpuMat dst; cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerBG2BGR, 4); @@ -1779,15 +1781,16 @@ TEST_P(CvtColor, BayerBG2BGR4) cv::Mat dst3; cv::cvtColor(dst4, dst3, cv::COLOR_BGRA2BGR); - EXPECT_MAT_NEAR(dst_gold, dst3, 0); + + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst3(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0); } TEST_P(CvtColor, BayerGB2BGR) { - if (depth != CV_8U && depth != CV_16U) + if ((depth != CV_8U && depth != CV_16U) || useRoi) return; - cv::Mat src = randomMat(size, CV_8UC1); + cv::Mat src = randomMat(size, depth); cv::gpu::GpuMat dst; cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGB2BGR); @@ -1795,15 +1798,15 @@ TEST_P(CvtColor, BayerGB2BGR) cv::Mat dst_gold; cv::cvtColor(src, dst_gold, cv::COLOR_BayerGB2BGR); - EXPECT_MAT_NEAR(dst_gold, dst, 0); + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0); } TEST_P(CvtColor, BayerGB2BGR4) { - if (depth != CV_8U && depth != CV_16U) + if ((depth != CV_8U && depth != CV_16U) || useRoi) return; - cv::Mat src = randomMat(size, CV_8UC1); + cv::Mat src = randomMat(size, depth); cv::gpu::GpuMat dst; cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGB2BGR, 4); @@ -1817,15 +1820,15 @@ TEST_P(CvtColor, BayerGB2BGR4) cv::Mat dst3; cv::cvtColor(dst4, dst3, cv::COLOR_BGRA2BGR); - EXPECT_MAT_NEAR(dst_gold, dst3, 0); + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst3(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0); } TEST_P(CvtColor, BayerRG2BGR) { - if (depth != CV_8U && depth != CV_16U) + if ((depth != CV_8U && depth != CV_16U) || useRoi) return; - cv::Mat src = randomMat(size, CV_8UC1); + cv::Mat src = randomMat(size, depth); cv::gpu::GpuMat dst; cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerRG2BGR); @@ -1833,15 +1836,15 @@ TEST_P(CvtColor, BayerRG2BGR) cv::Mat dst_gold; cv::cvtColor(src, dst_gold, cv::COLOR_BayerRG2BGR); - EXPECT_MAT_NEAR(dst_gold, dst, 0); + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0); } TEST_P(CvtColor, BayerRG2BGR4) { - if (depth != CV_8U && depth != CV_16U) + if ((depth != CV_8U && depth != CV_16U) || useRoi) return; - cv::Mat src = randomMat(size, CV_8UC1); + cv::Mat src = randomMat(size, depth); cv::gpu::GpuMat dst; cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerRG2BGR, 4); @@ -1855,15 +1858,15 @@ TEST_P(CvtColor, BayerRG2BGR4) cv::Mat dst3; cv::cvtColor(dst4, dst3, cv::COLOR_BGRA2BGR); - EXPECT_MAT_NEAR(dst_gold, dst3, 0); + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst3(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0); } TEST_P(CvtColor, BayerGR2BGR) { - if (depth != CV_8U && depth != CV_16U) + if ((depth != CV_8U && depth != CV_16U) || useRoi) return; - cv::Mat src = randomMat(size, CV_8UC1); + cv::Mat src = randomMat(size, depth); cv::gpu::GpuMat dst; cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGR2BGR); @@ -1871,15 +1874,15 @@ TEST_P(CvtColor, BayerGR2BGR) cv::Mat dst_gold; cv::cvtColor(src, dst_gold, cv::COLOR_BayerGR2BGR); - EXPECT_MAT_NEAR(dst_gold, dst, 0); + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0); } TEST_P(CvtColor, BayerGR2BGR4) { - if (depth != CV_8U && depth != CV_16U) + if ((depth != CV_8U && depth != CV_16U) || useRoi) return; - cv::Mat src = randomMat(size, CV_8UC1); + cv::Mat src = randomMat(size, depth); cv::gpu::GpuMat dst; cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGR2BGR, 4); @@ -1893,7 +1896,7 @@ TEST_P(CvtColor, BayerGR2BGR4) cv::Mat dst3; cv::cvtColor(dst4, dst3, cv::COLOR_BGRA2BGR); - EXPECT_MAT_NEAR(dst_gold, dst3, 0); + EXPECT_MAT_NEAR(dst_gold(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), dst3(cv::Rect(1, 1, dst.cols - 2, dst.rows - 2)), 0); } INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CvtColor, testing::Combine( @@ -1943,3 +1946,5 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, SwapChannels, testing::Combine( WHOLE_SUBMAT)); } // namespace + +#endif // HAVE_CUDA