Merge remote-tracking branch 'upstream/sycl' into urlog

Author: mateuszpn

sycl/source/detail/context_impl.cpp

@@ -338,16 +338,23 @@ void context_impl::removeAssociatedDeviceGlobal(const void *DeviceGlobalPtr) {
 void context_impl::addDeviceGlobalInitializer(
     ur_program_handle_t Program, const std::vector<device> &Devs,
     const RTDeviceBinaryImage *BinImage) {
+  if (BinImage->getDeviceGlobals().empty())
+    return;
   std::lock_guard<std::mutex> Lock(MDeviceGlobalInitializersMutex);
   for (const device &Dev : Devs) {
     auto Key = std::make_pair(Program, getSyclObjImpl(Dev)->getHandleRef());
-    MDeviceGlobalInitializers.emplace(Key, BinImage);
+    auto [Iter, Inserted] = MDeviceGlobalInitializers.emplace(Key, BinImage);
+    if (Inserted && !Iter->second.MDeviceGlobalsFullyInitialized)
+      ++MDeviceGlobalNotInitializedCnt;
   }
 }
 
 std::vector<ur_event_handle_t> context_impl::initializeDeviceGlobals(
     ur_program_handle_t NativePrg,
     const std::shared_ptr<queue_impl> &QueueImpl) {
+  if (!MDeviceGlobalNotInitializedCnt.load(std::memory_order_acquire))
+    return {};
+
   const AdapterPtr &Adapter = getAdapter();
   const DeviceImplPtr &DeviceImpl = QueueImpl->getDeviceImplPtr();
   std::lock_guard<std::mutex> NativeProgramLock(MDeviceGlobalInitializersMutex);
@@ -369,16 +376,17 @@ std::vector<ur_event_handle_t> context_impl::initializeDeviceGlobals(
           [&Adapter](const ur_event_handle_t &Event) {
             return get_event_info<info::event::command_execution_status>(
                        Event, Adapter) == info::event_command_status::complete;
-            return false;
           });
       // Release the removed events.
       for (auto EventIt = NewEnd; EventIt != InitEventsRef.end(); ++EventIt)
         Adapter->call<UrApiKind::urEventRelease>(*EventIt);
       // Remove them from the collection.
       InitEventsRef.erase(NewEnd, InitEventsRef.end());
       // If there are no more events, we can mark it as fully initialized.
-      if (InitEventsRef.empty())
+      if (InitEventsRef.empty()) {
         InitRef.MDeviceGlobalsFullyInitialized = true;
+        --MDeviceGlobalNotInitializedCnt;
+      }
       return InitEventsRef;
     } else if (InitRef.MDeviceGlobalsFullyInitialized) {
       // MDeviceGlobalsFullyInitialized could have been set while we were
@@ -387,7 +395,7 @@ std::vector<ur_event_handle_t> context_impl::initializeDeviceGlobals(
     }
 
     // There were no events and it was not set as fully initialized, so this is
-    // responsible for intializing the device globals.
+    // responsible for initializing the device globals.
     auto DeviceGlobals = InitRef.MBinImage->getDeviceGlobals();
     std::vector<std::string> DeviceGlobalIds;
     DeviceGlobalIds.reserve(DeviceGlobals.size());
@@ -402,6 +410,7 @@ std::vector<ur_event_handle_t> context_impl::initializeDeviceGlobals(
     // globals are trivially fully initialized and we can end early.
     if (DeviceGlobalEntries.empty()) {
       InitRef.MDeviceGlobalsFullyInitialized = true;
+      --MDeviceGlobalNotInitializedCnt;
       return {};
     }
 

sycl/source/detail/context_impl.hpp

@@ -307,10 +307,21 @@ class context_impl {
     std::vector<ur_event_handle_t> MDeviceGlobalInitEvents;
   };
 
-  std::map<std::pair<ur_program_handle_t, ur_device_handle_t>,
-           DeviceGlobalInitializer>
+  using HandleDevicePair = std::pair<ur_program_handle_t, ur_device_handle_t>;
+
+  struct HandleDevicePairHash {
+    std::size_t operator()(const HandleDevicePair &Key) const {
+      return std::hash<ur_program_handle_t>{}(Key.first) ^
+             std::hash<ur_device_handle_t>{}(Key.second);
+    }
+  };
+
+  std::unordered_map<HandleDevicePair, DeviceGlobalInitializer,
+                     HandleDevicePairHash>
       MDeviceGlobalInitializers;
   std::mutex MDeviceGlobalInitializersMutex;
+  // The number of device globals that have not been initialized yet.
+  std::atomic<size_t> MDeviceGlobalNotInitializedCnt = 0;
 
   // For device_global variables that are not used in any kernel code we still
   // allow copy operations on them. MDeviceGlobalUnregisteredData stores the

sycl/source/detail/graph_impl.cpp

@@ -1474,18 +1474,18 @@ void exec_graph_impl::populateURKernelUpdateStructs(
   ur_kernel_handle_t UrKernel = nullptr;
   auto Kernel = ExecCG.MSyclKernel;
   auto KernelBundleImplPtr = ExecCG.MKernelBundle;
-  std::shared_ptr<sycl::detail::kernel_impl> SyclKernelImpl = nullptr;
   const sycl::detail::KernelArgMask *EliminatedArgMask = nullptr;
 
-  if (auto SyclKernelImpl = KernelBundleImplPtr
-                                ? KernelBundleImplPtr->tryGetKernel(
-                                      ExecCG.MKernelName, KernelBundleImplPtr)
-                                : std::shared_ptr<kernel_impl>{nullptr}) {
-    UrKernel = SyclKernelImpl->getHandleRef();
-    EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
-  } else if (Kernel != nullptr) {
+  if (Kernel != nullptr) {
     UrKernel = Kernel->getHandleRef();
     EliminatedArgMask = Kernel->getKernelArgMask();
+  } else if (auto SyclKernelImpl =
+                 KernelBundleImplPtr
+                     ? KernelBundleImplPtr->tryGetKernel(ExecCG.MKernelName,
+                                                         KernelBundleImplPtr)
+                     : std::shared_ptr<kernel_impl>{nullptr}) {
+    UrKernel = SyclKernelImpl->getHandleRef();
+    EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
   } else {
     ur_program_handle_t UrProgram = nullptr;
     std::tie(UrKernel, std::ignore, EliminatedArgMask, UrProgram) =

sycl/source/detail/helpers.cpp

@@ -72,16 +72,16 @@ retrieveKernelBinary(const QueueImplPtr &Queue, const char *KernelName,
   const RTDeviceBinaryImage *DeviceImage = nullptr;
   ur_program_handle_t Program = nullptr;
   auto KernelBundleImpl = KernelCG->getKernelBundle();
-  if (auto SyclKernelImpl =
-          KernelBundleImpl
-              ? KernelBundleImpl->tryGetKernel(KernelName, KernelBundleImpl)
-              : std::shared_ptr<kernel_impl>{nullptr}) {
+  if (KernelCG->MSyclKernel != nullptr) {
+    DeviceImage = KernelCG->MSyclKernel->getDeviceImage()->get_bin_image_ref();
+    Program = KernelCG->MSyclKernel->getDeviceImage()->get_ur_program_ref();
+  } else if (auto SyclKernelImpl =
+                 KernelBundleImpl ? KernelBundleImpl->tryGetKernel(
+                                        KernelName, KernelBundleImpl)
+                                  : std::shared_ptr<kernel_impl>{nullptr}) {
     // Retrieve the device image from the kernel bundle.
     DeviceImage = SyclKernelImpl->getDeviceImage()->get_bin_image_ref();
     Program = SyclKernelImpl->getDeviceImage()->get_ur_program_ref();
-  } else if (KernelCG->MSyclKernel != nullptr) {
-    DeviceImage = KernelCG->MSyclKernel->getDeviceImage()->get_bin_image_ref();
-    Program = KernelCG->MSyclKernel->getDeviceImage()->get_ur_program_ref();
   } else {
     auto ContextImpl = Queue->getContextImplPtr();
     auto DeviceImpl = Queue->getDeviceImplPtr();

sycl/source/detail/program_manager/program_manager.cpp

@@ -765,35 +765,42 @@ ProgramManager::collectDependentDeviceImagesForVirtualFunctions(
     }
   }
 
-  while (!WorkList.empty()) {
-    std::string SetName = WorkList.front();
-    WorkList.pop();
+  if (!WorkList.empty()) {
+    // Guard read access to m_VFSet2BinImage:
+    // TODO: a better solution should be sought in the future, i.e. a different
+    // mutex than m_KernelIDsMutex, check lock check pattern, etc.
+    std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
 
-    // There could be more than one device image that uses the same set
-    // of virtual functions, or provides virtual funtions from the same
-    // set.
-    for (RTDeviceBinaryImage *BinImage : m_VFSet2BinImage[SetName]) {
-      // Here we can encounter both uses-virtual-functions-set and
-      // virtual-functions-set properties, but their handling is the same: we
-      // just grab all sets they reference and add them for consideration if
-      // we haven't done so already.
-      for (const sycl_device_binary_property &VFProp :
-           BinImage->getVirtualFunctions()) {
-        std::string StrValue = DeviceBinaryProperty(VFProp).asCString();
-        for (const auto &SetName : detail::split_string(StrValue, ',')) {
-          if (HandledSets.insert(SetName).second)
-            WorkList.push(SetName);
+    while (!WorkList.empty()) {
+      std::string SetName = WorkList.front();
+      WorkList.pop();
+
+      // There could be more than one device image that uses the same set
+      // of virtual functions, or provides virtual funtions from the same
+      // set.
+      for (RTDeviceBinaryImage *BinImage : m_VFSet2BinImage.at(SetName)) {
+        // Here we can encounter both uses-virtual-functions-set and
+        // virtual-functions-set properties, but their handling is the same: we
+        // just grab all sets they reference and add them for consideration if
+        // we haven't done so already.
+        for (const sycl_device_binary_property &VFProp :
+             BinImage->getVirtualFunctions()) {
+          std::string StrValue = DeviceBinaryProperty(VFProp).asCString();
+          for (const auto &SetName : detail::split_string(StrValue, ',')) {
+            if (HandledSets.insert(SetName).second)
+              WorkList.push(SetName);
+          }
         }
-      }
 
-      // TODO: Complete this part about handling of incompatible device images.
-      // If device image uses the same virtual function set, then we only
-      // link it if it is compatible.
-      // However, if device image provides virtual function set and it is
-      // incompatible, then we should link its "dummy" version to avoid link
-      // errors about unresolved external symbols.
-      if (doesDevSupportDeviceRequirements(Dev, *BinImage))
-        DeviceImagesToLink.insert(BinImage);
+        // TODO: Complete this part about handling of incompatible device
+        // images. If device image uses the same virtual function set, then we
+        // only link it if it is compatible. However, if device image provides
+        // virtual function set and it is incompatible, then we should link its
+        // "dummy" version to avoid link errors about unresolved external
+        // symbols.
+        if (doesDevSupportDeviceRequirements(Dev, *BinImage))
+          DeviceImagesToLink.insert(BinImage);
+      }
     }
   }
 
@@ -2163,8 +2170,14 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
 }
 
 void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
+  if (DeviceBinary->NumDeviceBinaries == 0)
+    return;
+  // Acquire lock to read and modify maps for kernel bundles
+  std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
+
   for (int I = 0; I < DeviceBinary->NumDeviceBinaries; I++) {
     sycl_device_binary RawImg = &(DeviceBinary->DeviceBinaries[I]);
+
     auto DevImgIt = m_DeviceImages.find(RawImg);
     if (DevImgIt == m_DeviceImages.end())
       continue;
@@ -2178,9 +2191,6 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     // Drop the kernel argument mask map
     m_EliminatedKernelArgMasks.erase(Img);
 
-    // Acquire lock to modify maps for kernel bundles
-    std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
-
     // Unmap the unique kernel IDs for the offload entries
     for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
          EntriesIt = EntriesIt->Increment()) {

sycl/source/detail/program_manager/program_manager.hpp

@@ -459,6 +459,7 @@ class ProgramManager {
 
   /// Keeps all device images we are refering to during program lifetime. Used
   /// for proper cleanup.
+  /// Access must be guarded by the m_KernelIDsMutex mutex.
   std::unordered_map<sycl_device_binary, RTDeviceBinaryImageUPtr>
       m_DeviceImages;
 
@@ -468,6 +469,7 @@ class ProgramManager {
 
   /// Caches list of device images that use or provide virtual functions from
   /// the same set. Used to simplify access.
+  /// Access must be guarded by the m_KernelIDsMutex mutex.
   std::unordered_map<std::string, std::set<RTDeviceBinaryImage *>>
       m_VFSet2BinImage;
 

sycl/source/detail/scheduler/commands.cpp

@@ -1997,16 +1997,16 @@ void instrumentationAddExtraKernelMetadata(
   std::mutex *KernelMutex = nullptr;
   const KernelArgMask *EliminatedArgMask = nullptr;
 
-  if (auto SyclKernelImpl = KernelBundleImplPtr
-                                ? KernelBundleImplPtr->tryGetKernel(
-                                      KernelName, KernelBundleImplPtr)
-                                : std::shared_ptr<kernel_impl>{nullptr}) {
-    EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
-    Program = SyclKernelImpl->getDeviceImage()->get_ur_program_ref();
-  } else if (nullptr != SyclKernel) {
+  if (nullptr != SyclKernel) {
     Program = SyclKernel->getProgramRef();
     if (!SyclKernel->isCreatedFromSource())
       EliminatedArgMask = SyclKernel->getKernelArgMask();
+  } else if (auto SyclKernelImpl =
+                 KernelBundleImplPtr ? KernelBundleImplPtr->tryGetKernel(
+                                           KernelName, KernelBundleImplPtr)
+                                     : std::shared_ptr<kernel_impl>{nullptr}) {
+    EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
+    Program = SyclKernelImpl->getDeviceImage()->get_ur_program_ref();
   } else if (Queue) {
     // NOTE: Queue can be null when kernel is directly enqueued to a command
     // buffer
@@ -2521,17 +2521,17 @@ getCGKernelInfo(const CGExecKernel &CommandGroup, ContextImplPtr ContextImpl,
   const KernelArgMask *EliminatedArgMask = nullptr;
   auto &KernelBundleImplPtr = CommandGroup.MKernelBundle;
 
-  if (auto SyclKernelImpl =
-          KernelBundleImplPtr
-              ? KernelBundleImplPtr->tryGetKernel(CommandGroup.MKernelName,
-                                                  KernelBundleImplPtr)
-              : std::shared_ptr<kernel_impl>{nullptr}) {
+  if (auto Kernel = CommandGroup.MSyclKernel; Kernel != nullptr) {
+    UrKernel = Kernel->getHandleRef();
+    EliminatedArgMask = Kernel->getKernelArgMask();
+  } else if (auto SyclKernelImpl =
+                 KernelBundleImplPtr
+                     ? KernelBundleImplPtr->tryGetKernel(
+                           CommandGroup.MKernelName, KernelBundleImplPtr)
+                     : std::shared_ptr<kernel_impl>{nullptr}) {
     UrKernel = SyclKernelImpl->getHandleRef();
     DeviceImageImpl = SyclKernelImpl->getDeviceImage();
     EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
-  } else if (auto Kernel = CommandGroup.MSyclKernel; Kernel != nullptr) {
-    UrKernel = Kernel->getHandleRef();
-    EliminatedArgMask = Kernel->getKernelArgMask();
   } else {
     ur_program_handle_t UrProgram = nullptr;
     std::tie(UrKernel, std::ignore, EliminatedArgMask, UrProgram) =
@@ -2678,18 +2678,7 @@ void enqueueImpKernel(
   std::shared_ptr<kernel_impl> SyclKernelImpl;
   std::shared_ptr<device_image_impl> DeviceImageImpl;
 
-  if ((SyclKernelImpl = KernelBundleImplPtr
-                            ? KernelBundleImplPtr->tryGetKernel(
-                                  KernelName, KernelBundleImplPtr)
-                            : std::shared_ptr<kernel_impl>{nullptr})) {
-    Kernel = SyclKernelImpl->getHandleRef();
-    DeviceImageImpl = SyclKernelImpl->getDeviceImage();
-
-    Program = DeviceImageImpl->get_ur_program_ref();
-
-    EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
-    KernelMutex = SyclKernelImpl->getCacheMutex();
-  } else if (nullptr != MSyclKernel) {
+  if (nullptr != MSyclKernel) {
     assert(MSyclKernel->get_info<info::kernel::context>() ==
            Queue->get_context());
     Kernel = MSyclKernel->getHandleRef();
@@ -2703,6 +2692,17 @@ void enqueueImpKernel(
     // their duplication in such cases.
     KernelMutex = &MSyclKernel->getNoncacheableEnqueueMutex();
     EliminatedArgMask = MSyclKernel->getKernelArgMask();
+  } else if ((SyclKernelImpl = KernelBundleImplPtr
+                                   ? KernelBundleImplPtr->tryGetKernel(
+                                         KernelName, KernelBundleImplPtr)
+                                   : std::shared_ptr<kernel_impl>{nullptr})) {
+    Kernel = SyclKernelImpl->getHandleRef();
+    DeviceImageImpl = SyclKernelImpl->getDeviceImage();
+
+    Program = DeviceImageImpl->get_ur_program_ref();
+
+    EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
+    KernelMutex = SyclKernelImpl->getCacheMutex();
   } else {
     std::tie(Kernel, KernelMutex, EliminatedArgMask, Program) =
         detail::ProgramManager::getInstance().getOrCreateKernel(
@@ -3511,6 +3511,18 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
     const AdapterPtr &Adapter = MQueue->getAdapter();
     if (MEvent != nullptr)
       MEvent->setHostEnqueueTime();
+    // User can specify explicit dependencies via depends_on call that we should
+    // honor here. It is very important for cross queue dependencies. We wait
+    // them explicitly since barrier w/o wait list waits for all commands
+    // submitted before and we can't add new dependencies to its wait list.
+    // Output event for wait operation is not requested since barrier is
+    // submitted immediately after and should synchronize it internally.
+    if (RawEvents.size()) {
+      auto Result = Adapter->call_nocheck<UrApiKind::urEnqueueEventsWait>(
+          MQueue->getHandleRef(), RawEvents.size(), &RawEvents[0], nullptr);
+      if (Result != UR_RESULT_SUCCESS)
+        return Result;
+    }
     if (auto Result =
             Adapter->call_nocheck<UrApiKind::urEnqueueEventsWaitWithBarrierExt>(
                 MQueue->getHandleRef(), &Properties, 0, nullptr, Event);
@@ -3545,6 +3557,12 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
     const AdapterPtr &Adapter = MQueue->getAdapter();
     if (MEvent != nullptr)
       MEvent->setHostEnqueueTime();
+    // User can specify explicit dependencies via depends_on call that we should
+    // honor here. It is very important for cross queue dependencies. Adding
+    // them to the barrier wait list since barrier w/ wait list waits only for
+    // the events provided in wait list and we can just extend the list.
+    UrEvents.insert(UrEvents.end(), RawEvents.begin(), RawEvents.end());
+
     if (auto Result =
             Adapter->call_nocheck<UrApiKind::urEnqueueEventsWaitWithBarrierExt>(
                 MQueue->getHandleRef(), &Properties, UrEvents.size(),

sycl/test-e2e/Graph/RecordReplay/barrier_multi_graph.cpp

@@ -5,6 +5,8 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 //
+// UNSUPPORTED: opencl
+// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
 
 #include "../graph_common.hpp"
 

sycl/test-e2e/Graph/RecordReplay/barrier_multi_queue.cpp

@@ -5,6 +5,8 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 //
+// UNSUPPORTED: opencl
+// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
 
 #include "../graph_common.hpp"
 

sycl/test-e2e/Graph/RecordReplay/ext_oneapi_enqueue_functions.cpp

@@ -7,6 +7,9 @@
 
 // Tests the enqueue free function kernel shortcuts.
 
+// UNSUPPORTED: opencl
+// UNSUPPORTED-INTENDED: USM memcpy command not supported for OpenCL
+
 #include "../graph_common.hpp"
 #include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
 #include <sycl/properties/all_properties.hpp>