Add hashconfig_enforce_kernel_loops ()

include/interface.h

@@ -1537,13 +1537,14 @@ void to_hccap_t (hccap_t *hccap, const uint salt_pos, const uint digest_pos, con
 
 void ascii_digest (char *out_buf, const uint salt_pos, const uint digest_pos, const hashconfig_t *hashconfig, const hashes_t *hashes);
 
-int   hashconfig_init                (hashconfig_t *hashconfig, const user_options_t *user_options);
-void  hashconfig_destroy             (hashconfig_t *hashconfig);
-u32   hashconfig_kernel_thread_force (const hashconfig_t *hashconfig, const hc_device_param_t *device_param);
-uint  hashconfig_general_pw_min      (hashconfig_t *hashconfig);
-uint  hashconfig_general_pw_max      (hashconfig_t *hashconfig);
-void  hashconfig_general_defaults    (hashconfig_t *hashconfig, hashes_t *hashes, const user_options_t *user_options);
-void  hashconfig_benchmark_defaults  (const hashconfig_t *hashconfig, salt_t *salt, void *esalt);
-char *hashconfig_benchmark_mask      (const hashconfig_t *hashconfig);
+int   hashconfig_init                   (hashconfig_t *hashconfig, const user_options_t *user_options);
+void  hashconfig_destroy                (hashconfig_t *hashconfig);
+u32   hashconfig_enforce_kernel_threads (const hashconfig_t *hashconfig, const hc_device_param_t *device_param);
+u32   hashconfig_enforce_kernel_loops   (const hashconfig_t *hashconfig, const user_options_t *user_options);
+uint  hashconfig_general_pw_min         (hashconfig_t *hashconfig);
+uint  hashconfig_general_pw_max         (hashconfig_t *hashconfig);
+void  hashconfig_general_defaults       (hashconfig_t *hashconfig, hashes_t *hashes, const user_options_t *user_options);
+void  hashconfig_benchmark_defaults     (const hashconfig_t *hashconfig, salt_t *salt, void *esalt);
+char *hashconfig_benchmark_mask         (const hashconfig_t *hashconfig);
 
 #endif // _INTERFACE_H

src/interface.c

@@ -19919,7 +19919,7 @@ void hashconfig_destroy (hashconfig_t *hashconfig)
   myfree (hashconfig);
 }
 
-u32 hashconfig_kernel_thread_force (const hashconfig_t *hashconfig, const hc_device_param_t *device_param)
+u32 hashconfig_enforce_kernel_threads (const hashconfig_t *hashconfig, const hc_device_param_t *device_param)
 {
   u32 kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size);
 
@@ -19952,6 +19952,48 @@ u32 hashconfig_kernel_thread_force (const hashconfig_t *hashconfig, const hc_dev
   return kernel_threads;
 }
 
+u32 hashconfig_enforce_kernel_loops (const hashconfig_t *hashconfig, const user_options_t *user_options)
+{
+  u32 kernel_loops_fixed = 0;
+
+  if (hashconfig->hash_mode == 1500 && user_options->attack_mode == ATTACK_MODE_BF)
+  {
+    kernel_loops_fixed = 1024;
+  }
+
+  if (hashconfig->hash_mode == 3000 && user_options->attack_mode == ATTACK_MODE_BF)
+  {
+    kernel_loops_fixed = 1024;
+  }
+
+  if (hashconfig->hash_mode == 8900)
+  {
+    kernel_loops_fixed = 1;
+  }
+
+  if (hashconfig->hash_mode == 9300)
+  {
+    kernel_loops_fixed = 1;
+  }
+
+  if (hashconfig->hash_mode == 12500)
+  {
+    kernel_loops_fixed = ROUNDS_RAR3 / 16;
+  }
+
+  if (hashconfig->hash_mode == 14000 && user_options->attack_mode == ATTACK_MODE_BF)
+  {
+    kernel_loops_fixed = 1024;
+  }
+
+  if (hashconfig->hash_mode == 14100 && user_options->attack_mode == ATTACK_MODE_BF)
+  {
+    kernel_loops_fixed = 1024;
+  }
+
+  return kernel_loops_fixed;
+}
+
 uint hashconfig_general_pw_min (hashconfig_t *hashconfig)
 {
   uint pw_min = PW_MIN;

src/opencl.c

@@ -2641,7 +2641,7 @@ int opencl_session_begin (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconf
      *                 there needs to be some upper limit, otherwise there's too much overhead
      */
 
-    u32 kernel_threads = hashconfig_kernel_thread_force (hashconfig, device_param);
+    u32 kernel_threads = hashconfig_enforce_kernel_threads (hashconfig, device_param);
 
     device_param->kernel_threads = kernel_threads;
 
@@ -2795,58 +2795,10 @@ int opencl_session_begin (opencl_ctx_t *opencl_ctx, const hashconfig_t *hashconf
      * some algorithms need a fixed kernel-loops count
      */
 
-    if (hashconfig->hash_mode == 1500 && user_options->attack_mode == ATTACK_MODE_BF)
-    {
-      const u32 kernel_loops_fixed = 1024;
-
-      device_param->kernel_loops_min = kernel_loops_fixed;
-      device_param->kernel_loops_max = kernel_loops_fixed;
-    }
-
-    if (hashconfig->hash_mode == 3000 && user_options->attack_mode == ATTACK_MODE_BF)
-    {
-      const u32 kernel_loops_fixed = 1024;
-
-      device_param->kernel_loops_min = kernel_loops_fixed;
-      device_param->kernel_loops_max = kernel_loops_fixed;
-    }
-
-    if (hashconfig->hash_mode == 8900)
-    {
-      const u32 kernel_loops_fixed = 1;
-
-      device_param->kernel_loops_min = kernel_loops_fixed;
-      device_param->kernel_loops_max = kernel_loops_fixed;
-    }
-
-    if (hashconfig->hash_mode == 9300)
-    {
-      const u32 kernel_loops_fixed = 1;
-
-      device_param->kernel_loops_min = kernel_loops_fixed;
-      device_param->kernel_loops_max = kernel_loops_fixed;
-    }
-
-    if (hashconfig->hash_mode == 12500)
-    {
-      const u32 kernel_loops_fixed = ROUNDS_RAR3 / 16;
-
-      device_param->kernel_loops_min = kernel_loops_fixed;
-      device_param->kernel_loops_max = kernel_loops_fixed;
-    }
+    const u32 kernel_loops_fixed = hashconfig_enforce_kernel_loops (hashconfig, user_options);
 
-    if (hashconfig->hash_mode == 14000 && user_options->attack_mode == ATTACK_MODE_BF)
+    if (kernel_loops_fixed != 0)
     {
-      const u32 kernel_loops_fixed = 1024;
-
-      device_param->kernel_loops_min = kernel_loops_fixed;
-      device_param->kernel_loops_max = kernel_loops_fixed;
-    }
-
-    if (hashconfig->hash_mode == 14100 && user_options->attack_mode == ATTACK_MODE_BF)
-    {
-      const u32 kernel_loops_fixed = 1024;
-
       device_param->kernel_loops_min = kernel_loops_fixed;
       device_param->kernel_loops_max = kernel_loops_fixed;
     }