[clang] [SYCL] The sycl_kernel_entry_point attribute. (PR #111389)

Ronan Keryell via cfe-commits cfe-commits at lists.llvm.org
Wed Oct 30 18:52:58 PDT 2024


================
@@ -455,6 +455,174 @@ The SYCL kernel in the previous code sample meets these expectations.
   }];
 }
 
+def SYCLKernelEntryPointDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+The ``sycl_kernel_entry_point`` attribute facilitates the generation of an
+offload kernel entry point, sometimes called a SYCL kernel caller function,
+suitable for invoking a SYCL kernel on an offload device. The attribute is
+intended for use in the implementation of SYCL kernel invocation functions
+like the ``single_task`` and ``parallel_for`` member functions of the
+``sycl::handler`` class specified in section 4.9.4, "Command group ``handler``
+class", of the SYCL 2020 specification.
+
+The attribute requires a single type argument that specifies a class type that
+meets the requirements for a SYCL kernel name as described in section 5.2,
+"Naming of kernels", of the SYCL 2020 specification. A unique kernel name type
+is required for each function declared with the attribute. The attribute may
+not first appear on a declaration that follows a definition of the function.
+
+The attribute only appertains to functions and only those that meet the
+following requirements.
+
+* Has a ``void`` return type.
+* Is not a non-static member function, constructor, or destructor.
+* Is not a C variadic function.
+* Is not a coroutine.
+* Is not defined as deleted or as defaulted.
+* Is not declared with the ``constexpr`` or ``consteval`` specifiers.
----------------
keryell wrote:

Just rethinking about this, since I have worked today on `constexpr` in a related context https://github.com/KhronosGroup/SYCL-CTS/pull/976
This looks like a pessimistic interpretation of the specification or something we could clarify in the SYCL committee.
It is not really important to consider `consteval` as it will always be in a constant evaluated context, so it will never reach SYCL codegen (perhaps in the future if we want to speed up the slooow constant evaluator of Clang itself on some accelerators? :rofl: ).
And why not having a kernel which is `constexpr` function too? It will only matter when in a non-constant evaluated context.


https://github.com/llvm/llvm-project/pull/111389


More information about the cfe-commits mailing list