Small refactoring
authorAlexander Karsakov <alexander.karsakov@itseez.com>
Tue, 1 Apr 2014 07:27:43 +0000 (11:27 +0400)
committerAlexander Karsakov <alexander.karsakov@itseez.com>
Wed, 2 Apr 2014 12:38:32 +0000 (16:38 +0400)
modules/imgproc/src/filter.cpp
modules/imgproc/src/opencl/filterSepCol.cl
modules/imgproc/src/opencl/filterSepRow.cl
modules/imgproc/src/opencl/gaussian_blur_8u.cl [deleted file]
modules/imgproc/src/smooth.cpp
modules/imgproc/test/ocl/test_filters.cpp

index d81f8af..7d870d7 100644 (file)
@@ -3269,8 +3269,10 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
     return k.run(2, globalsize, localsize, false);
 }
 
+const int shift_bits = 8;
+
 static bool ocl_sepRowFilter2D(const UMat & src, UMat & buf, const Mat & kernelX, int anchor,
-                               int borderType, int ddepth, bool fast8uc1)
+                               int borderType, int ddepth, bool fast8uc1, bool int_arithm)
 {
     int type = src.type(), cn = CV_MAT_CN(type), sdepth = CV_MAT_DEPTH(type);
     bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
@@ -3303,14 +3305,15 @@ static bool ocl_sepRowFilter2D(const UMat & src, UMat & buf, const Mat & kernelX
 
     char cvt[40];
     cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s"
-                                          " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s%s",
+                                          " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s%s%s",
                                           radiusX, (int)localsize[0], (int)localsize[1], cn, btype,
                                           extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
                                           isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED",
                                           ocl::typeToStr(type), ocl::typeToStr(buf_type),
                                           ocl::convertTypeStr(sdepth, bdepth, cn, cvt),
                                           ocl::typeToStr(sdepth), ocl::typeToStr(bdepth),
-                                          doubleSupport ? " -D DOUBLE_SUPPORT" : "");
+                                          doubleSupport ? " -D DOUBLE_SUPPORT" : "",
+                                          int_arithm ? " -D INTEGER_ARITHMETIC" : "");
     build_options += ocl::kernelToStr(kernelX, bdepth);
 
     Size srcWholeSize; Point srcOffset;
@@ -3338,7 +3341,7 @@ static bool ocl_sepRowFilter2D(const UMat & src, UMat & buf, const Mat & kernelX
     return k.run(2, globalsize, localsize, false);
 }
 
-static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY, double delta, int anchor, int bits)
+static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY, double delta, int anchor, bool int_arithm)
 {
     bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0;
     if (dst.depth() == CV_64F && !doubleSupport)
@@ -3361,12 +3364,13 @@ static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY
     char cvt[40];
     cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d"
                                           " -D srcT=%s -D dstT=%s -D convertToDstT=%s"
-                                          " -D srcT1=%s -D dstT1=%s -D BITS=%d%s",
+                                          " -D srcT1=%s -D dstT1=%s -D SHIFT_BITS=%d%s%s",
                                           anchor, (int)localsize[0], (int)localsize[1], cn,
                                           ocl::typeToStr(buf_type), ocl::typeToStr(dtype),
                                           ocl::convertTypeStr(bdepth, ddepth, cn, cvt),
                                           ocl::typeToStr(bdepth), ocl::typeToStr(ddepth),
-                                          bits, doubleSupport ? " -D DOUBLE_SUPPORT" : "");
+                                          2*shift_bits, doubleSupport ? " -D DOUBLE_SUPPORT" : "",
+                                          int_arithm ? " -D INTEGER_ARITHMETIC" : "");
     build_options += ocl::kernelToStr(kernelY, bdepth);
 
     ocl::Kernel k("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc,
@@ -3459,62 +3463,55 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth,
     if (ddepth < 0)
         ddepth = sdepth;
 
-    //CV_OCL_RUN_(kernelY.cols <= 21 && kernelX.cols <= 21 &&
-    //            imgSize.width > optimizedSepFilterLocalSize + (kernelX.cols >> 1) &&
-    //            imgSize.height > optimizedSepFilterLocalSize + (kernelY.cols >> 1) &&
-    //            (!(borderType & BORDER_ISOLATED) || _src.offset() == 0) && anchor == Point(-1, -1) &&
-    //            (d.isIntel() || (d.isAMD() && !d.hostUnifiedMemory())),
-    //            ocl_sepFilter2D_SinglePass(_src, _dst, kernelX, kernelY, delta,
-    //                                       borderType & ~BORDER_ISOLATED, ddepth), true)
-
     if (anchor.x < 0)
         anchor.x = kernelX.cols >> 1;
     if (anchor.y < 0)
         anchor.y = kernelY.cols >> 1;
 
-    UMat src = _src.getUMat();
-    Size srcWholeSize; Point srcOffset;
-    src.locateROI(srcWholeSize, srcOffset);
-
-    //bool fast8uc1 = type == CV_8UC1 && srcOffset.x % 4 == 0 &&
-    //        src.cols % 4 == 0 && src.step % 4 == 0;
-    bool fast8uc1 = false;
-
     int rtype = getKernelType(kernelX,
         kernelX.rows == 1 ? Point(anchor.x, 0) : Point(0, anchor.x));
     int ctype = getKernelType(kernelY,
         kernelY.rows == 1 ? Point(anchor.y, 0) : Point(0, anchor.y));
 
     int bdepth = CV_32F;
-    int bits = 0;
-
+    bool int_arithm = false;
     if( sdepth == CV_8U &&
         ((rtype == KERNEL_SMOOTH+KERNEL_SYMMETRICAL &&
           ctype == KERNEL_SMOOTH+KERNEL_SYMMETRICAL &&
           ddepth == CV_8U)))
     {
         bdepth = CV_32S;
-        bits = 8;
-        _kernelX.getMat().convertTo( kernelX, CV_32S, 1 << bits );
-        _kernelY.getMat().convertTo( kernelY, CV_32S, 1 << bits );
-        kernelX = kernelX.reshape(1,1);
-        kernelY = kernelY.reshape(1,1);
-        bits *= 2;
-        delta *= (1 << bits);
+        _kernelX.getMat().reshape(1,1).convertTo( kernelX, CV_32S, 1 << shift_bits );
+        _kernelY.getMat().reshape(1,1).convertTo( kernelY, CV_32S, 1 << shift_bits );
+        int_arithm = true;
     }
 
+    CV_OCL_RUN_(kernelY.cols <= 21 && kernelX.cols <= 21 && !int_arithm &&
+                imgSize.width > optimizedSepFilterLocalSize + anchor.x &&
+                imgSize.height > optimizedSepFilterLocalSize + anchor.y &&
+                (!(borderType & BORDER_ISOLATED) || _src.offset() == 0) &&
+                anchor == Point(kernelX.cols >> 1, kernelY.cols >> 1) &&
+                (d.isIntel() || (d.isAMD() && !d.hostUnifiedMemory())),
+                ocl_sepFilter2D_SinglePass(_src, _dst, kernelX, kernelY, delta,
+                                           borderType & ~BORDER_ISOLATED, ddepth), true)
+
+    UMat src = _src.getUMat();
+    Size srcWholeSize; Point srcOffset;
+    src.locateROI(srcWholeSize, srcOffset);
+
+    bool fast8uc1 = type == CV_8UC1 && srcOffset.x % 4 == 0 &&
+            src.cols % 4 == 0 && src.step % 4 == 0 && !int_arithm;
+
     Size srcSize = src.size();
     Size bufSize(srcSize.width, srcSize.height + kernelY.cols - 1);
     UMat buf(bufSize, CV_MAKETYPE(bdepth, cn));
-    if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, ddepth, fast8uc1))
+    if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, ddepth, fast8uc1, int_arithm))
         return false;
 
-    Mat buffer = buf.getMat(ACCESS_READ);
-
     _dst.create(srcSize, CV_MAKETYPE(ddepth, cn));
     UMat dst = _dst.getUMat();
 
-    return ocl_sepColFilter2D(buf, dst, kernelY, delta, anchor.y, bits);
+    return ocl_sepColFilter2D(buf, dst, kernelY, delta, anchor.y, int_arithm);
 }
 
 #endif
index 94730d8..1359505 100644 (file)
@@ -3,6 +3,7 @@
 //
 // 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) 2014, Itseez, Inc, all rights reserved.
 // Third party copyrights are property of their respective owners.
 //
 // @Authors
@@ -96,12 +97,15 @@ __kernel void col_filter(__global const uchar * src, int src_step, int src_offse
     {
         temp[0] = LDS_DAT[l_y + RADIUSY - i][l_x];
         temp[1] = LDS_DAT[l_y + RADIUSY + i][l_x];
+#ifndef INTEGER_ARITHMETIC
         sum += mad(temp[0], mat_kernel[RADIUSY - i], temp[1] * mat_kernel[RADIUSY + i]);
-        //sum += temp[0]*mat_kernel[RADIUSY - i] + temp[1] * mat_kernel[RADIUSY + i];
+#else
+        sum += mad24(temp[0],mat_kernel[RADIUSY - i], temp[1] * mat_kernel[RADIUSY + i]);
+#endif
     }
 
-#if BITS > 0
-    sum = sum >> BITS;
+#ifdef INTEGER_ARITHMETIC
+    sum = (sum + (1 << (SHIFT_BITS-1))) >> SHIFT_BITS;
 #endif
 
     // write the result to dst
index 8deec35..890eeb8 100644 (file)
@@ -3,6 +3,7 @@
 //
 // 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) 2014, Itseez, Inc, all rights reserved.
 // Third party copyrights are property of their respective owners.
 //
 // @Authors
@@ -355,8 +356,11 @@ __kernel void row_filter(__global const uchar * src, int src_step, int src_offse
     {
         temp[0] = LDS_DAT[l_y][l_x + RADIUSX - i];
         temp[1] = LDS_DAT[l_y][l_x + RADIUSX + i];
+#ifndef INTEGER_ARITHMETIC
         sum += mad(convertToDstT(temp[0]), mat_kernel[RADIUSX - i], convertToDstT(temp[1]) * mat_kernel[RADIUSX + i]);
-        //sum += convertToDstT(temp[0])*mat_kernel[RADIUSX - i] + convertToDstT(temp[1]) * mat_kernel[RADIUSX + i];
+#else
+        sum += mad24(convertToDstT(temp[0]), mat_kernel[RADIUSX - i], convertToDstT(temp[1]) * mat_kernel[RADIUSX + i]);
+#endif
     }
 
     // write the result to dst
diff --git a/modules/imgproc/src/opencl/gaussian_blur_8u.cl b/modules/imgproc/src/opencl/gaussian_blur_8u.cl
deleted file mode 100644 (file)
index 268d8b7..0000000
+++ /dev/null
@@ -1,189 +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) 2014, Intel Corporation, 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*/
-
-///////////////////////////////////////////////////////////////////////////////////////////////////
-/////////////////////////////////Macro for border type////////////////////////////////////////////
-/////////////////////////////////////////////////////////////////////////////////////////////////
-
-#ifdef BORDER_CONSTANT
-// CCCCCC|abcdefgh|CCCCCCC
-#define EXTRAPOLATE(x, maxV)
-#elif defined BORDER_REPLICATE
-// aaaaaa|abcdefgh|hhhhhhh
-#define EXTRAPOLATE(x, maxV) \
-    { \
-        (x) = max(min((x), (maxV) - 1), 0); \
-    }
-#elif defined BORDER_WRAP
-// cdefgh|abcdefgh|abcdefg
-#define EXTRAPOLATE(x, maxV) \
-    { \
-        (x) = ( (x) + (maxV) ) % (maxV); \
-    }
-#elif defined BORDER_REFLECT
-// fedcba|abcdefgh|hgfedcb
-#define EXTRAPOLATE(x, maxV) \
-    { \
-        (x) = min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) ); \
-    }
-#elif defined BORDER_REFLECT_101 || defined BORDER_REFLECT101
-// gfedcb|abcdefgh|gfedcba
-#define EXTRAPOLATE(x, maxV) \
-    { \
-        (x) = min(((maxV)-1)*2-(x), max((x),-(x)) ); \
-    }
-#else
-#error No extrapolation method
-#endif
-
-#if CN != 3
-#define loadpix(addr) *(__global const srcT *)(addr)
-#define storepix(val, addr)  *(__global dstT *)(addr) = val
-#define SRCSIZE (int)sizeof(srcT)
-#define DSTSIZE (int)sizeof(dstT)
-#else
-#define loadpix(addr)  vload3(0, (__global const srcT1 *)(addr))
-#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
-#define SRCSIZE (int)sizeof(srcT1)*3
-#define DSTSIZE (int)sizeof(dstT1)*3
-#endif
-
-#define SRC(_x,_y) convertToWT(loadpix(Src + mad24(_y, src_step, SRCSIZE * _x)))
-
-#ifdef BORDER_CONSTANT
-// CCCCCC|abcdefgh|CCCCCCC
-#define ELEM(_x,_y,r_edge,t_edge,const_v) (_x)<0 | (_x) >= (r_edge) | (_y)<0 | (_y) >= (t_edge) ? (const_v) : SRC((_x),(_y))
-#else
-#define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y))
-#endif
-
-#define noconvert
-
-// horizontal and vertical filter kernels
-// should be defined on host during compile time to avoid overhead
-#define DIG(a) a,
-__constant int mat_kernelX[] = { KERNEL_MATRIX_X };
-__constant int mat_kernelY[] = { KERNEL_MATRIX_Y };
-
-__kernel void gaussian_blur_8u(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width,
-                         __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
-{
-    // RADIUSX, RADIUSY are filter dimensions
-    // BLK_X, BLK_Y are local wrogroup sizes
-    // all these should be defined on host during compile time
-    // first lsmem array for source pixels used in first pass,
-    // second lsmemDy for storing first pass results
-    __local WT lsmem[BLK_Y + 2 * RADIUSY][BLK_X + 2 * RADIUSX];
-    __local WT lsmemDy[BLK_Y][BLK_X + 2 * RADIUSX];
-
-    // get local and global ids - used as image and local memory array indexes
-    int lix = get_local_id(0);
-    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;
-    int xb = srcX;
-    int yb = srcY;
-
-    // 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;
-    do
-    {
-        int yb = cSrcY;
-        EXTRAPOLATE(yb, (height));
-
-        int clocX = lix;
-        int cSrcX = srcX;
-        do
-        {
-            int xb = cSrcX;
-            EXTRAPOLATE(xb,(width));
-            lsmem[clocY][clocX] = ELEM(xb, yb, (width), (height), 0 );
-
-            clocX += BLK_X;
-            cSrcX += BLK_X;
-        }
-        while(clocX < BLK_X+(RADIUSX*2));
-
-        clocY += BLK_Y;
-        cSrcY += BLK_Y;
-    }
-    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 = 0;
-    do
-    {
-        sum = 0;
-        for (i=0; i<=2*RADIUSY; i++)
-            sum = mad(lsmem[liy+i][clocX], mat_kernelY[i], sum);
-        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;
-    for (i=0; i<=2*RADIUSX; i++)
-        sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum);
-    
-    sum = sum >> (GAUSSIAN_COEF_BITS * 2);
-
-    //store result into destination image
-    storepix(convertToDstT(sum), Dst + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset)));
-}
index e2365cd..6a18af5 100644 (file)
@@ -42,7 +42,6 @@
 
 #include "precomp.hpp"
 #include "opencl_kernels.hpp"
-#include <iostream>
 
 /*
  * This file includes the code, contributed by Simon Perreault
@@ -1070,73 +1069,6 @@ static void createGaussianKernels( Mat & kx, Mat & ky, int type, Size ksize,
         ky = getGaussianKernel( ksize.height, sigma2, std::max(depth, CV_32F) );
 }
 
-#define GAUSSIAN_COEF_BITS 11
-
-static bool GaussianBlur_8u(InputArray _src, OutputArray _dst, Size ksize,
-                   double sigma1, double sigma2,
-                   int borderType)
-{
-    int type = _src.type();
-    Mat kx, ky;
-    createGaussianKernels(kx, ky, CV_64F, ksize, sigma1, sigma2);
-    Mat kx_8u, ky_8u;
-
-    int scale_coef = 1 << GAUSSIAN_COEF_BITS;
-    kx.convertTo(kx_8u, CV_32S, scale_coef);
-    ky.convertTo(ky_8u, CV_32S, scale_coef);
-
-    kx_8u.reshape(1, 1);
-    ky_8u.reshape(1, 1);
-
-    Size size = _src.size(), wholeSize;
-    Point origin;
-    int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype),
-            esz = CV_ELEM_SIZE(stype), wdepth = CV_32S,
-            ddepth = sdepth;
-    size_t src_step = _src.step(), src_offset = _src.offset();
-
-    if ((src_offset % src_step) % esz != 0 || !(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE ||
-              borderType == BORDER_REFLECT || borderType == BORDER_WRAP ||
-              borderType == BORDER_REFLECT_101))
-        return false;
-
-    size_t lt2[2] = { 16, 16 };
-    size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), lt2[1] * (1 + (size.height - 1) / lt2[1]) };
-
-    char cvt[2][40];
-    const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP",
-                                       "BORDER_REFLECT_101" };
-
-    String opts = cv::format("-D BLK_X=%d -D BLK_Y=%d -D RADIUSX=%d -D RADIUSY=%d%s%s"
-                             " -D srcT=%s -D convertToWT=%s -D WT=%s -D dstT=%s -D convertToDstT=%s"
-                             " -D %s -D srcT1=%s -D dstT1=%s -D CN=%d -D GAUSSIAN_COEF_BITS=%d", (int)lt2[0], (int)lt2[1],
-                             kx.rows / 2, kx.rows / 2,
-                             ocl::kernelToStr(kx_8u, CV_32S, "KERNEL_MATRIX_X").c_str(),
-                             ocl::kernelToStr(ky_8u, CV_32S, "KERNEL_MATRIX_Y").c_str(),
-                             ocl::typeToStr(stype), ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]),
-                             ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)), ocl::typeToStr(stype),
-                             ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]), borderMap[borderType],
-                             ocl::typeToStr(sdepth), ocl::typeToStr(ddepth), cn, GAUSSIAN_COEF_BITS);
-
-    ocl::Kernel k("gaussian_blur_8u", ocl::imgproc::gaussian_blur_8u_oclsrc, opts);
-    if (k.empty())
-        return false;
-
-    UMat src = _src.getUMat();
-    _dst.create(size, stype);
-    UMat dst = _dst.getUMat();
-
-    int src_offset_x = static_cast<int>((src_offset % src_step) / esz);
-    int src_offset_y = static_cast<int>(src_offset / src_step);
-
-    src.locateROI(wholeSize, origin);
-
-    k.args(ocl::KernelArg::PtrReadOnly(src), (int)src_step, src_offset_x, src_offset_y,
-        wholeSize.height, wholeSize.width, ocl::KernelArg::WriteOnly(dst));
-
-    return k.run(2, gt2, lt2, false);
-}
-
 }
 
 cv::Ptr<cv::FilterEngine> cv::createGaussianFilter( int type, Size ksize,
@@ -1150,8 +1082,6 @@ cv::Ptr<cv::FilterEngine> cv::createGaussianFilter( int type, Size ksize,
 }
 
 
-
-
 void cv::GaussianBlur( InputArray _src, OutputArray _dst, Size ksize,
                    double sigma1, double sigma2,
                    int borderType )
@@ -1196,13 +1126,6 @@ void cv::GaussianBlur( InputArray _src, OutputArray _dst, Size ksize,
     }
 #endif
 
-    //if (type == CV_8U)
-    //{
-    //    CV_OCL_RUN_(_dst.isUMat() && _src.dims() <= 2 && 
-    //            (!(borderType & BORDER_ISOLATED) || _src.offset() == 0),
-    //            GaussianBlur_8u(_src, _dst, ksize, sigma1, sigma2, borderType))
-    //}
-
     Mat kx, ky;
     createGaussianKernels(kx, ky, type, ksize, sigma1, sigma2);
     sepFilter2D(_src, _dst, CV_MAT_DEPTH(type), kx, ky, Point(-1,-1), 0, borderType );
index aee1f08..bf45a72 100644 (file)
@@ -209,7 +209,7 @@ typedef FilterTestBase GaussianBlurTest;
 
 OCL_TEST_P(GaussianBlurTest, Mat)
 {
-    for (int j = 0; j < test_loop_times + 100; j++)
+    for (int j = 0; j < test_loop_times; j++)
     {
         random_roi();
 
@@ -219,28 +219,7 @@ OCL_TEST_P(GaussianBlurTest, Mat)
         OCL_OFF(cv::GaussianBlur(src_roi, dst_roi, Size(ksize, ksize), sigma1, sigma2, borderType));
         OCL_ON(cv::GaussianBlur(usrc_roi, udst_roi, Size(ksize, ksize), sigma1, sigma2, borderType));
 
-
-        if (checkNorm2(dst_roi, udst_roi) > 2 && CV_MAT_DEPTH(type) == CV_8U)
-        {
-            std::cout << "i = " << j << std::endl;
-            Mat uudst = udst_roi.getMat(ACCESS_READ);
-            Mat diff; 
-            absdiff(dst_roi, udst, diff);
-            int nonZero = countNonZero(diff);
-            double max;
-            Point maxn;
-            minMaxLoc(diff, (double*)0, &max, (Point*) 0, &maxn);
-
-            uchar a = dst_roi.at<uchar>(maxn);
-            uchar b = uudst.at<uchar>(maxn);
-
-            std::cout << "dst_roi" << dst_roi << std::endl;
-            std::cout << "udst_roi" << uudst << std::endl;
-        }
-
-        
-
-        Near(CV_MAT_DEPTH(type) == CV_8U ? 1 : 5e-5, false);
+        Near(CV_MAT_DEPTH(type) >= CV_32F ? 5e-5 : 1, false);
     }
 }