12/11/12 13:31:46.97 ClnWJME80
khronosの姿勢として標準のカーネルコンパイラを用意しないのはわかるんだけど
やっぱりglslの轍をちょっとは生かしてほしかったってのが個人的な思い
meriken氏乙
601:名無しさん＠お腹いっぱい。
12/11/12 13:37:18.53 /PHpLzn40
>>594
当方では一応4M/sぐらいで動くので、
IntelHD4000を使うか否かをチェックボックスとかで決めればいいと思いまーす
602:名無しさん＠お腹いっぱい。
12/11/12 13:56:47.61 FaMyVn9Z0
>>598
そうなんだ
ありがとう
603:名無しさん＠お腹いっぱい。
12/11/12 21:58:25.06 +66fUhHT0
OpenCLで盛り上がっているところにCPUのみの結果を報告。
【GPU】Quadro FX 3800
【CPU】Xeon X5680@3.33GHz x2CPU
【OS】MS Windows 7 Pro 64bit
【バージョン】0.07 Alpha 3 CUI64
【トリップの種類】12桁
【Display Driver】307.32
【その他】HT on
【その他のオプション】-c -t 24
【60時間の平均速度】80.51M TPS
【その他】HT off
【その他のオプション】-c -t 12
【2時間の平均速度】79.04M TPS
CPUだけで実行してもHTは殆ど効きません。NehalemとSandy Bridgeでは全然違うのかな?
ちなみにHT on の状態で、"-c -t 12"と指定すると、2CPU12コアに割り当てられずに、1CPU6コア12スレッドに割り当てられてスピードが出ません。
604:名無しさん＠お腹いっぱい。
12/11/12 21:59:57.08 +66fUhHT0
Alpha 6に上げて再度実行してみましたが、NVIDIAコントロールパネルの"3D設定"→"3D設定の管理"で"CUDA-GPU"を"なし"に設定すると、CUI64で"-c"オプションをつけても下記エラーが出て落ちます。
MERIKENsTripcodeFinderCUI: OPENCL FUNCTION FALL FAILED: Unknown (file 'Source Files\MTF_CUI_Main.cpp', line 715)
605:名無しさん＠お腹いっぱい。
12/11/12 23:26:48.48 ODHrB/Fw0
NVIDIA，第2世代Kepler「GK110」搭載の「Tesla K20」を正式発表。CUDA Core数は最大2688基に
URLﾘﾝｸ(www.4gamer.net)
606:名無しさん＠お腹いっぱい。
12/11/12 23:34:15.80 cklfqCrp0
X5680はOCすりゃいいじゃん
607:名無しさん＠お腹いっぱい。
12/11/13 02:39:21.60 FPgsAJYQ0
DualCPUにQuadro突っ込んでるようなガチWS機でOCとかあり得んでしょ
608:名無しさん＠お腹いっぱい。
12/11/13 02:50:45.59 l+bGYcyn0
倍率ロックフリーだろ？
609: ◆MERIKEN4.k
12/11/13 06:21:00.78 FpPqufE20
今気づいたんですけど「1CUあたりのワークグループの数」じゃなくて
「1CUあたりのワークアイテムの数」ですね、これ。
こりゃ当分の間安定版は出せないな…
610: ◆MERIKEN4.k
12/11/13 06:26:29.04 FpPqufE20
>>603-604
報告ありがとうございます。CPU検索ももうちょっと何とかしたいですね～
"Unknown"のエラーが出ているのは謎ですが、そこのエラーは無視するように
直しておきます。
611: ◆MERIKEN4.k
12/11/13 06:38:08.95 FpPqufE20
予想通りというべきか、10桁トリップ検索はなかなかスピードが出てくれません。
まじめにプロファイラを使わないと駄目ですね、こりゃ。
まあCUDAのときもそうだったので、のんびり時間をかけて取り組むことにします。
612: ◆MERIKEN4.k
12/11/13 07:26:06.97 FpPqufE20
>>605
GK110も試してみたいんですけどね～
Amazon Cluster GPU Instancesで使えるようにならないかしらん。
613: ◆MERIKEN4.k
12/11/13 09:29:27.60 FpPqufE20
Bitslice DES用の一時変数をどのメモリ空間に置けばいいのかいまいち
よくわからないので、とりあえず#ifdefで切り替えられるようにしておきます。
あと、一回のBitslice DESを複数のスレッドで同時に処理するかどうかも
CPU側で設定できるようにする予定です。こういうところは実行時にカーネルを
ビルドできるOpenCLはいいですねえ。
614:名無しさん＠お腹いっぱい。
12/11/13 09:34:06.33 ATY55mX00
【GPU】HD7970 CFX 2GUPｓ
【CPU】FX-8350
【OS】Win7 64bit
【バージョン】MERIKEN's Tripcode Finder 0.07 Alpha 6
【トリップの種類】12桁
【1CUあたりのワークグループの数】5120
【1WGあたりのワークアイテムの数】40
【その他のオプション】-c -g -t 6 -a 8 （-aオプションは有効なのか？）
【Display Driver】Catalyst 12.10
【10分間の平均速度】2614.21tripcodes/s
【GPUの平均速度】2575.40tripcodes/s
【CPUの平均速度】38.31tripcodes/s
【GPU使用率】95%
【その他】7完1タゲ
いろいろ調整したら瞬間最大風速では3000M程度出るようにはなりました
ラデはCPU負荷高いですね
フルにCPU8スレッドで回すと却って速度が出ないです
CPU単体でなら50M程出るんですが
それと、CFXの有効無効では速度は変わらないですよ
URLﾘﾝｸ(www.rupan.net)
615:名無しさん＠お腹いっぱい。
12/11/13 14:32:35.77 1xHvqvP20
ついにデュアルTahitiカードがAMDから登場。エーキューブ，サーバー向けグラフィックスカード「FirePro S10000」を国内発売
URLﾘﾝｸ(www.4gamer.net)
616:名無しさん＠お腹いっぱい。
12/11/13 14:34:05.53 1xHvqvP20
XeonPhiってどうなんですかねえ
617:名無しさん＠お腹いっぱい。
12/11/13 15:01:20.15 PlVYlKIA0
負荷で思い出したけど、同じGPU使用率でもmtyのGPU版は95℃までいくけどMERIKENは89℃までしか上がらないね。
618:名無しさん＠お腹いっぱい。
12/11/13 17:19:59.24 VU2bD6Zz0
>>615
サーバ向けのFireProだし、3599ドルらしい。
>>616
OpenMPが使えて既存アプリの移植が楽らしいけど、既にOpenCLになっている場合はメリット薄そう。
619: ◆MERIKEN4.k
12/11/13 19:10:11.95 FpPqufE20
>>614
なかなか良い感じに仕上がってますね。それだけOCした7970なら単体で1500M TPS近く
いくはずなので、単純に2枚で2倍の速度という訳にはいかないみたいですが…
Alpha 6では-aオプションは有効なはずです。CFXの話は別のところで見かけたんですが、
関係無かったみたいですね。
620: ◆MERIKEN4.k
12/11/13 19:22:18.89 FpPqufE20
>>617
Bitslice DESはSHA-1の処理に比べてメモリへのアクセスの量が段違いですからねえ。
621: ◆MERIKEN4.k
12/11/13 19:30:43.61 FpPqufE20
>>615
いいですね～ いつかこういうのをたくさんラックマウントサーバーに乗っけて
Tripcode Finderを動かしてみたいですｗ
622: ◆MERIKEN4.k
12/11/13 19:59:41.68 FpPqufE20
Southern Islandsだとコンスタントメモリは場合によってはグローバルメモリと
同じぐらい遅くなるそうで…こりゃCUDAと同じコードじゃ遅くなるわけだわ。
頻繁に使うのは最初にローカルメモリに移しておいたほうがいいな。
> 3. Varying Index
> More sophisticated addressing patterns, including the case where each work- item
> accesses different indices, are not hardware accelerated and deliver the same
> performance as a global memory read.
623: ◆MERIKEN4.k
12/11/13 20:10:06.56 FpPqufE20
あとローカルメモリにアクセスする際はuint2を使うといいみたいです。
> Currently, the native format of LDS is a 32-bit word. The theoretical
> LDS peak bandwidth is achieved when each thread operates on a
> two-vector of 32-bit words (16 threads per clock operate on 32 banks).
vector data typesの使い方はここに書いてありました。
Programming with OpenCL C
URLﾘﾝｸ(www.informit.com)
624:名無しさん＠お腹いっぱい。
12/11/13 20:28:02.76 ES3128Qj0
>>614
壁紙についてｋｗｓｋ
MERIKENさんの公式記録が越される日も近いか……
625:名無しさん＠お腹いっぱい。
12/11/13 23:12:02.50 ATY55mX00
>>624
この辺で拾ってます
これがいつのものだったかは忘れましたw
URLﾘﾝｸ(www.smashingmagazine.com)
626:名無しさん＠お腹いっぱい。
12/11/13 23:43:31.28 ES3128Qj0
>>625
そのまま2012/11の記事を見ても出てこないorz
URLﾘﾝｸ(www.smashingmagazine.com)
よろしければ画像アップローダに上げてくれませんか？
627: ◆MERIKEN4.k
12/11/14 00:24:53.55 mDY3eRDW0
>>624
それはどうでしょうね… (￣ー￣)ﾆﾔﾘ
628:名無しさん＠お腹いっぱい。
12/11/14 00:35:11.76 E/rQ2cMp0
>>626
ほい
URLﾘﾝｸ(www.rupan.net)
629:名無しさん＠お腹いっぱい。
12/11/14 00:38:28.17 eHqSRvz80
>>621
どうせならHD7970 X2に行きませんか？
消費電力が凄まじいのと、スロット占有が問題ですけどｗ
なんかリンクが貼れないので詳細は検索してください
デスクトップ向けにHD7950のデュアルが出てくれれば一番ですけどね。
HD7950のCFはグラフィックでも割りと良いというレビューもあったので、需要もある程度ありそうですし。
630:名無しさん＠お腹いっぱい。
12/11/14 00:50:05.56 eHqSRvz80
>>622-623
OpenCLは以前よりは情報も増えたようですが、まだ茨の道なのでしょうかね・・・
631:626
12/11/14 01:34:12.64 vuLXlPiG0
>>628
ありがとうございます！
632: ◆MERIKEN4.k
12/11/14 01:39:17.60 mDY3eRDW0
>>630
>>622-623はJohn the Ripperのメーリングリストを見てて気づきました。
URLﾘﾝｸ(www.openwall.com)
URLﾘﾝｸ(www.openwall.com)
このAlexanderという方はJohn the RipperとBitslice DESの偉い人です。
流石です。
URLﾘﾝｸ(www.openwall.com)
633: ◆MERIKEN4.k
12/11/14 01:40:56.22 mDY3eRDW0
>>630
まあGPGPUの不条理な制約にはCUDAで慣れっこになっているので
どうということはありませんｗ
634: ◆MERIKEN4.k
12/11/14 01:45:30.91 mDY3eRDW0
>>629
うちの検索用マシンにはGTX 580が2枚と590が1枚載っているので、
7970 2枚は余裕ですｗ 今590を売っぱらって6990を買おうかどうか
考えているところです。
635: ◆supernova.rT
12/11/14 02:04:56.65 Bf0HEkX10 BE:1020114162-DIA(123421)
僕はもうラデ2枚構成にしたのでゲフォ売ります
10桁検索対応が楽しみですよー
636:名無しさん＠お腹いっぱい。
12/11/14 02:16:09.00 eHqSRvz80
>>633
頼もしいです、頑張ってください。
>>634
HD7970を1ボードに2基載せたもので8ピンx3で3スロット占有という
モンスターというかクレイジーな代物が出るらしいですｗ
それの複数枚挿しは電源だけでなくマザボもかなり選びそうです。
HD7950のデュアルで8ピンx2で2スロットであればまだマシなのですけどねえ。
637:名無しさん＠お腹いっぱい。
12/11/14 03:47:50.59 peEcrqnb0
やっぱりさよならゲフォの流れになったね
638:名無しさん＠お腹いっぱい。
12/11/14 06:49:17.05 AbSbupmCP
RADEONは普及用チップでも倍精度が高速なのがいい
639:名無しさん＠お腹いっぱい。
12/11/14 07:58:59.38 vuLXlPiG0
mtyGPUがRadeonしか対応してないから、むしろゲフォ対応検索は(10桁では)貴重なんだが
640:名無しさん＠お腹いっぱい。
12/11/14 13:24:08.27 85Ooiiep0
>>638
マジレスすると倍精度演算が速いのは7970だけだしトリップ検索に倍精度演算の出番は無いぞ
641: ◆MERIKEN4.k
12/11/14 14:42:36.69 mDY3eRDW0
>>635
10桁トリップ検索は12桁よりかなり難しいので、実際どこまで速度を出せるかは
わかりませんけどね～ というか12桁検索の移植は正直うまくいきすぎでしたｗ
地道に取り組む予定なので、のんびり待っていて下さい。
642: ◆MERIKEN4.k
12/11/14 15:16:17.63 mDY3eRDW0
で、あれから色々試してみて、Bitslice DES用の一時変数はローカルメモリに
おかないと全く速度が出ないことが分かりました。ローカルメモリは
ワークグループ内で共有されるので、Bitslice DESを8個のスレッドで
並列処理するように書き換えてやりました。
その後、さらに性能を上げるためにAMD APP Profilerで解析してみました。
あんまり期待してなかったｗのですが、非常に使いやすいです。
で、気になっていたOccupancy Analysisを行なってみたら、
案の定ローカルメモリ(LDS)の使い過ぎであることが判明しました。
URLﾘﾝｸ(www.meriken2ch.com)
643:名無しさん＠お腹いっぱい。
12/11/14 15:31:41.20 AbSbupmCP
>>642
へぇ～
人目でボトルネックがLDSにあることが示されてる
凄いな
644: ◆MERIKEN4.k
12/11/14 15:51:16.97 mDY3eRDW0
同じ問題はCUDA版でも起きていたので思わず頭を抱えてしまったのですが、
ソースを眺めていたら解決方法を思いつきました。Bitslice DESの
一時変数は次の構造体にまとめられています。
> typedef struct {
> DES_Vector keys[56]; // 224 bytes
> DES_Vector dataBlocks[64]; // 256 bytes
> unsigned int dummy[1];
> } DESContext;
で、56bitのDESのキーが32個keys[]に収められているのですが、
これらのキーは実際にはほとんど同じです。
というわけで、キーの生成の方法を工夫してやれば、32個のキーの共通部分
51bitだけを保持して、残りは5bitのインデックス(0～31)から生成して
やればいいことに気づきました。
645: ◆MERIKEN4.k
12/11/14 16:01:44.57 mDY3eRDW0
これで使用するメモリの量は半分近くに減って、うまくいけば
CUDA版ともども10桁検索の速度が倍になることになります。
アルゴリズムはかなり複雑になりますが、試してみる価値は十分にあります。
hip2の話を聞いて、キーの生成方法にかなりの工夫の余地があることに
気づいたのは僥倖でしたｗ
646: ◆MERIKEN4.k
12/11/14 16:04:10.35 mDY3eRDW0
>>643
実際かなり便利です。CUDAのときはなんせExcelのスプレッドシートを
使わないとOccupancyのグラフが見れませんでしたからねｗ
647:名無しさん＠お腹いっぱい。
12/11/14 19:07:46.29 vuLXlPiG0
>>645
>速度が倍
うおおおおお！？頑張って下さい！
648:名無しさん＠お腹いっぱい。
12/11/14 19:20:22.61 dspeEFEK0
GTX670では470Mt/sくらいしか出ません。倍精度を使うわけでもないのになんでだろう。
649:前スレ927 ◆JouJaku.HzIz
12/11/14 20:16:05.54 HHBBdob70
ゲフォはさよならですかそうですか。
GTX480が何とか復活したので速度計測。
【GPU】GeForce GTX 480
【CPU】Xeon X5680@3.33GHz x2CPU
【OS】Win7Pro 64 SP1
【バージョン】0.07a6 CUI64
【トリップの種類】12桁
【1CUあたりのワークグループの数】N/A
【1WGあたりのワークアイテムの数】N/A
【その他のオプション】-c -g -x 128
【Display Driver】306.97
【10分間の平均速度】648.27M TPS
【GPUの平均速度】578.39M TPS
【CPUの平均速度】69.89M TPS
【GPU使用率】100%
【その他】"TEST/", HT off, GPU 92℃
Quadroをぶっちぎっているのですが・・・うるさい。とにかくうるさい。
常用は無理です。
650:名無しさん＠お腹いっぱい。
12/11/15 00:03:05.86 Gr7998EA0
>>642
これは便利そうですね。
>>644
DESは歴史もあり奥が深いですね。
>>648
レジスタ数がネックになって演算ユニットを使いきれていないのだと思います。
651:648
12/11/15 02:21:08.90 aNTlQCIF0
レジスタの仕様が違うのか。最適化しなおさないといけないわけね。
652: ◆MERIKEN4.k
12/11/15 03:09:09.94 dQ9rq2KX0
>>648
>>651
トリップ検索の速度は整数演算の性能に大きく影響されるんですけど、
GTX 600シリーズで使われているKeplerコアは残念ながら整数演算が
かなり遅いのです。この点は次の記事の「命令別スループット」の
項目で詳しく解説されています。
GTX680のグラフィック・GPGPU性能を調べる ≪ dokumaru
URLﾘﾝｸ(dokumaru.wordpress.com)
Keplerではゲームで使われる単精度演算以外はほとんど無視して
性能を稼いでいるので、GPGPU的にはかなり残念なことになっています。
653: ◆MERIKEN4.k
12/11/15 03:13:17.23 dQ9rq2KX0
>>649
580とあまり遜色のない速度が出ていますね。
自分の部屋ではGeForceが3枚24時間フル稼働してますｗ
CUDA版の開発も続けるので安心して下さい。
654: ◆MERIKEN4.k
12/11/15 09:58:39.20 dQ9rq2KX0
ここ数カ月のjohn-devでのOpenCLでのBitslice DESの実装についての
やり取りを追って見たのですが、なかなか面白かったです。
URLﾘﾝｸ(www.openwall.com)
URLﾘﾝｸ(www.openwall.com)
URLﾘﾝｸ(www.openwall.com)
現在John the Ripperは7970で20M c/sしか出せていないのですが、
OpenCLの実装を担当しているSayantan氏に対して、
Alexander氏が7970なら300M c/sは出るはずからボトルネックを探せ、
と言っているのが非常に興味深いです。
> Something like 300M c/s at DES-based crypt(3) on HD 7970. Maybe more
> than that if we hard-code E (generate or patch code on the fly).
URLﾘﾝｸ(www.openwall.com)
手元のTripcode Finderのコードは現在のJtRの実装より大分速いのですが、
それでも300M TPSには遠く及びません。レジスタ数にもまだ大分余裕があるし、
工夫の余地はいろいろあるのでしょう。実に奥が深いです。
655: ◆MERIKEN4.k
12/11/15 13:08:08.52 dQ9rq2KX0
Bitslice DESをマルチスレッド化したときにエンバクした模様。
結構な確率で間違ったトリップが出力されます。
CUDAと同じコードのはずなんですけど、barrier()がうまく動作してないの
かしらん。
しかしこれ、どうやってデバッグするんだろう…
656:ののたん ◆KiwamonoL.
12/11/15 14:35:25.04 et60Xlt20
>>655
昔ながらの printf でおｋ。
手段として美しくないのは嫌いとかなら知らん。
657: ◆MERIKEN4.k
12/11/15 14:47:56.84 dQ9rq2KX0
やっぱりそれしかないんですねorz
658:ののたん ◆KiwamonoL.
12/11/15 15:05:07.72 et60Xlt20
>>657
私が hip2 つくってた頃は printf すらなかったのに。
贅沢ね。
659: ◆MERIKEN4.k
12/11/15 15:15:33.54 dQ9rq2KX0
>>658
まあそりゃそうなんですけどね…
660: ◆MERIKEN4.k
12/11/15 15:18:14.38 dQ9rq2KX0
あ、原因分かったかも。CUDA版を書いてたときに適当だったところが
今になって問題になっているのかもしれません。
661: ◆MERIKEN4.k
12/11/15 15:27:59.89 dQ9rq2KX0
う～ん、違うな… もうちょっと全体的に腐ってる感じです。
662: ◆MERIKEN4.k
12/11/15 15:38:19.52 dQ9rq2KX0
まあいいや。マルチスレッド化の作業はまた明日やり直すことにしよっと。
663: ◆MERIKEN4.k
12/11/15 15:54:46.53 dQ9rq2KX0
コードをロールバックしたらちゃんと動作するようなのでやっぱり
マルチスレッド化が原因のようです。マルチスレッド化すると
速度が倍近くになるので次はなんとか成功させたいところです。
664:前スレ927 ◆JouJaku.HzIz
12/11/15 21:27:47.25 etuoVGYM0
480が余りにもうるさいので、590に交換。
【GPU】GeForce GTX 590
【CPU】Xeon X5680@3.33GHz x2CPU
【OS】Win7Pro 64 SP1
【バージョン】0.07a6 CUI64
【トリップの種類】12桁
【1CUあたりのワークグループの数】N/A
【1WGあたりのワークアイテムの数】N/A
【その他のオプション】-c -g -x 128
【Display Driver】306.97
【10分間の平均速度】978.15M TPS
【GPUの平均速度】922.60M TPS
【CPUの平均速度】55.55M TPS
【GPU使用率】0-100%
【その他】"TEST/", HT off, GPU 85℃
CPUの負荷変動がかなり激しいです。6コアx2が100％になることはまず無く、全コアが完全にストールすることも良く起こりました。
>>170 でもある程度CPUの負荷は変動しましたが、ここまで酷くは無かったです。
おまけにGPUもたまに完全にストールする始末。これは>>170 のマシンでは無かった。
タゲを増やすと負荷変動は落ち着きます。ここまで負荷がふらつく理由がさっぱり分かりません。
665:名無しさん＠お腹いっぱい。
12/11/16 01:59:34.00 QPHBSAhn0
電源容量が足りないんじゃ
666: ◆MERIKEN4.k
12/11/16 04:42:53.68 eP2LlovM0
OpenCLの10桁検索ですが、もうちょっと調べたらどうも移植した直後から
問題があったようです。APP Profilerがメモリリークを報告しているので
もうちょっと調べてみます。
667: ◆MERIKEN4.k
12/11/16 04:45:44.48 eP2LlovM0
>>664
温度に問題がないなら電源の可能性が高いですね。
電源は何を使われていますか?
668: ◆MERIKEN4.k
12/11/16 06:40:48.74 eP2LlovM0
どうやら問題はBitslice DESの処理そのものではなく
他の処理にある模様。ちゃんと出力をチェックするルーチンを
作りこんで、徹底的にテストするしかないようです。
やなよかんはしてたけど、やはり10桁検索は楽ではないですねえ。
669:前スレ927 ◆JouJaku.HzIz
12/11/16 23:39:50.31 SdQXCd/P0
電源が届くのを待ちきれなくて、無理矢理繋げて実行していました。
電力不足でこんな挙動をするとは初体験で全然知らず。お恥ずかしい限りです。
素直に電源届くまで待っています。
670: ◆MERIKEN4.k
12/11/17 09:52:44.80 Kz7friKn0
>>669
そりゃそこにカードがあれば試したくなりますよね。
その気持、わかりますｗ
電源が届いたらまたぜひ報告して下さい。
671: ◆MERIKEN4.k
12/11/17 10:07:54.25 Kz7friKn0
OpenCLの10桁検索の出力が腐っていた問題ですが、カーネルをすこしづつ削って
原因を探ったところ、結果を書き込む__globalの配列へのアクセスの前後に
barrier()を入れてやると問題が出なくなることが分かりました。
Bitslice DES用の一時変数を__privateに置いても直らなかったし、
CUDA版やOpenCLの12桁検索では全く問題がなかった部分なので、
AMDのOpenCLの実装のバグの可能性が非常に高いです。
AMDの実装は性能は出るのにいちいち造りが甘くて非常にもったいない
感じがします。ここらへんもCUDAのほうが任期がある理由なんでしょうねえ。
672: ◆MERIKEN4.k
12/11/17 12:33:38.78 Kz7friKn0
この件でコードをロールバックした時に気がついたのですが、
Bitslice DESの一時変数を__private空間においても割と速度が出ることが
わかりました。こっちのほうが__localよりもベクトル化しやすいので、
このまま__localを使わずに最適化をすすめることにします。
Bitslice DESの深さを32bitから128bitにして速度も4倍といきたい
ところですが…
673:名無しさん＠お腹いっぱい。
12/11/17 12:53:15.66 CDs2gwHh0
>>672
>ベクトル化
よく知らないのですが、GPUってベクトル演算なんですか……？
ベクトル化の意味は知っているのですが、なぜか「昔のスパコン」ってイメージが……ｗ
674:名無しさん＠お腹いっぱい。
12/11/17 17:05:57.82 RbPdKj5Y0
GPUはベクトル演算の極地だし、今のスパコンはほぼ全てベクトル演算ですが
675:名無しさん＠お腹いっぱい。
12/11/17 17:08:42.00 lv9DVzeD0
もの自体がベクタプロセッサの集合体
676: ◆MERIKEN4.k
12/11/18 01:33:23.65 7lmxdB8G0
>>673
そこがGPGPUの一番美味しいところですｗ
性能を引き出すのはなかなか難しいですけどね～
677: ◆MERIKEN4.k
12/11/18 01:59:41.80 7lmxdB8G0
あの後色々調べてみたんですけど、単純にDES_Vectorをuint2やuint4で置き換えて
やれば性能が出るというわけでもないようで、もうちょっと調べる必要が
あるみたいです。
あと、localなメモリに書き込んだ後は必ずbarrier()を呼び出さないと、
ちゃんとメモリ操作の結果が反映されないようです。おかしいなと思って
OpenCLの仕様書を見ると確かにこう書いてあります。
> The barrier function also queues a memory fence (reads and writes) to
> ensure correct ordering of memory operations to local or global memory.
URLﾘﾝｸ(www.khronos.org)
CUDAの場合は動機が必要なところで__syncthreads()を呼び出してやれば
後はなにも考えずに共有メモリとグローバルメモリに読み書きできたのですが、
どうも勝手が違うようです。
678: ◆MERIKEN4.k
12/11/18 12:44:30.93 7lmxdB8G0
OpenCLでの10桁検索の話の続きです。
>>545の案を実際に実装してメモリの使用量を半分に抑えることで、
速度を50%ほど向上させることができました。キーを動的に生成することに
よるペナルティが割と大きく2倍とはいきませんでしたが、
まあそれでもかなりの進歩です。Kernel Occupancyはこんな感じです。
URLﾘﾝｸ(www.meriken2ch.com)
ローカルメモリを使うと出力が化けまくるので、とりあえず
Bitslice DES用の一時変数はすべてレジスタ上においています。
このままレジスタの数を削ってOccupancyを上げてもいいし、
またローカルメモリに戻してみてもいいし、これでようやく先がすこし
見えてきた感じです。
679: ◆MERIKEN4.k
12/11/18 13:50:23.81 7lmxdB8G0
一応ローカルメモリに戻して速度を測ってみたのですが、
思ったほど速度は出ませんでした。というわけで
一時変数はこのまま__private空間においたまま
最適化をすすめることにします。
カーネルをなるべく簡単にして、キーの生成の準備をすべて
CPU側で行うことにします。
またレジスタの数を削る日々がはじまるお…
680: ◆MERIKEN4.k
12/11/18 18:56:32.90 7lmxdB8G0
あの後ちょこちょことカーネルをいじっていたんですけど、
適当なところにbarrier()を入れるとレジスタ数が減ったり
スピードが上がったりと不思議なことの連続でした。
色々実験してみるもんですね。こんなことは流石にマニュアルには
書いてあるわけないしｗ
681:名無しさん＠お腹いっぱい。
12/11/18 19:05:57.54 hHNMwY9r0
奇妙すぎる仕様だ……
682: ◆MERIKEN4.k
12/11/18 19:11:47.92 7lmxdB8G0
これは4日前に公開されたばかりのRadeon用のアセンブラです。
GCN ISA Assembler
URLﾘﾝｸ(devgurus.amd.com)
HetPas
URLﾘﾝｸ(realhet.wordpress.com)
GCNの命令セットについてなかなか面白いことが書いてありました。
こういうので最適化したらものすごい速度が出るんでしょうねえ。
JtRのAlexsander氏は動的にコードを生成しろなんて言ってたけど…
683: ◆MERIKEN4.k
12/11/19 16:30:19.77 FRx7NJvu0
>>681
まったく謎だらけですｗ カーネルアナライザを使えばもうちょっと詳しく
分かるんでしょうけど、goto文を使っているとエラーが出て動かないんですよね…
684: ◆MERIKEN4.k
12/11/19 16:36:30.60 FRx7NJvu0
気分転換で、前から欲しかったトリップの自動保存と自動検索実行の機能を
つけてみました。ブレーカーが落ちるたびにうんざりしながら検索君1号を
立ち上げなおしていたのですが、これで再起動もボタンを押すだけで済んで
検索結果が失われることもなくなりました。この機能は次の開発版から
利用できるようになる予定です。
685:名無しさん＠お腹いっぱい。
12/11/19 19:13:07.00 rYPhWgPq0
>>684
そんなにブレーカーが落ちる環境だったとは……
(開発以外)休んでも、いいのよ？
686:名無しさん＠お腹いっぱい。
12/11/20 07:59:44.51 8BgQYrDr0
海を越えると電気も日本みたいに高品質じゃないんだよ
687: ◆MERIKEN4.k
12/11/20 14:37:37.34 TS/gXHXx0
>>685-686
グラボ4枚で検索するようになってから急に落ちるようになりました。
ブレーカーがどうも古いみたいで、大家さんに言ったんですけど
ちっとも変えてくれません。まあでも消費電力に常に気を付けるように
したら大分ましになりました。
688: ◆MERIKEN4.k
12/11/20 14:48:46.83 TS/gXHXx0
レジスタ数を107から90まで頑張って減らしました。
目標の84まであともうちょっとなんですけど、
コンパイラの挙動が全く予想できないのでなかなか難しいです。
689: ◆MERIKEN4.k
12/11/21 12:12:44.49 eiBTExc50
カーネルアナライザが動かなくてカーネルのILとISAが見られなかったん
ですが、次のページを参考にしてようやく見れるようになりました。
カーネルをビルドするときにオプションで"-save-temps=[prefix]"と
していしてやればビルド時の一時ファイルが保存されます。
Looking for specific details of GPU_DUMP_DEVICE_KERNEL
URLﾘﾝｸ(devgurus.amd.com)
Kernelanalyzer refuses to compile anything
URLﾘﾝｸ(devgurus.amd.com)
最近のドライバではカーネルアナライザは動かないそうで…
AMDにはもうちょっと頑張ってもらいたいものです。
690: ◆MERIKEN4.k
12/11/23 14:34:10.31 ctiDd+QK0
あの後レジスタ数を減らすためにいろいろと試してみたのですが、
どうやっても90から更に減らすことはできませんでした。
どうも本気でレジスタ数の割付を最適化するためには
GCNのコードを直接書く以外ないようです。
仕方が無いので、割と時間がかかっているカーネルへの入出力の処理を
効率よく行うようにするための作業にとりかかりました。
とりあえずオーバーヘッドの大きいclEnqueueWriteBufferを1つにまとめたら、
なぜか未だに完全に消えてなかった出力が化けるバグが綺麗さっぱり
なくなりました。やれやれです。
691: ◆MERIKEN4.k
12/11/23 15:32:19.14 ctiDd+QK0
さっきjohn-devの11月のポストを読んでたんですけど、
何か問題が起きるとすぐにAMDのOpenCLドライバのバグが疑われてて
笑ってしまいましたｗ これは相当評判が悪いみたいですね…
> > All my accusations about driver bugs were... well they were based
> > on statistics, what can I say? :-)
> >
> > magnum
> We saw some craziness that justify our accusations.
URLﾘﾝｸ(www.openwall.com)
692: ◆MERIKEN4.k
12/11/23 15:44:36.45 ctiDd+QK0
なんにせよドライバのバグを華麗に避けつつOpenCLの10桁検索を使い物に
するには相当時間がかかりそうなので、とりあえず12桁検索のほうを
先に仕上げてしまうことにしました。今週末に次の開発版を公開する予定です。
693:名無しさん＠お腹いっぱい。
12/11/23 15:52:23.16 21daIx+z0
>>691
ドライバの完成度の問題ですか、厳しいですねえ・・・
694:名無しさん＠お腹いっぱい。
12/11/23 19:01:10.75 ixPLPIhe0
鳥屋は凄腕だな。
695: ◆MERIKEN4.k
12/11/23 20:06:39.28 ctiDd+QK0
鳥屋氏が凄腕なのは間違い無いですね。mtyのGPU版の速度は異常です。
ただCAL ILで書かれたmtyと同じ速度をOpenCLで出すのも無理な気がしますけどね～
JtRの20M c/sは論外にしても、Hashcatですら7970で79M c/sしか出せていない
ですからねえ。もうちょっとJtRのSayantan氏に頑張ってもらいたいものですけど、
メーリングリストのやり取りを見ている限りではとても期待できそうにありませんorz
696: ◆MERIKEN4.k
12/11/23 21:53:30.00 ctiDd+QK0
なにか10桁検索の参考にならないかと思ってJtRのソースを眺めていたら、
全然関係ない12桁検索の高速化のネタを見つけましたｗ
といってもハッシュ作成の際にbitselect()とrotate()を使うというだけの
話なんですけど、効果は抜群でOCした7970単体で1600M TPSを軽く超える
速度が出ています。いまだにこんなおいしいネタが転がっていたとは驚きです。
697:名無しさん＠お腹いっぱい。
12/11/23 23:24:15.72 6zLs77TA0
>>695-696
ということは12桁最高記録が300M/s以上増えることに！？
ところでmtyGPU版の10桁最高記録ってどれほどなのでしょう？
自分で(2chソースを)ググって分かったのは237M/s(1枚で)、枚数差しても～750M/sぐらいだったのですが……
698:ののたん ◆KiwamonoL.
12/11/23 23:33:28.75 34hnWziW0 BE:276537427-DIA(289888)
>>697
スレチだな。ｗ
こっちいけ。
URLﾘﾝｸ(yy43.60.kg)
699:名無しさん＠お腹いっぱい。
12/11/23 23:40:53.47 6zLs77TA0
>>698
後半荒らされ放題じゃないですか………‥
なるほど、少なくとも>>79で714M/sという記録が出ていたんですね。失礼しました
700:名無しさん＠お腹いっぱい。
12/11/23 23:46:32.74 NBSrQskj0
最大公約数的なプログラミングじゃなくて、自分の持ってるカードに絞ってゴリゴリ書いていけばいいんじゃないの？
その方が速度も出ると思うんだけど
701:名無しさん＠お腹いっぱい。
12/11/24 01:55:09.02 gl+pZhPfP
MERIKENさんってTOEIC満点とれる超人だったんですね・・・
702: ◆MERIKEN4.k
12/11/24 05:00:21.29 Npt3JpjG0
>>697
今でも3.5G TPSあたりなら堅いでしょう。いろいろ弾を仕込んでいる最中なので、
次に記録を狙うときには目標は4.5～5G TPSあたりになると思います。
703: ◆MERIKEN4.k
12/11/24 05:02:22.68 Npt3JpjG0
>>700
最大公約数的なプログラミングはとっくの昔に諦めて7970にターゲットを絞って
ますけど、それでもなかなか難しいです。
704: ◆MERIKEN4.k
12/11/24 05:04:48.04 Npt3JpjG0
>>701
私は大学からアメリカなのであれはいろんな意味で「おまけ」なのですｗ
705:名無しさん＠お腹いっぱい。
12/11/24 07:00:28.07 d1lnl00J0
1台のPCに積載できるGPUの量には限りがありますし、
そのうちサーバプログラム用意して検索条件の配布、検索結果の集計みたいな疎結合クラスタになりますん？
706: ◆MERIKEN4.k
12/11/24 08:25:58.55 Npt3JpjG0
>>705
そのうちそうなるでしょうねえ。スタンドアロンでの性能がちゃんと出るようになって
からということになるので相当先の話だと思いますけど…
707:名無しさん＠お腹いっぱい。
12/11/24 10:09:34.84 E4ne9Ljb0
>>705
トリップ検索クラスタ(物理)か……
GPUが絡まないと有り難みが薄いですねｗ
708:名無しさん＠お腹いっぱい。
12/11/24 11:01:21.01 NFbcJaLE0
トリップ検索p2pネットワークか‥胸熱
709:名無しさん＠お腹いっぱい。
12/11/24 19:19:01.23 TYsqoQfh0
>>698
スレチと、言ってるののたんに （はぁはぁ
710:名無しさん＠お腹いっぱい。
12/11/24 19:43:27.54 d1lnl00J0
>>708
個人でクラスタするのは有りだけど、
参加フリーでみんなの検索条件を合算するようになると生成されたトリップの判定にパワー食っちゃって……
711: ◆MERIKEN4.k
12/11/25 02:30:46.54 tDxdpeED0 BE:3591054296-2BP(12)
サーバーから検索条件をダウンロードしてみんなで12連とかのレアトリップを
探すというのも面白いかもしれませんねｗ
712: ◆MERIKEN4.k
12/11/25 02:38:56.68 tDxdpeED0
10桁検索のほうはAlexander氏の言っていた、動的にカーネルを書き換えて
DESのexpansion functionをソースに埋め込むという方法で以前に比べると
大分速くなりました。が、それと同時にドライバのバグによる出力が化ける問題が
再発生した模様。まったく地雷原を歩いているようです。
713: ◆MERIKEN4.k
12/11/25 09:08:59.10 tDxdpeED0
出力が化ける問題はなんとか解決できました。いや～、まいったまいった。
というわけで実行時のカーネルの書き換えでようやくHashcatとほぼ同じ速度が
出るようになりました。Tripcode FinderのCUDA版の10桁検索はHashcatよりも
ちょっと速いぐらいなので、もうそろそろ限界のような気もしないでも
ないです。あとはGCNのコードを手書きしてS-Boxを最適化して
レジスタ数を削るぐらいしか思いつきません。とりあえず10桁検索は
しばらく置いておいて、次の開発版を用意することにします。
714: ◆MERIKEN4.k
12/11/25 22:50:54.20 tDxdpeED0
2週間ぶりの開発版です。
MERIKEN's Tripcode Finder 0.07 Alpha 7
URLﾘﾝｸ(www.meriken2ch.com)
Alpha 6からの変更点は以下になります。
・OpenGLの12桁トリップ検索の高速化。
・検索スレッドと検索プロセスの数を指定するオプションの追加。
・検索の自動実行と検索結果の自動保存を行うオプションの追加。
・様々なバグの修正。
検索スレッドと検索プロセスの数を増やすことによって、
複数のGPUを使用する場合のGPUの使用率を増やすことができます。
おいしいです（＾ｑ＾）
715:名無しさん＠お腹いっぱい。
12/11/25 23:37:10.26 wZsqacQO0
おつおつ
回してみるべ
716:482
12/11/25 23:54:26.46 wZsqacQO0
この構成であるふぁ7
URLﾘﾝｸ(www.dotup.org)
717: ◆MERIKEN4.k
12/11/25 23:57:46.50 tDxdpeED0
Alpha 7用の新しい報告用のテンプレです。
【GPU】
【CPU】
【OS】
【バージョン】MERIKEN's Tripcode Finder 0.07 Alpha 7
【トリップの種類】12桁・10桁
【1SMあたりのブロックの数(CUDA)】
【1CUあたりのワークアイテムの数(OpenCL)】
【1WGあたりのワークアイテムの数(OpenCL)】
【1GPUあたりの検索プロセスの数(OpenCL)】
【1検索プロセスあたりの検索スレッドの数(OpenCL)】
【その他のオプション】
【Display Driver】
【10分間の平均速度】 tripcodes/s
【GPUの平均速度】 tripcodes/s
【CPUの平均速度】 tripcodes/s
【GPUの使用率】
【GPUの温度】
【その他】
718: ◆MERIKEN4.k
12/11/25 23:58:37.50 tDxdpeED0
>>715
ぜひ色々試してみてくださいｗ
719: ◆MERIKEN4.k
12/11/26 00:02:11.37 ICZhOGh80
>>716
これは1枚ですか? かなり出てますね～
720:名無しさん＠お腹いっぱい。
12/11/26 00:07:18.79 SipgVuQ10
書き忘れー
解凍したまんまで
GPUの温度は室温20度で41度まで上がった、負荷は100%
水冷だしこんなもんだね、ゲームだと36度くらいしか上がんないからいかにGPUが仕事してるかわかるw
721:名無しさん＠お腹いっぱい。
12/11/26 00:08:15.66 SipgVuQ10
>>719
ですよー、リファの7970
722:名無しさん＠お腹いっぱい。
12/11/26 00:12:39.54 AQz+o+st0
唐突だけどコマンドラインオプションの私的まとめ(☆はデフォルトでは自動設定される項目):
--redirection ?
-f [inputfile] 入力ファイル名
-r [inputfile] 入力ファイル名(正規表現)
-o [outputfile] 出力ファイル名
-l [length] 検索するトリップ長(12 or 10)
-g 検索にGPUを使用 (デフォルト)
-d [device] CUDAデバイス番号(0～) (デフォルトは全て使用)
-x [block/SM] ブロック/SM(CUDA) ☆
-y [workgroup] ワークグループ/CU(OpenCL) ☆
-z [workitem] ワークアイテム/WG(OpenCL) ☆
※workgroup mod workitem=0、workitem mod 8=0とすること
-c 検索にCPUを使用(-gと併用可)
-t [threads] CPUにおける検索スレッドの数 ☆
-a [threads] 1つのAMDのGPUに対する検索スレッドの数(OpenCL) ☆(～0.07Alpha6)
1検索プロセスあたりの検索スレッドの数(OpenCL) ☆(0.07Alpha7～)
-b [processes] 1GPUあたりの検索プロセスの数(0.07Alpha7～)
-m MutexForMERIKENsTripcodeFinder-4648 GUI版とCUI版が通信するときに使うおまじない(～0.07Alpha6)
-m MutexForMER GUI版とCUI版が通信するときに使うおまじない(0.07Alpha7～)
-i 2ちゃんねるで直接使用できないトリップを16進形式で出力
-w 検索スピードの急激な低下を警告
723: ◆MERIKEN4.k
12/11/26 00:18:40.01 ICZhOGh80
>>722
あ、-yは「ワークグループ」じゃなくて「ワークアイテム」です。
最初に書いたときに間違えちゃったんですよね～
724: ◆MERIKEN4.k
12/11/26 00:22:36.54 ICZhOGh80
>>721
う～ん、水冷は素晴らしいですね。空冷での温度を見慣れていると
別世界のようですｗ
725:名無しさん＠お腹いっぱい。
12/11/26 00:34:50.17 AQz+o+st0
>>723
つまりこうですね、分かります。
>-y [workitem1]ワークグループ/CU(OpenCL)(デフォルトは自動設定)
>-z [workitem2]ワークアイテム/WG(OpenCL)(デフォルトは自動設定)
>※workitem1 mod workitem2=0、workitem2 mod 8=0とすること
ところで--redirectionって何をリダイレクトしているんですか？
726:名無しさん＠お腹いっぱい。
12/11/26 00:37:16.39 AQz+o+st0
俺おっちょこちょいの素質あるのかな……
>-y [workitem1]ワークアイテム/CU(OpenCL)(デフォルトは自動設定)
>-z [workitem2]ワークアイテム/WG(OpenCL)(デフォルトは自動設定)
>※workitem1 mod workitem2=0、workitem2 mod 8=0とすること
次のVerからはREADMEに訂正が必要なようですね……＞MERIKENさん
727: ◆MERIKEN4.k
12/11/26 00:43:33.42 ICZhOGh80
>>725-726
リダイレクトしているのは標準出力です。
訂正はもう入ってますよ。
728:名無しさん＠お腹いっぱい。
12/11/26 00:43:54.12 SipgVuQ10
【GPU】Xeon E5-2687W×2
【CPU】HD6990×2
【OS】Windows８ Pro
【バージョン】MERIKEN's Tripcode Finder 0.07 Alpha 7
【トリップの種類】12桁
【1SMあたりのブロックの数(CUDA)】-
【1CUあたりのワークアイテムの数(OpenCL)】解凍時のまま
【1WGあたりのワークアイテムの数(OpenCL)】解凍時のまま
【1GPUあたりの検索プロセスの数(OpenCL)】解凍時のまま
【1検索プロセスあたりの検索スレッドの数(OpenCL)】解凍時のまま
【その他のオプション】-
【Display Driver】Catalyst12.11β
【5分間の平均速度】 4816.85tripcodes/s
【GPUの平均速度】 4711.99tripcodes/s
【CPUの平均速度】 104.86tripcodes/s
【GPUの使用率】100%
【GPUの温度】一番高いコアで46℃
【その他】GPUはTDP450Wモード定格
URLﾘﾝｸ(www.dotup.org)
これはもしやメインも仕事してくれるのではと思ったら案の定
時間ないんでどちらも5分でスマヌ
729: ◆MERIKEN4.k
12/11/26 01:03:17.64 ICZhOGh80
>>728
これは最高速の記録ですね。素晴らしいです。
私も次に記録を狙うときにはもうちょっと弾を揃えないと…
730:名無しさん＠お腹いっぱい。
12/11/26 01:12:41.03 AQz+o+st0
>>728
脳内での 最 速 記 録 が 塗 り 替 え ら れ た 瞬間であった
期待できないけどノートで回してくるー
731:名無しさん＠お腹いっぱい。
12/11/26 01:20:20.30 SipgVuQ10
因みにこれで1160W前後の消費電力
732:名無しさん＠お腹いっぱい。
12/11/26 01:30:34.82 Scm6xI9q0
>>569 です。Alpha7公開お疲れ様です。
【GPU】SAPPHIRE VAPOR-X HD5770 1G (OC: GPU 960MHz MEM 1265MHz)
【CPU】Intel Core i7-3770(無印)
【OS】Microsoft Windows 7 64bit SP1
【バージョン】MERIKEN's Tripcode Finder 0.07 Alpha 7
【トリップの種類】12桁
【1CUあたりのワークアイテムの数(OpenCL)】3200
【1WGあたりのワークアイテムの数(OpenCL)】64
【1GPUあたりの検索プロセスの数(OpenCL)】1
【その他のオプション】
【Display Driver】Catalyst 12.10
【10分間の平均速度】586.05M tripcodes/s
【GPUの平均速度】550.44M tripcodes/s
【CPUの平均速度】35.62M tripcodes/s
【GPU使用率】99%
【GPUの温度】72℃ （室温22℃）
【その他】テスト時間10分08秒、7完1タゲ
733: ◆MERIKEN4.k
12/11/26 01:30:58.70 ICZhOGh80
>>731
450W x 2 + αですか。こりゃすごいｗ
734: ◆MERIKEN4.k
12/11/26 01:34:21.82 ICZhOGh80
>>732
5770でもかなり速度が出てますね。
今回はかなり内部をいじったので、ちゃんと動いているようでほっとしました。
735:名無しさん＠お腹いっぱい。
12/11/26 02:46:46.01 qf13XQqh0
【GPU】HD7970 CFX 2GPUｓ ＠1150MHｚ
【CPU】FX-8350 ＠5GHｚ
【OS】Windows7 64bit
【バージョン】MERIKEN's Tripcode Finder 0.07 Alpha 7
【トリップの種類】12桁
【1SMあたりのブロックの数(CUDA)】
【1CUあたりのワークアイテムの数(OpenCL)】960
【1WGあたりのワークアイテムの数(OpenCL)】64
【1GPUあたりの検索プロセスの数(OpenCL)】default
【1検索プロセスあたりの検索スレッドの数(OpenCL)】default
【その他のオプション】-g -c -t 6
【Display Driver】Catalyst12.11 beta6
【10分間の平均速度】5277.77 tripcodes/s
【GPUの平均速度】5243.39 tripcodes/s
【CPUの平均速度】34.39 tripcodes/s
【GPUの使用率】99%
【GPUの温度】76℃
【その他】7完1タゲ
効率が上がったためか-t 6で回したら強制シャットダウン、恐らく冷却不足か電源容量不足
とりま、ぬるい設定で解凍したまま
※今までは検索始めるとマウスカーソルがカクカクになり、USB音源を見失っていましたが、そういった現象はなくなりました
URLﾘﾝｸ(www.rupan.net)
736:名無しさん＠お腹いっぱい。
12/11/26 02:47:44.61 qf13XQqh0
あ、-t 4 の間違いです
737:名無しさん＠お腹いっぱい。
12/11/26 03:26:15.30 AQz+o+st0
ノーパソから計測実験。デスクトップでグラボぶん回すのと比べると雑魚レベルだが許してくれ。
【GPU】NVIDIA GeForce 610M(、Intel HD Graphics 4000)
【CPU】Intel Core i5-3210M
【OS】Windows Vista Home Ultimate SP1 64bit
【その他のオプション】-g -c -l 10か-g -c -l 12での計測(速度が安定した時点で記録)
【Display Driver】見方を教えて下さい……
↑の条件で、ソフトのVerと桁数を変更しながら計算するとこうなった↓
0.07Alpha6 0.07Alpha6 0.07Alpha7 0.07Alpha7
10桁 12桁 10桁 12桁
----------------------------------------------
使用不可 160 使用不可 128 ←blocks/SM
使用不可 64 使用不可 使用不可 ←items/CU
使用不可 32 使用不可 使用不可 ←items/WG
4 2 3 3 ←CPU演算スレッド数
使用不可 48.9M/s 3.54M/s 48.96M/s ←速度(CUDA)
使用不可 4.9M/s 使用不可 使用不可 ←速度(OpenCL)
使用不可 9.66M/s 4.03M/s 10.91M/s ←速度(CPU)
5.27/s 63.43M/s 7.57M/s 59.87M/s ←合計速度
----------------------------------------------
確かに改良は効いているが、な ぜ ア ホ の 子 を 外 し た し
738: ◆MERIKEN4.k
12/11/26 09:08:45.23 ICZhOGh80
>>735
これはすごい数字ですねえ。いくらなんでも速すぎだろうと思って
Catalyst 12.11 Beta 8を試してみたら、うちの7970 1枚でも2497M TPS
出てて吹きましたｗ 12.9 Betaではここまでのスピードは出なかったので、
ここ2ヶ月でAMDのドライバにかなり手が入ってますね～
739:名無しさん＠お腹いっぱい。
12/11/26 10:37:04.72 7cZ4LG5Ni
7970の4wayやれば10Gか…
コンセントの端子が熱くなるな
740:名無しさん＠お腹いっぱい。
12/11/26 18:30:45.89 AQz+o+st0
公式サイト(URLﾘﾝｸ(www.meriken2ch.com))とか見ていると
OpenGLとOpenCLが脳内でごっちゃになりそうなのでまとめ：
OpenGL……シリコングラフィックスが開発していたクロスプラットフォームな3DグラフィックスのAPI。
ハードウェアに近い低水準な機能も使えるので高速だが、文字列描画が苦手。
GPGPUの利用法は、OpenCLよりもグラフィックス寄り。
OpenCL……アップルのKhronos Groupが開発した、クロスプラットフォームな並列コンピューティング用のAPI。
要するに、「CPUやGPUなどの計算資源を、並列演算用にまとめて扱えるようにするよ！」
といったもの。GPGPUの利用法は、OpenGLよりは演算寄り。
741: ◆MERIKEN4.k
12/11/26 19:27:18.93 ICZhOGh80
>>740
あ、あれはOpenCLの間違いで、OpenGLは一切関係ないですｗ
ご自分用のまとめはここに書き込まないでいただけると有難いです。
742: ◆MERIKEN4.k
12/11/26 19:36:47.88 ICZhOGh80
>>737
OpenCL以外の検索ルーチンはいじってないので速度は変わっていないはずです。
Intelのはドライバのバージョンによってアプリケーションが落ちるろいう報告が
あったのでやむなしです。
743:名無しさん＠お腹いっぱい。
12/11/26 19:56:04.00 AQz+o+st0
>>741
了解しました。
>>742
そうだったんですか……。チェックボックス対応でも、というのは無茶でしょうか。
10桁の演算速度が上がっているのは確実な気がするのですが、
単に自環境ではAlpha6でGPU演算が使えなかっただけ(デバイスが対応していない)
なのかもしれません。次買うのはRadeonGPU搭載PCにするかな…‥
744:名無しさん＠お腹いっぱい。
12/11/26 20:09:11.33 gXr4FeWs0
>>740
geforce君はもう書き込まないでくれるかな？
745: ◆MERIKEN4.k
12/11/26 20:44:41.34 ICZhOGh80
>>743
Intelのはドライバの出来がイマイチで性能が全く出ないのに
メンテの手間だけかかって、おいしいところが全くないんですよね。
Intel対応はXeon Phiが消費者向けに発売されたら考えますｗ
746:名無しさん＠お腹いっぱい。
12/11/26 20:49:46.82 AQz+o+st0
>>745
確かに、グラボが出す速度を考えたらIntelのは誤差の範囲ですよねｗ
もうその件については触れないことにします。回答ありがとうございました。
747: ◆MERIKEN4.k
12/11/26 21:13:07.19 ICZhOGh80
ドライバといえば、Catalystの新しいβ版で10桁検索を試してみたら、
速度が1/3になっていましたorz CUDAでもそうでしたけど、
GPGPUは開発環境やドライバによってアプリケーションの性能が
乱高下する傾向がありますねえ。ドライバの次のバージョンアップで
直っているといいんですけど…
748: ◆MERIKEN4.k
12/11/27 01:57:27.60 qHzcgcY70
新しいAMDのドライバで12桁トリップ検索のプロファイリングを行って見たのですが、
ベクターレジスタ(VGPR)の数が40まで減っていて、Occupancyが10から60にまで
上がっていました。どうりで検索速度が上がっているわけです。
どうやらAMDのコンパイラの最適化のアルゴリズムが、命令の数を増やしてでもレジスタ数を
減らすことを優先するものに変更されているようで、それが12桁の場合はうまく働いたけど
10桁の場合は完全に裏目に出ている、ということらしいです。やっぱり本気で10桁トリップ検索で
性能を出そうと思ったらILかGCNのコードを自分で書くしかないみたいですが、とりあえず
以前のドライバでOpenCLバイナリを生成して、実行時にはそれを使うように変えておくことにします。
749: ◆MERIKEN4.k
12/11/27 03:59:49.95 qHzcgcY70
AMDのOpenCLドライバをAMD APP 2.7のものにロールバックしたら
ようやく10桁検索の速度が元に戻りました。次のファイルは
ドライバのアンインストールでは削除されずに直接手で削除する
必要がありました。
SlotMaximizerBe.dll
SlotMaximizerAg.dll
amdocl.dll
OpenVideo.dll
OVDecode.dll
これがわかるまでエラく手間取りましたが、これでようやくOpenCLバイナリの
作成に取り掛かれます。
750: ◆MERIKEN4.k
12/11/27 06:43:01.11 qHzcgcY70
10桁トリップ検索のコードですが、なんとCatalyst 12.8以前のドライバでは
出力が化けることが判明しました。ドライバのバクにしても
いくらなんでもひどすぎるorz
751: ◆MERIKEN4.k
12/11/27 19:01:37.31 qHzcgcY70
>>746
手間がかからないならサポートしてもいいんですけど、テストの量が倍以上に
なりますからねえ。残念です。
752: ◆MERIKEN4.k
12/11/27 19:18:07.56 qHzcgcY70
>>739
これ2枚積めば10G TPS超できそうです。TDP 500Wの化物だけど、
普通に検索君1号に2枚収まりそうなんだけど、流石に20万は払えないよな…
あぶないあぶないｗ
Radeon最上位のデュアルGPU版「HD 7990」発売、重さを支える支柱付き カードは弩級
URLﾘﾝｸ(akiba-pc.watch.impress.co.jp)
Club 3D Radeon HD 7990 Dual GPU
URLﾘﾝｸ(www.club-3d.com)
753:名無しさん＠お腹いっぱい。
12/11/27 19:19:19.94 hGvQPEA10
>>752
そんなに電源虐めたいかｗｗ
754: ◆MERIKEN4.k
12/11/27 19:34:03.64 qHzcgcY70
>>753
そこにハードウエアがあれば限界まで性能を出したくなるのが
男のさがというものですｗ
755: ◆supernova.rT
12/11/27 19:36:20.84 3f/efQ6N0 BE:5355599279-DIA(123422)
10桁酉が割られる日も近いな…ｺﾞｸﾘ
756: ◆MERIKEN4.k
12/11/27 19:38:49.66 qHzcgcY70
10桁検索ですけど、crypt()のseedの値に基づいてカーネルを動的に
書き換えていたことをすっかり忘れていましたｗ
これって実行時にOpenCLバイナリを書き換えるか、seedの数だけバイナリを
用意しなきゃいけないってことだよな…
757:名無しさん＠お腹いっぱい。
12/11/27 19:59:40.97 oOatxZVu0
>>756
最適化スゲェ……
でも、10桁のシード(ソルト)って確か2バイト分(最大256^2=65536通り)あるんじゃ
758: ◆MERIKEN4.k
12/11/27 20:35:03.96 qHzcgcY70
>>757
実際には2chの仕様のせいで65^2=4225通りなんですけど、
それでも結構な数です。とりあえず実験的に作ってみますけど、
さすがにこれを配布パッケージに含めるのは考えちゃいますねｗ
759:名無しさん＠お腹いっぱい。
12/11/27 20:39:00.84 oOatxZVu0
>>758
単純に考えて、3.5MB×2×4225≒30GBかぁ……
動的書き換えでお願いします(切望)
760: ◆MERIKEN4.k
12/11/27 21:00:31.98 qHzcgcY70
>>759
書き換えが必要なのはOpenCLのカーネルのバイナリだけなので
そこまでひどくはならないですｗ せいぜい数十MBのオーダーでしょう。
圧縮がかなり効くはずなので配布パッケージ自体はそこまで大きくならない
はずですけど、こればっかりは試してみないとわかりません。
761: ◆MERIKEN4.k
12/11/27 22:03:05.43 qHzcgcY70
新しいドライバで10桁検索をプロファイリングしてみたのですが、
SALBusyが80.84%なのに比べてVALUBusyが28.91%と妙に低いのに
気づきました。MemUnitBusyが66.81%とかなり高いのも気になります。
これは実際にS-Boxで費やされている実行時間は全体の3割程度ということで、
かなり効率が悪いことになります。ちょっとドライバのバージョンを落として
比較してみます。
762:名無しさん＠お腹いっぱい。
12/11/27 22:41:52.20 X0Buxy760
>>747-748
バージョンによって最適化がかなり違うのですか、面倒ですねえ。
>>756
saltに応じてカーネルの動的書き換えとかできるのですか。
できるにしても実際にやるのが凄いですｗ
>>758
crypt(3)の仕様で64^2=4096通りではないのですか？
763: ◆MERIKEN4.k
12/11/27 23:10:13.31 qHzcgcY70
>>762
あれれ、そうでしたっけ? もうちょっと調べてみます。
764: ◆MERIKEN4.k
12/11/27 23:17:19.40 qHzcgcY70
>>762
CUDAでも開発環境のバージョンによってかなり速度差が出てましたけど、
OpenCLではドライバのバージョンで違ってくるので頭が痛いです。
HashcatはカーネルをLLVM IRで配布してるみたいですけど、
似たようなことをしたほうがいいのかもしれません。
765: ◆MERIKEN4.k
12/11/28 00:08:01.45 v1ASRvbE0
Catalyst 12.9 Betaに戻してみたら、こんな感じでした。
VALUBusy: 28.91% -> 36.15%
SALUBusy: 80.84% -> 113.88%
MemUnitBusy: 66.81% -> 63.67%
VALUBusyがちょっと上がっただけで速度は3倍になってるので、
ベクターユニットが遊んでいるせいで7970は相当余力を残している
ことになります。かなりの性能向上が期待できそうなので、
OpenCLの実装が一段落したら、自分でGCNのコードをいじってみようかな…
766:名無しさん＠お腹いっぱい。
12/11/28 00:12:07.31 JI44h4XK0
>>765
> SALUBusy: 80.84% -> 113.88%
100%越えってどゆことー？
767: ◆MERIKEN4.k
12/11/28 05:40:16.34 v1ASRvbE0
>>766
それはまったく謎ですｗ
768: ◆MERIKEN4.k
12/11/28 06:28:09.26 v1ASRvbE0
>>762
調べてみましたけど、2chのトリップ生成でのsaltの扱いはcrypt (3)の仕様から
ずれているので、やっぱり4225通りであってました。
2ch トリップ仕様
URLﾘﾝｸ(sourceforge.jp)
最初に10桁トリップ検索の実装をした時にも思ったことですが、
見れば見るほどひどい仕様ですｗ
769:名無しさん＠お腹いっぱい。
12/11/28 08:12:30.61 ttD8PkvV0
>>768
面妖な！
……ひょっとして10桁検索がどうしても遅くなるのはここにも理由があるんじゃ
770: ◆MERIKEN4.k
12/11/28 19:13:59.54 v1ASRvbE0
7990ですけど別の会社からも発売されてますね。
PowerColor AX7990 6GBD5-2DHJ Radeon HD 7990 6GB
URLﾘﾝｸ(www.newegg.com)
一枚で$899なので>>752のカードより大分安いですが、がまんがまん…
771: ◆MERIKEN4.k
12/11/28 19:29:28.10 v1ASRvbE0
>>769
10桁検索が遅くなるのはBitslice DESでメモリへのランダムアクセスが
大量に発生するのが大きいです。こればっかりは仕方ないですね。
772: ◆MERIKEN4.k
12/11/28 20:09:39.45 v1ASRvbE0
isaファイルを出力させてGCNのコードを眺めてたんですが、
register spillsが発生している模様。"ScratchSize = 140;"なる記述が
isaファイルにありました。道理でなかなか速度が出ないわけです。
プロファイラのScratchRegsの欄がNAになってたので完全に油断してました。
NAはnot applicableじゃなくてnot availableの略だったのね…
なんにせよこれでMemUnitBusyやMemUnitStalledが高いのも、VALUBusyが
低いのも説明がつきます。これってCUDAのときみたいにS-Boxを書き換えたら
なんとかなるのかしらん。
773: ◆MERIKEN4.k
12/11/29 00:31:56.88 VD1AV4Df0
S-Boxとおぼしき場所に倫理演算の命令に混じってbuffer_store_dwordと
s_buffer_load_dwordx4という命令が大量にあったので、
たぶんこれが速度が出ない原因なんでしょう。
ちょっとすっきりしたけど、これってコンパイラのレジスタの割付が
全然うまく行っていないということですよね。やれやれです。
774: ◆MERIKEN4.k
12/11/29 00:36:21.42 VD1AV4Df0
倫理演算じゃなくて論理演算でした。
775: ◆MERIKEN4.k
12/12/01 23:53:39.30 g8/dTHR/0
S-Boxの数を変えてISAファイルを調べてみたら、コンパイラがレジスタを
きちんと再利用していないことが判明。
S-Boxes: 1
Kernel occupancy: 10
NumVgprs = 180;
ScratchSize = 0;
S-Boxes: 7
Kernel occupancy: 10
NumVgprs = 239;
ScratchSize = 0;
S-Boxes: 8
Kernel occupancy: 20
NumVgprs = 105;
ScratchSize = 140;
register spillsが起きるとメモリアクセスが枷になって遅くなるし、
起きなければoccupancyが半分になるしでなかなかうまく行きません。
Bitslice DESに必要なレジスタの数は64 + 17 = 81ぐらいなので、
180～245というのはいくらなんでも多すぎです。
CUDAだったら直接PTXのコードを書けばいいんだけど、OpenCLだと
そういうわけにもいかないので実に難しいです。使用するレジスタの数も
CUDAみたいにコンパイル時に指定できたらいいんですけどねえ。
776:名無しさん＠お腹いっぱい。
12/12/02 13:44:02.57 E9WK095v0
駄目元でAMDのフォーラムに報告してみるとか
777:名無しさん＠お腹いっぱい。
12/12/03 19:33:58.87 VDyT7kE/0
URLﾘﾝｸ(www.meriken2ch.com)
そんなにPC酷使したいならこれで12桁の酉でも探してろ
778:名無しさん＠お腹いっぱい。
12/12/03 19:34:41.74 VDyT7kE/0
すまん間違えたｗちゃんと生贄連れてくるわ
779:名無しさん＠お腹いっぱい。
12/12/03 19:36:00.76 Q+462s2K0
よりによってこのスレに誤爆ｗｗｗ
780:名無しさん＠お腹いっぱい。
12/12/04 14:07:03.07 OIUiTKsY0
Catalyst 12.11 Beta11が出たな
781: ◆MERIKEN4.k
12/12/05 13:40:50.64 YhHPYAwa0
>>776
う～ん、どうなんでしょうねえ。レジスタ割り付けを改善すれば
速度が上がるのは自明なので、特に報告するまでもない気もします。
実際12桁検索は倍近く速くなったので、今後に期待といったところです。
782: ◆MERIKEN4.k
12/12/05 13:42:17.40 YhHPYAwa0
>>777-779
ぜひ活きのいいのをお願いしますｗ
783: ◆MERIKEN4.k
12/12/05 13:44:23.29 YhHPYAwa0
>>780
かなり頻繁に更新してますね。現在ダウンロード中です。
784: ◆MERIKEN4.k
12/12/05 13:58:19.89 YhHPYAwa0
>>287のPCIe用の延長ケーブルを使って、空冷用のスペースを
確保しつつ検索君1号にグラボを3枚積めることを確認しました。
見た目は最悪ｗですが、ちゃんと動いているので結果オーライです。
弾も色々揃えたので、帰省するまでに最高速の記録を更新できるかも
しれません。
785:名無しさん＠お腹いっぱい。
12/12/05 18:55:54.39 jmQ8Rzeo0
>>784
6G級あるか!?
786: ◆MERIKEN4.k
12/12/06 14:44:56.38 LfRKvPte0
>>785
さあ、どうでしょうねえ… (￣ー￣)ﾆﾔﾘ
787: ◆MERIKEN4.k
12/12/06 14:57:33.42 LfRKvPte0
ターゲットが長くなるとヒットするまでの平均時間をいまいち正確に
出せなかった問題ですが、次のライブラリを使うことで解決できることが
わかりました。
Multiple Precision Integers and Rationals
URLﾘﾝｸ(www.mpir.org)
Visual C++だとlong doubleがdoubleと同じ精度なので困ってたのですが、
これなら全く問題ないでしょう。
788: ◆MERIKEN4.k
12/12/06 16:26:33.08 LfRKvPte0
MPIRのビルドはあっさり成功して、ちゃんとTripcode Finderに
リンクすることができました。サンプルで2の120乗を計算してみましたが、
ちゃんと正しい結果が出ています。このライブラリには分数計算のルーチンも
含まれているので、非常に正確に確率計算ができるはずです。わくわく…
789: ◆MERIKEN4.k
12/12/06 17:18:12.37 LfRKvPte0
おっと、間違えた。サンプルで計算したのは2の1920乗でした。
このライブラリ、logが計算出来ないから使うの結構面倒そうだな。
どうしたものか…
790:名無しさん＠お腹いっぱい。
12/12/06 20:46:13.83 nOh2Wtf90
>>787-789
>ヒットするまでの平均時間をいまいち正確に出せなかった
そうだったのですか！？
ひょっとして有効桁数が2桁表示なのはそのせい……？
↓ところで、トリップ確率を計算するソフトを作っていたのですが、
URLﾘﾝｸ(up3.viploader.net)
桁数が変わる「.」とかが入った時や準X連な時の正確な組み合わせ数を計算するのが難しいのデス……
どういった計算アルゴリズムで出しているのですか？大雑把でいいので教えて下さい！
791: ◆MERIKEN4.k
12/12/07 08:23:57.13 G1/OJRD00
>>790
基本的な流れは以下のとおりです。
(1) 正規表現のパターンを位置と固定長文字列の組み合わせに展開する。
(2) 各組み合わせごとの確率を計算する。
(3) (2)の確率の合計を求める。
注意しなければならないのは、各文字が特定の位置に出現する確率は
通常は1/64ですが、特殊文字の場合は違うということです。
例えば"."と"[:digit:]"がヒットする確率はそれぞれ64/64と10/64と
しておかなければ正確な結果が出ません。
具体的な例を挙げると、12桁トリップ検索における"^test./"の出現確率は
p = (1/64)*(1/64)*(1/64)*(1/64)*(64/64)*(1/64)
となります。
また、位置指定をしていない"/test[:digit:]/"の場合、出現位置が
0～5の6通りなので、
p = (1/64)*(1/64)*(1/64)*(1/64)*(1/64)*(10/64)*(1/64)*6
になります。
792: ◆MERIKEN4.k
12/12/07 08:32:34.86 G1/OJRD00
MPIRの分数の型であるmpq_tを使って確率計算をすると、
遅くて使いものにならないことが判明orz
厳密にしすぎるのも考えものですね…
仕方ないので浮動小数点数の型のmpf_tを使うことにします。
任意の精度を指定できるのでこれで十分でしょう。
793: ◆MERIKEN4.k
12/12/07 10:59:12.72 G1/OJRD00
MPIRを使ってヒットまでの時間を予測するルーチンを書き直しましたが、
結局doubleを使った元のルーチンに比べて数パーセント精度が
向上しただけでした。元のルーチンもわりと正確だったということですが、
前からだいぶ気になっていた部分だったのでまあ良しとします。
794: ◆MERIKEN4.k
12/12/07 20:35:40.18 G1/OJRD00 BE:3192048386-2BP(12)
>>790
あ、あと書き忘れてたけど、準x連の場合は該当する文字が出現する確率は
大文字と小文字をあわせて2/64になります。例えば"^[Aa]*$"のような
準12連が出現する確率は、
p = pow(2.0/64.0, 12)
となります。
795:名無しさん＠お腹いっぱい。
12/12/07 22:19:25.47 1HdVOJHZ0
>>791
>>基本的な流れ
これだと、あるパターンが複数行で当てはまる際重複して数えてしまうような……
「当てはまる全パターン」を正確に計算するのはカナリ厳しいことがよく分かりました
>位置と固定長文字列の組み合わせ
ほほう、なるほど。パーサを見直せば出来そうです
ただ、実際にトリップ検索スレに出てくる案件を見る限りでは、
「.」とか「*」とかとかを使う機会は無さそうですね……
>>794
あーいや、こちらが言うところの「準X連」とは、正規表現では「*[Aa][Aa][Aa]*」みたいなもののことです
(これが「純X連」になると、「*AAA*」となります)
もちろん「^[Aa][Aa][Aa]*」から「*[Aa][Aa][Aa]$」まで虱潰しに出して合計してみてもいいのですが、
そうすると「BGCAAAAAAfgt」みたいなものが重複ヒットしてしまうようで……
足し引きしてなんとかすることにします
確率計算での参考:
URLﾘﾝｸ(www.geocities.jp)
796: ◆MERIKEN4.k
12/12/08 03:04:57.17 vyeW7s150
>>795
> これだと、あるパターンが複数行で当てはまる際重複して数えてしまうような……
この問題はパターンを固定文字列に展開したあとで重複するものを
取り除くことでほとんどの場合回避できます。Tripcode Finderでは
qsort()とuniq()の組み合わせで対処しています。
> あーいや、こちらが言うところの「準X連」とは、正規表現では
> 「*[Aa][Aa][Aa]*」みたいなもののことです
正規表現では"*"は先頭に来ないのでいまいちよくわからないですが、
"^[^Aa]*[Aa][Aa][Aa][^Aa]*$"のことでしょうか。
> もちろん「^[Aa][Aa][Aa]*」から「*[Aa][Aa][Aa]$」まで虱潰しに出して合計してみてもいいのですが、
> そうすると「BGCAAAAAAfgt」みたいなものが重複ヒットしてしまうようで……
確かにそうなんですけど、実際には上の処理さえ施しておけば
重複ヒットは無視できる確率でしか発生しないので、Tripcode Finderでは
そこまで厳密に処理はしていません。あまり気にしなくてもいいんじゃないで
しょうかｗ
797:名無しさん＠お腹いっぱい。
12/12/08 03:13:22.49 rwOPHj120
>>796
なるほど……固定文字列に展開する作戦ですか。勉強になります。
「トリップ検索人のための便利ツール」的なものを、頑張って完成させようと思います。それでは。
798: ◆JouJaku.HzIz
12/12/08 11:00:59.32 lc8WRVoJ0
ご無沙汰しております。
電源が届いた後、色々試してみましたがどうも上手く行きません。
Quadro FX 3800, GTX480, GTX590をPCに挿してNVIDIAコンパネでQuadroだけCUDA offにして0.07a7 CUI64を[-c -g -x 128]で走らせると、下記エラーが発生して落ちます。
MERIKENsTripcodeFinderCUI: CUDA FUNCTION FALL FAILED: the launch timed out and was terminated (file 'C:/Users/Nullpo/Documents/Visual Studio 2010/Projects/MERIKENsTripcodeFinderCUI/Source Files/MTF_CUI_CUDA_SHA-1.cu', line 554)
MERIKENsTripcodeFinderCUI: CUDA FUNCTION FALL FAILED: all CUDA-capable devices are busy or unavailable (file 'C:/Users/Nullpo/Documents/Visual Studio 2010/Projects/MERIKENsTripcodeFinderCUI/Source Files/MTF_CUI_CUDA_SHA-1.cu', line 461)
MERIKENsTripcodeFinderCUI: CUDA FUNCTION FALL FAILED: all CUDA-capable devices are busy or unavailable (file 'C:/Users/Nullpo/Documents/Visual Studio 2010/Projects/MERIKENsTripcodeFinderCUI/Source Files/MTF_CUI_CUDA_SHA-1.cu', line 461)
MERIKENsTripcodeFinderCUI: CUDA FUNCTION FALL FAILED: all CUDA-capable devices are busy or unavailable (file 'C:/Users/Nullpo/Documents/Visual Studio 2010/Projects/MERIKENsTripcodeFinderCUI/Source Files/MTF_CUI_CUDA_SHA-1.cu', line 461)
Quadro+GTX590だと発生しません。三枚挿すと発生します。仕方が無いので、現在はGTX480+GTX590で運用しています。
とりあえず
ガッ!
799: ◆MERIKEN4.k
12/12/08 11:52:16.32 vyeW7s150
>>798
||// ∧＿∧|∧＿∧
||/ r( (n´・ω・`n) ぬるぽついてないのに「がっ」される
|| ヽﾟﾎllﾇ）|（ ）
￣￣￣￣￣ ｕ―ｕ'
line 554とline 461はそれぞれ
> CUDA_ERROR(cudaMemcpy(outputArray, CUDA_outputArray, sizeOutputArray * sizeof(GPUOutput), cudaMemcpyDeviceToHost));
と
> cudaError = cudaMalloc((void **)&CUDA_outputArray, sizeof(GPUOutput) * sizeOutputArray);
> ERROR0(cudaError == cudaErrorMemoryAllocation, ERROR_NO_MEMORY, "Not enough memory.");
> CUDA_ERROR(cudaError);
なので、両方共CUDA側のメモリの処理ですね。480と590のCCが2.0で、
Quadro FX 3800のCCが1.3なのでそれが原因かとも思ったのですが、
Quadro + GTX 590で発生しないみたいなのでそうでもないようですねえ。
エラーメッセージを見る限りではCUDAが無効担っているにもかかわらず
APIからQuadroが見えているようです。NVIDIAコンパネでQuadroの
CUDAをonにした場合はちゃんと動作しますか?
800:名無しさん＠お腹いっぱい。
12/12/08 11:53:33.23 rwOPHj120
>>798
ユーザー名がぬるぽなのかガッ!と言いたいためにぬるぽにしたのか……
エラーメッセージでググる限りでは、
>the launch timed out and was terminated
「Primary Device(ディスプレイデバイス)に指定されているGPUで長時間カーネル関数を実行しすぎている」
(探したページではPrimary Deviceを切り替えて対処していたが、基本全部使うGPU検索ではどうか……)
>all CUDA-capable devices are busy or unavailable
「ゾンビプロセスがGPUを占有している」(1つ目のエラーのせいで発生したエラーってことか？)
「fork()する前にcudaThreadExit()すればいいんじゃね」(要するに処理のミス？)
「ドライバを少し古いものに戻してみるのはどうか」(GPUあるある)
てなかんじかな。
参考URL:
URLﾘﾝｸ(d.hatena.ne.jp)
URLﾘﾝｸ(septieme-sens.blogspot.jp)
URLﾘﾝｸ(tsubame.gsic.titech.ac.jp)
URLﾘﾝｸ(devtalk.nvidia.com)
801: ◆MERIKEN4.k
12/12/08 12:42:23.91 vyeW7s150
>>800
ユーザー名がもともとNullpoなのですｗ
本名にしておかなくてよかった…
普通はlaunch time outはカーネルの処理時間が長すぎて発生する
エラーなんですけど、このケースではCUDAが無効になっているはずの
Quadroに対して検索スレッドが実行されているようなので、ドライバーの
バグ臭いです。Quadroが無効になっていて480と590だけで検索が実行されて
いるなら、エラーの数(=検索スレッドの数)は3個のはずなので…
時間ができたらこちらで再現できないか試してみます。
802: ◆MERIKEN4.k
12/12/08 13:08:25.51 vyeW7s150
>>800
もうちょっと調べてみたら、特定のGPUでCUDAが無効になっている場合、
cudaDeviceProp::computeModeをいちいちチェックして
そのGPUが有効かどうか確認しなければいけないことがわかりましたorz
URLﾘﾝｸ(stackoverflow.com)
URLﾘﾝｸ(www.clear.rice.edu)
直すのにちょっと時間がかかりますが、作業が終わったらここで報告するので
しばらくお待ちください。
803:名無しさん＠お腹いっぱい。
12/12/08 19:37:38.89 rwOPHj120
>>801
別に恨みはないが言わせてもらおう……
（ ・∀・） | | ｶﾞｯ
と ） | |
Ｙ /ノ 人
/ ） < >__Λ∩
＿/し' ／／. Ｖ｀Д´）/ ←>>801
（＿フ彡 /
話は飛びますが、検索していると、トリップキーの発見予定時間が
「it takes 2.3 days」などと表示されますよね？
あれが単純に、「出現確率の逆数÷検索速度」だとした場合、
検索し始めて表示時間だけ待ってトリップキーが出現する確率は
せ い ぜ い 6 3 % ぐ ら い し か な い
ことを最近発見しました。要するに、「1/XのくじをN回引く間に1回でも当たる確率」ということですが。
この確率は、Nが極端に大きいと二項展開やテイラー展開で近似することができ、それによると
確率E=1-EXP(-N/X)。1/Xを「出現確率」、Nを「検索速度(毎秒)×時間(秒)」とすれば、
上記の値が出るということです。しかもこの値は比で考えることができるため、
「予想時間までに出てくる確率は63.2%」
「予想時間の半分の時間で出てくる確率は39.3%」
「予想時間の倍掛けて出る確率は86.5%」
などといったことが分かります。分かりやすくグラフにしてみました。
URLﾘﾝｸ(up3.viploader.net)
……いや別になんとなく思いついただけなのですが(震え声)
804: ◆MERIKEN4.k
12/12/08 20:27:31.63 vyeW7s150
>>803
表示されているのはあくまでも「平均の」待ち時間なので、
「検索し始めて表示時間だけ待ってトリップキーが出現する確率」は
50%になるように調整されています。
> 単純に、「出現確率の逆数÷検索速度」だとした場合
これだと上の確率がちゃんと50%にならないので次のように計算しています。
pをパターンの出現確率とすると、n回のトリップの生成で
パターンが出現*しない*確率q_nは、
q_n = (1 - p)^n
になります。これから50%の確率でパターンが出現するのに必要な
トリップ生成の回数n'は、
0.5 ≒ (1 - p) ^ n' ⇔ n' = ceiling(ln(0.5)/ln(1 - p))
となります。これから発見予定時間sは、次の式で求められます。
s = n' / [平均速度(TPS)]
この計算はMTF_CUI_Patterns.cpp内のLoadTargetPatterns()の
後半で行われています。詳しくはソースを参照してくださいと言いたい
ところですが、公開されているソースのこの計算の部分は非常にわかり
にくいですｗ MPIRを使って書きなおしたので次のバージョンでは
前よりわかりやすくなったはずですが、大して変わらないかもしれません。
805: ◆MERIKEN4.k
12/12/08 21:35:28.17 vyeW7s150 BE:3258549577-2BP(12)
>>800
580+590の組み合わせでは問題は再現できませんでした。
バージョン306.97のディスプレイドライバで
NVIDIA Control Panelで580でCUDAを使用しないように設定してやると、
ちゃんとCUDAのAPIからは580は隠蔽されるようになっていました。
というわけで、この問題はディスプレイドライバのバグである可能性が高いです。
一応cudaDeviceProp::computeModeをチェックする処理を追加しておいたので、
次の開発版を試してみて下さい。
806:名無しさん＠お腹いっぱい。
12/12/08 21:48:52.17 rwOPHj120
>>804
それぐらい折込済み、だと……!? おみそれいたしました。
でも、その場合でも、q_nは、「発見予定時間だけ経つと0.5である」「発見予定時間のX倍経つと0.5のX乗になる」
ことから、発見確率の予測はそれほど難しくないようです(X=2だと発見確率が75%、X=0.5だと29.3%ほど)。
当該ソースは「// Calculate the matching probability etc.」あたりでしょうか。一度読んでみます。
807: ◆MERIKEN4.k
12/12/08 22:45:58.29 vyeW7s150
というわけでバージョン0.07のβ版を用意しました。
MERIKEN's Tripcode Finder 0.07 Beta 1
URLﾘﾝｸ(www.meriken2ch.com)
主な変更点はヒットまでの待ち時間の予測の改善と>>798で報告された
問題への対処です。
808: ◆MERIKEN4.k
12/12/08 22:53:58.49 vyeW7s150
>>806
たしかにその場所ですけど、n'を計算する部分を書いたときには
うごかすことしか考えていなかったので本当に分かりにくいですよｗ
809: ◆JouJaku.HzIz
12/12/09 11:00:55.85 VG0S6xiH0
>>807
対応して頂きありがとうございます。これから試してみます。
そもそもGeForceとQuadroではドライバが別パッケージになっているので、同時差しでバグが発生する可能性は大きそうですね。
Quadro使うやつはTesla使えってことか・・・。ついていけねぇ。
810: ◆MERIKEN4.k
12/12/09 18:29:54.81 D9EB7VO00
12桁トリップ検索のRadeonへの対応の作業もほぼ終了したので、
最高速を測定してみました。オクでお安く手に入れた中古の6990を2枚使って
速度を稼いでいます。真ん中の7970は延長ケーブルでマザボにつなげて
2枚の6990の上に乗っけています。温度の心配はしなくても良くなったので
ギリギリまでOCしています。動くかどうか半信半疑だったのですが
なんとかなるもんですねｗ
【GPU0】DIAMOND 6990PE54G Radeon HD 6990 4GB @ 900MHz (OC)
【GPU1】Gigabyte GV-R7970C-3GD Radeon HD 7970 @ 1120MHz (OC)
【GPU2】DIAMOND 6990PE54G Radeon HD 6990 4GB @ 900MHz (OC)
【CPU】AMD Phenom II X6 1100T (定格)
【OS】 Microsoft Windows 7 64bit SP1
【バージョン】MERIKEN's Tripcode Finder 0.07 Beta 1
【トリップの種類】12桁
【1SMあたりのブロックの数(CUDA)】N/A
【1CUあたりのワークアイテムの数(OpenCL)】自動
【1WGあたりのワークアイテムの数(OpenCL)】自動
【1GPUあたりの検索プロセスの数(OpenCL)】1
【1検索プロセスあたりの検索スレッドの数(OpenCL)】2
【その他のオプション】-g
【Display Driver】Catalyst 12.11 Beta8
【10分間の平均速度】7428.97 tripcodes/s
【GPUの平均速度】7428.97 tripcodes/s
【CPUの平均速度】N/A
【GPUの使用率】97～99%
【GPUの温度】83～93℃
【その他】GPUのみ。
811:名無しさん＠お腹いっぱい。
12/12/09 18:40:27.95 HKJ77yRt0
6990×2に5870付けて待て屋やったときは1500W超えたな（ワットチェッカー上限超えたｗ
そんときはCPUも使ってたけど同等に電気食ってそうだｗｗ
812:名無しさん＠お腹いっぱい。
12/12/09 19:13:06.04 38oGO8IR0
>>810
ぐおおおおお！
CPUが空気ｗｗｗｗｗ
813:名無しさん＠お腹いっぱい。
12/12/09 20:50:13.83 MhsAJkOg0
最速記録の塗り替えか
814:名無しさん＠お腹いっぱい。
12/12/10 10:34:47.93 NpT5XAETi
6990って水冷にすれば1スロット化出来るよな
でPCIex16スロット7本有るマザー結構な数有るよな
7枚刺したらいいんじゃないかな～
815: ◆MERIKEN4.k
12/12/10 17:54:08.67 FmksHTb00
>>811
CPUには負荷はほとんどかかっていないのでそこまではいってないはずです。
恐らく検索君1号だけで1100～1200Wぐらいです。
>>812
ここまでGPUが速いとCPU検索を同時実行すると却って速度が落ちるのです。
>>813
前スレを立てたときにくらべて10倍以上の速度が出せたので満足ですｗ
>>814
お金があればもっと色々試したいんですけど、自分はさすがにもう限界ですねえ。
勇者の登場を待ちましょうｗ
816: ◆MERIKEN4.k
12/12/10 18:59:36.16 FmksHTb00
あ、そうそう。Beta 1に問題がなければ今週の金曜日ぐらいに
バージョン0.07の正式版をうｐする予定なので、
不具合があればそれまでに報告していただけると有難いです。
817:☆☆勇者さま☆☆☆━━━╋━⊂(￣▽￣∩)
12/12/10 19:36:17.47 vm9IVZbG0
|￣￣￣￣￣￣￣|
| 速くなったな |
| |
| | ,. . _
|＿＿＿＿＿＿＿| --' 、 ￣￣ヽー- ､
| | ヽ￣7 , , ＼ 、 「￣ 7
| | ヽ / /_ /ハ |ヽ、＼ V ./
| | i il/ ヽl ＼ヽ. V
,. -{-､ __ .| ii i! ｏ ｏ | il |
{ Y/ l il |､ Д | li |
`t-く ヽN ` --- ＜ﾘiレ'
| | `ー-- ､ / II - ２ ヽ ｀丶、
| | ￣ !.ギ 子_ﾉ >-' !
| | ,r`''ｰ─''｡r'^ヽ､_,／- 、
| | / ｀ヽ、 , '~~｀V-─ ､ ）
| | / ／´｀､ ! (_ノ
i_j. / .／ ﾞ､ !
/＿／ ﾞ､ !
:::`ー':::::::::::::::::::::::::::::ヽこﾉ:::
818: ◆..//.//./5Hv
12/12/10 20:41:56.24 Era62auz0
スレ発見しましたー。
MERIKENさんなら./の10完12桁出そうな予感！
酉ありがとうございます(ﾉ^^)ﾉ
819:名無しさん＠お腹いっぱい。
12/12/10 22:20:03.40 LbISDnqB0
>>816
WinXP 32bit、GPUなしでver0.07 beta1の.exeを起動させると、「OpenCL.dllが見つかりませんでした…。」と出て起動できない（検索出来ない）。
ver0.06の安定版では起動させることが出来る
820: ◆JouJaku.HzIz
12/12/10 22:41:16.10 astkHfvt0
>>807
対応ありがとうございます。
最初にQuadro, 480, 590を繋げて"CUI64 -c -g"で実行。エラーも出ずに実行されました。自動ブロック数設定は相変わらず安定しませんが・・・
次にNVIDIAコンパネでQuadroだけCUDA offにして"CUI64 -c -g -x 192"で実行。下記エラーが出るも、検索自体は実行されます。
MERIKENsTripcodeFinderCUI: CUDA FUNCTION FALL FAILED: unknown error (file 'C:/Users/Nullpo/Documents/Visual Studio 2010/Projects/MERIKENsTripcodeFinderCUI/Source Files/MTF_CUI_CUDA_SHA-1.cu', line 560)
画面の表示はこんな感じです。
CUDA0: (Quadro?)
CUDA1: 560.5M TPS, 192 blocks/SM (480)
CUDA2: 518.7M TPS, 192 blocks/SM (590)
CUDA3: 518.6M TPS, 192 blocks/SM (590)
^Cで強制終了させて、もう一度実行させると、例のエラーが三行出てCPUでのみ検索が実行されます。
挙動が良く分からない・・・
OpenGL用にQuadroを残しておきたいけど、熱的にやばそうなので480と590だけで運用することにします。
821:名無しさん＠お腹いっぱい。
12/12/10 23:34:40.12 Ya8wVC3a0
>>819
GPUでOpenCLかCUDA扱えないと使いづらいってのが俺の中でのこのソフトの認識
CPUだけなら待て屋とかSHArpとかがあるし(探索空間が違うから一緒にしてはいけない気もするが)
822: ◆MERIKEN4.k
12/12/11 07:36:11.77 G8KcgggZ0
>>819
報告ありがとうございます。こちらでも確認できました。
取りあえずOpenCLを添付することで対処したいと思います。
823: ◆MERIKEN4.k
12/12/11 08:50:25.36 G8KcgggZ0
>>821
実際Tripcode FinderのCPU検索は待て屋やSHArp Tripperほど速度は出ないですからねえ。
GPUが使用出来ないと警告が毎回出るのはさすがにやりすぎなのでこれは直しておきます。
824: ◆MERIKEN4.k
12/12/11 09:04:22.97 G8KcgggZ0
>>818
有難うございます。正規表現でいろいろパターンを指定できるので、
結構遊べますよw
825: ◆MERIKEN4.k
12/12/11 10:05:57.67 G8KcgggZ0
>>820
やっぱりドライバのバグみたいですねえ。
今度試す機会があったら"CUDA DEVICES"の"Compute Mode"の値を
調べてみて下さい。問題を回避できるかもしれません。
826:名無しさん＠お腹いっぱい。
12/12/11 15:41:44.59 l2lR+Gjg0
なんかやってます
WindowsのパスワードはGPUを25個使えば約6分から6時間で突破が可能、
毎秒3500億通りもの総当たりが可能な方法とは？
URLﾘﾝｸ(gigazine.net)
827: ◆MERIKEN4.k
12/12/11 16:11:15.02 G8KcgggZ0
>>819
ついさっき修正が完了しました。次の安定版では直っているはずです。
828: ◆MERIKEN4.k
12/12/11 16:31:35.11 G8KcgggZ0
>>826
これ5台のラックマウントサーバーですよね。グラボが25枚だそうですけど、
サーバーによって構成が違うみたいです。8枚載っているサーバーの
写真があるので、8枚+5枚+4枚*3という構成でしょうか。他のサーバーの
GPUを仮想化してHashcatで利用しているのは非常に興味深いです。
いつか自分でもこんな豪勢なクラスターを組み立ててみたいですねえ。
829:名無しさん＠お腹いっぱい。
12/12/11 16:40:36.00 l2lR+Gjg0
>>828
やろうと思えば、個人レベルでも出来てしまう辺りがおもしろいですね
830:名無しさん＠お腹いっぱい。
12/12/11 17:23:04.08 KG0LrKw40
古いPCが沢山あるのでネットワーク対応型MTFを待ってます
831: ◆MERIKEN4.k
12/12/11 19:39:54.02 G8KcgggZ0
>>826の記事のグラボが8枚載ったラックマウントサーバーはどうやら
これのようです。
URLﾘﾝｸ(www.advancedhpc.com)
しかしこうやってみると壮観ですねえ。
URLﾘﾝｸ(gigazine.jp)
832:名無しさん＠お腹いっぱい。
12/12/11 21:55:11.42 eYtNkyH+T
はりにきたらすでにはられてたか>>826
833:名無しさん＠お腹いっぱい。
12/12/11 22:56:32.68 6gmHNGHj0
>>821
常用しているのはうにだけど、
このソフトはCPUのみでも動くようになっているから、動かないのは問題かなと思って報告した。
>>827
早い対応ありがとうございます。
OpenCL.dllをいれようと思ったものの、検索してもよく分からなかったもので……。
834:名無しさん＠お腹いっぱい。
12/12/11 23:05:15.63 AXhxlsuZ0
>>828
控えめに一枚500M/sだとしても×25で12.5G/sか・・・
8完が(ln(0.5)/ln(1-1/64^8))/(12.5*10^9)≒4.3時間で出てくる計算に
835: ◆MERIKEN4.k
12/12/11 23:13:37.93 G8KcgggZ0
>>830
とりあえず10桁トリップ検索とコードの整理をするのが先ですけど、
ネットワーク対応はいずれぜひやりたいですねえ。
836: ◆MERIKEN4.k
12/12/11 23:29:58.01 G8KcgggZ0
>>834
研究発表のスライドにはSHA-1で63G hashes/sでているとありましたよ。
URLﾘﾝｸ(passwords12.at.ifi.uio.no)
これはパスワード解析での数字なので、トリップ検索ならもうちょっと
速くなるでしょう。なかなか豪気ですねえｗ
837:名無しさん＠お腹いっぱい。
12/12/11 23:31:14.25 AXhxlsuZ0
>>830
ネットワーク対応の暁には学校のPCルーム総動員で検索させてみたいな・・・
いやGPU買えよと言われそうだが
838:名無しさん＠お腹いっぱい。
12/12/11 23:33:44.76 AXhxlsuZ0
>>836
>トリップ検索ならもうちょっと速くなるでしょう
要するに単にハッシュ出して比較、だけじゃない最適化が掛かっているのか……
8完が1時間切るとかどんなモンスターだｗｗ
839: ◆JouJaku.HzIz
12/12/12 00:21:33.94 gPuKMjn30
>>825
Compute Modeは全てcudaComputeModeDefaultでした。
違うのはCompute Capabilityだけで、Quadroは1.3、他は2.0です。
他の手を考えてみます。
840: ◆MERIKEN4.k
12/12/12 06:17:27.59 FX/ZJoUj0
>>839
そうですか。それは残念… 将来的には各GPUを使用するかしないかを個別に
設定できるようにする予定なのでいずれ解決できるかもしれませんが、
今の段階では難しいですねえ。
841:名無しさん＠お腹いっぱい。
12/12/12 14:55:15.28 /XRCYi610
>>343のteslaがGTX5シリーズに負けてるのが印象的です
fermiコアの解析速度はプロセッサクロック×メモリバンド幅ですかね？
うちの560tiが580の報告の半分の速度しか出ないもので
842: ◆MERIKEN4.k
12/12/12 16:13:34.77 FX/ZJoUj0
>>841
メモリバンド幅は関係ないです。
580と560tiはそれぞれGF110とGF114なので単純には比較できないですけど
半分だとちょっと遅すぎるような気がしますね。ちゃんとCC 2.1用のバイナリは
入ってるはずだけど…
843:名無しさん＠お腹いっぱい。
12/12/12 16:49:47.66 EU7chw1W0
GF114はSMあたりのコア数はGF110の32コアから48コアに増えていますが、
レジスタ数は増えていなくて、GF110は16SMでGF114は8SMなので
GF114ではレジスタがボトルネックになりがちだったと思います。
とはいえSMあたりのコア数が増えている分少しは向上しているようでしたし、
リファレンスではクロックもGTX560Tiの方が上なので、半分となると遅すぎる気もしますが、
OCされたGTX580との比較でしょうか？
844:841
12/12/12 17:12:49.24 SeK148sf0
【GPU】Geforce GTX560ti ×2
【CPU】core i5 3470
【OS】Microsoft Windows 7 64bit SP1
【バージョン】MERIKEN's Tripcode Finder 0.07 Beta 1
【トリップの種類】12桁
【1SMあたりのブロックの数(CUDA)】192
【その他のオプション】-g -x 128
【Display Driver】306.97
【10分間の平均速度】 762.15Mtripcodes/s
【GPUの使用率】99％
【GPUの温度】71～80℃
【その他】 CUDA0,1:約381M TPS
845: ◆MERIKEN4.k
12/12/12 18:58:07.81 FX/ZJoUj0
さっき測ったら定格の580が683M TPSぐらいなので560tiの速度は
55%ぐらいですか。CUDA GPU Occupancy Calculatorで調べてみても
特にCC 2.1でOccupancyが下がるということもなかったので、ちょっと
原因がよくわからないですねえ。
846:名無しさん＠お腹いっぱい。
12/12/12 19:21:29.29 SeK148sf0
GF114はGPGPUには向いてないのですかねー。
現在最速はやはりGF110かな？
847:名無しさん＠お腹いっぱい。
12/12/12 19:37:38.59 jCx6f4p80
つURLﾘﾝｸ(dokumaru.wordpress.com)
848:名無しさん＠お腹いっぱい。
12/12/12 20:44:44.70 EU7chw1W0
55%ですか・・・もう少し出てもよさそうな気もしますが、おかしいというほどではないかと思います。
単精度や32ビット整数の演算性能自体は、GTX560Tiはコア数とクロック的にGTX580の80%近くありますが、
それはピーク性能であって、SHA-1ハッシュの演算ではレジスタがそれなりに必要になります。
SM数とクロック的にはGTX560TiはGTX580の53%程度であり、
それぞれのSMの違いはコア数（と倍精度や特殊関数など）でレジスタ数に変化は無いので
レジスタがネックでコアを使いきれていないのだと思います。
GF114はグラフィックよりではあると思いますが、GPGPUでもレジスタを大量に使うものばかりではないでしょうし
消費電力や値段を考えると、GPGPUにはベストではないけどそれなりにではないでしょうかね。
GK104はGPGPUにはピーキー過ぎてお勧めしませんけど・・・
849:名無しさん＠お腹いっぱい。
12/12/13 04:38:55.18 Fj613XFy0
GK110買えそう
楽しみ
850: ◆MERIKEN4.k
12/12/13 05:13:49.96 sid26Nen0
>>848
なるほどなるほど… CUDA Toolkit 5.0に添付されているOccupancy Calculatorでは
このあたりの事情が反映されていないようです。カーネルのレジスタ数は46～48で
Occupancyは42%なのでレジスタ数が特に多いというわけではないのですが、
これがボトルネックになっているのは確実ですね。
851: ◆MERIKEN4.k
12/12/13 05:17:50.58 sid26Nen0
>>849
Tesla K20ですか? いいな～ 買えたら是非報告をお願いします。
852: ◆MERIKEN4.k
12/12/13 05:44:43.66 sid26Nen0
>>838
パスワード解析に比べてトリップ検索ではキーの生成が比較的単純なので、
それをうまく利用してやれば速度は1～2割上がる傾向があります。
GPUクラスタの場合はノード間通信がボトルネックにならないので
更に速くなるものと思われます。しかしもう12桁トリップだと9完以上でないと
危ないですねえ。
853:名無しさん＠お腹いっぱい。
12/12/13 05:52:42.00 q8Aa1QZH0
>>852
いやいやいや
あくまでも12桁ですから、キーを割られる危険という意味では何完であろうと関係はないです
我々のような好き者にとっては問題なんですが
854: ◆MERIKEN4.k
12/12/13 07:25:01.97 sid26Nen0
>>853
> あくまでも12桁ですから、キーを割られる危険という意味では何完であろうと関係はないです
あ、「危険」と書いたのはそういう意味じゃないです。
トリップの場合はある程度一致すればなりすましができるので
キーが割られなくても十分危ないんですよね。トリップが一致しているか
どうかを判断しているのは一般のユーザーで、普通の人はわざわざ
12桁目まで細かく確認しているわけではありません。ここらへんは普通の
パスワードとはぜんぜん違うところです。
855: ◆MERIKEN4.k
12/12/13 08:47:43.45 sid26Nen0
今唐突に12桁トリップのCPU検索を高速化するアイディアを思いついたん
ですけど、1月の中旬まで帰省しているので実装はそれまでおあずけです。
残念…
なんでMTFのCPU検索がSHArp Tripperやhip2に比べて遅かったのか
不思議で仕方がなかったんですけど、よく考えたら普通のSHA-1の
ルーチンを使いまわしてたせいで、SSE2のレジスタをトリップ検索に
特化した形で効率的に使用していなかっただけでしたｗ
1個のハッシュの生成を高速化するより、SSE2の128bitレジスタを使って
4個同時に生成したほうが速いに決まってますよねえ。
856: ◆MERIKEN4.k
12/12/13 09:05:24.37 sid26Nen0
あと、よく考えたらキーの動的生成とBitslice DESのルーチンの動的書き換え
( >>712-713 )で10桁トリップのCPU検索も高速化できることに気づきました。
なんで時間のないときに限って面白い考えを思いつくんだろうorz
857:ののたん ◆KiwamonoL.
12/12/13 13:36:10.93 rNLBcKX70
>>855
えっ！？
SIMD ってなかったの！(SIMD るってなんだよ。ｗ
まさか、Radeon でもやってないとか・・・・・。
ソースを読んでみる気は無い。ｗｗｗ
あとまあ Hashcat 知ってるんなら知ってるかもしれんが。
URLﾘﾝｸ(hashcat.net)
858:名無しさん＠お腹いっぱい。
12/12/13 16:42:37.91 tgXDqPZ80
もうこれはMERIKENさんにメチャクチャ頑張ってもらうしかない展開
859: ◆MERIKEN4.k
12/12/13 17:01:04.69 sid26Nen0
>>857
SSE2を使ってるルーチンを拾ってきたんですけど、
ベクター化されてないのであんまり速度が出てなかったみたいです。
RadeonのほうはCUDA版のベタ移植なのでそれこそなにもしていませんｗ
OpenCLドライバが頑張ってるのでせう。Southern Islandsだとベクトル化しても
あんまり意味ないみたいですし… 資料のほうはあとでありがたく読ませて頂きます。
これでさらなる高速化が出来るかもしれないですね。ぐへへへへ…
860: ◆MERIKEN4.k
12/12/13 17:06:59.50 sid26Nen0
>>858
明日の朝の飛行機の便に間に合わせるのに徹夜で荷物をつめはじめたところなので
さすがに帰省前は無理ですｗ 来月を楽しみにしていて下さい。
家を出る前に0.07の安定版はうｐしておきます。
861: ◆YSRKENkO6Y
12/12/13 19:27:33.63 tgXDqPZ80
>>806です。
検索作業をサポートするソフトをリリースしてみます(実験版だけど)。
依頼を検索パターンに変換したり、特定パターンを自動生成したりできます。
良かったらどうぞ。
URLﾘﾝｸ(www1.axfc.net)
862:名無しさん＠お腹いっぱい。
12/12/13 19:51:21.07 DyqVV5mA0
レジューム機能がほしいです
863:名無しさん＠お腹いっぱい。
12/12/13 21:20:16.88 tgXDqPZ80
>>862
なんで検索空間＞＞酉空間なのにみんなレジューム機能が欲しくなるんだろうな……いや俺も思ってたことあったけど
自動実行と自動保存はAlpha 7で既に実装されてるから除くとして
864:名無しさん＠お腹いっぱい。
12/12/13 21:44:28.47 sR2+e44BP
Radeon HD8000シリーズ楽しみすぎる
865: ◆MERIKEN4.k
12/12/13 23:04:23.15 sid26Nen0
バージョン0.07の安定版です。
MERIKEN's Tripcode Finder 0.07
URLﾘﾝｸ(www.meriken2ch.com)
Alpha 7からの変更点は以下になります。
・OpenCLドライバがインストールされていないと起動できないバグの修正。