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

SWDEV-300655 - Added thread ID to hip trace

Browse Source 
Change-Id: I9234d4ec93e7687cd0a5d1bd930bd4f80936311b


[ROCm/clr commit: 42d33029dc]
This commit is contained in:
Sarbojit Sarkar
2021-09-01 11:56:28 +00:00
committed by Sarbojit Sarkar
parent d93c8fc413
commit 45953e81dd
3 changed files with 34 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/rocm/rocblit.cpp
+8 -13
View File
@@ -456,9 +456,8 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
// Copy memory line by line
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
"[%zx]!\t HSA Asycn Copy Rect wait_event=0x%zx, completion_signal=0x%zx",
std::this_thread::get_id(), (wait_events.size() != 0) ? wait_events[0].handle : 0,
active.handle);
"HSA Asycn Copy Rect wait_event=0x%zx, completion_signal=0x%zx",
(wait_events.size() != 0) ? wait_events[0].handle : 0, active.handle);
hsa_status_t status = hsa_amd_memory_async_copy_rect(&dstMem, &offset,
&srcMem, &offset, &dim, agent, direction, wait_events.size(), &wait_events[0], active);
if (status != HSA_STATUS_SUCCESS) {
@@ -478,9 +477,8 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
// Copy memory line by line
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
"[%zx]!\t HSA Asycn Copy wait_event=0x%zx, completion_signal=0x%zx",
std::this_thread::get_id(), (wait_events.size() != 0) ? wait_events[0].handle : 0,
active.handle);
"HSA Asycn Copy wait_event=0x%zx, completion_signal=0x%zx",
(wait_events.size() != 0) ? wait_events[0].handle : 0, active.handle);
hsa_status_t status = hsa_amd_memory_async_copy(
(reinterpret_cast<address>(dst) + dstOffset), dstAgent,
(reinterpret_cast<const_address>(src) + srcOffset), srcAgent,
@@ -666,9 +664,8 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
// Use SDMA to transfer the data
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
"[%zx]!\t HSA Asycn Copy wait_event=0x%zx, completion_signal=0x%zx",
std::this_thread::get_id(), (wait_events.size() != 0) ? wait_events[0].handle : 0,
active.handle);
"HSA Asycn Copy wait_event=0x%zx, completion_signal=0x%zx",
(wait_events.size() != 0) ? wait_events[0].handle : 0, active.handle);
status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent,
size[0], wait_events.size(), &wait_events[0], active);
@@ -724,8 +721,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
memcpy(hsaBuffer, hostSrc + offset, size);
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
"[%zx]!\t HSA Async Copy completion_signal=0x%zx",
std::this_thread::get_id(), active.handle);
"HSA Async Copy completion_signal=0x%zx", active.handle);
status = hsa_amd_memory_async_copy(hostDst + offset, dev().getBackendDevice(), hsaBuffer,
srcAgent, size, 0, nullptr, active);
if (status != HSA_STATUS_SUCCESS) {
@@ -754,8 +750,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
// Copy data from Device to Host
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
"[%zx]!\t HSA Async Copy completion_signal=0x%zx",
std::this_thread::get_id(), active.handle);
"HSA Async Copy completion_signal=0x%zx", active.handle);
status = hsa_amd_memory_async_copy(hsaBuffer, dstAgent, hostSrc + offset,
dev().getBackendDevice(), size, 0, nullptr, active);
if (status == HSA_STATUS_SUCCESS) {
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+13 -16
View File
@@ -509,8 +509,8 @@ bool VirtualGPU::HwQueueTracker::CpuWaitForSignal(ProfilingSignal* signal) {
signal->ts_ = nullptr;
} else if (!signal->done_) {
amd::ScopedLock lock(signal->LockSignalOps());
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "[%zx]!\t Host wait on completion_signal=0x%zx",
std::this_thread::get_id(), signal->signal_.handle);
ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Host wait on completion_signal=0x%zx",
signal->signal_.handle);
if (!WaitForSignal(signal->signal_, gpu_.ActiveWait())) {
LogPrintfError("Failed signal [0x%lx] wait", signal->signal_);
return false;
@@ -663,9 +663,8 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para
const void* globalAddress = *reinterpret_cast<const void* const*>(params + desc.offset_);
ClPrint(amd::LOG_INFO, amd::LOG_KERN,
"[%zx]!\tArg%d: %s %s = ptr:%p obj:[%p-%p]",
std::this_thread::get_id(), i, desc.typeName_.c_str(), desc.name_.c_str(),
globalAddress, gpuMem->getDeviceMemory(),
"Arg%d: %s %s = ptr:%p obj:[%p-%p]", i, desc.typeName_.c_str(),
desc.name_.c_str(), globalAddress, gpuMem->getDeviceMemory(),
reinterpret_cast<address>(gpuMem->getDeviceMemory()) + mem->getSize());
// Validate memory for a dependency in the queue
@@ -743,8 +742,7 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para
WriteAqlArgAt(const_cast<address>(params), &mem, sizeof(void*), it->second);
}
ClPrint(amd::LOG_INFO, amd::LOG_KERN,
"[%zx]!\tArg%d: %s %s = val:%lld",
std::this_thread::get_id(), i, desc.typeName_.c_str(), desc.name_.c_str(),
"Arg%d: %s %s = val:%lld", i, desc.typeName_.c_str(), desc.name_.c_str(),
*reinterpret_cast<const long long*>(srcArgPtr));
}
else if (desc.type_ == T_SAMPLER) {
@@ -816,12 +814,12 @@ bool VirtualGPU::dispatchGenericAqlPacket(
packet_store_release(reinterpret_cast<uint32_t*>(aql_loc), header, rest);
}
ClPrint(amd::LOG_DEBUG, amd::LOG_AQL,
"[%zx] HWq=0x%zx, Dispatch Header = "
"HWq=0x%zx, 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",
std::this_thread::get_id(), gpu_queue_,
header, extractAqlBits(header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
gpu_queue_, 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,
@@ -953,10 +951,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,
"[%zx] HWq=0x%zx, BarrierAND Header = 0x%x (type=%d, barrier=%d, acquire=%d,"
"HWq=0x%zx, 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",
std::this_thread::get_id(), gpu_queue_, packetHeader,
gpu_queue_, packetHeader,
extractAqlBits(packetHeader, HSA_PACKET_HEADER_TYPE,
HSA_PACKET_HEADER_WIDTH_TYPE),
extractAqlBits(packetHeader, HSA_PACKET_HEADER_BARRIER,
@@ -2239,10 +2237,10 @@ void VirtualGPU::dispatchBarrierValuePacket(const hsa_amd_barrier_value_packet_t
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
ClPrint(amd::LOG_DEBUG, amd::LOG_AQL,
"[%zx] HWq=0x%zx, BarrierValue Header = 0x%x AmdFormat = 0x%x ",
"HWq=0x%zx, BarrierValue Header = 0x%x AmdFormat = 0x%x ",
"(type=%d, barrier=%d, acquire=%d, release=%d), "
"completion_signal=0x%zx value = 0x%llx mask = 0x%llx cond: %d (GTE: %d EQ: %d NE: %d)",
std::this_thread::get_id(), gpu_queue_, header.header, header.AmdFormat,
gpu_queue_, header.header, header.AmdFormat,
extractAqlBits(header.header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
extractAqlBits(header.header, HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER),
extractAqlBits(header.header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE,
@@ -2655,8 +2653,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
return false;
}
ClPrint(amd::LOG_INFO, amd::LOG_KERN, "[%zx]!\tShaderName : %s",
std::this_thread::get_id(), gpuKernel.name().c_str());
ClPrint(amd::LOG_INFO, amd::LOG_KERN, "ShaderName : %s", gpuKernel.name().c_str());
// Check if runtime has to setup hidden arguments
for (uint32_t i = signature.numParameters(); i < signature.numParametersAll(); ++i) {
projects/clr/rocclr/utils/debug.cpp
+13 -6
View File
@@ -29,6 +29,8 @@
#include <cstdlib>
#include <cstdio>
#include <cstdarg>
#include <thread>
#include <sstream>
#ifdef _WIN32
#include <windows.h>
@@ -96,13 +98,15 @@ void log_timestamped(LogLevel level, const char* file, int line, const char* mes
// ================================================================================================
void log_printf(LogLevel level, const char* file, int line, const char* format, ...) {
va_list ap;
std::stringstream str_thrd_id;
str_thrd_id << std::hex << std::this_thread::get_id();
va_start(ap, format);
char message[4096];
vsnprintf(message, sizeof(message), format, ap);
va_end(ap);
uint64_t timeUs = Os::timeNanos() / 1000ULL;
fprintf(outFile, ":%d:%-25s:%-4d: %010lld us: %s\n", level, file, line, timeUs/1ULL, message);
fprintf(outFile, ":%d:%-25s:%-4d: %010lld us: %-5d: [tid:0x%s] %s\n", level, file, line,
timeUs/1ULL, getpid(), str_thrd_id.str().c_str(), message);
fflush(outFile);
}
@@ -110,17 +114,20 @@ void log_printf(LogLevel level, const char* file, int line, const char* format,
void log_printf(LogLevel level, const char* file, int line, uint64_t* start,
const char* format, ...) {
va_list ap;
std::stringstream str_thrd_id;
str_thrd_id << std::hex << std::this_thread::get_id();
va_start(ap, format);
char message[4096];
vsnprintf(message, sizeof(message), format, ap);
va_end(ap);
uint64_t timeUs = Os::timeNanos() / 1000ULL;
if (start == 0 || *start == 0) {
fprintf(outFile, ":%d:%-25s:%-4d: %010lld us: %s\n", level, file, line, timeUs/1ULL, message);
fprintf(outFile, ":%d:%-25s:%-4d: %010lld us: %-5d: [tid:0x%s] %s\n", level, file, line,
timeUs/1ULL, getpid(), str_thrd_id.str().c_str(), message);
} else {
fprintf(outFile, ":%d:%-25s:%-4d: %010lld us: %s: duration: %lld us\n", level, file, line,
timeUs/1ULL, message, (timeUs - *start)/1ULL);
fprintf(outFile, ":%d:%-25s:%-4d: %010lld us: %-5d: [tid:0x%s] %s: duration: %lld us\n",
level, file, line, timeUs/1ULL, getpid(), str_thrd_id.str().c_str(), message,
(timeUs - *start)/1ULL);
}
fflush(outFile);
if (*start == 0) {