DPCT1130

目次

DPCT1130#

メッセージ#

SYCL* 2020 標準は、動的並列処理 (デバイスコードでのカーネルの起動) をサポートしていません。コードを書き換えてください。

詳細な説明#

SYCL* はデバイスコードでのカーネル起動をサポートしていません。親カーネルと子カーネルをマージする必要があります。

修正方法の提案#

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

  __global__ void childKernel() { 
  ...
  } 
  __global__ void parentKernel() { 
   ... 
   childKernel<<<4, 4>>>(); 
   ... 
  } 
  void foo() { 
  ... 
  parentKernel<<<8, 8>>>(); 
  ... 
 }

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

  void childKernel() { 
   ... 
  } 
  void parentKernel() { 
   ... 
   /* 
   DPCT1130:0: SYCL 2020 standard does not support dynamic parallelism (launching 
   kernel in device code). Please rewrite the code. 
   */ 
  childKernel<<<4, 4>>>(); 
  ... 
 } 
 void foo() { 
  ... 
  dpct::get_in_order_queue().parallel_for( 
  sycl::nd_range<3>(sycl::range<3>(1, 1, 8) * sycl::range<3>(1, 1, 8), 
  sycl::range<3>(1, 1, 8)), 
  [=](sycl::nd_item<3> item_ct1) { 
  parentKernel(); 
  }); 
  ... 
 }

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

  void childKernel() { 
   ... 
  } 
  void parentKernel() { 
   ... 
   childKernel(); // childKernel() をデバイス関数として呼び出し、ワークを調整する必要があります 
   for each work item.
      ... 
  } 
 void foo() { 
  ... 
  dpct::get_in_order_queue().parallel_for( 
  sycl::nd_range<3>(sycl::range<3>(1, 1, placeholder /* スレッドモデルに基づいて、 
  親カーネルと子カーネル間のグローバル範囲を調整します */), 
  sycl::range<3>(1, 1, placeholder /* スレッドモデルに基づいて、 
  親カーネルと子カーネル間のローカル範囲を 
  調整します */)), 
  [=](sycl::nd_item<3> item_ct1) { 
  parentKernel(); 
  }); 
  ... 
 }