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

added RGB[A] <-> BGR[A] conversion to ocl::cvtColor

Browse Source
pull/1775/head
Ilya Lavrenov 11 years ago
parent 581a3e444d
commit 3cc9502c90
3 changed files with 136 additions and 175 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. 59
      modules/ocl/src/color.cpp
  2. 51
      modules/ocl/src/opencl/cvt_color.cl
  3. 201
      modules/ocl/test/test_color.cpp

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

@ -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<pair<size_t , const void *> > 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 );

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

@ -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
}
}

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

@ -121,149 +121,65 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool)
#define CVTCODE(name) COLOR_ ## name
// RGB <-> Gray
OCL_TEST_P(CvtColor, RGB2GRAY)
{
doTest(3, 1, CVTCODE(RGB2GRAY));
}
OCL_TEST_P(CvtColor, GRAY2RGB)
{
doTest(1, 3, CVTCODE(GRAY2RGB));
}
// 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)); }
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));
}
// RGB <-> Gray
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(

Write Preview
Loading…
Cancel
Save