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* コードについて考えてみます。

  __global__ void kernel() { 
     ...
  value = __shfl_down(x, delta); 
     ... 
  } 

  void foo() { 
    ... 
    auto GridSize = dim3(2); 
   auto BlockSize = dim3(8, 8, 1); 
   kernel<<<GridSize, BlockSize>>>(); 
    ... 
}

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

  void kernel(const sycl::nd_item<3> &item_ct1) { 
    ... 
    /* 
    DPCT1096:0: The right-most dimension of the work-group used in the SYCL kernel 
    that calls this function may be less than "32".The function 
    "dpct::shift_sub_group_left" may return an unexpected result on the CPU 
    device.Modify the size of the work-group to ensure that the value of the 
    right-most dimension is a multiple of "32".
    */ 
   value = dpct::shift_sub_group_left(item_ct1.get_sub_group(), x, delta); // May return unexpected result on CPU 
   ... 
 } 

 void foo() { 
   ... 
   auto GridSize = sycl::range<3>(1, 1, 2); 
   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.   dpct::get_in_order_queue().parallel_for( 
   sycl::nd_range<3>(GridSize * BlockSize, BlockSize), 
     [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { 
     kernel(item_ct1); 
   }); 
   ... 
}

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

  void kernel(const sycl::nd_item<3> &item_ct1) { 
    ... 
    value = dpct::shift_sub_group_left(item_ct1.get_sub_group(), x, delta); 
    ... 
  } 

  void foo() { 
    ... 
    auto GridSize = sycl::range<3>(1, 1, 2); 
   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.   dpct::get_in_order_queue().parallel_for( 
   sycl::nd_range<3>(GridSize * BlockSize, BlockSize), 
     [=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] { 
     kernel(item_ct1); 
     }); 
   ... 
}

修正方法の提案

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