Prepare for a more dynamic #pragma unroll use

[hashcat.git] / OpenCL / m01500_a1.cl

diff --git a/OpenCL/m01500_a1.cl b/OpenCL/m01500_a1.cl
index b4f4d67..135cf96 100644 (file)
--- a/OpenCL/m01500_a1.cl
+++ b/OpenCL/m01500_a1.cl
@@ -350,7 +350,7 @@ __constant u32 c_skb[8][64] =
 #define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
 #endif
 
-static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
+void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
 {
   u32x tt;
 
@@ -368,7 +368,9 @@ static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __loc
 
   c = c & 0x0fffffff;
 
-  #pragma unroll 16
+  #ifdef _unroll
+  #pragma unroll
+  #endif
   for (u32 i = 0; i < 16; i++)
   {
     if ((i < 2) || (i == 8) || (i == 15))
@@ -417,7 +419,7 @@ static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __loc
   }
 }
 
-static void _des_crypt_encrypt (u32x iv[2], u32 mask, u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
+void _des_crypt_encrypt (u32x iv[2], u32 mask, u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
 {
   const u32 E1 = (mask >> 2) & 0x3f0;
   const u32 E0 = mask & 0x3f;
@@ -427,6 +429,9 @@ static void _des_crypt_encrypt (u32x iv[2], u32 mask, u32x Kc[16], u32x Kd[16],
 
   for (u32 i = 0; i < 25; i++)
   {
+    #ifdef _unroll
+    #pragma unroll
+    #endif
     for (u32 j = 0; j < 16; j += 2)
     {
       u32x t;
@@ -641,10 +646,9 @@ __kernel void m01500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
 
-    u32x c = 0;
-    u32x d = 0;
+    u32x z = 0;
 
-    COMPARE_M_SIMD (iv[0], iv[1], c, d);
+    COMPARE_M_SIMD (iv[0], iv[1], z, z);
   }
 }
 
@@ -730,8 +734,8 @@ __kernel void m01500_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
   };
 
   /**
@@ -825,10 +829,9 @@ __kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
 
-    u32x c = 0;
-    u32x d = 0;
+    u32x z = 0;
 
-    COMPARE_S_SIMD (iv[0], iv[1], c, d);
+    COMPARE_S_SIMD (iv[0], iv[1], z, z);
   }
 }