diff --git a/rocclr/device/rocm/rocblit.cpp b/rocclr/device/rocm/rocblit.cpp index 9009b443c5..716790c837 100644 --- a/rocclr/device/rocm/rocblit.cpp +++ b/rocclr/device/rocm/rocblit.cpp @@ -379,6 +379,7 @@ bool DmaBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& dstMe return true; } +// ================================================================================================ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& dstMemory, const amd::BufferRect& srcRect, const amd::BufferRect& dstRect, const amd::Coord3D& size, bool entire) const { @@ -435,8 +436,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d } if (isSubwindowRectCopy ) { - const hsa_signal_value_t kInitVal = 1; - hsa_signal_store_relaxed(completion_signal_, kInitVal); + hsa_signal_store_relaxed(completion_signal_, kInitSignalValueOne); // Copy memory line by line hsa_status_t status = @@ -447,10 +447,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d return false; } - - hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, - uint64_t(-1), HSA_WAIT_STATE_BLOCKED); - if (val != 0) { + if (!WaitForSignal(completion_signal_)) { LogError("Async copy failed"); return false; } @@ -476,9 +473,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d } } - hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, - uint64_t(-1), HSA_WAIT_STATE_BLOCKED); - if (val != 0) { + if (!WaitForSignal(completion_signal_)) { LogError("Async copy failed"); return false; } @@ -488,6 +483,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d return true; } +// ================================================================================================ bool DmaBlitManager::copyImageToBuffer(device::Memory& srcMemory, device::Memory& dstMemory, const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, bool entire, size_t rowPitch, @@ -598,6 +594,7 @@ bool DmaBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dstMem return result; } +// ================================================================================================ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin, const amd::Coord3D& size, bool enableCopyRect, bool flushDMA) const { @@ -639,8 +636,7 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, srcAgent = dstAgent = dev().getBackendDevice(); } - const hsa_signal_value_t kInitVal = 1; - hsa_signal_store_relaxed(completion_signal_, kInitVal); + hsa_signal_store_relaxed(completion_signal_, kInitSignalValueOne); // Use SDMA to transfer the data status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, size[0], 0, nullptr, @@ -649,21 +645,7 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, if (status == HSA_STATUS_SUCCESS) { hsa_signal_value_t val; - // Use ACTIVE wait for small transfers. - // Might want to be dependent on also having an idle GPU - // or, if queue is busy, may want to enqueue a blank barrier - // before this and wait BLOCKED on its completion signal, followed - // by ACTIVE on this. - - constexpr size_t small_transfer_size = 4 * Mi; - if (size[0] < small_transfer_size) { - val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, - std::numeric_limits::max(), HSA_WAIT_STATE_ACTIVE); - } else { - val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, - std::numeric_limits::max(), HSA_WAIT_STATE_BLOCKED); - } - if (val != (kInitVal - 1)) { + if (!WaitForSignal(completion_signal_)) { LogError("Async copy failed"); status = HSA_STATUS_ERROR; } else { @@ -676,6 +658,7 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, return (status == HSA_STATUS_SUCCESS); } +// ================================================================================================ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_t size, address staging, bool hostToDev) const { // No allocation is necessary for Full Profile @@ -693,12 +676,10 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_ address hsaBuffer = staging; - const hsa_signal_value_t kInitVal = 1; - // Allocate requested size of memory while (totalSize > 0) { size = std::min(totalSize, dev().settings().stagedXferSize_); - hsa_signal_silent_store_relaxed(completion_signal_, kInitVal); + hsa_signal_silent_store_relaxed(completion_signal_, kInitSignalValueOne); // Copy data from Host to Device if (hostToDev) { @@ -712,10 +693,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_ status = hsa_amd_memory_async_copy(hostDst + offset, dev().getBackendDevice(), hsaBuffer, srcAgent, size, 0, nullptr, completion_signal_); if (status == HSA_STATUS_SUCCESS) { - hsa_signal_value_t val = hsa_signal_wait_scacquire( - completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED); - - if (val != (kInitVal - 1)) { + if (!WaitForSignal(completion_signal_)) { LogError("Async copy failed"); return false; } @@ -739,10 +717,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_ hsa_amd_memory_async_copy(hsaBuffer, dstAgent, hostSrc + offset, dev().getBackendDevice(), size, 0, nullptr, completion_signal_); if (status == HSA_STATUS_SUCCESS) { - hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, - 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED); - - if (val != (kInitVal - 1)) { + if (!WaitForSignal(completion_signal_)) { LogError("Async copy failed"); return false; } @@ -760,6 +735,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_ return true; } +// ================================================================================================ KernelBlitManager::KernelBlitManager(VirtualGPU& gpu, Setup setup) : DmaBlitManager(gpu, setup), program_(nullptr), @@ -1659,6 +1635,7 @@ bool KernelBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory return result; } +// ================================================================================================ bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost, const amd::Coord3D& origin, const amd::Coord3D& size, bool entire) const { @@ -1667,12 +1644,13 @@ bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost, if (dev().info().largeBar_ && size[0] <= kMaxD2hMemcpySize) { if ((srcMemory.owner()->getHostMem() == nullptr) && (srcMemory.owner()->getSvmPtr() != nullptr)) { - // CPU read ahead, hence release GPU memory - gpu().releaseGpuMemoryFence(); + // CPU read ahead, hence release GPU memory and force barrier to make sure L2 flush + constexpr bool ForceBarrier = true; + gpu().releaseGpuMemoryFence(ForceBarrier); char* src = reinterpret_cast(srcMemory.owner()->getSvmPtr()); std::memcpy(dstHost, src + origin[0], size[0]); - // Set HASPENDINGDISPATCH_ FLAG. That will force L2 invalidation on flush - gpu().hasPendingDispatch(); + // The first dispatch will invalidate L2 + gpu().addSystemScope(); return true; } } @@ -1717,6 +1695,7 @@ bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost, return result; } +// ================================================================================================ bool KernelBlitManager::readBufferRect(device::Memory& srcMemory, void* dstHost, const amd::BufferRect& bufRect, const amd::BufferRect& hostRect, const amd::Coord3D& size, @@ -1764,6 +1743,7 @@ bool KernelBlitManager::readBufferRect(device::Memory& srcMemory, void* dstHost, return result; } +// ================================================================================================ bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemory, const amd::Coord3D& origin, const amd::Coord3D& size, bool entire) const { @@ -1773,12 +1753,13 @@ bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemo if (dev().info().largeBar_ && size[0] <= kMaxH2dMemcpySize) { if ((dstMemory.owner()->getHostMem() == nullptr) && (dstMemory.owner()->getSvmPtr() != nullptr)) { // CPU read ahead, hence release GPU memory - gpu().releaseGpuMemoryFence(); + constexpr bool ForceBarrier = true; + gpu().releaseGpuMemoryFence(ForceBarrier); char* dst = reinterpret_cast(dstMemory.owner()->getSvmPtr()); std::memcpy(dst + origin[0], srcHost, size[0]); // Set HASPENDINGDISPATCH_ FLAG. Then releaseGpuMemoryFence() will use barrier to invalidate cache gpu().hasPendingDispatch(); - gpu().releaseGpuMemoryFence(); + gpu().releaseGpuMemoryFence(ForceBarrier); return true; } } @@ -1825,6 +1806,7 @@ bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemo return result; } +// ================================================================================================ bool KernelBlitManager::writeBufferRect(const void* srcHost, device::Memory& dstMemory, const amd::BufferRect& hostRect, const amd::BufferRect& bufRect, const amd::Coord3D& size, @@ -2284,6 +2266,7 @@ address KernelBlitManager::captureArguments(const amd::Kernel* kernel) const { void KernelBlitManager::releaseArguments(address args) const { } +// ================================================================================================ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, hsa_queue_t* schedulerQueue, hsa_signal_t& schedulerSignal, @@ -2310,7 +2293,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, sp->child_queue = reinterpret_cast(schedulerQueue); sp->complete_signal = schedulerSignal; - hsa_signal_store_relaxed(schedulerSignal, 1); + hsa_signal_store_relaxed(schedulerSignal, kInitSignalValueOne); sp->scheduler_aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | @@ -2346,8 +2329,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, } releaseArguments(parameters); - if (hsa_signal_wait_scacquire(schedulerSignal, HSA_SIGNAL_CONDITION_LT, 1, (-1), - HSA_WAIT_STATE_BLOCKED) != 0) { + if (!WaitForSignal(schedulerSignal)) { LogWarning("Failed schedulerSignal wait"); return false; } @@ -2355,6 +2337,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, return true; } +// ================================================================================================ bool KernelBlitManager::RunGwsInit( uint32_t value) const { amd::ScopedLock k(lockXferOps_); diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index 1331600124..0ff244fb1a 100755 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -718,7 +718,7 @@ bool Device::create() { } // Create signal for HMM prefetch operation on device - if (HSA_STATUS_SUCCESS != hsa_signal_create(InitSignalValue, 0, nullptr, &prefetch_signal_)) { + if (HSA_STATUS_SUCCESS != hsa_signal_create(kInitSignalValueOne, 0, nullptr, &prefetch_signal_)) { return false; } @@ -1476,6 +1476,12 @@ bool Device::populateOCLDeviceConstants() { device::VirtualDevice* Device::createVirtualDevice(amd::CommandQueue* queue) { amd::ScopedLock lock(vgpusAccess()); + // If barrier is disabled, then profiling should be enabled to make sure HSA signal is + // attached for every dispatch + if (!settings().barrier_sync_) { + queue->properties().set(CL_QUEUE_PROFILING_ENABLE); + } + bool profiling = (queue != nullptr) && queue->properties().test(CL_QUEUE_PROFILING_ENABLE); bool cooperative = false; @@ -1732,9 +1738,11 @@ device::Memory* Device::createMemory(amd::Memory& owner) const { return memory; } +// ================================================================================================ void* Device::hostAlloc(size_t size, size_t alignment, bool atomics) const { void* ptr = nullptr; - const hsa_amd_memory_pool_t segment = (!atomics) + // If runtime disables barrier, then all host allocations must have L2 disabled + const hsa_amd_memory_pool_t segment = (!atomics && settings().barrier_sync_) ? (system_coarse_segment_.handle != 0) ? system_coarse_segment_ : system_segment_ : system_segment_; assert(segment.handle != 0); @@ -1754,10 +1762,12 @@ void* Device::hostAlloc(size_t size, size_t alignment, bool atomics) const { return ptr; } +// ================================================================================================ void* Device::hostAgentAlloc(size_t size, const AgentInfo& agentInfo, bool atomics) const { void* ptr = nullptr; const hsa_amd_memory_pool_t segment = - (!atomics) ? + // If runtime disables barrier, then all host allocations must have L2 disabled + (!atomics && settings().barrier_sync_) ? (agentInfo.coarse_grain_pool.handle != 0) ? agentInfo.coarse_grain_pool : agentInfo.fine_grain_pool : agentInfo.fine_grain_pool; @@ -1778,6 +1788,7 @@ void* Device::hostAgentAlloc(size_t size, const AgentInfo& agentInfo, bool atomi return ptr; } +// ================================================================================================ void* Device::hostNumaAlloc(size_t size, size_t alignment, bool atomics) const { void* ptr = nullptr; #ifndef ROCCLR_SUPPORT_NUMA_POLICY @@ -2215,7 +2226,7 @@ bool Device::SvmAllocInit(void* memory, size_t size) const { #if AMD_HMM_SUPPORT // Initialize signal for the barrier - hsa_signal_store_relaxed(prefetch_signal_, InitSignalValue); + hsa_signal_store_relaxed(prefetch_signal_, kInitSignalValueOne); // Initiate a prefetch command which should force memory update in HMM hsa_status_t status = hsa_amd_svm_prefetch_async(memory, size, getBackendDevice(), @@ -2226,8 +2237,7 @@ bool Device::SvmAllocInit(void* memory, size_t size) const { } // Wait for the prefetch - if (hsa_signal_wait_scacquire(prefetch_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), - HSA_WAIT_STATE_BLOCKED) != 0) { + if (!WaitForSignal(prefetch_signal_)) { LogError("Barrier packet submission failed"); return false; } diff --git a/rocclr/device/rocm/rocsettings.cpp b/rocclr/device/rocm/rocsettings.cpp index 5fcdd5df3f..c264d06976 100755 --- a/rocclr/device/rocm/rocsettings.cpp +++ b/rocclr/device/rocm/rocsettings.cpp @@ -90,6 +90,7 @@ Settings::Settings() { hmmFlags_ = (!flagIsDefault(ROC_HMM_FLAGS)) ? ROC_HMM_FLAGS : Hmm::EnableSystemMemory; rocr_backend_ = true; + barrier_sync_ = (!flagIsDefault(ROC_BARRIER_SYNC)) ? ROC_BARRIER_SYNC : true; } bool Settings::create(bool fullProfile, int gfxipMajor, int gfxipMinor, bool coop_groups) { diff --git a/rocclr/device/rocm/rocsettings.hpp b/rocclr/device/rocm/rocsettings.hpp index d57a590f04..c12aaf2aa4 100644 --- a/rocclr/device/rocm/rocsettings.hpp +++ b/rocclr/device/rocm/rocsettings.hpp @@ -51,7 +51,8 @@ class Settings : public device::Settings { uint stagedXferRead_ : 1; //!< Uses a staged buffer read uint stagedXferWrite_ : 1; //!< Uses a staged buffer write uint imageBufferWar_ : 1; //!< Image buffer workaround for Gfx10 - uint reserved_ : 24; + uint barrier_sync_ : 1; //!< Use AQL barrier command to sync with CPU + uint reserved_ : 23; }; uint value_; }; @@ -82,7 +83,7 @@ class Settings : public device::Settings { size_t sdmaCopyThreshold_; //!< Use SDMA to copy above this size - uint32_t hmmFlags_; //!< HMM functionality control flags + uint32_t hmmFlags_; //!< HMM functionality control flags //! Default constructor Settings(); diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index da2e0ef69c..98a74fbac7 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -436,10 +436,12 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para return true; } +// ================================================================================================ static inline void packet_store_release(uint32_t* packet, uint16_t header, uint16_t rest) { __atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE); } +// ================================================================================================ template bool VirtualGPU::dispatchGenericAqlPacket( AqlPacket* packet, uint16_t header, uint16_t rest, bool blocking, size_t size) { @@ -478,10 +480,19 @@ bool VirtualGPU::dispatchGenericAqlPacket( } signal = packet->completion_signal; // Initialize signal for a wait - hsa_signal_store_relaxed(signal, InitSignalValue); + hsa_signal_store_relaxed(signal, kInitSignalValueOne); blocking = true; } + // If runtime doesn't use the barrier, then make sure it tracks the last submitted command + if (!dev().settings().barrier_sync_) { + // Initialize signal for a wait + assert(packet->completion_signal.handle != 0 && + "There is no HSA signal associated with the last command!"); + hsa_signal_store_relaxed(packet->completion_signal, kInitSignalValueOne); + last_signal_ = packet->completion_signal; + } + // Insert packet(s) // NOTE: need multiple packets to dispatch the performance counter // packet blob of the legacy devices (gfx8) @@ -521,8 +532,7 @@ bool VirtualGPU::dispatchGenericAqlPacket( // Wait on signal ? if (blocking) { - if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, uint64_t(-1), - HSA_WAIT_STATE_BLOCKED) != 0) { + if (!WaitForSignal(signal)) { LogPrintfError("Failed signal [0x%lx] wait", signal.handle); return false; } @@ -531,16 +541,19 @@ bool VirtualGPU::dispatchGenericAqlPacket( return true; } +// ================================================================================================ bool VirtualGPU::dispatchAqlPacket( hsa_kernel_dispatch_packet_t* packet, uint16_t header, uint16_t rest, bool blocking) { return dispatchGenericAqlPacket(packet, header, rest, blocking); } +// ================================================================================================ bool VirtualGPU::dispatchAqlPacket( hsa_barrier_and_packet_t* packet, uint16_t header, uint16_t rest, bool blocking) { return dispatchGenericAqlPacket(packet, header, rest, blocking); } +// ================================================================================================ bool VirtualGPU::dispatchCounterAqlPacket(hsa_ext_amd_aql_pm4_packet_t* packet, const uint32_t gfxVersion, bool blocking, const hsa_ven_amd_aqlprofile_1_00_pfn_t* extApi) { @@ -599,32 +612,8 @@ void VirtualGPU::dispatchBarrierPacket(const hsa_barrier_and_packet_t* packet) { packet->dep_signal[3], packet->dep_signal[4], packet->completion_signal); } -/** - * @brief Waits on an outstanding kernel without regard to how - * it was dispatched - with or without a signal - * - * @return bool true if Wait returned successfully, false - * otherwise - */ -bool VirtualGPU::releaseGpuMemoryFence() { - // Return if there is no pending dispatch - if (!hasPendingDispatch_) { - return false; - } - - // Initialize signal for the barrier packet. - hsa_signal_store_relaxed(barrier_signal_, InitSignalValue); - - // Dispatch barrier packet into the queue and wait till it finishes. - dispatchBarrierPacket(&barrier_packet_); - if (hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), - HSA_WAIT_STATE_BLOCKED) != 0) { - LogError("Barrier packet submission failed"); - return false; - } - - hasPendingDispatch_ = false; - +// ================================================================================================ +void VirtualGPU::ResetQueueStates() { // Release all transfer buffers on this command queue releaseXferWrite(); @@ -633,10 +622,42 @@ bool VirtualGPU::releaseGpuMemoryFence() { // Release the pool, since runtime just completed a barrier resetKernArgPool(); +} +// ================================================================================================ +bool VirtualGPU::releaseGpuMemoryFence(bool force_barrier) { + // Return if there is no pending dispatch + if (!hasPendingDispatch_) { + return false; + } + hsa_signal_t wait_signal = barrier_signal_; + + // If barrier sync was requested or runtime didn't provide the last signal + if (dev().settings().barrier_sync_ || force_barrier) { + // Initialize signal for the barrier packet. + hsa_signal_store_relaxed(barrier_signal_, kInitSignalValueOne); + + // Dispatch barrier packet into the queue and wait till it finishes. + dispatchBarrierPacket(&barrier_packet_); + } + else { + // Take the signal of the last submitted dispatch + wait_signal = last_signal_; + } + + // Wait for compute work previously submitted + if (!WaitForSignal(wait_signal)) { + LogError("Waiting for compute work failed!"); + return false; + } + + hasPendingDispatch_ = false; + + ResetQueueStates(); return true; } +// ================================================================================================ VirtualGPU::VirtualGPU(Device& device, bool profiling, bool cooperative, const std::vector& cuMask, amd::CommandQueue::Priority priority) @@ -772,7 +793,7 @@ bool VirtualGPU::create() { // Create signal for the barrier packet. hsa_signal_t signal = {0}; - if (HSA_STATUS_SUCCESS != hsa_signal_create(InitSignalValue, 0, nullptr, &signal)) { + if (HSA_STATUS_SUCCESS != hsa_signal_create(kInitSignalValueOne, 0, nullptr, &signal)) { return false; } barrier_signal_ = signal; @@ -851,12 +872,11 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) { //! We can issue a barrier to avoid expensive extra memory allocations. // Initialize signal for the barrier packet. - hsa_signal_store_relaxed(barrier_signal_, InitSignalValue); + hsa_signal_store_relaxed(barrier_signal_, kInitSignalValueOne); // Dispatch barrier packet into the queue and wait till it finishes. dispatchBarrierPacket(&barrier_packet_); - if (hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), - HSA_WAIT_STATE_BLOCKED) != 0) { + if (!WaitForSignal(barrier_signal_)) { LogError("Kernel arguments reset failed"); } @@ -879,7 +899,8 @@ void VirtualGPU::profilingBegin(amd::Command& command, bool drmProfiling) { This could have unintended consequences."); return; } - timestamp_ = new Timestamp; + // Without barrier profiling will wait for each individual signal + timestamp_ = new Timestamp(!dev().settings().barrier_sync_); timestamp_->start(); } } @@ -1202,7 +1223,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) { void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { #if AMD_HMM_SUPPORT // Initialize signal for the barrier - hsa_signal_store_relaxed(barrier_signal_, InitSignalValue); + hsa_signal_store_relaxed(barrier_signal_, kInitSignalValueOne); // Find the requested agent for the transfer hsa_agent_t agent = (cmd.cpu_access() || @@ -1214,9 +1235,7 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { const_cast(cmd.dev_ptr()), cmd.count(), agent, 0, nullptr, barrier_signal_); // Wait for the prefetch - if ((status != HSA_STATUS_SUCCESS) || - hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), - HSA_WAIT_STATE_BLOCKED) != 0) { + if ((status != HSA_STATUS_SUCCESS) || !WaitForSignal(barrier_signal_)) { LogError("hsa_amd_svm_prefetch_async failed"); cmd.setStatus(CL_INVALID_OPERATION); } @@ -2458,20 +2477,38 @@ void VirtualGPU::submitAcquireExtObjects(amd::AcquireExtObjectsCommand& vcmd) { profilingEnd(vcmd); } +// ================================================================================================ void VirtualGPU::submitReleaseExtObjects(amd::ReleaseExtObjectsCommand& vcmd) { // Make sure VirtualGPU has an exclusive access to the resources amd::ScopedLock lock(execution()); profilingBegin(vcmd); + if (!dev().settings().barrier_sync_) { + // Force barrier to make sure L2 flush, since interop can be in sysmem + constexpr bool ForceBarrier = true; + releaseGpuMemoryFence(ForceBarrier); + } profilingEnd(vcmd); } +// ================================================================================================ void VirtualGPU::flush(amd::Command* list, bool wait) { - releaseGpuMemoryFence(); + // If barrier is requested, then wait for everything, otherwise + // a per disaptch wait will occur later in updateCommandsState() + if (dev().settings().barrier_sync_) { + releaseGpuMemoryFence(); + } updateCommandsState(list); + + // Add extra clean up for resources if releaseGpuMemoryFence() was skipped + if (!dev().settings().barrier_sync_) { + ResetQueueStates(); + } + // Release all pinned memory releasePinnedMem(); } +// ================================================================================================ void VirtualGPU::addXferWrite(Memory& memory) { if (xferWriteBuffers_.size() > 7) { dev().xferWrite().release(*this, *xferWriteBuffers_.front()); diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index 862a2dd593..b05a787e3b 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -42,6 +42,26 @@ struct ProfilingSignal : public amd::HeapObject { ProfilingSignal() : ts_(nullptr) { signal_.handle = 0; } }; +// Initial HSA signal value +constexpr hsa_signal_value_t kInitSignalValueOne = 1; + +inline bool WaitForSignal(hsa_signal_t signal) { + constexpr uint64_t Timeout30us = 30000; + constexpr uint64_t UnlimitedWait = std::numeric_limits::max(); + uint64_t timeout = (ROC_ACTIVE_WAIT) ? UnlimitedWait : Timeout30us; + + // Active wait with a timeout + if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, + timeout, HSA_WAIT_STATE_ACTIVE) != 0) { + // Wait until the completion with CPU suspend + if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne, + UnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) { + return false; + } + } + return true; +} + // Timestamp for keeping track of some profiling information for various commands // including EnqueueNDRangeKernel and clEnqueueCopyBuffer. class Timestamp { @@ -53,6 +73,7 @@ class Timestamp { static double ticksToTime_; bool splittedDispatch_; std::vector splittedSignals_; + bool wait_for_signal_; //!< Wait for signal before gathering the timestamp values public: uint64_t getStart() { @@ -75,7 +96,12 @@ class Timestamp { void setAgent(hsa_agent_t agent) { agent_ = agent; } - Timestamp() : start_(0), end_(0), profilingSignal_(nullptr), splittedDispatch_(false) { + Timestamp(bool wait_for_signal = false) + : start_(0) + , end_(0) + , profilingSignal_(nullptr) + , splittedDispatch_(false) + , wait_for_signal_(wait_for_signal) { agent_.handle = 0; } @@ -90,6 +116,9 @@ class Timestamp { uint64_t start = UINT64_MAX; uint64_t end = 0; for (auto it = splittedSignals_.begin(); it < splittedSignals_.end(); it++) { + if (wait_for_signal_) { + WaitForSignal(*it); + } hsa_amd_profiling_get_dispatch_time(agent_, *it, &time); if (time.start < start) { start = time.start; @@ -101,6 +130,9 @@ class Timestamp { start_ = start * ticksToTime_; end_ = end * ticksToTime_; } else { + if (wait_for_signal_) { + WaitForSignal(profilingSignal_->signal_); + } hsa_amd_profiling_get_dispatch_time(agent_, profilingSignal_->signal_, &time); start_ = time.start * ticksToTime_; end_ = time.end * ticksToTime_; @@ -125,9 +157,6 @@ class Timestamp { class VirtualGPU : public device::VirtualDevice { public: - //! Initial signal value - static const hsa_signal_value_t InitSignalValue = 1; - class MemoryDependency : public amd::EmbeddedObject { public: //! Default constructor @@ -227,10 +256,9 @@ class VirtualGPU : public device::VirtualDevice { * @brief Waits on an outstanding kernel without regard to how * it was dispatched - with or without a signal * - * @return bool true if Wait returned successfully, false - * otherwise + * @return bool true if Wait returned successfully, false otherwise */ - bool releaseGpuMemoryFence(); + bool releaseGpuMemoryFence(bool force_barrier = false); hsa_agent_t gpu_device() { return gpu_device_; } hsa_queue_t* gpu_queue() { return gpu_queue_; } @@ -314,6 +342,9 @@ class VirtualGPU : public device::VirtualDevice { //! Updates AQL header for the upcomming dispatch void setAqlHeader(uint16_t header) { aqlHeader_ = header; } + //! Resets the current queue state. Note: should be called after AQL queue becomes idle + void ResetQueueStates(); + std::vector xferWriteBuffers_; //!< Stage write buffers std::vector pinnedMems_; //!< Pinned memory list @@ -336,6 +367,7 @@ class VirtualGPU : public device::VirtualDevice { hsa_queue_t* gpu_queue_; //!< Queue associated with a gpu hsa_barrier_and_packet_t barrier_packet_; hsa_signal_t barrier_signal_; + hsa_signal_t last_signal_ = {}; //!< Last submitted signal uint32_t dispatch_id_; //!< This variable must be updated atomically. Device& roc_device_; //!< roc device object PrintfDbg* printfdbg_; diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp index d4d5398fef..98c0f432f1 100644 --- a/rocclr/utils/flags.hpp +++ b/rocclr/utils/flags.hpp @@ -236,6 +236,10 @@ release(uint, HIP_HIDDEN_FREE_MEM, 0, \ "0 = Disable") \ release(size_t, GPU_FORCE_BLIT_COPY_SIZE, 0, \ "Size in KB of the threshold below which to force blit instead for sdma") \ +release(bool, ROC_BARRIER_SYNC, true, \ + "Enable AQL barrier packet for synchronization") \ +release(bool, ROC_ACTIVE_WAIT, false, \ + "Forces unconditional active wait for GPU") \ release(bool, ROC_ENABLE_LARGE_BAR, true, \ "Enable Large Bar if supported by the device") \ release(bool, HIP_FORCE_QUEUE_PROFILING, false, \