diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index ed2d05cab29fb..72bc18e1a1385 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/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. diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_composite_device.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_composite_device.asciidoc similarity index 96% rename from sycl/doc/extensions/proposed/sycl_ext_oneapi_composite_device.asciidoc rename to sycl/doc/extensions/experimental/sycl_ext_oneapi_composite_device.asciidoc index 27aa6ea0e1b79..2f81c4a808783 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_composite_device.asciidoc +++ b/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 diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 0c02a7705b51d..5059125da7646 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/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 { diff --git a/sycl/include/sycl/device_aspect_macros.hpp b/sycl/include/sycl/device_aspect_macros.hpp index 0620257778ed4..489163c556223 100644 --- a/sycl/include/sycl/device_aspect_macros.hpp +++ b/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 diff --git a/sycl/include/sycl/ext/oneapi/experimental/composite_device.hpp b/sycl/include/sycl/ext/oneapi/experimental/composite_device.hpp new file mode 100644 index 0000000000000..de33fadfcb521 --- /dev/null +++ b/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 + +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +__SYCL_EXPORT std::vector get_composite_devices(); +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index e1517b59ed250..a4a2296609e56 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/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) diff --git a/sycl/include/sycl/info/ext_oneapi_device_traits.def b/sycl/include/sycl/info/ext_oneapi_device_traits.def index 7a4668dbbdb6a..1842d66760e85 100644 --- a/sycl/include/sycl/info/ext_oneapi_device_traits.def +++ b/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, + 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 diff --git a/sycl/include/sycl/platform.hpp b/sycl/include/sycl/platform.hpp index 81a07ba3ab1f0..8cffdcd99e67e 100644 --- a/sycl/include/sycl/platform.hpp +++ b/sycl/include/sycl/platform.hpp @@ -188,6 +188,8 @@ class __SYCL_EXPORT platform : public detail::OwnerLessBase { /// \return the default context context ext_oneapi_get_default_context() const; + std::vector ext_oneapi_get_composite_devices() const; + private: pi_native_handle getNative() const; diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index 7a62f12b38335..be923a7bfc14a 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -85,6 +85,7 @@ #include #include #include +#include #include #include #include diff --git a/sycl/plugins/level_zero/CMakeLists.txt b/sycl/plugins/level_zero/CMakeLists.txt index 5c24988813c1f..3c41fe7953c14 100644 --- a/sycl/plugins/level_zero/CMakeLists.txt +++ b/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) diff --git a/sycl/plugins/unified_runtime/pi2ur.hpp b/sycl/plugins/unified_runtime/pi2ur.hpp index ed232c74e3c73..3de674126eaab 100644 --- a/sycl/plugins/unified_runtime/pi2ur.hpp +++ b/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(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; diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index d5789be86bde9..0da1c5fa0cea2 100644 --- a/sycl/source/CMakeLists.txt +++ b/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" diff --git a/sycl/source/detail/composite_device/composite_device.cpp b/sycl/source/detail/composite_device/composite_device.cpp new file mode 100644 index 0000000000000..6c57eb3015df1 --- /dev/null +++ b/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 +#include + +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +std::vector 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 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(); + Composites.insert(Composite); + } + } + std::vector Result; + std::copy_if( + Composites.begin(), Composites.end(), std::back_inserter(Result), + [&](const device &Composite) { + auto Components = Composite.get_info(); + // 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 diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 1aa7f506e4fea..a78daf5fe0f28 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/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::type Result; + getPlugin()->call( + 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); diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp index bd2cc473d24b3..8434fb79e4a88 100644 --- a/sycl/source/detail/device_info.hpp +++ b/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, + ext::oneapi::experimental::info::device::component_devices> { + static std::vector 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( + 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 Devs(DevCount); + Dev->getPlugin()->call( + Dev->getHandleRef(), + PiInfoCode< + ext::oneapi::experimental::info::device::component_devices>::value, + ResultSize, Devs.data(), nullptr); + std::vector Result; + const auto &Platform = Dev->getPlatformImpl(); + for (const auto &d : Devs) + Result.push_back(createSyclObjFromImpl( + 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::type Result; + Dev->getPlugin()->call( + Dev->getHandleRef(), + PiInfoCode< + ext::oneapi::experimental::info::device::composite_device>::value, + sizeof(Result), &Result, nullptr); + + if (Result) { + const auto &Platform = Dev->getPlatformImpl(); + return createSyclObjFromImpl( + 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::return_type get_device_info(const DeviceImplPtr &Dev) { static_assert(is_device_info_desc::value, @@ -2041,6 +2102,20 @@ inline float get_device_info_host< PI_ERROR_INVALID_DEVICE); } +template <> +inline std::vector 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 diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index f53459d7f9bee..f4575c22b8c6c 100644 --- a/sycl/source/feature_test.hpp.in +++ b/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) diff --git a/sycl/source/platform.cpp b/sycl/source/platform.cpp index 170dc08785b5a..5cc2a49801902 100644 --- a/sycl/source/platform.cpp +++ b/sycl/source/platform.cpp @@ -95,6 +95,41 @@ context platform::ext_oneapi_get_default_context() const { return detail::createSyclObjFromImpl(It->second); } +std::vector platform::ext_oneapi_get_composite_devices() const { + // Only GPU architectures can be composite devices. + auto GPUDevices = get_devices(info::device_type::gpu); + // Using ZE_FLAT_DEVICE_HIERARCHY=COMBINED, we receive tiles as devices, which + // are component devices. Thus, we need to get the composite device for each + // of the component devices, and filter out duplicates. + std::vector Composites; + std::vector Result; + for (auto &Dev : GPUDevices) { + if (!Dev.has(sycl::aspect::ext_oneapi_is_component)) + continue; + + auto Composite = Dev.get_info< + sycl::ext::oneapi::experimental::info::device::composite_device>(); + if (std::find(Result.begin(), Result.end(), Composite) == Result.end()) + Composites.push_back(Composite); + } + for (const auto &Composite : Composites) { + auto Components = Composite.get_info< + sycl::ext::oneapi::experimental::info::device::component_devices>(); + // Checking whether Components are GPU device is not enough, we need to + // check if they are in the list of available devices returned by + // `get_devices()`, because we cannot return a Composite device unless all + // of its components are available too. + size_t ComponentsFound = std::count_if( + Components.begin(), Components.end(), [&](const device &d) { + return std::find(GPUDevices.begin(), GPUDevices.end(), d) != + GPUDevices.end(); + }); + if (ComponentsFound == Components.size()) + Result.push_back(Composite); + } + return Result; +} + namespace detail { void enable_ext_oneapi_default_context(bool Val) { diff --git a/sycl/test-e2e/CompositeDevice/composite_device.cpp b/sycl/test-e2e/CompositeDevice/composite_device.cpp new file mode 100644 index 0000000000000..5ffbdb05442d3 --- /dev/null +++ b/sycl/test-e2e/CompositeDevice/composite_device.cpp @@ -0,0 +1,235 @@ +// RUN: %{build} -o %t.out +// RUN: env ZE_FLAT_DEVICE_HIERARCHY=COMBINED %{run} %t.out +// RUN: env ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE %{run} %t.out +// RUN: env ZE_FLAT_DEVICE_HIERARCHY=FLAT %{run} %t.out +// UNSUPPORTED: (windows && level_zero) + +#include + +#ifdef SYCL_EXT_ONEAPI_COMPOSITE_DEVICE + +using namespace sycl::ext::oneapi::experimental; + +bool isL0Backend(sycl::backend backend) { + return (backend == sycl::backend::ext_oneapi_level_zero); +} + +bool isCombinedMode() { + char *Mode = nullptr; + bool Res = false; +#ifdef _WIN32 + size_t Size = 0; + auto Err = _dupenv_s(&Mode, &Size, "ZE_FLAT_DEVICE_HIERARCHY"); + Res = (Mode != nullptr) && (std::strcmp(Mode, "COMBINED") == 0); + free(Mode); +#else + Mode = std::getenv("ZE_FLAT_DEVICE_HIERARCHY"); + Res = (Mode != nullptr) && (std::strcmp(Mode, "COMBINED") == 0); +#endif + return Res; +} + +int main() { + sycl::queue q; + bool IsCombined = isCombinedMode(); + auto Platforms = sycl::platform::get_platforms(); + + // Check that device::get_devices() and platform::get_devices() do not return + // composite devices. + { + auto Devs = sycl::device::get_devices(); + for (const auto &D : Devs) { + assert(!D.has(sycl::aspect::ext_oneapi_is_composite)); + + // If ZE_FLAT_DEVICE_HIERARCHY != COMBINED, + // sycl::aspect::ext_oneapi_is_component must be false. + assert(IsCombined || !D.has(sycl::aspect::ext_oneapi_is_component)); + } + + for (const auto &P : Platforms) { + bool IsL0 = isL0Backend(P.get_backend()); + if (!IsL0) + continue; + + Devs = P.get_devices(); + for (const auto &D : Devs) { + assert(!D.has(sycl::aspect::ext_oneapi_is_composite)); + + // If ZE_FLAT_DEVICE_HIERARCHY != COMBINED, + // sycl::aspect::ext_oneapi_is_component must be false. + assert(!(!IsCombined && D.has(sycl::aspect::ext_oneapi_is_component))); + } + } + } + + // Check that: + // A. The free function get_composite_devices returns all of the composite + // devices across all platforms. + // B. The member function platform::ext_oneapi_get_composite_devices returns + // the composite devices within the given platform. + // C. The APIs defined in this extension are only useful when using the + // Level Zero backend, and they are only useful when the Level Zero + // environment variable ZE_FLAT_DEVICE_HIERARCHY=COMBINED is set. The + // APIs may be called even when using other backends, but they will + // return an empty list of composite devices. + // D. The execution environment for a SYCL application has a fixed number of + // composite devices which does not vary as the application executes. As + // a result, each call to these functions returns the same set of device + // objects, and the order of those objects does not vary between calls. + { + std::vector AllCompositeDevs = get_composite_devices(); + std::vector CombinedCompositeDevs; + for (const auto &P : Platforms) { + auto CompositeDevs = P.ext_oneapi_get_composite_devices(); + bool IsL0 = isL0Backend(P.get_backend()); + // Check C. + assert(CompositeDevs.empty() || (IsL0 && IsCombined)); + + for (const auto &D : CompositeDevs) { + if (std::find(CombinedCompositeDevs.begin(), + CombinedCompositeDevs.end(), + D) == CombinedCompositeDevs.end()) + CombinedCompositeDevs.push_back(D); + } + } + // Check A. and B. + assert(AllCompositeDevs.size() == CombinedCompositeDevs.size()); + assert(std::all_of(AllCompositeDevs.begin(), AllCompositeDevs.end(), + [&](const sycl::device &D) { + const bool Found = + std::find(CombinedCompositeDevs.begin(), + CombinedCompositeDevs.end(), + D) != CombinedCompositeDevs.end(); + return Found && + D.has(sycl::aspect::ext_oneapi_is_composite); + })); + + // Check D. + std::vector AllCompositeDevs2 = get_composite_devices(); + std::vector CombinedCompositeDevs2; + for (const auto &P : Platforms) { + auto CompositeDevs = P.ext_oneapi_get_composite_devices(); + bool IsL0 = isL0Backend(P.get_backend()); + // Check C. + assert(CompositeDevs.empty() || (IsL0 && IsCombined)); + + for (const auto &D : CompositeDevs) { + if (std::find(CombinedCompositeDevs2.begin(), + CombinedCompositeDevs2.end(), + D) == CombinedCompositeDevs2.end()) + CombinedCompositeDevs2.push_back(D); + } + } + assert(AllCompositeDevs.size() == AllCompositeDevs2.size()); + assert(CombinedCompositeDevs.size() == CombinedCompositeDevs2.size()); + for (size_t i = 0; i < AllCompositeDevs.size(); ++i) { + assert(AllCompositeDevs[i] == AllCompositeDevs2[i]); + assert(CombinedCompositeDevs[i] == CombinedCompositeDevs2[i]); + } + } + + // Check that device::info::component_devices: + // A. Returns the set of component devices that are contained by a composite + // device (at least 2). + // B. If "this" device is not a composite device, returns an empty vector. + { + auto Devs = sycl::device::get_devices(); + for (const auto &D : Devs) { + // Check B. + assert(!D.has(sycl::aspect::ext_oneapi_is_composite)); + auto Components = D.get_info(); + assert(Components.empty()); + + // Check A. + auto IsComponent = D.has(sycl::aspect::ext_oneapi_is_component); + // A device can be neither composite nor component. This happens when + // there are not multiple tiles in a single card. + if (IsComponent) { + auto Composite = D.get_info(); + Components = Composite.get_info(); + assert(Components.size() >= 2); + } + } + } + + // Check that device::info::composite_device: + // A. Returns the composite device which contains this component device. + // B. Since the set of composite devices if fixed, returns a device object + // which is a copy of one of the device objects returned by + // get_composite_devices. + // C. Throws a synchronous exception with the errc::invalid error code if + // "this" device does not have aspect::ext_oneapi_is_component. + { + auto Devs = sycl::device::get_devices(); + for (const auto &D : Devs) { + bool IsL0 = isL0Backend(D.get_backend()); + if (!IsL0 || !IsCombined) + continue; + // Check A. + assert(!D.has(sycl::aspect::ext_oneapi_is_composite)); + auto IsComponent = D.has(sycl::aspect::ext_oneapi_is_component); + // A device can be neither composite nor component. This happens when + // there are not multiple tiles in a single card. + if (IsComponent) { + auto Composite = D.get_info(); + assert(Composite.has(sycl::aspect::ext_oneapi_is_composite)); + // Check B. + std::vector AllCompositeDevs = get_composite_devices(); + assert(std::find(AllCompositeDevs.begin(), AllCompositeDevs.end(), + Composite) != AllCompositeDevs.end()); + // Check C. + assert(!Composite.has(sycl::aspect::ext_oneapi_is_component)); + try { + auto Invalid = Composite.get_info(); + assert(false && "Exception expected."); + } catch (sycl::exception &E) { + assert(E.code() == sycl::errc::invalid && + "errc should be errc::invalid"); + } + } + } + } + + // Check that ext_oneapi_is_component applies only to a root device that is a + // direct component of some composite device. A sub-device will not have this + // aspect even if its parent is a component device. + { + auto Devs = sycl::device::get_devices(); + for (const auto &D : Devs) { + bool IsL0 = isL0Backend(D.get_backend()); + if (!IsL0 || !IsCombined) + continue; + + auto PartitionProperties = + D.get_info(); + if (PartitionProperties.empty()) + continue; + + std::vector SubDevices; + for (const auto &PartitionProperty : PartitionProperties) { + if (PartitionProperty == + sycl::info::partition_property::partition_equally) { + size_t CompUnits = 2; + SubDevices = D.create_sub_devices< + sycl::info::partition_property::partition_equally>(CompUnits); + } else if (PartitionProperty == + sycl::info::partition_property::partition_by_counts) { + SubDevices = D.create_sub_devices< + sycl::info::partition_property::partition_by_counts>( + std::vector{2}); + } else if (PartitionProperty == sycl::info::partition_property:: + partition_by_affinity_domain) { + SubDevices = D.create_sub_devices< + sycl::info::partition_property::partition_by_affinity_domain>( + sycl::info::partition_affinity_domain::numa); + } + } + + for (const auto &SubDevice : SubDevices) { + assert(!SubDevice.has(sycl::aspect::ext_oneapi_is_component)); + } + } + } +} + +#endif // SYCL_EXT_ONEAPI_COMPOSITE_DEVICE diff --git a/sycl/test-e2e/CompositeDevice/device_selector_test.cpp b/sycl/test-e2e/CompositeDevice/device_selector_test.cpp new file mode 100644 index 0000000000000..548403a7c2175 --- /dev/null +++ b/sycl/test-e2e/CompositeDevice/device_selector_test.cpp @@ -0,0 +1,27 @@ +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:0 ZE_FLAT_DEVICE_HIERARCHY=COMBINED %t.out +// REQUIRES: level_zero + +#include + +#ifdef SYCL_EXT_ONEAPI_COMPOSITE_DEVICE + +using namespace sycl::ext::oneapi::experimental; + +int main() { + sycl::queue q; + auto Platforms = sycl::platform::get_platforms(); + + // Check that we do not expose a composite device unless it represents all of + // the tiles on a card. Since we are setting ONEAPI_DEVICE_SELECTOR to use + // only a single tile, both get_composite_devices() and + // platform::ext_oneapi_get_composite_devices() should return an empty vector. + std::vector AllCompositeDevs = get_composite_devices(); + assert(AllCompositeDevs.empty()); + for (const auto &P : Platforms) { + auto CompositeDevs = P.ext_oneapi_get_composite_devices(); + assert(CompositeDevs.empty()); + } +} + +#endif // SYCL_EXT_ONEAPI_COMPOSITE_DEVICE diff --git a/sycl/test-e2e/CompositeDevice/run_on_composite_device.cpp b/sycl/test-e2e/CompositeDevice/run_on_composite_device.cpp new file mode 100644 index 0000000000000..4b6df96cf4bf0 --- /dev/null +++ b/sycl/test-e2e/CompositeDevice/run_on_composite_device.cpp @@ -0,0 +1,36 @@ +// RUN: %{build} -o %t.out +// RUN: env ZE_FLAT_DEVICE_HIERARCHY=COMBINED %{run} %t.out + +#include + +#ifdef SYCL_EXT_ONEAPI_COMPOSITE_DEVICE + +using namespace sycl::ext::oneapi::experimental; + +int main() { + std::vector CompositeDevs = get_composite_devices(); + for (const auto &Composite : CompositeDevs) { + // Check that `Composite` is indeed a composite device. + assert(Composite.has(sycl::aspect::ext_oneapi_is_composite)); + + // Create a new context and queue with `Composite` and run a test kernel. + sycl::context CompositeContext(Composite); + sycl::queue q(CompositeContext, Composite); + constexpr size_t N = 1024; + std::vector TestData(N, 0); + { + sycl::buffer TestData_b(TestData.data(), sycl::range<1>{TestData.size()}); + q.submit([&](sycl::handler &cgh) { + sycl::accessor TestData_acc{TestData_b, cgh}; + cgh.single_task([=]() { + for (size_t i = 0; i < N; ++i) + TestData_acc[i] = i; + }); + }); + } + for (size_t i = 0; i < N; ++i) + assert(TestData[i] == i); + } +} + +#endif // SYCL_EXT_ONEAPI_COMPOSITE_DEVICE diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index ac0385ca3bb58..8ea9beb5b0cb4 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3691,6 +3691,7 @@ _ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmRKNS3_16image_desc _ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmRKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmmmjRKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmmmjRKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental21get_composite_devicesEv _ZN4sycl3_V13ext6oneapi12experimental22get_image_channel_typeENS3_16image_mem_handleERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental22get_image_channel_typeENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental22get_image_num_channelsENS3_16image_mem_handleERKNS0_5queueE @@ -4307,6 +4308,8 @@ _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6de _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi1EEEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi2EEEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi3EEEEENT_11return_typeEv +_ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device16composite_deviceEEENT_11return_typeEv +_ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device17component_devicesEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device19matrix_combinationsEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device21image_row_pitch_alignEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device21mipmap_max_anisotropyEEENT_11return_typeEv @@ -4487,6 +4490,8 @@ _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device13graph_s _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi1EEEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi2EEEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi3EEEEENS0_6detail19is_device_info_descIT_E11return_typeEv +_ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device16composite_deviceEEENS0_6detail19is_device_info_descIT_E11return_typeEv +_ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device17component_devicesEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device19matrix_combinationsEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device21image_row_pitch_alignEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device21mipmap_max_anisotropyEEENS0_6detail19is_device_info_descIT_E11return_typeEv @@ -4763,6 +4768,7 @@ _ZNK4sycl3_V18platform11get_backendEv _ZNK4sycl3_V18platform11get_devicesENS0_4info11device_typeE _ZNK4sycl3_V18platform13has_extensionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZNK4sycl3_V18platform30ext_oneapi_get_default_contextEv +_ZNK4sycl3_V18platform32ext_oneapi_get_composite_devicesEv _ZNK4sycl3_V18platform3getEv _ZNK4sycl3_V18platform3hasENS0_6aspectE _ZNK4sycl3_V18platform7is_hostEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index b5d0b3f3de707..7b1d18f198967 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -49,6 +49,10 @@ ??$get_info@Ucompile_num_sub_groups@kernel_device_specific@info@_V1@sycl@@@kernel@_V1@sycl@@QEBAIAEBVdevice@12@@Z ??$get_info@Ucompile_sub_group_size@kernel_device_specific@info@_V1@sycl@@@kernel@_V1@sycl@@QEBAIAEBVdevice@12@@Z ??$get_info@Ucompile_work_group_size@kernel_device_specific@info@_V1@sycl@@@kernel@_V1@sycl@@QEBA?AV?$range@$02@12@AEBVdevice@12@@Z +??$get_info@Ucomponent_devices@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ +??$get_info@Ucomponent_devices@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ +??$get_info@Ucomposite_device@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBA?AV012@XZ +??$get_info@Ucomposite_device@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AVdevice@23@XZ ??$get_info@Ucontext@kernel@info@_V1@sycl@@@kernel@_V1@sycl@@QEBA?AVcontext@12@XZ ??$get_info@Ucontext@queue@info@_V1@sycl@@@queue@_V1@sycl@@QEBA?AVcontext@12@XZ ??$get_info@Udevice@queue@info@_V1@sycl@@@queue@_V1@sycl@@QEBA?AVdevice@12@XZ @@ -1047,6 +1051,7 @@ ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_fill_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEAVSYCLMemObjI@234@PEAX_KPEBDIV?$range@$02@34@6V?$id@$02@34@IV?$vector@IV?$allocator@I@std@@@6@PEAI@Z ?ext_oneapi_fill_usm_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEAX_KHV?$vector@IV?$allocator@I@std@@@6@PEAI@Z +?ext_oneapi_get_composite_devices@platform@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ ?ext_oneapi_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ ?ext_oneapi_get_kernel@kernel_bundle_plain@detail@_V1@sycl@@QEAA?AVkernel@34@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_get_last_event@queue@_V1@sycl@@QEBA?AVevent@23@XZ @@ -1233,6 +1238,7 @@ ?get_channel_order@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AW4image_channel_order@56@XZ ?get_channel_type@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AW4image_channel_type@56@XZ ?get_cl_code@exception@_V1@sycl@@QEBAHXZ +?get_composite_devices@experimental@oneapi@ext@_V1@sycl@@YA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ ?get_context@exception@_V1@sycl@@QEBA?AVcontext@23@XZ ?get_context@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVcontext@56@XZ ?get_context@kernel@_V1@sycl@@QEBA?AVcontext@23@XZ