From af7c6144384f4b26a94575eac7ac97f70d32aa50 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 13 Nov 2013 17:09:05 +0400 Subject: [PATCH] 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 b01dcef1f..c3f58c92c 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 210d1b766..697526157 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 732a0f8e4..83a648530 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)),