Skip to content

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

RGB[A] <-> HSV

parent 727a5e6d
Hide whitespace changes
Inline Side-by-side
Showing with 329 additions and 87 deletions
modules/core/include/opencv2/core/ocl.hpp
View file @ 0b900b54
......@@ -245,11 +245,13 @@ protected:
class CV_EXPORTS KernelArg
{
public:
enum { LOCAL=1, READ_ONLY=2, WRITE_ONLY=4, READ_WRITE=6, CONSTANT=8, NO_SIZE=256 };
enum { LOCAL=1, READ_ONLY=2, WRITE_ONLY=4, READ_WRITE=6, CONSTANT=8, PTR_ONLY = 16, NO_SIZE=256 };
KernelArg(int _flags, UMat* _m, int wscale=1, const void* _obj=0, size_t _sz=0);
KernelArg();
static KernelArg Local() { return KernelArg(LOCAL, 0); }
static KernelArg PtrOnly(const UMat & m)
{ return KernelArg(PTR_ONLY, (UMat*)&m); }
static KernelArg ReadWrite(const UMat& m, int wscale=1)
{ return KernelArg(READ_WRITE, (UMat*)&m, wscale); }
static KernelArg ReadWriteNoSize(const UMat& m, int wscale=1)
......
modules/core/src/ocl.cpp
View file @ 0b900b54
......@@ -41,6 +41,7 @@
#include "precomp.hpp"
#include <map>
#include <iostream>
/*
Part of the file is an extract from the standard OpenCL headers from Khronos site.
......@@ -2205,9 +2206,12 @@ int Kernel::set(int i, const KernelArg& arg)
{
int accessFlags = ((arg.flags & KernelArg::READ_ONLY) ? ACCESS_READ : 0) +
((arg.flags & KernelArg::WRITE_ONLY) ? ACCESS_WRITE : 0);
bool ptronly = (arg.flags & KernelArg::PTR_ONLY) != 0;
cl_mem h = (cl_mem)arg.m->handle(accessFlags);
if( arg.m->dims <= 2 )
if (ptronly)
clSetKernelArg(p->handle, (cl_uint)i++, sizeof(h), &h);
else if( arg.m->dims <= 2 )
{
UMat2D u2d(*arg.m);
clSetKernelArg(p->handle, (cl_uint)i, sizeof(h), &h);
......@@ -2350,7 +2354,7 @@ struct Program::Impl
retval = clBuildProgram(handle, n,
(const cl_device_id*)deviceList,
buildflags.c_str(), 0, 0);
if( retval == CL_BUILD_PROGRAM_FAILURE )
if( retval < 0 /*== CL_BUILD_PROGRAM_FAILURE*/ )
{
char buf[1<<16];
size_t retsz = 0;
......
modules/imgproc/src/color.cpp
View file @ 0b900b54
......@@ -2876,7 +2876,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
k.create("RGB2XYZ", ocl::imgproc::cvtcolor_oclsrc,
format("-D depth=%d -D scn=%d -D dcn=3 -D bidx=%d", depth, scn, bidx));
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::ReadOnlyNoSize(c));
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrOnly(c));
return k.run(2, globalsize, 0, false);
}
case COLOR_XYZ2BGR: case COLOR_XYZ2RGB:
......@@ -2925,19 +2925,91 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
k.create("XYZ2RGB", ocl::imgproc::cvtcolor_oclsrc,
format("-D depth=%d -D scn=3 -D dcn=%d -D bidx=%d", depth, dcn, bidx));
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::ReadOnlyNoSize(c));
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrOnly(c));
return k.run(2, globalsize, 0, false);
}
/*
case COLOR_BGR5652GRAY: case COLOR_BGR5552GRAY:
case COLOR_GRAY2BGR565: case COLOR_GRAY2BGR555:
case COLOR_BGR2XYZ: case COLOR_RGB2XYZ:
case COLOR_XYZ2BGR: case COLOR_XYZ2RGB:
case COLOR_BGR2HSV: case COLOR_RGB2HSV: case COLOR_BGR2HSV_FULL: case COLOR_RGB2HSV_FULL:
case COLOR_BGR2HLS: case COLOR_RGB2HLS: case COLOR_BGR2HLS_FULL: case COLOR_RGB2HLS_FULL:
case COLOR_HSV2BGR: case COLOR_HSV2RGB: case COLOR_HSV2BGR_FULL: case COLOR_HSV2RGB_FULL:
case COLOR_HLS2BGR: case COLOR_HLS2RGB: case COLOR_HLS2BGR_FULL: case COLOR_HLS2RGB_FULL:
*/
case COLOR_BGR2HSV: case COLOR_RGB2HSV: case COLOR_BGR2HSV_FULL: case COLOR_RGB2HSV_FULL:
case COLOR_BGR2HLS: case COLOR_RGB2HLS: case COLOR_BGR2HLS_FULL: case COLOR_RGB2HLS_FULL:
{
CV_Assert((scn == 3 || scn == 4) && (depth == CV_8U || depth == CV_32F));
bidx = code == COLOR_BGR2HSV || code == COLOR_BGR2HLS ||
code == COLOR_BGR2HSV_FULL || code == COLOR_BGR2HLS_FULL ? 0 : 2;
int hrange = depth == CV_32F ? 360 : code == COLOR_BGR2HSV || code == COLOR_RGB2HSV ||
code == COLOR_BGR2HLS || code == COLOR_RGB2HLS ? 180 : 256;
bool is_hsv = code == COLOR_BGR2HSV || code == COLOR_RGB2HSV || code == COLOR_BGR2HSV_FULL || code == COLOR_RGB2HSV_FULL;
String kernelName = String("RGB2") + (is_hsv ? "HSV" : "HLS");
dcn = 3;
if (is_hsv && depth == CV_8U)
{
static UMat sdiv_data;
static UMat hdiv_data180;
static UMat hdiv_data256;
static int sdiv_table[256];
static int hdiv_table180[256];
static int hdiv_table256[256];
static volatile bool initialized180 = false, initialized256 = false;
volatile bool & initialized = hrange == 180 ? initialized180 : initialized256;
if (!initialized)
{
int * const hdiv_table = hrange == 180 ? hdiv_table180 : hdiv_table256, hsv_shift = 12;
UMat & hdiv_data = hrange == 180 ? hdiv_data180 : hdiv_data256;
sdiv_table[0] = hdiv_table180[0] = hdiv_table256[0] = 0;
int v = 255 << hsv_shift;
if (!initialized180 && !initialized256)
{
for(int i = 1; i < 256; i++ )
sdiv_table[i] = saturate_cast<int>(v/(1.*i));
Mat(1, 256, CV_32SC1, sdiv_table).copyTo(sdiv_data);
}
v = hrange << hsv_shift;
for (int i = 1; i < 256; i++ )
hdiv_table[i] = saturate_cast<int>(v/(6.*i));
Mat(1, 256, CV_32SC1, hdiv_table).copyTo(hdiv_data);
initialized = true;
}
_dst.create(dstSz, CV_8UC3);
dst = _dst.getUMat();
k.create("RGB2HSV", ocl::imgproc::cvtcolor_oclsrc, format("-D depth=%d -D hrange=%d -D bidx=%d -D dcn=3 -D scn=%d",
depth, hrange, bidx, scn));
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst),
ocl::KernelArg::PtrOnly(sdiv_data), hrange == 256 ? ocl::KernelArg::PtrOnly(hdiv_data256) :
ocl::KernelArg::PtrOnly(hdiv_data180));
return k.run(2, globalsize, NULL, false);
}
else
k.create(kernelName.c_str(), ocl::imgproc::cvtcolor_oclsrc,
format("-D depth=%d -D hscale=%f -D bidx=%d -D scn=%d -D dcn=3", depth, hrange*(1.f/360.f), bidx, scn));
break;
}
case COLOR_HSV2BGR: case COLOR_HSV2RGB: case COLOR_HSV2BGR_FULL: case COLOR_HSV2RGB_FULL:
case COLOR_HLS2BGR: case COLOR_HLS2RGB: case COLOR_HLS2BGR_FULL: case COLOR_HLS2RGB_FULL:
{
if (dcn <= 0)
dcn = 3;
CV_Assert(scn == 3 && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));
bidx = code == COLOR_HSV2BGR || code == COLOR_HLS2BGR ||
code == COLOR_HSV2BGR_FULL || code == COLOR_HLS2BGR_FULL ? 0 : 2;
int hrange = depth == CV_32F ? 360 : code == COLOR_HSV2BGR || code == COLOR_HSV2RGB ||
code == COLOR_HLS2BGR || code == COLOR_HLS2RGB ? 180 : 255;
bool is_hsv = code == COLOR_HSV2BGR || code == COLOR_HSV2RGB ||
code == COLOR_HSV2BGR_FULL || code == COLOR_HSV2RGB_FULL;
String kernelName = String(is_hsv ? "HSV" : "HLS") + "2RGB";
k.create(kernelName.c_str(), ocl::imgproc::cvtcolor_oclsrc,
format("-D depth=%d -D dcn=%d -D scn=3 -D bidx=%d -D hrange=%d -D hscale=%f",
depth, dcn, bidx, hrange, 6.f/hrange));
break;
}
default:
;
}
......
modules/imgproc/src/opencl/cvtcolor.cl
View file @ 0b900b54
......@@ -81,6 +81,7 @@ enum
{
yuv_shift = 14,
xyz_shift = 12,
hsv_shift = 12,
R2Y = 4899,
G2Y = 9617,
B2Y = 1868,
......@@ -90,6 +91,14 @@ enum
#define scnbytes ((int)sizeof(DATA_TYPE)*scn)
#define dcnbytes ((int)sizeof(DATA_TYPE)*dcn)
#ifndef hscale
#define hscale 0
#endif
#ifndef hrange
#define hrange 0
#endif
///////////////////////////////////// RGB <-> GRAY //////////////////////////////////////
__kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset,
......@@ -352,7 +361,7 @@ __kernel void YCrCb2RGB(__global const uchar* src, int src_step, int src_offset,
__kernel void RGB2XYZ(__global const uchar * srcptr, int src_step, int src_offset,
__global uchar * dstptr, int dst_step, int dst_offset,
int rows, int cols, __constant COEFF_TYPE * coeffs, int a1, int a2)
int rows, int cols, __constant COEFF_TYPE * coeffs)
{
int dx = get_global_id(0);
int dy = get_global_id(1);
......@@ -384,7 +393,7 @@ __kernel void RGB2XYZ(__global const uchar * srcptr, int src_step, int src_offse
__kernel void XYZ2RGB(__global const uchar * srcptr, int src_step, int src_offset,
__global uchar * dstptr, int dst_step, int dst_offset,
int rows, int cols, __constant COEFF_TYPE * coeffs, int a1, int a2)
int rows, int cols, __constant COEFF_TYPE * coeffs)
{
int dx = get_global_id(0);
int dy = get_global_id(1);
......@@ -454,7 +463,6 @@ __kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset,
}
}
///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////
__kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset,
......@@ -562,7 +570,212 @@ __kernel void Gray2BGR5x5(__global const uchar* src, int src_step, int src_offse
}
}
//////////////////////////////////// RGB <-> HSV //////////////////////////////////////
__constant int sector_data[][3] = { { 1, 3, 0 },
{ 1, 0, 2 },
{ 3, 0, 1 },
{ 0, 2, 1 },
{ 0, 1, 3 },
{ 2, 1, 0 } };
#ifdef DEPTH_0
__kernel void RGB2HSV(__global const uchar* src, int src_step, int src_offset,
__global uchar* dst, int dst_step, int dst_offset,
int rows, int cols,
__constant int * sdiv_table, __constant int * hdiv_table)
{
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 * scnbytes);
int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
int b = src[src_idx + bidx], g = src[src_idx + 1], r = src[src_idx + (bidx^2)];
int h, s, v = b;
int vmin = b, diff;
int vr, vg;
v = max( v, g );
v = max( v, r );
vmin = min( vmin, g );
vmin = min( vmin, r );
diff = v - vmin;
vr = v == r ? -1 : 0;
vg = v == g ? -1 : 0;
s = (diff * sdiv_table[v] + (1 << (hsv_shift-1))) >> hsv_shift;
h = (vr & (g - b)) +
(~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff))));
h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift;
h += h < 0 ? hrange : 0;
dst[dst_idx] = convert_uchar_sat_rte(h);
dst[dst_idx + 1] = (uchar)s;
dst[dst_idx + 2] = (uchar)v;
}
}
__kernel void HSV2RGB(__global const uchar* src, int src_step, int src_offset,
__global uchar* dst, int dst_step, int dst_offset,
int rows, int cols)
{
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 * scnbytes);
int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
float h = src[src_idx], s = src[src_idx + 1]*(1/255.f), v = src[src_idx + 2]*(1/255.f);
float b, g, r;
if (s != 0)
{
float tab[4];
int sector;
h *= hscale;
if( h < 0 )
do h += 6; while( h < 0 );
else if( h >= 6 )
do h -= 6; while( h >= 6 );
sector = convert_int_sat_rtn(h);
h -= sector;
if( (unsigned)sector >= 6u )
{
sector = 0;
h = 0.f;
}
tab[0] = v;
tab[1] = v*(1.f - s);
tab[2] = v*(1.f - s*h);
tab[3] = v*(1.f - s*(1.f - h));
b = tab[sector_data[sector][0]];
g = tab[sector_data[sector][1]];
r = tab[sector_data[sector][2]];
}
else
b = g = r = v;
dst[dst_idx + bidx] = convert_uchar_sat_rte(b*255.f);
dst[dst_idx + 1] = convert_uchar_sat_rte(g*255.f);
dst[dst_idx + (bidx^2)] = convert_uchar_sat_rte(r*255.f);
#if dcn == 4
dst[dst_idx + 3] = MAX_NUM;
#endif
}
}
#elif defined DEPTH_5
__kernel void RGB2HSV(__global const uchar* srcptr, int src_step, int src_offset,
__global uchar* dstptr, int dst_step, int dst_offset,
int rows, int cols)
{
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 * scnbytes);
int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
__global const float * src = (__global const float *)(srcptr + src_idx);
__global float * dst = (__global float *)(dstptr + dst_idx);
float b = src[bidx], g = src[1], r = src[bidx^2];
float h, s, v;
float vmin, diff;
v = vmin = r;
if( v < g ) v = g;
if( v < b ) v = b;
if( vmin > g ) vmin = g;
if( vmin > b ) vmin = b;
diff = v - vmin;
s = diff/(float)(fabs(v) + FLT_EPSILON);
diff = (float)(60./(diff + FLT_EPSILON));
if( v == r )
h = (g - b)*diff;
else if( v == g )
h = (b - r)*diff + 120.f;
else
h = (r - g)*diff + 240.f;
if( h < 0 ) h += 360.f;
dst[0] = h*hscale;
dst[1] = s;
dst[2] = v;
}
}
__kernel void HSV2RGB(__global const uchar* srcptr, int src_step, int src_offset,
__global uchar* dstptr, int dst_step, int dst_offset,
int rows, int cols)
{
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 * scnbytes);
int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
__global const float * src = (__global const float *)(srcptr + src_idx);
__global float * dst = (__global float *)(dstptr + dst_idx);
float h = src[0], s = src[1], v = src[2];
float b, g, r;
if (s != 0)
{
float tab[4];
int sector;
h *= hscale;
if(h < 0)
do h += 6; while (h < 0);
else if (h >= 6)
do h -= 6; while (h >= 6);
sector = convert_int_sat_rtn(h);
h -= sector;
if ((unsigned)sector >= 6u)
{
sector = 0;
h = 0.f;
}
tab[0] = v;
tab[1] = v*(1.f - s);
tab[2] = v*(1.f - s*h);
tab[3] = v*(1.f - s*(1.f - h));
b = tab[sector_data[sector][0]];
g = tab[sector_data[sector][1]];
r = tab[sector_data[sector][2]];
}
else
b = g = r = v;
dst[bidx] = b;
dst[1] = g;
dst[bidx^2] = r;
#if dcn == 4
dst[3] = MAX_NUM;
#endif
}
}
#endif
......
modules/imgproc/test/ocl/test_color.cpp
View file @ 0b900b54
......@@ -171,25 +171,25 @@ OCL_TEST_P(CvtColor, XYZ2BGRA) { performTest(3, 4, CVTCODE(XYZ2BGR)); }
typedef CvtColor CvtColor8u32f;
//OCL_TEST_P(CvtColor8u32f, RGB2HSV) { performTest(3, 3, CVTCODE(RGB2HSV)); }
//OCL_TEST_P(CvtColor8u32f, BGR2HSV) { performTest(3, 3, CVTCODE(BGR2HSV)); }
//OCL_TEST_P(CvtColor8u32f, RGBA2HSV) { performTest(4, 3, CVTCODE(RGB2HSV)); }
//OCL_TEST_P(CvtColor8u32f, BGRA2HSV) { performTest(4, 3, CVTCODE(BGR2HSV)); }
//OCL_TEST_P(CvtColor8u32f, RGB2HSV_FULL) { performTest(3, 3, CVTCODE(RGB2HSV_FULL)); }
//OCL_TEST_P(CvtColor8u32f, BGR2HSV_FULL) { performTest(3, 3, CVTCODE(BGR2HSV_FULL)); }
//OCL_TEST_P(CvtColor8u32f, RGBA2HSV_FULL) { performTest(4, 3, CVTCODE(RGB2HSV_FULL)); }
//OCL_TEST_P(CvtColor8u32f, BGRA2HSV_FULL) { performTest(4, 3, CVTCODE(BGR2HSV_FULL)); }
//OCL_TEST_P(CvtColor8u32f, HSV2RGB) { performTest(3, 3, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); }
//OCL_TEST_P(CvtColor8u32f, HSV2BGR) { performTest(3, 3, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); }
//OCL_TEST_P(CvtColor8u32f, HSV2RGBA) { performTest(3, 4, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); }
//OCL_TEST_P(CvtColor8u32f, HSV2BGRA) { performTest(3, 4, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); }
//OCL_TEST_P(CvtColor8u32f, HSV2RGB_FULL) { performTest(3, 3, CVTCODE(HSV2RGB_FULL), depth == CV_8U ? 1 : 4e-1); }
//OCL_TEST_P(CvtColor8u32f, HSV2BGR_FULL) { performTest(3, 3, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); }
//OCL_TEST_P(CvtColor8u32f, HSV2RGBA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); }
//OCL_TEST_P(CvtColor8u32f, HSV2BGRA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); }
OCL_TEST_P(CvtColor8u32f, RGB2HSV) { performTest(3, 3, CVTCODE(RGB2HSV)); }
OCL_TEST_P(CvtColor8u32f, BGR2HSV) { performTest(3, 3, CVTCODE(BGR2HSV)); }
OCL_TEST_P(CvtColor8u32f, RGBA2HSV) { performTest(4, 3, CVTCODE(RGB2HSV)); }
OCL_TEST_P(CvtColor8u32f, BGRA2HSV) { performTest(4, 3, CVTCODE(BGR2HSV)); }
OCL_TEST_P(CvtColor8u32f, RGB2HSV_FULL) { performTest(3, 3, CVTCODE(RGB2HSV_FULL)); }
OCL_TEST_P(CvtColor8u32f, BGR2HSV_FULL) { performTest(3, 3, CVTCODE(BGR2HSV_FULL)); }
OCL_TEST_P(CvtColor8u32f, RGBA2HSV_FULL) { performTest(4, 3, CVTCODE(RGB2HSV_FULL)); }
OCL_TEST_P(CvtColor8u32f, BGRA2HSV_FULL) { performTest(4, 3, CVTCODE(BGR2HSV_FULL)); }
OCL_TEST_P(CvtColor8u32f, HSV2RGB) { performTest(3, 3, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); }
OCL_TEST_P(CvtColor8u32f, HSV2BGR) { performTest(3, 3, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); }
OCL_TEST_P(CvtColor8u32f, HSV2RGBA) { performTest(3, 4, CVTCODE(HSV2RGB), depth == CV_8U ? 1 : 4e-1); }
OCL_TEST_P(CvtColor8u32f, HSV2BGRA) { performTest(3, 4, CVTCODE(HSV2BGR), depth == CV_8U ? 1 : 4e-1); }
OCL_TEST_P(CvtColor8u32f, HSV2RGB_FULL) { performTest(3, 3, CVTCODE(HSV2RGB_FULL), depth == CV_8U ? 1 : 4e-1); }
OCL_TEST_P(CvtColor8u32f, HSV2BGR_FULL) { performTest(3, 3, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); }
OCL_TEST_P(CvtColor8u32f, HSV2RGBA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); }
OCL_TEST_P(CvtColor8u32f, HSV2BGRA_FULL) { performTest(3, 4, CVTCODE(HSV2BGR_FULL), depth == CV_8U ? 1 : 4e-1); }
// RGB <-> HLS
......
modules/ocl/src/opencl/cvt_color.cl
View file @ 0b900b54
......@@ -91,56 +91,7 @@ enum
BLOCK_SIZE = 256
};
///////////////////////////////////// RGB5x5 <-> Gray //////////////////////////////////////
__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
}
}
///////////////////////////////////// RGB <-> HSV //////////////////////////////////////
//////////////////////////////////// RGB <-> HSV //////////////////////////////////////
__constant int sector_data[][3] = { {1, 3, 0}, { 1, 0, 2 }, { 3, 0, 1 }, { 0, 2, 1 }, { 0, 1, 3 }, { 2, 1, 0 } };
......
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