diff --git a/rocclr/device/rocm/rocblit.cpp b/rocclr/device/rocm/rocblit.cpp index c89da0cf3b..d3b295e49e 100644 --- a/rocclr/device/rocm/rocblit.cpp +++ b/rocclr/device/rocm/rocblit.cpp @@ -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
(dst) + dstOffset), dstAgent, (reinterpret_cast(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) { diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 95d92d0b8e..f616c2fd02 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -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(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
(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
(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(srcArgPtr)); } else if (desc.type_ == T_SAMPLER) { @@ -816,12 +814,12 @@ bool VirtualGPU::dispatchGenericAqlPacket( packet_store_release(reinterpret_cast(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) { diff --git a/rocclr/utils/debug.cpp b/rocclr/utils/debug.cpp index 6a43a0d942..8127128030 100644 --- a/rocclr/utils/debug.cpp +++ b/rocclr/utils/debug.cpp @@ -29,6 +29,8 @@ #include #include #include +#include +#include #ifdef _WIN32 #include @@ -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) {