Replace the substring GPU to a more appropriate "device" or "kernel" substring depend...

[hashcat.git] / OpenCL / m11600.cl

diff --git a/OpenCL/m11600.cl b/OpenCL/m11600.cl
index 95d1022..5d64a77 100644 (file)
--- a/OpenCL/m11600.cl
+++ b/OpenCL/m11600.cl
@@ -14,12 +14,11 @@
 #define DGST_R3 3
 
 #include "include/kernel_functions.c"
-#include "types_ocl.c"
-#include "common.c"
+#include "OpenCL/types_ocl.c"
+#include "OpenCL/common.c"
 
-#ifdef  VECT_SIZE1
-#define COMPARE_M "check_multi_vect1_comp4.c"
-#endif
+#define COMPARE_S "OpenCL/check_single_comp4.c"
+#define COMPARE_M "OpenCL/check_multi_comp4.c"
 
 __constant u32 te0[256] =
 {
@@ -925,22 +924,22 @@ static void sha256_transform (const u32 w[16], u32 digest[8])
   u32 g = digest[6];
   u32 h = digest[7];
 
-  u32 w0_t = swap_workaround (w[ 0]);
-  u32 w1_t = swap_workaround (w[ 1]);
-  u32 w2_t = swap_workaround (w[ 2]);
-  u32 w3_t = swap_workaround (w[ 3]);
-  u32 w4_t = swap_workaround (w[ 4]);
-  u32 w5_t = swap_workaround (w[ 5]);
-  u32 w6_t = swap_workaround (w[ 6]);
-  u32 w7_t = swap_workaround (w[ 7]);
-  u32 w8_t = swap_workaround (w[ 8]);
-  u32 w9_t = swap_workaround (w[ 9]);
-  u32 wa_t = swap_workaround (w[10]);
-  u32 wb_t = swap_workaround (w[11]);
-  u32 wc_t = swap_workaround (w[12]);
-  u32 wd_t = swap_workaround (w[13]);
-  u32 we_t = swap_workaround (w[14]);
-  u32 wf_t = swap_workaround (w[15]);
+  u32 w0_t = swap32 (w[ 0]);
+  u32 w1_t = swap32 (w[ 1]);
+  u32 w2_t = swap32 (w[ 2]);
+  u32 w3_t = swap32 (w[ 3]);
+  u32 w4_t = swap32 (w[ 4]);
+  u32 w5_t = swap32 (w[ 5]);
+  u32 w6_t = swap32 (w[ 6]);
+  u32 w7_t = swap32 (w[ 7]);
+  u32 w8_t = swap32 (w[ 8]);
+  u32 w9_t = swap32 (w[ 9]);
+  u32 wa_t = swap32 (w[10]);
+  u32 wb_t = swap32 (w[11]);
+  u32 wc_t = swap32 (w[12]);
+  u32 wd_t = swap32 (w[13]);
+  u32 we_t = swap32 (w[14]);
+  u32 wf_t = swap32 (w[15]);
 
   #define ROUND_EXPAND()                            \
   {                                                 \
@@ -1074,21 +1073,7 @@ static u32 round_crc32 (u32 a, const u32 v)
 
   const u32 s = a >> 8;
 
-  #ifdef VECT_SIZE1
-  a    = crc32tab[k];
-  #endif
-
-  #ifdef VECT_SIZE2
-  a.s0 = crc32tab[k.s0];
-  a.s1 = crc32tab[k.s1];
-  #endif
-
-  #ifdef VECT_SIZE4
-  a.s0 = crc32tab[k.s0];
-  a.s1 = crc32tab[k.s1];
-  a.s2 = crc32tab[k.s2];
-  a.s3 = crc32tab[k.s3];
-  #endif
+  a = crc32tab[k];
 
   a ^= s;
 
@@ -1144,6 +1129,17 @@ static u32 memcat8c (u32 block[16], const u32 block_len, const u32 append[2], co
   u32 tmp1;
   u32 tmp2;
 
+  #ifdef IS_NV
+  const int offset_minus_4 = 4 - (block_len & 3);
+
+  const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
+
+  tmp0 = __byte_perm (        0, append[0], selector);
+  tmp1 = __byte_perm (append[0], append[1], selector);
+  tmp2 = __byte_perm (append[1],         0, selector);
+  #endif
+
+  #ifdef IS_AMD
   const int offset_minus_4 = 4 - block_len;
 
   tmp0 = amd_bytealign (append[0],         0, offset_minus_4);
@@ -1156,6 +1152,7 @@ static u32 memcat8c (u32 block[16], const u32 block_len, const u32 append[2], co
     tmp1 = tmp2;
     tmp2 = 0;
   }
+  #endif
 
   u32 carry[2] = { 0, 0 };
 
@@ -1259,6 +1256,23 @@ static u32 memcat32c (u32 block[16], const u32 block_len, const u32 append[8], c
   u32 tmp7;
   u32 tmp8;
 
+  #ifdef IS_NV
+  const int offset_minus_4 = 4 - (block_len & 3);
+
+  const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
+
+  tmp0 = __byte_perm (        0, append[0], selector);
+  tmp1 = __byte_perm (append[0], append[1], selector);
+  tmp2 = __byte_perm (append[1], append[2], selector);
+  tmp3 = __byte_perm (append[2], append[3], selector);
+  tmp4 = __byte_perm (append[3], append[4], selector);
+  tmp5 = __byte_perm (append[4], append[5], selector);
+  tmp6 = __byte_perm (append[5], append[6], selector);
+  tmp7 = __byte_perm (append[6], append[7], selector);
+  tmp8 = __byte_perm (append[7],         0, selector);
+  #endif
+
+  #ifdef IS_AMD
   const int offset_minus_4 = 4 - block_len;
 
   tmp0 = amd_bytealign (append[0],         0, offset_minus_4);
@@ -1283,6 +1297,7 @@ static u32 memcat32c (u32 block[16], const u32 block_len, const u32 append[8], c
     tmp7 = tmp8;
     tmp8 = 0;
   }
+  #endif
 
   u32 carry[8] = { 0, 0, 0, 0, 0, 0, 0, 0 };
 
@@ -1473,7 +1488,7 @@ static u32 memcat32c (u32 block[16], const u32 block_len, const u32 append[8], c
   return new_len;
 }
 
-__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_init (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global seven_zip_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global seven_zip_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
   /**
    * base
@@ -1539,7 +1554,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_init (__gl
   tmps[gid].final_len = final_len;
 }
 
-__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_loop (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global seven_zip_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global seven_zip_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
   /**
    * base
@@ -1652,7 +1667,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_loop (__gl
   tmps[gid].final_len = final_len;
 }
 
-__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global seven_zip_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global seven_zip_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
   /**
    * base
@@ -1770,7 +1785,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__gl
   u32 block_len = tmps[gid].block_len;
   u32 final_len = tmps[gid].final_len;
 
-  append_0x80_4 (block, block_len);
+  append_0x80_1x16 (block, block_len);
 
   if (block_len >= 56)
   {
@@ -1779,7 +1794,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__gl
     bzero16 (block);
   }
 
-  block[15] = swap_workaround (final_len * 8);
+  block[15] = swap32 (final_len * 8);
 
   sha256_transform (block, dgst);
 
@@ -1825,10 +1840,10 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__gl
   {
     u32 data[4];
 
-    data[0] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 0]);
-    data[1] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 1]);
-    data[2] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 2]);
-    data[3] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 3]);
+    data[0] = swap32 (esalt_bufs[salt_pos].data_buf[j + 0]);
+    data[1] = swap32 (esalt_bufs[salt_pos].data_buf[j + 1]);
+    data[2] = swap32 (esalt_bufs[salt_pos].data_buf[j + 2]);
+    data[3] = swap32 (esalt_bufs[salt_pos].data_buf[j + 3]);
 
     u32 out[4];
 
@@ -1844,20 +1859,20 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__gl
     iv[2] = data[2];
     iv[3] = data[3];
 
-    out[0] = swap_workaround (out[0]);
-    out[1] = swap_workaround (out[1]);
-    out[2] = swap_workaround (out[2]);
-    out[3] = swap_workaround (out[3]);
+    out[0] = swap32 (out[0]);
+    out[1] = swap32 (out[1]);
+    out[2] = swap32 (out[2]);
+    out[3] = swap32 (out[3]);
 
     crc = crc32 (out, 16, crc);
   }
 
   u32 data[4];
 
-  data[0] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 0]);
-  data[1] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 1]);
-  data[2] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 2]);
-  data[3] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 3]);
+  data[0] = swap32 (esalt_bufs[salt_pos].data_buf[j + 0]);
+  data[1] = swap32 (esalt_bufs[salt_pos].data_buf[j + 1]);
+  data[2] = swap32 (esalt_bufs[salt_pos].data_buf[j + 2]);
+  data[3] = swap32 (esalt_bufs[salt_pos].data_buf[j + 3]);
 
   u32 out[4];
 
@@ -1873,10 +1888,10 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__gl
   iv[2] = data[2];
   iv[3] = data[3];
 
-  out[0] = swap_workaround (out[0]);
-  out[1] = swap_workaround (out[1]);
-  out[2] = swap_workaround (out[2]);
-  out[3] = swap_workaround (out[3]);
+  out[0] = swap32 (out[0]);
+  out[1] = swap32 (out[1]);
+  out[2] = swap32 (out[2]);
+  out[3] = swap32 (out[3]);
 
   const u32 margin = data_len - unpack_size;