From 0567c3b720b08851b4d209f617ec8a7cf31a56d7 Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Wed, 31 Jan 2024 05:56:51 +0000 Subject: [PATCH] SWDEV-301667 - Better log Display queue base pointer in the log. This can be co-related with AQL packets Change-Id: I544f9b6db6ae01c85e57e4b3f0b3fffefcd7c2ed --- rocclr/device/rocm/rocdevice.cpp | 35 ++++++++++++++++--------------- rocclr/device/rocm/rocvirtual.cpp | 2 +- 2 files changed, 19 insertions(+), 18 deletions(-) diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index 8758c62913..15478fe5ad 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -273,11 +273,12 @@ Device::~Device() { auto& qInfo = qIter->second; if (qInfo.hostcallBuffer_) { ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hostcall buffer %p for hardware queue %p", - qInfo.hostcallBuffer_, qIter->first); + qInfo.hostcallBuffer_, qIter->first->base_address); disableHostcalls(qInfo.hostcallBuffer_); context().svmFree(qInfo.hostcallBuffer_); } - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hardware queue %p with refCount 0", queue); + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hardware queue %p with refCount 0", + queue->base_address); qIter = it.erase(qIter); hsa_queue_destroy(queue); } @@ -2350,7 +2351,7 @@ void Device::updateFreeMemory(size_t size, bool free) { } freeMem_ -= size; } - ClPrint(amd::LOG_INFO, amd::LOG_MEM, "device=0x%lx, freeMem_ = 0x%zx", this, freeMem_.load()); + ClPrint(amd::LOG_INFO, amd::LOG_MEM, "Device=0x%lx, freeMem_ = 0x%zx", this, freeMem_.load()); } // ================================================================================================ @@ -2895,8 +2896,8 @@ hsa_queue_t* Device::getQueueFromPool(const uint qIndex) { for (auto it = queuePool_[qIndex].begin(); it != queuePool_[qIndex].end(); it++) { if (it->second.refCount == 0) { it->second.refCount++; - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "selected queue refCount: %p (%d)", it->first, - it->second.refCount); + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Selected queue refCount: %p (%d)", + it->first->base_address, it->second.refCount); return it->first; } } @@ -2907,8 +2908,8 @@ hsa_queue_t* Device::getQueueFromPool(const uint qIndex) { queuePool_[qIndex].begin(), queuePool_[qIndex].end(), [](PoolRef A, PoolRef B) { return A.second.refCount < B.second.refCount; }); lowest->second.refCount++; - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "selected queue refCount: %p (%d)", lowest->first, - lowest->second.refCount); + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Selected queue refCount: %p (%d)", + lowest->first->base_address, lowest->second.refCount); return lowest->first; } } @@ -2922,7 +2923,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, queuePool_[QueuePriority::Normal].size() <= GPU_MAX_HW_QUEUES || queuePool_[QueuePriority::High].size() <= GPU_MAX_HW_QUEUES); - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "number of allocated hardware queues with low priority: %d," + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Number of allocated hardware queues with low priority: %d," " with normal priority: %d, with high priority: %d, maximum per priority is: %d", queuePool_[QueuePriority::Low].size(), queuePool_[QueuePriority::Normal].size(), @@ -3000,8 +3001,8 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, } } - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "created hardware queue %p with size %d with priority %d," - " cooperative: %i", queue, queue_size, queue_priority, coop_queue); + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Created hardware queue %p with size %d with priority %d," + " cooperative: %i", queue->base_address, queue_size, queue_priority, coop_queue); hsa_amd_profiling_set_profiler_enabled(queue, 1); if (cuMask.size() != 0 || info_.globalCUMask_.size() != 0) { @@ -3038,7 +3039,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, ss << mask[i]; } ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Setting CU mask 0x%s for hardware queue %p", - ss.str().c_str(), queue); + ss.str().c_str(), queue->base_address); std::vector final_mask = {}; // hsa_amd_queue_cu_set_mask expects each bit in cuMask to represent each CU @@ -3092,8 +3093,8 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, assert(result.second && "QueueInfo already exists"); auto &qInfo = result.first->second; qInfo.refCount = 1; - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "acquireQueue refCount: %p (%d)", result.first->first, - result.first->second.refCount); + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "acquireQueue refCount: %p (%d)", + result.first->first->base_address, result.first->second.refCount); return queue; } @@ -3104,8 +3105,8 @@ void Device::releaseQueue(hsa_queue_t* queue, const std::vector& cuMas auto &qInfo = qIter->second; assert(qInfo.refCount > 0); qInfo.refCount--; - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "releaseQueue refCount:%p (%d)", qIter->first, - qIter->second.refCount); + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "releaseQueue refCount:%p (%d)", + qIter->first->base_address, qIter->second.refCount); } } } @@ -3147,11 +3148,11 @@ void* Device::getOrCreateHostcallBuffer(hsa_queue_t* queue, bool coop_queue, void* buffer = context().svmAlloc(size, align, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS); if (!buffer) { ClPrint(amd::LOG_ERROR, amd::LOG_QUEUE, - "Failed to create hostcall buffer for hardware queue %p", queue); + "Failed to create hostcall buffer for hardware queue %p", queue->base_address); return nullptr; } ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Created hostcall buffer %p for hardware queue %p", buffer, - queue); + queue->base_address); if (!coop_queue) { qIter->second.hostcallBuffer_ = buffer; } else { diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 5abbb6c280..5a504feb09 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -1152,7 +1152,7 @@ void VirtualGPU::dispatchBarrierValuePacket(uint16_t packetHeader, bool resolveD "HWq=0x%zx, 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_, packetHeader, rest, + gpu_queue_->base_address, 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,