[SYCL] Add support for ext_intel_matrix aspect

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.

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

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)

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<arch> 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);

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();
   }
 

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

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 <iostream>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using arch = sycl::ext::oneapi::experimental::architecture;
+int main() {
+  const std::vector<arch> 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;
+}