Tools Programming Guide¶

Initialization¶

Environment Variables¶

The following environment variables are required to be enabled during zeInit for the respective feature.

Category	Name	Values	Description
Tools	ZET_ENABLE_API_TRACING	{0, 1}	Enables driver instrumentation for API tracing
	ZET_ENABLE_METRICS	{0, 1}	Enables driver instrumentation and dependencies for device metrics
	ZET_ENABLE_PROGRAM_INSTRUMENTATION	{0, 1}	Enables driver instrumentation and dependencies for program instrumentation
	ZET_ENABLE_PROGRAM_DEBUGGING	{0, 1}	Enables driver instrumentation and dependencies for program debugging

API Tracing¶

Introduction¶

API tracing provides a way for tools to receive notifications of API calls made by an application. The callbacks provide direct access to the input and output parameters for viewing and modification. Tools may also use these notifications as triggers to block and inject new API calls into the command stream, such as metrics.

Registration¶

Tools may independently register for enter and exit callbacks for individual API calls, per driver instance.

zetTracerSetPrologues is used to specify all the enter callbacks
zetTracerSetEpilogues is used to specify all the exit callbacks
If the value of a callback is nullptr, then it will be ignored.

The callbacks are defined as a collection of per-API function pointers, with 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 : the user’s pointer for the tracer’s data
ppTracerInstanceUserData : a per-tracer, per-instance thread-local storage location; typically used for passing data from the prologue to the epilogue

Note: since the creation of a tracer requires a device, on first glance it appears that zeInit, zeDriverGet and zeDeviceGet are not traceable. However, these APIs are traceable for all calls subsequent from the creation and enabling of the tracer itself.

Enabling/Disabling and Destruction¶

The tracer is created in a disabled state and must be explicitly enabled by calling zetTracerSetEnabled. The implementation guarantees that prologues and epilogues will always be executed in pairs; i.e.

if the prologue was called then the epilogue is guaranteed to be called, even if another thread disabled the tracer between execution
if the prologue was not called then the epilogue is guaranteed not to be called, even if another thread enabled the tracer between execution

The tracer should be disabled by the application before the tracer is destroyed. If multiple threads are in-flight, then it is still possible that callbacks will continue to execute even after the tracer is disabled; specifically, due to the pairing rules above. Due to the complexity involved in ensuring no threads are still or will be executing a callback even after its been disabled, the implementation will stall and wait for any outstanding threads during zetTracerDestroy.

The following pseudo-code demonstrates a basic usage of API tracing:

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_function_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_function_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);
}

void TracingExample( ... )
{
    my_tracer_data_t tracer_data = {};
    zet_tracer_desc_t tracer_desc;
    tracer_desc.stype = ZET_STRUCTURE_TYPE_TRACER_DESC;
    tracer_desc.pUserData = &tracer_data;
    zet_tracer_handle_t hTracer;
    zetTracerCreate(hDevice, &tracer_desc, &hTracer);

    // Set all callbacks
    zet_core_callbacks_t prologCbs = {};
    zet_core_callbacks_t epilogCbs = {};
    prologCbs.CommandList.pfnAppendLaunchFunction = OnEnterCommandListAppendLaunchKernel;
    epilogCbs.CommandList.pfnAppendLaunchFunction = OnExitCommandListAppendLaunchKernel;

    zetTracerSetPrologues(hTracer, &prologCbs);
    zetTracerSetEpilogues(hTracer, &epilogCbs);

    zetTracerSetEnabled(hTracer, true);

    zeCommandListAppendLaunchKernel(hCommandList, hFunction, &launchArgs, nullptr, 0, nullptr);
    zeCommandListAppendLaunchKernel(hCommandList, hFunction, &launchArgs, nullptr, 0, nullptr);
    zeCommandListAppendLaunchKernel(hCommandList, hFunction, &launchArgs, nullptr, 0, nullptr);

    zetTracerSetEnabled(hTracer, false);
    zetTracerDestroy(hTracer);
}

Metrics¶

Introduction¶

Devices provide programmable infrastructure designed to support performance debugging. The API described in this document provides access to these device metrics. The intention of this API is to support performance debug and it is not advised to use it in regular execution as overall performance can be negatively impacted.

The following diagram illustrates the relationship between the metric objects described in this document.

Most of the detailed metrics require the device to be properly programmed before use. It is important to understand that the device programming is in most cases global. This generally means that if a software tool or an application is using the metrics then no other application can reliably use the same device resources.

Metric Groups¶

The device infrastructure consists of a pre-defined set of non-programmable counters, and a programmable network of connections that work with a separate set of counters as well as other types of counters. For sake of simplicity, the smallest unit of configuration is a Metric Group. Metric Groups are sets of metrics that provide certain perspective on workload’s performance. The groups aggregate metrics, define device programming and available collection methods. An application may choose to collect data from several Metric Groups provided that they all belong to different domains. Domains are used as a software representation of independent device resources that can safely be used concurrently.

Sampling Types¶

Sampling types are a software representation of device capabilities in terms of reading metric values. Each Metric Group provides information on which sampling types it supports. There are separate sets of APIs supporting each of the sampling types Time-based and Event-based.

All available sampling types are defined in zet_metric_group_sampling_type_t.

Information about supported sampling types for a given Metric Group is provided in zet_metric_group_properties_t.samplingType.
It’s possible that a device provides multiple Metric Groups with the same names but different sampling types.
When enumerating, it’s important to choose a Metric Group which supports the desired sampling type.

Domains¶

Every Metric Group belongs to a given domain (zet_metric_group_properties_t.domain).

The Metric Group typically define a uniform device counter configuration used for measurements.
Each domain represents an exclusive resource used by the Metric Group.
It’s possible to simultaneously gather data for two different Metric Groups, only if they belong to a different domain i.e. Metric Groups that can be collected concurrently will have different domain values.

Enumeration¶

All available metrics are organized into Metric Groups.

During data collection, data for the whole Metric Group is gathered.
The list of available Metric Groups and their Metrics is device-specific.

The following APIs provide all the information needed for identification and usage.

Metric Group properties are accessed through the function zetMetricGroupGetProperties which returns zet_metric_group_properties_t.
Metric properties are accessed through the function zetMetricGetProperties which returns zet_metric_properties_t.

A common tool flow is to enumerate metrics looking for a specific Metric Group. Depending on the metrics required for a specific scenario a tool may choose to run the workload multiple times, recording different set of Metric Groups each time. Usually care must be taken to ensure run-to-run stability and result repeatability if metrics from different runs are meant to be used together. When enumerating metrics to find a desired Metric Group, it’s important to know in advance which sampling type will be used.

To enumerate through the Metric tree:

Call zetMetricGroupGet to obtain Metric Group count.
Call zetMetricGroupGet to obtain all Metric Groups.
Iterate over all available Metric Groups.
- At this point it’s possible to check Metric Group name, domain or sampling type.
- Metric Group names may not be unique.
Obtain the metric count for each Metric Group by calling zetMetricGroupGetProperties with Metric Group handle (zet_metric_group_handle_t) and checking zet_metric_group_properties_t.metricCount.
Iterate over available Metrics using zetMetricGet with parent Metric Group (zet_metric_group_handle_t).
Check Metric properties (e.g. name, description) calling zetMetricGetProperties with parent Metric (zet_metric_handle_t).

The following pseudo-code demonstrates a basic enumeration over all available metric groups and their metrics. Additionally, it returns a metric group with a chosen name and sampling type. Similar code could be used for selecting a preferred metric group for a specific type of measurements.

ze_result_t FindMetricGroup( ze_device_handle_t hDevice,
                               char* pMetricGroupName,
                               uint32_t desiredSamplingType,
                               zet_metric_group_handle_t* phMetricGroup )
{
    // Obtain available metric groups for the specific device
    uint32_t metricGroupCount = 0;
    zetMetricGroupGet( hDevice, &metricGroupCount, nullptr );

    zet_metric_group_handle_t* phMetricGroups = malloc(metricGroupCount * sizeof(zet_metric_group_handle_t));
    zetMetricGroupGet( hDevice, &metricGroupCount, phMetricGroups );

    // Iterate over all metric groups available
    for( i = 0; i < metricGroupCount; i++ )
    {
        // Get metric group under index 'i' and its properties
        zet_metric_group_properties_t metricGroupProperties;
        zetMetricGroupGetProperties( phMetricGroups[i], &metricGroupProperties );

        printf("Metric Group: %sn", metricGroupProperties.name);

        // Check whether the obtained metric group supports the desired sampling type
        if((metricGroupProperties.samplingType & desiredSamplingType) == desiredSamplingType)
        {
            // Check whether the obtained metric group has the desired name
            if( strcmp( pMetricGroupName, metricGroupProperties.name ) == 0 )
            {
                *phMetricGroup = phMetricGroups[i];
                break;
            }
        }
    }

    free(phMetricGroups);
}

Configuration¶

Use the zetDeviceActivateMetricGroups API call to configure the device for data collection.

Subsequent calls to the function will disable device programming for the metric groups not selected for activation.
To avoid undefined results only call the zetDeviceActivateMetricGroups between experiments i.e. while not collecting data.

Programming restrictions:

Any combination of metric groups can be configured simultaneously provided that all of them have a different zet_metric_group_properties_t.domain.
MetricGroup must be active until zetMetricStreamerClose and the last zetCommandListAppendMetricQueryEnd completes.

Collection¶

There are two modes of metrics collection supported: time-based and event-based.

Time-based collection is using a timer as well as other events to store data samples. A metric streamer interface is the software interface for configuration and collection.
Event-based collection is based on a pair of Begin/End events appended to command lists. A metric query interface is the software interface for configuration and collection.

Metric Streamer¶

Time-based collection uses a simple Open, Read, Close scheme:

zetMetricStreamerOpen opens the streamer.
zetMetricStreamerReadData reads the raw data to be later processed by zetMetricGroupCalculateMetricValues.
zetMetricStreamerClose closes the streamer.

The following pseudo-code demonstrates a basic sequence for time-based collection:

ze_result_t TimeBasedUsageExample( ze_driver_handle_t hDriver,
                                     ze_device_handle_t hDevice )
{
    zet_metric_group_handle_t     hMetricGroup           = nullptr;
    ze_event_handle_t            hNotificationEvent     = nullptr;
    ze_event_pool_handle_t       hEventPool             = nullptr;
    ze_event_pool_desc_t         eventPoolDesc          = {ZE_STRUCTURE_TYPE_EVENT_POOL_DESC, nullptr, ZE_EVENT_POOL_FLAG_DEFAULT , 1};
    ze_event_desc_t              eventDesc              = {ZE_STRUCTURE_TYPE_EVENT_DESC};
    zet_metric_streamer_handle_t  hMetricStreamer        = nullptr;
    zet_metric_streamer_desc_t    metricStreamerDesc     = {ZET_STRUCTURE_TYPE_METRIC_STREAMER_DESC};

    // Find a "ComputeBasic" metric group suitable for Time Based collection
    FindMetricGroup( hDevice, "ComputeBasic", ZET_METRIC_GROUP_SAMPLING_TYPE_TIME_BASED, &hMetricGroup );

    // Configure the HW
    zetDeviceActivateMetricGroups( hDevice, /* count= / 1, &hMetricGroup );

    // Create notification event
    :ref:`zeEventPoolCreate`( hDriver, &eventPoolDesc, 1, &hDevice, &hEventPool );
    eventDesc.index  = 0;
    eventDesc.signal = :ref:`ZE_EVENT_SCOPE_FLAG_HOST <ze-event-scope-flag-t>`;
    eventDesc.wait   = :ref:`ZE_EVENT_SCOPE_FLAG_HOST <ze-event-scope-flag-t>`;
    :ref:`zeEventCreate`( hEventPool, &eventDesc, &hNotificationEvent );

    // Open metric streamer
    metricStreamerDesc.samplingPeriod       = 1000;
    metricStreamerDesc.notifyEveryNReports  = 32768;
    :ref:`zetMetricStreamerOpen`( hDevice, hMetricGroup, &metricStreamerDesc, hNotificationEvent, &hMetricStreamer );

    // Run your workload, in this example we assume the data for the whole experiment fits in the device buffer
    Workload(hDevice);
    // Optionally insert markers during workload execution
    //:ref:`zetCommandListAppendMetricStreamerMarker`( hCommandList, hMetricStreamer, tool_marker_value );

    // Wait for data, optional in this example since the whole workload has already been executed by now
    //:ref:`zeEventHostSynchronize`( hNotificationEvent, 1000 /*timeout/ );
    // reset the event if it fired

    // Read raw data
    size_t rawSize = 0;
    zetMetricStreamerReadData( hMetricStreamer, UINT32_MAX, &rawSize, nullptr );
    uint8_t* rawData = malloc(rawSize);
    zetMetricStreamerReadData( hMetricStreamer, UINT32_MAX, &rawSize, rawData );

    // Close metric streamer
    zetMetricStreamerClose( hMetricStreamer );
    zeEventDestroy( hNotificationEvent );
    zeEventPoolDestroy( hEventPool );

    // Deconfigure the device
    zetDeviceActivateMetricGroups( hDevice, 0, nullptr );

    // Calculate metric data
    CalculateMetricsExample( hMetricGroup, rawSize, rawData );
    free(rawData);
}

Metric Query¶

Event-based collection uses a simple Begin, End, GetData scheme:

zetCommandListAppendMetricQueryBegin defines the start counting event
zetCommandListAppendMetricQueryEnd defines the finish counting event
zetMetricQueryGetData reads the raw data to be later processed by zetMetricGroupCalculateMetricValues.

Typically, multiple queries are used and recycled to characterize a workload. A Query Pool is used to efficiently use and reuse device memory for multiple queries.

zetMetricQueryPoolCreate creates a pool of homogeneous queries.
zetMetricQueryPoolDestroy frees the pool. The application must ensure no queries within the pool are in-use before freeing the pool.
zetMetricQueryCreate obtains a handle to a unique location in the pool.
zetMetricQueryReset allows for low-cost recycling of a location in the pool.

The following pseudo-code demonstrates a basic sequence for query-based collection:

ze_result_t MetricQueryUsageExample( ze_driver_handle_t hDriver,
                                       ze_device_handle_t hDevice )
{
    zet_metric_group_handle_t      hMetricGroup          = nullptr;
    ze_event_handle_t             hCompletionEvent      = nullptr;
    ze_event_pool_desc_t          eventPoolDesc         = {ZE_STRUCTURE_TYPE_EVENT_POOL_DESC};
    ze_event_desc_t               eventDesc             = {ZE_STRUCTURE_TYPE_EVENT_DESC};
    ze_event_pool_handle_t        hEventPool            = nullptr;
    zet_metric_query_pool_handle_t hMetricQueryPool      = nullptr;
    zet_metric_query_handle_t      hMetricQuery          = nullptr;
    zet_metric_query_pool_desc_t   queryPoolDesc         = {ZET_STRUCTURE_TYPE_METRIC_QUERY_POOL_DESC};

    // Find a "ComputeBasic" metric group suitable for Event Based collection
    FindMetricGroup( hDevice, "ComputeBasic", ZET_METRIC_GROUP_SAMPLING_TYPE_EVENT_BASED, &hMetricGroup );

    // Configure HW
    zetDeviceActivateMetricGroups( hDevice, 1 /* count /, &hMetricGroup );

    // Create metric query pool & completion event
    queryPoolDesc.flags        = :ref:`ZET_METRIC_QUERY_POOL_FLAG_PERFORMANCE <zet-metric-query-pool-flag-t>`;
    queryPoolDesc.count        = 1000;
    :ref:`zetMetricQueryPoolCreate`( hDevice, hMetricGroup, &queryPoolDesc, &hMetricQueryPool );
    eventPoolDesc.flags = :ref:`ZE_EVENT_POOL_FLAG_DEFAULT <ze-event-pool-flag-t>`;
    eventPoolDesc.count = 1000;
    :ref:`zeEventPoolCreate`( hDriver, &eventPoolDesc, 1, &hDevice, &hEventPool );

    // Write BEGIN metric query to command list
    :ref:`zetMetricQueryCreate`( hMetricQueryPool, 0 /*slot/, &hMetricQuery );
    zetCommandListAppendMetricQueryBegin( hCommandList, hMetricQuery );

    // build your command list
    ...

    // Write END metric query to command list, use an event to determine if the data is available
    eventDesc.index  = 0;
    eventDesc.signal = ZE_EVENT_SCOPE_FLAG_HOST;
    eventDesc.wait   = ZE_EVENT_SCOPE_FLAG_HOST;
    zeEventCreate( hEventPool, &eventDesc, &hCompletionEvent);
    zetCommandListAppendMetricQueryEnd( hCommandList, hMetricQuery, hCompletionEvent, 0, nullptr );

    // use zeCommandQueueExecuteCommandLists( , , , ) to submit your workload to the device

    // Wait for data
    zeEventHostSynchronize( hCompletionEvent, 1000 /timeout/ );

    // Read raw data
    size_t rawSize = 0;
    zetMetricQueryGetData( hMetricQuery, &rawSize, nullptr );
    uint8_t* rawData = malloc(rawSize);
    zetMetricQueryGetData( hMetricQuery, &rawSize, rawData );

    // Free the resources
    zeEventDestroy( hCompletionEvent );
    zeEventPoolDestroy( hEventPool );
    zetMetricQueryPoolDestroy( hMetricQueryPool );

    // Deconfigure HW
    zetDeviceActivateMetricGroups( hDevice, 0, nullptr );

    // Calculate metric data
    CalculateMetricsExample( hMetricGroup, rawSize, rawData );
    free(rawData);
}

Calculation¶

Both MetricStreamer and MetricQuery collect the data in device specific, raw form that is not suitable for application processing. To calculate metric values use zetMetricGroupCalculateMetricValues.

The following pseudo-code demonstrates a basic sequence for metric calculation and interpretation:

ze_result_t CalculateMetricsExample( zet_metric_group_handle_t hMetricGroup,
                                       size_t rawSize, uint8_t* rawData )
{
    // Calculate metric data
    uint32_t numMetricValues = 0;
    zetMetricGroupCalculateMetricValues( hMetricGroup, rawSize, rawData, &numMetricValues, nullptr );
    zet_typed_value_t* metricValues = malloc( numMetricValues * sizeof(zet_typed_value_t) );
    zetMetricGroupCalculateMetricValues( hMetricGroup, rawSize, rawData, &numMetricValues, metricValues );

    // Obtain available metrics for the specific metric group
    uint32_t metricCount = 0;
    zetMetricGet( hMetricGroup, &metricCount, nullptr );

    zet_metric_handle_t* phMetrics = malloc(metricCount * sizeof(zet_metric_handle_t));
    zetMetricGet( hMetricGroup, &metricCount, phMetrics );

    // Print metric results
    uint32_t numReports = numMetricValues / metricCount;
    for( uint32_t report = 0; report < numReports; ++report )
    {
        printf("Report: %dn", report);

        for( uint32_t metric = 0; metric < metricCount; ++metric )
        {
            zet_typed_value_t data = metricValues[report * metricCount + metric];

            zet_metric_properties_t metricProperties;
            zetMetricGetProperties( phMetrics[ metric ], &metricProperties );

            printf("Metric: %sn", metricProperties.name );

            switch( data.type )
            {
            case ZET_VALUE_TYPE_UINT32:
                printf(" Value: %lun", data.value.ui32 );
                break;
            case ZET_VALUE_TYPE_UINT64:
                printf(" Value: %llun", data.value.ui64 );
                break;
            case ZET_VALUE_TYPE_FLOAT32:
                printf(" Value: %fn", data.value.fp32 );
                break;
            case ZET_VALUE_TYPE_FLOAT64:
                printf(" Value: %fn", data.value.fp64 );
                break;
            case ZET_VALUE_TYPE_BOOL8:
                if( data.value.ui32 )
                    printf(" Value: truen" );
                else
                    printf(" Value: falsen" );
                break;
            default:
                break;
            };
        }
    }

    free(metricValues);
    free(phMetrics);
}

Program Instrumentation¶

Introduction¶

The program instrumentation APIs provide tools a basic framework for low-level profiling of device kernels, by allowing direct instrumentation of those programs. These capabilities, in combination with those already provided, and in combination with API tracing, are sufficient for more advanced frameworks to be developed.

There are two types of instrumentation available:

Inter-Function Instrumentation - intercepting and redirecting function calls
Intra-Function Instrumentation - injecting new instructions within a function

Inter-Function Instrumentation¶

The following capabilities allow for a tool to intercept and redirect function calls:

Inter-module function calls - the ability to call functions between different modules; e.g., the application’s module and a tool’s module
API-Tracing

For example, a tool may use API Tracing in any of the following ways:

zeModuleCreate - replace a module handle with instrumented module handle for all functions
zeKernelCreate - replace a kernel handle with instrumented kernel handle for all call sites
zeModuleGetFunctionPointer - replace a function pointer with instrumented function pointer for all call sites
zeCommandListAppendLaunchKernel - replace a kernel handle with instrumented kernel handle at call site

Intra-Function Instrumentation¶

The following capabilities allow for a tool to inject instructions within a kernel:

zetModuleGetDebugInfo - allows a tool to query standard debug info for an application’s module
zetKernelGetProfileInfo - allows a tool to query detailed information on aspects of a kernel
zeModuleGetNativeBinary - allows for a tool to retrieve the native binary of the application’s module, instrument it, then create a new module using the instrumented version
API-Tracing - same usage as Inter-Function Instrumentation above

Compilation¶

A module must be compiled with foreknowledge that instrumentation will be performed for the compiler to generate the proper profiling meta-data. Therefore, when the instrumentation layer is enabled, a new build flag is supported: “-zet-profile-flags <value>”, where <value> must be a combination of zet_profile_flag_t, in hexidecimal.

As an example, a tool could use API Tracing to inject this build flag on each zeModuleCreate call that the tool wishes to instrument. In another example, a tool could recompile a Module using the build flag and use API Tracing to replace the application’s Module handle with it’s own.

Instrumentation¶

Once the module has been compiled with instrumentation enabled, a tool may use zetModuleGetDebugInfo and zetKernelGetProfileInfo in order to decode the application’s instructions and register usage for each function in the module.

If a tool requires additional functions to be used, it may create other module(s) and use zeModuleGetFunctionPointer to call functions between the application and tool modules. A tool may use zeModuleGetFunctionPointer to retrieve the Host and device address of each function in the module.

There are no APIs provided for the actual instrumentation. Instead this is left up to the tool itself to decode the application module’s native binary and inject native instructions. This model prevents the instrumentation from being manipulated by the compiler.

Execution¶

If a tool requires changing the address of an application’s function, then it should use API Tracing. For example, zeModuleGetFunctionPointer and all flavors of zeCommandListAppendLaunchKernel.