Skip to content

O
opencv
Commit eda6360f authored by Ilya Lavrenov's avatar Ilya Lavrenov
Browse files
Options

added RGB5x5 <-> Gray

parent 1f421fce
Hide whitespace changes
Inline Side-by-side
Showing with 80 additions and 7 deletions
modules/ocl/src/color.cpp
View file @ eda6360f
......@@ -122,12 +122,12 @@ static void RGB_caller(const oclMat &src, oclMat &dst, bool reverse)
openCLExecuteKernel(src.clCxt, &cvt_color, "RGB", gt, lt, args, -1, -1, build_options.c_str());
}
static void RGB5x52RGB_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits)
static void fromRGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits, const std::string & kernelName)
{
std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D dcn=%d",
src.depth(), greenbits, dst.channels());
int src_offset = src.offset >> 1, src_step = src.step >> 1;
int dst_offset = (int)dst.offset, dst_step = (int)dst.step;
int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step / dst.elemSize1();
vector<pair<size_t , const void *> > args;
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols));
......@@ -141,10 +141,10 @@ static void RGB5x52RGB_caller(const oclMat &src, oclMat &dst, int bidx, int gree
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset ));
size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 };
openCLExecuteKernel(src.clCxt, &cvt_color, "RGB5x52RGB", gt, lt, args, -1, -1, build_options.c_str());
openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str());
}
static void RGB2RGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits)
static void toRGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits, const std::string & kernelName)
{
std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D scn=%d",
src.depth(), greenbits, src.channels());
......@@ -163,7 +163,7 @@ static void RGB2RGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int gree
args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset ));
size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 };
openCLExecuteKernel(src.clCxt, &cvt_color, "RGB2RGB5x5", gt, lt, args, -1, -1, build_options.c_str());
openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str());
}
static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
......@@ -194,7 +194,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
int greenbits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565 ||
code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5;
dst.create(sz, CV_8UC2);
RGB2RGB5x5_caller(src, dst, bidx, greenbits);
toRGB5x5_caller(src, dst, bidx, greenbits, "RGB2RGB5x5");
break;
}
case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
......@@ -207,7 +207,23 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
int greenbits = code == CV_BGR5652BGR || code == CV_BGR5652RGB ||
code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5;
dst.create(sz, CV_MAKETYPE(depth, dcn));
RGB5x52RGB_caller(src, dst, bidx, greenbits);
fromRGB5x5_caller(src, dst, bidx, greenbits, "RGB5x52RGB");
break;
}
case CV_BGR5652GRAY: case CV_BGR5552GRAY:
{
CV_Assert(scn == 2 && depth == CV_8U);
dst.create(sz, CV_8UC1);
int greenbits = code == CV_BGR5652GRAY ? 6 : 5;
fromRGB5x5_caller(src, dst, -1, greenbits, "BGR5x52Gray");
break;
}
case CV_GRAY2BGR565: case CV_GRAY2BGR555:
{
CV_Assert(scn == 1 && depth == CV_8U);
dst.create(sz, CV_8UC2);
int greenbits = code == CV_GRAY2BGR565 ? 6 : 5;
toRGB5x5_caller(src, dst, -1, greenbits, "Gray2BGR5x5");
break;
}
case CV_RGB2GRAY: case CV_BGR2GRAY:
......
modules/ocl/src/opencl/cvt_color.cl
View file @ eda6360f
......@@ -495,3 +495,52 @@ __kernel void RGB2RGB5x5(int cols, int rows, int src_step, int dst_step, int bid
#endif
}
}
///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////
__kernel void BGR5x52Gray(int cols, int rows, int src_step, int dst_step, int bidx,
__global const ushort * src, __global uchar * dst,
int src_offset, int dst_offset)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (y < rows && x < cols)
{
int src_idx = mad24(y, src_step, src_offset + x);
int dst_idx = mad24(y, dst_step, dst_offset + x);
int t = src[src_idx];
#if greenbits == 6
dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y +
((t >> 3) & 0xfc)*G2Y +
((t >> 8) & 0xf8)*R2Y, yuv_shift);
#else
dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y +
((t >> 2) & 0xf8)*G2Y +
((t >> 7) & 0xf8)*R2Y, yuv_shift);
#endif
}
}
__kernel void Gray2BGR5x5(int cols, int rows, int src_step, int dst_step, int bidx,
__global const uchar * src, __global ushort * dst,
int src_offset, int dst_offset)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (y < rows && x < cols)
{
int src_idx = mad24(y, src_step, src_offset + x);
int dst_idx = mad24(y, dst_step, dst_offset + x);
int t = src[src_idx];
#if greenbits == 6
dst[dst_idx] = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8));
#else
t >>= 3;
dst[dst_idx] = (ushort)(t|(t << 5)|(t << 10));
#endif
}
}
modules/ocl/test/test_color.cpp
View file @ eda6360f
......@@ -205,6 +205,14 @@ OCL_TEST_P(CvtColor8u, RGB2BGR555) { doTest(3, 2, CVTCODE(RGB2BGR555)); }
OCL_TEST_P(CvtColor8u, BGRA2BGR555) { doTest(4, 2, CVTCODE(BGRA2BGR555)); }
OCL_TEST_P(CvtColor8u, RGBA2BGR555) { doTest(4, 2, CVTCODE(RGBA2BGR555)); }
// RGB5x5 <-> Gray
OCL_TEST_P(CvtColor8u, BGR5652GRAY) { doTest(2, 1, CVTCODE(BGR5652GRAY)); }
OCL_TEST_P(CvtColor8u, BGR5552GRAY) { doTest(2, 1, CVTCODE(BGR5552GRAY)); }
OCL_TEST_P(CvtColor8u, GRAY2BGR565) { doTest(1, 2, CVTCODE(GRAY2BGR565)); }
OCL_TEST_P(CvtColor8u, GRAY2BGR555) { doTest(1, 2, CVTCODE(GRAY2BGR555)); }
// YUV -> RGBA_NV12
struct CvtColor_YUV420 :
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment