From caa15d7186805a356224b23e3adce8174d8d4136 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Fri, 16 Mar 2018 10:34:44 -0500 Subject: [PATCH] dispatch vs profiling packets signals race fixed Change-Id: I7f250d01f3d45f3f8ec2539b2187ab99503850aa --- inc/rocprofiler.h | 2 +- src/core/context.h | 6 +++--- src/core/intercept_queue.cpp | 1 + src/core/intercept_queue.h | 15 +++++++++----- src/core/rocprofiler.cpp | 6 ++++++ src/core/tracker.h | 11 ++++++++++ test/ctrl/tool.cpp | 39 +++++++++++++++++++++++++----------- test/gfx_metrics.xml | 3 +-- test/metrics.xml | 2 +- 9 files changed, 61 insertions(+), 24 deletions(-) diff --git a/inc/rocprofiler.h b/inc/rocprofiler.h index 0a2be56714..85061b9dc1 100644 --- a/inc/rocprofiler.h +++ b/inc/rocprofiler.h @@ -169,7 +169,7 @@ typedef enum { } rocprofiler_mode_t; // Profiling handler, calling on profiling completion -typedef void (*rocprofiler_handler_t)(rocprofiler_group_t group, void* arg); +typedef bool (*rocprofiler_handler_t)(rocprofiler_group_t group, void* arg); // Profiling preperties typedef struct { diff --git a/src/core/context.h b/src/core/context.h index 901cb811c1..e07fa2927b 100644 --- a/src/core/context.h +++ b/src/core/context.h @@ -6,6 +6,7 @@ #include #include #include +#include #include #include "core/metrics.h" @@ -125,8 +126,7 @@ class Group { } void ResetRefs() { refs_ = n_profiles_; } uint32_t DecrRefs() { - --refs_; - return refs_; + return (refs_ > 0) ? --refs_ : 0; } private: @@ -395,7 +395,7 @@ class Context { uint32_t r = group->DecrRefs(); group->GetContext()->mutex_.unlock(); if (r == 0) { - group->GetContext()->handler_(group->GetGroup(), group->GetContext()->handler_arg_); + return group->GetContext()->handler_(group->GetGroup(), group->GetContext()->handler_arg_); } return false; } diff --git a/src/core/intercept_queue.cpp b/src/core/intercept_queue.cpp index 6af4b84b0f..9c29a714e8 100644 --- a/src/core/intercept_queue.cpp +++ b/src/core/intercept_queue.cpp @@ -14,4 +14,5 @@ InterceptQueue::obj_map_t* InterceptQueue::obj_map_ = NULL; const char* InterceptQueue::kernel_none_ = ""; uint64_t InterceptQueue::timeout_ = UINT64_MAX; Tracker* InterceptQueue::tracker_ = NULL; +bool InterceptQueue::tracker_on_ = true; } // namespace rocprofiler diff --git a/src/core/intercept_queue.h b/src/core/intercept_queue.h index 8e6543c5ea..53ec083644 100644 --- a/src/core/intercept_queue.h +++ b/src/core/intercept_queue.h @@ -41,7 +41,7 @@ class InterceptQueue { group_segment_size, queue, &status); if (status != HSA_STATUS_SUCCESS) abort(); - if (!tracker_) tracker_ = new Tracker(timeout_); + if (tracker_on_ && (tracker_ == NULL)) tracker_ = new Tracker(timeout_); status = hsa_amd_profiling_set_profiler_enabled(*queue, true); if (status != HSA_STATUS_SUCCESS) abort(); @@ -89,15 +89,19 @@ class InterceptQueue { const hsa_kernel_dispatch_packet_t* dispatch_packet = reinterpret_cast(packet); const char* kernel_name = GetKernelName(dispatch_packet); - const auto* entry = tracker_->Add(obj->agent_info_->dev_id, dispatch_packet->completion_signal); - const_cast(dispatch_packet)->completion_signal = entry->signal; + const rocprofiler_dispatch_record_t* record = NULL; + if (tracker_ != NULL) { + const auto* entry = tracker_->Add(obj->agent_info_->dev_id, dispatch_packet->completion_signal); + const_cast(dispatch_packet)->completion_signal = entry->signal; + record = entry->record; + } rocprofiler_callback_data_t data = {obj->agent_info_->dev_id, obj->agent_info_->dev_index, obj->queue_, user_que_idx, dispatch_packet->kernel_object, kernel_name, - entry->record}; + record}; hsa_status_t status = dispatch_callback_(&data, callback_data_, &group); free(const_cast(kernel_name)); if ((status == HSA_STATUS_SUCCESS) && (group.context != NULL)) { @@ -137,6 +141,7 @@ class InterceptQueue { } static void SetTimeout(uint64_t timeout) { timeout_ = timeout; } + static void TrackerOn(bool on) { tracker_on_ = on; } private: InterceptQueue(const hsa_agent_t& agent, hsa_queue_t* const queue, ProxyQueue* proxy) : @@ -188,7 +193,7 @@ class InterceptQueue { static const char* kernel_none_; static uint64_t timeout_; static Tracker* tracker_; - static const bool tracker_on_ = true; + static bool tracker_on_; hsa_queue_t* const queue_; ProxyQueue* const proxy_; diff --git a/src/core/rocprofiler.cpp b/src/core/rocprofiler.cpp index ca3cb5d980..69075eeb30 100644 --- a/src/core/rocprofiler.cpp +++ b/src/core/rocprofiler.cpp @@ -148,6 +148,11 @@ CONSTRUCTOR_API void constructor() { Context::SetTimeout(timeout_val); InterceptQueue::SetTimeout(timeout_val); } + const char* tracker_on_str = getenv("ROCP_TRACKER_ON"); + if (tracker_on_str != NULL) { + if (strncmp(tracker_on_str, "true", 4) == 0) InterceptQueue::TrackerOn(true); + if (strncmp(tracker_on_str, "false", 4) == 0) InterceptQueue::TrackerOn(false); + } } DESTRUCTOR_API void destructor() { @@ -173,6 +178,7 @@ const MetricsDict* GetMetrics(const hsa_agent_t& agent) { return metrics; } +rocprofiler::Tracker::mutex_t rocprofiler::Tracker::mutex_; util::Logger::mutex_t util::Logger::mutex_; util::Logger* util::Logger::instance_ = NULL; uint64_t Context::timeout_ = UINT64_MAX; diff --git a/src/core/tracker.h b/src/core/tracker.h index 2dcf36c50c..98ae3fe19e 100644 --- a/src/core/tracker.h +++ b/src/core/tracker.h @@ -6,6 +6,7 @@ #include #include +#include #include "inc/rocprofiler.h" #include "util/exception.h" @@ -15,6 +16,7 @@ namespace rocprofiler { class Tracker { public: + typedef std::mutex mutex_t; typedef rocprofiler_dispatch_record_t record_t; struct entry_t; typedef std::list sig_list_t; @@ -29,6 +31,7 @@ class Tracker { Tracker(uint64_t timeout = UINT64_MAX) : timeout_(timeout) {} ~Tracker() { + mutex_.lock(); for (entry_t* entry : sig_list_) { assert(entry != NULL); while (1) { @@ -43,6 +46,7 @@ class Tracker { } Del(entry); } + mutex_.unlock(); } // Add tracker entry @@ -50,7 +54,9 @@ class Tracker { entry_t* entry = new entry_t{}; assert(entry); entry->tracker = this; + mutex_.lock(); entry->it = sig_list_.insert(sig_list_.begin(), entry); + mutex_.unlock(); entry->agent = agent; entry->orig = orig; @@ -73,7 +79,9 @@ class Tracker { // Delete tracker entry void Del(entry_t* entry) { hsa_signal_destroy(entry->signal); + mutex_.lock(); sig_list_.erase(entry->it); + mutex_.unlock(); delete entry; } @@ -84,6 +92,7 @@ class Tracker { hsa_status_t status = hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &record->complete); if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP)"); + if (record->complete == 0) EXC_RAISING(status, "hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP), time is zero"); hsa_amd_profiling_dispatch_time_t dispatch_time{}; status = hsa_amd_profiling_get_dispatch_time(entry->agent, entry->signal, &dispatch_time); @@ -106,6 +115,8 @@ class Tracker { uint64_t timeout_; // Tracked signals list sig_list_t sig_list_; + // Inter-thread synchronization + static mutex_t mutex_; }; } // namespace rocprofiler diff --git a/test/ctrl/tool.cpp b/test/ctrl/tool.cpp index 79fd6fffca..b53324d7c0 100644 --- a/test/ctrl/tool.cpp +++ b/test/ctrl/tool.cpp @@ -8,6 +8,7 @@ #include #include #include +#include #include #include #include @@ -243,10 +244,15 @@ void output_group(FILE* file, const rocprofiler_group_t* group, const char* str) } // Dump stored context profiling output data -void dump_context(context_entry_t* entry) { +bool dump_context(context_entry_t* entry) { hsa_status_t status = HSA_STATUS_ERROR; if (entry->valid) { + const rocprofiler_dispatch_record_t* record = entry->data.record; + if (record) { + if (record->complete == 0) return false; + } + ++context_collected; entry->valid = 0; const uint32_t index = entry->index; @@ -254,18 +260,22 @@ void dump_context(context_entry_t* entry) { const rocprofiler_feature_t* features = entry->features; const unsigned feature_count = entry->feature_count; - fprintf(file_handle, "dispatch[%u], queue_index(%lu), kernel_name(\"%s\"), time(%lu,%lu,%lu,%lu)\n", + fprintf(file_handle, "dispatch[%u], queue_index(%lu), kernel_name(\"%s\")", index, entry->data.queue_index, - entry->data.kernel_name, - entry->data.record->dispatch, - entry->data.record->begin, - entry->data.record->end, - entry->data.record->complete); + entry->data.kernel_name); + if (record) fprintf(file_handle, ", time(%lu,%lu,%lu,%lu)", + record->dispatch, + record->begin, + record->end, + record->complete); + fprintf(file_handle, "\n"); fflush(file_handle); - delete entry->data.record; - entry->data.record = NULL; + if (record) { + delete record; + entry->data.record = NULL; + } rocprofiler_group_t& group = entry->group; if (group.context != NULL) { @@ -285,6 +295,8 @@ void dump_context(context_entry_t* entry) { rocprofiler_close(group.context); } } + + return true; } // Dump all stored contexts profiling output data @@ -303,8 +315,9 @@ void dump_context_array() { } // Profiling completion handler -void handler(rocprofiler_group_t group, void* arg) { +bool handler(rocprofiler_group_t group, void* arg) { context_entry_t* entry = reinterpret_cast(arg); + bool ret = false; if (pthread_mutex_lock(&mutex) != 0) { perror("pthread_mutex_lock"); @@ -312,14 +325,16 @@ void handler(rocprofiler_group_t group, void* arg) { } if (context_array->find(entry->index) != context_array->end()) { - dump_context(entry); - dealloc_context_entry(entry); + if (dump_context(entry)) dealloc_context_entry(entry); + else ret = true; } if (pthread_mutex_unlock(&mutex) != 0) { perror("pthread_mutex_unlock"); exit(1); } + + return ret; } // Kernel disoatch callback diff --git a/test/gfx_metrics.xml b/test/gfx_metrics.xml index ca5a8009ef..1e05cefc27 100644 --- a/test/gfx_metrics.xml +++ b/test/gfx_metrics.xml @@ -40,7 +40,6 @@ - @@ -86,7 +85,7 @@ - + diff --git a/test/metrics.xml b/test/metrics.xml index 5aed1b7288..8f0e28e5b4 100644 --- a/test/metrics.xml +++ b/test/metrics.xml @@ -178,7 +178,7 @@ # WriteUnitStalled The percentage of GPUTime the Write unit is stalled. Value range: 0% to 100% (bad).