CUDA Compiler Driver NVCC - Options for Steering GPU code generationには、あいまいな行があります。
ABI が必要とする最小レジスター未満の値は、コンパイラーによって ABI の最小制限まで引き上げられます。
__global__
ABI には、__device__
関数が使用するレジスタの数に関する標準または制限はありますか?
CUDA Compiler Driver NVCC - Options for Steering GPU code generationには、あいまいな行があります。
ABI が必要とする最小レジスター未満の値は、コンパイラーによって ABI の最小制限まで引き上げられます。
__global__
ABI には、__device__
関数が使用するレジスタの数に関する標準または制限はありますか?
CUDA ABIには少なくとも16個のレジスタが必要だと思います(現在は参照が見つかりません)。そのため、(-maxrregcount などを使用して) より低いレジスタ カウントを指定すると、コンパイラは、指定された制限を ABI が必要とする最小値まで引き上げ、そのことを示すアドバイス メッセージを出力します。スレッドごとに使用できる 32 ビット レジスタの最大数については、GPU アーキテクチャに依存します。sm_1x では 124 レジスタ、sm_2x では 63 レジスタ、sm_3x では 254 レジスタです。
一般的に言えば、ABI (アプリケーション バイナリ インターフェイス) は、ストレージ レイアウト、関数への引数の受け渡し、呼び出し元への関数結果の受け渡しなどに関するアーキテクチャ固有の規則です。ABI (x86_64、ARM を含む) は、多くの場合、特定のタスクに対して特定のレジスタを指定します。スタック ポインター、関数の戻り値、関数の引数など。GPU アーキテクチャではスレッドごとに可変数のレジスターが許可されているため、ABI を使用するには、これらの定義された役割を満たすために最小限の数のレジスターが存在する必要があります。私の記憶が正しければ、CUDA はバージョン 3.0 で ABI を導入しました。これは Fermi クラスの GPU をサポートする最初のバージョンでした。
ABI には、コンピューティング機能 2.0 以上が必要です。古い GPU アーキテクチャには、ABI に必要なハードウェア機能がありませんでした。デバイス側の printf() や malloc() などの新しい CUDA 機能のほとんどは、関数の呼び出し、個別のコンパイルなど、ABI に依存しており、ABI を使用する必要があり、sm_20 およびその上。-Xptxas -abi=no で ABI の使用を無効にすることができます。そうしないことを強くお勧めします。