Merge branch 'sycl' into rm-xfail-jm

Author: dkhaldi

sycl/CMakeLists.txt

@@ -455,7 +455,8 @@ set( SYCL_TOOLCHAIN_DEPLOY_COMPONENTS
      sycl-headers-extras
      sycl
      libsycldevice
-     level-zero-sycl-dev
+     unified-memory-framework
+     unified-runtime-loader
      ${XPTIFW_LIBS}
      ${SYCL_TOOLCHAIN_DEPS}
 )

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -117,11 +117,13 @@ if(SYCL_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 568a96aabc6edabe8514ae163aecc64cd5a41878
-  # Author: Mateusz P. Nowak <[email protected]>
-  # Date:   Tue Oct 15 13:57:26 2024 +0200
-  #     Benchmark updates for faster run and more reliable results (#2164)
-  set(UNIFIED_RUNTIME_TAG 568a96aabc6edabe8514ae163aecc64cd5a41878)
+  # commit 80fdb0261d4a3623b05d1795f2f59ace9f718d76
+  # Merge: 568a96aa 77b6c4c7
+  # Author: aarongreig <[email protected]>
+  # Date:   Wed Oct 16 13:53:51 2024 +0100
+  #     Merge pull request #2206 from oneapi-src/revert-2138-counter-based-3
+  #     Revert "[L0] Phase 2 of Counter-Based Event Implementation"
+  set(UNIFIED_RUNTIME_TAG 80fdb0261d4a3623b05d1795f2f59ace9f718d76)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
   # Due to the use of dependentloadflag and no installer for UMF and hwloc we need
@@ -230,14 +232,10 @@ find_package(Threads REQUIRED)
 
 if(TARGET UnifiedRuntimeLoader)
   # Install the UR loader.
-  # TODO: this is piggy-backing on the existing target component level-zero-sycl-dev
-  # When UR is moved to its separate repo perhaps we should introduce new component,
-  # e.g. unified-runtime-sycl-dev.
-  # See github issue #14598
   install(TARGETS ur_loader
-    LIBRARY DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT level-zero-sycl-dev
-    ARCHIVE DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT level-zero-sycl-dev
-    RUNTIME DESTINATION "bin" COMPONENT level-zero-sycl-dev
+    LIBRARY DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT unified-runtime-loader
+    ARCHIVE DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT unified-runtime-loader
+    RUNTIME DESTINATION "bin" COMPONENT unified-runtime-loader
   )
 endif()
 
@@ -301,8 +299,7 @@ if("native_cpu" IN_LIST SYCL_ENABLE_BACKENDS)
   endif()
 endif()
 
-# TODO: this is piggy-backing on the existing target component level-zero-sycl-dev
 install(TARGETS umf
-  LIBRARY DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT level-zero-sycl-dev
-  ARCHIVE DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT level-zero-sycl-dev
-  RUNTIME DESTINATION "bin" COMPONENT level-zero-sycl-dev)
+  LIBRARY DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT unified-memory-framework
+  ARCHIVE DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT unified-memory-framework
+  RUNTIME DESTINATION "bin" COMPONENT unified-memory-framework)

sycl/include/sycl/accessor_image.hpp

@@ -12,20 +12,21 @@
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
-template <int Dim, typename T> struct IsValidCoordDataT;
-template <typename T> struct IsValidCoordDataT<1, T> {
-  constexpr static bool value = detail::is_contained<
-      T, detail::type_list<opencl::cl_int, opencl::cl_float>>::type::value;
+template <int Dim, typename T, bool AllowFP = true> struct IsValidCoordDataT;
+template <typename T, bool AllowFP> struct IsValidCoordDataT<1, T, AllowFP> {
+  constexpr static bool value =
+      std::is_same_v<T, opencl::cl_int> ||
+      (AllowFP && std::is_same_v<T, opencl::cl_float>);
 };
-template <typename T> struct IsValidCoordDataT<2, T> {
-  constexpr static bool value = detail::is_contained<
-      T, detail::type_list<vec<opencl::cl_int, 2>,
-                           vec<opencl::cl_float, 2>>>::type::value;
+template <typename T, bool AllowFP> struct IsValidCoordDataT<2, T, AllowFP> {
+  constexpr static bool value =
+      std::is_same_v<T, vec<opencl::cl_int, 2>> ||
+      (AllowFP && std::is_same_v<T, vec<opencl::cl_float, 2>>);
 };
-template <typename T> struct IsValidCoordDataT<3, T> {
-  constexpr static bool value = detail::is_contained<
-      T, detail::type_list<vec<opencl::cl_int, 4>,
-                           vec<opencl::cl_float, 4>>>::type::value;
+template <typename T, bool AllowFP> struct IsValidCoordDataT<3, T, AllowFP> {
+  constexpr static bool value =
+      std::is_same_v<T, vec<opencl::cl_int, 4>> ||
+      (AllowFP && std::is_same_v<T, vec<opencl::cl_float, 4>>);
 };
 
 template <int Dim, typename T> struct IsValidUnsampledCoord2020DataT;
@@ -448,12 +449,12 @@ class image_accessor
   // (accessTarget == access::target::image && accessMode == access::mode::read)
   // || (accessTarget == access::target::host_image && ( accessMode ==
   // access::mode::read || accessMode == access::mode::read_write))
-  template <typename CoordT, int Dims = Dimensions,
-            typename = std::enable_if_t<
-                (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
-                (detail::is_genint_v<CoordT>) &&
-                ((IsImageAcc && IsImageAccessReadOnly) ||
-                 (IsHostImageAcc && IsImageAccessAnyRead))>>
+  template <
+      typename CoordT, int Dims = Dimensions,
+      typename = std::enable_if_t<
+          (IsValidCoordDataT<Dims, CoordT, /* AllowFP = */ false>::value) &&
+          ((IsImageAcc && IsImageAccessReadOnly) ||
+           (IsHostImageAcc && IsImageAccessAnyRead))>>
   DataT read(const CoordT &Coords) const {
 #ifdef __SYCL_DEVICE_ONLY__
     return __invoke__ImageRead<DataT, OCLImageTy, CoordT>(MImageObj, Coords);
@@ -470,7 +471,7 @@ class image_accessor
   // access::mode::read || accessMode == access::mode::read_write))
   template <typename CoordT, int Dims = Dimensions,
             typename = std::enable_if_t<
-                (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value) &&
+                (IsValidCoordDataT<Dims, CoordT>::value) &&
                 ((IsImageAcc && IsImageAccessReadOnly) ||
                  (IsHostImageAcc && IsImageAccessAnyRead))>>
   DataT read(const CoordT &Coords, const sampler &Smpl) const {
@@ -494,10 +495,10 @@ class image_accessor
   // accessMode == access::mode::read_write))
   template <
       typename CoordT, int Dims = Dimensions,
-      typename = std::enable_if_t<(Dims > 0) && (detail::is_genint_v<CoordT>) &&
-                                  (IsValidCoordDataT<Dims, CoordT>::value) &&
-                                  ((IsImageAcc && IsImageAccessWriteOnly) ||
-                                   (IsHostImageAcc && IsImageAccessAnyWrite))>>
+      typename = std::enable_if_t<
+          (IsValidCoordDataT<Dims, CoordT, /* AllowFP = */ false>::value) &&
+          ((IsImageAcc && IsImageAccessWriteOnly) ||
+           (IsHostImageAcc && IsImageAccessAnyWrite))>>
   void write(const CoordT &Coords, const DataT &Color) const {
 #ifdef __SYCL_DEVICE_ONLY__
     __invoke__ImageWrite<OCLImageTy, CoordT, DataT>(MImageObj, Coords, Color);
@@ -546,23 +547,21 @@ class __image_array_slice__ {
       size_t Idx)
       : MBaseAcc(BaseAcc), MIdx(Idx) {}
 
-  template <typename CoordT, int Dims = Dimensions,
-            typename = std::enable_if_t<
-                (Dims > 0) && (IsValidCoordDataT<Dims, CoordT>::value)>>
+  template <
+      typename CoordT, int Dims = Dimensions,
+      typename = std::enable_if_t<(IsValidCoordDataT<Dims, CoordT>::value)>>
   DataT read(const CoordT &Coords) const {
     return MBaseAcc.read(getAdjustedCoords(Coords));
   }
 
   template <typename CoordT, int Dims = Dimensions,
-            typename = std::enable_if_t<(Dims > 0) &&
-                                        IsValidCoordDataT<Dims, CoordT>::value>>
+            typename = std::enable_if_t<IsValidCoordDataT<Dims, CoordT>::value>>
   DataT read(const CoordT &Coords, const sampler &Smpl) const {
     return MBaseAcc.read(getAdjustedCoords(Coords), Smpl);
   }
 
   template <typename CoordT, int Dims = Dimensions,
-            typename = std::enable_if_t<(Dims > 0) &&
-                                        IsValidCoordDataT<Dims, CoordT>::value>>
+            typename = std::enable_if_t<IsValidCoordDataT<Dims, CoordT>::value>>
   void write(const CoordT &Coords, const DataT &Color) const {
     return MBaseAcc.write(getAdjustedCoords(Coords), Color);
   }

sycl/include/sycl/detail/generic_type_lists.hpp

@@ -70,8 +70,6 @@ using scalar_vector_bfloat16_list =
 using bfloat16_list =
     tl_append<scalar_bfloat16_list, vector_bfloat16_list, marray_bfloat16_list>;
 
-using half_bfloat16_list = tl_append<scalar_half_list, scalar_bfloat16_list>;
-
 using scalar_float_list = type_list<float>;
 
 using vector_float_list =

sycl/include/sycl/detail/generic_type_traits.hpp

@@ -31,17 +31,6 @@ template <typename T>
 inline constexpr bool is_svgenfloatf_v =
     is_contained_v<T, gtl::scalar_vector_float_list>;
 
-template <typename T>
-inline constexpr bool is_half_v = is_contained_v<T, gtl::scalar_half_list>;
-
-template <typename T>
-inline constexpr bool is_bfloat16_v =
-    is_contained_v<T, gtl::scalar_bfloat16_list>;
-
-template <typename T>
-inline constexpr bool is_half_or_bf16_v =
-    is_contained_v<T, gtl::half_bfloat16_list>;
-
 template <typename T>
 inline constexpr bool is_svgenfloath_v =
     is_contained_v<T, gtl::scalar_vector_half_list>;
@@ -57,9 +46,6 @@ template <typename T>
 inline constexpr bool is_vgenfloat_v =
     is_contained_v<T, gtl::vector_floating_list>;
 
-template <typename T>
-inline constexpr bool is_genint_v = is_contained_v<T, gtl::signed_int_list>;
-
 template <typename T>
 inline constexpr bool is_geninteger_v = is_contained_v<T, gtl::integer_list>;
 
@@ -141,10 +127,11 @@ template <typename T> auto convertToOpenCLType(T &&x) {
     // sycl::half may convert to _Float16, and we would try to instantiate
     // vec class with _Float16 DataType, which is not expected there. As
     // such, leave vector<half, N> as-is.
-    using MatchingVec = vec<std::conditional_t<is_half_v<ElemTy>, ElemTy,
-                                               decltype(convertToOpenCLType(
-                                                   std::declval<ElemTy>()))>,
-                            no_ref::size()>;
+    using MatchingVec =
+        vec<std::conditional_t<std::is_same_v<ElemTy, half>, ElemTy,
+                               decltype(convertToOpenCLType(
+                                   std::declval<ElemTy>()))>,
+            no_ref::size()>;
 #ifdef __SYCL_DEVICE_ONLY__
     return sycl::bit_cast<typename MatchingVec::vector_t>(x);
 #else
@@ -160,11 +147,11 @@ template <typename T> auto convertToOpenCLType(T &&x) {
                                           fixed_width_unsigned<sizeof(no_ref)>>;
     static_assert(sizeof(OpenCLType) == sizeof(T));
     return static_cast<OpenCLType>(x);
-  } else if constexpr (is_half_v<no_ref>) {
+  } else if constexpr (std::is_same_v<no_ref, half>) {
     using OpenCLType = sycl::detail::half_impl::BIsRepresentationT;
     static_assert(sizeof(OpenCLType) == sizeof(T));
     return static_cast<OpenCLType>(x);
-  } else if constexpr (is_bfloat16_v<no_ref>) {
+  } else if constexpr (std::is_same_v<no_ref, ext::oneapi::bfloat16>) {
     // On host, don't interpret BF16 as uint16.
 #ifdef __SYCL_DEVICE_ONLY__
     using OpenCLType = sycl::ext::oneapi::detail::Bfloat16StorageT;

sycl/include/sycl/detail/helpers.hpp

@@ -240,11 +240,11 @@ getSPIRVMemorySemanticsMask(const access::fence_space AccessSpace,
 
 // To ensure loop unrolling is done when processing dimensions.
 template <size_t... Inds, class F>
-void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
+constexpr void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
   (f(std::integral_constant<size_t, Inds>{}), ...);
 }
 
-template <size_t count, class F> void loop(F &&f) {
+template <size_t count, class F> constexpr void loop(F &&f) {
   loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
 }
 inline constexpr bool is_power_of_two(int x) { return (x & (x - 1)) == 0; }

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

@@ -14,7 +14,6 @@
 #include <sycl/ext/oneapi/experimental/detail/invoke_simd_types.hpp>
 #include <sycl/ext/oneapi/experimental/uniform.hpp>
 
-#include <sycl/detail/boost/mp11.hpp>
 #include <sycl/sub_group.hpp>
 
 #include <functional>
@@ -71,8 +70,6 @@ namespace ext::oneapi::experimental {
 // --- Helpers
 namespace detail {
 
-namespace __MP11_NS = sycl::detail::boost::mp11;
-
 // This structure performs the SPMD-to-SIMD parameter type conversion as defined
 // by the spec.
 template <class T, int N, class = void> struct spmd2simd;
@@ -154,8 +151,7 @@ struct is_simd_or_mask_type<simd_mask<T, N>> : std::true_type {};
 // Checks if all the types in the parameter pack are uniform<T>.
 template <class... SpmdArgs> struct all_uniform_types {
   constexpr operator bool() {
-    using TypeList = __MP11_NS::mp_list<SpmdArgs...>;
-    return __MP11_NS::mp_all_of<TypeList, is_uniform_type>::value;
+    return ((is_uniform_type<SpmdArgs>::value && ...));
   }
 };
 
@@ -193,26 +189,32 @@ constexpr void verify_return_type_matches_sg_size() {
 // as prescribed by the spec assuming this subgroup size. One and only one
 // subgroup size should conform.
 template <class SimdCallable, class... SpmdArgs> struct sg_size {
-  template <class N>
-  using IsInvocableSgSize = __MP11_NS::mp_bool<std::is_invocable_v<
-      SimdCallable, typename spmd2simd<SpmdArgs, N::value>::type...>>;
-
   __DPCPP_SYCL_EXTERNAL constexpr operator int() {
-    using SupportedSgSizes = __MP11_NS::mp_list_c<int, 1, 2, 4, 8, 16, 32>;
-    using InvocableSgSizes =
-        __MP11_NS::mp_copy_if<SupportedSgSizes, IsInvocableSgSize>;
-    constexpr auto found_invoke_simd_target =
-        __MP11_NS::mp_empty<InvocableSgSizes>::value != 1;
-    if constexpr (found_invoke_simd_target) {
-      static_assert((__MP11_NS::mp_size<InvocableSgSizes>::value == 1) &&
-                    "multiple invoke_simd targets found");
-      return __MP11_NS::mp_front<InvocableSgSizes>::value;
-    }
-    static_assert(
-        found_invoke_simd_target,
-        "No callable invoke_simd target found. Confirm the "
-        "invoke_simd invocation argument types are convertible to the "
-        "invoke_simd target argument types");
+    constexpr auto x = []() constexpr {
+      constexpr int supported_sg_sizes[] = {1, 2, 4, 8, 16, 32};
+      int num_found = 0;
+      int found_sg_size = 0;
+      sycl::detail::loop<std::size(supported_sg_sizes)>([&](auto idx) {
+        constexpr auto sg_size = supported_sg_sizes[idx];
+        if (std::is_invocable_v<
+                SimdCallable, typename spmd2simd<SpmdArgs, sg_size>::type...>) {
+          ++num_found;
+          found_sg_size = sg_size;
+        }
+      });
+      return std::pair{num_found, found_sg_size};
+    }();
+
+    constexpr auto num_found = x.first;
+    constexpr auto found_sg_size = x.second;
+
+    static_assert(num_found != 0,
+                  "No callable invoke_simd target found. Confirm the "
+                  "invoke_simd invocation argument types are convertible to "
+                  "the invoke_simd target argument types");
+    static_assert(num_found == 1, "Multiple invoke_simd targets found!");
+
+    return found_sg_size;
   }
 };
 

sycl/include/sycl/ext/oneapi/memcpy2d.hpp

@@ -18,6 +18,7 @@ template <typename T, typename>
 void handler::ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src,
                                   size_t SrcPitch, size_t Width,
                                   size_t Height) {
+#ifndef __SYCL_DEVICE_ONLY__
   throwIfGraphAssociated<
       ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
           sycl_ext_oneapi_memcpy2d>();
@@ -30,6 +31,7 @@ void handler::ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src,
     throw sycl::exception(sycl::make_error_code(errc::invalid),
                           "Source pitch must be greater than or equal "
                           "to the width specified in 'ext_oneapi_memcpy2d'");
+#endif
 
   // Get the type of the pointers.
   context Ctx = detail::createSyclObjFromImpl<context>(getContextImplPtr());