From eba6754b061efdcf9941fb2aed88cffa556e1f76 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 8 Nov 2013 19:07:06 +0400 Subject: [PATCH 1/8] fixed ocl::cvtColor for CV_YUV2BGRA and CV_YUV2RGBA --- modules/ocl/src/color.cpp | 94 ++++++++++----------- modules/ocl/src/opencl/cvt_color.cl | 123 +++++++++++++--------------- modules/ocl/test/test_color.cpp | 85 +++++++++++++------ 3 files changed, 162 insertions(+), 140 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 27c2fd5f05..766e992dd1 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -50,52 +50,41 @@ using namespace cv; using namespace cv::ocl; -#ifndef CV_DESCALE -#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n)) -#endif - -#ifndef FLT_EPSILON -#define FLT_EPSILON 1.192092896e-07F -#endif - -namespace +static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx) { - -void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx) -{ - int channels = src.oclchannels(); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); std::string build_options = format("-D DEPTH_%d", src.depth()); + vector > args; + args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); + + size_t gt[3] = { src.cols, src.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2Gray", gt, lt, args, -1, -1, build_options.c_str()); +} + +static void Gray2RGB_caller(const oclMat &src, oclMat &dst) +{ + int channels = dst.channels(); + std::string build_options = format("-D DEPTH_%d", src.depth()); + int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); + int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); + vector > args; args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2Gray", gt, lt, args, -1, -1, build_options.c_str()); -} - -void Gray2RGB_caller(const oclMat &src, oclMat &dst) -{ - std::string build_options = format("-D DEPTH_%d", src.depth()); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); @@ -105,9 +94,8 @@ void Gray2RGB_caller(const oclMat &src, oclMat &dst) openCLExecuteKernel(src.clCxt, &cvt_color, "Gray2RGB", gt, lt, args, -1, -1, build_options.c_str()); } -void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx) +static void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx) { - int channels = src.oclchannels(); std::string build_options = format("-D DEPTH_%d", src.depth()); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -117,7 +105,6 @@ void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx) args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); @@ -128,9 +115,9 @@ void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx) openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YUV", gt, lt, args, -1, -1, build_options.c_str()); } -void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx) +static void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx) { - int channels = src.oclchannels(); + int channels = dst.channels(); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -152,7 +139,7 @@ void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx) openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGB", gt, lt, args, -1, -1, buildOptions.c_str()); } -void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx) +static void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx) { std::string build_options = format("-D DEPTH_%d", src.depth()); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); @@ -175,9 +162,8 @@ void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx) openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGBA_NV12", gt, lt, args, -1, -1, build_options.c_str()); } -void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx) +static void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx) { - int channels = src.oclchannels(); std::string build_options = format("-D DEPTH_%d", src.depth()); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -187,7 +173,6 @@ void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx) args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); @@ -198,10 +183,10 @@ void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx) openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YCrCb", gt, lt, args, -1, -1, build_options.c_str()); } -void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) +static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) { Size sz = src.size(); - int scn = src.oclchannels(), depth = src.depth(), bidx; + int scn = src.channels(), depth = src.depth(), bidx; CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32F); @@ -239,7 +224,7 @@ void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) case CV_RGB2YUV: { CV_Assert(scn == 3 || scn == 4); - bidx = code == CV_RGB2YUV ? 0 : 2; + bidx = code == CV_BGR2YUV ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, 3)); RGB2YUV_caller(src, dst, bidx); break; @@ -247,9 +232,11 @@ void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) case CV_YUV2BGR: case CV_YUV2RGB: { - CV_Assert(scn == 3 || scn == 4); - bidx = code == CV_YUV2RGB ? 0 : 2; - dst.create(sz, CV_MAKETYPE(depth, 3)); + if( dcn <= 0 ) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); + bidx = code == CV_YUV2BGR ? 0 : 2; + dst.create(sz, CV_MAKETYPE(depth, dcn)); YUV2RGB_caller(src, dst, bidx); break; } @@ -260,7 +247,7 @@ void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) { CV_Assert(scn == 1); CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); - dcn = code == CV_YUV2BGRA_NV12 || code == CV_YUV2RGBA_NV12 ? 4 : 3; + dcn = code == CV_YUV2BGRA_NV12 || code == CV_YUV2RGBA_NV12 ? 4 : 3; bidx = code == CV_YUV2BGRA_NV12 || code == CV_YUV2BGR_NV12 ? 0 : 2; Size dstSz(sz.width, sz.height * 2 / 3); @@ -280,6 +267,12 @@ void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) case CV_YCrCb2BGR: case CV_YCrCb2RGB: { + if( dcn <= 0 ) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); + bidx = code == CV_YCrCb2RGB ? 0 : 2; + dst.create(sz, CV_MAKETYPE(depth, dcn)); +// YUV2RGB_caller(src, dst, bidx); break; } /* @@ -297,7 +290,6 @@ void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" ); } } -} void cv::ocl::cvtColor(const oclMat &src, oclMat &dst, int code, int dcn) { diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 01286f7ad7..2ba7739c22 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -46,22 +46,18 @@ /**************************************PUBLICFUNC*************************************/ -#if defined (DOUBLE_SUPPORT) -#pragma OPENCL EXTENSION cl_khr_fp64:enable -#endif - #if defined (DEPTH_0) #define DATA_TYPE uchar #define MAX_NUM 255 #define HALF_MAX 128 -#define SAT_CAST(num) convert_uchar_sat(num) +#define SAT_CAST(num) convert_uchar_sat_rte(num) #endif #if defined (DEPTH_2) #define DATA_TYPE ushort #define MAX_NUM 65535 #define HALF_MAX 32768 -#define SAT_CAST(num) convert_ushort_sat(num) +#define SAT_CAST(num) convert_ushort_sat_rte(num) #endif #if defined (DEPTH_5) @@ -71,11 +67,7 @@ #define SAT_CAST(num) (num) #endif -#ifndef DATA_TYPE - #define DATA_TYPE UNDEFINED -#endif - -#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n)) +#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n)) enum { @@ -89,16 +81,16 @@ enum ///////////////////////////////////// RGB <-> GRAY ////////////////////////////////////// -__kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step, int channels, +__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) { - const int x = get_global_id(0); - const int y = get_global_id(1); + 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 * channels); + int src_idx = mad24(y, src_step, src_offset + (x << 2)); int dst_idx = mad24(y, dst_step, dst_offset + x); #if defined (DEPTH_5) dst[dst_idx] = src[src_idx + bidx] * 0.114f + src[src_idx + 1] * 0.587f + src[src_idx + (bidx^2)] * 0.299f; @@ -108,22 +100,24 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step, int chann } } -__kernel void Gray2RGB(int cols,int rows,int src_step,int dst_step, +__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int channels, __global const DATA_TYPE* src, __global DATA_TYPE* dst, int src_offset, int dst_offset) { - const int x = get_global_id(0); - const int y = get_global_id(1); + 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 * 4); + int dst_idx = mad24(y, dst_step, dst_offset + (x << 2)); + DATA_TYPE val = src[src_idx]; - dst[dst_idx++] = val; - dst[dst_idx++] = val; - dst[dst_idx++] = val; - dst[dst_idx] = MAX_NUM; + dst[dst_idx] = val; + dst[dst_idx + 1] = val; + dst[dst_idx + 2] = val; + if (channels == 4) + dst[dst_idx + 3] = MAX_NUM; } } @@ -132,7 +126,7 @@ __kernel void Gray2RGB(int cols,int rows,int src_step,int dst_step, __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 channels, +__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) { @@ -141,35 +135,34 @@ __kernel void RGB2YUV(int cols,int rows,int src_step,int dst_step,int channels, if (y < rows && x < cols) { - x *= channels; + x <<= 2; int src_idx = mad24(y, src_step, src_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x); - dst += dst_idx; - const DATA_TYPE rgb[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]}; + DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; #if defined (DEPTH_5) __constant float * coeffs = c_RGB2YUVCoeffs_f; - const DATA_TYPE Y = rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2]; - const DATA_TYPE Cr = (rgb[bidx] - Y) * coeffs[3] + HALF_MAX; - const DATA_TYPE Cb = (rgb[bidx^2] - Y) * coeffs[4] + HALF_MAX; + 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; - const int delta = HALF_MAX * (1 << yuv_shift); - const int Y = CV_DESCALE(rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2], yuv_shift); - const int Cr = CV_DESCALE((rgb[bidx] - Y) * coeffs[3] + delta, yuv_shift); - const int Cb = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[4] + delta, yuv_shift); + 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[0] = SAT_CAST( Y ); - dst[1] = SAT_CAST( Cr ); - dst[2] = SAT_CAST( Cb ); + 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 channels, +__kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, int channels, int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst, int src_offset, int dst_offset) { @@ -178,27 +171,28 @@ __kernel void YUV2RGB(int cols,int rows,int src_step,int dst_step,int channels, if (y < rows && x < cols) { - x *= channels; + x <<= 2; int src_idx = mad24(y, src_step, src_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x); - dst += dst_idx; - const DATA_TYPE yuv[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]}; + DATA_TYPE yuv[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; #if defined (DEPTH_5) __constant float * coeffs = c_YUV2RGBCoeffs_f; - const float b = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[3]; - const float g = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1]; - const float r = yuv[0] + (yuv[1] - HALF_MAX) * coeffs[0]; + 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; - const int b = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[3], yuv_shift); - const int g = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1], yuv_shift); - const int r = yuv[0] + CV_DESCALE((yuv[1] - HALF_MAX) * coeffs[0], yuv_shift); + 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[bidx^2] = SAT_CAST( b ); - dst[1] = SAT_CAST( g ); - dst[bidx] = SAT_CAST( r ); + dst[dst_idx + bidx] = SAT_CAST( b ); + dst[dst_idx + 1] = SAT_CAST( g ); + dst[dst_idx + (bidx^2)] = SAT_CAST( r ); + if (channels == 4) + dst[dst_idx + 3] = MAX_NUM; } } @@ -261,12 +255,12 @@ __kernel void YUV2RGBA_NV12(int cols,int rows,int src_step,int dst_step, } } -///////////////////////////////////// RGB <-> YUV ////////////////////////////////////// +///////////////////////////////////// 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 channels, +__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) { @@ -275,28 +269,27 @@ __kernel void RGB2YCrCb(int cols,int rows,int src_step,int dst_step,int channels if (y < rows && x < cols) { - x *= channels; + x <<= 2; int src_idx = mad24(y, src_step, src_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x); - dst += dst_idx; - const DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; + DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; #if defined (DEPTH_5) __constant float * coeffs = c_RGB2YCrCbCoeffs_f; - const DATA_TYPE Y = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx]; - const DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX; - const DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX; + 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; - const int delta = HALF_MAX * (1 << yuv_shift); - const int Y = CV_DESCALE(rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx], yuv_shift); - const int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift); - const int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift); + 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[0] = SAT_CAST( Y ); - dst[1] = SAT_CAST( Cr ); - dst[2] = SAT_CAST( Cb ); + dst[dst_idx] = SAT_CAST( Y ); + dst[dst_idx + 1] = SAT_CAST( Cr ); + dst[dst_idx + 2] = SAT_CAST( Cb ); } } diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 4b6a6e41b9..f5f9f43177 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -46,8 +46,6 @@ #include "test_precomp.hpp" #ifdef HAVE_OPENCL -namespace -{ using namespace testing; /////////////////////////////////////////////////////////////////////////////////////////////////////// @@ -108,8 +106,8 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) { random_roi(channelsIn, channelsOut); - cv::cvtColor(src1_roi, dst1_roi, code); - cv::ocl::cvtColor(gsrc1_roi, gdst1_roi, code); + cv::cvtColor(src1_roi, dst1_roi, code, channelsOut); + cv::ocl::cvtColor(gsrc1_roi, gdst1_roi, code, channelsOut); Near(); } @@ -125,7 +123,7 @@ OCL_TEST_P(CvtColor, RGB2GRAY) OCL_TEST_P(CvtColor, GRAY2RGB) { doTest(1, 3, CVTCODE(GRAY2RGB)); -}; +} OCL_TEST_P(CvtColor, BGR2GRAY) { @@ -134,25 +132,26 @@ OCL_TEST_P(CvtColor, BGR2GRAY) OCL_TEST_P(CvtColor, GRAY2BGR) { doTest(1, 3, CVTCODE(GRAY2BGR)); -}; +} OCL_TEST_P(CvtColor, RGBA2GRAY) { - doTest(3, 1, CVTCODE(RGBA2GRAY)); + doTest(4, 1, CVTCODE(RGBA2GRAY)); } OCL_TEST_P(CvtColor, GRAY2RGBA) { - doTest(1, 3, CVTCODE(GRAY2RGBA)); -}; + doTest(1, 4, CVTCODE(GRAY2RGBA)); +} OCL_TEST_P(CvtColor, BGRA2GRAY) { - doTest(3, 1, CVTCODE(BGRA2GRAY)); + doTest(4, 1, CVTCODE(BGRA2GRAY)); } OCL_TEST_P(CvtColor, GRAY2BGRA) { - doTest(1, 3, CVTCODE(GRAY2BGRA)); -}; + doTest(1, 4, CVTCODE(GRAY2BGRA)); +} + OCL_TEST_P(CvtColor, RGB2YUV) { @@ -162,6 +161,14 @@ OCL_TEST_P(CvtColor, BGR2YUV) { doTest(3, 3, CVTCODE(BGR2YUV)); } +OCL_TEST_P(CvtColor, RGBA2YUV) +{ + doTest(4, 3, CVTCODE(RGB2YUV)); +} +OCL_TEST_P(CvtColor, BGRA2YUV) +{ + doTest(4, 3, CVTCODE(BGR2YUV)); +} OCL_TEST_P(CvtColor, YUV2RGB) { doTest(3, 3, CVTCODE(YUV2RGB)); @@ -170,6 +177,16 @@ OCL_TEST_P(CvtColor, YUV2BGR) { doTest(3, 3, CVTCODE(YUV2BGR)); } +OCL_TEST_P(CvtColor, YUV2RGBA) +{ + doTest(3, 4, CVTCODE(YUV2RGB)); +} +OCL_TEST_P(CvtColor, YUV2BGRA) +{ + doTest(3, 4, CVTCODE(YUV2BGR)); +} + + OCL_TEST_P(CvtColor, RGB2YCrCb) { doTest(3, 3, CVTCODE(RGB2YCrCb)); @@ -178,8 +195,33 @@ OCL_TEST_P(CvtColor, BGR2YCrCb) { doTest(3, 3, CVTCODE(BGR2YCrCb)); } +OCL_TEST_P(CvtColor, RGBA2YCrCb) +{ + doTest(4, 3, CVTCODE(RGB2YCrCb)); +} +OCL_TEST_P(CvtColor, BGRA2YCrCb) +{ + doTest(4, 3, CVTCODE(BGR2YCrCb)); +} +//OCL_TEST_P(CvtColor, YCrCb2RGB) +//{ +// doTest(3, 3, CVTCODE(YCrCb2RGB)); +//} +//OCL_TEST_P(CvtColor, YCrCb2BGR) +//{ +// doTest(3, 3, CVTCODE(YCrCb2BGR)); +//} +//OCL_TEST_P(CvtColor, YCrCb2RGBA) +//{ +// doTest(3, 4, CVTCODE(YCrCb2RGB)); +//} +//OCL_TEST_P(CvtColor, YCrCb2BGRA) +//{ +// doTest(3, 4, CVTCODE(YCrCb2BGR)); +//} -struct CvtColor_YUV420 : CvtColor +struct CvtColor_YUV420 : + public CvtColor { void random_roi(int channelsIn, int channelsOut) { @@ -203,37 +245,32 @@ struct CvtColor_YUV420 : CvtColor OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { doTest(1, 4, CV_YUV2RGBA_NV12); -}; +} OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { doTest(1, 4, CV_YUV2BGRA_NV12); -}; +} OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { doTest(1, 3, CV_YUV2RGB_NV12); -}; +} OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { doTest(1, 3, CV_YUV2BGR_NV12); -}; +} INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine( testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)), - Bool() - ) - ); + Bool())); INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_YUV420, testing::Combine( testing::Values(MatDepth(CV_8U)), - Bool() - ) - ); + Bool())); -} #endif From a57030a0cda4a144c1b800ce5e7efb6c127aa571 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sat, 9 Nov 2013 17:03:30 +0400 Subject: [PATCH 2/8] added YCrCb to RGB, BGR, RGBA, BGRA modes to ocl::cvtColor --- modules/ocl/src/color.cpp | 28 +++++++++++++++++++-- modules/ocl/src/opencl/cvt_color.cl | 38 +++++++++++++++++++++++++++++ modules/ocl/test/test_color.cpp | 32 ++++++++++++------------ 3 files changed, 80 insertions(+), 18 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 766e992dd1..b25b71feee 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -162,6 +162,30 @@ static void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx) openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGBA_NV12", gt, lt, args, -1, -1, build_options.c_str()); } +static void YCrCb2RGB_caller(const oclMat &src, oclMat &dst, int bidx) +{ + int channels = dst.channels(); + int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); + int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); + + std::string buildOptions = format("-D DEPTH_%d", src.depth()); + + vector > args; + args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); + + size_t gt[3] = { src.cols, src.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, "YCrCb2RGB", gt, lt, args, -1, -1, buildOptions.c_str()); +} + static void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx) { std::string build_options = format("-D DEPTH_%d", src.depth()); @@ -270,9 +294,9 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) if( dcn <= 0 ) dcn = 3; CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); - bidx = code == CV_YCrCb2RGB ? 0 : 2; + bidx = code == CV_YCrCb2BGR ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, dcn)); -// YUV2RGB_caller(src, dst, bidx); + YCrCb2RGB_caller(src, dst, bidx); break; } /* diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 2ba7739c22..de53da52ee 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -293,3 +293,41 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step, 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 channels, + 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 (channels == 4) + dst[dst_idx + 3] = MAX_NUM; + } +} diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index f5f9f43177..b2f5c6fb2d 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -203,22 +203,22 @@ OCL_TEST_P(CvtColor, BGRA2YCrCb) { doTest(4, 3, CVTCODE(BGR2YCrCb)); } -//OCL_TEST_P(CvtColor, YCrCb2RGB) -//{ -// doTest(3, 3, CVTCODE(YCrCb2RGB)); -//} -//OCL_TEST_P(CvtColor, YCrCb2BGR) -//{ -// doTest(3, 3, CVTCODE(YCrCb2BGR)); -//} -//OCL_TEST_P(CvtColor, YCrCb2RGBA) -//{ -// doTest(3, 4, CVTCODE(YCrCb2RGB)); -//} -//OCL_TEST_P(CvtColor, YCrCb2BGRA) -//{ -// doTest(3, 4, CVTCODE(YCrCb2BGR)); -//} +OCL_TEST_P(CvtColor, YCrCb2RGB) +{ + doTest(3, 3, CVTCODE(YCrCb2RGB)); +} +OCL_TEST_P(CvtColor, YCrCb2BGR) +{ + doTest(3, 3, CVTCODE(YCrCb2BGR)); +} +OCL_TEST_P(CvtColor, YCrCb2RGBA) +{ + doTest(3, 4, CVTCODE(YCrCb2RGB)); +} +OCL_TEST_P(CvtColor, YCrCb2BGRA) +{ + doTest(3, 4, CVTCODE(YCrCb2BGR)); +} struct CvtColor_YUV420 : public CvtColor From 33ae64201c024ed3af1cc9eb38e64e58842c3d87 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sat, 9 Nov 2013 19:14:38 +0400 Subject: [PATCH 3/8] color.cpp refactoring: created generic interface for toRGB and fromRGB callers --- modules/ocl/src/color.cpp | 150 ++++------------------------ modules/ocl/src/opencl/cvt_color.cl | 14 +-- modules/ocl/test/test_color.cpp | 7 +- 3 files changed, 32 insertions(+), 139 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index b25b71feee..aed9a7f0cf 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -50,7 +50,7 @@ using namespace cv; using namespace cv::ocl; -static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx) +static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName) { int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -58,121 +58,30 @@ static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx) std::string build_options = format("-D DEPTH_%d", src.depth()); vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = { src.cols, src.rows, 1 }, lt[3] = { 16, 16, 1 }; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2Gray", gt, lt, args, -1, -1, build_options.c_str()); -} - -static void Gray2RGB_caller(const oclMat &src, oclMat &dst) -{ - int channels = dst.channels(); - std::string build_options = format("-D DEPTH_%d", src.depth()); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "Gray2RGB", gt, lt, args, -1, -1, build_options.c_str()); -} - -static void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx) -{ - std::string build_options = format("-D DEPTH_%d", src.depth()); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YUV", gt, lt, args, -1, -1, build_options.c_str()); -} - -static void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx) -{ - int channels = dst.channels(); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - std::string buildOptions = format("-D DEPTH_%d", src.depth()); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGB", gt, lt, args, -1, -1, buildOptions.c_str()); -} - -static void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx) -{ - std::string build_options = format("-D DEPTH_%d", src.depth()); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - size_t gt[3] = {dst.cols / 2, dst.rows / 2, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "YUV2RGBA_NV12", gt, lt, args, -1, -1, build_options.c_str()); + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } -static void YCrCb2RGB_caller(const oclMat &src, oclMat &dst, int bidx) +static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName) { int channels = dst.channels(); + std::string build_options = format("-D DEPTH_%d", src.depth()); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - std::string buildOptions = format("-D DEPTH_%d", src.depth()); - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); @@ -182,29 +91,8 @@ static void YCrCb2RGB_caller(const oclMat &src, oclMat &dst, int bidx) args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - size_t gt[3] = { src.cols, src.rows, 1 }, lt[3] = { 16, 16, 1 }; - openCLExecuteKernel(src.clCxt, &cvt_color, "YCrCb2RGB", gt, lt, args, -1, -1, buildOptions.c_str()); -} - -static void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx) -{ - std::string build_options = format("-D DEPTH_%d", src.depth()); - int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); - int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); - - vector > args; - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); - args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); - - size_t gt[3] = {src.cols, src.rows, 1}, lt[3] = {16, 16, 1}; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2YCrCb", gt, lt, args, -1, -1, build_options.c_str()); + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) @@ -232,7 +120,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, 1)); - RGB2Gray_caller(src, dst, bidx); + fromRGB_caller(src, dst, bidx, "RGB2Gray"); break; } case CV_GRAY2BGR: @@ -241,7 +129,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 1); dcn = code == CV_GRAY2BGRA ? 4 : 3; dst.create(sz, CV_MAKETYPE(depth, dcn)); - Gray2RGB_caller(src, dst); + toRGB_caller(src, dst, 0, "Gray2RGB"); break; } case CV_BGR2YUV: @@ -250,7 +138,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2YUV ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, 3)); - RGB2YUV_caller(src, dst, bidx); + fromRGB_caller(src, dst, bidx, "RGB2YUV"); break; } case CV_YUV2BGR: @@ -261,7 +149,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); bidx = code == CV_YUV2BGR ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, dcn)); - YUV2RGB_caller(src, dst, bidx); + toRGB_caller(src, dst, bidx, "YUV2RGB"); break; } case CV_YUV2RGB_NV12: @@ -276,7 +164,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) Size dstSz(sz.width, sz.height * 2 / 3); dst.create(dstSz, CV_MAKETYPE(depth, dcn)); - YUV2RGB_NV12_caller(src, dst, bidx); + toRGB_caller(src, dst, bidx, "YUV2RGBA_NV12"); break; } case CV_BGR2YCrCb: @@ -285,7 +173,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2YCrCb ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, 3)); - RGB2YCrCb_caller(src, dst, bidx); + fromRGB_caller(src, dst, bidx, "RGB2YCrCb"); break; } case CV_YCrCb2BGR: @@ -296,7 +184,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); bidx = code == CV_YCrCb2BGR ? 0 : 2; dst.create(sz, CV_MAKETYPE(depth, dcn)); - YCrCb2RGB_caller(src, dst, bidx); + toRGB_caller(src, dst, bidx, "YCrCb2RGB"); break; } /* diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index de53da52ee..d9426b28b1 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -100,7 +100,7 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step, } } -__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int channels, +__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int channels, int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst, int src_offset, int dst_offset) { @@ -203,17 +203,17 @@ __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, int width, int height, __global const uchar* src, __global uchar* dst, +__kernel void YUV2RGBA_NV12(int cols, int rows, int src_step, int dst_step, int channels, + int bidx, __global const uchar* src, __global uchar* dst, int src_offset, int dst_offset) { - const int x = get_global_id(0); // max_x = width / 2 - const int y = get_global_id(1); // max_y = height/ 2 + const int x = get_global_id(0); + const int y = get_global_id(1); - if (y < height / 2 && x < width / 2 ) + 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(height + y, 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); diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index b2f5c6fb2d..df52b94b27 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -116,6 +116,8 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) #define CVTCODE(name) cv::COLOR_ ## name +// RGB <-> Gray + OCL_TEST_P(CvtColor, RGB2GRAY) { doTest(3, 1, CVTCODE(RGB2GRAY)); @@ -152,6 +154,7 @@ OCL_TEST_P(CvtColor, GRAY2BGRA) doTest(1, 4, CVTCODE(GRAY2BGRA)); } +// RGB <-> YUV OCL_TEST_P(CvtColor, RGB2YUV) { @@ -186,6 +189,7 @@ OCL_TEST_P(CvtColor, YUV2BGRA) doTest(3, 4, CVTCODE(YUV2BGR)); } +// RGB <-> YCrCb OCL_TEST_P(CvtColor, RGB2YCrCb) { @@ -220,6 +224,8 @@ OCL_TEST_P(CvtColor, YCrCb2BGRA) doTest(3, 4, CVTCODE(YCrCb2BGR)); } +// YUV -> RGBA_NV12 + struct CvtColor_YUV420 : public CvtColor { @@ -262,7 +268,6 @@ OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) doTest(1, 3, CV_YUV2BGR_NV12); } - INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine( testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)), From 5e02b20482fff0954344f80d3603733c75a7ff72 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sun, 10 Nov 2013 13:38:09 +0400 Subject: [PATCH 4/8] added RGB -> XYZ conversion to ocl::cvtColor --- modules/ocl/src/color.cpp | 71 ++++++++++++++++++++++++++--- modules/ocl/src/opencl/cvt_color.cl | 37 +++++++++++++++ modules/ocl/test/test_color.cpp | 70 ++++++++++++++++++---------- 3 files changed, 149 insertions(+), 29 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index aed9a7f0cf..02e7bf3c8a 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -50,7 +50,8 @@ using namespace cv; using namespace cv::ocl; -static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName) +static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName, + const oclMat & data = oclMat()) { int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -68,6 +69,9 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std:: args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); + if (!data.empty()) + args.push_back( make_pair( sizeof(cl_mem) , (void *)&data.data )); + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } @@ -112,10 +116,10 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: */ - case CV_BGR2GRAY: - case CV_BGRA2GRAY: case CV_RGB2GRAY: + case CV_BGR2GRAY: case CV_RGBA2GRAY: + case CV_BGRA2GRAY: { CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2; @@ -190,9 +194,64 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) /* case CV_BGR5652GRAY: case CV_BGR5552GRAY: case CV_GRAY2BGR565: case CV_GRAY2BGR555: - case CV_BGR2YCrCb: case CV_RGB2YCrCb: - case CV_BGR2XYZ: case CV_RGB2XYZ: - case CV_XYZ2BGR: case CV_XYZ2RGB: + */ + case CV_BGR2XYZ: + case CV_RGB2XYZ: + { + CV_Assert(scn == 3 || scn == 4); + bidx = code == CV_BGR2XYZ ? 0 : 2; + dst.create(sz, CV_MAKE_TYPE(depth, 3)); + + void * pdata = NULL; + 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]); + } + pdata = coeffs; + } + 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]); + } + pdata = coeffs; + } + oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32F : CV_32S, pdata); + + fromRGB_caller(src, dst, bidx, "RGB2XYZ", oclCoeffs); + break; + } + case CV_XYZ2BGR: + case CV_XYZ2RGB: + { + if (dcn <= 0) + dcn = 3; + CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); + bidx = code == CV_XYZ2BGR ? 0 : 2; + dst.create(sz, CV_MAKE_TYPE(depth, dcn)); + toRGB_caller(src, dst, bidx, "XYZ2RGB"); + break; + } + /* case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL: case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL: case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL: diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index d9426b28b1..775628bab1 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -48,6 +48,7 @@ #if defined (DEPTH_0) #define DATA_TYPE uchar +#define COEFF_TYPE int #define MAX_NUM 255 #define HALF_MAX 128 #define SAT_CAST(num) convert_uchar_sat_rte(num) @@ -55,6 +56,7 @@ #if defined (DEPTH_2) #define DATA_TYPE ushort +#define COEFF_TYPE int #define MAX_NUM 65535 #define HALF_MAX 32768 #define SAT_CAST(num) convert_ushort_sat_rte(num) @@ -62,6 +64,7 @@ #if defined (DEPTH_5) #define DATA_TYPE float +#define COEFF_TYPE float #define MAX_NUM 1.0f #define HALF_MAX 0.5f #define SAT_CAST(num) (num) @@ -331,3 +334,37 @@ __kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step, int chan dst[dst_idx + 3] = MAX_NUM; } } + +///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// + +#pragma OPENCL EXTENSION cl_amd_printf:enable + +__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, __global COEFF_TYPE * coeffs) +{ + 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 r = src[src_idx], g = src[src_idx + 1], b = src[src_idx + 2]; + +#ifdef DEPTH_5 + DATA_TYPE x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2]; + DATA_TYPE y = r * coeffs[3] + g * coeffs[4] + b * coeffs[5]; + DATA_TYPE z = r * coeffs[6] + g * coeffs[7] + b * coeffs[8]; +#else + DATA_TYPE x = CV_DESCALE(r * coeffs[0] + g * coeffs[1] + b * coeffs[2], xyz_shift); + DATA_TYPE y = CV_DESCALE(r * coeffs[3] + g * coeffs[4] + b * coeffs[5], xyz_shift); + DATA_TYPE 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); + } +} diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index df52b94b27..0d9eb8f133 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -47,6 +47,7 @@ #ifdef HAVE_OPENCL using namespace testing; +using namespace cv; /////////////////////////////////////////////////////////////////////////////////////////////////////// // cvtColor @@ -57,20 +58,20 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) bool use_roi; // src mat - cv::Mat src1; - cv::Mat dst1; + Mat src; + Mat dst; // src mat with roi - cv::Mat src1_roi; - cv::Mat dst1_roi; + Mat src_roi; + Mat dst_roi; // ocl dst mat for testing - cv::ocl::oclMat gsrc1_whole; - cv::ocl::oclMat gdst1_whole; + ocl::oclMat gsrc_whole; + ocl::oclMat gdst_whole; // ocl mat with roi - cv::ocl::oclMat gsrc1_roi; - cv::ocl::oclMat gdst1_roi; + ocl::oclMat gsrc_roi; + ocl::oclMat gdst_roi; virtual void SetUp() { @@ -85,19 +86,23 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) Size roiSize = randomSize(1, MAX_VALUE); Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); - randomSubMat(src1, src1_roi, roiSize, srcBorder, srcType, 2, 100); + randomSubMat(src, src_roi, roiSize, srcBorder, srcType, 2, 100); - Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); - randomSubMat(dst1, dst1_roi, roiSize, dst1Border, dstType, 5, 16); + Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, dstType, 5, 16); - generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, srcBorder); - generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border); + generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder); + generateOclMat(gdst_whole, gdst_roi, dst, roiSize, dstBorder); } void Near(double threshold = 1e-3) { - EXPECT_MAT_NEAR(dst1, gdst1_whole, threshold); - EXPECT_MAT_NEAR(dst1_roi, gdst1_roi, threshold); + Mat whole, roi; + gdst_whole.download(whole); + gdst_roi.download(roi); + + EXPECT_MAT_NEAR(dst, whole, threshold); + EXPECT_MAT_NEAR(dst_roi, roi, threshold); } void doTest(int channelsIn, int channelsOut, int code) @@ -106,15 +111,15 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) { random_roi(channelsIn, channelsOut); - cv::cvtColor(src1_roi, dst1_roi, code, channelsOut); - cv::ocl::cvtColor(gsrc1_roi, gdst1_roi, code, channelsOut); + cvtColor(src_roi, dst_roi, code, channelsOut); + ocl::cvtColor(gsrc_roi, gdst_roi, code, channelsOut); Near(); } } }; -#define CVTCODE(name) cv::COLOR_ ## name +#define CVTCODE(name) COLOR_ ## name // RGB <-> Gray @@ -224,6 +229,25 @@ OCL_TEST_P(CvtColor, YCrCb2BGRA) doTest(3, 4, CVTCODE(YCrCb2BGR)); } +// RGB <-> XYZ + +OCL_TEST_P(CvtColor, RGB2XYZ) +{ + doTest(3, 3, CVTCODE(RGB2XYZ)); +} +OCL_TEST_P(CvtColor, BGR2XYZ) +{ + doTest(3, 3, CVTCODE(BGR2XYZ)); +} +OCL_TEST_P(CvtColor, RGBA2XYZ) +{ + doTest(4, 3, CVTCODE(RGB2XYZ)); +} +OCL_TEST_P(CvtColor, BGRA2XYZ) +{ + doTest(4, 3, CVTCODE(BGR2XYZ)); +} + // YUV -> RGBA_NV12 struct CvtColor_YUV420 : @@ -238,13 +262,13 @@ struct CvtColor_YUV420 : roiSize.width *= 2; roiSize.height *= 3; Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); - randomSubMat(src1, src1_roi, roiSize, srcBorder, srcType, 2, 100); + randomSubMat(src, src_roi, roiSize, srcBorder, srcType, 2, 100); - Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); - randomSubMat(dst1, dst1_roi, roiSize, dst1Border, dstType, 5, 16); + Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, dstType, 5, 16); - generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, srcBorder); - generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border); + generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder); + generateOclMat(gdst_whole, gdst_roi, dst, roiSize, dstBorder); } }; From 581a3e444d4f9be398a0517dbe57ad913556a1d3 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sun, 10 Nov 2013 14:30:37 +0400 Subject: [PATCH 5/8] added XYZ to RGB conversion to ocl::cvtColor --- modules/ocl/src/color.cpp | 50 +++++++++++++-- modules/ocl/src/opencl/cvt_color.cl | 94 ++++++++++++++++++++--------- modules/ocl/test/test_color.cpp | 17 ++++++ 3 files changed, 125 insertions(+), 36 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 02e7bf3c8a..9688be08d8 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -76,10 +76,10 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std:: openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } -static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName) +static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName, + const oclMat & data = oclMat()) { - int channels = dst.channels(); - std::string build_options = format("-D DEPTH_%d", src.depth()); + std::string build_options = format("-D DEPTH_%d -D channels=%d", src.depth(), dst.channels()); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -88,13 +88,15 @@ static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::st args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); - args.push_back( make_pair( sizeof(cl_int) , (void *)&channels)); args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); + if (!data.empty()) + args.push_back( make_pair( sizeof(cl_mem) , (void *)&data.data )); + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } @@ -235,7 +237,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) } pdata = coeffs; } - oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32F : CV_32S, pdata); + oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32FC1 : CV_32SC1, pdata); fromRGB_caller(src, dst, bidx, "RGB2XYZ", oclCoeffs); break; @@ -248,7 +250,43 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) CV_Assert(scn == 3 && (dcn == 3 || dcn == 4)); bidx = code == CV_XYZ2BGR ? 0 : 2; dst.create(sz, CV_MAKE_TYPE(depth, dcn)); - toRGB_caller(src, dst, bidx, "XYZ2RGB"); + + void * pdata = NULL; + 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]); + } + pdata = coeffs; + } + 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]); + } + pdata = coeffs; + } + oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32FC1 : CV_32SC1, pdata); + + toRGB_caller(src, dst, bidx, "XYZ2RGB", oclCoeffs); break; } /* diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 775628bab1..0f44897a10 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -46,7 +46,7 @@ /**************************************PUBLICFUNC*************************************/ -#if defined (DEPTH_0) +#ifdef DEPTH_0 #define DATA_TYPE uchar #define COEFF_TYPE int #define MAX_NUM 255 @@ -54,7 +54,7 @@ #define SAT_CAST(num) convert_uchar_sat_rte(num) #endif -#if defined (DEPTH_2) +#ifdef DEPTH_2 #define DATA_TYPE ushort #define COEFF_TYPE int #define MAX_NUM 65535 @@ -62,7 +62,7 @@ #define SAT_CAST(num) convert_ushort_sat_rte(num) #endif -#if defined (DEPTH_5) +#ifdef DEPTH_5 #define DATA_TYPE float #define COEFF_TYPE float #define MAX_NUM 1.0f @@ -95,7 +95,7 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step, { int src_idx = mad24(y, src_step, src_offset + (x << 2)); int dst_idx = mad24(y, dst_step, dst_offset + x); -#if defined (DEPTH_5) +#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); @@ -103,7 +103,7 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step, } } -__kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int channels, int bidx, +__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) { @@ -119,8 +119,9 @@ __kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int chann dst[dst_idx] = val; dst[dst_idx + 1] = val; dst[dst_idx + 2] = val; - if (channels == 4) +#if channels == 4 dst[dst_idx + 3] = MAX_NUM; +#endif } } @@ -143,7 +144,7 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step, int dst_idx = mad24(y, dst_step, dst_offset + x); DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; -#if defined (DEPTH_5) +#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; @@ -165,7 +166,7 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step, __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 channels, +__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) { @@ -179,7 +180,7 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, int channe int dst_idx = mad24(y, dst_step, dst_offset + x); DATA_TYPE yuv[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; -#if defined (DEPTH_5) +#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]; @@ -194,8 +195,9 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, int channe dst[dst_idx + bidx] = SAT_CAST( b ); dst[dst_idx + 1] = SAT_CAST( g ); dst[dst_idx + (bidx^2)] = SAT_CAST( r ); - if (channels == 4) +#if channels == 4 dst[dst_idx + 3] = MAX_NUM; +#endif } } @@ -206,7 +208,7 @@ __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 channels, +__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) { @@ -278,7 +280,7 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step, DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] }; -#if defined (DEPTH_5) +#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; @@ -300,7 +302,7 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step, __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 channels, +__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) { @@ -330,41 +332,73 @@ __kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step, int chan dst[dst_idx + (bidx^2)] = SAT_CAST(r); dst[dst_idx + 1] = SAT_CAST(g); dst[dst_idx + bidx] = SAT_CAST(b); - if (channels == 4) +#if channels == 4 dst[dst_idx + 3] = MAX_NUM; +#endif } } ///////////////////////////////////// RGB <-> XYZ ////////////////////////////////////// -#pragma OPENCL EXTENSION cl_amd_printf:enable - __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, __global COEFF_TYPE * coeffs) + int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst, + int src_offset, int dst_offset, __constant COEFF_TYPE * coeffs) { - int x = get_global_id(0); - int y = get_global_id(1); + int dx = get_global_id(0); + int dy = get_global_id(1); - if (y < rows && x < cols) + if (dy < rows && dx < cols) { - x <<= 2; - int src_idx = mad24(y, src_step, src_offset + x); - int dst_idx = mad24(y, dst_step, dst_offset + x); + 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 - DATA_TYPE x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2]; - DATA_TYPE y = r * coeffs[3] + g * coeffs[4] + b * coeffs[5]; - DATA_TYPE z = r * coeffs[6] + g * coeffs[7] + b * coeffs[8]; + 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 - DATA_TYPE x = CV_DESCALE(r * coeffs[0] + g * coeffs[1] + b * coeffs[2], xyz_shift); - DATA_TYPE y = CV_DESCALE(r * coeffs[3] + g * coeffs[4] + b * coeffs[5], xyz_shift); - DATA_TYPE z = CV_DESCALE(r * coeffs[6] + g * coeffs[7] + b * coeffs[8], xyz_shift); + 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 channels == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 0d9eb8f133..977fd206da 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -248,6 +248,23 @@ OCL_TEST_P(CvtColor, BGRA2XYZ) doTest(4, 3, CVTCODE(BGR2XYZ)); } +OCL_TEST_P(CvtColor, XYZ2RGB) +{ + doTest(3, 3, CVTCODE(XYZ2RGB)); +} +OCL_TEST_P(CvtColor, XYZ2BGR) +{ + doTest(3, 3, CVTCODE(XYZ2BGR)); +} +OCL_TEST_P(CvtColor, XYZ2RGBA) +{ + doTest(3, 4, CVTCODE(XYZ2RGB)); +} +OCL_TEST_P(CvtColor, XYZ2BGRA) +{ + doTest(3, 4, CVTCODE(XYZ2BGR)); +} + // YUV -> RGBA_NV12 struct CvtColor_YUV420 : From 3cc9502c90836c3fcda86edfb29dbf301477ae6b Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sun, 10 Nov 2013 20:56:18 +0400 Subject: [PATCH 6/8] added RGB[A] <-> BGR[A] conversion to ocl::cvtColor --- modules/ocl/src/color.cpp | 59 +++++--- modules/ocl/src/opencl/cvt_color.cl | 51 +++++-- modules/ocl/test/test_color.cpp | 201 +++++++--------------------- 3 files changed, 136 insertions(+), 175 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 9688be08d8..f27366a6ae 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -79,7 +79,7 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std:: static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName, const oclMat & data = oclMat()) { - std::string build_options = format("-D DEPTH_%d -D channels=%d", src.depth(), dst.channels()); + std::string build_options = format("-D DEPTH_%d -D dcn=%d", src.depth(), dst.channels()); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -101,6 +101,27 @@ static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::st openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } +static void RGB_caller(const oclMat &src, oclMat &dst, bool reverse) +{ + std::string build_options = format("-D DEPTH_%d -D dcn=%d -D scn=%d -D %s", src.depth(), + dst.channels(), src.channels(), reverse ? "REVERSE" : "ORDER"); + int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); + int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); + + vector > args; + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); + + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, "RGB", gt, lt, args, -1, -1, build_options.c_str()); +} + static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) { Size sz = src.size(); @@ -110,18 +131,24 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) switch (code) { - /* - case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR: - case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA: - case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: - case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555: - case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: - case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: - */ - case CV_RGB2GRAY: - case CV_BGR2GRAY: - case CV_RGBA2GRAY: - case CV_BGRA2GRAY: + case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR: + case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA: + { + CV_Assert(scn == 3 || scn == 4); + dcn = code == CV_BGR2BGRA || code == CV_RGB2BGRA || code == CV_BGRA2RGBA ? 4 : 3; + bool reverse = !(code == CV_BGR2BGRA || code == CV_BGRA2BGR); + dst.create(sz, CV_MAKE_TYPE(depth, dcn)); + RGB_caller(src, dst, reverse); + break; + } + /* + case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: + case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555: + case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: + case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: + */ + case CV_RGB2GRAY: case CV_BGR2GRAY: + case CV_RGBA2GRAY: case CV_BGRA2GRAY: { CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2; @@ -158,10 +185,8 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) toRGB_caller(src, dst, bidx, "YUV2RGB"); break; } - case CV_YUV2RGB_NV12: - case CV_YUV2BGR_NV12: - case CV_YUV2RGBA_NV12: - case CV_YUV2BGRA_NV12: + case CV_YUV2RGB_NV12: case CV_YUV2BGR_NV12: + case CV_YUV2RGBA_NV12: case CV_YUV2BGRA_NV12: { CV_Assert(scn == 1); CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 0f44897a10..916d44bf93 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -119,8 +119,8 @@ __kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int bidx, dst[dst_idx] = val; dst[dst_idx + 1] = val; dst[dst_idx + 2] = val; -#if channels == 4 - dst[dst_idx + 3] = MAX_NUM; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; #endif } } @@ -195,8 +195,8 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, dst[dst_idx + bidx] = SAT_CAST( b ); dst[dst_idx + 1] = SAT_CAST( g ); dst[dst_idx + (bidx^2)] = SAT_CAST( r ); -#if channels == 4 - dst[dst_idx + 3] = MAX_NUM; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; #endif } } @@ -332,8 +332,8 @@ __kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step, dst[dst_idx + (bidx^2)] = SAT_CAST(r); dst[dst_idx + 1] = SAT_CAST(g); dst[dst_idx + bidx] = SAT_CAST(b); -#if channels == 4 - dst[dst_idx + 3] = MAX_NUM; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; #endif } } @@ -397,8 +397,43 @@ __kernel void XYZ2RGB(int cols, int rows, int src_step, int dst_step, dst[dst_idx] = SAT_CAST(b); dst[dst_idx + 1] = SAT_CAST(g); dst[dst_idx + 2] = SAT_CAST(r); -#if channels == 4 - dst[dst_idx + 3] = MAX_NUM; +#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 } } diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 977fd206da..e98268425d 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -121,149 +121,65 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) #define CVTCODE(name) COLOR_ ## name +// RGB[A] <-> BGR[A] + +OCL_TEST_P(CvtColor, BGR2BGRA) { doTest(3, 4, CVTCODE(BGR2BGRA)); } +OCL_TEST_P(CvtColor, RGB2RGBA) { doTest(3, 4, CVTCODE(BGR2BGRA)); } +OCL_TEST_P(CvtColor, BGRA2BGR) { doTest(4, 3, CVTCODE(BGRA2BGR)); } +OCL_TEST_P(CvtColor, RGBA2RGB) { doTest(4, 3, CVTCODE(BGRA2BGR)); } +OCL_TEST_P(CvtColor, BGR2RGBA) { doTest(3, 4, CVTCODE(BGR2RGBA)); } +OCL_TEST_P(CvtColor, RGB2BGRA) { doTest(3, 4, CVTCODE(BGR2RGBA)); } +OCL_TEST_P(CvtColor, RGBA2BGR) { doTest(4, 3, CVTCODE(RGBA2BGR)); } +OCL_TEST_P(CvtColor, BGRA2RGB) { doTest(4, 3, CVTCODE(RGBA2BGR)); } +OCL_TEST_P(CvtColor, BGR2RGB) { doTest(3, 3, CVTCODE(BGR2RGB)); } +OCL_TEST_P(CvtColor, RGB2BGR) { doTest(3, 3, CVTCODE(BGR2RGB)); } +OCL_TEST_P(CvtColor, BGRA2RGBA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); } +OCL_TEST_P(CvtColor, RGBA2BGRA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); } + // RGB <-> Gray -OCL_TEST_P(CvtColor, RGB2GRAY) -{ - doTest(3, 1, CVTCODE(RGB2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2RGB) -{ - doTest(1, 3, CVTCODE(GRAY2RGB)); -} - -OCL_TEST_P(CvtColor, BGR2GRAY) -{ - doTest(3, 1, CVTCODE(BGR2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2BGR) -{ - doTest(1, 3, CVTCODE(GRAY2BGR)); -} - -OCL_TEST_P(CvtColor, RGBA2GRAY) -{ - doTest(4, 1, CVTCODE(RGBA2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2RGBA) -{ - doTest(1, 4, CVTCODE(GRAY2RGBA)); -} - -OCL_TEST_P(CvtColor, BGRA2GRAY) -{ - doTest(4, 1, CVTCODE(BGRA2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2BGRA) -{ - doTest(1, 4, CVTCODE(GRAY2BGRA)); -} +OCL_TEST_P(CvtColor, RGB2GRAY) { doTest(3, 1, CVTCODE(RGB2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2RGB) { doTest(1, 3, CVTCODE(GRAY2RGB)); } +OCL_TEST_P(CvtColor, BGR2GRAY) { doTest(3, 1, CVTCODE(BGR2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2BGR) { doTest(1, 3, CVTCODE(GRAY2BGR)); } +OCL_TEST_P(CvtColor, RGBA2GRAY) { doTest(4, 1, CVTCODE(RGBA2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2RGBA) { doTest(1, 4, CVTCODE(GRAY2RGBA)); } +OCL_TEST_P(CvtColor, BGRA2GRAY) { doTest(4, 1, CVTCODE(BGRA2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2BGRA) { doTest(1, 4, CVTCODE(GRAY2BGRA)); } // RGB <-> YUV -OCL_TEST_P(CvtColor, RGB2YUV) -{ - doTest(3, 3, CVTCODE(RGB2YUV)); -} -OCL_TEST_P(CvtColor, BGR2YUV) -{ - doTest(3, 3, CVTCODE(BGR2YUV)); -} -OCL_TEST_P(CvtColor, RGBA2YUV) -{ - doTest(4, 3, CVTCODE(RGB2YUV)); -} -OCL_TEST_P(CvtColor, BGRA2YUV) -{ - doTest(4, 3, CVTCODE(BGR2YUV)); -} -OCL_TEST_P(CvtColor, YUV2RGB) -{ - doTest(3, 3, CVTCODE(YUV2RGB)); -} -OCL_TEST_P(CvtColor, YUV2BGR) -{ - doTest(3, 3, CVTCODE(YUV2BGR)); -} -OCL_TEST_P(CvtColor, YUV2RGBA) -{ - doTest(3, 4, CVTCODE(YUV2RGB)); -} -OCL_TEST_P(CvtColor, YUV2BGRA) -{ - doTest(3, 4, CVTCODE(YUV2BGR)); -} +OCL_TEST_P(CvtColor, RGB2YUV) { doTest(3, 3, CVTCODE(RGB2YUV)); } +OCL_TEST_P(CvtColor, BGR2YUV) { doTest(3, 3, CVTCODE(BGR2YUV)); } +OCL_TEST_P(CvtColor, RGBA2YUV) { doTest(4, 3, CVTCODE(RGB2YUV)); } +OCL_TEST_P(CvtColor, BGRA2YUV) { doTest(4, 3, CVTCODE(BGR2YUV)); } +OCL_TEST_P(CvtColor, YUV2RGB) { doTest(3, 3, CVTCODE(YUV2RGB)); } +OCL_TEST_P(CvtColor, YUV2BGR) { doTest(3, 3, CVTCODE(YUV2BGR)); } +OCL_TEST_P(CvtColor, YUV2RGBA) { doTest(3, 4, CVTCODE(YUV2RGB)); } +OCL_TEST_P(CvtColor, YUV2BGRA) { doTest(3, 4, CVTCODE(YUV2BGR)); } // RGB <-> YCrCb -OCL_TEST_P(CvtColor, RGB2YCrCb) -{ - doTest(3, 3, CVTCODE(RGB2YCrCb)); -} -OCL_TEST_P(CvtColor, BGR2YCrCb) -{ - doTest(3, 3, CVTCODE(BGR2YCrCb)); -} -OCL_TEST_P(CvtColor, RGBA2YCrCb) -{ - doTest(4, 3, CVTCODE(RGB2YCrCb)); -} -OCL_TEST_P(CvtColor, BGRA2YCrCb) -{ - doTest(4, 3, CVTCODE(BGR2YCrCb)); -} -OCL_TEST_P(CvtColor, YCrCb2RGB) -{ - doTest(3, 3, CVTCODE(YCrCb2RGB)); -} -OCL_TEST_P(CvtColor, YCrCb2BGR) -{ - doTest(3, 3, CVTCODE(YCrCb2BGR)); -} -OCL_TEST_P(CvtColor, YCrCb2RGBA) -{ - doTest(3, 4, CVTCODE(YCrCb2RGB)); -} -OCL_TEST_P(CvtColor, YCrCb2BGRA) -{ - doTest(3, 4, CVTCODE(YCrCb2BGR)); -} +OCL_TEST_P(CvtColor, RGB2YCrCb) { doTest(3, 3, CVTCODE(RGB2YCrCb)); } +OCL_TEST_P(CvtColor, BGR2YCrCb) { doTest(3, 3, CVTCODE(BGR2YCrCb)); } +OCL_TEST_P(CvtColor, RGBA2YCrCb) { doTest(4, 3, CVTCODE(RGB2YCrCb)); } +OCL_TEST_P(CvtColor, BGRA2YCrCb) { doTest(4, 3, CVTCODE(BGR2YCrCb)); } +OCL_TEST_P(CvtColor, YCrCb2RGB) { doTest(3, 3, CVTCODE(YCrCb2RGB)); } +OCL_TEST_P(CvtColor, YCrCb2BGR) { doTest(3, 3, CVTCODE(YCrCb2BGR)); } +OCL_TEST_P(CvtColor, YCrCb2RGBA) { doTest(3, 4, CVTCODE(YCrCb2RGB)); } +OCL_TEST_P(CvtColor, YCrCb2BGRA) { doTest(3, 4, CVTCODE(YCrCb2BGR)); } // RGB <-> XYZ -OCL_TEST_P(CvtColor, RGB2XYZ) -{ - doTest(3, 3, CVTCODE(RGB2XYZ)); -} -OCL_TEST_P(CvtColor, BGR2XYZ) -{ - doTest(3, 3, CVTCODE(BGR2XYZ)); -} -OCL_TEST_P(CvtColor, RGBA2XYZ) -{ - doTest(4, 3, CVTCODE(RGB2XYZ)); -} -OCL_TEST_P(CvtColor, BGRA2XYZ) -{ - doTest(4, 3, CVTCODE(BGR2XYZ)); -} +OCL_TEST_P(CvtColor, RGB2XYZ) { doTest(3, 3, CVTCODE(RGB2XYZ)); } +OCL_TEST_P(CvtColor, BGR2XYZ) { doTest(3, 3, CVTCODE(BGR2XYZ)); } +OCL_TEST_P(CvtColor, RGBA2XYZ) { doTest(4, 3, CVTCODE(RGB2XYZ)); } +OCL_TEST_P(CvtColor, BGRA2XYZ) { doTest(4, 3, CVTCODE(BGR2XYZ)); } -OCL_TEST_P(CvtColor, XYZ2RGB) -{ - doTest(3, 3, CVTCODE(XYZ2RGB)); -} -OCL_TEST_P(CvtColor, XYZ2BGR) -{ - doTest(3, 3, CVTCODE(XYZ2BGR)); -} -OCL_TEST_P(CvtColor, XYZ2RGBA) -{ - doTest(3, 4, CVTCODE(XYZ2RGB)); -} -OCL_TEST_P(CvtColor, XYZ2BGRA) -{ - doTest(3, 4, CVTCODE(XYZ2BGR)); -} +OCL_TEST_P(CvtColor, XYZ2RGB) { doTest(3, 3, CVTCODE(XYZ2RGB)); } +OCL_TEST_P(CvtColor, XYZ2BGR) { doTest(3, 3, CVTCODE(XYZ2BGR)); } +OCL_TEST_P(CvtColor, XYZ2RGBA) { doTest(3, 4, CVTCODE(XYZ2RGB)); } +OCL_TEST_P(CvtColor, XYZ2BGRA) { doTest(3, 4, CVTCODE(XYZ2BGR)); } // YUV -> RGBA_NV12 @@ -289,25 +205,10 @@ struct CvtColor_YUV420 : } }; -OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) -{ - doTest(1, 4, CV_YUV2RGBA_NV12); -} - -OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) -{ - doTest(1, 4, CV_YUV2BGRA_NV12); -} - -OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) -{ - doTest(1, 3, CV_YUV2RGB_NV12); -} - -OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) -{ - doTest(1, 3, CV_YUV2BGR_NV12); -} +OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { doTest(1, 4, CV_YUV2RGBA_NV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { doTest(1, 4, CV_YUV2BGRA_NV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { doTest(1, 3, CV_YUV2RGB_NV12); } +OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { doTest(1, 3, CV_YUV2BGR_NV12); } INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine( From 1f421fce01efbfc2472d0b191cfccdd98da2f708 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Mon, 11 Nov 2013 13:04:55 +0400 Subject: [PATCH 7/8] added RGB5x5 <-> RGB conversion --- modules/ocl/src/color.cpp | 67 ++++++++++++++++++++++++++++- modules/ocl/src/opencl/cvt_color.cl | 58 +++++++++++++++++++++++++ modules/ocl/test/test_color.cpp | 28 ++++++++++++ 3 files changed, 151 insertions(+), 2 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index f27366a6ae..f9bccc89b6 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -122,6 +122,50 @@ static void RGB_caller(const oclMat &src, oclMat &dst, bool reverse) openCLExecuteKernel(src.clCxt, &cvt_color, "RGB", gt, lt, args, -1, -1, build_options.c_str()); } +static void RGB5x52RGB_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits) +{ + std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D dcn=%d", + src.depth(), greenbits, dst.channels()); + int src_offset = src.offset >> 1, src_step = src.step >> 1; + int dst_offset = (int)dst.offset, dst_step = (int)dst.step; + + vector > args; + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); + + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, "RGB5x52RGB", gt, lt, args, -1, -1, build_options.c_str()); +} + +static void RGB2RGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits) +{ + std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D scn=%d", + src.depth(), greenbits, src.channels()); + int src_offset = (int)src.offset, src_step = (int)src.step; + int dst_offset = dst.offset >> 1, dst_step = dst.step >> 1; + + vector > args; + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&bidx)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data)); + args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); + + size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; + openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2RGB5x5", gt, lt, args, -1, -1, build_options.c_str()); +} + static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) { Size sz = src.size(); @@ -141,12 +185,31 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) RGB_caller(src, dst, reverse); break; } - /* case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555: + { + CV_Assert((scn == 3 || scn == 4) && depth == CV_8U ); + bidx = code == CV_BGR2BGR565 || code == CV_BGR2BGR555 || + code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2; + int greenbits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565 || + code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5; + dst.create(sz, CV_8UC2); + RGB2RGB5x5_caller(src, dst, bidx, greenbits); + break; + } case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: - */ + { + dcn = code == CV_BGR5652BGRA || code == CV_BGR5552BGRA || code == CV_BGR5652RGBA || code == CV_BGR5552RGBA ? 4 : 3; + CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U); + bidx = code == CV_BGR5652BGR || code == CV_BGR5552BGR || + code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2; + int greenbits = code == CV_BGR5652BGR || code == CV_BGR5652RGB || + code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5; + dst.create(sz, CV_MAKETYPE(depth, dcn)); + RGB5x52RGB_caller(src, dst, bidx, greenbits); + break; + } case CV_RGB2GRAY: case CV_BGR2GRAY: case CV_RGBA2GRAY: case CV_BGRA2GRAY: { diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 916d44bf93..b0bd7d0848 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -437,3 +437,61 @@ __kernel void RGB(int cols, int rows, int src_step, int dst_step, #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 + } +} diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index e98268425d..105ee89d90 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -181,6 +181,30 @@ OCL_TEST_P(CvtColor, XYZ2BGR) { doTest(3, 3, CVTCODE(XYZ2BGR)); } OCL_TEST_P(CvtColor, XYZ2RGBA) { doTest(3, 4, CVTCODE(XYZ2RGB)); } OCL_TEST_P(CvtColor, XYZ2BGRA) { doTest(3, 4, CVTCODE(XYZ2BGR)); } +// RGB5x5 <-> RGB + +typedef CvtColor CvtColor8u; + +OCL_TEST_P(CvtColor8u, BGR5652BGR) { doTest(2, 3, CVTCODE(BGR5652BGR)); } +OCL_TEST_P(CvtColor8u, BGR5652RGB) { doTest(2, 3, CVTCODE(BGR5652RGB)); } +OCL_TEST_P(CvtColor8u, BGR5652BGRA) { doTest(2, 4, CVTCODE(BGR5652BGRA)); } +OCL_TEST_P(CvtColor8u, BGR5652RGBA) { doTest(2, 4, CVTCODE(BGR5652RGBA)); } + +OCL_TEST_P(CvtColor8u, BGR5552BGR) { doTest(2, 3, CVTCODE(BGR5552BGR)); } +OCL_TEST_P(CvtColor8u, BGR5552RGB) { doTest(2, 3, CVTCODE(BGR5552RGB)); } +OCL_TEST_P(CvtColor8u, BGR5552BGRA) { doTest(2, 4, CVTCODE(BGR5552BGRA)); } +OCL_TEST_P(CvtColor8u, BGR5552RGBA) { doTest(2, 4, CVTCODE(BGR5552RGBA)); } + +OCL_TEST_P(CvtColor8u, BGR2BGR565) { doTest(3, 2, CVTCODE(BGR2BGR565)); } +OCL_TEST_P(CvtColor8u, RGB2BGR565) { doTest(3, 2, CVTCODE(RGB2BGR565)); } +OCL_TEST_P(CvtColor8u, BGRA2BGR565) { doTest(4, 2, CVTCODE(BGRA2BGR565)); } +OCL_TEST_P(CvtColor8u, RGBA2BGR565) { doTest(4, 2, CVTCODE(RGBA2BGR565)); } + +OCL_TEST_P(CvtColor8u, BGR2BGR555) { doTest(3, 2, CVTCODE(BGR2BGR555)); } +OCL_TEST_P(CvtColor8u, RGB2BGR555) { doTest(3, 2, CVTCODE(RGB2BGR555)); } +OCL_TEST_P(CvtColor8u, BGRA2BGR555) { doTest(4, 2, CVTCODE(BGRA2BGR555)); } +OCL_TEST_P(CvtColor8u, RGBA2BGR555) { doTest(4, 2, CVTCODE(RGBA2BGR555)); } + // YUV -> RGBA_NV12 struct CvtColor_YUV420 : @@ -210,6 +234,10 @@ OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { doTest(1, 4, CV_YUV2BGRA_NV12); } OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { doTest(1, 3, CV_YUV2RGB_NV12); } OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { doTest(1, 3, CV_YUV2BGR_NV12); } + +INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor8u, + testing::Combine(testing::Values(MatDepth(CV_8U)), Bool())); + INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, testing::Combine( testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)), From eda6360fa329c12edbdec55ef0bbfffd63f6c590 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Mon, 11 Nov 2013 16:30:23 +0400 Subject: [PATCH 8/8] added RGB5x5 <-> Gray --- modules/ocl/src/color.cpp | 30 +++++++++++++----- modules/ocl/src/opencl/cvt_color.cl | 49 +++++++++++++++++++++++++++++ modules/ocl/test/test_color.cpp | 8 +++++ 3 files changed, 80 insertions(+), 7 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index f9bccc89b6..b01dcef1f4 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -122,12 +122,12 @@ static void RGB_caller(const oclMat &src, oclMat &dst, bool reverse) openCLExecuteKernel(src.clCxt, &cvt_color, "RGB", gt, lt, args, -1, -1, build_options.c_str()); } -static void RGB5x52RGB_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits) +static void fromRGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits, const std::string & kernelName) { std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D dcn=%d", src.depth(), greenbits, dst.channels()); int src_offset = src.offset >> 1, src_step = src.step >> 1; - int dst_offset = (int)dst.offset, dst_step = (int)dst.step; + int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step / dst.elemSize1(); vector > args; args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols)); @@ -141,10 +141,10 @@ static void RGB5x52RGB_caller(const oclMat &src, oclMat &dst, int bidx, int gree args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB5x52RGB", gt, lt, args, -1, -1, build_options.c_str()); + openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } -static void RGB2RGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits) +static void toRGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits, const std::string & kernelName) { std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D scn=%d", src.depth(), greenbits, src.channels()); @@ -163,7 +163,7 @@ static void RGB2RGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int gree args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset )); size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 }; - openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2RGB5x5", gt, lt, args, -1, -1, build_options.c_str()); + openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) @@ -194,7 +194,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) int greenbits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565 || code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5; dst.create(sz, CV_8UC2); - RGB2RGB5x5_caller(src, dst, bidx, greenbits); + toRGB5x5_caller(src, dst, bidx, greenbits, "RGB2RGB5x5"); break; } case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: @@ -207,7 +207,23 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) int greenbits = code == CV_BGR5652BGR || code == CV_BGR5652RGB || code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5; dst.create(sz, CV_MAKETYPE(depth, dcn)); - RGB5x52RGB_caller(src, dst, bidx, greenbits); + fromRGB5x5_caller(src, dst, bidx, greenbits, "RGB5x52RGB"); + break; + } + case CV_BGR5652GRAY: case CV_BGR5552GRAY: + { + CV_Assert(scn == 2 && depth == CV_8U); + dst.create(sz, CV_8UC1); + int greenbits = code == CV_BGR5652GRAY ? 6 : 5; + fromRGB5x5_caller(src, dst, -1, greenbits, "BGR5x52Gray"); + break; + } + case CV_GRAY2BGR565: case CV_GRAY2BGR555: + { + CV_Assert(scn == 1 && depth == CV_8U); + dst.create(sz, CV_8UC2); + int greenbits = code == CV_GRAY2BGR565 ? 6 : 5; + toRGB5x5_caller(src, dst, -1, greenbits, "Gray2BGR5x5"); break; } case CV_RGB2GRAY: case CV_BGR2GRAY: diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index b0bd7d0848..210d1b766e 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -495,3 +495,52 @@ __kernel void RGB2RGB5x5(int cols, int rows, int src_step, int dst_step, int bid #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 + } +} diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 105ee89d90..732a0f8e45 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -205,6 +205,14 @@ OCL_TEST_P(CvtColor8u, RGB2BGR555) { doTest(3, 2, CVTCODE(RGB2BGR555)); } OCL_TEST_P(CvtColor8u, BGRA2BGR555) { doTest(4, 2, CVTCODE(BGRA2BGR555)); } OCL_TEST_P(CvtColor8u, RGBA2BGR555) { doTest(4, 2, CVTCODE(RGBA2BGR555)); } +// RGB5x5 <-> Gray + +OCL_TEST_P(CvtColor8u, BGR5652GRAY) { doTest(2, 1, CVTCODE(BGR5652GRAY)); } +OCL_TEST_P(CvtColor8u, BGR5552GRAY) { doTest(2, 1, CVTCODE(BGR5552GRAY)); } + +OCL_TEST_P(CvtColor8u, GRAY2BGR565) { doTest(1, 2, CVTCODE(GRAY2BGR565)); } +OCL_TEST_P(CvtColor8u, GRAY2BGR555) { doTest(1, 2, CVTCODE(GRAY2BGR555)); } + // YUV -> RGBA_NV12 struct CvtColor_YUV420 :