From 3895aadba6d947b1a12f4987c28201411d99ee86 Mon Sep 17 00:00:00 2001 From: German Andryeyev <56892148+gandryey@users.noreply.github.com> Date: Wed, 10 Dec 2025 12:37:10 -0500 Subject: [PATCH] SWDEV-558849 - Make ROCR path in Windows more stable (#2181) --- projects/clr/hipamd/src/hip_runtime.cpp | 4 ++- projects/clr/hipamd/src/hip_stream.cpp | 11 +++++-- projects/clr/rocclr/device/device.cpp | 18 ++++++++++-- projects/clr/rocclr/device/rocm/rocdevice.cpp | 13 ++++++++- projects/clr/rocclr/device/rocm/rocdevice.hpp | 4 +++ .../clr/rocclr/device/rocm/rocvirtual.cpp | 29 ++++++++++++++----- .../clr/rocclr/device/rocm/rocvirtual.hpp | 16 +++++----- projects/clr/rocclr/platform/commandqueue.cpp | 13 ++++++--- 8 files changed, 83 insertions(+), 25 deletions(-) diff --git a/projects/clr/hipamd/src/hip_runtime.cpp b/projects/clr/hipamd/src/hip_runtime.cpp index 50b98ec237..2427b56340 100644 --- a/projects/clr/hipamd/src/hip_runtime.cpp +++ b/projects/clr/hipamd/src/hip_runtime.cpp @@ -52,7 +52,9 @@ extern "C" BOOL WINAPI DllMain(HINSTANCE hinst, DWORD reason, LPVOID reserved) { hip::PlatformState::instance().setDynamicLibraryHandle(static_cast(hinst)); break; case DLL_PROCESS_DETACH: - hip::ihipDestroyDevice(); + if (GPU_ENABLE_PAL != 0) { + hip::ihipDestroyDevice(); + } break; case DLL_THREAD_DETACH: break; diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 10e6e81273..bb087adff0 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -207,7 +207,6 @@ static hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, } // ================================================================================================ - stream_per_thread::stream_per_thread() { m_streams.resize(g_devices.size()); for (auto& stream : m_streams) { @@ -215,15 +214,21 @@ stream_per_thread::stream_per_thread() { } } +// ================================================================================================ stream_per_thread::~stream_per_thread() { for (auto& stream : m_streams) { if (stream != nullptr && hip::isValid(stream)) { - hip::Stream::Destroy(reinterpret_cast(stream)); + // @note: Global variables in hip runtime will be destroyed after ROCR's global variables. + // Any calls to rocr may cause invalid object access. Hence, avoid the stream destruction. + if (IS_LINUX || (GPU_ENABLE_PAL != 0)) { + hip::Stream::Destroy(reinterpret_cast(stream)); + } stream = nullptr; } } } +// ================================================================================================ hipStream_t stream_per_thread::get() { hip::Device* device = hip::getCurrentDevice(); int currDev = device->deviceId(); @@ -246,13 +251,13 @@ hipStream_t stream_per_thread::get() { return m_streams[currDev]; } +// ================================================================================================ void stream_per_thread::clear_spt() { if (!m_streams.empty()) { m_streams[getCurrentDevice()->deviceId()] = nullptr; } } - // ================================================================================================ void getStreamPerThread(hipStream_t& stream) { if (stream == hipStreamPerThread) { diff --git a/projects/clr/rocclr/device/device.cpp b/projects/clr/rocclr/device/device.cpp index 45758c1e62..e50bdf6517 100644 --- a/projects/clr/rocclr/device/device.cpp +++ b/projects/clr/rocclr/device/device.cpp @@ -666,6 +666,7 @@ bool Device::BlitProgram::create(amd::Device* device, const std::string& extraKe return true; } +// ================================================================================================ bool Device::init() { assert(!Runtime::initialized() && "initialize only once"); bool ret = false; @@ -690,7 +691,17 @@ bool Device::init() { // abort(); ClPrint(amd::LOG_DETAIL_DEBUG, amd::LOG_INIT, "KFD is not installed \n"); // Disable direct dispatch if ROC initialization wasn't successful - AMD_DIRECT_DISPATCH = flagIsDefault(AMD_DIRECT_DISPATCH) ? false : AMD_DIRECT_DISPATCH; + if (flagIsDefault(AMD_DIRECT_DISPATCH)) { + AMD_DIRECT_DISPATCH = false; + } + GPU_ENABLE_PAL = 1; + } else { + // ROC initialization successful, enable direct dispatch + if (flagIsDefault(AMD_DIRECT_DISPATCH)) { + AMD_DIRECT_DISPATCH = true; + } + // Disable PAL path + GPU_ENABLE_PAL = 0; } if (!amd::IS_HIP) { ret |= roc::NullDevice::init(); @@ -701,7 +712,9 @@ bool Device::init() { if (GPU_ENABLE_PAL != 0) { if (GPU_ENABLE_PAL == 1) { // PAL path can't support direct dispatch, unless it's forced - AMD_DIRECT_DISPATCH = flagIsDefault(AMD_DIRECT_DISPATCH) ? false : AMD_DIRECT_DISPATCH; + if (flagIsDefault(AMD_DIRECT_DISPATCH)) { + AMD_DIRECT_DISPATCH = false; + } } ret |= PalDeviceLoad(); } @@ -709,6 +722,7 @@ bool Device::init() { return ret; } +// ================================================================================================ void Device::tearDown() { if (devices_ != nullptr) { for (uint i = 0; i < devices_->size(); ++i) { diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index 40c06b2ba1..2c907e2f19 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -441,7 +441,11 @@ bool Device::init() { if (amd::IS_HIP && ROC_GLOBAL_CU_MASK[0] != '\0') { roc_device->getGlobalCUMask(ROC_GLOBAL_CU_MASK); } - + // Note: for now disable HSA path by default except for gfx942 + if (IS_WINDOWS && (GPU_ENABLE_PAL == 2) && + (std::strcmp(roc_device->info().name_, "gfx942") != 0)) { + return false; + } roc_device.release()->registerDevice(); } @@ -1626,6 +1630,13 @@ bool Device::populateOCLDeviceConstants() { LogError("HSA_AMD_AGENT_INFO_NUM_XCC query failed."); } + if (HSA_STATUS_SUCCESS != + Hsa::agent_get_info(bkendDevice_, + static_cast(HSA_AMD_AGENT_INFO_PM4_EMULATION), + &pm4_emulation_)) { + LogError("HSA_AMD_AGENT_INFO_PM4_EMULATION query failed."); + } + ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Gfx Major/Minor/Stepping: %d/%d/%d", isa().versionMajor(), isa().versionMinor(), isa().versionStepping()); ClPrint(amd::LOG_INFO, amd::LOG_INIT, "HMM support: %d, XNACK: %d, Direct host access: %d", diff --git a/projects/clr/rocclr/device/rocm/rocdevice.hpp b/projects/clr/rocclr/device/rocm/rocdevice.hpp index 9c30ef36b8..b6f42cf97c 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.hpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.hpp @@ -615,6 +615,9 @@ class Device : public NullDevice { // Returns the number of allocated normal queues on this device uint32_t NumNormalQueues() const { return num_normal_queues_.load(); } + //! Returns true if PM4 emulation is enabled + bool IsPm4Emulation() const { return pm4_emulation_; } + private: bool create(); @@ -701,6 +704,7 @@ class Device : public NullDevice { uint32_t maxSdmaReadMask_; uint32_t maxSdmaWriteMask_; bool isXgmi_; //!< Flag to indicate if there is XGMI between CPU<->GPU + bool pm4_emulation_ = false; //!< Flag to indicate if PM4 emulation is enabled //! Code object to kernel info map (used in the crash dump analysis) mutable std::map kernel_map_; diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index ff3e548454..f0314a2071 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -559,19 +559,20 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(hsa_signal_value_t init_va (ts->command().Callback() != nullptr || ts->command().GetBatchHead() != nullptr) && !ts->command().CpuWaitRequested(); } + bool use_irq = enqueHandler || (IS_WINDOWS && gpu_.ForceIrq()); // Check if the signal doesn't match the requested one. // Note: runtime needs the interrupts for the callbacks in DD mode - if ((signal_list_[current_id_]->flags_.interrupt_ != enqueHandler) && gpu_.dev().ActiveWait()) { + if ((signal_list_[current_id_]->flags_.interrupt_ != use_irq) && gpu_.dev().ActiveWait()) { // Use different stacks if an interrupt is required or not. // @note: if runtime needs an interrupt, then the tracking list replaces the original signal // with the interrupt signal and saves the signal without interrupt, or vise versa - auto& pool_get = (enqueHandler) ? signal_pool_irq_ : signal_pool_; - auto& pool_save = (enqueHandler) ? signal_pool_ : signal_pool_irq_; + auto& pool_get = (use_irq) ? signal_pool_irq_ : signal_pool_; + auto& pool_save = (use_irq) ? signal_pool_ : signal_pool_irq_; // Check if a free signal in the pop stack isn't available if (pool_get.empty()) { std::unique_ptr signal(new ProfilingSignal()); - if ((signal != nullptr) && CreateSignal(signal.get(), enqueHandler)) { + if ((signal != nullptr) && CreateSignal(signal.get(), use_irq)) { pool_get.push(signal.release()); } } @@ -1200,8 +1201,12 @@ bool VirtualGPU::dispatchGenericAqlPacket(AqlPacket* packet, uint16_t header, ui reinterpret_cast(packet)->completion_signal, reinterpret_cast(packet)->reserved2, Hsa::queue_load_read_index_scacquire(gpu_queue_), index); - - Hsa::signal_store_screlease(gpu_queue_->doorbell_signal, index); + // Optimization for native AQL path in windows has problems with PM4 emulation, + // skipping the doorbel will not wake up the AQL worker thread + //if (IS_WINDOWS && !dev().IsPm4Emulation() && (blocking || !hasPendingDispatch_)) + { + Hsa::signal_store_screlease(gpu_queue_->doorbell_signal, index); + } // Mark the flag indicating if a dispatch is outstanding. // We are not waiting after every dispatch. @@ -1793,6 +1798,11 @@ VirtualGPU::~VirtualGPU() { if (gpu_queue_ == nullptr) { gpu_queue_ = roc_device_.AcquireActiveNormalQueue(); } + // Windows requires an interrupt in more cases than Linux for OS fence updates + force_irq_ = IS_WINDOWS; + // Force extra barrier to make sure OS gets an interrupt, + // but avoid if the PM4 emulation, since PM4 path can deadlock during device destruction + hasPendingDispatch_ |= IS_WINDOWS && !dev().IsPm4Emulation(); // Release the resources of signal releaseGpuMemoryFence(); } @@ -4038,6 +4048,7 @@ void VirtualGPU::submitMarker(amd::Marker& vcmd) { // Make sure VirtualGPU has an exclusive access to the resources amd::ScopedLock lock(execution()); if (vcmd.CpuWaitRequested()) { + force_irq_ = IS_WINDOWS; // It should be safe to call flush directly if there are not pending dispatches without // HSA signal callback if (gpu_queue_ == nullptr) { @@ -4058,6 +4069,7 @@ void VirtualGPU::submitMarker(amd::Marker& vcmd) { } } else { // Submit a barrier with a cache flushes. + force_irq_ = IS_WINDOWS; if (settings.barrier_value_packet_ && vcmd.profilingInfo().marker_ts_) { dispatchBarrierValuePacket(kBarrierVendorPacketHeader, true); } else { @@ -4068,6 +4080,7 @@ void VirtualGPU::submitMarker(amd::Marker& vcmd) { } profilingEnd(); } + force_irq_ = false; } } @@ -4132,7 +4145,9 @@ void VirtualGPU::addPinnedMem(amd::Memory* mem) { pinnedMems_.push_back(mem); } } else { - if (command_ != nullptr) { + // Optimize pinning path for Linux only(KFD has special tracking for pinned memory) or OpenCL, + // since OpenCL always waits for completion on CPU + if ((command_ != nullptr) && (IS_LINUX || !amd::IS_HIP)) { command_->AddPinnedMemory(mem); } else { //! @note: ROCr backend doesn't have per resource busy tracking, hence runtime has to wait diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index 90cfe26399..de3185fd06 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.hpp @@ -109,7 +109,7 @@ class Timestamp : public amd::ReferenceCountedObject { uint64_t start_; uint64_t end_; VirtualGPU* gpu_; //!< Virtual GPU, associated with this timestamp - amd::Command& command_; ///!< Command, associated with this timestamp + amd::Command& command_; //!< Command, associated with this timestamp amd::Command* parsedCommand_; //!< Command down the list, considering command_ as head std::vector signals_; //!< The list of all signals, associated with the TS hsa_signal_t callback_signal_; //!< Signal associated with a callback for possible later update @@ -237,7 +237,7 @@ class VirtualGPU : public device::VirtualDevice { ~MemoryDependency() { delete[] memObjectsInQueue_; } - //! Creates memory dependecy structure + //! Creates memory dependency structure bool create(size_t numMemObj); //! Notify the tracker about new kernel @@ -274,7 +274,7 @@ class VirtualGPU : public device::VirtualDevice { //! Creates a pool of signals for tracking of HW operations on the queue bool Create(); - //! Finds a free signal for the upcomming operation + //! Finds a free signal for the upcoming operation hsa_signal_t ActiveSignal(hsa_signal_value_t init_val = kInitSignalValueOne, Timestamp* ts = nullptr, bool attach_signal = true); @@ -457,6 +457,7 @@ class VirtualGPU : public device::VirtualDevice { //! Analyzes a crashed AQL queue to find a broken AQL packet void AnalyzeAqlQueue() const; + bool ForceIrq() const { return force_irq_; } private: //! Dispatches a barrier with blocking HSA signals @@ -498,7 +499,7 @@ class VirtualGPU : public device::VirtualDevice { bool createSchedulerParam(); - //! Returns TRUE if virtual queue was successfully allocatted + //! Returns TRUE if virtual queue was successfully allocated bool createVirtualQueue(uint deviceQueueSize); //! Common function for fill memory used by both svm Fill and non-svm fill @@ -525,7 +526,7 @@ class VirtualGPU : public device::VirtualDevice { amd::CopyMetadata copyMetadata = amd::CopyMetadata() //!< Memory copy MetaData ); - //! Updates AQL header for the upcomming dispatch + //! Updates AQL header for the upcoming dispatch void setAqlHeader(uint16_t header) { aqlHeader_ = header; } //! Resets the current queue state. Note: should be called after AQL queue becomes idle @@ -569,6 +570,7 @@ class VirtualGPU : public device::VirtualDevice { uint32_t addSystemScope_ : 1; //!< Insert a system scope to the next aql uint32_t tracking_created_ : 1; //!< Enabled if tracking object was properly initialized uint32_t retainExternalSignals_ : 1; //!< Indicate to retain external signal array + uint32_t force_irq_ : 1; //!< Forces interrupt on the signal completion }; uint32_t state_; }; @@ -615,7 +617,7 @@ class VirtualGPU : public device::VirtualDevice { amd::CommandQueue::Priority priority_; //!< The priority for the hsa queue cl_command_type copy_command_type_; //!< Type of the copy command, used for ROC profiler - //!< OCL doesn't distinguish diffrent copy types, + //!< OCL doesn't distinguish different copy types, //!< but ROC profiler expects D2H or H2D detection int fence_state_; //!< Fence scope //!< kUnknown/kFlushedToDevice/kFlushedToSystem @@ -623,7 +625,7 @@ class VirtualGPU : public device::VirtualDevice { uint64_t last_write_index_ = 0; //!< The last HW queue write index for any packet uint64_t last_barrier_index_ = 0; //!< The last HW queue write index for a packet - //!< with a complition signal + //!< with a completion signal hsa_signal_t last_completion_signal_{}; //!< The last completion signal using KernelArgImpl = device::Settings::KernelArgImpl; diff --git a/projects/clr/rocclr/platform/commandqueue.cpp b/projects/clr/rocclr/platform/commandqueue.cpp index 34ccc0ebb8..98c80506b8 100644 --- a/projects/clr/rocclr/platform/commandqueue.cpp +++ b/projects/clr/rocclr/platform/commandqueue.cpp @@ -62,7 +62,7 @@ HostQueue::HostQueue(Context& context, Device& device, cl_command_queue_properti } bool HostQueue::terminate() { - // incase of force destroy skip checking on the last command + // In case of force destroy skip checking on the last command if (AMD_DIRECT_DISPATCH) { if (!forceDestroy_ && vdev() != nullptr) { // If the queue still has the last command, then wait and release it @@ -201,7 +201,7 @@ void HostQueue::finish(bool cpu_wait) { command->enqueue(); } - // Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status + // Check HW status of the ROCclr event. Note: not all ROCclr modes support HW status static constexpr bool kWaitCompletion = true; if (cpu_wait || !device().IsHwEventReady(command->event(), kWaitCompletion, GetSyncPolicy())) { ClPrint(LOG_DETAIL_DEBUG, LOG_CMD, @@ -217,8 +217,13 @@ void HostQueue::finish(bool cpu_wait) { // during finish() if (command == lastEnqueueCommand_) { device_.removeFromActiveQueues(this); - lastEnqueueCommand_->release(); - lastEnqueueCommand_ = nullptr; + // Under Windows runtime can't destroy objects in the callback thread. + // Also runtime should force interrupt before any destroy. Hence, if it was just gpu wait, + // then keep the lastEnqueueCommand_ for the interrupt handling. + if (IS_LINUX || cpu_wait) { + lastEnqueueCommand_->release(); + lastEnqueueCommand_ = nullptr; + } } } // Release all HW queues, which are idle or nearly idle