Skip to content

Commit 65498a6

Browse files
kurapov-petergmlueckaelovikov-intel
authored
[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]>
1 parent 48bb501 commit 65498a6

File tree

18 files changed

+283
-5
lines changed

18 files changed

+283
-5
lines changed

‎llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -86,6 +86,7 @@ def AspectExt_oneapi_cuda_cluster_group : Aspect<"ext_oneapi_cuda_cluster_group"
8686
def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
8787
def AspectExt_oneapi_atomic16 : Aspect<"ext_oneapi_atomic16">;
8888
def AspectExt_oneapi_virtual_functions : Aspect<"ext_oneapi_virtual_functions">;
89+
def AspectExt_intel_spill_memory_size : Aspect<"ext_intel_spill_memory_size">;
8990
// Deprecated aspects
9091
def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
9192
def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -150,7 +151,8 @@ def : TargetInfo<"__TestAspectList",
150151
AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group,
151152
AspectExt_intel_fpga_task_sequence,
152153
AspectExt_oneapi_atomic16,
153-
AspectExt_oneapi_virtual_functions],
154+
AspectExt_oneapi_virtual_functions,
155+
AspectExt_intel_spill_memory_size],
154156
[]>;
155157
// This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
156158
// match.

‎sycl/doc/extensions/supported/sycl_ext_intel_kernel_queries.asciidoc

Lines changed: 139 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,139 @@
1+
= sycl_ext_intel_kernel_queries
2+
3+
:source-highlighter: coderay
4+
:coderay-linenums-mode: table
5+
6+
// This section needs to be after the document title.
7+
:doctype: book
8+
:toc2:
9+
:toc: left
10+
:encoding: utf-8
11+
:lang: en
12+
:dpcpp: pass:[DPC++]
13+
:endnote: &#8212;{nbsp}end{nbsp}note
14+
15+
// Set the default source code type in this document to C++,
16+
// for syntax highlighting purposes. This is needed because
17+
// docbook uses c++ and html5 uses cpp.
18+
:language: {basebackend@docbook:c++:cpp}
19+
20+
21+
== Notice
22+
23+
[%hardbreaks]
24+
Copyright (C) 2025 Intel Corporation. All rights reserved.
25+
26+
Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
27+
of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by
28+
permission by Khronos.
29+
30+
31+
== Contact
32+
33+
To report problems with this extension, please open a new issue at:
34+
35+
https://github.com/intel/llvm/issues
36+
37+
38+
== Dependencies
39+
40+
This extension is written against the SYCL 2020 revision 9 specification.
41+
All references below to the "core SYCL specification" or to section numbers in
42+
the SYCL specification refer to that revision.
43+
44+
45+
== Status
46+
47+
This extension is implemented and fully supported by {dpcpp}.
48+
49+
50+
== Overview
51+
52+
This extension contains a collection of queries that provide low-level
53+
information about kernels.
54+
These queries generally forward directly to the backend and expose concepts that
55+
are specific to a particular implementation.
56+
As a result, these queries may not be supported for all devices.
57+
Each query has an associate device aspect, which tells whether the query is
58+
supported on that device.
59+
60+
61+
== Specification
62+
63+
=== Feature test macro
64+
65+
This extension provides a feature-test macro as described in the core SYCL
66+
specification.
67+
An implementation supporting this extension must predefine the macro
68+
`SYCL_EXT_INTEL_KERNEL_QUERIES` to one of the values defined in the table below.
69+
Applications can test for the existence of this macro to determine if the
70+
implementation supports this feature, or applications can test the macro's value
71+
to determine which of the extension's features the implementation supports.
72+
73+
[%header,cols="1,5"]
74+
|===
75+
|Value
76+
|Description
77+
78+
|1
79+
|Initial version of this extension.
80+
|===
81+
82+
=== Spill memory size
83+
84+
This query returns the kernel's spill memory size that is allocated by the
85+
compiler, as reported by Level Zero.
86+
87+
==== New device aspect
88+
89+
This extension adds the following new device aspect.
90+
91+
[source,c++]
92+
----
93+
namespace sycl {
94+
95+
enum class aspect {
96+
ext_intel_spill_memory_size
97+
98+
// ...
99+
};
100+
101+
}
102+
----
103+
104+
'''
105+
106+
`*ext_intel_spill_memory_size*`
107+
108+
Indicates that the `spill_memory_size` kernel information descriptor may be used
109+
to query kernels for this device.
110+
111+
'''
112+
113+
==== New device specific kernel information descriptor
114+
115+
This extension adds the following information descriptor that can be used with
116+
`kernel::get_info(const device&)`.
117+
118+
'''
119+
120+
`*ext::intel::info::kernel_device_specific::spill_memory_size*`
121+
122+
[source,c++]
123+
----
124+
namespace sycl::ext::intel::info::kernel_device_specific {
125+
struct spill_memory_size {
126+
using return_type = size_t;
127+
};
128+
} // namespace sycl::ext::intel::info::kernel_device_specific
129+
----
130+
131+
_Remarks:_ Template parameter to `kernel::get_info(const device&)`.
132+
133+
_Returns:_ The spill memory size that is allocated by the compiler for this
134+
kernel for the given device.
135+
136+
_Throws:_ An `exception` with the `errc::feature_not_supported` error code if
137+
the device does not have `aspect::ext_intel_spill_memory_size`.
138+
139+
'''

‎sycl/include/sycl/detail/info_desc_helpers.hpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -88,6 +88,15 @@ struct IsKernelInfo<info::kernel_device_specific::ext_codeplay_num_regs>
8888
#include <sycl/info/kernel_device_specific_traits.def>
8989
#undef __SYCL_PARAM_TRAITS_SPEC
9090

91+
#define __SYCL_PARAM_TRAITS_SPEC(Namespace, DescType, Desc, ReturnT, UrCode) \
92+
template <> \
93+
struct is_##DescType##_info_desc<Namespace::info::DescType::Desc> \
94+
: std::true_type { \
95+
using return_type = Namespace::info::DescType::Desc::return_type; \
96+
};
97+
#include <sycl/info/ext_intel_kernel_info_traits.def>
98+
#undef __SYCL_PARAM_TRAITS_SPEC
99+
91100
#define __SYCL_PARAM_TRAITS_SPEC(DescType, Desc, ReturnT, UrCode) \
92101
template <> \
93102
struct is_##DescType##_info_desc<info::DescType::Desc> : std::true_type { \

‎sycl/include/sycl/info/aspects.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,3 +72,4 @@ __SYCL_ASPECT(ext_oneapi_bindless_images_sample_1d_usm, 78)
7272
__SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 79)
7373
__SYCL_ASPECT(ext_oneapi_atomic16, 80)
7474
__SYCL_ASPECT(ext_oneapi_virtual_functions, 81)
75+
__SYCL_ASPECT(ext_intel_spill_memory_size, 82)

‎sycl/include/sycl/info/ext_intel_kernel_info_traits.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -261,6 +261,7 @@ struct work_item_progress_capabilities;
261261
} // namespace ext::oneapi::experimental::info::device
262262
#include <sycl/info/ext_codeplay_device_traits.def>
263263
#include <sycl/info/ext_intel_device_traits.def>
264+
#include <sycl/info/ext_intel_kernel_info_traits.def>
264265
#include <sycl/info/ext_oneapi_device_traits.def>
265266
#include <sycl/info/ext_oneapi_kernel_queue_specific_traits.def>
266267

‎sycl/source/detail/device_impl.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -786,6 +786,11 @@ bool device_impl::has(aspect Aspect) const {
786786
BE == sycl::backend::opencl;
787787
return (is_cpu() || is_gpu()) && isCompatibleBE;
788788
}
789+
case aspect::ext_intel_spill_memory_size: {
790+
backend BE = getBackend();
791+
bool isCompatibleBE = BE == sycl::backend::ext_oneapi_level_zero;
792+
return is_gpu() && isCompatibleBE;
793+
}
789794
}
790795

791796
return false; // This device aspect has not been implemented yet.

‎sycl/source/detail/kernel_impl.hpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -416,6 +416,23 @@ inline typename syclex::info::kernel_queue_specific::max_num_work_groups::
416416
DynamicLocalMemorySize);
417417
}
418418

419+
template <>
420+
inline typename ext::intel::info::kernel_device_specific::spill_memory_size::
421+
return_type
422+
kernel_impl::get_info<
423+
ext::intel::info::kernel_device_specific::spill_memory_size>(
424+
const device &Device) const {
425+
if (!Device.has(aspect::ext_intel_spill_memory_size))
426+
throw exception(
427+
make_error_code(errc::feature_not_supported),
428+
"This device does not have the ext_intel_spill_memory_size aspect");
429+
430+
return get_kernel_device_specific_info<
431+
ext::intel::info::kernel_device_specific::spill_memory_size>(
432+
this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(),
433+
getAdapter());
434+
}
435+
419436
template <>
420437
inline typename syclex::info::kernel_queue_specific::max_work_group_size::
421438
return_type

‎sycl/source/detail/kernel_info.hpp

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -146,6 +146,58 @@ uint32_t get_kernel_device_specific_info_with_input(ur_kernel_handle_t Kernel,
146146

147147
return Result;
148148
}
149+
150+
template <>
151+
inline ext::intel::info::kernel_device_specific::spill_memory_size::return_type
152+
get_kernel_device_specific_info<
153+
ext::intel::info::kernel_device_specific::spill_memory_size>(
154+
ur_kernel_handle_t Kernel, ur_device_handle_t Device,
155+
const AdapterPtr &Adapter) {
156+
size_t ResultSize = 0;
157+
158+
// First call to get the number of device images
159+
Adapter->call<UrApiKind::urKernelGetInfo>(
160+
Kernel, UR_KERNEL_INFO_SPILL_MEM_SIZE, 0, nullptr, &ResultSize);
161+
162+
size_t DeviceCount = ResultSize / sizeof(uint32_t);
163+
164+
// Second call to retrieve the data
165+
std::vector<uint32_t> Device2SpillMap(DeviceCount);
166+
Adapter->call<UrApiKind::urKernelGetInfo>(
167+
Kernel, UR_KERNEL_INFO_SPILL_MEM_SIZE, ResultSize, Device2SpillMap.data(),
168+
nullptr);
169+
170+
ur_program_handle_t Program;
171+
Adapter->call<UrApiKind::urKernelGetInfo>(Kernel, UR_KERNEL_INFO_PROGRAM,
172+
sizeof(ur_program_handle_t),
173+
&Program, nullptr);
174+
// Retrieve the associated device list
175+
size_t URDevicesSize = 0;
176+
Adapter->call<UrApiKind::urProgramGetInfo>(Program, UR_PROGRAM_INFO_DEVICES,
177+
0, nullptr, &URDevicesSize);
178+
179+
std::vector<ur_device_handle_t> URDevices(URDevicesSize /
180+
sizeof(ur_device_handle_t));
181+
Adapter->call<UrApiKind::urProgramGetInfo>(Program, UR_PROGRAM_INFO_DEVICES,
182+
URDevicesSize, URDevices.data(),
183+
nullptr);
184+
assert(Device2SpillMap.size() == URDevices.size());
185+
186+
// Map the result back to the program devices. UR provides the following
187+
// guarantee:
188+
// The order of the devices is guaranteed (i.e., the same as queried by
189+
// urDeviceGet) by the UR within a single application even if the runtime is
190+
// reinitialized.
191+
for (size_t idx = 0; idx < URDevices.size(); ++idx) {
192+
if (URDevices[idx] == Device)
193+
return size_t{Device2SpillMap[idx]};
194+
}
195+
throw exception(
196+
make_error_code(errc::runtime),
197+
"ext::intel::info::kernel::spill_memory_size failed to retrieve "
198+
"the requested value");
199+
}
200+
149201
} // namespace detail
150202
} // namespace _V1
151203
} // namespace sycl

‎sycl/source/detail/ur_info_code.hpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,15 @@ template <typename T> struct UrInfoCode;
7272
#include <sycl/info/ext_oneapi_device_traits.def>
7373
#undef __SYCL_PARAM_TRAITS_SPEC
7474

75+
#define __SYCL_PARAM_TRAITS_SPEC(Namespace, DescType, Desc, ReturnT, UrCode) \
76+
template <> struct UrInfoCode<Namespace::info::DescType::Desc> { \
77+
static constexpr ur_kernel_info_t value = \
78+
static_cast<ur_kernel_info_t>(UrCode); \
79+
};
80+
81+
#include <sycl/info/ext_intel_kernel_info_traits.def>
82+
#undef __SYCL_PARAM_TRAITS_SPEC
83+
7584
} // namespace detail
7685
} // namespace _V1
7786
} // namespace sycl

‎sycl/source/kernel.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,6 +101,14 @@ kernel::get_info(const device &Device, const range<3> &WGSize) const {
101101

102102
#undef __SYCL_PARAM_TRAITS_SPEC
103103

104+
#define __SYCL_PARAM_TRAITS_SPEC(Namespace, DescType, Desc, ReturnT, UrCode) \
105+
template __SYCL_EXPORT ReturnT \
106+
kernel::get_info<Namespace::info::DescType::Desc>(const device &) const;
107+
108+
#include <sycl/info/ext_intel_kernel_info_traits.def>
109+
110+
#undef __SYCL_PARAM_TRAITS_SPEC
111+
104112
template __SYCL_EXPORT uint32_t
105113
kernel::get_info<info::kernel_device_specific::max_sub_group_size>(
106114
const device &, const sycl::range<3> &) const;

‎sycl/test-e2e/Basic/aspects.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -90,6 +90,9 @@ int main() {
9090
if (plt.has(aspect::ext_oneapi_virtual_functions)) {
9191
std::cout << " ext_oneapi_virtual_functions" << std::endl;
9292
}
93+
if (plt.has(aspect::ext_intel_spill_memory_size)) {
94+
std::cout << " ext_intel_spill_memory_size" << std::endl;
95+
}
9396
}
9497
std::cout << "Passed." << std::endl;
9598
return 0;

‎sycl/test-e2e/Basic/kernel_info.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,6 +101,12 @@ int main() {
101101
krn.get_info<info::kernel_device_specific::compile_num_sub_groups>(dev);
102102
assert(compileNumSg <= maxNumSg);
103103

104+
size_t spillMemSz = 0;
105+
if (dev.has(aspect::ext_intel_spill_memory_size)) {
106+
spillMemSz = krn.get_info<
107+
ext::intel::info::kernel_device_specific::spill_memory_size>(dev);
108+
}
109+
104110
// Use ext_oneapi_get_kernel_info extension and check that answers match.
105111
const size_t wgSizeExt = syclex::get_kernel_info<
106112
SingleTask, info::kernel_device_specific::work_group_size>(ctx, dev);
@@ -125,6 +131,13 @@ int main() {
125131
dev);
126132
assert(compileNumSgExt == compileNumSg);
127133

134+
if (dev.has(aspect::ext_intel_spill_memory_size)) {
135+
const size_t spillMemSizeExt = syclex::get_kernel_info<
136+
SingleTask,
137+
ext::intel::info::kernel_device_specific::spill_memory_size>(ctx, dev);
138+
assert(spillMemSizeExt == spillMemSz);
139+
}
140+
128141
// Use ext_oneapi_get_kernel_info extension with queue parameter and check the
129142
// result.
130143
const size_t wgSizeExtQ =

‎sycl/test/abi/sycl_symbols_linux.dump

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3976,6 +3976,7 @@ _ZNK4sycl3_V16kernel19ext_oneapi_get_infoINS0_3ext6oneapi12experimental4info21ke
39763976
_ZNK4sycl3_V16kernel19ext_oneapi_get_infoINS0_3ext6oneapi12experimental4info21kernel_queue_specific14num_sub_groupsEEENS0_6detail34is_kernel_queue_specific_info_descIT_E11return_typeENS0_5queueERKNS0_5rangeILi2EEE
39773977
_ZNK4sycl3_V16kernel19ext_oneapi_get_infoINS0_3ext6oneapi12experimental4info21kernel_queue_specific18max_sub_group_sizeEEENS0_6detail34is_kernel_queue_specific_info_descIT_E11return_typeENS0_5queueERKNS0_5rangeILi3EEE
39783978
_ZNK4sycl3_V16kernel3getEv
3979+
_ZNK4sycl3_V16kernel8get_infoINS0_3ext5intel4info22kernel_device_specific17spill_memory_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
39793980
_ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific15work_group_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
39803981
_ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific16global_work_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE
39813982
_ZNK4sycl3_V16kernel8get_infoINS0_4info22kernel_device_specific16private_mem_sizeEEENS0_6detail35is_kernel_device_specific_info_descIT_E11return_typeERKNS0_6deviceE

‎sycl/test/abi/sycl_symbols_windows.dump

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -101,6 +101,7 @@
101101
??$get_info@Ureference_count@context@info@_V1@sycl@@@context@_V1@sycl@@QEBAIXZ
102102
??$get_info@Ureference_count@event@info@_V1@sycl@@@event@_V1@sycl@@QEBAIXZ
103103
??$get_info@Ureference_count@queue@info@_V1@sycl@@@queue@_V1@sycl@@QEBAIXZ
104+
??$get_info@Uspill_memory_size@kernel_device_specific@info@intel@ext@_V1@sycl@@@kernel@_V1@sycl@@QEBA_KAEBVdevice@12@@Z
104105
??$get_info@Usupports_fusion@device@info@experimental@codeplay@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA_NXZ
105106
??$get_info@Uuuid@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$array@E$0BA@@std@@XZ
106107
??$get_info@Uwork_group_size@kernel_device_specific@info@_V1@sycl@@@kernel@_V1@sycl@@QEBA_KAEBVdevice@12@@Z

‎sycl/test/include_deps/sycl_accessor.hpp.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -78,6 +78,7 @@
7878
// CHECK-NEXT: info/event_profiling_traits.def
7979
// CHECK-NEXT: info/ext_codeplay_device_traits.def
8080
// CHECK-NEXT: info/ext_intel_device_traits.def
81+
// CHECK-NEXT: info/ext_intel_kernel_info_traits.def
8182
// CHECK-NEXT: info/ext_oneapi_device_traits.def
8283
// CHECK-NEXT: info/ext_oneapi_kernel_queue_specific_traits.def
8384
// CHECK-NEXT: info/sycl_backend_traits.def

‎sycl/test/include_deps/sycl_detail_core.hpp.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -79,6 +79,7 @@
7979
// CHECK-NEXT: info/event_profiling_traits.def
8080
// CHECK-NEXT: info/ext_codeplay_device_traits.def
8181
// CHECK-NEXT: info/ext_intel_device_traits.def
82+
// CHECK-NEXT: info/ext_intel_kernel_info_traits.def
8283
// CHECK-NEXT: info/ext_oneapi_device_traits.def
8384
// CHECK-NEXT: info/ext_oneapi_kernel_queue_specific_traits.def
8485
// CHECK-NEXT: info/sycl_backend_traits.def

0 commit comments

Comments
 (0)