From 51dc1c7db33d3ea7345edf259c0712d8871b564c Mon Sep 17 00:00:00 2001
From: jsteube <jens.steube@gmail.com>
Date: Thu, 24 Aug 2017 19:53:29 +0200
Subject: [PATCH] Use truncate_block_4x4_le_S() instead of
 truncate_block_4x4_le() in -m 6800

---
 OpenCL/inc_common.cl | 110 +++++++++++++++++++++++++++++++++++++++++++
 OpenCL/m06800.cl     |   2 +-
 2 files changed, 111 insertions(+), 1 deletion(-)

diff --git a/OpenCL/inc_common.cl b/OpenCL/inc_common.cl
index 9efbec58b..20c69899d 100644
--- a/OpenCL/inc_common.cl
+++ b/OpenCL/inc_common.cl
@@ -31506,6 +31506,116 @@ static void switch_buffer_by_offset_1x64_be (u32x w[64], const u32 offset)
  * vector functions as scalar (for outer loop usage)
  */
 
+static void truncate_block_4x4_le_S (u32 w0[4], const u32 len)
+{
+  switch (len)
+  {
+    case  0:
+      w0[0]  = 0;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  1:
+      w0[0] &= 0x000000ff;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  2:
+      w0[0] &= 0x0000ffff;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  3:
+      w0[0] &= 0x00ffffff;
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  4:
+      w0[1]  = 0;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  5:
+      w0[1] &= 0x000000ff;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  6:
+      w0[1] &= 0x0000ffff;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  7:
+      w0[1] &= 0x00ffffff;
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  8:
+      w0[2]  = 0;
+      w0[3]  = 0;
+
+      break;
+
+    case  9:
+      w0[2] &= 0x000000ff;
+      w0[3]  = 0;
+
+      break;
+
+    case 10:
+      w0[2] &= 0x0000ffff;
+      w0[3]  = 0;
+
+      break;
+
+    case 11:
+      w0[2] &= 0x00ffffff;
+      w0[3]  = 0;
+
+      break;
+
+    case 12:
+      w0[3]  = 0;
+
+      break;
+
+    case 13:
+      w0[3] &= 0x000000ff;
+
+      break;
+
+    case 14:
+      w0[3] &= 0x0000ffff;
+
+      break;
+
+    case 15:
+      w0[3] &= 0x00ffffff;
+
+      break;
+  }
+}
+
 static void append_0x01_2x4_S (u32 w0[4], u32 w1[4], const u32 offset)
 {
   const u32 tmp = 0x01 << ((offset & 3) * 8);
diff --git a/OpenCL/m06800.cl b/OpenCL/m06800.cl
index 2347307af..25104cd4e 100644
--- a/OpenCL/m06800.cl
+++ b/OpenCL/m06800.cl
@@ -356,7 +356,7 @@ __kernel void m06800_comp (__global pw_t *pws, __constant const kernel_rule_t *r
     out[2] = swap32_S (out[2]);
     out[3] = swap32_S (out[3]);
 
-    truncate_block_4x4_le (out, salt_len);
+    truncate_block_4x4_le_S (out, salt_len);
 
     if ((out[0] == salt_buf[0])
      && (out[1] == salt_buf[1])