mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-22 16:18:09 +00:00
All remarks treated:
1) done + got rid of all u8 datatypes in shared struct 2) cf. previous 3) necessary as this is computed in _init then used in _comp 4) done 5) done 6) done => switch to 16 7) done
This commit is contained in:
Fist0urs
parent
80927f6f66
commit
a78dce94db
@ -1038,12 +1038,15 @@ typedef struct
|
||||
u32 version;
|
||||
u32 context;
|
||||
|
||||
u8 SID_tmp[64];
|
||||
u32 SID[32];
|
||||
u32 SID_len;
|
||||
u32 SID_offset;
|
||||
char cipher_algo[16];
|
||||
char hash_algo[16];
|
||||
|
||||
/* here only for possible
|
||||
forward compatibiliy
|
||||
*/
|
||||
//char cipher_algo[16];
|
||||
//char hash_algo[16];
|
||||
|
||||
u32 iv[4];
|
||||
u32 contents_len;
|
||||
|
||||
233
OpenCL/m15300.cl
233
OpenCL/m15300.cl
@ -15,6 +15,9 @@
|
||||
#include "inc_rp.cl"
|
||||
#include "inc_cipher_aes.cl"
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
|
||||
#define PERM_OP(a,b,tt,n,m) \
|
||||
{ \
|
||||
tt = a >> n; \
|
||||
@ -357,52 +360,31 @@ __constant u32a c_skb[8][64] =
|
||||
}
|
||||
};
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX(i,n,S) (S)[(n)][(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
|
||||
#endif
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX1(i,S) (S)[(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX1(i,S) (u32x) ((S)[(i).s0], (S)[(i).s1], (S)[(i).s2], (S)[(i).s3], (S)[(i).s4], (S)[(i).s5], (S)[(i).s6], (S)[(i).s7], (S)[(i).s8], (S)[(i).s9], (S)[(i).sa], (S)[(i).sb], (S)[(i).sc], (S)[(i).sd], (S)[(i).se], (S)[(i).sf])
|
||||
#endif
|
||||
|
||||
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x tt;
|
||||
u32 tt;
|
||||
|
||||
u32x r = data[0];
|
||||
u32x l = data[1];
|
||||
u32 r = data[0];
|
||||
u32 l = data[1];
|
||||
|
||||
IP (r, l, tt);
|
||||
|
||||
r = rotl32 (r, 3u);
|
||||
l = rotl32 (l, 3u);
|
||||
r = rotl32_S (r, 3u);
|
||||
l = rotl32_S (l, 3u);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
u32 u;
|
||||
u32 t;
|
||||
|
||||
u = Kc[i + 0] ^ r;
|
||||
t = Kd[i + 0] ^ rotl32 (r, 28u);
|
||||
t = Kd[i + 0] ^ rotl32_S (r, 28u);
|
||||
|
||||
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
@ -414,7 +396,7 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i + 1] ^ l;
|
||||
t = Kd[i + 1] ^ rotl32 (l, 28u);
|
||||
t = Kd[i + 1] ^ rotl32_S (l, 28u);
|
||||
|
||||
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
@ -426,8 +408,8 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
l = rotl32 (l, 29u);
|
||||
r = rotl32 (r, 29u);
|
||||
l = rotl32_S (l, 29u);
|
||||
r = rotl32_S (r, 29u);
|
||||
|
||||
FP (r, l, tt);
|
||||
|
||||
@ -435,28 +417,28 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
|
||||
iv[1] = r;
|
||||
}
|
||||
|
||||
void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
void _des_crypt_decrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x tt;
|
||||
u32 tt;
|
||||
|
||||
u32x r = data[0];
|
||||
u32x l = data[1];
|
||||
u32 r = data[0];
|
||||
u32 l = data[1];
|
||||
|
||||
IP (r, l, tt);
|
||||
|
||||
r = rotl32 (r, 3u);
|
||||
l = rotl32 (l, 3u);
|
||||
r = rotl32_S (r, 3u);
|
||||
l = rotl32_S (l, 3u);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 16; i > 0; i -= 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
u32 u;
|
||||
u32 t;
|
||||
|
||||
u = Kc[i - 1] ^ r;
|
||||
t = Kd[i - 1] ^ rotl32 (r, 28u);
|
||||
t = Kd[i - 1] ^ rotl32_S (r, 28u);
|
||||
|
||||
l ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
@ -468,7 +450,7 @@ void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i - 2] ^ l;
|
||||
t = Kd[i - 2] ^ rotl32 (l, 28u);
|
||||
t = Kd[i - 2] ^ rotl32_S (l, 28u);
|
||||
|
||||
r ^= BOX (((u >> 2) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
|
||||
@ -480,8 +462,8 @@ void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
|
||||
| BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
l = rotl32 (l, 29u);
|
||||
r = rotl32 (r, 29u);
|
||||
l = rotl32_S (l, 29u);
|
||||
r = rotl32_S (r, 29u);
|
||||
|
||||
FP (r, l, tt);
|
||||
|
||||
@ -489,9 +471,9 @@ void _des_crypt_decrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
|
||||
iv[1] = r;
|
||||
}
|
||||
|
||||
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
|
||||
void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64])
|
||||
{
|
||||
u32x tt;
|
||||
u32 tt;
|
||||
|
||||
PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
|
||||
HPERM_OP (c, tt, 2, 0xcccc0000);
|
||||
@ -526,13 +508,13 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
c = c & 0x0fffffff;
|
||||
d = d & 0x0fffffff;
|
||||
|
||||
const u32x c00 = (c >> 0) & 0x0000003f;
|
||||
const u32x c06 = (c >> 6) & 0x00383003;
|
||||
const u32x c07 = (c >> 7) & 0x0000003c;
|
||||
const u32x c13 = (c >> 13) & 0x0000060f;
|
||||
const u32x c20 = (c >> 20) & 0x00000001;
|
||||
const u32 c00 = (c >> 0) & 0x0000003f;
|
||||
const u32 c06 = (c >> 6) & 0x00383003;
|
||||
const u32 c07 = (c >> 7) & 0x0000003c;
|
||||
const u32 c13 = (c >> 13) & 0x0000060f;
|
||||
const u32 c20 = (c >> 20) & 0x00000001;
|
||||
|
||||
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
| BOX (((c06 >> 0) & 0xff)
|
||||
|((c07 >> 0) & 0xff), 1, s_skb)
|
||||
| BOX (((c13 >> 0) & 0xff)
|
||||
@ -541,12 +523,12 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
|((c13 >> 8) & 0xff)
|
||||
|((c06 >> 16) & 0xff), 3, s_skb);
|
||||
|
||||
const u32x d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32x d07 = (d >> 7) & 0x00003f03;
|
||||
const u32x d21 = (d >> 21) & 0x0000000f;
|
||||
const u32x d22 = (d >> 22) & 0x00000030;
|
||||
const u32 d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32 d07 = (d >> 7) & 0x00003f03;
|
||||
const u32 d21 = (d >> 21) & 0x0000000f;
|
||||
const u32 d22 = (d >> 22) & 0x00000030;
|
||||
|
||||
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
| BOX (((d07 >> 0) & 0xff)
|
||||
|((d00 >> 8) & 0xff), 5, s_skb)
|
||||
| BOX (((d07 >> 8) & 0xff), 6, s_skb)
|
||||
@ -556,8 +538,8 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
Kc[i] = ((t << 16) | (s & 0x0000ffff));
|
||||
Kd[i] = ((s >> 16) | (t & 0xffff0000));
|
||||
|
||||
Kc[i] = rotl32 (Kc[i], 2u);
|
||||
Kd[i] = rotl32 (Kd[i], 2u);
|
||||
Kc[i] = rotl32_S (Kc[i], 2u);
|
||||
Kd[i] = rotl32_S (Kd[i], 2u);
|
||||
}
|
||||
}
|
||||
|
||||
@ -1107,6 +1089,59 @@ void sha1_transform_V (const u32x w0[4], const u32x w1[4], const u32x w2[4], con
|
||||
digest[4] += E;
|
||||
}
|
||||
|
||||
void hmac_sha1_pad_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5])
|
||||
{
|
||||
w0[0] = w0[0] ^ 0x36363636;
|
||||
w0[1] = w0[1] ^ 0x36363636;
|
||||
w0[2] = w0[2] ^ 0x36363636;
|
||||
w0[3] = w0[3] ^ 0x36363636;
|
||||
w1[0] = w1[0] ^ 0x36363636;
|
||||
w1[1] = w1[1] ^ 0x36363636;
|
||||
w1[2] = w1[2] ^ 0x36363636;
|
||||
w1[3] = w1[3] ^ 0x36363636;
|
||||
w2[0] = w2[0] ^ 0x36363636;
|
||||
w2[1] = w2[1] ^ 0x36363636;
|
||||
w2[2] = w2[2] ^ 0x36363636;
|
||||
w2[3] = w2[3] ^ 0x36363636;
|
||||
w3[0] = w3[0] ^ 0x36363636;
|
||||
w3[1] = w3[1] ^ 0x36363636;
|
||||
w3[2] = w3[2] ^ 0x36363636;
|
||||
w3[3] = w3[3] ^ 0x36363636;
|
||||
|
||||
ipad[0] = SHA1M_A;
|
||||
ipad[1] = SHA1M_B;
|
||||
ipad[2] = SHA1M_C;
|
||||
ipad[3] = SHA1M_D;
|
||||
ipad[4] = SHA1M_E;
|
||||
|
||||
sha1_transform_V (w0, w1, w2, w3, ipad);
|
||||
|
||||
w0[0] = w0[0] ^ 0x6a6a6a6a;
|
||||
w0[1] = w0[1] ^ 0x6a6a6a6a;
|
||||
w0[2] = w0[2] ^ 0x6a6a6a6a;
|
||||
w0[3] = w0[3] ^ 0x6a6a6a6a;
|
||||
w1[0] = w1[0] ^ 0x6a6a6a6a;
|
||||
w1[1] = w1[1] ^ 0x6a6a6a6a;
|
||||
w1[2] = w1[2] ^ 0x6a6a6a6a;
|
||||
w1[3] = w1[3] ^ 0x6a6a6a6a;
|
||||
w2[0] = w2[0] ^ 0x6a6a6a6a;
|
||||
w2[1] = w2[1] ^ 0x6a6a6a6a;
|
||||
w2[2] = w2[2] ^ 0x6a6a6a6a;
|
||||
w2[3] = w2[3] ^ 0x6a6a6a6a;
|
||||
w3[0] = w3[0] ^ 0x6a6a6a6a;
|
||||
w3[1] = w3[1] ^ 0x6a6a6a6a;
|
||||
w3[2] = w3[2] ^ 0x6a6a6a6a;
|
||||
w3[3] = w3[3] ^ 0x6a6a6a6a;
|
||||
|
||||
opad[0] = SHA1M_A;
|
||||
opad[1] = SHA1M_B;
|
||||
opad[2] = SHA1M_C;
|
||||
opad[3] = SHA1M_D;
|
||||
opad[4] = SHA1M_E;
|
||||
|
||||
sha1_transform_V (w0, w1, w2, w3, opad);
|
||||
}
|
||||
|
||||
void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5], u32x digest[5])
|
||||
{
|
||||
digest[0] = ipad[0];
|
||||
@ -2136,7 +2171,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
key[4] = swap32_S (tmps[gid].out[4]);
|
||||
key[5] = swap32_S (tmps[gid].out[5]);
|
||||
|
||||
u32x iv[2];
|
||||
u32 iv[2];
|
||||
|
||||
iv[0] = swap32_S (tmps[gid].out[6]);
|
||||
iv[1] = swap32_S (tmps[gid].out[7]);
|
||||
@ -2145,27 +2180,27 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
/* Construct 3DES keys */
|
||||
|
||||
const u32x a = (key[0]);
|
||||
const u32x b = (key[1]);
|
||||
const u32 a = (key[0]);
|
||||
const u32 b = (key[1]);
|
||||
|
||||
u32x Ka[16];
|
||||
u32x Kb[16];
|
||||
u32 Ka[16];
|
||||
u32 Kb[16];
|
||||
|
||||
_des_crypt_keysetup (a, b, Ka, Kb, s_skb);
|
||||
|
||||
const u32x c = (key[2]);
|
||||
const u32x d = (key[3]);
|
||||
const u32 c = (key[2]);
|
||||
const u32 d = (key[3]);
|
||||
|
||||
u32x Kc[16];
|
||||
u32x Kd[16];
|
||||
u32 Kc[16];
|
||||
u32 Kd[16];
|
||||
|
||||
_des_crypt_keysetup (c, d, Kc, Kd, s_skb);
|
||||
|
||||
const u32x e = (key[4]);
|
||||
const u32x f = (key[5]);
|
||||
const u32 e = (key[4]);
|
||||
const u32 f = (key[5]);
|
||||
|
||||
u32x Ke[16];
|
||||
u32x Kf[16];
|
||||
u32 Ke[16];
|
||||
u32 Kf[16];
|
||||
|
||||
_des_crypt_keysetup (e, f, Ke, Kf, s_skb);
|
||||
|
||||
@ -2177,24 +2212,24 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
{
|
||||
/* First Pass */
|
||||
|
||||
u32x data[2];
|
||||
u32 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]);
|
||||
data[0] = swap32_S (esalt_bufs[digests_offset].contents[contents_off + 0]);
|
||||
data[1] = swap32_S (esalt_bufs[digests_offset].contents[contents_off + 1]);
|
||||
|
||||
u32x p1[2];
|
||||
u32 p1[2];
|
||||
|
||||
_des_crypt_decrypt (p1, data, Ke, Kf, s_SPtrans);
|
||||
|
||||
/* Second Pass */
|
||||
|
||||
u32x p2[2];
|
||||
u32 p2[2];
|
||||
|
||||
_des_crypt_encrypt (p2, p1, Kc, Kd, s_SPtrans);
|
||||
|
||||
/* Third Pass */
|
||||
|
||||
u32x out[2];
|
||||
u32 out[2];
|
||||
|
||||
_des_crypt_decrypt (out, p2, Ka, Kb, s_SPtrans);
|
||||
|
||||
@ -2403,21 +2438,21 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
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]);
|
||||
key[3] = l32_from_64 (tmps[gid].out64[1]);
|
||||
key[4] = h32_from_64 (tmps[gid].out64[2]);
|
||||
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]);
|
||||
key[0] = h32_from_64_S (tmps[gid].out64[0]);
|
||||
key[1] = l32_from_64_S (tmps[gid].out64[0]);
|
||||
key[2] = h32_from_64_S (tmps[gid].out64[1]);
|
||||
key[3] = l32_from_64_S (tmps[gid].out64[1]);
|
||||
key[4] = h32_from_64_S (tmps[gid].out64[2]);
|
||||
key[5] = l32_from_64_S (tmps[gid].out64[2]);
|
||||
key[6] = h32_from_64_S (tmps[gid].out64[3]);
|
||||
key[7] = l32_from_64_S (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_S (tmps[gid].out64[4]);
|
||||
iv[1] = l32_from_64_S (tmps[gid].out64[4]);
|
||||
iv[2] = h32_from_64_S (tmps[gid].out64[5]);
|
||||
iv[3] = l32_from_64_S (tmps[gid].out64[5]);
|
||||
|
||||
#define KEYLEN 60
|
||||
|
||||
@ -2584,14 +2619,14 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
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]))
|
||||
if ( expectedHmac[ 0] == h32_from_64_S (dgst64[0])
|
||||
&& expectedHmac[ 1] == l32_from_64_S (dgst64[0])
|
||||
&& expectedHmac[ 2] == h32_from_64_S (dgst64[1])
|
||||
&& expectedHmac[ 3] == l32_from_64_S (dgst64[1])
|
||||
&& expectedHmac[12] == h32_from_64_S (dgst64[6])
|
||||
&& expectedHmac[13] == l32_from_64_S (dgst64[6])
|
||||
&& expectedHmac[14] == h32_from_64_S (dgst64[7])
|
||||
&& expectedHmac[15] == l32_from_64_S (dgst64[7]))
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
|
||||
}
|
||||
|
||||
@ -49,28 +49,28 @@ static const char LM_MASKED_PLAIN[] = "[notfound]";
|
||||
#define LUKS_ALIGN_KEYSLOTS 4096
|
||||
|
||||
struct luks_phdr {
|
||||
char magic[LUKS_MAGIC_L];
|
||||
uint16_t version;
|
||||
char cipherName[LUKS_CIPHERNAME_L];
|
||||
char cipherMode[LUKS_CIPHERMODE_L];
|
||||
char hashSpec[LUKS_HASHSPEC_L];
|
||||
uint32_t payloadOffset;
|
||||
uint32_t keyBytes;
|
||||
char mkDigest[LUKS_DIGESTSIZE];
|
||||
char mkDigestSalt[LUKS_SALTSIZE];
|
||||
uint32_t mkDigestIterations;
|
||||
char uuid[UUID_STRING_L];
|
||||
struct {
|
||||
uint32_t active;
|
||||
/* parameters used for password processing */
|
||||
uint32_t passwordIterations;
|
||||
char passwordSalt[LUKS_SALTSIZE];
|
||||
/* parameters used for AF store/load */
|
||||
uint32_t keyMaterialOffset;
|
||||
uint32_t stripes;
|
||||
} keyblock[LUKS_NUMKEYS];
|
||||
/* Align it to 512 sector size */
|
||||
char _padding[432];
|
||||
char magic[LUKS_MAGIC_L];
|
||||
uint16_t version;
|
||||
char cipherName[LUKS_CIPHERNAME_L];
|
||||
char cipherMode[LUKS_CIPHERMODE_L];
|
||||
char hashSpec[LUKS_HASHSPEC_L];
|
||||
uint32_t payloadOffset;
|
||||
uint32_t keyBytes;
|
||||
char mkDigest[LUKS_DIGESTSIZE];
|
||||
char mkDigestSalt[LUKS_SALTSIZE];
|
||||
uint32_t mkDigestIterations;
|
||||
char uuid[UUID_STRING_L];
|
||||
struct {
|
||||
uint32_t active;
|
||||
/* parameters used for password processing */
|
||||
uint32_t passwordIterations;
|
||||
char passwordSalt[LUKS_SALTSIZE];
|
||||
/* parameters used for AF store/load */
|
||||
uint32_t keyMaterialOffset;
|
||||
uint32_t stripes;
|
||||
} keyblock[LUKS_NUMKEYS];
|
||||
/* Align it to 512 sector size */
|
||||
char _padding[432];
|
||||
};
|
||||
|
||||
// not from original headers start with hc_
|
||||
@ -425,12 +425,15 @@ typedef struct dpapimk
|
||||
u32 version;
|
||||
u32 context;
|
||||
|
||||
u8 SID_tmp[64];
|
||||
u32 SID[32];
|
||||
u32 SID_len;
|
||||
u32 SID_offset;
|
||||
u8 cipher_algo[16];
|
||||
u8 hash_algo[16];
|
||||
|
||||
/* here only for possible
|
||||
forward compatibiliy
|
||||
*/
|
||||
// u8 cipher_algo[16];
|
||||
// u8 hash_algo[16];
|
||||
|
||||
u32 iv[4];
|
||||
u32 contents_len;
|
||||
@ -1226,7 +1229,7 @@ typedef enum display_len
|
||||
DISPLAY_LEN_MAX_15100 = 6 + 6 + 1 + 8 + 1 + 28,
|
||||
DISPLAY_LEN_MIN_15200 = 1 + 10 + 1 + 2 + 1 + 1 + 1 + 1 + 1 + 64,
|
||||
DISPLAY_LEN_MAX_15200 = 1 + 10 + 1 + 2 + 1 + 8 + 1 + 5 + 1 + 20000,
|
||||
DISPLAY_LEN_MIN_15300 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 10 + 1 + 4 + 1 + 4 + 1 + 1 + 1 + 32 + 1 + 3 + 128,
|
||||
DISPLAY_LEN_MIN_15300 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 10 + 1 + 4 + 1 + 4 + 1 + 1 + 1 + 32 + 1 + 3 + 1 + 128,
|
||||
DISPLAY_LEN_MAX_15300 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 100 + 1 + 6 + 1 + 6 + 1 + 10 + 1 + 32 + 1 + 4 + 1 + 512,
|
||||
DISPLAY_LEN_MIN_99999 = 1,
|
||||
DISPLAY_LEN_MAX_99999 = 55,
|
||||
|
||||
@ -6,7 +6,7 @@
|
||||
#include "common.h"
|
||||
#include "benchmark.h"
|
||||
|
||||
const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_CNT = 150;
|
||||
const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_CNT = 151;
|
||||
|
||||
const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_BUF[] =
|
||||
{
|
||||
|
||||
130
src/interface.c
130
src/interface.c
@ -2862,16 +2862,12 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
|
||||
version_pos = input_buf + 8 + 1;
|
||||
|
||||
dpapimk->version = atoll ((const char *) version_pos);
|
||||
|
||||
context_pos = (u8 *) strchr ((const char *) version_pos, '*');
|
||||
|
||||
if (context_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
context_pos++;
|
||||
|
||||
dpapimk->context = atoll ((const char *) context_pos);
|
||||
|
||||
SID_pos = (u8 *) strchr ((const char *) context_pos, '*');
|
||||
|
||||
if (SID_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
@ -2882,36 +2878,20 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
|
||||
if (cipher_algo_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
for (int i = 0; i < cipher_algo_pos - SID_pos; i++)
|
||||
dpapimk->SID_tmp[i] = SID_pos[i];
|
||||
dpapimk->SID_tmp[cipher_algo_pos - SID_pos] = '\0';
|
||||
|
||||
cipher_algo_pos++;
|
||||
|
||||
hash_algo_pos = (u8 *) strchr ((const char *) cipher_algo_pos, '*');
|
||||
|
||||
if (hash_algo_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
for (int i = 0; i < hash_algo_pos - cipher_algo_pos; i++)
|
||||
{
|
||||
dpapimk->cipher_algo[i] = cipher_algo_pos[i];
|
||||
}
|
||||
dpapimk->cipher_algo[hash_algo_pos - cipher_algo_pos] = '\0';
|
||||
|
||||
hash_algo_pos++;
|
||||
|
||||
rounds_pos = (u8 *) strchr ((const char *) hash_algo_pos, '*');
|
||||
|
||||
if (rounds_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
for (int i = 0; i < rounds_pos - hash_algo_pos; i++)
|
||||
dpapimk->hash_algo[i] = hash_algo_pos[i];
|
||||
dpapimk->hash_algo[rounds_pos - hash_algo_pos] = '\0';
|
||||
|
||||
rounds_pos++;
|
||||
|
||||
salt->salt_iter = (atoll ((const char *) rounds_pos)) - 1;
|
||||
|
||||
iv_pos = (u8 *) strchr ((const char *) rounds_pos, '*');
|
||||
|
||||
if (iv_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
@ -2926,6 +2906,32 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
|
||||
if (is_valid_hex_string (iv_pos, 32) == false) return (PARSER_SALT_ENCODING);
|
||||
|
||||
contents_len_pos++;
|
||||
|
||||
contents_pos = (u8 *) strchr ((const char *) contents_len_pos, '*');
|
||||
|
||||
if (contents_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
contents_pos++;
|
||||
|
||||
u32 version = atoll ((const char *) version_pos);
|
||||
u32 contents_len = atoll ((const char *) contents_len_pos);
|
||||
|
||||
if (version == 1 && contents_len != 208) return (PARSER_SALT_LENGTH);
|
||||
if (version == 2 && contents_len != 288) return (PARSER_SALT_LENGTH);
|
||||
|
||||
if (is_valid_hex_string (contents_pos, contents_len) == false) return (PARSER_SALT_ENCODING);
|
||||
|
||||
u8 *end_line = (u8 *) strchr ((const char *) contents_pos, 0);
|
||||
|
||||
if (end_line - contents_pos != contents_len) return (PARSER_SALT_LENGTH);
|
||||
|
||||
dpapimk->version = version;
|
||||
|
||||
dpapimk->context = atoll ((const char *) context_pos);
|
||||
|
||||
salt->salt_iter = (atoll ((const char *) rounds_pos)) - 1;
|
||||
|
||||
dpapimk->iv[0] = hex_to_u32 ((const u8 *) &iv_pos[ 0]);
|
||||
dpapimk->iv[1] = hex_to_u32 ((const u8 *) &iv_pos[ 8]);
|
||||
dpapimk->iv[2] = hex_to_u32 ((const u8 *) &iv_pos[16]);
|
||||
@ -2936,24 +2942,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
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));
|
||||
|
||||
contents_pos = (u8 *) strchr ((const char *) contents_len_pos, '*');
|
||||
|
||||
if (contents_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
contents_pos++;
|
||||
|
||||
if (dpapimk->version == 1 && dpapimk->contents_len != 208) return (PARSER_SALT_LENGTH);
|
||||
if (dpapimk->version == 2 && dpapimk->contents_len != 288) return (PARSER_SALT_LENGTH);
|
||||
|
||||
if (is_valid_hex_string (contents_pos, dpapimk->contents_len) == false) return (PARSER_SALT_ENCODING);
|
||||
|
||||
u8 *end_line = (u8 *) strchr ((const char *) contents_pos, '\0');
|
||||
|
||||
if (end_line - contents_pos != dpapimk->contents_len) return (PARSER_SALT_LENGTH);
|
||||
dpapimk->contents_len = contents_len;
|
||||
|
||||
for (u32 i = 0; i < dpapimk->contents_len / 4; i++)
|
||||
{
|
||||
@ -2995,7 +2984,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
salt->salt_buf[2] = dpapimk->iv[2];
|
||||
salt->salt_buf[3] = dpapimk->iv[3];
|
||||
|
||||
salt->salt_len = 32;
|
||||
salt->salt_len = 16;
|
||||
|
||||
hcfree(SID_unicode);
|
||||
|
||||
@ -18511,24 +18500,52 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
u32 version = (u32) dpapimk->version;
|
||||
u32 context = (u32) dpapimk->context;
|
||||
u32 rounds = salt.salt_iter + 1;
|
||||
u32 iv_len = 32;
|
||||
u32 contents_len = (u32) dpapimk->contents_len;
|
||||
u32 SID_len = (u32) dpapimk->SID_len;
|
||||
u32 iv_len = 32;
|
||||
|
||||
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;
|
||||
u8 cipher_algorithm[8] = { 0 };
|
||||
u8 hash_algorithm[8] = { 0 };
|
||||
u8 SID[512] = { 0 };
|
||||
u8* SID_tmp;
|
||||
|
||||
u32 *ptr_SID = (u32 *) dpapimk->SID;
|
||||
u32 *ptr_iv = (u32 *) dpapimk->iv;
|
||||
u32 *ptr_contents = (u32 *) dpapimk->contents;
|
||||
|
||||
u32 u32_iv[4];
|
||||
u8 iv[32 + 1];
|
||||
|
||||
/* convert back SID */
|
||||
|
||||
SID_tmp = (u8 *) hcmalloc ((SID_len + 1) * sizeof(u8));
|
||||
|
||||
for (u32 i = 0; i < (SID_len / 4) + 1; i++)
|
||||
{
|
||||
u8 hex[8] = { 0 };
|
||||
u32_to_hex_lower (byte_swap_32 (ptr_SID[i]), hex);
|
||||
|
||||
for (u32 j = 0, k = 0; j < 8; j += 2, k++)
|
||||
{
|
||||
SID_tmp[i * 4 + k] = hex_to_u8 (&hex[j]);
|
||||
}
|
||||
}
|
||||
/* overwrite trailing 0x80 */
|
||||
SID_tmp[SID_len] = 0;
|
||||
|
||||
for (u32 i = 0, j = 0 ; j < SID_len ; i++, j += 2)
|
||||
{
|
||||
SID[i] = SID_tmp[j];
|
||||
}
|
||||
|
||||
hcfree(SID_tmp);
|
||||
|
||||
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);
|
||||
}
|
||||
iv[32] = '\0';
|
||||
iv[32] = 0;
|
||||
|
||||
u32 u32_contents[36];
|
||||
u8 contents[288 + 1];
|
||||
@ -18541,20 +18558,31 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
|
||||
if (version == 1)
|
||||
{
|
||||
contents[208] = '\0';
|
||||
contents[208] = 0;
|
||||
}
|
||||
else
|
||||
{
|
||||
contents[288] = '\0';
|
||||
contents[288] = 0;
|
||||
}
|
||||
|
||||
if (contents_len == 288 && version == 2)
|
||||
{
|
||||
memcpy(cipher_algorithm, "aes256", strlen("aes256"));
|
||||
memcpy(hash_algorithm, "sha512", strlen("sha512"));
|
||||
}
|
||||
else if (contents_len == 208 && version == 1)
|
||||
{
|
||||
memcpy(cipher_algorithm, "des3", strlen("des3"));
|
||||
memcpy(hash_algorithm, "sha1", strlen("sha1"));
|
||||
}
|
||||
|
||||
snprintf (out_buf, out_len - 1, "%s%d*%d*%s*%s*%s*%d*%s*%d*%s",
|
||||
SIGNATURE_DPAPIMK,
|
||||
version,
|
||||
context,
|
||||
ptr_SID,
|
||||
ptr_cipher_algorithm,
|
||||
ptr_hash_algorithm,
|
||||
SID,
|
||||
cipher_algorithm,
|
||||
hash_algorithm,
|
||||
rounds,
|
||||
iv,
|
||||
contents_len,
|
||||
|
||||
Loading…
Reference in New Issue
Block a user