pythonwebcoder
diff --git a/‎modules/core/include/opencv2/core/ocl.hpp
Lines changed: 16 additions & 0 deletions b/‎modules/core/include/opencv2/core/ocl.hpp
Lines changed: 16 additions & 0 deletions
diff --git a/‎modules/core/include/opencv2/core/utils/configuration.private.hpp
Lines changed: 16 additions & 0 deletions b/‎modules/core/include/opencv2/core/utils/configuration.private.hpp
Lines changed: 16 additions & 0 deletions
diff --git a/‎modules/core/src/ocl.cpp
Lines changed: 270 additions & 0 deletions b/‎modules/core/src/ocl.cpp
Lines changed: 270 additions & 0 deletions
diff --git a/‎modules/core/src/opencl/benchmark.cl
Lines changed: 45 additions & 0 deletions b/‎modules/core/src/opencl/benchmark.cl
Lines changed: 45 additions & 0 deletions
diff --git a/‎modules/core/src/precomp.hpp
Lines changed: 0 additions & 6 deletions b/‎modules/core/src/precomp.hpp
Lines changed: 0 additions & 6 deletions
@@ -665,6 +665,7 @@ CV_EXPORTS const char* convertTypeStr(int sdepth, int ddepth, int cn, char* buf)
 CV_EXPORTS const char* typeToStr(int t);
 CV_EXPORTS const char* memopTypeToStr(int t);
 CV_EXPORTS const char* vecopTypeToStr(int t);
+CV_EXPORTS const char* getOpenCLErrorString(int errorCode);
 CV_EXPORTS String kernelToStr(InputArray _kernel, int ddepth = -1, const char * name = NULL);
 CV_EXPORTS void getPlatfomsInfo(std::vector<PlatformInfo>& platform_info);
 
@@ -731,6 +732,21 @@ class CV_EXPORTS Image2D
     Impl* p;
 };
 
+class CV_EXPORTS Timer
+{
+public:
+    Timer(const Queue& q);
+    ~Timer();
+    void start();
+    void stop();
+    float milliSeconds();
+    float microSeconds();
+    float seconds();
+
+protected:
+    struct Impl;
+    Impl* p;
+};
 
 CV_EXPORTS MatAllocator* getOpenCLAllocator();
 
 
@@ -0,0 +1,16 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+#ifndef OPENCV_CONFIGURATION_PRIVATE_HPP
+#define OPENCV_CONFIGURATION_PRIVATE_HPP
+
+namespace cv { namespace utils {
+
+CV_EXPORTS bool getConfigurationParameterBool(const char* name, bool defaultValue);
+CV_EXPORTS size_t getConfigurationParameterSizeT(const char* name, size_t defaultValue);
+CV_EXPORTS cv::String getConfigurationParameterString(const char* name, const char* defaultValue);
+
+}} // namespace
+
+#endif // OPENCV_CONFIGURATION_PRIVATE_HPP
@@ -51,7 +51,10 @@
 #include <inttypes.h>
 #endif
 
+#include <opencv2/core/utils/configuration.private.hpp>
+
 #include "opencv2/core/ocl_genbase.hpp"
+#include "opencl_kernels_core.hpp"
 
 #define CV_OPENCL_ALWAYS_SHOW_BUILD_LOG 0
 #define CV_OPENCL_SHOW_RUN_ERRORS       0
@@ -4718,6 +4721,102 @@ const char* convertTypeStr(int sdepth, int ddepth, int cn, char* buf)
     return buf;
 }
 
+const char* getOpenCLErrorString(int errorCode)
+{
+    switch (errorCode)
+    {
+    case   0: return "CL_SUCCESS";
+    case  -1: return "CL_DEVICE_NOT_FOUND";
+    case  -2: return "CL_DEVICE_NOT_AVAILABLE";
+    case  -3: return "CL_COMPILER_NOT_AVAILABLE";
+    case  -4: return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
+    case  -5: return "CL_OUT_OF_RESOURCES";
+    case  -6: return "CL_OUT_OF_HOST_MEMORY";
+    case  -7: return "CL_PROFILING_INFO_NOT_AVAILABLE";
+    case  -8: return "CL_MEM_COPY_OVERLAP";
+    case  -9: return "CL_IMAGE_FORMAT_MISMATCH";
+    case -10: return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
+    case -11: return "CL_BUILD_PROGRAM_FAILURE";
+    case -12: return "CL_MAP_FAILURE";
+    case -13: return "CL_MISALIGNED_SUB_BUFFER_OFFSET";
+    case -14: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST";
+    case -15: return "CL_COMPILE_PROGRAM_FAILURE";
+    case -16: return "CL_LINKER_NOT_AVAILABLE";
+    case -17: return "CL_LINK_PROGRAM_FAILURE";
+    case -18: return "CL_DEVICE_PARTITION_FAILED";
+    case -19: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE";
+    case -30: return "CL_INVALID_VALUE";
+    case -31: return "CL_INVALID_DEVICE_TYPE";
+    case -32: return "CL_INVALID_PLATFORM";
+    case -33: return "CL_INVALID_DEVICE";
+    case -34: return "CL_INVALID_CONTEXT";
+    case -35: return "CL_INVALID_QUEUE_PROPERTIES";
+    case -36: return "CL_INVALID_COMMAND_QUEUE";
+    case -37: return "CL_INVALID_HOST_PTR";
+    case -38: return "CL_INVALID_MEM_OBJECT";
+    case -39: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
+    case -40: return "CL_INVALID_IMAGE_SIZE";
+    case -41: return "CL_INVALID_SAMPLER";
+    case -42: return "CL_INVALID_BINARY";
+    case -43: return "CL_INVALID_BUILD_OPTIONS";
+    case -44: return "CL_INVALID_PROGRAM";
+    case -45: return "CL_INVALID_PROGRAM_EXECUTABLE";
+    case -46: return "CL_INVALID_KERNEL_NAME";
+    case -47: return "CL_INVALID_KERNEL_DEFINITION";
+    case -48: return "CL_INVALID_KERNEL";
+    case -49: return "CL_INVALID_ARG_INDEX";
+    case -50: return "CL_INVALID_ARG_VALUE";
+    case -51: return "CL_INVALID_ARG_SIZE";
+    case -52: return "CL_INVALID_KERNEL_ARGS";
+    case -53: return "CL_INVALID_WORK_DIMENSION";
+    case -54: return "CL_INVALID_WORK_GROUP_SIZE";
+    case -55: return "CL_INVALID_WORK_ITEM_SIZE";
+    case -56: return "CL_INVALID_GLOBAL_OFFSET";
+    case -57: return "CL_INVALID_EVENT_WAIT_LIST";
+    case -58: return "CL_INVALID_EVENT";
+    case -59: return "CL_INVALID_OPERATION";
+    case -60: return "CL_INVALID_GL_OBJECT";
+    case -61: return "CL_INVALID_BUFFER_SIZE";
+    case -62: return "CL_INVALID_MIP_LEVEL";
+    case -63: return "CL_INVALID_GLOBAL_WORK_SIZE";
+    case -64: return "CL_INVALID_PROPERTY";
+    case -65: return "CL_INVALID_IMAGE_DESCRIPTOR";
+    case -66: return "CL_INVALID_COMPILER_OPTIONS";
+    case -67: return "CL_INVALID_LINKER_OPTIONS";
+    case -68: return "CL_INVALID_DEVICE_PARTITION_COUNT";
+    case -69: return "CL_INVALID_PIPE_SIZE";
+    case -70: return "CL_INVALID_DEVICE_QUEUE";
+    case -1000: return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR";
+    case -1001: return "CL_PLATFORM_NOT_FOUND_KHR";
+    case -1002: return "CL_INVALID_D3D10_DEVICE_KHR";
+    case -1003: return "CL_INVALID_D3D10_RESOURCE_KHR";
+    case -1004: return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR";
+    case -1005: return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR";
+    case -1024: return "clBLAS: Functionality is not implemented";
+    case -1023: return "clBLAS: Library is not initialized yet";
+    case -1022: return "clBLAS: Matrix A is not a valid memory object";
+    case -1021: return "clBLAS: Matrix B is not a valid memory object";
+    case -1020: return "clBLAS: Matrix C is not a valid memory object";
+    case -1019: return "clBLAS: Vector X is not a valid memory object";
+    case -1018: return "clBLAS: Vector Y is not a valid memory object";
+    case -1017: return "clBLAS: An input dimension (M:N:K) is invalid";
+    case -1016: return "clBLAS: Leading dimension A must not be less than the "
+                       "size of the first dimension";
+    case -1015: return "clBLAS: Leading dimension B must not be less than the "
+                       "size of the second dimension";
+    case -1014: return "clBLAS: Leading dimension C must not be less than the "
+                       "size of the third dimension";
+    case -1013: return "clBLAS: The increment for a vector X must not be 0";
+    case -1012: return "clBLAS: The increment for a vector Y must not be 0";
+    case -1011: return "clBLAS: The memory object for Matrix A is too small";
+    case -1010: return "clBLAS: The memory object for Matrix B is too small";
+    case -1009: return "clBLAS: The memory object for Matrix C is too small";
+    case -1008: return "clBLAS: The memory object for Vector X is too small";
+    case -1007: return "clBLAS: The memory object for Vector Y is too small";
+    default: return "Unknown OpenCL error";
+    }
+}
+
 template <typename T>
 static std::string kerToStr(const Mat & k)
 {
@@ -5134,4 +5233,175 @@ bool internal::isCLBuffer(UMat& u)
     return true;
 }
 
+struct Timer::Impl
+{
+    const Queue queue;
+
+    Impl(const Queue& q)
+        : queue(q)
+        , initted_(false)
+        , running_(false)
+        , has_run_at_least_once_(false)
+    {
+        init();
+    }
+
+    ~Impl()
+    {
+        clWaitForEvents(1, &start_gpu_cl_);
+        clWaitForEvents(1, &stop_gpu_cl_);
+        clReleaseEvent(start_gpu_cl_);
+        clReleaseEvent(stop_gpu_cl_);
+    }
+
+    void start()
+    {
+#ifdef HAVE_OPENCL
+        if (!running())
+        {
+            clWaitForEvents(1, &start_gpu_cl_);
+            clReleaseEvent(start_gpu_cl_);
+            ocl::Kernel kernel("null_kernel_float", ocl::core::benchmark_oclsrc);
+            float arg = 0;
+            clSetKernelArg((cl_kernel)kernel.ptr(), 0, sizeof(arg), &arg);
+            clEnqueueTask((cl_command_queue)queue.ptr(), (cl_kernel)kernel.ptr(), 0,
+                          NULL, &start_gpu_cl_);
+            clFinish((cl_command_queue)queue.ptr());
+            running_ = true;
+            has_run_at_least_once_ = true;
+        }
+#endif
+    }
+
+    void stop()
+    {
+#ifdef HAVE_OPENCL
+        if (running())
+        {
+            clWaitForEvents(1, &stop_gpu_cl_);
+            clReleaseEvent(stop_gpu_cl_);
+            ocl::Kernel kernel("null_kernel_float", ocl::core::benchmark_oclsrc);
+            float arg = 0;
+            clSetKernelArg((cl_kernel)kernel.ptr(), 0, sizeof(arg), &arg);
+            clEnqueueTask((cl_command_queue)queue.ptr(), (cl_kernel)kernel.ptr(), 0,
+                          NULL, &stop_gpu_cl_);
+            clFinish((cl_command_queue)queue.ptr());
+            running_ = false;
+        }
+#endif
+    }
+
+    float microSeconds()
+    {
+#ifdef HAVE_OPENCL
+        if (!has_run_at_least_once())
+        {
+            return 0;
+        }
+        if (running())
+        {
+            stop();
+        }
+        cl_ulong startTime, stopTime;
+        clWaitForEvents(1, &stop_gpu_cl_);
+        clGetEventProfilingInfo(start_gpu_cl_, CL_PROFILING_COMMAND_END,
+                                sizeof startTime, &startTime, NULL);
+        clGetEventProfilingInfo(stop_gpu_cl_, CL_PROFILING_COMMAND_START,
+                                sizeof stopTime, &stopTime, NULL);
+        double us = static_cast<double>(stopTime - startTime) / 1000.0;
+        elapsed_microseconds_ = static_cast<float>(us);
+        return elapsed_microseconds_;
+#else
+        return 0;
+#endif
+    }
+
+    float milliSeconds()
+    {
+#ifdef HAVE_OPENCL
+        if (!has_run_at_least_once())
+        {
+            return 0;
+        }
+        if (running())
+        {
+            stop();
+        }
+        cl_ulong startTime = 0, stopTime = 0;
+        clGetEventProfilingInfo(start_gpu_cl_, CL_PROFILING_COMMAND_END,
+                                sizeof startTime, &startTime, NULL);
+        clGetEventProfilingInfo(stop_gpu_cl_, CL_PROFILING_COMMAND_START,
+                                sizeof stopTime, &stopTime, NULL);
+        double ms = static_cast<double>(stopTime - startTime) / 1000000.0;
+        elapsed_milliseconds_ = static_cast<float>(ms);
+        return elapsed_milliseconds_;
+#else
+        return 0;
+#endif
+    }
+
+    float seconds()
+    {
+        return milliSeconds() / 1000.f;
+    }
+
+    void init()
+    {
+        CV_Assert(queue.getImpl() && queue.getImpl()->isProfilingQueue_);
+        if (!initted())
+        {
+            start_gpu_cl_ = 0;
+            stop_gpu_cl_ = 0;
+            initted_ = true;
+        }
+    }
+
+    inline bool initted() { return initted_; }
+    inline bool running() { return running_; }
+    inline bool has_run_at_least_once() { return has_run_at_least_once_; }
+
+    bool initted_;
+    bool running_;
+    bool has_run_at_least_once_;
+    float elapsed_milliseconds_;
+    float elapsed_microseconds_;
+    cl_event start_gpu_cl_;
+    cl_event stop_gpu_cl_;
+};
+
+Timer::Timer(const Queue& q)
+{
+    p = new Impl(q);
+}
+
+Timer::~Timer()
+{
+    if(p)
+    {
+        delete p;
+        p = 0;
+    }
+}
+
+void Timer::start()
+{
+    if(p)
+        p->start();
+}
+
+void Timer::stop()
+{
+    if(p)
+        p->stop();
+}
+
+float Timer::microSeconds()
+{ return p ? p->microSeconds() : 0; }
+
+float Timer::milliSeconds()
+{ return p ? p->milliSeconds() : 0; }
+
+float Timer::seconds()
+{ return p ? p->seconds() : 0; }
+
 }}
@@ -0,0 +1,45 @@
+/*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) 2017, Intel Corporation, all rights reserved.
+// Copyright (c) 2016-2017 Fabian David Tschopp, 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*/
+
+__kernel void null_kernel_float(float arg) {
+  float out = arg;
+}
@@ -297,12 +297,6 @@ TLSData<CoreTLSData>& getCoreTlsData();
 #define CL_RUNTIME_EXPORT
 #endif
 
-namespace utils {
-bool getConfigurationParameterBool(const char* name, bool defaultValue);
-size_t getConfigurationParameterSizeT(const char* name, size_t defaultValue);
-cv::String getConfigurationParameterString(const char* name, const char* defaultValue);
-}
-
 extern bool __termination; // skip some cleanups, because process is terminating
                            // (for example, if ExitProcess() was already called)