arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-22 16:18:09 +00:00
Code Issues Releases Wiki Activity

Get rid of hashconfig_get_kernel_threads() and hashconfig_get_kernel_loops()

Browse Source
This commit is contained in:
jsteube 2018-12-17 13:03:32 +01:00
parent af22a9d1d9
commit 074947c6db
7 changed files with 91 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
include/interface.h
View File

@ -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);

1
include/modules.h
View File

@ -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);

2
include/types.h
View File

@ -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 *);

1
modules/m01000.c
View File

@ -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;

15
src/interface.c
View File

@ -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;

71
src/opencl.c
View File

@ -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];

6
src/selftest.c
View File

@ -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;