Remove MD4/MD5 *H1/*H2 functions and use original H functions. Modern compilers will...

[hashcat.git] / OpenCL / m01000_a1.cl

diff --git a/OpenCL/m01000_a1.cl b/OpenCL/m01000_a1.cl
index 4220109..280c459 100644 (file)
--- a/OpenCL/m01000_a1.cl
+++ b/OpenCL/m01000_a1.cl
@@ -20,7 +20,7 @@
 #define COMPARE_S "OpenCL/check_single_comp4.c"
 #define COMPARE_M "OpenCL/check_multi_comp4.c"
 
-__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m01000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
   /**
    * modifier
@@ -70,7 +70,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m04 (__glo
   {
     append_0x80_2x4 (wordl0, wordl1, pw_l_len);
 
-    switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
+    switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
   }
 
   /**
@@ -113,7 +113,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m04 (__glo
 
     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
     {
-      switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
+      switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
     }
 
     u32 w0[4];
@@ -154,8 +154,6 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m04 (__glo
 
     w3_t[2] = pw_len * 8 * 2;
 
-    u32 tmp2;
-
     u32 a = MD4M_A;
     u32 b = MD4M_B;
     u32 c = MD4M_C;
@@ -195,22 +193,22 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m04 (__glo
     MD4_STEP (MD4_Go, c, d, a, b, w2_t[3], MD4C01, MD4S12);
     MD4_STEP (MD4_Go, b, c, d, a, w3_t[3], MD4C01, MD4S13);
 
-    MD4_STEP (MD4_H1, a, b, c, d, w0_t[0], MD4C02, MD4S20);
-    MD4_STEP (MD4_H2, d, a, b, c, w2_t[0], MD4C02, MD4S21);
-    MD4_STEP (MD4_H1, c, d, a, b, w1_t[0], MD4C02, MD4S22);
-    MD4_STEP (MD4_H2, b, c, d, a, w3_t[0], MD4C02, MD4S23);
-    MD4_STEP (MD4_H1, a, b, c, d, w0_t[2], MD4C02, MD4S20);
-    MD4_STEP (MD4_H2, d, a, b, c, w2_t[2], MD4C02, MD4S21);
-    MD4_STEP (MD4_H1, c, d, a, b, w1_t[2], MD4C02, MD4S22);
-    MD4_STEP (MD4_H2, b, c, d, a, w3_t[2], MD4C02, MD4S23);
-    MD4_STEP (MD4_H1, a, b, c, d, w0_t[1], MD4C02, MD4S20);
-    MD4_STEP (MD4_H2, d, a, b, c, w2_t[1], MD4C02, MD4S21);
-    MD4_STEP (MD4_H1, c, d, a, b, w1_t[1], MD4C02, MD4S22);
-    MD4_STEP (MD4_H2, b, c, d, a, w3_t[1], MD4C02, MD4S23);
-    MD4_STEP (MD4_H1, a, b, c, d, w0_t[3], MD4C02, MD4S20);
-    MD4_STEP (MD4_H2, d, a, b, c, w2_t[3], MD4C02, MD4S21);
-    MD4_STEP (MD4_H1, c, d, a, b, w1_t[3], MD4C02, MD4S22);
-    MD4_STEP (MD4_H2, b, c, d, a, w3_t[3], MD4C02, MD4S23);
+    MD4_STEP (MD4_H , a, b, c, d, w0_t[0], MD4C02, MD4S20);
+    MD4_STEP (MD4_H , d, a, b, c, w2_t[0], MD4C02, MD4S21);
+    MD4_STEP (MD4_H , c, d, a, b, w1_t[0], MD4C02, MD4S22);
+    MD4_STEP (MD4_H , b, c, d, a, w3_t[0], MD4C02, MD4S23);
+    MD4_STEP (MD4_H , a, b, c, d, w0_t[2], MD4C02, MD4S20);
+    MD4_STEP (MD4_H , d, a, b, c, w2_t[2], MD4C02, MD4S21);
+    MD4_STEP (MD4_H , c, d, a, b, w1_t[2], MD4C02, MD4S22);
+    MD4_STEP (MD4_H , b, c, d, a, w3_t[2], MD4C02, MD4S23);
+    MD4_STEP (MD4_H , a, b, c, d, w0_t[1], MD4C02, MD4S20);
+    MD4_STEP (MD4_H , d, a, b, c, w2_t[1], MD4C02, MD4S21);
+    MD4_STEP (MD4_H , c, d, a, b, w1_t[1], MD4C02, MD4S22);
+    MD4_STEP (MD4_H , b, c, d, a, w3_t[1], MD4C02, MD4S23);
+    MD4_STEP (MD4_H , a, b, c, d, w0_t[3], MD4C02, MD4S20);
+    MD4_STEP (MD4_H , d, a, b, c, w2_t[3], MD4C02, MD4S21);
+    MD4_STEP (MD4_H , c, d, a, b, w1_t[3], MD4C02, MD4S22);
+    MD4_STEP (MD4_H , b, c, d, a, w3_t[3], MD4C02, MD4S23);
 
     const u32 r0 = a;
     const u32 r1 = d;
@@ -221,15 +219,15 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m04 (__glo
   }
 }
 
-__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_m08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m01000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_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))) m01000_m16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m01000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_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))) m01000_s04 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m01000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
   /**
    * modifier
@@ -279,7 +277,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_s04 (__glo
   {
     append_0x80_2x4 (wordl0, wordl1, pw_l_len);
 
-    switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
+    switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
   }
 
   /**
@@ -334,7 +332,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_s04 (__glo
 
     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
     {
-      switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
+      switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
     }
 
     u32 w0[4];
@@ -375,8 +373,6 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_s04 (__glo
 
     w3_t[2] = pw_len * 8 * 2;
 
-    u32 tmp2;
-
     u32 a = MD4M_A;
     u32 b = MD4M_B;
     u32 c = MD4M_C;
@@ -416,22 +412,22 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_s04 (__glo
     MD4_STEP (MD4_Go, c, d, a, b, w2_t[3], MD4C01, MD4S12);
     MD4_STEP (MD4_Go, b, c, d, a, w3_t[3], MD4C01, MD4S13);
 
-    MD4_STEP (MD4_H1, a, b, c, d, w0_t[0], MD4C02, MD4S20);
-    MD4_STEP (MD4_H2, d, a, b, c, w2_t[0], MD4C02, MD4S21);
-    MD4_STEP (MD4_H1, c, d, a, b, w1_t[0], MD4C02, MD4S22);
-    MD4_STEP (MD4_H2, b, c, d, a, w3_t[0], MD4C02, MD4S23);
-    MD4_STEP (MD4_H1, a, b, c, d, w0_t[2], MD4C02, MD4S20);
-    MD4_STEP (MD4_H2, d, a, b, c, w2_t[2], MD4C02, MD4S21);
-    MD4_STEP (MD4_H1, c, d, a, b, w1_t[2], MD4C02, MD4S22);
-    MD4_STEP (MD4_H2, b, c, d, a, w3_t[2], MD4C02, MD4S23);
-    MD4_STEP (MD4_H1, a, b, c, d, w0_t[1], MD4C02, MD4S20);
-    MD4_STEP (MD4_H2, d, a, b, c, w2_t[1], MD4C02, MD4S21);
-    MD4_STEP (MD4_H1, c, d, a, b, w1_t[1], MD4C02, MD4S22);
-    MD4_STEP (MD4_H2, b, c, d, a, w3_t[1], MD4C02, MD4S23);
-    MD4_STEP (MD4_H1, a, b, c, d, w0_t[3], MD4C02, MD4S20);
-    MD4_STEP (MD4_H2, d, a, b, c, w2_t[3], MD4C02, MD4S21);
-    MD4_STEP (MD4_H1, c, d, a, b, w1_t[3], MD4C02, MD4S22);
-    MD4_STEP (MD4_H2, b, c, d, a, w3_t[3], MD4C02, MD4S23);
+    MD4_STEP (MD4_H , a, b, c, d, w0_t[0], MD4C02, MD4S20);
+    MD4_STEP (MD4_H , d, a, b, c, w2_t[0], MD4C02, MD4S21);
+    MD4_STEP (MD4_H , c, d, a, b, w1_t[0], MD4C02, MD4S22);
+    MD4_STEP (MD4_H , b, c, d, a, w3_t[0], MD4C02, MD4S23);
+    MD4_STEP (MD4_H , a, b, c, d, w0_t[2], MD4C02, MD4S20);
+    MD4_STEP (MD4_H , d, a, b, c, w2_t[2], MD4C02, MD4S21);
+    MD4_STEP (MD4_H , c, d, a, b, w1_t[2], MD4C02, MD4S22);
+    MD4_STEP (MD4_H , b, c, d, a, w3_t[2], MD4C02, MD4S23);
+    MD4_STEP (MD4_H , a, b, c, d, w0_t[1], MD4C02, MD4S20);
+    MD4_STEP (MD4_H , d, a, b, c, w2_t[1], MD4C02, MD4S21);
+    MD4_STEP (MD4_H , c, d, a, b, w1_t[1], MD4C02, MD4S22);
+    MD4_STEP (MD4_H , b, c, d, a, w3_t[1], MD4C02, MD4S23);
+    MD4_STEP (MD4_H , a, b, c, d, w0_t[3], MD4C02, MD4S20);
+    MD4_STEP (MD4_H , d, a, b, c, w2_t[3], MD4C02, MD4S21);
+    MD4_STEP (MD4_H , c, d, a, b, w1_t[3], MD4C02, MD4S22);
+    MD4_STEP (MD4_H , b, c, d, a, w3_t[3], MD4C02, MD4S23);
 
     const u32 r0 = a;
     const u32 r1 = d;
@@ -442,10 +438,10 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_s04 (__glo
   }
 }
 
-__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m01000_s08 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m01000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_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))) m01000_s16 (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m01000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *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 void *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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
 }