Twitter

最近のトラックバック

無料ブログはココログ

2014年10月19日 (日)

移転します

bloggerに移転することにしました。
http://mfutr.blogspot.jp/

タイトルも
「プログラミングなど2」
にしてみました。

移転しても更新間隔はかわらないかも?

2013年11月21日 (木)

GistのTest

最近githubに手を出しました.
とりあえずgistが便利そうなので貼ってみます.

ちゃんと出るかな?

2013年8月29日 (木)

ココログフリー

久しぶりに見てみたら広告が出ていますね.
ココログは広告も少なく,すっきりしたデザインで好きだったのですが.
ちょっとだけ残念です.

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年5月 3日 (木)

C言語

最近、http://www.open-std.org/jtc1/sc22/wg14/で見つけたC言語の規格書を読んでいます。
Draftと書いてはいますが、ダウンロードできて検索もできるので、正直JISを苦労して見るより、かなり便利だと思います。

ところで、Draftって草案って意味ですよね?
草案だから公開してるのでしょうか?
でもこの草案を参考文献にしている所もあるし・・・いったいどういうことなのでしょう?
ただいずれにしても、個人で使う分には十分だと思います。

この規格書、今までいろいろと疑問に思っていたことが、一瞬で解決されてしまう魔法の書です。
英語で500ページ程度なので、そんなに大量というわけでもありません。

少しでもプログラミングを経験した今だからこそ読めるのかもしれませんが、最初から読みたかった気もします。

色々と複雑な気がする言語なので、規格書の内容も参照しながら入門させてくれるC言語の入門書があればなぁ、と思います。

しかしC言語ですらほとんど理解出来ていないのに、C++となるとどうしていいのか・・・という気分になってしまいますね・・・。

2012年3月 8日 (木)

fortRun更新 ( v2.11 )

fortRunを更新しました。

最新版は以下からダウンロードできます。

「fortRun.lzh」をダウンロード

■ 変更部分

・位置とサイズの保存の復元ができないバグを修正
(僕の環境でしかテストできていないので、もし問題があれば連絡ください)

・まとめてインデント・アンインデント機能追加

・フォートラン固定書式の場合、まとめてコメントアウトと、まとめてコメントアウト解除機能追加

2012年2月27日 (月)

パソコン

最近身近な人たちは就職や進学でパソコンを買ったりしてます.

そんなの見てると欲しくなるけど,学生の身分ですので,動くうちは今のパソコンで我慢します(笑)
派手な計算は研究室のパソコンで出来るし.

あ,でもBuilderのXE2は欲しいです.

6でも全然問題ないけど,そろそろ未来がないかなぁ,なんて思ったりします.

僕は趣味でやってるので,全然問題ないのですが,やっぱり仕事でとかなると,C#とかJavaとか,そういった技術を使いこなせないと話にならないのかなぁ

マイコンな部分でも,Cコンパイラーと高度なライブラリーが使えるのが普通みたいだし,Arduinoみたいな流行の部分は全然分からないし.
い,い,・・・一応AVRのアセンブリならちょこっとかける・・・よ?・・・

アセンブリで書けなきゃ楽しくないみたいな妙なこだわりが,結局何にもできない原因かもしれません
色々勉強しないと色々やばそうです・・・

Ubuntuどうしよう

今更ながらUbuntu11.10にアップグレードしてみたのです.
でも残念ながら,少なくともUnityは重くて使い物になりませんでした.

Core2Duo1.6GHzでメモリー1GB,Radeon200M搭載のEpsonのノートです.
考えれば5年くらい前のパソコンなので当たり前かもしれません.

Chromeで動画を見ながら,ソフトウェアセンターを起動するという,当たり前レベルの動作で急激に重くなり,なんとウィンドウマネージャー(?)ごと落っこちてしまいました.(ただ,そのあと端末に戻らずログイン画面に戻ってきたのは素晴らしかったです)

全体的にはいろいろ進化してていい感じです.
以前までなら終了がうまく出来ないことが多かったのですが,11.10からは問題なく終了できるようになりました(この辺はカーネルのバージョンの影響かもしれません)
UIも,「これでいいの?」感はあるけれど,工夫している感じがあります.

全体的に明らかにmacのパクリとなっているのも,ここ最近の流行と思えば仕方ない気もします.
ただ設定画面,あれはどうかと思います(笑)

起動してしばらく使った段階でメモリーを400MB以上消費していたので,恐らく想定されているスペックが11.04のgnome2の時より高いのだと思います.

今後もパソコンやスマートフォン,家電の変化は止まるはずもないので,Ubuntuがそれを見据えて進化していこうというのは素晴らしいと思います.

ただ,「友人に気兼ねなく勧められるOS」じゃなくなったかなぁ,という気もします.
僕のPCのスペックは今では「古いパソコン」に分類されるんだと思うのですが,そういうパソコンがやっぱり今でも現役な気がするのです.
そういうパソコンにそのまま入れて動く保証がなくなってしまったと思うと,ちょっと残念な気もするのです.

そういえば,ここ数年パソコン関連はいろんな変化がある気がします.
64bitが当たり前になったり,メモリーは2GB程度が当たり前になったり,MBRがGPTになったり.
スマートフォンが(日本でやっと)使われ始めたり,ネットワークにファイルを置くのが当たり前になったり,携帯のモデム化が当たり前になったり.

今まで少しずつ積み重なった変化が一気に現れ始めたという気もします.

そんなわけで,Ubuntuが変わっていくのも仕方ないのですが,僕の使えるOSがなくなってしまって困っています.
11.04のgnome2のままで行けるのも時間の問題だし.
Xubuntuという選択肢もあるのでしょうが,なんとなく,UbuntuはUbuntuで使えた方が楽な気がします.

Unity2Dも何故かもっさりとしか動いてくれません.
何かスペック以外の問題があるのかなぁ・・・

メモリ不足は確かにあるので,とりあえずメモリを買い足そうかと思います.

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年12月14日 (水)

歯が口内炎が

痛い!

«CUDAデバッグ