Initial support for ext_oneapi_composite_device.

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -67,6 +67,8 @@ def AspectExt_oneapi_fixed_size_group : Aspect<"ext_oneapi_fixed_size_group">;
 def AspectExt_oneapi_opportunistic_group : Aspect<"ext_oneapi_opportunistic_group">;
 def AspectExt_oneapi_tangle_group : Aspect<"ext_oneapi_tangle_group">;
 def AspectExt_intel_matrix : Aspect<"ext_intel_matrix">;
+def AspectExt_oneapi_is_composite : Aspect<"ext_oneapi_is_composite">;
+def AspectExt_oneapi_is_component : Aspect<"ext_oneapi_is_component">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -116,7 +118,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_interop_memory_import, AspectExt_oneapi_interop_memory_export,
     AspectExt_oneapi_interop_semaphore_import, AspectExt_oneapi_interop_semaphore_export,
     AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_intel_esimd,
-    AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group, AspectExt_oneapi_tangle_group, AspectExt_intel_matrix],
+    AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
+    AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

sycl/doc/extensions/proposed/sycl_ext_oneapi_composite_device.asciidoc → sycl/doc/extensions/experimental/sycl_ext_oneapi_composite_device.asciidoc

@@ -43,11 +43,12 @@ SYCL specification refer to that revision.
 
 == 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.*
 
 
 == Backend support status

sycl/include/sycl/detail/pi.h

@@ -437,6 +437,10 @@ typedef enum {
   PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_SEMAPHORE_EXPORT_SUPPORT = 0x2010F,
 
   PI_EXT_ONEAPI_DEVICE_INFO_MATRIX_COMBINATIONS = 0x20110,
+
+  // Composite device
+  PI_EXT_ONEAPI_DEVICE_INFO_COMPONENT_DEVICES = 0x20111,
+  PI_EXT_ONEAPI_DEVICE_INFO_COMPOSITE_DEVICE = 0x20112,
 } _pi_device_info;
 
 typedef enum {

sycl/include/sycl/device_aspect_macros.hpp

@@ -303,6 +303,16 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_intel_matrix__ 0
 #endif
 
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_is_composite__
+// __SYCL_ASPECT(ext_oneapi_is_composite, 59)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_is_composite__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_is_component__
+// __SYCL_ASPECT(ext_oneapi_is_component, 60)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_is_component__ 0
+#endif
+
 #ifndef __SYCL_ANY_DEVICE_HAS_host__
 // __SYCL_ASPECT(host, 0)
 #define __SYCL_ANY_DEVICE_HAS_host__ 0
@@ -597,3 +607,13 @@
 // __SYCL_ASPECT(ext_intel_matrix, 58)
 #define __SYCL_ANY_DEVICE_HAS_ext_intel_matrix__ 0
 #endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_is_composite__
+// __SYCL_ASPECT(ext_oneapi_is_composite, 59)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_is_composite__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_is_component__
+// __SYCL_ASPECT(ext_oneapi_is_component, 60)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_is_component__ 0
+#endif

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

@@ -0,0 +1,21 @@
+//==---------- composite_device.hpp - SYCL Composite Device ----------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/device.hpp>
+
+#include <vector>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+__SYCL_EXPORT std::vector<device> get_composite_devices();
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/info/aspects.def

@@ -53,3 +53,5 @@ __SYCL_ASPECT(ext_oneapi_fixed_size_group, 55)
 __SYCL_ASPECT(ext_oneapi_opportunistic_group, 56)
 __SYCL_ASPECT(ext_oneapi_tangle_group, 57)
 __SYCL_ASPECT(ext_intel_matrix, 58)
+__SYCL_ASPECT(ext_oneapi_is_composite, 59)
+__SYCL_ASPECT(ext_oneapi_is_component, 60)

sycl/include/sycl/info/ext_oneapi_device_traits.def

@@ -36,6 +36,15 @@ __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device,
 __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device,
                          mipmap_max_anisotropy, float,
                          PI_EXT_ONEAPI_DEVICE_INFO_MIPMAP_MAX_ANISOTROPY)
+
+// Composite devices
+__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device,
+                         component_devices, std::vector<sycl::device>,
+                         PI_EXT_ONEAPI_DEVICE_INFO_COMPONENT_DEVICES)
+__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device,
+                         composite_device, sycl::device,
+                         PI_EXT_ONEAPI_DEVICE_INFO_COMPOSITE_DEVICE)
+
 #ifdef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC
 #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF

sycl/include/sycl/platform.hpp

@@ -188,6 +188,8 @@ class __SYCL_EXPORT platform : public detail::OwnerLessBase<platform> {
   /// \return the default context
   context ext_oneapi_get_default_context() const;
 
+  std::vector<device> ext_oneapi_get_composite_devices() const;
+
 private:
   pi_native_handle getNative() const;
 

sycl/include/sycl/sycl.hpp

@@ -85,6 +85,7 @@
 #include <sycl/ext/oneapi/experimental/ballot_group.hpp>
 #include <sycl/ext/oneapi/experimental/bfloat16_math.hpp>
 #include <sycl/ext/oneapi/experimental/builtins.hpp>
+#include <sycl/ext/oneapi/experimental/composite_device.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
 #include <sycl/ext/oneapi/experimental/fixed_size_group.hpp>
 #include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>

sycl/plugins/level_zero/CMakeLists.txt

@@ -4,7 +4,7 @@ if (NOT DEFINED LEVEL_ZERO_LIBRARY OR NOT DEFINED LEVEL_ZERO_INCLUDE_DIR)
   message(STATUS "Download Level Zero loader and headers from github.com")
 
   set(LEVEL_ZERO_LOADER_REPO "https://github.com/oneapi-src/level-zero.git")
-  set(LEVEL_ZERO_LOADER_TAG v1.11.0)
+  set(LEVEL_ZERO_LOADER_TAG v1.15.1)
 
   # Disable due to a bug https://github.com/oneapi-src/level-zero/issues/104
   set(CMAKE_INCLUDE_CURRENT_DIR OFF)

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -608,6 +608,12 @@ inline pi_result ur2piDeviceInfoValue(ur_device_info_t ParamName,
      * No need to convert since types are compatible
      */
     *ParamValueSizeRet = sizeof(pi_device_fp_config);
+  } else if (ParamName == UR_DEVICE_INFO_COMPONENT_DEVICES) {
+    if (ParamValueSizeRet && *ParamValueSizeRet != 0) {
+      const uint32_t UrNumberElements =
+          *ParamValueSizeRet / sizeof(ur_device_handle_t);
+      *ParamValueSizeRet = UrNumberElements * sizeof(pi_device);
+    }
   } else {
 
     // TODO: what else needs a UR-PI translation?
@@ -974,7 +980,6 @@ inline pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType,
 
 inline pi_result piDeviceRetain(pi_device Device) {
   PI_ASSERT(Device, PI_ERROR_INVALID_DEVICE);
-
   auto UrDevice = reinterpret_cast<ur_device_handle_t>(Device);
   HANDLE_ERRORS(urDeviceRetain(UrDevice));
   return PI_SUCCESS;
@@ -1008,7 +1013,6 @@ inline pi_result piPluginGetLastError(char **Message) {
 inline pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
                                  size_t ParamValueSize, void *ParamValue,
                                  size_t *ParamValueSizeRet) {
-
   ur_device_info_t InfoType;
   switch (ParamName) {
 #define PI_TO_UR_MAP_DEVICE_INFO(FROM, TO)                                     \
@@ -1270,6 +1274,10 @@ inline pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
         UR_DEVICE_INFO_INTEROP_SEMAPHORE_EXPORT_SUPPORT_EXP)
     PI_TO_UR_MAP_DEVICE_INFO(PI_EXT_INTEL_DEVICE_INFO_ESIMD_SUPPORT,
                              UR_DEVICE_INFO_ESIMD_SUPPORT)
+    PI_TO_UR_MAP_DEVICE_INFO(PI_EXT_ONEAPI_DEVICE_INFO_COMPONENT_DEVICES,
+                             UR_DEVICE_INFO_COMPONENT_DEVICES)
+    PI_TO_UR_MAP_DEVICE_INFO(PI_EXT_ONEAPI_DEVICE_INFO_COMPOSITE_DEVICE,
+                             UR_DEVICE_INFO_COMPOSITE_DEVICE)
 #undef PI_TO_UR_MAP_DEVICE_INFO
   default:
     return PI_ERROR_UNKNOWN;

sycl/source/CMakeLists.txt

@@ -171,6 +171,7 @@ set(SYCL_COMMON_SOURCES
     "detail/buffer_impl.cpp"
     "detail/pi.cpp"
     "detail/common.cpp"
+    "detail/composite_device/composite_device.cpp"
     "detail/config.cpp"
     "detail/context_impl.cpp"
     "detail/device_binary_image.cpp"

sycl/source/detail/composite_device/composite_device.cpp

@@ -0,0 +1,44 @@
+//==---------- composite_device.cpp - SYCL Composite Device ----------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <sycl/ext/oneapi/experimental/composite_device.hpp>
+#include <sycl/platform.hpp>
+
+#include <unordered_set>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+std::vector<device> get_composite_devices() {
+  // We use set to filter out duplicates, and unordered because we don't need it
+  // to be sorted, and unordered provides faster insertion.
+  std::unordered_set<device> Composites;
+  auto Devs = sycl::device::get_devices();
+  for (const auto &D : Devs) {
+    if (D.has(sycl::aspect::ext_oneapi_is_component)) {
+      auto Composite = D.get_info<info::device::composite_device>();
+      Composites.insert(Composite);
+    }
+  }
+  std::vector<device> Result;
+  std::copy_if(
+      Composites.begin(), Composites.end(), std::back_inserter(Result),
+      [&](const device &Composite) {
+        auto Components = Composite.get_info<info::device::component_devices>();
+        // Only return composite devices if all of its component
+        // devices are available.
+        return std::all_of(
+            Components.begin(), Components.end(), [&](const device &d) {
+              return std::find(Devs.begin(), Devs.end(), d) != Devs.end();
+            });
+      });
+  return Result;
+}
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/source/detail/device_impl.cpp

@@ -582,6 +582,26 @@ bool device_impl::has(aspect Aspect) const {
       return false;
     }
   }
+  case aspect::ext_oneapi_is_composite: {
+    auto components = get_info<
+        sycl::ext::oneapi::experimental::info::device::component_devices>();
+    // Any device with ext_oneapi_is_composite aspect will have at least two
+    // constituent component devices.
+    return components.size() >= 2;
+  }
+  case aspect::ext_oneapi_is_component: {
+    if (getBackend() != backend::ext_oneapi_level_zero)
+      return false;
+
+    typename sycl_to_pi<device>::type Result;
+    getPlugin()->call<PiApiKind::piDeviceGetInfo>(
+        getHandleRef(),
+        PiInfoCode<
+            ext::oneapi::experimental::info::device::composite_device>::value,
+        sizeof(Result), &Result, nullptr);
+
+    return Result != nullptr;
+  }
   }
   throw runtime_error("This device aspect has not been implemented yet.",
                       PI_ERROR_INVALID_DEVICE);

sycl/source/detail/device_info.hpp

@@ -1085,6 +1085,67 @@ struct get_device_info_impl<
   }
 };
 
+// Specialization for composite devices extension.
+template <>
+struct get_device_info_impl<
+    std::vector<sycl::device>,
+    ext::oneapi::experimental::info::device::component_devices> {
+  static std::vector<sycl::device> get(const DeviceImplPtr &Dev) {
+    if (Dev->getBackend() != backend::ext_oneapi_level_zero)
+      return {};
+    size_t ResultSize = 0;
+    // First call to get DevCount.
+    Dev->getPlugin()->call<PiApiKind::piDeviceGetInfo>(
+        Dev->getHandleRef(),
+        PiInfoCode<
+            ext::oneapi::experimental::info::device::component_devices>::value,
+        0, nullptr, &ResultSize);
+    size_t DevCount = ResultSize / sizeof(pi_device);
+    // Second call to get the list.
+    std::vector<pi_device> Devs(DevCount);
+    Dev->getPlugin()->call<PiApiKind::piDeviceGetInfo>(
+        Dev->getHandleRef(),
+        PiInfoCode<
+            ext::oneapi::experimental::info::device::component_devices>::value,
+        ResultSize, Devs.data(), nullptr);
+    std::vector<sycl::device> Result;
+    const auto &Platform = Dev->getPlatformImpl();
+    for (const auto &d : Devs)
+      Result.push_back(createSyclObjFromImpl<device>(
+          Platform->getOrMakeDeviceImpl(d, Platform)));
+
+    return Result;
+  }
+};
+template <>
+struct get_device_info_impl<
+    sycl::device, ext::oneapi::experimental::info::device::composite_device> {
+  static sycl::device get(const DeviceImplPtr &Dev) {
+    if (Dev->getBackend() != backend::ext_oneapi_level_zero)
+      return {};
+    if (!Dev->has(sycl::aspect::ext_oneapi_is_component))
+      throw sycl::exception(make_error_code(errc::invalid),
+                            "Only devices with aspect::ext_oneapi_is_component "
+                            "can call this function.");
+
+    typename sycl_to_pi<device>::type Result;
+    Dev->getPlugin()->call<PiApiKind::piDeviceGetInfo>(
+        Dev->getHandleRef(),
+        PiInfoCode<
+            ext::oneapi::experimental::info::device::composite_device>::value,
+        sizeof(Result), &Result, nullptr);
+
+    if (Result) {
+      const auto &Platform = Dev->getPlatformImpl();
+      return createSyclObjFromImpl<device>(
+          Platform->getOrMakeDeviceImpl(Result, Platform));
+    }
+    throw sycl::exception(make_error_code(errc::invalid),
+                          "A component with aspect::ext_oneapi_is_component "
+                          "must have a composite device.");
+  }
+};
+
 template <typename Param>
 typename Param::return_type get_device_info(const DeviceImplPtr &Dev) {
   static_assert(is_device_info_desc<Param>::value,
@@ -2041,6 +2102,20 @@ inline float get_device_info_host<
                       PI_ERROR_INVALID_DEVICE);
 }
 
+template <>
+inline std::vector<sycl::device> get_device_info_host<
+    ext::oneapi::experimental::info::device::component_devices>() {
+  throw runtime_error("Host devices cannot be component devices.",
+                      PI_ERROR_INVALID_DEVICE);
+}
+
+template <>
+inline sycl::device get_device_info_host<
+    ext::oneapi::experimental::info::device::composite_device>() {
+  throw runtime_error("Host devices cannot be composite devices.",
+                      PI_ERROR_INVALID_DEVICE);
+}
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/source/feature_test.hpp.in

@@ -55,6 +55,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_PROPERTIES 1
 #define SYCL_EXT_ONEAPI_NATIVE_MATH 1
 #define SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS 1
+#define SYCL_EXT_ONEAPI_COMPOSITE_DEVICE 1
 #define SYCL_EXT_INTEL_DATAFLOW_PIPES 1
 #ifdef __clang__
 #if __has_extension(sycl_extended_atomics)