2

カーネルは以下を使用します: ( --ptxas-options=-v)
0 バイトのスタック フレーム、0 バイトのスピル sotes、0 バイトのスピル ロード
ptxas 情報: 45 個のレジスタを使用、49152+0 バイトの smem、64 バイトの cmem[0]、12 バイトの cmem[16]

次のコマンドで起動するkernelA<<<20,512>>>(float parmA, int paramB);と、問題なく実行されます。
次で起動するとkernelA<<<20,513>>>(float parmA, int paramB);、リソース不足エラーが発生します。(起動に要求されたリソースが多すぎます)。

Fermi デバイスのプロパティ: SM ごとに 48KB の共有メモリ、定数メモリ 64KB、SM ごとに 32K レジスタ、ブロックごとに最大 1024 スレッド、comp 対応 2.1 (sm_21)

共有メモリ領域をすべて使用しています。約 700 スレッド/ブロックでブロック レジスタ スペースが不足します。MAX_threads/block の半分以上の数を要求すると、カーネルが起動しません。単なる偶然かもしれませんが、私はそれを疑っています。

  1. スレッドの完全なブロック (1024) を使用できないのはなぜですか?
  2. どのリソースが不足しているのか推測できますか?
  3. ストールしたスレッドのデータ/状態がワープ間でどこに行くのか、よく疑問に思います。これらを保持しているリソースは何ですか?
4

1 に答える 1

2

reg カウントを行ったとき、printf をコメントアウトしました。Reg count= 45
実行中は、printf がコーディングされていました。Reg カウント = 63 w/たくさんのこぼれ "reg's"。
各スレッドには実際には 64 個の reg があり、プログラムで使用できるのは 63 個だけだと思います。
64 reg * 512 スレッド = 32K - 1 つのブロックで使用できる最大値。

したがって、ブロックに使用可能な「コード」reg の数 = cudaDeviceProp::regsPerBlock - blockDim をお勧めします。つまり、カーネルは 32K レジスタすべてにアクセスできるわけではありません。コンパイラは現在、スレッドあたりの reg の数を 63 に制限しています (または lmem にスピルオーバーします)。この 63 は、HW アドレッシングの制限であると思われます。

そのため、レジスタスペースが不足しているようです。

于 2012-09-07T21:26:30.607 に答える