From a673aee037ab2c2f78acb5518228e6567c9376b1 Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 17 Jun 2017 17:57:30 +0200 Subject: [PATCH] Very hot commit, continue reading here: This is a test commit using buffers large enough to handle both passwords and salts up to length 256. It requires changes to the kernel code, which is not included in here. It also requires some of the host code to be modified. Before we're going to modify kernel code to support the larger lengths I want to be sure of: 1. Host code modification is ok (no overflows or underflows) 2. Passwords and Salts are printed correctly to status, outfile, show, left, etc. 3. Performance does not change (or only very minimal) This is not a patch that supports actual cracking both passwords and salts up to length 256, but it can not fail anyway. If if it does, there's no reason to continue to add support for both passwords and salts up to length 256. --- OpenCL/amp_a0.cl | 26 ++------------------------ OpenCL/amp_a3.cl | 43 +------------------------------------------ OpenCL/inc_types.cl | 42 +++++++++++++++++++++--------------------- include/types.h | 34 +++++++++++----------------------- src/convert.c | 2 +- src/hashes.c | 2 +- src/interface.c | 21 +++++++++------------ src/opencl.c | 4 ++++ src/outfile.c | 8 ++++---- src/status.c | 4 ++-- src/stdout.c | 10 +++++----- 11 files changed, 61 insertions(+), 135 deletions(-) diff --git a/OpenCL/amp_a0.cl b/OpenCL/amp_a0.cl index 9720ba485..4cef595ad 100644 --- a/OpenCL/amp_a0.cl +++ b/OpenCL/amp_a0.cl @@ -16,32 +16,10 @@ __kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global const ke if (gid >= gid_max) return; + if (rules_buf[0].cmds[0] == RULE_OP_MANGLE_NOOP && rules_buf[0].cmds[1] == 0) return; + const u32 pw_len = pws[gid].pw_len; - if (rules_buf[0].cmds[0] == RULE_OP_MANGLE_NOOP && rules_buf[0].cmds[1] == 0) - { - pws_amp[gid].i[ 0] = pws[gid].i[ 0]; - pws_amp[gid].i[ 1] = pws[gid].i[ 1]; - pws_amp[gid].i[ 2] = pws[gid].i[ 2]; - pws_amp[gid].i[ 3] = pws[gid].i[ 3]; - pws_amp[gid].i[ 4] = pws[gid].i[ 4]; - pws_amp[gid].i[ 5] = pws[gid].i[ 5]; - pws_amp[gid].i[ 6] = pws[gid].i[ 6]; - pws_amp[gid].i[ 7] = pws[gid].i[ 7]; - pws_amp[gid].i[ 8] = pws[gid].i[ 8]; - pws_amp[gid].i[ 9] = pws[gid].i[ 9]; - pws_amp[gid].i[10] = pws[gid].i[10]; - pws_amp[gid].i[11] = pws[gid].i[11]; - pws_amp[gid].i[12] = pws[gid].i[12]; - pws_amp[gid].i[13] = pws[gid].i[13]; - pws_amp[gid].i[14] = pws[gid].i[14]; - pws_amp[gid].i[15] = pws[gid].i[15]; - - pws_amp[gid].pw_len = pws[gid].pw_len; - - return; - } - u32 w0[4]; u32 w1[4]; diff --git a/OpenCL/amp_a3.cl b/OpenCL/amp_a3.cl index dfafdbd2f..576e0bd49 100644 --- a/OpenCL/amp_a3.cl +++ b/OpenCL/amp_a3.cl @@ -13,48 +13,7 @@ __kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global const ke if (gid >= gid_max) return; - const u32 pw_len = pws[gid].pw_len; - - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; - - w0[0] = pws[gid].i[ 0]; - w0[1] = pws[gid].i[ 1]; - w0[2] = pws[gid].i[ 2]; - w0[3] = pws[gid].i[ 3]; - w1[0] = pws[gid].i[ 4]; - w1[1] = pws[gid].i[ 5]; - w1[2] = pws[gid].i[ 6]; - w1[3] = pws[gid].i[ 7]; - w2[0] = pws[gid].i[ 8]; - w2[1] = pws[gid].i[ 9]; - w2[2] = pws[gid].i[10]; - w2[3] = pws[gid].i[11]; - w3[0] = pws[gid].i[12]; - w3[1] = pws[gid].i[13]; - w3[2] = pws[gid].i[14]; - w3[3] = pws[gid].i[15]; - const u32 w0r = bfs_buf[0].i; - pws_amp[gid].i[ 0] = w0[0] | w0r; - pws_amp[gid].i[ 1] = w0[1]; - pws_amp[gid].i[ 2] = w0[2]; - pws_amp[gid].i[ 3] = w0[3]; - pws_amp[gid].i[ 4] = w1[0]; - pws_amp[gid].i[ 5] = w1[1]; - pws_amp[gid].i[ 6] = w1[2]; - pws_amp[gid].i[ 7] = w1[3]; - pws_amp[gid].i[ 8] = w2[0]; - pws_amp[gid].i[ 9] = w2[1]; - pws_amp[gid].i[10] = w2[2]; - pws_amp[gid].i[11] = w2[3]; - pws_amp[gid].i[12] = w3[0]; - pws_amp[gid].i[13] = w3[1]; - pws_amp[gid].i[14] = w3[2]; - pws_amp[gid].i[15] = w3[3]; - - pws_amp[gid].pw_len = pw_len; + pws_amp[gid].i[0] |= w0r; } diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index c9e9dc0b6..79af1b077 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -674,8 +674,8 @@ typedef struct digest typedef struct salt { - u32 salt_buf[16]; - u32 salt_buf_pc[16]; + u32 salt_buf[64]; + u32 salt_buf_pc[64]; u32 salt_len; u32 salt_len_pc; @@ -1494,19 +1494,9 @@ typedef struct } kernel_rule_t; -typedef struct +typedef struct pw { - u32 salt_pos; - u32 digest_pos; - u32 hash_pos; - u32 gidvid; - u32 il_pos; - -} plain_t; - -typedef struct -{ - u32 i[16]; + u32 i[64]; u32 pw_len; @@ -1516,26 +1506,36 @@ typedef struct } pw_t; -typedef struct +typedef struct bf { - u32 i; + u32 i; } bf_t; -typedef struct +typedef struct comb { - u32 i[8]; + u32 i[64]; - u32 pw_len; + u32 pw_len; } comb_t; -typedef struct +typedef struct bs_word { - u32 b[32]; + u32 b[32]; } bs_word_t; +typedef struct +{ + u32 salt_pos; + u32 digest_pos; + u32 hash_pos; + u32 gidvid; + u32 il_pos; + +} plain_t; + typedef struct { #ifndef SCRYPT_TMP_ELEM diff --git a/include/types.h b/include/types.h index 755961cee..c67712ee7 100644 --- a/include/types.h +++ b/include/types.h @@ -673,8 +673,8 @@ typedef enum user_options_map typedef struct salt { - u32 salt_buf[16]; - u32 salt_buf_pc[16]; + u32 salt_buf[64]; + u32 salt_buf_pc[64]; u32 salt_len; u32 salt_len_pc; @@ -838,7 +838,7 @@ typedef struct hashconfig hashconfig_t; typedef struct pw { - u32 i[16]; + u32 i[64]; u32 pw_len; @@ -854,20 +854,20 @@ typedef struct bf } bf_t; +typedef struct comb +{ + u32 i[64]; + + u32 pw_len; + +} comb_t; + typedef struct bs_word { u32 b[32]; } bs_word_t; -typedef struct comb -{ - u32 i[8]; - - u32 pw_len; - -} comb_t; - typedef struct cpt { u32 cracked; @@ -885,18 +885,6 @@ typedef struct plain } plain_t; -typedef struct wordl -{ - u32 word_buf[16]; - -} wordl_t; - -typedef struct wordr -{ - u32 word_buf[1]; - -} wordr_t; - #include "ext_OpenCL.h" typedef struct hc_device_param diff --git a/src/convert.c b/src/convert.c index 347ca4e91..3c2f591ff 100644 --- a/src/convert.c +++ b/src/convert.c @@ -157,7 +157,7 @@ bool need_hexify (const u8 *buf, const int len, const char separator, bool alway void exec_hexify (const u8 *buf, const int len, u8 *out) { - const int max_len = (len >= 31) ? 31 : len; + const int max_len = (len >= PW_MAX) ? PW_MAX : len; for (int i = max_len - 1, j = i * 2; i >= 0; i -= 1, j -= 2) { diff --git a/src/hashes.c b/src/hashes.c index 3f7ef6f47..a9130c235 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -248,7 +248,7 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl // plain - u32 plain_buf[16] = { 0 }; + u32 plain_buf[64] = { 0 }; u8 *plain_ptr = (u8 *) plain_buf; int plain_len = 0; diff --git a/src/interface.c b/src/interface.c index a31b5f672..ee2744092 100644 --- a/src/interface.c +++ b/src/interface.c @@ -2414,7 +2414,7 @@ static void drupal7_encode (u8 digest[64], u8 buf[43]) static u32 parse_and_store_salt (u8 *out, u8 *in, u32 salt_len, MAYBE_UNUSED const hashconfig_t *hashconfig) { - u32 tmp_u32[64] = { 0 }; + u32 tmp_u32[(64 * 2) + 1] = { 0 }; u8 *tmp = (u8 *) tmp_u32; @@ -2456,18 +2456,15 @@ static u32 parse_and_store_salt (u8 *out, u8 *in, u32 salt_len, MAYBE_UNUSED con if (hashconfig->opts_type & OPTS_TYPE_ST_UTF16LE) { - if (salt_len < 20) + if (salt_len <= 128) { - tmp_u32[9] = ((tmp_u32[4] >> 8) & 0x00FF0000) | ((tmp_u32[4] >> 16) & 0x000000FF); - tmp_u32[8] = ((tmp_u32[4] << 8) & 0x00FF0000) | ((tmp_u32[4] >> 0) & 0x000000FF); - tmp_u32[7] = ((tmp_u32[3] >> 8) & 0x00FF0000) | ((tmp_u32[3] >> 16) & 0x000000FF); - tmp_u32[6] = ((tmp_u32[3] << 8) & 0x00FF0000) | ((tmp_u32[3] >> 0) & 0x000000FF); - tmp_u32[5] = ((tmp_u32[2] >> 8) & 0x00FF0000) | ((tmp_u32[2] >> 16) & 0x000000FF); - tmp_u32[4] = ((tmp_u32[2] << 8) & 0x00FF0000) | ((tmp_u32[2] >> 0) & 0x000000FF); - tmp_u32[3] = ((tmp_u32[1] >> 8) & 0x00FF0000) | ((tmp_u32[1] >> 16) & 0x000000FF); - tmp_u32[2] = ((tmp_u32[1] << 8) & 0x00FF0000) | ((tmp_u32[1] >> 0) & 0x000000FF); - tmp_u32[1] = ((tmp_u32[0] >> 8) & 0x00FF0000) | ((tmp_u32[0] >> 16) & 0x000000FF); - tmp_u32[0] = ((tmp_u32[0] << 8) & 0x00FF0000) | ((tmp_u32[0] >> 0) & 0x000000FF); + for (int i = 64 - 1; i >= 1; i -= 2) + { + const u32 v = tmp_u32[i / 2]; + + tmp_u32[i - 0] = ((v >> 8) & 0x00FF0000) | ((v >> 16) & 0x000000FF); + tmp_u32[i - 1] = ((v << 8) & 0x00FF0000) | ((v >> 0) & 0x000000FF); + } salt_len = salt_len * 2; } diff --git a/src/opencl.c b/src/opencl.c index 7700d23fc..ac417fe54 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -1152,6 +1152,10 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, if (run_init == true) { + CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->command_queue, device_param->d_pws_buf, device_param->d_pws_amp_buf, 0, 0, pws_cnt * sizeof (pw_t), 0, NULL, NULL); + + if (CL_rc == -1) return -1; + CL_rc = run_kernel_amp (hashcat_ctx, device_param, pws_cnt); if (CL_rc == -1) return -1; diff --git a/src/outfile.c b/src/outfile.c index f305e2b76..94843d676 100644 --- a/src/outfile.c +++ b/src/outfile.c @@ -42,7 +42,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl if (rc == -1) return -1; - for (int i = 0; i < 16; i++) + for (int i = 0; i < 64; i++) { plain_buf[i] = pw.i[i]; } @@ -63,7 +63,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl if (rc == -1) return -1; - for (int i = 0; i < 16; i++) + for (int i = 0; i < 64; i++) { plain_buf[i] = pw.i[i]; } @@ -123,7 +123,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl if (rc == -1) return -1; - for (int i = 0; i < 16; i++) + for (int i = 0; i < 64; i++) { plain_buf[i] = pw.i[i]; } @@ -149,7 +149,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl if (rc == -1) return -1; - for (int i = 0; i < 16; i++) + for (int i = 0; i < 64; i++) { plain_buf[i] = pw.i[i]; } diff --git a/src/status.c b/src/status.c index 3070b2a6d..7b1b793bd 100644 --- a/src/status.c +++ b/src/status.c @@ -751,8 +751,8 @@ char *status_get_guess_candidates_dev (const hashcat_ctx_t *hashcat_ctx, const i plain_t plain1 = { 0, 0, 0, outerloop_first, innerloop_first }; plain_t plain2 = { 0, 0, 0, outerloop_last, innerloop_last }; - u32 plain_buf1[40] = { 0 }; - u32 plain_buf2[40] = { 0 }; + u32 plain_buf1[(64 * 2) + 2] = { 0 }; + u32 plain_buf2[(64 * 2) + 2] = { 0 }; u8 *plain_ptr1 = (u8 *) plain_buf1; u8 *plain_ptr2 = (u8 *) plain_buf2; diff --git a/src/stdout.c b/src/stdout.c index 5f64c533e..b24b00878 100644 --- a/src/stdout.c +++ b/src/stdout.c @@ -88,7 +88,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, out.len = 0; - u32 plain_buf[16] = { 0 }; + u32 plain_buf[64] = { 0 }; u8 *plain_ptr = (u8 *) plain_buf; @@ -115,7 +115,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - for (int i = 0; i < 8; i++) + for (int i = 0; i < 64; i++) { plain_buf[i] = pw.i[i]; } @@ -147,7 +147,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - for (int i = 0; i < 8; i++) + for (int i = 0; i < 64; i++) { plain_buf[i] = pw.i[i]; } @@ -217,7 +217,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - for (int i = 0; i < 8; i++) + for (int i = 0; i < 64; i++) { plain_buf[i] = pw.i[i]; } @@ -254,7 +254,7 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - for (int i = 0; i < 8; i++) + for (int i = 0; i < 64; i++) { plain_buf[i] = pw.i[i]; }