github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-12-23 14:13:43 +01:00
Code Releases Activity

OpenCL Runtime: Fall back to 64 threads default (from 256) on AMD GPU to prevent creating too many workitems

Browse Source
This commit is contained in:
Jens Steube 2017-08-29 17:36:50 +02:00
parent a378abee66
commit fcd20fbacb
3 changed files with 9 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
docs/changes.txt
View File

@ -68,7 +68,7 @@
- OpenCL Kernels: Declared all include functions as static to reduce binary kernel cache size
- OpenCL Kernels: On AMD GPU, optimized kernels for use with AMD ROCm driver
- OpenCL Kernels: Removed some include functions that are no longer needed to reduce compile time
- OpenCL Runtime: Fall back to 64 threads default (from 256) because modern GPU create too many workitems otherwise
- OpenCL Runtime: Fall back to 64 threads default (from 256) on AMD GPU to prevent creating too many workitems
- OpenCL Runtime: Forcing OpenCL 1.2 no longer needed. Option removed from build options
- Restore: Fixed the version number used in the restore file header
- Time: added new type for time measurements hc_time_t and related functions to force the use of 64 bit times

2
include/types.h
View File

@ -155,7 +155,7 @@ typedef enum amplifier_count
KERNEL_BFS = 1024,
KERNEL_COMBS = 1024,
KERNEL_RULES = 256,
KERNEL_THREADS_MAX = 64,
KERNEL_THREADS_MAX = 256,
KERNEL_THREADS_MAX_CPU = 1
} amplifier_count_t;

9
src/interface.c
View File

@ -24920,6 +24920,11 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p
u32 kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size);
if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
kernel_threads = 64;
}
if ((hashconfig->hash_mode == 8900) || (hashconfig->hash_mode == 9300) || (hashconfig->hash_mode == 15700))
{
const hashes_t *hashes = hashcat_ctx->hashes;
@ -24930,11 +24935,11 @@ u32 hashconfig_get_kernel_threads (hashcat_ctx_t *hashcat_ctx, const hc_device_p
if (scrypt_l)
{
kernel_threads = KERNEL_THREADS_MAX / scrypt_l;
kernel_threads = 64 / scrypt_l;
}
else
{
kernel_threads = KERNEL_THREADS_MAX;
kernel_threads = 64;
}
}