mirror of
https://github.com/hashcat/hashcat
synced 2025-03-19 06:14:22 +01:00
Improve performance of bitsliced algorithms on ROCm
This commit is contained in:
parent
1a022217d0
commit
955bfeaa14
119
src/backend.c
119
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)
|
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_min = device_param->kernel_threads_min;
|
||||||
u32 kernel_threads_max = device_param->kernel_threads_max;
|
u32 kernel_threads_max = device_param->kernel_threads_max;
|
||||||
|
|
||||||
|
// the changes we do here are just optimizations, since the module always has priority.
|
||||||
|
|
||||||
|
const u32 device_maxworkgroup_size = (const u32) device_param->device_maxworkgroup_size;
|
||||||
|
|
||||||
|
kernel_threads_max = MIN (kernel_threads_max, device_maxworkgroup_size);
|
||||||
|
|
||||||
// for CPU we just do 1 ...
|
// for CPU we just do 1 ...
|
||||||
|
|
||||||
if (device_param->opencl_device_type & CL_DEVICE_TYPE_CPU)
|
if (device_param->opencl_device_type & CL_DEVICE_TYPE_CPU)
|
||||||
{
|
{
|
||||||
if ((1 >= kernel_threads_min) && (1 <= kernel_threads_max))
|
const u32 cpu_prefered_thread_count = 1;
|
||||||
{
|
|
||||||
kernel_threads_min = 1;
|
kernel_threads_max = MIN (kernel_threads_max, cpu_prefered_thread_count);
|
||||||
kernel_threads_max = 1;
|
|
||||||
}
|
|
||||||
}
|
}
|
||||||
|
|
||||||
// this is an upper limit, a good start, since our strategy is to reduce thread counts only
|
// 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.
|
||||||
|
|
||||||
const u32 device_maxworkgroup_size = (u32) device_param->device_maxworkgroup_size;
|
const u32 kernel_threads = MAX (kernel_threads_min, kernel_threads_max);
|
||||||
|
|
||||||
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
|
|
||||||
|
|
||||||
/*
|
|
||||||
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)
|
|
||||||
{
|
|
||||||
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;
|
|
||||||
|
|
||||||
if ((kernel_preferred_wgs_multiple2 >= kernel_threads_min) && (kernel_preferred_wgs_multiple2 <= kernel_threads_max))
|
|
||||||
{
|
|
||||||
kernel_threads = 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_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;
|
|
||||||
|
|
||||||
if ((kernel_preferred_wgs_multiple2 >= kernel_threads_min) && (kernel_preferred_wgs_multiple2 <= kernel_threads_max))
|
|
||||||
{
|
|
||||||
kernel_threads = kernel_preferred_wgs_multiple2;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
|
||||||
}
|
|
||||||
*/
|
|
||||||
|
|
||||||
return kernel_threads;
|
return kernel_threads;
|
||||||
}
|
}
|
||||||
|
@ -54,6 +54,13 @@ typedef struct blake2
|
|||||||
|
|
||||||
static const char *SIGNATURE_BLAKE2B = "$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)
|
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);
|
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);
|
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)
|
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;
|
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_accel_min = MODULE_DEFAULT;
|
||||||
module_ctx->module_kernel_loops_max = MODULE_DEFAULT;
|
module_ctx->module_kernel_loops_max = MODULE_DEFAULT;
|
||||||
module_ctx->module_kernel_loops_min = 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_kernel_threads_min = MODULE_DEFAULT;
|
||||||
module_ctx->module_kern_type = module_kern_type;
|
module_ctx->module_kern_type = module_kern_type;
|
||||||
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
||||||
|
@ -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_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; }
|
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 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;
|
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_accel_min = MODULE_DEFAULT;
|
||||||
module_ctx->module_kernel_loops_max = module_kernel_loops_max;
|
module_ctx->module_kernel_loops_max = module_kernel_loops_max;
|
||||||
module_ctx->module_kernel_loops_min = module_kernel_loops_min;
|
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_kernel_threads_min = MODULE_DEFAULT;
|
||||||
module_ctx->module_kern_type = module_kern_type;
|
module_ctx->module_kern_type = module_kern_type;
|
||||||
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
||||||
|
@ -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_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; }
|
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 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;
|
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_accel_min = MODULE_DEFAULT;
|
||||||
module_ctx->module_kernel_loops_max = module_kernel_loops_max;
|
module_ctx->module_kernel_loops_max = module_kernel_loops_max;
|
||||||
module_ctx->module_kernel_loops_min = module_kernel_loops_min;
|
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_kernel_threads_min = MODULE_DEFAULT;
|
||||||
module_ctx->module_kern_type = module_kern_type;
|
module_ctx->module_kern_type = module_kern_type;
|
||||||
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
||||||
|
Loading…
x
Reference in New Issue
Block a user