Fix -m 12500 by limiting max. length to 20; Limit fake rule copy in autotune() to it's max size

OpenCL/m12500.cl

@@ -31,6 +31,8 @@
 #define PUTCHAR_BE(a,p,c) ((u8 *)(a))[(p) ^ 3] = (u8) (c)
 #define GETCHAR_BE(a,p)   ((u8 *)(a))[(p) ^ 3]
 
+#define MIN(a,b) (((a) < (b)) ? (a) : (b))
+
 __constant u32 te0[256] =
 {
   0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
@@ -1017,7 +1019,7 @@ __kernel void m12500_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf
   pw_buf[3] = pws[gid].i[3];
   pw_buf[4] = pws[gid].i[4];
 
-  const u32 pw_len = pws[gid].pw_len;
+  const u32 pw_len = MIN (pws[gid].pw_len, 20);
 
   u32 salt_buf[2];
 
@@ -1150,7 +1152,7 @@ __kernel void m12500_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len;
+  const u32 pw_len = MIN (pws[gid].pw_len, 20);
 
   const u32 salt_len = 8;
 
@@ -1226,16 +1228,16 @@ __kernel void m12500_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf
     pw_buf[3] = pws[gid].i[3];
     pw_buf[4] = pws[gid].i[4];
 
-    const u32 pw_len = pws[gid].pw_len;
+    //const u32 pw_len = pws[gid].pw_len;
 
     u32 salt_buf[2];
 
     salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
     salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 
-    const u32 salt_len = 8;
+    //const u32 salt_len = 8;
 
-    const u32 p3 = (pw_len * 2) + salt_len + 3;
+    //const u32 p3 = (pw_len * 2) + salt_len + 3;
 
     u32 w[16];
 

src/hashcat.c

@@ -2977,15 +2977,17 @@ static void autotune (hc_device_param_t *device_param)
       device_param->pws_buf[i].pw_len = 7 + (i & 7);
     }
 
+    hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, kernel_power_max * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
+  }
+
+  if (data.attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
+  {
     if (data.kernel_rules_cnt > 1)
     {
-      hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, 0, 0, kernel_loops_max * sizeof (kernel_rule_t), 0, NULL, NULL);
+      hc_clEnqueueCopyBuffer (data.ocl, device_param->command_queue, device_param->d_rules, device_param->d_rules_c, 0, 0, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL);
     }
-
-    hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, kernel_power_max * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
   }
-
-  if (data.attack_exec == ATTACK_EXEC_OUTSIDE_KERNEL)
+  else
   {
     run_kernel_amp (device_param, kernel_power_max);
   }