diff --git a/OpenCL/m00050_a1.cl b/OpenCL/m00050_a1.cl
index 7d75b6e4f..de0739fb6 100644
--- a/OpenCL/m00050_a1.cl
+++ b/OpenCL/m00050_a1.cl
@@ -82,13 +82,13 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
       c[i] |= w[i];
     }
 
-    md5_hmac_ctx_vector_t ctx;
+    md5_hmac_ctx_t ctx;
 
-    md5_hmac_init_vector (&ctx, c, pw_len + comb_len);
+    md5_hmac_init (&ctx, c, pw_len + comb_len);
 
-    md5_hmac_update_vector (&ctx, s, salt_len);
+    md5_hmac_update (&ctx, s, salt_len);
 
-    md5_hmac_final_vector (&ctx);
+    md5_hmac_final (&ctx);
 
     const u32 r0 = ctx.opad.h[DGST_R0];
     const u32 r1 = ctx.opad.h[DGST_R1];
@@ -180,13 +180,13 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
       c[i] |= w[i];
     }
 
-    md5_hmac_ctx_vector_t ctx;
+    md5_hmac_ctx_t ctx;
 
-    md5_hmac_init_vector (&ctx, c, pw_len + comb_len);
+    md5_hmac_init (&ctx, c, pw_len + comb_len);
 
-    md5_hmac_update_vector (&ctx, s, salt_len);
+    md5_hmac_update (&ctx, s, salt_len);
 
-    md5_hmac_final_vector (&ctx);
+    md5_hmac_final (&ctx);
 
     const u32 r0 = ctx.opad.h[DGST_R0];
     const u32 r1 = ctx.opad.h[DGST_R1];
diff --git a/OpenCL/m00150_a1.cl b/OpenCL/m00150_a1.cl
index 15ab612bf..96153a25c 100644
--- a/OpenCL/m00150_a1.cl
+++ b/OpenCL/m00150_a1.cl
@@ -82,13 +82,13 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
       c[i] |= w[i];
     }
 
-    sha1_hmac_ctx_vector_t ctx;
+    sha1_hmac_ctx_t ctx;
 
-    sha1_hmac_init_vector (&ctx, c, pw_len + comb_len);
+    sha1_hmac_init (&ctx, c, pw_len + comb_len);
 
-    sha1_hmac_update_vector (&ctx, s, salt_len);
+    sha1_hmac_update (&ctx, s, salt_len);
 
-    sha1_hmac_final_vector (&ctx);
+    sha1_hmac_final (&ctx);
 
     const u32 r0 = ctx.opad.h[DGST_R0];
     const u32 r1 = ctx.opad.h[DGST_R1];
@@ -180,13 +180,13 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
       c[i] |= w[i];
     }
 
-    sha1_hmac_ctx_vector_t ctx;
+    sha1_hmac_ctx_t ctx;
 
-    sha1_hmac_init_vector (&ctx, c, pw_len + comb_len);
+    sha1_hmac_init (&ctx, c, pw_len + comb_len);
 
-    sha1_hmac_update_vector (&ctx, s, salt_len);
+    sha1_hmac_update (&ctx, s, salt_len);
 
-    sha1_hmac_final_vector (&ctx);
+    sha1_hmac_final (&ctx);
 
     const u32 r0 = ctx.opad.h[DGST_R0];
     const u32 r1 = ctx.opad.h[DGST_R1];
diff --git a/OpenCL/m01100_a0.cl b/OpenCL/m01100_a0.cl
index e257ed697..085e29e18 100644
--- a/OpenCL/m01100_a0.cl
+++ b/OpenCL/m01100_a0.cl
@@ -47,7 +47,7 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -140,7 +140,7 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m01100_a1.cl b/OpenCL/m01100_a1.cl
index 56fa3bfda..6c7ef2d9e 100644
--- a/OpenCL/m01100_a1.cl
+++ b/OpenCL/m01100_a1.cl
@@ -32,7 +32,7 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -114,7 +114,7 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m01450_a1.cl b/OpenCL/m01450_a1.cl
index c482c7626..25e81d266 100644
--- a/OpenCL/m01450_a1.cl
+++ b/OpenCL/m01450_a1.cl
@@ -82,13 +82,13 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
       c[i] |= w[i];
     }
 
-    sha256_hmac_ctx_vector_t ctx;
+    sha256_hmac_ctx_t ctx;
 
-    sha256_hmac_init_vector (&ctx, c, pw_len + comb_len);
+    sha256_hmac_init (&ctx, c, pw_len + comb_len);
 
-    sha256_hmac_update_vector (&ctx, s, salt_len);
+    sha256_hmac_update (&ctx, s, salt_len);
 
-    sha256_hmac_final_vector (&ctx);
+    sha256_hmac_final (&ctx);
 
     const u32 r0 = ctx.opad.h[DGST_R0];
     const u32 r1 = ctx.opad.h[DGST_R1];
@@ -180,13 +180,13 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
       c[i] |= w[i];
     }
 
-    sha256_hmac_ctx_vector_t ctx;
+    sha256_hmac_ctx_t ctx;
 
-    sha256_hmac_init_vector (&ctx, c, pw_len + comb_len);
+    sha256_hmac_init (&ctx, c, pw_len + comb_len);
 
-    sha256_hmac_update_vector (&ctx, s, salt_len);
+    sha256_hmac_update (&ctx, s, salt_len);
 
-    sha256_hmac_final_vector (&ctx);
+    sha256_hmac_final (&ctx);
 
     const u32 r0 = ctx.opad.h[DGST_R0];
     const u32 r1 = ctx.opad.h[DGST_R1];
diff --git a/OpenCL/m01500_a0.cl b/OpenCL/m01500_a0-optimized.cl
similarity index 100%
rename from OpenCL/m01500_a0.cl
rename to OpenCL/m01500_a0-optimized.cl
diff --git a/OpenCL/m01500_a1.cl b/OpenCL/m01500_a1-optimized.cl
similarity index 100%
rename from OpenCL/m01500_a1.cl
rename to OpenCL/m01500_a1-optimized.cl
diff --git a/OpenCL/m01500_a3.cl b/OpenCL/m01500_a3-optimized.cl
similarity index 100%
rename from OpenCL/m01500_a3.cl
rename to OpenCL/m01500_a3-optimized.cl
diff --git a/OpenCL/m01750_a1.cl b/OpenCL/m01750_a1.cl
index 27be72d5e..d6d9da0fd 100644
--- a/OpenCL/m01750_a1.cl
+++ b/OpenCL/m01750_a1.cl
@@ -82,13 +82,13 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
       c[i] |= w[i];
     }
 
-    sha512_hmac_ctx_vector_t ctx;
+    sha512_hmac_ctx_t ctx;
 
-    sha512_hmac_init_vector (&ctx, c, pw_len + comb_len);
+    sha512_hmac_init (&ctx, c, pw_len + comb_len);
 
-    sha512_hmac_update_vector (&ctx, s, salt_len);
+    sha512_hmac_update (&ctx, s, salt_len);
 
-    sha512_hmac_final_vector (&ctx);
+    sha512_hmac_final (&ctx);
 
     const u32 r0 = l32_from_64_S (ctx.opad.h[7]);
     const u32 r1 = h32_from_64_S (ctx.opad.h[7]);
@@ -180,13 +180,13 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
       c[i] |= w[i];
     }
 
-    sha512_hmac_ctx_vector_t ctx;
+    sha512_hmac_ctx_t ctx;
 
-    sha512_hmac_init_vector (&ctx, c, pw_len + comb_len);
+    sha512_hmac_init (&ctx, c, pw_len + comb_len);
 
-    sha512_hmac_update_vector (&ctx, s, salt_len);
+    sha512_hmac_update (&ctx, s, salt_len);
 
-    sha512_hmac_final_vector (&ctx);
+    sha512_hmac_final (&ctx);
 
     const u32 r0 = l32_from_64_S (ctx.opad.h[7]);
     const u32 r1 = h32_from_64_S (ctx.opad.h[7]);
diff --git a/OpenCL/m02000_a0.cl b/OpenCL/m02000_a0-optimized.cl
similarity index 100%
rename from OpenCL/m02000_a0.cl
rename to OpenCL/m02000_a0-optimized.cl
diff --git a/OpenCL/m02000_a1.cl b/OpenCL/m02000_a1-optimized.cl
similarity index 100%
rename from OpenCL/m02000_a1.cl
rename to OpenCL/m02000_a1-optimized.cl
diff --git a/OpenCL/m02000_a3.cl b/OpenCL/m02000_a3-optimized.cl
similarity index 100%
rename from OpenCL/m02000_a3.cl
rename to OpenCL/m02000_a3-optimized.cl
diff --git a/OpenCL/m02610_a0.cl b/OpenCL/m02610_a0.cl
index 633ea6f13..a09b5b41e 100644
--- a/OpenCL/m02610_a0.cl
+++ b/OpenCL/m02610_a0.cl
@@ -77,7 +77,7 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -205,7 +205,7 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m02610_a1.cl b/OpenCL/m02610_a1.cl
index 3818eb960..3fbdd72f9 100644
--- a/OpenCL/m02610_a1.cl
+++ b/OpenCL/m02610_a1.cl
@@ -62,7 +62,7 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -179,7 +179,7 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m02810_a0.cl b/OpenCL/m02810_a0.cl
index 02dfa50fd..502105952 100644
--- a/OpenCL/m02810_a0.cl
+++ b/OpenCL/m02810_a0.cl
@@ -77,7 +77,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -228,7 +228,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m02810_a1.cl b/OpenCL/m02810_a1.cl
index 28b4e271d..df010bf8d 100644
--- a/OpenCL/m02810_a1.cl
+++ b/OpenCL/m02810_a1.cl
@@ -62,7 +62,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -202,7 +202,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m02810_a3.cl b/OpenCL/m02810_a3.cl
index 0a8eff53b..48fa8911f 100644
--- a/OpenCL/m02810_a3.cl
+++ b/OpenCL/m02810_a3.cl
@@ -25,7 +25,7 @@
 #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7], l_bin2asc[(i).s8], l_bin2asc[(i).s9], l_bin2asc[(i).sa], l_bin2asc[(i).sb], l_bin2asc[(i).sc], l_bin2asc[(i).sd], l_bin2asc[(i).se], l_bin2asc[(i).sf])
 #endif
 
-__kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32 *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
   /**
    * modifier
@@ -62,7 +62,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 pw_lenv = ceil ((float) pw_len / 4);
 
-  u32x w[64] = { 0 };
+  u32 w[64] = { 0 };
 
   for (int idx = 0; idx < pw_lenv; idx++)
   {
@@ -75,7 +75,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -88,13 +88,13 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
    * loop
    */
 
-  u32x w0l = w[0];
+  u32 w0l = w[0];
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
   {
-    const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
+    const u32 w0r = words_buf_r[il_pos / VECT_SIZE];
 
-    const u32x w0lr = w0l | w0r;
+    const u32 w0lr = w0l | w0r;
 
     w[0] = w0lr;
 
@@ -106,10 +106,10 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     md5_final_vector (&ctx0);
 
-    const u32x a = ctx0.h[0];
-    const u32x b = ctx0.h[1];
-    const u32x c = ctx0.h[2];
-    const u32x d = ctx0.h[3];
+    const u32 a = ctx0.h[0];
+    const u32 b = ctx0.h[1];
+    const u32 c = ctx0.h[2];
+    const u32 d = ctx0.h[3];
 
     md5_ctx_vector_t ctx;
 
@@ -161,16 +161,16 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     md5_transform_vector (ctx.w0, ctx.w1, ctx.w2, ctx.w3, ctx.h);
 
-    const u32x r0 = ctx.h[DGST_R0];
-    const u32x r1 = ctx.h[DGST_R1];
-    const u32x r2 = ctx.h[DGST_R2];
-    const u32x r3 = ctx.h[DGST_R3];
+    const u32 r0 = ctx.h[DGST_R0];
+    const u32 r1 = ctx.h[DGST_R1];
+    const u32 r2 = ctx.h[DGST_R2];
+    const u32 r3 = ctx.h[DGST_R3];
 
     COMPARE_M_SIMD (r0, r1, r2, r3);
   }
 }
 
-__kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32 *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
   /**
    * modifier
@@ -219,7 +219,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 pw_lenv = ceil ((float) pw_len / 4);
 
-  u32x w[64] = { 0 };
+  u32 w[64] = { 0 };
 
   for (int idx = 0; idx < pw_lenv; idx++)
   {
@@ -232,7 +232,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -245,13 +245,13 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
    * loop
    */
 
-  u32x w0l = w[0];
+  u32 w0l = w[0];
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
   {
-    const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
+    const u32 w0r = words_buf_r[il_pos / VECT_SIZE];
 
-    const u32x w0lr = w0l | w0r;
+    const u32 w0lr = w0l | w0r;
 
     w[0] = w0lr;
 
@@ -263,10 +263,10 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     md5_final_vector (&ctx0);
 
-    const u32x a = ctx0.h[0];
-    const u32x b = ctx0.h[1];
-    const u32x c = ctx0.h[2];
-    const u32x d = ctx0.h[3];
+    const u32 a = ctx0.h[0];
+    const u32 b = ctx0.h[1];
+    const u32 c = ctx0.h[2];
+    const u32 d = ctx0.h[3];
 
     md5_ctx_vector_t ctx;
 
@@ -318,10 +318,10 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     md5_transform_vector (ctx.w0, ctx.w1, ctx.w2, ctx.w3, ctx.h);
 
-    const u32x r0 = ctx.h[DGST_R0];
-    const u32x r1 = ctx.h[DGST_R1];
-    const u32x r2 = ctx.h[DGST_R2];
-    const u32x r3 = ctx.h[DGST_R3];
+    const u32 r0 = ctx.h[DGST_R0];
+    const u32 r1 = ctx.h[DGST_R1];
+    const u32 r2 = ctx.h[DGST_R2];
+    const u32 r3 = ctx.h[DGST_R3];
 
     COMPARE_S_SIMD (r0, r1, r2, r3);
   }
diff --git a/OpenCL/m03000_a0.cl b/OpenCL/m03000_a0-optimized.cl
similarity index 100%
rename from OpenCL/m03000_a0.cl
rename to OpenCL/m03000_a0-optimized.cl
diff --git a/OpenCL/m03000_a1.cl b/OpenCL/m03000_a1-optimized.cl
similarity index 100%
rename from OpenCL/m03000_a1.cl
rename to OpenCL/m03000_a1-optimized.cl
diff --git a/OpenCL/m03000_a3.cl b/OpenCL/m03000_a3-optimized.cl
similarity index 100%
rename from OpenCL/m03000_a3.cl
rename to OpenCL/m03000_a3-optimized.cl
diff --git a/OpenCL/m03710_a0.cl b/OpenCL/m03710_a0.cl
index 054465f2d..89b894741 100644
--- a/OpenCL/m03710_a0.cl
+++ b/OpenCL/m03710_a0.cl
@@ -77,7 +77,7 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -218,7 +218,7 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m03710_a1.cl b/OpenCL/m03710_a1.cl
index b74fa951f..16e7caa5e 100644
--- a/OpenCL/m03710_a1.cl
+++ b/OpenCL/m03710_a1.cl
@@ -62,7 +62,7 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -192,7 +192,7 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m03910_a0.cl b/OpenCL/m03910_a0.cl
index ab08e1a87..1139e0cb6 100644
--- a/OpenCL/m03910_a0.cl
+++ b/OpenCL/m03910_a0.cl
@@ -77,7 +77,7 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -228,7 +228,7 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m03910_a1.cl b/OpenCL/m03910_a1.cl
index c4186f01e..40c9071dc 100644
--- a/OpenCL/m03910_a1.cl
+++ b/OpenCL/m03910_a1.cl
@@ -62,7 +62,7 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -202,7 +202,7 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[8] = { 0 };
+  u32 s[8] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m04310_a0.cl b/OpenCL/m04310_a0.cl
index 8ff005fa0..31d1af318 100644
--- a/OpenCL/m04310_a0.cl
+++ b/OpenCL/m04310_a0.cl
@@ -77,7 +77,7 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -205,7 +205,7 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m04310_a1.cl b/OpenCL/m04310_a1.cl
index aea272f93..bf33a9155 100644
--- a/OpenCL/m04310_a1.cl
+++ b/OpenCL/m04310_a1.cl
@@ -62,7 +62,7 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
@@ -179,7 +179,7 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 salt_lenv = ceil ((float) salt_len / 4);
 
-  u32x s[64] = { 0 };
+  u32 s[64] = { 0 };
 
   for (int idx = 0; idx < salt_lenv; idx++)
   {
diff --git a/OpenCL/m04800_a0.cl b/OpenCL/m04800_a0.cl
index 02c9903fa..cd7c9cc07 100644
--- a/OpenCL/m04800_a0.cl
+++ b/OpenCL/m04800_a0.cl
@@ -72,7 +72,7 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     md5_update (&ctx, w, pw_len);
 
-    md5_update_vector (&ctx, s, salt_len);
+    md5_update (&ctx, s, salt_len);
 
     md5_final (&ctx);
 
@@ -154,7 +154,7 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     md5_update (&ctx, w, pw_len);
 
-    md5_update_vector (&ctx, s, salt_len);
+    md5_update (&ctx, s, salt_len);
 
     md5_final (&ctx);
 
diff --git a/OpenCL/m04800_a1.cl b/OpenCL/m04800_a1.cl
index d7e968d1c..b213e74c2 100644
--- a/OpenCL/m04800_a1.cl
+++ b/OpenCL/m04800_a1.cl
@@ -57,7 +57,7 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
 
-    md5_update_vector (&ctx, s, salt_len);
+    md5_update (&ctx, s, salt_len);
 
     md5_final (&ctx);
 
@@ -126,7 +126,7 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
 
-    md5_update_vector (&ctx, s, salt_len);
+    md5_update (&ctx, s, salt_len);
 
     md5_final (&ctx);
 
diff --git a/OpenCL/m05400_a0.cl b/OpenCL/m05400_a0.cl
index 34b78379f..97d759c03 100644
--- a/OpenCL/m05400_a0.cl
+++ b/OpenCL/m05400_a0.cl
@@ -89,10 +89,10 @@ __kernel void m05400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     sha1_hmac_final (&ctx);
 
-    const u32x r0 = ctx.opad.h[DGST_R0];
-    const u32x r1 = ctx.opad.h[DGST_R1];
-    const u32x r2 = ctx.opad.h[DGST_R2];
-    const u32x r3 = ctx.opad.h[DGST_R3];
+    const u32 r0 = ctx.opad.h[DGST_R0];
+    const u32 r1 = ctx.opad.h[DGST_R1];
+    const u32 r2 = ctx.opad.h[DGST_R2];
+    const u32 r3 = ctx.opad.h[DGST_R3];
 
     COMPARE_M_SCALAR (r0, r1, r2, r3);
   }
@@ -184,10 +184,10 @@ __kernel void m05400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     sha1_hmac_final (&ctx);
 
-    const u32x r0 = ctx.opad.h[DGST_R0];
-    const u32x r1 = ctx.opad.h[DGST_R1];
-    const u32x r2 = ctx.opad.h[DGST_R2];
-    const u32x r3 = ctx.opad.h[DGST_R3];
+    const u32 r0 = ctx.opad.h[DGST_R0];
+    const u32 r1 = ctx.opad.h[DGST_R1];
+    const u32 r2 = ctx.opad.h[DGST_R2];
+    const u32 r3 = ctx.opad.h[DGST_R3];
 
     COMPARE_S_SCALAR (r0, r1, r2, r3);
   }
diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl
index 3c8a790ae..3780e2e1b 100644
--- a/OpenCL/m05500_a0.cl
+++ b/OpenCL/m05500_a0.cl
@@ -344,18 +344,18 @@ __constant u32a c_skb[8][64] =
 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
 #endif
 
-void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
+void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
 {
-  u32x r = data[0];
-  u32x l = data[1];
+  u32 r = data[0];
+  u32 l = data[1];
 
   #ifdef _unroll
   #pragma unroll
   #endif
   for (u32 i = 0; i < 16; i += 2)
   {
-    u32x u;
-    u32x t;
+    u32 u;
+    u32 t;
 
     u = Kc[i + 0] ^ rotl32 (r, 30u);
     t = Kd[i + 0] ^ rotl32 (r, 26u);
@@ -386,9 +386,9 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
   iv[1] = r;
 }
 
-void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
+void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64])
 {
-  u32x tt;
+  u32 tt;
 
   PERM_OP  (d, c, tt, 4, 0x0f0f0f0f);
   HPERM_OP (c,    tt, 2, 0xcccc0000);
@@ -423,13 +423,13 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
     c = c & 0x0fffffff;
     d = d & 0x0fffffff;
 
-    const u32x c00 = (c >>  0) & 0x0000003f;
-    const u32x c06 = (c >>  6) & 0x00383003;
-    const u32x c07 = (c >>  7) & 0x0000003c;
-    const u32x c13 = (c >> 13) & 0x0000060f;
-    const u32x c20 = (c >> 20) & 0x00000001;
+    const u32 c00 = (c >>  0) & 0x0000003f;
+    const u32 c06 = (c >>  6) & 0x00383003;
+    const u32 c07 = (c >>  7) & 0x0000003c;
+    const u32 c13 = (c >> 13) & 0x0000060f;
+    const u32 c20 = (c >> 20) & 0x00000001;
 
-    u32x s = BOX (((c00 >>  0) & 0xff), 0, s_skb)
+    u32 s = BOX (((c00 >>  0) & 0xff), 0, s_skb)
            | BOX (((c06 >>  0) & 0xff)
                  |((c07 >>  0) & 0xff), 1, s_skb)
            | BOX (((c13 >>  0) & 0xff)
@@ -438,12 +438,12 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
                  |((c13 >>  8) & 0xff)
                  |((c06 >> 16) & 0xff), 3, s_skb);
 
-    const u32x d00 = (d >>  0) & 0x00003c3f;
-    const u32x d07 = (d >>  7) & 0x00003f03;
-    const u32x d21 = (d >> 21) & 0x0000000f;
-    const u32x d22 = (d >> 22) & 0x00000030;
+    const u32 d00 = (d >>  0) & 0x00003c3f;
+    const u32 d07 = (d >>  7) & 0x00003f03;
+    const u32 d21 = (d >> 21) & 0x0000000f;
+    const u32 d22 = (d >> 22) & 0x00000030;
 
-    u32x t = BOX (((d00 >>  0) & 0xff), 4, s_skb)
+    u32 t = BOX (((d00 >>  0) & 0xff), 4, s_skb)
            | BOX (((d07 >>  0) & 0xff)
                  |((d00 >>  8) & 0xff), 5, s_skb)
            | BOX (((d07 >>  8) & 0xff), 6, s_skb)
@@ -455,9 +455,9 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
   }
 }
 
-void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
+void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
 {
-  u32x t[8];
+  u32 t[8];
 
   t[0] = (w0 >>  0) & 0xff;
   t[1] = (w0 >>  8) & 0xff;
@@ -468,7 +468,7 @@ void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
   t[6] = (w1 >> 16) & 0xff;
   t[7] = (w1 >> 24) & 0xff;
 
-  u32x k[8];
+  u32 k[8];
 
   k[0] =               (t[0] >> 0);
   k[1] = (t[0] << 7) | (t[1] >> 1);
diff --git a/OpenCL/m05500_a1.cl b/OpenCL/m05500_a1.cl
index 34bc4711f..be64670a1 100644
--- a/OpenCL/m05500_a1.cl
+++ b/OpenCL/m05500_a1.cl
@@ -341,18 +341,18 @@ __constant u32a c_skb[8][64] =
 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
 #endif
 
-void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
+void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
 {
-  u32x r = data[0];
-  u32x l = data[1];
+  u32 r = data[0];
+  u32 l = data[1];
 
   #ifdef _unroll
   #pragma unroll
   #endif
   for (u32 i = 0; i < 16; i += 2)
   {
-    u32x u;
-    u32x t;
+    u32 u;
+    u32 t;
 
     u = Kc[i + 0] ^ rotl32 (r, 30u);
     t = Kd[i + 0] ^ rotl32 (r, 26u);
@@ -383,9 +383,9 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
   iv[1] = r;
 }
 
-void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
+void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64])
 {
-  u32x tt;
+  u32 tt;
 
   PERM_OP  (d, c, tt, 4, 0x0f0f0f0f);
   HPERM_OP (c,    tt, 2, 0xcccc0000);
@@ -420,13 +420,13 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
     c = c & 0x0fffffff;
     d = d & 0x0fffffff;
 
-    const u32x c00 = (c >>  0) & 0x0000003f;
-    const u32x c06 = (c >>  6) & 0x00383003;
-    const u32x c07 = (c >>  7) & 0x0000003c;
-    const u32x c13 = (c >> 13) & 0x0000060f;
-    const u32x c20 = (c >> 20) & 0x00000001;
+    const u32 c00 = (c >>  0) & 0x0000003f;
+    const u32 c06 = (c >>  6) & 0x00383003;
+    const u32 c07 = (c >>  7) & 0x0000003c;
+    const u32 c13 = (c >> 13) & 0x0000060f;
+    const u32 c20 = (c >> 20) & 0x00000001;
 
-    u32x s = BOX (((c00 >>  0) & 0xff), 0, s_skb)
+    u32 s = BOX (((c00 >>  0) & 0xff), 0, s_skb)
            | BOX (((c06 >>  0) & 0xff)
                  |((c07 >>  0) & 0xff), 1, s_skb)
            | BOX (((c13 >>  0) & 0xff)
@@ -435,12 +435,12 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
                  |((c13 >>  8) & 0xff)
                  |((c06 >> 16) & 0xff), 3, s_skb);
 
-    const u32x d00 = (d >>  0) & 0x00003c3f;
-    const u32x d07 = (d >>  7) & 0x00003f03;
-    const u32x d21 = (d >> 21) & 0x0000000f;
-    const u32x d22 = (d >> 22) & 0x00000030;
+    const u32 d00 = (d >>  0) & 0x00003c3f;
+    const u32 d07 = (d >>  7) & 0x00003f03;
+    const u32 d21 = (d >> 21) & 0x0000000f;
+    const u32 d22 = (d >> 22) & 0x00000030;
 
-    u32x t = BOX (((d00 >>  0) & 0xff), 4, s_skb)
+    u32 t = BOX (((d00 >>  0) & 0xff), 4, s_skb)
            | BOX (((d07 >>  0) & 0xff)
                  |((d00 >>  8) & 0xff), 5, s_skb)
            | BOX (((d07 >>  8) & 0xff), 6, s_skb)
@@ -452,9 +452,9 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
   }
 }
 
-void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
+void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
 {
-  u32x t[8];
+  u32 t[8];
 
   t[0] = (w0 >>  0) & 0xff;
   t[1] = (w0 >>  8) & 0xff;
@@ -465,7 +465,7 @@ void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
   t[6] = (w1 >> 16) & 0xff;
   t[7] = (w1 >> 24) & 0xff;
 
-  u32x k[8];
+  u32 k[8];
 
   k[0] =               (t[0] >> 0);
   k[1] = (t[0] << 7) | (t[1] >> 1);
diff --git a/OpenCL/m10800_a0.cl b/OpenCL/m10800_a0.cl
index 1ae43d030..5e0780118 100644
--- a/OpenCL/m10800_a0.cl
+++ b/OpenCL/m10800_a0.cl
@@ -59,10 +59,10 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     sha384_final (&ctx);
 
-    const u32x r0 = l32_from_64 (ctx.h[3]);
-    const u32x r1 = h32_from_64 (ctx.h[3]);
-    const u32x r2 = l32_from_64 (ctx.h[2]);
-    const u32x r3 = h32_from_64 (ctx.h[2]);
+    const u32 r0 = l32_from_64 (ctx.h[3]);
+    const u32 r1 = h32_from_64 (ctx.h[3]);
+    const u32 r2 = l32_from_64 (ctx.h[2]);
+    const u32 r3 = h32_from_64 (ctx.h[2]);
 
     COMPARE_M_SCALAR (r0, r1, r2, r3);
   }
@@ -124,10 +124,10 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     sha384_final (&ctx);
 
-    const u32x r0 = l32_from_64 (ctx.h[3]);
-    const u32x r1 = h32_from_64 (ctx.h[3]);
-    const u32x r2 = l32_from_64 (ctx.h[2]);
-    const u32x r3 = h32_from_64 (ctx.h[2]);
+    const u32 r0 = l32_from_64 (ctx.h[3]);
+    const u32 r1 = h32_from_64 (ctx.h[3]);
+    const u32 r2 = l32_from_64 (ctx.h[2]);
+    const u32 r3 = h32_from_64 (ctx.h[2]);
 
     COMPARE_S_SCALAR (r0, r1, r2, r3);
   }
diff --git a/OpenCL/m10800_a1.cl b/OpenCL/m10800_a1.cl
index 9f1eded3a..a666a6019 100644
--- a/OpenCL/m10800_a1.cl
+++ b/OpenCL/m10800_a1.cl
@@ -46,10 +46,10 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     sha384_final (&ctx);
 
-    const u32x r0 = l32_from_64 (ctx.h[3]);
-    const u32x r1 = h32_from_64 (ctx.h[3]);
-    const u32x r2 = l32_from_64 (ctx.h[2]);
-    const u32x r3 = h32_from_64 (ctx.h[2]);
+    const u32 r0 = l32_from_64 (ctx.h[3]);
+    const u32 r1 = h32_from_64 (ctx.h[3]);
+    const u32 r2 = l32_from_64 (ctx.h[2]);
+    const u32 r3 = h32_from_64 (ctx.h[2]);
 
     COMPARE_M_SCALAR (r0, r1, r2, r3);
   }
@@ -100,10 +100,10 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
     sha384_final (&ctx);
 
-    const u32x r0 = l32_from_64 (ctx.h[3]);
-    const u32x r1 = h32_from_64 (ctx.h[3]);
-    const u32x r2 = l32_from_64 (ctx.h[2]);
-    const u32x r3 = h32_from_64 (ctx.h[2]);
+    const u32 r0 = l32_from_64 (ctx.h[3]);
+    const u32 r1 = h32_from_64 (ctx.h[3]);
+    const u32 r2 = l32_from_64 (ctx.h[2]);
+    const u32 r3 = h32_from_64 (ctx.h[2]);
 
     COMPARE_S_SCALAR (r0, r1, r2, r3);
   }