diff --git a/projects/roctracer/src/CMakeLists.txt b/projects/roctracer/src/CMakeLists.txt index f2068d65a4..6f596085f5 100644 --- a/projects/roctracer/src/CMakeLists.txt +++ b/projects/roctracer/src/CMakeLists.txt @@ -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} ) diff --git a/projects/roctracer/src/core/roctracer.cpp b/projects/roctracer/src/core/roctracer.cpp index 186a318b47..57c679c95d 100644 --- a/projects/roctracer/src/core/roctracer.cpp +++ b/projects/roctracer/src/core/roctracer.cpp @@ -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 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::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_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(); } diff --git a/projects/roctracer/src/proxy/tracker.h b/projects/roctracer/src/core/tracker.h similarity index 93% rename from projects/roctracer/src/proxy/tracker.h rename to projects/roctracer/src/core/tracker.h index 1bf98ce32c..a005f213d8 100644 --- a/projects/roctracer/src/proxy/tracker.h +++ b/projects/roctracer/src/core/tracker.h @@ -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 #include @@ -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_ diff --git a/projects/roctracer/src/proxy/hsa_proxy_queue.h b/projects/roctracer/src/proxy/hsa_proxy_queue.h deleted file mode 100644 index 4093e6807c..0000000000 --- a/projects/roctracer/src/proxy/hsa_proxy_queue.h +++ /dev/null @@ -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 -#include -#include -#include - -#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 diff --git a/projects/roctracer/src/proxy/hsa_queue.h b/projects/roctracer/src/proxy/hsa_queue.h deleted file mode 100644 index d8086615e6..0000000000 --- a/projects/roctracer/src/proxy/hsa_queue.h +++ /dev/null @@ -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 - -#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 diff --git a/projects/roctracer/src/proxy/intercept_queue.cpp b/projects/roctracer/src/proxy/intercept_queue.cpp deleted file mode 100644 index f67dfc5ea0..0000000000 --- a/projects/roctracer/src/proxy/intercept_queue.cpp +++ /dev/null @@ -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 diff --git a/projects/roctracer/src/proxy/intercept_queue.h b/projects/roctracer/src/proxy/intercept_queue.h deleted file mode 100644 index 0cc49de1e2..0000000000 --- a/projects/roctracer/src/proxy/intercept_queue.h +++ /dev/null @@ -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 -#include -#include -#include - -#include -#include -#include -#include - -#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_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 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 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 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(in_packets); - - // Submitting the original packets if profiling was not enabled - if (writer != NULL) { - writer(packets_arr, count); - } else { - InterceptQueue* obj = reinterpret_cast(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(in_packets); - InterceptQueue* obj = reinterpret_cast(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(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(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 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(packet); - return static_cast((*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(dispatch_packet->kernel_object), - reinterpret_cast(&kernel_code)); - if (HSA_STATUS_SUCCESS != status) { - kernel_code = reinterpret_cast(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(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 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 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 diff --git a/projects/roctracer/src/proxy/proxy_queue.cpp b/projects/roctracer/src/proxy/proxy_queue.cpp deleted file mode 100644 index 3c14b530fa..0000000000 --- a/projects/roctracer/src/proxy/proxy_queue.cpp +++ /dev/null @@ -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 diff --git a/projects/roctracer/src/proxy/proxy_queue.h b/projects/roctracer/src/proxy/proxy_queue.h deleted file mode 100644 index 8db8965504..0000000000 --- a/projects/roctracer/src/proxy/proxy_queue.h +++ /dev/null @@ -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 -#include -#include -#include -#include - -#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 diff --git a/projects/roctracer/src/proxy/queue.h b/projects/roctracer/src/proxy/queue.h deleted file mode 100644 index 60a7c3581e..0000000000 --- a/projects/roctracer/src/proxy/queue.h +++ /dev/null @@ -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 diff --git a/projects/roctracer/src/proxy/simple_proxy_queue.cpp b/projects/roctracer/src/proxy/simple_proxy_queue.cpp deleted file mode 100644 index a24259eca9..0000000000 --- a/projects/roctracer/src/proxy/simple_proxy_queue.cpp +++ /dev/null @@ -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 diff --git a/projects/roctracer/src/proxy/simple_proxy_queue.h b/projects/roctracer/src/proxy/simple_proxy_queue.h deleted file mode 100644 index 4c588b22db..0000000000 --- a/projects/roctracer/src/proxy/simple_proxy_queue.h +++ /dev/null @@ -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 -#include -#include -#include - -#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(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(packet); - packet_word_t* dst = reinterpret_cast(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* header_atomic_ptr = - reinterpret_cast*>(&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 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(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 diff --git a/projects/roctracer/src/proxy/types.h b/projects/roctracer/src/proxy/types.h deleted file mode 100644 index 50231a55f3..0000000000 --- a/projects/roctracer/src/proxy/types.h +++ /dev/null @@ -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 - -#include - -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_