Rocprofiler does not allow to change metrics when using intercept mode #71
Description
Currently, rocprofiler does not allow to change metrics at runtime for intercepted kernels, so the following example won't work:
rocprofiler_feature_t features[4];
features[0].kind = ROCPROFILER_FEATURE_KIND_METRIC;
features[0].name = "SQ_WAVES";
unsigned feature_count = 1;
init_intercept(features, feature_count);
start_intercept();
hipLaunchKernelGGL(vectoradd_float,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT);
hipDeviceSynchronize();
stop_intercept();
shutdown_intercept();
features[1].kind = ROCPROFILER_FEATURE_KIND_METRIC;
features[1].name = "SQ_INSTS_VALU";
feature_count += 1;
init_intercept(features, feature_count);
start_intercept();
hipLaunchKernelGGL(vectoradd_float,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT);
hipDeviceSynchronize();
stop_intercept();
shutdown_intercept();
Above, init_intercept() initializes the queue callbacks for intercept mode and calls rocprofiler_set_queue_callbacks(). start_intercept() and stop_intercept() call rocprofiler_start_queue_callbacks() and rocprofiler_stop_queue_callbacks(), respectively, and shutdown_intercept() calls rocprofiler_remove_queue_callbacks().
Rocprofiler does not allow users to call rocprofiler_set_queue_callbacks() if this has been already called. Thus, the second call to init_intercept() in the example code above causes the following error message:
> error(4096) "SetCallbacks(), reassigning queue callbacks - not supported”
The ability to change metrics at runtime (while using intercept mode) is a feature highly desirable for tools like PAPI. With the current implementation of rocprofiler PAPI users would have to define metrics once and have them applied to all the kernels being intercepted.