DPCT1128#

Message#

The type “<type name>” is not device copyable for <reason>. It is used in the SYCL kernel, please rewrite the code.

Detailed Help#

SYCL* 2020 needs the type of argument passed to the kernel function to be device copyable (Spec 3.13.1 device copyable), while CUDA* just does memory copy for all those arguments directly without emitting any error/warning. If a type does not meet the requirements of device copyable, the tool will try to add the specialization of the sycl::is_device_copyable for this type, as well as warn the user. The warning will specify which part of that type does not meet the requirements.

Suggestions to Fix#

For example, this original CUDA code:

 1template <class T1, class T2> struct UserStruct {
 2  UserStruct() {}
 3  UserStruct(UserStruct const &other) : a(other.a) {}
 4  int a;
 5};
 6
 7template <class V1, class V2> __global__ void k(UserStruct<V1, V2>) {}
 8
 9void foo() {
10  UserStruct<float, int> us;
11  k<<<1, 1>>>(us);
12}

results in the following migrated SYCL code:

 1/*
 2DPCT1128:0: The type "UserStruct<float, int>" is not device copyable for copy
 3constructor breaking the device copyable requirement. It is used in the SYCL
 4kernel, please rewrite the code.
 5*/
 6template <class T1, class T2> struct UserStruct {
 7  UserStruct() {}
 8  UserStruct(UserStruct const &other) : a(other.a) {}
 9  int a;
10};
11template <class T1, class T2>
12struct sycl::is_device_copyable<UserStruct<T1, T2>> : std::true_type {};
13
14template <class V1, class V2> void k(UserStruct<V1, V2>) {}
15
16void foo() {
17  UserStruct<float, int> us;
18  /*
19  DPCT1129:1: The type "UserStruct<float, int>" is used in the SYCL kernel, but
20  it is not device copyable. The sycl::is_device_copyable specialization has
21  been added for this type. Please review the code.
22  */
23  dpct::get_in_order_queue().parallel_for(
24      sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
25      [=](sycl::nd_item<3> item_ct1) {
26        k(us);
27      });
28}

which needs to be rewritten to:

 1template <class T1, class T2> struct UserStruct {
 2  UserStruct() {}
 3  UserStruct(UserStruct const &other) : a(other.a) {}
 4  int a;
 5};
 6template <class T1, class T2>
 7struct sycl::is_device_copyable<UserStruct<T1, T2>> : std::true_type {};
 8
 9template <class V1, class V2> void k(UserStruct<V1, V2>) {}
10
11void foo() {
12  UserStruct<float, int> us;
13  dpct::get_in_order_queue().parallel_for(
14      sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
15      [=](sycl::nd_item<3> item_ct1) {
16        k(us);
17      });
18}

 1template <class T1, class T2> struct UserStruct {
 2  UserStruct() {}
 3  int a;
 4};
 5
 6template <class V1, class V2> void k(UserStruct<V1, V2>) {}
 7
 8void foo() {
 9  UserStruct<float, int> us;
10  dpct::get_in_order_queue().parallel_for(
11      sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
12      [=](sycl::nd_item<3> item_ct1) {
13        k(us);
14      });
15}