diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index c3f58c92c4..17dcaa31fc 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -440,11 +440,11 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) initialized = true; } - fromRGB_caller(src, dst, bidx, kernelName, format(" -D hrange=%d", hrange), sdiv_data, hrange == 256 ? hdiv_data256 : hdiv_data180); + fromRGB_caller(src, dst, bidx, kernelName, format(" -D hrange=%d -D hscale=0", hrange), sdiv_data, hrange == 256 ? hdiv_data256 : hdiv_data180); return; } - fromRGB_caller(src, dst, bidx, kernelName, format(" -D hscale=%f", hrange*(1.f/360.f))); + fromRGB_caller(src, dst, bidx, kernelName, format(" -D hscale=%f -D hrange=0", hrange*(1.f/360.f))); break; } /* diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 6975261571..928f46559e 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -636,35 +636,104 @@ __kernel void RGB2HSV(int cols, int rows, int src_step, int dst_step, int bidx, #endif +///////////////////////////////////// RGB <-> HLS ////////////////////////////////////// +#ifdef DEPTH_0 +__kernel void RGB2HLS(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 b = src[src_idx + bidx]*(1/255.f), g = src[src_idx + 1]*(1/255.f), r = src[src_idx + (bidx^2)]*(1/255.f); + float h = 0.f, s = 0.f, l; + float vmin, vmax, diff; + + vmax = vmin = r; + if (vmax < g) vmax = g; + if (vmax < b) vmax = b; + if (vmin > g) vmin = g; + if (vmin > b) vmin = b; + + diff = vmax - vmin; + l = (vmax + vmin)*0.5f; + + if (diff > FLT_EPSILON) + { + s = l < 0.5f ? diff/(vmax + vmin) : diff/(2 - vmax - vmin); + diff = 60.f/diff; + + if( vmax == r ) + h = (g - b)*diff; + else if( vmax == g ) + h = (b - r)*diff + 120.f; + else + h = (r - g)*diff + 240.f; + + if( h < 0.f ) h += 360.f; + } + + dst[dst_idx] = convert_uchar_sat_rte(h*hscale); + dst[dst_idx + 1] = convert_uchar_sat_rte(l*255.f); + dst[dst_idx + 2] = convert_uchar_sat_rte(s*255.f); + } +} +#elif defined DEPTH_5 +__kernel void RGB2HLS(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 b = src[src_idx + bidx], g = src[src_idx + 1], r = src[src_idx + (bidx^2)]; + float h = 0.f, s = 0.f, l; + float vmin, vmax, diff; + vmax = vmin = r; + if (vmax < g) vmax = g; + if (vmax < b) vmax = b; + if (vmin > g) vmin = g; + if (vmin > b) vmin = b; + diff = vmax - vmin; + l = (vmax + vmin)*0.5f; + if (diff > FLT_EPSILON) + { + s = l < 0.5f ? diff/(vmax + vmin) : diff/(2 - vmax - vmin); + diff = 60.f/diff; + if( vmax == r ) + h = (g - b)*diff; + else if( vmax == g ) + h = (b - r)*diff + 120.f; + else + h = (r - g)*diff + 240.f; + if( h < 0.f ) h += 360.f; + } + dst[dst_idx] = h*hscale; + dst[dst_idx + 1] = l; + dst[dst_idx + 2] = s; + } +} - - - - - - - - - - - - - - - - - +#endif diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 83a6485309..b1128c79c1 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -95,7 +95,7 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) generateOclMat(gdst_whole, gdst_roi, dst, roiSize, dstBorder); } - void Near(double threshold = 1e-3) + void Near(double threshold) { Mat whole, roi; gdst_whole.download(whole); @@ -105,7 +105,7 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) EXPECT_MAT_NEAR(dst_roi, roi, threshold); } - void doTest(int channelsIn, int channelsOut, int code) + void doTest(int channelsIn, int channelsOut, int code, double threshold = 1e-3) { for (int j = 0; j < LOOP_TIMES; j++) { @@ -114,7 +114,7 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) cvtColor(src_roi, dst_roi, code, channelsOut); ocl::cvtColor(gsrc_roi, gdst_roi, code, channelsOut); - Near(); + Near(threshold); } } }; @@ -195,6 +195,18 @@ OCL_TEST_P(CvtColor8u32f, BGR2HSV_FULL) { doTest(3, 3, CVTCODE(BGR2HSV_FULL)); } OCL_TEST_P(CvtColor8u32f, RGBA2HSV_FULL) { doTest(4, 3, CVTCODE(RGB2HSV_FULL)); } OCL_TEST_P(CvtColor8u32f, BGRA2HSV_FULL) { doTest(4, 3, CVTCODE(BGR2HSV_FULL)); } +// RGB <-> HLS + +OCL_TEST_P(CvtColor8u32f, RGB2HLS) { doTest(3, 3, CVTCODE(RGB2HLS), 1); } +OCL_TEST_P(CvtColor8u32f, BGR2HLS) { doTest(3, 3, CVTCODE(BGR2HLS), 1); } +OCL_TEST_P(CvtColor8u32f, RGBA2HLS) { doTest(4, 3, CVTCODE(RGB2HLS), 1); } +OCL_TEST_P(CvtColor8u32f, BGRA2HLS) { doTest(4, 3, CVTCODE(BGR2HLS), 1); } + +OCL_TEST_P(CvtColor8u32f, RGB2HLS_FULL) { doTest(3, 3, CVTCODE(RGB2HLS_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, BGR2HLS_FULL) { doTest(3, 3, CVTCODE(BGR2HLS_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, RGBA2HLS_FULL) { doTest(4, 3, CVTCODE(RGB2HLS_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, BGRA2HLS_FULL) { doTest(4, 3, CVTCODE(BGR2HLS_FULL), 1); } + // RGB5x5 <-> RGB typedef CvtColor CvtColor8u;