From 3dd0a7140d4cab152283795980da393d74d0d20c Mon Sep 17 00:00:00 2001
From: jsteube <jens.steube@gmail.com>
Date: Fri, 22 Feb 2019 12:43:22 +0100
Subject: [PATCH] Respect combs_mode in -a 1/6/7 attack in -m 18700

---
 OpenCL/m18700_a1-optimized.cl | 30 +++++++++++++++++++++------
 OpenCL/m18700_a1-pure.cl      | 38 ++++++++++++++++++++++-------------
 2 files changed, 48 insertions(+), 20 deletions(-)

diff --git a/OpenCL/m18700_a1-optimized.cl b/OpenCL/m18700_a1-optimized.cl
index 45a555b81..60242880e 100644
--- a/OpenCL/m18700_a1-optimized.cl
+++ b/OpenCL/m18700_a1-optimized.cl
@@ -54,15 +54,24 @@ __kernel void m18700_m04 (KERN_ATTR_BASIC ())
 
   if (gid >= gid_max) return;
 
-  const u32 base = hashCode_g (0, pws[gid].i, pws[gid].pw_len);
-
   /**
    * loop
    */
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
   {
-    const u32 hash = hashCode_g (base, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    u32 hash = 0;
+
+    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
+    {
+      hash = hashCode_g (hash, pws[gid].i,          pws[gid].pw_len);
+      hash = hashCode_g (hash, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    }
+    else
+    {
+      hash = hashCode_g (hash, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+      hash = hashCode_g (hash, pws[gid].i,          pws[gid].pw_len);
+    }
 
     const u32x r0 = hash;
     const u32x r1 = 0;
@@ -97,8 +106,6 @@ __kernel void m18700_s04 (KERN_ATTR_BASIC ())
 
   if (gid >= gid_max) return;
 
-  const u32 base = hashCode_g (0, pws[gid].i, pws[gid].pw_len);
-
   /**
    * digest
    */
@@ -117,7 +124,18 @@ __kernel void m18700_s04 (KERN_ATTR_BASIC ())
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
   {
-    const u32 hash = hashCode_g (base, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    u32 hash = 0;
+
+    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
+    {
+      hash = hashCode_g (hash, pws[gid].i,          pws[gid].pw_len);
+      hash = hashCode_g (hash, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    }
+    else
+    {
+      hash = hashCode_g (hash, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+      hash = hashCode_g (hash, pws[gid].i,          pws[gid].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 ca15401ec..ef998501b 100644
--- a/OpenCL/m18700_a1-pure.cl
+++ b/OpenCL/m18700_a1-pure.cl
@@ -48,19 +48,24 @@ __kernel void m18700_mxx (KERN_ATTR_BASIC ())
 
   if (gid >= gid_max) return;
 
-  /**
-   * base
-   */
-
-  const u32 base = hashCode_g (0, pws[gid].i, pws[gid].pw_len);
-
   /**
    * loop
    */
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
   {
-    const u32 hash = hashCode_g (base, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    u32 hash = 0;
+
+    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
+    {
+      hash = hashCode_g (hash, pws[gid].i,          pws[gid].pw_len);
+      hash = hashCode_g (hash, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    }
+    else
+    {
+      hash = hashCode_g (hash, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+      hash = hashCode_g (hash, pws[gid].i,          pws[gid].pw_len);
+    }
 
     const u32x r0 = hash;
     const u32x r1 = 0;
@@ -94,19 +99,24 @@ __kernel void m18700_sxx (KERN_ATTR_BASIC ())
     0
   };
 
-  /**
-   * base
-   */
-
-  const u32 base = hashCode_g (0, pws[gid].i, pws[gid].pw_len);
-
   /**
    * loop
    */
 
   for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
   {
-    const u32 hash = hashCode_g (base, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    u32 hash = 0;
+
+    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
+    {
+      hash = hashCode_g (hash, pws[gid].i,          pws[gid].pw_len);
+      hash = hashCode_g (hash, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    }
+    else
+    {
+      hash = hashCode_g (hash, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+      hash = hashCode_g (hash, pws[gid].i,          pws[gid].pw_len);
+    }
 
     const u32x r0 = hash;
     const u32x r1 = 0;