Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

Merge pull request #633 from jet47:gpu-debayer-gray

Browse Source
pull/632/head^2
Andrey Kamaev 12 years ago committed by OpenCV Buildbot
parent 632211eb27 4ddf634c30
commit 1d3219747a
4 changed files with 248 additions and 89 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 7
      modules/gpu/perf/perf_imgproc.cpp
  2. 45
      modules/gpu/src/color.cpp
  3. 221
      modules/gpu/src/cuda/debayer.cu
  4. 64
      modules/gpu/test/test_color.cpp

7
modules/gpu/perf/perf_imgproc.cpp
Unescape View File

@ -1341,7 +1341,12 @@ PERF_TEST_P(Sz_Depth_Code, ImgProc_CvtColorBayer,
Values(CvtColorInfo(1, 3, cv::COLOR_BayerBG2BGR), Values(CvtColorInfo(1, 3, cv::COLOR_BayerBG2BGR),
CvtColorInfo(1, 3, cv::COLOR_BayerGB2BGR), CvtColorInfo(1, 3, cv::COLOR_BayerGB2BGR),
CvtColorInfo(1, 3, cv::COLOR_BayerRG2BGR), 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 cv::Size size = GET_PARAM(0);
const int depth = GET_PARAM(1); const int depth = GET_PARAM(1);

45
modules/gpu/src/color.cpp
Unescape View File

@ -1640,6 +1640,43 @@ namespace
{ {
bayer_to_bgr(src, dst, dcn, true, true, 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);
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) 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_bgr, // CV_YUV2BGR = 84
yuv_to_rgb, // CV_YUV2RGB = 85 yuv_to_rgb, // CV_YUV2RGB = 85
0, // CV_BayerBG2GRAY = 86 bayerBG_to_gray, // CV_BayerBG2GRAY = 86
0, // CV_BayerGB2GRAY = 87 bayerGB_to_gray, // CV_BayerGB2GRAY = 87
0, // CV_BayerRG2GRAY = 88 bayerRG_to_gray, // CV_BayerRG2GRAY = 88
0, // CV_BayerGR2GRAY = 89 bayerGR_to_gray, // CV_BayerGR2GRAY = 89
//YUV 4:2:0 formats family //YUV 4:2:0 formats family
0, // CV_YUV2RGB_NV12 = 90, 0, // CV_YUV2RGB_NV12 = 90,

221
modules/gpu/src/cuda/debayer.cu
Unescape View File

@ -42,42 +42,37 @@
#if !defined CUDA_DISABLER #if !defined CUDA_DISABLER
#include <opencv2/gpu/device/common.hpp> #include "opencv2/gpu/device/common.hpp"
#include <opencv2/gpu/device/vec_traits.hpp> #include "opencv2/gpu/device/vec_traits.hpp"
#include <opencv2/gpu/device/vec_math.hpp> #include "opencv2/gpu/device/vec_math.hpp"
#include <opencv2/gpu/device/limits.hpp> #include "opencv2/gpu/device/limits.hpp"
#include "opencv2/gpu/device/color.hpp"
namespace cv { namespace gpu { namespace cv { namespace gpu { namespace device
namespace device {
{ template <typename T> struct Bayer2BGR;
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;
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]; uchar4 patch[3][3];
patch[0][1] = ((const uchar4*) src.ptr(s_y - 1))[s_x]; 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][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][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][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][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][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)]; patch[2][2] = ((const uchar4*) src.ptr(s_y + 1))[::min(s_x + 1, ((src.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());
if ((s_y & 1) ^ start_with_green) if ((s_y & 1) ^ start_with_green)
{ {
@ -181,45 +176,69 @@ namespace cv { namespace gpu {
res3.z = t7; 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; if (s_y >= src.rows || (s_x << 2) >= src.cols)
const int d_y = blockIdx.y * blockDim.y + threadIdx.y; return;
dst(d_y, d_x) = res0; s_y = ::min(::max(s_y, 1), src.rows - 2);
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 <typename D> Bayer2BGR<uchar> bayer;
__global__ void Bayer2BGR_16u(const PtrStepb src, PtrStepSz<D> dst, const bool blue_last, const bool start_with_green) bayer.apply(src, s_x, s_y, blue_last, 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 << 1) >= dst.cols) const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
return; 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]; ushort2 patch[3][3];
patch[0][1] = ((const ushort2*) src.ptr(s_y - 1))[s_x]; 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][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][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][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][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][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)]; patch[2][2] = ((const ushort2*) src.ptr(s_y + 1))[::min(s_x + 1, ((src.cols + 1) >> 1) - 1)];
D res0 = VecTraits<D>::all(numeric_limits<ushort>::max());
D res1 = VecTraits<D>::all(numeric_limits<ushort>::max());
if ((s_y & 1) ^ start_with_green) if ((s_y & 1) ^ start_with_green)
{ {
@ -279,53 +298,87 @@ namespace cv { namespace gpu {
res1.z = t3; res1.z = t3;
} }
} }
}
};
const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 1; template <typename D> __device__ __forceinline__ D toDst(const ushort3& pix);
const int d_y = blockIdx.y * blockDim.y + threadIdx.y; 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; template <typename D>
if (d_x + 1 < dst.cols) __global__ void Bayer2BGR_16u(const PtrStepSzb src, PtrStep<D> dst, const bool blue_last, const bool start_with_green)
dst(d_y, d_x + 1) = res1; {
} const int s_x = blockIdx.x * blockDim.x + threadIdx.x;
int s_y = blockIdx.y * blockDim.y + threadIdx.y;
template <int cn> if (s_y >= src.rows || (s_x << 1) >= src.cols)
void Bayer2BGR_8u_gpu(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream) return;
{
typedef typename TypeVec<uchar, cn>::vec_type dst_t;
const dim3 block(32, 8); s_y = ::min(::max(s_y, 1), src.rows - 2);
const dim3 grid(divUp(dst.cols, 4 * block.x), divUp(dst.rows, block.y));
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); const int d_x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
cudaSafeCall( cudaGetLastError() ); const int d_y = blockIdx.y * blockDim.y + threadIdx.y;
if (stream == 0) dst(d_y, d_x) = toDst<D>(bayer.res0);
cudaSafeCall( cudaDeviceSynchronize() ); if (d_x + 1 < src.cols)
} dst(d_y, d_x + 1) = toDst<D>(bayer.res1);
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); template <int cn>
const dim3 grid(divUp(dst.cols, 2 * block.x), divUp(dst.rows, block.y)); 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( cudaFuncSetCacheConfig(Bayer2BGR_8u<dst_t>, cudaFuncCachePreferL1) );
cudaSafeCall( cudaGetLastError() );
if (stream == 0) Bayer2BGR_8u<dst_t><<<grid, block, 0, stream>>>(src, (PtrStepSz<dst_t>)dst, blue_last, start_with_green);
cudaSafeCall( cudaDeviceSynchronize() ); cudaSafeCall( cudaGetLastError() );
}
template void Bayer2BGR_8u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); if (stream == 0)
template void Bayer2BGR_8u_gpu<4>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); cudaSafeCall( cudaDeviceSynchronize() );
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 */ 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 */

64
modules/gpu/test/test_color.cpp
Unescape View File

@ -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); 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( INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CvtColor, testing::Combine(
ALL_DEVICES, ALL_DEVICES,
DIFFERENT_SIZES, DIFFERENT_SIZES,

Write Preview
Loading…
Cancel
Save