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

OpenCL Kernels: Thread-count is set to hardware native count except if -w 4 is used then OpenCL maximum is used

Browse Source
This commit is contained in:
Jens Steube 2017-10-22 14:16:32 +02:00
parent ecbf053639
commit 77f3eb2864
2 changed files with 7 additions and 16 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
docs/changes.txt
View File

@ -60,6 +60,7 @@
- OpenCL Kernels: Replace bitwise swaps with rotate() versions for AMD - OpenCL Kernels: Replace bitwise swaps with rotate() versions for AMD
- OpenCL Kernels: Rewritten Keccak kernel to run fully on registers and partially reversed last round - OpenCL Kernels: Rewritten Keccak kernel to run fully on registers and partially reversed last round
- OpenCL Kernels: Rewritten SIP kernel from scratch - OpenCL Kernels: Rewritten SIP kernel from scratch
- OpenCL Kernels: Thread-count is set to hardware native count except if -w 4 is used then OpenCL maximum is used
- OpenCL Kernels: Updated default scrypt TMTO to be ideal for latest NVidia and AMD top models - OpenCL Kernels: Updated default scrypt TMTO to be ideal for latest NVidia and AMD top models
- OpenCL Kernels: Vectorized tons of slow kernels to improve CPU cracking speed - OpenCL Kernels: Vectorized tons of slow kernels to improve CPU cracking speed
- OpenCL Runtime: Improved detection for AMD and NV devices on macOS - OpenCL Runtime: Improved detection for AMD and NV devices on macOS

20
src/interface.c
View File

@ -25076,7 +25076,6 @@ 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_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_param_t *device_param)
{ {
const hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
const user_options_t *user_options = hashcat_ctx->user_options; const user_options_t *user_options = hashcat_ctx->user_options;
// a kernel can force a fixed value // a kernel can force a fixed value
@ -25096,23 +25095,9 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p
else if (device_param->device_type & CL_DEVICE_TYPE_GPU) else if (device_param->device_type & CL_DEVICE_TYPE_GPU)
{ {
if (device_param->device_vendor_id == VENDOR_ID_NV) if (device_param->device_vendor_id == VENDOR_ID_NV)
{
if (hashconfig->attack_exec == ATTACK_EXEC_OUTSIDE_KERNEL)
{
if (user_options->workload_profile < 4)
{ {
kernel_threads = KERNEL_THREADS_MAX_GPU_NV; kernel_threads = KERNEL_THREADS_MAX_GPU_NV;
} }
else
{
kernel_threads = device_param->device_maxworkgroup_size;
}
}
else
{
kernel_threads = device_param->device_maxworkgroup_size;
}
}
else if (device_param->device_vendor_id == VENDOR_ID_AMD) else if (device_param->device_vendor_id == VENDOR_ID_AMD)
{ {
kernel_threads = KERNEL_THREADS_MAX_GPU_AMD; kernel_threads = KERNEL_THREADS_MAX_GPU_AMD;
@ -25127,6 +25112,11 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p
kernel_threads = KERNEL_THREADS_MAX_OTHER; kernel_threads = KERNEL_THREADS_MAX_OTHER;
} }
if (user_options->workload_profile == 4)
{
kernel_threads = device_param->device_maxworkgroup_size;
}
// and (2) an opencl device can force an lower value (limited resources on device) // and (2) an opencl device can force an lower value (limited resources on device)
kernel_threads = MIN (kernel_threads, device_param->device_maxworkgroup_size); kernel_threads = MIN (kernel_threads, device_param->device_maxworkgroup_size);