From c59432a760852bcfbda50a5790d5cb8976d3fce2 Mon Sep 17 00:00:00 2001 From: jsteube Date: Fri, 23 Jun 2017 12:13:51 +0200 Subject: [PATCH] Add hcstat2 support to enable masks of length up to 256, also adds a filetype header --- OpenCL/markov_be.cl | 59 +++++++++++--------------------------------- OpenCL/markov_le.cl | 59 +++++++++++--------------------------------- hashcat.hcstat | Bin 33685504 -> 0 bytes include/mpsp.h | 5 ++-- src/mpsp.c | 50 +++++++++++++++++++++++++++++++++++++ 5 files changed, 81 insertions(+), 92 deletions(-) delete mode 100644 hashcat.hcstat diff --git a/OpenCL/markov_be.cl b/OpenCL/markov_be.cl index 4f337eb13..955d44c56 100644 --- a/OpenCL/markov_be.cl +++ b/OpenCL/markov_be.cl @@ -9,25 +9,8 @@ #include "inc_types.cl" -inline void generate_pw (u32 pw_buf[16], __global const cs_t *root_css_buf, __global const cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val) +inline void generate_pw (u32 pw_buf[64], __global const cs_t *root_css_buf, __global const cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val) { - pw_buf[ 0] = 0; - pw_buf[ 1] = 0; - pw_buf[ 2] = 0; - pw_buf[ 3] = 0; - pw_buf[ 4] = 0; - pw_buf[ 5] = 0; - pw_buf[ 6] = 0; - pw_buf[ 7] = 0; - pw_buf[ 8] = 0; - pw_buf[ 9] = 0; - pw_buf[10] = 0; - pw_buf[11] = 0; - pw_buf[12] = 0; - pw_buf[13] = 0; - pw_buf[14] = 0; - pw_buf[15] = 0; - __global const cs_t *cs = &root_css_buf[pw_r_len]; u32 i; @@ -67,26 +50,15 @@ __kernel void l_markov (__global pw_t *pws_buf_l, __global const cs_t *root_css_ if (gid >= gid_max) return; - u32 pw_buf[16]; + u32 pw_buf[64] = { 0 }; generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid); - pws_buf_l[gid].i[ 0] = pw_buf[ 0]; - pws_buf_l[gid].i[ 1] = pw_buf[ 1]; - pws_buf_l[gid].i[ 2] = pw_buf[ 2]; - pws_buf_l[gid].i[ 3] = pw_buf[ 3]; - pws_buf_l[gid].i[ 4] = pw_buf[ 4]; - pws_buf_l[gid].i[ 5] = pw_buf[ 5]; - pws_buf_l[gid].i[ 6] = pw_buf[ 6]; - pws_buf_l[gid].i[ 7] = pw_buf[ 7]; - pws_buf_l[gid].i[ 8] = pw_buf[ 8]; - pws_buf_l[gid].i[ 9] = pw_buf[ 9]; - pws_buf_l[gid].i[10] = pw_buf[10]; - pws_buf_l[gid].i[11] = pw_buf[11]; - pws_buf_l[gid].i[12] = pw_buf[12]; - pws_buf_l[gid].i[13] = pw_buf[13]; - pws_buf_l[gid].i[14] = pw_buf[14]; - pws_buf_l[gid].i[15] = pw_buf[15]; + #pragma unroll + for (int idx = 0; idx < 64; idx++) + { + pws_buf_l[gid].i[idx] = pw_buf[idx]; + } pws_buf_l[gid].pw_len = pw_l_len + pw_r_len; } @@ -97,7 +69,7 @@ __kernel void r_markov (__global bf_t *pws_buf_r, __global const cs_t *root_css_ if (gid >= gid_max) return; - u32 pw_buf[16]; + u32 pw_buf[64] = { 0 }; generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid); @@ -110,18 +82,15 @@ __kernel void C_markov (__global comb_t *pws_buf, __global const cs_t *root_css_ if (gid >= gid_max) return; - u32 pw_buf[16]; + u32 pw_buf[64] = { 0 }; generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid); - pws_buf[gid].i[ 0] = pw_buf[ 0]; - pws_buf[gid].i[ 1] = pw_buf[ 1]; - pws_buf[gid].i[ 2] = pw_buf[ 2]; - pws_buf[gid].i[ 3] = pw_buf[ 3]; - pws_buf[gid].i[ 4] = pw_buf[ 4]; - pws_buf[gid].i[ 5] = pw_buf[ 5]; - pws_buf[gid].i[ 6] = pw_buf[ 6]; - pws_buf[gid].i[ 7] = pw_buf[ 7]; + #pragma unroll + for (int idx = 0; idx < 64; idx++) + { + pws_buf[gid].i[idx] = pw_buf[idx]; + } pws_buf[gid].pw_len = pw_len; } diff --git a/OpenCL/markov_le.cl b/OpenCL/markov_le.cl index f20f4770a..cb17da083 100644 --- a/OpenCL/markov_le.cl +++ b/OpenCL/markov_le.cl @@ -9,25 +9,8 @@ #include "inc_types.cl" -inline void generate_pw (u32 pw_buf[16], __global const cs_t *root_css_buf, __global const cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val) +inline void generate_pw (u32 pw_buf[64], __global const cs_t *root_css_buf, __global const cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val) { - pw_buf[ 0] = 0; - pw_buf[ 1] = 0; - pw_buf[ 2] = 0; - pw_buf[ 3] = 0; - pw_buf[ 4] = 0; - pw_buf[ 5] = 0; - pw_buf[ 6] = 0; - pw_buf[ 7] = 0; - pw_buf[ 8] = 0; - pw_buf[ 9] = 0; - pw_buf[10] = 0; - pw_buf[11] = 0; - pw_buf[12] = 0; - pw_buf[13] = 0; - pw_buf[14] = 0; - pw_buf[15] = 0; - __global const cs_t *cs = &root_css_buf[pw_r_len]; u32 i; @@ -67,26 +50,15 @@ __kernel void l_markov (__global pw_t *pws_buf_l, __global const cs_t *root_css_ if (gid >= gid_max) return; - u32 pw_buf[16]; + u32 pw_buf[64] = { 0 }; generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid); - pws_buf_l[gid].i[ 0] = pw_buf[ 0]; - pws_buf_l[gid].i[ 1] = pw_buf[ 1]; - pws_buf_l[gid].i[ 2] = pw_buf[ 2]; - pws_buf_l[gid].i[ 3] = pw_buf[ 3]; - pws_buf_l[gid].i[ 4] = pw_buf[ 4]; - pws_buf_l[gid].i[ 5] = pw_buf[ 5]; - pws_buf_l[gid].i[ 6] = pw_buf[ 6]; - pws_buf_l[gid].i[ 7] = pw_buf[ 7]; - pws_buf_l[gid].i[ 8] = pw_buf[ 8]; - pws_buf_l[gid].i[ 9] = pw_buf[ 9]; - pws_buf_l[gid].i[10] = pw_buf[10]; - pws_buf_l[gid].i[11] = pw_buf[11]; - pws_buf_l[gid].i[12] = pw_buf[12]; - pws_buf_l[gid].i[13] = pw_buf[13]; - pws_buf_l[gid].i[14] = pw_buf[14]; - pws_buf_l[gid].i[15] = pw_buf[15]; + #pragma unroll + for (int idx = 0; idx < 64; idx++) + { + pws_buf_l[gid].i[idx] = pw_buf[idx]; + } pws_buf_l[gid].pw_len = pw_l_len + pw_r_len; } @@ -97,7 +69,7 @@ __kernel void r_markov (__global bf_t *pws_buf_r, __global const cs_t *root_css_ if (gid >= gid_max) return; - u32 pw_buf[16]; + u32 pw_buf[64] = { 0 }; generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid); @@ -110,18 +82,15 @@ __kernel void C_markov (__global comb_t *pws_buf, __global const cs_t *root_css_ if (gid >= gid_max) return; - u32 pw_buf[16]; + u32 pw_buf[64] = { 0 }; generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid); - pws_buf[gid].i[ 0] = pw_buf[ 0]; - pws_buf[gid].i[ 1] = pw_buf[ 1]; - pws_buf[gid].i[ 2] = pw_buf[ 2]; - pws_buf[gid].i[ 3] = pw_buf[ 3]; - pws_buf[gid].i[ 4] = pw_buf[ 4]; - pws_buf[gid].i[ 5] = pw_buf[ 5]; - pws_buf[gid].i[ 6] = pw_buf[ 6]; - pws_buf[gid].i[ 7] = pw_buf[ 7]; + #pragma unroll + for (int idx = 0; idx < 64; idx++) + { + pws_buf[gid].i[idx] = pw_buf[idx]; + } pws_buf[gid].pw_len = pw_len; } diff --git a/hashcat.hcstat b/hashcat.hcstat deleted file mode 100644 index 983d73638c00198f2d2a1564eec1a69b520c0bbf..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 33685504 zcmeF4b-Y&9_Vt%yU=Sim!{N{^-QC@tg3^t&BB2O^fP#Px5@LXgpqSVw7A6KN79th` zcD=uM&+pj&?&h8syn63@?>&=07|)z*t{VG%)|hMU{p=?M|M!0x$lvz3h=?IsjE^VY zFg(=Z_~^s8h0XpspW*gX471fRJPEsCacgIKQS*<;Z8#{^aP4`9?P;g(N0lPSoV+U| z?0I{S2=hNsCBjomm5;E_HDx2be6?ZCkIF^%uTu;kTwFS`uRke*a24~oE34Usu`jUc zgtB?epKHC1w`{&MBS)qU*&;k+^%)VqSNHS?Z~e&{K7M(|$j*A*sS$2`(I!6j`4c01 z{jf|Cwrpzt2i`P$S;!Unx6Zcujeg4z*^P_a09y^Wag6J3{nh;0;%9fU`onLu@;B~1 zA*y#-X&X_r>n6+t*Pmwoxm=%LZ5rY2oHL?=L;FWZxZ<=) z5uRASU4$b`8SgcgRg3JdXEluQqn7puP@%v5ffUK#KJxFqqi%#J)#?=Ce6DYX$E^Mp zPuT^Sa{r*H{I!8LkUF2YiR^YQ?}%{hxg8^HJH^VI3{Q#drCn?S+b=iXdla+_+_9cv zDWfUuGTOqps8$FsX6A!~GvEw31L>TBsV1A@{AotSc@Uxc%_roj8nO}_HRK6LXMt8m3@CK&CW zHY_{E@VarPGfU9hiz}Kx^Xo>~o_bdPb8hoLklpa}Gp#&hf9vOkPt1Qh`9kG!X79Mv z#(Ca!i~qj2#ecWR%F8w}f96dVx8-NU5+e+!>^2-d%CM^qCwy$xLq4k>M#IP8S{By^ zK6+HSEs8w*6}xcj8yAQIo-Awj_fNhV*}FFyaX+7Igl2BsEAsE0|4f7jvz!06O_o2+ zd%s)cPkFXfgr7Z9HNvdx4oA3Si{$~I-C!5)>kShk|G|01BV4e)P=pJ68bR~kx9c=J zbN0x8_-C7Fr6rbVmb=;VqNk(M5XvzB?%3IwFQ$Vt;0!ne>6U@jw%raBIvWLzSs&N7 zGJD1x!^^WlQ4zx{Z95howe4)!ik-Ekm5)aEpDb$jS4#~q-D=pcf#EqjMn)eFePap2 z;o8k2`>DFNZ`}7_g~{!IIQ0K9uFu}DjL*kb8Qy0y7_PMN zuo_71^1KnvUn)mumKX z<0Cvbr~RQUUsfct=iFkt{X@QUB71Rn`vX{8YCvSaQLcQ1KU`$<&b0Eb$lm+i#0bAW zt$&0qc3M6-@JsuH`usLK;CO3`UFcfxS-vtSrzNWQMYTeh%lvy{e~69g;0!ne&Oo|l zpv(%J`eQQizkF-$E?aDIyE!iy z-PH8!d-mb?FSYVZU$XkeY8wuG%dqQg!!8$E+}TD$m~gxexY=)3zICTz=c_EPcrd*B z0UO7Vd1l{4KCsc`GZaMsH(7D$PJUJa%=Nz2KV;2^8~SaI9Cv(PCc@Ht8%B7>+ifDe zZr|+@X1m6S9CKgg$bRVrPFmJe?JZggbVzQg)E|I~t!z4nYVBW&2-5}G`m6CD4x^}j7D4PgTF zx86b^4Wrb-8E^)if%MBj+1G6Dc{dt zzGK_;aE-}HcU)>B2dpVm zH?k+3Q$NCc_F3E|zZm9X99eC97@C-@gpK*4@Q{c3tv2}x50^2=HOo!^uejUBxAK7D zmPZX&?X!j^eqrqtY+(LP94JiOX6=+ihkLfM@-Df||5Hm7+B|v9F7k!h!`fNgx6Q2G z8lRhg$SD@L=z6o;@q9+54HnnsCBy8!uxD61O=p<>atX7)HweG7{~w;o%6C_=f-9_> zu)EXts36zmViCSH&@N!URfdalzZdxzuCfciWPs^#iD|YkzkA~iQTfZw?GNJQPb>j^ zY|ntmUv*X02yg}-%c+=Xu;nD#SYGZGvExQTLxbK+@`*j`_}o; ziB<2L|NL7G*A+2*>s)N~vjus;TecjB%g86*sbO)mDjH_}$uQ?bwjej=K(<7sX;Hyl z&2NeDW=4e*@Q} zau+KvInMl>xsFS%nUDwjIlG-7J|KQN{-WjXiz?=Qu0VuMA2$6yu+sFu?KnGFX!M5Z z`>xY1Zr|0G;B|O#YE(hlb0av+`?9w!^NRbE1N|-R%G(!xI;nU9)3^UHdhO@Un)@B0M;@ zd4y|uVBzMirVrJBwD@X_rvdrME6nHV!cn!bfP5fBC9~__YI^&22`f1N1RGzWl7