From 04dfe6e89e69ace2494418ba27441452d328f985 Mon Sep 17 00:00:00 2001 From: jsteube Date: Wed, 1 Jun 2016 22:56:33 +0200 Subject: [PATCH] Fix -m 12500 by limiting max. length to 20; Limit fake rule copy in autotune() to it's max size --- OpenCL/m12500.cl | 12 +++++++----- src/hashcat.c | 14 ++++++++------ 2 files changed, 15 insertions(+), 11 deletions(-) diff --git a/OpenCL/m12500.cl b/OpenCL/m12500.cl index 8c6b6e371..403acd38a 100644 --- a/OpenCL/m12500.cl +++ b/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]; diff --git a/src/hashcat.c b/src/hashcat.c index 508ed458f..d6ecc231e 100644 --- a/src/hashcat.c +++ b/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); } - 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_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) + 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, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL); + } + } + else { run_kernel_amp (device_param, kernel_power_max); }