DPCT1118#

メッセージ#

SYCL* グループ関数とアルゴリズムは、収束制御フローで検出される必要があります。

詳細な説明#

SYCL* グループ関数とアルゴリズムは、work-ghroup 内のすべてのワーク項目の収束制御フローで使用する必要があります。ワーク項目がカーネルを完了し、barrier などの work-group の集合操作に到達することなく終了した場合、集合操作に到達する work-group 内のほかのすべてのワーク項目は、完了したワーク項目を持ちます。

追加情報については、条件文で使用されるグループバリアなど、work-group レベルの同期が原因で SYCL* コードがハングする問題を修正するにはどうすればいいですか? を参照してください。

修正方法の提案

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

1  __global__ void kernel(float *data) { 
2   int tid = threadIdx.x; 
3   if (tid < 32) { 
4     if (data[tid] < data[tid + 32]) { 
5       data[tid] = data[tid + 32]; 
6     } 
7     __syncthreads(); 
8     ...
9   } 
10 }

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

1  void kernel(float *data, const sycl::nd_item<3> &item_ct1) { 
2   int tid = item_ct1.get_local_id(2); 
3   if (tid < 32) { 
4     if (data[tid] < data[tid + 32]) { 
5       data[tid] = data[tid + 32]; 
6     } 
7   /* 
8   DPCT1118:0: SYCL group functions and algorithms must be encountered in converged control flow.You should check this condition holds.
9   */ 
10  /* 
11  DPCT1065:1: Consider replacing sycl::nd_item::barrier() with sycl::nd_item::barrier(sycl::access::fence_space::local_space) for better performance if there is no access to global memory.
12  */ 
13  item_ct1.barrier(); 
14  } 
15 }

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

1  void kernel(float *data, const sycl::nd_item<3> &item_ct1) { 
2   int tid = item_ct1.get_local_id(2); 
3 
4   if (tid < 32) { 
5     if (data[tid] < data[tid + 32]) { 
6       data[tid] = data[tid + 32]; 
7     } 
8   } 
9   item_ct1.barrier(); 
10      ... 
11 }