良いもの。悪いもの。

以前、 複数のGPUでマンデルブロ集合を並列計算 というブログ記事を書いた。このときはTeslaが1枚で7.45秒、Teslaが2枚で4.26秒で計算できたと報告したが、マルチGPU化は多少面倒だったし、ある程度のオーバーヘッドが掛かっていた。 しかし、 CUDA 4.0 がリリースされてから状況が一変した。自由にcudaSetDeviceが呼び出せるようになり、またUVA (Unified Virtual Addressing)を利用することにより、それはそれは簡単にマルチGPU化や、ホスト・複数GPU間での一元的なメモリの参照を実現できることになった。 というわけで、CUDA 4.0を使ったマンデルブロ集合のプログラムを以下に示す。シングルGPUからマルチGPUへ変更した部分については赤字で示してある。ただし、簡便のために、GPUデバイスは2枚で決め打ちとし、マンデルブロ集合の幅は2の倍数のピクセルとなるようにしてある。また、UVAは使わなくても良かったのだが、使い方を示すために利用している。 CUDA 4.0によるマルチGPU化によりマンデルブロ集合計算の実行時間(計算条件は 前回 と同じ)は 7.45秒から3.73秒 となった。ほぼ2倍の速度である。シングルGPUからの修正箇所はほんの僅か。カーネル関数に至っては一文字たりとも変更していない。良い時代になったものだ。 mandelbrot_multigpu.cu // madelbrot using multi GPU devices by nox, 2011.06.06 // nvcc -lcutil_x86_64 -arch sm_13 -use_fast_math -prec-sqrt=false -keep -L ~/NVIDIA_GPU_Computing_SDK/C/lib -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc -g mandelbrot_multigpu.cu -o mandelbrot_multigpu #include <iostream> #include <fstream> #include "cutil_inline.h" using namespace std; ...

範囲(-0.005, -0.005, 0.005, 0.005)のマンデルブロ集合を描画サイズ4,800x4,800、繰り返しの上限10,000として計算させてみた。C++でコードを書いてXeon X5650 1コアで走らせたところ942.97秒かかった(SIMD最適化はしていない)。TBBを利用した12論理コアの並列計算では88.68秒だった。次いでCUDAを使ってTeslaで走らせてみたが7.45秒まで短縮された。Teslaパない。そして贅沢にも2枚のTeslaをスレッドで並列化して使ってみたら4.26秒で計算できた。実にXeon 1コアの220倍の速度が出たわけだ。カリカリチューンをしなくてもこの程度の高速化ができるということが重要だ。 因みにマンデルブロ集合の上記の範囲は数値が発散しない集合部分であり、描画させてみても真っ黒なので面白くはない。これは並列化した際に計算が偏らないようにしたかったのでこのような範囲を指定している。プログラムで出力させたファイルは色の生データなのでそのままでは表示できない。一応、画像表示するためのPythonコードも付け加えておく。 以下に、GPU2枚を利用してマンデルブロ集合を計算するCUDAコードを示しておく。 mandelbrot_thread.cu // Mandelbrot set using GPGPU by nox, 2011.02.12 // nvcc -lcutil_x86_64 -arch sm_13 -use_fast_math -prec-sqrt=false -keep -L ~/NVIDIA_GPU_Computing_SDK/C/lib -I ~/NVIDIA_GPU_Computing_SDK/C/common/inc -g mandelbrot_thread.cu -o mandelbrot_thread #include <iostream> #include <fstream> #include <cutil_inline.h> #include <multithreading.h> using namespace std; const int BLOCK_SIZE_X = 16; const int BLOCK_SIZ...

以前に、 はじめてのCUDAプログラミングで分子動力学計算 というブログ記事を書いたことがある。最近このなんちゃって分子動力学計算(MD)プログラムのソースコードを読み直してみたのだが、かなりひどい。しかも、「CUDAプログラミング」でググってみると、この記事が2番目にくる。こんないい加減なコードを参考にされたら読んだ人にも迷惑が掛かるので修正することにした。それほどCUDAに慣れているわけではないが、前回のコードよりはましだと思う。 それにしてもコードの内容がひどい。意味もなく__syncthreads()が入っているし、複数のスレッドから同じグローバルメモリに書き込みしてるし、レジスタを活用してないし、計算の順序は非効率だし、ダメダメだ。 そこでまず、CUDA Visual Profilerで関数のパフォーマンス測定を行った。因みにこのプロファイラは CUDAプログラミングツールキット に含まれているもので、CUDAでプログラミングを行うには必須だと思う。使い方は簡単で、Windowsであればcudaprof.exeを実行して、FileメニューからNew...を選び、プロジェクトの名前と場所を設定して、実行するCUDAプログラムを指定するだけだ。引数が必要ならそれも指定しておく。デフォルトでは4回実行され解析される。それぞれの関数がどれだけ時間がかかったのかプロットされるので、どこがボトルネックになっているのか一目瞭然だ。 まあ、予想していた通り、Calc関数が全体の99.8%を占めていたので、ここから修正を行った。まず、iとjを使っていたループをjだけにした。iはスレッド数で分割されていたが、ブロックと合わせてそれを消した。次に、jのループ内で不必要にグローバルメモリにアクセスしないようにした。例えばchg[i]というグローバル変数はループ外でローカル変数に入れてそれを使うようにする。あとは、できるだけ除算などの演算を減らすようにした。 次に、系全体を中心に戻す関数があるのだが、中心座標を求める部分はホスト側の関数で実装し、系全体を戻す関数のみをCUDAの関数とした。もともとこの処理は毎回やらなくてもよいものなので、指定したステップ毎に1回行うようにした。今回は100ステップに1回としている。なので、全体の処理時間から見れば無視しても良く、頑張...

追記 (2009/12/12): ここで示してるコードはかなりひどいので少し手直しして、 CUDAで作成した分子動力学計算プログラムを書き直してみた という記事にまとめた。こちらも参照して欲しい。 最近、自分の周りでは CUDA を使ったGPGPUが流行っているので、どんなものかこの週末にいじってみた。 一日目はNVIDIA GeForce 8800 GTSを搭載したWindows XPマシンに CUDAをインストール をしてリファレンスやサンプルプログラムをパラパラと読んでみた。二日目になんちゃって分子動力学(MD)計算プログラムを書いた。擬似原子を扱っており、全原子のvdW力とクーロン力のみを考慮して分子内結合については考慮しない、液滴中のMD計算だ。取り敢えず、半径20Åの液滴中に1229原子を入れて1000ステップ計算させてみた。C++で書いたプログラムでは4分30秒かかった計算がCUDAでは32秒で計算できた。約9倍ほどの速度が出ているようだ。 まあ、C++のコードも併せて一日で書いたので最適化や高速化については考慮していない。後でスレッドでの分割がやりやすいかもと思ってiとjをそのまま全部計算しているし。気が向いたらちゃんとしたプログラムを書いてみるかも。 CUDAアーキテクチャについてはまだまだ理解が足りない。効率のよいメモリの扱い方を覚えないとなぁ。あと、sharedメモリをうまく扱えてないっぽい。さすがに一日二日じゃキツイか。 初心者が一日で書いたコードなので参考になるかわからないが、以下にソースコードとサンプルの入力データを示しておく。入力データについては、上記で用いた1229原子ではブログ記事するには大きすぎるので154原子の小さなデータを用意した。また、記事の表示量が多くなりすぎるのでC++のソースは割愛させてもらった。Windows上でのコンパイルは下記の通り。 nvcc -lcutil32 -lkernel32 --host-compilation c++ -Xcompiler /EHsc,/W3,/nologo,/O2,/Zi,/MT simple_md_gpu.cu -o simple_md_gpu.exe 入力ファイルをmd.datとして、100ステップ計算させる場合は、以下のように実行する。 si...