From 3d0f1626ed9951282766f385c4c8516454427bce Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Tue, 2 Oct 2018 11:01:54 -0400 Subject: [PATCH 01/32] Add support for TOTP (RFC 6238) --- OpenCL/m17300_a3-pure.cl | 225 +++++++++++++++++++++++++++++++++++++++ include/interface.h | 1 + src/hashes.c | 15 +++ src/interface.c | 98 +++++++++++++++++ src/usage.c | 1 + 5 files changed, 340 insertions(+) create mode 100644 OpenCL/m17300_a3-pure.cl diff --git a/OpenCL/m17300_a3-pure.cl b/OpenCL/m17300_a3-pure.cl new file mode 100644 index 000000000..eba26ad84 --- /dev/null +++ b/OpenCL/m17300_a3-pure.cl @@ -0,0 +1,225 @@ +/** + * 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_simd.cl" +#include "inc_hash_sha1.cl" + +__kernel void m17300_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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 u64 gid_max) +{ + /** + * modifier + */ + + const u64 lid = get_local_id (0); + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * base + */ + + const u32 pw_len = pws[gid].pw_len; + + u32x w[64] = { 0 }; + + for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1) + { + w[idx] = pws[gid].i[idx]; + } + + const u32 salt_len = salt_bufs[salt_pos].salt_len; + + u32x s[64] = { 0 }; + + for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1) + { + s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); + } + + /** + * loop + */ + + u32x w0l = w[0]; + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) + { + const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + + const u32x w0 = w0l | w0r; + + w[0] = w0; + + sha1_hmac_ctx_vector_t ctx; + + sha1_hmac_init_vector (&ctx, w, pw_len); + + sha1_hmac_update_vector (&ctx, s, salt_len); + + sha1_hmac_final_vector (&ctx); + + // ------- PUT TOTP HERE ------- // + // calculate the offset using the least 4 bits of the last byte of our hash + const int otp_offset = ctx.opad.h[4] & 0xf; + + // initialize a buffer for the otp code + unsigned int otp_code = 0; + + // grab 4 consecutive bytes of the hash, starting at offset + switch(otp_offset%4) + { + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + break; + case 2: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + break; + case 3: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + break; + default: + otp_code = ctx.opad.h[otp_offset/4]; + break; + } + // take only the lower 31 bits + otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code + otp_code %= 1000000; + + const u32x r0 = ctx.opad.h[DGST_R0]; + const u32x r1 = ctx.opad.h[DGST_R1]; + const u32x r2 = ctx.opad.h[DGST_R2]; + const u32x r3 = ctx.opad.h[DGST_R3]; + + COMPARE_M_SIMD (otp_code, 0, 0, 0); + //COMPARE_M_SIMD (r0, r1, r2, r3); + } +} + +__kernel void m17300_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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 u64 gid_max) +{ + /** + * modifier + */ + + const u64 lid = get_local_id (0); + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * digest + */ + + const u32 search[4] = + { + digests_buf[digests_offset].digest_buf[DGST_R0], + digests_buf[digests_offset].digest_buf[DGST_R1], + digests_buf[digests_offset].digest_buf[DGST_R2], + digests_buf[digests_offset].digest_buf[DGST_R3] + }; + + /** + * base + */ + + const u32 pw_len = pws[gid].pw_len; + + u32x w[64] = { 0 }; + + for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1) + { + w[idx] = pws[gid].i[idx]; + } + + const u32 salt_len = salt_bufs[salt_pos].salt_len; + + u32x s[64] = { 0 }; + + for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1) + { + s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); + } + + /** + * loop + */ + + u32x w0l = w[0]; + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) + { + const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + + const u32x w0 = w0l | w0r; + + w[0] = w0; + + sha1_hmac_ctx_vector_t ctx; + + sha1_hmac_init_vector (&ctx, w, pw_len); + + sha1_hmac_update_vector (&ctx, s, salt_len); + + sha1_hmac_final_vector (&ctx); + //printf("%d ", sizeof(ctx.opad.h)); + + // ------- PUT TOTP HERE ------- // + // calculate the offset using the least 4 bits of the last byte of our hash + const int otp_offset = ctx.opad.h[4] & 0xf; + + // initialize a buffer for the otp code + unsigned int otp_code = 0; + + // grab 4 consecutive bytes of the hash, starting at offset + switch(otp_offset%4) + { + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + break; + case 2: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + break; + case 3: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + break; + default: + otp_code = ctx.opad.h[otp_offset/4]; + break; + } + // take only the lower 31 bits + otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code + otp_code %= 1000000; + + const u32x r0 = ctx.opad.h[DGST_R0]; + const u32x r1 = ctx.opad.h[DGST_R1]; + const u32x r2 = ctx.opad.h[DGST_R2]; + const u32x r3 = ctx.opad.h[DGST_R3]; + + if( ctx.opad.h[0] == 0xc085d274) + //if( ctx.opad.h[0] == 0x50184678) + //if( ctx.opad.h[0] == 0x8e664b2e) + { + //printf(" [[ %d %d %d %d ]]\n ", ctx.opad.h[otp_offset], ctx.opad.h[otp_offset +1], ctx.opad.h[otp_offset +2], ctx.opad.h[otp_offset +3]); + //printf(" MAXX:: %d %d\n", pw_len, salt_len); + //printf(" SRCH:: %x\n", search[0]); + //printf(" SRCH:: %x\n", search[1]); + //printf(" SRCH:: %x\n", search[2]); + //printf(" SRCH:: %x\n", search[3]); + //printf(" CODE:: %x\n", otp_code); + } + COMPARE_S_SIMD (otp_code, 0, 0, 0); + //COMPARE_S_SIMD (r0, r1, r2, r3); + } +} diff --git a/include/interface.h b/include/interface.h index aac00b010..7d11a37e1 100644 --- a/include/interface.h +++ b/include/interface.h @@ -1325,6 +1325,7 @@ typedef enum kern_type KERN_TYPE_WPA_PMKID_PBKDF2 = 16800, KERN_TYPE_WPA_PMKID_PMK = 16801, KERN_TYPE_ANSIBLE_VAULT = 16900, + KERN_TYPE_TOTP_HMACSHA1 = 17300, KERN_TYPE_PLAINTEXT = 99999, } kern_type_t; diff --git a/src/hashes.c b/src/hashes.c index 6ec376dea..7ade9faac 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -275,6 +275,21 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl build_plain (hashcat_ctx, device_param, plain, plain_buf, &plain_len); + // TOTP should be base32 encoded + if (hashcat_ctx->hashconfig->hash_mode == KERN_TYPE_TOTP_HMACSHA1) + { + // we need a temp buffer for the base32 encoding + u32 temp_buf[64] = { 0 }; + u8 *temp_ptr = (u8 *) temp_buf; + + // encode our plain + base32_encode (int_to_base32, (const u8 *) plain_ptr, plain_len, (u8 *) temp_buf); + plain_len = strlen((const char *) temp_buf); + + // copy the base32 content into our plain buffer + plain_ptr = (u8 *) strdup((const char *) temp_ptr); + } + // crackpos u64 crackpos = 0; diff --git a/src/interface.c b/src/interface.c index f00829970..1ad8cbeb2 100644 --- a/src/interface.c +++ b/src/interface.c @@ -281,6 +281,7 @@ static const char *ST_HASH_16700 = "$fvde$1$16$84286044060108438487434858307513$ static const char *ST_HASH_16800 = "2582a8281bf9d4308d6f5731d0e61c61*4604ba734d4e*89acf0e761f4*ed487162465a774bfba60eb603a39f3a"; static const char *ST_HASH_16801 = "2582a8281bf9d4308d6f5731d0e61c61*4604ba734d4e*89acf0e761f4"; static const char *ST_HASH_16900 = "$ansible$0*0*6b761adc6faeb0cc0bf197d3d4a4a7d3f1682e4b169cae8fa6b459b3214ed41e*426d313c5809d4a80a4b9bc7d4823070*d8bad190c7fbc7c3cb1c60a27abfb0ff59d6fb73178681c7454d94a0f56a4360"; +static const char *ST_HASH_17300 = "597056:3600"; static const char *ST_HASH_99999 = "hashcat"; static const char *OPTI_STR_OPTIMIZED_KERNEL = "Optimized-Kernel"; @@ -529,6 +530,7 @@ static const char *HT_16700 = "FileVault 2"; static const char *HT_16800 = "WPA-PMKID-PBKDF2"; static const char *HT_16801 = "WPA-PMKID-PMK"; static const char *HT_16900 = "Ansible Vault"; +static const char *HT_17300 = "TOTP (HMAC-SHA1)"; static const char *HT_99999 = "Plaintext"; static const char *HT_00011 = "Joomla < 2.5.18"; @@ -5142,6 +5144,74 @@ int sha1s_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUS return (PARSER_OK); } +int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig) +{ + // this is going to start off like HMAC-SHA1 + u32 *digest = (u32 *) hash_buf->digest; + + salt_t *salt = hash_buf->salt; + + token_t token; + + token.token_cnt = 2; + + token.sep[0] = hashconfig->separator; + token.len_min[0] = 6; + token.len_max[0] = 6; + token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.len_min[1] = SALT_MIN; + token.len_max[1] = SALT_MAX; + token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH; + + if (hashconfig->opts_type & OPTS_TYPE_ST_HEX) + { + token.len_min[1] *= 2; + token.len_max[1] *= 2; + + token.attr[1] |= TOKEN_ATTR_VERIFY_HEX; + } + + const int rc_tokenizer = input_tokenizer (input_buf, input_len, &token); + + // now we need to reduce our hash into a token + int otp_code = 0; + for (int i = 0; i < 6; i++) + { + otp_code = otp_code * 10 + itoa32_to_int(input_buf[i]); + } + if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); + + u8 *hash_pos = token.buf[0]; + + digest[1] = otp_code; + + u8 *salt_pos = token.buf[1]; + int salt_len = token.len[1]; + + // convert ascii timestamp to ulong timestamp + unsigned long timestamp = 0; + for(int i = 0; i>= 32) + { + salt->salt_buf[i] = byte_swap_32(timestamp); + } + + // our salt will always be 8 bytes + salt->salt_len = 8; + + return (PARSER_OK); +} + int pstoken_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig) { u32 *digest = (u32 *) hash_buf->digest; @@ -18387,6 +18457,7 @@ const char *strhashtype (const u32 hash_mode) case 16800: return HT_16800; case 16801: return HT_16801; case 16900: return HT_16900; + case 17300: return HT_17300; case 99999: return HT_99999; } @@ -22102,6 +22173,14 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le byte_swap_32 (digest_buf[6]), byte_swap_32 (digest_buf[7])); } + else if (hash_mode == 17300) + { + // salt_buf[1] holds our 32 bit value. salt_buf[0] and salt_buf[1] would be 64 bits. + // for now, we only need to worry about 32 bit counters. + // we also need to multiply salt by our step to see the floor of our original timestamp range. + // again, we will use the default RFC 6238 step of 30. + snprintf (out_buf, out_len - 1, "%d:%d", digest_buf[1], byte_swap_32(salt.salt_buf[1]) * 30); + } else if (hash_mode == 99999) { char *ptr = (char *) digest_buf; @@ -27299,6 +27378,25 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN; break; + case 17300: hashconfig->hash_type = HASH_TYPE_SHA1; + hashconfig->salt_type = SALT_TYPE_EMBEDDED; + hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL; + hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE + | OPTS_TYPE_PT_ADD80 + | OPTS_TYPE_PT_ADDBITS15; + hashconfig->kern_type = KERN_TYPE_TOTP_HMACSHA1; + hashconfig->dgst_size = DGST_SIZE_4_5; + hashconfig->parse_func = totp_parse_hash; + hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_NOT_ITERATED; + hashconfig->dgst_pos0 = 1; + hashconfig->dgst_pos1 = 2; + hashconfig->dgst_pos2 = 3; + hashconfig->dgst_pos3 = 4; + hashconfig->st_hash = ST_HASH_17300; + 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; diff --git a/src/usage.c b/src/usage.c index e272e903d..8f8eb0a12 100644 --- a/src/usage.c +++ b/src/usage.c @@ -170,6 +170,7 @@ static const char *const USAGE_BIG[] = " 60 | HMAC-MD5 (key = $salt) | Raw Hash, Authenticated", " 150 | HMAC-SHA1 (key = $pass) | Raw Hash, Authenticated", " 160 | HMAC-SHA1 (key = $salt) | Raw Hash, Authenticated", + " 17300 | TOTP (HMAC-SHA1) | Raw Hash, Authenticated", " 1450 | HMAC-SHA256 (key = $pass) | Raw Hash, Authenticated", " 1460 | HMAC-SHA256 (key = $salt) | Raw Hash, Authenticated", " 1750 | HMAC-SHA512 (key = $pass) | Raw Hash, Authenticated", From 977b560bb45bf6bf3e8124b364a45429cd6d9dd4 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Tue, 2 Oct 2018 11:01:54 -0400 Subject: [PATCH 02/32] Add support for TOTP (RFC 6238) --- OpenCL/m17300_a3-pure.cl | 225 +++++++++++++++++++++++++++++++++++++++ include/interface.h | 1 + src/hashes.c | 15 +++ src/interface.c | 98 +++++++++++++++++ src/usage.c | 1 + 5 files changed, 340 insertions(+) create mode 100644 OpenCL/m17300_a3-pure.cl diff --git a/OpenCL/m17300_a3-pure.cl b/OpenCL/m17300_a3-pure.cl new file mode 100644 index 000000000..eba26ad84 --- /dev/null +++ b/OpenCL/m17300_a3-pure.cl @@ -0,0 +1,225 @@ +/** + * 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_simd.cl" +#include "inc_hash_sha1.cl" + +__kernel void m17300_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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 u64 gid_max) +{ + /** + * modifier + */ + + const u64 lid = get_local_id (0); + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * base + */ + + const u32 pw_len = pws[gid].pw_len; + + u32x w[64] = { 0 }; + + for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1) + { + w[idx] = pws[gid].i[idx]; + } + + const u32 salt_len = salt_bufs[salt_pos].salt_len; + + u32x s[64] = { 0 }; + + for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1) + { + s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); + } + + /** + * loop + */ + + u32x w0l = w[0]; + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) + { + const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + + const u32x w0 = w0l | w0r; + + w[0] = w0; + + sha1_hmac_ctx_vector_t ctx; + + sha1_hmac_init_vector (&ctx, w, pw_len); + + sha1_hmac_update_vector (&ctx, s, salt_len); + + sha1_hmac_final_vector (&ctx); + + // ------- PUT TOTP HERE ------- // + // calculate the offset using the least 4 bits of the last byte of our hash + const int otp_offset = ctx.opad.h[4] & 0xf; + + // initialize a buffer for the otp code + unsigned int otp_code = 0; + + // grab 4 consecutive bytes of the hash, starting at offset + switch(otp_offset%4) + { + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + break; + case 2: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + break; + case 3: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + break; + default: + otp_code = ctx.opad.h[otp_offset/4]; + break; + } + // take only the lower 31 bits + otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code + otp_code %= 1000000; + + const u32x r0 = ctx.opad.h[DGST_R0]; + const u32x r1 = ctx.opad.h[DGST_R1]; + const u32x r2 = ctx.opad.h[DGST_R2]; + const u32x r3 = ctx.opad.h[DGST_R3]; + + COMPARE_M_SIMD (otp_code, 0, 0, 0); + //COMPARE_M_SIMD (r0, r1, r2, r3); + } +} + +__kernel void m17300_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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 u64 gid_max) +{ + /** + * modifier + */ + + const u64 lid = get_local_id (0); + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * digest + */ + + const u32 search[4] = + { + digests_buf[digests_offset].digest_buf[DGST_R0], + digests_buf[digests_offset].digest_buf[DGST_R1], + digests_buf[digests_offset].digest_buf[DGST_R2], + digests_buf[digests_offset].digest_buf[DGST_R3] + }; + + /** + * base + */ + + const u32 pw_len = pws[gid].pw_len; + + u32x w[64] = { 0 }; + + for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1) + { + w[idx] = pws[gid].i[idx]; + } + + const u32 salt_len = salt_bufs[salt_pos].salt_len; + + u32x s[64] = { 0 }; + + for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1) + { + s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); + } + + /** + * loop + */ + + u32x w0l = w[0]; + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) + { + const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + + const u32x w0 = w0l | w0r; + + w[0] = w0; + + sha1_hmac_ctx_vector_t ctx; + + sha1_hmac_init_vector (&ctx, w, pw_len); + + sha1_hmac_update_vector (&ctx, s, salt_len); + + sha1_hmac_final_vector (&ctx); + //printf("%d ", sizeof(ctx.opad.h)); + + // ------- PUT TOTP HERE ------- // + // calculate the offset using the least 4 bits of the last byte of our hash + const int otp_offset = ctx.opad.h[4] & 0xf; + + // initialize a buffer for the otp code + unsigned int otp_code = 0; + + // grab 4 consecutive bytes of the hash, starting at offset + switch(otp_offset%4) + { + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + break; + case 2: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + break; + case 3: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + break; + default: + otp_code = ctx.opad.h[otp_offset/4]; + break; + } + // take only the lower 31 bits + otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code + otp_code %= 1000000; + + const u32x r0 = ctx.opad.h[DGST_R0]; + const u32x r1 = ctx.opad.h[DGST_R1]; + const u32x r2 = ctx.opad.h[DGST_R2]; + const u32x r3 = ctx.opad.h[DGST_R3]; + + if( ctx.opad.h[0] == 0xc085d274) + //if( ctx.opad.h[0] == 0x50184678) + //if( ctx.opad.h[0] == 0x8e664b2e) + { + //printf(" [[ %d %d %d %d ]]\n ", ctx.opad.h[otp_offset], ctx.opad.h[otp_offset +1], ctx.opad.h[otp_offset +2], ctx.opad.h[otp_offset +3]); + //printf(" MAXX:: %d %d\n", pw_len, salt_len); + //printf(" SRCH:: %x\n", search[0]); + //printf(" SRCH:: %x\n", search[1]); + //printf(" SRCH:: %x\n", search[2]); + //printf(" SRCH:: %x\n", search[3]); + //printf(" CODE:: %x\n", otp_code); + } + COMPARE_S_SIMD (otp_code, 0, 0, 0); + //COMPARE_S_SIMD (r0, r1, r2, r3); + } +} diff --git a/include/interface.h b/include/interface.h index aac00b010..7d11a37e1 100644 --- a/include/interface.h +++ b/include/interface.h @@ -1325,6 +1325,7 @@ typedef enum kern_type KERN_TYPE_WPA_PMKID_PBKDF2 = 16800, KERN_TYPE_WPA_PMKID_PMK = 16801, KERN_TYPE_ANSIBLE_VAULT = 16900, + KERN_TYPE_TOTP_HMACSHA1 = 17300, KERN_TYPE_PLAINTEXT = 99999, } kern_type_t; diff --git a/src/hashes.c b/src/hashes.c index 6ec376dea..7ade9faac 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -275,6 +275,21 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl build_plain (hashcat_ctx, device_param, plain, plain_buf, &plain_len); + // TOTP should be base32 encoded + if (hashcat_ctx->hashconfig->hash_mode == KERN_TYPE_TOTP_HMACSHA1) + { + // we need a temp buffer for the base32 encoding + u32 temp_buf[64] = { 0 }; + u8 *temp_ptr = (u8 *) temp_buf; + + // encode our plain + base32_encode (int_to_base32, (const u8 *) plain_ptr, plain_len, (u8 *) temp_buf); + plain_len = strlen((const char *) temp_buf); + + // copy the base32 content into our plain buffer + plain_ptr = (u8 *) strdup((const char *) temp_ptr); + } + // crackpos u64 crackpos = 0; diff --git a/src/interface.c b/src/interface.c index f00829970..1ad8cbeb2 100644 --- a/src/interface.c +++ b/src/interface.c @@ -281,6 +281,7 @@ static const char *ST_HASH_16700 = "$fvde$1$16$84286044060108438487434858307513$ static const char *ST_HASH_16800 = "2582a8281bf9d4308d6f5731d0e61c61*4604ba734d4e*89acf0e761f4*ed487162465a774bfba60eb603a39f3a"; static const char *ST_HASH_16801 = "2582a8281bf9d4308d6f5731d0e61c61*4604ba734d4e*89acf0e761f4"; static const char *ST_HASH_16900 = "$ansible$0*0*6b761adc6faeb0cc0bf197d3d4a4a7d3f1682e4b169cae8fa6b459b3214ed41e*426d313c5809d4a80a4b9bc7d4823070*d8bad190c7fbc7c3cb1c60a27abfb0ff59d6fb73178681c7454d94a0f56a4360"; +static const char *ST_HASH_17300 = "597056:3600"; static const char *ST_HASH_99999 = "hashcat"; static const char *OPTI_STR_OPTIMIZED_KERNEL = "Optimized-Kernel"; @@ -529,6 +530,7 @@ static const char *HT_16700 = "FileVault 2"; static const char *HT_16800 = "WPA-PMKID-PBKDF2"; static const char *HT_16801 = "WPA-PMKID-PMK"; static const char *HT_16900 = "Ansible Vault"; +static const char *HT_17300 = "TOTP (HMAC-SHA1)"; static const char *HT_99999 = "Plaintext"; static const char *HT_00011 = "Joomla < 2.5.18"; @@ -5142,6 +5144,74 @@ int sha1s_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUS return (PARSER_OK); } +int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig) +{ + // this is going to start off like HMAC-SHA1 + u32 *digest = (u32 *) hash_buf->digest; + + salt_t *salt = hash_buf->salt; + + token_t token; + + token.token_cnt = 2; + + token.sep[0] = hashconfig->separator; + token.len_min[0] = 6; + token.len_max[0] = 6; + token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.len_min[1] = SALT_MIN; + token.len_max[1] = SALT_MAX; + token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH; + + if (hashconfig->opts_type & OPTS_TYPE_ST_HEX) + { + token.len_min[1] *= 2; + token.len_max[1] *= 2; + + token.attr[1] |= TOKEN_ATTR_VERIFY_HEX; + } + + const int rc_tokenizer = input_tokenizer (input_buf, input_len, &token); + + // now we need to reduce our hash into a token + int otp_code = 0; + for (int i = 0; i < 6; i++) + { + otp_code = otp_code * 10 + itoa32_to_int(input_buf[i]); + } + if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); + + u8 *hash_pos = token.buf[0]; + + digest[1] = otp_code; + + u8 *salt_pos = token.buf[1]; + int salt_len = token.len[1]; + + // convert ascii timestamp to ulong timestamp + unsigned long timestamp = 0; + for(int i = 0; i>= 32) + { + salt->salt_buf[i] = byte_swap_32(timestamp); + } + + // our salt will always be 8 bytes + salt->salt_len = 8; + + return (PARSER_OK); +} + int pstoken_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig) { u32 *digest = (u32 *) hash_buf->digest; @@ -18387,6 +18457,7 @@ const char *strhashtype (const u32 hash_mode) case 16800: return HT_16800; case 16801: return HT_16801; case 16900: return HT_16900; + case 17300: return HT_17300; case 99999: return HT_99999; } @@ -22102,6 +22173,14 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le byte_swap_32 (digest_buf[6]), byte_swap_32 (digest_buf[7])); } + else if (hash_mode == 17300) + { + // salt_buf[1] holds our 32 bit value. salt_buf[0] and salt_buf[1] would be 64 bits. + // for now, we only need to worry about 32 bit counters. + // we also need to multiply salt by our step to see the floor of our original timestamp range. + // again, we will use the default RFC 6238 step of 30. + snprintf (out_buf, out_len - 1, "%d:%d", digest_buf[1], byte_swap_32(salt.salt_buf[1]) * 30); + } else if (hash_mode == 99999) { char *ptr = (char *) digest_buf; @@ -27299,6 +27378,25 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN; break; + case 17300: hashconfig->hash_type = HASH_TYPE_SHA1; + hashconfig->salt_type = SALT_TYPE_EMBEDDED; + hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL; + hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE + | OPTS_TYPE_PT_ADD80 + | OPTS_TYPE_PT_ADDBITS15; + hashconfig->kern_type = KERN_TYPE_TOTP_HMACSHA1; + hashconfig->dgst_size = DGST_SIZE_4_5; + hashconfig->parse_func = totp_parse_hash; + hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_NOT_ITERATED; + hashconfig->dgst_pos0 = 1; + hashconfig->dgst_pos1 = 2; + hashconfig->dgst_pos2 = 3; + hashconfig->dgst_pos3 = 4; + hashconfig->st_hash = ST_HASH_17300; + 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; diff --git a/src/usage.c b/src/usage.c index e272e903d..8f8eb0a12 100644 --- a/src/usage.c +++ b/src/usage.c @@ -170,6 +170,7 @@ static const char *const USAGE_BIG[] = " 60 | HMAC-MD5 (key = $salt) | Raw Hash, Authenticated", " 150 | HMAC-SHA1 (key = $pass) | Raw Hash, Authenticated", " 160 | HMAC-SHA1 (key = $salt) | Raw Hash, Authenticated", + " 17300 | TOTP (HMAC-SHA1) | Raw Hash, Authenticated", " 1450 | HMAC-SHA256 (key = $pass) | Raw Hash, Authenticated", " 1460 | HMAC-SHA256 (key = $salt) | Raw Hash, Authenticated", " 1750 | HMAC-SHA512 (key = $pass) | Raw Hash, Authenticated", From 2249ab4c13bb8d43539dd5f6c3bd8eda50fb4fb9 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Tue, 16 Oct 2018 14:53:02 -0400 Subject: [PATCH 03/32] Cleanup debug code --- OpenCL/m17300_a3-pure.cl | 17 ----------------- 1 file changed, 17 deletions(-) diff --git a/OpenCL/m17300_a3-pure.cl b/OpenCL/m17300_a3-pure.cl index eba26ad84..97cab3e76 100644 --- a/OpenCL/m17300_a3-pure.cl +++ b/OpenCL/m17300_a3-pure.cl @@ -68,7 +68,6 @@ __kernel void m17300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final_vector (&ctx); - // ------- PUT TOTP HERE ------- // // calculate the offset using the least 4 bits of the last byte of our hash const int otp_offset = ctx.opad.h[4] & 0xf; @@ -102,7 +101,6 @@ __kernel void m17300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule const u32x r3 = ctx.opad.h[DGST_R3]; COMPARE_M_SIMD (otp_code, 0, 0, 0); - //COMPARE_M_SIMD (r0, r1, r2, r3); } } @@ -172,9 +170,7 @@ __kernel void m17300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_update_vector (&ctx, s, salt_len); sha1_hmac_final_vector (&ctx); - //printf("%d ", sizeof(ctx.opad.h)); - // ------- PUT TOTP HERE ------- // // calculate the offset using the least 4 bits of the last byte of our hash const int otp_offset = ctx.opad.h[4] & 0xf; @@ -207,19 +203,6 @@ __kernel void m17300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule const u32x r2 = ctx.opad.h[DGST_R2]; const u32x r3 = ctx.opad.h[DGST_R3]; - if( ctx.opad.h[0] == 0xc085d274) - //if( ctx.opad.h[0] == 0x50184678) - //if( ctx.opad.h[0] == 0x8e664b2e) - { - //printf(" [[ %d %d %d %d ]]\n ", ctx.opad.h[otp_offset], ctx.opad.h[otp_offset +1], ctx.opad.h[otp_offset +2], ctx.opad.h[otp_offset +3]); - //printf(" MAXX:: %d %d\n", pw_len, salt_len); - //printf(" SRCH:: %x\n", search[0]); - //printf(" SRCH:: %x\n", search[1]); - //printf(" SRCH:: %x\n", search[2]); - //printf(" SRCH:: %x\n", search[3]); - //printf(" CODE:: %x\n", otp_code); - } COMPARE_S_SIMD (otp_code, 0, 0, 0); - //COMPARE_S_SIMD (r0, r1, r2, r3); } } From 1ecc5e55595948e7a404e3c513d2746980629a21 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Tue, 16 Oct 2018 15:01:58 -0400 Subject: [PATCH 04/32] Add TOTP to test.pl --- tools/test.pl | 31 +++++++++++++++++++++++++++---- 1 file changed, 27 insertions(+), 4 deletions(-) diff --git a/tools/test.pl b/tools/test.pl index c58145060..852e3159d 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -40,6 +40,7 @@ use Crypt::Skip32; use Crypt::OpenSSH::ChachaPoly; use JSON; use MIME::Base64 qw (encode_base64 decode_base64 encode_base64url decode_base64url); +use MIME::Base32 qw (encode_base32 decode_base32); use Authen::Passphrase::NTHash; use Authen::Passphrase::MySQL323; use Authen::Passphrase::PHPass; @@ -57,7 +58,7 @@ my $hashcat = "./hashcat"; my $MAX_LEN = 55; -my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7701, 7800, 7801, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 15400, 15500, 15600, 15700, 15900, 16000, 16100, 16200, 16300, 16400, 16500, 16600, 16700, 16800, 16900, 99999); +my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7701, 7800, 7801, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 15400, 15500, 15600, 15700, 15900, 16000, 16100, 16200, 16300, 16400, 16500, 16600, 16700, 16800, 16900, 17300, 99999); my %is_utf16le = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800); my %less_fifteen = map { $_ => 1 } qw (500 1600 1800 3200 6300 7400 10500 10700); @@ -238,7 +239,7 @@ sub verify $word = substr ($line, $index + 1); } # hash:salt - elsif ($mode == 10 || $mode == 11 || $mode == 12 || $mode == 20 || $mode == 21 || $mode == 22 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 112 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1100 || $mode == 1410 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 2410 || $mode == 2611 || $mode == 2711 || $mode == 2811 || $mode == 3100 || $mode == 3610 || $mode == 3710 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4520 || $mode == 4521 || $mode == 4522 || $mode == 4900 || $mode == 5800 || $mode == 8400 || $mode == 11000 || $mode == 12600 || $mode == 13500 || $mode == 13800 || $mode == 13900 || $mode == 14000 || $mode == 14100 || $mode == 14400 || $mode == 14900 || $mode == 15000) + elsif ($mode == 10 || $mode == 11 || $mode == 12 || $mode == 20 || $mode == 21 || $mode == 22 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 112 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1100 || $mode == 1410 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 2410 || $mode == 2611 || $mode == 2711 || $mode == 2811 || $mode == 3100 || $mode == 3610 || $mode == 3710 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4520 || $mode == 4521 || $mode == 4522 || $mode == 4900 || $mode == 5800 || $mode == 8400 || $mode == 11000 || $mode == 12600 || $mode == 13500 || $mode == 13800 || $mode == 13900 || $mode == 14000 || $mode == 14100 || $mode == 14400 || $mode == 14900 || $mode == 15000 || $mode == 17300) { # get hash my $index1 = index ($line, ":"); @@ -3550,7 +3551,7 @@ sub passthrough { $tmp_hash = gen_hash ($mode, $word_buf, ""); } - elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4900 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100) + elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4900 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 17300) { my $salt_len = get_random_num (1, 15); @@ -4090,7 +4091,7 @@ sub single } } } - elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 16500) + elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 16500 || $mode == 17300) { my $salt_len = get_random_num (1, 15); @@ -9928,6 +9929,28 @@ END_CODE $tmp_hash = sprintf ("%s*%s*%s*%s", substr ($pmkid, 0, 32), $macap, $macsta, $essid); } + elsif ($mode == 17300) + { + my $paddedTime = sprintf("%016x", int(int($salt_buf) / 30)); + my $data = pack('H*', $paddedTime); + my $key = $word_buf; + + #my $b32_salt_buf = encode_base32($salt_buf); + #print "SECRET: ". $b32_salt_buf . ('=' x (8 - (length($b32_salt_buf) % 8))) ." "; + $hash_buf = hmac_hex ($data, $key, \&sha1, 64); + #$hash_buf = hmac ($word_buf, $salt_buf, \&sha1, 64); + + my $offset = hex (substr ($hash_buf, -8)) & 0xf; + $offset *= 2; + my $token = hex (substr ($hash_buf, $offset, 8)); + $token &= 0x7fffffff; + $token %= 1000000; + #print "CODE: " . $token . "\n"; + + #$tmp_hash = sprintf ("%s:%s", $hash_buf, $salt_buf); + ## token must be leading zero padded, and salt leading zero stripped + $tmp_hash = sprintf ("%06d:%d", $token, int($salt_buf)); + } elsif ($mode == 99999) { $tmp_hash = sprintf ("%s", $word_buf); From 73aae1a734a80660679b67b255be107e90200482 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Tue, 16 Oct 2018 15:07:41 -0400 Subject: [PATCH 05/32] Add a0 kernel for TOTP --- OpenCL/m17300_a0-pure.cl | 188 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 188 insertions(+) create mode 100644 OpenCL/m17300_a0-pure.cl diff --git a/OpenCL/m17300_a0-pure.cl b/OpenCL/m17300_a0-pure.cl new file mode 100644 index 000000000..1dbaa4e54 --- /dev/null +++ b/OpenCL/m17300_a0-pure.cl @@ -0,0 +1,188 @@ +/** + * 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_sha1.cl" + +__kernel void m17300_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 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 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]); + + const u32 salt_len = salt_bufs[salt_pos].salt_len; + + u32 s[64] = { 0 }; + + for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1) + { + s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); + } + + /** + * 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); + + sha1_hmac_ctx_t ctx; + + sha1_hmac_init_swap (&ctx, tmp.i, tmp.pw_len); + + sha1_hmac_update (&ctx, s, salt_len); + + sha1_hmac_final (&ctx); + + // calculate the offset using the least 4 bits of the last byte of our hash + const int otp_offset = ctx.opad.h[4] & 0xf; + + // initialize a buffer for the otp code + unsigned int otp_code = 0; + + // grab 4 consecutive bytes of the hash, starting at offset + switch(otp_offset%4) + { + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + break; + case 2: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + break; + case 3: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + break; + default: + otp_code = ctx.opad.h[otp_offset/4]; + break; + } + // take only the lower 31 bits + otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code + otp_code %= 1000000; + + const u32 r0 = ctx.opad.h[DGST_R0]; + const u32 r1 = ctx.opad.h[DGST_R1]; + const u32 r2 = ctx.opad.h[DGST_R2]; + const u32 r3 = ctx.opad.h[DGST_R3]; + + COMPARE_M_SCALAR (otp_code, 0, 0, 0); + } +} + +__kernel void m17300_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 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 u64 gid_max) +{ + /** + * modifier + */ + + const u64 lid = get_local_id (0); + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * digest + */ + + const u32 search[4] = + { + digests_buf[digests_offset].digest_buf[DGST_R0], + digests_buf[digests_offset].digest_buf[DGST_R1], + digests_buf[digests_offset].digest_buf[DGST_R2], + digests_buf[digests_offset].digest_buf[DGST_R3] + }; + + /** + * base + */ + + COPY_PW (pws[gid]); + + const u32 salt_len = salt_bufs[salt_pos].salt_len; + + u32 s[64] = { 0 }; + + for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1) + { + s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); + } + + /** + * 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); + + sha1_hmac_ctx_t ctx; + + sha1_hmac_init_swap (&ctx, tmp.i, tmp.pw_len); + + sha1_hmac_update (&ctx, s, salt_len); + + sha1_hmac_final (&ctx); + + // calculate the offset using the least 4 bits of the last byte of our hash + const int otp_offset = ctx.opad.h[4] & 0xf; + + // initialize a buffer for the otp code + unsigned int otp_code = 0; + + // grab 4 consecutive bytes of the hash, starting at offset + switch(otp_offset%4) + { + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + break; + case 2: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + break; + case 3: + otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + break; + default: + otp_code = ctx.opad.h[otp_offset/4]; + break; + } + // take only the lower 31 bits + otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code + otp_code %= 1000000; + + const u32 r0 = ctx.opad.h[DGST_R0]; + const u32 r1 = ctx.opad.h[DGST_R1]; + const u32 r2 = ctx.opad.h[DGST_R2]; + const u32 r3 = ctx.opad.h[DGST_R3]; + + COMPARE_S_SCALAR (otp_code, 0, 0, 0); + } +} From 6cda8f707796542cb8ba19aa359d4b23e6d8e48d Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Tue, 16 Oct 2018 15:33:09 -0400 Subject: [PATCH 06/32] Change TOTP index from 17300 to 18100 --- OpenCL/{m17300_a0-pure.cl => m18100_a0-pure.cl} | 4 ++-- OpenCL/{m17300_a3-pure.cl => m18100_a3-pure.cl} | 4 ++-- include/interface.h | 2 +- src/interface.c | 12 ++++++------ src/usage.c | 2 +- tools/test.pl | 10 +++++----- 6 files changed, 17 insertions(+), 17 deletions(-) rename OpenCL/{m17300_a0-pure.cl => m18100_a0-pure.cl} (97%) rename OpenCL/{m17300_a3-pure.cl => m18100_a3-pure.cl} (98%) diff --git a/OpenCL/m17300_a0-pure.cl b/OpenCL/m18100_a0-pure.cl similarity index 97% rename from OpenCL/m17300_a0-pure.cl rename to OpenCL/m18100_a0-pure.cl index 1dbaa4e54..4de835280 100644 --- a/OpenCL/m17300_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -15,7 +15,7 @@ #include "inc_scalar.cl" #include "inc_hash_sha1.cl" -__kernel void m17300_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 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 u64 gid_max) +__kernel void m18100_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 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 u64 gid_max) { /** * modifier @@ -95,7 +95,7 @@ __kernel void m17300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru } } -__kernel void m17300_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 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 u64 gid_max) +__kernel void m18100_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 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 u64 gid_max) { /** * modifier diff --git a/OpenCL/m17300_a3-pure.cl b/OpenCL/m18100_a3-pure.cl similarity index 98% rename from OpenCL/m17300_a3-pure.cl rename to OpenCL/m18100_a3-pure.cl index 97cab3e76..a106acb99 100644 --- a/OpenCL/m17300_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -13,7 +13,7 @@ #include "inc_simd.cl" #include "inc_hash_sha1.cl" -__kernel void m17300_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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 u64 gid_max) +__kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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 u64 gid_max) { /** * modifier @@ -104,7 +104,7 @@ __kernel void m17300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule } } -__kernel void m17300_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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 u64 gid_max) +__kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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 u64 gid_max) { /** * modifier diff --git a/include/interface.h b/include/interface.h index 7d11a37e1..92f1ffd07 100644 --- a/include/interface.h +++ b/include/interface.h @@ -1325,7 +1325,7 @@ typedef enum kern_type KERN_TYPE_WPA_PMKID_PBKDF2 = 16800, KERN_TYPE_WPA_PMKID_PMK = 16801, KERN_TYPE_ANSIBLE_VAULT = 16900, - KERN_TYPE_TOTP_HMACSHA1 = 17300, + KERN_TYPE_TOTP_HMACSHA1 = 18100, KERN_TYPE_PLAINTEXT = 99999, } kern_type_t; diff --git a/src/interface.c b/src/interface.c index 1ad8cbeb2..40e4a2b2b 100644 --- a/src/interface.c +++ b/src/interface.c @@ -281,7 +281,7 @@ static const char *ST_HASH_16700 = "$fvde$1$16$84286044060108438487434858307513$ static const char *ST_HASH_16800 = "2582a8281bf9d4308d6f5731d0e61c61*4604ba734d4e*89acf0e761f4*ed487162465a774bfba60eb603a39f3a"; static const char *ST_HASH_16801 = "2582a8281bf9d4308d6f5731d0e61c61*4604ba734d4e*89acf0e761f4"; static const char *ST_HASH_16900 = "$ansible$0*0*6b761adc6faeb0cc0bf197d3d4a4a7d3f1682e4b169cae8fa6b459b3214ed41e*426d313c5809d4a80a4b9bc7d4823070*d8bad190c7fbc7c3cb1c60a27abfb0ff59d6fb73178681c7454d94a0f56a4360"; -static const char *ST_HASH_17300 = "597056:3600"; +static const char *ST_HASH_18100 = "597056:3600"; static const char *ST_HASH_99999 = "hashcat"; static const char *OPTI_STR_OPTIMIZED_KERNEL = "Optimized-Kernel"; @@ -530,7 +530,7 @@ static const char *HT_16700 = "FileVault 2"; static const char *HT_16800 = "WPA-PMKID-PBKDF2"; static const char *HT_16801 = "WPA-PMKID-PMK"; static const char *HT_16900 = "Ansible Vault"; -static const char *HT_17300 = "TOTP (HMAC-SHA1)"; +static const char *HT_18100 = "TOTP (HMAC-SHA1)"; static const char *HT_99999 = "Plaintext"; static const char *HT_00011 = "Joomla < 2.5.18"; @@ -18457,7 +18457,7 @@ const char *strhashtype (const u32 hash_mode) case 16800: return HT_16800; case 16801: return HT_16801; case 16900: return HT_16900; - case 17300: return HT_17300; + case 18100: return HT_18100; case 99999: return HT_99999; } @@ -22173,7 +22173,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le byte_swap_32 (digest_buf[6]), byte_swap_32 (digest_buf[7])); } - else if (hash_mode == 17300) + else if (hash_mode == 18100) { // salt_buf[1] holds our 32 bit value. salt_buf[0] and salt_buf[1] would be 64 bits. // for now, we only need to worry about 32 bit counters. @@ -27378,7 +27378,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN; break; - case 17300: hashconfig->hash_type = HASH_TYPE_SHA1; + case 18100: hashconfig->hash_type = HASH_TYPE_SHA1; hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL; hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE @@ -27393,7 +27393,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->dgst_pos1 = 2; hashconfig->dgst_pos2 = 3; hashconfig->dgst_pos3 = 4; - hashconfig->st_hash = ST_HASH_17300; + hashconfig->st_hash = ST_HASH_18100; hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN; break; diff --git a/src/usage.c b/src/usage.c index 8f8eb0a12..a3698c314 100644 --- a/src/usage.c +++ b/src/usage.c @@ -170,7 +170,7 @@ static const char *const USAGE_BIG[] = " 60 | HMAC-MD5 (key = $salt) | Raw Hash, Authenticated", " 150 | HMAC-SHA1 (key = $pass) | Raw Hash, Authenticated", " 160 | HMAC-SHA1 (key = $salt) | Raw Hash, Authenticated", - " 17300 | TOTP (HMAC-SHA1) | Raw Hash, Authenticated", + " 18100 | TOTP (HMAC-SHA1) | Raw Hash, Authenticated", " 1450 | HMAC-SHA256 (key = $pass) | Raw Hash, Authenticated", " 1460 | HMAC-SHA256 (key = $salt) | Raw Hash, Authenticated", " 1750 | HMAC-SHA512 (key = $pass) | Raw Hash, Authenticated", diff --git a/tools/test.pl b/tools/test.pl index 852e3159d..acd577df2 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -58,7 +58,7 @@ my $hashcat = "./hashcat"; my $MAX_LEN = 55; -my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7701, 7800, 7801, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 15400, 15500, 15600, 15700, 15900, 16000, 16100, 16200, 16300, 16400, 16500, 16600, 16700, 16800, 16900, 17300, 99999); +my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7701, 7800, 7801, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 15400, 15500, 15600, 15700, 15900, 16000, 16100, 16200, 16300, 16400, 16500, 16600, 16700, 16800, 16900, 18100, 99999); my %is_utf16le = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800); my %less_fifteen = map { $_ => 1 } qw (500 1600 1800 3200 6300 7400 10500 10700); @@ -239,7 +239,7 @@ sub verify $word = substr ($line, $index + 1); } # hash:salt - elsif ($mode == 10 || $mode == 11 || $mode == 12 || $mode == 20 || $mode == 21 || $mode == 22 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 112 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1100 || $mode == 1410 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 2410 || $mode == 2611 || $mode == 2711 || $mode == 2811 || $mode == 3100 || $mode == 3610 || $mode == 3710 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4520 || $mode == 4521 || $mode == 4522 || $mode == 4900 || $mode == 5800 || $mode == 8400 || $mode == 11000 || $mode == 12600 || $mode == 13500 || $mode == 13800 || $mode == 13900 || $mode == 14000 || $mode == 14100 || $mode == 14400 || $mode == 14900 || $mode == 15000 || $mode == 17300) + elsif ($mode == 10 || $mode == 11 || $mode == 12 || $mode == 20 || $mode == 21 || $mode == 22 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 112 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1100 || $mode == 1410 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 2410 || $mode == 2611 || $mode == 2711 || $mode == 2811 || $mode == 3100 || $mode == 3610 || $mode == 3710 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4520 || $mode == 4521 || $mode == 4522 || $mode == 4900 || $mode == 5800 || $mode == 8400 || $mode == 11000 || $mode == 12600 || $mode == 13500 || $mode == 13800 || $mode == 13900 || $mode == 14000 || $mode == 14100 || $mode == 14400 || $mode == 14900 || $mode == 15000 || $mode == 18100) { # get hash my $index1 = index ($line, ":"); @@ -3551,7 +3551,7 @@ sub passthrough { $tmp_hash = gen_hash ($mode, $word_buf, ""); } - elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4900 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 17300) + elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4900 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 18100) { my $salt_len = get_random_num (1, 15); @@ -4091,7 +4091,7 @@ sub single } } } - elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 16500 || $mode == 17300) + elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 16500 || $mode == 18100) { my $salt_len = get_random_num (1, 15); @@ -9929,7 +9929,7 @@ END_CODE $tmp_hash = sprintf ("%s*%s*%s*%s", substr ($pmkid, 0, 32), $macap, $macsta, $essid); } - elsif ($mode == 17300) + elsif ($mode == 18100) { my $paddedTime = sprintf("%016x", int(int($salt_buf) / 30)); my $data = pack('H*', $paddedTime); From 37983de4b219ae0f50196b295b30492910decad6 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Tue, 16 Oct 2018 17:39:46 -0400 Subject: [PATCH 07/32] Fix compile warnings in totp_parse_hash() --- include/interface.h | 1 + src/interface.c | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/include/interface.h b/include/interface.h index 0ccfd8eb5..7df40ec72 100644 --- a/include/interface.h +++ b/include/interface.h @@ -1597,6 +1597,7 @@ int filevault2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu int wpa_pmkid_pbkdf2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int wpa_pmkid_pmk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int ansible_vault_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); +int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); /** * hook functions diff --git a/src/interface.c b/src/interface.c index fcd665a34..a5b684e24 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5197,7 +5197,7 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE } if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); - u8 *hash_pos = token.buf[0]; + //u8 *hash_pos = token.buf[0]; digest[1] = otp_code; From 83c78cdf0bc6d79cbb94d5566ca84cd2637c06f5 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 08:15:49 -0400 Subject: [PATCH 08/32] Cleanup unused code --- src/interface.c | 2 -- 1 file changed, 2 deletions(-) diff --git a/src/interface.c b/src/interface.c index a5b684e24..b82827c62 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5197,8 +5197,6 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE } if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); - //u8 *hash_pos = token.buf[0]; - digest[1] = otp_code; u8 *salt_pos = token.buf[1]; From d66200a406250cbe5ce4ab5b8e0879bc4123ca18 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 08:17:48 -0400 Subject: [PATCH 09/32] Force --keep-guessing on mode 18100 --- src/interface.c | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/interface.c b/src/interface.c index b82827c62..35a50df8c 100644 --- a/src/interface.c +++ b/src/interface.c @@ -27636,7 +27636,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL; hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE | OPTS_TYPE_PT_ADD80 - | OPTS_TYPE_PT_ADDBITS15; + | OPTS_TYPE_PT_ADDBITS15 + | OPTS_TYPE_PT_NEVERCRACK; hashconfig->kern_type = KERN_TYPE_TOTP_HMACSHA1; hashconfig->dgst_size = DGST_SIZE_4_5; hashconfig->parse_func = totp_parse_hash; From b29b7b8188ed97c2a0b5d49fb04bf1818dfa39d7 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 08:54:52 -0400 Subject: [PATCH 10/32] Convert arithmetic ops to logical ops in byte alignment --- OpenCL/m18100_a0-pure.cl | 18 ++++++++++-------- OpenCL/m18100_a3-pure.cl | 18 ++++++++++-------- 2 files changed, 20 insertions(+), 16 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index 4de835280..76a7f2705 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -66,16 +66,17 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru unsigned int otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - switch(otp_offset%4) + // on some systems, &3 is faster than %4, so we will use it in our switch() + switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; @@ -158,16 +159,17 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru unsigned int otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - switch(otp_offset%4) + // on some systems, &3 is faster than %4, so we will use it in our switch() + switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index a106acb99..8f90361e9 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -75,16 +75,17 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule unsigned int otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - switch(otp_offset%4) + // on some systems, &3 is faster than %4, so we will use it in our switch() + switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfeffffff) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; @@ -178,16 +179,17 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule unsigned int otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - switch(otp_offset%4) + // on some systems, &3 is faster than %4, so we will use it in our switch() + switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xffff) << 16) | ((ctx.opad.h[otp_offset/4+1] % 0xffff0000) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0xff) << 24) | ((ctx.opad.h[otp_offset/4+1] % 0xffffff00) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; From 4724e1f5897f8ffa879c43fc8a1d8b9e53423ad5 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 10:42:22 -0400 Subject: [PATCH 11/32] Add 18100 to changes.txt --- docs/changes.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/docs/changes.txt b/docs/changes.txt index dc4be4838..840f93014 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -18,6 +18,7 @@ - Added hash-mode 17800 = Keccak-256 - Added hash-mode 17900 = Keccak-384 - Added hash-mode 18000 = Keccak-512 +- Added hash-mode 18100 = TOTP (HMAC-SHA1) - Removed hash-mode 5000 = SHA-3 (Keccak) ## From 1d43540fc4fae07af5b65b4f465bc905fc8431e8 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 11:03:20 -0400 Subject: [PATCH 12/32] Simplify alignment masks for mode 18100 --- OpenCL/m18100_a0-pure.cl | 12 ++++++------ OpenCL/m18100_a3-pure.cl | 14 +++++++------- 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index 76a7f2705..7ab09d5d1 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -70,13 +70,13 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; @@ -163,13 +163,13 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index 8f90361e9..1352338a1 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -79,13 +79,13 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfeffffff) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; @@ -182,14 +182,14 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule // on some systems, &3 is faster than %4, so we will use it in our switch() switch(otp_offset&3) { - case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; From 27366c3effea724a1440d47b07041948de9ba30c Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 11:44:34 -0400 Subject: [PATCH 13/32] Don't duplicate existing functionality --- src/interface.c | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/src/interface.c b/src/interface.c index 35a50df8c..8a9daa559 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5200,14 +5200,10 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE digest[1] = otp_code; u8 *salt_pos = token.buf[1]; - int salt_len = token.len[1]; // convert ascii timestamp to ulong timestamp unsigned long timestamp = 0; - for(int i = 0; i Date: Wed, 17 Oct 2018 11:47:04 -0400 Subject: [PATCH 14/32] Fix zero-padding on token output for matches --- src/interface.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/interface.c b/src/interface.c index 8a9daa559..0e4ac545a 100644 --- a/src/interface.c +++ b/src/interface.c @@ -22337,7 +22337,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le // for now, we only need to worry about 32 bit counters. // we also need to multiply salt by our step to see the floor of our original timestamp range. // again, we will use the default RFC 6238 step of 30. - snprintf (out_buf, out_len - 1, "%d:%d", digest_buf[1], byte_swap_32(salt.salt_buf[1]) * 30); + snprintf (out_buf, out_len - 1, "%06d:%d", digest_buf[1], byte_swap_32(salt.salt_buf[1]) * 30); } else if (hash_mode == 99999) { From db4ec8ed2c7159242f6b9aae3af5684645fbd962 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 16:34:34 -0400 Subject: [PATCH 15/32] Fix formatting to comply with hashcat coding guidelines --- OpenCL/m18100_a0-pure.cl | 4 ++-- OpenCL/m18100_a3-pure.cl | 4 ++-- src/hashes.c | 4 ++-- src/interface.c | 12 +++++------- tools/test.pl | 6 +++--- 5 files changed, 14 insertions(+), 16 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index 7ab09d5d1..e577d62ae 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -67,7 +67,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset&3) + switch(otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); @@ -160,7 +160,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset&3) + switch(otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index 1352338a1..9838c34aa 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -76,7 +76,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset&3) + switch(otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); @@ -180,7 +180,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset&3) + switch(otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); diff --git a/src/hashes.c b/src/hashes.c index 7ade9faac..1681a585e 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -284,10 +284,10 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl // encode our plain base32_encode (int_to_base32, (const u8 *) plain_ptr, plain_len, (u8 *) temp_buf); - plain_len = strlen((const char *) temp_buf); + plain_len = strlen ((const char *) temp_buf); // copy the base32 content into our plain buffer - plain_ptr = (u8 *) strdup((const char *) temp_ptr); + plain_ptr = (u8 *) strdup ((const char *) temp_ptr); } // crackpos diff --git a/src/interface.c b/src/interface.c index 0e4ac545a..1ee32e416 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5193,7 +5193,7 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE int otp_code = 0; for (int i = 0; i < 6; i++) { - otp_code = otp_code * 10 + itoa32_to_int(input_buf[i]); + otp_code = otp_code * 10 + itoa32_to_int (input_buf[i]); } if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); @@ -5203,16 +5203,14 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE // convert ascii timestamp to ulong timestamp unsigned long timestamp = 0; - timestamp = hc_strtoul((const char *) salt_pos, NULL, 10); + timestamp = hc_strtoul ((const char *) salt_pos, NULL, 10); // divide our timestamp by our step. We will use the RFC 6238 default of 30 for now timestamp /= 30; // convert counter to 8-byte salt - for (int i = 2; i--; timestamp >>= 32) - { - salt->salt_buf[i] = byte_swap_32(timestamp); - } + salt->salt_buf[1] = byte_swap_32 (timestamp); + salt->salt_buf[0] = byte_swap_32 (timestamp >> 32); // our salt will always be 8 bytes salt->salt_len = 8; @@ -22337,7 +22335,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le // for now, we only need to worry about 32 bit counters. // we also need to multiply salt by our step to see the floor of our original timestamp range. // again, we will use the default RFC 6238 step of 30. - snprintf (out_buf, out_len - 1, "%06d:%d", digest_buf[1], byte_swap_32(salt.salt_buf[1]) * 30); + snprintf (out_buf, out_len - 1, "%06d:%d", digest_buf[1], byte_swap_32 (salt.salt_buf[1]) * 30); } else if (hash_mode == 99999) { diff --git a/tools/test.pl b/tools/test.pl index c5620d27b..f10ae9b79 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -9974,8 +9974,8 @@ END_CODE } elsif ($mode == 18100) { - my $paddedTime = sprintf("%016x", int(int($salt_buf) / 30)); - my $data = pack('H*', $paddedTime); + my $paddedTime = sprintf ("%016x", int (int ($salt_buf) / 30)); + my $data = pack ('H*', $paddedTime); my $key = $word_buf; $hash_buf = hmac_hex ($data, $key, \&sha1, 64); @@ -9987,7 +9987,7 @@ END_CODE $token %= 1000000; ## token must be leading zero padded, and salt leading zero stripped - $tmp_hash = sprintf ("%06d:%d", $token, int($salt_buf)); + $tmp_hash = sprintf ("%06d:%d", $token, int ($salt_buf)); } elsif ($mode == 99999) { From 24ab7cae2a0b741194eefb9fd29d06901ab740a2 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 16:47:58 -0400 Subject: [PATCH 16/32] Add a1 kernel for mode 18100 --- OpenCL/m18100_a1-pure.cl | 238 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 238 insertions(+) create mode 100644 OpenCL/m18100_a1-pure.cl diff --git a/OpenCL/m18100_a1-pure.cl b/OpenCL/m18100_a1-pure.cl new file mode 100644 index 000000000..03bd729f1 --- /dev/null +++ b/OpenCL/m18100_a1-pure.cl @@ -0,0 +1,238 @@ +/** + * 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_scalar.cl" +#include "inc_hash_sha1.cl" + +__kernel void m18100_mxx (__global pw_t *pws, __global 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 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 u64 gid_max) +{ + /** + * modifier + */ + + const u64 lid = get_local_id (0); + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * base + */ + + const u32 pw_len = pws[gid].pw_len; + + u32 w[64] = { 0 }; + + for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1) + { + w[idx] = swap32_S (pws[gid].i[idx]); + } + + const u32 salt_len = salt_bufs[salt_pos].salt_len; + + u32 s[64] = { 0 }; + + for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1) + { + s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); + } + + /** + * loop + */ + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) + { + const u32 comb_len = combs_buf[il_pos].pw_len; + + u32 c[64]; + + #ifdef _unroll + #pragma unroll + #endif + for (int idx = 0; idx < 64; idx++) + { + c[idx] = swap32_S (combs_buf[il_pos].i[idx]); + } + + switch_buffer_by_offset_1x64_be_S (c, pw_len); + + #ifdef _unroll + #pragma unroll + #endif + for (int i = 0; i < 64; i++) + { + c[i] |= w[i]; + } + + sha1_hmac_ctx_t ctx; + + sha1_hmac_init (&ctx, c, pw_len + comb_len); + + sha1_hmac_update (&ctx, s, salt_len); + + sha1_hmac_final (&ctx); + + // calculate the offset using the least 4 bits of the last byte of our hash + const int otp_offset = ctx.opad.h[4] & 0xf; + + // initialize a buffer for the otp code + unsigned int otp_code = 0; + + // grab 4 consecutive bytes of the hash, starting at offset + // on some systems, &3 is faster than %4, so we will use it in our switch() + switch(otp_offset & 3) + { + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); + break; + case 2: + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); + break; + case 3: + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); + break; + default: + otp_code = ctx.opad.h[otp_offset/4]; + break; + } + // take only the lower 31 bits + otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code + otp_code %= 1000000; + + const u32 r0 = ctx.opad.h[DGST_R0]; + const u32 r1 = ctx.opad.h[DGST_R1]; + const u32 r2 = ctx.opad.h[DGST_R2]; + const u32 r3 = ctx.opad.h[DGST_R3]; + + COMPARE_M_SCALAR (otp_code, 0, 0, 0); + } +} + +__kernel void m18100_sxx (__global pw_t *pws, __global 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 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 u64 gid_max) +{ + /** + * modifier + */ + + const u64 lid = get_local_id (0); + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * digest + */ + + const u32 search[4] = + { + digests_buf[digests_offset].digest_buf[DGST_R0], + digests_buf[digests_offset].digest_buf[DGST_R1], + digests_buf[digests_offset].digest_buf[DGST_R2], + digests_buf[digests_offset].digest_buf[DGST_R3] + }; + + /** + * base + */ + + const u32 pw_len = pws[gid].pw_len; + + u32 w[64] = { 0 }; + + for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1) + { + w[idx] = swap32_S (pws[gid].i[idx]); + } + + const u32 salt_len = salt_bufs[salt_pos].salt_len; + + u32 s[64] = { 0 }; + + for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1) + { + s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); + } + + /** + * loop + */ + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) + { + const u32 comb_len = combs_buf[il_pos].pw_len; + + u32 c[64]; + + #ifdef _unroll + #pragma unroll + #endif + for (int idx = 0; idx < 64; idx++) + { + c[idx] = swap32_S (combs_buf[il_pos].i[idx]); + } + + switch_buffer_by_offset_1x64_be_S (c, pw_len); + + #ifdef _unroll + #pragma unroll + #endif + for (int i = 0; i < 64; i++) + { + c[i] |= w[i]; + } + + sha1_hmac_ctx_t ctx; + + sha1_hmac_init (&ctx, c, pw_len + comb_len); + + sha1_hmac_update (&ctx, s, salt_len); + + sha1_hmac_final (&ctx); + + // calculate the offset using the least 4 bits of the last byte of our hash + const int otp_offset = ctx.opad.h[4] & 0xf; + + // initialize a buffer for the otp code + unsigned int otp_code = 0; + + // grab 4 consecutive bytes of the hash, starting at offset + // on some systems, &3 is faster than %4, so we will use it in our switch() + switch(otp_offset & 3) + { + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); + break; + case 2: + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); + break; + case 3: + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); + break; + default: + otp_code = ctx.opad.h[otp_offset/4]; + break; + } + // take only the lower 31 bits + otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code + otp_code %= 1000000; + + const u32 r0 = ctx.opad.h[DGST_R0]; + const u32 r1 = ctx.opad.h[DGST_R1]; + const u32 r2 = ctx.opad.h[DGST_R2]; + const u32 r3 = ctx.opad.h[DGST_R3]; + + COMPARE_S_SCALAR (otp_code, 0, 0, 0); + } +} From b7e8fc9015d4f0802cfb12d796e55304241881f4 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 16:49:45 -0400 Subject: [PATCH 17/32] Add mode 18100 index to tab complete --- extra/tab_completion/hashcat.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/extra/tab_completion/hashcat.sh b/extra/tab_completion/hashcat.sh index df833b085..0a683d467 100644 --- a/extra/tab_completion/hashcat.sh +++ b/extra/tab_completion/hashcat.sh @@ -176,7 +176,7 @@ _hashcat () { local VERSION=4.2.1 - local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 200 300 400 500 501 600 900 1000 1100 1400 1410 1411 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2501 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12001 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14700 14800 14900 15000 15100 15200 15300 15400 15500 15600 15700 15900 16000 16100 16200 16300 16400 16500 16600 16700 16800 16801 16900 17300 17400 17500 17600 17700 17800 17900 18000" + local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 200 300 400 500 501 600 900 1000 1100 1400 1410 1411 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2501 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12001 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14700 14800 14900 15000 15100 15200 15300 15400 15500 15600 15700 15900 16000 16100 16200 16300 16400 16500 16600 16700 16800 16801 16900 17300 17400 17500 17600 17700 17800 17900 18000 18100" local ATTACK_MODES="0 1 3 6 7" local HCCAPX_MESSAGE_PAIRS="0 1 2 3 4 5" local OUTFILE_FORMATS="1 2 3 4 5 6 7 8 9 10 11 12 13 14 15" From ac2b811941e674ad0b728e6676d2c6e95adc9daf Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 19:27:06 -0400 Subject: [PATCH 18/32] Categorize TOTP under "One-Time Passwords" --- src/usage.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/usage.c b/src/usage.c index c5a82955b..87181385e 100644 --- a/src/usage.c +++ b/src/usage.c @@ -177,7 +177,6 @@ static const char *const USAGE_BIG[] = " 60 | HMAC-MD5 (key = $salt) | Raw Hash, Authenticated", " 150 | HMAC-SHA1 (key = $pass) | Raw Hash, Authenticated", " 160 | HMAC-SHA1 (key = $salt) | Raw Hash, Authenticated", - " 18100 | TOTP (HMAC-SHA1) | Raw Hash, Authenticated", " 1450 | HMAC-SHA256 (key = $pass) | Raw Hash, Authenticated", " 1460 | HMAC-SHA256 (key = $salt) | Raw Hash, Authenticated", " 1750 | HMAC-SHA512 (key = $pass) | Raw Hash, Authenticated", @@ -381,6 +380,7 @@ static const char *const USAGE_BIG[] = " 15700 | Ethereum Wallet, SCRYPT | Password Managers", " 16300 | Ethereum Pre-Sale Wallet, PBKDF2-HMAC-SHA256 | Password Managers", " 16900 | Ansible Vault | Password Managers", + " 18100 | TOTP (HMAC-SHA1) | One-Time Passwords", " 99999 | Plaintext | Plaintext", "", "- [ Outfile Formats ] -", From 3869ce9246d0bf415afbd453d026f64af02eb0dd Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Thu, 18 Oct 2018 08:55:55 -0400 Subject: [PATCH 19/32] More coding style fixes --- OpenCL/m18100_a0-pure.cl | 4 ++-- OpenCL/m18100_a1-pure.cl | 4 ++-- OpenCL/m18100_a3-pure.cl | 6 +++--- src/interface.c | 16 ++++++---------- 4 files changed, 13 insertions(+), 17 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index e577d62ae..20b1dc5af 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -67,7 +67,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset & 3) + switch (otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); @@ -160,7 +160,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset & 3) + switch (otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); diff --git a/OpenCL/m18100_a1-pure.cl b/OpenCL/m18100_a1-pure.cl index 03bd729f1..d82915ac5 100644 --- a/OpenCL/m18100_a1-pure.cl +++ b/OpenCL/m18100_a1-pure.cl @@ -90,7 +90,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset & 3) + switch (otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); @@ -208,7 +208,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset & 3) + switch (otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index 9838c34aa..b80ac8dce 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -76,7 +76,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset & 3) + switch (otp_offset & 3) { case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); @@ -180,9 +180,9 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() - switch(otp_offset & 3) + switch (otp_offset & 3) { - case 1: + case 1: otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: diff --git a/src/interface.c b/src/interface.c index 1ee32e416..e7574a3d7 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5184,17 +5184,14 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE token.len_min[1] *= 2; token.len_max[1] *= 2; - token.attr[1] |= TOKEN_ATTR_VERIFY_HEX; + token.attr[1] |= TOKEN_ATTR_VERIFY_DIGIT; } const int rc_tokenizer = input_tokenizer (input_buf, input_len, &token); // now we need to reduce our hash into a token - int otp_code = 0; - for (int i = 0; i < 6; i++) - { - otp_code = otp_code * 10 + itoa32_to_int (input_buf[i]); - } + int otp_code = hc_strtoul ((const char *) input_buf, NULL, 10); + if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); digest[1] = otp_code; @@ -5202,15 +5199,14 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE u8 *salt_pos = token.buf[1]; // convert ascii timestamp to ulong timestamp - unsigned long timestamp = 0; - timestamp = hc_strtoul ((const char *) salt_pos, NULL, 10); + u64 timestamp = hc_strtoul ((const char *) salt_pos, NULL, 10); // divide our timestamp by our step. We will use the RFC 6238 default of 30 for now timestamp /= 30; // convert counter to 8-byte salt - salt->salt_buf[1] = byte_swap_32 (timestamp); - salt->salt_buf[0] = byte_swap_32 (timestamp >> 32); + salt->salt_buf[1] = byte_swap_32 ((u32) (timestamp >> 0)); + salt->salt_buf[0] = byte_swap_32 ((u32) (timestamp >> 32)); // our salt will always be 8 bytes salt->salt_len = 8; From b657c75583f3a750e0837ac9f44b9bd124b62172 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Thu, 18 Oct 2018 09:36:57 -0400 Subject: [PATCH 20/32] Explicity cast otp_offset --- OpenCL/m18100_a0-pure.cl | 4 ++-- OpenCL/m18100_a1-pure.cl | 4 ++-- OpenCL/m18100_a3-pure.cl | 4 ++-- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index 20b1dc5af..7903b9cb3 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -60,7 +60,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru sha1_hmac_final (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = ctx.opad.h[4] & 0xf; + const int otp_offset = (int) ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code unsigned int otp_code = 0; @@ -153,7 +153,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru sha1_hmac_final (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = ctx.opad.h[4] & 0xf; + const int otp_offset = (int) ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code unsigned int otp_code = 0; diff --git a/OpenCL/m18100_a1-pure.cl b/OpenCL/m18100_a1-pure.cl index d82915ac5..d03f1a459 100644 --- a/OpenCL/m18100_a1-pure.cl +++ b/OpenCL/m18100_a1-pure.cl @@ -83,7 +83,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = ctx.opad.h[4] & 0xf; + const int otp_offset = (int) ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code unsigned int otp_code = 0; @@ -201,7 +201,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = ctx.opad.h[4] & 0xf; + const int otp_offset = (int) ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code unsigned int otp_code = 0; diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index b80ac8dce..22083acf8 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -69,7 +69,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final_vector (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = ctx.opad.h[4] & 0xf; + const int otp_offset = (int) ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code unsigned int otp_code = 0; @@ -173,7 +173,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final_vector (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = ctx.opad.h[4] & 0xf; + const int otp_offset = (int) ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code unsigned int otp_code = 0; From 4cb93ba9bfc6c1c30d6054b51293ead31118bf5a Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Thu, 18 Oct 2018 14:34:28 -0400 Subject: [PATCH 21/32] Add 18100 to test.sh --- tools/test.sh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tools/test.sh b/tools/test.sh index 005acaafa..40c6388e9 100755 --- a/tools/test.sh +++ b/tools/test.sh @@ -9,7 +9,7 @@ TDIR="$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )" # missing hash types: 5200,6251,6261,6271,6281 -HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 125 130 131 132 133 140 141 150 160 200 300 400 500 600 900 1000 1100 1300 1400 1410 1411 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5100 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7701 7800 7801 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12001 12100 12200 12300 12400 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 15400 15500 15600 15700 15900 16000 16100 16200 16300 16400 16500 16600 16700 16800 16900 17300 17400 17500 17600 17700 17800 17900 18000 99999" +HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 125 130 131 132 133 140 141 150 160 200 300 400 500 600 900 1000 1100 1300 1400 1410 1411 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5100 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7000 7100 7200 7300 7400 7500 7700 7701 7800 7801 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12001 12100 12200 12300 12400 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 15400 15500 15600 15700 15900 16000 16100 16200 16300 16400 16500 16600 16700 16800 16900 17300 17400 17500 17600 17700 17800 17900 18000 18100 99999" #ATTACK_MODES="0 1 3 6 7" ATTACK_MODES="0 1 3 7" @@ -20,7 +20,7 @@ MATCH_PASS_ONLY="2500 5300 5400 6600 6800 8200" HASHFILE_ONLY="2500" -NEVER_CRACK="11600 14900" +NEVER_CRACK="11600 14900 18100" SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12001 12100 12200 12300 12400 12500 12700 12800 12900 13000 13200 13400 13600 13751 13752 13753 14600 14611 14612 14613 14621 14622 14623 14631 14632 14633 14641 14642 14643 14700 14800 15100 15200 15300 15600 15700 15900 16000 16200 16300 16800 16900" From 0e5704c77e80d0e70f0953a5f44ef7216cd99043 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Thu, 18 Oct 2018 15:58:21 -0400 Subject: [PATCH 22/32] Disable NEW_SIMD_CODE for 18100 (it's not compatible) --- OpenCL/m18100_a3-pure.cl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index 22083acf8..dcba543b5 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" From 7904b9ae059bc0bad0c4aca5083684e0ce28f823 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Fri, 19 Oct 2018 07:56:51 -0400 Subject: [PATCH 23/32] Fix kernel types to align with style guide --- OpenCL/m18100_a0-pure.cl | 8 ++++---- OpenCL/m18100_a1-pure.cl | 8 ++++---- OpenCL/m18100_a3-pure.cl | 8 ++++---- 3 files changed, 12 insertions(+), 12 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index 7903b9cb3..420c6f9c4 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -60,10 +60,10 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru sha1_hmac_final (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = (int) ctx.opad.h[4] & 0xf; + const u32x otp_offset = ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code - unsigned int otp_code = 0; + u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() @@ -153,10 +153,10 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru sha1_hmac_final (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = (int) ctx.opad.h[4] & 0xf; + const u32x otp_offset = ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code - unsigned int otp_code = 0; + u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() diff --git a/OpenCL/m18100_a1-pure.cl b/OpenCL/m18100_a1-pure.cl index d03f1a459..b58389975 100644 --- a/OpenCL/m18100_a1-pure.cl +++ b/OpenCL/m18100_a1-pure.cl @@ -83,10 +83,10 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = (int) ctx.opad.h[4] & 0xf; + const u32x otp_offset = ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code - unsigned int otp_code = 0; + u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() @@ -201,10 +201,10 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = (int) ctx.opad.h[4] & 0xf; + const u32x otp_offset = ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code - unsigned int otp_code = 0; + u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index dcba543b5..15e32de93 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -69,10 +69,10 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final_vector (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = (int) ctx.opad.h[4] & 0xf; + const u32x otp_offset = ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code - unsigned int otp_code = 0; + u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() @@ -173,10 +173,10 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final_vector (&ctx); // calculate the offset using the least 4 bits of the last byte of our hash - const int otp_offset = (int) ctx.opad.h[4] & 0xf; + const u32x otp_offset = ctx.opad.h[4] & 0xf; // initialize a buffer for the otp code - unsigned int otp_code = 0; + u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset // on some systems, &3 is faster than %4, so we will use it in our switch() From 8b2af6b207fa9a91cfb0ba441351c66bf309a871 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Fri, 19 Oct 2018 08:45:44 -0400 Subject: [PATCH 24/32] More code style changes --- src/hashes.c | 3 ++- src/interface.c | 4 ++-- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/src/hashes.c b/src/hashes.c index 1681a585e..11676630d 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -254,6 +254,7 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl debugfile_ctx_t *debugfile_ctx = hashcat_ctx->debugfile_ctx; loopback_ctx_t *loopback_ctx = hashcat_ctx->loopback_ctx; hashes_t *hashes = hashcat_ctx->hashes; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; const u32 salt_pos = plain->salt_pos; const u32 digest_pos = plain->digest_pos; // relative @@ -276,7 +277,7 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl build_plain (hashcat_ctx, device_param, plain, plain_buf, &plain_len); // TOTP should be base32 encoded - if (hashcat_ctx->hashconfig->hash_mode == KERN_TYPE_TOTP_HMACSHA1) + if (hashconfig->hash_mode == KERN_TYPE_TOTP_HMACSHA1) { // we need a temp buffer for the base32 encoding u32 temp_buf[64] = { 0 }; diff --git a/src/interface.c b/src/interface.c index e7574a3d7..18a043b66 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5189,11 +5189,11 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE const int rc_tokenizer = input_tokenizer (input_buf, input_len, &token); + if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); + // now we need to reduce our hash into a token int otp_code = hc_strtoul ((const char *) input_buf, NULL, 10); - if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); - digest[1] = otp_code; u8 *salt_pos = token.buf[1]; From 770e21e5eef86de9e98e1ebac136b0643ade0bbd Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Fri, 19 Oct 2018 15:16:46 -0400 Subject: [PATCH 25/32] Support 64 bit timestamps properly --- src/interface.c | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/src/interface.c b/src/interface.c index 18a043b66..83bfb912b 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5199,7 +5199,7 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE u8 *salt_pos = token.buf[1]; // convert ascii timestamp to ulong timestamp - u64 timestamp = hc_strtoul ((const char *) salt_pos, NULL, 10); + u64 timestamp = hc_strtoull ((const char *) salt_pos, NULL, 10); // divide our timestamp by our step. We will use the RFC 6238 default of 30 for now timestamp /= 30; @@ -22328,10 +22328,13 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le else if (hash_mode == 18100) { // salt_buf[1] holds our 32 bit value. salt_buf[0] and salt_buf[1] would be 64 bits. - // for now, we only need to worry about 32 bit counters. // we also need to multiply salt by our step to see the floor of our original timestamp range. // again, we will use the default RFC 6238 step of 30. - snprintf (out_buf, out_len - 1, "%06d:%d", digest_buf[1], byte_swap_32 (salt.salt_buf[1]) * 30); + + u64 tmp_salt_buf = (((u64) byte_swap_32 (salt.salt_buf[0])) << 32) | ((u64) byte_swap_32 (salt.salt_buf[1])); + tmp_salt_buf *= 30; + + snprintf (out_buf, out_len - 1, "%06d:%llu", digest_buf[1], tmp_salt_buf); } else if (hash_mode == 99999) { From fddfd835d24f1dc0fc578309611320754d4d10ff Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Fri, 19 Oct 2018 15:35:52 -0400 Subject: [PATCH 26/32] Support 64 bit timestamps properly --- OpenCL/m18100_a0-pure.cl | 4 ++-- OpenCL/m18100_a1-pure.cl | 4 ++-- OpenCL/m18100_a3-pure.cl | 4 ++-- src/interface.c | 11 +++++++---- 4 files changed, 13 insertions(+), 10 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index 420c6f9c4..bb0894e19 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -32,7 +32,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru COPY_PW (pws[gid]); - const u32 salt_len = salt_bufs[salt_pos].salt_len; + const u32 salt_len = 8; u32 s[64] = { 0 }; @@ -125,7 +125,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru COPY_PW (pws[gid]); - const u32 salt_len = salt_bufs[salt_pos].salt_len; + const u32 salt_len = 8; u32 s[64] = { 0 }; diff --git a/OpenCL/m18100_a1-pure.cl b/OpenCL/m18100_a1-pure.cl index b58389975..815eec778 100644 --- a/OpenCL/m18100_a1-pure.cl +++ b/OpenCL/m18100_a1-pure.cl @@ -37,7 +37,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule w[idx] = swap32_S (pws[gid].i[idx]); } - const u32 salt_len = salt_bufs[salt_pos].salt_len; + const u32 salt_len = 8; u32 s[64] = { 0 }; @@ -155,7 +155,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule w[idx] = swap32_S (pws[gid].i[idx]); } - const u32 salt_len = salt_bufs[salt_pos].salt_len; + const u32 salt_len = 8; u32 s[64] = { 0 }; diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index 15e32de93..a71516aa8 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -37,7 +37,7 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule w[idx] = pws[gid].i[idx]; } - const u32 salt_len = salt_bufs[salt_pos].salt_len; + const u32 salt_len = 8; u32x s[64] = { 0 }; @@ -141,7 +141,7 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule w[idx] = pws[gid].i[idx]; } - const u32 salt_len = salt_bufs[salt_pos].salt_len; + const u32 salt_len = 8; u32x s[64] = { 0 }; diff --git a/src/interface.c b/src/interface.c index 83bfb912b..5e6c97742 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5201,6 +5201,10 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE // convert ascii timestamp to ulong timestamp u64 timestamp = hc_strtoull ((const char *) salt_pos, NULL, 10); + // store the original salt value. Step division will destroy granularity for output + salt->salt_buf[3] = ((u32) (timestamp >> 0)); + salt->salt_buf[2] = ((u32) (timestamp >> 32)); + // divide our timestamp by our step. We will use the RFC 6238 default of 30 for now timestamp /= 30; @@ -5208,8 +5212,8 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE salt->salt_buf[1] = byte_swap_32 ((u32) (timestamp >> 0)); salt->salt_buf[0] = byte_swap_32 ((u32) (timestamp >> 32)); - // our salt will always be 8 bytes - salt->salt_len = 8; + // our salt will always be 8 bytes, but we are going to cheat and store it twice, so... + salt->salt_len = 16; return (PARSER_OK); } @@ -22331,8 +22335,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le // we also need to multiply salt by our step to see the floor of our original timestamp range. // again, we will use the default RFC 6238 step of 30. - u64 tmp_salt_buf = (((u64) byte_swap_32 (salt.salt_buf[0])) << 32) | ((u64) byte_swap_32 (salt.salt_buf[1])); - tmp_salt_buf *= 30; + u64 tmp_salt_buf = (((u64) (salt.salt_buf[2])) << 32) | ((u64) (salt.salt_buf[3])); snprintf (out_buf, out_len - 1, "%06d:%llu", digest_buf[1], tmp_salt_buf); } From 6558253d7200d8221bb7f3f3ac29e48a68223335 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Fri, 19 Oct 2018 15:37:02 -0400 Subject: [PATCH 27/32] Don't cause memory leaks... --- src/hashes.c | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/hashes.c b/src/hashes.c index 11676630d..b451c4b23 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -288,7 +288,7 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl plain_len = strlen ((const char *) temp_buf); // copy the base32 content into our plain buffer - plain_ptr = (u8 *) strdup ((const char *) temp_ptr); + strncpy ((char *) plain_ptr, (char *) temp_ptr, sizeof (plain_buf)); } // crackpos From 3705ebfbf60134499f6d0f7685e3a569cc7080a6 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Fri, 19 Oct 2018 16:00:26 -0400 Subject: [PATCH 28/32] Remove extra whitespace in test.pl --- tools/test.pl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tools/test.pl b/tools/test.pl index f10ae9b79..57d1800cc 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -9980,7 +9980,7 @@ END_CODE $hash_buf = hmac_hex ($data, $key, \&sha1, 64); - my $offset = hex (substr ($hash_buf, -8)) & 0xf; + my $offset = hex (substr ($hash_buf, -8)) & 0xf; $offset *= 2; my $token = hex (substr ($hash_buf, $offset, 8)); $token &= 0x7fffffff; From cef3c1cb06cbe349a66a49ea3cccd98ce88a89dc Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Sat, 20 Oct 2018 22:00:15 -0400 Subject: [PATCH 29/32] Use C99 format specifier --- src/interface.c | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/interface.c b/src/interface.c index 5e6c97742..c4f08b6b4 100644 --- a/src/interface.c +++ b/src/interface.c @@ -3,6 +3,7 @@ * License.....: MIT */ +#include #include "common.h" #include "types.h" #include "bitops.h" @@ -22337,7 +22338,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le u64 tmp_salt_buf = (((u64) (salt.salt_buf[2])) << 32) | ((u64) (salt.salt_buf[3])); - snprintf (out_buf, out_len - 1, "%06d:%llu", digest_buf[1], tmp_salt_buf); + snprintf (out_buf, out_len - 1, "%06d:%" PRIu64, digest_buf[1], tmp_salt_buf); } else if (hash_mode == 99999) { From 1809794299602f72446214646cb45c6bb3d60e6c Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Mon, 22 Oct 2018 08:02:57 -0400 Subject: [PATCH 30/32] Simplify code in check_hash() --- src/hashes.c | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/hashes.c b/src/hashes.c index b451c4b23..3f7b9d5e5 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -284,8 +284,7 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl u8 *temp_ptr = (u8 *) temp_buf; // encode our plain - base32_encode (int_to_base32, (const u8 *) plain_ptr, plain_len, (u8 *) temp_buf); - plain_len = strlen ((const char *) temp_buf); + plain_len = base32_encode (int_to_base32, (const u8 *) plain_ptr, plain_len, (u8 *) temp_buf); // copy the base32 content into our plain buffer strncpy ((char *) plain_ptr, (char *) temp_ptr, sizeof (plain_buf)); From 55d56baaa58fe8dfb05283c2b354ce87f223b7f1 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Mon, 22 Oct 2018 08:25:05 -0400 Subject: [PATCH 31/32] Fix digest to allow auto optimized code by compiler --- src/interface.c | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/src/interface.c b/src/interface.c index c4f08b6b4..ff2275270 100644 --- a/src/interface.c +++ b/src/interface.c @@ -5195,7 +5195,7 @@ int totp_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE // now we need to reduce our hash into a token int otp_code = hc_strtoul ((const char *) input_buf, NULL, 10); - digest[1] = otp_code; + digest[0] = otp_code; u8 *salt_pos = token.buf[1]; @@ -22338,7 +22338,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le u64 tmp_salt_buf = (((u64) (salt.salt_buf[2])) << 32) | ((u64) (salt.salt_buf[3])); - snprintf (out_buf, out_len - 1, "%06d:%" PRIu64, digest_buf[1], tmp_salt_buf); + snprintf (out_buf, out_len - 1, "%06d:%" PRIu64, digest_buf[0], tmp_salt_buf); } else if (hash_mode == 99999) { @@ -27636,14 +27636,14 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) | OPTS_TYPE_PT_ADDBITS15 | OPTS_TYPE_PT_NEVERCRACK; hashconfig->kern_type = KERN_TYPE_TOTP_HMACSHA1; - hashconfig->dgst_size = DGST_SIZE_4_5; + hashconfig->dgst_size = DGST_SIZE_4_4; hashconfig->parse_func = totp_parse_hash; hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_NOT_ITERATED; - hashconfig->dgst_pos0 = 1; - hashconfig->dgst_pos1 = 2; - hashconfig->dgst_pos2 = 3; - hashconfig->dgst_pos3 = 4; + hashconfig->dgst_pos0 = 0; + hashconfig->dgst_pos1 = 1; + hashconfig->dgst_pos2 = 2; + hashconfig->dgst_pos3 = 3; hashconfig->st_hash = ST_HASH_18100; hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN; break; From 8c5c225d8f009148decdf7097365e44d9928fe0a Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Mon, 22 Oct 2018 13:27:35 -0400 Subject: [PATCH 32/32] Optimize performance on NVIDIA GTX --- OpenCL/m18100_a0-pure.cl | 70 +++++++++++++++++----------------- OpenCL/m18100_a1-pure.cl | 70 +++++++++++++++++----------------- OpenCL/m18100_a3-pure.cl | 82 +++++++++++++++++++--------------------- 3 files changed, 110 insertions(+), 112 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index bb0894e19..9c589b1b3 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -59,29 +59,30 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru sha1_hmac_final (&ctx); - // calculate the offset using the least 4 bits of the last byte of our hash - const u32x otp_offset = ctx.opad.h[4] & 0xf; - // initialize a buffer for the otp code u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - // on some systems, &3 is faster than %4, so we will use it in our switch() - switch (otp_offset & 3) + switch (ctx.opad.h[4] & 15) { - case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); - break; - case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); - break; - case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); - break; - default: - otp_code = ctx.opad.h[otp_offset/4]; - break; + case 0: otp_code = ctx.opad.h[0]; break; + case 1: otp_code = ctx.opad.h[0] << 8 | ctx.opad.h[1] >> 24; break; + case 2: otp_code = ctx.opad.h[0] << 16 | ctx.opad.h[1] >> 16; break; + case 3: otp_code = ctx.opad.h[0] << 24 | ctx.opad.h[1] >> 8; break; + case 4: otp_code = ctx.opad.h[1]; break; + case 5: otp_code = ctx.opad.h[1] << 8 | ctx.opad.h[2] >> 24; break; + case 6: otp_code = ctx.opad.h[1] << 16 | ctx.opad.h[2] >> 16; break; + case 7: otp_code = ctx.opad.h[1] << 24 | ctx.opad.h[2] >> 8; break; + case 8: otp_code = ctx.opad.h[2]; break; + case 9: otp_code = ctx.opad.h[2] << 8 | ctx.opad.h[3] >> 24; break; + case 10: otp_code = ctx.opad.h[2] << 16 | ctx.opad.h[3] >> 16; break; + case 11: otp_code = ctx.opad.h[2] << 24 | ctx.opad.h[3] >> 8; break; + case 12: otp_code = ctx.opad.h[3]; break; + case 13: otp_code = ctx.opad.h[3] << 8 | ctx.opad.h[4] >> 24; break; + case 14: otp_code = ctx.opad.h[3] << 16 | ctx.opad.h[4] >> 16; break; + case 15: otp_code = ctx.opad.h[3] << 24 | ctx.opad.h[4] >> 8; break; } + // take only the lower 31 bits otp_code &= 0x7fffffff; // we want to generate only 6 digits of code @@ -152,29 +153,30 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru sha1_hmac_final (&ctx); - // calculate the offset using the least 4 bits of the last byte of our hash - const u32x otp_offset = ctx.opad.h[4] & 0xf; - // initialize a buffer for the otp code u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - // on some systems, &3 is faster than %4, so we will use it in our switch() - switch (otp_offset & 3) + switch (ctx.opad.h[4] & 15) { - case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); - break; - case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); - break; - case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); - break; - default: - otp_code = ctx.opad.h[otp_offset/4]; - break; + case 0: otp_code = ctx.opad.h[0]; break; + case 1: otp_code = ctx.opad.h[0] << 8 | ctx.opad.h[1] >> 24; break; + case 2: otp_code = ctx.opad.h[0] << 16 | ctx.opad.h[1] >> 16; break; + case 3: otp_code = ctx.opad.h[0] << 24 | ctx.opad.h[1] >> 8; break; + case 4: otp_code = ctx.opad.h[1]; break; + case 5: otp_code = ctx.opad.h[1] << 8 | ctx.opad.h[2] >> 24; break; + case 6: otp_code = ctx.opad.h[1] << 16 | ctx.opad.h[2] >> 16; break; + case 7: otp_code = ctx.opad.h[1] << 24 | ctx.opad.h[2] >> 8; break; + case 8: otp_code = ctx.opad.h[2]; break; + case 9: otp_code = ctx.opad.h[2] << 8 | ctx.opad.h[3] >> 24; break; + case 10: otp_code = ctx.opad.h[2] << 16 | ctx.opad.h[3] >> 16; break; + case 11: otp_code = ctx.opad.h[2] << 24 | ctx.opad.h[3] >> 8; break; + case 12: otp_code = ctx.opad.h[3]; break; + case 13: otp_code = ctx.opad.h[3] << 8 | ctx.opad.h[4] >> 24; break; + case 14: otp_code = ctx.opad.h[3] << 16 | ctx.opad.h[4] >> 16; break; + case 15: otp_code = ctx.opad.h[3] << 24 | ctx.opad.h[4] >> 8; break; } + // take only the lower 31 bits otp_code &= 0x7fffffff; // we want to generate only 6 digits of code diff --git a/OpenCL/m18100_a1-pure.cl b/OpenCL/m18100_a1-pure.cl index 815eec778..940a7c5da 100644 --- a/OpenCL/m18100_a1-pure.cl +++ b/OpenCL/m18100_a1-pure.cl @@ -82,29 +82,30 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final (&ctx); - // calculate the offset using the least 4 bits of the last byte of our hash - const u32x otp_offset = ctx.opad.h[4] & 0xf; - // initialize a buffer for the otp code u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - // on some systems, &3 is faster than %4, so we will use it in our switch() - switch (otp_offset & 3) + switch (ctx.opad.h[4] & 15) { - case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); - break; - case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); - break; - case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); - break; - default: - otp_code = ctx.opad.h[otp_offset/4]; - break; + case 0: otp_code = ctx.opad.h[0]; break; + case 1: otp_code = ctx.opad.h[0] << 8 | ctx.opad.h[1] >> 24; break; + case 2: otp_code = ctx.opad.h[0] << 16 | ctx.opad.h[1] >> 16; break; + case 3: otp_code = ctx.opad.h[0] << 24 | ctx.opad.h[1] >> 8; break; + case 4: otp_code = ctx.opad.h[1]; break; + case 5: otp_code = ctx.opad.h[1] << 8 | ctx.opad.h[2] >> 24; break; + case 6: otp_code = ctx.opad.h[1] << 16 | ctx.opad.h[2] >> 16; break; + case 7: otp_code = ctx.opad.h[1] << 24 | ctx.opad.h[2] >> 8; break; + case 8: otp_code = ctx.opad.h[2]; break; + case 9: otp_code = ctx.opad.h[2] << 8 | ctx.opad.h[3] >> 24; break; + case 10: otp_code = ctx.opad.h[2] << 16 | ctx.opad.h[3] >> 16; break; + case 11: otp_code = ctx.opad.h[2] << 24 | ctx.opad.h[3] >> 8; break; + case 12: otp_code = ctx.opad.h[3]; break; + case 13: otp_code = ctx.opad.h[3] << 8 | ctx.opad.h[4] >> 24; break; + case 14: otp_code = ctx.opad.h[3] << 16 | ctx.opad.h[4] >> 16; break; + case 15: otp_code = ctx.opad.h[3] << 24 | ctx.opad.h[4] >> 8; break; } + // take only the lower 31 bits otp_code &= 0x7fffffff; // we want to generate only 6 digits of code @@ -200,29 +201,30 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final (&ctx); - // calculate the offset using the least 4 bits of the last byte of our hash - const u32x otp_offset = ctx.opad.h[4] & 0xf; - // initialize a buffer for the otp code u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - // on some systems, &3 is faster than %4, so we will use it in our switch() - switch (otp_offset & 3) + switch (ctx.opad.h[4] & 15) { - case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); - break; - case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); - break; - case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); - break; - default: - otp_code = ctx.opad.h[otp_offset/4]; - break; + case 0: otp_code = ctx.opad.h[0]; break; + case 1: otp_code = ctx.opad.h[0] << 8 | ctx.opad.h[1] >> 24; break; + case 2: otp_code = ctx.opad.h[0] << 16 | ctx.opad.h[1] >> 16; break; + case 3: otp_code = ctx.opad.h[0] << 24 | ctx.opad.h[1] >> 8; break; + case 4: otp_code = ctx.opad.h[1]; break; + case 5: otp_code = ctx.opad.h[1] << 8 | ctx.opad.h[2] >> 24; break; + case 6: otp_code = ctx.opad.h[1] << 16 | ctx.opad.h[2] >> 16; break; + case 7: otp_code = ctx.opad.h[1] << 24 | ctx.opad.h[2] >> 8; break; + case 8: otp_code = ctx.opad.h[2]; break; + case 9: otp_code = ctx.opad.h[2] << 8 | ctx.opad.h[3] >> 24; break; + case 10: otp_code = ctx.opad.h[2] << 16 | ctx.opad.h[3] >> 16; break; + case 11: otp_code = ctx.opad.h[2] << 24 | ctx.opad.h[3] >> 8; break; + case 12: otp_code = ctx.opad.h[3]; break; + case 13: otp_code = ctx.opad.h[3] << 8 | ctx.opad.h[4] >> 24; break; + case 14: otp_code = ctx.opad.h[3] << 16 | ctx.opad.h[4] >> 16; break; + case 15: otp_code = ctx.opad.h[3] << 24 | ctx.opad.h[4] >> 8; break; } + // take only the lower 31 bits otp_code &= 0x7fffffff; // we want to generate only 6 digits of code diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index a71516aa8..d622be763 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -68,39 +68,36 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final_vector (&ctx); - // calculate the offset using the least 4 bits of the last byte of our hash - const u32x otp_offset = ctx.opad.h[4] & 0xf; - // initialize a buffer for the otp code u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - // on some systems, &3 is faster than %4, so we will use it in our switch() - switch (otp_offset & 3) + switch (ctx.opad.h[4] & 15) { - case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); - break; - case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); - break; - case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); - break; - default: - otp_code = ctx.opad.h[otp_offset/4]; - break; + case 0: otp_code = ctx.opad.h[0]; break; + case 1: otp_code = ctx.opad.h[0] << 8 | ctx.opad.h[1] >> 24; break; + case 2: otp_code = ctx.opad.h[0] << 16 | ctx.opad.h[1] >> 16; break; + case 3: otp_code = ctx.opad.h[0] << 24 | ctx.opad.h[1] >> 8; break; + case 4: otp_code = ctx.opad.h[1]; break; + case 5: otp_code = ctx.opad.h[1] << 8 | ctx.opad.h[2] >> 24; break; + case 6: otp_code = ctx.opad.h[1] << 16 | ctx.opad.h[2] >> 16; break; + case 7: otp_code = ctx.opad.h[1] << 24 | ctx.opad.h[2] >> 8; break; + case 8: otp_code = ctx.opad.h[2]; break; + case 9: otp_code = ctx.opad.h[2] << 8 | ctx.opad.h[3] >> 24; break; + case 10: otp_code = ctx.opad.h[2] << 16 | ctx.opad.h[3] >> 16; break; + case 11: otp_code = ctx.opad.h[2] << 24 | ctx.opad.h[3] >> 8; break; + case 12: otp_code = ctx.opad.h[3]; break; + case 13: otp_code = ctx.opad.h[3] << 8 | ctx.opad.h[4] >> 24; break; + case 14: otp_code = ctx.opad.h[3] << 16 | ctx.opad.h[4] >> 16; break; + case 15: otp_code = ctx.opad.h[3] << 24 | ctx.opad.h[4] >> 8; break; } + // take only the lower 31 bits otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code otp_code %= 1000000; - const u32x r0 = ctx.opad.h[DGST_R0]; - const u32x r1 = ctx.opad.h[DGST_R1]; - const u32x r2 = ctx.opad.h[DGST_R2]; - const u32x r3 = ctx.opad.h[DGST_R3]; - COMPARE_M_SIMD (otp_code, 0, 0, 0); } } @@ -172,39 +169,36 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha1_hmac_final_vector (&ctx); - // calculate the offset using the least 4 bits of the last byte of our hash - const u32x otp_offset = ctx.opad.h[4] & 0xf; - // initialize a buffer for the otp code u32 otp_code = 0; // grab 4 consecutive bytes of the hash, starting at offset - // on some systems, &3 is faster than %4, so we will use it in our switch() - switch (otp_offset & 3) + switch (ctx.opad.h[4] & 15) { - case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); - break; - case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); - break; - case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); - break; - default: - otp_code = ctx.opad.h[otp_offset/4]; - break; + case 0: otp_code = ctx.opad.h[0]; break; + case 1: otp_code = ctx.opad.h[0] << 8 | ctx.opad.h[1] >> 24; break; + case 2: otp_code = ctx.opad.h[0] << 16 | ctx.opad.h[1] >> 16; break; + case 3: otp_code = ctx.opad.h[0] << 24 | ctx.opad.h[1] >> 8; break; + case 4: otp_code = ctx.opad.h[1]; break; + case 5: otp_code = ctx.opad.h[1] << 8 | ctx.opad.h[2] >> 24; break; + case 6: otp_code = ctx.opad.h[1] << 16 | ctx.opad.h[2] >> 16; break; + case 7: otp_code = ctx.opad.h[1] << 24 | ctx.opad.h[2] >> 8; break; + case 8: otp_code = ctx.opad.h[2]; break; + case 9: otp_code = ctx.opad.h[2] << 8 | ctx.opad.h[3] >> 24; break; + case 10: otp_code = ctx.opad.h[2] << 16 | ctx.opad.h[3] >> 16; break; + case 11: otp_code = ctx.opad.h[2] << 24 | ctx.opad.h[3] >> 8; break; + case 12: otp_code = ctx.opad.h[3]; break; + case 13: otp_code = ctx.opad.h[3] << 8 | ctx.opad.h[4] >> 24; break; + case 14: otp_code = ctx.opad.h[3] << 16 | ctx.opad.h[4] >> 16; break; + case 15: otp_code = ctx.opad.h[3] << 24 | ctx.opad.h[4] >> 8; break; } + // take only the lower 31 bits otp_code &= 0x7fffffff; + // we want to generate only 6 digits of code otp_code %= 1000000; - const u32x r0 = ctx.opad.h[DGST_R0]; - const u32x r1 = ctx.opad.h[DGST_R1]; - const u32x r2 = ctx.opad.h[DGST_R2]; - const u32x r3 = ctx.opad.h[DGST_R3]; - COMPARE_S_SIMD (otp_code, 0, 0, 0); } }