Commit b33a62be authored by Ilya Lavrenov's avatar Ilya Lavrenov

fixed separable filter extrapolation

parent f177e658
...@@ -1058,74 +1058,39 @@ template <> struct index_and_sizeof<float> ...@@ -1058,74 +1058,39 @@ template <> struct index_and_sizeof<float>
template <typename T> template <typename T>
void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel, int ksize, int anchor, int bordertype) void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel, int ksize, int anchor, int bordertype)
{ {
Context *clCxt = src.clCxt; CV_Assert(bordertype <= BORDER_REFLECT_101);
CV_Assert(ksize == (anchor << 1) + 1);
int channels = src.oclchannels(); int channels = src.oclchannels();
size_t localThreads[3] = {16, 16, 1}; size_t localThreads[3] = { 16, 16, 1 };
string kernelName = "row_filter"; size_t globalThreads[3] = { dst.cols, dst.rows, 1 };
char btype[30];
switch (bordertype) const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" };
{ std::string buildOptions = format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s",
case 0: anchor, (int)localThreads[0], (int)localThreads[1], channels, borderMap[bordertype]);
sprintf(btype, "BORDER_CONSTANT");
break;
case 1:
sprintf(btype, "BORDER_REPLICATE");
break;
case 2:
sprintf(btype, "BORDER_REFLECT");
break;
case 3:
sprintf(btype, "BORDER_WRAP");
break;
case 4:
sprintf(btype, "BORDER_REFLECT_101");
break;
}
char compile_option[128];
sprintf(compile_option, "-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s", anchor, (int)localThreads[0], (int)localThreads[1], channels, btype);
size_t globalThreads[3];
globalThreads[1] = (dst.rows + localThreads[1] - 1) / localThreads[1] * localThreads[1];
globalThreads[2] = (1 + localThreads[2] - 1) / localThreads[2] * localThreads[2];
if (src.depth() == CV_8U) if (src.depth() == CV_8U)
{ {
switch (channels) switch (channels)
{ {
case 1: case 1:
case 3: globalThreads[0] = (dst.cols + 3) >> 2;
globalThreads[0] = ((dst.cols + 4) / 4 + localThreads[0] - 1) / localThreads[0] * localThreads[0];
break; break;
case 2: case 2:
globalThreads[0] = ((dst.cols + 1) / 2 + localThreads[0] - 1) / localThreads[0] * localThreads[0]; globalThreads[0] = (dst.cols + 1) >> 1;
break; break;
case 4: case 4:
globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0]; globalThreads[0] = dst.cols;
break; break;
} }
} }
else
{
globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0];
}
//sanity checks int src_pix_per_row = src.step / src.elemSize();
CV_Assert(clCxt == dst.clCxt); int src_offset_x = (src.offset % src.step) / src.elemSize();
CV_Assert(src.cols == dst.cols); int src_offset_y = src.offset / src.step;
CV_Assert(src.oclchannels() == dst.oclchannels()); int dst_pix_per_row = dst.step / dst.elemSize();
CV_Assert(ksize == (anchor << 1) + 1);
int src_pix_per_row, dst_pix_per_row;
int src_offset_x, src_offset_y;//, dst_offset_in_pixel;
src_pix_per_row = src.step / src.elemSize();
src_offset_x = (src.offset % src.step) / src.elemSize();
src_offset_y = src.offset / src.step;
dst_pix_per_row = dst.step / dst.elemSize();
//dst_offset_in_pixel = dst.offset / dst.elemSize();
int ridusy = (dst.rows - src.rows) >> 1; int ridusy = (dst.rows - src.rows) >> 1;
vector<pair<size_t , const void *> > args; vector<pair<size_t , const void *> > args;
args.push_back(make_pair(sizeof(cl_mem), &src.data)); args.push_back(make_pair(sizeof(cl_mem), &src.data));
args.push_back(make_pair(sizeof(cl_mem), &dst.data)); args.push_back(make_pair(sizeof(cl_mem), &dst.data));
...@@ -1140,7 +1105,8 @@ void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel ...@@ -1140,7 +1105,8 @@ void linearRowFilter_gpu(const oclMat &src, const oclMat &dst, oclMat mat_kernel
args.push_back(make_pair(sizeof(cl_int), (void *)&ridusy)); args.push_back(make_pair(sizeof(cl_int), (void *)&ridusy));
args.push_back(make_pair(sizeof(cl_mem), (void *)&mat_kernel.data)); args.push_back(make_pair(sizeof(cl_mem), (void *)&mat_kernel.data));
openCLExecuteKernel(clCxt, &filter_sep_row, kernelName, globalThreads, localThreads, args, channels, src.depth(), compile_option); openCLExecuteKernel(src.clCxt, &filter_sep_row, "row_filter", globalThreads, localThreads,
args, channels, src.depth(), buildOptions.c_str());
} }
Ptr<BaseRowFilter_GPU> cv::ocl::getLinearRowFilter_GPU(int srcType, int /*bufType*/, const Mat &rowKernel, int anchor, int bordertype) Ptr<BaseRowFilter_GPU> cv::ocl::getLinearRowFilter_GPU(int srcType, int /*bufType*/, const Mat &rowKernel, int anchor, int bordertype)
......
...@@ -47,36 +47,6 @@ ...@@ -47,36 +47,6 @@
#define READ_TIMES_ROW ((2*(RADIUS+LSIZE0)-1)/LSIZE0) #define READ_TIMES_ROW ((2*(RADIUS+LSIZE0)-1)/LSIZE0)
#endif #endif
#ifdef BORDER_CONSTANT
//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii
#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
#endif
#ifdef BORDER_REPLICATE
//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
#define ADDR_L(i,l_edge,r_edge) (i) < (l_edge) ? (l_edge) : (i)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (r_edge)-1 : (addr)
#endif
#ifdef BORDER_REFLECT
//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
#define ADDR_L(i,l_edge,r_edge) (i) < (l_edge) ? -(i)-1 : (i)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)
#endif
#ifdef BORDER_REFLECT_101
//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
#define ADDR_L(i,l_edge,r_edge) (i) < (l_edge) ? -(i) : (i)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)
#endif
#ifdef BORDER_WRAP
//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
#define ADDR_L(i,l_edge,r_edge) (i) < (l_edge) ? (i)+(r_edge) : (i)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (i)-(r_edge) : (addr)
#endif
/********************************************************************************** /**********************************************************************************
These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur. These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur.
Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle
...@@ -107,15 +77,16 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter ...@@ -107,15 +77,16 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter
{ {
int x = get_global_id(0); int x = get_global_id(0);
int y = get_global_id(1); int y = get_global_id(1);
int l_x = get_local_id(0); int l_x = get_local_id(0);
int l_y = get_local_id(1); int l_y = get_local_id(1);
int start_addr = mad24(y,src_step_in_pixel,x);
int end_addr = mad24(src_whole_rows - 1,src_step_in_pixel,src_whole_cols);
int i;
GENTYPE_SRC sum;
GENTYPE_SRC temp[READ_TIMES_COL];
__local GENTYPE_SRC LDS_DAT[LSIZE1*READ_TIMES_COL][LSIZE0+1]; int start_addr = mad24(y, src_step_in_pixel, x);
int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols);
int i;
GENTYPE_SRC sum, temp[READ_TIMES_COL];
__local GENTYPE_SRC LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1];
//read pixels from src //read pixels from src
for(i = 0;i<READ_TIMES_COL;i++) for(i = 0;i<READ_TIMES_COL;i++)
......
This diff is collapsed.
...@@ -403,7 +403,7 @@ INSTANTIATE_TEST_CASE_P(Filter, SobelTest, Combine( ...@@ -403,7 +403,7 @@ INSTANTIATE_TEST_CASE_P(Filter, SobelTest, Combine(
Bool())); Bool()));
INSTANTIATE_TEST_CASE_P(Filter, ScharrTest, Combine( INSTANTIATE_TEST_CASE_P(Filter, ScharrTest, Combine(
Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC4), Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4),
Values(0), // not used Values(0), // not used
Values(Size(0, 1), Size(1, 0)), Values(Size(0, 1), Size(1, 0)),
Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101, Values((int)BORDER_CONSTANT, (int)BORDER_REFLECT101,
......
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