AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

SWDEV-301667 - Better log

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

Change-Id: Ib8db7a00e4e7c98f570e71029e61d86f5dccc2ed
This commit is contained in:
Saleel Kudchadker
2024-05-10 22:07:14 +00:00
committed by Maneesh Gupta
orang tua 1ba74c3ce3
melakukan 72d23a02c5
2 mengubah file dengan 7 tambahan dan 4 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_graph_internal.cpp
+1 -1
Melihat File
@@ -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);
rocclr/device/rocm/rocvirtual.cpp
+6 -3
Melihat File
@@ -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);