diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index ff06de04089c1..ed2d05cab29fb 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -66,6 +66,7 @@ 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_intel_matrix : Aspect<"ext_intel_matrix">; // Deprecated aspects def AspectInt64_base_atomics : Aspect<"int64_base_atomics">; def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">; @@ -115,7 +116,7 @@ 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_intel_matrix], []>; // 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/include/sycl/device_aspect_macros.hpp b/sycl/include/sycl/device_aspect_macros.hpp index 3f515ddc564af..0620257778ed4 100644 --- a/sycl/include/sycl/device_aspect_macros.hpp +++ b/sycl/include/sycl/device_aspect_macros.hpp @@ -298,6 +298,11 @@ #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_tangle_group__ 0 #endif +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_intel_matrix__ +// __SYCL_ASPECT(ext_intel_matrix, 58) +#define __SYCL_ALL_DEVICES_HAVE_ext_intel_matrix__ 0 +#endif + #ifndef __SYCL_ANY_DEVICE_HAS_host__ // __SYCL_ASPECT(host, 0) #define __SYCL_ANY_DEVICE_HAS_host__ 0 @@ -587,3 +592,8 @@ // __SYCL_ASPECT(ext_oneapi_tangle_group, 57) #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_tangle_group__ 0 #endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_intel_matrix__ +// __SYCL_ASPECT(ext_intel_matrix, 58) +#define __SYCL_ANY_DEVICE_HAS_ext_intel_matrix__ 0 +#endif diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index 31ccd0d038c4a..e1517b59ed250 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/sycl/include/sycl/info/aspects.def @@ -52,3 +52,4 @@ __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_intel_matrix, 58) diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 5d5b98e4cf02b..1aa7f506e4fea 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -567,6 +567,21 @@ bool device_impl::has(aspect Aspect) const { return (this->getBackend() == backend::ext_oneapi_level_zero) || (this->getBackend() == backend::opencl); } + case aspect::ext_intel_matrix: { + using arch = sycl::ext::oneapi::experimental::architecture; + const std::vector supported_archs = { + arch::intel_cpu_spr, arch::intel_gpu_pvc, arch::intel_gpu_dg2_g10, + arch::intel_gpu_dg2_g11, arch::intel_gpu_dg2_g12}; + try { + return std::any_of( + supported_archs.begin(), supported_archs.end(), + [=](const arch a) { return this->extOneapiArchitectureIs(a); }); + } catch (const sycl::exception &) { + // If we're here it means the device does not support architecture + // querying + return false; + } + } } throw runtime_error("This device aspect has not been implemented yet.", PI_ERROR_INVALID_DEVICE); diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index ba102c2477061..8ecaa6a2bfc62 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -234,7 +234,8 @@ class device_impl { std::string getDeviceName() const; - bool extOneapiArchitectureIs(ext::oneapi::experimental::architecture Arch) { + bool + extOneapiArchitectureIs(ext::oneapi::experimental::architecture Arch) const { return Arch == getDeviceArch(); } diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 3168dc2d75839..f53459d7f9bee 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -99,6 +99,7 @@ inline namespace _V1 { #define SYCL_EXT_INTEL_FP_CONTROL 1 #define SYCL_EXT_ONEAPI_NON_UNIFORM_GROUPS 1 #define SYCL_EXT_ONEAPI_IN_ORDER_QUEUE_EVENTS 1 +#define SYCL_EXT_INTEL_MATRIX 1 #ifndef __has_include #define __has_include(x) 0 diff --git a/sycl/test-e2e/Basic/AMX_aspect.cpp b/sycl/test-e2e/Basic/AMX_aspect.cpp new file mode 100644 index 0000000000000..09bcb89111061 --- /dev/null +++ b/sycl/test-e2e/Basic/AMX_aspect.cpp @@ -0,0 +1,40 @@ +// RUN: %{build} -o %t.out +// RUN: %t.out +// + +//==--------------- AMX_aspect.cpp - SYCL device test +//------------------------==// +// +// Checks that the has(aspect) method on a device returns the correct answer +// when queried about ext_intel_matrix AMX aspect. +// +// 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 + +using namespace sycl; +using arch = sycl::ext::oneapi::experimental::architecture; +int main() { + const std::vector supported_archs = { + arch::intel_cpu_spr, arch::intel_gpu_pvc, arch::intel_gpu_dg2_g10, + arch::intel_gpu_dg2_g11, arch::intel_gpu_dg2_g12}; + for (const auto &plt : platform::get_platforms()) { + for (auto &dev : plt.get_devices()) { + try { + if (std::any_of(supported_archs.begin(), supported_archs.end(), + [&](const auto &a) { + return dev.ext_oneapi_architecture_is(a); + })) { + assert(dev.has(sycl::aspect::ext_intel_matrix)); + } + } catch (sycl::exception &) { + } + } + } + return 0; +}