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

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" );
     }

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

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