Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[OpenCL] Change of OpenCL profiling logic #11180

Merged
merged 9 commits into from
May 10, 2022
Merged

[OpenCL] Change of OpenCL profiling logic #11180

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
4157ea0
Enable profiling only when it is used explicitly
argrento Apr 25, 2022
5619902
Change logic of clCommandQueue create/destroy
argrento Apr 27, 2022
4d22c06
Update comments
argrento Apr 29, 2022
8246261
Linter fix
argrento Apr 29, 2022
d94c6d0
Refactor queue create
argrento May 5, 2022
5a34acb
Move queue recreation logic to function
argrento May 5, 2022
eadcec2
Replace profiling flag by the queue info request
argrento May 6, 2022
0616943
Enhance readability
argrento May 8, 2022
e3bd930
Fix linter errors
argrento May 9, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
52 changes: 44 additions & 8 deletions src/runtime/opencl/opencl_common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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};
argrento marked this conversation as resolved.
Show resolved Hide resolved
// the device type
std::string device_type;
// the devices
Expand Down Expand Up @@ -422,23 +424,57 @@ 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. 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;
}
argrento marked this conversation as resolved.
Show resolved Hide resolved
}
// Timer stop
virtual void Stop() {
std::vector<cl_event> 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
virtual ~OpenCLTimerNode() {}
virtual ~OpenCLTimerNode() {
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;
}
}
// constructor
OpenCLTimerNode() {}
explicit OpenCLTimerNode(Device dev) : dev_(dev) {}
Expand Down
5 changes: 0 additions & 5 deletions src/runtime/opencl/opencl_device_api.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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());
Expand Down
18 changes: 9 additions & 9 deletions src/runtime/opencl/opencl_module.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down