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

introducing hsa api duplicating table

Browse Source 
Change-Id: I4bbf933def018ad7b282fa08e587c53f5a44b336
This commit is contained in:
Evgeny
2019-03-05 19:42:26 -06:00
parent 5b266dc794
commit 89ff4ace93
4 changed files with 137 additions and 74 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/core/hsa_queue.h
+1 -25
View File
@@ -35,31 +35,7 @@ class HsaQueue : public Queue {
HsaQueue(const util::AgentInfo* agent_info, hsa_queue_t* queue) : queue_(queue) {}
void Submit(const packet_t* packet) {
// Compute the write index of queue and copy Aql packet into it
const uint64_t que_idx = hsa_queue_load_write_index_relaxed(queue_);
// Increment the write index
hsa_queue_store_write_index_relaxed(queue_, que_idx + 1);
const uint32_t mask = queue_->size - 1;
// Copy packet to the queue
const packet_word_t* src = reinterpret_cast<const packet_word_t*>(packet);
packet_t* slot = reinterpret_cast<packet_t*>(queue_->base_address) + (que_idx & mask);
packet_word_t* dst = reinterpret_cast<packet_word_t*>(slot);
const uint32_t nwords = sizeof(packet_t) / sizeof(packet_word_t);
for (unsigned i = 1; i < nwords; ++i) {
dst[i] = src[i];
}
// To maintain global order to ensure the prior copy of the packet contents is made visible
// before the header is updated.
// With in-order CP it will wait until the first packet in the blob will be valid
std::atomic<packet_word_t>* header_atomic_ptr =
reinterpret_cast<std::atomic<packet_word_t>*>(&dst[0]);
header_atomic_ptr->store(src[0], std::memory_order_release);
// Doorbell signaling
hsa_signal_store_relaxed(queue_->doorbell_signal, que_idx);
rocprofiler::util::HsaRsrcFactory::Instance().Submit(queue_, packet);
}
private:
src/core/tracker.h
+10 -8
View File
@@ -99,9 +99,9 @@ class Tracker {
entry->record = record;
// Creating a proxy signal
status = hsa_signal_create(1, 0, NULL, &(entry->signal));
status = hsa_api_.hsa_signal_create(1, 0, NULL, &(entry->signal));
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_signal_create");
status = hsa_amd_signal_async_handler(entry->signal, HSA_SIGNAL_CONDITION_LT, 1, Handler, entry);
status = hsa_api_.hsa_amd_signal_async_handler(entry->signal, HSA_SIGNAL_CONDITION_LT, 1, Handler, entry);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_signal_async_handler");
// Adding antry to the list
@@ -115,7 +115,7 @@ class Tracker {
// Delete tracker entry
void Delete(entry_t* entry) {
hsa_signal_destroy(entry->signal);
hsa_api_.hsa_signal_destroy(entry->signal);
mutex_.lock();
sig_list_.erase(entry->it);
mutex_.unlock();
@@ -151,7 +151,8 @@ class Tracker {
private:
Tracker() :
outstanding_(0),
hsa_rsrc_(&(util::HsaRsrcFactory::Instance()))
hsa_rsrc_(&(util::HsaRsrcFactory::Instance())),
hsa_api_(*(hsa_rsrc_->HsaApi()))
{}
~Tracker() {
@@ -181,13 +182,13 @@ class Tracker {
// Query begin/end and complete timestamps
if (entry->is_memcopy) {
hsa_amd_profiling_async_copy_time_t async_copy_time{};
hsa_status_t status = hsa_amd_profiling_get_async_copy_time(entry->signal, &async_copy_time);
hsa_status_t status = hsa_api_.hsa_amd_profiling_get_async_copy_time(entry->signal, &async_copy_time);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_async_copy_time");
record->begin = hsa_rsrc_->SysclockToNs(async_copy_time.start);
record->end = hsa_rsrc_->SysclockToNs(async_copy_time.end);
} else {
hsa_amd_profiling_dispatch_time_t dispatch_time{};
hsa_status_t status = hsa_amd_profiling_get_dispatch_time(entry->agent, entry->signal, &dispatch_time);
hsa_status_t status = hsa_api_.hsa_amd_profiling_get_dispatch_time(entry->agent, entry->signal, &dispatch_time);
if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_dispatch_time");
record->begin = hsa_rsrc_->SysclockToNs(dispatch_time.start);
record->end = hsa_rsrc_->SysclockToNs(dispatch_time.end);
@@ -204,9 +205,9 @@ class Tracker {
orig_signal_ptr->start_ts = prof_signal_ptr->start_ts;
orig_signal_ptr->end_ts = prof_signal_ptr->end_ts;
const hsa_signal_value_t new_value = hsa_signal_load_relaxed(orig) - 1;
const hsa_signal_value_t new_value = hsa_api_.hsa_signal_load_relaxed(orig) - 1;
if (signal_value != new_value) EXC_ABORT(HSA_STATUS_ERROR, "Tracker::Complete bad signal value");
hsa_signal_store_screlease(orig, signal_value);
hsa_api_.hsa_signal_store_screlease(orig, signal_value);
}
}
@@ -273,6 +274,7 @@ class Tracker {
std::atomic<uint64_t> outstanding_;
// HSA resources factory
util::HsaRsrcFactory* hsa_rsrc_;
const util::hsa_pfn_t& hsa_api_;
// Handling ordering enabled
static const bool ordering_enabled_ = false;
// Enable tracing