github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2025-01-10 17:16:22 +01:00
Code Releases Activity

Extended IKE PSK md5/sha1 (-m 5300/5400) to print hashes correctly

Browse Source
This commit is contained in:
Sein Coray 2018-12-05 12:57:54 +01:00
parent 0076f4075b
commit c941e55a35
No known key found for this signature in database
GPG Key ID: 44C4180EA69758EC
15 changed files with 58 additions and 48 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
OpenCL/inc_types.cl
View File

@ -1259,7 +1259,7 @@ typedef struct ikepsk
u32 nr_len;
u32 msg_buf[128];
u32 msg_len;
u32 msg_len[6];
} ikepsk_t;

8
OpenCL/m05300_a0-optimized.cl
View File

@ -155,7 +155,7 @@ __kernel void m05300_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* loop
@ -222,7 +222,7 @@ __kernel void m05300_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0[0] = s_msg_buf[off + 0];
w0[1] = s_msg_buf[off + 1];
@ -330,7 +330,7 @@ __kernel void m05300_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* digest
@ -409,7 +409,7 @@ __kernel void m05300_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0[0] = s_msg_buf[off + 0];
w0[1] = s_msg_buf[off + 1];

4
OpenCL/m05300_a0-pure.cl
View File

@ -76,7 +76,7 @@ __kernel void m05300_mxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
md5_hmac_final (&ctx);
@ -162,7 +162,7 @@ __kernel void m05300_sxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
md5_hmac_final (&ctx);

8
OpenCL/m05300_a1-optimized.cl
View File

@ -153,7 +153,7 @@ __kernel void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* loop
@ -280,7 +280,7 @@ __kernel void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0[0] = s_msg_buf[off + 0];
w0[1] = s_msg_buf[off + 1];
@ -388,7 +388,7 @@ __kernel void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* digest
@ -527,7 +527,7 @@ __kernel void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0[0] = s_msg_buf[off + 0];
w0[1] = s_msg_buf[off + 1];

4
OpenCL/m05300_a1-pure.cl
View File

@ -99,7 +99,7 @@ __kernel void m05300_mxx (KERN_ATTR_ESALT (ikepsk_t))
md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
md5_hmac_final (&ctx);
@ -210,7 +210,7 @@ __kernel void m05300_sxx (KERN_ATTR_ESALT (ikepsk_t))
md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
md5_hmac_final (&ctx);

8
OpenCL/m05300_a3-optimized.cl
View File

@ -112,7 +112,7 @@ DECLSPEC void m05300m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* loop
@ -200,7 +200,7 @@ DECLSPEC void m05300m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0_t[0] = s_msg_buf[off + 0];
w0_t[1] = s_msg_buf[off + 1];
@ -259,7 +259,7 @@ DECLSPEC void m05300s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* digest
@ -359,7 +359,7 @@ DECLSPEC void m05300s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0_t[0] = s_msg_buf[off + 0];
w0_t[1] = s_msg_buf[off + 1];

4
OpenCL/m05300_a3-pure.cl
View File

@ -85,7 +85,7 @@ __kernel void m05300_mxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
md5_hmac_final (&ctx);
@ -182,7 +182,7 @@ __kernel void m05300_sxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
md5_hmac_final (&ctx);

8
OpenCL/m05400_a0-optimized.cl
View File

@ -159,7 +159,7 @@ __kernel void m05400_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* loop
@ -235,7 +235,7 @@ __kernel void m05400_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0[0] = s_msg_buf[off + 0];
w0[1] = s_msg_buf[off + 1];
@ -343,7 +343,7 @@ __kernel void m05400_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* digest
@ -431,7 +431,7 @@ __kernel void m05400_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0[0] = s_msg_buf[off + 0];
w0[1] = s_msg_buf[off + 1];

4
OpenCL/m05400_a0-pure.cl
View File

@ -76,7 +76,7 @@ __kernel void m05400_mxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
sha1_hmac_final (&ctx);
@ -162,7 +162,7 @@ __kernel void m05400_sxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
sha1_hmac_final (&ctx);

8
OpenCL/m05400_a1-optimized.cl
View File

@ -157,7 +157,7 @@ __kernel void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* loop
@ -301,7 +301,7 @@ __kernel void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0[0] = s_msg_buf[off + 0];
w0[1] = s_msg_buf[off + 1];
@ -409,7 +409,7 @@ __kernel void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* digest
@ -565,7 +565,7 @@ __kernel void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0[0] = s_msg_buf[off + 0];
w0[1] = s_msg_buf[off + 1];

4
OpenCL/m05400_a1-pure.cl
View File

@ -99,7 +99,7 @@ __kernel void m05400_mxx (KERN_ATTR_ESALT (ikepsk_t))
sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
sha1_hmac_final (&ctx);
@ -210,7 +210,7 @@ __kernel void m05400_sxx (KERN_ATTR_ESALT (ikepsk_t))
sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
sha1_hmac_final (&ctx);

8
OpenCL/m05400_a3-optimized.cl
View File

@ -116,7 +116,7 @@ DECLSPEC void m05400m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* loop
@ -204,7 +204,7 @@ DECLSPEC void m05400m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0_t[0] = s_msg_buf[off + 0];
w0_t[1] = s_msg_buf[off + 1];
@ -263,7 +263,7 @@ DECLSPEC void m05400s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
*/
const u32 nr_len = esalt_bufs[digests_offset].nr_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len;
const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
/**
* digest
@ -363,7 +363,7 @@ DECLSPEC void m05400s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
int left;
int off;
for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
{
w0_t[0] = s_msg_buf[off + 0];
w0_t[1] = s_msg_buf[off + 1];

4
OpenCL/m05400_a3-pure.cl
View File

@ -85,7 +85,7 @@ __kernel void m05400_mxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
sha1_hmac_final (&ctx);
@ -182,7 +182,7 @@ __kernel void m05400_sxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
sha1_hmac_final (&ctx);

2
include/interface.h
View File

@ -250,7 +250,7 @@ typedef struct ikepsk
u32 nr_len;
u32 msg_buf[128];
u32 msg_len;
u32 msg_len[6];
} ikepsk_t;

30
src/interface.c
View File

@ -6705,10 +6705,15 @@ int ikepsk_md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE
if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
ikepsk->msg_len = (token.len[0] + token.len[1] + token.len[2] + token.len[3] + token.len[4] + token.len[5]) / 2;
ikepsk->msg_len[0] = token.len[0] / 2;
ikepsk->msg_len[1] = ikepsk->msg_len[0] + token.len[1] / 2;
ikepsk->msg_len[2] = ikepsk->msg_len[1] + token.len[2] / 2;
ikepsk->msg_len[3] = ikepsk->msg_len[2] + token.len[3] / 2;
ikepsk->msg_len[4] = ikepsk->msg_len[3] + token.len[4] / 2;
ikepsk->msg_len[5] = ikepsk->msg_len[4] + token.len[5] / 2;
ikepsk->nr_len = (token.len[6] + token.len[7]) / 2;
if (ikepsk->msg_len > 512) return (PARSER_SALT_LENGTH);
if (ikepsk->msg_len[5] > 512) return (PARSER_SALT_LENGTH);
if (ikepsk->nr_len > 64) return (PARSER_SALT_LENGTH);
u8 *ptr1 = (u8 *) ikepsk->msg_buf;
@ -6821,10 +6826,15 @@ int ikepsk_sha1_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYB
if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
ikepsk->msg_len = (token.len[0] + token.len[1] + token.len[2] + token.len[3] + token.len[4] + token.len[5]) / 2;
ikepsk->msg_len[0] = token.len[0] / 2;
ikepsk->msg_len[1] = ikepsk->msg_len[0] + token.len[1] / 2;
ikepsk->msg_len[2] = ikepsk->msg_len[1] + token.len[2] / 2;
ikepsk->msg_len[3] = ikepsk->msg_len[2] + token.len[3] / 2;
ikepsk->msg_len[4] = ikepsk->msg_len[3] + token.len[4] / 2;
ikepsk->msg_len[5] = ikepsk->msg_len[4] + token.len[5] / 2;
ikepsk->nr_len = (token.len[6] + token.len[7]) / 2;
if (ikepsk->msg_len > 512) return (PARSER_SALT_LENGTH);
if (ikepsk->msg_len[5] > 512) return (PARSER_SALT_LENGTH);
if (ikepsk->nr_len > 64) return (PARSER_SALT_LENGTH);
u8 *ptr1 = (u8 *) ikepsk->msg_buf;
@ -20037,11 +20047,11 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
// msg_buf
u32 ikepsk_msg_len = ikepsk->msg_len / 4;
u32 ikepsk_msg_len = ikepsk->msg_len[5] / 4;
for (u32 i = 0; i < ikepsk_msg_len; i++)
{
if ((i == 32) || (i == 64) || (i == 66) || (i == 68) || (i == 108))
if ((i == ikepsk->msg_len[0] / 4) || (i == ikepsk->msg_len[1] / 4) || (i == ikepsk->msg_len[2] / 4) || (i == ikepsk->msg_len[3] / 4) || (i == ikepsk->msg_len[4] / 4))
{
snprintf (out_buf, buf_len, ":");
@ -20103,11 +20113,11 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
// msg_buf
u32 ikepsk_msg_len = ikepsk->msg_len / 4;
u32 ikepsk_msg_len = ikepsk->msg_len[5] / 4;
for (u32 i = 0; i < ikepsk_msg_len; i++)
{
if ((i == 32) || (i == 64) || (i == 66) || (i == 68) || (i == 108))
if ((i == ikepsk->msg_len[0] / 4) || (i == ikepsk->msg_len[1] / 4) || (i == ikepsk->msg_len[2] / 4) || (i == ikepsk->msg_len[3] / 4) || (i == ikepsk->msg_len[4] / 4))
{
snprintf (out_buf, buf_len, ":");
@ -29509,10 +29519,10 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo
case 2501: ((wpa_eapol_t *) esalt)->eapol_len = 128;
break;
case 5300: ((ikepsk_t *) esalt)->nr_len = 1;
((ikepsk_t *) esalt)->msg_len = 1;
((ikepsk_t *) esalt)->msg_len[5] = 1;
break;
case 5400: ((ikepsk_t *) esalt)->nr_len = 1;
((ikepsk_t *) esalt)->msg_len = 1;
((ikepsk_t *) esalt)->msg_len[5] = 1;
break;
case 5500: ((netntlm_t *) esalt)->user_len = 1;
((netntlm_t *) esalt)->domain_len = 1;