[OpenCL] Change of OpenCL profiling logic

[argrento]

Profiling in TVM is enabled or disable in compile time by the USE_PROFILER switch. It means that if we enable profile in config.cmake, but do not use any profiling features in the app, OpenCL is forced to collect cl_events objects.

Build TVM with set(USE_PROFILER ON).
Consider simple app, where we create module from the .so file:

tvm::runtime::Module mod_factory = tvm::runtime::Module::LoadFromFile("model.so");
tvm::runtime::Module gmod = mod_factory.GetFunction("default")(ctx);
tvm::runtime::PackedFunc set_input = gmod.GetFunction("set_input");
tvm::runtime::PackedFunc get_input = gmod.GetFunction("get_input");
tvm::runtime::PackedFunc get_output = gmod.GetFunction("get_output");
tvm::runtime::PackedFunc run = gmod.GetFunction("run");

// set inputs and outputs

size_t niterations = 5000;
for (size_t i = 0; i < niterations; i++) {
  run();
}

Then we collect memory usage info with Valgrind.

    MB
818.5^                                                                       #
     |                                                                    @@@#
     |                                                                @@@@@@@#
     |                                                            @@@@@@@@@@@#
     |                                                         @@@@@@@@@@@@@@#
     |                                                     ::::@@@@@@@@@@@@@@#
     |                                                  @:::: :@@@@@@@@@@@@@@#
     |                                             @@@:@@:::: :@@@@@@@@@@@@@@#
     |                                           @@@ @:@@:::: :@@@@@@@@@@@@@@#
     |                                      :@@@@@@@ @:@@:::: :@@@@@@@@@@@@@@#
     |                                  :@@::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     |                               ::::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     |                           :@@@: ::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     |                       :::@:@ @: ::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     |                   ::@@: :@:@ @: ::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     |                @@@: @@: :@:@ @: ::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     |           :::::@@ : @@: :@:@ @: ::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     | @@:::::::@:: ::@@ : @@: :@:@ @: ::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     | @ :::::::@:: ::@@ : @@: :@:@ @: ::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
     | @ :::::::@:: ::@@ : @@: :@:@ @: ::@ ::@@@ @@@ @:@@:::: :@@@@@@@@@@@@@@#
   0 +----------------------------------------------------------------------->Gi
     0                                                                   75.95

We do not use any profiling info, but it is collected implicitly because of the compile-time switches:

1. tvm/src/runtime/opencl/opencl_device_api.cc

   Line 431 in 6babb89

   clCreateCommandQueue(this->context, did, CL_QUEUE_PROFILING_ENABLE, &err_code));

2. tvm/src/runtime/opencl/opencl_module.cc

   Line 84 in 6babb89

   OPENCL_CALL(clEnqueueNDRangeKernel(queue, kernel, work_dim, nullptr, wl.work_size,

With the proposed modifications this behavior is changed: clCommandQueue by default is created in the normal mode and is recreated with profiling capabilities when user calls profiler explicitly. When a profiling session is finished, the queue is recreated again in normal mode, which allows to mix profile() calls and run() calls.

With the proposed changes valgrind shows no abnormal memory usage for the example above.

    MB
148.9^#                                                                       
     |#::::::::::::::::::::::::::@:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
     |#: ::::::::::::::: ::::::: @:::::::::::::@:::@:::::@::::::@:::::@:::::@:
   0 +----------------------------------------------------------------------->Gi
     0                                                                   83.07

[echuraev]

In general LGTM. Thanks

[echuraev]

@masahi could you please take a look at this PR?

[srkreddy1238]

@argrento

This PR causes CLML profiling failure. Reason explained below

In general the workspaces can be accessed and shared via “device_api.opencl”. CLML integration shares the workspace created by default OpenCL and it has a reference to the command queue. Changing the command queue in between makes them invalid.

Why do we enable profiler in compilation when we don't want to profile any thing ?
Or
Is there any case where we enable profiling some thing else but not OpenCL / dynamically enable/disable profiling? In such cases we could think of having different compilation flag for OpenCL or environment varible to control at runtime launch.

At any case dynamically recreating the queue will cause issues for other components.

@masahi & @valmat07 comment pls.