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

Move trace_buffer.h to the tool directory

Browse Source 
A trace buffer is used to efficiently store synchronous event records
so that they can be processed later, possibly in a different thread,
when the buffer is flushed. This helps reduce the latency added by
tracing API calls.

The API does not need to use trace buffers as synchronous events are
directly reported to the client with callbacks, and asynchronous events
(activities) are saved in memory pools.

The implentation of HSA asynchronous memory copy activities was using
a trace buffer shared with the tracer tool to write the records to a
file (async_copy_trace.txt), instead of using a memory pool and
reporting the activity to the client.

Removed the asynchronous memory copies trace buffer, and updated
hsa_async_copy_handler to use the pool specified when the activity
was enabled.

Updated the tracer tool to read HSA_OP_ID_COPY records out of the
default memory pool and write them to async_copy_trace.txt.

Move trace_buffer.h to test/tool as tracer_tool.cpp is now the only
file using it.

Change-Id: Ida95aba2eaf3c3f2a979ed6c2b060374017b7424
This commit is contained in:
Laurent Morichetti
2022-04-22 13:36:41 -07:00
committed by Laurent Morichetti
parent 48f4c82685
commit 61f35b0204
5 changed files with 77 additions and 77 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/core/roctracer.cpp
+20 -30
View File
@@ -40,7 +40,6 @@
#include "core/journal.h"
#include "core/loader.h"
#include "core/memory_pool.h"
#include "core/trace_buffer.h"
#include "core/tracker.h"
#include "ext/hsa_rt_utils.hpp"
#include "util/exception.h"
@@ -229,19 +228,12 @@ class hip_act_cb_tracker_t {
std::unordered_map<uint32_t, uint32_t> data_;
};
void hsa_async_copy_handler(Tracker::entry_t* entry);
constexpr TraceBuffer<trace_entry_t>::flush_prm_t trace_buffer_prm[] = {
{COPY_ENTRY_TYPE, hsa_async_copy_handler}};
TraceBuffer<trace_entry_t>* trace_buffer = NULL;
namespace hsa_support {
// callbacks table
cb_table_t cb_table;
// async copy activity callback
bool async_copy_callback_enabled = false;
activity_async_callback_t async_copy_callback_fun = NULL;
void* async_copy_callback_arg = NULL;
const char* output_prefix = NULL;
MemoryPool* async_copy_callback_memory_pool = nullptr;
// Table of function pointers to HSA Core Runtime
CoreApiTable CoreApiTable_saved{};
// Table of function pointers to AMD extensions
@@ -546,15 +538,14 @@ void close_output_file(FILE* file_handle) {
if ((file_handle != NULL) && (file_handle != stdout)) fclose(file_handle);
}
void hsa_async_copy_handler(Tracker::entry_t* entry) {
void hsa_async_copy_handler(const Tracker::entry_t* entry) {
activity_record_t record{};
record.domain = ACTIVITY_DOMAIN_HSA_OPS; // activity domain id
record.begin_ns = entry->begin; // host begin timestamp
record.end_ns = entry->end; // host end timestamp
record.device_id = 0; // device id
hsa_support::async_copy_callback_fun(hsa_support::HSA_OP_ID_async_copy, &record,
hsa_support::async_copy_callback_arg);
record.domain = ACTIVITY_DOMAIN_HSA_OPS;
record.op = HSA_OP_ID_COPY;
record.begin_ns = entry->begin;
record.end_ns = entry->end;
record.device_id = 0;
entry->pool->Write(record);
}
hsa_status_t hsa_amd_memory_async_copy_interceptor(void* dst, hsa_agent_t dst_agent,
@@ -564,8 +555,10 @@ hsa_status_t hsa_amd_memory_async_copy_interceptor(void* dst, hsa_agent_t dst_ag
hsa_signal_t completion_signal) {
hsa_status_t status = HSA_STATUS_SUCCESS;
if (hsa_support::async_copy_callback_enabled) {
trace_entry_t* entry = trace_buffer->GetEntry();
Tracker::Enable(COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
Tracker::entry_t* entry = new Tracker::entry_t();
entry->handler = hsa_async_copy_handler;
entry->pool = hsa_support::async_copy_callback_memory_pool;
Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
status = hsa_amd_memory_async_copy_fn(dst, dst_agent, src, src_agent, size, num_dep_signals,
dep_signals, entry->signal);
if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry);
@@ -583,8 +576,10 @@ hsa_status_t hsa_amd_memory_async_copy_rect_interceptor(
hsa_signal_t completion_signal) {
hsa_status_t status = HSA_STATUS_SUCCESS;
if (hsa_support::async_copy_callback_enabled) {
trace_entry_t* entry = trace_buffer->GetEntry();
Tracker::Enable(COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
Tracker::entry_t* entry = new Tracker::entry_t();
entry->handler = hsa_async_copy_handler;
entry->pool = hsa_support::async_copy_callback_memory_pool;
Tracker::Enable(Tracker::COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
status = hsa_amd_memory_async_copy_rect_fn(dst, dst_offset, src, src_offset, range, copy_agent,
dir, num_dep_signals, dep_signals, entry->signal);
if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry);
@@ -625,7 +620,6 @@ unsigned set_stopped(unsigned val) {
} // namespace roctracer
LOADER_INSTANTIATE();
TRACE_BUFFER_INSTANTIATE();
///////////////////////////////////////////////////////////////////////////////////////////////////
// Public library methods
@@ -968,6 +962,8 @@ static roctracer_status_t roctracer_enable_activity_fun(roctracer_domain_t domai
case ACTIVITY_DOMAIN_HSA_OPS: {
if (op == HSA_OP_ID_COPY) {
roctracer::hsa_support::async_copy_callback_enabled = true;
roctracer::hsa_support::async_copy_callback_memory_pool =
reinterpret_cast<roctracer::MemoryPool*>(pool);
} else {
const bool init_phase = (roctracer::RocpLoader::GetRef() == NULL);
if (roctracer::RocpLoader::GetRef() == NULL) break;
@@ -1059,7 +1055,8 @@ static roctracer_status_t roctracer_disable_activity_fun(roctracer_domain_t doma
switch (domain) {
case ACTIVITY_DOMAIN_HSA_OPS: {
if (op == HSA_OP_ID_COPY) {
roctracer::hsa_support::async_copy_callback_enabled = true;
roctracer::hsa_support::async_copy_callback_enabled = false;
roctracer::hsa_support::async_copy_callback_memory_pool = nullptr;
} else {
if (roctracer::RocpLoader::GetRef() == NULL) break;
const bool succ = roctracer::RocpLoader::Instance().EnableActivityCallback(op, false);
@@ -1145,7 +1142,6 @@ PUBLIC_API roctracer_status_t roctracer_flush_activity_expl(roctracer_pool_t* po
if (pool == NULL) pool = roctracer_default_pool();
roctracer::MemoryPool* memory_pool = reinterpret_cast<roctracer::MemoryPool*>(pool);
if (memory_pool != NULL) memory_pool->Flush();
roctracer::TraceBufferBase::FlushAll();
API_METHOD_SUFFIX
}
@@ -1223,9 +1219,6 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(roctracer_domain_t domain
roctracer::hsa_ops_properties_t* ops_properties =
reinterpret_cast<roctracer::hsa_ops_properties_t*>(properties);
HsaApiTable* table = reinterpret_cast<HsaApiTable*>(ops_properties->table);
roctracer::hsa_support::async_copy_callback_fun = ops_properties->async_copy_callback_fun;
roctracer::hsa_support::async_copy_callback_arg = ops_properties->async_copy_callback_arg;
roctracer::hsa_support::output_prefix = ops_properties->output_prefix;
#if 0
// HSA dispatches intercepting
@@ -1314,15 +1307,12 @@ PUBLIC_API void roctracer_unload() {
PUBLIC_API void roctracer_flush_buf() {
ONLOAD_TRACE_BEG();
roctracer::trace_buffer->Flush();
ONLOAD_TRACE_END();
}
CONSTRUCTOR_API void constructor() {
ONLOAD_TRACE_BEG();
roctracer::util::Logger::Create();
roctracer::trace_buffer = new roctracer::TraceBuffer<roctracer::trace_entry_t>(
"HSA GPU", 0x200000, roctracer::trace_buffer_prm, 2);
roctracer_load();
ONLOAD_TRACE_END();
}
src/core/trace_buffer.h
-328
View File
@@ -1,328 +0,0 @@
/* Copyright (c) 2018-2022 Advanced Micro Devices, Inc.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
#ifndef SRC_CORE_TRACE_BUFFER_H_
#define SRC_CORE_TRACE_BUFFER_H_
#include <atomic>
#include <iostream>
#include <list>
#include <mutex>
#include <sstream>
#include <pthread.h>
#include <string.h>
#include <unistd.h>
#define FATAL(stream) \
do { \
std::ostringstream oss; \
oss << __FUNCTION__ << "(), " << stream; \
std::cout << oss.str() << std::endl; \
abort(); \
} while (0)
#define PTHREAD_CALL(call) \
do { \
int err = call; \
if (err != 0) { \
errno = err; \
perror(#call); \
abort(); \
} \
} while (0)
namespace roctracer {
enum { TRACE_ENTRY_INV = 0, TRACE_ENTRY_INIT = 1, TRACE_ENTRY_COMPL = 2 };
enum entry_type_t {
DFLT_ENTRY_TYPE = 0,
API_ENTRY_TYPE = 1,
COPY_ENTRY_TYPE = 2,
KERNEL_ENTRY_TYPE = 3,
NUM_ENTRY_TYPE = 4
};
struct trace_entry_t {
std::atomic<uint32_t> valid;
entry_type_t type;
uint64_t dispatch;
uint64_t begin; // kernel begin timestamp, ns
uint64_t end; // kernel end timestamp, ns
uint64_t complete;
hsa_agent_t agent;
uint32_t dev_index;
hsa_signal_t orig;
hsa_signal_t signal;
union {
struct {
} copy;
struct {
const char* name;
hsa_agent_t agent;
uint32_t tid;
} kernel;
};
};
template <class T> struct push_element_fun {
T* const elem_;
T** prev_;
bool fun(T* node) {
if (node->priority_ > elem_->priority_) {
*prev_ = elem_;
elem_->next_elem_ = node;
} else if (node->next_elem_ == NULL) {
node->next_elem_ = elem_;
} else {
prev_ = &(node->next_elem_);
return false;
}
return true;
}
push_element_fun(T* elem, T** prev) : elem_(elem), prev_(prev) {}
};
template <class T> struct call_element_fun {
void (T::*fptr_)();
bool fun(T* node) const {
(node->*fptr_)();
return false;
}
call_element_fun(void (T::*f)()) : fptr_(f) {}
};
struct TraceBufferBase {
typedef std::mutex mutex_t;
virtual void StartWorkerThread() = 0;
virtual void Flush() = 0;
static void StartWorkerThreadAll() {
foreach (call_element_fun<TraceBufferBase>(&TraceBufferBase::StartWorkerThread))
;
}
static void FlushAll() {
foreach (call_element_fun<TraceBufferBase>(&TraceBufferBase::Flush))
;
}
static void Push(TraceBufferBase* elem) {
if (head_elem_ == NULL)
head_elem_ = elem;
else
foreach (push_element_fun<TraceBufferBase>(elem, &head_elem_))
;
}
TraceBufferBase(const uint32_t& prior) : priority_(prior), next_elem_(NULL) {}
template <class F> static void foreach (const F& f_in) {
std::lock_guard<mutex_t> lck(mutex_);
F f = f_in;
TraceBufferBase* p = head_elem_;
while (p != NULL) {
TraceBufferBase* next = p->next_elem_;
if (f.fun(p) == true) break;
p = next;
}
}
const uint32_t priority_;
TraceBufferBase* next_elem_;
static TraceBufferBase* head_elem_;
static mutex_t mutex_;
};
template <typename Entry> class TraceBuffer : protected TraceBufferBase {
public:
typedef void (*callback_t)(Entry*);
typedef TraceBuffer<Entry> Obj;
typedef uint64_t pointer_t;
typedef std::recursive_mutex mutex_t;
typedef typename std::list<Entry*> buf_list_t;
typedef typename buf_list_t::iterator buf_list_it_t;
struct flush_prm_t {
entry_type_t type;
callback_t fun;
};
TraceBuffer(const char* name, uint32_t size, const flush_prm_t* flush_prm_arr,
uint32_t flush_prm_count, uint32_t prior = 0)
: TraceBufferBase(prior), size_(size), work_thread_started_(false) {
name_ = strdup(name);
data_ = allocate_fun();
next_ = allocate_fun();
read_pointer_ = 0;
write_pointer_ = 0;
end_pointer_ = size;
buf_list_.push_back(data_);
memset(f_array_, 0, sizeof(f_array_));
for (const flush_prm_t* prm = flush_prm_arr; prm < flush_prm_arr + flush_prm_count; prm++) {
const entry_type_t type = prm->type;
if (type >= NUM_ENTRY_TYPE) FATAL("out of f_array bounds (" << type << ")");
if (f_array_[type] != NULL) FATAL("handler function ptr redefinition (" << type << ")");
f_array_[type] = prm->fun;
}
TraceBufferBase::Push(this);
}
~TraceBuffer() {
StopWorkerThread();
Flush();
}
void StartWorkerThread() {
std::lock_guard<mutex_t> lck(mutex_);
if (work_thread_started_ == false) {
PTHREAD_CALL(pthread_mutex_init(&work_mutex_, NULL));
PTHREAD_CALL(pthread_cond_init(&work_cond_, NULL));
PTHREAD_CALL(pthread_create(&work_thread_, NULL, allocate_worker, this));
work_thread_started_ = true;
}
}
void StopWorkerThread() {
std::lock_guard<mutex_t> lck(mutex_);
if (work_thread_started_ == true) {
PTHREAD_CALL(pthread_cancel(work_thread_));
void* res;
PTHREAD_CALL(pthread_join(work_thread_, &res));
if (res != PTHREAD_CANCELED) FATAL("consumer thread wasn't stopped correctly");
work_thread_started_ = false;
}
}
Entry* GetEntry() {
const pointer_t pointer = write_pointer_.fetch_add(1);
if (pointer >= end_pointer_) wrap_buffer(pointer);
if (pointer >= end_pointer_) FATAL("pointer >= end_pointer_ after buffer wrap");
Entry* entry = data_ + (size_ + pointer - end_pointer_);
entry->valid = TRACE_ENTRY_INV;
entry->type = DFLT_ENTRY_TYPE;
return entry;
}
void Flush() { flush_buf(); }
private:
void flush_buf() {
std::lock_guard<mutex_t> lck(mutex_);
pointer_t pointer = read_pointer_;
pointer_t curr_pointer = write_pointer_.load(std::memory_order_relaxed);
buf_list_it_t it = buf_list_.begin();
buf_list_it_t end_it = buf_list_.end();
while (it != end_it) {
Entry* buf = *it;
Entry* ptr = buf + (pointer % size_);
Entry* end_ptr = buf + size_;
while ((ptr < end_ptr) && (pointer < curr_pointer)) {
if (ptr->valid != TRACE_ENTRY_COMPL) break;
entry_type_t type = ptr->type;
if (type >= NUM_ENTRY_TYPE) FATAL("out of f_array bounds (" << type << ")");
callback_t f_ptr = f_array_[type];
if (f_ptr == NULL) FATAL("f_ptr == NULL");
(*f_ptr)(ptr);
ptr++;
pointer++;
}
buf_list_it_t prev = it;
it++;
if (ptr == end_ptr) {
free_fun(*prev);
buf_list_.erase(prev);
}
if (pointer == curr_pointer) break;
}
read_pointer_ = pointer;
}
inline Entry* allocate_fun() {
Entry* ptr = (Entry*)malloc(size_ * sizeof(Entry));
if (ptr == NULL) FATAL("malloc failed");
// memset(ptr, 0, size_ * sizeof(Entry));
return ptr;
}
inline void free_fun(void* ptr) { free(ptr); }
static void* allocate_worker(void* arg) {
Obj* obj = (Obj*)arg;
while (1) {
PTHREAD_CALL(pthread_mutex_lock(&(obj->work_mutex_)));
while (obj->next_ != NULL) {
PTHREAD_CALL(pthread_cond_wait(&(obj->work_cond_), &(obj->work_mutex_)));
}
obj->next_ = obj->allocate_fun();
PTHREAD_CALL(pthread_mutex_unlock(&(obj->work_mutex_)));
}
return NULL;
}
void wrap_buffer(const pointer_t pointer) {
std::lock_guard<mutex_t> lck(mutex_);
if (work_thread_started_ == false) StartWorkerThread();
PTHREAD_CALL(pthread_mutex_lock(&work_mutex_));
if (pointer >= end_pointer_) {
data_ = next_;
next_ = NULL;
PTHREAD_CALL(pthread_cond_signal(&work_cond_));
end_pointer_ += size_;
if (end_pointer_ == 0) FATAL("pointer overflow");
buf_list_.push_back(data_);
}
PTHREAD_CALL(pthread_mutex_unlock(&work_mutex_));
}
const char* name_;
const uint32_t size_;
Entry* data_;
Entry* next_;
pointer_t read_pointer_;
volatile std::atomic<pointer_t> write_pointer_;
volatile std::atomic<pointer_t> end_pointer_;
buf_list_t buf_list_;
callback_t f_array_[NUM_ENTRY_TYPE];
pthread_t work_thread_;
pthread_mutex_t work_mutex_;
pthread_cond_t work_cond_;
bool work_thread_started_;
mutex_t mutex_;
};
} // namespace roctracer
#define TRACE_BUFFER_INSTANTIATE() \
roctracer::TraceBufferBase* roctracer::TraceBufferBase::head_elem_ = NULL; \
roctracer::TraceBufferBase::mutex_t roctracer::TraceBufferBase::mutex_;
#endif // SRC_CORE_TRACE_BUFFER_H_
src/core/tracker.h
+43 -9
View File
@@ -31,14 +31,45 @@
#include "util/hsa_rsrc_factory.h"
#include "util/exception.h"
#include "util/logger.h"
#include "core/trace_buffer.h"
namespace roctracer {
class Tracker {
public:
typedef ::util::HsaRsrcFactory::timestamp_t timestamp_t;
typedef roctracer::trace_entry_t entry_t;
typedef roctracer::entry_type_t entry_type_t;
enum { ENTRY_INV = 0, ENTRY_INIT = 1, ENTRY_COMPL = 2 };
enum entry_type_t {
DFLT_ENTRY_TYPE = 0,
API_ENTRY_TYPE = 1,
COPY_ENTRY_TYPE = 2,
KERNEL_ENTRY_TYPE = 3,
NUM_ENTRY_TYPE = 4
};
struct entry_t {
std::atomic<uint32_t> valid;
entry_type_t type;
uint64_t dispatch;
uint64_t begin; // kernel begin timestamp, ns
uint64_t end; // kernel end timestamp, ns
uint64_t complete;
hsa_agent_t agent;
uint32_t dev_index;
hsa_signal_t orig;
hsa_signal_t signal;
void (*handler)(const entry_t*);
MemoryPool* pool;
union {
struct {
} copy;
struct {
const char* name;
hsa_agent_t agent;
uint32_t tid;
} kernel;
};
};
// Add tracker entry
inline static void Enable(entry_type_t type, const hsa_agent_t& agent, const hsa_signal_t& signal,
@@ -52,7 +83,7 @@ class Tracker {
entry->dev_index = 0; // hsa_rsrc->GetAgentInfo(agent)->dev_index;
entry->orig = signal;
entry->dispatch = hsa_rsrc->TimestampNs();
entry->valid.store(roctracer::TRACE_ENTRY_INIT, std::memory_order_release);
entry->valid.store(ENTRY_INIT, std::memory_order_release);
// Creating a proxy signal
status = hsa_signal_create(1, 0, NULL, &(entry->signal));
@@ -67,7 +98,7 @@ class Tracker {
// Delete tracker entry
inline static void Disable(entry_t* entry) {
hsa_signal_destroy(entry->signal);
entry->valid.store(roctracer::TRACE_ENTRY_INV, std::memory_order_release);
entry->valid.store(ENTRY_INV, std::memory_order_release);
}
private:
@@ -75,7 +106,7 @@ class Tracker {
inline static void Complete(hsa_signal_value_t signal_value, entry_t* entry) {
// Query begin/end and complete timestamps
::util::HsaRsrcFactory* hsa_rsrc = &(::util::HsaRsrcFactory::Instance());
if (entry->type == roctracer::COPY_ENTRY_TYPE) {
if (entry->type == COPY_ENTRY_TYPE) {
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);
if (status != HSA_STATUS_SUCCESS)
@@ -98,7 +129,10 @@ class Tracker {
hsa_signal_t signal = entry->signal;
// Releasing completed entry
entry->valid.store(roctracer::TRACE_ENTRY_COMPL, std::memory_order_release);
entry->valid.store(ENTRY_COMPL, std::memory_order_release);
assert(entry->handler != nullptr);
entry->handler(entry);
// Original intercepted signal completion
if (orig.handle) {
@@ -112,14 +146,14 @@ class Tracker {
hsa_signal_store_screlease(orig, signal_value);
}
hsa_signal_destroy(signal);
delete entry;
}
// Handler for packet completion
static bool Handler(hsa_signal_value_t signal_value, void* arg) {
// Acquire entry
entry_t* entry = reinterpret_cast<entry_t*>(arg);
while (entry->valid.load(std::memory_order_acquire) != roctracer::TRACE_ENTRY_INIT)
sched_yield();
while (entry->valid.load(std::memory_order_acquire) != ENTRY_INIT) sched_yield();
// Complete entry
Tracker::Complete(signal_value, entry);