DPCT1049#

メッセージ#

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

詳細な説明#

SYCL* デバイスの SYCL* カーネルに渡される work-group サイズには制限があります (SYCL* 2020 仕様の「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  }