Merge pull request #348 from anthraxx/master

[hashcat.git] / OpenCL / m11700_a1.cl

diff --git a/OpenCL/m11700_a1.cl b/OpenCL/m11700_a1.cl
index c5b6808..36ce0b9 100644 (file)
--- a/OpenCL/m11700_a1.cl
+++ b/OpenCL/m11700_a1.cl
@@ -2241,13 +2241,15 @@ __constant u64 sbob_rc64[12][8] =
   },
 };
 
-static void streebog_g (u64x h[8], const u64x m[8], __local u64 (*s_sbob_sl64)[256])
+void streebog_g (u64x h[8], const u64x m[8], __local u64 (*s_sbob_sl64)[256])
 {
   u64x k[8];
   u64x s[8];
   u64x t[8];
 
+  #ifdef _unroll
   #pragma unroll
+  #endif
   for (int i = 0; i < 8; i++)
   {
     t[i] = h[i];
@@ -2258,7 +2260,9 @@ static void streebog_g (u64x h[8], const u64x m[8], __local u64 (*s_sbob_sl64)[2
     k[i] = SBOG_LPSti64;
   }
 
+  #ifdef _unroll
   #pragma unroll
+  #endif
   for (int i = 0; i < 8; i++)
   {
     s[i] = m[i];
@@ -2266,13 +2270,17 @@ static void streebog_g (u64x h[8], const u64x m[8], __local u64 (*s_sbob_sl64)[2
 
   for (int r = 0; r < 12; r++)
   {
+    #ifdef _unroll
     #pragma unroll
+    #endif
     for (int i = 0; i < 8; i++)
     {
       t[i] = s[i] ^ k[i];
     }
 
+    #ifdef _unroll
     #pragma unroll
+    #endif
     for (int i = 0; i < 8; i++)
     {
       s[i] = SBOG_LPSti64;
@@ -2283,14 +2291,18 @@ static void streebog_g (u64x h[8], const u64x m[8], __local u64 (*s_sbob_sl64)[2
       t[i] = k[i] ^ sbob_rc64[r][i];
     }
 
+    #ifdef _unroll
     #pragma unroll
+    #endif
     for (int i = 0; i < 8; i++)
     {
       k[i] = SBOG_LPSti64;
     }
   }
 
+  #ifdef _unroll
   #pragma unroll
+  #endif
   for (int i = 0; i < 8; i++)
   {
     h[i] ^= s[i] ^ k[i] ^ m[i];
@@ -2434,14 +2446,14 @@ __kernel void m11700_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     w[ 5] = w1[1];
     w[ 6] = w1[2];
     w[ 7] = w1[3];
-    w[ 8] = 0;
-    w[ 9] = 0;
-    w[10] = 0;
-    w[11] = 0;
-    w[12] = 0;
-    w[13] = 0;
-    w[14] = 0;
-    w[15] = 0;
+    w[ 8] = w2[0];
+    w[ 9] = w2[1];
+    w[10] = w2[2];
+    w[11] = w2[3];
+    w[12] = w3[0];
+    w[13] = w3[1];
+    w[14] = w3[2];
+    w[15] = w3[3];
 
     /**
      * reverse message block
@@ -2662,14 +2674,14 @@ __kernel void m11700_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     w[ 5] = w1[1];
     w[ 6] = w1[2];
     w[ 7] = w1[3];
-    w[ 8] = 0;
-    w[ 9] = 0;
-    w[10] = 0;
-    w[11] = 0;
-    w[12] = 0;
-    w[13] = 0;
-    w[14] = 0;
-    w[15] = 0;
+    w[ 8] = w2[0];
+    w[ 9] = w2[1];
+    w[10] = w2[2];
+    w[11] = w2[3];
+    w[12] = w3[0];
+    w[13] = w3[1];
+    w[14] = w3[2];
+    w[15] = w3[3];
 
     /**
      * reverse message block