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)
#include "precomp.hpp"
#include <map>
+#include <iostream>
/*
Part of the file is an extract from the standard OpenCL headers from Khronos site.
{
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);
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;
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:
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:
;
}
{
yuv_shift = 14,
xyz_shift = 12,
+ hsv_shift = 12,
R2Y = 4899,
G2Y = 9617,
B2Y = 1868,
#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,
__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);
__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);
}
}
-
///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////
__kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset,
}
}
+//////////////////////////////////// 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
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
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 } };
projects / platform / upstream / opencv.git / commitdiff