From fcd20fbacb9b3723253e3790fcc56e4f88ee099a Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Tue, 29 Aug 2017 17:36:50 +0200 Subject: [PATCH] OpenCL Runtime: Fall back to 64 threads default (from 256) on AMD GPU to prevent creating too many workitems --- docs/changes.txt | 2 +- include/types.h | 2 +- src/interface.c | 9 +++++++-- 3 files changed, 9 insertions(+), 4 deletions(-) diff --git a/docs/changes.txt b/docs/changes.txt index 9056714a2..7dae65ace 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -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 diff --git a/include/types.h b/include/types.h index 14014220c..c5182ca17 100644 --- a/include/types.h +++ b/include/types.h @@ -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; diff --git a/src/interface.c b/src/interface.c index 7d02607ee..e877e54f3 100644 --- a/src/interface.c +++ b/src/interface.c @@ -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; } }