Merge pull request #1958 from vbystricky:ocl_filter2D
This commit is contained in:
Andrey Pavlenko
OpenCV Buildbot
commit
129b715e3d
@ -41,6 +41,7 @@
|
|||||||
//M*/
|
//M*/
|
||||||
|
|
||||||
#include "precomp.hpp"
|
#include "precomp.hpp"
|
||||||
|
#include "opencl_kernels.hpp"
|
||||||
|
|
||||||
/****************************************************************************************\
|
/****************************************************************************************\
|
||||||
Base Image Filter
|
Base Image Filter
|
||||||
@ -3115,6 +3116,206 @@ template<typename ST, class CastOp, class VecOp> 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 <typename T>
|
||||||
|
static int _prepareKernelFilter2D(std::vector<T>& data, const Mat &kernel)
|
||||||
|
{
|
||||||
|
Mat _kernel; kernel.convertTo(_kernel, DataDepth<T>::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<T>(y, x);
|
||||||
|
data[x * size_y_aligned + y + kernel.rows] = _kernel.at<T>(y, x);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
return size_y_aligned;
|
||||||
|
}
|
||||||
|
|
||||||
|
static 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<float> kernelMatDataFloat;
|
||||||
|
std::vector<double> kernelMatDataDouble;
|
||||||
|
int kernel_size_y2_aligned = useDouble ?
|
||||||
|
_prepareKernelFilter2D<double>(kernelMatDataDouble, kernelMat)
|
||||||
|
: _prepareKernelFilter2D<float>(kernelMatDataFloat, kernelMat);
|
||||||
|
|
||||||
|
|
||||||
|
cv::Size sz = _src.size();
|
||||||
|
size_t globalsize[2] = {sz.width, sz.height};
|
||||||
|
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];
|
||||||
|
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 = 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))
|
||||||
|
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();
|
||||||
|
if (src.empty())
|
||||||
|
src = _src.getUMat();
|
||||||
|
|
||||||
|
int idxArg = 0;
|
||||||
|
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};
|
||||||
|
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::BaseFilter> cv::getLinearFilter(int srcType, int dstType,
|
cv::Ptr<cv::BaseFilter> cv::getLinearFilter(int srcType, int dstType,
|
||||||
InputArray filter_kernel, Point anchor,
|
InputArray filter_kernel, Point anchor,
|
||||||
double delta, int bits)
|
double delta, int bits)
|
||||||
@ -3230,6 +3431,10 @@ void cv::filter2D( InputArray _src, OutputArray _dst, int ddepth,
|
|||||||
InputArray _kernel, Point anchor,
|
InputArray _kernel, Point anchor,
|
||||||
double delta, int borderType )
|
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();
|
Mat src = _src.getMat(), kernel = _kernel.getMat();
|
||||||
|
|
||||||
if( ddepth < 0 )
|
if( ddepth < 0 )
|
||||||
|
|||||||
375
modules/imgproc/src/opencl/filter2D.cl
Normal file
375
modules/imgproc/src/opencl/filter2D.cl
Normal file
@ -0,0 +1,375 @@
|
|||||||
|
/*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 = minX - (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 = minY - (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, 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 + 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 + 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 srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY,
|
||||||
|
__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 = {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;
|
||||||
|
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)(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 >= 0) && (pos.x < cols));
|
||||||
|
|
||||||
|
#if BLOCK_SIZE_Y > 1
|
||||||
|
bool readAllpixels = true;
|
||||||
|
int sy_index = 0; // current index in data[] array
|
||||||
|
|
||||||
|
dstRowsMax = min(rows, pos.y + BLOCK_SIZE_Y);
|
||||||
|
for (;
|
||||||
|
pos.y < dstRowsMax;
|
||||||
|
pos.y++,
|
||||||
|
dstPtr = (__global TYPE*)((__global char*)dstptr + dststep))
|
||||||
|
#endif
|
||||||
|
{
|
||||||
|
ASSERT(pos.y < dstRowsMax);
|
||||||
|
|
||||||
|
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, 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)
|
||||||
|
{
|
||||||
|
*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
|
||||||
|
}
|
||||||
|
}
|
||||||
135
modules/imgproc/test/ocl/test_filter2d.cpp
Normal file
135
modules/imgproc/test/ocl/test_filter2d.cpp
Normal file
@ -0,0 +1,135 @@
|
|||||||
|
/*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 {
|
||||||
|
|
||||||
|
/////////////////////////////////////////////////////////////////////////////////////////////////
|
||||||
|
// Filter2D
|
||||||
|
PARAM_TEST_CASE(Filter2D, MatDepth, Channels, BorderType, bool, bool)
|
||||||
|
{
|
||||||
|
static const int kernelMinSize = 2;
|
||||||
|
static const int kernelMaxSize = 10;
|
||||||
|
|
||||||
|
int type;
|
||||||
|
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) | (GET_PARAM(3) ? BORDER_ISOLATED : 0);
|
||||||
|
useRoi = GET_PARAM(4);
|
||||||
|
}
|
||||||
|
|
||||||
|
void random_roi()
|
||||||
|
{
|
||||||
|
dsize = randomSize(1, MAX_VALUE);
|
||||||
|
|
||||||
|
Size ksize = randomSize(kernelMinSize, kernelMaxSize);
|
||||||
|
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);
|
||||||
|
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 = randomInt(-1, ksize.width);
|
||||||
|
anchor.y = randomInt(-1, ksize.height);
|
||||||
|
|
||||||
|
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(), // BORDER_ISOLATED
|
||||||
|
Bool() // ROI
|
||||||
|
)
|
||||||
|
);
|
||||||
|
|
||||||
|
|
||||||
|
} } // namespace cvtest::ocl
|
||||||
|
|
||||||
|
#endif // HAVE_OPENCL
|
||||||
Loading…
Reference in New Issue
Block a user