[SYCL] Add sycl ext intel kernel queries extension (#16834)

This adds a `sycl_ext_intel_kernel_queries` extension for
target-specific low-level information retrieval. In a nutshell, the
extension adds compiled kernel traits for the intel backend. The traits
can be queried using the usual `kernel::get_info()` calls (and
`get_kernel_info`). Each query gets its own aspect to tell whether the
query is supported on the device. This PR adds
`ext_intel_spill_memory_size` aspect and an implementation for querying
the spilled memory size as reported by L0.

---------

Co-authored-by: Greg Lueck <[email protected]>
Co-authored-by: aelovikov-intel <[email protected]>

Author: 3 people

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -86,6 +86,7 @@ def AspectExt_oneapi_cuda_cluster_group : Aspect<"ext_oneapi_cuda_cluster_group"
 def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
 def AspectExt_oneapi_atomic16 : Aspect<"ext_oneapi_atomic16">;
 def AspectExt_oneapi_virtual_functions : Aspect<"ext_oneapi_virtual_functions">;
+def AspectExt_intel_spill_memory_size : Aspect<"ext_intel_spill_memory_size">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -150,7 +151,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group, 
     AspectExt_intel_fpga_task_sequence,
     AspectExt_oneapi_atomic16,
-    AspectExt_oneapi_virtual_functions],
+    AspectExt_oneapi_virtual_functions,
+    AspectExt_intel_spill_memory_size],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

sycl/doc/extensions/supported/sycl_ext_intel_kernel_queries.asciidoc

@@ -0,0 +1,139 @@
+= sycl_ext_intel_kernel_queries
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
+:endnote: —{nbsp}end{nbsp}note
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2025 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
+permission by Khronos.
+
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 9 specification.
+All references below to the "core SYCL specification" or to section numbers in
+the SYCL specification refer to that revision.
+
+
+== Status
+
+This extension is implemented and fully supported by {dpcpp}.
+
+
+== Overview
+
+This extension contains a collection of queries that provide low-level
+information about kernels.
+These queries generally forward directly to the backend and expose concepts that
+are specific to a particular implementation.
+As a result, these queries may not be supported for all devices.
+Each query has an associate device aspect, which tells whether the query is
+supported on that device.
+
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.
+An implementation supporting this extension must predefine the macro
+`SYCL_EXT_INTEL_KERNEL_QUERIES` to one of the values defined in the table below.
+Applications can test for the existence of this macro to determine if the
+implementation supports this feature, or applications can test the macro's value
+to determine which of the extension's features the implementation supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|Initial version of this extension.
+|===
+
+=== Spill memory size
+
+This query returns the kernel's spill memory size that is allocated by the
+compiler, as reported by Level Zero.
+
+==== New device aspect
+
+This extension adds the following new device aspect.
+
+[source,c++]
+----
+namespace sycl {
+
+enum class aspect {
+  ext_intel_spill_memory_size
+
+  // ...
+};
+
+}
+----
+
+'''
+
+`*ext_intel_spill_memory_size*`
+
+Indicates that the `spill_memory_size` kernel information descriptor may be used
+to query kernels for this device.
+
+'''
+
+==== New device specific kernel information descriptor
+
+This extension adds the following information descriptor that can be used with
+`kernel::get_info(const device&)`.
+
+'''
+
+`*ext::intel::info::kernel_device_specific::spill_memory_size*`
+
+[source,c++]
+----
+namespace sycl::ext::intel::info::kernel_device_specific {
+struct spill_memory_size {
+  using return_type = size_t;
+};
+} // namespace sycl::ext::intel::info::kernel_device_specific
+----
+
+_Remarks:_ Template parameter to `kernel::get_info(const device&)`.
+
+_Returns:_ The spill memory size that is allocated by the compiler for this
+kernel for the given device.
+
+_Throws:_ An `exception` with the `errc::feature_not_supported` error code if
+the device does not have `aspect::ext_intel_spill_memory_size`.
+
+'''

sycl/include/sycl/detail/info_desc_helpers.hpp

@@ -88,6 +88,15 @@ struct IsKernelInfo<info::kernel_device_specific::ext_codeplay_num_regs>
 #include <sycl/info/kernel_device_specific_traits.def>
 #undef __SYCL_PARAM_TRAITS_SPEC
 
+#define __SYCL_PARAM_TRAITS_SPEC(Namespace, DescType, Desc, ReturnT, UrCode)   \
+  template <>                                                                  \
+  struct is_##DescType##_info_desc<Namespace::info::DescType::Desc>            \
+      : std::true_type {                                                       \
+    using return_type = Namespace::info::DescType::Desc::return_type;          \
+  };
+#include <sycl/info/ext_intel_kernel_info_traits.def>
+#undef __SYCL_PARAM_TRAITS_SPEC
+
 #define __SYCL_PARAM_TRAITS_SPEC(DescType, Desc, ReturnT, UrCode)              \
   template <>                                                                  \
   struct is_##DescType##_info_desc<info::DescType::Desc> : std::true_type {    \

sycl/include/sycl/info/aspects.def

@@ -72,3 +72,4 @@ __SYCL_ASPECT(ext_oneapi_bindless_images_sample_1d_usm, 78)
 __SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 79)
 __SYCL_ASPECT(ext_oneapi_atomic16, 80)
 __SYCL_ASPECT(ext_oneapi_virtual_functions, 81)
+__SYCL_ASPECT(ext_intel_spill_memory_size, 82)

sycl/include/sycl/info/ext_intel_kernel_info_traits.def

@@ -0,0 +1 @@
+__SYCL_PARAM_TRAITS_SPEC(ext::intel, kernel_device_specific, spill_memory_size, size_t, UR_KERNEL_INFO_SPILL_MEM_SIZE)

sycl/include/sycl/info/info_desc.hpp

@@ -261,6 +261,7 @@ struct work_item_progress_capabilities;
 } // namespace ext::oneapi::experimental::info::device
 #include <sycl/info/ext_codeplay_device_traits.def>
 #include <sycl/info/ext_intel_device_traits.def>
+#include <sycl/info/ext_intel_kernel_info_traits.def>
 #include <sycl/info/ext_oneapi_device_traits.def>
 #include <sycl/info/ext_oneapi_kernel_queue_specific_traits.def>
 

sycl/source/detail/device_impl.cpp

@@ -786,6 +786,11 @@ bool device_impl::has(aspect Aspect) const {
                           BE == sycl::backend::opencl;
     return (is_cpu() || is_gpu()) && isCompatibleBE;
   }
+  case aspect::ext_intel_spill_memory_size: {
+    backend BE = getBackend();
+    bool isCompatibleBE = BE == sycl::backend::ext_oneapi_level_zero;
+    return is_gpu() && isCompatibleBE;
+  }
   }
 
   return false; // This device aspect has not been implemented yet.

sycl/source/detail/kernel_impl.hpp

@@ -416,6 +416,23 @@ inline typename syclex::info::kernel_queue_specific::max_num_work_groups::
                                DynamicLocalMemorySize);
 }
 
+template <>
+inline typename ext::intel::info::kernel_device_specific::spill_memory_size::
+    return_type
+    kernel_impl::get_info<
+        ext::intel::info::kernel_device_specific::spill_memory_size>(
+        const device &Device) const {
+  if (!Device.has(aspect::ext_intel_spill_memory_size))
+    throw exception(
+        make_error_code(errc::feature_not_supported),
+        "This device does not have the ext_intel_spill_memory_size aspect");
+
+  return get_kernel_device_specific_info<
+      ext::intel::info::kernel_device_specific::spill_memory_size>(
+      this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(),
+      getAdapter());
+}
+
 template <>
 inline typename syclex::info::kernel_queue_specific::max_work_group_size::
     return_type

sycl/source/detail/kernel_info.hpp

@@ -146,6 +146,58 @@ uint32_t get_kernel_device_specific_info_with_input(ur_kernel_handle_t Kernel,
 
   return Result;
 }
+
+template <>
+inline ext::intel::info::kernel_device_specific::spill_memory_size::return_type
+get_kernel_device_specific_info<
+    ext::intel::info::kernel_device_specific::spill_memory_size>(
+    ur_kernel_handle_t Kernel, ur_device_handle_t Device,
+    const AdapterPtr &Adapter) {
+  size_t ResultSize = 0;
+
+  // First call to get the number of device images
+  Adapter->call<UrApiKind::urKernelGetInfo>(
+      Kernel, UR_KERNEL_INFO_SPILL_MEM_SIZE, 0, nullptr, &ResultSize);
+
+  size_t DeviceCount = ResultSize / sizeof(uint32_t);
+
+  // Second call to retrieve the data
+  std::vector<uint32_t> Device2SpillMap(DeviceCount);
+  Adapter->call<UrApiKind::urKernelGetInfo>(
+      Kernel, UR_KERNEL_INFO_SPILL_MEM_SIZE, ResultSize, Device2SpillMap.data(),
+      nullptr);
+
+  ur_program_handle_t Program;
+  Adapter->call<UrApiKind::urKernelGetInfo>(Kernel, UR_KERNEL_INFO_PROGRAM,
+                                            sizeof(ur_program_handle_t),
+                                            &Program, nullptr);
+  // Retrieve the associated device list
+  size_t URDevicesSize = 0;
+  Adapter->call<UrApiKind::urProgramGetInfo>(Program, UR_PROGRAM_INFO_DEVICES,
+                                             0, nullptr, &URDevicesSize);
+
+  std::vector<ur_device_handle_t> URDevices(URDevicesSize /
+                                            sizeof(ur_device_handle_t));
+  Adapter->call<UrApiKind::urProgramGetInfo>(Program, UR_PROGRAM_INFO_DEVICES,
+                                             URDevicesSize, URDevices.data(),
+                                             nullptr);
+  assert(Device2SpillMap.size() == URDevices.size());
+
+  // Map the result back to the program devices. UR provides the following
+  // guarantee:
+  //   The order of the devices is guaranteed (i.e., the same as queried by
+  //   urDeviceGet) by the UR within a single application even if the runtime is
+  //   reinitialized.
+  for (size_t idx = 0; idx < URDevices.size(); ++idx) {
+    if (URDevices[idx] == Device)
+      return size_t{Device2SpillMap[idx]};
+  }
+  throw exception(
+      make_error_code(errc::runtime),
+      "ext::intel::info::kernel::spill_memory_size failed to retrieve "
+      "the requested value");
+}
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/source/detail/ur_info_code.hpp

@@ -72,6 +72,15 @@ template <typename T> struct UrInfoCode;
 #include <sycl/info/ext_oneapi_device_traits.def>
 #undef __SYCL_PARAM_TRAITS_SPEC
 
+#define __SYCL_PARAM_TRAITS_SPEC(Namespace, DescType, Desc, ReturnT, UrCode)   \
+  template <> struct UrInfoCode<Namespace::info::DescType::Desc> {             \
+    static constexpr ur_kernel_info_t value =                                  \
+        static_cast<ur_kernel_info_t>(UrCode);                                 \
+  };
+
+#include <sycl/info/ext_intel_kernel_info_traits.def>
+#undef __SYCL_PARAM_TRAITS_SPEC
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl