diff --git a/bin/rpl_run.sh b/bin/rpl_run.sh index 2ab65e8cd6..6f9af7ce65 100755 --- a/bin/rpl_run.sh +++ b/bin/rpl_run.sh @@ -433,8 +433,8 @@ while [ 1 ] ; do errck "Option '$ARG_IN', rate value" export ROCP_FLUSH_RATE="$period_rate" elif [ "$1" = "--obj-tracking" ] ; then - if [ "$2" = "on" ] ; then - export ROCP_OBJ_TRACKING=1 + if [ "$2" = "off" ] ; then + export ROCP_OBJ_TRACKING=0 fi elif [ "$1" = "--verbose" ] ; then ARG_VAL=0 diff --git a/src/core/rocprofiler.cpp b/src/core/rocprofiler.cpp index 5299e9ba33..618edf2397 100644 --- a/src/core/rocprofiler.cpp +++ b/src/core/rocprofiler.cpp @@ -203,6 +203,7 @@ uint32_t LoadTool() { settings.trace_local = TraceProfile::IsLocal() ? 1: 0; settings.timeout = util::HsaRsrcFactory::GetTimeoutNs(); settings.timestamp_on = InterceptQueue::IsTrackerOn() ? 1 : 0; + settings.code_obj_tracking = 1; if (handler) handler(); else if (handler_prop) handler_prop(&settings); diff --git a/src/util/hsa_rsrc_factory.cpp b/src/util/hsa_rsrc_factory.cpp index e42ca2c0a9..78833284fd 100644 --- a/src/util/hsa_rsrc_factory.cpp +++ b/src/util/hsa_rsrc_factory.cpp @@ -155,6 +155,9 @@ HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize // System timeout timeout_ = (timeout_ns_ == HsaTimer::TIMESTAMP_MAX) ? timeout_ns_ : timer_->ns_to_sysclock(timeout_ns_); + + // To dump code objects + to_dump_code_obj_ = getenv("ROCP_DUMP_CODEOBJ"); } // Destructor of the class @@ -195,6 +198,7 @@ void HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table) { hsa_api_.hsa_executable_create_alt = table->core_->hsa_executable_create_alt_fn; hsa_api_.hsa_executable_load_agent_code_object = table->core_->hsa_executable_load_agent_code_object_fn; hsa_api_.hsa_executable_freeze = table->core_->hsa_executable_freeze_fn; + hsa_api_.hsa_executable_destroy = table->core_->hsa_executable_destroy_fn; hsa_api_.hsa_executable_get_symbol = table->core_->hsa_executable_get_symbol_fn; hsa_api_.hsa_executable_symbol_get_info = table->core_->hsa_executable_symbol_get_info_fn; hsa_api_.hsa_executable_iterate_symbols = table->core_->hsa_executable_iterate_symbols_fn; @@ -235,6 +239,7 @@ void HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table) { hsa_api_.hsa_executable_create_alt = hsa_executable_create_alt; hsa_api_.hsa_executable_load_agent_code_object = hsa_executable_load_agent_code_object; hsa_api_.hsa_executable_freeze = hsa_executable_freeze; + hsa_api_.hsa_executable_destroy = hsa_executable_destroy; hsa_api_.hsa_executable_get_symbol = hsa_executable_get_symbol; hsa_api_.hsa_executable_symbol_get_info = hsa_executable_symbol_get_info; hsa_api_.hsa_executable_iterate_symbols = hsa_executable_iterate_symbols; @@ -710,6 +715,7 @@ void HsaRsrcFactory::EnableExecutableTracking(HsaApiTable* table) { std::lock_guard lck(mutex_); executable_tracking_on_ = true; table->core_->hsa_executable_freeze_fn = hsa_executable_freeze_interceptor; + table->core_->hsa_executable_destroy_fn = hsa_executable_destroy_interceptor; } hsa_status_t HsaRsrcFactory::executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data) { @@ -727,10 +733,14 @@ hsa_status_t HsaRsrcFactory::executable_symbols_cb(hsa_executable_t exec, hsa_ex status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name); CHECK_STATUS("Error in getting kernel name", status); name[len] = 0; - auto ret = symbols_map_->insert({addr, name}); - if (ret.second == false) { - delete[] ret.first->second; - ret.first->second = name; + if (data == NULL) { + auto ret = symbols_map_->insert({addr, name}); + if (ret.second == false) { + delete[] ret.first->second; + ret.first->second = name; + } + } else { + symbols_map_->erase(addr); } } return HSA_STATUS_SUCCESS; @@ -741,7 +751,16 @@ hsa_status_t HsaRsrcFactory::hsa_executable_freeze_interceptor(hsa_executable_t if (symbols_map_ == NULL) symbols_map_ = new symbols_map_t; hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(executable, executable_symbols_cb, NULL); CHECK_STATUS("Error in iterating executable symbols", status); - return hsa_api_.hsa_executable_freeze(executable, options);; + return hsa_api_.hsa_executable_freeze(executable, options); +} + +hsa_status_t HsaRsrcFactory::hsa_executable_destroy_interceptor(hsa_executable_t executable) { + std::lock_guard lck(mutex_); + if (symbols_map_ != NULL) { + hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(executable, executable_symbols_cb, (void*)1); + CHECK_STATUS("Error in iterating executable symbols", status); + } + return hsa_api_.hsa_executable_destroy(executable); } std::atomic HsaRsrcFactory::instance_{}; @@ -750,6 +769,7 @@ HsaRsrcFactory::timestamp_t HsaRsrcFactory::timeout_ns_ = HsaTimer::TIMESTAMP_MA hsa_pfn_t HsaRsrcFactory::hsa_api_{}; bool HsaRsrcFactory::executable_tracking_on_ = false; HsaRsrcFactory::symbols_map_t* HsaRsrcFactory::symbols_map_ = NULL; +void* HsaRsrcFactory::to_dump_code_obj_ = NULL; } // namespace util } // namespace rocprofiler diff --git a/src/util/hsa_rsrc_factory.h b/src/util/hsa_rsrc_factory.h index 38f9ceb83b..a8e392aa9e 100644 --- a/src/util/hsa_rsrc_factory.h +++ b/src/util/hsa_rsrc_factory.h @@ -97,6 +97,7 @@ struct hsa_pfn_t { decltype(hsa_executable_create_alt)* hsa_executable_create_alt; decltype(hsa_executable_load_agent_code_object)* hsa_executable_load_agent_code_object; decltype(hsa_executable_freeze)* hsa_executable_freeze; + decltype(hsa_executable_destroy)* hsa_executable_destroy; decltype(hsa_executable_get_symbol)* hsa_executable_get_symbol; decltype(hsa_executable_symbol_get_info)* hsa_executable_symbol_get_info; decltype(hsa_executable_iterate_symbols)* hsa_executable_iterate_symbols; @@ -498,7 +499,9 @@ class HsaRsrcFactory { typedef std::map symbols_map_t; static symbols_map_t* symbols_map_; static bool executable_tracking_on_; + static void* to_dump_code_obj_; static hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable, const char *options); + static hsa_status_t hsa_executable_destroy_interceptor(hsa_executable_t executable); static hsa_status_t executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data); // HSA runtime API table