From 4157ea0a7281c995d490a302f4f6eedc63b44624 Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Mon, 25 Apr 2022 14:27:44 +0300 Subject: [PATCH 1/9] Enable profiling only when it is used explicitly --- src/runtime/opencl/opencl_common.h | 37 +++++++++++++++++++------ src/runtime/opencl/opencl_device_api.cc | 5 ---- src/runtime/opencl/opencl_module.cc | 18 ++++++------ 3 files changed, 38 insertions(+), 22 deletions(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index 18061a7aeeb5..f9fb70503fdc 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -229,6 +229,8 @@ class OpenCLWorkspace : public DeviceAPI { cl_context context{nullptr}; // whether the workspace it initialized. bool initialized_{false}; + // whether the workspace is in profiling mode. + bool profiling{false}; // the device type std::string device_type; // the devices @@ -422,19 +424,38 @@ class OpenCLTimerNode : public TimerNode { virtual void Start() { cl::OpenCLWorkspace::Global()->GetEventQueue(dev_).clear(); this->duration = 0; + + if (!cl::OpenCLWorkspace::Global()->profiling) { + // Very first call of Start() leads to the recreation of + // OpenCL command queue in profiling mode + OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + OPENCL_CALL(clFinish(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + OPENCL_CALL(clReleaseCommandQueue(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + + cl_int err_code; + cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; + auto profiling_queue = clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, + did, + CL_QUEUE_PROFILING_ENABLE, + &err_code); + OPENCL_CHECK_ERROR(err_code); + cl::OpenCLWorkspace::Global()->queues[dev_.device_id] = profiling_queue; + cl::OpenCLWorkspace::Global()->profiling = true; + } } // Timer stop virtual void Stop() { std::vector evt_queue = cl::OpenCLWorkspace::Global()->GetEventQueue(dev_); cl_ulong start, end; - OPENCL_CALL(clWaitForEvents(1, &(cl::OpenCLWorkspace::Global()->GetEventQueue(dev_).back()))); - for (auto& kevt : evt_queue) { - OPENCL_CALL(clGetEventProfilingInfo(kevt, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), - &start, nullptr)); - OPENCL_CALL( - clGetEventProfilingInfo(kevt, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, nullptr)); - this->duration += (end - start); - } + if (cl::OpenCLWorkspace::Global()->GetEventQueue(dev_).size() > 0) { + OPENCL_CALL(clWaitForEvents(1, &(cl::OpenCLWorkspace::Global()->GetEventQueue(dev_).back()))); + for (auto& kevt : evt_queue) { + OPENCL_CALL(clGetEventProfilingInfo(kevt, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), + &start, nullptr)); + OPENCL_CALL(clGetEventProfilingInfo(kevt, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, + nullptr)); + this->duration += (end - start); + } } virtual int64_t SyncAndGetElapsedNanos() { return this->duration; } // destructor diff --git a/src/runtime/opencl/opencl_device_api.cc b/src/runtime/opencl/opencl_device_api.cc index c3527160429f..80b95a6ebfe9 100644 --- a/src/runtime/opencl/opencl_device_api.cc +++ b/src/runtime/opencl/opencl_device_api.cc @@ -426,12 +426,7 @@ void OpenCLWorkspace::Init(const std::string& type_key, const std::string& devic ICHECK_EQ(this->queues.size(), 0U); for (size_t i = 0; i < this->devices.size(); ++i) { cl_device_id did = this->devices[i]; -#ifdef USE_PROFILER - this->queues.push_back( - clCreateCommandQueue(this->context, did, CL_QUEUE_PROFILING_ENABLE, &err_code)); -#else this->queues.push_back(clCreateCommandQueue(this->context, did, 0, &err_code)); -#endif OPENCL_CHECK_ERROR(err_code); } this->events.resize(this->devices.size()); diff --git a/src/runtime/opencl/opencl_module.cc b/src/runtime/opencl/opencl_module.cc index e08c6070bc88..641ddc702c85 100644 --- a/src/runtime/opencl/opencl_module.cc +++ b/src/runtime/opencl/opencl_module.cc @@ -79,15 +79,15 @@ class OpenCLWrappedFunc { wl.work_size[i] *= wl.work_size[i + 3]; } // launch kernel -#ifdef USE_PROFILER - w_->GetEventQueue(t->device).resize(w_->GetEventQueue(t->device).size() + 1); - OPENCL_CALL(clEnqueueNDRangeKernel(queue, kernel, work_dim, nullptr, wl.work_size, - wl.work_size + 3, 0, nullptr, - &(w_->GetEventQueue(t->device).back()))); -#else - OPENCL_CALL(clEnqueueNDRangeKernel(queue, kernel, work_dim, nullptr, wl.work_size, - wl.work_size + 3, 0, nullptr, nullptr)); -#endif + if (w_->profiling) { + w_->GetEventQueue(t->device).resize(w_->GetEventQueue(t->device).size() + 1); + OPENCL_CALL(clEnqueueNDRangeKernel(queue, kernel, work_dim, nullptr, wl.work_size, + wl.work_size + 3, 0, nullptr, + &(w_->GetEventQueue(t->device).back()))); + } else { + OPENCL_CALL(clEnqueueNDRangeKernel(queue, kernel, work_dim, nullptr, wl.work_size, + wl.work_size + 3, 0, nullptr, nullptr)); + } } private: From 5619902c5adf3e3ce2daf28a67f429ea78f22f61 Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Wed, 27 Apr 2022 14:45:44 +0300 Subject: [PATCH 2/9] Change logic of clCommandQueue create/destroy --- src/runtime/opencl/opencl_common.h | 20 +++++++++++++++++++- 1 file changed, 19 insertions(+), 1 deletion(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index f9fb70503fdc..13a38a6bf3f6 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -456,10 +456,28 @@ class OpenCLTimerNode : public TimerNode { nullptr)); this->duration += (end - start); } + } } virtual int64_t SyncAndGetElapsedNanos() { return this->duration; } // destructor - virtual ~OpenCLTimerNode() {} + virtual ~OpenCLTimerNode() { + if (cl::OpenCLWorkspace::Global()->profiling) { + // Profiling session ends, recreate clCommandQueue + OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + OPENCL_CALL(clFinish(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + OPENCL_CALL(clReleaseCommandQueue(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + + cl_int err_code; + cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; + auto normal_queue = clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, + did, + 0, + &err_code); + OPENCL_CHECK_ERROR(err_code); + cl::OpenCLWorkspace::Global()->queues[dev_.device_id] = normal_queue; + cl::OpenCLWorkspace::Global()->profiling = false; + } + } // constructor OpenCLTimerNode() {} explicit OpenCLTimerNode(Device dev) : dev_(dev) {} From 4d22c0602b2f2d0d38a35a47d8ea0e24cbd8d0de Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Fri, 29 Apr 2022 13:49:45 +0300 Subject: [PATCH 3/9] Update comments --- src/runtime/opencl/opencl_common.h | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index 13a38a6bf3f6..14ff673382b3 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -427,7 +427,7 @@ class OpenCLTimerNode : public TimerNode { if (!cl::OpenCLWorkspace::Global()->profiling) { // Very first call of Start() leads to the recreation of - // OpenCL command queue in profiling mode + // OpenCL command queue in profiling mode. This allows to run profile after inference. OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); OPENCL_CALL(clFinish(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); OPENCL_CALL(clReleaseCommandQueue(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); @@ -462,7 +462,8 @@ class OpenCLTimerNode : public TimerNode { // destructor virtual ~OpenCLTimerNode() { if (cl::OpenCLWorkspace::Global()->profiling) { - // Profiling session ends, recreate clCommandQueue + // Profiling session ends, recreate clCommandQueue in non-profiling mode + // This will disable collection of cl_events in case of executing inference after profile OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); OPENCL_CALL(clFinish(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); OPENCL_CALL(clReleaseCommandQueue(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); From 82462619cfd5e3e5d8679f4b358a9b4a94fbe264 Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Fri, 29 Apr 2022 15:11:49 +0300 Subject: [PATCH 4/9] Linter fix --- src/runtime/opencl/opencl_common.h | 12 ++++-------- 1 file changed, 4 insertions(+), 8 deletions(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index 14ff673382b3..cda57605c8c3 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -434,10 +434,8 @@ class OpenCLTimerNode : public TimerNode { cl_int err_code; cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; - auto profiling_queue = clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, - did, - CL_QUEUE_PROFILING_ENABLE, - &err_code); + auto profiling_queue = clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, did, + CL_QUEUE_PROFILING_ENABLE, &err_code); OPENCL_CHECK_ERROR(err_code); cl::OpenCLWorkspace::Global()->queues[dev_.device_id] = profiling_queue; cl::OpenCLWorkspace::Global()->profiling = true; @@ -470,10 +468,8 @@ class OpenCLTimerNode : public TimerNode { cl_int err_code; cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; - auto normal_queue = clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, - did, - 0, - &err_code); + auto normal_queue = + clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, did, 0, &err_code); OPENCL_CHECK_ERROR(err_code); cl::OpenCLWorkspace::Global()->queues[dev_.device_id] = normal_queue; cl::OpenCLWorkspace::Global()->profiling = false; From d94c6d019a3d22d263ea4f0c0f81805ee8713726 Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Thu, 5 May 2022 13:00:22 +0300 Subject: [PATCH 5/9] Refactor queue create --- src/runtime/opencl/opencl_common.h | 45 +++++++++++++++--------------- 1 file changed, 23 insertions(+), 22 deletions(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index cda57605c8c3..a2c41960e124 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -428,17 +428,7 @@ class OpenCLTimerNode : public TimerNode { if (!cl::OpenCLWorkspace::Global()->profiling) { // Very first call of Start() leads to the recreation of // OpenCL command queue in profiling mode. This allows to run profile after inference. - OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); - OPENCL_CALL(clFinish(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); - OPENCL_CALL(clReleaseCommandQueue(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); - - cl_int err_code; - cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; - auto profiling_queue = clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, did, - CL_QUEUE_PROFILING_ENABLE, &err_code); - OPENCL_CHECK_ERROR(err_code); - cl::OpenCLWorkspace::Global()->queues[dev_.device_id] = profiling_queue; - cl::OpenCLWorkspace::Global()->profiling = true; + recreateCommandQueue(true); } } // Timer stop @@ -462,17 +452,7 @@ class OpenCLTimerNode : public TimerNode { if (cl::OpenCLWorkspace::Global()->profiling) { // Profiling session ends, recreate clCommandQueue in non-profiling mode // This will disable collection of cl_events in case of executing inference after profile - OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); - OPENCL_CALL(clFinish(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); - OPENCL_CALL(clReleaseCommandQueue(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); - - cl_int err_code; - cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; - auto normal_queue = - clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, did, 0, &err_code); - OPENCL_CHECK_ERROR(err_code); - cl::OpenCLWorkspace::Global()->queues[dev_.device_id] = normal_queue; - cl::OpenCLWorkspace::Global()->profiling = false; + recreateCommandQueue(false); } } // constructor @@ -485,6 +465,27 @@ class OpenCLTimerNode : public TimerNode { private: int64_t duration; Device dev_; + + void recreateCommandQueue(bool profiling) { + cl_command_queue_properties prop; + if (profiling) { + prop = CL_QUEUE_PROFILING_ENABLE; + } else { + prop = CL_QUEUE_PROFILING_ENABLE; + } + + OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + OPENCL_CALL(clFinish(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + OPENCL_CALL(clReleaseCommandQueue(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + + cl_int err_code; + cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; + auto profiling_queue = clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, did, + prop, &err_code); + OPENCL_CHECK_ERROR(err_code); + cl::OpenCLWorkspace::Global()->queues[dev_.device_id] = profiling_queue; + cl::OpenCLWorkspace::Global()->profiling = profiling; + } }; } // namespace runtime } // namespace tvm From 5a34acb8b1d08ac83f8dd71abba19f1ad04d66a9 Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Thu, 5 May 2022 22:40:26 +0300 Subject: [PATCH 6/9] Move queue recreation logic to function --- src/runtime/opencl/opencl_common.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index a2c41960e124..7f231edea501 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -471,7 +471,7 @@ class OpenCLTimerNode : public TimerNode { if (profiling) { prop = CL_QUEUE_PROFILING_ENABLE; } else { - prop = CL_QUEUE_PROFILING_ENABLE; + prop = 0; } OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); From eadcec206fb56a223620dfae5f8c6e6f554601a8 Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Fri, 6 May 2022 12:38:24 +0300 Subject: [PATCH 7/9] Replace profiling flag by the queue info request --- src/runtime/opencl/opencl_common.h | 25 ++++++++++++++++++------- src/runtime/opencl/opencl_module.cc | 3 ++- 2 files changed, 20 insertions(+), 8 deletions(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index 7f231edea501..07bb9ceec44d 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -229,8 +229,6 @@ class OpenCLWorkspace : public DeviceAPI { cl_context context{nullptr}; // whether the workspace it initialized. bool initialized_{false}; - // whether the workspace is in profiling mode. - bool profiling{false}; // the device type std::string device_type; // the devices @@ -276,6 +274,20 @@ class OpenCLWorkspace : public DeviceAPI { << "Invalid OpenCL device_id=" << dev.device_id; return events[dev.device_id]; } + // is current clCommandQueue in profiling mode + bool IsProfiling(Device dev) { + ICHECK(IsOpenCLDevice(dev)); + this->Init(); + ICHECK(dev.device_id >= 0 && static_cast(dev.device_id) < queues.size()) + << "Invalid OpenCL device_id=" << dev.device_id; + cl_command_queue queue = queues[dev.device_id]; + cl_command_queue_properties prop; + + OPENCL_CALL(clGetCommandQueueInfo(queue, CL_QUEUE_PROPERTIES, + sizeof(cl_command_queue_properties), &prop, nullptr)); + + return prop & CL_QUEUE_PROFILING_ENABLE; + } // override device API void SetDevice(Device dev) final; @@ -425,7 +437,7 @@ class OpenCLTimerNode : public TimerNode { cl::OpenCLWorkspace::Global()->GetEventQueue(dev_).clear(); this->duration = 0; - if (!cl::OpenCLWorkspace::Global()->profiling) { + if (!cl::OpenCLWorkspace::Global()->IsProfiling(dev_)) { // Very first call of Start() leads to the recreation of // OpenCL command queue in profiling mode. This allows to run profile after inference. recreateCommandQueue(true); @@ -449,7 +461,7 @@ class OpenCLTimerNode : public TimerNode { virtual int64_t SyncAndGetElapsedNanos() { return this->duration; } // destructor virtual ~OpenCLTimerNode() { - if (cl::OpenCLWorkspace::Global()->profiling) { + if (cl::OpenCLWorkspace::Global()->IsProfiling(dev_)) { // Profiling session ends, recreate clCommandQueue in non-profiling mode // This will disable collection of cl_events in case of executing inference after profile recreateCommandQueue(false); @@ -480,11 +492,10 @@ class OpenCLTimerNode : public TimerNode { cl_int err_code; cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; - auto profiling_queue = clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, did, - prop, &err_code); + auto profiling_queue = + clCreateCommandQueue(cl::OpenCLWorkspace::Global()->context, did, prop, &err_code); OPENCL_CHECK_ERROR(err_code); cl::OpenCLWorkspace::Global()->queues[dev_.device_id] = profiling_queue; - cl::OpenCLWorkspace::Global()->profiling = profiling; } }; } // namespace runtime diff --git a/src/runtime/opencl/opencl_module.cc b/src/runtime/opencl/opencl_module.cc index 641ddc702c85..9ae80d59d565 100644 --- a/src/runtime/opencl/opencl_module.cc +++ b/src/runtime/opencl/opencl_module.cc @@ -79,7 +79,8 @@ class OpenCLWrappedFunc { wl.work_size[i] *= wl.work_size[i + 3]; } // launch kernel - if (w_->profiling) { + + if (w_->IsProfiling(t->device)) { w_->GetEventQueue(t->device).resize(w_->GetEventQueue(t->device).size() + 1); OPENCL_CALL(clEnqueueNDRangeKernel(queue, kernel, work_dim, nullptr, wl.work_size, wl.work_size + 3, 0, nullptr, From 0616943a540beb62afee45fc6c8c3718f1541603 Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Sun, 8 May 2022 20:43:26 +0300 Subject: [PATCH 8/9] Enhance readability --- src/runtime/opencl/opencl_common.h | 31 ++++++++++++------------------ 1 file changed, 12 insertions(+), 19 deletions(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index 07bb9ceec44d..e04ed3003b7f 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -276,11 +276,7 @@ class OpenCLWorkspace : public DeviceAPI { } // is current clCommandQueue in profiling mode bool IsProfiling(Device dev) { - ICHECK(IsOpenCLDevice(dev)); - this->Init(); - ICHECK(dev.device_id >= 0 && static_cast(dev.device_id) < queues.size()) - << "Invalid OpenCL device_id=" << dev.device_id; - cl_command_queue queue = queues[dev.device_id]; + cl_command_queue queue = GetQueue(dev); cl_command_queue_properties prop; OPENCL_CALL(clGetCommandQueueInfo(queue, CL_QUEUE_PROPERTIES, @@ -436,12 +432,9 @@ class OpenCLTimerNode : public TimerNode { virtual void Start() { cl::OpenCLWorkspace::Global()->GetEventQueue(dev_).clear(); this->duration = 0; - - if (!cl::OpenCLWorkspace::Global()->IsProfiling(dev_)) { - // Very first call of Start() leads to the recreation of - // OpenCL command queue in profiling mode. This allows to run profile after inference. - recreateCommandQueue(true); - } + // Very first call of Start() leads to the recreation of + // OpenCL command queue in profiling mode. This allows to run profile after inference. + recreateCommandQueue(); } // Timer stop virtual void Stop() { @@ -461,11 +454,9 @@ class OpenCLTimerNode : public TimerNode { virtual int64_t SyncAndGetElapsedNanos() { return this->duration; } // destructor virtual ~OpenCLTimerNode() { - if (cl::OpenCLWorkspace::Global()->IsProfiling(dev_)) { // Profiling session ends, recreate clCommandQueue in non-profiling mode // This will disable collection of cl_events in case of executing inference after profile - recreateCommandQueue(false); - } + recreateCommandQueue(); } // constructor OpenCLTimerNode() {} @@ -478,17 +469,19 @@ class OpenCLTimerNode : public TimerNode { int64_t duration; Device dev_; - void recreateCommandQueue(bool profiling) { + void recreateCommandQueue() { cl_command_queue_properties prop; - if (profiling) { + if (!cl::OpenCLWorkspace::Global()->IsProfiling(dev_)) { prop = CL_QUEUE_PROFILING_ENABLE; } else { prop = 0; } - OPENCL_CALL(clFlush(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); - OPENCL_CALL(clFinish(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); - OPENCL_CALL(clReleaseCommandQueue(cl::OpenCLWorkspace::Global()->GetQueue(dev_))); + auto queue = cl::OpenCLWorkspace::Global()->GetQueue(dev_); + + OPENCL_CALL(clFlush(queue)); + OPENCL_CALL(clFinish(queue)); + OPENCL_CALL(clReleaseCommandQueue(queue)); cl_int err_code; cl_device_id did = cl::OpenCLWorkspace::Global()->devices[dev_.device_id]; From e3bd9307be8ae58f4412f2008346f8cb9cebce9b Mon Sep 17 00:00:00 2001 From: Kirill Snezhko <4477094+argrento@users.noreply.github.com> Date: Mon, 9 May 2022 14:19:16 +0300 Subject: [PATCH 9/9] Fix linter errors --- src/runtime/opencl/opencl_common.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/runtime/opencl/opencl_common.h b/src/runtime/opencl/opencl_common.h index e04ed3003b7f..c2905b432764 100644 --- a/src/runtime/opencl/opencl_common.h +++ b/src/runtime/opencl/opencl_common.h @@ -454,9 +454,9 @@ class OpenCLTimerNode : public TimerNode { virtual int64_t SyncAndGetElapsedNanos() { return this->duration; } // destructor virtual ~OpenCLTimerNode() { - // Profiling session ends, recreate clCommandQueue in non-profiling mode - // This will disable collection of cl_events in case of executing inference after profile - recreateCommandQueue(); + // Profiling session ends, recreate clCommandQueue in non-profiling mode + // This will disable collection of cl_events in case of executing inference after profile + recreateCommandQueue(); } // constructor OpenCLTimerNode() {}