color.cpp refactoring: created generic interface for toRGB and fromRGB callers

modules/ocl/src/color.cpp

@@ -50,7 +50,7 @@
 using namespace cv;
 using namespace cv::ocl;
 
-static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx)
+static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName)
 {
     int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
     int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
@@ -58,8 +58,8 @@ static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx)
     std::string build_options = format("-D DEPTH_%d", src.depth());
 
     vector<pair<size_t , const void *> > args;
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows));
+    args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols));
+    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 *)&bidx));
@@ -68,11 +68,11 @@ static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx)
     args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset ));
     args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset ));
 
-    size_t gt[3] = { src.cols, src.rows, 1 }, lt[3] = { 16, 16, 1 };
-    openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2Gray", gt, lt, args, -1, -1, build_options.c_str());
+    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());
 }
 
-static void Gray2RGB_caller(const oclMat &src, oclMat &dst)
+static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName)
 {
     int channels = dst.channels();
     std::string build_options = format("-D DEPTH_%d", src.depth());
@@ -80,99 +80,8 @@ static void Gray2RGB_caller(const oclMat &src, oclMat &dst)
     int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
 
     vector<pair<size_t , const void *> > args;
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.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_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 ));
-
-    size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
-    openCLExecuteKernel(src.clCxt, &cvt_color, "Gray2RGB", gt, lt, args, -1, -1, build_options.c_str());
-}
-
-static void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx)
-{
-    std::string build_options = format("-D DEPTH_%d", src.depth());
-    int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
-    int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
-
-    vector<pair<size_t , const void *> > args;
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.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 *)&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 ));
-
-    size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
-    openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YUV", gt, lt, args, -1, -1, build_options.c_str());
-}
-
-static void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx)
-{
-    int channels = dst.channels();
-    int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
-    int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
-
-    std::string buildOptions = format("-D DEPTH_%d", src.depth());
-
-    vector<pair<size_t , const void *> > args;
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.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 ));
-
-    size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
-    openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGB", gt, lt, args, -1, -1, buildOptions.c_str());
-}
-
-static void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx)
-{
-    std::string build_options = format("-D DEPTH_%d", src.depth());
-    int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
-    int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
-
-    vector<pair<size_t , const void *> > args;
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.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 *)&bidx));
     args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols));
     args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows));
-    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 ));
-
-    size_t gt[3] = {dst.cols / 2, dst.rows / 2, 1}, lt[3] = {16, 16, 1};
-    openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGBA_NV12", gt, lt, args, -1, -1, build_options.c_str());
-}
-
-static void YCrCb2RGB_caller(const oclMat &src, oclMat &dst, int bidx)
-{
-    int channels = dst.channels();
-    int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
-    int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
-
-    std::string buildOptions = format("-D DEPTH_%d", src.depth());
-
-    vector<pair<size_t , const void *> > args;
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.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));
@@ -182,29 +91,8 @@ static void YCrCb2RGB_caller(const oclMat &src, oclMat &dst, int bidx)
     args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset ));
     args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset ));
 
-    size_t gt[3] = { src.cols, src.rows, 1 }, lt[3] = { 16, 16, 1 };
-    openCLExecuteKernel(src.clCxt, &cvt_color, "YCrCb2RGB", gt, lt, args, -1, -1, buildOptions.c_str());
-}
-
-static void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx)
-{
-    std::string build_options = format("-D DEPTH_%d", src.depth());
-    int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
-    int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
-
-    vector<pair<size_t , const void *> > args;
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols));
-    args.push_back( make_pair( sizeof(cl_int) , (void *)&src.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 *)&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 ));
-
-    size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1};
-    openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YCrCb", gt, lt, args, -1, -1, build_options.c_str());
+    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());
 }
 
 static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
@@ -232,7 +120,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         CV_Assert(scn == 3 || scn == 4);
         bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
         dst.create(sz, CV_MAKETYPE(depth, 1));
-        RGB2Gray_caller(src, dst, bidx);
+        fromRGB_caller(src, dst, bidx, "RGB2Gray");
         break;
     }
     case CV_GRAY2BGR:
@@ -241,7 +129,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         CV_Assert(scn == 1);
         dcn  = code == CV_GRAY2BGRA ? 4 : 3;
         dst.create(sz, CV_MAKETYPE(depth, dcn));
-        Gray2RGB_caller(src, dst);
+        toRGB_caller(src, dst, 0, "Gray2RGB");
         break;
     }
     case CV_BGR2YUV:
@@ -250,7 +138,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         CV_Assert(scn == 3 || scn == 4);
         bidx = code == CV_BGR2YUV ? 0 : 2;
         dst.create(sz, CV_MAKETYPE(depth, 3));
-        RGB2YUV_caller(src, dst, bidx);
+        fromRGB_caller(src, dst, bidx, "RGB2YUV");
         break;
     }
     case CV_YUV2BGR:
@@ -261,7 +149,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         CV_Assert(scn == 3 && (dcn == 3 || dcn == 4));
         bidx = code == CV_YUV2BGR ? 0 : 2;
         dst.create(sz, CV_MAKETYPE(depth, dcn));
-        YUV2RGB_caller(src, dst, bidx);
+        toRGB_caller(src, dst, bidx, "YUV2RGB");
         break;
     }
     case CV_YUV2RGB_NV12:
@@ -276,7 +164,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
 
         Size dstSz(sz.width, sz.height * 2 / 3);
         dst.create(dstSz, CV_MAKETYPE(depth, dcn));
-        YUV2RGB_NV12_caller(src, dst, bidx);
+        toRGB_caller(src, dst, bidx, "YUV2RGBA_NV12");
         break;
     }
     case CV_BGR2YCrCb:
@@ -285,7 +173,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         CV_Assert(scn == 3 || scn == 4);
         bidx = code == CV_BGR2YCrCb ? 0 : 2;
         dst.create(sz, CV_MAKETYPE(depth, 3));
-        RGB2YCrCb_caller(src, dst, bidx);
+        fromRGB_caller(src, dst, bidx, "RGB2YCrCb");
         break;
     }
     case CV_YCrCb2BGR:
@@ -296,7 +184,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         CV_Assert(scn == 3 && (dcn == 3 || dcn == 4));
         bidx = code == CV_YCrCb2BGR ? 0 : 2;
         dst.create(sz, CV_MAKETYPE(depth, dcn));
-        YCrCb2RGB_caller(src, dst, bidx);
+        toRGB_caller(src, dst, bidx, "YCrCb2RGB");
         break;
     }
     /*

modules/ocl/src/opencl/cvt_color.cl

@@ -100,7 +100,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,
+__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int channels, int bidx,
                        __global const DATA_TYPE* src, __global DATA_TYPE* dst,
                        int src_offset, int dst_offset)
 {
@@ -203,17 +203,17 @@ __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 bidx, int width, int height, __global const uchar* src, __global uchar* dst,
+__kernel void YUV2RGBA_NV12(int cols, int rows, int src_step, int dst_step, int channels,
+                            int bidx, __global const uchar* src, __global uchar* dst,
                             int src_offset, int dst_offset)
 {
-    const int x = get_global_id(0); // max_x = width / 2
-    const int y = get_global_id(1); // max_y = height/ 2
+    const int x = get_global_id(0);
+    const int y = get_global_id(1);
 
-    if (y < height / 2 && x < width / 2 )
+    if (y < rows / 2 && x < cols / 2 )
     {
         __global const uchar* ysrc = src + mad24(y << 1, src_step, (x << 1) + src_offset);
-        __global const uchar* usrc = src + mad24(height + y, src_step, (x << 1) + src_offset);
+        __global const uchar* usrc = src + mad24(rows + y, src_step, (x << 1) + src_offset);
         __global uchar*       dst1 = dst + mad24(y << 1, dst_step, (x << 3) + dst_offset);
         __global uchar*       dst2 = dst + mad24((y << 1) + 1, dst_step, (x << 3) + dst_offset);
 

modules/ocl/test/test_color.cpp

@@ -116,6 +116,8 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool)
 
 #define CVTCODE(name) cv::COLOR_ ## name
 
+// RGB <-> Gray
+
 OCL_TEST_P(CvtColor, RGB2GRAY)
 {
     doTest(3, 1, CVTCODE(RGB2GRAY));
@@ -152,6 +154,7 @@ OCL_TEST_P(CvtColor, GRAY2BGRA)
     doTest(1, 4, CVTCODE(GRAY2BGRA));
 }
 
+// RGB <-> YUV
 
 OCL_TEST_P(CvtColor, RGB2YUV)
 {
@@ -186,6 +189,7 @@ OCL_TEST_P(CvtColor, YUV2BGRA)
     doTest(3, 4, CVTCODE(YUV2BGR));
 }
 
+// RGB <-> YCrCb
 
 OCL_TEST_P(CvtColor, RGB2YCrCb)
 {
@@ -220,6 +224,8 @@ OCL_TEST_P(CvtColor, YCrCb2BGRA)
     doTest(3, 4, CVTCODE(YCrCb2BGR));
 }
 
+// YUV -> RGBA_NV12
+
 struct CvtColor_YUV420 :
         public CvtColor
 {
@@ -262,7 +268,6 @@ OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12)
     doTest(1, 3, CV_YUV2BGR_NV12);
 }
 
-
 INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor,
                             testing::Combine(
                                 testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),