minor ocl.cpp refactoring

fix for cv::LUT and cv::transpose

modules/core/include/opencv2/core/ocl.hpp

@@ -90,7 +90,8 @@ public:
     String vendor() const;
     String OpenCL_C_Version() const;
     String OpenCLVersion() const;
-    String deviceVersion() const;
+    int deviceVersionMajor() const;
+    int deviceVersionMinor() const;
     String driverVersion() const;
     void* ptr() const;
 
@@ -224,16 +225,12 @@ public:
     static Context2& getDefault(bool initialize = true);
     void* ptr() const;
 
-    struct Impl;
+    friend void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device);
-    inline struct Impl* _getImpl() const { return p; }
 protected:
+    struct Impl;
     Impl* p;
 };
 
-
-// TODO Move to internal header
-void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device);
-
 class CV_EXPORTS Platform
 {
 public:
@@ -245,12 +242,14 @@ public:
     void* ptr() const;
     static Platform& getDefault();
 
-    struct Impl;
+    friend void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device);
-    inline struct Impl* _getImpl() const { return p; }
 protected:
+    struct Impl;
     Impl* p;
 };
 
+// TODO Move to internal header
+void initializeContextFromHandle(Context2& ctx, void* platform, void* context, void* device);
 
 class CV_EXPORTS Queue
 {
@@ -585,9 +584,12 @@ class CV_EXPORTS Image2D
 {
 public:
     Image2D();
-    Image2D(const UMat &src);
+    explicit Image2D(const UMat &src);
+    Image2D(const Image2D & i);
     ~Image2D();
 
+    Image2D & operator = (const Image2D & i);
+
     void* ptr() const;
 protected:
     struct Impl;

modules/core/src/convert.cpp

@@ -1505,6 +1505,9 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst)
                   format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s%s", dcn, lcn,
                          ocl::typeToStr(src.depth()), ocl::typeToStr(ddepth),
                          doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
+    if (k.empty())
+        return false;
+
     k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut),
            ocl::KernelArg::WriteOnly(dst));
 

modules/core/src/matrix.cpp

@@ -2909,6 +2909,9 @@ static bool ocl_transpose( InputArray _src, OutputArray _dst )
     ocl::Kernel k(kernelName.c_str(), ocl::core::transpose_oclsrc,
                   format("-D T=%s -D TILE_DIM=%d -D BLOCK_ROWS=%d",
                          ocl::memopTypeToStr(type), TILE_DIM, BLOCK_ROWS));
+    if (k.empty())
+        return false;
+
     if (inplace)
         k.args(ocl::KernelArg::ReadWriteNoSize(dst), dst.rows);
     else

modules/ts/src/ocl_test.cpp

@@ -160,17 +160,10 @@ void dumpOpenCLDevice()
         DUMP_MESSAGE_STDOUT("    Max memory allocation size = "<< maxMemAllocSizeStr);
         DUMP_PROPERTY_XML("cv_ocl_current_maxMemAllocSize", device.maxMemAllocSize());
 
-#if 0
-        const char* doubleSupportStr = device.haveDoubleSupport() ? "Yes" : "No";
-        DUMP_MESSAGE_STDOUT("    Double support = "<< doubleSupportStr);
-        DUMP_PROPERTY_XML("cv_ocl_current_haveDoubleSupport", device.haveDoubleSupport());
-#else
         const char* doubleSupportStr = device.doubleFPConfig() > 0 ? "Yes" : "No";
         DUMP_MESSAGE_STDOUT("    Double support = "<< doubleSupportStr);
         DUMP_PROPERTY_XML("cv_ocl_current_haveDoubleSupport", device.doubleFPConfig() > 0);
 
-#endif
-
         const char* isUnifiedMemoryStr = device.hostUnifiedMemory() ? "Yes" : "No";
         DUMP_MESSAGE_STDOUT("    Host unified memory = "<< isUnifiedMemoryStr);
         DUMP_PROPERTY_XML("cv_ocl_current_hostUnifiedMemory", device.hostUnifiedMemory());

modules/video/src/opencl/optical_flow_farneback.cl

@@ -142,11 +142,6 @@ inline int idx_row_high(const int y, const int last_row)
     return abs(last_row - abs(last_row - y)) % (last_row + 1);
 }
 
-inline int idx_row(const int y, const int last_row)
-{
-    return idx_row_low(idx_row_high(y, last_row), last_row);
-}
-
 inline int idx_col_low(const int x, const int last_col)
 {
     return abs(x) % (last_col + 1);
@@ -431,4 +426,4 @@ __kernel void updateFlow(__global const float * M, int mStep,
         flowx[mad24(y, xStep, x)] = (g11*h2 - g12*h1) * detInv;
         flowy[mad24(y, yStep, x)] = (g22*h1 - g12*h2) * detInv;
     }
-}
+}