From 6cbb5e4992504c7765b8f5a4a773cc4a6c573fda Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 5 Jan 2019 18:24:37 +0100 Subject: [PATCH] Refactor kernel_accel, kernel_loops and kernel_threads management in combination with modules --- include/interface.h | 19 +++- include/modules.h | 9 +- include/types.h | 21 +++-- modules/module_01000.c | 10 +- src/interface.c | 57 +++++++++--- src/interface_migrate.c | 48 +++++++++- src/opencl.c | 202 +++++++++++++++++++++++++++++----------- src/selftest.c | 6 +- 8 files changed, 281 insertions(+), 91 deletions(-) diff --git a/include/interface.h b/include/interface.h index 30f0eaed6..a33b71f78 100644 --- a/include/interface.h +++ b/include/interface.h @@ -257,6 +257,16 @@ typedef struct luks } luks_t; +typedef enum kernel_workload +{ + KERNEL_ACCEL_MIN = 1, + KERNEL_ACCEL_MAX = 1024, + KERNEL_LOOPS_MIN = 1, + KERNEL_LOOPS_MAX = 1024, + KERNEL_THREADS_MIN = 1, + KERNEL_THREADS_MAX = 1024, + +} kernel_workload_t; /** * output functions @@ -285,9 +295,6 @@ u32 default_dgst_pos2 (MAYBE_UNUSED const hashconfig_t *ha u32 default_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u32 default_dgst_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 default_esalt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); -u32 default_forced_kernel_accel (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); -u32 default_forced_kernel_loops (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); -u32 default_forced_kernel_threads (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u32 default_forced_outfile_format (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); const char *default_hash_name (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u32 default_hash_mode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); @@ -295,6 +302,12 @@ u32 default_hash_type (MAYBE_UNUSED const hashconfig_t *ha bool default_hlfmt_disable (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 default_hook_salt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 default_hook_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 default_kernel_accel_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 default_kernel_accel_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 default_kernel_loops_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 default_kernel_loops_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 default_kernel_threads_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 default_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 default_kern_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u32 default_opti_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 default_opts_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); diff --git a/include/modules.h b/include/modules.h index e97c7d5a1..56113fa57 100644 --- a/include/modules.h +++ b/include/modules.h @@ -16,9 +16,6 @@ u32 module_dgst_pos2 (MAYBE_UNUSED const hashconfig_t *ha u32 module_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u32 module_dgst_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 module_esalt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); -u32 module_forced_kernel_accel (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); -u32 module_forced_kernel_loops (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); -u32 module_forced_kernel_threads (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u32 module_forced_outfile_format (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); const char *module_hash_name (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u32 module_hash_mode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); @@ -26,6 +23,12 @@ u32 module_hash_type (MAYBE_UNUSED const hashconfig_t *ha bool module_hlfmt_disable (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 module_hook_salt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 module_hook_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 module_kernel_accel_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 module_kernel_accel_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 module_kernel_loops_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 module_kernel_loops_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 module_kernel_threads_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); +u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 module_kern_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u32 module_opti_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); u64 module_opts_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra); diff --git a/include/types.h b/include/types.h index c85ab4542..cdfbe0012 100644 --- a/include/types.h +++ b/include/types.h @@ -953,9 +953,13 @@ struct hashconfig void *benchmark_esalt; void *benchmark_hook_salt; - u32 forced_kernel_accel; - u32 forced_kernel_loops; - u32 forced_kernel_threads; + u32 kernel_accel_min; + u32 kernel_accel_max; + u32 kernel_loops_min; + u32 kernel_loops_max; + u32 kernel_threads_min; + u32 kernel_threads_max; + u32 forced_outfile_format; bool dictstat_disable; @@ -1148,6 +1152,8 @@ typedef struct hc_device_param u32 kernel_loops_min_sav; // the _sav are required because each -i iteration u32 kernel_loops_max_sav; // needs to recalculate the kernel_loops_min/max based on the current amplifier count u32 kernel_threads; + u32 kernel_threads_min; + u32 kernel_threads_max; u64 kernel_power; u64 hardware_power; @@ -2255,9 +2261,6 @@ typedef struct module_ctx u32 (*module_dgst_pos3) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u32 (*module_dgst_size) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u64 (*module_esalt_size) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); - u32 (*module_forced_kernel_accel) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); - u32 (*module_forced_kernel_loops) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); - u32 (*module_forced_kernel_threads) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u32 (*module_forced_outfile_format) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); const char *(*module_hash_name) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u32 (*module_hash_mode) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); @@ -2265,6 +2268,12 @@ typedef struct module_ctx bool (*module_hlfmt_disable) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u64 (*module_hook_salt_size) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u64 (*module_hook_size) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); + u32 (*module_kernel_accel_min) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); + u32 (*module_kernel_accel_max) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); + u32 (*module_kernel_loops_min) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); + u32 (*module_kernel_loops_max) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); + u32 (*module_kernel_threads_min) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); + u32 (*module_kernel_threads_max) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u64 (*module_kern_type) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u32 (*module_opti_type) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); u64 (*module_opts_type) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); diff --git a/modules/module_01000.c b/modules/module_01000.c index 2fd7d4613..ce7d4d513 100644 --- a/modules/module_01000.c +++ b/modules/module_01000.c @@ -115,6 +115,7 @@ void module_init (module_ctx_t *module_ctx) { // undefined functions automatically call corresponding default functions + module_ctx->module_attack_exec = module_attack_exec; module_ctx->module_benchmark_esalt = NULL; module_ctx->module_benchmark_hook_salt = NULL; @@ -129,9 +130,6 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_dictstat_disable = NULL; module_ctx->module_esalt_size = NULL; module_ctx->module_extra_buffer_size = NULL; - module_ctx->module_forced_kernel_accel = NULL; - module_ctx->module_forced_kernel_loops = NULL; - module_ctx->module_forced_kernel_threads = NULL; module_ctx->module_forced_outfile_format = NULL; module_ctx->module_hash_decode = module_hash_decode; module_ctx->module_hash_decode_outfile = NULL; @@ -146,6 +144,12 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_hook_salt_size = NULL; module_ctx->module_hook_size = NULL; module_ctx->module_jit_build_options = NULL; + module_ctx->module_kernel_accel_max = NULL; + module_ctx->module_kernel_accel_min = NULL; + module_ctx->module_kernel_loops_max = NULL; + module_ctx->module_kernel_loops_min = NULL; + module_ctx->module_kernel_threads_max = NULL; + module_ctx->module_kernel_threads_min = NULL; module_ctx->module_kern_type = module_kern_type; module_ctx->module_opti_type = module_opti_type; module_ctx->module_opts_type = module_opts_type; diff --git a/src/interface.c b/src/interface.c index 2b40c6708..1eb9f4776 100644 --- a/src/interface.c +++ b/src/interface.c @@ -674,9 +674,6 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->dgst_size = default_dgst_size (hashconfig, user_options, user_options_extra); hashconfig->dictstat_disable = default_dictstat_disable (hashconfig, user_options, user_options_extra); hashconfig->esalt_size = default_esalt_size (hashconfig, user_options, user_options_extra); - hashconfig->forced_kernel_accel = default_forced_kernel_accel (hashconfig, user_options, user_options_extra); - hashconfig->forced_kernel_loops = default_forced_kernel_loops (hashconfig, user_options, user_options_extra); - hashconfig->forced_kernel_threads = default_forced_kernel_threads (hashconfig, user_options, user_options_extra); hashconfig->forced_outfile_format = default_forced_outfile_format (hashconfig, user_options, user_options_extra); hashconfig->hash_mode = default_hash_mode (hashconfig, user_options, user_options_extra); hashconfig->hash_name = default_hash_name (hashconfig, user_options, user_options_extra); @@ -684,6 +681,12 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->hlfmt_disable = default_hlfmt_disable (hashconfig, user_options, user_options_extra); hashconfig->hook_salt_size = default_hook_salt_size (hashconfig, user_options, user_options_extra); hashconfig->hook_size = default_hook_size (hashconfig, user_options, user_options_extra); + hashconfig->kernel_accel_min = default_kernel_accel_min (hashconfig, user_options, user_options_extra); + hashconfig->kernel_accel_max = default_kernel_accel_max (hashconfig, user_options, user_options_extra); + hashconfig->kernel_loops_min = default_kernel_loops_min (hashconfig, user_options, user_options_extra); + hashconfig->kernel_loops_max = default_kernel_loops_max (hashconfig, user_options, user_options_extra); + hashconfig->kernel_threads_min = default_kernel_threads_min (hashconfig, user_options, user_options_extra); + hashconfig->kernel_threads_max = default_kernel_threads_max (hashconfig, user_options, user_options_extra); hashconfig->kern_type = default_kern_type (hashconfig, user_options, user_options_extra); hashconfig->opti_type = default_opti_type (hashconfig, user_options, user_options_extra); hashconfig->opts_type = default_opts_type (hashconfig, user_options, user_options_extra); @@ -723,9 +726,6 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) if (module_ctx->module_dgst_size) hashconfig->dgst_size = module_ctx->module_dgst_size (hashconfig, user_options, user_options_extra); if (module_ctx->module_dictstat_disable) hashconfig->dictstat_disable = module_ctx->module_dictstat_disable (hashconfig, user_options, user_options_extra); if (module_ctx->module_esalt_size) hashconfig->esalt_size = module_ctx->module_esalt_size (hashconfig, user_options, user_options_extra); - if (module_ctx->module_forced_kernel_accel) hashconfig->forced_kernel_accel = module_ctx->module_forced_kernel_accel (hashconfig, user_options, user_options_extra); - if (module_ctx->module_forced_kernel_loops) hashconfig->forced_kernel_loops = module_ctx->module_forced_kernel_loops (hashconfig, user_options, user_options_extra); - if (module_ctx->module_forced_kernel_threads) hashconfig->forced_kernel_threads = module_ctx->module_forced_kernel_threads (hashconfig, user_options, user_options_extra); if (module_ctx->module_forced_outfile_format) hashconfig->forced_outfile_format = module_ctx->module_forced_outfile_format (hashconfig, user_options, user_options_extra); if (module_ctx->module_hash_mode) hashconfig->hash_mode = module_ctx->module_hash_mode (hashconfig, user_options, user_options_extra); if (module_ctx->module_hash_name) hashconfig->hash_name = module_ctx->module_hash_name (hashconfig, user_options, user_options_extra); @@ -733,6 +733,12 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) if (module_ctx->module_hlfmt_disable) hashconfig->hlfmt_disable = module_ctx->module_hlfmt_disable (hashconfig, user_options, user_options_extra); if (module_ctx->module_hook_salt_size) hashconfig->hook_salt_size = module_ctx->module_hook_salt_size (hashconfig, user_options, user_options_extra); if (module_ctx->module_hook_size) hashconfig->hook_size = module_ctx->module_hook_size (hashconfig, user_options, user_options_extra); + if (module_ctx->module_kernel_accel_min) hashconfig->kernel_accel_min = module_ctx->module_kernel_accel_min (hashconfig, user_options, user_options_extra); + if (module_ctx->module_kernel_accel_max) hashconfig->kernel_accel_max = module_ctx->module_kernel_accel_max (hashconfig, user_options, user_options_extra); + if (module_ctx->module_kernel_loops_min) hashconfig->kernel_loops_min = module_ctx->module_kernel_loops_min (hashconfig, user_options, user_options_extra); + if (module_ctx->module_kernel_loops_max) hashconfig->kernel_loops_max = module_ctx->module_kernel_loops_max (hashconfig, user_options, user_options_extra); + if (module_ctx->module_kernel_threads_min) hashconfig->kernel_threads_min = module_ctx->module_kernel_threads_min (hashconfig, user_options, user_options_extra); + if (module_ctx->module_kernel_threads_max) hashconfig->kernel_threads_max = module_ctx->module_kernel_threads_max (hashconfig, user_options, user_options_extra); if (module_ctx->module_kern_type) hashconfig->kern_type = module_ctx->module_kern_type (hashconfig, user_options, user_options_extra); if (module_ctx->module_opti_type) hashconfig->opti_type = module_ctx->module_opti_type (hashconfig, user_options, user_options_extra); if (module_ctx->module_opts_type) hashconfig->opts_type = module_ctx->module_opts_type (hashconfig, user_options, user_options_extra); @@ -1306,25 +1312,46 @@ u64 default_esalt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSE return esalt_size; } -u32 default_forced_kernel_threads (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +u32 default_kernel_accel_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { - const u32 forced_kernel_threads = user_options->kernel_threads; + const u32 kernel_accel_min = KERNEL_ACCEL_MIN; - return forced_kernel_threads; + return kernel_accel_min; } -u32 default_forced_kernel_loops (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +u32 default_kernel_accel_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { - const u32 forced_kernel_loops = user_options->kernel_loops; + const u32 kernel_accel_max = KERNEL_ACCEL_MAX; - return forced_kernel_loops; + return kernel_accel_max; } -u32 default_forced_kernel_accel (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +u32 default_kernel_loops_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { - const u32 forced_kernel_accel = user_options->kernel_accel; + const u32 kernel_loops_min = KERNEL_LOOPS_MIN; - return forced_kernel_accel; + return kernel_loops_min; +} + +u32 default_kernel_loops_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_loops_max = KERNEL_LOOPS_MAX; + + return kernel_loops_max; +} + +u32 default_kernel_threads_min (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_threads_min = KERNEL_THREADS_MIN; + + return kernel_threads_min; +} + +u32 default_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) +{ + const u32 kernel_threads_max = KERNEL_THREADS_MAX; + + return kernel_threads_max; } u32 default_forced_outfile_format (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) diff --git a/src/interface_migrate.c b/src/interface_migrate.c index 51701458f..296087945 100644 --- a/src/interface_migrate.c +++ b/src/interface_migrate.c @@ -18094,7 +18094,7 @@ int apfs_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE -u32 hashconfig_forced_kernel_threads (hashcat_ctx_t *hashcat_ctx) +u32 kernel_threads_mxx (hashcat_ctx_t *hashcat_ctx) { @@ -18115,13 +18115,17 @@ u32 hashconfig_forced_kernel_threads (hashcat_ctx_t *hashcat_ctx) if (hashconfig->hash_mode == 18200) kernel_threads = 64; // RC4 if (hashconfig->hash_mode == 18600) kernel_threads = 8; // Blowfish + // let the module decide if it allows user-defined values over module defined valaues + if (user_options->kernel_threads_chgd == true) + { + device_param->kernel_threads_min = user_options->kernel_threads; + device_param->kernel_threads_max = user_options->kernel_threads; + } } - - -u32 hashconfig_get_kernel_loops (hashcat_ctx_t *hashcat_ctx) +u32 kernel_loops_mxx (hashcat_ctx_t *hashcat_ctx) { if (user_options->slow_candidates == true) @@ -18170,8 +18174,44 @@ u32 hashconfig_get_kernel_loops (hashcat_ctx_t *hashcat_ctx) kernel_loops_fixed = 1; } + + // let the module decide if it allows user-defined values over module defined valaues + + // commandline parameters overwrite tuningdb entries + + if (user_options->kernel_loops_chgd == true) + { + device_param->kernel_loops_min = user_options->kernel_loops; + device_param->kernel_loops_max = user_options->kernel_loops; + } + } +u32 kernel_accel_mxx (hashcat_ctx_t *hashcat_ctx) +{ + // limit scrypt accel otherwise we hurt ourself when calculating the scrypt tmto + + #define SCRYPT_MAX_ACCEL 16 + + if ((hashconfig->hash_mode == 8900) || (hashconfig->hash_mode == 9300) || (hashconfig->hash_mode == 15700)) + { + // 16 is actually a bit low, we may need to change this depending on user response + + device_param->kernel_accel_max = MIN (device_param->kernel_accel_max, SCRYPT_MAX_ACCEL); + } + + + // let the module decide if it allows user-defined values over module defined valaues + + // commandline parameters overwrite tuningdb entries + + if (user_options->kernel_accel_chgd == true) + { + device_param->kernel_accel_min = user_options->kernel_accel; + device_param->kernel_accel_max = user_options->kernel_accel; + } + +} void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, void *esalt, void *hook_salt) { diff --git a/src/opencl.c b/src/opencl.c index 706d3925b..be0fe2f94 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -4073,8 +4073,13 @@ void opencl_ctx_devices_sync_tuning (hashcat_ctx_t *hashcat_ctx) if (is_same_device_type (device_param_src, device_param_dst) == false) continue; - device_param_dst->kernel_accel = device_param_src->kernel_accel; - device_param_dst->kernel_loops = device_param_src->kernel_loops; + device_param_dst->kernel_accel = device_param_src->kernel_accel; + device_param_dst->kernel_loops = device_param_src->kernel_loops; + device_param_dst->kernel_threads = device_param_src->kernel_threads; + + const u32 hardware_power = device_param_dst->device_processors * device_param_dst->kernel_threads; + + device_param_dst->hardware_power = hardware_power; const u32 kernel_power = device_param_dst->hardware_power * device_param_dst->kernel_accel; @@ -4242,19 +4247,30 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param { const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; - // a kernel can force a fixed value - - const u32 forced_kernel_threads = hashconfig->forced_kernel_threads; + // a module can force a fixed value - if (forced_kernel_threads) return forced_kernel_threads; + u32 kernel_threads_min = device_param->kernel_threads_min; + u32 kernel_threads_max = device_param->kernel_threads_max; // for CPU we just do 1 ... - if (device_param->device_type & CL_DEVICE_TYPE_CPU) return 1; + if (device_param->device_type & CL_DEVICE_TYPE_CPU) + { + if ((kernel_threads_min >= 1) && (kernel_threads_max <= 1)) + { + kernel_threads_min = 1; + kernel_threads_max = 1; + } + } // this is an upper limit, a good start, since our strategy is to reduce thread counts only - u32 kernel_threads = (u32) device_param->device_maxworkgroup_size; + const u32 device_maxworkgroup_size = (u32) device_param->device_maxworkgroup_size; + + if (device_maxworkgroup_size < kernel_threads_max) + { + kernel_threads_max = device_maxworkgroup_size; + } // complicated kernel tend to confuse OpenCL runtime suggestions for maximum thread size // let's workaround that by sticking to their device specific preferred thread size @@ -4265,16 +4281,43 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param { if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { - if (device_param->kernel_preferred_wgs_multiple1) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple1); + if (device_param->kernel_preferred_wgs_multiple1) + { + const u32 kernel_preferred_wgs_multiple1 = device_param->kernel_preferred_wgs_multiple1; + + if ((kernel_preferred_wgs_multiple1 >= kernel_threads_min) && (kernel_preferred_wgs_multiple1 <= kernel_threads_max)) + { + kernel_threads_min = kernel_preferred_wgs_multiple1; + kernel_threads_max = kernel_preferred_wgs_multiple1; + } + } } else { - if (device_param->kernel_preferred_wgs_multiple4) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple4); + if (device_param->kernel_preferred_wgs_multiple4) + { + const u32 kernel_preferred_wgs_multiple4 = device_param->kernel_preferred_wgs_multiple4; + + if ((kernel_preferred_wgs_multiple4 >= kernel_threads_min) && (kernel_preferred_wgs_multiple4 <= kernel_threads_max)) + { + kernel_threads_min = kernel_preferred_wgs_multiple4; + kernel_threads_max = kernel_preferred_wgs_multiple4; + } + } } } else { - if (device_param->kernel_preferred_wgs_multiple2) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple2); + if (device_param->kernel_preferred_wgs_multiple2) + { + const u32 kernel_preferred_wgs_multiple2 = device_param->kernel_preferred_wgs_multiple2; + + if ((kernel_preferred_wgs_multiple2 >= kernel_threads_min) && (kernel_preferred_wgs_multiple2 <= kernel_threads_max)) + { + kernel_threads_min = kernel_preferred_wgs_multiple2; + kernel_threads_max = kernel_preferred_wgs_multiple2; + } + } } } else @@ -4283,20 +4326,47 @@ static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param { if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) { - if (device_param->kernel_wgs1) kernel_threads = MIN (kernel_threads, device_param->kernel_wgs1); + if (device_param->kernel_preferred_wgs_multiple1) + { + const u32 kernel_preferred_wgs_multiple1 = device_param->kernel_preferred_wgs_multiple1; + + if ((kernel_preferred_wgs_multiple1 >= kernel_threads_min) && (kernel_preferred_wgs_multiple1 <= kernel_threads_max)) + { + kernel_threads_min = kernel_preferred_wgs_multiple1; + kernel_threads_max = kernel_preferred_wgs_multiple1; + } + } } else { - if (device_param->kernel_wgs4) kernel_threads = MIN (kernel_threads, device_param->kernel_wgs4); + if (device_param->kernel_preferred_wgs_multiple4) + { + const u32 kernel_preferred_wgs_multiple4 = device_param->kernel_preferred_wgs_multiple4; + + if ((kernel_preferred_wgs_multiple4 >= kernel_threads_min) && (kernel_preferred_wgs_multiple4 <= kernel_threads_max)) + { + kernel_threads_min = kernel_preferred_wgs_multiple4; + kernel_threads_max = kernel_preferred_wgs_multiple4; + } + } } } else { - if (device_param->kernel_wgs2) kernel_threads = MIN (kernel_threads, device_param->kernel_wgs2); + if (device_param->kernel_preferred_wgs_multiple2) + { + const u32 kernel_preferred_wgs_multiple2 = device_param->kernel_preferred_wgs_multiple2; + + if ((kernel_preferred_wgs_multiple2 >= kernel_threads_min) && (kernel_preferred_wgs_multiple2 <= kernel_threads_max)) + { + kernel_threads_min = kernel_preferred_wgs_multiple2; + kernel_threads_max = kernel_preferred_wgs_multiple2; + } + } } } - return kernel_threads; + return kernel_threads_max; } int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) @@ -4407,13 +4477,14 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) * kernel accel and loops tuning db adjustment */ - device_param->kernel_accel_min = 1; - device_param->kernel_accel_max = 1024; - - device_param->kernel_loops_min = 1; - device_param->kernel_loops_max = 1024; + device_param->kernel_accel_min = hashconfig->kernel_accel_min; + device_param->kernel_accel_max = hashconfig->kernel_accel_max; + device_param->kernel_loops_min = hashconfig->kernel_loops_min; + device_param->kernel_loops_max = hashconfig->kernel_loops_max; + device_param->kernel_threads_min = hashconfig->kernel_threads_min; + device_param->kernel_threads_max = hashconfig->kernel_threads_max; - tuning_db_entry_t *tuningdb_entry; + tuning_db_entry_t *tuningdb_entry = NULL; if (user_options->slow_candidates == true) { @@ -4424,56 +4495,83 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) tuningdb_entry = tuning_db_search (hashcat_ctx, device_param->device_name, device_param->device_type, user_options->attack_mode, hashconfig->hash_mode); } - if (tuningdb_entry != NULL) + // user commandline option override tuning db + // but both have to stay inside the boundaries of the module + + if (user_options->kernel_accel_chgd == true) { - u32 _kernel_accel = tuningdb_entry->kernel_accel; - u32 _kernel_loops = tuningdb_entry->kernel_loops; + const u32 _kernel_accel = user_options->kernel_accel; - if (_kernel_accel) + if ((_kernel_accel >= device_param->kernel_accel_min) && (_kernel_accel <= device_param->kernel_accel_max)) { device_param->kernel_accel_min = _kernel_accel; device_param->kernel_accel_max = _kernel_accel; } - - if (_kernel_loops) + } + else + { + if (tuningdb_entry != NULL) { - if (user_options->workload_profile == 1) - { - _kernel_loops = (_kernel_loops > 8) ? _kernel_loops / 8 : 1; - } - else if (user_options->workload_profile == 2) + const u32 _kernel_accel = tuningdb_entry->kernel_accel; + + if (_kernel_accel) { - _kernel_loops = (_kernel_loops > 4) ? _kernel_loops / 4 : 1; + if ((_kernel_accel >= device_param->kernel_accel_min) && (_kernel_accel <= device_param->kernel_accel_max)) + { + device_param->kernel_accel_min = _kernel_accel; + device_param->kernel_accel_max = _kernel_accel; + } } + } + } + + if (user_options->kernel_loops_chgd == true) + { + const u32 _kernel_loops = user_options->kernel_loops; + if ((_kernel_loops >= device_param->kernel_loops_min) && (_kernel_loops <= device_param->kernel_loops_max)) + { device_param->kernel_loops_min = _kernel_loops; device_param->kernel_loops_max = _kernel_loops; } } - - // commandline parameters overwrite tuningdb entries - - if (user_options->kernel_accel_chgd == true) + else { - device_param->kernel_accel_min = user_options->kernel_accel; - device_param->kernel_accel_max = user_options->kernel_accel; - } + if (tuningdb_entry != NULL) + { + u32 _kernel_loops = tuningdb_entry->kernel_loops; - if (user_options->kernel_loops_chgd == true) - { - device_param->kernel_loops_min = user_options->kernel_loops; - device_param->kernel_loops_max = user_options->kernel_loops; - } + if (_kernel_loops) + { + if (user_options->workload_profile == 1) + { + _kernel_loops = (_kernel_loops > 8) ? _kernel_loops / 8 : 1; + } + else if (user_options->workload_profile == 2) + { + _kernel_loops = (_kernel_loops > 4) ? _kernel_loops / 4 : 1; + } - // limit scrypt accel otherwise we hurt ourself when calculating the scrypt tmto + if ((_kernel_loops >= device_param->kernel_loops_min) && (_kernel_loops <= device_param->kernel_loops_max)) + { + device_param->kernel_loops_min = _kernel_loops; + device_param->kernel_loops_max = _kernel_loops; + } + } + } + } - #define SCRYPT_MAX_ACCEL 16 + // there's not thread column in tuning db, stick to commandline if defined - if ((hashconfig->hash_mode == 8900) || (hashconfig->hash_mode == 9300) || (hashconfig->hash_mode == 15700)) + if (user_options->kernel_threads_chgd == true) { - // 16 is actually a bit low, we may need to change this depending on user response + const u32 _kernel_threads = user_options->kernel_threads; - device_param->kernel_accel_max = MIN (device_param->kernel_accel_max, SCRYPT_MAX_ACCEL); + if ((_kernel_threads >= device_param->kernel_threads_min) && (_kernel_threads <= device_param->kernel_threads_max)) + { + device_param->kernel_threads_min = _kernel_threads; + device_param->kernel_threads_max = _kernel_threads; + } } if (user_options->slow_candidates == true) @@ -4664,8 +4762,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) } /** - * some algorithms need a fixed kernel-loops count - */ + * some algorithms need a fixed kernel-loops count -- not longer needed with new concept const u32 forced_kernel_loops = hashconfig->forced_kernel_loops; @@ -4674,6 +4771,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_loops_min = forced_kernel_loops; device_param->kernel_loops_max = forced_kernel_loops; } + */ device_param->kernel_loops_min_sav = device_param->kernel_loops_min; device_param->kernel_loops_max_sav = device_param->kernel_loops_max; diff --git a/src/selftest.c b/src/selftest.c index 85f56ffc6..8e064a72f 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -384,11 +384,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param salt_t *salt_buf = &hashes->st_salts_buf[salt_pos]; - const u32 forced_kernel_loops = hashconfig->forced_kernel_loops; - - //const u32 loop_step = (forced_kernel_loops) ? forced_kernel_loops : 1; - // test mode, not sure if this is safe - const u32 loop_step = (forced_kernel_loops) ? forced_kernel_loops : 1024; + const u32 loop_step = hashconfig->kernel_loops_max; const u32 iter = salt_buf->salt_iter;