Skip to content

Commit a03dc0d

Browse files
gmlueckgmlueck
authored
[SYCL][Doc] Propose "get_kernel_info" extension (#14472)
Add a proposed extension to query a kernel's information descriptor without creating a kernel bundle.
1 parent 2df7d3f commit a03dc0d

File tree

2 files changed

+340
-1
lines changed

2 files changed

+340
-1
lines changed

‎sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc

Lines changed: 111 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -329,7 +329,6 @@ Otherwise `value` is `false`.
329329
The helper trait `is_kernel_v` provides the value of `value`.
330330
|====
331331

332-
333332
=== New kernel bundle member functions
334333

335334
This extension adds the following new functions which add kernel bundle support
@@ -494,6 +493,117 @@ _Throws_: An `exception` with the error code `errc::invalid` if the kernel with
494493
address `Func` does not reside in this kernel bundle.
495494
|====
496495

496+
=== New free functions to query kernel information descriptors
497+
498+
This extension adds the following new free functions, which allow an application
499+
to query the kernel information descriptors for a free function kernel without
500+
first creating a kernel bundle.
501+
502+
[frame=all,grid=none,separator="@"]
503+
!====
504+
a@
505+
[source,c++]
506+
----
507+
namespace sycl::ext::oneapi::experimental {
508+
509+
template<auto *Func, typename Param>
510+
typename Param::return_type get_kernel_info(const context& ctxt);
511+
512+
} // namespace sycl::ext::oneapi::experimental
513+
----
514+
!====
515+
516+
_Constraints_: Available only if `is_kernel_v<Func>` is `true`.
517+
Available only if `Param` is an information descriptor for the `kernel` class,
518+
which can be used by the `kernel::get_info()` overload.
519+
520+
_Returns:_ The same value `ret` that would be computed by:
521+
522+
[source,c++]
523+
----
524+
auto bundle =
525+
sycl::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
526+
auto ret = bundle.ext_oneapi_get_kernel<Func>().get_info<Param>();
527+
----
528+
529+
_Remarks:_ Each information descriptor may specify additional preconditions,
530+
exceptions that are thrown, etc.
531+
532+
'''
533+
534+
[frame=all,grid=none,separator="@"]
535+
!====
536+
a@
537+
[source,c++]
538+
----
539+
namespace sycl::ext::oneapi::experimental {
540+
541+
template<auto *Func, typename Param>
542+
typename Param::return_type get_kernel_info(const context& ctxt,
543+
const device& dev);
544+
545+
} // namespace sycl::ext::oneapi::experimental
546+
----
547+
!====
548+
549+
_Constraints_: Available only if `is_kernel_v<Func>` is `true`.
550+
Available only if `Param` is an information descriptor for the `kernel` class,
551+
which can be used by the `kernel::get_info(const device&)` overload.
552+
553+
_Preconditions:_ The device `dev` must be one of the devices contained by `ctxt`
554+
or must be a descendent device of some device in `ctxt`.
555+
The kernel `Func` must be compatible with the device `dev` as defined by
556+
`is_compatible`.
557+
558+
_Returns:_ The same value `ret` that would be computed by:
559+
560+
[source,c++]
561+
----
562+
auto bundle =
563+
sycl::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
564+
auto ret = bundle.ext_oneapi_get_kernel<Func>().get_info<Param>(dev);
565+
----
566+
567+
_Remarks:_ Each information descriptor may specify additional preconditions,
568+
exceptions that are thrown, etc.
569+
570+
'''
571+
572+
[frame=all,grid=none,separator="@"]
573+
!====
574+
a@
575+
[source,c++]
576+
----
577+
namespace sycl::ext::oneapi::experimental {
578+
579+
template<typename Func, typename Param>
580+
typename Param::return_type get_kernel_info(const queue& q);
581+
582+
} // namespace sycl::ext::oneapi::experimental
583+
----
584+
!====
585+
586+
_Constraints_: Available only if `is_kernel_v<Func>` is `true`.
587+
Available only if `Param` is an information descriptor for the `kernel` class,
588+
which can be used by the `kernel::get_info(const device&)` overload.
589+
590+
_Preconditions:_ The kernel `Func` must be compatible with the device associated
591+
with `q` as defined by `is_compatible`.
592+
593+
_Returns:_ The same value `ret` that would be computed by:
594+
595+
[source,c++]
596+
----
597+
sycl::context ctxt = q.get_context();
598+
sycl::device dev = q.get_device();
599+
auto bundle =
600+
sycl::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
601+
auto ret = bundle.ext_oneapi_get_kernel<Func>().get_info<Param>(dev);
602+
----
603+
604+
_Remarks:_ Each information descriptor may specify additional preconditions,
605+
exceptions that are thrown, etc.
606+
497607
=== Behavior with kernel bundle functions in the core SYCL specification
498608

499609
Free function kernels that are defined by the application have a corresponding

‎sycl/doc/extensions/proposed/sycl_ext_oneapi_get_kernel_info.asciidoc

Lines changed: 229 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,229 @@
1+
= sycl_ext_oneapi_get_kernel_info
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+
14+
// Set the default source code type in this document to C++,
15+
// for syntax highlighting purposes. This is needed because
16+
// docbook uses c++ and html5 uses cpp.
17+
:language: {basebackend@docbook:c++:cpp}
18+
19+
20+
== Notice
21+
22+
[%hardbreaks]
23+
Copyright (C) 2024 Intel Corporation. All rights reserved.
24+
25+
Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
26+
of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by
27+
permission by Khronos.
28+
29+
30+
== Contact
31+
32+
To report problems with this extension, please open a new issue at:
33+
34+
https://github.com/intel/llvm/issues
35+
36+
37+
== Dependencies
38+
39+
This extension is written against the SYCL 2020 revision 8 specification. All
40+
references below to the "core SYCL specification" or to section numbers in the
41+
SYCL specification refer to that revision.
42+
43+
44+
== Status
45+
46+
This is a proposed extension specification, intended to gather community
47+
feedback. Interfaces defined in this specification may not be implemented yet
48+
or may be in a preliminary state. The specification itself may also change in
49+
incompatible ways before it is finalized. *Shipping software products should
50+
not rely on APIs defined in this specification.*
51+
52+
53+
== Overview
54+
55+
Applications sometimes need to query a kernel's information descriptor in order
56+
to decide how to launch the kernel.
57+
For example, an application may need to query
58+
`info::kernel_device_specific::work_group_size` in order to determine the
59+
nd-range to use when launching the kernel.
60+
61+
Currently, the only way to do this is to create a kernel bundle, get the
62+
`kernel` object from that bundle, and then query the `kernel` object.
63+
This is very verbose, especially when the application doesn't need any of the
64+
other facilities provided by the kernel bundle API.
65+
66+
This extension provides a less verbose way to query a kernel's information
67+
descriptor without creating a kernel bundle.
68+
69+
70+
== Specification
71+
72+
=== Feature test macro
73+
74+
This extension provides a feature-test macro as described in the core SYCL
75+
specification. An implementation supporting this extension must predefine the
76+
macro `SYCL_EXT_ONEAPI_GET_KERNEL_INFO` to one of the values defined in the
77+
table below.
78+
Applications can test for the existence of this macro to determine if the
79+
implementation supports this feature, or applications can test the macro's value
80+
to determine which of the extension's features the implementation supports.
81+
82+
[%header,cols="1,5"]
83+
|===
84+
|Value
85+
|Description
86+
87+
|1
88+
|Initial version of this extension.
89+
|===
90+
91+
=== New free functions
92+
93+
This extension adds the following new free functions for querying a kernel's
94+
information descriptors.
95+
96+
'''
97+
98+
[frame=all,grid=none,separator="@"]
99+
!====
100+
a@
101+
[source,c++]
102+
----
103+
namespace sycl::ext::oneapi {
104+
105+
template<typename KernelName, typename Param>
106+
typename Param::return_type get_kernel_info(const context& ctxt);
107+
108+
} // namespace sycl::ext::oneapi
109+
----
110+
!====
111+
112+
_Constraints:_ Available only if `Param` is an information descriptor for the
113+
`kernel` class, which can be used by the `kernel::get_info()` overload.
114+
115+
_Preconditions:_ The `KernelName` must be the type kernel name of a kernel that
116+
is defined in the application.
117+
118+
_Returns:_ The same value `ret` that would be computed by:
119+
120+
[source,c++]
121+
----
122+
auto bundle =
123+
sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(ctxt);
124+
auto ret = bundle.get_kernel<KernelName>().get_info<Param>();
125+
----
126+
127+
_Remarks:_ Each information descriptor may specify additional preconditions,
128+
exceptions that are thrown, etc.
129+
130+
'''
131+
132+
[frame=all,grid=none,separator="@"]
133+
!====
134+
a@
135+
[source,c++]
136+
----
137+
namespace sycl::ext::oneapi {
138+
139+
template<typename KernelName, typename Param>
140+
typename Param::return_type get_kernel_info(const context& ctxt,
141+
const device& dev);
142+
143+
} // namespace sycl::ext::oneapi
144+
----
145+
!====
146+
147+
_Constraints:_ Available only if `Param` is an information descriptor for the
148+
`kernel` class, which can be used by the `kernel::get_info(const device &)`
149+
overload.
150+
151+
_Preconditions:_ The `KernelName` must be the type kernel name of a kernel that
152+
is defined in the application.
153+
The device `dev` must be one of the devices contained by `ctxt` or must be a
154+
descendent device of some device in `ctxt`.
155+
The kernel `KernelName` must be compatible with the device `dev` as defined by
156+
`is_compatible`.
157+
158+
_Returns:_ The same value `ret` that would be computed by:
159+
160+
[source,c++]
161+
----
162+
auto bundle =
163+
sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(ctxt);
164+
auto ret = bundle.get_kernel<KernelName>().get_info<Param>(dev);
165+
----
166+
167+
_Remarks:_ Each information descriptor may specify additional preconditions,
168+
exceptions that are thrown, etc.
169+
170+
'''
171+
172+
[frame=all,grid=none,separator="@"]
173+
!====
174+
a@
175+
[source,c++]
176+
----
177+
namespace sycl::ext::oneapi {
178+
179+
template<typename KernelName, typename Param>
180+
typename Param::return_type get_kernel_info(const queue& q);
181+
182+
} // namespace sycl::ext::oneapi
183+
----
184+
!====
185+
186+
_Constraints:_ Available only if `Param` is an information descriptor for the
187+
`kernel` class, which can be used by the `kernel::get_info(const device &)`
188+
overload.
189+
190+
_Preconditions:_ The `KernelName` must be the type kernel name of a kernel that
191+
is defined in the application.
192+
The kernel `KernelName` must be compatible with the device associated with `q`
193+
as defined by `is_compatible`.
194+
195+
_Returns:_ The same value `ret` that would be computed by:
196+
197+
[source,c++]
198+
----
199+
sycl::context ctxt = q.get_context();
200+
sycl::device dev = q.get_device();
201+
auto bundle =
202+
sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(ctxt);
203+
auto ret = bundle.get_kernel<KernelName>().get_info<Param>(dev);
204+
----
205+
206+
_Remarks:_ Each information descriptor may specify additional preconditions,
207+
exceptions that are thrown, etc.
208+
209+
210+
== Issues
211+
212+
* I purposely reduced the exceptions that are required to be thrown for certain
213+
error conditions and instead listed these as preconditions.
214+
An implementation can still diagnose these error conditions by throwing an
215+
exception, but it is not required.
216+
Since these APIs are likely on the critical path for launching a kernel, I
217+
don't think we want to mandate an error check at runtime.
218+
In retrospect, I think this is the right behavior for the core SYCL spec also,
219+
and we should consider changing the specified behavior.
220+
Thoughts?
221+
222+
* I'm not sure how to formally specify the requirements for `KernelName`.
223+
I think an implementation should be able to fail with a link-time error if
224+
`KernelName` is not the type-name of some kernel that is defined in the
225+
application.
226+
However, this seems different from a _Constraint_, which is expected to result
227+
in a compile-time error.
228+
For now, I just listed it as a _Precondition_, so there is no formal
229+
requirement for an implementation to diagnose this error.

0 commit comments

Comments
 (0)