for hashcat #1437, keepass WUA

m13400_cl.diff

diff --git a/OpenCL/m13400.cl b/OpenCL/m13400.cl
index 77320ab..b45a054 100644
--- a/OpenCL/m13400.cl
+++ b/OpenCL/m13400.cl
@@ -44,7 +44,12 @@ __kernel void m13400_init (__global pw_t *pws, __global const kernel_rule_t *rul
   digest[6] = ctx.h[6];
   digest[7] = ctx.h[7];
 
-  if (esalt_bufs[digests_offset].version == 2 && esalt_bufs[digests_offset].keyfile_len == 0)
+  u32 keyfile_len        = esalt_bufs[digests_offset].keyfile_len;
+  u32 protected_data_len = 64; // TODO: this must depend on the user supplied data
+
+  u32 extra_bytes = keyfile_len + protected_data_len;
+
+  if (esalt_bufs[digests_offset].version == 2 && extra_bytes == 0)
   {
     u32 w0[4];
     u32 w1[4];
@@ -84,33 +89,86 @@ __kernel void m13400_init (__global pw_t *pws, __global const kernel_rule_t *rul
     digest[7] = ctx.h[7];
   }
 
-  if (esalt_bufs[digests_offset].keyfile_len != 0)
+  if (extra_bytes != 0)
   {
     u32 w0[4];
     u32 w1[4];
     u32 w2[4];
     u32 w3[4];
 
-    w0[0] = digest[0];
-    w0[1] = digest[1];
-    w0[2] = digest[2];
-    w0[3] = digest[3];
-    w1[0] = digest[4];
-    w1[1] = digest[5];
-    w1[2] = digest[6];
-    w1[3] = digest[7];
-    w2[0] = esalt_bufs[digests_offset].keyfile[0];
-    w2[1] = esalt_bufs[digests_offset].keyfile[1];
-    w2[2] = esalt_bufs[digests_offset].keyfile[2];
-    w2[3] = esalt_bufs[digests_offset].keyfile[3];
-    w3[0] = esalt_bufs[digests_offset].keyfile[4];
-    w3[1] = esalt_bufs[digests_offset].keyfile[5];
-    w3[2] = esalt_bufs[digests_offset].keyfile[6];
-    w3[3] = esalt_bufs[digests_offset].keyfile[7];
-
     sha256_init (&ctx);
 
-    sha256_update_64 (&ctx, w0, w1, w2, w3, 64);
+    // the previous sha256 checksum of the password is *always* used (first) in case of extra bytes
+
+    if (keyfile_len > 0)
+    {
+      w0[0] = digest[0];
+      w0[1] = digest[1];
+      w0[2] = digest[2];
+      w0[3] = digest[3];
+      w1[0] = digest[4];
+      w1[1] = digest[5];
+      w1[2] = digest[6];
+      w1[3] = digest[7];
+      w2[0] = esalt_bufs[digests_offset].keyfile[0];
+      w2[1] = esalt_bufs[digests_offset].keyfile[1];
+      w2[2] = esalt_bufs[digests_offset].keyfile[2];
+      w2[3] = esalt_bufs[digests_offset].keyfile[3];
+      w3[0] = esalt_bufs[digests_offset].keyfile[4];
+      w3[1] = esalt_bufs[digests_offset].keyfile[5];
+      w3[2] = esalt_bufs[digests_offset].keyfile[6];
+      w3[3] = esalt_bufs[digests_offset].keyfile[7];
+
+      sha256_update_64 (&ctx, w0, w1, w2, w3, 64);
+    }
+    else
+    {
+      w0[0] = digest[0];
+      w0[1] = digest[1];
+      w0[2] = digest[2];
+      w0[3] = digest[3];
+      w1[0] = digest[4];
+      w1[1] = digest[5];
+      w1[2] = digest[6];
+      w1[3] = digest[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;
+
+      sha256_update_64 (&ctx, w0, w1, w2, w3, 32);
+    }
+
+    if (protected_data_len > 0)
+    {
+      // TODO: the 64 bytes unprotected data blob comes here (this should later on depend on the "hash" information, user-supplied):
+
+      // this is the output of protected_data_unprotect.exe:
+      // The decrypted data is: 591674cc5fb167aa5f9e55c0a96cd66fadeb02366215d731364b19603a36b3b779c6e3b9dbf3027fd34c074b4c15c4a6efaf172933afd46dcaf2e6f9d911fcfc
+
+      w0[0] = 0x591674cc;
+      w0[1] = 0x5fb167aa;
+      w0[2] = 0x5f9e55c0;
+      w0[3] = 0xa96cd66f;
+      w1[0] = 0xadeb0236;
+      w1[1] = 0x6215d731;
+      w1[2] = 0x364b1960;
+      w1[3] = 0x3a36b3b7;
+      w2[0] = 0x79c6e3b9;
+      w2[1] = 0xdbf3027f;
+      w2[2] = 0xd34c074b;
+      w2[3] = 0x4c15c4a6;
+      w3[0] = 0xefaf1729;
+      w3[1] = 0x33afd46d;
+      w3[2] = 0xcaf2e6f9;
+      w3[3] = 0xd911fcfc;
+
+      sha256_update_64 (&ctx, w0, w1, w2, w3, 64);
+    }
 
     sha256_final (&ctx);