arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-23 00:28:11 +00:00
Code Issues Releases Wiki Activity

Fixed -m 600 in -a 3 mode for passwords > 16

Browse Source
This commit is contained in:
jsteube 2019-02-13 13:48:31 +01:00
parent bb8827f2e0
commit dd293f7a93
1 changed files with 440 additions and 16 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

456
OpenCL/m00600_a3-optimized.cl
View File

@ -152,6 +152,214 @@ __kernel void m00600_m04 (KERN_ATTR_VECTOR_ESALT (blake2_t))
u32 w0l = pws[gid].i[0];
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
const u32x w0x = w0l | w0r;
u32x w0[4];
u32x w1[4];
u32x w2[4];
u32x w3[4];
w0[0] = w0x;
w0[1] = pws[gid].i[ 1];
w0[2] = pws[gid].i[ 2];
w0[3] = pws[gid].i[ 3];
w1[0] = 0;
w1[1] = 0;
w1[2] = 0;
w1[3] = 0;
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
u32x out_len = pws[gid].pw_len;
u64x digest[8];
u64x m[16];
u64x v[16];
u64x h[8];
u64x t[2];
u64x f[2];
h[0] = tmp_h[0];
h[1] = tmp_h[1];
h[2] = tmp_h[2];
h[3] = tmp_h[3];
h[4] = tmp_h[4];
h[5] = tmp_h[5];
h[6] = tmp_h[6];
h[7] = tmp_h[7];
t[0] = tmp_t[0];
t[1] = tmp_t[1];
f[0] = tmp_f[0];
f[1] = tmp_f[1];
blake2b_transform(h, t, f, m, v, w0, w1, w2, w3, out_len, BLAKE2B_FINAL);
digest[0] = h[0];
digest[1] = h[1];
digest[2] = h[2];
digest[3] = h[3];
digest[4] = h[4];
digest[5] = h[5];
digest[6] = h[6];
digest[7] = h[7];
const u32x r0 = h32_from_64(digest[0]);
const u32x r1 = l32_from_64(digest[0]);
const u32x r2 = h32_from_64(digest[1]);
const u32x r3 = l32_from_64(digest[1]);
COMPARE_M_SIMD(r0, r1, r2, r3);
}
}
__kernel void m00600_m08 (KERN_ATTR_VECTOR_ESALT (blake2_t))
{
/**
* modifier
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
u64 tmp_h[8];
u64 tmp_t[2];
u64 tmp_f[2];
tmp_h[0] = esalt_bufs[digests_offset].h[0];
tmp_h[1] = esalt_bufs[digests_offset].h[1];
tmp_h[2] = esalt_bufs[digests_offset].h[2];
tmp_h[3] = esalt_bufs[digests_offset].h[3];
tmp_h[4] = esalt_bufs[digests_offset].h[4];
tmp_h[5] = esalt_bufs[digests_offset].h[5];
tmp_h[6] = esalt_bufs[digests_offset].h[6];
tmp_h[7] = esalt_bufs[digests_offset].h[7];
tmp_t[0] = esalt_bufs[digests_offset].t[0];
tmp_t[1] = esalt_bufs[digests_offset].t[1];
tmp_f[0] = esalt_bufs[digests_offset].f[0];
tmp_f[1] = esalt_bufs[digests_offset].f[1];
/**
* loop
*/
u32 w0l = pws[gid].i[0];
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
const u32x w0x = w0l | w0r;
u32x w0[4];
u32x w1[4];
u32x w2[4];
u32x w3[4];
w0[0] = w0x;
w0[1] = pws[gid].i[ 1];
w0[2] = pws[gid].i[ 2];
w0[3] = pws[gid].i[ 3];
w1[0] = pws[gid].i[ 4];
w1[1] = pws[gid].i[ 5];
w1[2] = pws[gid].i[ 6];
w1[3] = pws[gid].i[ 7];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
u32x out_len = pws[gid].pw_len;
u64x digest[8];
u64x m[16];
u64x v[16];
u64x h[8];
u64x t[2];
u64x f[2];
h[0] = tmp_h[0];
h[1] = tmp_h[1];
h[2] = tmp_h[2];
h[3] = tmp_h[3];
h[4] = tmp_h[4];
h[5] = tmp_h[5];
h[6] = tmp_h[6];
h[7] = tmp_h[7];
t[0] = tmp_t[0];
t[1] = tmp_t[1];
f[0] = tmp_f[0];
f[1] = tmp_f[1];
blake2b_transform(h, t, f, m, v, w0, w1, w2, w3, out_len, BLAKE2B_FINAL);
digest[0] = h[0];
digest[1] = h[1];
digest[2] = h[2];
digest[3] = h[3];
digest[4] = h[4];
digest[5] = h[5];
digest[6] = h[6];
digest[7] = h[7];
const u32x r0 = h32_from_64(digest[0]);
const u32x r1 = l32_from_64(digest[0]);
const u32x r2 = h32_from_64(digest[1]);
const u32x r3 = l32_from_64(digest[1]);
COMPARE_M_SIMD(r0, r1, r2, r3);
}
}
__kernel void m00600_m16 (KERN_ATTR_VECTOR_ESALT (blake2_t))
{
/**
* modifier
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
u64 tmp_h[8];
u64 tmp_t[2];
u64 tmp_f[2];
tmp_h[0] = esalt_bufs[digests_offset].h[0];
tmp_h[1] = esalt_bufs[digests_offset].h[1];
tmp_h[2] = esalt_bufs[digests_offset].h[2];
tmp_h[3] = esalt_bufs[digests_offset].h[3];
tmp_h[4] = esalt_bufs[digests_offset].h[4];
tmp_h[5] = esalt_bufs[digests_offset].h[5];
tmp_h[6] = esalt_bufs[digests_offset].h[6];
tmp_h[7] = esalt_bufs[digests_offset].h[7];
tmp_t[0] = esalt_bufs[digests_offset].t[0];
tmp_t[1] = esalt_bufs[digests_offset].t[1];
tmp_f[0] = esalt_bufs[digests_offset].f[0];
tmp_f[1] = esalt_bufs[digests_offset].f[1];
/**
* loop
*/
u32 w0l = pws[gid].i[0];
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
@ -223,14 +431,6 @@ __kernel void m00600_m04 (KERN_ATTR_VECTOR_ESALT (blake2_t))
}
}
__kernel void m00600_m08 (KERN_ATTR_ESALT (blake2_t))
{
}
__kernel void m00600_m16 (KERN_ATTR_ESALT (blake2_t))
{
}
__kernel void m00600_s04 (KERN_ATTR_VECTOR_ESALT (blake2_t))
{
/**
@ -276,6 +476,238 @@ __kernel void m00600_s04 (KERN_ATTR_VECTOR_ESALT (blake2_t))
u32 w0l = pws[gid].i[0];
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
const u32x w0x = w0l | w0r;
u32x w0[4];
u32x w1[4];
u32x w2[4];
u32x w3[4];
w0[0] = w0x;
w0[1] = pws[gid].i[ 1];
w0[2] = pws[gid].i[ 2];
w0[3] = pws[gid].i[ 3];
w1[0] = 0;
w1[1] = 0;
w1[2] = 0;
w1[3] = 0;
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
u32x out_len = pws[gid].pw_len;
u64x digest[8];
u64x m[16];
u64x v[16];
u64x h[8];
u64x t[2];
u64x f[2];
h[0] = tmp_h[0];
h[1] = tmp_h[1];
h[2] = tmp_h[2];
h[3] = tmp_h[3];
h[4] = tmp_h[4];
h[5] = tmp_h[5];
h[6] = tmp_h[6];
h[7] = tmp_h[7];
t[0] = tmp_t[0];
t[1] = tmp_t[1];
f[0] = tmp_f[0];
f[1] = tmp_f[1];
blake2b_transform(h, t, f, m, v, w0, w1, w2, w3, out_len, BLAKE2B_FINAL);
digest[0] = h[0];
digest[1] = h[1];
digest[2] = h[2];
digest[3] = h[3];
digest[4] = h[4];
digest[5] = h[5];
digest[6] = h[6];
digest[7] = h[7];
const u32x r0 = h32_from_64(digest[0]);
const u32x r1 = l32_from_64(digest[0]);
const u32x r2 = h32_from_64(digest[1]);
const u32x r3 = l32_from_64(digest[1]);
COMPARE_S_SIMD(r0, r1, r2, r3);
}
}
__kernel void m00600_s08 (KERN_ATTR_VECTOR_ESALT (blake2_t))
{
/**
* modifier
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
u64 tmp_h[8];
u64 tmp_t[2];
u64 tmp_f[2];
tmp_h[0] = esalt_bufs[digests_offset].h[0];
tmp_h[1] = esalt_bufs[digests_offset].h[1];
tmp_h[2] = esalt_bufs[digests_offset].h[2];
tmp_h[3] = esalt_bufs[digests_offset].h[3];
tmp_h[4] = esalt_bufs[digests_offset].h[4];
tmp_h[5] = esalt_bufs[digests_offset].h[5];
tmp_h[6] = esalt_bufs[digests_offset].h[6];
tmp_h[7] = esalt_bufs[digests_offset].h[7];
tmp_t[0] = esalt_bufs[digests_offset].t[0];
tmp_t[1] = esalt_bufs[digests_offset].t[1];
tmp_f[0] = esalt_bufs[digests_offset].f[0];
tmp_f[1] = esalt_bufs[digests_offset].f[1];
/**
* digest
*/
const u32 search[4] =
{
digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2],
digests_buf[digests_offset].digest_buf[DGST_R3]
};
/**
* loop
*/
u32 w0l = pws[gid].i[0];
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
const u32x w0x = w0l | w0r;
u32x w0[4];
u32x w1[4];
u32x w2[4];
u32x w3[4];
w0[0] = w0x;
w0[1] = pws[gid].i[ 1];
w0[2] = pws[gid].i[ 2];
w0[3] = pws[gid].i[ 3];
w1[0] = pws[gid].i[ 4];
w1[1] = pws[gid].i[ 5];
w1[2] = pws[gid].i[ 6];
w1[3] = pws[gid].i[ 7];
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = 0;
u32x out_len = pws[gid].pw_len;
u64x digest[8];
u64x m[16];
u64x v[16];
u64x h[8];
u64x t[2];
u64x f[2];
h[0] = tmp_h[0];
h[1] = tmp_h[1];
h[2] = tmp_h[2];
h[3] = tmp_h[3];
h[4] = tmp_h[4];
h[5] = tmp_h[5];
h[6] = tmp_h[6];
h[7] = tmp_h[7];
t[0] = tmp_t[0];
t[1] = tmp_t[1];
f[0] = tmp_f[0];
f[1] = tmp_f[1];
blake2b_transform(h, t, f, m, v, w0, w1, w2, w3, out_len, BLAKE2B_FINAL);
digest[0] = h[0];
digest[1] = h[1];
digest[2] = h[2];
digest[3] = h[3];
digest[4] = h[4];
digest[5] = h[5];
digest[6] = h[6];
digest[7] = h[7];
const u32x r0 = h32_from_64(digest[0]);
const u32x r1 = l32_from_64(digest[0]);
const u32x r2 = h32_from_64(digest[1]);
const u32x r3 = l32_from_64(digest[1]);
COMPARE_S_SIMD(r0, r1, r2, r3);
}
}
__kernel void m00600_s16 (KERN_ATTR_VECTOR_ESALT (blake2_t))
{
/**
* modifier
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
u64 tmp_h[8];
u64 tmp_t[2];
u64 tmp_f[2];
tmp_h[0] = esalt_bufs[digests_offset].h[0];
tmp_h[1] = esalt_bufs[digests_offset].h[1];
tmp_h[2] = esalt_bufs[digests_offset].h[2];
tmp_h[3] = esalt_bufs[digests_offset].h[3];
tmp_h[4] = esalt_bufs[digests_offset].h[4];
tmp_h[5] = esalt_bufs[digests_offset].h[5];
tmp_h[6] = esalt_bufs[digests_offset].h[6];
tmp_h[7] = esalt_bufs[digests_offset].h[7];
tmp_t[0] = esalt_bufs[digests_offset].t[0];
tmp_t[1] = esalt_bufs[digests_offset].t[1];
tmp_f[0] = esalt_bufs[digests_offset].f[0];
tmp_f[1] = esalt_bufs[digests_offset].f[1];
/**
* digest
*/
const u32 search[4] =
{
digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2],
digests_buf[digests_offset].digest_buf[DGST_R3]
};
/**
* loop
*/
u32 w0l = pws[gid].i[0];
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
@ -346,11 +778,3 @@ __kernel void m00600_s04 (KERN_ATTR_VECTOR_ESALT (blake2_t))
COMPARE_S_SIMD(r0, r1, r2, r3);
}
}
__kernel void m00600_s08 (KERN_ATTR_ESALT (blake2_t))
{
}
__kernel void m00600_s16 (KERN_ATTR_ESALT (blake2_t))
{
}