it's possible to crack scrypt on GPU even with higher scrypt setting

the trick is to intentionally under-utilize the GPU warps
120H/s on my 4*1080: https://pastebin.com/z124G2cN

hashcat.hctune

@@ -473,6 +473,6 @@ DEVICE_TYPE_CPU                                 *       8900    1       1
 DEVICE_TYPE_CPU                                 *       9300    1       1       1
 DEVICE_TYPE_CPU                                 *       15700   1       1       1
 
-DEVICE_TYPE_GPU                                 *       8900    1       8       1
+DEVICE_TYPE_GPU                                 *       8900    1       16      1
-DEVICE_TYPE_GPU                                 *       9300    1       8       1
+DEVICE_TYPE_GPU                                 *       9300    1       16      1
 DEVICE_TYPE_GPU                                 *       15700   1       1       1

src/interface.c

@@ -23515,9 +23515,23 @@ 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 (hashconfig->hash_mode ==  8900) kernel_threads = 64; // Scrypt
+  if ((hashconfig->hash_mode == 8900) || (hashconfig->hash_mode == 9300) || (hashconfig->hash_mode == 15700))
-  if (hashconfig->hash_mode ==  9300) kernel_threads = 64; // Scrypt
+  {
-  if (hashconfig->hash_mode == 15700) kernel_threads = 64; // Scrypt
+    const hashes_t *hashes = hashcat_ctx->hashes;
+
+    const u32 scrypt_r = hashes->salts_buf[0].scrypt_r;
+    const u32 scrypt_p = hashes->salts_buf[0].scrypt_p;
+    const u32 scrypt_l = scrypt_r * scrypt_p;
+
+    if (scrypt_l)
+    {
+      kernel_threads = 256 / scrypt_l;
+    }
+    else
+    {
+      kernel_threads = 256;
+    }
+  }
 
   if (device_param->device_type & CL_DEVICE_TYPE_CPU)
   {

src/opencl.c

@@ -3685,11 +3685,12 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
       hashconfig->tmp_size = scrypt_tmp_size;
 
       u32 tmto_start = 0;
-      u32 tmto_stop  = 10;
+      u32 tmto_stop  = 5;
 
       if (user_options->scrypt_tmto)
       {
         tmto_start = user_options->scrypt_tmto;
+        tmto_stop  = user_options->scrypt_tmto;
       }
       else
       {
@@ -3718,17 +3719,6 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
             tmto_start = 4;
           }
         }
-        else if (hashconfig->hash_mode == 15700)
-        {
-          if (device_param->device_vendor_id == VENDOR_ID_AMD)
-          {
-            tmto_start = 5;
-          }
-          else if (device_param->device_vendor_id == VENDOR_ID_NV)
-          {
-            tmto_start = 6;
-          }
-        }
       }
 
       const u32 kernel_power_max = device_param->hardware_power * device_param->kernel_accel_max;
@@ -3772,9 +3762,11 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
         + size_tm
         + size_tmps;
 
+      bool not_enough_memory = true;
+
       u32 tmto;
 
-      for (tmto = tmto_start; tmto < tmto_stop; tmto++)
+      for (tmto = tmto_start; tmto <= tmto_stop; tmto++)
       {
         size_scrypt = (128 * scrypt_r) * scrypt_N;
 
@@ -3801,10 +3793,12 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx)
           scrypt_tmto_final = tmto;
         }
 
+        not_enough_memory = false;
+
         break;
       }
 
-      if (tmto == tmto_stop)
+      if (not_enough_memory == true)
       {
         event_log_error (hashcat_ctx, "Cannot allocate enough device memory.");