generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Merge pull request #2823 from ilya-lavrenov:tapi_warpaffine

Browse Source
This commit is contained in:
Alexander Alekhin 2014-06-11 14:41:29 +04:00 committed by OpenCV Buildbot
commit 14671e0cd6
2 changed files with 97 additions and 70 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

21
modules/imgproc/src/imgwarp.cpp
View File

@ -4167,11 +4167,12 @@ static bool ocl_warpTransform(InputArray _src, OutputArray _dst, InputArray _M0,
int op_type) int op_type)
{ {
CV_Assert(op_type == OCL_OP_AFFINE || op_type == OCL_OP_PERSPECTIVE); CV_Assert(op_type == OCL_OP_AFFINE || op_type == OCL_OP_PERSPECTIVE);
const ocl::Device & dev = ocl::Device::getDefault();
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
double doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; double doubleSupport = dev.doubleFPConfig() > 0;
int interpolation = flags & INTER_MAX; int interpolation = flags & INTER_MAX, rowsPerWI = dev.isIntel() && interpolation <= INTER_LINEAR ? 4 : 1;
if( interpolation == INTER_AREA ) if( interpolation == INTER_AREA )
interpolation = INTER_LINEAR; interpolation = INTER_LINEAR;
@ -4193,30 +4194,30 @@ static bool ocl_warpTransform(InputArray _src, OutputArray _dst, InputArray _M0,
String opts; String opts;
if (interpolation == INTER_NEAREST) if (interpolation == INTER_NEAREST)
{ {
opts = format("-D INTER_NEAREST -D T=%s%s -D T1=%s -D ST=%s -D cn=%d", ocl::typeToStr(type), opts = format("-D INTER_NEAREST -D T=%s%s -D T1=%s -D ST=%s -D cn=%d -D rowsPerWI=%d",
doubleSupport ? " -D DOUBLE_SUPPORT" : "", ocl::typeToStr(type), doubleSupport ? " -D DOUBLE_SUPPORT" : "",
ocl::typeToStr(CV_MAT_DEPTH(type)), ocl::typeToStr(CV_MAT_DEPTH(type)),
ocl::typeToStr(sctype), ocl::typeToStr(sctype), cn, rowsPerWI);
cn);
} }
else else
{ {
char cvt[2][50]; char cvt[2][50];
opts = format("-D INTER_%s -D T=%s -D T1=%s -D ST=%s -D WT=%s -D depth=%d -D convertToWT=%s -D convertToT=%s%s -D cn=%d", opts = format("-D INTER_%s -D T=%s -D T1=%s -D ST=%s -D WT=%s -D depth=%d"
" -D convertToWT=%s -D convertToT=%s%s -D cn=%d -D rowsPerWI=%d",
interpolationMap[interpolation], ocl::typeToStr(type), interpolationMap[interpolation], ocl::typeToStr(type),
ocl::typeToStr(CV_MAT_DEPTH(type)), ocl::typeToStr(CV_MAT_DEPTH(type)),
ocl::typeToStr(sctype), ocl::typeToStr(sctype),
ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)), depth, ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)), depth,
ocl::convertTypeStr(depth, wdepth, cn, cvt[0]), ocl::convertTypeStr(depth, wdepth, cn, cvt[0]),
ocl::convertTypeStr(wdepth, depth, cn, cvt[1]), ocl::convertTypeStr(wdepth, depth, cn, cvt[1]),
doubleSupport ? " -D DOUBLE_SUPPORT" : "", cn); doubleSupport ? " -D DOUBLE_SUPPORT" : "", cn, rowsPerWI);
} }
k.create(kernelName, program, opts); k.create(kernelName, program, opts);
if (k.empty()) if (k.empty())
return false; return false;
double borderBuf[] = {0, 0, 0, 0}; double borderBuf[] = { 0, 0, 0, 0 };
scalarToRawData(borderValue, borderBuf, sctype); scalarToRawData(borderValue, borderBuf, sctype);
UMat src = _src.getUMat(), M0; UMat src = _src.getUMat(), M0;
@ -4251,7 +4252,7 @@ static bool ocl_warpTransform(InputArray _src, OutputArray _dst, InputArray _M0,
k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrReadOnly(M0), k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrReadOnly(M0),
ocl::KernelArg(0, 0, 0, 0, borderBuf, CV_ELEM_SIZE(sctype))); ocl::KernelArg(0, 0, 0, 0, borderBuf, CV_ELEM_SIZE(sctype)));
size_t globalThreads[2] = { dst.cols, dst.rows }; size_t globalThreads[2] = { dst.cols, (dst.rows + rowsPerWI - 1) / rowsPerWI };
return k.run(2, globalThreads, NULL, false); return k.run(2, globalThreads, NULL, false);
} }

146
modules/imgproc/src/opencl/warp_affine.cl
View File

@ -91,29 +91,32 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
__constant CT * M, ST scalar_) __constant CT * M, ST scalar_)
{ {
int dx = get_global_id(0); int dx = get_global_id(0);
int dy = get_global_id(1); int dy0 = get_global_id(1) * rowsPerWI;
if (dx < dst_cols && dy < dst_rows) if (dx < dst_cols)
{ {
int round_delta = (AB_SCALE >> 1); int round_delta = (AB_SCALE >> 1);
int X0 = rint(M[0] * dx * AB_SCALE); int X0_ = rint(M[0] * dx * AB_SCALE);
int Y0 = rint(M[3] * dx * AB_SCALE); int Y0_ = rint(M[3] * dx * AB_SCALE);
X0 += rint((M[1]*dy + M[2]) * AB_SCALE) + round_delta; int dst_index = mad24(dy0, dst_step, mad24(dx, pixsize, dst_offset));
Y0 += rint((M[4]*dy + M[5]) * AB_SCALE) + round_delta;
short sx = convert_short_sat(X0 >> AB_BITS); for (int dy = dy0, dy1 = min(dst_rows, dy0 + rowsPerWI); dy < dy1; ++dy, dst_index += dst_step)
short sy = convert_short_sat(Y0 >> AB_BITS);
int dst_index = mad24(dy, dst_step, dst_offset + dx * pixsize);
if (sx >= 0 && sx < src_cols && sy >= 0 && sy < src_rows)
{ {
int src_index = mad24(sy, src_step, src_offset + sx * pixsize); int X0 = X0_ + rint(fma(M[1], dy, M[2]) * AB_SCALE) + round_delta;
storepix(loadpix(srcptr + src_index), dstptr + dst_index); int Y0 = Y0_ + rint(fma(M[4], dy, M[5]) * AB_SCALE) + round_delta;
short sx = convert_short_sat(X0 >> AB_BITS);
short sy = convert_short_sat(Y0 >> AB_BITS);
if (sx >= 0 && sx < src_cols && sy >= 0 && sy < src_rows)
{
int src_index = mad24(sy, src_step, mad24(sx, pixsize, src_offset));
storepix(loadpix(srcptr + src_index), dstptr + dst_index);
}
else
storepix(scalar, dstptr + dst_index);
} }
else
storepix(scalar, dstptr + dst_index);
} }
} }
@ -124,52 +127,63 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
__constant CT * M, ST scalar_) __constant CT * M, ST scalar_)
{ {
int dx = get_global_id(0); int dx = get_global_id(0);
int dy = get_global_id(1); int dy0 = get_global_id(1) * rowsPerWI;
if (dx < dst_cols && dy < dst_rows) if (dx < dst_cols)
{ {
int round_delta = AB_SCALE/INTER_TAB_SIZE/2; int round_delta = AB_SCALE/INTER_TAB_SIZE/2;
int tmp = (dx << AB_BITS); int tmp = (dx << AB_BITS);
int X0 = rint(M[0] * tmp); int X0_ = rint(M[0] * tmp);
int Y0 = rint(M[3] * tmp); int Y0_ = rint(M[3] * tmp);
X0 += rint((M[1]*dy + M[2]) * AB_SCALE) + round_delta;
Y0 += rint((M[4]*dy + M[5]) * AB_SCALE) + round_delta;
X0 = X0 >> (AB_BITS - INTER_BITS);
Y0 = Y0 >> (AB_BITS - INTER_BITS);
short sx = convert_short_sat(X0 >> INTER_BITS); for (int dy = dy0, dy1 = min(dst_rows, dy0 + rowsPerWI); dy < dy1; ++dy)
short sy = convert_short_sat(Y0 >> INTER_BITS); {
short ax = convert_short(X0 & (INTER_TAB_SIZE-1)); int X0 = X0_ + rint(fma(M[1], dy, M[2]) * AB_SCALE) + round_delta;
short ay = convert_short(Y0 & (INTER_TAB_SIZE-1)); int Y0 = Y0_ + rint(fma(M[4], dy, M[5]) * AB_SCALE) + round_delta;
X0 = X0 >> (AB_BITS - INTER_BITS);
Y0 = Y0 >> (AB_BITS - INTER_BITS);
WT v0 = (sx >= 0 && sx < src_cols && sy >= 0 && sy < src_rows) ? short sx = convert_short_sat(X0 >> INTER_BITS);
convertToWT(loadpix(srcptr + mad24(sy, src_step, src_offset + sx * pixsize))) : scalar; short sy = convert_short_sat(Y0 >> INTER_BITS);
WT v1 = (sx+1 >= 0 && sx+1 < src_cols && sy >= 0 && sy < src_rows) ? short ax = convert_short(X0 & (INTER_TAB_SIZE-1));
convertToWT(loadpix(srcptr + mad24(sy, src_step, src_offset + (sx+1) * pixsize))) : scalar; short ay = convert_short(Y0 & (INTER_TAB_SIZE-1));
WT v2 = (sx >= 0 && sx < src_cols && sy+1 >= 0 && sy+1 < src_rows) ?
convertToWT(loadpix(srcptr + mad24(sy+1, src_step, src_offset + sx * pixsize))) : scalar;
WT v3 = (sx+1 >= 0 && sx+1 < src_cols && sy+1 >= 0 && sy+1 < src_rows) ?
convertToWT(loadpix(srcptr + mad24(sy+1, src_step, src_offset + (sx+1) * pixsize))) : scalar;
float taby = 1.f/INTER_TAB_SIZE*ay; WT v0 = scalar, v1 = scalar, v2 = scalar, v3 = scalar;
float tabx = 1.f/INTER_TAB_SIZE*ax; if (sx >= 0 && sx < src_cols)
{
if (sy >= 0 && sy < src_rows)
v0 = convertToWT(loadpix(srcptr + mad24(sy, src_step, mad24(sx, pixsize, src_offset))));
if (sy+1 >= 0 && sy+1 < src_rows)
v2 = convertToWT(loadpix(srcptr + mad24(sy+1, src_step, mad24(sx, pixsize, src_offset))));
}
if (sx+1 >= 0 && sx+1 < src_cols)
{
if (sy >= 0 && sy < src_rows)
v1 = convertToWT(loadpix(srcptr + mad24(sy, src_step, mad24(sx+1, pixsize, src_offset))));
if (sy+1 >= 0 && sy+1 < src_rows)
v3 = convertToWT(loadpix(srcptr + mad24(sy+1, src_step, mad24(sx+1, pixsize, src_offset))));
}
int dst_index = mad24(dy, dst_step, dst_offset + dx * pixsize); float taby = 1.f/INTER_TAB_SIZE*ay;
float tabx = 1.f/INTER_TAB_SIZE*ax;
int dst_index = mad24(dy, dst_step, mad24(dx, pixsize, dst_offset));
#if depth <= 4 #if depth <= 4
int itab0 = convert_short_sat_rte( (1.0f-taby)*(1.0f-tabx) * INTER_REMAP_COEF_SCALE ); int itab0 = convert_short_sat_rte( (1.0f-taby)*(1.0f-tabx) * INTER_REMAP_COEF_SCALE );
int itab1 = convert_short_sat_rte( (1.0f-taby)*tabx * INTER_REMAP_COEF_SCALE ); int itab1 = convert_short_sat_rte( (1.0f-taby)*tabx * INTER_REMAP_COEF_SCALE );
int itab2 = convert_short_sat_rte( taby*(1.0f-tabx) * INTER_REMAP_COEF_SCALE ); int itab2 = convert_short_sat_rte( taby*(1.0f-tabx) * INTER_REMAP_COEF_SCALE );
int itab3 = convert_short_sat_rte( taby*tabx * INTER_REMAP_COEF_SCALE ); int itab3 = convert_short_sat_rte( taby*tabx * INTER_REMAP_COEF_SCALE );
WT val = v0 * itab0 + v1 * itab1 + v2 * itab2 + v3 * itab3; WT val = mad24(v0, itab0, mad24(v1, itab1, mad24(v2, itab2, v3 * itab3)));
storepix(convertToT((val + (1 << (INTER_REMAP_COEF_BITS-1))) >> INTER_REMAP_COEF_BITS), dstptr + dst_index); storepix(convertToT((val + (1 << (INTER_REMAP_COEF_BITS-1))) >> INTER_REMAP_COEF_BITS), dstptr + dst_index);
#else #else
float tabx2 = 1.0f - tabx, taby2 = 1.0f - taby; float tabx2 = 1.0f - tabx, taby2 = 1.0f - taby;
WT val = v0 * tabx2 * taby2 + v1 * tabx * taby2 + v2 * tabx2 * taby + v3 * tabx * taby; WT val = fma(v0, tabx2 * taby2, fma(v1, tabx * taby2, fma(v2, tabx2 * taby, v3 * tabx * taby)));
storepix(convertToT(val), dstptr + dst_index); storepix(convertToT(val), dstptr + dst_index);
#endif #endif
}
} }
} }
@ -179,9 +193,9 @@ inline void interpolateCubic( float x, float* coeffs )
{ {
const float A = -0.75f; const float A = -0.75f;
coeffs[0] = ((A*(x + 1.f) - 5.0f*A)*(x + 1.f) + 8.0f*A)*(x + 1.f) - 4.0f*A; coeffs[0] = fma(fma(fma(A, (x + 1.f), - 5.0f*A), (x + 1.f), 8.0f*A), x + 1.f, - 4.0f*A);
coeffs[1] = ((A + 2.f)*x - (A + 3.f))*x*x + 1.f; coeffs[1] = fma(fma(A + 2.f, x, - (A + 3.f)), x*x, 1.f);
coeffs[2] = ((A + 2.f)*(1.f - x) - (A + 3.f))*(1.f - x)*(1.f - x) + 1.f; coeffs[2] = fma(fma(A + 2.f, 1.f - x, - (A + 3.f)), (1.f - x)*(1.f - x), 1.f);
coeffs[3] = 1.f - coeffs[0] - coeffs[1] - coeffs[2]; coeffs[3] = 1.f - coeffs[0] - coeffs[1] - coeffs[2];
} }
@ -199,8 +213,9 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
int tmp = (dx << AB_BITS); int tmp = (dx << AB_BITS);
int X0 = rint(M[0] * tmp); int X0 = rint(M[0] * tmp);
int Y0 = rint(M[3] * tmp); int Y0 = rint(M[3] * tmp);
X0 += rint((M[1]*dy + M[2]) * AB_SCALE) + round_delta;
Y0 += rint((M[4]*dy + M[5]) * AB_SCALE) + round_delta; X0 += rint(fma(M[1], dy, M[2]) * AB_SCALE) + round_delta;
Y0 += rint(fma(M[4], dy, M[5]) * AB_SCALE) + round_delta;
X0 = X0 >> (AB_BITS - INTER_BITS); X0 = X0 >> (AB_BITS - INTER_BITS);
Y0 = Y0 >> (AB_BITS - INTER_BITS); Y0 = Y0 >> (AB_BITS - INTER_BITS);
@ -212,10 +227,21 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
WT v[16]; WT v[16];
#pragma unroll #pragma unroll
for (int y = 0; y < 4; y++) for (int y = 0; y < 4; y++)
#pragma unroll {
for (int x = 0; x < 4; x++) if (sy+y >= 0 && sy+y < src_rows)
v[mad24(y, 4, x)] = (sx+x >= 0 && sx+x < src_cols && sy+y >= 0 && sy+y < src_rows) ? {
convertToWT(loadpix(srcptr + mad24(sy+y, src_step, src_offset + (sx+x) * pixsize))) : scalar; #pragma unroll
for (int x = 0; x < 4; x++)
v[mad24(y, 4, x)] = sx+x >= 0 && sx+x < src_cols ?
convertToWT(loadpix(srcptr + mad24(sy+y, src_step, mad24(sx+x, pixsize, src_offset)))) : scalar;
}
else
{
#pragma unroll
for (int x = 0; x < 4; x++)
v[mad24(y, 4, x)] = scalar;
}
}
float tab1y[4], tab1x[4]; float tab1y[4], tab1x[4];
@ -224,7 +250,7 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
interpolateCubic(ayy, tab1y); interpolateCubic(ayy, tab1y);
interpolateCubic(axx, tab1x); interpolateCubic(axx, tab1x);
int dst_index = mad24(dy, dst_step, dst_offset + dx * pixsize); int dst_index = mad24(dy, dst_step, mad24(dx, pixsize, dst_offset));
WT sum = (WT)(0); WT sum = (WT)(0);
#if depth <= 4 #if depth <= 4
@ -236,12 +262,12 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
#pragma unroll #pragma unroll
for (int i = 0; i < 16; i++) for (int i = 0; i < 16; i++)
sum += v[i] * itab[i]; sum = mad24(v[i], itab[i], sum);
storepix(convertToT( (sum + (1 << (INTER_REMAP_COEF_BITS-1))) >> INTER_REMAP_COEF_BITS ), dstptr + dst_index); storepix(convertToT( (sum + (1 << (INTER_REMAP_COEF_BITS-1))) >> INTER_REMAP_COEF_BITS ), dstptr + dst_index);
#else #else
#pragma unroll #pragma unroll
for (int i = 0; i < 16; i++) for (int i = 0; i < 16; i++)
sum += v[i] * tab1y[(i>>2)] * tab1x[(i&3)]; sum = fma(v[i], tab1y[(i>>2)] * tab1x[(i&3)], sum);
storepix(convertToT( sum ), dstptr + dst_index); storepix(convertToT( sum ), dstptr + dst_index);
#endif #endif
} }