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;