diff --git a/src/opencl.c b/src/opencl.c index 4062e8e65..43bfbc92b 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -2381,7 +2381,6 @@ 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; @@ -2701,14 +2700,6 @@ 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; @@ -3700,6 +3691,47 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) device_param->kernel_accel_min = 1; device_param->kernel_accel_max = 8; + const u32 kernel_power_max = device_param->hardware_power * device_param->kernel_accel_max; + + // size_pws + + const size_t size_pws = kernel_power_max * sizeof (pw_t); + + // size_tmps + + const size_t size_tmps = kernel_power_max * hashconfig->tmp_size; + + // size_hooks + + const size_t size_hooks = kernel_power_max * hashconfig->hook_size; + + const u64 scrypt_extra_space + = bitmap_ctx->bitmap_size + + bitmap_ctx->bitmap_size + + bitmap_ctx->bitmap_size + + bitmap_ctx->bitmap_size + + bitmap_ctx->bitmap_size + + bitmap_ctx->bitmap_size + + bitmap_ctx->bitmap_size + + bitmap_ctx->bitmap_size + + size_bfs + + size_combs + + size_digests + + size_esalts + + size_hooks + + size_markov_css + + size_plains + + size_pws + + size_pws // not a bug + + size_results + + size_root_css + + size_rules + + size_rules_c + + size_salts + + size_shown + + size_tm + + size_tmps; + u32 tmto; for (tmto = tmto_start; tmto < tmto_stop; tmto++) @@ -3717,7 +3749,7 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) continue; } - if (size_scrypt > device_param->device_global_mem) + if ((size_scrypt + scrypt_extra_space) > device_param->device_global_mem) { if (user_options->quiet == false) event_log_warning (hashcat_ctx, "Not enough total device memory allocatable to use --scrypt-tmto %u, increasing...", tmto);