improved autotune engine

[hashcat.git] / src / oclHashcat.c

diff --git a/src/oclHashcat.c b/src/oclHashcat.c
index 4b1ec08..035a97d 100644 (file)
--- a/src/oclHashcat.c
+++ b/src/oclHashcat.c
@@ -2424,8 +2424,18 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co
 
     hc_clGetKernelWorkGroupInfo (data.ocl, kernel, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &workgroup_size, NULL);
 
+    if (kern_run == KERN_RUN_2)
+    {
+      if (data.opti_type & OPTI_TYPE_SLOW_HASH_SIMD)
+      {
+        num_elements = CEIL ((float) num_elements / device_param->vector_width);
+      }
+    }
+
     if (kernel_threads > workgroup_size) kernel_threads = workgroup_size;
 
+    while (num_elements % kernel_threads) num_elements++;
+
     const size_t global_work_size[3] = { num_elements,   1, 1 };
     const size_t local_work_size[3]  = { kernel_threads, 1, 1 };
 
@@ -3008,6 +3018,20 @@ 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 accel_left = kernel_accel_max / kernel_accel_best;
+
+  const int exec_accel_min = MIN (exec_left, accel_left);
+
+  if (exec_accel_min)
+  {
+    kernel_accel_best *= exec_accel_min;
+  }
+
   // reset timer
 
   device_param->exec_pos = 0;
@@ -5313,6 +5337,9 @@ int main (int argc, char **argv)
   if (getenv ("CUDA_CACHE_DISABLE") == NULL)
     putenv ((char *) "CUDA_CACHE_DISABLE=1");
 
+  if (getenv ("POCL_KERNEL_CACHE") == NULL)
+    putenv ((char *) "POCL_KERNEL_CACHE=0");
+
   /**
    * Real init
    */
@@ -6320,13 +6347,7 @@ int main (int argc, char **argv)
 
   if (loopback == 1)
   {
-    if (attack_mode == ATTACK_MODE_BF)
-    {
-      log_error ("ERROR: Parameter loopback not allowed in attack-mode 3");
-
-      return (-1);
-    }
-    else if (attack_mode == ATTACK_MODE_STRAIGHT)
+    if (attack_mode == ATTACK_MODE_STRAIGHT)
     {
       if ((rp_files_cnt == 0) && (rp_gen == 0))
       {
@@ -6335,6 +6356,12 @@ int main (int argc, char **argv)
         return (-1);
       }
     }
+    else
+    {
+      log_error ("ERROR: Parameter loopback allowed in attack-mode 0 only");
+
+      return (-1);
+    }
   }
 
   if (debug_mode > 0)
@@ -7611,7 +7638,8 @@ int main (int argc, char **argv)
                    dgst_size   = DGST_SIZE_4_4;
                    parse_func  = phpass_parse_hash;
                    sort_by_digest = sort_by_digest_4_4;
-                   opti_type   = OPTI_TYPE_ZERO_BYTE;
+                   opti_type   = OPTI_TYPE_ZERO_BYTE
+                               | OPTI_TYPE_SLOW_HASH_SIMD;
                    dgst_pos0   = 0;
                    dgst_pos1   = 1;
                    dgst_pos2   = 2;
@@ -13929,7 +13957,7 @@ int main (int argc, char **argv)
 
       // we don't have sm_* on vendors not NV but it doesn't matter
 
-      snprintf (build_opts, sizeof (build_opts) - 1, "-I\"%s/\" -DVENDOR_ID=%u -DCUDA_ARCH=%d -DVECT_SIZE=%u -DDEVICE_TYPE=%u", shared_dir, device_param->vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type);
+      snprintf (build_opts, sizeof (build_opts) - 1, "-cl-std=CL1.1 -I\"%s/\" -DVENDOR_ID=%u -DCUDA_ARCH=%d -DVECT_SIZE=%u -DDEVICE_TYPE=%u", shared_dir, device_param->vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type);
 
       /**
        * main kernel