Twitter

最近のトラックバック

無料ブログはココログ

卒研

2012年7月21日 (土)

CUDAでポインター配列など

以下の内容は自分用のメモとして作成しました.
間違った内容があったらごめんなさい.

この夏休みはかなり忙しくなりそうです.
とにかく研究しないと,大変なことに・・・
夏休みだっけ?

CUDA上でのポインター配列



CUDA上で「デバイスメモリーのポインター配列」を使えるかどうか気になり試してみました.
無理に使う必要はない機能ではありますが,ブロックIDやスレッドIDで扱う配列を変えることが出来るのでプログラミングしやすく,かつ効率の良いデバイスコードを書ける気がします.

一応以前からポインター配列の確保自体はエラーにならなかったのですが,正しいポインターをどのようにセットしたら良いのか,セットしたところで動作するのかを試してみました.

なお,試した環境は
Ubuntu 11.04 64bit
GTX580 CC2.0
GTX260 CC1.3

Mac OSX 10.6.8
MacBookAir
GeForce 320M
です.
開発環境はCUDA5.0です.

まず,デバイスメモリーであるポインター配列に正しいデバイスメモリーのポインターを書き込む方法は次の二通り考えられます.

  • cudaMallocで得られたポインターをホスト上のポインター配列に書き込み,cudaMemcpyする方法
  • ポインター配列に値を書き込むデバイスコードを作成し,デバイスコードにポインター配列へのポインターと配列の先頭要素へのポインターを渡して,デバイス上で書き込んでもらう方法

Stackoverflowなどでは,最初の方法では,ホスト上でのデバイスメモリーのポインターの値がデバイス上での正確なポインター値である保証はないのでうまくいかないのではないか,というような記述が見られます.
確かに,ホスト上でただのポインターとして扱われるのは非常に違和感がありますし,デバイス上でのリアルなポインターではなく,配列番号のようなものである可能性は充分あります.

もしかするとマニュアルのどこかに記述があるのかもしれません.
しかし,試してみるのも面白いと思い,とりあえず試してみました.

実験に使ったのは以下のコードです.
http://ideone.com/rShxo

配列を二つ作り,その配列へのポインターをもつ要素数2のポインター配列を作ります.

ひとつの配列にはランダムな値を入れて初期化し,デバイスコードでその値に定数を加算し,二つ目の配列へ書き込みます.
このとき,配列へのアクセスはポインター配列を用いて行います.

配列の設定方法はどちらも試すことが出来ます.

コンパイル

まず,コンパイルに関しては,-arch=sm_10を指定すると以下の警告が出ます

./main.cu(23): Warning: Cannot tell what pointer points to, assuming global memory space

「ポインターがどこを指してると言うことが出来ない.グローバル空間を仮定」と直訳できますが,要するにグローバル領域を指してるとは限らないから気をつけなさい,という意味で解釈してよいのでしょうか?

-arch=sm_20では何の警告もなく通ります ( ただし当然CC2.0以上のデバイスでなければ動作しません )
恐らく,UVAがサポートされたので,グローバルメモリーを指すとは限らなくなったからでしょう.

実行

実際に実行してみると,試した環境の全てで,適切な-archを設定すれば問題なく動作しました.
さらに,先ほど示した,ポインターを書き込む方法のどちらでも問題なく動作しました.

-arch=sm_10を指定すると,どのデバイスでも問題なく動作するので,CUDAの初期の段階からこの機能はサポートされていたのかもしれませんね.

GeForce 320Mの場合


nvcc -arch=sm_10 -use_fast_math -m64 -O2 --compiler-options -fpermissive -c main.cu
cc1: warning: command line option "-fpermissive" is valid for C++/ObjC++ but not for C
cc1: warning: command line option "-fpermissive" is valid for C++/ObjC++ but not for C
./main.cu(23): Warning: Cannot tell what pointer points to, assuming global memory space
./main.cu(23): Warning: Cannot tell what pointer points to, assuming global memory space
g++ -m64 -O2 -o main main.o -L/usr/local/cuda/lib -lcudart

$ ./main
device > 0
set pointer ( 0 : from host, 1 : from device ) > 0
dev : GeForce 320M

set pointer from host
i:0.676658 o:7.676658
i:0.588739 o:7.588739
i:0.942956 o:7.942956
i:0.263872 o:7.263873
i:0.904508 o:7.904508
i:0.060301 o:7.060301
i:0.475828 o:7.475828
i:0.243280 o:7.243279
i:0.798542 o:7.798542
i:0.089129 o:7.089129
$ ./main
device > 0
set pointer ( 0 : from host, 1 : from device ) > 1
dev : GeForce 320M

set pointer from device
i:0.676681 o:7.676682
i:0.983353 o:7.983353
i:0.209772 o:7.209772
i:0.639823 o:7.639823
i:0.502809 o:7.502810
i:0.717178 o:7.717178
i:0.616962 o:7.616962
i:0.279874 o:7.279874
i:0.836431 o:7.836431
i:0.893208 o:7.893208

GTX580と260の場合 ( 加算値は3 )

nvcc -arch=sm_10 -use_fast_math -m64 -O2 --compiler-options "-Wall" -c main.cu
./main.cu(18): Warning: Cannot tell what pointer points to, assuming global memory space
./main.cu(18): Warning: Cannot tell what pointer points to, assuming global memory space
g++ -O2 -o main main.o -L/usr/local/cuda/lib64 -lcudart -lopencv_core -lopencv_highgui -lopencv_video -lopencv_objdetect -lopencv_ml -lopencv_imgproc

$ ./main
dev : GeForce GTX 260

i:0.840188 o:3.840188
i:0.394383 o:3.394383
i:0.783099 o:3.783099
i:0.798440 o:3.798440
i:0.911647 o:3.911647
i:0.197551 o:3.197551
i:0.335223 o:3.335223
i:0.768230 o:3.768229
i:0.277775 o:3.277775
i:0.553970 o:3.553970

nvcc -arch=sm_10 -use_fast_math -m64 -O2 --compiler-options "-Wall" -c main.cu
./main.cu(18): Warning: Cannot tell what pointer points to, assuming global memory space
./main.cu(18): Warning: Cannot tell what pointer points to, assuming global memory space
g++ -O2 -o main main.o -L/usr/local/cuda/lib64 -lcudart -lopencv_core -lopencv_highgui -lopencv_video -lopencv_objdetect -lopencv_ml -lopencv_imgproc

$ ./main
dev : GeForce GTX 580

i:0.840188 o:3.840188
i:0.394383 o:3.394383
i:0.783099 o:3.783099
i:0.798440 o:3.798440
i:0.911647 o:3.911647
i:0.197551 o:3.197551
i:0.335223 o:3.335223
i:0.768230 o:3.768229
i:0.277775 o:3.277775
i:0.553970 o:3.553970

nvcc -arch=sm_20 -use_fast_math -m64 -O2 --compiler-options "-Wall" -c main.cu
g++ -O2 -o main main.o -L/usr/local/cuda/lib64 -lcudart -lopencv_core -lopencv_highgui -lopencv_video -lopencv_objdetect -lopencv_ml -lopencv_imgproc

$ ./main
dev : GeForce GTX 580

i:0.840188 o:3.840188
i:0.394383 o:3.394383
i:0.783099 o:3.783099
i:0.798440 o:3.798440
i:0.911647 o:3.911647
i:0.197551 o:3.197551
i:0.335223 o:3.335223
i:0.768230 o:3.768229
i:0.277775 o:3.277775
i:0.553970 o:3.553970

まとめ

CUDA上でのポインター配列は「普通のポインター配列」のようにアクセスすることが出来るようです.
また,ある程度古いデバイスでも問題なく動作しています.
ポインターの書き込みは,ホストからでも,デバイスからでも問題なく動作するようです.
ただ,今回使ったデバイスコードは単純な物ですし,複雑なポインター操作をしたわけではないので,全ての環境で問題なく使用できるのかは分かりません.
かなり限定的な環境だけで動作するのかもしれません.

しかし,便利な機能であることは間違いないので,結果を確認しつつ,安全であれば使ってみたいと思います.

マニュアルでの記述も調べてみたいと思います.

2012年1月21日 (土)

CUDA上の定数

CUDAでプログラムを作るときに,double型を使っていないにもかかわらず

ptxas /tmp/tmpxft_000008ef_00000000-8_fdirdet_dev.ptx, line 95; warning : Double is not supported. Demoting to float

という警告が消えませんでした.
nvccに渡したオプションは,macでは

-m64 -O2 -lcutil_x86_64 --compiler-options -fpermissive

という感じです.

しょうがないのでptxを出力させたところ,こんな結果となりました(関連部分を抜粋)

ld.global.f32 %f1, [%r12+0];
cvt.f64.f32 %fd1, %f1;
mov.f64 %fd2, 0dbfe0f5c28f5c28f6; // -0.53
mul.f64 %fd3, %fd1, %fd2;
cvt.rn.f32.f64 %f3, %fd3;

-0.53の部分は,ソースファイルの中で#defineした定数です.
何故か-0.53とかけ算するために,一度doubleに変換し,その後floatに戻しているように見えます(多分)
たったこれだけの計算でいろんなレジスタを使っていて,いかにも遅そうです.

ちなみに95行目のコードは

mul.f64 %fd3, %fd1, %fd2;

です.サポートしてないという割には64bitっぽい命令ですが,恐らくcubinに変換されるときに,float型の計算になるのではないのでしょうか.

問題は,-0.53がdouble型となっていることが推測されるので,ソースコードを

-0.53f

と変更してみました.
すると,警告はなくなり,周辺のコードはすっきりしました.
やってることも一目瞭然でいい感じです.

ld.global.f32 %f1, [%r12+0];
mov.f32 %f3, 0fbf07ae14; // -0.53
mul.f32 %f4, %f1, %f3;

そういえば,Cの実数定数の型はdoubleだった気がします.
意外と律儀に型を守ってくれるみたいですね.その方が嬉しいです.

オプションでアーキテクチャの指定をすると,この警告がなくなるというのは読んだことがあるので,
これでは今まで効率の悪いcubinが出力されていた可能性がありますね.

ただ,速度は全く上がりませんでしたw
computeprofレベルで見れば変わっているかもしれませんね.

ところで,アーキテクチャの指定をしていなくても,64bitからの変換や,64bitレジスタの使用には警告が出ないと言うことは,どんなに古いCUDAの動くGPUでも64bitの値を保持する仕組みはあったということでしょうか?

いずれにしても,警告が0になって一安心です.
これで堂々と論文に貼れますw

ちなみに,この警告を直している間にほかのバグを二つ発見しました・・・危なかった・・・

2011年10月11日 (火)

CUDAデバッグ

卒業研究でCUDAを使っています。

自分で書いたデバイスコードがまともなのか心配なので、CPUで同じ動作をするコードを書いて比較するのですが、どうにも微妙、非常に微妙に値がずれます。

計算方法をチェックしても間違ってるとは思えないし・・・

で、いろいろ調べてるとやっと見つけることができました。(マニュアル読んでないだけかも)

CUDAのデバッグをCPUと比較して行う場合、次のことに注意するとある程度うまく行くようです。

1.積和演算器による誤差の変化

2.実数の演算順序

1について、GPUには、頻繁に行われる実数の積和演算( 1.4. * 2.2 + 3.3のような演算 )を行う装置があり、積和演算器を通す場合と、それぞれ乗算器と加算器を使う場合では結果が異なる。

2について、たくさんの加算を実行する場合、コアレッシングやバンクコンフリクトの影響を考えて加算の順番をCPUと異なる順番で実行すると、結果が異なる。

・・・といった感じです。

1については__fmul_rnと__fadd_rnという関数を使えば最適化を防止し、CPUと同じ動作をさせることができます。

2については多分、なのですが、CPUでGPUの動作と同じになるようにプログラムすると、ちゃんと結果が一致したので多分こういうことだと思います。

いずれにしても、CPUと同じ結果に持って行くには、複雑なプログラムだと結構大変だと思います。

無理に同じにしようと努力するより、浮動小数点の仕様からどの程度までのずれが誤差と言えるのか調べたほうがいいかもしれません。

ただ、ロックの位置やメモリ操作のミスで値がずれたりもするのでチェックしたい誘惑にもかられます(笑)

便利なことには違いないですが、気を付けていないとひょんなところから穴にはまりそうで怖いです。

しかしCUDAは速いな・・・

2011年5月 5日 (木)

恐怖Cの乱数

卒業研究で僕はニューロンのモデル(?)を扱う研究室に入りました。

とりあえず、練習用に作ってみるプログラム・・・のはずだったのですが、ニューロン数という値が1024の時だけおかしな挙動を示しました。

もう1024ってだけでバグの匂いがしますね(笑)

通常、ニューロンが覚える記憶パターン数が増えると、あるところで初期状態から正しいパターンまで戻せなくなるそうです。

しかし、ニューロン数1024の時だけ、なぜか全く忘れず、安定して八割方回復するのです。

さすがにおかしいと思い色々調べた結果、Cの標準関数の乱数によって生成した「バラバラなはずのパターン」が結局バラバラになっていなかったようです。

このモデルでは、ニューロンの状態に発火と不発火の2状態があり、入力パターンでのその発生確率は1/2にして実験すると書いてあったので、最初は

rand()%2

を使って、入力パターン(ニューロン1024個で1パターン)を作成していました。

しかし、実際にニューロン数1024で作成したパターン間の相関をとってみると、相関が1の物がそれなりの数現れ、さらに相関の値の変化に周期性が見られたりしました。

Asoc_n1024_p2p_rand

詳しく見てみると、例えば400パターンでは、あるパターンについて、負の相関も入れるとそれぞれ6個ずつ相関1のものが存在します。つまり、実際には全体で60個程度(!)のパターンしか存在していないことになります。
これだと、実際の記憶容量は0.05程度なので回復するのも当然です。

そこで、パターン生成に使う乱数をメルセンヌ・ツイスタと呼ばれるものに変えてみると・・・

Asoc_n1024_p2p_mt

相関が1の物がなくなって、周期性も見られなくなりました。

同じものを覚え続ければ、忘れなくなるのも当然な気がします。

とにかく、乱数というのはとても大事だなと改めて実感し、「Cの乱数は使ってはならない」という話は、場合によっては真実なのだなと痛感しました。

乱数って大事ですね。