From 29ba7730ed6664ff9d52889eba376e97719df035 Mon Sep 17 00:00:00 2001 From: Vladimir Bystricky Date: Mon, 9 Dec 2013 10:45:23 +0400 Subject: [PATCH 1/5] Add ocl version of filter2D to the module img_proc --- modules/imgproc/src/filter.cpp | 186 ++++++++++ modules/imgproc/src/opencl/filter2D.cl | 377 +++++++++++++++++++++ modules/imgproc/test/ocl/test_filter2d.cpp | 138 ++++++++ 3 files changed, 701 insertions(+) create mode 100644 modules/imgproc/src/opencl/filter2D.cl create mode 100644 modules/imgproc/test/ocl/test_filter2d.cpp diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index fbf1de2479..80a9584165 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#include "opencl_kernels.hpp" /****************************************************************************************\ Base Image Filter @@ -3115,6 +3116,187 @@ template struct Filter2D : public BaseFi } +namespace cv +{ + +#define DIVUP(total, grain) (((total) + (grain) - 1) / (grain)) +#define ROUNDUP(sz, n) ((sz) + (n) - 1 - (((sz) + (n) - 1) % (n))) + +// prepare kernel: transpose and make double rows (+align). Returns size of aligned row +// Samples: +// a b c +// Input: d e f +// g h i +// Output, last two zeros is the alignment: +// a d g a d g 0 0 +// b e h b e h 0 0 +// c f i c f i 0 0 +template +static int _prepareKernelFilter2D(std::vector& data, const Mat &kernel) +{ + Mat _kernel; kernel.convertTo(_kernel, DataDepth::value); + int size_y_aligned = ROUNDUP(kernel.rows * 2, 4); + data.clear(); data.resize(size_y_aligned * kernel.cols, 0); + for (int x = 0; x < kernel.cols; x++) + { + for (int y = 0; y < kernel.rows; y++) + { + data[x * size_y_aligned + y] = _kernel.at(y, x); + data[x * size_y_aligned + y + kernel.rows] = _kernel.at(y, x); + } + } + return size_y_aligned; +} + +bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, + InputArray _kernel, Point anchor, + double delta, int borderType ) +{ + if (abs(delta) > FLT_MIN) + return false; + + int type = _src.type(); + int cn = CV_MAT_CN(type); + if ((1 != cn) && (2 != cn) && (4 != cn)) + return false;//TODO + + int sdepth = CV_MAT_DEPTH(type); + Size ksize = _kernel.size(); + if( anchor.x < 0 ) + anchor.x = ksize.width / 2; + if( anchor.y < 0 ) + anchor.y = ksize.height / 2; + if( ddepth < 0 ) + ddepth = sdepth; + else if (ddepth != sdepth) + return false; + + bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0; + bool useDouble = (CV_64F == sdepth); + const cv::ocl::Device &device = cv::ocl::Device::getDefault(); + int doubleFPConfig = device.doubleFPConfig(); + if (useDouble && (0 == doubleFPConfig)) + return false; + + const char* btype = NULL; + switch (borderType & ~BORDER_ISOLATED) + { + case BORDER_CONSTANT: + btype = "BORDER_CONSTANT"; + break; + case BORDER_REPLICATE: + btype = "BORDER_REPLICATE"; + break; + case BORDER_REFLECT: + btype = "BORDER_REFLECT"; + break; + case BORDER_WRAP: + return false; + case BORDER_REFLECT101: + btype = "BORDER_REFLECT_101"; + break; + } + + cv::Mat kernelMat = _kernel.getMat(); + std::vector kernelMatDataFloat; + std::vector kernelMatDataDouble; + int kernel_size_y2_aligned = useDouble ? + _prepareKernelFilter2D(kernelMatDataDouble, kernelMat) + : _prepareKernelFilter2D(kernelMatDataFloat, kernelMat); + + + cv::Size sz = _src.size(); + size_t globalsize[2] = {sz.width, sz.height}; + size_t localsize[2] = {0, 1}; + + ocl::Kernel kernel; + + size_t maxWorkItemSizes[32]; device.maxWorkItemSizes(maxWorkItemSizes); + size_t tryWorkItems = maxWorkItemSizes[0]; + for (;;) + { + size_t BLOCK_SIZE = tryWorkItems; + while (BLOCK_SIZE > 32 && BLOCK_SIZE >= (size_t)ksize.width * 2 && BLOCK_SIZE > (size_t)sz.width * 2) + BLOCK_SIZE /= 2; +#if 1 // TODO Mode with several blocks requires a much more VGPRs, so this optimization is not actual for the current devices + size_t BLOCK_SIZE_Y = 1; +#else + size_t BLOCK_SIZE_Y = 8; // TODO Check heuristic value on devices + while (BLOCK_SIZE_Y < BLOCK_SIZE / 8 && BLOCK_SIZE_Y * src.clCxt->getDeviceInfo().maxComputeUnits * 32 < (size_t)src.rows) + BLOCK_SIZE_Y *= 2; +#endif + + if ((size_t)ksize.width > BLOCK_SIZE) + return false; + + int requiredTop = anchor.y; + int requiredLeft = (int)BLOCK_SIZE; // not this: anchor.x; + int requiredBottom = ksize.height - 1 - anchor.y; + int requiredRight = (int)BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x; + int h = sz.height; + int w = sz.width; + bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight; + + if ((w < ksize.width) || (h < ksize.height)) + return false; + + char build_options[1024]; + sprintf(build_options, "-D LOCAL_SIZE=%d -D BLOCK_SIZE_Y=%d -D DATA_DEPTH=%d -D DATA_CHAN=%d -D USE_DOUBLE=%d " + "-D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d -D KERNEL_SIZE_Y2_ALIGNED=%d " + "-D %s -D %s -D %s", + (int)BLOCK_SIZE, (int)BLOCK_SIZE_Y, + sdepth, cn, useDouble ? 1 : 0, + anchor.x, anchor.y, ksize.width, ksize.height, kernel_size_y2_aligned, + btype, + extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION", + isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED"); + + localsize[0] = BLOCK_SIZE; + globalsize[0] = DIVUP(sz.width, BLOCK_SIZE - (ksize.width - 1)) * BLOCK_SIZE; + globalsize[1] = DIVUP(sz.height, BLOCK_SIZE_Y); + + cv::String errmsg; + if (!kernel.create("filter2D", cv::ocl::imgproc::filter2D_oclsrc, build_options)) + return false; + size_t kernelWorkGroupSize = kernel.workGroupSize(); + if (localsize[0] <= kernelWorkGroupSize) + break; + if (BLOCK_SIZE < kernelWorkGroupSize) + return false; + tryWorkItems = kernelWorkGroupSize; + } + + _dst.create(sz, CV_MAKETYPE(ddepth, cn)); + UMat dst = _dst.getUMat(); + UMat src = _src.getUMat(); + + int idxArg = 0; + idxArg = kernel.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(src)); + idxArg = kernel.set(idxArg, ocl::KernelArg::WriteOnly(dst)); + float borderValue[4] = {0, 0, 0, 0}; + double borderValueDouble[4] = {0, 0, 0, 0}; + if ((borderType & ~BORDER_ISOLATED) == BORDER_CONSTANT) + { + int cnocl = (3 == cn) ? 4 : cn; + if (useDouble) + idxArg = kernel.set(idxArg, (void *)&borderValueDouble[0], sizeof(double) * cnocl); + else + idxArg = kernel.set(idxArg, (void *)&borderValue[0], sizeof(float) * cnocl); + } + if (useDouble) + { + UMat kernalDataUMat(kernelMatDataDouble, true); + idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(kernalDataUMat)); + } + else + { + UMat kernalDataUMat(kernelMatDataFloat, true); + idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(kernalDataUMat)); + } + return kernel.run(2, globalsize, localsize, true); +} +} + cv::Ptr cv::getLinearFilter(int srcType, int dstType, InputArray filter_kernel, Point anchor, double delta, int bits) @@ -3230,6 +3412,10 @@ void cv::filter2D( InputArray _src, OutputArray _dst, int ddepth, InputArray _kernel, Point anchor, double delta, int borderType ) { + bool use_opencl = ocl::useOpenCL() && _dst.isUMat(); + if( use_opencl && ocl_filter2D(_src, _dst, ddepth, _kernel, anchor, delta, borderType)) + return; + Mat src = _src.getMat(), kernel = _kernel.getMat(); if( ddepth < 0 ) diff --git a/modules/imgproc/src/opencl/filter2D.cl b/modules/imgproc/src/opencl/filter2D.cl new file mode 100644 index 0000000000..1225be93fc --- /dev/null +++ b/modules/imgproc/src/opencl/filter2D.cl @@ -0,0 +1,377 @@ +/*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-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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 BORDER_REPLICATE +//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh +#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i)) +#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr)) +#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (t_edge) :(i)) +#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (b_edge)-1 :(addr)) +#endif + +#ifdef BORDER_REFLECT +//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb +#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i)) +#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)) +#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i)-1 : (i)) +#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr)) +#endif + +#ifdef BORDER_REFLECT_101 +//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba +#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i)) +#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)) +#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i)) +#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr)) +#endif + +//blur function does not support BORDER_WRAP +#ifdef BORDER_WRAP +//BORDER_WRAP: cdefgh|abcdefgh|abcdefg +#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i)) +#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr)) +#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? (i)+(b_edge) : (i)) +#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? (i)-(b_edge) : (addr)) +#endif + +#ifdef EXTRA_EXTRAPOLATION // border > src image size +#ifdef BORDER_CONSTANT +// None +#elif defined BORDER_REPLICATE +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + x = max(min(x, maxX - 1), minX); \ + y = max(min(y, maxY - 1), minY); \ + } +#elif defined BORDER_WRAP +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + if (x < minX) \ + x -= ((x - maxX + 1) / maxX) * maxX; \ + if (x >= maxX) \ + x %= maxX; \ + if (y < minY) \ + y -= ((y - maxY + 1) / maxY) * maxY; \ + if (y >= maxY) \ + y %= maxY; \ + } +#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) +#define EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, delta) \ + { \ + if (maxX - minX == 1) \ + x = minX; \ + else \ + do \ + { \ + if (x < minX) \ + x = -(x - minX) - 1 + delta; \ + else \ + x = maxX - 1 - (x - maxX) - delta; \ + } \ + while (x >= maxX || x < minX); \ + \ + if (maxY - minY == 1) \ + y = minY; \ + else \ + do \ + { \ + if (y < minY) \ + y = -(y - minY) - 1 + delta; \ + else \ + y = maxY - 1 - (y - maxY) - delta; \ + } \ + while (y >= maxY || y < minY); \ + } +#ifdef BORDER_REFLECT +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 0) +#elif defined(BORDER_REFLECT_101) +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 1) +#endif +#else +#error No extrapolation method +#endif +#else +#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \ + { \ + int _row = y - minY, _col = x - minX; \ + _row = ADDR_H(_row, 0, maxY - minY); \ + _row = ADDR_B(_row, maxY - minY, _row); \ + y = _row + minY; \ + \ + _col = ADDR_L(_col, 0, maxX - minX); \ + _col = ADDR_R(_col, maxX - minX, _col); \ + x = _col + minX; \ + } +#endif + +#if USE_DOUBLE +#ifdef cl_amd_fp64 +#pragma OPENCL EXTENSION cl_amd_fp64:enable +#elif defined (cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#endif +#define FPTYPE double +#define CONVERT_TO_FPTYPE CAT(convert_double, VEC_SIZE) +#else +#define FPTYPE float +#define CONVERT_TO_FPTYPE CAT(convert_float, VEC_SIZE) +#endif + +#if DATA_DEPTH == 0 +#define BASE_TYPE uchar +#elif DATA_DEPTH == 1 +#define BASE_TYPE char +#elif DATA_DEPTH == 2 +#define BASE_TYPE ushort +#elif DATA_DEPTH == 3 +#define BASE_TYPE short +#elif DATA_DEPTH == 4 +#define BASE_TYPE int +#elif DATA_DEPTH == 5 +#define BASE_TYPE float +#elif DATA_DEPTH == 6 +#define BASE_TYPE double +#else +#error data_depth +#endif + +#define __CAT(x, y) x##y +#define CAT(x, y) __CAT(x, y) + +#define uchar1 uchar +#define char1 char +#define ushort1 ushort +#define short1 short +#define int1 int +#define float1 float +#define double1 double + +#define convert_uchar1_sat_rte convert_uchar_sat_rte +#define convert_char1_sat_rte convert_char_sat_rte +#define convert_ushort1_sat_rte convert_ushort_sat_rte +#define convert_short1_sat_rte convert_short_sat_rte +#define convert_int1_sat_rte convert_int_sat_rte +#define convert_float1 +#define convert_double1 + +#if DATA_DEPTH == 5 || DATA_DEPTH == 6 +#define CONVERT_TO_TYPE CAT(CAT(convert_, BASE_TYPE), VEC_SIZE) +#else +#define CONVERT_TO_TYPE CAT(CAT(CAT(convert_, BASE_TYPE), VEC_SIZE), _sat_rte) +#endif + +#define VEC_SIZE DATA_CHAN + +#define VEC_TYPE CAT(BASE_TYPE, VEC_SIZE) +#define TYPE VEC_TYPE + +#define SCALAR_TYPE CAT(FPTYPE, VEC_SIZE) + +#define INTERMEDIATE_TYPE CAT(FPTYPE, VEC_SIZE) + +struct RectCoords +{ + int x1, y1, x2, y2; +}; + +//#define DEBUG +#ifdef DEBUG +#define DEBUG_ONLY(x) x +#define ASSERT(condition) do { if (!(condition)) { printf("BUG in boxFilter kernel (global=%d,%d): " #condition "\n", get_global_id(0), get_global_id(1)); } } while (0) +#else +#define DEBUG_ONLY(x) (void)0 +#define ASSERT(condition) (void)0 +#endif + + +inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, int srcstep, int srcoffset, const struct RectCoords srcCoords +#ifdef BORDER_CONSTANT + , SCALAR_TYPE borderValue +#endif + ) +{ +#ifdef BORDER_ISOLATED + if(pos.x >= srcCoords.x1 && pos.y >= srcCoords.y1 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) +#else + if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) +#endif + { + //__global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes); + __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + srcoffset + pos.x * sizeof(TYPE)); + return CONVERT_TO_FPTYPE(*ptr); + } + else + { +#ifdef BORDER_CONSTANT + return borderValue; +#else + int selected_col = pos.x; + int selected_row = pos.y; + + EXTRAPOLATE(selected_col, selected_row, +#ifdef BORDER_ISOLATED + srcCoords.x1, srcCoords.y1, +#else + 0, 0, +#endif + srcCoords.x2, srcCoords.y2 + ); + + // debug border mapping + //printf("pos=%d,%d --> %d, %d\n", pos.x, pos.y, selected_col, selected_row); + + pos = (int2)(selected_col, selected_row); + if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) + { + //__global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes); + __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + srcoffset + pos.x * sizeof(TYPE)); + return CONVERT_TO_FPTYPE(*ptr); + } + else + { + // for debug only + DEBUG_ONLY(printf("BUG in boxFilter kernel\n")); + return (FPTYPE)(0.0f); + } +#endif + } +} + +// INPUT PARAMETER: BLOCK_SIZE_Y (via defines) + +__kernel +__attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1))) +void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, + __global uchar* dstptr, int dststep, int dstoffset, + int rows, int cols, +#ifdef BORDER_CONSTANT + SCALAR_TYPE borderValue, +#endif + __constant FPTYPE* kernelData // transposed: [KERNEL_SIZE_X][KERNEL_SIZE_Y2_ALIGNED] + ) +{ + const struct RectCoords srcCoords = {0, 0, cols, rows}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY + const struct RectCoords dstCoords = {0, 0, cols, rows}; + + const int local_id = get_local_id(0); + const int x = local_id + (LOCAL_SIZE - (KERNEL_SIZE_X - 1)) * get_group_id(0) - ANCHOR_X; + const int y = get_global_id(1) * BLOCK_SIZE_Y; + + INTERMEDIATE_TYPE data[KERNEL_SIZE_Y]; + __local INTERMEDIATE_TYPE sumOfCols[LOCAL_SIZE]; + + int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y); + + int2 pos = (int2)(dstCoords.x1 + x, dstCoords.y1 + y); + __global TYPE* dstPtr = (__global TYPE*)((__global char*)dstptr + pos.y * dststep + dstoffset + pos.x * sizeof(TYPE)); // Pointer can be out of bounds! + bool writeResult = (local_id >= ANCHOR_X && local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X) && + pos.x >= dstCoords.x1 && pos.x < dstCoords.x2); + +#if BLOCK_SIZE_Y > 1 + bool readAllpixels = true; + int sy_index = 0; // current index in data[] array + + dstCoords.y2 = min(dstCoords.y2, pos.y + BLOCK_SIZE_Y); + for (; + pos.y < dstCoords.y2; + pos.y++, + dstPtr = (__global TYPE*)((__global char*)dstptr + dststep)) +#endif + { + ASSERT(pos.y < dstCoords.y2); + + for ( +#if BLOCK_SIZE_Y > 1 + int sy = readAllpixels ? 0 : -1; sy < (readAllpixels ? KERNEL_SIZE_Y : 0); +#else + int sy = 0, sy_index = 0; sy < KERNEL_SIZE_Y; +#endif + sy++, srcPos.y++) + { + data[sy + sy_index] = readSrcPixel(srcPos, srcptr, srcstep, srcoffset, srcCoords +#ifdef BORDER_CONSTANT + , borderValue +#endif + ); + } + + INTERMEDIATE_TYPE total_sum = 0; + for (int sx = 0; sx < KERNEL_SIZE_X; sx++) + { + { + __constant FPTYPE* k = &kernelData[KERNEL_SIZE_Y2_ALIGNED * sx +#if BLOCK_SIZE_Y > 1 + + KERNEL_SIZE_Y - sy_index +#endif + ]; + INTERMEDIATE_TYPE tmp_sum = 0; + for (int sy = 0; sy < KERNEL_SIZE_Y; sy++) + { + tmp_sum += data[sy] * k[sy]; + } + + sumOfCols[local_id] = tmp_sum; + barrier(CLK_LOCAL_MEM_FENCE); + } + + int id = local_id + sx - ANCHOR_X; + if (id >= 0 && id < LOCAL_SIZE) + total_sum += sumOfCols[id]; + + barrier(CLK_LOCAL_MEM_FENCE); + } + + if (writeResult) + { + ASSERT(pos.y >= dstCoords.y1 && pos.y < dstCoords.y2); + *dstPtr = CONVERT_TO_TYPE(total_sum); + } + +#if BLOCK_SIZE_Y > 1 + readAllpixels = false; +#if BLOCK_SIZE_Y > KERNEL_SIZE_Y + sy_index = (sy_index + 1 <= KERNEL_SIZE_Y) ? sy_index + 1 : 1; +#else + sy_index++; +#endif +#endif // BLOCK_SIZE_Y == 1 + } +} diff --git a/modules/imgproc/test/ocl/test_filter2d.cpp b/modules/imgproc/test/ocl/test_filter2d.cpp new file mode 100644 index 0000000000..a3c3a23609 --- /dev/null +++ b/modules/imgproc/test/ocl/test_filter2d.cpp @@ -0,0 +1,138 @@ +/*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. +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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*/ + +#include "test_precomp.hpp" +#include "opencv2/ts/ocl_test.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +enum +{ + noType = -1, +}; + + +///////////////////////////////////////////////////////////////////////////////////////////////// +// Filter2D +PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool) +{ + static const int kernelMinSize = 1; + static const int kernelMaxSize = 10; + + int type; + Size ksize; + Size dsize; + Point anchor; + int borderType; + bool useRoi; + Mat kernel; + + TEST_DECLARE_INPUT_PARAMETER(src) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1)); + borderType = GET_PARAM(2); + useRoi = GET_PARAM(3); + } + + void random_roi() + { + dsize = randomSize(1, MAX_VALUE); + + Size ksize = randomSize(kernelMinSize, kernelMaxSize); + kernel = randomMat(ksize, CV_MAKE_TYPE(((CV_64F == CV_MAT_DEPTH(type)) ? CV_64F : CV_32F), 1), -MAX_VALUE, MAX_VALUE); + + Size roiSize = randomSize(1, MAX_VALUE); + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, dsize, dstBorder, type, -MAX_VALUE, MAX_VALUE); + + anchor.x = anchor.y = -1; + + UMAT_UPLOAD_INPUT_PARAMETER(src) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + } + + void Near(double threshold = 0.0) + { + EXPECT_MAT_NEAR(dst, udst, threshold); + EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold); + } +}; + +OCL_TEST_P(Filter2D, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + OCL_OFF(cv::filter2D(src_roi, dst_roi, -1, kernel, anchor, 0.0, borderType)); + OCL_ON(cv::filter2D(usrc_roi, udst_roi, -1, kernel, anchor, 0.0, borderType)); + + Near(1.0); + } +} + + +OCL_INSTANTIATE_TEST_CASE_P(ImageProc, Filter2D, + Combine( + Values(CV_8U, CV_16U, CV_16S, CV_32F, CV_64F), + Values(1, 2, 4), + Values((BorderType)BORDER_CONSTANT, + (BorderType)BORDER_REPLICATE, + (BorderType)BORDER_REFLECT, + (BorderType)BORDER_REFLECT_101), + Bool()) + ); + + +} } // namespace cvtest::ocl + +#endif // HAVE_OPENCL From 179a43ef63745d5d4d58734324bf54202ccb2158 Mon Sep 17 00:00:00 2001 From: Vladimir Bystricky Date: Mon, 9 Dec 2013 11:10:48 +0400 Subject: [PATCH 2/5] Fix build errors on Linux system --- modules/imgproc/src/filter.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index 80a9584165..832b0d0381 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3148,7 +3148,7 @@ static int _prepareKernelFilter2D(std::vector& data, const Mat &kernel) return size_y_aligned; } -bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, +static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, InputArray _kernel, Point anchor, double delta, int borderType ) { From 8d8f5665f1a015e761811875df9937942e3effaf Mon Sep 17 00:00:00 2001 From: Vladimir Bystricky Date: Mon, 9 Dec 2013 11:31:55 +0400 Subject: [PATCH 3/5] Fix build errors on Linux system --- modules/imgproc/test/ocl/test_filter2d.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/modules/imgproc/test/ocl/test_filter2d.cpp b/modules/imgproc/test/ocl/test_filter2d.cpp index a3c3a23609..ca1450952c 100644 --- a/modules/imgproc/test/ocl/test_filter2d.cpp +++ b/modules/imgproc/test/ocl/test_filter2d.cpp @@ -63,7 +63,6 @@ PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool) static const int kernelMaxSize = 10; int type; - Size ksize; Size dsize; Point anchor; int borderType; From e7227d3e4b4b0e816b17a57b7ade9e39030c18dc Mon Sep 17 00:00:00 2001 From: Vladimir Bystricky Date: Wed, 11 Dec 2013 14:26:33 +0400 Subject: [PATCH 4/5] Fix problems with border extrapolation in kernel. Add Isolated/Nonisolated borders. --- modules/imgproc/src/filter.cpp | 27 +++++++++++++++--- modules/imgproc/src/opencl/filter2D.cl | 32 ++++++++++------------ modules/imgproc/test/ocl/test_filter2d.cpp | 27 +++++++++--------- 3 files changed, 51 insertions(+), 35 deletions(-) diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index 832b0d0381..d548168491 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3210,6 +3210,13 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, size_t localsize[2] = {0, 1}; ocl::Kernel kernel; + UMat src; Size wholeSize; + if (!isIsolatedBorder) + { + src = _src.getUMat(); + Point ofs; + src.locateROI(wholeSize, ofs); + } size_t maxWorkItemSizes[32]; device.maxWorkItemSizes(maxWorkItemSizes); size_t tryWorkItems = maxWorkItemSizes[0]; @@ -3233,8 +3240,8 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, int requiredLeft = (int)BLOCK_SIZE; // not this: anchor.x; int requiredBottom = ksize.height - 1 - anchor.y; int requiredRight = (int)BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x; - int h = sz.height; - int w = sz.width; + int h = isIsolatedBorder ? sz.height : wholeSize.height; + int w = isIsolatedBorder ? sz.width : wholeSize.width; bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight; if ((w < ksize.width) || (h < ksize.height)) @@ -3268,10 +3275,22 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, _dst.create(sz, CV_MAKETYPE(ddepth, cn)); UMat dst = _dst.getUMat(); - UMat src = _src.getUMat(); + if (src.empty()) + src = _src.getUMat(); int idxArg = 0; - idxArg = kernel.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(src)); + idxArg = kernel.set(idxArg, ocl::KernelArg::PtrReadOnly(src)); + idxArg = kernel.set(idxArg, (int)src.step); + + int srcOffsetX = (int)((src.offset % src.step) / src.elemSize()); + int srcOffsetY = (int)(src.offset / src.step); + int srcEndX = (isIsolatedBorder ? (srcOffsetX + sz.width) : wholeSize.width); + int srcEndY = (isIsolatedBorder ? (srcOffsetY + sz.height) : wholeSize.height); + idxArg = kernel.set(idxArg, srcOffsetX); + idxArg = kernel.set(idxArg, srcOffsetY); + idxArg = kernel.set(idxArg, srcEndX); + idxArg = kernel.set(idxArg, srcEndY); + idxArg = kernel.set(idxArg, ocl::KernelArg::WriteOnly(dst)); float borderValue[4] = {0, 0, 0, 0}; double borderValueDouble[4] = {0, 0, 0, 0}; diff --git a/modules/imgproc/src/opencl/filter2D.cl b/modules/imgproc/src/opencl/filter2D.cl index 1225be93fc..d360714971 100644 --- a/modules/imgproc/src/opencl/filter2D.cl +++ b/modules/imgproc/src/opencl/filter2D.cl @@ -102,7 +102,7 @@ do \ { \ if (x < minX) \ - x = -(x - minX) - 1 + delta; \ + x = minX - (x - minX) - 1 + delta; \ else \ x = maxX - 1 - (x - maxX) - delta; \ } \ @@ -114,7 +114,7 @@ do \ { \ if (y < minY) \ - y = -(y - minY) - 1 + delta; \ + y = minY - (y - minY) - 1 + delta; \ else \ y = maxY - 1 - (y - maxY) - delta; \ } \ @@ -222,7 +222,7 @@ struct RectCoords #endif -inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, int srcstep, int srcoffset, const struct RectCoords srcCoords +inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, int srcstep, const struct RectCoords srcCoords #ifdef BORDER_CONSTANT , SCALAR_TYPE borderValue #endif @@ -235,7 +235,7 @@ inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, in #endif { //__global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes); - __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + srcoffset + pos.x * sizeof(TYPE)); + __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + pos.x * sizeof(TYPE)); return CONVERT_TO_FPTYPE(*ptr); } else @@ -262,7 +262,7 @@ inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, in if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2) { //__global TYPE* ptr = (__global TYPE*)((__global char*)src + pos.x * sizeof(TYPE) + pos.y * srcStepBytes); - __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + srcoffset + pos.x * sizeof(TYPE)); + __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + pos.x * sizeof(TYPE)); return CONVERT_TO_FPTYPE(*ptr); } else @@ -279,8 +279,8 @@ inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, in __kernel __attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1))) -void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, - __global uchar* dstptr, int dststep, int dstoffset, +void filter2D(__global const uchar* srcptr, int srcstep, int srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY, + __global uchar* dstptr, int dststep, int dstoffset, int rows, int cols, #ifdef BORDER_CONSTANT SCALAR_TYPE borderValue, @@ -288,8 +288,7 @@ void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, __constant FPTYPE* kernelData // transposed: [KERNEL_SIZE_X][KERNEL_SIZE_Y2_ALIGNED] ) { - const struct RectCoords srcCoords = {0, 0, cols, rows}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY - const struct RectCoords dstCoords = {0, 0, cols, rows}; + const struct RectCoords srcCoords = {srcOffsetX, srcOffsetY, srcEndX, srcEndY}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY const int local_id = get_local_id(0); const int x = local_id + (LOCAL_SIZE - (KERNEL_SIZE_X - 1)) * get_group_id(0) - ANCHOR_X; @@ -300,23 +299,23 @@ void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y); - int2 pos = (int2)(dstCoords.x1 + x, dstCoords.y1 + y); + int2 pos = (int2)(x, y); __global TYPE* dstPtr = (__global TYPE*)((__global char*)dstptr + pos.y * dststep + dstoffset + pos.x * sizeof(TYPE)); // Pointer can be out of bounds! - bool writeResult = (local_id >= ANCHOR_X && local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X) && - pos.x >= dstCoords.x1 && pos.x < dstCoords.x2); + bool writeResult = ((local_id >= ANCHOR_X) && (local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X)) && + (pos.x >= 0) && (pos.x < cols)); #if BLOCK_SIZE_Y > 1 bool readAllpixels = true; int sy_index = 0; // current index in data[] array - dstCoords.y2 = min(dstCoords.y2, pos.y + BLOCK_SIZE_Y); + dstRowsMax = min(rows, pos.y + BLOCK_SIZE_Y); for (; - pos.y < dstCoords.y2; + pos.y < dstRowsMax; pos.y++, dstPtr = (__global TYPE*)((__global char*)dstptr + dststep)) #endif { - ASSERT(pos.y < dstCoords.y2); + ASSERT(pos.y < dstRowsMax); for ( #if BLOCK_SIZE_Y > 1 @@ -326,7 +325,7 @@ void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, #endif sy++, srcPos.y++) { - data[sy + sy_index] = readSrcPixel(srcPos, srcptr, srcstep, srcoffset, srcCoords + data[sy + sy_index] = readSrcPixel(srcPos, srcptr, srcstep, srcCoords #ifdef BORDER_CONSTANT , borderValue #endif @@ -361,7 +360,6 @@ void filter2D(__global const uchar* srcptr, int srcstep, int srcoffset, if (writeResult) { - ASSERT(pos.y >= dstCoords.y1 && pos.y < dstCoords.y2); *dstPtr = CONVERT_TO_TYPE(total_sum); } diff --git a/modules/imgproc/test/ocl/test_filter2d.cpp b/modules/imgproc/test/ocl/test_filter2d.cpp index ca1450952c..484080df44 100644 --- a/modules/imgproc/test/ocl/test_filter2d.cpp +++ b/modules/imgproc/test/ocl/test_filter2d.cpp @@ -49,17 +49,11 @@ namespace cvtest { namespace ocl { -enum -{ - noType = -1, -}; - - ///////////////////////////////////////////////////////////////////////////////////////////////// // Filter2D -PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool) +PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool, bool) { - static const int kernelMinSize = 1; + static const int kernelMinSize = 2; static const int kernelMaxSize = 10; int type; @@ -75,8 +69,8 @@ PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool) virtual void SetUp() { type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1)); - borderType = GET_PARAM(2); - useRoi = GET_PARAM(3); + borderType = GET_PARAM(2) | (GET_PARAM(3) ? BORDER_ISOLATED : 0); + useRoi = GET_PARAM(4); } void random_roi() @@ -84,16 +78,19 @@ PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool) dsize = randomSize(1, MAX_VALUE); Size ksize = randomSize(kernelMinSize, kernelMaxSize); - kernel = randomMat(ksize, CV_MAKE_TYPE(((CV_64F == CV_MAT_DEPTH(type)) ? CV_64F : CV_32F), 1), -MAX_VALUE, MAX_VALUE); + Mat temp = randomMat(ksize, CV_MAKE_TYPE(((CV_64F == CV_MAT_DEPTH(type)) ? CV_64F : CV_32F), 1), -MAX_VALUE, MAX_VALUE); + cv::normalize(temp, kernel, 1.0, 0.0, NORM_L1); - Size roiSize = randomSize(1, MAX_VALUE); + //Size roiSize = randomSize(ksize.width, MAX_VALUE, ksize.height, MAX_VALUE); + Size roiSize(1024, 1024); Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); randomSubMat(dst, dst_roi, dsize, dstBorder, type, -MAX_VALUE, MAX_VALUE); - anchor.x = anchor.y = -1; + anchor.x = randomInt(-1, ksize.width); + anchor.y = randomInt(-1, ksize.height); UMAT_UPLOAD_INPUT_PARAMETER(src) UMAT_UPLOAD_OUTPUT_PARAMETER(dst) @@ -128,7 +125,9 @@ OCL_INSTANTIATE_TEST_CASE_P(ImageProc, Filter2D, (BorderType)BORDER_REPLICATE, (BorderType)BORDER_REFLECT, (BorderType)BORDER_REFLECT_101), - Bool()) + Bool(), // BORDER_ISOLATED + Bool() // ROI + ) ); From f76bf8b425dd5af8a94517a35868f7f3514e34f3 Mon Sep 17 00:00:00 2001 From: Vladimir Bystricky Date: Wed, 11 Dec 2013 14:33:22 +0400 Subject: [PATCH 5/5] Fix error in test --- modules/imgproc/test/ocl/test_filter2d.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/modules/imgproc/test/ocl/test_filter2d.cpp b/modules/imgproc/test/ocl/test_filter2d.cpp index 484080df44..54d65459c5 100644 --- a/modules/imgproc/test/ocl/test_filter2d.cpp +++ b/modules/imgproc/test/ocl/test_filter2d.cpp @@ -81,8 +81,7 @@ PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool, bool) Mat temp = randomMat(ksize, CV_MAKE_TYPE(((CV_64F == CV_MAT_DEPTH(type)) ? CV_64F : CV_32F), 1), -MAX_VALUE, MAX_VALUE); cv::normalize(temp, kernel, 1.0, 0.0, NORM_L1); - //Size roiSize = randomSize(ksize.width, MAX_VALUE, ksize.height, MAX_VALUE); - Size roiSize(1024, 1024); + Size roiSize = randomSize(ksize.width, MAX_VALUE, ksize.height, MAX_VALUE); Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);