- 21日: 調子が優れず。しかし野郎だらけのクリスマスパーティに出席。
その後、次の日のゼミの発表資料を作る - 22日: なんか熱っぽいと思いつつ、朝6時ぐらいからゼミ資料を仕上げる。
資料が出来上がったので、10:30に健康サポートセンタに行ったら37.5℃ぐらい。
今ならすいているし、明日は祝日だからすぐに病院に行けという注意を受け、近くの内科へ
内科では37.8℃。インフル陰性で解熱剤と咳止めなどを処方して貰う。
その後、学校に戻りゼミに出席。ゼミ終了後、すぐに家に戻る。
家で熱を測ると38℃越え。氷枕を準備し、寝る。
夜: 39℃を超える。解熱剤を飲み、38.5℃まで下がる。 - 23日: 39.8℃あたりをうろうろ。10時ぐらいに母親から今から行くから！という電話が来る。
えー、まぢかよと返すが、背景の音を聞くと新幹線ホームで手遅れだと言うことを知り諦める。
母親襲来。解熱剤を再度のみ、37.5℃まで熱を落とすが、急に体温が下がったので貧血状態に。
トイレに立った時に貧血状態だったことが祟って倒れる。前歯が少し欠け、唇を切り血まみれに。
母親に促され、そのまま新幹線で実家に帰省。某駅を怪しい姿でふらふら。
夜には40.2℃ (ピーク) に達し、再び解熱剤を飲む。今度は貧血になりませんように。 - 24日: 38℃程度に熱が落ち着く。病院に行くと今度はA型インフル陽性。リレンザと咳止めを処方して貰う。
夜には37℃程度に落ち着く。 - 25日: 36.5℃と平熱になる。PC解禁でtwitterへ投稿をぽつぽつ。
学校に報告したら、健康サポートセンタのねーちゃんから電話を貰い、月・火曜日の行動を聴かれる。なんか答えづらいよね。
いやぁ。A型インフル菌と戯れるという素敵なクリスマスでした。あはははは…
一般教養が駄目な人間なので午前1がぼろぼろでしたが(自己採点でぎりぎり通過したことは解っていたけど)
なんとか合格して良かった。というか午後2もだめじゃん…
だいぶ前からThinkPad T43のCPUファンが完全に止まっていたのですが、
カーネル入れ替えのためリブートしたらFAN_ERRORで起動しなくなってしまったため
暫く死んでました。IBMよりパーツを購入し修理しましたので、12月3日午前4時より復活しました。
もし、死んでいる期間内にメールを送った方などは、再度送り直して貰った方が安全かもしれません
- rotate関数には対応していないみたいなので、手動巡回シフトに置き換える
- 何があっても、uint未満 (char, short) は使わない方が吉
(借り物だけど)GeForce GTX260を追加。HD5750との同時利用をすることになりました。
電源も500Wから650Wにパワーアップ(こっちは私物...orz)。だけど、Core i7 920 を3.5GHzにオーバークロックした
状態では、電源が足りないみたいで動作が凄い不安定でした。なのでオーバークロックを解除。
今のところ、安定して動作中。
Luffaは256bit単位で攪拌的な事をする仕様で、Luffa-256では256bitステートが3本あるので
この部分を並列化出来ます。というわけで、この部分を並列化する実装を書いてみました。
下に評価結果のグラフを載せます。左側のグラフは4KBのデータのハッシュを計算するときのスループットです。
横軸が並列数なので、右端では2048個の4KBのデータのハッシュ値を計算するという意味になります。
右側の図は効率で、スループットを並列数で割った値です。
なお、青いラインが今日作った実装で、赤いラインが11日に作った実装になります。
効率を見るとなんか面白いですよね。Radeon HD 5750なので、
1つのComputeUnit(=SIMD Engine)には16個のProcessingElement(=SIMD Unit, Thread Processor)が付いています。
んで、攪拌部分は3つ並列と先頭に書きましたが、諸事情によりLocalSize=4nとなっています。
なので、並列数が4(16÷4)～16付近がもっとも効率が良くなっている感じですね～
ちなみに、オフィシャルの64bit最適実装はCore i7 920(3.5GHz OC)では、980Mbpsでました。
なので、クアッドコア+HT補正(x1.5)とすると、8並列で5.8Gbpsでる計算でしょうか。
ビットスライスAESの方はさんざんな結果になっていますが、こっちは少しだけGPUで実行する意味がありそう？
(ただし、512並列というシチュエーションが存在するのか謎ですが)
openCLの.NETバインディングであるopenCL.NETの公開を開始しました。
ATI Stream PowerToysにPCIe Speed Testという物があって、
それによるとPCとGPU間の転送は4GB/sぐらい出てるみたい。
だけど、OpenCLのclEnqueueWriteBufferを使うとがんばっても2GB/sしか出ない。
ふとしたきっかけでリモートデスクトップ経由で使ってみたら、3GB/sも出た。
(だけど、リモートデスクトップ経由だとkernelの処理効率が三割ぐらい落ちる)
OpenCL周りの最適化不足が原因...? 結構基本的なところだと思うけど。
AESのビットスライス実装はwork-itemに細かく振り分けすぎたので、
ちょっと作り直しが必要な感じ。なので、気分転換?にSHA-3候補であるLuffa-256を
OpenCLで実装してRadon HD 5750で動かしてみた。(所謂、GPGPUってやつ)
ハッシュ関数は内部状態が膨大なため、ビットスライス実装をするときっと死ねるので、
普通の実装です。
ハッシュ関数はブロック暗号と異なりブロック間に依存性があるので並列計算できません。
なのでここでは別々のデータのハッシュ値を計算する場合を想定し、
並列数 (parallels) = 計算するデータ数 としました。
そして、各データの情報量は全て同じで、それはブロック数(blocks)で示しました。
(Luffa-256においてブロックあたりのバイト数は32バイトです)
| blocks | parallels | throughput |
| 1 | 1 | 0.014Mbps |
| 1 | 512 | 98.49Mbps |
| 1 | 1024 | 216.8Mbps |
| 1 | 2048 | 424.56Mbps |
| 1 | 4096 | 774.89Mbps |
| 1 | 8192 | 1337.68Mbps |
| 1024 | 1 | 0.014Mbps |
| 1024 | 512 | 5.41Mbps |
| 1024 | 1024 | 1915.94Mbps |
| 1024 | 2048 | 6630.51Mbps |
| 1024 | 4096 | 7446.99Mbps |
| 1024 | 8192 | 8399.05Mbps |
| blocks | parallels | throughput |
| 2048 | 1 | 6.275Mbps |
| 2048 | 512 | 2473.63Mbps |
| 2048 | 1024 | 4589.41Mbps |
| 2048 | 2048 | 7447.34Mbps |
| 2048 | 4096 | 8449.14Mbps |
| 4096 | 1 | 6.810Mbps |
| 4096 | 512 | 2810.42Mbps |
| 4096 | 1024 | 5000.69Mbps |
| 4096 | 2048 | 8316.72Mbps |
| 8192 | 1 | 7.138Mbps |
| 8192 | 512 | 3020.44Mbps |
| 8192 | 1024 | 5399.49Mbps |
CPUの場合はCore i7 3.5GHzにおいて、1スレッドで、540Mbpsほど出ました。
(オフィシャルの最適化版実装ではなく、openCrypto.NETに含まれる実装を用いた場合の評価結果です。おそらくオフィシャルの方が速いと思います)
なので、4-threadで2160Mbps。SMTを使って...3000Mbpsぐらいでしょうか？
CPUの方が速いですね (笑
しかも、同時に1000個とか2000個のデータのハッシュ値を計算するとか
そういうシチュエーションが思い浮かばない...
なんて無意味な実装なんでしょう。
if文とかswitch文とかがあると、きっとストールするんだろうなぁと勝手に思っていて、
使わない代わりに 32x32 の行列演算をさせていたんだけど、
行列の要素の70%が0となってしまっていたので、そう言った場合は、
素直に switch文で切り替えた方が良い感じ。
というわけで、926.7Mbps -> 1.01Gbps に高速化！
__constantや__globalなメモリを参照する箇所を減らすために、
参照回数が多い場合は __local なキャッシュに一度コピーしてから参照するように改良。
16MBの時のスループットが、779.8Mbps ->926.7Mbpsに向上
いろいろ計測してたら、16MBのデータをGPUにコピーし、結果をPCにコピーするだけで50msかかってる。
これは全体の30%に相当するため、この部分を隠蔽できれば、1Gbpsを超えるスループットが出る計算になる。
そこで、clEnqueueWriteBufferでメモリの一部を非同期で書き込んで、
その書き込みが終わったらclEnqueueNDRangeKernelのglobal_work_offsetに
書き込んだ部分のオフセットを指定すれば、メモリ転送時間を隠蔽できる！
と思いましたが、global_work_offsetは現在0 (NULL)以外指定できないそうで…
実装レベルで未対応じゃなくて、仕様書レベルで未対応なのかYO
MixColumnsでは行列演算をさせているのですが、
行列を__constant (readonlyな__global)から参照してました。
なので、一度 __local にコピーしてから利用するように変更。
その結果、16MBの時のスループットが482.1Mbpsから779.8Mbpsに向上
というわけで、OpenCLでAESのビットスライス実装をし、昨日買ったRadeon HD 5750で動かしてみた。
実装の指針としては、Processing Elementで1bit * 128bitを処理し、Compute Unitではそれを128個管理する。
なので、work-itemではuint4 (32bit*4のvector型) の処理を行い、work-groupで128個のAESブロックを処理する。
こうすることで、ラウンド処理に必要なデータは全てローカルメモリに収まるようにしたつもり。
(実際はGPUのローカルメモリサイズによって収まらない場合があるかも。そう言う場合はエラーで動かないのかな？)
| 2KB | 64KB | 512KB | 1MB | 2MB | 16MB | 64MB | |
| AMDのAES実装 | 9.84Mbps | 32.6Mbps | 64.28Mbps | 14.765Mbps | 14.75Mbps | 14.8Mbps | - |
| ビットスライス実装 | 8.43Mbps | 185.77Mbps | 257.85Mbps | 271.35Mbps | 352.36Mbps | 482.1Mbps | 505.6Mbps |
AMDのAES実装というのは、ATI Stream 2.0 Beta.4のサンプルに付属する実装で、
ProcessingElementに32bitずつ割り当てて、SBOXのテーブル参照を行う普通の実装。
何故か、512KBをピークにスループットが悪くなる。
 |
| Radeon HD 5750 を上から撮ってみた |
 |
| 一番手前のPCIがONKYOのSE-90PCI。次がHD5750。奥がHD2600Pro |
 |
| ケース内を俯瞰 |
しかし、何故かOpenCLのプログラムはGPU上で動かない…
Radeon HD 5750はまだ届いていませんが、デスクトップマシンのGPUは現在、HD 2600XTなので、
一応、OpenCL対応版のドライバを導入できるんですね。
そうすると、CPUをデバイスとしてOpenCLプログラムを動作させることが出来るのですが、
ノートPCでもOpenCLプログラムを書きたいという場合、SDKを導入するだけじゃ、DLLが足りないって怒られて
コンパイルは出来るんだけど、実行できません。
そこで、Radeonドライバが導入されているマシンのSystem32フォルダから、
aticalcl64.dll と aticalrt64.dll をコピーして exe と同じ場所に置けば、
CPUをデバイスとして動いてくれます。
これで、どこでも開発できますね。
(早く、Radeon HD 5750 届かないかなー)
後輩に頼んで車を出して貰って、PCショップ2店を回ったけど、どちらも在庫切れ。
発売してから一ヶ月もたってるのに…なんでだろ？
しかし、GPGPUをしたいという欲求は抑えられず、寝不足＋深夜という何時ものテンションで
上納金を納めてヤフオクで落札w
今は、寝て起きて次の日なんだけど、後悔は・・・してない！
というわけで、よさげ。
残念なことに受信機を一人暮らし先では所有していないので、NHKに税は納めていませんが、
実家ではBSも含めて納めているので、なんか良い用途が思いついたら使ってみよう。
Googleのウェブマスターツールを使って、
クロールエラー(HTTP404)となっているページのうち、
URLが変わってしまって、リンクエラーとなっている物を
新しいURLへリダイレクトするように調整した。
| リンク切れURL | リダイレクト先 |
| /audioencoder.php | /software/audioencoder |
| /blog/ | /memo |
| /default.aspx/software/audioencoder-old/downloads.xml | /software/obsolete/audioencoder/downloads |
| /default.aspx/software/audioencoder/overview.xml | /software/obsolete/audioencoder |
| /default.aspx/software/daily.xml | HTTP 404 |
| /default.aspx/software/obsolete/audioencoder/summary.xml | /software/obsolete/audioencoder |
| /dotopen/pictable.php | HTTP 404 |
| /index.php?AudioEncoder | / |
| /index.php?Mono%2FPatches%2F0009 | / |
| /index.php?News | / |
| /mocean/TaskbarEx.htm | HTTP 404 |
| /mono/mono.xml | HTTP 404 |
| /old/ | /software/obsolete |
| /old/audioencoder.php | /software/obsolete/audioencoder |
| /software/crypto | /software/opencrypto |
| /software/index | /software |
| /software/obsolete/audioencoder/top | /software/obsolete/audioencoder |
| /trac/wiki/FAQ/0001 | HTTP 404 |
/index.php?hoge というページは、/index.php でしかハンドリング出来ないので、
全部、トップページにリダイレクトするようにした。
というわけで、高校時代(3年後半あたりからかな？)から使っていたけど、
EmEditorとおさらばすることになるかも…64bit版がすぐに用意されたりと良い感じだったんだけどね。
EmEditor + Meadowという謎の使い分けをしないで、いい加減どちらか片方に統一しろと言うことなのかな。
しかし、MeadowはWindowsとの統合具合が微妙だし、EmacsはIMEとの連携がいまいちだし…
むぅ…w
長年お世話になったんだし、ライセンス買っとくかな…どうすっかな…