DPCT1110

メッセージ

デバイス関数 <function name> で宣言されたローカス変数の合計サイズが 128 バイトを超えているため、レジスター・プレッシャーが高くなる可能性があります。ハードウェア・ベンダーに問い合わせて、利用可能な合計レジスターサイズを調べてコードを調整するか、小さな sub-group サイズを使用してレジスタープレッシャーを回避します。

詳細な説明

特定のハードウェア構成では、各ワーク項目が利用できるレジスター数が制限されます。例えば、インテル® X^e-LP GPU アーキテクチャーでは、各ハードウェア・スレッドには 4KB のレジスターがあります。従って、sub-group サイズが 32 の場合、各ワーク項目は 128 バイトのレジスターを利用できます (4KB/32)。デバイス関数で宣言されたローカル変数のサイズが 128 バイトを超える場合、一部の変数がローカルメモリーやグローバルメモリーに格納される可能性があるため、頻繁にアクセスされるとパフォーマンスが低下することがあります。この問題に対処するには、sub-group サイズを減らして各ワーク項目で利用できるレジスター数を増やすか、oneAPI GPU 最適化ガイドの「レジスタースピルの最適化」の推奨事項を参照してください。他のハードウェアについては、ハードウェア・ベンダーに問い合わせて構成情報を入手してください。

修正方法の提案

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

  __global__ void Kernel(){ 
   int result[50]; 
    ...
  } 
  int main{ 
     ... 
   Kernel<<<1, 100>>>(); 
  }

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

 /* 
 DPCT1110:0: The total declared local variable size in device function "Kernel" exceeds 128 bytes and may cause high register pressure (デバイス関数 "Kernel" で宣言されたローカス変数の合計サイズが 128 バイトを超えているため、レジスター・プレッシャーが高くなる可能性があります).Consult with your hardware vendor to find the total register size available and adjust the code or use smaller sub-group size to avoid high register pressure (ハードウェア・ベンダーに問い合わせて、利用可能な合計レジスターサイズを調べてコードを調整するか、小さな sub-group サイズを使用してレジスタープレッシャーを回避します).
 */ 
 void Kernel(){ 
  int result[50]; 
  ... 
 } 
 int main{ 
  ... 
 q.parallel_for(sycl::range(100), [=](sycl::nd_item<3> item) [[intel::reqd_sub_group_size(32)]] { Kernel(); }); 
}

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

 void Kernel(){ 
  int result[50]; 
     ... 
 } 

 int main{ 
     ... 
  /* 
  Reduce sub_group size to make more registers available for each work-item, which may help to avoid high register pressure.
 */ 
 q.parallel_for(sycl::range(100), [=](sycl::nd_item<3> item) [[intel::reqd_sub_group_size(16)]] { Kernel(); }); 
}