あかすくぱるふぇ

同人サークル「あかすくぱるふぇ」のブログです。

CUDA

CUDA samplesのbilateralFilterの中身を解析してみました。

・メモリ領域の確保とデバイスへのコピー(.cuの118行目)
二次元配列を連続領域として確保したいのでcudaMallocPitch()とcudaMemcpy2D()を用いている。
http://www.slis.tsukuba.ac.jp/~fujisawa.makoto.fu/cgi-bin/wiki/index.php?%A5%EA%A5%CB%A5%A2%A5%E1%A5%E2%A5%EA%A4%C8CUDA%C7%DB%CE%F3

・テクスチャメモリ
テクスチャメモリの生成(.cuの17行目)
テクスチャメモリと画像の関連付け(cudaBindTexture2D(), .cuの186行目)
テクスチャメモリの参照(tex2D(), .cuの99行目)
http://gpu.fixstars.com/index.php/%E3%83%86%E3%82%AF%E3%82%B9%E3%83%81%E3%83%A3%E3%83%A6%E3%83%8B%E3%83%83%E3%83%88%E3%82%92%E4%BD%BF%E3%81%86

・PBOをCUDAからいじる
PBOをCUDAへ登録(cudaGraphicsGLRegisterBuffer(), .cppの401行目)
PBOのポインタを取得(cudaGraphicsResourceGetMappedPointer(), .cppの167行目)
PBOの値をいじる(.cuの114行目)
http://shouyu.hatenablog.com/entry/2011/12/05/192410

・PBOの描画(.cppの178行目)
PBOをテクスチャに設定して、テクスチャマッピング。

工学社の"はじめてのCUDAプログラミング"のまとめ。
内容が古かったり、分からなかったりしたところを補完します。

・GPUのアーキテクチャ(p.26~28)
演算性能[Flops] = SP数 * SPクロック * 2(積和演算)
メモリバンド幅[bit/sec] = メモリインターフェース[bit] * メモリクロック

・スレッドの階層構造(p.44~46)
1つのGPUで1グリッド。
グリッドは複数のブロックで構成される。1SMで1ブロック。ブロック数はGPUごとに異なる。
ブロックは複数のスレッドで構成される。1SP(CUDAコア)で1スレッド。スレッド数は世代ごとに異なる。

・高速アクセスのコツ
ウォープ:連続するスレッドでは同じ分岐を通るようにする。
コアレッシング:アライメントずれはSharedMemoryを使って解決する。
バンクコンフリクト:SharedMemoryの同じバンクに同時アクセスしない。

CUDAプログラムの最小構成とそれをVIsualStudioでコンパイルする際の注意点を示します。

ソースコードは以下の通り。ファイル名は仮にSource.cuとします。
以下のページのソースコードを一部変更したものです。
http://seesaawiki.jp/w/mikk_ni3_92/d/CUDA%ca%d402

無題
注意点は以下の通り。

・ビルドカスタマイズ
コンソールアプリケーションテンプレートから作る場合はビルドカスタマイズが必要です。
ソリューションエクスプローラーでプロジェクトを右クリック > ビルド依存関係 > ビルドのカスタマイズをクリック。
CUDA(.targets, .props)にチェックを入れます。

・ビルドツールの選択
ソリューションエクスプローラーでSource.cuを右クリック > プロパティをクリック。
構成プロパティ > 全般 > 項目の種類をCUDA C/C++に変更します。

・cutil.h問題
古いCUDAでよく使われていたcutil.hが新しいCUDAには入っていません。
CUDA SDKをmakeするとその中に入ってるとか、CudaSamplesに入っているhelper_cuda.hを代わりに使うとか色々方法はあるようですが、どうもスマートじゃない。
Toolkitに入っているcuda_runtime.hやdevice_launch_parameters.hで代用するのが正攻法かと思います。
http://saito-175.hatenablog.com/entry/2013/07/25/191519

・コンパイルできるけどエラーの赤線が出る問題
__syncthreads()などに赤線が引かれてしまう問題です。
__CUDACC__マクロをdefineすると赤線が出なくなるようですが、インテリセンスが効かなくなるので辛い。
今のところ根本的な解決策が見つかっていません。
http://www.momo86.net/amahodo/contents/cuda/entry002

↑このページのトップヘ