diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 49fabbb96e..ab6806f8e2 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 37dab79a8b..e8711cbfce 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 541c5f698f..dc81b31f83 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