Speedup -m 5300

[hashcat.git] / OpenCL / m05400_a3.cl

diff --git a/OpenCL/m05400_a3.cl b/OpenCL/m05400_a3.cl
index d5664b1..215d98b 100644 (file)
--- a/OpenCL/m05400_a3.cl
+++ b/OpenCL/m05400_a3.cl
@@ -765,26 +765,13 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05400_m16 (__glo
 
   __local u32 w_s[16];
 
-  if (lid < 1)
+  if (lid < 16)
   {
-    w_s[ 0] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 0]);
-    w_s[ 1] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 1]);
-    w_s[ 2] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 2]);
-    w_s[ 3] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 3]);
-    w_s[ 4] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 4]);
-    w_s[ 5] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 5]);
-    w_s[ 6] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 6]);
-    w_s[ 7] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 7]);
-    w_s[ 8] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 8]);
-    w_s[ 9] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 9]);
-    w_s[10] = swap32 (ikepsk_bufs[salt_pos].nr_buf[10]);
-    w_s[11] = swap32 (ikepsk_bufs[salt_pos].nr_buf[11]);
-    w_s[12] = swap32 (ikepsk_bufs[salt_pos].nr_buf[12]);
-    w_s[13] = swap32 (ikepsk_bufs[salt_pos].nr_buf[13]);
-    w_s[14] = 0;
-    w_s[15] = 0;
+    w_s[lid] = swap32 (ikepsk_bufs[salt_pos].nr_buf[lid]);
   }
 
+  barrier (CLK_LOCAL_MEM_FENCE);
+
   __local u32 s_msg_buf[128];
 
   const u32 lid2 = lid * 2;
@@ -853,26 +840,13 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05400_s04 (__glo
 
   __local u32 w_s[16];
 
-  if (lid < 1)
+  if (lid < 16)
   {
-    w_s[ 0] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 0]);
-    w_s[ 1] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 1]);
-    w_s[ 2] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 2]);
-    w_s[ 3] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 3]);
-    w_s[ 4] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 4]);
-    w_s[ 5] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 5]);
-    w_s[ 6] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 6]);
-    w_s[ 7] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 7]);
-    w_s[ 8] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 8]);
-    w_s[ 9] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 9]);
-    w_s[10] = swap32 (ikepsk_bufs[salt_pos].nr_buf[10]);
-    w_s[11] = swap32 (ikepsk_bufs[salt_pos].nr_buf[11]);
-    w_s[12] = swap32 (ikepsk_bufs[salt_pos].nr_buf[12]);
-    w_s[13] = swap32 (ikepsk_bufs[salt_pos].nr_buf[13]);
-    w_s[14] = 0;
-    w_s[15] = 0;
+    w_s[lid] = swap32 (ikepsk_bufs[salt_pos].nr_buf[lid]);
   }
 
+  barrier (CLK_LOCAL_MEM_FENCE);
+
   __local u32 s_msg_buf[128];
 
   const u32 lid2 = lid * 2;
@@ -941,26 +915,13 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05400_s08 (__glo
 
   __local u32 w_s[16];
 
-  if (lid < 1)
+  if (lid < 16)
   {
-    w_s[ 0] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 0]);
-    w_s[ 1] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 1]);
-    w_s[ 2] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 2]);
-    w_s[ 3] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 3]);
-    w_s[ 4] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 4]);
-    w_s[ 5] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 5]);
-    w_s[ 6] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 6]);
-    w_s[ 7] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 7]);
-    w_s[ 8] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 8]);
-    w_s[ 9] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 9]);
-    w_s[10] = swap32 (ikepsk_bufs[salt_pos].nr_buf[10]);
-    w_s[11] = swap32 (ikepsk_bufs[salt_pos].nr_buf[11]);
-    w_s[12] = swap32 (ikepsk_bufs[salt_pos].nr_buf[12]);
-    w_s[13] = swap32 (ikepsk_bufs[salt_pos].nr_buf[13]);
-    w_s[14] = 0;
-    w_s[15] = 0;
+    w_s[lid] = swap32 (ikepsk_bufs[salt_pos].nr_buf[lid]);
   }
 
+  barrier (CLK_LOCAL_MEM_FENCE);
+
   __local u32 s_msg_buf[128];
 
   const u32 lid2 = lid * 2;
@@ -1029,26 +990,13 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05400_s16 (__glo
 
   __local u32 w_s[16];
 
-  if (lid < 1)
+  if (lid < 16)
   {
-    w_s[ 0] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 0]);
-    w_s[ 1] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 1]);
-    w_s[ 2] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 2]);
-    w_s[ 3] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 3]);
-    w_s[ 4] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 4]);
-    w_s[ 5] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 5]);
-    w_s[ 6] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 6]);
-    w_s[ 7] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 7]);
-    w_s[ 8] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 8]);
-    w_s[ 9] = swap32 (ikepsk_bufs[salt_pos].nr_buf[ 9]);
-    w_s[10] = swap32 (ikepsk_bufs[salt_pos].nr_buf[10]);
-    w_s[11] = swap32 (ikepsk_bufs[salt_pos].nr_buf[11]);
-    w_s[12] = swap32 (ikepsk_bufs[salt_pos].nr_buf[12]);
-    w_s[13] = swap32 (ikepsk_bufs[salt_pos].nr_buf[13]);
-    w_s[14] = 0;
-    w_s[15] = 0;
+    w_s[lid] = swap32 (ikepsk_bufs[salt_pos].nr_buf[lid]);
   }
 
+  barrier (CLK_LOCAL_MEM_FENCE);
+
   __local u32 s_msg_buf[128];
 
   const u32 lid2 = lid * 2;