DPCT1128#

メッセージ#

<reason> により、<type name> はデバイスでコピーできません。SYCL* カーネルで使用されているため、コードを書き直してください。

詳細な説明#

SYCL* 2020 では、カーネル関数に渡される引数のタイプがデバイスでコピー可能である必要があります (仕様 3.13.1 デバイスでコピー可能)。一方、CUDA* では、エラーや警告を出力することなく、すべての引数のメモリーコピーを直接実行します。タイプがデバイスのコピー可能な要件を満たしていない場合、ツールはこのタイプに対して sycl::is_device_copyable の特殊化を追加することをユーザーに警告します。警告では、そのタイプのどの部分が要件を満たしていないか示されます。

修正方法の提案#

例えば、以下のオリジナル CUDA* コードについて考えてみます。

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

このコードは、以下の SYCL* コードに移行されます。

1   /* 
2   DPCT1128:0: The type "UserStruct<float, int>" is not device copyable for copy 
3   constructor breaking the device copyable requirement.It is used in the SYCL 
4   kernel, please rewrite the code.
5  */ 
6   template <class T1, class T2> struct UserStruct { 
7     UserStruct() {} 
8     UserStruct(UserStruct const &other) : a(other.a) {} 
9     int a; 
10 }; 
11 template <class T1, class T2> 
12 struct sycl::is_device_copyable<UserStruct<T1, T2>> : std::true_type {}; 
13 
14 template <class V1, class V2> void k(UserStruct<V1, V2>) {} 
15 
16 void 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 }

上記は次のように書き換える必要があります。

1  template <class T1, class T2> struct UserStruct { 
2     UserStruct() {} 
3     UserStruct(UserStruct const &other) : a(other.a) {} 
4     int a; 
5  }; 
6  template <class T1, class T2> 
7  struct sycl::is_device_copyable<UserStruct<T1, T2>> : std::true_type {}; 
8 
9  template <class V1, class V2> void k(UserStruct<V1, V2>) {} 
10 
11 void 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 }
1  template <class T1, class T2> struct UserStruct { 
2   UserStruct() {} 
3   int a; 
4  }; 
5 
6  template <class V1, class V2> void k(UserStruct<V1, V2>) {} 
7 
8  void 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 }