mirror of
https://github.com/hashcat/hashcat.git
synced 2024-12-27 00:48:13 +00:00
Preparing PR
This commit is contained in:
Fist0urs
parent
40bbb0023c
commit
7ff09c6710
@ -1042,7 +1042,6 @@ typedef struct
|
||||
u32 SID[32];
|
||||
u32 SID_len;
|
||||
u32 SID_offset;
|
||||
u32 userKey[5];
|
||||
char cipher_algo[16];
|
||||
char hash_algo[16];
|
||||
|
||||
@ -1399,12 +1398,14 @@ typedef struct
|
||||
u32 dgst[10];
|
||||
u32 out[10];
|
||||
|
||||
u32 userKey[5];
|
||||
|
||||
/* dedicated to hmac-sha512 */
|
||||
u64 ipad64[8];
|
||||
u64 opad64[8];
|
||||
u64 dgst64[16];
|
||||
u64 out64[16];
|
||||
|
||||
|
||||
} dpapimk_tmp_t;
|
||||
|
||||
typedef struct
|
||||
|
||||
391
OpenCL/m15300.cl
391
OpenCL/m15300.cl
@ -698,9 +698,6 @@ void AES256_decrypt (const u32 *in, u32 *out, const u32 *rdk, SHM_TYPE u32 *s_td
|
||||
^ rdk[59];
|
||||
}
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
|
||||
void md4_transform_S (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
|
||||
{
|
||||
u32 a = digest[0];
|
||||
@ -1501,38 +1498,38 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
u32 w0[4];
|
||||
|
||||
w0[0] = pws[gid].i[ 0];
|
||||
w0[1] = pws[gid].i[ 1];
|
||||
w0[2] = pws[gid].i[ 2];
|
||||
w0[3] = pws[gid].i[ 3];
|
||||
w0[0] = pws[gid].i[0];
|
||||
w0[1] = pws[gid].i[1];
|
||||
w0[2] = pws[gid].i[2];
|
||||
w0[3] = pws[gid].i[3];
|
||||
|
||||
u32 w1[4];
|
||||
|
||||
w1[0] = pws[gid].i[ 4];
|
||||
w1[1] = pws[gid].i[ 5];
|
||||
w1[2] = pws[gid].i[ 6];
|
||||
w1[3] = pws[gid].i[ 7];
|
||||
w1[0] = pws[gid].i[4];
|
||||
w1[1] = pws[gid].i[5];
|
||||
w1[2] = pws[gid].i[6];
|
||||
w1[3] = pws[gid].i[7];
|
||||
|
||||
u32 w2[4];
|
||||
|
||||
w2[0] = pws[gid].i[ 8];
|
||||
w2[1] = pws[gid].i[ 9];
|
||||
w2[2] = pws[gid].i[10];
|
||||
w2[3] = pws[gid].i[11];
|
||||
w2[0] = 0;
|
||||
w2[1] = 0;
|
||||
w2[2] = 0;
|
||||
w2[3] = 0;
|
||||
|
||||
u32 w3[4];
|
||||
|
||||
w3[0] = pws[gid].i[12];
|
||||
w3[1] = pws[gid].i[13];
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
w3[0] = 0;
|
||||
w3[1] = 0;
|
||||
w3[2] = 0;
|
||||
w3[3] = 0;
|
||||
|
||||
u32 pw_len = pws[gid].pw_len;
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
append_0x80_4x4_S (w0, w1, w2, w3, pw_len);
|
||||
append_0x80_2x4_S (w0, w1, pw_len);
|
||||
|
||||
make_unicode (w1, w2, w3);
|
||||
make_unicode (w0, w0, w1);
|
||||
make_unicode_S (w1, w2, w3);
|
||||
make_unicode_S (w0, w0, w1);
|
||||
|
||||
/**
|
||||
* main
|
||||
@ -1555,23 +1552,23 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
digest_context[3] = SHA1M_D;
|
||||
digest_context[4] = SHA1M_E;
|
||||
|
||||
w0[0] = swap32 (w0[0]);
|
||||
w0[1] = swap32 (w0[1]);
|
||||
w0[2] = swap32 (w0[2]);
|
||||
w0[3] = swap32 (w0[3]);
|
||||
w0[0] = swap32_S (w0[0]);
|
||||
w0[1] = swap32_S (w0[1]);
|
||||
w0[2] = swap32_S (w0[2]);
|
||||
w0[3] = swap32_S (w0[3]);
|
||||
|
||||
w1[0] = swap32 (w1[0]);
|
||||
w1[1] = swap32 (w1[1]);
|
||||
w1[2] = swap32 (w1[2]);
|
||||
w1[3] = swap32 (w1[3]);
|
||||
w1[0] = swap32_S (w1[0]);
|
||||
w1[1] = swap32_S (w1[1]);
|
||||
w1[2] = swap32_S (w1[2]);
|
||||
w1[3] = swap32_S (w1[3]);
|
||||
|
||||
w2[0] = swap32 (w2[0]);
|
||||
w2[1] = swap32 (w2[1]);
|
||||
w2[2] = swap32 (w2[2]);
|
||||
w2[3] = swap32 (w2[3]);
|
||||
w2[0] = swap32_S (w2[0]);
|
||||
w2[1] = swap32_S (w2[1]);
|
||||
w2[2] = swap32_S (w2[2]);
|
||||
w2[3] = swap32_S (w2[3]);
|
||||
|
||||
w3[0] = swap32 (w3[0]);
|
||||
w3[1] = swap32 (w3[1]);
|
||||
w3[0] = swap32_S (w3[0]);
|
||||
w3[1] = swap32_S (w3[1]);
|
||||
w3[2] = 0;
|
||||
w3[3] = pw_len * 2 * 8;
|
||||
|
||||
@ -1595,7 +1592,6 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
digest_context[3] = swap32_S (digest_context[3]);
|
||||
digest_context[4] = 0;
|
||||
}
|
||||
|
||||
/* initialize hmac-sha1 */
|
||||
|
||||
/**
|
||||
@ -1681,13 +1677,12 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
hmac_sha1_run_S (w0, w1, w2, w3, ipad, opad, key);
|
||||
/* this key is used as password for pbkdf2-hmac-* */
|
||||
|
||||
esalt_bufs[digests_offset].userKey[0] = key[0];
|
||||
esalt_bufs[digests_offset].userKey[1] = key[1];
|
||||
esalt_bufs[digests_offset].userKey[2] = key[2];
|
||||
esalt_bufs[digests_offset].userKey[3] = key[3];
|
||||
esalt_bufs[digests_offset].userKey[4] = key[4];
|
||||
|
||||
|
||||
tmps[gid].userKey[0] = key[0];
|
||||
tmps[gid].userKey[1] = key[1];
|
||||
tmps[gid].userKey[2] = key[2];
|
||||
tmps[gid].userKey[3] = key[3];
|
||||
tmps[gid].userKey[4] = key[4];
|
||||
|
||||
/* if DPAPImk version 1, pbkdf-hmac-sha1 is used */
|
||||
if (esalt_bufs[digests_offset].version == 1)
|
||||
{
|
||||
@ -1757,13 +1752,13 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
tmps[gid].dgst[3] = digest[3];
|
||||
tmps[gid].dgst[4] = digest[4];
|
||||
|
||||
tmps[gid].out[0] = digest[0];
|
||||
tmps[gid].out[1] = digest[1];
|
||||
tmps[gid].out[2] = digest[2];
|
||||
tmps[gid].out[3] = digest[3];
|
||||
tmps[gid].out[4] = digest[4];
|
||||
tmps[gid].out[0] = digest[0];
|
||||
tmps[gid].out[1] = digest[1];
|
||||
tmps[gid].out[2] = digest[2];
|
||||
tmps[gid].out[3] = digest[3];
|
||||
tmps[gid].out[4] = digest[4];
|
||||
|
||||
/* We need bigger output ! */
|
||||
/* We need bigger output! */
|
||||
w0[0] = esalt_bufs[digests_offset].iv[0];
|
||||
w0[1] = esalt_bufs[digests_offset].iv[1];
|
||||
w0[2] = esalt_bufs[digests_offset].iv[2];
|
||||
@ -1789,11 +1784,11 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
tmps[gid].dgst[8] = digest[3];
|
||||
tmps[gid].dgst[9] = digest[4];
|
||||
|
||||
tmps[gid].out[5] = digest[0];
|
||||
tmps[gid].out[6] = digest[1];
|
||||
tmps[gid].out[7] = digest[2];
|
||||
tmps[gid].out[8] = digest[3];
|
||||
tmps[gid].out[9] = digest[4];
|
||||
tmps[gid].out[5] = digest[0];
|
||||
tmps[gid].out[6] = digest[1];
|
||||
tmps[gid].out[7] = digest[2];
|
||||
tmps[gid].out[8] = digest[3];
|
||||
tmps[gid].out[9] = digest[4];
|
||||
}
|
||||
/* if DPAPImk version 2, pbkdf-hmac-sha512 is used*/
|
||||
else if (esalt_bufs[digests_offset].version == 2)
|
||||
@ -1873,14 +1868,14 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
tmps[gid].dgst64[6] = dgst64[6];
|
||||
tmps[gid].dgst64[7] = dgst64[7];
|
||||
|
||||
tmps[gid].out64[0] = dgst64[0];
|
||||
tmps[gid].out64[1] = dgst64[1];
|
||||
tmps[gid].out64[2] = dgst64[2];
|
||||
tmps[gid].out64[3] = dgst64[3];
|
||||
tmps[gid].out64[4] = dgst64[4];
|
||||
tmps[gid].out64[5] = dgst64[5];
|
||||
tmps[gid].out64[6] = dgst64[6];
|
||||
tmps[gid].out64[7] = dgst64[7];
|
||||
tmps[gid].out64[0] = dgst64[0];
|
||||
tmps[gid].out64[1] = dgst64[1];
|
||||
tmps[gid].out64[2] = dgst64[2];
|
||||
tmps[gid].out64[3] = dgst64[3];
|
||||
tmps[gid].out64[4] = dgst64[4];
|
||||
tmps[gid].out64[5] = dgst64[5];
|
||||
tmps[gid].out64[6] = dgst64[6];
|
||||
tmps[gid].out64[7] = dgst64[7];
|
||||
}
|
||||
}
|
||||
|
||||
@ -1919,18 +1914,18 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
{
|
||||
u32x dgst[5];
|
||||
u32x out[5];
|
||||
|
||||
|
||||
dgst[0] = packv (tmps, dgst, gid, i + 0);
|
||||
dgst[1] = packv (tmps, dgst, gid, i + 1);
|
||||
dgst[2] = packv (tmps, dgst, gid, i + 2);
|
||||
dgst[3] = packv (tmps, dgst, gid, i + 3);
|
||||
dgst[4] = packv (tmps, dgst, gid, i + 4);
|
||||
|
||||
out[0] = packv (tmps, out, gid, i + 0);
|
||||
out[1] = packv (tmps, out, gid, i + 1);
|
||||
out[2] = packv (tmps, out, gid, i + 2);
|
||||
out[3] = packv (tmps, out, gid, i + 3);
|
||||
out[4] = packv (tmps, out, gid, i + 4);
|
||||
|
||||
out[0] = packv (tmps, out, gid, i + 0);
|
||||
out[1] = packv (tmps, out, gid, i + 1);
|
||||
out[2] = packv (tmps, out, gid, i + 2);
|
||||
out[3] = packv (tmps, out, gid, i + 3);
|
||||
out[4] = packv (tmps, out, gid, i + 4);
|
||||
|
||||
for (u32 j = 0; j < loop_cnt; j++)
|
||||
{
|
||||
@ -1938,7 +1933,7 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
u32x w1[4];
|
||||
u32x w2[4];
|
||||
u32x w3[4];
|
||||
|
||||
|
||||
/* Microsoft PBKDF2 implementation. On purpose?
|
||||
Misunderstanding of them? Dunno...
|
||||
*/
|
||||
@ -1958,22 +1953,22 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[1] = 0;
|
||||
w3[2] = 0;
|
||||
w3[3] = (64 + 20) * 8;
|
||||
|
||||
|
||||
hmac_sha1_run_V (w0, w1, w2, w3, ipad, opad, dgst);
|
||||
|
||||
|
||||
out[0] ^= dgst[0];
|
||||
out[1] ^= dgst[1];
|
||||
out[2] ^= dgst[2];
|
||||
out[3] ^= dgst[3];
|
||||
out[4] ^= dgst[4];
|
||||
}
|
||||
|
||||
|
||||
unpackv (tmps, dgst, gid, i + 0, dgst[0]);
|
||||
unpackv (tmps, dgst, gid, i + 1, dgst[1]);
|
||||
unpackv (tmps, dgst, gid, i + 2, dgst[2]);
|
||||
unpackv (tmps, dgst, gid, i + 3, dgst[3]);
|
||||
unpackv (tmps, dgst, gid, i + 4, dgst[4]);
|
||||
|
||||
|
||||
unpackv (tmps, out, gid, i + 0, out[0]);
|
||||
unpackv (tmps, out, gid, i + 1, out[1]);
|
||||
unpackv (tmps, out, gid, i + 2, out[2]);
|
||||
@ -2016,14 +2011,14 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
dgst[6] = pack64v (tmps, dgst64, gid, 6);
|
||||
dgst[7] = pack64v (tmps, dgst64, gid, 7);
|
||||
|
||||
out[0] = pack64v (tmps, out64, gid, 0);
|
||||
out[1] = pack64v (tmps, out64, gid, 1);
|
||||
out[2] = pack64v (tmps, out64, gid, 2);
|
||||
out[3] = pack64v (tmps, out64, gid, 3);
|
||||
out[4] = pack64v (tmps, out64, gid, 4);
|
||||
out[5] = pack64v (tmps, out64, gid, 5);
|
||||
out[6] = pack64v (tmps, out64, gid, 6);
|
||||
out[7] = pack64v (tmps, out64, gid, 7);
|
||||
out[0] = pack64v (tmps, out64, gid, 0);
|
||||
out[1] = pack64v (tmps, out64, gid, 1);
|
||||
out[2] = pack64v (tmps, out64, gid, 2);
|
||||
out[3] = pack64v (tmps, out64, gid, 3);
|
||||
out[4] = pack64v (tmps, out64, gid, 4);
|
||||
out[5] = pack64v (tmps, out64, gid, 5);
|
||||
out[6] = pack64v (tmps, out64, gid, 6);
|
||||
out[7] = pack64v (tmps, out64, gid, 7);
|
||||
|
||||
for (u32 j = 0; j < loop_cnt; j++)
|
||||
{
|
||||
@ -2032,7 +2027,7 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
u64x w2[4];
|
||||
u64x w3[4];
|
||||
|
||||
/* Microsoft PBKDF2 implementation. On purpose?
|
||||
/* Microsoft PBKDF2 implementation. On purpose?
|
||||
Misunderstanding of them? Dunno...
|
||||
*/
|
||||
w0[0] = out[0];
|
||||
@ -2096,7 +2091,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
|
||||
u32 ipad[5];
|
||||
u32 opad[5];
|
||||
|
||||
@ -2106,7 +2101,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
__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];
|
||||
@ -2117,7 +2112,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
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];
|
||||
@ -2127,13 +2122,13 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
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;
|
||||
|
||||
u32 key[6];
|
||||
|
||||
|
||||
key[0] = swap32_S (tmps[gid].out[0]);
|
||||
key[1] = swap32_S (tmps[gid].out[1]);
|
||||
key[2] = swap32_S (tmps[gid].out[2]);
|
||||
@ -2142,102 +2137,102 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
key[5] = swap32_S (tmps[gid].out[5]);
|
||||
|
||||
u32x iv[2];
|
||||
|
||||
iv[0] = swap32_S (tmps[gid].out[6]);
|
||||
iv[1] = swap32_S (tmps[gid].out[7]);
|
||||
|
||||
iv[0] = swap32_S (tmps[gid].out[6]);
|
||||
iv[1] = swap32_S (tmps[gid].out[7]);
|
||||
|
||||
u32 decrypted[26];
|
||||
|
||||
|
||||
/* Construct 3DES keys */
|
||||
|
||||
|
||||
const u32x a = (key[0]);
|
||||
const u32x b = (key[1]);
|
||||
|
||||
|
||||
u32x Ka[16];
|
||||
u32x Kb[16];
|
||||
|
||||
|
||||
_des_crypt_keysetup (a, b, Ka, Kb, s_skb);
|
||||
|
||||
|
||||
const u32x c = (key[2]);
|
||||
const u32x d = (key[3]);
|
||||
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
|
||||
const u32x e = (key[4]);
|
||||
const u32x f = (key[5]);
|
||||
|
||||
u32x Ke[16];
|
||||
u32x Kf[16];
|
||||
|
||||
|
||||
_des_crypt_keysetup (e, f, Ke, Kf, s_skb);
|
||||
|
||||
u32 contents_pos;
|
||||
u32 contents_off;
|
||||
u32 wx_off;
|
||||
|
||||
|
||||
for (wx_off = 0, contents_pos = 0, contents_off = 0; contents_pos < esalt_bufs[digests_offset].contents_len; wx_off += 2, contents_pos += 8, contents_off += 2)
|
||||
{
|
||||
/* First Pass */
|
||||
|
||||
|
||||
u32x data[2];
|
||||
|
||||
|
||||
data[0] = (u32x) swap32_S (esalt_bufs[digests_offset].contents[contents_off + 0]);
|
||||
data[1] = (u32x) swap32_S (esalt_bufs[digests_offset].contents[contents_off + 1]);
|
||||
|
||||
u32x p1[2];
|
||||
|
||||
|
||||
_des_crypt_decrypt (p1, data, Ke, Kf, s_SPtrans);
|
||||
|
||||
/* Second Pass */
|
||||
|
||||
|
||||
u32x p2[2];
|
||||
|
||||
|
||||
_des_crypt_encrypt (p2, p1, Kc, Kd, s_SPtrans);
|
||||
|
||||
|
||||
/* Third Pass */
|
||||
|
||||
u32x out[2];
|
||||
|
||||
|
||||
_des_crypt_decrypt (out, p2, Ka, Kb, s_SPtrans);
|
||||
|
||||
|
||||
out[0] ^= iv[0];
|
||||
out[1] ^= iv[1];
|
||||
|
||||
decrypted[wx_off + 0] = out[0];
|
||||
decrypted[wx_off + 1] = out[1];
|
||||
|
||||
|
||||
iv[0] = data[0];
|
||||
iv[1] = data[1];
|
||||
}
|
||||
|
||||
u32 hmacSalt[4];
|
||||
u32 hmacTemoin[5];
|
||||
u32 expectedHmac[5];
|
||||
u32 lastKey[16];
|
||||
|
||||
hmacSalt[0] = swap32_S (decrypted[0]);
|
||||
hmacSalt[1] = swap32_S (decrypted[1]);
|
||||
hmacSalt[2] = swap32_S (decrypted[2]);
|
||||
hmacSalt[3] = swap32_S (decrypted[3]);
|
||||
hmacSalt[0] = swap32_S (decrypted[0]);
|
||||
hmacSalt[1] = swap32_S (decrypted[1]);
|
||||
hmacSalt[2] = swap32_S (decrypted[2]);
|
||||
hmacSalt[3] = swap32_S (decrypted[3]);
|
||||
|
||||
hmacTemoin[0] = swap32_S (decrypted[0 + 4]);
|
||||
hmacTemoin[1] = swap32_S (decrypted[1 + 4]);
|
||||
hmacTemoin[2] = swap32_S (decrypted[2 + 4]);
|
||||
hmacTemoin[3] = swap32_S (decrypted[3 + 4]);
|
||||
hmacTemoin[4] = swap32_S (decrypted[4 + 4]);
|
||||
expectedHmac[0] = swap32_S (decrypted[4 + 0]);
|
||||
expectedHmac[1] = swap32_S (decrypted[4 + 1]);
|
||||
expectedHmac[2] = swap32_S (decrypted[4 + 2]);
|
||||
expectedHmac[3] = swap32_S (decrypted[4 + 3]);
|
||||
expectedHmac[4] = swap32_S (decrypted[4 + 4]);
|
||||
|
||||
for(int i = 0; i < 16; i++)
|
||||
{
|
||||
lastKey[i] = decrypted[i + 26 - 16];
|
||||
lastKey[i] = decrypted[i + 26 - 16];
|
||||
}
|
||||
|
||||
w0[0] = esalt_bufs[digests_offset].userKey[0];
|
||||
w0[1] = esalt_bufs[digests_offset].userKey[1];
|
||||
w0[2] = esalt_bufs[digests_offset].userKey[2];
|
||||
w0[3] = esalt_bufs[digests_offset].userKey[3];
|
||||
w1[0] = esalt_bufs[digests_offset].userKey[4];
|
||||
w0[0] = tmps[gid].userKey[0];
|
||||
w0[1] = tmps[gid].userKey[1];
|
||||
w0[2] = tmps[gid].userKey[2];
|
||||
w0[3] = tmps[gid].userKey[3];
|
||||
w1[0] = tmps[gid].userKey[4];
|
||||
w1[1] = 0;
|
||||
w1[2] = 0;
|
||||
w1[3] = 0;
|
||||
@ -2293,7 +2288,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3[1] = 0;
|
||||
w3[2] = 0;
|
||||
w3[3] = 0;
|
||||
|
||||
|
||||
hmac_sha1_pad_S (w0, w1, w2, w3, ipad, opad);
|
||||
|
||||
/**
|
||||
@ -2340,11 +2335,13 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
if ((hmacTemoin[0] == digest[0]) && (hmacTemoin[1] == digest[1])
|
||||
&& (hmacTemoin[2] == digest[2]) && (hmacTemoin[3] == digest[3])
|
||||
&& (hmacTemoin[4] == digest[4]))
|
||||
if ( expectedHmac[0] == digest[0]
|
||||
&& expectedHmac[1] == digest[1]
|
||||
&& expectedHmac[2] == digest[2]
|
||||
&& expectedHmac[3] == digest[3]
|
||||
&& expectedHmac[4] == digest[4])
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
|
||||
}
|
||||
}
|
||||
else if (esalt_bufs[digests_offset].version == 2)
|
||||
@ -2352,21 +2349,21 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
/**
|
||||
* aes shared
|
||||
*/
|
||||
|
||||
|
||||
#ifdef REAL_SHM
|
||||
|
||||
|
||||
__local u32 s_td0[256];
|
||||
__local u32 s_td1[256];
|
||||
__local u32 s_td2[256];
|
||||
__local u32 s_td3[256];
|
||||
__local u32 s_td4[256];
|
||||
|
||||
|
||||
__local u32 s_te0[256];
|
||||
__local u32 s_te1[256];
|
||||
__local u32 s_te2[256];
|
||||
__local u32 s_te3[256];
|
||||
__local u32 s_te4[256];
|
||||
|
||||
|
||||
for (u32 i = lid; i < 256; i += lsz)
|
||||
{
|
||||
s_td0[i] = td0[i];
|
||||
@ -2374,38 +2371,38 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
s_td2[i] = td2[i];
|
||||
s_td3[i] = td3[i];
|
||||
s_td4[i] = td4[i];
|
||||
|
||||
|
||||
s_te0[i] = te0[i];
|
||||
s_te1[i] = te1[i];
|
||||
s_te2[i] = te2[i];
|
||||
s_te3[i] = te3[i];
|
||||
s_te4[i] = te4[i];
|
||||
}
|
||||
|
||||
|
||||
barrier (CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
|
||||
#else
|
||||
|
||||
|
||||
__constant u32a *s_td0 = td0;
|
||||
__constant u32a *s_td1 = td1;
|
||||
__constant u32a *s_td2 = td2;
|
||||
__constant u32a *s_td3 = td3;
|
||||
__constant u32a *s_td4 = td4;
|
||||
|
||||
|
||||
__constant u32a *s_te0 = te0;
|
||||
__constant u32a *s_te1 = te1;
|
||||
__constant u32a *s_te2 = te2;
|
||||
__constant u32a *s_te3 = te3;
|
||||
__constant u32a *s_te4 = te4;
|
||||
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
|
||||
/* Construct AES key */
|
||||
|
||||
|
||||
u32 key[8];
|
||||
|
||||
|
||||
key[0] = h32_from_64 (tmps[gid].out64[0]);
|
||||
key[1] = l32_from_64 (tmps[gid].out64[0]);
|
||||
key[2] = h32_from_64 (tmps[gid].out64[1]);
|
||||
@ -2414,88 +2411,88 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
key[5] = l32_from_64 (tmps[gid].out64[2]);
|
||||
key[6] = h32_from_64 (tmps[gid].out64[3]);
|
||||
key[7] = l32_from_64 (tmps[gid].out64[3]);
|
||||
|
||||
|
||||
u32 iv[4];
|
||||
|
||||
iv[0] = h32_from_64 (tmps[gid].out64[4]);
|
||||
iv[1] = l32_from_64 (tmps[gid].out64[4]);
|
||||
iv[2] = h32_from_64 (tmps[gid].out64[5]);
|
||||
iv[3] = l32_from_64 (tmps[gid].out64[5]);
|
||||
|
||||
|
||||
iv[0] = h32_from_64 (tmps[gid].out64[4]);
|
||||
iv[1] = l32_from_64 (tmps[gid].out64[4]);
|
||||
iv[2] = h32_from_64 (tmps[gid].out64[5]);
|
||||
iv[3] = l32_from_64 (tmps[gid].out64[5]);
|
||||
|
||||
#define KEYLEN 60
|
||||
|
||||
|
||||
u32 rek[KEYLEN];
|
||||
|
||||
|
||||
AES256_ExpandKey (key, rek, s_te0, s_te1, s_te2, s_te3, s_te4);
|
||||
|
||||
|
||||
u32 rdk[KEYLEN];
|
||||
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < KEYLEN; i++) rdk[i] = rek[i];
|
||||
|
||||
|
||||
AES256_InvertKey (rdk, s_td0, s_td1, s_td2, s_td3, s_td4, s_te0, s_te1, s_te2, s_te3, s_te4);
|
||||
|
||||
|
||||
/* 144 bytes */
|
||||
u32 decrypted[36] = { 0 };
|
||||
|
||||
|
||||
u32 contents_pos;
|
||||
u32 contents_off;
|
||||
u32 wx_off;
|
||||
|
||||
|
||||
for (wx_off = 0, contents_pos = 0, contents_off = 0; contents_pos < esalt_bufs[digests_offset].contents_len; wx_off += 4, contents_pos += 16, contents_off += 4)
|
||||
{
|
||||
u32 data[4];
|
||||
|
||||
|
||||
data[0] = esalt_bufs[digests_offset].contents[contents_off + 0];
|
||||
data[1] = esalt_bufs[digests_offset].contents[contents_off + 1];
|
||||
data[2] = esalt_bufs[digests_offset].contents[contents_off + 2];
|
||||
data[3] = esalt_bufs[digests_offset].contents[contents_off + 3];
|
||||
|
||||
|
||||
u32 out[4];
|
||||
|
||||
|
||||
AES256_decrypt (data, out, rdk, s_td0, s_td1, s_td2, s_td3, s_td4);
|
||||
|
||||
|
||||
out[0] ^= iv[0];
|
||||
out[1] ^= iv[1];
|
||||
out[2] ^= iv[2];
|
||||
out[3] ^= iv[3];
|
||||
|
||||
|
||||
decrypted[wx_off + 0] = out[0];
|
||||
decrypted[wx_off + 1] = out[1];
|
||||
decrypted[wx_off + 2] = out[2];
|
||||
decrypted[wx_off + 3] = out[3];
|
||||
|
||||
|
||||
iv[0] = data[0];
|
||||
iv[1] = data[1];
|
||||
iv[2] = data[2];
|
||||
iv[3] = data[3];
|
||||
}
|
||||
|
||||
|
||||
u32 hmacSalt[4];
|
||||
u32 hmacTemoin[16];
|
||||
u32 expectedHmac[16];
|
||||
u32 lastKey[16];
|
||||
|
||||
|
||||
hmacSalt[0] = decrypted[0];
|
||||
hmacSalt[1] = decrypted[1];
|
||||
hmacSalt[2] = decrypted[2];
|
||||
hmacSalt[3] = decrypted[3];
|
||||
|
||||
|
||||
for(int i = 0; i < 16; i++)
|
||||
{
|
||||
hmacTemoin[i] = decrypted[i + 4];
|
||||
lastKey[i] = decrypted[i + 36 - 16];
|
||||
expectedHmac[i] = decrypted[i + 4];
|
||||
lastKey[i] = decrypted[i + 36 - 16];
|
||||
}
|
||||
|
||||
|
||||
u64 w0_x64[4];
|
||||
u64 w1_x64[4];
|
||||
u64 w2_x64[4];
|
||||
u64 w3_x64[4];
|
||||
|
||||
w0_x64[0] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[0], esalt_bufs[digests_offset].userKey[1]);
|
||||
w0_x64[1] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[2], esalt_bufs[digests_offset].userKey[3]);
|
||||
w0_x64[2] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[4], 0);
|
||||
|
||||
w0_x64[0] = hl32_to_64_S (tmps[gid].userKey[0], tmps[gid].userKey[1]);
|
||||
w0_x64[1] = hl32_to_64_S (tmps[gid].userKey[2], tmps[gid].userKey[3]);
|
||||
w0_x64[2] = hl32_to_64_S (tmps[gid].userKey[4], 0);
|
||||
w0_x64[3] = 0;
|
||||
w1_x64[0] = 0;
|
||||
w1_x64[1] = 0;
|
||||
@ -2509,12 +2506,12 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3_x64[1] = 0;
|
||||
w3_x64[2] = 0;
|
||||
w3_x64[3] = 0;
|
||||
|
||||
|
||||
u64 ipad64[8];
|
||||
u64 opad64[8];
|
||||
|
||||
|
||||
hmac_sha512_pad_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64);
|
||||
|
||||
|
||||
w0_x64[0] = hl32_to_64_S (hmacSalt[0], hmacSalt[1]);
|
||||
w0_x64[1] = hl32_to_64_S (hmacSalt[2], hmacSalt[3]);
|
||||
w0_x64[2] = hl32_to_64_S (0x80000000, 0);
|
||||
@ -2531,13 +2528,13 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3_x64[1] = 0;
|
||||
w3_x64[2] = 0;
|
||||
w3_x64[3] = (128 + 16) * 8;
|
||||
|
||||
|
||||
u64 dgst64[8];
|
||||
|
||||
|
||||
hmac_sha512_run_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64, dgst64);
|
||||
|
||||
|
||||
u64 encKey[8];
|
||||
|
||||
|
||||
encKey[0] = dgst64[0];
|
||||
encKey[1] = dgst64[1];
|
||||
encKey[2] = dgst64[2];
|
||||
@ -2546,7 +2543,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
encKey[5] = dgst64[5];
|
||||
encKey[6] = dgst64[6];
|
||||
encKey[7] = dgst64[7];
|
||||
|
||||
|
||||
w0_x64[0] = encKey[0];
|
||||
w0_x64[1] = encKey[1];
|
||||
w0_x64[2] = encKey[2];
|
||||
@ -2563,9 +2560,9 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3_x64[1] = 0;
|
||||
w3_x64[2] = 0;
|
||||
w3_x64[3] = 0;
|
||||
|
||||
|
||||
hmac_sha512_pad_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64);
|
||||
|
||||
|
||||
w0_x64[0] = hl32_to_64_S (lastKey[ 0], lastKey[ 1]);
|
||||
w0_x64[1] = hl32_to_64_S (lastKey[ 2], lastKey[ 3]);
|
||||
w0_x64[2] = hl32_to_64_S (lastKey[ 4], lastKey[ 5]);
|
||||
@ -2582,15 +2579,21 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
w3_x64[1] = 0;
|
||||
w3_x64[2] = 0;
|
||||
w3_x64[3] = (128 + 16 * 4) * 8;
|
||||
|
||||
|
||||
hmac_sha512_run_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64, dgst64);
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
if ((hmacTemoin[0] == h32_from_64 (dgst64[0])) && (hmacTemoin[1] == l32_from_64 (dgst64[0]))
|
||||
&& (hmacTemoin[14] == h32_from_64 (dgst64[7])) && (hmacTemoin[15] == l32_from_64 (dgst64[7])))
|
||||
if ( expectedHmac[ 0] == h32_from_64 (dgst64[0])
|
||||
&& expectedHmac[ 1] == l32_from_64 (dgst64[0])
|
||||
&& expectedHmac[ 2] == h32_from_64 (dgst64[1])
|
||||
&& expectedHmac[ 3] == l32_from_64 (dgst64[1])
|
||||
&& expectedHmac[12] == h32_from_64 (dgst64[6])
|
||||
&& expectedHmac[13] == l32_from_64 (dgst64[6])
|
||||
&& expectedHmac[14] == h32_from_64 (dgst64[7])
|
||||
&& expectedHmac[15] == l32_from_64 (dgst64[7]))
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -6,6 +6,7 @@
|
||||
|
||||
- Added hash-mode 600 = Blake2-512
|
||||
- Added hash-mode 15200 = Blockchain, My Wallet, V2
|
||||
- Added hash-mode 15300 = DPAPI masterkey file v1 and v2
|
||||
|
||||
##
|
||||
## Features
|
||||
|
||||
@ -31,6 +31,7 @@ Jean-Christophe "Fist0urs" Delaunay <jean-christophe.delaunay@synacktiv.com> (@F
|
||||
* Kerberos TGS Rep enctype 23 kernel module
|
||||
* AxCrypt kernel module
|
||||
* KeePass kernel module
|
||||
* DPAPImk v1 and v2 kernel module
|
||||
|
||||
Other contributors to hashcat
|
||||
|
||||
|
||||
@ -145,6 +145,7 @@ NVIDIA GPUs require "NVIDIA Driver" (367.x or later)
|
||||
- NTLM
|
||||
- Domain Cached Credentials (DCC), MS Cache
|
||||
- Domain Cached Credentials 2 (DCC2), MS Cache 2
|
||||
- DPAPI masterkey file v1 and v2
|
||||
- MS-AzureSync PBKDF2-HMAC-SHA256
|
||||
- descrypt
|
||||
- bsdicrypt
|
||||
|
||||
@ -176,7 +176,7 @@ _hashcat ()
|
||||
{
|
||||
local VERSION=3.5.0
|
||||
|
||||
local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 200 300 400 500 501 600 900 1000 1100 1400 1410 1411 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12001 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14700 14800 14900 15000 15100 15200"
|
||||
local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 200 300 400 500 501 600 900 1000 1100 1400 1410 1411 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12001 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14700 14800 14900 15000 15100 15200 15300"
|
||||
local ATTACK_MODES="0 1 3 6 7"
|
||||
local HCCAPX_MESSAGE_PAIR="0 1 2 3 4 5"
|
||||
local OUTFILE_FORMATS="1 2 3 4 5 6 7 8 9 10 11 12 13 14 15"
|
||||
|
||||
@ -429,7 +429,6 @@ typedef struct dpapimk
|
||||
u32 SID[32];
|
||||
u32 SID_len;
|
||||
u32 SID_offset;
|
||||
u32 userKey[5];
|
||||
u8 cipher_algo[16];
|
||||
u8 hash_algo[16];
|
||||
|
||||
@ -837,12 +836,14 @@ typedef struct dpapimk_tmp
|
||||
u32 dgst[10];
|
||||
u32 out[10];
|
||||
|
||||
u32 userKey[5];
|
||||
|
||||
/* dedicated to hmac-sha512 */
|
||||
u64 ipad64[8];
|
||||
u64 opad64[8];
|
||||
u64 dgst64[16];
|
||||
u64 out64[16];
|
||||
|
||||
|
||||
} dpapimk_tmp_t;
|
||||
|
||||
typedef struct seven_zip_hook
|
||||
@ -1366,7 +1367,8 @@ typedef enum hash_type
|
||||
HASH_TYPE_ITUNES_BACKUP_10 = 57,
|
||||
HASH_TYPE_SKIP32 = 58,
|
||||
HASH_TYPE_BLAKE2B = 59,
|
||||
HASH_TYPE_DPAPIMK = 60,
|
||||
HASH_TYPE_CHACHA20 = 60,
|
||||
HASH_TYPE_DPAPIMK = 61,
|
||||
|
||||
} hash_type_t;
|
||||
|
||||
|
||||
@ -85,6 +85,7 @@ const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_BUF[] =
|
||||
1000,
|
||||
1100,
|
||||
2100,
|
||||
15300,
|
||||
12800,
|
||||
1500,
|
||||
12400,
|
||||
|
||||
@ -2208,16 +2208,16 @@ static u32 parse_and_store_salt (u8 *out, u8 *in, u32 salt_len, MAYBE_UNUSED con
|
||||
{
|
||||
if (salt_len < 20)
|
||||
{
|
||||
tmp_u32[ 9] = ((tmp_u32[ 4] >> 8) & 0x00FF0000) | ((tmp_u32[ 4] >> 16) & 0x000000FF);
|
||||
tmp_u32[ 8] = ((tmp_u32[ 4] << 8) & 0x00FF0000) | ((tmp_u32[ 4] >> 0) & 0x000000FF);
|
||||
tmp_u32[ 7] = ((tmp_u32[ 3] >> 8) & 0x00FF0000) | ((tmp_u32[ 3] >> 16) & 0x000000FF);
|
||||
tmp_u32[ 6] = ((tmp_u32[ 3] << 8) & 0x00FF0000) | ((tmp_u32[ 3] >> 0) & 0x000000FF);
|
||||
tmp_u32[ 5] = ((tmp_u32[ 2] >> 8) & 0x00FF0000) | ((tmp_u32[ 2] >> 16) & 0x000000FF);
|
||||
tmp_u32[ 4] = ((tmp_u32[ 2] << 8) & 0x00FF0000) | ((tmp_u32[ 2] >> 0) & 0x000000FF);
|
||||
tmp_u32[ 3] = ((tmp_u32[ 1] >> 8) & 0x00FF0000) | ((tmp_u32[ 1] >> 16) & 0x000000FF);
|
||||
tmp_u32[ 2] = ((tmp_u32[ 1] << 8) & 0x00FF0000) | ((tmp_u32[ 1] >> 0) & 0x000000FF);
|
||||
tmp_u32[ 1] = ((tmp_u32[ 0] >> 8) & 0x00FF0000) | ((tmp_u32[ 0] >> 16) & 0x000000FF);
|
||||
tmp_u32[ 0] = ((tmp_u32[ 0] << 8) & 0x00FF0000) | ((tmp_u32[ 0] >> 0) & 0x000000FF);
|
||||
tmp_u32[9] = ((tmp_u32[4] >> 8) & 0x00FF0000) | ((tmp_u32[4] >> 16) & 0x000000FF);
|
||||
tmp_u32[8] = ((tmp_u32[4] << 8) & 0x00FF0000) | ((tmp_u32[4] >> 0) & 0x000000FF);
|
||||
tmp_u32[7] = ((tmp_u32[3] >> 8) & 0x00FF0000) | ((tmp_u32[3] >> 16) & 0x000000FF);
|
||||
tmp_u32[6] = ((tmp_u32[3] << 8) & 0x00FF0000) | ((tmp_u32[3] >> 0) & 0x000000FF);
|
||||
tmp_u32[5] = ((tmp_u32[2] >> 8) & 0x00FF0000) | ((tmp_u32[2] >> 16) & 0x000000FF);
|
||||
tmp_u32[4] = ((tmp_u32[2] << 8) & 0x00FF0000) | ((tmp_u32[2] >> 0) & 0x000000FF);
|
||||
tmp_u32[3] = ((tmp_u32[1] >> 8) & 0x00FF0000) | ((tmp_u32[1] >> 16) & 0x000000FF);
|
||||
tmp_u32[2] = ((tmp_u32[1] << 8) & 0x00FF0000) | ((tmp_u32[1] >> 0) & 0x000000FF);
|
||||
tmp_u32[1] = ((tmp_u32[0] >> 8) & 0x00FF0000) | ((tmp_u32[0] >> 16) & 0x000000FF);
|
||||
tmp_u32[0] = ((tmp_u32[0] << 8) & 0x00FF0000) | ((tmp_u32[0] >> 0) & 0x000000FF);
|
||||
|
||||
salt_len = salt_len * 2;
|
||||
}
|
||||
@ -2839,9 +2839,9 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
|
||||
if (memcmp (SIGNATURE_DPAPIMK, input_buf, 9)) return (PARSER_SIGNATURE_UNMATCHED);
|
||||
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
|
||||
salt_t *salt = hash_buf->salt;
|
||||
salt_t *salt = hash_buf->salt;
|
||||
|
||||
dpapimk_t *dpapimk = (dpapimk_t *) hash_buf->esalt;
|
||||
|
||||
@ -2873,7 +2873,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
dpapimk->context = atoll ((const char *) context_pos);
|
||||
|
||||
SID_pos = (u8 *) strchr ((const char *) context_pos, '*');
|
||||
|
||||
|
||||
if (SID_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
SID_pos++;
|
||||
@ -2922,7 +2922,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
|
||||
if (contents_len_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
if (contents_len_pos - iv_pos != iv_len) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
if (contents_len_pos - iv_pos != iv_len) return (PARSER_SALT_LENGTH);
|
||||
|
||||
if (is_valid_hex_string (iv_pos, 32) == false) return (PARSER_SALT_ENCODING);
|
||||
|
||||
@ -2935,7 +2935,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
dpapimk->iv[1] = byte_swap_32 (dpapimk->iv[1]);
|
||||
dpapimk->iv[2] = byte_swap_32 (dpapimk->iv[2]);
|
||||
dpapimk->iv[3] = byte_swap_32 (dpapimk->iv[3]);
|
||||
|
||||
|
||||
contents_len_pos++;
|
||||
|
||||
dpapimk->contents_len = (atoll ((const char *) contents_len_pos));
|
||||
@ -2962,13 +2962,8 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
dpapimk->contents[i] = byte_swap_32 (dpapimk->contents[i]);
|
||||
}
|
||||
|
||||
digest[0] = dpapimk->iv[0];
|
||||
digest[1] = dpapimk->iv[1];
|
||||
digest[2] = dpapimk->iv[2];
|
||||
digest[3] = dpapimk->iv[3];
|
||||
|
||||
u32 SID_len = cipher_algo_pos - 1 - SID_pos;
|
||||
|
||||
|
||||
/* maximum size of SID supported */
|
||||
u8 *SID_unicode = (u8 *) hcmalloc (32 * 4);
|
||||
memset (SID_unicode, 0, 32 * 4);
|
||||
@ -2977,12 +2972,12 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
{
|
||||
SID_unicode[i * 2] = SID_pos[i];
|
||||
}
|
||||
|
||||
|
||||
SID_unicode[(SID_len + 1) * 2] = 0x80;
|
||||
|
||||
/* Specific to DPAPI: needs trailing '\0' while computing hash */
|
||||
dpapimk->SID_len = (SID_len + 1) * 2;
|
||||
|
||||
|
||||
memcpy ((u8 *) dpapimk->SID, SID_unicode, 32 * 4);
|
||||
|
||||
for (u32 i = 0; i < 32; i++)
|
||||
@ -2990,8 +2985,20 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
dpapimk->SID[i] = byte_swap_32 (dpapimk->SID[i]);
|
||||
}
|
||||
|
||||
digest[0] = dpapimk->iv[0];
|
||||
digest[1] = dpapimk->iv[1];
|
||||
digest[2] = dpapimk->iv[2];
|
||||
digest[3] = dpapimk->iv[3];
|
||||
|
||||
salt->salt_buf[0] = dpapimk->iv[0];
|
||||
salt->salt_buf[1] = dpapimk->iv[1];
|
||||
salt->salt_buf[2] = dpapimk->iv[2];
|
||||
salt->salt_buf[3] = dpapimk->iv[3];
|
||||
|
||||
salt->salt_len = 32;
|
||||
|
||||
hcfree(SID_unicode);
|
||||
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
@ -12611,7 +12618,7 @@ int krb5tgs_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
|
||||
u8 *edata_ptr = (u8 *) krb5tgs->edata2;
|
||||
|
||||
krb5tgs->edata2_len = (data_len - 32) / 2 ;
|
||||
krb5tgs->edata2_len = (data_len - 32) / 2;
|
||||
|
||||
/* skip '$' */
|
||||
for (u32 i = 16 * 2 + 1; i < (krb5tgs->edata2_len * 2) + (16 * 2 + 1); i += 2)
|
||||
@ -18117,11 +18124,11 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
u32 algorithm = (u32) keepass->algorithm;
|
||||
u32 keyfile_len = (u32) keepass->keyfile_len;
|
||||
|
||||
u32 *ptr_final_random_seed = (u32 *) keepass->final_random_seed ;
|
||||
u32 *ptr_transf_random_seed = (u32 *) keepass->transf_random_seed ;
|
||||
u32 *ptr_enc_iv = (u32 *) keepass->enc_iv ;
|
||||
u32 *ptr_contents_hash = (u32 *) keepass->contents_hash ;
|
||||
u32 *ptr_keyfile = (u32 *) keepass->keyfile ;
|
||||
u32 *ptr_final_random_seed = (u32 *) keepass->final_random_seed;
|
||||
u32 *ptr_transf_random_seed = (u32 *) keepass->transf_random_seed;
|
||||
u32 *ptr_enc_iv = (u32 *) keepass->enc_iv;
|
||||
u32 *ptr_contents_hash = (u32 *) keepass->contents_hash;
|
||||
u32 *ptr_keyfile = (u32 *) keepass->keyfile;
|
||||
|
||||
/* specific to version 2 */
|
||||
u32 expected_bytes_len;
|
||||
@ -18206,7 +18213,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
else if (version == 2)
|
||||
{
|
||||
expected_bytes_len = 8;
|
||||
ptr_expected_bytes = (u32 *) keepass->expected_bytes ;
|
||||
ptr_expected_bytes = (u32 *) keepass->expected_bytes;
|
||||
|
||||
for (u32 i = 0; i < expected_bytes_len; i++, ptr_data += 8)
|
||||
sprintf (ptr_data, "%08x", ptr_expected_bytes[i]);
|
||||
@ -18499,7 +18506,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
{
|
||||
dpapimk_t *dpapimks = (dpapimk_t *) esalts_buf;
|
||||
|
||||
dpapimk_t *dpapimk = &dpapimks[digest_cur];
|
||||
dpapimk_t *dpapimk = &dpapimks[digest_cur];
|
||||
|
||||
u32 version = (u32) dpapimk->version;
|
||||
u32 context = (u32) dpapimk->context;
|
||||
@ -18510,13 +18517,13 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
char *ptr_SID = (char *) dpapimk->SID_tmp;
|
||||
char *ptr_cipher_algorithm = (char *) dpapimk->cipher_algo;
|
||||
char *ptr_hash_algorithm = (char *) dpapimk->hash_algo;
|
||||
u32 *ptr_iv = (u32 *) dpapimk->iv;
|
||||
u32 *ptr_contents = (u32 *) dpapimk->contents;
|
||||
u32 *ptr_iv = (u32 *) dpapimk->iv;
|
||||
u32 *ptr_contents = (u32 *) dpapimk->contents;
|
||||
|
||||
u32 u32_iv[4];
|
||||
u8 iv[32 + 1];
|
||||
|
||||
for (u32 i = 0 ; i < iv_len / 8; i++)
|
||||
for (u32 i = 0; i < iv_len / 8; i++)
|
||||
{
|
||||
u32_iv[i] = byte_swap_32 (ptr_iv[i]);
|
||||
u32_to_hex_lower (u32_iv[i], iv + i * 8);
|
||||
@ -18524,14 +18531,22 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
iv[32] = '\0';
|
||||
|
||||
u32 u32_contents[36];
|
||||
u8 contents[288 + 1];
|
||||
u8 contents[288 + 1];
|
||||
|
||||
for (u32 i = 0 ; i < contents_len / 8; i++)
|
||||
for (u32 i = 0; i < contents_len / 8; i++)
|
||||
{
|
||||
u32_contents[i] = byte_swap_32 (ptr_contents[i]);
|
||||
u32_to_hex_lower (u32_contents[i], contents + i * 8);
|
||||
}
|
||||
contents[288] = '\0';
|
||||
|
||||
if (version == 1)
|
||||
{
|
||||
contents[208] = '\0';
|
||||
}
|
||||
else
|
||||
{
|
||||
contents[288] = '\0';
|
||||
}
|
||||
|
||||
snprintf (out_buf, out_len - 1, "%s%d*%d*%s*%s*%s*%d*%s*%d*%s",
|
||||
SIGNATURE_DPAPIMK,
|
||||
@ -22847,7 +22862,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
|
||||
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE;
|
||||
hashconfig->kern_type = KERN_TYPE_DPAPIMK;
|
||||
hashconfig->dgst_size = DGST_SIZE_4_5;
|
||||
hashconfig->dgst_size = DGST_SIZE_4_4;
|
||||
hashconfig->parse_func = dpapimk_parse_hash;
|
||||
hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE;
|
||||
hashconfig->dgst_pos0 = 0;
|
||||
|
||||
@ -241,9 +241,7 @@ static const char *USAGE_BIG[] =
|
||||
" 1000 | NTLM | Operating Systems",
|
||||
" 1100 | Domain Cached Credentials (DCC), MS Cache | Operating Systems",
|
||||
" 2100 | Domain Cached Credentials 2 (DCC2), MS Cache 2 | Operating Systems",
|
||||
/* Fist0urs */
|
||||
" 15300 | DPAPI masterkey file v1 and v2 | Operating Systems",
|
||||
/* Fist0urs_end */
|
||||
" 12800 | MS-AzureSync PBKDF2-HMAC-SHA256 | Operating Systems",
|
||||
" 1500 | descrypt, DES (Unix), Traditional DES | Operating Systems",
|
||||
" 12400 | BSDiCrypt, Extended DES | Operating Systems",
|
||||
|
||||
436
tools/test.pl
436
tools/test.pl
@ -7,19 +7,19 @@
|
||||
|
||||
use strict;
|
||||
use warnings;
|
||||
use Digest::MD4 qw (md4 md4_hex);
|
||||
use Digest::MD5 qw (md5 md5_hex);
|
||||
use Digest::SHA qw (sha1 sha256 sha384 sha512 sha1_hex sha224_hex sha256_hex sha384_hex sha512_hex);
|
||||
use Digest::HMAC qw (hmac hmac_hex);
|
||||
use Digest::Keccak qw (keccak_256_hex);
|
||||
use Digest::BLAKE2 qw (blake2b_hex);
|
||||
use Crypt::MySQL qw (password41);
|
||||
use Digest::GOST qw (gost gost_hex);
|
||||
use Digest::HMAC_MD5 qw (hmac_md5);
|
||||
use Digest::CRC qw (crc32);
|
||||
use Digest::MD4 qw (md4 md4_hex);
|
||||
use Digest::MD5 qw (md5 md5_hex);
|
||||
use Digest::SHA qw (sha1 sha256 sha384 sha512 sha1_hex sha224_hex sha256_hex sha384_hex sha512_hex hmac_sha1 hmac_sha512);
|
||||
use Digest::HMAC qw (hmac hmac_hex);
|
||||
use Digest::Keccak qw (keccak_256_hex);
|
||||
use Digest::BLAKE2 qw (blake2b_hex);
|
||||
use Crypt::MySQL qw (password41);
|
||||
use Digest::GOST qw (gost gost_hex);
|
||||
use Digest::HMAC_MD5 qw (hmac_md5);
|
||||
use Digest::CRC qw (crc32);
|
||||
use Crypt::PBKDF2;
|
||||
use Crypt::DES;
|
||||
use Crypt::ECB qw (encrypt);
|
||||
use Crypt::ECB qw (encrypt);
|
||||
use Crypt::CBC;
|
||||
use Crypt::Eksblowfish::Bcrypt qw (bcrypt en_base64);
|
||||
use Crypt::Digest::RIPEMD160 qw (ripemd160_hex);
|
||||
@ -47,7 +47,7 @@ my $hashcat = "./hashcat";
|
||||
|
||||
my $MAX_LEN = 55;
|
||||
|
||||
my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 99999);
|
||||
my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 99999);
|
||||
|
||||
my %is_unicode = map { $_ => 1 } qw(30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800);
|
||||
my %less_fifteen = map { $_ => 1 } qw(500 1600 1800 2400 2410 3200 6300 7400 10500 10700);
|
||||
@ -2621,6 +2621,61 @@ sub verify
|
||||
|
||||
next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
|
||||
}
|
||||
elsif ($mode == 15300)
|
||||
{
|
||||
($hash_in, $word) = split ":", $line;
|
||||
|
||||
next unless defined $hash_in;
|
||||
next unless defined $word;
|
||||
|
||||
my @tmp_data = split ('\$', $hash_in);
|
||||
|
||||
my $signature = $tmp_data[1];
|
||||
|
||||
next unless ($signature eq 'DPAPImk');
|
||||
|
||||
my @data = split ('\*', $tmp_data[2]);
|
||||
|
||||
next unless (scalar @data == 9);
|
||||
|
||||
my $version = shift @data;
|
||||
|
||||
next unless ($version == 1 || $version == 2);
|
||||
|
||||
my $context = shift @data;
|
||||
|
||||
my $SID = shift @data;
|
||||
|
||||
my $cipher_algorithm = shift @data;
|
||||
|
||||
my $hash_algorithm = shift @data;
|
||||
|
||||
my $iteration = shift @data;
|
||||
|
||||
my $iv = shift @data;
|
||||
|
||||
my $cipher_len = shift @data;
|
||||
|
||||
my $cipher = shift @data;
|
||||
|
||||
next unless (length ($cipher) == $cipher_len);
|
||||
|
||||
if ($version == 1)
|
||||
{
|
||||
next unless ($cipher_len == 208);
|
||||
}
|
||||
elsif ($version == 2)
|
||||
{
|
||||
next unless ($cipher_len == 288);
|
||||
}
|
||||
|
||||
$salt = substr ($hash_in, length ('$DPAPImk$'));
|
||||
|
||||
$param = $iv;
|
||||
$param2 = $cipher;
|
||||
|
||||
next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
|
||||
}
|
||||
else
|
||||
{
|
||||
print "ERROR: hash mode is not supported\n";
|
||||
@ -2996,6 +3051,14 @@ sub verify
|
||||
|
||||
$hash_out = $hash_in;
|
||||
}
|
||||
elsif ($mode == 15300)
|
||||
{
|
||||
$hash_out = gen_hash ($mode, $word, $salt, $iter, $param, $param2);
|
||||
|
||||
$len = length $hash_out;
|
||||
|
||||
return unless (substr ($line, 0, $len) eq $hash_out);
|
||||
}
|
||||
else
|
||||
{
|
||||
$hash_out = gen_hash ($mode, $word, $salt, $iter);
|
||||
@ -3510,6 +3573,12 @@ sub passthrough
|
||||
{
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32));
|
||||
}
|
||||
elsif ($mode == 15300)
|
||||
{
|
||||
$salt_buf = get_random_dpapimk_salt ();
|
||||
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, $salt_buf);
|
||||
}
|
||||
else
|
||||
{
|
||||
print "ERROR: Unsupported hash type\n";
|
||||
@ -4450,6 +4519,20 @@ sub single
|
||||
}
|
||||
}
|
||||
}
|
||||
elsif ($mode == 15300)
|
||||
{
|
||||
for (my $i = 1; $i < 16; $i++)
|
||||
{
|
||||
if ($len != 0)
|
||||
{
|
||||
rnd ($mode, $len, 16);
|
||||
}
|
||||
else
|
||||
{
|
||||
rnd ($mode, $i, 16);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -8153,6 +8236,270 @@ END_CODE
|
||||
|
||||
$tmp_hash = sprintf ("\$blockchain\$v2\$%d\$%s\$%s", $iterations, length ($salt_buf . $encrypted) / 2, $salt_buf . $encrypted);
|
||||
}
|
||||
elsif ($mode == 15300)
|
||||
{
|
||||
my @salt_arr = split ('\*', $salt_buf);
|
||||
|
||||
my $version = $salt_arr[0];
|
||||
|
||||
my $context = $salt_arr[1];
|
||||
|
||||
my $SID = $salt_arr[2];
|
||||
|
||||
my $cipher_algorithm;
|
||||
|
||||
my $hash_algorithm;
|
||||
|
||||
my $iterations = $salt_arr[3];
|
||||
|
||||
my $salt = pack ("H*", $salt_arr[4]);
|
||||
|
||||
my $cipher_len;
|
||||
|
||||
if ($version == 1)
|
||||
{
|
||||
$cipher_algorithm = "des3";
|
||||
|
||||
$hash_algorithm = "sha1";
|
||||
|
||||
$cipher_len = 208;
|
||||
}
|
||||
else
|
||||
{
|
||||
$cipher_algorithm = "aes256";
|
||||
|
||||
$hash_algorithm = "sha512";
|
||||
|
||||
$cipher_len = 288;
|
||||
}
|
||||
|
||||
my $cipher;
|
||||
|
||||
# intermediate values
|
||||
|
||||
my $user_hash;
|
||||
my $user_derivationKey;
|
||||
my $encKey;
|
||||
my $expected_hmac;
|
||||
my $cleartext;
|
||||
|
||||
if ($context == 1)
|
||||
{
|
||||
$user_hash = sha1 (encode ("UTF-16LE", $word_buf));
|
||||
}
|
||||
elsif ($context == 2)
|
||||
{
|
||||
$user_hash = md4 (encode ("UTF-16LE", $word_buf));
|
||||
}
|
||||
|
||||
$user_derivationKey = hmac_sha1 (encode ("UTF-16LE", $SID . "\x00"), $user_hash);
|
||||
|
||||
my $hmacSalt = randbytes (16);
|
||||
my $last_key = randbytes (64);
|
||||
|
||||
if ($version == 1)
|
||||
{
|
||||
$encKey = hmac_sha1 ($hmacSalt, $user_derivationKey);
|
||||
$expected_hmac = hmac_sha1 ($last_key, $encKey);
|
||||
|
||||
# need padding because keyLen is 24 and hashLen 20
|
||||
$expected_hmac = $expected_hmac . randbytes(4);
|
||||
}
|
||||
elsif ($version == 2)
|
||||
{
|
||||
$encKey = hmac_sha512 ($hmacSalt, $user_derivationKey);
|
||||
$expected_hmac = hmac_sha512 ($last_key, $encKey);
|
||||
}
|
||||
|
||||
$cleartext = $hmacSalt . $expected_hmac . $last_key;
|
||||
|
||||
my $derived_key;
|
||||
my $key;
|
||||
my $iv;
|
||||
|
||||
my $pbkdf2;
|
||||
|
||||
if ($version == 1)
|
||||
{
|
||||
$derived_key = dpapi_pbkdf2 ($user_derivationKey, $salt, $iterations, 32, \&hmac_sha1);
|
||||
}
|
||||
elsif ($version == 2)
|
||||
{
|
||||
$derived_key = dpapi_pbkdf2 ($user_derivationKey, $salt, $iterations, 48, \&hmac_sha512);
|
||||
}
|
||||
|
||||
if (defined $additional_param)
|
||||
{
|
||||
$salt = pack ("H*", $additional_param);
|
||||
}
|
||||
|
||||
if (defined $additional_param2)
|
||||
{
|
||||
$cipher = $additional_param2;
|
||||
|
||||
if ($version == 1)
|
||||
{
|
||||
$key = substr ($derived_key, 0, 24);
|
||||
$iv = substr ($derived_key, 24, 8);
|
||||
|
||||
my $p1 = Crypt::ECB->new ({
|
||||
key => substr ($key, 0, 8),
|
||||
cipher => "DES",
|
||||
literal_key => 1,
|
||||
header => "none",
|
||||
keysize => 8,
|
||||
padding => "null",
|
||||
});
|
||||
|
||||
my $p2 = Crypt::ECB->new ({
|
||||
key => substr ($key, 8, 8),
|
||||
cipher => "DES",
|
||||
literal_key => 1,
|
||||
header => "none",
|
||||
keysize => 8,
|
||||
padding => "null",
|
||||
});
|
||||
|
||||
my $p3 = Crypt::ECB->new ({
|
||||
key => substr ($key, 16, 8),
|
||||
cipher => "DES",
|
||||
literal_key => 1,
|
||||
header => "none",
|
||||
keysize => 8,
|
||||
padding => "null",
|
||||
});
|
||||
|
||||
# let's compute a 3DES-EDE-CBC decryption
|
||||
$iv = substr ($cipher, 0, 8);
|
||||
|
||||
my $out1;
|
||||
my $out2;
|
||||
my $out3;
|
||||
my $expected_cleartext = "";
|
||||
|
||||
# size of cipherlen is 104 bytes
|
||||
for (my $k = 0; $k < 13; $k++)
|
||||
{
|
||||
$out1 = $p3->decrypt (substr ($cipher, $k * 8, 8));
|
||||
$out2 = $p2->encrypt ($out1);
|
||||
$out3 = $p1->decrypt ($out2);
|
||||
|
||||
$expected_cleartext .= substr ($out3, 0, 8) ^ $iv;
|
||||
|
||||
$iv = substr ($cipher, $k * 8, 8);
|
||||
}
|
||||
|
||||
if ($expected_cleartext != $cleartext)
|
||||
{
|
||||
$cleartext = "0" x 104;
|
||||
}
|
||||
}
|
||||
elsif ($version == 2)
|
||||
{
|
||||
$key = substr ($derived_key, 0, 32);
|
||||
$iv = substr ($derived_key, 32, 16);
|
||||
|
||||
my $aes = Crypt::CBC->new ({
|
||||
key => $key,
|
||||
cipher => "Crypt::Rijndael",
|
||||
iv => $iv,
|
||||
literal_key => 1,
|
||||
header => "none",
|
||||
keysize => 32,
|
||||
padding => "null",
|
||||
});
|
||||
|
||||
my $expected_cleartext = $aes->decrypt(pack ("H*", $cipher));
|
||||
|
||||
if ($expected_cleartext != $cleartext)
|
||||
{
|
||||
$cleartext = "0" x 144;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if ($version == 1)
|
||||
{
|
||||
$key = substr ($derived_key, 0, 24);
|
||||
$iv = substr ($derived_key, 24, 8);
|
||||
|
||||
my $p1 = Crypt::ECB->new ({
|
||||
key => substr ($key, 0, 8),
|
||||
cipher => "DES",
|
||||
literal_key => 1,
|
||||
header => "none",
|
||||
keysize => 8,
|
||||
padding => "null",
|
||||
});
|
||||
|
||||
my $p2 = Crypt::ECB->new ({
|
||||
key => substr ($key, 8, 8),
|
||||
cipher => "DES",
|
||||
literal_key => 1,
|
||||
header => "none",
|
||||
keysize => 8,
|
||||
padding => "null",
|
||||
});
|
||||
|
||||
my $p3 = Crypt::ECB->new ({
|
||||
key => substr ($key, 16, 8),
|
||||
cipher => "DES",
|
||||
literal_key => 1,
|
||||
header => "none",
|
||||
keysize => 8,
|
||||
padding => "null",
|
||||
});
|
||||
|
||||
# let's compute a 3DES-EDE-CBC encryption
|
||||
|
||||
# compute first block
|
||||
my $out1 = $p1->encrypt (substr ($cleartext, 0, 8) ^ $iv);
|
||||
my $out2 = $p2->decrypt ($out1);
|
||||
my $out3 = $p3->encrypt ($out2);
|
||||
|
||||
$cipher = substr ($out3, 0, 8);
|
||||
|
||||
# size of cipherlen is 104 bytes
|
||||
for (my $k = 1; $k < 13; $k++)
|
||||
{
|
||||
$iv = $out3;
|
||||
|
||||
$out1 = $p1->encrypt (substr ($cleartext, $k * 8, 8) ^ $iv);
|
||||
$out2 = $p2->decrypt ($out1);
|
||||
$out3 = $p3->encrypt ($out2);
|
||||
|
||||
$cipher .= substr ($out3, 0, 8);
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
$key = substr ($derived_key, 0, 32);
|
||||
$iv = substr ($derived_key, 32, 16);
|
||||
|
||||
my $aes = Crypt::CBC->new ({
|
||||
key => $key,
|
||||
cipher => "Crypt::Rijndael",
|
||||
iv => $iv,
|
||||
literal_key => 1,
|
||||
header => "none",
|
||||
keysize => 32,
|
||||
padding => "null",
|
||||
});
|
||||
|
||||
$cipher = $aes->encrypt ($cleartext);
|
||||
}
|
||||
|
||||
$tmp_hash = sprintf ('$DPAPImk$%d*%d*%s*%s*%s*%d*%s*%d*%s',
|
||||
$version,
|
||||
$context,
|
||||
$SID,
|
||||
$cipher_algorithm,
|
||||
$hash_algorithm,
|
||||
$iterations,
|
||||
unpack ("H*", $salt),
|
||||
$cipher_len,
|
||||
unpack ("H*", $cipher));
|
||||
}
|
||||
elsif ($mode == 99999)
|
||||
{
|
||||
$tmp_hash = sprintf ("%s", $word_buf);
|
||||
@ -8161,6 +8508,30 @@ END_CODE
|
||||
return ($tmp_hash);
|
||||
}
|
||||
|
||||
#Thanks to Jochen Hoenicke <hoenicke@gmail.com>
|
||||
# (one of the authors of Palm Keyring)
|
||||
# for these next two subs.
|
||||
sub dpapi_pbkdf2
|
||||
{
|
||||
my ($password, $salt, $iter, $keylen, $prf) = @_;
|
||||
my ($k, $t, $u, $ui, $i);
|
||||
$t = "";
|
||||
for ($k = 1; length($t) < $keylen; $k++)
|
||||
{
|
||||
$u = $ui = &$prf($salt.pack('N', $k), $password);
|
||||
for ($i = 1; $i < $iter; $i++)
|
||||
{
|
||||
# modification to fit Microsoft
|
||||
# weird pbkdf2 implementation...
|
||||
$ui = &$prf($u, $password);
|
||||
$u ^= $ui;
|
||||
}
|
||||
$t .= $u;
|
||||
}
|
||||
return substr($t, 0, $keylen);
|
||||
}
|
||||
|
||||
|
||||
sub rnd
|
||||
{
|
||||
my $mode = shift;
|
||||
@ -8267,6 +8638,10 @@ sub rnd
|
||||
{
|
||||
$salt_buf = get_pstoken_salt ();
|
||||
}
|
||||
elsif ($mode == 15300)
|
||||
{
|
||||
$salt_buf = get_random_dpapimk_salt ();
|
||||
}
|
||||
else
|
||||
{
|
||||
my @salt_arr;
|
||||
@ -9707,6 +10082,43 @@ sub get_random_dnssec_salt
|
||||
return $salt_buf;
|
||||
}
|
||||
|
||||
sub get_random_dpapimk_salt
|
||||
{
|
||||
my $salt_buf = "";
|
||||
|
||||
my $version = get_random_num (1, 3);
|
||||
|
||||
my $context = get_random_num (1, 3);
|
||||
|
||||
my $iterations;
|
||||
|
||||
my $SID = sprintf ('S-15-21-%d-%d-%d-%d',
|
||||
get_random_num(400000000,490000000),
|
||||
get_random_num(400000000,490000000),
|
||||
get_random_num(400000000,490000000),
|
||||
get_random_num(1000,1999));
|
||||
|
||||
if ($version == 1)
|
||||
{
|
||||
$iterations = get_random_num (4000, 24000);
|
||||
}
|
||||
elsif ($version == 2)
|
||||
{
|
||||
$iterations = get_random_num (8000, 17000);
|
||||
}
|
||||
|
||||
my $iv = randbytes (16);
|
||||
$iv = unpack ("H*", $iv);
|
||||
|
||||
$salt_buf = $version . '*' .
|
||||
$context . '*' .
|
||||
$SID . '*' .
|
||||
$iterations . '*' .
|
||||
$iv . '*';
|
||||
|
||||
return $salt_buf;
|
||||
}
|
||||
|
||||
sub md5bit
|
||||
{
|
||||
my $digest = shift;
|
||||
|
||||
@ -9,7 +9,7 @@ TDIR="$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )"
|
||||
|
||||
# missing hash types: 5200,6251,6261,6271,6281
|
||||
|
||||
HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 125 130 131 132 133 140 141 150 160 200 300 400 500 600 900 1000 1100 1300 1400 1410 1411 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12001 12100 12200 12300 12400 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 99999"
|
||||
HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 125 130 131 132 133 140 141 150 160 200 300 400 500 600 900 1000 1100 1300 1400 1410 1411 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12001 12100 12200 12300 12400 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 99999"
|
||||
|
||||
#ATTACK_MODES="0 1 3 6 7"
|
||||
ATTACK_MODES="0 1 3 7"
|
||||
@ -22,7 +22,7 @@ HASHFILE_ONLY="2500"
|
||||
|
||||
NEVER_CRACK="11600 14900"
|
||||
|
||||
SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12001 12100 12200 12300 12400 12500 12700 12800 12900 13000 13200 13400 13600 14600 14700 14800 15100 15200"
|
||||
SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12001 12100 12200 12300 12400 12500 12700 12800 12900 13000 13200 13400 13600 14600 14700 14800 15100 15200 15300"
|
||||
|
||||
OPTS="--quiet --force --potfile-disable --runtime 400 --gpu-temp-disable --weak-hash-threshold=0"
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user