Upgrade kernel to support dynamic local work sizes

[hashcat.git] / OpenCL / m11600.cl

diff --git a/OpenCL/m11600.cl b/OpenCL/m11600.cl
index 51ae126..0a9c077 100644 (file)
--- a/OpenCL/m11600.cl
+++ b/OpenCL/m11600.cl
@@ -1488,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 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)
+__kernel void 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
@@ -1554,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 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)
+__kernel void 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
@@ -1667,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 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)
+__kernel void 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
@@ -1675,13 +1675,12 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__gl
 
   const u32 gid = get_global_id (0);
   const u32 lid = get_local_id (0);
+  const u32 lsz = get_local_size (0);
 
   /**
    * aes shared
    */
 
-  const u32 lid4 = lid * 4;
-
   __local u32 s_td0[256];
   __local u32 s_td1[256];
   __local u32 s_td2[256];
@@ -1694,55 +1693,20 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m11600_comp (__gl
   __local u32 s_te3[256];
   __local u32 s_te4[256];
 
-  s_td0[lid4 + 0] = td0[lid4 + 0];
-  s_td0[lid4 + 1] = td0[lid4 + 1];
-  s_td0[lid4 + 2] = td0[lid4 + 2];
-  s_td0[lid4 + 3] = td0[lid4 + 3];
-
-  s_td1[lid4 + 0] = td1[lid4 + 0];
-  s_td1[lid4 + 1] = td1[lid4 + 1];
-  s_td1[lid4 + 2] = td1[lid4 + 2];
-  s_td1[lid4 + 3] = td1[lid4 + 3];
-
-  s_td2[lid4 + 0] = td2[lid4 + 0];
-  s_td2[lid4 + 1] = td2[lid4 + 1];
-  s_td2[lid4 + 2] = td2[lid4 + 2];
-  s_td2[lid4 + 3] = td2[lid4 + 3];
-
-  s_td3[lid4 + 0] = td3[lid4 + 0];
-  s_td3[lid4 + 1] = td3[lid4 + 1];
-  s_td3[lid4 + 2] = td3[lid4 + 2];
-  s_td3[lid4 + 3] = td3[lid4 + 3];
-
-  s_td4[lid4 + 0] = td4[lid4 + 0];
-  s_td4[lid4 + 1] = td4[lid4 + 1];
-  s_td4[lid4 + 2] = td4[lid4 + 2];
-  s_td4[lid4 + 3] = td4[lid4 + 3];
-
-  s_te0[lid4 + 0] = te0[lid4 + 0];
-  s_te0[lid4 + 1] = te0[lid4 + 1];
-  s_te0[lid4 + 2] = te0[lid4 + 2];
-  s_te0[lid4 + 3] = te0[lid4 + 3];
-
-  s_te1[lid4 + 0] = te1[lid4 + 0];
-  s_te1[lid4 + 1] = te1[lid4 + 1];
-  s_te1[lid4 + 2] = te1[lid4 + 2];
-  s_te1[lid4 + 3] = te1[lid4 + 3];
-
-  s_te2[lid4 + 0] = te2[lid4 + 0];
-  s_te2[lid4 + 1] = te2[lid4 + 1];
-  s_te2[lid4 + 2] = te2[lid4 + 2];
-  s_te2[lid4 + 3] = te2[lid4 + 3];
-
-  s_te3[lid4 + 0] = te3[lid4 + 0];
-  s_te3[lid4 + 1] = te3[lid4 + 1];
-  s_te3[lid4 + 2] = te3[lid4 + 2];
-  s_te3[lid4 + 3] = te3[lid4 + 3];
-
-  s_te4[lid4 + 0] = te4[lid4 + 0];
-  s_te4[lid4 + 1] = te4[lid4 + 1];
-  s_te4[lid4 + 2] = te4[lid4 + 2];
-  s_te4[lid4 + 3] = te4[lid4 + 3];
+  for (u32 i = lid; i < 256; i += lsz)
+  {
+    s_td0[i] = td0[i];
+    s_td1[i] = td1[i];
+    s_td2[i] = td2[i];
+    s_td3[i] = td3[i];
+    s_td4[i] = td4[i];
+
+    s_te0[i] = te0[i];
+    s_te1[i] = te1[i];
+    s_te2[i] = te2[i];
+    s_te3[i] = te3[i];
+    s_te4[i] = te4[i];
+  }
 
   barrier (CLK_LOCAL_MEM_FENCE);