Initial support for ext_oneapi_composite_device.

devops/dependencies.json

@@ -19,9 +19,9 @@
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     },
     "level_zero": {
-      "github_tag": "v1.14.0",
-      "version": "v1.14.0",
-      "url": "https://github.com/oneapi-src/level-zero/releases/tag/v1.14.0",
+      "github_tag": "v1.15.1",
+      "version": "v1.15.1",
+      "url": "https://github.com/oneapi-src/level-zero/releases/tag/v1.15.1",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     },
     "tbb": {

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -66,6 +66,8 @@ def AspectExt_oneapi_ballot_group : Aspect<"ext_oneapi_ballot_group">;
 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_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">;
@@ -115,7 +117,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_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
+    AspectExt_oneapi_tangle_group, 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/include/sycl/detail/pi.h

@@ -433,6 +433,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

@@ -298,6 +298,16 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_tangle_group__ 0
 #endif
 
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_is_composite__
+// __SYCL_ASPECT(ext_oneapi_is_composite, 58)
+#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, 59)
+#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
@@ -587,3 +597,13 @@
 // __SYCL_ASPECT(ext_oneapi_tangle_group, 57)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_tangle_group__ 0
 #endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_is_composite__
+// __SYCL_ASPECT(ext_oneapi_is_composite, 58)
+#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, 59)
+#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

@@ -52,3 +52,5 @@ __SYCL_ASPECT(ext_oneapi_ballot_group, 54)
 __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_oneapi_is_composite, 58)
+__SYCL_ASPECT(ext_oneapi_is_component, 59)

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

@@ -84,6 +84,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/CMakeLists.txt

@@ -56,15 +56,13 @@ endif()
 if(SYCL_PI_UR_USE_FETCH_CONTENT)
   include(FetchContent)
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit d06ba9d9c539dc712444c6db95dfe1629bd5f7d8
-  # Merge:  e1414e1 35b6a5e
-  # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Thu Jan 11 15:09:01 2024 +0000
-  #     Merge pull request #1035 from Bensuo/cmd-buffer-profiling-l0
-  #     [EXP][CMDBUF] Add support for recovering CommandBuffer profiling
-  #     information
-  set(UNIFIED_RUNTIME_TAG d06ba9d9c539dc712444c6db95dfe1629bd5f7d8)
+  set(UNIFIED_RUNTIME_REPO "https://github.com/maarquitos14/unified-runtime.git")
+  # commit 75648295df39de3027c989299a0cadb018ea26c8 (HEAD -> maronas/ext_composite_device, origin/maronas/ext_composite_device)
+  # Merge: a9746c21 c63ad9b2
+  # Author: Marcos Maronas <marcos.maronas@intel.com>
+  # Date:   Tue Jan 16 09:04:27 2024 -0800
+  #     Merge remote-tracking branch 'intel/origin/main' into maronas/ext_composite_device
+  set(UNIFIED_RUNTIME_TAG 75648295df39de3027c989299a0cadb018ea26c8)
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
     set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")

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?
@@ -958,7 +964,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;
@@ -992,7 +997,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)                                     \
@@ -1254,6 +1258,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

@@ -170,6 +170,7 @@ set(SYCL_SOURCES
     "detail/builtins_relational.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,42 @@
+//==---------- 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>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+std::vector<device> get_composite_devices() {
+  std::vector<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>();
+      // Filter out duplicates.
+      if (std::find(Composites.begin(), Composites.end(), Composite) ==
+          Composites.end())
+        Composites.push_back(Composite);
+    }
+  }
+  std::vector<device> Result;
+  for (const auto &Composite : Composites) {
+    auto Components = Composite.get_info<info::device::component_devices>();
+    // Only return composite devices if all of its component devices are
+    // available.
+    if (std::all_of(Components.begin(), Components.end(), [&](const device &d) {
+          return std::find(Devs.begin(), Devs.end(), d) != Devs.end();
+        })) {
+      Result.push_back(Composite);
+    }
+  }
+  return Result;
+}
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/source/detail/device_impl.cpp

@@ -567,6 +567,28 @@ bool device_impl::has(aspect Aspect) const {
     return (this->getBackend() == backend::ext_oneapi_level_zero) ||
            (this->getBackend() == backend::opencl);
   }
+  case aspect::ext_oneapi_is_composite: {
+    if (getBackend() != backend::ext_oneapi_level_zero)
+      return false;
+    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,66 @@ 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));
+    }
+    assert(false && "Something went wrong.");
+    return {};
+  }
+};
+
 template <typename Param>
 typename Param::return_type get_device_info(const DeviceImplPtr &Dev) {
   static_assert(is_device_info_desc<Param>::value,
@@ -2041,6 +2101,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

@@ -54,6 +54,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)