From a57030a0cda4a144c1b800ce5e7efb6c127aa571 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sat, 9 Nov 2013 17:03:30 +0400 Subject: [PATCH] added YCrCb to RGB, BGR, RGBA, BGRA modes to ocl::cvtColor --- modules/ocl/src/color.cpp | 28 +++++++++++++++++++++++++-- modules/ocl/src/opencl/cvt_color.cl | 38 +++++++++++++++++++++++++++++++++++++ modules/ocl/test/test_color.cpp | 32 +++++++++++++++---------------- 3 files changed, 80 insertions(+), 18 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 766e992..b25b71f 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -162,6 +162,30 @@ static void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx) 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 > 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, "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()); @@ -270,9 +294,9 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) if( dcn <= 0 ) dcn = 3; CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); - bidx = code == CV_YCrCb2RGB ? 0 : 2; + bidx = code == CV_YCrCb2BGR ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, dcn)); -// YUV2RGB_caller(src, dst, bidx); + YCrCb2RGB_caller(src, dst, bidx); break; } /* diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 2ba7739..de53da5 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -293,3 +293,41 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step, dst[dst_idx + 2] = SAT_CAST( Cb ); } } + +__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, + int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* 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); + + DATA_TYPE ycrcb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; + +#ifdef DEPTH_5 + __constant float * coeff = c_YCrCb2RGBCoeffs_f; + float r = ycrcb[0] + coeff[0] * (ycrcb[1] - HALF_MAX); + float g = ycrcb[0] + coeff[1] * (ycrcb[1] - HALF_MAX) + coeff[2] * (ycrcb[2] - HALF_MAX); + float b = ycrcb[0] + coeff[3] * (ycrcb[2] - HALF_MAX); +#else + __constant int * coeff = c_YCrCb2RGBCoeffs_i; + int r = ycrcb[0] + CV_DESCALE(coeff[0] * (ycrcb[1] - HALF_MAX), yuv_shift); + int g = ycrcb[0] + CV_DESCALE(coeff[1] * (ycrcb[1] - HALF_MAX) + coeff[2] * (ycrcb[2] - HALF_MAX), yuv_shift); + int b = ycrcb[0] + CV_DESCALE(coeff[3] * (ycrcb[2] - HALF_MAX), yuv_shift); +#endif + + 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) + dst[dst_idx + 3] = MAX_NUM; + } +} diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index f5f9f43..b2f5c6f 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -203,22 +203,22 @@ OCL_TEST_P(CvtColor, BGRA2YCrCb) { doTest(4, 3, CVTCODE(BGR2YCrCb)); } -//OCL_TEST_P(CvtColor, YCrCb2RGB) -//{ -// doTest(3, 3, CVTCODE(YCrCb2RGB)); -//} -//OCL_TEST_P(CvtColor, YCrCb2BGR) -//{ -// doTest(3, 3, CVTCODE(YCrCb2BGR)); -//} -//OCL_TEST_P(CvtColor, YCrCb2RGBA) -//{ -// doTest(3, 4, CVTCODE(YCrCb2RGB)); -//} -//OCL_TEST_P(CvtColor, YCrCb2BGRA) -//{ -// doTest(3, 4, CVTCODE(YCrCb2BGR)); -//} +OCL_TEST_P(CvtColor, YCrCb2RGB) +{ + doTest(3, 3, CVTCODE(YCrCb2RGB)); +} +OCL_TEST_P(CvtColor, YCrCb2BGR) +{ + doTest(3, 3, CVTCODE(YCrCb2BGR)); +} +OCL_TEST_P(CvtColor, YCrCb2RGBA) +{ + doTest(3, 4, CVTCODE(YCrCb2RGB)); +} +OCL_TEST_P(CvtColor, YCrCb2BGRA) +{ + doTest(3, 4, CVTCODE(YCrCb2BGR)); +} struct CvtColor_YUV420 : public CvtColor -- 2.7.4