diff --git a/OpenCL/inc_hash_md5.cl b/OpenCL/inc_hash_md5.cl index 0a9080ede..36a6fe6e4 100644 --- a/OpenCL/inc_hash_md5.cl +++ b/OpenCL/inc_hash_md5.cl @@ -567,7 +567,7 @@ typedef struct md5_hmac_ctx } md5_hmac_ctx_t; -void md5_hmac_init (md5_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4]) +void md5_hmac_init_64 (md5_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4]) { u32 t0[4]; u32 t1[4]; @@ -621,6 +621,120 @@ void md5_hmac_init (md5_hmac_ctx_t *ctx, const u32 w0[4], const u32 w1[4], const md5_update_64 (&ctx->opad, t0, t1, t2, t3, 64); } +void md5_hmac_init (md5_hmac_ctx_t *ctx, const u32 *w, const int len) +{ + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + if (len > 64) + { + md5_ctx_t tmp; + + md5_init (&tmp); + + md5_update (&tmp, w, len); + + md5_final (&tmp); + + w0[0] = tmp.h[0]; + w0[1] = tmp.h[1]; + w0[2] = tmp.h[2]; + w0[3] = tmp.h[3]; + 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; + } + else + { + w0[0] = w[ 0]; + w0[1] = w[ 1]; + w0[2] = w[ 2]; + w0[3] = w[ 3]; + w1[0] = w[ 4]; + w1[1] = w[ 5]; + w1[2] = w[ 6]; + w1[3] = w[ 7]; + w2[0] = w[ 8]; + w2[1] = w[ 9]; + w2[2] = w[10]; + w2[3] = w[11]; + w3[0] = w[12]; + w3[1] = w[13]; + w3[2] = w[14]; + w3[3] = w[15]; + } + + md5_hmac_init_64 (ctx, w0, w1, w2, w3); +} + +void md5_hmac_init_global (md5_hmac_ctx_t *ctx, __global const u32 *w, const int len) +{ + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + + if (len > 64) + { + md5_ctx_t tmp; + + md5_init (&tmp); + + md5_update_global (&tmp, w, len); + + md5_final (&tmp); + + w0[0] = tmp.h[0]; + w0[1] = tmp.h[1]; + w0[2] = tmp.h[2]; + w0[3] = tmp.h[3]; + 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; + } + else + { + w0[0] = w[ 0]; + w0[1] = w[ 1]; + w0[2] = w[ 2]; + w0[3] = w[ 3]; + w1[0] = w[ 4]; + w1[1] = w[ 5]; + w1[2] = w[ 6]; + w1[3] = w[ 7]; + w2[0] = w[ 8]; + w2[1] = w[ 9]; + w2[2] = w[10]; + w2[3] = w[11]; + w3[0] = w[12]; + w3[1] = w[13]; + w3[2] = w[14]; + w3[3] = w[15]; + } + + md5_hmac_init_64 (ctx, w0, w1, w2, w3); +} + void md5_hmac_update_64 (md5_hmac_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len) { md5_update_64 (&ctx->ipad, w0, w1, w2, w3, len); diff --git a/OpenCL/m02500.cl b/OpenCL/m02500.cl index 50227d969..b979aa871 100644 --- a/OpenCL/m02500.cl +++ b/OpenCL/m02500.cl @@ -404,7 +404,7 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul md5_hmac_ctx_t ctx2; - md5_hmac_init (&ctx2, t0, t1, t2, t3); + md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); @@ -554,7 +554,7 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul md5_hmac_ctx_t ctx2; - md5_hmac_init (&ctx2, t0, t1, t2, t3); + md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); diff --git a/OpenCL/m02501.cl b/OpenCL/m02501.cl index 275667494..eead090fa 100644 --- a/OpenCL/m02501.cl +++ b/OpenCL/m02501.cl @@ -245,7 +245,7 @@ __kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rul md5_hmac_ctx_t ctx2; - md5_hmac_init (&ctx2, t0, t1, t2, t3); + md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); @@ -395,7 +395,7 @@ __kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rul md5_hmac_ctx_t ctx2; - md5_hmac_init (&ctx2, t0, t1, t2, t3); + md5_hmac_init_64 (&ctx2, t0, t1, t2, t3); md5_hmac_update_global (&ctx2, wpa->eapol, wpa->eapol_len); diff --git a/OpenCL/m11900.cl b/OpenCL/m11900.cl index 8c52db1a0..ec7c4ad25 100644 --- a/OpenCL/m11900.cl +++ b/OpenCL/m11900.cl @@ -60,31 +60,9 @@ __kernel void m11900_init (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - u32 w0[4]; - u32 w1[4]; - u32 w2[4]; - u32 w3[4]; - - w0[0] = pws[gid].i[ 0]; - w0[1] = pws[gid].i[ 1]; - w0[2] = pws[gid].i[ 2]; - w0[3] = pws[gid].i[ 3]; - w1[0] = pws[gid].i[ 4]; - w1[1] = pws[gid].i[ 5]; - w1[2] = pws[gid].i[ 6]; - w1[3] = pws[gid].i[ 7]; - w2[0] = pws[gid].i[ 8]; - w2[1] = pws[gid].i[ 9]; - w2[2] = pws[gid].i[10]; - w2[3] = pws[gid].i[11]; - w3[0] = pws[gid].i[12]; - w3[1] = pws[gid].i[13]; - w3[2] = pws[gid].i[14]; - w3[3] = pws[gid].i[15]; - md5_hmac_ctx_t md5_hmac_ctx; - md5_hmac_init (&md5_hmac_ctx, w0, w1, w2, w3); + md5_hmac_init_global (&md5_hmac_ctx, pws[gid].i, pws[gid].pw_len); tmps[gid].ipad[0] = md5_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = md5_hmac_ctx.ipad.h[1]; @@ -102,6 +80,11 @@ __kernel void m11900_init (__global pw_t *pws, __global const kernel_rule_t *rul { md5_hmac_ctx_t md5_hmac_ctx2 = md5_hmac_ctx; + u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[4]; + w0[0] = j << 24; w0[1] = 0; w0[2] = 0; diff --git a/src/interface.c b/src/interface.c index 6377a3282..b7058832b 100644 --- a/src/interface.c +++ b/src/interface.c @@ -24675,6 +24675,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) switch (hashconfig->hash_mode) { + case 2100: hashconfig->pw_max = PW_MAX; + break; case 5200: hashconfig->pw_max = PW_MAX; break; case 7900: hashconfig->pw_max = PW_MAX; @@ -24691,6 +24693,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) break; case 11600: hashconfig->pw_max = PW_MAX; break; + case 11900: hashconfig->pw_max = PW_MAX; + break; case 12200: hashconfig->pw_max = PW_MAX; break; case 12400: hashconfig->pw_max = PW_MAX; @@ -24783,8 +24787,6 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) break; case 10900: hashconfig->pw_max = 64; // PBKDF2-HMAC-SHA256 max break; - case 11900: hashconfig->pw_max = 64; // PBKDF2-HMAC-MD5 max - break; case 12000: hashconfig->pw_max = 64; // PBKDF2-HMAC-SHA1 max break; case 12001: hashconfig->pw_max = 64; // PBKDF2-HMAC-SHA1 max