[SYCL] Eliminate XPTI overhead when it is disabled (#18334)

Even if the XPTI collector is not used, we still have some overhead
related to XPTI instrumentation when we call `xptiRegisterStream`
function.

With these changes, we cache the result of the `xptiTraceEnabled`
function in the static variables. It is safe to cache it because XPTI
collector is enabled by the `XPTI_FRAMEWORK_DISPATCHER` environment
variable before the process starts.

VTune shows ~2% decrease in the number of instructions retired on a
submit kernel path.

---------

Signed-off-by: Sergei Vinogradov <[email protected]>

Author: vinser52

sycl/source/detail/graph_impl.cpp

@@ -800,19 +800,25 @@ exec_graph_impl::enqueueNodeDirect(sycl::context Ctx,
   ur_exp_command_buffer_command_handle_t NewCommand = 0;
 
 #ifdef XPTI_ENABLE_INSTRUMENTATION
-  int32_t StreamID = xptiRegisterStream(sycl::detail::SYCL_STREAM_NAME);
-  sycl::detail::CGExecKernel *CGExec =
-      static_cast<sycl::detail::CGExecKernel *>(Node->MCommandGroup.get());
-  sycl::detail::code_location CodeLoc(CGExec->MFileName.c_str(),
-                                      CGExec->MFunctionName.c_str(),
-                                      CGExec->MLine, CGExec->MColumn);
-  auto [CmdTraceEvent, InstanceID] = emitKernelInstrumentationData(
-      StreamID, CGExec->MSyclKernel, CodeLoc, CGExec->MIsTopCodeLoc,
-      CGExec->MKernelName.data(), nullptr, CGExec->MNDRDesc,
-      CGExec->MKernelBundle, CGExec->MArgs);
-  if (CmdTraceEvent)
-    sycl::detail::emitInstrumentationGeneral(
-        StreamID, InstanceID, CmdTraceEvent, xpti::trace_task_begin, nullptr);
+  const bool xptiEnabled = xptiTraceEnabled();
+  int32_t StreamID = xpti::invalid_id;
+  xpti_td *CmdTraceEvent = nullptr;
+  uint64_t InstanceID = 0;
+  if (xptiEnabled) {
+    StreamID = xptiRegisterStream(sycl::detail::SYCL_STREAM_NAME);
+    sycl::detail::CGExecKernel *CGExec =
+        static_cast<sycl::detail::CGExecKernel *>(Node->MCommandGroup.get());
+    sycl::detail::code_location CodeLoc(CGExec->MFileName.c_str(),
+                                        CGExec->MFunctionName.c_str(),
+                                        CGExec->MLine, CGExec->MColumn);
+    std::tie(CmdTraceEvent, InstanceID) = emitKernelInstrumentationData(
+        StreamID, CGExec->MSyclKernel, CodeLoc, CGExec->MIsTopCodeLoc,
+        CGExec->MKernelName.data(), nullptr, CGExec->MNDRDesc,
+        CGExec->MKernelBundle, CGExec->MArgs);
+    if (CmdTraceEvent)
+      sycl::detail::emitInstrumentationGeneral(
+          StreamID, InstanceID, CmdTraceEvent, xpti::trace_task_begin, nullptr);
+  }
 #endif
 
   ur_result_t Res = sycl::detail::enqueueImpCommandBufferKernel(
@@ -830,7 +836,7 @@ exec_graph_impl::enqueueNodeDirect(sycl::context Ctx,
   }
 
 #ifdef XPTI_ENABLE_INSTRUMENTATION
-  if (CmdTraceEvent)
+  if (xptiEnabled && CmdTraceEvent)
     sycl::detail::emitInstrumentationGeneral(
         StreamID, InstanceID, CmdTraceEvent, xpti::trace_task_end, nullptr);
 #endif

sycl/source/detail/queue_impl.cpp

@@ -603,11 +603,15 @@ void queue_impl::instrumentationEpilog(void *TelemetryEvent, std::string &Name,
 void queue_impl::wait(const detail::code_location &CodeLoc) {
   (void)CodeLoc;
 #ifdef XPTI_ENABLE_INSTRUMENTATION
+  const bool xptiEnabled = xptiTraceEnabled();
   void *TelemetryEvent = nullptr;
   uint64_t IId;
   std::string Name;
-  int32_t StreamID = xptiRegisterStream(SYCL_STREAM_NAME);
-  TelemetryEvent = instrumentationProlog(CodeLoc, Name, StreamID, IId);
+  int32_t StreamID = xpti::invalid_id;
+  if (xptiEnabled) {
+    StreamID = xptiRegisterStream(SYCL_STREAM_NAME);
+    TelemetryEvent = instrumentationProlog(CodeLoc, Name, StreamID, IId);
+  }
 #endif
 
   if (MGraph.lock()) {
@@ -675,7 +679,9 @@ void queue_impl::wait(const detail::code_location &CodeLoc) {
     Event->wait(Event);
 
 #ifdef XPTI_ENABLE_INSTRUMENTATION
-  instrumentationEpilog(TelemetryEvent, Name, StreamID, IId);
+  if (xptiEnabled) {
+    instrumentationEpilog(TelemetryEvent, Name, StreamID, IId);
+  }
 #endif
 }
 

sycl/source/handler.cpp

@@ -533,19 +533,23 @@ event handler::finalize() {
       }
 
 #ifdef XPTI_ENABLE_INSTRUMENTATION
-      // uint32_t StreamID, uint64_t InstanceID, xpti_td* TraceEvent,
-      int32_t StreamID = xptiRegisterStream(detail::SYCL_STREAM_NAME);
-      auto [CmdTraceEvent, InstanceID] = emitKernelInstrumentationData(
-          StreamID, MKernel, MCodeLoc, impl->MIsTopCodeLoc, MKernelName.data(),
-          MQueue, impl->MNDRDesc, KernelBundleImpPtr, impl->MArgs);
-      auto EnqueueKernel = [&, CmdTraceEvent = CmdTraceEvent,
-                            InstanceID = InstanceID]() {
-#else
-      auto EnqueueKernel = [&]() {
+      const bool xptiEnabled = xptiTraceEnabled();
 #endif
+      auto EnqueueKernel = [&]() {
 #ifdef XPTI_ENABLE_INSTRUMENTATION
-        detail::emitInstrumentationGeneral(StreamID, InstanceID, CmdTraceEvent,
-                                           xpti::trace_task_begin, nullptr);
+        int32_t StreamID = xpti::invalid_id;
+        xpti_td *CmdTraceEvent = nullptr;
+        uint64_t InstanceID = 0;
+        if (xptiEnabled) {
+          StreamID = xptiRegisterStream(detail::SYCL_STREAM_NAME);
+          std::tie(CmdTraceEvent, InstanceID) = emitKernelInstrumentationData(
+              StreamID, MKernel, MCodeLoc, impl->MIsTopCodeLoc,
+              MKernelName.data(), MQueue, impl->MNDRDesc, KernelBundleImpPtr,
+              impl->MArgs);
+          detail::emitInstrumentationGeneral(StreamID, InstanceID,
+                                             CmdTraceEvent,
+                                             xpti::trace_task_begin, nullptr);
+        }
 #endif
         const detail::RTDeviceBinaryImage *BinImage = nullptr;
         if (detail::SYCLConfig<detail::SYCL_JIT_AMDGCN_PTX_KERNELS>::get()) {
@@ -561,14 +565,17 @@ event handler::finalize() {
                          impl->MKernelUsesClusterLaunch,
                          impl->MKernelWorkGroupMemorySize, BinImage);
 #ifdef XPTI_ENABLE_INSTRUMENTATION
-        // Emit signal only when event is created
-        if (!DiscardEvent) {
-          detail::emitInstrumentationGeneral(
-              StreamID, InstanceID, CmdTraceEvent, xpti::trace_signal,
-              static_cast<const void *>(LastEventImpl->getHandle()));
+        if (xptiEnabled) {
+          // Emit signal only when event is created
+          if (!DiscardEvent) {
+            detail::emitInstrumentationGeneral(
+                StreamID, InstanceID, CmdTraceEvent, xpti::trace_signal,
+                static_cast<const void *>(LastEventImpl->getHandle()));
+          }
+          detail::emitInstrumentationGeneral(StreamID, InstanceID,
+                                             CmdTraceEvent,
+                                             xpti::trace_task_end, nullptr);
         }
-        detail::emitInstrumentationGeneral(StreamID, InstanceID, CmdTraceEvent,
-                                           xpti::trace_task_end, nullptr);
 #endif
       };