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

[Tom Honermann via cfe-commits]

================
@@ -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