DPCT1115#

メッセージ#

sycl::ext::oneapi::group_local_memory_for_overwrite は、work-group のデータ並列カーネルの非カーネル・ファンクター・スコープでグループのローカルメモリーを割り当てる際に使用されます。ソースコードの調整が必要な場合があります。

詳細な説明#

sycl::ext::oneapi::group_local_memory_for_overwrite は、work-group のデータ並列カーネルの非カーネル・ファンクター・スコープでグループのローカルメモリーを割り当てる際に使用されます。グループローカル変数をカーネルファンクターのスコープで定義する必要がある制限は、この拡張機能の将来のバージョンで改善される可能性があります。

詳細については、sycl_ext_oneapi_local_memory.asciidoc をご覧ください。

修正方法の提案#

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

1  template <int S> __device__ void devfun() { 
2   __shared__ int slm1[32 * S]; 
3     ...
4  } 
5 
6  template <int S> __global__ void kernel() { 
7   __shared__ int slm2[S]; 
8   devfun<S>(); 
9  } 
10 
11 void hostfun() { kernel<256><<<1, 1>>>(); }

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

1  template <int S> inline void devfun(int *p, const sycl::nd_item<3> &item_ct1) { 
2   /* 
3   DPCT1115:0: The sycl::ext::oneapi::group_local_memory_for_overwrite is used to allocate 
4   group-local memory at the none kernel functor scope of a work-group data 
5   parallel kernel. You may need to adjust the code. 
6   */ 
7   auto &slm1 = 
8   *sycl::ext::oneapi::group_local_memory_for_overwrite<int[32 * S]>(item_ct1.get_group()); 
9   ... 
10 } 
11 
12 template <int S> __dpct_inline__ void kernel(const sycl::nd_item<3> &item_ct1) { 
13  auto &slm2 = 
14  *sycl::ext::oneapi::group_local_memory_for_overwrite<int[S]>(item_ct1.get_group()); 
15  devfun<S>(item_ct1); 
16 } 
17 
18 void hostfun() { dpct::get_default_queue().parallel_for( 
19  sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
20  [=](sycl::nd_item<3> item_ct1) { 
21  kernel<256>(item_ct1); 
22  }); 
23 }

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

1  template <int S> inline void devfun(int *slm1) { 
2      ... 
3  } 
4 
5  template <int S> __dpct_inline__ void kernel(int *slm1, int *slm2) { 
6 
7   devfun<S>(slm1); 
8  } 
9 
10 void hostfun() { dpct::get_default_queue().submit( 
11  [&](sycl::handler &cgh) { 
12  sycl::local_accessor<int, 1> slm1_acc_ct1(sycl::range<1>(32 * 256), cgh); 
13  sycl::local_accessor<int, 1> slm2_acc_ct1(sycl::range<1>(256), cgh); 
14 
15  cgh.parallel_for( 
16  sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
17  [=](sycl::nd_item<3> item_ct1) { 
18  kernel<256>(slm1_acc_ct1.get_pointer(), slm2_acc_ct1.get_pointer()); 
19  }); 
20 });