Skip to content

Commit

Permalink
[SYCL][XPTI] Improvements to allow framework/app software level layer…
Browse files Browse the repository at this point in the history 
…s 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]>
  • Loading branch information
@gzadica1 @gzadicario @sergey-semenov
4 people authored Oct 7, 2024
1 parent a4f74a9 commit ea95271
Show file tree
Hide file tree
Showing 19 changed files with 388 additions and 109 deletions.
3 changes: 3 additions & 0 deletions sycl/include/sycl/detail/common.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
64 changes: 32 additions & 32 deletions sycl/include/sycl/ext/oneapi/bindless_images.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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
Expand Down
6 changes: 5 additions & 1 deletion sycl/include/sycl/handler.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
Loading

0 comments on commit ea95271

Please sign in to comment.