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

[hashcat.git] / OpenCL / m09000.cl

diff --git a/OpenCL/m09000.cl b/OpenCL/m09000.cl
index 9e12988..7bbbeb4 100644 (file)
--- a/OpenCL/m09000.cl
+++ b/OpenCL/m09000.cl
@@ -5,20 +5,14 @@
 
 #define _PSAFE2_
 
-#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"
 
 // http://www.schneier.com/code/constants.txt
 
@@ -320,16 +314,16 @@ __constant u32 c_pbox[18] =
 #endif
 
 #ifdef IS_NV
-#define BF_ROUND(L,R,N)           \
-{                                 \
-  u32 tmp;                        \
-                                  \
-  tmp  = S0[__bfe ((L), 24, 8)];  \
-  tmp += S1[__bfe ((L), 16, 8)];  \
-  tmp ^= S2[__bfe ((L),  8, 8)];  \
-  tmp += S3[__bfe ((L),  0, 8)];  \
-                                  \
-  (R) ^= tmp ^ P[(N)];            \
+#define BF_ROUND(L,R,N)             \
+{                                   \
+  u32 tmp;                          \
+                                    \
+  tmp  = S0[__bfe_S ((L), 24, 8)];  \
+  tmp += S1[__bfe_S ((L), 16, 8)];  \
+  tmp ^= S2[__bfe_S ((L),  8, 8)];  \
+  tmp += S3[__bfe_S ((L),  0, 8)];  \
+                                    \
+  (R) ^= tmp ^ P[(N)];              \
 }
 #endif
 
@@ -747,6 +741,7 @@ __kernel void __attribute__((reqd_work_group_size (8, 1, 1))) m09000_loop (__glo
 
   u32 P[18];
 
+  #pragma unroll
   for (u32 i = 0; i < 18; i++)
   {
     P[i] = tmps[gid].P[i];
@@ -762,6 +757,7 @@ __kernel void __attribute__((reqd_work_group_size (8, 1, 1))) m09000_loop (__glo
   __local u32 *S2 = S2_all[lid];
   __local u32 *S3 = S3_all[lid];
 
+  #pragma unroll
   for (u32 i = 0; i < 256; i++)
   {
     S0[i] = tmps[gid].S0[i];