15

カーネルに未使用のレジスタがたくさんあります。必要なときに毎回グローバルデータを読み取るのではなく、いくつかのレジスタを使用してデータを保持するようにCUDAに指示したいと思います。（共有メモリは使用できません。）

__global__ void simple(float *gData) {
float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

コンパイルw/：nvcc -arch sm_20 --ptxas-options = -v simple.cu、
0バイトのスタックフレーム、0バイトのスピルストア、0バイトのスピルロード
使用2レジスタ、40バイトcmem [0]

__global__ void simple(float *gData) {
register float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

レジスタ宣言は何もしません。
0バイトのスタックフレーム、0バイトのスピルストア、0バイトのスピルロード
使用された2つのレジスタ、40バイトのcmem [0]

__global__ void simple(float *gData) {
volatile float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

volatile宣言はスタックストレージを作成します：
4096バイトのスタックフレーム、0バイトのスピルストア、0バイトのスピルロード
使用21レジスタ、40バイトcmem [0]

1）変数にレジスタスペースを使用するようにコンパイラに指示する簡単な方法はありますか？
2）「スタックフレーム」はどこにありますか：レジスタ、グローバルmem、ローカルmem、...？スタックフレームとは何ですか？（GPUにはいつスタックがありますか？仮想スタックですか？）
3）simple.ptxファイルは基本的に空です：（nvcc -arch sm_20 -ptx simple.cu）

.loc 2 14 2
ret;

実際のマシン/コンパイルされたコードをどこで見つけることができるか考えていますか？

4

2 に答える 2

22
  • GPUレジスタファイルは動的にアドレス指定できないため、動的にインデックス付けされた配列をレジスタに格納することはできません。
  • スカラー変数は、コンパイラーによってレジスターに自動的に保管されます。
  • 静的にインデックス付けされた（つまり、コンパイル時にインデックスを決定できる場合）、小さな配列（たとえば、16未満のfloat）は、コンパイラーによってレジスターに格納される場合があります。

SM 2.0 GPU（Fermi）は、スレッドごとに最大63個のレジスタのみをサポートします。これを超えると、レジスタ値がローカル（オフチップ）メモリからスピル/フィルされ、キャッシュ階層でサポートされます。SM 3.5 GPUは、これをスレッドあたり最大255レジスタに拡張します。

一般に、Jaredが言及しているように、スレッドごとにあまりにも多くのレジスターを使用すると、占有率が低下し、カーネルのレイテンシー隠蔽機能が低下するため、望ましくありません。GPUは並列処理で繁栄し、他のスレッドからの作業でメモリの待ち時間をカバーすることでそれを実現します。

したがって、配列をレジスタに最適化するべきではありません。代わりに、スレッド間でこれらの配列へのメモリアクセスが可能な限りシーケンシャルに近いことを確認して、合体を最大化します（つまり、メモリトランザクションを最小化します）。

あなたが与える例は、次の場合に共有メモリの場合かもしれません：

  1. ブロック内の多くのスレッドが同じデータを使用する、または
  2. スレッドごとの配列サイズは、複数のスレッドブロック内のすべてのスレッドに十分なスペースを割り当てるのに十分なほど小さいです（スレッドごとに1024フロートははるかに大きいです）。

njuffaが述べたように、カーネルが2つのレジスタのみを使用する理由は、カーネル内のデータに対して何も有用なことを行わず、デッドコードがすべてコンパイラによって排除されたためです。

于 2012-08-29T01:04:01.587 に答える
6

すでに述べたように、レジスタ（およびPTX「パラメータスペース」）は動的にインデックス付けできません。そのためには、コンパイラはswitch...case動的インデックスをイミディエートに変換するためのブロックに関するコードを出力する必要があります。自動的に実行されるかどうかはわかりません。固定サイズのタプル構造とを使用して、それが発生するのを助けることができますswitch...case。C / C ++メタプログラミングは、このようなコードを管理しやすくするための武器になる可能性があります。

また、CUDA 4.0の場合、コマンドラインスイッチ-Xopencc=-O3を使用して、プレーンなスカラー（データ構造など）以外のものをレジスターにマップします（この投稿を参照）。CUDA> 4.0の場合、デバッグサポートを無効にする必要があります（-Gコマンドラインオプションなし-最適化はデバッグが無効になっている場合にのみ発生します）。

PTXレベルでは、ハードウェアよりもはるかに多くの仮想レジスタを使用できます。これらは、ロード時にハードウェアレジスタにマップされます。指定するレジスタ制限により、生成されたバイナリによって使用されるハードウェアリソースに上限を設定できます。これは、コンパイラがPTXにコンパイルするときにレジスタをいつスピルするか（以下を参照）を決定するヒューリスティックとして機能するため、特定の同時実行のニーズを満たすことができます（CUDAドキュメントの「起動境界」、「占有」、および「同時カーネル実行」を参照）。 -この最も興味深いプレゼンテーションもお楽しみいただけます）。

Fermi GPUの場合、最大64個のハードウェアレジスタがあります。64番目（または最後-ハードウェアの最大値未満を使用する場合）は、ABIによってスタックポインターとして使用され、したがって「レジスタスピル」（スタックに値を一時的に格納することによってレジスタを解放することを意味し、より多くのレジスタが発生した場合に発生します）に使用されます利用可能よりも必要です）ので、触れられません。

于 2012-08-29T20:28:04.227 に答える