[SYCL][XPTI] Emit XPTI events for kernel nodes added to user graph.

When graph is finalized kernels can be directly enqueued to a command
buffer without using the scheduler, in this case XPTI events for the
node_create/task_begin/task_end were missing.

Save code location from tls when CGF is added to a graph through
graph.add API. This captures user stored code location in TLS, if one
was set before calling graph.add

Signed-off-by: Guy Zadicario <[email protected]>

Author: gzadicario

sycl/source/detail/graph_impl.cpp

@@ -13,6 +13,7 @@
 #include <detail/queue_impl.hpp>
 #include <detail/scheduler/commands.hpp>
 #include <detail/sycl_mem_obj_t.hpp>
+#include <sycl/detail/common.hpp>
 #include <sycl/feature_test.hpp>
 #include <sycl/queue.hpp>
 
@@ -372,6 +373,15 @@ graph_impl::add(const std::shared_ptr<graph_impl> &Impl,
                 const std::vector<std::shared_ptr<node_impl>> &Dep) {
   (void)Args;
   sycl::handler Handler{Impl};
+
+  // save code location if one was set in TLS.
+  // idealy it would be nice to capture user's call code location
+  // by adding a parameter to the graph.add function, but this will
+  // break the API. At least capture code location from TLS, user
+  // can set it before calling graph.add
+  sycl::detail::tls_code_loc_t Tls;
+  Handler.saveCodeLoc(Tls.query(), Tls.isToplevel());
+
   CGF(Handler);
 
   if (Handler.getType() == sycl::detail::CGType::Barrier) {
@@ -667,6 +677,23 @@ exec_graph_impl::enqueueNodeDirect(sycl::context Ctx,
   }
   ur_exp_command_buffer_sync_point_t NewSyncPoint;
   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.c_str(), 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(
       Ctx, DeviceImpl, CommandBuffer,
       *static_cast<sycl::detail::CGExecKernel *>((Node->MCommandGroup.get())),
@@ -679,6 +706,12 @@ exec_graph_impl::enqueueNodeDirect(sycl::context Ctx,
                           "Failed to add kernel to UR command-buffer");
   }
 
+#ifdef XPTI_ENABLE_INSTRUMENTATION
+  if (CmdTraceEvent)
+    sycl::detail::emitInstrumentationGeneral(
+        StreamID, InstanceID, CmdTraceEvent, xpti::trace_task_end, nullptr);
+#endif
+
   return NewSyncPoint;
 }
 

sycl/source/detail/scheduler/commands.cpp

@@ -2007,8 +2007,10 @@ void instrumentationAddExtraKernelMetadata(
     Program = SyclKernel->getProgramRef();
     if (!SyclKernel->isCreatedFromSource())
       EliminatedArgMask = SyclKernel->getKernelArgMask();
-  } else {
-    assert(Queue && "Kernel submissions should have an associated queue");
+  } else if (Queue) {
+    // NOTE: Queue can be null when kerner is directly enqueued to a command
+    // buffer
+    //       by graph API, when a midifiable graph is finalized.
     std::tie(Kernel, KernelMutex, EliminatedArgMask, Program) =
         detail::ProgramManager::getInstance().getOrCreateKernel(
             Queue->getContextImplPtr(), Queue->getDeviceImplPtr(), KernelName);
@@ -2134,7 +2136,12 @@ std::pair<xpti_td *, uint64_t> emitKernelInstrumentationData(
 
   if (CmdTraceEvent) {
     // Stash the queue_id mutable metadata in TLS
-    xpti::framework::stash_tuple(XPTI_QUEUE_INSTANCE_ID_KEY, getQueueID(Queue));
+    // NOTE: Queue can be null when kerner is directly enqueued to a command
+    // buffer
+    //       by graph API, when a midifiable graph is finalized.
+    if (Queue.get())
+      xpti::framework::stash_tuple(XPTI_QUEUE_INSTANCE_ID_KEY,
+                                   getQueueID(Queue));
     instrumentationAddExtraKernelMetadata(CmdTraceEvent, NDRDesc,
                                           KernelBundleImplPtr, SyclKernelName,
                                           SyclKernel, Queue, CGArgs);