From 29d331ee171fa51680a98128c6c250bffdafc036 Mon Sep 17 00:00:00 2001 From: Fist0urs Date: Tue, 2 May 2017 23:56:00 +0200 Subject: [PATCH] hmac-sha1 + SID almost working, padding problem --- OpenCL/m15300.cl | 50 ++++++++++++++++++++++++++++-------------- src/interface.c | 57 ++++++++++++++++++++++-------------------------- 2 files changed, 60 insertions(+), 47 deletions(-) diff --git a/OpenCL/m15300.cl b/OpenCL/m15300.cl index 8cf3bd430..05a46e74b 100644 --- a/OpenCL/m15300.cl +++ b/OpenCL/m15300.cl @@ -1071,7 +1071,6 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul w0[2] = pws[gid].i[ 2]; w0[3] = pws[gid].i[ 3]; - u32 w1[4]; w1[0] = pws[gid].i[ 4]; @@ -1226,22 +1225,41 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul * hmac1 */ - w0[0] = 0xdddddddd; - w0[1] = 0xdddddddd; - w0[2] = 0xdddddddd; - w0[3] = 0xdddddddd; - w1[0] = 0xdddddddd; - w1[1] = 0x80000000; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; + w0[0] = esalt_bufs[digests_offset].SID[ 0]; + w0[1] = esalt_bufs[digests_offset].SID[ 1]; + w0[2] = esalt_bufs[digests_offset].SID[ 2]; + w0[3] = esalt_bufs[digests_offset].SID[ 3]; + w1[0] = esalt_bufs[digests_offset].SID[ 4]; + w1[1] = esalt_bufs[digests_offset].SID[ 5]; + w1[2] = esalt_bufs[digests_offset].SID[ 6]; + w1[3] = esalt_bufs[digests_offset].SID[ 7]; + w2[0] = esalt_bufs[digests_offset].SID[ 8]; + w2[1] = esalt_bufs[digests_offset].SID[ 9]; + w2[2] = esalt_bufs[digests_offset].SID[10]; + w2[3] = esalt_bufs[digests_offset].SID[11]; + w3[0] = esalt_bufs[digests_offset].SID[12]; + w3[1] = esalt_bufs[digests_offset].SID[13]; + w3[2] = esalt_bufs[digests_offset].SID[14]; + w3[3] = esalt_bufs[digests_offset].SID[15]; + + sha1_transform_S (w0, w1, w2, w3, ipad); + + w0[0] = esalt_bufs[digests_offset].SID[16 + 0]; + w0[1] = esalt_bufs[digests_offset].SID[16 + 1]; + w0[2] = esalt_bufs[digests_offset].SID[16 + 2]; + w0[3] = esalt_bufs[digests_offset].SID[16 + 3]; + w1[0] = esalt_bufs[digests_offset].SID[16 + 4]; + w1[1] = esalt_bufs[digests_offset].SID[16 + 5]; + w1[2] = esalt_bufs[digests_offset].SID[16 + 6]; + w1[3] = esalt_bufs[digests_offset].SID[16 + 7]; + w2[0] = esalt_bufs[digests_offset].SID[16 + 8]; + w2[1] = esalt_bufs[digests_offset].SID[16 + 9]; + w2[2] = esalt_bufs[digests_offset].SID[16 + 10]; + w2[3] = esalt_bufs[digests_offset].SID[16 + 11]; + w3[0] = esalt_bufs[digests_offset].SID[16 + 12]; + w3[1] = esalt_bufs[digests_offset].SID[16 + 13]; w3[2] = 0; - w3[3] = (64 + 20) * 8; + w3[3] = (64 + salt_len) * 8; u32 key[5]; diff --git a/src/interface.c b/src/interface.c index 2994ca83a..4cb0b5a64 100644 --- a/src/interface.c +++ b/src/interface.c @@ -2208,24 +2208,8 @@ static u32 parse_and_store_salt (u8 *out, u8 *in, u32 salt_len, MAYBE_UNUSED con if (hashconfig->opts_type & OPTS_TYPE_ST_UNICODE) { - if (salt_len < 52) + if (salt_len < 20) { - tmp_u32[25] = ((tmp_u32[12] >> 8) & 0x00FF0000) | ((tmp_u32[12] >> 16) & 0x000000FF); - tmp_u32[24] = ((tmp_u32[12] << 8) & 0x00FF0000) | ((tmp_u32[12] >> 0) & 0x000000FF); - tmp_u32[23] = ((tmp_u32[11] >> 8) & 0x00FF0000) | ((tmp_u32[11] >> 16) & 0x000000FF); - tmp_u32[22] = ((tmp_u32[11] << 8) & 0x00FF0000) | ((tmp_u32[11] >> 0) & 0x000000FF); - tmp_u32[21] = ((tmp_u32[10] >> 8) & 0x00FF0000) | ((tmp_u32[10] >> 16) & 0x000000FF); - tmp_u32[20] = ((tmp_u32[10] << 8) & 0x00FF0000) | ((tmp_u32[10] >> 0) & 0x000000FF); - tmp_u32[19] = ((tmp_u32[ 9] >> 8) & 0x00FF0000) | ((tmp_u32[ 9] >> 16) & 0x000000FF); - tmp_u32[18] = ((tmp_u32[ 9] << 8) & 0x00FF0000) | ((tmp_u32[ 9] >> 0) & 0x000000FF); - tmp_u32[17] = ((tmp_u32[ 8] >> 8) & 0x00FF0000) | ((tmp_u32[ 8] >> 16) & 0x000000FF); - tmp_u32[16] = ((tmp_u32[ 8] << 8) & 0x00FF0000) | ((tmp_u32[ 8] >> 0) & 0x000000FF); - tmp_u32[15] = ((tmp_u32[ 7] >> 8) & 0x00FF0000) | ((tmp_u32[ 7] >> 16) & 0x000000FF); - tmp_u32[14] = ((tmp_u32[ 7] << 8) & 0x00FF0000) | ((tmp_u32[ 7] >> 0) & 0x000000FF); - tmp_u32[13] = ((tmp_u32[ 6] >> 8) & 0x00FF0000) | ((tmp_u32[ 6] >> 16) & 0x000000FF); - tmp_u32[12] = ((tmp_u32[ 6] << 8) & 0x00FF0000) | ((tmp_u32[ 6] >> 0) & 0x000000FF); - tmp_u32[11] = ((tmp_u32[ 5] >> 8) & 0x00FF0000) | ((tmp_u32[ 5] >> 16) & 0x000000FF); - tmp_u32[10] = ((tmp_u32[ 5] << 8) & 0x00FF0000) | ((tmp_u32[ 5] >> 0) & 0x000000FF); tmp_u32[ 9] = ((tmp_u32[ 4] >> 8) & 0x00FF0000) | ((tmp_u32[ 4] >> 16) & 0x000000FF); tmp_u32[ 8] = ((tmp_u32[ 4] << 8) & 0x00FF0000) | ((tmp_u32[ 4] >> 0) & 0x000000FF); tmp_u32[ 7] = ((tmp_u32[ 3] >> 8) & 0x00FF0000) | ((tmp_u32[ 3] >> 16) & 0x000000FF); @@ -2900,10 +2884,8 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN u8 *contents_len_pos; u8 *contents_pos; - u32 salt_len; - u8 *salt_pos; - int dbg = 1; + version_pos = input_buf + 8 + 1; dpapimk->version = atoll ((const char *) version_pos); @@ -3024,26 +3006,39 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN digest[2] = dpapimk->iv[2]; digest[3] = dpapimk->iv[3]; - salt_len = cipher_algo_pos - 1 - SID_pos; + u32 SID_len = cipher_algo_pos - 1 - SID_pos; + + u8 *SID_unicode = (u8 *) hcmalloc ((SID_len + 1) * 2); + + memset (SID_unicode, 0, (SID_len + 1) * 2); if (dbg) - printf("SID_len_before: %d\n", salt_len); + printf("SID_len_before: %d\n", SID_len); /* Specific to DPAPI, SID + '\0' */ - salt_len = parse_and_store_salt (dpapimk->SID, dpapimk->SID_tmp, salt_len + 1, hashconfig); - - if (dbg) + for (u32 i = 0; i < SID_len; i += 1) { - printf("SID_hex :%02x%02x%02x%02x\n", dpapimk->SID[0],dpapimk->SID[1],dpapimk->SID[2],dpapimk->SID[3]); - printf("SID_hex :%c%c%c%c\n", dpapimk->SID[0],dpapimk->SID[1],dpapimk->SID[2],dpapimk->SID[3]); + SID_unicode[i*2] = SID_pos[i]; } - if (salt_len == UINT_MAX) return (PARSER_SALT_LENGTH); + //SID_unicode[SID_len*2] = '\0'; + //SID_unicode[SID_len*2+1] = '\0'; - salt->salt_len = salt_len; + salt->salt_len = (SID_len) * 2; + + memcpy ((u8 *) dpapimk->SID, SID_unicode, salt->salt_len); - if (dbg) - printf("SID_len_after: %d\n", salt_len); + u32 max = salt->salt_len / 4; + + if (salt->salt_len % 4) max++; + + for (u32 i = 0; i < max; i++) + { + dpapimk->SID[i] = byte_swap_32 (dpapimk->SID[i]); + } + dpapimk->SID[max] = 0x80000000; + + hcfree(SID_unicode); return (PARSER_OK); }