Initial commit

OpenCL/m13100_a0.cl

@@ -0,0 +1,776 @@
+/**
+ * Authors......: Jens Steube <jens.steube@gmail.com>
+ * Authors......: Fist0urs   <eddy.maaalou@gmail.com>
+
+ * License.....: MIT
+ */
+
+#define _KRB5TGS_
+
+#include "include/constants.h"
+#include "include/kernel_vendor.h"
+
+#define DGST_R0 0
+#define DGST_R1 1
+#define DGST_R2 2
+#define DGST_R3 3
+
+#include "include/kernel_functions.c"
+#include "OpenCL/types_ocl.c"
+#include "OpenCL/common.c"
+#include "include/rp_kernel.h"
+#include "OpenCL/rp.c"
+
+typedef struct
+{
+  u8 S[256];
+
+  u32 wtf_its_faster;
+
+} RC4_KEY;
+
+static void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
+{
+  u8 tmp;
+
+  tmp           = rc4_key->S[i];
+  rc4_key->S[i] = rc4_key->S[j];
+  rc4_key->S[j] = tmp;
+}
+
+static void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
+{
+  u32 v = 0x03020100;
+  u32 a = 0x04040404;
+
+  __local u32 *ptr = (__local u32 *) rc4_key->S;
+
+  #pragma unroll
+  for (u32 i = 0; i < 64; i++)
+  {
+    *ptr++ = v; v += a;
+  }
+
+  u32 j = 0;
+
+  for (u32 i = 0; i < 16; i++)
+  {
+    u32 idx = i * 16;
+
+    u32 v;
+
+    v = data[0];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[1];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[2];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[3];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+  }
+}
+
+static u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 in[4], u32 out[4])
+{
+  #pragma unroll
+  for (u32 k = 0; k < 4; k++)
+  {
+    u32 xor4 = 0;
+
+    u8 idx;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] <<  0;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] <<  8;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] << 16;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] << 24;
+
+    out[k] = in[k] ^ xor4;
+  }
+
+  return j;
+}
+
+static int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32* edata2 /*, u32 edata2_len */)
+{
+  rc4_init_16 (rc4_key, data);
+
+  u32 out[8];
+
+  u8 j = 0;
+
+  /* 
+    8 first bytes are nonce, then ASN1 structs (DER encoding: type-length-data)
+
+    if length >= 128 bytes:
+        length is on 2 bytes and type is \x63\x82 (encode_krb5_enc_tkt_part) and data is an ASN1 sequence \x30\x82
+    else:
+        length is on 1 byte and type is \x63\x81 and data is an ASN1 sequence \x30\x81
+
+    next headers follow the same ASN1 "type-length-data" scheme
+  */
+
+  j = rc4_next_16 (rc4_key,  0, 0, edata2, out);
+
+  if (((out[2] & 0xff00ffff) != 0x30008163) && ((out[2] & 0x0000ffff) != 0x00008263)) return 0;
+
+  j = rc4_next_16 (rc4_key,  16, j, edata2 + 4, out + 4);
+
+  if (((out[4] & 0x00ffffff) != 0x00000503) && (out[4] != 0x050307A0)) return 0;
+
+  // TODO (or not): add RC4'ing of all edata2 then hmac-md5 and compare with 
+  // checksum to be definitely sure that this is the correct pass (even if 
+  // collisions must be very rare)
+
+  return 1;
+}
+
+static void md4_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
+{
+  u32 a = digest[0];
+  u32 b = digest[1];
+  u32 c = digest[2];
+  u32 d = digest[3];
+
+  MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
+
+  MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
+
+  MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
+
+  digest[0] += a;
+  digest[1] += b;
+  digest[2] += c;
+  digest[3] += d;
+}
+
+static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
+{
+  u32 a = digest[0];
+  u32 b = digest[1];
+  u32 c = digest[2];
+  u32 d = digest[3];
+
+  u32 w0_t = w0[0];
+  u32 w1_t = w0[1];
+  u32 w2_t = w0[2];
+  u32 w3_t = w0[3];
+  u32 w4_t = w1[0];
+  u32 w5_t = w1[1];
+  u32 w6_t = w1[2];
+  u32 w7_t = w1[3];
+  u32 w8_t = w2[0];
+  u32 w9_t = w2[1];
+  u32 wa_t = w2[2];
+  u32 wb_t = w2[3];
+  u32 wc_t = w3[0];
+  u32 wd_t = w3[1];
+  u32 we_t = w3[2];
+  u32 wf_t = w3[3];
+
+  MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
+
+  MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
+
+  MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
+
+  MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
+
+  digest[0] += a;
+  digest[1] += b;
+  digest[2] += c;
+  digest[3] += d;
+}
+
+static void hmac_md5_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4])
+{
+  w0[0] = w0[0] ^ 0x36363636;
+  w0[1] = w0[1] ^ 0x36363636;
+  w0[2] = w0[2] ^ 0x36363636;
+  w0[3] = w0[3] ^ 0x36363636;
+  w1[0] = w1[0] ^ 0x36363636;
+  w1[1] = w1[1] ^ 0x36363636;
+  w1[2] = w1[2] ^ 0x36363636;
+  w1[3] = w1[3] ^ 0x36363636;
+  w2[0] = w2[0] ^ 0x36363636;
+  w2[1] = w2[1] ^ 0x36363636;
+  w2[2] = w2[2] ^ 0x36363636;
+  w2[3] = w2[3] ^ 0x36363636;
+  w3[0] = w3[0] ^ 0x36363636;
+  w3[1] = w3[1] ^ 0x36363636;
+  w3[2] = w3[2] ^ 0x36363636;
+  w3[3] = w3[3] ^ 0x36363636;
+
+  ipad[0] = MD5M_A;
+  ipad[1] = MD5M_B;
+  ipad[2] = MD5M_C;
+  ipad[3] = MD5M_D;
+
+  md5_transform (w0, w1, w2, w3, ipad);
+
+  w0[0] = w0[0] ^ 0x6a6a6a6a;
+  w0[1] = w0[1] ^ 0x6a6a6a6a;
+  w0[2] = w0[2] ^ 0x6a6a6a6a;
+  w0[3] = w0[3] ^ 0x6a6a6a6a;
+  w1[0] = w1[0] ^ 0x6a6a6a6a;
+  w1[1] = w1[1] ^ 0x6a6a6a6a;
+  w1[2] = w1[2] ^ 0x6a6a6a6a;
+  w1[3] = w1[3] ^ 0x6a6a6a6a;
+  w2[0] = w2[0] ^ 0x6a6a6a6a;
+  w2[1] = w2[1] ^ 0x6a6a6a6a;
+  w2[2] = w2[2] ^ 0x6a6a6a6a;
+  w2[3] = w2[3] ^ 0x6a6a6a6a;
+  w3[0] = w3[0] ^ 0x6a6a6a6a;
+  w3[1] = w3[1] ^ 0x6a6a6a6a;
+  w3[2] = w3[2] ^ 0x6a6a6a6a;
+  w3[3] = w3[3] ^ 0x6a6a6a6a;
+
+  opad[0] = MD5M_A;
+  opad[1] = MD5M_B;
+  opad[2] = MD5M_C;
+  opad[3] = MD5M_D;
+
+  md5_transform (w0, w1, w2, w3, opad);
+}
+
+static void hmac_md5_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4], u32 digest[4])
+{
+  digest[0] = ipad[0];
+  digest[1] = ipad[1];
+  digest[2] = ipad[2];
+  digest[3] = ipad[3];
+
+  md5_transform (w0, w1, w2, w3, digest);
+
+  w0[0] = digest[0];
+  w0[1] = digest[1];
+  w0[2] = digest[2];
+  w0[3] = digest[3];
+  w1[0] = 0x80;
+  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] = (64 + 16) * 8;
+  w3[3] = 0;
+
+  digest[0] = opad[0];
+  digest[1] = opad[1];
+  digest[2] = opad[2];
+  digest[3] = opad[3];
+
+  md5_transform (w0, w1, w2, w3, digest);
+}
+
+
+static void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32 checksum[4], u32 digest[4])
+{
+  /**
+   * pads
+   */
+
+  u32 w0_t[4];
+  u32 w1_t[4];
+  u32 w2_t[4];
+  u32 w3_t[4];
+
+  w0_t[0] = w0[0];
+  w0_t[1] = w0[1];
+  w0_t[2] = w0[2];
+  w0_t[3] = w0[3];
+  w1_t[0] = w1[0];
+  w1_t[1] = w1[1];
+  w1_t[2] = w1[2];
+  w1_t[3] = w1[3];
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  // K=MD4(Little_indian(UNICODE(pwd))
+
+  append_0x80_2x4 (w0_t, w1_t, pw_len);
+
+  make_unicode (w1_t, w2_t, w3_t);
+  make_unicode (w0_t, w0_t, w1_t);
+
+  w3_t[2] = pw_len * 8 * 2;
+  w3_t[3] = 0;
+
+  digest[0] = MD4M_A;
+  digest[1] = MD4M_B;
+  digest[2] = MD4M_C;
+  digest[3] = MD4M_D;
+
+  md4_transform (w0_t, w1_t, w2_t, w3_t, digest);
+
+  // K1=MD5_HMAC(K,1); with 2 encoded as little indian on 4 bytes (02000000 in hexa);
+
+  w0_t[0] = digest[0];
+  w0_t[1] = digest[1];
+  w0_t[2] = digest[2];
+  w0_t[3] = digest[3];
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  u32 ipad[4];
+  u32 opad[4];
+
+  hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
+
+  w0_t[0] = 2;
+  w0_t[1] = 0x80;
+  w0_t[2] = 0;
+  w0_t[3] = 0;
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = (64 + 4) * 8;
+  w3_t[3] = 0;
+
+  hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
+
+  // K3=MD5_HMAC(K1,checksum);  
+
+  w0_t[0] = digest[0];
+  w0_t[1] = digest[1];
+  w0_t[2] = digest[2];
+  w0_t[3] = digest[3];
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
+
+  w0_t[0] = checksum[0];
+  w0_t[1] = checksum[1];
+  w0_t[2] = checksum[2];
+  w0_t[3] = checksum[3];
+  w1_t[0] = 0x80;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = (64 + 16) * 8;
+  w3_t[3] = 0;
+
+  hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
+}
+
+__kernel void m13100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * modifier
+   */
+
+  __local RC4_KEY rc4_keys[64];
+
+  const u32 lid = get_local_id (0);
+  const u32 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  u32 pw_buf0[4];
+
+  pw_buf0[0] = pws[gid].i[ 0];
+  pw_buf0[1] = pws[gid].i[ 1];
+  pw_buf0[2] = pws[gid].i[ 2];
+  pw_buf0[3] = pws[gid].i[ 3];
+
+  u32 pw_buf1[4];
+
+  pw_buf1[0] = pws[gid].i[ 4];
+  pw_buf1[1] = pws[gid].i[ 5];
+  pw_buf1[2] = pws[gid].i[ 6];
+  pw_buf1[3] = pws[gid].i[ 7];
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  /**
+   * salt
+   */
+
+  u32 checksum[4];
+
+  checksum[0] = krb5tgs_bufs[salt_pos].checksum[0];
+  checksum[1] = krb5tgs_bufs[salt_pos].checksum[1];
+  checksum[2] = krb5tgs_bufs[salt_pos].checksum[2];
+  checksum[3] = krb5tgs_bufs[salt_pos].checksum[3];
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
+  {
+    u32 w0[4];
+
+    w0[0] = pw_buf0[0];
+    w0[1] = pw_buf0[1];
+    w0[2] = pw_buf0[2];
+    w0[3] = pw_buf0[3];
+
+    u32 w1[4];
+
+    w1[0] = pw_buf1[0];
+    w1[1] = pw_buf1[1];
+    w1[2] = pw_buf1[2];
+    w1[3] = pw_buf1[3];
+
+    u32 w2[4];
+
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+
+    u32 w3[4];
+
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+
+    const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
+
+    /**
+     * kerberos
+     */
+
+    u32 digest[4];
+
+    kerb_prepare (w0, w1, out_len, checksum, digest);
+
+    u32 tmp[4];
+
+    tmp[0] = digest[0];
+    tmp[1] = digest[1];
+    tmp[2] = digest[2];
+    tmp[3] = digest[3];
+
+    if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2) == 1)
+    {
+      mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos);
+
+      d_return_buf[lid] = 1;
+    }
+  }
+}
+
+__kernel void m13100_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m13100_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m13100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * modifier
+   */
+
+  __local RC4_KEY rc4_keys[64];
+
+  const u32 lid = get_local_id (0);
+  const u32 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  u32 pw_buf0[4];
+
+  pw_buf0[0] = pws[gid].i[ 0];
+  pw_buf0[1] = pws[gid].i[ 1];
+  pw_buf0[2] = pws[gid].i[ 2];
+  pw_buf0[3] = pws[gid].i[ 3];
+
+  u32 pw_buf1[4];
+
+  pw_buf1[0] = pws[gid].i[ 4];
+  pw_buf1[1] = pws[gid].i[ 5];
+  pw_buf1[2] = pws[gid].i[ 6];
+  pw_buf1[3] = pws[gid].i[ 7];
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  /**
+   * salt
+   */
+
+  u32 checksum[4];
+
+  checksum[0] = krb5tgs_bufs[salt_pos].checksum[0];
+  checksum[1] = krb5tgs_bufs[salt_pos].checksum[1];
+  checksum[2] = krb5tgs_bufs[salt_pos].checksum[2];
+  checksum[3] = krb5tgs_bufs[salt_pos].checksum[3];
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
+  {
+    u32 w0[4];
+
+    w0[0] = pw_buf0[0];
+    w0[1] = pw_buf0[1];
+    w0[2] = pw_buf0[2];
+    w0[3] = pw_buf0[3];
+
+    u32 w1[4];
+
+    w1[0] = pw_buf1[0];
+    w1[1] = pw_buf1[1];
+    w1[2] = pw_buf1[2];
+    w1[3] = pw_buf1[3];
+
+    u32 w2[4];
+
+    w2[0] = 0;
+    w2[1] = 0;
+    w2[2] = 0;
+    w2[3] = 0;
+
+    u32 w3[4];
+
+    w3[0] = 0;
+    w3[1] = 0;
+    w3[2] = 0;
+    w3[3] = 0;
+
+    const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
+
+    /**
+     * kerberos
+     */
+
+    u32 digest[4];
+
+    kerb_prepare (w0, w1, out_len, checksum, digest);
+
+    u32 tmp[4];
+
+    tmp[0] = digest[0];
+    tmp[1] = digest[1];
+    tmp[2] = digest[2];
+    tmp[3] = digest[3];
+
+    if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2) == 1)
+    {
+      mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos);
+
+      d_return_buf[lid] = 1;
+    }
+  }
+}
+
+__kernel void m13100_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m13100_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}

OpenCL/m13100_a1.cl

@@ -0,0 +1,881 @@
+/**
+ * Authors......: Jens Steube <jens.steube@gmail.com>
+ * Authors......: Fist0urs   <eddy.maaalou@gmail.com>
+
+ * License.....: MIT
+ */
+
+#define _KRB5TGS_
+
+#include "include/constants.h"
+#include "include/kernel_vendor.h"
+
+#define DGST_R0 0
+#define DGST_R1 1
+#define DGST_R2 2
+#define DGST_R3 3
+
+#include "include/kernel_functions.c"
+#include "OpenCL/types_ocl.c"
+#include "OpenCL/common.c"
+
+typedef struct
+{
+  u8 S[256];
+
+  u32 wtf_its_faster;
+
+} RC4_KEY;
+
+static void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
+{
+  u8 tmp;
+
+  tmp           = rc4_key->S[i];
+  rc4_key->S[i] = rc4_key->S[j];
+  rc4_key->S[j] = tmp;
+}
+
+static void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
+{
+  u32 v = 0x03020100;
+  u32 a = 0x04040404;
+
+  __local u32 *ptr = (__local u32 *) rc4_key->S;
+
+  #pragma unroll
+  for (u32 i = 0; i < 64; i++)
+  {
+    *ptr++ = v; v += a;
+  }
+
+  u32 j = 0;
+
+  for (u32 i = 0; i < 16; i++)
+  {
+    u32 idx = i * 16;
+
+    u32 v;
+
+    v = data[0];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[1];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[2];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[3];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+  }
+}
+
+static u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 in[4], u32 out[4])
+{
+  #pragma unroll
+  for (u32 k = 0; k < 4; k++)
+  {
+    u32 xor4 = 0;
+
+    u8 idx;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] <<  0;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] <<  8;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] << 16;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] << 24;
+
+    out[k] = in[k] ^ xor4;
+  }
+
+  return j;
+}
+
+static int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32* edata2)
+{
+  rc4_init_16 (rc4_key, data);
+
+  u32 out[8];
+
+  u8 j = 0;
+
+  /* 
+    8 first bytes are nonce, then ASN1 structs (DER encoding: type-length-data)
+
+    if length >= 128 bytes:
+        length is on 2 bytes and type is \x63\x82 (encode_krb5_enc_tkt_part) and data is an ASN1 sequence \x30\x82
+    else:
+        length is on 1 byte and type is \x63\x81 and data is an ASN1 sequence \x30\x81
+
+    next headers follow the same ASN1 "type-length-data" scheme
+  */
+
+  j = rc4_next_16 (rc4_key,  0, 0, edata2, out);
+
+  if (((out[2] & 0xff00ffff) != 0x30008163) && ((out[2] & 0x0000ffff) != 0x00008263)) return 0;
+
+  j = rc4_next_16 (rc4_key,  16, j, edata2 + 4, out + 4);
+
+  if (((out[4] & 0x00ffffff) != 0x00000503) && (out[4] != 0x050307A0)) return 0;
+
+  // TODO (or not): add RC4'ing of all edata2 then hmac-md5 and compare with 
+  // checksum to be definitely sure that this is the correct pass (even if 
+  // collisions must be very rare)
+
+  return 1;
+}
+
+static void md4_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
+{
+  u32 a = digest[0];
+  u32 b = digest[1];
+  u32 c = digest[2];
+  u32 d = digest[3];
+
+  MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
+
+  MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
+
+  MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
+
+  digest[0] += a;
+  digest[1] += b;
+  digest[2] += c;
+  digest[3] += d;
+}
+
+static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
+{
+  u32 a = digest[0];
+  u32 b = digest[1];
+  u32 c = digest[2];
+  u32 d = digest[3];
+
+  u32 w0_t = w0[0];
+  u32 w1_t = w0[1];
+  u32 w2_t = w0[2];
+  u32 w3_t = w0[3];
+  u32 w4_t = w1[0];
+  u32 w5_t = w1[1];
+  u32 w6_t = w1[2];
+  u32 w7_t = w1[3];
+  u32 w8_t = w2[0];
+  u32 w9_t = w2[1];
+  u32 wa_t = w2[2];
+  u32 wb_t = w2[3];
+  u32 wc_t = w3[0];
+  u32 wd_t = w3[1];
+  u32 we_t = w3[2];
+  u32 wf_t = w3[3];
+
+  MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
+
+  MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
+
+  MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
+
+  MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
+
+  digest[0] += a;
+  digest[1] += b;
+  digest[2] += c;
+  digest[3] += d;
+}
+
+static void hmac_md5_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4])
+{
+  w0[0] = w0[0] ^ 0x36363636;
+  w0[1] = w0[1] ^ 0x36363636;
+  w0[2] = w0[2] ^ 0x36363636;
+  w0[3] = w0[3] ^ 0x36363636;
+  w1[0] = w1[0] ^ 0x36363636;
+  w1[1] = w1[1] ^ 0x36363636;
+  w1[2] = w1[2] ^ 0x36363636;
+  w1[3] = w1[3] ^ 0x36363636;
+  w2[0] = w2[0] ^ 0x36363636;
+  w2[1] = w2[1] ^ 0x36363636;
+  w2[2] = w2[2] ^ 0x36363636;
+  w2[3] = w2[3] ^ 0x36363636;
+  w3[0] = w3[0] ^ 0x36363636;
+  w3[1] = w3[1] ^ 0x36363636;
+  w3[2] = w3[2] ^ 0x36363636;
+  w3[3] = w3[3] ^ 0x36363636;
+
+  ipad[0] = MD5M_A;
+  ipad[1] = MD5M_B;
+  ipad[2] = MD5M_C;
+  ipad[3] = MD5M_D;
+
+  md5_transform (w0, w1, w2, w3, ipad);
+
+  w0[0] = w0[0] ^ 0x6a6a6a6a;
+  w0[1] = w0[1] ^ 0x6a6a6a6a;
+  w0[2] = w0[2] ^ 0x6a6a6a6a;
+  w0[3] = w0[3] ^ 0x6a6a6a6a;
+  w1[0] = w1[0] ^ 0x6a6a6a6a;
+  w1[1] = w1[1] ^ 0x6a6a6a6a;
+  w1[2] = w1[2] ^ 0x6a6a6a6a;
+  w1[3] = w1[3] ^ 0x6a6a6a6a;
+  w2[0] = w2[0] ^ 0x6a6a6a6a;
+  w2[1] = w2[1] ^ 0x6a6a6a6a;
+  w2[2] = w2[2] ^ 0x6a6a6a6a;
+  w2[3] = w2[3] ^ 0x6a6a6a6a;
+  w3[0] = w3[0] ^ 0x6a6a6a6a;
+  w3[1] = w3[1] ^ 0x6a6a6a6a;
+  w3[2] = w3[2] ^ 0x6a6a6a6a;
+  w3[3] = w3[3] ^ 0x6a6a6a6a;
+
+  opad[0] = MD5M_A;
+  opad[1] = MD5M_B;
+  opad[2] = MD5M_C;
+  opad[3] = MD5M_D;
+
+  md5_transform (w0, w1, w2, w3, opad);
+}
+
+static void hmac_md5_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4], u32 digest[4])
+{
+  digest[0] = ipad[0];
+  digest[1] = ipad[1];
+  digest[2] = ipad[2];
+  digest[3] = ipad[3];
+
+  md5_transform (w0, w1, w2, w3, digest);
+
+  w0[0] = digest[0];
+  w0[1] = digest[1];
+  w0[2] = digest[2];
+  w0[3] = digest[3];
+  w1[0] = 0x80;
+  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] = (64 + 16) * 8;
+  w3[3] = 0;
+
+  digest[0] = opad[0];
+  digest[1] = opad[1];
+  digest[2] = opad[2];
+  digest[3] = opad[3];
+
+  md5_transform (w0, w1, w2, w3, digest);
+}
+
+static void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32 checksum[4], u32 digest[4])
+{
+  /**
+   * pads
+   */
+
+  u32 w0_t[4];
+  u32 w1_t[4];
+  u32 w2_t[4];
+  u32 w3_t[4];
+
+  w0_t[0] = w0[0];
+  w0_t[1] = w0[1];
+  w0_t[2] = w0[2];
+  w0_t[3] = w0[3];
+  w1_t[0] = w1[0];
+  w1_t[1] = w1[1];
+  w1_t[2] = w1[2];
+  w1_t[3] = w1[3];
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  // K=MD4(Little_indian(UNICODE(pwd))
+
+  append_0x80_2x4 (w0_t, w1_t, pw_len);
+
+  make_unicode (w1_t, w2_t, w3_t);
+  make_unicode (w0_t, w0_t, w1_t);
+
+  w3_t[2] = pw_len * 8 * 2;
+  w3_t[3] = 0;
+
+  digest[0] = MD4M_A;
+  digest[1] = MD4M_B;
+  digest[2] = MD4M_C;
+  digest[3] = MD4M_D;
+
+  md4_transform (w0_t, w1_t, w2_t, w3_t, digest);
+
+  // K1=MD5_HMAC(K,1); with 2 encoded as little indian on 4 bytes (02000000 in hexa);
+
+  w0_t[0] = digest[0];
+  w0_t[1] = digest[1];
+  w0_t[2] = digest[2];
+  w0_t[3] = digest[3];
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  u32 ipad[4];
+  u32 opad[4];
+
+  hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
+
+  w0_t[0] = 2;
+  w0_t[1] = 0x80;
+  w0_t[2] = 0;
+  w0_t[3] = 0;
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = (64 + 4) * 8;
+  w3_t[3] = 0;
+
+  hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
+
+  // K3=MD5_HMAC(K1,checksum);
+
+  w0_t[0] = digest[0];
+  w0_t[1] = digest[1];
+  w0_t[2] = digest[2];
+  w0_t[3] = digest[3];
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
+
+  w0_t[0] = checksum[0];
+  w0_t[1] = checksum[1];
+  w0_t[2] = checksum[2];
+  w0_t[3] = checksum[3];
+  w1_t[0] = 0x80;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = (64 + 16) * 8;
+  w3_t[3] = 0;
+
+  hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
+}
+
+__kernel void m13100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * modifier
+   */
+
+  __local RC4_KEY rc4_keys[64];
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  u32 wordl0[4];
+
+  wordl0[0] = pws[gid].i[ 0];
+  wordl0[1] = pws[gid].i[ 1];
+  wordl0[2] = pws[gid].i[ 2];
+  wordl0[3] = pws[gid].i[ 3];
+
+  u32 wordl1[4];
+
+  wordl1[0] = pws[gid].i[ 4];
+  wordl1[1] = pws[gid].i[ 5];
+  wordl1[2] = pws[gid].i[ 6];
+  wordl1[3] = pws[gid].i[ 7];
+
+  u32 wordl2[4];
+
+  wordl2[0] = 0;
+  wordl2[1] = 0;
+  wordl2[2] = 0;
+  wordl2[3] = 0;
+
+  u32 wordl3[4];
+
+  wordl3[0] = 0;
+  wordl3[1] = 0;
+  wordl3[2] = 0;
+  wordl3[3] = 0;
+
+  const u32 pw_l_len = pws[gid].pw_len;
+
+  if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
+  {
+    switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
+  }
+
+  /**
+   * salt
+   */
+
+  u32 checksum[4];
+
+  checksum[0] = krb5tgs_bufs[salt_pos].checksum[0];
+  checksum[1] = krb5tgs_bufs[salt_pos].checksum[1];
+  checksum[2] = krb5tgs_bufs[salt_pos].checksum[2];
+  checksum[3] = krb5tgs_bufs[salt_pos].checksum[3];
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
+  {
+    const u32 pw_r_len = combs_buf[il_pos].pw_len;
+
+    const u32 pw_len = pw_l_len + pw_r_len;
+
+    u32 wordr0[4];
+
+    wordr0[0] = combs_buf[il_pos].i[0];
+    wordr0[1] = combs_buf[il_pos].i[1];
+    wordr0[2] = combs_buf[il_pos].i[2];
+    wordr0[3] = combs_buf[il_pos].i[3];
+
+    u32 wordr1[4];
+
+    wordr1[0] = combs_buf[il_pos].i[4];
+    wordr1[1] = combs_buf[il_pos].i[5];
+    wordr1[2] = combs_buf[il_pos].i[6];
+    wordr1[3] = combs_buf[il_pos].i[7];
+
+    u32 wordr2[4];
+
+    wordr2[0] = 0;
+    wordr2[1] = 0;
+    wordr2[2] = 0;
+    wordr2[3] = 0;
+
+    u32 wordr3[4];
+
+    wordr3[0] = 0;
+    wordr3[1] = 0;
+    wordr3[2] = 0;
+    wordr3[3] = 0;
+
+    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
+    {
+      switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
+    }
+
+    u32 w0[4];
+
+    w0[0] = wordl0[0] | wordr0[0];
+    w0[1] = wordl0[1] | wordr0[1];
+    w0[2] = wordl0[2] | wordr0[2];
+    w0[3] = wordl0[3] | wordr0[3];
+
+    u32 w1[4];
+
+    w1[0] = wordl1[0] | wordr1[0];
+    w1[1] = wordl1[1] | wordr1[1];
+    w1[2] = wordl1[2] | wordr1[2];
+    w1[3] = wordl1[3] | wordr1[3];
+
+    u32 w2[4];
+
+    w2[0] = wordl2[0] | wordr2[0];
+    w2[1] = wordl2[1] | wordr2[1];
+    w2[2] = wordl2[2] | wordr2[2];
+    w2[3] = wordl2[3] | wordr2[3];
+
+    u32 w3[4];
+
+    w3[0] = wordl3[0] | wordr3[0];
+    w3[1] = wordl3[1] | wordr3[1];
+    w3[2] = 0;
+    w3[3] = 0;
+
+    /**
+     * kerberos
+     */
+
+    u32 digest[4];
+
+    kerb_prepare (w0, w1, pw_len, checksum, digest);
+
+    u32 tmp[4];
+
+    tmp[0] = digest[0];
+    tmp[1] = digest[1];
+    tmp[2] = digest[2];
+    tmp[3] = digest[3];
+
+    if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2) == 1)
+    {
+      mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos);
+
+      d_return_buf[lid] = 1;
+    }
+  }
+}
+
+__kernel void m13100_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m13100_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m13100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * modifier
+   */
+
+  __local RC4_KEY rc4_keys[64];
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  u32 wordl0[4];
+
+  wordl0[0] = pws[gid].i[ 0];
+  wordl0[1] = pws[gid].i[ 1];
+  wordl0[2] = pws[gid].i[ 2];
+  wordl0[3] = pws[gid].i[ 3];
+
+  u32 wordl1[4];
+
+  wordl1[0] = pws[gid].i[ 4];
+  wordl1[1] = pws[gid].i[ 5];
+  wordl1[2] = pws[gid].i[ 6];
+  wordl1[3] = pws[gid].i[ 7];
+
+  u32 wordl2[4];
+
+  wordl2[0] = 0;
+  wordl2[1] = 0;
+  wordl2[2] = 0;
+  wordl2[3] = 0;
+
+  u32 wordl3[4];
+
+  wordl3[0] = 0;
+  wordl3[1] = 0;
+  wordl3[2] = 0;
+  wordl3[3] = 0;
+
+  const u32 pw_l_len = pws[gid].pw_len;
+
+  if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
+  {
+    switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, combs_buf[0].pw_len);
+  }
+
+  /**
+   * salt
+   */
+
+  u32 checksum[4];
+
+  checksum[0] = krb5tgs_bufs[salt_pos].checksum[0];
+  checksum[1] = krb5tgs_bufs[salt_pos].checksum[1];
+  checksum[2] = krb5tgs_bufs[salt_pos].checksum[2];
+  checksum[3] = krb5tgs_bufs[salt_pos].checksum[3];
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
+  {
+    const u32 pw_r_len = combs_buf[il_pos].pw_len;
+
+    const u32 pw_len = pw_l_len + pw_r_len;
+
+    u32 wordr0[4];
+
+    wordr0[0] = combs_buf[il_pos].i[0];
+    wordr0[1] = combs_buf[il_pos].i[1];
+    wordr0[2] = combs_buf[il_pos].i[2];
+    wordr0[3] = combs_buf[il_pos].i[3];
+
+    u32 wordr1[4];
+
+    wordr1[0] = combs_buf[il_pos].i[4];
+    wordr1[1] = combs_buf[il_pos].i[5];
+    wordr1[2] = combs_buf[il_pos].i[6];
+    wordr1[3] = combs_buf[il_pos].i[7];
+
+    u32 wordr2[4];
+
+    wordr2[0] = 0;
+    wordr2[1] = 0;
+    wordr2[2] = 0;
+    wordr2[3] = 0;
+
+    u32 wordr3[4];
+
+    wordr3[0] = 0;
+    wordr3[1] = 0;
+    wordr3[2] = 0;
+    wordr3[3] = 0;
+
+    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
+    {
+      switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
+    }
+
+    u32 w0[4];
+
+    w0[0] = wordl0[0] | wordr0[0];
+    w0[1] = wordl0[1] | wordr0[1];
+    w0[2] = wordl0[2] | wordr0[2];
+    w0[3] = wordl0[3] | wordr0[3];
+
+    u32 w1[4];
+
+    w1[0] = wordl1[0] | wordr1[0];
+    w1[1] = wordl1[1] | wordr1[1];
+    w1[2] = wordl1[2] | wordr1[2];
+    w1[3] = wordl1[3] | wordr1[3];
+
+    u32 w2[4];
+
+    w2[0] = wordl2[0] | wordr2[0];
+    w2[1] = wordl2[1] | wordr2[1];
+    w2[2] = wordl2[2] | wordr2[2];
+    w2[3] = wordl2[3] | wordr2[3];
+
+    u32 w3[4];
+
+    w3[0] = wordl3[0] | wordr3[0];
+    w3[1] = wordl3[1] | wordr3[1];
+    w3[2] = 0;
+    w3[3] = 0;
+
+    /**
+     * kerberos
+     */
+
+    u32 digest[4];
+
+    kerb_prepare (w0, w1, pw_len, checksum, digest);
+
+    u32 tmp[4];
+
+    tmp[0] = digest[0];
+    tmp[1] = digest[1];
+    tmp[2] = digest[2];
+    tmp[3] = digest[3];
+
+    if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2) == 1)
+    {
+      mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos);
+
+      d_return_buf[lid] = 1;
+    }
+  }
+}
+
+__kernel void m13100_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m13100_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}

OpenCL/m13100_a3.cl

@@ -0,0 +1,809 @@
+/**
+ * Authors......: Jens Steube <jens.steube@gmail.com>
+ * Authors......: Fist0urs   <eddy.maaalou@gmail.com>
+ *
+ * License.....: MIT
+ */
+
+#define _KRB5TGS_
+
+#include "include/constants.h"
+#include "include/kernel_vendor.h"
+
+#define DGST_R0 0
+#define DGST_R1 1
+#define DGST_R2 2
+#define DGST_R3 3
+
+#include "include/kernel_functions.c"
+#include "OpenCL/types_ocl.c"
+#include "OpenCL/common.c"
+
+typedef struct
+{
+  u8 S[256];
+
+  u32 wtf_its_faster;
+
+} RC4_KEY;
+
+static void swap (__local RC4_KEY *rc4_key, const u8 i, const u8 j)
+{
+  u8 tmp;
+
+  tmp           = rc4_key->S[i];
+  rc4_key->S[i] = rc4_key->S[j];
+  rc4_key->S[j] = tmp;
+}
+
+static void rc4_init_16 (__local RC4_KEY *rc4_key, const u32 data[4])
+{
+  u32 v = 0x03020100;
+  u32 a = 0x04040404;
+
+  __local u32 *ptr = (__local u32 *) rc4_key->S;
+
+  #pragma unroll
+  for (u32 i = 0; i < 64; i++)
+  {
+    *ptr++ = v; v += a;
+  }
+
+  u32 j = 0;
+
+  for (u32 i = 0; i < 16; i++)
+  {
+    u32 idx = i * 16;
+
+    u32 v;
+
+    v = data[0];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[1];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[2];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+
+    v = data[3];
+
+    j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
+    j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
+  }
+}
+
+static u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 in[4], u32 out[4])
+{
+  #pragma unroll
+  for (u32 k = 0; k < 4; k++)
+  {
+    u32 xor4 = 0;
+
+    u8 idx;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] <<  0;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] <<  8;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] << 16;
+
+    i += 1;
+    j += rc4_key->S[i];
+
+    swap (rc4_key, i, j);
+
+    idx = rc4_key->S[i] + rc4_key->S[j];
+
+    xor4 |= rc4_key->S[idx] << 24;
+
+    out[k] = in[k] ^ xor4;
+  }
+
+  return j;
+}
+
+static int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32* edata2)
+{
+  rc4_init_16 (rc4_key, data);
+
+  u32 out[8];
+
+  u8 j = 0;
+
+  /* 
+    8 first bytes are nonce, then ASN1 structs (DER encoding: type-length-data)
+
+    if length >= 128 bytes:
+        length is on 2 bytes and type is \x63\x82 (encode_krb5_enc_tkt_part) and data is an ASN1 sequence \x30\x82
+    else:
+        length is on 1 byte and type is \x63\x81 and data is an ASN1 sequence \x30\x81
+
+    next headers follow the same ASN1 "type-length-data" scheme
+  */
+
+  j = rc4_next_16 (rc4_key,  0, 0, edata2, out);
+
+  if (((out[2] & 0xff00ffff) != 0x30008163) && ((out[2] & 0x0000ffff) != 0x00008263)) return 0;
+
+  j = rc4_next_16 (rc4_key,  16, j, edata2 + 4, out + 4);
+
+  if (((out[4] & 0x00ffffff) != 0x00000503) && (out[4] != 0x050307A0)) return 0;
+
+  // TODO (or not): add RC4'ing of all edata2 then hmac-md5 and compare with 
+  // checksum to be definitely sure that this is the correct pass (even if 
+  // collisions must be very rare)
+
+  return 1;
+}
+
+static void md4_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
+{
+  u32 a = digest[0];
+  u32 b = digest[1];
+  u32 c = digest[2];
+  u32 d = digest[3];
+
+  MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
+  MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
+  MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
+  MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
+  MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
+
+  MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
+  MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
+  MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
+  MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
+  MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
+
+  MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
+  MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
+  MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
+  MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
+  MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
+
+  digest[0] += a;
+  digest[1] += b;
+  digest[2] += c;
+  digest[3] += d;
+}
+
+static void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
+{
+  u32 a = digest[0];
+  u32 b = digest[1];
+  u32 c = digest[2];
+  u32 d = digest[3];
+
+  u32 w0_t = w0[0];
+  u32 w1_t = w0[1];
+  u32 w2_t = w0[2];
+  u32 w3_t = w0[3];
+  u32 w4_t = w1[0];
+  u32 w5_t = w1[1];
+  u32 w6_t = w1[2];
+  u32 w7_t = w1[3];
+  u32 w8_t = w2[0];
+  u32 w9_t = w2[1];
+  u32 wa_t = w2[2];
+  u32 wb_t = w2[3];
+  u32 wc_t = w3[0];
+  u32 wd_t = w3[1];
+  u32 we_t = w3[2];
+  u32 wf_t = w3[3];
+
+  MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
+  MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
+  MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
+  MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
+  MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
+
+  MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
+  MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
+  MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
+  MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
+  MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
+
+  MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
+  MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
+  MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
+  MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
+  MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
+
+  MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
+  MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
+  MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
+  MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
+  MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
+
+  digest[0] += a;
+  digest[1] += b;
+  digest[2] += c;
+  digest[3] += d;
+}
+
+static void hmac_md5_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4])
+{
+  w0[0] = w0[0] ^ 0x36363636;
+  w0[1] = w0[1] ^ 0x36363636;
+  w0[2] = w0[2] ^ 0x36363636;
+  w0[3] = w0[3] ^ 0x36363636;
+  w1[0] = w1[0] ^ 0x36363636;
+  w1[1] = w1[1] ^ 0x36363636;
+  w1[2] = w1[2] ^ 0x36363636;
+  w1[3] = w1[3] ^ 0x36363636;
+  w2[0] = w2[0] ^ 0x36363636;
+  w2[1] = w2[1] ^ 0x36363636;
+  w2[2] = w2[2] ^ 0x36363636;
+  w2[3] = w2[3] ^ 0x36363636;
+  w3[0] = w3[0] ^ 0x36363636;
+  w3[1] = w3[1] ^ 0x36363636;
+  w3[2] = w3[2] ^ 0x36363636;
+  w3[3] = w3[3] ^ 0x36363636;
+
+  ipad[0] = MD5M_A;
+  ipad[1] = MD5M_B;
+  ipad[2] = MD5M_C;
+  ipad[3] = MD5M_D;
+
+  md5_transform (w0, w1, w2, w3, ipad);
+
+  w0[0] = w0[0] ^ 0x6a6a6a6a;
+  w0[1] = w0[1] ^ 0x6a6a6a6a;
+  w0[2] = w0[2] ^ 0x6a6a6a6a;
+  w0[3] = w0[3] ^ 0x6a6a6a6a;
+  w1[0] = w1[0] ^ 0x6a6a6a6a;
+  w1[1] = w1[1] ^ 0x6a6a6a6a;
+  w1[2] = w1[2] ^ 0x6a6a6a6a;
+  w1[3] = w1[3] ^ 0x6a6a6a6a;
+  w2[0] = w2[0] ^ 0x6a6a6a6a;
+  w2[1] = w2[1] ^ 0x6a6a6a6a;
+  w2[2] = w2[2] ^ 0x6a6a6a6a;
+  w2[3] = w2[3] ^ 0x6a6a6a6a;
+  w3[0] = w3[0] ^ 0x6a6a6a6a;
+  w3[1] = w3[1] ^ 0x6a6a6a6a;
+  w3[2] = w3[2] ^ 0x6a6a6a6a;
+  w3[3] = w3[3] ^ 0x6a6a6a6a;
+
+  opad[0] = MD5M_A;
+  opad[1] = MD5M_B;
+  opad[2] = MD5M_C;
+  opad[3] = MD5M_D;
+
+  md5_transform (w0, w1, w2, w3, opad);
+}
+
+static void hmac_md5_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4], u32 digest[4])
+{
+  digest[0] = ipad[0];
+  digest[1] = ipad[1];
+  digest[2] = ipad[2];
+  digest[3] = ipad[3];
+
+  md5_transform (w0, w1, w2, w3, digest);
+
+  w0[0] = digest[0];
+  w0[1] = digest[1];
+  w0[2] = digest[2];
+  w0[3] = digest[3];
+  w1[0] = 0x80;
+  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] = (64 + 16) * 8;
+  w3[3] = 0;
+
+  digest[0] = opad[0];
+  digest[1] = opad[1];
+  digest[2] = opad[2];
+  digest[3] = opad[3];
+
+  md5_transform (w0, w1, w2, w3, digest);
+}
+
+static void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32 checksum[4], u32 digest[4])
+{
+  /**
+   * pads
+   */
+
+  u32 w0_t[4];
+  u32 w1_t[4];
+  u32 w2_t[4];
+  u32 w3_t[4];
+
+  w0_t[0] = w0[0];
+  w0_t[1] = w0[1];
+  w0_t[2] = w0[2];
+  w0_t[3] = w0[3];
+  w1_t[0] = w1[0];
+  w1_t[1] = w1[1];
+  w1_t[2] = w1[2];
+  w1_t[3] = w1[3];
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  // K=MD4(Little_indian(UNICODE(pwd))
+
+  append_0x80_2x4 (w0_t, w1_t, pw_len);
+
+  make_unicode (w1_t, w2_t, w3_t);
+  make_unicode (w0_t, w0_t, w1_t);
+
+  w3_t[2] = pw_len * 8 * 2;
+  w3_t[3] = 0;
+
+  digest[0] = MD4M_A;
+  digest[1] = MD4M_B;
+  digest[2] = MD4M_C;
+  digest[3] = MD4M_D;
+
+  md4_transform (w0_t, w1_t, w2_t, w3_t, digest);
+
+  // K1=MD5_HMAC(K,1); with 2 encoded as little indian on 4 bytes (02000000 in hexa);
+
+  w0_t[0] = digest[0];
+  w0_t[1] = digest[1];
+  w0_t[2] = digest[2];
+  w0_t[3] = digest[3];
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  u32 ipad[4];
+  u32 opad[4];
+
+  hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
+
+  w0_t[0] = 2;
+  w0_t[1] = 0x80;
+  w0_t[2] = 0;
+  w0_t[3] = 0;
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = (64 + 4) * 8;
+  w3_t[3] = 0;
+
+  hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
+
+  // K3=MD5_HMAC(K1,checksum);
+
+  w0_t[0] = digest[0];
+  w0_t[1] = digest[1];
+  w0_t[2] = digest[2];
+  w0_t[3] = digest[3];
+  w1_t[0] = 0;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = 0;
+  w3_t[3] = 0;
+
+  hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
+
+  w0_t[0] = checksum[0];
+  w0_t[1] = checksum[1];
+  w0_t[2] = checksum[2];
+  w0_t[3] = checksum[3];
+  w1_t[0] = 0x80;
+  w1_t[1] = 0;
+  w1_t[2] = 0;
+  w1_t[3] = 0;
+  w2_t[0] = 0;
+  w2_t[1] = 0;
+  w2_t[2] = 0;
+  w2_t[3] = 0;
+  w3_t[0] = 0;
+  w3_t[1] = 0;
+  w3_t[2] = (64 + 16) * 8;
+  w3_t[3] = 0;
+
+  hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
+}
+
+static void m13100 (__local RC4_KEY *rc4_keys, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
+{
+  /**
+   * modifier
+   */
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  /**
+   * salt
+   */
+
+  u32 checksum[4];
+
+  checksum[0] = krb5tgs_bufs[salt_pos].checksum[0];
+  checksum[1] = krb5tgs_bufs[salt_pos].checksum[1];
+  checksum[2] = krb5tgs_bufs[salt_pos].checksum[2];
+  checksum[3] = krb5tgs_bufs[salt_pos].checksum[3];
+
+  /**
+   * loop
+   */
+
+  u32 w0l = w0[0];
+
+  for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
+  {
+    const u32 w0r = bfs_buf[il_pos].i;
+
+    w0[0] = w0l | w0r;
+
+    u32 digest[4];
+
+    kerb_prepare (w0, w1, pw_len, checksum, digest);
+
+    u32 tmp[4];
+
+    tmp[0] = digest[0];
+    tmp[1] = digest[1];
+    tmp[2] = digest[2];
+    tmp[3] = digest[3];
+
+    if (decrypt_and_check (&rc4_keys[lid], tmp, krb5tgs_bufs[salt_pos].edata2) == 1)
+    {
+      mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos);
+
+      d_return_buf[lid] = 1;
+    }
+  }
+}
+
+__kernel void m13100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * base
+   */
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  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];
+
+  u32 w1[4];
+
+  w1[0] = 0;
+  w1[1] = 0;
+  w1[2] = 0;
+  w1[3] = 0;
+
+  u32 w2[4];
+
+  w2[0] = 0;
+  w2[1] = 0;
+  w2[2] = 0;
+  w2[3] = 0;
+
+  u32 w3[4];
+
+  w3[0] = 0;
+  w3[1] = 0;
+  w3[2] = 0;
+  w3[3] = 0;
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  /**
+   * main
+   */
+
+  __local RC4_KEY rc4_keys[64];
+
+  m13100 (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, krb5tgs_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
+}
+
+__kernel void m13100_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * base
+   */
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  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];
+
+  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];
+
+  u32 w2[4];
+
+  w2[0] = 0;
+  w2[1] = 0;
+  w2[2] = 0;
+  w2[3] = 0;
+
+  u32 w3[4];
+
+  w3[0] = 0;
+  w3[1] = 0;
+  w3[2] = 0;
+  w3[3] = 0;
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  /**
+   * main
+   */
+
+  __local RC4_KEY rc4_keys[64];
+
+  m13100 (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, krb5tgs_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
+}
+
+__kernel void m13100_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m13100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * base
+   */
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  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];
+
+  u32 w1[4];
+
+  w1[0] = 0;
+  w1[1] = 0;
+  w1[2] = 0;
+  w1[3] = 0;
+
+  u32 w2[4];
+
+  w2[0] = 0;
+  w2[1] = 0;
+  w2[2] = 0;
+  w2[3] = 0;
+
+  u32 w3[4];
+
+  w3[0] = 0;
+  w3[1] = 0;
+  w3[2] = 0;
+  w3[3] = 0;
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  /**
+   * main
+   */
+
+  __local RC4_KEY rc4_keys[64];
+
+  m13100 (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, krb5tgs_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
+}
+
+__kernel void m13100_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * base
+   */
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  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];
+
+  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];
+
+  u32 w2[4];
+
+  w2[0] = 0;
+  w2[1] = 0;
+  w2[2] = 0;
+  w2[3] = 0;
+
+  u32 w3[4];
+
+  w3[0] = 0;
+  w3[1] = 0;
+  w3[2] = 0;
+  w3[3] = 0;
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  /**
+   * main
+   */
+
+  __local RC4_KEY rc4_keys[64];
+
+  m13100 (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, krb5tgs_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
+}
+
+__kernel void m13100_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global krb5tgs_t *krb5tgs_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}

OpenCL/types_ocl.c

@@ -1151,6 +1151,8 @@ typedef struct
   u32  digest_buf[8];
   #elif defined _RAR5_
   u32  digest_buf[4];
+  #elif defined _KRB5TGS_
+  u32  digest_buf[4];
   #endif
 
 } digest_t;
@@ -1270,6 +1272,15 @@ typedef struct
 
 } krb5pa_t;
 
+typedef struct
+{
+  u32 account_info[512];
+  u32 checksum[4];
+  u32 edata2[2560];
+  u32 edata2_len;
+
+} krb5tgs_t;
+
 typedef struct
 {
   u32 salt_buf[16];

docs/readme.txt

@@ -138,6 +138,7 @@ NVidia users require ForceWare 346.59 or later (recommended 358.09 or later)
 - NetNTLMv1 + ESS
 - NetNTLMv2
 - Kerberos 5 AS-REQ Pre-Auth etype 23
+- Kerberos 5 TGS-REP etype 23
 - Netscape LDAP SHA/SSHA
 - LM
 - NTLM

include/constants.h

@@ -151,7 +151,7 @@
 #define SHA256C3f 0xc67178f2u
 #endif
 
-#if defined _MD4_ || defined _DCC2_ || defined _NETNTLMV2_ || defined _KRB5PA_ || defined _MS_DRSR_
+#if defined _MD4_ || defined _DCC2_ || defined _NETNTLMV2_ || defined _KRB5PA_ || defined _MS_DRSR_ || defined _KRB5TGS_
 /**
  * MD4 Constants
  */
@@ -179,7 +179,7 @@
 #define MD4C02 0x6ed9eba1u
 #endif
 
-#if defined _MD5_ || defined _MD5H_ || defined _SAPB_ || defined _OLDOFFICE01_ || defined _WPA_ || defined _MD5_SHA1_ || defined _SHA1_MD5_ || defined _NETNTLMV2_ || defined _KRB5PA_  || defined _PBKDF2_MD5_
+#if defined _MD5_ || defined _MD5H_ || defined _SAPB_ || defined _OLDOFFICE01_ || defined _WPA_ || defined _MD5_SHA1_ || defined _SHA1_MD5_ || defined _NETNTLMV2_ || defined _KRB5PA_  || defined _PBKDF2_MD5_ || defined _KRB5TGS_
 /**
  * MD5 Constants
  */

include/kernel_functions.c

@@ -3,7 +3,7 @@
  * License.....: MIT
  */
 
-#if defined _MD4_ || defined _DCC2_ || defined _NETNTLMV2_ || defined _KRB5PA_ || defined _MS_DRSR_
+#if defined _MD4_ || defined _DCC2_ || defined _NETNTLMV2_ || defined _KRB5PA_ || defined _MS_DRSR_ || defined _KRB5TGS_
 
 #define MD4_F_S(x,y,z)  (((x) & (y)) | ((~(x)) & (z)))
 #define MD4_G_S(x,y,z)  (((x) & (y)) | ((x) & (z)) | ((y) & (z)))
@@ -56,7 +56,7 @@
 
 #endif
 
-#if defined _MD5_ || defined _MD5H_ || defined _SAPB_ || defined _OLDOFFICE01_ || defined _WPA_ || defined _MD5_SHA1_ || defined _SHA1_MD5_ || defined _NETNTLMV2_ || defined _KRB5PA_ || defined _PBKDF2_MD5_
+#if defined _MD5_ || defined _MD5H_ || defined _SAPB_ || defined _OLDOFFICE01_ || defined _WPA_ || defined _MD5_SHA1_ || defined _SHA1_MD5_ || defined _NETNTLMV2_ || defined _KRB5PA_ || defined _PBKDF2_MD5_ || defined _KRB5TGS_
 
 #define MD5_F_S(x,y,z)  ((z) ^ ((x) & ((y) ^ (z))))
 #define MD5_G_S(x,y,z)  ((y) ^ ((z) & ((x) ^ (y))))

include/shared.h

@@ -361,6 +361,7 @@ extern hc_thread_mutex_t mux_display;
 #define HT_12800  "MS-AzureSync PBKDF2-HMAC-SHA256"
 #define HT_12900  "Android FDE (Samsung DEK)"
 #define HT_13000  "RAR5"
+#define HT_13100  "Kerberos 5 TGS-REP etype 23"
 
 #define HT_00011  "Joomla < 2.5.18"
 #define HT_00012  "PostgreSQL"
@@ -694,6 +695,8 @@ extern hc_thread_mutex_t mux_display;
 #define DISPLAY_LEN_MAX_12900 64 + 64 + 32
 #define DISPLAY_LEN_MIN_13000 1 + 4 + 1 + 2 + 1 + 32 + 1 + 2 + 1 + 32 + 1 + 1 + 1 + 16
 #define DISPLAY_LEN_MAX_13000 1 + 4 + 1 + 2 + 1 + 32 + 1 + 2 + 1 + 32 + 1 + 1 + 1 + 16
+#define DISPLAY_LEN_MIN_13100  1 + 7 + 1 + 2 + 1 + 0 + 0 + 32 + 1 + 64
+#define DISPLAY_LEN_MAX_13100  1 + 7 + 1 + 2 + 1 + 2 + 512 + 1 + 32 + 1 + 20480
 
 #define DISPLAY_LEN_MIN_11    32 + 1 + 16
 #define DISPLAY_LEN_MAX_11    32 + 1 + 32
@@ -807,6 +810,7 @@ extern hc_thread_mutex_t mux_display;
 #define HASH_TYPE_ORACLET        47
 #define HASH_TYPE_BSDICRYPT      48
 #define HASH_TYPE_RAR3HP         49
+#define HASH_TYPE_KRB5TGS        50
 
 #define KERN_TYPE_MD5                 0
 #define KERN_TYPE_MD5_PWSLT           10
@@ -952,6 +956,7 @@ extern hc_thread_mutex_t mux_display;
 #define KERN_TYPE_MS_DRSR             12800
 #define KERN_TYPE_ANDROIDFDE_SAMSUNG  12900
 #define KERN_TYPE_RAR5                13000
+#define KERN_TYPE_KRB5TGS             13100
 
 /**
  * signatures
@@ -1021,6 +1026,7 @@ extern hc_thread_mutex_t mux_display;
 #define SIGNATURE_MYWALLET        "$blockchain$"
 #define SIGNATURE_MS_DRSR         "v1;PPH1_MD4"
 #define SIGNATURE_RAR5            "$rar5$"
+#define SIGNATURE_KRB5TGS         "$krb5tgs$23"
 
 /**
  * Default iteration numbers
@@ -1509,6 +1515,7 @@ int hmacsha256_parse_hash         (char *input_buf, uint input_len, hash_t *hash
 int hmacsha512_parse_hash         (char *input_buf, uint input_len, hash_t *hash_buf);
 int hmacmd5_parse_hash            (char *input_buf, uint input_len, hash_t *hash_buf);
 int krb5pa_parse_hash             (char *input_buf, uint input_len, hash_t *hash_buf);
+int krb5tgs_parse_hash            (char *input_buf, uint input_len, hash_t *hash_buf);
 int sapb_parse_hash               (char *input_buf, uint input_len, hash_t *hash_buf);
 int sapg_parse_hash               (char *input_buf, uint input_len, hash_t *hash_buf);
 int drupal7_parse_hash            (char *input_buf, uint input_len, hash_t *hash_buf);

include/types.h

@@ -127,6 +127,15 @@ typedef struct
 
 } krb5pa_t;
 
+typedef struct
+{
+  uint account_info[512];
+  uint checksum[4];
+  uint edata2[2560];
+  uint edata2_len;
+
+} krb5tgs_t;
+
 typedef struct
 {
   uint salt_buf[16];

src/oclHashcat.c

@@ -136,7 +136,7 @@ double TARGET_MS_PROFILE[3]     = { 8, 16, 96 };
 
 #define MAX_DICTSTAT            10000
 
-#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 130
+#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 131
 
 #define global_free(attr)       \
 {                               \
@@ -182,6 +182,7 @@ static uint default_benchmark_algorithms[NUM_DEFAULT_BENCHMARK_ALGORITHMS] =
   5600,
   7300,
   7500,
+  13100,
   8300,
   11100,
   11200,
@@ -577,6 +578,7 @@ const char *USAGE_BIG[] =
   "  11100 = PostgreSQL Challenge-Response Authentication (MD5)",
   "  11200 = MySQL Challenge-Response Authentication (SHA1)",
   "  11400 = SIP digest authentication (MD5)",
+  "  13100 = Kerberos 5 TGS-REP etype 23",
   "",
   "[[ Forums, CMS, E-Commerce, Frameworks, Middleware, Wiki, Management ]]",
   "",
@@ -6109,7 +6111,7 @@ int main (int argc, char **argv)
     return (-1);
   }
 
-  if (hash_mode_chgd && hash_mode > 13000) // just added to remove compiler warnings for hash_mode_chgd
+  if (hash_mode_chgd && hash_mode > 13100) // just added to remove compiler warnings for hash_mode_chgd
   {
     log_error ("ERROR: Invalid hash-type specified");
 
@@ -10357,6 +10359,22 @@ int main (int argc, char **argv)
                    dgst_pos3   = 3;
                    break;
 
+      case 13100:  hash_type   = HASH_TYPE_KRB5TGS;
+                   salt_type   = SALT_TYPE_EMBEDDED;
+                   attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
+                   opts_type   = OPTS_TYPE_PT_GENERATE_LE;
+                   kern_type   = KERN_TYPE_KRB5TGS;
+                   dgst_size   = DGST_SIZE_4_4;
+                   parse_func  = krb5tgs_parse_hash;
+                   sort_by_digest = sort_by_digest_4_4;
+                   opti_type   = OPTI_TYPE_ZERO_BYTE
+                               | OPTI_TYPE_NOT_ITERATED;
+                   dgst_pos0   = 0;
+                   dgst_pos1   = 1;
+                   dgst_pos2   = 2;
+                   dgst_pos3   = 3;
+                   break;
+
       default:     usage_mini_print (PROGNAME); return (-1);
     }
 
@@ -10460,6 +10478,7 @@ int main (int argc, char **argv)
       case 12000:  esalt_size = sizeof (pbkdf2_sha1_t);   break;
       case 12100:  esalt_size = sizeof (pbkdf2_sha512_t); break;
       case 13000:  esalt_size = sizeof (rar5_t);          break;
+      case 13100:  esalt_size = sizeof (krb5tgs_t);       break;
     }
 
     data.esalt_size = esalt_size;

src/shared.c

@@ -5834,6 +5834,7 @@ char *strhashtype (const uint hash_mode)
     case 12800: return ((char *) HT_12800); break;
     case 12900: return ((char *) HT_12900); break;
     case 13000: return ((char *) HT_13000); break;
+    case 13100: return ((char *) HT_13100); break;
   }
 
   return ((char *) "Unknown");
@@ -8278,6 +8279,36 @@ void ascii_digest (char out_buf[4096], uint salt_pos, uint digest_pos)
       byte_swap_32 (digest_buf[1])
     );
   }
+  else if (hash_mode == 13100)
+  {
+    krb5tgs_t *krb5tgss = (krb5tgs_t *) data.esalts_buf;
+
+    krb5tgs_t *krb5tgs = &krb5tgss[salt_pos];
+
+    u8 *ptr_checksum  = (u8 *) krb5tgs->checksum;
+    u8 *ptr_edata2 = (u8 *) krb5tgs->edata2;
+
+    char data[256] = { 0 };
+
+    char *ptr_data = data;
+
+    for (uint i = 0; i < 16; i++, ptr_data += 2)
+      sprintf (ptr_data, "%02x", ptr_checksum[i]);
+
+    /* skip '$' */
+    ptr_data++;
+
+    for (uint i = 0; i < 32; i++, ptr_data += 2)
+      sprintf (ptr_data, "%02x", ptr_edata2[i]);
+
+    *ptr_data = 0;
+
+    snprintf (out_buf, len-1, "%s$%s$%s$%s",
+      SIGNATURE_KRB5TGS,
+      (char *) krb5tgs->account_info,
+      data,
+      data + 33);
+  }
   else
   {
     if (hash_type == HASH_TYPE_MD4)
@@ -18671,6 +18702,102 @@ int rar5_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
   return (PARSER_OK);
 }
 
+int krb5tgs_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
+{
+  if ((input_len < DISPLAY_LEN_MIN_13100) || (input_len > DISPLAY_LEN_MAX_13100)) return (PARSER_GLOBAL_LENGTH);
+
+  if (memcmp (SIGNATURE_KRB5TGS, input_buf, 11)) return (PARSER_SIGNATURE_UNMATCHED);
+
+  u32 *digest = (u32 *) hash_buf->digest;
+
+  salt_t *salt = hash_buf->salt;
+
+  krb5tgs_t *krb5tgs = (krb5tgs_t *) hash_buf->esalt;
+
+  /**
+   * parse line
+   */
+
+  /* Skip '$' */
+  char *account_pos = input_buf + 11 + 1; 
+  
+  char *data_pos;
+
+  uint data_len;
+
+  if (account_pos[0] == '*')
+  {
+    account_pos++;
+
+    data_pos = strchr (account_pos, '*');
+
+    /* Skip '*' */
+    data_pos++;
+
+    if (data_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+    uint account_len = data_pos - account_pos + 1;
+
+    if (account_len >= 512) return (PARSER_SALT_LENGTH);
+
+    /* Skip '$' */
+    data_pos++;
+
+    data_len = input_len - 11 - 1 - account_len - 2;
+
+    memcpy (krb5tgs->account_info, account_pos - 1, account_len);
+  }
+  else
+  {
+    /* assume $krb5tgs$23$checksum$edata2 */
+    data_pos = account_pos;
+
+    memcpy (krb5tgs->account_info, "**", 3);
+
+    data_len = input_len - 11 - 1 - 1;
+  }
+
+  if (data_len < ((16 + 32) * 2)) return (PARSER_SALT_LENGTH);
+
+  char *checksum_ptr = (char *) krb5tgs->checksum;
+
+  for (uint i = 0; i < 16 * 2; i += 2)
+  {
+    const char p0 = data_pos[i + 0];
+    const char p1 = data_pos[i + 1];
+
+    *checksum_ptr++ = hex_convert (p1) << 0
+                     | hex_convert (p0) << 4;
+  }
+
+  char *edata_ptr = (char *) krb5tgs->edata2;
+
+  /* skip '$' */
+  for (uint i = 16 * 2 + 1; i < input_len; i += 2)
+  {
+    const char p0 = data_pos[i + 0];
+    const char p1 = data_pos[i + 1];
+    *edata_ptr++ = hex_convert (p1) << 0
+                    | hex_convert (p0) << 4;
+  }
+  
+  krb5tgs->edata2_len = strlen(edata_ptr - input_len)/(2 * 4);
+
+  salt->salt_buf[0] = krb5tgs->checksum[0];
+  salt->salt_buf[1] = krb5tgs->checksum[1];
+  salt->salt_buf[2] = krb5tgs->checksum[2];
+  salt->salt_buf[3] = krb5tgs->checksum[3];
+  
+  salt->salt_len = 32;
+
+  digest[0] = krb5tgs->checksum[0];
+  digest[1] = krb5tgs->checksum[1];
+  digest[2] = krb5tgs->checksum[2];
+  digest[3] = krb5tgs->checksum[3];
+
+  return (PARSER_OK);
+}
+
 int cf10_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
 {
   if ((input_len < DISPLAY_LEN_MIN_12600) || (input_len > DISPLAY_LEN_MAX_12600)) return (PARSER_GLOBAL_LENGTH);

tools/test.pl

@@ -44,7 +44,7 @@ my $hashcat = "./oclHashcat";
 
 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, 130, 131, 132, 140, 141, 150, 160, 190, 200, 300, 400, 500, 900, 1000, 1100, 1400, 1410, 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, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7100, 7200, 7300, 7400, 7500, 7600, 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, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000);
+my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 130, 131, 132, 140, 141, 150, 160, 190, 200, 300, 400, 500, 900, 1000, 1100, 1400, 1410, 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, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7100, 7200, 7300, 7400, 7500, 7600, 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, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100);
 
 my %is_unicode      = map { $_ => 1 } qw(30 40 130 131 132 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800);
 my %less_fifteen    = map { $_ => 1 } qw(500 1600 1800 2400 2410 3200 6300 7400 10500 10700);
@@ -2209,6 +2209,37 @@ sub verify
 
       next unless (exists ($db->{$hash_line}) and (! defined ($db->{$hash_line})));
     }
+    elsif ($mode == 13100 )
+    {
+      ($hash_in, $word) = split ":", $line;
+
+      next unless defined $hash_in;
+      next unless defined $word;
+
+      my @data = split ('\$', $hash_in);
+
+      next unless scalar @data == 8;
+
+      shift @data;
+
+      my $signature = shift @data;
+      my $algorithm = shift @data;
+      my $user      = shift @data;
+      $user         = substr ($user, 1);
+      my $realm     = shift @data;
+      my $spn       = shift @data;
+      $spn          = substr ($spn, 0, length($spn) - 1);
+      my $checksum  = shift @data;
+      my $edata2    = shift @data;
+      
+      next unless ($signature eq "krb5tgs");
+      next unless (length($checksum) == 32);
+      next unless (length($edata2) == 64);
+
+      $salt = $user.'$'.$realm.'$'.$spn.'$'.substr($edata2, 0, 16);
+
+      next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
+    }
     else
     {
       print "ERROR: hash mode is not supported\n";
@@ -2491,6 +2522,14 @@ sub verify
 
       return unless (substr ($line, 0, $len) eq $hash_out);
     }
+    elsif ($mode == 13100)
+    {
+      $hash_out = gen_hash ($mode, $word, $salt);
+
+      $len = length $hash_out;
+
+      return unless (substr ($line, 0, $len) eq $hash_out);
+    }
     else
     {
       $hash_out = gen_hash ($mode, $word, $salt, $iter);
@@ -2911,6 +2950,12 @@ sub passthrough
     {
       $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32));
     }
+    elsif ($mode == 13100)
+    {
+      $salt_buf = get_random_kerberos5_tgs_salt();
+
+      $tmp_hash = gen_hash ($mode, $word_buf, $salt_buf);
+    }
     else
     {
       print "ERROR: Unsupported hash type\n";
@@ -3638,6 +3683,20 @@ sub single
         }
       }
     }
+    elsif ($mode == 13100)
+    {
+      for (my $i = 1; $i < 27; $i++)
+      {
+        if ($len != 0)
+        {
+          rnd ($mode, $len, 16);
+        }
+        else
+        {
+          rnd ($mode, $i, 16);
+        }
+      }
+    }
   }
 }
 
@@ -6612,6 +6671,41 @@ END_CODE
 
     $tmp_hash = sprintf ('$rar5$16$%s$%d$%s$8$%s', $salt_buf, $iterations, $iv, unpack ("H*", $hash_final));
   }
+  elsif ($mode == 13100)
+  {
+    my @salt_arr = split ('\$', $salt_buf);
+
+    my $user  = $salt_arr[0];
+
+    my $realm = $salt_arr[1];
+
+    my $spn   = $salt_arr[2];
+
+    my $nonce = $salt_arr[3];
+
+    my $cleartext_ticket = '6381b03081ada00703050050a00000a11b3019a003020117a1'.
+    '12041058e0d77776e8b8e03991f2966939222aa2171b154d594b5242544553542e434f4e5'.
+    '44f534f2e434f4da3133011a003020102a10a30081b067472616e6365a40b3009a0030201'.
+    '01a1020400a511180f32303136303231353134343735305aa611180f32303136303231353'.
+    '134343735305aa711180f32303136303231363030343735305aa811180f32303136303232'.
+    '323134343735305a';
+
+    $cleartext_ticket = $nonce . $cleartext_ticket;
+
+    my $k = md4 (encode ("UTF-16LE", $word_buf));
+
+    my $k1 = hmac_md5 ("\x02\x00\x00\x00", $k);
+
+    my $checksum = hmac_md5 (pack ("H*", $cleartext_ticket), $k1);
+
+    my $k3 = hmac_md5 ($checksum, $k1);
+
+    my $cipher = Crypt::RC4->new ($k3);
+
+    my $edata2 = $cipher->RC4 (pack ("H*", $cleartext_ticket));
+
+    $tmp_hash = sprintf ('$krb5tgs$23$*%s$%s$%s*$%s$%s', $user, $realm, $spn, unpack ("H*", $checksum), substr(unpack ("H*", $edata2), 0, 64));
+  }
 
   return ($tmp_hash);
 }
@@ -6706,6 +6800,10 @@ sub rnd
   {
     $salt_buf = get_random_dnssec_salt ();
   }
+  elsif ($mode == 13100)
+  {
+    $salt_buf = get_random_kerberos5_tgs_salt ();
+  }
   else
   {
     my @salt_arr;
@@ -7973,6 +8071,19 @@ sub get_random_kerberos5_salt
   return $salt_buf;
 }
 
+sub get_random_kerberos5_tgs_salt
+{
+  my $nonce = randbytes (8);
+
+  my $user  = "user";
+  my $realm = "realm";
+  my $spn   = "test/spn";
+
+  my $salt_buf = $user . "\$" . $realm . "\$" . $spn . "\$" . unpack ("H*",$nonce);
+
+  return $salt_buf;
+}
+
 sub get_random_md5chap_salt
 {
   my $salt_buf = shift;

tools/test.sh

@@ -9,7 +9,7 @@
 
 # missing hash types: 5200,6211,6221,6231,6241,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 130 131 132 140 141 150 160 190 200 300 400 500 900 1000 1100 1400 1410 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 4300 4400 4500 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 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 12100 12200 12300 12400 12600 12800 12900 13000"
+HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 130 131 132 140 141 150 160 190 200 300 400 500 900 1000 1100 1400 1410 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 4300 4400 4500 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 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 12100 12200 12300 12400 12600 12800 12900 13000 13100"
 
 #ATTACK_MODES="0 1 3 6 7"
 ATTACK_MODES="0 1 3 7"