Callback tracing

Callback tracing#

enum rocprofiler_code_object_storage_type_t#

ROCProfiler Enumeration for code object storage types (identical values to hsa_ven_amd_loader_code_object_storage_type_t enumeration)

Values:

enumerator ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_NONE#

enumerator ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE#

enumerator ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY#

enumerator ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_LAST#

typedef void (*rocprofiler_callback_tracing_cb_t)(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t *user_data, void *callback_data)#

API Tracing callback function. This function is invoked twice per API function: once before the function is invoked and once after the function is invoked. The external correlation id value within the record is assigned the value at the top of the external correlation id stack. It is permissible to invoke rocprofiler_push_external_correlation_id within the enter phase; when a new external correlation id is pushed during the enter phase, rocprofiler will use that external correlation id for any async events and provide the new external correlation id during the exit callback… In other words, pushing a new external correlation id within the enter callback will result in that external correlation id value in the exit callback (which may or may not be different from the external correlation id value in the enter callback). If a tool pushes new external correlation ids in the enter phase, it is recommended to pop the external correlation id in the exit callback.

Param record:: [in] Callback record data
Param user_data:: [inout] This parameter can be used to retain information in between the enter and exit phases.
Param callback_data:: [in] User data provided when configuring the callback tracing service

typedef int (*rocprofiler_callback_tracing_kind_cb_t)(rocprofiler_callback_tracing_kind_t kind, void *data)#

Callback function for mapping rocprofiler_callback_tracing_kind_t ids to string names.

See also

rocprofiler_iterate_callback_tracing_kind_names.

typedef int (*rocprofiler_callback_tracing_kind_operation_cb_t)(rocprofiler_callback_tracing_kind_t kind, rocprofiler_tracing_operation_t operation, void *data)#

Callback function for mapping the operations of a given rocprofiler_callback_tracing_kind_t to string names.

See also

rocprofiler_iterate_callback_tracing_kind_operation_names.

typedef int (*rocprofiler_callback_tracing_operation_args_cb_t)(rocprofiler_callback_tracing_kind_t kind, rocprofiler_tracing_operation_t operation, uint32_t arg_number, const void *const arg_value_addr, int32_t arg_indirection_count, const char *arg_type, const char *arg_name, const char *arg_value_str, int32_t arg_dereference_count, void *data)#

Callback function for iterating over the function arguments to a traced function. This function will be invoked for each argument.

See also

rocprofiler_iterate_callback_tracing_operation_args

Param kind:: [in] domain
Param operation:: [in] associated domain operation
Param arg_number:: [in] the argument number, starting at zero
Param arg_value_addr:: [in] the address of the argument stored by rocprofiler.
Param arg_indirection_count:: [in] the total number of indirection levels for the argument, e.g. int == 0, int* == 1, int** == 2
Param arg_type:: [in] the typeid name of the argument
Param arg_name:: [in] the name of the argument in the prototype (or rocprofiler union)
Param arg_value_str:: [in] conversion of the argument to a string, e.g. operator<< overload
Param arg_dereference_count:: [in] the number of times the argument was dereferenced when it was converted to a string
Param data:: [in] user data

typedef uint64_t (*rocprofiler_kernel_replay_pass_count_cb_t)(rocprofiler_kernel_dispatch_info_t dispatch_info, rocprofiler_user_data_t user_data)#: Tool-provided callback returning the number of replay passes for a dispatch.

typedef int (*rocprofiler_kernel_replay_continue_cb_t)(rocprofiler_kernel_dispatch_info_t dispatch_info, uint64_t current_pass, uint64_t total_passes, rocprofiler_user_data_t user_data)#: Tool-provided callback invoked after each pass to decide whether the loop continues.

typedef rocprofiler_status_t (*rocprofiler_kernel_replay_context_cb_t)(rocprofiler_context_id_t context_id)#: SDK-provided callback marking a context enabled or disabled for the current replay loop.

rocprofiler_status_t rocprofiler_configure_callback_tracing_service(rocprofiler_context_id_t context_id, rocprofiler_callback_tracing_kind_t kind, const rocprofiler_tracing_operation_t *operations, unsigned long operations_count, rocprofiler_callback_tracing_cb_t callback, void *callback_args)#

Configure Callback Tracing Service. The callback tracing service provides two synchronous callbacks around an API function on the same thread as the application which is invoking the API function. This function can only be invoked once per rocprofiler_callback_tracing_kind_t value, i.e. it can be invoked once for the HSA API, once for the HIP API, and so on but it will fail if it is invoked for the HSA API twice. Please note, the callback API does have the potentially non-trivial overhead of copying the function arguments into the record. If you are willing to let rocprofiler record the timestamps, do not require synchronous notifications of the API calls, and want to lowest possible overhead, use the.

See also

Asynchronous Tracing Service.

Parameters:

context_id – [in] Context to associate the service with
kind – [in] The domain of the callback tracing service
operations – [in] Array of operations in the domain (i.e. enum values which identify specific API functions). If this is null, all API functions in the domain will be traced
operations_count – [in] If the operations array is non-null, set this to the size of the array.
callback – [in] The function to invoke before and after an API function
callback_args – [in] Data provided to every invocation of the callback function

Return values:

ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED – Invoked outside of the initialization function in rocprofiler_tool_configure_result_t provided to rocprofiler via rocprofiler_configure function
ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND – The provided context is not valid/registered
ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED – if the same rocprofiler_callback_tracing_kind_t value is provided more than once (per context) — in other words, we do not support overriding or combining the operations in separate function calls.

Returns:

rocprofiler_status_t

rocprofiler_status_t rocprofiler_query_callback_tracing_kind_name(rocprofiler_callback_tracing_kind_t kind, const char **name, uint64_t *name_len)#

Query the name of the callback tracing kind. The name retrieved from this function is a string literal that is encoded in the read-only section of the binary (i.e. it is always “allocated” and never “deallocated”).

Parameters:

kind – [in] Callback tracing domain
name – [out] If non-null and the name is a constant string that does not require dynamic allocation, this parameter will be set to the address of the string literal, otherwise it will be set to nullptr
name_len – [out] If non-null, this will be assigned the length of the name (regardless of the name is a constant string or requires dynamic allocation)

Returns:

rocprofiler_status_t

rocprofiler_status_t rocprofiler_query_callback_tracing_kind_operation_name(rocprofiler_callback_tracing_kind_t kind, rocprofiler_tracing_operation_t operation, const char **name, uint64_t *name_len)#

Query the name of the callback tracing kind. The name retrieved from this function is a string literal that is encoded in the read-only section of the binary (i.e. it is always “allocated” and never “deallocated”).

Parameters:

kind – [in] Callback tracing domain
operation – [in] Enumeration id value which maps to a specific API function or event type
name – [out] If non-null and the name is a constant string that does not require dynamic allocation, this parameter will be set to the address of the string literal, otherwise it will be set to nullptr
name_len – [out] If non-null, this will be assigned the length of the name (regardless of the name is a constant string or requires dynamic allocation)

Return values:

ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND – Domain id is not valid
ROCPROFILER_STATUS_SUCCESS – Valid domain provided, regardless if there is a constant string or not.

Returns:

rocprofiler_status_t

rocprofiler_status_t rocprofiler_iterate_callback_tracing_kinds(rocprofiler_callback_tracing_kind_cb_t callback, void *data)#

Iterate over all the mappings of the callback tracing kinds and get a callback for each kind.

Parameters:

callback – [in] Callback function invoked for each enumeration value in rocprofiler_callback_tracing_kind_t with the exception of the NONE and LAST values.
data – [in] User data passed back into the callback

Returns:

rocprofiler_status_t

rocprofiler_status_t rocprofiler_iterate_callback_tracing_kind_operations(rocprofiler_callback_tracing_kind_t kind, rocprofiler_callback_tracing_kind_operation_cb_t callback, void *data)#

Iterates over all the mappings of the operations for a given rocprofiler_callback_tracing_kind_t and invokes the callback with the kind id, operation id, and user-provided data.

Parameters:

kind – [in] which tracing callback kind operations to iterate over
callback – [in] Callback function invoked for each operation associated with rocprofiler_callback_tracing_kind_t with the exception of the NONE and LAST values.
data – [in] User data passed back into the callback

Return values:

ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND – Invalid domain id
ROCPROFILER_STATUS_SUCCESS – Valid domain

Returns:

rocprofiler_status_t

rocprofiler_status_t rocprofiler_iterate_callback_tracing_kind_operation_args(rocprofiler_callback_tracing_record_t record, rocprofiler_callback_tracing_operation_args_cb_t callback, int32_t max_dereference_count, void *user_data)#

ROCPROFILER_CODE_OBJECT_ID_NONE#: The NULL value of a code object id. Used when code object is unknown.

struct rocprofiler_callback_tracing_hsa_api_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler HSA API Callback Data.

struct rocprofiler_callback_tracing_hip_api_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler HIP runtime and compiler API Tracer Callback Data.

struct rocprofiler_callback_tracing_ompt_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler OMPT Callback Data.

struct rocprofiler_callback_tracing_marker_api_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Marker Tracer Callback Data.

struct rocprofiler_callback_tracing_rccl_api_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler RCCL API Callback Data.

struct rocprofiler_callback_tracing_rocdecode_api_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler rocDecode API Callback Data.

struct rocprofiler_callback_tracing_rocjpeg_api_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler rocJPEG API Callback Data.

struct rocprofiler_callback_tracing_rocshmem_api_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler rocSHMEM API Callback Data.

struct rocprofiler_callback_tracing_hipfile_api_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler hipFILE API Callback Data.

struct rocprofiler_callback_tracing_code_object_load_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Code Object Load Tracer Callback Record.

struct rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Code Object Kernel Symbol Tracer Callback Record.

struct rocprofiler_callback_tracing_code_object_host_kernel_symbol_register_data_t#

struct rocprofiler_callback_tracing_kernel_dispatch_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Kernel Dispatch Callback Tracer Record.

struct rocprofiler_callback_tracing_hip_event_data_t#

#include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler HIP Event Callback Tracer Record.

Timestamps are only non-zero in the ROCPROFILER_CALLBACK_PHASE_NONE callback (barrier completion). For WAIT operations, source_queue_id identifies where the event was originally recorded; for RECORD operations it equals queue_id.

struct rocprofiler_callback_tracing_memory_copy_data_t#

#include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Memory Copy Callback Tracer Record.

The timestamps in this record will only be non-zero in the ROCPROFILER_CALLBACK_PHASE_EXIT callback

struct rocprofiler_callback_tracing_memory_allocation_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Memory Allocation Tracer Record.

struct rocprofiler_callback_tracing_scratch_memory_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Scratch Memory Callback Data.

struct rocprofiler_callback_tracing_runtime_initialization_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Runtime Initialization Data.

struct rocprofiler_callback_tracing_hip_stream_data_t#: #include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler Stream Handle Callback Data.

struct rocprofiler_callback_tracing_hip_graph_data_t#

#include <rocprofiler-sdk/callback_tracing.h>

ROCProfiler HIP Graph Callback Data.

Delivered with callbacks fired for hipGraphInstantiate*, hipGraphExecDestroy, and hipGraphLaunch{,_spt} via the ROCPROFILER_CALLBACK_TRACING_HIP_GRAPH domain.

Sub-operations (rocprofiler_hip_graph_operation_t):

EXEC_CREATE fires after a successful hipGraphInstantiate*
EXEC_DESTROY fires after a successful hipGraphExecDestroy
EXEC_LAUNCH fires at ENTER and EXIT phases of hipGraphLaunch / hipGraphLaunch_spt

Tools that want to associate kernel dispatches with their producing graph node should:

Subscribe to ROCPROFILER_CALLBACK_TRACING_HIP_GRAPH and maintain a per-thread stack of (graph_exec_id, node_counter) by handling EXEC_LAUNCH ENTER (push) and EXIT (pop).
Register an external correlation id request callback that, when invoked for KERNEL_DISPATCH / MEMORY_COPY, captures the current top-of-stack graph attribution into the external correlation id and increments the node_counter.
At record consumption time, extract the captured attribution from rocprofiler_correlation_id_t::external on the dispatch/copy record.

This mirrors the ROCPROFILER_CALLBACK_TRACING_HIP_STREAM design.

struct rocprofiler_callback_tracing_kernel_replay_data_t#

#include <rocprofiler-sdk/experimental/kernel_replay.h>

ROCProfiler Kernel Replay Callback Tracer Record.

Payload for ROCPROFILER_CALLBACK_TRACING_KERNEL_REPLAY callbacks. All members are present in the struct (no unions). Which members are meaningful depends on the current operation:

ROCPROFILER_KERNEL_REPLAY_CONFIG: dispatch_info is populated by the SDK. The tool sets replay_pass_count and optionally replay_continue during ROCPROFILER_CALLBACK_PHASE_ENTER. Pass-info fields are zero.
ROCPROFILER_KERNEL_REPLAY_PASS: dispatch_info, current_pass, and total_passes are populated by the SDK. replay_pass_count reads as NULL and must not be modified (the pass count is fixed at CONFIG PHASE_ENTER). replay_continue is NULL during PASS PHASE_ENTER; during PASS PHASE_EXIT the SDK populates it with the callback in effect for this pass, and the tool may overwrite it there to change the continue decision for this pass only (the next pass reverts to the CONFIG default).

The SDK maintains a single rocprofiler_user_data_t for the entire replay sequence (CONFIG + all PASS operations). A tool can write per-dispatch state into user_data during CONFIG PHASE_ENTER; the same value is delivered to every subsequent PASS callback and to replay_pass_count and replay_continue for the same dispatch.

Each PASS receives its own copy of that value, so a write during a PASS callback is scoped to the pass that performed it: that pass’s PASS PHASE_EXIT and its replay_continue decision observe the write, while the next PASS and CONFIG PHASE_EXIT still observe the value set during CONFIG PHASE_ENTER. replay_continue runs after PASS PHASE_EXIT and receives that pass’s copy (the value its PHASE_EXIT left), not the CONFIG value. For state that must persist across the whole sequence regardless of per-pass writes, place it behind user_data.ptr: mutating the pointed-to object is visible from every callback, whereas overwriting the union value is per-pass. See samples/kernel_replay/basic_client_with_user_data.cpp for a worked example.

Exactly one context may configure a KERNEL_REPLAY service; a second rocprofiler_configure_callback_tracing_service for this domain returns ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED.

See also

docs/how-to/using-kernel-replay.rst for the full limitation list and docs/conceptual/kernel_replay/kernel_replay_memory_snapshot.md for what the snapshot covers and why.

Warning

Beta. A dispatch is replayed only when its submission is a single packet containing a single dispatch. HIP graph launches are not replayed, and a multi-packet submission runs once without replay; each case warns once. Repeatability rests on a snapshot covering coarse-grained device allocations owned by the agent plus module-scope __device__ / __constant__ variables. Unified or managed memory, hipMallocAsync and other virtual-memory-mapped allocations, and host, fine-grained and kernarg memory are not captured, so a kernel writing to them observes values accumulated across passes rather than identical inputs. Allocations carrying HSA_AMD_MEMORY_POOL_EXECUTABLE_FLAG are excluded from the snapshot (HIP kernarg pools / profiler buffers share that flag). A direct-HSA application that puts ordinary writable device data behind the same flag sees the same omission — an unsupported allocation class for beta. Declining every replay while any such trackable allocation is live is not viable because the HIP runtime itself keeps them live under interception. A failed restore aborts the process rather than continuing with partially restored memory.

rocprofiler_callback_tracing_code_object_load_data_t.__unnamed4__

Public Members

rocprofiler_agent_id_t rocp_agent#: Deprecated. Renamed to agent_id.

rocprofiler_agent_id_t agent_id#: The agent on which this loaded code object is loaded.

rocprofiler_callback_tracing_code_object_load_data_t.__unnamed6__

Public Members

struct rocprofiler_callback_tracing_code_object_load_data_t

struct rocprofiler_callback_tracing_code_object_load_data_t

rocprofiler_callback_tracing_code_object_load_data_t.__unnamed6__.__unnamed8__

rocprofiler_callback_tracing_code_object_load_data_t.__unnamed6__.__unnamed10__

Callback tracing

Contents

Callback tracing#