From eae9a63570f9e1b5dceabc12b7402c795aff481d Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Sat, 29 Oct 2022 01:01:50 +0000 Subject: [PATCH 01/41] Reimplemented tracing in a session/client aware manner --- .../onnxruntime/core/common/profiler_common.h | 41 +- onnxruntime/core/common/profiler_common.cc | 27 +- .../core/providers/rocm/RoctracerLogger.cc | 348 --------- .../core/providers/rocm/RoctracerLogger.h | 170 ---- .../core/providers/rocm/roctracer_manager.cc | 735 ++++++++++++++++++ .../core/providers/rocm/roctracer_manager.h | 271 +++++++ 6 files changed, 1065 insertions(+), 527 deletions(-) delete mode 100644 onnxruntime/core/providers/rocm/RoctracerLogger.cc delete mode 100644 onnxruntime/core/providers/rocm/RoctracerLogger.h create mode 100644 onnxruntime/core/providers/rocm/roctracer_manager.cc create mode 100644 onnxruntime/core/providers/rocm/roctracer_manager.h diff --git a/include/onnxruntime/core/common/profiler_common.h b/include/onnxruntime/core/common/profiler_common.h index a2f31442b16b3..58e20fdb878a8 100644 --- a/include/onnxruntime/core/common/profiler_common.h +++ b/include/onnxruntime/core/common/profiler_common.h @@ -27,19 +27,44 @@ static constexpr const char* event_categor_names_[EVENT_CATEGORY_MAX] = { // Timing record for all events. struct EventRecord { + EventRecord() = default; EventRecord(EventCategory category, int process_id, int thread_id, - std::string event_name, + std::string&& event_name, long long time_stamp, long long duration, - std::unordered_map&& event_args) : cat(category), - pid(process_id), - tid(thread_id), - name(std::move(event_name)), - ts(time_stamp), - dur(duration), - args(event_args) {} + std::unordered_map&& event_args) + : cat(category), + pid(process_id), + tid(thread_id), + name(std::move(event_name)), + ts(time_stamp), + dur(duration), + args(std::move(event_args)) {} + + EventRecord(EventCategory category, + int process_id, + int thread_id, + const std::string& event_name, + long long time_stamp, + long long duration, + const std::unordered_map& event_args) + : cat(category), + pid(process_id), + tid(thread_id), + name(event_name), + ts(time_stamp), + dur(duration), + args(event_args) {} + + EventRecord(EventRecord&& other) + : cat(other.cat), pid(other.pid), tid(other.tid), name(std::move(other.name)), + ts(other.ts), dur(other.dur), args(std::move(other.args)) {} + + EventRecord& operator = (const EventRecord& other); + EventRecord& operator = (EventRecord&& other); + EventCategory cat; int pid; int tid; diff --git a/onnxruntime/core/common/profiler_common.cc b/onnxruntime/core/common/profiler_common.cc index ee14533b2ad93..25d64970a287c 100644 --- a/onnxruntime/core/common/profiler_common.cc +++ b/onnxruntime/core/common/profiler_common.cc @@ -48,6 +48,31 @@ std::string demangle(const std::string& name) { return demangle(name.c_str()); } +EventRecord& EventRecord::operator = (const EventRecord& other) { + if (&other == this) { + return *this; + } + cat = other.cat; + pid = other.pid; + tid = other.tid; + name = other.name; + ts = other.ts; + dur = other.dur; + args = other.args; + return *this; +} + +EventRecord& EventRecord::operator = (EventRecord&& other) { + if (&other == this) { + return *this; + } + cat = other.cat; + pid = other.pid; + tid = other.tid; + std::swap(name, other.name); + dur = other.dur; + std::swap(args, other.args); +} + } // namespace profiling } // namespace onnxruntime - diff --git a/onnxruntime/core/providers/rocm/RoctracerLogger.cc b/onnxruntime/core/providers/rocm/RoctracerLogger.cc deleted file mode 100644 index ab48ad71a2162..0000000000000 --- a/onnxruntime/core/providers/rocm/RoctracerLogger.cc +++ /dev/null @@ -1,348 +0,0 @@ -#include "RoctracerLogger.h" - -#include -#include -#include - -#include "ThreadUtil.h" - -typedef uint64_t timestamp_t; - -static timestamp_t timespec_to_ns(const timespec& time) { - return ((timestamp_t)time.tv_sec * 1000000000) + time.tv_nsec; - } - -//using namespace std::chrono; - -RoctracerLogger& RoctracerLogger::singleton() { - static RoctracerLogger instance; - return instance; -} - -RoctracerLogger::RoctracerLogger() { - gpuTraceBuffers_ = std::make_unique>(); -} - -RoctracerLogger::~RoctracerLogger() { - stopLogging(); - endTracing(); -} - -namespace { - thread_local std::deque t_externalIds[RoctracerLogger::CorrelationDomain::size]; -} - -void RoctracerLogger::pushCorrelationID(uint64_t id, CorrelationDomain type) { - if (!singleton().externalCorrelationEnabled_) { - return; - } - t_externalIds[type].push_back(id); -} - -void RoctracerLogger::popCorrelationID(CorrelationDomain type) { - if (!singleton().externalCorrelationEnabled_) { - return; - } - t_externalIds[type].pop_back(); -} - -void RoctracerLogger::clearLogs() { - rows_.clear(); - kernelRows_.clear(); - copyRows_.clear(); - mallocRows_.clear(); - gpuTraceBuffers_->clear(); - for (int i = 0; i < CorrelationDomain::size; ++i) { - externalCorrelations_[i].clear(); - } -} - -void RoctracerLogger::api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) { - RoctracerLogger *dis = &singleton(); - - if (domain == ACTIVITY_DOMAIN_HIP_API && dis->loggedIds_.contains(cid)) { - const hip_api_data_t* data = (const hip_api_data_t*)(callback_data); - - // Pack callbacks into row structures - - thread_local timespec timestamp; - - if (data->phase == ACTIVITY_API_PHASE_ENTER) { - clock_gettime(CLOCK_MONOTONIC, ×tamp); // record proper clock - } else { // (data->phase == ACTIVITY_API_PHASE_EXIT) - timespec endTime; - timespec startTime { timestamp }; - clock_gettime(CLOCK_MONOTONIC, &endTime); // record proper clock - - switch (cid) { - case HIP_API_ID_hipLaunchKernel: - case HIP_API_ID_hipExtLaunchKernel: - case HIP_API_ID_hipLaunchCooperativeKernel: // Should work here - { - auto &args = data->args.hipLaunchKernel; - dis->kernelRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - args.function_address, - nullptr, - args.numBlocks.x, - args.numBlocks.y, - args.numBlocks.z, - args.dimBlocks.x, - args.dimBlocks.y, - args.dimBlocks.z, - args.sharedMemBytes, - args.stream - ); - } - break; - case HIP_API_ID_hipHccModuleLaunchKernel: - case HIP_API_ID_hipModuleLaunchKernel: - case HIP_API_ID_hipExtModuleLaunchKernel: - { - auto &args = data->args.hipModuleLaunchKernel; - dis->kernelRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - nullptr, - args.f, - args.gridDimX, - args.gridDimY, - args.gridDimZ, - args.blockDimX, - args.blockDimY, - args.blockDimZ, - args.sharedMemBytes, - args.stream - ); - } - break; - case HIP_API_ID_hipLaunchCooperativeKernelMultiDevice: - case HIP_API_ID_hipExtLaunchMultiKernelMultiDevice: -#if 0 - { - auto &args = data->args.hipLaunchCooperativeKernelMultiDevice.launchParamsList__val; - dis->kernelRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - args.function_address, - nullptr, - args.numBlocks.x, - args.numBlocks.y, - args.numBlocks.z, - args.dimBlocks.x, - args.dimBlocks.y, - args.dimBlocks.z, - args.sharedMemBytes, - args.stream - ); - } -#endif - break; - case HIP_API_ID_hipMalloc: - dis->mallocRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - data->args.hipMalloc.ptr__val, - data->args.hipMalloc.size - ); - break; - case HIP_API_ID_hipFree: - dis->mallocRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - data->args.hipFree.ptr, - 0 - ); - break; - case HIP_API_ID_hipMemcpy: - { - auto &args = data->args.hipMemcpy; - dis->copyRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - args.src, - args.dst, - args.sizeBytes, - args.kind, - static_cast(0) // use placeholder? - ); - } - break; - case HIP_API_ID_hipMemcpyAsync: - case HIP_API_ID_hipMemcpyWithStream: - { - auto &args = data->args.hipMemcpyAsync; - dis->copyRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - args.src, - args.dst, - args.sizeBytes, - args.kind, - args.stream - ); - } - break; - default: - dis->rows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime) - ); - break; - } // switch - // External correlation - for (int it = CorrelationDomain::begin; it < CorrelationDomain::end; ++it) { - if (t_externalIds[it].size() > 0) { - dis->externalCorrelations_[it][data->correlation_id] = t_externalIds[it].back(); - } - } - } // phase exit - } -} - -void RoctracerLogger::activity_callback(const char* begin, const char* end, void* arg) -{ - size_t size = end - begin; - uint8_t *buffer = (uint8_t*) malloc(size); - auto &gpuTraceBuffers = singleton().gpuTraceBuffers_; - memcpy(buffer, begin, size); - gpuTraceBuffers->emplace_back(buffer, size); -} - -void RoctracerLogger::startLogging() { - if (!registered_) { - roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, nullptr); // Magic encantation - - // Set some api calls to ignore - loggedIds_.setInvertMode(true); // Omit the specified api - loggedIds_.add("hipGetDevice"); - loggedIds_.add("hipSetDevice"); - loggedIds_.add("hipGetLastError"); - loggedIds_.add("__hipPushCallConfiguration"); - loggedIds_.add("__hipPopCallConfiguration"); - loggedIds_.add("hipCtxSetCurrent"); - loggedIds_.add("hipEventRecord"); - loggedIds_.add("hipEventQuery"); - loggedIds_.add("hipGetDeviceProperties"); - loggedIds_.add("hipPeekAtLastError"); - loggedIds_.add("hipModuleGetFunction"); - loggedIds_.add("hipEventCreateWithFlags"); - - // Enable API callbacks - if (loggedIds_.invertMode() == true) { - // exclusion list - enable entire domain and turn off things in list - roctracer_enable_domain_callback(ACTIVITY_DOMAIN_HIP_API, api_callback, nullptr); - const std::unordered_map &filter = loggedIds_.filterList(); - for (auto it = filter.begin(); it != filter.end(); ++it) { - roctracer_disable_op_callback(ACTIVITY_DOMAIN_HIP_API, it->first); - } - } - else { - // inclusion list - only enable things in the list - const std::unordered_map &filter = loggedIds_.filterList(); - roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); - for (auto it = filter.begin(); it != filter.end(); ++it) { - roctracer_enable_op_callback(ACTIVITY_DOMAIN_HIP_API, it->first, api_callback, nullptr); - } - } - //roctracer_enable_domain_callback(ACTIVITY_DOMAIN_ROCTX, api_callback, nullptr); - - // Allocate default tracing pool - roctracer_properties_t properties; - memset(&properties, 0, sizeof(roctracer_properties_t)); - properties.buffer_size = 0x1000; - roctracer_open_pool(&properties); - - // Enable async op collection - roctracer_properties_t hcc_cb_properties; - memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); - hcc_cb_properties.buffer_size = 0x4000; - hcc_cb_properties.buffer_callback_fun = activity_callback; - roctracer_open_pool_expl(&hcc_cb_properties, &hccPool_); - roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HCC_OPS, hccPool_); - - registered_ = true; - } - - externalCorrelationEnabled_ = true; - roctracer_start(); -} - -void RoctracerLogger::stopLogging() { - roctracer_stop(); - roctracer_flush_activity_expl(hccPool_); -} - -void RoctracerLogger::endTracing() { - if (registered_ == true) { - roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); - //roctracer_disable_domain_callback(ACTIVITY_DOMAIN_ROCTX); - - roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS); - roctracer_close_pool_expl(hccPool_); - } -} - - -ApiIdList::ApiIdList() -: invert_(true) -{ -} - -void ApiIdList::add(const std::string &apiName) -{ - uint32_t cid = 0; - if (roctracer_op_code(ACTIVITY_DOMAIN_HIP_API, apiName.c_str(), &cid, nullptr) == ROCTRACER_STATUS_SUCCESS) { - filter_[cid] = 1; - } -} -void ApiIdList::remove(const std::string &apiName) -{ - uint32_t cid = 0; - if (roctracer_op_code(ACTIVITY_DOMAIN_HIP_API, apiName.c_str(), &cid, nullptr) == ROCTRACER_STATUS_SUCCESS) { - filter_.erase(cid); - } -} - -bool ApiIdList::loadUserPrefs() -{ - // placeholder - return false; -} -bool ApiIdList::contains(uint32_t apiId) -{ - return (filter_.find(apiId) != filter_.end()) ? !invert_ : invert_; // XOR -} diff --git a/onnxruntime/core/providers/rocm/RoctracerLogger.h b/onnxruntime/core/providers/rocm/RoctracerLogger.h deleted file mode 100644 index 55c952a9b9d8a..0000000000000 --- a/onnxruntime/core/providers/rocm/RoctracerLogger.h +++ /dev/null @@ -1,170 +0,0 @@ -#pragma once - -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include -#include - - -namespace onnxruntime{ -namespace profiling { -class RocmProfiler; -} -} - -class RoctracerActivityBuffer { -public: - // data must be allocated using malloc. - // Ownership is transferred to this object. - RoctracerActivityBuffer(uint8_t* data, size_t validSize) - : data_(data), validSize_(validSize) {} - - ~RoctracerActivityBuffer() { - free(data_); - } - - // Allocated by malloc - uint8_t* data_{nullptr}; - - // Number of bytes used - size_t validSize_; -}; - - -class ApiIdList -{ -public: - ApiIdList(); - bool invertMode() { return invert_; } - void setInvertMode(bool invert) { invert_ = invert; } - void add(const std::string &apiName); - void remove(const std::string &apiName); - bool loadUserPrefs(); - bool contains(uint32_t apiId); - const std::unordered_map &filterList() { return filter_; } - -private: - std::unordered_map filter_; - bool invert_; -}; - -struct roctracerRow { - roctracerRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid - , uint32_t tid, uint64_t begin, uint64_t end) - : id(id), domain(domain), cid(cid), pid(pid), tid(tid), begin(begin), end(end) {} - uint64_t id; // correlation_id - uint32_t domain; - uint32_t cid; - uint32_t pid; - uint32_t tid; - uint64_t begin; - uint64_t end; -}; - -struct kernelRow : public roctracerRow { - kernelRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid - , uint32_t tid, uint64_t begin, uint64_t end - , const void *faddr, hipFunction_t function - , unsigned int gx, unsigned int gy, unsigned int gz - , unsigned int wx, unsigned int wy, unsigned int wz - , size_t gss, hipStream_t stream) - : roctracerRow(id, domain, cid, pid, tid, begin, end), functionAddr(faddr) - , function(function), gridX(gx), gridY(gy), gridZ(gz) - , workgroupX(wx), workgroupY(wy), workgroupZ(wz), groupSegmentSize(gss) - , stream(stream) {} - const void* functionAddr; - hipFunction_t function; - unsigned int gridX; - unsigned int gridY; - unsigned int gridZ; - unsigned int workgroupX; - unsigned int workgroupY; - unsigned int workgroupZ; - size_t groupSegmentSize; - hipStream_t stream; -}; - -struct copyRow : public roctracerRow { - copyRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid - , uint32_t tid, uint64_t begin, uint64_t end - , const void* src, const void *dst, size_t size, hipMemcpyKind kind - , hipStream_t stream) - : roctracerRow(id, domain, cid, pid, tid, begin, end) - , src(src), dst(dst), size(size), kind(kind), stream(stream) {} - const void *src; - const void *dst; - size_t size; - hipMemcpyKind kind; - hipStream_t stream; -}; - -struct mallocRow : public roctracerRow { - mallocRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid - , uint32_t tid, uint64_t begin, uint64_t end - , const void* ptr, size_t size) - : roctracerRow(id, domain, cid, pid, tid, begin, end) - , ptr(ptr), size(size) {} - const void *ptr; - size_t size; -}; - - -class RoctracerLogger { - public: - enum CorrelationDomain { - begin, - Default = begin, - Domain0 = begin, - Domain1, - end, - size = end - }; - - RoctracerLogger(); - RoctracerLogger(const RoctracerLogger&) = delete; - RoctracerLogger& operator=(const RoctracerLogger&) = delete; - - virtual ~RoctracerLogger(); - - static RoctracerLogger& singleton(); - - static void pushCorrelationID(uint64_t id, CorrelationDomain type); - static void popCorrelationID(CorrelationDomain type); - - void startLogging(); - void stopLogging(); - void clearLogs(); - - private: - bool registered_{false}; - void endTracing(); - - roctracer_pool_t *hccPool_{NULL}; - static void api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg); - static void activity_callback(const char* begin, const char* end, void* arg); - - ApiIdList loggedIds_; - - // Api callback data - std::deque rows_; - std::deque kernelRows_; - std::deque copyRows_; - std::deque mallocRows_; - std::map externalCorrelations_[CorrelationDomain::size]; // tracer -> ext - - std::unique_ptr> gpuTraceBuffers_; - bool externalCorrelationEnabled_{true}; - - friend class onnxruntime::profiling::RocmProfiler; -}; diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc new file mode 100644 index 0000000000000..38f913fff8b39 --- /dev/null +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -0,0 +1,735 @@ +#include +#include +#include + +#include "roctracer_manager.h" + +namespace onnxruntime { +namespace profiling { + +// allocate a 16K buffer for recording async activities +static constexpr size_t sc_activity_buffer_size = 0x4000; + +// Implementation of RoctracerActivityBuffer +RoctracerActivityBuffer& RoctracerActivityBuffer::operator = (const RoctracerActivityBuffer& other) { + if (&other == this) { + return *this; + } + if (data_ != nullptr) { + free(data_); + } + + size_ = other.size_; + data_ = (uint8_t*)malloc(size_); + memcpy(data_, other.data_, size_); + return *this; +} + +RoctracerActivityBuffer& RoctracerActivityBuffer::operator = (RoctracerActivityBuffer&& other) { + if (&other == this) { + return *this; + } + std::swap(data_, other.data_); + std::swap(size_, other.size_); +} + +RoctracerActivityBuffer::~RoctracerActivityBuffer() { + if (data_ != nullptr) { + free(data_); + data_ = nullptr; + } +} + + +// Implementation of RoctracerManager +RoctracerManager& RoctracerManager::GetInstance() { + static RoctracerManager instance; + return instance; +} + +RoctracerManager::~RoctracerManager() { + StopLogging(); +} + +uint64_t RoctracerManager::RegisterClient() { + std::lock_guard lock(global_op_mutex_); + auto res = next_client_id_++; + per_client_locks_[res] = std::make_unique(); + per_client_events_by_ext_correlation_.insert({res, {}}); + return res; +} + +void RoctracerManager::DeregisterClient(uint64_t client_handle) { + std::lock_guard lock(global_op_mutex_); + auto it = per_client_locks_.find(client_handle); + per_client_events_by_ext_correlation_.erase(client_handle); + per_client_locks_.erase(client_handle); +} + +void RoctracerManager::StartLogging() { + std::lock_guard lock(global_op_mutex_); + if (logging_enabled_) { + return; + } + + roctracer_properties_t hcc_cb_properties; + memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); + hcc_cb_properties.buffer_size = sc_activity_buffer_size; + hcc_cb_properties.buffer_callback_fun = ActivityCallback; + roctracer_open_pool_expl(&hcc_cb_properties, &activity_pool_); + roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HCC_OPS, activity_pool_); + roctracer_start(); + logging_enabled_ = true; +} + +void RoctracerManager::StopLogging() { + std::lock_guard lock(global_op_mutex_); + if (!logging_enabled_) { + return; + } + + roctracer_stop(); + roctracer_flush_activity_expl(activity_pool_); + roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS); + roctracer_close_pool_expl(activity_pool_); + logging_enabled_ = false; +} + +void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_time, Events& events) { + std::vector activity_buffers; + { + std::lock_guard lock(unprocessed_activity_buffers_lock_); + std::swap(unprocessed_activity_buffers_, activity_buffers); + unprocessed_activity_buffers_.clear(); + } + + ProcessActivityBuffers(activity_buffers); +} + +bool RoctracerManager::PushCorrelation(uint64_t client_handle, uint64_t external_correlation_id) { + auto it = per_client_events_by_ext_correlation_.find(client_handle); + if (it == per_client_events_by_ext_correlation_.end()) { + return false; + } + roctracer_activity_push_external_correlation_id(external_correlation_id); + // assumption: correlation_ids are unique. This should generally work, + // because we use timestamps as correlation ids. + external_correlation_id_to_client_[external_correlation_id] = client_handle; + return true; +} + +void RoctracerManager::PopCorrelation(uint64_t& popped_external_correlation_id) { + roctracer_activity_pop_external_correlation_id(&popped_external_correlation_id); + external_correlation_id_to_client_.erase(popped_external_correlation_id); +} + +void RoctracerManager::ActivityCallback(const char* begin, const char* end, void* arg) { + size_t size = end - begin; + RoctracerActivityBuffer activity_buffer{reinterpret_cast(begin), size}; + auto& instance = GetInstance(); + { + std::lock_guard lock(instance.unprocessed_activity_buffers_lock_); + instance.unprocessed_activity_buffers_.emplace_back(std::move(activity_buffer)); + } +} + +void RoctracerManager::ApiCallback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) { + if (domain != ACTIVITY_DOMAIN_HIP_API) { + return; + } + const hip_api_data_t* data = reinterpret_cast(callback_data); + if (data->phase == ACTIVITY_API_PHASE_EXIT) { + // We only save args for async launches on the ACTIVITY_API_PHASE_ENTER phase + return; + } + + auto& instance = GetInstance(); + { + std::lock_guard lock(instance.api_call_args_lock_); + auto& record = instance.api_call_args_[data->correlation_id]; + record.domain_ = domain; + record.cid_ = cid; + record.api_data_ = *data; + } +} + +void RoctracerManager::CreateEventForKernelRecord(const roctracer_record_t* record, + uint64_t start_time_ns, + const ApiCallRecord& call_record, + EventRecord& event) { + auto const& launch_args = call_record.api_data_; + auto name = demangle(hipKernelNameRefByPtr(launch_args.args.hipLaunchKernel.function_address, + launch_args.args.hipLaunchKernel.stream)); + std::unordered_map args { + {"stream", std::to_string(launch_args.args.hipLaunchKernel.stream)}, + {"grid_x", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.x)}, + {"grid_y", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.y)}, + {"grid_z", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.z)}, + {"block_x", std::to_string(launch_args.args.hipLaunchKernel.dimBlocks.x)}, + {"block_y", std::to_string(launch_args.args.hipLaunchKernel.dimBlocks.y)}, + {"block_z", std::to_string(launch_args.args.hipLaunchKernel.dimBlocks.z)} + }; + + new (&event) EventRecord { + /* cat = */ EventCategory::KERNEL_EVENT, + /* pid = */ record->device_id, + /* tid = */ record->queue_id, + /* name = */ std::move(name), + /* ts = */ (record->begin_ns - start_time_ns) / 1000, + /* dur = */ (record->end_ns - record->begin_ns) / 1000, + /* args = */ std::move(args) + }; +} + +static inline std::string PointerToHexString(const void* ptr) +{ + std::ostringstream sstr; + sstr << std::hex << ptr; + return sstr.str(); +} + +void RoctracerManager::CreateEventForMemsetRecord(const roctracer_record_t* record, + uint64_t start_time_ns, + const ApiCallRecord& call_record, + EventRecord& event) { + auto const& launch_args = call_record.api_data_; + auto dst_string = PointerToHexString(launch_args.args.hipMemset.dst); + std::string name {roctracer_op_string(call_record.domain_, call_record.cid_, 0)}; + + std::unordered_map args { + {"stream", call_record.cid_ == HIP_API_ID_hipMemset + ? "0" + : std::to_string(launch_args.args.hipMemsetAsync.stream)}, + {"dst", dst_string}, + {"size", std::to_string(launch_args.args.hipMemset.sizeBytes)}, + {"value", std::to_string(launch_args.args.hipMemset.value)} + }; + new (&event) EventRecord { + /* cat = */ EventCategory::KERNEL_EVENT, + /* pid = */ record->device_id, + /* tid = */ record->queue_id, + /* name = */ std::move(name), + /* ts = */ (record->begin_ns - start_time_ns) / 1000, + /* dur = */ (record->end_ns - record->begin_ns) / 1000, + /* args = */ std::move(args) + }; +} + +static inline std::string MemcpyKindToString(hipMemcpyKind kind) { + switch(launch_args.args.hipMemcpy.kind) { + case hipMemcpyHostToHost: + return "H2H"; + case hipMemcpyHostToDevice: + return "H2D"; + case hipMemcpyDeviceToHost: + return "D2H"; + case hipMemcpyDeviceToDevice: + return "D2D"; + default: + return "Default"; + } +} + +void RoctracerManager::CreateEventForMemcpyRecord(const roctracer_record_t* record, uint64_t start_time_ns, + const ApiCallRecord& call_record, EventRecord& event) { + auto const& launch_args = call_record.api_data_; + auto src_string = PointerToHexString(launch_args.args.hipMemcpy.src); + auto dst_string = PointerToHexString(launch_args.args.hipMemcpy.dst); + std::string name {roctracer_op_string(call_record.domain_, call_record.cid_, 0)}; + + std::string memcpy_kind_string = MemcpyKindToString(launch_args.args.hipMemcpy.kind); + + std::unordered_map args { + {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy + ? "0" + : std::to_string(launch_args.args.hipMemcpyAsync.stream)}, + {"src", src_string}, + {"dst", dst_string}, + {"kind", memcpy_kind_string} + }; + + new (&event) EventRecord { + /* cat = */ EventCategory::KERNEL_EVENT, + /* pid = */ record->device_id, + /* tid = */ record->queue_id, + /* name = */ std::move(name), + /* ts = */ (record->begin_ns - start_time_ns) / 1000, + /* dur = */ (record->end_ns - record->begin_ns) / 1000, + /* args = */ std::move(args) + }; +} + +void RoctracerManager::CreateEventForMemcpy2DRecord(const roctracer_record_t* record, uint64_t start_time_ns, + const ApiCallRecord& call_record, EventRecord& event) { + auto const& launch_args = call_record.api_data_; + auto src_string = PointerToHexString(launch_args.args.hipMemcpy2D.src); + auto dst_string = PointerToHexString(launch_args.args.hipMemcpy2D.dst); + std::string name {roctracer_op_string(call_record.domain_, call_record.cid_, 0)}; + + std::string memcpy_kind_string = MemcpyKindToString(launch_args.args.hipMemcpy2D.kind); + + std::unordered_map args { + {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy2D + ? "0" + : std::to_string(launch_args.args.hipMemcpy2DAsync.stream)}, + {"src", src_string}, + {"dst", dst_string}, + {"spitch", std::to_string(launch_args.args.hipMemcpy2D.spitch)}, + {"dpitch", std::to_string(launch_args.args.hipMemcpy2D.dpitch)}, + {"width", std::to_string(launch_args.args.hipMemcpy2D.width)}, + {"height", std::to_string(launch_args.args.hipMemcpy2D.height)}, + {"kind", memcpy_kind_string} + }; + + new (&event) EventRecord { + /* cat = */ EventCategory::KERNEL_EVENT, + /* pid = */ record->device_id, + /* tid = */ record->queue_id, + /* name = */ std::move(name), + /* ts = */ (record->begin_ns - start_time_ns) / 1000, + /* dur = */ (record->end_ns - record->begin_ns) / 1000, + /* args = */ std::move(args) + }; +} + +void RoctracerManager::MapEventsToClient(uint64_t external_correlation_id, Events&& events) { + auto client_it = external_correlation_id_to_client_.find(external_correlation_id); + if (client_it == external_correlation_id_to_client_.end()) { + // :-( well, we tried really, really hard to map this event to a client. + return; + } + auto& event_list = per_client_events_by_ext_correlation_[client_it->second][client_it->first]; + event_list.insert(event_list.end(), + std::make_move_iterator(events.begin()), + std::make_move_iterator(events.end())); +} + +void RoctracerManager::MapEventToClient(uint64_t external_correlation_id, EventRecord&& event) { + auto client_it = external_correlation_id_to_client_.find(external_correlation_id); + if (client_it == external_correlation_id_to_client_.end()) { + // :-( well, we tried really, really hard to map this event to a client. + return; + } + per_client_events_by_ext_correlation_[client_it->second][client_it->first].emplace_back(std::move(event)); +} + +void RoctracerManager::ProcessActivityBuffers(const std::vector& buffers, + const TimePoint& start_time) { + std::unordered_map> events_pending_client_mapping; + auto start_time_ns = std::chrono::duration_cast(start_time.time_since_epoch()).count(); + + for (auto const& buffer : buffers) { + auto current_record = reinterpret_cast(buffer.GetData()); + auto data_end = reinterpret_cast(buffer.GetData() + buffer.GetSize()); + for ( ; current_record < data_end; roctracer_next_record(current_record, ¤t_record)) { + EventRecord event; + if (current_record->domain == ACTIVITY_DOMAIN_EXT_API) { + roctracer_correlation_to_external_correlation_[current_record->correlation_id] = current_record->external_id; + + // check for any events pending client mapping on this correlation + auto pending_it = events_pending_client_mapping.find(current_record->correlation_id); + if (pending_it == events_pending_client_mapping.end()) { + continue; + } + + // we have one or more pending events, map them to the client + MapEventsToClient(current_record->external_id, std::move(pending_it->second)); + events_pending_client_mapping.erase(current_record->correlation_id); + + } else if (current_record->domain == ACTIVITY_DOMAIN_HIP_OPS) { + if (current_record->op == 1 && current_record->kind == HipOpMarker) { + // this is just a marker, ignore it. + continue; + } + + auto api_it = api_call_args_.find(current_record->correlation_id); + if (api_it == api_call_args_.end()) { + // we're not tracking this activity, ignore it + continue; + } + + auto const& call_record = api_it->second; + switch (call_record.cid_) + { + case HIP_API_ID_hipLaunchKernel: + CreateEventForKernelRecord(current_record, start_time_ns, call_record, event); + break; + + case HIP_API_ID_hipMemset: + case HIP_API_ID_hipMemsetAsync: + CreateEventForMemsetRecord(current_record, start_time_ns, call_record, event); + break; + + case HIP_API_ID_hipMemcpy: + case HIP_API_ID_hipMemcpyAsync: + CreateEventForMemcpyRecord(current_record, start_time_ns, call_record, event); + break; + + case HIP_API_ID_hipMemcpy2D: + case HIP_API_ID_hipMemcpy2DAsync: + CreateEventForMemcpy2DRecord(current_record, start_time_ns, call_record, event); + break; + + default: + break; + } + } + + // map the event to the right client + auto ext_corr_it = roctracer_correlation_to_external_correlation_.find(current_record->correlation_id); + if (ext_corr_it == roctracer_correlation_to_external_correlation_.end()) { + // defer the processing of this event + events_pending_client_mapping[current_record->correlation_id].emplace_back(std::move(event)); + continue; + } + MapEventToClient(ext_corr_it->second, std::move(event)); + } + } +} + +} /* end namespace profiling */ +} /* end namespace onnxruntime */ + + +/* +typedef uint64_t timestamp_t; + +static timestamp_t timespec_to_ns(const timespec& time) { + return ((timestamp_t)time.tv_sec * 1000000000) + time.tv_nsec; +} + +//using namespace std::chrono; + +RoctracerLogger& RoctracerLogger::singleton() { + static RoctracerLogger instance; + return instance; +} + +RoctracerLogger::RoctracerLogger() { + gpuTraceBuffers_ = std::make_unique>(); +} + +RoctracerLogger::~RoctracerLogger() { + stopLogging(); + endTracing(); +} + +namespace { + thread_local std::deque t_externalIds[RoctracerLogger::CorrelationDomain::size]; +} + +void RoctracerLogger::pushCorrelationID(uint64_t id, CorrelationDomain type) { + if (!singleton().externalCorrelationEnabled_) { + return; + } + t_externalIds[type].push_back(id); +} + +void RoctracerLogger::popCorrelationID(CorrelationDomain type) { + if (!singleton().externalCorrelationEnabled_) { + return; + } + t_externalIds[type].pop_back(); +} + +void RoctracerLogger::clearLogs() { + rows_.clear(); + kernelRows_.clear(); + copyRows_.clear(); + mallocRows_.clear(); + gpuTraceBuffers_->clear(); + for (int i = 0; i < CorrelationDomain::size; ++i) { + externalCorrelations_[i].clear(); + } +} + +void RoctracerLogger::api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) { + RoctracerLogger *dis = &singleton(); + + if (domain == ACTIVITY_DOMAIN_HIP_API && dis->loggedIds_.contains(cid)) { + const hip_api_data_t* data = (const hip_api_data_t*)(callback_data); + + // Pack callbacks into row structures + + thread_local timespec timestamp; + + if (data->phase == ACTIVITY_API_PHASE_ENTER) { + clock_gettime(CLOCK_MONOTONIC, ×tamp); // record proper clock + } else { // (data->phase == ACTIVITY_API_PHASE_EXIT) + timespec endTime; + timespec startTime { timestamp }; + clock_gettime(CLOCK_MONOTONIC, &endTime); // record proper clock + + switch (cid) { + case HIP_API_ID_hipLaunchKernel: + case HIP_API_ID_hipExtLaunchKernel: + case HIP_API_ID_hipLaunchCooperativeKernel: // Should work here + { + auto &args = data->args.hipLaunchKernel; + dis->kernelRows_.emplace_back(data->correlation_id, + domain, + cid, + processId(), + systemThreadId(), + timespec_to_ns(startTime), + timespec_to_ns(endTime), + args.function_address, + nullptr, + args.numBlocks.x, + args.numBlocks.y, + args.numBlocks.z, + args.dimBlocks.x, + args.dimBlocks.y, + args.dimBlocks.z, + args.sharedMemBytes, + args.stream + ); + } + break; + case HIP_API_ID_hipHccModuleLaunchKernel: + case HIP_API_ID_hipModuleLaunchKernel: + case HIP_API_ID_hipExtModuleLaunchKernel: + { + auto &args = data->args.hipModuleLaunchKernel; + dis->kernelRows_.emplace_back(data->correlation_id, + domain, + cid, + processId(), + systemThreadId(), + timespec_to_ns(startTime), + timespec_to_ns(endTime), + nullptr, + args.f, + args.gridDimX, + args.gridDimY, + args.gridDimZ, + args.blockDimX, + args.blockDimY, + args.blockDimZ, + args.sharedMemBytes, + args.stream + ); + } + break; + case HIP_API_ID_hipLaunchCooperativeKernelMultiDevice: + case HIP_API_ID_hipExtLaunchMultiKernelMultiDevice: +#if 0 + { + auto &args = data->args.hipLaunchCooperativeKernelMultiDevice.launchParamsList__val; + dis->kernelRows_.emplace_back(data->correlation_id, + domain, + cid, + processId(), + systemThreadId(), + timespec_to_ns(startTime), + timespec_to_ns(endTime), + args.function_address, + nullptr, + args.numBlocks.x, + args.numBlocks.y, + args.numBlocks.z, + args.dimBlocks.x, + args.dimBlocks.y, + args.dimBlocks.z, + args.sharedMemBytes, + args.stream + ); + } +#endif + break; + case HIP_API_ID_hipMalloc: + dis->mallocRows_.emplace_back(data->correlation_id, + domain, + cid, + processId(), + systemThreadId(), + timespec_to_ns(startTime), + timespec_to_ns(endTime), + data->args.hipMalloc.ptr__val, + data->args.hipMalloc.size + ); + break; + case HIP_API_ID_hipFree: + dis->mallocRows_.emplace_back(data->correlation_id, + domain, + cid, + processId(), + systemThreadId(), + timespec_to_ns(startTime), + timespec_to_ns(endTime), + data->args.hipFree.ptr, + 0 + ); + break; + case HIP_API_ID_hipMemcpy: + { + auto &args = data->args.hipMemcpy; + dis->copyRows_.emplace_back(data->correlation_id, + domain, + cid, + processId(), + systemThreadId(), + timespec_to_ns(startTime), + timespec_to_ns(endTime), + args.src, + args.dst, + args.sizeBytes, + args.kind, + static_cast(0) // use placeholder? + ); + } + break; + case HIP_API_ID_hipMemcpyAsync: + case HIP_API_ID_hipMemcpyWithStream: + { + auto &args = data->args.hipMemcpyAsync; + dis->copyRows_.emplace_back(data->correlation_id, + domain, + cid, + processId(), + systemThreadId(), + timespec_to_ns(startTime), + timespec_to_ns(endTime), + args.src, + args.dst, + args.sizeBytes, + args.kind, + args.stream + ); + } + break; + default: + dis->rows_.emplace_back(data->correlation_id, + domain, + cid, + processId(), + systemThreadId(), + timespec_to_ns(startTime), + timespec_to_ns(endTime) + ); + break; + } // switch + // External correlation + for (int it = CorrelationDomain::begin; it < CorrelationDomain::end; ++it) { + if (t_externalIds[it].size() > 0) { + dis->externalCorrelations_[it][data->correlation_id] = t_externalIds[it].back(); + } + } + } // phase exit + } +} + +void RoctracerLogger::activity_callback(const char* begin, const char* end, void* arg) +{ + size_t size = end - begin; + uint8_t *buffer = (uint8_t*) malloc(size); + auto &gpuTraceBuffers = singleton().gpuTraceBuffers_; + memcpy(buffer, begin, size); + gpuTraceBuffers->emplace_back(buffer, size); +} + +void RoctracerLogger::startLogging() { + if (!registered_) { + roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, nullptr); // Magic encantation + + // Set some api calls to ignore + loggedIds_.setInvertMode(true); // Omit the specified api + loggedIds_.add("hipGetDevice"); + loggedIds_.add("hipSetDevice"); + loggedIds_.add("hipGetLastError"); + loggedIds_.add("__hipPushCallConfiguration"); + loggedIds_.add("__hipPopCallConfiguration"); + loggedIds_.add("hipCtxSetCurrent"); + loggedIds_.add("hipEventRecord"); + loggedIds_.add("hipEventQuery"); + loggedIds_.add("hipGetDeviceProperties"); + loggedIds_.add("hipPeekAtLastError"); + loggedIds_.add("hipModuleGetFunction"); + loggedIds_.add("hipEventCreateWithFlags"); + + // Enable API callbacks + if (loggedIds_.invertMode() == true) { + // exclusion list - enable entire domain and turn off things in list + roctracer_enable_domain_callback(ACTIVITY_DOMAIN_HIP_API, api_callback, nullptr); + const std::unordered_map &filter = loggedIds_.filterList(); + for (auto it = filter.begin(); it != filter.end(); ++it) { + roctracer_disable_op_callback(ACTIVITY_DOMAIN_HIP_API, it->first); + } + } + else { + // inclusion list - only enable things in the list + const std::unordered_map &filter = loggedIds_.filterList(); + roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); + for (auto it = filter.begin(); it != filter.end(); ++it) { + roctracer_enable_op_callback(ACTIVITY_DOMAIN_HIP_API, it->first, api_callback, nullptr); + } + } + //roctracer_enable_domain_callback(ACTIVITY_DOMAIN_ROCTX, api_callback, nullptr); + + // Allocate default tracing pool + roctracer_properties_t properties; + memset(&properties, 0, sizeof(roctracer_properties_t)); + properties.buffer_size = 0x1000; + roctracer_open_pool(&properties); + + // Enable async op collection + roctracer_properties_t hcc_cb_properties; + memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); + hcc_cb_properties.buffer_size = 0x4000; + hcc_cb_properties.buffer_callback_fun = activity_callback; + roctracer_open_pool_expl(&hcc_cb_properties, &hccPool_); + roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HCC_OPS, hccPool_); + + registered_ = true; + } + + externalCorrelationEnabled_ = true; + roctracer_start(); +} + +void RoctracerLogger::stopLogging() { + roctracer_stop(); + roctracer_flush_activity_expl(hccPool_); +} + +void RoctracerLogger::endTracing() { + if (registered_ == true) { + roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); + //roctracer_disable_domain_callback(ACTIVITY_DOMAIN_ROCTX); + + roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS); + roctracer_close_pool_expl(hccPool_); + } +} + + +ApiIdList::ApiIdList() +: invert_(true) +{ +} + +void ApiIdList::add(const std::string &apiName) +{ + uint32_t cid = 0; + if (roctracer_op_code(ACTIVITY_DOMAIN_HIP_API, apiName.c_str(), &cid, nullptr) == ROCTRACER_STATUS_SUCCESS) { + filter_[cid] = 1; + } +} +void ApiIdList::remove(const std::string &apiName) +{ + uint32_t cid = 0; + if (roctracer_op_code(ACTIVITY_DOMAIN_HIP_API, apiName.c_str(), &cid, nullptr) == ROCTRACER_STATUS_SUCCESS) { + filter_.erase(cid); + } +} + +bool ApiIdList::loadUserPrefs() +{ + // placeholder + return false; +} +bool ApiIdList::contains(uint32_t apiId) +{ + return (filter_.find(apiId) != filter_.end()) ? !invert_ : invert_; // XOR +} +*/ diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h new file mode 100644 index 0000000000000..d42998d6ed73f --- /dev/null +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -0,0 +1,271 @@ +#pragma once + +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include "core/common/profiler_common.h" + + +namespace onnxruntime{ +namespace profiling { + +class RoctracerActivityBuffer { +public: + RoctracerActivityBuffer() + : data_(nullptr), size_(0) {} + + RoctracerActivityBuffer(const uint8_t* data, size_t size) + : data_((uint8_t*)malloc(size)), size_(size) { + memcpy(data_, data, size); + } + + RoctracerActivityBuffer(const RoctracerActivityBuffer& other) + : RoctracerActivityBuffer(other.data_, other.size_) {} + + RoctracerActivityBuffer(RoctracerActivityBuffer&& other) + : RoctracerActivityBuffer() { + std::swap(data_, other.data_); + std::swap(size_, other.size_); + } + + RoctracerActivityBuffer& operator = (const RoctracerActivityBuffer& other); + RoctracerActivityBuffer& operator = (RoctracerActivityBuffer&& other); + + ~RoctracerActivityBuffer(); + + // accessors + uint8_t* GetData() { return data_; } + const uint8_t* GetData() const { return data_; } + size_t GetSize() const { return size_; } + +private: + uint8_t* data_; + size_t size_; +}; + +static const char* hip_api_calls_to_trace[] = { + "hipMemcpy", + "hipMemcpy2D", + "hipMemcpyAsync", + "hipMemcpy2DAsync", + "hipLaunchKernel", + "hipMemset", + "hipMemsetAsync", +}; + +struct ApiCallRecord { + uint32_t domain_; + uint32_t cid_; + hip_api_data_t api_data_; +}; + +class RoctracerManager +{ +public: + RoctracerManager(const RoctracerManager&) = delete; + RoctracerManager& operator = (const RoctracerManager&) = delete; + RoctracerManager() = default; + ~RoctracerManager(); + + static RoctracerManager& GetInstance(); + + uint64_t RegisterClient(); + void DeregisterClient(uint64_t client_handle); + + void StartLogging(); + void StopLogging(); + void Consume(uint64_t client_handle, const TimePoint& start_time, Events& events); + + bool PushCorrelation(uint64_t client_handle, uint64_t external_correlation_id); + void PopCorrelation(uint64_t& popped_correlation_id); + bool PopCorrelation(); + +private: + static void ActivityCallback(const char* begin, const char* end, void* arg); + static void ApiCallback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg); + void ProcessActivityBuffers(const std::vector& buffers, const TimePoint& start_time); + + // Per-API (roughly) helpers for event construction. + void CreateEventForKernelRecord(const roctracer_record_t* record, uint64_t start_time_ns, + const ApiCallRecord& call_record, EventRecord& event); + void CreateEventForMemsetRecord(const roctracer_record_t* record, uint64_t start_time_ns, + const ApiCallRecord& call_record, EventRecord& event); + void CreateEventForMemcpyRecord(const roctracer_record_t* record, uint64_t start_time_ns, + const ApiCallRecord& call_record, EventRecord& event); + void CreateEventForMemcpy2DRecord(const roctracer_record_t* record, uint64_t start_time_ns, + const ApiCallRecord& call_record, EventRecord& event); + void MapEventToClient(uint64_t external_correlation_id, EventRecord&& event); + void MapEventsToClient(uint64_t external_correlation_id, Events&& events); + + // Some useful constants for processing activity buffers + static constexpr uint32_t HipOpMarker = 4606; + + std::mutex unprocessed_activity_buffers_lock_; + std::vector unprocessed_activity_buffers_; + std::mutex api_call_args_lock_; + std::unordered_map api_call_args_; + + // Keyed on client_id / client_handle + std::unordered_map> per_client_locks_; + + // Keyed on external_correlation_id -> client_id/client_handle + std::unordered_map external_correlation_id_to_client_; + + // Keyed on roctracer_correlation_id -> external_correlation_id + std::unordered_map roctracer_correlation_to_external_correlation_; + + // client_id/client_handle -> external_correlation_id -> events + std::unordered_map> per_client_events_by_ext_correlation_; + uint64_t next_client_id_ = 1; + bool logging_enabled_ = false; + std::mutex global_op_mutex_; + roctracer_pool_t* activity_pool_; +}; + +} /* end namespace profiling */ +} /* end namespace onnxruntime */ + + +/* +} +} + +class RoctracerActivityBuffer { +public: + // data must be allocated using malloc. + // Ownership is transferred to this object. + RoctracerActivityBuffer(uint8_t* data, size_t validSize) + : data_(data), validSize_(validSize) {} + + ~RoctracerActivityBuffer() { + free(data_); + } + + // Allocated by malloc + uint8_t* data_{nullptr}; + + // Number of bytes used + size_t validSize_; +}; + +struct roctracerRow { + roctracerRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid + , uint32_t tid, uint64_t begin, uint64_t end) + : id(id), domain(domain), cid(cid), pid(pid), tid(tid), begin(begin), end(end) {} + uint64_t id; // correlation_id + uint32_t domain; + uint32_t cid; + uint32_t pid; + uint32_t tid; + uint64_t begin; + uint64_t end; +}; + +struct kernelRow : public roctracerRow { + kernelRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid + , uint32_t tid, uint64_t begin, uint64_t end + , const void *faddr, hipFunction_t function + , unsigned int gx, unsigned int gy, unsigned int gz + , unsigned int wx, unsigned int wy, unsigned int wz + , size_t gss, hipStream_t stream) + : roctracerRow(id, domain, cid, pid, tid, begin, end), functionAddr(faddr) + , function(function), gridX(gx), gridY(gy), gridZ(gz) + , workgroupX(wx), workgroupY(wy), workgroupZ(wz), groupSegmentSize(gss) + , stream(stream) {} + const void* functionAddr; + hipFunction_t function; + unsigned int gridX; + unsigned int gridY; + unsigned int gridZ; + unsigned int workgroupX; + unsigned int workgroupY; + unsigned int workgroupZ; + size_t groupSegmentSize; + hipStream_t stream; +}; + +struct copyRow : public roctracerRow { + copyRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid + , uint32_t tid, uint64_t begin, uint64_t end + , const void* src, const void *dst, size_t size, hipMemcpyKind kind + , hipStream_t stream) + : roctracerRow(id, domain, cid, pid, tid, begin, end) + , src(src), dst(dst), size(size), kind(kind), stream(stream) {} + const void *src; + const void *dst; + size_t size; + hipMemcpyKind kind; + hipStream_t stream; +}; + +struct mallocRow : public roctracerRow { + mallocRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid + , uint32_t tid, uint64_t begin, uint64_t end + , const void* ptr, size_t size) + : roctracerRow(id, domain, cid, pid, tid, begin, end) + , ptr(ptr), size(size) {} + const void *ptr; + size_t size; +}; + + +class RoctracerLogger { + public: + enum CorrelationDomain { + begin, + Default = begin, + Domain0 = begin, + Domain1, + end, + size = end + }; + + RoctracerLogger(); + RoctracerLogger(const RoctracerLogger&) = delete; + RoctracerLogger& operator=(const RoctracerLogger&) = delete; + + virtual ~RoctracerLogger(); + + static RoctracerLogger& singleton(); + + static void pushCorrelationID(uint64_t id, CorrelationDomain type); + static void popCorrelationID(CorrelationDomain type); + + void startLogging(); + void stopLogging(); + void clearLogs(); + + private: + bool registered_{false}; + void endTracing(); + + roctracer_pool_t *hccPool_{NULL}; + static void api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg); + static void activity_callback(const char* begin, const char* end, void* arg); + + ApiIdList loggedIds_; + + // Api callback data + std::deque rows_; + std::deque kernelRows_; + std::deque copyRows_; + std::deque mallocRows_; + std::map externalCorrelations_[CorrelationDomain::size]; // tracer -> ext + + std::unique_ptr> gpuTraceBuffers_; + bool externalCorrelationEnabled_{true}; + + friend class onnxruntime::profiling::RocmProfiler; +}; +*/ From 037add468e7ca42fe0c052e4b603b2cc44a69bf9 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Sat, 29 Oct 2022 01:03:15 +0000 Subject: [PATCH 02/41] Delete duplicated code --- onnxruntime/core/providers/rocm/ThreadUtil.cc | 37 ------------------- onnxruntime/core/providers/rocm/ThreadUtil.h | 16 -------- 2 files changed, 53 deletions(-) delete mode 100644 onnxruntime/core/providers/rocm/ThreadUtil.cc delete mode 100644 onnxruntime/core/providers/rocm/ThreadUtil.h diff --git a/onnxruntime/core/providers/rocm/ThreadUtil.cc b/onnxruntime/core/providers/rocm/ThreadUtil.cc deleted file mode 100644 index 2ae95256ce639..0000000000000 --- a/onnxruntime/core/providers/rocm/ThreadUtil.cc +++ /dev/null @@ -1,37 +0,0 @@ -#include "ThreadUtil.h" - -#include "core/common/logging/logging.h" - -namespace onnxruntime { - -namespace profiling { - -namespace { -thread_local int32_t _pid = 0; -thread_local int32_t _tid = 0; -thread_local int32_t _sysTid = 0; -} - -int32_t processId() { - if (!_pid) { - _pid = onnxruntime::logging::GetProcessId(); - } - return _pid; -} - -int32_t systemThreadId() { - if (!_sysTid) { - _sysTid = onnxruntime::logging::GetThreadId(); - } - return _sysTid; -} - -int32_t threadId() { - if (!_tid) { - _tid = 0; //defeat for now - } - return _tid; -} - -} // namespace profiling -} // namespace onnxruntime diff --git a/onnxruntime/core/providers/rocm/ThreadUtil.h b/onnxruntime/core/providers/rocm/ThreadUtil.h deleted file mode 100644 index 2ec6aac722688..0000000000000 --- a/onnxruntime/core/providers/rocm/ThreadUtil.h +++ /dev/null @@ -1,16 +0,0 @@ -#pragma once - -#include - -namespace onnxruntime { - -namespace profiling { - -int32_t systemThreadId(); -int32_t threadId(); -int32_t processId(); - -} -} - -using namespace onnxruntime::profiling; From 40fed8b13a5323d31a5fc413cd2c06d4e434cf16 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 31 Oct 2022 19:33:54 +0000 Subject: [PATCH 03/41] WIP --- .../core/providers/rocm/rocm_profiler.cc | 385 +++++++++--------- .../core/providers/rocm/rocm_profiler.h | 21 +- .../core/providers/rocm/roctracer_manager.cc | 29 +- .../core/providers/rocm/roctracer_manager.h | 13 +- 4 files changed, 224 insertions(+), 224 deletions(-) diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.cc b/onnxruntime/core/providers/rocm/rocm_profiler.cc index 1b3b364d7f307..919cad17fb223 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.cc +++ b/onnxruntime/core/providers/rocm/rocm_profiler.cc @@ -7,7 +7,7 @@ #include "core/common/profiler_common.h" #include "core/providers/rocm/rocm_profiler.h" -#include "core/providers/rocm/RoctracerLogger.h" +#include "core/providers/rocm/roctracer_manager.h" #define BSIZE 4096 @@ -17,221 +17,228 @@ static timestamp_t timespec_to_ns(const timespec& time) { } namespace onnxruntime { - namespace profiling { -RocmProfiler::RocmProfiler() -: d(&RoctracerLogger::singleton()) -{ +RocmProfiler::RocmProfiler() { + auto& manager = RoctracerManager::GetInstance(); + client_handle_ = manager.RegisterClient(); } -RocmProfiler::~RocmProfiler() -{ +RocmProfiler::~RocmProfiler() { + auto& manager = RoctracerManager::GetInstance(); + manager.DeregisterClient(client_handle_); } bool RocmProfiler::StartProfiling() { - d->clearLogs(); - d->startLogging(); - return true; + auto& manager = RoctracerManager::GetInstance(); + manager.StartLogging(); } void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) { - d->stopLogging(); - - std::map> event_map; - std::map kernelLaunches; // correlation id -> kernel info - std::map copyLaunches; // correlation id -> copy info - - // Generate EventRecords - int64_t profiling_start = std::chrono::duration_cast(start_time.time_since_epoch()).count(); - - // Wrong clock again - all the cool kids are doing it - timespec t0, t1, t00; - clock_gettime(CLOCK_REALTIME, &t0); - clock_gettime(CLOCK_MONOTONIC, &t1); - clock_gettime(CLOCK_REALTIME, &t00); - const uint64_t toffset = (timespec_to_ns(t0) >> 1) + (timespec_to_ns(t00) >> 1) - timespec_to_ns(t1); - profiling_start = profiling_start - toffset; - - char buff[BSIZE]; - - for (auto &item : d->rows_) { - std::initializer_list> args = {{"op_name", ""}}; - addEventRecord(item, profiling_start, args, event_map); - } - - for (auto &item : d->mallocRows_) { - snprintf(buff, BSIZE, "%p", item.ptr); - const std::string arg_ptr{buff}; - - std::initializer_list> args = {{"op_name", ""}, - {"ptr", arg_ptr}, - {"size", std::to_string(item.size)} - }; - addEventRecord(item, profiling_start, args, event_map); - } - - for (auto &item : d->copyRows_) { - snprintf(buff, BSIZE, "%p", item.stream); - const std::string arg_stream{buff}; - snprintf(buff, BSIZE, "%p", item.src); - const std::string arg_src{buff}; - snprintf(buff, BSIZE, "%p", item.dst); - const std::string arg_dst{buff}; - - std::initializer_list> args = {{"op_name", ""}, - {"stream", arg_stream}, - {"src", arg_src}, - {"dst", arg_dst}, - {"size", std::to_string(item.size)}, - {"kind", std::to_string(item.kind)}, - }; - addEventRecord(item, profiling_start, args, event_map); - copyLaunches[item.id] = &item; - } - - for (auto &item : d->kernelRows_) { - snprintf(buff, BSIZE, "%p", item.stream); - const std::string arg_stream{buff}; - if (item.functionAddr) - snprintf(buff, BSIZE, "%s", demangle(hipKernelNameRefByPtr(item.functionAddr, item.stream)).c_str()); - else if (item.function) - snprintf(buff, BSIZE, "%s", demangle(hipKernelNameRef(item.function)).c_str()); - else - snprintf(buff, BSIZE, " "); - const std::string arg_kernel{buff}; - - std::initializer_list> args = {{"op_name", ""}, - {"stream", arg_stream}, - {"kernel", arg_kernel}, - {"grid_x", std::to_string(item.gridX)}, - {"grid_y", std::to_string(item.gridY)}, - {"grid_z", std::to_string(item.gridZ)}, - {"block_x", std::to_string(item.workgroupX)}, - {"block_y", std::to_string(item.workgroupY)}, - {"block_z", std::to_string(item.workgroupZ)}, - }; - addEventRecord(item, profiling_start, args, event_map); - kernelLaunches[item.id] = &item; - } + auto& manager = RoctracerManager::GetInstance(); + std::map event_map; + manager.Consume(client_handle_, start_time, event_map); - // Async Ops - e.g. "Kernel" - - for (auto& buffer : *d->gpuTraceBuffers_) { - const roctracer_record_t* record = (const roctracer_record_t*)(buffer.data_); - const roctracer_record_t* end_record = (const roctracer_record_t*)(buffer.data_ + buffer.validSize_); - while (record < end_record) { - std::unordered_map args; - std::string name = roctracer_op_string(record->domain, record->op, record->kind); - - // Add kernel args if we have them - auto kit = kernelLaunches.find(record->correlation_id); - if (kit != kernelLaunches.end()) { - auto &item = *(*kit).second; - snprintf(buff, BSIZE, "%p", item.stream); - args["stream"] = std::string(buff); - args["grid_x"] = std::to_string(item.gridX); - args["grid_y"] = std::to_string(item.gridY); - args["grid_z"] = std::to_string(item.gridZ); - args["block_x"] = std::to_string(item.workgroupX); - args["block_y"] = std::to_string(item.workgroupY); - args["block_z"] = std::to_string(item.workgroupZ); - if (item.functionAddr != nullptr) { - name = demangle(hipKernelNameRefByPtr(item.functionAddr, item.stream)).c_str(); - } - else if (item.function != nullptr) { - name = demangle(hipKernelNameRef(item.function)).c_str(); - } - } - - // Add copy args if we have them - auto cit = copyLaunches.find(record->correlation_id); - if (cit != copyLaunches.end()) { - auto &item = *(*cit).second; - snprintf(buff, BSIZE, "%p", item.stream); - args["stream"] = std::string(buff); - snprintf(buff, BSIZE, "%p", item.src); - args["dst"] = std::string(buff); - snprintf(buff, BSIZE, "%p", item.dst); - args["src"] = std::string(buff); - args["size"] = std::to_string(item.size); - args["kind"] = std::to_string(item.kind); - } - - EventRecord event{ - onnxruntime::profiling::KERNEL_EVENT, - static_cast(record->device_id), - static_cast(record->queue_id), - name, - static_cast((record->begin_ns - profiling_start) / 1000), - static_cast((record->end_ns - record->begin_ns) / 1000), - std::move(args)}; - - // FIXME: deal with missing ext correlation - auto extId = d->externalCorrelations_[RoctracerLogger::CorrelationDomain::Default][record->correlation_id]; - if (event_map.find(extId) == event_map.end()) { - event_map.insert({extId, {event}}); - } - else { - event_map[extId].push_back(std::move(event)); - } - - roctracer_next_record(record, &record); - } - } + Events merged_events; - // General - auto insert_iter = events.begin(); + auto event_iter = std::make_move_iterator(events.begin()); + auto event_end = std::make_move_iterator(events.end()); for (auto& map_iter : event_map) { auto ts = static_cast(map_iter.first); - while (insert_iter != events.end() && insert_iter->ts < ts) { - insert_iter++; + while (event_iter != event_end && event_iter->ts < ts) { + merged_events.emplace_back(*event_iter); + ++event_iter; } - if (insert_iter != events.end() && insert_iter->ts == ts) { - for (auto& evt_iter : map_iter.second) { - evt_iter.args["op_name"] = insert_iter->args["op_name"]; + + if (event_iter != event_end && event_iter->ts == ts) { + for (auto& evt : map_iter.second) { + evt.args["op_name"] = event_iter->args["op_name"]; } - insert_iter = events.insert(insert_iter+1, map_iter.second.begin(), map_iter.second.end()); - } else { - insert_iter = events.insert(insert_iter, map_iter.second.begin(), map_iter.second.end()); - } - while (insert_iter != events.end() && insert_iter->cat == EventCategory::KERNEL_EVENT) { - insert_iter++; + merged_events.emplace_back(*event_iter); + ++event_iter; } + + merged_events.insert(merged_events.end(), + std::make_move_iterator(map_iter.second.begin()), + std::make_move_iterator(map_iter.second.end())); } -} -void RocmProfiler::Start(uint64_t id) -{ - d->pushCorrelationID(id, RoctracerLogger::CorrelationDomain::Default); + // move any remaining events + merged_events.insert(merged_events.end(), event_iter, event_end); + std::swap(events, merged_events); + + // // Generate EventRecords + // int64_t profiling_start = std::chrono::duration_cast(start_time.time_since_epoch()).count(); + + // // Wrong clock again - all the cool kids are doing it + // timespec t0, t1, t00; + // clock_gettime(CLOCK_REALTIME, &t0); + // clock_gettime(CLOCK_MONOTONIC, &t1); + // clock_gettime(CLOCK_REALTIME, &t00); + // const uint64_t toffset = (timespec_to_ns(t0) >> 1) + (timespec_to_ns(t00) >> 1) - timespec_to_ns(t1); + // profiling_start = profiling_start - toffset; + + // char buff[BSIZE]; + + // for (auto &item : d->rows_) { + // std::initializer_list> args = {{"op_name", ""}}; + // addEventRecord(item, profiling_start, args, event_map); + // } + + // for (auto &item : d->mallocRows_) { + // snprintf(buff, BSIZE, "%p", item.ptr); + // const std::string arg_ptr{buff}; + + // std::initializer_list> args = {{"op_name", ""}, + // {"ptr", arg_ptr}, + // {"size", std::to_string(item.size)} + // }; + // addEventRecord(item, profiling_start, args, event_map); + // } + + // for (auto &item : d->copyRows_) { + // snprintf(buff, BSIZE, "%p", item.stream); + // const std::string arg_stream{buff}; + // snprintf(buff, BSIZE, "%p", item.src); + // const std::string arg_src{buff}; + // snprintf(buff, BSIZE, "%p", item.dst); + // const std::string arg_dst{buff}; + + // std::initializer_list> args = {{"op_name", ""}, + // {"stream", arg_stream}, + // {"src", arg_src}, + // {"dst", arg_dst}, + // {"size", std::to_string(item.size)}, + // {"kind", std::to_string(item.kind)}, + // }; + // addEventRecord(item, profiling_start, args, event_map); + // copyLaunches[item.id] = &item; + // } + + // for (auto &item : d->kernelRows_) { + // snprintf(buff, BSIZE, "%p", item.stream); + // const std::string arg_stream{buff}; + // if (item.functionAddr) + // snprintf(buff, BSIZE, "%s", demangle(hipKernelNameRefByPtr(item.functionAddr, item.stream)).c_str()); + // else if (item.function) + // snprintf(buff, BSIZE, "%s", demangle(hipKernelNameRef(item.function)).c_str()); + // else + // snprintf(buff, BSIZE, " "); + // const std::string arg_kernel{buff}; + + // std::initializer_list> args = {{"op_name", ""}, + // {"stream", arg_stream}, + // {"kernel", arg_kernel}, + // {"grid_x", std::to_string(item.gridX)}, + // {"grid_y", std::to_string(item.gridY)}, + // {"grid_z", std::to_string(item.gridZ)}, + // {"block_x", std::to_string(item.workgroupX)}, + // {"block_y", std::to_string(item.workgroupY)}, + // {"block_z", std::to_string(item.workgroupZ)}, + // }; + // addEventRecord(item, profiling_start, args, event_map); + // kernelLaunches[item.id] = &item; + // } + + // // Async Ops - e.g. "Kernel" + + // for (auto& buffer : *d->gpuTraceBuffers_) { + // const roctracer_record_t* record = (const roctracer_record_t*)(buffer.data_); + // const roctracer_record_t* end_record = (const roctracer_record_t*)(buffer.data_ + buffer.validSize_); + // while (record < end_record) { + // std::unordered_map args; + // std::string name = roctracer_op_string(record->domain, record->op, record->kind); + + // // Add kernel args if we have them + // auto kit = kernelLaunches.find(record->correlation_id); + // if (kit != kernelLaunches.end()) { + // auto &item = *(*kit).second; + // snprintf(buff, BSIZE, "%p", item.stream); + // args["stream"] = std::string(buff); + // args["grid_x"] = std::to_string(item.gridX); + // args["grid_y"] = std::to_string(item.gridY); + // args["grid_z"] = std::to_string(item.gridZ); + // args["block_x"] = std::to_string(item.workgroupX); + // args["block_y"] = std::to_string(item.workgroupY); + // args["block_z"] = std::to_string(item.workgroupZ); + // if (item.functionAddr != nullptr) { + // name = demangle(hipKernelNameRefByPtr(item.functionAddr, item.stream)).c_str(); + // } + // else if (item.function != nullptr) { + // name = demangle(hipKernelNameRef(item.function)).c_str(); + // } + // } + + // // Add copy args if we have them + // auto cit = copyLaunches.find(record->correlation_id); + // if (cit != copyLaunches.end()) { + // auto &item = *(*cit).second; + // snprintf(buff, BSIZE, "%p", item.stream); + // args["stream"] = std::string(buff); + // snprintf(buff, BSIZE, "%p", item.src); + // args["dst"] = std::string(buff); + // snprintf(buff, BSIZE, "%p", item.dst); + // args["src"] = std::string(buff); + // args["size"] = std::to_string(item.size); + // args["kind"] = std::to_string(item.kind); + // } + + // EventRecord event{ + // onnxruntime::profiling::KERNEL_EVENT, + // static_cast(record->device_id), + // static_cast(record->queue_id), + // name, + // static_cast((record->begin_ns - profiling_start) / 1000), + // static_cast((record->end_ns - record->begin_ns) / 1000), + // std::move(args)}; + + // // FIXME: deal with missing ext correlation + // auto extId = d->externalCorrelations_[RoctracerLogger::CorrelationDomain::Default][record->correlation_id]; + // if (event_map.find(extId) == event_map.end()) { + // event_map.insert({extId, {event}}); + // } + // else { + // event_map[extId].push_back(std::move(event)); + // } + + // roctracer_next_record(record, &record); + // } + // } + + // // General + // auto insert_iter = events.begin(); + // for (auto& map_iter : event_map) { + // auto ts = static_cast(map_iter.first); + // while (insert_iter != events.end() && insert_iter->ts < ts) { + // insert_iter++; + // } + // if (insert_iter != events.end() && insert_iter->ts == ts) { + // for (auto& evt_iter : map_iter.second) { + // evt_iter.args["op_name"] = insert_iter->args["op_name"]; + // } + // insert_iter = events.insert(insert_iter+1, map_iter.second.begin(), map_iter.second.end()); + // } else { + // insert_iter = events.insert(insert_iter, map_iter.second.begin(), map_iter.second.end()); + // } + // while (insert_iter != events.end() && insert_iter->cat == EventCategory::KERNEL_EVENT) { + // insert_iter++; + // } + // } } -void RocmProfiler::Stop(uint64_t) +void RocmProfiler::Start(uint64_t id) { - d->popCorrelationID(RoctracerLogger::CorrelationDomain::Default); + auto& manager = RoctracerManager::GetInstance(); + manager.PushCorrelation(client_handle_, id); } - -void RocmProfiler::addEventRecord(const roctracerRow &item, int64_t pstart, const std::initializer_list> &args, std::map> &event_map) +void RocmProfiler::Stop(uint64_t id) { - EventRecord event{onnxruntime::profiling::API_EVENT, - static_cast(item.pid), - static_cast(item.tid), - std::string(roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, item.cid, 0)), - static_cast((item.begin - pstart) / 1000), - static_cast((item.end - item.begin) / 1000), - {args.begin(), args.end()}}; - - // FIXME: deal with missing ext correlation - auto extId = d->externalCorrelations_[RoctracerLogger::CorrelationDomain::Default][item.id]; - if (event_map.find(extId) == event_map.end()) { - event_map.insert({extId, {event}}); - } - else { - event_map[extId].push_back(std::move(event)); - } + auto& manager = RoctracerManager::GetInstance(); + uint64_t unused; + manager.PopCorrelation(unused); } } // namespace profiling diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.h b/onnxruntime/core/providers/rocm/rocm_profiler.h index 593436a701082..f6be59bfba0ee 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.h +++ b/onnxruntime/core/providers/rocm/rocm_profiler.h @@ -7,11 +7,7 @@ #if defined(USE_ROCM) && defined(ENABLE_ROCM_PROFILING) -class RoctracerLogger; -class roctracerRow; - namespace onnxruntime { - namespace profiling { using Events = std::vector; @@ -21,27 +17,14 @@ class RocmProfiler final : public EpProfiler { RocmProfiler(); RocmProfiler(const RocmProfiler&) = delete; RocmProfiler& operator=(const RocmProfiler&) = delete; -#if 0 - RocmProfiler(RocmProfiler&& rocm_profiler) noexcept { - initialized_ = rocm_profiler.initialized_; - rocm_profiler.initialized_ = false; - } - RocmProfiler& operator=(RocmProfiler&& rocm_profiler) noexcept { - initialized_ = rocm_profiler.initialized_; - rocm_profiler.initialized_ = false; - return *this; - } -#endif ~RocmProfiler(); bool StartProfiling() override; void EndProfiling(TimePoint start_time, Events& events) override; void Start(uint64_t) override; void Stop(uint64_t) override; - private: - void addEventRecord(const roctracerRow &item, int64_t pstart, const std::initializer_list> &args, std::map> &event_map); - - RoctracerLogger *d; +private: + uint64_t client_handle_; }; } // namespace profiling diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 38f913fff8b39..5aabb360e9ddf 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -52,22 +52,19 @@ RoctracerManager::~RoctracerManager() { } uint64_t RoctracerManager::RegisterClient() { - std::lock_guard lock(global_op_mutex_); + std::lock_guard lock(roctracer_manager_mutex_); auto res = next_client_id_++; - per_client_locks_[res] = std::make_unique(); per_client_events_by_ext_correlation_.insert({res, {}}); return res; } void RoctracerManager::DeregisterClient(uint64_t client_handle) { - std::lock_guard lock(global_op_mutex_); - auto it = per_client_locks_.find(client_handle); + std::lock_guard lock(roctracer_manager_mutex_); per_client_events_by_ext_correlation_.erase(client_handle); - per_client_locks_.erase(client_handle); } void RoctracerManager::StartLogging() { - std::lock_guard lock(global_op_mutex_); + std::lock_guard lock(roctracer_manager_mutex_); if (logging_enabled_) { return; } @@ -83,7 +80,7 @@ void RoctracerManager::StartLogging() { } void RoctracerManager::StopLogging() { - std::lock_guard lock(global_op_mutex_); + std::lock_guard lock(roctracer_manager_mutex_); if (!logging_enabled_) { return; } @@ -95,7 +92,9 @@ void RoctracerManager::StopLogging() { logging_enabled_ = false; } -void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_time, Events& events) { +void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_time, + std::map& events) { + events.clear(); std::vector activity_buffers; { std::lock_guard lock(unprocessed_activity_buffers_lock_); @@ -103,10 +102,21 @@ void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_ti unprocessed_activity_buffers_.clear(); } - ProcessActivityBuffers(activity_buffers); + { + // Ensure that at most one thread is working through the activity buffers at any time. + std::lock_guard lock(activity_buffer_processor_mutex_); + ProcessActivityBuffers(activity_buffers, start_time); + std::lock_guard lock(event_list_mutex_); + auto it = per_client_events_by_ext_correlation_.find(client_handle); + if (it == per_client_events_by_ext_correlation_.end()) { + return; + } + std::swap(events, it->second); + } } bool RoctracerManager::PushCorrelation(uint64_t client_handle, uint64_t external_correlation_id) { + std::lock_guard lock(roctracer_manager_mutex_); auto it = per_client_events_by_ext_correlation_.find(client_handle); if (it == per_client_events_by_ext_correlation_.end()) { return false; @@ -119,6 +129,7 @@ bool RoctracerManager::PushCorrelation(uint64_t client_handle, uint64_t external } void RoctracerManager::PopCorrelation(uint64_t& popped_external_correlation_id) { + std::lock_guard lock(roctracer_manager_mutex_); roctracer_activity_pop_external_correlation_id(&popped_external_correlation_id); external_correlation_id_to_client_.erase(popped_external_correlation_id); } diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index d42998d6ed73f..45822575469f4 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -1,9 +1,9 @@ #pragma once #include +#include #include #include -#include #include #include @@ -85,7 +85,7 @@ class RoctracerManager void StartLogging(); void StopLogging(); - void Consume(uint64_t client_handle, const TimePoint& start_time, Events& events); + void Consume(uint64_t client_handle, const TimePoint& start_time, std::map& events); bool PushCorrelation(uint64_t client_handle, uint64_t external_correlation_id); void PopCorrelation(uint64_t& popped_correlation_id); @@ -113,12 +113,10 @@ class RoctracerManager std::mutex unprocessed_activity_buffers_lock_; std::vector unprocessed_activity_buffers_; + std::mutex activity_buffer_processor_mutex_; std::mutex api_call_args_lock_; std::unordered_map api_call_args_; - // Keyed on client_id / client_handle - std::unordered_map> per_client_locks_; - // Keyed on external_correlation_id -> client_id/client_handle std::unordered_map external_correlation_id_to_client_; @@ -126,10 +124,11 @@ class RoctracerManager std::unordered_map roctracer_correlation_to_external_correlation_; // client_id/client_handle -> external_correlation_id -> events - std::unordered_map> per_client_events_by_ext_correlation_; + std::mutex event_list_mutex_; + std::unordered_map> per_client_events_by_ext_correlation_; uint64_t next_client_id_ = 1; bool logging_enabled_ = false; - std::mutex global_op_mutex_; + std::mutex roctracer_manager_mutex_; roctracer_pool_t* activity_pool_; }; From e06f9fc5255d9cd3913d7a9d6cdfd562a3c849a0 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 31 Oct 2022 22:01:03 +0000 Subject: [PATCH 04/41] Fixes --- include/onnxruntime/core/common/profiler_common.h | 4 ++++ onnxruntime/core/common/profiler.cc | 4 ++-- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/include/onnxruntime/core/common/profiler_common.h b/include/onnxruntime/core/common/profiler_common.h index 58e20fdb878a8..ccbc350a79b58 100644 --- a/include/onnxruntime/core/common/profiler_common.h +++ b/include/onnxruntime/core/common/profiler_common.h @@ -58,6 +58,10 @@ struct EventRecord { dur(duration), args(event_args) {} + EventRecord(const EventRecord& other) + : cat(other.cat), pid(other.pid), tid(other.tid), name(other.name), + ts(other.ts), dur(other.dur), args(other.args) {} + EventRecord(EventRecord&& other) : cat(other.cat), pid(other.pid), tid(other.tid), name(std::move(other.name)), ts(other.ts), dur(other.dur), args(std::move(other.args)) {} diff --git a/onnxruntime/core/common/profiler.cc b/onnxruntime/core/common/profiler.cc index 05900fe126b4c..b64b0b9ebade8 100644 --- a/onnxruntime/core/common/profiler.cc +++ b/onnxruntime/core/common/profiler.cc @@ -25,7 +25,7 @@ ::onnxruntime::TimePoint profiling::Profiler::Start() { auto ts = TimeDiffMicroSeconds(profiling_start_time_, start_time); for (const auto& ep_profiler : ep_profilers_) { ep_profiler->Start(ts); - } + } return start_time; } @@ -87,7 +87,7 @@ void Profiler::EndTimeAndRecordEvent(EventCategory category, //TODO: sync_gpu if needed. std::lock_guard lock(mutex_); if (events_.size() < max_num_events_) { - events_.emplace_back(event); + events_.emplace_back(std::move(event)); } else { if (session_logger_ && !max_events_reached) { LOGS(*session_logger_, ERROR) From e908085e60cae7bbd92d9b4af8df00dbf280c8d4 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 31 Oct 2022 22:02:00 +0000 Subject: [PATCH 05/41] Fixes --- onnxruntime/core/common/profiler_common.cc | 1 + 1 file changed, 1 insertion(+) diff --git a/onnxruntime/core/common/profiler_common.cc b/onnxruntime/core/common/profiler_common.cc index 25d64970a287c..c286c844456ce 100644 --- a/onnxruntime/core/common/profiler_common.cc +++ b/onnxruntime/core/common/profiler_common.cc @@ -72,6 +72,7 @@ EventRecord& EventRecord::operator = (EventRecord&& other) { std::swap(name, other.name); dur = other.dur; std::swap(args, other.args); + return *this; } } // namespace profiling From 41a973c1162c323d4cf88b50e3581f1825672445 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 31 Oct 2022 22:07:09 +0000 Subject: [PATCH 06/41] Fixes --- onnxruntime/core/providers/rocm/rocm_profiler.cc | 1 + onnxruntime/core/providers/rocm/roctracer_manager.cc | 4 ++-- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.cc b/onnxruntime/core/providers/rocm/rocm_profiler.cc index 919cad17fb223..29079b5001b52 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.cc +++ b/onnxruntime/core/providers/rocm/rocm_profiler.cc @@ -33,6 +33,7 @@ bool RocmProfiler::StartProfiling() { auto& manager = RoctracerManager::GetInstance(); manager.StartLogging(); + return true; } void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 5aabb360e9ddf..1225f906aacc5 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -104,7 +104,7 @@ void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_ti { // Ensure that at most one thread is working through the activity buffers at any time. - std::lock_guard lock(activity_buffer_processor_mutex_); + std::lock_guard lock_two(activity_buffer_processor_mutex_); ProcessActivityBuffers(activity_buffers, start_time); std::lock_guard lock(event_list_mutex_); auto it = per_client_events_by_ext_correlation_.find(client_handle); @@ -172,7 +172,7 @@ void RoctracerManager::CreateEventForKernelRecord(const roctracer_record_t* reco auto name = demangle(hipKernelNameRefByPtr(launch_args.args.hipLaunchKernel.function_address, launch_args.args.hipLaunchKernel.stream)); std::unordered_map args { - {"stream", std::to_string(launch_args.args.hipLaunchKernel.stream)}, + {"stream", PointerToHexString((void*)(launch_args.args.hipLaunchKernel.stream))}, {"grid_x", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.x)}, {"grid_y", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.y)}, {"grid_z", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.z)}, From b660fc742bffac41e325028b79232b42a11f107d Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 31 Oct 2022 22:37:08 +0000 Subject: [PATCH 07/41] Fixes --- .../core/providers/rocm/roctracer_manager.cc | 50 +++++++++---------- 1 file changed, 25 insertions(+), 25 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 1225f906aacc5..ecac69c3f4a91 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -164,6 +164,13 @@ void RoctracerManager::ApiCallback(uint32_t domain, uint32_t cid, const void* ca } } +static inline std::string PointerToHexString(const void* ptr) +{ + std::ostringstream sstr; + sstr << std::hex << ptr; + return sstr.str(); +} + void RoctracerManager::CreateEventForKernelRecord(const roctracer_record_t* record, uint64_t start_time_ns, const ApiCallRecord& call_record, @@ -183,22 +190,15 @@ void RoctracerManager::CreateEventForKernelRecord(const roctracer_record_t* reco new (&event) EventRecord { /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ record->device_id, - /* tid = */ record->queue_id, + /* pid = */ -1, + /* tid = */ -1, /* name = */ std::move(name), - /* ts = */ (record->begin_ns - start_time_ns) / 1000, - /* dur = */ (record->end_ns - record->begin_ns) / 1000, + /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, + /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, /* args = */ std::move(args) }; } -static inline std::string PointerToHexString(const void* ptr) -{ - std::ostringstream sstr; - sstr << std::hex << ptr; - return sstr.str(); -} - void RoctracerManager::CreateEventForMemsetRecord(const roctracer_record_t* record, uint64_t start_time_ns, const ApiCallRecord& call_record, @@ -210,24 +210,24 @@ void RoctracerManager::CreateEventForMemsetRecord(const roctracer_record_t* reco std::unordered_map args { {"stream", call_record.cid_ == HIP_API_ID_hipMemset ? "0" - : std::to_string(launch_args.args.hipMemsetAsync.stream)}, + : PointerToHexString(launch_args.args.hipMemsetAsync.stream)}, {"dst", dst_string}, {"size", std::to_string(launch_args.args.hipMemset.sizeBytes)}, {"value", std::to_string(launch_args.args.hipMemset.value)} }; new (&event) EventRecord { /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ record->device_id, - /* tid = */ record->queue_id, + /* pid = */ -1, + /* tid = */ -1, /* name = */ std::move(name), - /* ts = */ (record->begin_ns - start_time_ns) / 1000, - /* dur = */ (record->end_ns - record->begin_ns) / 1000, + /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, + /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, /* args = */ std::move(args) }; } static inline std::string MemcpyKindToString(hipMemcpyKind kind) { - switch(launch_args.args.hipMemcpy.kind) { + switch(kind) { case hipMemcpyHostToHost: return "H2H"; case hipMemcpyHostToDevice: @@ -261,11 +261,11 @@ void RoctracerManager::CreateEventForMemcpyRecord(const roctracer_record_t* reco new (&event) EventRecord { /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ record->device_id, - /* tid = */ record->queue_id, + /* pid = */ -1, + /* tid = */ -1, /* name = */ std::move(name), - /* ts = */ (record->begin_ns - start_time_ns) / 1000, - /* dur = */ (record->end_ns - record->begin_ns) / 1000, + /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, + /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, /* args = */ std::move(args) }; } @@ -294,11 +294,11 @@ void RoctracerManager::CreateEventForMemcpy2DRecord(const roctracer_record_t* re new (&event) EventRecord { /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ record->device_id, - /* tid = */ record->queue_id, + /* pid = */ -1, + /* tid = */ -1, /* name = */ std::move(name), - /* ts = */ (record->begin_ns - start_time_ns) / 1000, - /* dur = */ (record->end_ns - record->begin_ns) / 1000, + /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, + /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, /* args = */ std::move(args) }; } From 8dc917d75687fe23f223bb30e0262cef1d430673 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 31 Oct 2022 22:40:05 +0000 Subject: [PATCH 08/41] Delete unnecessary code --- .../core/providers/rocm/rocm_profiler.cc | 165 ------------------ 1 file changed, 165 deletions(-) diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.cc b/onnxruntime/core/providers/rocm/rocm_profiler.cc index 29079b5001b52..8708154a136bf 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.cc +++ b/onnxruntime/core/providers/rocm/rocm_profiler.cc @@ -9,13 +9,6 @@ #include "core/providers/rocm/rocm_profiler.h" #include "core/providers/rocm/roctracer_manager.h" -#define BSIZE 4096 - -typedef uint64_t timestamp_t; -static timestamp_t timespec_to_ns(const timespec& time) { - return ((timestamp_t)time.tv_sec * 1000000000) + time.tv_nsec; -} - namespace onnxruntime { namespace profiling { @@ -69,164 +62,6 @@ void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) // move any remaining events merged_events.insert(merged_events.end(), event_iter, event_end); std::swap(events, merged_events); - - // // Generate EventRecords - // int64_t profiling_start = std::chrono::duration_cast(start_time.time_since_epoch()).count(); - - // // Wrong clock again - all the cool kids are doing it - // timespec t0, t1, t00; - // clock_gettime(CLOCK_REALTIME, &t0); - // clock_gettime(CLOCK_MONOTONIC, &t1); - // clock_gettime(CLOCK_REALTIME, &t00); - // const uint64_t toffset = (timespec_to_ns(t0) >> 1) + (timespec_to_ns(t00) >> 1) - timespec_to_ns(t1); - // profiling_start = profiling_start - toffset; - - // char buff[BSIZE]; - - // for (auto &item : d->rows_) { - // std::initializer_list> args = {{"op_name", ""}}; - // addEventRecord(item, profiling_start, args, event_map); - // } - - // for (auto &item : d->mallocRows_) { - // snprintf(buff, BSIZE, "%p", item.ptr); - // const std::string arg_ptr{buff}; - - // std::initializer_list> args = {{"op_name", ""}, - // {"ptr", arg_ptr}, - // {"size", std::to_string(item.size)} - // }; - // addEventRecord(item, profiling_start, args, event_map); - // } - - // for (auto &item : d->copyRows_) { - // snprintf(buff, BSIZE, "%p", item.stream); - // const std::string arg_stream{buff}; - // snprintf(buff, BSIZE, "%p", item.src); - // const std::string arg_src{buff}; - // snprintf(buff, BSIZE, "%p", item.dst); - // const std::string arg_dst{buff}; - - // std::initializer_list> args = {{"op_name", ""}, - // {"stream", arg_stream}, - // {"src", arg_src}, - // {"dst", arg_dst}, - // {"size", std::to_string(item.size)}, - // {"kind", std::to_string(item.kind)}, - // }; - // addEventRecord(item, profiling_start, args, event_map); - // copyLaunches[item.id] = &item; - // } - - // for (auto &item : d->kernelRows_) { - // snprintf(buff, BSIZE, "%p", item.stream); - // const std::string arg_stream{buff}; - // if (item.functionAddr) - // snprintf(buff, BSIZE, "%s", demangle(hipKernelNameRefByPtr(item.functionAddr, item.stream)).c_str()); - // else if (item.function) - // snprintf(buff, BSIZE, "%s", demangle(hipKernelNameRef(item.function)).c_str()); - // else - // snprintf(buff, BSIZE, " "); - // const std::string arg_kernel{buff}; - - // std::initializer_list> args = {{"op_name", ""}, - // {"stream", arg_stream}, - // {"kernel", arg_kernel}, - // {"grid_x", std::to_string(item.gridX)}, - // {"grid_y", std::to_string(item.gridY)}, - // {"grid_z", std::to_string(item.gridZ)}, - // {"block_x", std::to_string(item.workgroupX)}, - // {"block_y", std::to_string(item.workgroupY)}, - // {"block_z", std::to_string(item.workgroupZ)}, - // }; - // addEventRecord(item, profiling_start, args, event_map); - // kernelLaunches[item.id] = &item; - // } - - // // Async Ops - e.g. "Kernel" - - // for (auto& buffer : *d->gpuTraceBuffers_) { - // const roctracer_record_t* record = (const roctracer_record_t*)(buffer.data_); - // const roctracer_record_t* end_record = (const roctracer_record_t*)(buffer.data_ + buffer.validSize_); - // while (record < end_record) { - // std::unordered_map args; - // std::string name = roctracer_op_string(record->domain, record->op, record->kind); - - // // Add kernel args if we have them - // auto kit = kernelLaunches.find(record->correlation_id); - // if (kit != kernelLaunches.end()) { - // auto &item = *(*kit).second; - // snprintf(buff, BSIZE, "%p", item.stream); - // args["stream"] = std::string(buff); - // args["grid_x"] = std::to_string(item.gridX); - // args["grid_y"] = std::to_string(item.gridY); - // args["grid_z"] = std::to_string(item.gridZ); - // args["block_x"] = std::to_string(item.workgroupX); - // args["block_y"] = std::to_string(item.workgroupY); - // args["block_z"] = std::to_string(item.workgroupZ); - // if (item.functionAddr != nullptr) { - // name = demangle(hipKernelNameRefByPtr(item.functionAddr, item.stream)).c_str(); - // } - // else if (item.function != nullptr) { - // name = demangle(hipKernelNameRef(item.function)).c_str(); - // } - // } - - // // Add copy args if we have them - // auto cit = copyLaunches.find(record->correlation_id); - // if (cit != copyLaunches.end()) { - // auto &item = *(*cit).second; - // snprintf(buff, BSIZE, "%p", item.stream); - // args["stream"] = std::string(buff); - // snprintf(buff, BSIZE, "%p", item.src); - // args["dst"] = std::string(buff); - // snprintf(buff, BSIZE, "%p", item.dst); - // args["src"] = std::string(buff); - // args["size"] = std::to_string(item.size); - // args["kind"] = std::to_string(item.kind); - // } - - // EventRecord event{ - // onnxruntime::profiling::KERNEL_EVENT, - // static_cast(record->device_id), - // static_cast(record->queue_id), - // name, - // static_cast((record->begin_ns - profiling_start) / 1000), - // static_cast((record->end_ns - record->begin_ns) / 1000), - // std::move(args)}; - - // // FIXME: deal with missing ext correlation - // auto extId = d->externalCorrelations_[RoctracerLogger::CorrelationDomain::Default][record->correlation_id]; - // if (event_map.find(extId) == event_map.end()) { - // event_map.insert({extId, {event}}); - // } - // else { - // event_map[extId].push_back(std::move(event)); - // } - - // roctracer_next_record(record, &record); - // } - // } - - // // General - // auto insert_iter = events.begin(); - // for (auto& map_iter : event_map) { - // auto ts = static_cast(map_iter.first); - // while (insert_iter != events.end() && insert_iter->ts < ts) { - // insert_iter++; - // } - // if (insert_iter != events.end() && insert_iter->ts == ts) { - // for (auto& evt_iter : map_iter.second) { - // evt_iter.args["op_name"] = insert_iter->args["op_name"]; - // } - // insert_iter = events.insert(insert_iter+1, map_iter.second.begin(), map_iter.second.end()); - // } else { - // insert_iter = events.insert(insert_iter, map_iter.second.begin(), map_iter.second.end()); - // } - // while (insert_iter != events.end() && insert_iter->cat == EventCategory::KERNEL_EVENT) { - // insert_iter++; - // } - // } } void RocmProfiler::Start(uint64_t id) From 367fcfc571061fccf325b7c8b08ac0c25c96d537 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 31 Oct 2022 22:41:41 +0000 Subject: [PATCH 09/41] Minor fixes --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index ecac69c3f4a91..e561f978213b2 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -210,7 +210,7 @@ void RoctracerManager::CreateEventForMemsetRecord(const roctracer_record_t* reco std::unordered_map args { {"stream", call_record.cid_ == HIP_API_ID_hipMemset ? "0" - : PointerToHexString(launch_args.args.hipMemsetAsync.stream)}, + : PointerToHexString((void*)launch_args.args.hipMemsetAsync.stream)}, {"dst", dst_string}, {"size", std::to_string(launch_args.args.hipMemset.sizeBytes)}, {"value", std::to_string(launch_args.args.hipMemset.value)} @@ -253,7 +253,7 @@ void RoctracerManager::CreateEventForMemcpyRecord(const roctracer_record_t* reco std::unordered_map args { {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy ? "0" - : std::to_string(launch_args.args.hipMemcpyAsync.stream)}, + : PointerToHexString((void*)launch_args.args.hipMemcpyAsync.stream)}, {"src", src_string}, {"dst", dst_string}, {"kind", memcpy_kind_string} @@ -282,7 +282,7 @@ void RoctracerManager::CreateEventForMemcpy2DRecord(const roctracer_record_t* re std::unordered_map args { {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy2D ? "0" - : std::to_string(launch_args.args.hipMemcpy2DAsync.stream)}, + : PointerToHexString((void*)launch_args.args.hipMemcpy2DAsync.stream)}, {"src", src_string}, {"dst", dst_string}, {"spitch", std::to_string(launch_args.args.hipMemcpy2D.spitch)}, From 2ae9dad6be7104407fa24b17802371940c2af73d Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 31 Oct 2022 23:15:30 +0000 Subject: [PATCH 10/41] Fixes + startup --- .../core/providers/rocm/roctracer_manager.cc | 29 +++++++++++++++++-- .../core/providers/rocm/roctracer_manager.h | 15 +++------- 2 files changed, 31 insertions(+), 13 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index e561f978213b2..3cca38594d79f 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -10,6 +10,16 @@ namespace profiling { // allocate a 16K buffer for recording async activities static constexpr size_t sc_activity_buffer_size = 0x4000; +const std::vector RoctracerManager::hip_api_calls_to_trace = { + "hipMemcpy", + "hipMemcpy2D", + "hipMemcpyAsync", + "hipMemcpy2DAsync", + "hipLaunchKernel", + "hipMemset", + "hipMemsetAsync", +}; + // Implementation of RoctracerActivityBuffer RoctracerActivityBuffer& RoctracerActivityBuffer::operator = (const RoctracerActivityBuffer& other) { if (&other == this) { @@ -31,6 +41,7 @@ RoctracerActivityBuffer& RoctracerActivityBuffer::operator = (RoctracerActivityB } std::swap(data_, other.data_); std::swap(size_, other.size_); + return *this; } RoctracerActivityBuffer::~RoctracerActivityBuffer() { @@ -69,12 +80,24 @@ void RoctracerManager::StartLogging() { return; } + // Enable selective activity and API callbacks for the HIP APIs + roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); + roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_API); + + for (auto const& logged_api : hip_api_calls_to_trace) { + uint32_t cid = 0; + roctracer_op_code(ACTIVITY_DOMAIN_HIP_API, logged_api.c_str(), &cid, nullptr); + roctracer_enable_op_callback(ACTIVITY_DOMAIN_HIP_API, cid, ApiCallback, nullptr); + roctracer_enable_op_activity(ACTIVITY_DOMAIN_HIP_API, cid); + } + + // Enable activity logging in the HCC domain. roctracer_properties_t hcc_cb_properties; memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); hcc_cb_properties.buffer_size = sc_activity_buffer_size; hcc_cb_properties.buffer_callback_fun = ActivityCallback; roctracer_open_pool_expl(&hcc_cb_properties, &activity_pool_); - roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HCC_OPS, activity_pool_); + roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HIP_OPS, activity_pool_); roctracer_start(); logging_enabled_ = true; } @@ -87,8 +110,10 @@ void RoctracerManager::StopLogging() { roctracer_stop(); roctracer_flush_activity_expl(activity_pool_); - roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS); + roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_API); + roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS); roctracer_close_pool_expl(activity_pool_); + roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); logging_enabled_ = false; } diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 45822575469f4..9d2b1af392f7b 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -54,20 +54,10 @@ class RoctracerActivityBuffer { size_t size_; }; -static const char* hip_api_calls_to_trace[] = { - "hipMemcpy", - "hipMemcpy2D", - "hipMemcpyAsync", - "hipMemcpy2DAsync", - "hipLaunchKernel", - "hipMemset", - "hipMemsetAsync", -}; - struct ApiCallRecord { uint32_t domain_; uint32_t cid_; - hip_api_data_t api_data_; + hip_api_data_t api_data_ {}; }; class RoctracerManager @@ -130,6 +120,9 @@ class RoctracerManager bool logging_enabled_ = false; std::mutex roctracer_manager_mutex_; roctracer_pool_t* activity_pool_; + + // The api calls to track + static const std::vector hip_api_calls_to_trace; }; } /* end namespace profiling */ From bd3e7a7345cf1f1a09a59ab68aa26f5708c37ad9 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 00:02:53 +0000 Subject: [PATCH 11/41] Fixes --- .../onnxruntime/core/common/profiler_common.h | 28 +++++++++++++++++-- onnxruntime/core/common/profiler_common.cc | 27 ------------------ .../core/providers/rocm/roctracer_manager.cc | 16 +++++------ 3 files changed, 34 insertions(+), 37 deletions(-) diff --git a/include/onnxruntime/core/common/profiler_common.h b/include/onnxruntime/core/common/profiler_common.h index ccbc350a79b58..f0e55b5093bce 100644 --- a/include/onnxruntime/core/common/profiler_common.h +++ b/include/onnxruntime/core/common/profiler_common.h @@ -66,8 +66,32 @@ struct EventRecord { : cat(other.cat), pid(other.pid), tid(other.tid), name(std::move(other.name)), ts(other.ts), dur(other.dur), args(std::move(other.args)) {} - EventRecord& operator = (const EventRecord& other); - EventRecord& operator = (EventRecord&& other); + EventRecord& operator = (const EventRecord& other) { + if (&other == this) { + return *this; + } + cat = other.cat; + pid = other.pid; + tid = other.tid; + name = other.name; + ts = other.ts; + dur = other.dur; + args = other.args; + return *this; + } + + EventRecord& operator = (EventRecord&& other) { + if (&other == this) { + return *this; + } + cat = other.cat; + pid = other.pid; + tid = other.tid; + std::swap(name, other.name); + dur = other.dur; + std::swap(args, other.args); + return *this; + } EventCategory cat; int pid; diff --git a/onnxruntime/core/common/profiler_common.cc b/onnxruntime/core/common/profiler_common.cc index c286c844456ce..03a06adcfff2f 100644 --- a/onnxruntime/core/common/profiler_common.cc +++ b/onnxruntime/core/common/profiler_common.cc @@ -48,32 +48,5 @@ std::string demangle(const std::string& name) { return demangle(name.c_str()); } -EventRecord& EventRecord::operator = (const EventRecord& other) { - if (&other == this) { - return *this; - } - cat = other.cat; - pid = other.pid; - tid = other.tid; - name = other.name; - ts = other.ts; - dur = other.dur; - args = other.args; - return *this; -} - -EventRecord& EventRecord::operator = (EventRecord&& other) { - if (&other == this) { - return *this; - } - cat = other.cat; - pid = other.pid; - tid = other.tid; - std::swap(name, other.name); - dur = other.dur; - std::swap(args, other.args); - return *this; -} - } // namespace profiling } // namespace onnxruntime diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 3cca38594d79f..66d1834d2886f 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -51,7 +51,6 @@ RoctracerActivityBuffer::~RoctracerActivityBuffer() { } } - // Implementation of RoctracerManager RoctracerManager& RoctracerManager::GetInstance() { static RoctracerManager instance; @@ -80,6 +79,14 @@ void RoctracerManager::StartLogging() { return; } + // Enable activity logging in the HIP_OPS/HCC_OPS domain. + roctracer_properties_t hcc_cb_properties; + memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); + hcc_cb_properties.buffer_size = sc_activity_buffer_size; + hcc_cb_properties.buffer_callback_fun = ActivityCallback; + roctracer_open_pool_expl(&hcc_cb_properties, &activity_pool_); + roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HIP_OPS, activity_pool_); + // Enable selective activity and API callbacks for the HIP APIs roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_API); @@ -91,13 +98,6 @@ void RoctracerManager::StartLogging() { roctracer_enable_op_activity(ACTIVITY_DOMAIN_HIP_API, cid); } - // Enable activity logging in the HCC domain. - roctracer_properties_t hcc_cb_properties; - memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); - hcc_cb_properties.buffer_size = sc_activity_buffer_size; - hcc_cb_properties.buffer_callback_fun = ActivityCallback; - roctracer_open_pool_expl(&hcc_cb_properties, &activity_pool_); - roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HIP_OPS, activity_pool_); roctracer_start(); logging_enabled_ = true; } From f1f53b09c82a3b5f852c0cbc9f9cba9fb108556d Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 01:56:10 +0000 Subject: [PATCH 12/41] Fixes --- .../providers/rocm/rocm_execution_provider.cc | 17 +++++++++++------ .../providers/rocm/rocm_execution_provider.h | 4 +++- .../core/providers/rocm/roctracer_manager.cc | 15 ++++++++++----- .../core/providers/rocm/roctracer_manager.h | 1 - 4 files changed, 24 insertions(+), 13 deletions(-) diff --git a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc index ed8a9afd42261..1a9ad3d6a85b9 100644 --- a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc +++ b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc @@ -376,7 +376,7 @@ void ReleaseCpuBufferCallback(hipStream_t /*stream*/, hipError_t /*status*/, voi } } -Status ROCMExecutionProvider::EnqueueDeferredRelease() { +Status ROCMExecutionProvider::EnqueueDeferredRelease(bool actually_defer) { // Release CPU buffers allocated for GPU kernels (type: RocmKernel). // They have to be released outside GPU kernels because they must be alive // during asynchronous GPU computation even after the CPU part (e.g, @@ -404,8 +404,12 @@ Status ROCMExecutionProvider::EnqueueDeferredRelease() { // we should replace the following line with // hipLaunchHostFunc(stream, ReleaseCpuBufferCallback, cpu_buffers_info); if (cpu_buffers_info->allocator->Info().alloc_type == OrtArenaAllocator) { - // Release memory asynchronously to avoid blocking the compute stream. - HIP_RETURN_IF_ERROR(hipStreamAddCallback(stream, ReleaseCpuBufferCallback, cpu_buffers_info.release(), 0)); + if (actually_defer) { + // Release memory asynchronously to avoid blocking the compute stream. + HIP_RETURN_IF_ERROR(hipStreamAddCallback(stream, ReleaseCpuBufferCallback, cpu_buffers_info.release(), 0)); + } else { + ReleaseCpuBufferCallback(nullptr, hipSuccess, cpu_buffers_info.release()); + } } else { // Per @@ -436,9 +440,10 @@ Status ROCMExecutionProvider::OnRunEnd(bool sync_stream) { // Enqueue deferred CPU memory release on related streams. // This will release all deferred-release CPU buffers allocated // before calling OnRunEnd. - ORT_RETURN_IF_ERROR(EnqueueDeferredRelease()); - - if (sync_stream) { + if (!sync_stream) { + ORT_RETURN_IF_ERROR(EnqueueDeferredRelease()); + } else { + ORT_RETURN_IF_ERROR(EnqueueDeferredRelease(/* actually_defer = */ false)); HIP_RETURN_IF_ERROR(hipStreamSynchronize(static_cast(GetComputeStream()))); } diff --git a/onnxruntime/core/providers/rocm/rocm_execution_provider.h b/onnxruntime/core/providers/rocm/rocm_execution_provider.h index fe15bb441fe1c..29569234857fb 100644 --- a/onnxruntime/core/providers/rocm/rocm_execution_provider.h +++ b/onnxruntime/core/providers/rocm/rocm_execution_provider.h @@ -17,6 +17,8 @@ namespace onnxruntime { +void RunOnUnload(std::function function); + // Logical device representation. class ROCMExecutionProvider : public IExecutionProvider { public: @@ -69,7 +71,7 @@ class ROCMExecutionProvider : public IExecutionProvider { // Release all buffers added by // AddDeferredReleaseCPUPtr using // GPU callback (so it's async). - Status EnqueueDeferredRelease(); + Status EnqueueDeferredRelease(bool actually_defer = true); template IAllocatorUniquePtr GetScratchBuffer(size_t count_or_bytes) const { diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 66d1834d2886f..74ceeda9915a9 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -79,13 +79,15 @@ void RoctracerManager::StartLogging() { return; } - // Enable activity logging in the HIP_OPS/HCC_OPS domain. + // The following line shows up in all the samples, I do not know + // what the point is, but without it, the roctracer APIs don't work. + roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, nullptr); + roctracer_properties_t hcc_cb_properties; memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); hcc_cb_properties.buffer_size = sc_activity_buffer_size; hcc_cb_properties.buffer_callback_fun = ActivityCallback; - roctracer_open_pool_expl(&hcc_cb_properties, &activity_pool_); - roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HIP_OPS, activity_pool_); + roctracer_open_pool(&hcc_cb_properties); // Enable selective activity and API callbacks for the HIP APIs roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); @@ -98,6 +100,9 @@ void RoctracerManager::StartLogging() { roctracer_enable_op_activity(ACTIVITY_DOMAIN_HIP_API, cid); } + // Enable activity logging in the HIP_OPS/HCC_OPS domain. + roctracer_enable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS); + roctracer_start(); logging_enabled_ = true; } @@ -109,10 +114,10 @@ void RoctracerManager::StopLogging() { } roctracer_stop(); - roctracer_flush_activity_expl(activity_pool_); + roctracer_flush_activity(); roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_API); roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS); - roctracer_close_pool_expl(activity_pool_); + roctracer_close_pool(); roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); logging_enabled_ = false; } diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 9d2b1af392f7b..cef717b2acc29 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -119,7 +119,6 @@ class RoctracerManager uint64_t next_client_id_ = 1; bool logging_enabled_ = false; std::mutex roctracer_manager_mutex_; - roctracer_pool_t* activity_pool_; // The api calls to track static const std::vector hip_api_calls_to_trace; From 4c8e46813773548a413768ae5518eb1e2da614ec Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 02:22:39 +0000 Subject: [PATCH 13/41] Fixes --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 1 - 1 file changed, 1 deletion(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 74ceeda9915a9..410d836e391d8 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -161,7 +161,6 @@ bool RoctracerManager::PushCorrelation(uint64_t client_handle, uint64_t external void RoctracerManager::PopCorrelation(uint64_t& popped_external_correlation_id) { std::lock_guard lock(roctracer_manager_mutex_); roctracer_activity_pop_external_correlation_id(&popped_external_correlation_id); - external_correlation_id_to_client_.erase(popped_external_correlation_id); } void RoctracerManager::ActivityCallback(const char* begin, const char* end, void* arg) { From 69c7c8eab67d1e1eca455f3eeded26d5ed8051f0 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 20:08:34 +0000 Subject: [PATCH 14/41] Add tracking for hipExtModuleLaunchKernel --- .../core/providers/rocm/roctracer_manager.cc | 35 +++++++++++++++++-- .../core/providers/rocm/roctracer_manager.h | 2 ++ 2 files changed, 35 insertions(+), 2 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 410d836e391d8..917f4a32495bf 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -15,9 +15,11 @@ const std::vector RoctracerManager::hip_api_calls_to_trace = { "hipMemcpy2D", "hipMemcpyAsync", "hipMemcpy2DAsync", + "hipMemcpyWithStream" "hipLaunchKernel", "hipMemset", "hipMemsetAsync", + "hipExtModuleLaunchKernel" }; // Implementation of RoctracerActivityBuffer @@ -193,8 +195,7 @@ void RoctracerManager::ApiCallback(uint32_t domain, uint32_t cid, const void* ca } } -static inline std::string PointerToHexString(const void* ptr) -{ +static inline std::string PointerToHexString(const void* ptr) { std::ostringstream sstr; sstr << std::hex << ptr; return sstr.str(); @@ -332,6 +333,31 @@ void RoctracerManager::CreateEventForMemcpy2DRecord(const roctracer_record_t* re }; } +void RoctracerManager::CreateEventForExtModuleLaunchKernel(const roctracer_record_t* record, uint64_t start_time_ns, + const ApiCallRecord& call_record, EventRecord& event) { + auto const& launch_args = call_record.api_data_; + auto name = hipKernelNameRef(launch_args.args.hipExtModuleLaunchKernel.f); + std::unordered_map args { + {"stream", PointerToHexString((void*)launch_args.args.hipExtModuleLaunchKernel.hStream)}, + {"grid_x", std::to_string(launch_args.args.hipExtModuleLaunchKernel.globalWorkSizeX)}, + {"grid_y", std::to_string(launch_args.args.hipExtModuleLaunchKernel.globalWorkSizeY)}, + {"grid_z", std::to_string(launch_args.args.hipExtModuleLaunchKernel.globalWorkSizeZ)}, + {"block_x", std::to_string(launch_args.args.hipExtModuleLaunchKernel.localWorkSizeX)}, + {"block_y", std::to_string(launch_args.args.hipExtModuleLaunchKernel.localWorkSizeY)}, + {"block_z", std::to_string(launch_args.args.hipExtModuleLaunchKernel.localWorkSizeZ)}, + }; + + new (&event) EventRecord { + /* cat = */ EventCategory::KERNEL_EVENT, + /* pid = */ -1, + /* tid = */ -1, + /* name = */ std::move(name), + /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, + /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, + /* args = */ std::move(args) + }; +} + void RoctracerManager::MapEventsToClient(uint64_t external_correlation_id, Events&& events) { auto client_it = external_correlation_id_to_client_.find(external_correlation_id); if (client_it == external_correlation_id_to_client_.end()) { @@ -402,6 +428,7 @@ void RoctracerManager::ProcessActivityBuffers(const std::vector Date: Tue, 1 Nov 2022 21:43:44 +0000 Subject: [PATCH 15/41] Fixes + ts adjustments --- .../core/providers/rocm/rocm_profiler.cc | 19 ++- .../core/providers/rocm/roctracer_manager.cc | 20 +++ .../core/providers/rocm/roctracer_manager.h | 140 +----------------- 3 files changed, 34 insertions(+), 145 deletions(-) diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.cc b/onnxruntime/core/providers/rocm/rocm_profiler.cc index 8708154a136bf..b9d7e8efc7406 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.cc +++ b/onnxruntime/core/providers/rocm/rocm_profiler.cc @@ -22,15 +22,13 @@ RocmProfiler::~RocmProfiler() { manager.DeregisterClient(client_handle_); } -bool RocmProfiler::StartProfiling() -{ +bool RocmProfiler::StartProfiling() { auto& manager = RoctracerManager::GetInstance(); manager.StartLogging(); return true; } -void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) -{ +void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) { auto& manager = RoctracerManager::GetInstance(); std::map event_map; manager.Consume(client_handle_, start_time, event_map); @@ -47,8 +45,15 @@ void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) } if (event_iter != event_end && event_iter->ts == ts) { + uint64_t increment = 1; for (auto& evt : map_iter.second) { evt.args["op_name"] = event_iter->args["op_name"]; + + // roctracer timestamps don't use Jan 1 1970 as an epoch, + // not sure what epoch it uses, but we adjust the timestamp + // here to something sensible. + evt.ts = event_iter->ts + increment; + ++increment; } merged_events.emplace_back(*event_iter); ++event_iter; @@ -64,14 +69,12 @@ void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) std::swap(events, merged_events); } -void RocmProfiler::Start(uint64_t id) -{ +void RocmProfiler::Start(uint64_t id) { auto& manager = RoctracerManager::GetInstance(); manager.PushCorrelation(client_handle_, id); } -void RocmProfiler::Stop(uint64_t id) -{ +void RocmProfiler::Stop(uint64_t id) { auto& manager = RoctracerManager::GetInstance(); uint64_t unused; manager.PopCorrelation(unused); diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 917f4a32495bf..23c0ba54e956b 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -67,12 +67,17 @@ uint64_t RoctracerManager::RegisterClient() { std::lock_guard lock(roctracer_manager_mutex_); auto res = next_client_id_++; per_client_events_by_ext_correlation_.insert({res, {}}); + ++num_active_clients_; return res; } void RoctracerManager::DeregisterClient(uint64_t client_handle) { std::lock_guard lock(roctracer_manager_mutex_); per_client_events_by_ext_correlation_.erase(client_handle); + --num_active_clients_; + if (num_active_clients_ == 0) { + StopLogging(); + } } void RoctracerManager::StartLogging() { @@ -109,6 +114,15 @@ void RoctracerManager::StartLogging() { logging_enabled_ = true; } +void RoctracerManager::Clear() +{ + unprocessed_activity_buffers_.clear(); + api_call_args_.clear(); + external_correlation_id_to_client_.clear(); + roctracer_correlation_to_external_correlation_.clear(); + per_client_events_by_ext_correlation_.clear(); +} + void RoctracerManager::StopLogging() { std::lock_guard lock(roctracer_manager_mutex_); if (!logging_enabled_) { @@ -121,7 +135,9 @@ void RoctracerManager::StopLogging() { roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS); roctracer_close_pool(); roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); + logging_enabled_ = false; + Clear(); } void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_time, @@ -444,6 +460,10 @@ void RoctracerManager::ProcessActivityBuffers(const std::vector> per_client_events_by_ext_correlation_; uint64_t next_client_id_ = 1; + uint64_t num_active_clients_ = 0; bool logging_enabled_ = false; std::mutex roctracer_manager_mutex_; // The api calls to track static const std::vector hip_api_calls_to_trace; }; - -} /* end namespace profiling */ -} /* end namespace onnxruntime */ - - -/* -} -} - -class RoctracerActivityBuffer { -public: - // data must be allocated using malloc. - // Ownership is transferred to this object. - RoctracerActivityBuffer(uint8_t* data, size_t validSize) - : data_(data), validSize_(validSize) {} - - ~RoctracerActivityBuffer() { - free(data_); - } - - // Allocated by malloc - uint8_t* data_{nullptr}; - - // Number of bytes used - size_t validSize_; -}; - -struct roctracerRow { - roctracerRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid - , uint32_t tid, uint64_t begin, uint64_t end) - : id(id), domain(domain), cid(cid), pid(pid), tid(tid), begin(begin), end(end) {} - uint64_t id; // correlation_id - uint32_t domain; - uint32_t cid; - uint32_t pid; - uint32_t tid; - uint64_t begin; - uint64_t end; -}; - -struct kernelRow : public roctracerRow { - kernelRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid - , uint32_t tid, uint64_t begin, uint64_t end - , const void *faddr, hipFunction_t function - , unsigned int gx, unsigned int gy, unsigned int gz - , unsigned int wx, unsigned int wy, unsigned int wz - , size_t gss, hipStream_t stream) - : roctracerRow(id, domain, cid, pid, tid, begin, end), functionAddr(faddr) - , function(function), gridX(gx), gridY(gy), gridZ(gz) - , workgroupX(wx), workgroupY(wy), workgroupZ(wz), groupSegmentSize(gss) - , stream(stream) {} - const void* functionAddr; - hipFunction_t function; - unsigned int gridX; - unsigned int gridY; - unsigned int gridZ; - unsigned int workgroupX; - unsigned int workgroupY; - unsigned int workgroupZ; - size_t groupSegmentSize; - hipStream_t stream; -}; - -struct copyRow : public roctracerRow { - copyRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid - , uint32_t tid, uint64_t begin, uint64_t end - , const void* src, const void *dst, size_t size, hipMemcpyKind kind - , hipStream_t stream) - : roctracerRow(id, domain, cid, pid, tid, begin, end) - , src(src), dst(dst), size(size), kind(kind), stream(stream) {} - const void *src; - const void *dst; - size_t size; - hipMemcpyKind kind; - hipStream_t stream; -}; - -struct mallocRow : public roctracerRow { - mallocRow(uint64_t id, uint32_t domain, uint32_t cid, uint32_t pid - , uint32_t tid, uint64_t begin, uint64_t end - , const void* ptr, size_t size) - : roctracerRow(id, domain, cid, pid, tid, begin, end) - , ptr(ptr), size(size) {} - const void *ptr; - size_t size; -}; - - -class RoctracerLogger { - public: - enum CorrelationDomain { - begin, - Default = begin, - Domain0 = begin, - Domain1, - end, - size = end - }; - - RoctracerLogger(); - RoctracerLogger(const RoctracerLogger&) = delete; - RoctracerLogger& operator=(const RoctracerLogger&) = delete; - - virtual ~RoctracerLogger(); - - static RoctracerLogger& singleton(); - - static void pushCorrelationID(uint64_t id, CorrelationDomain type); - static void popCorrelationID(CorrelationDomain type); - - void startLogging(); - void stopLogging(); - void clearLogs(); - - private: - bool registered_{false}; - void endTracing(); - - roctracer_pool_t *hccPool_{NULL}; - static void api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg); - static void activity_callback(const char* begin, const char* end, void* arg); - - ApiIdList loggedIds_; - - // Api callback data - std::deque rows_; - std::deque kernelRows_; - std::deque copyRows_; - std::deque mallocRows_; - std::map externalCorrelations_[CorrelationDomain::size]; // tracer -> ext - - std::unique_ptr> gpuTraceBuffers_; - bool externalCorrelationEnabled_{true}; - - friend class onnxruntime::profiling::RocmProfiler; -}; -*/ From 22fefa5ef9e687a263816d855a3dcbd6aac12855 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 21:44:31 +0000 Subject: [PATCH 16/41] Delete old code --- .../core/providers/rocm/roctracer_manager.cc | 344 ------------------ 1 file changed, 344 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 23c0ba54e956b..4cbd4cb883cd4 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -480,347 +480,3 @@ void RoctracerManager::ProcessActivityBuffers(const std::vector>(); -} - -RoctracerLogger::~RoctracerLogger() { - stopLogging(); - endTracing(); -} - -namespace { - thread_local std::deque t_externalIds[RoctracerLogger::CorrelationDomain::size]; -} - -void RoctracerLogger::pushCorrelationID(uint64_t id, CorrelationDomain type) { - if (!singleton().externalCorrelationEnabled_) { - return; - } - t_externalIds[type].push_back(id); -} - -void RoctracerLogger::popCorrelationID(CorrelationDomain type) { - if (!singleton().externalCorrelationEnabled_) { - return; - } - t_externalIds[type].pop_back(); -} - -void RoctracerLogger::clearLogs() { - rows_.clear(); - kernelRows_.clear(); - copyRows_.clear(); - mallocRows_.clear(); - gpuTraceBuffers_->clear(); - for (int i = 0; i < CorrelationDomain::size; ++i) { - externalCorrelations_[i].clear(); - } -} - -void RoctracerLogger::api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg) { - RoctracerLogger *dis = &singleton(); - - if (domain == ACTIVITY_DOMAIN_HIP_API && dis->loggedIds_.contains(cid)) { - const hip_api_data_t* data = (const hip_api_data_t*)(callback_data); - - // Pack callbacks into row structures - - thread_local timespec timestamp; - - if (data->phase == ACTIVITY_API_PHASE_ENTER) { - clock_gettime(CLOCK_MONOTONIC, ×tamp); // record proper clock - } else { // (data->phase == ACTIVITY_API_PHASE_EXIT) - timespec endTime; - timespec startTime { timestamp }; - clock_gettime(CLOCK_MONOTONIC, &endTime); // record proper clock - - switch (cid) { - case HIP_API_ID_hipLaunchKernel: - case HIP_API_ID_hipExtLaunchKernel: - case HIP_API_ID_hipLaunchCooperativeKernel: // Should work here - { - auto &args = data->args.hipLaunchKernel; - dis->kernelRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - args.function_address, - nullptr, - args.numBlocks.x, - args.numBlocks.y, - args.numBlocks.z, - args.dimBlocks.x, - args.dimBlocks.y, - args.dimBlocks.z, - args.sharedMemBytes, - args.stream - ); - } - break; - case HIP_API_ID_hipHccModuleLaunchKernel: - case HIP_API_ID_hipModuleLaunchKernel: - case HIP_API_ID_hipExtModuleLaunchKernel: - { - auto &args = data->args.hipModuleLaunchKernel; - dis->kernelRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - nullptr, - args.f, - args.gridDimX, - args.gridDimY, - args.gridDimZ, - args.blockDimX, - args.blockDimY, - args.blockDimZ, - args.sharedMemBytes, - args.stream - ); - } - break; - case HIP_API_ID_hipLaunchCooperativeKernelMultiDevice: - case HIP_API_ID_hipExtLaunchMultiKernelMultiDevice: -#if 0 - { - auto &args = data->args.hipLaunchCooperativeKernelMultiDevice.launchParamsList__val; - dis->kernelRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - args.function_address, - nullptr, - args.numBlocks.x, - args.numBlocks.y, - args.numBlocks.z, - args.dimBlocks.x, - args.dimBlocks.y, - args.dimBlocks.z, - args.sharedMemBytes, - args.stream - ); - } -#endif - break; - case HIP_API_ID_hipMalloc: - dis->mallocRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - data->args.hipMalloc.ptr__val, - data->args.hipMalloc.size - ); - break; - case HIP_API_ID_hipFree: - dis->mallocRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - data->args.hipFree.ptr, - 0 - ); - break; - case HIP_API_ID_hipMemcpy: - { - auto &args = data->args.hipMemcpy; - dis->copyRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - args.src, - args.dst, - args.sizeBytes, - args.kind, - static_cast(0) // use placeholder? - ); - } - break; - case HIP_API_ID_hipMemcpyAsync: - case HIP_API_ID_hipMemcpyWithStream: - { - auto &args = data->args.hipMemcpyAsync; - dis->copyRows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime), - args.src, - args.dst, - args.sizeBytes, - args.kind, - args.stream - ); - } - break; - default: - dis->rows_.emplace_back(data->correlation_id, - domain, - cid, - processId(), - systemThreadId(), - timespec_to_ns(startTime), - timespec_to_ns(endTime) - ); - break; - } // switch - // External correlation - for (int it = CorrelationDomain::begin; it < CorrelationDomain::end; ++it) { - if (t_externalIds[it].size() > 0) { - dis->externalCorrelations_[it][data->correlation_id] = t_externalIds[it].back(); - } - } - } // phase exit - } -} - -void RoctracerLogger::activity_callback(const char* begin, const char* end, void* arg) -{ - size_t size = end - begin; - uint8_t *buffer = (uint8_t*) malloc(size); - auto &gpuTraceBuffers = singleton().gpuTraceBuffers_; - memcpy(buffer, begin, size); - gpuTraceBuffers->emplace_back(buffer, size); -} - -void RoctracerLogger::startLogging() { - if (!registered_) { - roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, nullptr); // Magic encantation - - // Set some api calls to ignore - loggedIds_.setInvertMode(true); // Omit the specified api - loggedIds_.add("hipGetDevice"); - loggedIds_.add("hipSetDevice"); - loggedIds_.add("hipGetLastError"); - loggedIds_.add("__hipPushCallConfiguration"); - loggedIds_.add("__hipPopCallConfiguration"); - loggedIds_.add("hipCtxSetCurrent"); - loggedIds_.add("hipEventRecord"); - loggedIds_.add("hipEventQuery"); - loggedIds_.add("hipGetDeviceProperties"); - loggedIds_.add("hipPeekAtLastError"); - loggedIds_.add("hipModuleGetFunction"); - loggedIds_.add("hipEventCreateWithFlags"); - - // Enable API callbacks - if (loggedIds_.invertMode() == true) { - // exclusion list - enable entire domain and turn off things in list - roctracer_enable_domain_callback(ACTIVITY_DOMAIN_HIP_API, api_callback, nullptr); - const std::unordered_map &filter = loggedIds_.filterList(); - for (auto it = filter.begin(); it != filter.end(); ++it) { - roctracer_disable_op_callback(ACTIVITY_DOMAIN_HIP_API, it->first); - } - } - else { - // inclusion list - only enable things in the list - const std::unordered_map &filter = loggedIds_.filterList(); - roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); - for (auto it = filter.begin(); it != filter.end(); ++it) { - roctracer_enable_op_callback(ACTIVITY_DOMAIN_HIP_API, it->first, api_callback, nullptr); - } - } - //roctracer_enable_domain_callback(ACTIVITY_DOMAIN_ROCTX, api_callback, nullptr); - - // Allocate default tracing pool - roctracer_properties_t properties; - memset(&properties, 0, sizeof(roctracer_properties_t)); - properties.buffer_size = 0x1000; - roctracer_open_pool(&properties); - - // Enable async op collection - roctracer_properties_t hcc_cb_properties; - memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); - hcc_cb_properties.buffer_size = 0x4000; - hcc_cb_properties.buffer_callback_fun = activity_callback; - roctracer_open_pool_expl(&hcc_cb_properties, &hccPool_); - roctracer_enable_domain_activity_expl(ACTIVITY_DOMAIN_HCC_OPS, hccPool_); - - registered_ = true; - } - - externalCorrelationEnabled_ = true; - roctracer_start(); -} - -void RoctracerLogger::stopLogging() { - roctracer_stop(); - roctracer_flush_activity_expl(hccPool_); -} - -void RoctracerLogger::endTracing() { - if (registered_ == true) { - roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); - //roctracer_disable_domain_callback(ACTIVITY_DOMAIN_ROCTX); - - roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS); - roctracer_close_pool_expl(hccPool_); - } -} - - -ApiIdList::ApiIdList() -: invert_(true) -{ -} - -void ApiIdList::add(const std::string &apiName) -{ - uint32_t cid = 0; - if (roctracer_op_code(ACTIVITY_DOMAIN_HIP_API, apiName.c_str(), &cid, nullptr) == ROCTRACER_STATUS_SUCCESS) { - filter_[cid] = 1; - } -} -void ApiIdList::remove(const std::string &apiName) -{ - uint32_t cid = 0; - if (roctracer_op_code(ACTIVITY_DOMAIN_HIP_API, apiName.c_str(), &cid, nullptr) == ROCTRACER_STATUS_SUCCESS) { - filter_.erase(cid); - } -} - -bool ApiIdList::loadUserPrefs() -{ - // placeholder - return false; -} -bool ApiIdList::contains(uint32_t apiId) -{ - return (filter_.find(apiId) != filter_.end()) ? !invert_ : invert_; // XOR -} -*/ From 4190935395b6e1f2d5324c06241a66fdece9b1d6 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 21:45:02 +0000 Subject: [PATCH 17/41] Delete blank lines --- onnxruntime/core/providers/rocm/rocm_profiler.h | 3 --- 1 file changed, 3 deletions(-) diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.h b/onnxruntime/core/providers/rocm/rocm_profiler.h index f6be59bfba0ee..96caba2085079 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.h +++ b/onnxruntime/core/providers/rocm/rocm_profiler.h @@ -47,7 +47,4 @@ class RocmProfiler final : public EpProfiler { } } - - - #endif From 88493356a91fff2ba59e54e96467ac0ae5219db1 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 21:57:51 +0000 Subject: [PATCH 18/41] Fixes --- onnxruntime/core/providers/rocm/roctracer_manager.h | 3 +++ 1 file changed, 3 insertions(+) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index d6ee0198b0a53..e33bf3cc4d04d 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -128,3 +128,6 @@ class RoctracerManager // The api calls to track static const std::vector hip_api_calls_to_trace; }; + +} /* end namespace profiling */ +} /* end namespace onnxruntime*/ From 1eaaa8ef77616761987408ff097ae11fc64479bd Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 22:19:29 +0000 Subject: [PATCH 19/41] Fix re-entrancy deadlock --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 1 - onnxruntime/core/providers/rocm/roctracer_manager.h | 3 +-- 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 4cbd4cb883cd4..2d79267e0ce97 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -124,7 +124,6 @@ void RoctracerManager::Clear() } void RoctracerManager::StopLogging() { - std::lock_guard lock(roctracer_manager_mutex_); if (!logging_enabled_) { return; } diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index e33bf3cc4d04d..2762cb98e7f65 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -74,7 +74,6 @@ class RoctracerManager void DeregisterClient(uint64_t client_handle); void StartLogging(); - void StopLogging(); void Consume(uint64_t client_handle, const TimePoint& start_time, std::map& events); bool PushCorrelation(uint64_t client_handle, uint64_t external_correlation_id); @@ -99,7 +98,7 @@ class RoctracerManager const ApiCallRecord& call_record, EventRecord& event); void MapEventToClient(uint64_t external_correlation_id, EventRecord&& event); void MapEventsToClient(uint64_t external_correlation_id, Events&& events); - + void StopLogging(); void Clear(); // Some useful constants for processing activity buffers From dd21a0c1395e081f3704b34023e915b346837023 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 23:22:54 +0000 Subject: [PATCH 20/41] Fixes --- .../core/providers/rocm/roctracer_manager.cc | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 2d79267e0ce97..a772e0972c1f7 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -15,11 +15,11 @@ const std::vector RoctracerManager::hip_api_calls_to_trace = { "hipMemcpy2D", "hipMemcpyAsync", "hipMemcpy2DAsync", - "hipMemcpyWithStream" + "hipMemcpyWithStream", "hipLaunchKernel", "hipMemset", "hipMemsetAsync", - "hipExtModuleLaunchKernel" + "hipExtModuleLaunchKernel", }; // Implementation of RoctracerActivityBuffer @@ -128,12 +128,12 @@ void RoctracerManager::StopLogging() { return; } - roctracer_stop(); - roctracer_flush_activity(); roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_API); roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS); - roctracer_close_pool(); roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API); + roctracer_stop(); + roctracer_flush_activity(); + roctracer_close_pool(); logging_enabled_ = false; Clear(); @@ -416,7 +416,8 @@ void RoctracerManager::ProcessActivityBuffers(const std::vectorexternal_id, std::move(pending_it->second)); events_pending_client_mapping.erase(current_record->correlation_id); - + // no additional events to be mapped for this record + continue; } else if (current_record->domain == ACTIVITY_DOMAIN_HIP_OPS) { if (current_record->op == 1 && current_record->kind == HipOpMarker) { // this is just a marker, ignore it. From 0646193f8462daa7178b463588d226a69401c858 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 1 Nov 2022 23:30:18 +0000 Subject: [PATCH 21/41] Minor fix --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index a772e0972c1f7..5420288598560 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -458,7 +458,7 @@ void RoctracerManager::ProcessActivityBuffers(const std::vector Date: Wed, 2 Nov 2022 01:48:21 +0000 Subject: [PATCH 22/41] Cleanups --- .../core/providers/rocm/roctracer_manager.cc | 250 +++++++----------- .../core/providers/rocm/roctracer_manager.h | 15 +- 2 files changed, 100 insertions(+), 165 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 5420288598560..64614eb195ee4 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -216,61 +216,6 @@ static inline std::string PointerToHexString(const void* ptr) { return sstr.str(); } -void RoctracerManager::CreateEventForKernelRecord(const roctracer_record_t* record, - uint64_t start_time_ns, - const ApiCallRecord& call_record, - EventRecord& event) { - auto const& launch_args = call_record.api_data_; - auto name = demangle(hipKernelNameRefByPtr(launch_args.args.hipLaunchKernel.function_address, - launch_args.args.hipLaunchKernel.stream)); - std::unordered_map args { - {"stream", PointerToHexString((void*)(launch_args.args.hipLaunchKernel.stream))}, - {"grid_x", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.x)}, - {"grid_y", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.y)}, - {"grid_z", std::to_string(launch_args.args.hipLaunchKernel.numBlocks.z)}, - {"block_x", std::to_string(launch_args.args.hipLaunchKernel.dimBlocks.x)}, - {"block_y", std::to_string(launch_args.args.hipLaunchKernel.dimBlocks.y)}, - {"block_z", std::to_string(launch_args.args.hipLaunchKernel.dimBlocks.z)} - }; - - new (&event) EventRecord { - /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ -1, - /* tid = */ -1, - /* name = */ std::move(name), - /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, - /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, - /* args = */ std::move(args) - }; -} - -void RoctracerManager::CreateEventForMemsetRecord(const roctracer_record_t* record, - uint64_t start_time_ns, - const ApiCallRecord& call_record, - EventRecord& event) { - auto const& launch_args = call_record.api_data_; - auto dst_string = PointerToHexString(launch_args.args.hipMemset.dst); - std::string name {roctracer_op_string(call_record.domain_, call_record.cid_, 0)}; - - std::unordered_map args { - {"stream", call_record.cid_ == HIP_API_ID_hipMemset - ? "0" - : PointerToHexString((void*)launch_args.args.hipMemsetAsync.stream)}, - {"dst", dst_string}, - {"size", std::to_string(launch_args.args.hipMemset.sizeBytes)}, - {"value", std::to_string(launch_args.args.hipMemset.value)} - }; - new (&event) EventRecord { - /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ -1, - /* tid = */ -1, - /* name = */ std::move(name), - /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, - /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, - /* args = */ std::move(args) - }; -} - static inline std::string MemcpyKindToString(hipMemcpyKind kind) { switch(kind) { case hipMemcpyHostToHost: @@ -286,81 +231,103 @@ static inline std::string MemcpyKindToString(hipMemcpyKind kind) { } } -void RoctracerManager::CreateEventForMemcpyRecord(const roctracer_record_t* record, uint64_t start_time_ns, - const ApiCallRecord& call_record, EventRecord& event) { - auto const& launch_args = call_record.api_data_; - auto src_string = PointerToHexString(launch_args.args.hipMemcpy.src); - auto dst_string = PointerToHexString(launch_args.args.hipMemcpy.dst); - std::string name {roctracer_op_string(call_record.domain_, call_record.cid_, 0)}; - - std::string memcpy_kind_string = MemcpyKindToString(launch_args.args.hipMemcpy.kind); - - std::unordered_map args { - {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy - ? "0" - : PointerToHexString((void*)launch_args.args.hipMemcpyAsync.stream)}, - {"src", src_string}, - {"dst", dst_string}, - {"kind", memcpy_kind_string} - }; +bool RoctracerManager::CreateEventForActivityRecord(const roctracer_record_t* record, + uint64_t start_time_ns, + const ApiCallRecord& call_record, + EventRecord& event) { + std::string name; + std::unordered_map args; + + switch(call_record.cid_) { + case HIP_API_ID_hipLaunchKernel: { + name = demangle(hipKernelNameRefByPtr(launch_args.args.hipLaunchKernel.function_address, + launch_args.args.hipLaunchKernel.stream)); + auto const& launch_args = call_record.api_data_.args.hipLaunchKernel; + + args = { + {"stream", PointerToHexString((void*)(launch_args.stream))}, + {"grid_x", std::to_string(launch_args.numBlocks.x)}, + {"grid_y", std::to_string(launch_args.numBlocks.y)}, + {"grid_z", std::to_string(launch_args.numBlocks.z)}, + {"block_x", std::to_string(launch_args.dimBlocks.x)}, + {"block_y", std::to_string(launch_args.dimBlocks.y)}, + {"block_z", std::to_string(launch_args.dimBlocks.z)} + }; + break; + } - new (&event) EventRecord { - /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ -1, - /* tid = */ -1, - /* name = */ std::move(name), - /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, - /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, - /* args = */ std::move(args) - }; -} + case HIP_API_ID_hipMemset: + case HIP_API_ID_hipMemsetAsync: { + auto const& launch_args = call_record.api_data_.args; + name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); + + args = { + {"stream", call_record.cid_ == HIP_API_ID_hipMemset + ? "0" + : PointerToHexString((void*)launch_args.hipMemsetAsync.stream)}, + {"dst", PointerToHexString(launch_args.hipMemset.dst)}, + {"size", std::to_string(launch_args.hipMemset.sizeBytes)}, + {"value", std::to_string(launch_args.hipMemset.value)} + }; + break; + } -void RoctracerManager::CreateEventForMemcpy2DRecord(const roctracer_record_t* record, uint64_t start_time_ns, - const ApiCallRecord& call_record, EventRecord& event) { - auto const& launch_args = call_record.api_data_; - auto src_string = PointerToHexString(launch_args.args.hipMemcpy2D.src); - auto dst_string = PointerToHexString(launch_args.args.hipMemcpy2D.dst); - std::string name {roctracer_op_string(call_record.domain_, call_record.cid_, 0)}; - - std::string memcpy_kind_string = MemcpyKindToString(launch_args.args.hipMemcpy2D.kind); - - std::unordered_map args { - {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy2D - ? "0" - : PointerToHexString((void*)launch_args.args.hipMemcpy2DAsync.stream)}, - {"src", src_string}, - {"dst", dst_string}, - {"spitch", std::to_string(launch_args.args.hipMemcpy2D.spitch)}, - {"dpitch", std::to_string(launch_args.args.hipMemcpy2D.dpitch)}, - {"width", std::to_string(launch_args.args.hipMemcpy2D.width)}, - {"height", std::to_string(launch_args.args.hipMemcpy2D.height)}, - {"kind", memcpy_kind_string} - }; + case HIP_API_ID_hipMemcpy: + case HIP_API_ID_hipMemcpyAsync: + case HIP_API_ID_hipMemcpyWithStream: { + auto const& launch_args = call_record.api_data_.args; + name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); + + args = { + {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy + ? "0" + : PointerToHexString((void*)launch_args.hipMemcpyAsync.stream)}, + {"src", PointerToHexString(launch_args.hipMemcpy.src)}, + {"dst", PointerToHexString(launch_args.hipMemcpy.dst)}, + {"kind", MemcpyKindToString(launch_args.hipMemcpy.kind)} + }; + break; + } - new (&event) EventRecord { - /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ -1, - /* tid = */ -1, - /* name = */ std::move(name), - /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, - /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, - /* args = */ std::move(args) - }; -} + case HIP_API_ID_hipMemcpy2D: + case HIP_API_ID_hipMemcpy2DAsync: { + auto const& launch_args = call_record.api_data_.args; + name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); + + args = { + {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy2D + ? "0" + : PointerToHexString((void*)launch_args.hipMemcpy2DAsync.stream)}, + {"src", PointerToHexString(launch_args.hipMemcpy2D.src)}, + {"dst", PointerToHexString(launch_args.hipMemcpy2D.dst)}, + {"spitch", std::to_string(launch_args.hipMemcpy2D.spitch)}, + {"dpitch", std::to_string(launch_args.hipMemcpy2D.dpitch)}, + {"width", std::to_string(launch_args.hipMemcpy2D.width)}, + {"height", std::to_string(launch_args.hipMemcpy2D.height)}, + {"kind", MemcpyKindToString(launch_args.hipMemcpy2D.kind)} + }; + break; + } -void RoctracerManager::CreateEventForExtModuleLaunchKernel(const roctracer_record_t* record, uint64_t start_time_ns, - const ApiCallRecord& call_record, EventRecord& event) { - auto const& launch_args = call_record.api_data_; - auto name = hipKernelNameRef(launch_args.args.hipExtModuleLaunchKernel.f); - std::unordered_map args { - {"stream", PointerToHexString((void*)launch_args.args.hipExtModuleLaunchKernel.hStream)}, - {"grid_x", std::to_string(launch_args.args.hipExtModuleLaunchKernel.globalWorkSizeX)}, - {"grid_y", std::to_string(launch_args.args.hipExtModuleLaunchKernel.globalWorkSizeY)}, - {"grid_z", std::to_string(launch_args.args.hipExtModuleLaunchKernel.globalWorkSizeZ)}, - {"block_x", std::to_string(launch_args.args.hipExtModuleLaunchKernel.localWorkSizeX)}, - {"block_y", std::to_string(launch_args.args.hipExtModuleLaunchKernel.localWorkSizeY)}, - {"block_z", std::to_string(launch_args.args.hipExtModuleLaunchKernel.localWorkSizeZ)}, - }; + case HIP_API_ID_hipExtModuleLaunchKernel: { + auto const& launch_args = call_record.api_data_.args.hipExtModuleLaunchKernel; + name = demangle(hipKernelNameRef(launch_args.f)); + + args = { + {"stream", PointerToHexString((void*)launch_args.hStream)}, + {"grid_x", std::to_string(launch_args.globalWorkSizeX)}, + {"grid_y", std::to_string(launch_args.globalWorkSizeY)}, + {"grid_z", std::to_string(launch_args.globalWorkSizeZ)}, + {"block_x", std::to_string(launch_args.localWorkSizeX)}, + {"block_y", std::to_string(launch_args.localWorkSizeY)}, + {"block_z", std::to_string(launch_args.localWorkSizeZ)}, + }; + break; + } + + default: + return false; + } new (&event) EventRecord { /* cat = */ EventCategory::KERNEL_EVENT, @@ -371,6 +338,7 @@ void RoctracerManager::CreateEventForExtModuleLaunchKernel(const roctracer_recor /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, /* args = */ std::move(args) }; + return true; } void RoctracerManager::MapEventsToClient(uint64_t external_correlation_id, Events&& events) { @@ -431,33 +399,9 @@ void RoctracerManager::ProcessActivityBuffers(const std::vectorsecond; - switch (call_record.cid_) - { - case HIP_API_ID_hipLaunchKernel: - CreateEventForKernelRecord(current_record, start_time_ns, call_record, event); - break; - - case HIP_API_ID_hipMemset: - case HIP_API_ID_hipMemsetAsync: - CreateEventForMemsetRecord(current_record, start_time_ns, call_record, event); - break; - - case HIP_API_ID_hipMemcpy: - case HIP_API_ID_hipMemcpyAsync: - case HIP_API_ID_hipMemcpyWithStream: - CreateEventForMemcpyRecord(current_record, start_time_ns, call_record, event); - break; - - case HIP_API_ID_hipMemcpy2D: - case HIP_API_ID_hipMemcpy2DAsync: - CreateEventForMemcpy2DRecord(current_record, start_time_ns, call_record, event); - break; - - case HIP_API_ID_hipExtModuleLaunchKernel: - CreateEventForExtModuleLaunchKernel(current_record, start_time_ns, call_record, event); - break; - - default: + if (!CreateEventForActivityRecord(current_record, start_time_ns, call_record, event)) { + // No event created, skip to the next record to avoid associating an empty + // event with a client continue; } } else { diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 2762cb98e7f65..852055572f0cd 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -83,19 +83,10 @@ class RoctracerManager private: static void ActivityCallback(const char* begin, const char* end, void* arg); static void ApiCallback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg); - void ProcessActivityBuffers(const std::vector& buffers, const TimePoint& start_time); - - // Per-API (roughly) helpers for event construction. - void CreateEventForKernelRecord(const roctracer_record_t* record, uint64_t start_time_ns, - const ApiCallRecord& call_record, EventRecord& event); - void CreateEventForMemsetRecord(const roctracer_record_t* record, uint64_t start_time_ns, - const ApiCallRecord& call_record, EventRecord& event); - void CreateEventForMemcpyRecord(const roctracer_record_t* record, uint64_t start_time_ns, - const ApiCallRecord& call_record, EventRecord& event); - void CreateEventForMemcpy2DRecord(const roctracer_record_t* record, uint64_t start_time_ns, + void ProcessActivityBuffers(const std::vector& buffers, + const TimePoint& start_time); + bool CreateEventForActivityRecord(const roctracer_record_t* record, uint64_t start_time_ns, const ApiCallRecord& call_record, EventRecord& event); - void CreateEventForExtModuleLaunchKernel(const roctracer_record_t* record, uint64_t start_time_ns, - const ApiCallRecord& call_record, EventRecord& event); void MapEventToClient(uint64_t external_correlation_id, EventRecord&& event); void MapEventsToClient(uint64_t external_correlation_id, Events&& events); void StopLogging(); From ce97928a4632135f27bed65d9cb5110c24fd1749 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Wed, 2 Nov 2022 01:50:09 +0000 Subject: [PATCH 23/41] Fix --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 64614eb195ee4..5e2122c3cc4b6 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -240,9 +240,9 @@ bool RoctracerManager::CreateEventForActivityRecord(const roctracer_record_t* re switch(call_record.cid_) { case HIP_API_ID_hipLaunchKernel: { - name = demangle(hipKernelNameRefByPtr(launch_args.args.hipLaunchKernel.function_address, - launch_args.args.hipLaunchKernel.stream)); auto const& launch_args = call_record.api_data_.args.hipLaunchKernel; + name = demangle(hipKernelNameRefByPtr(launch_args.function_address, + launch_args.stream)); args = { {"stream", PointerToHexString((void*)(launch_args.stream))}, From b89bcd65669d4d17e2a4072ff7e839e6ea738033 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Wed, 2 Nov 2022 19:01:56 +0000 Subject: [PATCH 24/41] Minor fix --- onnxruntime/core/providers/rocm/rocm_profiler.cc | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.cc b/onnxruntime/core/providers/rocm/rocm_profiler.cc index b9d7e8efc7406..ba4f37de8db72 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.cc +++ b/onnxruntime/core/providers/rocm/rocm_profiler.cc @@ -44,6 +44,11 @@ void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) { ++event_iter; } + // find the last event with the same timestamp. + while (event_iter != event_end && event_iter->ts == ts && (event_iter + 1)->ts == ts) { + ++event_iter; + } + if (event_iter != event_end && event_iter->ts == ts) { uint64_t increment = 1; for (auto& evt : map_iter.second) { From b95d47032c9caeb597c2f2d718eaa3d9595b19f9 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Wed, 2 Nov 2022 19:59:12 +0000 Subject: [PATCH 25/41] Flush activity records on Consume --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 5e2122c3cc4b6..e73ee8aa53ece 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -114,6 +114,7 @@ void RoctracerManager::StartLogging() { logging_enabled_ = true; } +// Requires: roctracer_manager_mutex_ must be held void RoctracerManager::Clear() { unprocessed_activity_buffers_.clear(); @@ -123,6 +124,7 @@ void RoctracerManager::Clear() per_client_events_by_ext_correlation_.clear(); } +// Requires: roctracer_manager_mutex_ must be held void RoctracerManager::StopLogging() { if (!logging_enabled_) { return; @@ -142,6 +144,13 @@ void RoctracerManager::StopLogging() { void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_time, std::map& events) { events.clear(); + { + // Flush any pending activity records before starting + // to process the accumulated activity records. + std::lock_guard lock_manager(roctracer_manager_mutex_); + roctracer_flush_activity(); + } + std::vector activity_buffers; { std::lock_guard lock(unprocessed_activity_buffers_lock_); From c7064778b81b5da04ce673aa07379195eee25d9c Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Wed, 2 Nov 2022 21:58:00 +0000 Subject: [PATCH 26/41] Implement unique correlation ids between sessions --- .../onnxruntime/core/common/profiler_common.h | 2 +- .../core/providers/cuda/cuda_profiler.cc | 3 +- .../core/providers/cuda/cuda_profiler.h | 4 +- .../core/providers/rocm/rocm_profiler.cc | 5 +- .../core/providers/rocm/rocm_profiler.h | 5 +- .../core/providers/rocm/roctracer_manager.cc | 98 ++++++++++++++----- .../core/providers/rocm/roctracer_manager.h | 12 ++- 7 files changed, 91 insertions(+), 38 deletions(-) diff --git a/include/onnxruntime/core/common/profiler_common.h b/include/onnxruntime/core/common/profiler_common.h index f0e55b5093bce..e1a87ad31f3c0 100644 --- a/include/onnxruntime/core/common/profiler_common.h +++ b/include/onnxruntime/core/common/profiler_common.h @@ -108,7 +108,7 @@ using Events = std::vector; class EpProfiler { public: virtual ~EpProfiler() = default; - virtual bool StartProfiling() = 0; // called when profiling starts + virtual bool StartProfiling(TimePoint profiling_start_time) = 0; // called when profiling starts virtual void EndProfiling(TimePoint start_time, Events& events) = 0; // called when profiling ends, save all captures numbers to "events" virtual void Start(uint64_t){}; // called before op start, accept an id as argument to identify the op virtual void Stop(uint64_t){}; // called after op stop, accept an id as argument to identify the op diff --git a/onnxruntime/core/providers/cuda/cuda_profiler.cc b/onnxruntime/core/providers/cuda/cuda_profiler.cc index 817816f43b163..c086f4b8db3a9 100644 --- a/onnxruntime/core/providers/cuda/cuda_profiler.cc +++ b/onnxruntime/core/providers/cuda/cuda_profiler.cc @@ -93,7 +93,8 @@ void CUPTIAPI CudaProfiler::BufferCompleted(CUcontext, uint32_t, uint8_t* buffer free(buffer); } -bool CudaProfiler::StartProfiling() { +bool CudaProfiler::StartProfiling(TimePoint profiling_start_time) { + (void) profiling_start_time; if (!enabled.test_and_set()) { if (cuptiActivityEnable(CUPTI_ACTIVITY_KIND_RUNTIME) == CUPTI_SUCCESS && cuptiActivityEnable(CUPTI_ACTIVITY_KIND_DRIVER) == CUPTI_SUCCESS && diff --git a/onnxruntime/core/providers/cuda/cuda_profiler.h b/onnxruntime/core/providers/cuda/cuda_profiler.h index bd625a7c6ac3e..c089705570751 100644 --- a/onnxruntime/core/providers/cuda/cuda_profiler.h +++ b/onnxruntime/core/providers/cuda/cuda_profiler.h @@ -31,7 +31,7 @@ class CudaProfiler final : public EpProfiler { return *this; } ~CudaProfiler(); - bool StartProfiling() override; + bool StartProfiling(TimePoint profiling_start_time) override; void EndProfiling(TimePoint start_time, Events& events) override; void Start(uint64_t) override; void Stop(uint64_t) override; @@ -72,7 +72,7 @@ namespace profiling { class CudaProfiler final : public EpProfiler { public: - bool StartProfiling() override { return true; } + bool StartProfiling(TimePoint) override { return true; } void EndProfiling(TimePoint, Events&) override{}; void Start(uint64_t) override{}; void Stop(uint64_t) override{}; diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.cc b/onnxruntime/core/providers/rocm/rocm_profiler.cc index ba4f37de8db72..6559ac59f0c9e 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.cc +++ b/onnxruntime/core/providers/rocm/rocm_profiler.cc @@ -22,9 +22,10 @@ RocmProfiler::~RocmProfiler() { manager.DeregisterClient(client_handle_); } -bool RocmProfiler::StartProfiling() { +bool RocmProfiler::StartProfiling(TimePoint profiling_start_time) { auto& manager = RoctracerManager::GetInstance(); manager.StartLogging(); + profiling_start_time_ = profiling_start_time; return true; } @@ -76,7 +77,7 @@ void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) { void RocmProfiler::Start(uint64_t id) { auto& manager = RoctracerManager::GetInstance(); - manager.PushCorrelation(client_handle_, id); + manager.PushCorrelation(client_handle_, id, profiling_start_time_); } void RocmProfiler::Stop(uint64_t id) { diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.h b/onnxruntime/core/providers/rocm/rocm_profiler.h index 96caba2085079..249709634b183 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.h +++ b/onnxruntime/core/providers/rocm/rocm_profiler.h @@ -18,13 +18,14 @@ class RocmProfiler final : public EpProfiler { RocmProfiler(const RocmProfiler&) = delete; RocmProfiler& operator=(const RocmProfiler&) = delete; ~RocmProfiler(); - bool StartProfiling() override; + bool StartProfiling(TimePoint profiling_start_time) override; void EndProfiling(TimePoint start_time, Events& events) override; void Start(uint64_t) override; void Stop(uint64_t) override; private: uint64_t client_handle_; + TimePoint profiling_start_time_; }; } // namespace profiling @@ -38,7 +39,7 @@ namespace profiling { class RocmProfiler final : public EpProfiler { public: - bool StartProfiling() override { return true; } + bool StartProfiling(TimePoint) override { return true; } void EndProfiling(TimePoint, Events&) override{}; void Start(uint64_t) override{}; void Stop(uint64_t) override{}; diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index e73ee8aa53ece..397450029f751 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -119,8 +119,8 @@ void RoctracerManager::Clear() { unprocessed_activity_buffers_.clear(); api_call_args_.clear(); - external_correlation_id_to_client_.clear(); - roctracer_correlation_to_external_correlation_.clear(); + unique_correlation_id_to_client_offset_.clear(); + roctracer_correlation_to_unique_correlation_.clear(); per_client_events_by_ext_correlation_.clear(); } @@ -171,22 +171,58 @@ void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_ti } } -bool RoctracerManager::PushCorrelation(uint64_t client_handle, uint64_t external_correlation_id) { +bool RoctracerManager::PushCorrelation(uint64_t client_handle, + uint64_t external_correlation_id, + TimePoint profiling_start_time) { std::lock_guard lock(roctracer_manager_mutex_); + auto it = per_client_events_by_ext_correlation_.find(client_handle); if (it == per_client_events_by_ext_correlation_.end()) { + // not a registered client, do nothing return false; } - roctracer_activity_push_external_correlation_id(external_correlation_id); - // assumption: correlation_ids are unique. This should generally work, - // because we use timestamps as correlation ids. - external_correlation_id_to_client_[external_correlation_id] = client_handle; + + // external_correlation_id is simply the timestamp of this event, + // relative to profiling_start_time. i.e., it was computed as: + // external_correlation_id = + // std::chrono::duration_cast(event_start_time - profiling_start_time).count() + // + // Because of the relative nature of the external_correlation_id, the same + // external_correlation_id can be reused across different clients, which then makes it + // impossible to recover the client from the external_correlation_id, which in turn + // makes it impossible to map events (which are tagged with external_correlation_id) to clients. + // + // To address these difficulties, we construct a new correlation_id (let's call it unique_cid) + // as follows: + // unique_cid = + // external_correlation_id + + // std::chrono::duration_cast(profiling_start_time.time_since_epoch()).count() + // now, unique_cid is monotonically increasing with time, so it can be used to reliably map events to clients. + // + // Of course, clients expect lists of events to be returned (on a call to Consume()), that are + // still keyed on the external_correlation_id that they've specified here, so we need to remember the + // offset to be subtracted + + uint64_t offset = + std::chrono::duration_cast(profiling_start_time.time_since_epoch()).count(); + auto unique_cid = external_correlation_id + offset; + roctracer_activity_push_external_correlation_id(unique_cid); + + unique_correlation_id_to_client_offset_[unique_cid] = std::make_pair(client_handle, offset); return true; } void RoctracerManager::PopCorrelation(uint64_t& popped_external_correlation_id) { std::lock_guard lock(roctracer_manager_mutex_); - roctracer_activity_pop_external_correlation_id(&popped_external_correlation_id); + uint64_t unique_cid; + roctracer_activity_pop_external_correlation_id(&unique_cid); + // lookup the offset and subtract it before returning popped_external_correlation_id to the client + auto client_it = unique_correlation_id_to_client_offset_.find(unique_cid); + if (client_it == unique_correlation_id_to_client_offset_.end()) { + popped_external_correlation_id = 0; + return; + } + popped_external_correlation_id = unique_cid - client_it->second.second; } void RoctracerManager::ActivityCallback(const char* begin, const char* end, void* arg) { @@ -350,25 +386,37 @@ bool RoctracerManager::CreateEventForActivityRecord(const roctracer_record_t* re return true; } -void RoctracerManager::MapEventsToClient(uint64_t external_correlation_id, Events&& events) { - auto client_it = external_correlation_id_to_client_.find(external_correlation_id); - if (client_it == external_correlation_id_to_client_.end()) { +Events* RoctracerManager::GetEventListForUniqueCorrelationId(uint64_t unique_correlation_id) { + auto client_it = unique_correlation_id_to_client_offset_.find(unique_correlation_id); + if (client_it == unique_correlation_id_to_client_offset_.end()) { // :-( well, we tried really, really hard to map this event to a client. - return; + return nullptr; } - auto& event_list = per_client_events_by_ext_correlation_[client_it->second][client_it->first]; - event_list.insert(event_list.end(), - std::make_move_iterator(events.begin()), - std::make_move_iterator(events.end())); + + // See the comments on the PushCorrelation method for an explanation of + // of this offset computation and why it's required. + auto const& client_handle_offset = client_it->second; + auto external_correlation = unique_correlation_id - client_handle_offset.second; + + auto& event_list = per_client_events_by_ext_correlation_[client_handle_offset.first][external_correlation]; + return &event_list; } -void RoctracerManager::MapEventToClient(uint64_t external_correlation_id, EventRecord&& event) { - auto client_it = external_correlation_id_to_client_.find(external_correlation_id); - if (client_it == external_correlation_id_to_client_.end()) { - // :-( well, we tried really, really hard to map this event to a client. - return; +void RoctracerManager::MapEventsToClient(uint64_t unique_correlation_id, Events&& events) { + auto p_event_list = GetEventListForUniqueCorrelationId(unique_correlation_id); + if (p_event_list != nullptr) { + p_event_list->insert(p_event_list->end(), + std::make_move_iterator(events.begin()), + std::make_move_iterator(events.end())); + } +} + +void RoctracerManager::MapEventToClient(uint64_t unique_correlation_id, EventRecord&& event) { + auto p_event_list = GetEventListForUniqueCorrelationId(unique_correlation_id); + + if (p_event_list != nullptr) { + p_event_list->emplace_back(std::move(event)); } - per_client_events_by_ext_correlation_[client_it->second][client_it->first].emplace_back(std::move(event)); } void RoctracerManager::ProcessActivityBuffers(const std::vector& buffers, @@ -382,7 +430,7 @@ void RoctracerManager::ProcessActivityBuffers(const std::vectordomain == ACTIVITY_DOMAIN_EXT_API) { - roctracer_correlation_to_external_correlation_[current_record->correlation_id] = current_record->external_id; + roctracer_correlation_to_unique_correlation_[current_record->correlation_id] = current_record->external_id; // check for any events pending client mapping on this correlation auto pending_it = events_pending_client_mapping.find(current_record->correlation_id); @@ -420,8 +468,8 @@ void RoctracerManager::ProcessActivityBuffers(const std::vectorcorrelation_id); - if (ext_corr_it == roctracer_correlation_to_external_correlation_.end()) { + auto ext_corr_it = roctracer_correlation_to_unique_correlation_.find(current_record->correlation_id); + if (ext_corr_it == roctracer_correlation_to_unique_correlation_.end()) { // defer the processing of this event events_pending_client_mapping[current_record->correlation_id].emplace_back(std::move(event)); continue; diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 852055572f0cd..24cf2ad3fdd20 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -76,7 +76,7 @@ class RoctracerManager void StartLogging(); void Consume(uint64_t client_handle, const TimePoint& start_time, std::map& events); - bool PushCorrelation(uint64_t client_handle, uint64_t external_correlation_id); + bool PushCorrelation(uint64_t client_handle, uint64_t external_correlation_id, TimePoint profiling_start_time); void PopCorrelation(uint64_t& popped_correlation_id); bool PopCorrelation(); @@ -87,6 +87,7 @@ class RoctracerManager const TimePoint& start_time); bool CreateEventForActivityRecord(const roctracer_record_t* record, uint64_t start_time_ns, const ApiCallRecord& call_record, EventRecord& event); + Events* GetEventListForUniqueCorrelationId(uint64_t unique_correlation_id); void MapEventToClient(uint64_t external_correlation_id, EventRecord&& event); void MapEventsToClient(uint64_t external_correlation_id, Events&& events); void StopLogging(); @@ -101,11 +102,12 @@ class RoctracerManager std::mutex api_call_args_lock_; std::unordered_map api_call_args_; - // Keyed on external_correlation_id -> client_id/client_handle - std::unordered_map external_correlation_id_to_client_; + // Keyed on unique_correlation_id -> (client_id/client_handle, offset) + // unique_correlation_id - offset == external_correlation_id + std::unordered_map> unique_correlation_id_to_client_offset_; - // Keyed on roctracer_correlation_id -> external_correlation_id - std::unordered_map roctracer_correlation_to_external_correlation_; + // Keyed on roctracer_correlation_id -> unique_correlation_id + std::unordered_map roctracer_correlation_to_unique_correlation_; // client_id/client_handle -> external_correlation_id -> events std::mutex event_list_mutex_; From 589cecce78e66224ea9f2ad7e1710640f3ccc604 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Wed, 2 Nov 2022 22:01:42 +0000 Subject: [PATCH 27/41] Bug --- onnxruntime/core/common/profiler.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/common/profiler.cc b/onnxruntime/core/common/profiler.cc index b64b0b9ebade8..e32d33dc22ae9 100644 --- a/onnxruntime/core/common/profiler.cc +++ b/onnxruntime/core/common/profiler.cc @@ -49,7 +49,7 @@ void Profiler::StartProfiling(const logging::Logger* custom_logger) { custom_logger_ = custom_logger; profiling_start_time_ = std::chrono::high_resolution_clock::now(); for (const auto& ep_profiler : ep_profilers_) { - ep_profiler->StartProfiling(); + ep_profiler->StartProfiling(profiling_start_time_); } } From c7672553379efdacff5180f51f7c9ebaa67f61e8 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Wed, 2 Nov 2022 22:03:08 +0000 Subject: [PATCH 28/41] Bug fix --- onnxruntime/core/common/profiler.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/onnxruntime/core/common/profiler.h b/onnxruntime/core/common/profiler.h index 9c6f3ed0e1006..80417e5aaef19 100644 --- a/onnxruntime/core/common/profiler.h +++ b/onnxruntime/core/common/profiler.h @@ -108,12 +108,12 @@ class Profiler { static void SetGlobalMaxNumEvents(size_t new_max_num_events) { global_max_num_events_.store(new_max_num_events); } - + void AddEpProfilers(std::unique_ptr ep_profiler) { if (ep_profiler) { ep_profilers_.push_back(std::move(ep_profiler)); if (enabled_) { - ep_profilers_.back()->StartProfiling(); + ep_profilers_.back()->StartProfiling(profiling_start_time_); } } } From 29a3aa712a8877f3b5022be2029dd0d796f492d3 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Wed, 2 Nov 2022 22:04:47 +0000 Subject: [PATCH 29/41] Bug fix --- onnxruntime/core/common/profiler.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/common/profiler.cc b/onnxruntime/core/common/profiler.cc index e32d33dc22ae9..89dc5a667a8fb 100644 --- a/onnxruntime/core/common/profiler.cc +++ b/onnxruntime/core/common/profiler.cc @@ -62,7 +62,7 @@ void Profiler::StartProfiling(const std::basic_string& file_name) { profile_stream_file_ = ToUTF8String(file_name); profiling_start_time_ = std::chrono::high_resolution_clock::now(); for (const auto& ep_profiler : ep_profilers_) { - ep_profiler->StartProfiling(); + ep_profiler->StartProfiling(profiling_start_time_); } } From a0e9a2008caefe43ada789dcae68149155ccb8cd Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Wed, 2 Nov 2022 22:51:42 +0000 Subject: [PATCH 30/41] Remove unnecessary mute --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 2 -- onnxruntime/core/providers/rocm/roctracer_manager.h | 1 - 2 files changed, 3 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 397450029f751..4d8917b1f7760 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -162,7 +162,6 @@ void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_ti // Ensure that at most one thread is working through the activity buffers at any time. std::lock_guard lock_two(activity_buffer_processor_mutex_); ProcessActivityBuffers(activity_buffers, start_time); - std::lock_guard lock(event_list_mutex_); auto it = per_client_events_by_ext_correlation_.find(client_handle); if (it == per_client_events_by_ext_correlation_.end()) { return; @@ -413,7 +412,6 @@ void RoctracerManager::MapEventsToClient(uint64_t unique_correlation_id, Events& void RoctracerManager::MapEventToClient(uint64_t unique_correlation_id, EventRecord&& event) { auto p_event_list = GetEventListForUniqueCorrelationId(unique_correlation_id); - if (p_event_list != nullptr) { p_event_list->emplace_back(std::move(event)); } diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 24cf2ad3fdd20..22fda2b00e4f6 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -110,7 +110,6 @@ class RoctracerManager std::unordered_map roctracer_correlation_to_unique_correlation_; // client_id/client_handle -> external_correlation_id -> events - std::mutex event_list_mutex_; std::unordered_map> per_client_events_by_ext_correlation_; uint64_t next_client_id_ = 1; uint64_t num_active_clients_ = 0; From 25be99fbe495f263006736b64fa1f5dec1b46c31 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Thu, 3 Nov 2022 20:59:44 +0000 Subject: [PATCH 31/41] Ran clang-format on all changed source files --- .../onnxruntime/core/common/profiler_common.h | 42 ++-- onnxruntime/core/common/profiler.h | 3 +- .../core/providers/cuda/cuda_profiler.cc | 6 +- .../core/providers/cuda/cuda_profiler.h | 4 +- .../providers/rocm/rocm_execution_provider.cc | 11 +- .../core/providers/rocm/rocm_profiler.h | 6 +- .../core/providers/rocm/roctracer_manager.cc | 214 +++++++++--------- .../core/providers/rocm/roctracer_manager.h | 30 ++- 8 files changed, 155 insertions(+), 161 deletions(-) diff --git a/include/onnxruntime/core/common/profiler_common.h b/include/onnxruntime/core/common/profiler_common.h index e1a87ad31f3c0..038751f5a6e65 100644 --- a/include/onnxruntime/core/common/profiler_common.h +++ b/include/onnxruntime/core/common/profiler_common.h @@ -35,13 +35,13 @@ struct EventRecord { long long time_stamp, long long duration, std::unordered_map&& event_args) - : cat(category), - pid(process_id), - tid(thread_id), - name(std::move(event_name)), - ts(time_stamp), - dur(duration), - args(std::move(event_args)) {} + : cat(category), + pid(process_id), + tid(thread_id), + name(std::move(event_name)), + ts(time_stamp), + dur(duration), + args(std::move(event_args)) {} EventRecord(EventCategory category, int process_id, @@ -50,23 +50,25 @@ struct EventRecord { long long time_stamp, long long duration, const std::unordered_map& event_args) - : cat(category), - pid(process_id), - tid(thread_id), - name(event_name), - ts(time_stamp), - dur(duration), - args(event_args) {} + : cat(category), + pid(process_id), + tid(thread_id), + name(event_name), + ts(time_stamp), + dur(duration), + args(event_args) {} EventRecord(const EventRecord& other) - : cat(other.cat), pid(other.pid), tid(other.tid), name(other.name), - ts(other.ts), dur(other.dur), args(other.args) {} + : cat(other.cat), pid(other.pid), tid(other.tid), + name(other.name), ts(other.ts), dur(other.dur), + args(other.args) {} EventRecord(EventRecord&& other) - : cat(other.cat), pid(other.pid), tid(other.tid), name(std::move(other.name)), - ts(other.ts), dur(other.dur), args(std::move(other.args)) {} + : cat(other.cat), pid(other.pid), tid(other.tid), + name(std::move(other.name)), ts(other.ts), dur(other.dur), + args(std::move(other.args)) {} - EventRecord& operator = (const EventRecord& other) { + EventRecord& operator=(const EventRecord& other) { if (&other == this) { return *this; } @@ -80,7 +82,7 @@ struct EventRecord { return *this; } - EventRecord& operator = (EventRecord&& other) { + EventRecord& operator=(EventRecord&& other) { if (&other == this) { return *this; } diff --git a/onnxruntime/core/common/profiler.h b/onnxruntime/core/common/profiler.h index 80417e5aaef19..516ae367fc414 100644 --- a/onnxruntime/core/common/profiler.h +++ b/onnxruntime/core/common/profiler.h @@ -67,7 +67,8 @@ class Profiler { */ uint64_t GetStartTimeNs() const { return std::chrono::duration_cast( - profiling_start_time_.time_since_epoch()).count(); + profiling_start_time_.time_since_epoch()) + .count(); } /* Record a single event. Time is measured till the call of this function from diff --git a/onnxruntime/core/providers/cuda/cuda_profiler.cc b/onnxruntime/core/providers/cuda/cuda_profiler.cc index c086f4b8db3a9..6a21f299c8f8c 100644 --- a/onnxruntime/core/providers/cuda/cuda_profiler.cc +++ b/onnxruntime/core/providers/cuda/cuda_profiler.cc @@ -94,7 +94,7 @@ void CUPTIAPI CudaProfiler::BufferCompleted(CUcontext, uint32_t, uint8_t* buffer } bool CudaProfiler::StartProfiling(TimePoint profiling_start_time) { - (void) profiling_start_time; + (void)profiling_start_time; if (!enabled.test_and_set()) { if (cuptiActivityEnable(CUPTI_ACTIVITY_KIND_RUNTIME) == CUPTI_SUCCESS && cuptiActivityEnable(CUPTI_ACTIVITY_KIND_DRIVER) == CUPTI_SUCCESS && @@ -147,7 +147,7 @@ void CudaProfiler::EndProfiling(TimePoint start_time, Events& events) { for (auto& evt_iter : map_iter.second) { evt_iter.args["op_name"] = insert_iter->args["op_name"]; } - insert_iter = events.insert(insert_iter+1, map_iter.second.begin(), map_iter.second.end()); + insert_iter = events.insert(insert_iter + 1, map_iter.second.begin(), map_iter.second.end()); } else { insert_iter = events.insert(insert_iter, map_iter.second.begin(), map_iter.second.end()); } @@ -198,7 +198,7 @@ void CudaProfiler::Clear() { } } -#else // for cuda 10.x, no profiling +#else // for cuda 10.x, no profiling void CUPTIAPI CudaProfiler::BufferRequested(uint8_t**, size_t*, size_t*) {} void CUPTIAPI CudaProfiler::BufferCompleted(CUcontext, uint32_t, uint8_t*, size_t, size_t) {} diff --git a/onnxruntime/core/providers/cuda/cuda_profiler.h b/onnxruntime/core/providers/cuda/cuda_profiler.h index c089705570751..d98d53d9bc17b 100644 --- a/onnxruntime/core/providers/cuda/cuda_profiler.h +++ b/onnxruntime/core/providers/cuda/cuda_profiler.h @@ -78,7 +78,7 @@ class CudaProfiler final : public EpProfiler { void Stop(uint64_t) override{}; }; -} -} +} // namespace profiling +} // namespace onnxruntime #endif diff --git a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc index 1a9ad3d6a85b9..334a90a27549f 100644 --- a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc +++ b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc @@ -161,17 +161,17 @@ std::optional LoadEnv(const std::string& name, const std::unordered return std::nullopt; } - LOGS_DEFAULT(WARNING) << "Environment variable "<< name << " is used. It is reserved for internal testing prupose. " - "End users should opt for provider options or session options and must not rely on it."; + LOGS_DEFAULT(WARNING) << "Environment variable " << name << " is used. It is reserved for internal testing prupose. " + "End users should opt for provider options or session options and must not rely on it."; if (valid_values.find(env) == valid_values.cend()) { std::ostringstream oss; auto it = valid_values.cbegin(); oss << *it++; - while(it != valid_values.cend()) { + while (it != valid_values.cend()) { oss << ", " << *it++; } - ORT_THROW("Value of environment variable ", name," must be ",oss.str(), ", but got ", env); + ORT_THROW("Value of environment variable ", name, " must be ", oss.str(), ", but got ", env); } return env; @@ -371,7 +371,7 @@ void ReleaseCpuBufferCallback(hipStream_t /*stream*/, hipError_t /*status*/, voi // it will be cleaned up automatically at the end of this function. std::unique_ptr cpu_buffers_info = std::make_unique(); cpu_buffers_info.reset(reinterpret_cast(raw_info)); - for (auto ptr: cpu_buffers_info->buffers) { + for (auto ptr : cpu_buffers_info->buffers) { cpu_buffers_info->allocator->Free(ptr); } } @@ -411,7 +411,6 @@ Status ROCMExecutionProvider::EnqueueDeferredRelease(bool actually_defer) { ReleaseCpuBufferCallback(nullptr, hipSuccess, cpu_buffers_info.release()); } } else { - // Per // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#implicit-synchronization // cudaHostFree doesn't block stream, so a synchronitation is needed to make sure no kernels diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.h b/onnxruntime/core/providers/rocm/rocm_profiler.h index 249709634b183..40516a184c795 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.h +++ b/onnxruntime/core/providers/rocm/rocm_profiler.h @@ -23,7 +23,7 @@ class RocmProfiler final : public EpProfiler { void Start(uint64_t) override; void Stop(uint64_t) override; -private: + private: uint64_t client_handle_; TimePoint profiling_start_time_; }; @@ -45,7 +45,7 @@ class RocmProfiler final : public EpProfiler { void Stop(uint64_t) override{}; }; -} -} +} // namespace profiling +} // namespace onnxruntime #endif diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 4d8917b1f7760..7e78516665db7 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -11,19 +11,19 @@ namespace profiling { static constexpr size_t sc_activity_buffer_size = 0x4000; const std::vector RoctracerManager::hip_api_calls_to_trace = { - "hipMemcpy", - "hipMemcpy2D", - "hipMemcpyAsync", - "hipMemcpy2DAsync", - "hipMemcpyWithStream", - "hipLaunchKernel", - "hipMemset", - "hipMemsetAsync", - "hipExtModuleLaunchKernel", + "hipMemcpy", + "hipMemcpy2D", + "hipMemcpyAsync", + "hipMemcpy2DAsync", + "hipMemcpyWithStream", + "hipLaunchKernel", + "hipMemset", + "hipMemsetAsync", + "hipExtModuleLaunchKernel", }; // Implementation of RoctracerActivityBuffer -RoctracerActivityBuffer& RoctracerActivityBuffer::operator = (const RoctracerActivityBuffer& other) { +RoctracerActivityBuffer& RoctracerActivityBuffer::operator=(const RoctracerActivityBuffer& other) { if (&other == this) { return *this; } @@ -37,7 +37,7 @@ RoctracerActivityBuffer& RoctracerActivityBuffer::operator = (const RoctracerAct return *this; } -RoctracerActivityBuffer& RoctracerActivityBuffer::operator = (RoctracerActivityBuffer&& other) { +RoctracerActivityBuffer& RoctracerActivityBuffer::operator=(RoctracerActivityBuffer&& other) { if (&other == this) { return *this; } @@ -115,8 +115,7 @@ void RoctracerManager::StartLogging() { } // Requires: roctracer_manager_mutex_ must be held -void RoctracerManager::Clear() -{ +void RoctracerManager::Clear() { unprocessed_activity_buffers_.clear(); api_call_args_.clear(); unique_correlation_id_to_client_offset_.clear(); @@ -203,7 +202,7 @@ bool RoctracerManager::PushCorrelation(uint64_t client_handle, // offset to be subtracted uint64_t offset = - std::chrono::duration_cast(profiling_start_time.time_since_epoch()).count(); + std::chrono::duration_cast(profiling_start_time.time_since_epoch()).count(); auto unique_cid = external_correlation_id + offset; roctracer_activity_push_external_correlation_id(unique_cid); @@ -261,7 +260,7 @@ static inline std::string PointerToHexString(const void* ptr) { } static inline std::string MemcpyKindToString(hipMemcpyKind kind) { - switch(kind) { + switch (kind) { case hipMemcpyHostToHost: return "H2H"; case hipMemcpyHostToDevice: @@ -282,106 +281,101 @@ bool RoctracerManager::CreateEventForActivityRecord(const roctracer_record_t* re std::string name; std::unordered_map args; - switch(call_record.cid_) { - case HIP_API_ID_hipLaunchKernel: { - auto const& launch_args = call_record.api_data_.args.hipLaunchKernel; - name = demangle(hipKernelNameRefByPtr(launch_args.function_address, - launch_args.stream)); - - args = { - {"stream", PointerToHexString((void*)(launch_args.stream))}, - {"grid_x", std::to_string(launch_args.numBlocks.x)}, - {"grid_y", std::to_string(launch_args.numBlocks.y)}, - {"grid_z", std::to_string(launch_args.numBlocks.z)}, - {"block_x", std::to_string(launch_args.dimBlocks.x)}, - {"block_y", std::to_string(launch_args.dimBlocks.y)}, - {"block_z", std::to_string(launch_args.dimBlocks.z)} - }; - break; - } + switch (call_record.cid_) { + case HIP_API_ID_hipLaunchKernel: { + auto const& launch_args = call_record.api_data_.args.hipLaunchKernel; + name = demangle(hipKernelNameRefByPtr(launch_args.function_address, + launch_args.stream)); + + args = { + {"stream", PointerToHexString((void*)(launch_args.stream))}, + {"grid_x", std::to_string(launch_args.numBlocks.x)}, + {"grid_y", std::to_string(launch_args.numBlocks.y)}, + {"grid_z", std::to_string(launch_args.numBlocks.z)}, + {"block_x", std::to_string(launch_args.dimBlocks.x)}, + {"block_y", std::to_string(launch_args.dimBlocks.y)}, + {"block_z", std::to_string(launch_args.dimBlocks.z)}}; + break; + } - case HIP_API_ID_hipMemset: - case HIP_API_ID_hipMemsetAsync: { - auto const& launch_args = call_record.api_data_.args; - name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); - - args = { - {"stream", call_record.cid_ == HIP_API_ID_hipMemset - ? "0" - : PointerToHexString((void*)launch_args.hipMemsetAsync.stream)}, - {"dst", PointerToHexString(launch_args.hipMemset.dst)}, - {"size", std::to_string(launch_args.hipMemset.sizeBytes)}, - {"value", std::to_string(launch_args.hipMemset.value)} - }; - break; - } + case HIP_API_ID_hipMemset: + case HIP_API_ID_hipMemsetAsync: { + auto const& launch_args = call_record.api_data_.args; + name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); + + args = { + {"stream", call_record.cid_ == HIP_API_ID_hipMemset + ? "0" + : PointerToHexString((void*)launch_args.hipMemsetAsync.stream)}, + {"dst", PointerToHexString(launch_args.hipMemset.dst)}, + {"size", std::to_string(launch_args.hipMemset.sizeBytes)}, + {"value", std::to_string(launch_args.hipMemset.value)}}; + break; + } - case HIP_API_ID_hipMemcpy: - case HIP_API_ID_hipMemcpyAsync: - case HIP_API_ID_hipMemcpyWithStream: { - auto const& launch_args = call_record.api_data_.args; - name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); - - args = { - {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy - ? "0" - : PointerToHexString((void*)launch_args.hipMemcpyAsync.stream)}, - {"src", PointerToHexString(launch_args.hipMemcpy.src)}, - {"dst", PointerToHexString(launch_args.hipMemcpy.dst)}, - {"kind", MemcpyKindToString(launch_args.hipMemcpy.kind)} - }; - break; - } + case HIP_API_ID_hipMemcpy: + case HIP_API_ID_hipMemcpyAsync: + case HIP_API_ID_hipMemcpyWithStream: { + auto const& launch_args = call_record.api_data_.args; + name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); + + args = { + {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy + ? "0" + : PointerToHexString((void*)launch_args.hipMemcpyAsync.stream)}, + {"src", PointerToHexString(launch_args.hipMemcpy.src)}, + {"dst", PointerToHexString(launch_args.hipMemcpy.dst)}, + {"kind", MemcpyKindToString(launch_args.hipMemcpy.kind)}}; + break; + } - case HIP_API_ID_hipMemcpy2D: - case HIP_API_ID_hipMemcpy2DAsync: { - auto const& launch_args = call_record.api_data_.args; - name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); - - args = { - {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy2D - ? "0" - : PointerToHexString((void*)launch_args.hipMemcpy2DAsync.stream)}, - {"src", PointerToHexString(launch_args.hipMemcpy2D.src)}, - {"dst", PointerToHexString(launch_args.hipMemcpy2D.dst)}, - {"spitch", std::to_string(launch_args.hipMemcpy2D.spitch)}, - {"dpitch", std::to_string(launch_args.hipMemcpy2D.dpitch)}, - {"width", std::to_string(launch_args.hipMemcpy2D.width)}, - {"height", std::to_string(launch_args.hipMemcpy2D.height)}, - {"kind", MemcpyKindToString(launch_args.hipMemcpy2D.kind)} - }; - break; - } + case HIP_API_ID_hipMemcpy2D: + case HIP_API_ID_hipMemcpy2DAsync: { + auto const& launch_args = call_record.api_data_.args; + name = roctracer_op_string(call_record.domain_, call_record.cid_, 0); + + args = { + {"stream", call_record.cid_ == HIP_API_ID_hipMemcpy2D + ? "0" + : PointerToHexString((void*)launch_args.hipMemcpy2DAsync.stream)}, + {"src", PointerToHexString(launch_args.hipMemcpy2D.src)}, + {"dst", PointerToHexString(launch_args.hipMemcpy2D.dst)}, + {"spitch", std::to_string(launch_args.hipMemcpy2D.spitch)}, + {"dpitch", std::to_string(launch_args.hipMemcpy2D.dpitch)}, + {"width", std::to_string(launch_args.hipMemcpy2D.width)}, + {"height", std::to_string(launch_args.hipMemcpy2D.height)}, + {"kind", MemcpyKindToString(launch_args.hipMemcpy2D.kind)}}; + break; + } - case HIP_API_ID_hipExtModuleLaunchKernel: { - auto const& launch_args = call_record.api_data_.args.hipExtModuleLaunchKernel; - name = demangle(hipKernelNameRef(launch_args.f)); - - args = { - {"stream", PointerToHexString((void*)launch_args.hStream)}, - {"grid_x", std::to_string(launch_args.globalWorkSizeX)}, - {"grid_y", std::to_string(launch_args.globalWorkSizeY)}, - {"grid_z", std::to_string(launch_args.globalWorkSizeZ)}, - {"block_x", std::to_string(launch_args.localWorkSizeX)}, - {"block_y", std::to_string(launch_args.localWorkSizeY)}, - {"block_z", std::to_string(launch_args.localWorkSizeZ)}, - }; - break; - } + case HIP_API_ID_hipExtModuleLaunchKernel: { + auto const& launch_args = call_record.api_data_.args.hipExtModuleLaunchKernel; + name = demangle(hipKernelNameRef(launch_args.f)); + + args = { + {"stream", PointerToHexString((void*)launch_args.hStream)}, + {"grid_x", std::to_string(launch_args.globalWorkSizeX)}, + {"grid_y", std::to_string(launch_args.globalWorkSizeY)}, + {"grid_z", std::to_string(launch_args.globalWorkSizeZ)}, + {"block_x", std::to_string(launch_args.localWorkSizeX)}, + {"block_y", std::to_string(launch_args.localWorkSizeY)}, + {"block_z", std::to_string(launch_args.localWorkSizeZ)}, + }; + break; + } - default: - return false; + default: + return false; } - new (&event) EventRecord { - /* cat = */ EventCategory::KERNEL_EVENT, - /* pid = */ -1, - /* tid = */ -1, - /* name = */ std::move(name), - /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, - /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, - /* args = */ std::move(args) - }; + new (&event) EventRecord{ + /* cat = */ EventCategory::KERNEL_EVENT, + /* pid = */ -1, + /* tid = */ -1, + /* name = */ std::move(name), + /* ts = */ (int64_t)(record->begin_ns - start_time_ns) / 1000, + /* dur = */ (int64_t)(record->end_ns - record->begin_ns) / 1000, + /* args = */ std::move(args)}; return true; } @@ -425,7 +419,7 @@ void RoctracerManager::ProcessActivityBuffers(const std::vector(buffer.GetData()); auto data_end = reinterpret_cast(buffer.GetData() + buffer.GetSize()); - for ( ; current_record < data_end; roctracer_next_record(current_record, ¤t_record)) { + for (; current_record < data_end; roctracer_next_record(current_record, ¤t_record)) { EventRecord event; if (current_record->domain == ACTIVITY_DOMAIN_EXT_API) { roctracer_correlation_to_unique_correlation_[current_record->correlation_id] = current_record->external_id; @@ -449,7 +443,7 @@ void RoctracerManager::ProcessActivityBuffers(const std::vectorcorrelation_id); if (api_it == api_call_args_.end()) { - // we're not tracking this activity, ignore it + // we're not tracking this activity, ignore it continue; } diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 22fda2b00e4f6..af60f4d5bc38f 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -16,31 +16,30 @@ #include "core/common/profiler_common.h" - -namespace onnxruntime{ +namespace onnxruntime { namespace profiling { class RoctracerActivityBuffer { -public: + public: RoctracerActivityBuffer() - : data_(nullptr), size_(0) {} + : data_(nullptr), size_(0) {} RoctracerActivityBuffer(const uint8_t* data, size_t size) - : data_((uint8_t*)malloc(size)), size_(size) { + : data_((uint8_t*)malloc(size)), size_(size) { memcpy(data_, data, size); } RoctracerActivityBuffer(const RoctracerActivityBuffer& other) - : RoctracerActivityBuffer(other.data_, other.size_) {} + : RoctracerActivityBuffer(other.data_, other.size_) {} RoctracerActivityBuffer(RoctracerActivityBuffer&& other) - : RoctracerActivityBuffer() { + : RoctracerActivityBuffer() { std::swap(data_, other.data_); std::swap(size_, other.size_); } - RoctracerActivityBuffer& operator = (const RoctracerActivityBuffer& other); - RoctracerActivityBuffer& operator = (RoctracerActivityBuffer&& other); + RoctracerActivityBuffer& operator=(const RoctracerActivityBuffer& other); + RoctracerActivityBuffer& operator=(RoctracerActivityBuffer&& other); ~RoctracerActivityBuffer(); @@ -49,7 +48,7 @@ class RoctracerActivityBuffer { const uint8_t* GetData() const { return data_; } size_t GetSize() const { return size_; } -private: + private: uint8_t* data_; size_t size_; }; @@ -57,14 +56,13 @@ class RoctracerActivityBuffer { struct ApiCallRecord { uint32_t domain_; uint32_t cid_; - hip_api_data_t api_data_ {}; + hip_api_data_t api_data_{}; }; -class RoctracerManager -{ -public: +class RoctracerManager { + public: RoctracerManager(const RoctracerManager&) = delete; - RoctracerManager& operator = (const RoctracerManager&) = delete; + RoctracerManager& operator=(const RoctracerManager&) = delete; RoctracerManager() = default; ~RoctracerManager(); @@ -80,7 +78,7 @@ class RoctracerManager void PopCorrelation(uint64_t& popped_correlation_id); bool PopCorrelation(); -private: + private: static void ActivityCallback(const char* begin, const char* end, void* arg); static void ApiCallback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg); void ProcessActivityBuffers(const std::vector& buffers, From 67a3508e02fa97b7a767371f45e6cafa2a551a98 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Thu, 3 Nov 2022 21:02:05 +0000 Subject: [PATCH 32/41] Use ORT_UNUSED_PARAMETER --- onnxruntime/core/providers/cuda/cuda_profiler.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/providers/cuda/cuda_profiler.cc b/onnxruntime/core/providers/cuda/cuda_profiler.cc index 6a21f299c8f8c..7abb086fd141b 100644 --- a/onnxruntime/core/providers/cuda/cuda_profiler.cc +++ b/onnxruntime/core/providers/cuda/cuda_profiler.cc @@ -94,7 +94,7 @@ void CUPTIAPI CudaProfiler::BufferCompleted(CUcontext, uint32_t, uint8_t* buffer } bool CudaProfiler::StartProfiling(TimePoint profiling_start_time) { - (void)profiling_start_time; + ORT_UNUSED_PARAMETER(profiling_start_time); if (!enabled.test_and_set()) { if (cuptiActivityEnable(CUPTI_ACTIVITY_KIND_RUNTIME) == CUPTI_SUCCESS && cuptiActivityEnable(CUPTI_ACTIVITY_KIND_DRIVER) == CUPTI_SUCCESS && From a2fd04d78322b71791a0d5a7b2add1d192044240 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Mon, 7 Nov 2022 23:32:49 +0000 Subject: [PATCH 33/41] review rework --- .../onnxruntime/core/common/profiler_common.h | 50 ++++--------------- .../core/providers/cuda/cuda_profiler.cc | 3 +- .../core/providers/rocm/rocm_profiler.cc | 5 +- .../core/providers/rocm/roctracer_manager.h | 13 +++-- 4 files changed, 20 insertions(+), 51 deletions(-) diff --git a/include/onnxruntime/core/common/profiler_common.h b/include/onnxruntime/core/common/profiler_common.h index 038751f5a6e65..635041761c6f5 100644 --- a/include/onnxruntime/core/common/profiler_common.h +++ b/include/onnxruntime/core/common/profiler_common.h @@ -63,45 +63,17 @@ struct EventRecord { name(other.name), ts(other.ts), dur(other.dur), args(other.args) {} - EventRecord(EventRecord&& other) - : cat(other.cat), pid(other.pid), tid(other.tid), - name(std::move(other.name)), ts(other.ts), dur(other.dur), - args(std::move(other.args)) {} - - EventRecord& operator=(const EventRecord& other) { - if (&other == this) { - return *this; - } - cat = other.cat; - pid = other.pid; - tid = other.tid; - name = other.name; - ts = other.ts; - dur = other.dur; - args = other.args; - return *this; - } - - EventRecord& operator=(EventRecord&& other) { - if (&other == this) { - return *this; - } - cat = other.cat; - pid = other.pid; - tid = other.tid; - std::swap(name, other.name); - dur = other.dur; - std::swap(args, other.args); - return *this; - } - - EventCategory cat; - int pid; - int tid; - std::string name; - long long ts; - long long dur; - std::unordered_map args; + EventRecord(EventRecord&& other) = default; + EventRecord& operator=(const EventRecord& other) = default; + EventRecord& operator=(EventRecord&& other) = default; + + EventCategory cat = EventCategory::API_EVENT; + int pid = -1; + int tid = -1; + std::string name{}; + long long ts = 0; + long long dur = 0; + std::unordered_map args{}; }; using Events = std::vector; diff --git a/onnxruntime/core/providers/cuda/cuda_profiler.cc b/onnxruntime/core/providers/cuda/cuda_profiler.cc index 7abb086fd141b..01048451748f2 100644 --- a/onnxruntime/core/providers/cuda/cuda_profiler.cc +++ b/onnxruntime/core/providers/cuda/cuda_profiler.cc @@ -93,8 +93,7 @@ void CUPTIAPI CudaProfiler::BufferCompleted(CUcontext, uint32_t, uint8_t* buffer free(buffer); } -bool CudaProfiler::StartProfiling(TimePoint profiling_start_time) { - ORT_UNUSED_PARAMETER(profiling_start_time); +bool CudaProfiler::StartProfiling(TimePoint /* profiling_start_time */) { if (!enabled.test_and_set()) { if (cuptiActivityEnable(CUPTI_ACTIVITY_KIND_RUNTIME) == CUPTI_SUCCESS && cuptiActivityEnable(CUPTI_ACTIVITY_KIND_DRIVER) == CUPTI_SUCCESS && diff --git a/onnxruntime/core/providers/rocm/rocm_profiler.cc b/onnxruntime/core/providers/rocm/rocm_profiler.cc index 6559ac59f0c9e..6fb0221a378c9 100644 --- a/onnxruntime/core/providers/rocm/rocm_profiler.cc +++ b/onnxruntime/core/providers/rocm/rocm_profiler.cc @@ -55,9 +55,8 @@ void RocmProfiler::EndProfiling(TimePoint start_time, Events& events) { for (auto& evt : map_iter.second) { evt.args["op_name"] = event_iter->args["op_name"]; - // roctracer timestamps don't use Jan 1 1970 as an epoch, - // not sure what epoch it uses, but we adjust the timestamp - // here to something sensible. + // roctracer doesn't use Jan 1 1970 as an epoch for its timestamps. + // So, we adjust the timestamp here to something sensible. evt.ts = event_iter->ts + increment; ++increment; } diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index af60f4d5bc38f..821bedb110aaf 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -4,8 +4,6 @@ #include #include #include -#include -#include #include #include @@ -15,6 +13,7 @@ #include #include "core/common/profiler_common.h" +#include "core/common/inlined_containers.h" namespace onnxruntime { namespace profiling { @@ -95,20 +94,20 @@ class RoctracerManager { static constexpr uint32_t HipOpMarker = 4606; std::mutex unprocessed_activity_buffers_lock_; - std::vector unprocessed_activity_buffers_; + InlinedVector unprocessed_activity_buffers_; std::mutex activity_buffer_processor_mutex_; std::mutex api_call_args_lock_; - std::unordered_map api_call_args_; + InlinedHashMap api_call_args_; // Keyed on unique_correlation_id -> (client_id/client_handle, offset) // unique_correlation_id - offset == external_correlation_id - std::unordered_map> unique_correlation_id_to_client_offset_; + InlinedHashMap> unique_correlation_id_to_client_offset_; // Keyed on roctracer_correlation_id -> unique_correlation_id - std::unordered_map roctracer_correlation_to_unique_correlation_; + InlinedHashMap roctracer_correlation_to_unique_correlation_; // client_id/client_handle -> external_correlation_id -> events - std::unordered_map> per_client_events_by_ext_correlation_; + InlinedHashMap> per_client_events_by_ext_correlation_; uint64_t next_client_id_ = 1; uint64_t num_active_clients_ = 0; bool logging_enabled_ = false; From 2161b02262f1f4eb902a35277101dbf113518208 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 8 Nov 2022 00:09:48 +0000 Subject: [PATCH 34/41] review rework --- .../core/providers/rocm/roctracer_manager.cc | 24 ++++++------------- .../core/providers/rocm/roctracer_manager.h | 16 ++++++------- 2 files changed, 14 insertions(+), 26 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 7e78516665db7..d11d65152bd96 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -27,13 +27,10 @@ RoctracerActivityBuffer& RoctracerActivityBuffer::operator=(const RoctracerActiv if (&other == this) { return *this; } - if (data_ != nullptr) { - free(data_); - } size_ = other.size_; - data_ = (uint8_t*)malloc(size_); - memcpy(data_, other.data_, size_); + data_ = std::make_unique(other.size_); + memcpy(data_.get(), other.data_, size_); return *this; } @@ -46,13 +43,6 @@ RoctracerActivityBuffer& RoctracerActivityBuffer::operator=(RoctracerActivityBuf return *this; } -RoctracerActivityBuffer::~RoctracerActivityBuffer() { - if (data_ != nullptr) { - free(data_); - data_ = nullptr; - } -} - // Implementation of RoctracerManager RoctracerManager& RoctracerManager::GetInstance() { static RoctracerManager instance; @@ -150,7 +140,7 @@ void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_ti roctracer_flush_activity(); } - std::vector activity_buffers; + InlinedVector activity_buffers; { std::lock_guard lock(unprocessed_activity_buffers_lock_); std::swap(unprocessed_activity_buffers_, activity_buffers); @@ -225,7 +215,7 @@ void RoctracerManager::PopCorrelation(uint64_t& popped_external_correlation_id) void RoctracerManager::ActivityCallback(const char* begin, const char* end, void* arg) { size_t size = end - begin; - RoctracerActivityBuffer activity_buffer{reinterpret_cast(begin), size}; + RoctracerActivityBuffer activity_buffer{reinterpret_cast(begin), size}; auto& instance = GetInstance(); { std::lock_guard lock(instance.unprocessed_activity_buffers_lock_); @@ -395,7 +385,7 @@ Events* RoctracerManager::GetEventListForUniqueCorrelationId(uint64_t unique_cor return &event_list; } -void RoctracerManager::MapEventsToClient(uint64_t unique_correlation_id, Events&& events) { +void RoctracerManager::MapEventsToClient(uint64_t unique_correlation_id, InlinedVector&& events) { auto p_event_list = GetEventListForUniqueCorrelationId(unique_correlation_id); if (p_event_list != nullptr) { p_event_list->insert(p_event_list->end(), @@ -411,9 +401,9 @@ void RoctracerManager::MapEventToClient(uint64_t unique_correlation_id, EventRec } } -void RoctracerManager::ProcessActivityBuffers(const std::vector& buffers, +void RoctracerManager::ProcessActivityBuffers(const InlinedVector& buffers, const TimePoint& start_time) { - std::unordered_map> events_pending_client_mapping; + InlinedHashMap> events_pending_client_mapping; auto start_time_ns = std::chrono::duration_cast(start_time.time_since_epoch()).count(); for (auto const& buffer : buffers) { diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 821bedb110aaf..5400875e471f7 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -23,13 +23,13 @@ class RoctracerActivityBuffer { RoctracerActivityBuffer() : data_(nullptr), size_(0) {} - RoctracerActivityBuffer(const uint8_t* data, size_t size) - : data_((uint8_t*)malloc(size)), size_(size) { - memcpy(data_, data, size); + RoctracerActivityBuffer(const char* data, size_t size) + : data_(std::make_unique(new char[size])), size_(size) { + memcpy(data_.get(), data, size); } RoctracerActivityBuffer(const RoctracerActivityBuffer& other) - : RoctracerActivityBuffer(other.data_, other.size_) {} + : RoctracerActivityBuffer(other.data_.get(), other.size_) {} RoctracerActivityBuffer(RoctracerActivityBuffer&& other) : RoctracerActivityBuffer() { @@ -40,15 +40,13 @@ class RoctracerActivityBuffer { RoctracerActivityBuffer& operator=(const RoctracerActivityBuffer& other); RoctracerActivityBuffer& operator=(RoctracerActivityBuffer&& other); - ~RoctracerActivityBuffer(); - // accessors - uint8_t* GetData() { return data_; } - const uint8_t* GetData() const { return data_; } + char* GetData() { return data_.get(); } + const char* GetData() const { return data_.get(); } size_t GetSize() const { return size_; } private: - uint8_t* data_; + std::unique_ptr data_; size_t size_; }; From 2971c9e166b942d11742ee3e43f75d7576d41e0d Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 8 Nov 2022 01:02:27 +0000 Subject: [PATCH 35/41] Review rework --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 10 +++++----- onnxruntime/core/providers/rocm/roctracer_manager.h | 4 ++-- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index d11d65152bd96..406740d448165 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -30,7 +30,7 @@ RoctracerActivityBuffer& RoctracerActivityBuffer::operator=(const RoctracerActiv size_ = other.size_; data_ = std::make_unique(other.size_); - memcpy(data_.get(), other.data_, size_); + memcpy(data_.get(), other.data_.get(), size_); return *this; } @@ -140,7 +140,7 @@ void RoctracerManager::Consume(uint64_t client_handle, const TimePoint& start_ti roctracer_flush_activity(); } - InlinedVector activity_buffers; + std::vector activity_buffers; { std::lock_guard lock(unprocessed_activity_buffers_lock_); std::swap(unprocessed_activity_buffers_, activity_buffers); @@ -385,7 +385,7 @@ Events* RoctracerManager::GetEventListForUniqueCorrelationId(uint64_t unique_cor return &event_list; } -void RoctracerManager::MapEventsToClient(uint64_t unique_correlation_id, InlinedVector&& events) { +void RoctracerManager::MapEventsToClient(uint64_t unique_correlation_id, std::vector&& events) { auto p_event_list = GetEventListForUniqueCorrelationId(unique_correlation_id); if (p_event_list != nullptr) { p_event_list->insert(p_event_list->end(), @@ -401,9 +401,9 @@ void RoctracerManager::MapEventToClient(uint64_t unique_correlation_id, EventRec } } -void RoctracerManager::ProcessActivityBuffers(const InlinedVector& buffers, +void RoctracerManager::ProcessActivityBuffers(const std::vector& buffers, const TimePoint& start_time) { - InlinedHashMap> events_pending_client_mapping; + InlinedHashMap> events_pending_client_mapping; auto start_time_ns = std::chrono::duration_cast(start_time.time_since_epoch()).count(); for (auto const& buffer : buffers) { diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 5400875e471f7..92cbffd6f9ec7 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -24,7 +24,7 @@ class RoctracerActivityBuffer { : data_(nullptr), size_(0) {} RoctracerActivityBuffer(const char* data, size_t size) - : data_(std::make_unique(new char[size])), size_(size) { + : data_(std::make_unique(size)), size_(size) { memcpy(data_.get(), data, size); } @@ -92,7 +92,7 @@ class RoctracerManager { static constexpr uint32_t HipOpMarker = 4606; std::mutex unprocessed_activity_buffers_lock_; - InlinedVector unprocessed_activity_buffers_; + std::vector unprocessed_activity_buffers_; std::mutex activity_buffer_processor_mutex_; std::mutex api_call_args_lock_; InlinedHashMap api_call_args_; From 2a44480f1157fa45279a622b7672c5e3c01d3b50 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 8 Nov 2022 01:07:44 +0000 Subject: [PATCH 36/41] Minor fix --- include/onnxruntime/core/common/profiler_common.h | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/include/onnxruntime/core/common/profiler_common.h b/include/onnxruntime/core/common/profiler_common.h index 635041761c6f5..07d6cc101b31c 100644 --- a/include/onnxruntime/core/common/profiler_common.h +++ b/include/onnxruntime/core/common/profiler_common.h @@ -58,11 +58,7 @@ struct EventRecord { dur(duration), args(event_args) {} - EventRecord(const EventRecord& other) - : cat(other.cat), pid(other.pid), tid(other.tid), - name(other.name), ts(other.ts), dur(other.dur), - args(other.args) {} - + EventRecord(const EventRecord& other) = default; EventRecord(EventRecord&& other) = default; EventRecord& operator=(const EventRecord& other) = default; EventRecord& operator=(EventRecord&& other) = default; From 65578c39b9cc7a1ec4c42624ff25ad2f0d1409ee Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 8 Nov 2022 02:12:55 +0000 Subject: [PATCH 37/41] Review rework --- onnxruntime/core/providers/rocm/rocm_execution_provider.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc index 334a90a27549f..2f83299e2cc41 100644 --- a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc +++ b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc @@ -406,7 +406,7 @@ Status ROCMExecutionProvider::EnqueueDeferredRelease(bool actually_defer) { if (cpu_buffers_info->allocator->Info().alloc_type == OrtArenaAllocator) { if (actually_defer) { // Release memory asynchronously to avoid blocking the compute stream. - HIP_RETURN_IF_ERROR(hipStreamAddCallback(stream, ReleaseCpuBufferCallback, cpu_buffers_info.release(), 0)); + HIP_CALL_THROW(hipStreamAddCallback(stream, ReleaseCpuBufferCallback, cpu_buffers_info.release(), 0)); } else { ReleaseCpuBufferCallback(nullptr, hipSuccess, cpu_buffers_info.release()); } From c04652d3a4ac4de6d437ada163db8ab4c94b08b3 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 8 Nov 2022 02:27:26 +0000 Subject: [PATCH 38/41] Revert prev change --- onnxruntime/core/providers/rocm/rocm_execution_provider.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc index 2f83299e2cc41..334a90a27549f 100644 --- a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc +++ b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc @@ -406,7 +406,7 @@ Status ROCMExecutionProvider::EnqueueDeferredRelease(bool actually_defer) { if (cpu_buffers_info->allocator->Info().alloc_type == OrtArenaAllocator) { if (actually_defer) { // Release memory asynchronously to avoid blocking the compute stream. - HIP_CALL_THROW(hipStreamAddCallback(stream, ReleaseCpuBufferCallback, cpu_buffers_info.release(), 0)); + HIP_RETURN_IF_ERROR(hipStreamAddCallback(stream, ReleaseCpuBufferCallback, cpu_buffers_info.release(), 0)); } else { ReleaseCpuBufferCallback(nullptr, hipSuccess, cpu_buffers_info.release()); } From ac86163e1f739ae2be9e6f2bc2300d3c5cadcb4b Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 8 Nov 2022 03:49:24 +0000 Subject: [PATCH 39/41] Review rework --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 406740d448165..581f0e1e4aaa7 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -8,7 +8,7 @@ namespace onnxruntime { namespace profiling { // allocate a 16K buffer for recording async activities -static constexpr size_t sc_activity_buffer_size = 0x4000; +static constexpr size_t kActivityBufferSize = 0x4000; const std::vector RoctracerManager::hip_api_calls_to_trace = { "hipMemcpy", @@ -82,7 +82,7 @@ void RoctracerManager::StartLogging() { roctracer_properties_t hcc_cb_properties; memset(&hcc_cb_properties, 0, sizeof(roctracer_properties_t)); - hcc_cb_properties.buffer_size = sc_activity_buffer_size; + hcc_cb_properties.buffer_size = kActivityBufferSize; hcc_cb_properties.buffer_callback_fun = ActivityCallback; roctracer_open_pool(&hcc_cb_properties); From ab2d864473aa08282b1eadf268fc2b914f13ff25 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 8 Nov 2022 18:21:57 +0000 Subject: [PATCH 40/41] Review rework --- onnxruntime/core/providers/rocm/roctracer_manager.cc | 9 ++++----- onnxruntime/core/providers/rocm/roctracer_manager.h | 9 ++++++--- 2 files changed, 10 insertions(+), 8 deletions(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.cc b/onnxruntime/core/providers/rocm/roctracer_manager.cc index 581f0e1e4aaa7..e02abd93af3be 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.cc +++ b/onnxruntime/core/providers/rocm/roctracer_manager.cc @@ -403,7 +403,6 @@ void RoctracerManager::MapEventToClient(uint64_t unique_correlation_id, EventRec void RoctracerManager::ProcessActivityBuffers(const std::vector& buffers, const TimePoint& start_time) { - InlinedHashMap> events_pending_client_mapping; auto start_time_ns = std::chrono::duration_cast(start_time.time_since_epoch()).count(); for (auto const& buffer : buffers) { @@ -415,14 +414,14 @@ void RoctracerManager::ProcessActivityBuffers(const std::vectorcorrelation_id] = current_record->external_id; // check for any events pending client mapping on this correlation - auto pending_it = events_pending_client_mapping.find(current_record->correlation_id); - if (pending_it == events_pending_client_mapping.end()) { + auto pending_it = events_pending_client_mapping_.find(current_record->correlation_id); + if (pending_it == events_pending_client_mapping_.end()) { continue; } // we have one or more pending events, map them to the client MapEventsToClient(current_record->external_id, std::move(pending_it->second)); - events_pending_client_mapping.erase(current_record->correlation_id); + events_pending_client_mapping_.erase(pending_it); // no additional events to be mapped for this record continue; } else if (current_record->domain == ACTIVITY_DOMAIN_HIP_OPS) { @@ -453,7 +452,7 @@ void RoctracerManager::ProcessActivityBuffers(const std::vectorcorrelation_id); if (ext_corr_it == roctracer_correlation_to_unique_correlation_.end()) { // defer the processing of this event - events_pending_client_mapping[current_record->correlation_id].emplace_back(std::move(event)); + events_pending_client_mapping_[current_record->correlation_id].emplace_back(std::move(event)); continue; } MapEventToClient(ext_corr_it->second, std::move(event)); diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 92cbffd6f9ec7..33c467518b0cb 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -58,9 +58,7 @@ struct ApiCallRecord { class RoctracerManager { public: - RoctracerManager(const RoctracerManager&) = delete; - RoctracerManager& operator=(const RoctracerManager&) = delete; - RoctracerManager() = default; + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(); ~RoctracerManager(); static RoctracerManager& GetInstance(); @@ -76,6 +74,7 @@ class RoctracerManager { bool PopCorrelation(); private: + RoctracerManager() = default; static void ActivityCallback(const char* begin, const char* end, void* arg); static void ApiCallback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg); void ProcessActivityBuffers(const std::vector& buffers, @@ -111,6 +110,10 @@ class RoctracerManager { bool logging_enabled_ = false; std::mutex roctracer_manager_mutex_; + // Keyed on roctracer correlation_id, keeps track of activity records + // for which we haven't established the external_correlation_id yet. + InlinedHashMap> events_pending_client_mapping_; + // The api calls to track static const std::vector hip_api_calls_to_trace; }; From 5474c0ae592f2478b1458233e99d5404781ee702 Mon Sep 17 00:00:00 2001 From: Abhishek Udupa Date: Tue, 8 Nov 2022 18:47:02 +0000 Subject: [PATCH 41/41] Include typename in macro --- onnxruntime/core/providers/rocm/roctracer_manager.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/providers/rocm/roctracer_manager.h b/onnxruntime/core/providers/rocm/roctracer_manager.h index 33c467518b0cb..31ba862d5d65c 100644 --- a/onnxruntime/core/providers/rocm/roctracer_manager.h +++ b/onnxruntime/core/providers/rocm/roctracer_manager.h @@ -58,7 +58,7 @@ struct ApiCallRecord { class RoctracerManager { public: - ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(); + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(RoctracerManager); ~RoctracerManager(); static RoctracerManager& GetInstance();