Update to better scrypt tmto defaults

[hashcat.git] / src / oclHashcat.c

diff --git a/src/oclHashcat.c b/src/oclHashcat.c
index 97b09a9..55525f3 100644 (file)
--- a/src/oclHashcat.c
+++ b/src/oclHashcat.c
@@ -84,6 +84,8 @@ double TARGET_MS_PROFILE[3]     = { 8, 16, 96 };
 #define KERNEL_RULES            1024
 #define KERNEL_COMBS            1024
 #define KERNEL_BFS              1024
+#define KERNEL_THREADS_MAX      256
+#define KERNEL_THREADS_MAX_CPU  16
 #define POWERTUNE_ENABLE        0
 #define LOGFILE_DISABLE         0
 #define SCRYPT_TMTO             0
@@ -3024,9 +3026,8 @@ static void autotune (hc_device_param_t *device_param)
   }
 
   // because of the balance we may have some free space left!
-  // at this point, allow a small variance to overdrive the limit
 
-  const int exec_left = (target_ms * 1.2) / exec_best;
+  const int exec_left = target_ms / exec_best;
 
   const int accel_left = kernel_accel_max / kernel_accel_best;
 
@@ -13090,14 +13091,19 @@ int main (int argc, char **argv)
                 if (data.quiet == 0) log_info ("           See the wiki on how to disable it: https://hashcat.net/wiki/doku.php?id=timeout_patch");
               }
             }
-            else if (vendor_id == VENDOR_ID_POCL)
+          }
+
+          if (device_type & CL_DEVICE_TYPE_CPU)
+          {
+            if (vendor_id == VENDOR_ID_AMD)
             {
               if (force == 0)
               {
                 log_info ("");
-                log_info ("ATTENTION! All pocl drivers are known to be broken due to broken LLVM <= 3.7");
+                log_info ("ATTENTION! OpenCL support for CPU of catalyst driver is not reliable.");
                 log_info ("You are STRONGLY encouraged not to use it");
                 log_info ("You can use --force to override this but do not post error reports if you do so");
+                log_info ("A good alternative is the free pocl, but make sure to use a version >= 3.8");
                 log_info ("");
 
                 return (-1);
@@ -13583,15 +13589,21 @@ int main (int argc, char **argv)
       /**
        * kernel threads: some algorithms need a fixed kernel-threads count
        *                 because of shared memory usage or bitslice
+       *                 there needs to be some upper limit, otherwise there's too much overhead
        */
 
-      uint kernel_threads = device_param->device_maxworkgroup_size;
+      uint kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size);
+
+      if (device_param->device_type & CL_DEVICE_TYPE_CPU)
+      {
+        kernel_threads = KERNEL_THREADS_MAX_CPU;
+      }
 
       if (hash_mode ==  1500) kernel_threads = 64; // DES
       if (hash_mode ==  3000) kernel_threads = 64; // DES
-      if (hash_mode ==  3200) kernel_threads = 8;  // blowfish
+      if (hash_mode ==  3200) kernel_threads = 8;  // Blowfish
       if (hash_mode ==  7500) kernel_threads = 64; // RC4
-      if (hash_mode ==  9000) kernel_threads = 8;  // blowfish
+      if (hash_mode ==  9000) kernel_threads = 8;  // Blowfish
       if (hash_mode ==  9700) kernel_threads = 64; // RC4
       if (hash_mode ==  9710) kernel_threads = 64; // RC4
       if (hash_mode ==  9800) kernel_threads = 64; // RC4
@@ -13658,18 +13670,18 @@ int main (int argc, char **argv)
             }
             else if (device_param->vendor_id == VENDOR_ID_NV)
             {
-              tmto_start = 3;
+              tmto_start = 2;
             }
           }
           else if (hash_mode == 9300)
           {
             if (device_param->vendor_id == VENDOR_ID_AMD)
             {
-              tmto_start = 3;
+              tmto_start = 2;
             }
             else if (device_param->vendor_id == VENDOR_ID_NV)
             {
-              tmto_start = 5;
+              tmto_start = 2;
             }
           }
         }