github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-12-09 02:13:10 +01:00
Code Releases Activity

Increase 1Password, cloudkeychain cracking performance

Browse Source
This commit is contained in:
jsteube 2016-05-27 23:59:52 +02:00
parent b5fca5f223
commit 2ec7429dea
2 changed files with 83 additions and 75 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
OpenCL/inc_vendor.cl
View File

@ -101,9 +101,6 @@
#if KERN_TYPE == 7900 #if KERN_TYPE == 7900
#undef _unroll #undef _unroll
#endif #endif
#if KERN_TYPE == 8200
#undef _unroll
#endif
#if KERN_TYPE == 10400 #if KERN_TYPE == 10400
#undef _unroll #undef _unroll
#endif #endif

155
OpenCL/m08200.cl
View File

@ -385,6 +385,59 @@ void hmac_sha512_run (const u64 w1[16], const u64 ipad[8], const u64 opad[8], u6
sha512_transform (w, dgst); sha512_transform (w, dgst);
} }
void hmac_sha512_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_sha512_init (u64 w[16], u64 ipad[8], u64 opad[8]) void hmac_sha512_init (u64 w[16], u64 ipad[8], u64 opad[8])
{ {
w[ 0] ^= 0x3636363636363636; w[ 0] ^= 0x3636363636363636;
@ -617,81 +670,39 @@ __kernel void m08200_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf
opad[6] = tmps[gid].opad[6]; opad[6] = tmps[gid].opad[6];
opad[7] = tmps[gid].opad[7]; 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_sha512_run_x (ipad, opad, dgst);
dgst[0] = tmps[gid].dgst[i + 0]; tmps[gid].out[0] ^= dgst[0];
dgst[1] = tmps[gid].dgst[i + 1]; tmps[gid].out[1] ^= dgst[1];
dgst[2] = tmps[gid].dgst[i + 2]; tmps[gid].out[2] ^= dgst[2];
dgst[3] = tmps[gid].dgst[i + 3]; tmps[gid].out[3] ^= dgst[3];
dgst[4] = tmps[gid].dgst[i + 4]; tmps[gid].out[4] ^= dgst[4];
dgst[5] = tmps[gid].dgst[i + 5]; tmps[gid].out[5] ^= dgst[5];
dgst[6] = tmps[gid].dgst[i + 6]; tmps[gid].out[6] ^= dgst[6];
dgst[7] = tmps[gid].dgst[i + 7]; tmps[gid].out[7] ^= dgst[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_sha512_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].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 m08200_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha512_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 cloudkey_t *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) __kernel void m08200_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha512_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 cloudkey_t *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)