mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-26 09:58:16 +00:00
First working -m 16100 kernel
This commit is contained in:
jsteube
parent
7a54490da2
commit
5847067c96
@ -1467,6 +1467,17 @@ typedef struct ethereum_scrypt
|
||||
|
||||
} ethereum_scrypt_t;
|
||||
|
||||
typedef struct tacacs_plus
|
||||
{
|
||||
u32 session_buf[16];
|
||||
|
||||
u32 ct_data_buf[64];
|
||||
u32 ct_data_len;
|
||||
|
||||
u32 sequence_buf[16];
|
||||
|
||||
} tacacs_plus_t;
|
||||
|
||||
typedef struct pdf14_tmp
|
||||
{
|
||||
u32 digest[4];
|
||||
|
||||
131
OpenCL/m16100_a0.cl
Normal file
131
OpenCL/m16100_a0.cl
Normal file
@ -0,0 +1,131 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
//#define NEW_SIMD_CODE
|
||||
|
||||
#include "inc_vendor.cl"
|
||||
#include "inc_hash_constants.h"
|
||||
#include "inc_hash_functions.cl"
|
||||
#include "inc_types.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
__kernel void m16100_mxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tacacs_plus_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m16100_sxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const tacacs_plus_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
COPY_PW (pws[gid]);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
u32 es0[4];
|
||||
u32 es1[4];
|
||||
u32 es2[4];
|
||||
u32 es3[4];
|
||||
|
||||
es0[ 0] = esalt_bufs[digests_offset].session_buf[0];
|
||||
es0[ 1] = 0;
|
||||
es0[ 2] = 0;
|
||||
es0[ 3] = 0;
|
||||
es1[ 0] = 0;
|
||||
es1[ 1] = 0;
|
||||
es1[ 2] = 0;
|
||||
es1[ 3] = 0;
|
||||
es2[ 0] = 0;
|
||||
es2[ 1] = 0;
|
||||
es2[ 2] = 0;
|
||||
es2[ 3] = 0;
|
||||
es3[ 0] = 0;
|
||||
es3[ 1] = 0;
|
||||
es3[ 2] = 0;
|
||||
es3[ 3] = 0;
|
||||
|
||||
md5_update_64 (&ctx0, es0, es1, es2, es3, 4);
|
||||
|
||||
es0[ 0] = esalt_bufs[digests_offset].sequence_buf[0];
|
||||
es0[ 1] = 0;
|
||||
es0[ 2] = 0;
|
||||
es0[ 3] = 0;
|
||||
es1[ 0] = 0;
|
||||
es1[ 1] = 0;
|
||||
es1[ 2] = 0;
|
||||
es1[ 3] = 0;
|
||||
es2[ 0] = 0;
|
||||
es2[ 1] = 0;
|
||||
es2[ 2] = 0;
|
||||
es2[ 3] = 0;
|
||||
es3[ 0] = 0;
|
||||
es3[ 1] = 0;
|
||||
es3[ 2] = 0;
|
||||
es3[ 3] = 0;
|
||||
|
||||
u32 ct_buf[2];
|
||||
|
||||
ct_buf[0] = esalt_bufs[digests_offset].ct_data_buf[0];
|
||||
ct_buf[1] = esalt_bufs[digests_offset].ct_data_buf[1];
|
||||
|
||||
u32 cl_len = esalt_bufs[digests_offset].ct_data_len;
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
pw_t tmp = PASTE_PW;
|
||||
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update_64 (&ctx, es0, es1, es2, es3, 2);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
u32 test[2];
|
||||
|
||||
test[0] = ctx.h[0] ^ ct_buf[0];
|
||||
test[1] = ctx.h[1] ^ ct_buf[1];
|
||||
|
||||
const u32 status = ((test[0] >> 0) & 0xff);
|
||||
const u32 flags = ((test[0] >> 8) & 0xff);
|
||||
const u32 msg_len = ((test[0] >> 16) & 0xff) << 8
|
||||
| ((test[0] >> 24) & 0xff) << 0;
|
||||
const u32 data_len = ((test[1] >> 0) & 0xff) << 8
|
||||
| ((test[1] >> 8) & 0xff) << 0;
|
||||
|
||||
if (((status >= 0x01 && status <= 0x07) || status == 0x21)
|
||||
&& (flags == 0x01 || flags == 0x00)
|
||||
&& (6 + msg_len + data_len == cl_len))
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -469,6 +469,17 @@ typedef struct ethereum_scrypt
|
||||
|
||||
} ethereum_scrypt_t;
|
||||
|
||||
typedef struct tacacs_plus
|
||||
{
|
||||
u32 session_buf[16];
|
||||
|
||||
u32 ct_data_buf[64];
|
||||
u32 ct_data_len;
|
||||
|
||||
u32 sequence_buf[16];
|
||||
|
||||
} tacacs_plus_t;
|
||||
|
||||
typedef struct luks_tmp
|
||||
{
|
||||
u32 ipad32[8];
|
||||
@ -1283,6 +1294,8 @@ typedef enum display_len
|
||||
DISPLAY_LEN_MAX_15900 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 100 + 1 + 6 + 1 + 6 + 1 + 10 + 1 + 32 + 1 + 4 + 1 + 512,
|
||||
DISPLAY_LEN_MIN_16000 = 10,
|
||||
DISPLAY_LEN_MAX_16000 = 10,
|
||||
DISPLAY_LEN_MIN_16100 = 1 + 11 + 1 + 1 + 1 + 8 + 12 + 1 + 4,
|
||||
DISPLAY_LEN_MAX_16100 = 1 + 11 + 1 + 1 + 1 + 8 + 256 + 1 + 4,
|
||||
DISPLAY_LEN_MIN_99999 = 1,
|
||||
DISPLAY_LEN_MAX_99999 = 55,
|
||||
|
||||
@ -1411,6 +1424,7 @@ typedef enum hash_type
|
||||
HASH_TYPE_CHACHA20 = 60,
|
||||
HASH_TYPE_DPAPIMK = 61,
|
||||
HASH_TYPE_JKS_SHA1 = 62,
|
||||
HASH_TYPE_TACACS_PLUS = 63,
|
||||
|
||||
} hash_type_t;
|
||||
|
||||
@ -1608,6 +1622,7 @@ typedef enum kern_type
|
||||
KERN_TYPE_ETHEREUM_SCRYPT = 15700,
|
||||
KERN_TYPE_DPAPIMK_V2 = 15900,
|
||||
KERN_TYPE_TRIPCODE = 16000,
|
||||
KERN_TYPE_TACACS_PLUS = 16100,
|
||||
KERN_TYPE_PLAINTEXT = 99999,
|
||||
|
||||
} kern_type_t;
|
||||
@ -1866,6 +1881,8 @@ int jks_sha1_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu
|
||||
int ethereum_pbkdf2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int ethereum_scrypt_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int tripcode_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int tacacs_plus_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
|
||||
/**
|
||||
* hook functions
|
||||
*/
|
||||
|
||||
149
src/interface.c
149
src/interface.c
@ -266,6 +266,7 @@ static const char *ST_HASH_15600 = "$ethereum$p*1024*383531313538313333383131383
|
||||
static const char *ST_HASH_15700 = "$ethereum$s*1024*1*1*3033363133373132373638333437323331383637383437333631373038323434*69eaf081695cf971ef7ee5a49997c1a3922e7efef59068109e83853755ee31c3*64a1adec1750ee4416b22b81111dd2a3c2fede820d6da8bf788dca2641d5b181";
|
||||
static const char *ST_HASH_15900 = "$DPAPImk$2*1*S-15-21-439882973-489230393-482956683-1522*aes256*sha512*12900*79f7ca399f2626e21aad108c3922af7c*288*c47bc8a985ca6aa708b01c97b004bff20cc52379dc2635b4acf59ce17970a2cb47ace98c7e8de977f265243c5c03d0a97e4b954b494d9e38d9158d0c1e729d16a28ba69e2e7c6c3bc0e3afc9c9b6306b83372ccb35d89b98925728fd36315b8ee95b4d4eccdcb31564769f9a4b9ee10828184e16d4af336675d5e31d987dd87233d34fbbb98880c5e1f64cbb9b043ad8";
|
||||
static const char *ST_HASH_16000 = "pfaRCwDe0U";
|
||||
static const char *ST_HASH_16100 = "$tacacs-plus$0$5fde8e68$4e13e8fb33df$c006";
|
||||
static const char *ST_HASH_99999 = "hashcat";
|
||||
|
||||
static const char *OPTI_STR_OPTIMIZED_KERNEL = "Optimized-Kernel";
|
||||
@ -499,6 +500,7 @@ static const char *HT_15600 = "Ethereum Wallet, PBKDF2-HMAC-SHA256";
|
||||
static const char *HT_15700 = "Ethereum Wallet, SCRYPT";
|
||||
static const char *HT_15900 = "DPAPI masterkey file v2";
|
||||
static const char *HT_16000 = "Tripcode";
|
||||
static const char *HT_16100 = "TACACS+";
|
||||
static const char *HT_99999 = "Plaintext";
|
||||
|
||||
static const char *HT_00011 = "Joomla < 2.5.18";
|
||||
@ -643,6 +645,7 @@ static const char *SIGNATURE_CHACHA20 = "$chacha20$";
|
||||
static const char *SIGNATURE_JKS_SHA1 = "$jksprivk$";
|
||||
static const char *SIGNATURE_ETHEREUM_PBKDF2 = "$ethereum$p";
|
||||
static const char *SIGNATURE_ETHEREUM_SCRYPT = "$ethereum$s";
|
||||
static const char *SIGNATURE_TACACS_PLUS = "$tacacs-plus$0$";
|
||||
|
||||
/**
|
||||
* decoder / encoder
|
||||
@ -15723,6 +15726,106 @@ int tripcode_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_U
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int tacacs_plus_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
|
||||
{
|
||||
if ((input_len < DISPLAY_LEN_MIN_16100) || (input_len > DISPLAY_LEN_MAX_16100)) return (PARSER_GLOBAL_LENGTH);
|
||||
|
||||
if (memcmp (SIGNATURE_TACACS_PLUS, input_buf, 15) != 0) return (PARSER_SIGNATURE_UNMATCHED);
|
||||
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
|
||||
salt_t *salt = hash_buf->salt;
|
||||
|
||||
tacacs_plus_t *tacacs_plus = (tacacs_plus_t *) hash_buf->esalt;
|
||||
|
||||
/**
|
||||
* parse line
|
||||
*/
|
||||
|
||||
// session
|
||||
|
||||
u8 *session_pos = input_buf + 15;
|
||||
|
||||
// ct_buf
|
||||
|
||||
u8 *ct_buf_pos = (u8 *) strchr ((const char *) session_pos, '$');
|
||||
|
||||
if (ct_buf_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
u32 session_len = ct_buf_pos - session_pos;
|
||||
|
||||
ct_buf_pos++;
|
||||
|
||||
// sequence
|
||||
|
||||
u8 *sequence_pos = (u8 *) strchr ((const char *) ct_buf_pos, '$');
|
||||
|
||||
if (sequence_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
u32 ct_buf_len = sequence_pos - ct_buf_pos;
|
||||
|
||||
sequence_pos++;
|
||||
|
||||
u32 sequence_len = input_len - 15 - session_len - 1 - ct_buf_len - 1;
|
||||
|
||||
/**
|
||||
* verify some data
|
||||
*/
|
||||
|
||||
if (session_len != 8) return (PARSER_SALT_LENGTH);
|
||||
if (ct_buf_len >= 256) return (PARSER_SALT_LENGTH);
|
||||
if (sequence_len != 4) return (PARSER_SALT_LENGTH);
|
||||
|
||||
if (is_valid_hex_string (session_pos, session_len) == false) return (PARSER_SALT_ENCODING);
|
||||
if (is_valid_hex_string (ct_buf_pos, ct_buf_len) == false) return (PARSER_HASH_ENCODING);
|
||||
if (is_valid_hex_string (sequence_pos, sequence_len) == false) return (PARSER_HASH_ENCODING);
|
||||
|
||||
/**
|
||||
* store data
|
||||
*/
|
||||
|
||||
// ciphtertext
|
||||
|
||||
u8 *session_ptr = (u8 *) tacacs_plus->session_buf;
|
||||
|
||||
session_ptr[0] = hex_to_u8 ((const u8 *) session_pos + 0);
|
||||
session_ptr[1] = hex_to_u8 ((const u8 *) session_pos + 2);
|
||||
session_ptr[2] = hex_to_u8 ((const u8 *) session_pos + 4);
|
||||
session_ptr[3] = hex_to_u8 ((const u8 *) session_pos + 6);
|
||||
|
||||
u8 *sequence_ptr = (u8 *) tacacs_plus->sequence_buf;
|
||||
|
||||
sequence_ptr[0] = hex_to_u8 ((const u8 *) sequence_pos + 0);
|
||||
sequence_ptr[1] = hex_to_u8 ((const u8 *) sequence_pos + 2);
|
||||
|
||||
u8 *ct_data_ptr = (u8 *) tacacs_plus->ct_data_buf;
|
||||
|
||||
for (u32 i = 0, j = 0; j < ct_buf_len; i += 1, j += 2)
|
||||
{
|
||||
ct_data_ptr[i] = hex_to_u8 ((const u8 *) &ct_buf_pos[j]);
|
||||
|
||||
tacacs_plus->ct_data_len++;
|
||||
}
|
||||
|
||||
// fake salt
|
||||
|
||||
salt->salt_buf[0] = tacacs_plus->session_buf[0];
|
||||
salt->salt_buf[1] = tacacs_plus->sequence_buf[0];
|
||||
salt->salt_buf[2] = tacacs_plus->ct_data_buf[0];
|
||||
salt->salt_buf[3] = tacacs_plus->ct_data_buf[1];
|
||||
|
||||
salt->salt_len = 16;
|
||||
|
||||
// fake hash
|
||||
|
||||
digest[0] = tacacs_plus->ct_data_buf[2];
|
||||
digest[1] = tacacs_plus->ct_data_buf[3];
|
||||
digest[2] = tacacs_plus->ct_data_buf[4];
|
||||
digest[3] = tacacs_plus->ct_data_buf[5];
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
/**
|
||||
* hook functions
|
||||
*/
|
||||
@ -16154,6 +16257,7 @@ const char *strhashtype (const u32 hash_mode)
|
||||
case 15700: return HT_15700;
|
||||
case 15900: return HT_15900;
|
||||
case 16000: return HT_16000;
|
||||
case 16100: return HT_16100;
|
||||
case 99999: return HT_99999;
|
||||
}
|
||||
|
||||
@ -19559,6 +19663,34 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
|
||||
out_buf[10] = 0;
|
||||
}
|
||||
else if (hash_mode == 16100)
|
||||
{
|
||||
tacacs_plus_t *tacacs_pluss = (tacacs_plus_t *) esalts_buf;
|
||||
|
||||
tacacs_plus_t *tacacs_plus = &tacacs_pluss[digest_cur];
|
||||
|
||||
char ct_data[256 + 1] = { 0 };
|
||||
|
||||
u8 *ct_data_ptr = (u8 *) tacacs_plus->ct_data_buf;
|
||||
|
||||
for (u32 i = 0, j = 0; i < tacacs_plus->ct_data_len; i += 1, j += 2)
|
||||
{
|
||||
sprintf (ct_data + j, "%02x", ct_data_ptr[i]);
|
||||
}
|
||||
|
||||
u8 *session_ptr = (u8 *) tacacs_plus->session_buf;
|
||||
u8 *sequence_ptr = (u8 *) tacacs_plus->sequence_buf;
|
||||
|
||||
snprintf (out_buf, out_len - 1, "%s%02x%02x%02x%02x$%s$%02x%02x",
|
||||
SIGNATURE_TACACS_PLUS,
|
||||
session_ptr[0],
|
||||
session_ptr[1],
|
||||
session_ptr[2],
|
||||
session_ptr[3],
|
||||
ct_data,
|
||||
sequence_ptr[0],
|
||||
sequence_ptr[1]);
|
||||
}
|
||||
else if (hash_mode == 99999)
|
||||
{
|
||||
char *ptr = (char *) digest_buf;
|
||||
@ -24519,6 +24651,22 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN;
|
||||
break;
|
||||
|
||||
case 16100: hashconfig->hash_type = HASH_TYPE_TACACS_PLUS;
|
||||
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
|
||||
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
|
||||
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE;
|
||||
hashconfig->kern_type = KERN_TYPE_TACACS_PLUS;
|
||||
hashconfig->dgst_size = DGST_SIZE_4_4; // originally DGST_SIZE_4_2
|
||||
hashconfig->parse_func = tacacs_plus_parse_hash;
|
||||
hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE;
|
||||
hashconfig->dgst_pos0 = 0;
|
||||
hashconfig->dgst_pos1 = 1;
|
||||
hashconfig->dgst_pos2 = 2;
|
||||
hashconfig->dgst_pos3 = 3;
|
||||
hashconfig->st_hash = ST_HASH_16100;
|
||||
hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN;
|
||||
break;
|
||||
|
||||
case 99999: hashconfig->hash_type = HASH_TYPE_PLAINTEXT;
|
||||
hashconfig->salt_type = SALT_TYPE_NONE;
|
||||
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
|
||||
@ -24727,6 +24875,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
case 15600: hashconfig->esalt_size = sizeof (ethereum_pbkdf2_t); break;
|
||||
case 15700: hashconfig->esalt_size = sizeof (ethereum_scrypt_t); break;
|
||||
case 15900: hashconfig->esalt_size = sizeof (dpapimk_t); break;
|
||||
case 16100: hashconfig->esalt_size = sizeof (tacacs_plus_t); break;
|
||||
}
|
||||
|
||||
// hook_salt_size
|
||||
|
||||
Loading…
Reference in New Issue
Block a user