================
@@ -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:
Ok, thanks. Note that this design doesn't require the `#ifdef
__SYCL_DEVICE_ONLY__` dance that was needed with the `sycl_kernel` attribute.
This is an intentional divergence made to accommodate a future single pass
compilation model. So, the body of `skep()` does exist during both host and
device compilation, it just isn't evaluated when called. These are details for
a later PR.
> However, the 'faster/more parallel' is a bit of a QoI thing, so 'call this
> slower, but at compile time' should work as well.
I agree, but that needs to be handled at a higher level.
`sycl_kernel_entry_point` attributed functions can't fill the role of a user
facing interface since they don't include the mechanism needed for device
dispatch. These functions can only be used as part of the implementation detail
for higher level interfaces.
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
