08/07/14 11:56:25
問題は、自宅が規制に巻き込まれていてなかなか書き込めない辺り。
取り敢えず、注意点を列挙しておく。
・共有メモリを確保するのはglobalFunc<<<blocks, threads, sharedMemorySize, streamNo>>>(parameters)の
三番目のパラメータでサイズが指定されたときだけ。
・共有メモリは一回の<<<>>>の呼び出しの間だけしか有効じゃない。
# つまり、次の回には残っていない。
・共有メモリをハンドリングするには、extern __shared__ anyType * nameで宣言するだけ。
# つまり、コンパイラは型のマッチングやサイズのチェックをしないので自分で管理しないといけない。
・共有メモリはblock間で独立、block内ではthread数に関わらず共有。
# つまり、実際のデバイスにそぐわないthread数を指定した場合はCUDA側で同期処理が入るのか、遅くなってしまう。
・あるthreadが共有メモリに書いた後、別threadが読む前には__syncthreads()で同期を取らないといけない。
# ある意味当然なんだけど、その所為で遅くなるのも事実。
あー、ついでにメモリの違いを簡単に。
※グローバルメモリ
・読み書きできる。coalescedにアクセスできれば結構速い。消えない。広い。
・ホスト側スレッドごとに独立している。ホスト側から見ると、毎回同じアドレスになるのでどのくらい使えるか判りにくい。
※共有メモリ
・読み書きできる。遅くない。
・呼び出しごとに消えてしまう。余り広く取れない。事実上同期を取る必要があって使い難い。
※定数メモリ
・速い。消えない。そこそこ広い。
・例えばfloat2を読み込むインストラクションがないので実はグローバルメモリからfloat2を読むより遅くなる場合もある。
・デバイス側から書き込めない。ホスト側スレッドごとに独立している。複数スレッドからCUDAを使うと毎回転送しなおすのか?
※レジスタ
・読み書きできる。速い。厳密に型チェックされる。つーか、型ごとに違うインストラクションが使われるからptxファイルで追える。
・呼び出しごとに消える。他のメモリに較べれば狭い。
# ローカルな配列は宣言したことないからよく判らん。