Remove PBKDF2-HMAC-MD5 includes password length limit

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);

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);
 

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);
 

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;

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