ocl: copyMakeBorder 3-channel

This commit is contained in:
Alexander Alekhin
2014-02-27 14:49:36 +04:00
parent 70e22b682d
commit e4608229f0
3 changed files with 42 additions and 15 deletions

View File

@@ -42,6 +42,18 @@
#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))
#define convertScalar(a) (a)
#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)
#define convertScalar(a) (T)(a.x, a.y, a.z)
#endif
#ifdef BORDER_CONSTANT
#define EXTRAPOLATE(x, y, v) v = scalar;
#elif defined BORDER_REPLICATE
@@ -49,7 +61,7 @@
{ \
x = clamp(x, 0, src_cols - 1); \
y = clamp(y, 0, src_rows - 1); \
v = *(__global const T *)(srcptr + mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset))); \
v = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset))); \
}
#elif defined BORDER_WRAP
#define EXTRAPOLATE(x, y, v) \
@@ -63,7 +75,7 @@
y -= ((y - src_rows + 1) / src_rows) * src_rows; \
if( y >= src_rows ) \
y %= src_rows; \
v = *(__global const T *)(srcptr + mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset))); \
v = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset))); \
}
#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
#ifdef BORDER_REFLECT
@@ -97,7 +109,7 @@
y = src_rows - 1 - (y - src_rows) - delta; \
} \
while (y >= src_rows || y < 0); \
v = *(__global const T *)(srcptr + mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset))); \
v = loadpix(srcptr + mad24(y, src_step, mad24(x, TSIZE, src_offset))); \
}
#else
#error No extrapolation method
@@ -107,26 +119,33 @@
__kernel void copyMakeBorder(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,
__global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols,
int top, int left, T scalar)
int top, int left, ST nVal)
{
int x = get_global_id(0);
int y = get_global_id(1);
#ifdef BORDER_CONSTANT
T scalar = convertScalar(nVal);
#endif
if (x < dst_cols && y < dst_rows)
{
int src_x = x - left;
int src_y = y - top;
int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(T), dst_offset));
int dst_index = mad24(y, dst_step, mad24(x, (int)TSIZE, dst_offset));
__global T * dst = (__global T *)(dstptr + dst_index);
T v;
if (NEED_EXTRAPOLATION(src_x, src_y))
EXTRAPOLATE(src_x, src_y, dst[0])
{
EXTRAPOLATE(src_x, src_y, v)
}
else
{
int src_index = mad24(src_y, src_step, mad24(src_x, (int)sizeof(T), src_offset));
__global const T * src = (__global const T *)(srcptr + src_index);
dst[0] = src[0];
int src_index = mad24(src_y, src_step, mad24(src_x, TSIZE, src_offset));
v = loadpix(srcptr + src_index);
}
storepix(v, dst);
}
}