diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index ab6806f8e28377a33898e22ed5cdc7f8c886844f..c375b05332adbe0a50eb240e2144d5a7f3bb62d5 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 e8711cbfce350a431f14499822ea86a3a1b58523..e6ceada8e2b03eed9ce84a4e8df3ecb4d8d93692 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 dc81b31f83722763990062b533fb5764de7c39cf..d65c6f24012848b35581e38773016c8b14206bec 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 0e058a791ff077f7c730555f7ce20fa0a6baaa8e..5786505bdad3ca8e66d261c28017c865f170ac46 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 //////////////////////////////////////