Rewrite some code to workaround strict aliasing rule violation for older compilers
[hashcat.git] / src / hashcat.c
index 7c01f7a..59e682a 100644 (file)
@@ -6,7 +6,7 @@
  * License.....: MIT
  */
 
-#ifdef DARWIN
+#if defined(DARWIN) || defined(__FreeBSD__)
 #include <stdio.h>
 #endif
 
@@ -86,7 +86,7 @@ double TARGET_MS_PROFILE[4]     = { 2, 12, 96, 480 };
 #define KERNEL_COMBS            1024
 #define KERNEL_BFS              1024
 #define KERNEL_THREADS_MAX      256
-#define KERNEL_THREADS_MAX_CPU  16
+#define KERNEL_THREADS_MAX_CPU  1
 #define POWERTUNE_ENABLE        0
 #define LOGFILE_DISABLE         0
 #define SCRYPT_TMTO             0
@@ -468,14 +468,14 @@ const char *USAGE_BIG[] =
   "     40 | md5($salt.unicode($pass))                        | Raw Hash, Salted and / or Iterated",
   "   3800 | md5($salt.$pass.$salt)                           | Raw Hash, Salted and / or Iterated",
   "   3710 | md5($salt.md5($pass))                            | Raw Hash, Salted and / or Iterated",
-  "   2600 | md5(md5($pass)                                   | Raw Hash, Salted and / or Iterated",
+  "   2600 | md5(md5($pass))                                  | Raw Hash, Salted and / or Iterated",
   "   4300 | md5(strtoupper(md5($pass)))                      | Raw Hash, Salted and / or Iterated",
   "   4400 | md5(sha1($pass))                                 | Raw Hash, Salted and / or Iterated",
   "    110 | sha1($pass.$salt)                                | Raw Hash, Salted and / or Iterated",
   "    120 | sha1($salt.$pass)                                | Raw Hash, Salted and / or Iterated",
   "    130 | sha1(unicode($pass).$salt)                       | Raw Hash, Salted and / or Iterated",
   "    140 | sha1($salt.unicode($pass))                       | Raw Hash, Salted and / or Iterated",
-  "   4500 | sha1(sha1($pass)                                 | Raw Hash, Salted and / or Iterated",
+  "   4500 | sha1(sha1($pass))                                | Raw Hash, Salted and / or Iterated",
   "   4700 | sha1(md5($pass))                                 | Raw Hash, Salted and / or Iterated",
   "   4900 | sha1($salt.$pass.$salt)                          | Raw Hash, Salted and / or Iterated",
   "   1410 | sha256($pass.$salt)                              | Raw Hash, Salted and / or Iterated",
@@ -925,6 +925,9 @@ void status_display ()
   if (data.devices_status == STATUS_INIT)     return;
   if (data.devices_status == STATUS_STARTING) return;
 
+  // in this case some required buffers are free'd, ascii_digest() would run into segfault
+  if (data.shutdown_inner == 1) return;
+
   if (data.machine_readable == 1)
   {
     status_display_machine_readable ();
@@ -1339,26 +1342,22 @@ void status_display ()
       }
       else
       {
-        char display_etc[32] = { 0 };
+        char display_etc[32]     = { 0 };
+        char display_runtime[32] = { 0 };
 
         struct tm tm_etc;
+        struct tm tm_runtime;
 
         struct tm *tmp = NULL;
 
         #ifdef WIN
-
         tmp = _gmtime64 (&sec_etc);
-
         #else
-
         tmp = gmtime (&sec_etc);
-
         #endif
 
         if (tmp != NULL)
         {
-          memset (&tm_etc, 0, sizeof (tm_etc));
-
           memcpy (&tm_etc, tmp, sizeof (tm_etc));
 
           format_timer_display (&tm_etc, display_etc, sizeof (display_etc));
@@ -1376,7 +1375,43 @@ void status_display ()
           if (etc[etc_len - 1] == '\n') etc[etc_len - 1] = 0;
           if (etc[etc_len - 2] == '\r') etc[etc_len - 2] = 0;
 
-          log_info ("Time.Estimated.: %s (%s)", etc, display_etc);
+          if (data.runtime)
+          {
+            time_t runtime_cur;
+
+            time (&runtime_cur);
+
+            #ifdef WIN
+
+            __time64_t runtime_left = data.proc_start + data.runtime - runtime_cur;
+
+            tmp = _gmtime64 (&runtime_left);
+
+            #else
+
+            time_t runtime_left = data.proc_start + data.runtime - runtime_cur;
+
+            tmp = gmtime (&runtime_left);
+
+            #endif
+
+            if ((tmp != NULL) && (runtime_left > 0) && (runtime_left < sec_etc))
+            {
+              memcpy (&tm_runtime, tmp, sizeof (tm_runtime));
+
+              format_timer_display (&tm_runtime, display_runtime, sizeof (display_runtime));
+
+              log_info ("Time.Estimated.: %s (%s), but limited (%s)", etc, display_etc, display_runtime);
+            }
+            else
+            {
+              log_info ("Time.Estimated.: %s (%s), but limit exceeded", etc, display_etc);
+            }
+          }
+          else
+          {
+            log_info ("Time.Estimated.: %s (%s)", etc, display_etc);
+          }
         }
       }
     }
@@ -1683,6 +1718,8 @@ static void status_benchmark ()
   if (data.devices_status == STATUS_INIT)     return;
   if (data.devices_status == STATUS_STARTING) return;
 
+  if (data.shutdown_inner == 1) return;
+
   if (data.machine_readable == 1)
   {
     status_benchmark_automate ();
@@ -2390,7 +2427,7 @@ static void process_stdout (hc_device_param_t *device_param, const uint pws_cnt)
 
   uint plain_len = 0;
 
-  const uint il_cnt = device_param->kernel_params_buf32[27]; // ugly, i know
+  const uint il_cnt = device_param->kernel_params_buf32[30]; // ugly, i know
 
   if (data.attack_mode == ATTACK_MODE_STRAIGHT)
   {
@@ -2649,8 +2686,8 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co
 {
   uint num_elements = num;
 
-  device_param->kernel_params_buf32[30] = data.combs_mode;
-  device_param->kernel_params_buf32[31] = num;
+  device_param->kernel_params_buf32[33] = data.combs_mode;
+  device_param->kernel_params_buf32[34] = num;
 
   uint kernel_threads = device_param->kernel_threads;
 
@@ -2667,9 +2704,6 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co
     case KERN_RUN_3:    kernel = device_param->kernel3;     break;
   }
 
-  hc_clSetKernelArg (data.ocl, kernel, 21, sizeof (cl_uint), device_param->kernel_params[21]);
-  hc_clSetKernelArg (data.ocl, kernel, 22, sizeof (cl_uint), device_param->kernel_params[22]);
-  hc_clSetKernelArg (data.ocl, kernel, 23, sizeof (cl_uint), device_param->kernel_params[23]);
   hc_clSetKernelArg (data.ocl, kernel, 24, sizeof (cl_uint), device_param->kernel_params[24]);
   hc_clSetKernelArg (data.ocl, kernel, 25, sizeof (cl_uint), device_param->kernel_params[25]);
   hc_clSetKernelArg (data.ocl, kernel, 26, sizeof (cl_uint), device_param->kernel_params[26]);
@@ -2678,6 +2712,9 @@ static void run_kernel (const uint kern_run, hc_device_param_t *device_param, co
   hc_clSetKernelArg (data.ocl, kernel, 29, sizeof (cl_uint), device_param->kernel_params[29]);
   hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
   hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
+  hc_clSetKernelArg (data.ocl, kernel, 32, sizeof (cl_uint), device_param->kernel_params[32]);
+  hc_clSetKernelArg (data.ocl, kernel, 33, sizeof (cl_uint), device_param->kernel_params[33]);
+  hc_clSetKernelArg (data.ocl, kernel, 34, sizeof (cl_uint), device_param->kernel_params[34]);
 
   cl_event event;
 
@@ -3029,8 +3066,8 @@ static void choose_kernel (hc_device_param_t *device_param, const uint attack_ex
 
       loop_left = MIN (loop_left, loop_step);
 
-      device_param->kernel_params_buf32[25] = loop_pos;
-      device_param->kernel_params_buf32[26] = loop_left;
+      device_param->kernel_params_buf32[28] = loop_pos;
+      device_param->kernel_params_buf32[29] = loop_left;
 
       run_kernel (KERN_RUN_2, device_param, pws_cnt, true, slow_iteration);
 
@@ -3170,9 +3207,9 @@ static double try_run (hc_device_param_t *device_param, const u32 kernel_accel,
 {
   const u32 kernel_power_try = device_param->device_processors * device_param->kernel_threads * kernel_accel;
 
-  device_param->kernel_params_buf32[25] = 0;
-  device_param->kernel_params_buf32[26] = kernel_loops; // not a bug, both need to be set
-  device_param->kernel_params_buf32[27] = kernel_loops; // because there's two variables for inner iters for slow and fast hashes
+  device_param->kernel_params_buf32[28] = 0;
+  device_param->kernel_params_buf32[29] = kernel_loops; // not a bug, both need to be set
+  device_param->kernel_params_buf32[30] = kernel_loops; // because there's two variables for inner iters for slow and fast hashes
 
   if (data.attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
   {
@@ -3496,9 +3533,9 @@ static void run_cracker (hc_device_param_t *device_param, const uint pws_cnt)
 
     salt_t *salt_buf = &data.salts_buf[salt_pos];
 
-    device_param->kernel_params_buf32[24] = salt_pos;
-    device_param->kernel_params_buf32[28] = salt_buf->digests_cnt;
-    device_param->kernel_params_buf32[29] = salt_buf->digests_offset;
+    device_param->kernel_params_buf32[27] = salt_pos;
+    device_param->kernel_params_buf32[31] = salt_buf->digests_cnt;
+    device_param->kernel_params_buf32[32] = salt_buf->digests_offset;
 
     FILE *combs_fp = device_param->combs_fp;
 
@@ -3534,7 +3571,7 @@ static void run_cracker (hc_device_param_t *device_param, const uint pws_cnt)
       device_param->innerloop_pos  = innerloop_pos;
       device_param->innerloop_left = innerloop_left;
 
-      device_param->kernel_params_buf32[27] = innerloop_left;
+      device_param->kernel_params_buf32[30] = innerloop_left;
 
       // i think we can get rid of this
       if (innerloop_left == 0)
@@ -4179,8 +4216,11 @@ static void *thread_monitor (void *p)
         {
           if (data.hm_nvapi)
           {
-            NV_GPU_PERF_POLICIES_INFO_PARAMS_V1   perfPolicies_info   = { 0 };
-            NV_GPU_PERF_POLICIES_STATUS_PARAMS_V1 perfPolicies_status = { 0 };
+            NV_GPU_PERF_POLICIES_INFO_PARAMS_V1   perfPolicies_info;
+            NV_GPU_PERF_POLICIES_STATUS_PARAMS_V1 perfPolicies_status;
+
+            memset (&perfPolicies_info,   0, sizeof (NV_GPU_PERF_POLICIES_INFO_PARAMS_V1));
+            memset (&perfPolicies_status, 0, sizeof (NV_GPU_PERF_POLICIES_STATUS_PARAMS_V1));
 
             perfPolicies_info.version   = MAKE_NVAPI_VERSION (NV_GPU_PERF_POLICIES_INFO_PARAMS_V1, 1);
             perfPolicies_status.version = MAKE_NVAPI_VERSION (NV_GPU_PERF_POLICIES_STATUS_PARAMS_V1, 1);
@@ -5244,12 +5284,12 @@ static void weak_hash_check (hc_device_param_t *device_param, const uint salt_po
 
   salt_t *salt_buf = &data.salts_buf[salt_pos];
 
-  device_param->kernel_params_buf32[24] = salt_pos;
-  device_param->kernel_params_buf32[27] = 1;
-  device_param->kernel_params_buf32[28] = salt_buf->digests_cnt;
-  device_param->kernel_params_buf32[29] = salt_buf->digests_offset;
-  device_param->kernel_params_buf32[30] = 0;
-  device_param->kernel_params_buf32[31] = 1;
+  device_param->kernel_params_buf32[27] = salt_pos;
+  device_param->kernel_params_buf32[30] = 1;
+  device_param->kernel_params_buf32[31] = salt_buf->digests_cnt;
+  device_param->kernel_params_buf32[32] = salt_buf->digests_offset;
+  device_param->kernel_params_buf32[33] = 0;
+  device_param->kernel_params_buf32[34] = 1;
 
   char *dictfile_old = data.dictfile;
 
@@ -5283,8 +5323,8 @@ static void weak_hash_check (hc_device_param_t *device_param, const uint salt_po
 
       loop_left = MIN (loop_left, loop_step);
 
-      device_param->kernel_params_buf32[25] = loop_pos;
-      device_param->kernel_params_buf32[26] = loop_left;
+      device_param->kernel_params_buf32[28] = loop_pos;
+      device_param->kernel_params_buf32[29] = loop_left;
 
       run_kernel (KERN_RUN_2, device_param, 1, false, 0);
     }
@@ -5302,14 +5342,14 @@ static void weak_hash_check (hc_device_param_t *device_param, const uint salt_po
    * cleanup
    */
 
-  device_param->kernel_params_buf32[24] = 0;
-  device_param->kernel_params_buf32[25] = 0;
-  device_param->kernel_params_buf32[26] = 0;
   device_param->kernel_params_buf32[27] = 0;
   device_param->kernel_params_buf32[28] = 0;
   device_param->kernel_params_buf32[29] = 0;
   device_param->kernel_params_buf32[30] = 0;
   device_param->kernel_params_buf32[31] = 0;
+  device_param->kernel_params_buf32[32] = 0;
+  device_param->kernel_params_buf32[33] = 0;
+  device_param->kernel_params_buf32[34] = 0;
 
   data.dictfile = dictfile_old;
 
@@ -6134,11 +6174,26 @@ int main (int argc, char **argv)
 
   char *exec_path = get_exec_path ();
 
-  #ifdef LINUX
+
+  #if defined(LINUX) || defined(__APPLE__) || defined(__FreeBSD__)
 
   char *resolved_install_folder = realpath (INSTALL_FOLDER, NULL);
   char *resolved_exec_path      = realpath (exec_path, NULL);
 
+  if (resolved_install_folder == NULL)
+  {
+    log_error ("ERROR: %s: %s", resolved_install_folder, strerror (errno));
+
+    return (-1);
+  }
+
+  if (resolved_exec_path == NULL)
+  {
+    log_error ("ERROR: %s: %s", resolved_exec_path, strerror (errno));
+
+    return (-1);
+  }
+
   char *install_dir = get_install_dir (resolved_exec_path);
   char *profile_dir = NULL;
   char *session_dir = NULL;
@@ -12966,6 +13021,7 @@ int main (int argc, char **argv)
                   // here we have in line_buf: ESSID:MAC1:MAC2   (without the plain)
                   // manipulate salt_buf
 
+                  memset (line_buf_cpy, 0, HCBUFSIZ);
                   memcpy (line_buf_cpy, line_buf, i);
 
                   char *mac2_pos = strrchr (line_buf_cpy, ':');
@@ -13864,6 +13920,8 @@ int main (int argc, char **argv)
 
         device_param->platform_devices_id = platform_devices_id;
 
+        device_param->platform = platform;
+
         // device_type
 
         cl_device_type device_type;
@@ -14464,10 +14522,15 @@ int main (int argc, char **argv)
      */
 
     #ifdef HAVE_HWMON
-    hm_attrs_t hm_adapters_adl[DEVICES_MAX]     = { { 0 } };
-    hm_attrs_t hm_adapters_nvapi[DEVICES_MAX]   = { { 0 } };
-    hm_attrs_t hm_adapters_nvml[DEVICES_MAX]    = { { 0 } };
-    hm_attrs_t hm_adapters_xnvctrl[DEVICES_MAX] = { { 0 } };
+    hm_attrs_t hm_adapters_adl[DEVICES_MAX];
+    hm_attrs_t hm_adapters_nvapi[DEVICES_MAX];
+    hm_attrs_t hm_adapters_nvml[DEVICES_MAX];
+    hm_attrs_t hm_adapters_xnvctrl[DEVICES_MAX];
+
+    memset (hm_adapters_adl,     0, sizeof (hm_adapters_adl));
+    memset (hm_adapters_nvapi,   0, sizeof (hm_adapters_nvapi));
+    memset (hm_adapters_nvml,    0, sizeof (hm_adapters_nvml));
+    memset (hm_adapters_xnvctrl, 0, sizeof (hm_adapters_xnvctrl));
 
     if (gpu_temp_disable == 0)
     {
@@ -14983,7 +15046,13 @@ int main (int argc, char **argv)
        * create context for each device
        */
 
-      device_param->context = hc_clCreateContext (data.ocl, NULL, 1, &device_param->device, NULL, NULL);
+      cl_context_properties properties[3];
+
+      properties[0] = CL_CONTEXT_PLATFORM;
+      properties[1] = (cl_context_properties) device_param->platform;
+      properties[2] = 0;
+
+      device_param->context = hc_clCreateContext (data.ocl, properties, 1, &device_param->device, NULL, NULL);
 
       /**
        * create command-queue
@@ -15002,6 +15071,9 @@ int main (int argc, char **argv)
 
       uint kernel_threads = MIN (KERNEL_THREADS_MAX, device_param->device_maxworkgroup_size);
 
+      if (hash_mode ==  8900) kernel_threads = 64; // Scrypt
+      if (hash_mode ==  9300) kernel_threads = 64; // Scrypt
+
       if (device_param->device_type & CL_DEVICE_TYPE_CPU)
       {
         kernel_threads = KERNEL_THREADS_MAX_CPU;
@@ -15057,7 +15129,7 @@ int main (int argc, char **argv)
 
       // scryptV stuff
 
-      size_t size_scryptV = 1;
+      size_t size_scrypt = 4;
 
       if ((hash_mode == 8900) || (hash_mode == 9300))
       {
@@ -15089,14 +15161,13 @@ int main (int argc, char **argv)
         else
         {
           // in case the user did not specify the tmto manually
-          // use some values known to run best (tested on 290x for AMD and 980ti for NV)
-          // but set the lower end only in case the user has a device with too less memory
+          // use some values known to run best (tested on 290x for AMD and GTX1080 for NV)
 
           if (hash_mode == 8900)
           {
             if (device_param->device_vendor_id == VENDOR_ID_AMD)
             {
-              tmto_start = 1;
+              tmto_start = 3;
             }
             else if (device_param->device_vendor_id == VENDOR_ID_NV)
             {
@@ -15111,12 +15182,12 @@ int main (int argc, char **argv)
             }
             else if (device_param->device_vendor_id == VENDOR_ID_NV)
             {
-              tmto_start = 2;
+              tmto_start = 4;
             }
           }
         }
 
-        data.scrypt_tmp_size = (128 * scrypt_r);
+        data.scrypt_tmp_size = (128 * scrypt_r * scrypt_p);
 
         device_param->kernel_accel_min = 1;
         device_param->kernel_accel_max = 8;
@@ -15125,15 +15196,22 @@ int main (int argc, char **argv)
 
         for (tmto = tmto_start; tmto < tmto_stop; tmto++)
         {
-          size_scryptV = (128 * scrypt_r) * scrypt_N;
+          size_scrypt = (128 * scrypt_r) * scrypt_N;
+
+          size_scrypt /= 1 << tmto;
 
-          size_scryptV /= 1 << tmto;
+          size_scrypt *= device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel_max;
 
-          size_scryptV *= device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel_max;
+          if ((size_scrypt / 4) > device_param->device_maxmem_alloc)
+          {
+            if (quiet == 0) log_info ("WARNING: Not enough single-block device memory allocatable to use --scrypt-tmto %d, increasing...", tmto);
+
+            continue;
+          }
 
-          if (size_scryptV > device_param->device_maxmem_alloc)
+          if (size_scrypt > device_param->device_global_mem)
           {
-            if (quiet == 0) log_info ("WARNING: Not enough device memory allocatable to use --scrypt-tmto %d, increasing...", tmto);
+            if (quiet == 0) log_info ("WARNING: Not enough total device memory allocatable to use --scrypt-tmto %d, increasing...", tmto);
 
             continue;
           }
@@ -15153,9 +15231,11 @@ int main (int argc, char **argv)
           return -1;
         }
 
-        if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %u\n", data.scrypt_tmto_final, size_scryptV);
+        if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %llu\n", data.scrypt_tmto_final, (unsigned long long int) size_scrypt);
       }
 
+      size_t size_scrypt4 = size_scrypt / 4;
+
       /**
        * some algorithms need a fixed kernel-loops count
        */
@@ -15372,7 +15452,10 @@ int main (int argc, char **argv)
           + size_rules
           + size_rules_c
           + size_salts
-          + size_scryptV
+          + size_scrypt4
+          + size_scrypt4
+          + size_scrypt4
+          + size_scrypt4
           + size_shown
           + size_tm
           + size_tmps;
@@ -15932,7 +16015,10 @@ int main (int argc, char **argv)
       device_param->d_digests_shown = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_shown,   NULL);
       device_param->d_salt_bufs     = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY,   size_salts,   NULL);
       device_param->d_result        = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_results, NULL);
-      device_param->d_scryptV_buf   = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_scryptV, NULL);
+      device_param->d_scryptV0_buf  = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_scrypt4, NULL);
+      device_param->d_scryptV1_buf  = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_scrypt4, NULL);
+      device_param->d_scryptV2_buf  = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_scrypt4, NULL);
+      device_param->d_scryptV3_buf  = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_WRITE,  size_scrypt4, NULL);
 
       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_a,    CL_TRUE, 0, bitmap_size,  bitmap_s1_a,        0, NULL, NULL);
       hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_bitmap_s1_b,    CL_TRUE, 0, bitmap_size,  bitmap_s1_b,        0, NULL, NULL);
@@ -16000,17 +16086,17 @@ int main (int argc, char **argv)
        * kernel args
        */
 
-      device_param->kernel_params_buf32[21] = bitmap_mask;
-      device_param->kernel_params_buf32[22] = bitmap_shift1;
-      device_param->kernel_params_buf32[23] = bitmap_shift2;
-      device_param->kernel_params_buf32[24] = 0; // salt_pos
-      device_param->kernel_params_buf32[25] = 0; // loop_pos
-      device_param->kernel_params_buf32[26] = 0; // loop_cnt
-      device_param->kernel_params_buf32[27] = 0; // kernel_rules_cnt
-      device_param->kernel_params_buf32[28] = 0; // digests_cnt
-      device_param->kernel_params_buf32[29] = 0; // digests_offset
-      device_param->kernel_params_buf32[30] = 0; // combs_mode
-      device_param->kernel_params_buf32[31] = 0; // gid_max
+      device_param->kernel_params_buf32[24] = bitmap_mask;
+      device_param->kernel_params_buf32[25] = bitmap_shift1;
+      device_param->kernel_params_buf32[26] = bitmap_shift2;
+      device_param->kernel_params_buf32[27] = 0; // salt_pos
+      device_param->kernel_params_buf32[28] = 0; // loop_pos
+      device_param->kernel_params_buf32[29] = 0; // loop_cnt
+      device_param->kernel_params_buf32[30] = 0; // kernel_rules_cnt
+      device_param->kernel_params_buf32[31] = 0; // digests_cnt
+      device_param->kernel_params_buf32[32] = 0; // digests_offset
+      device_param->kernel_params_buf32[33] = 0; // combs_mode
+      device_param->kernel_params_buf32[34] = 0; // gid_max
 
       device_param->kernel_params[ 0] = (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
                                       ? &device_param->d_pws_buf
@@ -16034,10 +16120,10 @@ int main (int argc, char **argv)
       device_param->kernel_params[17] = &device_param->d_salt_bufs;
       device_param->kernel_params[18] = &device_param->d_esalt_bufs;
       device_param->kernel_params[19] = &device_param->d_result;
-      device_param->kernel_params[20] = &device_param->d_scryptV_buf;
-      device_param->kernel_params[21] = &device_param->kernel_params_buf32[21];
-      device_param->kernel_params[22] = &device_param->kernel_params_buf32[22];
-      device_param->kernel_params[23] = &device_param->kernel_params_buf32[23];
+      device_param->kernel_params[20] = &device_param->d_scryptV0_buf;
+      device_param->kernel_params[21] = &device_param->d_scryptV1_buf;
+      device_param->kernel_params[22] = &device_param->d_scryptV2_buf;
+      device_param->kernel_params[23] = &device_param->d_scryptV3_buf;
       device_param->kernel_params[24] = &device_param->kernel_params_buf32[24];
       device_param->kernel_params[25] = &device_param->kernel_params_buf32[25];
       device_param->kernel_params[26] = &device_param->kernel_params_buf32[26];
@@ -16046,6 +16132,9 @@ int main (int argc, char **argv)
       device_param->kernel_params[29] = &device_param->kernel_params_buf32[29];
       device_param->kernel_params[30] = &device_param->kernel_params_buf32[30];
       device_param->kernel_params[31] = &device_param->kernel_params_buf32[31];
+      device_param->kernel_params[32] = &device_param->kernel_params_buf32[32];
+      device_param->kernel_params[33] = &device_param->kernel_params_buf32[33];
+      device_param->kernel_params[34] = &device_param->kernel_params_buf32[34];
 
       device_param->kernel_params_mp_buf64[3] = 0;
       device_param->kernel_params_mp_buf32[4] = 0;
@@ -16209,7 +16298,7 @@ int main (int argc, char **argv)
       hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel2, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
       hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel3, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
 
-      for (uint i = 0; i <= 20; i++)
+      for (uint i = 0; i <= 23; i++)
       {
         hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
         hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_mem), device_param->kernel_params[i]);
@@ -16219,7 +16308,7 @@ int main (int argc, char **argv)
         if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_mem), device_param->kernel_params[i]);
       }
 
-      for (uint i = 21; i <= 31; i++)
+      for (uint i = 24; i <= 34; i++)
       {
         hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_uint), device_param->kernel_params[i]);
         hc_clSetKernelArg (data.ocl, device_param->kernel2, i, sizeof (cl_uint), device_param->kernel_params[i]);
@@ -17429,6 +17518,36 @@ int main (int argc, char **argv)
 
             mask = mask + str_pos + 1;
           }
+
+          /**
+           * What follows is a very special case where "\," is within the mask field of a line in a .hcmask file only because otherwise (without the "\")
+           * it would be interpreted as a custom charset definition.
+           *
+           * We need to replace all "\," with just "," within the mask (but allow the special case "\\," which means "\" followed by ",")
+           * Note: "\\" is not needed to replace all "\" within the mask! The meaning of "\\" within a line containing the string "\\," is just to allow "\" followed by ","
+           */
+
+          uint mask_len_cur = strlen (mask);
+
+          uint mask_out_pos = 0;
+          char mask_prev = 0;
+
+          for (uint mask_iter = 0; mask_iter < mask_len_cur; mask_iter++, mask_out_pos++)
+          {
+            if (mask[mask_iter] == ',')
+            {
+              if (mask_prev == '\\')
+              {
+                mask_out_pos -= 1; // this means: skip the previous "\"
+              }
+            }
+
+            mask_prev = mask[mask_iter];
+
+            mask[mask_out_pos] = mask[mask_iter];
+          }
+
+          mask[mask_out_pos] = '\0';
         }
 
         if ((attack_mode == ATTACK_MODE_HYBRID1) || (attack_mode == ATTACK_MODE_HYBRID2))
@@ -18276,15 +18395,26 @@ int main (int argc, char **argv)
 
             log_info ("");
 
-            if (stdout_flag == 0) status_display ();
+            status_display ();
 
             log_info ("");
           }
+          else
+          {
+            if (status == 1)
+            {
+              status_display ();
+            }
+          }
         }
 
         if (induction_dictionaries_cnt)
         {
           qsort (induction_dictionaries, induction_dictionaries_cnt, sizeof (char *), sort_by_mtime);
+
+          // yeah, this next statement is a little hack to make sure that --loopback runs correctly (because with it we guarantee that the loop iterates one more time)
+
+          dictpos--;
         }
 
         time_t runtime_stop;
@@ -18454,7 +18584,10 @@ int main (int argc, char **argv)
       if (device_param->d_tmps)             hc_clReleaseMemObject     (data.ocl, device_param->d_tmps);
       if (device_param->d_hooks)            hc_clReleaseMemObject     (data.ocl, device_param->d_hooks);
       if (device_param->d_result)           hc_clReleaseMemObject     (data.ocl, device_param->d_result);
-      if (device_param->d_scryptV_buf)      hc_clReleaseMemObject     (data.ocl, device_param->d_scryptV_buf);
+      if (device_param->d_scryptV0_buf)     hc_clReleaseMemObject     (data.ocl, device_param->d_scryptV0_buf);
+      if (device_param->d_scryptV1_buf)     hc_clReleaseMemObject     (data.ocl, device_param->d_scryptV1_buf);
+      if (device_param->d_scryptV2_buf)     hc_clReleaseMemObject     (data.ocl, device_param->d_scryptV2_buf);
+      if (device_param->d_scryptV3_buf)     hc_clReleaseMemObject     (data.ocl, device_param->d_scryptV3_buf);
       if (device_param->d_root_css_buf)     hc_clReleaseMemObject     (data.ocl, device_param->d_root_css_buf);
       if (device_param->d_markov_css_buf)   hc_clReleaseMemObject     (data.ocl, device_param->d_markov_css_buf);
       if (device_param->d_tm_c)             hc_clReleaseMemObject     (data.ocl, device_param->d_tm_c);