arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-01-23 14:10:57 +00:00
Code Issues Releases Wiki Activity

Fixed cracking of NetNTLMv1 passwords in mask-attack mode if mask > length 16 (optimized kernels only)

Browse Source
This commit is contained in:
jsteube 2019-03-05 05:14:34 +01:00
parent 277db2ecee
commit eb07d2108c
2 changed files with 141 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

140
OpenCL/m05500_a3-optimized.cl
View File

@ -969,6 +969,76 @@ __kernel void m05500_m08 (KERN_ATTR_VECTOR ())
__kernel void m05500_m16 (KERN_ATTR_VECTOR ())
{
/**
* modifier
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
/**
* sbox, kbox
*/
__local u32 s_SPtrans[8][64];
__local u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
s_SPtrans[0][i] = c_SPtrans[0][i];
s_SPtrans[1][i] = c_SPtrans[1][i];
s_SPtrans[2][i] = c_SPtrans[2][i];
s_SPtrans[3][i] = c_SPtrans[3][i];
s_SPtrans[4][i] = c_SPtrans[4][i];
s_SPtrans[5][i] = c_SPtrans[5][i];
s_SPtrans[6][i] = c_SPtrans[6][i];
s_SPtrans[7][i] = c_SPtrans[7][i];
s_skb[0][i] = c_skb[0][i];
s_skb[1][i] = c_skb[1][i];
s_skb[2][i] = c_skb[2][i];
s_skb[3][i] = c_skb[3][i];
s_skb[4][i] = c_skb[4][i];
s_skb[5][i] = c_skb[5][i];
s_skb[6][i] = c_skb[6][i];
s_skb[7][i] = c_skb[7][i];
}
barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
/**
* base
*/
u32 w[16];
w[ 0] = pws[gid].i[ 0];
w[ 1] = pws[gid].i[ 1];
w[ 2] = pws[gid].i[ 2];
w[ 3] = pws[gid].i[ 3];
w[ 4] = pws[gid].i[ 4];
w[ 5] = pws[gid].i[ 5];
w[ 6] = pws[gid].i[ 6];
w[ 7] = pws[gid].i[ 7];
w[ 8] = pws[gid].i[ 8];
w[ 9] = pws[gid].i[ 9];
w[10] = pws[gid].i[10];
w[11] = pws[gid].i[11];
w[12] = pws[gid].i[12];
w[13] = pws[gid].i[13];
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
*/
m05500m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}
__kernel void m05500_s04 (KERN_ATTR_VECTOR ())
@ -1121,4 +1191,74 @@ __kernel void m05500_s08 (KERN_ATTR_VECTOR ())
__kernel void m05500_s16 (KERN_ATTR_VECTOR ())
{
/**
* modifier
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
/**
* sbox, kbox
*/
__local u32 s_SPtrans[8][64];
__local u32 s_skb[8][64];
for (u32 i = lid; i < 64; i += lsz)
{
s_SPtrans[0][i] = c_SPtrans[0][i];
s_SPtrans[1][i] = c_SPtrans[1][i];
s_SPtrans[2][i] = c_SPtrans[2][i];
s_SPtrans[3][i] = c_SPtrans[3][i];
s_SPtrans[4][i] = c_SPtrans[4][i];
s_SPtrans[5][i] = c_SPtrans[5][i];
s_SPtrans[6][i] = c_SPtrans[6][i];
s_SPtrans[7][i] = c_SPtrans[7][i];
s_skb[0][i] = c_skb[0][i];
s_skb[1][i] = c_skb[1][i];
s_skb[2][i] = c_skb[2][i];
s_skb[3][i] = c_skb[3][i];
s_skb[4][i] = c_skb[4][i];
s_skb[5][i] = c_skb[5][i];
s_skb[6][i] = c_skb[6][i];
s_skb[7][i] = c_skb[7][i];
}
barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
/**
* base
*/
u32 w[16];
w[ 0] = pws[gid].i[ 0];
w[ 1] = pws[gid].i[ 1];
w[ 2] = pws[gid].i[ 2];
w[ 3] = pws[gid].i[ 3];
w[ 4] = pws[gid].i[ 4];
w[ 5] = pws[gid].i[ 5];
w[ 6] = pws[gid].i[ 6];
w[ 7] = pws[gid].i[ 7];
w[ 8] = pws[gid].i[ 8];
w[ 9] = pws[gid].i[ 9];
w[10] = pws[gid].i[10];
w[11] = pws[gid].i[11];
w[12] = pws[gid].i[12];
w[13] = pws[gid].i[13];
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
*/
m05500s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
}

1
docs/changes.txt
View File

@ -22,6 +22,7 @@
- Fixed output of IKE PSK (mode 5300 and 5400) hashes to have separators at right position
- Fixed the validation of the --brain-client-features command line argument (only values 1, 2 or 3 are allowed)
- Fixed cracking of Cisco-PIX and Cisco-ASA MD5 passwords in mask-attack mode if mask > length 16
- Fixed cracking of NetNTLMv1 passwords in mask-attack mode if mask > length 16 (optimized kernels only)
- Fixed the 7-Zip parser to allow the entire supported range of encrypted and decrypted data lengths
- Fixed incorrect progress-only result in a special race condition
- Fixed maximum password length limit which was announced as 256 but actually was 255