From 506c19616d1888bdd4e771dc349e3e7f6268ed0f Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 19:24:27 +0400 Subject: [PATCH 01/11] 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 0cbfb381b..2711855cb 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 06ff7d4e5..e21114218 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 cef0c9239..93ede912b 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 bf3b6cfa7..be9aa99db 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 }; From 65ee06eb2b6fd73e6770debb97253fa828a73db0 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 19:25:26 +0400 Subject: [PATCH 02/11] RGB[A] <-> XYZ --- modules/imgproc/src/opencl/cvtcolor.cl | 65 +++++++++++++++++++++++++- modules/ocl/src/opencl/cvt_color.cl | 40 ---------------- 2 files changed, 64 insertions(+), 41 deletions(-) diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index e21114218..0822b5607 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -345,7 +345,70 @@ __kernel void YCrCb2RGB(__global const uchar* src, int src_step, int src_offset, } } - +///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// + +__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, __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 r = src[src_idx], g = src[src_idx + 1], b = src[src_idx + 2]; + +#ifdef DEPTH_5 + 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 + 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 dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index be9aa99db..8078eda41 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -91,46 +91,6 @@ enum BLOCK_SIZE = 256 }; - -__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 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 dcn == 4 - dst[dst_idx + 3] = MAX_NUM; -#endif - } -} - ///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// __kernel void RGB2XYZ(int cols, int rows, int src_step, int dst_step, From ab9b883c69bf94428bce17aaf2a49a055fc244c9 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 19:37:27 +0400 Subject: [PATCH 03/11] RGB[A] <-> XYZ --- modules/imgproc/src/color.cpp | 96 +++++++++++++++++++++++++ modules/imgproc/src/opencl/cvtcolor.cl | 49 +++++++------ modules/imgproc/test/ocl/test_color.cpp | 16 ++--- modules/ocl/src/opencl/cvt_color.cl | 65 ----------------- 4 files changed, 132 insertions(+), 94 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 2711855cb..49fabbb96 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -2784,6 +2784,102 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) format("-D depth=%d -D scn=%d -D dcn=%d -D bidx=%d", depth, scn, dcn, bidx)); break; } + case COLOR_BGR2XYZ: case COLOR_RGB2XYZ: + { + CV_Assert(scn == 3 || scn == 4); + bidx = code == COLOR_BGR2XYZ ? 0 : 2; + + UMat c; + if (depth == CV_32F) + { + float coeffs[] = + { + 0.412453f, 0.357580f, 0.180423f, + 0.212671f, 0.715160f, 0.072169f, + 0.019334f, 0.119193f, 0.950227f + }; + if (bidx == 0) + { + std::swap(coeffs[0], coeffs[2]); + std::swap(coeffs[3], coeffs[5]); + std::swap(coeffs[6], coeffs[8]); + } + Mat(1, 9, CV_32FC1, &coeffs[0]).copyTo(c); + } + else + { + int coeffs[] = + { + 1689, 1465, 739, + 871, 2929, 296, + 79, 488, 3892 + }; + if (bidx == 0) + { + std::swap(coeffs[0], coeffs[2]); + std::swap(coeffs[3], coeffs[5]); + std::swap(coeffs[6], coeffs[8]); + } + Mat(1, 9, CV_32SC1, &coeffs[0]).copyTo(c); + } + + _dst.create(dstSz, CV_MAKETYPE(depth, 3)); + dst = _dst.getUMat(); + + k.create("RGB2XYZ", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D scn=%d -D dcn=3 -D bidx=%d", depth, scn, bidx)); + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::ReadOnlyNoSize(c)); + return k.run(2, globalsize, 0, false); + } + case COLOR_XYZ2BGR: case COLOR_XYZ2RGB: + { + if (dcn <= 0) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); + bidx = code == COLOR_XYZ2BGR ? 0 : 2; + + UMat c; + 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]); + } + Mat(1, 9, CV_32FC1, &coeffs[0]).copyTo(c); + } + 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]); + } + Mat(1, 9, CV_32SC1, &coeffs[0]).copyTo(c); + } + + _dst.create(dstSz, CV_MAKETYPE(depth, dcn)); + dst = _dst.getUMat(); + + k.create("XYZ2RGB", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D scn=3 -D dcn=%d -D bidx=%d", depth, dcn, bidx)); + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::ReadOnlyNoSize(c)); + return k.run(2, globalsize, 0, false); + } /* case COLOR_BGR5652GRAY: case COLOR_BGR5552GRAY: case COLOR_GRAY2BGR565: case COLOR_GRAY2BGR555: diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 0822b5607..37dab79a8 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -54,18 +54,21 @@ #define DATA_TYPE uchar #define MAX_NUM 255 #define HALF_MAX 128 + #define COEFF_TYPE int #define SAT_CAST(num) convert_uchar_sat(num) #define DEPTH_0 #elif depth == 2 #define DATA_TYPE ushort #define MAX_NUM 65535 #define HALF_MAX 32768 + #define COEFF_TYPE int #define SAT_CAST(num) convert_ushort_sat(num) #define DEPTH_2 #elif depth == 5 #define DATA_TYPE float #define MAX_NUM 1.0f #define HALF_MAX 0.5f + #define COEFF_TYPE float #define SAT_CAST(num) (num) #define DEPTH_5 #else @@ -347,20 +350,22 @@ __kernel void YCrCb2RGB(__global const uchar* src, int src_step, int src_offset, ///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// -__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, __constant COEFF_TYPE * coeffs) +__kernel void RGB2XYZ(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, + int rows, int cols, __constant COEFF_TYPE * coeffs, int a1, int a2) { 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); + int src_idx = mad24(dy, src_step, src_offset + dx * scnbytes); + int dst_idx = mad24(dy, dst_step, dst_offset + dx * dcnbytes); - DATA_TYPE r = src[src_idx], g = src[src_idx + 1], b = src[src_idx + 2]; + __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_idx); + __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_idx); + + DATA_TYPE r = src[0], g = src[1], b = src[2]; #ifdef DEPTH_5 float x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2]; @@ -371,26 +376,28 @@ __kernel void RGB2XYZ(int cols, int rows, int src_step, int dst_step, 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); + dst[0] = SAT_CAST(x); + dst[1] = SAT_CAST(y); + dst[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) +__kernel void XYZ2RGB(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, + int rows, int cols, __constant COEFF_TYPE * coeffs, int a1, int a2) { 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); + int src_idx = mad24(dy, src_step, src_offset + dx * scnbytes); + int dst_idx = mad24(dy, dst_step, dst_offset + dx * dcnbytes); - DATA_TYPE x = src[src_idx], y = src[src_idx + 1], z = src[src_idx + 2]; + __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_idx); + __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_idx); + + DATA_TYPE x = src[0], y = src[1], z = src[2]; #ifdef DEPTH_5 float b = x * coeffs[0] + y * coeffs[1] + z * coeffs[2]; @@ -401,11 +408,11 @@ __kernel void XYZ2RGB(int cols, int rows, int src_step, int dst_step, 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); + dst[0] = SAT_CAST(b); + dst[1] = SAT_CAST(g); + dst[2] = SAT_CAST(r); #if dcn == 4 - dst[dst_idx + 3] = MAX_NUM; + dst[3] = MAX_NUM; #endif } } diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 93ede912b..541c5f698 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -157,15 +157,15 @@ OCL_TEST_P(CvtColor, YCrCb2BGRA) { performTest(3, 4, CVTCODE(YCrCb2BGR)); } // RGB <-> XYZ -//OCL_TEST_P(CvtColor, RGB2XYZ) { performTest(3, 3, CVTCODE(RGB2XYZ)); } -//OCL_TEST_P(CvtColor, BGR2XYZ) { performTest(3, 3, CVTCODE(BGR2XYZ)); } -//OCL_TEST_P(CvtColor, RGBA2XYZ) { performTest(4, 3, CVTCODE(RGB2XYZ)); } -//OCL_TEST_P(CvtColor, BGRA2XYZ) { performTest(4, 3, CVTCODE(BGR2XYZ)); } +OCL_TEST_P(CvtColor, RGB2XYZ) { performTest(3, 3, CVTCODE(RGB2XYZ)); } +OCL_TEST_P(CvtColor, BGR2XYZ) { performTest(3, 3, CVTCODE(BGR2XYZ)); } +OCL_TEST_P(CvtColor, RGBA2XYZ) { performTest(4, 3, CVTCODE(RGB2XYZ)); } +OCL_TEST_P(CvtColor, BGRA2XYZ) { performTest(4, 3, CVTCODE(BGR2XYZ)); } -//OCL_TEST_P(CvtColor, XYZ2RGB) { performTest(3, 3, CVTCODE(XYZ2RGB)); } -//OCL_TEST_P(CvtColor, XYZ2BGR) { performTest(3, 3, CVTCODE(XYZ2BGR)); } -//OCL_TEST_P(CvtColor, XYZ2RGBA) { performTest(3, 4, CVTCODE(XYZ2RGB)); } -//OCL_TEST_P(CvtColor, XYZ2BGRA) { performTest(3, 4, CVTCODE(XYZ2BGR)); } +OCL_TEST_P(CvtColor, XYZ2RGB) { performTest(3, 3, CVTCODE(XYZ2RGB)); } +OCL_TEST_P(CvtColor, XYZ2BGR) { performTest(3, 3, CVTCODE(XYZ2BGR)); } +OCL_TEST_P(CvtColor, XYZ2RGBA) { performTest(3, 4, CVTCODE(XYZ2RGB)); } +OCL_TEST_P(CvtColor, XYZ2BGRA) { performTest(3, 4, CVTCODE(XYZ2BGR)); } // RGB <-> HSV diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 8078eda41..0e058a791 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -91,71 +91,6 @@ enum BLOCK_SIZE = 256 }; -///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// - -__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, __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 r = src[src_idx], g = src[src_idx + 1], b = src[src_idx + 2]; - -#ifdef DEPTH_5 - 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 - 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 dcn == 4 - dst[dst_idx + 3] = MAX_NUM; -#endif - } -} - ///////////////////////////////////// RGB[A] <-> BGR[A] ////////////////////////////////////// __kernel void RGB(int cols, int rows, int src_step, int dst_step, From af367c4f8588e1a68488595c5e1be05e9dd63be0 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 19:44:37 +0400 Subject: [PATCH 04/11] RGB[A] <-> BGR[A] --- modules/imgproc/src/color.cpp | 13 +++++++-- modules/imgproc/src/opencl/cvtcolor.cl | 37 ++++++++++++++++++++++++- modules/imgproc/test/ocl/test_color.cpp | 24 ++++++++-------- 3 files changed, 59 insertions(+), 15 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 49fabbb96..ab6806f8e 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -2704,13 +2704,22 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) switch (code) { /* - case COLOR_BGR2BGRA: case COLOR_RGB2BGRA: case COLOR_BGRA2BGR: - case COLOR_RGBA2BGR: case COLOR_RGB2BGR: case COLOR_BGRA2RGBA: case COLOR_BGR2BGR565: case COLOR_BGR2BGR555: case COLOR_RGB2BGR565: case COLOR_RGB2BGR555: case COLOR_BGRA2BGR565: case COLOR_BGRA2BGR555: case COLOR_RGBA2BGR565: case COLOR_RGBA2BGR555: 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_BGR2BGRA: case COLOR_RGB2BGRA: case COLOR_BGRA2BGR: + case COLOR_RGBA2BGR: case COLOR_RGB2BGR: case COLOR_BGRA2RGBA: + { + CV_Assert(scn == 3 || scn == 4); + dcn = code == COLOR_BGR2BGRA || code == COLOR_RGB2BGRA || code == COLOR_BGRA2RGBA ? 4 : 3; + bool reverse = !(code == COLOR_BGR2BGRA || code == COLOR_BGRA2BGR); + k.create("RGB", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D scn=%d -D dcn=%d -D bidx=0 -D %s", depth, scn, dcn, + reverse ? "REVERSE" : "ORDER")); + break; + } case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY: case COLOR_RGB2GRAY: case COLOR_RGBA2GRAY: { diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 37dab79a8..e8711cbfc 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -417,7 +417,42 @@ __kernel void XYZ2RGB(__global const uchar * srcptr, int src_step, int src_offse } } - +///////////////////////////////////// RGB[A] <-> BGR[A] ////////////////////////////////////// + +__kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, 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 * src = (__global const DATA_TYPE *)(srcptr + src_idx); + __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_idx); + +#ifdef REVERSE + dst[0] = src[2]; + dst[1] = src[1]; + dst[2] = src[0]; +#elif defined ORDER + dst[0] = src[0]; + dst[1] = src[1]; + dst[2] = src[2]; +#endif + +#if dcn == 4 +#if scn == 3 + dst[3] = MAX_NUM; +#else + dst[3] = src[3]; +#endif +#endif + } +} diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 541c5f698..dc81b31f8 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -109,18 +109,18 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) // RGB[A] <-> BGR[A] -//OCL_TEST_P(CvtColor, BGR2BGRA) { performTest(3, 4, CVTCODE(BGR2BGRA)); } -//OCL_TEST_P(CvtColor, RGB2RGBA) { performTest(3, 4, CVTCODE(RGB2RGBA)); } -//OCL_TEST_P(CvtColor, BGRA2BGR) { performTest(4, 3, CVTCODE(BGRA2BGR)); } -//OCL_TEST_P(CvtColor, RGBA2RGB) { performTest(4, 3, CVTCODE(RGBA2RGB)); } -//OCL_TEST_P(CvtColor, BGR2RGBA) { performTest(3, 4, CVTCODE(BGR2RGBA)); } -//OCL_TEST_P(CvtColor, RGB2BGRA) { performTest(3, 4, CVTCODE(RGB2BGRA)); } -//OCL_TEST_P(CvtColor, RGBA2BGR) { performTest(4, 3, CVTCODE(RGBA2BGR)); } -//OCL_TEST_P(CvtColor, BGRA2RGB) { performTest(4, 3, CVTCODE(BGRA2RGB)); } -//OCL_TEST_P(CvtColor, BGR2RGB) { performTest(3, 3, CVTCODE(BGR2RGB)); } -//OCL_TEST_P(CvtColor, RGB2BGR) { performTest(3, 3, CVTCODE(RGB2BGR)); } -//OCL_TEST_P(CvtColor, BGRA2RGBA) { performTest(4, 4, CVTCODE(BGRA2RGBA)); } -//OCL_TEST_P(CvtColor, RGBA2BGRA) { performTest(4, 4, CVTCODE(RGBA2BGRA)); } +OCL_TEST_P(CvtColor, BGR2BGRA) { performTest(3, 4, CVTCODE(BGR2BGRA)); } +OCL_TEST_P(CvtColor, RGB2RGBA) { performTest(3, 4, CVTCODE(RGB2RGBA)); } +OCL_TEST_P(CvtColor, BGRA2BGR) { performTest(4, 3, CVTCODE(BGRA2BGR)); } +OCL_TEST_P(CvtColor, RGBA2RGB) { performTest(4, 3, CVTCODE(RGBA2RGB)); } +OCL_TEST_P(CvtColor, BGR2RGBA) { performTest(3, 4, CVTCODE(BGR2RGBA)); } +OCL_TEST_P(CvtColor, RGB2BGRA) { performTest(3, 4, CVTCODE(RGB2BGRA)); } +OCL_TEST_P(CvtColor, RGBA2BGR) { performTest(4, 3, CVTCODE(RGBA2BGR)); } +OCL_TEST_P(CvtColor, BGRA2RGB) { performTest(4, 3, CVTCODE(BGRA2RGB)); } +OCL_TEST_P(CvtColor, BGR2RGB) { performTest(3, 3, CVTCODE(BGR2RGB)); } +OCL_TEST_P(CvtColor, RGB2BGR) { performTest(3, 3, CVTCODE(RGB2BGR)); } +OCL_TEST_P(CvtColor, BGRA2RGBA) { performTest(4, 4, CVTCODE(BGRA2RGBA)); } +OCL_TEST_P(CvtColor, RGBA2BGRA) { performTest(4, 4, CVTCODE(RGBA2BGRA)); } // RGB <-> Gray From 81b9c9c104724fde1667074bbcf56350aaea8657 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 19:52:42 +0400 Subject: [PATCH 05/11] RGB[A] -> RGB5x5 --- modules/imgproc/src/color.cpp | 15 ++++++- modules/imgproc/src/opencl/cvtcolor.cl | 57 +++++++++++++++++++++++++ modules/imgproc/test/ocl/test_color.cpp | 18 ++++---- modules/ocl/src/opencl/cvt_color.cl | 34 --------------- 4 files changed, 79 insertions(+), 45 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index ab6806f8e..c375b0533 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -2706,8 +2706,6 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) /* case COLOR_BGR2BGR565: case COLOR_BGR2BGR555: case COLOR_RGB2BGR565: case COLOR_RGB2BGR555: case COLOR_BGRA2BGR565: case COLOR_BGRA2BGR555: case COLOR_RGBA2BGR565: case COLOR_RGBA2BGR555: - 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_BGR2BGRA: case COLOR_RGB2BGRA: case COLOR_BGRA2BGR: case COLOR_RGBA2BGR: case COLOR_RGB2BGR: case COLOR_BGRA2RGBA: @@ -2720,6 +2718,19 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) reverse ? "REVERSE" : "ORDER")); break; } + case COLOR_BGR5652BGR: case COLOR_BGR5552BGR: case COLOR_BGR5652RGB: case COLOR_BGR5552RGB: + case COLOR_BGR5652BGRA: case COLOR_BGR5552BGRA: case COLOR_BGR5652RGBA: case COLOR_BGR5552RGBA: + { + dcn = code == COLOR_BGR5652BGRA || code == COLOR_BGR5552BGRA || code == COLOR_BGR5652RGBA || code == COLOR_BGR5552RGBA ? 4 : 3; + CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U); + bidx = code == COLOR_BGR5652BGR || code == COLOR_BGR5552BGR || + code == COLOR_BGR5652BGRA || code == COLOR_BGR5552BGRA ? 0 : 2; + int greenbits = code == COLOR_BGR5652BGR || code == COLOR_BGR5652RGB || + code == COLOR_BGR5652BGRA || code == COLOR_BGR5652RGBA ? 6 : 5; + k.create("RGB5x52RGB", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D scn=2 -D dcn=%s -D bidx=%d -D greenbits=%d", depth, scn, dcn, bidx, greenbits)); + break; + } case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY: case COLOR_RGB2GRAY: case COLOR_RGBA2GRAY: { diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index e8711cbfc..e6ceada8e 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -455,6 +455,63 @@ __kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset, } +///////////////////////////////////// RGB5x5 <-> RGB ////////////////////////////////////// + +__kernel void RGB5x52RGB(__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); + ushort t = *((__global const ushort*)(src + src_idx)); + +#if greenbits == 6 + dst[dst_idx + bidx] = (uchar)(t << 3); + dst[dst_idx + 1] = (uchar)((t >> 3) & ~3); + dst[dst_idx + (bidx^2)] = (uchar)((t >> 8) & ~7); +#else + dst[dst_idx + bidx] = (uchar)(t << 3); + dst[dst_idx + 1] = (uchar)((t >> 2) & ~7); + dst[dst_idx + (bidx^2)] = (uchar)((t >> 7) & ~7); +#endif + +#if dcn == 4 +#if greenbits == 6 + dst[dst_idx + 3] = 255; +#else + dst[dst_idx + 3] = t & 0x8000 ? 255 : 0; +#endif +#endif + } +} + +//__kernel void RGB2RGB5x5(int cols, int rows, int src_step, int dst_step, int bidx, +// __global const uchar * src, __global ushort * 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); + +//#if greenbits == 6 +// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~3) << 3)|((src[src_idx + (bidx^2)]&~7) << 8)); +//#elif scn == 3 +// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|((src[src_idx + (bidx^2)]&~7) << 7)); +//#else +// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)| +// ((src[src_idx + (bidx^2)]&~7) << 7)|(src[src_idx + 3] ? 0x8000 : 0)); +//#endif +// } +//} diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index dc81b31f8..d65c6f240 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -215,17 +215,17 @@ OCL_TEST_P(CvtColor, XYZ2BGRA) { performTest(3, 4, CVTCODE(XYZ2BGR)); } // RGB5x5 <-> RGB -//typedef CvtColor CvtColor8u; +typedef CvtColor CvtColor8u; -//OCL_TEST_P(CvtColor8u, BGR5652BGR) { performTest(2, 3, CVTCODE(BGR5652BGR)); } -//OCL_TEST_P(CvtColor8u, BGR5652RGB) { performTest(2, 3, CVTCODE(BGR5652RGB)); } -//OCL_TEST_P(CvtColor8u, BGR5652BGRA) { performTest(2, 4, CVTCODE(BGR5652BGRA)); } -//OCL_TEST_P(CvtColor8u, BGR5652RGBA) { performTest(2, 4, CVTCODE(BGR5652RGBA)); } +OCL_TEST_P(CvtColor8u, BGR5652BGR) { performTest(2, 3, CVTCODE(BGR5652BGR)); } +OCL_TEST_P(CvtColor8u, BGR5652RGB) { performTest(2, 3, CVTCODE(BGR5652RGB)); } +OCL_TEST_P(CvtColor8u, BGR5652BGRA) { performTest(2, 4, CVTCODE(BGR5652BGRA)); } +OCL_TEST_P(CvtColor8u, BGR5652RGBA) { performTest(2, 4, CVTCODE(BGR5652RGBA)); } -//OCL_TEST_P(CvtColor8u, BGR5552BGR) { performTest(2, 3, CVTCODE(BGR5552BGR)); } -//OCL_TEST_P(CvtColor8u, BGR5552RGB) { performTest(2, 3, CVTCODE(BGR5552RGB)); } -//OCL_TEST_P(CvtColor8u, BGR5552BGRA) { performTest(2, 4, CVTCODE(BGR5552BGRA)); } -//OCL_TEST_P(CvtColor8u, BGR5552RGBA) { performTest(2, 4, CVTCODE(BGR5552RGBA)); } +OCL_TEST_P(CvtColor8u, BGR5552BGR) { performTest(2, 3, CVTCODE(BGR5552BGR)); } +OCL_TEST_P(CvtColor8u, BGR5552RGB) { performTest(2, 3, CVTCODE(BGR5552RGB)); } +OCL_TEST_P(CvtColor8u, BGR5552BGRA) { performTest(2, 4, CVTCODE(BGR5552BGRA)); } +OCL_TEST_P(CvtColor8u, BGR5552RGBA) { performTest(2, 4, CVTCODE(BGR5552RGBA)); } //OCL_TEST_P(CvtColor8u, BGR2BGR565) { performTest(3, 2, CVTCODE(BGR2BGR565)); } //OCL_TEST_P(CvtColor8u, RGB2BGR565) { performTest(3, 2, CVTCODE(RGB2BGR565)); } diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 0e058a791..5786505bd 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -91,40 +91,6 @@ enum BLOCK_SIZE = 256 }; -///////////////////////////////////// RGB[A] <-> BGR[A] ////////////////////////////////////// - -__kernel void RGB(int cols, int rows, int src_step, int dst_step, - __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); - -#ifdef REVERSE - dst[dst_idx] = src[src_idx + 2]; - dst[dst_idx + 1] = src[src_idx + 1]; - dst[dst_idx + 2] = src[src_idx]; -#elif defined ORDER - dst[dst_idx] = src[src_idx]; - dst[dst_idx + 1] = src[src_idx + 1]; - dst[dst_idx + 2] = src[src_idx + 2]; -#endif - -#if dcn == 4 -#if scn == 3 - dst[dst_idx + 3] = MAX_NUM; -#else - dst[dst_idx + 3] = src[src_idx + 3]; -#endif -#endif - } -} ///////////////////////////////////// RGB5x5 <-> RGB ////////////////////////////////////// From 8a236468973be7817231c2da73f61e0d27da0eaa Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 20:00:35 +0400 Subject: [PATCH 06/11] RGB[A] <- RGB5x5 --- modules/imgproc/src/color.cpp | 16 ++++++- modules/imgproc/src/opencl/cvtcolor.cl | 40 ++++++++-------- modules/imgproc/test/ocl/test_color.cpp | 26 +++++------ modules/ocl/src/opencl/cvt_color.cl | 61 +------------------------ 4 files changed, 49 insertions(+), 94 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index c375b0533..88655757c 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -92,6 +92,7 @@ #include "precomp.hpp" #include "opencl_kernels.hpp" #include +#include #define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n)) @@ -2728,7 +2729,20 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) int greenbits = code == COLOR_BGR5652BGR || code == COLOR_BGR5652RGB || code == COLOR_BGR5652BGRA || code == COLOR_BGR5652RGBA ? 6 : 5; k.create("RGB5x52RGB", ocl::imgproc::cvtcolor_oclsrc, - format("-D depth=%d -D scn=2 -D dcn=%s -D bidx=%d -D greenbits=%d", depth, scn, dcn, bidx, greenbits)); + format("-D depth=%d -D scn=2 -D dcn=%d -D bidx=%d -D greenbits=%d", depth, dcn, bidx, greenbits)); + break; + } + case COLOR_BGR2BGR565: case COLOR_BGR2BGR555: case COLOR_RGB2BGR565: case COLOR_RGB2BGR555: + case COLOR_BGRA2BGR565: case COLOR_BGRA2BGR555: case COLOR_RGBA2BGR565: case COLOR_RGBA2BGR555: + { + CV_Assert((scn == 3 || scn == 4) && depth == CV_8U ); + bidx = code == COLOR_BGR2BGR565 || code == COLOR_BGR2BGR555 || + code == COLOR_BGRA2BGR565 || code == COLOR_BGRA2BGR555 ? 0 : 2; + int greenbits = code == COLOR_BGR2BGR565 || code == COLOR_RGB2BGR565 || + code == COLOR_BGRA2BGR565 || code == COLOR_RGBA2BGR565 ? 6 : 5; + dcn = 2; + k.create("RGB2RGB5x5", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D scn=%d -D dcn=2 -D bidx=%d -D greenbits=%d", depth, scn, bidx, greenbits)); break; } case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY: diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index e6ceada8e..0d8272fae 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -490,28 +490,28 @@ __kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset } } -//__kernel void RGB2RGB5x5(int cols, int rows, int src_step, int dst_step, int bidx, -// __global const uchar * src, __global ushort * dst, -// int src_offset, int dst_offset) -//{ -// int x = get_global_id(0); -// int y = get_global_id(1); +__kernel void RGB2RGB5x5(__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 << 2)); -// int dst_idx = mad24(y, dst_step, dst_offset + x); + 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); -//#if greenbits == 6 -// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~3) << 3)|((src[src_idx + (bidx^2)]&~7) << 8)); -//#elif scn == 3 -// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|((src[src_idx + (bidx^2)]&~7) << 7)); -//#else -// dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)| -// ((src[src_idx + (bidx^2)]&~7) << 7)|(src[src_idx + 3] ? 0x8000 : 0)); -//#endif -// } -//} +#if greenbits == 6 + *((__global ushort*)(dst + dst_idx)) = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~3) << 3)|((src[src_idx + (bidx^2)]&~7) << 8)); +#elif scn == 3 + *((__global ushort*)(dst + dst_idx)) = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|((src[src_idx + (bidx^2)]&~7) << 7)); +#else + *((__global ushort*)(dst + dst_idx)) = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)| + ((src[src_idx + (bidx^2)]&~7) << 7)|(src[src_idx + 3] ? 0x8000 : 0)); +#endif + } +} diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index d65c6f240..695d3d164 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -169,7 +169,7 @@ OCL_TEST_P(CvtColor, XYZ2BGRA) { performTest(3, 4, CVTCODE(XYZ2BGR)); } // RGB <-> HSV -//typedef CvtColor CvtColor8u32f; +typedef CvtColor CvtColor8u32f; //OCL_TEST_P(CvtColor8u32f, RGB2HSV) { performTest(3, 3, CVTCODE(RGB2HSV)); } //OCL_TEST_P(CvtColor8u32f, BGR2HSV) { performTest(3, 3, CVTCODE(BGR2HSV)); } @@ -227,15 +227,15 @@ OCL_TEST_P(CvtColor8u, BGR5552RGB) { performTest(2, 3, CVTCODE(BGR5552RGB)); } OCL_TEST_P(CvtColor8u, BGR5552BGRA) { performTest(2, 4, CVTCODE(BGR5552BGRA)); } OCL_TEST_P(CvtColor8u, BGR5552RGBA) { performTest(2, 4, CVTCODE(BGR5552RGBA)); } -//OCL_TEST_P(CvtColor8u, BGR2BGR565) { performTest(3, 2, CVTCODE(BGR2BGR565)); } -//OCL_TEST_P(CvtColor8u, RGB2BGR565) { performTest(3, 2, CVTCODE(RGB2BGR565)); } -//OCL_TEST_P(CvtColor8u, BGRA2BGR565) { performTest(4, 2, CVTCODE(BGRA2BGR565)); } -//OCL_TEST_P(CvtColor8u, RGBA2BGR565) { performTest(4, 2, CVTCODE(RGBA2BGR565)); } +OCL_TEST_P(CvtColor8u, BGR2BGR565) { performTest(3, 2, CVTCODE(BGR2BGR565)); } +OCL_TEST_P(CvtColor8u, RGB2BGR565) { performTest(3, 2, CVTCODE(RGB2BGR565)); } +OCL_TEST_P(CvtColor8u, BGRA2BGR565) { performTest(4, 2, CVTCODE(BGRA2BGR565)); } +OCL_TEST_P(CvtColor8u, RGBA2BGR565) { performTest(4, 2, CVTCODE(RGBA2BGR565)); } -//OCL_TEST_P(CvtColor8u, BGR2BGR555) { performTest(3, 2, CVTCODE(BGR2BGR555)); } -//OCL_TEST_P(CvtColor8u, RGB2BGR555) { performTest(3, 2, CVTCODE(RGB2BGR555)); } -//OCL_TEST_P(CvtColor8u, BGRA2BGR555) { performTest(4, 2, CVTCODE(BGRA2BGR555)); } -//OCL_TEST_P(CvtColor8u, RGBA2BGR555) { performTest(4, 2, CVTCODE(RGBA2BGR555)); } +OCL_TEST_P(CvtColor8u, BGR2BGR555) { performTest(3, 2, CVTCODE(BGR2BGR555)); } +OCL_TEST_P(CvtColor8u, RGB2BGR555) { performTest(3, 2, CVTCODE(RGB2BGR555)); } +OCL_TEST_P(CvtColor8u, BGRA2BGR555) { performTest(4, 2, CVTCODE(BGRA2BGR555)); } +OCL_TEST_P(CvtColor8u, RGBA2BGR555) { performTest(4, 2, CVTCODE(RGBA2BGR555)); } // RGB5x5 <-> Gray @@ -280,11 +280,11 @@ OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { performTest(1, 3, COLOR_YUV2RGB_NV12 OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { performTest(1, 3, COLOR_YUV2BGR_NV12); } -//OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u, -// testing::Combine(testing::Values(MatDepth(CV_8U)), Bool())); +OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u, + testing::Combine(testing::Values(MatDepth(CV_8U)), Bool())); -//OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u32f, -// testing::Combine(testing::Values(MatDepth(CV_8U), MatDepth(CV_32F)), Bool())); +OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u32f, + testing::Combine(testing::Values(MatDepth(CV_8U), MatDepth(CV_32F)), Bool())); OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor, testing::Combine( diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 5786505bd..3b2e84c0a 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -91,66 +91,7 @@ enum BLOCK_SIZE = 256 }; - -///////////////////////////////////// RGB5x5 <-> RGB ////////////////////////////////////// - -__kernel void RGB5x52RGB(int cols, int rows, int src_step, int dst_step, int bidx, - __global const ushort * 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) - { - int src_idx = mad24(y, src_step, src_offset + x); - int dst_idx = mad24(y, dst_step, dst_offset + (x << 2)); - ushort t = src[src_idx]; - -#if greenbits == 6 - dst[dst_idx + bidx] = (uchar)(t << 3); - dst[dst_idx + 1] = (uchar)((t >> 3) & ~3); - dst[dst_idx + (bidx^2)] = (uchar)((t >> 8) & ~7); -#else - dst[dst_idx + bidx] = (uchar)(t << 3); - dst[dst_idx + 1] = (uchar)((t >> 2) & ~7); - dst[dst_idx + (bidx^2)] = (uchar)((t >> 7) & ~7); -#endif - -#if dcn == 4 -#if greenbits == 6 - dst[dst_idx + 3] = 255; -#else - dst[dst_idx + 3] = t & 0x8000 ? 255 : 0; -#endif -#endif - } -} - -__kernel void RGB2RGB5x5(int cols, int rows, int src_step, int dst_step, int bidx, - __global const uchar * src, __global ushort * 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); - -#if greenbits == 6 - dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~3) << 3)|((src[src_idx + (bidx^2)]&~7) << 8)); -#elif scn == 3 - dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|((src[src_idx + (bidx^2)]&~7) << 7)); -#else - dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)| - ((src[src_idx + (bidx^2)]&~7) << 7)|(src[src_idx + 3] ? 0x8000 : 0)); -#endif - } -} - -///////////////////////////////////// RGB5x5 <-> RGB ////////////////////////////////////// +///////////////////////////////////// RGB5x5 <-> Gray ////////////////////////////////////// __kernel void BGR5x52Gray(int cols, int rows, int src_step, int dst_step, int bidx, __global const ushort * src, __global uchar * dst, From 727a5e6df471111fb0627504371b5f9591800b5e Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 22:19:44 +0400 Subject: [PATCH 07/11] BGR5x5 <-> Gray --- modules/imgproc/src/color.cpp | 22 +++++++++-- modules/imgproc/src/opencl/cvtcolor.cl | 50 ++++++++++++++++++++++++- modules/imgproc/test/ocl/test_color.cpp | 8 ++-- 3 files changed, 70 insertions(+), 10 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 88655757c..f5f8118f4 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -2704,10 +2704,6 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) switch (code) { - /* - case COLOR_BGR2BGR565: case COLOR_BGR2BGR555: case COLOR_RGB2BGR565: case COLOR_RGB2BGR555: - case COLOR_BGRA2BGR565: case COLOR_BGRA2BGR555: case COLOR_RGBA2BGR565: case COLOR_RGBA2BGR555: - */ case COLOR_BGR2BGRA: case COLOR_RGB2BGRA: case COLOR_BGRA2BGR: case COLOR_RGBA2BGR: case COLOR_RGB2BGR: case COLOR_BGRA2RGBA: { @@ -2745,6 +2741,24 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) format("-D depth=%d -D scn=%d -D dcn=2 -D bidx=%d -D greenbits=%d", depth, scn, bidx, greenbits)); break; } + case COLOR_BGR5652GRAY: case COLOR_BGR5552GRAY: + { + CV_Assert(scn == 2 && depth == CV_8U); + dcn = 1; + int greenbits = code == COLOR_BGR5652GRAY ? 6 : 5; + k.create("BGR5x52Gray", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D scn=2 -D dcn=1 -D bidx=0 -D greenbits=%d", depth, greenbits)); + break; + } + case COLOR_GRAY2BGR565: case COLOR_GRAY2BGR555: + { + CV_Assert(scn == 1 && depth == CV_8U); + dcn = 2; + int greenbits = code == COLOR_GRAY2BGR565 ? 6 : 5; + k.create("Gray2BGR5x5", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D scn=1 -D dcn=2 -D bidx=0 -D greenbits=%d", depth, greenbits)); + break; + } case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY: case COLOR_RGB2GRAY: case COLOR_RGBA2GRAY: { diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 0d8272fae..846574160 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -513,8 +513,54 @@ __kernel void RGB2RGB5x5(__global const uchar* src, int src_step, int src_offset } } - - +///////////////////////////////////// RGB5x5 <-> Gray ////////////////////////////////////// + +__kernel void BGR5x52Gray(__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); + int t = *((__global const ushort*)(src + src_idx)); + +#if greenbits == 6 + dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + + ((t >> 3) & 0xfc)*G2Y + + ((t >> 8) & 0xf8)*R2Y, yuv_shift); +#else + dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + + ((t >> 2) & 0xf8)*G2Y + + ((t >> 7) & 0xf8)*R2Y, yuv_shift); +#endif + } +} + +__kernel void Gray2BGR5x5(__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); + int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); + int t = src[src_idx]; + +#if greenbits == 6 + *((__global ushort*)(dst + dst_idx)) = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8)); +#else + t >>= 3; + *((__global ushort*)(dst + dst_idx)) = (ushort)(t|(t << 5)|(t << 10)); +#endif + } +} diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 695d3d164..0dfcae84a 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -239,11 +239,11 @@ OCL_TEST_P(CvtColor8u, RGBA2BGR555) { performTest(4, 2, CVTCODE(RGBA2BGR555)); } // RGB5x5 <-> Gray -//OCL_TEST_P(CvtColor8u, BGR5652GRAY) { performTest(2, 1, CVTCODE(BGR5652GRAY)); } -//OCL_TEST_P(CvtColor8u, BGR5552GRAY) { performTest(2, 1, CVTCODE(BGR5552GRAY)); } +OCL_TEST_P(CvtColor8u, BGR5652GRAY) { performTest(2, 1, CVTCODE(BGR5652GRAY)); } +OCL_TEST_P(CvtColor8u, BGR5552GRAY) { performTest(2, 1, CVTCODE(BGR5552GRAY)); } -//OCL_TEST_P(CvtColor8u, GRAY2BGR565) { performTest(1, 2, CVTCODE(GRAY2BGR565)); } -//OCL_TEST_P(CvtColor8u, GRAY2BGR555) { performTest(1, 2, CVTCODE(GRAY2BGR555)); } +OCL_TEST_P(CvtColor8u, GRAY2BGR565) { performTest(1, 2, CVTCODE(GRAY2BGR565)); } +OCL_TEST_P(CvtColor8u, GRAY2BGR555) { performTest(1, 2, CVTCODE(GRAY2BGR555)); } // RGBA <-> mRGBA From 0b900b54e54e8cbd19a65cfe69d20838f142e6b3 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 23:30:29 +0400 Subject: [PATCH 08/11] RGB[A] <-> HSV --- modules/core/include/opencv2/core/ocl.hpp | 4 +- modules/core/src/ocl.cpp | 8 +- modules/imgproc/src/color.cpp | 96 ++++++++-- modules/imgproc/src/opencl/cvtcolor.cl | 221 +++++++++++++++++++++- modules/imgproc/test/ocl/test_color.cpp | 32 ++-- modules/ocl/src/opencl/cvt_color.cl | 51 +---- 6 files changed, 327 insertions(+), 85 deletions(-) diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index 9a3096206..971e4de8a 100644 --- a/modules/core/include/opencv2/core/ocl.hpp +++ b/modules/core/include/opencv2/core/ocl.hpp @@ -245,11 +245,13 @@ protected: class CV_EXPORTS KernelArg { public: - enum { LOCAL=1, READ_ONLY=2, WRITE_ONLY=4, READ_WRITE=6, CONSTANT=8, NO_SIZE=256 }; + enum { LOCAL=1, READ_ONLY=2, WRITE_ONLY=4, READ_WRITE=6, CONSTANT=8, PTR_ONLY = 16, NO_SIZE=256 }; KernelArg(int _flags, UMat* _m, int wscale=1, const void* _obj=0, size_t _sz=0); KernelArg(); static KernelArg Local() { return KernelArg(LOCAL, 0); } + static KernelArg PtrOnly(const UMat & m) + { return KernelArg(PTR_ONLY, (UMat*)&m); } static KernelArg ReadWrite(const UMat& m, int wscale=1) { return KernelArg(READ_WRITE, (UMat*)&m, wscale); } static KernelArg ReadWriteNoSize(const UMat& m, int wscale=1) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 3133a069f..8eed5fbe7 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -41,6 +41,7 @@ #include "precomp.hpp" #include +#include /* Part of the file is an extract from the standard OpenCL headers from Khronos site. @@ -2205,9 +2206,12 @@ int Kernel::set(int i, const KernelArg& arg) { int accessFlags = ((arg.flags & KernelArg::READ_ONLY) ? ACCESS_READ : 0) + ((arg.flags & KernelArg::WRITE_ONLY) ? ACCESS_WRITE : 0); + bool ptronly = (arg.flags & KernelArg::PTR_ONLY) != 0; cl_mem h = (cl_mem)arg.m->handle(accessFlags); - if( arg.m->dims <= 2 ) + if (ptronly) + clSetKernelArg(p->handle, (cl_uint)i++, sizeof(h), &h); + else if( arg.m->dims <= 2 ) { UMat2D u2d(*arg.m); clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h); @@ -2350,7 +2354,7 @@ struct Program::Impl retval = clBuildProgram(handle, n, (const cl_device_id*)deviceList, buildflags.c_str(), 0, 0); - if( retval == CL_BUILD_PROGRAM_FAILURE ) + if( retval < 0 /*== CL_BUILD_PROGRAM_FAILURE*/ ) { char buf[1<<16]; size_t retsz = 0; diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index f5f8118f4..5ce28114f 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -2876,7 +2876,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) k.create("RGB2XYZ", ocl::imgproc::cvtcolor_oclsrc, format("-D depth=%d -D scn=%d -D dcn=3 -D bidx=%d", depth, scn, bidx)); - k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::ReadOnlyNoSize(c)); + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrOnly(c)); return k.run(2, globalsize, 0, false); } case COLOR_XYZ2BGR: case COLOR_XYZ2RGB: @@ -2925,19 +2925,91 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) k.create("XYZ2RGB", ocl::imgproc::cvtcolor_oclsrc, format("-D depth=%d -D scn=3 -D dcn=%d -D bidx=%d", depth, dcn, bidx)); - k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::ReadOnlyNoSize(c)); + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrOnly(c)); return k.run(2, globalsize, 0, false); } - /* - case COLOR_BGR5652GRAY: case COLOR_BGR5552GRAY: - case COLOR_GRAY2BGR565: case COLOR_GRAY2BGR555: - 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: - case COLOR_BGR2HLS: case COLOR_RGB2HLS: case COLOR_BGR2HLS_FULL: case COLOR_RGB2HLS_FULL: - case COLOR_HSV2BGR: case COLOR_HSV2RGB: case COLOR_HSV2BGR_FULL: case COLOR_HSV2RGB_FULL: - case COLOR_HLS2BGR: case COLOR_HLS2RGB: case COLOR_HLS2BGR_FULL: case COLOR_HLS2RGB_FULL: - */ + case COLOR_BGR2HSV: case COLOR_RGB2HSV: case COLOR_BGR2HSV_FULL: case COLOR_RGB2HSV_FULL: + case COLOR_BGR2HLS: case COLOR_RGB2HLS: case COLOR_BGR2HLS_FULL: case COLOR_RGB2HLS_FULL: + { + CV_Assert((scn == 3 || scn == 4) && (depth == CV_8U || depth == CV_32F)); + bidx = code == COLOR_BGR2HSV || code == COLOR_BGR2HLS || + code == COLOR_BGR2HSV_FULL || code == COLOR_BGR2HLS_FULL ? 0 : 2; + int hrange = depth == CV_32F ? 360 : code == COLOR_BGR2HSV || code == COLOR_RGB2HSV || + code == COLOR_BGR2HLS || code == COLOR_RGB2HLS ? 180 : 256; + bool is_hsv = code == COLOR_BGR2HSV || code == COLOR_RGB2HSV || code == COLOR_BGR2HSV_FULL || code == COLOR_RGB2HSV_FULL; + String kernelName = String("RGB2") + (is_hsv ? "HSV" : "HLS"); + dcn = 3; + + if (is_hsv && depth == CV_8U) + { + static UMat sdiv_data; + static UMat hdiv_data180; + static UMat hdiv_data256; + static int sdiv_table[256]; + static int hdiv_table180[256]; + static int hdiv_table256[256]; + static volatile bool initialized180 = false, initialized256 = false; + volatile bool & initialized = hrange == 180 ? initialized180 : initialized256; + + if (!initialized) + { + int * const hdiv_table = hrange == 180 ? hdiv_table180 : hdiv_table256, hsv_shift = 12; + UMat & hdiv_data = hrange == 180 ? hdiv_data180 : hdiv_data256; + + sdiv_table[0] = hdiv_table180[0] = hdiv_table256[0] = 0; + + int v = 255 << hsv_shift; + if (!initialized180 && !initialized256) + { + for(int i = 1; i < 256; i++ ) + sdiv_table[i] = saturate_cast(v/(1.*i)); + Mat(1, 256, CV_32SC1, sdiv_table).copyTo(sdiv_data); + } + + v = hrange << hsv_shift; + for (int i = 1; i < 256; i++ ) + hdiv_table[i] = saturate_cast(v/(6.*i)); + + Mat(1, 256, CV_32SC1, hdiv_table).copyTo(hdiv_data); + initialized = true; + } + + _dst.create(dstSz, CV_8UC3); + dst = _dst.getUMat(); + + k.create("RGB2HSV", ocl::imgproc::cvtcolor_oclsrc, format("-D depth=%d -D hrange=%d -D bidx=%d -D dcn=3 -D scn=%d", + depth, hrange, bidx, scn)); + + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), + ocl::KernelArg::PtrOnly(sdiv_data), hrange == 256 ? ocl::KernelArg::PtrOnly(hdiv_data256) : + ocl::KernelArg::PtrOnly(hdiv_data180)); + + return k.run(2, globalsize, NULL, false); + } + else + k.create(kernelName.c_str(), ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D hscale=%f -D bidx=%d -D scn=%d -D dcn=3", depth, hrange*(1.f/360.f), bidx, scn)); + break; + } + case COLOR_HSV2BGR: case COLOR_HSV2RGB: case COLOR_HSV2BGR_FULL: case COLOR_HSV2RGB_FULL: + case COLOR_HLS2BGR: case COLOR_HLS2RGB: case COLOR_HLS2BGR_FULL: case COLOR_HLS2RGB_FULL: + { + if (dcn <= 0) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F)); + bidx = code == COLOR_HSV2BGR || code == COLOR_HLS2BGR || + code == COLOR_HSV2BGR_FULL || code == COLOR_HLS2BGR_FULL ? 0 : 2; + int hrange = depth == CV_32F ? 360 : code == COLOR_HSV2BGR || code == COLOR_HSV2RGB || + code == COLOR_HLS2BGR || code == COLOR_HLS2RGB ? 180 : 255; + bool is_hsv = code == COLOR_HSV2BGR || code == COLOR_HSV2RGB || + code == COLOR_HSV2BGR_FULL || code == COLOR_HSV2RGB_FULL; + + String kernelName = String(is_hsv ? "HSV" : "HLS") + "2RGB"; + k.create(kernelName.c_str(), ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D dcn=%d -D scn=3 -D bidx=%d -D hrange=%d -D hscale=%f", + depth, dcn, bidx, hrange, 6.f/hrange)); + break; + } default: ; } diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 846574160..b34e8b621 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -81,6 +81,7 @@ enum { yuv_shift = 14, xyz_shift = 12, + hsv_shift = 12, R2Y = 4899, G2Y = 9617, B2Y = 1868, @@ -90,6 +91,14 @@ enum #define scnbytes ((int)sizeof(DATA_TYPE)*scn) #define dcnbytes ((int)sizeof(DATA_TYPE)*dcn) +#ifndef hscale +#define hscale 0 +#endif + +#ifndef hrange +#define hrange 0 +#endif + ///////////////////////////////////// RGB <-> GRAY ////////////////////////////////////// __kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset, @@ -352,7 +361,7 @@ __kernel void YCrCb2RGB(__global const uchar* src, int src_step, int src_offset, __kernel void RGB2XYZ(__global const uchar * srcptr, int src_step, int src_offset, __global uchar * dstptr, int dst_step, int dst_offset, - int rows, int cols, __constant COEFF_TYPE * coeffs, int a1, int a2) + int rows, int cols, __constant COEFF_TYPE * coeffs) { int dx = get_global_id(0); int dy = get_global_id(1); @@ -384,7 +393,7 @@ __kernel void RGB2XYZ(__global const uchar * srcptr, int src_step, int src_offse __kernel void XYZ2RGB(__global const uchar * srcptr, int src_step, int src_offset, __global uchar * dstptr, int dst_step, int dst_offset, - int rows, int cols, __constant COEFF_TYPE * coeffs, int a1, int a2) + int rows, int cols, __constant COEFF_TYPE * coeffs) { int dx = get_global_id(0); int dy = get_global_id(1); @@ -454,7 +463,6 @@ __kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset, } } - ///////////////////////////////////// RGB5x5 <-> RGB ////////////////////////////////////// __kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset, @@ -562,7 +570,212 @@ __kernel void Gray2BGR5x5(__global const uchar* src, int src_step, int src_offse } } - +//////////////////////////////////// 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(__global const uchar* src, int src_step, int src_offset, + __global uchar* dst, int dst_step, int dst_offset, + int rows, int cols, + __constant int * sdiv_table, __constant int * hdiv_table) +{ + 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); + + int b = src[src_idx + bidx], g = src[src_idx + 1], r = src[src_idx + (bidx^2)]; + int h, s, v = b; + int vmin = b, diff; + int vr, vg; + + v = max( v, g ); + v = max( v, r ); + vmin = min( vmin, g ); + vmin = min( vmin, r ); + + diff = v - vmin; + vr = v == r ? -1 : 0; + vg = v == g ? -1 : 0; + + s = (diff * sdiv_table[v] + (1 << (hsv_shift-1))) >> hsv_shift; + h = (vr & (g - b)) + + (~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff)))); + h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift; + h += h < 0 ? hrange : 0; + + dst[dst_idx] = convert_uchar_sat_rte(h); + dst[dst_idx + 1] = (uchar)s; + dst[dst_idx + 2] = (uchar)v; + } +} + +__kernel void HSV2RGB(__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); + + 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(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, 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 float * src = (__global const float *)(srcptr + src_idx); + __global float * dst = (__global float *)(dstptr + dst_idx); + + float b = src[bidx], g = src[1], r = src[bidx^2]; + float h, s, v; + + float vmin, diff; + + v = vmin = r; + if( v < g ) v = g; + if( v < b ) v = b; + if( vmin > g ) vmin = g; + if( vmin > b ) vmin = b; + + diff = v - vmin; + s = diff/(float)(fabs(v) + FLT_EPSILON); + diff = (float)(60./(diff + FLT_EPSILON)); + if( v == r ) + h = (g - b)*diff; + else if( v == g ) + h = (b - r)*diff + 120.f; + else + h = (r - g)*diff + 240.f; + + if( h < 0 ) h += 360.f; + + dst[0] = h*hscale; + dst[1] = s; + dst[2] = v; + } +} + +__kernel void HSV2RGB(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, 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 float * src = (__global const float *)(srcptr + src_idx); + __global float * dst = (__global float *)(dstptr + dst_idx); + + float h = src[0], s = src[1], v = src[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[bidx] = b; + dst[1] = g; + dst[bidx^2] = r; +#if dcn == 4 + dst[3] = MAX_NUM; +#endif + } +} + +#endif diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 0dfcae84a..5af5b98a6 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -171,25 +171,25 @@ OCL_TEST_P(CvtColor, XYZ2BGRA) { performTest(3, 4, CVTCODE(XYZ2BGR)); } typedef CvtColor CvtColor8u32f; -//OCL_TEST_P(CvtColor8u32f, RGB2HSV) { performTest(3, 3, CVTCODE(RGB2HSV)); } -//OCL_TEST_P(CvtColor8u32f, BGR2HSV) { performTest(3, 3, CVTCODE(BGR2HSV)); } -//OCL_TEST_P(CvtColor8u32f, RGBA2HSV) { performTest(4, 3, CVTCODE(RGB2HSV)); } -//OCL_TEST_P(CvtColor8u32f, BGRA2HSV) { performTest(4, 3, CVTCODE(BGR2HSV)); } +OCL_TEST_P(CvtColor8u32f, RGB2HSV) { performTest(3, 3, CVTCODE(RGB2HSV)); } +OCL_TEST_P(CvtColor8u32f, BGR2HSV) { performTest(3, 3, CVTCODE(BGR2HSV)); } +OCL_TEST_P(CvtColor8u32f, RGBA2HSV) { performTest(4, 3, CVTCODE(RGB2HSV)); } +OCL_TEST_P(CvtColor8u32f, BGRA2HSV) { performTest(4, 3, CVTCODE(BGR2HSV)); } -//OCL_TEST_P(CvtColor8u32f, RGB2HSV_FULL) { performTest(3, 3, CVTCODE(RGB2HSV_FULL)); } -//OCL_TEST_P(CvtColor8u32f, BGR2HSV_FULL) { performTest(3, 3, CVTCODE(BGR2HSV_FULL)); } -//OCL_TEST_P(CvtColor8u32f, RGBA2HSV_FULL) { performTest(4, 3, CVTCODE(RGB2HSV_FULL)); } -//OCL_TEST_P(CvtColor8u32f, BGRA2HSV_FULL) { performTest(4, 3, CVTCODE(BGR2HSV_FULL)); } +OCL_TEST_P(CvtColor8u32f, RGB2HSV_FULL) { performTest(3, 3, CVTCODE(RGB2HSV_FULL)); } +OCL_TEST_P(CvtColor8u32f, BGR2HSV_FULL) { performTest(3, 3, CVTCODE(BGR2HSV_FULL)); } +OCL_TEST_P(CvtColor8u32f, RGBA2HSV_FULL) { performTest(4, 3, CVTCODE(RGB2HSV_FULL)); } +OCL_TEST_P(CvtColor8u32f, BGRA2HSV_FULL) { performTest(4, 3, CVTCODE(BGR2HSV_FULL)); } -//OCL_TEST_P(CvtColor8u32f, HSV2RGB) { performTest(3, 3, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); } -//OCL_TEST_P(CvtColor8u32f, HSV2BGR) { performTest(3, 3, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); } -//OCL_TEST_P(CvtColor8u32f, HSV2RGBA) { performTest(3, 4, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); } -//OCL_TEST_P(CvtColor8u32f, HSV2BGRA) { performTest(3, 4, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2RGB) { performTest(3, 3, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2BGR) { performTest(3, 3, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2RGBA) { performTest(3, 4, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2BGRA) { performTest(3, 4, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); } -//OCL_TEST_P(CvtColor8u32f, HSV2RGB_FULL) { performTest(3, 3, CVTCODE(HSV2RGB_FULL), depth == CV_8U ? 1 : 4e-1); } -//OCL_TEST_P(CvtColor8u32f, HSV2BGR_FULL) { performTest(3, 3, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } -//OCL_TEST_P(CvtColor8u32f, HSV2RGBA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } -//OCL_TEST_P(CvtColor8u32f, HSV2BGRA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2RGB_FULL) { performTest(3, 3, CVTCODE(HSV2RGB_FULL), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2BGR_FULL) { performTest(3, 3, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2RGBA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } +OCL_TEST_P(CvtColor8u32f, HSV2BGRA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); } // RGB <-> HLS diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 3b2e84c0a..75b23825b 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -91,56 +91,7 @@ enum BLOCK_SIZE = 256 }; -///////////////////////////////////// RGB5x5 <-> Gray ////////////////////////////////////// - -__kernel void BGR5x52Gray(int cols, int rows, int src_step, int dst_step, int bidx, - __global const ushort * 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) - { - int src_idx = mad24(y, src_step, src_offset + x); - int dst_idx = mad24(y, dst_step, dst_offset + x); - int t = src[src_idx]; - -#if greenbits == 6 - dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + - ((t >> 3) & 0xfc)*G2Y + - ((t >> 8) & 0xf8)*R2Y, yuv_shift); -#else - dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + - ((t >> 2) & 0xf8)*G2Y + - ((t >> 7) & 0xf8)*R2Y, yuv_shift); -#endif - } -} - -__kernel void Gray2BGR5x5(int cols, int rows, int src_step, int dst_step, int bidx, - __global const uchar * src, __global ushort * 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); - int t = src[src_idx]; - -#if greenbits == 6 - dst[dst_idx] = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8)); -#else - t >>= 3; - dst[dst_idx] = (ushort)(t|(t << 5)|(t << 10)); -#endif - } -} - -///////////////////////////////////// RGB <-> HSV ////////////////////////////////////// +//////////////////////////////////// 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 } }; From f771a0ba8116014cddfcbea4ac3a0359d4b9b48e Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 23:37:58 +0400 Subject: [PATCH 09/11] RGB[A] <-> HLS --- modules/imgproc/src/opencl/cvtcolor.cl | 211 +++++++++++++++++++++++- modules/imgproc/test/ocl/test_color.cpp | 32 ++-- modules/ocl/src/opencl/cvt_color.cl | 198 ---------------------- 3 files changed, 226 insertions(+), 215 deletions(-) diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index b34e8b621..53bb5a95b 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -777,7 +777,216 @@ __kernel void HSV2RGB(__global const uchar* srcptr, int src_step, int src_offset #endif - +///////////////////////////////////// RGB <-> HLS ////////////////////////////////////// + +#ifdef DEPTH_0 + +__kernel void RGB2HLS(__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); + + float b = src[src_idx + bidx]*(1/255.f), g = src[src_idx + 1]*(1/255.f), r = src[src_idx + (bidx^2)]*(1/255.f); + float h = 0.f, s = 0.f, l; + float vmin, vmax, diff; + + vmax = vmin = r; + if (vmax < g) vmax = g; + if (vmax < b) vmax = b; + if (vmin > g) vmin = g; + if (vmin > b) vmin = b; + + diff = vmax - vmin; + l = (vmax + vmin)*0.5f; + + if (diff > FLT_EPSILON) + { + s = l < 0.5f ? diff/(vmax + vmin) : diff/(2 - vmax - vmin); + diff = 60.f/diff; + + if( vmax == r ) + h = (g - b)*diff; + else if( vmax == g ) + h = (b - r)*diff + 120.f; + else + h = (r - g)*diff + 240.f; + + if( h < 0.f ) h += 360.f; + } + + dst[dst_idx] = convert_uchar_sat_rte(h*hscale); + dst[dst_idx + 1] = convert_uchar_sat_rte(l*255.f); + dst[dst_idx + 2] = convert_uchar_sat_rte(s*255.f); + } +} + +__kernel void HLS2RGB(__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); + + float h = src[src_idx], l = src[src_idx + 1]*(1.f/255.f), s = src[src_idx + 2]*(1.f/255.f); + float b, g, r; + + if (s != 0) + { + float tab[4]; + + float p2 = l <= 0.5f ? l*(1 + s) : l + s - l*s; + float p1 = 2*l - p2; + + h *= hscale; + if( h < 0 ) + do h += 6; while( h < 0 ); + else if( h >= 6 ) + do h -= 6; while( h >= 6 ); + + int sector = convert_int_sat_rtn(h); + h -= sector; + + tab[0] = p2; + tab[1] = p1; + tab[2] = p1 + (p2 - p1)*(1-h); + tab[3] = p1 + (p2 - p1)*h; + + b = tab[sector_data[sector][0]]; + g = tab[sector_data[sector][1]]; + r = tab[sector_data[sector][2]]; + } + else + b = g = r = l; + + 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 RGB2HLS(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, 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 float * src = (__global const float *)(srcptr + src_idx); + __global float * dst = (__global float *)(dstptr + dst_idx); + + float b = src[bidx], g = src[1], r = src[bidx^2]; + float h = 0.f, s = 0.f, l; + float vmin, vmax, diff; + + vmax = vmin = r; + if (vmax < g) vmax = g; + if (vmax < b) vmax = b; + if (vmin > g) vmin = g; + if (vmin > b) vmin = b; + + diff = vmax - vmin; + l = (vmax + vmin)*0.5f; + + if (diff > FLT_EPSILON) + { + s = l < 0.5f ? diff/(vmax + vmin) : diff/(2 - vmax - vmin); + diff = 60.f/diff; + + if( vmax == r ) + h = (g - b)*diff; + else if( vmax == g ) + h = (b - r)*diff + 120.f; + else + h = (r - g)*diff + 240.f; + + if( h < 0.f ) h += 360.f; + } + + dst[0] = h*hscale; + dst[1] = l; + dst[2] = s; + } +} + +__kernel void HLS2RGB(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, 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 float * src = (__global const float *)(srcptr + src_idx); + __global float * dst = (__global float *)(dstptr + dst_idx); + + float h = src[0], l = src[1], s = src[2]; + float b, g, r; + + if (s != 0) + { + float tab[4]; + int sector; + + float p2 = l <= 0.5f ? l*(1 + s) : l + s - l*s; + float p1 = 2*l - p2; + + 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; + + tab[0] = p2; + tab[1] = p1; + tab[2] = p1 + (p2 - p1)*(1-h); + tab[3] = p1 + (p2 - p1)*h; + + b = tab[sector_data[sector][0]]; + g = tab[sector_data[sector][1]]; + r = tab[sector_data[sector][2]]; + } + else + b = g = r = l; + + dst[bidx] = b; + dst[1] = g; + dst[bidx^2] = r; +#if dcn == 4 + dst[3] = MAX_NUM; +#endif + } +} + +#endif diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index 5af5b98a6..aad04b619 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -193,25 +193,25 @@ OCL_TEST_P(CvtColor8u32f, HSV2BGRA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FUL // RGB <-> HLS -//OCL_TEST_P(CvtColor8u32f, RGB2HLS) { performTest(3, 3, CVTCODE(RGB2HLS), depth == CV_8U ? 1 : 1e-3); } -//OCL_TEST_P(CvtColor8u32f, BGR2HLS) { performTest(3, 3, CVTCODE(BGR2HLS), depth == CV_8U ? 1 : 1e-3); } -//OCL_TEST_P(CvtColor8u32f, RGBA2HLS) { performTest(4, 3, CVTCODE(RGB2HLS), depth == CV_8U ? 1 : 1e-3); } -//OCL_TEST_P(CvtColor8u32f, BGRA2HLS) { performTest(4, 3, CVTCODE(BGR2HLS), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, RGB2HLS) { performTest(3, 3, CVTCODE(RGB2HLS), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, BGR2HLS) { performTest(3, 3, CVTCODE(BGR2HLS), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, RGBA2HLS) { performTest(4, 3, CVTCODE(RGB2HLS), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, BGRA2HLS) { performTest(4, 3, CVTCODE(BGR2HLS), depth == CV_8U ? 1 : 1e-3); } -//OCL_TEST_P(CvtColor8u32f, RGB2HLS_FULL) { performTest(3, 3, CVTCODE(RGB2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } -//OCL_TEST_P(CvtColor8u32f, BGR2HLS_FULL) { performTest(3, 3, CVTCODE(BGR2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } -//OCL_TEST_P(CvtColor8u32f, RGBA2HLS_FULL) { performTest(4, 3, CVTCODE(RGB2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } -//OCL_TEST_P(CvtColor8u32f, BGRA2HLS_FULL) { performTest(4, 3, CVTCODE(BGR2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, RGB2HLS_FULL) { performTest(3, 3, CVTCODE(RGB2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, BGR2HLS_FULL) { performTest(3, 3, CVTCODE(BGR2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, RGBA2HLS_FULL) { performTest(4, 3, CVTCODE(RGB2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } +OCL_TEST_P(CvtColor8u32f, BGRA2HLS_FULL) { performTest(4, 3, CVTCODE(BGR2HLS_FULL), depth == CV_8U ? 1 : 1e-3); } -//OCL_TEST_P(CvtColor8u32f, HLS2RGB) { performTest(3, 3, CVTCODE(HLS2RGB), 1); } -//OCL_TEST_P(CvtColor8u32f, HLS2BGR) { performTest(3, 3, CVTCODE(HLS2BGR), 1); } -//OCL_TEST_P(CvtColor8u32f, HLS2RGBA) { performTest(3, 4, CVTCODE(HLS2RGB), 1); } -//OCL_TEST_P(CvtColor8u32f, HLS2BGRA) { performTest(3, 4, CVTCODE(HLS2BGR), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2RGB) { performTest(3, 3, CVTCODE(HLS2RGB), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2BGR) { performTest(3, 3, CVTCODE(HLS2BGR), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2RGBA) { performTest(3, 4, CVTCODE(HLS2RGB), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2BGRA) { performTest(3, 4, CVTCODE(HLS2BGR), 1); } -//OCL_TEST_P(CvtColor8u32f, HLS2RGB_FULL) { performTest(3, 3, CVTCODE(HLS2RGB_FULL), 1); } -//OCL_TEST_P(CvtColor8u32f, HLS2BGR_FULL) { performTest(3, 3, CVTCODE(HLS2BGR_FULL), 1); } -//OCL_TEST_P(CvtColor8u32f, HLS2RGBA_FULL) { performTest(3, 4, CVTCODE(HLS2RGB_FULL), 1); } -//OCL_TEST_P(CvtColor8u32f, HLS2BGRA_FULL) { performTest(3, 4, CVTCODE(HLS2BGR_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2RGB_FULL) { performTest(3, 3, CVTCODE(HLS2RGB_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2BGR_FULL) { performTest(3, 3, CVTCODE(HLS2BGR_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2RGBA_FULL) { performTest(3, 4, CVTCODE(HLS2RGB_FULL), 1); } +OCL_TEST_P(CvtColor8u32f, HLS2BGRA_FULL) { performTest(3, 4, CVTCODE(HLS2BGR_FULL), 1); } // RGB5x5 <-> RGB diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 75b23825b..6ef08d307 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -91,206 +91,8 @@ enum BLOCK_SIZE = 256 }; -//////////////////////////////////// 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, - __global const uchar * src, __global uchar * dst, - int src_offset, int dst_offset, - __constant int * sdiv_table, __constant int * hdiv_table) -{ - 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); - - int b = src[src_idx + bidx], g = src[src_idx + 1], r = src[src_idx + (bidx^2)]; - int h, s, v = b; - int vmin = b, diff; - int vr, vg; - - v = max( v, g ); - v = max( v, r ); - vmin = min( vmin, g ); - vmin = min( vmin, r ); - - diff = v - vmin; - vr = v == r ? -1 : 0; - vg = v == g ? -1 : 0; - - s = (diff * sdiv_table[v] + (1 << (hsv_shift-1))) >> hsv_shift; - h = (vr & (g - b)) + - (~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff)))); - h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift; - h += h < 0 ? hrange : 0; - - dst[dst_idx] = convert_uchar_sat_rte(h); - dst[dst_idx + 1] = (uchar)s; - dst[dst_idx + 2] = (uchar)v; - } -} - -__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, - __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 b = src[src_idx + bidx], g = src[src_idx + 1], r = src[src_idx + (bidx^2)]; - float h, s, v; - - float vmin, diff; - - v = vmin = r; - if( v < g ) v = g; - if( v < b ) v = b; - if( vmin > g ) vmin = g; - if( vmin > b ) vmin = b; - - diff = v - vmin; - s = diff/(float)(fabs(v) + FLT_EPSILON); - diff = (float)(60./(diff + FLT_EPSILON)); - if( v == r ) - h = (g - b)*diff; - else if( v == g ) - h = (b - r)*diff + 120.f; - else - h = (r - g)*diff + 240.f; - - if( h < 0 ) h += 360.f; - - dst[dst_idx] = h*hscale; - dst[dst_idx + 1] = s; - dst[dst_idx + 2] = v; - } -} - -__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 ////////////////////////////////////// #ifdef DEPTH_0 From d2e1318341a4361378bae92e9898bcec852febe8 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 23:42:10 +0400 Subject: [PATCH 10/11] RGBA <-> mRGBA --- modules/imgproc/src/color.cpp | 9 ++++++ modules/imgproc/src/opencl/cvtcolor.cl | 43 +++++++++++++++++++++++++ modules/imgproc/test/ocl/test_color.cpp | 4 +-- 3 files changed, 54 insertions(+), 2 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 5ce28114f..264490c6a 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -3010,6 +3010,15 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) depth, dcn, bidx, hrange, 6.f/hrange)); break; } + case COLOR_RGBA2mRGBA: case COLOR_mRGBA2RGBA: + { + CV_Assert(scn == 4 && depth == CV_8U); + dcn = 4; + + k.create(code == COLOR_RGBA2mRGBA ? "RGBA2mRGBA" : "mRGBA2RGBA", ocl::imgproc::cvtcolor_oclsrc, + format("-D depth=%d -D dcn=4 -D scn=4 -D bidx=3", depth)); + break; + } default: ; } diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 53bb5a95b..ca696290d 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -988,12 +988,55 @@ __kernel void HLS2RGB(__global const uchar* srcptr, int src_step, int src_offset #endif +/////////////////////////// RGBA <-> mRGBA (alpha premultiplied) ////////////// +#ifdef DEPTH_0 +__kernel void RGBA2mRGBA(__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) + { + x <<= 2; + int src_idx = mad24(y, src_step, src_offset + x); + int dst_idx = mad24(y, dst_step, dst_offset + x); + uchar v0 = src[src_idx], v1 = src[src_idx + 1]; + uchar v2 = src[src_idx + 2], v3 = src[src_idx + 3]; + dst[dst_idx] = (v0 * v3 + HALF_MAX) / MAX_NUM; + dst[dst_idx + 1] = (v1 * v3 + HALF_MAX) / MAX_NUM; + dst[dst_idx + 2] = (v2 * v3 + HALF_MAX) / MAX_NUM; + dst[dst_idx + 3] = v3; + } +} +__kernel void mRGBA2RGBA(__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) + { + x <<= 2; + int src_idx = mad24(y, src_step, src_offset + x); + int dst_idx = mad24(y, dst_step, dst_offset + x); + uchar v0 = src[src_idx], v1 = src[src_idx + 1]; + uchar v2 = src[src_idx + 2], v3 = src[src_idx + 3]; + uchar v3_half = v3 / 2; + dst[dst_idx] = v3 == 0 ? 0 : (v0 * MAX_NUM + v3_half) / v3; + dst[dst_idx + 1] = v3 == 0 ? 0 : (v1 * MAX_NUM + v3_half) / v3; + dst[dst_idx + 2] = v3 == 0 ? 0 : (v2 * MAX_NUM + v3_half) / v3; + dst[dst_idx + 3] = v3; + } +} + +#endif diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index aad04b619..22b6b5cda 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -247,8 +247,8 @@ OCL_TEST_P(CvtColor8u, GRAY2BGR555) { performTest(1, 2, CVTCODE(GRAY2BGR555)); } // RGBA <-> mRGBA -//OCL_TEST_P(CvtColor8u, RGBA2mRGBA) { performTest(4, 4, CVTCODE(RGBA2mRGBA)); } -//OCL_TEST_P(CvtColor8u, mRGBA2RGBA) { performTest(4, 4, CVTCODE(mRGBA2RGBA)); } +OCL_TEST_P(CvtColor8u, RGBA2mRGBA) { performTest(4, 4, CVTCODE(RGBA2mRGBA)); } +OCL_TEST_P(CvtColor8u, mRGBA2RGBA) { performTest(4, 4, CVTCODE(mRGBA2RGBA)); } // YUV -> RGBA_NV12 From 41d046a2db3d722b5f8095754513ab609b2f5ed5 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 23:42:49 +0400 Subject: [PATCH 11/11] restored ocl_cvtcolor.cl --- modules/core/src/ocl.cpp | 3 +- modules/imgproc/src/color.cpp | 5 +- modules/ocl/src/opencl/cvt_color.cl | 661 ++++++++++++++++++++++++++++ 3 files changed, 664 insertions(+), 5 deletions(-) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 8eed5fbe7..48f44a600 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -41,7 +41,6 @@ #include "precomp.hpp" #include -#include /* Part of the file is an extract from the standard OpenCL headers from Khronos site. @@ -2354,7 +2353,7 @@ struct Program::Impl retval = clBuildProgram(handle, n, (const cl_device_id*)deviceList, buildflags.c_str(), 0, 0); - if( retval < 0 /*== CL_BUILD_PROGRAM_FAILURE*/ ) + if( retval == CL_BUILD_PROGRAM_FAILURE ) { char buf[1<<16]; size_t retsz = 0; diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 264490c6a..13f474d80 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -92,7 +92,6 @@ #include "precomp.hpp" #include "opencl_kernels.hpp" #include -#include #define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n)) @@ -3045,8 +3044,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 (void)ocl_cvtColor(_src, _dst, code, dcn); + if( use_opencl && ocl_cvtColor(_src, _dst, code, dcn) ) + return; Mat src = _src.getMat(), dst; Size sz = src.size(); diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 6ef08d307..bf3b6cfa7 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -91,8 +91,669 @@ 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 }; + +__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) +{ + 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 dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} + +///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// + +__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, __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 r = src[src_idx], g = src[src_idx + 1], b = src[src_idx + 2]; + +#ifdef DEPTH_5 + 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 + 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 dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} + +///////////////////////////////////// RGB[A] <-> BGR[A] ////////////////////////////////////// + +__kernel void RGB(int cols, int rows, int src_step, int dst_step, + __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); + +#ifdef REVERSE + dst[dst_idx] = src[src_idx + 2]; + dst[dst_idx + 1] = src[src_idx + 1]; + dst[dst_idx + 2] = src[src_idx]; +#elif defined ORDER + dst[dst_idx] = src[src_idx]; + dst[dst_idx + 1] = src[src_idx + 1]; + dst[dst_idx + 2] = src[src_idx + 2]; +#endif + +#if dcn == 4 +#if scn == 3 + dst[dst_idx + 3] = MAX_NUM; +#else + dst[dst_idx + 3] = src[src_idx + 3]; +#endif +#endif + } +} + +///////////////////////////////////// RGB5x5 <-> RGB ////////////////////////////////////// + +__kernel void RGB5x52RGB(int cols, int rows, int src_step, int dst_step, int bidx, + __global const ushort * 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) + { + int src_idx = mad24(y, src_step, src_offset + x); + int dst_idx = mad24(y, dst_step, dst_offset + (x << 2)); + ushort t = src[src_idx]; + +#if greenbits == 6 + dst[dst_idx + bidx] = (uchar)(t << 3); + dst[dst_idx + 1] = (uchar)((t >> 3) & ~3); + dst[dst_idx + (bidx^2)] = (uchar)((t >> 8) & ~7); +#else + dst[dst_idx + bidx] = (uchar)(t << 3); + dst[dst_idx + 1] = (uchar)((t >> 2) & ~7); + dst[dst_idx + (bidx^2)] = (uchar)((t >> 7) & ~7); +#endif + +#if dcn == 4 +#if greenbits == 6 + dst[dst_idx + 3] = 255; +#else + dst[dst_idx + 3] = t & 0x8000 ? 255 : 0; +#endif +#endif + } +} + +__kernel void RGB2RGB5x5(int cols, int rows, int src_step, int dst_step, int bidx, + __global const uchar * src, __global ushort * 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); + +#if greenbits == 6 + dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~3) << 3)|((src[src_idx + (bidx^2)]&~7) << 8)); +#elif scn == 3 + dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|((src[src_idx + (bidx^2)]&~7) << 7)); +#else + dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)| + ((src[src_idx + (bidx^2)]&~7) << 7)|(src[src_idx + 3] ? 0x8000 : 0)); +#endif + } +} + +///////////////////////////////////// RGB5x5 <-> RGB ////////////////////////////////////// + +__kernel void BGR5x52Gray(int cols, int rows, int src_step, int dst_step, int bidx, + __global const ushort * 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) + { + int src_idx = mad24(y, src_step, src_offset + x); + int dst_idx = mad24(y, dst_step, dst_offset + x); + int t = src[src_idx]; + +#if greenbits == 6 + dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + + ((t >> 3) & 0xfc)*G2Y + + ((t >> 8) & 0xf8)*R2Y, yuv_shift); +#else + dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + + ((t >> 2) & 0xf8)*G2Y + + ((t >> 7) & 0xf8)*R2Y, yuv_shift); +#endif + } +} + +__kernel void Gray2BGR5x5(int cols, int rows, int src_step, int dst_step, int bidx, + __global const uchar * src, __global ushort * 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); + int t = src[src_idx]; + +#if greenbits == 6 + dst[dst_idx] = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8)); +#else + t >>= 3; + dst[dst_idx] = (ushort)(t|(t << 5)|(t << 10)); +#endif + } +} + +///////////////////////////////////// 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, + __global const uchar * src, __global uchar * dst, + int src_offset, int dst_offset, + __constant int * sdiv_table, __constant int * hdiv_table) +{ + 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); + + int b = src[src_idx + bidx], g = src[src_idx + 1], r = src[src_idx + (bidx^2)]; + int h, s, v = b; + int vmin = b, diff; + int vr, vg; + + v = max( v, g ); + v = max( v, r ); + vmin = min( vmin, g ); + vmin = min( vmin, r ); + + diff = v - vmin; + vr = v == r ? -1 : 0; + vg = v == g ? -1 : 0; + + s = (diff * sdiv_table[v] + (1 << (hsv_shift-1))) >> hsv_shift; + h = (vr & (g - b)) + + (~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff)))); + h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift; + h += h < 0 ? hrange : 0; + + dst[dst_idx] = convert_uchar_sat_rte(h); + dst[dst_idx + 1] = (uchar)s; + dst[dst_idx + 2] = (uchar)v; + } +} + +__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, + __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 b = src[src_idx + bidx], g = src[src_idx + 1], r = src[src_idx + (bidx^2)]; + float h, s, v; + + float vmin, diff; + + v = vmin = r; + if( v < g ) v = g; + if( v < b ) v = b; + if( vmin > g ) vmin = g; + if( vmin > b ) vmin = b; + + diff = v - vmin; + s = diff/(float)(fabs(v) + FLT_EPSILON); + diff = (float)(60./(diff + FLT_EPSILON)); + if( v == r ) + h = (g - b)*diff; + else if( v == g ) + h = (b - r)*diff + 120.f; + else + h = (r - g)*diff + 240.f; + + if( h < 0 ) h += 360.f; + + dst[dst_idx] = h*hscale; + dst[dst_idx + 1] = s; + dst[dst_idx + 2] = v; + } +} + +__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 ////////////////////////////////////// #ifdef DEPTH_0