diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index 9a30962061..971e4de8a9 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 3133a069f1..48f44a600d 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -2205,9 +2205,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); diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 0cbfb381b3..13f474d80b 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -2703,18 +2703,63 @@ 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_BGR2GRAY: - case COLOR_BGRA2GRAY: - case COLOR_RGB2GRAY: - case COLOR_RGBA2GRAY: + 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_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=%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_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: { CV_Assert(scn == 3 || scn == 4); bidx = code == COLOR_BGR2GRAY || code == COLOR_BGRA2GRAY ? 0 : 2; @@ -2752,10 +2797,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,18 +2822,202 @@ 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_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::PtrOnly(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::PtrOnly(c)); + return k.run(2, globalsize, 0, false); + } + 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; - /* - 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: - 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_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 06ff7d4e58..ca696290d9 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 @@ -78,6 +81,7 @@ enum { yuv_shift = 14, xyz_shift = 12, + hsv_shift = 12, R2Y = 4899, G2Y = 9617, B2Y = 1868, @@ -87,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, @@ -268,7 +280,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 +316,727 @@ __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 + } +} + +///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// + +__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 dx = get_global_id(0); + int dy = get_global_id(1); + + if (dy < rows && dx < cols) + { + int src_idx = mad24(dy, src_step, src_offset + dx * scnbytes); + int dst_idx = mad24(dy, dst_step, dst_offset + dx * dcnbytes); + + __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]; + 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[0] = SAT_CAST(x); + dst[1] = SAT_CAST(y); + dst[2] = SAT_CAST(z); + } +} + +__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 dx = get_global_id(0); + int dy = get_global_id(1); + + if (dy < rows && dx < cols) + { + int src_idx = mad24(dy, src_step, src_offset + dx * scnbytes); + int dst_idx = mad24(dy, dst_step, dst_offset + dx * dcnbytes); + + __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]; + 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[0] = SAT_CAST(b); + dst[1] = SAT_CAST(g); + dst[2] = SAT_CAST(r); +#if dcn == 4 + dst[3] = MAX_NUM; +#endif + } +} + +///////////////////////////////////// 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 + } +} + +///////////////////////////////////// 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(__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); + +#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 + } +} + +///////////////////////////////////// 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 + } +} + +//////////////////////////////////// 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 + +///////////////////////////////////// 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 + +/////////////////////////// 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 cef0c92396..22b6b5cdad 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 @@ -150,105 +150,105 @@ 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 -//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 -//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)); } -//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 -//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 -//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)); } -//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 -//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 -//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 @@ -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(