DPCT1108#

メッセージ#

<original API> は、実験的な機能 <feature name> と共に移行されましたが、これはすべてのコンパイラーでサポートされるわけではありません。ソースコードの調整が必要な場合があります。

詳細な説明#

より適切な移行をサポートするため、インテル® DPC++ 互換性ツールには、移行中に実験的な機能を有効にする -use-experimental-features オプションが用意されています。<feature name> は、すべての SYCL* コンパイラーやランタイムでサポートされるわけではありません。ターゲットの SYCL* コンパイラーやランタイムが <feature name> をサポートしない場合、コードを調整して <feature name> を使用しないようにします。

修正方法の提案#

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

1  __global__ void kernel() { 
2   int lane_id = threadIdx.x % 32; 
3   int foo = 0, result = 0; 
4   int mask = 0xf; 
5   if (lane_id == 0) { 
6     result = 10; 
7   } 
8   if (lane_id & mask) { 
9     foo = __shfl_sync(mask, result, 0); 
10  } 
11 }

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

1  void kernel(const sycl::nd_item<3> &item_ct1) { 
2   int lane_id = item_ct1.get_local_id(2) % 32; 
3   int foo = 0, result = 0; 
4   int mask = 0xf; 
5   if (lane_id == 0) { 
6     result = 10; 
7   } 
8   if (lane_id & mask) { 
9   /* 
10  DPCT1108:0: '__shfl_sync' was migrated with the experimental feature masked 
11  sub_group function which may not be supported by all compilers or runtimes.
12  You may need to adjust the code.
13  */ 
14    foo = dpct::experimental::select_from_sub_group( 
15    mask, item_ct1.get_sub_group(), result, 0); 
16  } 
17 }

このコードは次のように書き換えられます。

1  void kernel(const sycl::nd_item<3> &item_ct1) { 
2   int lane_id = item_ct1.get_local_id(2) % 32; 
3   int foo = 0, result = 0; 
4   int mask = 0xf; 
5   if (lane_id == 0) { 
6     result = 10; 
7   } 
8   // If not use experimental feature masked sub-group function 
9   int foo_tmp = dpct::select_from_sub_group(item_ct1.get_sub_group(), result, 0); 
10  if (lane_id & mask) { 
11    foo = foo_tmp; 
12  } 
13 }