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.");