【トリップ検索】MERIKEN's Tripcode Finder

1 名前： ◆MERIKEN4.k [2012/09/25(火) 18:24:38.09 ID:BDWiD/680] こちらはMERIKEN's Tripcode Finderの本スレです。動作報告・質問・要望等は こちらでどうぞ。 MERIKEN's Tripcode Finder(旧CUDA SHA-1 Tripper MERIKEN's Branch)は 2012年9月現在で最速の12桁トリップ検索プログラムです(最高速の記録は 1689.88M tripcodes/s)。CPUのみでも検索できますが、NVIDIA GeForce シリーズのビデオカードを使用すれば非常に高速に検索を行うことが できます。特徴は以下の通りです。 ・ビデオカードのGPUによる高速検索(CPU検索にも対応)。 ・GUIによる簡単な操作(コマンドラインからの使用も可能)。 ・正規表現によるターゲットの指定。 ・漢字等のShift-JIS文字を含むキーの探索。 ・ヒット率、ヒットまでの平均時間等のさまざまな情報の表示。 ・ターゲットの数の制限の撤廃。 ・10桁トリップ検索への暫定的対応。 ・検索速度の実行時の最適化。 ・GPLv3で公開されたソースコード。 ■入手先 www.meriken2ch.com/programming/merikens-tripcode-finder ■前スレ anago.2ch.net/test/read.cgi/software/1311428038/ 577 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 15:10:23.29 ID:jethYJ0v0 BE:1197019229-2BP(12)] >>576 > -a 2←1つのAMDのGPUに対する検索スレッドの数？ これはあってます。2番目のはGUI版とCUI版が通信するときに使うおまじないです。 578 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 18:53:29.24 ID:jethYJ0v0] これまで使っていた検索君1号のFermi軍団に加えて、開発用PCの7970でも同時に 検索をしているのですが、ここ数日で3回ブレーカーが落ちましたｗ GTX 590の電圧を絞ることでなんとか対処しましたが、 消費電力のほうもそろそろ限界です。 579 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/11(日) 19:25:42.41 ID:6A5ebb7J0] 単相200V契約しよう 580 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/11(日) 20:09:00.62 ID:tLApF2aS0] 前スレだったと思いますが、ドライバを触ったらこのソフトからGPUが認識されなくなったと書き込んだものです。 相変わらず認識されないままですが、今日system32に入っているclinfo.exeというプログラムでOpenCLの情報を取得できると知ったので、 取得できたテキストをアップロードしておきます。 cyclotron.moe.hm/up/small/src/cyclotron_s4856.zip 念のため正常に使用できるPCのぶんも入れておきます。 581 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/11(日) 20:10:29.57 ID:kFGFk5jZ0] USAは、110ボルトですね。 582 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 20:22:46.82 ID:jethYJ0v0] >>580 ひょっとしてこのスレの>>395さんですか? たしかに両方ともRadeonが見えていますね。 OSが64bit版ならCUI64ならうまく動くかもしれません。 583 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 20:29:24.47 ID:jethYJ0v0] >>579 したいのはやまやまなんですけど、今のアパートだと無理なんです… >>581 120Vです。15Aなので1800Wまで大丈夫なんですが、グラボ4枚で1100Wぐらい いってます。やばいです。 584 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/11(日) 20:42:11.43 ID:tLApF2aS0] >>582 あ、このスレでしたか。 Alpha5をダウンロードしてCUI64を起動してみましたが MERIKENsTripcodeFinderCUI: OPENCL FUNCTION FALL FAILED: CL_DEVICE_NOT_FOUND (file 'Source Files\MTF_CUI_Main.cpp', line 732) と表示されてそこから進みません。 c++は門外漢でソースちらっと眺めただけですけど、プラットフォームが2個あって、最初の片方がCPUのみってところでなんかコケたりしてません？ 585 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 20:47:49.07 ID:jethYJ0v0] >>584 CPUは無視するようにしているので問題はありません。 32bit版のMERIKENsTripcodeFinderCUI.exeではどうですか? 586 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/11(日) 20:51:41.93 ID:tLApF2aS0] >>585 32bitでも同じエラーです。 587 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 21:06:57.97 ID:jethYJ0v0] う〜ん、やっぱりCL_DEVICE_NOT_FOUNDが返されているのかなあ。 うちのIntelのドライバではエラーは出なかったんですが… これから修正して新しい開発版をうｐするので、そちらを試してみてください。 588 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 21:57:58.95 ID:jethYJ0v0] >>580で報告していただいた、環境によっては起動できないバグを修正した 開発版をうｐしました。というか、これで直っているといいんですけど… MERIKEN's Tripcode Finder 0.07 Alpha 6 www.meriken2ch.com/programming/merikens-tripcode-finder 589 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/11(日) 22:02:50.51 ID:tLApF2aS0] >>588 ありがとうございます。 起動できるようになりました。 590 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 22:08:25.91 ID:jethYJ0v0] >>589 それはよかった! こちらこそバグ報告をありがとうございました。 591 名前： ◆MERIKEN4.k mailto:sage [2012/11/11(日) 22:12:41.07 ID:jethYJ0v0] もうそろそろ安定してきたと思って10桁トリップ検索の移植の作業を 始めてたんですけど、まだ結構不具合が残っていますねえ。 592 名前： ◆GTX680Mcys3u mailto:sage [2012/11/12(月) 06:26:56.37 ID:2Urt7HcK0] >>588 全グラフィックチップ(680M+iHD4000)指定だと 検索開始後エラーメッセージなしでソフトごと落ちます HD4000はバッサリ切った方がいいかと思われます 593 名前： ◆MERIKEN4.k mailto:sage [2012/11/12(月) 07:13:21.49 ID:x4jnfC130] >>592 ありゃりゃ… こりゃいかんですねえ。テスト用には便利だったんですけど しょうがないですね。次の開発版からは無視するようにします。 594 名前： ◆MERIKEN4.k mailto:sage [2012/11/12(月) 07:34:25.71 ID:x4jnfC130] というわけで面倒くさいのでIntelのプラットフォームは最初から無視することに しちゃいました。OpenCLはオープンスタンダードな分だけそれに伴う 問題も多いですね。 595 名前： ◆MERIKEN4.k mailto:sage [2012/11/12(月) 08:51:49.76 ID:x4jnfC130] 気を取り直して10桁トリップ検索の移植作業を続けます。 CPU側のコードは10桁の場合とほとんど同じなのですぐに終わりました。 問題はOpenCLのコードですが、バグが紛れ込むと見つけるのが 非常に困難になるので、慎重に作業を進めてます。 596 名前： ◆MERIKEN4.k mailto:sage [2012/11/12(月) 10:00:14.78 ID:x4jnfC130] カーネルの入り口の部分の書き換えは終了しました。 あとはBitslice DESの本体だけですが、CUDAのコードをコピペするだけなので 問題はないでしょう。うまく動いてくれるといいんだけど、どうでしょうね〜 597 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/12(月) 10:29:57.20 ID:FaMyVn9Z0] geforceでopenCL版って動くの？ 動いてもcudaよりは遅い？ 598 名前： ◆MERIKEN4.k mailto:sage [2012/11/12(月) 13:13:54.13 ID:x4jnfC130] >>597 いまはNVIDIAのカードでは強制的にCUDAを使うようにしています。 OpenCLでも動くことは動くと思いますけど、基本的に全く同じコードなので 速度は変わらないでしょう。 599 名前： ◆MERIKEN4.k mailto:sage [2012/11/12(月) 13:27:00.53 ID:x4jnfC130] OpenCLの10桁検索のコードは1発で動いたんですけど、Bitslice DESで使う変数を 何も考えずに全部__privateメモリ空間に突っ込んだら、案の定というか まったく速度が出ていませんｗ まあでもコードの書き換え自体は問題なかったよう なので、とりあえず一安心です。これから__globalと__localを試してみます。 600 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/12(月) 13:31:46.97 ID:ClnWJME80] khronosの姿勢として標準のカーネルコンパイラを用意しないのはわかるんだけど やっぱりglslの轍をちょっとは生かしてほしかったってのが個人的な思い meriken氏乙 601 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/12(月) 13:37:18.53 ID:/PHpLzn40] >>594 当方では一応4M/sぐらいで動くので、 IntelHD4000を使うか否かをチェックボックスとかで決めればいいと思いまーす 602 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/12(月) 13:56:47.61 ID:FaMyVn9Z0] >>598 そうなんだ ありがとう 603 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/12(月) 21:58:25.06 ID:+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 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/12(月) 21:59:57.08 ID:+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 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/12(月) 23:26:48.48 ID:ODHrB/Fw0] NVIDIA，第2世代Kepler「GK110」搭載の「Tesla K20」を正式発表。CUDA Core数は最大2688基に www.4gamer.net/games/121/G012181/20121110004/ 606 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/12(月) 23:34:15.80 ID:cklfqCrp0] X5680はOCすりゃいいじゃん 607 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 02:39:21.60 ID:FPgsAJYQ0] DualCPUにQuadro突っ込んでるようなガチWS機でOCとかあり得んでしょ 608 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 02:50:45.59 ID:l+bGYcyn0] 倍率ロックフリーだろ？ 609 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 06:21:00.78 ID:FpPqufE20] 今気づいたんですけど「1CUあたりのワークグループの数」じゃなくて 「1CUあたりのワークアイテムの数」ですね、これ。 こりゃ当分の間安定版は出せないな… 610 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 06:26:29.04 ID:FpPqufE20] >>603-604 報告ありがとうございます。CPU検索ももうちょっと何とかしたいですね〜 "Unknown"のエラーが出ているのは謎ですが、そこのエラーは無視するように 直しておきます。 611 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 06:38:08.95 ID:FpPqufE20] 予想通りというべきか、10桁トリップ検索はなかなかスピードが出てくれません。 まじめにプロファイラを使わないと駄目ですね、こりゃ。 まあCUDAのときもそうだったので、のんびり時間をかけて取り組むことにします。 612 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 07:26:06.97 ID:FpPqufE20] >>605 GK110も試してみたいんですけどね〜 Amazon Cluster GPU Instancesで使えるようにならないかしらん。 613 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 09:29:27.60 ID:FpPqufE20] Bitslice DES用の一時変数をどのメモリ空間に置けばいいのかいまいち よくわからないので、とりあえず#ifdefで切り替えられるようにしておきます。 あと、一回のBitslice DESを複数のスレッドで同時に処理するかどうかも CPU側で設定できるようにする予定です。こういうところは実行時にカーネルを ビルドできるOpenCLはいいですねえ。 614 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 09:34:06.33 ID: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の有効無効では速度は変わらないですよ www.rupan.net/uploader/download/1352766539.png 615 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 14:32:35.77 ID:1xHvqvP20] ついにデュアルTahitiカードがAMDから登場。エーキューブ，サーバー向けグラフィックスカード「FirePro S10000」を国内発売 www.4gamer.net/games/133/G013322/20121113023/ 616 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 14:34:05.53 ID:1xHvqvP20] XeonPhiってどうなんですかねえ 617 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 15:01:20.15 ID:PlVYlKIA0] 負荷で思い出したけど、同じGPU使用率でもmtyのGPU版は95℃までいくけどMERIKENは89℃までしか上がらないね。 618 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 17:19:59.24 ID:VU2bD6Zz0] >>615 サーバ向けのFireProだし、3599ドルらしい。 >>616 OpenMPが使えて既存アプリの移植が楽らしいけど、既にOpenCLになっている場合はメリット薄そう。 619 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 19:10:11.95 ID:FpPqufE20] >>614 なかなか良い感じに仕上がってますね。それだけOCした7970なら単体で1500M TPS近く いくはずなので、単純に2枚で2倍の速度という訳にはいかないみたいですが… Alpha 6では-aオプションは有効なはずです。CFXの話は別のところで見かけたんですが、 関係無かったみたいですね。 620 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 19:22:18.89 ID:FpPqufE20] >>617 Bitslice DESはSHA-1の処理に比べてメモリへのアクセスの量が段違いですからねえ。 621 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 19:30:43.61 ID:FpPqufE20] >>615 いいですね〜 いつかこういうのをたくさんラックマウントサーバーに乗っけて Tripcode Finderを動かしてみたいですｗ 622 名前： ◆MERIKEN4.k mailto:sage [2012/11/13(火) 19:59:41.68 ID: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 mailto:sage [2012/11/13(火) 20:10:06.56 ID: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 www.informit.com/articles/article.aspx?p=1732873&seqNum=3 624 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 20:28:02.76 ID:ES3128Qj0] >>614 壁紙についてｋｗｓｋ MERIKENさんの公式記録が越される日も近いか…… 625 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 23:12:02.50 ID:ATY55mX00] >>624 この辺で拾ってます これがいつのものだったかは忘れましたw www.smashingmagazine.com/tag/wallpapers/ 626 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/13(火) 23:43:31.28 ID:ES3128Qj0] >>625 そのまま2012/11の記事を見ても出てこないorz www.smashingmagazine.com/2012/10/31/desktop-wallpaper-calendar-november-2012/ よろしければ画像アップローダに上げてくれませんか？ 627 名前： ◆MERIKEN4.k mailto:sage [2012/11/14(水) 00:24:53.55 ID:mDY3eRDW0] >>624 それはどうでしょうね… (￣ー￣)ﾆﾔﾘ 628 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 00:35:11.76 ID:E/rQ2cMp0] >>626 ほい www.rupan.net/uploader/download/1352820886.jpg 629 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 00:38:28.17 ID:eHqSRvz80] >>621 どうせならHD7970 X2に行きませんか？ 消費電力が凄まじいのと、スロット占有が問題ですけどｗ なんかリンクが貼れないので詳細は検索してください デスクトップ向けにHD7950のデュアルが出てくれれば一番ですけどね。 HD7950のCFはグラフィックでも割りと良いというレビューもあったので、需要もある程度ありそうですし。 630 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 00:50:05.56 ID:eHqSRvz80] >>622-623 OpenCLは以前よりは情報も増えたようですが、まだ茨の道なのでしょうかね・・・ 631 名前：626 mailto:sage [2012/11/14(水) 01:34:12.64 ID:vuLXlPiG0] >>628 ありがとうございます！ 632 名前： ◆MERIKEN4.k mailto:sage [2012/11/14(水) 01:39:17.60 ID:mDY3eRDW0] >>630 >>622-623はJohn the Ripperのメーリングリストを見てて気づきました。 www.openwall.com/lists/john-dev/2012/08/13/8 www.openwall.com/lists/john-dev/2012/09/10/5 このAlexanderという方はJohn the RipperとBitslice DESの偉い人です。 流石です。 www.openwall.com/lists/john-users/2011/06/22/1 633 名前： ◆MERIKEN4.k mailto:sage [2012/11/14(水) 01:40:56.22 ID:mDY3eRDW0] >>630 まあGPGPUの不条理な制約にはCUDAで慣れっこになっているので どうということはありませんｗ 634 名前： ◆MERIKEN4.k mailto:sage [2012/11/14(水) 01:45:30.91 ID:mDY3eRDW0] >>629 うちの検索用マシンにはGTX 580が2枚と590が1枚載っているので、 7970 2枚は余裕ですｗ 今590を売っぱらって6990を買おうかどうか 考えているところです。 635 名前： ◆supernova.rT mailto:sage [2012/11/14(水) 02:04:56.65 ID:Bf0HEkX10 BE:1020114162-DIA(123421)] 僕はもうラデ2枚構成にしたのでゲフォ売ります 10桁検索対応が楽しみですよー 636 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 02:16:09.00 ID:eHqSRvz80] >>633 頼もしいです、頑張ってください。 >>634 HD7970を1ボードに2基載せたもので8ピンx3で3スロット占有という モンスターというかクレイジーな代物が出るらしいですｗ それの複数枚挿しは電源だけでなくマザボもかなり選びそうです。 HD7950のデュアルで8ピンx2で2スロットであればまだマシなのですけどねえ。 637 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 03:47:50.59 ID:peEcrqnb0] やっぱりさよならゲフォの流れになったね 638 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 06:49:17.05 ID:AbSbupmCP] RADEONは普及用チップでも倍精度が高速なのがいい 639 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 07:58:59.38 ID:vuLXlPiG0] mtyGPUがRadeonしか対応してないから、むしろゲフォ対応検索は(10桁では)貴重なんだが 640 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 13:24:08.27 ID:85Ooiiep0] >>638 マジレスすると倍精度演算が速いのは7970だけだしトリップ検索に倍精度演算の出番は無いぞ 641 名前： ◆MERIKEN4.k mailto:sage [2012/11/14(水) 14:42:36.69 ID:mDY3eRDW0] >>635 10桁トリップ検索は12桁よりかなり難しいので、実際どこまで速度を出せるかは わかりませんけどね〜 というか12桁検索の移植は正直うまくいきすぎでしたｗ 地道に取り組む予定なので、のんびり待っていて下さい。 642 名前： ◆MERIKEN4.k mailto:sage [2012/11/14(水) 15:16:17.63 ID:mDY3eRDW0] で、あれから色々試してみて、Bitslice DES用の一時変数はローカルメモリに おかないと全く速度が出ないことが分かりました。ローカルメモリは ワークグループ内で共有されるので、Bitslice DESを8個のスレッドで 並列処理するように書き換えてやりました。 その後、さらに性能を上げるためにAMD APP Profilerで解析してみました。 あんまり期待してなかったｗのですが、非常に使いやすいです。 で、気になっていたOccupancy Analysisを行なってみたら、 案の定ローカルメモリ(LDS)の使い過ぎであることが判明しました。 www.meriken2ch.com/files/2012-11-13-AMD-APP-profiler.jpg 643 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 15:31:41.20 ID:AbSbupmCP] >>642 へぇ〜 人目でボトルネックがLDSにあることが示されてる 凄いな 644 名前： ◆MERIKEN4.k mailto:sage [2012/11/14(水) 15:51:16.97 ID: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 mailto:sage [2012/11/14(水) 16:01:44.57 ID:mDY3eRDW0] これで使用するメモリの量は半分近くに減って、うまくいけば CUDA版ともども10桁検索の速度が倍になることになります。 アルゴリズムはかなり複雑になりますが、試してみる価値は十分にあります。 hip2の話を聞いて、キーの生成方法にかなりの工夫の余地があることに 気づいたのは僥倖でしたｗ 646 名前： ◆MERIKEN4.k mailto:sage [2012/11/14(水) 16:04:10.35 ID:mDY3eRDW0] >>643 実際かなり便利です。CUDAのときはなんせExcelのスプレッドシートを 使わないとOccupancyのグラフが見れませんでしたからねｗ 647 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 19:07:46.29 ID:vuLXlPiG0] >>645 >速度が倍 うおおおおお！？頑張って下さい！ 648 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/14(水) 19:20:22.61 ID:dspeEFEK0] GTX670では470Mt/sくらいしか出ません。倍精度を使うわけでもないのになんでだろう。 649 名前：前スレ927 ◆JouJaku.HzIz mailto:sage [2012/11/14(水) 20:16:05.54 ID: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 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/15(木) 00:03:05.86 ID:Gr7998EA0] >>642 これは便利そうですね。 >>644 DESは歴史もあり奥が深いですね。 >>648 レジスタ数がネックになって演算ユニットを使いきれていないのだと思います。 651 名前：648 mailto:sage [2012/11/15(木) 02:21:08.90 ID:aNTlQCIF0] レジスタの仕様が違うのか。最適化しなおさないといけないわけね。 652 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 03:09:09.94 ID:dQ9rq2KX0] >>648 >>651 トリップ検索の速度は整数演算の性能に大きく影響されるんですけど、 GTX 600シリーズで使われているKeplerコアは残念ながら整数演算が かなり遅いのです。この点は次の記事の「命令別スループット」の 項目で詳しく解説されています。 GTX680のグラフィック・GPGPU性能を調べる ≪ dokumaru dokumaru.wordpress.com/2012/03/27/gtx680-spec/ Keplerではゲームで使われる単精度演算以外はほとんど無視して 性能を稼いでいるので、GPGPU的にはかなり残念なことになっています。 653 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 03:13:17.23 ID:dQ9rq2KX0] >>649 580とあまり遜色のない速度が出ていますね。 自分の部屋ではGeForceが3枚24時間フル稼働してますｗ CUDA版の開発も続けるので安心して下さい。 654 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 09:58:39.20 ID:dQ9rq2KX0] ここ数カ月のjohn-devでのOpenCLでのBitslice DESの実装についての やり取りを追って見たのですが、なかなか面白かったです。 www.openwall.com/lists/john-dev/2012/09 www.openwall.com/lists/john-dev/2012/10 www.openwall.com/lists/john-dev/2012/11 現在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). www.openwall.com/lists/john-dev/2012/10/14/1 手元のTripcode Finderのコードは現在のJtRの実装より大分速いのですが、 それでも300M TPSには遠く及びません。レジスタ数にもまだ大分余裕があるし、 工夫の余地はいろいろあるのでしょう。実に奥が深いです。 655 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 13:08:08.52 ID:dQ9rq2KX0] Bitslice DESをマルチスレッド化したときにエンバクした模様。 結構な確率で間違ったトリップが出力されます。 CUDAと同じコードのはずなんですけど、barrier()がうまく動作してないの かしらん。 しかしこれ、どうやってデバッグするんだろう… 656 名前：ののたん ◆KiwamonoL. mailto:sage [2012/11/15(木) 14:35:25.04 ID:et60Xlt20] >>655 昔ながらの printf でおｋ。 手段として美しくないのは嫌いとかなら知らん。 657 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 14:47:56.84 ID:dQ9rq2KX0] やっぱりそれしかないんですねorz 658 名前：ののたん ◆KiwamonoL. mailto:sage [2012/11/15(木) 15:05:07.72 ID:et60Xlt20] >>657 私が hip2 つくってた頃は printf すらなかったのに。 贅沢ね。 659 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 15:15:33.54 ID:dQ9rq2KX0] >>658 まあそりゃそうなんですけどね… 660 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 15:18:14.38 ID:dQ9rq2KX0] あ、原因分かったかも。CUDA版を書いてたときに適当だったところが 今になって問題になっているのかもしれません。 661 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 15:27:59.89 ID:dQ9rq2KX0] う〜ん、違うな… もうちょっと全体的に腐ってる感じです。 662 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 15:38:19.52 ID:dQ9rq2KX0] まあいいや。マルチスレッド化の作業はまた明日やり直すことにしよっと。 663 名前： ◆MERIKEN4.k mailto:sage [2012/11/15(木) 15:54:46.53 ID:dQ9rq2KX0] コードをロールバックしたらちゃんと動作するようなのでやっぱり マルチスレッド化が原因のようです。マルチスレッド化すると 速度が倍近くになるので次はなんとか成功させたいところです。 664 名前：前スレ927 ◆JouJaku.HzIz mailto:sage [2012/11/15(木) 21:27:47.25 ID: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 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/16(金) 01:59:34.00 ID:QPHBSAhn0] 電源容量が足りないんじゃ 666 名前： ◆MERIKEN4.k mailto:sage [2012/11/16(金) 04:42:53.68 ID:eP2LlovM0] OpenCLの10桁検索ですが、もうちょっと調べたらどうも移植した直後から 問題があったようです。APP Profilerがメモリリークを報告しているので もうちょっと調べてみます。 667 名前： ◆MERIKEN4.k mailto:sage [2012/11/16(金) 04:45:44.48 ID:eP2LlovM0] >>664 温度に問題がないなら電源の可能性が高いですね。 電源は何を使われていますか? 668 名前： ◆MERIKEN4.k mailto:sage [2012/11/16(金) 06:40:48.74 ID:eP2LlovM0] どうやら問題はBitslice DESの処理そのものではなく 他の処理にある模様。ちゃんと出力をチェックするルーチンを 作りこんで、徹底的にテストするしかないようです。 やなよかんはしてたけど、やはり10桁検索は楽ではないですねえ。 669 名前：前スレ927 ◆JouJaku.HzIz mailto:sage [2012/11/16(金) 23:39:50.31 ID:SdQXCd/P0] 電源が届くのを待ちきれなくて、無理矢理繋げて実行していました。 電力不足でこんな挙動をするとは初体験で全然知らず。お恥ずかしい限りです。 素直に電源届くまで待っています。 670 名前： ◆MERIKEN4.k mailto:sage [2012/11/17(土) 09:52:44.80 ID:Kz7friKn0] >>669 そりゃそこにカードがあれば試したくなりますよね。 その気持、わかりますｗ 電源が届いたらまたぜひ報告して下さい。 671 名前： ◆MERIKEN4.k mailto:sage [2012/11/17(土) 10:07:54.25 ID:Kz7friKn0] OpenCLの10桁検索の出力が腐っていた問題ですが、カーネルをすこしづつ削って 原因を探ったところ、結果を書き込む__globalの配列へのアクセスの前後に barrier()を入れてやると問題が出なくなることが分かりました。 Bitslice DES用の一時変数を__privateに置いても直らなかったし、 CUDA版やOpenCLの12桁検索では全く問題がなかった部分なので、 AMDのOpenCLの実装のバグの可能性が非常に高いです。 AMDの実装は性能は出るのにいちいち造りが甘くて非常にもったいない 感じがします。ここらへんもCUDAのほうが任期がある理由なんでしょうねえ。 672 名前： ◆MERIKEN4.k mailto:sage [2012/11/17(土) 12:33:38.78 ID:Kz7friKn0] この件でコードをロールバックした時に気がついたのですが、 Bitslice DESの一時変数を__private空間においても割と速度が出ることが わかりました。こっちのほうが__localよりもベクトル化しやすいので、 このまま__localを使わずに最適化をすすめることにします。 Bitslice DESの深さを32bitから128bitにして速度も4倍といきたい ところですが… 673 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/17(土) 12:53:15.66 ID:CDs2gwHh0] >>672 >ベクトル化 よく知らないのですが、GPUってベクトル演算なんですか……？ ベクトル化の意味は知っているのですが、なぜか「昔のスパコン」ってイメージが……ｗ 674 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/17(土) 17:05:57.82 ID:RbPdKj5Y0] GPUはベクトル演算の極地だし、今のスパコンはほぼ全てベクトル演算ですが 675 名前：名無しさん＠お腹いっぱい。 mailto:sage [2012/11/17(土) 17:08:42.00 ID:lv9DVzeD0] もの自体がベクタプロセッサの集合体 676 名前： ◆MERIKEN4.k mailto:sage [2012/11/18(日) 01:33:23.65 ID:7lmxdB8G0] >>673 そこがGPGPUの一番美味しいところですｗ 性能を引き出すのはなかなか難しいですけどね〜 677 名前： ◆MERIKEN4.k mailto:sage [2012/11/18(日) 01:59:41.80 ID: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. www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/barrier.html CUDAの場合は動機が必要なところで__syncthreads()を呼び出してやれば 後はなにも考えずに共有メモリとグローバルメモリに読み書きできたのですが、 どうも勝手が違うようです。

---

---

[ 続きを読む ] / [ 携帯版 ]

read.cgi ver5.27 [feat.BBS2 +1.6] / e.0.2 (02/09/03) / eucaly.net products.
担当:undef