From dbe9ee09247cf5dddde4363b53d7875f2f20dd86 Mon Sep 17 00:00:00 2001 From: Wu Zhiwen Date: Mon, 2 Oct 2017 19:22:28 +0800 Subject: [PATCH] ocl: simplify ocl::Timer Use clFinish to gurantee commands completed, instead of waiting for events. Signed-off-by: Wu Zhiwen --- modules/core/src/ocl.cpp | 106 +++++------------------------------ modules/core/src/opencl/benchmark.cl | 45 --------------- modules/dnn/src/opencl/benchmark.cl | 45 --------------- 3 files changed, 13 insertions(+), 183 deletions(-) delete mode 100644 modules/core/src/opencl/benchmark.cl delete mode 100644 modules/dnn/src/opencl/benchmark.cl diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 60dca8b..a432c8d 100644 --- 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(stopTime - startTime) / 1000.0; - elapsed_microseconds_ = static_cast(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(stopTime - startTime) / 1000000.0; - elapsed_milliseconds_ = static_cast(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 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 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; -} -- 2.7.4