diff --git a/OpenCL/markov_be.cl b/OpenCL/markov_be.cl
index 4f337eb13..955d44c56 100644
--- a/OpenCL/markov_be.cl
+++ b/OpenCL/markov_be.cl
@@ -9,25 +9,8 @@
 
 #include "inc_types.cl"
 
-inline void generate_pw (u32 pw_buf[16], __global const cs_t *root_css_buf, __global const cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val)
+inline void generate_pw (u32 pw_buf[64], __global const cs_t *root_css_buf, __global const cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val)
 {
-  pw_buf[ 0] = 0;
-  pw_buf[ 1] = 0;
-  pw_buf[ 2] = 0;
-  pw_buf[ 3] = 0;
-  pw_buf[ 4] = 0;
-  pw_buf[ 5] = 0;
-  pw_buf[ 6] = 0;
-  pw_buf[ 7] = 0;
-  pw_buf[ 8] = 0;
-  pw_buf[ 9] = 0;
-  pw_buf[10] = 0;
-  pw_buf[11] = 0;
-  pw_buf[12] = 0;
-  pw_buf[13] = 0;
-  pw_buf[14] = 0;
-  pw_buf[15] = 0;
-
   __global const cs_t *cs = &root_css_buf[pw_r_len];
 
   u32 i;
@@ -67,26 +50,15 @@ __kernel void l_markov (__global pw_t *pws_buf_l, __global const cs_t *root_css_
 
   if (gid >= gid_max) return;
 
-  u32 pw_buf[16];
+  u32 pw_buf[64] = { 0 };
 
   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid);
 
-  pws_buf_l[gid].i[ 0] = pw_buf[ 0];
-  pws_buf_l[gid].i[ 1] = pw_buf[ 1];
-  pws_buf_l[gid].i[ 2] = pw_buf[ 2];
-  pws_buf_l[gid].i[ 3] = pw_buf[ 3];
-  pws_buf_l[gid].i[ 4] = pw_buf[ 4];
-  pws_buf_l[gid].i[ 5] = pw_buf[ 5];
-  pws_buf_l[gid].i[ 6] = pw_buf[ 6];
-  pws_buf_l[gid].i[ 7] = pw_buf[ 7];
-  pws_buf_l[gid].i[ 8] = pw_buf[ 8];
-  pws_buf_l[gid].i[ 9] = pw_buf[ 9];
-  pws_buf_l[gid].i[10] = pw_buf[10];
-  pws_buf_l[gid].i[11] = pw_buf[11];
-  pws_buf_l[gid].i[12] = pw_buf[12];
-  pws_buf_l[gid].i[13] = pw_buf[13];
-  pws_buf_l[gid].i[14] = pw_buf[14];
-  pws_buf_l[gid].i[15] = pw_buf[15];
+  #pragma unroll
+  for (int idx = 0; idx < 64; idx++)
+  {
+    pws_buf_l[gid].i[idx] = pw_buf[idx];
+  }
 
   pws_buf_l[gid].pw_len = pw_l_len + pw_r_len;
 }
@@ -97,7 +69,7 @@ __kernel void r_markov (__global bf_t *pws_buf_r, __global const cs_t *root_css_
 
   if (gid >= gid_max) return;
 
-  u32 pw_buf[16];
+  u32 pw_buf[64] = { 0 };
 
   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid);
 
@@ -110,18 +82,15 @@ __kernel void C_markov (__global comb_t *pws_buf, __global const cs_t *root_css_
 
   if (gid >= gid_max) return;
 
-  u32 pw_buf[16];
+  u32 pw_buf[64] = { 0 };
 
   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid);
 
-  pws_buf[gid].i[ 0] = pw_buf[ 0];
-  pws_buf[gid].i[ 1] = pw_buf[ 1];
-  pws_buf[gid].i[ 2] = pw_buf[ 2];
-  pws_buf[gid].i[ 3] = pw_buf[ 3];
-  pws_buf[gid].i[ 4] = pw_buf[ 4];
-  pws_buf[gid].i[ 5] = pw_buf[ 5];
-  pws_buf[gid].i[ 6] = pw_buf[ 6];
-  pws_buf[gid].i[ 7] = pw_buf[ 7];
+  #pragma unroll
+  for (int idx = 0; idx < 64; idx++)
+  {
+    pws_buf[gid].i[idx] = pw_buf[idx];
+  }
 
   pws_buf[gid].pw_len = pw_len;
 }
diff --git a/OpenCL/markov_le.cl b/OpenCL/markov_le.cl
index f20f4770a..cb17da083 100644
--- a/OpenCL/markov_le.cl
+++ b/OpenCL/markov_le.cl
@@ -9,25 +9,8 @@
 
 #include "inc_types.cl"
 
-inline void generate_pw (u32 pw_buf[16], __global const cs_t *root_css_buf, __global const cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val)
+inline void generate_pw (u32 pw_buf[64], __global const cs_t *root_css_buf, __global const cs_t *markov_css_buf, const u32 pw_l_len, const u32 pw_r_len, const u32 mask80, const u32 bits14, const u32 bits15, u64 val)
 {
-  pw_buf[ 0] = 0;
-  pw_buf[ 1] = 0;
-  pw_buf[ 2] = 0;
-  pw_buf[ 3] = 0;
-  pw_buf[ 4] = 0;
-  pw_buf[ 5] = 0;
-  pw_buf[ 6] = 0;
-  pw_buf[ 7] = 0;
-  pw_buf[ 8] = 0;
-  pw_buf[ 9] = 0;
-  pw_buf[10] = 0;
-  pw_buf[11] = 0;
-  pw_buf[12] = 0;
-  pw_buf[13] = 0;
-  pw_buf[14] = 0;
-  pw_buf[15] = 0;
-
   __global const cs_t *cs = &root_css_buf[pw_r_len];
 
   u32 i;
@@ -67,26 +50,15 @@ __kernel void l_markov (__global pw_t *pws_buf_l, __global const cs_t *root_css_
 
   if (gid >= gid_max) return;
 
-  u32 pw_buf[16];
+  u32 pw_buf[64] = { 0 };
 
   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_l_len, pw_r_len, mask80, bits14, bits15, off + gid);
 
-  pws_buf_l[gid].i[ 0] = pw_buf[ 0];
-  pws_buf_l[gid].i[ 1] = pw_buf[ 1];
-  pws_buf_l[gid].i[ 2] = pw_buf[ 2];
-  pws_buf_l[gid].i[ 3] = pw_buf[ 3];
-  pws_buf_l[gid].i[ 4] = pw_buf[ 4];
-  pws_buf_l[gid].i[ 5] = pw_buf[ 5];
-  pws_buf_l[gid].i[ 6] = pw_buf[ 6];
-  pws_buf_l[gid].i[ 7] = pw_buf[ 7];
-  pws_buf_l[gid].i[ 8] = pw_buf[ 8];
-  pws_buf_l[gid].i[ 9] = pw_buf[ 9];
-  pws_buf_l[gid].i[10] = pw_buf[10];
-  pws_buf_l[gid].i[11] = pw_buf[11];
-  pws_buf_l[gid].i[12] = pw_buf[12];
-  pws_buf_l[gid].i[13] = pw_buf[13];
-  pws_buf_l[gid].i[14] = pw_buf[14];
-  pws_buf_l[gid].i[15] = pw_buf[15];
+  #pragma unroll
+  for (int idx = 0; idx < 64; idx++)
+  {
+    pws_buf_l[gid].i[idx] = pw_buf[idx];
+  }
 
   pws_buf_l[gid].pw_len = pw_l_len + pw_r_len;
 }
@@ -97,7 +69,7 @@ __kernel void r_markov (__global bf_t *pws_buf_r, __global const cs_t *root_css_
 
   if (gid >= gid_max) return;
 
-  u32 pw_buf[16];
+  u32 pw_buf[64] = { 0 };
 
   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_r_len, 0, 0, 0, 0, off + gid);
 
@@ -110,18 +82,15 @@ __kernel void C_markov (__global comb_t *pws_buf, __global const cs_t *root_css_
 
   if (gid >= gid_max) return;
 
-  u32 pw_buf[16];
+  u32 pw_buf[64] = { 0 };
 
   generate_pw (pw_buf, root_css_buf, markov_css_buf, pw_len, 0, mask80, bits14, bits15, off + gid);
 
-  pws_buf[gid].i[ 0] = pw_buf[ 0];
-  pws_buf[gid].i[ 1] = pw_buf[ 1];
-  pws_buf[gid].i[ 2] = pw_buf[ 2];
-  pws_buf[gid].i[ 3] = pw_buf[ 3];
-  pws_buf[gid].i[ 4] = pw_buf[ 4];
-  pws_buf[gid].i[ 5] = pw_buf[ 5];
-  pws_buf[gid].i[ 6] = pw_buf[ 6];
-  pws_buf[gid].i[ 7] = pw_buf[ 7];
+  #pragma unroll
+  for (int idx = 0; idx < 64; idx++)
+  {
+    pws_buf[gid].i[idx] = pw_buf[idx];
+  }
 
   pws_buf[gid].pw_len = pw_len;
 }
diff --git a/hashcat.hcstat b/hashcat.hcstat
deleted file mode 100644
index 983d73638..000000000
Binary files a/hashcat.hcstat and /dev/null differ
diff --git a/include/mpsp.h b/include/mpsp.h
index 797f63d09..ee2e223ec 100644
--- a/include/mpsp.h
+++ b/include/mpsp.h
@@ -12,9 +12,10 @@
 
 #define CHARSIZ       0x100
 
-#define SP_HCSTAT     "hashcat.hcstat"
+#define SP_HCSTAT     "hashcat.hcstat2"
+#define SP_VERSION    (0x6863737461740000 | 0x0002)
 #define SP_PW_MIN     2
-#define SP_PW_MAX     64
+#define SP_PW_MAX     256
 #define SP_ROOT_CNT   (SP_PW_MAX * CHARSIZ)
 #define SP_MARKOV_CNT (SP_PW_MAX * CHARSIZ * CHARSIZ)
 
diff --git a/src/mpsp.c b/src/mpsp.c
index 6473cbe34..a2472b3b0 100644
--- a/src/mpsp.c
+++ b/src/mpsp.c
@@ -7,6 +7,7 @@
 #include "types.h"
 #include "memory.h"
 #include "event.h"
+#include "bitops.h"
 #include "logfile.h"
 #include "convert.h"
 #include "filehandling.h"
@@ -687,6 +688,27 @@ static int sp_setup_tbl (hashcat_ctx_t *hashcat_ctx)
     return -1;
   }
 
+  u64 v = 0;
+  u64 z = 0;
+
+  if (fread (&v, sizeof (u64), 1, fd) != 1)
+  {
+    event_log_error (hashcat_ctx, "%s: Could not load data.", hcstat);
+
+    fclose (fd);
+
+    return -1;
+  }
+
+  if (fread (&z, sizeof (u64), 1, fd) != 1)
+  {
+    event_log_error (hashcat_ctx, "%s: Could not load data.", hcstat);
+
+    fclose (fd);
+
+    return -1;
+  }
+
   if (fread (root_stats_buf, sizeof (u64), SP_ROOT_CNT, fd) != SP_ROOT_CNT)
   {
     event_log_error (hashcat_ctx, "%s: Could not load data.", hcstat);
@@ -707,6 +729,34 @@ static int sp_setup_tbl (hashcat_ctx_t *hashcat_ctx)
 
   fclose (fd);
 
+  /**
+   * switch endianess
+   */
+
+  v = byte_swap_64 (v);
+  z = byte_swap_64 (z);
+
+  for (int i = 0; i < SP_ROOT_CNT; i++)   root_stats_buf[i]   = byte_swap_64 (root_stats_buf[i]);
+  for (int i = 0; i < SP_MARKOV_CNT; i++) markov_stats_buf[i] = byte_swap_64 (markov_stats_buf[i]);
+
+  /**
+   * verify header
+   */
+
+  if (v != SP_VERSION)
+  {
+    event_log_error (hashcat_ctx, "%s: Invalid header", hcstat);
+
+    return -1;
+  }
+
+  if (z != 0)
+  {
+    event_log_error (hashcat_ctx, "%s: Invalid header", hcstat);
+
+    return -1;
+  }
+
   /**
    * Markov modifier of hcstat_table on user request
    */