その手の平は尻もつかめるさ

ギジュツ的な事をメーンで書く予定です

トップ > CUDA

CUDA で 2 枚以上の GPU を使って cufft する時は FFT Plan を各々の device に乗せる必要がある

cuda

以下は間違い．

cudaSetDevice(devices[0]);
cufftHandle fft_plan;
cufftPlan1d(fft_plan, SIZE, CUFFT_C2C, BATCH); // ここで作られる FFT Plan は device 0 にしか確保されない

cufftComplex *ffted_data[2], *orig_data[2];
cudaSetDevice(devices[0]);
cudaMalloc((void **)&orig_data[0], SIZE);
cudaMalloc((void **)&ffted_data[0], SIZE);
cudaSetDevice(devices[1]);
cudaMalloc((void **)&orig_data[1], SIZE);
cudaMalloc((void **)&ffted_data[1], SIZE);

cudaSetDevice(devices[0]);
cufftExecC2C(fft_plan, ffted_data[0], orig_data[0], CUFFT_FORWARD);
cudaSetDevice(devices[1]);
cufftExecC2C(fft_plan, ffted_data[1], orig_data[1], CUFFT_FORWARD); // ここ (device 1 上) で使っている fft_plan は device 0 にしか乗っていないのでマズい

この場合，fft_plan が device 0 にだけ乗っている為，device 1 で FFT を走らせようとすると変な結果が得られる．
エラーではなく変な結果が得られる！！！！
つらい！！！！！！！！！！！
原因がわかりにくい！！！！！！
あああああああああああああああああああああああああああああああああああああああああああああ！！！！！！


正しくは以下のように，FFT を実行する各々の device 上で FFT Plan を確保しなければならない．

/*
 * ここで各 device 上で FFT Plan を確保する
 */
cufftHandle fft_plan[2];
cudaSetDevice(devices[0]);
cufftPlan1d(fft_plan[0], SIZE, CUFFT_C2C, BATCH);
cudaSetDevice(devices[1]);
cufftPlan1d(fft_plan[1], SIZE, CUFFT_C2C, BATCH);

cufftComplex *ffted_data[2], *orig_data[2];
cudaSetDevice(devices[0]);
cudaMalloc((void **)&orig_data[0], SIZE);
cudaMalloc((void **)&ffted_data[0], SIZE);
cudaSetDevice(devices[1]);
cudaMalloc((void **)&orig_data[1], SIZE);
cudaMalloc((void **)&ffted_data[1], SIZE);

/*
 * 各 device に応じて使用する FFT Plan を差し替える
 */
cudaSetDevice(devices[0]);
cufftExecC2C(fft_plan[0], ffted_data[0], orig_data[0], CUFFT_FORWARD);
cudaSetDevice(devices[1]);
cufftExecC2C(fft_plan[1], ffted_data[1], orig_data[1], CUFFT_FORWARD);

こうするとちゃんと動く．ハマった……

あああああああああああああつらい！！！！！！！！！！！！！！！　勘弁してれ！！！！！！！！！！！！！

CUDAのエラーチェックを楽にする

CUDA

CUDAの組み込み関数はほとんどcudaError_tというエラーコードのenumの値を返してくるので，これを適宜見てエラー処理をする必要があります．最近ではそこまででもないですが，CUDAの組み込み関数はカジュアルにエラーを吐くのでちゃんと見てやらなければならない．
見てやると言っても，CUDAの組み込み関数がエラーを吐いてきたらほとんどの場合そのまま潔く死んで欲しいので (個人の意見)，以下の様に戻り値をチェックしてexit()するというコードを大量に書くことになります．

int *foo;
cudaError_t err;
err = cudaMalloc((void **)&foo, sizeof(int) * 42);
if (err != cudaSuccess) {
    exit(err);
}
...

まあ，素直に書けばええやん！　という意見は尤もなのですが，いかんせんこういうエラーチェックが大量に入ると本質的なコードが追いにくくなってつらい．
あとカジュアルにエラーチェックを忘れてしまって，バグッた時にハマる (しかしまあ，これはもうどうしようもない気がしている)．golangとかだと，返ってきたerrが使われてないと「err使われてないよ！ (つまりエラー処理してないんじゃねえの！)」と処理系が叱ってくれるからナイスですね．羨ましい．あとgolangは「エラーを受け取っても受け取らなくても良いよ」みたいなものが無いから，そこら辺のコードを規約化出来て良いという印象があります．


さて，そうしたリッチな機構が無い上でどうするかというと，僕は以下の様なマクロを書いてしのいでいます．

でもって，こう書く

int *foo;
CUDA_SAFE_CALL( cudaMalloc((void **)&foo, sizeof(int) * 42) );

こうしておくと，CUDA組み込み関数からエラーが吐かれた時にエラーメッセージ・ファイル名・行番号と共に，エラーコードでexit()してくれるのでまあまあ便利．
ちなみにcudaGetErrorString()っていう関数はcudaError_tの値をもとにエラーメッセージを引っ張ってきてくれる君です．


しかし，これも野蛮な方法であることには違いないし (なにより，CUDAの組み込み関数がCUDA_SAFE_CALLに渡される事が保証されない)，エラーチェックが忘れがちになってしまうという問題については何一つ解決されていません．
現段階では，ちょっとだけコードが見やすくなって便利，くらいの感覚です．


もっと良い方法知ってる人いたら是非教えて下さい．

CUDAで特定の条件に合致したGPUのIDを持ってきたい

CUDA

CUDAで，特にマルチGPUプログラミングなどをやっておりますと，特定の条件に合致したGPUのIDを持ってきたいという要求に高確率でぶち当たる事となると存じます．俺はぶち当たる．


GPUがマザーボードに5枚刺さっていて，そのうち4枚は映像のアウトプット端子がないGPGPU専用の板で，残る1枚は映像を出力するためだけの貧弱なボード，という構成のサーバはこの世の中少なくありません *1．
そうした構成の時に，不慮の事故によりcudaSetDevice()によってGPGPU専用のグラボではなく映像出力専用の貧弱な板がアサインされてしまったばかりにパフォーマンスが死ぬほど落ちて死ぬ，というのは割とよくある事例であります．そうした事故は未然に防がなくてはなりません．


原因を考えてみましょう．
なぜこういう悲しい事故が起こるかと言うと，cudaSetDevice()に食べさせるGPU IDをハードコードしているからこういう事が起こるわけです．
GPUのIDはハードウェアの構成が変わるとそれと共に変化します (あと良くわからんけどこの前マシンをリブートしたらIDが変わってハマった，ふざけんな)．
つまり，書いたプログラムを違う環境に持って行くと動かなくなったりパフォーマンスが下がる可能性が出てくるわけですね．
GPU IDをハードコードしてはならない．


で，どうするかと言うと，

こういう感じのプログラムを書くことで解決を試みています．
cudaGetDeviceCount()でGPUボードの数を持ってきて，それをもとに全部のGPUボードを舐めつつcudaGetDeviceProperties()でGPUのプロパティを引き出し，そのプロパティのデータを使って条件と一致するボードかどうか (今回はボードの名前) を見てやるという感じですね．
基本的にcudaGetDeviceProperties()で取れるcudaDeviceProp構造体はボードのほとんどの情報を持ってるので，これを使えば大体なんとかなります．


こういう感じで持ってきたIDをcudaSetDevice()に渡してやると事故が起きなくて便利．

追記

つらい

*1:とは言え，「GPUなのに映像出力できないとはけしからん！！」と怒る怖い人がいるので，最近ではGPGPU専用の板でも映像を出力できたりします

nvcc の最適化オプションについて

CUDA

nvcc の最適化オプションについて、不正確な情報をしばしば耳にするのでそのことについて書いておきます。

nvcc の最適化オプションとは

$ nvcc foo.cu -O2
の'-O2' のような、'O' というスタイルのオプションです。
gcc とかで見覚えがある事と思います。

nvcc の最適化オプションが最適化するのは

ホストコードのみです。
カーネルコードは最適化しません。

実際に確認してみる

まず、object ファイルを出力して差分を取ってみましょう。
object ファイルを出力するにはgcc 等のコンパイラと同様に'-c' オプションを付加します。
$ nvcc foo.cu -c
最適化オプションを付加したものについてもobject ファイルを出力しましょう。
$ nvcc foo.cu -c -O2
これら2つのobject ファイルについて差分を取ります。
object ファイルはバイナリファイルなので普通にdiff を取ると何の事やらさっぱり分かりませんが、
バイナリエディタ等で差分を取ってみると全くの別モノである事が分かるので、最適化されている事が分かります。*1

次に、ptx ファイル (カーネルコードをGPU アセンブリ言語にしたもの) を出力してdiff を取ってみましょう。
ptx ファイルを生成するには以下のように'--ptx' オプションを付加します。
$ nvcc --ptx foo.cu
最適化オプションをくっつけたものについてもptx ファイルを出力します。
$ nvcc --ptx foo.cu -O2
この2つのdiff を取ります。*2
恐らくここで得られる差分はTemp ファイルの違いくらいで、処理に関わる部分については差が無いことが読み取れます。
つまり、最適化されていない事がわかります。

結論

nvcc の最適化オプションによって最適化されるのはホストコードのみです。
カーネルコードは最適化されませんので！！

*1:暴論ですが……

*2:ptx ファイルはテキストなので普通にdiff で良いです

CUDA 初学者は「効率的なプログラミング」を一旦忘れた方が良いと思った話

CUDA

ある事情でCUDA 初学者が集まる勉強会で (恐れ多くも) TA をやってきて、その時に色々と思うところがあったので書きます。

「CUDA を始める動機」が与える影響

CUDA (GPGPU プログラミング)を始めるにあたっての動機の大半は「処理を高速化したいから」だと思います。
中には「GPGPU を使ったプログラミングに興味があったから」とか「CUDA が面白そうだったから」とか、
そういう知的好奇心からCUDA を始める強者もいるとは思いますが、
「GPU 使えば処理にかかる時間を短縮できるんでしょ？」という即物的な動機からCUDA の学習や導入を検討する例が多いような印象を受けます。

で、その世相を反映してか、CUDA 初学者向けの書籍や記事には「高速に動作する処理の書き方」だとか「効率的なプログラムの書き方」という
発展的な内容にカジュアルに触れているものが多く、今回参加した初学者向け勉強会でもそういった内容がガンガン出ていたので、
「これって初学者にとってハードルが高いのでは？」と感じました。

とりあえず「動いたわーい」を経験した方が良いのでは

始めてプログラムを学んだ時、あるいは始めて触るプログラミング言語を学習した時のことを思い出して下さい。
果たして、その時に「効率的な書き方」に気を配ったりしたでしょうか？*1

例えば、
my $str = 'something';
my $length = length($str);
for (my $i = 0; $i < $length; $i++) {
    # Do Something
}
と
my $str = 'something';
for (my $i = 0; $i < length($str); $i++) {
    # Do Something
}
とでは、どちらが効率的か？　という事を、最初の内は意識せずに書いていたのでは無いでしょうか。*2
で、「効率を意識」だとか「高速にチューン」だとかは、ある程度その言語でプログラムが書けるようになってからの次のステップというか、アドバンスな内容である訳です。

しかし、CUDA の入門書籍や記事はそのアドバンスな内容が先行している為に、来るものを拒んでいる感じが否めません。
しかもCUDA のチューニングといえば……
  • コアレッシング
  • バンクコンフリクトの解消
  • ワープダイバージェントの回避
  • Block の分割
  • etc...
と、初学者には難しい内容が多いので、最初からこれらを意識してプログラムを書こうとすると心がグシャッと折れる可能性が高まります。
「効率的なプログラミング？　あーなんか難しいな……やめるか……」みたいな感じになりかねません。
なので、とりあえず最初は難しい内容 (チューニングとか) をガン無視して「とりあえず動くもの」を書いてみて
「わーい、動いた！」を経験した方がモチベーションを殺さずに学習を続けられて良いのでは無いかと思います。

(最近は) 普通に書いても早くなる場合が多い

つーか最近は、GPU のアーキテクチャやCUDA 処理系が賢くなっていて色々と世話を焼いてくれるので、*3
(CUDA の作法に則っていれば) 効率化を特別意識してプログラムを書かなくても、CPU よりも高速に動作したりします。
限界までギンギンにチューンする必要がある場合はさておき、「現状 (CPU ベース) よりも高速に動けば嬉しいなー」という程度の要求であれば
ぐだぐだ難しい事を考えずにプログラムを書いても大丈夫だと思います。

GPGPUプログラミングは独自のプログラミングパラダイムである

CUDA の入門記事等を読むと「CUDA はC言語ベースの言語だから、C言語の知識さえあれば簡単にGPGPU プログラミングが出来る」
という記述をよく見かけますが、その認識は捨て去った方が良いと思います。
「GPGPU プログラミング」は、「オブジェクト指向プログラミング」や「関数プログラミング」や「論理プログラミング」のように独立したプログラミングパラダイムだと感じています。

C言語のように、繰り返し処理 (for, while等) と分岐処理 (if, case等)を組み合わせて任意の結果を得る「手続き型プログラミング」と、
繰り返し処理と分岐処理を (なるべく) 排除して、莫大な数のThread を同時並列で実行させる「GPGPU プログラミング」とでは
全く異なるプログラミングパラダイムであると言えるでしょう。そんな具合にCUDA とC言語は異なるパラダイムなのにも関わらず、
「CUDA とC言語はほぼ同じ言語だ」という不要な前知識が存在する所為で理解や学習を妨げる可能性があります。
(例えば、C言語で効率の良い書き方が必ずしもCUDA で効率の良い書き方であるとは限らないのです)

WEB+DB PRESS Vol.67 の「入門　関数プログラミング」において、筆者の山本和彦さんは
「関数プログラミングを習得するには、これまで命令プログラミングで培った技術は一旦忘れ、
真っ白な気持ちで臨む必要があります。関数型の山を登るためには、命令型の山を降りなければなりません」
と書かれています。GPGPU プログラミングも同様だと感じています。GPGPU プログラミングをマスターする為には、命令型の山を一度降りねばならないと思います。
なので、CUDA は「構文こそC言語と似ているものの、中身は全く別の言語」という認識の下で学習した方が身に付くのでは無いでしょうか。

まとめ

  • 最初のうちはCUDA で効率的なプログラムを意識しないで書いた方が良い
    • とりあえず動くものを作ってみよう
  • というか、最近は特別効率を意識してプログラムを書かなくてもある程度高速に処理される
  • GPGPU プログラミングは独自のプログラミングパラダイムである
    • C言語の知識がCUDA の高速化・効率化に適用できるとは限らない
という感じでしょうか。


ここから先は

特に効率化とかそこら辺は関係無い話です。
ただ書きたいから書いているだけです。

(補足1) 今出ている入門書はモダンではない

2012.10.30 現在、出版されているCUDA の入門書籍は古い内容のものが多く、
モダンなアーキテクチャや書式に対応していなかったりしているので注意が必要です。*4
例えば、「はじめてのCUDAプログラミング」はCUDA の入門には良い書籍だと思いますが*5、
ご多分に漏れず内容が若干古いので適宜インターネット等で情報収集を行った方が良いと思います。
はじめてのCUDAプログラミング―驚異の開発環境[GPU+CUDA]を使いこなす! (I・O BOOKS)

はじめてのCUDAプログラミング―驚異の開発環境[GPU+CUDA]を使いこなす! (I・O BOOKS)


(補足2) "GPU = 高速" とは限らない

「GPU が得意な処理かどうか」というのは見極めなければなりません。
GPU が得意な処理 (分岐させる必要が無い若しくは少ない、等)であれば、GPU で実装したときにCPU よりも高速に処理されると思いますが、
逆に不得意な処理をGPU に行わせようとするとCPU よりも低速になる可能性があります。
(ワープダイバージェントや、スレッド間の同期や、ホスト・デバイス間メモリ転送等でオーバヘッドが生じる為)
なので、GPU は必ずしも高速であるとは限らないという事を頭の隅に留めておいた方が良いと思います。
(「GPU にすれば何でも早くなるんでしょ？」みたいな事を良く訊かれたので)

*1:「おうともよ！」というツワモノはオッケイです。最初から高速化を意識したCUDA プログラミングをなさった方が幸せになれると思います

*2:とは言うものの、この例の$str の文字列長では有意な差は出ないと思いますが

*3:自動的にメモリキャッシュが効いてくれたり、メモリのアライメントをよしなにやってくれたり……

*4:特にFermi 以降のアーキテクチャは色々と劇的に変わっているので

*5:ただ、割と応用的な内容にも食い込んでいますが