generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

added OpenCL RGB <-> Luv conversions

Browse Source
This commit is contained in:
Ilya Lavrenov 2014-05-05 20:00:05 +04:00
parent dacf63a888
commit 3ccaa5294e
4 changed files with 299 additions and 33 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
modules/imgproc/perf/opencl/perf_color.cpp
View File

@ -59,7 +59,7 @@ using std::tr1::make_tuple;
CV_ENUM(ConversionTypes, COLOR_RGB2GRAY, COLOR_RGB2BGR, COLOR_RGB2YUV, COLOR_YUV2RGB, COLOR_RGB2YCrCb,
COLOR_YCrCb2RGB, COLOR_RGB2XYZ, COLOR_XYZ2RGB, COLOR_RGB2HSV, COLOR_HSV2RGB, COLOR_RGB2HLS,
COLOR_HLS2RGB, COLOR_BGR5652BGR, COLOR_BGR2BGR565, COLOR_RGBA2mRGBA, COLOR_mRGBA2RGBA, COLOR_YUV2RGB_NV12,
COLOR_RGB2Lab, COLOR_Lab2BGR)
COLOR_RGB2Lab, COLOR_Lab2BGR, COLOR_RGB2Luv, COLOR_Luv2LBGR)
typedef tuple<Size, tuple<ConversionTypes, int, int> > CvtColorParams;
typedef TestBaseWithParam<CvtColorParams> CvtColorFixture;
@ -85,7 +85,9 @@ OCL_PERF_TEST_P(CvtColorFixture, CvtColor, testing::Combine(
make_tuple(ConversionTypes(COLOR_mRGBA2RGBA), 4, 4),
make_tuple(ConversionTypes(COLOR_YUV2RGB_NV12), 1, 3),
make_tuple(ConversionTypes(COLOR_RGB2Lab), 3, 3),
make_tuple(ConversionTypes(COLOR_Lab2BGR), 3, 4)
make_tuple(ConversionTypes(COLOR_Lab2BGR), 3, 4),
make_tuple(ConversionTypes(COLOR_RGB2Luv), 3, 3),
make_tuple(ConversionTypes(COLOR_Luv2LBGR), 3, 4)
)))
{
CvtColorParams params = GetParam();

102
modules/imgproc/src/color.cpp
View File

@ -2749,12 +2749,13 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
ocl::Device dev = ocl::Device::getDefault();
int pxPerWIy = 1;
if (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU))
{
if (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU) &&
!(code == CV_BGR2Luv || code == CV_RGB2Luv || code == CV_LBGR2Luv || code == CV_LRGB2Luv ||
code == CV_Luv2BGR || code == CV_Luv2RGB || code == CV_Luv2LBGR || code == CV_Luv2LRGB))
pxPerWIy = 4;
}
globalsize[1] = DIVUP(globalsize[1], pxPerWIy);
opts += format("-D PIX_PER_WI_Y=%d ", pxPerWIy);
opts += format("-D PIX_PER_WI_Y=%d ", pxPerWIy);
switch (code)
{
@ -3084,16 +3085,20 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
break;
}
case CV_BGR2Lab: case CV_RGB2Lab: case CV_LBGR2Lab: case CV_LRGB2Lab:
case CV_BGR2Luv: case CV_RGB2Luv: case CV_LBGR2Luv: case CV_LRGB2Luv:
{
CV_Assert( (scn == 3 || scn == 4) && (depth == CV_8U || depth == CV_32F) );
bidx = code == CV_BGR2Lab || code == CV_LBGR2Lab ? 0 : 2;
bool srgb = code == CV_BGR2Lab || code == CV_RGB2Lab;
bidx = code == CV_BGR2Lab || code == CV_LBGR2Lab || code == CV_BGR2Luv || code == CV_LBGR2Luv ? 0 : 2;
bool srgb = code == CV_BGR2Lab || code == CV_RGB2Lab || code == CV_RGB2Luv || code == CV_BGR2Luv;
bool lab = code == CV_BGR2Lab || code == CV_RGB2Lab || code == CV_LBGR2Lab || code == CV_LRGB2Lab;
float un, vn;
dcn = 3;
k.create("BGR2Lab", ocl::imgproc::cvtcolor_oclsrc,
opts + format("-D dcn=3 -D bidx=%d%s",
bidx, srgb ? " -D SRGB" : ""));
k.create(format("BGR2%s", lab ? "Lab" : "Luv").c_str(),
ocl::imgproc::cvtcolor_oclsrc,
opts + format("-D dcn=%d -D bidx=%d%s",
dcn, bidx, srgb ? " -D SRGB" : ""));
if (k.empty())
return false;
@ -3105,7 +3110,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
dstarg = ocl::KernelArg::WriteOnly(dst);
if (depth == CV_8U)
if (depth == CV_8U && lab)
{
static UMat usRGBGammaTab, ulinearGammaTab, uLabCbrtTab, ucoeffs;
@ -3148,10 +3153,12 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
}
else
{
static UMat usRGBGammaTab, ucoeffs;
static UMat usRGBGammaTab, ucoeffs, uLabCbrtTab;
if (srgb && usRGBGammaTab.empty())
Mat(1, GAMMA_TAB_SIZE * 4, CV_32FC1, sRGBGammaTab).copyTo(usRGBGammaTab);
if (!lab && uLabCbrtTab.empty())
Mat(1, LAB_CBRT_TAB_SIZE * 4, CV_32FC1, LabCbrtTab).copyTo(uLabCbrtTab);
{
float coeffs[9];
@ -3161,39 +3168,59 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
for (int i = 0; i < 3; i++)
{
int j = i * 3;
coeffs[j + (bidx ^ 2)] = _coeffs[j] * scale[i];
coeffs[j + 1] = _coeffs[j + 1] * scale[i];
coeffs[j + bidx] = _coeffs[j + 2] * scale[i];
coeffs[j + (bidx ^ 2)] = _coeffs[j] * (lab ? scale[i] : 1);
coeffs[j + 1] = _coeffs[j + 1] * (lab ? scale[i] : 1);
coeffs[j + bidx] = _coeffs[j + 2] * (lab ? scale[i] : 1);
CV_Assert( coeffs[j] >= 0 && coeffs[j + 1] >= 0 && coeffs[j + 2] >= 0 &&
coeffs[j] + coeffs[j + 1] + coeffs[j + 2] < 1.5f*LabCbrtTabScale );
coeffs[j] + coeffs[j + 1] + coeffs[j + 2] < 1.5f*(lab ? LabCbrtTabScale : 1) );
}
float d = 1.f/(_whitept[0] + _whitept[1]*15 + _whitept[2]*3);
un = 13*4*_whitept[0]*d;
vn = 13*9*_whitept[1]*d;
Mat(1, 9, CV_32FC1, coeffs).copyTo(ucoeffs);
}
float _1_3 = 1.0f / 3.0f, _a = 16.0f / 116.0f;
ocl::KernelArg ucoeffsarg = ocl::KernelArg::PtrReadOnly(ucoeffs);
if (srgb)
k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBGammaTab),
ucoeffsarg, _1_3, _a);
if (lab)
{
if (srgb)
k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBGammaTab),
ucoeffsarg, _1_3, _a);
else
k.args(srcarg, dstarg, ucoeffsarg, _1_3, _a);
}
else
k.args(srcarg, dstarg, ucoeffsarg, _1_3, _a);
{
ocl::KernelArg LabCbrtTabarg = ocl::KernelArg::PtrReadOnly(uLabCbrtTab);
if (srgb)
k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBGammaTab),
LabCbrtTabarg, ucoeffsarg, un, vn);
else
k.args(srcarg, dstarg, LabCbrtTabarg, ucoeffsarg, un, vn);
}
}
return k.run(dims, globalsize, NULL, false);
}
case CV_Lab2BGR: case CV_Lab2RGB: case CV_Lab2LBGR: case CV_Lab2LRGB:
case CV_Luv2BGR: case CV_Luv2RGB: case CV_Luv2LBGR: case CV_Luv2LRGB:
{
if( dcn <= 0 )
dcn = 3;
CV_Assert( scn == 3 && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F) );
bidx = code == CV_Lab2BGR || code == CV_Lab2LBGR ? 0 : 2;
bool srgb = code == CV_Lab2BGR || code == CV_Lab2RGB;
bidx = code == CV_Lab2BGR || code == CV_Lab2LBGR || code == CV_Luv2BGR || code == CV_Luv2LBGR ? 0 : 2;
bool srgb = code == CV_Lab2BGR || code == CV_Lab2RGB || code == CV_Luv2BGR || code == CV_Luv2RGB;
bool lab = code == CV_Lab2BGR || code == CV_Lab2RGB || code == CV_Lab2LBGR || code == CV_Lab2LRGB;
float un, vn;
k.create("Lab2BGR", ocl::imgproc::cvtcolor_oclsrc,
k.create(format("%s2BGR", lab ? "Lab" : "Luv").c_str(),
ocl::imgproc::cvtcolor_oclsrc,
opts + format("-D dcn=%d -D bidx=%d%s",
dcn, bidx, srgb ? " -D SRGB" : ""));
if (k.empty())
@ -3211,11 +3238,15 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
for( int i = 0; i < 3; i++ )
{
coeffs[i+(bidx^2)*3] = _coeffs[i]*_whitept[i];
coeffs[i+3] = _coeffs[i+3]*_whitept[i];
coeffs[i+bidx*3] = _coeffs[i+6]*_whitept[i];
coeffs[i+(bidx^2)*3] = _coeffs[i] * (lab ? _whitept[i] : 1);
coeffs[i+3] = _coeffs[i+3] * (lab ? _whitept[i] : 1);
coeffs[i+bidx*3] = _coeffs[i+6] * (lab ? _whitept[i] : 1);
}
float d = 1.f/(_whitept[0] + _whitept[1]*15 + _whitept[2]*3);
un = 4*_whitept[0]*d;
vn = 9*_whitept[1]*d;
Mat(1, 9, CV_32FC1, coeffs).copyTo(ucoeffs);
}
@ -3229,11 +3260,22 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
dstarg = ocl::KernelArg::WriteOnly(dst),
coeffsarg = ocl::KernelArg::PtrReadOnly(ucoeffs);
if (srgb)
k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBInvGammaTab),
coeffsarg, lThresh, fThresh);
if (lab)
{
if (srgb)
k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBInvGammaTab),
coeffsarg, lThresh, fThresh);
else
k.args(srcarg, dstarg, coeffsarg, lThresh, fThresh);
}
else
k.args(srcarg, dstarg, coeffsarg, lThresh, fThresh);
{
if (srgb)
k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBInvGammaTab),
coeffsarg, un, vn);
else
k.args(srcarg, dstarg, coeffsarg, un, vn);
}
return k.run(dims, globalsize, NULL, false);
}
@ -3264,7 +3306,7 @@ void cv::cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
int stype = _src.type();
int scn = CV_MAT_CN(stype), depth = CV_MAT_DEPTH(stype), bidx;
CV_OCL_RUN( _src.dims() <= 2 && _dst.isUMat(),
CV_OCL_RUN( _src.dims() <= 2 && _dst.isUMat() && !(depth == CV_8U && (code == CV_Luv2BGR || code == CV_Luv2RGB)),
ocl_cvtColor(_src, _dst, code, dcn) )
Mat src = _src.getMat(), dst;

202
modules/imgproc/src/opencl/cvtcolor.cl
View File

@ -1537,3 +1537,205 @@ __kernel void Lab2BGR(__global const uchar * srcptr, int src_step, int src_offse
}
#endif
/////////////////////////////////// [l|s]RGB <-> Luv ///////////////////////////
#define LAB_CBRT_TAB_SIZE 1024
#define LAB_CBRT_TAB_SIZE_B (256*3/2*(1<<gamma_shift))
__constant float LabCbrtTabScale = LAB_CBRT_TAB_SIZE/1.5f;
#ifdef DEPTH_5
__kernel void BGR2Luv(__global const uchar * srcptr, int src_step, int src_offset,
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
#ifdef SRGB
__global const float * gammaTab,
#endif
__global const float * LabCbrtTab, __constant float * coeffs, float _un, float _vn)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < cols && y < rows)
{
int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset));
int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
__global const float * src = (__global const float *)(srcptr + src_idx);
__global float * dst = (__global float *)(dstptr + dst_idx);
float R = src[0], G = src[1], B = src[2];
#ifdef SRGB
R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
#endif
float X = R*coeffs[0] + G*coeffs[1] + B*coeffs[2];
float Y = R*coeffs[3] + G*coeffs[4] + B*coeffs[5];
float Z = R*coeffs[6] + G*coeffs[7] + B*coeffs[8];
float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE);
L = 116.f*L - 16.f;
float d = (4*13) / max(X + 15 * Y + 3 * Z, FLT_EPSILON);
float u = L*(X*d - _un);
float v = L*((9*0.25f)*Y*d - _vn);
dst[0] = L;
dst[1] = u;
dst[2] = v;
}
}
#elif defined DEPTH_0
__kernel void BGR2Luv(__global const uchar * src, int src_step, int src_offset,
__global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
#ifdef SRGB
__global const float * gammaTab,
#endif
__global const float * LabCbrtTab, __constant float * coeffs, float _un, float _vn)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < cols && y < rows)
{
int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset));
int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
src += src_idx;
dst += dst_idx;
float scale = 1.0f / 255.0f;
float R = src[0]*scale, G = src[1]*scale, B = src[2]*scale;
#ifdef SRGB
R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
#endif
float X = R*coeffs[0] + G*coeffs[1] + B*coeffs[2];
float Y = R*coeffs[3] + G*coeffs[4] + B*coeffs[5];
float Z = R*coeffs[6] + G*coeffs[7] + B*coeffs[8];
float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE);
L = 116.f*L - 16.f;
float d = (4*13) / max(X + 15 * Y + 3 * Z, FLT_EPSILON);
float u = L*(X*d - _un);
float v = L*((9*0.25f)*Y*d - _vn);
dst[0] = SAT_CAST(L * 2.55f);
dst[1] = SAT_CAST(mad(u, 0.72033898305084743f, 96.525423728813564f));
dst[2] = SAT_CAST(mad(v, 0.99609375f, 139.453125f));
}
}
#endif
#ifdef DEPTH_5
__kernel void Luv2BGR(__global const uchar * srcptr, int src_step, int src_offset,
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
#ifdef SRGB
__global const float * gammaTab,
#endif
__constant float * coeffs, float _un, float _vn)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < cols && y < rows)
{
int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset));
int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
__global const float * src = (__global const float *)(srcptr + src_idx);
__global float * dst = (__global float *)(dstptr + dst_idx);
float L = src[0], u = src[1], v = src[2], d, X, Y, Z;
Y = (L + 16.f) * (1.f/116.f);
Y = Y*Y*Y;
d = (1.f/13.f)/L;
u = u*d + _un;
v = v*d + _vn;
float iv = 1.f/v;
X = 2.25f * u * Y * iv ;
Z = (12 - 3 * u - 20 * v) * Y * 0.25f * iv;
float R = X*coeffs[0] + Y*coeffs[1] + Z*coeffs[2];
float G = X*coeffs[3] + Y*coeffs[4] + Z*coeffs[5];
float B = X*coeffs[6] + Y*coeffs[7] + Z*coeffs[8];
#ifdef SRGB
R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
#endif
dst[0] = R;
dst[1] = G;
dst[2] = B;
#if dcn == 4
dst[3] = MAX_NUM;
#endif
}
}
#elif defined DEPTH_0
__kernel void Luv2BGR(__global const uchar * src, int src_step, int src_offset,
__global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
#ifdef SRGB
__global const float * gammaTab,
#endif
__constant float * coeffs, float _un, float _vn)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < cols && y < rows)
{
int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset));
int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
src += src_idx;
dst += dst_idx;
float d, X, Y, Z;
float L = src[0]*(100.f/255.f);
float u = (float)(src[1]*1.388235294117647f - 134.f);
float v = (float)(src[2]*1.003921568627451f - 140.f);
Y = (L + 16.f) * (1.f/116.f);
Y = Y*Y*Y;
d = (1.f/13.f)/L;
u = u*d + _un;
v = v*d + _vn;
float iv = 1.f/v;
X = 2.25f * u * Y * iv ;
Z = (12 - 3 * u - 20 * v) * Y * 0.25f * iv;
float R = X*coeffs[0] + Y*coeffs[1] + Z*coeffs[2];
float G = X*coeffs[3] + Y*coeffs[4] + Z*coeffs[5];
float B = X*coeffs[6] + Y*coeffs[7] + Z*coeffs[8];
#ifdef SRGB
R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
#endif
dst[0] = SAT_CAST(R * 255.0f);
dst[1] = SAT_CAST(G * 255.0f);
dst[2] = SAT_CAST(B * 255.0f);
#if dcn == 4
dst[3] = MAX_NUM;
#endif
}
}
#endif

22
modules/imgproc/test/ocl/test_color.cpp
View File

@ -156,7 +156,7 @@ OCL_TEST_P(CvtColor, YCrCb2BGRA) { performTest(3, 4, CVTCODE(YCrCb2BGR)); }
// RGB <-> XYZ
#if IPP_VERSION_X100 > 0
#define IPP_EPS depth <= CV_32S ? 1 : 4e-5
#define IPP_EPS depth <= CV_32S ? 1 : 5e-5
#else
#define IPP_EPS 1e-3
#endif
@ -300,6 +300,26 @@ OCL_TEST_P(CvtColor8u32f, Lab2RGBA) { performTest(3, 4, CVTCODE(Lab2RGB), depth
OCL_TEST_P(CvtColor8u32f, Lab2LBGRA) { performTest(3, 4, CVTCODE(Lab2LBGR), depth == CV_8U ? 1 : 1e-5); }
OCL_TEST_P(CvtColor8u32f, Lab2LRGBA) { performTest(3, 4, CVTCODE(Lab2LRGB), depth == CV_8U ? 1 : 1e-5); }
// RGB -> Luv
OCL_TEST_P(CvtColor8u32f, BGR2Luv) { performTest(3, 3, CVTCODE(BGR2Luv), depth == CV_8U ? 1 : 1e-2); }
OCL_TEST_P(CvtColor8u32f, RGB2Luv) { performTest(3, 3, CVTCODE(RGB2Luv), depth == CV_8U ? 1 : 1e-2); }
OCL_TEST_P(CvtColor8u32f, LBGR2Luv) { performTest(3, 3, CVTCODE(LBGR2Luv), depth == CV_8U ? 1 : 3e-3); }
OCL_TEST_P(CvtColor8u32f, LRGB2Luv) { performTest(3, 3, CVTCODE(LRGB2Luv), depth == CV_8U ? 1 : 4e-3); }
OCL_TEST_P(CvtColor8u32f, BGRA2Luv) { performTest(4, 3, CVTCODE(BGR2Luv), depth == CV_8U ? 1 : 8e-3); }
OCL_TEST_P(CvtColor8u32f, RGBA2Luv) { performTest(4, 3, CVTCODE(RGB2Luv), depth == CV_8U ? 1 : 7e-3); }
OCL_TEST_P(CvtColor8u32f, LBGRA2Luv) { performTest(4, 3, CVTCODE(LBGR2Luv), depth == CV_8U ? 1 : 4e-3); }
OCL_TEST_P(CvtColor8u32f, LRGBA2Luv) { performTest(4, 3, CVTCODE(LRGB2Luv), depth == CV_8U ? 1 : 4e-3); }
OCL_TEST_P(CvtColor8u32f, Luv2BGR) { performTest(3, 3, CVTCODE(Luv2BGR), depth == CV_8U ? 1 : 7e-5); }
OCL_TEST_P(CvtColor8u32f, Luv2RGB) { performTest(3, 3, CVTCODE(Luv2RGB), depth == CV_8U ? 1 : 7e-5); }
OCL_TEST_P(CvtColor8u32f, Luv2LBGR) { performTest(3, 3, CVTCODE(Luv2LBGR), depth == CV_8U ? 1 : 1e-5); }
OCL_TEST_P(CvtColor8u32f, Luv2LRGB) { performTest(3, 3, CVTCODE(Luv2LRGB), depth == CV_8U ? 1 : 1e-5); }
OCL_TEST_P(CvtColor8u32f, Luv2BGRA) { performTest(3, 4, CVTCODE(Luv2BGR), depth == CV_8U ? 1 : 7e-5); }
OCL_TEST_P(CvtColor8u32f, Luv2RGBA) { performTest(3, 4, CVTCODE(Luv2RGB), depth == CV_8U ? 1 : 7e-5); }
OCL_TEST_P(CvtColor8u32f, Luv2LBGRA) { performTest(3, 4, CVTCODE(Luv2LBGR), depth == CV_8U ? 1 : 1e-5); }
OCL_TEST_P(CvtColor8u32f, Luv2LRGBA) { performTest(3, 4, CVTCODE(Luv2LRGB), depth == CV_8U ? 1 : 1e-5); }
// YUV -> RGBA_NV12
struct CvtColor_YUV420 :