DPCT1049

メッセージ

SYCL* カーネルに渡される work-group サイズが制限を超えている可能性があります。デバイスの上限値は、info::device::max_work_group_size で照会できます。必要に応じて、work-group サイズを調整します。

詳細な説明

SYCL* デバイスの SYCL* カーネルに渡される work-group サイズには制限があります (SYCL* 1.2.1 仕様の「4.6.4.2 デバイス情報記述子」を参照)。

この警告は、ローカルレンジの次元をすべて評価できなかった場合や、ローカルレンジの次元の積が 256 以上の場合に表示されます。

修正方法の提案

info::device::max_work_group_size を照会して、使用するデバイスの work-group サイズの上限を定義します。コードで使用されている work-group サイズが制限値を下回っている場合は、この警告を無視できます。そうでない場合は、 work-group サイズを小さくする必要があります。

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


1 __global__ void k() {} 
2 
3 void foo() { 
4     k<<<1, 2048>>>(); 
5 }

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


1 void k() {} 
2 
3 void foo() { 
4   /* 
5   DPCT1049:0: The work-group size passed to the SYCL kernel may exceed the 
6   limit. To get the device limit, query info::device::max_work_group_size. 
7   Adjust the work-group size if needed. 
8   */ 
9   dpct::get_default_queue().parallel_for( 
10       sycl::nd_range<3>(sycl::range<3>(1, 1, 2048), sycl::range<3>(1, 1, 2048)), 
11       [=](sycl::nd_item<3> item_ct1) { 
12         k(); 
13       }); 
14 }

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


1 void k() {} 
2 
3 void foo() { 
4   size_t max_work_group_size = 
5      dpct::get_default_queue() 
6           .get_device() 
7           .get_info<sycl::info::device::max_work_group_size>(); 
8   size_t work_group_size = 2048; 
9   if (work_group_size > max_work_group_size) { 
10     work_group_size = max_work_group_size; 
11   } 
12   size_t work_group_num = std::ceil((float)2048 / (float)work_group_size); 
13   dpct::get_default_queue().parallel_for( 
14      sycl::nd_range<3>(sycl::range<3>(1, 1, work_group_num * work_group_size), 
15                sycl::range<3>(1, 1, work_group_size)), 
16      [=](sycl::nd_item<3> item_ct1) { k(); }); 
17   }