Skip to content

O
opencv
Commit 4ddf634c authored by Vladislav Vinogradov's avatar Vladislav Vinogradov
Browse files
Options

gpu : implement Bayer* -> Gray color conversion

parent 13f402a5
Hide whitespace changes
Inline Side-by-side
Showing with 248 additions and 90 deletions
modules/gpu/perf/perf_imgproc.cpp
View file @ 4ddf634c
......@@ -1341,7 +1341,12 @@ PERF_TEST_P(Sz_Depth_Code, ImgProc_CvtColorBayer,
Values(CvtColorInfo(1, 3, cv::COLOR_BayerBG2BGR),
CvtColorInfo(1, 3, cv::COLOR_BayerGB2BGR),
CvtColorInfo(1, 3, cv::COLOR_BayerRG2BGR),
CvtColorInfo(1, 3, cv::COLOR_BayerGR2BGR))))
CvtColorInfo(1, 3, cv::COLOR_BayerGR2BGR),
CvtColorInfo(1, 1, cv::COLOR_BayerBG2GRAY),
CvtColorInfo(1, 1, cv::COLOR_BayerGB2GRAY),
CvtColorInfo(1, 1, cv::COLOR_BayerRG2GRAY),
CvtColorInfo(1, 1, cv::COLOR_BayerGR2GRAY))))
{
const cv::Size size = GET_PARAM(0);
const int depth = GET_PARAM(1);
......
modules/gpu/src/color.cpp
View file @ 4ddf634c
......@@ -1640,6 +1640,43 @@ namespace
{
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);
static const func_t funcs[3] =
{
Bayer2BGR_8u_gpu<1>,
0,
Bayer2BGR_16u_gpu<1>,
};
CV_Assert(src.type() == CV_8UC1 || src.type() == CV_16UC1);
CV_Assert(src.rows > 2 && src.cols > 2);
dst.create(src.size(), CV_MAKETYPE(src.depth(), 1));
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);
}
}
void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream)
......@@ -1756,10 +1793,10 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
yuv_to_bgr, // CV_YUV2BGR = 84
yuv_to_rgb, // CV_YUV2RGB = 85
0, // CV_BayerBG2GRAY = 86
0, // CV_BayerGB2GRAY = 87
0, // CV_BayerRG2GRAY = 88
0, // CV_BayerGR2GRAY = 89
bayerBG_to_gray, // CV_BayerBG2GRAY = 86
bayerGB_to_gray, // CV_BayerGB2GRAY = 87
bayerRG_to_gray, // CV_BayerRG2GRAY = 88
bayerGR_to_gray, // CV_BayerGR2GRAY = 89
//YUV 4:2:0 formats family
0, // CV_YUV2RGB_NV12 = 90,
......
modules/gpu/src/cuda/debayer.cu
View file @ 4ddf634c
......@@ -42,42 +42,37 @@
#if !defined CUDA_DISABLER
#include <opencv2/gpu/device/common.hpp>
#include <opencv2/gpu/device/vec_traits.hpp>
#include <opencv2/gpu/device/vec_math.hpp>
#include <opencv2/gpu/device/limits.hpp>
#include "opencv2/gpu/device/common.hpp"
#include "opencv2/gpu/device/vec_traits.hpp"
#include "opencv2/gpu/device/vec_math.hpp"
#include "opencv2/gpu/device/limits.hpp"
#include "opencv2/gpu/device/color.hpp"
namespace cv { namespace gpu {
namespace device
{
template <typename D>
__global__ void Bayer2BGR_8u(const PtrStepb src, PtrStepSz<D> dst, const bool blue_last, const bool start_with_green)
{
const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
if (s_y >= dst.rows || (s_x << 2) >= dst.cols)
return;
namespace cv { namespace gpu { namespace device
{
template <typename T> struct Bayer2BGR;
s_y = ::min(::max(s_y, 1), dst.rows - 2);
template <> struct Bayer2BGR<uchar>
{
uchar3 res0;
uchar3 res1;
uchar3 res2;
uchar3 res3;
__device__ void apply(const PtrStepSzb& src, int s_x, int s_y, bool blue_last, bool start_with_green)
{
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)];
patch[0][2] = ((const uchar4*) src.ptr(s_y - 1))[::min(s_x + 1, ((src.cols + 3) >> 2) - 1)];
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[1][2] = ((const uchar4*) src.ptr(s_y))[::min(s_x + 1, ((src.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<D>::all(numeric_limits<uchar>::max());
D res1 = VecTraits<D>::all(numeric_limits<uchar>::max());
D res2 = VecTraits<D>::all(numeric_limits<uchar>::max());
D res3 = VecTraits<D>::all(numeric_limits<uchar>::max());
patch[2][2] = ((const uchar4*) src.ptr(s_y + 1))[::min(s_x + 1, ((src.cols + 3) >> 2) - 1)];
if ((s_y & 1) ^ start_with_green)
{
......@@ -181,45 +176,69 @@ namespace cv { namespace gpu {
res3.z = t7;
}
}
}
};
template <typename D> __device__ __forceinline__ D toDst(const uchar3& pix);
template <> __device__ __forceinline__ uchar toDst<uchar>(const uchar3& pix)
{
typename bgr_to_gray_traits<uchar>::functor_type f = bgr_to_gray_traits<uchar>::create_functor();
return f(pix);
}
template <> __device__ __forceinline__ uchar3 toDst<uchar3>(const uchar3& pix)
{
return pix;
}
template <> __device__ __forceinline__ uchar4 toDst<uchar4>(const uchar3& pix)
{
return make_uchar4(pix.x, pix.y, pix.z, 255);
}
template <typename D>
__global__ void Bayer2BGR_8u(const PtrStepSzb src, PtrStep<D> dst, const bool blue_last, const bool start_with_green)
{
const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
if (s_y >= src.rows || (s_x << 2) >= src.cols)
return;
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;
}
s_y = ::min(::max(s_y, 1), src.rows - 2);
template <typename D>
__global__ void Bayer2BGR_16u(const PtrStepb src, PtrStepSz<D> dst, const bool blue_last, const bool start_with_green)
{
const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
Bayer2BGR<uchar> bayer;
bayer.apply(src, s_x, s_y, blue_last, start_with_green);
if (s_y >= dst.rows || (s_x << 1) >= dst.cols)
return;
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
s_y = ::min(::max(s_y, 1), dst.rows - 2);
dst(d_y, d_x) = toDst<D>(bayer.res0);
if (d_x + 1 < src.cols)
dst(d_y, d_x + 1) = toDst<D>(bayer.res1);
if (d_x + 2 < src.cols)
dst(d_y, d_x + 2) = toDst<D>(bayer.res2);
if (d_x + 3 < src.cols)
dst(d_y, d_x + 3) = toDst<D>(bayer.res3);
}
template <> struct Bayer2BGR<ushort>
{
ushort3 res0;
ushort3 res1;
__device__ void apply(const PtrStepSzb& src, int s_x, int s_y, bool blue_last, bool start_with_green)
{
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[0][2] = ((const ushort2*) src.ptr(s_y - 1))[::min(s_x + 1, ((src.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[1][2] = ((const ushort2*) src.ptr(s_y))[::min(s_x + 1, ((src.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<D>::all(numeric_limits<ushort>::max());
D res1 = VecTraits<D>::all(numeric_limits<ushort>::max());
patch[2][2] = ((const ushort2*) src.ptr(s_y + 1))[::min(s_x + 1, ((src.cols + 1) >> 1) - 1)];
if ((s_y & 1) ^ start_with_green)
{
......@@ -279,53 +298,87 @@ namespace cv { namespace gpu {
res1.z = t3;
}
}
}
};
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
template <typename D> __device__ __forceinline__ D toDst(const ushort3& pix);
template <> __device__ __forceinline__ ushort toDst<ushort>(const ushort3& pix)
{
typename bgr_to_gray_traits<ushort>::functor_type f = bgr_to_gray_traits<ushort>::create_functor();
return f(pix);
}
template <> __device__ __forceinline__ ushort3 toDst<ushort3>(const ushort3& pix)
{
return pix;
}
template <> __device__ __forceinline__ ushort4 toDst<ushort4>(const ushort3& pix)
{
return make_ushort4(pix.x, pix.y, pix.z, numeric_limits<ushort>::max());
}
dst(d_y, d_x) = res0;
if (d_x + 1 < dst.cols)
dst(d_y, d_x + 1) = res1;
}
template <typename D>
__global__ void Bayer2BGR_16u(const PtrStepSzb src, PtrStep<D> dst, const bool blue_last, const bool start_with_green)
{
const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
template <int cn>
void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream)
{
typedef typename TypeVec<uchar, cn>::vec_type dst_t;
if (s_y >= src.rows || (s_x << 1) >= src.cols)
return;
const dim3 block(32, 8);
const dim3 grid(divUp(dst.cols, 4 * block.x), divUp(dst.rows, block.y));
s_y = ::min(::max(s_y, 1), src.rows - 2);
cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u<dst_t>, cudaFuncCachePreferL1) );
Bayer2BGR<ushort> bayer;
bayer.apply(src, s_x, s_y, blue_last, start_with_green);
Bayer2BGR_8u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
cudaSafeCall( cudaGetLastError() );
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
template <int cn>
void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream)
{
typedef typename TypeVec<ushort, cn>::vec_type dst_t;
dst(d_y, d_x) = toDst<D>(bayer.res0);
if (d_x + 1 < src.cols)
dst(d_y, d_x + 1) = toDst<D>(bayer.res1);
}
const dim3 block(32, 8);
const dim3 grid(divUp(dst.cols, 2 * block.x), divUp(dst.rows, block.y));
template <int cn>
void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream)
{
typedef typename TypeVec<uchar, cn>::vec_type dst_t;
cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u<dst_t>, cudaFuncCachePreferL1) );
const dim3 block(32, 8);
const dim3 grid(divUp(src.cols, 4 * block.x), divUp(src.rows, block.y));
Bayer2BGR_16u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
cudaSafeCall( cudaGetLastError() );
cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_8u<dst_t>, cudaFuncCachePreferL1) );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
Bayer2BGR_8u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
cudaSafeCall( cudaGetLastError() );
template void Bayer2BGR_8u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
template void Bayer2BGR_8u_gpu<4>(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);
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
}}
#endif /* CUDA_DISABLER */
\ No newline at end of file
template <int cn>
void Bayer2BGR_16u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream)
{
typedef typename TypeVec<ushort, cn>::vec_type dst_t;
const dim3 block(32, 8);
const dim3 grid(divUp(src.cols, 2 * block.x), divUp(src.rows, block.y));
cudaSafeCall( cudaFuncSetCacheConfig(Bayer2BGR_16u<dst_t>, cudaFuncCachePreferL1) );
Bayer2BGR_16u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
template void Bayer2BGR_8u_gpu<1>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
template void Bayer2BGR_8u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
template void Bayer2BGR_8u_gpu<4>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream);
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);
}}}
#endif /* CUDA_DISABLER */
modules/gpu/test/test_color.cpp
View file @ 4ddf634c
......@@ -2218,6 +2218,70 @@ GPU_TEST_P(CvtColor, BayerGR2BGR4)
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);
}
GPU_TEST_P(CvtColor, BayerBG2Gray)
{
if ((depth != CV_8U && depth != CV_16U) || useRoi)
return;
cv::Mat src = randomMat(size, depth);
cv::gpu::GpuMat dst;
cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerBG2GRAY);
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_BayerBG2GRAY);
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)), 2);
}
GPU_TEST_P(CvtColor, BayerGB2Gray)
{
if ((depth != CV_8U && depth != CV_16U) || useRoi)
return;
cv::Mat src = randomMat(size, depth);
cv::gpu::GpuMat dst;
cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGB2GRAY);
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_BayerGB2GRAY);
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)), 2);
}
GPU_TEST_P(CvtColor, BayerRG2Gray)
{
if ((depth != CV_8U && depth != CV_16U) || useRoi)
return;
cv::Mat src = randomMat(size, depth);
cv::gpu::GpuMat dst;
cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerRG2GRAY);
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_BayerRG2GRAY);
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)), 2);
}
GPU_TEST_P(CvtColor, BayerGR2Gray)
{
if ((depth != CV_8U && depth != CV_16U) || useRoi)
return;
cv::Mat src = randomMat(size, depth);
cv::gpu::GpuMat dst;
cv::gpu::cvtColor(loadMat(src, useRoi), dst, cv::COLOR_BayerGR2GRAY);
cv::Mat dst_gold;
cv::cvtColor(src, dst_gold, cv::COLOR_BayerGR2GRAY);
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)), 2);
}
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CvtColor, testing::Combine(
ALL_DEVICES,
DIFFERENT_SIZES,
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment