Updated the prohibited list, and fixed a type casting in hough_lines,

Added Filter2D with Border Replicate and No extra extrapolation to the prohibited list to prevent Laplacian Filter accuracy failures.

Added LUT to the prohibited list to address the LUT failures in the core accuracy test.

diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index e662cacb9c76ec9e2b84ae42b676f70ae4e4301b..6a150ec8f751c4ecd704a8956c79cd881f75b05b 100644 (file)
--- a/modules/core/src/ocl.cpp
+++ b/modules/core/src/ocl.cpp
@@ -69,7 +69,7 @@
 
 #ifdef CV_TIOPENCL
 
-#define MAX_PROHIBITED_SIZE  12
+#define MAX_PROHIBITED_SIZE  17
 
 const char * const prohibitedList[]={
          "mixChannels", "", // Module: Core // repeated errors in accuracy tests. //Need to investigate why accuracy tests fail
@@ -83,7 +83,12 @@ const char * const prohibitedList[]={
       "warpBackwardKernel", "", //Module: Video // Image datatype is not supported by TI OpenCL
          "classify_hists_180_kernel", "", //Module: objdetect // Accuracy tests fail OCL_ObjDetect/HOG.Detect/0, where GetParam() = (64x128, 8UC1), & (48x96, 8UC1)
          "classify_hists_252_kernel", "", //Module: objdetect // Accuracy tests fail OCL_ObjDetect/HOG.Detect/0, where GetParam() = (64x128, 8UC1), & (48x96, 8UC1)
-         "calcLUT", "" // Module: imgproc // Performance failures with OCL_EqualizeHistFixture_EqualizeHist.EqualizeHist
+         "calcLUT", "", // Module: imgproc // Performance failures with OCL_EqualizeHistFixture_EqualizeHist.EqualizeHist
+         "filter2D", "-D BORDER_REPLICATE -D EXTRA_EXTRAPOLATION -D srcT=float",
+         "filter2D", "-D BORDER_REPLICATE -D NO_EXTRA_EXTRAPOLATION -D srcT=float",
+         "filter2D", "-D BORDER_REFLECT -D NO_EXTRA_EXTRAPOLATION -D srcT=float",
+         "filter2D", "-D BORDER_REFLECT_101 -D NO_EXTRA_EXTRAPOLATION -D srcT=float",
+         "LUT", "-D dcn=4 -D lcn=1 -D srcT=uchar -D dstT=uchar"
 /*       "split","", // Module: Core
       "mixChannels", "", // Module: Core // repeated errors in accuracy tests
       "fastNlMeansDenoising", "cn=3", // Module: Photo //OpenCL runtime bug while setting kernel argument for float3 datatype.
@@ -3598,6 +3603,15 @@ int Kernel::set(int i, const KernelArg& arg)
 
         if (!h)
         {
+               /* Dnparikh: When an allocation fails when a UMat is created on the DSP side.
+                * Say 2 buffers are allocated as Umat and the 3rd buffer is allocated as Mat
+                * When the third buffer is set as an argument to the kernel, since the 3rd
+                * buffer is a Mat and not UMat the kernel is deleted. However before deleting
+                * the kernel, the arguments that are set need to be cleaned up so that Urefcount
+                * is decremented so that the buffer is actually deallocated when it is cleaned up!
+                * Dnparikh added the p->cleanupUMats line.
+                */
+               p->cleanupUMats();
             p->release();
             p = 0;
             return -1;
@@ -3670,7 +3684,7 @@ bool Kernel::run(int dims, size_t _globalsize[], size_t _localsize[],
                  bool sync, const Queue& q)
 {
     if(!p || !p->handle || p->e != 0)
-        return false;
+       return false;
 
     cl_command_queue qq = getQueue(q);
     size_t offset[CV_MAX_DIM] = {0}, globalsize[CV_MAX_DIM] = {1,1,1};

diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp
index f70389b07a33f6c7f6529b2f69ea1924c397c127..f61ea015e28c2ce117a0ef607191bd15d5b54630 100644 (file)
--- a/modules/imgproc/src/filter.cpp
+++ b/modules/imgproc/src/filter.cpp
@@ -43,6 +43,8 @@
 #include "precomp.hpp"
 #include "opencl_kernels_imgproc.hpp"
 
+#include <iostream>
+
 /****************************************************************************************\
                                     Base Image Filter
 \****************************************************************************************/
@@ -4050,11 +4052,13 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
 
     // For smaller filter kernels, there is a special kernel that is more
     // efficient than the general one.
+
     UMat kernalDataUMat;
     if (device.isIntel() && (device.type() & ocl::Device::TYPE_GPU) &&
         ((ksize.width < 5 && ksize.height < 5) ||
         (ksize.width == 5 && ksize.height == 5 && cn == 1)))
     {
+
         kernelMat = kernelMat.reshape(0, 1);
         String kerStr = ocl::kernelToStr(kernelMat, CV_32F);
         int h = isolated ? sz.height : wholeSize.height;
@@ -4139,7 +4143,24 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
 
             if ((w < ksize.width) || (h < ksize.height))
                 return false;
-
+#ifdef CV_TIOPENCL
+            String opts = format("-D LOCAL_SIZE=%d -D cn=%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 srcT=%s -D srcT1=%s -D dstT=%s -D dstT1=%s -D WT=%s -D WT1=%s "
+                                                "-D %s%s%s "
+                                 "-D convertToWT=%s -D convertToDstT=%s",
+                                 (int)BLOCK_SIZE, cn, anchor.x, anchor.y,
+                                 ksize.width, ksize.height, kernel_size_y2_aligned, borderMap[borderType],
+                                 extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
+                                 ocl::typeToStr(type), ocl::typeToStr(sdepth), ocl::typeToStr(dtype),
+                                 ocl::typeToStr(ddepth), ocl::typeToStr(wtype), ocl::typeToStr(wdepth),
+                                 isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED",
+                                 doubleSupport ? " -D DOUBLE_SUPPORT" : "", kerStr.c_str(),
+                                 ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]),
+                                 ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]));
+#else
             String opts = format("-D LOCAL_SIZE=%d -D cn=%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%s%s "
@@ -4154,6 +4175,7 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth,
                                  ocl::typeToStr(ddepth), ocl::typeToStr(wtype), ocl::typeToStr(wdepth),
                                  ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]),
                                  ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]));
+#endif
 
             localsize[0] = BLOCK_SIZE;
             globalsize[0] = DIVUP(sz.width, BLOCK_SIZE - (ksize.width - 1)) * BLOCK_SIZE;

diff --git a/modules/imgproc/src/opencl/hough_lines.cl b/modules/imgproc/src/opencl/hough_lines.cl
index 4c2d0a94cd5493accdba1cdfa0c6fa226d9d4592..3329b3001224c5d820ce27f3182ee1abdd7eabe7 100644 (file)
--- a/modules/imgproc/src/opencl/hough_lines.cl
+++ b/modules/imgproc/src/opencl/hough_lines.cl
@@ -271,9 +271,9 @@ __kernel void get_lines(__global const uchar * accum_ptr, int accum_step, int ac
                 return;
 
             for (;;)
-            {
-                if (*(src_ptr + mad24(p0.y, src_step, p0.x + src_offset)))
-                {
+            {                  
+                if (*(src_ptr + mad24((int) p0.y, src_step, ((int) p0.x) + src_offset)))
+                {                
                     gap = 0;
 
                     if (!inLine)

diff --git a/modules/ts/src/ocl_test.cpp b/modules/ts/src/ocl_test.cpp
index 69d10aac0dab0279ea613192a65b7c82a8bb129d..acf4c32eb54a899c6a3f4d505e6623fe49d8ccf8 100644 (file)
--- a/modules/ts/src/ocl_test.cpp
+++ b/modules/ts/src/ocl_test.cpp
@@ -144,9 +144,13 @@ void dumpOpenCLDevice()
         DUMP_MESSAGE_STDOUT("    Platform = "<< device.getPlatform().name());
         DUMP_PROPERTY_XML("cv_ocl_current_platformName", device.getPlatform().name());
 #endif
-
+#ifdef CV_TIOPENCL
+        const char* deviceTypeStr = device.type() == Device::TYPE_CPU
+                    ? ("CPU") : (device.type() == Device::TYPE_GPU ? device.hostUnifiedMemory() ? "iGPU" : "dGPU" : (device.type() == Device::TYPE_ACCELERATOR ? "TI_DSP" : "unknown"));
+#else
         const char* deviceTypeStr = device.type() == Device::TYPE_CPU
             ? ("CPU") : (device.type() == Device::TYPE_GPU ? device.hostUnifiedMemory() ? "iGPU" : "dGPU" : "unknown");
+#endif
         DUMP_MESSAGE_STDOUT("    Type = "<< deviceTypeStr);
         DUMP_PROPERTY_XML("cv_ocl_current_deviceType", deviceTypeStr);
 

diff --git a/platforms/linux/arm-gnueabi.toolchain.cmake b/platforms/linux/arm-gnueabi.toolchain.cmake
index 1715354b5a07c1ba086832e7a41caf6099f1ff75..47075c2e874a47836e59dea221e51dedd13e3483 100644 (file)
--- a/platforms/linux/arm-gnueabi.toolchain.cmake
+++ b/platforms/linux/arm-gnueabi.toolchain.cmake
@@ -13,6 +13,7 @@ find_program(CMAKE_C_COMPILER NAMES arm-linux-gnueabi${FLOAT_ABI_SUFFIX}-gcc)
 find_program(CMAKE_CXX_COMPILER NAMES arm-linux-gnueabi${FLOAT_ABI_SUFFIX}-g++)
 find_program(CMAKE_AR NAMES arm-linux-gnueabi${FLOAT_ABI_SUFFIX}-ar)
 find_program(CMAKE_RANLIB NAMES arm-linux-gnueabi${FLOAT_ABI_SUFFIX}-ranlib)
+find_program(PKG_CONFIG_EXECUTABLE NAMES pkg-config)
 set(ARM_LINUX_SYSROOT /usr/arm-linux-gnueabi${FLOAT_ABI_SUFFIX} CACHE PATH "ARM cross compilation system root")
 
 set(CMAKE_CXX_FLAGS           ""                    CACHE STRING "c++ flags")