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_groupdpct::shift_sub_group_leftdpct::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 int main() { 
8 ... 9 auto GridSize = dim3(2); 
10 auto BlockSize = dim3(8, 8, 1); 
11 kernel<<<GridSize, BlockSize>>>(); 
12 ... 13 }

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

1 void kernel(sycl::nd_item<3> item) { 
2 ... 3 value = dpct::shift_sub_group_left(item.get_sub_group(), x, delta); // May return unexpected result on CPU 
4 ... 5 } 
6 
7 int main() { 
8 ... 9 auto GridSize = sycl::range<3>(1, 1, 2); 
10 auto BlockSize = sycl::range<3>(1, 8, 8); // Problem: value of the right-most dimension 8 is less than the kernel sub group size attribute 32. 11 queue.parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item)[[intel::reqd_sub_group_size(32)]] { 
12 kernel(item); 
13 }); 
14 ... 15 }

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

1 void kernel(sycl::nd_item<3> item) { 
2 ... 3 value = dpct::shift_sub_group_left(item.get_sub_group(), x, delta); 
4 ... 5 } 
6 
7 int main() { 
8 auto GridSize = sycl::range<3>(1, 1, 2); 
9 auto BlockSize = sycl::range<3>(1, 2, 32); // Fix: modified work group size to make the right-most dimension to be multiple of the kernel sub group size attribute value, which is 32. 10 queue.parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item)[[intel::reqd_sub_group_size(32)]] { 
11 kernel(item); 
12 }); 
13 ... 14 }

修正方法の提案#

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