From 8abbf9475b6809d4dce303651331e0e6cda4e18c Mon Sep 17 00:00:00 2001 From: David Yat Sin Date: Wed, 29 Nov 2023 19:13:33 +0000 Subject: [PATCH] PC Sampling: Implement flush Flush is used by the client to retrieve data that are currently stored in the buffers. This is used by the client to retrieve current data when the buffers are not full. Change-Id: Ib8304dcdfb2797cb060ec72df4970d95cf6be348 --- runtime/hsa-runtime/core/inc/amd_gpu_agent.h | 1 + .../core/runtime/amd_gpu_agent.cpp | 60 ++++++++++++++++++- 2 files changed, 60 insertions(+), 1 deletion(-) diff --git a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h index d5abfbbf3b..ebf2e33259 100644 --- a/runtime/hsa-runtime/core/inc/amd_gpu_agent.h +++ b/runtime/hsa-runtime/core/inc/amd_gpu_agent.h @@ -734,6 +734,7 @@ class GpuAgent : public GpuAgentInt { uint8_t* host_buffer_wrap_pos; uint8_t* host_write_ptr; uint8_t* host_read_ptr; + std::mutex host_buffer_mutex; uint32_t which_buffer; uint64_t* old_val; diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index a4c470a97d..208028a6f0 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -3030,6 +3030,7 @@ void GpuAgent::PcSamplingThread() { } while (true); HSA::hsa_signal_store_screlease(done_sig[which_buffer], 1); + std::lock_guard lock(ht_data.host_buffer_mutex); if (PcSamplingFlushHostTrapDeviceBuffers(session) != HSA_STATUS_SUCCESS) goto thread_exit; @@ -3094,7 +3095,64 @@ void GpuAgent::PcSamplingThreadRun(void* _agent) { } hsa_status_t GpuAgent::PcSamplingFlush(pcs::PcsRuntime::PcSamplingSession& session) { - // TODO: implement me + pcs_hosttrap_t& ht_data = pcs_hosttrap_data_; + + uint8_t* host_buffer_begin = ht_data.host_buffer; + uint8_t* host_buffer_end = ht_data.host_buffer + ht_data.host_buffer_size; + + size_t bytes_before_wrap; + size_t bytes_after_wrap; + + std::lock_guard lock(ht_data.host_buffer_mutex); + if (PcSamplingFlushHostTrapDeviceBuffers(session) != HSA_STATUS_SUCCESS) + return HSA_STATUS_ERROR; + + assert(ht_data.host_read_ptr >= host_buffer_begin && ht_data.host_read_ptr < host_buffer_end); + assert(ht_data.host_write_ptr >= host_buffer_begin && ht_data.host_write_ptr < host_buffer_end); + assert(ht_data.host_buffer_wrap_pos ? (ht_data.host_read_ptr > ht_data.host_write_ptr) + : (ht_data.host_read_ptr <= ht_data.host_write_ptr)); + + if (ht_data.host_buffer_wrap_pos) { + assert(ht_data.host_buffer_wrap_pos <= host_buffer_end && + ht_data.host_buffer_wrap_pos > host_buffer_begin); + assert(ht_data.host_read_ptr <= ht_data.host_buffer_wrap_pos); + + // Wrapped around + bytes_before_wrap = ht_data.host_buffer_wrap_pos - ht_data.host_read_ptr; + bytes_after_wrap = ht_data.host_write_ptr - host_buffer_begin; + + while (bytes_before_wrap > 0) { + size_t bytes_to_copy = std::min(bytes_before_wrap, session.buffer_size()); + + session.HandleSampleData(ht_data.host_read_ptr, bytes_to_copy, NULL, 0, 0); + ht_data.host_read_ptr += bytes_to_copy; + bytes_before_wrap = ht_data.host_buffer_wrap_pos - ht_data.host_read_ptr; + } + + assert(ht_data.host_read_ptr == ht_data.host_buffer_wrap_pos); + ht_data.host_buffer_wrap_pos = 0; + ht_data.host_read_ptr = host_buffer_begin; + + while (bytes_after_wrap > 0) { + size_t bytes_to_copy = std::min(bytes_after_wrap, session.buffer_size()); + + session.HandleSampleData(ht_data.host_read_ptr, bytes_to_copy, NULL, 0, 0); + ht_data.host_read_ptr += bytes_to_copy; + bytes_after_wrap = ht_data.host_write_ptr - ht_data.host_read_ptr; + } + } else { + bytes_before_wrap = ht_data.host_write_ptr - ht_data.host_read_ptr; + + while (bytes_before_wrap) { + size_t bytes_to_copy = std::min(bytes_before_wrap, session.buffer_size()); + assert(ht_data.host_read_ptr >= host_buffer_begin && + ht_data.host_read_ptr + bytes_to_copy <= host_buffer_end); + + session.HandleSampleData(ht_data.host_read_ptr, bytes_to_copy, NULL, 0, 0); + ht_data.host_read_ptr += bytes_to_copy; + bytes_before_wrap = ht_data.host_write_ptr - ht_data.host_read_ptr; + } + } return HSA_STATUS_SUCCESS; }