diff --git a/sycl/doc/EnvironmentVariables.md b/sycl/doc/EnvironmentVariables.md index aecb6ddbab134..044f2b70ef54b 100755 --- a/sycl/doc/EnvironmentVariables.md +++ b/sycl/doc/EnvironmentVariables.md @@ -150,6 +150,7 @@ variables in production code. | `SYCL_PI_LEVEL_ZERO_USE_COMPUTE_ENGINE` | Integer | It can be set to an integer (>=0) in which case all compute commands will be submitted to the command-queue with the given index in the compute command group. If it is instead set to a negative value then all available compute engines may be used. The default value is "0" | | `SYCL_PI_LEVEL_ZERO_USE_COPY_ENGINE_FOR_D2D_COPY` (experimental) | Integer | Allows the use of copy engine, if available in the device, in Level Zero plugin for device to device copy operations. The default is 0. This option is experimental and will be removed once heuristics are added to make a decision about use of copy engine for device to device copy operations. | | `SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS` | Any(\*) | Enable support of device-scope events whose state is not visible to the host. If enabled mode is SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS=1 the Level Zero plugin would create all events having device-scope only and create proxy host-visible events for them when their status is needed (wait/query) on the host. If enabled mode is SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS=2 the Level Zero plugin would create all events having device-scope and add proxy host-visible event at the end of each command-list submission. The default is 0, meaning all events are host-visible. | +| `SYCL_PI_LEVEL_ZERO_ENABLE_TRACING` | Any(\*) | Enable XPTI-based tracing in L0 plugin | ## Debugging variables for CUDA Plugin diff --git a/sycl/doc/design/SYCLInstrumentationUsingXPTI.md b/sycl/doc/design/SYCLInstrumentationUsingXPTI.md index 581aa9b9e25ea..8a9d3f83b14a6 100644 --- a/sycl/doc/design/SYCLInstrumentationUsingXPTI.md +++ b/sycl/doc/design/SYCLInstrumentationUsingXPTI.md @@ -299,3 +299,23 @@ All trace point types in bold provide semantic information about the graph, node | `mem_alloc_end` |
  • **trace_type**: `xpti::trace_point_type_t::mem_alloc_end` that marks the end of memory allocation process
  • **parent**: Event ID created for all functions in the `oneapi.level_zero.experimental.mem_alloc` layer.
  • **event**: `nullptr` - since the stream of data just captures functions being called.
  • **instance**: Unique ID to allow the correlation of the `mem_alloc_begin` event with the `mem_alloc_end` event. This value is guaranteed to be the same value received by the trace event for the corresponding `mem_alloc_begin`.
  • **user_data**: A pointer to `mem_alloc_data_t` object, that includes memory object ID (if any), allocated pointer, allocation size, and guard zone size (if any).
    • | None | | `mem_release_begin` |
  • **trace_type**: `xpti::trace_point_type_t::mem_release_begin` that marks the beginning of memory allocation process
  • **parent**: Event ID created for all functions in the `oneapi.level_zero.experimental.mem_alloc` layer.
  • **event**: `nullptr` - since the stream of data just captures functions being called.
  • **instance**: Unique ID to allow the correlation of the `mem_release_begin` event with the `mem_release_end` event.
  • **user_data**: A pointer to `mem_alloc_data_t` object, that includes memory object ID (if any) and released pointer.
    • | None | | `mem_release_end` |
  • **trace_type**: `xpti::trace_point_type_t::mem_release_end` that marks the end of memory allocation process
  • **parent**: Event ID created for all functions in the `oneapi.level_zero.experimental.mem_alloc` layer.
  • **event**: `nullptr` - since the stream of data just captures functions being called.
  • **instance**: Unique ID to allow the correlation of the `mem_release_begin` event with the `mem_release_end` event. This value is guaranteed to be the same value received by the trace event for the corresponding `mem_release_begin`.
  • **user_data**: A pointer to `mem_alloc_data_t` object, that includes memory object ID (if any) and released pointer.
    • | None | + +## SYCL Stream `"sycl.experimental.level_zero.call"` Notification Signatures + +This stream transfers events about Level Zero API calls made by SYCL +application. + +| Trace Point Type | Parameter Description | Metadata | +| :--------------: | :-------------------- | :------- | +| `function_begin` |
  • **trace_type**: `xpti::trace_point_type_t::function_begin` that marks the beginning of a function
  • **parent**: Event ID created for all functions in the `sycl.pi` layer.
  • **event**: `nullptr` - since the stream of data just captures functions being called.
  • **instance**: Unique ID to allow the correlation of the `function_begin` event with the `function_end` event.
  • **user_data**: Name of the function being called sent in as `const char *`
    • | None | +| `function_end` |
  • **trace_type**: `xpti::trace_point_type_t::function_end` that marks the beginning of a function
  • **parent**: Event ID created for all functions in the `sycl.pi` layer.
  • **event**: `nullptr` - since the stream of data just captures functions being called.
  • **instance**: Unique ID to allow the correlation of the `function_begin` event with the `function_end` event. This value is guaranteed to be the same value received by the trace event for the corresponding `function_begin`
  • **user_data**: Name of the function being called sent in as `const char *`
    • | None | + +## SYCL Stream `"sycl.experimental.level_zero.debug"` Notification Signatures + +This stream transfers events about Level Zero API calls and their function +arguments made by SYCL application. + +| Trace Point Type | Parameter Description | Metadata | +| :------------------------: | :-------------------- | :------- | +| `function_with_args_begin` |
  • **trace_type**: `xpti::trace_point_type_t::function_with_args_begin` that marks the beginning of a function
  • **parent**: Event ID created for all functions in the `sycl.pi.debug` layer.
  • **event**: `nullptr` - since the stream of data just captures functions being called.
  • **instance**: Unique ID to allow the correlation of the `function_with_args_begin` event with the `function_with_args_end` event.
  • **user_data**: A pointer to `function_with_args_t` object, that includes function ID, name, and arguments.
    • | None | +| `function_with_args_end` |
  • **trace_type**: `xpti::trace_point_type_t::function_with_args_end` that marks the beginning of a function
  • **parent**: Event ID created for all functions in the `sycl.pi.debug` layer.
  • **event**: `nullptr` - since the stream of data just captures functions being called.
  • **instance**: Unique ID to allow the correlation of the `function_with_args_begin` event with the `function_with_args_end` event. This value is guaranteed to be the same value received by the trace event for the corresponding `function_with_args_begin`
  • **user_data**: A pointer to `function_with_args_t` object, that includes function ID, name, arguments, and return value.
    • | None | diff --git a/sycl/plugins/level_zero/CMakeLists.txt b/sycl/plugins/level_zero/CMakeLists.txt index 950d60394f7e2..80de3fa4666b8 100755 --- a/sycl/plugins/level_zero/CMakeLists.txt +++ b/sycl/plugins/level_zero/CMakeLists.txt @@ -99,6 +99,10 @@ target_include_directories(LevelZeroLoader-Headers INTERFACE "${LEVEL_ZERO_INCLUDE_DIR}" ) +if (SYCL_ENABLE_XPTI_TRACING) + set(XPTI_PROXY_SRC "${CMAKE_SOURCE_DIR}/../xpti/src/xpti_proxy.cpp") +endif() + find_package(Threads REQUIRED) add_sycl_plugin(level_zero SOURCES @@ -107,12 +111,36 @@ add_sycl_plugin(level_zero "${CMAKE_CURRENT_SOURCE_DIR}/pi_level_zero.hpp" "${CMAKE_CURRENT_SOURCE_DIR}/usm_allocator.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/usm_allocator.hpp" + "${CMAKE_CURRENT_SOURCE_DIR}/tracing.cpp" + ${XPTI_PROXY_SRC} LIBRARIES "${LEVEL_ZERO_LOADER}" Threads::Threads ) +find_package(Python3 REQUIRED) + +add_custom_target(ze-api + COMMAND ${Python3_EXECUTABLE} + ${CMAKE_CURRENT_SOURCE_DIR}/ze_api_generator.py + ${LEVEL_ZERO_INCLUDE_DIR}/level_zero/ze_api.h + BYPRODUCTS + ${CMAKE_CURRENT_BINARY_DIR}/ze_api.def + ) +target_include_directories(pi_level_zero PRIVATE ${CMAKE_CURRENT_BINARY_DIR}) +add_dependencies(pi_level_zero ze-api) + +if (SYCL_ENABLE_XPTI_TRACING) + target_compile_definitions(pi_level_zero PRIVATE + XPTI_ENABLE_INSTRUMENTATION + XPTI_STATIC_LIBRARY + ) + target_include_directories(pi_level_zero PRIVATE "${CMAKE_SOURCE_DIR}/../xpti/include") + target_link_libraries(pi_level_zero PRIVATE ${CMAKE_DL_LIBS}) +endif() + if (TARGET level-zero-loader) + add_dependencies(ze-api level-zero-loader) add_dependencies(pi_level_zero level-zero-loader) endif() diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 3e6dc6f8ee6b7..3eb240a7631b5 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -36,6 +36,8 @@ static pi_result EventCreate(pi_context Context, pi_queue Queue, bool HostVisible, pi_event *RetEvent); } +void enableL0Tracing(); + namespace { // Controls Level Zero calls serialization to w/a Level Zero driver being not MT @@ -7664,6 +7666,10 @@ pi_result piPluginInit(pi_plugin *PluginInit) { (PluginInit->PiFunctionTable).api = (decltype(&::api))(&api); #include + if (std::getenv("SYCL_PI_LEVEL_ZERO_ENABLE_TRACING") != nullptr) { + enableL0Tracing(); + } + return PI_SUCCESS; } diff --git a/sycl/plugins/level_zero/tracing.cpp b/sycl/plugins/level_zero/tracing.cpp new file mode 100644 index 0000000000000..3429cd7fed53e --- /dev/null +++ b/sycl/plugins/level_zero/tracing.cpp @@ -0,0 +1,145 @@ +//===-------------- tracing.cpp - L0 Host API Tracing ----------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "xpti/xpti_data_types.h" +#include +#include +#include +#include + +#include + +constexpr auto L0_CALL_STREAM_NAME = "sycl.experimental.level_zero.call"; +constexpr auto L0_DEBUG_STREAM_NAME = "sycl.experimental.level_zero.debug"; + +thread_local uint64_t CallCorrelationID = 0; +thread_local uint64_t DebugCorrelationID = 0; + +constexpr auto GVerStr = "0.1"; +constexpr int GMajVer = 0; +constexpr int GMinVer = 1; + +#ifdef XPTI_ENABLE_INSTRUMENTATION +static xpti_td *GCallEvent = nullptr; +static xpti_td *GDebugEvent = nullptr; +#endif // XPTI_ENABLE_INSTRUMENTATION + +enum class ZEApiKind { +#define _ZE_API(call, domain, cb, params_type) call, +#include "ze_api.def" +#undef _ZE_API +}; + +void enableL0Tracing() { +#ifdef XPTI_ENABLE_INSTRUMENTATION + if (!xptiTraceEnabled()) + return; + + xptiRegisterStream(L0_CALL_STREAM_NAME); + xptiInitialize(L0_CALL_STREAM_NAME, GMajVer, GMinVer, GVerStr); + xptiRegisterStream(L0_DEBUG_STREAM_NAME); + xptiInitialize(L0_DEBUG_STREAM_NAME, GMajVer, GMinVer, GVerStr); + + uint64_t Dummy; + xpti::payload_t L0Payload("Level Zero Plugin Layer"); + GCallEvent = + xptiMakeEvent("L0 Plugin Layer", &L0Payload, xpti::trace_algorithm_event, + xpti_at::active, &Dummy); + + xpti::payload_t L0DebugPayload("L0 Plugin Debug Layer"); + GDebugEvent = + xptiMakeEvent("L0 Plugin Debug Layer", &L0DebugPayload, + xpti::trace_algorithm_event, xpti_at::active, &Dummy); + + ze_result_t Status = zeInit(0); + if (Status != ZE_RESULT_SUCCESS) { + // Most likey there are no Level Zero devices. + return; + } + + int Foo = 0; + zel_tracer_desc_t TracerDesc = {ZEL_STRUCTURE_TYPE_TRACER_EXP_DESC, nullptr, + &Foo}; + zel_tracer_handle_t Tracer = nullptr; + + Status = zelTracerCreate(&TracerDesc, &Tracer); + + if (Status != ZE_RESULT_SUCCESS || Tracer == nullptr) { + std::cerr << "[WARNING] Failed to create L0 tracer: " << Status << "\n"; + return; + } + + zel_core_callbacks_t Prologue = {}; + zel_core_callbacks_t Epilogue = {}; + +#define _ZE_API(call, domain, cb, params_type) \ + Prologue.domain.cb = [](params_type *Params, ze_result_t, void *, void **) { \ + if (xptiTraceEnabled()) { \ + uint8_t CallStreamID = xptiRegisterStream(L0_CALL_STREAM_NAME); \ + uint8_t DebugStreamID = xptiRegisterStream(L0_DEBUG_STREAM_NAME); \ + CallCorrelationID = xptiGetUniqueId(); \ + DebugCorrelationID = xptiGetUniqueId(); \ + const char *FuncName = #call; \ + xptiNotifySubscribers( \ + CallStreamID, (uint16_t)xpti::trace_point_type_t::function_begin, \ + GCallEvent, nullptr, CallCorrelationID, FuncName); \ + uint32_t FuncID = static_cast(ZEApiKind::call); \ + xpti::function_with_args_t Payload{FuncID, FuncName, Params, nullptr, \ + nullptr}; \ + xptiNotifySubscribers( \ + DebugStreamID, \ + (uint16_t)xpti::trace_point_type_t::function_with_args_begin, \ + GDebugEvent, nullptr, DebugCorrelationID, &Payload); \ + } \ + }; \ + Epilogue.domain.cb = [](params_type *Params, ze_result_t Result, void *, \ + void **) { \ + if (xptiTraceEnabled()) { \ + uint8_t CallStreamID = xptiRegisterStream(L0_CALL_STREAM_NAME); \ + uint8_t DebugStreamID = xptiRegisterStream(L0_DEBUG_STREAM_NAME); \ + const char *FuncName = #call; \ + xptiNotifySubscribers(CallStreamID, \ + (uint16_t)xpti::trace_point_type_t::function_end, \ + GCallEvent, nullptr, CallCorrelationID, FuncName); \ + uint32_t FuncID = static_cast(ZEApiKind::call); \ + xpti::function_with_args_t Payload{FuncID, FuncName, Params, &Result, \ + nullptr}; \ + xptiNotifySubscribers( \ + DebugStreamID, \ + (uint16_t)xpti::trace_point_type_t::function_with_args_end, \ + GDebugEvent, nullptr, DebugCorrelationID, &Payload); \ + } \ + }; + +#include "ze_api.def" + +#undef _ZE_API + + Status = zelTracerSetPrologues(Tracer, &Prologue); + if (Status != ZE_RESULT_SUCCESS) { + std::cerr << "Failed to enable L0 tracing\n"; + std::terminate(); + } + Status = zelTracerSetEpilogues(Tracer, &Epilogue); + if (Status != ZE_RESULT_SUCCESS) { + std::cerr << "Failed to enable L0 tracing\n"; + std::terminate(); + } + + Status = zelTracerSetEnabled(Tracer, true); + if (Status != ZE_RESULT_SUCCESS) { + std::cerr << "Failed to enable L0 tracing\n"; + std::terminate(); + } +#endif +} + +void disableL0Tracing() { + xptiFinalize(L0_CALL_STREAM_NAME); + xptiFinalize(L0_DEBUG_STREAM_NAME); +} diff --git a/sycl/plugins/level_zero/ze_api_generator.py b/sycl/plugins/level_zero/ze_api_generator.py new file mode 100644 index 0000000000000..ca70341c93a41 --- /dev/null +++ b/sycl/plugins/level_zero/ze_api_generator.py @@ -0,0 +1,40 @@ +import re +import sys + +def camel_to_snake(src): + return re.sub(r'(?init(); } - if (std::string_view(StreamName) == "sycl.pi") { + std::string_view NameView{StreamName}; + if (NameView == "sycl.pi") { uint8_t StreamID = xptiRegisterStream(StreamName); xptiRegisterCallback(StreamID, xpti::trace_function_begin, - piBeginEndCallback); + apiBeginEndCallback); xptiRegisterCallback(StreamID, xpti::trace_function_end, - piBeginEndCallback); - } else if (std::string_view(StreamName) == "sycl") { + apiBeginEndCallback); + } else if (NameView == "sycl") { uint8_t StreamID = xptiRegisterStream(StreamName); xptiRegisterCallback(StreamID, xpti::trace_task_begin, taskBeginEndCallback); @@ -89,23 +90,28 @@ XPTI_CALLBACK_API void xptiTraceInit(unsigned int /*major_version*/, waitBeginEndCallback); xptiRegisterCallback(StreamID, xpti::trace_barrier_end, waitBeginEndCallback); + } else if (NameView == "sycl.experimental.level_zero.call") { + uint8_t StreamID = xptiRegisterStream(StreamName); + xptiRegisterCallback(StreamID, xpti::trace_function_begin, + apiBeginEndCallback); + xptiRegisterCallback(StreamID, xpti::trace_function_end, + apiBeginEndCallback); } } XPTI_CALLBACK_API void xptiTraceFinish(const char *) { GWriter->finalize(); } -XPTI_CALLBACK_API void piBeginEndCallback(uint16_t TraceType, - xpti::trace_event_data_t *, - xpti::trace_event_data_t *, - uint64_t /*Instance*/, - const void *UserData) { +XPTI_CALLBACK_API void apiBeginEndCallback(uint16_t TraceType, + xpti::trace_event_data_t *, + xpti::trace_event_data_t *, + uint64_t /*Instance*/, + const void *UserData) { auto [TID, PID, TS] = measure(); if (TraceType == xpti::trace_function_begin) { - GWriter->writeBegin(static_cast(UserData), "Plugin", PID, TID, + GWriter->writeBegin(static_cast(UserData), "API", PID, TID, TS); } else { - GWriter->writeEnd(static_cast(UserData), "Plugin", PID, TID, - TS); + GWriter->writeEnd(static_cast(UserData), "API", PID, TID, TS); } } diff --git a/sycl/tools/sycl-prof/main.cpp b/sycl/tools/sycl-prof/main.cpp index 6b1a96080f25e..37c9041ab2038 100644 --- a/sycl/tools/sycl-prof/main.cpp +++ b/sycl/tools/sycl-prof/main.cpp @@ -45,6 +45,8 @@ int main(int argc, char **argv, char *env[]) { NewEnv.push_back("XPTI_FRAMEWORK_DISPATCHER=libxptifw.so"); NewEnv.push_back("XPTI_SUBSCRIBERS=libsycl_profiler_collector.so"); NewEnv.push_back("XPTI_TRACE_ENABLE=1"); + NewEnv.push_back("SYCL_PI_LEVEL_ZERO_ENABLE_TRACING=1"); + NewEnv.push_back("ZE_ENABLE_TRACING_LAYER=1"); std::vector Args; diff --git a/sycl/tools/sycl-trace/CMakeLists.txt b/sycl/tools/sycl-trace/CMakeLists.txt index e04db01eeecd2..3ebdcb2e8144f 100644 --- a/sycl/tools/sycl-trace/CMakeLists.txt +++ b/sycl/tools/sycl-trace/CMakeLists.txt @@ -10,7 +10,44 @@ link_llvm_libs(sycl-trace LLVMSupport ) -add_library(sycl_pi_trace_collector SHARED pi_trace_collector.cpp) +add_library(sycl_pi_trace_collector SHARED + collector.cpp + pi_trace_collector.cpp + ze_trace_collector.cpp +) + +find_package(Python3 REQUIRED) + +add_custom_target(pi-pretty-printers + COMMAND ${Python3_EXECUTABLE} + ${CMAKE_CURRENT_SOURCE_DIR}/generate_pi_pretty_printers.py + ${sycl_inc_dir}/CL/sycl/detail/pi.h + SOURCES ${sycl_inc_dir}/CL/sycl/detail/pi.h + BYPRODUCTS + ${CMAKE_CURRENT_BINARY_DIR}/pi_structs.hpp + ${CMAKE_CURRENT_BINARY_DIR}/pi_printers.def + ) + +# To get L0 loader +add_dependencies(sycl_pi_trace_collector pi_level_zero) + +if (NOT DEFINED LEVEL_ZERO_INCLUDE_DIR) + set(LEVEL_ZERO_INCLUDE_DIR + "${CMAKE_CURRENT_BINARY_DIR}/../../plugins/level_zero/level_zero_loader/include") +endif() +target_include_directories(sycl_pi_trace_collector PRIVATE + ${LEVEL_ZERO_INCLUDE_DIR} +) + +add_custom_target(ze-pretty-printers + COMMAND ${Python3_EXECUTABLE} + ${CMAKE_CURRENT_SOURCE_DIR}/generate_ze_pretty_printers.py + ${LEVEL_ZERO_INCLUDE_DIR}/level_zero/ze_api.h + DEPENDS pi_level_zero + BYPRODUCTS + ${CMAKE_CURRENT_BINARY_DIR}/ze_printers.def + ) + target_compile_definitions(sycl_pi_trace_collector PRIVATE XPTI_CALLBACK_API_EXPORTS) target_link_libraries(sycl_pi_trace_collector PRIVATE xptifw) if (TARGET OpenCL-Headers) @@ -21,8 +58,10 @@ target_include_directories(sycl_pi_trace_collector PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/../xpti_helpers/" "${sycl_inc_dir}" "${sycl_src_dir}" + "${CMAKE_CURRENT_BINARY_DIR}" ) +add_dependencies(sycl_pi_trace_collector pi-pretty-printers ze-pretty-printers) add_dependencies(sycl-trace sycl_pi_trace_collector) add_dependencies(sycl-toolchain sycl-trace) diff --git a/sycl/tools/sycl-trace/collector.cpp b/sycl/tools/sycl-trace/collector.cpp new file mode 100644 index 0000000000000..0d44b32b30b58 --- /dev/null +++ b/sycl/tools/sycl-trace/collector.cpp @@ -0,0 +1,64 @@ +//==---------------------- collector.cpp -----------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "xpti/xpti_trace_framework.h" + +#include + +sycl::detail::SpinLock GlobalLock; + +bool HasZEPrinter = false; +bool HasPIPrinter = false; + +void zePrintersInit(); +void zePrintersFinish(); +void piPrintersInit(); +void piPrintersFinish(); + +XPTI_CALLBACK_API void piCallback(uint16_t TraceType, + xpti::trace_event_data_t *Parent, + xpti::trace_event_data_t *Event, + uint64_t Instance, const void *UserData); +XPTI_CALLBACK_API void zeCallback(uint16_t TraceType, + xpti::trace_event_data_t *Parent, + xpti::trace_event_data_t *Event, + uint64_t Instance, const void *UserData); + +XPTI_CALLBACK_API void xptiTraceInit(unsigned int /*major_version*/, + unsigned int /*minor_version*/, + const char * /*version_str*/, + const char *StreamName) { + if (std::string_view(StreamName) == "sycl.pi.debug" && + std::getenv("SYCL_TRACE_PI_ENABLE")) { + piPrintersInit(); + uint16_t StreamID = xptiRegisterStream(StreamName); + xptiRegisterCallback(StreamID, xpti::trace_function_with_args_begin, + piCallback); + xptiRegisterCallback(StreamID, xpti::trace_function_with_args_end, + piCallback); + } else if (std::string_view(StreamName) == + "sycl.experimental.level_zero.debug" && + std::getenv("SYCL_TRACE_ZE_ENABLE")) { + zePrintersInit(); + uint16_t StreamID = xptiRegisterStream(StreamName); + xptiRegisterCallback(StreamID, xpti::trace_function_with_args_begin, + zeCallback); + xptiRegisterCallback(StreamID, xpti::trace_function_with_args_end, + zeCallback); + } +} + +XPTI_CALLBACK_API void xptiTraceFinish(const char *StreamName) { + if (std::string_view(StreamName) == "sycl.pi.debug" && + std::getenv("SYCL_TRACE_PI_ENABLE")) + piPrintersFinish(); + else if (std::string_view(StreamName) == + "sycl.experimental.level_zero.debug" && + std::getenv("SYCL_TRACE_ZE_ENABLE")) + zePrintersFinish(); +} diff --git a/sycl/tools/sycl-trace/generate_pi_pretty_printers.py b/sycl/tools/sycl-trace/generate_pi_pretty_printers.py new file mode 100644 index 0000000000000..07b0c3a41c945 --- /dev/null +++ b/sycl/tools/sycl-trace/generate_pi_pretty_printers.py @@ -0,0 +1,54 @@ +import re +import sys + +def generate_pi_pretty_printers(header): + hdr = open("pi_structs.hpp", "w") + hdr.write("//===-------------- pi_structs.hpp - PI Trace Structs ----------------------==//\n") + hdr.write("//\n") + hdr.write("// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.\n") + hdr.write("// See https://llvm.org/LICENSE.txt for license information.\n") + hdr.write("// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception\n") + hdr.write("//\n") + hdr.write("//===----------------------------------------------------------------------===//\n") + hdr.write("// clang-format off\n") + hdr.write("// This file is auto-generated! Do not modify!\n") + hdr.write("#pragma once\n") + printers = open("pi_printers.def", "w") + + matches = re.finditer(r'(pi[a-zA-Z]+)\(\n?\r?([\sa-zA-Z_,\*,=0-9]+)\);', header) + + for match in matches: + api_name = str(match.group(1)) + + if api_name == 'piPluginInit': + continue + + all_args = match.group(2).replace('\n', '').split(',') + + hdr.write("struct __attribute__((packed)) " + api_name + "_args {\n") + + for arg in all_args: + hdr.write(arg.strip() + ";\n") + + hdr.write("};\n") + + arg_names = [] + + for arg in all_args: + name = arg.split("=")[0].strip().split(" ")[-1].replace('*', '') + arg_names.append(name) + + printers.write("case static_cast(sycl::detail::PiApiKind::{}): {{\n".format(api_name)) + printers.write("const auto *Args = reinterpret_cast<{}_args*>(Data->args_data);\n".format(api_name)) + for name in arg_names: + printers.write('std::cout << " {}: " << Args->{} << "\\n";\n'.format(name, name)) + printers.write("break;\n") + printers.write("}\n") + +if __name__ == "__main__": + """ + Usage: python generate_pi_pretty_printers.py path/to/pi.h + """ + with open(sys.argv[1], 'r') as f: + header = f.read() + generate_pi_pretty_printers(header) diff --git a/sycl/tools/sycl-trace/generate_ze_pretty_printers.py b/sycl/tools/sycl-trace/generate_ze_pretty_printers.py new file mode 100644 index 0000000000000..40d1b600b9c5f --- /dev/null +++ b/sycl/tools/sycl-trace/generate_ze_pretty_printers.py @@ -0,0 +1,59 @@ +import re +import sys + +def camel_to_snake(src): + return re.sub(r'(?(ZEApiKind::{}): {{\n".format(api_name)) + printers.write("const auto *Args = reinterpret_cast<{}*>(Data->args_data);\n".format(param_type)) + for arg in args: + arg_name = arg.strip().split(" ")[-1].replace('*', '') + arg_types = [ x.strip() for x in arg.strip().split(" ")[:-1]] + printers.write("PrintOffset();\n") + scalar = ["size_t*", "void**", "uint32_t*", "uint64_t*"] + if any(item in scalar for item in arg_types): + printers.write('std::cout << "{}: " << *(Args->{}) << "\\n";\n'.format(arg_name[1:], arg_name)) + else: + printers.write(' std::cout << "{}: " << Args->{} << "\\n";\n'.format(arg_name, arg_name)) + printers.write("break;\n") + printers.write("}\n") + + printers.close() + +if __name__ == "__main__": + """ + Usage: python generate_pi_pretty_printers.py path/to/ze_api.h + """ + with open(sys.argv[1], 'r') as f: + header = f.read() + generate_ze_pretty_printers(header) + diff --git a/sycl/tools/sycl-trace/main.cpp b/sycl/tools/sycl-trace/main.cpp index 13a9524a54c3c..433f8954a080b 100644 --- a/sycl/tools/sycl-trace/main.cpp +++ b/sycl/tools/sycl-trace/main.cpp @@ -14,14 +14,24 @@ using namespace llvm; -enum ModeKind { PI }; +enum ModeKind { PI, ZE }; +enum PrintFormatKind { PRETTY_COMPACT, PRETTY_VERBOSE, CLASSIC }; int main(int argc, char **argv, char *env[]) { - cl::opt Mode( - "mode", cl::desc("Set tracing mode:"), + cl::list Modes( + cl::desc("Available tracing modes:"), cl::values( // TODO graph dot - clEnumValN(PI, "plugin", "Trace Plugin Interface calls"))); + clEnumValN(PI, "plugin", "Trace Plugin Interface calls"), + clEnumValN(ZE, "level_zero", "Trace Level Zero calls"))); + cl::opt PrintFormat( + "print-format", cl::desc("Print format"), + cl::values( + clEnumValN(PRETTY_COMPACT, "compact", "Human readable compact"), + clEnumValN(PRETTY_VERBOSE, "verbose", "Human readable verbose"), + clEnumValN( + CLASSIC, "classic", + "Similar to SYCL_PI_TRACE, only compatible with PI layer"))); cl::opt TargetExecutable( cl::Positional, cl::desc(""), cl::Required); cl::list Argv(cl::ConsumeAfter, @@ -41,6 +51,39 @@ int main(int argc, char **argv, char *env[]) { NewEnv.push_back("XPTI_SUBSCRIBERS=libsycl_pi_trace_collector.so"); NewEnv.push_back("XPTI_TRACE_ENABLE=1"); + const auto EnablePITrace = [&]() { + NewEnv.push_back("SYCL_TRACE_PI_ENABLE=1"); + }; + const auto EnableZETrace = [&]() { + NewEnv.push_back("SYCL_TRACE_ZE_ENABLE=1"); + NewEnv.push_back("SYCL_PI_LEVEL_ZERO_ENABLE_TRACING=1"); + NewEnv.push_back("ZE_ENABLE_TRACING_LAYER=1"); + }; + + for (auto Mode : Modes) { + switch (Mode) { + case PI: + EnablePITrace(); + break; + case ZE: + EnableZETrace(); + break; + } + } + + if (PrintFormat == CLASSIC) { + NewEnv.push_back("SYCL_TRACE_PRINT_FORMAT=classic"); + } else if (PrintFormat == PRETTY_VERBOSE) { + NewEnv.push_back("SYCL_TRACE_PRINT_FORMAT=verbose"); + } else { + NewEnv.push_back("SYCL_TRACE_PRINT_FORMAT=compact"); + } + + if (Modes.size() == 0) { + EnablePITrace(); + EnableZETrace(); + } + std::vector Args; Args.push_back(TargetExecutable); diff --git a/sycl/tools/sycl-trace/pi_trace_collector.cpp b/sycl/tools/sycl-trace/pi_trace_collector.cpp index 2f277ef3e4598..9cce40a005765 100644 --- a/sycl/tools/sycl-trace/pi_trace_collector.cpp +++ b/sycl/tools/sycl-trace/pi_trace_collector.cpp @@ -1,4 +1,4 @@ -//==----------- pi_trace.cpp.cpp -------------------------------------------==// +//==---------------------- pi_trace_collector.cpp --------------------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -6,14 +6,15 @@ // //===----------------------------------------------------------------------===// -/// \file pi_trace.cpp -/// A sample XPTI subscriber to demonstrate how to collect PI function call -/// arguments. +/// \file pi_trace_collector.cpp +/// Routines to collect and print Plugin Interface calls. #include "xpti/xpti_trace_framework.h" #include "pi_arguments_handler.hpp" +#include "pi_structs.hpp" +#include #include #include @@ -22,63 +23,196 @@ #include #include -static uint8_t GStreamID = 0; -std::mutex GIOMutex; - -sycl::xpti_helpers::PiArgumentsHandler ArgHandler; - -// The lone callback function we are going to use to demonstrate how to attach -// the collector to the running executable -XPTI_CALLBACK_API void tpCallback(uint16_t trace_type, - xpti::trace_event_data_t *parent, - xpti::trace_event_data_t *event, - uint64_t instance, const void *user_data); - -// Based on the documentation, every subscriber MUST implement the -// xptiTraceInit() and xptiTraceFinish() APIs for their subscriber collector to -// be loaded successfully. -XPTI_CALLBACK_API void xptiTraceInit(unsigned int /*major_version*/, - unsigned int /*minor_version*/, - const char * /*version_str*/, - const char *stream_name) { - if (std::string_view(stream_name) == "sycl.pi.debug") { - GStreamID = xptiRegisterStream(stream_name); - xptiRegisterCallback(GStreamID, xpti::trace_function_with_args_begin, - tpCallback); - xptiRegisterCallback(GStreamID, xpti::trace_function_with_args_end, - tpCallback); +extern sycl::detail::SpinLock GlobalLock; +extern bool HasZEPrinter; +extern bool HasPIPrinter; + +using HeaderPrinterT = + std::function; + +static sycl::xpti_helpers::PiArgumentsHandler *ArgHandler = nullptr; +static HeaderPrinterT *HeaderPrinter = nullptr; +static std::function *ResultPrinter = nullptr; + +static std::string getResult(pi_result Res) { + switch (Res) { + case PI_SUCCESS: + return "PI_SUCCESS"; + case PI_INVALID_KERNEL_NAME: + return "PI_INVALID_KERNEL_NAME"; + case PI_INVALID_OPERATION: + return "CL_INVALID_OPERATION"; + case PI_INVALID_KERNEL: + return "PI_INVALID_KERNEL"; + case PI_INVALID_QUEUE_PROPERTIES: + return "PI_INVALID_QUEUE_PROPERTIES"; + case PI_INVALID_VALUE: + return "PI_INVALID_VALUE"; + case PI_INVALID_CONTEXT: + return "PI_INVALID_CONTEXT"; + case PI_INVALID_PLATFORM: + return "PI_INVALID_PLATFORM"; + case PI_INVALID_DEVICE: + return "PI_INVALID_DEVICE"; + case PI_INVALID_BINARY: + return "PI_INVALID_BINARY"; + case PI_INVALID_QUEUE: + return "PI_INVALID_COMMAND_QUEUE"; + case PI_OUT_OF_HOST_MEMORY: + return "PI_OUT_OF_HOST_MEMORY"; + case PI_INVALID_PROGRAM: + return "PI_INVALID_PROGRAM"; + case PI_INVALID_PROGRAM_EXECUTABLE: + return "PI_INVALID_PROGRAM_EXECUTABLE"; + case PI_INVALID_SAMPLER: + return "PI_INVALID_SAMPLER"; + case PI_INVALID_BUFFER_SIZE: + return "PI_INVALID_BUFFER_SIZE"; + case PI_INVALID_MEM_OBJECT: + return "PI_INVALID_MEM_OBJECT"; + case PI_OUT_OF_RESOURCES: + return "PI_OUT_OF_RESOURCES"; + case PI_INVALID_EVENT: + return "PI_INVALID_EVENT"; + case PI_INVALID_EVENT_WAIT_LIST: + return "PI_INVALID_EVENT_WAIT_LIST"; + case PI_MISALIGNED_SUB_BUFFER_OFFSET: + return "PI_MISALIGNED_SUB_BUFFER_OFFSET"; + case PI_BUILD_PROGRAM_FAILURE: + return "PI_BUILD_PROGRAM_FAILURE"; + case PI_INVALID_WORK_GROUP_SIZE: + return "PI_INVALID_WORK_GROUP_SIZE"; + case PI_COMPILER_NOT_AVAILABLE: + return "PI_COMPILER_NOT_AVAILABLE"; + case PI_PROFILING_INFO_NOT_AVAILABLE: + return "PI_PROFILING_INFO_NOT_AVAILABLE"; + case PI_DEVICE_NOT_FOUND: + return "PI_DEVICE_NOT_FOUND"; + case PI_INVALID_WORK_ITEM_SIZE: + return "PI_INVALID_WORK_ITEM_SIZE"; + case PI_INVALID_WORK_DIMENSION: + return "PI_INVALID_WORK_DIMENSION"; + case PI_INVALID_KERNEL_ARGS: + return "PI_INVALID_KERNEL_ARGS"; + case PI_INVALID_IMAGE_SIZE: + return "PI_INVALID_IMAGE_SIZE"; + case PI_INVALID_ARG_VALUE: + return "PI_INVALID_ARG_VALUE"; + case PI_INVALID_IMAGE_FORMAT_DESCRIPTOR: + return "PI_INVALID_IMAGE_FORMAT_DESCRIPTOR"; + case PI_IMAGE_FORMAT_NOT_SUPPORTED: + return "PI_IMAGE_FORMAT_NOT_SUPPORTED"; + case PI_MEM_OBJECT_ALLOCATION_FAILURE: + return "PI_MEM_OBJECT_ALLOCATION_FAILURE"; + case PI_LINK_PROGRAM_FAILURE: + return "PI_LINK_PROGRAM_FAILURE"; + case PI_COMMAND_EXECUTION_FAILURE: + return "PI_COMMAND_EXECUTION_FAILURE"; + case PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE: + return "PI_FUNCTION_ADDRESS_IS_NOT_AVAILABLE"; + case PI_ERROR_UNKNOWN: + return "PI_ERROR_UNKNOWN"; + } + + return "UNKNOWN RESULT"; +} + +static void setupClassicPrinter() { + ArgHandler = new sycl::xpti_helpers::PiArgumentsHandler(); #define _PI_API(api) \ - ArgHandler.set##_##api( \ + ArgHandler->set##_##api( \ [](const pi_plugin &, std::optional, auto &&...Args) { \ std::cout << "---> " << #api << "(" \ << "\n"; \ sycl::detail::pi::printArgs(Args...); \ - std::cout << ") ---> "; \ }); #include #undef _PI_API + + ResultPrinter = new std::function( + [](pi_result Res) { std::cout << ") ---> " << Res << std::endl; }); + HeaderPrinter = new std::function( + [](const pi_plugin &Plugin, const xpti::function_with_args_t *Data) { + ArgHandler->handle(Data->function_id, Plugin, std::nullopt, + Data->args_data); + }); +} + +static void setupPrettyPrinter(bool Verbose) { + HeaderPrinter = new std::function( + [Verbose](const pi_plugin &, const xpti::function_with_args_t *Data) { + if (Verbose) { + std::string Source = ""; + size_t Line = 0; + + auto *Payload = xptiQueryPayloadByUID(xptiGetUniversalId()); + + if (Payload) { + if (Payload->source_file != nullptr) { + Source = Payload->source_file; + Line = Payload->line_no; + } + } + + auto TID = std::this_thread::get_id(); + std::cout << "[PI:TID " << TID << ":"; + std::cout << Source << ":" << Line << "]\n"; + } else { + std::cout << "[PI] "; + } + std::cout << Data->function_name << "(\n"; + switch (Data->function_id) { +#include "pi_printers.def" + } + std::cout << ")"; + + if (HasZEPrinter) { + std::cout << " {" << std::endl; + } + }); + ResultPrinter = new std::function([](pi_result Res) { + if (HasZEPrinter) { + std::cout << "}"; + } + std::cout << " ---> " << getResult(Res) << "\n" << std::endl; + }); +} + +void piPrintersInit() { + HasPIPrinter = true; + std::string_view PrinterType(std::getenv("SYCL_TRACE_PRINT_FORMAT")); + + if (PrinterType == "classic") { + setupClassicPrinter(); + } else if (PrinterType == "verbose") { + setupPrettyPrinter(/*verbose*/ true); + } else if (PrinterType == "compact") { + setupPrettyPrinter(/*verbose*/ false); } } -XPTI_CALLBACK_API void xptiTraceFinish(const char * /*stream_name*/) { - // NOP +void piPrintersFinish() { + if (ArgHandler) + delete ArgHandler; + delete HeaderPrinter; + delete ResultPrinter; } -XPTI_CALLBACK_API void tpCallback(uint16_t TraceType, +XPTI_CALLBACK_API void piCallback(uint16_t TraceType, xpti::trace_event_data_t * /*Parent*/, xpti::trace_event_data_t * /*Event*/, uint64_t /*Instance*/, const void *UserData) { - if (TraceType == xpti::trace_function_with_args_end) { - // Lock while we print information - std::lock_guard Lock(GIOMutex); + if (!HeaderPrinter || !ResultPrinter) + return; - const auto *Data = - static_cast(UserData); + // Lock while we print information + std::lock_guard _{GlobalLock}; + const auto *Data = static_cast(UserData); + if (TraceType == xpti::trace_function_with_args_begin) { const auto *Plugin = static_cast(Data->user_data); - - ArgHandler.handle(Data->function_id, *Plugin, std::nullopt, - Data->args_data); - std::cout << *static_cast(Data->ret_data) << "\n"; + (*HeaderPrinter)(*Plugin, Data); + } else if (TraceType == xpti::trace_function_with_args_end) { + (*ResultPrinter)(*static_cast(Data->ret_data)); } } diff --git a/sycl/tools/sycl-trace/ze_trace_collector.cpp b/sycl/tools/sycl-trace/ze_trace_collector.cpp new file mode 100644 index 0000000000000..c9be5b9c20d07 --- /dev/null +++ b/sycl/tools/sycl-trace/ze_trace_collector.cpp @@ -0,0 +1,323 @@ +//==----------- ze_trace_collector.cpp -------------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +/// \file ze_trace_collector.cpp +/// Routines to collect and print Level Zero API calls. + +#include "xpti/xpti_trace_framework.h" + +#include + +#include + +#include +#include +#include +#include +#include + +extern sycl::detail::SpinLock GlobalLock; + +extern bool HasZEPrinter; +extern bool HasPIPrinter; + +enum class ZEApiKind { +#define _ZE_API(call, domain, cb, params_type) call, +#include "../../plugins/level_zero/ze_api.def" +#undef _ZE_API +}; + +bool PrintVerbose = false; + +static std::string getResult(ze_result_t Res) { + std::string ResultStr; + switch (Res) { + case ZE_RESULT_SUCCESS: + ResultStr = "ZE_RESULT_SUCCESS"; + if (PrintVerbose) + ResultStr += " (success)"; + break; + case ZE_RESULT_NOT_READY: + ResultStr = "ZE_RESULT_NOT_READY"; + if (PrintVerbose) + ResultStr += " (synchronization primitive not signaled)"; + break; + case ZE_RESULT_ERROR_DEVICE_LOST: + ResultStr = "ZE_RESULT_ERROR_DEVICE_LOST"; + if (PrintVerbose) + ResultStr += + " (device hung, reset, was removed, or driver update occurred)"; + break; + case ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY: + ResultStr = "ZE_RESULT_ERROR_OUT_OF_HOST_MEMORY"; + if (PrintVerbose) + ResultStr += " (insufficient host memory to satisfy call)"; + break; + case ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY: + ResultStr = "ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY"; + if (PrintVerbose) + ResultStr += " (insufficient device memory to satisfy call)"; + break; + case ZE_RESULT_ERROR_MODULE_BUILD_FAILURE: + ResultStr = "ZE_RESULT_ERROR_MODULE_BUILD_FAILURE"; + if (PrintVerbose) + ResultStr += + " (error occurred when building module, see build log for details)"; + break; + case ZE_RESULT_ERROR_MODULE_LINK_FAILURE: + ResultStr = "ZE_RESULT_ERROR_MODULE_LINK_FAILURE"; + if (PrintVerbose) + ResultStr += + " (error occurred when linking modules, see build log for details)"; + break; + case ZE_RESULT_ERROR_DEVICE_REQUIRES_RESET: + ResultStr = "ZE_RESULT_ERROR_DEVICE_REQUIRES_RESET"; + if (PrintVerbose) + ResultStr += " (device requires a reset)"; + break; + case ZE_RESULT_ERROR_DEVICE_IN_LOW_POWER_STATE: + ResultStr = "ZE_RESULT_ERROR_DEVICE_IN_LOW_POWER_STATE"; + if (PrintVerbose) + ResultStr += " (device currently in low power state)"; + break; + case ZE_RESULT_ERROR_INSUFFICIENT_PERMISSIONS: + ResultStr = "ZE_RESULT_ERROR_INSUFFICIENT_PERMISSIONS"; + if (PrintVerbose) + ResultStr += " (access denied due to permission level)"; + break; + case ZE_RESULT_ERROR_NOT_AVAILABLE: + ResultStr = "ZE_RESULT_ERROR_NOT_AVAILABLE"; + if (PrintVerbose) + ResultStr += " (resource already in use and simultaneous access not " + "allowed or resource was removed)"; + break; + case ZE_RESULT_ERROR_DEPENDENCY_UNAVAILABLE: + ResultStr = "ZE_RESULT_ERROR_DEPENDENCY_UNAVAILABLE"; + if (PrintVerbose) + ResultStr += " (external required dependency is unavailable or missing)"; + break; + case ZE_RESULT_ERROR_UNINITIALIZED: + ResultStr = "ZE_RESULT_ERROR_UNINITIALIZED"; + if (PrintVerbose) + ResultStr += " (driver is not initialized)"; + break; + case ZE_RESULT_ERROR_UNSUPPORTED_VERSION: + ResultStr = "ZE_RESULT_ERROR_UNSUPPORTED_VERSION"; + if (PrintVerbose) + ResultStr += " (generic error code for unsupported versions)"; + break; + case ZE_RESULT_ERROR_UNSUPPORTED_FEATURE: + ResultStr = "ZE_RESULT_ERROR_UNSUPPORTED_FEATURE"; + if (PrintVerbose) + ResultStr += " (generic error code for unsupported features)"; + break; + case ZE_RESULT_ERROR_INVALID_ARGUMENT: + ResultStr = "ZE_RESULT_ERROR_INVALID_ARGUMENT"; + if (PrintVerbose) + ResultStr += " (generic error code for invalid arguments)"; + break; + case ZE_RESULT_ERROR_INVALID_NULL_HANDLE: + ResultStr = "ZE_RESULT_ERROR_INVALID_NULL_HANDLE"; + if (PrintVerbose) + ResultStr += " (handle argument is not valid)"; + break; + case ZE_RESULT_ERROR_HANDLE_OBJECT_IN_USE: + ResultStr = "ZE_RESULT_ERROR_HANDLE_OBJECT_IN_USE"; + if (PrintVerbose) + ResultStr += " (object pointed to by handle still in-use by device)"; + break; + case ZE_RESULT_ERROR_INVALID_NULL_POINTER: + ResultStr = "ZE_RESULT_ERROR_INVALID_NULL_POINTER"; + if (PrintVerbose) + ResultStr += " (pointer argument may not be nullptr)"; + break; + case ZE_RESULT_ERROR_INVALID_SIZE: + ResultStr = "ZE_RESULT_ERROR_INVALID_SIZE"; + if (PrintVerbose) + ResultStr += " (size argument is invalid (e.g., must not be zero))"; + break; + case ZE_RESULT_ERROR_UNSUPPORTED_SIZE: + ResultStr = "ZE_RESULT_ERROR_UNSUPPORTED_SIZE"; + if (PrintVerbose) + ResultStr += + " (size argument is not supported by the device (e.g., too large))"; + break; + case ZE_RESULT_ERROR_UNSUPPORTED_ALIGNMENT: + ResultStr = "ZE_RESULT_ERROR_UNSUPPORTED_ALIGNMENT"; + if (PrintVerbose) + ResultStr += " (alignment argument is not supported by the device (e.g., " + "too small))"; + break; + case ZE_RESULT_ERROR_INVALID_SYNCHRONIZATION_OBJECT: + ResultStr = "ZE_RESULT_ERROR_INVALID_SYNCHRONIZATION_OBJECT"; + if (PrintVerbose) + ResultStr += " (synchronization object in invalid state)"; + break; + case ZE_RESULT_ERROR_INVALID_ENUMERATION: + ResultStr = "ZE_RESULT_ERROR_INVALID_ENUMERATION"; + if (PrintVerbose) + ResultStr += " (enumerator argument is not valid)"; + break; + case ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION: + ResultStr = "ZE_RESULT_ERROR_UNSUPPORTED_ENUMERATION"; + if (PrintVerbose) + ResultStr += " (enumerator argument is not supported by the device)"; + break; + case ZE_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT: + ResultStr = "ZE_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT"; + if (PrintVerbose) + ResultStr += " (image format is not supported by the device)"; + break; + case ZE_RESULT_ERROR_INVALID_NATIVE_BINARY: + ResultStr += "ZE_RESULT_ERROR_INVALID_NATIVE_BINARY"; + if (PrintVerbose) + ResultStr += " (native binary is not supported by the device)"; + break; + case ZE_RESULT_ERROR_INVALID_GLOBAL_NAME: + ResultStr = "ZE_RESULT_ERROR_INVALID_GLOBAL_NAME"; + if (PrintVerbose) + ResultStr += " (global variable is not found in the module)"; + break; + case ZE_RESULT_ERROR_INVALID_KERNEL_NAME: + ResultStr = "ZE_RESULT_ERROR_INVALID_KERNEL_NAME"; + if (PrintVerbose) + ResultStr += " (kernel name is not found in the module)"; + break; + case ZE_RESULT_ERROR_INVALID_FUNCTION_NAME: + ResultStr = "ZE_RESULT_ERROR_INVALID_FUNCTION_NAME"; + if (PrintVerbose) + ResultStr += " (function name is not found in the module)"; + break; + case ZE_RESULT_ERROR_INVALID_GROUP_SIZE_DIMENSION: + ResultStr = "ZE_RESULT_ERROR_INVALID_GROUP_SIZE_DIMENSION"; + if (PrintVerbose) + ResultStr += + " (group size dimension is not valid for the kernel or device)"; + break; + case ZE_RESULT_ERROR_INVALID_GLOBAL_WIDTH_DIMENSION: + ResultStr = "ZE_RESULT_ERROR_INVALID_GLOBAL_WIDTH_DIMENSION"; + if (PrintVerbose) + ResultStr += + " (global width dimension is not valid for the kernel or device)"; + break; + case ZE_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX: + ResultStr = "ZE_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_INDEX"; + if (PrintVerbose) + ResultStr += " (kernel argument index is not valid for kernel)"; + break; + case ZE_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE: + ResultStr = "ZE_RESULT_ERROR_INVALID_KERNEL_ARGUMENT_SIZE"; + if (PrintVerbose) + ResultStr += " (kernel argument size does not match kernel)"; + break; + case ZE_RESULT_ERROR_INVALID_KERNEL_ATTRIBUTE_VALUE: + ResultStr = "ZE_RESULT_ERROR_INVALID_KERNEL_ATTRIBUTE_VALUE"; + if (PrintVerbose) + ResultStr += + " (value of kernel attribute is not valid for the kernel or device)"; + break; + case ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED: + ResultStr = "ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED"; + if (PrintVerbose) + ResultStr += " (module with imports needs to be linked before kernels " + "can be created from it)"; + break; + case ZE_RESULT_ERROR_INVALID_COMMAND_LIST_TYPE: + ResultStr = "ZE_RESULT_ERROR_INVALID_COMMAND_LIST_TYPE"; + if (PrintVerbose) + ResultStr += " (command list type does not match command queue type)"; + break; + case ZE_RESULT_ERROR_OVERLAPPING_REGIONS: + ResultStr = "ZE_RESULT_ERROR_OVERLAPPING_REGIONS"; + if (PrintVerbose) + ResultStr += + " (copy operations do not support overlapping regions of memory)"; + break; + default: + ResultStr = "UNKNOWN ERROR"; + break; + } + + return ResultStr; +} + +XPTI_CALLBACK_API void zeCallback(uint16_t TraceType, + xpti::trace_event_data_t * /*Parent*/, + xpti::trace_event_data_t * /*Event*/, + uint64_t /*Instance*/, const void *UserData) { + std::lock_guard _{GlobalLock}; + const auto *Data = static_cast(UserData); + const auto PrintPrefix = [] { + if (HasPIPrinter) + std::cout << "* "; + }; + if (TraceType == xpti::trace_function_with_args_begin) { + + const auto PrintOffset = [PrintPrefix]() { + PrintPrefix(); + std::cout << " "; + }; + + PrintPrefix(); + if (PrintVerbose) { + std::string Source = ""; + size_t Line = 0; + + auto *Payload = xptiQueryPayloadByUID(xptiGetUniversalId()); + + if (Payload) { + if (Payload->source_file != nullptr) { + Source = Payload->source_file; + Line = Payload->line_no; + } + } + + auto TID = std::this_thread::get_id(); + std::cout << "[L0:TID " << TID << ":"; + std::cout << Source << ":" << Line << "]\n"; + PrintPrefix(); + } else { + std::cout << "[L0] "; + } + + std::cout << Data->function_name << "(\n"; + + switch (Data->function_id) { +#include "ze_printers.def" + default: + break; // unknown API + } + + if (HasPIPrinter) { + std::cout << "* "; + } + std::cout << std::flush; + } else if (TraceType == xpti::trace_function_with_args_end) { + std::cout << ") ---> " + << getResult(*static_cast(Data->ret_data)) + << std::endl; + PrintPrefix(); + std::cout << std::endl; + } +} + +void zePrintersInit() { + HasZEPrinter = true; + + std::string_view PrinterType(std::getenv("SYCL_TRACE_PRINT_FORMAT")); + if (PrinterType == "classic") { + std::cerr << "Classic output is unsupported for Level Zero\n"; + } else if (PrinterType == "verbose") { + PrintVerbose = true; + } else if (PrinterType == "compact") { + PrintVerbose = false; + } +} + +// For unification purpose +void zePrintersFinish() {}