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();

unified-runtime/source/adapters/offload/enqueue.cpp

@@ -20,30 +20,43 @@
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     ur_queue_handle_t hQueue, ur_kernel_handle_t hKernel, uint32_t workDim,
     const size_t *pGlobalWorkOffset, const size_t *pGlobalWorkSize,
-    const size_t *pLocalWorkSize, uint32_t numEventsInWaitList,
-    const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
+    const size_t *pLocalWorkSize, uint32_t, const ur_kernel_launch_property_t *,
+    uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList,
+    ur_event_handle_t *phEvent) {
   // Ignore wait list for now
   (void)numEventsInWaitList;
   (void)phEventWaitList;
   //
 
   (void)pGlobalWorkOffset;
-  (void)pLocalWorkSize;
 
-  if (workDim == 1) {
-    std::cerr
-        << "UR Offload adapter only supports 1d kernel launches at the moment";
-    return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
+  size_t GlobalSize[3] = {1, 1, 1};
+  for (uint32_t I = 0; I < workDim; I++) {
+    GlobalSize[I] = pGlobalWorkSize[I];
+  }
+
+  // TODO: We default to 1, 1, 1 here. In future if pLocalWorkSize is not
+  // specified, we should pick the "best" one
+  size_t GroupSize[3] = {1, 1, 1};
+  if (pLocalWorkSize) {
+    for (uint32_t I = 0; I < workDim; I++) {
+      GroupSize[I] = pLocalWorkSize[I];
+    }
+  }
+
+  if (GroupSize[0] > GlobalSize[0] || GroupSize[1] > GlobalSize[1] ||
+      GroupSize[2] > GlobalSize[2]) {
+    return UR_RESULT_ERROR_INVALID_WORK_GROUP_SIZE;
   }
 
   ol_kernel_launch_size_args_t LaunchArgs;
   LaunchArgs.Dimensions = workDim;
-  LaunchArgs.NumGroupsX = pGlobalWorkSize[0];
-  LaunchArgs.NumGroupsY = 1;
-  LaunchArgs.NumGroupsZ = 1;
-  LaunchArgs.GroupSizeX = 1;
-  LaunchArgs.GroupSizeY = 1;
-  LaunchArgs.GroupSizeZ = 1;
+  LaunchArgs.NumGroupsX = GlobalSize[0] / GroupSize[0];
+  LaunchArgs.NumGroupsY = GlobalSize[1] / GroupSize[1];
+  LaunchArgs.NumGroupsZ = GlobalSize[2] / GroupSize[2];
+  LaunchArgs.GroupSizeX = GroupSize[0];
+  LaunchArgs.GroupSizeY = GroupSize[1];
+  LaunchArgs.GroupSizeZ = GroupSize[2];
   LaunchArgs.DynSharedMemory = 0;
 
   ol_event_handle_t EventOut;

unified-runtime/source/adapters/offload/ur_interface_loader.cpp

@@ -381,25 +381,12 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEnqueueExpProcAddrTable(
     return result;
   }
 
-  pDdiTable->pfnCooperativeKernelLaunchExp = nullptr;
   pDdiTable->pfnTimestampRecordingExp = nullptr;
   pDdiTable->pfnNativeCommandExp = nullptr;
 
   return UR_RESULT_SUCCESS;
 }
 
-UR_DLLEXPORT ur_result_t UR_APICALL urGetKernelExpProcAddrTable(
-    ur_api_version_t version, ur_kernel_exp_dditable_t *pDdiTable) {
-  auto result = validateProcInputs(version, pDdiTable);
-  if (UR_RESULT_SUCCESS != result) {
-    return result;
-  }
-
-  pDdiTable->pfnSuggestMaxCooperativeGroupCountExp = nullptr;
-
-  return UR_RESULT_SUCCESS;
-}
-
 UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramExpProcAddrTable(
     ur_api_version_t version, ur_program_exp_dditable_t *pDdiTable) {
   auto result = validateProcInputs(version, pDdiTable);
@@ -424,7 +411,6 @@ UR_DLLEXPORT ur_result_t UR_APICALL urAllAddrTable(ur_api_version_t version,
   urGetEnqueueExpProcAddrTable(version, &pDdiTable->EnqueueExp);
   urGetEventProcAddrTable(version, &pDdiTable->Event);
   urGetKernelProcAddrTable(version, &pDdiTable->Kernel);
-  urGetKernelExpProcAddrTable(version, &pDdiTable->KernelExp);
   urGetMemProcAddrTable(version, &pDdiTable->Mem);
   urGetPhysicalMemProcAddrTable(version, &pDdiTable->PhysicalMem);
   urGetPlatformProcAddrTable(version, &pDdiTable->Platform);

unified-runtime/source/adapters/opencl/event.cpp

@@ -149,12 +149,25 @@ UR_APIEXPORT ur_result_t UR_APICALL urEventRetain(ur_event_handle_t hEvent) {
 
 UR_APIEXPORT ur_result_t UR_APICALL
 urEventWait(uint32_t numEvents, const ur_event_handle_t *phEventWaitList) {
-  std::vector<cl_event> CLEvents(numEvents);
+  ur_context_handle_t hContext = phEventWaitList[0]->Context;
+  std::vector<cl_event> CLEvents;
+  CLEvents.reserve(numEvents);
+
+  // clWaitForEvents can only be called on events from the same context.
+  // If the events are from different contexts, we need to wait for each
+  // set of events separately.
   for (uint32_t i = 0; i < numEvents; i++) {
-    CLEvents[i] = phEventWaitList[i]->CLEvent;
+    if (phEventWaitList[i]->Context != hContext) {
+      CL_RETURN_ON_FAILURE(clWaitForEvents(CLEvents.size(), CLEvents.data()));
+      CLEvents.clear();
+    }
+
+    CLEvents.push_back(phEventWaitList[i]->CLEvent);
+    hContext = phEventWaitList[i]->Context;
+  }
+  if (CLEvents.size()) {
+    CL_RETURN_ON_FAILURE(clWaitForEvents(CLEvents.size(), CLEvents.data()));
   }
-  cl_int RetErr = clWaitForEvents(numEvents, CLEvents.data());
-  CL_RETURN_ON_FAILURE(RetErr);
   return UR_RESULT_SUCCESS;
 }