Skip to content

[XPTI] Emit metadata to sycl stream based on requested detail level #21865

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
againull merged 5 commits into from
Apr 29, 2026
Merged

[XPTI] Emit metadata to sycl stream based on requested detail level #21865

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
0b837ea
[XPTI] Add per-subscriber stream detail level control
againull Apr 22, 2026
8547684
[XPTI] Emit metadata to sycl stream based on requested detail level
againull Apr 22, 2026
9000036
Address review
againull Apr 28, 2026
f1427e3
Fix missed checks
againull Apr 28, 2026
1cb15de
Merge branch 'sycl' into againull/sycl_xpti_stream_detail_level
againull Apr 29, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
56 changes: 34 additions & 22 deletions sycl/doc/design/SYCLInstrumentationUsingXPTI.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -244,6 +244,18 @@ All trace point types in bold provide semantic information about the graph, node
The `"sycl.debug"` stream emits the same notifications as the `"sycl"` stream, but with additional metadata. If toolchains want to keep the overhead low then subscribing to `"sycl"` stream is the right option, if toolchains want to get more data and keeping overheads low is not important then they should subscribe to `"sycl.debug"`.

If a tool subscribes to both `"sycl"` and `"sycl.debug"`, only notifications from `"sycl.debug"` will be delivered to avoid duplication.

### Stream Detail Level Control

The SYCL runtime implements XPTI stream detail level control for the `"sycl"` stream (see [XPTI Stream Detail Level Control](https://github.com/intel/llvm/tree/sycl/xptifw/doc/XPTI_Framework.md#stream-detail-level-control) for the XPTI framework feature description). This allows subscribers to control the amount of optional metadata emitted based on their needs, providing fine-grained control over tracing overhead.

The SYCL runtime respects the effective detail level when emitting metadata and uses the following strategy:
- Metadata is emitted at VERBOSE level if either:
- A subscriber requests VERBOSE level on the `"sycl"` stream, OR
- A subscriber is registered for the `"sycl.debug"` stream (legacy behavior)
- Otherwise, metadata is emitted according to the requested detail level (BASIC or NORMAL)

See the metadata table below for which fields are emitted at each detail level. Note that stream detail level control is only implemented for the `"sycl"` stream; the `"sycl.debug"` stream always emits all metadata (equivalent to VERBOSE level).
| Trace Point Type | Parameter Description | Metadata |
| :----------------: | :---------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- | :--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- |
| **`graph_create`** | <div style="text-align: left"><li>**trace_type**: `xpti::trace_point_type_t::graph_create` that marks the creation of an asynchronous graph.</li> <li> **parent**: `nullptr`</li> <li> **event**: The global asynchronous graph object ID. All other graph related events such as node and edge creation will always this ID as the parent ID. </li> <li> **instance**: Unique ID related to the event, but not a correlation ID as there are other events to correlate to. </li> <li> **user_data**: `nullptr`</li> <p></p> SYCL runtime will always have one instance of a graph object with many disjoint subgraphs that get created during the execution of an application. </div> | None |
Expand All @@ -262,28 +274,28 @@ If a tool subscribes to both `"sycl"` and `"sycl.debug"`, only notifications fro

### Metadata description

| Metadata | Type | Description |
| :--------------------: | :-------------------------------: | :----------------------------------------------------------------------------------------------------------------------------------------------------------------- |
| `access_mode` | `int` | Value of `sycl::access::mode` enum |
| `access_range_start` | `size_t` | Start of accessor range |
| `access_range_end` | `size_t` | End of accessor range |
| `allocation_type` | C-style string | Allocation type |
| `copy_from` | `size_t` | ID of source device |
| `copy_to` | `size_t` | ID of target device |
| `event` | `size_t` | Unique identifier of event |
| `from_source` | `bool` | `true` if kernel comes from user source |
| `kernel_name` | C-style string | Kernel name |
| `memory_object` | `size_t` | Unique identifier of memory object |
| `offset` | `size_t` | Accessor offset size in bytes |
| `sycl_device` | `size_t` | Unique identifier of SYCL device |
| `sycl_device_type` | C-style string | `CPU`, `GPU`, `ACC`, or `HOST` |
| `sycl_device_name` | C-style string | Result of `sycl::device::get_info<sycl::info::name>()` |
| `sym_function_name` | C-style string | Function name |
| `sym_source_file_name` | C-style string | Source file name |
| `sym_line_no` | `int32_t` | File line number |
| `sym_column_no` | `int32_t` | File column number |
| `enqueue_kernel_data` | `xpti::offload_kernel_arg_data_t` | Includes kernel execution parameters (global size, local size, offset) and number of kernel arguments |
| `argN` | `xpti::offload_kernel_arg_data_t` | Description for the Nth kernel argument. It includes argument kind (sycl::detail::kernel_param_kind_t), pointer to the value, size and index in the argument list. |
| Metadata | Type | Detail Level | Description |
| :--------------------: | :-------------------------------: | :----------------: | :----------------------------------------------------------------------------------------------------------------------------------------------------------------- |
| `kernel_name` | C-style string | BASIC | Kernel name |
| `memory_object` | `size_t` | BASIC | Unique identifier of memory object |
| `offset` | `size_t` | NORMAL | Accessor offset size in bytes |
| `access_mode` | `int` | NORMAL | Value of `sycl::access::mode` enum |
| `access_range_start` | `size_t` | NORMAL | Start of accessor range |
| `access_range_end` | `size_t` | NORMAL | End of accessor range |
| `allocation_type` | C-style string | NORMAL | Allocation type |
| `copy_from` | `size_t` | NORMAL | ID of source device |
| `copy_to` | `size_t` | NORMAL | ID of target device |
| `event` | `size_t` | NORMAL | Unique identifier of event |
| `sycl_device` | `size_t` | NORMAL | Unique identifier of SYCL device |
| `sycl_device_type` | C-style string | NORMAL | `CPU`, `GPU`, `ACC`, or `HOST` |
| `sycl_device_name` | C-style string | NORMAL | Result of `sycl::device::get_info<sycl::info::name>()` |
| `from_source` | `bool` | VERBOSE | `true` if kernel comes from user source |
| `sym_function_name` | C-style string | VERBOSE | Function name |
| `sym_source_file_name` | C-style string | VERBOSE | Source file name |
| `sym_line_no` | `int32_t` | VERBOSE | File line number |
| `sym_column_no` | `int32_t` | VERBOSE | File column number |
| `enqueue_kernel_data` | `xpti::offload_kernel_arg_data_t` | VERBOSE | Includes kernel execution parameters (global size, local size, offset) and number of kernel arguments |
| `argN` | `xpti::offload_kernel_arg_data_t` | VERBOSE | Description for the Nth kernel argument. It includes argument kind (sycl::detail::kernel_param_kind_t), pointer to the value, size and index in the argument list. |

## Buffer management stream `"sycl.experimental.buffer"` Notification Signatures

Expand Down
Loading
Loading