- Added inline declaration to functions from simd.c, common.c, rp.c and types_ocl...

[hashcat.git] / src / oclHashcat.c

diff --git a/src/oclHashcat.c b/src/oclHashcat.c
index e718e9b..f289bb1 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 };
 
@@ -5313,6 +5323,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 +6333,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 +6342,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 +7624,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;
@@ -10304,8 +10318,8 @@ int main (int argc, char **argv)
                    attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
                    opts_type   = OPTS_TYPE_PT_GENERATE_BE
                                | OPTS_TYPE_PT_UNICODE
-                              | OPTS_TYPE_PT_ADD80;
-                   kern_type   = KERN_TYPE_SHA1_SLTPWU;
+                                                | OPTS_TYPE_PT_ADD80;
+                   kern_type   = KERN_TYPE_PSTOKEN;
                    dgst_size   = DGST_SIZE_4_5;
                    parse_func  = pstoken_parse_hash;
                    sort_by_digest = sort_by_digest_4_5;
@@ -11532,6 +11546,8 @@ int main (int argc, char **argv)
                       break;
           case 13400: ((keepass_t *) hashes_buf[0].esalt)->version       = 2;
                       break;
+          case 13500: ((pstoken_t *) hashes_buf[0].esalt)->salt_len      = 113;
+                      break;
         }
       }
 
@@ -13927,7 +13943,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