API tracing provides a way for tools to receive notifications of L0 API calls made by an application. These notifications are provided via callback functions that provide direct access to the input and output parameters which can be viewed and modified. Tools may also use these notifications as triggers to block and inject new API calls into the command stream, such as metrics.
Tracing provides for a tool to create one or more tracers. A tracer is identified by a tracer handle, and has associated with it a set of prologue functions for each L0 API function that can be invoked at the beginning of these function calls, and a matching set of epilogue functions that can be invoked at the end of L0 API calls.
In summary, this tracing implementation provides functions to create one or more tracers, acquiring a tracer handle for each, then registering a set of prologue and epilogue functions for each tracer handle, enabling and disabling a tracer, and for destroying a tracer.
Tracing is implemented as a layer in the loader. This tracing layer must be enabled by setting the environment variable ZE_ENABLE_TRACING_LAYER to 1. This environment variable must be defined in the application process's context before that process calls zeInit().
The API for using this tracing implementation is this header file below. Please examine that header file for tracing API details.
include/level_zero/layers/zel_tracing_api.h
zelTracerCreate returns a tracer handle representing that tracer. A tracer represents a set of prologue callbacks and epilogue callbacks. See Creation, Registration, Enabling, Disabling and Destruction below.
The per-L0-API function structures used to pass arguments into callback handlers (..params_t) are defined in the ze_api.h. Only L0 API functions declared in ze_api.h can be traced.
include/level_zero/ze_api.h
The structure used to declare sets of prologue and epilogue callbacks (zel_core_callbacks_t) is defined in zel_tracing_api.h. This currently references the ze_core_callbacks_t declaration in ze_api.h. Please use the zel_core_callbacks_t definition.
include/level_zero/layers/zel_tracing_api.h
Users of tracing must first create one or more tracers, using zelTracerCreate.
Users of tracing may independently register for enter and exit callbacks for individual L0 API calls. These callbacks are associated with a tracer handle returned from the zelTracerCreate. The set of tracers applies across all drivers and devices. There are now TWO classes of interfaces for registering callbacks:
For this class of registration functions, the zet_core_callbacks_t argument is a structure of pointers to callback handler functions. A nullptr value for one of these entries means no callback handler is defined for the corresponding L0 API function.
This zet_core_callbacks_t structure is not extensible in a way that could support binary compatibility as new APIs are added to the L0 specification. As a consequence, these registration functions are deprecated and will be removed in the future. The definition of the zet_core_callbacks_t structure is frozen as of the L0 API 1.0 specification. Any new L0 API functions added since version 1.0 will not be traceable using these registration functions.
-
zelTracerSetPrologues is used to specify all the enter callbacks for a tracer handle.
-
zelTracerSetEpilogues is used to specify all the exit callbacks for a tracer handle.
These functions can be called only when the tracer specified by tracer handle is in the disabled state. See Enabling, Disabling and Destruction below.
A new set of registration functions has been added, one for each L0 API function. These registration functions have the general form:
-zelTracerXRegisterCallback(zel_tracer_handle_t hTracer, zel_tracer_reg_t callback_type, ze_pfnXCb_t callback_handler_function)
The zel_tracer_reg_t value can be either ZEL_REGISTER_PROLOGUE or ZEL_REGISTER_EPILOGUE, specifying whether a prologue or an epilogue handler is being registered.
These new registration functions are defined in the header file include/level_zero/layers/tracing/zel_tracing_register_cb.h. This header file also includes prototypes for callback handler functions and Xcb_t structure declarations for L0 API functions that have been added since specification version 1.0. When the older zelTracerSetPrologues and zelTracerSetEpilogues functions are removed, the zet_core_callbacks_t structure will also be removed, and all ze_pfnXCb_t and XCb_t declarations that are in ze_api.h will be relocated into this header file.
If the callback_handler_function pointer is NULL, then no callback handler will be registered for that API function.
These register callback functions can be called only when the hTracer argument references a tracer that is in the disabled state.
zelTracerResetAllCallbacks(zel_tracer_handle_t hTracer) can be used to set ALL prologue and epilogue callback handlers to NULL.
Callback handlers are functions that are implemented by the application, and registered through either the set epilogue/set prologue functions, or the RegisterCallback APIs. The ze_api.h header file or the zel_tracing_register_cb.h header file contain prototype declarations for these functions. Generally, these functions take the following parameters:
- __params__ : a structure capturing pointers to the input and output parameters of the current instance.
- __result__ : the current value of the return value.
- __pTracerUserData__ : a per-tracer per-API pointer to user's data that is passed into the callback handler functions.
- __ppTracerInstanceUserData__ : a per-tracer, per-instance, per-thread storage location; typically used for passing data from the prologue to the epilogue. See example below.
The tracer is created in a disabled state and must be explicitly enabled by calling zelTracerSetEnabled. The implementation guarantees that prologue and epilogue handlers for a given L0 API function will always be executed in pairs; i.e.
-
if the prologue function was called and there is a corresponding epilogue function, then the epilogue is guaranteed to be called, even if another thread disabled the tracer between execution
-
if the prologue function was not called then the epilogue function is guaranteed not to be called, even if another thread enabled the tracer between execution
The tracer must be disabled by the application before the tracer is destroyed. If multiple threads are in-flight, then callbacks that are in progress for that tracer will continue to execute even after the tracer is disabled; the implementation will stall and wait for any outstanding threads executing a tracer prologue or epilogue functions to complete those during zelTracerDestroy of that tracer.
The following pseudo-code demonstrates a basic usage of API tracing:
#include "level_zero/ze_api.h"
#include "level_zero/layers/zel_tracing_api.h"
typedef struct _my_tracer_data_t
{
uint32_t instance;
} my_tracer_data_t;
typedef struct _my_instance_data_t
{
clock_t start;
} my_instance_data_t;
void OnEnterCommandListAppendLaunchKernel(
ze_command_list_append_launch_kernel_params_t* params,
ze_result_t result,
void* pTracerUserData,
void** ppTracerInstanceUserData )
{
my_instance_data_t* instance_data = malloc( sizeof(my_instance_data_t) );
*ppTracerInstanceUserData = instance_data;
instance_data->start = clock();
}
void OnExitCommandListAppendLaunchKernel(
ze_command_list_append_launch_kernel_params_t* params,
ze_result_t result,
void* pTracerUserData,
void** ppTracerInstanceUserData )
{
clock_t end = clock();
my_tracer_data_t* tracer_data = (my_tracer_data_t*)pTracerUserData;
my_instance_data_t* instance_data = *(my_instance_data_t**)ppTracerInstanceUserData;
float time = 1000.f * ( end - instance_data->start ) / CLOCKS_PER_SEC;
printf("zeCommandListAppendLaunchKernel #%d takes %.4f msn", tracer_data->instance++, time);
free(instance_data);
}
// An example using deprecated setepilogue/setprologue functions
void TracingExample1( ... )
{
my_tracer_data_t tracer_data = {};
zel_tracer_desc_t tracer_desc;
tracer_desc.stype = ZEL_STRUCTURE_TYPE_TRACER_DESC;
tracer_desc.pUserData = &tracer_data;
zel_tracer_handle_t hTracer;
zelTracerCreate(hDevice, &tracer_desc, &hTracer);
// Set all callbacks
zel_core_callbacks_t prologCbs = {};
zel_core_callbacks_t epilogCbs = {};
prologCbs.CommandList.pfnAppendLaunchKernelCb = OnEnterCommandListAppendLaunchKernel;
epilogCbs.CommandList.pfnAppendLaunchKernelCb = OnExitCommandListAppendLaunchKernel;
zelTracerSetPrologues(hTracer, &prologCbs);
zelTracerSetEpilogues(hTracer, &epilogCbs);
zelTracerSetEnabled(hTracer, true);
zeCommandListAppendLaunchKernel(hCommandList, hFunction, &launchArgs, nullptr, 0, nullptr);
zelTracerSetEnabled(hTracer, false);
zelTracerDestroy(hTracer);
}
// an example using RegisterCallback functions
void TracingExample2( ... )
{
my_tracer_data_t tracer_data = {};
zel_tracer_desc_t tracer_desc;
tracer_desc.stype = ZEL_STRUCTURE_TYPE_TRACER_DESC;
tracer_desc.pUserData = &tracer_data;
zel_tracer_handle_t hTracer;
zelTracerCreate(hDevice, &tracer_desc, &hTracer);
zelTracerCommandListAppendLaunchKernelRegisterCallback(hTracer, ZEL_REGISTER_PROLOGUE, OnEnterCommandListAppendLaunchKernel);
zelTracerCommandListAppendLaunchKernelRegisterCallback(hTracer, ZEL_REGISTER_EPILOGUE, OnExitCommandListAppendLaunchKernel);
zelTracerSetEnabled(hTracer, true);
zeCommandListAppendLaunchKernel(hCommandList, hFunction, &launchArgs, nullptr, 0, nullptr);
zelTracerSetEnabled(hTracer, false);
zelTracerDestroy(hTracer);
}