diff --git a/OpenCL/inc_hash_constants.h b/OpenCL/inc_hash_constants.h index 654a8159f..521793405 100644 --- a/OpenCL/inc_hash_constants.h +++ b/OpenCL/inc_hash_constants.h @@ -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, diff --git a/OpenCL/m06200_a3.cl b/OpenCL/m06200_a3.cl new file mode 100644 index 000000000..646124ac8 --- /dev/null +++ b/OpenCL/m06200_a3.cl @@ -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) +{ +} diff --git a/b2test.hash b/b2test.hash new file mode 100644 index 000000000..59600fd82 --- /dev/null +++ b/b2test.hash @@ -0,0 +1 @@ +b3910b0f4b6f1aede44da90bb7705a868b265861b36e6f7f29dba7223f6f1ce7b10e0dd25e47deb70bd7f3b24f7da653409cd9014f8715e4013c15fee38ab418 diff --git a/b2test.sh b/b2test.sh new file mode 100755 index 000000000..d87ac2c91 --- /dev/null +++ b/b2test.sh @@ -0,0 +1 @@ +./hashcat -m 6200 -a 3 ?d?d?d?d?d?d b2test.hash diff --git a/include/interface.h b/include/interface.h index e824aa34d..e557a40ed 100644 --- a/include/interface.h +++ b/include/interface.h @@ -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); diff --git a/src/interface.c b/src/interface.c index 65b99a602..7b099527f 100644 --- a/src/interface.c +++ b/src/interface.c @@ -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; diff --git a/src/usage.c b/src/usage.c index 407807216..963a16ee0 100644 --- a/src/usage.c +++ b/src/usage.c @@ -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",