制限レジスタの使用：<a href="http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#ixzz4kk6SRKqY" rel="nofollow noreferrer">NVIDIA CUDA C Programming Guide</a>からmaxrregcount

Question

対__launch_bounds__：

レジスタの使用がmaxrregcountコンパイラ オプションを使用して制御または起動境界で説明したように境界を起動することができます。私の理解から

（と私が間違っているなら、私を修正）、-maxrregcount全体.cuファイルを使用することができ、レジスタの数を制限しながら、__launch_bounds__修飾子は、各__global__カーネル用maxThreadsPerBlockとminBlocksPerMultiprocessorを定義します。これらの2つは、同じタスクを2つの異なる方法で実行します。

私の使用方法では、性能を最大限に引き出すためにスレッドごとに40レジスタが必要です。したがって、-maxrregcount 40を使用することができます。また、__launch_bounds__(256, 6)を使用して40レジスタを強制的に実行することもできますが、これによりロードが発生します。&ストアレジスタのスピルが発生します。

この2つのレジスタスルールの違いは何ですか？

Answer 1

この質問の序文には、カーネルが使用するCUDA C Programming Guide、

少ないレジスタ、複数のスレッドとスレッドブロック を引用 パフォーマンスを向上させることができ、マルチプロセッサ、上に存在する可能性がある、ということです。

ここで、__launch_bounds__とmaxregcountは2つの異なるメカニズムによる制限レジスタの使用です。

__launch_bounds__

nvcc性能とカーネルの起動設定の一般性のバランスをとるを通じて__global__機能で使用するレジスタの数を決定します。別の言い方をすると、使用されるレジスタの数のこのような選択は、ブロック当たりのスレッド数とマルチプロセッサあたりのブロック数の違いに対して「有効性を保証します」。しかし、コンパイル時にブロックあたりのスレッドの最大数とマルチプロセッサあたりの最小ブロック数の可能な近似的な考え方が利用可能であれば、この情報を使ってそのような起動のためにカーネルを最適化することができます。 nvccが「最適」な方法でこのような起動設定用レジスタの数を選択することができるように言い換える

#define MAX_THREADS_PER_BLOCK 256 
#define MIN_BLOCKS_PER_MP  2 

__global__ void 
__launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP) 
fooKernel(int *inArr, int *outArr) 
{ 
    // ... Computation of kernel 
} 

は、おそらく起動構成のコンパイラに通知します。

パラメータは必須ですが、MIN_BLOCKS_PER_MPパラメータはオプションです。また、ブロックあたりのスレッド数がMAX_THREADS_PER_BLOCKより大きいカーネルを起動すると、カーネルの起動に失敗することにも注意してください。

次のように制限機構がProgramming Guideに記載されている：

を発射範囲が指定されている場合、コンパイラは、まずカーネルが確保する を使用すべきレジスタの数にそれらから 上限Lを導出しますブロック （minBlocksPerMultiprocessorが指定されていない場合）のminBlocksPerMultiprocessorブロックは、 maxThreadsPerBlockのスレッドをマルチプロセッサに置くことができます。 では通常、より多くの地元の費用を最初のレジスタの使用がLよりも高い場合、それは以下Lになるまで

は、コンパイラはそれをさらに低減し

• ： コンパイラは、次の方法での使用を登録最適化メモリ使用量および/またはより高い数の 命令;

したがって、__launch_bounds__は流出を登録するためにつながることができます。

maxrregcount

maxrregcount

単にレジスタの使用を再配置するようにコンパイラに強制することによって、__launch_bounds__と分散で、ユーザにより設定された数のために用いレジスタの数をhardlimitsコンパイラフラグです。コンパイラが課された制限値を下回ることができない場合、コンパイラはそれをローカルメモリに単に流し込みます。これは実際にはDRAMです。このローカル変数でさえも、グローバル変数DRAMに格納され、メモリ変数はL1、L2にキャッシュされます。