From 98915e06bc7d1bb93f29e0c99f03264606fba6a0 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 13 Nov 2013 19:08:37 +0400 Subject: [PATCH] added HSV -> RGB[A][FULL] conversion --- modules/ocl/src/color.cpp | 30 +++++-- modules/ocl/src/opencl/cvt_color.cl | 118 ++++++++++++++++++++++++++++ modules/ocl/test/test_color.cpp | 26 ++++-- 3 files changed, 160 insertions(+), 14 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 17dcaa31f..82ed8044f 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -82,9 +82,12 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std:: } static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName, - const oclMat & data = oclMat()) + const std::string & additionalOptions = std::string(), const oclMat & data = oclMat()) { std::string build_options = format("-D DEPTH_%d -D dcn=%d", src.depth(), dst.channels()); + if (!additionalOptions.empty()) + build_options += additionalOptions; + int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -391,7 +394,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) } oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32FC1 : CV_32SC1, pdata); - toRGB_caller(src, dst, bidx, "XYZ2RGB", oclCoeffs); + toRGB_caller(src, dst, bidx, "XYZ2RGB", "", oclCoeffs); break; } case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL: @@ -440,17 +443,32 @@ 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 -D hscale=0", hrange), sdiv_data, hrange == 256 ? hdiv_data256 : hdiv_data180); + fromRGB_caller(src, dst, bidx, kernelName, format(" -D hrange=%d", hrange), sdiv_data, hrange == 256 ? hdiv_data256 : hdiv_data180); return; } - fromRGB_caller(src, dst, bidx, kernelName, format(" -D hscale=%f -D hrange=0", hrange*(1.f/360.f))); + fromRGB_caller(src, dst, bidx, kernelName, format(" -D hscale=%f", hrange*(1.f/360.f))); break; } - /* case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL: case CV_HLS2BGR: case CV_HLS2RGB: case CV_HLS2BGR_FULL: case CV_HLS2RGB_FULL: - */ + { + if (dcn <= 0) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F)); + bidx = code == CV_HSV2BGR || code == CV_HLS2BGR || + code == CV_HSV2BGR_FULL || code == CV_HLS2BGR_FULL ? 0 : 2; + int hrange = depth == CV_32F ? 360 : code == CV_HSV2BGR || code == CV_HSV2RGB || + code == CV_HLS2BGR || code == CV_HLS2RGB ? 180 : 255; + bool is_hsv = code == CV_HSV2BGR || code == CV_HSV2RGB || + code == CV_HSV2BGR_FULL || code == CV_HSV2RGB_FULL; + + dst.create(sz, CV_MAKETYPE(depth, dcn)); + + std::string kernelName = std::string(is_hsv ? "HSV" : "HLS") + "2RGB"; + toRGB_caller(src, dst, bidx, kernelName, format(" -D hrange=%d -D hscale=%f", hrange, 6.f/hrange)); + break; + } default: CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" ); } diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 928f46559..ae271be1d 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -46,6 +46,14 @@ /**************************************PUBLICFUNC*************************************/ +#ifndef hscale +#define hscale 0 +#endif + +#ifndef hrange +#define hrange 0 +#endif + #ifdef DEPTH_0 #define DATA_TYPE uchar #define COEFF_TYPE int @@ -548,6 +556,8 @@ __kernel void Gray2BGR5x5(int cols, int rows, int src_step, int dst_step, int bi ///////////////////////////////////// RGB <-> HSV ////////////////////////////////////// +__constant int sector_data[][3] = { {1, 3, 0}, { 1, 0, 2 }, { 3, 0, 1 }, { 0, 2, 1 }, { 0, 1, 3 }, { 2, 1, 0 } }; + #ifdef DEPTH_0 __kernel void RGB2HSV(int cols, int rows, int src_step, int dst_step, int bidx, @@ -590,6 +600,60 @@ __kernel void RGB2HSV(int cols, int rows, int src_step, int dst_step, int bidx, } } +__kernel void HSV2RGB(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], s = src[src_idx + 1]*(1/255.f), v = src[src_idx + 2]*(1/255.f); + float b, g, r; + + if (s != 0) + { + float tab[4]; + int sector; + 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; + if( (unsigned)sector >= 6u ) + { + sector = 0; + h = 0.f; + } + + tab[0] = v; + tab[1] = v*(1.f - s); + tab[2] = v*(1.f - s*h); + tab[3] = v*(1.f - s*(1.f - h)); + + b = tab[sector_data[sector][0]]; + g = tab[sector_data[sector][1]]; + r = tab[sector_data[sector][2]]; + } + else + b = g = r = v; + + 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 RGB2HSV(int cols, int rows, int src_step, int dst_step, int bidx, @@ -634,6 +698,60 @@ __kernel void RGB2HSV(int cols, int rows, int src_step, int dst_step, int bidx, } } +__kernel void HSV2RGB(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], s = src[src_idx + 1], v = src[src_idx + 2]; + float b, g, r; + + if (s != 0) + { + float tab[4]; + int sector; + 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; + if ((unsigned)sector >= 6u) + { + sector = 0; + h = 0.f; + } + + tab[0] = v; + tab[1] = v*(1.f - s); + tab[2] = v*(1.f - s*h); + tab[3] = v*(1.f - s*(1.f - h)); + + b = tab[sector_data[sector][0]]; + g = tab[sector_data[sector][1]]; + r = tab[sector_data[sector][2]]; + } + else + b = g = r = v; + + 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 ///////////////////////////////////// RGB <-> HLS ////////////////////////////////////// diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index b1128c79c..2b798d5c5 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -195,17 +195,27 @@ 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)); } +OCL_TEST_P(CvtColor8u32f, HSV2RGB) { doTest(3, 3, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2BGR) { doTest(3, 3, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2RGBA) { doTest(3, 4, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2BGRA) { doTest(3, 4, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); } + +OCL_TEST_P(CvtColor8u32f, HSV2RGB_FULL) { doTest(3, 3, CVTCODE(HSV2RGB_FULL), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2BGR_FULL) { doTest(3, 3, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2RGBA_FULL) { doTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2BGRA_FULL) { doTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } + // 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) { doTest(3, 3, CVTCODE(RGB2HLS), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, BGR2HLS) { doTest(3, 3, CVTCODE(BGR2HLS), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, RGBA2HLS) { doTest(4, 3, CVTCODE(RGB2HLS), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, BGRA2HLS) { doTest(4, 3, CVTCODE(BGR2HLS), depth == CV_8U ? 1 : 1e-3); } -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); } +OCL_TEST_P(CvtColor8u32f, RGB2HLS_FULL) { doTest(3, 3, CVTCODE(RGB2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, BGR2HLS_FULL) { doTest(3, 3, CVTCODE(BGR2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } +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); } // RGB5x5 <-> RGB