From 88ebca40b8a52c16fd0d9d24f7a2f63d8d7f4400 Mon Sep 17 00:00:00 2001 From: jsteube Date: Wed, 25 Jul 2018 16:46:06 +0200 Subject: [PATCH] Added hash-mode 16800 = WPA-PMKID-PBKDF2 Added hash-mode 16801 = WPA-PMKID-PMK Renamed lot's of existing WPA related variables to WPA-EAPOL in order to distinguish them with WPA-PMKID variables Renamed WPA/WPA2 to WPA-EAPOL-PBKDF2 Renamed WPA/WPA2 PMK to WPA-EAPOL-PMK --- OpenCL/inc_types.cl | 23 +- OpenCL/m02500-pure.cl | 338 +++++++++--------- OpenCL/m02501-pure.cl | 338 +++++++++--------- OpenCL/m06400-pure.cl | 6 +- OpenCL/m16800-pure.cl | 254 ++++++++++++++ OpenCL/m16801-pure.cl | 129 +++++++ docs/changes.txt | 2 + docs/readme.txt | 6 +- extra/tab_completion/hashcat.sh | 2 +- include/interface.h | 41 ++- src/hashes.c | 42 +-- src/interface.c | 596 +++++++++++++++++++++++++------- src/opencl.c | 10 +- src/potfile.c | 2 +- src/status.c | 26 +- src/usage.c | 6 +- tools/test.pl | 118 ++++++- tools/test.sh | 159 ++++++++- 18 files changed, 1570 insertions(+), 528 deletions(-) create mode 100644 OpenCL/m16800-pure.cl create mode 100644 OpenCL/m16801-pure.cl diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index fd611e2f8..a2e9fcb47 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -1194,7 +1194,7 @@ typedef struct pdf } pdf_t; -typedef struct wpa +typedef struct wpa_eapol { u32 pke[32]; u32 eapol[64 + 16]; @@ -1215,7 +1215,18 @@ typedef struct wpa int detected_le; int detected_be; -} wpa_t; +} wpa_eapol_t; + +typedef struct wpa_pmkid +{ + u32 pmkid[4]; + u32 pmkid_data[16]; + u8 orig_mac_ap[6]; + u8 orig_mac_sta[6]; + u8 essid_len; + u32 essid_buf[16]; + +} wpa_pmkid_t; typedef struct bitcoin_wallet { @@ -1589,7 +1600,7 @@ typedef struct sha512crypt_tmp } sha512crypt_tmp_t; -typedef struct wpa_tmp +typedef struct wpa_pbkdf2_tmp { u32 ipad[5]; u32 opad[5]; @@ -1597,13 +1608,13 @@ typedef struct wpa_tmp u32 dgst[10]; u32 out[10]; -} wpa_tmp_t; +} wpa_pbkdf2_tmp_t; -typedef struct wpapmk_tmp +typedef struct wpa_pmk_tmp { u32 out[8]; -} wpapmk_tmp_t; +} wpa_pmk_tmp_t; typedef struct bitcoin_wallet_tmp { diff --git a/OpenCL/m02500-pure.cl b/OpenCL/m02500-pure.cl index 7bce2560c..219945cb5 100644 --- a/OpenCL/m02500-pure.cl +++ b/OpenCL/m02500-pure.cl @@ -85,7 +85,7 @@ DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipa sha1_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m02500_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_tmp_t *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 wpa_t *wpa_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 m02500_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { /** * base @@ -157,7 +157,7 @@ __kernel void m02500_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m02500_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_tmp_t *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 wpa_t *wpa_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 m02500_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); @@ -242,12 +242,12 @@ __kernel void m02500_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_tmp_t *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 wpa_t *wpa_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 m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { // not in use here, special case... } -__kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_tmp_t *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 wpa_t *wpa_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 m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); @@ -269,46 +269,46 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 digest_cur = digests_offset + digest_pos; - __global const wpa_t *wpa = &wpa_bufs[digest_cur]; + __global const wpa_eapol_t *wpa_eapol = &wpa_eapol_bufs[digest_cur]; u32 pke[32]; - pke[ 0] = wpa->pke[ 0]; - pke[ 1] = wpa->pke[ 1]; - pke[ 2] = wpa->pke[ 2]; - pke[ 3] = wpa->pke[ 3]; - pke[ 4] = wpa->pke[ 4]; - pke[ 5] = wpa->pke[ 5]; - pke[ 6] = wpa->pke[ 6]; - pke[ 7] = wpa->pke[ 7]; - pke[ 8] = wpa->pke[ 8]; - pke[ 9] = wpa->pke[ 9]; - pke[10] = wpa->pke[10]; - pke[11] = wpa->pke[11]; - pke[12] = wpa->pke[12]; - pke[13] = wpa->pke[13]; - pke[14] = wpa->pke[14]; - pke[15] = wpa->pke[15]; - pke[16] = wpa->pke[16]; - pke[17] = wpa->pke[17]; - pke[18] = wpa->pke[18]; - pke[19] = wpa->pke[19]; - pke[20] = wpa->pke[20]; - pke[21] = wpa->pke[21]; - pke[22] = wpa->pke[22]; - pke[23] = wpa->pke[23]; - pke[24] = wpa->pke[24]; - pke[25] = wpa->pke[25]; - pke[26] = wpa->pke[26]; - pke[27] = wpa->pke[27]; - pke[28] = wpa->pke[28]; - pke[29] = wpa->pke[29]; - pke[30] = wpa->pke[30]; - pke[31] = wpa->pke[31]; + pke[ 0] = wpa_eapol->pke[ 0]; + pke[ 1] = wpa_eapol->pke[ 1]; + pke[ 2] = wpa_eapol->pke[ 2]; + pke[ 3] = wpa_eapol->pke[ 3]; + pke[ 4] = wpa_eapol->pke[ 4]; + pke[ 5] = wpa_eapol->pke[ 5]; + pke[ 6] = wpa_eapol->pke[ 6]; + pke[ 7] = wpa_eapol->pke[ 7]; + pke[ 8] = wpa_eapol->pke[ 8]; + pke[ 9] = wpa_eapol->pke[ 9]; + pke[10] = wpa_eapol->pke[10]; + pke[11] = wpa_eapol->pke[11]; + pke[12] = wpa_eapol->pke[12]; + pke[13] = wpa_eapol->pke[13]; + pke[14] = wpa_eapol->pke[14]; + pke[15] = wpa_eapol->pke[15]; + pke[16] = wpa_eapol->pke[16]; + pke[17] = wpa_eapol->pke[17]; + pke[18] = wpa_eapol->pke[18]; + pke[19] = wpa_eapol->pke[19]; + pke[20] = wpa_eapol->pke[20]; + pke[21] = wpa_eapol->pke[21]; + pke[22] = wpa_eapol->pke[22]; + pke[23] = wpa_eapol->pke[23]; + pke[24] = wpa_eapol->pke[24]; + pke[25] = wpa_eapol->pke[25]; + pke[26] = wpa_eapol->pke[26]; + pke[27] = wpa_eapol->pke[27]; + pke[28] = wpa_eapol->pke[28]; + pke[29] = wpa_eapol->pke[29]; + pke[30] = wpa_eapol->pke[30]; + pke[31] = wpa_eapol->pke[31]; u32 to; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { to = pke[15] << 24 | pke[16] >> 8; @@ -319,9 +319,9 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul | pke[24] >> 8; } - const u32 nonce_error_corrections = wpa->nonce_error_corrections; + const u32 nonce_error_corrections = wpa_eapol->nonce_error_corrections; - if (wpa->detected_le == 1) + if (wpa_eapol->detected_le == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -330,7 +330,7 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul t -= nonce_error_corrections / 2; t += nonce_error_correction; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -447,7 +447,7 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); - md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + md5_hmac_update_global (&ctx2, wpa_eapol->eapol, wpa_eapol->eapol_len); md5_hmac_final (&ctx2); @@ -460,10 +460,10 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -473,7 +473,7 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul } } - if (wpa->detected_be == 1) + if (wpa_eapol->detected_be == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -486,7 +486,7 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul t = swap32_S (t); - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -603,7 +603,7 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); - md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + md5_hmac_update_global (&ctx2, wpa_eapol->eapol, wpa_eapol->eapol_len); md5_hmac_final (&ctx2); @@ -616,10 +616,10 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -630,7 +630,7 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_tmp_t *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 wpa_t *wpa_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 m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); @@ -652,46 +652,46 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 digest_cur = digests_offset + digest_pos; - __global const wpa_t *wpa = &wpa_bufs[digest_cur]; + __global const wpa_eapol_t *wpa_eapol = &wpa_eapol_bufs[digest_cur]; u32 pke[32]; - pke[ 0] = wpa->pke[ 0]; - pke[ 1] = wpa->pke[ 1]; - pke[ 2] = wpa->pke[ 2]; - pke[ 3] = wpa->pke[ 3]; - pke[ 4] = wpa->pke[ 4]; - pke[ 5] = wpa->pke[ 5]; - pke[ 6] = wpa->pke[ 6]; - pke[ 7] = wpa->pke[ 7]; - pke[ 8] = wpa->pke[ 8]; - pke[ 9] = wpa->pke[ 9]; - pke[10] = wpa->pke[10]; - pke[11] = wpa->pke[11]; - pke[12] = wpa->pke[12]; - pke[13] = wpa->pke[13]; - pke[14] = wpa->pke[14]; - pke[15] = wpa->pke[15]; - pke[16] = wpa->pke[16]; - pke[17] = wpa->pke[17]; - pke[18] = wpa->pke[18]; - pke[19] = wpa->pke[19]; - pke[20] = wpa->pke[20]; - pke[21] = wpa->pke[21]; - pke[22] = wpa->pke[22]; - pke[23] = wpa->pke[23]; - pke[24] = wpa->pke[24]; - pke[25] = wpa->pke[25]; - pke[26] = wpa->pke[26]; - pke[27] = wpa->pke[27]; - pke[28] = wpa->pke[28]; - pke[29] = wpa->pke[29]; - pke[30] = wpa->pke[30]; - pke[31] = wpa->pke[31]; + pke[ 0] = wpa_eapol->pke[ 0]; + pke[ 1] = wpa_eapol->pke[ 1]; + pke[ 2] = wpa_eapol->pke[ 2]; + pke[ 3] = wpa_eapol->pke[ 3]; + pke[ 4] = wpa_eapol->pke[ 4]; + pke[ 5] = wpa_eapol->pke[ 5]; + pke[ 6] = wpa_eapol->pke[ 6]; + pke[ 7] = wpa_eapol->pke[ 7]; + pke[ 8] = wpa_eapol->pke[ 8]; + pke[ 9] = wpa_eapol->pke[ 9]; + pke[10] = wpa_eapol->pke[10]; + pke[11] = wpa_eapol->pke[11]; + pke[12] = wpa_eapol->pke[12]; + pke[13] = wpa_eapol->pke[13]; + pke[14] = wpa_eapol->pke[14]; + pke[15] = wpa_eapol->pke[15]; + pke[16] = wpa_eapol->pke[16]; + pke[17] = wpa_eapol->pke[17]; + pke[18] = wpa_eapol->pke[18]; + pke[19] = wpa_eapol->pke[19]; + pke[20] = wpa_eapol->pke[20]; + pke[21] = wpa_eapol->pke[21]; + pke[22] = wpa_eapol->pke[22]; + pke[23] = wpa_eapol->pke[23]; + pke[24] = wpa_eapol->pke[24]; + pke[25] = wpa_eapol->pke[25]; + pke[26] = wpa_eapol->pke[26]; + pke[27] = wpa_eapol->pke[27]; + pke[28] = wpa_eapol->pke[28]; + pke[29] = wpa_eapol->pke[29]; + pke[30] = wpa_eapol->pke[30]; + pke[31] = wpa_eapol->pke[31]; u32 to; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { to = pke[15] << 24 | pke[16] >> 8; @@ -702,9 +702,9 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul | pke[24] >> 8; } - const u32 nonce_error_corrections = wpa->nonce_error_corrections; + const u32 nonce_error_corrections = wpa_eapol->nonce_error_corrections; - if (wpa->detected_le == 1) + if (wpa_eapol->detected_le == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -713,7 +713,7 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul t -= nonce_error_corrections / 2; t += nonce_error_correction; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -830,7 +830,7 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); - sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + sha1_hmac_update_global (&ctx2, wpa_eapol->eapol, wpa_eapol->eapol_len); sha1_hmac_final (&ctx2); @@ -843,10 +843,10 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -856,7 +856,7 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul } } - if (wpa->detected_be == 1) + if (wpa_eapol->detected_be == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -869,7 +869,7 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul t = swap32_S (t); - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -986,7 +986,7 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); - sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + sha1_hmac_update_global (&ctx2, wpa_eapol->eapol, wpa_eapol->eapol_len); sha1_hmac_final (&ctx2); @@ -999,10 +999,10 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -1013,7 +1013,7 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_tmp_t *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 wpa_t *wpa_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 m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); @@ -1087,46 +1087,46 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 digest_cur = digests_offset + digest_pos; - __global const wpa_t *wpa = &wpa_bufs[digest_cur]; + __global const wpa_eapol_t *wpa_eapol = &wpa_eapol_bufs[digest_cur]; u32 pke[32]; - pke[ 0] = wpa->pke[ 0]; - pke[ 1] = wpa->pke[ 1]; - pke[ 2] = wpa->pke[ 2]; - pke[ 3] = wpa->pke[ 3]; - pke[ 4] = wpa->pke[ 4]; - pke[ 5] = wpa->pke[ 5]; - pke[ 6] = wpa->pke[ 6]; - pke[ 7] = wpa->pke[ 7]; - pke[ 8] = wpa->pke[ 8]; - pke[ 9] = wpa->pke[ 9]; - pke[10] = wpa->pke[10]; - pke[11] = wpa->pke[11]; - pke[12] = wpa->pke[12]; - pke[13] = wpa->pke[13]; - pke[14] = wpa->pke[14]; - pke[15] = wpa->pke[15]; - pke[16] = wpa->pke[16]; - pke[17] = wpa->pke[17]; - pke[18] = wpa->pke[18]; - pke[19] = wpa->pke[19]; - pke[20] = wpa->pke[20]; - pke[21] = wpa->pke[21]; - pke[22] = wpa->pke[22]; - pke[23] = wpa->pke[23]; - pke[24] = wpa->pke[24]; - pke[25] = wpa->pke[25]; - pke[26] = wpa->pke[26]; - pke[27] = wpa->pke[27]; - pke[28] = wpa->pke[28]; - pke[29] = wpa->pke[29]; - pke[30] = wpa->pke[30]; - pke[31] = wpa->pke[31]; + pke[ 0] = wpa_eapol->pke[ 0]; + pke[ 1] = wpa_eapol->pke[ 1]; + pke[ 2] = wpa_eapol->pke[ 2]; + pke[ 3] = wpa_eapol->pke[ 3]; + pke[ 4] = wpa_eapol->pke[ 4]; + pke[ 5] = wpa_eapol->pke[ 5]; + pke[ 6] = wpa_eapol->pke[ 6]; + pke[ 7] = wpa_eapol->pke[ 7]; + pke[ 8] = wpa_eapol->pke[ 8]; + pke[ 9] = wpa_eapol->pke[ 9]; + pke[10] = wpa_eapol->pke[10]; + pke[11] = wpa_eapol->pke[11]; + pke[12] = wpa_eapol->pke[12]; + pke[13] = wpa_eapol->pke[13]; + pke[14] = wpa_eapol->pke[14]; + pke[15] = wpa_eapol->pke[15]; + pke[16] = wpa_eapol->pke[16]; + pke[17] = wpa_eapol->pke[17]; + pke[18] = wpa_eapol->pke[18]; + pke[19] = wpa_eapol->pke[19]; + pke[20] = wpa_eapol->pke[20]; + pke[21] = wpa_eapol->pke[21]; + pke[22] = wpa_eapol->pke[22]; + pke[23] = wpa_eapol->pke[23]; + pke[24] = wpa_eapol->pke[24]; + pke[25] = wpa_eapol->pke[25]; + pke[26] = wpa_eapol->pke[26]; + pke[27] = wpa_eapol->pke[27]; + pke[28] = wpa_eapol->pke[28]; + pke[29] = wpa_eapol->pke[29]; + pke[30] = wpa_eapol->pke[30]; + pke[31] = wpa_eapol->pke[31]; u32 to; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { to = pke[15] << 24 | pke[16] >> 8; @@ -1137,9 +1137,9 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul | pke[24] >> 8; } - const u32 nonce_error_corrections = wpa->nonce_error_corrections; + const u32 nonce_error_corrections = wpa_eapol->nonce_error_corrections; - if (wpa->detected_le == 1) + if (wpa_eapol->detected_le == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -1148,7 +1148,7 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul t -= nonce_error_corrections / 2; t += nonce_error_correction; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -1262,20 +1262,20 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul int eapol_left; int eapol_idx; - for (eapol_left = wpa->eapol_len, eapol_idx = 0; eapol_left > 16; eapol_left -= 16, eapol_idx += 4) + for (eapol_left = wpa_eapol->eapol_len, eapol_idx = 0; eapol_left > 16; eapol_left -= 16, eapol_idx += 4) { - m[0] = wpa->eapol[eapol_idx + 0] ^ iv[0]; - m[1] = wpa->eapol[eapol_idx + 1] ^ iv[1]; - m[2] = wpa->eapol[eapol_idx + 2] ^ iv[2]; - m[3] = wpa->eapol[eapol_idx + 3] ^ iv[3]; + m[0] = wpa_eapol->eapol[eapol_idx + 0] ^ iv[0]; + m[1] = wpa_eapol->eapol[eapol_idx + 1] ^ iv[1]; + m[2] = wpa_eapol->eapol[eapol_idx + 2] ^ iv[2]; + m[3] = wpa_eapol->eapol[eapol_idx + 3] ^ iv[3]; aes128_encrypt (ks, m, iv, s_te0, s_te1, s_te2, s_te3, s_te4); } - m[0] = wpa->eapol[eapol_idx + 0]; - m[1] = wpa->eapol[eapol_idx + 1]; - m[2] = wpa->eapol[eapol_idx + 2]; - m[3] = wpa->eapol[eapol_idx + 3]; + m[0] = wpa_eapol->eapol[eapol_idx + 0]; + m[1] = wpa_eapol->eapol[eapol_idx + 1]; + m[2] = wpa_eapol->eapol[eapol_idx + 2]; + m[3] = wpa_eapol->eapol[eapol_idx + 3]; u32 k[4]; @@ -1309,10 +1309,10 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -1322,7 +1322,7 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul } } - if (wpa->detected_be == 1) + if (wpa_eapol->detected_be == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -1335,7 +1335,7 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul t = swap32_S (t); - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -1449,20 +1449,20 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul int eapol_left; int eapol_idx; - for (eapol_left = wpa->eapol_len, eapol_idx = 0; eapol_left > 16; eapol_left -= 16, eapol_idx += 4) + for (eapol_left = wpa_eapol->eapol_len, eapol_idx = 0; eapol_left > 16; eapol_left -= 16, eapol_idx += 4) { - m[0] = wpa->eapol[eapol_idx + 0] ^ iv[0]; - m[1] = wpa->eapol[eapol_idx + 1] ^ iv[1]; - m[2] = wpa->eapol[eapol_idx + 2] ^ iv[2]; - m[3] = wpa->eapol[eapol_idx + 3] ^ iv[3]; + m[0] = wpa_eapol->eapol[eapol_idx + 0] ^ iv[0]; + m[1] = wpa_eapol->eapol[eapol_idx + 1] ^ iv[1]; + m[2] = wpa_eapol->eapol[eapol_idx + 2] ^ iv[2]; + m[3] = wpa_eapol->eapol[eapol_idx + 3] ^ iv[3]; aes128_encrypt (ks, m, iv, s_te0, s_te1, s_te2, s_te3, s_te4); } - m[0] = wpa->eapol[eapol_idx + 0]; - m[1] = wpa->eapol[eapol_idx + 1]; - m[2] = wpa->eapol[eapol_idx + 2]; - m[3] = wpa->eapol[eapol_idx + 3]; + m[0] = wpa_eapol->eapol[eapol_idx + 0]; + m[1] = wpa_eapol->eapol[eapol_idx + 1]; + m[2] = wpa_eapol->eapol[eapol_idx + 2]; + m[3] = wpa_eapol->eapol[eapol_idx + 3]; u32 k[4]; @@ -1496,10 +1496,10 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { diff --git a/OpenCL/m02501-pure.cl b/OpenCL/m02501-pure.cl index 01d6da884..cc98dc5d4 100644 --- a/OpenCL/m02501-pure.cl +++ b/OpenCL/m02501-pure.cl @@ -64,7 +64,7 @@ DECLSPEC void make_kn (u32 *k) k[3] ^= c * 0x87000000; } -__kernel void m02501_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *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 wpa_t *wpa_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 m02501_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); @@ -110,19 +110,19 @@ __kernel void m02501_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].out[7] = swap32_S (out[7]); } -__kernel void m02501_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *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 wpa_t *wpa_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 m02501_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); if (gid >= gid_max) return; } -__kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *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 wpa_t *wpa_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 m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { // not in use here, special case... } -__kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *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 wpa_t *wpa_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 m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); @@ -144,46 +144,46 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 digest_cur = digests_offset + digest_pos; - __global const wpa_t *wpa = &wpa_bufs[digest_cur]; + __global const wpa_eapol_t *wpa_eapol = &wpa_eapol_bufs[digest_cur]; u32 pke[32]; - pke[ 0] = wpa->pke[ 0]; - pke[ 1] = wpa->pke[ 1]; - pke[ 2] = wpa->pke[ 2]; - pke[ 3] = wpa->pke[ 3]; - pke[ 4] = wpa->pke[ 4]; - pke[ 5] = wpa->pke[ 5]; - pke[ 6] = wpa->pke[ 6]; - pke[ 7] = wpa->pke[ 7]; - pke[ 8] = wpa->pke[ 8]; - pke[ 9] = wpa->pke[ 9]; - pke[10] = wpa->pke[10]; - pke[11] = wpa->pke[11]; - pke[12] = wpa->pke[12]; - pke[13] = wpa->pke[13]; - pke[14] = wpa->pke[14]; - pke[15] = wpa->pke[15]; - pke[16] = wpa->pke[16]; - pke[17] = wpa->pke[17]; - pke[18] = wpa->pke[18]; - pke[19] = wpa->pke[19]; - pke[20] = wpa->pke[20]; - pke[21] = wpa->pke[21]; - pke[22] = wpa->pke[22]; - pke[23] = wpa->pke[23]; - pke[24] = wpa->pke[24]; - pke[25] = wpa->pke[25]; - pke[26] = wpa->pke[26]; - pke[27] = wpa->pke[27]; - pke[28] = wpa->pke[28]; - pke[29] = wpa->pke[29]; - pke[30] = wpa->pke[30]; - pke[31] = wpa->pke[31]; + pke[ 0] = wpa_eapol->pke[ 0]; + pke[ 1] = wpa_eapol->pke[ 1]; + pke[ 2] = wpa_eapol->pke[ 2]; + pke[ 3] = wpa_eapol->pke[ 3]; + pke[ 4] = wpa_eapol->pke[ 4]; + pke[ 5] = wpa_eapol->pke[ 5]; + pke[ 6] = wpa_eapol->pke[ 6]; + pke[ 7] = wpa_eapol->pke[ 7]; + pke[ 8] = wpa_eapol->pke[ 8]; + pke[ 9] = wpa_eapol->pke[ 9]; + pke[10] = wpa_eapol->pke[10]; + pke[11] = wpa_eapol->pke[11]; + pke[12] = wpa_eapol->pke[12]; + pke[13] = wpa_eapol->pke[13]; + pke[14] = wpa_eapol->pke[14]; + pke[15] = wpa_eapol->pke[15]; + pke[16] = wpa_eapol->pke[16]; + pke[17] = wpa_eapol->pke[17]; + pke[18] = wpa_eapol->pke[18]; + pke[19] = wpa_eapol->pke[19]; + pke[20] = wpa_eapol->pke[20]; + pke[21] = wpa_eapol->pke[21]; + pke[22] = wpa_eapol->pke[22]; + pke[23] = wpa_eapol->pke[23]; + pke[24] = wpa_eapol->pke[24]; + pke[25] = wpa_eapol->pke[25]; + pke[26] = wpa_eapol->pke[26]; + pke[27] = wpa_eapol->pke[27]; + pke[28] = wpa_eapol->pke[28]; + pke[29] = wpa_eapol->pke[29]; + pke[30] = wpa_eapol->pke[30]; + pke[31] = wpa_eapol->pke[31]; u32 to; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { to = pke[15] << 24 | pke[16] >> 8; @@ -194,9 +194,9 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul | pke[24] >> 8; } - const u32 nonce_error_corrections = wpa->nonce_error_corrections; + const u32 nonce_error_corrections = wpa_eapol->nonce_error_corrections; - if (wpa->detected_le == 1) + if (wpa_eapol->detected_le == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -205,7 +205,7 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul t -= nonce_error_corrections / 2; t += nonce_error_correction; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -286,7 +286,7 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); - md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + md5_hmac_update_global (&ctx2, wpa_eapol->eapol, wpa_eapol->eapol_len); md5_hmac_final (&ctx2); @@ -299,10 +299,10 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -312,7 +312,7 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul } } - if (wpa->detected_be == 1) + if (wpa_eapol->detected_be == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -325,7 +325,7 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul t = swap32_S (t); - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -406,7 +406,7 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); - md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + md5_hmac_update_global (&ctx2, wpa_eapol->eapol, wpa_eapol->eapol_len); md5_hmac_final (&ctx2); @@ -419,10 +419,10 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -433,7 +433,7 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *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 wpa_t *wpa_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 m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); @@ -455,46 +455,46 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 digest_cur = digests_offset + digest_pos; - __global const wpa_t *wpa = &wpa_bufs[digest_cur]; + __global const wpa_eapol_t *wpa_eapol = &wpa_eapol_bufs[digest_cur]; u32 pke[32]; - pke[ 0] = wpa->pke[ 0]; - pke[ 1] = wpa->pke[ 1]; - pke[ 2] = wpa->pke[ 2]; - pke[ 3] = wpa->pke[ 3]; - pke[ 4] = wpa->pke[ 4]; - pke[ 5] = wpa->pke[ 5]; - pke[ 6] = wpa->pke[ 6]; - pke[ 7] = wpa->pke[ 7]; - pke[ 8] = wpa->pke[ 8]; - pke[ 9] = wpa->pke[ 9]; - pke[10] = wpa->pke[10]; - pke[11] = wpa->pke[11]; - pke[12] = wpa->pke[12]; - pke[13] = wpa->pke[13]; - pke[14] = wpa->pke[14]; - pke[15] = wpa->pke[15]; - pke[16] = wpa->pke[16]; - pke[17] = wpa->pke[17]; - pke[18] = wpa->pke[18]; - pke[19] = wpa->pke[19]; - pke[20] = wpa->pke[20]; - pke[21] = wpa->pke[21]; - pke[22] = wpa->pke[22]; - pke[23] = wpa->pke[23]; - pke[24] = wpa->pke[24]; - pke[25] = wpa->pke[25]; - pke[26] = wpa->pke[26]; - pke[27] = wpa->pke[27]; - pke[28] = wpa->pke[28]; - pke[29] = wpa->pke[29]; - pke[30] = wpa->pke[30]; - pke[31] = wpa->pke[31]; + pke[ 0] = wpa_eapol->pke[ 0]; + pke[ 1] = wpa_eapol->pke[ 1]; + pke[ 2] = wpa_eapol->pke[ 2]; + pke[ 3] = wpa_eapol->pke[ 3]; + pke[ 4] = wpa_eapol->pke[ 4]; + pke[ 5] = wpa_eapol->pke[ 5]; + pke[ 6] = wpa_eapol->pke[ 6]; + pke[ 7] = wpa_eapol->pke[ 7]; + pke[ 8] = wpa_eapol->pke[ 8]; + pke[ 9] = wpa_eapol->pke[ 9]; + pke[10] = wpa_eapol->pke[10]; + pke[11] = wpa_eapol->pke[11]; + pke[12] = wpa_eapol->pke[12]; + pke[13] = wpa_eapol->pke[13]; + pke[14] = wpa_eapol->pke[14]; + pke[15] = wpa_eapol->pke[15]; + pke[16] = wpa_eapol->pke[16]; + pke[17] = wpa_eapol->pke[17]; + pke[18] = wpa_eapol->pke[18]; + pke[19] = wpa_eapol->pke[19]; + pke[20] = wpa_eapol->pke[20]; + pke[21] = wpa_eapol->pke[21]; + pke[22] = wpa_eapol->pke[22]; + pke[23] = wpa_eapol->pke[23]; + pke[24] = wpa_eapol->pke[24]; + pke[25] = wpa_eapol->pke[25]; + pke[26] = wpa_eapol->pke[26]; + pke[27] = wpa_eapol->pke[27]; + pke[28] = wpa_eapol->pke[28]; + pke[29] = wpa_eapol->pke[29]; + pke[30] = wpa_eapol->pke[30]; + pke[31] = wpa_eapol->pke[31]; u32 to; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { to = pke[15] << 24 | pke[16] >> 8; @@ -505,9 +505,9 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul | pke[24] >> 8; } - const u32 nonce_error_corrections = wpa->nonce_error_corrections; + const u32 nonce_error_corrections = wpa_eapol->nonce_error_corrections; - if (wpa->detected_le == 1) + if (wpa_eapol->detected_le == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -516,7 +516,7 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul t -= nonce_error_corrections / 2; t += nonce_error_correction; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -597,7 +597,7 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); - sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + sha1_hmac_update_global (&ctx2, wpa_eapol->eapol, wpa_eapol->eapol_len); sha1_hmac_final (&ctx2); @@ -610,10 +610,10 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -623,7 +623,7 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul } } - if (wpa->detected_be == 1) + if (wpa_eapol->detected_be == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -636,7 +636,7 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul t = swap32_S (t); - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -717,7 +717,7 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); - sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + sha1_hmac_update_global (&ctx2, wpa_eapol->eapol, wpa_eapol->eapol_len); sha1_hmac_final (&ctx2); @@ -730,10 +730,10 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -744,7 +744,7 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *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 wpa_t *wpa_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 m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_eapol_t *wpa_eapol_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) { const u64 gid = get_global_id (0); const u64 lid = get_local_id (0); @@ -818,46 +818,46 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 digest_cur = digests_offset + digest_pos; - __global const wpa_t *wpa = &wpa_bufs[digest_cur]; + __global const wpa_eapol_t *wpa_eapol = &wpa_eapol_bufs[digest_cur]; u32 pke[32]; - pke[ 0] = wpa->pke[ 0]; - pke[ 1] = wpa->pke[ 1]; - pke[ 2] = wpa->pke[ 2]; - pke[ 3] = wpa->pke[ 3]; - pke[ 4] = wpa->pke[ 4]; - pke[ 5] = wpa->pke[ 5]; - pke[ 6] = wpa->pke[ 6]; - pke[ 7] = wpa->pke[ 7]; - pke[ 8] = wpa->pke[ 8]; - pke[ 9] = wpa->pke[ 9]; - pke[10] = wpa->pke[10]; - pke[11] = wpa->pke[11]; - pke[12] = wpa->pke[12]; - pke[13] = wpa->pke[13]; - pke[14] = wpa->pke[14]; - pke[15] = wpa->pke[15]; - pke[16] = wpa->pke[16]; - pke[17] = wpa->pke[17]; - pke[18] = wpa->pke[18]; - pke[19] = wpa->pke[19]; - pke[20] = wpa->pke[20]; - pke[21] = wpa->pke[21]; - pke[22] = wpa->pke[22]; - pke[23] = wpa->pke[23]; - pke[24] = wpa->pke[24]; - pke[25] = wpa->pke[25]; - pke[26] = wpa->pke[26]; - pke[27] = wpa->pke[27]; - pke[28] = wpa->pke[28]; - pke[29] = wpa->pke[29]; - pke[30] = wpa->pke[30]; - pke[31] = wpa->pke[31]; + pke[ 0] = wpa_eapol->pke[ 0]; + pke[ 1] = wpa_eapol->pke[ 1]; + pke[ 2] = wpa_eapol->pke[ 2]; + pke[ 3] = wpa_eapol->pke[ 3]; + pke[ 4] = wpa_eapol->pke[ 4]; + pke[ 5] = wpa_eapol->pke[ 5]; + pke[ 6] = wpa_eapol->pke[ 6]; + pke[ 7] = wpa_eapol->pke[ 7]; + pke[ 8] = wpa_eapol->pke[ 8]; + pke[ 9] = wpa_eapol->pke[ 9]; + pke[10] = wpa_eapol->pke[10]; + pke[11] = wpa_eapol->pke[11]; + pke[12] = wpa_eapol->pke[12]; + pke[13] = wpa_eapol->pke[13]; + pke[14] = wpa_eapol->pke[14]; + pke[15] = wpa_eapol->pke[15]; + pke[16] = wpa_eapol->pke[16]; + pke[17] = wpa_eapol->pke[17]; + pke[18] = wpa_eapol->pke[18]; + pke[19] = wpa_eapol->pke[19]; + pke[20] = wpa_eapol->pke[20]; + pke[21] = wpa_eapol->pke[21]; + pke[22] = wpa_eapol->pke[22]; + pke[23] = wpa_eapol->pke[23]; + pke[24] = wpa_eapol->pke[24]; + pke[25] = wpa_eapol->pke[25]; + pke[26] = wpa_eapol->pke[26]; + pke[27] = wpa_eapol->pke[27]; + pke[28] = wpa_eapol->pke[28]; + pke[29] = wpa_eapol->pke[29]; + pke[30] = wpa_eapol->pke[30]; + pke[31] = wpa_eapol->pke[31]; u32 to; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { to = pke[15] << 24 | pke[16] >> 8; @@ -868,9 +868,9 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul | pke[24] >> 8; } - const u32 nonce_error_corrections = wpa->nonce_error_corrections; + const u32 nonce_error_corrections = wpa_eapol->nonce_error_corrections; - if (wpa->detected_le == 1) + if (wpa_eapol->detected_le == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -879,7 +879,7 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul t -= nonce_error_corrections / 2; t += nonce_error_correction; - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -957,20 +957,20 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul int eapol_left; int eapol_idx; - for (eapol_left = wpa->eapol_len, eapol_idx = 0; eapol_left > 16; eapol_left -= 16, eapol_idx += 4) + for (eapol_left = wpa_eapol->eapol_len, eapol_idx = 0; eapol_left > 16; eapol_left -= 16, eapol_idx += 4) { - m[0] = wpa->eapol[eapol_idx + 0] ^ iv[0]; - m[1] = wpa->eapol[eapol_idx + 1] ^ iv[1]; - m[2] = wpa->eapol[eapol_idx + 2] ^ iv[2]; - m[3] = wpa->eapol[eapol_idx + 3] ^ iv[3]; + m[0] = wpa_eapol->eapol[eapol_idx + 0] ^ iv[0]; + m[1] = wpa_eapol->eapol[eapol_idx + 1] ^ iv[1]; + m[2] = wpa_eapol->eapol[eapol_idx + 2] ^ iv[2]; + m[3] = wpa_eapol->eapol[eapol_idx + 3] ^ iv[3]; aes128_encrypt (ks, m, iv, s_te0, s_te1, s_te2, s_te3, s_te4); } - m[0] = wpa->eapol[eapol_idx + 0]; - m[1] = wpa->eapol[eapol_idx + 1]; - m[2] = wpa->eapol[eapol_idx + 2]; - m[3] = wpa->eapol[eapol_idx + 3]; + m[0] = wpa_eapol->eapol[eapol_idx + 0]; + m[1] = wpa_eapol->eapol[eapol_idx + 1]; + m[2] = wpa_eapol->eapol[eapol_idx + 2]; + m[3] = wpa_eapol->eapol[eapol_idx + 3]; u32 k[4]; @@ -1004,10 +1004,10 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { @@ -1017,7 +1017,7 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul } } - if (wpa->detected_be == 1) + if (wpa_eapol->detected_be == 1) { for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { @@ -1030,7 +1030,7 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul t = swap32_S (t); - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); pke[16] = (pke[16] & ~0xffffff00) | (t << 8); @@ -1108,20 +1108,20 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul int eapol_left; int eapol_idx; - for (eapol_left = wpa->eapol_len, eapol_idx = 0; eapol_left > 16; eapol_left -= 16, eapol_idx += 4) + for (eapol_left = wpa_eapol->eapol_len, eapol_idx = 0; eapol_left > 16; eapol_left -= 16, eapol_idx += 4) { - m[0] = wpa->eapol[eapol_idx + 0] ^ iv[0]; - m[1] = wpa->eapol[eapol_idx + 1] ^ iv[1]; - m[2] = wpa->eapol[eapol_idx + 2] ^ iv[2]; - m[3] = wpa->eapol[eapol_idx + 3] ^ iv[3]; + m[0] = wpa_eapol->eapol[eapol_idx + 0] ^ iv[0]; + m[1] = wpa_eapol->eapol[eapol_idx + 1] ^ iv[1]; + m[2] = wpa_eapol->eapol[eapol_idx + 2] ^ iv[2]; + m[3] = wpa_eapol->eapol[eapol_idx + 3] ^ iv[3]; aes128_encrypt (ks, m, iv, s_te0, s_te1, s_te2, s_te3, s_te4); } - m[0] = wpa->eapol[eapol_idx + 0]; - m[1] = wpa->eapol[eapol_idx + 1]; - m[2] = wpa->eapol[eapol_idx + 2]; - m[3] = wpa->eapol[eapol_idx + 3]; + m[0] = wpa_eapol->eapol[eapol_idx + 0]; + m[1] = wpa_eapol->eapol[eapol_idx + 1]; + m[2] = wpa_eapol->eapol[eapol_idx + 2]; + m[3] = wpa_eapol->eapol[eapol_idx + 3]; u32 k[4]; @@ -1155,10 +1155,10 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul * final compare */ - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) + if ((keymic[0] == wpa_eapol->keymic[0]) + && (keymic[1] == wpa_eapol->keymic[1]) + && (keymic[2] == wpa_eapol->keymic[2]) + && (keymic[3] == wpa_eapol->keymic[3])) { if (atomic_inc (&hashes_shown[digest_cur]) == 0) { diff --git a/OpenCL/m06400-pure.cl b/OpenCL/m06400-pure.cl index d1fb9d91b..4939c77f9 100644 --- a/OpenCL/m06400-pure.cl +++ b/OpenCL/m06400-pure.cl @@ -58,7 +58,7 @@ DECLSPEC void hmac_sha256_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *i sha256_transform_vector (w0, w1, w2, w3, digest); } -__kernel void m06400_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global sha256aix_tmp_t *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 wpa_t *wpa_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 m06400_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global sha256aix_tmp_t *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) { /** * base @@ -142,7 +142,7 @@ __kernel void m06400_init (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m06400_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global sha256aix_tmp_t *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 wpa_t *wpa_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 m06400_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global sha256aix_tmp_t *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) { const u64 gid = get_global_id (0); @@ -248,7 +248,7 @@ __kernel void m06400_loop (__global pw_t *pws, __global const kernel_rule_t *rul } } -__kernel void m06400_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global sha256aix_tmp_t *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 wpa_t *wpa_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 m06400_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global sha256aix_tmp_t *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) { /** * base diff --git a/OpenCL/m16800-pure.cl b/OpenCL/m16800-pure.cl new file mode 100644 index 000000000..9d13e669a --- /dev/null +++ b/OpenCL/m16800-pure.cl @@ -0,0 +1,254 @@ +/** + * 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" + +#define COMPARE_S "inc_comp_single.cl" +#define COMPARE_M "inc_comp_multi.cl" + +DECLSPEC void hmac_sha1_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *ipad, u32x *opad, u32x *digest) +{ + digest[0] = ipad[0]; + digest[1] = ipad[1]; + digest[2] = ipad[2]; + digest[3] = ipad[3]; + digest[4] = ipad[4]; + + sha1_transform_vector (w0, w1, w2, w3, digest); + + w0[0] = digest[0]; + w0[1] = digest[1]; + w0[2] = digest[2]; + w0[3] = digest[3]; + w1[0] = digest[4]; + 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; + w3[2] = 0; + w3[3] = (64 + 20) * 8; + + digest[0] = opad[0]; + digest[1] = opad[1]; + digest[2] = opad[2]; + digest[3] = opad[3]; + digest[4] = opad[4]; + + sha1_transform_vector (w0, w1, w2, w3, digest); +} + +__kernel void m16800_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_pmkid_t *wpa_pmkid_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) +{ + /** + * base + */ + + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + sha1_hmac_ctx_t sha1_hmac_ctx; + + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + + tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; + tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; + tmps[gid].ipad[2] = sha1_hmac_ctx.ipad.h[2]; + tmps[gid].ipad[3] = sha1_hmac_ctx.ipad.h[3]; + tmps[gid].ipad[4] = sha1_hmac_ctx.ipad.h[4]; + + tmps[gid].opad[0] = sha1_hmac_ctx.opad.h[0]; + tmps[gid].opad[1] = sha1_hmac_ctx.opad.h[1]; + tmps[gid].opad[2] = sha1_hmac_ctx.opad.h[2]; + tmps[gid].opad[3] = sha1_hmac_ctx.opad.h[3]; + tmps[gid].opad[4] = sha1_hmac_ctx.opad.h[4]; + + sha1_hmac_update_global_swap (&sha1_hmac_ctx, wpa_pmkid_bufs[digests_offset].essid_buf, wpa_pmkid_bufs[digests_offset].essid_len); + + for (u32 i = 0, j = 1; i < 8; i += 5, j += 1) + { + sha1_hmac_ctx_t sha1_hmac_ctx2 = sha1_hmac_ctx; + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = j; + w0[1] = 0; + w0[2] = 0; + w0[3] = 0; + w1[0] = 0; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + sha1_hmac_update_64 (&sha1_hmac_ctx2, w0, w1, w2, w3, 4); + + sha1_hmac_final (&sha1_hmac_ctx2); + + tmps[gid].dgst[i + 0] = sha1_hmac_ctx2.opad.h[0]; + tmps[gid].dgst[i + 1] = sha1_hmac_ctx2.opad.h[1]; + tmps[gid].dgst[i + 2] = sha1_hmac_ctx2.opad.h[2]; + tmps[gid].dgst[i + 3] = sha1_hmac_ctx2.opad.h[3]; + tmps[gid].dgst[i + 4] = sha1_hmac_ctx2.opad.h[4]; + + tmps[gid].out[i + 0] = tmps[gid].dgst[i + 0]; + tmps[gid].out[i + 1] = tmps[gid].dgst[i + 1]; + tmps[gid].out[i + 2] = tmps[gid].dgst[i + 2]; + tmps[gid].out[i + 3] = tmps[gid].dgst[i + 3]; + tmps[gid].out[i + 4] = tmps[gid].dgst[i + 4]; + } +} + +__kernel void m16800_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_pmkid_t *wpa_pmkid_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) +{ + const u64 gid = get_global_id (0); + + if ((gid * VECT_SIZE) >= gid_max) return; + + u32x ipad[5]; + u32x opad[5]; + + ipad[0] = packv (tmps, ipad, gid, 0); + ipad[1] = packv (tmps, ipad, gid, 1); + ipad[2] = packv (tmps, ipad, gid, 2); + ipad[3] = packv (tmps, ipad, gid, 3); + ipad[4] = packv (tmps, ipad, gid, 4); + + opad[0] = packv (tmps, opad, gid, 0); + opad[1] = packv (tmps, opad, gid, 1); + opad[2] = packv (tmps, opad, gid, 2); + opad[3] = packv (tmps, opad, gid, 3); + opad[4] = packv (tmps, opad, gid, 4); + + for (u32 i = 0; i < 8; i += 5) + { + u32x dgst[5]; + u32x out[5]; + + dgst[0] = packv (tmps, dgst, gid, i + 0); + dgst[1] = packv (tmps, dgst, gid, i + 1); + dgst[2] = packv (tmps, dgst, gid, i + 2); + dgst[3] = packv (tmps, dgst, gid, i + 3); + dgst[4] = packv (tmps, dgst, gid, i + 4); + + out[0] = packv (tmps, out, gid, i + 0); + out[1] = packv (tmps, out, gid, i + 1); + out[2] = packv (tmps, out, gid, i + 2); + out[3] = packv (tmps, out, gid, i + 3); + out[4] = packv (tmps, out, gid, i + 4); + + for (u32 j = 0; j < loop_cnt; j++) + { + u32x w0[4]; + u32x w1[4]; + u32x w2[4]; + u32x w3[4]; + + w0[0] = dgst[0]; + w0[1] = dgst[1]; + w0[2] = dgst[2]; + w0[3] = dgst[3]; + w1[0] = dgst[4]; + 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; + w3[2] = 0; + w3[3] = (64 + 20) * 8; + + hmac_sha1_run_V (w0, w1, w2, w3, ipad, opad, dgst); + + out[0] ^= dgst[0]; + out[1] ^= dgst[1]; + out[2] ^= dgst[2]; + out[3] ^= dgst[3]; + out[4] ^= dgst[4]; + } + + unpackv (tmps, dgst, gid, i + 0, dgst[0]); + unpackv (tmps, dgst, gid, i + 1, dgst[1]); + unpackv (tmps, dgst, gid, i + 2, dgst[2]); + unpackv (tmps, dgst, gid, i + 3, dgst[3]); + unpackv (tmps, dgst, gid, i + 4, dgst[4]); + + unpackv (tmps, out, gid, i + 0, out[0]); + unpackv (tmps, out, gid, i + 1, out[1]); + unpackv (tmps, out, gid, i + 2, out[2]); + unpackv (tmps, out, gid, i + 3, out[3]); + unpackv (tmps, out, gid, i + 4, out[4]); + } +} + +__kernel void m16800_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pbkdf2_tmp_t *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 wpa_pmkid_t *wpa_pmkid_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) +{ + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + + if (gid >= gid_max) return; + + u32 w[16]; + + w[ 0] = tmps[gid].out[0]; + w[ 1] = tmps[gid].out[1]; + w[ 2] = tmps[gid].out[2]; + w[ 3] = tmps[gid].out[3]; + w[ 4] = tmps[gid].out[4]; + w[ 5] = tmps[gid].out[5]; + w[ 6] = tmps[gid].out[6]; + w[ 7] = tmps[gid].out[7]; + w[ 8] = 0; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 0; + w[15] = 0; + + sha1_hmac_ctx_t sha1_hmac_ctx; + + sha1_hmac_init (&sha1_hmac_ctx, w, 32); + + sha1_hmac_update_global_swap (&sha1_hmac_ctx, wpa_pmkid_bufs[digests_offset].pmkid_data, 20); + + sha1_hmac_final (&sha1_hmac_ctx); + + const u32 r0 = sha1_hmac_ctx.opad.h[0]; + const u32 r1 = sha1_hmac_ctx.opad.h[1]; + const u32 r2 = sha1_hmac_ctx.opad.h[2]; + const u32 r3 = sha1_hmac_ctx.opad.h[3]; + + #define il_pos 0 + + #include COMPARE_M +} diff --git a/OpenCL/m16801-pure.cl b/OpenCL/m16801-pure.cl new file mode 100644 index 000000000..1ad301f97 --- /dev/null +++ b/OpenCL/m16801-pure.cl @@ -0,0 +1,129 @@ +/** + * 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" + +#define COMPARE_S "inc_comp_single.cl" +#define COMPARE_M "inc_comp_multi.cl" + +DECLSPEC u8 hex_convert (const u8 c) +{ + return (c & 15) + (c >> 6) * 9; +} + +DECLSPEC u8 hex_to_u8 (const u8 *hex) +{ + u8 v = 0; + + v |= ((u8) hex_convert (hex[1]) << 0); + v |= ((u8) hex_convert (hex[0]) << 4); + + return (v); +} + +__kernel void m16801_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_pmkid_t *wpa_pmkid_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) +{ + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; + + u32 in[16]; + + in[ 0] = pws[gid].i[ 0]; + in[ 1] = pws[gid].i[ 1]; + in[ 2] = pws[gid].i[ 2]; + in[ 3] = pws[gid].i[ 3]; + in[ 4] = pws[gid].i[ 4]; + in[ 5] = pws[gid].i[ 5]; + in[ 6] = pws[gid].i[ 6]; + in[ 7] = pws[gid].i[ 7]; + in[ 8] = pws[gid].i[ 8]; + in[ 9] = pws[gid].i[ 9]; + in[10] = pws[gid].i[10]; + in[11] = pws[gid].i[11]; + in[12] = pws[gid].i[12]; + in[13] = pws[gid].i[13]; + in[14] = pws[gid].i[14]; + in[15] = pws[gid].i[15]; + + u8 *in_ptr = (u8 *) in; + + u32 out[8]; + + u8 *out_ptr = (u8 *) out; + + for (int i = 0, j = 0; i < 32; i += 1, j += 2) + { + out_ptr[i] = hex_to_u8 (in_ptr + j); + } + + tmps[gid].out[0] = swap32_S (out[0]); + tmps[gid].out[1] = swap32_S (out[1]); + tmps[gid].out[2] = swap32_S (out[2]); + tmps[gid].out[3] = swap32_S (out[3]); + tmps[gid].out[4] = swap32_S (out[4]); + tmps[gid].out[5] = swap32_S (out[5]); + tmps[gid].out[6] = swap32_S (out[6]); + tmps[gid].out[7] = swap32_S (out[7]); +} + +__kernel void m16801_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_pmkid_t *wpa_pmkid_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) +{ + const u64 gid = get_global_id (0); + + if (gid >= gid_max) return; +} + +__kernel void m16801_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpa_pmk_tmp_t *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 wpa_pmkid_t *wpa_pmkid_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) +{ + const u64 gid = get_global_id (0); + const u64 lid = get_local_id (0); + + if (gid >= gid_max) return; + + u32 w[16]; + + w[ 0] = tmps[gid].out[0]; + w[ 1] = tmps[gid].out[1]; + w[ 2] = tmps[gid].out[2]; + w[ 3] = tmps[gid].out[3]; + w[ 4] = tmps[gid].out[4]; + w[ 5] = tmps[gid].out[5]; + w[ 6] = tmps[gid].out[6]; + w[ 7] = tmps[gid].out[7]; + w[ 8] = 0; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 0; + w[15] = 0; + + sha1_hmac_ctx_t sha1_hmac_ctx; + + sha1_hmac_init (&sha1_hmac_ctx, w, 32); + + sha1_hmac_update_global_swap (&sha1_hmac_ctx, wpa_pmkid_bufs[digests_offset].pmkid_data, 20); + + sha1_hmac_final (&sha1_hmac_ctx); + + const u32 r0 = sha1_hmac_ctx.opad.h[0]; + const u32 r1 = sha1_hmac_ctx.opad.h[1]; + const u32 r2 = sha1_hmac_ctx.opad.h[2]; + const u32 r3 = sha1_hmac_ctx.opad.h[3]; + + #define il_pos 0 + + #include COMPARE_M +} diff --git a/docs/changes.txt b/docs/changes.txt index 57e73a85d..1f5c6c136 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -5,6 +5,8 @@ ## - Added hash-mode 16700 = FileVault 2 +- Added hash-mode 16800 = WPA-PMKID-PBKDF2 +- Added hash-mode 16801 = WPA-PMKID-PMK ## ## Improvements diff --git a/docs/readme.txt b/docs/readme.txt index dd78b0039..7f1694d52 100644 --- a/docs/readme.txt +++ b/docs/readme.txt @@ -101,8 +101,10 @@ NVIDIA GPUs require "NVIDIA Driver" (367.x or later) - PBKDF2-HMAC-SHA256 - PBKDF2-HMAC-SHA512 - Skype -- WPA/WPA2 -- WPA/WPA2 PMK +- WPA-EAPOL-PBKDF2 +- WPA-EAPOL-PMK +- WPA-PMKID-PBKDF2 +- WPA-PMKID-PMK - iSCSI CHAP authentication, MD5(CHAP) - IKE-PSK MD5 - IKE-PSK SHA1 diff --git a/extra/tab_completion/hashcat.sh b/extra/tab_completion/hashcat.sh index f606ac9f7..006a7141f 100644 --- a/extra/tab_completion/hashcat.sh +++ b/extra/tab_completion/hashcat.sh @@ -176,7 +176,7 @@ _hashcat () { local VERSION=4.2.0 - 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 5000 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" + 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 5000 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" local ATTACK_MODES="0 1 3 6 7" local HCCAPX_MESSAGE_PAIR="0 1 2 3 4 5" local OUTFILE_FORMATS="1 2 3 4 5 6 7 8 9 10 11 12 13 14 15" diff --git a/include/interface.h b/include/interface.h index 646c0d9e0..521768cc7 100644 --- a/include/interface.h +++ b/include/interface.h @@ -176,7 +176,7 @@ typedef struct pdf } pdf_t; -typedef struct wpa +typedef struct wpa_eapol { u32 pke[32]; u32 eapol[64 + 16]; @@ -197,7 +197,18 @@ typedef struct wpa int detected_le; int detected_be; -} wpa_t; +} wpa_eapol_t; + +typedef struct wpa_pmkid +{ + u32 pmkid[4]; + u32 pmkid_data[16]; + u8 orig_mac_ap[6]; + u8 orig_mac_sta[6]; + u8 essid_len; + u32 essid_buf[16]; + +} wpa_pmkid_t; typedef struct bitcoin_wallet { @@ -586,7 +597,7 @@ typedef struct sha512crypt_tmp } sha512crypt_tmp_t; -typedef struct wpa_tmp +typedef struct wpa_pbkdf2_tmp { u32 ipad[5]; u32 opad[5]; @@ -594,13 +605,13 @@ typedef struct wpa_tmp u32 dgst[10]; u32 out[10]; -} wpa_tmp_t; +} wpa_pbkdf2_tmp_t; -typedef struct wpapmk_tmp +typedef struct wpa_pmk_tmp { u32 out[8]; -} wpapmk_tmp_t; +} wpa_pmk_tmp_t; typedef struct bitcoin_wallet_tmp { @@ -1029,7 +1040,7 @@ typedef enum hash_type HASH_TYPE_SHA384 = 7, HASH_TYPE_SHA512 = 8, HASH_TYPE_DCC2 = 9, - HASH_TYPE_WPA = 10, + HASH_TYPE_WPA_EAPOL = 10, HASH_TYPE_LM = 11, HASH_TYPE_DESCRYPT = 12, HASH_TYPE_ORACLEH = 13, @@ -1087,6 +1098,8 @@ typedef enum hash_type HASH_TYPE_CRAM_MD5_DOVECOT = 65, HASH_TYPE_JWT = 66, HASH_TYPE_ELECTRUM_WALLET = 67, + HASH_TYPE_WPA_PMKID_PBKDF2 = 68, + HASH_TYPE_WPA_PMKID_PMK = 69, } hash_type_t; @@ -1136,8 +1149,8 @@ typedef enum kern_type KERN_TYPE_DCC2 = 2100, KERN_TYPE_MD5PIX = 2400, KERN_TYPE_MD5ASA = 2410, - KERN_TYPE_WPA = 2500, - KERN_TYPE_WPAPMK = 2501, + KERN_TYPE_WPA_EAPOL_PBKDF2 = 2500, + KERN_TYPE_WPA_EAPOL_PMK = 2501, KERN_TYPE_MD55 = 2600, KERN_TYPE_MD55_PWSLT1 = 2610, KERN_TYPE_MD55_PWSLT2 = 2710, @@ -1294,6 +1307,8 @@ typedef enum kern_type KERN_TYPE_JWT_HS384 = 16512, KERN_TYPE_JWT_HS512 = 16513, KERN_TYPE_ELECTRUM_WALLET13 = 16600, + KERN_TYPE_WPA_PMKID_PBKDF2 = 16800, + KERN_TYPE_WPA_PMKID_PMK = 16801, KERN_TYPE_PLAINTEXT = 99999, } kern_type_t; @@ -1306,8 +1321,8 @@ typedef enum rounds_count { ROUNDS_PHPASS = (1 << 11), // $P$B ROUNDS_DCC2 = 10240, - ROUNDS_WPA = 4096, - ROUNDS_WPAPMK = 1, + ROUNDS_WPA_PBKDF2 = 4096, + ROUNDS_WPA_PMK = 1, ROUNDS_BCRYPT = (1 << 5), ROUNDS_PSAFE3 = 2048, ROUNDS_ANDROIDPIN = 1024, @@ -1424,7 +1439,7 @@ int sha512_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu int sha512s_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int sha512crypt_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int vb30_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); -int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); +int wpa_eapol_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int psafe2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int psafe3_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int ikepsk_md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); @@ -1552,6 +1567,8 @@ int ethereum_presale_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu int jwt_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int electrum_wallet13_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); int filevault2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig); +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); /** * hook functions diff --git a/src/hashes.c b/src/hashes.c index 0e59e6c5f..cd026f149 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -823,27 +823,27 @@ int hashes_init_stage1 (hashcat_ctx_t *hashcat_ctx) if ((user_options->hash_mode == 2500) || (user_options->hash_mode == 2501)) { - wpa_t *wpa = (wpa_t *) hashes_buf[hashes_cnt].esalt; + wpa_eapol_t *wpa_eapol = (wpa_eapol_t *) hashes_buf[hashes_cnt].esalt; if (user_options->hccapx_message_pair_chgd == true) { - wpa->message_pair_chgd = (int) user_options->hccapx_message_pair_chgd; - wpa->message_pair = (u8) user_options->hccapx_message_pair; + wpa_eapol->message_pair_chgd = (int) user_options->hccapx_message_pair_chgd; + wpa_eapol->message_pair = (u8) user_options->hccapx_message_pair; } - if (wpa->message_pair & (1 << 4)) + if (wpa_eapol->message_pair & (1 << 4)) { // ap-less attack detected, nc not needed - wpa->nonce_error_corrections = 0; + wpa_eapol->nonce_error_corrections = 0; } else { - if (wpa->message_pair & (1 << 7)) + if (wpa_eapol->message_pair & (1 << 7)) { // replaycount not checked, nc needed - wpa->nonce_error_corrections = user_options->nonce_error_corrections; + wpa_eapol->nonce_error_corrections = user_options->nonce_error_corrections; } else { @@ -851,11 +851,11 @@ int hashes_init_stage1 (hashcat_ctx_t *hashcat_ctx) if (user_options->nonce_error_corrections_chgd == true) { - wpa->nonce_error_corrections = user_options->nonce_error_corrections; + wpa_eapol->nonce_error_corrections = user_options->nonce_error_corrections; } else { - wpa->nonce_error_corrections = 0; + wpa_eapol->nonce_error_corrections = 0; } } } @@ -866,18 +866,18 @@ int hashes_init_stage1 (hashcat_ctx_t *hashcat_ctx) // this means that we check both even if both are not set! // however if one of them is set, we can assume that the endianess has been checked and the other one is not needed - wpa->detected_le = 1; - wpa->detected_be = 1; + wpa_eapol->detected_le = 1; + wpa_eapol->detected_be = 1; - if (wpa->message_pair & (1 << 5)) + if (wpa_eapol->message_pair & (1 << 5)) { - wpa->detected_le = 1; - wpa->detected_be = 0; + wpa_eapol->detected_le = 1; + wpa_eapol->detected_be = 0; } - else if (wpa->message_pair & (1 << 6)) + else if (wpa_eapol->message_pair & (1 << 6)) { - wpa->detected_le = 0; - wpa->detected_be = 1; + wpa_eapol->detected_le = 0; + wpa_eapol->detected_be = 1; } } } @@ -1697,12 +1697,12 @@ int hashes_init_selftest (hashcat_ctx_t *hashcat_ctx) hcfree (tmpdata); - wpa_t *wpa = (wpa_t *) st_esalts_buf; + wpa_eapol_t *wpa_eapol = (wpa_eapol_t *) st_esalts_buf; - wpa->detected_le = 1; - wpa->detected_be = 0; + wpa_eapol->detected_le = 1; + wpa_eapol->detected_be = 0; - wpa->nonce_error_corrections = 3; + wpa_eapol->nonce_error_corrections = 3; } else if (hashconfig->opts_type & OPTS_TYPE_BINARY_HASHFILE) { diff --git a/src/interface.c b/src/interface.c index 4f944deb7..9f542ffd5 100644 --- a/src/interface.c +++ b/src/interface.c @@ -32,6 +32,7 @@ static const char *ST_PASS_HEX_02501 = "7f620a599c445155935a35634638fa67b4aa static const char *ST_PASS_BIN_09710 = "\x91\xb2\xe0\x62\xb9"; static const char *ST_PASS_BIN_09810 = "\xb8\xf6\x36\x19\xca"; static const char *ST_PASS_BIN_10410 = "\x6a\x8a\xed\xcc\xb7"; +static const char *ST_PASS_HEX_16801 = "5b13d4babb3714ccc62c9f71864bc984efd6a55f237c7a87fc2151e1ca658a9d"; /** * Missing self-test hashes: @@ -277,6 +278,8 @@ static const char *ST_HASH_16300 = "$ethereum$w*e94a8e49deac2d62206bf9bfb7d2aaea static const char *ST_HASH_16400 = "{CRAM-MD5}5389b33b9725e5657cb631dc50017ff100000000000000000000000000000000"; static const char *ST_HASH_16600 = "$electrum$1*44358283104603165383613672586868*c43a6632d9f59364f74c395a03d8c2ea"; static const char *ST_HASH_16700 = "$fvde$1$16$84286044060108438487434858307513$20000$f1620ab93192112f0a23eea89b5d4df065661f974b704191"; +static const char *ST_HASH_16800 = "2582a8281bf9d4308d6f5731d0e61c61*4604ba734d4e*89acf0e761f4*ed487162465a774bfba60eb603a39f3a"; +static const char *ST_HASH_16801 = "2582a8281bf9d4308d6f5731d0e61c61*4604ba734d4e*89acf0e761f4"; static const char *ST_HASH_99999 = "hashcat"; static const char *OPTI_STR_OPTIMIZED_KERNEL = "Optimized-Kernel"; @@ -384,8 +387,8 @@ static const char *HT_01800 = "sha512crypt $6$, SHA512 (Unix)"; static const char *HT_02100 = "Domain Cached Credentials 2 (DCC2), MS Cache 2"; static const char *HT_02400 = "Cisco-PIX MD5"; static const char *HT_02410 = "Cisco-ASA MD5"; -static const char *HT_02500 = "WPA/WPA2"; -static const char *HT_02501 = "WPA/WPA2 PMK"; +static const char *HT_02500 = "WPA-EAPOL-PBKDF2"; +static const char *HT_02501 = "WPA-EAPOL-PMK"; static const char *HT_02600 = "md5(md5($pass))"; static const char *HT_03000 = "LM"; static const char *HT_03100 = "Oracle H: Type (Oracle 7+)"; @@ -521,6 +524,8 @@ static const char *HT_16400 = "CRAM-MD5 Dovecot"; static const char *HT_16500 = "JWT (JSON Web Token)"; static const char *HT_16600 = "Electrum Wallet (Salt-Type 1-3)"; 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_99999 = "Plaintext"; static const char *HT_00011 = "Joomla < 2.5.18"; @@ -3475,13 +3480,13 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN return (PARSER_OK); } -int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig) +int wpa_eapol_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig) { u32 *digest = (u32 *) hash_buf->digest; salt_t *salt = hash_buf->salt; - wpa_t *wpa = (wpa_t *) hash_buf->esalt; + wpa_eapol_t *wpa_eapol = (wpa_eapol_t *) hash_buf->esalt; // the *wpa was partially initialized beforehand, we can not simply memset it to zero @@ -3495,7 +3500,7 @@ int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED if (in.eapol_len < 1 || in.eapol_len > 255) return (PARSER_HCCAPX_EAPOL_LEN); - memcpy (wpa->keymic, in.keymic, 16); + memcpy (wpa_eapol->keymic, in.keymic, 16); /* http://www.one-net.eu/jsw/j_sec/m_ptype.html @@ -3514,21 +3519,21 @@ int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED salt->salt_len = salt_len; - salt->salt_iter = ROUNDS_WPA - 1; + salt->salt_iter = ROUNDS_WPA_PBKDF2 - 1; - memcpy (wpa->essid, in.essid, in.essid_len); + memcpy (wpa_eapol->essid, in.essid, in.essid_len); - wpa->essid_len = in.essid_len; + wpa_eapol->essid_len = in.essid_len; - wpa->keyver = in.keyver; + wpa_eapol->keyver = in.keyver; - if ((wpa->keyver != 1) && (wpa->keyver != 2) && (wpa->keyver != 3)) return (PARSER_SALT_VALUE); + if ((wpa_eapol->keyver != 1) && (wpa_eapol->keyver != 2) && (wpa_eapol->keyver != 3)) return (PARSER_SALT_VALUE); - u8 *pke_ptr = (u8 *) wpa->pke; + u8 *pke_ptr = (u8 *) wpa_eapol->pke; memset (pke_ptr, 0, 128); - if ((wpa->keyver == 1) || (wpa->keyver == 2)) + if ((wpa_eapol->keyver == 1) || (wpa_eapol->keyver == 2)) { memcpy (pke_ptr, "Pairwise key expansion", 23); @@ -3543,9 +3548,9 @@ int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED memcpy (pke_ptr + 29, in.mac_ap, 6); } - wpa->nonce_compare = memcmp (in.nonce_ap, in.nonce_sta, 32); + wpa_eapol->nonce_compare = memcmp (in.nonce_ap, in.nonce_sta, 32); - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { memcpy (pke_ptr + 35, in.nonce_ap, 32); memcpy (pke_ptr + 67, in.nonce_sta, 32); @@ -3556,7 +3561,7 @@ int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED memcpy (pke_ptr + 67, in.nonce_ap, 32); } } - else if (wpa->keyver == 3) + else if (wpa_eapol->keyver == 3) { pke_ptr[0] = 1; pke_ptr[1] = 0; @@ -3574,9 +3579,9 @@ int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED memcpy (pke_ptr + 30, in.mac_ap, 6); } - wpa->nonce_compare = memcmp (in.nonce_ap, in.nonce_sta, 32); + wpa_eapol->nonce_compare = memcmp (in.nonce_ap, in.nonce_sta, 32); - if (wpa->nonce_compare < 0) + if (wpa_eapol->nonce_compare < 0) { memcpy (pke_ptr + 36, in.nonce_ap, 32); memcpy (pke_ptr + 68, in.nonce_sta, 32); @@ -3593,52 +3598,52 @@ int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED for (int i = 0; i < 32; i++) { - wpa->pke[i] = byte_swap_32 (wpa->pke[i]); + wpa_eapol->pke[i] = byte_swap_32 (wpa_eapol->pke[i]); } - memcpy (wpa->orig_mac_ap, in.mac_ap, 6); - memcpy (wpa->orig_mac_sta, in.mac_sta, 6); - memcpy (wpa->orig_nonce_ap, in.nonce_ap, 32); - memcpy (wpa->orig_nonce_sta, in.nonce_sta, 32); + memcpy (wpa_eapol->orig_mac_ap, in.mac_ap, 6); + memcpy (wpa_eapol->orig_mac_sta, in.mac_sta, 6); + memcpy (wpa_eapol->orig_nonce_ap, in.nonce_ap, 32); + memcpy (wpa_eapol->orig_nonce_sta, in.nonce_sta, 32); u8 message_pair_orig = in.message_pair; in.message_pair &= 0x7f; // ignore the highest bit (it is used to indicate if the replay counters did match) - if (wpa->message_pair_chgd == true) + if (wpa_eapol->message_pair_chgd == true) { - if (wpa->message_pair != in.message_pair) return (PARSER_HCCAPX_MESSAGE_PAIR); + if (wpa_eapol->message_pair != in.message_pair) return (PARSER_HCCAPX_MESSAGE_PAIR); } - wpa->message_pair = message_pair_orig; + wpa_eapol->message_pair = message_pair_orig; - wpa->eapol_len = in.eapol_len; + wpa_eapol->eapol_len = in.eapol_len; - u8 *eapol_ptr = (u8 *) wpa->eapol; + u8 *eapol_ptr = (u8 *) wpa_eapol->eapol; - memcpy (eapol_ptr, in.eapol, wpa->eapol_len); + memcpy (eapol_ptr, in.eapol, wpa_eapol->eapol_len); - memset (eapol_ptr + wpa->eapol_len, 0, (256 + 64) - wpa->eapol_len); + memset (eapol_ptr + wpa_eapol->eapol_len, 0, (256 + 64) - wpa_eapol->eapol_len); - eapol_ptr[wpa->eapol_len] = 0x80; + eapol_ptr[wpa_eapol->eapol_len] = 0x80; - if (wpa->keyver == 1) + if (wpa_eapol->keyver == 1) { // nothing to do } - else if (wpa->keyver == 2) + else if (wpa_eapol->keyver == 2) { - wpa->keymic[0] = byte_swap_32 (wpa->keymic[0]); - wpa->keymic[1] = byte_swap_32 (wpa->keymic[1]); - wpa->keymic[2] = byte_swap_32 (wpa->keymic[2]); - wpa->keymic[3] = byte_swap_32 (wpa->keymic[3]); + wpa_eapol->keymic[0] = byte_swap_32 (wpa_eapol->keymic[0]); + wpa_eapol->keymic[1] = byte_swap_32 (wpa_eapol->keymic[1]); + wpa_eapol->keymic[2] = byte_swap_32 (wpa_eapol->keymic[2]); + wpa_eapol->keymic[3] = byte_swap_32 (wpa_eapol->keymic[3]); for (int i = 0; i < 64; i++) { - wpa->eapol[i] = byte_swap_32 (wpa->eapol[i]); + wpa_eapol->eapol[i] = byte_swap_32 (wpa_eapol->eapol[i]); } } - else if (wpa->keyver == 3) + else if (wpa_eapol->keyver == 3) { // nothing to do } @@ -3664,58 +3669,58 @@ int wpa_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED md5_64 (block, hash); - for (int i = 0; i < 16; i++) block[i] = wpa->pke[i + 0]; + for (int i = 0; i < 16; i++) block[i] = wpa_eapol->pke[i + 0]; md5_64 (block, hash); - for (int i = 0; i < 16; i++) block[i] = wpa->pke[i + 16]; + for (int i = 0; i < 16; i++) block[i] = wpa_eapol->pke[i + 16]; md5_64 (block, hash); - for (int i = 0; i < 16; i++) block[i] = wpa->eapol[i + 0]; + for (int i = 0; i < 16; i++) block[i] = wpa_eapol->eapol[i + 0]; md5_64 (block, hash); - for (int i = 0; i < 16; i++) block[i] = wpa->eapol[i + 16]; + for (int i = 0; i < 16; i++) block[i] = wpa_eapol->eapol[i + 16]; md5_64 (block, hash); - for (int i = 0; i < 16; i++) block[i] = wpa->eapol[i + 32]; + for (int i = 0; i < 16; i++) block[i] = wpa_eapol->eapol[i + 32]; md5_64 (block, hash); - for (int i = 0; i < 16; i++) block[i] = wpa->eapol[i + 48]; + for (int i = 0; i < 16; i++) block[i] = wpa_eapol->eapol[i + 48]; md5_64 (block, hash); - for (int i = 0; i < 6; i++) block_ptr[i + 0] = wpa->orig_mac_ap[i]; - for (int i = 0; i < 6; i++) block_ptr[i + 6] = wpa->orig_mac_sta[i]; + for (int i = 0; i < 6; i++) block_ptr[i + 0] = wpa_eapol->orig_mac_ap[i]; + for (int i = 0; i < 6; i++) block_ptr[i + 6] = wpa_eapol->orig_mac_sta[i]; md5_64 (block, hash); - for (int i = 0; i < 32; i++) block_ptr[i + 0] = wpa->orig_nonce_ap[i]; - for (int i = 0; i < 32; i++) block_ptr[i + 32] = wpa->orig_nonce_sta[i]; + for (int i = 0; i < 32; i++) block_ptr[i + 0] = wpa_eapol->orig_nonce_ap[i]; + for (int i = 0; i < 32; i++) block_ptr[i + 32] = wpa_eapol->orig_nonce_sta[i]; md5_64 (block, hash); - block[0] = wpa->keymic[0]; - block[1] = wpa->keymic[1]; - block[2] = wpa->keymic[2]; - block[3] = wpa->keymic[3]; + block[0] = wpa_eapol->keymic[0]; + block[1] = wpa_eapol->keymic[1]; + block[2] = wpa_eapol->keymic[2]; + block[3] = wpa_eapol->keymic[3]; md5_64 (block, hash); - wpa->hash[0] = hash[0]; - wpa->hash[1] = hash[1]; - wpa->hash[2] = hash[2]; - wpa->hash[3] = hash[3]; + wpa_eapol->hash[0] = hash[0]; + wpa_eapol->hash[1] = hash[1]; + wpa_eapol->hash[2] = hash[2]; + wpa_eapol->hash[3] = hash[3]; // make all this stuff unique - digest[0] = wpa->hash[0]; - digest[1] = wpa->hash[1]; - digest[2] = wpa->hash[2]; - digest[3] = wpa->hash[3]; + digest[0] = wpa_eapol->hash[0]; + digest[1] = wpa_eapol->hash[1]; + digest[2] = wpa_eapol->hash[2]; + digest[3] = wpa_eapol->hash[3]; return (PARSER_OK); } @@ -17519,6 +17524,248 @@ int filevault2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE return (PARSER_OK); } +int wpa_pmkid_pbkdf2_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig) +{ + u32 *digest = (u32 *) hash_buf->digest; + + salt_t *salt = hash_buf->salt; + + wpa_pmkid_t *wpa_pmkid = (wpa_pmkid_t *) hash_buf->esalt; + + token_t token; + + token.token_cnt = 4; + + token.sep[0] = '*'; + token.len_min[0] = 32; + token.len_max[0] = 32; + token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.sep[1] = '*'; + token.len_min[1] = 12; + token.len_max[1] = 12; + token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.sep[2] = '*'; + token.len_min[2] = 12; + token.len_max[2] = 12; + token.attr[2] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.sep[3] = '*'; + token.len_min[3] = 0; + token.len_max[3] = 64; + token.attr[3] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + const int rc_tokenizer = input_tokenizer (input_buf, input_len, &token); + + if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); + + // pmkid + + u8 *pmkid_buf = token.buf[0]; + + wpa_pmkid->pmkid[0] = hex_to_u32 (pmkid_buf + 0); + wpa_pmkid->pmkid[1] = hex_to_u32 (pmkid_buf + 8); + wpa_pmkid->pmkid[2] = hex_to_u32 (pmkid_buf + 16); + wpa_pmkid->pmkid[3] = hex_to_u32 (pmkid_buf + 24); + + // mac_ap + + u8 *macap_buf = token.buf[1]; + + wpa_pmkid->orig_mac_ap[0] = hex_to_u8 (macap_buf + 0); + wpa_pmkid->orig_mac_ap[1] = hex_to_u8 (macap_buf + 2); + wpa_pmkid->orig_mac_ap[2] = hex_to_u8 (macap_buf + 4); + wpa_pmkid->orig_mac_ap[3] = hex_to_u8 (macap_buf + 6); + wpa_pmkid->orig_mac_ap[4] = hex_to_u8 (macap_buf + 8); + wpa_pmkid->orig_mac_ap[5] = hex_to_u8 (macap_buf + 10); + + // mac_sta + + u8 *macsta_buf = token.buf[2]; + + wpa_pmkid->orig_mac_sta[0] = hex_to_u8 (macsta_buf + 0); + wpa_pmkid->orig_mac_sta[1] = hex_to_u8 (macsta_buf + 2); + wpa_pmkid->orig_mac_sta[2] = hex_to_u8 (macsta_buf + 4); + wpa_pmkid->orig_mac_sta[3] = hex_to_u8 (macsta_buf + 6); + wpa_pmkid->orig_mac_sta[4] = hex_to_u8 (macsta_buf + 8); + wpa_pmkid->orig_mac_sta[5] = hex_to_u8 (macsta_buf + 10); + + // essid + + u8 *essid_buf = token.buf[3]; + int essid_len = token.len[3]; + + u8 *essid_ptr = (u8 *) wpa_pmkid->essid_buf; + + for (int i = 0, j = 0; i < essid_len; i += 2, j += 1) + { + essid_ptr[j] = hex_to_u8 (essid_buf + i); + } + + wpa_pmkid->essid_len = essid_len / 2; + + // pmkid_data + + wpa_pmkid->pmkid_data[0] = 0x204b4d50; // "PMK " + wpa_pmkid->pmkid_data[1] = 0x656d614e; // "Name" + wpa_pmkid->pmkid_data[2] = (wpa_pmkid->orig_mac_ap[0] << 0) + | (wpa_pmkid->orig_mac_ap[1] << 8) + | (wpa_pmkid->orig_mac_ap[2] << 16) + | (wpa_pmkid->orig_mac_ap[3] << 24); + wpa_pmkid->pmkid_data[3] = (wpa_pmkid->orig_mac_ap[4] << 0) + | (wpa_pmkid->orig_mac_ap[5] << 8) + | (wpa_pmkid->orig_mac_sta[0] << 16) + | (wpa_pmkid->orig_mac_sta[1] << 24); + wpa_pmkid->pmkid_data[4] = (wpa_pmkid->orig_mac_sta[2] << 0) + | (wpa_pmkid->orig_mac_sta[3] << 8) + | (wpa_pmkid->orig_mac_sta[4] << 16) + | (wpa_pmkid->orig_mac_sta[5] << 24); + + // salt + + salt->salt_buf[0] = wpa_pmkid->pmkid_data[0]; + salt->salt_buf[1] = wpa_pmkid->pmkid_data[1]; + salt->salt_buf[2] = wpa_pmkid->pmkid_data[2]; + salt->salt_buf[3] = wpa_pmkid->pmkid_data[3]; + salt->salt_buf[4] = wpa_pmkid->pmkid_data[4]; + salt->salt_buf[5] = wpa_pmkid->pmkid_data[5]; + salt->salt_buf[6] = wpa_pmkid->pmkid_data[6]; + salt->salt_buf[7] = wpa_pmkid->pmkid_data[7]; + + salt->salt_len = 32; + salt->salt_iter = ROUNDS_WPA_PBKDF2 - 1; + + // hash + + digest[0] = wpa_pmkid->pmkid[0]; + digest[1] = wpa_pmkid->pmkid[1]; + digest[2] = wpa_pmkid->pmkid[2]; + digest[3] = wpa_pmkid->pmkid[3]; + + digest[0] = byte_swap_32 (digest[0]); + digest[1] = byte_swap_32 (digest[1]); + digest[2] = byte_swap_32 (digest[2]); + digest[3] = byte_swap_32 (digest[3]); + + return (PARSER_OK); +} + +int wpa_pmkid_pmk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig) +{ + u32 *digest = (u32 *) hash_buf->digest; + + salt_t *salt = hash_buf->salt; + + wpa_pmkid_t *wpa_pmkid = (wpa_pmkid_t *) hash_buf->esalt; + + token_t token; + + token.token_cnt = 3; + + token.sep[0] = '*'; + token.len_min[0] = 32; + token.len_max[0] = 32; + token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.sep[1] = '*'; + token.len_min[1] = 12; + token.len_max[1] = 12; + token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + token.sep[2] = '*'; + token.len_min[2] = 12; + token.len_max[2] = 12; + token.attr[2] = TOKEN_ATTR_VERIFY_LENGTH + | TOKEN_ATTR_VERIFY_HEX; + + const int rc_tokenizer = input_tokenizer (input_buf, input_len, &token); + + if (rc_tokenizer != PARSER_OK) return (rc_tokenizer); + + // pmkid + + u8 *pmkid_buf = token.buf[0]; + + wpa_pmkid->pmkid[0] = hex_to_u32 (pmkid_buf + 0); + wpa_pmkid->pmkid[1] = hex_to_u32 (pmkid_buf + 8); + wpa_pmkid->pmkid[2] = hex_to_u32 (pmkid_buf + 16); + wpa_pmkid->pmkid[3] = hex_to_u32 (pmkid_buf + 24); + + // mac_ap + + u8 *macap_buf = token.buf[1]; + + wpa_pmkid->orig_mac_ap[0] = hex_to_u8 (macap_buf + 0); + wpa_pmkid->orig_mac_ap[1] = hex_to_u8 (macap_buf + 2); + wpa_pmkid->orig_mac_ap[2] = hex_to_u8 (macap_buf + 4); + wpa_pmkid->orig_mac_ap[3] = hex_to_u8 (macap_buf + 6); + wpa_pmkid->orig_mac_ap[4] = hex_to_u8 (macap_buf + 8); + wpa_pmkid->orig_mac_ap[5] = hex_to_u8 (macap_buf + 10); + + // mac_sta + + u8 *macsta_buf = token.buf[2]; + + wpa_pmkid->orig_mac_sta[0] = hex_to_u8 (macsta_buf + 0); + wpa_pmkid->orig_mac_sta[1] = hex_to_u8 (macsta_buf + 2); + wpa_pmkid->orig_mac_sta[2] = hex_to_u8 (macsta_buf + 4); + wpa_pmkid->orig_mac_sta[3] = hex_to_u8 (macsta_buf + 6); + wpa_pmkid->orig_mac_sta[4] = hex_to_u8 (macsta_buf + 8); + wpa_pmkid->orig_mac_sta[5] = hex_to_u8 (macsta_buf + 10); + + // pmkid_data + + wpa_pmkid->pmkid_data[0] = 0x204b4d50; // "PMK " + wpa_pmkid->pmkid_data[1] = 0x656d614e; // "Name" + wpa_pmkid->pmkid_data[2] = (wpa_pmkid->orig_mac_ap[0] << 0) + | (wpa_pmkid->orig_mac_ap[1] << 8) + | (wpa_pmkid->orig_mac_ap[2] << 16) + | (wpa_pmkid->orig_mac_ap[3] << 24); + wpa_pmkid->pmkid_data[3] = (wpa_pmkid->orig_mac_ap[4] << 0) + | (wpa_pmkid->orig_mac_ap[5] << 8) + | (wpa_pmkid->orig_mac_sta[0] << 16) + | (wpa_pmkid->orig_mac_sta[1] << 24); + wpa_pmkid->pmkid_data[4] = (wpa_pmkid->orig_mac_sta[2] << 0) + | (wpa_pmkid->orig_mac_sta[3] << 8) + | (wpa_pmkid->orig_mac_sta[4] << 16) + | (wpa_pmkid->orig_mac_sta[5] << 24); + + // salt + + salt->salt_buf[0] = wpa_pmkid->pmkid_data[0]; + salt->salt_buf[1] = wpa_pmkid->pmkid_data[1]; + salt->salt_buf[2] = wpa_pmkid->pmkid_data[2]; + salt->salt_buf[3] = wpa_pmkid->pmkid_data[3]; + salt->salt_buf[4] = wpa_pmkid->pmkid_data[4]; + salt->salt_buf[5] = wpa_pmkid->pmkid_data[5]; + salt->salt_buf[6] = wpa_pmkid->pmkid_data[6]; + salt->salt_buf[7] = wpa_pmkid->pmkid_data[7]; + + salt->salt_len = 32; + salt->salt_iter = ROUNDS_WPA_PBKDF2 - 1; + + // hash + + digest[0] = wpa_pmkid->pmkid[0]; + digest[1] = wpa_pmkid->pmkid[1]; + digest[2] = wpa_pmkid->pmkid[2]; + digest[3] = wpa_pmkid->pmkid[3]; + + digest[0] = byte_swap_32 (digest[0]); + digest[1] = byte_swap_32 (digest[1]); + digest[2] = byte_swap_32 (digest[2]); + digest[3] = byte_swap_32 (digest[3]); + + return (PARSER_OK); +} + /** * hook functions */ @@ -17959,6 +18206,8 @@ const char *strhashtype (const u32 hash_mode) case 16500: return HT_16500; case 16600: return HT_16600; case 16700: return HT_16700; + case 16800: return HT_16800; + case 16801: return HT_16801; case 99999: return HT_99999; } @@ -18049,49 +18298,49 @@ void to_hccapx_t (hashcat_ctx_t *hashcat_ctx, hccapx_t *hccapx, const u32 salt_p memcpy (hccapx->essid, salt->salt_buf, hccapx->essid_len); - wpa_t *wpas = (wpa_t *) esalts_buf; - wpa_t *wpa = &wpas[digest_cur]; + wpa_eapol_t *wpa_eapols = (wpa_eapol_t *) esalts_buf; + wpa_eapol_t *wpa_eapol = &wpa_eapols[digest_cur]; - hccapx->message_pair = wpa->message_pair; - hccapx->keyver = wpa->keyver; + hccapx->message_pair = wpa_eapol->message_pair; + hccapx->keyver = wpa_eapol->keyver; - hccapx->eapol_len = wpa->eapol_len; + hccapx->eapol_len = wpa_eapol->eapol_len; - if (wpa->keyver != 1) + if (wpa_eapol->keyver != 1) { u32 eapol_tmp[64] = { 0 }; for (u32 i = 0; i < 64; i++) { - eapol_tmp[i] = byte_swap_32 (wpa->eapol[i]); + eapol_tmp[i] = byte_swap_32 (wpa_eapol->eapol[i]); } - memcpy (hccapx->eapol, eapol_tmp, wpa->eapol_len); + memcpy (hccapx->eapol, eapol_tmp, wpa_eapol->eapol_len); } else { - memcpy (hccapx->eapol, wpa->eapol, wpa->eapol_len); + memcpy (hccapx->eapol, wpa_eapol->eapol, wpa_eapol->eapol_len); } - memcpy (hccapx->mac_ap, wpa->orig_mac_ap, 6); - memcpy (hccapx->mac_sta, wpa->orig_mac_sta, 6); - memcpy (hccapx->nonce_ap, wpa->orig_nonce_ap, 32); - memcpy (hccapx->nonce_sta, wpa->orig_nonce_sta, 32); + memcpy (hccapx->mac_ap, wpa_eapol->orig_mac_ap, 6); + memcpy (hccapx->mac_sta, wpa_eapol->orig_mac_sta, 6); + memcpy (hccapx->nonce_ap, wpa_eapol->orig_nonce_ap, 32); + memcpy (hccapx->nonce_sta, wpa_eapol->orig_nonce_sta, 32); - if (wpa->keyver != 1) + if (wpa_eapol->keyver != 1) { u32 digest_tmp[4]; - digest_tmp[0] = byte_swap_32 (wpa->keymic[0]); - digest_tmp[1] = byte_swap_32 (wpa->keymic[1]); - digest_tmp[2] = byte_swap_32 (wpa->keymic[2]); - digest_tmp[3] = byte_swap_32 (wpa->keymic[3]); + digest_tmp[0] = byte_swap_32 (wpa_eapol->keymic[0]); + digest_tmp[1] = byte_swap_32 (wpa_eapol->keymic[1]); + digest_tmp[2] = byte_swap_32 (wpa_eapol->keymic[2]); + digest_tmp[3] = byte_swap_32 (wpa_eapol->keymic[3]); memcpy (hccapx->keymic, digest_tmp, 16); } else { - memcpy (hccapx->keymic, wpa->keymic, 16); + memcpy (hccapx->keymic, wpa_eapol->keymic, 16); } } @@ -18848,15 +19097,15 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le } else if ((hash_mode == 2500) || (hash_mode == 2501)) { - wpa_t *wpas = (wpa_t *) esalts_buf; + wpa_eapol_t *wpa_eapols = (wpa_eapol_t *) esalts_buf; - wpa_t *wpa = &wpas[digest_cur]; + wpa_eapol_t *wpa_eapol = &wpa_eapols[digest_cur]; - char *essid = (char *) wpa->essid; + char *essid = (char *) wpa_eapol->essid; int tmp_len = 0; - if (need_hexify (wpa->essid, wpa->essid_len, hashconfig->separator, 0) == true) + if (need_hexify (wpa_eapol->essid, wpa_eapol->essid_len, hashconfig->separator, 0) == true) { tmp_buf[tmp_len++] = '$'; tmp_buf[tmp_len++] = 'H'; @@ -18864,9 +19113,9 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le tmp_buf[tmp_len++] = 'X'; tmp_buf[tmp_len++] = '['; - exec_hexify (wpa->essid, wpa->essid_len, (u8 *) tmp_buf + tmp_len); + exec_hexify (wpa_eapol->essid, wpa_eapol->essid_len, (u8 *) tmp_buf + tmp_len); - tmp_len += wpa->essid_len * 2; + tmp_len += wpa_eapol->essid_len * 2; tmp_buf[tmp_len++] = ']'; @@ -18874,22 +19123,22 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le } snprintf (out_buf, out_len - 1, "%08x%08x%08x%08x:%02x%02x%02x%02x%02x%02x:%02x%02x%02x%02x%02x%02x:%s", - wpa->hash[0], - wpa->hash[1], - wpa->hash[2], - wpa->hash[3], - wpa->orig_mac_ap[0], - wpa->orig_mac_ap[1], - wpa->orig_mac_ap[2], - wpa->orig_mac_ap[3], - wpa->orig_mac_ap[4], - wpa->orig_mac_ap[5], - wpa->orig_mac_sta[0], - wpa->orig_mac_sta[1], - wpa->orig_mac_sta[2], - wpa->orig_mac_sta[3], - wpa->orig_mac_sta[4], - wpa->orig_mac_sta[5], + wpa_eapol->hash[0], + wpa_eapol->hash[1], + wpa_eapol->hash[2], + wpa_eapol->hash[3], + wpa_eapol->orig_mac_ap[0], + wpa_eapol->orig_mac_ap[1], + wpa_eapol->orig_mac_ap[2], + wpa_eapol->orig_mac_ap[3], + wpa_eapol->orig_mac_ap[4], + wpa_eapol->orig_mac_ap[5], + wpa_eapol->orig_mac_sta[0], + wpa_eapol->orig_mac_sta[1], + wpa_eapol->orig_mac_sta[2], + wpa_eapol->orig_mac_sta[3], + wpa_eapol->orig_mac_sta[4], + wpa_eapol->orig_mac_sta[5], essid); } else if (hash_mode == 4400) @@ -21581,6 +21830,61 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le byte_swap_32 (apple_secure_notes->ZCRYPTOWRAPPEDKEY[4]), byte_swap_32 (apple_secure_notes->ZCRYPTOWRAPPEDKEY[5])); } + else if (hash_mode == 16800) + { + wpa_pmkid_t *wpa_pmkids = (wpa_pmkid_t *) esalts_buf; + + wpa_pmkid_t *wpa_pmkid = &wpa_pmkids[digest_cur]; + + exec_hexify ((const u8*) wpa_pmkid->essid_buf, wpa_pmkid->essid_len, (u8 *) tmp_buf); + + int tmp_len = wpa_pmkid->essid_len * 2; + + tmp_buf[tmp_len] = 0; + + snprintf (out_buf, out_len - 1, "%08x%08x%08x%08x*%02x%02x%02x%02x%02x%02x*%02x%02x%02x%02x%02x%02x*%s", + byte_swap_32 (wpa_pmkid->pmkid[0]), + byte_swap_32 (wpa_pmkid->pmkid[1]), + byte_swap_32 (wpa_pmkid->pmkid[2]), + byte_swap_32 (wpa_pmkid->pmkid[3]), + wpa_pmkid->orig_mac_ap[0], + wpa_pmkid->orig_mac_ap[1], + wpa_pmkid->orig_mac_ap[2], + wpa_pmkid->orig_mac_ap[3], + wpa_pmkid->orig_mac_ap[4], + wpa_pmkid->orig_mac_ap[5], + wpa_pmkid->orig_mac_sta[0], + wpa_pmkid->orig_mac_sta[1], + wpa_pmkid->orig_mac_sta[2], + wpa_pmkid->orig_mac_sta[3], + wpa_pmkid->orig_mac_sta[4], + wpa_pmkid->orig_mac_sta[5], + tmp_buf); + } + else if (hash_mode == 16801) + { + wpa_pmkid_t *wpa_pmkids = (wpa_pmkid_t *) esalts_buf; + + wpa_pmkid_t *wpa_pmkid = &wpa_pmkids[digest_cur]; + + snprintf (out_buf, out_len - 1, "%08x%08x%08x%08x*%02x%02x%02x%02x%02x%02x*%02x%02x%02x%02x%02x%02x", + wpa_pmkid->pmkid[0], + wpa_pmkid->pmkid[1], + wpa_pmkid->pmkid[2], + wpa_pmkid->pmkid[3], + wpa_pmkid->orig_mac_ap[0], + wpa_pmkid->orig_mac_ap[1], + wpa_pmkid->orig_mac_ap[2], + wpa_pmkid->orig_mac_ap[3], + wpa_pmkid->orig_mac_ap[4], + wpa_pmkid->orig_mac_ap[5], + wpa_pmkid->orig_mac_sta[0], + wpa_pmkid->orig_mac_sta[1], + wpa_pmkid->orig_mac_sta[2], + wpa_pmkid->orig_mac_sta[3], + wpa_pmkid->orig_mac_sta[4], + wpa_pmkid->orig_mac_sta[5]); + } else if (hash_mode == 99999) { char *ptr = (char *) digest_buf; @@ -23451,7 +23755,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN; break; - case 2500: hashconfig->hash_type = HASH_TYPE_WPA; + case 2500: hashconfig->hash_type = HASH_TYPE_WPA_EAPOL; hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE @@ -23459,9 +23763,9 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) | OPTS_TYPE_AUX2 | OPTS_TYPE_AUX3 | OPTS_TYPE_BINARY_HASHFILE; - hashconfig->kern_type = KERN_TYPE_WPA; + hashconfig->kern_type = KERN_TYPE_WPA_EAPOL_PBKDF2; hashconfig->dgst_size = DGST_SIZE_4_4; - hashconfig->parse_func = wpa_parse_hash; + hashconfig->parse_func = wpa_eapol_parse_hash; hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; hashconfig->dgst_pos0 = 0; @@ -23472,7 +23776,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->st_pass = ST_PASS_HASHCAT_EXCL; break; - case 2501: hashconfig->hash_type = HASH_TYPE_WPA; + case 2501: hashconfig->hash_type = HASH_TYPE_WPA_EAPOL; hashconfig->salt_type = SALT_TYPE_EMBEDDED; hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE @@ -23480,9 +23784,9 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) | OPTS_TYPE_AUX2 | OPTS_TYPE_AUX3 | OPTS_TYPE_BINARY_HASHFILE; - hashconfig->kern_type = KERN_TYPE_WPAPMK; + hashconfig->kern_type = KERN_TYPE_WPA_EAPOL_PMK; hashconfig->dgst_size = DGST_SIZE_4_4; - hashconfig->parse_func = wpa_parse_hash; + hashconfig->parse_func = wpa_eapol_parse_hash; hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; hashconfig->dgst_pos0 = 0; @@ -26725,6 +27029,40 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN; break; + case 16800: hashconfig->hash_type = HASH_TYPE_WPA_PMKID_PBKDF2; + hashconfig->salt_type = SALT_TYPE_EMBEDDED; + hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; + hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE; + hashconfig->kern_type = KERN_TYPE_WPA_PMKID_PBKDF2; + hashconfig->dgst_size = DGST_SIZE_4_4; + hashconfig->parse_func = wpa_pmkid_pbkdf2_parse_hash; + hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + hashconfig->dgst_pos0 = 0; + hashconfig->dgst_pos1 = 1; + hashconfig->dgst_pos2 = 2; + hashconfig->dgst_pos3 = 3; + hashconfig->st_hash = ST_HASH_16800; + hashconfig->st_pass = ST_PASS_HASHCAT_EXCL; + break; + + case 16801: hashconfig->hash_type = HASH_TYPE_WPA_PMKID_PMK; + hashconfig->salt_type = SALT_TYPE_EMBEDDED; + hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL; + hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE; + hashconfig->kern_type = KERN_TYPE_WPA_PMKID_PMK; + hashconfig->dgst_size = DGST_SIZE_4_4; + hashconfig->parse_func = wpa_pmkid_pmk_parse_hash; + hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE + | OPTI_TYPE_SLOW_HASH_SIMD_LOOP; + hashconfig->dgst_pos0 = 0; + hashconfig->dgst_pos1 = 1; + hashconfig->dgst_pos2 = 2; + hashconfig->dgst_pos3 = 3; + hashconfig->st_hash = ST_HASH_16801; + hashconfig->st_pass = ST_PASS_HEX_16801; + break; + case 99999: hashconfig->hash_type = HASH_TYPE_PLAINTEXT; hashconfig->salt_type = SALT_TYPE_NONE; hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL; @@ -26857,8 +27195,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) switch (hashconfig->hash_mode) { case 600: hashconfig->esalt_size = sizeof (blake2_t); break; - case 2500: hashconfig->esalt_size = sizeof (wpa_t); break; - case 2501: hashconfig->esalt_size = sizeof (wpa_t); break; + case 2500: hashconfig->esalt_size = sizeof (wpa_eapol_t); break; + case 2501: hashconfig->esalt_size = sizeof (wpa_eapol_t); break; case 5300: hashconfig->esalt_size = sizeof (ikepsk_t); break; case 5400: hashconfig->esalt_size = sizeof (ikepsk_t); break; case 5500: hashconfig->esalt_size = sizeof (netntlm_t); break; @@ -26946,6 +27284,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) case 16500: hashconfig->esalt_size = sizeof (jwt_t); break; case 16600: hashconfig->esalt_size = sizeof (electrum_wallet_t); break; case 16700: hashconfig->esalt_size = sizeof (apple_secure_notes_t); break; + case 16800: hashconfig->esalt_size = sizeof (wpa_pmkid_t); break; + case 16801: hashconfig->esalt_size = sizeof (wpa_pmkid_t); break; } // hook_salt_size @@ -26969,8 +27309,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) case 1600: hashconfig->tmp_size = sizeof (md5crypt_tmp_t); break; case 1800: hashconfig->tmp_size = sizeof (sha512crypt_tmp_t); break; case 2100: hashconfig->tmp_size = sizeof (dcc2_tmp_t); break; - case 2500: hashconfig->tmp_size = sizeof (wpa_tmp_t); break; - case 2501: hashconfig->tmp_size = sizeof (wpapmk_tmp_t); break; + case 2500: hashconfig->tmp_size = sizeof (wpa_pbkdf2_tmp_t); break; + case 2501: hashconfig->tmp_size = sizeof (wpa_pmk_tmp_t); break; case 3200: hashconfig->tmp_size = sizeof (bcrypt_tmp_t); break; case 5200: hashconfig->tmp_size = sizeof (pwsafe3_tmp_t); break; case 5800: hashconfig->tmp_size = sizeof (androidpin_tmp_t); break; @@ -27056,6 +27396,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) case 16200: hashconfig->tmp_size = sizeof (apple_secure_notes_tmp_t); break; case 16300: hashconfig->tmp_size = sizeof (pbkdf2_sha256_tmp_t); break; case 16700: hashconfig->tmp_size = sizeof (apple_secure_notes_tmp_t); break; + case 16800: hashconfig->tmp_size = sizeof (wpa_pbkdf2_tmp_t); break; + case 16801: hashconfig->tmp_size = sizeof (wpa_pmk_tmp_t); break; }; // hook_size @@ -27243,8 +27585,8 @@ int hashconfig_get_pw_min (hashcat_ctx_t *hashcat_ctx, const bool optimized_kern switch (hashconfig->hash_mode) { - case 2500: pw_min = 8; break; // WPA min RFC - case 2501: pw_min = 64; break; // WPA PMK fixed + case 2500: pw_min = 8; break; // WPA-EAPOL-PBKDF2: min RFC + case 2501: pw_min = 64; break; // WPA-EAPOL-PMK: fixed case 9710: pw_min = 5; break; // RC4-40 fixed case 9810: pw_min = 5; break; // RC4-40 fixed case 10410: pw_min = 5; break; // RC4-40 fixed @@ -27252,6 +27594,8 @@ int hashconfig_get_pw_min (hashcat_ctx_t *hashcat_ctx, const bool optimized_kern case 14100: pw_min = 24; break; // 3DES fixed case 14900: pw_min = 10; break; // Skip32 fixed case 15400: pw_min = 32; break; // ChaCha20 fixed + case 16800: pw_min = 8; break; // WPA-PMKID-PBKDF2: min RFC + case 16801: pw_min = 64; break; // WPA-PMKID-PMK: fixed } return pw_min; @@ -27358,8 +27702,8 @@ int hashconfig_get_pw_max (hashcat_ctx_t *hashcat_ctx, const bool optimized_kern case 112: pw_max = 30; break; // https://www.toadworld.com/platforms/oracle/b/weblog/archive/2013/11/12/oracle-12c-passwords case 1500: pw_max = 8; break; // Underlaying DES max case 2100: pw_max = PW_MAX; break; - case 2500: pw_max = 63; break; // WPA/WPA2 limits itself to 63 by RFC - case 2501: pw_max = 64; break; // WPA/WPA2 PMK fixed length + case 2500: pw_max = 63; break; // WPA-EAPOL-PBKDF2: limits itself to 63 by RFC + case 2501: pw_max = 64; break; // WPA-EAPOL-PMK: fixed length case 3000: pw_max = 7; break; // LM max case 3100: pw_max = 30; break; // http://www.red-database-security.de/whitepaper/oracle_passwords.html case 3200: pw_max = 72; break; // Underlaying Blowfish max @@ -27471,6 +27815,8 @@ int hashconfig_get_pw_max (hashcat_ctx_t *hashcat_ctx, const bool optimized_kern case 15700: pw_max = PW_MAX; break; case 15900: pw_max = PW_MAX; break; case 16000: pw_max = 8; break; // Underlaying DES max + case 16800: pw_max = 63; break; // WPA-PMKID-PBKDF2: limits itself to 63 by RFC + case 16801: pw_max = 64; break; // WPA-PMKID-PMK: fixed length } return pw_max; @@ -27758,15 +28104,19 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo break; case 16700: salt->salt_len = 16; break; + case 16800: memcpy (salt->salt_buf, "hashcat.net", 11); + break; + case 16801: memcpy (salt->salt_buf, "hashcat.net", 11); + break; } // special esalt handling switch (hashconfig->hash_mode) { - case 2500: ((wpa_t *) esalt)->eapol_len = 128; + case 2500: ((wpa_eapol_t *) esalt)->eapol_len = 128; break; - case 2501: ((wpa_t *) esalt)->eapol_len = 128; + case 2501: ((wpa_eapol_t *) esalt)->eapol_len = 128; break; case 5300: ((ikepsk_t *) esalt)->nr_len = 1; ((ikepsk_t *) esalt)->msg_len = 1; @@ -27860,9 +28210,9 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo break; case 2100: salt->salt_iter = ROUNDS_DCC2; break; - case 2500: salt->salt_iter = ROUNDS_WPA; + case 2500: salt->salt_iter = ROUNDS_WPA_PBKDF2; break; - case 2501: salt->salt_iter = ROUNDS_WPAPMK; + case 2501: salt->salt_iter = ROUNDS_WPA_PMK; break; case 3200: salt->salt_iter = ROUNDS_BCRYPT; break; @@ -28041,6 +28391,10 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo break; case 16700: salt->salt_iter = ROUNDS_APPLE_SECURE_NOTES - 1; break; + case 16800: salt->salt_iter = ROUNDS_WPA_PBKDF2; + break; + case 16801: salt->salt_iter = ROUNDS_WPA_PMK; + break; } } @@ -28070,6 +28424,10 @@ const char *hashconfig_benchmark_mask (hashcat_ctx_t *hashcat_ctx) break; case 14900: mask = "?b?b?b?b?bxxxxx"; break; + case 16800: mask = "?a?a?a?a?a?a?a?a"; + break; + case 16801: mask = "?a?a?a?a?a?a?a?axxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"; + break; default: mask = "?b?b?b?b?b?b?b"; break; } diff --git a/src/opencl.c b/src/opencl.c index c73ed58cd..f6efed500 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -1412,23 +1412,23 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 digests_offset = hashes->salts_buf[salt_pos].digests_offset; - wpa_t *wpas = (wpa_t *) hashes->esalts_buf; + wpa_eapol_t *wpa_eapols = (wpa_eapol_t *) hashes->esalts_buf; - wpa_t *wpa = &wpas[digests_offset + loops_pos]; + wpa_eapol_t *wpa_eapol = &wpa_eapols[digests_offset + loops_pos]; - if (wpa->keyver == 1) + if (wpa_eapol->keyver == 1) { CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX1, pws_cnt, false, 0); if (CL_rc == -1) return -1; } - else if (wpa->keyver == 2) + else if (wpa_eapol->keyver == 2) { CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX2, pws_cnt, false, 0); if (CL_rc == -1) return -1; } - else if (wpa->keyver == 3) + else if (wpa_eapol->keyver == 3) { CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_AUX3, pws_cnt, false, 0); diff --git a/src/potfile.c b/src/potfile.c index 5e04a4f9b..9a6f5f452 100644 --- a/src/potfile.c +++ b/src/potfile.c @@ -614,7 +614,7 @@ int potfile_remove_parse (hashcat_ctx_t *hashcat_ctx) memcpy (hash_buf.salt->salt_buf, essid_pos, essid_len); hash_buf.salt->salt_len = (u32) essid_len; - hash_buf.salt->salt_iter = ROUNDS_WPA - 1; + hash_buf.salt->salt_iter = ROUNDS_WPA_PBKDF2 - 1; u32 hash[4]; diff --git a/src/status.c b/src/status.c index c2ffe2716..6a13c274a 100644 --- a/src/status.c +++ b/src/status.c @@ -259,22 +259,22 @@ const char *status_get_hash_target (const hashcat_ctx_t *hashcat_ctx) { char *tmp_buf; - wpa_t *wpa = (wpa_t *) hashes->esalts_buf; + wpa_eapol_t *wpa_eapol = (wpa_eapol_t *) hashes->esalts_buf; hc_asprintf (&tmp_buf, "%s (AP:%02x:%02x:%02x:%02x:%02x:%02x STA:%02x:%02x:%02x:%02x:%02x:%02x)", (char *) hashes->salts_buf[0].salt_buf, - wpa->orig_mac_ap[0], - wpa->orig_mac_ap[1], - wpa->orig_mac_ap[2], - wpa->orig_mac_ap[3], - wpa->orig_mac_ap[4], - wpa->orig_mac_ap[5], - wpa->orig_mac_sta[0], - wpa->orig_mac_sta[1], - wpa->orig_mac_sta[2], - wpa->orig_mac_sta[3], - wpa->orig_mac_sta[4], - wpa->orig_mac_sta[5]); + wpa_eapol->orig_mac_ap[0], + wpa_eapol->orig_mac_ap[1], + wpa_eapol->orig_mac_ap[2], + wpa_eapol->orig_mac_ap[3], + wpa_eapol->orig_mac_ap[4], + wpa_eapol->orig_mac_ap[5], + wpa_eapol->orig_mac_sta[0], + wpa_eapol->orig_mac_sta[1], + wpa_eapol->orig_mac_sta[2], + wpa_eapol->orig_mac_sta[3], + wpa_eapol->orig_mac_sta[4], + wpa_eapol->orig_mac_sta[5]); return tmp_buf; } diff --git a/src/usage.c b/src/usage.c index ac42f6ad5..1fd6b6f87 100644 --- a/src/usage.c +++ b/src/usage.c @@ -184,8 +184,10 @@ static const char *const USAGE_BIG[] = " 10900 | PBKDF2-HMAC-SHA256 | Generic KDF", " 12100 | PBKDF2-HMAC-SHA512 | Generic KDF", " 23 | Skype | Network Protocols", - " 2500 | WPA/WPA2 | Network Protocols", - " 2501 | WPA/WPA2 PMK | Network Protocols", + " 2500 | WPA-EAPOL-PBKDF2 | Network Protocols", + " 2501 | WPA-EAPOL-PMK | Network Protocols", + " 16800 | WPA-PMKID-PBKDF2 | Network Protocols", + " 16801 | WPA-PMKID-PMK | Network Protocols", " 4800 | iSCSI CHAP authentication, MD5(CHAP) | Network Protocols", " 5300 | IKE-PSK MD5 | Network Protocols", " 5400 | IKE-PSK SHA1 | Network Protocols", diff --git a/tools/test.pl b/tools/test.pl index 76df928e6..68bf4580e 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -57,7 +57,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, 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, 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); @@ -1675,10 +1675,10 @@ sub verify next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in}))); } - # WPA/WPA2 + # WPA-EAPOL-PBKDF2 elsif ($mode == 2500) { - print "ERROR: verify currently not supported for WPA/WPA2 (because of hashcat's output format)\n"; + print "ERROR: verify currently not supported for WPA-EAPOL-PBKDF2 (because of hashcat's output format)\n"; exit (1); } @@ -2954,6 +2954,26 @@ sub verify next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in}))); } + # WPA-PMKID-PMKDF2 + elsif ($mode == 16800) + { + ($hash_in, $word) = split ":", $line; + + next unless defined $hash_in; + next unless defined $word; + + my @data = split (/\*/, $hash_in); + + next unless scalar @data == 4; + + my ($pmkid, $macap, $macsta, $essid) = @data; + + $param = $macap; + $param2 = $macsta; + $param3 = $essid; + + next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in}))); + } else { print "ERROR: hash mode is not supported\n"; @@ -3409,6 +3429,14 @@ sub verify return unless (substr ($line, 0, $len) eq $hash_out); } + elsif ($mode == 16800) + { + $hash_out = gen_hash ($mode, $word, undef, 0, $param, $param2, $param3); + + $len = length $hash_out; + + return unless (substr ($line, 0, $len) eq $hash_out); + } else { $hash_out = gen_hash ($mode, $word, $salt, $iter); @@ -3984,6 +4012,12 @@ sub passthrough { $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32)); } + elsif ($mode == 16800) + { + next if length ($word_buf) < 8; + + $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32)); + } else { print "ERROR: Unsupported hash type\n"; @@ -5081,6 +5115,27 @@ sub single } } } + elsif ($mode == 16800) + { + my $salt_len = get_random_num (0, 32); + + for (my $i = 8; $i < 16; $i++) + { + if ($len != 0) + { + if ($len < 8) + { + $len += 7; + } + + rnd ($mode, $len, $salt_len); + } + else + { + rnd ($mode, $i, $salt_len); + } + } + } } } @@ -9739,6 +9794,63 @@ END_CODE $tmp_hash = sprintf ('$fvde$%d$%d$%s$%d$%s', $Z_PK, length ($salt_bin), unpack ("H*", $salt_bin), $iterations, unpack ("H*", $blob_bin)); } + elsif ($mode == 16800) + { + my $macap; + my $macsta; + my $essid; + + if (!defined ($additional_param)) + { + $macap = unpack ("H*", randbytes (6)); + } + else + { + $macap = $additional_param; + } + + if (!defined ($additional_param2)) + { + $macsta = unpack ("H*", randbytes (6)); + } + else + { + $macsta = $additional_param2; + } + + if (!defined ($additional_param3)) + { + $essid = unpack ("H*", randbytes (get_random_num (8, 32) & 0x1e)); + } + else + { + $essid = $additional_param3; + } + + # generate the Pairwise Master Key (PMK) + + my $iterations = 4096; + + my $pbkdf2 = Crypt::PBKDF2->new + ( + hash_class => 'HMACSHA1', + iterations => $iterations, + output_len => 32, + ); + + my $essid_bin = pack ("H*", $essid); + + my $pmk = $pbkdf2->PBKDF2 ($essid_bin, $word_buf); + + my $macap_bin = pack ("H*", $macap); + my $macsta_bin = pack ("H*", $macsta); + + my $data = "PMK Name" . $macap_bin . $macsta_bin; + + my $pmkid = hmac_hex ($data, $pmk, \&sha1); + + $tmp_hash = sprintf ("%s*%s*%s*%s", substr ($pmkid, 0, 32), $macap, $macsta, $essid); + } elsif ($mode == 99999) { $tmp_hash = sprintf ("%s", $word_buf); diff --git a/tools/test.sh b/tools/test.sh index 6c86ffa3d..44468bd0b 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 5000 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 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 5000 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 99999" #ATTACK_MODES="0 1 3 6 7" ATTACK_MODES="0 1 3 7" @@ -22,7 +22,7 @@ HASHFILE_ONLY="2500" NEVER_CRACK="11600 14900" -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 14600 14700 14800 15100 15200 15300 15600 15700 15900" +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" OPTS="--quiet --force --potfile-disable --runtime 400 --gpu-temp-disable" @@ -256,6 +256,8 @@ function init() elif [ "${hash_type}" -eq 15400 ]; then min=0 min_offset=3 + elif [ "${hash_type}" -eq 16800 ]; then + min_offset=7 # means length 8, since we start with 0 fi # foreach password entry split password in 2 (skip first entry, is len 1) @@ -310,6 +312,8 @@ function init() min_len=9 elif [ "${hash_type}" -eq 15400 ]; then min_len=31 + elif [ "${hash_type}" -eq 16800 ]; then + min_len=7 # means length 8, since we start with 0 fi # generate multiple pass/hash foreach len (2 to 8) @@ -726,6 +730,8 @@ function attack_1() offset=7 elif [ ${hash_type} -eq 16000 ]; then offset=7 + elif [ ${hash_type} -eq 16800 ]; then + offset=7 fi hash_file=${OUTD}/${hash_type}_multihash_combi.txt @@ -855,6 +861,9 @@ function attack_3() elif [ "${hash_type}" -eq 15400 ]; then mask_offset=3 max=1 + elif [ "${hash_type}" -eq 16800 ]; then + mask_offset=7 + max=7 fi # special case: we need to split the first line @@ -936,6 +945,28 @@ function attack_3() fi + if [ "${hash_type}" -eq 16800 ]; then + + pass=$(sed -n ${i}p ${dict}) + + mask=${pass} + + # replace the first x positions in the mask with ?d's + + # first: remove first i (== amount) chars + + mask=$(echo ${mask} | cut -b $((i + 1))-) + + # prepend the ?d's + + for i in $(seq 1 ${i}); do + + mask="?d${mask}" + + done + + fi + if [ "${mask_offset}" -ne 0 ]; then mask=${mask_custom} fi @@ -1029,6 +1060,11 @@ function attack_3() increment_max=9 fi + if [ "${hash_type}" -eq 16800 ]; then + increment_min=8 + increment_max=9 + fi + hash_file=${OUTD}/${hash_type}_multihash_bruteforce.txt head -n $((increment_max - ${increment_min} + 1)) ${OUTD}/${hash_type}_hashes.txt > ${hash_file} @@ -1146,6 +1182,91 @@ function attack_3() custom_charsets="-1 ${charset_1} -2 ${charset_2} -3 ${charset_3} -4 ${charset_4}" fi + if [ "${hash_type}" -eq 16800 ]; then + + mask="?d?d?d?d?d?1?2?3?4" + + charset_1="" + charset_2="" + charset_3="" + charset_4="" + + # check positions (here we assume that mask is always composed of non literal chars + # i.e. something like ?d?l?u?s?1 is possible, but ?d?dsuffix not + charset_1_pos=$(expr index "${mask}" 1) + charset_2_pos=$(expr index "${mask}" 2) + charset_3_pos=$(expr index "${mask}" 3) + charset_4_pos=$(expr index "${mask}" 4) + + # divide each charset position by 2 since each of them occupies 2 positions in the mask + + charset_1_pos=$((charset_1_pos / 2)) + charset_2_pos=$((charset_2_pos / 2)) + charset_3_pos=$((charset_3_pos / 2)) + charset_4_pos=$((charset_4_pos / 2)) + + i=1 + + while read -u 9 hash; do + + pass=$(sed -n ${i}p ${OUTD}/${hash_type}_passwords.txt) + + # charset 1 + char=$(echo "${pass}" | cut -b ${charset_1_pos}) + charset_1=$(echo -e "${charset_1}\n${char}") + + # charset 2 + char=$(echo "${pass}" | cut -b ${charset_2_pos}) + charset_2=$(echo -e "${charset_2}\n${char}") + + # charset 3 + char=$(echo "${pass}" | cut -b ${charset_3_pos}) + charset_3=$(echo -e "${charset_3}\n${char}") + + # charset 4 + char=$(echo "${pass}" | cut -b ${charset_4_pos}) + charset_4=$(echo -e "${charset_4}\n${char}") + + i=$((i + 1)) + + done 9< ${OUTD}/${hash_type}_multihash_bruteforce.txt + + # just make sure that all custom charset fields are initialized + + if [ -z "${charset_1}" ]; then + + charset_1="1" + + fi + + if [ -z "${charset_2}" ]; then + + charset_2="2" + + fi + + if [ -z "${charset_3}" ]; then + + charset_3="3" + + fi + + if [ -z "${charset_4}" ]; then + + charset_4="4" + + fi + + # unique and remove new lines + + charset_1=$(echo "${charset_1}" | sort -u | tr -d '\n') + charset_2=$(echo "${charset_2}" | sort -u | tr -d '\n') + charset_3=$(echo "${charset_3}" | sort -u | tr -d '\n') + charset_4=$(echo "${charset_4}" | sort -u | tr -d '\n') + + custom_charsets="-1 ${charset_1} -2 ${charset_2} -3 ${charset_3} -4 ${charset_4}" + fi + CMD="./${BIN} ${OPTS} -a 3 -m ${hash_type} --increment --increment-min ${increment_min} --increment-max ${increment_max} ${custom_charsets} ${hash_file} ${mask} " echo "> Testing hash type $hash_type with attack mode 3, markov ${MARKOV}, multi hash, Device-Type ${TYPE}, vector-width ${VECTOR}." &>> ${OUTD}/logfull.txt @@ -1247,6 +1368,8 @@ function attack_6() min=0 max=1 mask_offset=29 + elif [ "${hash_type}" -eq 16800 ]; then + max=6 fi # special case: we need to split the first line @@ -1418,6 +1541,8 @@ function attack_6() max=8 elif [ ${hash_type} -eq 8500 ]; then max=8 + elif [ ${hash_type} -eq 16800 ]; then + max=5 fi if ! contains ${hash_type} ${TIMEOUT_ALGOS}; then @@ -1560,6 +1685,8 @@ function attack_7() mask_offset=3 min=0 max=1 + elif [ "${hash_type}" -eq 16800 ]; then + max=5 fi # special case: we need to split the first line @@ -1642,6 +1769,32 @@ function attack_7() fi + if [ "${hash_type}" -eq 16800 ]; then + + line_nr=1 + + if [ "${i}" -gt 1 ]; then + line_nr=$((${i} - 1)) + fi + + pass_part_1=$(sed -n ${line_nr}p ${OUTD}/${hash_type}_dict1) + pass_part_2=$(sed -n ${line_nr}p ${OUTD}/${hash_type}_dict2) + + pass_part_2_len=${#pass_part_2} + + pass=${pass_part_1}${pass_part_2} + pass_len=${#pass} + + # add first x chars of password to mask and append the (old) mask + + mask_len=${#mask} + mask_len=$((mask_len / 2)) + + mask_prefix=$(echo ${pass} | cut -b -$((pass_len - ${mask_len} - ${pass_part_2_len}))) + mask=${mask_prefix}${mask} + + fi + dict1=${OUTD}/${hash_type}_dict1 dict2=${OUTD}/${hash_type}_dict2 @@ -1756,6 +1909,8 @@ function attack_7() max=5 elif [ ${hash_type} -eq 15400 ]; then max=5 + elif [ ${hash_type} -eq 16800 ]; then + max=5 fi if ! contains ${hash_type} ${TIMEOUT_ALGOS}; then