From 8ae30e036d0c61c91d67f5bd17b555bb8b04e13e Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Mon, 10 Apr 2023 15:19:57 -0700 Subject: [PATCH] SWDEV-384557 - Fix engine status query - Maintain a map of SDMA engine# to stream allocated following a greedy approach - Anything past that will query SDMA engine status always and go with a SDMA or Blit copy path Change-Id: Ibfaed7f951ab84d80cb0430596a4d11b5aec9202 [ROCm/clr commit: 5865c642d4463e01e955f5262b9759106264e59c] --- projects/clr/rocclr/device/device.hpp | 1 + projects/clr/rocclr/device/rocm/rocblit.cpp | 131 ++++++++++-------- projects/clr/rocclr/device/rocm/rocblit.hpp | 2 - projects/clr/rocclr/device/rocm/rocdevice.cpp | 85 ++++++++++-- projects/clr/rocclr/device/rocm/rocdevice.hpp | 10 ++ 5 files changed, 161 insertions(+), 68 deletions(-) diff --git a/projects/clr/rocclr/device/device.hpp b/projects/clr/rocclr/device/device.hpp index 3e65a3311a..25fcabdf87 100644 --- a/projects/clr/rocclr/device/device.hpp +++ b/projects/clr/rocclr/device/device.hpp @@ -611,6 +611,7 @@ struct Info : public amd::EmbeddedObject { //! Number of VGPRs per SIMD uint32_t vgprsPerSimd_; uint32_t vgprAllocGranularity_; + uint32_t numSDMAengines_; //!< Number of available SDMA engines }; //! Device settings diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index 38cf6d6109..d170fcaeb0 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -32,9 +32,7 @@ DmaBlitManager::DmaBlitManager(VirtualGPU& gpu, Setup setup) : HostBlitManager(gpu, setup), MinSizeForPinnedTransfer(dev().settings().pinnedMinXferSize_), completeOperation_(false), - context_(nullptr), - lastCopyMask_(0), - lastUsedCopyEngine_(HwQueueEngine::Unknown) {} + context_(nullptr) {} inline void DmaBlitManager::synchronize() const { if (syncOperation_) { @@ -118,8 +116,9 @@ bool DmaBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost, if (pinned != nullptr) { // Get device memory for this virtual device Memory* dstMemory = dev().getRocMemory(pinned); - - if (!hsaCopy(gpuMem(srcMemory), *dstMemory, srcPin, dst, copySizePin)) { + const KernelBlitManager *kb = dynamic_cast(this); + if (!kb->copyBuffer(gpuMem(srcMemory), *dstMemory, srcPin, dst, + copySizePin)) { LogWarning("DmaBlitManager::readBuffer failed a pinned copy!"); gpu().addPinnedMem(pinned); break; @@ -287,8 +286,9 @@ bool DmaBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemory, if (pinned != nullptr) { // Get device memory for this virtual device Memory* srcMemory = dev().getRocMemory(pinned); - - if (!hsaCopy(*srcMemory, gpuMem(dstMemory), src, dstPin, copySizePin)) { + const KernelBlitManager *kb = dynamic_cast(this); + if (!kb->copyBuffer(*srcMemory, gpuMem(dstMemory), src, dstPin, + copySizePin)) { LogWarning("DmaBlitManager::writeBuffer failed a pinned copy!"); gpu().addPinnedMem(pinned); break; @@ -677,63 +677,74 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, srcAgent = dstAgent = dev().getBackendDevice(); } + uint32_t copyMask = 0; + uint32_t freeEngineMask = 0; + HwQueueEngine engine = HwQueueEngine::Unknown; if ((srcAgent.handle == dev().getCpuAgent().handle) && (dstAgent.handle != dev().getCpuAgent().handle)) { engine = HwQueueEngine::SdmaWrite; + copyMask = dev().fetchSDMAMask(this, false); } else if ((srcAgent.handle != dev().getCpuAgent().handle) && (dstAgent.handle == dev().getCpuAgent().handle)) { engine = HwQueueEngine::SdmaRead; + copyMask = dev().fetchSDMAMask(this, true); } - auto wait_events = gpu().Barriers().WaitingSignal(engine); - hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp()); - uint32_t freeEngineMask = 0; - uint32_t copyMask = lastCopyMask_; + if (engine != HwQueueEngine::Unknown) { + if (copyMask == Device::kSkipQueryStatus) { + // Do not query engine status or take copy_on_engine path + status = HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } - if ((engine != lastUsedCopyEngine_)) { - // Check SDMA engine status - status = hsa_amd_memory_copy_engine_status(dstAgent, srcAgent, &freeEngineMask); - ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Query copy engine status %x, freemask %x", - status, freeEngineMask); - // Return a mask with the rightmost bit set - copyMask = freeEngineMask - (freeEngineMask & (freeEngineMask - 1)); - } + if (copyMask == 0) { + // Check SDMA engine status + status = hsa_amd_memory_copy_engine_status(dstAgent, srcAgent, &freeEngineMask); + ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Query copy engine status %x, free_engine mask %x", + status, freeEngineMask); + // Return a mask with the rightmost bit set + copyMask = freeEngineMask - (freeEngineMask & (freeEngineMask - 1)); + } - if (copyMask != 0 && engine != HwQueueEngine::Unknown) { - // Copy on the first available free engine if ROCr returns a valid mask - hsa_amd_sdma_engine_id_t copyEngine = static_cast(copyMask); + if (copyMask != 0 && status == HSA_STATUS_SUCCESS) { + auto wait_events = gpu().Barriers().WaitingSignal(engine); + hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp()); + // Copy on the first available free engine if ROCr returns a valid mask + hsa_amd_sdma_engine_id_t copyEngine = static_cast(copyMask); - ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, - "HSA Async Copy on copy_engine=%x, dst=0x%zx, src=0x%zx, " - "size=%ld, wait_event=0x%zx, completion_signal=0x%zx", copyEngine, - dst, src, size[0], (wait_events.size() != 0) ? wait_events[0].handle : 0, - active.handle); + ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, + "HSA Async Copy on copy_engine=%x, dst=0x%zx, src=0x%zx, " + "size=%ld, wait_event=0x%zx, completion_signal=0x%zx", copyEngine, + dst, src, size[0], (wait_events.size() != 0) ? wait_events[0].handle : 0, + active.handle); - status = hsa_amd_memory_async_copy_on_engine(dst, dstAgent, src, srcAgent, - size[0], wait_events.size(), - wait_events.data(), active, copyEngine, false); + status = hsa_amd_memory_async_copy_on_engine(dst, dstAgent, src, srcAgent, + size[0], wait_events.size(), + wait_events.data(), active, copyEngine, false); + if (status != HSA_STATUS_SUCCESS) { + gpu().Barriers().ResetCurrentSignal(); + } + } } else { - // Force copy with BLIT in ROCr. Forcing agents to the GPU device causes ROCr to take - // blit path internally - srcAgent = dstAgent = dev().getBackendDevice(); + auto wait_events = gpu().Barriers().WaitingSignal(engine); + hsa_signal_t active = gpu().Barriers().ActiveSignal(kInitSignalValueOne, gpu().timestamp()); ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, - "HSA Async Blit Copy dst=0x%zx, src=0x%zx, size=%ld, wait_event=0x%zx, " - "completion_signal=0x%zx", - dst, src, size[0], (wait_events.size() != 0) ? wait_events[0].handle : 0, - active.handle); + "HSA Async Blit Copy dst=0x%zx, src=0x%zx, size=%ld, wait_event=0x%zx, " + "completion_signal=0x%zx", + dst, src, size[0], (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.data(), active); + if (status != HSA_STATUS_SUCCESS) { + gpu().Barriers().ResetCurrentSignal(); + } } if (status == HSA_STATUS_SUCCESS) { - lastUsedCopyEngine_ = engine; - lastCopyMask_ = copyMask; gpu().addSystemScope(); } else { - gpu().Barriers().ResetCurrentSignal(); - LogPrintfError("HSA copy from host to device failed with code %d", status); + LogPrintfError("HSA copy failed with code %d, falling to Blit copy", status); } return (status == HSA_STATUS_SUCCESS); @@ -855,6 +866,9 @@ KernelBlitManager::~KernelBlitManager() { kernels_[i]->release(); } } + + dev().resetSDMAMask(this); + if (nullptr != program_) { program_->release(); } @@ -2264,9 +2278,26 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds #endif #endif - if (setup_.disableHwlCopyBuffer_ || - (!srcMemory.isHostMemDirectAccess() && !dstMemory.isHostMemDirectAccess() && - !(p2p || asan) && !ipcShared)) { + bool useShaderCopyPath = setup_.disableHwlCopyBuffer_ || + (!srcMemory.isHostMemDirectAccess() && + !dstMemory.isHostMemDirectAccess() && + !(p2p || asan) && !ipcShared); + + if (!useShaderCopyPath) { + if (amd::IS_HIP) { + // Update the command type for ROC profiler + if (srcMemory.isHostMemDirectAccess()) { + gpu().SetCopyCommandType(CL_COMMAND_WRITE_BUFFER); + } + if (dstMemory.isHostMemDirectAccess()) { + gpu().SetCopyCommandType(CL_COMMAND_READ_BUFFER); + } + } + result = DmaBlitManager::copyBuffer(srcMemory, dstMemory, srcOrigin, dstOrigin, sizeIn, entire, + copyMetadata); + } + + if (!result) { uint blitType = BlitCopyBuffer; size_t dim = 1; size_t globalWorkOffset[3] = {0, 0, 0}; @@ -2338,18 +2369,6 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds address parameters = captureArguments(kernels_[blitType]); result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr); releaseArguments(parameters); - } else { - if (amd::IS_HIP) { - // Update the command type for ROC profiler - if (srcMemory.isHostMemDirectAccess()) { - gpu().SetCopyCommandType(CL_COMMAND_WRITE_BUFFER); - } - if (dstMemory.isHostMemDirectAccess()) { - gpu().SetCopyCommandType(CL_COMMAND_READ_BUFFER); - } - } - result = DmaBlitManager::copyBuffer(srcMemory, dstMemory, srcOrigin, dstOrigin, sizeIn, entire, - copyMetadata); } synchronize(); diff --git a/projects/clr/rocclr/device/rocm/rocblit.hpp b/projects/clr/rocclr/device/rocm/rocblit.hpp index 545e205b37..3c755acae8 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.hpp +++ b/projects/clr/rocclr/device/rocm/rocblit.hpp @@ -237,8 +237,6 @@ class DmaBlitManager : public device::HostBlitManager { const size_t MinSizeForPinnedTransfer; bool completeOperation_; //!< DMA blit manager must complete operation amd::Context* context_; //!< A dummy context - mutable uint32_t lastCopyMask_; //!< Last used copy mask - mutable HwQueueEngine lastUsedCopyEngine_; //!< Last used copy engine private: //! Disable copy constructor diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index eacb2f4467..1c24279558 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -240,12 +240,12 @@ Device::~Device() { hsa_queue_t* queue = qIter->first; auto& qInfo = qIter->second; if (qInfo.hostcallBuffer_) { - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "deleting hostcall buffer %p for hardware queue %p", + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hostcall buffer %p for hardware queue %p", qInfo.hostcallBuffer_, qIter->first); 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); qIter = it.erase(qIter); hsa_queue_destroy(queue); } @@ -1183,6 +1183,17 @@ bool Device::populateOCLDeviceConstants() { return false; } + if (HSA_STATUS_SUCCESS != + hsa_agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_NUM_SDMA_ENG), + &info_.numSDMAengines_)) { + return false; + } + + for (uint32_t i = 0; i < info_.numSDMAengines_; i++) { + engineAssignMap_[1 << i] = 0; + } + setupCpuAgent(); checkAtomicSupport(); @@ -1200,7 +1211,8 @@ bool Device::populateOCLDeviceConstants() { hsa_status_t err; // Can another GPU (agent) have access to the current GPU memory pool (gpuvm_segment_)? hsa_amd_memory_pool_access_t access; - err = hsa_amd_agent_memory_pool_get_info(agent, gpuvm_segment_, HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); + err = hsa_amd_agent_memory_pool_get_info(agent, gpuvm_segment_, + HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); if (err != HSA_STATUS_SUCCESS) { continue; } @@ -1214,7 +1226,7 @@ bool Device::populateOCLDeviceConstants() { } } - /* Keep track of all P2P Agents in a Array including current device handle for IPC */ + // Keep track of all P2P Agents in a Array including current device handle for IPC p2p_agents_list_ = new hsa_agent_t[1 + p2p_agents_.size()]; p2p_agents_list_[0] = getBackendDevice(); for (size_t agent_idx = 0; agent_idx < p2p_agents_.size(); ++agent_idx) { @@ -1229,6 +1241,20 @@ bool Device::populateOCLDeviceConstants() { } assert(group_segment_size > 0); + // Find SDMA read mask + if (HSA_STATUS_SUCCESS != hsa_amd_memory_copy_engine_status(getCpuAgent(), getBackendDevice(), + &maxSdmaReadMask)) { + return false; + } + assert(maxSdmaReadMask > 0 && "No SDMA engines available for Read"); + + // Find SDMA write mask + if (HSA_STATUS_SUCCESS != hsa_amd_memory_copy_engine_status(getBackendDevice(), getCpuAgent(), + &maxSdmaWriteMask)) { + return false; + } + assert(maxSdmaWriteMask > 0 && "No SDMA engines available for Write"); + info_.localMemSizePerCU_ = group_segment_size; info_.localMemSize_ = group_segment_size; @@ -1631,7 +1657,7 @@ bool Device::populateOCLDeviceConstants() { LogError("HSA_AMD_AGENT_INFO_SVM_DIRECT_HOST_ACCESS query failed."); } - ClPrint(amd::LOG_INFO, amd::LOG_INIT, "HMM support: %d, xnack: %d, direct host access: %d\n", + ClPrint(amd::LOG_INFO, amd::LOG_INIT, "HMM support: %d, xnack: %d, direct host access: %d", info_.hmmSupported_, info_.hmmCpuMemoryAccessible_, info_.hmmDirectHostAccess_); info_.globalCUMask_ = {}; @@ -2185,7 +2211,7 @@ bool Device::IpcCreate(void* dev_ptr, size_t* mem_size, void* handle, size_t* me amd::Memory* amd_mem_obj = amd::MemObjMap::FindMemObj(dev_ptr); if (amd_mem_obj == nullptr) { - DevLogPrintfError("Cannot retrieve amd_mem_obj for dev_ptr: 0x%x \n", dev_ptr); + DevLogPrintfError("Cannot retrieve amd_mem_obj for dev_ptr: 0x%x", dev_ptr); return false; } @@ -2788,7 +2814,7 @@ 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)\n", it->first, + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "selected queue refCount: %p (%d)", it->first, it->second.refCount); return it->first; } @@ -2930,7 +2956,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue, for (int i = mask.size() - 1; i >= 0; i--) { ss << mask[i]; } - ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "setting CU mask 0x%s for hardware queue %p", + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Setting CU mask 0x%s for hardware queue %p", ss.str().c_str(), queue); hsa_status_t status = hsa_amd_queue_cu_set_mask(queue, mask.size() * 32, mask.data()); @@ -2960,7 +2986,7 @@ 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)\n", result.first->first, + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "acquireQueue refCount: %p (%d)", result.first->first, result.first->second.refCount); return queue; } @@ -2972,7 +2998,7 @@ 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)\n", qIter->first, + ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "releaseQueue refCount:%p (%d)", qIter->first, qIter->second.refCount); } } @@ -3292,6 +3318,45 @@ void Device::HiddenHeapAlloc(const VirtualGPU& gpu) { std::call_once(heap_initialized_, HeapAllocZeroOut); } +// ================================================================================================ +uint32_t Device::fetchSDMAMask(const device::BlitManager* handle, bool readEngine) const { + uint32_t engine = 0; + { + amd::ScopedLock lock(vgpusAccess()); + for (auto it = engineAssignMap_.rbegin(); it != engineAssignMap_.rend(); ++it) { + // If blitManager handle is in the map return the engine ID else + // add to the map + if (it->second == handle) { + engine = it->first; + break; + } else if (it->second == 0) { + it->second = handle; + engine = it->first; + break; + } + } + } + + uint32_t mask = (readEngine ? maxSdmaReadMask : maxSdmaWriteMask) & engine; + if (engine != 0 && mask == 0 ) { + return kSkipQueryStatus; + } else { + return mask; + } +} + +// ================================================================================================ +void Device::resetSDMAMask(const device::BlitManager* handle) const { + amd::ScopedLock lock(vgpusAccess()); + + for (auto& it : engineAssignMap_) { + if (it.second == handle) { + it.second = 0; + break; + } + } +} + // ================================================================================================ ProfilingSignal::~ProfilingSignal() { if (signal_.handle != 0) { diff --git a/projects/clr/rocclr/device/rocm/rocdevice.hpp b/projects/clr/rocclr/device/rocm/rocdevice.hpp index b36d21ef7a..348c6c0687 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.hpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.hpp @@ -551,6 +551,9 @@ class Device : public NullDevice { //! Allocates hidden heap for device memory allocations void HiddenHeapAlloc(const VirtualGPU& gpu); + uint32_t fetchSDMAMask(const device::BlitManager* handle, bool readEngine = true) const; + void resetSDMAMask(const device::BlitManager* handle) const ; + private: bool create(); @@ -619,7 +622,14 @@ class Device : public NullDevice { //! Pool of HSA queues with custom CU masks std::vector> queueWithCUMaskPool_; + //! Read and Write mask for device<->host + uint32_t maxSdmaReadMask; + uint32_t maxSdmaWriteMask; + //! Map of SDMA engineId<->stream + mutable std::map engineAssignMap_; + public: + constexpr static uint32_t kSkipQueryStatus = 1 << 31; std::atomic numOfVgpus_; //!< Virtual gpu unique index //! enum for keeping the total and available queue priorities