From 547025ec475c2aaa4edf16184e91c24fbefe08ff Mon Sep 17 00:00:00 2001 From: jsteube Date: Fri, 15 Jun 2018 17:00:41 +0200 Subject: [PATCH] HCCAPX management: Use advanced hints in message_pair stored by hcxtools about endian bitness of replay counter Fixed missing code section in -m 2500 and -m 2501 to crack corrupted handshakes with a LE endian bitness base --- OpenCL/inc_types.cl | 2 + OpenCL/m02500.cl | 1422 +++++++++++++++++++++++++++++-------------- OpenCL/m02501.cl | 1098 ++++++++++++++++++++++----------- docs/changes.txt | 2 + include/interface.h | 2 + include/types.h | 1 + src/hashes.c | 52 +- src/user_options.c | 3 +- 8 files changed, 1762 insertions(+), 820 deletions(-) diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index ca93b946f..90bb8621c 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -1207,6 +1207,8 @@ typedef struct wpa u32 hash[4]; int nonce_compare; int nonce_error_corrections; + int detected_le; + int detected_be; } wpa_t; diff --git a/OpenCL/m02500.cl b/OpenCL/m02500.cl index 468da9bd5..d174bb64c 100644 --- a/OpenCL/m02500.cl +++ b/OpenCL/m02500.cl @@ -321,155 +321,310 @@ __kernel void m02500_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 nonce_error_corrections = wpa->nonce_error_corrections; - for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + if (wpa->detected_le == 1) { - u32 t = to; - - t = swap32_S (t); - - t -= nonce_error_corrections / 2; - t += nonce_error_correction; - - t = swap32_S (t); - - if (wpa->nonce_compare < 0) + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { - pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); - pke[16] = (pke[16] & ~0xffffff00) | (t << 8); - } - else - { - pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); - pke[24] = (pke[24] & ~0xffffff00) | (t << 8); - } + u32 t = to; - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; + t -= nonce_error_corrections / 2; + t += nonce_error_correction; - w0[0] = out[0]; - w0[1] = out[1]; - w0[2] = out[2]; - w0[3] = out[3]; - w1[0] = out[4]; - w1[1] = out[5]; - w1[2] = out[6]; - w1[3] = out[7]; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - u32 keymic[4]; - - keymic[0] = 0; - keymic[1] = 0; - keymic[2] = 0; - keymic[3] = 0; - - sha1_hmac_ctx_t ctx1; - - sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); - - ctx1.ipad.w0[0] = pke[ 0]; - ctx1.ipad.w0[1] = pke[ 1]; - ctx1.ipad.w0[2] = pke[ 2]; - ctx1.ipad.w0[3] = pke[ 3]; - ctx1.ipad.w1[0] = pke[ 4]; - ctx1.ipad.w1[1] = pke[ 5]; - ctx1.ipad.w1[2] = pke[ 6]; - ctx1.ipad.w1[3] = pke[ 7]; - ctx1.ipad.w2[0] = pke[ 8]; - ctx1.ipad.w2[1] = pke[ 9]; - ctx1.ipad.w2[2] = pke[10]; - ctx1.ipad.w2[3] = pke[11]; - ctx1.ipad.w3[0] = pke[12]; - ctx1.ipad.w3[1] = pke[13]; - ctx1.ipad.w3[2] = pke[14]; - ctx1.ipad.w3[3] = pke[15]; - - sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); - - ctx1.ipad.w0[0] = pke[16]; - ctx1.ipad.w0[1] = pke[17]; - ctx1.ipad.w0[2] = pke[18]; - ctx1.ipad.w0[3] = pke[19]; - ctx1.ipad.w1[0] = pke[20]; - ctx1.ipad.w1[1] = pke[21]; - ctx1.ipad.w1[2] = pke[22]; - ctx1.ipad.w1[3] = pke[23]; - ctx1.ipad.w2[0] = pke[24]; - ctx1.ipad.w2[1] = pke[25]; - ctx1.ipad.w2[2] = pke[26]; - ctx1.ipad.w2[3] = pke[27]; - ctx1.ipad.w3[0] = pke[28]; - ctx1.ipad.w3[1] = pke[29]; - ctx1.ipad.w3[2] = pke[30]; - ctx1.ipad.w3[3] = pke[31]; - - ctx1.ipad.len += 100; - - sha1_hmac_final (&ctx1); - - u32 digest[4]; - - digest[0] = ctx1.opad.h[0]; - digest[1] = ctx1.opad.h[1]; - digest[2] = ctx1.opad.h[2]; - digest[3] = ctx1.opad.h[3]; - - u32 t0[4]; - u32 t1[4]; - u32 t2[4]; - u32 t3[4]; - - t0[0] = swap32_S (digest[0]); - t0[1] = swap32_S (digest[1]); - t0[2] = swap32_S (digest[2]); - t0[3] = swap32_S (digest[3]); - t1[0] = 0; - t1[1] = 0; - t1[2] = 0; - t1[3] = 0; - t2[0] = 0; - t2[1] = 0; - t2[2] = 0; - t2[3] = 0; - t3[0] = 0; - t3[1] = 0; - t3[2] = 0; - t3[3] = 0; - - md5_hmac_ctx_t ctx2; - - md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); - - md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); - - md5_hmac_final (&ctx2); - - keymic[0] = ctx2.opad.h[0]; - keymic[1] = ctx2.opad.h[1]; - keymic[2] = ctx2.opad.h[2]; - keymic[3] = ctx2.opad.h[3]; - - /** - * final compare - */ - - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) - { - if (atomic_inc (&hashes_shown[digest_cur]) == 0) + if (wpa->nonce_compare < 0) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha1_hmac_ctx_t ctx1; + + sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + ctx1.ipad.w0[0] = pke[ 0]; + ctx1.ipad.w0[1] = pke[ 1]; + ctx1.ipad.w0[2] = pke[ 2]; + ctx1.ipad.w0[3] = pke[ 3]; + ctx1.ipad.w1[0] = pke[ 4]; + ctx1.ipad.w1[1] = pke[ 5]; + ctx1.ipad.w1[2] = pke[ 6]; + ctx1.ipad.w1[3] = pke[ 7]; + ctx1.ipad.w2[0] = pke[ 8]; + ctx1.ipad.w2[1] = pke[ 9]; + ctx1.ipad.w2[2] = pke[10]; + ctx1.ipad.w2[3] = pke[11]; + ctx1.ipad.w3[0] = pke[12]; + ctx1.ipad.w3[1] = pke[13]; + ctx1.ipad.w3[2] = pke[14]; + ctx1.ipad.w3[3] = pke[15]; + + sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); + + ctx1.ipad.w0[0] = pke[16]; + ctx1.ipad.w0[1] = pke[17]; + ctx1.ipad.w0[2] = pke[18]; + ctx1.ipad.w0[3] = pke[19]; + ctx1.ipad.w1[0] = pke[20]; + ctx1.ipad.w1[1] = pke[21]; + ctx1.ipad.w1[2] = pke[22]; + ctx1.ipad.w1[3] = pke[23]; + ctx1.ipad.w2[0] = pke[24]; + ctx1.ipad.w2[1] = pke[25]; + ctx1.ipad.w2[2] = pke[26]; + ctx1.ipad.w2[3] = pke[27]; + ctx1.ipad.w3[0] = pke[28]; + ctx1.ipad.w3[1] = pke[29]; + ctx1.ipad.w3[2] = pke[30]; + ctx1.ipad.w3[3] = pke[31]; + + ctx1.ipad.len += 100; + + sha1_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = ctx1.opad.h[0]; + digest[1] = ctx1.opad.h[1]; + digest[2] = ctx1.opad.h[2]; + digest[3] = ctx1.opad.h[3]; + + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = swap32_S (digest[0]); + t0[1] = swap32_S (digest[1]); + t0[2] = swap32_S (digest[2]); + t0[3] = swap32_S (digest[3]); + t1[0] = 0; + t1[1] = 0; + t1[2] = 0; + t1[3] = 0; + t2[0] = 0; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = 0; + + md5_hmac_ctx_t ctx2; + + md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); + + md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + + md5_hmac_final (&ctx2); + + keymic[0] = ctx2.opad.h[0]; + keymic[1] = ctx2.opad.h[1]; + keymic[2] = ctx2.opad.h[2]; + keymic[3] = ctx2.opad.h[3]; + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } + } + } + } + + if (wpa->detected_be == 1) + { + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + { + u32 t = to; + + t = swap32_S (t); + + t -= nonce_error_corrections / 2; + t += nonce_error_correction; + + t = swap32_S (t); + + if (wpa->nonce_compare < 0) + { + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha1_hmac_ctx_t ctx1; + + sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + ctx1.ipad.w0[0] = pke[ 0]; + ctx1.ipad.w0[1] = pke[ 1]; + ctx1.ipad.w0[2] = pke[ 2]; + ctx1.ipad.w0[3] = pke[ 3]; + ctx1.ipad.w1[0] = pke[ 4]; + ctx1.ipad.w1[1] = pke[ 5]; + ctx1.ipad.w1[2] = pke[ 6]; + ctx1.ipad.w1[3] = pke[ 7]; + ctx1.ipad.w2[0] = pke[ 8]; + ctx1.ipad.w2[1] = pke[ 9]; + ctx1.ipad.w2[2] = pke[10]; + ctx1.ipad.w2[3] = pke[11]; + ctx1.ipad.w3[0] = pke[12]; + ctx1.ipad.w3[1] = pke[13]; + ctx1.ipad.w3[2] = pke[14]; + ctx1.ipad.w3[3] = pke[15]; + + sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); + + ctx1.ipad.w0[0] = pke[16]; + ctx1.ipad.w0[1] = pke[17]; + ctx1.ipad.w0[2] = pke[18]; + ctx1.ipad.w0[3] = pke[19]; + ctx1.ipad.w1[0] = pke[20]; + ctx1.ipad.w1[1] = pke[21]; + ctx1.ipad.w1[2] = pke[22]; + ctx1.ipad.w1[3] = pke[23]; + ctx1.ipad.w2[0] = pke[24]; + ctx1.ipad.w2[1] = pke[25]; + ctx1.ipad.w2[2] = pke[26]; + ctx1.ipad.w2[3] = pke[27]; + ctx1.ipad.w3[0] = pke[28]; + ctx1.ipad.w3[1] = pke[29]; + ctx1.ipad.w3[2] = pke[30]; + ctx1.ipad.w3[3] = pke[31]; + + ctx1.ipad.len += 100; + + sha1_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = ctx1.opad.h[0]; + digest[1] = ctx1.opad.h[1]; + digest[2] = ctx1.opad.h[2]; + digest[3] = ctx1.opad.h[3]; + + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = swap32_S (digest[0]); + t0[1] = swap32_S (digest[1]); + t0[2] = swap32_S (digest[2]); + t0[3] = swap32_S (digest[3]); + t1[0] = 0; + t1[1] = 0; + t1[2] = 0; + t1[3] = 0; + t2[0] = 0; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = 0; + + md5_hmac_ctx_t ctx2; + + md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); + + md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + + md5_hmac_final (&ctx2); + + keymic[0] = ctx2.opad.h[0]; + keymic[1] = ctx2.opad.h[1]; + keymic[2] = ctx2.opad.h[2]; + keymic[3] = ctx2.opad.h[3]; + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } } } } @@ -549,155 +704,310 @@ __kernel void m02500_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 nonce_error_corrections = wpa->nonce_error_corrections; - for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + if (wpa->detected_le == 1) { - u32 t = to; - - t = swap32_S (t); - - t -= nonce_error_corrections / 2; - t += nonce_error_correction; - - t = swap32_S (t); - - if (wpa->nonce_compare < 0) + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { - pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); - pke[16] = (pke[16] & ~0xffffff00) | (t << 8); - } - else - { - pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); - pke[24] = (pke[24] & ~0xffffff00) | (t << 8); - } + u32 t = to; - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; + t -= nonce_error_corrections / 2; + t += nonce_error_correction; - w0[0] = out[0]; - w0[1] = out[1]; - w0[2] = out[2]; - w0[3] = out[3]; - w1[0] = out[4]; - w1[1] = out[5]; - w1[2] = out[6]; - w1[3] = out[7]; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - u32 keymic[4]; - - keymic[0] = 0; - keymic[1] = 0; - keymic[2] = 0; - keymic[3] = 0; - - sha1_hmac_ctx_t ctx1; - - sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); - - ctx1.ipad.w0[0] = pke[ 0]; - ctx1.ipad.w0[1] = pke[ 1]; - ctx1.ipad.w0[2] = pke[ 2]; - ctx1.ipad.w0[3] = pke[ 3]; - ctx1.ipad.w1[0] = pke[ 4]; - ctx1.ipad.w1[1] = pke[ 5]; - ctx1.ipad.w1[2] = pke[ 6]; - ctx1.ipad.w1[3] = pke[ 7]; - ctx1.ipad.w2[0] = pke[ 8]; - ctx1.ipad.w2[1] = pke[ 9]; - ctx1.ipad.w2[2] = pke[10]; - ctx1.ipad.w2[3] = pke[11]; - ctx1.ipad.w3[0] = pke[12]; - ctx1.ipad.w3[1] = pke[13]; - ctx1.ipad.w3[2] = pke[14]; - ctx1.ipad.w3[3] = pke[15]; - - sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); - - ctx1.ipad.w0[0] = pke[16]; - ctx1.ipad.w0[1] = pke[17]; - ctx1.ipad.w0[2] = pke[18]; - ctx1.ipad.w0[3] = pke[19]; - ctx1.ipad.w1[0] = pke[20]; - ctx1.ipad.w1[1] = pke[21]; - ctx1.ipad.w1[2] = pke[22]; - ctx1.ipad.w1[3] = pke[23]; - ctx1.ipad.w2[0] = pke[24]; - ctx1.ipad.w2[1] = pke[25]; - ctx1.ipad.w2[2] = pke[26]; - ctx1.ipad.w2[3] = pke[27]; - ctx1.ipad.w3[0] = pke[28]; - ctx1.ipad.w3[1] = pke[29]; - ctx1.ipad.w3[2] = pke[30]; - ctx1.ipad.w3[3] = pke[31]; - - ctx1.ipad.len += 100; - - sha1_hmac_final (&ctx1); - - u32 digest[4]; - - digest[0] = ctx1.opad.h[0]; - digest[1] = ctx1.opad.h[1]; - digest[2] = ctx1.opad.h[2]; - digest[3] = ctx1.opad.h[3]; - - u32 t0[4]; - u32 t1[4]; - u32 t2[4]; - u32 t3[4]; - - t0[0] = digest[0]; - t0[1] = digest[1]; - t0[2] = digest[2]; - t0[3] = digest[3]; - t1[0] = 0; - t1[1] = 0; - t1[2] = 0; - t1[3] = 0; - t2[0] = 0; - t2[1] = 0; - t2[2] = 0; - t2[3] = 0; - t3[0] = 0; - t3[1] = 0; - t3[2] = 0; - t3[3] = 0; - - sha1_hmac_ctx_t ctx2; - - sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); - - sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); - - sha1_hmac_final (&ctx2); - - keymic[0] = ctx2.opad.h[0]; - keymic[1] = ctx2.opad.h[1]; - keymic[2] = ctx2.opad.h[2]; - keymic[3] = ctx2.opad.h[3]; - - /** - * final compare - */ - - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) - { - if (atomic_inc (&hashes_shown[digest_cur]) == 0) + if (wpa->nonce_compare < 0) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha1_hmac_ctx_t ctx1; + + sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + ctx1.ipad.w0[0] = pke[ 0]; + ctx1.ipad.w0[1] = pke[ 1]; + ctx1.ipad.w0[2] = pke[ 2]; + ctx1.ipad.w0[3] = pke[ 3]; + ctx1.ipad.w1[0] = pke[ 4]; + ctx1.ipad.w1[1] = pke[ 5]; + ctx1.ipad.w1[2] = pke[ 6]; + ctx1.ipad.w1[3] = pke[ 7]; + ctx1.ipad.w2[0] = pke[ 8]; + ctx1.ipad.w2[1] = pke[ 9]; + ctx1.ipad.w2[2] = pke[10]; + ctx1.ipad.w2[3] = pke[11]; + ctx1.ipad.w3[0] = pke[12]; + ctx1.ipad.w3[1] = pke[13]; + ctx1.ipad.w3[2] = pke[14]; + ctx1.ipad.w3[3] = pke[15]; + + sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); + + ctx1.ipad.w0[0] = pke[16]; + ctx1.ipad.w0[1] = pke[17]; + ctx1.ipad.w0[2] = pke[18]; + ctx1.ipad.w0[3] = pke[19]; + ctx1.ipad.w1[0] = pke[20]; + ctx1.ipad.w1[1] = pke[21]; + ctx1.ipad.w1[2] = pke[22]; + ctx1.ipad.w1[3] = pke[23]; + ctx1.ipad.w2[0] = pke[24]; + ctx1.ipad.w2[1] = pke[25]; + ctx1.ipad.w2[2] = pke[26]; + ctx1.ipad.w2[3] = pke[27]; + ctx1.ipad.w3[0] = pke[28]; + ctx1.ipad.w3[1] = pke[29]; + ctx1.ipad.w3[2] = pke[30]; + ctx1.ipad.w3[3] = pke[31]; + + ctx1.ipad.len += 100; + + sha1_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = ctx1.opad.h[0]; + digest[1] = ctx1.opad.h[1]; + digest[2] = ctx1.opad.h[2]; + digest[3] = ctx1.opad.h[3]; + + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = 0; + t1[1] = 0; + t1[2] = 0; + t1[3] = 0; + t2[0] = 0; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = 0; + + sha1_hmac_ctx_t ctx2; + + sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); + + sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + + sha1_hmac_final (&ctx2); + + keymic[0] = ctx2.opad.h[0]; + keymic[1] = ctx2.opad.h[1]; + keymic[2] = ctx2.opad.h[2]; + keymic[3] = ctx2.opad.h[3]; + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } + } + } + } + + if (wpa->detected_be == 1) + { + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + { + u32 t = to; + + t = swap32_S (t); + + t -= nonce_error_corrections / 2; + t += nonce_error_correction; + + t = swap32_S (t); + + if (wpa->nonce_compare < 0) + { + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha1_hmac_ctx_t ctx1; + + sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + ctx1.ipad.w0[0] = pke[ 0]; + ctx1.ipad.w0[1] = pke[ 1]; + ctx1.ipad.w0[2] = pke[ 2]; + ctx1.ipad.w0[3] = pke[ 3]; + ctx1.ipad.w1[0] = pke[ 4]; + ctx1.ipad.w1[1] = pke[ 5]; + ctx1.ipad.w1[2] = pke[ 6]; + ctx1.ipad.w1[3] = pke[ 7]; + ctx1.ipad.w2[0] = pke[ 8]; + ctx1.ipad.w2[1] = pke[ 9]; + ctx1.ipad.w2[2] = pke[10]; + ctx1.ipad.w2[3] = pke[11]; + ctx1.ipad.w3[0] = pke[12]; + ctx1.ipad.w3[1] = pke[13]; + ctx1.ipad.w3[2] = pke[14]; + ctx1.ipad.w3[3] = pke[15]; + + sha1_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); + + ctx1.ipad.w0[0] = pke[16]; + ctx1.ipad.w0[1] = pke[17]; + ctx1.ipad.w0[2] = pke[18]; + ctx1.ipad.w0[3] = pke[19]; + ctx1.ipad.w1[0] = pke[20]; + ctx1.ipad.w1[1] = pke[21]; + ctx1.ipad.w1[2] = pke[22]; + ctx1.ipad.w1[3] = pke[23]; + ctx1.ipad.w2[0] = pke[24]; + ctx1.ipad.w2[1] = pke[25]; + ctx1.ipad.w2[2] = pke[26]; + ctx1.ipad.w2[3] = pke[27]; + ctx1.ipad.w3[0] = pke[28]; + ctx1.ipad.w3[1] = pke[29]; + ctx1.ipad.w3[2] = pke[30]; + ctx1.ipad.w3[3] = pke[31]; + + ctx1.ipad.len += 100; + + sha1_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = ctx1.opad.h[0]; + digest[1] = ctx1.opad.h[1]; + digest[2] = ctx1.opad.h[2]; + digest[3] = ctx1.opad.h[3]; + + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = 0; + t1[1] = 0; + t1[2] = 0; + t1[3] = 0; + t2[0] = 0; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = 0; + + sha1_hmac_ctx_t ctx2; + + sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); + + sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + + sha1_hmac_final (&ctx2); + + keymic[0] = ctx2.opad.h[0]; + keymic[1] = ctx2.opad.h[1]; + keymic[2] = ctx2.opad.h[2]; + keymic[3] = ctx2.opad.h[3]; + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } } } } @@ -829,186 +1139,372 @@ __kernel void m02500_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 nonce_error_corrections = wpa->nonce_error_corrections; - for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + if (wpa->detected_le == 1) { - u32 t = to; - - t = swap32_S (t); - - t -= nonce_error_corrections / 2; - t += nonce_error_correction; - - t = swap32_S (t); - - if (wpa->nonce_compare < 0) + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { - pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); - pke[16] = (pke[16] & ~0xffffff00) | (t << 8); - } - else - { - pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); - pke[24] = (pke[24] & ~0xffffff00) | (t << 8); - } + u32 t = to; - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; + t -= nonce_error_corrections / 2; + t += nonce_error_correction; - w0[0] = out[0]; - w0[1] = out[1]; - w0[2] = out[2]; - w0[3] = out[3]; - w1[0] = out[4]; - w1[1] = out[5]; - w1[2] = out[6]; - w1[3] = out[7]; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - u32 keymic[4]; - - keymic[0] = 0; - keymic[1] = 0; - keymic[2] = 0; - keymic[3] = 0; - - sha256_hmac_ctx_t ctx1; - - sha256_hmac_init_64 (&ctx1, w0, w1, w2, w3); - - ctx1.ipad.w0[0] = pke[ 0]; - ctx1.ipad.w0[1] = pke[ 1]; - ctx1.ipad.w0[2] = pke[ 2]; - ctx1.ipad.w0[3] = pke[ 3]; - ctx1.ipad.w1[0] = pke[ 4]; - ctx1.ipad.w1[1] = pke[ 5]; - ctx1.ipad.w1[2] = pke[ 6]; - ctx1.ipad.w1[3] = pke[ 7]; - ctx1.ipad.w2[0] = pke[ 8]; - ctx1.ipad.w2[1] = pke[ 9]; - ctx1.ipad.w2[2] = pke[10]; - ctx1.ipad.w2[3] = pke[11]; - ctx1.ipad.w3[0] = pke[12]; - ctx1.ipad.w3[1] = pke[13]; - ctx1.ipad.w3[2] = pke[14]; - ctx1.ipad.w3[3] = pke[15]; - - sha256_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); - - ctx1.ipad.w0[0] = pke[16]; - ctx1.ipad.w0[1] = pke[17]; - ctx1.ipad.w0[2] = pke[18]; - ctx1.ipad.w0[3] = pke[19]; - ctx1.ipad.w1[0] = pke[20]; - ctx1.ipad.w1[1] = pke[21]; - ctx1.ipad.w1[2] = pke[22]; - ctx1.ipad.w1[3] = pke[23]; - ctx1.ipad.w2[0] = pke[24]; - ctx1.ipad.w2[1] = pke[25]; - ctx1.ipad.w2[2] = pke[26]; - ctx1.ipad.w2[3] = pke[27]; - ctx1.ipad.w3[0] = pke[28]; - ctx1.ipad.w3[1] = pke[29]; - ctx1.ipad.w3[2] = pke[30]; - ctx1.ipad.w3[3] = pke[31]; - - ctx1.ipad.len += 102; - - sha256_hmac_final (&ctx1); - - u32 digest[4]; - - digest[0] = swap32_S (ctx1.opad.h[0]); - digest[1] = swap32_S (ctx1.opad.h[1]); - digest[2] = swap32_S (ctx1.opad.h[2]); - digest[3] = swap32_S (ctx1.opad.h[3]); - - // AES CMAC - - u32 ks[44]; - - aes128_set_encrypt_key (ks, digest, s_te0, s_te1, s_te2, s_te3, s_te4); - - u32 m[4]; - - m[0] = 0; - m[1] = 0; - m[2] = 0; - m[3] = 0; - - u32 iv[4]; - - iv[0] = 0; - iv[1] = 0; - iv[2] = 0; - iv[3] = 0; - - int eapol_left; - int eapol_idx; - - for (eapol_left = wpa->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]; - - 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]; - - u32 k[4]; - - k[0] = 0; - k[1] = 0; - k[2] = 0; - k[3] = 0; - - aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); - - make_kn (k); - - if (eapol_left < 16) - { - make_kn (k); - } - - m[0] ^= k[0]; - m[1] ^= k[1]; - m[2] ^= k[2]; - m[3] ^= k[3]; - - m[0] ^= iv[0]; - m[1] ^= iv[1]; - m[2] ^= iv[2]; - m[3] ^= iv[3]; - - aes128_encrypt (ks, m, keymic, s_te0, s_te1, s_te2, s_te3, s_te4); - - /** - * final compare - */ - - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) - { - if (atomic_inc (&hashes_shown[digest_cur]) == 0) + if (wpa->nonce_compare < 0) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha256_hmac_ctx_t ctx1; + + sha256_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + ctx1.ipad.w0[0] = pke[ 0]; + ctx1.ipad.w0[1] = pke[ 1]; + ctx1.ipad.w0[2] = pke[ 2]; + ctx1.ipad.w0[3] = pke[ 3]; + ctx1.ipad.w1[0] = pke[ 4]; + ctx1.ipad.w1[1] = pke[ 5]; + ctx1.ipad.w1[2] = pke[ 6]; + ctx1.ipad.w1[3] = pke[ 7]; + ctx1.ipad.w2[0] = pke[ 8]; + ctx1.ipad.w2[1] = pke[ 9]; + ctx1.ipad.w2[2] = pke[10]; + ctx1.ipad.w2[3] = pke[11]; + ctx1.ipad.w3[0] = pke[12]; + ctx1.ipad.w3[1] = pke[13]; + ctx1.ipad.w3[2] = pke[14]; + ctx1.ipad.w3[3] = pke[15]; + + sha256_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); + + ctx1.ipad.w0[0] = pke[16]; + ctx1.ipad.w0[1] = pke[17]; + ctx1.ipad.w0[2] = pke[18]; + ctx1.ipad.w0[3] = pke[19]; + ctx1.ipad.w1[0] = pke[20]; + ctx1.ipad.w1[1] = pke[21]; + ctx1.ipad.w1[2] = pke[22]; + ctx1.ipad.w1[3] = pke[23]; + ctx1.ipad.w2[0] = pke[24]; + ctx1.ipad.w2[1] = pke[25]; + ctx1.ipad.w2[2] = pke[26]; + ctx1.ipad.w2[3] = pke[27]; + ctx1.ipad.w3[0] = pke[28]; + ctx1.ipad.w3[1] = pke[29]; + ctx1.ipad.w3[2] = pke[30]; + ctx1.ipad.w3[3] = pke[31]; + + ctx1.ipad.len += 102; + + sha256_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = swap32_S (ctx1.opad.h[0]); + digest[1] = swap32_S (ctx1.opad.h[1]); + digest[2] = swap32_S (ctx1.opad.h[2]); + digest[3] = swap32_S (ctx1.opad.h[3]); + + // AES CMAC + + u32 ks[44]; + + aes128_set_encrypt_key (ks, digest, s_te0, s_te1, s_te2, s_te3, s_te4); + + u32 m[4]; + + m[0] = 0; + m[1] = 0; + m[2] = 0; + m[3] = 0; + + u32 iv[4]; + + iv[0] = 0; + iv[1] = 0; + iv[2] = 0; + iv[3] = 0; + + int eapol_left; + int eapol_idx; + + for (eapol_left = wpa->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]; + + 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]; + + u32 k[4]; + + k[0] = 0; + k[1] = 0; + k[2] = 0; + k[3] = 0; + + aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); + + make_kn (k); + + if (eapol_left < 16) + { + make_kn (k); + } + + m[0] ^= k[0]; + m[1] ^= k[1]; + m[2] ^= k[2]; + m[3] ^= k[3]; + + m[0] ^= iv[0]; + m[1] ^= iv[1]; + m[2] ^= iv[2]; + m[3] ^= iv[3]; + + aes128_encrypt (ks, m, keymic, s_te0, s_te1, s_te2, s_te3, s_te4); + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } + } + } + } + + if (wpa->detected_be == 1) + { + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + { + u32 t = to; + + t = swap32_S (t); + + t -= nonce_error_corrections / 2; + t += nonce_error_correction; + + t = swap32_S (t); + + if (wpa->nonce_compare < 0) + { + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha256_hmac_ctx_t ctx1; + + sha256_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + ctx1.ipad.w0[0] = pke[ 0]; + ctx1.ipad.w0[1] = pke[ 1]; + ctx1.ipad.w0[2] = pke[ 2]; + ctx1.ipad.w0[3] = pke[ 3]; + ctx1.ipad.w1[0] = pke[ 4]; + ctx1.ipad.w1[1] = pke[ 5]; + ctx1.ipad.w1[2] = pke[ 6]; + ctx1.ipad.w1[3] = pke[ 7]; + ctx1.ipad.w2[0] = pke[ 8]; + ctx1.ipad.w2[1] = pke[ 9]; + ctx1.ipad.w2[2] = pke[10]; + ctx1.ipad.w2[3] = pke[11]; + ctx1.ipad.w3[0] = pke[12]; + ctx1.ipad.w3[1] = pke[13]; + ctx1.ipad.w3[2] = pke[14]; + ctx1.ipad.w3[3] = pke[15]; + + sha256_transform (ctx1.ipad.w0, ctx1.ipad.w1, ctx1.ipad.w2, ctx1.ipad.w3, ctx1.ipad.h); + + ctx1.ipad.w0[0] = pke[16]; + ctx1.ipad.w0[1] = pke[17]; + ctx1.ipad.w0[2] = pke[18]; + ctx1.ipad.w0[3] = pke[19]; + ctx1.ipad.w1[0] = pke[20]; + ctx1.ipad.w1[1] = pke[21]; + ctx1.ipad.w1[2] = pke[22]; + ctx1.ipad.w1[3] = pke[23]; + ctx1.ipad.w2[0] = pke[24]; + ctx1.ipad.w2[1] = pke[25]; + ctx1.ipad.w2[2] = pke[26]; + ctx1.ipad.w2[3] = pke[27]; + ctx1.ipad.w3[0] = pke[28]; + ctx1.ipad.w3[1] = pke[29]; + ctx1.ipad.w3[2] = pke[30]; + ctx1.ipad.w3[3] = pke[31]; + + ctx1.ipad.len += 102; + + sha256_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = swap32_S (ctx1.opad.h[0]); + digest[1] = swap32_S (ctx1.opad.h[1]); + digest[2] = swap32_S (ctx1.opad.h[2]); + digest[3] = swap32_S (ctx1.opad.h[3]); + + // AES CMAC + + u32 ks[44]; + + aes128_set_encrypt_key (ks, digest, s_te0, s_te1, s_te2, s_te3, s_te4); + + u32 m[4]; + + m[0] = 0; + m[1] = 0; + m[2] = 0; + m[3] = 0; + + u32 iv[4]; + + iv[0] = 0; + iv[1] = 0; + iv[2] = 0; + iv[3] = 0; + + int eapol_left; + int eapol_idx; + + for (eapol_left = wpa->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]; + + 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]; + + u32 k[4]; + + k[0] = 0; + k[1] = 0; + k[2] = 0; + k[3] = 0; + + aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); + + make_kn (k); + + if (eapol_left < 16) + { + make_kn (k); + } + + m[0] ^= k[0]; + m[1] ^= k[1]; + m[2] ^= k[2]; + m[3] ^= k[3]; + + m[0] ^= iv[0]; + m[1] ^= iv[1]; + m[2] ^= iv[2]; + m[3] ^= iv[3]; + + aes128_encrypt (ks, m, keymic, s_te0, s_te1, s_te2, s_te3, s_te4); + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } } } } diff --git a/OpenCL/m02501.cl b/OpenCL/m02501.cl index 032cdccc5..eada1a0ad 100644 --- a/OpenCL/m02501.cl +++ b/OpenCL/m02501.cl @@ -196,119 +196,238 @@ __kernel void m02501_aux1 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 nonce_error_corrections = wpa->nonce_error_corrections; - for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + if (wpa->detected_le == 1) { - u32 t = to; - - t = swap32_S (t); - - t -= nonce_error_corrections / 2; - t += nonce_error_correction; - - t = swap32_S (t); - - if (wpa->nonce_compare < 0) + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { - pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); - pke[16] = (pke[16] & ~0xffffff00) | (t << 8); - } - else - { - pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); - pke[24] = (pke[24] & ~0xffffff00) | (t << 8); - } + u32 t = to; - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; + t -= nonce_error_corrections / 2; + t += nonce_error_correction; - w0[0] = out[0]; - w0[1] = out[1]; - w0[2] = out[2]; - w0[3] = out[3]; - w1[0] = out[4]; - w1[1] = out[5]; - w1[2] = out[6]; - w1[3] = out[7]; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - u32 keymic[4]; - - keymic[0] = 0; - keymic[1] = 0; - keymic[2] = 0; - keymic[3] = 0; - - sha1_hmac_ctx_t ctx1; - - sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); - - sha1_hmac_update (&ctx1, pke, 100); - - sha1_hmac_final (&ctx1); - - u32 digest[4]; - - digest[0] = ctx1.opad.h[0]; - digest[1] = ctx1.opad.h[1]; - digest[2] = ctx1.opad.h[2]; - digest[3] = ctx1.opad.h[3]; - - u32 t0[4]; - u32 t1[4]; - u32 t2[4]; - u32 t3[4]; - - t0[0] = swap32_S (digest[0]); - t0[1] = swap32_S (digest[1]); - t0[2] = swap32_S (digest[2]); - t0[3] = swap32_S (digest[3]); - t1[0] = 0; - t1[1] = 0; - t1[2] = 0; - t1[3] = 0; - t2[0] = 0; - t2[1] = 0; - t2[2] = 0; - t2[3] = 0; - t3[0] = 0; - t3[1] = 0; - t3[2] = 0; - t3[3] = 0; - - md5_hmac_ctx_t ctx2; - - md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); - - md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); - - md5_hmac_final (&ctx2); - - keymic[0] = ctx2.opad.h[0]; - keymic[1] = ctx2.opad.h[1]; - keymic[2] = ctx2.opad.h[2]; - keymic[3] = ctx2.opad.h[3]; - - /** - * final compare - */ - - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) - { - if (atomic_inc (&hashes_shown[digest_cur]) == 0) + if (wpa->nonce_compare < 0) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha1_hmac_ctx_t ctx1; + + sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + sha1_hmac_update (&ctx1, pke, 100); + + sha1_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = ctx1.opad.h[0]; + digest[1] = ctx1.opad.h[1]; + digest[2] = ctx1.opad.h[2]; + digest[3] = ctx1.opad.h[3]; + + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = swap32_S (digest[0]); + t0[1] = swap32_S (digest[1]); + t0[2] = swap32_S (digest[2]); + t0[3] = swap32_S (digest[3]); + t1[0] = 0; + t1[1] = 0; + t1[2] = 0; + t1[3] = 0; + t2[0] = 0; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = 0; + + md5_hmac_ctx_t ctx2; + + md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); + + md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + + md5_hmac_final (&ctx2); + + keymic[0] = ctx2.opad.h[0]; + keymic[1] = ctx2.opad.h[1]; + keymic[2] = ctx2.opad.h[2]; + keymic[3] = ctx2.opad.h[3]; + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } + } + } + } + + if (wpa->detected_be == 1) + { + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + { + u32 t = to; + + t = swap32_S (t); + + t -= nonce_error_corrections / 2; + t += nonce_error_correction; + + t = swap32_S (t); + + if (wpa->nonce_compare < 0) + { + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha1_hmac_ctx_t ctx1; + + sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + sha1_hmac_update (&ctx1, pke, 100); + + sha1_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = ctx1.opad.h[0]; + digest[1] = ctx1.opad.h[1]; + digest[2] = ctx1.opad.h[2]; + digest[3] = ctx1.opad.h[3]; + + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = swap32_S (digest[0]); + t0[1] = swap32_S (digest[1]); + t0[2] = swap32_S (digest[2]); + t0[3] = swap32_S (digest[3]); + t1[0] = 0; + t1[1] = 0; + t1[2] = 0; + t1[3] = 0; + t2[0] = 0; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = 0; + + md5_hmac_ctx_t ctx2; + + md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); + + md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + + md5_hmac_final (&ctx2); + + keymic[0] = ctx2.opad.h[0]; + keymic[1] = ctx2.opad.h[1]; + keymic[2] = ctx2.opad.h[2]; + keymic[3] = ctx2.opad.h[3]; + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } } } } @@ -388,119 +507,238 @@ __kernel void m02501_aux2 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 nonce_error_corrections = wpa->nonce_error_corrections; - for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + if (wpa->detected_le == 1) { - u32 t = to; - - t = swap32_S (t); - - t -= nonce_error_corrections / 2; - t += nonce_error_correction; - - t = swap32_S (t); - - if (wpa->nonce_compare < 0) + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { - pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); - pke[16] = (pke[16] & ~0xffffff00) | (t << 8); - } - else - { - pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); - pke[24] = (pke[24] & ~0xffffff00) | (t << 8); - } + u32 t = to; - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; + t -= nonce_error_corrections / 2; + t += nonce_error_correction; - w0[0] = out[0]; - w0[1] = out[1]; - w0[2] = out[2]; - w0[3] = out[3]; - w1[0] = out[4]; - w1[1] = out[5]; - w1[2] = out[6]; - w1[3] = out[7]; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - u32 keymic[4]; - - keymic[0] = 0; - keymic[1] = 0; - keymic[2] = 0; - keymic[3] = 0; - - sha1_hmac_ctx_t ctx1; - - sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); - - sha1_hmac_update (&ctx1, pke, 100); - - sha1_hmac_final (&ctx1); - - u32 digest[4]; - - digest[0] = ctx1.opad.h[0]; - digest[1] = ctx1.opad.h[1]; - digest[2] = ctx1.opad.h[2]; - digest[3] = ctx1.opad.h[3]; - - u32 t0[4]; - u32 t1[4]; - u32 t2[4]; - u32 t3[4]; - - t0[0] = digest[0]; - t0[1] = digest[1]; - t0[2] = digest[2]; - t0[3] = digest[3]; - t1[0] = 0; - t1[1] = 0; - t1[2] = 0; - t1[3] = 0; - t2[0] = 0; - t2[1] = 0; - t2[2] = 0; - t2[3] = 0; - t3[0] = 0; - t3[1] = 0; - t3[2] = 0; - t3[3] = 0; - - sha1_hmac_ctx_t ctx2; - - sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); - - sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); - - sha1_hmac_final (&ctx2); - - keymic[0] = ctx2.opad.h[0]; - keymic[1] = ctx2.opad.h[1]; - keymic[2] = ctx2.opad.h[2]; - keymic[3] = ctx2.opad.h[3]; - - /** - * final compare - */ - - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) - { - if (atomic_inc (&hashes_shown[digest_cur]) == 0) + if (wpa->nonce_compare < 0) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha1_hmac_ctx_t ctx1; + + sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + sha1_hmac_update (&ctx1, pke, 100); + + sha1_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = ctx1.opad.h[0]; + digest[1] = ctx1.opad.h[1]; + digest[2] = ctx1.opad.h[2]; + digest[3] = ctx1.opad.h[3]; + + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = 0; + t1[1] = 0; + t1[2] = 0; + t1[3] = 0; + t2[0] = 0; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = 0; + + sha1_hmac_ctx_t ctx2; + + sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); + + sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + + sha1_hmac_final (&ctx2); + + keymic[0] = ctx2.opad.h[0]; + keymic[1] = ctx2.opad.h[1]; + keymic[2] = ctx2.opad.h[2]; + keymic[3] = ctx2.opad.h[3]; + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } + } + } + } + + if (wpa->detected_be == 1) + { + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + { + u32 t = to; + + t = swap32_S (t); + + t -= nonce_error_corrections / 2; + t += nonce_error_correction; + + t = swap32_S (t); + + if (wpa->nonce_compare < 0) + { + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha1_hmac_ctx_t ctx1; + + sha1_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + sha1_hmac_update (&ctx1, pke, 100); + + sha1_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = ctx1.opad.h[0]; + digest[1] = ctx1.opad.h[1]; + digest[2] = ctx1.opad.h[2]; + digest[3] = ctx1.opad.h[3]; + + u32 t0[4]; + u32 t1[4]; + u32 t2[4]; + u32 t3[4]; + + t0[0] = digest[0]; + t0[1] = digest[1]; + t0[2] = digest[2]; + t0[3] = digest[3]; + t1[0] = 0; + t1[1] = 0; + t1[2] = 0; + t1[3] = 0; + t2[0] = 0; + t2[1] = 0; + t2[2] = 0; + t2[3] = 0; + t3[0] = 0; + t3[1] = 0; + t3[2] = 0; + t3[3] = 0; + + sha1_hmac_ctx_t ctx2; + + sha1_hmac_init_64 (&ctx2, t0, t1, t2, t3); + + sha1_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); + + sha1_hmac_final (&ctx2); + + keymic[0] = ctx2.opad.h[0]; + keymic[1] = ctx2.opad.h[1]; + keymic[2] = ctx2.opad.h[2]; + keymic[3] = ctx2.opad.h[3]; + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } } } } @@ -632,150 +870,300 @@ __kernel void m02501_aux3 (__global pw_t *pws, __global const kernel_rule_t *rul const u32 nonce_error_corrections = wpa->nonce_error_corrections; - for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + if (wpa->detected_le == 1) { - u32 t = to; - - t = swap32_S (t); - - t -= nonce_error_corrections / 2; - t += nonce_error_correction; - - t = swap32_S (t); - - if (wpa->nonce_compare < 0) + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) { - pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); - pke[16] = (pke[16] & ~0xffffff00) | (t << 8); - } - else - { - pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); - pke[24] = (pke[24] & ~0xffffff00) | (t << 8); - } + u32 t = to; - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; + t -= nonce_error_corrections / 2; + t += nonce_error_correction; - w0[0] = out[0]; - w0[1] = out[1]; - w0[2] = out[2]; - w0[3] = out[3]; - w1[0] = out[4]; - w1[1] = out[5]; - w1[2] = out[6]; - w1[3] = out[7]; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - u32 keymic[4]; - - keymic[0] = 0; - keymic[1] = 0; - keymic[2] = 0; - keymic[3] = 0; - - sha256_hmac_ctx_t ctx1; - - sha256_hmac_init_64 (&ctx1, w0, w1, w2, w3); - - sha256_hmac_update (&ctx1, pke, 102); - - sha256_hmac_final (&ctx1); - - u32 digest[4]; - - digest[0] = swap32_S (ctx1.opad.h[0]); - digest[1] = swap32_S (ctx1.opad.h[1]); - digest[2] = swap32_S (ctx1.opad.h[2]); - digest[3] = swap32_S (ctx1.opad.h[3]); - - // AES CMAC - - u32 ks[44]; - - aes128_set_encrypt_key (ks, digest, s_te0, s_te1, s_te2, s_te3, s_te4); - - u32 m[4]; - - m[0] = 0; - m[1] = 0; - m[2] = 0; - m[3] = 0; - - u32 iv[4]; - - iv[0] = 0; - iv[1] = 0; - iv[2] = 0; - iv[3] = 0; - - int eapol_left; - int eapol_idx; - - for (eapol_left = wpa->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]; - - 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]; - - u32 k[4]; - - k[0] = 0; - k[1] = 0; - k[2] = 0; - k[3] = 0; - - aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); - - make_kn (k); - - if (eapol_left < 16) - { - make_kn (k); - } - - m[0] ^= k[0]; - m[1] ^= k[1]; - m[2] ^= k[2]; - m[3] ^= k[3]; - - m[0] ^= iv[0]; - m[1] ^= iv[1]; - m[2] ^= iv[2]; - m[3] ^= iv[3]; - - aes128_encrypt (ks, m, keymic, s_te0, s_te1, s_te2, s_te3, s_te4); - - /** - * final compare - */ - - if ((keymic[0] == wpa->keymic[0]) - && (keymic[1] == wpa->keymic[1]) - && (keymic[2] == wpa->keymic[2]) - && (keymic[3] == wpa->keymic[3])) - { - if (atomic_inc (&hashes_shown[digest_cur]) == 0) + if (wpa->nonce_compare < 0) { - mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha256_hmac_ctx_t ctx1; + + sha256_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + sha256_hmac_update (&ctx1, pke, 102); + + sha256_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = swap32_S (ctx1.opad.h[0]); + digest[1] = swap32_S (ctx1.opad.h[1]); + digest[2] = swap32_S (ctx1.opad.h[2]); + digest[3] = swap32_S (ctx1.opad.h[3]); + + // AES CMAC + + u32 ks[44]; + + aes128_set_encrypt_key (ks, digest, s_te0, s_te1, s_te2, s_te3, s_te4); + + u32 m[4]; + + m[0] = 0; + m[1] = 0; + m[2] = 0; + m[3] = 0; + + u32 iv[4]; + + iv[0] = 0; + iv[1] = 0; + iv[2] = 0; + iv[3] = 0; + + int eapol_left; + int eapol_idx; + + for (eapol_left = wpa->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]; + + 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]; + + u32 k[4]; + + k[0] = 0; + k[1] = 0; + k[2] = 0; + k[3] = 0; + + aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); + + make_kn (k); + + if (eapol_left < 16) + { + make_kn (k); + } + + m[0] ^= k[0]; + m[1] ^= k[1]; + m[2] ^= k[2]; + m[3] ^= k[3]; + + m[0] ^= iv[0]; + m[1] ^= iv[1]; + m[2] ^= iv[2]; + m[3] ^= iv[3]; + + aes128_encrypt (ks, m, keymic, s_te0, s_te1, s_te2, s_te3, s_te4); + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } + } + } + } + + if (wpa->detected_be == 1) + { + for (u32 nonce_error_correction = 0; nonce_error_correction <= nonce_error_corrections; nonce_error_correction++) + { + u32 t = to; + + t = swap32_S (t); + + t -= nonce_error_corrections / 2; + t += nonce_error_correction; + + t = swap32_S (t); + + if (wpa->nonce_compare < 0) + { + pke[15] = (pke[15] & ~0x000000ff) | (t >> 24); + pke[16] = (pke[16] & ~0xffffff00) | (t << 8); + } + else + { + pke[23] = (pke[23] & ~0x000000ff) | (t >> 24); + pke[24] = (pke[24] & ~0xffffff00) | (t << 8); + } + + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + w0[0] = out[0]; + w0[1] = out[1]; + w0[2] = out[2]; + w0[3] = out[3]; + w1[0] = out[4]; + w1[1] = out[5]; + w1[2] = out[6]; + w1[3] = out[7]; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 0; + w3[3] = 0; + + u32 keymic[4]; + + keymic[0] = 0; + keymic[1] = 0; + keymic[2] = 0; + keymic[3] = 0; + + sha256_hmac_ctx_t ctx1; + + sha256_hmac_init_64 (&ctx1, w0, w1, w2, w3); + + sha256_hmac_update (&ctx1, pke, 102); + + sha256_hmac_final (&ctx1); + + u32 digest[4]; + + digest[0] = swap32_S (ctx1.opad.h[0]); + digest[1] = swap32_S (ctx1.opad.h[1]); + digest[2] = swap32_S (ctx1.opad.h[2]); + digest[3] = swap32_S (ctx1.opad.h[3]); + + // AES CMAC + + u32 ks[44]; + + aes128_set_encrypt_key (ks, digest, s_te0, s_te1, s_te2, s_te3, s_te4); + + u32 m[4]; + + m[0] = 0; + m[1] = 0; + m[2] = 0; + m[3] = 0; + + u32 iv[4]; + + iv[0] = 0; + iv[1] = 0; + iv[2] = 0; + iv[3] = 0; + + int eapol_left; + int eapol_idx; + + for (eapol_left = wpa->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]; + + 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]; + + u32 k[4]; + + k[0] = 0; + k[1] = 0; + k[2] = 0; + k[3] = 0; + + aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); + + make_kn (k); + + if (eapol_left < 16) + { + make_kn (k); + } + + m[0] ^= k[0]; + m[1] ^= k[1]; + m[2] ^= k[2]; + m[3] ^= k[3]; + + m[0] ^= iv[0]; + m[1] ^= iv[1]; + m[2] ^= iv[2]; + m[3] ^= iv[3]; + + aes128_encrypt (ks, m, keymic, s_te0, s_te1, s_te2, s_te3, s_te4); + + /** + * final compare + */ + + if ((keymic[0] == wpa->keymic[0]) + && (keymic[1] == wpa->keymic[1]) + && (keymic[2] == wpa->keymic[2]) + && (keymic[3] == wpa->keymic[3])) + { + if (atomic_inc (&hashes_shown[digest_cur]) == 0) + { + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0); + } } } } diff --git a/docs/changes.txt b/docs/changes.txt index 3dcf79dfd..f987750df 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -4,6 +4,7 @@ ## Improvements ## +- HCCAPX management: Use advanced hints in message_pair stored by hcxtools about endian bitness of replay counter - OpenCL kernels: Abort session if kernel self-test failed - Added JtR-compatible support for hex notation in rules engine @@ -12,6 +13,7 @@ ## - Fixed a missing kernel in -m 5600 in combination with -a 3 and -O if mask is >= 16 characters +- Fixed missing code section in -m 2500 and -m 2501 to crack corrupted handshakes with a LE endian bitness base * changes v4.0.1 -> v4.1.0 diff --git a/include/interface.h b/include/interface.h index a00009b6e..5c0e115bf 100644 --- a/include/interface.h +++ b/include/interface.h @@ -194,6 +194,8 @@ typedef struct wpa u32 hash[4]; int nonce_compare; int nonce_error_corrections; + int detected_le; + int detected_be; } wpa_t; diff --git a/include/types.h b/include/types.h index 1a5f23ecb..c68b77d04 100644 --- a/include/types.h +++ b/include/types.h @@ -1581,6 +1581,7 @@ typedef struct user_options bool workload_profile_chgd; bool segment_size_chgd; bool hccapx_message_pair_chgd; + bool nonce_error_corrections_chgd; bool advice_disable; bool benchmark; diff --git a/src/hashes.c b/src/hashes.c index b5f1e2cde..0e59e6c5f 100644 --- a/src/hashes.c +++ b/src/hashes.c @@ -831,7 +831,54 @@ int hashes_init_stage1 (hashcat_ctx_t *hashcat_ctx) wpa->message_pair = (u8) user_options->hccapx_message_pair; } - wpa->nonce_error_corrections = user_options->nonce_error_corrections; + if (wpa->message_pair & (1 << 4)) + { + // ap-less attack detected, nc not needed + + wpa->nonce_error_corrections = 0; + } + else + { + if (wpa->message_pair & (1 << 7)) + { + // replaycount not checked, nc needed + + wpa->nonce_error_corrections = user_options->nonce_error_corrections; + } + else + { + // replaycount checked, nc not needed, but we allow user overwrites + + if (user_options->nonce_error_corrections_chgd == true) + { + wpa->nonce_error_corrections = user_options->nonce_error_corrections; + } + else + { + wpa->nonce_error_corrections = 0; + } + } + } + + // now some optimization related to replay counter endianess + // hcxtools has techniques to detect them + // since we can not guarantee to get our handshakes from hcxtools we enable both by default + // 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; + + if (wpa->message_pair & (1 << 5)) + { + wpa->detected_le = 1; + wpa->detected_be = 0; + } + else if (wpa->message_pair & (1 << 6)) + { + wpa->detected_le = 0; + wpa->detected_be = 1; + } } } @@ -1652,6 +1699,9 @@ int hashes_init_selftest (hashcat_ctx_t *hashcat_ctx) wpa_t *wpa = (wpa_t *) st_esalts_buf; + wpa->detected_le = 1; + wpa->detected_be = 0; + wpa->nonce_error_corrections = 3; } else if (hashconfig->opts_type & OPTS_TYPE_BINARY_HASHFILE) diff --git a/src/user_options.c b/src/user_options.c index 9c9f2e20d..dd0d237dd 100644 --- a/src/user_options.c +++ b/src/user_options.c @@ -396,7 +396,8 @@ int user_options_getopt (hashcat_ctx_t *hashcat_ctx, int argc, char **argv) case IDX_LOGFILE_DISABLE: user_options->logfile_disable = true; break; case IDX_HCCAPX_MESSAGE_PAIR: user_options->hccapx_message_pair = hc_strtoul (optarg, NULL, 10); user_options->hccapx_message_pair_chgd = true; break; - case IDX_NONCE_ERROR_CORRECTIONS: user_options->nonce_error_corrections = hc_strtoul (optarg, NULL, 10); break; + case IDX_NONCE_ERROR_CORRECTIONS: user_options->nonce_error_corrections = hc_strtoul (optarg, NULL, 10); + user_options->nonce_error_corrections_chgd = true; break; case IDX_TRUECRYPT_KEYFILES: user_options->truecrypt_keyfiles = optarg; break; case IDX_VERACRYPT_KEYFILES: user_options->veracrypt_keyfiles = optarg; break; case IDX_VERACRYPT_PIM: user_options->veracrypt_pim = hc_strtoul (optarg, NULL, 10); break;