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* コードに移行されます。

1void kernel(const sycl::nd_item<3> &item_ct1){ 
2 int target = 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 // CHECK: /* 
10 // CHECK: DPCT1108:{{[0-9]+}}: '__shfl_sync' was migrated with the experimental feature masked sub_group function which may not be supported by all compilers or runtimes. You may need to adjust the code. 11 // CHECK: */ 
12 foo = dpct::select_from_sub_group(mask, item_ct1.get_sub_group(), 10, 0); 
13 } 
14}

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

1void kernel(const sycl::nd_item<3> &item_ct1){ 
2 int target = 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 // Don't 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}