From 955bfeaa14edaa347cb179506329f4d32f2edfe5 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Wed, 19 Jun 2019 16:35:52 +0200 Subject: [PATCH] Improve performance of bitsliced algorithms on ROCm --- src/backend.c | 115 ++++--------------------------------- src/modules/module_00600.c | 10 +++- src/modules/module_03000.c | 9 ++- src/modules/module_14000.c | 9 ++- 4 files changed, 35 insertions(+), 108 deletions(-) diff --git a/src/backend.c b/src/backend.c index 694bd8061..aad024b58 100644 --- a/src/backend.c +++ b/src/backend.c @@ -6949,123 +6949,30 @@ static int get_opencl_kernel_local_mem_size (hashcat_ctx_t *hashcat_ctx, hc_devi static u32 get_kernel_threads (const hc_device_param_t *device_param) { - // a module can force a fixed value + // this is an upper limit, a good start, since our strategy is to reduce thread counts only. u32 kernel_threads_min = device_param->kernel_threads_min; u32 kernel_threads_max = device_param->kernel_threads_max; - // for CPU we just do 1 ... + // the changes we do here are just optimizations, since the module always has priority. - if (device_param->opencl_device_type & CL_DEVICE_TYPE_CPU) - { - if ((1 >= kernel_threads_min) && (1 <= kernel_threads_max)) - { - 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 + const u32 device_maxworkgroup_size = (const u32) device_param->device_maxworkgroup_size; - const u32 device_maxworkgroup_size = (u32) device_param->device_maxworkgroup_size; + kernel_threads_max = MIN (kernel_threads_max, device_maxworkgroup_size); - if (device_maxworkgroup_size < kernel_threads_max) - { - kernel_threads_max = device_maxworkgroup_size; - } - - u32 kernel_threads = kernel_threads_max; - - // 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 - // this section was replaced by autotune + // for CPU we just do 1 ... - /* - if (hashconfig->opts_type & OPTS_TYPE_PREFERED_THREAD) + if (device_param->opencl_device_type & CL_DEVICE_TYPE_CPU) { - if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) - { - 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 = kernel_preferred_wgs_multiple1; - } - } - } - else - { - 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 = kernel_preferred_wgs_multiple4; - } - } - } - } - else - { - if (device_param->kernel_preferred_wgs_multiple2) - { - const u32 kernel_preferred_wgs_multiple2 = device_param->kernel_preferred_wgs_multiple2; + const u32 cpu_prefered_thread_count = 1; - if ((kernel_preferred_wgs_multiple2 >= kernel_threads_min) && (kernel_preferred_wgs_multiple2 <= kernel_threads_max)) - { - kernel_threads = kernel_preferred_wgs_multiple2; - } - } - } + kernel_threads_max = MIN (kernel_threads_max, cpu_prefered_thread_count); } - else - { - if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) - { - if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL) - { - 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 = kernel_preferred_wgs_multiple1; - } - } - } - else - { - 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 = kernel_preferred_wgs_multiple4; - } - } - } - } - else - { - if (device_param->kernel_preferred_wgs_multiple2) - { - const u32 kernel_preferred_wgs_multiple2 = device_param->kernel_preferred_wgs_multiple2; + // this is intenionally! at this point, kernel_threads_min can be higher than kernel_threads_max. + // in this case we actually want kernel_threads_min selected. - if ((kernel_preferred_wgs_multiple2 >= kernel_threads_min) && (kernel_preferred_wgs_multiple2 <= kernel_threads_max)) - { - kernel_threads = kernel_preferred_wgs_multiple2; - } - } - } - } - */ + const u32 kernel_threads = MAX (kernel_threads_min, kernel_threads_max); return kernel_threads; } diff --git a/src/modules/module_00600.c b/src/modules/module_00600.c index 096fb3e85..6fd67c51f 100644 --- a/src/modules/module_00600.c +++ b/src/modules/module_00600.c @@ -54,6 +54,13 @@ typedef struct blake2 static const char *SIGNATURE_BLAKE2B = "$BLAKE2$"; +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) +{ + const u32 kernel_threads_max = 64; // performance only optimization + + return kernel_threads_max; +} + 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) { const u64 esalt_size = (const u64) sizeof (blake2_t); @@ -118,7 +125,6 @@ int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSE return (PARSER_OK); } - int module_hash_encode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const void *digest_buf, MAYBE_UNUSED const salt_t *salt, MAYBE_UNUSED const void *esalt_buf, MAYBE_UNUSED const void *hook_salt_buf, MAYBE_UNUSED const hashinfo_t *hash_info, char *line_buf, MAYBE_UNUSED const int line_size) { const u64 *digest = (const u64 *) digest_buf; @@ -192,7 +198,7 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_kernel_accel_min = MODULE_DEFAULT; module_ctx->module_kernel_loops_max = MODULE_DEFAULT; module_ctx->module_kernel_loops_min = MODULE_DEFAULT; - module_ctx->module_kernel_threads_max = MODULE_DEFAULT; + module_ctx->module_kernel_threads_max = module_kernel_threads_max; module_ctx->module_kernel_threads_min = MODULE_DEFAULT; module_ctx->module_kern_type = module_kern_type; module_ctx->module_kern_type_dynamic = MODULE_DEFAULT; diff --git a/src/modules/module_03000.c b/src/modules/module_03000.c index dcd62e787..e5c819941 100644 --- a/src/modules/module_03000.c +++ b/src/modules/module_03000.c @@ -46,6 +46,13 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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) +{ + const u32 kernel_threads_max = 64; // performance only optimization + + return kernel_threads_max; +} + 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 kernel_loops_max = KERNEL_LOOPS_MAX; @@ -209,7 +216,7 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_kernel_accel_min = MODULE_DEFAULT; module_ctx->module_kernel_loops_max = module_kernel_loops_max; module_ctx->module_kernel_loops_min = module_kernel_loops_min; - module_ctx->module_kernel_threads_max = MODULE_DEFAULT; + module_ctx->module_kernel_threads_max = module_kernel_threads_max; module_ctx->module_kernel_threads_min = MODULE_DEFAULT; module_ctx->module_kern_type = module_kern_type; module_ctx->module_kern_type_dynamic = MODULE_DEFAULT; diff --git a/src/modules/module_14000.c b/src/modules/module_14000.c index 9703ac82d..bb55b36ec 100644 --- a/src/modules/module_14000.c +++ b/src/modules/module_14000.c @@ -43,6 +43,13 @@ u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, const char *module_st_hash (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 ST_HASH; } const char *module_st_pass (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 ST_PASS; } +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) +{ + const u32 kernel_threads_max = 64; // performance only optimization + + return kernel_threads_max; +} + 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 kernel_loops_max = KERNEL_LOOPS_MAX; @@ -240,7 +247,7 @@ void module_init (module_ctx_t *module_ctx) module_ctx->module_kernel_accel_min = MODULE_DEFAULT; module_ctx->module_kernel_loops_max = module_kernel_loops_max; module_ctx->module_kernel_loops_min = module_kernel_loops_min; - module_ctx->module_kernel_threads_max = MODULE_DEFAULT; + module_ctx->module_kernel_threads_max = module_kernel_threads_max; module_ctx->module_kernel_threads_min = MODULE_DEFAULT; module_ctx->module_kern_type = module_kern_type; module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;