mirror of
https://github.com/opencv/opencv.git
synced 2024-11-25 11:40:44 +08:00
Merge pull request #428 from bitwangyaoyao:2.4_erode_dilate
This commit is contained in:
commit
3ed6c09485
@ -19,6 +19,7 @@
|
||||
// Jia Haipeng, jiahaipeng95@gmail.com
|
||||
// Zero Lin, Zero.Lin@amd.com
|
||||
// Zhang Ying, zhangying913@gmail.com
|
||||
// Yao Wang, bitwangyaoyao@gmail.com
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
@ -309,21 +310,22 @@ public:
|
||||
|
||||
namespace
|
||||
{
|
||||
typedef void (*GPUMorfFilter_t)(const oclMat & , oclMat & , oclMat & , Size &, const Point);
|
||||
typedef void (*GPUMorfFilter_t)(const oclMat & , oclMat & , oclMat & , Size &, const Point, bool rectKernel, bool usrROI);
|
||||
|
||||
class MorphFilter_GPU : public BaseFilter_GPU
|
||||
{
|
||||
public:
|
||||
MorphFilter_GPU(const Size &ksize_, const Point &anchor_, const oclMat &kernel_, GPUMorfFilter_t func_) :
|
||||
BaseFilter_GPU(ksize_, anchor_, BORDER_CONSTANT), kernel(kernel_), func(func_) {}
|
||||
BaseFilter_GPU(ksize_, anchor_, BORDER_CONSTANT), kernel(kernel_), func(func_), rectKernel(false) {}
|
||||
|
||||
virtual void operator()(const oclMat &src, oclMat &dst)
|
||||
{
|
||||
func(src, dst, kernel, ksize, anchor) ;
|
||||
func(src, dst, kernel, ksize, anchor, rectKernel, false) ;
|
||||
}
|
||||
|
||||
oclMat kernel;
|
||||
GPUMorfFilter_t func;
|
||||
bool rectKernel;
|
||||
};
|
||||
}
|
||||
|
||||
@ -332,7 +334,8 @@ public:
|
||||
**Extend this if necessary later.
|
||||
**Note that the kernel need to be further refined.
|
||||
*/
|
||||
static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel, Size &ksize, const Point anchor)
|
||||
static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
|
||||
Size &ksize, const Point anchor, bool rectKernel, bool useROI)
|
||||
{
|
||||
//Normalize the result by default
|
||||
//float alpha = ksize.height * ksize.width;
|
||||
@ -388,7 +391,11 @@ static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel, Size &k
|
||||
}
|
||||
|
||||
char compile_option[128];
|
||||
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D ERODE %s", anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], s);
|
||||
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D ERODE %s %s %s",
|
||||
anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1],
|
||||
rectKernel?"-D RECTKERNEL":"",
|
||||
useROI?"-D USEROI":"",
|
||||
s);
|
||||
vector< pair<size_t, const void *> > args;
|
||||
args.push_back(make_pair(sizeof(cl_mem), (void *)&src.data));
|
||||
args.push_back(make_pair(sizeof(cl_mem), (void *)&dst.data));
|
||||
@ -407,7 +414,8 @@ static void GPUErode(const oclMat &src, oclMat &dst, oclMat &mat_kernel, Size &k
|
||||
|
||||
|
||||
//! data type supported: CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4
|
||||
static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel, Size &ksize, const Point anchor)
|
||||
static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel,
|
||||
Size &ksize, const Point anchor, bool rectKernel, bool useROI)
|
||||
{
|
||||
//Normalize the result by default
|
||||
//float alpha = ksize.height * ksize.width;
|
||||
@ -426,12 +434,13 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel, Size &
|
||||
Context *clCxt = src.clCxt;
|
||||
string kernelName;
|
||||
size_t localThreads[3] = {16, 16, 1};
|
||||
size_t globalThreads[3] = {(src.cols + localThreads[0]) / localThreads[0] *localThreads[0], (src.rows + localThreads[1]) / localThreads[1] *localThreads[1], 1};
|
||||
size_t globalThreads[3] = {(src.cols + localThreads[0] - 1) / localThreads[0] *localThreads[0],
|
||||
(src.rows + localThreads[1] - 1) / localThreads[1] *localThreads[1], 1};
|
||||
|
||||
if (src.type() == CV_8UC1)
|
||||
{
|
||||
kernelName = "morph_C1_D0";
|
||||
globalThreads[0] = ((src.cols + 3) / 4 + localThreads[0]) / localThreads[0] * localThreads[0];
|
||||
globalThreads[0] = ((src.cols + 3) / 4 + localThreads[0] - 1) / localThreads[0] * localThreads[0];
|
||||
CV_Assert(localThreads[0]*localThreads[1] * 8 >= (localThreads[0] * 4 + ksize.width - 1) * (localThreads[1] + ksize.height - 1));
|
||||
}
|
||||
else
|
||||
@ -463,7 +472,11 @@ static void GPUDilate(const oclMat &src, oclMat &dst, oclMat &mat_kernel, Size &
|
||||
}
|
||||
|
||||
char compile_option[128];
|
||||
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D DILATE %s", anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], s);
|
||||
sprintf(compile_option, "-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D DILATE %s %s %s",
|
||||
anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1],
|
||||
rectKernel?"-D RECTKERNEL":"",
|
||||
useROI?"-D USEROI":"",
|
||||
s);
|
||||
vector< pair<size_t, const void *> > args;
|
||||
args.push_back(make_pair(sizeof(cl_mem), (void *)&src.data));
|
||||
args.push_back(make_pair(sizeof(cl_mem), (void *)&dst.data));
|
||||
@ -495,7 +508,14 @@ Ptr<BaseFilter_GPU> cv::ocl::getMorphologyFilter_GPU(int op, int type, const Mat
|
||||
normalizeKernel(kernel, gpu_krnl);
|
||||
normalizeAnchor(anchor, ksize);
|
||||
|
||||
return Ptr<BaseFilter_GPU>(new MorphFilter_GPU(ksize, anchor, gpu_krnl, GPUMorfFilter_callers[op][CV_MAT_CN(type)]));
|
||||
bool noZero = true;
|
||||
for(int i = 0; i < kernel.rows * kernel.cols; ++i)
|
||||
if(kernel.data[i] != 1)
|
||||
noZero = false;
|
||||
MorphFilter_GPU* mfgpu=new MorphFilter_GPU(ksize, anchor, gpu_krnl, GPUMorfFilter_callers[op][CV_MAT_CN(type)]);
|
||||
if(noZero)
|
||||
mfgpu->rectKernel = true;
|
||||
return Ptr<BaseFilter_GPU>(mfgpu);
|
||||
}
|
||||
|
||||
namespace
|
||||
|
@ -8,6 +8,7 @@
|
||||
// @Authors
|
||||
// Niko Li, newlife20080214@gmail.com
|
||||
// Zero Lin, zero.lin@amd.com
|
||||
// Yao Wang, bitwangyaoyao@gmail.com
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
@ -100,14 +101,26 @@ __kernel void morph_C1_D0(__global const uchar * restrict src,
|
||||
LDS_DAT[point2] = temp1;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
uchar4 res = (uchar4)VAL;
|
||||
for(int i=0;i<2*RADIUSY+1;i++)
|
||||
for(int j=0;j<2*RADIUSX+1;j++)
|
||||
|
||||
for(int i=0; i<2*RADIUSY+1; i++)
|
||||
for(int j=0; j<2*RADIUSX+1; j++)
|
||||
{
|
||||
res =mat_kernel[i*(2*RADIUSX+1)+j]? MORPH_OP(res,vload4(0,(__local uchar*)&LDS_DAT[mad24((l_y+i),width,l_x)]+offset+j)):res;
|
||||
res =
|
||||
#ifndef RECTKERNEL
|
||||
mat_kernel[i*(2*RADIUSX+1)+j] ?
|
||||
#endif
|
||||
MORPH_OP(res,vload4(0,(__local uchar*)&LDS_DAT[mad24((l_y+i),width,l_x)]+offset+j))
|
||||
#ifndef RECTKERNEL
|
||||
:res
|
||||
#endif
|
||||
;
|
||||
}
|
||||
|
||||
int gidx = get_global_id(0)<<2;
|
||||
int gidy = get_global_id(1);
|
||||
int out_addr = mad24(gidy,dst_step_in_pixel,gidx+dst_offset_in_pixel);
|
||||
|
||||
#ifdef USEROI
|
||||
if(gidx+3<cols && gidy<rows && (dst_offset_in_pixel&3==0))
|
||||
{
|
||||
*(__global uchar4*)&dst[out_addr] = res;
|
||||
@ -137,16 +150,19 @@ __kernel void morph_C1_D0(__global const uchar * restrict src,
|
||||
dst[out_addr] = res.x;
|
||||
}
|
||||
}
|
||||
#else
|
||||
*(__global uchar4*)&dst[out_addr] = res;
|
||||
#endif
|
||||
}
|
||||
#else
|
||||
__kernel void morph(__global const GENTYPE * restrict src,
|
||||
__global GENTYPE *dst,
|
||||
int src_offset_x, int src_offset_y,
|
||||
int cols, int rows,
|
||||
int src_step_in_pixel, int dst_step_in_pixel,
|
||||
__constant uchar * mat_kernel,
|
||||
int src_whole_cols, int src_whole_rows,
|
||||
int dst_offset_in_pixel)
|
||||
__global GENTYPE *dst,
|
||||
int src_offset_x, int src_offset_y,
|
||||
int cols, int rows,
|
||||
int src_step_in_pixel, int dst_step_in_pixel,
|
||||
__constant uchar * mat_kernel,
|
||||
int src_whole_cols, int src_whole_rows,
|
||||
int dst_offset_in_pixel)
|
||||
{
|
||||
int l_x = get_local_id(0);
|
||||
int l_y = get_local_id(1);
|
||||
@ -154,7 +170,7 @@ __kernel void morph(__global const GENTYPE * restrict src,
|
||||
int y = get_group_id(1)*LSIZE1;
|
||||
int start_x = x+src_offset_x-RADIUSX;
|
||||
int end_x = x + src_offset_x+LSIZE0+RADIUSX;
|
||||
int width = end_x -start_x+1;
|
||||
int width = end_x -(x+src_offset_x-RADIUSX)+1;
|
||||
int start_y = y+src_offset_y-RADIUSY;
|
||||
int point1 = mad24(l_y,LSIZE0,l_x);
|
||||
int point2 = point1 + LSIZE0*LSIZE1;
|
||||
@ -188,10 +204,18 @@ __kernel void morph(__global const GENTYPE * restrict src,
|
||||
LDS_DAT[point2] = temp1;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
GENTYPE res = (GENTYPE)VAL;
|
||||
for(int i=0;i<2*RADIUSY+1;i++)
|
||||
for(int j=0;j<2*RADIUSX+1;j++)
|
||||
for(int i=0; i<2*RADIUSY+1; i++)
|
||||
for(int j=0; j<2*RADIUSX+1; j++)
|
||||
{
|
||||
res =mat_kernel[i*(2*RADIUSX+1)+j]? MORPH_OP(res,LDS_DAT[mad24(l_y+i,width,l_x+j)]):res;
|
||||
res =
|
||||
#ifndef RECTKERNEL
|
||||
mat_kernel[i*(2*RADIUSX+1)+j] ?
|
||||
#endif
|
||||
MORPH_OP(res,LDS_DAT[mad24(l_y+i,width,l_x+j)])
|
||||
#ifndef RECTKERNEL
|
||||
:res
|
||||
#endif
|
||||
;
|
||||
}
|
||||
int gidx = get_global_id(0);
|
||||
int gidy = get_global_id(1);
|
||||
|
@ -365,10 +365,10 @@ TEST_P(Laplacian, Accuracy)
|
||||
/////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
// erode & dilate
|
||||
|
||||
PARAM_TEST_CASE(ErodeDilateBase, MatType, bool)
|
||||
PARAM_TEST_CASE(ErodeDilateBase, MatType, int)
|
||||
{
|
||||
int type;
|
||||
//int iterations;
|
||||
int iterations;
|
||||
|
||||
//erode or dilate kernel
|
||||
cv::Mat kernel;
|
||||
@ -399,7 +399,7 @@ PARAM_TEST_CASE(ErodeDilateBase, MatType, bool)
|
||||
virtual void SetUp()
|
||||
{
|
||||
type = GET_PARAM(0);
|
||||
// iterations = GET_PARAM(1);
|
||||
iterations = GET_PARAM(1);
|
||||
|
||||
cv::RNG &rng = TS::ptr()->get_rng();
|
||||
cv::Size size(MWIDTH, MHEIGHT);
|
||||
@ -409,10 +409,6 @@ PARAM_TEST_CASE(ErodeDilateBase, MatType, bool)
|
||||
// rng.fill(kernel, cv::RNG::UNIFORM, cv::Scalar::all(0), cv::Scalar::all(3));
|
||||
kernel = randomMat(rng, Size(3, 3), CV_8UC1, 0, 3, false);
|
||||
|
||||
//int devnums = getDevice(oclinfo, OPENCV_DEFAULT_OPENCL_DEVICE);
|
||||
//CV_Assert(devnums > 0);
|
||||
////if you want to use undefault device, set it here
|
||||
////setDevice(oclinfo[0]);
|
||||
}
|
||||
|
||||
void random_roi()
|
||||
@ -456,12 +452,9 @@ TEST_P(Erode, Mat)
|
||||
for(int j = 0; j < LOOP_TIMES; j++)
|
||||
{
|
||||
random_roi();
|
||||
//int iterations =3;
|
||||
//cv::erode(mat1_roi, dst_roi, kernel, Point(-1, -1), iterations);
|
||||
//cv::ocl::erode(gmat1, gdst, kernel, Point(-1, -1), iterations);
|
||||
|
||||
cv::erode(mat1_roi, dst_roi, kernel);
|
||||
cv::ocl::erode(gmat1, gdst, kernel);
|
||||
cv::erode(mat1_roi, dst_roi, kernel, Point(-1, -1), iterations);
|
||||
cv::ocl::erode(gmat1, gdst, kernel, Point(-1, -1), iterations);
|
||||
|
||||
cv::Mat cpu_dst;
|
||||
gdst_whole.download(cpu_dst);
|
||||
@ -486,12 +479,8 @@ TEST_P(Dilate, Mat)
|
||||
for(int j = 0; j < LOOP_TIMES; j++)
|
||||
{
|
||||
random_roi();
|
||||
//int iterations =3;
|
||||
// cv::erode(mat1_roi, dst_roi, kernel, Point(-1, -1), iterations);
|
||||
// cv::ocl::erode(gmat1, gdst, kernel, Point(-1, -1), iterations);
|
||||
|
||||
cv::dilate(mat1_roi, dst_roi, kernel);
|
||||
cv::ocl::dilate(gmat1, gdst, kernel);
|
||||
cv::erode(mat1_roi, dst_roi, kernel, Point(-1, -1), iterations);
|
||||
cv::ocl::erode(gmat1, gdst, kernel, Point(-1, -1), iterations);
|
||||
|
||||
cv::Mat cpu_dst;
|
||||
gdst_whole.download(cpu_dst);
|
||||
@ -831,13 +820,13 @@ INSTANTIATE_TEST_CASE_P(Filters, Laplacian, Combine(
|
||||
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
|
||||
Values(1, 3)));
|
||||
|
||||
//INSTANTIATE_TEST_CASE_P(Filter, ErodeDilate, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(1, 2, 3)));
|
||||
INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(1)));
|
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(Values(CV_8UC1, CV_8UC1), Values(false)));
|
||||
//INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(Values(CV_8UC1, CV_8UC1), Values(false)));
|
||||
|
||||
//INSTANTIATE_TEST_CASE_P(Filter, ErodeDilate, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(1, 2, 3)));
|
||||
INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(1)));
|
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(Values(CV_8UC1, CV_8UC1), Values(false)));
|
||||
//INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(Values(CV_8UC1, CV_8UC1), Values(false)));
|
||||
|
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filter, Sobel, Combine(Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
|
||||
|
Loading…
Reference in New Issue
Block a user