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 // バリアー後のグローバルメモリーのアクセスが、現在のワークグループ内のワークアイテム間のバリアー前の
11 // 同じグローバルメモリーのアクセス (読み取り後書き込み、
12 // 読み取り後書き込み、または書き込み後書き込み) に依存しない場合は、
13 // global_local_space を local_space に置き換えることができます。
14 item_ct1.barrier(sycl::access::fence_space::local_space);
15 only_read_data(data[tid].device_data);
16 }