AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

counters dumping optimization

Browse Source 
Change-Id: I8c694e5380e15179453148dd9ab3a3e51b6db861
This commit is contained in:
Evgeny
2020-07-07 16:18:29 -05:00
parent f96c9f49aa
commit 2a7f77b290
13 changed files with 415 additions and 90 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/core/context.h
+33 -6
View File
@@ -183,7 +183,7 @@ class Context {
uint32_t GetGroupCount() const { return set_.size(); }
inline rocprofiler_group_t GetGroupInfo(Group* g) {
inline rocprofiler_group_t GetGroupDescr(Group* g) {
rocprofiler::info_vector_t& info_vector = g->GetInfoVector();
rocprofiler_group_t group = {};
group.index = g->GetIndex();
@@ -192,12 +192,12 @@ class Context {
group.feature_count = info_vector.size();
return group;
}
inline rocprofiler_group_t GetGroupInfo(const uint32_t& index) {
inline rocprofiler_group_t GetGroupDescr(const uint32_t& index) {
rocprofiler_group_t group = {};
if (set_.empty()) {
group.context = reinterpret_cast<rocprofiler_t*>(this);
} else {
group = GetGroupInfo(&set_[index]);
group = GetGroupDescr(&set_[index]);
}
return group;
}
@@ -288,8 +288,8 @@ class Context {
Context* context = group->GetContext();
auto r = group->FetchDecrRefsCount();
if (r == 1) {
const rocprofiler_group_t group_info = context->GetGroupInfo(group);
context->handler_(group_info, context->handler_arg_);
const rocprofiler_group_t group_descr = context->GetGroupDescr(group);
context->handler_(group_descr, context->handler_arg_);
}
return false;
}
@@ -298,6 +298,22 @@ class Context {
Group* GetGroup(const uint32_t& index) { return &set_[index]; }
rocprofiler_handler_t GetHandler(void** arg) const { *arg = handler_arg_; return handler_; }
void SetDispatchSignal(const hsa_signal_t &signal) {
dispatch_signal_ = signal;
}
hsa_signal_t& GetDispatchSignal() {
return dispatch_signal_;
}
void SetOrigSignal(const hsa_signal_t &signal) {
orig_signal_ = signal;
}
const hsa_signal_t& GetOrigSignal() const {
return orig_signal_;
}
rocprofiler_dispatch_record_t* GetRecord() {
return &record_;
}
private:
Context(const util::AgentInfo* agent_info, Queue* queue, rocprofiler_feature_t* info,
const uint32_t info_count, rocprofiler_handler_t handler, void* handler_arg)
@@ -309,7 +325,10 @@ class Context {
metrics_(NULL),
handler_(handler),
handler_arg_(handler_arg),
pcsmp_mode_(false)
pcsmp_mode_(false),
dispatch_signal_{},
orig_signal_{},
record_{}
{}
~Context() { Destruct(); }
@@ -355,6 +374,9 @@ class Context {
}
}
}
hsa_status_t status = hsa_signal_create(1, 0, NULL, &dispatch_signal_);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "MetricsDict create failed");
}
// Initialize rocprofiler context
@@ -593,6 +615,11 @@ class Context {
// PC sampling mode
bool pcsmp_mode_;
// kernel packet dispatch copmletion signal
hsa_signal_t dispatch_signal_;
hsa_signal_t orig_signal_;
rocprofiler_dispatch_record_t record_;
};
} // namespace rocprofiler
src/core/intercept_queue.cpp
+1
View File
@@ -43,5 +43,6 @@ rocprofiler_hsa_callback_fun_t InterceptQueue::submit_callback_fun_ = NULL;
void* InterceptQueue::submit_callback_arg_ = NULL;
bool InterceptQueue::k_concurrent_ = false;
bool InterceptQueue::opt_mode_ = false;
} // namespace rocprofiler
src/core/intercept_queue.h
+82
View File
@@ -92,6 +92,8 @@ class InterceptQueue {
(*obj_map_)[(uint64_t)(*queue)] = obj;
if (k_concurrent_) {
status = proxy->SetInterceptCB(OnSubmitCB_SQTT, obj);
} else if (opt_mode_) {
status = proxy->SetInterceptCB(OnSubmitCB_opt, obj);
} else {
status = proxy->SetInterceptCB(OnSubmitCB, obj);
}
@@ -138,6 +140,85 @@ class InterceptQueue {
return status;
}
static void OnSubmitCB_opt(const void* in_packets, uint64_t count, uint64_t user_que_idx, void* data,
hsa_amd_queue_intercept_packet_writer writer) {
const packet_t* packets_arr = reinterpret_cast<const packet_t*>(in_packets);
InterceptQueue* obj = reinterpret_cast<InterceptQueue*>(data);
Queue* proxy = obj->proxy_;
// Travers input packets
for (uint64_t j = 0; j < count; ++j) {
const packet_t* packet = &packets_arr[j];
bool to_submit = true;
// Checking for dispatch packet type
if ((GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) &&
(dispatch_callback_.load(std::memory_order_acquire) != NULL)) {
const hsa_kernel_dispatch_packet_t* dispatch_packet =
reinterpret_cast<const hsa_kernel_dispatch_packet_t*>(packet);
const hsa_signal_t completion_signal = dispatch_packet->completion_signal;
#if 0
// Prepareing dispatch callback data
uint64_t kernel_object = dispatch_packet->kernel_object;
const amd_kernel_code_t* kernel_code = GetKernelCode(kernel_object);
const char* kernel_name = QueryKernelName(kernel_object, kernel_code);
#endif
rocprofiler_callback_data_t data = {obj->agent_info_->dev_id,
obj->agent_info_->dev_index,
obj->queue_,
user_que_idx,
obj->queue_id,
completion_signal,
dispatch_packet,
NULL, // kernel_name
0, // kernel_object
NULL, // kernel_code
0, // (uint32_t)syscall(__NR_gettid),
NULL};
// Calling dispatch callback
rocprofiler_group_t group = {};
hsa_status_t status = (dispatch_callback_.load())(&data, callback_data_, &group);
#if 0
free(const_cast<char*>(kernel_name));
#endif
Context* context = reinterpret_cast<Context*>(group.context);
// Injecting profiling start/stop packets
if ((status == HSA_STATUS_SUCCESS) && (context != NULL)) {
if (group.feature_count != 0) {
if (tracker_ != NULL) {
const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal = context->GetDispatchSignal();
Group* context_group = context->GetGroup(group.index);
Tracker::Enable_opt(context_group, completion_signal);
context_group->IncrRefsCount();
}
const pkt_vector_t& start_vector = context->StartPackets(group.index);
const pkt_vector_t& stop_vector = context->StopPackets(group.index);
pkt_vector_t packets = start_vector;
packets.insert(packets.end(), *packet);
packets.insert(packets.end(), stop_vector.begin(), stop_vector.end());
if (writer != NULL) {
writer(&packets[0], packets.size());
} else {
proxy->Submit(&packets[0], packets.size());
}
to_submit = false;
}
}
}
// Submitting the original packets if profiling was not enabled
if (to_submit) {
if (writer != NULL) {
writer(packet, 1);
} else {
proxy->Submit(packet, 1);
}
}
}
}
static void OnSubmitCB(const void* in_packets, uint64_t count, uint64_t user_que_idx, void* data,
hsa_amd_queue_intercept_packet_writer writer) {
const packet_t* packets_arr = reinterpret_cast<const packet_t*>(in_packets);
@@ -400,6 +481,7 @@ class InterceptQueue {
static bool IsTrackerOn() { return tracker_on_; }
static bool k_concurrent_;
static bool opt_mode_;
private:
static void queue_event_callback(hsa_status_t status, hsa_queue_t *queue, void *arg) {
src/core/rocprofiler.cpp
+3 -3
View File
@@ -217,6 +217,7 @@ uint32_t LoadTool() {
if (settings.memcopy_tracking) intercept_mode |= MEMCOPY_INTERCEPT_MODE;
if (settings.hsa_intercepting) intercept_mode |= HSA_INTERCEPT_MODE;
if (settings.k_concurrent) InterceptQueue::k_concurrent_ = true;
if (settings.opt_mode) InterceptQueue::opt_mode_ = true;
}
ONLOAD_TRACE("end intercept_mode(" << intercept_mode << ")");
@@ -537,8 +538,7 @@ PUBLIC_API hsa_status_t rocprofiler_open(hsa_agent_t agent, rocprofiler_feature_
if (mode != 0) {
if (mode & ROCPROFILER_MODE_STANDALONE) {
if (mode & ROCPROFILER_MODE_CREATEQUEUE) {
if (hsa_rsrc->CreateQueue(agent_info, properties->queue_depth, &(properties->queue)) ==
false) {
if (hsa_rsrc->CreateQueue(agent_info, properties->queue_depth, &(properties->queue)) == false) {
EXC_RAISING(HSA_STATUS_ERROR, "CreateQueue() failed");
}
}
@@ -592,7 +592,7 @@ PUBLIC_API hsa_status_t rocprofiler_get_group(rocprofiler_t* handle, uint32_t gr
rocprofiler_group_t* group) {
API_METHOD_PREFIX
rocprofiler::Context* context = reinterpret_cast<rocprofiler::Context*>(handle);
*group = context->GetGroupInfo(group_index);
*group = context->GetGroupDescr(group_index);
API_METHOD_SUFFIX
}
src/core/tracker.h
+43
View File
@@ -155,6 +155,49 @@ class Tracker {
Enable(entry, reinterpret_cast<void*>(handler), arg);
}
// Enable tracking
static void Enable_opt(Group* group, const hsa_signal_t& orig_signal) {
Context* context = group->GetContext();
context->SetOrigSignal(orig_signal);
context->GetRecord()->dispatch = util::HsaRsrcFactory::Instance().TimestampNs();
// Creating a proxy signal
const hsa_signal_value_t signal_value = (orig_signal.handle) ?
util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_load_relaxed(orig_signal) : 1;
hsa_signal_t& dispatch_signal = context->GetDispatchSignal();
util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_store_screlease(dispatch_signal, signal_value);
hsa_status_t status =
util::HsaRsrcFactory::Instance().HsaApi()->hsa_amd_signal_async_handler(dispatch_signal, HSA_SIGNAL_CONDITION_LT, signal_value, Handler, group);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_signal_async_handler");
}
// Tracker handler
static bool Handler_opt(hsa_signal_value_t signal_value, void* arg) {
Group* group = reinterpret_cast<Group*>(arg);
Context* context = group->GetContext();
hsa_signal_t dispatch_signal = context->GetDispatchSignal();
record_t* record = context->GetRecord();
hsa_amd_profiling_dispatch_time_t dispatch_time{};
hsa_status_t status =
util::HsaRsrcFactory::Instance().HsaApi()->hsa_amd_profiling_get_dispatch_time(context->GetAgent(), dispatch_signal, &dispatch_time);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_dispatch_time");
record->begin = util::HsaRsrcFactory::Instance().SysclockToNs(dispatch_time.start);
record->end = util::HsaRsrcFactory::Instance().SysclockToNs(dispatch_time.end);
record->complete = util::HsaRsrcFactory::Instance().TimestampNs();
// Original intercepted signal completion
const hsa_signal_t& orig_signal = context->GetOrigSignal();
if (orig_signal.handle) {
amd_signal_t* orig_signal_ptr = reinterpret_cast<amd_signal_t*>(orig_signal.handle);
amd_signal_t* prof_signal_ptr = reinterpret_cast<amd_signal_t*>(dispatch_signal.handle);
orig_signal_ptr->start_ts = prof_signal_ptr->start_ts;
orig_signal_ptr->end_ts = prof_signal_ptr->end_ts;
util::HsaRsrcFactory::Instance().HsaApi()->hsa_signal_store_screlease(orig_signal, signal_value);
}
Context::Handler(signal_value, arg);
}
private:
Tracker() :
outstanding_(0),