Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

added 3-channels support to morphology operations

Browse Source
pull/2496/head
Ilya Lavrenov 11 years ago
parent b70332d806
commit d1cfcfcafd
3 changed files with 86 additions and 101 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 54
      modules/imgproc/src/morph.cpp
  2. 117
      modules/imgproc/src/opencl/morph.cl
  3. 16
      modules/imgproc/test/ocl/test_filters.cpp

54
modules/imgproc/src/morph.cpp
Unescape View File

@ -42,7 +42,6 @@
#include "precomp.hpp" #include "precomp.hpp"
#include <limits.h> #include <limits.h>
#include <stdio.h>
#include "opencl_kernels.hpp" #include "opencl_kernels.hpp"
/****************************************************************************************\ /****************************************************************************************\
@ -1291,9 +1290,10 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
{ {
CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE); CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE);
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
if (_src.depth() == CV_64F && !doubleSupport) if (depth == CV_64F && !doubleSupport)
return false; return false;
UMat kernel8U; UMat kernel8U;
@ -1324,13 +1324,14 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
return false; return false;
static const char * const op2str[] = { "ERODE", "DILATE" }; static const char * const op2str[] = { "ERODE", "DILATE" };
String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s -D GENTYPE=%s -D DEPTH_%d", String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s"
anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], op2str[op], " -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s", anchor.x, anchor.y,
(int)localThreads[0], (int)localThreads[1], op2str[op],
doubleSupport ? " -D DOUBLE_SUPPORT" : "", rectKernel ? " -D RECTKERNEL" : "", doubleSupport ? " -D DOUBLE_SUPPORT" : "", rectKernel ? " -D RECTKERNEL" : "",
ocl::typeToStr(_src.type()), _src.depth() ); ocl::typeToStr(_src.type()), _src.depth(), cn, ocl::typeToStr(depth));
std::vector<ocl::Kernel> kernels; std::vector<ocl::Kernel> kernels;
for (int i = 0; i<iterations; i++) for (int i = 0; i < iterations; i++)
{ {
ocl::Kernel k("morph", ocl::imgproc::morph_oclsrc, buildOptions); ocl::Kernel k("morph", ocl::imgproc::morph_oclsrc, buildOptions);
if (k.empty()) if (k.empty())
@ -1341,33 +1342,27 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
_dst.create(src.size(), src.type()); _dst.create(src.size(), src.type());
UMat dst = _dst.getUMat(); UMat dst = _dst.getUMat();
if( iterations== 1 && src.u != dst.u) if (iterations == 1 && src.u != dst.u)
{ {
Size wholesize; Size wholesize;
Point ofs; Point ofs;
src.locateROI(wholesize, ofs); src.locateROI(wholesize, ofs);
int wholecols = wholesize.width, wholerows = wholesize.height; int wholecols = wholesize.width, wholerows = wholesize.height;
int idxArg = 0; kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst),
idxArg = kernels[0].set(idxArg, ocl::KernelArg::ReadOnlyNoSize(src)); ofs.x, ofs.y, src.cols, src.rows, ocl::KernelArg::PtrReadOnly(kernel8U),
idxArg = kernels[0].set(idxArg, ocl::KernelArg::WriteOnlyNoSize(dst)); wholecols, wholerows);
idxArg = kernels[0].set(idxArg, ofs.x);
idxArg = kernels[0].set(idxArg, ofs.y);
idxArg = kernels[0].set(idxArg, src.cols);
idxArg = kernels[0].set(idxArg, src.rows);
idxArg = kernels[0].set(idxArg, ocl::KernelArg::PtrReadOnly(kernel8U));
idxArg = kernels[0].set(idxArg, wholecols);
idxArg = kernels[0].set(idxArg, wholerows);
return kernels[0].run(2, globalThreads, localThreads, false); return kernels[0].run(2, globalThreads, localThreads, false);
} }
for(int i = 0; i< iterations; i++) for (int i = 0; i < iterations; i++)
{ {
UMat source; UMat source;
Size wholesize; Size wholesize;
Point ofs; Point ofs;
if( i == 0)
if (i == 0)
{ {
int cols = src.cols, rows = src.rows; int cols = src.cols, rows = src.rows;
src.locateROI(wholesize,ofs); src.locateROI(wholesize,ofs);
@ -1385,20 +1380,11 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel,
dst.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x); dst.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x);
source.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x); source.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x);
} }
source.locateROI(wholesize, ofs); source.locateROI(wholesize, ofs);
int wholecols = wholesize.width, wholerows = wholesize.height;
int idxArg = 0; kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst),
idxArg = kernels[i].set(idxArg, ocl::KernelArg::ReadOnlyNoSize(source)); ofs.x, ofs.y, source.cols, source.rows, ocl::KernelArg::PtrReadOnly(kernel8U),
idxArg = kernels[i].set(idxArg, ocl::KernelArg::WriteOnlyNoSize(dst)); wholesize.width, wholesize.height);
idxArg = kernels[i].set(idxArg, ofs.x);
idxArg = kernels[i].set(idxArg, ofs.y);
idxArg = kernels[i].set(idxArg, source.cols);
idxArg = kernels[i].set(idxArg, source.rows);
idxArg = kernels[i].set(idxArg, ocl::KernelArg::PtrReadOnly(kernel8U));
idxArg = kernels[i].set(idxArg, wholecols);
idxArg = kernels[i].set(idxArg, wholerows);
if (!kernels[i].run(2, globalThreads, localThreads, false)) if (!kernels[i].run(2, globalThreads, localThreads, false))
return false; return false;
@ -1427,13 +1413,13 @@ static void morphOp( int op, InputArray _src, OutputArray _dst,
return; return;
#endif #endif
if( iterations == 0 || kernel.rows*kernel.cols == 1 ) if (iterations == 0 || kernel.rows*kernel.cols == 1)
{ {
_src.copyTo(_dst); _src.copyTo(_dst);
return; return;
} }
if( !kernel.data ) if (!kernel.data)
{ {
kernel = getStructuringElement(MORPH_RECT, Size(1+iterations*2,1+iterations*2)); kernel = getStructuringElement(MORPH_RECT, Size(1+iterations*2,1+iterations*2));
anchor = Point(iterations, iterations); anchor = Point(iterations, iterations);
@ -1450,7 +1436,7 @@ static void morphOp( int op, InputArray _src, OutputArray _dst,
} }
CV_OCL_RUN(_dst.isUMat() && _src.size() == _dst.size() && src_type == dst_type && CV_OCL_RUN(_dst.isUMat() && _src.size() == _dst.size() && src_type == dst_type &&
_src.dims() <= 2 && (src_cn == 1 || src_cn == 4) && _src.dims() <= 2 && src_cn <= 4 &&
(src_depth == CV_8U || src_depth == CV_32F || src_depth == CV_64F ) && (src_depth == CV_8U || src_depth == CV_32F || src_depth == CV_64F ) &&
borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() && borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() &&
(op == MORPH_ERODE || op == MORPH_DILATE), (op == MORPH_ERODE || op == MORPH_DILATE),

117
modules/imgproc/src/opencl/morph.cl
Unescape View File

@ -43,6 +43,16 @@
#endif #endif
#endif #endif
#if cn != 3
#define loadpix(addr) *(__global const T *)(addr)
#define storepix(val, addr) *(__global T *)(addr) = val
#define TSIZE (int)sizeof(T)
#else
#define loadpix(addr) vload3(0, (__global const T1 *)(addr))
#define storepix(val, addr) vstore3(val, 0, (__global T1 *)(addr))
#define TSIZE ((int)sizeof(T1)*3)
#endif
#ifdef DEPTH_0 #ifdef DEPTH_0
#ifdef ERODE #ifdef ERODE
#define VAL 255 #define VAL 255
@ -50,16 +60,14 @@
#ifdef DILATE #ifdef DILATE
#define VAL 0 #define VAL 0
#endif #endif
#endif #elif defined DEPTH_5
#ifdef DEPTH_5
#ifdef ERODE #ifdef ERODE
#define VAL FLT_MAX #define VAL FLT_MAX
#endif #endif
#ifdef DILATE #ifdef DILATE
#define VAL -FLT_MAX #define VAL -FLT_MAX
#endif #endif
#endif #elif defined DEPTH_6
#ifdef DEPTH_6
#ifdef ERODE #ifdef ERODE
#define VAL DBL_MAX #define VAL DBL_MAX
#endif #endif
@ -74,79 +82,70 @@
#ifdef DILATE #ifdef DILATE
#define MORPH_OP(A,B) max((A),(B)) #define MORPH_OP(A,B) max((A),(B))
#endif #endif
//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii
#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
__kernel void morph(__global const uchar * restrict srcptr, int src_step, int src_offset, // BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii
#define ELEM(i, l_edge, r_edge, elem1, elem2) (i) < (l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
__kernel void morph(__global const uchar * srcptr, int src_step, int src_offset,
__global uchar * dstptr, int dst_step, int dst_offset, __global uchar * dstptr, int dst_step, int dst_offset,
int src_offset_x, int src_offset_y, int src_offset_x, int src_offset_y, int cols, int rows,
int cols, int rows, __constant uchar * mat_kernel, int src_whole_cols, int src_whole_rows)
__constant uchar * mat_kernel,
int src_whole_cols, int src_whole_rows)
{ {
int l_x = get_local_id(0); int gidx = get_global_id(0), gidy = get_global_id(1);
int l_y = get_local_id(1); int l_x = get_local_id(0), l_y = get_local_id(1);
int x = get_group_id(0)*LSIZE0; int x = get_group_id(0) * LSIZE0, y = get_group_id(1) * LSIZE1;
int y = get_group_id(1)*LSIZE1; int start_x = x + src_offset_x - RADIUSX;
int start_x = x+src_offset_x-RADIUSX; int end_x = x + src_offset_x + LSIZE0 + RADIUSX;
int end_x = x + src_offset_x+LSIZE0+RADIUSX; int width = end_x - (x + src_offset_x - RADIUSX) + 1;
int width = end_x -(x+src_offset_x-RADIUSX)+1; int start_y = y + src_offset_y - RADIUSY;
int start_y = y+src_offset_y-RADIUSY; int point1 = mad24(l_y, LSIZE0, l_x);
int point1 = mad24(l_y,LSIZE0,l_x); int point2 = point1 + LSIZE0 * LSIZE1;
int point2 = point1 + LSIZE0*LSIZE1; int tl_x = point1 % width, tl_y = point1 / width;
int tl_x = point1 % width; int tl_x2 = point2 % width, tl_y2 = point2 / width;
int tl_y = point1 / width; int cur_x = start_x + tl_x, cur_y = start_y + tl_y;
int tl_x2 = point2 % width; int cur_x2 = start_x + tl_x2, cur_y2 = start_y + tl_y2;
int tl_y2 = point2 / width; int start_addr = mad24(cur_y, src_step, cur_x * TSIZE);
int cur_x = start_x + tl_x; int start_addr2 = mad24(cur_y2, src_step, cur_x2 * TSIZE);
int cur_y = start_y + tl_y;
int cur_x2 = start_x + tl_x2; __local T LDS_DAT[2*LSIZE1*LSIZE0];
int cur_y2 = start_y + tl_y2;
int start_addr = mad24(cur_y,src_step, cur_x*(int)sizeof(GENTYPE));
int start_addr2 = mad24(cur_y2,src_step, cur_x2*(int)sizeof(GENTYPE));
GENTYPE temp0,temp1;
__local GENTYPE LDS_DAT[2*LSIZE1*LSIZE0];
int end_addr = mad24(src_whole_rows - 1,src_step,src_whole_cols*(int)sizeof(GENTYPE)); // read pixels from src
//read pixels from src int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * TSIZE);
start_addr = ((start_addr < end_addr) && (start_addr > 0)) ? start_addr : 0; start_addr = start_addr < end_addr && start_addr > 0 ? start_addr : 0;
start_addr2 = ((start_addr2 < end_addr) && (start_addr2 > 0)) ? start_addr2 : 0; start_addr2 = start_addr2 < end_addr && start_addr2 > 0 ? start_addr2 : 0;
__global const GENTYPE * src;
src = (__global const GENTYPE *)(srcptr+start_addr);
temp0 = src[0];
src = (__global const GENTYPE *)(srcptr+start_addr2);
temp1 = src[0];
//judge if read out of boundary
temp0= ELEM(cur_x,0,src_whole_cols,(GENTYPE)VAL,temp0);
temp0= ELEM(cur_y,0,src_whole_rows,(GENTYPE)VAL,temp0);
temp1= ELEM(cur_x2,0,src_whole_cols,(GENTYPE)VAL,temp1); T temp0 = loadpix(srcptr + start_addr);
temp1= ELEM(cur_y2,0,src_whole_rows,(GENTYPE)VAL,temp1); T temp1 = loadpix(srcptr + start_addr2);
// judge if read out of boundary
temp0 = ELEM(cur_x, 0, src_whole_cols, (T)(VAL),temp0);
temp0 = ELEM(cur_y, 0, src_whole_rows, (T)(VAL),temp0);
temp1 = ELEM(cur_x2, 0, src_whole_cols, (T)(VAL), temp1);
temp1 = ELEM(cur_y2, 0, src_whole_rows, (T)(VAL), temp1);
LDS_DAT[point1] = temp0; LDS_DAT[point1] = temp0;
LDS_DAT[point2] = temp1; LDS_DAT[point2] = temp1;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
GENTYPE res = (GENTYPE)VAL;
for(int i=0; i<2*RADIUSY+1; i++) T res = (T)(VAL);
for(int j=0; j<2*RADIUSX+1; j++) for (int i = 0, sizey = 2 * RADIUSY + 1; i < sizey; i++)
for (int j = 0, sizex = 2 * RADIUSX + 1; j < sizex; j++)
{ {
res = res =
#ifndef RECTKERNEL #ifndef RECTKERNEL
mat_kernel[i*(2*RADIUSX+1)+j] ? mat_kernel[i*(2*RADIUSX+1)+j] ?
#endif #endif
MORPH_OP(res,LDS_DAT[mad24(l_y+i,width,l_x+j)]) MORPH_OP(res, LDS_DAT[mad24(l_y + i, width, l_x + j)])
#ifndef RECTKERNEL #ifndef RECTKERNEL
:res : res
#endif #endif
; ;
} }
int gidx = get_global_id(0);
int gidy = get_global_id(1); if (gidx < cols && gidy < rows)
if(gidx<cols && gidy<rows)
{ {
int dst_index = mad24(gidy, dst_step, dst_offset + gidx * (int)sizeof(GENTYPE)); int dst_index = mad24(gidy, dst_step, mad24(gidx, TSIZE, dst_offset));
__global GENTYPE * dst = (__global GENTYPE *)(dstptr + dst_index); storepix(res, dstptr + dst_index);
dst[0] = res;
} }
} }

16
modules/imgproc/test/ocl/test_filters.cpp
Unescape View File

@ -355,28 +355,28 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine(
Bool())); Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine( OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine(
Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4), Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
Values(3, 5, 7), Values(3, 5, 7),
Values(Size(0,0)),//not used Values(Size(0,0)),//not used
Values((BorderType)BORDER_CONSTANT),//not used Values((BorderType)BORDER_CONSTANT),//not used
Values(1.0, 2.0, 3.0), Values(1.0, 2.0, 3.0),
Bool() ) ); Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine( OCL_INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine(
Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4), Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
Values(3, 5, 7), Values(3, 5, 7),
Values(Size(0,0)),//not used Values(Size(0,0)),//not used
Values((BorderType)BORDER_CONSTANT),//not used Values((BorderType)BORDER_CONSTANT),//not used
Values(1.0, 2.0, 3.0), Values(1.0, 2.0, 3.0),
Bool() ) ); Bool()));
OCL_INSTANTIATE_TEST_CASE_P(Filter, MorphologyEx, Combine( OCL_INSTANTIATE_TEST_CASE_P(Filter, MorphologyEx, Combine(
Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4), Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4),
Values(3, 5, 7), Values(3, 5, 7),
Values(Size(0,0), Size(0,1), Size(0,2), Size(0,3), Size(0,4), Size(0,5),Size(0,6)),//uses as generator of operations Values(Size(0, 0), Size(0, 1), Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations
Values((BorderType)BORDER_CONSTANT),//not used Values((BorderType)BORDER_CONSTANT),// not used
Values(1.0, 2.0, 3.0), Values(1.0, 2.0, 3.0),
Bool() ) ); Bool()));
} } // namespace cvtest::ocl } } // namespace cvtest::ocl

Write Preview
Loading…
Cancel
Save