Definition of "kernel argument" is wrong for named kernel object

[gmlueck]

Specification Version

SYCL 2020 (Revision 10)

Section Number(s)

section 4.12.4. "Rules for parameter passing to kernels"

Bug Description

During a conversation at the Philadelphia F2F, we decided that this statement in that section is not quite right:

If the kernel is a named function object, the operator() member function (or other member functions that it calls) may reference member variables inside the same named function object. Any such member variables become parameters to the kernel.

This should be changed to say that all member variables of the named function object are kernel arguments. If an implementation can prove that certain member variables are not used in the kernel, then it can avoid passing them. This is essentially dead argument elimination.

During our conversation, we noted that the compiler cannot always prove whether a member variable of the kernel object is used in device code. Consider the case when the operator() calls some SYCL_EXTERNAL member function that is defined in another translation unit. In such a case, the compiler cannot know which member variables are used in the SYCL_EXTERNAL function.

[gmlueck]

According to @tahonermann, we think that the DPC++ front-end already considers all member variables of a named kernel object to be kernel arguments when it checks their types. Therefore, the front-end diagnoses an error if a member variable has a type that is not legal for a kernel argument, regardless of whether the member variable is used by operator() (or the functions it calls).

Therefore, changing the spec language as described above should not impact the errors that DPC++ diagnoses.

[tahonermann]

Thanks for filing this, @gmlueck. The example I shared during the meeting is shown below with some slight alterations to add the static_assert declarations and remove sycl::is_device_copyable_v specializations; more on this below. This example demonstrates that DPC++ diagnoses the presence of kernel object type data members that do not satisfy the 3.13.1, "Device copyable" requirements, even when they are not used by the kernel.

https://godbolt.org/z/c9hbEfr37:

#include <type_traits>
#include <sycl/sycl.hpp>

struct move_only {
  move_only(const move_only&) = delete;
  move_only(move_only&&);
  int dm;
};
static_assert(!std::is_trivially_copyable_v<move_only>);

struct kernel {
  kernel();
  void operator()() const {}
  move_only mo;
};
static_assert(!std::is_trivially_copyable_v<move_only>);

int main() {
  sycl::queue q;
  kernel k{};
  q.single_task<struct KN>(k);
}

Curiously, DPC++ still rejects this example when sycl::is_device_copyable_v<...> specializations are added; https://godbolt.org/z/c6MbGzdKG. That seems wrong. It seems that DPC++ has a possibly unintentional dependency on a non-deleted copy constructor, even for types for which sycl::is_device_copyable_v<...> is true.

[gmlueck]

I volunteered to draft a PR for this.