diff --git a/events.cpp b/events.cpp index e28491f31c..7e4f80ce4d 100644 --- a/events.cpp +++ b/events.cpp @@ -123,7 +123,7 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtWaitOnMultipleEvents_Ext(HsaEvent *Events[], HSAKMT_STATUS HSAKMTAPI hsaKmtOpenSMI(HSAuint32 NodeId, int *fd) { CHECK_DXG_OPEN(); - pr_debug("[%s] node %d\n", __func__, NodeId); + pr_debug("node id %d\n", NodeId); assert(false); return HSAKMT_STATUS_SUCCESS; } diff --git a/memory.cpp b/memory.cpp index 260d0b0a30..effdd2a9bb 100644 --- a/memory.cpp +++ b/memory.cpp @@ -258,7 +258,7 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtRegisterMemoryWithFlags( if (!MemoryAddress) return HSAKMT_STATUS_INVALID_PARAMETER; - pr_debug("[%s] address %p\n", __func__, MemoryAddress); + pr_debug("address %p\n", MemoryAddress); if (MemFlags.ui32.ExtendedCoherent && MemFlags.ui32.CoarseGrain) return HSAKMT_STATUS_INVALID_PARAMETER; @@ -297,7 +297,7 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtRegisterGraphicsHandleToNodesExt(HSAuint64 Graphic uint32_t *gpu_id_array = NULL; HSAKMT_STATUS ret = HSAKMT_STATUS_SUCCESS; - pr_debug("[%s] number of nodes %lu\n", __func__, NumberOfNodes); + pr_debug("number of nodes %lu\n", NumberOfNodes); GraphicsResourceInfo->NodeId = 1; return hsaKmtImportDMABufHandle(GraphicsResourceHandle, &GraphicsResourceInfo->MemoryAddress); @@ -412,7 +412,7 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtDeregisterMemory(void *MemoryAddress) { if (!MemoryAddress) return HSAKMT_STATUS_INVALID_PARAMETER; - pr_debug("[%s] address %p\n", __func__, MemoryAddress); + pr_debug("address %p\n", MemoryAddress); return HSAKMT_STATUS_SUCCESS; } @@ -511,7 +511,7 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtUnmapMemoryToGPU(void *MemoryAddress) { return HSAKMT_STATUS_SUCCESS; } - pr_debug("[%s] address %p\n", __func__, MemoryAddress); + pr_debug("address %p\n", MemoryAddress); wsl::thunk::GpuMemoryHandle handle = nullptr; { @@ -574,7 +574,7 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtQueryPointerInfo(const void *Pointer, if (!Pointer || !PointerInfo) return HSAKMT_STATUS_INVALID_PARAMETER; - pr_debug("[%s] pointer %p\n", __func__, Pointer); + pr_debug("pointer %p\n", Pointer); memset(PointerInfo, 0, sizeof(HsaPointerInfo)); @@ -617,7 +617,7 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtSetMemoryUserData(const void *Pointer, HSAKMT_STATUS HSAKMTAPI hsaKmtReplaceAsanHeaderPage(void *addr) { assert(false); #ifdef SANITIZER_AMDGPU - pr_debug("[%s] address %p\n", __func__, addr); + pr_debug("address %p\n", addr); CHECK_DXG_OPEN(); return HSAKMT_STATUS_SUCCESS; @@ -629,7 +629,7 @@ HSAKMT_STATUS HSAKMTAPI hsaKmtReplaceAsanHeaderPage(void *addr) { HSAKMT_STATUS HSAKMTAPI hsaKmtReturnAsanHeaderPage(void *addr) { assert(false); #ifdef SANITIZER_AMDGPU - pr_debug("[%s] address %p\n", __func__, addr); + pr_debug("address %p\n", addr); CHECK_DXG_OPEN(); return HSAKMT_STATUS_SUCCESS; diff --git a/util/lnx/os_linux.cpp b/util/lnx/os_linux.cpp index a686f0f216..368951c494 100644 --- a/util/lnx/os_linux.cpp +++ b/util/lnx/os_linux.cpp @@ -152,7 +152,7 @@ class os_thread { } err = pthread_create(&thread, &attrib, ThreadTrampoline, args.get()); if (err != EINVAL) break; - debug_print("pthread_create returned EINVAL, doubling stack size\n"); + pr_debug("pthread_create returned EINVAL, doubling stack size\n"); } } @@ -308,7 +308,7 @@ static int callback(struct dl_phdr_info* info, size_t size, void* data) { if (dyn_section[j].d_tag == DT_STRSZ) limit = dyn_section[j].d_un.d_val; } - if (strings == nullptr) debug_print("String table not found"); + if (strings == nullptr) pr_debug("String table not found\n"); /* * Hacky lookup, if string and symbol tables are found, diff --git a/wddm/device.cpp b/wddm/device.cpp index f2c284a379..17a378732b 100644 --- a/wddm/device.cpp +++ b/wddm/device.cpp @@ -301,8 +301,8 @@ bool WDDMDevice::ReserveLocalHeapSpace(void) { match_index = i; local_heap_space_start_ = local_va; - debug_print("%s success to reserve gpu va %lx and va cpu %p in %d time\n", - __FUNCTION__, local_va, ptr, i); + pr_debug("success to reserve gpu va %lx and va cpu %p in %d time\n", + local_va, ptr, i); break; } else { debug_print("%s fail to reserve gpu va for cpu va %p in %d time!\n", @@ -365,7 +365,7 @@ void WDDMDevice::SetPowerOptimization(bool restore) { d3dkmt_escape.Flags.HardwareAccess = true; NTSTATUS status = D3DKMTEscape(&d3dkmt_escape); - debug_print("%s status %d restore %d\n", __FUNCTION__, status, restore); + pr_debug("status %d, restore %d\n", status, restore); thunk_proxy::DestroyPrivData(priv_data); } @@ -596,9 +596,10 @@ bool WDDMDevice::CreateSyncobj(D3DKMT_HANDLE *handle, uint64_t **addr) { if (ret == STATUS_SUCCESS) { *handle = args.hSyncObject; *addr = (uint64_t *)args.Info.MonitoredFence.FenceValueCPUVirtualAddress; - debug_print("create syncobj cpu addr=%p gpu addr=%" PRIx64 "\n", - args.Info.MonitoredFence.FenceValueCPUVirtualAddress, - args.Info.MonitoredFence.FenceValueGPUVirtualAddress); + pr_debug("create syncobj cpu addr=%p gpu addr=%" PRIx64 "\n", + args.Info.MonitoredFence.FenceValueCPUVirtualAddress, + args.Info.MonitoredFence.FenceValueGPUVirtualAddress); + return true; } @@ -730,7 +731,7 @@ void WDDMDevice::GetClockCounters(uint64_t *gpu, uint64_t *cpu) { NTSTATUS status = D3DKMTEscape(&d3dkmt_escape); if (status) { - debug_print("%s status %d \n", __FUNCTION__, status); + pr_debug("status %d \n", status); } else { thunk_proxy::QueryCalibratedTimestamps(priv_data, gpu, cpu); } diff --git a/wddm/gpu_memory.cpp b/wddm/gpu_memory.cpp index 41ab2aab51..795ced34c7 100644 --- a/wddm/gpu_memory.cpp +++ b/wddm/gpu_memory.cpp @@ -391,13 +391,13 @@ ErrorCode GpuMemory::ImportPhysicalHandle(int dmabuf_fd) { debug_print("%s query resource info from nt handle failed %d\n", __FUNCTION__, static_cast(ret)); return ErrorCode::InvalidateParams; } - debug_print("wsl-thunk: import from nt handle %d, get allocation number %d," - " runtime data size %#x total driver data size %#x resource data size=%#x\n", - dmabuf_fd, - query_args.NumAllocations, - query_args.PrivateRuntimeDataSize, - query_args.TotalPrivateDriverDataSize, - query_args.ResourcePrivateDriverDataSize); + pr_debug("wsl-thunk: import from nt handle %d, get allocation number %d," + " runtime data size %#x total driver data size %#x resource data size=%#x\n", + dmabuf_fd, + query_args.NumAllocations, + query_args.PrivateRuntimeDataSize, + query_args.TotalPrivateDriverDataSize, + query_args.ResourcePrivateDriverDataSize); SharedHandleInfo shared_info; if(sizeof(shared_info) != query_args.PrivateRuntimeDataSize) { diff --git a/wddm/queue.cpp b/wddm/queue.cpp index ab6ebe4c06..3294f75943 100644 --- a/wddm/queue.cpp +++ b/wddm/queue.cpp @@ -135,7 +135,7 @@ hsa_status_t WDDMQueue::SetPriority(hsa_amd_queue_priority_t priority) { if (prio == new_prio) return HSA_STATUS_SUCCESS; - debug_print("set prio %d -> %d\n", prio, new_prio); + pr_debug("set prio %d -> %d\n", prio, new_prio); device->DestroyHwQueue(this); prio = new_prio; @@ -219,13 +219,13 @@ void ComputeQueue::AqlToPm4Thread(ComputeQueue *queue) { (sleep && queue->IsInvalidPacket())) { if (queue->thread_stop_) break; - debug_print("wait %p wptr=%" PRIx64 " rptr=%" PRIx64 "\n", - queue->ring, queue->GetRingWptr()->load(), queue->GetRingRptr()->load()); + pr_debug("wait %p wptr=%" PRIx64 " rptr=%" PRIx64 "\n", + queue->ring, queue->GetRingWptr()->load(), queue->GetRingRptr()->load()); queue->thread_cond_.wait(lock); } } - debug_print("aql to pm4 thread %p exit\n", queue->ring); + pr_debug("aql to pm4 thread %p exit\n", queue->ring); } ComputeQueue::ComputeQueue(WDDMDevice *device, @@ -447,8 +447,8 @@ bool ComputeQueue::UpdateScratch(uint32_t private_segment_size, bool wave32) { if (scratch_size_ >= scratch_size) return true; - debug_print("need realloc scratch buffer, size %x -> %x\n", - scratch_size_, scratch_size); + pr_debug("need realloc scratch buffer, size %x -> %x\n", + scratch_size_, scratch_size); GpuMemoryCreateInfo create_info{}; create_info.size = scratch_size; @@ -537,8 +537,8 @@ uint64_t ComputeQueue::GetKernelObjAddr(uint64_t addr) const { void ComputeQueue::RingDoorbell() { thread_cond_lock_.lock(); thread_cond_lock_.unlock(); - debug_print("notify %p wptr=%" PRIx64 " rptr=%" PRIx64 "\n", - ring, GetRingWptr()->load(), GetRingRptr()->load()); + pr_debug("notify %p wptr=%" PRIx64 " rptr=%" PRIx64 "\n", + ring, GetRingWptr()->load(), GetRingRptr()->load()); thread_cond_.notify_one(); } @@ -599,7 +599,7 @@ hsa_status_t ComputeQueue::Submit(void) { hsa_status_t ComputeQueue::KernelDispatchAqlToPm4(char *cpu, hsa_kernel_dispatch_packet_t *packet) { - debug_print("queue %p kernel dispatch head=%x setup=%x wx=%x wy=%x wz=%x " + pr_debug("queue %p kernel dispatch head=%x setup=%x wx=%x wy=%x wz=%x " "gx=%x gy=%x gz=%x ps=%x gs=%x ko=%" PRIx64 " ka=%p cs=%" PRIx64 "\n", ring, packet->header, packet->setup, packet->workgroup_size_x, packet->workgroup_size_y, @@ -625,10 +625,10 @@ ComputeQueue::KernelDispatchAqlToPm4(char *cpu, hsa_kernel_dispatch_packet_t *pa void* entry = (void*)(packet->kernel_object + kernel_object->kernel_code_entry_byte_offset); assert((size_t)entry % AMD_ISA_ALIGN_BYTES == 0); - debug_print("kernel object property=%x entry=%p lds=%x+%x\n", - kernel_object->kernel_code_properties, entry, - kernel_object->workgroup_group_segment_byte_size, - packet->group_segment_size); + pr_debug("kernel object property=%x entry=%p lds=%x+%x\n", + kernel_object->kernel_code_properties, entry, + kernel_object->workgroup_group_segment_byte_size, + packet->group_segment_size); if (packet->setup == 0 || packet->setup > 3) return HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS; @@ -704,7 +704,7 @@ ComputeQueue::KernelDispatchAqlToPm4(char *cpu, hsa_kernel_dispatch_packet_t *pa assert(signal->kind == AMD_SIGNAL_KIND_USER); uint64_t *signal_addr = (uint64_t *)&signal->value; - debug_print("signal value=%" PRIx64 "\n", signal->value); + pr_debug("signal value=%" PRIx64 "\n", signal->value); if (platform_atomic_support_) i += cmd_util.BuildAtomicMem(signal_addr, TC_OP_ATOMIC_ADD_RTN_64, cpu + i, cache_policy__mec_atomic_mem__bypass, -1); @@ -729,7 +729,7 @@ ComputeQueue::KernelDispatchAqlToPm4(char *cpu, hsa_kernel_dispatch_packet_t *pa hsa_status_t ComputeQueue::BarrierGenericAqlToPm4(char *cpu, hsa_barrier_and_packet_t *packet, bool is_or) { - debug_print("queue %p %s head=%x dep %" PRIx64 " %" PRIx64 " %" PRIx64 + pr_debug("queue %p %s head=%x dep %" PRIx64 " %" PRIx64 " %" PRIx64 " %" PRIx64 " %" PRIx64 " cs=%" PRIx64"\n", ring, is_or ? "or" : "and", packet->header, packet->dep_signal[0].handle, @@ -776,7 +776,7 @@ ComputeQueue::BarrierGenericAqlToPm4(char *cpu, hsa_barrier_and_packet_t *packet amd_signal_t *signal = (amd_signal_t *)packet->completion_signal.handle; assert(signal->kind == AMD_SIGNAL_KIND_USER); uint64_t *signal_addr = (uint64_t *)&signal->value; - debug_print("signal value=%" PRIx64 "\n", signal->value); + pr_debug("signal value=%" PRIx64 "\n", signal->value); // Record start timestamp when enabling profiling if (EnableProfiling()) @@ -823,11 +823,11 @@ hsa_status_t ComputeQueue::VendorSpecificAqlToPm4(char *cpu, amd_aql_pm4_ib *pac assert(op == IT_INDIRECT_BUFFER); uint32_t* pm4_addr = reinterpret_cast((static_cast(packet->ib_jump_cmd[2]) << 32) | (static_cast(packet->ib_jump_cmd[1]) & ~3ull)); uint32_t pm4_size = packet->ib_jump_cmd[3]&0xfffff; - debug_print("queue %p %s VENDOR_SPECIFIC pkt pm4_addr %p pm4_size %#x cs=%" PRIx64"\n", + pr_debug("queue %p %s VENDOR_SPECIFIC pkt pm4_addr %p pm4_size %#x cs=%" PRIx64"\n", ring, vendor_packet_process ? "process" : "skip", pm4_addr, pm4_size, packet->completion_signal.handle); for (int i = 0; i < pm4_size; i++) { - debug_print("pm4_addr[%d]=%#x\n", i, pm4_addr[i]); + pr_debug("pm4_addr[%d]=%#x\n", i, pm4_addr[i]); } int i = ib_size; @@ -841,7 +841,7 @@ hsa_status_t ComputeQueue::VendorSpecificAqlToPm4(char *cpu, amd_aql_pm4_ib *pac amd_signal_t *signal = (amd_signal_t *)packet->completion_signal.handle; assert(signal->kind == AMD_SIGNAL_KIND_USER); uint64_t *signal_addr = (uint64_t *)&signal->value; - debug_print("signal value=%" PRIx64 "\n", signal->value); + pr_debug("signal value=%" PRIx64 "\n", signal->value); // Record start timestamp when enabling profiling if (EnableProfiling()) @@ -933,8 +933,8 @@ hsa_status_t ComputeQueue::Process(void) { while (cmdbuf_aql_frame_write_index < ring_wptr->load() && !IsInvalidPacket()) { - debug_print("process %p wptr=%" PRIx64 " rptr=%" PRIx64 "\n", - ring, ring_wptr->load(), ring_rptr->load()); + pr_debug("process %p wptr=%" PRIx64 " rptr=%" PRIx64 "\n", + ring, ring_wptr->load(), ring_rptr->load()); hsa_status_t ret; @@ -970,9 +970,8 @@ hsa_status_t ComputeQueue::Process(void) { ready_to_submit = false; - debug_print("done %p wptr=%" PRIx64 " rptr=%" PRIx64 "\n", - ring, ring_wptr->load(), ring_rptr->load()); - + pr_debug("done %p wptr=%" PRIx64 " rptr=%" PRIx64 "\n", + ring, ring_wptr->load(), ring_rptr->load()); } return HSA_STATUS_SUCCESS; @@ -994,7 +993,7 @@ void SDMAQueue::SdmaThread(SDMAQueue *queue) { uint64_t end = queue->wptr_queue_.front().second; queue->wptr_queue_.pop(); lock.unlock(); - debug_print("SDMA: wptr %lx %lx\n", start, end); + pr_debug("wptr %lx %lx\n", start, end); SDMA_PKT_POLL_REGMEM* poll_pkt = reinterpret_cast(queue->cmdbuf_addr + queue->WrapIntoRocrRing(start)); SDMA_PKT_POLL_REGMEM* poll_next_pkt = poll_pkt + 1; @@ -1019,7 +1018,7 @@ void SDMAQueue::SdmaThread(SDMAQueue *queue) { amd_signal_t* signal = (amd_signal_t*)((char*)poll_addr - offsetof(amd_signal_t, value)); uint64_t signal_handle = reinterpret_cast(signal); - debug_print("SDMA: poll signal %#lx addr %#lx val %ld\n", signal_handle, poll_addr, poll_val); + pr_debug("poll signal %#lx addr %#lx val %ld\n", signal_handle, poll_addr, poll_val); hsa_signal_t hsa_signal = {signal_handle}; hsa_signal_value_t value = hsakmt_hsa_signal_wait_relaxed(hsa_signal, HSA_SIGNAL_CONDITION_EQ, poll_val, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); @@ -1049,7 +1048,7 @@ void SDMAQueue::SdmaThread(SDMAQueue *queue) { start_time = std::chrono::steady_clock::now(); } } - debug_print("sdma thread exit\n"); + pr_debug("thread exit\n"); } SDMAQueue::SDMAQueue(WDDMDevice *device, @@ -1081,7 +1080,7 @@ SDMAQueue::~SDMAQueue() { } void SDMAQueue::RingDoorbell() { - debug_print("SDMA: ringdoorbell %#lx %#lx\n", wptr_pre_, wptr_next_); + pr_debug("ringdoorbell %#lx %#lx\n", wptr_pre_, wptr_next_); { std::lock_guard lock(wptr_queue_lock_);