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年8月11日 (木)

パーティションを切りたい

Linuxを使うとなると、パーティションの話やマルチブートの話からは逃れることができません。

パーティションについていろいろ調べてみましたが、やっぱり自分で切るのは難しいですね・・・

今回も、以下のサイト様の情報が非常に役に立ちました。本当にありがとうございます。
http://wikiwiki.jp/disklessfun/?hduse
http://www37.tok2.com/home/nobusan/

以下の内容は僕個人の、思ったこと、分かったことのメモなので不正確な情報である可能性が高いです。
読んでくださる方は、その上で読んでいただけると助かります。

いろんなサイトによる情報をまとめると・・・

CHSはもはや形だけ(?)
これからは物理的には4kセクタのHDDが出てくる。
よってパーティションの先頭は512バイトセクター番号で8の倍数にすべき?
SSDはページと言う単位やブロックという単位で読み書きされるので、パーティションをそれに合わせて切るべき。
fdisk -H 64 -S 32
昔は H255 S63 ( WinXPSP1以前等、古いOSではこれ + 下の2つのルールで切るのが安全? )
最初のパーティションは第二トラックから。
パーティションはシリンダー境界で切る。
GPTというパーティションテーブルではLBA33まで使われるらしい?
XPSP1以前では48bitLBAに対応していないので128GByte内のパーティションへ?

・・・という感じ?

シリンダー境界などはもうあまり意味が無いのでは、とも思いますが、不安なので揃えたいです。

たぶん-H64 -S32で第二トラックからパーティションを切れば先頭のパーティションの先頭のLBAは8の倍数な気がします。
ただそうすると、GPTというパーティションテーブルに微妙にかかっている気もします。
そもそもMBRでブートするHDDと共存する状況があり得るのかわかりませんが、なんとなく怖いです。
しかし、さらによく見ると、第二シリンダーからパーティションを切る、とあります。これなら確かに余裕ですね。

まとめると、

「H64 S32で第一パーティションを第二シリンダーから、以降、パーティションはシリンダー境界で切る」

・・・と言うのがベストな方法・・・なのかな?

やっぱりパーティションは大変です。


その他メモ:

とりあえずGpartedかfdiskで先頭1MiB開けて、以降1MiB単位で切るのが安全そう?

Ubuntuでgrubのインストール先を間違えたら

sudo dpkg-reconfigure grub-pc

というのが有効そう?

2011年8月 5日 (金)

OS起動の仕組み

最近Linuxをインストールしたときに、OSの起動やパーティションテーブルについて何も知らないことを痛感したので、いろいろ調べています。

そんなわけで備忘録(?)を

・・・というか、パーティションやブートに関する理解は、ほぼこちらのサイト様だけで十分です

http://www37.tok2.com/home/nobusan/index.html

それにしても、OSのブート周りは非常に複雑ですね。きっと規格の拡張を繰り返した結果こんな風になったのだとは思いますが、覚えることが多すぎてちょっと困ります。

□ その他

HDDの先頭部分には空き空間があって、そこにブートプログラムが入ったりする、こともある。
それは、MBRから続く1トラック分の領域。

ここを使ってPBRのブートローダーをロードできるブートローダー(?)を使うのが一番安全?

fdiskのHSの指定は物理的意味じゃなくて、パーティション区切りのルールを付加するという意味でいいのだろうか?

そういう意味で、SSD等のパーティションを切る際には注意が必要(区切りがハード的に意味を持ってるから)?

2011年3月27日 (日)

ブラウザ上でC言語

ブラウザ上でCがコンパイル、実行出来るなんて・・・

http://ideone.com/

こういうのをクラウドっていうのかな?

クラウドとかは良くわからないし、どうなのかな~とも思うのですが、これは純粋に楽しいです。

しかもです、scanfで入力もできるし、ctypeとかstringと言ったヘッダファイルをインクルードしたり出来ました。
どうなってんだこれ・・・?

それで、まさかとは思ってこの前作った電卓プログラムをコピペしたら・・・

input:
a=2*(3-1)+cos(pi),b=2,a+b

output:
5.000000
0x5
5.000000e+00

動く!?

数学系の関数も動くみたいです・・・恐ろしい・・・

いや~・・・すごい!

・・・ということはブラウザ上でCコンパイラをコンパイルしてコンパイルする(?)・・・とかできるかもですね。

あ、それと作ったコードは公開したり出来ちゃうみたいです。すげ~
http://ideone.com/y2xJv

こんなにわくわくするサービスは久しぶりです。
あと、とてもたくさんの言語に対応してるので、ある言語を試してみたくなった時環境をいちいち作らなくていいので楽でいいですね。

とにかくとても楽しいサービスでした!

2009年10月 6日 (火)

MinGWすげぇ・・・

tolsetから行き着いたのですが、

なかなか、MinGWもいい感じ!

bccと同等に、なんでもできます。

今回は、C++Builderに頼らないで作ってみようかなぁ・・・

リソースコンパイラ(?)まで付いてますw

2009年7月13日 (月)

picoPad

大分前に作ってたPIC用エディターをちょっぴり掘り起こし中です。

色分け機能はそれなりに便利かなぁ、なんて思ってます。
もうちょっと安定したらいつか公開してみたいです。

ダイアログとかはかなりTeraPadに似ているんですが、わざとですw
TeraPad使ったことあればすぐ使えるようにできたらいいなぁ~なんて思ってます。

いつ完成するやら・・・w

さて・・・そんなことより、勉強せねば!
してきますw

Pp

2009年6月29日 (月)

落ちる理由

デバッガを見てたら落ちる理由がわかったかもしれません・・・

・・・なんか0除算エラーが、マーカーをカメラの正面に向けると発生します・・・

なぜに!?
その周辺のコードを読んでも、それこそちんぷんかんぷん・・・

とにかく、arDetectMarkerでおっこちます。
具体的にはココです。

Zdiv

もう、if文で無理やり脱出させて事なきを得る手段に出ますw

・・・読めるようになれば、それこそ、いいんですけどね・・・

でも、かなり面白いことをしてそうなので、そのうち読んでみたいです。
なんせ、こここそ心臓部ですもんw

2009年6月28日 (日)

途中経過

かなりいい感じに動作してます。

Glut3DTVと言う、わけわかんないタイトルですw

ARToolkitのマーカーを頭にのっけるアイデア自体はニコニコからの転用です。

たまーに謎のエラーで落っこちます。

間に合うかな?

あ、ちなみにこれはかなりボツになる可能性が高いですw

当日は、少なくともgoogleearthの立体視をする予定です。
去年よりソフト自体は速度が上がっているんですが、問題はgoogleearthのモデル量が増えて、カックカック状態になってることです・・・

Mk

2009年6月23日 (火)

PTAM

マーカーレスARなんてものがあるんですね!!
http://kougaku.blog28.fc2.com/blog-entry-317.html

世の中すごい人はすごいなぁ~・・・

僕の知識じゃコンパイルすら不可能ですが、でもいつか使ってみたいです!