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

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

@@ -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/CMakeLists.txt

@@ -56,14 +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 79c28d0f0713f58358d5080653d95803fd131749
- # Merge: 25e0b603 45d76b78
- # Author: aarongreig <aaron.greig@codeplay.com>
- # Date: Fri Jan 12 16:14:44 2024 +0000
- # Merge pull request #1186 from hdelan/device-global-hip
- # [HIP] Add support for global variable read write
- set(UNIFIED_RUNTIME_TAG 79c28d0f0713f58358d5080653d95803fd131749)
+ set(UNIFIED_RUNTIME_REPO "https://github.com/maarquitos14/unified-runtime.git")
+ # commit 75648295df39de3027c989299a0cadb018ea26c8
+ # 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?
@@ -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

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

@@ -567,6 +567,26 @@ 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: {
+ 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)