DPCT1096

DPCT1096#

メッセージ#

この関数を呼び出す SYCL* カーネルで使用される work-group の右端の次元は、<value of kernel sub-group size attribute> (カーネル sub-group サイズ属性値) 未満である可能性があります。関数 <help function name> は、CPU デバイス上で予期しない結果を返す可能性があります。work-group サイズを変更して、右端の次元の値が <value of kernel sub-group size attribute> (カーネル sub-group サイズ属性値) の倍数になるようにします。

詳細な説明#

dpct::select_from_sub_group、dpct::shift_sub_group_left、dpct::shift_sub_group_right、および dpct::permute_sub_group_by_xor 関数は、work-group の右端の次元値が OpenCL* バックエンドを備えた CPU デバイスで実行すると、予期しない結果を返す場合があります。これらの関数を呼び出す SYCL* カーネルで使用されるサイズは、カーネルの sub-group サイズ属性値よりも小さくなります。実際の sub-group サイズは、カーネルの sub-group サイズ属性で指定された値ではない可能性があり、CPU デバイス上のヘルパー関数が予期しない結果を返す可能性があります。

work-group サイズを変更してコードを調整し、右端の次元の値がカーネルの sub-group サイズ属性値) の倍数になるようにします。

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


1   __global__ void kernel() { 
2     ...
3   value = __shfl_down(x, delta); 
4     ... 
5   } 
6 
7   void foo() { 
8     ... 
9   auto GridSize = dim3(2); 
10  auto BlockSize = dim3(8, 8, 1); 
11  kernel<<<GridSize, BlockSize>>>(); 
12    ... 
13  }

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


1   void kernel(const sycl::nd_item<3> &item_ct1) { 
2     ... 
3   /* 
4   DPCT1096:0: The right-most dimension of the work-group used in the SYCL kernel 
5   that calls this function may be less than "32".The function 
6   "dpct::shift_sub_group_left" may return an unexpected result on the CPU 
7   device.Modify the size of the work-group to ensure that the value of the 
8   right-most dimension is a multiple of "32".
9   */ 
10  value = dpct::shift_sub_group_left(item_ct1.get_sub_group(), x, delta); // CPU で予期しない結果が返される可能性があります 
11    ... 
12  } 
13 
14  void foo() { 
15    ... 
16  auto GridSize = sycl::range<3>(1, 1, 2); 
17  auto BlockSize = sycl::range<3>(1, 8, 8); // 問題: 右端の次元 8 の値がカーネル サブグループ サイズ属性 32 より小さいです 
18  dpct::get_in_order_queue().parallel_for( 
19  sycl::nd_range<3>(GridSize * BlockSize, BlockSize), 
20  [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { 
21  kernel(item_ct1); 
22  }); 
23    ... 
24  }

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


1   void kernel(const sycl::nd_item<3> &item_ct1) { 
2     ... 
3   value = dpct::shift_sub_group_left(item_ct1.get_sub_group(), x, delta); 
4     ... 
5   } 
6 
7   void foo() { 
8     ... 
9   auto GridSize = sycl::range<3>(1, 1, 2); 
10  auto BlockSize = sycl::range<3>(1, 2, 32); // 修正: work-group サイズを変更し、右端の次元がカーネルのサブグループ・サイズ属性値の倍数 (32) になるようにしました。 
11  dpct::get_in_order_queue().parallel_for( 
12  sycl::nd_range<3>(GridSize * BlockSize, BlockSize), 
13  [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { 
14  kernel(item_ct1); 
15  }); 
16    ... 
17  }

修正方法の提案#

プログラムを CPU デバイス上で実行する場合、コードの調整が必要になることがあります。