github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-12-23 14:13:43 +01:00
Code Releases Activity

Increase Oracle T: Type (Oracle 12+) cracking performance

Browse Source
This commit is contained in:
jsteube 2016-05-28 13:02:21 +02:00
parent c0e40632bb
commit 333ca26f21
1 changed files with 83 additions and 72 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

155
OpenCL/m12300.cl
View File

@ -176,6 +176,59 @@ void hmac_run (const u64 w1[16], const u64 ipad[8], const u64 opad[8], u64 dgst[
sha512_transform (w, dgst);
}
void hmac_run_x (const u64 ipad[8], const u64 opad[8], u64 dgst[8])
{
u64 w[16];
w[ 0] = dgst[0];
w[ 1] = dgst[1];
w[ 2] = dgst[2];
w[ 3] = dgst[3];
w[ 4] = dgst[4];
w[ 5] = dgst[5];
w[ 6] = dgst[6];
w[ 7] = dgst[7];
w[ 8] = 0x8000000000000000;
w[ 9] = 0;
w[10] = 0;
w[11] = 0;
w[12] = 0;
w[13] = 0;
w[14] = 0;
w[15] = (128 + 64) * 8;
dgst[0] = ipad[0];
dgst[1] = ipad[1];
dgst[2] = ipad[2];
dgst[3] = ipad[3];
dgst[4] = ipad[4];
dgst[5] = ipad[5];
dgst[6] = ipad[6];
dgst[7] = ipad[7];
sha512_transform (w, dgst);
w[ 0] = dgst[0];
w[ 1] = dgst[1];
w[ 2] = dgst[2];
w[ 3] = dgst[3];
w[ 4] = dgst[4];
w[ 5] = dgst[5];
w[ 6] = dgst[6];
w[ 7] = dgst[7];
dgst[0] = opad[0];
dgst[1] = opad[1];
dgst[2] = opad[2];
dgst[3] = opad[3];
dgst[4] = opad[4];
dgst[5] = opad[5];
dgst[6] = opad[6];
dgst[7] = opad[7];
sha512_transform (w, dgst);
}
void hmac_init (u64 w[16], u64 ipad[8], u64 opad[8])
{
w[ 0] ^= 0x3636363636363636;
@ -392,81 +445,39 @@ __kernel void m12300_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf
opad[6] = tmps[gid].opad[6];
opad[7] = tmps[gid].opad[7];
for (u32 i = 0; i < 8; i += 8)
u64 dgst[8];
dgst[0] = tmps[gid].dgst[0];
dgst[1] = tmps[gid].dgst[1];
dgst[2] = tmps[gid].dgst[2];
dgst[3] = tmps[gid].dgst[3];
dgst[4] = tmps[gid].dgst[4];
dgst[5] = tmps[gid].dgst[5];
dgst[6] = tmps[gid].dgst[6];
dgst[7] = tmps[gid].dgst[7];
for (u32 j = 0; j < loop_cnt; j++)
{
u64 dgst[8];
hmac_run_x (ipad, opad, dgst);
dgst[0] = tmps[gid].dgst[i + 0];
dgst[1] = tmps[gid].dgst[i + 1];
dgst[2] = tmps[gid].dgst[i + 2];
dgst[3] = tmps[gid].dgst[i + 3];
dgst[4] = tmps[gid].dgst[i + 4];
dgst[5] = tmps[gid].dgst[i + 5];
dgst[6] = tmps[gid].dgst[i + 6];
dgst[7] = tmps[gid].dgst[i + 7];
u64 out[8];
out[0] = tmps[gid].out[i + 0];
out[1] = tmps[gid].out[i + 1];
out[2] = tmps[gid].out[i + 2];
out[3] = tmps[gid].out[i + 3];
out[4] = tmps[gid].out[i + 4];
out[5] = tmps[gid].out[i + 5];
out[6] = tmps[gid].out[i + 6];
out[7] = tmps[gid].out[i + 7];
for (u32 j = 0; j < loop_cnt; j++)
{
u64 w[16];
w[ 0] = dgst[0];
w[ 1] = dgst[1];
w[ 2] = dgst[2];
w[ 3] = dgst[3];
w[ 4] = dgst[4];
w[ 5] = dgst[5];
w[ 6] = dgst[6];
w[ 7] = dgst[7];
w[ 8] = 0x8000000000000000;
w[ 9] = 0;
w[10] = 0;
w[11] = 0;
w[12] = 0;
w[13] = 0;
w[14] = 0;
w[15] = (128 + 64) * 8;
hmac_run (w, ipad, opad, dgst);
out[0] ^= dgst[0];
out[1] ^= dgst[1];
out[2] ^= dgst[2];
out[3] ^= dgst[3];
out[4] ^= dgst[4];
out[5] ^= dgst[5];
out[6] ^= dgst[6];
out[7] ^= dgst[7];
}
tmps[gid].dgst[i + 0] = dgst[0];
tmps[gid].dgst[i + 1] = dgst[1];
tmps[gid].dgst[i + 2] = dgst[2];
tmps[gid].dgst[i + 3] = dgst[3];
tmps[gid].dgst[i + 4] = dgst[4];
tmps[gid].dgst[i + 5] = dgst[5];
tmps[gid].dgst[i + 6] = dgst[6];
tmps[gid].dgst[i + 7] = dgst[7];
tmps[gid].out[i + 0] = out[0];
tmps[gid].out[i + 1] = out[1];
tmps[gid].out[i + 2] = out[2];
tmps[gid].out[i + 3] = out[3];
tmps[gid].out[i + 4] = out[4];
tmps[gid].out[i + 5] = out[5];
tmps[gid].out[i + 6] = out[6];
tmps[gid].out[i + 7] = out[7];
tmps[gid].out[0] ^= dgst[0];
tmps[gid].out[1] ^= dgst[1];
tmps[gid].out[2] ^= dgst[2];
tmps[gid].out[3] ^= dgst[3];
tmps[gid].out[4] ^= dgst[4];
tmps[gid].out[5] ^= dgst[5];
tmps[gid].out[6] ^= dgst[6];
tmps[gid].out[7] ^= dgst[7];
}
tmps[gid].dgst[0] = dgst[0];
tmps[gid].dgst[1] = dgst[1];
tmps[gid].dgst[2] = dgst[2];
tmps[gid].dgst[3] = dgst[3];
tmps[gid].dgst[4] = dgst[4];
tmps[gid].dgst[5] = dgst[5];
tmps[gid].dgst[6] = dgst[6];
tmps[gid].dgst[7] = dgst[7];
}
__kernel void m12300_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global oraclet_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)