From 136165173066c69f1c335eff4173f878d2352658 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Tue, 13 Feb 2018 10:27:43 +0100 Subject: [PATCH] Add OPTS_TYPE_PREFERED_THREAD marker --- include/types.h | 1 + src/interface.c | 36 ++++++++++++++++++++++++++++++++---- 2 files changed, 33 insertions(+), 4 deletions(-) diff --git a/include/types.h b/include/types.h index b998e66a4..fc589a1a9 100644 --- a/include/types.h +++ b/include/types.h @@ -385,6 +385,7 @@ typedef enum opts_type OPTS_TYPE_AUX3 = (1ULL << 36), OPTS_TYPE_AUX4 = (1ULL << 37), OPTS_TYPE_BINARY_HASHFILE = (1ULL << 38), + OPTS_TYPE_PREFERED_THREAD = (1ULL << 39), // some algorithms (complicated ones with many branches) benefit from this } opts_type_t; diff --git a/src/interface.c b/src/interface.c index d4629024f..4c135f67a 100644 --- a/src/interface.c +++ b/src/interface.c @@ -22278,7 +22278,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) case 1800: hashconfig->hash_type = HASH_TYPE_SHA512; hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; - hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE; + hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE + | OPTS_TYPE_PREFERED_THREAD; hashconfig->kern_type = KERN_TYPE_SHA512CRYPT; hashconfig->dgst_size = DGST_SIZE_8_8; hashconfig->parse_func = sha512crypt_parse_hash; @@ -24731,7 +24732,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) case 13400: hashconfig->hash_type = HASH_TYPE_AES; hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; - hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE; + hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE + | OPTS_TYPE_PREFERED_THREAD; hashconfig->kern_type = KERN_TYPE_KEEPASS; hashconfig->dgst_size = DGST_SIZE_4_4; hashconfig->parse_func = keepass_parse_hash; @@ -25435,7 +25437,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) case 15900: hashconfig->hash_type = HASH_TYPE_DPAPIMK; hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; - hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE; + hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE + | OPTS_TYPE_PREFERED_THREAD; hashconfig->kern_type = KERN_TYPE_DPAPIMK_V2; hashconfig->dgst_size = DGST_SIZE_4_4; hashconfig->parse_func = dpapimk_parse_hash; @@ -25990,7 +25993,32 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p u32 kernel_threads = (u32) device_param->device_maxworkgroup_size; - // for CPU we use a special path + // complicated kernel tend to confuse OpenCL runtime suggestions for maximum thread size + // let's workaround that by sticking to their preferred thread size + + if (hashconfig->opts_type & OPTS_TYPE_PREFERED_THREAD) + { + if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_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_multiple2) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple2); + if (device_param->kernel_preferred_wgs_multiple3) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple3); + if (device_param->kernel_preferred_wgs_multiple4) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple4); + if (device_param->kernel_preferred_wgs_multiple_tm) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple_tm); + } + else + { + if (device_param->kernel_preferred_wgs_multiple1) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple1); + if (device_param->kernel_preferred_wgs_multiple2) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple2); + if (device_param->kernel_preferred_wgs_multiple3) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple3); + if (device_param->kernel_preferred_wgs_multiple12) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple12); + if (device_param->kernel_preferred_wgs_multiple23) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple23); + if (device_param->kernel_preferred_wgs_multiple_init2) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple_init2); + if (device_param->kernel_preferred_wgs_multiple_loop2) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple_loop2); + } + } + + // for CPU we do the same, because some allow up to 8192 thread which seem to be a bit excessive if (device_param->device_type & CL_DEVICE_TYPE_CPU) {