diff --git a/OpenCL/m00600_a3.cl b/OpenCL/m00600_a3.cl index 79c0a2dde..252609030 100644 --- a/OpenCL/m00600_a3.cl +++ b/OpenCL/m00600_a3.cl @@ -12,13 +12,38 @@ #include "inc_common.cl" #include "inc_simd.cl" -#if defined(_MSC_VER) -#define BLAKE2_PACKED(x) __pragma(pack(push, 1)) x __pragma(pack(pop)) -#else -#define BLAKE2_PACKED(x) x __attribute__((packed)) -#endif +u64 rotr64_w( const u64x w, const u32 c ) +{ + return ( w >> c ) | ( w << ( 64 - c ) ); +} -BLAKE2_PACKED(struct blake2b_param__ +u64 load64( const void *src ) +{ + const u8 *p = ( const u8 * )src; + return (( u64 )( p[0] ) << 0) | + (( u64 )( p[1] ) << 8) | + (( u64 )( p[2] ) << 16) | + (( u64 )( p[3] ) << 24) | + (( u64 )( p[4] ) << 32) | + (( u64 )( p[5] ) << 40) | + (( u64 )( p[6] ) << 48) | + (( u64 )( p[7] ) << 56) ; +} + +u64 load64_inv( const void *src ) +{ + const u8 *p = ( const u8 * )src; + return (( u64 )( p[7] ) << 0) | + (( u64 )( p[6] ) << 8) | + (( u64 )( p[5] ) << 16) | + (( u64 )( p[4] ) << 24) | + (( u64 )( p[3] ) << 32) | + (( u64 )( p[2] ) << 40) | + (( u64 )( p[1] ) << 48) | + (( u64 )( p[0] ) << 56) ; +} + +typedef struct blake2b_param__ { u8 digest_length; /* 1 */ u8 key_length; /* 2 */ @@ -32,9 +57,7 @@ BLAKE2_PACKED(struct blake2b_param__ u8 reserved[14]; /* 32 */ u8 salt[BLAKE2B_SALTBYTES]; /* 48 */ u8 personal[BLAKE2B_PERSONALBYTES]; /* 64 */ -}); - -typedef struct blake2b_param__ blake2b_param; +} blake2b_param; typedef struct { @@ -47,7 +70,7 @@ typedef struct u8 last_node; } blake2b_state; -__constant u64 blake2b_IV[8] = +__constant u64a blake2b_IV[8] = { 0x6a09e667f3bcc908, 0xbb67ae8584caa73b, 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, @@ -71,20 +94,20 @@ __constant u8 blake2b_sigma[12][16] = { 14, 10, 4, 8, 9, 15, 13, 6, 1, 12, 0, 2, 11, 7, 5, 3 } }; -#define BLAKE2B_G(r,i,a,b,c,d) \ - do { \ - a = a + b + m[blake2b_sigma[r][2*i+0]]; \ - d = rotr64(d ^ a, 32); \ - c = c + d; \ - b = rotr64(b ^ c, 24); \ - a = a + b + m[blake2b_sigma[r][2*i+1]]; \ - d = rotr64(d ^ a, 16); \ - c = c + d; \ - b = rotr64(b ^ c, 63); \ +#define BLAKE2B_G(r,i,a,b,c,d) \ + do { \ + a = a + b + m[blake2b_sigma[r][2*i+0]]; \ + d = rotr64_w(d ^ a, 32); \ + c = c + d; \ + b = rotr64_w(b ^ c, 24); \ + a = a + b + m[blake2b_sigma[r][2*i+1]]; \ + d = rotr64_w(d ^ a, 16); \ + c = c + d; \ + b = rotr64_w(b ^ c, 63); \ } while(0) #define BLAKE2B_ROUND(r) \ - do { \ + do { \ BLAKE2B_G(r,0,v[ 0],v[ 4],v[ 8],v[12]); \ BLAKE2B_G(r,1,v[ 1],v[ 5],v[ 9],v[13]); \ BLAKE2B_G(r,2,v[ 2],v[ 6],v[10],v[14]); \ @@ -95,7 +118,7 @@ __constant u8 blake2b_sigma[12][16] = BLAKE2B_G(r,7,v[ 3],v[ 4],v[ 9],v[14]); \ } while(0) -void blake2b_compress (const u32x *pw, u64x digest[8]) +void blake2b_compress (const u32x pw[16], const u64 pw_len, u64x digest[8]) { /* @@ -114,21 +137,27 @@ void blake2b_compress (const u32x *pw, u64x digest[8]) P->xof_length = 0; P->node_depth = 0; P->inner_length = 0; - for (i = 0; i < 14; i++) P->reserved[i] = 0; - for (i = 0; i < BLAKE2B_SALTBYTES; i++) P->salt[i] = 0; - for (i = 0; i < BLAKE2B_PERSONALBYTES; i++) P->personal[i] = 0; + + for (i = 0; i < 14; i++) + P->reserved[i] = 0; + for (i = 0; i < BLAKE2B_SALTBYTES; i++) + P->salt[i] = 0; + for (i = 0; i < BLAKE2B_PERSONALBYTES; i++) + P->personal[i] = 0; /* * Blake2b Init State */ + blake2b_state S[1]; - for (i = 0; i < 8; i++) + for (i = 0; i < 8; ++i) S->h[i] = blake2b_IV[i]; - S->t[0] = 0; - S->t[1] = 0; - S->f[0] = 0; - S->f[1] = 0; + + S->t[0] = pw_len; + S->t[1] = 0; + S->f[0] = -1; + S->f[1] = 0; S->buflen = 0; S->outlen = 0; S->last_node = 0; @@ -137,15 +166,30 @@ void blake2b_compress (const u32x *pw, u64x digest[8]) /* IV XOR ParamBlock */ for( i = 0; i < 8; ++i ) - S->h[i] ^= *(p + sizeof(S->h[i]) * i); + S->h[i] ^= load64(p + sizeof(S->h[i]) * i); S->outlen = P->digest_length; /* * Compress */ + u64 v[16]; - u64 *m = pw; + u64 m[16]; + + p = (const u8 *)pw; + + for (i = 0; i < 8; ++i) + m[i] = load64_inv(p + sizeof(m[i]) * i); + + m[8] = 0; + m[9] = 0; + m[10] = 0; + m[11] = 0; + m[12] = 0; + m[13] = 0; + m[14] = 0; + m[15] = 0; for ( i = 0; i < 8; ++i ) v[i] = S->h[i]; @@ -172,15 +216,15 @@ void blake2b_compress (const u32x *pw, u64x digest[8]) BLAKE2B_ROUND( 10 ); BLAKE2B_ROUND( 11 ); - for ( i = 0; i < 8; i++ ) { + for ( i = 0; i < 8; ++i ) { S->h[i] = S->h[i] ^ v[i] ^ v[i + 8]; } - for ( i = 0; i < 8; i++ ) - digest[i] = S->h[i]; + for ( i = 0; i < 8; ++i ) + digest[i] = load64_inv(&(S->h[i])); } -void m00600s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) +void m00600s (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset) { /** * modifier @@ -205,7 +249,7 @@ void m00600s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke * loop */ - u32 w0l = w[0]; + u32 w0l = pws[gid].i[0]; for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { @@ -215,22 +259,24 @@ void m00600s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke u32x pw[16]; - pw[ 0] = w0; - pw[ 1] = w[ 1]; - pw[ 2] = w[ 2]; - pw[ 3] = w[ 3]; - pw[ 4] = w[ 4]; - pw[ 5] = w[ 5]; - pw[ 6] = w[ 6]; - pw[ 7] = w[ 7]; - pw[ 8] = w[ 8]; - pw[ 9] = w[ 9]; - pw[10] = w[10]; - pw[11] = w[11]; - pw[12] = w[12]; - pw[13] = w[13]; - pw[14] = w[14]; - pw[15] = w[15]; + pw[ 1] = w0; + pw[ 0] = pws[gid].i[ 1]; + pw[ 3] = pws[gid].i[ 2]; + pw[ 2] = pws[gid].i[ 3]; + pw[ 5] = pws[gid].i[ 4]; + pw[ 4] = pws[gid].i[ 5]; + pw[ 7] = pws[gid].i[ 6]; + pw[ 6] = pws[gid].i[ 7]; + pw[ 9] = pws[gid].i[ 8]; + pw[ 8] = pws[gid].i[ 9]; + pw[11] = pws[gid].i[10]; + pw[10] = pws[gid].i[11]; + pw[13] = pws[gid].i[12]; + pw[12] = pws[gid].i[13]; + pw[15] = pws[gid].i[14]; + pw[15] = pws[gid].i[15]; + + u32 pw_len = pws[gid].pw_len; u64x digest[8]; @@ -243,12 +289,12 @@ void m00600s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke digest[6] = 0; digest[7] = 0; - blake2b_compress (&pw, digest); + blake2b_compress (pw, pw_len, digest); - const u32x r0 = l32_from_64 (digest[7]); - const u32x r1 = h32_from_64 (digest[7]); - const u32x r2 = l32_from_64 (digest[3]); - const u32x r3 = h32_from_64 (digest[3]); + 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); } @@ -270,41 +316,8 @@ __kernel void m00600_m16 (__global pw_t *pws, __global const kernel_rule_t *rule __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { - - /* - * base - */ - - const u32 gid = get_global_id (0); - - if (gid >= gid_max) return; - - u32 w[16]; - - w[ 0] = pws[gid].i[ 0]; - w[ 1] = pws[gid].i[ 1]; - w[ 2] = pws[gid].i[ 2]; - w[ 3] = pws[gid].i[ 3]; - w[ 4] = 0; - w[ 5] = 0; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 0; - w[15] = pws[gid].i[15]; - - const u32 pw_len = pws[gid].pw_len; - - /* - * main - */ - - m00600s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); + + m00600s (pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); } __kernel void m00600_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) diff --git a/b2test.sh b/b2test.sh index f119a92a4..9c254bcfd 100755 --- a/b2test.sh +++ b/b2test.sh @@ -1 +1,2 @@ -./hashcat -m 600 -a 3 b2test.hash ?d?d?d?d?d?d +./hashcat -m 600 --potfile-disable --weak-hash-threshold 0 -u1 -n1 --force -a 3 b2test.hash ?d?d?d?d?d?d + diff --git a/src/Makefile b/src/Makefile index b572e771d..bcebe2a6e 100644 --- a/src/Makefile +++ b/src/Makefile @@ -4,7 +4,7 @@ ## SHARED := 0 -DEBUG := 0 +DEBUG := 1 PRODUCTION := 0 ## diff --git a/src/interface.c b/src/interface.c index f72129f24..5836be44b 100644 --- a/src/interface.c +++ b/src/interface.c @@ -19254,11 +19254,10 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) break; case 600: hashconfig->hash_type = HASH_TYPE_BLAKE2B; - hashconfig->salt_type = SALT_TYPE_NONE; + hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL; hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE - | OPTS_TYPE_PT_ADD80 - | OPTS_TYPE_PT_ADDBITS15; + | OPTS_TYPE_ST_ADDBITS15; hashconfig->kern_type = KERN_TYPE_BLAKE2B; hashconfig->dgst_size = DGST_SIZE_8_8; hashconfig->parse_func = blake2b_parse_hash; @@ -19267,10 +19266,10 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) | OPTI_TYPE_NOT_SALTED | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_RAW_HASH; - hashconfig->dgst_pos0 = 14; - hashconfig->dgst_pos1 = 15; - hashconfig->dgst_pos2 = 6; - hashconfig->dgst_pos3 = 7; + hashconfig->dgst_pos0 = 1; + hashconfig->dgst_pos1 = 0; + hashconfig->dgst_pos2 = 3; + hashconfig->dgst_pos3 = 2; break; case 900: hashconfig->hash_type = HASH_TYPE_MD4;