[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:

Thanks, Erich. I'd still like to understand what your mental model is for what `constexpr` would mean on one of these functions; just to ensure there isn't some misunderstanding.

I've been assuming that the model you had in mind would effectively implicitly inline the `if consteval` statement from my `single_task()` example into `skep()` such that it would behave when called as-if written as follows:
```
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k) {
  if consteval {
    k();
  } else {
    // Do nothing; an offload entry point function will have been generated
    // that the SYCL run-time library will arrange to call.
  }
}
```

This is the reason for my comments about inconsistent behavior in and out of constant evaluation context. Note that any caller would still have to be written to be sensitive to whether invocation is occurring during constant evaluation since it would have to arrange (or not) for the offload entry point to be invoked on the appropriate device.

If you have a different model in mind, please do share details.

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