Redesign time correlation in oneprof, improve DPC++ kernel demangling

VERSION

@@ -1 +1 @@
-0.42.1
+0.42.4

build_utils/CMakeLists.txt

@@ -166,6 +166,7 @@ macro(FindIGALibrary TARGET)
         "IGA is not found. "
         "You may need to install Intel(R) Processor Graphics Driver to fix this issue.")
     else()
+      list(GET IGA_DLL_PATH 0 IGA_DLL_PATH)
       message(STATUS
         "IGA is found at ${IGA_DLL_PATH}")
       find_library(IGA_LIB_PATH

loader/loader.cc

@@ -63,12 +63,10 @@ static bool CheckBitness(HANDLE parent, HANDLE child) {
 int main(int argc, char* argv[]) {
   std::string library_file_name = GetLibFileName();
   std::string executable_path = utils::GetExecutablePath();
-  std::string library_file_path = executable_path + library_file_name;
 
+  std::string library_file_path = executable_path + library_file_name;
   if (!IsFileExists(library_file_path.c_str())) {
-    std::cout << "[ERROR] Failed to find " <<
-      library_file_name << " near the loader" << std::endl;
-    return 0;
+    library_file_path = library_file_name;
   }
 
   SharedLibrary* lib = SharedLibrary::Create(library_file_path);

tests/run.py

@@ -35,6 +35,7 @@
          ["sysmon", "-p", "-l", "-d"],
          ["onetrace",
           "-c", "-h", "-d", "-v", "-t", "-s",
+          "--demangle",
           "--kernels-per-tile",
           "--chrome-call-logging",
           "--chrome-device-timeline",
@@ -44,6 +45,7 @@
           "cl", "ze", "omp"],
          ["cl_tracer",
           "-c", "-h", "-d", "-v", "-t", "-s",
+          "--demangle",
           "--chrome-call-logging",
           "--chrome-device-timeline",
           "--chrome-kernel-timeline",
@@ -52,6 +54,7 @@
           "gpu", "dpc", "omp"],
          ["ze_tracer",
           "-c", "-h", "-d", "-v", "-t", "-s",
+          "--demangle",
           "--kernels-per-tile",
           "--chrome-call-logging",
           "--chrome-device-timeline",

tests/tools/cl_tracer.py

@@ -51,7 +51,7 @@ def run(path, option):
     if option == "gpu":
       command = [file_name_prefix + "cl_tracer" + file_extention,\
         "-h", "-d", "-t", app_file, "gpu", "1024", "1"]
-    elif option == "-v" or option == "--conditional-collection":
+    elif option == "-v" or option == "--demangle" or option == "--conditional-collection":
       command = [file_name_prefix + "cl_tracer" + file_extention,\
         "-d", option, app_file, "cpu", "1024", "1"]
     else:
@@ -102,6 +102,8 @@ def main(option):
     option = "-s"
   if len(sys.argv) > 1 and sys.argv[1] == "-v":
     option = "-v"
+  if len(sys.argv) > 1 and sys.argv[1] == "--demangle":
+    option = "--demangle"
   if len(sys.argv) > 1 and sys.argv[1] == "--chrome-call-logging":
     option = "--chrome-call-logging"
   if len(sys.argv) > 1 and sys.argv[1] == "--chrome-device-timeline":

tests/tools/onetrace.py

@@ -35,7 +35,8 @@ def run(path, option):
     app_folder = utils.get_sample_executable_path("omp_gemm")
     app_file = os.path.join(app_folder, "omp_gemm")
     command = ["./onetrace", "-h", "-d", "-t", app_file, "gpu", "1024", "1"]
-  elif option == "-v" or option == "--kernels-per-tile" or option == "--conditional-collection":
+  elif option == "-v" or option == "--demangle" or\
+       option == "--kernels-per-tile" or option == "--conditional-collection":
     app_folder = utils.get_sample_executable_path("dpc_gemm")
     app_file = os.path.join(app_folder, "dpc_gemm")
     command = ["./onetrace", "-d", option, app_file, "gpu", "1024", "1"]
@@ -88,6 +89,8 @@ def main(option):
     option = "-s"
   if len(sys.argv) > 1 and sys.argv[1] == "-v":
     option = "-v"
+  if len(sys.argv) > 1 and sys.argv[1] == "--demangle":
+    option = "--demangle"
   if len(sys.argv) > 1 and sys.argv[1] == "--kernels-per-tile":
     option = "--kernels-per-tile"
   if len(sys.argv) > 1 and sys.argv[1] == "--chrome-call-logging":

tests/tools/ze_tracer.py

@@ -30,7 +30,8 @@ def run(path, option):
     app_folder = utils.get_sample_executable_path("omp_gemm")
     app_file = os.path.join(app_folder, "omp_gemm")
     command = ["./ze_tracer", "-h", "-d", "-t", app_file, "gpu", "1024", "1"]
-  elif option == "-v" or option == "--kernels-per-tile" or option == "--conditional-collection":
+  elif option == "-v" or option == "--demangle" or\
+       option == "--kernels-per-tile" or option == "--conditional-collection":
     app_folder = utils.get_sample_executable_path("ze_gemm")
     app_file = os.path.join(app_folder, "ze_gemm")
     command = ["./ze_tracer", "-d", option, app_file, "1024", "1"]
@@ -81,6 +82,8 @@ def main(option):
     option = "-s"
   if len(sys.argv) > 1 and sys.argv[1] == "-v":
     option = "-v"
+  if len(sys.argv) > 1 and sys.argv[1] == "--demangle":
+    option = "--demangle"
   if len(sys.argv) > 1 and sys.argv[1] == "--kernels-per-tile":
     option = "--kernels-per-tile"
   if len(sys.argv) > 1 and sys.argv[1] == "--chrome-call-logging":

tools/cl_tracer/README.md

@@ -16,6 +16,7 @@ Options:
 --chrome-kernel-timeline       Dump device activities to JSON file per kernel name
 --chrome-device-stages         Dump device activities by stages to JSON file
 --verbose [-v]                 Enable verbose mode to show more kernel information
+--demangle                     Demangle DPC++ kernel names
 --tid                          Print thread ID into host API trace
 --pid                          Print process ID into host API and device activity trace
 --output [-o] <filename>       Print console logs into the file

tools/cl_tracer/cl_api_callbacks.h

@@ -2931,7 +2931,8 @@ static void clEnqueueTaskOnEnter(
   stream << " commandQueue = " << *(params->commandQueue);
   stream << " kernel = " << *(params->kernel);
   if (*(params->kernel) != nullptr) {
-    std::string kernel_name = utils::cl::GetKernelName(*(params->kernel));
+    std::string kernel_name = utils::cl::GetKernelName(
+        *(params->kernel), collector->Demangle());
     if (!kernel_name.empty()) {
       stream << " (" << kernel_name << ")";
     }
@@ -3582,7 +3583,8 @@ static void clEnqueueNDRangeKernelOnEnter(
   stream << " commandQueue = " << *(params->commandQueue);
   stream << " kernel = " << *(params->kernel);
   if (*(params->kernel) != nullptr) {
-    std::string kernel_name = utils::cl::GetKernelName(*(params->kernel));
+    std::string kernel_name = utils::cl::GetKernelName(
+        *(params->kernel), collector->Demangle());
     if (!kernel_name.empty()) {
       stream << " (" << kernel_name << ")";
     }
@@ -5966,6 +5968,9 @@ static void clCreateKernelOnEnter(
     stream << " kernelName = \"\"";
   } else {
     stream << " kernelName = \"" << *(params->kernelName) << "\"";
+    if (collector->Demangle()) {
+      stream << " (" << utils::Demangle(*(params->kernelName)) << ")";
+    }
   }
   stream << " errcodeRet = " << *(params->errcodeRet);
   stream << std::endl;

tools/cl_tracer/cl_api_collector.h

@@ -60,7 +60,7 @@ class ClApiCollector {
   static ClApiCollector* Create(
       cl_device_id device,
       Correlator* correlator,
-      ApiCollectorOptions options = {false, false, false},
+      ApiCollectorOptions options,
       OnClFunctionFinishCallback callback = nullptr,
       void* callback_data = nullptr) {
     PTI_ASSERT(device != nullptr);
@@ -115,6 +115,10 @@ class ClApiCollector {
     return options_.need_pid;
   }
 
+  bool Demangle() const {
+    return options_.demangle;
+  }
+
   void Log(const std::string& text) {
     PTI_ASSERT(correlator_ != nullptr);
     correlator_->Log(text);
@@ -332,7 +336,7 @@ class ClApiCollector {
   ClApiTracer* tracer_ = nullptr;
 
   Correlator* correlator_ = nullptr;
-  ApiCollectorOptions options_ = {false, false, false};
+  ApiCollectorOptions options_;
   cl_device_type device_type_ = CL_DEVICE_TYPE_ALL;
 
   OnClFunctionFinishCallback callback_ = nullptr;

tools/cl_tracer/cl_kernel_collector.h

@@ -113,15 +113,15 @@ class ClKernelCollector {
   static ClKernelCollector* Create(
       cl_device_id device,
       Correlator* correlator,
-      bool verbose,
+      KernelCollectorOptions options,
       OnClKernelFinishCallback callback = nullptr,
       void* callback_data = nullptr) {
     PTI_ASSERT(device != nullptr);
     PTI_ASSERT(correlator != nullptr);
     TraceGuard guard;
 
     ClKernelCollector* collector = new ClKernelCollector(
-        device, correlator, verbose, callback, callback_data);
+        device, correlator, options, callback, callback_data);
     PTI_ASSERT(collector != nullptr);
 
     ClApiTracer* tracer = new ClApiTracer(device, Callback, collector);
@@ -281,12 +281,12 @@ class ClKernelCollector {
   ClKernelCollector(
       cl_device_id device,
       Correlator* correlator,
-      bool verbose,
+      KernelCollectorOptions options,
       OnClKernelFinishCallback callback,
       void* callback_data)
       : device_(device),
         correlator_(correlator),
-        verbose_(verbose),
+        options_(options),
         callback_(callback),
         callback_data_(callback_data),
         kernel_id_(1) {
@@ -443,6 +443,13 @@ class ClKernelCollector {
       PTI_ASSERT(device != nullptr);
       AddKernelInterval(instance, device, started, ended);
 #else // PTI_KERNEL_INTERVALS
+      std::string name = instance->props.name;
+      PTI_ASSERT(!name.empty());
+
+      if (options_.verbose) {
+        name = GetVerboseName(&(instance->props));
+      }
+
       uint64_t host_queued = 0, host_submitted = 0;
       uint64_t host_started = 0, host_ended = 0;
       ComputeHostTimestamps(
@@ -451,19 +458,12 @@ class ClKernelCollector {
           host_queued, host_submitted,
           host_started, host_ended);
       AddKernelInfo(
-        &instance->props,
+        name,
         host_submitted - host_queued,
         host_started - host_submitted,
         host_ended - host_started);
 
       if (callback_ != nullptr) {
-        std::string name = instance->props.name;
-        PTI_ASSERT(!name.empty());
-
-        if (verbose_) {
-          name = GetVerboseName(&(instance->props));
-        }
-
         std::stringstream stream;
         stream << std::hex << queue;
 
@@ -551,17 +551,10 @@ class ClKernelCollector {
   }
 
   void AddKernelInfo(
-      const ClKernelProps* props, uint64_t queued_time,
+      std::string name, uint64_t queued_time,
       uint64_t submit_time, uint64_t execute_time) {
-    PTI_ASSERT(props != nullptr);
-
-    std::string name = props->name;
     PTI_ASSERT(!name.empty());
 
-    if (verbose_) {
-      name = GetVerboseName(props);
-    }
-
     if (kernel_info_map_.count(name) == 0) {
       ClKernelInfo info;
       info.queued_time = queued_time;
@@ -606,7 +599,7 @@ class ClKernelCollector {
     std::string name = instance->props.name;
     PTI_ASSERT(!name.empty());
 
-    if (verbose_) {
+    if (options_.verbose) {
       name = GetVerboseName(&instance->props);
     }
 
@@ -755,11 +748,9 @@ class ClKernelCollector {
     cl_ulong host_timestamp = 0;
     utils::cl::GetTimestamps(
         collector->device_, &enqueue_data->host_sync, &enqueue_data->device_sync);
-#ifndef PTI_KERNEL_INTERVALS
     PTI_ASSERT(collector->correlator_ != nullptr);
     enqueue_data->host_sync =
       collector->correlator_->GetTimestamp(enqueue_data->host_sync);
-#endif
 
     const T* params = reinterpret_cast<const T*>(data->functionParams);
     PTI_ASSERT(params != nullptr);
@@ -797,7 +788,8 @@ class ClKernelCollector {
       instance->event = **(params->event);
 
       cl_kernel kernel = *(params->kernel);
-      instance->props.name = utils::cl::GetKernelName(kernel);
+      instance->props.name = utils::cl::GetKernelName(
+          kernel, collector->options_.demangle);
 
       cl_command_queue queue = *(params->commandQueue);
       PTI_ASSERT(queue != nullptr);
@@ -1401,7 +1393,7 @@ class ClKernelCollector {
   ClApiTracer* tracer_ = nullptr;
   Correlator* correlator_ = nullptr;
 
-  bool verbose_ = false;
+  KernelCollectorOptions options_;
 
   std::atomic<uint64_t> kernel_id_;
   cl_device_id device_ = nullptr;

tools/cl_tracer/cl_tracer.h

@@ -81,11 +81,14 @@ class ClTracer {
         callback = ChromeStagesCallback;
       }
 
+      KernelCollectorOptions kernel_options;
+      kernel_options.verbose = tracer->CheckOption(TRACE_VERBOSE);
+      kernel_options.demangle = tracer->CheckOption(TRACE_DEMANGLE);
+
       if (cpu_device != nullptr) {
         cpu_kernel_collector = ClKernelCollector::Create(
             cpu_device, &tracer->correlator_,
-            tracer->CheckOption(TRACE_VERBOSE),
-            callback, tracer);
+            kernel_options, callback, tracer);
         if (cpu_kernel_collector == nullptr) {
           std::cerr <<
             "[WARNING] Unable to create kernel collector for CPU backend" <<
@@ -97,8 +100,7 @@ class ClTracer {
       if (gpu_device != nullptr) {
         gpu_kernel_collector = ClKernelCollector::Create(
             gpu_device, &tracer->correlator_,
-            tracer->CheckOption(TRACE_VERBOSE),
-            callback, tracer);
+            kernel_options, callback, tracer);
         if (gpu_kernel_collector == nullptr) {
           std::cerr <<
             "[WARNING] Unable to create kernel collector for GPU backend" <<
@@ -125,15 +127,16 @@ class ClTracer {
         callback = ChromeLoggingCallback;
       }
 
-      ApiCollectorOptions cl_api_options{false, false, false};
-      cl_api_options.call_tracing = tracer->CheckOption(TRACE_CALL_LOGGING);
-      cl_api_options.need_tid = tracer->CheckOption(TRACE_TID);
-      cl_api_options.need_pid = tracer->CheckOption(TRACE_PID);
+      ApiCollectorOptions api_options;
+      api_options.call_tracing = tracer->CheckOption(TRACE_CALL_LOGGING);
+      api_options.need_tid = tracer->CheckOption(TRACE_TID);
+      api_options.need_pid = tracer->CheckOption(TRACE_PID);
+      api_options.demangle = tracer->CheckOption(TRACE_DEMANGLE);
 
       if (cpu_device != nullptr) {
         cpu_api_collector = ClApiCollector::Create(
             cpu_device, &tracer->correlator_,
-            cl_api_options, callback, tracer);
+            api_options, callback, tracer);
         if (cpu_api_collector == nullptr) {
           std::cerr <<
             "[WARNING] Unable to create API collector for CPU backend" <<
@@ -145,7 +148,7 @@ class ClTracer {
       if (gpu_device != nullptr) {
         gpu_api_collector = ClApiCollector::Create(
             gpu_device, &tracer->correlator_,
-            cl_api_options, callback, tracer);
+            api_options, callback, tracer);
         if (gpu_api_collector == nullptr) {
           std::cerr <<
             "[WARNING] Unable to create API collector for GPU backend" <<

tools/cl_tracer/tool.cc

@@ -60,6 +60,10 @@ void Usage() {
     "--verbose [-v]                 " <<
     "Enable verbose mode to show more kernel information" <<
     std::endl;
+  std::cout <<
+    "--demangle                     " <<
+    "Demangle DPC++ kernel names" <<
+    std::endl;
   std::cout <<
     "--tid                          " <<
     "Print thread ID into host API trace" <<
@@ -125,6 +129,9 @@ int ParseArgs(int argc, char* argv[]) {
                strcmp(argv[i], "-v") == 0) {
       utils::SetEnv("CLT_Verbose", "1");
       ++app_index;
+    } else if (strcmp(argv[i], "--demangle") == 0) {
+      utils::SetEnv("CLT_Demangle", "1");
+      ++app_index;
     } else if (strcmp(argv[i], "--tid") == 0) {
       utils::SetEnv("CLT_Tid", "1");
       ++app_index;
@@ -235,6 +242,11 @@ static TraceOptions ReadArgs() {
     flags |= (1 << TRACE_VERBOSE);
   }
 
+  value = utils::GetEnv("CLT_Demangle");
+  if (!value.empty() && value == "1") {
+    flags |= (1 << TRACE_DEMANGLE);
+  }
+
   value = utils::GetEnv("CLT_Tid");
   if (!value.empty() && value == "1") {
     flags |= (1 << TRACE_TID);