@Fznamznon, I added this just for temporary testing purposes so that I could validate that code generation was working as expected. Some of this might be useful to you; the creation of the implicit array-to-pointer-decay cast for example. Use whatever you find useful and remove the rest.

@Fznamznon, this is the major work that is left to do. Note that name lookup and overload resolution should be performed from the point of definition of the sycl_kernel_entry_point attributed function or, if that function is an instantiated function template specialization, the point of instantiation. I'm hopeful that SemaRef.CurContext reflects this point; see the associated assert in SemaSYCL::BuildSYCLKernelCallStmt() below.

We should consider whether to support implicit object parameter declarations or, at a minimum, add a test to ensure a proper diagnostic is issued for an example like this:

struct sycl_kernel_launcher {
  template<typename KN, typename... Ts>
  void sycl_enqueue_kernel_launch(this S* self, const char *name, Ts...);
  template<typename KN, typename KT>
  [[clang::sycl_kernel_entry_point(KN)]]
  void kernel_entry_point(KT kernel) { ... }
};

There is a design decision we can consider here too. Should we allow the sycl_enqueue_kernel_launch name to be a default name that can be customized by an additional argument to the sycl_kernel_entry_point attribute? Let me know if you have opinions on it. I'm otherwise content to leave such considerations to future work.

If we are going to go with sycl_enqueue_kernel_launch as a default name, it should either be reserved by SYCL spec or we need to add __ to it, otherwise it may/will clash with user's code. Same goes with class's name sycl_kernel_launcher. I don't have a strong preference on whether it should be possible to customize sycl_enqueue_kernel_launch name, I think people from Runtime library side may have a better understanding.

A reserved name isn't needed since the SYCL RT will provide the sycl_enqueue_kernel_launch declaration. We aren't implicitly declaring any entities with this approach; we're just performing lookup for entities declared by the SYCL RT.

The sycl_kernel_launcher name isn't special in any way; I just used that name as an example. In this example, lookup for sycl_enqueue_kernel_launch would be performed from the definition of sycl_kernel_launcher::kernel_entry_point() and sycl_kernel_launcher::sycl_enqueue_kernel_launch() is in scope for name lookup at that point.

We aren't implicitly declaring any entities with this approach; we're just performing lookup for entities declared by the SYCL RT.

I understand that. Still it may be unexpected for the users that SYCL RT declares that name.

Agreed. We'll need to ensure the diagnostic issued when overload resolution fails is clear. I hope we can do better than what Clang does for CUDA when the __cudaPushCallConfiguration() declaration is missing:

t.cu:6:9: error: use of undeclared identifier __cudaPushCallConfiguration
   6 |   kernel<<<1,1>>>([] __attribute__((device)) {});
     |         ^

@tahonermann , shouldn't that function and others like it accept kernel object parameter?

Per offline discussion, yes it should be kind of maybe. I think what is happening here is that the kernel object is passed by value and, since the kernel object type is an empty class, the target ABI specifies that no arguments actually be passed; the (stateless) parameter object is constructed in the body of the function.

We could add state to the kernel object types to force an argument to be provided. I don't think it matters much though; all we really need to validate here is that the synthesized call to the sycl_enqueue_kernel_launch function is present and looks correct.

Another thing that we should validate is exception handling. I don't know if any actual code changes will be required. Scenarios to be considered:

• The sycl_kernel_entry_point function is declared noexcept; an exception propagating from the call to the sycl_enqueue_kernel_launch() function (including from construction of arguments) should result in a call to std::terminate().
• The sycl_kernel_entry_point function is not declared noexcept; if the called sycl_enqueue_kernel_launch() function is also not declared noexcept, then exception handling scaffolding might be required (e.g., to destruct value parameters that have non-trivial destruction) and a thrown exception should propagate.
• The sycl_kernel_entry_point function is not declared noexcept; if the called sycl_enqueue_kernel_launch() function is declared noexcept, then exception handling scaffolding can be omitted.