cuda 2.1を試してみる

なんとなく遊んでいるサーバー用のマシンにPCI Expressが付いているので、"Quadro FX 370 Low Profile"を購入して載せてみた。
CUDAとしては中途半端でもったいない感じはしますが、そろそろメインデスクトップのマザーボードをAGP 8xなAM2NF3-VSTAから変更した時の事を考えました。 純粋にCUDAのテストに使うなら、G9x世代のミッドレンジよりハイエンドに属するボードでないと速さが実感できなくて、おもしろくないだろうなというのが使ってみた感想です。

しばらく前に出ていたCUDA 2.1 betaを入れていたのですが、今日になって2.1のリリースに気がついたのでドライバやらToolkitやらを入れ直しました。
環境はこんな感じです。

• CPU: AMD Turion 64 MT-37
• Memory: 1GB
• Graphic Card: Quadro FX 370 Low Profile
• OS: Ubuntu 8.04 LTS Server x86_64版

うーん、非力だなぁ。

さて、CUDA Zoneにある資料をみるとvecAddという関数の付いたサンプルプログラムが乗っています。これをちょっと変更しながら、いろいろ試してみました。
CUDA Zoneの学習ページにある「CUDAプログラミングモデルの概要」をみると、"cudaMalloc"と"CudaMemcpy"を使ってデータをやり取りする辺りが最低限必要な知識みたい。 サンプルを少し変更しながら、なんとか簡単なプログラムは書けるようになりました。

さて、お題は何にしようかなぁ、と考えてみて、いつも通り簡単なループを沢山回して素数をみつけるプログラムを書く事にしました。
アルゴリズムは「2より大きい自分自身以外の数で割れない事」を素直に確認します。 効率とか考えません。単純に自分自身以外の数で割れない事を計算から確認していきます。
このロジックは実用性は皆無だけれど、新しいプログラミング言語を学習する時は、適度にループと条件分岐と配列やらハッシュやらのデータ構造をなにか1つは使うし、時間もかからないので良く使います。
なにしろ結果を表示させて正しく動いているか確認が簡単ですから。 新しく勉強する言語では結果そのものが怪しいので、簡単に確認できる事が重要です。

とりあえず試した感想は、この程度の仕事だとCPUよりもちょっと早いかなぁぐらい。 単純に1関数で1つの素数判定(?そんなに大袈裟な事ではない)をさせた場合は、CPUよりも遅くなりました。 1回の呼び出しで適当な処理単位(#define WORKUNIT 256)分の数字について素数か判定させる処理を繰り返すと,ある程度はスケールするようでした。

idx = blockDim.x * blockIdx.x + threadIdx.x;
for(i=0; i<WORKUNIT; i++)
  target_number = (WORKUNIT * idx + i)

でもこの作業を割り振るためのユニットが最上位機種でFX 370の30倍ぐらいあるかと思うと、スケールする可能性は確かに感じました。
今回の例はシンプルすぎるから、メモリ空間にもっとアクセスするようなケースでないとCPUで処理をさせた時との差がわかりずらいかなと思います。 2次元の各画素に計算を適用するとか、音声信号に対するFFTみたいな操作で、ちゃんと動いているか検証も簡単な良いネタがあると良いですね。