From 0c2ed0d1991e97c1ae76006adfee656e637a93f2 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Sun, 29 Jun 2025 14:39:14 +0200 Subject: [PATCH] Update plugins that benefit from an artificially limited register count (NVIDIA). Update default hash settings to 64MiB:3:4 for Argon2 in -m 70000, following RFC 9106 recommendations. Add option OPTS_TYPE_THREAD_MULTI_DISABLE: allows plugin developers to disable scaling the password candidate batch size based on device thread count. This can be useful for super slow hash algorithms that utilize threads differently, e.g., when the algorithm allows parallelization. Note: thread count for the device can still be set normally. Add options OPTI_TYPE_SLOW_HASH_DIMY_INIT/LOOP/COMP: enable 2D launches for slow hash init/loop/comp kernel with dimensions X and Y. The Y value must be set via salt->salt_dimy attribute. Change autotune kernel-loops start value to the lowest multiple of the target hash iteration count, if kernel_loops_min permits. Fixed a bug in autotune where kernel_threads_max was not respected during initial init and loop-prepare kernel runs. --- OpenCL/inc_types.h | 1 + include/types.h | 18 ++++--- src/autotune.c | 22 ++++---- src/backend.c | 101 +++++++++++++++++++++++++++++------ src/modules/module_01800.c | 1 + src/modules/module_06400.c | 1 + src/modules/module_06800.c | 1 + src/modules/module_07700.c | 1 + src/modules/module_07701.c | 1 + src/modules/module_08300.c | 3 +- src/modules/module_13751.c | 3 +- src/modules/module_13752.c | 3 +- src/modules/module_13753.c | 3 +- src/modules/module_13761.c | 3 +- src/modules/module_13762.c | 3 +- src/modules/module_13763.c | 3 +- src/modules/module_14800.c | 1 + src/modules/module_14900.c | 3 +- src/modules/module_15600.c | 1 + src/modules/module_16300.c | 1 + src/modules/module_16700.c | 1 + src/modules/module_16900.c | 1 + src/modules/module_18100.c | 1 + src/modules/module_20800.c | 1 + src/modules/module_21300.c | 1 + src/modules/module_21500.c | 1 + src/modules/module_21700.c | 1 + src/modules/module_22100.c | 3 +- src/modules/module_22300.c | 1 + src/modules/module_22921.c | 3 +- src/modules/module_22941.c | 3 +- src/modules/module_23400.c | 1 + src/modules/module_23600.c | 1 + src/modules/module_23800.c | 3 +- src/modules/module_24200.c | 1 + src/modules/module_24420.c | 1 + src/modules/module_25500.c | 1 + src/modules/module_25900.c | 3 +- src/modules/module_26000.c | 1 + src/modules/module_26100.c | 1 + src/modules/module_26600.c | 1 + src/modules/module_26700.c | 3 +- src/modules/module_26800.c | 3 +- src/modules/module_26900.c | 1 + src/modules/module_27300.c | 1 + src/modules/module_27400.c | 1 + src/modules/module_27500.c | 1 + src/modules/module_27600.c | 1 + src/modules/module_29451.c | 3 +- src/modules/module_29452.c | 3 +- src/modules/module_29453.c | 3 +- src/modules/module_29461.c | 3 +- src/modules/module_29462.c | 3 +- src/modules/module_29463.c | 3 +- src/modules/module_70000.c | 2 +- src/shared.c | 6 +++ tools/test_modules/m70000.pm | 2 +- 57 files changed, 190 insertions(+), 54 deletions(-) diff --git a/OpenCL/inc_types.h b/OpenCL/inc_types.h index 233e28958..a13c89b8e 100644 --- a/OpenCL/inc_types.h +++ b/OpenCL/inc_types.h @@ -2008,6 +2008,7 @@ typedef struct salt u32 salt_len_pc; u32 salt_iter; u32 salt_iter2; + u32 salt_dimy; u32 salt_sign[2]; u32 salt_repeats; diff --git a/include/types.h b/include/types.h index 22c57d85a..600e7bd93 100644 --- a/include/types.h +++ b/include/types.h @@ -412,6 +412,9 @@ typedef enum opti_type OPTI_TYPE_REGISTER_LIMIT = (1 << 20), // We'll limit the register count to 128 OPTI_TYPE_SLOW_HASH_SIMD_INIT2 = (1 << 21), OPTI_TYPE_SLOW_HASH_SIMD_LOOP2 = (1 << 22), + OPTI_TYPE_SLOW_HASH_DIMY_INIT = (1 << 23), + OPTI_TYPE_SLOW_HASH_DIMY_LOOP = (1 << 24), + OPTI_TYPE_SLOW_HASH_DIMY_COMP = (1 << 25), } opti_type_t; @@ -476,14 +479,17 @@ typedef enum opts_type OPTS_TYPE_DYNAMIC_SHARED = (1ULL << 53), // use dynamic shared memory (note: needs special kernel changes) OPTS_TYPE_SELF_TEST_DISABLE = (1ULL << 54), // some algos use JiT in combinations with a salt or create too much startup time OPTS_TYPE_MP_MULTI_DISABLE = (1ULL << 55), // do not multiply the kernel-accel with the multiprocessor count per device to allow more fine-tuned workload settings - OPTS_TYPE_NATIVE_THREADS = (1ULL << 56), // forces "native" thread count: CPU=1, GPU-Intel=8, GPU-AMD=64 (wavefront), GPU-NV=32 (warps) - OPTS_TYPE_MAXIMUM_THREADS = (1ULL << 57), // disable else branch in pre-compilation thread count optimization setting - OPTS_TYPE_POST_AMP_UTF16LE = (1ULL << 58), // run the utf8 to utf16le conversion kernel after they have been processed from amplifiers + OPTS_TYPE_THREAD_MULTI_DISABLE // do not multiply the kernel-power with the thread count per device for super slow algos + = (1ULL << 56), + OPTS_TYPE_NATIVE_THREADS = (1ULL << 57), // forces "native" thread count: CPU=1, GPU-Intel=8, GPU-AMD=64 (wavefront), GPU-NV=32 (warps) + OPTS_TYPE_MAXIMUM_THREADS = (1ULL << 58), // disable else branch in pre-compilation thread count optimization setting + OPTS_TYPE_POST_AMP_UTF16LE = (1ULL << 59), // run the utf8 to utf16le conversion kernel after they have been processed from amplifiers OPTS_TYPE_AUTODETECT_DISABLE - = (1ULL << 59), // skip autodetect engine - OPTS_TYPE_STOCK_MODULE = (1ULL << 60), // module included with hashcat default distribution + = (1ULL << 60), // skip autodetect engine + OPTS_TYPE_STOCK_MODULE = (1ULL << 61), // module included with hashcat default distribution OPTS_TYPE_MULTIHASH_DESPITE_ESALT - = (1ULL << 61), // overrule multihash cracking check same salt but not same esalt + = (1ULL << 62), // overrule multihash cracking check same salt but not same esalt + OPTS_TYPE_MAXIMUM_ACCEL = (1ULL << 63) // try to maximize kernel-accel during autotune } opts_type_t; diff --git a/src/autotune.c b/src/autotune.c index 87637b29d..065c0a217 100644 --- a/src/autotune.c +++ b/src/autotune.c @@ -43,7 +43,8 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par device_param->kernel_param.loop_cnt = kernel_loops; // not a bug, both need to be set device_param->kernel_param.il_cnt = kernel_loops; // because there's two variables for inner iters for slow and fast hashes - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads); u32 kernel_power_try = hardware_power * kernel_accel; @@ -133,7 +134,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_accel = kernel_accel_min; device_param->kernel_loops = kernel_loops_min; device_param->kernel_threads = kernel_threads_min; - device_param->hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads_min; + device_param->hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads_min); device_param->kernel_power = device_param->hardware_power * kernel_accel_min; } @@ -212,7 +214,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param // from here it's clear we are allowed to autotune // so let's init some fake words - const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads_max; + const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : kernel_threads_max); u32 kernel_power_max = hardware_power_max * kernel_accel_max; @@ -298,13 +301,13 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param { const u32 kernel_threads_sav = device_param->kernel_threads; - device_param->kernel_threads = device_param->kernel_wgs1; + device_param->kernel_threads = MIN (device_param->kernel_wgs1, kernel_threads_max); 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; + device_param->kernel_threads = MIN (device_param->kernel_wgs2p, kernel_threads_max); run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, kernel_power_max, false, 0, true); } @@ -330,8 +333,6 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param // v7 autotuner is a lot more straight forward - u32 kernel_loops_min_start = kernel_loops_min; - if (hashes && hashes->st_salts_buf) { u32 start = kernel_loops_max; @@ -348,12 +349,12 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if ((start >= kernel_loops_min) && (start <= kernel_loops_max)) { - kernel_loops_min_start = start; + kernel_loops = start; } } } - for (u32 kernel_loops_test = kernel_loops_min_start; kernel_loops_test <= kernel_loops_max; kernel_loops_test <<= 1) + for (u32 kernel_loops_test = kernel_loops; kernel_loops_test <= kernel_loops_max; kernel_loops_test <<= 1) { double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_test, kernel_threads_min, 2); @@ -564,7 +565,8 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_loops = kernel_loops; device_param->kernel_threads = kernel_threads; - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * device_param->kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param->kernel_threads); device_param->hardware_power = hardware_power; diff --git a/src/backend.c b/src/backend.c index fc0051dd4..00cac2245 100644 --- a/src/backend.c +++ b/src/backend.c @@ -2598,7 +2598,10 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con if (kernel_threads == 0) kernel_threads = 1; - num_elements = CEILDIV (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = CEILDIV (num_elements, kernel_threads); + } if (kern_run == KERN_RUN_1) { @@ -2636,14 +2639,29 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + u32 gridDimX = num_elements; + u32 gridDimY = 1; + u32 gridDimZ = 1; + + u32 blockDimX = kernel_threads; + u32 blockDimY = 1; + u32 blockDimZ = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + 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_cuLaunchKernel (hashcat_ctx, cuda_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, 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; + if (hc_cuLaunchKernel (hashcat_ctx, cuda_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->cuda_stream, device_param->kernel_params, NULL) == -1) return -1; if (hc_cuEventRecord (hashcat_ctx, device_param->cuda_event2, device_param->cuda_stream) == -1) return -1; @@ -2699,7 +2717,10 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con if (kernel_threads == 0) kernel_threads = 1; - num_elements = CEILDIV (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = CEILDIV (num_elements, kernel_threads); + } if (kern_run == KERN_RUN_1) { @@ -2737,14 +2758,31 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } + u32 gridDimX = num_elements; + u32 gridDimY = 1; + u32 gridDimZ = 1; + + u32 blockDimX = kernel_threads; + u32 blockDimY = 1; + u32 blockDimZ = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + blockDimY = hashcat_ctx->hashes->salts_buf->salt_dimy; + + //printf ("%d %d %d %d %d %d %d\n", kern_run, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ); + 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_hipLaunchKernel (hashcat_ctx, hip_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, 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; + if (hc_hipLaunchKernel (hashcat_ctx, hip_function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, dynamic_shared_mem, device_param->hip_stream, device_param->kernel_params, NULL) == -1) return -1; if (hc_hipEventRecord (hashcat_ctx, device_param->hip_event2, device_param->hip_stream) == -1) return -1; @@ -2984,17 +3022,44 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con } } - num_elements = round_up_multiple_64 (num_elements, kernel_threads); + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + num_elements = round_up_multiple_64 (num_elements, kernel_threads); + } + else + { + num_elements = num_elements * kernel_threads; + } - const size_t global_work_size[3] = { num_elements, 1, 1 }; - const size_t local_work_size[3] = { kernel_threads, 1, 1 }; + size_t global_work_size[3] = { num_elements, 1, 1 }; + size_t local_work_size[3] = { kernel_threads, 1, 1 }; + + cl_uint work_dim = 1; + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_INIT) && (kern_run == KERN_RUN_1)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_LOOP) && (kern_run == KERN_RUN_2)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } + + if ((hashconfig->opti_type & OPTI_TYPE_SLOW_HASH_DIMY_COMP) && (kern_run == KERN_RUN_3)) + { + global_work_size[1] = local_work_size[1] = hashcat_ctx->hashes->salts_buf->salt_dimy; + work_dim = 2; + } 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, work_dim, 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; + if (hc_clEnqueueNDRangeKernel (hashcat_ctx, device_param->opencl_command_queue, opencl_kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &opencl_event) == -1) return -1; // spin damper section @@ -8952,7 +9017,8 @@ void backend_ctx_devices_sync_tuning (hashcat_ctx_t *hashcat_ctx) device_param_dst->kernel_loops = device_param_src->kernel_loops; device_param_dst->kernel_threads = device_param_src->kernel_threads; - const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param_dst->device_processors) * device_param_dst->kernel_threads; + const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param_dst->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param_dst->kernel_threads); device_param_dst->hardware_power = hardware_power; @@ -9522,7 +9588,11 @@ static bool load_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_p hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "-D MAX_THREADS_PER_BLOCK=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-architecture=%s", device_param->gcnArchName); - hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-max-threads-per-block=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); + + if ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) == 0) + { + hc_asprintf (&hiprtc_options[hiprtc_options_idx++], "--gpu-max-threads-per-block=%d", (user_options->kernel_threads_chgd == true) ? user_options->kernel_threads : device_param->kernel_threads_max); + } // untested but it should work #if defined (_WIN) || defined (__CYGWIN__) || defined (__MSYS__) @@ -10436,7 +10506,7 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) * device properties */ - const u32 device_processors = device_param->device_processors; + //const u32 device_processors = device_param->device_processors; /** * device threads @@ -16059,7 +16129,8 @@ int backend_session_begin (hashcat_ctx_t *hashcat_ctx) // device_param->kernel_threads = kernel_threads; device_param->kernel_threads = 0; - u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_processors) * device_param->kernel_threads_max; + const u32 hardware_power_max = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) + * ((hashconfig->opts_type & OPTS_TYPE_THREAD_MULTI_DISABLE) ? 1 : device_param->kernel_threads_max); u32 kernel_accel_min = device_param->kernel_accel_min; u32 kernel_accel_max = device_param->kernel_accel_max; diff --git a/src/modules/module_01800.c b/src/modules/module_01800.c index fea4851c6..f0d2e9623 100644 --- a/src/modules/module_01800.c +++ b/src/modules/module_01800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OS; static const char *HASH_NAME = "sha512crypt $6$, SHA512 (Unix)"; static const u64 KERN_TYPE = 1800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_06400.c b/src/modules/module_06400.c index b7ac9d87f..a553846c7 100644 --- a/src/modules/module_06400.c +++ b/src/modules/module_06400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OS; static const char *HASH_NAME = "AIX {ssha256}"; static const u64 KERN_TYPE = 6400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_06800.c b/src/modules/module_06800.c index 215f26992..e6cf96fd6 100644 --- a/src/modules/module_06800.c +++ b/src/modules/module_06800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "LastPass + LastPass sniffed"; static const u64 KERN_TYPE = 6800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_07700.c b/src/modules/module_07700.c index 2642a0bb7..1d783d351 100644 --- a/src/modules/module_07700.c +++ b/src/modules/module_07700.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SAP CODVN B (BCODE)"; static const u64 KERN_TYPE = 7700; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_PRECOMPUTE_INIT + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_07701.c b/src/modules/module_07701.c index bc852639b..232d7ca4d 100644 --- a/src/modules/module_07701.c +++ b/src/modules/module_07701.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SAP CODVN B (BCODE) from RFC_READ_TABLE"; static const u64 KERN_TYPE = 7701; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_PRECOMPUTE_INIT + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_08300.c b/src/modules/module_08300.c index c499667ae..80dd3475d 100644 --- a/src/modules/module_08300.c +++ b/src/modules/module_08300.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_5; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_SERVER; static const char *HASH_NAME = "DNSSEC (NSEC3)"; static const u64 KERN_TYPE = 8300; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +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_BE | OPTS_TYPE_ST_HEX diff --git a/src/modules/module_13751.c b/src/modules/module_13751.c index dac31a633..ee27acfd6 100644 --- a/src/modules/module_13751.c +++ b/src/modules/module_13751.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit (legacy)"; static const u64 KERN_TYPE = 13751; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13752.c b/src/modules/module_13752.c index e959a80af..095758dd2 100644 --- a/src/modules/module_13752.c +++ b/src/modules/module_13752.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit (legacy)"; static const u64 KERN_TYPE = 13752; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13753.c b/src/modules/module_13753.c index 277afaf02..4eee98625 100644 --- a/src/modules/module_13753.c +++ b/src/modules/module_13753.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit (legacy)"; static const u64 KERN_TYPE = 13753; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13761.c b/src/modules/module_13761.c index e33a693e2..5b1b82b27 100644 --- a/src/modules/module_13761.c +++ b/src/modules/module_13761.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13751; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13762.c b/src/modules/module_13762.c index 0f5f18545..6f1a27929 100644 --- a/src/modules/module_13762.c +++ b/src/modules/module_13762.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13752; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_13763.c b/src/modules/module_13763.c index c9e28d4c7..fb50002a2 100644 --- a/src/modules/module_13763.c +++ b/src/modules/module_13763.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit + boot-mode (legacy)"; static const u64 KERN_TYPE = 13753; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_BINARY_HASHFILE diff --git a/src/modules/module_14800.c b/src/modules/module_14800.c index 9f3c1ca0f..1e4b91a15 100644 --- a/src/modules/module_14800.c +++ b/src/modules/module_14800.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "iTunes backup >= 10.0"; static const u64 KERN_TYPE = 14800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_14900.c b/src/modules/module_14900.c index cbe009024..fcc639688 100644 --- a/src/modules/module_14900.c +++ b/src/modules/module_14900.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_CIPHER_KPA; static const char *HASH_NAME = "Skip32 (PT = $salt, key = $pass)"; static const u64 KERN_TYPE = 14900; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +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_SUGGEST_KG; diff --git a/src/modules/module_15600.c b/src/modules/module_15600.c index 1f2b7b9a3..66469d9bf 100644 --- a/src/modules/module_15600.c +++ b/src/modules/module_15600.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Ethereum Wallet, PBKDF2-HMAC-SHA256"; static const u64 KERN_TYPE = 15600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_16300.c b/src/modules/module_16300.c index e7507f1fa..6a12a1b39 100644 --- a/src/modules/module_16300.c +++ b/src/modules/module_16300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Ethereum Pre-Sale Wallet, PBKDF2-HMAC-SHA256"; static const u64 KERN_TYPE = 16300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_16700.c b/src/modules/module_16700.c index 9dce98e62..2fec2f557 100644 --- a/src/modules/module_16700.c +++ b/src/modules/module_16700.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "FileVault 2"; static const u64 KERN_TYPE = 16200; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_16900.c b/src/modules/module_16900.c index 443da6007..70eca3f16 100644 --- a/src/modules/module_16900.c +++ b/src/modules/module_16900.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Ansible Vault"; static const u64 KERN_TYPE = 16900; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_18100.c b/src/modules/module_18100.c index 46dadfe60..678cb8265 100644 --- a/src/modules/module_18100.c +++ b/src/modules/module_18100.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_OTP; static const char *HASH_NAME = "TOTP (HMAC-SHA1)"; static const u64 KERN_TYPE = 18100; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE diff --git a/src/modules/module_20800.c b/src/modules/module_20800.c index 6cece615d..1330b99bc 100644 --- a/src/modules/module_20800.c +++ b/src/modules/module_20800.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "sha256(md5($pass))"; static const u64 KERN_TYPE = 20800; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_21300.c b/src/modules/module_21300.c index 75e57c1f4..34ebff7f5 100644 --- a/src/modules/module_21300.c +++ b/src/modules/module_21300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "md5($salt.sha1($salt.$pass))"; static const u64 KERN_TYPE = 21300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_21500.c b/src/modules/module_21500.c index dba580a7c..8936be11c 100644 --- a/src/modules/module_21500.c +++ b/src/modules/module_21500.c @@ -21,6 +21,7 @@ static const char *HASH_NAME = "SolarWinds Orion"; static const u64 KERN_TYPE = 21500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_USES_BITS_64 + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_21700.c b/src/modules/module_21700.c index 65536027b..e00cf232f 100644 --- a/src/modules/module_21700.c +++ b/src/modules/module_21700.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Electrum Wallet (Salt-Type 4)"; static const u64 KERN_TYPE = 21700; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_22100.c b/src/modules/module_22100.c index 344218243..4d79d30bf 100644 --- a/src/modules/module_22100.c +++ b/src/modules/module_22100.c @@ -20,7 +20,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "BitLocker"; static const u64 KERN_TYPE = 22100; -static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP; +static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_MP_MULTI_DISABLE; diff --git a/src/modules/module_22300.c b/src/modules/module_22300.c index a85d1d753..04a4bc08c 100644 --- a/src/modules/module_22300.c +++ b/src/modules/module_22300.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_RAW_HASH_SALTED; static const char *HASH_NAME = "sha256($salt.$pass.$salt)"; static const u64 KERN_TYPE = 22300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_PRECOMPUTE_INIT | OPTI_TYPE_EARLY_SKIP | OPTI_TYPE_NOT_ITERATED diff --git a/src/modules/module_22921.c b/src/modules/module_22921.c index 163aacee9..c6471a232 100644 --- a/src/modules/module_22921.c +++ b/src/modules/module_22921.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "RSA/DSA/EC/OpenSSH Private Keys ($6$)"; static const u64 KERN_TYPE = 22921; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +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; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; diff --git a/src/modules/module_22941.c b/src/modules/module_22941.c index 4b4bf09ac..65e7d7c2b 100644 --- a/src/modules/module_22941.c +++ b/src/modules/module_22941.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "RSA/DSA/EC/OpenSSH Private Keys ($4$)"; static const u64 KERN_TYPE = 22941; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +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; static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED; diff --git a/src/modules/module_23400.c b/src/modules/module_23400.c index 13743815f..c2cc2a2e2 100644 --- a/src/modules/module_23400.c +++ b/src/modules/module_23400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Bitwarden"; static const u64 KERN_TYPE = 23400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_23600.c b/src/modules/module_23600.c index fdb3467d1..0b8994ff7 100644 --- a/src/modules/module_23600.c +++ b/src/modules/module_23600.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "AxCrypt 2 AES-256"; static const u64 KERN_TYPE = 23600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64 | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_23800.c b/src/modules/module_23800.c index c6c8fec61..c75a9aaaf 100644 --- a/src/modules/module_23800.c +++ b/src/modules/module_23800.c @@ -20,7 +20,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; // actually only DGST_SIZE_4_ static const u32 HASH_CATEGORY = HASH_CATEGORY_ARCHIVE; static const char *HASH_NAME = "RAR3-p (Compressed)"; static const u64 KERN_TYPE = 23800; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +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_HOOK23 diff --git a/src/modules/module_24200.c b/src/modules/module_24200.c index 8b6a8e5c8..20ea6ccec 100644 --- a/src/modules/module_24200.c +++ b/src/modules/module_24200.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_DATABASE_SERVER; static const char *HASH_NAME = "MongoDB ServerKey SCRAM-SHA-256"; static const u64 KERN_TYPE = 24200; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_24420.c b/src/modules/module_24420.c index 39820834a..544e07ee1 100644 --- a/src/modules/module_24420.c +++ b/src/modules/module_24420.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PRIVATE_KEY; static const char *HASH_NAME = "PKCS#8 Private Keys (PBKDF2-HMAC-SHA256 + 3DES/AES)"; static const u64 KERN_TYPE = 24420; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_25500.c b/src/modules/module_25500.c index afc451fb9..06a1e795d 100644 --- a/src/modules/module_25500.c +++ b/src/modules/module_25500.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "Stargazer Stellar Wallet XLM"; static const u64 KERN_TYPE = 25500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_25900.c b/src/modules/module_25900.c index 7ad951f6d..a460a3bd8 100644 --- a/src/modules/module_25900.c +++ b/src/modules/module_25900.c @@ -19,7 +19,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_SERVER; static const char *HASH_NAME = "KNX IP Secure - Device Authentication Code"; static const u64 KERN_TYPE = 25900; -static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP; +static const u32 OPTI_TYPE = OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_DEEP_COMP_KERNEL; diff --git a/src/modules/module_26000.c b/src/modules/module_26000.c index 35a6dd937..8ed90d5ee 100644 --- a/src/modules/module_26000.c +++ b/src/modules/module_26000.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Mozilla key3.db"; static const u64 KERN_TYPE = 26000; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_NOT_ITERATED; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_BE; diff --git a/src/modules/module_26100.c b/src/modules/module_26100.c index c42958289..3414fcbdb 100644 --- a/src/modules/module_26100.c +++ b/src/modules/module_26100.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_PASSWORD_MANAGER; static const char *HASH_NAME = "Mozilla key4.db"; static const u64 KERN_TYPE = 26100; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26600.c b/src/modules/module_26600.c index 926555e94..92b5c8efc 100644 --- a/src/modules/module_26600.c +++ b/src/modules/module_26600.c @@ -21,6 +21,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_CRYPTOCURRENCY_WALLET; static const char *HASH_NAME = "MetaMask Wallet (needs all data, checks AES-GCM tag)"; static const u64 KERN_TYPE = 26600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26700.c b/src/modules/module_26700.c index 63854c0d2..47427ba32 100644 --- a/src/modules/module_26700.c +++ b/src/modules/module_26700.c @@ -21,7 +21,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_4; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA224-128"; static const u64 KERN_TYPE = 26700; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26800.c b/src/modules/module_26800.c index 024bfa636..44d72f7df 100644 --- a/src/modules/module_26800.c +++ b/src/modules/module_26800.c @@ -21,7 +21,8 @@ static const u32 DGST_SIZE = DGST_SIZE_4_6; static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA256-192"; static const u64 KERN_TYPE = 26800; -static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE; +static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS | OPTS_TYPE_PT_GENERATE_LE; diff --git a/src/modules/module_26900.c b/src/modules/module_26900.c index 8d916ebf2..29e3cdc92 100644 --- a/src/modules/module_26900.c +++ b/src/modules/module_26900.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA384-256"; static const u64 KERN_TYPE = 26900; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS diff --git a/src/modules/module_27300.c b/src/modules/module_27300.c index ff20acc9d..4d607051e 100644 --- a/src/modules/module_27300.c +++ b/src/modules/module_27300.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_NETWORK_PROTOCOL; static const char *HASH_NAME = "SNMPv3 HMAC-SHA512-384"; static const u64 KERN_TYPE = 27300; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_USES_BITS_64; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_NATIVE_THREADS diff --git a/src/modules/module_27400.c b/src/modules/module_27400.c index 7d9f0fbd4..059b50cb5 100644 --- a/src/modules/module_27400.c +++ b/src/modules/module_27400.c @@ -20,6 +20,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VMware VMX (PBKDF2-HMAC-SHA1 + AES-256-CBC)"; static const u64 KERN_TYPE = 27400; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE diff --git a/src/modules/module_27500.c b/src/modules/module_27500.c index ebabe24e5..6b79eba3a 100644 --- a/src/modules/module_27500.c +++ b/src/modules/module_27500.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VirtualBox (PBKDF2-HMAC-SHA256 & AES-128-XTS)"; static const u64 KERN_TYPE = 27500; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_27600.c b/src/modules/module_27600.c index b4cdade54..3259ce05f 100644 --- a/src/modules/module_27600.c +++ b/src/modules/module_27600.c @@ -22,6 +22,7 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VirtualBox (PBKDF2-HMAC-SHA256 & AES-256-XTS)"; static const u64 KERN_TYPE = 27600; static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_REGISTER_LIMIT | OPTI_TYPE_SLOW_HASH_SIMD_LOOP | OPTI_TYPE_SLOW_HASH_SIMD_LOOP2; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE diff --git a/src/modules/module_29451.c b/src/modules/module_29451.c index c8e10bdcf..fe3c0737f 100644 --- a/src/modules/module_29451.c +++ b/src/modules/module_29451.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit"; static const u64 KERN_TYPE = 13751; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29452.c b/src/modules/module_29452.c index 395839f64..842f36f47 100644 --- a/src/modules/module_29452.c +++ b/src/modules/module_29452.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit"; static const u64 KERN_TYPE = 13752; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29453.c b/src/modules/module_29453.c index 97bf9a333..6162337de 100644 --- a/src/modules/module_29453.c +++ b/src/modules/module_29453.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit"; static const u64 KERN_TYPE = 13753; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29461.c b/src/modules/module_29461.c index be7dcfc2e..ff2705da7 100644 --- a/src/modules/module_29461.c +++ b/src/modules/module_29461.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 512 bit + boot-mode"; static const u64 KERN_TYPE = 13751; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29462.c b/src/modules/module_29462.c index 7e9a6535c..14b53d6ad 100644 --- a/src/modules/module_29462.c +++ b/src/modules/module_29462.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1024 bit + boot-mode"; static const u64 KERN_TYPE = 13752; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_29463.c b/src/modules/module_29463.c index e98c71289..7a5c8ac92 100644 --- a/src/modules/module_29463.c +++ b/src/modules/module_29463.c @@ -23,7 +23,8 @@ static const u32 HASH_CATEGORY = HASH_CATEGORY_FDE; static const char *HASH_NAME = "VeraCrypt SHA256 + XTS 1536 bit + boot-mode"; static const u64 KERN_TYPE = 13753; // old kernel used here static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE - | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP + | OPTI_TYPE_REGISTER_LIMIT; static const u64 OPTS_TYPE = OPTS_TYPE_STOCK_MODULE | OPTS_TYPE_PT_GENERATE_LE | OPTS_TYPE_LOOP_EXTENDED diff --git a/src/modules/module_70000.c b/src/modules/module_70000.c index 01504ea09..57bd77d12 100644 --- a/src/modules/module_70000.c +++ b/src/modules/module_70000.c @@ -29,7 +29,7 @@ static const u64 BRIDGE_TYPE = BRIDGE_TYPE_MATCH_TUNINGS // optional - impr | BRIDGE_TYPE_REPLACE_LOOP; static const char *BRIDGE_NAME = "argon2id_reference"; static const char *ST_PASS = "hashcat"; -static const char *ST_HASH = "$argon2id$v=19$m=4096,t=3,p=1$FoIjFnZlM2JSJWYXUgMFAw$eYKMzhbW8uyT1LLtKRdRcJj2CQeRrdr2pKv/Y71YbAQ"; +static const char *ST_HASH = "$argon2id$v=19$m=65536,t=3,p=1$FBMjI4RJBhIykCgol1KEJA$2ky5GAdhT1kH4kIgPN/oERE3Taiy43vNN70a3HpiKQU"; u32 module_attack_exec (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC; } u32 module_dgst_pos0 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0; } diff --git a/src/shared.c b/src/shared.c index a00025352..f7dbcd78d 100644 --- a/src/shared.c +++ b/src/shared.c @@ -104,6 +104,9 @@ static const char *const OPTI_STR_USES_BITS_8 = "Uses-8-Bit"; static const char *const OPTI_STR_USES_BITS_16 = "Uses-16-Bit"; static const char *const OPTI_STR_USES_BITS_32 = "Uses-32-Bit"; static const char *const OPTI_STR_USES_BITS_64 = "Uses-64-Bit"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_INIT = "Slow-Hash-DimensionY-INIT"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_COMP = "Slow-Hash-DimensionY-LOOP"; +static const char *const OPTI_STR_SLOW_HASH_DIMY_LOOP = "Slow-Hash-DimensionY-COMP"; static const char *const HASH_CATEGORY_UNDEFINED_STR = "Undefined"; static const char *const HASH_CATEGORY_RAW_HASH_STR = "Raw Hash"; @@ -1072,6 +1075,9 @@ const char *stroptitype (const u32 opti_type) case OPTI_TYPE_SLOW_HASH_SIMD_LOOP: return OPTI_STR_SLOW_HASH_SIMD_LOOP; case OPTI_TYPE_SLOW_HASH_SIMD_LOOP2: return OPTI_STR_SLOW_HASH_SIMD_LOOP2; case OPTI_TYPE_SLOW_HASH_SIMD_COMP: return OPTI_STR_SLOW_HASH_SIMD_COMP; + case OPTI_TYPE_SLOW_HASH_DIMY_INIT: return OPTI_STR_SLOW_HASH_DIMY_INIT; + case OPTI_TYPE_SLOW_HASH_DIMY_LOOP: return OPTI_STR_SLOW_HASH_DIMY_LOOP; + case OPTI_TYPE_SLOW_HASH_DIMY_COMP: return OPTI_STR_SLOW_HASH_DIMY_COMP; case OPTI_TYPE_USES_BITS_8: return OPTI_STR_USES_BITS_8; case OPTI_TYPE_USES_BITS_16: return OPTI_STR_USES_BITS_16; case OPTI_TYPE_USES_BITS_32: return OPTI_STR_USES_BITS_32; diff --git a/tools/test_modules/m70000.pm b/tools/test_modules/m70000.pm index 4b44334ac..72861ae49 100644 --- a/tools/test_modules/m70000.pm +++ b/tools/test_modules/m70000.pm @@ -17,7 +17,7 @@ sub module_generate_hash { my $word = shift; my $salt = shift; - my $m = shift // 4096; + my $m = shift // 65536; my $t = shift // 3; my $p = shift // 1; my $len = shift // random_number (1, 2) * 16;