potok/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2026-07-31 08:13:04 +04:00
Code Packages Releases Wiki Activity

fixed many bugs related to Mat::getUMat(), asynchronous kernel execution etc. Also, played a bit with ocl::cvtColor vs cv::cvtColor performance

Browse Source
This commit is contained in:
Vadim Pisarevsky
2013-12-01 03:12:19 +04:00
parent daefe6983f
commit 6da5d21331
10 changed files with 401 additions and 208 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/imgproc/src/color.cpp
+5 -2
View File
@@ -2695,6 +2695,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
UMat src = _src.getUMat(), dst;
Size sz = src.size(), dstSz = sz;
int scn = src.channels(), depth = src.depth(), bidx;
int dims = 2, stripeSize = 32;
size_t globalsize[] = { src.cols, src.rows };
ocl::Kernel k;
@@ -2765,7 +2766,9 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
bidx = code == COLOR_BGR2GRAY || code == COLOR_BGRA2GRAY ? 0 : 2;
dcn = 1;
k.create("RGB2Gray", ocl::imgproc::cvtcolor_oclsrc,
format("-D depth=%d -D scn=%d -D dcn=1 -D bidx=%d", depth, scn, bidx));
format("-D depth=%d -D scn=%d -D dcn=1 -D bidx=%d -D STRIPE_SIZE=%d",
depth, scn, bidx, stripeSize));
globalsize[0] = (src.cols + stripeSize-1)/stripeSize;
break;
}
case COLOR_GRAY2BGR:
@@ -3027,7 +3030,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
_dst.create(dstSz, CV_MAKETYPE(depth, dcn));
dst = _dst.getUMat();
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst));
ok = k.run(2, globalsize, 0, false);
ok = k.run(dims, globalsize, 0, false);
}
return ok;
}
modules/imgproc/src/opencl/cvtcolor.cl
+24
View File
@@ -75,6 +75,10 @@
#error "invalid depth: should be 0 (CV_8U), 2 (CV_16U) or 5 (CV_32F)"
#endif
#ifndef STRIPE_SIZE
#define STRIPE_SIZE 1
#endif
#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
enum
@@ -105,6 +109,7 @@ __kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset,
__global uchar* dstptr, int dststep, int dstoffset,
int rows, int cols)
{
#if 0
const int x = get_global_id(0);
const int y = get_global_id(1);
@@ -118,6 +123,25 @@ __kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset,
dst[0] = (DATA_TYPE)CV_DESCALE((src[bidx] * B2Y + src[1] * G2Y + src[(bidx^2)] * R2Y), yuv_shift);
#endif
}
#else
const int x0 = get_global_id(0)*STRIPE_SIZE;
const int x1 = min(x0 + STRIPE_SIZE, cols);
const int y = get_global_id(1);
if( y < rows )
{
__global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset)) + x0*scn;
__global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset));
int x;
for( x = x0; x < x1; x++, src += scn )
#ifdef DEPTH_5
dst[x] = src[bidx] * 0.114f + src[1] * 0.587f + src[(bidx^2)] * 0.299f;
#else
dst[x] = (DATA_TYPE)(mad24(src[bidx], B2Y, mad24(src[1], G2Y,
mad24(src[(bidx^2)], R2Y, 1 << (yuv_shift-1)))) >> yuv_shift);
#endif
}
#endif
}
__kernel void Gray2RGB(__global const uchar* srcptr, int srcstep, int srcoffset,