diff --git a/include/interface.h b/include/interface.h index 00c20fe5e..c17c41fc8 100644 --- a/include/interface.h +++ b/include/interface.h @@ -1714,8 +1714,8 @@ bool initialize_keyboard_layout_mapping (hashcat_ctx_t *hashcat_ctx, const char int hashconfig_init (hashcat_ctx_t *hashcat_ctx); void hashconfig_destroy (hashcat_ctx_t *hashcat_ctx); //u32 hashconfig_forced_kernel_threads (hashcat_ctx_t *hashcat_ctx); -u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param_t *device_param); -u32 hashconfig_get_kernel_loops (hashcat_ctx_t *hashcat_ctx); +//u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param_t *device_param); +//u32 hashconfig_get_kernel_loops (hashcat_ctx_t *hashcat_ctx); int hashconfig_general_defaults (hashcat_ctx_t *hashcat_ctx); void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, void *esalt, void *hook_salt); @@ -1731,6 +1731,7 @@ u32 default_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hash 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_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_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); 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); u32 default_hash_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 13b5aa062..27ec06339 100644 --- a/include/modules.h +++ b/include/modules.h @@ -14,6 +14,7 @@ u32 module_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hash 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_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_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); 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); u32 module_hash_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 b27ee4f1a..5318d3bc7 100644 --- a/include/types.h +++ b/include/types.h @@ -939,6 +939,7 @@ struct hashconfig void *benchmark_hook_salt; u32 forced_kernel_threads; + u32 forced_kernel_loops; }; typedef struct hashconfig hashconfig_t; @@ -2222,6 +2223,7 @@ typedef struct module_ctx 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_threads) (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 *); 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 *); u32 (*module_hash_type) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *); diff --git a/modules/m01000.c b/modules/m01000.c index bdca88584..e0ad7436f 100644 --- a/modules/m01000.c +++ b/modules/m01000.c @@ -124,6 +124,7 @@ void module_register (module_ctx_t *module_ctx) module_ctx->module_dgst_size = module_dgst_size; module_ctx->module_esalt_size = NULL; module_ctx->module_forced_kernel_threads = NULL; + module_ctx->module_forced_kernel_loops = NULL; module_ctx->module_hash_decode = module_hash_decode; module_ctx->module_hash_encode = module_hash_encode; module_ctx->module_hash_mode = NULL; diff --git a/src/interface.c b/src/interface.c index 936a74af8..78512ebe8 100644 --- a/src/interface.c +++ b/src/interface.c @@ -23612,6 +23612,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->dgst_size = default_dgst_size (hashconfig, user_options, user_options_extra); hashconfig->esalt_size = default_esalt_size (hashconfig, user_options, user_options_extra); hashconfig->forced_kernel_threads = default_forced_kernel_threads (hashconfig, user_options, user_options_extra); + hashconfig->forced_kernel_loops = default_forced_kernel_loops (hashconfig, user_options, user_options_extra); hashconfig->hash_name = default_hash_name (hashconfig, user_options, user_options_extra); hashconfig->hash_mode = default_hash_mode (hashconfig, user_options, user_options_extra); hashconfig->hash_type = default_hash_type (hashconfig, user_options, user_options_extra); @@ -23638,7 +23639,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) module_register (module_ctx); if (module_ctx->module_attack_exec) hashconfig->attack_exec = module_ctx->module_attack_exec (hashconfig, user_options, user_options_extra); - if (module_ctx->module_benchmark_esalt) hashconfig->benchmark_esalt = module_ctx->module_benchmark_salt (hashconfig, user_options, user_options_extra); + if (module_ctx->module_benchmark_esalt) hashconfig->benchmark_esalt = module_ctx->module_benchmark_esalt (hashconfig, user_options, user_options_extra); if (module_ctx->module_benchmark_hook_salt) hashconfig->benchmark_salt = module_ctx->module_benchmark_hook_salt (hashconfig, user_options, user_options_extra); if (module_ctx->module_benchmark_mask) hashconfig->benchmark_mask = module_ctx->module_benchmark_mask (hashconfig, user_options, user_options_extra); if (module_ctx->module_benchmark_salt) hashconfig->benchmark_salt = module_ctx->module_benchmark_salt (hashconfig, user_options, user_options_extra); @@ -23649,6 +23650,7 @@ 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_esalt_size) hashconfig->esalt_size = module_ctx->module_esalt_size (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_kernel_loops) hashconfig->forced_kernel_loops = module_ctx->module_forced_kernel_loops (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); 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_type) hashconfig->hash_type = module_ctx->module_hash_type (hashconfig, user_options, user_options_extra); @@ -29319,6 +29321,7 @@ u32 hashconfig_forced_kernel_threads (hashcat_ctx_t *hashcat_ctx) } */ +/* u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param_t *device_param) { const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; @@ -29379,7 +29382,9 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p return kernel_threads; } +*/ +/* u32 hashconfig_get_kernel_loops (hashcat_ctx_t *hashcat_ctx) { const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; @@ -29435,6 +29440,7 @@ u32 hashconfig_get_kernel_loops (hashcat_ctx_t *hashcat_ctx) return kernel_loops_fixed; } +*/ int hashconfig_general_defaults (hashcat_ctx_t *hashcat_ctx) { @@ -30890,6 +30896,13 @@ u32 default_forced_kernel_threads (MAYBE_UNUSED const hashconfig_t *hashconfig, return forced_kernel_threads; } +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) +{ + const u32 forced_kernel_loops = user_options->kernel_loops; + + return forced_kernel_loops; +} + u32 default_salt_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) { const u32 salt_type = SALT_TYPE_NONE; diff --git a/src/opencl.c b/src/opencl.c index 42d1d2ff2..921f064e9 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -4250,6 +4250,67 @@ static int get_kernel_local_mem_size (hashcat_ctx_t *hashcat_ctx, hc_device_para return 0; } +static u32 get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param_t *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; + + if (forced_kernel_threads) return forced_kernel_threads; + + // for CPU we just do 1 ... + + if (device_param->device_type & CL_DEVICE_TYPE_CPU) return 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; + + // 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 + + if (hashconfig->opts_type & OPTS_TYPE_PREFERED_THREAD) + { + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + 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); + } + else + { + if (device_param->kernel_preferred_wgs_multiple4) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple4); + } + } + else + { + if (device_param->kernel_preferred_wgs_multiple2) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple2); + } + } + else + { + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) + { + if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) + { + if (device_param->kernel_wgs1) kernel_threads = MIN (kernel_threads, device_param->kernel_wgs1); + } + else + { + if (device_param->kernel_wgs4) kernel_threads = MIN (kernel_threads, device_param->kernel_wgs4); + } + } + else + { + if (device_param->kernel_wgs2) kernel_threads = MIN (kernel_threads, device_param->kernel_wgs2); + } + } + + return kernel_threads; +} + int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) { bitmap_ctx_t *bitmap_ctx = hashcat_ctx->bitmap_ctx; @@ -4656,12 +4717,12 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) * some algorithms need a fixed kernel-loops count */ - const u32 kernel_loops_fixed = hashconfig_get_kernel_loops (hashcat_ctx); + const u32 forced_kernel_loops = hashconfig->forced_kernel_loops; - if (kernel_loops_fixed != 0) + if (forced_kernel_loops != 0) { - device_param->kernel_loops_min = kernel_loops_fixed; - device_param->kernel_loops_max = kernel_loops_fixed; + 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; @@ -6482,7 +6543,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) * now everything that depends on threads and accel, basically dynamic workload */ - u32 kernel_threads = hashconfig_get_kernel_threads (hashcat_ctx, device_param); + u32 kernel_threads = get_kernel_threads (hashcat_ctx, device_param); // this is required because inside the kernels there is this: // __local pw_t s_pws[64]; diff --git a/src/selftest.c b/src/selftest.c index 99a946991..0d30657f5 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -383,9 +383,11 @@ 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 kernel_loops_fixed = hashconfig_get_kernel_loops (hashcat_ctx); + const u32 forced_kernel_loops = hashconfig->forced_kernel_loops; - const u32 loop_step = (kernel_loops_fixed) ? kernel_loops_fixed : 1; + //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 iter = salt_buf->salt_iter;