AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

PC Sampling: Push data to PC Sampling client

Browse Source 
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
This commit is contained in:
David Yat Sin
2023-09-05 17:05:10 +00:00
parent 855e454671
commit 5177d17f5d
3 changed files with 97 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+52 -2
View File
@@ -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");