Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

fixed ocl::cvtColor for CV_YUV2BGRA and CV_YUV2RGBA

Browse Source
pull/1775/head
Ilya Lavrenov 12 years ago
parent 800d53f76b
commit eba6754b06
3 changed files with 145 additions and 123 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 58
      modules/ocl/src/color.cpp
  2. 123
      modules/ocl/src/opencl/cvt_color.cl
  3. 87
      modules/ocl/test/test_color.cpp

58
modules/ocl/src/color.cpp
Unescape View File

@ -50,20 +50,8 @@
using namespace cv; using namespace cv;
using namespace cv::ocl; using namespace cv::ocl;
#ifndef CV_DESCALE static void RGB2Gray_caller(const oclMat &src, oclMat &dst, int bidx)
#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n))
#endif
#ifndef FLT_EPSILON
#define FLT_EPSILON 1.192092896e-07F
#endif
namespace
{
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 src_offset = src.offset / src.elemSize1(), src_step = src.step1();
int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
@ -74,19 +62,19 @@ void RGB2Gray_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.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); 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 *)&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_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); 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_mem) , (void *)&dst.data));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset )); args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_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}; 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()); openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2Gray", gt, lt, args, -1, -1, build_options.c_str());
} }
void Gray2RGB_caller(const oclMat &src, oclMat &dst) static void Gray2RGB_caller(const oclMat &src, oclMat &dst)
{ {
int channels = dst.channels();
std::string build_options = format("-D DEPTH_%d", src.depth()); std::string build_options = format("-D DEPTH_%d", src.depth());
int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
@ -96,6 +84,7 @@ void Gray2RGB_caller(const oclMat &src, oclMat &dst)
args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows)); 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 *)&src_step));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_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 *)&src.data));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.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 *)&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()); 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()); std::string build_options = format("-D DEPTH_%d", src.depth());
int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.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.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); 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 *)&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_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); 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_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()); 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 src_offset = src.offset / src.elemSize1(), src_step = src.step1();
int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.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()); 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()); std::string build_options = format("-D DEPTH_%d", src.depth());
int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); 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()); 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()); std::string build_options = format("-D DEPTH_%d", src.depth());
int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.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.rows));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step)); 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 *)&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_int) , (void *)&bidx));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); 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_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()); 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(); 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); 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: case CV_RGB2YUV:
{ {
CV_Assert(scn == 3 || scn == 4); 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)); dst.create(sz, CV_MAKETYPE(depth, 3));
RGB2YUV_caller(src, dst, bidx); RGB2YUV_caller(src, dst, bidx);
break; break;
@ -247,9 +232,11 @@ void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
case CV_YUV2BGR: case CV_YUV2BGR:
case CV_YUV2RGB: case CV_YUV2RGB:
{ {
CV_Assert(scn == 3 || scn == 4); if( dcn <= 0 )
bidx = code == CV_YUV2RGB ? 0 : 2; dcn = 3;
dst.create(sz, CV_MAKETYPE(depth, 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); YUV2RGB_caller(src, dst, bidx);
break; break;
} }
@ -260,7 +247,7 @@ void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
{ {
CV_Assert(scn == 1); CV_Assert(scn == 1);
CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); 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; bidx = code == CV_YUV2BGRA_NV12 || code == CV_YUV2BGR_NV12 ? 0 : 2;
Size dstSz(sz.width, sz.height * 2 / 3); 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_YCrCb2BGR:
case CV_YCrCb2RGB: 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; 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" ); CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );
} }
} }
}
void cv::ocl::cvtColor(const oclMat &src, oclMat &dst, int code, int dcn) void cv::ocl::cvtColor(const oclMat &src, oclMat &dst, int code, int dcn)
{ {

123
modules/ocl/src/opencl/cvt_color.cl
Unescape View File

@ -46,22 +46,18 @@
/**************************************PUBLICFUNC*************************************/ /**************************************PUBLICFUNC*************************************/
#if defined (DOUBLE_SUPPORT)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#if defined (DEPTH_0) #if defined (DEPTH_0)
#define DATA_TYPE uchar #define DATA_TYPE uchar
#define MAX_NUM 255 #define MAX_NUM 255
#define HALF_MAX 128 #define HALF_MAX 128
#define SAT_CAST(num) convert_uchar_sat(num) #define SAT_CAST(num) convert_uchar_sat_rte(num)
#endif #endif
#if defined (DEPTH_2) #if defined (DEPTH_2)
#define DATA_TYPE ushort #define DATA_TYPE ushort
#define MAX_NUM 65535 #define MAX_NUM 65535
#define HALF_MAX 32768 #define HALF_MAX 32768
#define SAT_CAST(num) convert_ushort_sat(num) #define SAT_CAST(num) convert_ushort_sat_rte(num)
#endif #endif
#if defined (DEPTH_5) #if defined (DEPTH_5)
@ -71,11 +67,7 @@
#define SAT_CAST(num) (num) #define SAT_CAST(num) (num)
#endif #endif
#ifndef DATA_TYPE #define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n))
#define DATA_TYPE UNDEFINED
#endif
#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
enum enum
{ {
@ -89,16 +81,16 @@ enum
///////////////////////////////////// RGB <-> GRAY ////////////////////////////////////// ///////////////////////////////////// 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 bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset) int src_offset, int dst_offset)
{ {
const int x = get_global_id(0); int x = get_global_id(0);
const int y = get_global_id(1); int y = get_global_id(1);
if (y < rows && x < cols) 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); int dst_idx = mad24(y, dst_step, dst_offset + x);
#if defined (DEPTH_5) #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; 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, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset) int src_offset, int dst_offset)
{ {
const int x = get_global_id(0); int x = get_global_id(0);
const int y = get_global_id(1); int y = get_global_id(1);
if (y < rows && x < cols) if (y < rows && x < cols)
{ {
int src_idx = mad24(y, src_step, src_offset + x); 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]; DATA_TYPE val = src[src_idx];
dst[dst_idx++] = val; dst[dst_idx] = val;
dst[dst_idx++] = val; dst[dst_idx + 1] = val;
dst[dst_idx++] = val; dst[dst_idx + 2] = val;
dst[dst_idx] = MAX_NUM; 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 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 }; __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 bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset) 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) if (y < rows && x < cols)
{ {
x *= channels; x <<= 2;
int src_idx = mad24(y, src_step, src_offset + x); int src_idx = mad24(y, src_step, src_offset + x);
int dst_idx = mad24(y, dst_step, dst_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x);
dst += dst_idx; DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
const DATA_TYPE rgb[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]};
#if defined (DEPTH_5) #if defined (DEPTH_5)
__constant float * coeffs = c_RGB2YUVCoeffs_f; __constant float * coeffs = c_RGB2YUVCoeffs_f;
const DATA_TYPE Y = rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2]; DATA_TYPE Y = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx];
const DATA_TYPE Cr = (rgb[bidx] - Y) * coeffs[3] + HALF_MAX; DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX;
const DATA_TYPE Cb = (rgb[bidx^2] - Y) * coeffs[4] + HALF_MAX; DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX;
#else #else
__constant int * coeffs = c_RGB2YUVCoeffs_i; __constant int * coeffs = c_RGB2YUVCoeffs_i;
const int delta = HALF_MAX * (1 << yuv_shift); 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); 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] - Y) * coeffs[3] + delta, yuv_shift); int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift);
const int Cb = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[4] + delta, yuv_shift); int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift);
#endif #endif
dst[0] = SAT_CAST( Y ); dst[dst_idx] = SAT_CAST( Y );
dst[1] = SAT_CAST( Cr ); dst[dst_idx + 1] = SAT_CAST( Cr );
dst[2] = SAT_CAST( Cb ); dst[dst_idx + 2] = SAT_CAST( Cb );
} }
} }
__constant float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f }; __constant float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f };
__constant int c_YUV2RGBCoeffs_i[5] = { 33292, -6472, -9519, 18678 }; __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 bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset) 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) if (y < rows && x < cols)
{ {
x *= channels; x <<= 2;
int src_idx = mad24(y, src_step, src_offset + x); int src_idx = mad24(y, src_step, src_offset + x);
int dst_idx = mad24(y, dst_step, dst_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x);
dst += dst_idx; DATA_TYPE yuv[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
const DATA_TYPE yuv[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]};
#if defined (DEPTH_5) #if defined (DEPTH_5)
__constant float * coeffs = c_YUV2RGBCoeffs_f; __constant float * coeffs = c_YUV2RGBCoeffs_f;
const float b = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[3]; 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]; 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 r = yuv[0] + (yuv[1] - HALF_MAX) * coeffs[0];
#else #else
__constant int * coeffs = c_YUV2RGBCoeffs_i; __constant int * coeffs = c_YUV2RGBCoeffs_i;
const int b = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[3], yuv_shift); 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); 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 r = yuv[0] + CV_DESCALE((yuv[1] - HALF_MAX) * coeffs[0], yuv_shift);
#endif #endif
dst[bidx^2] = SAT_CAST( b ); dst[dst_idx + bidx] = SAT_CAST( b );
dst[1] = SAT_CAST( g ); dst[dst_idx + 1] = SAT_CAST( g );
dst[bidx] = SAT_CAST( r ); 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 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}; __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 bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset) 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) if (y < rows && x < cols)
{ {
x *= channels; x <<= 2;
int src_idx = mad24(y, src_step, src_offset + x); int src_idx = mad24(y, src_step, src_offset + x);
int dst_idx = mad24(y, dst_step, dst_offset + x); int dst_idx = mad24(y, dst_step, dst_offset + x);
dst += dst_idx; DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
const DATA_TYPE rgb[] = { src[src_idx], src[src_idx + 1], src[src_idx + 2] };
#if defined (DEPTH_5) #if defined (DEPTH_5)
__constant float * coeffs = c_RGB2YCrCbCoeffs_f; __constant float * coeffs = c_RGB2YCrCbCoeffs_f;
const DATA_TYPE Y = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx]; 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; DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX;
const DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX; DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX;
#else #else
__constant int * coeffs = c_RGB2YCrCbCoeffs_i; __constant int * coeffs = c_RGB2YCrCbCoeffs_i;
const int delta = HALF_MAX * (1 << yuv_shift); 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); 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); 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 Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift);
#endif #endif
dst[0] = SAT_CAST( Y ); dst[dst_idx] = SAT_CAST( Y );
dst[1] = SAT_CAST( Cr ); dst[dst_idx + 1] = SAT_CAST( Cr );
dst[2] = SAT_CAST( Cb ); dst[dst_idx + 2] = SAT_CAST( Cb );
} }
} }

87
modules/ocl/test/test_color.cpp
Unescape View File

@ -46,8 +46,6 @@
#include "test_precomp.hpp" #include "test_precomp.hpp"
#ifdef HAVE_OPENCL #ifdef HAVE_OPENCL
namespace
{
using namespace testing; using namespace testing;
/////////////////////////////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////////////////////////////////////////////
@ -108,8 +106,8 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool)
{ {
random_roi(channelsIn, channelsOut); random_roi(channelsIn, channelsOut);
cv::cvtColor(src1_roi, dst1_roi, code); cv::cvtColor(src1_roi, dst1_roi, code, channelsOut);
cv::ocl::cvtColor(gsrc1_roi, gdst1_roi, code); cv::ocl::cvtColor(gsrc1_roi, gdst1_roi, code, channelsOut);
Near(); Near();
} }
@ -125,7 +123,7 @@ OCL_TEST_P(CvtColor, RGB2GRAY)
OCL_TEST_P(CvtColor, GRAY2RGB) OCL_TEST_P(CvtColor, GRAY2RGB)
{ {
doTest(1, 3, CVTCODE(GRAY2RGB)); doTest(1, 3, CVTCODE(GRAY2RGB));
}; }
OCL_TEST_P(CvtColor, BGR2GRAY) OCL_TEST_P(CvtColor, BGR2GRAY)
{ {
@ -134,25 +132,26 @@ OCL_TEST_P(CvtColor, BGR2GRAY)
OCL_TEST_P(CvtColor, GRAY2BGR) OCL_TEST_P(CvtColor, GRAY2BGR)
{ {
doTest(1, 3, CVTCODE(GRAY2BGR)); doTest(1, 3, CVTCODE(GRAY2BGR));
}; }
OCL_TEST_P(CvtColor, RGBA2GRAY) OCL_TEST_P(CvtColor, RGBA2GRAY)
{ {
doTest(3, 1, CVTCODE(RGBA2GRAY)); doTest(4, 1, CVTCODE(RGBA2GRAY));
} }
OCL_TEST_P(CvtColor, GRAY2RGBA) OCL_TEST_P(CvtColor, GRAY2RGBA)
{ {
doTest(1, 3, CVTCODE(GRAY2RGBA)); doTest(1, 4, CVTCODE(GRAY2RGBA));
}; }
OCL_TEST_P(CvtColor, BGRA2GRAY) OCL_TEST_P(CvtColor, BGRA2GRAY)
{ {
doTest(3, 1, CVTCODE(BGRA2GRAY)); doTest(4, 1, CVTCODE(BGRA2GRAY));
} }
OCL_TEST_P(CvtColor, GRAY2BGRA) OCL_TEST_P(CvtColor, GRAY2BGRA)
{ {
doTest(1, 3, CVTCODE(GRAY2BGRA)); doTest(1, 4, CVTCODE(GRAY2BGRA));
}; }
OCL_TEST_P(CvtColor, RGB2YUV) OCL_TEST_P(CvtColor, RGB2YUV)
{ {
@ -162,6 +161,14 @@ OCL_TEST_P(CvtColor, BGR2YUV)
{ {
doTest(3, 3, CVTCODE(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) OCL_TEST_P(CvtColor, YUV2RGB)
{ {
doTest(3, 3, CVTCODE(YUV2RGB)); doTest(3, 3, CVTCODE(YUV2RGB));
@ -170,6 +177,16 @@ OCL_TEST_P(CvtColor, YUV2BGR)
{ {
doTest(3, 3, CVTCODE(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) OCL_TEST_P(CvtColor, RGB2YCrCb)
{ {
doTest(3, 3, CVTCODE(RGB2YCrCb)); doTest(3, 3, CVTCODE(RGB2YCrCb));
@ -178,8 +195,33 @@ OCL_TEST_P(CvtColor, BGR2YCrCb)
{ {
doTest(3, 3, CVTCODE(BGR2YCrCb)); doTest(3, 3, CVTCODE(BGR2YCrCb));
} }
OCL_TEST_P(CvtColor, RGBA2YCrCb)
struct CvtColor_YUV420 : CvtColor {
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 :
public CvtColor
{ {
void random_roi(int channelsIn, int channelsOut) void random_roi(int channelsIn, int channelsOut)
{ {
@ -203,37 +245,32 @@ struct CvtColor_YUV420 : CvtColor
OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12)
{ {
doTest(1, 4, CV_YUV2RGBA_NV12); doTest(1, 4, CV_YUV2RGBA_NV12);
}; }
OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12)
{ {
doTest(1, 4, CV_YUV2BGRA_NV12); doTest(1, 4, CV_YUV2BGRA_NV12);
}; }
OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12)
{ {
doTest(1, 3, CV_YUV2RGB_NV12); doTest(1, 3, CV_YUV2RGB_NV12);
}; }
OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12)
{ {
doTest(1, 3, CV_YUV2BGR_NV12); doTest(1, 3, CV_YUV2BGR_NV12);
}; }
INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor, INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor,
testing::Combine( testing::Combine(
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)), testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),
Bool() Bool()));
)
);
INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_YUV420, INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_YUV420,
testing::Combine( testing::Combine(
testing::Values(MatDepth(CV_8U)), testing::Values(MatDepth(CV_8U)),
Bool() Bool()));
)
);
}
#endif #endif

Write Preview
Loading…
Cancel
Save