From cfc3fa64c04ae07f59e7f95c1632b49d69210da3 Mon Sep 17 00:00:00 2001 From: DoZ10 Date: Mon, 15 May 2017 08:47:40 -0400 Subject: [PATCH] Implemented offset parameter to reach full ks block of 64 bytes --- OpenCL/inc_types.cl | 1 + OpenCL/m15400_a0.cl | 49 +++++++++++++++++++++++++++++++++++++-------- include/interface.h | 4 ++-- include/types.h | 1 + src/Makefile | 2 +- src/interface.c | 17 ++++++++++++---- tools/test.pl | 8 +++++--- 7 files changed, 64 insertions(+), 18 deletions(-) diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index b2910d403..d49ed5174 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -784,6 +784,7 @@ typedef struct chacha20 u32 iv[2]; u32 plain[2]; u32 position[2]; + u32 offset; } chacha20_t; diff --git a/OpenCL/m15400_a0.cl b/OpenCL/m15400_a0.cl index c6a414f74..ead3984f3 100644 --- a/OpenCL/m15400_a0.cl +++ b/OpenCL/m15400_a0.cl @@ -31,7 +31,7 @@ x[b] = rotl32(x[b] ^ x[c], 7); \ } while (0); -void chacha20_transform (const u32x w0[4], const u32x w1[4], const u32 position[2], const u32 iv[2], const u32 plain[4], u32x digest[4]) +void chacha20_transform (const u32x w0[4], const u32x w1[4], const u32 position[2], const u32 offset, const u32 iv[2], const u32 plain[4], u32x digest[4]) { u32x ctx[16]; @@ -102,11 +102,36 @@ void chacha20_transform (const u32x w0[4], const u32x w1[4], const u32 position[ x[13] += ctx[13]; x[14] += ctx[14]; x[15] += ctx[15]; - - digest[1] = plain[0] ^ x[0]; - digest[0] = plain[1] ^ x[1]; - digest[3] = plain[2] ^ x[2]; - digest[2] = plain[3] ^ x[3]; + + + u32 index = offset / 4; + u32 remain = offset % 4; + + //printf("index: %d, offset: %d, remain: %d\n", index, offset, remain); + + digest[0] = plain[1]; + digest[1] = plain[0]; + + if (remain > 0) + { + u32x tmp[3]; + tmp[0] = x[index + 0]; + tmp[1] = x[index + 1]; + tmp[2] = x[index + 2]; + + digest[1] ^= tmp[0] >> (remain * 8); + digest[1] ^= tmp[1] << (32 - remain * 8); + + digest[0] ^= tmp[1] >> (remain * 8); + digest[0] ^= tmp[2] << (32 - remain * 8); + } + else + { + digest[1] ^= x[index + 0]; + digest[0] ^= x[index + 1]; + } + + //printf("digest[0]: %08x, x[0]: %08x, digest[1]: %08x, x[1]: %08x\n", digest[0], x[0], digest[1], x[1]); } __kernel void m15400_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 chacha20_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 u32 gid_max) @@ -139,10 +164,13 @@ __kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule u32 iv[2] = { 0 }; u32 plain[2] = { 0 }; u32 position[2] = { 0 }; + u32 offset = 0; position[0] = esalt_bufs->position[0]; position[1] = esalt_bufs->position[1]; + offset = esalt_bufs->offset; + iv[0] = esalt_bufs->iv[0]; iv[1] = esalt_bufs->iv[1]; @@ -164,7 +192,7 @@ __kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule u32x digest[4] = { 0 }; - chacha20_transform (w0, w1, position, iv, plain, digest); + chacha20_transform (w0, w1, position, offset, iv, plain, digest); const u32x r0 = digest[0]; const u32x r1 = digest[1]; @@ -216,10 +244,13 @@ __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule u32 iv[2] = { 0 }; u32 plain[2] = { 0 }; u32 position[2] = { 0 }; + u32 offset = 0; position[0] = esalt_bufs->position[0]; position[1] = esalt_bufs->position[1]; + offset = esalt_bufs->offset; + iv[0] = esalt_bufs->iv[0]; iv[1] = esalt_bufs->iv[1]; @@ -253,13 +284,15 @@ __kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule u32x digest[4] = { 0 }; - chacha20_transform (w0, w1, position, iv, plain, digest); + chacha20_transform (w0, w1, position, offset, iv, plain, digest); const u32x r0 = digest[0]; const u32x r1 = digest[1]; const u32x r2 = digest[2]; const u32x r3 = digest[3]; + // printf("r0: %08x, search[0]: %08x, r1: %08x, search[1]: %08x, r2: %08x, search[2]: %08x, r3: %08x, search[3]: %08x\n", r0, search[0], r1, search[1], r2, search[2], r3, search[3]); + COMPARE_S_SIMD(r0, r1, r2, r3); } } diff --git a/include/interface.h b/include/interface.h index 8c4250f35..81f87b80d 100644 --- a/include/interface.h +++ b/include/interface.h @@ -1190,8 +1190,8 @@ typedef enum display_len DISPLAY_LEN_MAX_15100 = 6 + 6 + 1 + 8 + 1 + 28, DISPLAY_LEN_MIN_15200 = 1 + 10 + 1 + 2 + 1 + 1 + 1 + 1 + 1 + 64, DISPLAY_LEN_MAX_15200 = 1 + 10 + 1 + 2 + 1 + 8 + 1 + 5 + 1 + 20000, - DISPLAY_LEN_MIN_15400 = 10 + 1 + 16 + 1 + 16 + 1 + 16 + 1 + 16, - DISPLAY_LEN_MAX_15400 = 10 + 1 + 16 + 1 + 16 + 1 + 16 + 1 + 16, + DISPLAY_LEN_MIN_15400 = 10 + 1 + 16 + 1 + 1 + 1 + 16 + 1 + 16 + 1 + 16, + DISPLAY_LEN_MAX_15400 = 10 + 1 + 16 + 1 + 2 + 1 + 16 + 1 + 16 + 1 + 16, DISPLAY_LEN_MIN_99999 = 1, DISPLAY_LEN_MAX_99999 = 55, diff --git a/include/types.h b/include/types.h index 1bba55b60..71f93a5a0 100644 --- a/include/types.h +++ b/include/types.h @@ -697,6 +697,7 @@ typedef struct u32 iv[2]; u32 plain[2]; u32 position[2]; + u32 offset; } chacha20_t; diff --git a/src/Makefile b/src/Makefile index 2b64c798e..5f61ee2c6 100644 --- a/src/Makefile +++ b/src/Makefile @@ -4,7 +4,7 @@ ## SHARED := 0 -DEBUG := 0 +DEBUG := 1 PRODUCTION := 0 ## diff --git a/src/interface.c b/src/interface.c index 44631d015..00536c152 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5326,7 +5326,13 @@ int chacha20_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_U if (position_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED); if (is_valid_hex_string (position_marker, 16) == false) return (PARSER_SALT_ENCODING); - u8 *iv_marker = (u8 *) strchr ((const char *) position_marker, '*') + 1; + u8 *offset_marker = (u8 *) strchr ((const char *) position_marker, '*') + 1; + if (offset_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED); + + int offset = atoi ((char*) offset_marker); + if (offset > 36) return (PARSER_SALT_VALUE); + + u8 *iv_marker = (u8 *) strchr ((const char *) offset_marker, '*') + 1; if (iv_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED); if (is_valid_hex_string (iv_marker, 16) == false) return (PARSER_SALT_ENCODING); @@ -5347,6 +5353,8 @@ int chacha20_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_U chacha20->position[0] = byte_swap_32(hex_to_u32 ((const u8 *) position_marker + 8)); chacha20->position[1] = byte_swap_32(hex_to_u32 ((const u8 *) position_marker + 0)); + chacha20->offset = offset; + /* some fake salt for the sorting mechanisms */ salt->salt_buf[0] = chacha20->iv[0]; @@ -5355,9 +5363,9 @@ int chacha20_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_U salt->salt_buf[3] = chacha20->plain[1]; salt->salt_buf[4] = chacha20->position[0]; salt->salt_buf[5] = chacha20->position[1]; - salt->salt_buf[6] = 0; + salt->salt_buf[6] = chacha20->offset; salt->salt_buf[7] = 0; - salt->salt_len = 24; + salt->salt_len = 32; /* Store cipher for search mechanism */ digest[0] = hex_to_u32 ((const u8 *) cipher_marker + 8); @@ -18556,10 +18564,11 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le const chacha20_t *chacha20 = (const chacha20_t *) esalts_buf; - snprintf (out_buf, out_len - 1, "%s*%08x%08x*%08x%08x*%08x%08x*%08x%08x", + snprintf (out_buf, out_len - 1, "%s*%08x%08x*%d*%08x%08x*%08x%08x*%08x%08x", SIGNATURE_CHACHA20, chacha20->position[1], chacha20->position[0], + chacha20->offset, byte_swap_32(chacha20->iv[1]), byte_swap_32(chacha20->iv[0]), byte_swap_32(chacha20->plain[1]), diff --git a/tools/test.pl b/tools/test.pl index 7ff4dd0ad..60988d121 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -8158,16 +8158,18 @@ END_CODE { my $eight_byte_iv = pack("H*", "0000000000000000"); my $eight_byte_counter = pack("H*", "0100000000000000"); # little endian 64 bits + my $offset = int(rand(36)); my $pad_len = 32 - length $word_buf; my $key = $word_buf . "\0" x $pad_len; my $cipher = Crypt::OpenSSH::ChachaPoly->new($key); $cipher->ivsetup($eight_byte_iv, $eight_byte_counter); - my $enc = $cipher->encrypt("AAAAAAAA"); - $hash_buf = $enc; + my $enc = $cipher->encrypt("AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA"); + my $enc_offset = substr($enc, $offset, 8); + $hash_buf = $enc_offset; - $tmp_hash = sprintf ("\$Chacha20\$\*%08x%08x\*0000000000000000\*4141414141414141\*%s", (unpack("V*", $eight_byte_counter))[1], (unpack("V*", $eight_byte_counter))[0], unpack("H*", $enc)); + $tmp_hash = sprintf ("\$Chacha20\$\*%08x%08x\*%d\*0000000000000000\*4141414141414141\*%s", (unpack("V*", $eight_byte_counter))[1], (unpack("V*", $eight_byte_counter))[0], $offset, unpack("H*", $enc_offset)); } elsif ($mode == 99999) {