================
@@ -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.
----------------
tahonermann wrote:
With regard to `constexpr`, consider the following example. Is the lambda
actually invoked when `f()` is called?
```
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k) { k(); }
void f() {
skep<struct K>([]{});
}
```
The answer is no. As described in the documentation and quoted in an earlier
comment in this conversation, the call to the `sycl_kernel_entry_point`
attributed function happens, but it's apparent body is not actually executed.
The body serves as a pattern for the offload kernel entry point function.
With that in mind, what should the following do?
```
constexpr auto v = (skep<struct K>([]{}), 0);
```
The answer is that it should do nothing (invoking the lambda would be
inconsistent with what happens outside of constant evaluation).
The way to support constant evaluation of a SYCL kernel (which may be
`constexpr`) would be to write the SYCL invocation function to detect constant
evaluation and react accordingly.
```
template<typename KN, typename KT>
constexpr void single_task(KT k) {
if consteval {
k();
} else {
skep<KN>(k); // Trigger emission of the offload kernel entry point.
// FIXME: Do stuff to invoke the offload kernel entry point
}
}
constexpr auto v = (single_task<struct K>([]{}), 0); // Huzzah, constant
evaluated SYCL kernel.
```
https://github.com/llvm/llvm-project/pull/111389
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
