Diskcryptor password

[Turysta]

I am desperate to break the password of my backup. I encoded the CD image and recorded it. Unfortunately, I can not remember the password for anything, but I remember what could be the string of letters.

Is possible to use hash cat for crack password diskcryptor?

I tried mount ISO image as truecrypt contener, but is Is't working.

Here is other mounted volume (for obtain information)

  q1.JPG (Size: 15.63 KB / Downloads: 15)


Best wishes

[atom]

You can not use hashcat to crack this even if this exact combination of mode, cipher, and hash are used in TC/VC, too. It would require to add a new module for diskcryptor.

[Turysta]

Thank you very much for information.
I will live with the hope that one day the program will support Diskcryptor

Best wishes

[philsmd]

well, we had a similar request from here: https://hashcat.net/forum/thread-8012.html

and I tried to patch hashcat v5.1.0 (with commit: 1493bc01cfbf3fe0b7f4f639e43813f590439a4d) a little bit to support diskcrypto...

This is my working patch for 5.1.0 (applied to commit 1493bc01cfbf3fe0b7f4f639e43813f590439a4d), but I didn't have the time to adapt it to the new plugin-based branch (now in master):

Code:

diff --git a/OpenCL/inc_diskcrypto_modified_xts.cl b/OpenCL/inc_diskcrypto_modified_xts.cl
new file mode 100644
index 00000000..03cc5c58
--- /dev/null
+++ b/OpenCL/inc_diskcrypto_modified_xts.cl
@@ -0,0 +1,70 @@
+DECLSPEC void xts_mul2 (u32 *in, u32 *out)
+{
+  const u32 c = in[3] >> 31;
+
+  out[3] = (in[3] << 1) | (in[2] >> 31);
+  out[2] = (in[2] << 1) | (in[1] >> 31);
+  out[1] = (in[1] << 1) | (in[0] >> 31);
+  out[0] = (in[0] << 1);
+
+  out[0] ^= c * 0x87;
+}
+
+DECLSPEC void aes256_decrypt_xts (const u32 *ukey1, const u32 *ukey2, const u32 *in, u32 *out, u32 *S, u32 *T, u32 *ks, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+{
+  out[0] = in[0];
+  out[1] = in[1];
+  out[2] = in[2];
+  out[3] = in[3];
+
+  aes256_set_encrypt_key (ks, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4);
+  aes256_encrypt (ks, S, T, s_te0, s_te1, s_te2, s_te3, s_te4);
+
+  // skip four blocks (the starting position + 64 raw salt bytes that were replaced after encryption):
+
+  xts_mul2 (T, T);
+  xts_mul2 (T, T);
+  xts_mul2 (T, T);
+  xts_mul2 (T, T);
+
+  out[0] ^= T[0];
+  out[1] ^= T[1];
+  out[2] ^= T[2];
+  out[3] ^= T[3];
+
+  aes256_set_decrypt_key (ks, ukey1, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4);
+  aes256_decrypt (ks, out, out, s_td0, s_td1, s_td2, s_td3, s_td4);
+
+  out[0] ^= T[0];
+  out[1] ^= T[1];
+  out[2] ^= T[2];
+  out[3] ^= T[3];
+}
+
+DECLSPEC int decrypt_and_check (__global const u32 *encrypted_data, const u32 *ukey1, const u32 *ukey2, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
+{
+  u32 ks_aes[60];
+
+  u32 S[4] = { 1, 0, 0, 0 }; // this damn offset / sector ID, gosh damnit !!! found by lot of research by philsmd
+
+  u32 T_aes[4] = { 0 };
+
+  u32 data[4];
+
+  data[0] = encrypted_data[0];
+  data[1] = encrypted_data[1];
+  data[2] = encrypted_data[2];
+  data[3] = encrypted_data[3];
+
+  u32 out[4];
+
+  aes256_decrypt_xts (ukey1, ukey2, data, out, S, T_aes, ks_aes, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4);
+
+  if (out[0] != 0x50524344) return 0; // signature / magic: "DCRP"
+
+  if ((out[2] != 0x00040002) && (out[2] != 0x00050002)) return 0; // header version 0x0002 and flags either 0x04 or 0x05
+
+  if ((out[3] & 0xffff) != 0) return 0; // remaining 2 bytes of 0x00000004 / 0x00000005 => must be 0x0000
+
+  return 1;
+}
diff --git a/OpenCL/inc_hash_sha512.cl b/OpenCL/inc_hash_sha512.cl
index 869a5d99..2c099e2d 100644
--- a/OpenCL/inc_hash_sha512.cl
+++ b/OpenCL/inc_hash_sha512.cl
@@ -1710,6 +1710,123 @@ DECLSPEC void sha512_hmac_init_global_swap (sha512_hmac_ctx_t *ctx, __global con
   sha512_hmac_init_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7);
}

+DECLSPEC void sha512_hmac_init_global_utf16le_swap (sha512_hmac_ctx_t *ctx, __global const u32 *w, const int len)
+{
+  u32 w0[4];
+  u32 w1[4];
+  u32 w2[4];
+  u32 w3[4];
+  u32 w4[4];
+  u32 w5[4];
+  u32 w6[4];
+  u32 w7[4];
+
+  const int len_new = len * 2;
+
+  if (len_new > 128)
+  {
+    sha512_ctx_t tmp;
+
+    sha512_init (&tmp);
+
+    sha512_update_global_utf16le_swap (&tmp, w, len);
+
+    sha512_final (&tmp);
+
+    w0[0] = h32_from_64_S (tmp.h[0]);
+    w0[1] = l32_from_64_S (tmp.h[0]);
+    w0[2] = h32_from_64_S (tmp.h[1]);
+    w0[3] = l32_from_64_S (tmp.h[1]);
+    w1[0] = h32_from_64_S (tmp.h[2]);
+    w1[1] = l32_from_64_S (tmp.h[2]);
+    w1[2] = h32_from_64_S (tmp.h[3]);
+    w1[3] = l32_from_64_S (tmp.h[3]);
+    w2[0] = h32_from_64_S (tmp.h[4]);
+    w2[1] = l32_from_64_S (tmp.h[4]);
+    w2[2] = h32_from_64_S (tmp.h[5]);
+    w2[3] = l32_from_64_S (tmp.h[5]);
+    w3[0] = h32_from_64_S (tmp.h[6]);
+    w3[1] = l32_from_64_S (tmp.h[6]);
+    w3[2] = h32_from_64_S (tmp.h[7]);
+    w3[3] = l32_from_64_S (tmp.h[7]);
+    w4[0] = 0;
+    w4[1] = 0;
+    w4[2] = 0;
+    w4[3] = 0;
+    w5[0] = 0;
+    w5[1] = 0;
+    w5[2] = 0;
+    w5[3] = 0;
+    w6[0] = 0;
+    w6[1] = 0;
+    w6[2] = 0;
+    w6[3] = 0;
+    w7[0] = 0;
+    w7[1] = 0;
+    w7[2] = 0;
+    w7[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];
+
+    make_utf16le_S (w3, w6, w7);
+    make_utf16le_S (w2, w4, w5);
+    make_utf16le_S (w1, w2, w3);
+    make_utf16le_S (w0, w0, w1);
+
+    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_S (w1[0]);
+    w1[1] = swap32_S (w1[1]);
+    w1[2] = swap32_S (w1[2]);
+    w1[3] = swap32_S (w1[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_S (w3[0]);
+    w3[1] = swap32_S (w3[1]);
+    w3[2] = swap32_S (w3[2]);
+    w3[3] = swap32_S (w3[3]);
+    w4[0] = swap32_S (w4[0]);
+    w4[1] = swap32_S (w4[1]);
+    w4[2] = swap32_S (w4[2]);
+    w4[3] = swap32_S (w4[3]);
+    w5[0] = swap32_S (w5[0]);
+    w5[1] = swap32_S (w5[1]);
+    w5[2] = swap32_S (w5[2]);
+    w5[3] = swap32_S (w5[3]);
+    w6[0] = swap32_S (w6[0]);
+    w6[1] = swap32_S (w6[1]);
+    w6[2] = swap32_S (w6[2]);
+    w6[3] = swap32_S (w6[3]);
+    w7[0] = swap32_S (w7[0]);
+    w7[1] = swap32_S (w7[1]);
+    w7[2] = swap32_S (w7[2]);
+    w7[3] = swap32_S (w7[3]);
+  }
+
+  sha512_hmac_init_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7);
+}
+
DECLSPEC void sha512_hmac_update_128 (sha512_hmac_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, u32 *w4, u32 *w5, u32 *w6, u32 *w7, const int len)
{
   sha512_update_128 (&ctx->ipad, w0, w1, w2, w3, w4, w5, w6, w7, len);
diff --git a/OpenCL/m98765-pure.cl b/OpenCL/m98765-pure.cl
new file mode 100644
index 00000000..628856db
--- /dev/null
+++ b/OpenCL/m98765-pure.cl
@@ -0,0 +1,398 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+#define NEW_SIMD_CODE
+
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
+#include "inc_simd.cl"
+#include "inc_hash_sha512.cl"
+
+#include "inc_cipher_aes.cl"
+#include "inc_diskcrypto_modified_xts.cl"
+
+DECLSPEC void hmac_sha512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u32x *w4, u32x *w5, u32x *w6, u32x *w7, u64x *ipad, u64x *opad, u64x *digest)
+{
+  digest[0] = ipad[0];
+  digest[1] = ipad[1];
+  digest[2] = ipad[2];
+  digest[3] = ipad[3];
+  digest[4] = ipad[4];
+  digest[5] = ipad[5];
+  digest[6] = ipad[6];
+  digest[7] = ipad[7];
+
+  sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest);
+
+  w0[0] = h32_from_64 (digest[0]);
+  w0[1] = l32_from_64 (digest[0]);
+  w0[2] = h32_from_64 (digest[1]);
+  w0[3] = l32_from_64 (digest[1]);
+  w1[0] = h32_from_64 (digest[2]);
+  w1[1] = l32_from_64 (digest[2]);
+  w1[2] = h32_from_64 (digest[3]);
+  w1[3] = l32_from_64 (digest[3]);
+  w2[0] = h32_from_64 (digest[4]);
+  w2[1] = l32_from_64 (digest[4]);
+  w2[2] = h32_from_64 (digest[5]);
+  w2[3] = l32_from_64 (digest[5]);
+  w3[0] = h32_from_64 (digest[6]);
+  w3[1] = l32_from_64 (digest[6]);
+  w3[2] = h32_from_64 (digest[7]);
+  w3[3] = l32_from_64 (digest[7]);
+  w4[0] = 0x80000000;
+  w4[1] = 0;
+  w4[2] = 0;
+  w4[3] = 0;
+  w5[0] = 0;
+  w5[1] = 0;
+  w5[2] = 0;
+  w5[3] = 0;
+  w6[0] = 0;
+  w6[1] = 0;
+  w6[2] = 0;
+  w6[3] = 0;
+  w7[0] = 0;
+  w7[1] = 0;
+  w7[2] = 0;
+  w7[3] = (128 + 64) * 8;
+
+  digest[0] = opad[0];
+  digest[1] = opad[1];
+  digest[2] = opad[2];
+  digest[3] = opad[3];
+  digest[4] = opad[4];
+  digest[5] = opad[5];
+  digest[6] = opad[6];
+  digest[7] = opad[7];
+
+  sha512_transform_vector (w0, w1, w2, w3, w4, w5, w6, w7, digest);
+}
+
+__kernel void m98765_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, pbkdf2_sha512_t))
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  sha512_hmac_ctx_t sha512_hmac_ctx;
+
+  // Attention: we limit password length here for better speed, it's not needed, you can use the 255 bytes limit as usual, aswell
+  sha512_hmac_init_global_utf16le_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 63); // we limit the password length too
+
+  tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
+  tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
+  tmps[gid].ipad[2] = sha512_hmac_ctx.ipad.h[2];
+  tmps[gid].ipad[3] = sha512_hmac_ctx.ipad.h[3];
+  tmps[gid].ipad[4] = sha512_hmac_ctx.ipad.h[4];
+  tmps[gid].ipad[5] = sha512_hmac_ctx.ipad.h[5];
+  tmps[gid].ipad[6] = sha512_hmac_ctx.ipad.h[6];
+  tmps[gid].ipad[7] = sha512_hmac_ctx.ipad.h[7];
+
+  tmps[gid].opad[0] = sha512_hmac_ctx.opad.h[0];
+  tmps[gid].opad[1] = sha512_hmac_ctx.opad.h[1];
+  tmps[gid].opad[2] = sha512_hmac_ctx.opad.h[2];
+  tmps[gid].opad[3] = sha512_hmac_ctx.opad.h[3];
+  tmps[gid].opad[4] = sha512_hmac_ctx.opad.h[4];
+  tmps[gid].opad[5] = sha512_hmac_ctx.opad.h[5];
+  tmps[gid].opad[6] = sha512_hmac_ctx.opad.h[6];
+  tmps[gid].opad[7] = sha512_hmac_ctx.opad.h[7];
+
+  sha512_hmac_update_global_swap (&sha512_hmac_ctx, esalt_bufs[digests_offset].salt_buf, salt_bufs[salt_pos].salt_len);
+
+  for (u32 i = 0, j = 1; i < 8; i += 8, j += 1)
+  {
+    sha512_hmac_ctx_t sha512_hmac_ctx2 = sha512_hmac_ctx;
+
+    u32 w0[4];
+    u32 w1[4];
+    u32 w2[4];
+    u32 w3[4];
+    u32 w4[4];
+    u32 w5[4];
+    u32 w6[4];
+    u32 w7[4];
+
+    w0[0] = j;
+    w0[1] = 0;
+    w0[2] = 0;
+    w0[3] = 0;
+    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;
+    w4[0] = 0;
+    w4[1] = 0;
+    w4[2] = 0;
+    w4[3] = 0;
+    w5[0] = 0;
+    w5[1] = 0;
+    w5[2] = 0;
+    w5[3] = 0;
+    w6[0] = 0;
+    w6[1] = 0;
+    w6[2] = 0;
+    w6[3] = 0;
+    w7[0] = 0;
+    w7[1] = 0;
+    w7[2] = 0;
+    w7[3] = 0;
+
+    sha512_hmac_update_128 (&sha512_hmac_ctx2, w0, w1, w2, w3, w4, w5, w6, w7, 4);
+
+    sha512_hmac_final (&sha512_hmac_ctx2);
+
+    tmps[gid].dgst[i + 0] = sha512_hmac_ctx2.opad.h[0];
+    tmps[gid].dgst[i + 1] = sha512_hmac_ctx2.opad.h[1];
+    tmps[gid].dgst[i + 2] = sha512_hmac_ctx2.opad.h[2];
+    tmps[gid].dgst[i + 3] = sha512_hmac_ctx2.opad.h[3];
+    tmps[gid].dgst[i + 4] = sha512_hmac_ctx2.opad.h[4];
+    tmps[gid].dgst[i + 5] = sha512_hmac_ctx2.opad.h[5];
+    tmps[gid].dgst[i + 6] = sha512_hmac_ctx2.opad.h[6];
+    tmps[gid].dgst[i + 7] = sha512_hmac_ctx2.opad.h[7];
+
+    tmps[gid].out[i + 0] = tmps[gid].dgst[i + 0];
+    tmps[gid].out[i + 1] = tmps[gid].dgst[i + 1];
+    tmps[gid].out[i + 2] = tmps[gid].dgst[i + 2];
+    tmps[gid].out[i + 3] = tmps[gid].dgst[i + 3];
+    tmps[gid].out[i + 4] = tmps[gid].dgst[i + 4];
+    tmps[gid].out[i + 5] = tmps[gid].dgst[i + 5];
+    tmps[gid].out[i + 6] = tmps[gid].dgst[i + 6];
+    tmps[gid].out[i + 7] = tmps[gid].dgst[i + 7];
+  }
+}
+
+__kernel void m98765_loop (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, pbkdf2_sha512_t))
+{
+  const u64 gid = get_global_id (0);
+
+  if ((gid * VECT_SIZE) >= gid_max) return;
+
+  u64x ipad[8];
+  u64x opad[8];
+
+  ipad[0] = pack64v (tmps, ipad, gid, 0);
+  ipad[1] = pack64v (tmps, ipad, gid, 1);
+  ipad[2] = pack64v (tmps, ipad, gid, 2);
+  ipad[3] = pack64v (tmps, ipad, gid, 3);
+  ipad[4] = pack64v (tmps, ipad, gid, 4);
+  ipad[5] = pack64v (tmps, ipad, gid, 5);
+  ipad[6] = pack64v (tmps, ipad, gid, 6);
+  ipad[7] = pack64v (tmps, ipad, gid, 7);
+
+  opad[0] = pack64v (tmps, opad, gid, 0);
+  opad[1] = pack64v (tmps, opad, gid, 1);
+  opad[2] = pack64v (tmps, opad, gid, 2);
+  opad[3] = pack64v (tmps, opad, gid, 3);
+  opad[4] = pack64v (tmps, opad, gid, 4);
+  opad[5] = pack64v (tmps, opad, gid, 5);
+  opad[6] = pack64v (tmps, opad, gid, 6);
+  opad[7] = pack64v (tmps, opad, gid, 7);
+
+  for (u32 i = 0; i < 8; i += 8)
+  {
+    u64x dgst[8];
+    u64x out[8];
+
+    dgst[0] = pack64v (tmps, dgst, gid, i + 0);
+    dgst[1] = pack64v (tmps, dgst, gid, i + 1);
+    dgst[2] = pack64v (tmps, dgst, gid, i + 2);
+    dgst[3] = pack64v (tmps, dgst, gid, i + 3);
+    dgst[4] = pack64v (tmps, dgst, gid, i + 4);
+    dgst[5] = pack64v (tmps, dgst, gid, i + 5);
+    dgst[6] = pack64v (tmps, dgst, gid, i + 6);
+    dgst[7] = pack64v (tmps, dgst, gid, i + 7);
+
+    out[0] = pack64v (tmps, out, gid, i + 0);
+    out[1] = pack64v (tmps, out, gid, i + 1);
+    out[2] = pack64v (tmps, out, gid, i + 2);
+    out[3] = pack64v (tmps, out, gid, i + 3);
+    out[4] = pack64v (tmps, out, gid, i + 4);
+    out[5] = pack64v (tmps, out, gid, i + 5);
+    out[6] = pack64v (tmps, out, gid, i + 6);
+    out[7] = pack64v (tmps, out, gid, i + 7);
+
+    for (u32 j = 0; j < loop_cnt; j++)
+    {
+      u32x w0[4];
+      u32x w1[4];
+      u32x w2[4];
+      u32x w3[4];
+      u32x w4[4];
+      u32x w5[4];
+      u32x w6[4];
+      u32x w7[4];
+
+      w0[0] = h32_from_64 (dgst[0]);
+      w0[1] = l32_from_64 (dgst[0]);
+      w0[2] = h32_from_64 (dgst[1]);
+      w0[3] = l32_from_64 (dgst[1]);
+      w1[0] = h32_from_64 (dgst[2]);
+      w1[1] = l32_from_64 (dgst[2]);
+      w1[2] = h32_from_64 (dgst[3]);
+      w1[3] = l32_from_64 (dgst[3]);
+      w2[0] = h32_from_64 (dgst[4]);
+      w2[1] = l32_from_64 (dgst[4]);
+      w2[2] = h32_from_64 (dgst[5]);
+      w2[3] = l32_from_64 (dgst[5]);
+      w3[0] = h32_from_64 (dgst[6]);
+      w3[1] = l32_from_64 (dgst[6]);
+      w3[2] = h32_from_64 (dgst[7]);
+      w3[3] = l32_from_64 (dgst[7]);
+      w4[0] = 0x80000000;
+      w4[1] = 0;
+      w4[2] = 0;
+      w4[3] = 0;
+      w5[0] = 0;
+      w5[1] = 0;
+      w5[2] = 0;
+      w5[3] = 0;
+      w6[0] = 0;
+      w6[1] = 0;
+      w6[2] = 0;
+      w6[3] = 0;
+      w7[0] = 0;
+      w7[1] = 0;
+      w7[2] = 0;
+      w7[3] = (128 + 64) * 8;
+
+      hmac_sha512_run_V (w0, w1, w2, w3, w4, w5, w6, w7, ipad, opad, dgst);
+
+      out[0] ^= dgst[0];
+      out[1] ^= dgst[1];
+      out[2] ^= dgst[2];
+      out[3] ^= dgst[3];
+      out[4] ^= dgst[4];
+      out[5] ^= dgst[5];
+      out[6] ^= dgst[6];
+      out[7] ^= dgst[7];
+    }
+
+    unpack64v (tmps, dgst, gid, i + 0, dgst[0]);
+    unpack64v (tmps, dgst, gid, i + 1, dgst[1]);
+    unpack64v (tmps, dgst, gid, i + 2, dgst[2]);
+    unpack64v (tmps, dgst, gid, i + 3, dgst[3]);
+    unpack64v (tmps, dgst, gid, i + 4, dgst[4]);
+    unpack64v (tmps, dgst, gid, i + 5, dgst[5]);
+    unpack64v (tmps, dgst, gid, i + 6, dgst[6]);
+    unpack64v (tmps, dgst, gid, i + 7, dgst[7]);
+
+    unpack64v (tmps, out, gid, i + 0, out[0]);
+    unpack64v (tmps, out, gid, i + 1, out[1]);
+    unpack64v (tmps, out, gid, i + 2, out[2]);
+    unpack64v (tmps, out, gid, i + 3, out[3]);
+    unpack64v (tmps, out, gid, i + 4, out[4]);
+    unpack64v (tmps, out, gid, i + 5, out[5]);
+    unpack64v (tmps, out, gid, i + 6, out[6]);
+    unpack64v (tmps, out, gid, i + 7, out[7]);
+  }
+}
+
+__kernel void m98765_comp (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, pbkdf2_sha512_t))
+{
+  const u64 gid = get_global_id (0);
+  const u64 lid = get_local_id (0);
+  const u64 lsz = get_local_size (0);
+
+  /**
+   * 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 (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
+  {
+    s_td0[i] = td0[i];
+    s_td1[i] = td1[i];
+    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;
+
+  u32 ukey1[8];
+
+  ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0]));
+  ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0]));
+  ukey1[2] = swap32_S (h32_from_64_S (tmps[gid].out[1]));
+  ukey1[3] = swap32_S (l32_from_64_S (tmps[gid].out[1]));
+  ukey1[4] = swap32_S (h32_from_64_S (tmps[gid].out[2]));
+  ukey1[5] = swap32_S (l32_from_64_S (tmps[gid].out[2]));
+  ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3]));
+  ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3]));
+
+  u32 ukey2[8];
+
+  ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4]));
+  ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4]));
+  ukey2[2] = swap32_S (h32_from_64_S (tmps[gid].out[5]));
+  ukey2[3] = swap32_S (l32_from_64_S (tmps[gid].out[5]));
+  ukey2[4] = swap32_S (h32_from_64_S (tmps[gid].out[6]));
+  ukey2[5] = swap32_S (l32_from_64_S (tmps[gid].out[6]));
+  ukey2[6] = swap32_S (h32_from_64_S (tmps[gid].out[7]));
+  ukey2[7] = swap32_S (l32_from_64_S (tmps[gid].out[7]));
+
+  #define il_pos 0
+
+  if (decrypt_and_check (digests_buf[digests_offset].digest_buf, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
+  {
+    if (atomic_inc (&hashes_shown[digests_offset]) == 0)
+    {
+      mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset, gid, il_pos);
+    }
+  }
+}
diff --git a/diskcrypto2hashcat.py b/diskcrypto2hashcat.py
new file mode 100755
index 00000000..9b816093
--- /dev/null
+++ b/diskcrypto2hashcat.py
@@ -0,0 +1,56 @@
+#!/usr/bin/env python2
+
+#
+# Author: philsmd
+# Date released: 18th December 2018
+# License: public domains, credits go to philsmd of team hashcat
+#
+
+import sys
+
+#
+# Constants
+#
+
+SALT_LENGTH     = 64
+MIN_DATA_LEN    = (4 + 1) * 16 # = 80 i.e. 5 AES blocks (the first 4 are garbage, because of the replaced salt)
+HASH_SIGNATURE  = "$diskcrypto$"
+
+#
+# Start
+#
+
+# check command line arguments
+
+if len (sys.argv) != (1 + 1):
+  sys.stderr.write ("Usage: " + sys.argv[0] + " volume_header.txt\n")
+  sys.exit (1)
+
+# open encrypted volume header file
+
+data_file_name = sys.argv[1]
+
+data = ""
+
+try:
+  data_file = open (data_file_name, 'rb')
+
+  if data_file:
+    data = data_file.read (MIN_DATA_LEN)
+    data_file.close ()
+except:
+  sys.stderr.write ("[Error] Could not read from the file '" + data_file_name + "'\n")
+  sys.exit (1)
+
+if len (data) < MIN_DATA_LEN:
+  sys.stderr.write ("[Error] The file '" + data_file_name + "' must contain at least " + str (MIN_DATA_LEN) + " bytes\n")
+  sys.exit (1)
+
+# salt is part of the data/volume header (first 64 bytes)
+
+salt  = data[:SALT_LENGTH]
+block = data[SALT_LENGTH:] # only next 16 bytes (one block) of data needed
+
+# the final output:
+
+print HASH_SIGNATURE + salt.encode ('hex') + "$" + block.encode ('hex')
diff --git a/diskcrypto_crack.py b/diskcrypto_crack.py
new file mode 100755
index 00000000..60081bf7
--- /dev/null
+++ b/diskcrypto_crack.py
@@ -0,0 +1,152 @@
+#!/usr/bin/env python2
+
+#
+# Author: philsmd
+# Date released: 17th December 2018
+# License: public domains, credits go to philsmd of team hashcat
+#
+
+# How to install it on a linux system
+
+# sudo pip   install pycryptoplus
+# sudo pip uninstall pycryptodome
+
+import sys
+import time
+
+from CryptoPlus.Cipher        import AES
+from Crypto.Hash              import SHA512
+from CryptoPlus.Hash.pypbkdf2 import PBKDF2
+
+#
+# Constants
+#
+
+NUM_ITERATIONS  = 1000
+SALT_LENGTH     = 64
+MIN_DATA_LEN    = (4 + 1) * 16 # = 80 i.e. 5 AES blocks (the first 4 are garbage, because of the replaced salt)
+
+# bytes that we use to verify the data !
+
+SIGNATURE_START = 64 # the bytes after SALT_LENGTH (therefore SIGNATURE_START == SALT_LENGTH)
+SIGNATURE_END   = 68
+SIGNATURE       = "DCRP"
+
+KNOWN_HEADER_VERSION_START = 72
+KNOWN_HEADER_VERSION_END   = 74
+KNOWN_HEADER_VERSION       = "\x02\x00"
+
+KNOWN_FLAGS_START = 74
+KNOWN_FLAGS_END   = 78
+KNOWN_FLAG1       = "\x05\x00\x00\x00"
+KNOWN_FLAG2       = "\x04\x00\x00\x00"
+
+PROGRESS_UPDATE = 1000 # print a status update after every x password candidates
+
+#
+# Start
+#
+
+# check command line arguments
+
+if len (sys.argv) != (1 + 2):
+  sys.stderr.write ("Usage: " + sys.argv[0] + " volume_header.txt dict.txt\n")
+  sys.exit (1)
+
+
+# open encrypted volume header file
+
+data_file_name = sys.argv[1]
+
+data = ""
+
+try:
+  data_file = open (data_file_name, 'rb')
+
+  if data_file:
+    data = data_file.read (MIN_DATA_LEN)
+    data_file.close ()
+except:
+  sys.stderr.write ("[Error] Could not read from the file '" + data_file_name + "'\n")
+  sys.exit (1)
+
+if len (data) < MIN_DATA_LEN:
+  sys.stderr.write ("[Error] The file '" + data_file_name + "' must contain at least " + str (MIN_DATA_LEN) + " bytes\n")
+  sys.exit (1)
+
+# salt is part of the data/volume header (first 64 bytes)
+
+salt = data[:SALT_LENGTH]
+
+
+# open file containing the list of password candidates (dict)
+
+dict_file_name = sys.argv[2]
+
+try:
+  dict_file = open (dict_file_name, "rb")
+except:
+  sys.stderr.write ("[Error] Could not read from the file '" + dict_file_name + "'\n")
+
+
+# main loop
+
+progress_num = 0
+progress_mod = 0
+
+last_time = time.time ()
+
+for password in dict_file:
+  password = password.rstrip ()
+
+  if len (password) < 1:
+    continue
+
+  # generate the AES decryption key:
+
+  try:
+      password = password.encode ('UTF-16LE')
+  except:
+     pass
+
+  aes_key = PBKDF2 (password, salt, NUM_ITERATIONS, SHA512)
+
+  aes_key1 = aes_key.read (32) # "main" key
+  aes_key2 = aes_key.read (32) # tweak  key
+
+  xts_key = (aes_key1, aes_key2)
+
+  cipher = AES.new (xts_key, AES.MODE_XTS)
+
+  # decrypt the data and verify:
+
+  sequence = '01'.decode ('hex') # this was the main problem ! gosh dammit
+
+  data_decrypted = cipher.decrypt (data, sequence)
+
+  if data_decrypted[SIGNATURE_START:SIGNATURE_END] == SIGNATURE:
+    if data_decrypted[KNOWN_HEADER_VERSION_START:KNOWN_HEADER_VERSION_END] == KNOWN_HEADER_VERSION:
+      flags = data_decrypted[KNOWN_FLAGS_START:KNOWN_FLAGS_END]
+      if flags == KNOWN_FLAG1 or flags == KNOWN_FLAG2:
+        print "[Success] Found the password: '" + password + "' (flags: " + flags.encode ('hex') + ")"
+
+        dict_file.close ()
+        sys.exit (0) # if we expect many false positives we should/could continue here instead
+
+  # progress only (cosmetic):
+
+  if progress_mod == PROGRESS_UPDATE:
+    progress_num += 1
+
+    this_time = time.time ()
+    sys.stderr.write ("[Progress] Unsuccessful attempts: " + str (progress_mod * progress_num) + " (" + str (int (PROGRESS_UPDATE / (this_time - last_time))) + " H/s)\n")
+    last_time = this_time
+
+    progress_mod = 0
+
+  else:
+    progress_mod += 1
+
+dict_file.close ()
+
+sys.exit (1)
diff --git a/docs/changes.txt b/docs/changes.txt
index 24a3b7dd..975b120e 100644
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -7,6 +7,7 @@
- Added hash-mode 18400 (Open Document Format (ODF) 1.2 (SHA-256, AES))
- Added hash-mode 18500 sha1(md5(md5($pass)))
- Added hash-mode 18600 (Open Document Format (ODF) 1.1 (SHA-1, Blowfish))
+- Added hash-mode 98765 (DisKCrypto)

##
## Bugs
diff --git a/include/interface.h b/include/interface.h
index e8655eb6..519f59c2 100644
--- a/include/interface.h
+++ b/include/interface.h
@@ -1189,6 +1189,7 @@ typedef enum hash_type
   HASH_TYPE_KRB5ASREP           = 71,
   HASH_TYPE_ODF12               = 72,
   HASH_TYPE_ODF11               = 73,
+  HASH_TYPE_DISKCRYPTO          = 74,

} hash_type_t;

@@ -1419,6 +1420,7 @@ typedef enum kern_type
   KERN_TYPE_ODF12                   = 18400,
   KERN_TYPE_SHA1_DOUBLE_MD5         = 18500,
   KERN_TYPE_ODF11                   = 18600,
+  KERN_TYPE_DISKCRYPTO              = 98765,
   KERN_TYPE_PLAINTEXT               = 99999,

} kern_type_t;
@@ -1498,6 +1500,7 @@ typedef enum rounds_count
    ROUNDS_APPLE_SECURE_NOTES = 20000,
    ROUNDS_ETHEREUM_PRESALE   = 2000 - 1,
    ROUNDS_ANSIBLE_VAULT      = 10000,
+   ROUNDS_DISKCRYPTO         = 1000,
    ROUNDS_STDOUT             = 0

} rounds_count_t;
diff --git a/src/interface.c b/src/interface.c
index 793d3155..5a59c5ba 100644
--- a/src/interface.c
+++ b/src/interface.c
@@ -302,6 +302,7 @@ static const char *ST_HASH_18300 = "$fvde$2$16$58778104701476542047675521040224$
static const char *ST_HASH_18400 = "$odf$*1*1*100000*32*751854d8b90731ce0579f96bea6f0d4ac2fb2f546b31f1b6af9a5f66952a0bf4*16*2185a966155baa9e2fb597298febecbc*16*c18eaae34bcbbe9119be017fe5f8b52d*0*051e0f1ce0e866f2b771029e03a6c7119aad132af54c4e45824f16f61f357a40407ab82744fe6370c7b2346075fcd4c2e58ab244411b3ab1d532a46e2321599ef13c3d3472fc2f14d480d8c33215e473da67f90540279d3ef1f62dde314fa222796046e496c951235ddf88aa754620b7810d22ebc8835c90dce9276946f52b8ea7d95d2f86e4cc725366a8b3edacc2ce88518e535991a5f84d5ea8795dc02bfb731b5f202ecaf7d4b245d928c4248709fcdf3fba2acf1a08be0c1eee7dbeda07e8c3a6983565635e99952b8ad79d31c965f245ae90b5cc3dba6387898c66fa35cad9ac9595c41b62e68efcdd73185b38e220cf004269b77ec6974474b03b7569afc3b503a2bf8b2d035756f3f4cb880d9ba815e5c944508a0bde214076c35bf0e0814a96d21ccaa744c9056948ed935209f5c7933841d2ede3d28dd84da89d477d4a0041ce6d8ddab891d929340db6daa921d69b46fd5aee306d0bcef88c38acbb495d0466df7e2f744e3d10201081215c02db5dd479a4cda15a3338969c7baec9d3d2c378a8dd30449319b149dc3b4e7f00996a59fcb5f243d0df2cbaf749241033f7865aefa960adfeb8ebf205b270f90b1f82c34f80d5a8a0db7aec89972a32f5daa2a73c5895d1fced01b3ab8e576bd2630eff01cad97781f4966d4b528e1b15f011f28ae907a352073c96b203adc7742d2b79b2e2f440b17e7856ae119e08d15d8bdf951f6d4a3f9b516da2d9a8f9dd93488f8e0119f3da19138ab787f0d7098a652cccd914aa0ff81d375bd6a5a165acc936f591639059287975cfc3ca4342e5f9501b3249a76d14e56d6d56b319e036bc0449ac7b5afa24ffbea11babed8183edf8d4fdca1c3f0d23bfd4a02797627d556634f1a9304e03737604bd86f6b5a26aa687d6df73383e0f7dfe62a131e8dbb8c3f4f13d24857dd29d76984eac6c45df7428fc79323ffa1f4e7962d705df74320141ed1f16d1ad483b872168df60315ffadbfa1b7f4afaed8a0017421bf5e05348cb5c707a5e852d6fee6077ec1c33bc707bcd97b7701ee05a03d6fa78b0d31c8c97ea16e0edf434961bd5cc7cbb7eb2553730f0405c9bd21cee09b3f7c1bc57779fdfc15f3935985737a1b522004c4436b631a39a66e8577a03f5020e6aa41952c0662c8c57f66caa483b47af38b8cb5d457245fd3241749e17433e6f929233e8862d7c584111b1991b2d6e94278e7e6e1908cee5a83d94c78b75a84a695d25aeb9fdde72174fe6dd75e8d406671f44892a385a4a1e249f61ebc993e985607423a0a5742e668d52c1ebf5cecae7c2b7908f4627b92ec49354a9ccff8cb5763ad074a00e65a485a41bf4c25ce7e6fae49358a58547b1c0ca79713e297310c0a367c3de196f1dd685ca4be643bdf1e4f6b034211d020557e37a3b6614d061010b4a3416b6b279728c245d3322";
static const char *ST_HASH_18500 = "888a2ffcb3854fba0321110c5d0d434ad1aa2880";
static const char *ST_HASH_18600 = "$odf$*0*0*1024*16*bff753835f4ea15644b8a2f8e4b5be3d147b9576*8*ee371da34333b69d*16*a902eff54a4d782a26a899a31f97bef4*0*dae7e41fbc3a500d3ce152edd8876c4f38fb17d673ee2ac44ef1e0e283622cd2ae298a82d8d98f2ea737247881fc353e73a2f535c6e13e0cdc60821c1a61c53a4b0c46ff3a3b355d7b793fad50de15999fc7c1194321d1c54316c3806956c4a3ade7daabb912a2a36398eba883af088b3cb69b43365d9ba9fce3fb0c1524f73947a7e9fc1bf3adb5f85a367035feacb5d97c578b037144c2793f34aa09dcd04bdaa455aee0d4c52fe377248611dd56f2bd4eb294673525db905f5d905a28dec0909348e6bf94bcebf03ddd61a48797cd5728ce6dbb71037b268f526e806401abcf495f6edd0b5d87118671ec690d4627f86a43e51c7f6d42a75a56eec51204d47e115e813ed4425c97b16b195e02ce776c185194b9de43ae89f356e29face016cb393d6fb93af8ea305d921d5592dd184051ac790b9b90266f52b8d53ce1cb1d762942d6d5bbd0e3821be21af9fa6874ba0c60e64f41d3e5b6caca1c53b575afdc5d8f6a3edbf874dbe009c6cb296466fe9637aed4aed8a43a95ea7d26b4090ad33d4ee7a83844b0893e8bc0f04944205fb9576cb5720f019028cd75ca9ac47b3e5fa231354d74135564df43b659cfaea7e195c4a896e0e0e0c85dc9ce3a9ce9ba552bc2a6dbac4901c19558818e1957ed72d78662bb5ba53475ca584371f1825ae0c92322a4404e63c2baad92665aac29b5c6f96e1e6338d48fb0aef4d0b686063974f58b839484f8dcf0a02537cba67a7d2c4de13125d74820cb07ec72782035af1ea6c4db61c77016d1c021b63c8b07adb4e8510f5c41bbc501f60f3dd16462399b52eb146787e38e700147c7aa23ac4d5d22d9d1c93e67a01c92a197d4765cbf8d56a862a1205abb450a182913a69b8d5334a59924f86fb3ccd0dcfe7426053e26ba26b57c05f38d85863fff1f81135b0366e8cd8680663ae8aaf7d005317b849d5e08be882708fa0d8d02d47e89150124b507c34845c922b95e62aa0b3fef218773d7aeb572c67b35ad8787f31ecc6e1846b673b8ba6172223176eabf0020b6aa3aa71405b40b2fc2127bf9741a103f1d8eca21bf27328cdf15153f2f223eff7b831a72ed8ecacf4ea8df4ea44f3a3921e5a88fb2cfa355ece0f05cbc88fdd1ecd368d6e3b2dfabd999e5b708f1bccaeebb296c9d7b76659967742fe966aa6871cbbffe710b0cd838c6e02e6eb608cb5c81d066b60b5b3604396331d97d4a2c4c2317406e48c9f5387a2c72511d1e6899bd450e9ca88d535755bcfddb53a6df118cd9cdc7d8b4b814f7bc17684d8e5975defaa25d06f410ed0724c16b8f69ec3869bc1f05c71483666968d1c04509875dadd72c6182733d564eb1a7d555dc34f6b817c5418626214d0b2c3901c5a46f5b20fddfdf9f71a7dfd75b9928778a3f65e1832dff22be973c2b259744d500a3027c2a2e08972eaaad4c5c4ec871";
+static const char *ST_HASH_98765 = "$diskcrypto$37f6252cf81f8049f68deb41de5becfb46851909e5d4f41f8f5da4c4dc830992c5e29905fa6e0cb755e42c6cfc0509a751a2a4f01fb884968c9fd18bc9007c1e$e7e67d1b7cf0e23ba82235517b93db8e";
static const char *ST_HASH_99999 = "hashcat";

static const char *OPTI_STR_OPTIMIZED_KERNEL     = "Optimized-Kernel";
@@ -567,6 +568,7 @@ static const char *HT_18300 = "Apple File System (APFS)";
static const char *HT_18400 = "Open Document Format (ODF) 1.2 (SHA-256, AES)";
static const char *HT_18500 = "sha1(md5(md5($pass)))";
static const char *HT_18600 = "Open Document Format (ODF) 1.1 (SHA-1, Blowfish)";
+static const char *HT_98765 = "DiskCryto";
static const char *HT_99999 = "Plaintext";

static const char *HT_00011 = "Joomla < 2.5.18";
@@ -724,6 +726,7 @@ static const char *SIGNATURE_ELECTRUM_WALLET    = "$electrum$";
static const char *SIGNATURE_FILEVAULT2         = "$fvde$";
static const char *SIGNATURE_ANSIBLE_VAULT      = "$ansible$";
static const char *SIGNATURE_APFS               = "$fvde$";
+static const char *SIGNATURE_DISKCRYPTO         = "$diskcrypto$";

/**
  * decoder / encoder
@@ -18909,6 +18912,85 @@ int apfs_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSE
   return (PARSER_OK);
}

+int diskcrypto_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig)
+{
+  u64 *digest = (u64 *) hash_buf->digest;
+
+  pbkdf2_sha512_t *pbkdf2_sha512 = (pbkdf2_sha512_t *) hash_buf->esalt;
+
+  salt_t *salt = hash_buf->salt;
+
+  token_t token;
+
+  token.token_cnt = 3;
+
+  token.signatures_cnt    = 1;
+  token.signatures_buf[0] = SIGNATURE_DISKCRYPTO;
+
+  token.len[0]     = 12;
+  token.attr[0]    = TOKEN_ATTR_FIXED_LENGTH
+                   | TOKEN_ATTR_VERIFY_SIGNATURE;
+
+  token.len_min[1] = 128;
+  token.len_max[1] = 128;
+  token.sep[1]     = '$';
+  token.attr[1]    = TOKEN_ATTR_VERIFY_LENGTH
+                   | TOKEN_ATTR_VERIFY_HEX;
+
+  token.len[2]     = 32;
+  token.attr[2]    = TOKEN_ATTR_FIXED_LENGTH
+                   | TOKEN_ATTR_VERIFY_HEX;
+
+  const int rc_tokenizer = input_tokenizer (input_buf, input_len, &token);
+
+  if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
+
+  const u8 *salt_buf       = token.buf[1];
+  const u8 *encrypted_data = token.buf[2];
+
+  // esalt
+
+  pbkdf2_sha512->salt_buf[ 0] = hex_to_u32 (&salt_buf[  0]);
+  pbkdf2_sha512->salt_buf[ 1] = hex_to_u32 (&salt_buf[  8]);
+  pbkdf2_sha512->salt_buf[ 2] = hex_to_u32 (&salt_buf[ 16]);
+  pbkdf2_sha512->salt_buf[ 3] = hex_to_u32 (&salt_buf[ 24]);
+  pbkdf2_sha512->salt_buf[ 4] = hex_to_u32 (&salt_buf[ 32]);
+  pbkdf2_sha512->salt_buf[ 5] = hex_to_u32 (&salt_buf[ 40]);
+  pbkdf2_sha512->salt_buf[ 6] = hex_to_u32 (&salt_buf[ 48]);
+  pbkdf2_sha512->salt_buf[ 7] = hex_to_u32 (&salt_buf[ 56]);
+  pbkdf2_sha512->salt_buf[ 8] = hex_to_u32 (&salt_buf[ 64]);
+  pbkdf2_sha512->salt_buf[ 9] = hex_to_u32 (&salt_buf[ 72]);
+  pbkdf2_sha512->salt_buf[10] = hex_to_u32 (&salt_buf[ 80]);
+  pbkdf2_sha512->salt_buf[11] = hex_to_u32 (&salt_buf[ 88]);
+  pbkdf2_sha512->salt_buf[12] = hex_to_u32 (&salt_buf[ 96]);
+  pbkdf2_sha512->salt_buf[13] = hex_to_u32 (&salt_buf[104]);
+  pbkdf2_sha512->salt_buf[14] = hex_to_u32 (&salt_buf[112]);
+  pbkdf2_sha512->salt_buf[15] = hex_to_u32 (&salt_buf[120]);
+
+  // salt param
+
+  salt->salt_len = 64;
+
+  salt->salt_iter = 1000 - 1;
+
+  // fake salt
+
+  salt->salt_buf[0] = pbkdf2_sha512->salt_buf[0];
+  salt->salt_buf[1] = pbkdf2_sha512->salt_buf[1];
+  salt->salt_buf[2] = pbkdf2_sha512->salt_buf[2];
+  salt->salt_buf[3] = pbkdf2_sha512->salt_buf[3];
+  salt->salt_buf[4] = salt->salt_iter;
+
+  /**
+   * digest
+   */
+
+  digest[0] = hex_to_u64 (&encrypted_data[ 0]);
+  digest[1] = hex_to_u64 (&encrypted_data[16]);
+
+  return (PARSER_OK);
+}
+
/**
  * hook functions
  */
@@ -19372,6 +19454,7 @@ const char *strhashtype (const u32 hash_mode)
     case 18400: return HT_18400;
     case 18500: return HT_18500;
     case 18600: return HT_18600;
+    case 98765: return HT_98765;
     case 99999: return HT_99999;
   }

@@ -23307,6 +23390,35 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
           odf11->encrypted_data[i + 7]);
     }
   }
+  else if (hash_mode == 98765)
+  {
+    pbkdf2_sha512_t *pbkdf2_sha512s = (pbkdf2_sha512_t *) esalts_buf;
+
+    pbkdf2_sha512_t *pbkdf2_sha512  = &pbkdf2_sha512s[digest_cur];
+
+    snprintf (out_buf, out_len, "%s%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x$%08x%08x%08x%08x",
+      SIGNATURE_DISKCRYPTO,
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 0]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 1]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 2]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 3]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 4]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 5]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 6]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 7]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 8]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[ 9]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[10]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[11]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[12]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[13]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[14]),
+      byte_swap_32 (pbkdf2_sha512->salt_buf[15]),
+      digest_buf[0],
+      digest_buf[1],
+      digest_buf[2],
+      digest_buf[3]);
+  }
   else if (hash_mode == 99999)
   {
     char *ptr = (char *) digest_buf;
@@ -28833,6 +28945,25 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
                  hashconfig->st_pass        = ST_PASS_HASHCAT_PLAIN;
                  break;

+    case 98765:  hashconfig->hash_type      = HASH_TYPE_DISKCRYPTO;
+                 hashconfig->salt_type      = SALT_TYPE_EMBEDDED;
+                 hashconfig->attack_exec    = ATTACK_EXEC_OUTSIDE_KERNEL;
+                 hashconfig->opts_type      = OPTS_TYPE_PT_UTF16LE
+                                            | OPTS_TYPE_PT_GENERATE_LE;
+                 hashconfig->kern_type      = KERN_TYPE_DISKCRYPTO;
+                 hashconfig->dgst_size      = DGST_SIZE_4_4;
+                 hashconfig->parse_func     = diskcrypto_parse_hash;
+                 hashconfig->opti_type      = OPTI_TYPE_ZERO_BYTE
+                                            | OPTI_TYPE_USES_BITS_64
+                                            | OPTI_TYPE_SLOW_HASH_SIMD_LOOP;
+                 hashconfig->dgst_pos0      = 0;
+                 hashconfig->dgst_pos1      = 1;
+                 hashconfig->dgst_pos2      = 2;
+                 hashconfig->dgst_pos3      = 3;
+                 hashconfig->st_hash        = ST_HASH_98765;
+                 hashconfig->st_pass        = ST_PASS_HASHCAT_PLAIN;
+                 break;
+
     case 99999:  hashconfig->hash_type      = HASH_TYPE_PLAINTEXT;
                  hashconfig->salt_type      = SALT_TYPE_NONE;
                  hashconfig->attack_exec    = ATTACK_EXEC_INSIDE_KERNEL;
@@ -29075,6 +29206,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
     case 18300: hashconfig->esalt_size = sizeof (apple_secure_notes_t); break;
     case 18400: hashconfig->esalt_size = sizeof (odf12_t);              break;
     case 18600: hashconfig->esalt_size = sizeof (odf11_t);              break;
+    case 98765: hashconfig->esalt_size = sizeof (pbkdf2_sha512_t);      break;
   }

   // hook_salt_size
@@ -29194,6 +29326,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
     case 18300: hashconfig->tmp_size = sizeof (apple_secure_notes_tmp_t); break;
     case 18400: hashconfig->tmp_size = sizeof (odf12_tmp_t);              break;
     case 18600: hashconfig->tmp_size = sizeof (odf11_tmp_t);              break;
+    case 98765: hashconfig->tmp_size = sizeof (pbkdf2_sha512_tmp_t);      break;
   };

   // hook_size
@@ -30271,6 +30404,8 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo
                  break;
     case 18600:  salt->salt_iter  = ROUNDS_OPENOFFICE - 1;
                  break;
+    case 98765:  salt->salt_iter  = ROUNDS_DISKCRYPTO - 1;
+                 break;
   }
}

diff --git a/src/usage.c b/src/usage.c
index 9795b5ec..c2ec1ddd 100644
--- a/src/usage.c
+++ b/src/usage.c
@@ -378,6 +378,7 @@ static const char *const USAGE_BIG[] =
   "  14600 | LUKS                                             | Full-Disk Encryption (FDE)",
   "  16700 | FileVault 2                                      | Full-Disk Encryption (FDE)",
   "  18300 | Apple File System (APFS)                         | Full-Disk Encryption (FDE)",
+  "  98765 | DiskCrypto                                       | Full-Disk Encryption (FDE)",
   "   9700 | MS Office <= 2003 $0/$1, MD5 + RC4               | Documents",
   "   9710 | MS Office <= 2003 $0/$1, MD5 + RC4, collider #1  | Documents",
   "   9720 | MS Office <= 2003 $0/$1, MD5 + RC4, collider #2  | Documents",

as you can see: $diskcrypto$37f6252cf81f8049f68deb41de5becfb46851909e5d4f41f8f5da4c4dc830992c5e29905fa6e0cb755e42c6cfc0509a751a2a4f01fb884968c9fd18bc9007c1e$e7e67d1b7cf0e23ba82235517b93db8e
is an example "hashcat" (the password, without quotes) hash

if anybody wants to adjust it a little bit to work with newest master, feel free to do so (but please mention it here first, such that we avoid double work)

The hash type -m 98765 would of course also need to be renamed to the latest free hash type (for instance you could ask atom about the next free integer that should be used, if you are unsure).

but again, the patch should work AS-IS, just not yet with the latest master... I made this experiment during heavy refactoring etc in hashcat master and therefore it wasn't committed yet I'm sorry about that. didn't have the time for adapting it

[Turysta]

I would highly ask you to add Diskcryptor support in the future.

Now I must generate dictoniary for DiskCryptor. Windows CMD speed - 20password per second

  dc.jpg (Size: 71.83 KB / Downloads: 15)

Greetings.

[philsmd]

The first step, i.e. implementing the basic DiskCryptor AES module + kernel + test was successfully done: see https://github.com/hashcat/hashcat/commi...ce891dae22 and https://twitter.com/hashcat/status/1114947180626944002

I'm not sure how important other encryption algos (like serpent) are and how they can be generated and tested.
Some help with basic research on how to generate the other ISO-based volumes / hashes would be nice. Are some encrytion types only allowed for ISO images? can you only select pure AES for full disk encryption (physical disk) ? etc...

---

update: we added further OpenCL kernel code to support the twofish/serpent variant with https://github.com/hashcat/hashcat/commi...29ca4456ec . It should work, but we weren't able to test the cascade of 2+ encryption algorithms, becase we have no such test vectors/examples yet. we are trying to generate some ourself, but it's not clear where this encrypt ISO image menu can be found that we see here: https://github.com/hashcat/hashcat/issue...-475846478

[Turysta]

Really well done

[philsmd]

with our latest commit https://github.com/hashcat/hashcat/commi...27f4e51210 we added the support for all the 7 combinations of aes/twofish/serpent (3 pure, 3 with 2 algos and 1 with 3 algos in a cascade).
Test hashes (that we used to verify our code, that were generated directly with diskcryptor CD image option) will be uploaded to https://hashcat.net/wiki/example_hashes soon. The format that we now use (the above POC is different) is identical to the jtr format and you can use diskcryptor2john.py from the john the ripper magnumripper branch (file is in the run/ folder) to extract hashes from the iso/images.

you can test new hash modes -m 19811 (if you are sure it's a single algo, no cascade), -m 19812 (if you are sure that max 2 algos are used, or also just a single algo), -m 19813 (if you think that all 3 algos could be used, or 2, or 1):
https://hashcat.net/beta/

[Turysta]

It works very well. Thank you very much for your time.
I wish you all good