arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-01-22 13:40:56 +00:00
Code Issues Releases Wiki Activity

init

Browse Source
This commit is contained in:
DoZ10 2017-04-13 08:38:59 -04:00
parent fab4ede364
commit 4e9bb8b093
7 changed files with 194 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
OpenCL/inc_hash_constants.h
View File

@ -621,6 +621,15 @@ typedef enum keccak_constants
} keccak_constants_t;
typedef enum blake2b_constants
{
BLAKE2B_BLOCKBYTES = 128,
BLAKE2B_OUTBYTES = 64,
BLAKE2B_KEYBYTES = 64,
BLAKE2B_SALTBYTES = 16,
BLAKE2B_PERSONALBYTES = 16
} blake2b_constants_t;
typedef enum mysql323_constants
{
MYSQL323_A=0x50305735,

113
OpenCL/m06200_a3.cl Normal file
View File

@ -0,0 +1,113 @@
/**
* 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_simd.cl"
#ifndef BLAKE2B_ROUNDS
#define BLAKE2B_ROUNDS
#endif
#define G(r,i,a,b,c,d) \
do { \
a = a + b + m[blake2b_sigma[r][2*i+0]]; \
d = rotr64(d ^ a, 32); \
c = c + d; \
b = rotr64(b ^ c, 24); \
a = a + b + m[blake2b_sigma[r][2*i+1]]; \
d = rotr64(d ^ a, 16); \
c = c + d; \
b = rotr64(b ^ c, 63); \
} while(0)
#define ROUND(r) \
do { \
G(r,0,v[ 0],v[ 4],v[ 8],v[12]); \
G(r,1,v[ 1],v[ 5],v[ 9],v[13]); \
G(r,2,v[ 2],v[ 6],v[10],v[14]); \
G(r,3,v[ 3],v[ 7],v[11],v[15]); \
G(r,4,v[ 0],v[ 5],v[10],v[15]); \
G(r,5,v[ 1],v[ 6],v[11],v[12]); \
G(r,6,v[ 2],v[ 7],v[ 8],v[13]); \
G(r,7,v[ 3],v[ 4],v[ 9],v[14]); \
} while(0)
__kernel void m06200_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 void *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 u32 gid_max)
{
}
__kernel void m05000_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 void *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 u32 gid_max)
{
}
__kernel void m05000_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 void *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 u32 gid_max)
{
}
__kernel void m06200_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 void *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 u32 gid_max)
{
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{
/*
* Compress function
*/
u64 m[16];
u64 v[16];
u8 i;
for( i = 0; i < 16; ++i ) {
m[i] = load64( block + i * sizeof( m[i] ) );
}
for( i = 0; i < 8; ++i ) {
v[i] = S->h[i];
}
v[ 8] = blake2b_IV[0];
v[ 9] = blake2b_IV[1];
v[10] = blake2b_IV[2];
v[11] = blake2b_IV[3];
v[12] = blake2b_IV[4] ^ S->t[0];
v[13] = blake2b_IV[5] ^ S->t[1];
v[14] = blake2b_IV[6] ^ S->f[0];
v[15] = blake2b_IV[7] ^ S->f[1];
ROUND( 0 );
ROUND( 1 );
ROUND( 2 );
ROUND( 3 );
ROUND( 4 );
ROUND( 5 );
ROUND( 6 );
ROUND( 7 );
ROUND( 8 );
ROUND( 9 );
ROUND( 10 );
ROUND( 11 );
for( i = 0; i < 8; ++i ) {
S->h[i] = S->h[i] ^ v[i] ^ v[i + 8];
}
}
}
__kernel void m05000_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 void *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 u32 gid_max)
{
}
__kernel void m05000_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 void *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 u32 gid_max)
{
}

1
b2test.hash Normal file
View File

@ -0,0 +1 @@
b3910b0f4b6f1aede44da90bb7705a868b265861b36e6f7f29dba7223f6f1ce7b10e0dd25e47deb70bd7f3b24f7da653409cd9014f8715e4013c15fee38ab418

1
b2test.sh Executable file
View File

@ -0,0 +1 @@
./hashcat -m 6200 -a 3 ?d?d?d?d?d?d b2test.hash

5
include/interface.h
View File

@ -1016,6 +1016,8 @@ typedef enum display_len
DISPLAY_LEN_MAX_6000 = 40,
DISPLAY_LEN_MIN_6100 = 128,
DISPLAY_LEN_MAX_6100 = 128,
DISPLAY_LEN_MIN_6200 = 128,
DISPLAY_LEN_MAX_6200 = 128,
DISPLAY_LEN_MIN_6300 = 6 + 1 + 8 + 22,
DISPLAY_LEN_MAX_6300 = 6 + 1 + 48 + 22,
DISPLAY_LEN_MIN_6400 = 9 + 2 + 1 + 16 + 1 + 43,
@ -1324,6 +1326,7 @@ typedef enum hash_type
HASH_TYPE_ITUNES_BACKUP_9 = 56,
HASH_TYPE_ITUNES_BACKUP_10 = 57,
HASH_TYPE_SKIP32 = 58,
HASH_TYPE_BLAKE2G = 59,
} hash_type_t;
@ -1403,6 +1406,7 @@ typedef enum kern_type
KERN_TYPE_ANDROIDPIN = 5800,
KERN_TYPE_RIPEMD160 = 6000,
KERN_TYPE_WHIRLPOOL = 6100,
KERN_TYPE_BLAKE2G = 6200,
KERN_TYPE_TCRIPEMD160_XTS512 = 6211,
KERN_TYPE_TCRIPEMD160_XTS1024 = 6212,
KERN_TYPE_TCRIPEMD160_XTS1536 = 6213,
@ -1601,6 +1605,7 @@ int joomla_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu
int postgresql_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int netscreen_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int keccak_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int blake2g_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int lm_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int md4_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);

64
src/interface.c
View File

@ -148,6 +148,7 @@ static const char HT_05700[] = "Cisco-IOS type 4 (SHA256)";
static const char HT_05800[] = "Samsung Android Password/PIN";
static const char HT_06000[] = "RIPEMD-160";
static const char HT_06100[] = "Whirlpool";
static const char HT_06200[] = "Blake2g";
static const char HT_06300[] = "AIX {smd5}";
static const char HT_06400[] = "AIX {ssha256}";
static const char HT_06500[] = "AIX {ssha512}";
@ -5260,6 +5261,30 @@ int keccak_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNU
return (PARSER_OK);
}
int blake2g_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
{
if ((input_len < DISPLAY_LEN_MIN_6200) || (input_len > DISPLAY_LEN_MAX_6200)) return (PARSE_GLOBAL_LENGTH);
if (input_len % 16) return (PARSER_GLOBAL_LENGTH);
if (is_valid_hex_string (input_buf, input_len) == false) return (PARSER_HASH_ENCODING);
u64 *digest = (u64 *) hash_buf->digest;
salt_t *salt = hash_buf->salt;
/*
u32 keccak_mdlen = input_len / 2;
for (u32 i = 0; i < keccak_mdlen / 8; i++)
{
digest[i] = hex_to_u64 ((const u8 *) &input_buf[i * 16]);
}
salt->keccak_mdlen = keccak_mdlen;
*/
return (PARSER_OK);
}
int ikepsk_md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
{
if ((input_len < DISPLAY_LEN_MIN_5300) || (input_len > DISPLAY_LEN_MAX_5300)) return (PARSER_GLOBAL_LENGTH);
@ -18314,6 +18339,26 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
out_buf[salt.keccak_mdlen * 2] = 0;
}
else if (hash_type == HASH_TYPE_BLAKE2G)
{
snprintf (out_buf, out_len - 1, "%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x",
digest_buf[ 0],
digest_buf[ 1],
digest_buf[ 2],
digest_buf[ 3],
digest_buf[ 4],
digest_buf[ 5],
digest_buf[ 6],
digest_buf[ 7],
digest_buf[ 8],
digest_buf[ 9],
digest_buf[10],
digest_buf[11],
digest_buf[12],
digest_buf[13],
digest_buf[14],
digest_buf[15],
}
else if (hash_type == HASH_TYPE_RIPEMD160)
{
snprintf (out_buf, out_len - 1, "%08x%08x%08x%08x%08x",
@ -20492,6 +20537,23 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->dgst_pos3 = 3;
break;
case 6200: hashconfig->hash_type = HASH_TYPE_BLAKE2G;
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_ADD01;
hashconfig->kern_type = KERN_TYPE_BLAKE2G;
hashconfig->dgst_size = DGST_SIZE_8_25;
hashconfig->parse_func = blake2g_parse_hash;
hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE
| OPTI_TYPE_USES_BITS_64
| OPTI_TYPE_RAW_HASH;
hashconfig->dgst_pos0 = 2;
hashconfig->dgst_pos1 = 3;
hashconfig->dgst_pos2 = 4;
hashconfig->dgst_pos3 = 5;
break;
case 6211: hashconfig->hash_type = HASH_TYPE_RIPEMD160;
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
@ -22937,6 +22999,8 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo
break;
case 5800: salt->salt_len = 16;
break;
case 6200: salt->salt_len = 16;
break;
case 6800: salt->salt_len = 32;
break;
case 8400: salt->salt_len = 40;

1
src/usage.c
View File

@ -123,6 +123,7 @@ static const char *USAGE_BIG[] =
" 10800 | SHA-384 | Raw Hash",
" 1700 | SHA-512 | Raw Hash",
" 5000 | SHA-3 (Keccak) | Raw Hash",
" 6200 | BLAKE2G | Raw Hash",
" 10100 | SipHash | Raw Hash",
" 6000 | RIPEMD-160 | Raw Hash",
" 6100 | Whirlpool | Raw Hash",