Improve the logic for when to use funnelshift and when not to. Some algorithms, such as SHA1-HMAC and DCC1, do not work well with it, so it has been disabled for them.

Fix the automatic reduction of the kernel-accel maximum based on available memory per device by accounting for the additional size needed to handle register spilling.
Fix the tools/benchmark_deep.pl script to recognize benchmark masks more reliably.

OpenCL/inc_platform.cl

@@ -128,108 +128,128 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
 DECLSPEC u32x rotl32 (const u32x a, const int n)
 {
   #if VECT_SIZE == 1
-  return __funnelshift_l(a, a, n);
+
-  #endif
+  return rotl32_S (a, n);
+
+  #else
 
   u32x t = 0;
 
   #if VECT_SIZE >= 2
-  t.s0 = __funnelshift_l(a.s0, a.s0, n);
+  t.s0 = rotl32_S (a.s0, n);
-  t.s1 = __funnelshift_l(a.s1, a.s1, n);
+  t.s1 = rotl32_S (a.s1, n);
   #endif
 
   #if VECT_SIZE >= 4
-  t.s2 = __funnelshift_l(a.s2, a.s2, n);
+  t.s2 = rotl32_S (a.s2, n);
-  t.s3 = __funnelshift_l(a.s3, a.s3, n);
+  t.s3 = rotl32_S (a.s3, n);
   #endif
 
   #if VECT_SIZE >= 8
-  t.s4 = __funnelshift_l(a.s4, a.s4, n);
+  t.s4 = rotl32_S (a.s4, n);
-  t.s5 = __funnelshift_l(a.s5, a.s5, n);
+  t.s5 = rotl32_S (a.s5, n);
-  t.s6 = __funnelshift_l(a.s6, a.s6, n);
+  t.s6 = rotl32_S (a.s6, n);
-  t.s7 = __funnelshift_l(a.s7, a.s7, n);
+  t.s7 = rotl32_S (a.s7, n);
   #endif
 
   #if VECT_SIZE >= 16
-  t.s8 = __funnelshift_l(a.s8, a.s8, n);
+  t.s8 = rotl32_S (a.s8, n);
-  t.s9 = __funnelshift_l(a.s9, a.s9, n);
+  t.s9 = rotl32_S (a.s9, n);
-  t.sa = __funnelshift_l(a.sa, a.sa, n);
+  t.sa = rotl32_S (a.sa, n);
-  t.sb = __funnelshift_l(a.sb, a.sb, n);
+  t.sb = rotl32_S (a.sb, n);
-  t.sc = __funnelshift_l(a.sc, a.sc, n);
+  t.sc = rotl32_S (a.sc, n);
-  t.sd = __funnelshift_l(a.sd, a.sd, n);
+  t.sd = rotl32_S (a.sd, n);
-  t.se = __funnelshift_l(a.se, a.se, n);
+  t.se = rotl32_S (a.se, n);
-  t.sf = __funnelshift_l(a.sf, a.sf, n);
+  t.sf = rotl32_S (a.sf, n);
   #endif
 
   return t;
-  //return ((a << n) | ((a >> (32 - n))));
+
+  #endif
 }
 
 DECLSPEC u32x rotr32 (const u32x a, const int n)
 {
   #if VECT_SIZE == 1
-  return __funnelshift_r(a, a, n);
+
-  #endif
+  return rotr32_S (a, n);
+
+  #else
 
   u32x t = 0;
 
   #if VECT_SIZE >= 2
-  t.s0 = __funnelshift_r(a.s0, a.s0, n);
+  t.s0 = rotr32_S (a.s0, n);
-  t.s1 = __funnelshift_r(a.s1, a.s1, n);
+  t.s1 = rotr32_S (a.s1, n);
   #endif
 
   #if VECT_SIZE >= 4
-  t.s2 = __funnelshift_r(a.s2, a.s2, n);
+  t.s2 = rotr32_S (a.s2, n);
-  t.s3 = __funnelshift_r(a.s3, a.s3, n);
+  t.s3 = rotr32_S (a.s3, n);
   #endif
 
   #if VECT_SIZE >= 8
-  t.s4 = __funnelshift_r(a.s4, a.s4, n);
+  t.s4 = rotr32_S (a.s4, n);
-  t.s5 = __funnelshift_r(a.s5, a.s5, n);
+  t.s5 = rotr32_S (a.s5, n);
-  t.s6 = __funnelshift_r(a.s6, a.s6, n);
+  t.s6 = rotr32_S (a.s6, n);
-  t.s7 = __funnelshift_r(a.s7, a.s7, n);
+  t.s7 = rotr32_S (a.s7, n);
   #endif
 
   #if VECT_SIZE >= 16
-  t.s8 = __funnelshift_r(a.s8, a.s8, n);
+  t.s8 = rotr32_S (a.s8, n);
-  t.s9 = __funnelshift_r(a.s9, a.s9, n);
+  t.s9 = rotr32_S (a.s9, n);
-  t.sa = __funnelshift_r(a.sa, a.sa, n);
+  t.sa = rotr32_S (a.sa, n);
-  t.sb = __funnelshift_r(a.sb, a.sb, n);
+  t.sb = rotr32_S (a.sb, n);
-  t.sc = __funnelshift_r(a.sc, a.sc, n);
+  t.sc = rotr32_S (a.sc, n);
-  t.sd = __funnelshift_r(a.sd, a.sd, n);
+  t.sd = rotr32_S (a.sd, n);
-  t.se = __funnelshift_r(a.se, a.se, n);
+  t.se = rotr32_S (a.se, n);
-  t.sf = __funnelshift_r(a.sf, a.sf, n);
+  t.sf = rotr32_S (a.sf, n);
   #endif
 
   return t;
-  //return ((a >> n) | ((a << (32 - n))));
+
+  #endif
 }
 
 DECLSPEC u32 rotl32_S (const u32 a, const int n)
 {
+  #ifdef USE_FUNNELSHIFT
   return __funnelshift_l (a, a, n);
-  //return ((a << n) | ((a >> (32 - n))));
+  #else
+  return ((a << n) | ((a >> (32 - n))));
+  #endif
 }
 
 DECLSPEC u32 rotr32_S (const u32 a, const int n)
 {
+  #ifdef USE_FUNNELSHIFT
   return __funnelshift_r (a, a, n);
-  //return ((a >> n) | ((a << (32 - n))));
+  #else
+  return ((a >> n) | ((a << (32 - n))));
+  #endif
 }
 
 DECLSPEC u64x rotl64 (const u64x a, const int n)
 {
+  #if VECT_SIZE == 1
+  return rotl64_S (a, n);
+  #else
   return ((a << n) | ((a >> (64 - n))));
+  #endif
 }
 
 DECLSPEC u64x rotr64 (const u64x a, const int n)
 {
+  #if VECT_SIZE == 1
+  return rotr64_S (a, n);
+  #else
   return ((a >> n) | ((a << (64 - n))));
+  #endif
 }
 
 DECLSPEC u64 rotl64_S (const u64 a, const int n)
 {
-  return ((a << n) | ((a >> (64 - n))));
+  return rotr64_S (a, 64 - n);
 }
 
 DECLSPEC u64 rotr64_S (const u64 a, const int n)
@@ -309,93 +329,105 @@ DECLSPEC size_t get_local_size (const u32 dimindx __attribute__((unused)))
 DECLSPEC u32x rotl32 (const u32x a, const int n)
 {
   #if VECT_SIZE == 1
-  return __funnelshift_l(a, a, n);
-  #endif
 
-  u32x t;
+  return rotl32_S (a, n);
+
+  #else
+
+  u32x t = 0;
 
   #if VECT_SIZE >= 2
-  t.s0 = __funnelshift_l(a.s0, a.s0, n);
+  t.s0 = rotl32_S (a.s0, n);
-  t.s1 = __funnelshift_l(a.s1, a.s1, n);
+  t.s1 = rotl32_S (a.s1, n);
   #endif
 
   #if VECT_SIZE >= 4
-  t.s2 = __funnelshift_l(a.s2, a.s2, n);
+  t.s2 = rotl32_S (a.s2, n);
-  t.s3 = __funnelshift_l(a.s3, a.s3, n);
+  t.s3 = rotl32_S (a.s3, n);
   #endif
 
   #if VECT_SIZE >= 8
-  t.s4 = __funnelshift_l(a.s4, a.s4, n);
+  t.s4 = rotl32_S (a.s4, n);
-  t.s5 = __funnelshift_l(a.s5, a.s5, n);
+  t.s5 = rotl32_S (a.s5, n);
-  t.s6 = __funnelshift_l(a.s6, a.s6, n);
+  t.s6 = rotl32_S (a.s6, n);
-  t.s7 = __funnelshift_l(a.s7, a.s7, n);
+  t.s7 = rotl32_S (a.s7, n);
   #endif
 
   #if VECT_SIZE >= 16
-  t.s8 = __funnelshift_l(a.s8, a.s8, n);
+  t.s8 = rotl32_S (a.s8, n);
-  t.s9 = __funnelshift_l(a.s9, a.s9, n);
+  t.s9 = rotl32_S (a.s9, n);
-  t.sa = __funnelshift_l(a.sa, a.sa, n);
+  t.sa = rotl32_S (a.sa, n);
-  t.sb = __funnelshift_l(a.sb, a.sb, n);
+  t.sb = rotl32_S (a.sb, n);
-  t.sc = __funnelshift_l(a.sc, a.sc, n);
+  t.sc = rotl32_S (a.sc, n);
-  t.sd = __funnelshift_l(a.sd, a.sd, n);
+  t.sd = rotl32_S (a.sd, n);
-  t.se = __funnelshift_l(a.se, a.se, n);
+  t.se = rotl32_S (a.se, n);
-  t.sf = __funnelshift_l(a.sf, a.sf, n);
+  t.sf = rotl32_S (a.sf, n);
   #endif
 
   return t;
-  //return ((a << n) | ((a >> (32 - n))));
+
+  #endif
 }
 
 DECLSPEC u32x rotr32 (const u32x a, const int n)
 {
   #if VECT_SIZE == 1
-  return __funnelshift_r(a, a, n);
-  #endif
 
-  u32x t;
+  return rotr32_S (a, n);
+
+  #else
+
+  u32x t = 0;
 
   #if VECT_SIZE >= 2
-  t.s0 = __funnelshift_r(a.s0, a.s0, n);
+  t.s0 = rotr32_S (a.s0, n);
-  t.s1 = __funnelshift_r(a.s1, a.s1, n);
+  t.s1 = rotr32_S (a.s1, n);
   #endif
 
   #if VECT_SIZE >= 4
-  t.s2 = __funnelshift_r(a.s2, a.s2, n);
+  t.s2 = rotr32_S (a.s2, n);
-  t.s3 = __funnelshift_r(a.s3, a.s3, n);
+  t.s3 = rotr32_S (a.s3, n);
   #endif
 
   #if VECT_SIZE >= 8
-  t.s4 = __funnelshift_r(a.s4, a.s4, n);
+  t.s4 = rotr32_S (a.s4, n);
-  t.s5 = __funnelshift_r(a.s5, a.s5, n);
+  t.s5 = rotr32_S (a.s5, n);
-  t.s6 = __funnelshift_r(a.s6, a.s6, n);
+  t.s6 = rotr32_S (a.s6, n);
-  t.s7 = __funnelshift_r(a.s7, a.s7, n);
+  t.s7 = rotr32_S (a.s7, n);
   #endif
 
   #if VECT_SIZE >= 16
-  t.s8 = __funnelshift_r(a.s8, a.s8, n);
+  t.s8 = rotr32_S (a.s8, n);
-  t.s9 = __funnelshift_r(a.s9, a.s9, n);
+  t.s9 = rotr32_S (a.s9, n);
-  t.sa = __funnelshift_r(a.sa, a.sa, n);
+  t.sa = rotr32_S (a.sa, n);
-  t.sb = __funnelshift_r(a.sb, a.sb, n);
+  t.sb = rotr32_S (a.sb, n);
-  t.sc = __funnelshift_r(a.sc, a.sc, n);
+  t.sc = rotr32_S (a.sc, n);
-  t.sd = __funnelshift_r(a.sd, a.sd, n);
+  t.sd = rotr32_S (a.sd, n);
-  t.se = __funnelshift_r(a.se, a.se, n);
+  t.se = rotr32_S (a.se, n);
-  t.sf = __funnelshift_r(a.sf, a.sf, n);
+  t.sf = rotr32_S (a.sf, n);
   #endif
 
   return t;
-  //return ((a >> n) | ((a << (32 - n))));
+
+  #endif
 }
 
 DECLSPEC u32 rotl32_S (const u32 a, const int n)
 {
+  #ifdef USE_FUNNELSHIFT
   return __funnelshift_l (a, a, n);
-  //return ((a << n) | ((a >> (32 - n))));
+  #else
+  return ((a << n) | ((a >> (32 - n))));
+  #endif
 }
 
 DECLSPEC u32 rotr32_S (const u32 a, const int n)
 {
+  #ifdef USE_FUNNELSHIFT
   return __funnelshift_r (a, a, n);
-  //return ((a >> n) | ((a << (32 - n))));
+  #else
+  return ((a >> n) | ((a << (32 - n))));
+  #endif
 }
 
 DECLSPEC u64x rotl64 (const u64x a, const int n)
@@ -423,6 +455,7 @@ DECLSPEC u64 rotl64_S (const u64 a, const int n)
 
 DECLSPEC u64 rotr64_S (const u64 a, const int n)
 {
+  #ifdef USE_FUNNELSHIFT
   vconv64_t in;
 
   in.v64 = a;
@@ -444,6 +477,9 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
   }
 
   return out.v64;
+  #else
+  return ((a >> n) | ((a << (64 - n))));
+  #endif
 }
 
 #define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0)

OpenCL/inc_vendor.h

@@ -183,11 +183,13 @@ using namespace metal;
 #ifdef IS_CUDA
 #define USE_BITSELECT
 #define USE_ROTATE
+#define USE_FUNNELSHIFT
 #endif
 
 #ifdef IS_HIP
 #define USE_BITSELECT
 #define USE_ROTATE
+#define USE_FUNNELSHIFT
 #endif
 
 #ifdef IS_ROCM
@@ -218,4 +220,9 @@ using namespace metal;
 #define s3 w
 #endif
 
+// some algorithms do not like this, eg 150, 1100, ...
+#ifdef NO_FUNNELSHIFT
+#undef USE_FUNNELSHIFT
+#endif
+
 #endif // INC_VENDOR_H

src/backend.c

@@ -15769,6 +15769,19 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
     {
       const u64 kernel_power_max = hardware_power_max * kernel_accel_max;
 
+      // size_spilling: we cannot query this directly.
+      // Example output:
+      //   ptxas         .     4096 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
+      // However, this is very relevant for us. In theory, these numbers could reach gigabytes,
+      // but in practice, excessive spilling would make the kernel too slow,
+      // and the kernel developer would adapt accordingly. We'll assume a maximum spilling buffer
+      // size of 4 KiB per thread for now.
+      // This setting will reduce the available memory pool on a 4090:
+      //   4 * 1024 * 128 * 32 = 16 MiB per -n accel,
+      // which adds up to 2 GiB with -n 128.
+
+      size_t size_spilling = kernel_power_max * (4 * 1024);
+
       // size_pws
 
       size_pws = kernel_power_max * sizeof (pw_t);
@@ -15904,7 +15917,8 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx)
         + size_st_digests
         + size_st_salts
         + size_st_esalts
-        + size_kernel_params;
+        + size_kernel_params
+        + size_spilling;
 
       if ((size_total + EXTRA_SPACE) > device_param->device_available_mem) memory_limit_hit = 1;
 

src/modules/module_00150.c

@@ -44,6 +44,15 @@ u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig,
 const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
 const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  hc_asprintf (&jit_build_options, "-D NO_FUNNELSHIFT");
+
+  return jit_build_options;
+}
+
 int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
 {
   u32 *digest = (u32 *) digest_buf;
@@ -190,7 +199,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_01100.c

@@ -50,6 +50,15 @@ u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig,
 const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
 const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  hc_asprintf (&jit_build_options, "-D NO_FUNNELSHIFT");
+
+  return jit_build_options;
+}
+
 int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
 {
   u32 *digest = (u32 *) digest_buf;
@@ -197,7 +206,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_06000.c

@@ -42,6 +42,15 @@ u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig,
 const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
 const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  hc_asprintf (&jit_build_options, "-D NO_FUNNELSHIFT");
+
+  return jit_build_options;
+}
+
 int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
 {
   u32 *digest = (u32 *) digest_buf;
@@ -150,7 +159,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_06211.c

@@ -77,6 +77,15 @@ typedef struct tc
 static const int   ROUNDS_TRUECRYPT_2K         = 2000;
 static const float MIN_SUFFICIENT_ENTROPY_FILE = 7.0f;
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  hc_asprintf (&jit_build_options, "-D NO_FUNNELSHIFT");
+
+  return jit_build_options;
+}
+
 bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hc_device_param_t *device_param)
 {
   if ((device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE) && (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU))
@@ -298,7 +307,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_06212.c

@@ -77,6 +77,15 @@ typedef struct tc
 static const int   ROUNDS_TRUECRYPT_2K         = 2000;
 static const float MIN_SUFFICIENT_ENTROPY_FILE = 7.0f;
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  hc_asprintf (&jit_build_options, "-D NO_FUNNELSHIFT");
+
+  return jit_build_options;
+}
+
 bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hc_device_param_t *device_param)
 {
   if ((device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE) && (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU))
@@ -298,7 +307,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_06213.c

@@ -77,6 +77,15 @@ typedef struct tc
 static const int   ROUNDS_TRUECRYPT_2K         = 2000;
 static const float MIN_SUFFICIENT_ENTROPY_FILE = 7.0f;
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  hc_asprintf (&jit_build_options, "-D NO_FUNNELSHIFT");
+
+  return jit_build_options;
+}
+
 bool module_unstable_warning (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hc_device_param_t *device_param)
 {
   if ((device_param->opencl_platform_vendor_id == VENDOR_ID_APPLE) && (device_param->opencl_device_type & CL_DEVICE_TYPE_GPU))
@@ -304,7 +313,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_08600.c

@@ -44,6 +44,18 @@ u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig,
 const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
 const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  if (device_param->is_cuda == true)
+  {
+    hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size);
+  }
+
+  return jit_build_options;
+}
+
 u32 module_pw_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 pw_max = 16; // Lotus Notes/Domino 5 limits itself to 16
@@ -146,7 +158,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_08700.c

@@ -43,6 +43,18 @@ u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig,
 const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
 const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  if (device_param->is_cuda == true)
+  {
+    hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size);
+  }
+
+  return jit_build_options;
+}
+
 u32 module_pw_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 bool optimized_kernel = (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL);
@@ -188,7 +200,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_09710.c

@@ -28,6 +28,7 @@ static const u64   OPTS_TYPE      = OPTS_TYPE_STOCK_MODULE
                                   | OPTS_TYPE_PT_ALWAYS_HEXIFY
                                   | OPTS_TYPE_AUTODETECT_DISABLE;
 static const u32   SALT_TYPE      = SALT_TYPE_EMBEDDED;
+static const char *BENCHMARK_MASK = "?b?b?b?b?b";
 static const char *ST_PASS        = "\x91\xb2\xe0\x62\xb9";
 static const char *ST_HASH        = "$oldoffice$0*55045061647456688860411218030058*e7e24d163fbd743992d4b8892bf3f2f7*493410dbc832557d3fe1870ace8397e2";
 
@@ -134,7 +135,7 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con
 
 const char *module_benchmark_mask (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
-  return "?b?b?b?b?b";
+  return BENCHMARK_MASK;
 }
 
 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)

src/modules/module_09810.c

@@ -27,6 +27,7 @@ static const u64   OPTS_TYPE      = OPTS_TYPE_STOCK_MODULE
                                   | OPTS_TYPE_PT_ALWAYS_HEXIFY
                                   | OPTS_TYPE_AUTODETECT_DISABLE;
 static const u32   SALT_TYPE      = SALT_TYPE_EMBEDDED;
+static const char *BENCHMARK_MASK = "?b?b?b?b?b";
 static const char *ST_PASS        = "\xb8\xf6\x36\x19\xca";
 static const char *ST_HASH        = "$oldoffice$3*83328705222323020515404251156288*2855956a165ff6511bc7f4cd77b9e101*941861655e73a09c40f7b1e9dfd0c256ed285acd";
 
@@ -135,7 +136,7 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con
 
 const char *module_benchmark_mask (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
-  return "?b?b?b?b?b";
+  return BENCHMARK_MASK;
 }
 
 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)

src/modules/module_10410.c

@@ -26,6 +26,7 @@ static const u64   OPTS_TYPE      = OPTS_TYPE_STOCK_MODULE
                                   | OPTS_TYPE_PT_ALWAYS_HEXIFY
                                   | OPTS_TYPE_AUTODETECT_DISABLE;
 static const u32   SALT_TYPE      = SALT_TYPE_EMBEDDED;
+static const char *BENCHMARK_MASK = "?b?b?b?b?b";
 static const char *ST_PASS        = "\x6a\x8a\xed\xcc\xb7";
 static const char *ST_HASH        = "$pdf$1*2*40*-1*0*16*01221086741440841668371056103222*32*27c3fecef6d46a78eb61b8b4dbc690f5f8a2912bbb9afc842c12d79481568b74*32*0000000000000000000000000000000000000000000000000000000000000000";
 
@@ -144,7 +145,7 @@ u32 module_pw_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con
 
 const char *module_benchmark_mask (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
-  return "?b?b?b?b?b";
+  return BENCHMARK_MASK;
 }
 
 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)

src/modules/module_11000.c

@@ -45,6 +45,15 @@ u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig,
 const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
 const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  hc_asprintf (&jit_build_options, "-D NO_FUNNELSHIFT");
+
+  return jit_build_options;
+}
+
 u32 module_salt_min (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 salt_min = 56;
@@ -181,7 +190,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_14400.c

@@ -43,6 +43,15 @@ u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig,
 const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
 const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  hc_asprintf (&jit_build_options, "-D NO_FUNNELSHIFT");
+
+  return jit_build_options;
+}
+
 u32 module_pw_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 bool optimized_kernel = (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL);
@@ -188,7 +197,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_14900.c

@@ -24,6 +24,7 @@ static const u64   OPTS_TYPE      = OPTS_TYPE_STOCK_MODULE
                                   | OPTS_TYPE_PT_GENERATE_LE
                                   | OPTS_TYPE_SUGGEST_KG;
 static const u32   SALT_TYPE      = SALT_TYPE_EMBEDDED;
+static const char *BENCHMARK_MASK = "?b?b?b?b?bxxxxx";
 static const char *ST_PASS        = "hashcat!!!";
 static const char *ST_HASH        = "7090b6b9:04223875";
 
@@ -58,7 +59,7 @@ u32 module_pw_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED con
 
 const char *module_benchmark_mask (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
 {
-  return "?b?b?b?b?bxxxxx";
+  return BENCHMARK_MASK;
 }
 
 int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)

src/modules/module_16000.c

@@ -42,6 +42,19 @@ u32         module_salt_type      (MAYBE_UNUSED const hashconfig_t *hashconfig,
 const char *module_st_hash        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH;         }
 const char *module_st_pass        (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS;         }
 
+
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  if (device_param->is_cuda == true)
+  {
+    hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size);
+  }
+
+  return jit_build_options;
+}
+
 u32 module_pw_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 pw_max = 8; // Underlaying DES max
@@ -166,7 +179,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_16600.c

@@ -55,6 +55,18 @@ typedef struct electrum_wallet
 
 static const char *SIGNATURE_ELECTRUM_WALLET = "$electrum$";
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  if (device_param->is_cuda == true)
+  {
+    hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size);
+  }
+
+  return jit_build_options;
+}
+
 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 u64 esalt_size = (const u64) sizeof (electrum_wallet_t);
@@ -232,7 +244,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

src/modules/module_20510.c

@@ -108,6 +108,18 @@ typedef struct pkzip_extra
 
 } pkzip_extra_t;
 
+char *module_jit_build_options (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra, MAYBE_UNUSED const hashes_t *hashes, MAYBE_UNUSED const hc_device_param_t *device_param)
+{
+  char *jit_build_options = NULL;
+
+  if (device_param->is_cuda == true)
+  {
+    hc_asprintf (&jit_build_options, "-D FIXED_LOCAL_SIZE=%u", (u32) device_param->device_maxworkgroup_size);
+  }
+
+  return jit_build_options;
+}
+
 u32 module_hashes_count_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 u64 tmp_size = 1;
@@ -243,7 +255,7 @@ void module_init (module_ctx_t *module_ctx)
   module_ctx->module_hook23                   = MODULE_DEFAULT;
   module_ctx->module_hook_salt_size           = MODULE_DEFAULT;
   module_ctx->module_hook_size                = MODULE_DEFAULT;
-  module_ctx->module_jit_build_options        = MODULE_DEFAULT;
+  module_ctx->module_jit_build_options        = module_jit_build_options;
   module_ctx->module_jit_cache_disable        = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_max         = MODULE_DEFAULT;
   module_ctx->module_kernel_accel_min         = MODULE_DEFAULT;

tools/benchmark_deep.pl

@@ -548,7 +548,7 @@ sub get_module
 
   close (IN);
 
-  my $mask = (defined $benchmark_mask) ? $benchmark_mask : $default_mask;
+  my $mask = $default_mask;
 
   if ($pw_min != -1)
   {
@@ -571,6 +571,8 @@ sub get_module
     }
   }
 
+  $mask = (defined $benchmark_mask) ? $benchmark_mask : $mask;
+
   my $module =
   {
     "is_binary" => $is_binary,