ocl: simplify ocl::Timer

Use clFinish to gurantee commands completed, instead of waiting for events.

Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com>

diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp
index 60dca8b..a432c8d 100644 (file)
--- a/modules/core/src/ocl.cpp
+++ b/modules/core/src/ocl.cpp
@@ -5239,78 +5239,31 @@ struct Timer::Impl
 
     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_);
-    }
+    ~Impl(){}
 
     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;
-        }
+        clFinish((cl_command_queue)queue.ptr());
+        timer.start();
 #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;
-        }
+        clFinish((cl_command_queue)queue.ptr());
+        timer.stop();
 #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_;
+        return (float)timer.getTimeMicro();
 #else
         return 0;
 #endif
@@ -5319,22 +5272,7 @@ struct Timer::Impl
     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_;
+        return (float)timer.getTimeMilli();
 #else
         return 0;
 #endif
@@ -5342,31 +5280,13 @@ struct Timer::Impl
 
     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;
-        }
+#ifdef HAVE_OPENCL
+        return (float)timer.getTimeSec();
+#else
+        return 0;
+#endif
     }
-
-    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_;
+    TickMeter timer;
 };
 
 Timer::Timer(const Queue& q)

diff --git a/modules/core/src/opencl/benchmark.cl b/modules/core/src/opencl/benchmark.cl
deleted file mode 100644 (file)
index 22acb93..0000000
--- a/modules/core/src/opencl/benchmark.cl
+++ /dev/null
@@ -1,45 +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) 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;
-}

diff --git a/modules/dnn/src/opencl/benchmark.cl b/modules/dnn/src/opencl/benchmark.cl
deleted file mode 100644 (file)
index 22acb93..0000000
--- a/modules/dnn/src/opencl/benchmark.cl
+++ /dev/null
@@ -1,45 +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) 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;
-}