diff --git a/OpenCL/m05000_a0-optimized.cl b/OpenCL/m05000_a0-optimized.cl
index d91901727..dbca7398d 100644
--- a/OpenCL/m05000_a0-optimized.cl
+++ b/OpenCL/m05000_a0-optimized.cl
@@ -98,28 +98,18 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
    * constants
    */
 
-  const u8 keccakf_rotc[24] =
+  const int keccakf_rotc[24] =
   {
      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
   };
 
-  const u8 keccakf_piln[24] =
+  const int keccakf_piln[24] =
   {
     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
   };
 
-  /**
-   * 0x80 keccak, very special
-   */
-
-  const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
-
-  const u32 rsiz = 200 - (2 * mdlen);
-
-  const u32 add80w = (rsiz - 1) / 8;
-
   /**
    * loop
    */
@@ -157,7 +147,7 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
     st[13] = 0;
     st[14] = 0;
     st[15] = 0;
-    st[16] = 0;
+    st[16] = 0x8000000000000000;
     st[17] = 0;
     st[18] = 0;
     st[19] = 0;
@@ -167,8 +157,6 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
     st[23] = 0;
     st[24] = 0;
 
-    st[add80w] |= 0x8000000000000000;
-
     int round;
 
     for (round = 0; round < KECCAK_ROUNDS; round++)
@@ -282,28 +270,18 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
    * constants
    */
 
-  const u8 keccakf_rotc[24] =
+  const int keccakf_rotc[24] =
   {
      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
   };
 
-  const u8 keccakf_piln[24] =
+  const int keccakf_piln[24] =
   {
     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
   };
 
-  /**
-   * 0x80 keccak, very special
-   */
-
-  const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
-
-  const u32 rsiz = 200 - (2 * mdlen);
-
-  const u32 add80w = (rsiz - 1) / 8;
-
   /**
    * digest
    */
@@ -353,7 +331,7 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     st[13] = 0;
     st[14] = 0;
     st[15] = 0;
-    st[16] = 0;
+    st[16] = 0x8000000000000000;
     st[17] = 0;
     st[18] = 0;
     st[19] = 0;
@@ -363,8 +341,6 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     st[23] = 0;
     st[24] = 0;
 
-    st[add80w] |= 0x8000000000000000;
-
     int round;
 
     for (round = 0; round < KECCAK_ROUNDS; round++)
diff --git a/OpenCL/m05000_a1-optimized.cl b/OpenCL/m05000_a1-optimized.cl
index 08a39ee0b..0505870a2 100644
--- a/OpenCL/m05000_a1-optimized.cl
+++ b/OpenCL/m05000_a1-optimized.cl
@@ -96,28 +96,18 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
    * constants
    */
 
-  const u8 keccakf_rotc[24] =
+  const int keccakf_rotc[24] =
   {
      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
   };
 
-  const u8 keccakf_piln[24] =
+  const int keccakf_piln[24] =
   {
     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
   };
 
-  /**
-   * 0x80 keccak, very special
-   */
-
-  const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
-
-  const u32 rsiz = 200 - (2 * mdlen);
-
-  const u32 add80w = (rsiz - 1) / 8;
-
   /**
    * loop
    */
@@ -213,7 +203,7 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
     st[13] = 0;
     st[14] = 0;
     st[15] = 0;
-    st[16] = 0;
+    st[16] = 0x8000000000000000;
     st[17] = 0;
     st[18] = 0;
     st[19] = 0;
@@ -223,8 +213,6 @@ __kernel void m05000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
     st[23] = 0;
     st[24] = 0;
 
-    st[add80w] |= 0x8000000000000000;
-
     int round;
 
     for (round = 0; round < KECCAK_ROUNDS; round++)
@@ -338,28 +326,18 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
    * constants
    */
 
-  const u8 keccakf_rotc[24] =
+  const int keccakf_rotc[24] =
   {
      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
   };
 
-  const u8 keccakf_piln[24] =
+  const int keccakf_piln[24] =
   {
     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
   };
 
-  /**
-   * 0x80 keccak, very special
-   */
-
-  const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
-
-  const u32 rsiz = 200 - (2 * mdlen);
-
-  const u32 add80w = (rsiz - 1) / 8;
-
   /**
    * digest
    */
@@ -467,7 +445,7 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     st[13] = 0;
     st[14] = 0;
     st[15] = 0;
-    st[16] = 0;
+    st[16] = 0x8000000000000000;
     st[17] = 0;
     st[18] = 0;
     st[19] = 0;
@@ -477,8 +455,6 @@ __kernel void m05000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     st[23] = 0;
     st[24] = 0;
 
-    st[add80w] |= 0x8000000000000000;
-
     int round;
 
     for (round = 0; round < KECCAK_ROUNDS; round++)
diff --git a/OpenCL/m05000_a3-optimized.cl b/OpenCL/m05000_a3-optimized.cl
index 1fb2fa36f..d0598bd80 100644
--- a/OpenCL/m05000_a3-optimized.cl
+++ b/OpenCL/m05000_a3-optimized.cl
@@ -75,28 +75,18 @@ void m05000m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
    * constants
    */
 
-  const u8 keccakf_rotc[24] =
+  const int keccakf_rotc[24] =
   {
      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
   };
 
-  const u8 keccakf_piln[24] =
+  const int keccakf_piln[24] =
   {
     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
   };
 
-  /**
-   * 0x80 keccak, very special
-   */
-
-  const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
-
-  const u32 rsiz = 200 - (2 * mdlen);
-
-  const u32 add80w = (rsiz - 1) / 8;
-
   /**
    * loop
    */
@@ -153,7 +143,7 @@ void m05000m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
     st[13] = 0;
     st[14] = 0;
     st[15] = 0;
-    st[16] = 0;
+    st[16] = 0x8000000000000000;
     st[17] = 0;
     st[18] = 0;
     st[19] = 0;
@@ -163,8 +153,6 @@ void m05000m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
     st[23] = 0;
     st[24] = 0;
 
-    st[add80w] |= 0x8000000000000000;
-
     int round;
 
     for (round = 0; round < KECCAK_ROUNDS; round++)
@@ -238,28 +226,18 @@ void m05000s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
    * constants
    */
 
-  const u8 keccakf_rotc[24] =
+  const int keccakf_rotc[24] =
   {
      1,  3,  6, 10, 15, 21, 28, 36, 45, 55,  2, 14,
     27, 41, 56,  8, 25, 43, 62, 18, 39, 61, 20, 44
   };
 
-  const u8 keccakf_piln[24] =
+  const int keccakf_piln[24] =
   {
     10,  7, 11, 17, 18,  3,  5, 16,  8, 21, 24,  4,
     15, 23, 19, 13, 12,  2, 20, 14, 22,  9,  6,  1
   };
 
-  /**
-   * 0x80 keccak, very special
-   */
-
-  const u32 mdlen = salt_bufs[salt_pos].keccak_mdlen;
-
-  const u32 rsiz = 200 - (2 * mdlen);
-
-  const u32 add80w = (rsiz - 1) / 8;
-
   /**
    * digest
    */
@@ -328,7 +306,7 @@ void m05000s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
     st[13] = 0;
     st[14] = 0;
     st[15] = 0;
-    st[16] = 0;
+    st[16] = 0x8000000000000000;
     st[17] = 0;
     st[18] = 0;
     st[19] = 0;
@@ -338,8 +316,6 @@ void m05000s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
     st[23] = 0;
     st[24] = 0;
 
-    st[add80w] |= 0x8000000000000000;
-
     int round;
 
     for (round = 0; round < KECCAK_ROUNDS; round++)