Merge branch 'sycl' into e2e-split

.github/workflows/sycl-linux-precommit.yml

@@ -162,7 +162,6 @@ jobs:
             runner: '["Linux", "amdgpu"]'
             image: ghcr.io/intel/llvm/ubuntu2204_build:latest
             image_extra_opts: --device=/dev/dri --device=/dev/kfd
-            extra_cmake_args: -DHIP_PLATFORM="AMD" -DAMD_ARCH="gfx1031"
           - name: CUDA system
             runner: '["Linux", "cuda"]'
             image: ghcr.io/intel/llvm/ubuntu2204_build:latest

.github/workflows/sycl-linux-run-tests.yml

@@ -21,9 +21,6 @@ on:
         type: string
         required: True
       extra_cmake_args:
-        description: |
-          If empty, then HIP_PLATFORM and AMD_ARCH would be automatically added
-          if inputs.target_devices contains 'ext_oneapi_hip'
         type: string
         required: False
       tests_selector:
@@ -282,16 +279,6 @@ jobs:
       run: |
         if [ -n "$CMAKE_EXTRA_ARGS" ]; then
           echo "opts=$CMAKE_EXTRA_ARGS" >> $GITHUB_OUTPUT
-        else
-          if [ "${{ contains(inputs.target_devices, 'ext_oneapi_hip')  }}" == "true" ]; then
-            if [ "${{ runner.name }}" == "cp-amd-runner" ]; then
-              echo 'opts=-DHIP_PLATFORM="AMD" -DAMD_ARCH="gfx1030"' >> $GITHUB_OUTPUT
-            else
-              echo 'opts=-DHIP_PLATFORM="AMD" -DAMD_ARCH="gfx1031"' >> $GITHUB_OUTPUT
-            fi
-          else
-            echo 'opts=' >> $GITHUB_OUTPUT
-          fi
         fi
     - name: Configure E2E tests
       if: inputs.tests_selector == 'e2e'

clang/lib/Sema/SemaAccess.cpp

@@ -301,6 +301,7 @@ static AccessResult IsDerivedFromInclusive(const CXXRecordDecl *Derived,
                                            const CXXRecordDecl *Target) {
   assert(Derived->getCanonicalDecl() == Derived);
   assert(Target->getCanonicalDecl() == Target);
+  assert(Derived->getDefinition() && "Expecting a complete type");
 
   if (Derived == Target) return AR_accessible;
 
@@ -776,6 +777,8 @@ static AccessResult HasAccess(Sema &S,
     // [B3] and [M3]
     } else {
       assert(Access == AS_protected);
+      if (!ECRecord->getDefinition())
+        continue;
       switch (IsDerivedFromInclusive(ECRecord, NamingClass)) {
       case AR_accessible: break;
       case AR_inaccessible: continue;

clang/test/SemaSYCL/attr-add-ir-attributes.cpp

@@ -991,3 +991,16 @@ struct __attribute__((sycl_special_class)) InvalidSpecialClassStruct32 {
 struct [[__sycl_detail__::add_ir_attributes_kernel_parameter("Attr1", 1)]] InvalidKernelParameterSubjectStruct;     // expected-error {{'add_ir_attributes_kernel_parameter' attribute only applies to parameters}}
 [[__sycl_detail__::add_ir_attributes_kernel_parameter("Attr1", 1)]] void InvalidKernelParameterSubjectFunction() {} // expected-error {{'add_ir_attributes_kernel_parameter' attribute only applies to parameters}}
 [[__sycl_detail__::add_ir_attributes_kernel_parameter("Attr1", 1)]] int InvalidKernelParameterSubjectVar;           // expected-error {{'add_ir_attributes_kernel_parameter' attribute only applies to parameters}}
+
+struct A {
+  protected:
+  static constexpr const char *ir_attribute_name = ""; // expected-note {{declared protected here}}
+  static constexpr auto ir_attribute_value = nullptr;  // expected-note {{declared protected here}}
+};
+
+template <typename Ts>
+struct [[__sycl_detail__::add_ir_attributes_global_variable(
+             Ts::ir_attribute_name, Ts::ir_attribute_value)]] B {  // expected-error {{'ir_attribute_name' is a protected member of 'A'}} // expected-error {{'ir_attribute_value' is a protected member of 'A'}}
+};
+
+B<A> v; // expected-note {{in instantiation of template class 'B<A>' requested here}}

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

@@ -1359,6 +1359,9 @@ class BinaryWrapper {
     ObjcopyPath = *ObjcopyPathOrErr;
   }
 
+  BinaryWrapper(const BinaryWrapper &BW) = delete;
+  BinaryWrapper &operator=(const BinaryWrapper &BW) = delete;
+
   ~BinaryWrapper() {
     if (TempFiles.empty())
       return;

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-a20debc",
-      "version": "a20debc",
-      "updated_at": "2024-09-15T13:44:38Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/1934718090/zip",
+      "github_tag": "igc-dev-ad75a20",
+      "version": "ad75a20",
+      "updated_at": "2024-11-10T01:11:34Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2167439771/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

devops/scripts/install_drivers.sh

@@ -143,7 +143,9 @@ InstallIGFX () {
     echo "Download IGC dev git hash $IGC_DEV_VER"
     get_pre_release_igfx $IGC_DEV_URL $IGC_DEV_VER
     echo "Install IGC dev git hash $IGC_DEV_VER"
-    dpkg -i *.deb
+    # New dev IGC packaged iga64 conflicting with iga64 from intel-igc-media
+    # force overwrite to workaround it first.
+    dpkg -i --force-overwrite *.deb
     echo "Install libopencl-clang"
     # Workaround only, will download deb and install with dpkg once fixed.
     cp -d libopencl-clang.so.14*  /usr/local/lib/

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,7 +1,7 @@
-# commit cd92e72bbc4ebddef63c63c0f7e66a410f4b9552
-# Merge: 9a209aa5 b1222f08
-# Author: Callum Fare <callum@codeplay.com>
-# Date:   Wed Nov 13 09:57:16 2024 +0000
-#     Merge pull request #2254 from PietroGhg/pietro/events_rr
-#     [NATIVECPU] Implement events on Native CPU
-set(UNIFIED_RUNTIME_TAG cd92e72bbc4ebddef63c63c0f7e66a410f4b9552)
+# commit 3a5b23c8b475712f9107c1d5ab41f27a1465578e
+# Merge: f9f71f17 1696524d
+# Author: Piotr Balcer <piotr.balcer@intel.com>
+# Date:   Thu Nov 14 14:38:05 2024 +0100
+#    Merge pull request #2253 from pbalcer/low-power-events   
+#    add low-power events experimental extension spec
+set(UNIFIED_RUNTIME_TAG 3a5b23c8b475712f9107c1d5ab41f27a1465578e)

sycl/include/sycl/ext/intel/esimd/memory_properties.hpp

@@ -77,8 +77,8 @@ class properties
 // Deduction guides
 template <typename... PropertyValueTs>
 properties(PropertyValueTs... props)
-    -> properties<typename sycl::ext::oneapi::experimental::detail::Sorted<
-        PropertyValueTs...>::type>;
+    -> properties<typename sycl::ext::oneapi::experimental::detail::
+                      properties_sorter<PropertyValueTs...>::type>;
 #endif
 
 /// The 'alignment' property is used to specify the alignment of memory

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

@@ -54,16 +54,10 @@ template <typename... Ts>
 using contains_alignment =
     detail::ContainsProperty<alignment_key, std::tuple<Ts...>>;
 
-// properties filter
-template <typename property_list, template <class...> typename filter>
-using PropertiesFilter =
-    sycl::detail::boost::mp11::mp_copy_if<property_list, filter>;
-
 // filter properties that are applied on annotations
-template <typename... Props>
-using annotation_filter =
-    properties<PropertiesFilter<detail::properties_type_list<Props...>,
-                                propagateToPtrAnnotation>>;
+template <typename PropertyListTy>
+using annotation_filter = decltype(filter_properties<propagateToPtrAnnotation>(
+    std::declval<PropertyListTy>()));
 } // namespace detail
 
 template <typename I, typename P> struct annotationHelper {};
@@ -108,8 +102,8 @@ class annotated_ref<T, detail::properties_t<Props...>> {
   // implicit conversion with annotaion
   operator T() const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return annotationHelper<T, detail::annotation_filter<Props...>>::load(
-        m_Ptr);
+    return annotationHelper<
+        T, detail::annotation_filter<property_list_t>>::load(m_Ptr);
 #else
     return *m_Ptr;
 #endif
@@ -119,8 +113,8 @@ class annotated_ref<T, detail::properties_t<Props...>> {
   template <class O, typename = std::enable_if_t<!detail::is_ann_ref_v<O>>>
   T operator=(O &&Obj) const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return annotationHelper<T, detail::annotation_filter<Props...>>::store(
-        m_Ptr, Obj);
+    return annotationHelper<
+        T, detail::annotation_filter<property_list_t>>::store(m_Ptr, Obj);
 #else
     return *m_Ptr = std::forward<O>(Obj);
 #endif
@@ -385,8 +379,8 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
 
   T *get() const noexcept {
 #ifdef __SYCL_DEVICE_ONLY__
-    return annotationHelper<T, detail::annotation_filter<Props...>>::annotate(
-        m_Ptr);
+    return annotationHelper<
+        T, detail::annotation_filter<property_list_t>>::annotate(m_Ptr);
 #else
     return m_Ptr;
 #endif

sycl/include/sycl/ext/oneapi/properties/properties.hpp

@@ -132,6 +132,54 @@ constexpr bool properties_are_valid_for_ctad = []() constexpr {
 
 template <typename... property_tys> struct properties_type_list;
 template <typename... property_tys> struct invalid_properties_type_list {};
+
+template <typename... property_tys> struct properties_sorter {
+  // Not using "auto" due to MSVC bug in v19.36 and older. v19.37 and later is
+  // able to compile "auto" just fine. See https://godbolt.org/z/eW3rjjs7n.
+  static constexpr std::array<int, sizeof...(property_tys)> sorted_indices =
+      []() constexpr {
+        int idx = 0;
+        int N = sizeof...(property_tys);
+        // std::sort isn't constexpr until C++20. Also, it's possible there will
+        // be a compiler builtin to sort types, in which case we should start
+        // using that.
+        std::array to_sort{
+            std::pair{PropertyID<property_tys>::value, idx++}...};
+        auto swap_pair = [](auto &x, auto &y) constexpr {
+          auto tmp_first = x.first;
+          auto tmp_second = x.second;
+          x.first = y.first;
+          x.second = y.second;
+          y.first = tmp_first;
+          y.second = tmp_second;
+        };
+        for (int i = 0; i < N; ++i)
+          for (int j = i; j < N; ++j)
+            if (to_sort[j].first < to_sort[i].first)
+              swap_pair(to_sort[i], to_sort[j]);
+
+        std::array<int, sizeof...(property_tys)> sorted_indices{};
+        for (int i = 0; i < N; ++i)
+          sorted_indices[i] = to_sort[i].second;
+
+        return sorted_indices;
+      }();
+
+  template <typename> struct helper;
+  template <int... IdxSeq>
+  struct helper<std::integer_sequence<int, IdxSeq...>> {
+    using type = properties_type_list<
+        nth_type_t<sorted_indices[IdxSeq], property_tys...>...>;
+  };
+
+  using type = typename helper<
+      std::make_integer_sequence<int, sizeof...(property_tys)>>::type;
+};
+// Specialization to avoid zero-size array creation.
+template <> struct properties_sorter<> {
+  using type = properties_type_list<>;
+};
+
 } // namespace detail
 
 template <typename properties_type_list_ty> class __SYCL_EBO properties;
@@ -271,17 +319,19 @@ class __SYCL_EBO properties<detail::properties_type_list<property_tys...>>
 };
 
 // Deduction guides
-template <typename... PropertyValueTs,
+template <typename... unsorted_property_tys,
           typename = std::enable_if_t<
-              detail::properties_are_valid_for_ctad<PropertyValueTs...>>>
-properties(PropertyValueTs... props)
-    -> properties<typename detail::Sorted<PropertyValueTs...>::type>;
+              detail::properties_are_valid_for_ctad<unsorted_property_tys...>>>
+properties(unsorted_property_tys... props)
+    -> properties<
+        typename detail::properties_sorter<unsorted_property_tys...>::type>;
 
-template <typename... PropertyValueTs,
+template <typename... unsorted_property_tys,
           typename = std::enable_if_t<
-              !detail::properties_are_valid_for_ctad<PropertyValueTs...>>>
-properties(PropertyValueTs... props)
-    -> properties<detail::invalid_properties_type_list<PropertyValueTs...>>;
+              !detail::properties_are_valid_for_ctad<unsorted_property_tys...>>>
+properties(unsorted_property_tys... props)
+    -> properties<
+        detail::invalid_properties_type_list<unsorted_property_tys...>>;
 
 using empty_properties_t = decltype(properties{});
 

sycl/include/sycl/ext/oneapi/properties/property_utils.hpp

@@ -8,10 +8,6 @@
 
 #pragma once
 
-#include <sycl/detail/boost/mp11/algorithm.hpp>        // for mp_sort_q
-#include <sycl/detail/boost/mp11/detail/mp_list.hpp>   // for mp_list
-#include <sycl/detail/boost/mp11/detail/mp_rename.hpp> // for mp_rename
-#include <sycl/detail/boost/mp11/integral.hpp>         // for mp_bool
 #include <sycl/ext/oneapi/properties/property.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 
@@ -94,23 +90,6 @@ template <typename RHS> struct SelectNonVoid<void, RHS> {
   using type = RHS;
 };
 
-// Sort types accoring to their PropertyID.
-struct SortByPropertyId {
-  template <typename T1, typename T2>
-  using fn = sycl::detail::boost::mp11::mp_bool<(PropertyID<T1>::value <
-                                                 PropertyID<T2>::value)>;
-};
-template <typename... Ts> struct Sorted {
-  static_assert(detail::AllPropertyValues<std::tuple<Ts...>>::value,
-                "Unrecognized property in property list.");
-  using properties = sycl::detail::boost::mp11::mp_list<Ts...>;
-  using sortedProperties =
-      sycl::detail::boost::mp11::mp_sort_q<properties, SortByPropertyId>;
-  using type =
-      sycl::detail::boost::mp11::mp_rename<sortedProperties,
-                                           detail::properties_type_list>;
-};
-
 //******************************************************************************
 // Property merging
 //******************************************************************************
@@ -326,6 +305,50 @@ struct ConditionalPropertyMetaInfo
     : std::conditional_t<Condition, PropertyMetaInfo<PropT>,
                          IgnoredPropertyMetaInfo> {};
 
+template <template <typename> typename predicate, typename... property_tys>
+struct filter_properties_impl {
+  static constexpr auto idx_info = []() constexpr {
+    constexpr int N = sizeof...(property_tys);
+    std::array<int, N> indexes{};
+    int num_matched = 0;
+    int idx = 0;
+    (((predicate<property_tys>::value ? indexes[num_matched++] = idx++ : idx++),
+      ...));
+
+    return std::pair{indexes, num_matched};
+  }();
+
+  // Helper to convert constexpr indices values to an std::index_sequence type.
+  // Values -> type is the key here.
+  template <int... Idx>
+  static constexpr auto idx_seq(std::integer_sequence<int, Idx...>) {
+    return std::integer_sequence<int, idx_info.first[Idx]...>{};
+  }
+
+  using selected_idx_seq =
+      decltype(idx_seq(std::make_integer_sequence<int, idx_info.second>{}));
+
+  // Using prop_list_ty so that we don't need to explicitly spell out
+  //  `properties` template parameters' implementation-details.
+  template <typename prop_list_ty, int... Idxs>
+  static constexpr auto apply_impl(const prop_list_ty &props,
+                                   std::integer_sequence<int, Idxs...>) {
+    return properties{props.template get_property<
+        typename nth_type_t<Idxs, property_tys...>::key_t>()...};
+  }
+
+  template <typename prop_list_ty>
+  static constexpr auto apply(const prop_list_ty &props) {
+    return apply_impl(props, selected_idx_seq{});
+  }
+};
+
+template <template <typename> typename predicate, typename... property_tys>
+constexpr auto filter_properties(
+    const properties<properties_type_list<property_tys...>> &props) {
+  return filter_properties_impl<predicate, property_tys...>::apply(props);
+}
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1

sycl/source/detail/compiler.hpp

@@ -31,6 +31,9 @@
 #define __SYCL_DEVICE_BINARY_TARGET_NVPTX64 "nvptx64"
 #define __SYCL_DEVICE_BINARY_TARGET_AMDGCN "amdgcn"
 #define __SYCL_DEVICE_BINARY_TARGET_NATIVE_CPU "native_cpu"
+// JIT compilation targets for CUDA & HIP devices.
+#define __SYCL_DEVICE_BINARY_TARGET_LLVM_NVPTX64 "llvm_nvptx64"
+#define __SYCL_DEVICE_BINARY_TARGET_LLVM_AMDGCN "llvm_amdgcn"
 
 /// Device binary image property set names recognized by the SYCL runtime.
 /// Name must be consistent with