Skip to content

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

added XYZ to RGB conversion to ocl::cvtColor

parent 5e02b204
Hide whitespace changes
Inline Side-by-side
Showing with 125 additions and 36 deletions
modules/ocl/src/color.cpp
View file @ 581a3e44
......@@ -76,10 +76,10 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::
openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str());
}
static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName)
static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName,
const oclMat & data = oclMat())
{
int channels = dst.channels();
std::string build_options = format("-D DEPTH_%d", src.depth());
std::string build_options = format("-D DEPTH_%d -D channels=%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();
......@@ -88,13 +88,15 @@ static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::st
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 ));
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());
}
......@@ -235,7 +237,7 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
}
pdata = coeffs;
}
oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32F : CV_32S, pdata);
oclMat oclCoeffs(1, 9, depth == CV_32F ? CV_32FC1 : CV_32SC1, pdata);
fromRGB_caller(src, dst, bidx, "RGB2XYZ", oclCoeffs);
break;
......@@ -248,7 +250,43 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
CV_Assert(scn == 3 && (dcn == 3 || dcn == 4));
bidx = code == CV_XYZ2BGR ? 0 : 2;
dst.create(sz, CV_MAKE_TYPE(depth, dcn));
toRGB_caller(src, dst, bidx, "XYZ2RGB");
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;
}
/*
......
modules/ocl/src/opencl/cvt_color.cl
View file @ 581a3e44
......@@ -46,7 +46,7 @@
/**************************************PUBLICFUNC*************************************/
#if defined (DEPTH_0)
#ifdef DEPTH_0
#define DATA_TYPE uchar
#define COEFF_TYPE int
#define MAX_NUM 255
......@@ -54,7 +54,7 @@
#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
......@@ -62,7 +62,7 @@
#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
......@@ -95,7 +95,7 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step,
{
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);
......@@ -103,7 +103,7 @@ __kernel void RGB2Gray(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, int bidx,
__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)
{
......@@ -119,8 +119,9 @@ __kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int chann
dst[dst_idx] = val;
dst[dst_idx + 1] = val;
dst[dst_idx + 2] = val;
if (channels == 4)
#if channels == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
......@@ -143,7 +144,7 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step,
int dst_idx = mad24(y, dst_step, dst_offset + x);
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;
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;
......@@ -165,7 +166,7 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step,
__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)
{
......@@ -179,7 +180,7 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, int channe
int dst_idx = mad24(y, dst_step, dst_offset + x);
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;
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];
......@@ -194,8 +195,9 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, int channe
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)
#if channels == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
......@@ -206,7 +208,7 @@ __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 channels,
__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)
{
......@@ -278,7 +280,7 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step,
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;
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;
......@@ -300,7 +302,7 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step,
__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,
__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)
{
......@@ -330,41 +332,73 @@ __kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step, int chan
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)
#if channels == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
///////////////////////////////////// RGB <-> XYZ //////////////////////////////////////
#pragma OPENCL EXTENSION cl_amd_printf:enable
__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, __global COEFF_TYPE * coeffs)
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst,
int src_offset, int dst_offset, __constant COEFF_TYPE * coeffs)
{
int x = get_global_id(0);
int y = get_global_id(1);
int dx = get_global_id(0);
int dy = get_global_id(1);
if (y < rows && x < cols)
if (dy < rows && dx < cols)
{
x <<= 2;
int src_idx = mad24(y, src_step, src_offset + x);
int dst_idx = mad24(y, dst_step, dst_offset + x);
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
DATA_TYPE x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2];
DATA_TYPE y = r * coeffs[3] + g * coeffs[4] + b * coeffs[5];
DATA_TYPE z = r * coeffs[6] + g * coeffs[7] + b * coeffs[8];
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
DATA_TYPE x = CV_DESCALE(r * coeffs[0] + g * coeffs[1] + b * coeffs[2], xyz_shift);
DATA_TYPE y = CV_DESCALE(r * coeffs[3] + g * coeffs[4] + b * coeffs[5], xyz_shift);
DATA_TYPE z = CV_DESCALE(r * coeffs[6] + g * coeffs[7] + b * coeffs[8], xyz_shift);
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 channels == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
modules/ocl/test/test_color.cpp
View file @ 581a3e44
......@@ -248,6 +248,23 @@ 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));
}
// 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