X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm05100_a0.cl;h=3d75ad21fe7b3ff5639e6934d533c0ce2a88501b;hb=ed1863c2630bc8194d38587d753dde71a315402a;hp=15488680c363731fde71e24cc580ff01413228e9;hpb=6a2c2934576a2220e535c42a9382de3668f15677;p=hashcat.git

diff --git a/OpenCL/m05100_a0.cl b/OpenCL/m05100_a0.cl
index 1548868..3d75ad2 100644
--- a/OpenCL/m05100_a0.cl
+++ b/OpenCL/m05100_a0.cl
@@ -7,20 +7,14 @@
 
 #define NEW_SIMD_CODE
 
-#include "include/constants.h"
-#include "include/kernel_vendor.h"
-
-#define DGST_R0 0
-#define DGST_R1 1
-#define DGST_R2 2
-#define DGST_R3 3
-
-#include "include/kernel_functions.c"
-#include "OpenCL/types_ocl.c"
-#include "OpenCL/common.c"
-#include "include/rp_kernel.h"
-#include "OpenCL/rp.c"
-#include "OpenCL/simd.c"
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
+#include "inc_rp.h"
+#include "inc_rp.cl"
+#include "inc_simd.cl"
 
 __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
@@ -39,18 +33,16 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   if (gid >= gid_max) return;
 
   u32 pw_buf0[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];
-  pw_buf1[3] = pws[gid].i[ 7];
+  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_len = pws[gid].pw_len;
 
@@ -70,6 +62,11 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     append_0x80_2x4_VV (w0, w1, out_len);
 
     w3[2] = out_len * 8;
+    w3[3] = 0;
+
+    /**
+     * md5
+     */
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;
@@ -149,14 +146,11 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     c += MD5M_C;
     d += MD5M_D;
 
-    u32x e = 0;
-    u32x f = 0;
-
-    COMPARE_M_SIMD (a, b, e, f);
+    u32x z = 0;
 
-    COMPARE_M_SIMD (b, c, e, f);
-
-    COMPARE_M_SIMD (c, d, e, f);
+    COMPARE_M_SIMD (a, b, z, z);
+    COMPARE_M_SIMD (b, c, z, z);
+    COMPARE_M_SIMD (c, d, z, z);
   }
 }
 
@@ -185,18 +179,16 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   if (gid >= gid_max) return;
 
   u32 pw_buf0[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];
-  pw_buf1[3] = pws[gid].i[ 7];
+  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_len = pws[gid].pw_len;
 
@@ -208,8 +200,8 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   {
     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
   };
 
   /**
@@ -228,6 +220,11 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     append_0x80_2x4_VV (w0, w1, out_len);
 
     w3[2] = out_len * 8;
+    w3[3] = 0;
+
+    /**
+     * md5
+     */
 
     u32x a = MD5M_A;
     u32x b = MD5M_B;
@@ -307,14 +304,11 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     c += MD5M_C;
     d += MD5M_D;
 
-    u32x e = 0;
-    u32x f = 0;
-
-    COMPARE_S_SIMD (a, b, e, f);
-
-    COMPARE_S_SIMD (b, c, e, f);
+    u32x z = 0;
 
-    COMPARE_S_SIMD (c, d, e, f);
+    COMPARE_S_SIMD (a, b, z, z);
+    COMPARE_S_SIMD (b, c, z, z);
+    COMPARE_S_SIMD (c, d, z, z);
   }
 }