From 7ff09c6710e5d888f1d81abbcab7d8bd73837961 Mon Sep 17 00:00:00 2001
From: Fist0urs <eddy.maaalou@gmail.com>
Date: Sun, 7 May 2017 02:04:53 +0200
Subject: [PATCH] Preparing PR

---
 OpenCL/inc_types.cl             |   5 +-
 OpenCL/m15300.cl                | 391 ++++++++++++++--------------
 docs/changes.txt                |   1 +
 docs/credits.txt                |   1 +
 docs/readme.txt                 |   1 +
 extra/tab_completion/hashcat.sh |   2 +-
 include/interface.h             |   8 +-
 src/benchmark.c                 |   1 +
 src/interface.c                 |  93 ++++---
 src/usage.c                     |   2 -
 tools/test.pl                   | 436 +++++++++++++++++++++++++++++++-
 tools/test.sh                   |   4 +-
 12 files changed, 690 insertions(+), 255 deletions(-)

diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl
index 8452f5549..4c3999e6b 100644
--- a/OpenCL/inc_types.cl
+++ b/OpenCL/inc_types.cl
@@ -1042,7 +1042,6 @@ typedef struct
   u32 SID[32];
   u32 SID_len;
   u32 SID_offset;
-  u32 userKey[5];
   char cipher_algo[16];
   char hash_algo[16];
 
@@ -1399,12 +1398,14 @@ typedef struct
   u32 dgst[10];
   u32 out[10];
 
+  u32 userKey[5];
+
   /* dedicated to hmac-sha512 */
   u64 ipad64[8];
   u64 opad64[8];
   u64 dgst64[16];
   u64 out64[16];
-  
+
 } dpapimk_tmp_t;
 
 typedef struct
diff --git a/OpenCL/m15300.cl b/OpenCL/m15300.cl
index e0fb11dc0..fb741b617 100644
--- a/OpenCL/m15300.cl
+++ b/OpenCL/m15300.cl
@@ -698,9 +698,6 @@ void AES256_decrypt (const u32 *in, u32 *out, const u32 *rdk, SHM_TYPE u32 *s_td
          ^ rdk[59];
 }
 
-#define COMPARE_S "inc_comp_single.cl"
-#define COMPARE_M "inc_comp_multi.cl"
-
 void md4_transform_S (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
 {
   u32 a = digest[0];
@@ -1501,38 +1498,38 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
 
   u32 w0[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];
+  w0[0] = pws[gid].i[0];
+  w0[1] = pws[gid].i[1];
+  w0[2] = pws[gid].i[2];
+  w0[3] = pws[gid].i[3];
 
   u32 w1[4];
 
-  w1[0] = pws[gid].i[ 4];
-  w1[1] = pws[gid].i[ 5];
-  w1[2] = pws[gid].i[ 6];
-  w1[3] = pws[gid].i[ 7];
+  w1[0] = pws[gid].i[4];
+  w1[1] = pws[gid].i[5];
+  w1[2] = pws[gid].i[6];
+  w1[3] = pws[gid].i[7];
 
   u32 w2[4];
 
-  w2[0] = pws[gid].i[ 8];
-  w2[1] = pws[gid].i[ 9];
-  w2[2] = pws[gid].i[10];
-  w2[3] = pws[gid].i[11];
+  w2[0] = 0;
+  w2[1] = 0;
+  w2[2] = 0;
+  w2[3] = 0;
 
   u32 w3[4];
 
-  w3[0] = pws[gid].i[12];
-  w3[1] = pws[gid].i[13];
-  w3[2] = pws[gid].i[14];
-  w3[3] = pws[gid].i[15];
+  w3[0] = 0;
+  w3[1] = 0;
+  w3[2] = 0;
+  w3[3] = 0;
 
-  u32 pw_len = pws[gid].pw_len;
+  const u32 pw_len = pws[gid].pw_len;
 
-  append_0x80_4x4_S (w0, w1, w2, w3, pw_len);
+  append_0x80_2x4_S (w0, w1, pw_len);
 
-  make_unicode (w1, w2, w3);
-  make_unicode (w0, w0, w1);
+  make_unicode_S (w1, w2, w3);
+  make_unicode_S (w0, w0, w1);
 
   /**
    * main
@@ -1555,23 +1552,23 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
     digest_context[3] = SHA1M_D;
     digest_context[4] = SHA1M_E;
 
-    w0[0] = swap32 (w0[0]);
-    w0[1] = swap32 (w0[1]);
-    w0[2] = swap32 (w0[2]);
-    w0[3] = swap32 (w0[3]);
+    w0[0] = swap32_S (w0[0]);
+    w0[1] = swap32_S (w0[1]);
+    w0[2] = swap32_S (w0[2]);
+    w0[3] = swap32_S (w0[3]);
 
-    w1[0] = swap32 (w1[0]);
-    w1[1] = swap32 (w1[1]);
-    w1[2] = swap32 (w1[2]);
-    w1[3] = swap32 (w1[3]);
+    w1[0] = swap32_S (w1[0]);
+    w1[1] = swap32_S (w1[1]);
+    w1[2] = swap32_S (w1[2]);
+    w1[3] = swap32_S (w1[3]);
 
-    w2[0] = swap32 (w2[0]);
-    w2[1] = swap32 (w2[1]);
-    w2[2] = swap32 (w2[2]);
-    w2[3] = swap32 (w2[3]);
+    w2[0] = swap32_S (w2[0]);
+    w2[1] = swap32_S (w2[1]);
+    w2[2] = swap32_S (w2[2]);
+    w2[3] = swap32_S (w2[3]);
 
-    w3[0] = swap32 (w3[0]);
-    w3[1] = swap32 (w3[1]);
+    w3[0] = swap32_S (w3[0]);
+    w3[1] = swap32_S (w3[1]);
     w3[2] = 0;
     w3[3] = pw_len * 2 * 8;
 
@@ -1595,7 +1592,6 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
     digest_context[3] = swap32_S (digest_context[3]);
     digest_context[4] = 0;
   }
-
   /* initialize hmac-sha1 */
 
    /**
@@ -1681,13 +1677,12 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
   hmac_sha1_run_S (w0, w1, w2, w3, ipad, opad, key);
   /* this key is used as password for pbkdf2-hmac-* */
 
-  esalt_bufs[digests_offset].userKey[0] = key[0];
-  esalt_bufs[digests_offset].userKey[1] = key[1];
-  esalt_bufs[digests_offset].userKey[2] = key[2];
-  esalt_bufs[digests_offset].userKey[3] = key[3];
-  esalt_bufs[digests_offset].userKey[4] = key[4];
-  
-  
+  tmps[gid].userKey[0] = key[0];
+  tmps[gid].userKey[1] = key[1];
+  tmps[gid].userKey[2] = key[2];
+  tmps[gid].userKey[3] = key[3];
+  tmps[gid].userKey[4] = key[4];
+
   /* if DPAPImk version 1, pbkdf-hmac-sha1 is used */
   if (esalt_bufs[digests_offset].version == 1)
   {
@@ -1757,13 +1752,13 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
     tmps[gid].dgst[3] = digest[3];
     tmps[gid].dgst[4] = digest[4];
 
-    tmps[gid].out[0] = digest[0];
-    tmps[gid].out[1] = digest[1];
-    tmps[gid].out[2] = digest[2];
-    tmps[gid].out[3] = digest[3];
-    tmps[gid].out[4] = digest[4];
+    tmps[gid].out[0]  = digest[0];
+    tmps[gid].out[1]  = digest[1];
+    tmps[gid].out[2]  = digest[2];
+    tmps[gid].out[3]  = digest[3];
+    tmps[gid].out[4]  = digest[4];
 
-    /* We need bigger output ! */
+    /* We need bigger output! */
     w0[0] = esalt_bufs[digests_offset].iv[0];
     w0[1] = esalt_bufs[digests_offset].iv[1];
     w0[2] = esalt_bufs[digests_offset].iv[2];
@@ -1789,11 +1784,11 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
     tmps[gid].dgst[8] = digest[3];
     tmps[gid].dgst[9] = digest[4];
 
-    tmps[gid].out[5] = digest[0];
-    tmps[gid].out[6] = digest[1];
-    tmps[gid].out[7] = digest[2];
-    tmps[gid].out[8] = digest[3];
-    tmps[gid].out[9] = digest[4];
+    tmps[gid].out[5]  = digest[0];
+    tmps[gid].out[6]  = digest[1];
+    tmps[gid].out[7]  = digest[2];
+    tmps[gid].out[8]  = digest[3];
+    tmps[gid].out[9]  = digest[4];
   }
   /* if DPAPImk version 2, pbkdf-hmac-sha512 is used*/
   else if (esalt_bufs[digests_offset].version == 2)
@@ -1873,14 +1868,14 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
     tmps[gid].dgst64[6] = dgst64[6];
     tmps[gid].dgst64[7] = dgst64[7];
 
-    tmps[gid].out64[0] = dgst64[0];
-    tmps[gid].out64[1] = dgst64[1];
-    tmps[gid].out64[2] = dgst64[2];
-    tmps[gid].out64[3] = dgst64[3];
-    tmps[gid].out64[4] = dgst64[4];
-    tmps[gid].out64[5] = dgst64[5];
-    tmps[gid].out64[6] = dgst64[6];
-    tmps[gid].out64[7] = dgst64[7];
+    tmps[gid].out64[0]  = dgst64[0];
+    tmps[gid].out64[1]  = dgst64[1];
+    tmps[gid].out64[2]  = dgst64[2];
+    tmps[gid].out64[3]  = dgst64[3];
+    tmps[gid].out64[4]  = dgst64[4];
+    tmps[gid].out64[5]  = dgst64[5];
+    tmps[gid].out64[6]  = dgst64[6];
+    tmps[gid].out64[7]  = dgst64[7];
   }
 }
 
@@ -1919,18 +1914,18 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
     {
       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);
+
+      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++)
       {
@@ -1938,7 +1933,7 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
         u32x w1[4];
         u32x w2[4];
         u32x w3[4];
-    
+
         /* Microsoft PBKDF2 implementation. On purpose?
            Misunderstanding of them? Dunno...
         */
@@ -1958,22 +1953,22 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
         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]);
@@ -2016,14 +2011,14 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
     dgst[6] = pack64v (tmps, dgst64, gid, 6);
     dgst[7] = pack64v (tmps, dgst64, gid, 7);
 
-    out[0] = pack64v (tmps, out64, gid, 0);
-    out[1] = pack64v (tmps, out64, gid, 1);
-    out[2] = pack64v (tmps, out64, gid, 2);
-    out[3] = pack64v (tmps, out64, gid, 3);
-    out[4] = pack64v (tmps, out64, gid, 4);
-    out[5] = pack64v (tmps, out64, gid, 5);
-    out[6] = pack64v (tmps, out64, gid, 6);
-    out[7] = pack64v (tmps, out64, gid, 7);
+    out[0]  = pack64v (tmps, out64, gid, 0);
+    out[1]  = pack64v (tmps, out64, gid, 1);
+    out[2]  = pack64v (tmps, out64, gid, 2);
+    out[3]  = pack64v (tmps, out64, gid, 3);
+    out[4]  = pack64v (tmps, out64, gid, 4);
+    out[5]  = pack64v (tmps, out64, gid, 5);
+    out[6]  = pack64v (tmps, out64, gid, 6);
+    out[7]  = pack64v (tmps, out64, gid, 7);
 
     for (u32 j = 0; j < loop_cnt; j++)
     {
@@ -2032,7 +2027,7 @@ __kernel void m15300_loop (__global pw_t *pws, __global const kernel_rule_t *rul
       u64x w2[4];
       u64x w3[4];
 
-      /* Microsoft PBKDF2 implementation. On purpose? 
+      /* Microsoft PBKDF2 implementation. On purpose?
          Misunderstanding of them? Dunno...
       */
       w0[0] = out[0];
@@ -2096,7 +2091,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     u32 w1[4];
     u32 w2[4];
     u32 w3[4];
-    
+
     u32 ipad[5];
     u32 opad[5];
 
@@ -2106,7 +2101,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     __local u32 s_SPtrans[8][64];
     __local u32 s_skb[8][64];
-    
+
     for (u32 i = lid; i < 64; i += lsz)
     {
       s_SPtrans[0][i] = c_SPtrans[0][i];
@@ -2117,7 +2112,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
       s_SPtrans[5][i] = c_SPtrans[5][i];
       s_SPtrans[6][i] = c_SPtrans[6][i];
       s_SPtrans[7][i] = c_SPtrans[7][i];
-    
+
       s_skb[0][i] = c_skb[0][i];
       s_skb[1][i] = c_skb[1][i];
       s_skb[2][i] = c_skb[2][i];
@@ -2127,13 +2122,13 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
       s_skb[6][i] = c_skb[6][i];
       s_skb[7][i] = c_skb[7][i];
     }
-    
+
     barrier (CLK_LOCAL_MEM_FENCE);
-    
+
     if (gid >= gid_max) return;
 
     u32 key[6];
-    
+
     key[0] = swap32_S (tmps[gid].out[0]);
     key[1] = swap32_S (tmps[gid].out[1]);
     key[2] = swap32_S (tmps[gid].out[2]);
@@ -2142,102 +2137,102 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     key[5] = swap32_S (tmps[gid].out[5]);
 
     u32x iv[2];
-    
-    iv[0] = swap32_S (tmps[gid].out[6]);
-    iv[1] = swap32_S (tmps[gid].out[7]);
+
+    iv[0]  = swap32_S (tmps[gid].out[6]);
+    iv[1]  = swap32_S (tmps[gid].out[7]);
 
     u32 decrypted[26];
-    
+
     /* Construct 3DES keys */
-    
+
     const u32x a = (key[0]);
     const u32x b = (key[1]);
-    
+
     u32x Ka[16];
     u32x Kb[16];
-    
+
     _des_crypt_keysetup (a, b, Ka, Kb, s_skb);
-    
+
     const u32x c = (key[2]);
     const u32x d = (key[3]);
-    
+
     u32x Kc[16];
     u32x Kd[16];
-    
+
     _des_crypt_keysetup (c, d, Kc, Kd, s_skb);
-    
+
     const u32x e = (key[4]);
     const u32x f = (key[5]);
 
     u32x Ke[16];
     u32x Kf[16];
-    
+
     _des_crypt_keysetup (e, f, Ke, Kf, s_skb);
 
     u32 contents_pos;
     u32 contents_off;
     u32 wx_off;
-    
+
     for (wx_off = 0, contents_pos = 0, contents_off = 0; contents_pos < esalt_bufs[digests_offset].contents_len; wx_off += 2, contents_pos += 8, contents_off += 2)
     {
       /* First Pass */
-    
+
       u32x data[2];
-      
+
       data[0] = (u32x) swap32_S (esalt_bufs[digests_offset].contents[contents_off + 0]);
       data[1] = (u32x) swap32_S (esalt_bufs[digests_offset].contents[contents_off + 1]);
 
       u32x p1[2];
-    
+
       _des_crypt_decrypt (p1, data, Ke, Kf, s_SPtrans);
 
       /* Second Pass */
-      
+
       u32x p2[2];
-      
+
       _des_crypt_encrypt (p2, p1, Kc, Kd, s_SPtrans);
-      
+
       /* Third Pass */
 
       u32x out[2];
-      
+
       _des_crypt_decrypt (out, p2, Ka, Kb, s_SPtrans);
-      
+
       out[0] ^= iv[0];
       out[1] ^= iv[1];
 
       decrypted[wx_off + 0] = out[0];
       decrypted[wx_off + 1] = out[1];
-      
+
       iv[0] = data[0];
       iv[1] = data[1];
     }
 
     u32 hmacSalt[4];
-    u32 hmacTemoin[5];
+    u32 expectedHmac[5];
     u32 lastKey[16];
 
-    hmacSalt[0] = swap32_S (decrypted[0]);
-    hmacSalt[1] = swap32_S (decrypted[1]);
-    hmacSalt[2] = swap32_S (decrypted[2]);
-    hmacSalt[3] = swap32_S (decrypted[3]);
+    hmacSalt[0]     = swap32_S (decrypted[0]);
+    hmacSalt[1]     = swap32_S (decrypted[1]);
+    hmacSalt[2]     = swap32_S (decrypted[2]);
+    hmacSalt[3]     = swap32_S (decrypted[3]);
 
-    hmacTemoin[0] = swap32_S (decrypted[0 + 4]);
-    hmacTemoin[1] = swap32_S (decrypted[1 + 4]);
-    hmacTemoin[2] = swap32_S (decrypted[2 + 4]);
-    hmacTemoin[3] = swap32_S (decrypted[3 + 4]);
-    hmacTemoin[4] = swap32_S (decrypted[4 + 4]);
+    expectedHmac[0] = swap32_S (decrypted[4 + 0]);
+    expectedHmac[1] = swap32_S (decrypted[4 + 1]);
+    expectedHmac[2] = swap32_S (decrypted[4 + 2]);
+    expectedHmac[3] = swap32_S (decrypted[4 + 3]);
+    expectedHmac[4] = swap32_S (decrypted[4 + 4]);
 
     for(int i = 0; i < 16; i++)
     {
-      lastKey[i] = decrypted[i + 26 - 16]; 
+      lastKey[i] = decrypted[i + 26 - 16];
     }
 
-    w0[0] = esalt_bufs[digests_offset].userKey[0];
-    w0[1] = esalt_bufs[digests_offset].userKey[1];
-    w0[2] = esalt_bufs[digests_offset].userKey[2];
-    w0[3] = esalt_bufs[digests_offset].userKey[3];
-    w1[0] = esalt_bufs[digests_offset].userKey[4];
+    w0[0] = tmps[gid].userKey[0];
+    w0[1] = tmps[gid].userKey[1];
+    w0[2] = tmps[gid].userKey[2];
+    w0[3] = tmps[gid].userKey[3];
+    w1[0] = tmps[gid].userKey[4];
     w1[1] = 0;
     w1[2] = 0;
     w1[3] = 0;
@@ -2293,7 +2288,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     w3[1] = 0;
     w3[2] = 0;
     w3[3] = 0;
-    
+
     hmac_sha1_pad_S (w0, w1, w2, w3, ipad, opad);
 
     /**
@@ -2340,11 +2335,13 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
 
     #define il_pos 0
 
-    if ((hmacTemoin[0] == digest[0]) && (hmacTemoin[1] == digest[1])
-       && (hmacTemoin[2] == digest[2]) && (hmacTemoin[3] == digest[3])
-       && (hmacTemoin[4] == digest[4]))
+    if (  expectedHmac[0] == digest[0]
+      &&  expectedHmac[1] == digest[1]
+      &&  expectedHmac[2] == digest[2]
+      &&  expectedHmac[3] == digest[3]
+      &&  expectedHmac[4] == digest[4])
     {
-        mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
+      mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
     }
   }
   else if (esalt_bufs[digests_offset].version == 2)
@@ -2352,21 +2349,21 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     /**
      * aes shared
      */
-    
+
     #ifdef REAL_SHM
-    
+
     __local u32 s_td0[256];
     __local u32 s_td1[256];
     __local u32 s_td2[256];
     __local u32 s_td3[256];
     __local u32 s_td4[256];
-    
+
     __local u32 s_te0[256];
     __local u32 s_te1[256];
     __local u32 s_te2[256];
     __local u32 s_te3[256];
     __local u32 s_te4[256];
-    
+
     for (u32 i = lid; i < 256; i += lsz)
     {
       s_td0[i] = td0[i];
@@ -2374,38 +2371,38 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
       s_td2[i] = td2[i];
       s_td3[i] = td3[i];
       s_td4[i] = td4[i];
-    
+
       s_te0[i] = te0[i];
       s_te1[i] = te1[i];
       s_te2[i] = te2[i];
       s_te3[i] = te3[i];
       s_te4[i] = te4[i];
     }
-    
+
     barrier (CLK_LOCAL_MEM_FENCE);
-    
+
     #else
-    
+
     __constant u32a *s_td0 = td0;
     __constant u32a *s_td1 = td1;
     __constant u32a *s_td2 = td2;
     __constant u32a *s_td3 = td3;
     __constant u32a *s_td4 = td4;
-    
+
     __constant u32a *s_te0 = te0;
     __constant u32a *s_te1 = te1;
     __constant u32a *s_te2 = te2;
     __constant u32a *s_te3 = te3;
     __constant u32a *s_te4 = te4;
-    
+
     #endif
-    
+
     if (gid >= gid_max) return;
-    
+
     /* Construct AES key */
-    
+
     u32 key[8];
-    
+
     key[0] = h32_from_64 (tmps[gid].out64[0]);
     key[1] = l32_from_64 (tmps[gid].out64[0]);
     key[2] = h32_from_64 (tmps[gid].out64[1]);
@@ -2414,88 +2411,88 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     key[5] = l32_from_64 (tmps[gid].out64[2]);
     key[6] = h32_from_64 (tmps[gid].out64[3]);
     key[7] = l32_from_64 (tmps[gid].out64[3]);
-    
+
     u32 iv[4];
-    
-    iv[0] = h32_from_64 (tmps[gid].out64[4]);
-    iv[1] = l32_from_64 (tmps[gid].out64[4]);
-    iv[2] = h32_from_64 (tmps[gid].out64[5]);
-    iv[3] = l32_from_64 (tmps[gid].out64[5]);
-    
+
+    iv[0]  = h32_from_64 (tmps[gid].out64[4]);
+    iv[1]  = l32_from_64 (tmps[gid].out64[4]);
+    iv[2]  = h32_from_64 (tmps[gid].out64[5]);
+    iv[3]  = l32_from_64 (tmps[gid].out64[5]);
+
     #define KEYLEN 60
-    
+
     u32 rek[KEYLEN];
-    
+
     AES256_ExpandKey (key, rek, s_te0, s_te1, s_te2, s_te3, s_te4);
-    
+
     u32 rdk[KEYLEN];
-    
+
     #ifdef _unroll
     #pragma unroll
     #endif
     for (u32 i = 0; i < KEYLEN; i++) rdk[i] = rek[i];
-    
+
     AES256_InvertKey (rdk, s_td0, s_td1, s_td2, s_td3, s_td4, s_te0, s_te1, s_te2, s_te3, s_te4);
-    
+
     /* 144 bytes */
     u32 decrypted[36] = { 0 };
-    
+
     u32 contents_pos;
     u32 contents_off;
     u32 wx_off;
-    
+
     for (wx_off = 0, contents_pos = 0, contents_off = 0; contents_pos < esalt_bufs[digests_offset].contents_len; wx_off += 4, contents_pos += 16, contents_off += 4)
     {
       u32 data[4];
-    
+
       data[0] = esalt_bufs[digests_offset].contents[contents_off + 0];
       data[1] = esalt_bufs[digests_offset].contents[contents_off + 1];
       data[2] = esalt_bufs[digests_offset].contents[contents_off + 2];
       data[3] = esalt_bufs[digests_offset].contents[contents_off + 3];
-    
+
       u32 out[4];
-    
+
       AES256_decrypt (data, out, rdk, s_td0, s_td1, s_td2, s_td3, s_td4);
-    
+
       out[0] ^= iv[0];
       out[1] ^= iv[1];
       out[2] ^= iv[2];
       out[3] ^= iv[3];
-    
+
       decrypted[wx_off + 0] = out[0];
       decrypted[wx_off + 1] = out[1];
       decrypted[wx_off + 2] = out[2];
       decrypted[wx_off + 3] = out[3];
-    
+
       iv[0] = data[0];
       iv[1] = data[1];
       iv[2] = data[2];
       iv[3] = data[3];
     }
-    
+
     u32 hmacSalt[4];
-    u32 hmacTemoin[16];
+    u32 expectedHmac[16];
     u32 lastKey[16];
-      
+
     hmacSalt[0] = decrypted[0];
     hmacSalt[1] = decrypted[1];
     hmacSalt[2] = decrypted[2];
     hmacSalt[3] = decrypted[3];
-      
+
     for(int i = 0; i < 16; i++)
     {
-      hmacTemoin[i] = decrypted[i + 4];
-      lastKey[i] = decrypted[i + 36 - 16]; 
+      expectedHmac[i] = decrypted[i + 4];
+      lastKey[i]      = decrypted[i + 36 - 16];
     }
-    
+
     u64 w0_x64[4];
     u64 w1_x64[4];
     u64 w2_x64[4];
     u64 w3_x64[4];
-    
-    w0_x64[0] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[0], esalt_bufs[digests_offset].userKey[1]);
-    w0_x64[1] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[2], esalt_bufs[digests_offset].userKey[3]);
-    w0_x64[2] = hl32_to_64_S (esalt_bufs[digests_offset].userKey[4], 0);
+
+    w0_x64[0] = hl32_to_64_S (tmps[gid].userKey[0], tmps[gid].userKey[1]);
+    w0_x64[1] = hl32_to_64_S (tmps[gid].userKey[2], tmps[gid].userKey[3]);
+    w0_x64[2] = hl32_to_64_S (tmps[gid].userKey[4], 0);
     w0_x64[3] = 0;
     w1_x64[0] = 0;
     w1_x64[1] = 0;
@@ -2509,12 +2506,12 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     w3_x64[1] = 0;
     w3_x64[2] = 0;
     w3_x64[3] = 0;
-    
+
     u64 ipad64[8];
     u64 opad64[8];
-    
+
     hmac_sha512_pad_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64);
-    
+
     w0_x64[0] = hl32_to_64_S (hmacSalt[0], hmacSalt[1]);
     w0_x64[1] = hl32_to_64_S (hmacSalt[2], hmacSalt[3]);
     w0_x64[2] = hl32_to_64_S (0x80000000, 0);
@@ -2531,13 +2528,13 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     w3_x64[1] = 0;
     w3_x64[2] = 0;
     w3_x64[3] = (128 + 16) * 8;
-    
+
     u64 dgst64[8];
-    
+
     hmac_sha512_run_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64, dgst64);
-    
+
     u64 encKey[8];
-    
+
     encKey[0] = dgst64[0];
     encKey[1] = dgst64[1];
     encKey[2] = dgst64[2];
@@ -2546,7 +2543,7 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     encKey[5] = dgst64[5];
     encKey[6] = dgst64[6];
     encKey[7] = dgst64[7];
-    
+
     w0_x64[0] = encKey[0];
     w0_x64[1] = encKey[1];
     w0_x64[2] = encKey[2];
@@ -2563,9 +2560,9 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     w3_x64[1] = 0;
     w3_x64[2] = 0;
     w3_x64[3] = 0;
-    
+
     hmac_sha512_pad_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64);
-    
+
     w0_x64[0] = hl32_to_64_S (lastKey[ 0], lastKey[ 1]);
     w0_x64[1] = hl32_to_64_S (lastKey[ 2], lastKey[ 3]);
     w0_x64[2] = hl32_to_64_S (lastKey[ 4], lastKey[ 5]);
@@ -2582,15 +2579,21 @@ __kernel void m15300_comp (__global pw_t *pws, __global const kernel_rule_t *rul
     w3_x64[1] = 0;
     w3_x64[2] = 0;
     w3_x64[3] = (128 + 16 * 4) * 8;
-    
+
     hmac_sha512_run_S (w0_x64, w1_x64, w2_x64, w3_x64, ipad64, opad64, dgst64);
 
     #define il_pos 0
 
-    if ((hmacTemoin[0] == h32_from_64 (dgst64[0])) && (hmacTemoin[1] == l32_from_64 (dgst64[0]))
-       && (hmacTemoin[14] == h32_from_64 (dgst64[7])) && (hmacTemoin[15] == l32_from_64 (dgst64[7])))
+    if (  expectedHmac[ 0] == h32_from_64 (dgst64[0])
+       && expectedHmac[ 1] == l32_from_64 (dgst64[0])
+       && expectedHmac[ 2] == h32_from_64 (dgst64[1])
+       && expectedHmac[ 3] == l32_from_64 (dgst64[1])
+       && expectedHmac[12] == h32_from_64 (dgst64[6])
+       && expectedHmac[13] == l32_from_64 (dgst64[6])
+       && expectedHmac[14] == h32_from_64 (dgst64[7])
+       && expectedHmac[15] == l32_from_64 (dgst64[7]))
     {
-        mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
+      mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
     }
   }
 }
diff --git a/docs/changes.txt b/docs/changes.txt
index 7acfae9e8..3757de308 100644
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -6,6 +6,7 @@
 
 - Added hash-mode   600 = Blake2-512
 - Added hash-mode 15200 = Blockchain, My Wallet, V2
+- Added hash-mode 15300 = DPAPI masterkey file v1 and v2
 
 ##
 ## Features
diff --git a/docs/credits.txt b/docs/credits.txt
index dbb1b6c51..db0a6855d 100644
--- a/docs/credits.txt
+++ b/docs/credits.txt
@@ -31,6 +31,7 @@ Jean-Christophe "Fist0urs" Delaunay <jean-christophe.delaunay@synacktiv.com> (@F
 * Kerberos TGS Rep enctype 23 kernel module
 * AxCrypt kernel module
 * KeePass kernel module
+* DPAPImk v1 and v2 kernel module
 
 Other contributors to hashcat
 
diff --git a/docs/readme.txt b/docs/readme.txt
index d128e38e6..22007843a 100644
--- a/docs/readme.txt
+++ b/docs/readme.txt
@@ -145,6 +145,7 @@ NVIDIA GPUs require "NVIDIA Driver" (367.x or later)
 - NTLM
 - Domain Cached Credentials (DCC), MS Cache
 - Domain Cached Credentials 2 (DCC2), MS Cache 2
+- DPAPI masterkey file v1 and v2
 - MS-AzureSync PBKDF2-HMAC-SHA256
 - descrypt
 - bsdicrypt
diff --git a/extra/tab_completion/hashcat.sh b/extra/tab_completion/hashcat.sh
index 11d07c76a..261982863 100644
--- a/extra/tab_completion/hashcat.sh
+++ b/extra/tab_completion/hashcat.sh
@@ -176,7 +176,7 @@ _hashcat ()
 {
   local VERSION=3.5.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 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"
+  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 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"
   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 1dc1ffbf8..7e4ef2633 100644
--- a/include/interface.h
+++ b/include/interface.h
@@ -429,7 +429,6 @@ typedef struct dpapimk
   u32 SID[32];
   u32 SID_len;
   u32 SID_offset;
-  u32 userKey[5];
   u8 cipher_algo[16];
   u8 hash_algo[16];
 
@@ -837,12 +836,14 @@ typedef struct dpapimk_tmp
   u32 dgst[10];
   u32 out[10];
 
+  u32 userKey[5];
+
   /* dedicated to hmac-sha512 */
   u64 ipad64[8];
   u64 opad64[8];
   u64 dgst64[16];
   u64 out64[16];
-  
+
 } dpapimk_tmp_t;
 
 typedef struct seven_zip_hook
@@ -1366,7 +1367,8 @@ typedef enum hash_type
   HASH_TYPE_ITUNES_BACKUP_10    = 57,
   HASH_TYPE_SKIP32              = 58,
   HASH_TYPE_BLAKE2B             = 59,
-  HASH_TYPE_DPAPIMK             = 60,
+  HASH_TYPE_CHACHA20            = 60,
+  HASH_TYPE_DPAPIMK             = 61,
 
 } hash_type_t;
 
diff --git a/src/benchmark.c b/src/benchmark.c
index fd85e023b..e6bc32ab8 100644
--- a/src/benchmark.c
+++ b/src/benchmark.c
@@ -85,6 +85,7 @@ const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_BUF[] =
   1000,
   1100,
   2100,
+  15300,
   12800,
   1500,
   12400,
diff --git a/src/interface.c b/src/interface.c
index 40a17bb21..8cbe95a54 100644
--- a/src/interface.c
+++ b/src/interface.c
@@ -2208,16 +2208,16 @@ static u32 parse_and_store_salt (u8 *out, u8 *in, u32 salt_len, MAYBE_UNUSED con
   {
     if (salt_len < 20)
     {
-      tmp_u32[ 9] = ((tmp_u32[ 4] >> 8) & 0x00FF0000) | ((tmp_u32[ 4] >> 16) & 0x000000FF);
-      tmp_u32[ 8] = ((tmp_u32[ 4] << 8) & 0x00FF0000) | ((tmp_u32[ 4] >>  0) & 0x000000FF);
-      tmp_u32[ 7] = ((tmp_u32[ 3] >> 8) & 0x00FF0000) | ((tmp_u32[ 3] >> 16) & 0x000000FF);
-      tmp_u32[ 6] = ((tmp_u32[ 3] << 8) & 0x00FF0000) | ((tmp_u32[ 3] >>  0) & 0x000000FF);
-      tmp_u32[ 5] = ((tmp_u32[ 2] >> 8) & 0x00FF0000) | ((tmp_u32[ 2] >> 16) & 0x000000FF);
-      tmp_u32[ 4] = ((tmp_u32[ 2] << 8) & 0x00FF0000) | ((tmp_u32[ 2] >>  0) & 0x000000FF);
-      tmp_u32[ 3] = ((tmp_u32[ 1] >> 8) & 0x00FF0000) | ((tmp_u32[ 1] >> 16) & 0x000000FF);
-      tmp_u32[ 2] = ((tmp_u32[ 1] << 8) & 0x00FF0000) | ((tmp_u32[ 1] >>  0) & 0x000000FF);
-      tmp_u32[ 1] = ((tmp_u32[ 0] >> 8) & 0x00FF0000) | ((tmp_u32[ 0] >> 16) & 0x000000FF);
-      tmp_u32[ 0] = ((tmp_u32[ 0] << 8) & 0x00FF0000) | ((tmp_u32[ 0] >>  0) & 0x000000FF);
+      tmp_u32[9] = ((tmp_u32[4] >> 8) & 0x00FF0000) | ((tmp_u32[4] >> 16) & 0x000000FF);
+      tmp_u32[8] = ((tmp_u32[4] << 8) & 0x00FF0000) | ((tmp_u32[4] >>  0) & 0x000000FF);
+      tmp_u32[7] = ((tmp_u32[3] >> 8) & 0x00FF0000) | ((tmp_u32[3] >> 16) & 0x000000FF);
+      tmp_u32[6] = ((tmp_u32[3] << 8) & 0x00FF0000) | ((tmp_u32[3] >>  0) & 0x000000FF);
+      tmp_u32[5] = ((tmp_u32[2] >> 8) & 0x00FF0000) | ((tmp_u32[2] >> 16) & 0x000000FF);
+      tmp_u32[4] = ((tmp_u32[2] << 8) & 0x00FF0000) | ((tmp_u32[2] >>  0) & 0x000000FF);
+      tmp_u32[3] = ((tmp_u32[1] >> 8) & 0x00FF0000) | ((tmp_u32[1] >> 16) & 0x000000FF);
+      tmp_u32[2] = ((tmp_u32[1] << 8) & 0x00FF0000) | ((tmp_u32[1] >>  0) & 0x000000FF);
+      tmp_u32[1] = ((tmp_u32[0] >> 8) & 0x00FF0000) | ((tmp_u32[0] >> 16) & 0x000000FF);
+      tmp_u32[0] = ((tmp_u32[0] << 8) & 0x00FF0000) | ((tmp_u32[0] >>  0) & 0x000000FF);
 
       salt_len = salt_len * 2;
     }
@@ -2839,9 +2839,9 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
 
   if (memcmp (SIGNATURE_DPAPIMK, input_buf, 9)) return (PARSER_SIGNATURE_UNMATCHED);
 
-  u32 *digest = (u32 *) hash_buf->digest;
+  u32 *digest        = (u32 *) hash_buf->digest;
 
-  salt_t *salt = hash_buf->salt;
+  salt_t *salt       = hash_buf->salt;
 
   dpapimk_t *dpapimk = (dpapimk_t *) hash_buf->esalt;
 
@@ -2873,7 +2873,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
   dpapimk->context = atoll ((const char *) context_pos);
 
   SID_pos = (u8 *) strchr ((const char *) context_pos, '*');
-  
+
   if (SID_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
 
   SID_pos++;
@@ -2922,7 +2922,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
 
   if (contents_len_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
 
-  if (contents_len_pos - iv_pos != iv_len) return (PARSER_SEPARATOR_UNMATCHED);
+  if (contents_len_pos - iv_pos != iv_len) return (PARSER_SALT_LENGTH);
 
   if (is_valid_hex_string (iv_pos, 32) == false) return (PARSER_SALT_ENCODING);
 
@@ -2935,7 +2935,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
   dpapimk->iv[1] = byte_swap_32 (dpapimk->iv[1]);
   dpapimk->iv[2] = byte_swap_32 (dpapimk->iv[2]);
   dpapimk->iv[3] = byte_swap_32 (dpapimk->iv[3]);
-  
+
   contents_len_pos++;
 
   dpapimk->contents_len = (atoll ((const char *) contents_len_pos));
@@ -2962,13 +2962,8 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
     dpapimk->contents[i] = byte_swap_32 (dpapimk->contents[i]);
   }
 
-  digest[0] = dpapimk->iv[0];
-  digest[1] = dpapimk->iv[1];
-  digest[2] = dpapimk->iv[2];
-  digest[3] = dpapimk->iv[3];
-
   u32 SID_len = cipher_algo_pos - 1 - SID_pos;
-  
+
   /* maximum size of SID supported */
   u8 *SID_unicode = (u8 *) hcmalloc (32 * 4);
   memset (SID_unicode, 0, 32 * 4);
@@ -2977,12 +2972,12 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
   {
     SID_unicode[i * 2] = SID_pos[i];
   }
-  
+
   SID_unicode[(SID_len + 1) * 2] = 0x80;
 
   /* Specific to DPAPI: needs trailing '\0' while computing hash */
   dpapimk->SID_len = (SID_len + 1) * 2;
-  
+
   memcpy ((u8 *) dpapimk->SID, SID_unicode, 32 * 4);
 
   for (u32 i = 0; i < 32; i++)
@@ -2990,8 +2985,20 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
     dpapimk->SID[i] = byte_swap_32 (dpapimk->SID[i]);
   }
 
+  digest[0] = dpapimk->iv[0];
+  digest[1] = dpapimk->iv[1];
+  digest[2] = dpapimk->iv[2];
+  digest[3] = dpapimk->iv[3];
+
+  salt->salt_buf[0] = dpapimk->iv[0];
+  salt->salt_buf[1] = dpapimk->iv[1];
+  salt->salt_buf[2] = dpapimk->iv[2];
+  salt->salt_buf[3] = dpapimk->iv[3];
+
+  salt->salt_len = 32;
+
   hcfree(SID_unicode);
-  
+
   return (PARSER_OK);
 }
 
@@ -12611,7 +12618,7 @@ int krb5tgs_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
 
   u8 *edata_ptr = (u8 *) krb5tgs->edata2;
 
-  krb5tgs->edata2_len = (data_len - 32) / 2 ;
+  krb5tgs->edata2_len = (data_len - 32) / 2;
 
   /* skip '$' */
   for (u32 i = 16 * 2 + 1; i < (krb5tgs->edata2_len * 2) + (16 * 2 + 1); i += 2)
@@ -18117,11 +18124,11 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
     u32 algorithm   = (u32) keepass->algorithm;
     u32 keyfile_len = (u32) keepass->keyfile_len;
 
-    u32 *ptr_final_random_seed  = (u32 *) keepass->final_random_seed ;
-    u32 *ptr_transf_random_seed = (u32 *) keepass->transf_random_seed ;
-    u32 *ptr_enc_iv             = (u32 *) keepass->enc_iv ;
-    u32 *ptr_contents_hash      = (u32 *) keepass->contents_hash ;
-    u32 *ptr_keyfile            = (u32 *) keepass->keyfile ;
+    u32 *ptr_final_random_seed  = (u32 *) keepass->final_random_seed;
+    u32 *ptr_transf_random_seed = (u32 *) keepass->transf_random_seed;
+    u32 *ptr_enc_iv             = (u32 *) keepass->enc_iv;
+    u32 *ptr_contents_hash      = (u32 *) keepass->contents_hash;
+    u32 *ptr_keyfile            = (u32 *) keepass->keyfile;
 
     /* specific to version 2 */
     u32 expected_bytes_len;
@@ -18206,7 +18213,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
     else if (version == 2)
     {
       expected_bytes_len = 8;
-      ptr_expected_bytes = (u32 *) keepass->expected_bytes ;
+      ptr_expected_bytes = (u32 *) keepass->expected_bytes;
 
       for (u32 i = 0; i < expected_bytes_len; i++, ptr_data += 8)
         sprintf (ptr_data, "%08x", ptr_expected_bytes[i]);
@@ -18499,7 +18506,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
   {
     dpapimk_t *dpapimks = (dpapimk_t *) esalts_buf;
 
-    dpapimk_t *dpapimk = &dpapimks[digest_cur];
+    dpapimk_t *dpapimk  = &dpapimks[digest_cur];
 
     u32 version      = (u32) dpapimk->version;
     u32 context      = (u32) dpapimk->context;
@@ -18510,13 +18517,13 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
     char *ptr_SID               = (char *) dpapimk->SID_tmp;
     char *ptr_cipher_algorithm  = (char *) dpapimk->cipher_algo;
     char *ptr_hash_algorithm    = (char *) dpapimk->hash_algo;
-    u32 *ptr_iv                 = (u32 *)  dpapimk->iv;
-    u32 *ptr_contents           = (u32 *)  dpapimk->contents;
+    u32  *ptr_iv                = (u32 *)  dpapimk->iv;
+    u32  *ptr_contents          = (u32 *)  dpapimk->contents;
 
     u32 u32_iv[4];
     u8 iv[32 + 1];
 
-    for (u32 i = 0 ; i < iv_len / 8; i++)
+    for (u32 i = 0; i < iv_len / 8; i++)
     {
       u32_iv[i] = byte_swap_32 (ptr_iv[i]);
       u32_to_hex_lower (u32_iv[i], iv +  i * 8);
@@ -18524,14 +18531,22 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
     iv[32] = '\0';
 
     u32 u32_contents[36];
-    u8 contents[288 + 1];
+    u8  contents[288 + 1];
 
-    for (u32 i = 0 ; i < contents_len / 8; i++)
+    for (u32 i = 0; i < contents_len / 8; i++)
     {
       u32_contents[i] = byte_swap_32 (ptr_contents[i]);
       u32_to_hex_lower (u32_contents[i], contents +  i * 8);
     }
-    contents[288] = '\0';
+
+    if (version == 1)
+    {
+      contents[208] = '\0';
+    }
+    else
+    {
+      contents[288] = '\0';
+    }
 
     snprintf (out_buf, out_len - 1, "%s%d*%d*%s*%s*%s*%d*%s*%d*%s",
       SIGNATURE_DPAPIMK,
@@ -22847,7 +22862,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
                  hashconfig->attack_exec    = ATTACK_EXEC_OUTSIDE_KERNEL;
                  hashconfig->opts_type      = OPTS_TYPE_PT_GENERATE_LE;
                  hashconfig->kern_type      = KERN_TYPE_DPAPIMK;
-                 hashconfig->dgst_size      = DGST_SIZE_4_5;
+                 hashconfig->dgst_size      = DGST_SIZE_4_4;
                  hashconfig->parse_func     = dpapimk_parse_hash;
                  hashconfig->opti_type      = OPTI_TYPE_ZERO_BYTE;
                  hashconfig->dgst_pos0      = 0;
diff --git a/src/usage.c b/src/usage.c
index 491801a7c..f39f4fad8 100644
--- a/src/usage.c
+++ b/src/usage.c
@@ -241,9 +241,7 @@ static const char *USAGE_BIG[] =
   "   1000 | NTLM                                             | Operating Systems",
   "   1100 | Domain Cached Credentials (DCC), MS Cache        | Operating Systems",
   "   2100 | Domain Cached Credentials 2 (DCC2), MS Cache 2   | Operating Systems",
-/* Fist0urs */
   "  15300 | DPAPI masterkey file v1 and v2                   | Operating Systems",
-/* Fist0urs_end */
   "  12800 | MS-AzureSync  PBKDF2-HMAC-SHA256                 | Operating Systems",
   "   1500 | descrypt, DES (Unix), Traditional DES            | Operating Systems",
   "  12400 | BSDiCrypt, Extended DES                          | Operating Systems",
diff --git a/tools/test.pl b/tools/test.pl
index 9a9565daf..4315f4ada 100755
--- a/tools/test.pl
+++ b/tools/test.pl
@@ -7,19 +7,19 @@
 
 use strict;
 use warnings;
-use Digest::MD4      qw (md4 md4_hex);
-use Digest::MD5      qw (md5 md5_hex);
-use Digest::SHA      qw (sha1 sha256 sha384 sha512 sha1_hex sha224_hex sha256_hex sha384_hex sha512_hex);
-use Digest::HMAC     qw (hmac hmac_hex);
-use Digest::Keccak   qw (keccak_256_hex);
-use Digest::BLAKE2   qw (blake2b_hex);
-use Crypt::MySQL     qw (password41);
-use Digest::GOST     qw (gost gost_hex);
-use Digest::HMAC_MD5 qw (hmac_md5);
-use Digest::CRC      qw (crc32);
+use Digest::MD4       qw (md4 md4_hex);
+use Digest::MD5       qw (md5 md5_hex);
+use Digest::SHA       qw (sha1 sha256 sha384 sha512 sha1_hex sha224_hex sha256_hex sha384_hex sha512_hex hmac_sha1 hmac_sha512);
+use Digest::HMAC      qw (hmac hmac_hex);
+use Digest::Keccak    qw (keccak_256_hex);
+use Digest::BLAKE2    qw (blake2b_hex);
+use Crypt::MySQL      qw (password41);
+use Digest::GOST      qw (gost gost_hex);
+use Digest::HMAC_MD5  qw (hmac_md5);
+use Digest::CRC       qw (crc32);
 use Crypt::PBKDF2;
 use Crypt::DES;
-use Crypt::ECB       qw (encrypt);
+use Crypt::ECB        qw (encrypt);
 use Crypt::CBC;
 use Crypt::Eksblowfish::Bcrypt qw (bcrypt en_base64);
 use Crypt::Digest::RIPEMD160   qw (ripemd160_hex);
@@ -47,7 +47,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, 7800, 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, 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, 7800, 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, 99999);
 
 my %is_unicode      = 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 2400 2410 3200 6300 7400 10500 10700);
@@ -2621,6 +2621,61 @@ sub verify
 
       next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
     }
+    elsif ($mode == 15300)
+    {
+      ($hash_in, $word) = split ":", $line;
+
+      next unless defined $hash_in;
+      next unless defined $word;
+
+      my @tmp_data  = split ('\$', $hash_in);
+
+      my $signature = $tmp_data[1];
+
+      next unless ($signature eq 'DPAPImk');
+
+      my @data = split ('\*', $tmp_data[2]);
+
+      next unless (scalar @data == 9);
+
+      my $version = shift @data;
+
+      next unless ($version == 1 || $version == 2);
+
+      my $context          = shift @data;
+
+      my $SID              = shift @data;
+
+      my $cipher_algorithm = shift @data;
+
+      my $hash_algorithm   = shift @data;
+
+      my $iteration        = shift @data;
+
+      my $iv               = shift @data;
+
+      my $cipher_len       = shift @data;
+
+      my $cipher           = shift @data;
+
+      next unless (length ($cipher) == $cipher_len);
+
+      if ($version == 1)
+      {
+        next unless ($cipher_len == 208);
+      }
+      elsif ($version == 2)
+      {
+        next unless ($cipher_len == 288);
+      }
+
+      $salt   = substr ($hash_in, length ('$DPAPImk$'));
+
+      $param  = $iv;
+      $param2 = $cipher;
+
+      next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
+    }
     else
     {
       print "ERROR: hash mode is not supported\n";
@@ -2996,6 +3051,14 @@ sub verify
 
       $hash_out = $hash_in;
     }
+    elsif ($mode == 15300)
+    {
+      $hash_out = gen_hash ($mode, $word, $salt, $iter, $param, $param2);
+
+      $len      = length $hash_out;
+
+      return unless (substr ($line, 0, $len) eq $hash_out);
+    }
     else
     {
       $hash_out = gen_hash ($mode, $word, $salt, $iter);
@@ -3510,6 +3573,12 @@ sub passthrough
     {
       $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32));
     }
+    elsif ($mode == 15300)
+    {
+      $salt_buf = get_random_dpapimk_salt ();
+
+      $tmp_hash = gen_hash ($mode, $word_buf, $salt_buf);
+    }
     else
     {
       print "ERROR: Unsupported hash type\n";
@@ -4450,6 +4519,20 @@ sub single
         }
       }
     }
+    elsif ($mode == 15300)
+    {
+      for (my $i = 1; $i < 16; $i++)
+      {
+        if ($len != 0)
+        {
+          rnd ($mode, $len, 16);
+        }
+        else
+        {
+          rnd ($mode, $i, 16);
+        }
+      }
+    }
   }
 }
 
@@ -8153,6 +8236,270 @@ END_CODE
 
     $tmp_hash = sprintf ("\$blockchain\$v2\$%d\$%s\$%s", $iterations, length ($salt_buf . $encrypted) / 2, $salt_buf . $encrypted);
   }
+  elsif ($mode == 15300)
+  {
+    my @salt_arr = split ('\*', $salt_buf);
+
+    my $version   = $salt_arr[0];
+
+    my $context   = $salt_arr[1];
+
+    my $SID       = $salt_arr[2];
+
+    my $cipher_algorithm;
+
+    my $hash_algorithm;
+
+    my $iterations = $salt_arr[3];
+
+    my $salt       = pack ("H*", $salt_arr[4]);
+
+    my $cipher_len;
+
+    if ($version == 1)
+    {
+      $cipher_algorithm = "des3";
+
+      $hash_algorithm   = "sha1";
+
+      $cipher_len       = 208;
+    }
+    else
+    {
+      $cipher_algorithm = "aes256";
+
+      $hash_algorithm   = "sha512";
+
+      $cipher_len       = 288;
+    }
+
+    my $cipher;
+
+    # intermediate values
+
+    my $user_hash;
+    my $user_derivationKey;
+    my $encKey;
+    my $expected_hmac;
+    my $cleartext;
+
+    if ($context == 1)
+    {
+       $user_hash = sha1 (encode ("UTF-16LE", $word_buf));
+    }
+    elsif ($context == 2)
+    {
+       $user_hash = md4 (encode ("UTF-16LE", $word_buf));
+    }
+
+    $user_derivationKey = hmac_sha1 (encode ("UTF-16LE", $SID . "\x00"), $user_hash);
+
+    my $hmacSalt = randbytes (16);
+    my $last_key = randbytes (64);
+
+    if ($version == 1)
+    {
+      $encKey        = hmac_sha1 ($hmacSalt, $user_derivationKey);
+      $expected_hmac = hmac_sha1 ($last_key, $encKey);
+
+      # need padding because keyLen is 24 and hashLen 20
+      $expected_hmac = $expected_hmac . randbytes(4);
+    }
+    elsif ($version == 2)
+    {
+      $encKey        = hmac_sha512 ($hmacSalt, $user_derivationKey);
+      $expected_hmac = hmac_sha512 ($last_key, $encKey);
+    }
+
+    $cleartext = $hmacSalt . $expected_hmac . $last_key;
+
+    my $derived_key;
+    my $key;
+    my $iv;
+
+    my $pbkdf2;
+
+    if ($version == 1)
+    {
+      $derived_key = dpapi_pbkdf2 ($user_derivationKey, $salt, $iterations, 32, \&hmac_sha1);
+    }
+    elsif ($version == 2)
+    {
+      $derived_key = dpapi_pbkdf2 ($user_derivationKey, $salt, $iterations, 48, \&hmac_sha512);
+    }
+
+    if (defined $additional_param)
+    {
+      $salt = pack ("H*", $additional_param);
+    }
+
+    if (defined $additional_param2)
+    {
+      $cipher = $additional_param2;
+
+      if ($version == 1)
+      {
+        $key = substr ($derived_key,   0, 24);
+        $iv  = substr ($derived_key,  24,  8);
+
+        my $p1 = Crypt::ECB->new ({
+          key         => substr ($key, 0, 8),
+          cipher      => "DES",
+          literal_key => 1,
+          header      => "none",
+          keysize     => 8,
+          padding     => "null",
+        });
+
+        my $p2 = Crypt::ECB->new ({
+          key         => substr ($key, 8, 8),
+          cipher      => "DES",
+          literal_key => 1,
+          header      => "none",
+          keysize     => 8,
+          padding     => "null",
+        });
+
+        my $p3 = Crypt::ECB->new ({
+          key         => substr ($key, 16, 8),
+          cipher      => "DES",
+          literal_key => 1,
+          header      => "none",
+          keysize     => 8,
+          padding     => "null",
+        });
+
+        # let's compute a 3DES-EDE-CBC decryption
+        $iv = substr ($cipher, 0, 8);
+
+        my $out1;
+        my $out2;
+        my $out3;
+        my $expected_cleartext = "";
+
+        # size of cipherlen is 104 bytes
+        for (my $k = 0; $k < 13; $k++)
+        {
+          $out1 = $p3->decrypt (substr ($cipher, $k * 8, 8));
+          $out2 = $p2->encrypt ($out1);
+          $out3 = $p1->decrypt ($out2);
+
+          $expected_cleartext .= substr ($out3, 0, 8) ^ $iv;
+
+          $iv = substr ($cipher, $k * 8, 8);
+        }
+
+        if ($expected_cleartext != $cleartext)
+        {
+          $cleartext = "0" x 104;
+        }
+      }
+      elsif ($version == 2)
+      {
+        $key = substr ($derived_key,  0, 32);
+        $iv  = substr ($derived_key, 32, 16);
+
+        my $aes = Crypt::CBC->new ({
+          key         => $key,
+          cipher      => "Crypt::Rijndael",
+          iv          => $iv,
+          literal_key => 1,
+          header      => "none",
+          keysize     => 32,
+          padding     => "null",
+        });
+
+        my $expected_cleartext = $aes->decrypt(pack ("H*", $cipher));
+
+        if ($expected_cleartext != $cleartext)
+        {
+          $cleartext = "0" x 144;
+        }
+      }
+    }
+
+    if ($version == 1)
+    {
+      $key = substr ($derived_key,   0, 24);
+      $iv  = substr ($derived_key,  24,  8);
+
+      my $p1 = Crypt::ECB->new ({
+        key         => substr ($key, 0, 8),
+        cipher      => "DES",
+        literal_key => 1,
+        header      => "none",
+        keysize     => 8,
+        padding     => "null",
+      });
+
+      my $p2 = Crypt::ECB->new ({
+        key         => substr ($key, 8, 8),
+        cipher      => "DES",
+        literal_key => 1,
+        header      => "none",
+        keysize     => 8,
+        padding     => "null",
+      });
+
+      my $p3 = Crypt::ECB->new ({
+        key         => substr ($key, 16, 8),
+        cipher      => "DES",
+        literal_key => 1,
+        header      => "none",
+        keysize     => 8,
+        padding     => "null",
+      });
+
+      # let's compute a 3DES-EDE-CBC encryption
+
+      # compute first block
+      my $out1 = $p1->encrypt (substr ($cleartext, 0, 8) ^ $iv);
+      my $out2 = $p2->decrypt ($out1);
+      my $out3 = $p3->encrypt ($out2);
+
+      $cipher = substr ($out3, 0, 8);
+
+      # size of cipherlen is 104 bytes
+      for (my $k = 1; $k < 13; $k++)
+      {
+        $iv = $out3;
+
+        $out1 = $p1->encrypt (substr ($cleartext, $k * 8, 8) ^ $iv);
+        $out2 = $p2->decrypt ($out1);
+        $out3 = $p3->encrypt ($out2);
+
+        $cipher .= substr ($out3, 0, 8);
+      }
+    }
+    else
+    {
+      $key = substr ($derived_key,  0, 32);
+      $iv  = substr ($derived_key, 32, 16);
+
+      my $aes = Crypt::CBC->new ({
+        key         => $key,
+        cipher      => "Crypt::Rijndael",
+        iv          => $iv,
+        literal_key => 1,
+        header      => "none",
+        keysize     => 32,
+        padding     => "null",
+      });
+
+      $cipher = $aes->encrypt ($cleartext);
+    }
+
+   $tmp_hash = sprintf ('$DPAPImk$%d*%d*%s*%s*%s*%d*%s*%d*%s',
+                 $version,
+                 $context,
+                 $SID,
+                 $cipher_algorithm,
+                 $hash_algorithm,
+                 $iterations,
+                 unpack ("H*", $salt),
+                 $cipher_len,
+                 unpack ("H*", $cipher));
+  }
   elsif ($mode == 99999)
   {
     $tmp_hash = sprintf ("%s", $word_buf);
@@ -8161,6 +8508,30 @@ END_CODE
   return ($tmp_hash);
 }
 
+#Thanks to Jochen Hoenicke <hoenicke@gmail.com>
+# (one of the authors of Palm Keyring)
+# for these next two subs.
+sub dpapi_pbkdf2
+{
+    my ($password, $salt, $iter, $keylen, $prf) = @_;
+    my ($k, $t, $u, $ui, $i);
+    $t = "";
+    for ($k = 1; length($t) <  $keylen; $k++)
+    {
+      $u = $ui = &$prf($salt.pack('N', $k), $password);
+      for ($i = 1; $i < $iter; $i++)
+      {
+        # modification to fit Microsoft
+        # weird pbkdf2 implementation...
+        $ui = &$prf($u, $password);
+        $u ^= $ui;
+      }
+      $t .= $u;
+    }
+    return substr($t, 0, $keylen);
+}
+
+
 sub rnd
 {
   my $mode = shift;
@@ -8267,6 +8638,10 @@ sub rnd
   {
     $salt_buf = get_pstoken_salt ();
   }
+  elsif ($mode == 15300)
+  {
+    $salt_buf = get_random_dpapimk_salt ();
+  }
   else
   {
     my @salt_arr;
@@ -9707,6 +10082,43 @@ sub get_random_dnssec_salt
   return $salt_buf;
 }
 
+sub get_random_dpapimk_salt
+{
+  my $salt_buf = "";
+
+  my $version = get_random_num (1, 3);
+
+  my $context = get_random_num (1, 3);
+
+  my $iterations;
+
+  my $SID = sprintf ('S-15-21-%d-%d-%d-%d',
+             get_random_num(400000000,490000000),
+             get_random_num(400000000,490000000),
+             get_random_num(400000000,490000000),
+             get_random_num(1000,1999));
+
+  if ($version == 1)
+  {
+    $iterations = get_random_num (4000, 24000);
+  }
+  elsif ($version == 2)
+  {
+    $iterations = get_random_num (8000, 17000);
+  }
+
+  my $iv = randbytes (16);
+  $iv    = unpack ("H*", $iv);
+
+  $salt_buf = $version . '*' .
+              $context . '*' .
+              $SID     . '*' .
+              $iterations       . '*' .
+              $iv . '*';
+
+  return $salt_buf;
+}
+
 sub md5bit
 {
   my $digest = shift;
diff --git a/tools/test.sh b/tools/test.sh
index ea7893239..236e5e7ec 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 7800 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 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 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 7800 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 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 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"
+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"
 
 OPTS="--quiet --force --potfile-disable --runtime 400 --gpu-temp-disable --weak-hash-threshold=0"