Merge remote-tracking branch 'origin/sycl' into llvmspirv_pulldown

Author: jsji

devops/containers/ubuntu2404_base.Dockerfile

@@ -4,6 +4,11 @@ ENV DEBIAN_FRONTEND=noninteractive
 
 USER root
 
+# Configure LLVM nightly repo
+RUN apt-get update -qq && apt-get install --no-install-recommends -yqq curl ca-certificates
+RUN curl -sSL https://apt.llvm.org/llvm-snapshot.gpg.key -o /etc/apt/trusted.gpg.d/apt.llvm.org.asc
+RUN echo 'deb http://apt.llvm.org/noble/ llvm-toolchain-noble main' > /etc/apt/sources.list.d/llvm.list
+
 # Install SYCL prerequisites
 COPY scripts/install_build_tools.sh /install.sh
 RUN /install.sh

sycl/doc/extensions/proposed/sycl_ext_intel_event_mode.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_intel_event_mode.asciidoc

@@ -52,11 +52,12 @@ This extension also depends on the following other SYCL extensions:
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
+This is an experimental extension specification, intended to provide early
+access to features and gather community feedback.  Interfaces defined in this
+specification are implemented in {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
 
 
 == Overview

sycl/source/detail/kernel_name_based_cache_t.hpp

@@ -74,6 +74,10 @@ struct FastKernelSubcacheT {
 struct KernelNameBasedCacheT {
   FastKernelSubcacheT FastKernelSubcache;
   std::optional<bool> UsesAssert;
+  // Implicit local argument position is represented by an optional int, this
+  // uses another optional on top of that to represent lazy initialization of
+  // the cached value.
+  std::optional<std::optional<int>> ImplicitLocalArgPos;
 };
 
 } // namespace detail

sycl/source/detail/program_manager/program_manager.cpp

@@ -1841,12 +1841,24 @@ void ProgramManager::cacheKernelImplicitLocalArg(RTDeviceBinaryImage &Img) {
     }
 }
 
-std::optional<int>
-ProgramManager::kernelImplicitLocalArgPos(KernelNameStrRefT KernelName) const {
-  auto it = m_KernelImplicitLocalArgPos.find(KernelName);
-  if (it != m_KernelImplicitLocalArgPos.end())
-    return it->second;
-  return {};
+std::optional<int> ProgramManager::kernelImplicitLocalArgPos(
+    KernelNameStrRefT KernelName,
+    KernelNameBasedCacheT *KernelNameBasedCachePtr) const {
+  auto getLocalArgPos = [&]() -> std::optional<int> {
+    auto it = m_KernelImplicitLocalArgPos.find(KernelName);
+    if (it != m_KernelImplicitLocalArgPos.end())
+      return it->second;
+    return {};
+  };
+
+  if (!KernelNameBasedCachePtr)
+    return getLocalArgPos();
+  std::optional<std::optional<int>> &ImplicitLocalArgPos =
+      KernelNameBasedCachePtr->ImplicitLocalArgPos;
+  if (!ImplicitLocalArgPos.has_value()) {
+    ImplicitLocalArgPos = getLocalArgPos();
+  }
+  return ImplicitLocalArgPos.value();
 }
 
 static bool isBfloat16DeviceLibImage(sycl_device_binary RawImg,

sycl/source/detail/program_manager/program_manager.hpp

@@ -373,8 +373,9 @@ class ProgramManager {
 
   SanitizerType kernelUsesSanitizer() const { return m_SanitizerFoundInImage; }
 
-  std::optional<int>
-  kernelImplicitLocalArgPos(KernelNameStrRefT KernelName) const;
+  std::optional<int> kernelImplicitLocalArgPos(
+      KernelNameStrRefT KernelName,
+      KernelNameBasedCacheT *KernelNameBasedCachePtr) const;
 
   std::set<RTDeviceBinaryImage *>
   getRawDeviceImages(const std::vector<kernel_id> &KernelIDs);

sycl/source/detail/scheduler/commands.cpp

@@ -2389,8 +2389,9 @@ static ur_result_t SetKernelParamsAndLaunch(
     const std::function<void *(Requirement *Req)> &getMemAllocationFunc,
     bool IsCooperative, bool KernelUsesClusterLaunch,
     uint32_t WorkGroupMemorySize, const RTDeviceBinaryImage *BinImage,
-    KernelNameStrRefT KernelName, void *KernelFuncPtr = nullptr,
-    int KernelNumArgs = 0,
+    KernelNameStrRefT KernelName,
+    KernelNameBasedCacheT *KernelNameBasedCachePtr,
+    void *KernelFuncPtr = nullptr, int KernelNumArgs = 0,
     detail::kernel_param_desc_t (*KernelParamDescGetter)(int) = nullptr,
     bool KernelHasSpecialCaptures = true) {
   const AdapterPtr &Adapter = Queue.getAdapter();
@@ -2437,7 +2438,8 @@ static ur_result_t SetKernelParamsAndLaunch(
   }
 
   std::optional<int> ImplicitLocalArg =
-      ProgramManager::getInstance().kernelImplicitLocalArgPos(KernelName);
+      ProgramManager::getInstance().kernelImplicitLocalArgPos(
+          KernelName, KernelNameBasedCachePtr);
   // Set the implicit local memory buffer to support
   // get_work_group_scratch_memory. This is for backend not supporting
   // CUDA-style local memory setting. Note that we may have -1 as a position,
@@ -2752,8 +2754,8 @@ void enqueueImpKernel(
         *Queue, Args, DeviceImageImpl, Kernel, NDRDesc, EventsWaitList,
         OutEventImpl, EliminatedArgMask, getMemAllocationFunc,
         KernelIsCooperative, KernelUsesClusterLaunch, WorkGroupMemorySize,
-        BinImage, KernelName, KernelFuncPtr, KernelNumArgs,
-        KernelParamDescGetter, KernelHasSpecialCaptures);
+        BinImage, KernelName, KernelNameBasedCachePtr, KernelFuncPtr,
+        KernelNumArgs, KernelParamDescGetter, KernelHasSpecialCaptures);
   }
   if (UR_RESULT_SUCCESS != Error) {
     // If we have got non-success error code, let's analyze it to emit nice

sycl/source/feature_test.hpp.in

@@ -120,6 +120,7 @@ inline namespace _V1 {
 // In progress yet
 #define SYCL_EXT_ONEAPI_ATOMIC16 0
 #define SYCL_KHR_DEFAULT_CONTEXT 1
+#define SYCL_EXT_INTEL_EVENT_MODE 1
 
 #ifndef __has_include
 #define __has_include(x) 0

sycl/test-e2e/Basic/multi_context_wait.cpp

@@ -0,0 +1,100 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include <iostream>
+#include <vector>
+
+std::vector<sycl::event> submit_dependencies(sycl::queue q1, sycl::queue q2,
+                                             int *mem1, int *mem2) {
+  int delay_ops = 1024 * 1024;
+  auto delay = [=] {
+    volatile int value = delay_ops;
+    while (--value)
+      ;
+  };
+
+  auto ev1 =
+      q1.parallel_for(sycl::range<1>(1024), [=]([[maybe_unused]] auto u) {
+        delay();
+        mem1[u.get_id()] = 1;
+      });
+  auto ev2 =
+      q2.parallel_for(sycl::range<1>(1024), [=]([[maybe_unused]] auto u) {
+        delay();
+        mem2[u.get_id()] = 2;
+      });
+
+  return {ev1, ev2};
+}
+
+void test_host_task() {
+  sycl::context c1{};
+  sycl::context c2{};
+
+  sycl::queue q1(c1, sycl::default_selector_v);
+  sycl::queue q2(c2, sycl::default_selector_v);
+
+  auto mem1 = sycl::malloc_host<int>(1024, q1);
+  auto mem2 = sycl::malloc_host<int>(1024, q2);
+
+  auto events = submit_dependencies(q1, q2, mem1, mem2);
+
+  q2.submit([&](sycl::handler &cgh) {
+    cgh.depends_on(events[0]);
+    cgh.depends_on(events[1]);
+    cgh.host_task([=]() {
+      for (int i = 0; i < 1024; i++) {
+        assert(mem1[i] == 1);
+        assert(mem2[i] == 2);
+      }
+    });
+  });
+
+  q2.wait();
+
+  sycl::free(mem1, c1);
+  sycl::free(mem2, c2);
+}
+
+void test_kernel() {
+  sycl::context c1{};
+  sycl::context c2{};
+
+  sycl::queue q1(c1, sycl::default_selector_v);
+  sycl::queue q2(c2, sycl::default_selector_v);
+
+  auto mem1 = sycl::malloc_device<int>(1024, q1);
+  auto mem2 = sycl::malloc_device<int>(1024, q2);
+
+  auto events = submit_dependencies(q1, q2, mem1, mem2);
+
+  q1.submit([&](sycl::handler &cgh) {
+    cgh.depends_on(events[0]);
+    cgh.depends_on(events[1]);
+    cgh.parallel_for(sycl::range<1>(1024),
+                     [=](auto item) { assert(mem1[item.get_id()] == 1); });
+  });
+
+  q2.submit([&](sycl::handler &cgh) {
+    cgh.depends_on(events[0]);
+    cgh.depends_on(events[1]);
+    cgh.parallel_for(sycl::range<1>(1024),
+                     [=](auto item) { assert(mem2[item.get_id()] == 2); });
+  });
+
+  q1.wait();
+  q2.wait();
+
+  sycl::free(mem1, c1);
+  sycl::free(mem2, c2);
+}
+
+int main() {
+  test_host_task();
+  test_kernel();
+
+  return 0;
+}

unified-runtime/source/adapters/level_zero/v2/memory.hpp

@@ -247,7 +247,7 @@ struct ur_mem_image_t : ur_object {
   ZeStruct<ze_image_desc_t> zeImageDesc;
 };
 
-struct ur_mem_handle_t_ {
+struct ur_mem_handle_t_ : ur::handle_base<ur::level_zero::ddi_getter> {
   template <typename T, typename... Args>
   static ur_mem_handle_t_ *create(Args &&...args) {
     return new ur_mem_handle_t_(std::in_place_type<T>,
@@ -293,7 +293,8 @@ struct ur_mem_handle_t_ {
 private:
   template <typename T, typename... Args>
   ur_mem_handle_t_(std::in_place_type_t<T>, Args &&...args)
-      : mem(std::in_place_type<T>, std::forward<Args>(args)...) {}
+      : ur::handle_base<ur::level_zero::ddi_getter>(),
+        mem(std::in_place_type<T>, std::forward<Args>(args)...) {}
 
   std::variant<ur_usm_handle_t, ur_integrated_buffer_handle_t,
                ur_discrete_buffer_handle_t, ur_shared_buffer_handle_t,

unified-runtime/source/adapters/level_zero/v2/queue_immediate_in_order.cpp

@@ -1026,9 +1026,14 @@ ur_result_t ur_queue_immediate_in_order_t::enqueueCommandBufferExp(
   ur_event_handle_t executionEvent =
       hCommandBuffer->getExecutionEventUnlocked();
 
+  if (executionEvent != nullptr) {
+    ZE2UR_CALL(zeEventHostSynchronize,
+               (executionEvent->getZeEvent(), UINT64_MAX));
+  }
+
   UR_CALL(enqueueGenericCommandListsExp(
       1, &commandBufferCommandList, phEvent, numEventsInWaitList,
-      phEventWaitList, UR_COMMAND_ENQUEUE_COMMAND_BUFFER_EXP, executionEvent));
+      phEventWaitList, UR_COMMAND_ENQUEUE_COMMAND_BUFFER_EXP, nullptr));
   UR_CALL(hCommandBuffer->registerExecutionEventUnlocked(*phEvent));
   if (internalEvent != nullptr) {
     internalEvent->release();