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); 
   ... 
} 

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

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

void kernel(sycl::nd_item<3> item) { 
   ... 
   value = dpct::shift_sub_group_left(item.get_sub_group(), x, delta); // May return unexpected result on CPU 
   ... 
} 

int main() { 
   ... 
   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. 
   queue.parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item)[[intel::reqd_sub_group_size(32)]] { 
      kernel(item); 
   }); 
   ... 
}

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

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

int main() { 
   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. 
  queue.parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item)[[intel::reqd_sub_group_size(32)]] { 
      kernel(item); 
  }); 
   ... 
}

修正方法の提案

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