Update for scrypt patch

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);