From 506c19616d1888bdd4e771dc349e3e7f6268ed0f Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 19:24:27 +0400 Subject: [PATCH] YCrCb -> RGB[A] --- modules/imgproc/src/color.cpp | 25 ++-- modules/imgproc/src/opencl/cvtcolor.cl | 65 +++++++++- modules/imgproc/test/ocl/test_color.cpp | 8 +- modules/ocl/src/opencl/cvt_color.cl | 216 -------------------------------- 4 files changed, 82 insertions(+), 232 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 0cbfb38..2711855 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -2711,10 +2711,8 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) case COLOR_BGR5652BGR: case COLOR_BGR5552BGR: case COLOR_BGR5652RGB: case COLOR_BGR5552RGB: case COLOR_BGR5652BGRA: case COLOR_BGR5552BGRA: case COLOR_BGR5652RGBA: case COLOR_BGR5552RGBA: */ - case COLOR_BGR2GRAY: - case COLOR_BGRA2GRAY: - case COLOR_RGB2GRAY: - case COLOR_RGBA2GRAY: + case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY: + case COLOR_RGB2GRAY: case COLOR_RGBA2GRAY: { CV_Assert(scn == 3 || scn == 4); bidx = code == COLOR_BGR2GRAY || code == COLOR_BGRA2GRAY ? 0 : 2; @@ -2752,10 +2750,8 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) format("-D depth=%d -D scn=3 -D dcn=%d -D bidx=%d", depth, dcn, bidx)); break; } - case COLOR_YUV2RGB_NV12: - case COLOR_YUV2BGR_NV12: - case COLOR_YUV2RGBA_NV12: - case COLOR_YUV2BGRA_NV12: + case COLOR_YUV2RGB_NV12: case COLOR_YUV2BGR_NV12: + case COLOR_YUV2RGBA_NV12: case COLOR_YUV2BGRA_NV12: { CV_Assert( scn == 1 ); CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); @@ -2779,11 +2775,18 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) } case COLOR_YCrCb2BGR: case COLOR_YCrCb2RGB: + { + if( dcn <= 0 ) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); + bidx = code == COLOR_YCrCb2BGR ? 0 : 2; + k.create("YCrCb2RGB", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D scn=%d -D dcn=%d -D bidx=%d", depth, scn, dcn, bidx)); break; + } /* case COLOR_BGR5652GRAY: case COLOR_BGR5552GRAY: case COLOR_GRAY2BGR565: case COLOR_GRAY2BGR555: - case COLOR_BGR2YCrCb: case COLOR_RGB2YCrCb: case COLOR_BGR2XYZ: case COLOR_RGB2XYZ: case COLOR_XYZ2BGR: case COLOR_XYZ2RGB: case COLOR_BGR2HSV: case COLOR_RGB2HSV: case COLOR_BGR2HSV_FULL: case COLOR_RGB2HSV_FULL: @@ -2817,8 +2820,8 @@ void cv::cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) int stype = _src.type(); int scn = CV_MAT_CN(stype), depth = CV_MAT_DEPTH(stype), bidx; - if( use_opencl && ocl_cvtColor(_src, _dst, code, dcn) ) - return; + if( use_opencl /*&& ocl_cvtColor(_src, _dst, code, dcn)*/ ) + return (void)ocl_cvtColor(_src, _dst, code, dcn); Mat src = _src.getMat(), dst; Size sz = src.size(); diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 06ff7d4..e211142 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -268,7 +268,7 @@ __kernel void YUV2RGB_NV12(__global const uchar* srcptr, int srcstep, int srcoff } } -///////////////////////////////////// RGB -> YCrCb ////////////////////////////////////// +///////////////////////////////////// RGB <-> YCrCb ////////////////////////////////////// __constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f}; __constant int c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241}; @@ -304,3 +304,66 @@ __kernel void RGB2YCrCb(__global const uchar* srcptr, int srcstep, int srcoffset dst[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(__global const uchar* src, int src_step, int src_offset, + __global uchar* dst, int dst_step, int dst_offset, + int rows, int cols) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (y < rows && x < cols) + { + int src_idx = mad24(y, src_step, src_offset + x * scnbytes); + int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); + __global const DATA_TYPE * srcptr = (__global const DATA_TYPE*)(src + src_idx); + __global DATA_TYPE * dstptr = (__global DATA_TYPE*)(dst + dst_idx); + + DATA_TYPE y = srcptr[0], cr = srcptr[1], cb = srcptr[2]; + +#ifdef DEPTH_5 + __constant float * coeff = c_YCrCb2RGBCoeffs_f; + float r = y + coeff[0] * (cr - HALF_MAX); + float g = y + coeff[1] * (cr - HALF_MAX) + coeff[2] * (cb - HALF_MAX); + float b = y + coeff[3] * (cb - HALF_MAX); +#else + __constant int * coeff = c_YCrCb2RGBCoeffs_i; + int r = y + CV_DESCALE(coeff[0] * (cr - HALF_MAX), yuv_shift); + int g = y + CV_DESCALE(coeff[1] * (cr - HALF_MAX) + coeff[2] * (cb - HALF_MAX), yuv_shift); + int b = y + CV_DESCALE(coeff[3] * (cb - HALF_MAX), yuv_shift); +#endif + + dstptr[(bidx^2)] = SAT_CAST(r); + dstptr[1] = SAT_CAST(g); + dstptr[bidx] = SAT_CAST(b); +#if dcn == 4 + dstptr[3] = MAX_NUM; +#endif + } +} + + + + + + + + + + + + + + + + + + + + + + + diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index cef0c92..93ede91 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -150,10 +150,10 @@ OCL_TEST_P(CvtColor, RGB2YCrCb) { performTest(3, 3, CVTCODE(RGB2YCrCb)); } OCL_TEST_P(CvtColor, BGR2YCrCb) { performTest(3, 3, CVTCODE(BGR2YCrCb)); } OCL_TEST_P(CvtColor, RGBA2YCrCb) { performTest(4, 3, CVTCODE(RGB2YCrCb)); } OCL_TEST_P(CvtColor, BGRA2YCrCb) { performTest(4, 3, CVTCODE(BGR2YCrCb)); } -//OCL_TEST_P(CvtColor, YCrCb2RGB) { performTest(3, 3, CVTCODE(YCrCb2RGB)); } -//OCL_TEST_P(CvtColor, YCrCb2BGR) { performTest(3, 3, CVTCODE(YCrCb2BGR)); } -//OCL_TEST_P(CvtColor, YCrCb2RGBA) { performTest(3, 4, CVTCODE(YCrCb2RGB)); } -//OCL_TEST_P(CvtColor, YCrCb2BGRA) { performTest(3, 4, CVTCODE(YCrCb2BGR)); } +OCL_TEST_P(CvtColor, YCrCb2RGB) { performTest(3, 3, CVTCODE(YCrCb2RGB)); } +OCL_TEST_P(CvtColor, YCrCb2BGR) { performTest(3, 3, CVTCODE(YCrCb2BGR)); } +OCL_TEST_P(CvtColor, YCrCb2RGBA) { performTest(3, 4, CVTCODE(YCrCb2RGB)); } +OCL_TEST_P(CvtColor, YCrCb2BGRA) { performTest(3, 4, CVTCODE(YCrCb2BGR)); } // RGB <-> XYZ diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index bf3b6cf..be9aa99 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -91,222 +91,6 @@ enum BLOCK_SIZE = 256 }; -///////////////////////////////////// RGB <-> GRAY ////////////////////////////////////// - -__kernel void RGB2Gray(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) -{ - int x = get_global_id(0); - int y = get_global_id(1); - - if (y < rows && x < cols) - { - int src_idx = mad24(y, src_step, src_offset + (x << 2)); - int dst_idx = mad24(y, dst_step, dst_offset + x); -#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); -#endif - } -} - -__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) -{ - int x = get_global_id(0); - int y = get_global_id(1); - - if (y < rows && x < cols) - { - int src_idx = mad24(y, src_step, src_offset + x); - int dst_idx = mad24(y, dst_step, dst_offset + (x << 2)); - - DATA_TYPE val = src[src_idx]; - dst[dst_idx] = val; - dst[dst_idx + 1] = val; - dst[dst_idx + 2] = val; -#if dcn == 4 - dst[dst_idx + 3] = MAX_NUM; -#endif - } -} - -///////////////////////////////////// RGB <-> YUV ////////////////////////////////////// - -__constant float c_RGB2YUVCoeffs_f[5] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f }; -__constant int c_RGB2YUVCoeffs_i[5] = { B2Y, G2Y, R2Y, 8061, 14369 }; - -__kernel void RGB2YUV(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) -{ - 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 rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; - -#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; - DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX; -#else - __constant int * coeffs = c_RGB2YUVCoeffs_i; - int delta = HALF_MAX * (1 << yuv_shift); - int Y = CV_DESCALE(rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx], yuv_shift); - int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift); - int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift); -#endif - - dst[dst_idx] = SAT_CAST( Y ); - dst[dst_idx + 1] = SAT_CAST( Cr ); - dst[dst_idx + 2] = SAT_CAST( Cb ); - } -} - -__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 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 yuv[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; - -#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]; - float r = yuv[0] + (yuv[1] - HALF_MAX) * coeffs[0]; -#else - __constant int * coeffs = c_YUV2RGBCoeffs_i; - int b = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[3], yuv_shift); - int g = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1], yuv_shift); - int r = yuv[0] + CV_DESCALE((yuv[1] - HALF_MAX) * coeffs[0], yuv_shift); -#endif - - dst[dst_idx + bidx] = SAT_CAST( b ); - dst[dst_idx + 1] = SAT_CAST( g ); - dst[dst_idx + (bidx^2)] = SAT_CAST( r ); -#if dcn == 4 - dst[dst_idx + 3] = MAX_NUM; -#endif - } -} - -__constant int ITUR_BT_601_CY = 1220542; -__constant int ITUR_BT_601_CUB = 2116026; -__constant int ITUR_BT_601_CUG = 409993; -__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, __global const uchar* src, __global uchar* dst, - int src_offset, int dst_offset) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - - 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(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); - - int Y1 = ysrc[0]; - int Y2 = ysrc[1]; - int Y3 = ysrc[src_step]; - int Y4 = ysrc[src_step + 1]; - - int U = usrc[0] - 128; - int V = usrc[1] - 128; - - int ruv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CVR * V; - int guv = (1 << (ITUR_BT_601_SHIFT - 1)) - ITUR_BT_601_CVG * V - ITUR_BT_601_CUG * U; - int buv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CUB * U; - - Y1 = max(0, Y1 - 16) * ITUR_BT_601_CY; - dst1[2 - bidx] = convert_uchar_sat((Y1 + ruv) >> ITUR_BT_601_SHIFT); - dst1[1] = convert_uchar_sat((Y1 + guv) >> ITUR_BT_601_SHIFT); - dst1[bidx] = convert_uchar_sat((Y1 + buv) >> ITUR_BT_601_SHIFT); - dst1[3] = 255; - - Y2 = max(0, Y2 - 16) * ITUR_BT_601_CY; - dst1[6 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT); - dst1[5] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT); - dst1[4 + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT); - dst1[7] = 255; - - Y3 = max(0, Y3 - 16) * ITUR_BT_601_CY; - dst2[2 - bidx] = convert_uchar_sat((Y3 + ruv) >> ITUR_BT_601_SHIFT); - dst2[1] = convert_uchar_sat((Y3 + guv) >> ITUR_BT_601_SHIFT); - dst2[bidx] = convert_uchar_sat((Y3 + buv) >> ITUR_BT_601_SHIFT); - dst2[3] = 255; - - Y4 = max(0, Y4 - 16) * ITUR_BT_601_CY; - dst2[6 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT); - dst2[5] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT); - dst2[4 + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT); - dst2[7] = 255; - } -} - -///////////////////////////////////// RGB <-> YCrCb ////////////////////////////////////// - -__constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f}; -__constant int c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241}; - -__kernel void RGB2YCrCb(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) -{ - 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 rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; - -#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; - DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX; -#else - __constant int * coeffs = c_RGB2YCrCbCoeffs_i; - int delta = HALF_MAX * (1 << yuv_shift); - int Y = CV_DESCALE(rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx], yuv_shift); - int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift); - int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift); -#endif - - dst[dst_idx] = SAT_CAST( Y ); - dst[dst_idx + 1] = SAT_CAST( Cr ); - 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 }; -- 2.7.4