From ed10e6a913142b0ad0e31f70f9c35da06956de9f Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Sun, 22 Jun 2025 20:17:52 +0200 Subject: [PATCH] Autotune and Benchmark refactoring MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit This change affects three key areas, each improving autotuning: - Autotune refactoring itself The main autotune algorithm had become too complex to maintain and has now been rewritten from scratch. The engine is now closer to the old v6.0.0 version, using a much more straightforward approach. Additionally, the backend is now informed when the autotune engine runs its operations and runs an extra invisible kernel invocation. This significantly improves runtime accuracy because the same caching mechanisms which kick in normal cracking sessions now also apply during autotuning. This leads to more consistent and reliable automatic workload tuning. - Benchmarking and '--speed-only' accuracy bugs fixed Benchmark runtimes had become too short, especially since the default benchmark mask changed from '?b?b?b?b?b?b?b' to '?a?a?a?a?a?a?a?a'. For very fast hashes like NTLM, benchmarks often stopped immediately when base words needed to be regenerated, producing highly inaccurate results. This issue also misled users tuning '-n' values, as manually oversubscribing kernels could mask the problem, creating the impression that increasing '-n' had a larger impact on performance than it truly does. While '-n' still has an effect, it’s not as significant. With this fix, users achieve the same speed without needing to tune '-n' manually. The bug was fixed by enforcing a minimum benchmark runtime of 4 seconds, regardless of kernel runtime or kernel type. This ensures more stable and realistic benchmark results, but typically increasing the benchmark duration by up to 4 seconds. - Kernel-Threads set to 32 and plugin configuration cleanup Some plugin configurations existed solely to work around the old benchmarking bug and can now be removed. For example, 'OPTS_TYPE_MAXIMUM_THREADS' is no longer required and has been removed from all plugins, although the parameter itself remains to avoid breaking custom plugins. Because increasing threads beyond 32 no longer offers meaningful performance gains, the default is now capped at 32 (unless overridden with '-T'). This simplifies GPU memory management. Currently, work-item counts are indirectly limited by buffer sizes (e.g., 'pws_buf[]'), which must not exceed 4 GiB (a hard-coded limit). This buffer size depends on the product of 'kernel-accel', 'kernel-threads', and the device’s compute units. By reducing the default threads from 1024 to 32, there is now more space available for base words. --- include/backend.h | 4 +- src/autotune.c | 144 ++++--------------------------------- src/backend.c | 110 ++++++++++++++++------------ src/modules/module_01800.c | 13 +--- src/modules/module_06231.c | 3 +- src/modules/module_06232.c | 3 +- src/modules/module_06233.c | 3 +- src/modules/module_08700.c | 3 +- src/modules/module_09600.c | 11 ++- src/modules/module_10700.c | 3 +- src/modules/module_11700.c | 3 +- src/modules/module_11750.c | 3 +- src/modules/module_11760.c | 3 +- src/modules/module_11800.c | 3 +- src/modules/module_11850.c | 3 +- src/modules/module_11860.c | 3 +- src/modules/module_13100.c | 3 +- src/modules/module_13731.c | 3 +- src/modules/module_13732.c | 3 +- src/modules/module_13733.c | 3 +- src/modules/module_13771.c | 3 +- src/modules/module_13772.c | 3 +- src/modules/module_13773.c | 3 +- src/modules/module_13781.c | 3 +- src/modules/module_13782.c | 3 +- src/modules/module_13783.c | 3 +- src/modules/module_18600.c | 3 +- src/modules/module_19000.c | 3 +- src/modules/module_19300.c | 3 +- src/modules/module_20510.c | 1 - src/modules/module_21000.c | 3 +- src/modules/module_24500.c | 3 +- src/modules/module_26000.c | 3 +- src/modules/module_27000.c | 1 - src/modules/module_27100.c | 1 - src/modules/module_29431.c | 3 +- src/modules/module_29432.c | 3 +- src/modules/module_29433.c | 3 +- src/selftest.c | 36 +++++----- src/user_options.c | 2 +- 40 files changed, 141 insertions(+), 272 deletions(-) diff --git a/include/backend.h b/include/backend.h index 0e3fe2359..339fd6bf9 100644 --- a/include/backend.h +++ b/include/backend.h @@ -56,7 +56,7 @@ int gidd_to_pw_t (hashcat_ctx_t *hashcat_ctx, hc_devi int copy_pws_idx (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u64 gidd, const u64 cnt, pw_idx_t *dest); int copy_pws_comp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, u32 off, u32 cnt, u32 *dest); -int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 highest_pw_len, const u64 pws_pos, const u64 pws_cnt, const u32 fast_iteration, const u32 salt_pos); +int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 highest_pw_len, const u64 pws_pos, const u64 pws_cnt, const u32 fast_iteration, const u32 salt_pos, const bool is_autotune); 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); @@ -84,7 +84,7 @@ int run_opencl_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_devi 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); +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, const bool is_autotune); int run_kernel_mp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kern_run, const u64 num); int run_kernel_tm (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param); int run_kernel_amp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u64 num); diff --git a/src/autotune.c b/src/autotune.c index 64bfab721..754b4827a 100644 --- a/src/autotune.c +++ b/src/autotune.c @@ -47,16 +47,16 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par { if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { - run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_try, true, 0); + run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_try, true, 0, true); } else { - run_kernel (hashcat_ctx, device_param, KERN_RUN_4, 0, kernel_power_try, true, 0); + run_kernel (hashcat_ctx, device_param, KERN_RUN_4, 0, kernel_power_try, true, 0, true); } } else { - run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, kernel_power_try, true, 0); + run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, kernel_power_try, true, 0, true); } device_param->spin_damp = spin_damp_sav; @@ -142,28 +142,21 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param // this thread limiting is also performed inside 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; @@ -298,13 +291,13 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_threads = device_param->kernel_wgs1; - run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_max, false, 0); + run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_max, false, 0, true); 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); + run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, kernel_power_max, false, 0, true); } device_param->kernel_threads = kernel_threads_sav; @@ -312,8 +305,6 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param // Do a pre-autotune test run to find out if kernel runtime is above some TDR limit - u32 kernel_loops_max_reduced = kernel_loops_max; - if (true) { double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min, kernel_threads); @@ -326,146 +317,39 @@ 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, kernel_threads); - - const u32 mm = kernel_loops_max / kernel_loops_min; - - if ((exec_msec * mm) > target_msec) - { - const u32 loops_valid = (const u32) (target_msec / exec_msec); - - kernel_loops_max_reduced = kernel_loops_min * loops_valid; - } } - // first find out highest kernel-loops that stays below target_msec + // v7 autotuner is a lot more straight forward - if (kernel_loops_min < kernel_loops_max) + for (u32 kernel_loops_test = kernel_loops_min; kernel_loops_test <= kernel_loops_max; kernel_loops_test <<= 1) { - for (kernel_loops = kernel_loops_max; kernel_loops > kernel_loops_min; kernel_loops >>= 1) - { - if (kernel_loops > kernel_loops_max_reduced) continue; + double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_test, kernel_threads); - double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_min, kernel_loops, kernel_threads, 1); + if (exec_msec > target_msec) break; - if (exec_msec < target_msec) break; - } + kernel_loops = kernel_loops_test; } - #define STEPS_CNT 16 + #define STEPS_CNT 12 - // now the same for kernel-accel but with the new kernel-loops from previous loop set + // now we tune 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++) { - const u32 kernel_accel_try = 1U << i; + const u32 kernel_accel_try = kernel_accel * 2; if (kernel_accel_try < kernel_accel_min) continue; if (kernel_accel_try > kernel_accel_max) break; - double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_try, kernel_loops, kernel_threads, 1); + double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_try, kernel_loops, kernel_threads, 3); if (exec_msec > target_msec) break; kernel_accel = kernel_accel_try; } } - - // now find the middle balance between kernel_accel and kernel_loops - // while respecting allowed ranges at the same time - - if (kernel_accel < kernel_loops) - { - const u32 kernel_accel_orig = kernel_accel; - const u32 kernel_loops_orig = 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++) - { - const u32 kernel_accel_try = kernel_accel_orig * (1U << i); - const u32 kernel_loops_try = kernel_loops_orig / (1U << i); - - if (kernel_accel_try < kernel_accel_min) continue; - if (kernel_accel_try > kernel_accel_max) break; - - if (kernel_loops_try > kernel_loops_max) continue; - if (kernel_loops_try < kernel_loops_min) break; - - // do a real test - - 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; - - exec_msec_prev = exec_msec; - - // so far, so good! save - - kernel_accel = kernel_accel_try; - kernel_loops = kernel_loops_try; - - // too much if the next test is true - - if (kernel_loops_try < kernel_accel_try) break; - } - } - - 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); - - const u32 accel_left = kernel_accel_max / kernel_accel; - - const u32 exec_accel_min = MIN (exec_left, accel_left); // we want that to be int - - if (exec_accel_min >= 1) - { - // this is safe to not overflow kernel_accel_max because of accel_left - - kernel_accel *= exec_accel_min; - } - - // v6.2.4 new section: find 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 measurements 3 times - - if (kernel_threads_max > kernel_threads_min) - { - const u32 kernel_accel_orig = kernel_accel; - const u32 kernel_threads_orig = kernel_threads; - - double exec_msec_prev = try_run_times (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads, 3); - - 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); - - // since we do not modify total amount of workitems, we can (and need) to do increase kernel_accel_max - - const u32 kernel_accel_max_try = kernel_accel_max * (1U << i); - - if (kernel_accel_try > kernel_accel_max_try) break; - - 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 diff --git a/src/backend.c b/src/backend.c index 51375a56a..d9880c195 100644 --- a/src/backend.c +++ b/src/backend.c @@ -941,7 +941,7 @@ int copy_pws_comp (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, return 0; } -int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 highest_pw_len, const u64 pws_pos, const u64 pws_cnt, const u32 fast_iteration, const u32 salt_pos) +int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 highest_pw_len, const u64 pws_pos, const u64 pws_cnt, const u32 fast_iteration, const u32 salt_pos, const bool is_autotune) { bridge_ctx_t *bridge_ctx = hashcat_ctx->bridge_ctx; hashconfig_t *hashconfig = hashcat_ctx->hashconfig; @@ -1025,20 +1025,20 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (highest_pw_len < 16) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_1, pws_pos, pws_cnt, true, fast_iteration) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_1, pws_pos, pws_cnt, true, fast_iteration, is_autotune) == -1) return -1; } else if (highest_pw_len < 32) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2, pws_pos, pws_cnt, true, fast_iteration) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2, pws_pos, pws_cnt, true, fast_iteration, is_autotune) == -1) return -1; } else { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_pos, pws_cnt, true, fast_iteration) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_pos, pws_cnt, true, fast_iteration, is_autotune) == -1) return -1; } } else { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_4, pws_pos, pws_cnt, true, fast_iteration) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_4, pws_pos, pws_cnt, true, fast_iteration, is_autotune) == -1) return -1; } } else @@ -1159,12 +1159,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_INIT) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_1, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_1, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_HOOK12) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_12, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_12, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; if (device_param->is_cuda == true) { @@ -1259,7 +1259,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; } if (true) @@ -1279,12 +1279,12 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_LOOP) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2, pws_pos, pws_cnt, true, slow_iteration) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2, pws_pos, pws_cnt, true, slow_iteration, is_autotune) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_LOOP_EXTENDED) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2E, pws_pos, pws_cnt, true, slow_iteration) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2E, pws_pos, pws_cnt, true, slow_iteration, is_autotune) == -1) return -1; } //bug? @@ -1416,7 +1416,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_HOOK23) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_23, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_23, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; if (device_param->is_cuda == true) { @@ -1508,7 +1508,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_INIT2) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_INIT2, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_INIT2, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; } if (true) @@ -1521,7 +1521,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_LOOP2_PREPARE) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_LOOP2P, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_LOOP2P, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_LOOP2) @@ -1539,7 +1539,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, device_param->kernel_param.loop_pos = loop_pos; device_param->kernel_param.loop_cnt = loop_left; - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_LOOP2, pws_pos, pws_cnt, true, slow_iteration) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_LOOP2, pws_pos, pws_cnt, true, slow_iteration, is_autotune) == -1) return -1; //bug? //while (status_ctx->run_thread_level2 == false) break; @@ -1647,7 +1647,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_AUX1) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX1, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX1, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; if (status_ctx->run_thread_level2 == false) break; @@ -1656,7 +1656,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_AUX2) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX2, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX2, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; if (status_ctx->run_thread_level2 == false) break; @@ -1665,7 +1665,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_AUX3) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX3, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX3, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; if (status_ctx->run_thread_level2 == false) break; @@ -1674,7 +1674,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (hashconfig->opts_type & OPTS_TYPE_AUX4) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX4, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX4, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; if (status_ctx->run_thread_level2 == false) break; @@ -1685,7 +1685,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (hashconfig->opts_type & OPTS_TYPE_COMP) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; } if (status_ctx->run_thread_level2 == false) break; @@ -1703,7 +1703,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 deep_comp_kernel = module_ctx->module_deep_comp_kernel (hashes, salt_pos, loops_pos); - if (run_kernel (hashcat_ctx, device_param, deep_comp_kernel, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, deep_comp_kernel, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; if (status_ctx->run_thread_level2 == false) break; } @@ -1713,7 +1713,7 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, { if (hashconfig->opts_type & OPTS_TYPE_COMP) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_pos, pws_cnt, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_pos, pws_cnt, false, 0, is_autotune) == -1) return -1; } } } @@ -2268,7 +2268,7 @@ int run_opencl_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *devi return 0; } -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) +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, const bool is_autotune) { const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const status_ctx_t *status_ctx = hashcat_ctx->status_ctx; @@ -2431,6 +2431,11 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + if (is_autotune == true) + { + if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; + } + if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event1, device_param->cuda_stream) == -1) return -1; if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; @@ -2527,6 +2532,11 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + if (is_autotune == true) + { + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; + } + if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event1, device_param->hip_stream) == -1) return -1; if (hc_hipLaunchKernel (hashcat_ctx, hip_function, num_elements, 1, 1, kernel_threads, 1, 1, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; @@ -2650,6 +2660,11 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con const size_t global_work_size[3] = { num_elements, 1, 1 }; const size_t local_work_size[3] = { kernel_threads, 1, 1 }; + if (is_autotune == true) + { + hc_mtlEncodeComputeCommand (hashcat_ctx, metal_command_encoder, metal_command_buffer, global_work_size[0], local_work_size[0], &ms); + } + double ms = 0; const int rc_cc = hc_mtlEncodeComputeCommand (hashcat_ctx, metal_command_encoder, metal_command_buffer, global_work_size[0], local_work_size[0], &ms); @@ -2760,6 +2775,11 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con const size_t global_work_size[3] = { num_elements, 1, 1 }; const size_t local_work_size[3] = { kernel_threads, 1, 1 }; + if (is_autotune == true) + { + if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; + } + if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; // spin damper section @@ -3683,6 +3703,23 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } } + if (device_param->is_cuda == true) + { + if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1; + } + + if (device_param->is_hip == true) + { + if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1; + } + + #if defined (__APPLE__) + if (device_param->is_metal == true) + { + // what to do here? + } + #endif + if (device_param->is_opencl == true) { if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1; @@ -3756,13 +3793,6 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co device_param->outerloop_pos = 0; device_param->outerloop_left = pws_cnt; - // we ignore the time to copy data over pci bus in this case - - if (user_options->speed_only == true) - { - hc_timer_set (&device_param->timer_speed); - } - // loop start: most outer loop = salt iteration, then innerloops (if multi) u32 salts_cnt = hashes->salts_cnt; @@ -4290,7 +4320,7 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co } } - if (choose_kernel (hashcat_ctx, device_param, highest_pw_len, pws_pos, pws_cnt, fast_iteration, salt_pos) == -1) return -1; + if (choose_kernel (hashcat_ctx, device_param, highest_pw_len, pws_pos, pws_cnt, fast_iteration, salt_pos, false) == -1) return -1; /** * benchmark was aborted because too long kernel runtime (slow hashes only) @@ -4357,16 +4387,6 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co if (user_options->speed_only == true) { - // let's abort this so that the user doesn't have to wait too long on the result - // for slow hashes it's fine anyway as boost mode should be turned on - - if (hashconfig->attack_exec == ATTACK_EXEC_OUTSIDE_KERNEL) - { - device_param->speed_only_finish = true; - - break; - } - double total_msec = device_param->speed_msec[0]; for (u32 speed_pos = 1; speed_pos < device_param->speed_pos; speed_pos++) @@ -4443,7 +4463,7 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co device_param->outerloop_msec = total_msec * hashes->salts_cnt * device_param->outerloop_multi; - device_param->speed_only_finish = true; + //device_param->speed_only_finish = true; } if (iconv_enabled == true) @@ -10188,8 +10208,10 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) } else { - if (device_param->is_hip == true) - { + // v7 test, needs some larger test, but I think we will need to stick to this + + //if (device_param->is_hip == true) + //{ const u32 native_threads = device_param->kernel_preferred_wgs_multiple; if ((native_threads >= device_param->kernel_threads_min) && (native_threads <= device_param->kernel_threads_max)) @@ -10201,7 +10223,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) { // abort? } - } + //} } // this seems to work always diff --git a/src/modules/module_01800.c b/src/modules/module_01800.c index 3e47ca1b5..fea4851c6 100644 --- a/src/modules/module_01800.c +++ b/src/modules/module_01800.c @@ -22,9 +22,7 @@ static const u64 KERN_TYPE = 1800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE - | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$6$72820166$U4DVzpcYxgw7MVVDGGvB2/H5lRistD5.Ah4upwENR5UtffLR4X4SxSzfREv8z6wVl0jRFX40/KnYVvK4829kD1"; @@ -417,13 +415,6 @@ static void sha512crypt_encode (const u8 digest[64], u8 buf[86]) buf[85] = int_to_itoa64 (l & 0x3f); //l >>= 6; } -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 = 256; - - return kernel_threads_max; -} - 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)) @@ -616,7 +607,7 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_kernel_accel_min = MODULE_DEFAULT; module_ctx->module_kernel_loops_max = MODULE_DEFAULT; module_ctx->module_kernel_loops_min = MODULE_DEFAULT; - 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; diff --git a/src/modules/module_06231.c b/src/modules/module_06231.c index 72525f3a0..9dbc28864 100644 --- a/src/modules/module_06231.c +++ b/src/modules/module_06231.c @@ -26,8 +26,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_BINARY_HASHFILE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_BINARY_HASHFILE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "cf53d4153414b63285e701e52c2d99e148c6ccc4508132f82cb41862d0a0ac9ea16274285ac261c339c1508eec9fea54c33e382458662913678f2a88a84959a678e238973985ec670d50252677430587ee28b72bfa5edfb2f79c40b734ba8a54a3662642a6ab067e75f41154688ad4adb5d6decd891462dd537188195a51e06fa5baf22b69d0f472cfeeae77ab9a90091731863af1d8b5b380da179fa7d5227ef031732b1ae06e0fe34c0b28b7a64eac34e5a08e09d7001394b3afc804ac69bf819cdd2d383fe96a721f7c683628da8e529d84bdaa68d702573d8f7ef26f75d1bd5c91efa88cb33b1e9c006b87981c55ed3b8063ab7068f8e99b128bc56ea3e883efa55d6f340b2681e50405d91f5f6d76cdbeac404944164d329d3ee01311de0bc6547310f126b5a4c0e9fb74825f91faefa60b7ac828819d4544c1872ff5041e61d5cf093553f427358b2181046376d7b876e1bccf0774d5d251b7c922c214bb5c70c715165d028e1dca73e7adeca3396d77f6e597a10dd4c58f37fdbbdc1d04cd8890ba4c5025776a88a349bb925add13193becf1ca10fe32536db0c0b06a1ef799fb692e304b3716ca5a8a80859c4012ca3e06701b46b5a32f4d10e285a0cdaf6c24e0d98139e7f306e52503c9b503aa28f1fbbb236284907068074fcb3e267e3c4aab2bd3b79b24a7a08106bb55850fa2bb8e2f6d9919a6743cb822c164"; diff --git a/src/modules/module_06232.c b/src/modules/module_06232.c index 962f3541c..83088fd6b 100644 --- a/src/modules/module_06232.c +++ b/src/modules/module_06232.c @@ -26,8 +26,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_BINARY_HASHFILE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_BINARY_HASHFILE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "e9e503972b72dee996b0bfced2df003a54b42399e3586520cf1f69475ba32aff564e40e604a505af95ce15220f558ae815e94ce4953882a8299ee3fffb12e9bd62bf8e2c41c0a8337ce20d45715440cc83e394200d351c5b04be5b70fa11b8467320a091a1d703c88cc7b26fd114795c04a973b3266ba97f55d4b4e4771bb1b4a6aabc9d57e03f0ae7c8a77dfc3d37078efba45031e7d63bb514726e2f2dc6da8cce167a17e36b32c326a5bcaa2c4b445f6e10e1f899a9adcc2a698769f900b7909f7aec52fc9862d75286ffda67933f9c52e5c681d590ad0329b85f8db0f6bb6daa3b2d55b62c65da37e3e7fcb99954e0abe20c39724e8fb2c7f839ec67d35f151dfd8c4dd4bc8dc4393fab291efa08cc0099277d219a0ba4c6272af3684d8043ed3f502b98e196dc7aa0291627613179199976f28eff08649acf70aa0c0dc5896ed13eb18ea28fdd6c460a9c7cfedeab5ac80a3c195226cfca094a7590fa2ae5ed2133ba09b5466b2049b6291f8dcf345e5718a4c0ef3f9c8d8e07d0e5dddd07452b533fbf243ef063fb6d26759ae725d8ca430f8cf17b86665d23bdff1c9dbdfe601b88e87cb7c89f23abc4a8bb1f0b7375cc29b1d81c950ffe92e16e2080e1d6270bbb3ba753322d2b623caed87213e552c33e699d4010f0f61df2b7f460d7cd82e70a711388f1c0b591d424259d3de8b3628daf62c6c5b71864eb0e7d31"; diff --git a/src/modules/module_06233.c b/src/modules/module_06233.c index 4dbf860a8..5c928a646 100644 --- a/src/modules/module_06233.c +++ b/src/modules/module_06233.c @@ -26,8 +26,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_BINARY_HASHFILE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_BINARY_HASHFILE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "de7d6725cc4c910a7e96307df69d41335e64d17b4425ca5bf1730f27820f92df9f20f3e855d8566eb5255927153f987348789666c8e563e366a09e68a8126b11c25ac817b2706dde5cec3946e64332b21b41b928985c1a637559ead5b4fecac74ff0d625ef6d8be93dea3eaca05394f23ee9e079d3504a77b4c0b22d3cfcafa9c670966bfa3a5f30539250d97267a9e56b5a1437b1fd2ce58f4ab78b52ba61d01c28d7a6b726d92c8819711c70f820690cf2b9bbef75f196ba87fb5f72a29e213096a8be3b6e6d0ff3dc22563dc9e7d95be68ad169c233289fccfdc2f5528c658cb178b4e78d54e96cb452859b01dd756ca0245bdd586fb450e84988071428c80af0a6dc5f16dea8094da3acb51ac5d2a710414256b2423e0333584437ea9a65a07f06bd241103a478d137e9a274a78a19d3ca121f1bc10e4c9e5fc277d23107db1fb447f71ba0f92b20e3ead77cffaca25f772182705a75e500d9aab3996bfda042f4bdfe35a3a477e355c76a711ad0f64848d6144073ce6ec4152c87973fc3e69626523463812061c51f51fc08487e8a4dbae1ca7965c11f222c607688b3384c5c29d4fe91d14d2cc940a6a9d94486d1823261928d88f56fe00e206d7a31734de0217afd38afa3d2cf3499c2dcff13332a369c4b1f39867f6dfc83ec32d19b931b082f07acac7e70bdd537e8432245c11662d89ec3cc97e582de5d2cc6bde7"; diff --git a/src/modules/module_08700.c b/src/modules/module_08700.c index 8351a9dc0..39b157196 100644 --- a/src/modules/module_08700.c +++ b/src/modules/module_08700.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED | OPTI_TYPE_RAW_HASH; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE - | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "(GDJ0nDZI8l8RJzlRbemg)"; diff --git a/src/modules/module_09600.c b/src/modules/module_09600.c index d6b4b589d..e0f7727fe 100644 --- a/src/modules/module_09600.c +++ b/src/modules/module_09600.c @@ -24,9 +24,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_DEEP_COMP_KERNEL - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_DEEP_COMP_KERNEL; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$office$*2013*100000*256*16*67805436882475302087847656644837*0c392d3b9ca889656d1e615c54f9f3c9*612b79e33b96322c3253fc8a0f314463cd76bc4efe1352f7efffca0f374f7e4b"; @@ -61,6 +59,13 @@ typedef struct office2013_tmp static const char *SIGNATURE_OFFICE2013 = "$office$"; +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) +{ + const u32 kernel_loops_max = 1000; + + return kernel_loops_max; +} + 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)) diff --git a/src/modules/module_10700.c b/src/modules/module_10700.c index 3c5f332c6..5a4f52746 100644 --- a/src/modules/module_10700.c +++ b/src/modules/module_10700.c @@ -24,8 +24,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_HASH_COPY - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_HASH_COPY; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$pdf$5*6*256*-1028*1*16*62137640825124540503886403748430*127*0391647179352257f7181236ba371e540c2dbb82fac1c462313eb58b772a54956213764082512454050388640374843000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000*127*00000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000*32*0000000000000000000000000000000000000000000000000000000000000000*32*0000000000000000000000000000000000000000000000000000000000000000"; diff --git a/src/modules/module_11700.c b/src/modules/module_11700.c index 9de690ddc..9df95c265 100644 --- a/src/modules/module_11700.c +++ b/src/modules/module_11700.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_PT_ADD01 - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_ADD01; static const u32 SALT_TYPE = SALT_TYPE_NONE; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "57e9e50caec93d72e9498c211d6dc4f4d328248b48ecf46ba7abfa874f666e36"; diff --git a/src/modules/module_11750.c b/src/modules/module_11750.c index 1e3b779a7..afd277f68 100644 --- a/src/modules/module_11750.c +++ b/src/modules/module_11750.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_PT_ADD01 - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_ADD01; static const u32 SALT_TYPE = SALT_TYPE_GENERIC; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "0f71c7c82700c9094ca95eee3d804cc283b538bec49428a9ef8da7b34effb3ba:08151337"; diff --git a/src/modules/module_11760.c b/src/modules/module_11760.c index 25e931699..df888f5aa 100644 --- a/src/modules/module_11760.c +++ b/src/modules/module_11760.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_PT_ADD01 - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_ADD01; static const u32 SALT_TYPE = SALT_TYPE_GENERIC; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "d5c6b874338a492ac57ddc6871afc3c70dcfd264185a69d84cf839a07ef92b2c:08151337"; diff --git a/src/modules/module_11800.c b/src/modules/module_11800.c index d06f91bdc..7e8711f94 100644 --- a/src/modules/module_11800.c +++ b/src/modules/module_11800.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_PT_ADD01 - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_ADD01; static const u32 SALT_TYPE = SALT_TYPE_NONE; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "5d5bdba48c8f89ee6c0a0e11023540424283e84902de08013aeeb626e819950bb32842903593a1d2e8f71897ff7fe72e17ac9ba8ce1d1d2f7e9c4359ea63bdc3"; diff --git a/src/modules/module_11850.c b/src/modules/module_11850.c index ca04fbfaa..28ccdc94f 100644 --- a/src/modules/module_11850.c +++ b/src/modules/module_11850.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_PT_ADD01 - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_ADD01; static const u32 SALT_TYPE = SALT_TYPE_GENERIC; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "be4555415af4a05078dcf260bb3c0a35948135df3dbf93f7c8b80574ceb0d71ea4312127f839b7707bf39ccc932d9e7cb799671183455889e8dde3738dfab5b6:08151337"; diff --git a/src/modules/module_11860.c b/src/modules/module_11860.c index 835d5819f..f08306f3d 100644 --- a/src/modules/module_11860.c +++ b/src/modules/module_11860.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_PT_ADD01 - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_ADD01; static const u32 SALT_TYPE = SALT_TYPE_GENERIC; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "bebf6831b3f9f958acb345a88cb98f30cb0374cff13e6012818487c8dc8d5857f23bca2caed280195ad558b8ce393503e632e901e8d1eb2ccb349a544ac195fd:08151337"; diff --git a/src/modules/module_13100.c b/src/modules/module_13100.c index 6d18aa116..49b609d54 100644 --- a/src/modules/module_13100.c +++ b/src/modules/module_13100.c @@ -22,8 +22,7 @@ static const u64 KERN_TYPE = 13100; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE - | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$krb5tgs$23$*user$realm$test/spn*$b548e10f5694ae018d7ad63c257af7dc$35e8e45658860bc31a859b41a08989265f4ef8afd75652ab4d7a30ef151bf6350d879ae189a8cb769e01fa573c6315232b37e4bcad9105520640a781e5fd85c09615e78267e494f433f067cc6958200a82f70627ce0eebc2ac445729c2a8a0255dc3ede2c4973d2d93ac8c1a56b26444df300cb93045d05ff2326affaa3ae97f5cd866c14b78a459f0933a550e0b6507bf8af27c2391ef69fbdd649dd059a4b9ae2440edd96c82479645ccdb06bae0eead3b7f639178a90cf24d9a"; diff --git a/src/modules/module_13731.c b/src/modules/module_13731.c index 843ea9595..53e77b18c 100644 --- a/src/modules/module_13731.c +++ b/src/modules/module_13731.c @@ -29,8 +29,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_BINARY_HASHFILE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "48f79476aa0aa8327a8a9056e61450f4e2883c9e9669142f2e2f022c2f85303b897d088dea03d64329f6c402a56fed05b3919715929090a25c8ae84c67dbdb364ebfa3e9ccc0b391c130a4c3dd6495a1d6eb5d2eab72f8009096f7475ecb736bb3225b6da144e1596d859dad159fae5a739beea88ea074771e9d0b2d7c48ae302606a60d7cff6db54f3e460c548c06a4f47dc1ac203a8c8349fbff6a652219a63f27bc76327543e22be4f8dab8e4f90a4283fbf1552119fe24114ce8869eb20ce87dd72300f7aad3f7b4a26a355f16517725449151cf0373dbd0b281f6ac753485a14a5361cc75d40928e241a6b4684658801774843238048cf8c7f2fd88950abac040e12b0c41fdcaca3702907e951ec11c061a91b3050a4855abe6f3b50b4bd0b17c4be1f5b50b873eadc2d8446cd72c4fcac576bbce3acea769f740c5322ee8c927ffd4dd11c8a9e66f06e58df2e5d4d85c13b44c412bab839c9512b7a0acdd97b37dcccc4b70854eda0f36de12d62dd10cc13bc6154103d083bf6540bc78e5d0aad5d063cc74dad4cbe6e060febda2a9fd79c238f99dcb0766ff4addcfd0c03e619c765f65b1c75d5d22c6536958bcda78077ff44b64c4da741bf50154df310d4e0724238a777b524237b9478277e400ad8146dc3ca1da83e3d2f1c5115a4b7fcdc71dd7d56ba86a2f9b721c9a4137aabb07c3c5fedcf5342c4fae4898c9"; diff --git a/src/modules/module_13732.c b/src/modules/module_13732.c index 0dd660582..2435a0478 100644 --- a/src/modules/module_13732.c +++ b/src/modules/module_13732.c @@ -29,8 +29,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_BINARY_HASHFILE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "1b721942019ebe8cedddbed7744a0702c0e053281a467e0ed69bf875c7406407d72eb8f2aea21270e41898c0a2c14382f86e04c15e7bc019d1d9dd813eabee0ae5173e3cb1d927859d3e6de1006335a5184ae12b4c8dc2db2b1cd785063152a776f4dc5cacc1856a919b880d704b7450f5a0e0c9521bc9b4d67213c36a50e6664a1cbcea33f997b858e654111c7e9fca74f361528e85a28880381ec2600e3c1cd508c3833dd21cc91978185cba53caefd7b3c82d219d49f0b41e536d32e8d3ce194ad7923ca742213e19dcebdbd9687979d5a594654a5c611e8b829c4019e90a3cfb14e5fd7f8ed91e0fc79eed182399f02a3e3e202d4becaa6730e1f05f99ce06ce16dba7777ccddac72e85f2d3be5ecc9c808ac273f10ceb71cad666166abc327c4061a5f47424a5b6d9d093782f34b49924342a2e8cea663446ed4232a9a415ee2dfde988fa827b06d7438fec20ad0689543c3ee4602ce3ec3806fc7d668ef7e34330edd1e077b329a7627fa3ae5c89308258a17ecefbee114c80c2ab06f8271f14de8f2d13d1d6e5a119b71a6bae88ab151f76cdb2442284bc481d0df7e2163c3acfe763d3968195450d275af9034a00184a30cefed163e636626bffe6a35df3472508a49cb2b9b4c4a95d11c5d17e4e0539e9f13112125515778bcd1c2813c62a02673663062ad60583ec6a02c8a572865829e5b8c767b285728bea4907"; diff --git a/src/modules/module_13733.c b/src/modules/module_13733.c index 5cc512263..9c2729067 100644 --- a/src/modules/module_13733.c +++ b/src/modules/module_13733.c @@ -29,8 +29,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_BINARY_HASHFILE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "5eb128daef63eff7e6db6aa10a8858f89964f47844acca68df82ebb2e73866fa75e3b7a53f9d2ff1ecdd1f4dc90e9c0fdf51f60d11b1992cd2971b4889edfc8920bbf346fd7693f675b617cb9e4e9a43e6f445021068fc13453b130f2eb1d753ee83ecc61dabec293e88b62110cf6a8fab670e171f6aba2226550b54893263f5fa086b3cc41dd3db2eae07b585e5162c7a0d9723a426d408d83266c4d6018dc1b8b456d28a224033a30bfe62b1e58c2ddf596e07f7ff31849a6f5cfcc1c977b82d8484c270d44ededb0afdb781295e92968fc8cc69766af0ce1e72f02d6b4e124ba4b1af71519dcaade857bb3f371f93a350da6e65ee46c2ac782f134c75c10fe9d653fccc08c614dc362871911af8b83bdfc479f770dfe4b3c86b5d895842c53852fe4912738f848bf7c3e10b8189d25faceab9ef30b6fa0284edaa471752ac2b65335179b8d605417709f64fded7d94383618a921660d4cdb190bbb3769a8e56d2cd1ee07078ebc3b68ebeb016893f7099018e40cb326e32b29a62806eaf1a3fd382f4f876bf721eadfc019c5545813e81fd7168995f743663b136762b07910a63b6eec5b728a4ad07a689cceecb14c2802f334401a0a4fd2ec49e2da7f3cb24d6181f01ceed93ee73dedc3378133c83c9a71155c86785ff20dd5a64323d2fd4bf076bab3c17a1bb45edf81c30a7bd7dbbb097ece0dca83fff9138d56ae668"; diff --git a/src/modules/module_13771.c b/src/modules/module_13771.c index 74efd072e..b753d171f 100644 --- a/src/modules/module_13771.c +++ b/src/modules/module_13771.c @@ -30,8 +30,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_BINARY_HASHFILE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "444ec71554f0a2989b34bd8a5750ae7b5ed8b1ccdead29120fc030bd5186f312a7fa18ab4f4389d7798e43c073afd1e71dda2052db38dec04a700e8d6b488802ead0cf95d6e6cecc8eaf6464baf94a64acbbd1a86f826333115b6380bda18cf936150efd6ffc2a344bb78b0b4875781a8c5079772429ef50ddf148f35895496d2e39f32ffaf68a007b070b0beaad316c4b3adf43c0c58ad24430a34abf168ed455b64958ca5465cae0684adadc00f7b9c13fc7671b4520892d23aebff49ea92bc15e804cc650dc3bbd5b8f5122051636f0c576977d4b64ba355bf6e6a8e042fc5165f2a8affa51aa12ff718cee4c543976bf565997b4b57c74e79584e317f4bdb3920f2937c4251af87f432bb8ce78dcb30675246f0303db4aaea913c93be5a26d16dbf8d4d20773aa2a4608d2151491ca6593b51965baeaf9b58f78905df522bf88976fe9436a916c8de38d5a6ca7ca7f436e7982a36335a404298304322ebe194bb34e91e8f7ee7c6541679bb0ce9d80bf4431d1c475b1a785e943e57f8e27a4e665940389b6da2771bd27d943955185379f83ca6a124ec55b2b63d4ef2e2ad6ee27de25f959708f3a64facfe07f06e29459a14f02699751d530f258d0c744a759c188de4f9423f2bd21d3d999ea28df4f3a93a2c47a7e788fe43ccbfbe267277b048002da1ef8c1e7b26690230285675a3a8fdc0f2acf46a4cb24141b3ad1"; diff --git a/src/modules/module_13772.c b/src/modules/module_13772.c index 3a990cc06..6494ee637 100644 --- a/src/modules/module_13772.c +++ b/src/modules/module_13772.c @@ -30,8 +30,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_BINARY_HASHFILE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "0f5da0b17c60edcd392058752ec29c389b140b54cd1f94de43dccea703b1fd37936e75a500b7f9d4e94e7f214c4696c051be9697792a856ebf9c0f5a598cf8ba5621e49c7505eba3b4738acdc860b6ed648f52e5b673ae06bb04616de438a090ab19abea11c30984ead06859de9b7aec8e436c40816f67a56cb53d5f125e58c42225315a4bf494da8128f0df924bcf6ad4b91c9efc5cb0be67cb0cd753c392388d780f57aba39197513a191cc684e9ebee41bc901dd99e9a625141cf98e55e8f74d838baea3bf8f411b85c14eff8cddd1720c2539eef7a38a72c4ed9745a05476b6a16bcda2a5391c94b6f499e3bea64ff412d03d060741e938ed3dc905d8bd6dbb2420e9277251ebe3421be389ea8b02782baeb258b9ec7e0732b3817ee6da58209871aee4e16d57a132c6215782364570238157d8a7fdcd29f54ab2295f68d027dc9f2e0c951afad7500cafe3219e6530699918ac55f4fa1141bc3596155b05bae2fdc8b0a5438edeb5bb0cfac592565b20645be90b406a1fd59846957e7539fd8423bfd4c7ae7d608aacb084ae887baa1a83b14afff8d2063565086c66e293234a8667af39642b90a38c3a5bd4fa8a787c60f73882535c9b34cb7b243465dcc32aff29cee0e741ff059c6acd8ddcbdb3cfafecdcd0f45c84dd871be4fbffd5ac2ab9e01898009adcf7d932c37d6568ad875e4d6ea15db29a1e8ba5a4e86bd"; diff --git a/src/modules/module_13773.c b/src/modules/module_13773.c index fd9dd44ff..6f2740347 100644 --- a/src/modules/module_13773.c +++ b/src/modules/module_13773.c @@ -30,8 +30,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_BINARY_HASHFILE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "18d2e8314961850f8fc26d2bc6f896db9c4eee301b5fa7295615166552b2422042c6cf6212187ec9c0234908e7934009c23ceed0c4858a7a4deecbc59b50a303afdc7d583cde1b0c06f0bf56162ef1d6d8df8f194aadcbe395780b3d1d7127faf39910eb10f4805abdd1c3ef7a66972603124a475e2b9224699e60a9e12f4096597f20c5fb0528f590d7bd317e41dc6a2128cf5e58a99803a28c213feb8286350b1d7ab56d43bb52e511f3c860e5002472a4454a549509c8ce0c34f17ece23d5b61aa7c63389c8ca44ed10c2caae03e7ed30b3ef98565926d7e4f3a2a9abf03b278083bed7aaadd78d5bffb7cd45ffae92990c06d9e9f375a77a94226035d1f90e177c46a04dab416dfb7ed7c4ed9ee7e84580bed65c5fee9f4b1545b9a7cf6af533870d393eced609aebe308ec1eee3729da09eb7df7a8d1282b15c4a1b8266a456c06b4ea20c209c549d5d6b58a861f8e15cca3b6cef114accbf470ec76d717f6d7d416d7a32f064ab560c1167f9ef4e93310fbd927b088bffbb0cf5d5c2e271c9cad4c604e489e9983a990b23e1a2f973682fdfe38df385474f73ecdc9bce701d01d627192d3051240f4b96bbdcf2346b275e05aa75add4acb97b286cc00e830fee95d0f86a8b1e315ccb6f3f8642180392b3baac01ed2c97c200489b5e5ca4dcb0a6417e622b6196482a10e640b2b6b08e3f62acac3d45dfc6b88c666205"; diff --git a/src/modules/module_13781.c b/src/modules/module_13781.c index fcc2ed181..fbbb6b0bf 100644 --- a/src/modules/module_13781.c +++ b/src/modules/module_13781.c @@ -31,8 +31,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE | OPTS_TYPE_KEYBOARD_MAPPING - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "2bfe4a72e13388a9ce074bbe0711a48d62f123df85b09e0350771edc4a0e4f397038a49b900275c9158145a96b52f95e92f927b3f963c7eadb71a07518d643231041c457d2794d0aa505f794153b52b24441271185d386833fbabf0e880c51b544f583d0db2ab6a926ddd3cdd0b68a61d7f5fe3f0ac6aa06ca676a868f373d35073605cf9d521ff55862b5005213a881a7b9025afc3409fa34dc86496620835df072fecd5b501f15e08113835c510d9f0bfd09d2ef1ac0e7bd01f0523d74a54fe984eb497cb960cce5bb154e024dc0c6c61a61e20a45a8f8ef319c63ca9646fbe00930302a5910891a1bc84bd936c926ca535b3b40c9e0ab255363b24a28bb8216d3d32244a725774e6ebbd73d6d3f2a2adcbc28d5341679cbb747efd56db1a09ce80b24640583ffc6f7ca5bd60d59114afcc78601184ba8feadb8d472f86c32bebf70e8158aa56f9db3b3200ef432aa7b370aa4ba408ef11b70d6806f1a21aaa3b629fa06f71dac0ae3e0ce95c7e5b550fc8c46017e151cbbcdf64b3b62b1b846a08925a217227286acfdad35b28407d589bec9578c2a4e9a4320f4a78e1e590fdf53c0a20fe0a1bb6c7d693abcd0e991c449e569477980d4b8972d21e4abc917d897e48ca427c954c3a3e0c8465ef40de51ffc9188047c043224c4a18638f1d91cd88c36623a1d880f18fd0d6ca0b3bbfa7d5d795acfb63576e2c2d83772e8"; diff --git a/src/modules/module_13782.c b/src/modules/module_13782.c index 2772e0bcb..e1e6c89c4 100644 --- a/src/modules/module_13782.c +++ b/src/modules/module_13782.c @@ -31,8 +31,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE | OPTS_TYPE_KEYBOARD_MAPPING - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "af7a64c7c81f608527552532cc7049b0d369e2ce20202d7a41ffb94300cbc9c7ce2130247f49ace4c1512fc3d1b4289ca965e8eb065b35faee5b8479a4e43c4a269f4ee7f6d20f22fe61b2570d46df07b4307f44ba6926f3b44524f0a47be2a0d677d225e2c50ff618b2e5078a19f0613a856bb3145d765cc4c1726aef27b5f03648dcf421b040e7b4fde3193ad9f8a0ae6d91c079610f826e3d556776753d8ca11320632c16a2e49a4eec6e8371681b39be2d7bb826d81dea19eb1dda2e6c71c520a2ad9128b3209a7caf10c196a16ac6f4267ffea8e7be7ddb856976438e0e997773dab75e3dfe0c047f82e4ed0b6e107261b891c4b161fa3c29017428afaaabee5c2dc727fa23b4195265716d92d06e7b828535a77521113077e6f219d7ca721eb5dab66524a530ca3ceba52e3703ec3f435ad1dfee092b090174f4acd1546107db5b403a0ba4fa90c0b4ec19af92a54ebedfd28783dcd83c78499bd33baf3ed10af229ff29634110e2970b6c04232dc95120a365947688abe011f0c8de0e651d9bd112ce0bdf80c4e37c77b962d92f1418272e7484df5217f5f2f3ba1e9b965773ed2727c5d03938516dd789236479b5ff340335c92260b1ad82a313ffa568f912fac799f93b857aaff7b4d76cb525f120a0a9efc376d39c8e27866eff689be01f5adf693ae63ad4b2a77ca96ea985ab7931457f4d8d1afaeb7e423dd752"; diff --git a/src/modules/module_13783.c b/src/modules/module_13783.c index 8765fc6c8..b1339f7a4 100644 --- a/src/modules/module_13783.c +++ b/src/modules/module_13783.c @@ -31,8 +31,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE | OPTS_TYPE_KEYBOARD_MAPPING - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "0c9d7444e9e64a833e857163787b2f6349224bdb4bbf788ce25156c870514226674725be3eebc3f2a2c2ee8adbf8bb3ec1405a333e8e091cec0c5aa77fa9b65048ca01d954912bf3a3b1c38c00297a33ea0e014156ce08d9526150c085e5e2776a1faeb272c6e9539f466f4f93ffe6497c77d3aed54ffcdf1a3e6171cffac7b2ad96bd9e7cc553058894058def68beea05891b0ce734b6a166b8a5f24b4052fc7014b424bd6c33c9d710fb409cdf1a6c7567c1ba6a3010b03f9bda8aa2ef6733542d198a316da0c83106a6c31043f11ac191169db3db994493168ea996737355ccff84f27f6792b3dc87025d3594edb9e759ba3885980df17bc8c751ce3aba0df67aa7997906348729e81c4893cc654dc6b1da3ff7c588a327f45b8acff976d0593cc607dad48a25468d7c3ebc6dd49aa32fc526dd513852cdec4b36f3683b4998800afa25bb968c242d4c66b9b0c77b20d7bd40ffb403e9e087990d59c94ee7d36e9ebfa35a310bab963c253596e6bc89f67d5307823851c526ac789d0628a3eb81f2cdfd7d7612d8be1dade1b17f30aa2bb5d02eb8534caca0c334a269085939a5041c4ad112d325b1bfe3e9851bfdcad80bbc05ecbddc3f2ac09e2ad7182daf6ca5ccc510a100514f5d2dce1ff5046e0c8e7edf0bdc27f8fcdf4e9b3bce786c24bfa28dacee65ee8c913fc18eee5df61b8a43224f3a9c4e1b5de7b600d9e0"; diff --git a/src/modules/module_18600.c b/src/modules/module_18600.c index 86605e2e6..e97870d51 100644 --- a/src/modules/module_18600.c +++ b/src/modules/module_18600.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_DYNAMIC_SHARED - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_DYNAMIC_SHARED; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$odf$*0*0*1024*16*bff753835f4ea15644b8a2f8e4b5be3d147b9576*8*ee371da34333b69d*16*a902eff54a4d782a26a899a31f97bef4*0*dae7e41fbc3a500d3ce152edd8876c4f38fb17d673ee2ac44ef1e0e283622cd2ae298a82d8d98f2ea737247881fc353e73a2f535c6e13e0cdc60821c1a61c53a4b0c46ff3a3b355d7b793fad50de15999fc7c1194321d1c54316c3806956c4a3ade7daabb912a2a36398eba883af088b3cb69b43365d9ba9fce3fb0c1524f73947a7e9fc1bf3adb5f85a367035feacb5d97c578b037144c2793f34aa09dcd04bdaa455aee0d4c52fe377248611dd56f2bd4eb294673525db905f5d905a28dec0909348e6bf94bcebf03ddd61a48797cd5728ce6dbb71037b268f526e806401abcf495f6edd0b5d87118671ec690d4627f86a43e51c7f6d42a75a56eec51204d47e115e813ed4425c97b16b195e02ce776c185194b9de43ae89f356e29face016cb393d6fb93af8ea305d921d5592dd184051ac790b9b90266f52b8d53ce1cb1d762942d6d5bbd0e3821be21af9fa6874ba0c60e64f41d3e5b6caca1c53b575afdc5d8f6a3edbf874dbe009c6cb296466fe9637aed4aed8a43a95ea7d26b4090ad33d4ee7a83844b0893e8bc0f04944205fb9576cb5720f019028cd75ca9ac47b3e5fa231354d74135564df43b659cfaea7e195c4a896e0e0e0c85dc9ce3a9ce9ba552bc2a6dbac4901c19558818e1957ed72d78662bb5ba53475ca584371f1825ae0c92322a4404e63c2baad92665aac29b5c6f96e1e6338d48fb0aef4d0b686063974f58b839484f8dcf0a02537cba67a7d2c4de13125d74820cb07ec72782035af1ea6c4db61c77016d1c021b63c8b07adb4e8510f5c41bbc501f60f3dd16462399b52eb146787e38e700147c7aa23ac4d5d22d9d1c93e67a01c92a197d4765cbf8d56a862a1205abb450a182913a69b8d5334a59924f86fb3ccd0dcfe7426053e26ba26b57c05f38d85863fff1f81135b0366e8cd8680663ae8aaf7d005317b849d5e08be882708fa0d8d02d47e89150124b507c34845c922b95e62aa0b3fef218773d7aeb572c67b35ad8787f31ecc6e1846b673b8ba6172223176eabf0020b6aa3aa71405b40b2fc2127bf9741a103f1d8eca21bf27328cdf15153f2f223eff7b831a72ed8ecacf4ea8df4ea44f3a3921e5a88fb2cfa355ece0f05cbc88fdd1ecd368d6e3b2dfabd999e5b708f1bccaeebb296c9d7b76659967742fe966aa6871cbbffe710b0cd838c6e02e6eb608cb5c81d066b60b5b3604396331d97d4a2c4c2317406e48c9f5387a2c72511d1e6899bd450e9ca88d535755bcfddb53a6df118cd9cdc7d8b4b814f7bc17684d8e5975defaa25d06f410ed0724c16b8f69ec3869bc1f05c71483666968d1c04509875dadd72c6182733d564eb1a7d555dc34f6b817c5418626214d0b2c3901c5a46f5b20fddfdf9f71a7dfd75b9928778a3f65e1832dff22be973c2b259744d500a3027c2a2e08972eaaad4c5c4ec871"; diff --git a/src/modules/module_19000.c b/src/modules/module_19000.c index d8cd72d1c..d6c1787f1 100644 --- a/src/modules/module_19000.c +++ b/src/modules/module_19000.c @@ -23,8 +23,7 @@ static const char *HASH_NAME = "QNX /etc/shadow (MD5)"; static const u64 KERN_TYPE = 19000; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE - | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "@m@75f6f129f9c9e77b6b1b78f791ed764a@8741857532330050"; diff --git a/src/modules/module_19300.c b/src/modules/module_19300.c index aca7830a6..85b83b695 100644 --- a/src/modules/module_19300.c +++ b/src/modules/module_19300.c @@ -23,8 +23,7 @@ static const u64 KERN_TYPE = 19300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_RAW_HASH; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE - | OPTS_TYPE_PT_GENERATE_BE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_GENERATE_BE; static const u32 SALT_TYPE = SALT_TYPE_GENERIC; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "630d2e918ab98e5fad9c61c0e4697654c4c16d73:18463812876898603420835420139870031762867:4449516425193605979760642927684590668549584534278112685644182848763890902699756869283142014018311837025441092624864168514500447147373198033271040848851687108629922695275682773136540885737874252666804716579965812709728589952868736177317883550827482248620334"; diff --git a/src/modules/module_20510.c b/src/modules/module_20510.c index 4c8a340fc..394a69a3b 100644 --- a/src/modules/module_20510.c +++ b/src/modules/module_20510.c @@ -80,7 +80,6 @@ static const u64 KERN_TYPE = 20510; static const u32 OPTI_TYPE = 0; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS | OPTS_TYPE_AUTODETECT_DISABLE | OPTS_TYPE_SUGGEST_KG; static const u32 SALT_TYPE = SALT_TYPE_NONE; diff --git a/src/modules/module_21000.c b/src/modules/module_21000.c index 4a5137e54..bfb8067a7 100644 --- a/src/modules/module_21000.c +++ b/src/modules/module_21000.c @@ -29,8 +29,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE | OPTS_TYPE_PT_ADD80 - | OPTS_TYPE_PT_ADDBITS15 - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_ADDBITS15; static const u32 SALT_TYPE = SALT_TYPE_NONE; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "caec04bdf7c17f763a9ec7439f7c9abda112f1bfc9b1bb684fef9b6142636979b9896cfc236896d821a69a961a143dd19c96d59777258201f1bbe5ecc2a2ecf5"; diff --git a/src/modules/module_24500.c b/src/modules/module_24500.c index 17b67bf8f..3b4e6159b 100644 --- a/src/modules/module_24500.c +++ b/src/modules/module_24500.c @@ -23,8 +23,7 @@ static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE - | OPTS_TYPE_PT_GENERATE_LE - | OPTS_TYPE_MP_MULTI_DISABLE; + | OPTS_TYPE_PT_GENERATE_LE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$telegram$2*100000*77461dcb457ce9539f8e4235d33bd12455b4a38446e63b52ecdf2e7b65af4476*f705dda3247df6d690dfc7f44d8c666979737cae9505d961130071bcc18eeadaef0320ac6985e4a116834c0761e55314464aae56dadb8f80ab8886c16f72f8b95adca08b56a60c4303d84210f75cfd78a3e1a197c84a747988ce2e1b247397b61041823bdb33932714ba16ca7279e6c36b75d3f994479a469b50a7b2c7299a4d7aadb775fb030d3bb55ca77b7ce8ac2f5cf5eb7bdbcc10821b8953a4734b448060246e5bb93f130d6d3f2e28b9e04f2a064820be562274c040cd849f1473d45141559fc45da4c54abeaf5ca40d2d57f8f8e33bdb232c7279872f758b3fb452713b5d91c855383f7cec8376649a53b83951cf8edd519a99e91b8a6cb90153088e35d9fed332c7253771740f49f9dc40c7da50352656395bbfeae63e10f754d24a"; diff --git a/src/modules/module_26000.c b/src/modules/module_26000.c index 459bb3588..35a6dd937 100644 --- a/src/modules/module_26000.c +++ b/src/modules/module_26000.c @@ -22,8 +22,7 @@ static const u64 KERN_TYPE = 26000; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE - | OPTS_TYPE_PT_GENERATE_BE - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_PT_GENERATE_BE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$mozilla$*3DES*b735d19e6cadb5136376a98c2369f22819d08c79*2b36961682200a877f7d5550975b614acc9fefe3*f03f3575fd5bdbc9e32232316eab7623"; diff --git a/src/modules/module_27000.c b/src/modules/module_27000.c index 88f0695fa..672800eed 100644 --- a/src/modules/module_27000.c +++ b/src/modules/module_27000.c @@ -28,7 +28,6 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_ADDBITS14 | OPTS_TYPE_PT_UTF16LE | OPTS_TYPE_ST_HEX - | OPTS_TYPE_MAXIMUM_THREADS | OPTS_TYPE_AUTODETECT_DISABLE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "b4b9b02e6f09a9bd760f388b67351e2b"; diff --git a/src/modules/module_27100.c b/src/modules/module_27100.c index 65a1f8717..1c981fbc1 100644 --- a/src/modules/module_27100.c +++ b/src/modules/module_27100.c @@ -28,7 +28,6 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_ADDBITS14 | OPTS_TYPE_PT_UTF16LE | OPTS_TYPE_ST_HEX - | OPTS_TYPE_MAXIMUM_THREADS | OPTS_TYPE_AUTODETECT_DISABLE; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "b4b9b02e6f09a9bd760f388b67351e2b"; diff --git a/src/modules/module_29431.c b/src/modules/module_29431.c index eccc2beda..a22c24e3d 100644 --- a/src/modules/module_29431.c +++ b/src/modules/module_29431.c @@ -28,8 +28,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$veracrypt$48f79476aa0aa8327a8a9056e61450f4e2883c9e9669142f2e2f022c2f85303b897d088dea03d64329f6c402a56fed05b3919715929090a25c8ae84c67dbdb36$4ebfa3e9ccc0b391c130a4c3dd6495a1d6eb5d2eab72f8009096f7475ecb736bb3225b6da144e1596d859dad159fae5a739beea88ea074771e9d0b2d7c48ae302606a60d7cff6db54f3e460c548c06a4f47dc1ac203a8c8349fbff6a652219a63f27bc76327543e22be4f8dab8e4f90a4283fbf1552119fe24114ce8869eb20ce87dd72300f7aad3f7b4a26a355f16517725449151cf0373dbd0b281f6ac753485a14a5361cc75d40928e241a6b4684658801774843238048cf8c7f2fd88950abac040e12b0c41fdcaca3702907e951ec11c061a91b3050a4855abe6f3b50b4bd0b17c4be1f5b50b873eadc2d8446cd72c4fcac576bbce3acea769f740c5322ee8c927ffd4dd11c8a9e66f06e58df2e5d4d85c13b44c412bab839c9512b7a0acdd97b37dcccc4b70854eda0f36de12d62dd10cc13bc6154103d083bf6540bc78e5d0aad5d063cc74dad4cbe6e060febda2a9fd79c238f99dcb0766ff4addcfd0c03e619c765f65b1c75d5d22c6536958bcda78077ff44b64c4da741bf50154df310d4e0724238a777b524237b9478277e400ad8146dc3ca1da83e3d2f1c5115a4b7fcdc71dd7d56ba86a2f9b721c9a4137aabb07c3c5fedcf5342c4fae4898c9"; diff --git a/src/modules/module_29432.c b/src/modules/module_29432.c index 6e8de6aac..563b254c0 100644 --- a/src/modules/module_29432.c +++ b/src/modules/module_29432.c @@ -28,8 +28,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$veracrypt$1b721942019ebe8cedddbed7744a0702c0e053281a467e0ed69bf875c7406407d72eb8f2aea21270e41898c0a2c14382f86e04c15e7bc019d1d9dd813eabee0a$e5173e3cb1d927859d3e6de1006335a5184ae12b4c8dc2db2b1cd785063152a776f4dc5cacc1856a919b880d704b7450f5a0e0c9521bc9b4d67213c36a50e6664a1cbcea33f997b858e654111c7e9fca74f361528e85a28880381ec2600e3c1cd508c3833dd21cc91978185cba53caefd7b3c82d219d49f0b41e536d32e8d3ce194ad7923ca742213e19dcebdbd9687979d5a594654a5c611e8b829c4019e90a3cfb14e5fd7f8ed91e0fc79eed182399f02a3e3e202d4becaa6730e1f05f99ce06ce16dba7777ccddac72e85f2d3be5ecc9c808ac273f10ceb71cad666166abc327c4061a5f47424a5b6d9d093782f34b49924342a2e8cea663446ed4232a9a415ee2dfde988fa827b06d7438fec20ad0689543c3ee4602ce3ec3806fc7d668ef7e34330edd1e077b329a7627fa3ae5c89308258a17ecefbee114c80c2ab06f8271f14de8f2d13d1d6e5a119b71a6bae88ab151f76cdb2442284bc481d0df7e2163c3acfe763d3968195450d275af9034a00184a30cefed163e636626bffe6a35df3472508a49cb2b9b4c4a95d11c5d17e4e0539e9f13112125515778bcd1c2813c62a02673663062ad60583ec6a02c8a572865829e5b8c767b285728bea4907"; diff --git a/src/modules/module_29433.c b/src/modules/module_29433.c index b746e43f9..90e2033c9 100644 --- a/src/modules/module_29433.c +++ b/src/modules/module_29433.c @@ -28,8 +28,7 @@ static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED | OPTS_TYPE_MP_MULTI_DISABLE - | OPTS_TYPE_COPY_TMPS - | OPTS_TYPE_MAXIMUM_THREADS; + | OPTS_TYPE_COPY_TMPS; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; static const char *ST_PASS = "hashcat"; static const char *ST_HASH = "$veracrypt$5eb128daef63eff7e6db6aa10a8858f89964f47844acca68df82ebb2e73866fa75e3b7a53f9d2ff1ecdd1f4dc90e9c0fdf51f60d11b1992cd2971b4889edfc89$20bbf346fd7693f675b617cb9e4e9a43e6f445021068fc13453b130f2eb1d753ee83ecc61dabec293e88b62110cf6a8fab670e171f6aba2226550b54893263f5fa086b3cc41dd3db2eae07b585e5162c7a0d9723a426d408d83266c4d6018dc1b8b456d28a224033a30bfe62b1e58c2ddf596e07f7ff31849a6f5cfcc1c977b82d8484c270d44ededb0afdb781295e92968fc8cc69766af0ce1e72f02d6b4e124ba4b1af71519dcaade857bb3f371f93a350da6e65ee46c2ac782f134c75c10fe9d653fccc08c614dc362871911af8b83bdfc479f770dfe4b3c86b5d895842c53852fe4912738f848bf7c3e10b8189d25faceab9ef30b6fa0284edaa471752ac2b65335179b8d605417709f64fded7d94383618a921660d4cdb190bbb3769a8e56d2cd1ee07078ebc3b68ebeb016893f7099018e40cb326e32b29a62806eaf1a3fd382f4f876bf721eadfc019c5545813e81fd7168995f743663b136762b07910a63b6eec5b728a4ad07a689cceecb14c2802f334401a0a4fd2ec49e2da7f3cb24d6181f01ceed93ee73dedc3378133c83c9a71155c86785ff20dd5a64323d2fd4bf076bab3c17a1bb45edf81c30a7bd7dbbb097ece0dca83fff9138d56ae668"; diff --git a/src/selftest.c b/src/selftest.c index f237aa94b..f1b2b2dd5 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -516,20 +516,20 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param { if (highest_pw_len < 16) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, 1, false, 0, false) == -1) return -1; } else if (highest_pw_len < 32) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, 1, false, 0, false) == -1) return -1; } else { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, 0, 1, false, 0, false) == -1) return -1; } } else { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_4, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_4, 0, 1, false, 0, false) == -1) return -1; } } else @@ -563,12 +563,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->opts_type & OPTS_TYPE_INIT) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, 1, false, 0, false) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_HOOK12) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_12, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_12, 0, 1, false, 0, false) == -1) return -1; if (device_param->is_cuda == true) { @@ -636,7 +636,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, 1, false, 0, false) == -1) return -1; } const u32 iter = salt_buf->salt_iter; @@ -652,12 +652,12 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->opts_type & OPTS_TYPE_LOOP) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, 1, false, 0, false) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_LOOP_EXTENDED) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2E, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_2E, 0, 1, false, 0, false) == -1) return -1; } if (hashconfig->bridge_type & BRIDGE_TYPE_LAUNCH_LOOP) @@ -739,7 +739,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->opts_type & OPTS_TYPE_HOOK23) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_23, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_23, 0, 1, false, 0, false) == -1) return -1; if (device_param->is_cuda == true) { @@ -796,7 +796,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->opts_type & OPTS_TYPE_INIT2) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_INIT2, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_INIT2, 0, 1, false, 0, false) == -1) return -1; } for (u32 salt_repeat = 0; salt_repeat <= salt_repeats; salt_repeat++) @@ -805,7 +805,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->opts_type & OPTS_TYPE_LOOP2_PREPARE) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_LOOP2P, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_LOOP2P, 0, 1, false, 0, false) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_LOOP2) @@ -821,7 +821,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_param.loop_pos = loop_pos; device_param->kernel_param.loop_cnt = loop_left; - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_LOOP2, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_LOOP2, 0, 1, false, 0, false) == -1) return -1; if (hashconfig->bridge_type & BRIDGE_TYPE_LAUNCH_LOOP2) { @@ -904,28 +904,28 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->opts_type & OPTS_TYPE_AUX1) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX1, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX1, 0, 1, false, 0, false) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_AUX2) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX2, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX2, 0, 1, false, 0, false) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_AUX3) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX3, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX3, 0, 1, false, 0, false) == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_AUX4) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX4, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX4, 0, 1, false, 0, false) == -1) return -1; } } if (hashconfig->opts_type & OPTS_TYPE_COMP) { - if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, 0, 1, false, 0) == -1) return -1; + if (run_kernel (hashcat_ctx, device_param, KERN_RUN_3, 0, 1, false, 0, false) == -1) return -1; } } diff --git a/src/user_options.c b/src/user_options.c index 5de443725..23e195ee6 100644 --- a/src/user_options.c +++ b/src/user_options.c @@ -1042,7 +1042,7 @@ int user_options_sanity (hashcat_ctx_t *hashcat_ctx) return -1; } - if (user_options->kernel_loops > 1024) + if (user_options->kernel_loops > KERNEL_LOOPS_MAX) { event_log_error (hashcat_ctx, "Invalid kernel-loops specified.");