From e1e243588e73aa96a458f8f585f4ed055fc685d4 Mon Sep 17 00:00:00 2001 From: Sancho McCann Date: Wed, 11 Jun 2014 17:25:22 -0700 Subject: [PATCH 01/24] Bugfix: Memory leak in deletion of er_stack nodes of ERFilter. --- modules/objdetect/src/erfilter.cpp | 24 +++++++++++++++++++++++- 1 file changed, 23 insertions(+), 1 deletion(-) diff --git a/modules/objdetect/src/erfilter.cpp b/modules/objdetect/src/erfilter.cpp index e942094b1d..6651a00933 100644 --- a/modules/objdetect/src/erfilter.cpp +++ b/modules/objdetect/src/erfilter.cpp @@ -42,6 +42,7 @@ #include "precomp.hpp" #include +#include #if defined _MSC_VER && _MSC_VER == 1500 typedef int int_fast32_t; @@ -57,6 +58,27 @@ using namespace std; namespace cv { +// Deletes a tree of ERStat regions starting at root. Used only +// internally to this implementation. +static void deleteERStatTree(ERStat* root) { + queue to_delete; + to_delete.push(root); + while (!to_delete.empty()) { + ERStat* n = to_delete.front(); + to_delete.pop(); + ERStat* c = n->child; + if (c != NULL) { + to_delete.push(c); + ERStat* sibling = c->next; + while (sibling != NULL) { + to_delete.push(sibling); + sibling = sibling->next; + } + } + delete n; + } +} + ERStat::ERStat(int init_level, int init_pixel, int init_x, int init_y) : pixel(init_pixel), level(init_level), area(0), perimeter(0), euler(0), probability(1.0), parent(0), child(0), next(0), prev(0), local_maxima(0), @@ -497,7 +519,7 @@ void ERFilterNM::er_tree_extract( InputArray image ) delete(stat->crossings); stat->crossings = NULL; } - delete stat; + deleteERStatTree(stat); } er_stack.clear(); From 93712e913a2af63e4954df206659618d8af5ebba Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 12 Jun 2014 20:10:16 +0400 Subject: [PATCH 02/24] optimization of cv::warpAffine INTER_CUBIC --- modules/imgproc/src/imgwarp.cpp | 3 +- modules/imgproc/src/opencl/warp_affine.cl | 140 ++++++++++++++++++---- 2 files changed, 122 insertions(+), 21 deletions(-) diff --git a/modules/imgproc/src/imgwarp.cpp b/modules/imgproc/src/imgwarp.cpp index c946afc97a..34cd2c8eee 100644 --- a/modules/imgproc/src/imgwarp.cpp +++ b/modules/imgproc/src/imgwarp.cpp @@ -4187,7 +4187,8 @@ static bool ocl_warpTransform(InputArray _src, OutputArray _dst, InputArray _M0, const char * const kernelName = op_type == OCL_OP_AFFINE ? "warpAffine" : "warpPerspective"; int scalarcn = cn == 3 ? 4 : cn; - int wdepth = interpolation == INTER_NEAREST ? depth : std::max(CV_32S, depth); + bool is32f = !dev.isAMD() && (interpolation == INTER_CUBIC || interpolation == INTER_LINEAR); + int wdepth = interpolation == INTER_NEAREST ? depth : std::max(is32f ? CV_32F : CV_32S, depth); int sctype = CV_MAKETYPE(wdepth, scalarcn); ocl::Kernel k; diff --git a/modules/imgproc/src/opencl/warp_affine.cl b/modules/imgproc/src/opencl/warp_affine.cl index bb041d1601..8ee34d0d65 100644 --- a/modules/imgproc/src/opencl/warp_affine.cl +++ b/modules/imgproc/src/opencl/warp_affine.cl @@ -61,6 +61,7 @@ #define AB_SCALE (1 << AB_BITS) #define INTER_REMAP_COEF_BITS 15 #define INTER_REMAP_COEF_SCALE (1 << INTER_REMAP_COEF_BITS) +#define ROUND_DELTA (1 << (AB_BITS - INTER_BITS - 1)) #define noconvert @@ -122,6 +123,14 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of #elif defined INTER_LINEAR +__constant float coeffs[64] = +{ 1.000000f, 0.000000f, 0.968750f, 0.031250f, 0.937500f, 0.062500f, 0.906250f, 0.093750f, 0.875000f, 0.125000f, 0.843750f, 0.156250f, + 0.812500f, 0.187500f, 0.781250f, 0.218750f, 0.750000f, 0.250000f, 0.718750f, 0.281250f, 0.687500f, 0.312500f, 0.656250f, 0.343750f, + 0.625000f, 0.375000f, 0.593750f, 0.406250f, 0.562500f, 0.437500f, 0.531250f, 0.468750f, 0.500000f, 0.500000f, 0.468750f, 0.531250f, + 0.437500f, 0.562500f, 0.406250f, 0.593750f, 0.375000f, 0.625000f, 0.343750f, 0.656250f, 0.312500f, 0.687500f, 0.281250f, 0.718750f, + 0.250000f, 0.750000f, 0.218750f, 0.781250f, 0.187500f, 0.812500f, 0.156250f, 0.843750f, 0.125000f, 0.875000f, 0.093750f, 0.906250f, + 0.062500f, 0.937500f, 0.031250f, 0.968750f }; + __kernel void warpAffine(__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, __constant CT * M, ST scalar_) @@ -131,24 +140,21 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of if (dx < dst_cols) { - 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 Y0_ = rint(M[3] * tmp); for (int dy = dy0, dy1 = min(dst_rows, dy0 + rowsPerWI); dy < dy1; ++dy) { - int X0 = X0_ + rint(fma(M[1], dy, M[2]) * AB_SCALE) + round_delta; - int Y0 = Y0_ + rint(fma(M[4], dy, M[5]) * AB_SCALE) + round_delta; + int X0 = X0_ + rint(fma(M[1], dy, M[2]) * AB_SCALE) + ROUND_DELTA; + 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); - short sx = convert_short_sat(X0 >> INTER_BITS); - short sy = convert_short_sat(Y0 >> INTER_BITS); - short ax = convert_short(X0 & (INTER_TAB_SIZE-1)); - short ay = convert_short(Y0 & (INTER_TAB_SIZE-1)); + short sx = convert_short_sat(X0 >> INTER_BITS), sy = convert_short_sat(Y0 >> INTER_BITS); + short ax = convert_short(X0 & (INTER_TAB_SIZE-1)), ay = convert_short(Y0 & (INTER_TAB_SIZE-1)); +#if defined AMD_DEVICE || depth > 4 WT v0 = scalar, v1 = scalar, v2 = scalar, v3 = scalar; if (sx >= 0 && sx < src_cols) { @@ -180,8 +186,48 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of storepix(convertToT((val + (1 << (INTER_REMAP_COEF_BITS-1))) >> INTER_REMAP_COEF_BITS), dstptr + dst_index); #else float tabx2 = 1.0f - tabx, taby2 = 1.0f - taby; - WT val = fma(v0, tabx2 * taby2, fma(v1, tabx * taby2, fma(v2, tabx2 * taby, v3 * tabx * taby))); + WT val = fma(tabx2, fma(v0, taby2, v2 * taby), tabx * fma(v1, taby2, v3 * taby)); storepix(convertToT(val), dstptr + dst_index); +#endif +#else // INTEL_DEVICE + __constant float * coeffs_y = coeffs + (ay << 1), * coeffs_x = coeffs + (ax << 1); + + int src_index0 = mad24(sy, src_step, mad24(sx, pixsize, src_offset)), src_index; + int dst_index = mad24(dy, dst_step, mad24(dx, pixsize, dst_offset)); + + WT sum = (WT)(0), xsum; + #pragma unroll + for (int y = 0; y < 2; y++) + { + src_index = mad24(y, src_step, src_index0); + if (sy + y >= 0 && sy + y < src_rows) + { + xsum = (WT)(0); + if (sx >= 0 && sx + 2 < src_cols) + { +#if depth == 0 && cn == 1 + uchar2 value = vload2(0, srcptr + src_index); + xsum = dot(convert_float2(value), (float2)(coeffs_x[0], coeffs_x[1])); +#else + #pragma unroll + for (int x = 0; x < 2; x++) + xsum = fma(convertToWT(loadpix(srcptr + mad24(x, pixsize, src_index))), coeffs_x[x], xsum); +#endif + } + else + { + #pragma unroll + for (int x = 0; x < 2; x++) + xsum = fma(sx + x >= 0 && sx + x < src_cols ? + convertToWT(loadpix(srcptr + mad24(x, pixsize, src_index))) : scalar, coeffs_x[x], xsum); + } + sum = fma(xsum, coeffs_y[y], sum); + } + else + sum = fma(scalar, coeffs_y[y], sum); + } + + storepix(convertToT(sum), dstptr + dst_index); #endif } } @@ -189,6 +235,8 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of #elif defined INTER_CUBIC +#ifdef AMD_DEVICE + inline void interpolateCubic( float x, float* coeffs ) { const float A = -0.75f; @@ -199,6 +247,23 @@ inline void interpolateCubic( float x, float* coeffs ) coeffs[3] = 1.f - coeffs[0] - coeffs[1] - coeffs[2]; } +#else + +__constant float coeffs[128] = + { 0.000000f, 1.000000f, 0.000000f, 0.000000f, -0.021996f, 0.997841f, 0.024864f, -0.000710f, -0.041199f, 0.991516f, 0.052429f, -0.002747f, + -0.057747f, 0.981255f, 0.082466f, -0.005974f, -0.071777f, 0.967285f, 0.114746f, -0.010254f, -0.083427f, 0.949837f, 0.149040f, -0.015450f, + -0.092834f, 0.929138f, 0.185120f, -0.021423f, -0.100136f, 0.905418f, 0.222755f, -0.028038f, -0.105469f, 0.878906f, 0.261719f, -0.035156f, + -0.108971f, 0.849831f, 0.301781f, -0.042641f, -0.110779f, 0.818420f, 0.342712f, -0.050354f, -0.111031f, 0.784904f, 0.384285f, -0.058159f, + -0.109863f, 0.749512f, 0.426270f, -0.065918f, -0.107414f, 0.712471f, 0.468437f, -0.073494f, -0.103821f, 0.674011f, 0.510559f, -0.080750f, + -0.099220f, 0.634361f, 0.552406f, -0.087547f, -0.093750f, 0.593750f, 0.593750f, -0.093750f, -0.087547f, 0.552406f, 0.634361f, -0.099220f, + -0.080750f, 0.510559f, 0.674011f, -0.103821f, -0.073494f, 0.468437f, 0.712471f, -0.107414f, -0.065918f, 0.426270f, 0.749512f, -0.109863f, + -0.058159f, 0.384285f, 0.784904f, -0.111031f, -0.050354f, 0.342712f, 0.818420f, -0.110779f, -0.042641f, 0.301781f, 0.849831f, -0.108971f, + -0.035156f, 0.261719f, 0.878906f, -0.105469f, -0.028038f, 0.222755f, 0.905418f, -0.100136f, -0.021423f, 0.185120f, 0.929138f, -0.092834f, + -0.015450f, 0.149040f, 0.949837f, -0.083427f, -0.010254f, 0.114746f, 0.967285f, -0.071777f, -0.005974f, 0.082466f, 0.981255f, -0.057747f, + -0.002747f, 0.052429f, 0.991516f, -0.041199f, -0.000710f, 0.024864f, 0.997841f, -0.021996f }; + +#endif + __kernel void warpAffine(__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, __constant CT * M, ST scalar_) @@ -208,22 +273,17 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of if (dx < dst_cols && dy < dst_rows) { - int round_delta = ((AB_SCALE>>INTER_BITS)>>1); - int tmp = (dx << AB_BITS); - int X0 = rint(M[0] * tmp); - int Y0 = rint(M[3] * tmp); + int X0 = rint(M[0] * tmp) + rint(fma(M[1], dy, M[2]) * AB_SCALE) + ROUND_DELTA; + int Y0 = rint(M[3] * tmp) + rint(fma(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); Y0 = Y0 >> (AB_BITS - INTER_BITS); - int sx = (short)(X0 >> INTER_BITS) - 1; - int sy = (short)(Y0 >> INTER_BITS) - 1; - int ay = (short)(Y0 & (INTER_TAB_SIZE-1)); - int ax = (short)(X0 & (INTER_TAB_SIZE-1)); + int sx = (short)(X0 >> INTER_BITS) - 1, sy = (short)(Y0 >> INTER_BITS) - 1; + int ay = (short)(Y0 & (INTER_TAB_SIZE - 1)), ax = (short)(X0 & (INTER_TAB_SIZE - 1)); +#ifdef AMD_DEVICE WT v[16]; #pragma unroll for (int y = 0; y < 4; y++) @@ -269,6 +329,46 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of for (int i = 0; i < 16; i++) sum = fma(v[i], tab1y[(i>>2)] * tab1x[(i&3)], sum); storepix(convertToT( sum ), dstptr + dst_index); +#endif +#else // INTEL_DEVICE + __constant float * coeffs_y = coeffs + (ay << 2), * coeffs_x = coeffs + (ax << 2); + + int src_index0 = mad24(sy, src_step, mad24(sx, pixsize, src_offset)), src_index; + int dst_index = mad24(dy, dst_step, mad24(dx, pixsize, dst_offset)); + + WT sum = (WT)(0), xsum; + #pragma unroll + for (int y = 0; y < 4; y++) + { + src_index = mad24(y, src_step, src_index0); + if (sy + y >= 0 && sy + y < src_rows) + { + xsum = (WT)(0); + if (sx >= 0 && sx + 4 < src_cols) + { +#if depth == 0 && cn == 1 + uchar4 value = vload4(0, srcptr + src_index); + xsum = dot(convert_float4(value), (float4)(coeffs_x[0], coeffs_x[1], coeffs_x[2], coeffs_x[3])); +#else + #pragma unroll + for (int x = 0; x < 4; x++) + xsum = fma(convertToWT(loadpix(srcptr + mad24(x, pixsize, src_index))), coeffs_x[x], xsum); +#endif + } + else + { + #pragma unroll + for (int x = 0; x < 4; x++) + xsum = fma(sx + x >= 0 && sx + x < src_cols ? + convertToWT(loadpix(srcptr + mad24(x, pixsize, src_index))) : scalar, coeffs_x[x], xsum); + } + sum = fma(xsum, coeffs_y[y], sum); + } + else + sum = fma(scalar, coeffs_y[y], sum); + } + + storepix(convertToT(sum), dstptr + dst_index); #endif } } From 007593cab796a878cac5f24242b25f17daada932 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 4 Jun 2014 15:24:29 +0400 Subject: [PATCH 03/24] cvtColor - optimized index calculations; usage of build-in functions --- modules/imgproc/src/color.cpp | 16 +- modules/imgproc/src/opencl/cvtcolor.cl | 878 ++++++++++++++---------- modules/imgproc/test/ocl/test_color.cpp | 6 +- 3 files changed, 516 insertions(+), 384 deletions(-) diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp index 94401781ee..fe460ee75a 100644 --- a/modules/imgproc/src/color.cpp +++ b/modules/imgproc/src/color.cpp @@ -2730,8 +2730,6 @@ struct mRGBA2RGBA #ifdef HAVE_OPENCL -#define DIVUP(total, grain) (((total) + (grain) - 1) / (grain)) - static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) { bool ok = false; @@ -2739,23 +2737,17 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn ) Size sz = src.size(), dstSz = sz; int scn = src.channels(), depth = src.depth(), bidx; int dims = 2, stripeSize = 1; - size_t globalsize[] = { src.cols, src.rows }; ocl::Kernel k; if (depth != CV_8U && depth != CV_16U && depth != CV_32F) return false; - cv::String opts = format("-D depth=%d -D scn=%d ", depth, scn); - ocl::Device dev = ocl::Device::getDefault(); - int pxPerWIy = 1; - if (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU) && - !(code == CV_BGR2Luv || code == CV_RGB2Luv || code == CV_LBGR2Luv || code == CV_LRGB2Luv || - code == CV_Luv2BGR || code == CV_Luv2RGB || code == CV_Luv2LBGR || code == CV_Luv2LRGB)) - pxPerWIy = 4; + int pxPerWIy = dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU) ? 4 : 1; - globalsize[1] = DIVUP(globalsize[1], pxPerWIy); - opts += format("-D PIX_PER_WI_Y=%d ", pxPerWIy); + size_t globalsize[] = { src.cols, (src.rows + pxPerWIy - 1) / pxPerWIy }; + cv::String opts = format("-D depth=%d -D scn=%d -D PIX_PER_WI_Y=%d ", + depth, scn, pxPerWIy); switch (code) { diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 5bad3eee61..da835e08f3 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -71,10 +71,6 @@ #error "invalid depth: should be 0 (CV_8U), 2 (CV_16U) or 5 (CV_32F)" #endif -#ifndef STRIPE_SIZE -#define STRIPE_SIZE 1 -#endif - #define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n)) enum @@ -122,8 +118,8 @@ enum ///////////////////////////////////// RGB <-> GRAY ////////////////////////////////////// -__kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void RGB2Gray(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols) { int x = get_global_id(0); @@ -131,27 +127,32 @@ __kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index); DATA_TYPE_4 src_pix = vload4(0, src); #ifdef DEPTH_5 - dst[0] = src_pix.B_COMP * 0.114f + src_pix.G_COMP * 0.587f + src_pix.R_COMP * 0.299f; + dst[0] = fma(src_pix.B_COMP, 0.114f, fma(src_pix.G_COMP, 0.587f, src_pix.R_COMP * 0.299f)); #else - dst[0] = (DATA_TYPE)CV_DESCALE((src_pix.B_COMP * B2Y + src_pix.G_COMP * G2Y + src_pix.R_COMP * R2Y), yuv_shift); + dst[0] = (DATA_TYPE)CV_DESCALE(mad24(src_pix.B_COMP, B2Y, mad24(src_pix.G_COMP, G2Y, src_pix.R_COMP * R2Y)), yuv_shift); #endif + ++y; + src_index += src_step; + dst_index += dst_step; } - ++y; } } } -__kernel void Gray2RGB(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void Gray2RGB(__global const uchar * srcptr, int src_step, int src_offset, + __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols) { int x = get_global_id(0); @@ -159,20 +160,29 @@ __kernel void Gray2RGB(__global const uchar* srcptr, int srcstep, int srcoffset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index); DATA_TYPE val = src[0]; +#if dcn == 3 || defined DEPTH_5 dst[0] = dst[1] = dst[2] = val; #if dcn == 4 dst[3] = MAX_NUM; #endif +#else + *(__global DATA_TYPE_4 *)dst = (DATA_TYPE_4)(val, val, val, MAX_NUM); +#endif + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -182,8 +192,8 @@ __kernel void Gray2RGB(__global const uchar* srcptr, int srcstep, int srcoffset, __constant float c_RGB2YUVCoeffs_f[5] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f }; __constant int c_RGB2YUVCoeffs_i[5] = { B2Y, G2Y, R2Y, 8061, 14369 }; -__kernel void RGB2YUV(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void RGB2YUV(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, int dst_step, int dt_offset, int rows, int cols) { int x = get_global_id(0); @@ -191,34 +201,40 @@ __kernel void RGB2YUV(__global const uchar* srcptr, int srcstep, int srcoffset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dt_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index); DATA_TYPE_4 src_pix = vload4(0, src); - DATA_TYPE b=src_pix.B_COMP, g=src_pix.G_COMP, r=src_pix.R_COMP; + DATA_TYPE b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP; #ifdef DEPTH_5 __constant float * coeffs = c_RGB2YUVCoeffs_f; - const DATA_TYPE Y = b * coeffs[0] + g * coeffs[1] + r * coeffs[2]; - const DATA_TYPE U = (b - Y) * coeffs[3] + HALF_MAX; - const DATA_TYPE V = (r - Y) * coeffs[4] + HALF_MAX; + const DATA_TYPE Y = fma(b, coeffs[0], fma(g, coeffs[1], r * coeffs[2])); + const DATA_TYPE U = fma(b - Y, coeffs[3], HALF_MAX); + const DATA_TYPE V = fma(r - Y, coeffs[4], HALF_MAX); #else __constant int * coeffs = c_RGB2YUVCoeffs_i; const int delta = HALF_MAX * (1 << yuv_shift); - const int Y = CV_DESCALE(b * coeffs[0] + g * coeffs[1] + r * coeffs[2], yuv_shift); - const int U = CV_DESCALE((b - Y) * coeffs[3] + delta, yuv_shift); - const int V = CV_DESCALE((r - Y) * coeffs[4] + delta, yuv_shift); + const int Y = CV_DESCALE(mad24(b, coeffs[0], mad24(g, coeffs[1], r * coeffs[2])), yuv_shift); + const int U = CV_DESCALE(mad24(b - Y, coeffs[3], delta), yuv_shift); + const int V = CV_DESCALE(mad24(r - Y, coeffs[4], delta), yuv_shift); #endif dst[0] = SAT_CAST( Y ); dst[1] = SAT_CAST( U ); dst[2] = SAT_CAST( V ); + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -226,8 +242,8 @@ __kernel void RGB2YUV(__global const uchar* srcptr, int srcstep, int srcoffset, __constant float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f }; __constant int c_YUV2RGBCoeffs_i[5] = { 33292, -6472, -9519, 18678 }; -__kernel void YUV2RGB(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void YUV2RGB(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, int dst_step, int dt_offset, int rows, int cols) { int x = get_global_id(0); @@ -235,25 +251,28 @@ __kernel void YUV2RGB(__global const uchar* srcptr, int srcstep, int srcoffset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dt_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index); DATA_TYPE_4 src_pix = vload4(0, src); DATA_TYPE Y = src_pix.x, U = src_pix.y, V = src_pix.z; #ifdef DEPTH_5 __constant float * coeffs = c_YUV2RGBCoeffs_f; - const float r = Y + (V - HALF_MAX) * coeffs[3]; - const float g = Y + (V - HALF_MAX) * coeffs[2] + (U - HALF_MAX) * coeffs[1]; - const float b = Y + (U - HALF_MAX) * coeffs[0]; + float r = fma(V - HALF_MAX, coeffs[3], Y); + float g = fma(V - HALF_MAX, coeffs[2], fma(U - HALF_MAX, coeffs[1], Y)); + float b = fma(U - HALF_MAX, coeffs[0], Y); #else __constant int * coeffs = c_YUV2RGBCoeffs_i; const int r = Y + CV_DESCALE((V - HALF_MAX) * coeffs[3], yuv_shift); - const int g = Y + CV_DESCALE((V - HALF_MAX) * coeffs[2] + (U - HALF_MAX) * coeffs[1], yuv_shift); + const int g = Y + CV_DESCALE(mad24(V - HALF_MAX, coeffs[2], (U - HALF_MAX) * coeffs[1]), yuv_shift); const int b = Y + CV_DESCALE((U - HALF_MAX) * coeffs[0], yuv_shift); #endif @@ -263,8 +282,10 @@ __kernel void YUV2RGB(__global const uchar* srcptr, int srcstep, int srcoffset, #if dcn == 4 dst[3] = MAX_NUM; #endif + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -276,8 +297,8 @@ __constant int ITUR_BT_601_CVG = 852492; __constant int ITUR_BT_601_CVR = 1673527; __constant int ITUR_BT_601_SHIFT = 20; -__kernel void YUV2RGB_NV12(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void YUV2RGB_NV12(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, int dst_step, int dt_offset, int rows, int cols) { int x = get_global_id(0); @@ -290,15 +311,15 @@ __kernel void YUV2RGB_NV12(__global const uchar* srcptr, int srcstep, int srcoff { if (y < rows / 2 ) { - __global const uchar* ysrc = srcptr + mad24(y << 1, srcstep, (x << 1) + srcoffset); - __global const uchar* usrc = srcptr + mad24(rows + y, srcstep, (x << 1) + srcoffset); - __global uchar* dst1 = dstptr + mad24(y << 1, dststep, x * (dcn<<1) + dstoffset); - __global uchar* dst2 = dstptr + mad24((y << 1) + 1, dststep, x * (dcn<<1) + dstoffset); + __global const uchar* ysrc = srcptr + mad24(y << 1, src_step, (x << 1) + src_offset); + __global const uchar* usrc = srcptr + mad24(rows + y, src_step, (x << 1) + src_offset); + __global uchar* dst1 = dstptr + mad24(y << 1, dst_step, x * (dcn<<1) + dt_offset); + __global uchar* dst2 = dstptr + mad24((y << 1) + 1, dst_step, x * (dcn<<1) + dt_offset); int Y1 = ysrc[0]; int Y2 = ysrc[1]; - int Y3 = ysrc[srcstep]; - int Y4 = ysrc[srcstep + 1]; + int Y3 = ysrc[src_step]; + int Y4 = ysrc[src_step + 1]; int U = usrc[0] - 128; int V = usrc[1] - 128; @@ -349,8 +370,8 @@ __kernel void YUV2RGB_NV12(__global const uchar* srcptr, int srcstep, int srcoff __constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f}; __constant int c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241}; -__kernel void RGB2YCrCb(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +__kernel void RGB2YCrCb(__global const uchar* srcptr, int src_step, int src_offset, + __global uchar* dstptr, int dst_step, int dt_offset, int rows, int cols) { int x = get_global_id(0); @@ -358,34 +379,40 @@ __kernel void RGB2YCrCb(__global const uchar* srcptr, int srcstep, int srcoffset if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dt_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index); DATA_TYPE_4 src_pix = vload4(0, src); - DATA_TYPE b=src_pix.B_COMP, g=src_pix.G_COMP, r=src_pix.R_COMP; + DATA_TYPE b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP; #ifdef DEPTH_5 __constant float * coeffs = c_RGB2YCrCbCoeffs_f; - DATA_TYPE Y = b * coeffs[2] + g * coeffs[1] + r * coeffs[0]; - DATA_TYPE Cr = (r - Y) * coeffs[3] + HALF_MAX; - DATA_TYPE Cb = (b - Y) * coeffs[4] + HALF_MAX; + DATA_TYPE Y = fma(b, coeffs[2], fma(g, coeffs[1], r * coeffs[0])); + DATA_TYPE Cr = fma(r - Y, coeffs[3], HALF_MAX); + DATA_TYPE Cb = fma(b - Y, coeffs[4], HALF_MAX); #else __constant int * coeffs = c_RGB2YCrCbCoeffs_i; int delta = HALF_MAX * (1 << yuv_shift); - int Y = CV_DESCALE(b * coeffs[2] + g * coeffs[1] + r * coeffs[0], yuv_shift); - int Cr = CV_DESCALE((r - Y) * coeffs[3] + delta, yuv_shift); - int Cb = CV_DESCALE((b - Y) * coeffs[4] + delta, yuv_shift); + int Y = CV_DESCALE(mad24(b, coeffs[2], mad24(g, coeffs[1], r * coeffs[0])), yuv_shift); + int Cr = CV_DESCALE(mad24(r - Y, coeffs[3], delta), yuv_shift); + int Cb = CV_DESCALE(mad24(b - Y, coeffs[4], delta), yuv_shift); #endif dst[0] = SAT_CAST( Y ); dst[1] = SAT_CAST( Cr ); dst[2] = SAT_CAST( Cb ); + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -402,28 +429,29 @@ __kernel void YCrCb2RGB(__global const uchar* src, int src_step, int src_offset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - __global const DATA_TYPE * srcptr = (__global const DATA_TYPE*)(src + src_idx); - __global DATA_TYPE * dstptr = (__global DATA_TYPE*)(dst + dst_idx); + __global const DATA_TYPE * srcptr = (__global const DATA_TYPE*)(src + src_index); + __global DATA_TYPE * dstptr = (__global DATA_TYPE*)(dst + dst_index); DATA_TYPE_4 src_pix = vload4(0, srcptr); DATA_TYPE y = src_pix.x, cr = src_pix.y, cb = src_pix.z; #ifdef DEPTH_5 __constant float * coeff = c_YCrCb2RGBCoeffs_f; - float r = y + coeff[0] * (cr - HALF_MAX); - float g = y + coeff[1] * (cr - HALF_MAX) + coeff[2] * (cb - HALF_MAX); - float b = y + coeff[3] * (cb - HALF_MAX); + float r = fma(coeff[0], cr - HALF_MAX, y); + float g = fma(coeff[1], cr - HALF_MAX, fma(coeff[2], cb - HALF_MAX, y)); + float b = fma(coeff[3], cb - HALF_MAX, y); #else __constant int * coeff = c_YCrCb2RGBCoeffs_i; int r = y + CV_DESCALE(coeff[0] * (cr - HALF_MAX), yuv_shift); - int g = y + CV_DESCALE(coeff[1] * (cr - HALF_MAX) + coeff[2] * (cb - HALF_MAX), yuv_shift); + int g = y + CV_DESCALE(mad24(coeff[1], cr - HALF_MAX, coeff[2] * (cb - HALF_MAX)), yuv_shift); int b = y + CV_DESCALE(coeff[3] * (cb - HALF_MAX), yuv_shift); #endif @@ -433,8 +461,11 @@ __kernel void YCrCb2RGB(__global const uchar* src, int src_step, int src_offset, #if dcn == 4 dstptr[3] = MAX_NUM; #endif + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -450,34 +481,37 @@ __kernel void RGB2XYZ(__global const uchar * srcptr, int src_step, int src_offse if (dx < cols) { + int src_index = mad24(dy, src_step, mad24(dx, scnbytes, src_offset)); + int dst_index = mad24(dy, dst_step, mad24(dx, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (dy < rows) { - int src_idx = mad24(dy, src_step, src_offset + dx * scnbytes); - int dst_idx = mad24(dy, dst_step, dst_offset + dx * dcnbytes); - - __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_idx); - __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_idx); + __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_index); + __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_index); DATA_TYPE_4 src_pix = vload4(0, src); DATA_TYPE r = src_pix.x, g = src_pix.y, b = src_pix.z; #ifdef DEPTH_5 - float x = r * coeffs[0] + g * coeffs[1] + b * coeffs[2]; - float y = r * coeffs[3] + g * coeffs[4] + b * coeffs[5]; - float z = r * coeffs[6] + g * coeffs[7] + b * coeffs[8]; + float x = fma(r, coeffs[0], fma(g, coeffs[1], b * coeffs[2])); + float y = fma(r, coeffs[3], fma(g, coeffs[4], b * coeffs[5])); + float z = fma(r, coeffs[6], fma(g, coeffs[7], b * coeffs[8])); #else - int x = CV_DESCALE(r * coeffs[0] + g * coeffs[1] + b * coeffs[2], xyz_shift); - int y = CV_DESCALE(r * coeffs[3] + g * coeffs[4] + b * coeffs[5], xyz_shift); - int z = CV_DESCALE(r * coeffs[6] + g * coeffs[7] + b * coeffs[8], xyz_shift); + int x = CV_DESCALE(mad24(r, coeffs[0], mad24(g, coeffs[1], b * coeffs[2])), xyz_shift); + int y = CV_DESCALE(mad24(r, coeffs[3], mad24(g, coeffs[4], b * coeffs[5])), xyz_shift); + int z = CV_DESCALE(mad24(r, coeffs[6], mad24(g, coeffs[7], b * coeffs[8])), xyz_shift); #endif dst[0] = SAT_CAST(x); dst[1] = SAT_CAST(y); dst[2] = SAT_CAST(z); + + ++dy; + dst_index += dst_step; + src_index += src_step; } - ++dy; } } } @@ -491,37 +525,48 @@ __kernel void XYZ2RGB(__global const uchar * srcptr, int src_step, int src_offse if (dx < cols) { + int src_index = mad24(dy, src_step, mad24(dx, scnbytes, src_offset)); + int dst_index = mad24(dy, dst_step, mad24(dx, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (dy < rows) { - int src_idx = mad24(dy, src_step, src_offset + dx * scnbytes); - int dst_idx = mad24(dy, dst_step, dst_offset + dx * dcnbytes); - - __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_idx); - __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_idx); + __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_index); + __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_index); DATA_TYPE_4 src_pix = vload4(0, src); DATA_TYPE x = src_pix.x, y = src_pix.y, z = src_pix.z; #ifdef DEPTH_5 - float b = x * coeffs[0] + y * coeffs[1] + z * coeffs[2]; - float g = x * coeffs[3] + y * coeffs[4] + z * coeffs[5]; - float r = x * coeffs[6] + y * coeffs[7] + z * coeffs[8]; + float b = fma(x, coeffs[0], fma(y, coeffs[1], z * coeffs[2])); + float g = fma(x, coeffs[3], fma(y, coeffs[4], z * coeffs[5])); + float r = fma(x, coeffs[6], fma(y, coeffs[7], z * coeffs[8])); #else - int b = CV_DESCALE(x * coeffs[0] + y * coeffs[1] + z * coeffs[2], xyz_shift); - int g = CV_DESCALE(x * coeffs[3] + y * coeffs[4] + z * coeffs[5], xyz_shift); - int r = CV_DESCALE(x * coeffs[6] + y * coeffs[7] + z * coeffs[8], xyz_shift); + int b = CV_DESCALE(mad24(x, coeffs[0], mad24(y, coeffs[1], z * coeffs[2])), xyz_shift); + int g = CV_DESCALE(mad24(x, coeffs[3], mad24(y, coeffs[4], z * coeffs[5])), xyz_shift); + int r = CV_DESCALE(mad24(x, coeffs[6], mad24(y, coeffs[7], z * coeffs[8])), xyz_shift); #endif - dst[0] = SAT_CAST(b); - dst[1] = SAT_CAST(g); - dst[2] = SAT_CAST(r); + + DATA_TYPE dst0 = SAT_CAST(b); + DATA_TYPE dst1 = SAT_CAST(g); + DATA_TYPE dst2 = SAT_CAST(r); +#if dcn == 3 || defined DEPTH_5 + dst[0] = dst0; + dst[1] = dst1; + dst[2] = dst2; #if dcn == 4 dst[3] = MAX_NUM; #endif +#else + *(__global DATA_TYPE_4 *)dst = (DATA_TYPE_4)(dst0, dst1, dst2, MAX_NUM); +#endif + + ++dy; + dst_index += dst_step; + src_index += src_step; } - ++dy; } } } @@ -537,16 +582,16 @@ __kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - - __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_idx); - __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_idx); + __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_index); + __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_index); DATA_TYPE_4 src_pix = vload4(0, src); #ifdef REVERSE @@ -566,8 +611,11 @@ __kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset, dst[3] = src[3]; #endif #endif + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -583,34 +631,38 @@ __kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - ushort t = *((__global const ushort*)(src + src_idx)); + ushort t = *((__global const ushort*)(src + src_index)); #if greenbits == 6 - dst[dst_idx + bidx] = (uchar)(t << 3); - dst[dst_idx + 1] = (uchar)((t >> 3) & ~3); - dst[dst_idx + (bidx^2)] = (uchar)((t >> 8) & ~7); + dst[dst_index + bidx] = (uchar)(t << 3); + dst[dst_index + 1] = (uchar)((t >> 3) & ~3); + dst[dst_index + (bidx^2)] = (uchar)((t >> 8) & ~7); #else - dst[dst_idx + bidx] = (uchar)(t << 3); - dst[dst_idx + 1] = (uchar)((t >> 2) & ~7); - dst[dst_idx + (bidx^2)] = (uchar)((t >> 7) & ~7); + dst[dst_index + bidx] = (uchar)(t << 3); + dst[dst_index + 1] = (uchar)((t >> 2) & ~7); + dst[dst_index + (bidx^2)] = (uchar)((t >> 7) & ~7); #endif #if dcn == 4 #if greenbits == 6 - dst[dst_idx + 3] = 255; + dst[dst_index + 3] = 255; #else - dst[dst_idx + 3] = t & 0x8000 ? 255 : 0; + dst[dst_index + 3] = t & 0x8000 ? 255 : 0; #endif #endif + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -624,25 +676,29 @@ __kernel void RGB2RGB5x5(__global const uchar* src, int src_step, int src_offset if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - uchar4 src_pix = vload4(0, src + src_idx); + uchar4 src_pix = vload4(0, src + src_index); #if greenbits == 6 - *((__global ushort*)(dst + dst_idx)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~3) << 3)|((src_pix.R_COMP&~7) << 8)); + *((__global ushort*)(dst + dst_index)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~3) << 3)|((src_pix.R_COMP&~7) << 8)); #elif scn == 3 - *((__global ushort*)(dst + dst_idx)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~7) << 2)|((src_pix.R_COMP&~7) << 7)); + *((__global ushort*)(dst + dst_index)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~7) << 2)|((src_pix.R_COMP&~7) << 7)); #else - *((__global ushort*)(dst + dst_idx)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~7) << 2)| + *((__global ushort*)(dst + dst_index)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~7) << 2)| ((src_pix.R_COMP&~7) << 7)|(src_pix.w ? 0x8000 : 0)); #endif + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -658,26 +714,25 @@ __kernel void BGR5x52Gray(__global const uchar* src, int src_step, int src_offse if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, dst_offset + x); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x); - int t = *((__global const ushort*)(src + src_idx)); + int t = *((__global const ushort*)(src + src_index)); #if greenbits == 6 - dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + - ((t >> 3) & 0xfc)*G2Y + - ((t >> 8) & 0xf8)*R2Y, yuv_shift); + dst[dst_index] = (uchar)CV_DESCALE(mad24((t << 3) & 0xf8, B2Y, mad24((t >> 3) & 0xfc, G2Y, ((t >> 8) & 0xf8) * R2Y)), yuv_shift); #else - dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + - ((t >> 2) & 0xf8)*G2Y + - ((t >> 7) & 0xf8)*R2Y, yuv_shift); + dst[dst_index] = (uchar)CV_DESCALE(mad24((t << 3) & 0xf8, B2Y, mad24((t >> 2) & 0xf8, G2Y, ((t >> 7) & 0xf8) * R2Y)), yuv_shift); #endif + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -691,23 +746,26 @@ __kernel void Gray2BGR5x5(__global const uchar* src, int src_step, int src_offse if (x < cols) { + int src_index = mad24(y, src_step, src_offset + x); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - int t = src[src_idx]; + int t = src[src_index]; #if greenbits == 6 - *((__global ushort*)(dst + dst_idx)) = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8)); + *((__global ushort*)(dst + dst_index)) = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8)); #else t >>= 3; - *((__global ushort*)(dst + dst_idx)) = (ushort)(t|(t << 5)|(t << 10)); + *((__global ushort*)(dst + dst_index)) = (ushort)(t|(t << 5)|(t << 10)); #endif + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -733,40 +791,44 @@ __kernel void RGB2HSV(__global const uchar* src, int src_step, int src_offset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - uchar4 src_pix = vload4(0, src + src_idx); + uchar4 src_pix = vload4(0, src + src_index); int b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP; int h, s, v = b; int vmin = b, diff; int vr, vg; - v = max( v, g ); - v = max( v, r ); - vmin = min( vmin, g ); - vmin = min( vmin, r ); + v = max(v, g); + v = max(v, r); + vmin = min(vmin, g); + vmin = min(vmin, r); diff = v - vmin; vr = v == r ? -1 : 0; vg = v == g ? -1 : 0; - s = (diff * sdiv_table[v] + (1 << (hsv_shift-1))) >> hsv_shift; + s = mad24(diff, sdiv_table[v], (1 << (hsv_shift-1))) >> hsv_shift; h = (vr & (g - b)) + - (~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff)))); - h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift; + (~vr & ((vg & mad24(diff, 2, b - r)) + ((~vg) & mad24(4, diff, r - g)))); + h = mad24(h, hdiv_table[diff], (1 << (hsv_shift-1))) >> hsv_shift; h += h < 0 ? hrange : 0; - dst[dst_idx] = convert_uchar_sat_rte(h); - dst[dst_idx + 1] = (uchar)s; - dst[dst_idx + 2] = (uchar)v; + dst[dst_index] = convert_uchar_sat_rte(h); + dst[dst_index + 1] = (uchar)s; + dst[dst_index + 2] = (uchar)v; + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -780,14 +842,15 @@ __kernel void HSV2RGB(__global const uchar* src, int src_step, int src_offset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - uchar4 src_pix = vload4(0, src + src_idx); + uchar4 src_pix = vload4(0, src + src_index); float h = src_pix.x, s = src_pix.y*(1/255.f), v = src_pix.z*(1/255.f); float b, g, r; @@ -821,14 +884,17 @@ __kernel void HSV2RGB(__global const uchar* src, int src_step, int src_offset, else b = g = r = v; - dst[dst_idx + bidx] = convert_uchar_sat_rte(b*255.f); - dst[dst_idx + 1] = convert_uchar_sat_rte(g*255.f); - dst[dst_idx + (bidx^2)] = convert_uchar_sat_rte(r*255.f); + dst[dst_index + bidx] = convert_uchar_sat_rte(b*255.f); + dst[dst_index + 1] = convert_uchar_sat_rte(g*255.f); + dst[dst_index + (bidx^2)] = convert_uchar_sat_rte(r*255.f); #if dcn == 4 - dst[dst_idx + 3] = MAX_NUM; + dst[dst_index + 3] = MAX_NUM; #endif + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -844,16 +910,16 @@ __kernel void RGB2HSV(__global const uchar* srcptr, int src_step, int src_offset if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - - __global const float * src = (__global const float *)(srcptr + src_idx); - __global float * dst = (__global float *)(dstptr + dst_idx); + __global const float * src = (__global const float *)(srcptr + src_index); + __global float * dst = (__global float *)(dstptr + dst_index); float4 src_pix = vload4(0, src); float b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP; @@ -873,17 +939,21 @@ __kernel void RGB2HSV(__global const uchar* srcptr, int src_step, int src_offset if( v == r ) h = (g - b)*diff; else if( v == g ) - h = (b - r)*diff + 120.f; + h = fma(b - r, diff, 120.f); else - h = (r - g)*diff + 240.f; + h = fma(r - g, diff, 240.f); - if( h < 0 ) h += 360.f; + if( h < 0 ) + h += 360.f; dst[0] = h*hscale; dst[1] = s; dst[2] = v; + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -897,16 +967,17 @@ __kernel void HSV2RGB(__global const uchar* srcptr, int src_step, int src_offset if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - __global const float * src = (__global const float *)(srcptr + src_idx); - __global float * dst = (__global float *)(dstptr + dst_idx); + __global const float * src = (__global const float *)(srcptr + src_index); + __global float * dst = (__global float *)(dstptr + dst_index); float4 src_pix = vload4(0, src); float h = src_pix.x, s = src_pix.y, v = src_pix.z; @@ -947,8 +1018,11 @@ __kernel void HSV2RGB(__global const uchar* srcptr, int src_step, int src_offset #if dcn == 4 dst[3] = MAX_NUM; #endif + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -968,14 +1042,15 @@ __kernel void RGB2HLS(__global const uchar* src, int src_step, int src_offset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - uchar4 src_pix = vload4(0, src + src_idx); + uchar4 src_pix = vload4(0, src + src_index); float b = src_pix.B_COMP*(1/255.f), g = src_pix.G_COMP*(1/255.f), r = src_pix.R_COMP*(1/255.f); float h = 0.f, s = 0.f, l; @@ -998,18 +1073,22 @@ __kernel void RGB2HLS(__global const uchar* src, int src_step, int src_offset, if( vmax == r ) h = (g - b)*diff; else if( vmax == g ) - h = (b - r)*diff + 120.f; + h = fma(b - r, diff, 120.f); else - h = (r - g)*diff + 240.f; + h = fma(r - g, diff, 240.f); - if( h < 0.f ) h += 360.f; + if( h < 0.f ) + h += 360.f; } - dst[dst_idx] = convert_uchar_sat_rte(h*hscale); - dst[dst_idx + 1] = convert_uchar_sat_rte(l*255.f); - dst[dst_idx + 2] = convert_uchar_sat_rte(s*255.f); + dst[dst_index] = convert_uchar_sat_rte(h*hscale); + dst[dst_index + 1] = convert_uchar_sat_rte(l*255.f); + dst[dst_index + 2] = convert_uchar_sat_rte(s*255.f); + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1023,14 +1102,15 @@ __kernel void HLS2RGB(__global const uchar* src, int src_step, int src_offset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - uchar4 src_pix = vload4(0, src + src_idx); + uchar4 src_pix = vload4(0, src + src_index); float h = src_pix.x, l = src_pix.y*(1.f/255.f), s = src_pix.z*(1.f/255.f); float b, g, r; @@ -1053,8 +1133,8 @@ __kernel void HLS2RGB(__global const uchar* src, int src_step, int src_offset, tab[0] = p2; tab[1] = p1; - tab[2] = p1 + (p2 - p1)*(1-h); - tab[3] = p1 + (p2 - p1)*h; + tab[2] = fma(p2 - p1, 1-h, p1); + tab[3] = fma(p2 - p1, h, p1); b = tab[sector_data[sector][0]]; g = tab[sector_data[sector][1]]; @@ -1063,14 +1143,17 @@ __kernel void HLS2RGB(__global const uchar* src, int src_step, int src_offset, else b = g = r = l; - dst[dst_idx + bidx] = convert_uchar_sat_rte(b*255.f); - dst[dst_idx + 1] = convert_uchar_sat_rte(g*255.f); - dst[dst_idx + (bidx^2)] = convert_uchar_sat_rte(r*255.f); + dst[dst_index + bidx] = convert_uchar_sat_rte(b*255.f); + dst[dst_index + 1] = convert_uchar_sat_rte(g*255.f); + dst[dst_index + (bidx^2)] = convert_uchar_sat_rte(r*255.f); #if dcn == 4 - dst[dst_idx + 3] = MAX_NUM; + dst[dst_index + 3] = MAX_NUM; #endif + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1086,16 +1169,16 @@ __kernel void RGB2HLS(__global const uchar* srcptr, int src_step, int src_offset if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - - __global const float * src = (__global const float *)(srcptr + src_idx); - __global float * dst = (__global float *)(dstptr + dst_idx); + __global const float * src = (__global const float *)(srcptr + src_index); + __global float * dst = (__global float *)(dstptr + dst_index); float4 src_pix = vload4(0, src); float b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP; @@ -1119,9 +1202,9 @@ __kernel void RGB2HLS(__global const uchar* srcptr, int src_step, int src_offset if( vmax == r ) h = (g - b)*diff; else if( vmax == g ) - h = (b - r)*diff + 120.f; + h = fma(b - r, diff, 120.f); else - h = (r - g)*diff + 240.f; + h = fma(r - g, diff, 240.f); if( h < 0.f ) h += 360.f; } @@ -1129,8 +1212,11 @@ __kernel void RGB2HLS(__global const uchar* srcptr, int src_step, int src_offset dst[0] = h*hscale; dst[1] = l; dst[2] = s; + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1144,16 +1230,16 @@ __kernel void HLS2RGB(__global const uchar* srcptr, int src_step, int src_offset if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - - __global const float * src = (__global const float *)(srcptr + src_idx); - __global float * dst = (__global float *)(dstptr + dst_idx); + __global const float * src = (__global const float *)(srcptr + src_index); + __global float * dst = (__global float *)(dstptr + dst_index); float4 src_pix = vload4(0, src); float h = src_pix.x, l = src_pix.y, s = src_pix.z; @@ -1178,8 +1264,8 @@ __kernel void HLS2RGB(__global const uchar* srcptr, int src_step, int src_offset tab[0] = p2; tab[1] = p1; - tab[2] = p1 + (p2 - p1)*(1-h); - tab[3] = p1 + (p2 - p1)*h; + tab[2] = fma(p2 - p1, 1-h, p1); + tab[3] = fma(p2 - p1, h, p1); b = tab[sector_data[sector][0]]; g = tab[sector_data[sector][1]]; @@ -1194,8 +1280,11 @@ __kernel void HLS2RGB(__global const uchar* srcptr, int src_step, int src_offset #if dcn == 4 dst[3] = MAX_NUM; #endif + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1215,24 +1304,25 @@ __kernel void RGBA2mRGBA(__global const uchar* src, int src_step, int src_offset if (x < cols) { + int src_index = mad24(y, src_step, src_offset + (x << 2)); + int dst_index = mad24(y, dst_step, dst_offset + (x << 2)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + (x << 2)); - int dst_idx = mad24(y, dst_step, dst_offset + (x << 2)); - uchar4 src_pix = vload4(0, src + src_idx); + uchar4 src_pix = *(__global const uchar4 *)(src + src_index); - uchar v0 = src_pix.x, v1 = src_pix.y; - uchar v2 = src_pix.z, v3 = src_pix.w; + *(__global uchar4 *)(dst + dst_index) = + (uchar4)(mad24(src_pix.x, src_pix.w, HALF_MAX) / MAX_NUM, + mad24(src_pix.y, src_pix.w, HALF_MAX) / MAX_NUM, + mad24(src_pix.z, src_pix.w, HALF_MAX) / MAX_NUM, src_pix.w); - dst[dst_idx] = (v0 * v3 + HALF_MAX) / MAX_NUM; - dst[dst_idx + 1] = (v1 * v3 + HALF_MAX) / MAX_NUM; - dst[dst_idx + 2] = (v2 * v3 + HALF_MAX) / MAX_NUM; - dst[dst_idx + 3] = v3; + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1246,25 +1336,29 @@ __kernel void mRGBA2RGBA(__global const uchar* src, int src_step, int src_offset if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, 4, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, 4, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + (x << 2)); - int dst_idx = mad24(y, dst_step, dst_offset + (x << 2)); - uchar4 src_pix = vload4(0, src + src_idx); + uchar4 src_pix = *(__global const uchar4 *)(src + src_index); + uchar v3 = src_pix.w, v3_half = v3 / 2; - uchar v0 = src_pix.x, v1 = src_pix.y; - uchar v2 = src_pix.z, v3 = src_pix.w; - uchar v3_half = v3 / 2; + if (v3 == 0) + *(__global uchar4 *)(dst + dst_index) = (uchar4)(0, 0, 0, 0); + else + *(__global uchar4 *)(dst + dst_index) = + (uchar4)(mad24(src_pix.x, MAX_NUM, v3_half) / v3, + mad24(src_pix.y, MAX_NUM, v3_half) / v3, + mad24(src_pix.z, MAX_NUM, v3_half) / v3, v3); - dst[dst_idx] = v3 == 0 ? 0 : (v0 * MAX_NUM + v3_half) / v3; - dst[dst_idx + 1] = v3 == 0 ? 0 : (v1 * MAX_NUM + v3_half) / v3; - dst[dst_idx + 2] = v3 == 0 ? 0 : (v2 * MAX_NUM + v3_half) / v3; - dst[dst_idx + 3] = v3; + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1283,8 +1377,8 @@ inline float splineInterpolate(float x, __global const float * tab, int n) { int ix = clamp(convert_int_sat_rtn(x), 0, n-1); x -= ix; - tab += ix*4; - return ((tab[3]*x + tab[2])*x + tab[1])*x + tab[0]; + tab += ix << 2; + return fma(fma(fma(tab[3], x, tab[2]), x, tab[1]), x, tab[0]); } #ifdef DEPTH_0 @@ -1299,16 +1393,16 @@ __kernel void BGR2Lab(__global const uchar * src, int src_step, int src_offset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - - __global const uchar* src_ptr = src + src_idx; - __global uchar* dst_ptr = dst + dst_idx; + __global const uchar* src_ptr = src + src_index; + __global uchar* dst_ptr = dst + dst_index; uchar4 src_pix = vload4(0, src_ptr); int C0 = coeffs[0], C1 = coeffs[1], C2 = coeffs[2], @@ -1316,19 +1410,22 @@ __kernel void BGR2Lab(__global const uchar * src, int src_step, int src_offset, C6 = coeffs[6], C7 = coeffs[7], C8 = coeffs[8]; int R = gammaTab[src_pix.x], G = gammaTab[src_pix.y], B = gammaTab[src_pix.z]; - int fX = LabCbrtTab_b[CV_DESCALE(R*C0 + G*C1 + B*C2, lab_shift)]; - int fY = LabCbrtTab_b[CV_DESCALE(R*C3 + G*C4 + B*C5, lab_shift)]; - int fZ = LabCbrtTab_b[CV_DESCALE(R*C6 + G*C7 + B*C8, lab_shift)]; + int fX = LabCbrtTab_b[CV_DESCALE(mad24(R, C0, mad24(G, C1, B*C2)), lab_shift)]; + int fY = LabCbrtTab_b[CV_DESCALE(mad24(R, C3, mad24(G, C4, B*C5)), lab_shift)]; + int fZ = LabCbrtTab_b[CV_DESCALE(mad24(R, C6, mad24(G, C7, B*C8)), lab_shift)]; int L = CV_DESCALE( Lscale*fY + Lshift, lab_shift2 ); - int a = CV_DESCALE( 500*(fX - fY) + 128*(1 << lab_shift2), lab_shift2 ); - int b = CV_DESCALE( 200*(fY - fZ) + 128*(1 << lab_shift2), lab_shift2 ); + int a = CV_DESCALE( mad24(500, fX - fY, 128*(1 << lab_shift2)), lab_shift2 ); + int b = CV_DESCALE( mad24(200, fY - fZ, 128*(1 << lab_shift2)), lab_shift2 ); dst_ptr[0] = SAT_CAST(L); dst_ptr[1] = SAT_CAST(a); dst_ptr[2] = SAT_CAST(b); + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1347,16 +1444,16 @@ __kernel void BGR2Lab(__global const uchar * srcptr, int src_step, int src_offse if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - - __global const float * src = (__global const float *)(srcptr + src_idx); - __global float * dst = (__global float *)(dstptr + dst_idx); + __global const float * src = (__global const float *)(srcptr + src_index); + __global float * dst = (__global float *)(dstptr + dst_index); float4 src_pix = vload4(0, src); float C0 = coeffs[0], C1 = coeffs[1], C2 = coeffs[2], @@ -1373,23 +1470,26 @@ __kernel void BGR2Lab(__global const uchar * srcptr, int src_step, int src_offse B = splineInterpolate(B * GammaTabScale, gammaTab, GAMMA_TAB_SIZE); #endif - float X = R*C0 + G*C1 + B*C2; - float Y = R*C3 + G*C4 + B*C5; - float Z = R*C6 + G*C7 + B*C8; + float X = fma(R, C0, fma(G, C1, B*C2)); + float Y = fma(R, C3, fma(G, C4, B*C5)); + float Z = fma(R, C6, fma(G, C7, B*C8)); - float FX = X > 0.008856f ? pow(X, _1_3) : (7.787f * X + _a); - float FY = Y > 0.008856f ? pow(Y, _1_3) : (7.787f * Y + _a); - float FZ = Z > 0.008856f ? pow(Z, _1_3) : (7.787f * Z + _a); + float FX = X > 0.008856f ? rootn(X, 3) : fma(7.787f, X, _a); + float FY = Y > 0.008856f ? rootn(Y, 3) : fma(7.787f, Y, _a); + float FZ = Z > 0.008856f ? rootn(Z, 3) : fma(7.787f, Z, _a); - float L = Y > 0.008856f ? (116.f * FY - 16.f) : (903.3f * Y); + float L = Y > 0.008856f ? fma(116.f, FY, -16.f) : (903.3f * Y); float a = 500.f * (FX - FY); float b = 200.f * (FY - FZ); dst[0] = L; dst[1] = a; dst[2] = b; + + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1412,7 +1512,7 @@ inline void Lab2BGR_f(const float * srcbuf, float * dstbuf, if (li <= lThresh) { y = li / 903.3f; - fy = 7.787f * y + 16.0f / 116.0f; + fy = fma(7.787f, y, 16.0f / 116.0f); } else { @@ -1422,6 +1522,7 @@ inline void Lab2BGR_f(const float * srcbuf, float * dstbuf, float fxz[] = { ai / 500.0f + fy, fy - bi / 200.0f }; + #pragma unroll for (int j = 0; j < 2; j++) if (fxz[j] <= fThresh) fxz[j] = (fxz[j] - 16.0f / 116.0f) / 7.787f; @@ -1429,9 +1530,9 @@ inline void Lab2BGR_f(const float * srcbuf, float * dstbuf, fxz[j] = fxz[j] * fxz[j] * fxz[j]; float x = fxz[0], z = fxz[1]; - float ro = clamp(C0 * x + C1 * y + C2 * z, 0.0f, 1.0f); - float go = clamp(C3 * x + C4 * y + C5 * z, 0.0f, 1.0f); - float bo = clamp(C6 * x + C7 * y + C8 * z, 0.0f, 1.0f); + float ro = clamp(fma(C0, x, fma(C1, y, C2 * z)), 0.0f, 1.0f); + float go = clamp(fma(C3, x, fma(C4, y, C5 * z)), 0.0f, 1.0f); + float bo = clamp(fma(C6, x, fma(C7, y, C8 * z)), 0.0f, 1.0f); #ifdef SRGB ro = splineInterpolate(ro * GammaTabScale, gammaTab, GAMMA_TAB_SIZE); @@ -1456,16 +1557,16 @@ __kernel void Lab2BGR(__global const uchar * src, int src_step, int src_offset, if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - - __global const uchar* src_ptr = src + src_idx; - __global uchar* dst_ptr = dst + dst_idx; + __global const uchar* src_ptr = src + src_index; + __global uchar * dst_ptr = dst + dst_index; uchar4 src_pix = vload4(0, src_ptr); float srcbuf[3], dstbuf[3]; @@ -1479,14 +1580,18 @@ __kernel void Lab2BGR(__global const uchar * src, int src_step, int src_offset, #endif coeffs, lThresh, fThresh); +#if dcn == 3 dst_ptr[0] = SAT_CAST(dstbuf[0] * 255.0f); dst_ptr[1] = SAT_CAST(dstbuf[1] * 255.0f); dst_ptr[2] = SAT_CAST(dstbuf[2] * 255.0f); -#if dcn == 4 - dst_ptr[3] = MAX_NUM; +#else + *(__global uchar4 *)dst_ptr = (uchar4)(SAT_CAST(dstbuf[0] * 255.0f), + SAT_CAST(dstbuf[1] * 255.0f), SAT_CAST(dstbuf[2] * 255.0f), MAX_NUM); #endif + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1505,16 +1610,16 @@ __kernel void Lab2BGR(__global const uchar * srcptr, int src_step, int src_offse if (x < cols) { + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + #pragma unroll for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) { if (y < rows) { - int src_idx = mad24(y, src_step, src_offset + x * scnbytes); - int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes); - - __global const float * src = (__global const float *)(srcptr + src_idx); - __global float * dst = (__global float *)(dstptr + dst_idx); + __global const float * src = (__global const float *)(srcptr + src_index); + __global float * dst = (__global float *)(dstptr + dst_index); float4 src_pix = vload4(0, src); float srcbuf[3], dstbuf[3]; @@ -1530,8 +1635,10 @@ __kernel void Lab2BGR(__global const uchar * srcptr, int src_step, int src_offse #if dcn == 4 dst[3] = MAX_NUM; #endif + ++y; + dst_index += dst_step; + src_index += src_step; } - ++y; } } } @@ -1555,37 +1662,46 @@ __kernel void BGR2Luv(__global const uchar * srcptr, int src_step, int src_offse __global const float * LabCbrtTab, __constant float * coeffs, float _un, float _vn) { int x = get_global_id(0); - int y = get_global_id(1); + int y = get_global_id(1) * PIX_PER_WI_Y; - if (x < cols && y < rows) + if (x < cols) { - int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset)); - int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); - __global const float * src = (__global const float *)(srcptr + src_idx); - __global float * dst = (__global float *)(dstptr + dst_idx); + #pragma unroll + for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) + if (y < rows) + { + __global const float * src = (__global const float *)(srcptr + src_index); + __global float * dst = (__global float *)(dstptr + dst_index); - float R = src[0], G = src[1], B = src[2]; + float R = src[0], G = src[1], B = src[2]; #ifdef SRGB - R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); - G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); - B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); #endif - float X = R*coeffs[0] + G*coeffs[1] + B*coeffs[2]; - float Y = R*coeffs[3] + G*coeffs[4] + B*coeffs[5]; - float Z = R*coeffs[6] + G*coeffs[7] + B*coeffs[8]; + float X = fma(R, coeffs[0], fma(G, coeffs[1], B*coeffs[2])); + float Y = fma(R, coeffs[3], fma(G, coeffs[4], B*coeffs[5])); + float Z = fma(R, coeffs[6], fma(G, coeffs[7], B*coeffs[8])); - float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE); - L = 116.f*L - 16.f; + float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE); + L = fma(116.f, L, -16.f); - float d = (4*13) / max(X + 15 * Y + 3 * Z, FLT_EPSILON); - float u = L*(X*d - _un); - float v = L*((9*0.25f)*Y*d - _vn); + float d = 52.0f / fmax(fma(15.0f, Y, fma(3.0f, Z, X)), FLT_EPSILON); + float u = L*fma(X, d, -_un); + float v = L*fma(2.25f, Y*d, -_vn); - dst[0] = L; - dst[1] = u; - dst[2] = v; + dst[0] = L; + dst[1] = u; + dst[2] = v; + + ++y; + dst_index += dst_step; + src_index += src_step; + } } } @@ -1599,38 +1715,44 @@ __kernel void BGR2Luv(__global const uchar * src, int src_step, int src_offset, __global const float * LabCbrtTab, __constant float * coeffs, float _un, float _vn) { int x = get_global_id(0); - int y = get_global_id(1); + int y = get_global_id(1) * PIX_PER_WI_Y; - if (x < cols && y < rows) + if (x < cols) { - int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset)); - int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + src += mad24(y, src_step, mad24(x, scnbytes, src_offset)); + dst += mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); - src += src_idx; - dst += dst_idx; - - float scale = 1.0f / 255.0f; - float R = src[0]*scale, G = src[1]*scale, B = src[2]*scale; + #pragma unroll + for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) + if (y < rows) + { + float scale = 1.0f / 255.0f; + float R = src[0]*scale, G = src[1]*scale, B = src[2]*scale; #ifdef SRGB - R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); - G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); - B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); #endif - float X = R*coeffs[0] + G*coeffs[1] + B*coeffs[2]; - float Y = R*coeffs[3] + G*coeffs[4] + B*coeffs[5]; - float Z = R*coeffs[6] + G*coeffs[7] + B*coeffs[8]; + float X = fma(R, coeffs[0], fma(G, coeffs[1], B*coeffs[2])); + float Y = fma(R, coeffs[3], fma(G, coeffs[4], B*coeffs[5])); + float Z = fma(R, coeffs[6], fma(G, coeffs[7], B*coeffs[8])); - float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE); - L = 116.f*L - 16.f; + float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE); + L = 116.f*L - 16.f; - float d = (4*13) / max(X + 15 * Y + 3 * Z, FLT_EPSILON); - float u = L*(X*d - _un); - float v = L*((9*0.25f)*Y*d - _vn); + float d = (4*13) / fmax(fma(15.0f, Y, fma(3.0f, Z, X)), FLT_EPSILON); + float u = L*(X*d - _un); + float v = L*fma(2.25f, Y*d, -_vn); - dst[0] = SAT_CAST(L * 2.55f); - dst[1] = SAT_CAST(mad(u, 0.72033898305084743f, 96.525423728813564f)); - dst[2] = SAT_CAST(mad(v, 0.99609375f, 139.453125f)); + dst[0] = SAT_CAST(L * 2.55f); + dst[1] = SAT_CAST(fma(u, 0.72033898305084743f, 96.525423728813564f)); + dst[2] = SAT_CAST(fma(v, 0.99609375f, 139.453125f)); + + ++y; + dst += dst_step; + src += src_step; + } } } @@ -1646,42 +1768,50 @@ __kernel void Luv2BGR(__global const uchar * srcptr, int src_step, int src_offse __constant float * coeffs, float _un, float _vn) { int x = get_global_id(0); - int y = get_global_id(1); + int y = get_global_id(1) * PIX_PER_WI_Y; - if (x < cols && y < rows) + if (x < cols) { - int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset)); - int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset)); + int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); - __global const float * src = (__global const float *)(srcptr + src_idx); - __global float * dst = (__global float *)(dstptr + dst_idx); + #pragma unroll + for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) + if (y < rows) + { + __global const float * src = (__global const float *)(srcptr + src_index); + __global float * dst = (__global float *)(dstptr + dst_index); - float L = src[0], u = src[1], v = src[2], d, X, Y, Z; - Y = (L + 16.f) * (1.f/116.f); - Y = Y*Y*Y; - d = (1.f/13.f)/L; - u = u*d + _un; - v = v*d + _vn; - float iv = 1.f/v; - X = 2.25f * u * Y * iv ; - Z = (12 - 3 * u - 20 * v) * Y * 0.25f * iv; + float L = src[0], u = src[1], v = src[2], d, X, Y, Z; + Y = (L + 16.f) * (1.f/116.f); + Y = Y*Y*Y; + d = (1.f/13.f)/L; + u = fma(u, d, _un); + v = fma(v, d, _vn); + float iv = 1.f/v; + X = 2.25f * u * Y * iv; + Z = (12 - fma(3.0f, u, 20.0f * v)) * Y * 0.25f * iv; - float R = X*coeffs[0] + Y*coeffs[1] + Z*coeffs[2]; - float G = X*coeffs[3] + Y*coeffs[4] + Z*coeffs[5]; - float B = X*coeffs[6] + Y*coeffs[7] + Z*coeffs[8]; + float R = fma(X, coeffs[0], fma(Y, coeffs[1], Z * coeffs[2])); + float G = fma(X, coeffs[3], fma(Y, coeffs[4], Z * coeffs[5])); + float B = fma(X, coeffs[6], fma(Y, coeffs[7], Z * coeffs[8])); #ifdef SRGB - R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); - G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); - B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); #endif - dst[0] = R; - dst[1] = G; - dst[2] = B; + dst[0] = R; + dst[1] = G; + dst[2] = B; #if dcn == 4 - dst[3] = MAX_NUM; + dst[3] = MAX_NUM; #endif + ++y; + dst_index += dst_step; + src_index += src_step; + } } } @@ -1695,46 +1825,56 @@ __kernel void Luv2BGR(__global const uchar * src, int src_step, int src_offset, __constant float * coeffs, float _un, float _vn) { int x = get_global_id(0); - int y = get_global_id(1); + int y = get_global_id(1) * PIX_PER_WI_Y; - if (x < cols && y < rows) + if (x < cols) { - int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset)); - int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); + src += mad24(y, src_step, mad24(x, scnbytes, src_offset)); + dst += mad24(y, dst_step, mad24(x, dcnbytes, dst_offset)); - src += src_idx; - dst += dst_idx; + #pragma unroll + for (int cy = 0; cy < PIX_PER_WI_Y; ++cy) + if (y < rows) + { + float d, X, Y, Z; + float L = src[0]*(100.f/255.f); + float u = fma(convert_float(src[1]), 1.388235294117647f, -134.f); + float v = fma(convert_float(src[2]), 1.003921568627451f, - 140.f); + Y = (L + 16.f) * (1.f/116.f); + Y = Y*Y*Y; + d = (1.f/13.f)/L; + u = fma(u, d, _un); + v = fma(v, d, _vn); + float iv = 1.f/v; + X = 2.25f * u * Y * iv ; + Z = (12 - fma(3.0f, u, 20.0f * v)) * Y * 0.25f * iv; - float d, X, Y, Z; - float L = src[0]*(100.f/255.f); - float u = (float)(src[1]*1.388235294117647f - 134.f); - float v = (float)(src[2]*1.003921568627451f - 140.f); - Y = (L + 16.f) * (1.f/116.f); - Y = Y*Y*Y; - d = (1.f/13.f)/L; - u = u*d + _un; - v = v*d + _vn; - float iv = 1.f/v; - X = 2.25f * u * Y * iv ; - Z = (12 - 3 * u - 20 * v) * Y * 0.25f * iv; - - float R = X*coeffs[0] + Y*coeffs[1] + Z*coeffs[2]; - float G = X*coeffs[3] + Y*coeffs[4] + Z*coeffs[5]; - float B = X*coeffs[6] + Y*coeffs[7] + Z*coeffs[8]; + float R = fma(X, coeffs[0], fma(Y, coeffs[1], Z * coeffs[2])); + float G = fma(X, coeffs[3], fma(Y, coeffs[4], Z * coeffs[5])); + float B = fma(X, coeffs[6], fma(Y, coeffs[7], Z * coeffs[8])); #ifdef SRGB - R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); - G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); - B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); + B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE); #endif - dst[0] = SAT_CAST(R * 255.0f); - dst[1] = SAT_CAST(G * 255.0f); - dst[2] = SAT_CAST(B * 255.0f); + uchar dst0 = SAT_CAST(R * 255.0f); + uchar dst1 = SAT_CAST(G * 255.0f); + uchar dst2 = SAT_CAST(B * 255.0f); #if dcn == 4 - dst[3] = MAX_NUM; + *(__global uchar4 *)dst = (uchar4)(dst0, dst1, dst2, MAX_NUM); +#else + dst[0] = dst0; + dst[1] = dst1; + dst[2] = dst2; #endif + + ++y; + dst += dst_step; + src += src_step; + } } } diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp index e9fb0d3896..82bf2c06f1 100644 --- a/modules/imgproc/test/ocl/test_color.cpp +++ b/modules/imgproc/test/ocl/test_color.cpp @@ -305,11 +305,11 @@ OCL_TEST_P(CvtColor8u32f, Lab2LRGBA) { performTest(3, 4, CVTCODE(Lab2LRGB), dept OCL_TEST_P(CvtColor8u32f, BGR2Luv) { performTest(3, 3, CVTCODE(BGR2Luv), depth == CV_8U ? 1 : 1e-2); } OCL_TEST_P(CvtColor8u32f, RGB2Luv) { performTest(3, 3, CVTCODE(RGB2Luv), depth == CV_8U ? 1 : 1e-2); } OCL_TEST_P(CvtColor8u32f, LBGR2Luv) { performTest(3, 3, CVTCODE(LBGR2Luv), depth == CV_8U ? 1 : 4e-3); } -OCL_TEST_P(CvtColor8u32f, LRGB2Luv) { performTest(3, 3, CVTCODE(LRGB2Luv), depth == CV_8U ? 1 : 4e-3); } +OCL_TEST_P(CvtColor8u32f, LRGB2Luv) { performTest(3, 3, CVTCODE(LRGB2Luv), depth == CV_8U ? 1 : 5e-3); } OCL_TEST_P(CvtColor8u32f, BGRA2Luv) { performTest(4, 3, CVTCODE(BGR2Luv), depth == CV_8U ? 1 : 8e-3); } OCL_TEST_P(CvtColor8u32f, RGBA2Luv) { performTest(4, 3, CVTCODE(RGB2Luv), depth == CV_8U ? 1 : 9e-3); } -OCL_TEST_P(CvtColor8u32f, LBGRA2Luv) { performTest(4, 3, CVTCODE(LBGR2Luv), depth == CV_8U ? 1 : 4e-3); } -OCL_TEST_P(CvtColor8u32f, LRGBA2Luv) { performTest(4, 3, CVTCODE(LRGB2Luv), depth == CV_8U ? 1 : 4e-3); } +OCL_TEST_P(CvtColor8u32f, LBGRA2Luv) { performTest(4, 3, CVTCODE(LBGR2Luv), depth == CV_8U ? 1 : 5e-3); } +OCL_TEST_P(CvtColor8u32f, LRGBA2Luv) { performTest(4, 3, CVTCODE(LRGB2Luv), depth == CV_8U ? 1 : 5e-3); } OCL_TEST_P(CvtColor8u32f, Luv2BGR) { performTest(3, 3, CVTCODE(Luv2BGR), depth == CV_8U ? 1 : 7e-5); } OCL_TEST_P(CvtColor8u32f, Luv2RGB) { performTest(3, 3, CVTCODE(Luv2RGB), depth == CV_8U ? 1 : 7e-5); } From e89cee35e5821377a75b3de6d92d1c33391bc31e Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sat, 21 Jun 2014 19:17:03 +0400 Subject: [PATCH 04/24] optimized cv::inRange --- modules/core/src/arithm.cpp | 17 ++++++--- modules/core/src/opencl/inrange.cl | 57 +++++++++++++++++++++--------- 2 files changed, 53 insertions(+), 21 deletions(-) diff --git a/modules/core/src/arithm.cpp b/modules/core/src/arithm.cpp index 7ca8b4b48d..197ebc1de4 100644 --- a/modules/core/src/arithm.cpp +++ b/modules/core/src/arithm.cpp @@ -3130,9 +3130,16 @@ static bool ocl_inRange( InputArray _src, InputArray _lowerb, (!haveScalar && (sdepth != ldepth || sdepth != udepth)) ) return false; - ocl::Kernel ker("inrange", ocl::core::inrange_oclsrc, - format("%s-D cn=%d -D T=%s%s", haveScalar ? "-D HAVE_SCALAR " : "", - cn, ocl::typeToStr(sdepth), doubleSupport ? " -D DOUBLE_SUPPORT" : "")); + int kercn = haveScalar ? cn : std::max(std::min(ocl::predictOptimalVectorWidth(_src, _lowerb, _upperb, _dst), 4), cn); + if (kercn % cn != 0) + kercn = cn; + int colsPerWI = kercn / cn; + String opts = format("%s-D cn=%d -D srcT=%s -D srcT1=%s -D dstT=%s -D kercn=%d -D depth=%d%s -D colsPerWI=%d", + haveScalar ? "-D HAVE_SCALAR " : "", cn, ocl::typeToStr(CV_MAKE_TYPE(sdepth, kercn)), + ocl::typeToStr(sdepth), ocl::typeToStr(CV_8UC(colsPerWI)), kercn, sdepth, + doubleSupport ? " -D DOUBLE_SUPPORT" : "", colsPerWI); + + ocl::Kernel ker("inrange", ocl::core::inrange_oclsrc, opts); if (ker.empty()) return false; @@ -3180,7 +3187,7 @@ static bool ocl_inRange( InputArray _src, InputArray _lowerb, } ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src), - dstarg = ocl::KernelArg::WriteOnly(dst); + dstarg = ocl::KernelArg::WriteOnly(dst, 1, colsPerWI); if (haveScalar) { @@ -3194,7 +3201,7 @@ static bool ocl_inRange( InputArray _src, InputArray _lowerb, ker.args(srcarg, dstarg, ocl::KernelArg::ReadOnlyNoSize(lscalaru), ocl::KernelArg::ReadOnlyNoSize(uscalaru), rowsPerWI); - size_t globalsize[2] = { ssize.width, (ssize.height + rowsPerWI - 1) / rowsPerWI }; + size_t globalsize[2] = { ssize.width / colsPerWI, (ssize.height + rowsPerWI - 1) / rowsPerWI }; return ker.run(2, globalsize, NULL, false); } diff --git a/modules/core/src/opencl/inrange.cl b/modules/core/src/opencl/inrange.cl index 0de561f5f2..538259539a 100644 --- a/modules/core/src/opencl/inrange.cl +++ b/modules/core/src/opencl/inrange.cl @@ -52,7 +52,7 @@ __kernel void inrange(__global const uchar * src1ptr, int src1_step, int src1_offset, __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, #ifdef HAVE_SCALAR - __global const T * src2, __global const T * src3, + __global const srcT1 * src2, __global const srcT1 * src3, #else __global const uchar * src2ptr, int src2_step, int src2_offset, __global const uchar * src3ptr, int src3_step, int src3_offset, @@ -64,31 +64,56 @@ __kernel void inrange(__global const uchar * src1ptr, int src1_step, int src1_of if (x < dst_cols) { - int src1_index = mad24(y0, src1_step, mad24(x, (int)sizeof(T) * cn, src1_offset)); - int dst_index = mad24(y0, dst_step, x + dst_offset); + int src1_index = mad24(y0, src1_step, mad24(x, (int)sizeof(srcT1) * kercn, src1_offset)); + int dst_index = mad24(y0, dst_step, mad24(x, colsPerWI, dst_offset)); #ifndef HAVE_SCALAR - int src2_index = mad24(y0, src2_step, mad24(x, (int)sizeof(T) * cn, src2_offset)); - int src3_index = mad24(y0, src3_step, mad24(x, (int)sizeof(T) * cn, src3_offset)); + int src2_index = mad24(y0, src2_step, mad24(x, (int)sizeof(srcT1) * kercn, src2_offset)); + int src3_index = mad24(y0, src3_step, mad24(x, (int)sizeof(srcT1) * kercn, src3_offset)); #endif for (int y = y0, y1 = min(dst_rows, y0 + rowsPerWI); y < y1; ++y, src1_index += src1_step, dst_index += dst_step) { - __global const T * src1 = (__global const T *)(src1ptr + src1_index); +#if kercn >= cn && kercn == 4 && depth <= 4 && !defined HAVE_SCALAR + srcT src1 = *(__global const srcT *)(src1ptr + src1_index); + srcT src2 = *(__global const srcT *)(src2ptr + src2_index); + srcT src3 = *(__global const srcT *)(src3ptr + src3_index); + __global dstT * dst = (__global dstT *)(dstptr + dst_index); +#if cn == 1 + dst[0] = src2 > src1 || src3 < src1 ? (dstT)(0) : (dstT)(255); +#elif cn == 2 + dst[0] = (dstT)(src2.xy > src1.xy || src3.xy < src1.xy || + src2.zw > src1.zw || src3.zw < src1.zw ? (dstT)(0) : (dstT)(255); +#elif cn == 4 + dst[0] = (dstT)(src2.x > src1.x || src3.x < src1.x || + src2.y > src1.y || src3.y < src1.y || + src2.z > src1.z || src3.z < src1.z || + src2.w > src1.w || src3.w < src1.w ? 0 : 255); +#endif +#else + __global const srcT1 * src1 = (__global const srcT1 *)(src1ptr + src1_index); __global uchar * dst = dstptr + dst_index; #ifndef HAVE_SCALAR - __global const T * src2 = (__global const T *)(src2ptr + src2_index); - __global const T * src3 = (__global const T *)(src3ptr + src3_index); + __global const srcT1 * src2 = (__global const srcT1 *)(src2ptr + src2_index); + __global const srcT1 * src3 = (__global const srcT1 *)(src3ptr + src3_index); #endif - dst[0] = 255; - - for (int c = 0; c < cn; ++c) - if (src2[c] > src1[c] || src3[c] < src1[c]) - { - dst[0] = 0; - break; - } + #pragma unroll + for (int px = 0; px < colsPerWI; ++px, src1 += cn +#ifndef HAVE_SCALAR + , src2 += cn, src3 += cn +#endif + ) + { + dst[px] = 255; + for (int c = 0; c < cn; ++c) + if (src2[c] > src1[c] || src3[c] < src1[c]) + { + dst[px] = 0; + break; + } + } +#endif // kercn >= cn #ifndef HAVE_SCALAR src2_index += src2_step; src3_index += src3_step; From 09bcc061dd32361b4864c3b1e62ef5d7121571e5 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Thu, 19 Jun 2014 14:39:49 +0400 Subject: [PATCH 05/24] Change kernel for optimization. Remove restriction to align data Fix kernel compilation errors on AMD system Fix licanse information in cl file Support CV_64F destination type Change build options of the kernel Optimize sum of square Remove separate kernel for integral square Increase epsilon for perfomance tests Increase epsilon for perfomance tests Test double support on AMD devices Fix some issues Try to fix problems with AMD device Try to solve problem with AMD device Fix error of destination size in kernel Fix warnings --- modules/imgproc/perf/opencl/perf_imgproc.cpp | 6 +- modules/imgproc/src/opencl/integral_sqrsum.cl | 512 ------------------ modules/imgproc/src/opencl/integral_sum.cl | 384 +++++-------- modules/imgproc/src/sumpixels.cpp | 122 ++--- 4 files changed, 199 insertions(+), 825 deletions(-) delete mode 100644 modules/imgproc/src/opencl/integral_sqrsum.cl diff --git a/modules/imgproc/perf/opencl/perf_imgproc.cpp b/modules/imgproc/perf/opencl/perf_imgproc.cpp index 1b3ba7f19e..7f0770853d 100644 --- a/modules/imgproc/perf/opencl/perf_imgproc.cpp +++ b/modules/imgproc/perf/opencl/perf_imgproc.cpp @@ -231,7 +231,7 @@ OCL_PERF_TEST_P(IntegralFixture, Integral1, ::testing::Combine(OCL_TEST_SIZES, O OCL_TEST_CYCLE() cv::integral(src, dst, ddepth); - SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE); + SANITY_CHECK(dst, 2e-6, ERROR_RELATIVE); } OCL_PERF_TEST_P(IntegralFixture, Integral2, ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32S, CV_32F))) @@ -243,11 +243,11 @@ OCL_PERF_TEST_P(IntegralFixture, Integral2, ::testing::Combine(OCL_TEST_SIZES, O checkDeviceMaxMemoryAllocSize(srcSize, ddepth); UMat src(srcSize, CV_8UC1), sum(srcSize + Size(1, 1), ddepth), sqsum(srcSize + Size(1, 1), CV_32F); - declare.in(src, WARMUP_RNG).out(sum).out(sqsum); + declare.in(src, WARMUP_RNG).out(sum, sqsum); OCL_TEST_CYCLE() cv::integral(src, sum, sqsum, ddepth, CV_32F); - SANITY_CHECK(sum, 1e-6, ERROR_RELATIVE); + SANITY_CHECK(sum, 2e-4, ERROR_RELATIVE); SANITY_CHECK(sqsum, 5e-5, ERROR_RELATIVE); } diff --git a/modules/imgproc/src/opencl/integral_sqrsum.cl b/modules/imgproc/src/opencl/integral_sqrsum.cl deleted file mode 100644 index 8b5d2452b8..0000000000 --- a/modules/imgproc/src/opencl/integral_sqrsum.cl +++ /dev/null @@ -1,512 +0,0 @@ -/*M/////////////////////////////////////////////////////////////////////////////////////// -// -// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. -// -// By downloading, copying, installing or using the software you agree to this license. -// If you do not agree to this license, do not download, install, -// copy or use the software. -// -// -// License Agreement -// For Open Source Computer Vision Library -// -// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. -// Third party copyrights are property of their respective owners. -// -// @Authors -// Shengen Yan,yanshengen@gmail.com -// -// Redistribution and use in source and binary forms, with or without modification, -// are permitted provided that the following conditions are met: -// -// * Redistribution's of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// -// * Redistribution's in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// * The name of the copyright holders may not be used to endorse or promote products -// derived from this software without specific prior written permission. -// -// This software is provided by the copyright holders and contributors as is and -// any express or implied warranties, including, but not limited to, the implied -// warranties of merchantability and fitness for a particular purpose are disclaimed. -// In no event shall the Intel Corporation or contributors be liable for any direct, -// indirect, incidental, special, exemplary, or consequential damages -// (including, but not limited to, procurement of substitute goods or services; -// loss of use, data, or profits; or business interruption) however caused -// and on any theory of liability, whether in contract, strict liability, -// or tort (including negligence or otherwise) arising in any way out of -// the use of this software, even if advised of the possibility of such damage. -// -//M*/ - -#ifdef DOUBLE_SUPPORT -#ifdef cl_amd_fp64 -#pragma OPENCL EXTENSION cl_amd_fp64:enable -#elif defined (cl_khr_fp64) -#pragma OPENCL EXTENSION cl_khr_fp64:enable -#endif -#endif - -#if sqdepth == 6 -#define CONVERT(step) ((step)>>1) -#else -#define CONVERT(step) ((step)) -#endif - -#define LSIZE 256 -#define LSIZE_1 255 -#define LSIZE_2 254 -#define HF_LSIZE 128 -#define LOG_LSIZE 8 -#define LOG_NUM_BANKS 5 -#define NUM_BANKS 32 -#define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS) - -#define noconvert - -#if sdepth == 4 - -kernel void integral_cols(__global uchar4 *src, __global int *sum, __global TYPE *sqsum, - int src_offset, int pre_invalid, int rows, int cols, int src_step, int dst_step, int dst1_step) -{ - int lid = get_local_id(0); - int gid = get_group_id(0); - int4 src_t[2], sum_t[2]; - TYPE4 sqsum_t[2]; - __local int4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local TYPE4 lm_sqsum[2][LSIZE + LOG_LSIZE]; - __local int* sum_p; - __local TYPE* sqsum_p; - src_step = src_step >> 2; - gid = gid << 1; - for(int i = 0; i < rows; i =i + LSIZE_1) - { - src_t[0] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + min(gid, cols - 1)]) : 0); - src_t[1] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + min(gid + 1, cols - 1)]) : 0); - - sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sqsum_t[0] = (i == 0 ? (TYPE4)0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); - sqsum_t[1] = (i == 0 ? (TYPE4)0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]); - barrier(CLK_LOCAL_MEM_FENCE); - - int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - lm_sum[0][bf_loc] = src_t[0]; - lm_sqsum[0][bf_loc] = convert_TYPE4(src_t[0] * src_t[0]); - - lm_sum[1][bf_loc] = src_t[1]; - lm_sqsum[1][bf_loc] = convert_TYPE4(src_t[1] * src_t[1]); - - int offset = 1; - for(int d = LSIZE >> 1 ; d > 0; d>>=1) - { - barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai]; - } - offset <<= 1; - } - barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 2) - { - lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; - lm_sqsum[lid][LSIZE_2 + LOG_LSIZE] = 0; - } - for(int d = 1; d < LSIZE; d <<= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - offset >>= 1; - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; - - lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai]; - lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step /4, loc_s1 = loc_s0 + dst_step ; - int loc_sq0 = gid * CONVERT(dst1_step) + i + lid - 1 - pre_invalid * dst1_step / sizeof(TYPE),loc_sq1 = loc_sq0 + CONVERT(dst1_step); - if(lid > 0 && (i+lid) <= rows) - { - lm_sum[0][bf_loc] += sum_t[0]; - lm_sum[1][bf_loc] += sum_t[1]; - lm_sqsum[0][bf_loc] += sqsum_t[0]; - lm_sqsum[1][bf_loc] += sqsum_t[1]; - sum_p = (__local int*)(&(lm_sum[0][bf_loc])); - sqsum_p = (__local TYPE*)(&(lm_sqsum[0][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 4 + k >= cols + pre_invalid || gid * 4 + k < pre_invalid) continue; - sum[loc_s0 + k * dst_step / 4] = sum_p[k]; - sqsum[loc_sq0 + k * dst1_step / sizeof(TYPE)] = sqsum_p[k]; - } - sum_p = (__local int*)(&(lm_sum[1][bf_loc])); - sqsum_p = (__local TYPE*)(&(lm_sqsum[1][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 4 + k + 4 >= cols + pre_invalid) break; - sum[loc_s1 + k * dst_step / 4] = sum_p[k]; - sqsum[loc_sq1 + k * dst1_step / sizeof(TYPE)] = sqsum_p[k]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - } -} - -kernel void integral_rows(__global int4 *srcsum, __global TYPE4 * srcsqsum,__global int *sum, - __global TYPE *sqsum, int rows, int cols, int src_step, int src1_step, int sum_step, - int sqsum_step, int sum_offset, int sqsum_offset) -{ - int lid = get_local_id(0); - int gid = get_group_id(0); - int4 src_t[2], sum_t[2]; - TYPE4 sqsrc_t[2],sqsum_t[2]; - __local int4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local TYPE4 lm_sqsum[2][LSIZE + LOG_LSIZE]; - __local int *sum_p; - __local TYPE *sqsum_p; - src_step = src_step >> 4; - src1_step = (src1_step / sizeof(TYPE)) >> 2 ; - gid <<= 1; - for(int i = 0; i < rows; i =i + LSIZE_1) - { - src_t[0] = i + lid < rows ? srcsum[(lid+i) * src_step + gid ] : (int4)0; - sqsrc_t[0] = i + lid < rows ? srcsqsum[(lid+i) * src1_step + gid ] : (TYPE4)0; - src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid + 1] : (int4)0; - sqsrc_t[1] = i + lid < rows ? srcsqsum[(lid+i) * src1_step + gid + 1] : (TYPE4)0; - - sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sqsum_t[0] = (i == 0 ? (TYPE4)0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); - sqsum_t[1] = (i == 0 ? (TYPE4)0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]); - barrier(CLK_LOCAL_MEM_FENCE); - - int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - lm_sum[0][bf_loc] = src_t[0]; - lm_sqsum[0][bf_loc] = sqsrc_t[0]; - - lm_sum[1][bf_loc] = src_t[1]; - lm_sqsum[1][bf_loc] = sqsrc_t[1]; - - int offset = 1; - for(int d = LSIZE >> 1 ; d > 0; d>>=1) - { - barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai]; - } - offset <<= 1; - } - barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 2) - { - lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; - lm_sqsum[lid][LSIZE_2 + LOG_LSIZE] = 0; - } - for(int d = 1; d < LSIZE; d <<= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - offset >>= 1; - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; - - lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai]; - lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - if(gid == 0 && (i + lid) <= rows) - { - sum[sum_offset + i + lid] = 0; - sqsum[sqsum_offset + i + lid] = 0; - } - if(i + lid == 0) - { - int loc0 = gid * sum_step; - int loc1 = gid * CONVERT(sqsum_step); - for(int k = 1; k <= 8; k++) - { - if(gid * 4 + k > cols) break; - sum[sum_offset + loc0 + k * sum_step / 4] = 0; - sqsum[sqsum_offset + loc1 + k * sqsum_step / sizeof(TYPE)] = 0; - } - } - int loc_s0 = sum_offset + gid * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ; - int loc_sq0 = sqsum_offset + gid * CONVERT(sqsum_step) + sqsum_step / sizeof(TYPE) + i + lid, loc_sq1 = loc_sq0 + CONVERT(sqsum_step) ; - - if(lid > 0 && (i+lid) <= rows) - { - lm_sum[0][bf_loc] += sum_t[0]; - lm_sum[1][bf_loc] += sum_t[1]; - lm_sqsum[0][bf_loc] += sqsum_t[0]; - lm_sqsum[1][bf_loc] += sqsum_t[1]; - sum_p = (__local int*)(&(lm_sum[0][bf_loc])); - sqsum_p = (__local TYPE*)(&(lm_sqsum[0][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 4 + k >= cols) break; - sum[loc_s0 + k * sum_step / 4] = sum_p[k]; - sqsum[loc_sq0 + k * sqsum_step / sizeof(TYPE)] = sqsum_p[k]; - } - sum_p = (__local int*)(&(lm_sum[1][bf_loc])); - sqsum_p = (__local TYPE*)(&(lm_sqsum[1][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 4 + 4 + k >= cols) break; - sum[loc_s1 + k * sum_step / 4] = sum_p[k]; - sqsum[loc_sq1 + k * sqsum_step / sizeof(TYPE)] = sqsum_p[k]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - } -} - -#elif sdepth == 5 - -kernel void integral_cols(__global uchar4 *src, __global float *sum, __global TYPE *sqsum, - int src_offset, int pre_invalid, int rows, int cols, int src_step, int dst_step, int dst1_step) -{ - int lid = get_local_id(0); - int gid = get_group_id(0); - float4 src_t[2], sum_t[2]; - TYPE4 sqsum_t[2]; - __local float4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local TYPE4 lm_sqsum[2][LSIZE + LOG_LSIZE]; - __local float* sum_p; - __local TYPE* sqsum_p; - src_step = src_step >> 2; - gid = gid << 1; - for(int i = 0; i < rows; i =i + LSIZE_1) - { - src_t[0] = (i + lid < rows ? convert_float4(src[src_offset + (lid+i) * src_step + min(gid, cols - 1)]) : (float4)0); - src_t[1] = (i + lid < rows ? convert_float4(src[src_offset + (lid+i) * src_step + min(gid + 1, cols - 1)]) : (float4)0); - - sum_t[0] = (i == 0 ? (float4)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sqsum_t[0] = (i == 0 ? (TYPE4)0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? (float4)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); - sqsum_t[1] = (i == 0 ? (TYPE4)0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]); - barrier(CLK_LOCAL_MEM_FENCE); - - int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - lm_sum[0][bf_loc] = src_t[0]; - lm_sqsum[0][bf_loc] = convert_TYPE4(src_t[0] * src_t[0]); -// printf("%f\n", src_t[0].s0); - - lm_sum[1][bf_loc] = src_t[1]; - lm_sqsum[1][bf_loc] = convert_TYPE4(src_t[1] * src_t[1]); - - int offset = 1; - for(int d = LSIZE >> 1 ; d > 0; d>>=1) - { - barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai]; - } - offset <<= 1; - } - barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 2) - { - lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; - lm_sqsum[lid][LSIZE_2 + LOG_LSIZE] = 0; - } - for(int d = 1; d < LSIZE; d <<= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - offset >>= 1; - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; - - lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai]; - lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ; - int loc_sq0 = gid * CONVERT(dst1_step) + i + lid - 1 - pre_invalid * dst1_step / sizeof(TYPE), loc_sq1 = loc_sq0 + CONVERT(dst1_step); - if(lid > 0 && (i+lid) <= rows) - { - lm_sum[0][bf_loc] += sum_t[0]; - lm_sum[1][bf_loc] += sum_t[1]; - lm_sqsum[0][bf_loc] += sqsum_t[0]; - lm_sqsum[1][bf_loc] += sqsum_t[1]; - sum_p = (__local float*)(&(lm_sum[0][bf_loc])); - sqsum_p = (__local TYPE*)(&(lm_sqsum[0][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 4 + k >= cols + pre_invalid || gid * 4 + k < pre_invalid) continue; - sum[loc_s0 + k * dst_step / 4] = sum_p[k]; - sqsum[loc_sq0 + k * dst1_step / sizeof(TYPE)] = sqsum_p[k]; - } - sum_p = (__local float*)(&(lm_sum[1][bf_loc])); - sqsum_p = (__local TYPE*)(&(lm_sqsum[1][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 4 + k + 4 >= cols + pre_invalid) break; - sum[loc_s1 + k * dst_step / 4] = sum_p[k]; - sqsum[loc_sq1 + k * dst1_step / sizeof(TYPE)] = sqsum_p[k]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - } -} - -kernel void integral_rows(__global float4 *srcsum, __global TYPE4 * srcsqsum, __global float *sum , - __global TYPE *sqsum, int rows, int cols, int src_step, int src1_step, int sum_step, - int sqsum_step, int sum_offset, int sqsum_offset) -{ - int lid = get_local_id(0); - int gid = get_group_id(0); - float4 src_t[2], sum_t[2]; - TYPE4 sqsrc_t[2],sqsum_t[2]; - __local float4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local TYPE4 lm_sqsum[2][LSIZE + LOG_LSIZE]; - __local float *sum_p; - __local TYPE *sqsum_p; - src_step = src_step >> 4; - src1_step = (src1_step / sizeof(TYPE)) >> 2; - for(int i = 0; i < rows; i =i + LSIZE_1) - { - src_t[0] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2] : (float4)0; - sqsrc_t[0] = i + lid < rows ? srcsqsum[(lid+i) * src1_step + gid * 2] : (TYPE4)0; - src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2 + 1] : (float4)0; - sqsrc_t[1] = i + lid < rows ? srcsqsum[(lid+i) * src1_step + gid * 2 + 1] : (TYPE4)0; - - sum_t[0] = (i == 0 ? (float4)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sqsum_t[0] = (i == 0 ? (TYPE4)0 : lm_sqsum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? (float4)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); - sqsum_t[1] = (i == 0 ? (TYPE4)0 : lm_sqsum[1][LSIZE_2 + LOG_LSIZE]); - barrier(CLK_LOCAL_MEM_FENCE); - - int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - lm_sum[0][bf_loc] = src_t[0]; - lm_sqsum[0][bf_loc] = sqsrc_t[0]; - - lm_sum[1][bf_loc] = src_t[1]; - lm_sqsum[1][bf_loc] = sqsrc_t[1]; - - int offset = 1; - for(int d = LSIZE >> 1 ; d > 0; d>>=1) - { - barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai]; - } - offset <<= 1; - } - barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 2) - { - lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; - lm_sqsum[lid][LSIZE_2 + LOG_LSIZE] = 0; - } - for(int d = 1; d < LSIZE; d <<= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - offset >>= 1; - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; - - lm_sqsum[lid >> 7][bi] += lm_sqsum[lid >> 7][ai]; - lm_sqsum[lid >> 7][ai] = lm_sqsum[lid >> 7][bi] - lm_sqsum[lid >> 7][ai]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - if(gid == 0 && (i + lid) <= rows) - { - sum[sum_offset + i + lid] = 0; - sqsum[sqsum_offset + i + lid] = 0; - } - if(i + lid == 0) - { - int loc0 = gid * 2 * sum_step; - int loc1 = gid * 2 * CONVERT(sqsum_step); - for(int k = 1; k <= 8; k++) - { - if(gid * 8 + k > cols) break; - sum[sum_offset + loc0 + k * sum_step / 4] = 0; - sqsum[sqsum_offset + loc1 + k * sqsum_step / sizeof(TYPE)] = 0; - } - } - int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ; - int loc_sq0 = sqsum_offset + gid * 2 * CONVERT(sqsum_step) + sqsum_step / sizeof(TYPE) + i + lid, loc_sq1 = loc_sq0 + CONVERT(sqsum_step) ; - if(lid > 0 && (i+lid) <= rows) - { - lm_sum[0][bf_loc] += sum_t[0]; - lm_sum[1][bf_loc] += sum_t[1]; - lm_sqsum[0][bf_loc] += sqsum_t[0]; - lm_sqsum[1][bf_loc] += sqsum_t[1]; - sum_p = (__local float*)(&(lm_sum[0][bf_loc])); - sqsum_p = (__local TYPE*)(&(lm_sqsum[0][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 8 + k >= cols) break; - sum[loc_s0 + k * sum_step / 4] = sum_p[k]; - sqsum[loc_sq0 + k * sqsum_step / sizeof(TYPE)] = sqsum_p[k]; - } - sum_p = (__local float*)(&(lm_sum[1][bf_loc])); - sqsum_p = (__local TYPE*)(&(lm_sqsum[1][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 8 + 4 + k >= cols) break; - sum[loc_s1 + k * sum_step / 4] = sum_p[k]; - sqsum[loc_sq1 + k * sqsum_step / sizeof(TYPE)] = sqsum_p[k]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - } -} - -#endif diff --git a/modules/imgproc/src/opencl/integral_sum.cl b/modules/imgproc/src/opencl/integral_sum.cl index 333c7121cb..49a3bde955 100644 --- a/modules/imgproc/src/opencl/integral_sum.cl +++ b/modules/imgproc/src/opencl/integral_sum.cl @@ -1,46 +1,9 @@ /*M/////////////////////////////////////////////////////////////////////////////////////// -// -// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. -// -// By downloading, copying, installing or using the software you agree to this license. -// If you do not agree to this license, do not download, install, -// copy or use the software. -// -// -// License Agreement -// For Open Source Computer Vision Library -// -// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. +// Copyright (C) 2014, Itseez, Inc., all rights reserved. // Third party copyrights are property of their respective owners. -// -// @Authors -// Shengen Yan,yanshengen@gmail.com -// -// Redistribution and use in source and binary forms, with or without modification, -// are permitted provided that the following conditions are met: -// -// * Redistribution's of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// -// * Redistribution's in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// * The name of the copyright holders may not be used to endorse or promote products -// derived from this software without specific prior written permission. -// -// This software is provided by the copyright holders and contributors as is and -// any express or implied warranties, including, but not limited to, the implied -// warranties of merchantability and fitness for a particular purpose are disclaimed. -// In no event shall the Intel Corporation or contributors be liable for any direct, -// indirect, incidental, special, exemplary, or consequential damages -// (including, but not limited to, procurement of substitute goods or services; -// loss of use, data, or profits; or business interruption) however caused -// and on any theory of liability, whether in contract, strict liability, -// or tort (including negligence or otherwise) arising in any way out of -// the use of this software, even if advised of the possibility of such damage. -// //M*/ #ifdef DOUBLE_SUPPORT @@ -51,237 +14,170 @@ #endif #endif -#define LSIZE 256 -#define LSIZE_1 255 -#define LSIZE_2 254 -#define HF_LSIZE 128 -#define LOG_LSIZE 8 -#define LOG_NUM_BANKS 5 -#define NUM_BANKS 32 -#define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS) - -#if sdepth == 4 -#define sumT int -#define vecSumT int4 -#define convertToSum4 convert_int4 -#elif sdepth == 5 -#define sumT float -#define vecSumT float4 -#define convertToSum4 convert_float4 +#ifndef LOCAL_SUM_SIZE +#define LOCAL_SUM_SIZE 16 #endif +#define LOCAL_SUM_STRIDE (LOCAL_SUM_SIZE + 1) -kernel void integral_sum_cols(__global const uchar4 *src, __global uchar *sum_ptr, - int src_offset, int rows, int cols, int src_step, int dst_step) + +kernel void integral_sum_cols(__global const uchar *src_ptr, int src_step, int src_offset, int rows, int cols, + __global uchar *buf_ptr, int buf_step, int buf_offset +#ifdef SUM_SQUARE + ,__global uchar *buf_sq_ptr, int buf_sq_step, int buf_sq_offset +#endif + ) { - __global sumT *sum = (__global sumT *)sum_ptr; + __local sumT lm_sum[LOCAL_SUM_STRIDE * LOCAL_SUM_SIZE]; +#ifdef SUM_SQUARE + __local sumSQT lm_sum_sq[LOCAL_SUM_STRIDE * LOCAL_SUM_SIZE]; +#endif int lid = get_local_id(0); int gid = get_group_id(0); - vecSumT src_t[2], sum_t[2]; - __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE]; - __local sumT* sum_p; - src_step = src_step >> 2; - gid = gid << 1; - int lid_prim = ((lid & 127) << 1) + 1; - for (int i = 0; i < rows; i += LSIZE_1) + + int x = get_global_id(0); + int src_index = x + src_offset; + + sumT accum = 0; +#ifdef SUM_SQUARE + sumSQT accum_sq = 0; +#endif + for (int y = 0; y < rows; y += LOCAL_SUM_SIZE) { - if (i + lid < rows) + int lsum_index = lid; + #pragma unroll + for (int yin = 0; yin < LOCAL_SUM_SIZE; yin++, src_index+=src_step, lsum_index += LOCAL_SUM_STRIDE) { - int src_index = mad24((lid+i), src_step, gid + src_offset); - src_t[0] = convertToSum4(src[src_index]); - src_t[1] = convertToSum4(src[src_index + 1]); - } - else - { - src_t[0] = (vecSumT)0; - src_t[1] = (vecSumT)0; - } - - if (i == 0) - { - sum_t[0] = (vecSumT)0; - sum_t[1] = (vecSumT)0; - } - else - { - sum_t[0] = lm_sum[0][LSIZE_2 + LOG_LSIZE]; - sum_t[1] = lm_sum[1][LSIZE_2 + LOG_LSIZE]; + if ((x < cols) && (y + yin < rows)) + { + __global const uchar *src = src_ptr + src_index; + accum += src[0]; +#ifdef SUM_SQUARE + sumSQT temp = src[0] * src[0]; + accum_sq += temp; +#endif + } + lm_sum[lsum_index] = accum; +#ifdef SUM_SQUARE + lm_sum_sq[lsum_index] = accum_sq; +#endif } barrier(CLK_LOCAL_MEM_FENCE); - int bf_loc = lid + GET_CONFLICT_OFFSET(lid); + //int buf_index = buf_offset + buf_step * LOCAL_SUM_COLS * gid + sizeof(sumT) * y + sizeof(sumT) * lid; + int buf_index = mad24(buf_step, LOCAL_SUM_SIZE * gid, mad24((int)sizeof(sumT), y + lid, buf_offset)); +#ifdef SUM_SQUARE + int buf_sq_index = mad24(buf_sq_step, LOCAL_SUM_SIZE * gid, mad24((int)sizeof(sumSQT), y + lid, buf_sq_offset)); +#endif - lm_sum[0][bf_loc] = src_t[0]; - lm_sum[1][bf_loc] = src_t[1]; - - int offset = 1; - for (int d = LSIZE >> 1 ; d > 0; d>>=1) + lsum_index = LOCAL_SUM_STRIDE * lid; + #pragma unroll + for (int yin = 0; yin < LOCAL_SUM_SIZE; yin++, lsum_index ++) { - barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * lid_prim - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - } - offset <<= 1; - } - barrier(CLK_LOCAL_MEM_FENCE); - if (lid < 2) - { - lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; - } - for (int d = 1; d < LSIZE; d <<= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - offset >>= 1; - int ai = offset * lid_prim - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - if (lid > 0 && (i+lid) <= rows) - { - int loc_s0 = mad24(gid, dst_step, i + lid - 1), loc_s1 = loc_s0 + dst_step; - lm_sum[0][bf_loc] += sum_t[0]; - lm_sum[1][bf_loc] += sum_t[1]; - sum_p = (__local sumT*)(&(lm_sum[0][bf_loc])); - for (int k = 0; k < 4; k++) - { - if (gid * 4 + k >= cols) - break; - sum[loc_s0 + k * dst_step / 4] = sum_p[k]; - } - sum_p = (__local sumT*)(&(lm_sum[1][bf_loc])); - for (int k = 0; k < 4; k++) - { - if (gid * 4 + k + 4 >= cols) - break; - sum[loc_s1 + k * dst_step / 4] = sum_p[k]; - } + __global sumT *buf = (__global sumT *)(buf_ptr + buf_index); + buf[0] = lm_sum[lsum_index]; + buf_index += buf_step; +#ifdef SUM_SQUARE + __global sumSQT *bufsq = (__global sumSQT *)(buf_sq_ptr + buf_sq_index); + bufsq[0] = lm_sum_sq[lsum_index]; + buf_sq_index += buf_sq_step; +#endif } barrier(CLK_LOCAL_MEM_FENCE); } } - -kernel void integral_sum_rows(__global const uchar *srcsum_ptr, __global uchar *sum_ptr, - int rows, int cols, int src_step, int sum_step, int sum_offset) +kernel void integral_sum_rows(__global const uchar *buf_ptr, int buf_step, int buf_offset, +#ifdef SUM_SQUARE + __global uchar *buf_sq_ptr, int buf_sq_step, int buf_sq_offset, +#endif + __global uchar *dst_ptr, int dst_step, int dst_offset, int rows, int cols +#ifdef SUM_SQUARE + ,__global uchar *dst_sq_ptr, int dst_sq_step, int dst_sq_offset +#endif + ) { - __global const vecSumT *srcsum = (__global const vecSumT *)srcsum_ptr; - __global sumT *sum = (__global sumT *)sum_ptr; + __local sumT lm_sum[LOCAL_SUM_STRIDE * LOCAL_SUM_SIZE]; +#ifdef SUM_SQUARE + __local sumSQT lm_sum_sq[LOCAL_SUM_STRIDE * LOCAL_SUM_SIZE]; +#endif int lid = get_local_id(0); int gid = get_group_id(0); - vecSumT src_t[2], sum_t[2]; - __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE]; - __local sumT *sum_p; - src_step = src_step >> 4; - int lid_prim = ((lid & 127) << 1) + 1; - for (int i = 0; i < rows; i += LSIZE_1) + + int gs = get_global_size(0); + + int x = get_global_id(0); + + __global sumT *dst = (__global sumT *)(dst_ptr + dst_offset); + for (int xin = x; xin < cols; xin += gs) { - if (i + lid < rows) + dst[xin] = 0; + } + dst_offset += dst_step; + + if (x < rows - 1) + { + dst = (__global sumT *)(dst_ptr + mad24(x, dst_step, dst_offset)); + dst[0] = 0; + } + + int buf_index = mad24((int)sizeof(sumT), x, buf_offset); + sumT accum = 0; + +#ifdef SUM_SQUARE + __global sumSQT *dst_sq = (__global sumT *)(dst_sq_ptr + dst_sq_offset); + for (int xin = x; xin < cols; xin += gs) + { + dst_sq[xin] = 0; + } + dst_sq_offset += dst_sq_step; + + dst_sq = (__global sumSQT *)(dst_sq_ptr + mad24(x, dst_sq_step, dst_sq_offset)); + dst_sq[0] = 0; + + int buf_sq_index = mad24((int)sizeof(sumSQT), x, buf_sq_offset); + sumSQT accum_sq = 0; +#endif + + for (int y = 1; y < cols; y += LOCAL_SUM_SIZE) + { + int lsum_index = lid; + #pragma unroll + for (int yin = 0; yin < LOCAL_SUM_SIZE; yin++, lsum_index += LOCAL_SUM_STRIDE) { - int sum_idx = mad24(lid + i, src_step, gid * 2); - src_t[0] = srcsum[sum_idx]; - src_t[1] = srcsum[sum_idx + 1]; - } - else - { - src_t[0] = 0; - src_t[1] = 0; - } - if (i == 0) - { - sum_t[0] = 0; - sum_t[1] = 0; - } - else - { - sum_t[0] = lm_sum[0][LSIZE_2 + LOG_LSIZE]; - sum_t[1] = lm_sum[1][LSIZE_2 + LOG_LSIZE]; + __global const sumT *buf = (__global const sumT *)(buf_ptr + buf_index); + accum += buf[0]; + lm_sum[lsum_index] = accum; + buf_index += buf_step; +#ifdef SUM_SQUARE + __global const sumSQT *buf_sq = (__global const sumSQT *)(buf_sq_ptr + buf_sq_index); + accum_sq += buf_sq[0]; + lm_sum_sq[lsum_index] = accum_sq; + buf_sq_index += buf_sq_step; +#endif } barrier(CLK_LOCAL_MEM_FENCE); - int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - - lm_sum[0][bf_loc] = src_t[0]; - lm_sum[1][bf_loc] = src_t[1]; - - int offset = 1; - for (int d = LSIZE >> 1 ; d > 0; d>>=1) + if (y + lid < cols) { - barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * lid_prim - 1, bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) + //int dst_index = dst_offset + dst_step * LOCAL_SUM_COLS * gid + sizeof(sumT) * y + sizeof(sumT) * lid; + int dst_index = mad24(dst_step, LOCAL_SUM_SIZE * gid, mad24((int)sizeof(sumT), y + lid, dst_offset)); +#ifdef SUM_SQUARE + int dst_sq_index = mad24(dst_sq_step, LOCAL_SUM_SIZE * gid, mad24((int)sizeof(sumSQT), y + lid, dst_sq_offset)); +#endif + lsum_index = LOCAL_SUM_STRIDE * lid; + int yin_max = min(rows - 1 - LOCAL_SUM_SIZE * gid, LOCAL_SUM_SIZE); + #pragma unroll + for (int yin = 0; yin < yin_max; yin++, lsum_index++) { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - } - offset <<= 1; - } - barrier(CLK_LOCAL_MEM_FENCE); - if (lid < 2) - { - lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; - } - for (int d = 1; d < LSIZE; d <<= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - offset >>= 1; - int ai = offset * lid_prim - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if ((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - if (gid == 0 && (i + lid) <= rows) - { - sum[sum_offset + i + lid] = 0; - } - if (i + lid == 0) - { - int loc0 = gid * 2 * sum_step; - for(int k = 1; k <= 8; k++) - { - if (gid * 8 + k > cols) - break; - sum[sum_offset + loc0 + k * sum_step / 4] = 0; - } - } - - if (lid > 0 && (i+lid) <= rows) - { - int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ; - lm_sum[0][bf_loc] += sum_t[0]; - lm_sum[1][bf_loc] += sum_t[1]; - sum_p = (__local sumT*)(&(lm_sum[0][bf_loc])); - for(int k = 0; k < 4; k++) - { - if (gid * 8 + k >= cols) - break; - sum[loc_s0 + k * sum_step / 4] = sum_p[k]; - } - sum_p = (__local sumT*)(&(lm_sum[1][bf_loc])); - for(int k = 0; k < 4; k++) - { - if (gid * 8 + 4 + k >= cols) - break; - sum[loc_s1 + k * sum_step / 4] = sum_p[k]; + dst = (__global sumT *)(dst_ptr + dst_index); + dst[0] = lm_sum[lsum_index]; + dst_index += dst_step; +#ifdef SUM_SQUARE + dst_sq = (__global sumSQT *)(dst_sq_ptr + dst_sq_index); + dst_sq[0] = lm_sum_sq[lsum_index]; + dst_sq_index += dst_sq_step; +#endif } } barrier(CLK_LOCAL_MEM_FENCE); diff --git a/modules/imgproc/src/sumpixels.cpp b/modules/imgproc/src/sumpixels.cpp index 1d246ec7bc..e7694b01a9 100755 --- a/modules/imgproc/src/sumpixels.cpp +++ b/modules/imgproc/src/sumpixels.cpp @@ -235,97 +235,87 @@ typedef void (*IntegralFunc)(const uchar* src, size_t srcstep, uchar* sum, size_ #ifdef HAVE_OPENCL -enum { vlen = 4 }; - static bool ocl_integral( InputArray _src, OutputArray _sum, int sdepth ) { - if ( _src.type() != CV_8UC1 || _src.step() % vlen != 0 || _src.offset() % vlen != 0 || - !(sdepth == CV_32S || sdepth == CV_32F) ) + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if ( (_src.type() != CV_8UC1) || + !(sdepth == CV_32S || sdepth == CV_32F || (doubleSupport && sdepth == CV_64F))) return false; - ocl::Kernel k1("integral_sum_cols", ocl::imgproc::integral_sum_oclsrc, - format("-D sdepth=%d", sdepth)); - if (k1.empty()) + static const int tileSize = 16; + + String build_opt = format("-D sumT=%s -D LOCAL_SUM_SIZE=%d%s", + ocl::typeToStr(sdepth), tileSize, + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); + + ocl::Kernel kcols("integral_sum_cols", ocl::imgproc::integral_sum_oclsrc, build_opt); + if (kcols.empty()) return false; - Size size = _src.size(), t_size = Size(((size.height + vlen - 1) / vlen) * vlen, size.width), - ssize(size.width + 1, size.height + 1); - _sum.create(ssize, sdepth); - UMat src = _src.getUMat(), t_sum(t_size, sdepth), sum = _sum.getUMat(); - t_sum = t_sum(Range::all(), Range(0, size.height)); - - int offset = (int)src.offset / vlen; - int vcols = (src.cols + vlen - 1) / vlen; - int sum_offset = (int)sum.offset / vlen; - - k1.args(ocl::KernelArg::PtrReadOnly(src), ocl::KernelArg::PtrWriteOnly(t_sum), - offset, src.rows, src.cols, (int)src.step, (int)t_sum.step); - size_t gt = ((vcols + 1) / 2) * 256, lt = 256; - if (!k1.run(1, >, <, false)) + UMat src = _src.getUMat(); + Size src_size = src.size(); + Size bufsize(((src_size.height + tileSize - 1) / tileSize) * tileSize, ((src_size.width + tileSize - 1) / tileSize) * tileSize); + UMat buf(bufsize, sdepth); + kcols.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnlyNoSize(buf)); + size_t gt = src.cols, lt = tileSize; + if (!kcols.run(1, >, <, false)) return false; - ocl::Kernel k2("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc, - format("-D sdepth=%d", sdepth)); - k2.args(ocl::KernelArg::PtrReadOnly(t_sum), ocl::KernelArg::PtrWriteOnly(sum), - t_sum.rows, t_sum.cols, (int)t_sum.step, (int)sum.step, sum_offset); + ocl::Kernel krows("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc, build_opt); + if (krows.empty()) + return false; - size_t gt2 = t_sum.cols * 32, lt2 = 256; - return k2.run(1, >2, <2, false); + Size sumsize(src_size.width + 1, src_size.height + 1); + _sum.create(sumsize, sdepth); + UMat sum = _sum.getUMat(); + + krows.args(ocl::KernelArg::ReadOnlyNoSize(buf), ocl::KernelArg::WriteOnly(sum)); + gt = src.rows; + return krows.run(1, >, <, false); } static bool ocl_integral( InputArray _src, OutputArray _sum, OutputArray _sqsum, int sdepth, int sqdepth ) { bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; - if ( _src.type() != CV_8UC1 || _src.step() % vlen != 0 || _src.offset() % vlen != 0 || - (!doubleSupport && (sdepth == CV_64F || sqdepth == CV_64F)) ) + if ( _src.type() != CV_8UC1 || (!doubleSupport && (sdepth == CV_64F || sqdepth == CV_64F)) ) return false; - char cvt[40]; - String opts = format("-D sdepth=%d -D sqdepth=%d -D TYPE=%s -D TYPE4=%s4 -D convert_TYPE4=%s%s", - sdepth, sqdepth, ocl::typeToStr(sqdepth), ocl::typeToStr(sqdepth), - ocl::convertTypeStr(sdepth, sqdepth, 4, cvt), - doubleSupport ? " -D DOUBLE_SUPPORT" : ""); + static const int tileSize = 16; - ocl::Kernel k1("integral_cols", ocl::imgproc::integral_sqrsum_oclsrc, opts); - if (k1.empty()) + String build_opt = format("-D SUM_SQUARE -D sumT=%s -D sumSQT=%s -D LOCAL_SUM_SIZE=%d%s", + ocl::typeToStr(sdepth), ocl::typeToStr(sqdepth), + tileSize, + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); + + ocl::Kernel kcols("integral_sum_cols", ocl::imgproc::integral_sum_oclsrc, build_opt); + if (kcols.empty()) return false; - Size size = _src.size(), dsize = Size(size.width + 1, size.height + 1), - t_size = Size(((size.height + vlen - 1) / vlen) * vlen, size.width); - UMat src = _src.getUMat(), t_sum(t_size, sdepth), t_sqsum(t_size, sqdepth); - t_sum = t_sum(Range::all(), Range(0, size.height)); - t_sqsum = t_sqsum(Range::all(), Range(0, size.height)); - - _sum.create(dsize, sdepth); - _sqsum.create(dsize, sqdepth); - UMat sum = _sum.getUMat(), sqsum = _sqsum.getUMat(); - - int offset = (int)src.offset / vlen; - int pre_invalid = src.offset % vlen; - int vcols = (pre_invalid + src.cols + vlen - 1) / vlen; - int sum_offset = (int)(sum.offset / sum.elemSize()); - int sqsum_offset = (int)(sqsum.offset / sqsum.elemSize()); - - k1.args(ocl::KernelArg::PtrReadOnly(src), ocl::KernelArg::PtrWriteOnly(t_sum), - ocl::KernelArg::PtrWriteOnly(t_sqsum), offset, pre_invalid, src.rows, - src.cols, (int)src.step, (int)t_sum.step, (int)t_sqsum.step); - - size_t gt = ((vcols + 1) / 2) * 256, lt = 256; - if (!k1.run(1, >, <, false)) + UMat src = _src.getUMat(); + Size src_size = src.size(); + Size bufsize(((src_size.height + tileSize - 1) / tileSize) * tileSize, ((src_size.width + tileSize - 1) / tileSize) * tileSize); + UMat buf(bufsize, sdepth); + UMat buf_sq(bufsize, sqdepth); + kcols.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnlyNoSize(buf), ocl::KernelArg::WriteOnlyNoSize(buf_sq)); + size_t gt = src.cols, lt = tileSize; + if (!kcols.run(1, >, <, false)) return false; - ocl::Kernel k2("integral_rows", ocl::imgproc::integral_sqrsum_oclsrc, opts); - if (k2.empty()) + ocl::Kernel krows("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc, build_opt); + if (krows.empty()) return false; - k2.args(ocl::KernelArg::PtrReadOnly(t_sum), ocl::KernelArg::PtrReadOnly(t_sqsum), - ocl::KernelArg::PtrWriteOnly(sum), ocl::KernelArg::PtrWriteOnly(sqsum), - t_sum.rows, t_sum.cols, (int)t_sum.step, (int)t_sqsum.step, - (int)sum.step, (int)sqsum.step, sum_offset, sqsum_offset); + Size sumsize(src_size.width + 1, src_size.height + 1); + _sum.create(sumsize, sdepth); + UMat sum = _sum.getUMat(); + _sqsum.create(sumsize, sqdepth); + UMat sum_sq = _sqsum.getUMat(); - size_t gt2 = t_sum.cols * 32, lt2 = 256; - return k2.run(1, >2, <2, false); + krows.args(ocl::KernelArg::ReadOnlyNoSize(buf), ocl::KernelArg::ReadOnlyNoSize(buf_sq), ocl::KernelArg::WriteOnly(sum), ocl::KernelArg::WriteOnlyNoSize(sum_sq)); + gt = src.rows; + return krows.run(1, >, <, false); } #endif From a73809e6fc986ca84f1b7c58492ba7433396b5e6 Mon Sep 17 00:00:00 2001 From: "Fco. Javier Delgado del Hoyo" Date: Wed, 25 Jun 2014 11:46:05 +0200 Subject: [PATCH 06/24] Fix GCC 4.9 compiler warning --- modules/imgproc/src/morph.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index 3ab495d75f..f05c5afc15 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -1257,7 +1257,7 @@ static bool IPPMorphReplicate(int op, const Mat &src, Mat &dst, const Mat &kerne } #undef IPP_MORPH_CASE -#if defined(__GNUC__) && __GNUC__ == 4 && __GNUC_MINOR__ == 8 +#if defined(__GNUC__) && __GNUC__ == 4 && __GNUC_MINOR__ > 8 return false; /// It disables false positive warning in GCC 4.8.2 #endif } From 47e345bcb6177ee701baf27357abeb2d3ab6aa9a Mon Sep 17 00:00:00 2001 From: "Fco. Javier Delgado del Hoyo" Date: Wed, 25 Jun 2014 11:49:26 +0200 Subject: [PATCH 07/24] Change comment according to fix --- modules/imgproc/src/morph.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index f05c5afc15..eaae9b6d0c 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -1258,7 +1258,7 @@ static bool IPPMorphReplicate(int op, const Mat &src, Mat &dst, const Mat &kerne #undef IPP_MORPH_CASE #if defined(__GNUC__) && __GNUC__ == 4 && __GNUC_MINOR__ > 8 - return false; /// It disables false positive warning in GCC 4.8.2 + return false; /// It disables false positive warning in GCC 4.8 and further #endif } } From 730ead44fedeabb6f30f31f2219dc30007d7c6e2 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Thu, 26 Jun 2014 12:46:03 +0400 Subject: [PATCH 08/24] Optimize OpenCL version of sepFilter2D --- modules/imgproc/src/filter.cpp | 2 +- .../src/opencl/filterSep_singlePass.cl | 90 ++++++++++++------- 2 files changed, 59 insertions(+), 33 deletions(-) diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index e51986c394..d23de91ea1 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3492,7 +3492,7 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, return false; size_t lt2[2] = { optimizedSepFilterLocalSize, optimizedSepFilterLocalSize }; - size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), lt2[1] * (1 + (size.height - 1) / lt2[1]) }; + size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), optimizedSepFilterLocalSize}; char cvt[2][40]; const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", diff --git a/modules/imgproc/src/opencl/filterSep_singlePass.cl b/modules/imgproc/src/opencl/filterSep_singlePass.cl index 3952577d77..6c3bbdc161 100644 --- a/modules/imgproc/src/opencl/filterSep_singlePass.cl +++ b/modules/imgproc/src/opencl/filterSep_singlePass.cl @@ -119,17 +119,15 @@ __kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int int liy = get_local_id(1); int x = get_global_id(0); - int y = get_global_id(1); // calculate pixel position in source image taking image offset into account int srcX = x + srcOffsetX - RADIUSX; - int srcY = y + srcOffsetY - RADIUSY; // extrapolate coordinates, if needed // and read my own source pixel into local memory // with account for extra border pixels, which will be read by starting workitems int clocY = liy; - int cSrcY = srcY; + int cSrcY = liy + srcOffsetY - RADIUSY; do { int yb = cSrcY; @@ -154,48 +152,76 @@ __kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int while (clocY < BLK_Y+(RADIUSY*2)); barrier(CLK_LOCAL_MEM_FENCE); - // do vertical filter pass - // and store intermediate results to second local memory array - int i, clocX = lix; - WT sum = (WT) 0; - do + for (int y = 0; y < dst_rows; y+=BLK_Y) { - sum = (WT) 0; - for (i=0; i<=2*RADIUSY; i++) + // do vertical filter pass + // and store intermediate results to second local memory array + int i, clocX = lix; + WT sum = (WT) 0; + do + { + sum = (WT) 0; + for (i=0; i<=2*RADIUSY; i++) #if (defined(INTEGER_ARITHMETIC) && !INTEL_DEVICE) - sum = mad24(lsmem[liy+i][clocX], mat_kernelY[i], sum); + sum = mad24(lsmem[liy + i][clocX], mat_kernelY[i], sum); #else - sum = mad(lsmem[liy+i][clocX], mat_kernelY[i], sum); + sum = mad(lsmem[liy + i][clocX], mat_kernelY[i], sum); #endif - lsmemDy[liy][clocX] = sum; - clocX += BLK_X; - } - while(clocX < BLK_X+(RADIUSX*2)); - barrier(CLK_LOCAL_MEM_FENCE); + lsmemDy[liy][clocX] = sum; + clocX += BLK_X; + } + while(clocX < BLK_X+(RADIUSX*2)); + barrier(CLK_LOCAL_MEM_FENCE); - // if this pixel happened to be out of image borders because of global size rounding, - // then just return - if( x >= dst_cols || y >=dst_rows ) - return; - - // do second horizontal filter pass - // and calculate final result - sum = 0.0f; - for (i=0; i<=2*RADIUSX; i++) + // if this pixel happened to be out of image borders because of global size rounding, + // then just return + if ((x < dst_cols) && (y + liy < dst_rows)) + { + // do second horizontal filter pass + // and calculate final result + sum = 0.0f; + for (i=0; i<=2*RADIUSX; i++) #if (defined(INTEGER_ARITHMETIC) && !INTEL_DEVICE) - sum = mad24(lsmemDy[liy][lix+i], mat_kernelX[i], sum); + sum = mad24(lsmemDy[liy][lix+i], mat_kernelX[i], sum); #else - sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum); + sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum); #endif #ifdef INTEGER_ARITHMETIC #ifdef INTEL_DEVICE - sum = (sum + (1 << (SHIFT_BITS-1))) / (1 << SHIFT_BITS); + sum = (sum + (1 << (SHIFT_BITS-1))) / (1 << SHIFT_BITS); #else - sum = (sum + (1 << (SHIFT_BITS-1))) >> SHIFT_BITS; + sum = (sum + (1 << (SHIFT_BITS-1))) >> SHIFT_BITS; #endif #endif + // store result into destination image + storepix(convertToDstT(sum + (WT)(delta)), Dst + mad24(y + liy, dst_step, mad24(x, DSTSIZE, dst_offset))); + } + + for (int i = liy * BLK_X + lix; i < (RADIUSY*2) * (BLK_X+(RADIUSX*2)); i += BLK_X * BLK_Y) + { + int clocX = i % (BLK_X+(RADIUSX*2)); + int clocY = i / (BLK_X+(RADIUSX*2)); + lsmem[clocY][clocX] = lsmem[clocY + BLK_Y][clocX]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + int cSrcY = y + BLK_Y + liy + srcOffsetY + RADIUSY; + EXTRAPOLATE(cSrcY, (height)); + + clocX = lix; + int cSrcX = x + srcOffsetX - RADIUSX; + do + { + int xb = cSrcX; + EXTRAPOLATE(xb,(width)); + lsmem[liy + 2*RADIUSY][clocX] = ELEM(xb, cSrcY, (width), (height), 0 ); + + clocX += BLK_X; + cSrcX += BLK_X; + } + while(clocX < BLK_X+(RADIUSX*2)); + barrier(CLK_LOCAL_MEM_FENCE); + } - // store result into destination image - storepix(convertToDstT(sum + (WT)(delta)), Dst + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset))); } From 6f2fea7ea5f3f61cee860168700537165f3a77db Mon Sep 17 00:00:00 2001 From: mlyashko Date: Wed, 25 Jun 2014 16:07:33 +0400 Subject: [PATCH 09/24] added perftest for moments --- modules/imgproc/perf/perf_moments.cpp | 38 +++++++++++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 modules/imgproc/perf/perf_moments.cpp diff --git a/modules/imgproc/perf/perf_moments.cpp b/modules/imgproc/perf/perf_moments.cpp new file mode 100644 index 0000000000..9b3c5428f3 --- /dev/null +++ b/modules/imgproc/perf/perf_moments.cpp @@ -0,0 +1,38 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. + +// Copyright (C) 2014, Itseez, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. + +#include "perf_precomp.hpp" + +using namespace std; +using namespace cv; +using namespace perf; +using namespace testing; +using std::tr1::make_tuple; +using std::tr1::get; + +typedef std::tr1::tuple MomentsParams_t; +typedef perf::TestBaseWithParam MomentsFixture_val; + +PERF_TEST_P(MomentsFixture_val, Moments1, + ::testing::Combine( + testing::Values(TYPICAL_MAT_SIZES), + testing::Values(CV_16U, CV_16S, CV_32F, CV_64F), + testing::Bool())) +{ + const MomentsParams_t params = GetParam(); + const Size srcSize = get<0>(params); + const MatDepth srcDepth = get<1>(params); + const bool binaryImage = get<2>(params); + + cv::Moments m; + Mat src(srcSize, srcDepth); + declare.in(src, WARMUP_RNG); + + TEST_CYCLE() m = cv::moments(src, binaryImage); + + SANITY_CHECK_MOMENTS(m, 1e-4, ERROR_RELATIVE); +} From ad9272e83652f8c21eaaa89840e00898107b9778 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 26 Jun 2014 13:13:01 +0400 Subject: [PATCH 10/24] reverted to original plain C++ code --- modules/imgproc/src/morph.cpp | 93 +++++++++-------------------------- 1 file changed, 22 insertions(+), 71 deletions(-) diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index 3ab495d75f..6e31673a5f 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -1545,93 +1545,44 @@ void cv::morphologyEx( InputArray _src, OutputArray _dst, int op, InputArray kernel, Point anchor, int iterations, int borderType, const Scalar& borderValue ) { - int src_type = _src.type(), dst_type = _dst.type(), - src_cn = CV_MAT_CN(src_type), src_depth = CV_MAT_DEPTH(src_type); - - bool use_opencl = cv::ocl::useOpenCL() && _src.isUMat() && _src.size() == _dst.size() && src_type == dst_type && - _src.dims()<=2 && (src_cn == 1 || src_cn == 4) && (anchor.x == -1) && (anchor.y == -1) && - (src_depth == CV_8U || src_depth == CV_32F || src_depth == CV_64F ) && - (borderType == cv::BORDER_CONSTANT) && (borderValue == morphologyDefaultBorderValue()); - - _dst.create(_src.size(), _src.type()); - Mat src, dst, temp; - UMat usrc, udst, utemp; + Mat src = _src.getMat(), temp; + _dst.create(src.size(), src.type()); + Mat dst = _dst.getMat(); switch( op ) { case MORPH_ERODE: - erode( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + erode( src, dst, kernel, anchor, iterations, borderType, borderValue ); break; case MORPH_DILATE: - dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + dilate( src, dst, kernel, anchor, iterations, borderType, borderValue ); break; case MORPH_OPEN: - erode( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); - dilate( _dst, _dst, kernel, anchor, iterations, borderType, borderValue ); + erode( src, dst, kernel, anchor, iterations, borderType, borderValue ); + dilate( dst, dst, kernel, anchor, iterations, borderType, borderValue ); break; case CV_MOP_CLOSE: - dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); - erode( _dst, _dst, kernel, anchor, iterations, borderType, borderValue ); + dilate( src, dst, kernel, anchor, iterations, borderType, borderValue ); + erode( dst, dst, kernel, anchor, iterations, borderType, borderValue ); break; case CV_MOP_GRADIENT: - erode( _src, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, anchor, iterations, borderType, borderValue ); - dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); - if(use_opencl) - { - udst = _dst.getUMat(); - subtract(udst, utemp, udst); - } - else - { - dst = _dst.getMat(); - dst -= temp; - } + erode( src, temp, kernel, anchor, iterations, borderType, borderValue ); + dilate( src, dst, kernel, anchor, iterations, borderType, borderValue ); + dst -= temp; break; case CV_MOP_TOPHAT: - if(use_opencl) - { - usrc = _src.getUMat(); - udst = _dst.getUMat(); - if( usrc.u != udst.u ) - utemp = udst; - } - else - { - src = _src.getMat(); - dst = _dst.getMat(); - if( src.data != dst.data ) - temp = dst; - } - erode( _src, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, anchor, iterations, borderType, borderValue ); - dilate( use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, - anchor, iterations, borderType, borderValue ); - if(use_opencl) - subtract(usrc, utemp, udst); - else - dst = src - temp; + if( src.data != dst.data ) + temp = dst; + erode( src, temp, kernel, anchor, iterations, borderType, borderValue ); + dilate( temp, temp, kernel, anchor, iterations, borderType, borderValue ); + dst = src - temp; break; case CV_MOP_BLACKHAT: - if(use_opencl) - { - usrc = _src.getUMat(); - udst = _dst.getUMat(); - if( usrc.u != udst.u ) - utemp = udst; - } - else - { - src = _src.getMat(); - dst = _dst.getMat(); - if( src.data != dst.data ) - temp = dst; - } - dilate( _src, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, anchor, iterations, borderType, borderValue ); - erode( use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, - anchor, iterations, borderType, borderValue ); - if(use_opencl) - subtract(utemp, usrc, udst); - else - dst = temp - src; + if( src.data != dst.data ) + temp = dst; + dilate( src, temp, kernel, anchor, iterations, borderType, borderValue ); + erode( temp, temp, kernel, anchor, iterations, borderType, borderValue ); + dst = temp - src; break; default: CV_Error( CV_StsBadArg, "unknown morphological operation" ); From a3592cd068231d6cf7dfcfb092cfbcaaa2d4c41f Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 26 Jun 2014 13:18:03 +0400 Subject: [PATCH 11/24] added ocl_** function --- modules/imgproc/src/morph.cpp | 58 +++++++++++++++++++++++++++++++++++ 1 file changed, 58 insertions(+) diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index 6e31673a5f..097179f99e 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -1541,10 +1541,68 @@ void cv::dilate( InputArray src, OutputArray dst, InputArray kernel, morphOp( MORPH_DILATE, src, dst, kernel, anchor, iterations, borderType, borderValue ); } +#ifdef HAVE_OPENCL + +static void ocl_morphologyEx(InputArray _src, OutputArray _dst, int op, + InputArray kernel, Point anchor, int iterations, + int borderType, const Scalar& borderValue) +{ + int type = _src.type(), cn = CV_MAT_CN(type); + Size ksize = kernel.size(); + + _dst.create(_src.size(), type); + UMat temp; + + switch( op ) + { + case MORPH_ERODE: + erode( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + break; + case MORPH_DILATE: + dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + break; + case MORPH_OPEN: + erode( _src, temp, kernel, anchor, iterations, borderType, borderValue ); + dilate( temp, _dst, kernel, anchor, iterations, borderType, borderValue ); + break; + case CV_MOP_CLOSE: + dilate( _src, temp, kernel, anchor, iterations, borderType, borderValue ); + erode( temp, _dst, kernel, anchor, iterations, borderType, borderValue ); + break; + case CV_MOP_GRADIENT: + // ?? + erode( _src, temp, kernel, anchor, iterations, borderType, borderValue ); + dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + subtract(_dst, temp, _dst); + break; + case CV_MOP_TOPHAT: + // ?? + erode( _src, temp, kernel, anchor, iterations, borderType, borderValue ); + dilate( temp, _dst, kernel, anchor, iterations, borderType, borderValue ); + subtract(_src, _dst, _dst); + break; + case CV_MOP_BLACKHAT: + // ?? + dilate( _src, temp, kernel, anchor, iterations, borderType, borderValue ); + erode( temp, _dst, kernel, anchor, iterations, borderType, borderValue ); + subtract(_dst, _src, _dst); + break; + default: + CV_Error( CV_StsBadArg, "unknown morphological operation" ); + } +} + +#endif + void cv::morphologyEx( InputArray _src, OutputArray _dst, int op, InputArray kernel, Point anchor, int iterations, int borderType, const Scalar& borderValue ) { + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && cn <= 4 && + anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1 && + borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue(), + ocl_morphologyEx(_src, _dst, op, kernel, anchor, iterations, borderType, borderValue)) + Mat src = _src.getMat(), temp; _dst.create(src.size(), src.type()); Mat dst = _dst.getMat(); From 16ab6ec534a6c45b9d8090d3146d70b27c87348c Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 26 Jun 2014 13:25:24 +0400 Subject: [PATCH 12/24] ported changes from PR #2867 --- modules/imgproc/src/morph.cpp | 171 +++++++++++++++------------- modules/imgproc/src/opencl/morph.cl | 89 +++++++-------- 2 files changed, 137 insertions(+), 123 deletions(-) diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index 097179f99e..c1a5d28245 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -1342,58 +1342,80 @@ static bool IPPMorphOp(int op, InputArray _src, OutputArray _dst, #ifdef HAVE_OPENCL -static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, - const Size & ksize, const Point & anchor, int iterations, int op) +static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel, + Point anchor, int iterations, int op, int borderType, + const Scalar& borderValue) { - CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE); + if (borderType != BORDER_CONSTANT) + return false; + Mat kernel = _kernel.getMat(); + Size ksize = kernel.data ? kernel.size() : Size(3,3); + anchor = normalizeAnchor(anchor, ksize); + + if (iterations == 0 || kernel.rows*kernel.cols == 1) + { + _src.copyTo(_dst); + return true; + } + + if (!kernel.data) + { + kernel = getStructuringElement(MORPH_RECT, Size(1+iterations*2,1+iterations*2)); + anchor = Point(iterations, iterations); + iterations = 1; + } + else if( iterations > 1 && countNonZero(kernel) == kernel.rows*kernel.cols ) + { + anchor = Point(anchor.x*iterations, anchor.y*iterations); + kernel = getStructuringElement(MORPH_RECT, + Size(ksize.width + (iterations-1)*(ksize.width-1), + ksize.height + (iterations-1)*(ksize.height-1)), + anchor); + iterations = 1; + } + + const ocl::Device & dev = ocl::Device::getDefault(); int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); - bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + bool doubleSupport = dev.doubleFPConfig() > 0; if (depth == CV_64F && !doubleSupport) return false; - UMat kernel8U; - kernel.convertTo(kernel8U, CV_8U); - kernel8U = kernel8U.reshape(1, 1); - - bool rectKernel = true; - { - Mat m = kernel.reshape(1, 1); - for (int i = 0; i < m.size().area(); ++i) - if (m.at(i) != 1) - { - rectKernel = false; - break; - } - } - UMat src = _src.getUMat(); #ifdef ANDROID - size_t localThreads[3] = {16, 8, 1}; + size_t localThreads[2] = { 16, 8 }; #else - size_t localThreads[3] = {16, 16, 1}; + size_t localThreads[2] = { 16, 16 }; #endif - size_t globalThreads[3] = {(src.cols + localThreads[0] - 1) / localThreads[0] *localThreads[0], (src.rows + localThreads[1] - 1) / localThreads[1] *localThreads[1], 1}; + size_t globalThreads[2] = { src.cols, src.rows }; if (localThreads[0]*localThreads[1] * 2 < (localThreads[0] + ksize.width - 1) * (localThreads[1] + ksize.height - 1)) return false; - static const char * const op2str[] = { "ERODE", "DILATE" }; - String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s" - " -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s", anchor.x, anchor.y, - (int)localThreads[0], (int)localThreads[1], op2str[op], - doubleSupport ? " -D DOUBLE_SUPPORT" : "", rectKernel ? " -D RECTKERNEL" : "", - ocl::typeToStr(_src.type()), _src.depth(), cn, ocl::typeToStr(depth)); + // build processing + String processing; + Mat kernel8u; + kernel.convertTo(kernel8u, CV_8U); + for (int y = 0; y < kernel8u.rows; ++y) + for (int x = 0; x < kernel8u.cols; ++x) + if (kernel8u.at(y, x) != 0) + processing += format("PROCESS(%d,%d)", y, x); - std::vector kernels; + static const char * const op2str[] = { "ERODE", "DILATE" }; + String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s" + " -D PROCESS_ELEMS=%s -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s", anchor.x, anchor.y, + (int)localThreads[0], (int)localThreads[1], op2str[op], + doubleSupport ? " -D DOUBLE_SUPPORT" : "", processing.c_str(), + ocl::typeToStr(type), depth, cn, ocl::typeToStr(depth)); + + std::vector kernels(iterations); for (int i = 0; i < iterations; i++) { - ocl::Kernel k("morph", ocl::imgproc::morph_oclsrc, buildOptions); - if (k.empty()) + kernels[i].create("morph", ocl::imgproc::morph_oclsrc, buildOptions); + if (kernels[i].empty()) return false; - kernels.push_back(k); } _dst.create(src.size(), src.type()); @@ -1407,8 +1429,7 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, int wholecols = wholesize.width, wholerows = wholesize.height; kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst), - ofs.x, ofs.y, src.cols, src.rows, ocl::KernelArg::PtrReadOnly(kernel8U), - wholecols, wholerows); + ofs.x, ofs.y, src.cols, src.rows, wholecols, wholerows); return kernels[0].run(2, globalThreads, localThreads, false); } @@ -1422,19 +1443,20 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, if (i == 0) { int cols = src.cols, rows = src.rows; - src.locateROI(wholesize,ofs); + src.locateROI(wholesize, ofs); src.adjustROI(ofs.y, wholesize.height - rows - ofs.y, ofs.x, wholesize.width - cols - ofs.x); if(src.u != dst.u) source = src; else src.copyTo(source); + src.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x); source.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x); } else { int cols = dst.cols, rows = dst.rows; - dst.locateROI(wholesize,ofs); + dst.locateROI(wholesize, ofs); dst.adjustROI(ofs.y, wholesize.height - rows - ofs.y, ofs.x, wholesize.width - cols - ofs.x); dst.copyTo(source); dst.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x); @@ -1443,12 +1465,12 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, source.locateROI(wholesize, ofs); kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst), - ofs.x, ofs.y, source.cols, source.rows, ocl::KernelArg::PtrReadOnly(kernel8U), - wholesize.width, wholesize.height); + ofs.x, ofs.y, source.cols, source.rows, wholesize.width, wholesize.height); if (!kernels[i].run(2, globalThreads, localThreads, false)) return false; } + return true; } @@ -1459,10 +1481,10 @@ static void morphOp( int op, InputArray _src, OutputArray _dst, Point anchor, int iterations, int borderType, const Scalar& borderValue ) { -#ifdef HAVE_OPENCL - int src_type = _src.type(), - src_cn = CV_MAT_CN(src_type), src_depth = CV_MAT_DEPTH(src_type); -#endif + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && _src.channels() <= 4 && + borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() && + (op == MORPH_ERODE || op == MORPH_DILATE), + ocl_morphOp(_src, _dst, _kernel, anchor, iterations, op, borderType, borderValue) ) Mat kernel = _kernel.getMat(); Size ksize = kernel.data ? kernel.size() : Size(3,3); @@ -1490,12 +1512,6 @@ static void morphOp( int op, InputArray _src, OutputArray _dst, iterations = 1; } - CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && src_cn <= 4 && - (src_depth == CV_8U || src_depth == CV_32F || src_depth == CV_64F ) && - borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() && - (op == MORPH_ERODE || op == MORPH_DILATE), - ocl_morphology_op(_src, _dst, kernel, ksize, anchor, iterations, op) ) - #if IPP_VERSION_X100 >= 801 if( IPPMorphOp(op, _src, _dst, kernel, anchor, iterations, borderType, borderValue) ) return; @@ -1515,13 +1531,6 @@ static void morphOp( int op, InputArray _src, OutputArray _dst, parallel_for_(Range(0, nStripes), MorphologyRunner(src, dst, nStripes, iterations, op, kernel, anchor, borderType, borderType, borderValue)); - - //Ptr f = createMorphologyFilter(op, src.type(), - // kernel, anchor, borderType, borderType, borderValue ); - - //f->apply( src, dst ); - //for( int i = 1; i < iterations; i++ ) - // f->apply( dst, dst ); } } @@ -1543,53 +1552,56 @@ void cv::dilate( InputArray src, OutputArray dst, InputArray kernel, #ifdef HAVE_OPENCL -static void ocl_morphologyEx(InputArray _src, OutputArray _dst, int op, +namespace cv { + +static bool ocl_morphologyEx(InputArray _src, OutputArray _dst, int op, InputArray kernel, Point anchor, int iterations, int borderType, const Scalar& borderValue) { - int type = _src.type(), cn = CV_MAT_CN(type); - Size ksize = kernel.size(); - - _dst.create(_src.size(), type); + _dst.createSameSize(_src, _src.type()); UMat temp; switch( op ) { case MORPH_ERODE: - erode( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); break; case MORPH_DILATE: - dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); break; case MORPH_OPEN: - erode( _src, temp, kernel, anchor, iterations, borderType, borderValue ); - dilate( temp, _dst, kernel, anchor, iterations, borderType, borderValue ); + ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); + ocl_morphOp( temp, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); break; - case CV_MOP_CLOSE: - dilate( _src, temp, kernel, anchor, iterations, borderType, borderValue ); - erode( temp, _dst, kernel, anchor, iterations, borderType, borderValue ); + case MORPH_CLOSE: + ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); + ocl_morphOp( temp, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); break; - case CV_MOP_GRADIENT: + case MORPH_GRADIENT: // ?? - erode( _src, temp, kernel, anchor, iterations, borderType, borderValue ); - dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); + ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); subtract(_dst, temp, _dst); break; - case CV_MOP_TOPHAT: + case MORPH_TOPHAT: // ?? - erode( _src, temp, kernel, anchor, iterations, borderType, borderValue ); - dilate( temp, _dst, kernel, anchor, iterations, borderType, borderValue ); + ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); + ocl_morphOp( temp, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); subtract(_src, _dst, _dst); break; - case CV_MOP_BLACKHAT: + case MORPH_BLACKHAT: // ?? - dilate( _src, temp, kernel, anchor, iterations, borderType, borderValue ); - erode( temp, _dst, kernel, anchor, iterations, borderType, borderValue ); + ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); + ocl_morphOp( temp, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); subtract(_dst, _src, _dst); break; default: CV_Error( CV_StsBadArg, "unknown morphological operation" ); } + + return true; +} + } #endif @@ -1598,10 +1610,13 @@ void cv::morphologyEx( InputArray _src, OutputArray _dst, int op, InputArray kernel, Point anchor, int iterations, int borderType, const Scalar& borderValue ) { - CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && cn <= 4 && +#ifdef HAVE_OPENCL + Size ksize = kernel.size(); + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && _src.channels() <= 4 && anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1 && - borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue(), + borderType == cv::BORDER_CONSTANT, ocl_morphologyEx(_src, _dst, op, kernel, anchor, iterations, borderType, borderValue)) +#endif Mat src = _src.getMat(), temp; _dst.create(src.size(), src.type()); diff --git a/modules/imgproc/src/opencl/morph.cl b/modules/imgproc/src/opencl/morph.cl index a7611c50f9..7df09ec61e 100644 --- a/modules/imgproc/src/opencl/morph.cl +++ b/modules/imgproc/src/opencl/morph.cl @@ -54,59 +54,70 @@ #endif #ifdef DEPTH_0 -#ifdef ERODE -#define VAL 255 -#endif -#ifdef DILATE -#define VAL 0 -#endif +#define MIN_VAL 0 +#define MAX_VAL UCHAR_MAX +#elif defined DEPTH_1 +#define MIN_VAL SCHAR_MIN +#define MAX_VAL SCHAR_MAX +#elif defined DEPTH_2 +#define MIN_VAL 0 +#define MAX_VAL USHRT_MAX +#elif defined DEPTH_3 +#define MIN_VAL SHRT_MIN +#define MAX_VAL SHRT_MAX +#elif defined DEPTH_4 +#define MIN_VAL INT_MIN +#define MAX_VAL INT_MAX #elif defined DEPTH_5 -#ifdef ERODE -#define VAL FLT_MAX -#endif -#ifdef DILATE -#define VAL -FLT_MAX -#endif +#define MIN_VAL (-FLT_MAX) +#define MAX_VAL FLT_MAX #elif defined DEPTH_6 -#ifdef ERODE -#define VAL DBL_MAX -#endif -#ifdef DILATE -#define VAL -DBL_MAX -#endif +#define MIN_VAL (-DBL_MAX) +#define MAX_VAL DBL_MAX #endif -#ifdef ERODE -#if defined(INTEL_DEVICE) && (DEPTH_0) +#ifdef OP_ERODE +#define VAL MAX_VAL +#elif defined OP_DILATE +#define VAL MIN_VAL +#else +#error "Unknown operation" +#endif + +#ifdef OP_ERODE +#if defined INTEL_DEVICE && defined DEPTH_0 // workaround for bug in Intel HD graphics drivers (10.18.10.3496 or older) #define __CAT(x, y) x##y #define CAT(x, y) __CAT(x, y) #define WA_CONVERT_1 CAT(convert_uint, cn) #define WA_CONVERT_2 CAT(convert_, T) #define convert_uint1 convert_uint -#define MORPH_OP(A,B) WA_CONVERT_2(min(WA_CONVERT_1(A),WA_CONVERT_1(B))) +#define MORPH_OP(A, B) WA_CONVERT_2(min(WA_CONVERT_1(A), WA_CONVERT_1(B))) #else -#define MORPH_OP(A,B) min((A),(B)) +#define MORPH_OP(A, B) min((A), (B)) #endif #endif -#ifdef DILATE -#define MORPH_OP(A,B) max((A),(B)) +#ifdef OP_DILATE +#define MORPH_OP(A, B) max((A), (B)) #endif +#define PROCESS(y, x) \ + res = MORPH_OP(res, LDS_DAT[mad24(l_y + y, width, l_x + x)]); + // BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii #define ELEM(i, l_edge, r_edge, elem1, elem2) (i) < (l_edge) | (i) >= (r_edge) ? (elem1) : (elem2) + __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, __global uchar * dstptr, int dst_step, int dst_offset, int src_offset_x, int src_offset_y, int cols, int rows, - __constant uchar * mat_kernel, int src_whole_cols, int src_whole_rows) + int src_whole_cols, int src_whole_rows EXTRA_PARAMS) { int gidx = get_global_id(0), gidy = get_global_id(1); int l_x = get_local_id(0), l_y = get_local_id(1); int x = get_group_id(0) * LSIZE0, y = get_group_id(1) * LSIZE1; int start_x = x + src_offset_x - RADIUSX; - int end_x = x + src_offset_x + LSIZE0 + RADIUSX; - int width = end_x - (x + src_offset_x - RADIUSX) + 1; + int width = mad24(RADIUSX, 2, LSIZE0 + 1); int start_y = y + src_offset_y - RADIUSY; int point1 = mad24(l_y, LSIZE0, l_x); int point2 = point1 + LSIZE0 * LSIZE1; @@ -117,7 +128,7 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, int start_addr = mad24(cur_y, src_step, cur_x * TSIZE); int start_addr2 = mad24(cur_y2, src_step, cur_x2 * TSIZE); - __local T LDS_DAT[2*LSIZE1*LSIZE0]; + __local T LDS_DAT[2 * LSIZE1 * LSIZE0]; // read pixels from src int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * TSIZE); @@ -128,8 +139,8 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, T temp1 = loadpix(srcptr + start_addr2); // judge if read out of boundary - temp0 = ELEM(cur_x, 0, src_whole_cols, (T)(VAL),temp0); - temp0 = ELEM(cur_y, 0, src_whole_rows, (T)(VAL),temp0); + temp0 = ELEM(cur_x, 0, src_whole_cols, (T)(VAL), temp0); + temp0 = ELEM(cur_y, 0, src_whole_rows, (T)(VAL), temp0); temp1 = ELEM(cur_x2, 0, src_whole_cols, (T)(VAL), temp1); temp1 = ELEM(cur_y2, 0, src_whole_rows, (T)(VAL), temp1); @@ -138,23 +149,11 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, LDS_DAT[point2] = temp1; barrier(CLK_LOCAL_MEM_FENCE); - T res = (T)(VAL); - for (int i = 0, sizey = 2 * RADIUSY + 1; i < sizey; i++) - for (int j = 0, sizex = 2 * RADIUSX + 1; j < sizex; j++) - { - res = -#ifndef RECTKERNEL - mat_kernel[i*(2*RADIUSX+1)+j] ? -#endif - MORPH_OP(res, LDS_DAT[mad24(l_y + i, width, l_x + j)]) -#ifndef RECTKERNEL - : res -#endif - ; - } - if (gidx < cols && gidy < rows) { + T res = (T)(VAL); + PROCESS_ELEMS; + int dst_index = mad24(gidy, dst_step, mad24(gidx, TSIZE, dst_offset)); storepix(res, dstptr + dst_index); } From 1a73aa1f6a3e25cdbe749e4ceaf1a75ece0fb6fb Mon Sep 17 00:00:00 2001 From: vbystricky Date: Thu, 26 Jun 2014 15:43:40 +0400 Subject: [PATCH 13/24] Change local size --- modules/imgproc/src/filter.cpp | 11 ++++++----- modules/imgproc/src/opencl/filterSep_singlePass.cl | 10 ++++------ 2 files changed, 10 insertions(+), 11 deletions(-) diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index d23de91ea1..6c0da79ccf 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3471,7 +3471,8 @@ static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY return k.run(2, globalsize, localsize, false); } -const int optimizedSepFilterLocalSize = 16; +const int optimizedSepFilterLocalWidth = 16; +const int optimizedSepFilterLocalHeight = 8; static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, Mat row_kernel, Mat col_kernel, @@ -3491,8 +3492,8 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, borderType == BORDER_REFLECT_101)) return false; - size_t lt2[2] = { optimizedSepFilterLocalSize, optimizedSepFilterLocalSize }; - size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), optimizedSepFilterLocalSize}; + size_t lt2[2] = { optimizedSepFilterLocalWidth, optimizedSepFilterLocalHeight }; + size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), lt2[1]}; char cvt[2][40]; const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", @@ -3584,8 +3585,8 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, } CV_OCL_RUN_(kernelY.cols <= 21 && kernelX.cols <= 21 && - imgSize.width > optimizedSepFilterLocalSize + anchor.x && - imgSize.height > optimizedSepFilterLocalSize + anchor.y && + imgSize.width > optimizedSepFilterLocalWidth + anchor.x && + imgSize.height > optimizedSepFilterLocalHeight + anchor.y && (!(borderType & BORDER_ISOLATED) || _src.offset() == 0) && anchor == Point(kernelX.cols >> 1, kernelY.cols >> 1) && (d.isIntel() || (d.isAMD() && !d.hostUnifiedMemory())), diff --git a/modules/imgproc/src/opencl/filterSep_singlePass.cl b/modules/imgproc/src/opencl/filterSep_singlePass.cl index 6c3bbdc161..8c14f2d77e 100644 --- a/modules/imgproc/src/opencl/filterSep_singlePass.cl +++ b/modules/imgproc/src/opencl/filterSep_singlePass.cl @@ -127,10 +127,9 @@ __kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int // and read my own source pixel into local memory // with account for extra border pixels, which will be read by starting workitems int clocY = liy; - int cSrcY = liy + srcOffsetY - RADIUSY; do { - int yb = cSrcY; + int yb = clocY + srcOffsetY - RADIUSY; EXTRAPOLATE(yb, (height)); int clocX = lix; @@ -147,7 +146,6 @@ __kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int while(clocX < BLK_X+(RADIUSX*2)); clocY += BLK_Y; - cSrcY += BLK_Y; } while (clocY < BLK_Y+(RADIUSY*2)); barrier(CLK_LOCAL_MEM_FENCE); @@ -206,8 +204,8 @@ __kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int } barrier(CLK_LOCAL_MEM_FENCE); - int cSrcY = y + BLK_Y + liy + srcOffsetY + RADIUSY; - EXTRAPOLATE(cSrcY, (height)); + int yb = y + liy + BLK_Y + srcOffsetY + RADIUSY; + EXTRAPOLATE(yb, (height)); clocX = lix; int cSrcX = x + srcOffsetX - RADIUSX; @@ -215,7 +213,7 @@ __kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int { int xb = cSrcX; EXTRAPOLATE(xb,(width)); - lsmem[liy + 2*RADIUSY][clocX] = ELEM(xb, cSrcY, (width), (height), 0 ); + lsmem[liy + 2*RADIUSY][clocX] = ELEM(xb, yb, (width), (height), 0 ); clocX += BLK_X; cSrcX += BLK_X; From 36db85a94dac360d2eb1cd883b3f3693e3ee9453 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 26 Jun 2014 18:15:13 +0400 Subject: [PATCH 14/24] optimized some operations --- modules/imgproc/src/morph.cpp | 130 +++++++++++++--------- modules/imgproc/src/opencl/morph.cl | 21 ++++ modules/imgproc/test/ocl/test_filters.cpp | 32 +++--- 3 files changed, 116 insertions(+), 67 deletions(-) diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index c1a5d28245..5eec87cb26 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -1333,10 +1333,7 @@ static bool IPPMorphOp(int op, InputArray _src, OutputArray _dst, if( iterations > 1 ) return false; - if (IPPMorphReplicate( op, src, dst, kernel, ksize, anchor, rectKernel )) - return true; - - return false; + return IPPMorphReplicate( op, src, dst, kernel, ksize, anchor, rectKernel ); } #endif @@ -1344,14 +1341,19 @@ static bool IPPMorphOp(int op, InputArray _src, OutputArray _dst, static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel, Point anchor, int iterations, int op, int borderType, - const Scalar& borderValue) + const Scalar &, int actual_op = -1, InputArray _extraMat = noArray()) { - if (borderType != BORDER_CONSTANT) + const ocl::Device & dev = ocl::Device::getDefault(); + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + bool doubleSupport = dev.doubleFPConfig() > 0; + + if ((depth == CV_64F && !doubleSupport) || borderType != BORDER_CONSTANT) return false; Mat kernel = _kernel.getMat(); - Size ksize = kernel.data ? kernel.size() : Size(3,3); - anchor = normalizeAnchor(anchor, ksize); + bool haveExtraMat = !_extraMat.empty(); + Size ksize = kernel.data ? kernel.size() : Size(3, 3), ssize = _src.size(); + CV_Assert(actual_op <= 3 || haveExtraMat); if (iterations == 0 || kernel.rows*kernel.cols == 1) { @@ -1375,21 +1377,12 @@ static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel, iterations = 1; } - const ocl::Device & dev = ocl::Device::getDefault(); - int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); - bool doubleSupport = dev.doubleFPConfig() > 0; - - if (depth == CV_64F && !doubleSupport) - return false; - - UMat src = _src.getUMat(); - #ifdef ANDROID size_t localThreads[2] = { 16, 8 }; #else size_t localThreads[2] = { 16, 16 }; #endif - size_t globalThreads[2] = { src.cols, src.rows }; + size_t globalThreads[2] = { ssize.width, ssize.height }; if (localThreads[0]*localThreads[1] * 2 < (localThreads[0] + ksize.width - 1) * (localThreads[1] + ksize.height - 1)) return false; @@ -1403,21 +1396,35 @@ static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel, if (kernel8u.at(y, x) != 0) processing += format("PROCESS(%d,%d)", y, x); - static const char * const op2str[] = { "ERODE", "DILATE" }; - String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s" - " -D PROCESS_ELEMS=%s -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s", anchor.x, anchor.y, - (int)localThreads[0], (int)localThreads[1], op2str[op], - doubleSupport ? " -D DOUBLE_SUPPORT" : "", processing.c_str(), - ocl::typeToStr(type), depth, cn, ocl::typeToStr(depth)); + static const char * const op2str[] = { "OP_ERODE", "OP_DILATE", NULL, NULL, "OP_GRADIENT", "OP_TOPHAT", "OP_BLACKHAT" }; + + char cvt[2][50]; + int wdepth = std::max(depth, CV_32F), scalarcn = cn == 3 ? 4 : cn; + + if (actual_op < 0) + actual_op = op; std::vector kernels(iterations); for (int i = 0; i < iterations; i++) { + int current_op = iterations == i + 1 ? actual_op : op; + String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s" + " -D PROCESS_ELEMS=%s -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s" + " -D convertToWT=%s -D convertToT=%s -D ST=%s%s", + anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], op2str[op], + doubleSupport ? " -D DOUBLE_SUPPORT" : "", processing.c_str(), + ocl::typeToStr(type), depth, cn, ocl::typeToStr(depth), + ocl::convertTypeStr(depth, wdepth, cn, cvt[0]), + ocl::convertTypeStr(wdepth, depth, cn, cvt[1]), + ocl::typeToStr(CV_MAKE_TYPE(depth, scalarcn)), + current_op == op ? "" : cv::format(" -D %s", op2str[current_op]).c_str()); + kernels[i].create("morph", ocl::imgproc::morph_oclsrc, buildOptions); if (kernels[i].empty()) return false; } + UMat src = _src.getUMat(), extraMat = _extraMat.getUMat(); _dst.create(src.size(), src.type()); UMat dst = _dst.getUMat(); @@ -1428,7 +1435,12 @@ static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel, src.locateROI(wholesize, ofs); int wholecols = wholesize.width, wholerows = wholesize.height; - kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst), + if (haveExtraMat) + kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst), + ofs.x, ofs.y, src.cols, src.rows, wholecols, wholerows, + ocl::KernelArg::ReadOnlyNoSize(extraMat)); + else + kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst), ofs.x, ofs.y, src.cols, src.rows, wholecols, wholerows); return kernels[0].run(2, globalThreads, localThreads, false); @@ -1464,8 +1476,13 @@ static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel, } source.locateROI(wholesize, ofs); - kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst), - ofs.x, ofs.y, source.cols, source.rows, wholesize.width, wholesize.height); + if (haveExtraMat && iterations == i + 1) + kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst), + ofs.x, ofs.y, source.cols, source.rows, wholesize.width, wholesize.height, + ocl::KernelArg::ReadOnlyNoSize(extraMat)); + else + kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst), + ofs.x, ofs.y, source.cols, source.rows, wholesize.width, wholesize.height); if (!kernels[i].run(2, globalThreads, localThreads, false)) return false; @@ -1481,15 +1498,16 @@ static void morphOp( int op, InputArray _src, OutputArray _dst, Point anchor, int iterations, int borderType, const Scalar& borderValue ) { - CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && _src.channels() <= 4 && - borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() && - (op == MORPH_ERODE || op == MORPH_DILATE), - ocl_morphOp(_src, _dst, _kernel, anchor, iterations, op, borderType, borderValue) ) - Mat kernel = _kernel.getMat(); Size ksize = kernel.data ? kernel.size() : Size(3,3); anchor = normalizeAnchor(anchor, ksize); + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && _src.channels() <= 4 && + borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() && + (op == MORPH_ERODE || op == MORPH_DILATE) && + anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1, + ocl_morphOp(_src, _dst, kernel, anchor, iterations, op, borderType, borderValue) ) + if (iterations == 0 || kernel.rows*kernel.cols == 1) { _src.copyTo(_dst); @@ -1559,41 +1577,49 @@ static bool ocl_morphologyEx(InputArray _src, OutputArray _dst, int op, int borderType, const Scalar& borderValue) { _dst.createSameSize(_src, _src.type()); + bool submat = _dst.isSubmatrix(); UMat temp; + _OutputArray _temp = submat ? _dst : _OutputArray(temp); switch( op ) { case MORPH_ERODE: - ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); + if (!ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; break; case MORPH_DILATE: - ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); + if (!ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue )) + return false; break; case MORPH_OPEN: - ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); - ocl_morphOp( temp, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); + if (!ocl_morphOp( _src, _temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _temp, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue )) + return false; break; case MORPH_CLOSE: - ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); - ocl_morphOp( temp, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); + if (!ocl_morphOp( _src, _temp, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _temp, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; break; case MORPH_GRADIENT: - // ?? - ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); - ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); - subtract(_dst, temp, _dst); + if (!ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue, MORPH_GRADIENT, temp )) + return false; break; case MORPH_TOPHAT: - // ?? - ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); - ocl_morphOp( temp, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); - subtract(_src, _dst, _dst); + if (!ocl_morphOp( _src, _temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _temp, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue, MORPH_TOPHAT, _src )) + return false; break; case MORPH_BLACKHAT: - // ?? - ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue ); - ocl_morphOp( temp, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue ); - subtract(_dst, _src, _dst); + if (!ocl_morphOp( _src, _temp, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _temp, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue, MORPH_BLACKHAT, _src )) + return false; break; default: CV_Error( CV_StsBadArg, "unknown morphological operation" ); @@ -1612,9 +1638,11 @@ void cv::morphologyEx( InputArray _src, OutputArray _dst, int op, { #ifdef HAVE_OPENCL Size ksize = kernel.size(); + anchor = normalizeAnchor(anchor, ksize); + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && _src.channels() <= 4 && anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1 && - borderType == cv::BORDER_CONSTANT, + borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue(), ocl_morphologyEx(_src, _dst, op, kernel, anchor, iterations, borderType, borderValue)) #endif diff --git a/modules/imgproc/src/opencl/morph.cl b/modules/imgproc/src/opencl/morph.cl index 7df09ec61e..f78af89c9c 100644 --- a/modules/imgproc/src/opencl/morph.cl +++ b/modules/imgproc/src/opencl/morph.cl @@ -43,6 +43,8 @@ #endif #endif +#define noconvert + #if cn != 3 #define loadpix(addr) *(__global const T *)(addr) #define storepix(val, addr) *(__global T *)(addr) = val @@ -107,6 +109,11 @@ // BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii #define ELEM(i, l_edge, r_edge, elem1, elem2) (i) < (l_edge) | (i) >= (r_edge) ? (elem1) : (elem2) +#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT +#define EXTRA_PARAMS , __global const uchar * matptr, int mat_step, int mat_offset +#else +#define EXTRA_PARAMS +#endif __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, __global uchar * dstptr, int dst_step, int dst_offset, @@ -155,6 +162,20 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, PROCESS_ELEMS; int dst_index = mad24(gidy, dst_step, mad24(gidx, TSIZE, dst_offset)); + +#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT + int mat_index = mad24(gidy, mat_step, mad24(gidx, TSIZE, mat_offset)); + T value = loadpix(matptr + mat_index); + +#ifdef OP_GRADIENT + storepix(convertToT(convertToWT(res) - convertToWT(value)), dstptr + dst_index); +#elif defined OP_TOPHAT + storepix(convertToT(convertToWT(value) - convertToWT(res)), dstptr + dst_index); +#elif defined OP_BLACKHAT + storepix(convertToT(convertToWT(res) - convertToWT(value)), dstptr + dst_index); +#endif +#else // erode or dilate storepix(res, dstptr + dst_index); +#endif } } diff --git a/modules/imgproc/test/ocl/test_filters.cpp b/modules/imgproc/test/ocl/test_filters.cpp index 46d77285d7..1fe2927886 100644 --- a/modules/imgproc/test/ocl/test_filters.cpp +++ b/modules/imgproc/test/ocl/test_filters.cpp @@ -63,7 +63,7 @@ PARAM_TEST_CASE(FilterTestBase, MatType, BorderType, // border type double, // optional parameter bool, // roi or not - int) //width multiplier + int) // width multiplier { int type, borderType, ksize; Size size; @@ -244,8 +244,8 @@ OCL_TEST_P(Erode, Mat) random_roi(); Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3); - OCL_OFF(cv::erode(src_roi, dst_roi, kernel, Point(-1,-1), iterations) ); - OCL_ON(cv::erode(usrc_roi, udst_roi, kernel, Point(-1,-1), iterations) ); + OCL_OFF(cv::erode(src_roi, dst_roi, kernel, Point(-1, -1), iterations) ); + OCL_ON(cv::erode(usrc_roi, udst_roi, kernel, Point(-1, -1), iterations) ); Near(); } @@ -266,8 +266,8 @@ OCL_TEST_P(Dilate, Mat) random_roi(); Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3); - OCL_OFF(cv::dilate(src_roi, dst_roi, kernel, Point(-1,-1), iterations) ); - OCL_ON(cv::dilate(usrc_roi, udst_roi, kernel, Point(-1,-1), iterations) ); + OCL_OFF(cv::dilate(src_roi, dst_roi, kernel, Point(-1, -1), iterations) ); + OCL_ON(cv::dilate(usrc_roi, udst_roi, kernel, Point(-1, -1), iterations) ); Near(); } @@ -289,8 +289,8 @@ OCL_TEST_P(MorphologyEx, Mat) random_roi(); Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3); - OCL_OFF(cv::morphologyEx(src_roi, dst_roi, op, kernel, Point(-1,-1), iterations) ); - OCL_ON(cv::morphologyEx(usrc_roi, udst_roi, op, kernel, Point(-1,-1), iterations) ); + OCL_OFF(cv::morphologyEx(src_roi, dst_roi, op, kernel, Point(-1, -1), iterations) ); + OCL_ON(cv::morphologyEx(usrc_roi, udst_roi, op, kernel, Point(-1, -1), iterations) ); Near(); } @@ -360,8 +360,8 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine( OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4), Values(3, 5, 7), - Values(Size(0,0)),//not used - Values((BorderType)BORDER_CONSTANT),//not used + Values(Size(0, 0)), //not used + Values((BorderType)BORDER_CONSTANT), Values(1.0, 2.0, 3.0), Bool(), Values(1))); // not used @@ -369,20 +369,20 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine( OCL_INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4), Values(3, 5, 7), - Values(Size(0,0)),//not used - Values((BorderType)BORDER_CONSTANT),//not used + Values(Size(0, 0)), // not used + Values((BorderType)BORDER_CONSTANT), Values(1.0, 2.0, 3.0), Bool(), - Values(1))); //not used + Values(1))); // not used OCL_INSTANTIATE_TEST_CASE_P(Filter, MorphologyEx, Combine( - Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4), + Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), Values(3, 5, 7), - Values(Size(0, 0), Size(0, 1), Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations - Values((BorderType)BORDER_CONSTANT),// not used + Values(Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations + Values((BorderType)BORDER_CONSTANT), Values(1.0, 2.0, 3.0), Bool(), - Values(1))); //not used + Values(1))); // not used } } // namespace cvtest::ocl From 654bdde8ed5ca4fb27de61d63344055958e24ed2 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Mon, 30 Jun 2014 01:47:51 +0400 Subject: [PATCH 15/24] SSE2 optimization of cv::preCornerDetect --- modules/imgproc/src/corner.cpp | 26 +++++++++++++++++++++++++- 1 file changed, 25 insertions(+), 1 deletion(-) diff --git a/modules/imgproc/src/corner.cpp b/modules/imgproc/src/corner.cpp index eeb20fbc16..923d78b30f 100644 --- a/modules/imgproc/src/corner.cpp +++ b/modules/imgproc/src/corner.cpp @@ -608,6 +608,11 @@ void cv::preCornerDetect( InputArray _src, OutputArray _dst, int ksize, int bord factor *= 255; factor = 1./(factor * factor * factor); +#if CV_SSE2 + volatile bool haveSSE2 = cv::checkHardwareSupport(CV_CPU_SSE2); + __m128 v_factor = _mm_set1_ps((float)factor), v_m2 = _mm_set1_ps(-2.0f); +#endif + Size size = src.size(); int i, j; for( i = 0; i < size.height; i++ ) @@ -619,7 +624,26 @@ void cv::preCornerDetect( InputArray _src, OutputArray _dst, int ksize, int bord const float* d2ydata = (const float*)(D2y.data + i*D2y.step); const float* dxydata = (const float*)(Dxy.data + i*Dxy.step); - for( j = 0; j < size.width; j++ ) + j = 0; + +#if CV_SSE2 + if (haveSSE2) + { + for( ; j <= size.width - 4; j += 4 ) + { + __m128 v_dx = _mm_loadu_ps((const float *)(dxdata + j)); + __m128 v_dy = _mm_loadu_ps((const float *)(dydata + j)); + + __m128 v_s1 = _mm_mul_ps(_mm_mul_ps(v_dx, v_dx), _mm_loadu_ps((const float *)(d2ydata + j))); + __m128 v_s2 = _mm_mul_ps(_mm_mul_ps(v_dy, v_dy), _mm_loadu_ps((const float *)(d2xdata + j))); + __m128 v_s3 = _mm_mul_ps(_mm_mul_ps(v_dx, v_dy), _mm_loadu_ps((const float *)(dxydata + j))); + v_s1 = _mm_mul_ps(v_factor, _mm_add_ps(v_s1, _mm_add_ps(v_s2, _mm_mul_ps(v_s3, v_m2)))); + _mm_storeu_ps(dstdata + j, v_s1); + } + } +#endif + + for( ; j < size.width; j++ ) { float dx = dxdata[j]; float dy = dydata[j]; From 978f7eb44ae983e9689e210ad801bb852997617d Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 27 Jun 2014 14:06:32 +0400 Subject: [PATCH 16/24] added perf test for transpose inplace --- modules/core/perf/opencl/perf_arithm.cpp | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/modules/core/perf/opencl/perf_arithm.cpp b/modules/core/perf/opencl/perf_arithm.cpp index 8488c1799b..17badca765 100644 --- a/modules/core/perf/opencl/perf_arithm.cpp +++ b/modules/core/perf/opencl/perf_arithm.cpp @@ -308,6 +308,23 @@ OCL_PERF_TEST_P(TransposeFixture, Transpose, ::testing::Combine( SANITY_CHECK(dst); } +OCL_PERF_TEST_P(TransposeFixture, TransposeInplace, ::testing::Combine( + OCL_PERF_ENUM(Size(640, 640), Size(1280, 1280), Size(2160, 2160)), OCL_TEST_TYPES_134)) +{ + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + + checkDeviceMaxMemoryAllocSize(srcSize, type); + + UMat src(srcSize, type); + declare.in(src, WARMUP_RNG).out(src, WARMUP_NONE); + + OCL_TEST_CYCLE() cv::transpose(src, src); + + SANITY_CHECK_NOTHING(); +} + ///////////// Flip //////////////////////// enum From 9c8b9fc73384ca3945962a7850eb5e58837bab80 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 27 Jun 2014 12:44:32 +0400 Subject: [PATCH 17/24] cv::transpose --- modules/core/src/matrix.cpp | 4 ++-- modules/core/src/opencl/transpose.cl | 4 +++- 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index 653efe63ae..c27c961077 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -3002,8 +3002,8 @@ static bool ocl_transpose( InputArray _src, OutputArray _dst ) k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnlyNoSize(dst)); - size_t localsize[3] = { TILE_DIM, BLOCK_ROWS, 1 }; - size_t globalsize[3] = { src.cols, inplace ? src.rows : divUp(src.rows, TILE_DIM) * BLOCK_ROWS, 1 }; + size_t localsize[2] = { TILE_DIM, BLOCK_ROWS }; + size_t globalsize[2] = { src.cols, inplace ? src.rows : divUp(src.rows, TILE_DIM) * BLOCK_ROWS }; return k.run(2, globalsize, localsize, false); } diff --git a/modules/core/src/opencl/transpose.cl b/modules/core/src/opencl/transpose.cl index b5ec4b6f95..d56e499099 100644 --- a/modules/core/src/opencl/transpose.cl +++ b/modules/core/src/opencl/transpose.cl @@ -53,7 +53,7 @@ #define TSIZE ((int)sizeof(T1)*3) #endif -#define LDS_STEP TILE_DIM +#define LDS_STEP (TILE_DIM + 1) __kernel void transpose(__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) @@ -90,6 +90,7 @@ __kernel void transpose(__global const uchar * srcptr, int src_step, int src_off { int index_src = mad24(y, src_step, mad24(x, TSIZE, src_offset)); + #pragma unroll for (int i = 0; i < TILE_DIM; i += BLOCK_ROWS) if (y + i < src_rows) { @@ -103,6 +104,7 @@ __kernel void transpose(__global const uchar * srcptr, int src_step, int src_off { int index_dst = mad24(y_index, dst_step, mad24(x_index, TSIZE, dst_offset)); + #pragma unroll for (int i = 0; i < TILE_DIM; i += BLOCK_ROWS) if ((y_index + i) < src_cols) { From 54e4ef657c41a5bb794ebd2a34932cae16a8507b Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 27 Jun 2014 14:03:05 +0400 Subject: [PATCH 18/24] optimized cv::transpose inplace --- modules/core/src/matrix.cpp | 16 ++++++++++++---- modules/core/src/opencl/transpose.cl | 20 +++++++++++++------- 2 files changed, 25 insertions(+), 11 deletions(-) diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index c27c961077..f199cb2534 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -2973,8 +2973,10 @@ static inline int divUp(int a, int b) static bool ocl_transpose( InputArray _src, OutputArray _dst ) { + const ocl::Device & dev = ocl::Device::getDefault(); const int TILE_DIM = 32, BLOCK_ROWS = 8; - int type = _src.type(), cn = CV_MAT_CN(type), depth = CV_MAT_DEPTH(type); + int type = _src.type(), cn = CV_MAT_CN(type), depth = CV_MAT_DEPTH(type), + rowsPerWI = dev.isIntel() ? 4 : 1; UMat src = _src.getUMat(); _dst.create(src.cols, src.rows, type); @@ -2990,9 +2992,9 @@ static bool ocl_transpose( InputArray _src, OutputArray _dst ) } ocl::Kernel k(kernelName.c_str(), ocl::core::transpose_oclsrc, - format("-D T=%s -D T1=%s -D cn=%d -D TILE_DIM=%d -D BLOCK_ROWS=%d", + format("-D T=%s -D T1=%s -D cn=%d -D TILE_DIM=%d -D BLOCK_ROWS=%d -D rowsPerWI=%d", ocl::memopTypeToStr(type), ocl::memopTypeToStr(depth), - cn, TILE_DIM, BLOCK_ROWS)); + cn, TILE_DIM, BLOCK_ROWS, rowsPerWI)); if (k.empty()) return false; @@ -3003,7 +3005,13 @@ static bool ocl_transpose( InputArray _src, OutputArray _dst ) ocl::KernelArg::WriteOnlyNoSize(dst)); size_t localsize[2] = { TILE_DIM, BLOCK_ROWS }; - size_t globalsize[2] = { src.cols, inplace ? src.rows : divUp(src.rows, TILE_DIM) * BLOCK_ROWS }; + size_t globalsize[2] = { src.cols, inplace ? (src.rows + rowsPerWI - 1) / rowsPerWI : (divUp(src.rows, TILE_DIM) * BLOCK_ROWS) }; + + if (inplace && dev.isIntel()) + { + localsize[0] = 16; + localsize[1] = dev.maxWorkGroupSize() / localsize[0]; + } return k.run(2, globalsize, localsize, false); } diff --git a/modules/core/src/opencl/transpose.cl b/modules/core/src/opencl/transpose.cl index d56e499099..cad3616b5d 100644 --- a/modules/core/src/opencl/transpose.cl +++ b/modules/core/src/opencl/transpose.cl @@ -117,18 +117,24 @@ __kernel void transpose(__global const uchar * srcptr, int src_step, int src_off __kernel void transpose_inplace(__global uchar * srcptr, int src_step, int src_offset, int src_rows) { int x = get_global_id(0); - int y = get_global_id(1); + int y = get_global_id(1) * rowsPerWI; - if (y < src_rows && x < y) + if (x < y + rowsPerWI) { int src_index = mad24(y, src_step, mad24(x, TSIZE, src_offset)); int dst_index = mad24(x, src_step, mad24(y, TSIZE, src_offset)); + T tmp; - __global const uchar * src = srcptr + src_index; - __global uchar * dst = srcptr + dst_index; + #pragma unroll + for (int i = 0; i < rowsPerWI; ++i, ++y, src_index += src_step, dst_index += TSIZE) + if (y < src_rows && x < y) + { + __global uchar * src = srcptr + src_index; + __global uchar * dst = srcptr + dst_index; - T tmp = loadpix(dst); - storepix(loadpix(src), dst); - storepix(tmp, src); + tmp = loadpix(dst); + storepix(loadpix(src), dst); + storepix(tmp, src); + } } } From 972119bff3cae561ed9440eba1aa80c85a7a27ca Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sun, 29 Jun 2014 21:03:16 +0400 Subject: [PATCH 19/24] sse2 optimization of cv::convertScaleAbs --- modules/core/src/convert.cpp | 174 ++++++++++++++++++++++++++++++++++- 1 file changed, 172 insertions(+), 2 deletions(-) diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp index 64d24304cd..7b2684c85e 100644 --- a/modules/core/src/convert.cpp +++ b/modules/core/src/convert.cpp @@ -851,6 +851,175 @@ void cv::insertChannel(InputArray _src, InputOutputArray _dst, int coi) namespace cv { +template +struct cvtScaleAbs_SSE2 +{ + int operator () (const T *, DT *, int, WT, WT) const + { + return 0; + } +}; + +#if CV_SSE2 + +template <> +struct cvtScaleAbs_SSE2 +{ + int operator () (const uchar * src, uchar * dst, int width, + float scale, float shift) const + { + int x = 0; + + if (USE_SSE2) + { + __m128 v_scale = _mm_set1_ps(scale), v_shift = _mm_set1_ps(shift), + v_zero_f = _mm_setzero_ps(); + __m128i v_zero_i = _mm_setzero_si128(); + + for ( ; x <= width - 16; x += 16) + { + __m128i v_src = _mm_loadu_si128((const __m128i *)(src + x)); + __m128i v_src12 = _mm_unpacklo_epi8(v_src, v_zero_i), v_src_34 = _mm_unpackhi_epi8(v_src, v_zero_i); + __m128 v_dst1 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(_mm_unpacklo_epi16(v_src12, v_zero_i)), v_scale), v_shift); + v_dst1 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst1), v_dst1); + __m128 v_dst2 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(_mm_unpackhi_epi16(v_src12, v_zero_i)), v_scale), v_shift); + v_dst2 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst2), v_dst2); + __m128 v_dst3 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(_mm_unpacklo_epi16(v_src_34, v_zero_i)), v_scale), v_shift); + v_dst3 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst3), v_dst3); + __m128 v_dst4 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(_mm_unpackhi_epi16(v_src_34, v_zero_i)), v_scale), v_shift); + v_dst4 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst4), v_dst4); + + __m128i v_dst_i = _mm_packus_epi16(_mm_packs_epi32(_mm_cvtps_epi32(v_dst1), _mm_cvtps_epi32(v_dst2)), + _mm_packs_epi32(_mm_cvtps_epi32(v_dst3), _mm_cvtps_epi32(v_dst4))); + _mm_storeu_si128((__m128i *)(dst + x), v_dst_i); + } + } + + return x; + } +}; + +template <> +struct cvtScaleAbs_SSE2 +{ + int operator () (const ushort * src, uchar * dst, int width, + float scale, float shift) const + { + int x = 0; + + if (USE_SSE2) + { + __m128 v_scale = _mm_set1_ps(scale), v_shift = _mm_set1_ps(shift), + v_zero_f = _mm_setzero_ps(); + __m128i v_zero_i = _mm_setzero_si128(); + + for ( ; x <= width - 8; x += 8) + { + __m128i v_src = _mm_loadu_si128((const __m128i *)(src + x)); + __m128 v_dst1 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(_mm_unpacklo_epi16(v_src, v_zero_i)), v_scale), v_shift); + v_dst1 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst1), v_dst1); + __m128 v_dst2 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(_mm_unpackhi_epi16(v_src, v_zero_i)), v_scale), v_shift); + v_dst2 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst2), v_dst2); + + __m128i v_dst_i = _mm_packus_epi16(_mm_packs_epi32(_mm_cvtps_epi32(v_dst1), _mm_cvtps_epi32(v_dst2)), v_zero_i); + _mm_storel_epi64((__m128i *)(dst + x), v_dst_i); + } + } + + return x; + } +}; + +template <> +struct cvtScaleAbs_SSE2 +{ + int operator () (const short * src, uchar * dst, int width, + float scale, float shift) const + { + int x = 0; + + if (USE_SSE2) + { + __m128 v_scale = _mm_set1_ps(scale), v_shift = _mm_set1_ps(shift), + v_zero_f = _mm_setzero_ps(); + __m128i v_zero_i = _mm_setzero_si128(); + + for ( ; x <= width - 8; x += 8) + { + __m128i v_src = _mm_loadu_si128((const __m128i *)(src + x)); + __m128 v_dst1 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(_mm_srai_epi32(_mm_unpacklo_epi16(v_src, v_src), 16)), v_scale), v_shift); + v_dst1 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst1), v_dst1); + __m128 v_dst2 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(_mm_srai_epi32(_mm_unpackhi_epi16(v_src, v_src), 16)), v_scale), v_shift); + v_dst2 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst2), v_dst2); + + __m128i v_dst_i = _mm_packus_epi16(_mm_packs_epi32(_mm_cvtps_epi32(v_dst1), _mm_cvtps_epi32(v_dst2)), v_zero_i); + _mm_storel_epi64((__m128i *)(dst + x), v_dst_i); + } + } + + return x; + } +}; + +template <> +struct cvtScaleAbs_SSE2 +{ + int operator () (const int * src, uchar * dst, int width, + float scale, float shift) const + { + int x = 0; + + if (USE_SSE2) + { + __m128 v_scale = _mm_set1_ps(scale), v_shift = _mm_set1_ps(shift), + v_zero_f = _mm_setzero_ps(); + __m128i v_zero_i = _mm_setzero_si128(); + + for ( ; x <= width - 8; x += 4) + { + __m128i v_src = _mm_loadu_si128((const __m128i *)(src + x)); + __m128 v_dst1 = _mm_add_ps(_mm_mul_ps(_mm_cvtepi32_ps(v_src), v_scale), v_shift); + v_dst1 = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst1), v_dst1); + + __m128i v_dst_i = _mm_packus_epi16(_mm_packs_epi32(_mm_cvtps_epi32(v_dst1), v_zero_i), v_zero_i); + _mm_storel_epi64((__m128i *)(dst + x), v_dst_i); + } + } + + return x; + } +}; + +template <> +struct cvtScaleAbs_SSE2 +{ + int operator () (const float * src, uchar * dst, int width, + float scale, float shift) const + { + int x = 0; + + if (USE_SSE2) + { + __m128 v_scale = _mm_set1_ps(scale), v_shift = _mm_set1_ps(shift), + v_zero_f = _mm_setzero_ps(); + __m128i v_zero_i = _mm_setzero_si128(); + + for ( ; x <= width - 8; x += 4) + { + __m128 v_dst = _mm_add_ps(_mm_mul_ps(_mm_loadu_ps(src + x), v_scale), v_shift); + v_dst = _mm_max_ps(_mm_sub_ps(v_zero_f, v_dst), v_dst); + + __m128i v_dst_i = _mm_packs_epi32(_mm_cvtps_epi32(v_dst), v_zero_i); + _mm_storel_epi64((__m128i *)(dst + x), _mm_packus_epi16(v_dst_i, v_zero_i)); + } + } + + return x; + } +}; + +#endif + template static void cvtScaleAbs_( const T* src, size_t sstep, DT* dst, size_t dstep, Size size, @@ -858,10 +1027,12 @@ cvtScaleAbs_( const T* src, size_t sstep, { sstep /= sizeof(src[0]); dstep /= sizeof(dst[0]); + cvtScaleAbs_SSE2 vop; for( ; size.height--; src += sstep, dst += dstep ) { - int x = 0; + int x = vop(src, dst, size.width, scale, shift); + #if CV_ENABLE_UNROLLED for( ; x <= size.width - 4; x += 4 ) { @@ -879,7 +1050,6 @@ cvtScaleAbs_( const T* src, size_t sstep, } } - template static void cvtScale_( const T* src, size_t sstep, DT* dst, size_t dstep, Size size, From 157257b0723e7bfb9f7d929402772aef1966a752 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Tue, 1 Jul 2014 14:24:14 +0400 Subject: [PATCH 20/24] ocl: don't use OpenCL CPU device in default setup --- modules/core/src/ocl.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 24ca6ee4d3..0b196d9c52 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -2228,7 +2228,8 @@ static cl_device_id selectOpenCLDevice() if (!isID) { deviceTypes.push_back("GPU"); - deviceTypes.push_back("CPU"); + if (configuration) + deviceTypes.push_back("CPU"); } else deviceTypes.push_back("ALL"); From 46e2216165ada8ed3f038470c2b1bec687ad4845 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 2 Jul 2014 19:32:10 +0400 Subject: [PATCH 21/24] fixed cv::warpPerspective --- modules/imgproc/src/imgwarp.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/imgproc/src/imgwarp.cpp b/modules/imgproc/src/imgwarp.cpp index b69d75d6d4..e07641f6a1 100644 --- a/modules/imgproc/src/imgwarp.cpp +++ b/modules/imgproc/src/imgwarp.cpp @@ -4188,7 +4188,7 @@ static bool ocl_warpTransform(InputArray _src, OutputArray _dst, InputArray _M0, const char * const kernelName = op_type == OCL_OP_AFFINE ? "warpAffine" : "warpPerspective"; int scalarcn = cn == 3 ? 4 : cn; - bool is32f = !dev.isAMD() && (interpolation == INTER_CUBIC || interpolation == INTER_LINEAR); + bool is32f = !dev.isAMD() && (interpolation == INTER_CUBIC || interpolation == INTER_LINEAR) && op_type == OCL_OP_AFFINE; int wdepth = interpolation == INTER_NEAREST ? depth : std::max(is32f ? CV_32F : CV_32S, depth); int sctype = CV_MAKETYPE(wdepth, scalarcn); From 93f8581b6954233284dc9983a27ff808afac70d4 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Tue, 1 Jul 2014 14:31:25 +0400 Subject: [PATCH 22/24] fixed warnings --- modules/highgui/src/cap_qtkit.mm | 21 ++++++++++----------- 1 file changed, 10 insertions(+), 11 deletions(-) diff --git a/modules/highgui/src/cap_qtkit.mm b/modules/highgui/src/cap_qtkit.mm index b0800a1302..461bc1f33c 100644 --- a/modules/highgui/src/cap_qtkit.mm +++ b/modules/highgui/src/cap_qtkit.mm @@ -448,13 +448,12 @@ double CvCaptureCAM::getProperty(int property_id){ QTFormatDescription* format = [[connections objectAtIndex:0] formatDescription]; NSSize s1 = [[format attributeForKey:QTFormatDescriptionVideoCleanApertureDisplaySizeAttribute] sizeValue]; - int width=s1.width, height=s1.height; switch (property_id) { case CV_CAP_PROP_FRAME_WIDTH: - retval = width; + retval = s1.width; break; case CV_CAP_PROP_FRAME_HEIGHT: - retval = height; + retval = s1.height; break; default: retval = 0; @@ -1011,22 +1010,22 @@ bool CvVideoWriter_QT::writeFrame(const IplImage* image) { cvCvtColor(image, argbimage, CV_BGR2BGRA); - unsigned char* imagedata = (unsigned char*)argbimage->imageData; + unsigned char* imagedata_ = (unsigned char*)argbimage->imageData; //BGRA --> ARGB for (int j = 0; j < argbimage->height; j++) { int rowstart = argbimage->widthStep * j; for (int i = rowstart; i < rowstart+argbimage->widthStep; i+=4) { - unsigned char temp = imagedata[i]; - imagedata[i] = 255; - imagedata[i+3] = temp; - temp = imagedata[i+2]; - imagedata[i+2] = imagedata[i+1]; - imagedata[i+1] = temp; + unsigned char temp = imagedata_[i]; + imagedata_[i] = 255; + imagedata_[i+3] = temp; + temp = imagedata_[i+2]; + imagedata_[i+2] = imagedata_[i+1]; + imagedata_[i+1] = temp; } } - NSBitmapImageRep* imageRep = [[NSBitmapImageRep alloc] initWithBitmapDataPlanes:&imagedata + NSBitmapImageRep* imageRep = [[NSBitmapImageRep alloc] initWithBitmapDataPlanes:&imagedata_ pixelsWide:movieSize.width pixelsHigh:movieSize.height bitsPerSample:8 From c5d3c08256e06fab298e153f3bde89793eceb221 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Thu, 3 Jul 2014 12:18:19 +0400 Subject: [PATCH 23/24] ocl: add try-catch for OpenCL device getter --- modules/core/src/ocl.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 0b196d9c52..8ffd1d200a 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -1416,7 +1416,16 @@ bool useOpenCL() { CoreTLSData* data = coreTlsData.get(); if( data->useOpenCL < 0 ) - data->useOpenCL = (int)haveOpenCL() && Device::getDefault().ptr() != NULL; + { + try + { + data->useOpenCL = (int)haveOpenCL() && Device::getDefault().ptr() != NULL; + } + catch (...) + { + data->useOpenCL = 0; + } + } return data->useOpenCL > 0; } From 2fe07abf83360467a318406dbecb6b32d3b2522f Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 3 Jul 2014 13:45:55 +0400 Subject: [PATCH 24/24] use vectors for devices of all vendors --- modules/core/src/ocl.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 0b196d9c52..43e1ab2c95 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -4428,11 +4428,13 @@ int predictOptimalVectorWidth(InputArray src1, InputArray src2, InputArray src3, d.preferredVectorWidthShort(), d.preferredVectorWidthShort(), d.preferredVectorWidthInt(), d.preferredVectorWidthFloat(), d.preferredVectorWidthDouble(), -1 }, kercn = vectorWidths[depth]; - if (d.isIntel()) + + // if the device says don't use vectors + if (vectorWidths[0] == 1) { // it's heuristic - int vectorWidthsIntel[] = { 16, 16, 8, 8, 1, 1, 1, -1 }; - kercn = vectorWidthsIntel[depth]; + int vectorWidthsOthers[] = { 16, 16, 8, 8, 1, 1, 1, -1 }; + kercn = vectorWidthsOthers[depth]; } if (ssize.width * cn < kercn || kercn <= 0)