Move macros DGST_R0 - DGST_R3 to host, define dgst_size for opencl kernel from host...

[hashcat.git] / OpenCL / m00200_a1.cl

diff --git a/OpenCL/m00200_a1.cl b/OpenCL/m00200_a1.cl
index 856fb56..fbbc263 100644 (file)
--- a/OpenCL/m00200_a1.cl
+++ b/OpenCL/m00200_a1.cl
@@ -8,18 +8,12 @@
 //incompatible
 //#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 "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_simd.cl"
 
 __kernel void m00200_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)
 {
@@ -149,8 +143,6 @@ __kernel void m00200_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     u32x a = MYSQL323_A;
     u32x b = MYSQL323_B;
-    u32x c = 0;
-    u32x d = 0;
 
     u32x add = 7;
 
@@ -197,7 +189,9 @@ __kernel void m00200_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     a &= 0x7fffffff;
     b &= 0x7fffffff;
 
-    COMPARE_M_SIMD (a, b, c, d);
+    u32x z = 0;
+
+    COMPARE_M_SIMD (a, b, z, z);
   }
 }
 
@@ -247,8 +241,8 @@ __kernel void m00200_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
   };
 
   /**
@@ -397,7 +391,9 @@ __kernel void m00200_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     a &= 0x7fffffff;
     b &= 0x7fffffff;
 
-    COMPARE_S_SIMD (a, b, c, d);
+    u32x z = 0;
+
+    COMPARE_S_SIMD (a, b, z, z);
   }
 }