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


[ROCm/ROCR-Runtime commit: 5177d17f5d]

projects/rocr-runtime/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");

projects/rocr-runtime/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<PcsRuntime::PcSamplingSession*>(_session);
+
+  return session->DataCopyCallback(reinterpret_cast<uint8_t*>(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) {

projects/rocr-runtime/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(