fix: Kernel m00600_a3.cl function s04.

OpenCL/m00600_a3.cl

@@ -12,13 +12,38 @@
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
-#if defined(_MSC_VER)
+u64 rotr64_w( const u64x w, const u32 c )
-#define BLAKE2_PACKED(x) __pragma(pack(push, 1)) x __pragma(pack(pop))
+{
-#else
+  return ( w >> c ) | ( w << ( 64 - c ) );
-#define BLAKE2_PACKED(x) x __attribute__((packed))
+}
-#endif
 
-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 */
-});
+} blake2b_param;
-
-typedef struct 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)                      \
+#define BLAKE2B_G(r,i,a,b,c,d)                \
-  do {                                      \
+  do {                                        \
-    a = a + b + m[blake2b_sigma[r][2*i+0]]; \
+    a = a + b + m[blake2b_sigma[r][2*i+0]];   \
-    d = rotr64(d ^ a, 32);                  \
+    d = rotr64_w(d ^ a, 32);                  \
-    c = c + d;                              \
+    c = c + d;                                \
-    b = rotr64(b ^ c, 24);                  \
+    b = rotr64_w(b ^ c, 24);                  \
-    a = a + b + m[blake2b_sigma[r][2*i+1]]; \
+    a = a + b + m[blake2b_sigma[r][2*i+1]];   \
-    d = rotr64(d ^ a, 16);                  \
+    d = rotr64_w(d ^ a, 16);                  \
-    c = c + d;                              \
+    c = c + d;                                \
-    b = rotr64(b ^ c, 63);                  \
+    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 < 14; i++) 
-  for (i = 0; i < BLAKE2B_PERSONALBYTES; i++) P->personal[i] = 0;
+    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->t[0] =  pw_len;
-  S->f[0] = 0;
+  S->t[1] =  0;
-  S->f[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++ ) 
+  for ( i = 0; i < 8; ++i ) 
-    digest[i] = S->h[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] = w0;
-    pw[ 1] = w[ 1];
+    pw[ 0] = pws[gid].i[ 1];
-    pw[ 2] = w[ 2];
+    pw[ 3] = pws[gid].i[ 2];
-    pw[ 3] = w[ 3];
+    pw[ 2] = pws[gid].i[ 3];
-    pw[ 4] = w[ 4];
+    pw[ 5] = pws[gid].i[ 4];
-    pw[ 5] = w[ 5];
+    pw[ 4] = pws[gid].i[ 5];
-    pw[ 6] = w[ 6];
+    pw[ 7] = pws[gid].i[ 6];
-    pw[ 7] = w[ 7];
+    pw[ 6] = pws[gid].i[ 7];
-    pw[ 8] = w[ 8];
+    pw[ 9] = pws[gid].i[ 8];
-    pw[ 9] = w[ 9];
+    pw[ 8] = pws[gid].i[ 9];
-    pw[10] = w[10];
+    pw[11] = pws[gid].i[10];
-    pw[11] = w[11];
+    pw[10] = pws[gid].i[11];
-    pw[12] = w[12];
+    pw[13] = pws[gid].i[12];
-    pw[13] = w[13];
+    pw[12] = pws[gid].i[13];
-    pw[14] = w[14];
+    pw[15] = pws[gid].i[14];
-    pw[15] = w[15];
+    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 r0 = h32_from_64 (digest[0]);
-    const u32x r1 = h32_from_64 (digest[7]);
+    const u32x r1 = l32_from_64 (digest[0]);
-    const u32x r2 = l32_from_64 (digest[3]);
+    const u32x r2 = h32_from_64 (digest[1]);
-    const u32x r3 = h32_from_64 (digest[3]);
+    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)
 { 
-
+ 
-  /*
+  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);  
-   * 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);  
 }
 
 __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)

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
+

src/Makefile

@@ -4,7 +4,7 @@
 ##
 
 SHARED                  := 0
-DEBUG                   := 0
+DEBUG                   := 1
 PRODUCTION              := 0
 
 ##

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_ST_ADDBITS15;
-                                            | OPTS_TYPE_PT_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_pos0      = 1;
-                 hashconfig->dgst_pos1      = 15;
+                 hashconfig->dgst_pos1      = 0;
-                 hashconfig->dgst_pos2      = 6;
+                 hashconfig->dgst_pos2      = 3;
-                 hashconfig->dgst_pos3      = 7;
+                 hashconfig->dgst_pos3      = 2;
                  break;
 
     case   900:  hashconfig->hash_type      = HASH_TYPE_MD4;