DPCT1089
目次
DPCT1089#
メッセージ#
sub-group サイズ属性の引数 <argument name> の値を評価できません。dpct_placeholder
を整数定数式に置き換えます。
詳細な説明#
属性の引数は整数定数式である必要があります。この警告は、インテル® DPC++ 互換性ツールが sub-group サイズの引数値を整数定数式として評価できない場合に出力されます。sub-group サイズ属性の引数を整数定数式に置き換えられるかどうかを確認し、置き換えられない場合はコードロジックを再設計します。
例えば、以下のオリジナル CUDA* コードについて考えてみます。
1 _global_ void kernel(int WarpSize) { int Input, Output, Lane; ...Output = __shfl(Input, Lane, WarpSize); }
2 ...
3 if(condition)
4 { kernel<<<GirdSize, BlockSize>>>(16); }
5 else
6 { kernel<<<GirdSize, BlockSize>>>(32); }
このコードは、以下の SYCL* コードに移行されます。
1 void kernel(int WarpSize, sycl::nd_item<3> item_ct1) {
2 int Input, Output, Lane; ... Output = Item_ct1.get_sub_group().shuffle(Input, Lane);
3 }
4 ...
5 if(condition) {
6 /* DPCT1089 */
7 q_ct1.parallel_for(
8 sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
9 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(dpct_placeholder)]] { kernel(16, item_ct1); });
10 } else {
11 /* DPCT1089 */
12 q_ct1.parallel_for(
13 sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
14 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(dpct_placeholder)]] { kernel(32, item_ct1); }
15 }
このコードを以下のように手動で調整します。
1 void kernel(int WarpSize, sycl::nd_item<3> item_ct1) { int Input, Output, Lane; ... Output = Item_ct1.get_sub_group().shuffle(Input, Lane); }
2 ...
3 if(condition) {
4 /* DPCT1089 */
5 q_ct1.parallel_for(
6 sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
7 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(16)]]
8 { kernel(16, item_ct1); }
9 );
10 } else {
11 /* DPCT1089 */
12 q_ct1.parallel_for(
13 sycl::nd_range<3>(GridSize * BlockSize, BlockSize),
14 [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]]
15 { kernel(32, item_ct1); }
16 }
修正方法の提案#
コードを手動で修正する必要があります。“dpct_placeholder” を整数定数式に置き換えます。