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

Re-enable CPU optimizations and some CPU case in thread management

Browse Source
This commit is contained in:
Jens Steube 2018-02-16 18:56:21 +01:00
parent fe321105fe
commit aa82d8d34d
3 changed files with 18 additions and 42 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
OpenCL/inc_vendor.cl
View File

@ -82,7 +82,7 @@
*/ */
#if defined IS_CPU #if defined IS_CPU
#define DECLSPEC static #define DECLSPEC inline
#elif defined IS_GPU #elif defined IS_GPU
#if defined IS_AMD #if defined IS_AMD
#define DECLSPEC inline #define DECLSPEC inline

56
src/interface.c
View File

@ -25993,56 +25993,33 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p
if (forced_kernel_threads) return forced_kernel_threads; if (forced_kernel_threads) return forced_kernel_threads;
// it can also depends on the opencl device type // 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; u32 kernel_threads = (u32) device_param->device_maxworkgroup_size;
// complicated kernel tend to confuse OpenCL runtime suggestions for maximum thread size // complicated kernel tend to confuse OpenCL runtime suggestions for maximum thread size
// let's workaround that by sticking to their preferred 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->opts_type & OPTS_TYPE_PREFERED_THREAD)
{ {
if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) 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 (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
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_multiple1) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple1);
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_multiple4) kernel_threads = MIN (kernel_threads, device_param->kernel_preferred_wgs_multiple4);
}
} }
else 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_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)
{
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);
} }
} }
else else
@ -26064,9 +26041,6 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p
} }
} }
// we'll return a number power of two, makes future processing much more easy
// kernel_threads = power_of_two_floor_32 (kernel_threads);
return kernel_threads; return kernel_threads;
} }

2
src/opencl.c
View File

@ -4546,6 +4546,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -D VENDOR_ID=%u -D CUDA_ARCH=%u -D AMD_ROCM=%u -D VECT_SIZE=%u -D DEVICE_TYPE=%u -D DGST_R0=%u -D DGST_R1=%u -D DGST_R2=%u -D DGST_R3=%u -D DGST_ELEM=%u -D KERN_TYPE=%u -D _unroll -w", build_opts, device_param->platform_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->is_rocm, device_param->vector_width, (u32) device_param->device_type, hashconfig->dgst_pos0, hashconfig->dgst_pos1, hashconfig->dgst_pos2, hashconfig->dgst_pos3, hashconfig->dgst_size / 4, hashconfig->kern_type); snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -D VENDOR_ID=%u -D CUDA_ARCH=%u -D AMD_ROCM=%u -D VECT_SIZE=%u -D DEVICE_TYPE=%u -D DGST_R0=%u -D DGST_R1=%u -D DGST_R2=%u -D DGST_R3=%u -D DGST_ELEM=%u -D KERN_TYPE=%u -D _unroll -w", build_opts, device_param->platform_vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->is_rocm, device_param->vector_width, (u32) device_param->device_type, hashconfig->dgst_pos0, hashconfig->dgst_pos1, hashconfig->dgst_pos2, hashconfig->dgst_pos3, hashconfig->dgst_size / 4, hashconfig->kern_type);
#endif #endif
/*
if (device_param->device_type & CL_DEVICE_TYPE_CPU) if (device_param->device_type & CL_DEVICE_TYPE_CPU)
{ {
if (device_param->platform_vendor_id == VENDOR_ID_INTEL_SDK) if (device_param->platform_vendor_id == VENDOR_ID_INTEL_SDK)
@ -4553,6 +4554,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
strncat (build_opts_new, " -cl-opt-disable", 16); strncat (build_opts_new, " -cl-opt-disable", 16);
} }
} }
*/
strncpy (build_opts, build_opts_new, sizeof (build_opts) - 1); strncpy (build_opts, build_opts_new, sizeof (build_opts) - 1);