DPCT1118

メッセージ

SYCL* グループの関数とアルゴリズムは、統合された制御フローで使用される必要があります。

説明

SYCL* グループ関数とアルゴリズムは、work-ghroup 内のすべてのワーク項目の収束制御フローで使用する必要があります。ワーク項目がカーネルを完了し、バリアなどの 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* コードに移行されます。


1void 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}

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


1void 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}