From 7b025474e2977cadd33f73db1a0f5c6f461301ff Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Thu, 2 Aug 2012 15:04:00 +0400 Subject: [PATCH] added debayer to gpu::cvtColor --- modules/gpu/src/color.cpp | 58 ++++++++- modules/gpu/src/cuda/debayer.cu | 208 ++++++++++++++++++++++++++++++++ modules/gpu/test/test_color.cpp | 152 +++++++++++++++++++++++ 3 files changed, 414 insertions(+), 4 deletions(-) create mode 100644 modules/gpu/src/cuda/debayer.cu diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp index 543227aeb3..faebca6da9 100644 --- a/modules/gpu/src/color.cpp +++ b/modules/gpu/src/color.cpp @@ -54,6 +54,15 @@ void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_nog #else /* !defined (HAVE_CUDA) */ #include + +namespace cv { namespace gpu { + namespace device + { + template + void Bayer2BGR_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream); + } +}} + using namespace ::cv::gpu::device; namespace @@ -1302,6 +1311,47 @@ namespace nppSafeCall( nppiAlphaPremul_16u_AC4R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI) ); #endif } + + void bayer_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, bool blue_last, bool start_with_green, Stream& stream) + { + 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,0,0}, + {0,0,Bayer2BGR_gpu, Bayer2BGR_gpu} + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.type() == CV_8UC1 || src.type() == CV_16UC1); + CV_Assert(src.rows > 2 && src.cols > 2); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + 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 cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream) @@ -1366,10 +1416,10 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream bgr_to_lab, // CV_BGR2Lab =44 rgb_to_lab, // CV_RGB2Lab =45 - 0, // CV_BayerBG2BGR =46 - 0, // CV_BayerGB2BGR =47 - 0, // CV_BayerRG2BGR =48 - 0, // CV_BayerGR2BGR =49 + bayerBG_to_bgr, // CV_BayerBG2BGR =46 + bayerGB_to_bgr, // CV_BayerGB2BGR =47 + bayerRG_to_bgr, // CV_BayerRG2BGR =48 + bayerGR_to_bgr, // CV_BayerGR2BGR =49 bgr_to_luv, // CV_BGR2Luv =50 rgb_to_luv, // CV_RGB2Luv =51 diff --git a/modules/gpu/src/cuda/debayer.cu b/modules/gpu/src/cuda/debayer.cu new file mode 100644 index 0000000000..61adfd64b1 --- /dev/null +++ b/modules/gpu/src/cuda/debayer.cu @@ -0,0 +1,208 @@ +/*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. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., 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 the copyright holders 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 + +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) + { + const int tx = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (y >= height) + 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; + + int x = tx * 2; + + if (start_with_green) + { + --x; + + if (tx == 0) + { + 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; + + T res; + res.x = blue_last ? t0 : t1; + res.y = src(y + 1, 1); + res.z = blue_last ? t1 : t0; + + 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; + } + } + } + + 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; + } + else + { + res1.x = src(y + 1, x + 1); + res1.y = t1; + res1.z = t0; + + res2.x = t3; + res2.y = src(y + 1, x + 2); + res2.z = t2; + } + + dst(y + 1, x + 1) = res1; + dst(y + 1, x + 2) = res2; + + 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; + } + } + else if (x == width - 1) + { + 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; + + 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; + + dst(y + 1, x + 1) = dst(y + 1, x + 2) = res; + if (y == 0) + { + dst(0, x + 1) = dst(0, x + 2) = res; + } + else if (y == height - 1) + { + dst(height + 1, x + 1) = dst(height + 1, x + 2) = res; + } + } + } + + template + void Bayer2BGR_gpu(DevMem2Db src, DevMem2Db dst, bool blue_last, bool start_with_green, cudaStream_t stream) + { + typedef typename TypeVec::vec_type dst_t; + + const int width = src.cols - 2; + const int height = src.rows - 2; + + const dim3 total(divUp(width, 2), height); + + const dim3 block(32, 8); + const dim3 grid(divUp(total.x, block.x), divUp(total.y, block.y)); + + Bayer2BGR, dst_t><<>>((DevMem2D_)src, (DevMem2D_)dst, width, height, 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); + } +}} diff --git a/modules/gpu/test/test_color.cpp b/modules/gpu/test/test_color.cpp index f627326b2e..bb522eb140 100644 --- a/modules/gpu/test/test_color.cpp +++ b/modules/gpu/test/test_color.cpp @@ -1744,6 +1744,158 @@ TEST_P(CvtColor, RGBA2mRGBA) } } +TEST_P(CvtColor, BayerBG2BGR) +{ + if (depth != CV_8U && depth != CV_16U) + return; + + cv::Mat src = randomMat(size, CV_8UC1); + + cv::gpu::GpuMat dst; + cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerBG2BGR); + + cv::Mat dst_gold; + cv::cvtColor(src, dst_gold, cv::COLOR_BayerBG2BGR); + + EXPECT_MAT_NEAR(dst_gold, dst, 0); +} + +TEST_P(CvtColor, BayerBG2BGR4) +{ + if (depth != CV_8U && depth != CV_16U) + return; + + cv::Mat src = randomMat(size, CV_8UC1); + + cv::gpu::GpuMat dst; + cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerBG2BGR, 4); + + ASSERT_EQ(4, dst.channels()); + + cv::Mat dst_gold; + cv::cvtColor(src, dst_gold, cv::COLOR_BayerBG2BGR); + + cv::Mat dst4(dst); + cv::Mat dst3; + cv::cvtColor(dst4, dst3, cv::COLOR_BGRA2BGR); + + EXPECT_MAT_NEAR(dst_gold, dst3, 0); +} + +TEST_P(CvtColor, BayerGB2BGR) +{ + if (depth != CV_8U && depth != CV_16U) + return; + + cv::Mat src = randomMat(size, CV_8UC1); + + cv::gpu::GpuMat dst; + cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGB2BGR); + + cv::Mat dst_gold; + cv::cvtColor(src, dst_gold, cv::COLOR_BayerGB2BGR); + + EXPECT_MAT_NEAR(dst_gold, dst, 0); +} + +TEST_P(CvtColor, BayerGB2BGR4) +{ + if (depth != CV_8U && depth != CV_16U) + return; + + cv::Mat src = randomMat(size, CV_8UC1); + + cv::gpu::GpuMat dst; + cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGB2BGR, 4); + + ASSERT_EQ(4, dst.channels()); + + cv::Mat dst_gold; + cv::cvtColor(src, dst_gold, cv::COLOR_BayerGB2BGR); + + cv::Mat dst4(dst); + cv::Mat dst3; + cv::cvtColor(dst4, dst3, cv::COLOR_BGRA2BGR); + + EXPECT_MAT_NEAR(dst_gold, dst3, 0); +} + +TEST_P(CvtColor, BayerRG2BGR) +{ + if (depth != CV_8U && depth != CV_16U) + return; + + cv::Mat src = randomMat(size, CV_8UC1); + + cv::gpu::GpuMat dst; + cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerRG2BGR); + + cv::Mat dst_gold; + cv::cvtColor(src, dst_gold, cv::COLOR_BayerRG2BGR); + + EXPECT_MAT_NEAR(dst_gold, dst, 0); +} + +TEST_P(CvtColor, BayerRG2BGR4) +{ + if (depth != CV_8U && depth != CV_16U) + return; + + cv::Mat src = randomMat(size, CV_8UC1); + + cv::gpu::GpuMat dst; + cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerRG2BGR, 4); + + ASSERT_EQ(4, dst.channels()); + + cv::Mat dst_gold; + cv::cvtColor(src, dst_gold, cv::COLOR_BayerRG2BGR); + + cv::Mat dst4(dst); + cv::Mat dst3; + cv::cvtColor(dst4, dst3, cv::COLOR_BGRA2BGR); + + EXPECT_MAT_NEAR(dst_gold, dst3, 0); +} + +TEST_P(CvtColor, BayerGR2BGR) +{ + if (depth != CV_8U && depth != CV_16U) + return; + + cv::Mat src = randomMat(size, CV_8UC1); + + cv::gpu::GpuMat dst; + cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGR2BGR); + + cv::Mat dst_gold; + cv::cvtColor(src, dst_gold, cv::COLOR_BayerGR2BGR); + + EXPECT_MAT_NEAR(dst_gold, dst, 0); +} + +TEST_P(CvtColor, BayerGR2BGR4) +{ + if (depth != CV_8U && depth != CV_16U) + return; + + cv::Mat src = randomMat(size, CV_8UC1); + + cv::gpu::GpuMat dst; + cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGR2BGR, 4); + + ASSERT_EQ(4, dst.channels()); + + cv::Mat dst_gold; + cv::cvtColor(src, dst_gold, cv::COLOR_BayerGR2BGR); + + cv::Mat dst4(dst); + cv::Mat dst3; + cv::cvtColor(dst4, dst3, cv::COLOR_BGRA2BGR); + + EXPECT_MAT_NEAR(dst_gold, dst3, 0); +} + INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CvtColor, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, -- GitLab