From a0fba5fb11ae3e26409d4e2357b3e9141b83163e Mon Sep 17 00:00:00 2001
From: jsteube <jens.steube@gmail.com>
Date: Fri, 22 Feb 2019 12:33:16 +0100
Subject: [PATCH] Improve -m 18700 cracking speed

---
 OpenCL/m18700_a0-optimized.cl |  48 ++++----
 OpenCL/m18700_a0-pure.cl      |  48 ++++----
 OpenCL/m18700_a1-optimized.cl | 212 +++++-----------------------------
 OpenCL/m18700_a1-pure.cl      |  72 +++++-------
 OpenCL/m18700_a3-optimized.cl |  90 ++++++++-------
 OpenCL/m18700_a3-pure.cl      |  74 +++++++-----
 6 files changed, 200 insertions(+), 344 deletions(-)

diff --git a/OpenCL/m18700_a0-optimized.cl b/OpenCL/m18700_a0-optimized.cl
index 04ea3c589..7cacbb6c7 100644
--- a/OpenCL/m18700_a0-optimized.cl
+++ b/OpenCL/m18700_a0-optimized.cl
@@ -14,6 +14,30 @@
 #include "inc_rp_optimized.cl"
 #include "inc_simd.cl"
 
+DECLSPEC u32 hashCode (const u32 init, const u32 *w, const u32 pw_len)
+{
+  u32 hash = init;
+
+  for (u32 i = 0; i < pw_len; i += 4)
+  {
+    u32 tmp = w[i / 4];
+
+    const u32 left = pw_len - i;
+
+    const u32 c = (left > 4) ? 4 : left;
+
+    switch (c)
+    {
+      case 4: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 3: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 2: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 1: hash *= 31; hash += tmp & 0xff;
+    }
+  }
+
+  return hash;
+}
+
 __kernel void m18700_m04 (KERN_ATTR_RULES ())
 {
   /**
@@ -54,17 +78,7 @@ __kernel void m18700_m04 (KERN_ATTR_RULES ())
 
     const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w + 0, w + 4);
 
-    u32x hash = 0;
-
-    for (u32 i = 0; i < pw_len; i++)
-    {
-      const u32 c32 = w[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    u32x hash = hashCode (0, w, out_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
@@ -135,17 +149,7 @@ __kernel void m18700_s04 (KERN_ATTR_RULES ())
 
     const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w + 0, w + 4);
 
-    u32x hash = 0;
-
-    for (u32 i = 0; i < pw_len; i++)
-    {
-      const u32 c32 = w[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    u32x hash = hashCode (0, w, out_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
diff --git a/OpenCL/m18700_a0-pure.cl b/OpenCL/m18700_a0-pure.cl
index 3e646b521..ad7b6e647 100644
--- a/OpenCL/m18700_a0-pure.cl
+++ b/OpenCL/m18700_a0-pure.cl
@@ -15,6 +15,30 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md5.cl"
 
+DECLSPEC u32 hashCode (const u32 init, const u32 *w, const u32 pw_len)
+{
+  u32 hash = init;
+
+  for (u32 i = 0; i < pw_len; i += 4)
+  {
+    u32 tmp = w[i / 4];
+
+    const u32 left = pw_len - i;
+
+    const u32 c = (left > 4) ? 4 : left;
+
+    switch (c)
+    {
+      case 4: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 3: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 2: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 1: hash *= 31; hash += tmp & 0xff;
+    }
+  }
+
+  return hash;
+}
+
 __kernel void m18700_mxx (KERN_ATTR_RULES ())
 {
   /**
@@ -42,17 +66,7 @@ __kernel void m18700_mxx (KERN_ATTR_RULES ())
 
     tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
 
-    u32 hash = 0;
-
-    for (u32 i = 0; i < tmp.pw_len; i++)
-    {
-      const u32 c32 = tmp.i[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    const u32 hash = hashCode (0, tmp.i, tmp.pw_len);
 
     const u32 r0 = hash;
     const u32 r1 = 0;
@@ -102,17 +116,7 @@ __kernel void m18700_sxx (KERN_ATTR_RULES ())
 
     tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
 
-    u32 hash = 0;
-
-    for (u32 i = 0; i < tmp.pw_len; i++)
-    {
-      const u32 c32 = tmp.i[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    const u32 hash = hashCode (0, tmp.i, tmp.pw_len);
 
     const u32 r0 = hash;
     const u32 r1 = 0;
diff --git a/OpenCL/m18700_a1-optimized.cl b/OpenCL/m18700_a1-optimized.cl
index 8c78f841d..45a555b81 100644
--- a/OpenCL/m18700_a1-optimized.cl
+++ b/OpenCL/m18700_a1-optimized.cl
@@ -14,6 +14,30 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+DECLSPEC u32 hashCode_g (const u32 init, __global u32 * restrict w, const u32 pw_len)
+{
+  u32 hash = init;
+
+  for (u32 i = 0; i < pw_len; i += 4)
+  {
+    u32 tmp = w[i / 4];
+
+    const u32 left = pw_len - i;
+
+    const u32 c = (left > 4) ? 4 : left;
+
+    switch (c)
+    {
+      case 4: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 3: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 2: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 1: hash *= 31; hash += tmp & 0xff;
+    }
+  }
+
+  return hash;
+}
+
 __kernel void m18700_m04 (KERN_ATTR_BASIC ())
 {
   /**
@@ -30,19 +54,7 @@ __kernel void m18700_m04 (KERN_ATTR_BASIC ())
 
   if (gid >= gid_max) return;
 
-  u32 pw_buf0[4];
-  u32 pw_buf1[4];
-
-  pw_buf0[0] = pws[gid].i[0];
-  pw_buf0[1] = pws[gid].i[1];
-  pw_buf0[2] = pws[gid].i[2];
-  pw_buf0[3] = pws[gid].i[3];
-  pw_buf1[0] = pws[gid].i[4];
-  pw_buf1[1] = pws[gid].i[5];
-  pw_buf1[2] = pws[gid].i[6];
-  pw_buf1[3] = pws[gid].i[7];
-
-  const u32 pw_l_len = pws[gid].pw_len & 63;
+  const u32 base = hashCode_g (0, pws[gid].i, pws[gid].pw_len);
 
   /**
    * loop
@@ -50,85 +62,7 @@ __kernel void m18700_m04 (KERN_ATTR_BASIC ())
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
   {
-    const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63;
-
-    const u32x pw_len = (pw_l_len + pw_r_len) & 63;
-
-    /**
-     * concat password candidate
-     */
-
-    u32x wordl0[4] = { 0 };
-    u32x wordl1[4] = { 0 };
-    u32x wordl2[4] = { 0 };
-    u32x wordl3[4] = { 0 };
-
-    wordl0[0] = pw_buf0[0];
-    wordl0[1] = pw_buf0[1];
-    wordl0[2] = pw_buf0[2];
-    wordl0[3] = pw_buf0[3];
-    wordl1[0] = pw_buf1[0];
-    wordl1[1] = pw_buf1[1];
-    wordl1[2] = pw_buf1[2];
-    wordl1[3] = pw_buf1[3];
-
-    u32x wordr0[4] = { 0 };
-    u32x wordr1[4] = { 0 };
-    u32x wordr2[4] = { 0 };
-    u32x wordr3[4] = { 0 };
-
-    wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
-    wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
-    wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
-    wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
-    wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
-    wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
-    wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
-    wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
-
-    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
-    {
-      switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
-    }
-    else
-    {
-      switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
-    }
-
-    u32x w[16];
-
-    w[ 0] = wordl0[0] | wordr0[0];
-    w[ 1] = wordl0[1] | wordr0[1];
-    w[ 2] = wordl0[2] | wordr0[2];
-    w[ 3] = wordl0[3] | wordr0[3];
-    w[ 4] = wordl1[0] | wordr1[0];
-    w[ 5] = wordl1[1] | wordr1[1];
-    w[ 6] = wordl1[2] | wordr1[2];
-    w[ 7] = wordl1[3] | wordr1[3];
-    w[ 8] = wordl2[0] | wordr2[0];
-    w[ 9] = wordl2[1] | wordr2[1];
-    w[10] = wordl2[2] | wordr2[2];
-    w[11] = wordl2[3] | wordr2[3];
-    w[12] = wordl3[0] | wordr3[0];
-    w[13] = wordl3[1] | wordr3[1];
-    w[14] = wordl3[2] | wordr3[2];
-    w[15] = wordl3[3] | wordr3[3];
-
-    /**
-     * hashCode()
-     */
-
-    u32x hash = 0;
-
-    for (u32 i = 0; i < pw_len; i++)
-    {
-      const u32 c32 = w[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    const u32 hash = hashCode_g (base, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
@@ -163,19 +97,7 @@ __kernel void m18700_s04 (KERN_ATTR_BASIC ())
 
   if (gid >= gid_max) return;
 
-  u32 pw_buf0[4];
-  u32 pw_buf1[4];
-
-  pw_buf0[0] = pws[gid].i[0];
-  pw_buf0[1] = pws[gid].i[1];
-  pw_buf0[2] = pws[gid].i[2];
-  pw_buf0[3] = pws[gid].i[3];
-  pw_buf1[0] = pws[gid].i[4];
-  pw_buf1[1] = pws[gid].i[5];
-  pw_buf1[2] = pws[gid].i[6];
-  pw_buf1[3] = pws[gid].i[7];
-
-  const u32 pw_l_len = pws[gid].pw_len & 63;
+  const u32 base = hashCode_g (0, pws[gid].i, pws[gid].pw_len);
 
   /**
    * digest
@@ -195,85 +117,7 @@ __kernel void m18700_s04 (KERN_ATTR_BASIC ())
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
   {
-    const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63;
-
-    const u32x pw_len = (pw_l_len + pw_r_len) & 63;
-
-    /**
-     * concat password candidate
-     */
-
-    u32x wordl0[4] = { 0 };
-    u32x wordl1[4] = { 0 };
-    u32x wordl2[4] = { 0 };
-    u32x wordl3[4] = { 0 };
-
-    wordl0[0] = pw_buf0[0];
-    wordl0[1] = pw_buf0[1];
-    wordl0[2] = pw_buf0[2];
-    wordl0[3] = pw_buf0[3];
-    wordl1[0] = pw_buf1[0];
-    wordl1[1] = pw_buf1[1];
-    wordl1[2] = pw_buf1[2];
-    wordl1[3] = pw_buf1[3];
-
-    u32x wordr0[4] = { 0 };
-    u32x wordr1[4] = { 0 };
-    u32x wordr2[4] = { 0 };
-    u32x wordr3[4] = { 0 };
-
-    wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
-    wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
-    wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
-    wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
-    wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
-    wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
-    wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
-    wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
-
-    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
-    {
-      switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
-    }
-    else
-    {
-      switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
-    }
-
-    u32x w[16];
-
-    w[ 0] = wordl0[0] | wordr0[0];
-    w[ 1] = wordl0[1] | wordr0[1];
-    w[ 2] = wordl0[2] | wordr0[2];
-    w[ 3] = wordl0[3] | wordr0[3];
-    w[ 4] = wordl1[0] | wordr1[0];
-    w[ 5] = wordl1[1] | wordr1[1];
-    w[ 6] = wordl1[2] | wordr1[2];
-    w[ 7] = wordl1[3] | wordr1[3];
-    w[ 8] = wordl2[0] | wordr2[0];
-    w[ 9] = wordl2[1] | wordr2[1];
-    w[10] = wordl2[2] | wordr2[2];
-    w[11] = wordl2[3] | wordr2[3];
-    w[12] = wordl3[0] | wordr3[0];
-    w[13] = wordl3[1] | wordr3[1];
-    w[14] = wordl3[2] | wordr3[2];
-    w[15] = wordl3[3] | wordr3[3];
-
-    /**
-     * hashCode()
-     */
-
-    u32x hash = 0;
-
-    for (u32 i = 0; i < pw_len; i++)
-    {
-      const u32 c32 = w[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    const u32 hash = hashCode_g (base, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
diff --git a/OpenCL/m18700_a1-pure.cl b/OpenCL/m18700_a1-pure.cl
index 0995a919e..ca15401ec 100644
--- a/OpenCL/m18700_a1-pure.cl
+++ b/OpenCL/m18700_a1-pure.cl
@@ -13,6 +13,30 @@
 #include "inc_scalar.cl"
 #include "inc_hash_md5.cl"
 
+DECLSPEC u32 hashCode_g (const u32 init, __global u32 * restrict w, const u32 pw_len)
+{
+  u32 hash = init;
+
+  for (u32 i = 0; i < pw_len; i += 4)
+  {
+    u32 tmp = w[i / 4];
+
+    const u32 left = pw_len - i;
+
+    const u32 c = (left > 4) ? 4 : left;
+
+    switch (c)
+    {
+      case 4: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 3: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 2: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 1: hash *= 31; hash += tmp & 0xff;
+    }
+  }
+
+  return hash;
+}
+
 __kernel void m18700_mxx (KERN_ATTR_BASIC ())
 {
   /**
@@ -28,17 +52,7 @@ __kernel void m18700_mxx (KERN_ATTR_BASIC ())
    * base
    */
 
-  u32 base = 0;
-
-  for (u32 i = 0; i < pws[gid].pw_len; i++)
-  {
-    const u32 c32 = pws[gid].i[i / 4];
-
-    const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-    base *= 31;
-    base += c;
-  }
+  const u32 base = hashCode_g (0, pws[gid].i, pws[gid].pw_len);
 
   /**
    * loop
@@ -46,17 +60,7 @@ __kernel void m18700_mxx (KERN_ATTR_BASIC ())
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
   {
-    u32 hash = base;
-
-    for (u32 i = 0; i < combs_buf[il_pos].pw_len; i++)
-    {
-      const u32 c32 = combs_buf[il_pos].i[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    const u32 hash = hashCode_g (base, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
@@ -94,17 +98,7 @@ __kernel void m18700_sxx (KERN_ATTR_BASIC ())
    * base
    */
 
-  u32 base = 0;
-
-  for (u32 i = 0; i < pws[gid].pw_len; i++)
-  {
-    const u32 c32 = pws[gid].i[i / 4];
-
-    const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-    base *= 31;
-    base += c;
-  }
+  const u32 base = hashCode_g (0, pws[gid].i, pws[gid].pw_len);
 
   /**
    * loop
@@ -112,17 +106,7 @@ __kernel void m18700_sxx (KERN_ATTR_BASIC ())
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
   {
-    u32 hash = base;
-
-    for (u32 i = 0; i < combs_buf[il_pos].pw_len; i++)
-    {
-      const u32 c32 = combs_buf[il_pos].i[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    const u32 hash = hashCode_g (base, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
diff --git a/OpenCL/m18700_a3-optimized.cl b/OpenCL/m18700_a3-optimized.cl
index d3f112370..54561f6d2 100644
--- a/OpenCL/m18700_a3-optimized.cl
+++ b/OpenCL/m18700_a3-optimized.cl
@@ -13,7 +13,49 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
-DECLSPEC void m18700m (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
+DECLSPEC u32x hashCode_w0 (const u32x init, const u32x w0, const u32 *w, const u32 pw_len)
+{
+  u32x hash = init;
+
+  u32x tmp0 = w0;
+
+  const u32 c0 = (pw_len > 4) ? 4 : pw_len;
+
+  switch (c0)
+  {
+    case 1: hash += tmp0 & 0xff; tmp0 >>= 8; break;
+    case 2: hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; break;
+    case 3: hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; break;
+    case 4: hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; break;
+  }
+
+  for (u32 i = 4; i < pw_len; i += 4)
+  {
+    u32 tmp = w[i / 4];
+
+    const u32 left = pw_len - i;
+
+    const u32 c = (left > 4) ? 4 : left;
+
+    switch (c)
+    {
+      case 4: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 3: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 2: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 1: hash *= 31; hash += tmp & 0xff;
+    }
+  }
+
+  return hash;
+}
+
+DECLSPEC void m18700m (const u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
 {
   /**
    * modifier
@@ -34,27 +76,7 @@ DECLSPEC void m18700m (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
 
     const u32x w0 = w0l | w0r;
 
-    u32x hash = 0;
-
-    for (u32 i = 0; i < 4; i++)
-    {
-      if (i == pw_len) break;
-
-      const u32x c = (w0 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
-
-    for (u32 i = 4; i < pw_len; i++)
-    {
-      const u32 c32 = w[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    const u32x hash = hashCode_w0 (0, w0, w, pw_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
@@ -65,7 +87,7 @@ DECLSPEC void m18700m (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
   }
 }
 
-DECLSPEC void m18700s (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
+DECLSPEC void m18700s (const u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
 {
   /**
    * modifier
@@ -98,27 +120,7 @@ DECLSPEC void m18700s (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
 
     const u32x w0 = w0l | w0r;
 
-    u32x hash = 0;
-
-    for (u32 i = 0; i < 4; i++)
-    {
-      if (i == pw_len) break;
-
-      const u32x c = (w0 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
-
-    for (u32 i = 4; i < pw_len; i++)
-    {
-      const u32 c32 = w[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    const u32x hash = hashCode_w0 (0, w0, w, pw_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
diff --git a/OpenCL/m18700_a3-pure.cl b/OpenCL/m18700_a3-pure.cl
index ae32446ce..1d599dc69 100644
--- a/OpenCL/m18700_a3-pure.cl
+++ b/OpenCL/m18700_a3-pure.cl
@@ -13,6 +13,48 @@
 #include "inc_simd.cl"
 #include "inc_hash_md5.cl"
 
+DECLSPEC u32x hashCode_w0 (const u32x init, const u32x w0, const u32 *w, const u32 pw_len)
+{
+  u32x hash = init;
+
+  u32x tmp0 = w0;
+
+  const u32 c0 = (pw_len > 4) ? 4 : pw_len;
+
+  switch (c0)
+  {
+    case 1: hash += tmp0 & 0xff; tmp0 >>= 8; break;
+    case 2: hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; break;
+    case 3: hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; break;
+    case 4: hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; hash *= 31;
+            hash += tmp0 & 0xff; tmp0 >>= 8; break;
+  }
+
+  for (u32 i = 4; i < pw_len; i += 4)
+  {
+    u32 tmp = w[i / 4];
+
+    const u32 left = pw_len - i;
+
+    const u32 c = (left > 4) ? 4 : left;
+
+    switch (c)
+    {
+      case 4: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 3: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 2: hash *= 31; hash += tmp & 0xff; tmp >>= 8;
+      case 1: hash *= 31; hash += tmp & 0xff;
+    }
+  }
+
+  return hash;
+}
+
 __kernel void m18700_mxx (KERN_ATTR_VECTOR ())
 {
   /**
@@ -30,7 +72,7 @@ __kernel void m18700_mxx (KERN_ATTR_VECTOR ())
 
   const u32 pw_len = pws[gid].pw_len & 255;
 
-  u32x w[64] = { 0 };
+  u32 w[64] = { 0 };
 
   for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
   {
@@ -49,19 +91,7 @@ __kernel void m18700_mxx (KERN_ATTR_VECTOR ())
 
     const u32x w0 = w0l | w0r;
 
-    w[0] = w0;
-
-    u32x hash = 0;
-
-    for (u32 i = 0; i < pw_len; i++)
-    {
-      const u32 c32 = w[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    u32x hash = hashCode_w0 (0, w0, w, pw_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;
@@ -101,7 +131,7 @@ __kernel void m18700_sxx (KERN_ATTR_VECTOR ())
 
   const u32 pw_len = pws[gid].pw_len & 255;
 
-  u32x w[64] = { 0 };
+  u32 w[64] = { 0 };
 
   for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
   {
@@ -120,19 +150,7 @@ __kernel void m18700_sxx (KERN_ATTR_VECTOR ())
 
     const u32x w0 = w0l | w0r;
 
-    w[0] = w0;
-
-    u32x hash = 0;
-
-    for (u32 i = 0; i < pw_len; i++)
-    {
-      const u32 c32 = w[i / 4];
-
-      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
-
-      hash *= 31;
-      hash += c;
-    }
+    u32x hash = hashCode_w0 (0, w0, w, pw_len);
 
     const u32x r0 = hash;
     const u32x r1 = 0;