diff --git a/include/interface.h b/include/interface.h index 3da17c2e4..27bf7d3c3 100644 --- a/include/interface.h +++ b/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 diff --git a/src/interface.c b/src/interface.c index 7a9f4b78c..3c71e1de0 100644 --- a/src/interface.c +++ b/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; diff --git a/src/opencl.c b/src/opencl.c index 919f3bbd6..6ba905720 100644 --- a/src/opencl.c +++ b/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 = hashconfig_enforce_kernel_loops (hashconfig, user_options); + + 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 == 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; - } - - if (hashconfig->hash_mode == 14000 && 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 == 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; }