Merge branch 'hashcat:master' into master

OpenCL/inc_common.cl

@@ -2066,6 +2066,12 @@ DECLSPEC int hc_enc_next (hc_enc_t *hc_enc, const u32 *src_buf, const int src_le
 
   int src_pos = hc_enc->pos;
 
+  #if VENDOR_ID == 8
+  // Work around segmentation fault in Intel JiT
+  // Tested with 2021.12.6.0.19_160000
+  volatile
+  #endif
+
   int dst_pos = hc_enc->clen;
 
   dst_buf[0] = hc_enc->cbuf;
@@ -2197,6 +2203,12 @@ DECLSPEC int hc_enc_next_global (hc_enc_t *hc_enc, GLOBAL_AS const u32 *src_buf,
 
   int src_pos = hc_enc->pos;
 
+  #if VENDOR_ID == 8
+  // Work around segmentation fault in Intel JiT
+  // Tested with 2021.12.6.0.19_160000
+  volatile
+  #endif
+
   int dst_pos = hc_enc->clen;
 
   dst_buf[0] = hc_enc->cbuf;
@@ -2300,6 +2312,8 @@ DECLSPEC int hc_enc_next_global (hc_enc_t *hc_enc, GLOBAL_AS const u32 *src_buf,
 
       if ((dst_pos + 2) == dst_sz)
       {
+        // this section seems to break intel opencl runtime but is unknown why
+
         dst_ptr[dst_pos++] = (a >> 0) & 0xff;
         dst_ptr[dst_pos++] = (a >> 8) & 0xff;
 

OpenCL/inc_hash_blake2b.cl

@@ -113,7 +113,7 @@ DECLSPEC u64 blake2b_rot24_S (const u64 a)
 
   #else
 
-  return hc_rotr64_S (a, 16);
+  return hc_rotr64_S (a, 24);
 
   #endif
 }

OpenCL/inc_rp.cl

@@ -269,6 +269,35 @@ DECLSPEC int mangle_toggle_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p
   return (len);
 }
 
+DECLSPEC int mangle_toggle_at_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
+{
+  if (len >= RP_PASSWORD_SIZE) return (len);
+
+  u8 occurence = 0;
+
+  u32 rem = 0;
+
+  for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
+  {
+    const u32 t = buf[idx];
+
+    buf[idx] = t | generate_cmask (t);
+
+    u32 out = rem;
+
+    rem = 0;
+
+    if (((t >>  0) & 0xff) == p1) { if (occurence == p0) out = 0x0000ff00; occurence++; }
+    if (((t >>  8) & 0xff) == p1) { if (occurence == p0) out = 0x00ff0000; occurence++; }
+    if (((t >> 16) & 0xff) == p1) { if (occurence == p0) out = 0xff000000; occurence++; }
+    if (((t >> 24) & 0xff) == p1) { if (occurence == p0) rem = 0x000000ff; occurence++; }
+
+    buf[idx] = t ^ (generate_cmask (t) & out);
+  }
+
+  return (len);
+}
+
 DECLSPEC int mangle_reverse (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
 {
   for (int l = 0; l < len / 2; l++)
@@ -725,6 +754,7 @@ DECLSPEC int apply_rule (const u32 name, MAYBE_UNUSED const u8 p0, MAYBE_UNUSED
     case RULE_OP_MANGLE_UREST_LFIRST:     out_len = mangle_urest_lfirst     (p0, p1,        buf, out_len); break;
     case RULE_OP_MANGLE_TREST:            out_len = mangle_trest            (p0, p1,        buf, out_len); break;
     case RULE_OP_MANGLE_TOGGLE_AT:        out_len = mangle_toggle_at        (p0, p1,        buf, out_len); break;
+    case RULE_OP_MANGLE_TOGGLE_AT_SEP:    out_len = mangle_toggle_at_sep    (p0, p1,        buf, out_len); break;
     case RULE_OP_MANGLE_REVERSE:          out_len = mangle_reverse          (p0, p1,        buf, out_len); break;
     case RULE_OP_MANGLE_DUPEWORD:         out_len = mangle_dupeword         (p0, p1,        buf, out_len); break;
     case RULE_OP_MANGLE_DUPEWORD_TIMES:   out_len = mangle_dupeword_times   (p0, p1, (u8 *) buf, out_len); break;

OpenCL/inc_rp.h

@@ -21,6 +21,7 @@
 #define RULE_OP_MANGLE_UREST_LFIRST     'C'
 #define RULE_OP_MANGLE_TREST            't'
 #define RULE_OP_MANGLE_TOGGLE_AT        'T'
+#define RULE_OP_MANGLE_TOGGLE_AT_SEP    '3'
 #define RULE_OP_MANGLE_REVERSE          'r'
 #define RULE_OP_MANGLE_DUPEWORD         'd'
 #define RULE_OP_MANGLE_DUPEWORD_TIMES   'p'
@@ -83,6 +84,7 @@ DECLSPEC int mangle_urest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u
 DECLSPEC int mangle_urest_lfirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 DECLSPEC int mangle_trest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 DECLSPEC int mangle_toggle_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
+DECLSPEC int mangle_toggle_at_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 DECLSPEC int mangle_reverse (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 DECLSPEC int mangle_dupeword (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
 DECLSPEC int mangle_dupeword_times (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);

OpenCL/inc_rp_optimized.cl

@@ -1163,6 +1163,82 @@ DECLSPEC u32 rule_op_mangle_toggle_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED c
   return (in_len);
 }
 
+DECLSPEC u32 rule_op_mangle_toggle_at_sep (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
+{
+  if (in_len == 0) return in_len;
+
+  u32 r0 = search_on_register (buf0[0], p1);
+  u32 r1 = search_on_register (buf0[1], p1);
+  u32 r2 = search_on_register (buf0[2], p1);
+  u32 r3 = search_on_register (buf0[3], p1);
+  u32 r4 = search_on_register (buf1[0], p1);
+  u32 r5 = search_on_register (buf1[1], p1);
+  u32 r6 = search_on_register (buf1[2], p1);
+  u32 r7 = search_on_register (buf1[3], p1);
+
+  const u32 rn = (r0 <<  0)
+               | (r1 <<  4)
+               | (r2 <<  8)
+               | (r3 << 12)
+               | (r4 << 16)
+               | (r5 << 20)
+               | (r6 << 24)
+               | (r7 << 28);
+
+  if (rn == 0) return in_len;
+
+  u32 occurence = 0;
+
+  u32 ro = 0;
+
+  #ifdef _unroll
+  #pragma unroll
+  #endif
+  for (int i = 0; i < 32; i++)
+  {
+    if ((rn >> i) & 1)
+    {
+      if (occurence == p0)
+      {
+        ro = 1 << i;
+
+        break;
+      }
+
+      occurence++;
+    }
+  }
+
+  r0 = (ro >>  0) & 15;
+  r1 = (ro >>  4) & 15;
+  r2 = (ro >>  8) & 15;
+  r3 = (ro >> 12) & 15;
+  r4 = (ro >> 16) & 15;
+  r5 = (ro >> 20) & 15;
+  r6 = (ro >> 24) & 15;
+  r7 = (ro >> 28) & 15;
+
+  r0 <<= 1;
+  r1 <<= 1; r1 |= r0 >> 4;
+  r2 <<= 1; r2 |= r1 >> 4;
+  r3 <<= 1; r3 |= r2 >> 4;
+  r4 <<= 1; r4 |= r3 >> 4;
+  r5 <<= 1; r5 |= r4 >> 4;
+  r6 <<= 1; r6 |= r5 >> 4;
+  r7 <<= 1; r7 |= r6 >> 4;
+
+  buf0[0] = toggle_on_register (buf0[0], r0);
+  buf0[1] = toggle_on_register (buf0[1], r1);
+  buf0[2] = toggle_on_register (buf0[2], r2);
+  buf0[3] = toggle_on_register (buf0[3], r3);
+  buf1[0] = toggle_on_register (buf1[0], r4);
+  buf1[1] = toggle_on_register (buf1[1], r5);
+  buf1[2] = toggle_on_register (buf1[2], r6);
+  buf1[3] = toggle_on_register (buf1[3], r7);
+
+  return in_len;
+}
+
 DECLSPEC u32 rule_op_mangle_reverse (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
 {
   reverse_block_optimized (buf0, buf1, buf0, buf1, in_len);
@@ -2285,6 +2361,7 @@ DECLSPEC u32 apply_rule_optimized (const u32 name, const u32 p0, const u32 p1, u
     case RULE_OP_MANGLE_UREST_LFIRST:     out_len = rule_op_mangle_urest_lfirst     (p0, p1, buf0, buf1, out_len); break;
     case RULE_OP_MANGLE_TREST:            out_len = rule_op_mangle_trest            (p0, p1, buf0, buf1, out_len); break;
     case RULE_OP_MANGLE_TOGGLE_AT:        out_len = rule_op_mangle_toggle_at        (p0, p1, buf0, buf1, out_len); break;
+    case RULE_OP_MANGLE_TOGGLE_AT_SEP:    out_len = rule_op_mangle_toggle_at_sep    (p0, p1, buf0, buf1, out_len); break;
     case RULE_OP_MANGLE_REVERSE:          out_len = rule_op_mangle_reverse          (p0, p1, buf0, buf1, out_len); break;
     case RULE_OP_MANGLE_DUPEWORD:         out_len = rule_op_mangle_dupeword         (p0, p1, buf0, buf1, out_len); break;
     case RULE_OP_MANGLE_DUPEWORD_TIMES:   out_len = rule_op_mangle_dupeword_times   (p0, p1, buf0, buf1, out_len); break;

OpenCL/inc_rp_optimized.h

@@ -21,6 +21,7 @@
 #define RULE_OP_MANGLE_UREST_LFIRST     'C'
 #define RULE_OP_MANGLE_TREST            't'
 #define RULE_OP_MANGLE_TOGGLE_AT        'T'
+#define RULE_OP_MANGLE_TOGGLE_AT_SEP    '3'
 #define RULE_OP_MANGLE_REVERSE          'r'
 #define RULE_OP_MANGLE_DUPEWORD         'd'
 #define RULE_OP_MANGLE_DUPEWORD_TIMES   'p'
@@ -85,6 +86,7 @@ DECLSPEC u32 rule_op_mangle_lrest_ufirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSE
 DECLSPEC u32 rule_op_mangle_urest_lfirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 DECLSPEC u32 rule_op_mangle_trest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 DECLSPEC u32 rule_op_mangle_toggle_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
+DECLSPEC u32 rule_op_mangle_toggle_at_sep (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 DECLSPEC u32 rule_op_mangle_reverse (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 DECLSPEC u32 rule_op_mangle_dupeword (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
 DECLSPEC u32 rule_op_mangle_dupeword_times (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);

OpenCL/inc_vendor.h

@@ -16,7 +16,6 @@
 #define IS_OPENCL
 #endif
 
-
 #if defined IS_NATIVE
 #define CONSTANT_VK
 #define CONSTANT_AS

OpenCL/m12500-pure.cl

@@ -804,7 +804,7 @@ KERNEL_FQ void m12500_loop (KERN_ATTR_TMPS (rar3_tmp_t))
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len & 255;
+  const u32 pw_len = pws[gid].pw_len;
 
   const u32 salt_len = 8;
 
@@ -970,7 +970,7 @@ KERNEL_FQ void m12500_comp (KERN_ATTR_TMPS (rar3_tmp_t))
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len & 255;
+  const u32 pw_len = pws[gid].pw_len;
 
   const u32 salt_len = 8;
 

OpenCL/m13600-pure.cl

@@ -37,7 +37,7 @@ typedef struct zip2
   u32 verify_bytes;
   u32 compress_length;
   u32 data_len;
-  u32 data_buf[0x4000000];
+  u32 data_buf[0x200000];
   u32 auth_len;
   u32 auth_buf[4];
 

OpenCL/m23700-pure.cl

@@ -912,7 +912,7 @@ KERNEL_FQ void m23700_loop (KERN_ATTR_TMPS_ESALT (rar3_tmp_t, rar3_t))
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len & 255;
+  const u32 pw_len = pws[gid].pw_len;
 
   const u32 salt_len = 8;
 
@@ -1086,7 +1086,7 @@ KERNEL_FQ void m23700_comp (KERN_ATTR_TMPS_ESALT (rar3_tmp_t, rar3_t))
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len & 255;
+  const u32 pw_len = pws[gid].pw_len;
 
   const u32 salt_len = 8;
 

OpenCL/m23800-pure.cl

@@ -823,7 +823,7 @@ KERNEL_FQ void m23800_loop (KERN_ATTR_TMPS_HOOKS_ESALT (rar3_tmp_t, rar3_hook_t,
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len & 255;
+  const u32 pw_len = pws[gid].pw_len;
 
   const u32 salt_len = 8;
 
@@ -990,7 +990,7 @@ KERNEL_FQ void m23800_hook23 (KERN_ATTR_TMPS_HOOKS_ESALT (rar3_tmp_t, rar3_hook_
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len & 255;
+  const u32 pw_len = pws[gid].pw_len;
 
   const u32 salt_len = 8;
 

OpenCL/m25700_a3-optimized.cl

@@ -14,7 +14,7 @@
 #include "inc_hash_md5.cl"
 #endif
 
-DECLSPEC u32x MurmurHash_w0 (const u32 seed, const u32x w0, const u32 *w, const int pw_len)
+DECLSPEC u32x MurmurHash_w0 (const u32 seed, const u32x w0, const u32x *w, const int pw_len)
 {
   u32x hash = seed;
 

OpenCL/m27000-pure.cl

@@ -0,0 +1,697 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+//#define NEW_SIMD_CODE
+
+#ifdef KERNEL_STATIC
+#include "inc_vendor.h"
+#include "inc_types.h"
+#include "inc_platform.cl"
+#include "inc_common.cl"
+#include "inc_rp.h"
+#include "inc_rp.cl"
+#include "inc_scalar.cl"
+#include "inc_hash_md4.cl"
+#endif
+
+#define COMPARE_S "inc_comp_single.cl"
+#define COMPARE_M "inc_comp_multi.cl"
+
+#define PERM_OP(a,b,tt,n,m) \
+{                           \
+  tt = a >> n;              \
+  tt = tt ^ b;              \
+  tt = tt & m;              \
+  b = b ^ tt;               \
+  tt = tt << n;             \
+  a = a ^ tt;               \
+}
+
+#define HPERM_OP(a,tt,n,m)  \
+{                           \
+  tt = a << (16 + n);       \
+  tt = tt ^ a;              \
+  tt = tt & m;              \
+  a  = a ^ tt;              \
+  tt = tt >> (16 + n);      \
+  a  = a ^ tt;              \
+}
+
+CONSTANT_VK u32a c_SPtrans[8][64] =
+{
+  {
+    0x02080800, 0x00080000, 0x02000002, 0x02080802,
+    0x02000000, 0x00080802, 0x00080002, 0x02000002,
+    0x00080802, 0x02080800, 0x02080000, 0x00000802,
+    0x02000802, 0x02000000, 0x00000000, 0x00080002,
+    0x00080000, 0x00000002, 0x02000800, 0x00080800,
+    0x02080802, 0x02080000, 0x00000802, 0x02000800,
+    0x00000002, 0x00000800, 0x00080800, 0x02080002,
+    0x00000800, 0x02000802, 0x02080002, 0x00000000,
+    0x00000000, 0x02080802, 0x02000800, 0x00080002,
+    0x02080800, 0x00080000, 0x00000802, 0x02000800,
+    0x02080002, 0x00000800, 0x00080800, 0x02000002,
+    0x00080802, 0x00000002, 0x02000002, 0x02080000,
+    0x02080802, 0x00080800, 0x02080000, 0x02000802,
+    0x02000000, 0x00000802, 0x00080002, 0x00000000,
+    0x00080000, 0x02000000, 0x02000802, 0x02080800,
+    0x00000002, 0x02080002, 0x00000800, 0x00080802,
+  },
+  {
+    0x40108010, 0x00000000, 0x00108000, 0x40100000,
+    0x40000010, 0x00008010, 0x40008000, 0x00108000,
+    0x00008000, 0x40100010, 0x00000010, 0x40008000,
+    0x00100010, 0x40108000, 0x40100000, 0x00000010,
+    0x00100000, 0x40008010, 0x40100010, 0x00008000,
+    0x00108010, 0x40000000, 0x00000000, 0x00100010,
+    0x40008010, 0x00108010, 0x40108000, 0x40000010,
+    0x40000000, 0x00100000, 0x00008010, 0x40108010,
+    0x00100010, 0x40108000, 0x40008000, 0x00108010,
+    0x40108010, 0x00100010, 0x40000010, 0x00000000,
+    0x40000000, 0x00008010, 0x00100000, 0x40100010,
+    0x00008000, 0x40000000, 0x00108010, 0x40008010,
+    0x40108000, 0x00008000, 0x00000000, 0x40000010,
+    0x00000010, 0x40108010, 0x00108000, 0x40100000,
+    0x40100010, 0x00100000, 0x00008010, 0x40008000,
+    0x40008010, 0x00000010, 0x40100000, 0x00108000,
+  },
+  {
+    0x04000001, 0x04040100, 0x00000100, 0x04000101,
+    0x00040001, 0x04000000, 0x04000101, 0x00040100,
+    0x04000100, 0x00040000, 0x04040000, 0x00000001,
+    0x04040101, 0x00000101, 0x00000001, 0x04040001,
+    0x00000000, 0x00040001, 0x04040100, 0x00000100,
+    0x00000101, 0x04040101, 0x00040000, 0x04000001,
+    0x04040001, 0x04000100, 0x00040101, 0x04040000,
+    0x00040100, 0x00000000, 0x04000000, 0x00040101,
+    0x04040100, 0x00000100, 0x00000001, 0x00040000,
+    0x00000101, 0x00040001, 0x04040000, 0x04000101,
+    0x00000000, 0x04040100, 0x00040100, 0x04040001,
+    0x00040001, 0x04000000, 0x04040101, 0x00000001,
+    0x00040101, 0x04000001, 0x04000000, 0x04040101,
+    0x00040000, 0x04000100, 0x04000101, 0x00040100,
+    0x04000100, 0x00000000, 0x04040001, 0x00000101,
+    0x04000001, 0x00040101, 0x00000100, 0x04040000,
+  },
+  {
+    0x00401008, 0x10001000, 0x00000008, 0x10401008,
+    0x00000000, 0x10400000, 0x10001008, 0x00400008,
+    0x10401000, 0x10000008, 0x10000000, 0x00001008,
+    0x10000008, 0x00401008, 0x00400000, 0x10000000,
+    0x10400008, 0x00401000, 0x00001000, 0x00000008,
+    0x00401000, 0x10001008, 0x10400000, 0x00001000,
+    0x00001008, 0x00000000, 0x00400008, 0x10401000,
+    0x10001000, 0x10400008, 0x10401008, 0x00400000,
+    0x10400008, 0x00001008, 0x00400000, 0x10000008,
+    0x00401000, 0x10001000, 0x00000008, 0x10400000,
+    0x10001008, 0x00000000, 0x00001000, 0x00400008,
+    0x00000000, 0x10400008, 0x10401000, 0x00001000,
+    0x10000000, 0x10401008, 0x00401008, 0x00400000,
+    0x10401008, 0x00000008, 0x10001000, 0x00401008,
+    0x00400008, 0x00401000, 0x10400000, 0x10001008,
+    0x00001008, 0x10000000, 0x10000008, 0x10401000,
+  },
+  {
+    0x08000000, 0x00010000, 0x00000400, 0x08010420,
+    0x08010020, 0x08000400, 0x00010420, 0x08010000,
+    0x00010000, 0x00000020, 0x08000020, 0x00010400,
+    0x08000420, 0x08010020, 0x08010400, 0x00000000,
+    0x00010400, 0x08000000, 0x00010020, 0x00000420,
+    0x08000400, 0x00010420, 0x00000000, 0x08000020,
+    0x00000020, 0x08000420, 0x08010420, 0x00010020,
+    0x08010000, 0x00000400, 0x00000420, 0x08010400,
+    0x08010400, 0x08000420, 0x00010020, 0x08010000,
+    0x00010000, 0x00000020, 0x08000020, 0x08000400,
+    0x08000000, 0x00010400, 0x08010420, 0x00000000,
+    0x00010420, 0x08000000, 0x00000400, 0x00010020,
+    0x08000420, 0x00000400, 0x00000000, 0x08010420,
+    0x08010020, 0x08010400, 0x00000420, 0x00010000,
+    0x00010400, 0x08010020, 0x08000400, 0x00000420,
+    0x00000020, 0x00010420, 0x08010000, 0x08000020,
+  },
+  {
+    0x80000040, 0x00200040, 0x00000000, 0x80202000,
+    0x00200040, 0x00002000, 0x80002040, 0x00200000,
+    0x00002040, 0x80202040, 0x00202000, 0x80000000,
+    0x80002000, 0x80000040, 0x80200000, 0x00202040,
+    0x00200000, 0x80002040, 0x80200040, 0x00000000,
+    0x00002000, 0x00000040, 0x80202000, 0x80200040,
+    0x80202040, 0x80200000, 0x80000000, 0x00002040,
+    0x00000040, 0x00202000, 0x00202040, 0x80002000,
+    0x00002040, 0x80000000, 0x80002000, 0x00202040,
+    0x80202000, 0x00200040, 0x00000000, 0x80002000,
+    0x80000000, 0x00002000, 0x80200040, 0x00200000,
+    0x00200040, 0x80202040, 0x00202000, 0x00000040,
+    0x80202040, 0x00202000, 0x00200000, 0x80002040,
+    0x80000040, 0x80200000, 0x00202040, 0x00000000,
+    0x00002000, 0x80000040, 0x80002040, 0x80202000,
+    0x80200000, 0x00002040, 0x00000040, 0x80200040,
+  },
+  {
+    0x00004000, 0x00000200, 0x01000200, 0x01000004,
+    0x01004204, 0x00004004, 0x00004200, 0x00000000,
+    0x01000000, 0x01000204, 0x00000204, 0x01004000,
+    0x00000004, 0x01004200, 0x01004000, 0x00000204,
+    0x01000204, 0x00004000, 0x00004004, 0x01004204,
+    0x00000000, 0x01000200, 0x01000004, 0x00004200,
+    0x01004004, 0x00004204, 0x01004200, 0x00000004,
+    0x00004204, 0x01004004, 0x00000200, 0x01000000,
+    0x00004204, 0x01004000, 0x01004004, 0x00000204,
+    0x00004000, 0x00000200, 0x01000000, 0x01004004,
+    0x01000204, 0x00004204, 0x00004200, 0x00000000,
+    0x00000200, 0x01000004, 0x00000004, 0x01000200,
+    0x00000000, 0x01000204, 0x01000200, 0x00004200,
+    0x00000204, 0x00004000, 0x01004204, 0x01000000,
+    0x01004200, 0x00000004, 0x00004004, 0x01004204,
+    0x01000004, 0x01004200, 0x01004000, 0x00004004,
+  },
+  {
+    0x20800080, 0x20820000, 0x00020080, 0x00000000,
+    0x20020000, 0x00800080, 0x20800000, 0x20820080,
+    0x00000080, 0x20000000, 0x00820000, 0x00020080,
+    0x00820080, 0x20020080, 0x20000080, 0x20800000,
+    0x00020000, 0x00820080, 0x00800080, 0x20020000,
+    0x20820080, 0x20000080, 0x00000000, 0x00820000,
+    0x20000000, 0x00800000, 0x20020080, 0x20800080,
+    0x00800000, 0x00020000, 0x20820000, 0x00000080,
+    0x00800000, 0x00020000, 0x20000080, 0x20820080,
+    0x00020080, 0x20000000, 0x00000000, 0x00820000,
+    0x20800080, 0x20020080, 0x20020000, 0x00800080,
+    0x20820000, 0x00000080, 0x00800080, 0x20020000,
+    0x20820080, 0x00800000, 0x20800000, 0x20000080,
+    0x00820000, 0x00020080, 0x20020080, 0x20800000,
+    0x00000080, 0x20820000, 0x00820080, 0x00000000,
+    0x20000000, 0x20800080, 0x00020000, 0x00820080,
+  }
+};
+
+CONSTANT_VK u32a c_skb[8][64] =
+{
+  {
+    0x00000000, 0x00000010, 0x20000000, 0x20000010,
+    0x00010000, 0x00010010, 0x20010000, 0x20010010,
+    0x00000800, 0x00000810, 0x20000800, 0x20000810,
+    0x00010800, 0x00010810, 0x20010800, 0x20010810,
+    0x00000020, 0x00000030, 0x20000020, 0x20000030,
+    0x00010020, 0x00010030, 0x20010020, 0x20010030,
+    0x00000820, 0x00000830, 0x20000820, 0x20000830,
+    0x00010820, 0x00010830, 0x20010820, 0x20010830,
+    0x00080000, 0x00080010, 0x20080000, 0x20080010,
+    0x00090000, 0x00090010, 0x20090000, 0x20090010,
+    0x00080800, 0x00080810, 0x20080800, 0x20080810,
+    0x00090800, 0x00090810, 0x20090800, 0x20090810,
+    0x00080020, 0x00080030, 0x20080020, 0x20080030,
+    0x00090020, 0x00090030, 0x20090020, 0x20090030,
+    0x00080820, 0x00080830, 0x20080820, 0x20080830,
+    0x00090820, 0x00090830, 0x20090820, 0x20090830,
+  },
+  {
+    0x00000000, 0x02000000, 0x00002000, 0x02002000,
+    0x00200000, 0x02200000, 0x00202000, 0x02202000,
+    0x00000004, 0x02000004, 0x00002004, 0x02002004,
+    0x00200004, 0x02200004, 0x00202004, 0x02202004,
+    0x00000400, 0x02000400, 0x00002400, 0x02002400,
+    0x00200400, 0x02200400, 0x00202400, 0x02202400,
+    0x00000404, 0x02000404, 0x00002404, 0x02002404,
+    0x00200404, 0x02200404, 0x00202404, 0x02202404,
+    0x10000000, 0x12000000, 0x10002000, 0x12002000,
+    0x10200000, 0x12200000, 0x10202000, 0x12202000,
+    0x10000004, 0x12000004, 0x10002004, 0x12002004,
+    0x10200004, 0x12200004, 0x10202004, 0x12202004,
+    0x10000400, 0x12000400, 0x10002400, 0x12002400,
+    0x10200400, 0x12200400, 0x10202400, 0x12202400,
+    0x10000404, 0x12000404, 0x10002404, 0x12002404,
+    0x10200404, 0x12200404, 0x10202404, 0x12202404,
+  },
+  {
+    0x00000000, 0x00000001, 0x00040000, 0x00040001,
+    0x01000000, 0x01000001, 0x01040000, 0x01040001,
+    0x00000002, 0x00000003, 0x00040002, 0x00040003,
+    0x01000002, 0x01000003, 0x01040002, 0x01040003,
+    0x00000200, 0x00000201, 0x00040200, 0x00040201,
+    0x01000200, 0x01000201, 0x01040200, 0x01040201,
+    0x00000202, 0x00000203, 0x00040202, 0x00040203,
+    0x01000202, 0x01000203, 0x01040202, 0x01040203,
+    0x08000000, 0x08000001, 0x08040000, 0x08040001,
+    0x09000000, 0x09000001, 0x09040000, 0x09040001,
+    0x08000002, 0x08000003, 0x08040002, 0x08040003,
+    0x09000002, 0x09000003, 0x09040002, 0x09040003,
+    0x08000200, 0x08000201, 0x08040200, 0x08040201,
+    0x09000200, 0x09000201, 0x09040200, 0x09040201,
+    0x08000202, 0x08000203, 0x08040202, 0x08040203,
+    0x09000202, 0x09000203, 0x09040202, 0x09040203,
+  },
+  {
+    0x00000000, 0x00100000, 0x00000100, 0x00100100,
+    0x00000008, 0x00100008, 0x00000108, 0x00100108,
+    0x00001000, 0x00101000, 0x00001100, 0x00101100,
+    0x00001008, 0x00101008, 0x00001108, 0x00101108,
+    0x04000000, 0x04100000, 0x04000100, 0x04100100,
+    0x04000008, 0x04100008, 0x04000108, 0x04100108,
+    0x04001000, 0x04101000, 0x04001100, 0x04101100,
+    0x04001008, 0x04101008, 0x04001108, 0x04101108,
+    0x00020000, 0x00120000, 0x00020100, 0x00120100,
+    0x00020008, 0x00120008, 0x00020108, 0x00120108,
+    0x00021000, 0x00121000, 0x00021100, 0x00121100,
+    0x00021008, 0x00121008, 0x00021108, 0x00121108,
+    0x04020000, 0x04120000, 0x04020100, 0x04120100,
+    0x04020008, 0x04120008, 0x04020108, 0x04120108,
+    0x04021000, 0x04121000, 0x04021100, 0x04121100,
+    0x04021008, 0x04121008, 0x04021108, 0x04121108,
+  },
+  {
+    0x00000000, 0x10000000, 0x00010000, 0x10010000,
+    0x00000004, 0x10000004, 0x00010004, 0x10010004,
+    0x20000000, 0x30000000, 0x20010000, 0x30010000,
+    0x20000004, 0x30000004, 0x20010004, 0x30010004,
+    0x00100000, 0x10100000, 0x00110000, 0x10110000,
+    0x00100004, 0x10100004, 0x00110004, 0x10110004,
+    0x20100000, 0x30100000, 0x20110000, 0x30110000,
+    0x20100004, 0x30100004, 0x20110004, 0x30110004,
+    0x00001000, 0x10001000, 0x00011000, 0x10011000,
+    0x00001004, 0x10001004, 0x00011004, 0x10011004,
+    0x20001000, 0x30001000, 0x20011000, 0x30011000,
+    0x20001004, 0x30001004, 0x20011004, 0x30011004,
+    0x00101000, 0x10101000, 0x00111000, 0x10111000,
+    0x00101004, 0x10101004, 0x00111004, 0x10111004,
+    0x20101000, 0x30101000, 0x20111000, 0x30111000,
+    0x20101004, 0x30101004, 0x20111004, 0x30111004,
+  },
+  {
+    0x00000000, 0x08000000, 0x00000008, 0x08000008,
+    0x00000400, 0x08000400, 0x00000408, 0x08000408,
+    0x00020000, 0x08020000, 0x00020008, 0x08020008,
+    0x00020400, 0x08020400, 0x00020408, 0x08020408,
+    0x00000001, 0x08000001, 0x00000009, 0x08000009,
+    0x00000401, 0x08000401, 0x00000409, 0x08000409,
+    0x00020001, 0x08020001, 0x00020009, 0x08020009,
+    0x00020401, 0x08020401, 0x00020409, 0x08020409,
+    0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
+    0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
+    0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
+    0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
+    0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
+    0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
+    0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
+    0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
+  },
+  {
+    0x00000000, 0x00000100, 0x00080000, 0x00080100,
+    0x01000000, 0x01000100, 0x01080000, 0x01080100,
+    0x00000010, 0x00000110, 0x00080010, 0x00080110,
+    0x01000010, 0x01000110, 0x01080010, 0x01080110,
+    0x00200000, 0x00200100, 0x00280000, 0x00280100,
+    0x01200000, 0x01200100, 0x01280000, 0x01280100,
+    0x00200010, 0x00200110, 0x00280010, 0x00280110,
+    0x01200010, 0x01200110, 0x01280010, 0x01280110,
+    0x00000200, 0x00000300, 0x00080200, 0x00080300,
+    0x01000200, 0x01000300, 0x01080200, 0x01080300,
+    0x00000210, 0x00000310, 0x00080210, 0x00080310,
+    0x01000210, 0x01000310, 0x01080210, 0x01080310,
+    0x00200200, 0x00200300, 0x00280200, 0x00280300,
+    0x01200200, 0x01200300, 0x01280200, 0x01280300,
+    0x00200210, 0x00200310, 0x00280210, 0x00280310,
+    0x01200210, 0x01200310, 0x01280210, 0x01280310,
+  },
+  {
+    0x00000000, 0x04000000, 0x00040000, 0x04040000,
+    0x00000002, 0x04000002, 0x00040002, 0x04040002,
+    0x00002000, 0x04002000, 0x00042000, 0x04042000,
+    0x00002002, 0x04002002, 0x00042002, 0x04042002,
+    0x00000020, 0x04000020, 0x00040020, 0x04040020,
+    0x00000022, 0x04000022, 0x00040022, 0x04040022,
+    0x00002020, 0x04002020, 0x00042020, 0x04042020,
+    0x00002022, 0x04002022, 0x00042022, 0x04042022,
+    0x00000800, 0x04000800, 0x00040800, 0x04040800,
+    0x00000802, 0x04000802, 0x00040802, 0x04040802,
+    0x00002800, 0x04002800, 0x00042800, 0x04042800,
+    0x00002802, 0x04002802, 0x00042802, 0x04042802,
+    0x00000820, 0x04000820, 0x00040820, 0x04040820,
+    0x00000822, 0x04000822, 0x00040822, 0x04040822,
+    0x00002820, 0x04002820, 0x00042820, 0x04042820,
+    0x00002822, 0x04002822, 0x00042822, 0x04042822
+  }
+};
+
+#if   VECT_SIZE == 1
+#define BOX(i,n,S) (S)[(n)][(i)]
+#elif VECT_SIZE == 2
+#define BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
+#elif VECT_SIZE == 4
+#define BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
+#elif VECT_SIZE == 8
+#define BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
+#elif VECT_SIZE == 16
+#define BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
+#endif
+
+DECLSPEC void _des_crypt_encrypt (u32 *iv, u32 *data, u32 *Kc, u32 *Kd, SHM_TYPE u32 (*s_SPtrans)[64])
+{
+  u32 r = data[0];
+  u32 l = data[1];
+
+  #ifdef _unroll
+  #pragma unroll
+  #endif
+  for (u32 i = 0; i < 16; i += 2)
+  {
+    u32 u;
+    u32 t;
+
+    u = Kc[i + 0] ^ hc_rotl32 (r, 30u);
+    t = Kd[i + 0] ^ hc_rotl32 (r, 26u);
+
+    l ^= BOX (((u >>  0) & 0x3f), 0, s_SPtrans)
+       | BOX (((u >>  8) & 0x3f), 2, s_SPtrans)
+       | BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
+       | BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
+       | BOX (((t >>  0) & 0x3f), 1, s_SPtrans)
+       | BOX (((t >>  8) & 0x3f), 3, s_SPtrans)
+       | BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
+       | BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
+
+    u = Kc[i + 1] ^ hc_rotl32 (l, 30u);
+    t = Kd[i + 1] ^ hc_rotl32 (l, 26u);
+
+    r ^= BOX (((u >>  0) & 0x3f), 0, s_SPtrans)
+       | BOX (((u >>  8) & 0x3f), 2, s_SPtrans)
+       | BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
+       | BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
+       | BOX (((t >>  0) & 0x3f), 1, s_SPtrans)
+       | BOX (((t >>  8) & 0x3f), 3, s_SPtrans)
+       | BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
+       | BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
+  }
+
+  iv[0] = l;
+  iv[1] = r;
+}
+
+DECLSPEC void _des_crypt_keysetup (u32 c, u32 d, u32 *Kc, u32 *Kd, SHM_TYPE u32 (*s_skb)[64])
+{
+  u32 tt;
+
+  PERM_OP  (d, c, tt, 4, 0x0f0f0f0f);
+  HPERM_OP (c,    tt, 2, 0xcccc0000);
+  HPERM_OP (d,    tt, 2, 0xcccc0000);
+  PERM_OP  (d, c, tt, 1, 0x55555555);
+  PERM_OP  (c, d, tt, 8, 0x00ff00ff);
+  PERM_OP  (d, c, tt, 1, 0x55555555);
+
+  d = ((d & 0x000000ff) << 16)
+    | ((d & 0x0000ff00) <<  0)
+    | ((d & 0x00ff0000) >> 16)
+    | ((c & 0xf0000000) >>  4);
+
+  c = c & 0x0fffffff;
+
+  #ifdef _unroll
+  #pragma unroll
+  #endif
+  for (u32 i = 0; i < 16; i++)
+  {
+    if ((i < 2) || (i == 8) || (i == 15))
+    {
+      c = ((c >> 1) | (c << 27));
+      d = ((d >> 1) | (d << 27));
+    }
+    else
+    {
+      c = ((c >> 2) | (c << 26));
+      d = ((d >> 2) | (d << 26));
+    }
+
+    c = c & 0x0fffffff;
+    d = d & 0x0fffffff;
+
+    const u32 c00 = (c >>  0) & 0x0000003f;
+    const u32 c06 = (c >>  6) & 0x00383003;
+    const u32 c07 = (c >>  7) & 0x0000003c;
+    const u32 c13 = (c >> 13) & 0x0000060f;
+    const u32 c20 = (c >> 20) & 0x00000001;
+
+    u32 s = BOX (((c00 >>  0) & 0xff), 0, s_skb)
+           | BOX (((c06 >>  0) & 0xff)
+                 |((c07 >>  0) & 0xff), 1, s_skb)
+           | BOX (((c13 >>  0) & 0xff)
+                 |((c06 >>  8) & 0xff), 2, s_skb)
+           | BOX (((c20 >>  0) & 0xff)
+                 |((c13 >>  8) & 0xff)
+                 |((c06 >> 16) & 0xff), 3, s_skb);
+
+    const u32 d00 = (d >>  0) & 0x00003c3f;
+    const u32 d07 = (d >>  7) & 0x00003f03;
+    const u32 d21 = (d >> 21) & 0x0000000f;
+    const u32 d22 = (d >> 22) & 0x00000030;
+
+    u32 t = BOX (((d00 >>  0) & 0xff), 4, s_skb)
+           | BOX (((d07 >>  0) & 0xff)
+                 |((d00 >>  8) & 0xff), 5, s_skb)
+           | BOX (((d07 >>  8) & 0xff), 6, s_skb)
+           | BOX (((d21 >>  0) & 0xff)
+                 |((d22 >>  0) & 0xff), 7, s_skb);
+
+    Kc[i] = ((t << 16) | (s & 0x0000ffff));
+    Kd[i] = ((s >> 16) | (t & 0xffff0000));
+  }
+}
+
+DECLSPEC void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 *out)
+{
+  u32 t[8];
+
+  t[0] = (w0 >>  0) & 0xff;
+  t[1] = (w0 >>  8) & 0xff;
+  t[2] = (w0 >> 16) & 0xff;
+  t[3] = (w0 >> 24) & 0xff;
+  t[4] = (w1 >>  0) & 0xff;
+  t[5] = (w1 >>  8) & 0xff;
+  t[6] = (w1 >> 16) & 0xff;
+  t[7] = (w1 >> 24) & 0xff;
+
+  u32 k[8];
+
+  k[0] =               (t[0] >> 0);
+  k[1] = (t[0] << 7) | (t[1] >> 1);
+  k[2] = (t[1] << 6) | (t[2] >> 2);
+  k[3] = (t[2] << 5) | (t[3] >> 3);
+  k[4] = (t[3] << 4) | (t[4] >> 4);
+  k[5] = (t[4] << 3) | (t[5] >> 5);
+  k[6] = (t[5] << 2) | (t[6] >> 6);
+  k[7] = (t[6] << 1);
+
+  out[0] = ((k[0] & 0xff) <<  0)
+         | ((k[1] & 0xff) <<  8)
+         | ((k[2] & 0xff) << 16)
+         | ((k[3] & 0xff) << 24);
+
+  out[1] = ((k[4] & 0xff) <<  0)
+         | ((k[5] & 0xff) <<  8)
+         | ((k[6] & 0xff) << 16)
+         | ((k[7] & 0xff) << 24);
+}
+
+#ifdef KERNEL_STATIC
+DECLSPEC u8 hex_convert (const u8 c)
+{
+  return (c & 15) + (c >> 6) * 9;
+}
+
+DECLSPEC u8 hex_to_u8 (const u8 *hex)
+{
+  u8 v = 0;
+
+  v |= ((u8) hex_convert (hex[1]) << 0);
+  v |= ((u8) hex_convert (hex[0]) << 4);
+
+  return (v);
+}
+#endif
+
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
+typedef struct netntlm_tmp
+{
+  u32 digest_buf[4];
+
+} netntlm_tmp_t;
+
+KERNEL_FQ void m27000_init (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
+{
+  /**
+   * modifier
+   */
+
+  const u64 gid = get_global_id (0);
+  const u64 lid = get_local_id (0);
+  const u64 lsz = get_local_size (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * salt
+   */
+
+  u32 in[16];
+
+  in[ 0] = pws[gid].i[ 0];
+  in[ 1] = pws[gid].i[ 1];
+  in[ 2] = pws[gid].i[ 2];
+  in[ 3] = pws[gid].i[ 3];
+  in[ 4] = pws[gid].i[ 4];
+  in[ 5] = pws[gid].i[ 5];
+  in[ 6] = pws[gid].i[ 6];
+  in[ 7] = pws[gid].i[ 7];
+
+  u8 *in_ptr = (u8 *) in;
+
+  u32 out[4];
+
+  u8 *out_ptr = (u8 *) out;
+
+  for (int i = 0, j = 0; i < 16; i += 1, j += 2)
+  {
+    out_ptr[i] = hex_to_u8 (in_ptr + j);
+  }
+
+  tmps[gid].digest_buf[0] = out[ 0];
+  tmps[gid].digest_buf[1] = out[ 1];
+  tmps[gid].digest_buf[2] = out[ 2];
+  tmps[gid].digest_buf[3] = out[ 3];
+
+}
+
+KERNEL_FQ void m27000_loop (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
+{
+
+}
+
+KERNEL_FQ void m27000_comp (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
+{
+   /**
+   * modifier
+   */
+
+  const u64 gid = get_global_id (0);
+  const u64 lid = get_local_id (0);
+  const u64 lsz = get_local_size (0);
+
+  /**
+   * sbox, kbox
+   */
+
+  #ifdef REAL_SHM
+
+  LOCAL_VK u32 s_SPtrans[8][64];
+  LOCAL_VK u32 s_skb[8][64];
+
+  for (u32 i = lid; i < 64; i += lsz)
+  {
+    s_SPtrans[0][i] = c_SPtrans[0][i];
+    s_SPtrans[1][i] = c_SPtrans[1][i];
+    s_SPtrans[2][i] = c_SPtrans[2][i];
+    s_SPtrans[3][i] = c_SPtrans[3][i];
+    s_SPtrans[4][i] = c_SPtrans[4][i];
+    s_SPtrans[5][i] = c_SPtrans[5][i];
+    s_SPtrans[6][i] = c_SPtrans[6][i];
+    s_SPtrans[7][i] = c_SPtrans[7][i];
+
+    s_skb[0][i] = c_skb[0][i];
+    s_skb[1][i] = c_skb[1][i];
+    s_skb[2][i] = c_skb[2][i];
+    s_skb[3][i] = c_skb[3][i];
+    s_skb[4][i] = c_skb[4][i];
+    s_skb[5][i] = c_skb[5][i];
+    s_skb[6][i] = c_skb[6][i];
+    s_skb[7][i] = c_skb[7][i];
+  }
+
+  SYNC_THREADS ();
+
+  #else
+
+  CONSTANT_AS u32a (*s_SPtrans)[64] = c_SPtrans;
+  CONSTANT_AS u32a (*s_skb)[64]     = c_skb;
+
+  #endif
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  const u32 s0 = salt_bufs[SALT_POS].salt_buf[0];
+  const u32 s1 = salt_bufs[SALT_POS].salt_buf[1];
+  const u32 s2 = salt_bufs[SALT_POS].salt_buf[2];
+
+  const u32 a = tmps[gid].digest_buf[0];
+  const u32 b = tmps[gid].digest_buf[1];
+  const u32 c = tmps[gid].digest_buf[2];
+  const u32 d = tmps[gid].digest_buf[3];
+
+  // I believe this matches the last 2 bytes and throws away.
+  // Taken from 5500. 
+  if ((d >> 16) != s2) return;
+
+  /**
+    * DES1
+    */
+
+  u32 key[2];
+
+  transform_netntlmv1_key (a, b, key);
+
+  u32 Kc[16];
+  u32 Kd[16];
+
+  _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
+
+  u32 data[2];
+
+  data[0] = s0;
+  data[1] = s1;
+
+  u32 out1[2];
+
+  _des_crypt_encrypt (out1, data, Kc, Kd, s_SPtrans);
+
+  /**
+    * DES2
+    */
+
+  transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
+
+  _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
+
+  u32 out2[2];
+
+  _des_crypt_encrypt (out2, data, Kc, Kd, s_SPtrans);
+
+  /**
+   * digest
+   */
+
+  const u32 r0 = out1[0];
+  const u32 r1 = out1[1];
+  const u32 r2 = out2[0];
+  const u32 r3 = out2[1];
+
+  #define il_pos 0
+
+  #ifdef KERNEL_STATIC
+  #include COMPARE_M
+  #endif
+}

OpenCL/m27100-pure.cl

@@ -0,0 +1,197 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+// #define NEW_SIMD_CODE
+
+#ifdef KERNEL_STATIC
+#include "inc_vendor.h"
+#include "inc_types.h"
+#include "inc_platform.cl"
+#include "inc_common.cl"
+#include "inc_rp.h"
+#include "inc_rp.cl"
+#include "inc_scalar.cl"
+#include "inc_hash_md4.cl"
+#include "inc_hash_md5.cl"
+#endif
+
+#define COMPARE_S "inc_comp_single.cl"
+#define COMPARE_M "inc_comp_multi.cl"
+
+#ifdef KERNEL_STATIC
+DECLSPEC u8 hex_convert (const u8 c)
+{
+  return (c & 15) + (c >> 6) * 9;
+}
+
+DECLSPEC u8 hex_to_u8 (const u8 *hex)
+{
+  u8 v = 0;
+
+  v |= ((u8) hex_convert (hex[1]) << 0);
+  v |= ((u8) hex_convert (hex[0]) << 4);
+
+  return (v);
+}
+#endif
+
+typedef struct netntlm
+{
+  u32 user_len;
+  u32 domain_len;
+  u32 srvchall_len;
+  u32 clichall_len;
+
+  u32 userdomain_buf[64];
+  u32 chall_buf[256];
+
+} netntlm_t;
+
+typedef struct netntlmv2_tmp
+{
+  u32 digest_buf[4];
+
+} netntlm_tmp_t;
+
+
+KERNEL_FQ void m27100_init (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  u32 in[16];
+
+  in[ 0] = pws[gid].i[ 0];
+  in[ 1] = pws[gid].i[ 1];
+  in[ 2] = pws[gid].i[ 2];
+  in[ 3] = pws[gid].i[ 3];
+  in[ 4] = pws[gid].i[ 4];
+  in[ 5] = pws[gid].i[ 5];
+  in[ 6] = pws[gid].i[ 6];
+  in[ 7] = pws[gid].i[ 7];
+
+  u8 *in_ptr = (u8 *) in;
+
+  u32 out[4];
+
+  u8 *out_ptr = (u8 *) out;
+
+  for (int i = 0, j = 0; i < 16; i += 1, j += 2)
+  {
+    out_ptr[i] = hex_to_u8 (in_ptr + j);
+  }
+
+  tmps[gid].digest_buf[0] = out[ 0];
+  tmps[gid].digest_buf[1] = out[ 1];
+  tmps[gid].digest_buf[2] = out[ 2];
+  tmps[gid].digest_buf[3] = out[ 3];
+
+}
+
+
+KERNEL_FQ void m27100_loop (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
+{
+
+
+}
+
+KERNEL_FQ void m27100_comp (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
+{
+   /**
+   * modifier
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u64 lid = get_local_id (0);
+
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+
+  w0[0] = tmps[gid].digest_buf[0];
+  w0[1] = tmps[gid].digest_buf[1];
+  w0[2] = tmps[gid].digest_buf[2];
+  w0[3] = tmps[gid].digest_buf[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;
+
+  md5_hmac_ctx_t ctx0;
+
+  md5_hmac_init_64 (&ctx0, w0, w1, w2, w3);
+
+  md5_hmac_update_global (&ctx0, esalt_bufs[DIGESTS_OFFSET].userdomain_buf, esalt_bufs[DIGESTS_OFFSET].user_len + esalt_bufs[DIGESTS_OFFSET].domain_len);
+
+  md5_hmac_final (&ctx0);
+
+  w0[0] = ctx0.opad.h[0];
+  w0[1] = ctx0.opad.h[1];
+  w0[2] = ctx0.opad.h[2];
+  w0[3] = ctx0.opad.h[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;
+
+  md5_hmac_ctx_t ctx;
+
+  md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
+
+  md5_hmac_update_global (&ctx, esalt_bufs[DIGESTS_OFFSET].chall_buf, esalt_bufs[DIGESTS_OFFSET].srvchall_len + esalt_bufs[DIGESTS_OFFSET].clichall_len);
+
+  md5_hmac_final (&ctx);
+
+  tmps[gid].digest_buf[0] = ctx.opad.h[0];
+  tmps[gid].digest_buf[1] = ctx.opad.h[1];
+  tmps[gid].digest_buf[2] = ctx.opad.h[2];
+  tmps[gid].digest_buf[3] = ctx.opad.h[3];
+
+
+  /**
+   * digest
+   */
+
+  const u32 r0 = ctx.opad.h[DGST_R0];
+  const u32 r1 = ctx.opad.h[DGST_R1];
+  const u32 r2 = ctx.opad.h[DGST_R2];
+  const u32 r3 = ctx.opad.h[DGST_R3];
+
+  #define il_pos 0
+
+  #ifdef KERNEL_STATIC
+  #include COMPARE_M
+  #endif
+}

docs/changes.txt

@@ -11,9 +11,13 @@
 ## Bugs
 ##
 
+- Fixed buffer overflow in Stargazer Stellar Wallet XLM module in hash_encode() if a hash was cracked
 - Fixed autotune unitialized tmps variable for slow hashes by calling _init kernel before calling _loop kernel
 - Fixed datatype in function sha384_hmac_init_vector_128() that could come into effect if vector datatype was manually set
 - Fixed false negative in all VeraCrypt hash-modes if both conditions are met: 1. use CPU for cracking and 2. PIM range was used
+- Fixed multiple buffer overflow in DPAPI masterkey file v1 and v2 module
+- Fixed out-of-boundary read in input_tokenizer() if the signature in the hash is longer than the length of the plugins' signature constant
+- Fixed out-of-boundary read in Stuffit5 module in hash_decode()
 
 ##
 ## Improvements
@@ -41,6 +45,7 @@
 - Kernel Cache: Add kernel threads into hash computation which is later used in the kernel cache filename
 - Memory Management: Refactored the code responsible for limiting kernel accel in order to avoid out of -host- memory situations
 - SCRYPT Kernels: Add more optimized values for some new NV/AMD GPUs
+- 7-Zip Hook: Increase supported data length from 320kb to 8mb
 
 ##
 ## Algorithms

docs/rules.txt

@@ -5,6 +5,7 @@
 #define RULE_OP_MANGLE_UREST_LFIRST     'C' // upper case all chars, lower case 1st
 #define RULE_OP_MANGLE_TREST            't' // switch the case of each char
 #define RULE_OP_MANGLE_TOGGLE_AT        'T' // switch the case of each char on pos N
+#define RULE_OP_MANGLE_TOGGLE_AT_SEP    '3' // switch the case of the first letter after occurrence N of char X
 #define RULE_OP_MANGLE_REVERSE          'r' // reverse word
 #define RULE_OP_MANGLE_DUPEWORD         'd' // append word to itself
 #define RULE_OP_MANGLE_DUPEWORD_TIMES   'p' // append word to itself N times

hashcat.hctune

@@ -386,189 +386,13 @@ GeForce_GTX_TITAN                               3       9900    2       A
 DEVICE_TYPE_CPU                                 *       3200    1       N       A
 
 ##
-## SCRYPT
+## SCRYPT: Tunings for SCRYPT based hash-modes can be found inside the plugin source
+##         See function module_extra_tuningdb_block()
 ##
 
-DEVICE_TYPE_CPU                                 *       8900    1       N       A
-DEVICE_TYPE_CPU                                 *       9300    1       N       A
-DEVICE_TYPE_CPU                                 *       15700   1       N       A
-DEVICE_TYPE_CPU                                 *       22700   1       N       A
-
-DEVICE_TYPE_GPU                                 *       8900    1       N       A
-DEVICE_TYPE_GPU                                 *       9300    1       N       A
-DEVICE_TYPE_GPU                                 *       15700   1       1       A
-DEVICE_TYPE_GPU                                 *       22700   1       N       A
-
 ##
 ## CryptoAPI
 ##
 
 DEVICE_TYPE_CPU                                 *       14500   1       A       A
 DEVICE_TYPE_GPU                                 *       14500   1       A       A
-
-## Here's an example of how to manually tune SCRYPT algorithm kernels for your hardware.
-## Manually tuning the GPU will yield increased performance. There is typically no noticeable change to CPU performance.
-##
-## First, you need to know the parameters of your SCRYPT hash: N, r and p.
-##
-## The reference SCRYPT parameter values are N=14, r=8 and p=1, but these will likely not match the parameters used by real-world applications.
-## For reference, the N value represents an exponent (2^N, which we calculate by bit shifting 1 left by N bits).
-## Hashcat expects this N value in decimal format: 1 << 14 = 16384
-##
-## Now that you have the 3 configuration items in decimal format, multiply them by 128 (underlaying crypto primitive block size).
-## For example: 128 * 16384 * 8 * 1 = 16777216 = 16MB
-## This is the amount of memory required for the GPU to compute the hash of one password candidate.
-##
-## Hashcat computes multiple password candidates in parallel - this is what allows for full utilization of the device.
-## The number of password candidates that Hashcat can run in parallel is VRAM limited and depends on:
-##
-## 1. Compute devices' native compute units
-## 2. Compute devices' native thread count
-## 3. Artificial multiplier (--kernel-accel aka -n)
-##
-## In order to find these values:
-##
-## 1. On startup Hashcat will show: * Device #1: GeForce GTX 980, 3963/4043 MB, 16MCU. The 16 MCU is the number of compute units on that device.
-## 2. Native thread counts are fixed values: CPU=1, GPU-Intel=8, GPU-AMD=64 (wavefronts), GPU-NVIDIA=32 (warps)
-##
-## Now multiply them together. For my GTX980: 16 * 32 * 16777216 = 8589934592 = 8GB
-##
-## If we want to actually make use of all computing resources, this GPU would require 8GB of GPU RAM.
-## However, it doesn't have that:
-##
-## Device #1: GeForce GTX 980, 3963/4043 MB, 16MCU. We only have 4043 MB (4GB minus some overhead from the OS).
-##
-## How do we deal with this? This is where SCRYPT TMTO(time-memory trde off) kicks in. The SCRYPT algorithm is designed in such a way that we
-## can pre-compute that 16MB buffer from a self-choosen offset. Details on how this actually works are not important for this process.
-##
-## What's relevant to us is that we can halve the buffer size, but we pay with twice the computation time.
-## We can repeat this as often as we want. That's why it's a trade-off.
-##
-## This mechanic can be manually set using --scrypt-tmto on the commandline, but this is not the best way.
-##
-## Back to our problem. We need 8GB of memory but have only ~4GB.
-## It's not a full 4GB. The OS needs some of it and Hashcat needs some of it to store password candidates and other things.
-## If you run a headless server it should be safe to subtract a fixed value of 200MB from whatever you have in your GPU.
-##
-## So lets divide our required memory(8GB) by 2 until it fits in our VRAM - 200MB.
-##
-## (8GB >> 0) = 8GB < 3.8GB = No, Does not fit
-## (8GB >> 1) = 4GB < 3.8GB = No, Does not fit
-## (8GB >> 2) = 2GB < 3.8GB = Yes!
-##
-## This process is automated in Hashcat, but it is important to understand what's happening here.
-## Because of the light overhead from the OS and Hashcat, we pay a very high price.
-## Even though it is just 200MB, it forces us to increase the TMTO by another step.
-## In terms of speed, the speed is now only 1/4 of what we could archieve on that same GPU if it had only 8.2GB ram.
-## But now we end up in a situation that we waste 1.8GB RAM which costs us ((1.8GB/16MB)>>1) candidates/second.
-##
-## This is where manual tuning can come into play.
-## If we know that the resources we need are close to what we have (in this case 3.8GB <-> 4.0GB)
-## We could decide to throw away some of our compute units so that we will no longer need 4.0GB but only 3.8GB.
-## Therefore, we do not need to increase the TMTO by another step to fit in VRAM.
-##
-## If we cut down our 16 MCU to only 15 MCU or 14 MCU using --kernel-accel(-n), we end up with:
-##
-## 16 * 32 * 16777216 = 8589934592 / 2 = 4294967296 = 4.00GB < 3.80GB = Nope, next
-## 15 * 32 * 16777216 = 8053063680 / 2 = 4026531840 = 3.84GB < 3.80GB = Nope, next
-## 14 * 32 * 16777216 = 7516192768 / 2 = 3758096384 = 3.58GB < 3.80GB = Yes!
-##
-## So we can throw away 2/16 compute units, but save half of the computation trade-off on the rest of the compute device.
-## On my GTX980, this improves the performance from 163 H/s to 201 H/s.
-## You don't need to control --scrypt-tmto manually because now that the multiplier (-n) is smaller than the native value
-## Hashcat will automatically realize it can decrease the TMTO by one.
-##
-## At this point, you found the optimal base value for your compute device. In this case: 14.
-##
-## Depending on our hardware, especially hardware with very slow memory access like a GPU
-## there's a good chance that it's cheaper (faster) to compute an extra step on the GPU register.
-## So if we increase the TMTO again by one, this gives an extra speed boost.
-##
-## On my GTX980, this improves the performance from 201 H/s to 255 H/s.
-## Again, there's no need to control this with --scrypt-tmto. Hashcat will realize it has to increase the TMTO again.
-##
-## All together, you can control all of this by using the -n parameter in the command line.
-## This is not ideal in a production environment because you must use the --force flag.
-## The best way to set this is by using this Hashcat.hctune file to store it. This avoids the need to bypass any warnings.
-##
-## Find the ideal -n value, then store it here along with the proper compute device name.
-## Formatting guidelines are availabe at the top of this document.
-##
-## -------------------------------------------------
-##
-## You can also ignore all theoretical derivations and semi-automate the process in the real scenario (I prefer this approach):
-##
-## 1. For example, to find the value for 8900, first create a valid hash for 8900 as follows:
-##
-## $ ./hashcat --example-hashes -m 8900 | grep Example.Hash | grep -v Format | cut -b 25- > tmp.hash.8900
-##
-## 2. Now let it iterate through all -n values to a certain point. In this case, I'm using 200, but in general it's a value that is at least twice that of the multiprocessor. If you don't mind you can just leave it as it is, it just runs a little longer.
-##
-## $ export i=1; while [ $i -ne 201 ]; do echo $i; ./hashcat --quiet tmp.hash.8900 --keep-guessing --self-test-disable --markov-disable --restore-disable --outfile-autohex-disable --wordlist-autohex-disable --potfile-disable --logfile-disable --hwmon-disable --status --status-timer 1 --runtime 28 --machine-readable --optimized-kernel-enable --workload-profile 3 --hash-type 8900 --attack-mode 3 ?b?b?b?b?b?b?b --backend-devices 1 --force -n $i; i=$(($i+1)); done | tee x
-##
-## 3. Determine the highest measured H/s speed. But don't just use the highest value. Instead, use the number that seems most stable, usually at the beginning.
-##
-## $ grep "$(printf 'STATUS\t3')" x | cut -f4 -d$'\t' | sort -n | tail
-##
-## 4. To match the speed you have chosen to the correct value in the "x" file, simply search for it in it. Then go up a little on the block where you found him. The value -n is the single value that begins before the block start. If you have multiple blocks at the same speed, choose the lowest value for -n
-##
-
-## 4GB
-GeForce_GTX_980                                 *       8900    1      29       A
-GeForce_GTX_980                                 *       9300    1     128       A
-GeForce_GTX_980                                 *       15700   1      24       A
-GeForce_GTX_980                                 *       22700   1      29       A
-
-## 8GB
-GeForce_GTX_1080                                *       8900    1      15       A
-GeForce_GTX_1080                                *       9300    1     256       A
-GeForce_GTX_1080                                *       15700   1      28       A
-GeForce_GTX_1080                                *       22700   1      15       A
-
-## 11GB
-GeForce_RTX_2080_Ti                             *       8900    1      68       A
-GeForce_RTX_2080_Ti                             *       9300    1     528       A
-GeForce_RTX_2080_Ti                             *       15700   1      68       A
-GeForce_RTX_2080_Ti                             *       22700   1      68       A
-
-## 8GB
-GeForce_RTX_3060_Ti                             *       8900    1      51       A
-GeForce_RTX_3060_Ti                             *       9300    1     256       A
-GeForce_RTX_3060_Ti                             *       15700   1      11       A
-GeForce_RTX_3060_Ti                             *       22700   1      51       A
-
-## 8GB
-GeForce_RTX_3070                                *       8900    1      46       A
-GeForce_RTX_3070                                *       9300    1     368       A
-GeForce_RTX_3070                                *       15700   1      22       A
-GeForce_RTX_3070                                *       22700   1      46       A
-
-## 24GB
-GeForce_RTX_3090                                *       8900    1      82       A
-GeForce_RTX_3090                                *       9300    1     984       A
-GeForce_RTX_3090                                *       15700   1      82       A
-GeForce_RTX_3090                                *       22700   1      82       A
-
-## 4GB
-ALIAS_AMD_RX480                                 *       8900    1      15       A
-ALIAS_AMD_RX480                                 *       9300    1     232       A
-ALIAS_AMD_RX480                                 *       15700   1      58       A
-ALIAS_AMD_RX480                                 *       22700   1      15       A
-
-## 8GB
-ALIAS_AMD_Vega64                                *       8900    1      31       A
-ALIAS_AMD_Vega64                                *       9300    1     440       A
-ALIAS_AMD_Vega64                                *       15700   1      53       A
-ALIAS_AMD_Vega64                                *       22700   1      31       A
-
-## 32GB
-ALIAS_AMD_MI100                                 *       8900    1      79       A
-ALIAS_AMD_MI100                                 *       9300    1    1000       A
-ALIAS_AMD_MI100                                 *       15700   1     120       A
-ALIAS_AMD_MI100                                 *       22700   1      79       A
-
-## 16GB
-ALIAS_AMD_RX6900XT                              *       8900    1      59       A
-ALIAS_AMD_RX6900XT                              *       9300    1     720       A
-ALIAS_AMD_RX6900XT                              *       15700   1      56       A
-ALIAS_AMD_RX6900XT                              *       22700   1      59       A

include/backend.h

@@ -69,13 +69,12 @@ int hc_cuFuncSetAttribute        (hashcat_ctx_t *hashcat_ctx, CUfunction hfunc,
 int hc_cuInit                    (hashcat_ctx_t *hashcat_ctx, unsigned int Flags);
 int hc_cuLaunchKernel            (hashcat_ctx_t *hashcat_ctx, CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra);
 int hc_cuMemAlloc                (hashcat_ctx_t *hashcat_ctx, CUdeviceptr *dptr, size_t bytesize);
-int hc_cuMemcpyDtoD              (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount);
 int hc_cuMemcpyDtoDAsync         (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream);
-int hc_cuMemcpyDtoH              (hashcat_ctx_t *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount);
 int hc_cuMemcpyDtoHAsync         (hashcat_ctx_t *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream);
-int hc_cuMemcpyHtoD              (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount);
 int hc_cuMemcpyHtoDAsync         (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount, CUstream hStream);
 int hc_cuMemFree                 (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dptr);
+int hc_cuMemsetD32Async          (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N, CUstream hStream);
+int hc_cuMemsetD8Async           (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N, CUstream hStream);
 int hc_cuModuleGetFunction       (hashcat_ctx_t *hashcat_ctx, CUfunction *hfunc, CUmodule hmod, const char *name);
 int hc_cuModuleLoadDataEx        (hashcat_ctx_t *hashcat_ctx, CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues);
 int hc_cuModuleUnload            (hashcat_ctx_t *hashcat_ctx, CUmodule hmod);
@@ -120,12 +119,11 @@ int hc_hipInit                   (hashcat_ctx_t *hashcat_ctx, unsigned int Flags
 int hc_hipLaunchKernel           (hashcat_ctx_t *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra);
 int hc_hipMemAlloc               (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize);
 int hc_hipMemFree                (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dptr);
-int hc_hipMemcpyDtoD             (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount);
-int hc_hipMemcpyDtoHAsync        (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
-int hc_hipMemcpyDtoH             (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount);
 int hc_hipMemcpyDtoDAsync        (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
-int hc_hipMemcpyHtoD             (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount);
+int hc_hipMemcpyDtoHAsync        (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
 int hc_hipMemcpyHtoDAsync        (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream);
+int hc_hipMemsetD32Async         (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N, hipStream_t hStream);
+int hc_hipMemsetD8Async          (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N, hipStream_t hStream);
 int hc_hipModuleGetFunction      (hashcat_ctx_t *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name);
 int hc_hipModuleGetGlobal        (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name);
 int hc_hipModuleLoadDataEx       (hashcat_ctx_t *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues);
@@ -143,6 +141,7 @@ int hc_clCreateKernel            (hashcat_ctx_t *hashcat_ctx, cl_program program
 int hc_clCreateProgramWithBinary (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status, cl_program *program);
 int hc_clCreateProgramWithSource (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_program *program);
 int hc_clEnqueueCopyBuffer       (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_buffer, size_t src_offset, size_t dst_offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
+int hc_clEnqueueFillBuffer       (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_mem buffer, const void *pattern, size_t pattern_size, size_t offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
 int hc_clEnqueueMapBuffer        (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_map, cl_map_flags map_flags, size_t offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event, void **buf);
 int hc_clEnqueueNDRangeKernel    (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
 int hc_clEnqueueReadBuffer       (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
@@ -178,17 +177,20 @@ void rebuild_pws_compressed_append (hc_device_param_t *device_param, const u64 p
 
 int run_cuda_kernel_atinit          (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 num);
 int run_cuda_kernel_utf8toutf16le   (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 num);
-int run_cuda_kernel_memset          (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u32 value, const u64 size);
+int run_cuda_kernel_memset          (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u8  value, const u64 size);
+int run_cuda_kernel_memset32        (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u32 value, const u64 size);
 int run_cuda_kernel_bzero           (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 size);
 
 int run_hip_kernel_atinit           (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 num);
 int run_hip_kernel_utf8toutf16le    (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 num);
-int run_hip_kernel_memset           (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u32 value, const u64 size);
+int run_hip_kernel_memset           (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u8  value, const u64 size);
+int run_hip_kernel_memset32         (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u32 value, const u64 size);
 int run_hip_kernel_bzero            (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 size);
 
 int run_opencl_kernel_atinit        (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num);
 int run_opencl_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num);
-int run_opencl_kernel_memset        (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u32 value, const u64 size);
+int run_opencl_kernel_memset        (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 offset, const u8  value, const u64 size);
+int run_opencl_kernel_memset32      (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 offset, const u32 value, const u64 size);
 int run_opencl_kernel_bzero         (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 size);
 
 int run_kernel                (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kern_run, const u64 pws_pos, const u64 num, const u32 event_update, const u32 iteration);

include/common.h

@@ -63,6 +63,14 @@
 #define HC_API_CALL
 #endif
 
+#if defined (__GNUC__)
+#define HC_ALIGN(x) __attribute__((aligned(x)))
+#elif defined (_MSC_VER)
+#define HC_ALIGN(x) __declspec(align(x))
+#else
+#define HC_ALIGN(x)
+#endif
+
 #if defined (_WIN)
 #define WIN32_LEAN_AND_MEAN
 #endif

include/ext_OpenCL.h

@@ -46,6 +46,7 @@ typedef cl_context       (CL_API_CALL *OCL_CLCREATECONTEXT)           (const cl_
 typedef cl_kernel        (CL_API_CALL *OCL_CLCREATEKERNEL)            (cl_program, const char *, cl_int *);
 typedef cl_program       (CL_API_CALL *OCL_CLCREATEPROGRAMWITHBINARY) (cl_context, cl_uint, const cl_device_id *, const size_t *, const unsigned char **, cl_int *, cl_int *);
 typedef cl_program       (CL_API_CALL *OCL_CLCREATEPROGRAMWITHSOURCE) (cl_context, cl_uint, const char **, const size_t *, cl_int *);
+typedef cl_int           (CL_API_CALL *OCL_CLENQUEUEFILLBUFFER)       (cl_command_queue, cl_mem, const void *, size_t, size_t, size_t, cl_uint, const cl_event *, cl_event *);
 typedef cl_int           (CL_API_CALL *OCL_CLENQUEUECOPYBUFFER)       (cl_command_queue, cl_mem, cl_mem, size_t, size_t, size_t, cl_uint, const cl_event *, cl_event *);
 typedef void *           (CL_API_CALL *OCL_CLENQUEUEMAPBUFFER)        (cl_command_queue, cl_mem, cl_bool, cl_map_flags, size_t, size_t, cl_uint, const cl_event *, cl_event *, cl_int *);
 typedef cl_int           (CL_API_CALL *OCL_CLENQUEUENDRANGEKERNEL)    (cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *);
@@ -87,6 +88,7 @@ typedef struct hc_opencl_lib
   OCL_CLCREATEPROGRAMWITHBINARY clCreateProgramWithBinary;
   OCL_CLCREATEPROGRAMWITHSOURCE clCreateProgramWithSource;
   OCL_CLENQUEUECOPYBUFFER       clEnqueueCopyBuffer;
+  OCL_CLENQUEUEFILLBUFFER       clEnqueueFillBuffer;
   OCL_CLENQUEUEMAPBUFFER        clEnqueueMapBuffer;
   OCL_CLENQUEUENDRANGEKERNEL    clEnqueueNDRangeKernel;
   OCL_CLENQUEUEREADBUFFER       clEnqueueReadBuffer;

include/ext_cuda.h

@@ -1028,17 +1028,14 @@ typedef CUresult (CUDA_API_CALL *CUDA_CUINIT)                   (unsigned int);
 typedef CUresult (CUDA_API_CALL *CUDA_CULAUNCHKERNEL)           (CUfunction, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, CUstream, void **, void **);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMEMALLOC)               (CUdeviceptr *, size_t);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMEMALLOCHOST)           (void **, size_t);
-typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOD)             (CUdeviceptr, CUdeviceptr, size_t);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTODASYNC)        (CUdeviceptr, CUdeviceptr, size_t, CUstream);
-typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOH)             (void *, CUdeviceptr, size_t);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOHASYNC)        (void *, CUdeviceptr, size_t, CUstream);
-typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYHTOD)             (CUdeviceptr, const void *, size_t);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYHTODASYNC)        (CUdeviceptr, const void *, size_t, CUstream);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMEMFREE)                (CUdeviceptr);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMEMFREEHOST)            (void *);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMEMGETINFO)             (size_t *, size_t *);
-typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32)              (CUdeviceptr, unsigned int, size_t);
+typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32ASYNC)         (CUdeviceptr, unsigned int, size_t, CUstream);
-typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8)               (CUdeviceptr, unsigned char, size_t);
+typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8ASYNC)          (CUdeviceptr, unsigned char, size_t, CUstream);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMODULEGETFUNCTION)      (CUfunction *, CUmodule, const char *);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMODULEGETGLOBAL)        (CUdeviceptr *, size_t *, CUmodule, const char *);
 typedef CUresult (CUDA_API_CALL *CUDA_CUMODULELOAD)             (CUmodule *, const char *);
@@ -1093,17 +1090,14 @@ typedef struct hc_cuda_lib
   CUDA_CULAUNCHKERNEL           cuLaunchKernel;
   CUDA_CUMEMALLOC               cuMemAlloc;
   CUDA_CUMEMALLOCHOST           cuMemAllocHost;
-  CUDA_CUMEMCPYDTOD             cuMemcpyDtoD;
   CUDA_CUMEMCPYDTODASYNC        cuMemcpyDtoDAsync;
-  CUDA_CUMEMCPYDTOH             cuMemcpyDtoH;
   CUDA_CUMEMCPYDTOHASYNC        cuMemcpyDtoHAsync;
-  CUDA_CUMEMCPYHTOD             cuMemcpyHtoD;
   CUDA_CUMEMCPYHTODASYNC        cuMemcpyHtoDAsync;
   CUDA_CUMEMFREE                cuMemFree;
   CUDA_CUMEMFREEHOST            cuMemFreeHost;
   CUDA_CUMEMGETINFO             cuMemGetInfo;
-  CUDA_CUMEMSETD32              cuMemsetD32;
+  CUDA_CUMEMSETD32ASYNC         cuMemsetD32Async;
-  CUDA_CUMEMSETD8               cuMemsetD8;
+  CUDA_CUMEMSETD8ASYNC          cuMemsetD8Async;
   CUDA_CUMODULEGETFUNCTION      cuModuleGetFunction;
   CUDA_CUMODULEGETGLOBAL        cuModuleGetGlobal;
   CUDA_CUMODULELOAD             cuModuleLoad;

include/ext_hip.h

@@ -378,12 +378,11 @@ typedef hipError_t (HIP_API_CALL *HIP_HIPLAUNCHKERNEL)           (hipFunction_t,
 typedef hipError_t (HIP_API_CALL *HIP_HIPMEMALLOC)               (hipDeviceptr_t *, size_t);
 typedef hipError_t (HIP_API_CALL *HIP_HIPMEMFREE)                (hipDeviceptr_t);
 typedef hipError_t (HIP_API_CALL *HIP_HIPMEMGETINFO)             (size_t *, size_t *);
-typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOD)             (hipDeviceptr_t, hipDeviceptr_t, size_t);
 typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTODASYNC)        (hipDeviceptr_t, hipDeviceptr_t, size_t, hipStream_t);
-typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOH)             (void *, hipDeviceptr_t, size_t);
 typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOHASYNC)        (void *, hipDeviceptr_t, size_t, hipStream_t);
-typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTOD)             (hipDeviceptr_t, const void *, size_t);
 typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTODASYNC)        (hipDeviceptr_t, const void *, size_t, hipStream_t);
+typedef hipError_t (HIP_API_CALL *HIP_HIPMEMSETD32ASYNC)         (hipDeviceptr_t, unsigned int, size_t, hipStream_t);
+typedef hipError_t (HIP_API_CALL *HIP_HIPMEMSETD8ASYNC)          (hipDeviceptr_t, unsigned char, size_t, hipStream_t);
 typedef hipError_t (HIP_API_CALL *HIP_HIPMODULEGETFUNCTION)      (hipFunction_t *, hipModule_t, const char *);
 typedef hipError_t (HIP_API_CALL *HIP_HIPMODULEGETGLOBAL)        (hipDeviceptr_t *, size_t *, hipModule_t, const char *);
 typedef hipError_t (HIP_API_CALL *HIP_HIPMODULELOADDATAEX)       (hipModule_t *, const void *, unsigned int, hipJitOption *, void **);
@@ -421,12 +420,11 @@ typedef struct hc_hip_lib
   HIP_HIPMEMALLOC               hipMemAlloc;
   HIP_HIPMEMFREE                hipMemFree;
   HIP_HIPMEMGETINFO             hipMemGetInfo;
-  HIP_HIPMEMCPYDTOD             hipMemcpyDtoD;
   HIP_HIPMEMCPYDTODASYNC        hipMemcpyDtoDAsync;
-  HIP_HIPMEMCPYDTOH             hipMemcpyDtoH;
   HIP_HIPMEMCPYDTOHASYNC        hipMemcpyDtoHAsync;
-  HIP_HIPMEMCPYHTOD             hipMemcpyHtoD;
   HIP_HIPMEMCPYHTODASYNC        hipMemcpyHtoDAsync;
+  HIP_HIPMEMSETD32ASYNC         hipMemsetD32Async;
+  HIP_HIPMEMSETD8ASYNC          hipMemsetD8Async;
   HIP_HIPMODULEGETFUNCTION      hipModuleGetFunction;
   HIP_HIPMODULEGETGLOBAL        hipModuleGetGlobal;
   HIP_HIPMODULELOADDATAEX       hipModuleLoadDataEx;

include/hashes.h

@@ -15,7 +15,7 @@ int hash_encode (const hashconfig_t *hashconfig, const hashes_t *hashes, const m
 
 int save_hash (hashcat_ctx_t *hashcat_ctx);
 
-void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, plain_t *plain);
+int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, plain_t *plain);
 
 //int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 salt_pos);
 int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param);

include/interface.h

@@ -13,7 +13,7 @@
 #include <limits.h>
 #include <inttypes.h>
 
-static const int MODULE_INTERFACE_VERSION_MINIMUM = 520;
+static const int MODULE_INTERFACE_VERSION_MINIMUM = 630;
 
 static const int MODULE_HASH_MODES_MAXIMUM = 100000;
 

include/modules.h

@@ -18,6 +18,7 @@ u32         module_dgst_pos2                (MAYBE_UNUSED const hashconfig_t *ha
 u32         module_dgst_pos3                (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
 u32         module_dgst_size                (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
 u64         module_esalt_size               (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
+const char *module_extra_tuningdb_block     (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
 u32         module_forced_outfile_format    (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
 u32         module_hash_category            (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
 const char *module_hash_name                (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);

include/tuningdb.h

@@ -11,9 +11,13 @@
 
 #define TUNING_DB_FILE "hashcat.hctune"
 
+int sort_by_tuning_db_alias (const void *v1, const void *v2);
+int sort_by_tuning_db_entry (const void *v1, const void *v2);
+
 int  tuning_db_init    (hashcat_ctx_t *hashcat_ctx);
 void tuning_db_destroy (hashcat_ctx_t *hashcat_ctx);
 
+bool tuning_db_process_line (hashcat_ctx_t *hashcat_ctx, const char *line_buf, const int line_num);
 tuning_db_entry_t *tuning_db_search (hashcat_ctx_t *hashcat_ctx, const char *device_name, const cl_device_type device_type, int attack_mode, const int hash_mode);
 
 #endif // _TUNINGDB_H

include/types.h

@@ -295,6 +295,7 @@ typedef enum rule_functions
   RULE_OP_MANGLE_UREST_LFIRST    = 'C',
   RULE_OP_MANGLE_TREST           = 't',
   RULE_OP_MANGLE_TOGGLE_AT       = 'T',
+  RULE_OP_MANGLE_TOGGLE_AT_SEP   = '3',
   RULE_OP_MANGLE_REVERSE         = 'r',
   RULE_OP_MANGLE_DUPEWORD        = 'd',
   RULE_OP_MANGLE_DUPEWORD_TIMES  = 'p',
@@ -926,7 +927,6 @@ typedef struct hashes
 
   void        *digests_buf;
   u32         *digests_shown;
-  u32         *digests_shown_tmp;
 
   u32          salts_cnt;
   u32          salts_done;
@@ -1427,6 +1427,7 @@ typedef struct hc_device_param
 
   CUevent           cuda_event1;
   CUevent           cuda_event2;
+  CUevent           cuda_event3;
 
   CUmodule          cuda_module;
   CUmodule          cuda_module_shared;
@@ -1508,6 +1509,7 @@ typedef struct hc_device_param
 
   hipEvent_t        hip_event1;
   hipEvent_t        hip_event2;
+  hipEvent_t        hip_event3;
 
   hipModule_t       hip_module;
   hipModule_t       hip_module_shared;
@@ -2042,9 +2044,11 @@ typedef struct tuning_db
 
   tuning_db_alias_t *alias_buf;
   int                alias_cnt;
+  int                alias_alloc;
 
   tuning_db_entry_t *entry_buf;
   int                entry_cnt;
+  int                entry_alloc;
 
 } tuning_db_t;
 
@@ -2648,6 +2652,7 @@ typedef struct module_ctx
   u32         (*module_dgst_pos3)               (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
   u32         (*module_dgst_size)               (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
   u64         (*module_esalt_size)              (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
+  const char *(*module_extra_tuningdb_block)    (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
   u32         (*module_forced_outfile_format)   (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
   u32         (*module_hash_category)           (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
   const char *(*module_hash_name)               (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);

src/Makefile

@@ -53,7 +53,7 @@ endif
 ## Do not modify
 ##
 
-MODULE_INTERFACE_VERSION := 620
+MODULE_INTERFACE_VERSION := 630
 
 ##
 ## Native compiler paths
@@ -203,12 +203,11 @@ endif
 ## because UNRAR
 ifeq ($(ENABLE_UNRAR),1)
 ifeq ($(USE_SYSTEM_UNRAR),0)
-ifneq ($(UNAME),Darwin)
+ifneq ($(CC),clang)
-CFLAGS_UNRAR            += -Wno-misleading-indentation
 CFLAGS_UNRAR            += -Wno-class-memaccess
-else
+CFLAGS_UNRAR            += -Wno-misleading-indentation
-CFLAGS_UNRAR            += -Wno-missing-braces
 endif
+CFLAGS_UNRAR            += -Wno-missing-braces
 CFLAGS_UNRAR            += -Wno-unused-variable
 CFLAGS_UNRAR            += -Wno-unused-parameter
 CFLAGS_UNRAR            += -Wno-unused-function

src/autotune.c

@@ -10,7 +10,7 @@
 #include "status.h"
 #include "autotune.h"
 
-static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops)
+static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops, const u32 kernel_threads)
 {
   hashconfig_t   *hashconfig   = hashcat_ctx->hashconfig;
   user_options_t *user_options = hashcat_ctx->user_options;
@@ -19,7 +19,9 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
   device_param->kernel_params_buf32[29] = kernel_loops; // not a bug, both need to be set
   device_param->kernel_params_buf32[30] = kernel_loops; // because there's two variables for inner iters for slow and fast hashes
 
-  u32 kernel_power_try = device_param->hardware_power * kernel_accel;
+  const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads;
+
+  u32 kernel_power_try = hardware_power * kernel_accel;
 
   if (user_options->attack_mode == ATTACK_MODE_ASSOCIATION)
   {
@@ -33,6 +35,10 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
     }
   }
 
+  const u32 kernel_threads_sav = device_param->kernel_threads;
+
+  device_param->kernel_threads = kernel_threads;
+
   const double spin_damp_sav = device_param->spin_damp;
 
   device_param->spin_damp = 0;
@@ -50,71 +56,51 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
   }
   else
   {
-    run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_try, true, 0);
-
-    if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE)
-    {
-      run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, kernel_power_try, true, 0);
-    }
-
     run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, kernel_power_try, true, 0);
   }
 
   device_param->spin_damp = spin_damp_sav;
 
-  const double exec_msec_prev = get_avg_exec_time (device_param, 1);
-
-  return exec_msec_prev;
-}
-
-/*
-static double try_run_preferred (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops)
-{
-  hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
-
-  device_param->kernel_params_buf32[28] = 0;
-  device_param->kernel_params_buf32[29] = kernel_loops; // not a bug, both need to be set
-  device_param->kernel_params_buf32[30] = kernel_loops; // because there's two variables for inner iters for slow and fast hashes
-
-  const u32 kernel_power_try = device_param->hardware_power * kernel_accel;
-
-  const u32 kernel_threads_sav = device_param->kernel_threads;
-
-  const double spin_damp_sav = device_param->spin_damp;
-
-  device_param->spin_damp = 0;
-
-  if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
-  {
-    if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
-    {
-      device_param->kernel_threads = device_param->kernel_preferred_wgs_multiple1;
-
-      run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_try, true, 0);
-    }
-    else
-    {
-      device_param->kernel_threads = device_param->kernel_preferred_wgs_multiple4;
-
-      run_kernel (hashcat_ctx, device_param, KERN_RUN_4, 0, kernel_power_try, true, 0);
-    }
-  }
-  else
-  {
-    device_param->kernel_threads = device_param->kernel_preferred_wgs_multiple2;
-
-    run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, kernel_power_try, true, 0);
-  }
-
   device_param->kernel_threads = kernel_threads_sav;
 
-  device_param->spin_damp = spin_damp_sav;
-
   const double exec_msec_prev = get_avg_exec_time (device_param, 1);
 
   return exec_msec_prev;
 }
-*/
+
+static double try_run_times (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops, const u32 kernel_threads, const int times)
+{
+  double exec_msec_best = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
+
+  for (int i = 1; i < times; i++)
+  {
+    double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
+
+    if (exec_msec > exec_msec_best) continue;
+
+    exec_msec_best = exec_msec;
+  }
+
+  return exec_msec_best;
+}
+
+static u32 previous_power_of_two (const u32 x)
+{
+  // https://stackoverflow.com/questions/2679815/previous-power-of-2
+  // really cool!
+
+  if (x == 0) return 0;
+
+  u32 r = x;
+
+  r |= (r >>  1);
+  r |= (r >>  2);
+  r |= (r >>  4);
+  r |= (r >>  8);
+  r |= (r >> 16);
+
+  return r - (r >> 1);
+}
 
 static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
 {
@@ -131,9 +117,57 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
   const u32 kernel_loops_min = device_param->kernel_loops_min;
   const u32 kernel_loops_max = device_param->kernel_loops_max;
 
+  const u32 kernel_threads_min = device_param->kernel_threads_min;
+  const u32 kernel_threads_max = device_param->kernel_threads_max;
+
   u32 kernel_accel = kernel_accel_min;
   u32 kernel_loops = kernel_loops_min;
 
+  // for the threads we take as initial value what we receive from the runtime
+  // but is only to start with something, we will fine tune this value as soon as we have our workload specified
+  // this thread limiting is also performed insinde run_kernel() so we need to redo it here, too
+
+  u32 kernel_wgs = 0;
+  u32 kernel_wgs_multiple = 0;
+
+  if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
+  {
+    if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
+    {
+      kernel_wgs = device_param->kernel_wgs1;
+
+      kernel_wgs_multiple = device_param->kernel_preferred_wgs_multiple1;
+    }
+    else
+    {
+      kernel_wgs = device_param->kernel_wgs4;
+
+      kernel_wgs_multiple = device_param->kernel_preferred_wgs_multiple4;
+    }
+  }
+  else
+  {
+    kernel_wgs = device_param->kernel_wgs2;
+
+    kernel_wgs_multiple = device_param->kernel_preferred_wgs_multiple2;
+  }
+
+  u32 kernel_threads = kernel_threads_max;
+
+  if ((kernel_wgs >= kernel_threads_min) && (kernel_wgs <= kernel_threads_max))
+  {
+    kernel_threads = kernel_wgs;
+  }
+
+  // having a value power of 2 makes it easier to divide
+
+  const u32 kernel_threads_p2 = previous_power_of_two (kernel_threads);
+
+  if ((kernel_threads_p2 >= kernel_threads_min) && (kernel_threads_p2 <= kernel_threads_max))
+  {
+    kernel_threads = kernel_threads_p2;
+  }
+
   // in this case the user specified a fixed -n and -u on the commandline
   // no way to tune anything
   // but we need to run a few caching rounds
@@ -149,10 +183,10 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
 
     if (hashconfig->warmup_disable == false)
     {
-      try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
+      try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
-      try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
+      try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
-      try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
+      try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
-      try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
+      try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
     }
 
     #endif
@@ -164,29 +198,19 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
 
     const u32 kernel_power_max = device_param->hardware_power * kernel_accel_max;
 
-    int CU_rc;
-    int HIP_rc;
-    int CL_rc;
-
     if (device_param->is_cuda == true)
     {
-      CU_rc = run_cuda_kernel_atinit (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, kernel_power_max);
+      if (run_cuda_kernel_atinit (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, kernel_power_max) == -1) return -1;
-
-      if (CU_rc == -1) return -1;
     }
 
     if (device_param->is_hip == true)
     {
-      HIP_rc = run_hip_kernel_atinit (hashcat_ctx, device_param, device_param->hip_d_pws_buf, kernel_power_max);
+      if (run_hip_kernel_atinit (hashcat_ctx, device_param, device_param->hip_d_pws_buf, kernel_power_max) == -1) return -1;
-
-      if (HIP_rc == -1) return -1;
     }
 
     if (device_param->is_opencl == true)
     {
-      CL_rc = run_opencl_kernel_atinit (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, kernel_power_max);
+      if (run_opencl_kernel_atinit (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, kernel_power_max) == -1) return -1;
-
-      if (CL_rc == -1) return -1;
     }
 
     if (user_options->slow_candidates == true)
@@ -200,26 +224,44 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
         {
           if (device_param->is_cuda == true)
           {
-            CU_rc = hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t));
+            if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), device_param->cuda_stream) == -1) return -1;
-
-            if (CU_rc == -1) return -1;
           }
 
           if (device_param->is_hip == true)
           {
-            HIP_rc = hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t));
+            if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), device_param->hip_stream) == -1) return -1;
-
-            if (HIP_rc == -1) return -1;
           }
 
           if (device_param->is_opencl == true)
           {
-            CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, device_param->opencl_d_rules_c, 0, 0, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL);
+            if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, device_param->opencl_d_rules_c, 0, 0, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL) == -1) return -1;
+          }
+        }
+      }
+    }
 
-            if (CL_rc == -1) return -1;
+    // we also need to initialize some values using kernels
-          }
+
+    if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
+    {
+      // nothing to do
     }
+    else
+    {
+      const u32 kernel_threads_sav = device_param->kernel_threads;
+
+      device_param->kernel_threads = device_param->kernel_wgs1;
+
+      run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_max, false, 0);
+
+      if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE)
+      {
+        device_param->kernel_threads = device_param->kernel_wgs2p;
+
+        run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, kernel_power_max, false, 0);
       }
+
+      device_param->kernel_threads = kernel_threads_sav;
     }
 
     // Do a pre-autotune test run to find out if kernel runtime is above some TDR limit
@@ -228,7 +270,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
 
     if (true)
     {
-      double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min);
+      double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min, kernel_threads);
 
       if (exec_msec > 2000)
       {
@@ -237,7 +279,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
         return -1;
       }
 
-      exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min);
+      exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min, kernel_threads);
 
       const u32 mm = kernel_loops_max / kernel_loops_min;
 
@@ -257,16 +299,16 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
       {
         if (kernel_loops > kernel_loops_max_reduced) continue;
 
-        double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops);
+        double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_min, kernel_loops, kernel_threads, 1);
 
         if (exec_msec < target_msec) break;
       }
     }
 
-    // now the same for kernel-accel but with the new kernel-loops from previous loop set
-
     #define STEPS_CNT 16
 
+    // now the same for kernel-accel but with the new kernel-loops from previous loop set
+
     if (kernel_accel_min < kernel_accel_max)
     {
       for (int i = 0; i < STEPS_CNT; i++)
@@ -276,7 +318,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
         if (kernel_accel_try < kernel_accel_min) continue;
         if (kernel_accel_try > kernel_accel_max) break;
 
-        double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_try, kernel_loops);
+        double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_try, kernel_loops, kernel_threads, 1);
 
         if (exec_msec > target_msec) break;
 
@@ -292,7 +334,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
       const u32 kernel_accel_orig = kernel_accel;
       const u32 kernel_loops_orig = kernel_loops;
 
-      double exec_msec_prev = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
+      double exec_msec_prev = try_run_times (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads, 1);
 
       for (int i = 1; i < STEPS_CNT; i++)
       {
@@ -307,7 +349,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
 
         // do a real test
 
-        const double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_try, kernel_loops_try);
+        const double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_try, kernel_loops_try, kernel_threads, 1);
 
         if (exec_msec_prev < exec_msec) break;
 
@@ -324,7 +366,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
       }
     }
 
-    double exec_msec_pre_final = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
+    double exec_msec_pre_final = try_run_times (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads, 1);
 
     const u32 exec_left = (const u32) (target_msec / exec_msec_pre_final);
 
@@ -339,124 +381,87 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
       kernel_accel *= exec_accel_min;
     }
 
-    // start finding best thread count is easier.
+    // v6.2.4 new section: find thread count
-    // it's either the preferred or the maximum thread count
+    // This is not as effective as it could be because of inaccurate kernel return timers
+    // But is better than fixed values
+    // Timers in this section are critical, so we rerun meassurements 3 times
 
-    /*
+    if (kernel_threads_max > kernel_threads_min)
-    const u32 kernel_threads_min = device_param->kernel_threads_min;
+    {
-    const u32 kernel_threads_max = device_param->kernel_threads_max;
+      const u32 kernel_accel_orig   = kernel_accel;
+      const u32 kernel_threads_orig = kernel_threads;
 
-    if (kernel_threads_min < kernel_threads_max)
+      double exec_msec_prev = try_run_times (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads, 3);
-    {
-      const double exec_msec_max = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
 
-      u32 preferred_threads = 0;
+      for (int i = 1; i < STEPS_CNT; i++)
+      {
+        const u32 kernel_accel_try   = kernel_accel_orig   * (1U << i);
+        const u32 kernel_threads_try = kernel_threads_orig / (1U << i);
 
-      if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
+        // since we do not modify total amount of workitems, we can (and need) to do increase kernel_accel_max
-      {
-        if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
-        {
-          preferred_threads = device_param->kernel_preferred_wgs_multiple1;
-        }
-        else
-        {
-          preferred_threads = device_param->kernel_preferred_wgs_multiple4;
-        }
-      }
-      else
-      {
-        preferred_threads = device_param->kernel_preferred_wgs_multiple2;
-      }
 
-      if ((preferred_threads >= kernel_threads_min) && (preferred_threads <= kernel_threads_max))
+        const u32 kernel_accel_max_try = kernel_accel_max * (1U << i);
-      {
-        const double exec_msec_preferred = try_run_preferred (hashcat_ctx, device_param, kernel_accel, kernel_loops);
 
-        if (exec_msec_preferred < exec_msec_max)
+        if (kernel_accel_try > kernel_accel_max_try) break;
-        {
+
-          device_param->kernel_threads = preferred_threads;
+        if (kernel_threads_try < kernel_threads_min) break;
+
+        if (kernel_threads_try % kernel_wgs_multiple) break; // this would just be waste of time
+
+        double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_try, kernel_loops, kernel_threads_try, 3);
+
+        if (exec_msec > exec_msec_prev) continue;
+
+        exec_msec_prev = exec_msec;
+
+        kernel_accel   = kernel_accel_try;
+        kernel_threads = kernel_threads_try;
       }
     }
   }
-    */
-  }
 
   // reset them fake words
   // reset other buffers in case autotune cracked something
 
   if (device_param->is_cuda == true)
   {
-    int CU_rc;
+    if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, device_param->size_pws) == -1) return -1;
 
-    CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, device_param->size_pws);
+    if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_plain_bufs, device_param->size_plains) == -1) return -1;
 
-    if (CU_rc == -1) return -1;
+    if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_digests_shown, device_param->size_shown) == -1) return -1;
 
-    CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_plain_bufs, device_param->size_plains);
+    if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_result, device_param->size_results) == -1) return -1;
 
-    if (CU_rc == -1) return -1;
+    if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_tmps, device_param->size_tmps) == -1) return -1;
-
-    CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_digests_shown, device_param->size_shown);
-
-    if (CU_rc == -1) return -1;
-
-    CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_result, device_param->size_results);
-
-    if (CU_rc == -1) return -1;
-
-    CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_tmps, device_param->size_tmps);
-
-    if (CU_rc == -1) return -1;
   }
 
   if (device_param->is_hip == true)
   {
-    int HIP_rc;
+    if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_buf, device_param->size_pws) == -1) return -1;
 
-    HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_buf, device_param->size_pws);
+    if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_plain_bufs, device_param->size_plains) == -1) return -1;
 
-    if (HIP_rc == -1) return -1;
+    if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_digests_shown, device_param->size_shown) == -1) return -1;
 
-    HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_plain_bufs, device_param->size_plains);
+    if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_result, device_param->size_results) == -1) return -1;
 
-    if (HIP_rc == -1) return -1;
+    if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tmps, device_param->size_tmps) == -1) return -1;
-
-    HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_digests_shown, device_param->size_shown);
-
-    if (HIP_rc == -1) return -1;
-
-    HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_result, device_param->size_results);
-
-    if (HIP_rc == -1) return -1;
-
-    HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tmps, device_param->size_tmps);
-
-    if (HIP_rc == -1) return -1;
   }
 
   if (device_param->is_opencl == true)
   {
-    int CL_rc;
+    if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, device_param->size_pws) == -1) return -1;
 
-    CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, device_param->size_pws);
+    if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_plain_bufs, device_param->size_plains) == -1) return -1;
 
-    if (CL_rc == -1) return -1;
+    if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_digests_shown, device_param->size_shown) == -1) return -1;
 
-    CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_plain_bufs, device_param->size_plains);
+    if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_result, device_param->size_results) == -1) return -1;
 
-    if (CL_rc == -1) return -1;
+    if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_tmps, device_param->size_tmps) == -1) return -1;
 
-    CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_digests_shown, device_param->size_shown);
+    if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1;
-
-    if (CL_rc == -1) return -1;
-
-    CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_result, device_param->size_results);
-
-    if (CL_rc == -1) return -1;
-
-    CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_tmps, device_param->size_tmps);
-
-    if (CL_rc == -1) return -1;
   }
 
   // reset timer
@@ -480,6 +485,11 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
 
   device_param->kernel_accel   = kernel_accel;
   device_param->kernel_loops   = kernel_loops;
+  device_param->kernel_threads = kernel_threads;
+
+  const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * device_param->kernel_threads;
+
+  device_param->hardware_power = hardware_power;
 
   const u32 kernel_power = device_param->hardware_power * device_param->kernel_accel;
 

src/hashes.c

@@ -300,7 +300,7 @@ int save_hash (hashcat_ctx_t *hashcat_ctx)
   return 0;
 }
 
-void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, plain_t *plain)
+int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, plain_t *plain)
 {
   const debugfile_ctx_t *debugfile_ctx = hashcat_ctx->debugfile_ctx;
   const hashes_t        *hashes        = hashcat_ctx->hashes;
@@ -313,23 +313,62 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
   void *tmps = NULL;
 
+  cl_event opencl_event;
+  int rc;
+
   if (hashconfig->opts_type & OPTS_TYPE_COPY_TMPS)
   {
     tmps = hcmalloc (hashconfig->tmp_size);
 
     if (device_param->is_cuda == true)
     {
-      hc_cuMemcpyDtoH (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size);
+      rc = hc_cuMemcpyDtoHAsync (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size, device_param->cuda_stream);
+
+      if (rc == 0)
+      {
+        rc = hc_cuEventRecord (hashcat_ctx, device_param->cuda_event3, device_param->cuda_stream);
+      }
+
+      if (rc == -1)
+      {
+        hcfree (tmps);
+
+        return -1;
+      }
     }
 
     if (device_param->is_hip == true)
     {
-      hc_hipMemcpyDtoH (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size);
+      rc = hc_hipMemcpyDtoHAsync (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size, device_param->hip_stream);
+
+      if (rc == 0)
+      {
+        rc = hc_hipEventRecord (hashcat_ctx, device_param->hip_event3, device_param->hip_stream);
+      }
+
+      if (rc == -1)
+      {
+        hcfree (tmps);
+
+        return -1;
+      }
     }
 
     if (device_param->is_opencl == true)
     {
-      hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_TRUE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, NULL);
+      rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_FALSE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, &opencl_event);
+
+      if (rc == 0)
+      {
+        rc = hc_clFlush (hashcat_ctx, device_param->opencl_command_queue);
+      }
+
+      if (rc == -1)
+      {
+        hcfree (tmps);
+
+        return -1;
+      }
     }
   }
 
@@ -337,15 +376,14 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
   u8 *out_buf = hashes->out_buf;
 
-  int out_len = hash_encode (hashcat_ctx->hashconfig, hashcat_ctx->hashes, hashcat_ctx->module_ctx, (char *) out_buf, HCBUFSIZ_LARGE, salt_pos, digest_pos);
+  int out_len = hash_encode (hashconfig, hashes, module_ctx, (char *) out_buf, HCBUFSIZ_LARGE, salt_pos, digest_pos);
 
   out_buf[out_len] = 0;
 
   // plain
 
-  u8 plain_buf[0x1000]; // while the password itself can have only length 256, the module could encode it with something like base64 which inflates the requires buffer size
+  u8 plain_buf[HCBUFSIZ_TINY] = { 0 }; // while the password itself can have only length 256, the module could encode it with something like base64 which inflates the requires buffer size
-
+  u8 postprocess_buf[HCBUFSIZ_TINY] = { 0 };
-  memset (plain_buf, 0, sizeof (plain_buf));
 
   u8 *plain_ptr = plain_buf;
 
@@ -355,18 +393,27 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
   if (module_ctx->module_build_plain_postprocess != MODULE_DEFAULT)
   {
-    u8 temp_buf[0x1000];
+    if (hashconfig->opts_type & OPTS_TYPE_COPY_TMPS)
-
-    memset (temp_buf, 0, sizeof (temp_buf));
-
-    const int temp_len = module_ctx->module_build_plain_postprocess (hashcat_ctx->hashconfig, hashcat_ctx->hashes, tmps, (u32 *) plain_buf, sizeof (plain_buf), plain_len, (u32 *)temp_buf, sizeof (temp_buf));
-
-    if (temp_len < (int) sizeof (plain_buf))
     {
-      memcpy (plain_buf, temp_buf, temp_len);
+      if (device_param->is_cuda == true)
-
+      {
-      plain_len = temp_len;
+        if (hc_cuEventSynchronize (hashcat_ctx, device_param->cuda_event3) == -1) return -1;
       }
+
+      if (device_param->is_hip == true)
+      {
+        if (hc_hipEventSynchronize (hashcat_ctx, device_param->hip_event3) == -1) return -1;
+      }
+
+      if (device_param->is_opencl == true)
+      {
+        if (hc_clWaitForEvents (hashcat_ctx, 1, &opencl_event) == -1) return -1;
+      }
+    }
+
+    plain_len = module_ctx->module_build_plain_postprocess (hashconfig, hashes, tmps, (u32 *) plain_buf, sizeof (plain_buf), plain_len, (u32 *) postprocess_buf, sizeof (postprocess_buf));
+
+    plain_ptr = postprocess_buf;
   }
 
   // crackpos
@@ -407,6 +454,24 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
   if (module_ctx->module_hash_encode_potfile != MODULE_DEFAULT)
   {
+    if (hashconfig->opts_type & OPTS_TYPE_COPY_TMPS)
+    {
+      if (device_param->is_cuda == true)
+      {
+        if (hc_cuEventSynchronize (hashcat_ctx, device_param->cuda_event3) == -1) return -1;
+      }
+
+      if (device_param->is_hip == true)
+      {
+        if (hc_hipEventSynchronize (hashcat_ctx, device_param->hip_event3) == -1) return -1;
+      }
+
+      if (device_param->is_opencl == true)
+      {
+        if (hc_clWaitForEvents (hashcat_ctx, 1, &opencl_event) == -1) return -1;
+      }
+    }
+
     salt_t *salts_buf = hashes->salts_buf;
 
     salts_buf += salt_pos;
@@ -471,9 +536,16 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
   if (hashconfig->opts_type & OPTS_TYPE_COPY_TMPS)
   {
     hcfree (tmps);
+
+    if (device_param->is_opencl == true)
+    {
+      if (hc_clReleaseEvent (hashcat_ctx, opencl_event) == -1) return -1;
     }
   }
 
+  return 0;
+}
+
 //int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 salt_pos)
 int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
 {
@@ -484,63 +556,83 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
   user_options_t *user_options = hashcat_ctx->user_options;
 
   u32 num_cracked = 0;
-
+  int rc;
-  int CU_rc;
-  int HIP_rc;
-  int CL_rc;
 
   if (device_param->is_cuda == true)
   {
-    CU_rc = hc_cuMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32));
+    if (hc_cuMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32), device_param->cuda_stream) == -1) return -1;
 
-    if (CU_rc == -1) return -1;
+    if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1;
   }
 
   if (device_param->is_hip == true)
   {
-    HIP_rc = hc_hipMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32));
+    if (hc_hipMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32), device_param->hip_stream) == -1) return -1;
 
-    if (HIP_rc == -1) return -1;
+    if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1;
   }
 
   if (device_param->is_opencl == true)
   {
-    CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL);
+    /* blocking */
-
+    if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL) == -1) return -1;
-    if (CL_rc == -1) return -1;
   }
 
-  if (user_options->speed_only == true)
+  if (num_cracked == 0 || user_options->speed_only == true)
   {
-    // we want the hc_clEnqueueReadBuffer to run in benchmark mode because it has an influence in performance
+    // we want to get the num_cracked in benchmark mode because it has an influence in performance
     // however if the benchmark cracks the artificial hash used for benchmarks we don't want to see that!
 
     return 0;
   }
 
-  if (num_cracked)
+  plain_t *cracked = (plain_t *) hcmalloc (num_cracked * sizeof (plain_t));
-  {
-    plain_t *cracked = (plain_t *) hccalloc (num_cracked, sizeof (plain_t));
 
   if (device_param->is_cuda == true)
   {
-      CU_rc = hc_cuMemcpyDtoH (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t));
+    rc = hc_cuMemcpyDtoHAsync (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t), device_param->cuda_stream);
 
-      if (CU_rc == -1) return -1;
+    if (rc == 0)
+    {
+      rc = hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream);
+    }
+
+    if (rc == -1)
+    {
+      hcfree (cracked);
+
+      return -1;
+    }
   }
 
   if (device_param->is_hip == true)
   {
-      HIP_rc = hc_hipMemcpyDtoH (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t));
+    rc = hc_hipMemcpyDtoHAsync (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t), device_param->hip_stream);
 
-      if (HIP_rc == -1) return -1;
+    if (rc == 0)
+    {
+      rc = hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream);
+    }
+
+    if (rc == -1)
+    {
+      hcfree (cracked);
+
+      return -1;
+    }
   }
 
   if (device_param->is_opencl == true)
   {
-      CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL);
+    /* blocking */
+    rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL);
 
-      if (CL_rc == -1) return -1;
+    if (rc == -1)
+    {
+      hcfree (cracked);
+
+      return -1;
+    }
   }
 
   u32 cpt_cracked = 0;
@@ -576,7 +668,12 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
 
     if (hashes->salts_done == hashes->salts_cnt) mycracked (hashcat_ctx);
 
-      check_hash (hashcat_ctx, device_param, &cracked[i]);
+    rc = check_hash (hashcat_ctx, device_param, &cracked[i]);
+
+    if (rc == -1)
+    {
+      break;
+    }
 
     if (hashconfig->opts_type & OPTS_TYPE_PT_NEVERCRACK)
     {
@@ -584,27 +681,35 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
       // otherwise host thinks again and again the hash was cracked
       // and returns invalid password each time
 
-        memset (hashes->digests_shown_tmp, 0, salt_buf->digests_cnt * sizeof (u32));
-
       if (device_param->is_cuda == true)
       {
-          CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), &hashes->digests_shown_tmp[salt_buf->digests_offset], salt_buf->digests_cnt * sizeof (u32));
+        rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), salt_buf->digests_cnt * sizeof (u32));
 
-          if (CU_rc == -1) return -1;
+        if (rc == -1)
+        {
+          break;
+        }
       }
 
       if (device_param->is_hip == true)
       {
-          HIP_rc = hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), &hashes->digests_shown_tmp[salt_buf->digests_offset], salt_buf->digests_cnt * sizeof (u32));
+        rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), salt_buf->digests_cnt * sizeof (u32));
 
-          if (HIP_rc == -1) return -1;
+        if (rc == -1)
+        {
+          break;
+        }
       }
 
       if (device_param->is_opencl == true)
       {
-          CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (u32), salt_buf->digests_cnt * sizeof (u32), &hashes->digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
+        /* NOTE: run_opencl_kernel_bzero() does not handle buffer offset */
+        rc = run_opencl_kernel_memset32 (hashcat_ctx, device_param, device_param->opencl_d_digests_shown, salt_buf->digests_offset * sizeof (u32), 0, salt_buf->digests_cnt * sizeof (u32));
 
-          if (CL_rc == -1) return -1;
+        if (rc == -1)
+        {
+          break;
+        }
       }
     }
   }
@@ -613,6 +718,11 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
 
   hcfree (cracked);
 
+  if (rc == -1)
+  {
+    return -1;
+  }
+
   if (cpt_cracked > 0)
   {
     hc_thread_mutex_lock (status_ctx->mux_display);
@@ -629,28 +739,21 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
     hc_thread_mutex_unlock (status_ctx->mux_display);
   }
 
-    num_cracked = 0;
-
   if (device_param->is_cuda == true)
   {
-      CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_result, &num_cracked, sizeof (u32));
+    if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_result, sizeof (u32)) == -1) return -1;
-
-      if (CU_rc == -1) return -1;
   }
 
   if (device_param->is_hip == true)
   {
-      HIP_rc = hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_result, &num_cracked, sizeof (u32));
+    if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_result, sizeof (u32)) == -1) return -1;
-
-      if (HIP_rc == -1) return -1;
   }
 
   if (device_param->is_opencl == true)
   {
-      CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL);
+    if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_result, sizeof (u32)) == -1) return -1;
 
-      if (CL_rc == -1) return -1;
+    if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1;
-    }
   }
 
   return 0;
@@ -1569,7 +1672,6 @@ int hashes_init_stage2 (hashcat_ctx_t *hashcat_ctx)
   u32 digests_done = 0;
 
   u32 *digests_shown     = (u32 *) hccalloc (digests_cnt, sizeof (u32));
-  u32 *digests_shown_tmp = (u32 *) hccalloc (digests_cnt, sizeof (u32));
 
   u32 salts_cnt   = 0;
   u32 salts_done  = 0;
@@ -1706,7 +1808,6 @@ int hashes_init_stage2 (hashcat_ctx_t *hashcat_ctx)
   hashes->digests_done      = digests_done;
   hashes->digests_buf       = digests_buf_new;
   hashes->digests_shown     = digests_shown;
-  hashes->digests_shown_tmp = digests_shown_tmp;
 
   hashes->salts_cnt         = salts_cnt;
   hashes->salts_done        = salts_done;
@@ -2193,7 +2294,6 @@ void hashes_destroy (hashcat_ctx_t *hashcat_ctx)
 
   hcfree (hashes->digests_buf);
   hcfree (hashes->digests_shown);
-  hcfree (hashes->digests_shown_tmp);
 
   hcfree (hashes->salts_buf);
   hcfree (hashes->salts_shown);

src/interface.c

@@ -150,6 +150,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
   CHECK_DEFINED (module_ctx->module_esalt_size);
   CHECK_DEFINED (module_ctx->module_extra_buffer_size);
   CHECK_DEFINED (module_ctx->module_extra_tmp_size);
+  CHECK_DEFINED (module_ctx->module_extra_tuningdb_block);
   CHECK_DEFINED (module_ctx->module_forced_outfile_format);
   CHECK_DEFINED (module_ctx->module_hash_binary_count);
   CHECK_DEFINED (module_ctx->module_hash_binary_parse);

src/modules/module_00000.c

@@ -138,6 +138,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00010.c

@@ -161,6 +161,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00011.c

@@ -164,6 +164,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00012.c

@@ -156,6 +156,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00020.c

@@ -160,6 +160,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00021.c

@@ -155,6 +155,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00022.c

@@ -268,6 +268,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00023.c

@@ -170,6 +170,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00024.c

@@ -154,6 +154,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00030.c

@@ -162,6 +162,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00040.c

@@ -161,6 +161,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00050.c

@@ -140,6 +140,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00060.c

@@ -140,6 +140,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00070.c

@@ -136,6 +136,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00100.c

@@ -154,6 +154,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00101.c

@@ -158,6 +158,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00110.c

@@ -180,6 +180,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00111.c

@@ -184,6 +184,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00112.c

@@ -181,6 +181,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00120.c

@@ -180,6 +180,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00121.c

@@ -181,6 +181,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00122.c

@@ -160,6 +160,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00124.c

@@ -179,6 +179,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00125.c

@@ -162,6 +162,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00130.c

@@ -181,6 +181,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00131.c

@@ -176,6 +176,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00132.c

@@ -170,6 +170,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00133.c

@@ -153,6 +153,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00140.c

@@ -181,6 +181,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00141.c

@@ -188,6 +188,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00150.c

@@ -158,6 +158,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00160.c

@@ -158,6 +158,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00170.c

@@ -155,6 +155,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00200.c

@@ -122,6 +122,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00300.c

@@ -153,6 +153,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00400.c

@@ -249,6 +249,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00500.c

@@ -262,6 +262,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00501.c

@@ -341,6 +341,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00600.c

@@ -127,6 +127,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_00900.c

@@ -138,6 +138,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01000.c

@@ -141,6 +141,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01100.c

@@ -165,6 +165,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01300.c

@@ -168,6 +168,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01400.c

@@ -175,6 +175,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01410.c

@@ -201,6 +201,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01411.c

@@ -197,6 +197,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01420.c

@@ -201,6 +201,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01421.c

@@ -177,6 +177,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01430.c

@@ -202,6 +202,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01440.c

@@ -202,6 +202,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01441.c

@@ -203,6 +203,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01450.c

@@ -172,6 +172,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01460.c

@@ -204,6 +204,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01470.c

@@ -176,6 +176,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01500.c

@@ -95,13 +95,6 @@ int module_build_plain_postprocess (MAYBE_UNUSED const hashconfig_t *hashconfig,
   return src_len;
 }
 
-u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
-{
-  const u32 kernel_threads_max = 64; // performance only optimization
-
-  return kernel_threads_max;
-}
-
 u32 module_kernel_loops_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
   u32 kernel_loops_max = KERNEL_LOOPS_MAX;
@@ -315,6 +308,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;
@@ -345,7 +339,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;
   module_ctx->module_kernel_loops_max         = module_kernel_loops_max;
   module_ctx->module_kernel_loops_min         = module_kernel_loops_min;
-  module_ctx->module_kernel_threads_max       = module_kernel_threads_max;
+  module_ctx->module_kernel_threads_max       = MODULE_DEFAULT;
   module_ctx->module_kernel_threads_min       = MODULE_DEFAULT;
   module_ctx->module_kern_type                = module_kern_type;
   module_ctx->module_kern_type_dynamic        = MODULE_DEFAULT;

src/modules/module_01600.c

@@ -262,6 +262,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01700.c

@@ -202,6 +202,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01710.c

@@ -202,6 +202,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01711.c

@@ -198,6 +198,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01720.c

@@ -227,6 +227,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01722.c

@@ -204,6 +204,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01730.c

@@ -203,6 +203,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01731.c

@@ -189,6 +189,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01740.c

@@ -203,6 +203,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01750.c

@@ -173,6 +173,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01760.c

@@ -174,6 +174,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01770.c

@@ -191,6 +191,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;

src/modules/module_01800.c

@@ -567,6 +567,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_esalt_size               = MODULE_DEFAULT;
   module_ctx->module_extra_buffer_size        = MODULE_DEFAULT;
   module_ctx->module_extra_tmp_size           = MODULE_DEFAULT;
+  module_ctx->module_extra_tuningdb_block     = MODULE_DEFAULT;
   module_ctx->module_forced_outfile_format    = MODULE_DEFAULT;
   module_ctx->module_hash_binary_count        = MODULE_DEFAULT;
   module_ctx->module_hash_binary_parse        = MODULE_DEFAULT;