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

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