From 72d23a02c541fa65d7676beb9fb11e6c600f194a Mon Sep 17 00:00:00 2001
From: Saleel Kudchadker <Saleel.Kudchadker@amd.com>
Date: Fri, 10 May 2024 22:07:14 +0000
Subject: [PATCH] SWDEV-301667 - Better log

- Print kernelname for graph launches, its hard to correlate packets
otherwise
- Print correlation_id if any

Change-Id: Ib8db7a00e4e7c98f570e71029e61d86f5dccc2ed
---
 hipamd/src/hip_graph_internal.cpp | 2 +-
 rocclr/device/rocm/rocvirtual.cpp | 9 ++++++---
 2 files changed, 7 insertions(+), 4 deletions(-)

diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp
index e31146aebd..fb20b553eb 100644
--- a/hipamd/src/hip_graph_internal.cpp
+++ b/hipamd/src/hip_graph_internal.cpp
@@ -589,8 +589,8 @@ hipError_t GraphExec::Run(hipStream_t stream) {
           topoOrder_[i]->GetType() == hipGraphNodeTypeKernel &&
           !reinterpret_cast<hip::GraphKernelNode*>(topoOrder_[i])->HasHiddenHeap()) {
         if (topoOrder_[i]->GetEnabled()) {
-          hip_stream->vdev()->dispatchAqlPacket(topoOrder_[i]->GetAqlPacket(), accumulate);
           accumulate->addKernelName(topoOrder_[i]->GetKernelName());
+          hip_stream->vdev()->dispatchAqlPacket(topoOrder_[i]->GetAqlPacket(), accumulate);
         }
       } else {
         topoOrder_[i]->SetStream(hip_stream, this);
diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp
index 13ffe9b35b..24d8e7c4ed 100644
--- a/rocclr/device/rocm/rocvirtual.cpp
+++ b/rocclr/device/rocm/rocvirtual.cpp
@@ -906,8 +906,8 @@ bool VirtualGPU::dispatchGenericAqlPacket(
           "SWq=0x%zx, HWq=0x%zx, id=%d, Dispatch Header = "
           "0x%x (type=%d, barrier=%d, acquire=%d, release=%d), "
           "setup=%d, grid=[%zu, %zu, %zu], workgroup=[%zu, %zu, %zu], private_seg_size=%zu, "
-          "group_seg_size=%zu, kernel_obj=0x%zx, kernarg_address=0x%zx, completion_signal=0x%zx "
-          "rptr=%u, wptr=%u",
+          "group_seg_size=%zu, kernel_obj=0x%zx, kernarg_address=0x%zx, completion_signal=0x%zx, "
+          "correlation_id=%zu, rptr=%u, wptr=%u",
           gpu_queue_, gpu_queue_->base_address, gpu_queue_->id, header,
           extractAqlBits(header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
           extractAqlBits(header, HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER),
@@ -925,7 +925,8 @@ bool VirtualGPU::dispatchGenericAqlPacket(
           reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->group_segment_size,
           reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->kernel_object,
           reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->kernarg_address,
-          reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->completion_signal, read,
+          reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->completion_signal,
+          reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->reserved2, read,
           index);
 
   hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
@@ -989,6 +990,8 @@ inline bool VirtualGPU::dispatchAqlPacket(uint8_t* aqlpacket, amd::AccumulateCom
   }
   dispatchBlockingWait();
   auto packet = reinterpret_cast<hsa_kernel_dispatch_packet_t*>(aqlpacket);
+  ClPrint(amd::LOG_INFO, amd::LOG_KERN, "Graph shader name : %s",
+          vcmd->getKernelNames().back().c_str());
   dispatchGenericAqlPacket(packet, packet->header, packet->setup, false);
   if (vcmd != nullptr) {
     profilingEnd(*vcmd);