[SYCL][XPTI] Improvements to allow framework/app software level layers to provide code locations for sycl generated XPTI events. (#15190)

It is useful for framework software layers which uses sycl in their
implementation (like IPEX) to provide framework level code location
information for XPTI events generated by sycl. This allows a framework
specific instrumentation tool to capture sycl XPTI events with code
location
information coming from the framework level.
This allows the specific instrumentation tool, for example, to capture
sycl
stream task_begin and task_end events and correlate the specific
execution
with the upper layer graph node (or application level name of work) that
this
task represents by querying the payload attached to the events.

The change does not require any new APIs or ABI change, to capture a
code location the framework software layer should instantiate the
existing
sycl::detail::tls_code_loc_t object before calling a sycl entry point
(usually
queue.submit or graph.add).

There are 3 commits in this PR:
1) Change all sycl entry points that tries to set code location in TLS
to use the
code location that is already set in TLS, if one is set. Instead of
passing on
    the entry point code location at any case.
2) Payload for kernel execution commands uses the kernel name in place
of the
function name from code location. This changes this behavior in case
that the
upper layer software has captured code location in TLS before calling
sycl.
3) Fixes XPTI events in graph mode, some events were missing when
bypassing scheduler.

---------

Signed-off-by: Guy Zadicario <[email protected]>
Co-authored-by: Guy Zadicario <[email protected]>
Co-authored-by: Sergey Semenov <[email protected]>
Co-authored-by: Guy Zadickario <[email protected]>

Author: 4 people

sycl/include/sycl/detail/common.hpp

@@ -140,6 +140,9 @@ class __SYCL_EXPORT tls_code_loc_t {
   /// @return The code location information saved in the TLS slot. If not TLS
   /// entry has been set up, a default coe location is returned.
   const detail::code_location &query();
+  /// @brief Returns true if the TLS slot was cleared when this object was
+  /// constructed.
+  bool isToplevel() const { return !MLocalScope; }
 
 private:
   // The flag that is used to determine if the object is in a local scope or in

sycl/include/sycl/ext/oneapi/bindless_images.hpp

@@ -1369,7 +1369,7 @@ inline event queue::ext_oneapi_copy(
   detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
   return submit(
       [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, DestImgDesc); },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1383,7 +1383,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
                             DestImgDesc, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1396,7 +1396,7 @@ inline event queue::ext_oneapi_copy(
         CGH.depends_on(DepEvent);
         CGH.ext_oneapi_copy(Src, Dest, DestImgDesc);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1412,7 +1412,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
                             DestImgDesc, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1425,7 +1425,7 @@ inline event queue::ext_oneapi_copy(
         CGH.depends_on(DepEvents);
         CGH.ext_oneapi_copy(Src, Dest, DestImgDesc);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1441,7 +1441,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset,
                             DestImgDesc, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1451,7 +1451,7 @@ inline event queue::ext_oneapi_copy(
   detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
   return submit(
       [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc); },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1466,7 +1466,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
                             DestExtent, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1479,7 +1479,7 @@ inline event queue::ext_oneapi_copy(
         CGH.depends_on(DepEvent);
         CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1496,7 +1496,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
                             DestExtent, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1509,7 +1509,7 @@ inline event queue::ext_oneapi_copy(
         CGH.depends_on(DepEvents);
         CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1526,7 +1526,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
                             DestExtent, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1538,7 +1538,7 @@ inline event queue::ext_oneapi_copy(
       [&](handler &CGH) {
         CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1553,7 +1553,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
                             DeviceRowPitch, HostExtent, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1567,7 +1567,7 @@ inline event queue::ext_oneapi_copy(
         CGH.depends_on(DepEvent);
         CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1581,7 +1581,7 @@ inline event queue::ext_oneapi_copy(
         CGH.depends_on(DepEvent);
         CGH.ext_oneapi_copy(Src, Dest, ImageDesc);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1595,7 +1595,7 @@ inline event queue::ext_oneapi_copy(
         CGH.depends_on(DepEvents);
         CGH.ext_oneapi_copy(Src, Dest, ImageDesc);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1606,7 +1606,7 @@ inline event queue::ext_oneapi_copy(
   detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
   return submit(
       [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, ImageDesc); },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1622,7 +1622,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
                             DeviceRowPitch, HostExtent, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1636,7 +1636,7 @@ inline event queue::ext_oneapi_copy(
         CGH.depends_on(DepEvents);
         CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_copy(
@@ -1652,7 +1652,7 @@ inline event queue::ext_oneapi_copy(
         CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc,
                             DeviceRowPitch, HostExtent, CopyExtent);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_wait_external_semaphore(
@@ -1664,7 +1664,7 @@ inline event queue::ext_oneapi_wait_external_semaphore(
         CGH.depends_on(DepEvent);
         CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_wait_external_semaphore(
@@ -1676,7 +1676,7 @@ inline event queue::ext_oneapi_wait_external_semaphore(
         CGH.depends_on(DepEvents);
         CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_wait_external_semaphore(
@@ -1687,7 +1687,7 @@ inline event queue::ext_oneapi_wait_external_semaphore(
       [&](handler &CGH) {
         CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle, WaitValue);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_wait_external_semaphore(
@@ -1699,7 +1699,7 @@ inline event queue::ext_oneapi_wait_external_semaphore(
         CGH.depends_on(DepEvent);
         CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle, WaitValue);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_wait_external_semaphore(
@@ -1712,7 +1712,7 @@ inline event queue::ext_oneapi_wait_external_semaphore(
         CGH.depends_on(DepEvents);
         CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle, WaitValue);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_signal_external_semaphore(
@@ -1723,7 +1723,7 @@ inline event queue::ext_oneapi_signal_external_semaphore(
       [&](handler &CGH) {
         CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_signal_external_semaphore(
@@ -1735,7 +1735,7 @@ inline event queue::ext_oneapi_signal_external_semaphore(
         CGH.depends_on(DepEvent);
         CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_signal_external_semaphore(
@@ -1747,7 +1747,7 @@ inline event queue::ext_oneapi_signal_external_semaphore(
         CGH.depends_on(DepEvents);
         CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_signal_external_semaphore(
@@ -1758,7 +1758,7 @@ inline event queue::ext_oneapi_signal_external_semaphore(
       [&](handler &CGH) {
         CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle, SignalValue);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_signal_external_semaphore(
@@ -1771,7 +1771,7 @@ inline event queue::ext_oneapi_signal_external_semaphore(
         CGH.depends_on(DepEvent);
         CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle, SignalValue);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 inline event queue::ext_oneapi_signal_external_semaphore(
@@ -1784,7 +1784,7 @@ inline event queue::ext_oneapi_signal_external_semaphore(
         CGH.depends_on(DepEvents);
         CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle, SignalValue);
       },
-      CodeLoc);
+      TlsCodeLocCapture.query());
 }
 
 } // namespace _V1

sycl/include/sycl/handler.hpp

@@ -591,7 +591,11 @@ class __SYCL_EXPORT handler {
 
   /// Saves the location of user's code passed in \p CodeLoc for future usage in
   /// finalize() method.
-  void saveCodeLoc(detail::code_location CodeLoc) { MCodeLoc = CodeLoc; }
+  /// TODO: remove the first version of this func (the one without the IsTopCodeLoc arg)
+  ///   at the next ABI breaking window since removing it breaks ABI on windows.
+  void saveCodeLoc(detail::code_location CodeLoc);
+  void saveCodeLoc(detail::code_location CodeLoc, bool IsTopCodeLoc);
+  void copyCodeLoc(const handler &other);
 
   /// Constructs CG object of specific type, passes it to Scheduler and
   /// returns sycl::event object representing the command group.