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

SWDEV-301667 - Better log

Browse Source 
- Print SWq for AQL packets, this helps correlating a stream to the HWq
mapped

Change-Id: I610430c0872a1abc6636027c00163ec46983cd65


[ROCm/clr commit: 984c86f407]
This commit is contained in:
Saleel Kudchadker
2024-02-29 23:37:12 +00:00
parent 286599824e
commit 4ff5ec0a02
2 changed files with 9 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+6 -6
View File
@@ -897,12 +897,12 @@ bool VirtualGPU::dispatchGenericAqlPacket(
packet_store_release(reinterpret_cast<uint32_t*>(aql_loc), header, rest);
}
ClPrint(amd::LOG_DEBUG, amd::LOG_AQL,
"HWq=0x%zx id=%d, Dispatch Header = "
"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",
gpu_queue_->base_address, gpu_queue_->id, header,
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),
extractAqlBits(header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
@@ -1071,10 +1071,10 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal,
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
ClPrint(amd::LOG_DEBUG, amd::LOG_AQL,
"HWq=0x%zx, BarrierAND Header = 0x%x (type=%d, barrier=%d, acquire=%d,"
"SWq=0x%zx, HWq=0x%zx, id=%d, BarrierAND Header = 0x%x (type=%d, barrier=%d, acquire=%d,"
" release=%d), "
"dep_signal=[0x%zx, 0x%zx, 0x%zx, 0x%zx, 0x%zx], completion_signal=0x%zx",
gpu_queue_->base_address, packetHeader,
gpu_queue_, gpu_queue_->base_address, gpu_queue_->id, packetHeader,
extractAqlBits(packetHeader, HSA_PACKET_HEADER_TYPE,
HSA_PACKET_HEADER_WIDTH_TYPE),
extractAqlBits(packetHeader, HSA_PACKET_HEADER_BARRIER,
@@ -1150,10 +1150,10 @@ void VirtualGPU::dispatchBarrierValuePacket(uint16_t packetHeader, bool resolveD
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
ClPrint(amd::LOG_DEBUG, amd::LOG_AQL,
"HWq=0x%zx, BarrierValue Header = 0x%x AmdFormat = 0x%x "
"SWq=0x%zx, HWq=0x%zx, id=%d, BarrierValue Header = 0x%x AmdFormat = 0x%x "
"(type=%d, barrier=%d, acquire=%d, release=%d), "
"signal=0x%zx, value = 0x%llx mask = 0x%llx cond: %s, completion_signal=0x%zx",
gpu_queue_->base_address, packetHeader, rest,
gpu_queue_, gpu_queue_->base_address, gpu_queue_->id, packetHeader, rest,
extractAqlBits(packetHeader, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
extractAqlBits(packetHeader, HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER),
extractAqlBits(packetHeader, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,