AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Remove unused proxy utilities

Przeglądaj źródła 
The proxy queue implements packet interception to enable timestamps
collection. As it is, the roctracer is not intercepting packets, and
instead relies on the rocprofiler tool to collect the timestamps for
kernel dispatches.

This is an issue as the roctracer API does not implement HSA_OPS
activities for kernel dispatches. This will be addressed in a future
commit.

Change-Id: Ib6a778a513410bec4579f223a9d9e9fd9b6054df


[ROCm/roctracer commit: 6b06322578]
This commit is contained in:
Laurent Morichetti
2022-04-22 13:22:30 -07:00
zatwierdzone przez Laurent Morichetti
rodzic 4a50f3b88f
commit 159a56ffff
13 zmienionych plików z 16 dodań i 1043 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/roctracer/src/CMakeLists.txt
-3
Wyświetl plik
@@ -40,9 +40,6 @@ execute_process ( COMMAND sh -xc "ln -s ${ROOT_DIR}/../rocprofiler/src/core/acti
set ( TARGET_LIB ${TARGET_NAME} )
set ( LIB_SRC
${LIB_DIR}/core/roctracer.cpp
${LIB_DIR}/proxy/proxy_queue.cpp
${LIB_DIR}/proxy/simple_proxy_queue.cpp
${LIB_DIR}/proxy/intercept_queue.cpp
${LIB_DIR}/util/hsa_rsrc_factory.cpp
)
add_library ( ${TARGET_LIB} ${LIBRARY_TYPE} ${LIB_SRC} )
projects/roctracer/src/core/roctracer.cpp
+8 -53
Wyświetl plik
@@ -40,17 +40,12 @@
#include "core/loader.h"
#include "core/memory_pool.h"
#include "core/trace_buffer.h"
#include "proxy/tracker.h"
#include "core/tracker.h"
#include "ext/hsa_rt_utils.hpp"
#include "util/exception.h"
#include "util/hsa_rsrc_factory.h"
#include "util/logger.h"
#include "proxy/hsa_queue.h"
#include "proxy/intercept_queue.h"
#include "proxy/proxy_queue.h"
#include "proxy/simple_proxy_queue.h"
#define PUBLIC_API __attribute__((visibility("default")))
#define CONSTRUCTOR_API __attribute__((constructor))
#define DESTRUCTOR_API __attribute__((destructor))
@@ -233,10 +228,9 @@ class hip_act_cb_tracker_t {
std::unordered_map<uint32_t, uint32_t> data_;
};
void hsa_async_copy_handler(::proxy::Tracker::entry_t* entry);
void hsa_kernel_handler(::proxy::Tracker::entry_t* entry);
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}, {KERNEL_ENTRY_TYPE, hsa_kernel_handler}};
{COPY_ENTRY_TYPE, hsa_async_copy_handler}};
TraceBuffer<trace_entry_t>* trace_buffer = NULL;
namespace hsa_support {
@@ -554,45 +548,7 @@ void close_output_file(FILE* file_handle) {
if ((file_handle != NULL) && (file_handle != stdout)) fclose(file_handle);
}
FILE* kernel_file_handle = NULL;
void hsa_kernel_handler(::proxy::Tracker::entry_t* entry) {
static uint64_t index = 0;
if (index == 0) {
kernel_file_handle = open_output_file(hsa_support::output_prefix, "results.txt");
}
fprintf(kernel_file_handle,
"dispatch[%lu], gpu-id(%u), tid(%u), kernel-name(\"%s\"), time(%lu,%lu,%lu,%lu)\n", index,
//::util::HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index,
entry->dev_index, entry->kernel.tid, entry->kernel.name, entry->dispatch, entry->begin,
entry->end, entry->complete);
#if 0
fprintf(file_handle, "dispatch[%u], gpu-id(%u), queue-id(%u), queue-index(%lu), tid(%lu), grd(%u), wgr(%u), lds(%u), scr(%u), vgpr(%u), sgpr(%u), fbar(%u), sig(0x%lx), kernel-name(\"%s\")",
index,
HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index,
entry->data.queue_id,
entry->data.queue_index,
entry->data.thread_id,
entry->kernel_properties.grid_size,
entry->kernel_properties.workgroup_size,
entry->kernel_properties.lds_size,
entry->kernel_properties.scratch_size,
entry->kernel_properties.vgpr_count,
entry->kernel_properties.sgpr_count,
entry->kernel_properties.fbarrier_count,
entry->kernel_properties.signal.handle,
nik_name.c_str());
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);
#endif
index++;
}
void hsa_async_copy_handler(::proxy::Tracker::entry_t* entry) {
void hsa_async_copy_handler(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
@@ -611,10 +567,10 @@ hsa_status_t hsa_amd_memory_async_copy_interceptor(void* dst, hsa_agent_t dst_ag
hsa_status_t status = HSA_STATUS_SUCCESS;
if (hsa_support::async_copy_callback_enabled) {
trace_entry_t* entry = trace_buffer->GetEntry();
::proxy::Tracker::Enable(COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
Tracker::Enable(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) ::proxy::Tracker::Disable(entry);
if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry);
} else {
status = hsa_amd_memory_async_copy_fn(dst, dst_agent, src, src_agent, size, num_dep_signals,
dep_signals, completion_signal);
@@ -630,10 +586,10 @@ hsa_status_t hsa_amd_memory_async_copy_rect_interceptor(
hsa_status_t status = HSA_STATUS_SUCCESS;
if (hsa_support::async_copy_callback_enabled) {
trace_entry_t* entry = trace_buffer->GetEntry();
::proxy::Tracker::Enable(COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
Tracker::Enable(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) ::proxy::Tracker::Disable(entry);
if (status != HSA_STATUS_SUCCESS) Tracker::Disable(entry);
} else {
status =
hsa_amd_memory_async_copy_rect_fn(dst, dst_offset, src, src_offset, range, copy_agent, dir,
@@ -1355,7 +1311,6 @@ PUBLIC_API void roctracer_unload() {
roctracer::act_journal = NULL;
}
roctracer::close_output_file(roctracer::kernel_file_handle);
ONLOAD_TRACE_END();
}
projects/roctracer/src/proxy/tracker.h → projects/roctracer/src/core/tracker.h
+8 -8
Wyświetl plik
@@ -18,8 +18,8 @@
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
#ifndef SRC_PROXY_TRACKER_H_
#define SRC_PROXY_TRACKER_H_
#ifndef SRC_CORE_TRACKER_H_
#define SRC_CORE_TRACKER_H_
#include <amd_hsa_signal.h>
#include <assert.h>
@@ -33,10 +33,10 @@
#include "util/logger.h"
#include "core/trace_buffer.h"
namespace proxy {
namespace roctracer {
class Tracker {
public:
typedef util::HsaRsrcFactory::timestamp_t timestamp_t;
typedef ::util::HsaRsrcFactory::timestamp_t timestamp_t;
typedef roctracer::trace_entry_t entry_t;
typedef roctracer::entry_type_t entry_type_t;
@@ -44,7 +44,7 @@ class Tracker {
inline static void Enable(entry_type_t type, const hsa_agent_t& agent, const hsa_signal_t& signal,
entry_t* entry) {
hsa_status_t status = HSA_STATUS_ERROR;
util::HsaRsrcFactory* hsa_rsrc = &(util::HsaRsrcFactory::Instance());
::util::HsaRsrcFactory* hsa_rsrc = &(::util::HsaRsrcFactory::Instance());
// Creating a new tracker entry
entry->type = type;
@@ -72,7 +72,7 @@ class Tracker {
// Entry completion
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());
::util::HsaRsrcFactory* hsa_rsrc = &(::util::HsaRsrcFactory::Instance());
if (entry->type == roctracer::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);
@@ -125,6 +125,6 @@ class Tracker {
}
};
} // namespace proxy
} // namespace roctracer
#endif // SRC_PROXY_TRACKER_H_
#endif // SRC_CORE_TRACKER_H_
projects/roctracer/src/proxy/hsa_proxy_queue.h
-65
Wyświetl plik
@@ -1,65 +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_HSA_PROXY_QUEUE_H
#define _SRC_CORE_HSA_PROXY_QUEUE_H
#include <hsa.h>
#include <atomic>
#include <map>
#include <mutex>
#include "proxy/proxy_queue.h"
#include "util/exception.h"
namespace rocprofiler {
extern decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
extern decltype(hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_fn;
extern decltype(hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn;
class HsaProxyQueue : public ProxyQueue {
public:
hsa_status_t SetInterceptCB(on_submit_cb_t on_submit_cb, void* data) {
return hsa_amd_queue_intercept_register_fn(queue_, on_submit_cb, data);
}
void Submit(const packet_t* packet) {
EXC_RAISING(HSA_STATUS_ERROR, "HsaProxyQueue::Submit() is not supported");
}
private:
hsa_status_t Init(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
void* data, uint32_t private_segment_size, uint32_t group_segment_size,
hsa_queue_t** queue) {
const auto status = hsa_amd_queue_intercept_create_fn(
agent, size, type, callback, data, private_segment_size, group_segment_size, &queue_);
*queue = queue_;
return status;
}
hsa_status_t Cleanup() const { return hsa_queue_destroy_fn(queue_); }
hsa_queue_t* queue_;
};
} // namespace rocprofiler
#endif // _SRC_CORE_HSA_PROXY_QUEUE_H
projects/roctracer/src/proxy/hsa_queue.h
-42
Wyświetl plik
@@ -1,42 +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_HSA_QUEUE_H
#define _SRC_CORE_HSA_QUEUE_H
#include <atomic>
#include "proxy/queue.h"
namespace rocprofiler {
class HsaQueue : public Queue {
public:
HsaQueue(const util::AgentInfo* agent_info, hsa_queue_t* queue) : queue_(queue) {}
void Submit(const packet_t* packet) { util::HsaRsrcFactory::Instance().Submit(queue_, packet); }
private:
hsa_queue_t* queue_;
};
} // namespace rocprofiler
#endif // _SRC_CORE_HSA_QUEUE_H
projects/roctracer/src/proxy/intercept_queue.cpp
-40
Wyświetl plik
@@ -1,40 +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. */
#include "proxy/intercept_queue.h"
namespace rocprofiler {
void InterceptQueue::HsaIntercept(HsaApiTable* table) {
table->core_->hsa_queue_create_fn = rocprofiler::InterceptQueue::QueueCreate;
table->core_->hsa_queue_destroy_fn = rocprofiler::InterceptQueue::QueueDestroy;
}
InterceptQueue::mutex_t InterceptQueue::mutex_;
// rocprofiler_callback_t InterceptQueue::dispatch_callback_ = NULL;
// InterceptQueue::queue_callback_t InterceptQueue::create_callback_ = NULL;
// InterceptQueue::queue_callback_t InterceptQueue::destroy_callback_ = NULL;
// void* InterceptQueue::callback_data_ = NULL;
InterceptQueue::obj_map_t* InterceptQueue::obj_map_ = NULL;
const char* InterceptQueue::kernel_none_ = "";
bool InterceptQueue::in_create_call_ = false;
InterceptQueue::queue_id_t InterceptQueue::current_queue_id = 0;
bool InterceptQueue::is_enabled = false;
} // namespace rocprofiler
projects/roctracer/src/proxy/intercept_queue.h
-304
Wyświetl plik
@@ -1,304 +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_INTERCEPT_QUEUE_H
#define _SRC_CORE_INTERCEPT_QUEUE_H
#include <amd_hsa_kernel_code.h>
#include <cxxabi.h>
#include <dlfcn.h>
#include <sys/syscall.h>
#include <atomic>
#include <iostream>
#include <map>
#include <mutex>
#include "core/trace_buffer.h"
#include "proxy/tracker.h"
#include "proxy/proxy_queue.h"
#include "util/hsa_rsrc_factory.h"
#include "util/exception.h"
namespace roctracer {
extern TraceBuffer<trace_entry_t>* trace_buffer;
}
namespace rocprofiler {
extern decltype(hsa_queue_create)* hsa_queue_create_fn;
extern decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
class InterceptQueue {
public:
typedef std::recursive_mutex mutex_t;
typedef std::map<uint64_t, InterceptQueue*> obj_map_t;
typedef hsa_status_t (*queue_callback_t)(hsa_queue_t*, void* data);
typedef void (*queue_event_callback_t)(hsa_status_t status, hsa_queue_t* queue, void* arg);
typedef uint32_t queue_id_t;
static void HsaIntercept(HsaApiTable* table);
static hsa_status_t InterceptQueueCreate(
hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data,
uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue,
const bool& tracker_on) {
std::lock_guard<mutex_t> lck(mutex_);
hsa_status_t status = HSA_STATUS_ERROR;
if (in_create_call_) EXC_ABORT(status, "recursive InterceptQueueCreate()");
in_create_call_ = true;
ProxyQueue* proxy =
ProxyQueue::Create(agent, size, type, queue_event_callback, data, private_segment_size,
group_segment_size, queue, &status);
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "ProxyQueue::Create()");
status = util::HsaRsrcFactory::HsaApi()->hsa_amd_profiling_set_profiler_enabled(*queue, true);
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_profiling_set_profiler_enabled()");
if (!obj_map_) obj_map_ = new obj_map_t;
InterceptQueue* obj = new InterceptQueue(agent, *queue, proxy);
obj->queue_event_callback_ = callback;
obj->queue_id = current_queue_id;
(*obj_map_)[(uint64_t)(*queue)] = obj;
status = (is_enabled) ? proxy->SetInterceptCB(OnSubmitCB, obj)
: proxy->SetInterceptCB(OnSubmitCB_dummy, obj);
#if 0
if (create_callback_ != NULL) {
status = create_callback_(*queue, callback_data_);
}
#endif
++current_queue_id;
in_create_call_ = false;
return status;
}
static hsa_status_t QueueCreate(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source,
void* data),
void* data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t** queue) {
return InterceptQueueCreate(agent, size, type, callback, data, private_segment_size,
group_segment_size, queue, false);
}
static hsa_status_t QueueCreateTracked(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source,
void* data),
void* data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t** queue) {
return InterceptQueueCreate(agent, size, type, callback, data, private_segment_size,
group_segment_size, queue, true);
}
static hsa_status_t QueueDestroy(hsa_queue_t* queue) {
std::lock_guard<mutex_t> lck(mutex_);
hsa_status_t status = HSA_STATUS_SUCCESS;
#if 0
if (destroy_callback_ != NULL) {
status = destroy_callback_(queue, callback_data_);
}
#endif
if (status == HSA_STATUS_SUCCESS) {
status = DelObj(queue);
}
return status;
}
static void OnSubmitCB_dummy(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);
// Submitting the original packets if profiling was not enabled
if (writer != NULL) {
writer(packets_arr, count);
} else {
InterceptQueue* obj = reinterpret_cast<InterceptQueue*>(data);
Queue* proxy = obj->proxy_;
proxy->Submit(packets_arr, count);
}
}
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);
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];
// Checking for dispatch packet type
if (GetHeaderType(packet) == HSA_PACKET_TYPE_KERNEL_DISPATCH) {
const hsa_kernel_dispatch_packet_t* dispatch_packet =
reinterpret_cast<const hsa_kernel_dispatch_packet_t*>(packet);
// Prepareing dispatch callback data
const hsa_signal_t completion_signal = dispatch_packet->completion_signal;
const amd_kernel_code_t* kernel_code = GetKernelCode(dispatch_packet);
const uint64_t kernel_symbol = kernel_code->runtime_loader_kernel_symbol;
const char* kernel_name = GetKernelName(kernel_symbol);
// Adding kernel timing tracker
::proxy::Tracker::entry_t* entry = roctracer::trace_buffer->GetEntry();
entry->kernel.tid = syscall(__NR_gettid);
entry->kernel.name = kernel_name;
::proxy::Tracker::Enable(roctracer::KERNEL_ENTRY_TYPE, obj->agent_info_->dev_id,
completion_signal, entry);
const_cast<hsa_kernel_dispatch_packet_t*>(dispatch_packet)->completion_signal =
entry->signal;
}
}
// Submitting the original packets if profiling was not enabled
if (writer != NULL) {
writer(packets_arr, count);
} else {
proxy->Submit(packets_arr, count);
}
}
#if 0
static void SetCallbacks(rocprofiler_callback_t dispatch_callback,
queue_callback_t create_callback,
queue_callback_t destroy_callback,
void* data)
{
std::lock_guard<mutex_t> lck(mutex_);
callback_data_ = data;
dispatch_callback_ = dispatch_callback;
create_callback_ = create_callback;
destroy_callback_ = destroy_callback;
}
#endif
static void Enable(bool val) { is_enabled = val; }
private:
static void queue_event_callback(hsa_status_t status, hsa_queue_t* queue, void* arg) {
if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "queue error handling is not supported");
InterceptQueue* obj = GetObj(queue);
if (obj->queue_event_callback_) obj->queue_event_callback_(status, obj->queue_, arg);
}
static hsa_packet_type_t GetHeaderType(const packet_t* packet) {
const packet_word_t* header = reinterpret_cast<const packet_word_t*>(packet);
return static_cast<hsa_packet_type_t>((*header >> HSA_PACKET_HEADER_TYPE) & header_type_mask);
}
static const amd_kernel_code_t* GetKernelCode(
const hsa_kernel_dispatch_packet_t* dispatch_packet) {
const amd_kernel_code_t* kernel_code = NULL;
hsa_status_t status =
util::HsaRsrcFactory::Instance().LoaderApi()->hsa_ven_amd_loader_query_host_address(
reinterpret_cast<const void*>(dispatch_packet->kernel_object),
reinterpret_cast<const void**>(&kernel_code));
if (HSA_STATUS_SUCCESS != status) {
kernel_code = reinterpret_cast<amd_kernel_code_t*>(dispatch_packet->kernel_object);
}
return kernel_code;
}
static const char* GetKernelName(const uint64_t kernel_symbol) {
amd_runtime_loader_debug_info_t* dbg_info =
reinterpret_cast<amd_runtime_loader_debug_info_t*>(kernel_symbol);
const char* kernel_name = (dbg_info != NULL) ? dbg_info->kernel_name : NULL;
return (kernel_name != NULL) ? strdup(kernel_name) : strdup(kernel_none_);
#if 0
// Kernel name is mangled name
// apply __cxa_demangle() to demangle it
const char* funcname = NULL;
if (kernel_name != NULL) {
size_t funcnamesize = 0;
int status;
const char* ret = abi::__cxa_demangle(kernel_name, NULL, &funcnamesize, &status);
funcname = (ret != 0) ? ret : strdup(kernel_name);
}
if (funcname == NULL) funcname = strdup(kernel_none_);
return funcname;
#endif
}
// method to get an intercept queue object
static InterceptQueue* GetObj(const hsa_queue_t* queue) {
std::lock_guard<mutex_t> lck(mutex_);
InterceptQueue* obj = NULL;
obj_map_t::const_iterator it = obj_map_->find((uint64_t)queue);
if (it != obj_map_->end()) {
obj = it->second;
assert(queue == obj->queue_);
}
return obj;
}
// method to delete an intercept queue object
static hsa_status_t DelObj(const hsa_queue_t* queue) {
std::lock_guard<mutex_t> lck(mutex_);
hsa_status_t status = HSA_STATUS_ERROR;
obj_map_t::const_iterator it = obj_map_->find((uint64_t)queue);
if (it != obj_map_->end()) {
const InterceptQueue* obj = it->second;
assert(queue == obj->queue_);
delete obj;
obj_map_->erase(it);
status = HSA_STATUS_SUCCESS;
;
}
return status;
}
InterceptQueue(const hsa_agent_t& agent, hsa_queue_t* const queue, ProxyQueue* proxy)
: queue_(queue), proxy_(proxy) {
agent_info_ = util::HsaRsrcFactory::Instance().GetAgentInfo(agent);
queue_event_callback_ = NULL;
}
~InterceptQueue() { ProxyQueue::Destroy(proxy_); }
static bool is_enabled;
static mutex_t mutex_;
static const packet_word_t header_type_mask = (1ul << HSA_PACKET_HEADER_WIDTH_TYPE) - 1;
#if 0
static queue_callback_t create_callback_;
static queue_callback_t destroy_callback_;
static void* callback_data_;
#endif
static obj_map_t* obj_map_;
static const char* kernel_none_;
static bool in_create_call_;
static queue_id_t current_queue_id;
hsa_queue_t* const queue_;
ProxyQueue* const proxy_;
const util::AgentInfo* agent_info_;
queue_event_callback_t queue_event_callback_;
queue_id_t queue_id;
};
} // namespace rocprofiler
#endif // _SRC_CORE_INTERCEPT_QUEUE_H
projects/roctracer/src/proxy/proxy_queue.cpp
-61
Wyświetl plik
@@ -1,61 +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. */
#include "proxy/proxy_queue.h"
#include "proxy/hsa_proxy_queue.h"
#include "proxy/simple_proxy_queue.h"
namespace rocprofiler {
void ProxyQueue::HsaIntercept(HsaApiTable* table) {
if (rocp_type_) SimpleProxyQueue::HsaIntercept(table);
}
ProxyQueue* ProxyQueue::Create(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source,
void* data),
void* data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t** queue,
hsa_status_t* status) {
hsa_status_t suc = HSA_STATUS_ERROR;
ProxyQueue* instance =
(rocp_type_) ? (ProxyQueue*)new SimpleProxyQueue() : (ProxyQueue*)new HsaProxyQueue();
if (instance != NULL) {
suc = instance->Init(agent, size, type, callback, data, private_segment_size,
group_segment_size, queue);
if (suc != HSA_STATUS_SUCCESS) {
delete instance;
instance = NULL;
}
}
*status = suc;
assert(*status == HSA_STATUS_SUCCESS);
return instance;
}
hsa_status_t ProxyQueue::Destroy(const ProxyQueue* obj) {
assert(obj != NULL);
auto suc = obj->Cleanup();
delete obj;
return suc;
}
bool ProxyQueue::rocp_type_ = false;
} // namespace rocprofiler
projects/roctracer/src/proxy/proxy_queue.h
-75
Wyświetl plik
@@ -1,75 +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_PROXY_QUEUE_H
#define _SRC_CORE_PROXY_QUEUE_H
#include <hsa.h>
#include <hsa_api_trace.h>
#include <atomic>
#include <map>
#include <mutex>
#include "proxy/queue.h"
#include "proxy/types.h"
struct HsaApiTable;
namespace rocprofiler {
typedef void (*hsa_amd_queue_intercept_packet_writer)(const void* packets, uint64_t count);
typedef void (*on_submit_cb_t)(const void* packet, uint64_t count, uint64_t que_idx, void* data,
hsa_amd_queue_intercept_packet_writer writer);
class ProxyQueue : public Queue {
public:
static void InitFactory() {
const char* type = getenv("ROCP_PROXY_QUEUE");
if (type != NULL) {
if (strncmp(type, "rocp", 4) == 0) rocp_type_ = true;
}
}
static void HsaIntercept(HsaApiTable* table);
static ProxyQueue* Create(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
void* data, uint32_t private_segment_size, uint32_t group_segment_size,
hsa_queue_t** queue, hsa_status_t* status);
static hsa_status_t Destroy(const ProxyQueue* obj);
virtual hsa_status_t Init(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
void* data, uint32_t private_segment_size, uint32_t group_segment_size,
hsa_queue_t** queue) = 0;
virtual hsa_status_t Cleanup() const = 0;
virtual hsa_status_t SetInterceptCB(on_submit_cb_t on_submit_cb, void* data) = 0;
virtual void Submit(const packet_t* packet) = 0;
protected:
virtual ~ProxyQueue(){};
private:
static bool rocp_type_;
};
} // namespace rocprofiler
#endif // _SRC_CORE_PROXY_QUEUE_H
projects/roctracer/src/proxy/queue.h
-40
Wyświetl plik
@@ -1,40 +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_QUEUE_H
#define _SRC_CORE_QUEUE_H
#include "proxy/types.h"
namespace rocprofiler {
class Queue {
public:
Queue() {}
virtual ~Queue() {}
virtual void Submit(const packet_t* packet) = 0;
virtual void Submit(const packet_t* packet, const size_t& count) {
for (const packet_t* p = packet; p < packet + count; ++p) Submit(p);
}
};
} // namespace rocprofiler
#endif // _SRC_CORE_QUEUE_H
projects/roctracer/src/proxy/simple_proxy_queue.cpp
-44
Wyświetl plik
@@ -1,44 +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. */
#include "proxy/simple_proxy_queue.h"
namespace rocprofiler {
void SimpleProxyQueue::HsaIntercept(HsaApiTable* table) {
table->core_->hsa_signal_store_relaxed_fn = rocprofiler::SimpleProxyQueue::SignalStore;
table->core_->hsa_signal_store_screlease_fn = rocprofiler::SimpleProxyQueue::SignalStore;
table->core_->hsa_queue_load_write_index_relaxed_fn =
rocprofiler::SimpleProxyQueue::GetQueueIndex;
table->core_->hsa_queue_store_write_index_relaxed_fn =
rocprofiler::SimpleProxyQueue::SetQueueIndex;
table->core_->hsa_queue_load_read_index_relaxed_fn =
rocprofiler::SimpleProxyQueue::GetSubmitIndex;
table->core_->hsa_queue_load_write_index_scacquire_fn =
rocprofiler::SimpleProxyQueue::GetQueueIndex;
table->core_->hsa_queue_store_write_index_screlease_fn =
rocprofiler::SimpleProxyQueue::SetQueueIndex;
table->core_->hsa_queue_load_read_index_scacquire_fn =
rocprofiler::SimpleProxyQueue::GetSubmitIndex;
}
SimpleProxyQueue::queue_map_t* SimpleProxyQueue::queue_map_ = NULL;
} // namespace rocprofiler
projects/roctracer/src/proxy/simple_proxy_queue.h
-259
Wyświetl plik
@@ -1,259 +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_SIMPLE_PROXY_QUEUE_H
#define _SRC_CORE_SIMPLE_PROXY_QUEUE_H
#include <hsa.h>
#include <atomic>
#include <map>
#include <mutex>
#include "proxy/proxy_queue.h"
#include "util/hsa_rsrc_factory.h"
#ifndef ROCP_PROXY_LOCK
#define ROCP_PROXY_LOCK 1
#endif
namespace rocprofiler {
extern decltype(hsa_queue_create)* hsa_queue_create_fn;
extern decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
extern decltype(hsa_signal_store_relaxed)* hsa_signal_store_relaxed_fn;
extern decltype(hsa_signal_store_relaxed)* hsa_signal_store_screlease_fn;
extern decltype(hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed_fn;
extern decltype(hsa_queue_store_write_index_relaxed)* hsa_queue_store_write_index_relaxed_fn;
extern decltype(hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_fn;
extern decltype(hsa_queue_load_write_index_scacquire)* hsa_queue_load_write_index_scacquire_fn;
extern decltype(hsa_queue_store_write_index_screlease)* hsa_queue_store_write_index_screlease_fn;
extern decltype(hsa_queue_load_read_index_scacquire)* hsa_queue_load_read_index_scacquire_fn;
typedef decltype(hsa_signal_t::handle) signal_handle_t;
class SimpleProxyQueue : public ProxyQueue {
public:
static void HsaIntercept(HsaApiTable* table);
static void SignalStore(hsa_signal_t signal, hsa_signal_value_t que_idx) {
auto it = queue_map_->find(signal.handle);
if (it != queue_map_->end()) {
SimpleProxyQueue* instance = it->second;
instance->mutex_lock();
const uint64_t begin = instance->submit_index_;
const uint64_t end = que_idx + 1;
instance->submit_index_ = end;
instance->mutex_unlock();
for (uint64_t j = begin; j < end; ++j) {
// Submited packet
const uint32_t idx = j & instance->queue_mask_;
packet_t* packet = reinterpret_cast<packet_t*>(instance->queue_->base_address) + idx;
if (instance->on_submit_cb_ != NULL)
instance->on_submit_cb_(packet, 1, j, instance->on_submit_cb_data_, NULL);
else
instance->Submit(packet);
}
} else {
hsa_signal_store_relaxed_fn(signal, que_idx);
}
}
static uint64_t GetSubmitIndex(const hsa_queue_t* queue) {
uint64_t index = 0;
auto it = queue_map_->find(queue->doorbell_signal.handle);
if (it != queue_map_->end()) {
SimpleProxyQueue* instance = it->second;
index = instance->submit_index_;
} else {
index = hsa_queue_load_read_index_relaxed_fn(queue);
}
return index;
}
static uint64_t GetQueueIndex(const hsa_queue_t* queue) {
uint64_t index = 0;
auto it = queue_map_->find(queue->doorbell_signal.handle);
if (it != queue_map_->end()) {
SimpleProxyQueue* instance = it->second;
instance->mutex_lock();
index = instance->queue_index_;
} else {
index = hsa_queue_load_write_index_relaxed_fn(queue);
}
return index;
}
static void SetQueueIndex(const hsa_queue_t* queue, uint64_t value) {
auto it = queue_map_->find(queue->doorbell_signal.handle);
if (it != queue_map_->end()) {
SimpleProxyQueue* instance = it->second;
instance->queue_index_ = value;
instance->mutex_unlock();
} else {
hsa_queue_store_write_index_relaxed_fn(queue, value);
}
}
hsa_status_t SetInterceptCB(on_submit_cb_t on_submit_cb, void* data) {
on_submit_cb_ = on_submit_cb;
on_submit_cb_data_ = data;
return HSA_STATUS_SUCCESS;
}
void Submit(const packet_t* packet) {
// Compute the write index of queue
const uint64_t que_idx = hsa_queue_load_write_index_relaxed_fn(queue_);
// Waiting until there is a free space in the queue
while (que_idx >= (hsa_queue_load_read_index_relaxed_fn(queue_) + size_))
;
// Increment the write index
hsa_queue_store_write_index_relaxed_fn(queue_, que_idx + 1);
const uint32_t mask = queue_->size - 1;
const uint32_t idx = que_idx & mask;
// Copy packet to the queue
const packet_word_t* src = reinterpret_cast<const packet_word_t*>(packet);
packet_word_t* dst = reinterpret_cast<packet_word_t*>(base_address_ + idx);
for (unsigned i = 1; i < sizeof(packet_t) / sizeof(packet_word_t); ++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 to submit the packet
hsa_signal_store_relaxed_fn(doorbell_signal_, que_idx);
}
SimpleProxyQueue()
: agent_info_(NULL),
queue_(NULL),
base_address_(NULL),
doorbell_signal_({}),
queue_index_(0),
queue_mask_(0),
submit_index_(0),
on_submit_cb_(NULL),
on_submit_cb_data_(NULL) {
printf("ROCProfiler: SimpleProxyQueue is enabled\n");
fflush(stdout);
}
~SimpleProxyQueue() {}
private:
typedef std::map<signal_handle_t, SimpleProxyQueue*> queue_map_t;
hsa_status_t Init(hsa_agent_t agent, uint32_t size, hsa_queue_type32_t type,
void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data),
void* data, uint32_t private_segment_size, uint32_t group_segment_size,
hsa_queue_t** queue) {
size_ = size;
auto status = Init(agent, size);
*queue = queue_;
return status;
}
hsa_status_t Init(hsa_agent_t agent, uint32_t size) {
hsa_status_t status = HSA_STATUS_ERROR;
agent_info_ = util::HsaRsrcFactory::Instance().GetAgentInfo(agent);
if (agent_info_ != NULL) {
if (agent_info_->dev_type == HSA_DEVICE_TYPE_GPU) {
status = hsa_queue_create_fn(agent, size, HSA_QUEUE_TYPE_MULTI, NULL, NULL, UINT32_MAX,
UINT32_MAX, &queue_);
if (status == HSA_STATUS_SUCCESS) {
base_address_ = reinterpret_cast<packet_t*>(queue_->base_address);
doorbell_signal_ = queue_->doorbell_signal;
data_array_ = calloc(size + 1, sizeof(packet_t));
uintptr_t addr = (uintptr_t)data_array_;
queue_->base_address = (void*)((addr + align_mask_) & ~align_mask_);
status = hsa_signal_create(1, 0, NULL, &(queue_->doorbell_signal));
if (status != HSA_STATUS_SUCCESS) abort();
queue_mask_ = size - 1;
if (queue_map_ == NULL) queue_map_ = new queue_map_t;
(*queue_map_)[queue_->doorbell_signal.handle] = this;
} else
abort();
}
}
if (status != HSA_STATUS_SUCCESS) abort();
return status;
}
hsa_status_t Cleanup() const {
hsa_status_t status = HSA_STATUS_ERROR;
hsa_signal_t queue_signal = queue_->doorbell_signal;
// Destroy original HSA queue
queue_->base_address = base_address_;
queue_->doorbell_signal = doorbell_signal_;
status = hsa_queue_destroy_fn(queue_);
if (status != HSA_STATUS_SUCCESS) abort();
// Destroy overloaded virtual queue data and signal
free(data_array_);
status = hsa_signal_destroy(queue_signal);
if (status != HSA_STATUS_SUCCESS) abort();
return status;
}
void mutex_lock() {
#if ROCP_PROXY_LOCK
mutex_.lock();
#endif
}
void mutex_unlock() {
#if ROCP_PROXY_LOCK
mutex_.unlock();
#endif
}
uint32_t size_;
static queue_map_t* queue_map_;
const util::AgentInfo* agent_info_;
hsa_queue_t* queue_;
static const uintptr_t align_mask_ = sizeof(packet_t) - 1;
packet_t* base_address_;
hsa_signal_t doorbell_signal_;
uint64_t queue_index_;
uint64_t queue_mask_;
uint64_t submit_index_;
std::mutex mutex_;
on_submit_cb_t on_submit_cb_;
void* on_submit_cb_data_;
void* data_array_;
};
} // namespace rocprofiler
#endif // _SRC_CORE_SIMPLE_PROXY_QUEUE_H
projects/roctracer/src/proxy/types.h
-49
Wyświetl plik
@@ -1,49 +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_TYPES_H_
#define SRC_CORE_TYPES_H_
#include <iostream>
#include <hsa_ven_amd_aqlprofile.h>
namespace rocprofiler {
typedef hsa_ven_amd_aqlprofile_pfn_t pfn_t;
typedef hsa_ven_amd_aqlprofile_event_t event_t;
typedef hsa_ven_amd_aqlprofile_parameter_t parameter_t;
typedef hsa_ven_amd_aqlprofile_profile_t profile_t;
typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
typedef uint32_t packet_word_t;
typedef uint64_t timestamp_t;
inline std::ostream& operator<<(std::ostream& out, const event_t& event) {
out << "[block_name(" << event.block_name << "). block_index(" << event.block_index
<< "). counter_id(" << event.counter_id << ")]";
return out;
}
inline std::ostream& operator<<(std::ostream& out, const parameter_t& parameter) {
out << "[parameter_name(" << parameter.parameter_name << "). value(" << parameter.value << ")]";
return out;
}
} // namespace rocprofiler
#endif // SRC_CORE_TYPES_H_