From 5177d17f5d3e6101f68800b2e0b3df5cb6bde20b Mon Sep 17 00:00:00 2001 From: David Yat Sin Date: Tue, 5 Sep 2023 17:05:10 +0000 Subject: [PATCH] PC Sampling: Push data to PC Sampling client Each time there is enough data to fill the client session buffer, callback the client data ready function to transfer the buffer contents to the client. Change-Id: Id79775426fa6d22e00dc2ef6f55c439eacb9b2af --- .../core/runtime/amd_gpu_agent.cpp | 54 ++++++++++++++++++- runtime/hsa-runtime/pcs/pcs_runtime.cpp | 34 ++++++++++++ runtime/hsa-runtime/pcs/pcs_runtime.h | 11 ++++ 3 files changed, 97 insertions(+), 2 deletions(-) diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index d7d10a3d72..a4c470a97d 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -2867,7 +2867,6 @@ hsa_status_t GpuAgent::PcSamplingFlushHostTrapDeviceBuffers( next_buffer = (which_buffer + 1) % 2; reset_write_val = (uint64_t)next_buffer << 63; - // HSA::hsa_signal_store_screlease(done_sig[which_buffer], 1); /* * ATOMIC_MEM, perform atomic_exchange * We use a double-buffer mechanism so that trap handlers calls are writing to one buffer while @@ -3031,7 +3030,58 @@ void GpuAgent::PcSamplingThread() { } while (true); HSA::hsa_signal_store_screlease(done_sig[which_buffer], 1); - if (PcSamplingFlushHostTrapDeviceBuffers(session) != HSA_STATUS_SUCCESS) break; + if (PcSamplingFlushHostTrapDeviceBuffers(session) != HSA_STATUS_SUCCESS) + goto thread_exit; + + size_t bytes_before_wrap; + size_t bytes_after_wrap; + + 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 >= session.buffer_size()) { + session.HandleSampleData(ht_data.host_read_ptr, session.buffer_size(), NULL, 0, 0); + ht_data.host_read_ptr += session.buffer_size(); + bytes_before_wrap = ht_data.host_buffer_wrap_pos - ht_data.host_read_ptr; + } + + if (bytes_before_wrap + bytes_after_wrap >= session.buffer_size()) { + session.HandleSampleData(ht_data.host_read_ptr, bytes_before_wrap, host_buffer_begin, + (session.buffer_size() - bytes_before_wrap), 0); + ht_data.host_read_ptr = host_buffer_begin + (session.buffer_size() - bytes_before_wrap); + bytes_before_wrap = 0; + ht_data.host_buffer_wrap_pos = 0; + bytes_after_wrap = ht_data.host_write_ptr - ht_data.host_read_ptr; + } + + while (bytes_after_wrap >= session.buffer_size()) { + session.HandleSampleData(ht_data.host_read_ptr, session.buffer_size(), NULL, 0, 0); + ht_data.host_read_ptr += session.buffer_size(); + bytes_before_wrap = 0; + 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 >= session.buffer_size()) { + assert(ht_data.host_read_ptr >= host_buffer_begin && + ht_data.host_read_ptr + session.buffer_size() < host_buffer_end); + session.HandleSampleData(ht_data.host_read_ptr, session.buffer_size(), NULL, 0, 0); + ht_data.host_read_ptr += session.buffer_size(); + bytes_before_wrap = ht_data.host_write_ptr - ht_data.host_read_ptr; + } + } } thread_exit: debug_print("PcSamplingThread::Exiting\n"); diff --git a/runtime/hsa-runtime/pcs/pcs_runtime.cpp b/runtime/hsa-runtime/pcs/pcs_runtime.cpp index b0f57ca13e..dd787d6ef5 100644 --- a/runtime/hsa-runtime/pcs/pcs_runtime.cpp +++ b/runtime/hsa-runtime/pcs/pcs_runtime.cpp @@ -159,6 +159,40 @@ void PcsRuntime::PcSamplingSession::GetHsaKmtSamplingInfo(HsaPcSamplingInfo* sam } } +hsa_status_t PcSamplingDataCopyCallback(void* _session, size_t bytes_to_copy, void* destination) { + assert(_session); + assert(destination); + + PcsRuntime::PcSamplingSession* session = + reinterpret_cast(_session); + + return session->DataCopyCallback(reinterpret_cast(destination), bytes_to_copy); +} + +hsa_status_t PcsRuntime::PcSamplingSession::DataCopyCallback(uint8_t* buffer, + size_t bytes_to_copy) { + if (bytes_to_copy != (data_rdy.buf1_sz + data_rdy.buf2_sz)) return HSA_STATUS_ERROR_EXCEPTION; + + if (data_rdy.buf1_sz) memcpy(buffer, data_rdy.buf1, data_rdy.buf1_sz); + if (data_rdy.buf2_sz) memcpy(buffer + data_rdy.buf1_sz, data_rdy.buf2, data_rdy.buf2_sz); + + return HSA_STATUS_SUCCESS; +} + +hsa_status_t PcsRuntime::PcSamplingSession::HandleSampleData(uint8_t* buf1, size_t buf1_sz, + uint8_t* buf2, size_t buf2_sz, + size_t lost_sample_count) { + data_rdy.buf1 = buf1; + data_rdy.buf1_sz = buf1_sz; + data_rdy.buf2 = buf2; + data_rdy.buf2_sz = buf2_sz; + + csd.data_ready_callback(csd.client_callback_data, buf1_sz + buf2_sz, lost_sample_count, + &PcSamplingDataCopyCallback, + /* hsa_callback_data*/ this); + return HSA_STATUS_SUCCESS; +} + hsa_status_t PcsRuntime::PcSamplingIterateConfig( core::Agent* agent, hsa_ven_amd_pcs_iterate_configuration_callback_t configuration_callback, void* callback_data) { diff --git a/runtime/hsa-runtime/pcs/pcs_runtime.h b/runtime/hsa-runtime/pcs/pcs_runtime.h index 3fbac4650e..6d8db2127f 100644 --- a/runtime/hsa-runtime/pcs/pcs_runtime.h +++ b/runtime/hsa-runtime/pcs/pcs_runtime.h @@ -84,6 +84,9 @@ class PcsRuntime { const size_t sample_size() { return sample_size_; } void GetHsaKmtSamplingInfo(HsaPcSamplingInfo* sampleInfo); + hsa_status_t HandleSampleData(uint8_t* buf1, size_t buf1_sz, uint8_t* buf2, size_t buf2_sz, + size_t lost_sample_count); + hsa_status_t DataCopyCallback(uint8_t* buffer, size_t buffer_size); core::Agent* agent; void SetThunkId(HsaPcSamplingTraceId thunkId) { thunkId_ = thunkId; } @@ -109,6 +112,14 @@ class PcsRuntime { void* client_callback_data; }; struct client_session_data_t csd; + + struct data_ready_info_t { + uint8_t* buf1; + size_t buf1_sz; + uint8_t* buf2; + size_t buf2_sz; + }; + struct data_ready_info_t data_rdy; }; // class PcSamplingSession hsa_status_t PcSamplingIterateConfig(