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<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) {
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(