added XYZ to RGB conversion to ocl::cvtColor

diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp
index 02e7bf3..9688be0 100644 (file)
--- a/modules/ocl/src/color.cpp
+++ b/modules/ocl/src/color.cpp
@@ -76,10 +76,10 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::
     openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str());
 }
 
-static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName)
+static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName,
+                         const oclMat & data = oclMat())
 {
-    int channels = dst.channels();
-    std::string build_options = format("-D DEPTH_%d", src.depth());
+    std::string build_options = format("-D DEPTH_%d -D channels=%d", src.depth(), dst.channels());
     int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
     int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
 
@@ -88,13 +88,15 @@ static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::st
     args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows));
     args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step));
     args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step));
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&channels));
     args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx));
     args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
     args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
     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 ));
+
     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());
 }
@@ -235,7 +237,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
             }
             pdata = coeffs;
         }
-        oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32F : CV_32S, pdata);
+        oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32FC1 : CV_32SC1, pdata);
 
         fromRGB_caller(src, dst, bidx, "RGB2XYZ", oclCoeffs);
         break;
@@ -248,7 +250,43 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         CV_Assert(scn == 3 && (dcn == 3 || dcn == 4));
         bidx = code == CV_XYZ2BGR ? 0 : 2;
         dst.create(sz, CV_MAKE_TYPE(depth, dcn));
-        toRGB_caller(src, dst, bidx, "XYZ2RGB");
+
+        void * pdata = NULL;
+        if (depth == CV_32F)
+        {
+            float coeffs[] =
+            {
+                3.240479f, -1.53715f, -0.498535f,
+                -0.969256f, 1.875991f, 0.041556f,
+                0.055648f, -0.204043f, 1.057311f
+            };
+            if (bidx == 0)
+            {
+                std::swap(coeffs[0], coeffs[6]);
+                std::swap(coeffs[1], coeffs[7]);
+                std::swap(coeffs[2], coeffs[8]);
+            }
+            pdata = coeffs;
+        }
+        else
+        {
+            int coeffs[] =
+            {
+                13273,  -6296,  -2042,
+                -3970,   7684,    170,
+                  228,   -836,   4331
+            };
+            if (bidx == 0)
+            {
+                std::swap(coeffs[0], coeffs[6]);
+                std::swap(coeffs[1], coeffs[7]);
+                std::swap(coeffs[2], coeffs[8]);
+            }
+            pdata = coeffs;
+        }
+        oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32FC1 : CV_32SC1, pdata);
+
+        toRGB_caller(src, dst, bidx, "XYZ2RGB", oclCoeffs);
         break;
     }
     /*

diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl
index 775628b..0f44897 100644 (file)
--- a/modules/ocl/src/opencl/cvt_color.cl
+++ b/modules/ocl/src/opencl/cvt_color.cl
@@ -46,7 +46,7 @@
 
 /**************************************PUBLICFUNC*************************************/
 
-#if defined (DEPTH_0)
+#ifdef DEPTH_0
 #define DATA_TYPE uchar
 #define COEFF_TYPE int
 #define MAX_NUM  255
@@ -54,7 +54,7 @@
 #define SAT_CAST(num) convert_uchar_sat_rte(num)
 #endif
 
-#if defined (DEPTH_2)
+#ifdef DEPTH_2
 #define DATA_TYPE ushort
 #define COEFF_TYPE int
 #define MAX_NUM  65535
@@ -62,7 +62,7 @@
 #define SAT_CAST(num) convert_ushort_sat_rte(num)
 #endif
 
-#if defined (DEPTH_5)
+#ifdef DEPTH_5
 #define DATA_TYPE float
 #define COEFF_TYPE float
 #define MAX_NUM  1.0f
@@ -95,7 +95,7 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step,
     {
         int src_idx = mad24(y, src_step, src_offset + (x << 2));
         int dst_idx = mad24(y, dst_step, dst_offset + x);
-#if defined (DEPTH_5)
+#ifdef DEPTH_5
         dst[dst_idx] = src[src_idx + bidx] * 0.114f + src[src_idx + 1] * 0.587f + src[src_idx + (bidx^2)] * 0.299f;
 #else
         dst[dst_idx] = (DATA_TYPE)CV_DESCALE((src[src_idx + bidx] * B2Y + src[src_idx + 1] * G2Y + src[src_idx + (bidx^2)] * R2Y), yuv_shift);
@@ -103,7 +103,7 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step,
     }
 }
 
-__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int channels, int bidx,
+__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int bidx,
                        __global const DATA_TYPE* src, __global DATA_TYPE* dst,
                        int src_offset, int dst_offset)
 {
@@ -119,8 +119,9 @@ __kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int chann
         dst[dst_idx] = val;
         dst[dst_idx + 1] = val;
         dst[dst_idx + 2] = val;
-        if (channels == 4)
+#if channels == 4
             dst[dst_idx + 3] = MAX_NUM;
+#endif
     }
 }
 
@@ -143,7 +144,7 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step,
         int dst_idx = mad24(y, dst_step, dst_offset + x);
         DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
 
-#if defined (DEPTH_5)
+#ifdef DEPTH_5
         __constant float * coeffs = c_RGB2YUVCoeffs_f;
         DATA_TYPE Y  = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx];
         DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX;
@@ -165,7 +166,7 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step,
 __constant float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f };
 __constant int   c_YUV2RGBCoeffs_i[5] = { 33292, -6472, -9519, 18678 };
 
-__kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, int channels,
+__kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step,
                       int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
                       int src_offset, int dst_offset)
 {
@@ -179,7 +180,7 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, int channe
         int dst_idx = mad24(y, dst_step, dst_offset + x);
         DATA_TYPE yuv[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
 
-#if defined (DEPTH_5)
+#ifdef DEPTH_5
         __constant float * coeffs = c_YUV2RGBCoeffs_f;
         float b = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[3];
         float g = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1];
@@ -194,8 +195,9 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, int channe
         dst[dst_idx + bidx] = SAT_CAST( b );
         dst[dst_idx + 1]      = SAT_CAST( g );
         dst[dst_idx + (bidx^2)]   = SAT_CAST( r );
-        if (channels == 4)
+#if channels == 4
             dst[dst_idx + 3] = MAX_NUM;
+#endif
     }
 }
 
@@ -206,7 +208,7 @@ __constant int ITUR_BT_601_CVG = 852492;
 __constant int ITUR_BT_601_CVR = 1673527;
 __constant int ITUR_BT_601_SHIFT = 20;
 
-__kernel void YUV2RGBA_NV12(int cols, int rows, int src_step, int dst_step, int channels,
+__kernel void YUV2RGBA_NV12(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)
 {
@@ -278,7 +280,7 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step,
 
         DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
 
-#if defined (DEPTH_5)
+#ifdef DEPTH_5
         __constant float * coeffs = c_RGB2YCrCbCoeffs_f;
         DATA_TYPE Y  = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx];
         DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX;
@@ -300,7 +302,7 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step,
 __constant float c_YCrCb2RGBCoeffs_f[4] = { 1.403f, -0.714f, -0.344f, 1.773f };
 __constant int   c_YCrCb2RGBCoeffs_i[4] = { 22987, -11698, -5636, 29049 };
 
-__kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step, int channels,
+__kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step,
                         int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
                         int src_offset, int dst_offset)
 {
@@ -330,41 +332,73 @@ __kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step, int chan
         dst[dst_idx + (bidx^2)] = SAT_CAST(r);
         dst[dst_idx + 1] = SAT_CAST(g);
         dst[dst_idx + bidx] = SAT_CAST(b);
-        if (channels == 4)
+#if channels == 4
             dst[dst_idx + 3] = MAX_NUM;
+#endif
     }
 }
 
 ///////////////////////////////////// RGB <-> XYZ //////////////////////////////////////
 
-#pragma OPENCL EXTENSION cl_amd_printf:enable
-
 __kernel void RGB2XYZ(int cols, int rows, int src_step, int dst_step,
-                        int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
-                        int src_offset, int dst_offset, __global COEFF_TYPE * coeffs)
+                      int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
+                      int src_offset, int dst_offset, __constant COEFF_TYPE * coeffs)
 {
-    int x = get_global_id(0);
-    int y = get_global_id(1);
+    int dx = get_global_id(0);
+    int dy = get_global_id(1);
 
-    if (y < rows && x < cols)
+    if (dy < rows && dx < cols)
     {
-        x <<= 2;
-        int src_idx = mad24(y, src_step, src_offset + x);
-        int dst_idx = mad24(y, dst_step, dst_offset + x);
+        dx <<= 2;
+        int src_idx = mad24(dy, src_step, src_offset + dx);
+        int dst_idx = mad24(dy, dst_step, dst_offset + dx);
 
         DATA_TYPE r = src[src_idx], g = src[src_idx + 1], b = src[src_idx + 2];
 
 #ifdef DEPTH_5
-        DATA_TYPE x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2];
-        DATA_TYPE y = r * coeffs[3] + g * coeffs[4] + b * coeffs[5];
-        DATA_TYPE z = r * coeffs[6] + g * coeffs[7] + b * coeffs[8];
+        float x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2];
+        float y = r * coeffs[3] + g * coeffs[4] + b * coeffs[5];
+        float z = r * coeffs[6] + g * coeffs[7] + b * coeffs[8];
 #else
-        DATA_TYPE x = CV_DESCALE(r * coeffs[0] + g * coeffs[1] + b * coeffs[2], xyz_shift);
-        DATA_TYPE y = CV_DESCALE(r * coeffs[3] + g * coeffs[4] + b * coeffs[5], xyz_shift);
-        DATA_TYPE z = CV_DESCALE(r * coeffs[6] + g * coeffs[7] + b * coeffs[8], xyz_shift);
+        int x = CV_DESCALE(r * coeffs[0] + g * coeffs[1] + b * coeffs[2], xyz_shift);
+        int y = CV_DESCALE(r * coeffs[3] + g * coeffs[4] + b * coeffs[5], xyz_shift);
+        int z = CV_DESCALE(r * coeffs[6] + g * coeffs[7] + b * coeffs[8], xyz_shift);
 #endif
         dst[dst_idx] = SAT_CAST(x);
         dst[dst_idx + 1] = SAT_CAST(y);
         dst[dst_idx + 2] = SAT_CAST(z);
     }
 }
+
+__kernel void XYZ2RGB(int cols, int rows, int src_step, int dst_step,
+                      int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
+                      int src_offset, int dst_offset, __constant COEFF_TYPE * coeffs)
+{
+    int dx = get_global_id(0);
+    int dy = get_global_id(1);
+
+    if (dy < rows && dx < cols)
+    {
+        dx <<= 2;
+        int src_idx = mad24(dy, src_step, src_offset + dx);
+        int dst_idx = mad24(dy, dst_step, dst_offset + dx);
+
+        DATA_TYPE x = src[src_idx], y = src[src_idx + 1], z = src[src_idx + 2];
+
+#ifdef DEPTH_5
+        float b = x * coeffs[0] + y * coeffs[1] + z * coeffs[2];
+        float g = x * coeffs[3] + y * coeffs[4] + z * coeffs[5];
+        float r = x * coeffs[6] + y * coeffs[7] + z * coeffs[8];
+#else
+        int b = CV_DESCALE(x * coeffs[0] + y * coeffs[1] + z * coeffs[2], xyz_shift);
+        int g = CV_DESCALE(x * coeffs[3] + y * coeffs[4] + z * coeffs[5], xyz_shift);
+        int r = CV_DESCALE(x * coeffs[6] + y * coeffs[7] + z * coeffs[8], xyz_shift);
+#endif
+        dst[dst_idx] = SAT_CAST(b);
+        dst[dst_idx + 1] = SAT_CAST(g);
+        dst[dst_idx + 2] = SAT_CAST(r);
+#if channels == 4
+            dst[dst_idx + 3] = MAX_NUM;
+#endif
+    }
+}

diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp
index 0d9eb8f..977fd20 100644 (file)
--- a/modules/ocl/test/test_color.cpp
+++ b/modules/ocl/test/test_color.cpp
@@ -248,6 +248,23 @@ OCL_TEST_P(CvtColor, BGRA2XYZ)
     doTest(4, 3, CVTCODE(BGR2XYZ));
 }
 
+OCL_TEST_P(CvtColor, XYZ2RGB)
+{
+    doTest(3, 3, CVTCODE(XYZ2RGB));
+}
+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));
+}
+
 // YUV -> RGBA_NV12
 
 struct CvtColor_YUV420 :