Merge pull request #644 from jet47:gpu-debayer-mht
This commit is contained in:
cuda-geek
OpenCV Buildbot
commit
1d626194af
@ -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
|
||||
|
||||
@ -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
|
||||
|
||||
|
||||
@ -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);
|
||||
|
||||
@ -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 */
|
||||
|
||||
@ -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
|
||||
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user