#define noconvert
-
__kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols,
__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
+ 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];
+
+ __global uchar * dstData = dst + dst_offset;
+ __global const uchar * srcData = src + src_offset;
+
+ if( tidx < 10 && tidy < 10 )
+ {
+ 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);
+
+ s_srcPatch[tidy][tidx] = convertToFT(loadpix(srcData + srcy * src_step + srcx * PIXSIZE));
+ }
+
+ 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;
+
+ if (tidy < 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)];
+ }
+
+ 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];
+
+ if ((x < dst_cols) && (y < dst_rows))
+ storepix(convertToT(sum), dstData + y * dst_step + x * PIXSIZE);
+}
+
+
+__kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
+{
const int lx = 2*get_local_id(0);
const int ly = 2*get_local_id(1);
FT sum;
- const FT co1 = 0.375f;
- const FT co2 = 0.25f;
- const FT co3 = 0.0625f;
+ const FT co1 = 0.75f;
+ const FT co2 = 0.5f;
+ const FT co3 = 0.125f;
// (x,y)
sum = co3 * s_srcPatch[1 + (ly >> 1)][1 + ((lx - 2) >> 1)];
sum = sum + co3 * s_dstPatch[2 + ly + 2][lx];
if ((dst_x < dst_cols) && (dst_y < dst_rows))
- storepix(convertToT(4.0f * sum), dstData + dst_y * dst_step + dst_x * PIXSIZE);
+ 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 + co3 * s_dstPatch[2 + ly + 2][lx+1];
if ((dst_x+1 < dst_cols) && (dst_y < dst_rows))
- storepix(convertToT(4.0f * sum), dstData + dst_y * dst_step + (dst_x+1) * PIXSIZE);
+ 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 + co3 * s_dstPatch[2 + ly+1 + 2][lx];
if ((dst_x < dst_cols) && (dst_y+1 < dst_rows))
- storepix(convertToT(4.0f * sum), dstData + (dst_y+1) * dst_step + dst_x * PIXSIZE);
+ 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 + co3 * s_dstPatch[2 + ly+1 + 2][lx+1];
if ((dst_x+1 < dst_cols) && (dst_y+1 < dst_rows))
- storepix(convertToT(4.0f * sum), dstData + (dst_y+1) * dst_step + (dst_x+1) * PIXSIZE);
+ storepix(convertToT(sum), dstData + (dst_y+1) * dst_step + (dst_x+1) * PIXSIZE);
}
doubleSupport ? " -D DOUBLE_SUPPORT" : "",
ocl::typeToStr(depth), channels, local_size
);
- ocl::Kernel k("pyrUp", ocl::imgproc::pyr_up_oclsrc, buildOptions);
+ size_t globalThreads[2];
+ 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 globalThreads[2] = {dst.cols/2, dst.rows/2};
size_t localThreads[2] = {local_size, local_size};
return k.run(2, globalThreads, localThreads, false);