1
0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-25 17:38:23 +00:00

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.
This commit is contained in:
jsteube 2017-06-17 17:57:30 +02:00
parent 7905d79a28
commit a673aee037
11 changed files with 61 additions and 135 deletions

View File

@ -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];

View File

@ -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;
}

View File

@ -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

View File

@ -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

View File

@ -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)
{

View File

@ -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;

View File

@ -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;
}

View File

@ -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;

View File

@ -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];
}

View File

@ -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;

View File

@ -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];
}