SWDEV-558849 - Make ROCR path in Windows more stable (#2181)

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<void*>(hinst));
       break;
     case DLL_PROCESS_DETACH:
-      hip::ihipDestroyDevice();
+      if (GPU_ENABLE_PAL != 0) {
+        hip::ihipDestroyDevice();
+      }
       break;
     case DLL_THREAD_DETACH:
       break;

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<hip::Stream*>(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<hip::Stream*>(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) {

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) {

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_agent_info_t>(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",

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<uint64_t, Kernel&> kernel_map_;

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<ProfilingSignal> 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<hsa_kernel_dispatch_packet_t*>(packet)->completion_signal,
           reinterpret_cast<hsa_kernel_dispatch_packet_t*>(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

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<ProfilingSignal*> 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;

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