Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

Merge pull request #1775 from ilya-lavrenov:ocl_cvtColor

Browse Source
pull/1785/head
Roman Donchenko 11 years ago committed by OpenCV Buildbot
parent 834705490e eda6360fa3
commit a81efdbb25
3 changed files with 668 additions and 299 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 321
      modules/ocl/src/color.cpp
  2. 400
      modules/ocl/src/opencl/cvt_color.cl
  3. 246
      modules/ocl/test/test_color.cpp

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

@ -50,180 +50,189 @@
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 fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName,
const oclMat & data = oclMat())
{
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<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 *)&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));
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<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_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] = {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());
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());
}
void RGB2YUV_caller(const oclMat &src, oclMat &dst, int bidx)
static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName,
const oclMat & data = oclMat())
{
int channels = src.oclchannels();
std::string build_options = format("-D DEPTH_%d", src.depth());
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();
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 *)&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));
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());
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());
}
void YUV2RGB_caller(const oclMat &src, oclMat &dst, int bidx)
static void RGB_caller(const oclMat &src, oclMat &dst, bool reverse)
{
int channels = src.oclchannels();
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();
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 *)&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));
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());
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());
}
void YUV2RGB_NV12_caller(const oclMat &src, oclMat &dst, int bidx)
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", 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 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 = dst.offset / dst.elemSize1(), dst_step = dst.step / dst.elemSize1();
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 *)&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_int) , (void *)&dst.cols));
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows));
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());
}
void RGB2YCrCb_caller(const oclMat &src, oclMat &dst, int bidx)
static void toRGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits, const std::string & kernelName)
{
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();
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<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 *)&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));
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());
}
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);
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_BGR2GRAY:
case CV_BGRA2GRAY:
case CV_RGB2GRAY:
case CV_RGBA2GRAY:
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:
{
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);
toRGB5x5_caller(src, dst, bidx, greenbits, "RGB2RGB5x5");
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));
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:
case CV_RGBA2GRAY: case CV_BGRA2GRAY:
{
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:
@ -232,40 +241,40 @@ 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:
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);
fromRGB_caller(src, dst, bidx, "RGB2YUV");
break;
}
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));
YUV2RGB_caller(src, dst, bidx);
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));
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 );
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);
dst.create(dstSz, CV_MAKETYPE(depth, dcn));
YUV2RGB_NV12_caller(src, dst, bidx);
toRGB_caller(src, dst, bidx, "YUV2RGBA_NV12");
break;
}
case CV_BGR2YCrCb:
@ -274,20 +283,117 @@ 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:
case CV_YCrCb2RGB:
{
if( dcn <= 0 )
dcn = 3;
CV_Assert(scn == 3 && (dcn == 3 || dcn == 4));
bidx = code == CV_YCrCb2BGR ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, dcn));
toRGB_caller(src, dst, bidx, "YCrCb2RGB");
break;
}
/*
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_32FC1 : CV_32SC1, 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));
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;
}
/*
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:
@ -297,7 +403,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)
{

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

@ -46,36 +46,31 @@
/**************************************PUBLICFUNC*************************************/
#if defined (DOUBLE_SUPPORT)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#if defined (DEPTH_0)
#ifdef 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(num)
#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
#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)
#ifdef 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)
#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,18 +84,18 @@ 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)
#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);
@ -108,22 +103,25 @@ __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 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);
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 dcn == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
@ -132,7 +130,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 +139,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)
#ifdef 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 bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset)
{
@ -178,27 +175,29 @@ __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)
#ifdef 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 dcn == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
@ -209,17 +208,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 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);
@ -261,12 +260,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 +274,273 @@ __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)
#ifdef 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[dst_idx] = SAT_CAST( Y );
dst[dst_idx + 1] = SAT_CAST( Cr );
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 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 dcn == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
///////////////////////////////////// RGB <-> XYZ //////////////////////////////////////
__kernel void RGB2XYZ(int cols, int rows, int src_step, int dst_step,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset, __constant COEFF_TYPE * coeffs)
{
int dx = get_global_id(0);
int dy = get_global_id(1);
if (dy < rows && dx < cols)
{
dx <<= 2;
int src_idx = mad24(dy, src_step, src_offset + dx);
int dst_idx = mad24(dy, dst_step, dst_offset + dx);
DATA_TYPE r = src[src_idx], g = src[src_idx + 1], b = src[src_idx + 2];
#ifdef DEPTH_5
float x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2];
float y = r * coeffs[3] + g * coeffs[4] + b * coeffs[5];
float z = r * coeffs[6] + g * coeffs[7] + b * coeffs[8];
#else
int x = CV_DESCALE(r * coeffs[0] + g * coeffs[1] + b * coeffs[2], xyz_shift);
int y = CV_DESCALE(r * coeffs[3] + g * coeffs[4] + b * coeffs[5], xyz_shift);
int z = CV_DESCALE(r * coeffs[6] + g * coeffs[7] + b * coeffs[8], xyz_shift);
#endif
dst[dst_idx] = SAT_CAST(x);
dst[dst_idx + 1] = SAT_CAST(y);
dst[dst_idx + 2] = SAT_CAST(z);
}
}
__kernel void XYZ2RGB(int cols, int rows, int src_step, int dst_step,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset, __constant COEFF_TYPE * coeffs)
{
int dx = get_global_id(0);
int dy = get_global_id(1);
if (dy < rows && dx < cols)
{
dx <<= 2;
int src_idx = mad24(dy, src_step, src_offset + dx);
int dst_idx = mad24(dy, dst_step, dst_offset + dx);
DATA_TYPE x = src[src_idx], y = src[src_idx + 1], z = src[src_idx + 2];
#ifdef DEPTH_5
float b = x * coeffs[0] + y * coeffs[1] + z * coeffs[2];
float g = x * coeffs[3] + y * coeffs[4] + z * coeffs[5];
float r = x * coeffs[6] + y * coeffs[7] + z * coeffs[8];
#else
int b = CV_DESCALE(x * coeffs[0] + y * coeffs[1] + z * coeffs[2], xyz_shift);
int g = CV_DESCALE(x * coeffs[3] + y * coeffs[4] + z * coeffs[5], xyz_shift);
int r = CV_DESCALE(x * coeffs[6] + y * coeffs[7] + z * coeffs[8], xyz_shift);
#endif
dst[dst_idx] = SAT_CAST(b);
dst[dst_idx + 1] = SAT_CAST(g);
dst[dst_idx + 2] = SAT_CAST(r);
#if dcn == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
///////////////////////////////////// RGB[A] <-> BGR[A] //////////////////////////////////////
__kernel void RGB(int cols, int rows, int src_step, int dst_step,
__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
}
}
///////////////////////////////////// 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
}
}
dst[0] = SAT_CAST( Y );
dst[1] = SAT_CAST( Cr );
dst[2] = SAT_CAST( Cb );
///////////////////////////////////// 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
}
}

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

@ -46,9 +46,8 @@
#include "test_precomp.hpp"
#ifdef HAVE_OPENCL
namespace
{
using namespace testing;
using namespace cv;
///////////////////////////////////////////////////////////////////////////////////////////////////////
// cvtColor
@ -59,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()
{
@ -87,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)
@ -108,78 +111,112 @@ 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);
cvtColor(src_roi, dst_roi, code, channelsOut);
ocl::cvtColor(gsrc_roi, gdst_roi, code, channelsOut);
Near();
}
}
};
#define CVTCODE(name) cv::COLOR_ ## name
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(3, 1, CVTCODE(RGBA2GRAY));
}
OCL_TEST_P(CvtColor, GRAY2RGBA)
{
doTest(1, 3, CVTCODE(GRAY2RGBA));
};
OCL_TEST_P(CvtColor, BGRA2GRAY)
{
doTest(3, 1, CVTCODE(BGRA2GRAY));
}
OCL_TEST_P(CvtColor, GRAY2BGRA)
{
doTest(1, 3, CVTCODE(GRAY2BGRA));
};
OCL_TEST_P(CvtColor, RGB2YUV)
{
doTest(3, 3, CVTCODE(RGB2YUV));
}
OCL_TEST_P(CvtColor, BGR2YUV)
{
doTest(3, 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, RGB2YCrCb)
{
doTest(3, 3, CVTCODE(RGB2YCrCb));
}
OCL_TEST_P(CvtColor, BGR2YCrCb)
{
doTest(3, 3, CVTCODE(BGR2YCrCb));
}
struct CvtColor_YUV420 : CvtColor
#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)); }
// 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)); }
// 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)); }
// 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, 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)); }
// 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)); }
// 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 :
public CvtColor
{
void random_roi(int channelsIn, int channelsOut)
{
@ -190,50 +227,33 @@ struct CvtColor_YUV420 : CvtColor
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);
}
};
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, 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, 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)),
Bool()
)
);
Bool()));
INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor_YUV420,
testing::Combine(
testing::Values(MatDepth(CV_8U)),
Bool()
)
);
Bool()));
}
#endif

Write Preview
Loading…
Cancel
Save