Can pointer to data member types be used in device code? Are they device copyable?

[tahonermann]

Is the following code intended to be valid? The example has a SYCL kernel that captures a variable (proj) of pointer to data member type int P::*.

#include <sycl/sycl.hpp>
struct P {
  int x;
};
int main() {
  sycl::queue queue;
  queue.submit([&](sycl::handler &cgh) {
    auto proj = &P::x;
    cgh.single_task([proj]{});
  });
} 

There are no uses of the "pointer-to-member", "pointer to member", or "member pointer" terms in the SYCL 2020 specification. It is unclear whether that is intentional or due to an oversight.

Section 5.4, "Language restrictions for device functions", of the SYCL 2020 specification does not contain wording that restricts use of data member pointers in device code.

Section 3.13.1, "Device copyable", of the SYCL 2020 specification states:

Any type that is trivially copyable (as defined by the C++ core language) is implicitly device copyable.

C++17 [basic.types]p9 states:

Arithmetic types (6.9.1), enumeration types, pointer types, pointer to member types (6.9.2), std::nullptr_t, and cv-qualified (6.9.3) versions of these types are collectively called scalar types. Scalar types, POD classes (Clause 12), arrays of such types and cv-qualified versions of these types are collectively called POD types. Cv-unqualified scalar types, trivially copyable class types (Clause 12), arrays of such types, and cv-qualified versions of these types are collectively called trivially copyable types.

Section 4.12.4, "Rules for parameter passing to kernels", of the SYCL 2020 specification states:

Any device copyable type is a legal parameter type.

The above appears to indicate that data member pointers are allowed in device code, are trivially copyable, and are eligible for use as a kernel parameter.

Data member pointers are problematic for ABI reasons. Consider a compiler that targets the Microsoft ABI for host code on Windows but uses the Itanium ABI for device code. The Microsoft ABI does not use a uniform underlying data type, size, and alignment for all data member pointers; the data member pointer representation for a particular class depends on whether the class is incomplete or uses single, multiple, or virtual inheritance. See Microsoft's documentation for the /vmb and /vmg options, the inheritance keywords, and the pointers_to_members pragma. The Itanium ABI does use a uniform underlying data type of type ptrdiff_t for all data member pointers and thus all such types have the same size and alignment. See section 2.3.1, "Data Member Pointers", of the Itanium ABI. The existing wording appears to require such implementations to somehow coerce data member pointers from one representation to the other when passed as kernel arguments or otherwise copied between the host and devices. That might be feasible for single variables of pointer to data member type, but quickly becomes untenable when data member pointers are themselves data members of device copyable class types. In the example code above, the lambda that captures proj and is used as a SYCL kernel has a different size on the host vs the device for such implementations.

[nliber]

I think the reasoning in the comments on #373 apply here as well. It isn't sound if the size of a type can change between the host and the device, and can trivially lead to an ODR violation.

Take std::integral_constant<size_t, sizeof(proj)>. Should it take the sizeof(proj) on the host or device?