[SYCL] Add support for ext_intel_matrix aspect (#12512)

This PR adds the aspect for the sycl_ext_intel_matrix extension and
allows code to verify whether the device supports the advanced matrix
API. The aspect has been added to the aspect definition file and the
corresponding macro has been added to the feature test file set to 1
meaning supported.

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;
+}