Skip to content

Commit 9cce0d9

Browse files
againullagainull
authored
[XPTI] Emit metadata to sycl stream based on requested detail level (#21865)
Take into account subscriber requested detail level when emitting metadata into the "sycl" stream. Add metadata table showing which detail level provides each field. Existing tools using `sycl` or `sycl.debug` continue to work unchanged. Optionally tools can request BASIC level for minimal overhead (PTI use case). Depends on #21863
1 parent 2bb4cef commit 9cce0d9

5 files changed

Lines changed: 171 additions & 92 deletions

File tree

‎sycl/doc/design/SYCLInstrumentationUsingXPTI.md‎

Lines changed: 34 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -244,6 +244,18 @@ All trace point types in bold provide semantic information about the graph, node
244244
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"`.
245245

246246
If a tool subscribes to both `"sycl"` and `"sycl.debug"`, only notifications from `"sycl.debug"` will be delivered to avoid duplication.
247+
248+
### Stream Detail Level Control
249+
250+
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.
251+
252+
The SYCL runtime respects the effective detail level when emitting metadata and uses the following strategy:
253+
- Metadata is emitted at VERBOSE level if either:
254+
- A subscriber requests VERBOSE level on the `"sycl"` stream, OR
255+
- A subscriber is registered for the `"sycl.debug"` stream (legacy behavior)
256+
- Otherwise, metadata is emitted according to the requested detail level (BASIC or NORMAL)
257+
258+
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).
247259
| Trace Point Type | Parameter Description | Metadata |
248260
| :----------------: | :---------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- | :--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- |
249261
| **`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 |
@@ -262,28 +274,28 @@ If a tool subscribes to both `"sycl"` and `"sycl.debug"`, only notifications fro
262274

263275
### Metadata description
264276

265-
| Metadata | Type | Description |
266-
| :--------------------: | :-------------------------------: | :----------------------------------------------------------------------------------------------------------------------------------------------------------------- |
267-
| `access_mode` | `int` | Value of `sycl::access::mode` enum |
268-
| `access_range_start` | `size_t` | Start of accessor range |
269-
| `access_range_end` | `size_t` | End of accessor range |
270-
| `allocation_type` | C-style string | Allocation type |
271-
| `copy_from` | `size_t` | ID of source device |
272-
| `copy_to` | `size_t` | ID of target device |
273-
| `event` | `size_t` | Unique identifier of event |
274-
| `from_source` | `bool` | `true` if kernel comes from user source |
275-
| `kernel_name` | C-style string | Kernel name |
276-
| `memory_object` | `size_t` | Unique identifier of memory object |
277-
| `offset` | `size_t` | Accessor offset size in bytes |
278-
| `sycl_device` | `size_t` | Unique identifier of SYCL device |
279-
| `sycl_device_type` | C-style string | `CPU`, `GPU`, `ACC`, or `HOST` |
280-
| `sycl_device_name` | C-style string | Result of `sycl::device::get_info<sycl::info::name>()` |
281-
| `sym_function_name` | C-style string | Function name |
282-
| `sym_source_file_name` | C-style string | Source file name |
283-
| `sym_line_no` | `int32_t` | File line number |
284-
| `sym_column_no` | `int32_t` | File column number |
285-
| `enqueue_kernel_data` | `xpti::offload_kernel_arg_data_t` | Includes kernel execution parameters (global size, local size, offset) and number of kernel arguments |
286-
| `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. |
277+
| Metadata | Type | Detail Level | Description |
278+
| :--------------------: | :-------------------------------: | :----------------: | :----------------------------------------------------------------------------------------------------------------------------------------------------------------- |
279+
| `kernel_name` | C-style string | BASIC | Kernel name |
280+
| `memory_object` | `size_t` | BASIC | Unique identifier of memory object |
281+
| `offset` | `size_t` | NORMAL | Accessor offset size in bytes |
282+
| `access_mode` | `int` | NORMAL | Value of `sycl::access::mode` enum |
283+
| `access_range_start` | `size_t` | NORMAL | Start of accessor range |
284+
| `access_range_end` | `size_t` | NORMAL | End of accessor range |
285+
| `allocation_type` | C-style string | NORMAL | Allocation type |
286+
| `copy_from` | `size_t` | NORMAL | ID of source device |
287+
| `copy_to` | `size_t` | NORMAL | ID of target device |
288+
| `event` | `size_t` | NORMAL | Unique identifier of event |
289+
| `sycl_device` | `size_t` | NORMAL | Unique identifier of SYCL device |
290+
| `sycl_device_type` | C-style string | NORMAL | `CPU`, `GPU`, `ACC`, or `HOST` |
291+
| `sycl_device_name` | C-style string | NORMAL | Result of `sycl::device::get_info<sycl::info::name>()` |
292+
| `from_source` | `bool` | VERBOSE | `true` if kernel comes from user source |
293+
| `sym_function_name` | C-style string | VERBOSE | Function name |
294+
| `sym_source_file_name` | C-style string | VERBOSE | Source file name |
295+
| `sym_line_no` | `int32_t` | VERBOSE | File line number |
296+
| `sym_column_no` | `int32_t` | VERBOSE | File column number |
297+
| `enqueue_kernel_data` | `xpti::offload_kernel_arg_data_t` | VERBOSE | Includes kernel execution parameters (global size, local size, offset) and number of kernel arguments |
298+
| `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. |
287299

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

0 commit comments

Comments
 (0)