From ab9b883c69bf94428bce17aaf2a49a055fc244c9 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 27 Nov 2013 19:37:27 +0400 Subject: [PATCH] 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 2711855cb0..49fabbb96e 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 0822b56076..37dab79a8b 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 93ede912b2..541c5f698f 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 8078eda413..0e058a791f 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,