🛠️ 🧝🏻 👩🏼‍🎤 パスワードハッキングを使用したOpenCLの探索 👿 🚼 🧙🏾

エントリー

最近、GPGPUに関するさまざまな記事やプレゼンテーションを読んだ後、ビデオカードのプログラミングも自分で試してみることにしました。実際、この分野での技術の選択は素晴らしいものではありません-CUDA（nVidia独自の標準）とOpenCL（無料の標準、ATIのGPU、nVidia、および中央処理装置）のみがまだ生きており、開発中です。私のラップトップにはATIグラフィックカード（Mobility Radeon 5650 HD）があるため、選択肢は完全に1つのオプション-OpenCLに限定されました。この記事では、OpenCLをゼロから学習するプロセスと、その起源について説明します。

OpenCLおよびPyOpenClの概要

一見すると、Cの制御コードと、いわゆるカーネル（カーネル）のコードの両方が、非常に混乱しているように見えました。提供されているC APIでは、特に少なくともいくつかのエラーの処理では、単純なプログラムの起動でさえ多くの行を必要とするため、より便利で人道的なものを見つけたいと思いました。その選択はPyOpenCLライブラリにかかっていました。その名前から、制御コードはPythonで書かれていることが明らかです。 OpenCLコードを初めて見た人でも、すべてがより理解しやすいように見えます（もちろん、これは単純な例にのみ適用されます）。ただし、カーネル自体のコードはまだわずかに変更されたCで記述されているため、まだ学習する必要があります。それに関する完全なドキュメントは、標準の開発者のWebサイト（ Khronos ）で入手でき、特定の実装に関する情報は、それぞれATIおよびnVidia Webサイトで入手できます。

最も簡単な例（2つの配列を追加）を使用して、言語の第一印象を得ることができます。

__kernel void sum(__global const float *a, __global const float *b, __global float *c) { int gid = get_global_id(0); c[gid] = a[gid] + b[gid]; }

そして、この例を実行して検証するために必要な完全なコードは次のとおりです（PyOpenCLドキュメントから取得）。

非表示のテキスト

 import pyopencl as cl import numpy import numpy.linalg as la a = numpy.random.rand(50000).astype(numpy.float32) b = numpy.random.rand(50000).astype(numpy.float32) ctx = cl.create_some_context() queue = cl.CommandQueue(ctx) mf = cl.mem_flags a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a) b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b) dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes) prg = cl.Program(ctx, """ __kernel void sum(__global const float *a, __global const float *b, __global float *c) { int gid = get_global_id(0); c[gid] = a[gid] + b[gid]; } """).build() prg.sum(queue, a.shape, None, a_buf, b_buf, dest_buf) a_plus_b = numpy.empty_like(a) cl.enqueue_copy(queue, a_plus_b, dest_buf) print la.norm(a_plus_b - (a+b))

特定の行がすぐに表示されます：コンテキストの作成、実行キュー、デバイスへのバッファーの作成とコピー、カーネル自体のコンパイルと起動。 OpenCLのコンテキストとキューの詳細についてはドキュメントを参照してください。比較的単純なプログラムの場合、必要なキューとコンテキストは1つだけです。これらは、例と非常によく似た行で作成されます。一般に、OpenCLプログラムの計算の構造は次のようになります。

コンテキストの作成、キュー、プログラムのコンパイル
実行中に変更されないデバイスデータ（バッファ）へのコピー
サイクル
- この反復に固有のデータをデバイスにコピーする
- カーネル実行
- 計算されたデータをメインメモリにコピーして、何らかの処理が可能

SHA1ハッシュ

レベルを下げて、カーネルコードがどのように機能するかを理解するときが来ました。 OpenCL関数を外部から実行するには、__ kernel属性で指定する必要があり、値の型がvoidであり、直接値（int、float4、...）または領域へのポインターである特定の数の引数が必要ですメモリ__global、__ constant、__ local。また、便宜上、カーネルから呼び出すことができる他の関数をプログラムで宣言できますが、これはパフォーマンスに影響を与えません。すべての関数は自動的に置き換えられます（つまり、インラインディレクティブと同様）。これは、OpenCLの再帰がまったくサポートされていないためです。

OpenCLが変更されたC言語であるという事実を使用して、SHA1などのハッシュ関数の既製の実装を使用し、わずかな変更を使用できます。

非表示のテキスト

 #define K1 0x5A827999 #define K2 0x6ED9EBA1 #define K3 0x8F1BBCDC #define K4 0xCA62C1D6 #define a0 0x67452301; #define b0 0xEFCDAB89; #define c0 0x98BADCFE; #define d0 0x10325476; #define e0 0xC3D2E1F0; #define f1(x,y,z) ( z ^ ( x & ( y ^ z ) ) ) /* Rounds 0-19 */ #define f2(x,y,z) ( x ^ y ^ z ) /* Rounds 20-39 */ #define f3(x,y,z) ( ( x & y ) | ( z & ( x | y ) ) ) /* Rounds 40-59 */ #define f4(x,y,z) ( x ^ y ^ z ) /* Rounds 60-79 */ #define ROTL(n,X) ( ( ( X ) << n ) | ( ( X ) >> ( 32 - n ) ) ) #define expand(W,i) ( W[ i & 15 ] = ROTL( 1, ( W[ i & 15 ] ^ W[ (i - 14) & 15 ] ^ \ W[ (i - 8) & 15 ] ^ W[ (i - 3) & 15 ] ) ) ) #define subRound(a, b, c, d, e, f, k, data) \ ( e += ROTL( 5, a ) + f( b, c, d ) + k + data, b = ROTL( 30, b ) ) #define REVERSE(value) value = ((value & 0xFF000000) >> 24) | ((value & 0x00FF0000) >> 8) | ((value & 0x0000FF00) << 8) | ((value & 0x000000FF) << 24) long sha1(uint *eData, const int length) { unsigned int A = a0; unsigned int B = b0; unsigned int C = c0; unsigned int D = d0; unsigned int E = e0; ((__local char *)eData)[length] = 0x80; for (int i = 0; i <= length / 4; i++) { REVERSE(eData[i]); } eData[14] = 0; eData[15] = length * 8; subRound( A, B, C, D, E, f1, K1, eData[ 0 ] ); subRound( E, A, B, C, D, f1, K1, eData[ 1 ] ); subRound( D, E, A, B, C, f1, K1, eData[ 2 ] ); subRound( C, D, E, A, B, f1, K1, eData[ 3 ] ); subRound( B, C, D, E, A, f1, K1, eData[ 4 ] ); subRound( A, B, C, D, E, f1, K1, eData[ 5 ] ); subRound( E, A, B, C, D, f1, K1, eData[ 6 ] ); subRound( D, E, A, B, C, f1, K1, eData[ 7 ] ); subRound( C, D, E, A, B, f1, K1, eData[ 8 ] ); subRound( B, C, D, E, A, f1, K1, eData[ 9 ] ); subRound( A, B, C, D, E, f1, K1, eData[ 10 ] ); subRound( E, A, B, C, D, f1, K1, eData[ 11 ] ); subRound( D, E, A, B, C, f1, K1, eData[ 12 ] ); subRound( C, D, E, A, B, f1, K1, eData[ 13 ] ); subRound( B, C, D, E, A, f1, K1, eData[ 14 ] ); subRound( A, B, C, D, E, f1, K1, eData[ 15 ] ); subRound( E, A, B, C, D, f1, K1, expand( eData, 16 ) ); subRound( D, E, A, B, C, f1, K1, expand( eData, 17 ) ); subRound( C, D, E, A, B, f1, K1, expand( eData, 18 ) ); subRound( B, C, D, E, A, f1, K1, expand( eData, 19 ) ); subRound( A, B, C, D, E, f2, K2, expand( eData, 20 ) ); subRound( E, A, B, C, D, f2, K2, expand( eData, 21 ) ); subRound( D, E, A, B, C, f2, K2, expand( eData, 22 ) ); subRound( C, D, E, A, B, f2, K2, expand( eData, 23 ) ); subRound( B, C, D, E, A, f2, K2, expand( eData, 24 ) ); subRound( A, B, C, D, E, f2, K2, expand( eData, 25 ) ); subRound( E, A, B, C, D, f2, K2, expand( eData, 26 ) ); subRound( D, E, A, B, C, f2, K2, expand( eData, 27 ) ); subRound( C, D, E, A, B, f2, K2, expand( eData, 28 ) ); subRound( B, C, D, E, A, f2, K2, expand( eData, 29 ) ); subRound( A, B, C, D, E, f2, K2, expand( eData, 30 ) ); subRound( E, A, B, C, D, f2, K2, expand( eData, 31 ) ); subRound( D, E, A, B, C, f2, K2, expand( eData, 32 ) ); subRound( C, D, E, A, B, f2, K2, expand( eData, 33 ) ); subRound( B, C, D, E, A, f2, K2, expand( eData, 34 ) ); subRound( A, B, C, D, E, f2, K2, expand( eData, 35 ) ); subRound( E, A, B, C, D, f2, K2, expand( eData, 36 ) ); subRound( D, E, A, B, C, f2, K2, expand( eData, 37 ) ); subRound( C, D, E, A, B, f2, K2, expand( eData, 38 ) ); subRound( B, C, D, E, A, f2, K2, expand( eData, 39 ) ); subRound( A, B, C, D, E, f3, K3, expand( eData, 40 ) ); subRound( E, A, B, C, D, f3, K3, expand( eData, 41 ) ); subRound( D, E, A, B, C, f3, K3, expand( eData, 42 ) ); subRound( C, D, E, A, B, f3, K3, expand( eData, 43 ) ); subRound( B, C, D, E, A, f3, K3, expand( eData, 44 ) ); subRound( A, B, C, D, E, f3, K3, expand( eData, 45 ) ); subRound( E, A, B, C, D, f3, K3, expand( eData, 46 ) ); subRound( D, E, A, B, C, f3, K3, expand( eData, 47 ) ); subRound( C, D, E, A, B, f3, K3, expand( eData, 48 ) ); subRound( B, C, D, E, A, f3, K3, expand( eData, 49 ) ); subRound( A, B, C, D, E, f3, K3, expand( eData, 50 ) ); subRound( E, A, B, C, D, f3, K3, expand( eData, 51 ) ); subRound( D, E, A, B, C, f3, K3, expand( eData, 52 ) ); subRound( C, D, E, A, B, f3, K3, expand( eData, 53 ) ); subRound( B, C, D, E, A, f3, K3, expand( eData, 54 ) ); subRound( A, B, C, D, E, f3, K3, expand( eData, 55 ) ); subRound( E, A, B, C, D, f3, K3, expand( eData, 56 ) ); subRound( D, E, A, B, C, f3, K3, expand( eData, 57 ) ); subRound( C, D, E, A, B, f3, K3, expand( eData, 58 ) ); subRound( B, C, D, E, A, f3, K3, expand( eData, 59 ) ); subRound( A, B, C, D, E, f4, K4, expand( eData, 60 ) ); subRound( E, A, B, C, D, f4, K4, expand( eData, 61 ) ); subRound( D, E, A, B, C, f4, K4, expand( eData, 62 ) ); subRound( C, D, E, A, B, f4, K4, expand( eData, 63 ) ); subRound( B, C, D, E, A, f4, K4, expand( eData, 64 ) ); subRound( A, B, C, D, E, f4, K4, expand( eData, 65 ) ); subRound( E, A, B, C, D, f4, K4, expand( eData, 66 ) ); subRound( D, E, A, B, C, f4, K4, expand( eData, 67 ) ); subRound( C, D, E, A, B, f4, K4, expand( eData, 68 ) ); subRound( B, C, D, E, A, f4, K4, expand( eData, 69 ) ); subRound( A, B, C, D, E, f4, K4, expand( eData, 70 ) ); subRound( E, A, B, C, D, f4, K4, expand( eData, 71 ) ); subRound( D, E, A, B, C, f4, K4, expand( eData, 72 ) ); subRound( C, D, E, A, B, f4, K4, expand( eData, 73 ) ); subRound( B, C, D, E, A, f4, K4, expand( eData, 74 ) ); subRound( A, B, C, D, E, f4, K4, expand( eData, 75 ) ); subRound( E, A, B, C, D, f4, K4, expand( eData, 76 ) ); subRound( D, E, A, B, C, f4, K4, expand( eData, 77 ) ); subRound( C, D, E, A, B, f4, K4, expand( eData, 78 ) ); subRound( B, C, D, E, A, f4, K4, expand( eData, 79 ) ); A += a0; B += b0; C += c0; D += d0; E += e0; return as_ulong((uint2)(D, E)); }

ここでいくつかの明確化が必要です。クラックのための「実験的な」パスワードハッシュとして、私はリークされたLinkedInハッシュを取りました。これはほぼ600万（一意）です。リスト内の存在をかなり迅速に確認するためのいくつかのオプションがあります。ハッシュテーブルを使用しました（詳細は後ほど）。メモリ消費を削減し、SHA1の20バイト全体ではなく、最後の8バイト、つまり 1つのlong / ulong値。もちろん、これは誤った一致の可能性を高めますが、それは非常に小さいままです。私が持っていたすべてのパスワードの中で、私はそのようなケースを6つしか持っていませんでした。したがって、切り捨てられた値（最後の8バイト）は、上記の関数からすぐに返されます。それ以外の場合、すべてが標準であり、SHA1アルゴリズムは、56バイト未満の文字列の場合の仕様に従って直接実装されます。

検索の構成

次に、検索自体を整理する必要があります。最も単純なオプションは、すべての位置に対して同じ文字セットのブルートフォースであり、これは、たとえば同様の方法で直接実装できます。

非表示のテキスト

 __kernel void do_brute( __global const long *table, const ulong start_num, __global ulong *result, __global uint *res_ind) { char s[64]; uint *s_l = (__local uint *)s; int i, j; ulong _n, n; ulong start = start_num + get_global_id(0) * HASHES_PER_WORKITEM; for (j = 0; j < HASHES_PER_WORKITEM; j++) { n = _n = j + start; for (i = 15; i >= 0; i--) { s_l[i] = 0; } for (i = COMB_LEN - 1; i >= 0; i--) { s[i] = charset[n % CHARS_CNT]; n /= CHARS_CNT; } if (check_in_table(table, sha1(s_l, COMB_LEN))) { result[atomic_inc(res_ind)] = _n; } } }

ここで、HASHES_PER_WORKITEMは、1回の実行で1つの作業項目（ストリーム）で処理されるハッシュの数、COMB_LENは組み合わせの長さ、charsetは文字の配列、CHARS_CNTは配列内の文字の数です。ご覧のとおり、起動時に、ハッシュテーブルへのポインター、検索を開始するパスワード番号、および結果を表示する配列へのポインターとその中のインデックスがこのカーネルに渡されます。

OpenCLでは、一度に1つのスレッドが起動されるわけではありませんが、その一部はグローバルワークサイズと呼ばれ、すべてのスレッドは同じ引数を受け取ります。それらのそれぞれがキースペースの一部を整理するために、文字列ulong start = start_num + get_global_id(0) * HASHES_PER_WORKITEM;

特定のスレッドの番号を計算します（get_global_id（0）は、0から現在のグローバル作業サイズまでのスレッドインデックスを返す標準関数です）。

次に、各ストリームHASHES_PER_WORKITEMはパスワードを列挙します。各パスワードはsha1関数によってハッシュ化され、後述するcheck_in_table関数の存在を確認します。この形式で、ハッシュテーブルの最も単純な実装では、1秒あたり約2,000万のパスワードの結果が得られました。たとえば、ラップトップで3億以上のハッシュを生成するoclHascat（検証が多数のハッシュのリストを通過する場合でも）これを正当化するものではありません）。今後は、単純なブルートフォースで1秒間に1億6,000万の速度を達成できたと言います。これはoclHascatの半分の速度（ハッシュ1つ）です。

ハッシュテーブル

そのため、ハッシュの存在を確認します。実装された最初のオプションは、オープンアドレス指定の最も単純なハッシュテーブルでした。問題を複雑にしないように、プロセッサとビデオカードではなく、OpenCLでは、要求のみがありました。このケースは次のようになりました。

非表示のテキスト

 bool check_in_table( __global const long *table, const long value) { uint index = calc_index(value); uint step = calc_step(value); for (uint a = 1; ; a++) { index %= TABLE_SIZE; if (table[index] == 0) { return false; } if (table[index] == value) { return true; } index += a * step + 1; } }

さまざまなテーブルサイズとプローブ方法を試しましたが、速度はそれほど向上しませんでした。 GPUハッシュテーブルに関する資料を探して、特定のCuckoo Hashtable （ロシア語への翻訳が確立されているかどうかはわかりません）について言及しているVasily Volkovの記事「GPUでの効率的なハッシュテーブルの構築」に出会いました。聞かなかった。簡単に言えば、その本質は、1つではなくいくつかのハッシュ関数を使用することと、特別な塗りつぶし方法にあります。その後、k個以下のメモリアクセスに対して要素が見つかります。kはハッシュ関数の数です。占有メモリよりも速度が重要なので、k = 2を使用しました。それを埋めることはCPUでも発生します。

おわりに

また、もちろん、最適化はコードの別の部分、つまりパスワードの生成に影響しました。上記のバージョンでは、いくつかの最適ではない場所がすぐに表示されます。たとえば、次の各パスワードは最初から生成されますが、前のものは変更できます。特にOpenCLに固有の最適化には他の場所があります：より高速なローカルの代わりに、文字の配列にグローバルメモリまたは定数メモリを使用します（特定の実装の開発者から直接メモリ領域の詳細を読む方が良いです）ただし、カーネルコードのさまざまな最適化に関する個別の記事を書く価値はありますが、ここでは、GPU向けのプログラミングを行う場合、さまざまなオプションを試してその速度を調べる価値があると言います。目で言うと、動作が速くなるとは限りません。特定の命令を削除しても、実行速度が大幅に低下する場合があります。

将来的には、位置ごとに異なるアルファベットのサポートを追加しました。さらに重要なことは、個々の文字だけでなく、単語のアルファベットもサポートすることです。利便性と柔軟性のために、カーネルコードはMakoテンプレートエンジンによって処理されます。これらはすべてアーカイブにあります（以下を参照）。

結論

だから、私が最終的に取得したもの：

多くの微妙な点を備えたOpenCLの知識（最初から、かなり成功していると思います）
Pythonでのプログラミングのレベルを上げました（まったくゼロからではなく、かなり低いレベルから）
関連技術を学び、試してみました：Cythonコンパイラー、Makoテンプレートエンジン、VCS git

有形の結果：

プログラムは、ラップトップビデオカードで単純なブルートフォース（文字ではなく単語のアルファベットを使用する場合、速度は毎秒5,000万から）のために、毎秒約1億6千万の速度でパスワードを列挙します-一方、oclHascatで3億と比較できます1つはハッシュ、もう1つはCPUの最速ブルートフォースプロセッサ（i5 2.5 GHzを搭載）で3000万個、アセンブラーとSSE命令を使用して記述
それを使用して、トリックなしで、LinkedInハッシュから約250万のパスワードが選択されました（あまり多くはありませんが、ここでは記録を目指していませんでした）
ここからダウンロードできます（一部の場所には恐ろしい糞があります！）。開始するには、Python 2.7、numpy、PyOpenCL、Cython、Mako、およびLinkedInハッシュ付きのファイルが必要です。

印象

Pythonは非常に遅い言語ですが（書くのは非常に簡単です）、ループで大きなリストを処理するのは非常に遅く、さまざまな理解度です。これがCythonがプログラムの一部に使用される理由です-ところで、本当に便利なこと
プロファイリングとその後の最適化に対する適切なアプローチを備えたOpenCLは非常に迅速に機能するため、アセンブラーで何かを書き換える必要はありません。

PS：

さまざまなコメントや推奨事項を歓迎します。

パスワードハッキングを使用したOpenCLの探索

エントリー

OpenCLおよびPyOpenClの概要

SHA1ハッシュ

検索の構成

ハッシュテーブル

おわりに

結論

PS：

More articles: