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

[hashcat.git] / OpenCL / m01800.cl

diff --git a/OpenCL/m01800.cl b/OpenCL/m01800.cl
index e0c91f8..ed562fd 100644 (file)
--- a/OpenCL/m01800.cl
+++ b/OpenCL/m01800.cl
@@ -5,25 +5,20 @@
 
 #define _SHA512_
 
-#include "include/constants.h"
-#include "include/kernel_vendor.h"
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
 
-#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"
-
-#define COMPARE_S "OpenCL/check_single_comp4.c"
-#define COMPARE_M "OpenCL/check_multi_comp4.c"
+#define COMPARE_S "inc_comp_single.cl"
+#define COMPARE_M "inc_comp_multi.cl"
 
 // Buggy drivers...
 
 #ifdef IS_AMD
 #define STATE_DECL volatile
+//#define STATE_DECL
 #else
 #define STATE_DECL
 #endif
@@ -133,7 +128,9 @@ void sha512_transform (const u64 w[16], u64 digest[8])
 
   ROUND_STEP (0);
 
+  #ifdef _unroll
   #pragma unroll
+  #endif
   for (int i = 16; i < 80; i += 16)
   {
     ROUND_EXPAND (); ROUND_STEP (i);
@@ -499,7 +496,9 @@ __kernel void m01800_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf
     {
       const u32 block_len = wpc_len[pc];
 
-      #pragma unroll 64
+      #ifdef _unroll
+      #pragma unroll
+      #endif
       for (u32 k = 0, p = block_len - 64; k < 64; k++, p++)
       {
         PUTCHAR64_BE (block, p, GETCHAR64_BE (l_alt_result, k));