2013年08月30日

RegisterとLmem,PTX

以下の最大値を見つけるだけのデバイス関数をデバッグ実行しようとしたら,
lmemが192Byteも使われていて悩んでいたので,メモ.

まずは,関数で
7 too many resources requested for launch
30 Unknown Error
などが発生していた.
各GPUで使用できるRegister数などのメモリ量が異なる.
(デバイスqueryで確認できる http://blog.tsp.me.uk/article/72126659.html )

で,各関数で使用されているRegister数などを確認しないといけない.
VC++では,通常の設定だとPTXが見つからなかったので,
プロパティ->構成->CUDA C/C++->Device Verbose PTXAS Output をYes
(もし最大Register数を指定するのであれば,Max Used Register に値を指定. 例;1,024threads/block使用して,許容が65,536Register/Blockならば,64を指定.)
これで,コンパイル時の出力に,各グローバルデバイス関数での各メモリ使用量が分かる(/thread)

ついでに,CUDA C/C++->Host Additional Compiler Options に-keepを付けておくと,PTXが(消されずに?)作られています.


「出力」に出力されるPTXASを確認すると,以下の関数で,
10 registers, 8216Byte shared mem, 24B cmem, 192B lmem(Local mem)
使用していました.
GTX 780は65,536Registers/Blockなので,1024threadsでも64個使えます.ので,Local使用するのはおかしいなと.Registerなら1クロックなのに,ローカルは100かかえるし.
Local memは-maxrregcount=64での使用量を超えないと使用されないと思い込んでいたので.

で,仕方ないので()プログラミングガイドを検索.しかし,原因となりそうな記述は見つからず.
・許容Register数を超えていない
・数学関数を使用していない
ので.

ここで,デバッグ情報に気付いて,リリースビルドしてみると,
13regs, 8216B shared, 12B cmemで,ローカルの記述が消えた.


あと,関数の実行後にGetLastErrorして問題なくても,その後変数をHostにコピーしようとしてUnknown Errorする場合は,
変数を確保し忘れているなどのミスの可能性もある.
変数がMallocされていなくても動作する(Pointerで動作するから?)




__global__ void max_search(double *D_e)
{
//extern
__shared__ double error[1024]; //収束誤差の計算用
int id= RTN_threadIdx_inBlock(threadIdx);
int3 ijk= RTN_index_ijk(blockIdx, threadIdx, false);
int i= ijk.x;
int j= ijk.y;
int k= ijk.z;
if(i
posted by にゃんこ at 16:21| Comment(0) | CUDA C
この記事へのコメント
コメントを書く
お名前:

メールアドレス:

ホームページアドレス:

コメント: [必須入力]