2009-11-01から1ヶ月間の記事一覧

IBMが次世代Cellの開発中止?

少し前にIBMがCellの開発から撤退するという噂がネット上で流れていたが、それに対してIBMから開発中止を否定する公式コメントが出たらしい。IBM Says Cell Processor Technology Will Continue In the article, David Turek, vice president of deep comput…

CUDA最適化実習(2)

2048 x 2048のCUDAを使った行列の転置も完成。 最適化していないCPUでの実行の10倍ぐらいは速くなった。この処理の場合はHostとDevice間の転送がボトルネックになってくるから、その部分をどうするかが鍵になるな。 と言っても、選択肢はそれ程ないけど。

CUDA最適化実習

いろいろあって、次の3つの関数をCUDAで最適化しないといけない。 512x512の大きさの正方行列同士の乗算 2048x2048の正方行列の転置 128k個のfloatの値の総和 とりあえず乗算は終わった。 時間は毎回のぶれが大きいが速いときは3msぐらい。 最適化せずにcpu…

本購入

昔からずっとMakefileの文法をちゃんと覚えたいと思っていつかは本を買おうと思っていたが、思い切って買ってみた。make 改訂版 (A nutshell handbook)作者: Andrew Oram,Steve Talbott,矢吹道郎,菊池彰出版社/メーカー: オライリー・ジャパン発売日: 1997/0…

Teslaの信頼性は?

パフォーマンスの点では値段の高いTeslaよりも安いGTX285やGTX290の方が高い。 おまけにGTXは画像の出力機能まである。 Teslaの値段が高いのは信頼性が高いからと言われるがどのくらい違うのか。そのヒントになりそうなのが次の論文(右側のdownloadの所から…

global memoryからshared memoryへの大量の読み込み(補足)

昨日の説明は文章だけでわかりにくかったから、矩形領域をglobal memoryからshared memoryへ読み込むコードと結果の追加。 uint4で読むときには一列をhalf warp(16thread)の1回の読み込みで済ましている。(今回は横幅の最大を64要素に制限しているので) int…

Windows 7でCapsLockをESCと入れ替える

自分は秀丸使いなのでLinux上のファイルを編集するときにはsamba+秀丸で作業するが、ちょっとした編集ならviを使うこともある。Windows VistaではXKeymacsというソフトでキーの入れ替えを行っていたがWindows 7ではこのソフトは動作しなかった。viはESCを多…

nvccで-deviceemuの指定

はじめてのCUDAプログラミング―驚異の開発環境[GPU+CUDA]を使いこなす! (I・O BOOKS)作者: 青木尊之,額田彰,第二I O編集部出版社/メーカー: 工学社発売日: 2009/11メディア: 単行本購入: 11人 クリック: 310回この商品を含むブログ (29件) を見る11/20にCUDA…

global memoryからshared memoryへの大量の読み込み

複数のthreadを利用して画像のような矩形領域をglobal memoryからshared memoryへ読み込む場合、しかも全てのthreadがint型で読んでも1回じゃ読み切れない場合にはint型で複数回読むのとuint4型で読むのとどちらが良い速いのか。(uint4でなら1回で読み込…

CUDA Profileで値が0ばっかりで正しい値が取れていないときには

blockの数を増やしてみるとうまくいくことがあるらしい。

CUDAの初期化を速くする(ただし今が遅い場合に限る)

今日はCUDAについて、良い情報がいろいろ得られた。 しかし、どれも会社の別の人が編み出した事ばかりなのでどこまで書いて良いか微妙。 なので、ほどほどな感じで書く予定。CUDAのプログラムで初期化(最初に呼んだCUDA API)が1秒近くかかるのは改善可能。…

ばれた

知り合いに誰にも教えずにこのブログ書いていたけど、ばれてしまったようだ。 まあ、何でも知っている人だから間違ったことを書いたら指摘してくれるだろう。

社内勉強会

社内の持ち回りの勉強会で、今日の当番の人が忙しくて辞退したので、代わりに担当してみた。 CUDAでの最適化と画像に関するアルゴリズムの話をした。 自分で話しておいて何だが、アルゴリズムに関する数学的な質問は答えられないので一緒の仕事をしている頭…

健康診断

無事終わった。 本当に無事だったかは結果出るまで分からないが。 (前回も要経過観察があったし)聴覚検査は緊張した。うまく押せたか今回は自信がない。バリウムもやっぱり辛かった。 まあ、最近は胃の働きを弱める注射を打つから、昔よりは飲む量がずっと…

cpuspeed

Linux上でベンチマークを取る事が多い人はcpuspeedというserviceを止めるといいらしい。(root権限必要) cpuspeedはcpuのアイドル比率に応じてCPUのクロックを上げ下げしてくれるらしい。 エコだ。一時的に止める場合 #service cpuspeed stop ずっと止める場…

明日は健康診断

朝は早いし食事制限あるし、いろいろと億劫だ。 特にバリウムが嫌だ。

OpenMP

下のサイトとか読んでOpenMPについて軽く調べてみた。http://www.na.cse.nagoya-u.ac.jp/~reiji/lect/hpc02/OpenMPintro.html実際に使ってみないとよく分からないな。 説明を読むとOpenMPよりもMPIのが良さそうな事が書いてある。 どちらも使ったことはない…

ブラック会社に勤めてるんだが、もう俺は限界かもしれない

映画を見てきた。 電車男とかと同じで2chが元ネタ。まとめサイトで読んだ内容からは、少し映画用にアレンジされてた。 登場人物も整理されていたし。ブラック会社に勤めてるんだが、もう俺は限界かもしれない@まとめwiki - アットウィキ観客は、公開直後の…

Twitterデビューなう

つぶやいたら負けと思ったけど、Twitterでつぶやく知人が増えてきて、一生懸命グーグルリーダーで拾って読んでたけど、数が増えてきたのでTwitterデビューしてしまった。ためになるブログを書いていた人がTwitter始めてからブログをまったく更新しなくなった…

CUDA本読み終わった

11章の粒子計算や12章の偏微分を最適化しながら実効性能を乗せてくれていたのは良かった。 (自分には計算式に、ついていけない部分もあったけど)付録Aに載っていたマニュアルへのパスの通し方は早速.bashrcに追加しておいた。 export MANPATH=$MANPATH:/us…

SMあたりのBlock数

CUDA本はまだ読み途中。性能を引き出す上で6章に書いてあった、Occupancyを高めるよりもマルチプロセッサあたりのアクティブなブロック数の方が大事というのはいいヒントになりそうだ。ついついBlockあたりのThread数を上げようとしてしまうが、同時におよそ…

cudaMalloc()を利用しないでDeviceメモリを確保する方法

__device__をつけてバッファを用意する。 HOST側のプログラムからは次の関数を利用してアドレスを取得する。 cudaError_t cudaGetSymbolAddress (void **devPtr, const char * symbol); こんな感じになるかな __device__ int d_buff[256]; int main(void) { …

目薬

薬局で600円ぐらいの少し高めの目薬を買おうとしたら、店員に同じぐらいの値段でもっといいのがあると、別の商品を勧められた。その時に、いろいろアドバイスを聞けた。 充血に効く目薬は、慣れがあるので充血が治まったら使うのをやめた方がいいらしい。 疲…

火事

家のそばの居酒屋が燃えていた。使ったことは無い店だけど。 消防車と野次馬が大勢きていた。燃えている店に面した道路の端っこを通ったけど、かなり熱かった。 火に近づかないといけない消防士は大変だ。商店街の込み入った一角だから、周りに広がらなけれ…

IBMがCellの開発から撤退?

http://www.playstationuniversity.com/ibm-cancels-cell-processor-development-1295/ うーん。ショックだ。 Cellのプログラムも好きなのになぁ。細かいレベルの最適化が好きな人間としては非常に残念だ。 今後のIBMの別のチップに少しでもCellの魂が引き継…

映画 マクロスF

初日に見に行くのは、なんとなく負けな気がしたけど席が予約できたので見に行った。 綺麗に纏まって終わったTVシリーズから、どのように話を続けるのか気になっていたけど、時間を戻しての総集編のような感じになっていた。 (少し話が変わりそうなところも…

はじめてのCUDAプログラミング

秋葉ヨドバシで最後の一冊を手に入れた。 あまり仕入れていないのか、それとも凄い売れているのか?まだ5章までしか読んでいないけど、4章に書いてあった-deviceemuと-D__DEVICE_EMULATION__を 同時に指定するのは良い方法かも。 Kernel中にprintf()とかasse…

CUDAドライバのバグ?

doubleを利用するならcudaドライバとtool kitは最新の2.3を使った方が良いかもしれない。 doubleを利用して-arch=sm_13を指定したときに2.2でのみ発生する問題に遭遇した。 (もう少し切り分けのための検証が必要そうだけど)でも使っているFedora 9には2.3…

探査機「はやぶさ」

2chからのコピペ 日本の技術者すげーー ●本来ならリアクションホイール3個で姿勢制御 ⇒ ホイール1個壊れたのでホイール残り2個と化学スラスタで制御 (ここまでは普通の「こんなこともあろうかと」、な範囲) ⇒ ホイールさらに1個壊れた!化学スラスタ全損!燃…

CUDAでのKernel関数のポインタテーブル

CUDAでは最適化のテクニックとしてKernel関数をC++のtemplateを利用して量産して、適切な関数を呼び出すというのが結構使える。 たとえばKernel関数内部の最内のループ回数が1〜10の間で可変のときに、 ループする回数ごとに10種類の関数を用意して利用する。…