diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index e12240a8fd..9e32b9b6f3 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -1479,7 +1479,7 @@ device::VirtualDevice* Device::createVirtualDevice(amd::CommandQueue* queue) { // Initialization of heap and other resources occur during the command // queue creation time. - VirtualGPU* virtualDevice = new VirtualGPU(*this, profiling, cooperative); + VirtualGPU* virtualDevice = new VirtualGPU(*this, profiling, cooperative, queue->cuMask()); if (!virtualDevice->create()) { delete virtualDevice; @@ -1908,7 +1908,8 @@ static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) { } } -hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue) { +hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, + const std::vector& cuMask) { assert(queuePool_.size() <= GPU_MAX_HW_QUEUES); ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "number of allocated hardware queues: %d, maximum: %d", queuePool_.size(), GPU_MAX_HW_QUEUES); @@ -1916,7 +1917,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue) { // If we have reached the max number of queues, reuse an existing queue, // choosing the one with the least number of users. // Note: Don't attempt to reuse the cooperative queue, since it's single per device - if (!coop_queue && (queuePool_.size() == GPU_MAX_HW_QUEUES)) { + if (!coop_queue && (cuMask.size() == 0) && (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) { @@ -1958,6 +1959,24 @@ 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, cooperative: %i", queue, queue_size, coop_queue); hsa_amd_profiling_set_profiler_enabled(queue, 1); + if (cuMask.size() != 0) { + std::stringstream ss; + ss << std::hex; + for (int i = cuMask.size() - 1; i >= 0; i--) { + ss << cuMask[i]; + } + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "setting custom CU mask 0x%s for hardware queue %p", + ss.str().c_str(), queue); + hsa_status_t status = HSA_STATUS_SUCCESS; + status = hsa_amd_queue_cu_set_mask(queue, cuMask.size() * 32, cuMask.data()); + if (status != HSA_STATUS_SUCCESS) { + DevLogError("Device::acquireQueue: hsa_amd_queue_cu_set_mask failed!"); + hsa_queue_destroy(queue); + return nullptr; + } + // Skip queue recycling for queues with custom CU mask + return queue; + } if (coop_queue) { // Skip queue recycling for cooperative queues, since it should be just one // per device. diff --git a/rocclr/device/rocm/rocdevice.hpp b/rocclr/device/rocm/rocdevice.hpp index 0013213174..dd623a2d28 100644 --- a/rocclr/device/rocm/rocdevice.hpp +++ b/rocclr/device/rocm/rocdevice.hpp @@ -432,7 +432,8 @@ class Device : public NullDevice { //! Acquire HSA queue. This method can create a new HSA queue or //! share previously created - hsa_queue_t* acquireQueue(uint32_t queue_size_hint, bool coop_queue = false); + hsa_queue_t* acquireQueue(uint32_t queue_size_hint, bool coop_queue = false, + const std::vector& cuMask = {}); //! Release HSA queue void releaseQueue(hsa_queue_t*); diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index dafdc21f0d..b1a5843eab 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -590,7 +590,8 @@ bool VirtualGPU::releaseGpuMemoryFence() { return true; } -VirtualGPU::VirtualGPU(Device& device, bool profiling, bool cooperative) +VirtualGPU::VirtualGPU(Device& device, bool profiling, bool cooperative, + const std::vector& cuMask) : device::VirtualDevice(device), state_(0), gpu_queue_(nullptr), @@ -601,7 +602,8 @@ VirtualGPU::VirtualGPU(Device& device, bool profiling, bool cooperative) schedulerThreads_(0), schedulerParam_(nullptr), schedulerQueue_(nullptr), - schedulerSignal_({0}) + schedulerSignal_({0}), + cuMask_(cuMask) { index_ = device.numOfVgpus_++; gpu_device_ = device.getBackendDevice(); @@ -643,6 +645,7 @@ VirtualGPU::VirtualGPU(Device& device, bool profiling, bool cooperative) // Note: Virtual GPU device creation must be a thread safe operation roc_device_.vgpus_.resize(roc_device_.numOfVgpus_); roc_device_.vgpus_[index()] = this; + } VirtualGPU::~VirtualGPU() { @@ -702,7 +705,7 @@ VirtualGPU::~VirtualGPU() { bool VirtualGPU::create() { // Pick a reasonable queue size uint32_t queue_size = 1024; - gpu_queue_ = roc_device_.acquireQueue(queue_size, cooperative_); + gpu_queue_ = roc_device_.acquireQueue(queue_size, cooperative_, cuMask_); if (!gpu_queue_) return false; if (!initPool(dev().settings().kernargPoolSize_, (profiling_) ? queue_size : 0)) { diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index 90f7a11420..e14a5101d2 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -163,7 +163,8 @@ class VirtualGPU : public device::VirtualDevice { size_t maxMemObjectsInQueue_; //!< Maximum number of mem objects in the queue }; - VirtualGPU(Device& device, bool profiling = false, bool cooperative = false); + VirtualGPU(Device& device, bool profiling = false, bool cooperative = false, + const std::vector& cuMask = {}); ~VirtualGPU(); bool create(); @@ -362,6 +363,8 @@ class VirtualGPU : public device::VirtualDevice { uint16_t dispatchPacketHeaderNoSync_; uint16_t dispatchPacketHeader_; + + const std::vector& cuMask_; //!< The CU mask }; template diff --git a/rocclr/platform/commandqueue.cpp b/rocclr/platform/commandqueue.cpp index 911bdbe8ee..9d7f6b9884 100644 --- a/rocclr/platform/commandqueue.cpp +++ b/rocclr/platform/commandqueue.cpp @@ -34,9 +34,9 @@ namespace amd { HostQueue::HostQueue(Context& context, Device& device, cl_command_queue_properties properties, - uint queueRTCUs, Priority priority) + uint queueRTCUs, Priority priority, const std::vector& cuMask) : CommandQueue(context, device, properties, device.info().queueProperties_, - queueRTCUs, priority), + queueRTCUs, priority, cuMask), lastEnqueueCommand_(nullptr) { if (thread_.state() >= Thread::INITIALIZED) { ScopedLock sl(queueLock_); diff --git a/rocclr/platform/commandqueue.hpp b/rocclr/platform/commandqueue.hpp index f3f2a57aa5..9a58a68578 100644 --- a/rocclr/platform/commandqueue.hpp +++ b/rocclr/platform/commandqueue.hpp @@ -108,6 +108,9 @@ class CommandQueue : public RuntimeObject { //! Returns the queue priority Priority priority() const { return priority_; } + //! Returns the CU mask array + const std::vector& cuMask() const { return cuMask_; } + protected: //! CommandQueue constructor is protected //! to keep the CommandQueue class as a virtual interface @@ -116,14 +119,16 @@ class CommandQueue : public RuntimeObject { cl_command_queue_properties properties, //!< Queue properties cl_command_queue_properties propMask, //!< Queue properties mask uint rtCUs = RealTimeDisabled, //!< Avaialble real time compute units - Priority priority = Priority::Normal //!< Queue priority + Priority priority = Priority::Normal, //!< Queue priority + const std::vector& cuMask = {} //!< CU mask ) : properties_(propMask, properties), rtCUs_(rtCUs), priority_(priority), queueLock_("CommandQueue::queueLock"), device_(device), - context_(context) {} + context_(context), + cuMask_(cuMask){} Properties properties_; //!< Queue properties uint rtCUs_; //!< The number of used RT compute units @@ -131,6 +136,7 @@ class CommandQueue : public RuntimeObject { Monitor queueLock_; //!< Lock protecting the queue Device& device_; //!< The device SharedReference context_; //!< The context of this command queue + const std::vector& cuMask_; //!< The CU mask private: //! Disable copy constructor @@ -194,7 +200,8 @@ class HostQueue : public CommandQueue { * given device. */ HostQueue(Context& context, Device& device, cl_command_queue_properties properties, - uint queueRTCUs = 0, Priority priority = Priority::Normal); + uint queueRTCUs = 0, Priority priority = Priority::Normal, + const std::vector& cuMask = {}); //! Returns TRUE if this command queue can accept commands. virtual bool create() { return thread_.acceptingCommands_; }