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

[Erich Keane 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.
----------------
erichkeane wrote:

That is effectively  my mental model, except without the `if consteval`, since the body of the `skep` doesn't really exist when doing the offload codegen. The point is there is a lot of `device` stuff that gets arranged that is, in my mental model, implementation details for "call this function object, just faster/more parallel". 

However, the 'faster/more parallel' is a bit of a QoI thing, so 'call this slower, but at compile time' should work as well.

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