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

Fix -m 8100 in pure kernel mode for password length 256

Browse Source
This commit is contained in:
jsteube 2019-03-02 17:33:22 +01:00
parent 80bbb064e5
commit 93760dab34
3 changed files with 31 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
OpenCL/m08100_a0-pure.cl
View File

@ -30,6 +30,8 @@ __kernel void m08100_mxx (KERN_ATTR_RULES ())
* base * base
*/ */
u32 z[16] = { 0 };
COPY_PW (pws[gid]); COPY_PW (pws[gid]);
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
@ -50,7 +52,9 @@ __kernel void m08100_mxx (KERN_ATTR_RULES ())
sha1_ctx_t ctx = ctx0; sha1_ctx_t ctx = ctx0;
sha1_update_swap (&ctx, tmp.i, tmp.pw_len + 1); sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
sha1_update (&ctx, z, 1);
sha1_final (&ctx); sha1_final (&ctx);
@ -90,6 +94,8 @@ __kernel void m08100_sxx (KERN_ATTR_RULES ())
* base * base
*/ */
u32 z[16] = { 0 };
COPY_PW (pws[gid]); COPY_PW (pws[gid]);
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
@ -110,7 +116,9 @@ __kernel void m08100_sxx (KERN_ATTR_RULES ())
sha1_ctx_t ctx = ctx0; sha1_ctx_t ctx = ctx0;
sha1_update_swap (&ctx, tmp.i, tmp.pw_len + 1); sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
sha1_update (&ctx, z, 1);
sha1_final (&ctx); sha1_final (&ctx);

13
OpenCL/m08100_a1-pure.cl
View File

@ -28,6 +28,8 @@ __kernel void m08100_mxx (KERN_ATTR_BASIC ())
* base * base
*/ */
u32 z[16] = { 0 };
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
sha1_init (&ctx0); sha1_init (&ctx0);
@ -44,7 +46,9 @@ __kernel void m08100_mxx (KERN_ATTR_BASIC ())
{ {
sha1_ctx_t ctx = ctx0; sha1_ctx_t ctx = ctx0;
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len + 1); sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update (&ctx, z, 1);
sha1_final (&ctx); sha1_final (&ctx);
@ -57,6 +61,7 @@ __kernel void m08100_mxx (KERN_ATTR_BASIC ())
} }
} }
__kernel void m08100_sxx (KERN_ATTR_BASIC ()) __kernel void m08100_sxx (KERN_ATTR_BASIC ())
{ {
/** /**
@ -84,6 +89,8 @@ __kernel void m08100_sxx (KERN_ATTR_BASIC ())
* base * base
*/ */
u32 z[16] = { 0 };
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
sha1_init (&ctx0); sha1_init (&ctx0);
@ -100,7 +107,9 @@ __kernel void m08100_sxx (KERN_ATTR_BASIC ())
{ {
sha1_ctx_t ctx = ctx0; sha1_ctx_t ctx = ctx0;
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len + 1); sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update (&ctx, z, 1);
sha1_final (&ctx); sha1_final (&ctx);

12
OpenCL/m08100_a3-pure.cl
View File

@ -28,6 +28,8 @@ __kernel void m08100_mxx (KERN_ATTR_VECTOR ())
* base * base
*/ */
u32x z[16] = { 0 };
const u32 pw_len = pws[gid].pw_len; const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 }; u32x w[64] = { 0 };
@ -61,7 +63,9 @@ __kernel void m08100_mxx (KERN_ATTR_VECTOR ())
sha1_init_vector_from_scalar (&ctx, &ctx0); sha1_init_vector_from_scalar (&ctx, &ctx0);
sha1_update_vector (&ctx, w, pw_len + 1); sha1_update_vector (&ctx, w, pw_len);
sha1_update_vector (&ctx, z, 1);
sha1_final_vector (&ctx); sha1_final_vector (&ctx);
@ -101,6 +105,8 @@ __kernel void m08100_sxx (KERN_ATTR_VECTOR ())
* base * base
*/ */
u32x z[16] = { 0 };
const u32 pw_len = pws[gid].pw_len; const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 }; u32x w[64] = { 0 };
@ -134,7 +140,9 @@ __kernel void m08100_sxx (KERN_ATTR_VECTOR ())
sha1_init_vector_from_scalar (&ctx, &ctx0); sha1_init_vector_from_scalar (&ctx, &ctx0);
sha1_update_vector (&ctx, w, pw_len + 1); sha1_update_vector (&ctx, w, pw_len);
sha1_update_vector (&ctx, z, 1);
sha1_final_vector (&ctx); sha1_final_vector (&ctx);