DPCT1065#

メッセージ#

グローバルメモリーにアクセスしない場合、パフォーマンスを向上するには sycl::<...>::barrier() を sycl::<...>::barrier(sycl::access::fence_space::local_space) に置き換えることを検討してください。

詳細な説明#

sycl::<...>::barrier() 関数は、グローバルおよびローカルアドレス空間で適切なメモリーアクセス順序を保証します。カーネル関数がグローバルメモリー内のメモリーアクセスを行わない場合、パフォーマンスを向上するため sycl::<...>::barrier() を sycl::<...>::barrier(sycl::access::fence_space::local_space) に置き換えても安全です。

修正方法の提案#

sycl::<...>::barrier() を sycl::<...>::barrier(sycl::access::fence_space::local_space) に置き換えます。

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

1  struct Data_t { 
2   float *host_data; 
3   float *device_data; 
4  }; 
5 
6  __global__ void k(Data_t *data) { 
7   auto tid = threadIdx.x + blockDim.x * blockIdx.x; 
8   only_read_data(data[tid].device_data); 
9   __syncthreads(); 
10  only_read_data(data[tid].device_data); 
11 }

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

1  struct Data_t { 
2   float *host_data; 
3   float *device_data; 
4  }; 
5 
6  void k(Data_t *data, const sycl::nd_item<3> &item_ct1) { 
7   auto tid = item_ct1.get_local_id(2) + 
8   item_ct1.get_local_range(2) * item_ct1.get_group(2); 
9   only_read_data(data[tid].device_data); 
10  /* 
11  DPCT1065:0: Consider replacing sycl::nd_item::barrier() with 
12  sycl::nd_item::barrier(sycl::access::fence_space::local_space) for better 
13  performance if there is no access to global memory.
14  */ 
15  item_ct1.barrier(); 
16  only_read_data(data[tid].device_data); 
17 }

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

1  struct Data_t { 
2   float *host_data; 
3   float *device_data; 
4  }; 
5 
6  void k(Data_t *data, const sycl::nd_item<3> &item_ct1) { 
7   auto tid = item_ct1.get_local_id(2) + 
8   item_ct1.get_local_range(2) * item_ct1.get_group(2); 
9   only_read_data(data[tid].device_data); 
10  // global_local_space can be replaced with local_space if the access 
11  // of the global memory after the barrier does not depend on (read-after-write or 
12  // write-after-read or write-after-write) the access of the same global memory 
13  // before the barrier among work-items in the current work-group. 
14  item_ct1.barrier(sycl::access::fence_space::local_space); 
15  only_read_data(data[tid].device_data); 
16 }