Merge remote-tracking branch 'upstream/sycl' into bump-cb

Author: PatKamin

clang/test/Driver/clang-offload-extract.c

@@ -64,8 +64,8 @@ __declspec(align(sizeof(void*) * 2))
 const void* padding[2] = {0, 0};
 
 #ifdef _WIN32
-char __start_omp_offloading_entries = 1;
-char __stop_omp_offloading_entries = 1;
+char __start_llvm_offload_entries = 1;
+char __stop_llvm_offload_entries = 1;
 #endif
 
 void __tgt_register_lib(void *desc) {}

clang/test/Driver/clang-offload-wrapper.c

@@ -129,7 +129,7 @@
 // CHECK-IR: [[ENTBEGIN:@.+]] = external hidden constant [[ENTTY]]
 // CHECK-IR: [[ENTEND:@.+]] = external hidden constant [[ENTTY]]
 
-// CHECK-IR: [[DUMMY:@.+]] = hidden constant [0 x [[ENTTY]]] zeroinitializer, section "omp_offloading_entries"
+// CHECK-IR: [[DUMMY:@.+]] = hidden constant [0 x [[ENTTY]]] zeroinitializer, section "llvm_offload_entries"
 
 // CHECK-IR: [[OMP_BIN:@.+]] = internal unnamed_addr constant [[OMP_BINTY:\[[0-9]+ x i8\]]] c"Content of device file3{{.+}}"
 // CHECK-IR: [[OMP_INFO:@.+]] = internal local_unnamed_addr constant [2 x i64] [i64 ptrtoint (ptr [[OMP_BIN]] to i64), i64 24], section ".tgtimg", align 16

clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp

@@ -1006,9 +1006,9 @@ class BinaryWrapper {
   /// library. It is defined as follows
   ///
   /// __attribute__((visibility("hidden")))
-  /// extern __tgt_offload_entry *__start_omp_offloading_entries;
+  /// extern __tgt_offload_entry *__start_llvm_offload_entries;
   /// __attribute__((visibility("hidden")))
-  /// extern __tgt_offload_entry *__stop_omp_offloading_entries;
+  /// extern __tgt_offload_entry *__stop_llvm_offload_entries;
   ///
   /// static const char Image0[] = { <Bufs.front() contents> };
   ///  ...
@@ -1018,23 +1018,23 @@ class BinaryWrapper {
   ///   {
   ///     Image0,                            /*ImageStart*/
   ///     Image0 + sizeof(Image0),           /*ImageEnd*/
-  ///     __start_omp_offloading_entries,    /*EntriesBegin*/
-  ///     __stop_omp_offloading_entries      /*EntriesEnd*/
+  ///     __start_llvm_offload_entries,      /*EntriesBegin*/
+  ///     __stop_llvm_offload_entries        /*EntriesEnd*/
   ///   },
   ///   ...
   ///   {
   ///     ImageN,                            /*ImageStart*/
   ///     ImageN + sizeof(ImageN),           /*ImageEnd*/
-  ///     __start_omp_offloading_entries,    /*EntriesBegin*/
-  ///     __stop_omp_offloading_entries      /*EntriesEnd*/
+  ///     __start_llvm_offload_entries,      /*EntriesBegin*/
+  ///     __stop_llvm_offload_entries        /*EntriesEnd*/
   ///   }
   /// };
   ///
   /// static const __tgt_bin_desc BinDesc = {
   ///   sizeof(Images) / sizeof(Images[0]),  /*NumDeviceImages*/
   ///   Images,                              /*DeviceImages*/
-  ///   __start_omp_offloading_entries,      /*HostEntriesBegin*/
-  ///   __stop_omp_offloading_entries        /*HostEntriesEnd*/
+  ///   __start_llvm_offload_entries,        /*HostEntriesBegin*/
+  ///   __stop_llvm_offload_entries          /*HostEntriesEnd*/
   /// };
   ///
   /// Global variable that represents BinDesc is returned.
@@ -1049,24 +1049,24 @@ class BinaryWrapper {
       // Create external begin/end symbols for the offload entries table.
       auto *EntriesStart = new GlobalVariable(
           M, getEntryTy(), /*isConstant*/ true, GlobalValue::ExternalLinkage,
-          /*Initializer*/ nullptr, "__start_omp_offloading_entries");
+          /*Initializer*/ nullptr, "__start_llvm_offload_entries");
       EntriesStart->setVisibility(GlobalValue::HiddenVisibility);
       auto *EntriesStop = new GlobalVariable(
           M, getEntryTy(), /*isConstant*/ true, GlobalValue::ExternalLinkage,
-          /*Initializer*/ nullptr, "__stop_omp_offloading_entries");
+          /*Initializer*/ nullptr, "__stop_llvm_offload_entries");
       EntriesStop->setVisibility(GlobalValue::HiddenVisibility);
 
       // We assume that external begin/end symbols that we have created above
       // will be defined by the linker. But linker will do that only if linker
-      // inputs have section with "omp_offloading_entries" name which is not
+      // inputs have section with "llvm_offload_entries" name which is not
       // guaranteed. So, we just create dummy zero sized object in the offload
       // entries section to force linker to define those symbols.
       auto *DummyInit =
           ConstantAggregateZero::get(ArrayType::get(getEntryTy(), 0u));
       auto *DummyEntry = new GlobalVariable(
           M, DummyInit->getType(), true, GlobalVariable::ExternalLinkage,
           DummyInit, "__dummy.omp_offloading.entry");
-      DummyEntry->setSection("omp_offloading_entries");
+      DummyEntry->setSection("llvm_offload_entries");
       DummyEntry->setVisibility(GlobalValue::HiddenVisibility);
 
       EntriesB = EntriesStart;

devops/scripts/benchmarks/utils/utils.py

@@ -49,31 +49,37 @@ def run(
             command = command.split()
 
         env = os.environ.copy()
-
         for ldlib in ld_library:
             if os.path.isdir(ldlib):
-                env["LD_LIBRARY_PATH"] = (
-                    ldlib + os.pathsep + env.get("LD_LIBRARY_PATH", "")
+                env_vars["LD_LIBRARY_PATH"] = os.pathsep.join(
+                    filter(None, [ldlib, env_vars.get("LD_LIBRARY_PATH", "")])
                 )
             else:
                 log.warning(f"LD_LIBRARY_PATH component does not exist: {ldlib}")
 
         # order is important, we want provided sycl rt libraries to be first
         if add_sycl:
             sycl_bin_path = os.path.join(options.sycl, "bin")
-            env["PATH"] = sycl_bin_path + os.pathsep + env.get("PATH", "")
+            env_vars["PATH"] = os.pathsep.join(
+                filter(None, [sycl_bin_path, env_vars.get("PATH", "")])
+            )
             sycl_lib_path = os.path.join(options.sycl, "lib")
-            env["LD_LIBRARY_PATH"] = (
-                sycl_lib_path + os.pathsep + env.get("LD_LIBRARY_PATH", "")
+            env_vars["LD_LIBRARY_PATH"] = os.pathsep.join(
+                filter(None, [sycl_lib_path, env_vars.get("LD_LIBRARY_PATH", "")])
             )
 
-        env.update(env_vars)
-
         command_str = " ".join(command)
         env_str = " ".join(f"{key}={value}" for key, value in env_vars.items())
         full_command_str = f"{env_str} {command_str}".strip()
         log.debug(f"Running: {full_command_str}")
 
+        for key, value in env_vars.items():
+            # Only PATH and LD_LIBRARY_PATH should be prepended to existing values
+            if key in ("PATH", "LD_LIBRARY_PATH") and (old := env.get(key)):
+                env[key] = os.pathsep.join([value, old])
+            else:
+                env[key] = value
+
         # Normalize input to bytes if it's a str
         if isinstance(input, str):
             input_bytes = input.encode()

llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp

@@ -66,11 +66,12 @@ bool isModuleUsingTsan(const Module &M) {
 // Optional.
 // Otherwise, it returns an Optional containing a list of reached
 // SPIR kernel function's names.
-static std::optional<std::vector<StringRef>>
-traverseCGToFindSPIRKernels(const Function *StartingFunction) {
+static std::optional<std::vector<StringRef>> traverseCGToFindSPIRKernels(
+    const std::vector<Function *> &StartingFunctionVec) {
   std::queue<const Function *> FunctionsToVisit;
   std::unordered_set<const Function *> VisitedFunctions;
-  FunctionsToVisit.push(StartingFunction);
+  for (const Function *FPtr : StartingFunctionVec)
+    FunctionsToVisit.push(FPtr);
   std::vector<StringRef> KernelNames;
 
   while (!FunctionsToVisit.empty()) {
@@ -106,13 +107,20 @@ traverseCGToFindSPIRKernels(const Function *StartingFunction) {
   return {std::move(KernelNames)};
 }
 
-static std::vector<StringRef> getKernelNamesUsingAssert(const Module &M) {
-  auto *DevicelibAssertFailFunction = M.getFunction("__devicelib_assert_fail");
-  if (!DevicelibAssertFailFunction)
+static std::vector<StringRef>
+getKernelNamesUsingSpecialFunctions(const Module &M,
+                                    const std::vector<StringRef> &FNames) {
+  std::vector<Function *> SpecialFunctionVec;
+  for (const auto Fn : FNames) {
+    Function *FPtr = M.getFunction(Fn);
+    if (FPtr)
+      SpecialFunctionVec.push_back(FPtr);
+  }
+
+  if (SpecialFunctionVec.size() == 0)
     return {};
 
-  auto TraverseResult =
-      traverseCGToFindSPIRKernels(DevicelibAssertFailFunction);
+  auto TraverseResult = traverseCGToFindSPIRKernels(SpecialFunctionVec);
 
   if (TraverseResult.has_value())
     return std::move(*TraverseResult);
@@ -442,7 +450,9 @@ PropSetRegTy computeModuleProperties(const Module &M,
       PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "optLevel", OptLevel);
   }
   {
-    std::vector<StringRef> FuncNames = getKernelNamesUsingAssert(M);
+    std::vector<StringRef> AssertFuncNames{"__devicelib_assert_fail"};
+    std::vector<StringRef> FuncNames =
+        getKernelNamesUsingSpecialFunctions(M, AssertFuncNames);
     for (const StringRef &FName : FuncNames)
       PropSet.add(PropSetRegTy::SYCL_ASSERT_USED, FName, true);
   }

sycl/CMakeLists.txt

@@ -363,6 +363,13 @@ if (NOT WIN32)
     RENAME "libsycl.so.${SYCL_VERSION_STRING}-gdb.py"
     DESTINATION "lib${LLVM_LIBDIR_SUFFIX}/"
     COMPONENT sycl-headers-extras)
+  if (SYCL_ENABLE_MAJOR_RELEASE_PREVIEW_LIB)
+    install(FILES
+      "${CMAKE_CURRENT_SOURCE_DIR}/gdb/libsycl.so-gdb.py"
+      RENAME "libsycl-preview.so.${SYCL_VERSION_STRING}-gdb.py"
+      DESTINATION "lib${LLVM_LIBDIR_SUFFIX}/"
+      COMPONENT sycl-headers-extras)
+  endif()
 endif()
 
 if(SYCL_ENABLE_XPTI_TRACING AND

sycl/include/sycl/detail/kernel_launch_helper.hpp

@@ -12,8 +12,12 @@
 #include <sycl/detail/compile_time_kernel_info.hpp>
 #include <sycl/detail/helpers.hpp>
 #include <sycl/detail/is_device_copyable.hpp>
+#include <sycl/detail/type_traits.hpp>
 #include <sycl/ext/intel/experimental/fp_control_kernel_properties.hpp>
 #include <sycl/ext/intel/experimental/kernel_execution_properties.hpp>
+#include <sycl/ext/oneapi/experimental/cluster_group_prop.hpp>
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+#include <sycl/ext/oneapi/experimental/use_root_sync_prop.hpp>
 #include <sycl/ext/oneapi/experimental/virtual_functions.hpp>
 #include <sycl/ext/oneapi/kernel_properties/properties.hpp>
 #include <sycl/ext/oneapi/work_group_scratch_memory.hpp>
@@ -253,23 +257,130 @@ struct KernelWrapper<
   }
 }; // KernelWrapper struct
 
-struct KernelLaunchPropertyWrapper {
-  template <typename KernelName, typename PropertyProcessor,
-            typename KernelType>
-  static void parseProperties([[maybe_unused]] PropertyProcessor h,
-                              [[maybe_unused]] const KernelType &KernelFunc) {
-#ifndef __SYCL_DEVICE_ONLY__
-    // If there are properties provided by get method then process them.
-    if constexpr (ext::oneapi::experimental::detail::
-                      HasKernelPropertiesGetMethod<const KernelType &>::value) {
-
-      h->template processProperties<
-          detail::CompileTimeKernelInfo<KernelName>.IsESIMD>(
-          KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
+// This namespace encapsulates everything related to parsing kernel launch
+// properties.
+inline namespace kernel_launch_properties_v1 {
+
+template <typename key, typename = void> struct MarshalledProperty;
+
+// Generic implementation for runtime properties.
+template <typename PropertyTy>
+struct MarshalledProperty<
+    PropertyTy,
+    std::enable_if_t<!std::is_empty_v<PropertyTy> &&
+                     std::is_same_v<PropertyTy, typename PropertyTy::key_t>>> {
+  std::optional<PropertyTy> MProperty;
+
+  template <typename InputPropertyTy>
+  MarshalledProperty(const InputPropertyTy &Props) {
+    (void)Props;
+    if constexpr (InputPropertyTy::template has_property<PropertyTy>())
+      MProperty = Props.template get_property<PropertyTy>();
+  }
+
+  MarshalledProperty() = default;
+};
+
+// Generic implementation for properties with non-template value_t.
+template <typename PropertyTy>
+struct MarshalledProperty<PropertyTy,
+                          std::void_t<typename PropertyTy::value_t>> {
+  bool MPresent = false;
+
+  template <typename InputPropertyTy>
+  MarshalledProperty(const InputPropertyTy &) {
+    using namespace sycl::ext::oneapi::experimental;
+    MPresent = InputPropertyTy::template has_property<
+        sycl::ext::oneapi::experimental::use_root_sync_key>();
+  }
+
+  MarshalledProperty() = default;
+};
+
+// Specialization for work group progress property.
+template <typename PropertyTy>
+struct MarshalledProperty<
+    PropertyTy,
+    std::enable_if_t<check_type_in_v<
+        PropertyTy, sycl::ext::oneapi::experimental::work_group_progress_key,
+        sycl::ext::oneapi::experimental::sub_group_progress_key,
+        sycl::ext::oneapi::experimental::work_item_progress_key>>> {
+
+  using forward_progress_guarantee =
+      sycl::ext::oneapi::experimental::forward_progress_guarantee;
+  using execution_scope = sycl::ext::oneapi::experimental::execution_scope;
+
+  std::optional<forward_progress_guarantee> MFPGuarantee;
+  std::optional<execution_scope> MFPCoordinationScope;
+
+  template <typename InputPropertyTy>
+  MarshalledProperty(const InputPropertyTy &Props) {
+    (void)Props;
+
+    if constexpr (InputPropertyTy::template has_property<PropertyTy>()) {
+      MFPGuarantee = Props.template get_property<PropertyTy>().guarantee;
+      MFPCoordinationScope =
+          Props.template get_property<PropertyTy>().coordinationScope;
     }
-#endif
   }
-}; // KernelLaunchPropertyWrapper struct
+
+  MarshalledProperty() = default;
+};
+
+template <typename... keys> struct PropsHolder : MarshalledProperty<keys>... {
+  bool MEmpty = true;
+
+  template <typename PropertiesT,
+            class = typename std::enable_if_t<
+                ext::oneapi::experimental::is_property_list_v<PropertiesT>>>
+  PropsHolder(PropertiesT Props)
+      : MarshalledProperty<keys>(Props)...,
+        MEmpty(((!PropertiesT::template has_property<keys>() && ...))) {}
+
+  PropsHolder() = default;
+
+  constexpr bool isEmpty() const { return MEmpty; }
+
+  template <typename PropertyCastKey> constexpr auto get() const {
+    return static_cast<const MarshalledProperty<PropertyCastKey> *>(this);
+  }
+};
+
+using KernelPropertyHolderStructTy =
+    PropsHolder<sycl::ext::oneapi::experimental::work_group_scratch_size,
+                sycl::ext::intel::experimental::cache_config_key,
+                sycl::ext::oneapi::experimental::use_root_sync_key,
+                sycl::ext::oneapi::experimental::work_group_progress_key,
+                sycl::ext::oneapi::experimental::sub_group_progress_key,
+                sycl::ext::oneapi::experimental::work_item_progress_key,
+                sycl::ext::oneapi::experimental::cuda::cluster_size_key<1>,
+                sycl::ext::oneapi::experimental::cuda::cluster_size_key<2>,
+                sycl::ext::oneapi::experimental::cuda::cluster_size_key<3>>;
+
+/// Note: it is important that this function *does not* depend on kernel
+/// name or kernel type, because then it will be instantiated for every
+/// kernel, even though body of those instantiated functions could be almost
+/// the same, thus unnecessary increasing compilation time.
+template <bool IsESIMDKernel = false, typename PropertiesT,
+          class = typename std::enable_if_t<
+              ext::oneapi::experimental::is_property_list_v<PropertiesT>>>
+constexpr KernelPropertyHolderStructTy
+extractKernelProperties(PropertiesT Props) {
+  static_assert(
+      !PropertiesT::template has_property<
+          sycl::ext::intel::experimental::fp_control_key>() ||
+          (PropertiesT::template has_property<
+               sycl::ext::intel::experimental::fp_control_key>() &&
+           IsESIMDKernel),
+      "Floating point control property is supported for ESIMD kernels only.");
+  static_assert(
+      !PropertiesT::template has_property<
+          sycl::ext::oneapi::experimental::indirectly_callable_key>(),
+      "indirectly_callable property cannot be applied to SYCL kernels");
+
+  return KernelPropertyHolderStructTy(Props);
+}
+} // namespace kernel_launch_properties_v1
 
 } // namespace detail
 } // namespace _V1

sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp

@@ -22,7 +22,7 @@ struct cluster_size
           cluster_size<Dim>,
           ::sycl::ext::oneapi::experimental::detail::ClusterLaunch> {
   cluster_size(const range<Dim> &size) : size(size) {}
-  sycl::range<Dim> get_cluster_size() { return size; }
+  sycl::range<Dim> get_cluster_size() const { return size; }
 
 private:
   range<Dim> size;