Fix -m 1500 multihash and benchmark

[hashcat.git] / OpenCL / m11500_a0.cl

diff --git a/OpenCL/m11500_a0.cl b/OpenCL/m11500_a0.cl
index 8e9005c..3bcc15b 100644 (file)
--- a/OpenCL/m11500_a0.cl
+++ b/OpenCL/m11500_a0.cl
@@ -7,7 +7,8 @@
 
 #define _CRC32_
 
-#define NEW_SIMD_CODE
+//incompatible because of branches
+//#define NEW_SIMD_CODE
 
 #include "include/constants.h"
 #include "include/kernel_vendor.h"
@@ -92,7 +93,7 @@ __constant u32 crc32tab[0x100] =
   0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d
 };
 
-static u32x round_crc32 (u32x a, const u32x v)
+u32x round_crc32 (u32x a, const u32x v)
 {
   const u32x k = (a ^ v) & 0xff;
 
@@ -115,7 +116,7 @@ static u32x round_crc32 (u32x a, const u32x v)
   return a;
 }
 
-static u32x crc32 (const u32x w[16], const u32 pw_len, const u32 iv)
+u32x crc32 (const u32x w[16], const u32 pw_len, const u32 iv)
 {
   u32x a = iv ^ ~0;
 
@@ -152,14 +153,12 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   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];
-
-  u32 pw_buf1[4];
-
   pw_buf1[0] = pws[gid].i[ 4];
   pw_buf1[1] = pws[gid].i[ 5];
   pw_buf1[2] = pws[gid].i[ 6];
@@ -168,7 +167,7 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   const u32 pw_len = pws[gid].pw_len;
 
   /**
-   * digest
+   * salt
    */
 
   const u32 iv = salt_bufs[salt_pos].salt_buf[0];
@@ -186,6 +185,10 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
+    /**
+     * crc32
+     */
+
     u32x w[16];
 
     w[ 0] = w0[0];
@@ -205,12 +208,11 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     w[14] = 0;
     w[15] = 0;
 
-    u32x a = crc32_VV (w, out_len, iv);
-    u32x b = 0;
-    u32x c = 0;
-    u32x d = 0;
+    u32x a = crc32 (w, out_len, iv);
+
+    u32x z = 0;
 
-    COMPARE_M_SIMD (a, b, c, d);
+    COMPARE_M_SIMD (a, z, z, z);
   }
 }
 
@@ -239,14 +241,12 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   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];
-
-  u32 pw_buf1[4];
-
   pw_buf1[0] = pws[gid].i[ 4];
   pw_buf1[1] = pws[gid].i[ 5];
   pw_buf1[2] = pws[gid].i[ 6];
@@ -255,17 +255,21 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   const u32 pw_len = pws[gid].pw_len;
 
   /**
-   * digest
+   * salt
    */
 
   const u32 iv = salt_bufs[salt_pos].salt_buf[0];
 
+  /**
+   * digest
+   */
+
   const u32 search[4] =
   {
     digests_buf[digests_offset].digest_buf[DGST_R0],
-    digests_buf[digests_offset].digest_buf[DGST_R1],
-    digests_buf[digests_offset].digest_buf[DGST_R2],
-    digests_buf[digests_offset].digest_buf[DGST_R3]
+    0,
+    0,
+    0
   };
 
   /**
@@ -281,6 +285,10 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
+    /**
+     * crc32
+     */
+
     u32x w[16];
 
     w[ 0] = w0[0];
@@ -300,12 +308,11 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     w[14] = 0;
     w[15] = 0;
 
-    u32x a = crc32_VV (w, out_len, iv);
-    u32x b = 0;
-    u32x c = 0;
-    u32x d = 0;
+    u32x a = crc32 (w, out_len, iv);
+
+    u32x z = 0;
 
-    COMPARE_S_SIMD (a, b, c, d);
+    COMPARE_S_SIMD (a, z, z, z);
   }
 }