- 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 91f4506..f289bb1 100644 (file)
--- a/src/oclHashcat.c
+++ b/src/oclHashcat.c
@@ -148,7 +148,7 @@ double TARGET_MS_PROFILE[3]     = { 8, 16, 96 };
 
 #define MAX_DICTSTAT            10000
 
-#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 135
+#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 136
 
 #define global_free(attr)       \
 {                               \
@@ -268,6 +268,7 @@ static uint default_benchmark_algorithms[NUM_DEFAULT_BENCHMARK_ALGORITHMS] =
   8700,
   9100,
   133,
+  13500,
   11600,
   12500,
   13000,
@@ -696,6 +697,7 @@ const char *USAGE_BIG[] =
   "   8700 = Lotus Notes/Domino 6",
   "   9100 = Lotus Notes/Domino 8",
   "    133 = PeopleSoft",
+  "  13500 = PeopleSoft Token",
   "",
   "[[ Archives ]]",
   "",
@@ -2422,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 };
 
@@ -5311,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
    */
@@ -5971,7 +5986,7 @@ int main (int argc, char **argv)
     return (-1);
   }
 
-  if (hash_mode_chgd && hash_mode > 13400) // just added to remove compiler warnings for hash_mode_chgd
+  if (hash_mode_chgd && hash_mode > 13500) // just added to remove compiler warnings for hash_mode_chgd
   {
     log_error ("ERROR: Invalid hash-type specified");
 
@@ -6318,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))
       {
@@ -6333,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)
@@ -7609,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;
@@ -10297,6 +10313,28 @@ int main (int argc, char **argv)
                    dgst_pos3   = 3;
                    break;
 
+      case 13500:  hash_type   = HASH_TYPE_SHA1;
+                   salt_type   = SALT_TYPE_EMBEDDED;
+                   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_PSTOKEN;
+                   dgst_size   = DGST_SIZE_4_5;
+                   parse_func  = pstoken_parse_hash;
+                   sort_by_digest = sort_by_digest_4_5;
+                   opti_type   = OPTI_TYPE_ZERO_BYTE
+                               | OPTI_TYPE_PRECOMPUTE_INIT
+                               | OPTI_TYPE_EARLY_SKIP
+                               | OPTI_TYPE_NOT_ITERATED
+                               | OPTI_TYPE_PREPENDED_SALT
+                               | OPTI_TYPE_RAW_HASH;
+                   dgst_pos0   = 3;
+                   dgst_pos1   = 4;
+                   dgst_pos2   = 2;
+                   dgst_pos3   = 1;
+                   break;
+
       default:     usage_mini_print (PROGNAME); return (-1);
     }
 
@@ -10402,6 +10440,7 @@ int main (int argc, char **argv)
       case 13000:  esalt_size = sizeof (rar5_t);          break;
       case 13100:  esalt_size = sizeof (krb5tgs_t);       break;
       case 13400:  esalt_size = sizeof (keepass_t);       break;
+      case 13500:  esalt_size = sizeof (pstoken_t);       break;
     }
 
     data.esalt_size = esalt_size;
@@ -11507,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;
         }
       }
 
@@ -13902,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