added RGB[A] -> HSV[FULL] conversion

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<pair<size_t , const void *> > 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())
+    if (!data1.empty())
-        args.push_back( make_pair( sizeof(cl_mem) , (void *)&data.data ));
+        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<int>(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<int>(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:
     */

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
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+

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)),