fc2ブログ

CUDA by Example 汎用GPUプログラミング入門

CUDA by Example 汎用GPUプログラミング入門を読んだので忘れないようにまとめておきます。一度本を読んだ状態で、復習したいときに読むように書いてます。2回目を読まずに済むといいなあ。



本のサポートページ
https://developer.nvidia.com/content/cuda-example-introduction-general-purpose-gpu-programming-0

これはCUDA CでのGPGPUプログラミングについて初歩から書かれている本です。
訳本ですが、とても読みやすく書かれていてわかりやすいです。初心者にお勧め。
CUDAを使って画像処理の計算時間を短くしたいと思って読んだので、その視点でまとめました。もちろん、それ以外にも有用な内容がたくさんあります。

ここからはこの本で述べられている内容について簡単に列挙していきます。


・GPGPUプログラミングとは
グラフィックス以外の汎用的な計算にGPUを使って超並列に計算することで計算時間を減らすというもの

・CUDA Cとは
NVIDIAのGPUを使ってGPGPUプログラミングを実現するための言語


・お勧めのドキュメント
NVIDIA CUDA Programming Guide
NVIDIA CUDA Best Practices Guide

・サンプルコード
下記のHPのDownload source code for the book's examples (.zip)から取得できます。
https://developer.nvidia.com/content/cuda-example-introduction-general-purpose-gpu-programming-0

・開発に必要なもの
CUDA対応のGPU
NVIDIAデバイスドライバ
CUDA開発ツールキット
標準Cコンパイラ


ここからプログラミングの内容です。
・デバイスの確認
cudaGetDeviceProperties()を使うと、GPUの情報を取得することができます。

・ブロック・スレッドの制限
ブロック:1次元につき65535ブロック
スレッド:512スレッド
要素数がブロックやスレッドの制限を超えた場合は以下のようにして超えた要素へ繰り返し処理を適用します。
while( id < N )
{
// なんか処理する
id += blocDim.x * threadDim.x;
}

・共有メモリ
同じブロックのスレッド間で共有されるメモリ。
変数宣言時に__shared__で修飾するとその変数を共有メモリに配置できる。

・縮約
配列の総和を計算する例で説明されている。
すべてのスレッドで一度に同じ変数に対して足し算をしようとするとアクセスが集中してしまい遅くなってしまう。そこで、配列の要素を2つ足すという計算を繰り返してアクセスを分散する方法。

・コンスタントメモリ
カーネル実行中に変化しないデータを使うメモリ。
メモリサイズは64KB。
コンスタントメモリからの同じデータの読み取りではメモリ帯域幅を節約できる。
ただし、スレッドがどれも異なるアドレスを読み取るときは遅くなってしまうことに注意。
コンスタントメモリの領域は静的に割り当てられなければならない(コンパイル時にサイズが確定されていなければならない)ため、cudaMalloc()やcudaFree()は必要ない。

・イベント
GPUのタイムスタンプ。
以下のようにGPU実行時間の計測に利用可能。
cudaEvent_t start, stop;
cudaEventCreate( &start );
cudaEventCreate( &stop );
cudaEventRecord( start, 0 );
//なんか処理
cudaEventRecord( stop, 0 );
cudaEventSynchronize();

・テクスチャメモリ
読み取り専用のメモリで1次元テクスチャメモリと2次元テクスチャメモリがある。
読み取りアクセスが局所的な場合にメモリトラフィックを削減できる。
画像処理なんかでは空間局所性を利用することが多いのでうまく利用すると速くなりそう。
以下のように使う・
texture< float > tex1;
texture< float, 2 > tex2;
cudaBindTexture(...);

// 1次元アクセス
tex1Dfetch(tex1, x)
// 2次元アクセス
tex2D( tex2, x, y )
cudaUnbindTexture( tex );

・OpenGLとCUDAの相互運用
glutを利用したOpenGLの例が載っている。
OpenGLのバッファ名GLuintとCUDA Cのバッファ名 cudaGraphicsResource*を対応付けることで、同じデータを共有することができ、CUDAで処理したデータを転送なしにそのまま表示することができる。

・アトミック
各スレッドで同じ変数に演算をしようとすると読み取りから書き込みの間に別のスレッドに書き換えられてしまうという問題が発生する。
そこで、読み取りから書き込みまでの一連の処理の間に別のスレッドから処理できないようにする必要がある。
atomicAdd()などを使う。

・ページロックホストメモリ
ページがロックされたメモリ。ピンメモリとも。
物理メモリに常駐される。
利用例
cudaHostAlloc( ... , cudaHostAllocDefault )

・ストリーム
GPU上のタスク。
これを2つ以上使えばGPU上で並列に別々のタスクを処理することができる。
同じ処理を複数のストリームに分割して処理することで、メモリの転送とカーネル実行を並列に行うことができるため計算時間を短縮することが可能。

・マルチGPUプログラミング
CPUをマルチスレッドにして各スレッドで別々のGPUを設定することで複数のGPUを利用することができる。
GPUのセットはcudaSetDevice()でできる。

・CUDAツール
CUFFT: 高速フーリエ変換ライブラリ
CUBLAS: 線形代数ライブラリBLASのGUDA版
NVIDIA GPU Computing SDK: CUDAを使ったいろんなサンプルがある。画像・映像処理もある
NPP: 画像・映像処理などの関数ライブラリ

・参考になるコード
CULATOOLS: 線形代数ライブラリLAPACKのCUDA実装。LU分解、QR分解、線形ソルバ、特異値分解、最小二乗ソルバ、制約付き最小二乗ソルバなど。無償のBasicパッケージ、優勝のPremiumライセンス、Commercialライセンスがある。

・デバッグ
NVIDIA Parallel Nsight: Visual Studio用デバッガ
CUDA Visual Profiler: パフォーマンスをあげたいときに

・CUDA公式サイト
http://www.nvidia.co.jp/object/cuda-jp.html

・GPU勉強時に見つけた参考資料
GPU による高速画像処理(名古屋大学大学院情報科学研究科),出口 大輔,井手 一郎,村瀬 洋
http://www.murase.nuie.nagoya-u.ac.jp/~ide/res/paper/J09-talk-ddeguchi-2.pdf">http://www.murase.nuie.nagoya-u.ac.jp/~ide/res/paper/J09-talk-ddeguchi-2.pdf
http://www.murase.nuie.nagoya-u.ac.jp/publications/543-pdf.pdf

これからの並列計算のためのGPGPU連載講座
http://www.cc.u-tokyo.ac.jp/support/press/news/

理研「CUDAプログラミング入門」
http://accc.riken.jp/hpc/training/
http://accc.riken.jp/secure/4467/cuda-programming_main.pdf

・GPUの本
GPUプログラミング入門 CUDA5による実装
スポンサーサイト



スポンサーリンク

テーマ:プログラミング - ジャンル:コンピュータ

コメントの投稿

非公開コメント