Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

Removed storing of zeros in local memory

Browse Source
pull/2790/head
Alexander Karsakov 11 years ago
parent 8e54845049
commit 06fb5da7c8
2 changed files with 47 additions and 120 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 153
      modules/imgproc/src/opencl/pyr_up.cl
  2. 14
      modules/imgproc/src/pyramids.cpp

153
modules/imgproc/src/opencl/pyr_up.cl
Unescape View File

@ -78,28 +78,19 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in
const int x = get_global_id(0);
const int y = get_global_id(1);
const int lsizex = get_local_size(0);
const int lsizey = get_local_size(1);
const int tidx = get_local_id(0);
const int tidy = get_local_id(1);
__local FT s_srcPatch[10][10];
__local FT s_dstPatch[20][16];
__local FT s_srcPatch[LOCAL_SIZE/2 + 2][LOCAL_SIZE/2 + 2];
__local FT s_dstPatch[LOCAL_SIZE/2 + 2][LOCAL_SIZE];
__global uchar * dstData = dst + dst_offset;
__global const uchar * srcData = src + src_offset;
if( tidx < 10 && tidy < 10 )
if( tidx < (LOCAL_SIZE/2 + 2) && tidy < LOCAL_SIZE/2 + 2 )
{
int srcx = mad24((int)get_group_id(0), lsizex>>1, tidx) - 1;
int srcy = mad24((int)get_group_id(1), lsizey>>1, tidy) - 1;
srcx = abs(srcx);
srcx = min(src_cols - 1, srcx);
srcy = abs(srcy);
srcy = min(src_rows - 1, srcy);
int srcx = EXTRAPOLATE(mad24((int)get_group_id(0), LOCAL_SIZE/2, tidx) - 1, src_cols);
int srcy = EXTRAPOLATE(mad24((int)get_group_id(1), LOCAL_SIZE/2, tidy) - 1, src_rows);
s_srcPatch[tidy][tidx] = convertToFT(loadpix(srcData + srcy * src_step + srcx * PIXSIZE));
}
@ -107,63 +98,30 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in
barrier(CLK_LOCAL_MEM_FENCE);
FT sum = 0.f;
const FT evenFlag = (FT)((tidx & 1) == 0);
const FT oddFlag = (FT)((tidx & 1) != 0);
const bool eveny = ((tidy & 1) == 0);
const FT co1 = 0.75f;
const FT co2 = 0.5f;
const FT co3 = 0.125f;
if(eveny)
{
sum = ( evenFlag* co3 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx - 2) >> 1)];
sum = sum + ( oddFlag * co2 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx - 1) >> 1)];
sum = sum + ( evenFlag* co1 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx ) >> 1)];
sum = sum + ( oddFlag * co2 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx + 1) >> 1)];
sum = sum + ( evenFlag* co3 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx + 2) >> 1)];
}
s_dstPatch[2 + tidy][tidx] = sum;
const FT coef1 = (tidx & 1) == 0 ? co1 : (FT) 0;
const FT coef2 = (tidx & 1) == 0 ? co3 : co2;
const FT coefy1 = (tidy & 1) == 0 ? co1 : (FT) 0;
const FT coefy2 = (tidy & 1) == 0 ? co3 : co2;
if (tidy < 2)
if(tidy < LOCAL_SIZE/2 + 2)
{
sum = 0;
if (eveny)
{
sum = (evenFlag * co3 ) * s_srcPatch[lsizey-16][1 + ((tidx - 2) >> 1)];
sum = sum + ( oddFlag * co2 ) * s_srcPatch[lsizey-16][1 + ((tidx - 1) >> 1)];
sum = sum + (evenFlag * co1 ) * s_srcPatch[lsizey-16][1 + ((tidx ) >> 1)];
sum = sum + ( oddFlag * co2 ) * s_srcPatch[lsizey-16][1 + ((tidx + 1) >> 1)];
sum = sum + (evenFlag * co3 ) * s_srcPatch[lsizey-16][1 + ((tidx + 2) >> 1)];
}
sum = coef2* s_srcPatch[tidy][1 + ((tidx - 1) >> 1)];
sum = mad(coef1, s_srcPatch[tidy][1 + ((tidx ) >> 1)], sum);
sum = mad(coef2, s_srcPatch[tidy][1 + ((tidx + 2) >> 1)], sum);
s_dstPatch[tidy][tidx] = sum;
}
if (tidy > 13)
{
sum = 0;
if (eveny)
{
sum = (evenFlag * co3) * s_srcPatch[lsizey-7][1 + ((tidx - 2) >> 1)];
sum = sum + ( oddFlag * co2) * s_srcPatch[lsizey-7][1 + ((tidx - 1) >> 1)];
sum = sum + (evenFlag * co1) * s_srcPatch[lsizey-7][1 + ((tidx ) >> 1)];
sum = sum + ( oddFlag * co2) * s_srcPatch[lsizey-7][1 + ((tidx + 1) >> 1)];
sum = sum + (evenFlag * co3) * s_srcPatch[lsizey-7][1 + ((tidx + 2) >> 1)];
}
s_dstPatch[4 + tidy][tidx] = sum;
}
barrier(CLK_LOCAL_MEM_FENCE);
sum = co3 * s_dstPatch[2 + tidy - 2][tidx];
sum = sum + co2 * s_dstPatch[2 + tidy - 1][tidx];
sum = sum + co1 * s_dstPatch[2 + tidy ][tidx];
sum = sum + co2 * s_dstPatch[2 + tidy + 1][tidx];
sum = sum + co3 * s_dstPatch[2 + tidy + 2][tidx];
sum = coefy2* s_dstPatch[1 + ((tidy - 1) >> 1)][tidx];
sum = mad(coefy1, s_dstPatch[1 + ((tidy ) >> 1)][tidx], sum);
sum = mad(coefy2, s_dstPatch[1 + ((tidy + 2) >> 1)][tidx], sum);
if ((x < dst_cols) && (y < dst_rows))
storepix(convertToT(sum), dstData + y * dst_step + x * PIXSIZE);
@ -177,12 +135,12 @@ __kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_o
const int ly = 2*get_local_id(1);
__local FT s_srcPatch[LOCAL_SIZE+2][LOCAL_SIZE+2];
__local FT s_dstPatch[2*LOCAL_SIZE+4][2*LOCAL_SIZE];
__local FT s_dstPatch[LOCAL_SIZE+2][2*LOCAL_SIZE];
__global uchar * dstData = dst + dst_offset;
__global const uchar * srcData = src + src_offset;
if( lx < (LOCAL_SIZE+2) && lx < (LOCAL_SIZE+2) )
if( lx < (LOCAL_SIZE+2) && ly < (LOCAL_SIZE+2) )
{
int srcx = mad24((int)get_group_id(0), LOCAL_SIZE, lx) - 1;
int srcy = mad24((int)get_group_id(1), LOCAL_SIZE, ly) - 1;
@ -209,17 +167,13 @@ __kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_o
sum = co3 * s_srcPatch[1 + (ly >> 1)][1 + ((lx - 2) >> 1)];
sum = sum + co1 * s_srcPatch[1 + (ly >> 1)][1 + ((lx ) >> 1)];
sum = sum + co3 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 2) >> 1)];
s_dstPatch[2 + ly][lx] = sum;
s_dstPatch[1 + get_local_id(1)][lx] = sum;
// (x+1,y)
sum = co2 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 1 - 1) >> 1)];
sum = sum + co2 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 1 + 1) >> 1)];
s_dstPatch[2 + ly][lx+1] = sum;
// (x, y+1) (x+1, y+1)
s_dstPatch[2 + ly+1][lx] = 0.f;
s_dstPatch[2 + ly+1][lx+1] = 0.f;
s_dstPatch[1 + get_local_id(1)][lx+1] = sum;
if (ly < 1)
{
@ -227,16 +181,12 @@ __kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_o
sum = co3 * s_srcPatch[0][1 + ((lx - 2) >> 1)];
sum = sum + co1 * s_srcPatch[0][1 + ((lx ) >> 1)];
sum = sum + co3 * s_srcPatch[0][1 + ((lx + 2) >> 1)];
s_dstPatch[ly][lx] = sum;
s_dstPatch[0][lx] = sum;
// (x+1,y)
sum = co2 * s_srcPatch[0][1 + ((lx + 1 - 1) >> 1)];
sum = sum + co2 * s_srcPatch[0][1 + ((lx + 1 + 1) >> 1)];
s_dstPatch[ly][lx+1] = sum;
// (x, y+1) (x+1, y+1)
s_dstPatch[ly+1][lx] = 0.f;
s_dstPatch[ly+1][lx+1] = 0.f;
s_dstPatch[0][lx+1] = sum;
}
if (ly > 2*LOCAL_SIZE-3)
@ -245,59 +195,40 @@ __kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_o
sum = co3 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx - 2) >> 1)];
sum = sum + co1 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx ) >> 1)];
sum = sum + co3 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 2) >> 1)];
s_dstPatch[4 + ly][lx] = sum;
s_dstPatch[LOCAL_SIZE+1][lx] = sum;
// (x+1,y)
sum = co2 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 1 - 1) >> 1)];
sum = sum + co2 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 1 + 1) >> 1)];
s_dstPatch[4 + ly][lx+1] = sum;
// (x, y+1) (x+1, y+1)
s_dstPatch[4 + ly+1][lx] = 0.f;
s_dstPatch[4 + ly+1][lx+1] = 0.f;
s_dstPatch[LOCAL_SIZE+1][lx+1] = sum;
}
barrier(CLK_LOCAL_MEM_FENCE);
int dst_x = 2*get_global_id(0);
int dst_y = 2*get_global_id(1);
// (x,y)
sum = co3 * s_dstPatch[2 + ly - 2][lx];
sum = sum + co2 * s_dstPatch[2 + ly - 1][lx];
sum = sum + co1 * s_dstPatch[2 + ly ][lx];
sum = sum + co2 * s_dstPatch[2 + ly + 1][lx];
sum = sum + co3 * s_dstPatch[2 + ly + 2][lx];
if ((dst_x < dst_cols) && (dst_y < dst_rows))
{
// (x,y)
sum = co3 * s_dstPatch[1 + get_local_id(1) - 1][lx];
sum = sum + co1 * s_dstPatch[1 + get_local_id(1) ][lx];
sum = sum + co3 * s_dstPatch[1 + get_local_id(1) + 1][lx];
storepix(convertToT(sum), dstData + dst_y * dst_step + dst_x * PIXSIZE);
// (x+1,y)
sum = co3 * s_dstPatch[2 + ly - 2][lx+1];
sum = sum + co2 * s_dstPatch[2 + ly - 1][lx+1];
sum = sum + co1 * s_dstPatch[2 + ly ][lx+1];
sum = sum + co2 * s_dstPatch[2 + ly + 1][lx+1];
sum = sum + co3 * s_dstPatch[2 + ly + 2][lx+1];
if ((dst_x+1 < dst_cols) && (dst_y < dst_rows))
// (x+1,y)
sum = co3 * s_dstPatch[1 + get_local_id(1) - 1][lx+1];
sum = sum + co1 * s_dstPatch[1 + get_local_id(1) ][lx+1];
sum = sum + co3 * s_dstPatch[1 + get_local_id(1) + 1][lx+1];
storepix(convertToT(sum), dstData + dst_y * dst_step + (dst_x+1) * PIXSIZE);
// (x,y+1)
sum = co3 * s_dstPatch[2 + ly+1 - 2][lx];
sum = sum + co2 * s_dstPatch[2 + ly+1 - 1][lx];
sum = sum + co1 * s_dstPatch[2 + ly+1 ][lx];
sum = sum + co2 * s_dstPatch[2 + ly+1 + 1][lx];
sum = sum + co3 * s_dstPatch[2 + ly+1 + 2][lx];
if ((dst_x < dst_cols) && (dst_y+1 < dst_rows))
// (x,y+1)
sum = co2 * s_dstPatch[1 + get_local_id(1) ][lx];
sum = sum + co2 * s_dstPatch[1 + get_local_id(1) + 1][lx];
storepix(convertToT(sum), dstData + (dst_y+1) * dst_step + dst_x * PIXSIZE);
// (x+1,y+1)
sum = co3 * s_dstPatch[2 + ly+1 - 2][lx+1];
sum = sum + co2 * s_dstPatch[2 + ly+1 - 1][lx+1];
sum = sum + co1 * s_dstPatch[2 + ly+1 ][lx+1];
sum = sum + co2 * s_dstPatch[2 + ly+1 + 1][lx+1];
sum = sum + co3 * s_dstPatch[2 + ly+1 + 2][lx+1];
if ((dst_x+1 < dst_cols) && (dst_y+1 < dst_rows))
// (x+1,y+1)
sum = co2 * s_dstPatch[1 + get_local_id(1) ][lx+1];
sum = sum + co2 * s_dstPatch[1 + get_local_id(1) + 1][lx+1];
storepix(convertToT(sum), dstData + (dst_y+1) * dst_step + (dst_x+1) * PIXSIZE);
}
}

14
modules/imgproc/src/pyramids.cpp
Unescape View File

@ -467,7 +467,7 @@ static bool ocl_pyrUp( InputArray _src, OutputArray _dst, const Size& _dsz, int
UMat dst = _dst.getUMat();
int float_depth = depth == CV_64F ? CV_64F : CV_32F;
int local_size = 8;
const int local_size = 16;
char cvt[2][50];
String buildOptions = format(
"-D T=%s -D FT=%s -D convertToT=%s -D convertToFT=%s%s "
@ -478,25 +478,21 @@ static bool ocl_pyrUp( InputArray _src, OutputArray _dst, const Size& _dsz, int
doubleSupport ? " -D DOUBLE_SUPPORT" : "",
ocl::typeToStr(depth), channels, local_size
);
size_t globalThreads[2];
size_t globalThreads[2] = { dst.cols, dst.rows };
size_t localThreads[2] = { local_size, local_size };
ocl::Kernel k;
if (ocl::Device::getDefault().isIntel() && channels == 1)
{
k.create("pyrUp_unrolled", ocl::imgproc::pyr_up_oclsrc, buildOptions);
globalThreads[0] = dst.cols/2; globalThreads[1] = dst.rows/2;
}
}
else
{
k.create("pyrUp", ocl::imgproc::pyr_up_oclsrc, buildOptions);
local_size = 16;
globalThreads[0] = dst.cols; globalThreads[1] = dst.rows;
}
if (k.empty())
return false;
k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst));
size_t localThreads[2] = {local_size, local_size};
return k.run(2, globalThreads, localThreads, false);
}

Write Preview
Loading…
Cancel
Save