From 6f4f575ac700b7dc45113a13830aeae05ce9cf6e Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 8 Oct 2019 10:27:38 -0500 Subject: [PATCH] P4 to Git Change 2009907 by ssahasra@ssahasra-hip-vdi on 2019/10/08 11:19:25 SWDEV-204782 - store extra information per HSA queue The new struct QueueInfo is used to store metadata about each HSA queue. For hostcall, this structure will eventually contain a pointer to the hostcall buffer allocated to each HSA queue. Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#135 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.hpp#41 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#88 edit [ROCm/clr commit: 1820fe21cfcdc967c1b81f72a81e77909892465c] --- .../rocclr/runtime/device/rocm/rocdevice.cpp | 64 +++++++++++++++++++ .../rocclr/runtime/device/rocm/rocdevice.hpp | 11 +++- .../rocclr/runtime/device/rocm/rocvirtual.cpp | 42 ++---------- 3 files changed, 78 insertions(+), 39 deletions(-) diff --git a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp index 91a5eeced2..ced55b2c01 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocdevice.cpp @@ -1821,5 +1821,69 @@ bool Device::SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeI bool result = true; return result; } + +hsa_queue_t *Device::acquireQueue(uint32_t queue_size_hint) { + assert(queuePool_.size() <= GPU_MAX_HW_QUEUES); + LogPrintfInfo("number of allocated hardware queues: %d, maximum: %d", + queuePool_.size(), GPU_MAX_HW_QUEUES); + + // If we have reached the max number of queues, reuse an existing queue, + // choosing the one with the least number of users. + if (queuePool_.size() == GPU_MAX_HW_QUEUES) { + typedef decltype(queuePool_)::const_reference PoolRef; + auto lowest = std::min_element(queuePool_.begin(), queuePool_.end(), + [] (PoolRef A, PoolRef B) { + return A.second.refCount < B.second.refCount; + }); + LogPrintfInfo("selected queue with least refCount: %p (%d)", + lowest->first, lowest->second.refCount); + lowest->second.refCount++; + return lowest->first; + } + + // Else create a new queue. This also includes the initial state where there + // is no queue. + uint32_t queue_max_packets = 0; + if (HSA_STATUS_SUCCESS != + hsa_agent_get_info(_bkendDevice, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_max_packets)) { + return nullptr; + } + auto queue_size = (queue_max_packets < queue_size_hint) ? queue_max_packets : queue_size_hint; + + hsa_queue_t *queue; + while (hsa_queue_create(_bkendDevice, queue_size, HSA_QUEUE_TYPE_MULTI, nullptr, nullptr, + std::numeric_limits::max(), std::numeric_limits::max(), + &queue) != HSA_STATUS_SUCCESS) { + queue_size >>= 1; + if (queue_size < 64) { + return nullptr; + } + } + LogPrintfInfo("created hardware queue %p with size %d", + queue, queue_size); + hsa_amd_profiling_set_profiler_enabled(queue, 1); + auto result = queuePool_.emplace(std::make_pair(queue, QueueInfo())); + assert(result.second && "QueueInfo already exists"); + auto &qInfo = result.first->second; + qInfo.refCount = 1; + return queue; } + +void Device::releaseQueue(hsa_queue_t* queue) { + auto qIter = queuePool_.find(queue); + assert(qIter != queuePool_.end()); + + auto &qInfo = qIter->second; + assert(qInfo.refCount > 0); + qInfo.refCount--; + if (qInfo.refCount != 0) { + return; + } + LogPrintfInfo("deleting hardware queue %p with refCount 0", queue); + + hsa_queue_destroy(queue); + queuePool_.erase(qIter); +} + +} // namespace roc #endif // WITHOUT_HSA_BACKEND diff --git a/projects/clr/rocclr/runtime/device/rocm/rocdevice.hpp b/projects/clr/rocclr/runtime/device/rocm/rocdevice.hpp index b63d772161..378952501e 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocdevice.hpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocdevice.hpp @@ -406,8 +406,6 @@ class Device : public NullDevice { VirtualGPU* xferQueue() const; - std::map& QueuePool() { return queue_pool_; } - hsa_amd_memory_pool_t SystemSegment() const { return system_segment_; } hsa_amd_memory_pool_t SystemCoarseSegment() const { return system_coarse_segment_; } @@ -442,10 +440,17 @@ class Device : public NullDevice { std::atomic freeMem_; //!< Total of free memory available mutable amd::Monitor vgpusAccess_; //!< Lock to serialise virtual gpu list access bool hsa_exclusive_gpu_access_; //!< TRUE if current device was moved into exclusive GPU access mode - std::map queue_pool_; //!< Pool of HSA queues for recycling + + struct QueueInfo { + int refCount; + }; + std::map queuePool_; //!< Pool of HSA queues for recycling public: amd::Atomic numOfVgpus_; //!< Virtual gpu unique index + + hsa_queue_t *acquireQueue(uint32_t queue_size_hint); + void releaseQueue(hsa_queue_t*); }; // class roc::Device } // namespace roc diff --git a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp index fda6732fff..91e8f43daa 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp @@ -555,6 +555,7 @@ bool VirtualGPU::releaseGpuMemoryFence() { VirtualGPU::VirtualGPU(Device& device) : device::VirtualDevice(device), + gpu_queue_(nullptr), roc_device_(device), virtualQueue_(nullptr), deviceQueueSize_(0), @@ -631,12 +632,9 @@ VirtualGPU::~VirtualGPU() { for (uint idx = index(); idx < roc_device_.vgpus().size(); ++idx) { roc_device_.vgpus()[idx]->index_--; } - // Decrement the counter - roc_device_.QueuePool()[gpu_queue_]--; - // Release the queue if the counter is 0 - if (roc_device_.QueuePool()[gpu_queue_] == 0) { - hsa_status_t err = hsa_queue_destroy(gpu_queue_); - roc_device_.QueuePool().erase(gpu_queue_); + + if (gpu_queue_) { + roc_device_.releaseQueue(gpu_queue_); } } @@ -646,38 +644,10 @@ bool VirtualGPU::create(bool profilingEna) { return false; } - uint32_t queue_max_packets = 0; - if (HSA_STATUS_SUCCESS != - hsa_agent_get_info(gpu_device_, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_max_packets)) { - return false; - } - // Pick a reasonable queue size uint32_t queue_size = 1024; - queue_size = (queue_max_packets < queue_size) ? queue_max_packets : queue_size; - if (roc_device_.QueuePool().size() < GPU_MAX_HW_QUEUES) { - while (hsa_queue_create(gpu_device_, queue_size, HSA_QUEUE_TYPE_MULTI, nullptr, nullptr, - std::numeric_limits::max(), std::numeric_limits::max(), - &gpu_queue_) != HSA_STATUS_SUCCESS) { - queue_size >>= 1; - if (queue_size < 64) { - return false; - } - } - hsa_amd_profiling_set_profiler_enabled(gpu_queue(), 1); - roc_device_.QueuePool().insert({gpu_queue_, 1}); - } else { - int usage = std::numeric_limits::max(); - // Loop through all allocated queues and find the lowest usage - for (const auto it : roc_device_.QueuePool()) { - if (it.second < usage) { - gpu_queue_ = it.first; - usage = it.second; - } - } - // Increment the usage of the current queue - roc_device_.QueuePool()[gpu_queue_]++; - } + gpu_queue_ = roc_device_.acquireQueue(queue_size); + if (!gpu_queue_) return false; if (!initPool(dev().settings().kernargPoolSize_, (profilingEna) ? queue_size : 0)) { LogError("Couldn't allocate arguments/signals for the queue");