diff --git a/OpenCL/amp_a0.cl b/OpenCL/amp_a0.cl
index 9720ba485..4cef595ad 100644
--- a/OpenCL/amp_a0.cl
+++ b/OpenCL/amp_a0.cl
@@ -16,32 +16,10 @@ __kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global const ke
 
   if (gid >= gid_max) return;
 
+  if (rules_buf[0].cmds[0] == RULE_OP_MANGLE_NOOP && rules_buf[0].cmds[1] == 0) return;
+
   const u32 pw_len = pws[gid].pw_len;
 
-  if (rules_buf[0].cmds[0] == RULE_OP_MANGLE_NOOP && rules_buf[0].cmds[1] == 0)
-  {
-    pws_amp[gid].i[ 0] = pws[gid].i[ 0];
-    pws_amp[gid].i[ 1] = pws[gid].i[ 1];
-    pws_amp[gid].i[ 2] = pws[gid].i[ 2];
-    pws_amp[gid].i[ 3] = pws[gid].i[ 3];
-    pws_amp[gid].i[ 4] = pws[gid].i[ 4];
-    pws_amp[gid].i[ 5] = pws[gid].i[ 5];
-    pws_amp[gid].i[ 6] = pws[gid].i[ 6];
-    pws_amp[gid].i[ 7] = pws[gid].i[ 7];
-    pws_amp[gid].i[ 8] = pws[gid].i[ 8];
-    pws_amp[gid].i[ 9] = pws[gid].i[ 9];
-    pws_amp[gid].i[10] = pws[gid].i[10];
-    pws_amp[gid].i[11] = pws[gid].i[11];
-    pws_amp[gid].i[12] = pws[gid].i[12];
-    pws_amp[gid].i[13] = pws[gid].i[13];
-    pws_amp[gid].i[14] = pws[gid].i[14];
-    pws_amp[gid].i[15] = pws[gid].i[15];
-
-    pws_amp[gid].pw_len = pws[gid].pw_len;
-
-    return;
-  }
-
   u32 w0[4];
   u32 w1[4];
 
diff --git a/OpenCL/amp_a3.cl b/OpenCL/amp_a3.cl
index dfafdbd2f..576e0bd49 100644
--- a/OpenCL/amp_a3.cl
+++ b/OpenCL/amp_a3.cl
@@ -13,48 +13,7 @@ __kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global const ke
 
   if (gid >= gid_max) return;
 
-  const u32 pw_len = pws[gid].pw_len;
-
-  u32 w0[4];
-  u32 w1[4];
-  u32 w2[4];
-  u32 w3[4];
-
-  w0[0] = pws[gid].i[ 0];
-  w0[1] = pws[gid].i[ 1];
-  w0[2] = pws[gid].i[ 2];
-  w0[3] = pws[gid].i[ 3];
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
-
   const u32 w0r = bfs_buf[0].i;
 
-  pws_amp[gid].i[ 0] = w0[0] | w0r;
-  pws_amp[gid].i[ 1] = w0[1];
-  pws_amp[gid].i[ 2] = w0[2];
-  pws_amp[gid].i[ 3] = w0[3];
-  pws_amp[gid].i[ 4] = w1[0];
-  pws_amp[gid].i[ 5] = w1[1];
-  pws_amp[gid].i[ 6] = w1[2];
-  pws_amp[gid].i[ 7] = w1[3];
-  pws_amp[gid].i[ 8] = w2[0];
-  pws_amp[gid].i[ 9] = w2[1];
-  pws_amp[gid].i[10] = w2[2];
-  pws_amp[gid].i[11] = w2[3];
-  pws_amp[gid].i[12] = w3[0];
-  pws_amp[gid].i[13] = w3[1];
-  pws_amp[gid].i[14] = w3[2];
-  pws_amp[gid].i[15] = w3[3];
-
-  pws_amp[gid].pw_len = pw_len;
+  pws_amp[gid].i[0] |= w0r;
 }
diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl
index c9e9dc0b6..79af1b077 100644
--- a/OpenCL/inc_types.cl
+++ b/OpenCL/inc_types.cl
@@ -674,8 +674,8 @@ typedef struct digest
 
 typedef struct salt
 {
-  u32 salt_buf[16];
-  u32 salt_buf_pc[16];
+  u32 salt_buf[64];
+  u32 salt_buf_pc[64];
 
   u32 salt_len;
   u32 salt_len_pc;
@@ -1494,19 +1494,9 @@ typedef struct
 
 } kernel_rule_t;
 
-typedef struct
+typedef struct pw
 {
-  u32 salt_pos;
-  u32 digest_pos;
-  u32 hash_pos;
-  u32 gidvid;
-  u32 il_pos;
-
-} plain_t;
-
-typedef struct
-{
-  u32 i[16];
+  u32 i[64];
 
   u32 pw_len;
 
@@ -1516,26 +1506,36 @@ typedef struct
 
 } pw_t;
 
-typedef struct
+typedef struct bf
 {
-  u32 i;
+  u32  i;
 
 } bf_t;
 
-typedef struct
+typedef struct comb
 {
-  u32 i[8];
+  u32  i[64];
 
-  u32 pw_len;
+  u32  pw_len;
 
 } comb_t;
 
-typedef struct
+typedef struct bs_word
 {
-  u32 b[32];
+  u32  b[32];
 
 } bs_word_t;
 
+typedef struct
+{
+  u32 salt_pos;
+  u32 digest_pos;
+  u32 hash_pos;
+  u32 gidvid;
+  u32 il_pos;
+
+} plain_t;
+
 typedef struct
 {
   #ifndef SCRYPT_TMP_ELEM
diff --git a/include/types.h b/include/types.h
index 755961cee..c67712ee7 100644
--- a/include/types.h
+++ b/include/types.h
@@ -673,8 +673,8 @@ typedef enum user_options_map
 
 typedef struct salt
 {
-  u32 salt_buf[16];
-  u32 salt_buf_pc[16];
+  u32 salt_buf[64];
+  u32 salt_buf_pc[64];
 
   u32 salt_len;
   u32 salt_len_pc;
@@ -838,7 +838,7 @@ typedef struct hashconfig hashconfig_t;
 
 typedef struct pw
 {
-  u32 i[16];
+  u32 i[64];
 
   u32 pw_len;
 
@@ -854,20 +854,20 @@ typedef struct bf
 
 } bf_t;
 
+typedef struct comb
+{
+  u32  i[64];
+
+  u32  pw_len;
+
+} comb_t;
+
 typedef struct bs_word
 {
   u32  b[32];
 
 } bs_word_t;
 
-typedef struct comb
-{
-  u32  i[8];
-
-  u32  pw_len;
-
-} comb_t;
-
 typedef struct cpt
 {
   u32    cracked;
@@ -885,18 +885,6 @@ typedef struct plain
 
 } plain_t;
 
-typedef struct wordl
-{
-  u32  word_buf[16];
-
-} wordl_t;
-
-typedef struct wordr
-{
-  u32  word_buf[1];
-
-} wordr_t;
-
 #include "ext_OpenCL.h"
 
 typedef struct hc_device_param
diff --git a/src/convert.c b/src/convert.c
index 347ca4e91..3c2f591ff 100644
--- a/src/convert.c
+++ b/src/convert.c
@@ -157,7 +157,7 @@ bool need_hexify (const u8 *buf, const int len, const char separator, bool alway
 
 void exec_hexify (const u8 *buf, const int len, u8 *out)
 {
-  const int max_len = (len >= 31) ? 31 : len;
+  const int max_len = (len >= PW_MAX) ? PW_MAX : len;
 
   for (int i = max_len - 1, j = i * 2; i >= 0; i -= 1, j -= 2)
   {
diff --git a/src/hashes.c b/src/hashes.c
index 3f7ef6f47..a9130c235 100644
--- a/src/hashes.c
+++ b/src/hashes.c
@@ -248,7 +248,7 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
   // plain
 
-  u32 plain_buf[16] = { 0 };
+  u32 plain_buf[64] = { 0 };
 
   u8 *plain_ptr = (u8 *) plain_buf;
   int plain_len = 0;
diff --git a/src/interface.c b/src/interface.c
index a31b5f672..ee2744092 100644
--- a/src/interface.c
+++ b/src/interface.c
@@ -2414,7 +2414,7 @@ static void drupal7_encode (u8 digest[64], u8 buf[43])
 
 static u32 parse_and_store_salt (u8 *out, u8 *in, u32 salt_len, MAYBE_UNUSED const hashconfig_t *hashconfig)
 {
-  u32 tmp_u32[64] = { 0 };
+  u32 tmp_u32[(64 * 2) + 1] = { 0 };
 
   u8 *tmp = (u8 *) tmp_u32;
 
@@ -2456,18 +2456,15 @@ static u32 parse_and_store_salt (u8 *out, u8 *in, u32 salt_len, MAYBE_UNUSED con
 
   if (hashconfig->opts_type & OPTS_TYPE_ST_UTF16LE)
   {
-    if (salt_len < 20)
+    if (salt_len <= 128)
     {
-      tmp_u32[9] = ((tmp_u32[4] >> 8) & 0x00FF0000) | ((tmp_u32[4] >> 16) & 0x000000FF);
-      tmp_u32[8] = ((tmp_u32[4] << 8) & 0x00FF0000) | ((tmp_u32[4] >>  0) & 0x000000FF);
-      tmp_u32[7] = ((tmp_u32[3] >> 8) & 0x00FF0000) | ((tmp_u32[3] >> 16) & 0x000000FF);
-      tmp_u32[6] = ((tmp_u32[3] << 8) & 0x00FF0000) | ((tmp_u32[3] >>  0) & 0x000000FF);
-      tmp_u32[5] = ((tmp_u32[2] >> 8) & 0x00FF0000) | ((tmp_u32[2] >> 16) & 0x000000FF);
-      tmp_u32[4] = ((tmp_u32[2] << 8) & 0x00FF0000) | ((tmp_u32[2] >>  0) & 0x000000FF);
-      tmp_u32[3] = ((tmp_u32[1] >> 8) & 0x00FF0000) | ((tmp_u32[1] >> 16) & 0x000000FF);
-      tmp_u32[2] = ((tmp_u32[1] << 8) & 0x00FF0000) | ((tmp_u32[1] >>  0) & 0x000000FF);
-      tmp_u32[1] = ((tmp_u32[0] >> 8) & 0x00FF0000) | ((tmp_u32[0] >> 16) & 0x000000FF);
-      tmp_u32[0] = ((tmp_u32[0] << 8) & 0x00FF0000) | ((tmp_u32[0] >>  0) & 0x000000FF);
+      for (int i = 64 - 1; i >= 1; i -= 2)
+      {
+        const u32 v = tmp_u32[i / 2];
+
+        tmp_u32[i - 0] = ((v >> 8) & 0x00FF0000) | ((v >> 16) & 0x000000FF);
+        tmp_u32[i - 1] = ((v << 8) & 0x00FF0000) | ((v >>  0) & 0x000000FF);
+      }
 
       salt_len = salt_len * 2;
     }
diff --git a/src/opencl.c b/src/opencl.c
index 7700d23fc..ac417fe54 100644
--- a/src/opencl.c
+++ b/src/opencl.c
@@ -1152,6 +1152,10 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
 
     if (run_init == true)
     {
+      CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->command_queue, device_param->d_pws_buf, device_param->d_pws_amp_buf, 0, 0, pws_cnt * sizeof (pw_t), 0, NULL, NULL);
+
+      if (CL_rc == -1) return -1;
+
       CL_rc = run_kernel_amp (hashcat_ctx, device_param, pws_cnt);
 
       if (CL_rc == -1) return -1;
diff --git a/src/outfile.c b/src/outfile.c
index f305e2b76..94843d676 100644
--- a/src/outfile.c
+++ b/src/outfile.c
@@ -42,7 +42,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
     if (rc == -1) return -1;
 
-    for (int i = 0; i < 16; i++)
+    for (int i = 0; i < 64; i++)
     {
       plain_buf[i] = pw.i[i];
     }
@@ -63,7 +63,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
     if (rc == -1) return -1;
 
-    for (int i = 0; i < 16; i++)
+    for (int i = 0; i < 64; i++)
     {
       plain_buf[i] = pw.i[i];
     }
@@ -123,7 +123,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
     if (rc == -1) return -1;
 
-    for (int i = 0; i < 16; i++)
+    for (int i = 0; i < 64; i++)
     {
       plain_buf[i] = pw.i[i];
     }
@@ -149,7 +149,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
 
     if (rc == -1) return -1;
 
-    for (int i = 0; i < 16; i++)
+    for (int i = 0; i < 64; i++)
     {
       plain_buf[i] = pw.i[i];
     }
diff --git a/src/status.c b/src/status.c
index 3070b2a6d..7b1b793bd 100644
--- a/src/status.c
+++ b/src/status.c
@@ -751,8 +751,8 @@ char *status_get_guess_candidates_dev (const hashcat_ctx_t *hashcat_ctx, const i
   plain_t plain1 = { 0, 0, 0, outerloop_first, innerloop_first };
   plain_t plain2 = { 0, 0, 0, outerloop_last,  innerloop_last  };
 
-  u32 plain_buf1[40] = { 0 };
-  u32 plain_buf2[40] = { 0 };
+  u32 plain_buf1[(64 * 2) + 2] = { 0 };
+  u32 plain_buf2[(64 * 2) + 2] = { 0 };
 
   u8 *plain_ptr1 = (u8 *) plain_buf1;
   u8 *plain_ptr2 = (u8 *) plain_buf2;
diff --git a/src/stdout.c b/src/stdout.c
index 5f64c533e..b24b00878 100644
--- a/src/stdout.c
+++ b/src/stdout.c
@@ -88,7 +88,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
 
   out.len = 0;
 
-  u32 plain_buf[16] = { 0 };
+  u32 plain_buf[64] = { 0 };
 
   u8 *plain_ptr = (u8 *) plain_buf;
 
@@ -115,7 +115,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
 
       for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
       {
-        for (int i = 0; i < 8; i++)
+        for (int i = 0; i < 64; i++)
         {
           plain_buf[i] = pw.i[i];
         }
@@ -147,7 +147,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
 
       for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
       {
-        for (int i = 0; i < 8; i++)
+        for (int i = 0; i < 64; i++)
         {
           plain_buf[i] = pw.i[i];
         }
@@ -217,7 +217,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
 
       for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
       {
-        for (int i = 0; i < 8; i++)
+        for (int i = 0; i < 64; i++)
         {
           plain_buf[i] = pw.i[i];
         }
@@ -254,7 +254,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
 
       for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
       {
-        for (int i = 0; i < 8; i++)
+        for (int i = 0; i < 64; i++)
         {
           plain_buf[i] = pw.i[i];
         }