roc_profiler.c File Reference

Include dependency graph for roc_profiler.c:

Go to the source code of this file.

Data Structures
struct	ntv_event_t
struct	ntv_event_table_t
struct	rocp_ctx_t
struct	event_info_t
struct	ntv_arg
struct	cb_context_node_t
struct	cb_dispatch_arg_t
struct	cb_context_arg_t
struct	cb_context_payload_t

Macros
#define	EVENTS_WIDTH   (sizeof(uint64_t) * 8)
#define	DEVICE_WIDTH   ( 7)
#define	INSTAN_WIDTH   ( 7)
#define	QLMASK_WIDTH   ( 2)
#define	NAMEID_WIDTH   (12)
#define	UNUSED_WIDTH   (EVENTS_WIDTH - DEVICE_WIDTH - INSTAN_WIDTH - QLMASK_WIDTH - NAMEID_WIDTH)
#define	DEVICE_SHIFT   (EVENTS_WIDTH - UNUSED_WIDTH - DEVICE_WIDTH)
#define	INSTAN_SHIFT   (DEVICE_SHIFT - INSTAN_WIDTH)
#define	QLMASK_SHIFT   (INSTAN_SHIFT - QLMASK_WIDTH)
#define	NAMEID_SHIFT   (QLMASK_SHIFT - NAMEID_WIDTH)
#define	DEVICE_MASK   ((0xFFFFFFFFFFFFFFFF >> (EVENTS_WIDTH - DEVICE_WIDTH)) << DEVICE_SHIFT)
#define	INSTAN_MASK   ((0xFFFFFFFFFFFFFFFF >> (EVENTS_WIDTH - INSTAN_WIDTH)) << INSTAN_SHIFT)
#define	QLMASK_MASK   ((0xFFFFFFFFFFFFFFFF >> (EVENTS_WIDTH - QLMASK_WIDTH)) << QLMASK_SHIFT)
#define	NAMEID_MASK   ((0xFFFFFFFFFFFFFFFF >> (EVENTS_WIDTH - NAMEID_WIDTH)) << NAMEID_SHIFT)
#define	DEVICE_FLAG   (0x2)
#define	INSTAN_FLAG   (0x1)

Functions
static int	load_rocp_sym (void)
static int	init_rocp_env (void)
static int	init_event_table (void)
static int	unload_rocp_sym (void)
static int	sampling_ctx_open (uint64_t *, int, rocp_ctx_t *)
static int	intercept_ctx_open (uint64_t *, int, rocp_ctx_t *)
static int	sampling_ctx_close (rocp_ctx_t)
static int	intercept_ctx_close (rocp_ctx_t)
static int	sampling_ctx_start (rocp_ctx_t)
static int	intercept_ctx_start (rocp_ctx_t)
static int	sampling_ctx_stop (rocp_ctx_t)
static int	intercept_ctx_stop (rocp_ctx_t)
static int	sampling_ctx_read (rocp_ctx_t, long long **)
static int	intercept_ctx_read (rocp_ctx_t, long long **)
static int	sampling_ctx_reset (rocp_ctx_t)
static int	intercept_ctx_reset (rocp_ctx_t)
static int	sampling_shutdown (void)
static int	intercept_shutdown (void)
static int	evt_code_to_name (uint64_t event_code, char *name, int len)
static int	evt_id_create (event_info_t *info, uint64_t *event_id)
static int	evt_id_to_info (uint64_t event_id, event_info_t *info)
static int	evt_name_to_device (const char *name, int *device)
static int	evt_name_to_instance (const char *name, int *instance)
static int	evt_name_to_basename (const char *name, char *base, int len)
int	rocp_init_environment (void)
int	rocp_init (void)
int	rocp_evt_enum (uint64_t *event_code, int modifier)
int	rocp_evt_code_to_descr (uint64_t event_code, char *descr, int len)
int	rocp_evt_name_to_code (const char *name, uint64_t *event_code)
int	rocp_evt_code_to_name (uint64_t event_code, char *name, int len)
int	rocp_evt_code_to_info (uint64_t event_code, PAPI_event_info_t *info)
int	rocp_ctx_open (uint64_t *events_id, int num_events, rocp_ctx_t *rocp_ctx)
int	rocp_ctx_close (rocp_ctx_t rocp_ctx)
int	rocp_ctx_start (rocp_ctx_t rocp_ctx)
int	rocp_ctx_stop (rocp_ctx_t rocp_ctx)
int	rocp_ctx_read (rocp_ctx_t rocp_ctx, long long **counts)
int	rocp_ctx_reset (rocp_ctx_t rocp_ctx)
int	rocp_shutdown (void)
static hsa_status_t	count_ntv_events_cb (const rocprofiler_info_data_t, void *)
static hsa_status_t	get_ntv_events_cb (const rocprofiler_info_data_t, void *)
hsa_status_t	count_ntv_events_cb (const rocprofiler_info_data_t info __attribute__((unused)), void *count)
static int	init_features (uint64_t *, int, rocprofiler_feature_t *)
static int	finalize_features (rocprofiler_feature_t *, int)
static int	sampling_ctx_init (uint64_t *, int, rocp_ctx_t *)
static int	sampling_ctx_finalize (rocp_ctx_t *)
static int	ctx_open (rocp_ctx_t)
static int	ctx_close (rocp_ctx_t)
static int	ctx_init (uint64_t *, int, rocp_ctx_t *)
static int	ctx_finalize (rocp_ctx_t *)
static int	ctx_get_dev_feature_count (rocp_ctx_t, int)
static int	shutdown_event_table (void)
static int	event_id_to_dev_id_cb (uint64_t event_id, int *device)
static int	sampling_ctx_get_dev_feature_count (rocp_ctx_t, int)
static int	intercept_ctx_get_dev_feature_count (rocp_ctx_t, int)
static int	verify_events (uint64_t *, int)
static int	init_callbacks (rocprofiler_feature_t *, int)
static int	register_dispatch_counter (unsigned long, int *)
static int	increment_and_fetch_dispatch_counter (unsigned long)
static int	decrement_and_fetch_dispatch_counter (unsigned long)
static int	unregister_dispatch_counter (unsigned long)
static int	fetch_dispatch_counter (unsigned long)
static cb_context_node_t *	alloc_context_node (int)
static void	free_context_node (cb_context_node_t *)
static int	get_context_node (int, cb_context_node_t **)
static int	get_context_counters (int, cb_context_node_t *, rocp_ctx_t)
static void	put_context_counters (rocprofiler_feature_t *, int, cb_context_node_t *)
static void	put_context_node (int, cb_context_node_t *)
static int	intercept_ctx_init (uint64_t *, int, rocp_ctx_t *)
static int	intercept_ctx_finalize (rocp_ctx_t *)
static int	count_unique_events (uint64_t *events_id, int num_events, int *num_unique)
static int	copy_unique_events (uint64_t *target, uint64_t *source, int source_len)
static int	save_callback_features (rocprofiler_feature_t *features, int feature_count)
static int	cleanup_callback_features (rocprofiler_feature_t *features, int feature_count)
static bool	context_handler_cb (const rocprofiler_pool_entry_t *, void *)
static hsa_status_t	dispatch_cb (const rocprofiler_callback_data_t *, void *, rocprofiler_group_t *)
static void	process_context_entry (cb_context_payload_t *, rocprofiler_feature_t *, int)
void	__attribute__ ((visibility("default")))

Variables
unsigned int	rocm_prof_mode
unsigned int	_rocm_lock
static hsa_status_t(*	rocp_get_info_p )(const hsa_agent_t *, rocprofiler_info_kind_t, void *)
static hsa_status_t(*	rocp_iterate_info_p )(const hsa_agent_t *, rocprofiler_info_kind_t, hsa_status_t(*)(const rocprofiler_info_data_t, void *), void *)
static hsa_status_t(*	rocp_error_string_p )(const char **)
static hsa_status_t(*	rocp_open_p )(hsa_agent_t, rocprofiler_feature_t *, uint32_t, rocprofiler_t **, uint32_t, rocprofiler_properties_t *)
static hsa_status_t(*	rocp_close_p )(rocprofiler_t *)
static hsa_status_t(*	rocp_group_count_p )(const rocprofiler_t *, uint32_t *)
static hsa_status_t(*	rocp_start_p )(rocprofiler_t *, uint32_t)
static hsa_status_t(*	rocp_read_p )(rocprofiler_t *, uint32_t)
static hsa_status_t(*	rocp_stop_p )(rocprofiler_t *, uint32_t)
static hsa_status_t(*	rocp_get_group_p )(rocprofiler_t *, uint32_t, rocprofiler_group_t *)
static hsa_status_t(*	rocp_get_data_p )(rocprofiler_t *, uint32_t)
static hsa_status_t(*	rocp_group_get_data_p )(rocprofiler_group_t *)
static hsa_status_t(*	rocp_get_metrics_p )(const rocprofiler_t *)
static hsa_status_t(*	rocp_reset_p )(rocprofiler_t *, uint32_t)
static hsa_status_t(*	rocp_pool_open_p )(hsa_agent_t, rocprofiler_feature_t *, uint32_t, rocprofiler_pool_t **, uint32_t, rocprofiler_pool_properties_t *)
static hsa_status_t(*	rocp_pool_close_p )(rocprofiler_pool_t *)
static hsa_status_t(*	rocp_pool_fetch_p )(rocprofiler_pool_t *, rocprofiler_pool_entry_t *)
static hsa_status_t(*	rocp_pool_flush_p )(rocprofiler_pool_t *)
static hsa_status_t(*	rocp_set_queue_cbs_p )(rocprofiler_queue_callbacks_t, void *)
static hsa_status_t(*	rocp_start_queue_cbs_p )(void)
static hsa_status_t(*	rocp_stop_queue_cbs_p )(void)
static hsa_status_t(*	rocp_remove_queue_cbs_p )(void)
static void *	rocp_dlp = NULL
static ntv_event_table_t	ntv_table
static ntv_event_table_t *	ntv_table_p
static void *	htable
static void *	htable_intercept
struct {
uint64_t *   events_id
rocprofiler_feature_t *   features
int   feature_count
int   active_thread_count
int   kernel_count
}	intercept_global_state
static cb_dispatch_arg_t	cb_dispatch_arg
static cb_context_node_t *	cb_ctx_list_heads [PAPI_ROCM_MAX_DEV_COUNT]

Detailed Description

Author

Giuseppe Congiu gcong.nosp@m.iu@i.nosp@m.cl.ut.nosp@m.k.ed.nosp@m.u

Definition in file roc_profiler.c.

Macro Definition Documentation

DEVICE_FLAG

#define DEVICE_FLAG   (0x2)

Definition at line 39 of file roc_profiler.c.

DEVICE_MASK

#define DEVICE_MASK   ((0xFFFFFFFFFFFFFFFF >> (EVENTS_WIDTH - DEVICE_WIDTH)) << DEVICE_SHIFT)

Definition at line 35 of file roc_profiler.c.

DEVICE_SHIFT

#define DEVICE_SHIFT   (EVENTS_WIDTH - UNUSED_WIDTH - DEVICE_WIDTH)

Definition at line 31 of file roc_profiler.c.

DEVICE_WIDTH

#define DEVICE_WIDTH   ( 7)

Definition at line 26 of file roc_profiler.c.

EVENTS_WIDTH

#define EVENTS_WIDTH   (sizeof(uint64_t) * 8)

Event identifier encoding format: +------------------------------—+----—+----—+–+---------—+ | unused | dev | inst | | nameid | +------------------------------—+----—+----—+–+---------—+

unused : 36 bits device : 7 bits ([0 - 127] devices) instance : 7 bits ([0 - 127] instances) qlmask : 2 bits (qualifier mask) nameid : 12 bits ([0 - 4095] event names)

Definition at line 25 of file roc_profiler.c.

INSTAN_FLAG

#define INSTAN_FLAG   (0x1)

Definition at line 40 of file roc_profiler.c.

INSTAN_MASK

#define INSTAN_MASK   ((0xFFFFFFFFFFFFFFFF >> (EVENTS_WIDTH - INSTAN_WIDTH)) << INSTAN_SHIFT)

Definition at line 36 of file roc_profiler.c.

INSTAN_SHIFT

#define INSTAN_SHIFT   (DEVICE_SHIFT - INSTAN_WIDTH)

Definition at line 32 of file roc_profiler.c.

INSTAN_WIDTH

#define INSTAN_WIDTH   ( 7)

Definition at line 27 of file roc_profiler.c.

NAMEID_MASK

#define NAMEID_MASK   ((0xFFFFFFFFFFFFFFFF >> (EVENTS_WIDTH - NAMEID_WIDTH)) << NAMEID_SHIFT)

Definition at line 38 of file roc_profiler.c.

NAMEID_SHIFT

#define NAMEID_SHIFT   (QLMASK_SHIFT - NAMEID_WIDTH)

Definition at line 34 of file roc_profiler.c.

NAMEID_WIDTH

#define NAMEID_WIDTH   (12)

Definition at line 29 of file roc_profiler.c.

QLMASK_MASK

#define QLMASK_MASK   ((0xFFFFFFFFFFFFFFFF >> (EVENTS_WIDTH - QLMASK_WIDTH)) << QLMASK_SHIFT)

Definition at line 37 of file roc_profiler.c.

QLMASK_SHIFT

#define QLMASK_SHIFT   (INSTAN_SHIFT - QLMASK_WIDTH)

Definition at line 33 of file roc_profiler.c.

QLMASK_WIDTH

#define QLMASK_WIDTH   ( 2)

Definition at line 28 of file roc_profiler.c.

UNUSED_WIDTH

#define UNUSED_WIDTH   (EVENTS_WIDTH - DEVICE_WIDTH - INSTAN_WIDTH - QLMASK_WIDTH - NAMEID_WIDTH)

Definition at line 30 of file roc_profiler.c.

Function Documentation

__attribute__()

void __attribute__

(

(visibility("default"))

)

Definition at line 2374 of file roc_profiler.c.

{
    init_rocp_env();
}

Here is the call graph for this function:

alloc_context_node()

cb_context_node_t * alloc_context_node ( int  num_events )

static

Definition at line 2211 of file roc_profiler.c.

{
    cb_context_node_t *n = papi_malloc(sizeof(*n));
    if (n == NULL) {
        return NULL;
    }
 
    n->counters = papi_malloc(num_events * sizeof(long long));
    if (n->counters == NULL) {
        papi_free(n);
        return NULL;
    }
 
    return n;
}

Here is the caller graph for this function:

cleanup_callback_features()

int cleanup_callback_features ( rocprofiler_feature_t *  features, int  feature_count  )

static

Definition at line 1916 of file roc_profiler.c.

{
    int i;
    for (i = 0; i < feature_count; ++i) {
        htable_delete(htable_intercept, features[i].name);
    }
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

context_handler_cb()

bool context_handler_cb ( const rocprofiler_pool_entry_t *  entry, void *  arg  )

static

Definition at line 2156 of file roc_profiler.c.

{
    cb_context_payload_t *payload = (cb_context_payload_t *) entry->payload;
    cb_context_arg_t *context_arg = (cb_context_arg_t *) arg;
 
    process_context_entry(payload, context_arg->features, context_arg->feature_count);
 
    return false;
}

Here is the call graph for this function:

Here is the caller graph for this function:

copy_unique_events()

int copy_unique_events ( uint64_t *  target, uint64_t *  source, int  source_len  )

static

Definition at line 1875 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    char name[PAPI_MAX_STR_LEN] = { 0 };
    int i, j;
    void *count_table, *p;
 
    htable_init(&count_table);
 
    for (i = 0, j = 0; i < source_len; ++i) {
        event_info_t info;
        papi_errno = evt_id_to_info(source[i], &info);
        if (papi_errno) {
            return papi_errno;
        }
        if (ntv_table_p->events[info.nameid].instances > 1) {
            sprintf(name, "%s[%i]", ntv_table_p->events[info.nameid].name, info.instance);
        } else {
            sprintf(name, "%s", ntv_table_p->events[info.nameid].name);
        }
        if (htable_find(count_table, name, &p) != HTABLE_SUCCESS) {
            htable_insert(count_table, name, NULL);
            target[j++] = source[i];
        }
    }
 
    htable_shutdown(count_table);
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

count_ntv_events_cb() [1/2]

hsa_status_t count_ntv_events_cb	(	const rocprofiler_info_data_t info	__attribute__(unused),
		void *	count
	)

init_event_table utility functions

Definition at line 873 of file roc_profiler.c.

{
    (*(int *) count) += 1;
    return HSA_STATUS_SUCCESS;
}

count_ntv_events_cb() [2/2]

static hsa_status_t count_ntv_events_cb ( const  rocprofiler_info_data_t, void *    )

static

Here is the caller graph for this function:

count_unique_events()

int count_unique_events ( uint64_t *  events_id, int  num_events, int *  num_unique  )

static

Definition at line 1841 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    char name[PAPI_MAX_STR_LEN] = { 0 };
    int i;
    int count = 0;
    void *count_table, *p;
 
    htable_init(&count_table);
 
    for (i = 0; i < num_events; ++i) {
        event_info_t info;
        papi_errno = evt_id_to_info(events_id[i], &info);
        if (papi_errno != PAPI_OK) {
            return papi_errno;
        }
        if (ntv_table_p->events[info.nameid].instances > 1) {
            sprintf(name, "%s[%i]", ntv_table_p->events[info.nameid].name, info.instance);
        } else {
            sprintf(name, "%s", ntv_table_p->events[info.nameid].name);
        }
        if (htable_find(count_table, name, &p) != HTABLE_SUCCESS) {
            htable_insert(count_table, name, NULL);
            ++count;
        }
    }
 
    *num_unique = count;
 
    htable_shutdown(count_table);
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

ctx_close()

int ctx_close ( rocp_ctx_t  rocp_ctx )

static

Definition at line 1324 of file roc_profiler.c.

{
    int papi_errno;
    int i, devs_count;
    rocc_dev_get_count(rocp_ctx->u.sampling.device_map, &devs_count);
 
    for (i = 0; i < devs_count; ++i) {
        if (rocp_close_p(rocp_ctx->u.sampling.contexts[i]) != HSA_STATUS_SUCCESS) {
            papi_errno = PAPI_EMISC;
        }
 
        if (hsa_queue_destroy_p(rocp_ctx->u.sampling.ctx_prop[i].queue) != HSA_STATUS_SUCCESS) {
            papi_errno = PAPI_EMISC;
        }
    }
 
    papi_errno = rocc_dev_release(rocp_ctx->u.sampling.device_map);
 
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

ctx_finalize()

int ctx_finalize ( rocp_ctx_t *  rocp_ctx )

static

Definition at line 1958 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_finalize(rocp_ctx);
    }
 
    return intercept_ctx_finalize(rocp_ctx);
}

Here is the call graph for this function:

Here is the caller graph for this function:

ctx_get_dev_feature_count()

int ctx_get_dev_feature_count ( rocp_ctx_t  rocp_ctx, int  i  )

static

Definition at line 1384 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_get_dev_feature_count(rocp_ctx, i);
    }
 
    return intercept_ctx_get_dev_feature_count(rocp_ctx, i);
}

Here is the call graph for this function:

Here is the caller graph for this function:

ctx_init()

int ctx_init ( uint64_t *  events_id, int  num_events, rocp_ctx_t *  rocp_ctx  )

static

Context init and finalize

Definition at line 1947 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_init(events_id, num_events, rocp_ctx);
    }
 
    return intercept_ctx_init(events_id, num_events, rocp_ctx);
}

Here is the call graph for this function:

Here is the caller graph for this function:

ctx_open()

int ctx_open ( rocp_ctx_t  rocp_ctx )

static

Definition at line 1268 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    int i, j;
    rocprofiler_feature_t *features = rocp_ctx->u.sampling.features;
    int dev_feature_offset = 0;
    int dev_count;
    rocprofiler_t **contexts = rocp_ctx->u.sampling.contexts;
    rocprofiler_properties_t *ctx_prop = rocp_ctx->u.sampling.ctx_prop;
 
    papi_errno = rocc_dev_get_count(rocp_ctx->u.sampling.device_map, &dev_count);
    if (papi_errno != PAPI_OK) {
        return papi_errno;
    }
 
    for (i = 0; i < dev_count; ++i) {
        int dev_id;
        papi_errno = rocc_dev_get_id(rocp_ctx->u.sampling.device_map, i, &dev_id);
        if (papi_errno != PAPI_OK) {
            goto fn_fail;
        }
 
        int dev_feature_count = ctx_get_dev_feature_count(rocp_ctx, dev_id);
        rocprofiler_feature_t *dev_features = features + dev_feature_offset;
 
        const uint32_t mode =
            ROCPROFILER_MODE_STANDALONE | ROCPROFILER_MODE_CREATEQUEUE | ROCPROFILER_MODE_SINGLEGROUP;
 
        ctx_prop[i].queue_depth = 128;
        hsa_status_t rocp_errno = rocp_open_p(device_table_p->devices[dev_id], dev_features,
                                              dev_feature_count, &contexts[i], mode,
                                              &ctx_prop[i]);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            papi_errno = PAPI_EMISC;
            goto fn_fail;
        }
 
        dev_feature_offset += dev_feature_count;
    }
 
    papi_errno = rocc_dev_acquire(rocp_ctx->u.sampling.device_map);
    if (papi_errno != PAPI_OK) {
       goto fn_fail;
    }
 
  fn_exit:
    return papi_errno;
  fn_fail:
    for (j = 0; j < i; ++j) {
        rocp_close_p(contexts[j]);
        hsa_queue_destroy_p(ctx_prop[j].queue);
    }
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

decrement_and_fetch_dispatch_counter()

int decrement_and_fetch_dispatch_counter ( unsigned long  tid )

static

Definition at line 2327 of file roc_profiler.c.

{
    int htable_errno = HTABLE_SUCCESS;
    char key[PAPI_MIN_STR_LEN] = { 0 };
 
    sprintf(key, "%lu", tid);
    int *counter_p;
    htable_errno = htable_find(htable, (const char *) key, (void **) &counter_p);
    if (htable_errno != HTABLE_SUCCESS) {
        return 0;
    }
 
    return --(*counter_p);
}

Here is the call graph for this function:

Here is the caller graph for this function:

dispatch_cb()

hsa_status_t dispatch_cb ( const rocprofiler_callback_data_t *  callback_data, void *  arg, rocprofiler_group_t *  group  )

static

Definition at line 2114 of file roc_profiler.c.

{
    hsa_agent_t agent = callback_data->agent;
    hsa_status_t status = HSA_STATUS_SUCCESS;
 
    int dev_id;
    rocc_dev_get_agent_id(agent, &dev_id);
 
    cb_dispatch_arg_t *dispatch_arg = (cb_dispatch_arg_t *) arg;
    rocprofiler_pool_t *pool = dispatch_arg->pools[dev_id];
    rocprofiler_pool_entry_t pool_entry;
    hsa_status_t rocp_errno = rocp_pool_fetch_p(pool, &pool_entry);
    if (rocp_errno != HSA_STATUS_SUCCESS) {
        status = rocp_errno;
        goto fn_exit;
    }
 
    rocprofiler_t *context = pool_entry.context;
    cb_context_payload_t *payload = (cb_context_payload_t *) pool_entry.payload;
 
    rocp_errno = rocp_get_group_p(context, 0, group);
    if (rocp_errno != HSA_STATUS_SUCCESS) {
        status = rocp_errno;
        goto fn_exit;
    }
 
    unsigned long tid;
    rocc_thread_get_id(&tid);
    payload->tid = tid;
    payload->agent = agent;
    payload->group = *group;
    payload->data = *callback_data;
 
    _papi_hwi_lock(_rocm_lock);
    payload->valid = true;
    _papi_hwi_unlock(_rocm_lock);
 
  fn_exit:
    return status;
}

Here is the call graph for this function:

Here is the caller graph for this function:

event_id_to_dev_id_cb()

static int event_id_to_dev_id_cb ( uint64_t  event_id, int *  device  )

static

sampling_ctx_open utility functions

Definition at line 1145 of file roc_profiler.c.

{
    event_info_t info;
    int papi_errno = evt_id_to_info(event_id, &info);
    *device = info.device;
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

evt_code_to_name()

int evt_code_to_name ( uint64_t  event_code, char *  name, int  len  )

static

Definition at line 739 of file roc_profiler.c.

{
    int papi_errno;
 
    event_info_t info;
    papi_errno = evt_id_to_info(event_code, &info);
    if (papi_errno != PAPI_OK) {
        return papi_errno;
    }
 
    switch (info.flags) {
        case (DEVICE_FLAG | INSTAN_FLAG):
            snprintf(name, len, "%s:device=%i:instance=%i", ntv_table_p->events[info.nameid].name, info.device, info.instance);
            break;
        case (DEVICE_FLAG):
            snprintf(name, len, "%s:device=%i", ntv_table_p->events[info.nameid].name, info.device);
            break;
        default:
            snprintf(name, len, "%s", ntv_table_p->events[info.nameid].name);
    }
 
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

evt_id_create()

int evt_id_create ( event_info_t *  info, uint64_t *  event_id  )

static

Definition at line 764 of file roc_profiler.c.

{
    *event_id  = (uint64_t)(info->device   << DEVICE_SHIFT);
    *event_id |= (uint64_t)(info->instance << INSTAN_SHIFT);
    *event_id |= (uint64_t)(info->flags    << QLMASK_SHIFT);
    *event_id |= (uint64_t)(info->nameid   << NAMEID_SHIFT);
    return PAPI_OK;
}

Here is the caller graph for this function:

evt_id_to_info()

int evt_id_to_info ( uint64_t  event_id, event_info_t *  info  )

static

Definition at line 774 of file roc_profiler.c.

{
    info->device   = (int)((event_id & DEVICE_MASK) >> DEVICE_SHIFT);
    info->instance = (int)((event_id & INSTAN_MASK) >> INSTAN_SHIFT);
    info->flags    = (int)((event_id & QLMASK_MASK) >> QLMASK_SHIFT);
    info->nameid   = (int)((event_id & NAMEID_MASK) >> NAMEID_SHIFT);
 
    if (info->device >= device_table_p->count) {
        return PAPI_ENOEVNT;
    }
 
    if (0 == (info->flags & DEVICE_FLAG) && info->device > 0) {
        return PAPI_ENOEVNT;
    }
 
    if (rocc_dev_check(ntv_table_p->events[info->nameid].device_map, info->device) == 0) {
        return PAPI_ENOEVNT;
    }
 
    if (info->nameid >= ntv_table_p->count) {
        return PAPI_ENOEVNT;
    }
 
    if (ntv_table_p->events[info->nameid].instances > 1 && 0 == (info->flags & INSTAN_FLAG) && info->instance > 0) {
        return PAPI_ENOEVNT;
    }
 
    if (info->instance >= ntv_table_p->events[info->nameid].instances) {
        return PAPI_ENOEVNT;
    }
 
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

evt_name_to_basename()

int evt_name_to_basename ( const char *  name, char *  base, int  len  )

static

Definition at line 851 of file roc_profiler.c.

{
    char *p = strstr(name, ":");
    if (p) {
        if (len < (int)(p - name)) {
            return PAPI_EBUF;
        }
        strncpy(base, name, (size_t)(p - name));
    } else {
        if (len < (int) strlen(name)) {
            return PAPI_EBUF;
        }
        strncpy(base, name, (size_t) len);
    }
    return PAPI_OK;
}

Here is the caller graph for this function:

evt_name_to_device()

int evt_name_to_device ( const char *  name, int *  device  )

static

Definition at line 809 of file roc_profiler.c.

{
    char *p = strstr(name, ":device=");
    if (!p) {
        return PAPI_ENOEVNT;
    }
    *device = (int) strtol(p + strlen(":device="), NULL, 10);
    return PAPI_OK;
}

Here is the caller graph for this function:

evt_name_to_instance()

int evt_name_to_instance ( const char *  name, int *  instance  )

static

Definition at line 820 of file roc_profiler.c.

{
    *instance = 0;
 
    char basename[PAPI_MAX_STR_LEN] = { 0 };
    int papi_errno = evt_name_to_basename(name, basename, PAPI_MAX_STR_LEN);
    if (papi_errno != PAPI_OK) {
        return papi_errno;
    }
 
    ntv_event_t *event;
    if (htable_find(htable, basename, (void **) &event) != HTABLE_SUCCESS) {
        return PAPI_ENOEVNT;
    }
 
    char *p = strstr(name, ":instance=");
    if (event->instances > 1) {
        if (!p) {
            return PAPI_ENOEVNT;
        }
        *instance = (int) strtol(p + strlen(":instance="), NULL, 10);
    } else {
        if (p) {
            return PAPI_ENOEVNT;
        }
    }
 
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

fetch_dispatch_counter()

int fetch_dispatch_counter ( unsigned long  tid )

static

Definition at line 2281 of file roc_profiler.c.

{
    int htable_errno = HTABLE_SUCCESS;
    char key[PAPI_MIN_STR_LEN] = { 0 };
 
    sprintf(key, "%lu", tid);
    int *counter_p;
    htable_errno = htable_find(htable, (const char *) key, (void **) &counter_p);
    if (htable_errno != HTABLE_SUCCESS) {
        return 0;
    }
 
    return (*counter_p);
}

Here is the call graph for this function:

Here is the caller graph for this function:

finalize_features()

int finalize_features ( rocprofiler_feature_t *  features, int  feature_count  )

static

Definition at line 1371 of file roc_profiler.c.

{
    int i;
    for (i = 0; i < feature_count; ++i) {
        papi_free((char *) features[i].name);
    }
    return PAPI_OK;
}

Here is the caller graph for this function:

free_context_node()

void free_context_node ( cb_context_node_t *  n )

static

Definition at line 2368 of file roc_profiler.c.

{
    papi_free(n->counters);
    papi_free(n);
}

Here is the caller graph for this function:

get_context_counters()

int get_context_counters ( int  dev_id, cb_context_node_t *  n, rocp_ctx_t  rocp_ctx  )

static

Definition at line 2343 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    uint64_t *events_id = rocp_ctx->u.intercept.events_id;
 
    /* Here we get events_id ordered according to user's viewpoint and we want
     * to map these to events_id ordered according to callbacks' viewpoint. We
     * compare events from the user and the callbacks using a brute force
     * approach as the number of events is typically small. */
    int i, j;
    for (i = 0; i < intercept_global_state.feature_count; ++i) {
        uint64_t event_id = intercept_global_state.events_id[i] | (dev_id << DEVICE_SHIFT);
 
        for (j = 0; j < rocp_ctx->u.intercept.feature_count; ++j) {
            if (event_id == events_id[j]) {
                rocp_ctx->u.intercept.counters[j] += n->counters[i];
                break;
            }
        }
    }
 
    return papi_errno;
}

Here is the caller graph for this function:

get_context_node()

int get_context_node ( int  dev_id, cb_context_node_t **  n  )

static

Definition at line 2297 of file roc_profiler.c.

{
    cb_context_node_t *curr = cb_ctx_list_heads[dev_id];
    cb_context_node_t *flag = NULL;
    cb_context_node_t *prev = curr;
    cb_context_node_t *flag_prev;
 
    while (curr) {
        unsigned long tid;
        rocc_thread_get_id(&tid);
        if (curr->tid == tid) {
            flag_prev = prev;
            flag = curr;
        }
        prev = curr;
        curr = curr->next;
    }
 
    if (flag != NULL) {
        flag_prev->next = flag->next;
        if (cb_ctx_list_heads[dev_id] == flag) {
            cb_ctx_list_heads[dev_id] = NULL;
        }
    }
 
    *n = flag;
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

get_ntv_events_cb()

hsa_status_t get_ntv_events_cb ( const rocprofiler_info_data_t  info, void *  ntv_arg  )

static

Definition at line 880 of file roc_profiler.c.

{
    struct ntv_arg *arg = (struct ntv_arg *) ntv_arg;
    int capacity = ntv_table.count;
    int *count = &arg->count;
    ntv_event_t *events = ntv_table.events;
 
    if (*count > capacity) {
        snprintf(error_string, PAPI_MAX_STR_LEN, "Number of events exceeds detected count.");
        return HSA_STATUS_ERROR;
    }
 
    ntv_event_t *event;
    if (htable_find(htable, info.metric.name, (void **) &event) != HTABLE_SUCCESS) {
        event = &events[(*count)++];
        event->name = papi_strdup(info.metric.name);
        event->descr = papi_strdup(info.metric.description);
        event->instances = info.metric.instances;
        htable_insert(htable, info.metric.name, event);
    }
 
    rocc_dev_set(&event->device_map, arg->dev_id);
 
    return HSA_STATUS_SUCCESS;
}

Here is the call graph for this function:

Here is the caller graph for this function:

increment_and_fetch_dispatch_counter()

int increment_and_fetch_dispatch_counter ( unsigned long  tid )

static

Definition at line 2265 of file roc_profiler.c.

{
    int htable_errno = HTABLE_SUCCESS;
    char key[PAPI_MIN_STR_LEN] = { 0 };
 
    sprintf(key, "%lu", tid);
    int *counter_p;
    htable_errno = htable_find(htable, (const char *) key, (void **) &counter_p);
    if (htable_errno != HTABLE_SUCCESS) {
        return 0;
    }
 
    return ++(*counter_p);
}

Here is the call graph for this function:

Here is the caller graph for this function:

init_callbacks()

int init_callbacks ( rocprofiler_feature_t *  features, int  feature_count  )

static

Definition at line 1993 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    cb_context_arg_t *context_arg = papi_calloc(1, sizeof(cb_context_arg_t));
    if (context_arg == NULL) {
        papi_errno = PAPI_ENOMEM;
        goto fn_fail;
    }
 
    context_arg->features = features;
    context_arg->feature_count = feature_count;
 
    rocprofiler_pool_properties_t properties;
    properties.num_entries = 128;
    properties.payload_bytes = sizeof(cb_context_payload_t);
    properties.handler = context_handler_cb;
    properties.handler_arg = context_arg;
 
    /* FIXME: the intercept code initializes callbacks for every device
     *        regardless what the user asked for. Moreover, every device
     *        is initialized with the same callback events (features).
     *        The intercept code should eventually be changed to allow
     *        user to initialize different callbacks on different devices
     *        and also to reinitialize already initialized callbacks on
     *        any given device. Rocm 5.3.0 still does not support this
     *        callback initialization mechanism.
     */
    int i;
    for (i = 0; i < device_table_p->count; ++i) {
        hsa_agent_t agent = device_table_p->devices[i];
 
        rocprofiler_pool_t *pool = NULL;
        if (rocp_pool_open_p(agent, features, feature_count, &pool, 0, &properties) != HSA_STATUS_SUCCESS) {
            papi_errno = PAPI_EMISC;
            goto fn_fail;
        }
 
        cb_dispatch_arg.pools[i] = pool;
    }
 
    rocprofiler_queue_callbacks_t dispatch_ptrs = { 0 };
    dispatch_ptrs.dispatch = dispatch_cb;
 
    if (rocp_set_queue_cbs_p(dispatch_ptrs, &cb_dispatch_arg) != HSA_STATUS_SUCCESS) {
        papi_errno = PAPI_EMISC;
        goto fn_fail;
    }
 
  fn_exit:
    return papi_errno;
  fn_fail:
    if (context_arg) {
        papi_free(context_arg);
    }
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

init_event_table()

int init_event_table ( void  )

static

Definition at line 682 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    int i;
 
    for (i = 0; i < device_table_p->count; ++i) {
        hsa_status_t rocp_errno = rocp_iterate_info_p(&device_table_p->devices[i],
                                                      ROCPROFILER_INFO_KIND_METRIC,
                                                      &count_ntv_events_cb,
                                                      &ntv_table.count);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            const char *error_string_p;
            hsa_status_string_p(rocp_errno, &error_string_p);
            snprintf(error_string, PAPI_MAX_STR_LEN, "%s", error_string_p);
            papi_errno = PAPI_EMISC;
            goto fn_fail;
        }
    }
 
    ntv_table.events = papi_calloc(ntv_table.count, sizeof(ntv_event_t));
    if (ntv_table.events == NULL) {
        papi_errno = PAPI_ENOMEM;
        goto fn_fail;
    }
 
    struct ntv_arg arg;
    arg.count = 0;
 
    for (i = 0; i < device_table_p->count; ++i) {
        arg.dev_id = i;
        hsa_status_t rocp_errno = rocp_iterate_info_p(&device_table_p->devices[i],
                                                      ROCPROFILER_INFO_KIND_METRIC,
                                                      &get_ntv_events_cb,
                                                      &arg);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            const char *error_string_p;
            hsa_status_string_p(rocp_errno, &error_string_p);
            snprintf(error_string, PAPI_MAX_STR_LEN, "%s", error_string_p);
            papi_errno = PAPI_EMISC;
            goto fn_fail;
        }
    }
 
    ntv_table.events = papi_realloc(ntv_table.events, arg.count * sizeof(ntv_event_t));
    if (ntv_table.events == NULL) {
        papi_errno = PAPI_ENOMEM;
    }
 
    ntv_table.count = arg.count;
 
  fn_exit:
    return papi_errno;
  fn_fail:
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

init_features()

int init_features ( uint64_t *  events_id, int  num_events, rocprofiler_feature_t *  features  )

static

rocp_ctx_{open,close,start,stop,read,reset} sampling mode utility functions

Definition at line 1346 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    int i;
    for (i = 0; i < num_events; ++i) {
        char name[PAPI_MAX_STR_LEN] = { 0 };
        event_info_t info;
        papi_errno = evt_id_to_info(events_id[i], &info);
        if (papi_errno != PAPI_OK) {
            break;
        }
        if (ntv_table_p->events[info.nameid].instances > 1) {
            sprintf(name, "%s[%i]", ntv_table_p->events[info.nameid].name, info.instance);
        } else {
            strcpy(name, ntv_table_p->events[info.nameid].name);
        }
        features[i].kind = (rocprofiler_feature_kind_t) ROCPROFILER_INFO_KIND_METRIC;
        features[i].name = papi_strdup(name);
    }
 
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

init_rocp_env()

int init_rocp_env ( void  )

static

Definition at line 580 of file roc_profiler.c.

{
    static int rocp_env_initialized;
 
    if (rocp_env_initialized) {
        return PAPI_OK;
    }
 
    const char *rocp_mode = getenv("ROCP_HSA_INTERCEPT");
    rocm_prof_mode = (rocp_mode != NULL) ?
        atoi(rocp_mode) : ROCM_PROFILE_SAMPLING_MODE;
 
    char pathname[PATH_MAX];
    char *rocm_root = getenv("PAPI_ROCM_ROOT");
    if (rocm_root == NULL) {
        snprintf(error_string, PAPI_MAX_STR_LEN, "Can't set HSA_TOOLS_LIB. PAPI_ROCM_ROOT not set.");
        return PAPI_EMISC;
    }
 
    int err;
    int override_hsa_tools_lib = 1;
    struct stat stat_info;
    char *hsa_tools_lib = getenv("HSA_TOOLS_LIB");
    if (hsa_tools_lib) {
        err = stat(hsa_tools_lib, &stat_info);
        if (err == 0 && S_ISREG(stat_info.st_mode)) {
            override_hsa_tools_lib = 0;
        }
    }
 
    if (override_hsa_tools_lib) {
        /* Account for change of librocprofiler64.so file location in rocm-5.2.0
         * directory structure */
 
        /* prefer .so.1 as .so might not be available in 5.7 anymore, in 5.6 it
         * was a linker script. */
        const char *candidates[] = {
            "lib/librocprofiler64.so.1",
            "lib/librocprofiler64.so",
            "rocprofiler/lib/libprofiler64.so.1",
            "rocprofiler/lib/libprofiler64.so",
            NULL
        };
        const char **candidate = candidates;
        while (*candidate) {
            sprintf(pathname, "%s/%s", rocm_root, *candidate);
 
            err = stat(pathname, &stat_info);
            if (err == 0) {
                break;
            }
            ++candidate;
        }
        if (!*candidate) {
            snprintf(error_string, PAPI_MAX_STR_LEN, "Rocprofiler librocprofiler64.so file not found.");
            return PAPI_EMISC;
        }
 
        setenv("HSA_TOOLS_LIB", pathname, 1);
    }
 
    int override_rocp_metrics = 1;
    char *rocp_metrics = getenv("ROCP_METRICS");
    if (rocp_metrics) {
        err = stat(rocp_metrics, &stat_info);
        if (err == 0 && S_ISREG(stat_info.st_mode)) {
            override_rocp_metrics = 0;
        }
    }
 
    if (override_rocp_metrics) {
        /* Account for change of metrics file location in rocm-5.2.0
         * directory structure */
        sprintf(pathname, "%s/lib/rocprofiler/metrics.xml", rocm_root);
 
        err = stat(pathname, &stat_info);
        if (err < 0) {
            sprintf(pathname, "%s/rocprofiler/lib/metrics.xml", rocm_root);
 
            err = stat(pathname, &stat_info);
            if (err < 0) {
                snprintf(error_string, PAPI_MAX_STR_LEN, "Rocprofiler metrics.xml file not found.");
                return PAPI_EMISC;
            }
        }
 
        setenv("ROCP_METRICS", pathname, 1);
    }
 
    rocp_env_initialized = 1;
    return PAPI_OK;
}

Here is the caller graph for this function:

intercept_ctx_close()

int intercept_ctx_close ( rocp_ctx_t  rocp_ctx )

static

Definition at line 1512 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    _papi_hwi_lock(_rocm_lock);
 
    if (intercept_global_state.active_thread_count == 0) {
        goto fn_exit;
    }
 
    unsigned long tid;
    rocc_thread_get_id(&tid);
    papi_errno = unregister_dispatch_counter(tid);
    if (papi_errno != PAPI_OK) {
        goto fn_exit;
    }
 
    ctx_finalize(&rocp_ctx);
 
  fn_exit:
    _papi_hwi_unlock(_rocm_lock);
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

intercept_ctx_finalize()

int intercept_ctx_finalize ( rocp_ctx_t *  rocp_ctx )

static

Definition at line 1926 of file roc_profiler.c.

{
    if (*rocp_ctx == NULL) {
        return PAPI_OK;
    }
 
    if ((*rocp_ctx)->u.intercept.counters) {
        papi_free((*rocp_ctx)->u.intercept.counters);
    }
 
    papi_free(*rocp_ctx);
    *rocp_ctx = NULL;
 
    return PAPI_OK;
}

Here is the caller graph for this function:

intercept_ctx_get_dev_feature_count()

int intercept_ctx_get_dev_feature_count ( rocp_ctx_t  rocp_ctx, int  i  )

static

Definition at line 1416 of file roc_profiler.c.

{
    int start, stop, j = 0;
    int num_events = rocp_ctx->u.intercept.feature_count;
    uint64_t *events_id = rocp_ctx->u.intercept.events_id;
 
    while (j < num_events && (events_id[j] & DEVICE_MASK) >> DEVICE_SHIFT != (uint64_t) i) {
        ++j;
    }
 
    start = j;
 
    while (j < num_events && (events_id[j] & DEVICE_MASK) >> DEVICE_SHIFT == (uint64_t) i) {
        ++j;
    }
 
    stop = j;
 
    return stop - start;
}

Here is the caller graph for this function:

intercept_ctx_init()

int intercept_ctx_init ( uint64_t *  events_id, int  num_events, rocp_ctx_t *  rocp_ctx  )

static

Definition at line 1755 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    long long *counters = NULL;
    *rocp_ctx = NULL;
 
    rocc_bitmap_t bitmap;
    papi_errno = rocc_dev_get_map(event_id_to_dev_id_cb, events_id, num_events, &bitmap);
    if (papi_errno != PAPI_OK) {
        return papi_errno;
    }
 
    if (intercept_global_state.events_id == NULL) {
        int num_unique_events = 0;
        count_unique_events(events_id, num_events, &num_unique_events);
        intercept_global_state.events_id = papi_calloc(num_unique_events, sizeof(uint64_t));
        if (intercept_global_state.events_id == NULL) {
            papi_errno = PAPI_ENOMEM;
            goto fn_fail_undo;
        }
        copy_unique_events(intercept_global_state.events_id, events_id, num_events);
 
        intercept_global_state.features = papi_calloc(num_unique_events, sizeof(*intercept_global_state.features));
        if (intercept_global_state.features == NULL) {
            papi_errno = PAPI_ENOMEM;
            goto fn_fail_undo;
        }
 
        papi_errno = init_features(intercept_global_state.events_id, num_unique_events, intercept_global_state.features);
        if (papi_errno != PAPI_OK) {
            goto fn_fail_undo;
        }
 
        intercept_global_state.feature_count = num_unique_events;
 
        papi_errno = save_callback_features(intercept_global_state.features, intercept_global_state.feature_count);
        if (papi_errno != PAPI_OK) {
            goto fn_fail_undo;
        }
 
        papi_errno = init_callbacks(intercept_global_state.features, intercept_global_state.feature_count);
        if (papi_errno != PAPI_OK) {
            goto fn_fail;
        }
    }
 
    counters = papi_calloc(num_events, sizeof(*counters));
    if (counters == NULL) {
        papi_errno = PAPI_ENOMEM;
        goto fn_fail;
    }
 
    *rocp_ctx = papi_calloc(1, sizeof(**rocp_ctx));
    if (*rocp_ctx == NULL) {
        return PAPI_ENOMEM;
    }
 
    (*rocp_ctx)->u.intercept.events_id = events_id;
    (*rocp_ctx)->u.intercept.counters = counters;
    (*rocp_ctx)->u.intercept.dispatch_count = 0;
    (*rocp_ctx)->u.intercept.device_map = bitmap;
    (*rocp_ctx)->u.intercept.feature_count = num_events;
 
  fn_exit:
    return papi_errno;
  fn_fail:
    if (counters) {
        papi_free(counters);
    }
    if (*rocp_ctx) {
        papi_free(*rocp_ctx);
    }
    *rocp_ctx = NULL;
    goto fn_exit;
  fn_fail_undo:
    cleanup_callback_features(intercept_global_state.features, intercept_global_state.feature_count);
    if (intercept_global_state.events_id) {
        papi_free(intercept_global_state.events_id);
    }
    if (intercept_global_state.features) {
        papi_free(intercept_global_state.features);
    }
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

intercept_ctx_open()

int intercept_ctx_open ( uint64_t *  events_id, int  num_events, rocp_ctx_t *  rocp_ctx  )

static

Definition at line 1472 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    if (num_events <= 0) {
        return PAPI_ENOEVNT;
    }
 
    _papi_hwi_lock(_rocm_lock);
 
    papi_errno = verify_events(events_id, num_events);
    if (papi_errno != PAPI_OK) {
        SUBDBG("[ROCP intercept mode] Can only monitor one set of events "
               "per application run.");
        goto fn_fail;
    }
 
    papi_errno = ctx_init(events_id, num_events, rocp_ctx);
    if (papi_errno != PAPI_OK) {
        goto fn_fail;
    }
 
    unsigned long tid;
    rocc_thread_get_id(&tid);
    papi_errno = register_dispatch_counter(tid, &(*rocp_ctx)->u.intercept.dispatch_count);
    if (papi_errno != PAPI_OK) {
        goto fn_fail;
    }
 
    (*rocp_ctx)->u.intercept.state |= ROCM_EVENTS_OPENED;
 
  fn_exit:
    _papi_hwi_unlock(_rocm_lock);
    return papi_errno;
  fn_fail:
    ctx_finalize(rocp_ctx);
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

intercept_ctx_read()

int intercept_ctx_read ( rocp_ctx_t  rocp_ctx, long long **  counts  )

static

Definition at line 1607 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    _papi_hwi_lock(_rocm_lock);
 
    unsigned long tid;
    rocc_thread_get_id(&tid);
    int dispatch_count = fetch_dispatch_counter(tid);
    if (dispatch_count == 0) {
        *counts = rocp_ctx->u.intercept.counters;
        goto fn_exit;
    }
 
    cb_context_node_t *n = NULL;
 
    int devs_count;
    papi_errno = rocc_dev_get_count(rocp_ctx->u.intercept.device_map, &devs_count);
    if (papi_errno != PAPI_OK) {
        goto fn_exit;
    }
 
    int i;
    for (i = 0; i < devs_count; ++i) {
        int dev_id;
        papi_errno = rocc_dev_get_id(rocp_ctx->u.intercept.device_map, i, &dev_id);
        if (papi_errno != PAPI_OK) {
            goto fn_exit;
        }
 
        while (dispatch_count > 0) {
            get_context_node(dev_id, &n);
            if (n == NULL) {
                break;
            }
 
            get_context_counters(dev_id, n, rocp_ctx);
            dispatch_count = decrement_and_fetch_dispatch_counter(tid);
            free_context_node(n);
        }
    }
 
    if (dispatch_count > 0) {
        SUBDBG("[ROCP intercept mode] User monitoring GPU i but running on j.");
    }
 
    *counts = rocp_ctx->u.intercept.counters;
 
  fn_exit:
    _papi_hwi_unlock(_rocm_lock);
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

intercept_ctx_reset()

int intercept_ctx_reset ( rocp_ctx_t  rocp_ctx )

static

Definition at line 1661 of file roc_profiler.c.

{
    int i;
 
    for (i = 0; i < rocp_ctx->u.intercept.feature_count; ++i) {
        rocp_ctx->u.intercept.counters[i] = 0;
    }
 
    return PAPI_OK;
}

Here is the caller graph for this function:

intercept_ctx_start()

int intercept_ctx_start ( rocp_ctx_t  rocp_ctx )

static

Definition at line 1537 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    _papi_hwi_lock(_rocm_lock);
 
    if (!(rocp_ctx->u.sampling.state & ROCM_EVENTS_OPENED)) {
        SUBDBG("[ROCP intercept mode] Cannot start eventset, not opened.");
        papi_errno = PAPI_EINVAL;
        goto fn_fail;
    }
 
    if (rocp_ctx->u.intercept.state & ROCM_EVENTS_RUNNING) {
        SUBDBG("[ROCP intercept mode] Cannot start eventset, already running.");
        papi_errno = PAPI_EINVAL;
        goto fn_fail;
    }
 
    if (intercept_global_state.kernel_count++ == 0) {
        if (rocp_start_queue_cbs_p() != HSA_STATUS_SUCCESS) {
            papi_errno = PAPI_EMISC;
            goto fn_fail;
        }
    }
 
    rocp_ctx->u.intercept.state |= ROCM_EVENTS_RUNNING;
 
  fn_exit:
    _papi_hwi_unlock(_rocm_lock);
    return papi_errno;
  fn_fail:
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

intercept_ctx_stop()

int intercept_ctx_stop ( rocp_ctx_t  rocp_ctx )

static

Definition at line 1572 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    _papi_hwi_lock(_rocm_lock);
 
    if (!(rocp_ctx->u.sampling.state & ROCM_EVENTS_OPENED)) {
        SUBDBG("[ROCP intercept mode] Cannot stop eventset, not opened.");
        papi_errno = PAPI_EINVAL;
        goto fn_fail;
    }
 
    if (!(rocp_ctx->u.intercept.state & ROCM_EVENTS_RUNNING)) {
        SUBDBG("[ROCP intercept mode] Cannot stop eventset, not running.");
        papi_errno = PAPI_EINVAL;
        goto fn_fail;
    }
 
    if (--intercept_global_state.kernel_count == 0) {
        if (rocp_stop_queue_cbs_p() != HSA_STATUS_SUCCESS) {
            papi_errno = PAPI_EMISC;
            goto fn_fail;
        }
    }
 
    rocp_ctx->u.intercept.state &= ~ROCM_EVENTS_RUNNING;
 
  fn_exit:
    _papi_hwi_unlock(_rocm_lock);
    return papi_errno;
  fn_fail:
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

intercept_shutdown()

int intercept_shutdown ( void  )

static

Definition at line 1688 of file roc_profiler.c.

{
    /* calling rocprofiler_pool_close() here would cause
     * a double free runtime error. */
 
    shutdown_event_table();
    htable_shutdown(htable);
    htable_shutdown(htable_intercept);
 
    (*hsa_shut_down_p)();
 
    unload_rocp_sym();
 
    if (intercept_global_state.features) {
        finalize_features(intercept_global_state.features, intercept_global_state.feature_count);
        papi_free(intercept_global_state.features);
    }
 
    if (intercept_global_state.events_id) {
        papi_free(intercept_global_state.events_id);
    }
 
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

load_rocp_sym()

int load_rocp_sym ( void  )

static

rocp_{init,shutdown} and rocp_ctx_{open,close,start,stop,read,reset} functions

rocp_init utility functions

Definition at line 470 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    char *pathname = getenv("HSA_TOOLS_LIB");
    if (pathname == NULL) {
        snprintf(error_string, PAPI_MAX_STR_LEN, "Can't load librocprofiler64.so, neither PAPI_ROCM_ROOT nor HSA_TOOLS_LIB are set.");
        goto fn_fail;
    }
 
    rocp_dlp = dlopen(pathname, RTLD_NOW | RTLD_GLOBAL);
    if (rocp_dlp == NULL) {
        sprintf(error_string, "%s", dlerror());
        goto fn_fail;
    }
 
    rocp_get_info_p        = dlsym(rocp_dlp, "rocprofiler_get_info");
    rocp_iterate_info_p    = dlsym(rocp_dlp, "rocprofiler_iterate_info");
    rocp_error_string_p    = dlsym(rocp_dlp, "rocprofiler_error_string");
    rocp_open_p            = dlsym(rocp_dlp, "rocprofiler_open");
    rocp_close_p           = dlsym(rocp_dlp, "rocprofiler_close");
    rocp_group_count_p     = dlsym(rocp_dlp, "rocprofiler_group_count");
    rocp_start_p           = dlsym(rocp_dlp, "rocprofiler_start");
    rocp_read_p            = dlsym(rocp_dlp, "rocprofiler_read");
    rocp_stop_p            = dlsym(rocp_dlp, "rocprofiler_stop");
    rocp_get_group_p       = dlsym(rocp_dlp, "rocprofiler_get_group");
    rocp_get_data_p        = dlsym(rocp_dlp, "rocprofiler_get_data");
    rocp_group_get_data_p  = dlsym(rocp_dlp, "rocprofiler_group_get_data");
    rocp_get_metrics_p     = dlsym(rocp_dlp, "rocprofiler_get_metrics");
    rocp_reset_p           = dlsym(rocp_dlp, "rocprofiler_reset");
    rocp_pool_open_p       = dlsym(rocp_dlp, "rocprofiler_pool_open");
    rocp_pool_close_p      = dlsym(rocp_dlp, "rocprofiler_pool_close");
    rocp_pool_fetch_p      = dlsym(rocp_dlp, "rocprofiler_pool_fetch");
    rocp_pool_flush_p      = dlsym(rocp_dlp, "rocprofiler_pool_flush");
    rocp_set_queue_cbs_p   = dlsym(rocp_dlp, "rocprofiler_set_queue_callbacks");
    rocp_start_queue_cbs_p = dlsym(rocp_dlp, "rocprofiler_start_queue_callbacks");
    rocp_stop_queue_cbs_p  = dlsym(rocp_dlp, "rocprofiler_stop_queue_callbacks");
    rocp_remove_queue_cbs_p= dlsym(rocp_dlp, "rocprofiler_remove_queue_callbacks");
 
    int rocp_not_initialized = (!rocp_get_info_p       ||
                                !rocp_iterate_info_p   ||
                                !rocp_error_string_p   ||
                                !rocp_open_p           ||
                                !rocp_close_p          ||
                                !rocp_group_count_p    ||
                                !rocp_start_p          ||
                                !rocp_read_p           ||
                                !rocp_stop_p           ||
                                !rocp_get_group_p      ||
                                !rocp_get_data_p       ||
                                !rocp_group_get_data_p ||
                                !rocp_get_metrics_p    ||
                                !rocp_reset_p          ||
                                !rocp_pool_open_p      ||
                                !rocp_pool_close_p     ||
                                !rocp_pool_fetch_p     ||
                                !rocp_pool_flush_p     ||
                                !rocp_set_queue_cbs_p  ||
                                !rocp_start_queue_cbs_p||
                                !rocp_stop_queue_cbs_p ||
                                !rocp_remove_queue_cbs_p);
 
    papi_errno = (rocp_not_initialized) ? PAPI_EMISC : PAPI_OK;
    if (papi_errno != PAPI_OK) {
        snprintf(error_string, PAPI_MAX_STR_LEN, "Error while loading rocprofiler symbols.");
    }
 
  fn_exit:
    return papi_errno;
  fn_fail:
    papi_errno = PAPI_ENOSUPP;
    goto fn_exit;
}

Here is the caller graph for this function:

process_context_entry()

void process_context_entry ( cb_context_payload_t *  payload, rocprofiler_feature_t *  features, int  feature_count  )

static

intercept mode counter read infrastructure

Definition at line 2167 of file roc_profiler.c.

{
  fn_check_again:
    _papi_hwi_lock(_rocm_lock);
    if (payload->valid == false) {
        _papi_hwi_unlock(_rocm_lock);
        goto fn_check_again;
    }
 
    if (feature_count < 1) {
        goto fn_exit;
    }
 
    if (rocp_group_get_data_p(&payload->group) != HSA_STATUS_SUCCESS) {
        goto fn_exit;
    }
 
    if (rocp_get_metrics_p(payload->group.context)) {
        goto fn_exit;
    }
 
    if (increment_and_fetch_dispatch_counter(payload->tid) < 0) {
        /* thread not registered, ignore counters */
        goto fn_exit;
    }
 
    cb_context_node_t *n = alloc_context_node(feature_count);
    if (n == NULL) {
        decrement_and_fetch_dispatch_counter(payload->tid);
        goto fn_exit;
    }
 
    int dev_id;
    rocc_dev_get_agent_id(payload->agent, &dev_id);
 
    n->tid = payload->tid;
    put_context_counters(features, feature_count, n);
    put_context_node(dev_id, n);
 
  fn_exit:
    _papi_hwi_unlock(_rocm_lock);
}

Here is the call graph for this function:

Here is the caller graph for this function:

put_context_counters()

void put_context_counters ( rocprofiler_feature_t *  features, int  feature_count, cb_context_node_t *  n  )

static

Definition at line 2228 of file roc_profiler.c.

{
    int i;
    for (i = 0; i < feature_count; ++i) {
        const rocprofiler_feature_t *f = &features[i];
        switch(f->data.kind) {
            case ROCPROFILER_DATA_KIND_INT32:
                n->counters[i] = (long long) f->data.result_int32;
                break;
            case ROCPROFILER_DATA_KIND_INT64:
                n->counters[i] = f->data.result_int64;
                break;
            case ROCPROFILER_DATA_KIND_FLOAT:
                n->counters[i] = (long long) f->data.result_float;
                break;
            case ROCPROFILER_DATA_KIND_DOUBLE:
                n->counters[i] = (long long) f->data.result_double;
                break;
            default:
                SUBDBG("Unsupported data kind from rocprofiler");
        }
    }
}

Here is the call graph for this function:

Here is the caller graph for this function:

put_context_node()

void put_context_node ( int  dev_id, cb_context_node_t *  n  )

static

Definition at line 2253 of file roc_profiler.c.

{
    n->next = NULL;
 
    if (cb_ctx_list_heads[dev_id] != NULL) {
        n->next = cb_ctx_list_heads[dev_id];
    }
 
    cb_ctx_list_heads[dev_id] = n;
}

Here is the caller graph for this function:

register_dispatch_counter()

int register_dispatch_counter ( unsigned long  tid, int *  counter  )

static

Definition at line 2052 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    int htable_errno = HTABLE_SUCCESS;
    char key[PAPI_MIN_STR_LEN] = { 0 };
 
    /* FIXME: probably better using a different hash table for this */
    sprintf(key, "%lu", tid);
    int *counter_p;
    htable_errno = htable_find(htable, key, (void **) &counter_p);
    if (htable_errno == HTABLE_SUCCESS) {
        papi_errno = PAPI_EMISC;
        goto fn_exit;
    }
 
    htable_insert(htable, (const char *) key, counter);
    ++intercept_global_state.active_thread_count;
 
  fn_exit:
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_ctx_close()

int rocp_ctx_close

(

rocp_ctx_t

rocp_ctx

)

Definition at line 401 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_close(rocp_ctx);
    }
 
    return intercept_ctx_close(rocp_ctx);
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_ctx_open()

int rocp_ctx_open	(	uint64_t *	events_id,
		int	num_events,
		rocp_ctx_t *	rocp_ctx
	)

Definition at line 390 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_open(events_id, num_events, rocp_ctx);
    }
 
    return intercept_ctx_open(events_id, num_events, rocp_ctx);
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_ctx_read()

int rocp_ctx_read	(	rocp_ctx_t	rocp_ctx,
		long long **	counts
	)

Definition at line 434 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_read(rocp_ctx, counts);
    }
 
    return intercept_ctx_read(rocp_ctx, counts);
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_ctx_reset()

int rocp_ctx_reset

(

rocp_ctx_t

rocp_ctx

)

Definition at line 445 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_reset(rocp_ctx);
    }
 
    return intercept_ctx_reset(rocp_ctx);
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_ctx_start()

int rocp_ctx_start

(

rocp_ctx_t

rocp_ctx

)

Definition at line 412 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_start(rocp_ctx);
    }
 
    return intercept_ctx_start(rocp_ctx);
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_ctx_stop()

int rocp_ctx_stop

(

rocp_ctx_t

rocp_ctx

)

Definition at line 423 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_ctx_stop(rocp_ctx);
    }
 
    return intercept_ctx_stop(rocp_ctx);
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_evt_code_to_descr()

int rocp_evt_code_to_descr	(	uint64_t	event_code,
		char *	descr,
		int	len
	)

Definition at line 260 of file roc_profiler.c.

{
    int papi_errno;
 
    event_info_t info;
    papi_errno = evt_id_to_info(event_code, &info);
    if (papi_errno != PAPI_OK) {
        return papi_errno;
    }
 
    snprintf(descr, (size_t) len, "%s", ntv_table_p->events[info.nameid].descr);
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_evt_code_to_info()

int rocp_evt_code_to_info	(	uint64_t	event_code,
		PAPI_event_info_t *	info
	)

Definition at line 331 of file roc_profiler.c.

{
    int papi_errno;
 
    event_info_t inf;
    papi_errno = evt_id_to_info(event_code, &inf);
    if (papi_errno != PAPI_OK) {
        return papi_errno;
    }
 
    switch (inf.flags) {
        case 0:
            sprintf(info->symbol, "%s", ntv_table_p->events[inf.nameid].name);
            sprintf(info->long_descr, "%s", ntv_table_p->events[inf.nameid].descr);
            break;
        case (DEVICE_FLAG | INSTAN_FLAG):
        {
            int i;
            char devices[PAPI_MAX_STR_LEN] = { 0 };
            for (i = 0; i < device_table_p->count; ++i) {
                if (rocc_dev_check(ntv_table_p->events[inf.nameid].device_map, i)) {
                    sprintf(devices + strlen(devices), "%i,", i);
                }
            }
            *(devices + strlen(devices) - 1) = 0;
            sprintf(info->symbol, "%s:device=%i:instance=%i", ntv_table_p->events[inf.nameid].name, inf.device, inf.instance);
            sprintf(info->long_descr, "%s, masks:Mandatory device qualifier [%s]:Mandatory instance qualifier in range [0-%i]",
                    ntv_table_p->events[inf.nameid].descr, devices, ntv_table_p->events[inf.nameid].instances - 1);
            break;
        }
        case DEVICE_FLAG:
        {
            int i;
            char devices[PAPI_MAX_STR_LEN] = { 0 };
            for (i = 0; i < device_table_p->count; ++i) {
                if (rocc_dev_check(ntv_table_p->events[inf.nameid].device_map, i)) {
                    sprintf(devices + strlen(devices), "%i,", i);
                }
            }
            *(devices + strlen(devices) - 1) = 0;
            sprintf(info->symbol, "%s:device=%i", ntv_table_p->events[inf.nameid].name, inf.device);
            sprintf(info->long_descr, "%s, masks:Mandatory device qualifier [%s]",
                    ntv_table_p->events[inf.nameid].descr, devices);
            break;
        }
        case INSTAN_FLAG:
            sprintf(info->symbol, "%s:instance=%i", ntv_table_p->events[inf.nameid].name, inf.instance);
            sprintf(info->long_descr, "%s, masks:Mandatory instance qualifier in range [0-%i]",
                    ntv_table_p->events[inf.nameid].descr, ntv_table_p->events[inf.nameid].instances - 1);
            break;
        default:
            papi_errno = PAPI_EINVAL;
    }
 
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_evt_code_to_name()

int rocp_evt_code_to_name	(	uint64_t	event_code,
		char *	name,
		int	len
	)

Definition at line 324 of file roc_profiler.c.

{
    return evt_code_to_name(event_code, name, len);
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_evt_enum()

int rocp_evt_enum	(	uint64_t *	event_code,
		int	modifier
	)

Definition at line 193 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    event_info_t info;
    SUBDBG("ENTER: event_code: %lu, modifier: %d\n", *event_code, modifier);
 
 
    switch(modifier) {
        case PAPI_ENUM_FIRST:
            if (ntv_table_p->count == 0) {
                papi_errno = PAPI_ENOEVNT;
                break;
            }
            info.device = 0;
            info.instance = 0;
            info.flags = 0;
            info.nameid = 0;
            papi_errno = evt_id_create(&info, event_code);
            break;
        case PAPI_ENUM_EVENTS:
            papi_errno = evt_id_to_info(*event_code, &info);
            if (papi_errno != PAPI_OK) {
                break;
            }
            if (ntv_table_p->count > info.nameid + 1) {
                info.device = 0;
                info.instance = 0;
                info.flags = 0;
                info.nameid++;
                papi_errno = evt_id_create(&info, event_code);
                break;
            }
            papi_errno = PAPI_END;
            break;
        case PAPI_NTV_ENUM_UMASKS:
            papi_errno = evt_id_to_info(*event_code, &info);
            if (papi_errno != PAPI_OK) {
                break;
            }
            if (info.flags == 0) {
                info.device = 0;
                info.instance = 0;
                info.flags = DEVICE_FLAG;
                papi_errno = evt_id_create(&info, event_code);
                break;
            }
            if (info.flags & DEVICE_FLAG) {
                if (ntv_table_p->events[info.nameid].instances > 1) {
                    info.device = 0;
                    info.instance = 0;
                    info.flags = INSTAN_FLAG;
                    papi_errno = evt_id_create(&info, event_code);
                    break;
                }
            }
            papi_errno = PAPI_END;
            break;
        default:
            papi_errno = PAPI_EINVAL;
    }
 
    SUBDBG("EXIT: %s\n", PAPI_strerror(papi_errno));
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_evt_name_to_code()

int rocp_evt_name_to_code	(	const char *	name,
		uint64_t *	event_code
	)

Definition at line 276 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    int htable_errno;
    SUBDBG("ENTER: name: %s, event_code: %p\n", name, event_code);
 
    int device;
    papi_errno = evt_name_to_device(name, &device);
    if (papi_errno != PAPI_OK) {
        goto fn_exit;
    }
 
    int instance;
    papi_errno = evt_name_to_instance(name, &instance);
    if (papi_errno != PAPI_OK) {
        goto fn_exit;
    }
 
    char base[PAPI_MAX_STR_LEN] = { 0 };
    papi_errno = evt_name_to_basename(name, base, PAPI_MAX_STR_LEN);
    if (papi_errno != PAPI_OK) {
        goto fn_exit;
    }
 
    ntv_event_t *event;
    htable_errno = htable_find(htable, base, (void **) &event);
    if (htable_errno != HTABLE_SUCCESS) {
        papi_errno = (htable_errno == HTABLE_ENOVAL) ? PAPI_ENOEVNT : PAPI_ECMP;
        goto fn_exit;
    }
 
    int flags = (event->instances > 1) ? (DEVICE_FLAG | INSTAN_FLAG) : DEVICE_FLAG;
    int nameid = (int) (event - ntv_table_p->events);
    event_info_t info = { device, instance, flags, nameid };
    papi_errno = evt_id_create(&info, event_code);
    if (papi_errno != PAPI_OK) {
        goto fn_exit;
    }
 
    papi_errno = evt_id_to_info(*event_code, &info);
 
  fn_exit:
    SUBDBG("EXIT: %s\n", PAPI_strerror(papi_errno));
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_init()

int rocp_init

(

void

)

Definition at line 159 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    SUBDBG("ENTER\n");
 
    papi_errno = load_rocp_sym();
    if (papi_errno != PAPI_OK) {
        goto fn_fail;
    }
 
    htable_init(&htable);
 
    if (rocm_prof_mode != ROCM_PROFILE_SAMPLING_MODE) {
        htable_init(&htable_intercept);
    }
 
    papi_errno = init_event_table();
    if (papi_errno != PAPI_OK) {
        (*hsa_shut_down_p)();
        goto fn_fail;
    }
 
    ntv_table_p = &ntv_table;
 
  fn_exit:
    SUBDBG("EXIT: %s\n", PAPI_strerror(papi_errno));
    return papi_errno;
  fn_fail:
    unload_rocp_sym();
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_init_environment()

int rocp_init_environment

(

void

)

Definition at line 152 of file roc_profiler.c.

{
    return init_rocp_env();
}

Here is the call graph for this function:

Here is the caller graph for this function:

rocp_shutdown()

int rocp_shutdown

(

void

)

Definition at line 456 of file roc_profiler.c.

{
    if (rocm_prof_mode == ROCM_PROFILE_SAMPLING_MODE) {
        return sampling_shutdown();
    }
 
    return intercept_shutdown();
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_ctx_close()

int sampling_ctx_close ( rocp_ctx_t  rocp_ctx )

static

Definition at line 952 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    _papi_hwi_lock(_rocm_lock);
 
    papi_errno = ctx_close(rocp_ctx);
    if (papi_errno != PAPI_OK) {
        goto fn_fail;
    }
 
    ctx_finalize(&rocp_ctx);
 
  fn_exit:
    _papi_hwi_unlock(_rocm_lock);
    return papi_errno;
  fn_fail:
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_ctx_finalize()

int sampling_ctx_finalize ( rocp_ctx_t *  rocp_ctx )

static

Definition at line 1238 of file roc_profiler.c.

{
    if (*rocp_ctx == NULL) {
        return PAPI_OK;
    }
 
    if ((*rocp_ctx)->u.sampling.features) {
        finalize_features((*rocp_ctx)->u.sampling.features, (*rocp_ctx)->u.sampling.feature_count);
        papi_free((*rocp_ctx)->u.sampling.features);
    }
 
    if ((*rocp_ctx)->u.sampling.contexts) {
        papi_free((*rocp_ctx)->u.sampling.contexts);
    }
 
    if ((*rocp_ctx)->u.sampling.ctx_prop) {
        papi_free((*rocp_ctx)->u.sampling.ctx_prop);
    }
 
    if ((*rocp_ctx)->u.sampling.counters) {
        papi_free((*rocp_ctx)->u.sampling.counters);
    }
 
    papi_free(*rocp_ctx);
    *rocp_ctx = NULL;
 
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_ctx_get_dev_feature_count()

int sampling_ctx_get_dev_feature_count ( rocp_ctx_t  rocp_ctx, int  i  )

static

Definition at line 1394 of file roc_profiler.c.

{
    int start, stop, j = 0;
    int num_events = rocp_ctx->u.sampling.feature_count;
    uint64_t *events_id = rocp_ctx->u.sampling.events_id;
 
    while (j < num_events && (events_id[j] & DEVICE_MASK) >> DEVICE_SHIFT != (uint64_t) i) {
        ++j;
    }
 
    start = j;
 
    while (j < num_events && (events_id[j] & DEVICE_MASK) >> DEVICE_SHIFT == (uint64_t) i) {
        ++j;
    }
 
    stop = j;
 
    return stop - start;
}

Here is the caller graph for this function:

sampling_ctx_init()

int sampling_ctx_init ( uint64_t *  events_id, int  num_events, rocp_ctx_t *  rocp_ctx  )

static

Definition at line 1154 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    int num_devs;
    rocprofiler_feature_t *features = NULL;
    rocprofiler_t **contexts = NULL;
    rocprofiler_properties_t *ctx_prop = NULL;
    long long *counters = NULL;
    *rocp_ctx = NULL;
 
    rocc_bitmap_t bitmap;
    papi_errno = rocc_dev_get_map(event_id_to_dev_id_cb, events_id, num_events, &bitmap);
    if (papi_errno != PAPI_OK) {
        return papi_errno;
    }
 
    papi_errno = rocc_dev_get_count(bitmap, &num_devs);
    if (papi_errno != PAPI_OK) {
        goto fn_fail;
    }
 
    contexts = papi_calloc(num_devs, sizeof(*contexts));
    if (contexts == NULL) {
        papi_errno = PAPI_ENOMEM;
        goto fn_fail;
    }
 
    ctx_prop = papi_calloc(num_devs, sizeof(*ctx_prop));
    if (ctx_prop == NULL) {
        papi_errno = PAPI_ENOMEM;
        goto fn_fail;
    }
 
    features = papi_calloc(num_events, sizeof(*features));
    if (features == NULL) {
        papi_errno = PAPI_ENOMEM;
        goto fn_fail;
    }
 
    counters = papi_malloc(num_events * sizeof(*counters));
    if (counters == NULL) {
        papi_errno = PAPI_ENOMEM;
        goto fn_fail;
    }
 
    papi_errno = init_features(events_id, num_events, features);
    if (papi_errno != PAPI_OK) {
        goto fn_fail;
    }
 
    *rocp_ctx = papi_calloc(1, sizeof(**rocp_ctx));
    if (*rocp_ctx == NULL) {
        papi_errno = PAPI_ENOMEM;
        goto fn_fail;
    }
 
    (*rocp_ctx)->u.sampling.events_id = events_id;
    (*rocp_ctx)->u.sampling.features = features;
    (*rocp_ctx)->u.sampling.feature_count = num_events;
    (*rocp_ctx)->u.sampling.contexts = contexts;
    (*rocp_ctx)->u.sampling.counters = counters;
    (*rocp_ctx)->u.sampling.device_map = bitmap;
    (*rocp_ctx)->u.sampling.ctx_prop = ctx_prop;
 
  fn_exit:
    return papi_errno;
  fn_fail:
    if (contexts) {
        papi_free(contexts);
    }
    if (features) {
        papi_free(features);
    }
    if (counters) {
        papi_free(counters);
    }
    if (*rocp_ctx) {
        papi_free(*rocp_ctx);
    }
    *rocp_ctx = NULL;
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_ctx_open()

int sampling_ctx_open ( uint64_t *  events_id, int  num_events, rocp_ctx_t *  rocp_ctx  )

static

Definition at line 921 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
 
    if (num_events <= 0) {
        return PAPI_ENOEVNT;
    }
 
    _papi_hwi_lock(_rocm_lock);
 
    papi_errno = ctx_init(events_id, num_events, rocp_ctx);
    if (papi_errno != PAPI_OK) {
        goto fn_fail;
    }
 
    papi_errno = ctx_open(*rocp_ctx);
    if (papi_errno != PAPI_OK) {
        goto fn_fail;
    }
 
    (*rocp_ctx)->u.sampling.state |= ROCM_EVENTS_OPENED;
 
  fn_exit:
    _papi_hwi_unlock(_rocm_lock);
    return papi_errno;
  fn_fail:
    ctx_finalize(rocp_ctx);
    goto fn_exit;
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_ctx_read()

int sampling_ctx_read ( rocp_ctx_t  rocp_ctx, long long **  counts  )

static

Definition at line 1029 of file roc_profiler.c.

{
    int i, j, k = 0;
    int dev_feature_offset = 0;
    int dev_id, dev_count;
    rocprofiler_feature_t *features = rocp_ctx->u.sampling.features;
 
    rocc_dev_get_count(rocp_ctx->u.sampling.device_map, &dev_count);
 
    for (i = 0; i < dev_count; ++i) {
        hsa_status_t rocp_errno = rocp_read_p(rocp_ctx->u.sampling.contexts[i], 0);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            return PAPI_EMISC;
        }
 
        rocp_errno = rocp_get_data_p(rocp_ctx->u.sampling.contexts[i], 0);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            return PAPI_EMISC;
        }
 
        rocp_errno = rocp_get_metrics_p(rocp_ctx->u.sampling.contexts[i]);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            return PAPI_EMISC;
        }
 
        int papi_errno = rocc_dev_get_id(rocp_ctx->u.sampling.device_map, i, &dev_id);
        if (papi_errno != PAPI_OK) {
            return papi_errno;
        }
 
        int dev_feature_count = ctx_get_dev_feature_count(rocp_ctx, dev_id);
        rocprofiler_feature_t *dev_features = features + dev_feature_offset;
        long long *counters = rocp_ctx->u.sampling.counters;
 
        for (j = 0; j < dev_feature_count; ++j) {
            switch(dev_features[j].data.kind) {
                case ROCPROFILER_DATA_KIND_INT32:
                    counters[k++] = (long long) dev_features[j].data.result_int32;
                    break;
                case ROCPROFILER_DATA_KIND_INT64:
                    counters[k++] = dev_features[j].data.result_int64;
                    break;
                case ROCPROFILER_DATA_KIND_FLOAT:
                    counters[k++] = (long long) dev_features[j].data.result_float;
                    break;
                case ROCPROFILER_DATA_KIND_DOUBLE:
                    counters[k++] = (long long) dev_features[j].data.result_double;
                    break;
                default:
                    return PAPI_EINVAL;
            }
        }
        dev_feature_offset += dev_feature_count;
    }
    *counts = rocp_ctx->u.sampling.counters;
 
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_ctx_reset()

int sampling_ctx_reset ( rocp_ctx_t  rocp_ctx )

static

Definition at line 1089 of file roc_profiler.c.

{
    int i, devs_count;
    rocc_dev_get_count(rocp_ctx->u.sampling.device_map, &devs_count);
 
    for (i = 0; i < devs_count; ++i) {
        hsa_status_t rocp_errno = rocp_reset_p(rocp_ctx->u.sampling.contexts[i], 0);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            return PAPI_EMISC;
        }
    }
    for (i = 0; i < rocp_ctx->u.sampling.feature_count; ++i) {
        rocp_ctx->u.sampling.counters[i] = 0;
    }
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_ctx_start()

int sampling_ctx_start ( rocp_ctx_t  rocp_ctx )

static

Definition at line 973 of file roc_profiler.c.

{
    if (!(rocp_ctx->u.sampling.state & ROCM_EVENTS_OPENED)) {
        SUBDBG("[ROCP sampling mode] Cannot start eventset, not opened.");
        return PAPI_EINVAL;
    }
 
    if (rocp_ctx->u.sampling.state & ROCM_EVENTS_RUNNING) {
        SUBDBG("[ROCP sampling mode] Cannot start eventset, already running.");
        return PAPI_EINVAL;
    }
 
    int devs_count;
    rocc_dev_get_count(rocp_ctx->u.sampling.device_map, &devs_count);
 
    int i;
    for (i = 0; i < devs_count; ++i) {
        hsa_status_t rocp_errno = rocp_start_p(rocp_ctx->u.sampling.contexts[i], 0);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            return PAPI_EMISC;
        }
    }
 
    rocp_ctx->u.sampling.state |= ROCM_EVENTS_RUNNING;
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_ctx_stop()

int sampling_ctx_stop ( rocp_ctx_t  rocp_ctx )

static

Definition at line 1001 of file roc_profiler.c.

{
    if (!(rocp_ctx->u.sampling.state & ROCM_EVENTS_OPENED)) {
        SUBDBG("[ROCP sampling mode] Cannot stop eventset, not opened.");
        return PAPI_EINVAL;
    }
 
    if (!(rocp_ctx->u.sampling.state & ROCM_EVENTS_RUNNING)) {
        SUBDBG("[ROCP sampling mode] Cannot stop eventset, not running.");
        return PAPI_EINVAL;
    }
 
    int devs_count;
    rocc_dev_get_count(rocp_ctx->u.sampling.device_map, &devs_count);
 
    int i;
    for (i = 0; i < devs_count; ++i) {
        hsa_status_t rocp_errno = rocp_stop_p(rocp_ctx->u.sampling.contexts[i], 0);
        if (rocp_errno != HSA_STATUS_SUCCESS) {
            return PAPI_EMISC;
        }
    }
 
    rocp_ctx->u.sampling.state &= ~ROCM_EVENTS_RUNNING;
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

sampling_shutdown()

int sampling_shutdown ( void  )

static

Definition at line 1113 of file roc_profiler.c.

{
    shutdown_event_table();
    htable_shutdown(htable);
 
    unload_rocp_sym();
 
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

save_callback_features()

int save_callback_features ( rocprofiler_feature_t *  features, int  feature_count  )

static

Definition at line 1906 of file roc_profiler.c.

{
    int i;
    for (i = 0; i < feature_count; ++i) {
        htable_insert(htable_intercept, features[i].name, NULL);
    }
    return PAPI_OK;
}

Here is the call graph for this function:

Here is the caller graph for this function:

shutdown_event_table()

int shutdown_event_table ( void  )

static

rocp_shutdown sampling mode utility functions

Definition at line 1124 of file roc_profiler.c.

{
    int i;
 
    for (i = 0; i < ntv_table_p->count; ++i) {
        papi_free(ntv_table_p->events[i].name);
        papi_free(ntv_table_p->events[i].descr);
    }
 
    ntv_table_p->count = 0;
 
    papi_free(ntv_table_p->events);
 
    return PAPI_OK;
}

Here is the caller graph for this function:

unload_rocp_sym()

int unload_rocp_sym ( void  )

static

Definition at line 545 of file roc_profiler.c.

{
    if (rocp_dlp == NULL) {
        return PAPI_OK;
    }
 
    rocp_get_info_p        = NULL;
    rocp_iterate_info_p    = NULL;
    rocp_error_string_p    = NULL;
    rocp_open_p            = NULL;
    rocp_close_p           = NULL;
    rocp_group_count_p     = NULL;
    rocp_start_p           = NULL;
    rocp_read_p            = NULL;
    rocp_stop_p            = NULL;
    rocp_get_group_p       = NULL;
    rocp_get_data_p        = NULL;
    rocp_group_get_data_p  = NULL;
    rocp_get_metrics_p     = NULL;
    rocp_reset_p           = NULL;
    rocp_pool_open_p       = NULL;
    rocp_pool_close_p      = NULL;
    rocp_pool_fetch_p      = NULL;
    rocp_pool_flush_p      = NULL;
    rocp_set_queue_cbs_p   = NULL;
    rocp_start_queue_cbs_p = NULL;
    rocp_stop_queue_cbs_p  = NULL;
    rocp_remove_queue_cbs_p= NULL;
 
    dlclose(rocp_dlp);
 
    return PAPI_OK;
}

Here is the caller graph for this function:

unregister_dispatch_counter()

int unregister_dispatch_counter ( unsigned long  tid )

static

Definition at line 2075 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    int htable_errno = HTABLE_SUCCESS;
    char key[PAPI_MIN_STR_LEN] = { 0 };
 
    sprintf(key, "%lu", tid);
    int *counter_p;
    htable_errno = htable_find(htable, (const char *) key, (void **) &counter_p);
    if (htable_errno != HTABLE_SUCCESS) {
        papi_errno = PAPI_ECMP;
        goto fn_exit;
    }
 
    htable_delete(htable, (const char *) key);
    --intercept_global_state.active_thread_count;
 
  fn_exit:
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

verify_events()

int verify_events ( uint64_t *  events_id, int  num_events  )

static

intercept_ctx_{open,close} utility functions

Definition at line 1718 of file roc_profiler.c.

{
    int papi_errno = PAPI_OK;
    int i;
    char name[PAPI_MAX_STR_LEN] = { 0 };
 
    if (intercept_global_state.events_id == NULL) {
        return papi_errno;
    }
 
    for (i = 0; i < num_events; ++i) {
        event_info_t info;
        papi_errno = evt_id_to_info(events_id[i], &info);
        if (papi_errno != PAPI_OK) {
            break;
        }
        if (ntv_table_p->events[info.nameid].instances > 1) {
            sprintf(name, "%s[%i]", ntv_table_p->events[info.nameid].name, info.instance);
        } else {
            sprintf(name, "%s", ntv_table_p->events[info.nameid].name);
        }
        void *p;
        if (htable_find(htable_intercept, name, &p) != HTABLE_SUCCESS) {
            papi_errno = PAPI_ECNFLCT;
            break;
        }
    }
 
    return papi_errno;
}

Here is the call graph for this function:

Here is the caller graph for this function:

Variable Documentation

_rocm_lock

unsigned int _rocm_lock

Definition at line 85 of file roc_profiler.c.

active_thread_count

int active_thread_count

Definition at line 1451 of file roc_profiler.c.

cb_ctx_list_heads

cb_context_node_t* cb_ctx_list_heads[PAPI_ROCM_MAX_DEV_COUNT]

static

The context handler prepares a node for every processes entry. The node is associated with the thread that generated the monitoring request and contains the value of the counters read by rocprofiler. Each node is then added to the corresponding device queue and is eventually read by intercept_ctx_read

Definition at line 2111 of file roc_profiler.c.

cb_dispatch_arg

cb_dispatch_arg_t cb_dispatch_arg

static

Definition at line 1685 of file roc_profiler.c.

events_id

uint64_t* events_id

Definition at line 1448 of file roc_profiler.c.

feature_count

int feature_count

Definition at line 1450 of file roc_profiler.c.

features

rocprofiler_feature_t* features

Definition at line 1449 of file roc_profiler.c.

htable

void* htable

static

Definition at line 147 of file roc_profiler.c.

htable_intercept

void* htable_intercept

static

Definition at line 148 of file roc_profiler.c.

struct { ... } intercept_global_state

kernel_count

int kernel_count

Definition at line 1452 of file roc_profiler.c.

ntv_table

ntv_event_table_t ntv_table

static

Definition at line 145 of file roc_profiler.c.

ntv_table_p

ntv_event_table_t* ntv_table_p

static

Definition at line 146 of file roc_profiler.c.

rocm_prof_mode

unsigned int rocm_prof_mode

Definition at line 84 of file roc_profiler.c.

rocp_close_p

hsa_status_t(* rocp_close_p) (rocprofiler_t *) ( rocprofiler_t *  )

static

Definition at line 94 of file roc_profiler.c.

rocp_dlp

void* rocp_dlp = NULL

static

Definition at line 144 of file roc_profiler.c.

rocp_error_string_p

hsa_status_t(* rocp_error_string_p) (const char **) ( const char **  )

static

Definition at line 90 of file roc_profiler.c.

rocp_get_data_p

hsa_status_t(* rocp_get_data_p) (rocprofiler_t *, uint32_t) ( rocprofiler_t *  , uint32_t    )

static

Definition at line 100 of file roc_profiler.c.

rocp_get_group_p

hsa_status_t(* rocp_get_group_p) (rocprofiler_t *, uint32_t, rocprofiler_group_t *) ( rocprofiler_t *  , uint32_t  , rocprofiler_group_t *    )

static

Definition at line 99 of file roc_profiler.c.

rocp_get_info_p

hsa_status_t(* rocp_get_info_p) (const hsa_agent_t *, rocprofiler_info_kind_t, void *) ( const hsa_agent_t *  , rocprofiler_info_kind_t  , void *    )

static

Definition at line 88 of file roc_profiler.c.

rocp_get_metrics_p

hsa_status_t(* rocp_get_metrics_p) (const rocprofiler_t *) ( const rocprofiler_t *  )

static

Definition at line 102 of file roc_profiler.c.

rocp_group_count_p

hsa_status_t(* rocp_group_count_p) (const rocprofiler_t *, uint32_t *) ( const rocprofiler_t *  , uint32_t *    )

static

Definition at line 95 of file roc_profiler.c.

rocp_group_get_data_p

hsa_status_t(* rocp_group_get_data_p) (rocprofiler_group_t *) ( rocprofiler_group_t *  )

static

Definition at line 101 of file roc_profiler.c.

rocp_iterate_info_p

hsa_status_t(* rocp_iterate_info_p) (const hsa_agent_t *, rocprofiler_info_kind_t, hsa_status_t(*)(const rocprofiler_info_data_t, void *), void *) ( const hsa_agent_t *  , rocprofiler_info_kind_t  , hsa_status_t(*)(const rocprofiler_info_data_t, void *)  , void *    )

static

Definition at line 89 of file roc_profiler.c.

rocp_open_p

hsa_status_t(* rocp_open_p) (hsa_agent_t, rocprofiler_feature_t *, uint32_t, rocprofiler_t **, uint32_t, rocprofiler_properties_t *) ( hsa_agent_t  , rocprofiler_feature_t *  , uint32_t  , rocprofiler_t **  , uint32_t  , rocprofiler_properties_t *    )

static

Definition at line 93 of file roc_profiler.c.

rocp_pool_close_p

hsa_status_t(* rocp_pool_close_p) (rocprofiler_pool_t *) ( rocprofiler_pool_t *  )

static

Definition at line 107 of file roc_profiler.c.

rocp_pool_fetch_p

hsa_status_t(* rocp_pool_fetch_p) (rocprofiler_pool_t *, rocprofiler_pool_entry_t *) ( rocprofiler_pool_t *  , rocprofiler_pool_entry_t *    )

static

Definition at line 108 of file roc_profiler.c.

rocp_pool_flush_p

hsa_status_t(* rocp_pool_flush_p) (rocprofiler_pool_t *) ( rocprofiler_pool_t *  )

static

Definition at line 109 of file roc_profiler.c.

rocp_pool_open_p

hsa_status_t(* rocp_pool_open_p) (hsa_agent_t, rocprofiler_feature_t *, uint32_t, rocprofiler_pool_t **, uint32_t, rocprofiler_pool_properties_t *) ( hsa_agent_t  , rocprofiler_feature_t *  , uint32_t  , rocprofiler_pool_t **  , uint32_t  , rocprofiler_pool_properties_t *    )

static

Definition at line 106 of file roc_profiler.c.

rocp_read_p

hsa_status_t(* rocp_read_p) (rocprofiler_t *, uint32_t) ( rocprofiler_t *  , uint32_t    )

static

Definition at line 97 of file roc_profiler.c.

rocp_remove_queue_cbs_p

hsa_status_t(* rocp_remove_queue_cbs_p) (void) ( void  )

static

Definition at line 113 of file roc_profiler.c.

rocp_reset_p

hsa_status_t(* rocp_reset_p) (rocprofiler_t *, uint32_t) ( rocprofiler_t *  , uint32_t    )

static

Definition at line 103 of file roc_profiler.c.

rocp_set_queue_cbs_p

hsa_status_t(* rocp_set_queue_cbs_p) (rocprofiler_queue_callbacks_t, void *) ( rocprofiler_queue_callbacks_t  , void *    )

static

Definition at line 110 of file roc_profiler.c.

rocp_start_p

hsa_status_t(* rocp_start_p) (rocprofiler_t *, uint32_t) ( rocprofiler_t *  , uint32_t    )

static

Definition at line 96 of file roc_profiler.c.

rocp_start_queue_cbs_p

hsa_status_t(* rocp_start_queue_cbs_p) (void) ( void  )

static

Definition at line 111 of file roc_profiler.c.

rocp_stop_p

hsa_status_t(* rocp_stop_p) (rocprofiler_t *, uint32_t) ( rocprofiler_t *  , uint32_t    )

static

Definition at line 98 of file roc_profiler.c.

rocp_stop_queue_cbs_p

hsa_status_t(* rocp_stop_queue_cbs_p) (void) ( void  )

static

Definition at line 112 of file roc_profiler.c.