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

[hashcat.git] / OpenCL / m08000_a0.cl

diff --git a/OpenCL/m08000_a0.cl b/OpenCL/m08000_a0.cl
index 50104e8..cbe1584 100644 (file)
--- a/OpenCL/m08000_a0.cl
+++ b/OpenCL/m08000_a0.cl
@@ -9,20 +9,14 @@
 
 #define NEW_SIMD_CODE
 
-#include "include/constants.h"
-#include "include/kernel_vendor.h"
-
-#define DGST_R0 3
-#define DGST_R1 7
-#define DGST_R2 2
-#define DGST_R3 6
-
-#include "include/kernel_functions.c"
-#include "OpenCL/types_ocl.c"
-#include "OpenCL/common.c"
-#include "include/rp_kernel.h"
-#include "OpenCL/rp.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_rp.h"
+#include "inc_rp.cl"
+#include "inc_simd.cl"
 
 __constant u32 k_sha256[64] =
 {
@@ -49,7 +43,7 @@ __constant u32 k_sha256[64] =
 
 #define SHA256_EXPAND_S(x,y,z,w) (SHA256_S1_S (x) + y + SHA256_S0_S (z) + w)
 
-static void sha256_transform (u32x digest[8], const u32x w[16])
+void sha256_transform (u32x digest[8], const u32x w[16])
 {
   u32x a = digest[0];
   u32x b = digest[1];
@@ -119,7 +113,9 @@ static void sha256_transform (u32x digest[8], const u32x w[16])
 
   ROUND_STEP (0);
 
+  #ifdef _unroll
   #pragma unroll
+  #endif
   for (int i = 16; i < 64; i += 16)
   {
     ROUND_EXPAND (); ROUND_STEP (i);
@@ -135,7 +131,7 @@ static void sha256_transform (u32x digest[8], const u32x w[16])
   digest[7] += h;
 }
 
-static void sha256_transform_z (u32x digest[8])
+void sha256_transform_z (u32x digest[8])
 {
   u32x a = digest[0];
   u32x b = digest[1];
@@ -168,7 +164,9 @@ static void sha256_transform_z (u32x digest[8])
 
   ROUND_STEP_Z (0);
 
+  #ifdef _unroll
   #pragma unroll
+  #endif
   for (int i = 16; i < 64; i += 16)
   {
     ROUND_STEP_Z (i);
@@ -184,7 +182,7 @@ static void sha256_transform_z (u32x digest[8])
   digest[7] += h;
 }
 
-static void sha256_transform_s (u32x digest[8], __local u32 *w)
+void sha256_transform_s (u32x digest[8], __local u32 *w)
 {
   u32x a = digest[0];
   u32x b = digest[1];
@@ -217,7 +215,9 @@ static void sha256_transform_s (u32x digest[8], __local u32 *w)
 
   ROUND_STEP_S (0);
 
+  #ifdef _unroll
   #pragma unroll
+  #endif
   for (int i = 16; i < 64; i += 16)
   {
     ROUND_STEP_S (i);
@@ -270,7 +270,9 @@ __kernel void m08000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   {
     w_s1[15] =               0 | salt_buf0 >> 16;
 
+    #ifdef _unroll
     #pragma unroll
+    #endif
     for (int i = 16; i < 64; i++)
     {
       w_s1[i] = SHA256_EXPAND_S (w_s1[i - 2], w_s1[i - 7], w_s1[i - 15], w_s1[i - 16]);
@@ -281,7 +283,9 @@ __kernel void m08000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     w_s2[ 2] = salt_buf2 << 16 | 0;
     w_s2[15] = (510 + 8) * 8;
 
+    #ifdef _unroll
     #pragma unroll
+    #endif
     for (int i = 16; i < 64; i++)
     {
       w_s2[i] = SHA256_EXPAND_S (w_s2[i - 2], w_s2[i - 7], w_s2[i - 15], w_s2[i - 16]);
@@ -297,14 +301,12 @@ __kernel void m08000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
    */
 
   u32 pw_buf0[4];
+  u32 pw_buf1[4];
 
   pw_buf0[0] = pws[gid].i[ 0];
   pw_buf0[1] = pws[gid].i[ 1];
   pw_buf0[2] = pws[gid].i[ 2];
   pw_buf0[3] = pws[gid].i[ 3];
-
-  u32 pw_buf1[4];
-
   pw_buf1[0] = pws[gid].i[ 4];
   pw_buf1[1] = pws[gid].i[ 5];
   pw_buf1[2] = pws[gid].i[ 6];
@@ -323,7 +325,7 @@ __kernel void m08000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     u32x w2[4] = { 0 };
     u32x w3[4] = { 0 };
 
-    const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
+    apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
     u32x w0_t[4];
     u32x w1_t[4];
@@ -331,7 +333,6 @@ __kernel void m08000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     u32x w3_t[4];
 
     make_unicode (w0, w0_t, w1_t);
-
     make_unicode (w1, w2_t, w3_t);
 
     u32x w_t[16];
@@ -440,7 +441,9 @@ __kernel void m08000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
   {
     w_s1[15] =               0 | salt_buf0 >> 16;
 
+    #ifdef _unroll
     #pragma unroll
+    #endif
     for (int i = 16; i < 64; i++)
     {
       w_s1[i] = SHA256_EXPAND_S (w_s1[i - 2], w_s1[i - 7], w_s1[i - 15], w_s1[i - 16]);
@@ -451,7 +454,9 @@ __kernel void m08000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     w_s2[ 2] = salt_buf2 << 16 | 0;
     w_s2[15] = (510 + 8) * 8;
 
+    #ifdef _unroll
     #pragma unroll
+    #endif
     for (int i = 16; i < 64; i++)
     {
       w_s2[i] = SHA256_EXPAND_S (w_s2[i - 2], w_s2[i - 7], w_s2[i - 15], w_s2[i - 16]);
@@ -467,14 +472,12 @@ __kernel void m08000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
    */
 
   u32 pw_buf0[4];
+  u32 pw_buf1[4];
 
   pw_buf0[0] = pws[gid].i[ 0];
   pw_buf0[1] = pws[gid].i[ 1];
   pw_buf0[2] = pws[gid].i[ 2];
   pw_buf0[3] = pws[gid].i[ 3];
-
-  u32 pw_buf1[4];
-
   pw_buf1[0] = pws[gid].i[ 4];
   pw_buf1[1] = pws[gid].i[ 5];
   pw_buf1[2] = pws[gid].i[ 6];
@@ -505,7 +508,7 @@ __kernel void m08000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     u32x w2[4] = { 0 };
     u32x w3[4] = { 0 };
 
-    const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
+    apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
     u32x w0_t[4];
     u32x w1_t[4];
@@ -513,7 +516,6 @@ __kernel void m08000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     u32x w3_t[4];
 
     make_unicode (w0, w0_t, w1_t);
-
     make_unicode (w1, w2_t, w3_t);
 
     u32x w_t[16];