DPCT1121#

メッセージ#

SYCL* グループ関数/アルゴリズムで使用される <引数> が初期化されていることを確認します。

詳細な説明#

SYCL* 仕様 4.17.2 および 4.17.3 の SYCL* グループ関数/アルゴリズムの場合、入力スカラー値はグループ内のすべての work-item によって初期化される必要があります。これは、標準の C++ ルールとの互換性を保つため、またコンパイラーが過度に最適化しないようにするために必要です。

修正方法の提案#

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

1  __device__ void d(int var) { 
2   if (threadIdx.x == 0) 
3     var = 123; 
4   var = __shfl_sync(0xffffffff, var, 0); 
5  } 
6 
7  __global__ void kernel() { 
8   int var; 
9   d(var); 
10 }

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

1  void d(int var, const sycl::nd_item<3> &item_ct1) { 
2   if (item_ct1.get_local_id(2) == 0) 
3     var = 123; 
4   /* 
5   DPCT1121:0: Make sure that the "var" which is used in the SYCL group 
6   function/algorithm is initialized.7   */ 
8   var = dpct::select_from_sub_group(item_ct1.get_sub_group(), var, 0); 
9  } 
10 void kernel(const sycl::nd_item<3> &item_ct1) { 
11  int var; 
12  d(var, item_ct1); 
13 }

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

1  void d(int var, const sycl::nd_item<3> &item_ct1) { 
2   if (item_ct1.get_local_id(2) == 0) 
3     var = 123; 
4   var = dpct::select_from_sub_group(item_ct1.get_sub_group(), var, 0); 
5  } 
6  void kernel(const sycl::nd_item<3> &item_ct1) { 
7   int var = 0; 
8   d(var, item_ct1); 
9  }