From b2684553f094378f124bcd7eae9ba8eb517d2e66 Mon Sep 17 00:00:00 2001 From: jsteube Date: Tue, 28 Feb 2017 12:12:06 +0100 Subject: [PATCH] In scrypt we really need to get all memory we can get, thus we workaround the 1/4 limit by allocating 4 blocks but this can eventually be larger than the total maximum available memory --- src/opencl.c | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/src/opencl.c b/src/opencl.c index 2d936f743..4062e8e65 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -2381,6 +2381,7 @@ void opencl_ctx_destroy (hashcat_ctx_t *hashcat_ctx) int opencl_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) { + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; opencl_ctx_t *opencl_ctx = hashcat_ctx->opencl_ctx; user_options_t *user_options = hashcat_ctx->user_options; @@ -2700,6 +2701,14 @@ int opencl_ctx_devices_init (hashcat_ctx_t *hashcat_ctx, const int comptime) // testwise disabling that //device_param->device_maxmem_alloc = MIN (device_maxmem_alloc, 0x7fffffff); + if ((hashconfig->hash_mode == 8900) || (hashconfig->hash_mode == 9300)) + { + // in scrypt we really need to get all memory we can get, thus we workaround the 1/4 limit + // by allocating 4 blocks but this can eventually be larger than the total maximum available memory + + device_param->device_maxmem_alloc = MIN (device_maxmem_alloc, device_param->device_global_mem / 4); + } + // max_work_group_size size_t device_maxworkgroup_size;