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.

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.
../_images/tools_metric_hierarchy.png

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.

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:

  1. Call zetMetricGroupGet to obtain Metric Group count.

  2. Call zetMetricGroupGet to obtain all Metric Groups.

  3. 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.

  4. 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.

  5. Iterate over available Metrics using zetMetricGet with parent Metric Group (zet_metric_group_handle_t).

  6. 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:

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:

../_images/tools_metric_streamer.png

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:

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.

../_images/tools_metric_query.png

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:

  1. Inter-Function Instrumentation - intercepting and redirecting function calls

  2. 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:

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.