added YCrCb to RGB, BGR, RGBA, BGRA modes to ocl::cvtColor
This commit is contained in:
parent
eba6754b06
commit
a57030a0cd
@ -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());
|
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<pair<size_t , const void *> > 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)
|
static void RGB2YCrCb_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());
|
||||||
@ -270,9 +294,9 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
|
|||||||
if( dcn <= 0 )
|
if( dcn <= 0 )
|
||||||
dcn = 3;
|
dcn = 3;
|
||||||
CV_Assert(scn == 3 && (dcn == 3 || dcn == 4));
|
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));
|
dst.create(sz, CV_MAKETYPE(depth, dcn));
|
||||||
// YUV2RGB_caller(src, dst, bidx);
|
YCrCb2RGB_caller(src, dst, bidx);
|
||||||
break;
|
break;
|
||||||
}
|
}
|
||||||
/*
|
/*
|
||||||
|
@ -293,3 +293,41 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step,
|
|||||||
dst[dst_idx + 2] = SAT_CAST( Cb );
|
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;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
@ -203,22 +203,22 @@ OCL_TEST_P(CvtColor, BGRA2YCrCb)
|
|||||||
{
|
{
|
||||||
doTest(4, 3, CVTCODE(BGR2YCrCb));
|
doTest(4, 3, CVTCODE(BGR2YCrCb));
|
||||||
}
|
}
|
||||||
//OCL_TEST_P(CvtColor, YCrCb2RGB)
|
OCL_TEST_P(CvtColor, YCrCb2RGB)
|
||||||
//{
|
{
|
||||||
// doTest(3, 3, CVTCODE(YCrCb2RGB));
|
doTest(3, 3, CVTCODE(YCrCb2RGB));
|
||||||
//}
|
}
|
||||||
//OCL_TEST_P(CvtColor, YCrCb2BGR)
|
OCL_TEST_P(CvtColor, YCrCb2BGR)
|
||||||
//{
|
{
|
||||||
// doTest(3, 3, CVTCODE(YCrCb2BGR));
|
doTest(3, 3, CVTCODE(YCrCb2BGR));
|
||||||
//}
|
}
|
||||||
//OCL_TEST_P(CvtColor, YCrCb2RGBA)
|
OCL_TEST_P(CvtColor, YCrCb2RGBA)
|
||||||
//{
|
{
|
||||||
// doTest(3, 4, CVTCODE(YCrCb2RGB));
|
doTest(3, 4, CVTCODE(YCrCb2RGB));
|
||||||
//}
|
}
|
||||||
//OCL_TEST_P(CvtColor, YCrCb2BGRA)
|
OCL_TEST_P(CvtColor, YCrCb2BGRA)
|
||||||
//{
|
{
|
||||||
// doTest(3, 4, CVTCODE(YCrCb2BGR));
|
doTest(3, 4, CVTCODE(YCrCb2BGR));
|
||||||
//}
|
}
|
||||||
|
|
||||||
struct CvtColor_YUV420 :
|
struct CvtColor_YUV420 :
|
||||||
public CvtColor
|
public CvtColor
|
||||||
|
Loading…
x
Reference in New Issue
Block a user