From af7c6144384f4b26a94575eac7ac97f70d32aa50 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 13 Nov 2013 17:09:05 +0400 Subject: [PATCH 1/4] added RGB[A] -> HSV[FULL] conversion --- modules/ocl/src/color.cpp | 70 ++++++++++++++-- modules/ocl/src/opencl/cvt_color.cl | 124 ++++++++++++++++++++++++++++ modules/ocl/test/test_color.cpp | 17 ++++ 3 files changed, 202 insertions(+), 9 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index b01dcef1f4..c3f58c92c4 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -51,12 +51,15 @@ using namespace cv; using namespace cv::ocl; static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName, - const oclMat & data = oclMat()) + const std::string & additionalOptions = std::string(), + const oclMat & data1 = oclMat(), const oclMat & data2 = oclMat()) { int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); std::string build_options = format("-D DEPTH_%d", src.depth()); + if (!additionalOptions.empty()) + build_options += additionalOptions; vector > args; args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); @@ -69,8 +72,10 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std:: args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - if (!data.empty()) - args.push_back( make_pair( sizeof(cl_mem) , (void *)&data.data )); + if (!data1.empty()) + args.push_back( make_pair( sizeof(cl_mem) , (void *)&data1.data )); + if (!data2.empty()) + args.push_back( make_pair( sizeof(cl_mem) , (void *)&data2.data )); size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); @@ -297,10 +302,6 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) toRGB_caller(src, dst, bidx, "YCrCb2RGB"); break; } - /* - case CV_BGR5652GRAY: case CV_BGR5552GRAY: - case CV_GRAY2BGR565: case CV_GRAY2BGR555: - */ case CV_BGR2XYZ: case CV_RGB2XYZ: { @@ -343,7 +344,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); - fromRGB_caller(src, dst, bidx, "RGB2XYZ", oclCoeffs); + fromRGB_caller(src, dst, bidx, "RGB2XYZ", "", oclCoeffs); break; } case CV_XYZ2BGR: @@ -393,9 +394,60 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) toRGB_caller(src, dst, bidx, "XYZ2RGB", oclCoeffs); break; } - /* case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL: case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL: + { + CV_Assert((scn == 3 || scn == 4) && (depth == CV_8U || depth == CV_32F)); + bidx = code == CV_BGR2HSV || code == CV_BGR2HLS || + code == CV_BGR2HSV_FULL || code == CV_BGR2HLS_FULL ? 0 : 2; + int hrange = depth == CV_32F ? 360 : code == CV_BGR2HSV || code == CV_RGB2HSV || + code == CV_BGR2HLS || code == CV_RGB2HLS ? 180 : 256; + bool is_hsv = code == CV_BGR2HSV || code == CV_RGB2HSV || code == CV_BGR2HSV_FULL || code == CV_RGB2HSV_FULL; + dst.create(sz, CV_MAKETYPE(depth, 3)); + std::string kernelName = std::string("RGB2") + (is_hsv ? "HSV" : "HLS"); + + if (is_hsv && depth == CV_8U) + { + static oclMat sdiv_data; + static oclMat hdiv_data180; + static oclMat hdiv_data256; + static int sdiv_table[256]; + static int hdiv_table180[256]; + static int hdiv_table256[256]; + static volatile bool initialized180 = false, initialized256 = false; + volatile bool & initialized = hrange == 180 ? initialized180 : initialized256; + + if (!initialized) + { + int * const hdiv_table = hrange == 180 ? hdiv_table180 : hdiv_table256, hsv_shift = 12; + oclMat & hdiv_data = hrange == 180 ? hdiv_data180 : hdiv_data256; + + sdiv_table[0] = hdiv_table180[0] = hdiv_table256[0] = 0; + + int v = 255 << hsv_shift; + if (!initialized180 && !initialized256) + { + for(int i = 1; i < 256; i++ ) + sdiv_table[i] = saturate_cast(v/(1.*i)); + sdiv_data.upload(Mat(1, 256, CV_32SC1, sdiv_table)); + } + + v = hrange << hsv_shift; + for (int i = 1; i < 256; i++ ) + hdiv_table[i] = saturate_cast(v/(6.*i)); + + hdiv_data.upload(Mat(1, 256, CV_32SC1, hdiv_table)); + initialized = true; + } + + 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", 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: */ diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 210d1b766e..6975261571 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -76,6 +76,7 @@ enum { yuv_shift = 14, xyz_shift = 12, + hsv_shift = 12, R2Y = 4899, G2Y = 9617, B2Y = 1868, @@ -544,3 +545,126 @@ __kernel void Gray2BGR5x5(int cols, int rows, int src_step, int dst_step, int bi #endif } } + +///////////////////////////////////// RGB <-> HSV ////////////////////////////////////// + +#ifdef DEPTH_0 + +__kernel void RGB2HSV(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, + __constant int * sdiv_table, __constant int * hdiv_table) +{ + 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); + + int b = src[src_idx + bidx], g = src[src_idx + 1], r = src[src_idx + (bidx^2)]; + int h, s, v = b; + int vmin = b, diff; + int vr, vg; + + v = max( v, g ); + v = max( v, r ); + vmin = min( vmin, g ); + vmin = min( vmin, r ); + + diff = v - vmin; + vr = v == r ? -1 : 0; + vg = v == g ? -1 : 0; + + s = (diff * sdiv_table[v] + (1 << (hsv_shift-1))) >> hsv_shift; + h = (vr & (g - b)) + + (~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff)))); + h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift; + h += h < 0 ? hrange : 0; + + dst[dst_idx] = convert_uchar_sat_rte(h); + dst[dst_idx + 1] = (uchar)s; + dst[dst_idx + 2] = (uchar)v; + } +} + +#elif defined DEPTH_5 + +__kernel void RGB2HSV(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, s, v; + + float vmin, diff; + + v = vmin = r; + if( v < g ) v = g; + if( v < b ) v = b; + if( vmin > g ) vmin = g; + if( vmin > b ) vmin = b; + + diff = v - vmin; + s = diff/(float)(fabs(v) + FLT_EPSILON); + diff = (float)(60./(diff + FLT_EPSILON)); + if( v == r ) + h = (g - b)*diff; + else if( v == g ) + h = (b - r)*diff + 120.f; + else + h = (r - g)*diff + 240.f; + + if( h < 0 ) h += 360.f; + + dst[dst_idx] = h*hscale; + dst[dst_idx + 1] = s; + dst[dst_idx + 2] = v; + } +} + +#endif + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 732a0f8e45..83a6485309 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -181,6 +181,20 @@ OCL_TEST_P(CvtColor, XYZ2BGR) { doTest(3, 3, CVTCODE(XYZ2BGR)); } OCL_TEST_P(CvtColor, XYZ2RGBA) { doTest(3, 4, CVTCODE(XYZ2RGB)); } OCL_TEST_P(CvtColor, XYZ2BGRA) { doTest(3, 4, CVTCODE(XYZ2BGR)); } +// RGB <-> HSV + +typedef CvtColor CvtColor8u32f; + +OCL_TEST_P(CvtColor8u32f, RGB2HSV) { doTest(3, 3, CVTCODE(RGB2HSV)); } +OCL_TEST_P(CvtColor8u32f, BGR2HSV) { doTest(3, 3, CVTCODE(BGR2HSV)); } +OCL_TEST_P(CvtColor8u32f, RGBA2HSV) { doTest(4, 3, CVTCODE(RGB2HSV)); } +OCL_TEST_P(CvtColor8u32f, BGRA2HSV) { doTest(4, 3, CVTCODE(BGR2HSV)); } + +OCL_TEST_P(CvtColor8u32f, RGB2HSV_FULL) { doTest(3, 3, CVTCODE(RGB2HSV_FULL)); } +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)); } + // RGB5x5 <-> RGB typedef CvtColor CvtColor8u; @@ -246,6 +260,9 @@ OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { doTest(1, 3, CV_YUV2BGR_NV12); } INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor8u, testing::Combine(testing::Values(MatDepth(CV_8U)), Bool())); +INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor8u32f, + testing::Combine(testing::Values(MatDepth(CV_8U), MatDepth(CV_32F)), Bool())); + INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine( testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)), From 3bdd9626f3ad0e8ee9719e19e125292b22ad3594 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 13 Nov 2013 17:41:40 +0400 Subject: [PATCH 2/4] added RGB[A] -> HLS[FULL] conversion --- modules/ocl/src/color.cpp | 4 +- modules/ocl/src/opencl/cvt_color.cl | 103 +++++++++++++++++++++++----- modules/ocl/test/test_color.cpp | 18 ++++- 3 files changed, 103 insertions(+), 22 deletions(-) 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; From 98915e06bc7d1bb93f29e0c99f03264606fba6a0 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 13 Nov 2013 19:08:37 +0400 Subject: [PATCH 3/4] 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 17dcaa31fc..82ed8044ff 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 928f46559e..ae271be1da 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 b1128c79c1..2b798d5c57 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 From 1b7c5b201d2d9718068da9676bd2e066dbc0584d Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 13 Nov 2013 19:35:24 +0400 Subject: [PATCH 4/4] added HLS -> RGB[A][FULL] conversion --- modules/ocl/src/color.cpp | 24 +++---- modules/ocl/src/opencl/cvt_color.cl | 107 ++++++++++++++++++++++++++++ modules/ocl/test/test_color.cpp | 22 ++++-- 3 files changed, 131 insertions(+), 22 deletions(-) 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;