diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 82ed8044ff..b807afd41c 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -234,8 +234,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) toRGB5x5_caller(src, dst, -1, greenbits, "Gray2BGR5x5"); break; } - case CV_RGB2GRAY: case CV_BGR2GRAY: - case CV_RGBA2GRAY: case CV_BGRA2GRAY: + case CV_RGB2GRAY: case CV_BGR2GRAY: case CV_RGBA2GRAY: case CV_BGRA2GRAY: { CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2; @@ -243,8 +242,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) fromRGB_caller(src, dst, bidx, "RGB2Gray"); break; } - case CV_GRAY2BGR: - case CV_GRAY2BGRA: + case CV_GRAY2BGR: case CV_GRAY2BGRA: { CV_Assert(scn == 1); dcn = code == CV_GRAY2BGRA ? 4 : 3; @@ -252,8 +250,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) toRGB_caller(src, dst, 0, "Gray2RGB"); break; } - case CV_BGR2YUV: - case CV_RGB2YUV: + case CV_BGR2YUV: case CV_RGB2YUV: { CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2YUV ? 0 : 2; @@ -261,8 +258,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) fromRGB_caller(src, dst, bidx, "RGB2YUV"); break; } - case CV_YUV2BGR: - case CV_YUV2RGB: + case CV_YUV2BGR: case CV_YUV2RGB: { if( dcn <= 0 ) dcn = 3; @@ -285,8 +281,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) toRGB_caller(src, dst, bidx, "YUV2RGBA_NV12"); break; } - case CV_BGR2YCrCb: - case CV_RGB2YCrCb: + case CV_BGR2YCrCb: case CV_RGB2YCrCb: { CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2YCrCb ? 0 : 2; @@ -294,8 +289,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) fromRGB_caller(src, dst, bidx, "RGB2YCrCb"); break; } - case CV_YCrCb2BGR: - case CV_YCrCb2RGB: + case CV_YCrCb2BGR: case CV_YCrCb2RGB: { if( dcn <= 0 ) dcn = 3; @@ -305,8 +299,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) toRGB_caller(src, dst, bidx, "YCrCb2RGB"); break; } - case CV_BGR2XYZ: - case CV_RGB2XYZ: + case CV_BGR2XYZ: case CV_RGB2XYZ: { CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2XYZ ? 0 : 2; @@ -350,8 +343,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) fromRGB_caller(src, dst, bidx, "RGB2XYZ", "", oclCoeffs); break; } - case CV_XYZ2BGR: - case CV_XYZ2RGB: + case CV_XYZ2BGR: case CV_XYZ2RGB: { if (dcn <= 0) dcn = 3; diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index ae271be1da..70922c3ff9 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -805,6 +805,59 @@ __kernel void RGB2HLS(int cols, int rows, int src_step, int dst_step, int bidx, } } +__kernel void HLS2RGB(int cols, int rows, int src_step, int dst_step, int bidx, + __global const uchar * src, __global uchar * dst, + int src_offset, int dst_offset) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (y < rows && x < cols) + { + x <<= 2; + int src_idx = mad24(y, src_step, src_offset + x); + int dst_idx = mad24(y, dst_step, dst_offset + x); + + float h = src[src_idx], l = src[src_idx + 1]*(1.f/255.f), s = src[src_idx + 2]*(1.f/255.f); + float b, g, r; + + if (s != 0) + { + float tab[4]; + + float p2 = l <= 0.5f ? l*(1 + s) : l + s - l*s; + float p1 = 2*l - p2; + + h *= hscale; + if( h < 0 ) + do h += 6; while( h < 0 ); + else if( h >= 6 ) + do h -= 6; while( h >= 6 ); + + int sector = convert_int_sat_rtn(h); + h -= sector; + + tab[0] = p2; + tab[1] = p1; + tab[2] = p1 + (p2 - p1)*(1-h); + tab[3] = p1 + (p2 - p1)*h; + + b = tab[sector_data[sector][0]]; + g = tab[sector_data[sector][1]]; + r = tab[sector_data[sector][2]]; + } + else + b = g = r = l; + + dst[dst_idx + bidx] = convert_uchar_sat_rte(b*255.f); + dst[dst_idx + 1] = convert_uchar_sat_rte(g*255.f); + dst[dst_idx + (bidx^2)] = convert_uchar_sat_rte(r*255.f); +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} + #elif defined DEPTH_5 __kernel void RGB2HLS(int cols, int rows, int src_step, int dst_step, int bidx, @@ -854,4 +907,58 @@ __kernel void RGB2HLS(int cols, int rows, int src_step, int dst_step, int bidx, } } +__kernel void HLS2RGB(int cols, int rows, int src_step, int dst_step, int bidx, + __global const float * src, __global float * dst, + int src_offset, int dst_offset) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (y < rows && x < cols) + { + x <<= 2; + int src_idx = mad24(y, src_step, src_offset + x); + int dst_idx = mad24(y, dst_step, dst_offset + x); + + float h = src[src_idx], l = src[src_idx + 1], s = src[src_idx + 2]; + float b, g, r; + + if (s != 0) + { + float tab[4]; + int sector; + + float p2 = l <= 0.5f ? l*(1 + s) : l + s - l*s; + float p1 = 2*l - p2; + + h *= hscale; + if( h < 0 ) + do h += 6; while( h < 0 ); + else if( h >= 6 ) + do h -= 6; while( h >= 6 ); + + sector = convert_int_sat_rtn(h); + h -= sector; + + tab[0] = p2; + tab[1] = p1; + tab[2] = p1 + (p2 - p1)*(1-h); + tab[3] = p1 + (p2 - p1)*h; + + b = tab[sector_data[sector][0]]; + g = tab[sector_data[sector][1]]; + r = tab[sector_data[sector][2]]; + } + else + b = g = r = l; + + dst[dst_idx + bidx] = b; + dst[dst_idx + 1] = g; + dst[dst_idx + (bidx^2)] = r; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} + #endif diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 2b798d5c57..935e974628 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -124,17 +124,17 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) // RGB[A] <-> BGR[A] OCL_TEST_P(CvtColor, BGR2BGRA) { doTest(3, 4, CVTCODE(BGR2BGRA)); } -OCL_TEST_P(CvtColor, RGB2RGBA) { doTest(3, 4, CVTCODE(BGR2BGRA)); } +OCL_TEST_P(CvtColor, RGB2RGBA) { doTest(3, 4, CVTCODE(RGB2RGBA)); } OCL_TEST_P(CvtColor, BGRA2BGR) { doTest(4, 3, CVTCODE(BGRA2BGR)); } -OCL_TEST_P(CvtColor, RGBA2RGB) { doTest(4, 3, CVTCODE(BGRA2BGR)); } +OCL_TEST_P(CvtColor, RGBA2RGB) { doTest(4, 3, CVTCODE(RGBA2RGB)); } OCL_TEST_P(CvtColor, BGR2RGBA) { doTest(3, 4, CVTCODE(BGR2RGBA)); } -OCL_TEST_P(CvtColor, RGB2BGRA) { doTest(3, 4, CVTCODE(BGR2RGBA)); } +OCL_TEST_P(CvtColor, RGB2BGRA) { doTest(3, 4, CVTCODE(RGB2BGRA)); } OCL_TEST_P(CvtColor, RGBA2BGR) { doTest(4, 3, CVTCODE(RGBA2BGR)); } -OCL_TEST_P(CvtColor, BGRA2RGB) { doTest(4, 3, CVTCODE(RGBA2BGR)); } +OCL_TEST_P(CvtColor, BGRA2RGB) { doTest(4, 3, CVTCODE(BGRA2RGB)); } OCL_TEST_P(CvtColor, BGR2RGB) { doTest(3, 3, CVTCODE(BGR2RGB)); } -OCL_TEST_P(CvtColor, RGB2BGR) { doTest(3, 3, CVTCODE(BGR2RGB)); } +OCL_TEST_P(CvtColor, RGB2BGR) { doTest(3, 3, CVTCODE(RGB2BGR)); } OCL_TEST_P(CvtColor, BGRA2RGBA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); } -OCL_TEST_P(CvtColor, RGBA2BGRA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); } +OCL_TEST_P(CvtColor, RGBA2BGRA) { doTest(4, 4, CVTCODE(RGBA2BGRA)); } // RGB <-> Gray @@ -217,6 +217,16 @@ OCL_TEST_P(CvtColor8u32f, BGR2HLS_FULL) { doTest(3, 3, CVTCODE(BGR2HLS_FULL), de OCL_TEST_P(CvtColor8u32f, RGBA2HLS_FULL) { doTest(4, 3, CVTCODE(RGB2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } OCL_TEST_P(CvtColor8u32f, BGRA2HLS_FULL) { doTest(4, 3, CVTCODE(BGR2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, HLS2RGB) { doTest(3, 3, CVTCODE(HLS2RGB), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2BGR) { doTest(3, 3, CVTCODE(HLS2BGR), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2RGBA) { doTest(3, 4, CVTCODE(HLS2RGB), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2BGRA) { doTest(3, 4, CVTCODE(HLS2BGR), 1); } + +OCL_TEST_P(CvtColor8u32f, HLS2RGB_FULL) { doTest(3, 3, CVTCODE(HLS2RGB_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2BGR_FULL) { doTest(3, 3, CVTCODE(HLS2BGR_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2RGBA_FULL) { doTest(3, 4, CVTCODE(HLS2RGB_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2BGRA_FULL) { doTest(3, 4, CVTCODE(HLS2BGR_FULL), 1); } + // RGB5x5 <-> RGB typedef CvtColor CvtColor8u;