/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprofiler-docs/checkouts/docs-5.6.0/inc/rocprofiler.h Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprofiler-docs/checkouts/docs-5.6.0/inc/rocprofiler.h Source File#

ROCProfilerV1 API: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprofiler-docs/checkouts/docs-5.6.0/inc/rocprofiler.h Source File
inc
Go to the documentation of this file.
 /******************************************************************************
 Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
  
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
 in the Software without restriction, including without limitation the rights
 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
 copies of the Software, and to permit persons to whom the Software is
 furnished to do so, subject to the following conditions:
  
 The above copyright notice and this permission notice shall be included in
 all copies or substantial portions of the Software.
  
 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 *******************************************************************************/
  
 //
 // ROC Profiler API
 //
 // The goal of the implementation is to provide a HW specific low-level
 // performance analysis interface for profiling of GPU compute applications.
 // The profiling includes HW performance counters (PMC) with complex
 // performance metrics and traces.
 //
 // The library can be used by a tool library loaded by HSA runtime or by
 // higher level HW independent performance analysis API like PAPI.
 //
 // The library is written on C and will be based on AQLprofile AMD specific
 // HSA extension. The library implementation requires HSA API intercepting and
 // a profiling queue supporting a submit callback interface.
 //
 //
 //
  
 #ifndef INC_ROCPROFILER_H_
 #define INC_ROCPROFILER_H_
  
 /* Placeholder for calling convention and import/export macros */
 #if !defined(ROCPROFILER_CALL)
 #define ROCPROFILER_CALL
 #endif /* !defined (ROCPROFILER_CALL) */
  
 #if !defined(ROCPROFILER_EXPORT_DECORATOR)
 #if defined(__GNUC__)
 #define ROCPROFILER_EXPORT_DECORATOR __attribute__((visibility("default")))
 #elif defined(_MSC_VER)
 #define ROCPROFILER_EXPORT_DECORATOR __declspec(dllexport)
 #endif /* defined (_MSC_VER) */
 #endif /* !defined (ROCPROFILER_EXPORT_DECORATOR) */
  
 #if !defined(ROCPROFILER_IMPORT_DECORATOR)
 #if defined(__GNUC__)
 #define ROCPROFILER_IMPORT_DECORATOR
 #elif defined(_MSC_VER)
 #define ROCPROFILER_IMPORT_DECORATOR __declspec(dllimport)
 #endif /* defined (_MSC_VER) */
 #endif /* !defined (ROCPROFILER_IMPORT_DECORATOR) */
  
 #define ROCPROFILER_EXPORT ROCPROFILER_EXPORT_DECORATOR ROCPROFILER_CALL
 #define ROCPROFILER_IMPORT ROCPROFILER_IMPORT_DECORATOR ROCPROFILER_CALL
  
 #if !defined(ROCPROFILER)
 #if defined(ROCPROFILER_EXPORTS)
 #define ROCPROFILER_API ROCPROFILER_EXPORT
 #else /* !defined (ROCPROFILER_EXPORTS) */
 #define ROCPROFILER_API ROCPROFILER_IMPORT
 #endif /* !defined (ROCPROFILER_EXPORTS) */
 #endif /* !defined (ROCPROFILER) */
  
 #include <stddef.h>
 #include <stdint.h>
  
 #ifdef __cplusplus
 extern "C" {
 #endif /* __cplusplus */
  
 #define ROCPROFILER_VERSION_9_0
  
 #define ROCPROFILER_VERSION_MAJOR 9
  
 #define ROCPROFILER_VERSION_MINOR 0
  
 ROCPROFILER_API uint32_t rocprofiler_version_major();
  
 ROCPROFILER_API uint32_t rocprofiler_version_minor();
  
 // TODO(aelwazir): Fix them to use the new Error codes
 typedef enum {
   ROCPROFILER_STATUS_SUCCESS = 0,
   ROCPROFILER_STATUS_ERROR = -1,
   ROCPROFILER_STATUS_ERROR_ALREADY_INITIALIZED = -2,
   ROCPROFILER_STATUS_ERROR_NOT_INITIALIZED = -3,
   ROCPROFILER_STATUS_ERROR_SESSION_MISSING_BUFFER = -4,
   ROCPROFILER_STATUS_ERROR_TIMESTAMP_NOT_APPLICABLE = -5,
   ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND = -6,
   ROCPROFILER_STATUS_ERROR_AGENT_INFORMATION_MISSING = -7,
   ROCPROFILER_STATUS_ERROR_QUEUE_NOT_FOUND = -8,
   ROCPROFILER_STATUS_ERROR_QUEUE_INFORMATION_MISSING = -9,
   ROCPROFILER_STATUS_ERROR_KERNEL_NOT_FOUND = -10,
   ROCPROFILER_STATUS_ERROR_KERNEL_INFORMATION_MISSING = -11,
   ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND = -12,
   ROCPROFILER_STATUS_ERROR_COUNTER_INFORMATION_MISSING = -13,
   ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_NOT_FOUND = -14,
   ROCPROFILER_STATUS_ERROR_TRACER_API_DATA_INFORMATION_MISSING = -15,
   ROCPROFILER_STATUS_ERROR_INCORRECT_DOMAIN = -16,
   ROCPROFILER_STATUS_ERROR_SESSION_NOT_FOUND = -17,
   ROCPROFILER_STATUS_ERROR_CORRUPTED_SESSION_BUFFER = -18,
   ROCPROFILER_STATUS_ERROR_RECORD_CORRUPTED = -19,
   ROCPROFILER_STATUS_ERROR_INCORRECT_REPLAY_MODE = -20,
   ROCPROFILER_STATUS_ERROR_SESSION_MISSING_FILTER = -21,
   ROCPROFILER_STATUS_ERROR_INCORRECT_SIZE = -22,
   ROCPROFILER_STATUS_ERROR_INCORRECT_FLUSH_INTERVAL = -23,
   ROCPROFILER_STATUS_ERROR_SESSION_FILTER_DATA_MISMATCH = -24,
   ROCPROFILER_STATUS_ERROR_FILTER_DATA_CORRUPTED = -25,
   ROCPROFILER_STATUS_ERROR_CORRUPTED_LABEL_DATA = -26,
   ROCPROFILER_STATUS_ERROR_RANGE_STACK_IS_EMPTY = -27,
   ROCPROFILER_STATUS_ERROR_PASS_NOT_STARTED = -28,
   ROCPROFILER_STATUS_ERROR_HAS_ACTIVE_SESSION = -29,
   ROCPROFILER_STATUS_ERROR_SESSION_NOT_ACTIVE = -30,
   ROCPROFILER_STATUS_ERROR_FILTER_NOT_FOUND = -31,
   ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND = -32,
   ROCPROFILER_STATUS_ERROR_FILTER_NOT_SUPPORTED = -33
 } rocprofiler_status_t;
  
 ROCPROFILER_API const char* rocprofiler_error_str(rocprofiler_status_t status) ROCPROFILER_VERSION_9_0;
  
 // TODO(aelwazir): More clear description, (think about nested!!??)
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_initialize() ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_finalize() ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   uint64_t handle;
 } rocprofiler_session_id_t;
  
 typedef struct {
   uint64_t value;
 } rocprofiler_timestamp_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_get_timestamp(rocprofiler_timestamp_t* timestamp) ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   rocprofiler_timestamp_t begin;
   rocprofiler_timestamp_t end;
 } rocprofiler_record_header_timestamp_t;
  
 typedef struct {
   uint64_t handle;
 } rocprofiler_record_id_t;
  
 typedef enum {
   ROCPROFILER_PROFILER_RECORD = 0,
   ROCPROFILER_TRACER_RECORD = 1,
   ROCPROFILER_ATT_TRACER_RECORD = 2,
   ROCPROFILER_PC_SAMPLING_RECORD = 3,
   ROCPROFILER_SPM_RECORD = 4,
   ROCPROFILER_COUNTERS_SAMPLER_RECORD = 5
 } rocprofiler_record_kind_t;
  
 typedef struct {
   rocprofiler_record_kind_t kind;
   rocprofiler_record_id_t id;
 } rocprofiler_record_header_t;
  
 typedef struct {
   uint64_t handle;  // Topology folder serial number
 } rocprofiler_agent_id_t;
  
 typedef enum {
   ROCPROFILER_CPU_AGENT = 0,
   ROCPROFILER_GPU_AGENT = 1
 } rocprofiler_agent_type_t;
  
 // TODO(aelwazir): check if we need to report the family name as well!!?? OR
 // return the agent itself so that they can use HSA API
 typedef enum {
   ROCPROFILER_AGENT_NAME = 0,
   ROCPROFILER_AGENT_TYPE = 1
 } rocprofiler_agent_info_kind_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_agent_info_size(rocprofiler_agent_info_kind_t kind,
                                                                  rocprofiler_agent_id_t agent_id,
                                                                  size_t* data_size) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_agent_info(rocprofiler_agent_info_kind_t kind,
                                                             rocprofiler_agent_id_t descriptor,
                                                             const char** name) ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   uint64_t handle;
 } rocprofiler_queue_id_t;
  
 // TODO(aelwazir): Check if there is anymore Queue Information needed
 typedef enum {
   ROCPROFILER_QUEUE_SIZE = 0
 } rocprofiler_queue_info_kind_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_queue_info_size(rocprofiler_queue_info_kind_t kind,
                                                                  rocprofiler_queue_id_t agent_id,
                                                                  size_t* data_size) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_queue_info(rocprofiler_queue_info_kind_t kind,
                                                             rocprofiler_queue_id_t descriptor,
                                                             const char** name) ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   uint64_t handle;
 } rocprofiler_kernel_id_t;
  
 typedef enum {
   ROCPROFILER_KERNEL_NAME = 0
 } rocprofiler_kernel_info_kind_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_kernel_info_size(rocprofiler_kernel_info_kind_t kind,
                                                                   rocprofiler_kernel_id_t kernel_id,
                                                                   size_t* data_size) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_kernel_info(rocprofiler_kernel_info_kind_t kind,
                                                              rocprofiler_kernel_id_t kernel_id,
                                                              const char** data) ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   uint32_t value;
 } rocprofiler_thread_id_t;
  
 typedef struct {
   const char* name;
   const char* description;
   const char* expression;
   uint32_t instances_count;
   const char* block_name;
   uint32_t block_counters;
 } rocprofiler_counter_info_t;
  
 typedef int (*rocprofiler_counters_info_callback_t)(rocprofiler_counter_info_t counter,
                                                   const char* gpu_name, uint32_t gpu_index) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t
 rocprofiler_iterate_counters(rocprofiler_counters_info_callback_t counters_info_callback) ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   uint64_t handle;
 } rocprofiler_counter_id_t;
  
 typedef enum {
   ROCPROFILER_COUNTER_NAME = 0,
   ROCPROFILER_COUNTER_BLOCK_ID = 2,
   ROCPROFILER_COUNTER_HIERARCHY_LEVEL = 3
 } rocprofiler_counter_info_kind_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_counter_info_size(
     rocprofiler_session_id_t session_id, rocprofiler_counter_info_kind_t counter_info_type,
     rocprofiler_counter_id_t counter_id, size_t* data_size) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_counter_info(rocprofiler_session_id_t session_id,
                                                               rocprofiler_counter_info_kind_t kind,
                                                               rocprofiler_counter_id_t counter_id,
                                                               const char** data) ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   uint64_t value;
 } rocprofiler_queue_index_t;
  
 // TODO(aelwazir): add more types to the values should we use unions??!!
 typedef struct {
   double value;
 } rocprofiler_record_counter_value_t;
  
 typedef struct {
   rocprofiler_counter_id_t counter_handler;  // Counter Handler
   rocprofiler_record_counter_value_t value;  // Counter Value
 } rocprofiler_record_counter_instance_t;
  
 typedef struct {
   uint64_t value;
 } rocprofiler_record_counters_instances_count_t;
  
 typedef struct {
   uint64_t grid_size;
   uint64_t workgroup_size;
   uint64_t lds_size;
   uint64_t scratch_size;
   uint64_t arch_vgpr_count;
   uint64_t accum_vgpr_count;
   uint64_t sgpr_count;
   uint64_t wave_size;
   uint64_t signal_handle;
  
 } rocprofiler_kernel_properties_t;
 typedef struct {
   rocprofiler_record_header_t header;
   rocprofiler_kernel_id_t kernel_id;
   rocprofiler_agent_id_t gpu_id;
   rocprofiler_queue_id_t queue_id;
   rocprofiler_record_header_timestamp_t timestamps;
   rocprofiler_record_counter_instance_t* counters;
   rocprofiler_kernel_properties_t kernel_properties;
   rocprofiler_thread_id_t thread_id;
   rocprofiler_queue_index_t queue_idx;
   rocprofiler_record_counters_instances_count_t counters_count; /* Counters Count */
 } rocprofiler_record_profiler_t;
  
 typedef struct {
   uint32_t value;
  
 } rocprofiler_event_id_t;
  
 typedef struct {
   uint16_t value;  // Counter Value
  
 } rocprofiler_record_spm_counters_instances_count_t;
  
 typedef struct {
   rocprofiler_record_spm_counters_instances_count_t counters_data[32];
  
 } rocprofiler_record_se_spm_data_t;
  
  
 typedef struct {
   rocprofiler_record_header_t header;
  
   rocprofiler_record_header_timestamp_t timestamps;
   rocprofiler_record_se_spm_data_t shader_engine_data[4];
  
 } rocprofiler_record_spm_t;
  
 typedef struct {
   void* buffer_ptr;
   uint32_t buffer_size;
 } rocprofiler_record_se_att_data_t;
  
 typedef struct {
   rocprofiler_record_header_t header;
   rocprofiler_kernel_id_t kernel_id;
   rocprofiler_agent_id_t gpu_id;
   rocprofiler_queue_id_t queue_id;
   rocprofiler_kernel_properties_t kernel_properties;
   rocprofiler_thread_id_t thread_id;
   rocprofiler_queue_index_t queue_idx;
   rocprofiler_record_se_att_data_t* shader_engine_data;
   uint64_t shader_engine_data_count;
 } rocprofiler_record_att_tracer_t;
  
  
  
 typedef enum {
   ACTIVITY_DOMAIN_HSA_API = 0,
   ACTIVITY_DOMAIN_HSA_OPS = 1,
   ACTIVITY_DOMAIN_HIP_OPS = 2,
   ACTIVITY_DOMAIN_HIP_API = 3,
   ACTIVITY_DOMAIN_KFD_API = 4,
   ACTIVITY_DOMAIN_EXT_API = 5,
   ACTIVITY_DOMAIN_ROCTX = 6,
   // TODO(aelwazir): Used in kernel Info, memcpy, ..etc, refer to hsa_support
   // TODO(aelwazir): Move HSA Events to hsa_support
   ACTIVITY_DOMAIN_HSA_EVT = 7,
   ACTIVITY_DOMAIN_NUMBER
 } rocprofiler_tracer_activity_domain_t;
  
 typedef struct {
   uint32_t id;
 } rocprofiler_tracer_operation_id_t;
  
 typedef struct {
   uint64_t value;
 } rocprofiler_tracer_activity_correlation_id_t;
  
 typedef struct {
   const void* handle;
   size_t size;
 } rocprofiler_tracer_api_data_handle_t;
  
 typedef enum {
   ROCPROFILER_ROCTX_MESSAGE = 0,
   ROCPROFILER_ROCTX_ID = 1
 } rocprofiler_tracer_roctx_api_data_info_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_roctx_tracer_api_data_info_size(
     rocprofiler_session_id_t session_id, rocprofiler_tracer_roctx_api_data_info_t kind,
     rocprofiler_tracer_api_data_handle_t api_data_id, rocprofiler_tracer_operation_id_t operation_id,
     size_t* data_size) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_roctx_tracer_api_data_info(
     rocprofiler_session_id_t session_id, rocprofiler_tracer_roctx_api_data_info_t kind,
     rocprofiler_tracer_api_data_handle_t api_data_id, rocprofiler_tracer_operation_id_t operation_id,
     char** data) ROCPROFILER_VERSION_9_0;
  
 typedef enum {
   ROCPROFILER_HSA_FUNCTION_NAME = 0,
   ROCPROFILER_HSA_API_DATA_STR = 1,
   ROCPROFILER_HSA_ACTIVITY_NAME = 2,
   ROCPROFILER_HSA_API_DATA = 3
 } rocprofiler_tracer_hsa_api_data_info_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hsa_tracer_api_data_info_size(
     rocprofiler_session_id_t session_id, rocprofiler_tracer_hsa_api_data_info_t kind,
     rocprofiler_tracer_api_data_handle_t api_data_id, rocprofiler_tracer_operation_id_t operation_id,
     size_t* data_size) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hsa_tracer_api_data_info(
     rocprofiler_session_id_t session_id, rocprofiler_tracer_hsa_api_data_info_t kind,
     rocprofiler_tracer_api_data_handle_t api_data_id, rocprofiler_tracer_operation_id_t operation_id,
     char** data) ROCPROFILER_VERSION_9_0;
  
 typedef enum {
   // TODO(aelwazir): Get the data from hip_api_data_t
   ROCPROFILER_HIP_FUNCTION_NAME = 0,
   ROCPROFILER_HIP_KERNEL_NAME = 1,
   ROCPROFILER_HIP_MEM_COPY_SRC = 2,
   ROCPROFILER_HIP_MEM_COPY_DST = 3,
   ROCPROFILER_HIP_MEM_COPY_SIZE = 4,
   ROCPROFILER_HIP_API_DATA_STR = 5,
   ROCPROFILER_HIP_ACTIVITY_NAME = 6,
   ROCPROFILER_HIP_STREAM_ID = 7,
   ROCPROFILER_HIP_API_DATA = 8
 } rocprofiler_tracer_hip_api_data_info_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hip_tracer_api_data_info_size(
     rocprofiler_session_id_t session_id, rocprofiler_tracer_hip_api_data_info_t kind,
     rocprofiler_tracer_api_data_handle_t api_data_id, rocprofiler_tracer_operation_id_t operation_id,
     size_t* data_size) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_query_hip_tracer_api_data_info(
     rocprofiler_session_id_t session_id, rocprofiler_tracer_hip_api_data_info_t kind,
     rocprofiler_tracer_api_data_handle_t api_data_id, rocprofiler_tracer_operation_id_t operation_id,
     char** data) ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   uint64_t id;
 } rocprofiler_tracer_external_id_t;
  
 typedef struct {
   rocprofiler_record_header_t header;
   rocprofiler_tracer_external_id_t external_id;
   rocprofiler_tracer_activity_domain_t domain;
   rocprofiler_tracer_operation_id_t operation_id;
   rocprofiler_tracer_api_data_handle_t api_data_handle;
   rocprofiler_tracer_activity_correlation_id_t correlation_id;
   rocprofiler_record_header_timestamp_t timestamps;
   rocprofiler_agent_id_t agent_id;
   rocprofiler_queue_id_t queue_id;
   rocprofiler_thread_id_t thread_id;
 } rocprofiler_record_tracer_t;
  
 typedef struct {
   uint64_t value;
 } rocprofiler_kernel_dispatch_id_t;
  
 typedef struct {
   rocprofiler_kernel_dispatch_id_t dispatch_id;
   union {
     rocprofiler_timestamp_t timestamp;
     uint64_t cycle;
   };
   uint64_t pc;
   uint32_t se;
   rocprofiler_agent_id_t gpu_id;
 } rocprofiler_pc_sample_t;
  
 typedef struct {
   rocprofiler_record_header_t header;
   rocprofiler_pc_sample_t pc_sample;
 } rocprofiler_record_pc_sample_t;
  
 typedef enum {
   ROCPROFILER_BUFFER_PROPERTY_KIND_INTERVAL_FLUSH = 0,
   // Periodic Flush
   // Size
   // Think of using the kind as an end of the array!!??
 } rocprofiler_buffer_property_kind_t;
  
 typedef struct {
   rocprofiler_buffer_property_kind_t kind;
   uint64_t value;
 } rocprofiler_buffer_property_t;
  
 typedef struct {
   uint64_t value;
 } rocprofiler_buffer_id_t;
  
 typedef struct {
   uint64_t value;
 } rocprofiler_filter_id_t;
  
 typedef void (*rocprofiler_buffer_callback_t)(const rocprofiler_record_header_t* begin,
                                             const rocprofiler_record_header_t* end,
                                             rocprofiler_session_id_t session_id,
                                             rocprofiler_buffer_id_t buffer_id);
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_flush_data(rocprofiler_session_id_t session_id,
                                                       rocprofiler_buffer_id_t buffer_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_next_record(const rocprofiler_record_header_t* record,
                                                        const rocprofiler_record_header_t** next,
                                                        rocprofiler_session_id_t session_id,
                                                        rocprofiler_buffer_id_t buffer_id) ROCPROFILER_VERSION_9_0;
  
 // TODO(aelwazir): Replay mode naming !!?? (If changed, reflect on start&stop)
 typedef enum {
   ROCPROFILER_NONE_REPLAY_MODE = -1,
   ROCPROFILER_APPLICATION_REPLAY_MODE = 0,
   ROCPROFILER_KERNEL_REPLAY_MODE = 1,
   ROCPROFILER_USER_REPLAY_MODE = 2
 } rocprofiler_replay_mode_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_create_session(rocprofiler_replay_mode_t replay_mode,
                                                           rocprofiler_session_id_t* session_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_destroy_session(rocprofiler_session_id_t session_id) ROCPROFILER_VERSION_9_0;
  
 typedef enum {
   ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION = 1,
   ROCPROFILER_COUNTERS_COLLECTION = 2,
   ROCPROFILER_PC_SAMPLING_COLLECTION = 3,
   ROCPROFILER_ATT_TRACE_COLLECTION = 4,
   ROCPROFILER_SPM_COLLECTION = 5,
   ROCPROFILER_API_TRACE = 6,
   ROCPROFILER_COUNTERS_SAMPLER = 7
 } rocprofiler_filter_kind_t;
  
 typedef enum {
   ROCPROFILER_FILTER_HSA_TRACER_API_FUNCTIONS = 1,
   ROCPROFILER_FILTER_HIP_TRACER_API_FUNCTIONS = 2,
   ROCPROFILER_FILTER_GPU_NAME = 3,
   // TODO(aelwazir): Add more clear description on how to use?
   ROCPROFILER_FILTER_RANGE = 4,
   ROCPROFILER_FILTER_KERNEL_NAMES = 5
 } rocprofiler_filter_property_kind_t;
  
 // TODO(aelwazir): Another way to define this as needed
 typedef const char* rocprofiler_hip_function_name_t;
 typedef const char* rocprofiler_hsa_function_name_t;
  
 // ATT tracing parameter names
 typedef enum {
   ROCPROFILER_ATT_COMPUTE_UNIT_TARGET = 0,
   ROCPROFILER_ATT_VM_ID_MASK = 1,
   ROCPROFILER_ATT_MASK = 2,
   ROCPROFILER_ATT_TOKEN_MASK = 3,
   ROCPROFILER_ATT_TOKEN_MASK2 = 4,
   ROCPROFILER_ATT_SE_MASK = 5,
   ROCPROFILER_ATT_SAMPLE_RATE = 6,
   ROCPROFILER_ATT_PERF_MASK = 240,
   ROCPROFILER_ATT_PERF_CTRL = 241,
   ROCPROFILER_ATT_PERFCOUNTER = 242,
   ROCPROFILER_ATT_PERFCOUNTER_NAME = 243,
   ROCPROFILER_ATT_MAXVALUE
 } rocprofiler_att_parameter_name_t;
  
 // att tracing parameters object
 typedef struct {
   rocprofiler_att_parameter_name_t parameter_name;
   union {
     uint32_t value;
     const char* counter_name;
   };
 } rocprofiler_att_parameter_t;
  
 typedef struct {
   rocprofiler_filter_property_kind_t kind;
   // TODO(aelwazir): get HIP or HSA or counters as enums
   union {
     const char** name_regex;
     rocprofiler_hip_function_name_t* hip_functions_names;
     rocprofiler_hsa_function_name_t* hsa_functions_names;
     uint32_t range[2];
   };
   uint64_t data_count;
 } rocprofiler_filter_property_t;
  
 typedef struct {
   const char** counters_names;
   int counters_count;
   uint32_t sampling_rate;
   rocprofiler_agent_id_t* gpu_agent_id;
  
 } rocprofiler_spm_parameter_t;
  
 typedef enum{
   ROCPROFILER_COUNTERS_SAMPLER_PCIE_COUNTERS = 0,
   ROCPROFILER_COUNTERS_SAMPLER_XGMI_COUNTERS = 1
 } rocprofiler_counters_sampler_counter_type_t;
  
 typedef struct{
   char* name;
   rocprofiler_counters_sampler_counter_type_t type;
 } rocprofiler_counters_sampler_counter_input_t;
  
 typedef struct{
   rocprofiler_counters_sampler_counter_type_t type;
   rocprofiler_record_counter_value_t value;
 } rocprofiler_counters_sampler_counter_output_t;
  
 typedef struct{
   rocprofiler_counters_sampler_counter_input_t* counters;
   int counters_num;
   uint32_t sampling_rate;
   uint32_t sampling_duration;
   uint32_t initial_delay;
   int gpu_agent_index;
 }rocprofiler_counters_sampler_parameters_t;
  
 typedef struct{
   rocprofiler_record_header_t header;
   rocprofiler_agent_id_t gpu_id;
   rocprofiler_counters_sampler_counter_output_t* counters;
   uint32_t num_counters;
 }rocprofiler_record_counters_sampler_t;
  
 typedef union {
   rocprofiler_tracer_activity_domain_t* trace_apis;
   const char** counters_names;
   rocprofiler_att_parameter_t* att_parameters;
   rocprofiler_spm_parameter_t* spm_parameters;
   rocprofiler_counters_sampler_parameters_t counters_sampler_parameters;
 } rocprofiler_filter_data_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_create_filter(rocprofiler_session_id_t session_id,
                                                          rocprofiler_filter_kind_t filter_kind,
                                                          rocprofiler_filter_data_t data,
                                                          uint64_t data_count,
                                                          rocprofiler_filter_id_t* filter_id,
                                                          rocprofiler_filter_property_t property) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_set_filter_buffer(rocprofiler_session_id_t session_id,
                                                              rocprofiler_filter_id_t filter_id,
                                                              rocprofiler_buffer_id_t buffer_id) ROCPROFILER_VERSION_9_0;
  
 typedef void (*rocprofiler_sync_callback_t)(rocprofiler_record_tracer_t record,
                                           rocprofiler_session_id_t session_id);
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_set_api_trace_sync_callback(
     rocprofiler_session_id_t session_id, rocprofiler_filter_id_t filter_id,
     rocprofiler_sync_callback_t callback) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_destroy_filter(rocprofiler_session_id_t session_id,
                                                           rocprofiler_filter_id_t filter_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_create_buffer(
     rocprofiler_session_id_t session_id, rocprofiler_buffer_callback_t buffer_callback,
     size_t buffer_size, rocprofiler_buffer_id_t* buffer_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_set_buffer_properties(
     rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id,
     rocprofiler_buffer_property_t* buffer_properties, uint32_t buffer_properties_count) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_destroy_buffer(rocprofiler_session_id_t session_id,
                                                           rocprofiler_buffer_id_t buffer_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_create_ready_session(
     rocprofiler_replay_mode_t replay_mode, rocprofiler_filter_kind_t filter_kind,
     rocprofiler_filter_data_t data, uint64_t data_count, size_t buffer_size,
     rocprofiler_buffer_callback_t buffer_callback, rocprofiler_session_id_t* session_id,
     rocprofiler_filter_property_t property, rocprofiler_sync_callback_t callback) ROCPROFILER_VERSION_9_0;
  
 // TODO(aelwazir): Multiple sessions activate for different set of filters
 ROCPROFILER_API rocprofiler_status_t rocprofiler_start_session(rocprofiler_session_id_t session_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_terminate_session(rocprofiler_session_id_t session_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_push_range(const char* label) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_pop_range() ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_start_replay_pass() ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_end_replay_pass() ROCPROFILER_VERSION_9_0;
  
 typedef struct {
   double value;
 } rocprofiler_counter_value_t;
  
 typedef struct {
   char metric_name[64];
   rocprofiler_counter_value_t value;
 } rocprofiler_device_profile_metric_t;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_create(
     const char** counter_names, uint64_t num_counters, rocprofiler_session_id_t* session_id,
     int cpu_index, int gpu_index) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t
 rocprofiler_device_profiling_session_start(rocprofiler_session_id_t session_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_poll(
     rocprofiler_session_id_t session_id, rocprofiler_device_profile_metric_t* data) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t
 rocprofiler_device_profiling_session_stop(rocprofiler_session_id_t session_id) ROCPROFILER_VERSION_9_0;
  
 ROCPROFILER_API rocprofiler_status_t
 rocprofiler_device_profiling_session_destroy(rocprofiler_session_id_t session_id) ROCPROFILER_VERSION_9_0;
  
 #ifdef __cplusplus
 }  // extern "C" block
 #endif  // __cplusplus
  
 #endif  // INC_ROCPROFILER_H_