Moved init params in CPU (interface.c). Fixed vector-type problem in kernel a0.

OpenCL/inc_hash_constants.h

@@ -638,3 +638,17 @@ typedef enum fortigate_constants
   FORTIGATE_F=0x705b81a9
 
 } fortigate_constants_t;
+
+typedef enum blake2b_constants
+{
+  BLAKE2B_IV_00=0x6a09e667f3bcc908,
+  BLAKE2B_IV_01=0xbb67ae8584caa73b,
+  BLAKE2B_IV_02=0x3c6ef372fe94f82b,
+  BLAKE2B_IV_03=0xa54ff53a5f1d36f1,
+  BLAKE2B_IV_04=0x510e527fade682d1,
+  BLAKE2B_IV_05=0x9b05688c2b3e6c1f,
+  BLAKE2B_IV_06=0x1f83d9abfb41bd6b,
+  BLAKE2B_IV_07=0x5be0cd19137e2179
+
+} blake2b_constants_t;
+

OpenCL/inc_types.cl

@@ -752,6 +752,17 @@ typedef struct itunes_backup
 
 } itunes_backup_t;
 
+typedef struct
+{
+  u64 h[8];
+  u64 t[2];
+  u64 f[2];
+  u32 buflen;
+  u32 outlen;
+  u8  last_node;
+
+} blake2_state_t;
+
 typedef struct luks_tmp
 {
   u32 ipad32[8];
@@ -1469,12 +1480,3 @@ typedef enum combinator_mode
 
 } combinator_mode_t;
 
-typedef enum blake2b_constants__
-{
-  BLAKE2B_BLOCKBYTES = 128,
-  BLAKE2B_OUTBYTES   = 64,
-  BLAKE2B_KEYBYTES   = 64,
-  BLAKE2B_SALTBYTES  = 16,
-  BLAKE2B_PERSONALBYTES = 16
-} blake2b_constants;
-

OpenCL/m00600_a0.cl

@@ -14,19 +14,8 @@
 #include "inc_rp.cl"
 #include "inc_simd.cl"
 
-
+#define BLAKE2B_FINAL   1
-inline u64 hl8_to_32(const u8 a, const u8 b, const u8 c, const u8 d)
+#define BLAKE2B_UPDATE  0
-{
-  return as_uint((uchar4) (a, b, c, d));
-}
-
-__constant u64a blake2b_IV[8] =
-{
-  0x6a09e667f3bcc908, 0xbb67ae8584caa73b,
-  0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1,
-  0x510e527fade682d1, 0x9b05688c2b3e6c1f,
-  0x1f83d9abfb41bd6b, 0x5be0cd19137e2179
-};
 
 __constant u8a blake2b_sigma[12][16] =
 {
@@ -68,65 +57,21 @@ __constant u8a blake2b_sigma[12][16] =
     BLAKE2B_G(r,7,v[ 3],v[ 4],v[ 9],v[14]); \
 } while(0)
 
-void blake2b_transform(const u32x pw[16], const u32x out_len, const u64 p_salt[2], const u64 p_key[16], const u8 key_length, const u8 digest_length, u64x digest[8])
+void blake2b_compress(u64x h[8], u64x t[2], u64x f[2], u64x m[16], u64x v[16], const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], const u32x out_len, const u8 isFinal)
 {
-  /*
+  if (isFinal)
-   * Blake2b Init Param
+    f[0] = -1;
-   */
 
-  u8  p_digest_length = digest_length;
+  t[0] += hl32_to_64(0, out_len);
-  u8  p_key_length    = key_length;
-  u8  p_fanout        = 1;
-  u8  p_depth         = 1;
-  u32 p_leaf_length   = 0;
-  u32 p_node_offset   = 0;
-  u32 p_xof_length    = 0;
-  u8  p_node_depth    = 0;
-  u8  p_inner_length  = 0;
-  u8  p_reserved[14];
-  u8  p_personnel[BLAKE2B_PERSONALBYTES];
 
-  /*
+  m[0] = hl32_to_64(w0[1], w0[0]);
-   * Blake2b Init State
+  m[1] = hl32_to_64(w0[3], w0[2]);
-   */
+  m[2] = hl32_to_64(w1[1], w1[0]);
-
+  m[3] = hl32_to_64(w1[3], w1[2]);
-  u64x s_h[8];
+  m[4] = hl32_to_64(w2[1], w2[0]);
-  u64x s_t[2];
+  m[5] = hl32_to_64(w2[3], w2[2]);
-  u64x s_f[2];
+  m[6] = hl32_to_64(w3[1], w3[0]);
-  u32x s_buflen;
+  m[7] = hl32_to_64(w3[3], w3[2]);
-  u32x s_outlen;
-  u8x  s_last_node;
-
-  s_h[0] = blake2b_IV[0] ^ hl8_to_32(p_digest_length, p_key_length, p_fanout, p_depth);
-  s_h[1] = blake2b_IV[1];
-  s_h[2] = blake2b_IV[2];
-  s_h[3] = blake2b_IV[3];
-  s_h[4] = blake2b_IV[4] ^ p_salt[0];
-  s_h[5] = blake2b_IV[5] ^ p_salt[1];
-  s_h[6] = blake2b_IV[6];
-  s_h[7] = blake2b_IV[7];
-  s_t[0] = hl32_to_64(0, out_len);
-  s_t[1] =  0;
-  s_f[0] = -1;
-  s_f[1] =  0;
-  s_outlen = 0;
-  s_last_node = 0;
-
-  /*
-   * Compress
-   */
-
-  u64x v[16];
-  u64x m[16];
-
-  m[0] = swap64(hl32_to_64(pw[ 1], pw[ 0]));
-  m[1] = swap64(hl32_to_64(pw[ 3], pw[ 2]));
-  m[2] = swap64(hl32_to_64(pw[ 5], pw[ 4]));
-  m[3] = swap64(hl32_to_64(pw[ 7], pw[ 6]));
-  m[4] = swap64(hl32_to_64(pw[ 9], pw[ 8]));
-  m[5] = swap64(hl32_to_64(pw[11], pw[10]));
-  m[6] = swap64(hl32_to_64(pw[13], pw[12]));
-  m[7] = swap64(hl32_to_64(pw[15], pw[14]));
   m[8] = 0;
   m[9] = 0;
   m[10] = 0;
@@ -136,22 +81,22 @@ void blake2b_transform(const u32x pw[16], const u32x out_len, const u64 p_salt[2
   m[14] = 0;
   m[15] = 0;
 
-  v[ 0] = s_h[0];
+  v[ 0] = h[0];
-  v[ 1] = s_h[1];
+  v[ 1] = h[1];
-  v[ 2] = s_h[2];
+  v[ 2] = h[2];
-  v[ 3] = s_h[3];
+  v[ 3] = h[3];
-  v[ 4] = s_h[4];
+  v[ 4] = h[4];
-  v[ 5] = s_h[5];
+  v[ 5] = h[5];
-  v[ 6] = s_h[6];
+  v[ 6] = h[6];
-  v[ 7] = s_h[7];
+  v[ 7] = h[7];
-  v[ 8] = blake2b_IV[0];
+  v[ 8] = BLAKE2B_IV_00;
-  v[ 9] = blake2b_IV[1];
+  v[ 9] = BLAKE2B_IV_01;
-  v[10] = blake2b_IV[2];
+  v[10] = BLAKE2B_IV_02;
-  v[11] = blake2b_IV[3];
+  v[11] = BLAKE2B_IV_03;
-  v[12] = blake2b_IV[4] ^ s_t[0];
+  v[12] = BLAKE2B_IV_04 ^ t[0];
-  v[13] = blake2b_IV[5] ^ s_t[1];
+  v[13] = BLAKE2B_IV_05 ^ t[1];
-  v[14] = blake2b_IV[6] ^ s_f[0];
+  v[14] = BLAKE2B_IV_06 ^ f[0];
-  v[15] = blake2b_IV[7] ^ s_f[1];
+  v[15] = BLAKE2B_IV_07 ^ f[1];
 
   BLAKE2B_ROUND( 0);
   BLAKE2B_ROUND( 1);
@@ -166,13 +111,17 @@ void blake2b_transform(const u32x pw[16], const u32x out_len, const u64 p_salt[2
   BLAKE2B_ROUND(10);
   BLAKE2B_ROUND(11);
 
-  for (int i = 0; i < 8; ++i) {
+  h[0] = h[0] ^ v[0] ^ v[ 8];
-    s_h[i] = s_h[i] ^ v[i] ^ v[i + 8];
+  h[1] = h[1] ^ v[1] ^ v[ 9];
-    digest[i] = s_h[i];
+  h[2] = h[2] ^ v[2] ^ v[10];
-  }
+  h[3] = h[3] ^ v[3] ^ v[11];
+  h[4] = h[4] ^ v[4] ^ v[12];
+  h[5] = h[5] ^ v[5] ^ v[13];
+  h[6] = h[6] ^ v[6] ^ v[14];
+  h[7] = h[7] ^ v[7] ^ v[15];
 }
 
-__kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const blake2_state_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 { 
   /**
    * modifier
@@ -208,39 +157,39 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
 
     const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
-    u32x pw[16];
-
-    pw[ 1] = swap32(w0[0]);
-    pw[ 0] = swap32(w0[1]);
-    pw[ 3] = swap32(w0[2]);
-    pw[ 2] = swap32(w0[3]);
-    pw[ 5] = swap32(w1[0]);
-    pw[ 4] = swap32(w1[1]);
-    pw[ 7] = swap32(w1[2]);
-    pw[ 6] = swap32(w1[3]);
-    pw[ 9] = swap32(w2[0]);
-    pw[ 8] = swap32(w2[1]);
-    pw[11] = swap32(w2[2]);
-    pw[10] = swap32(w2[3]);
-    pw[13] = swap32(w3[0]);
-    pw[12] = swap32(w3[1]);
-    pw[15] = swap32(w3[2]);
-    pw[14] = swap32(w3[3]);
-
     u64x digest[8];
 
-    digest[0] = 0;
+    u64x m[16];
-    digest[1] = 0;
+    u64x v[16];
-    digest[2] = 0;
-    digest[3] = 0;
-    digest[4] = 0;
-    digest[5] = 0;
-    digest[6] = 0;
-    digest[7] = 0;
 
-    u64 p_salt[2] = { 0, 0 };
+    u64x h[8];
-    u64 p_key[16] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
+    u64x t[2];
-    blake2b_transform(pw, out_len, p_salt, p_key, 0, BLAKE2B_OUTBYTES, digest);
+    u64x f[2];
+
+    h[0] = esalt_bufs->h[0];
+    h[1] = esalt_bufs->h[1];
+    h[2] = esalt_bufs->h[2];
+    h[3] = esalt_bufs->h[3];
+    h[4] = esalt_bufs->h[4];
+    h[5] = esalt_bufs->h[5];
+    h[6] = esalt_bufs->h[6];
+    h[7] = esalt_bufs->h[7];
+
+    t[0] = esalt_bufs->t[0];
+    t[1] = esalt_bufs->t[1];
+    f[0] = esalt_bufs->f[0];
+    f[1] = esalt_bufs->f[1];
+
+    blake2b_compress(h, t, f, m, v, w0, w1, w2, w3, out_len, BLAKE2B_FINAL);
+
+    digest[0] = h[0];
+    digest[1] = h[1];
+    digest[2] = h[2];
+    digest[3] = h[3];
+    digest[4] = h[4];
+    digest[5] = h[5];
+    digest[6] = h[6];
+    digest[7] = h[7];
 
     const u32x r0 = h32_from_64(digest[0]);
     const u32x r1 = l32_from_64(digest[0]);
@@ -251,15 +200,15 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
   } 
 }
 
-__kernel void m00600_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m00600_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const blake2_state_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
 }
 
-__kernel void m00600_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m00600_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const blake2_state_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
 }
 
-__kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const blake2_state_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 { 
   /**
    * modifier
@@ -310,39 +259,38 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
     
     const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
-    u32x pw[16];
-
-    pw[ 1] = swap32(w0[0]);
-    pw[ 0] = swap32(w0[1]);
-    pw[ 3] = swap32(w0[2]);
-    pw[ 2] = swap32(w0[3]);
-    pw[ 5] = swap32(w1[0]);
-    pw[ 4] = swap32(w1[1]);
-    pw[ 7] = swap32(w1[2]);
-    pw[ 6] = swap32(w1[3]);
-    pw[ 9] = swap32(w2[0]);
-    pw[ 8] = swap32(w2[1]);
-    pw[11] = swap32(w2[2]);
-    pw[10] = swap32(w2[3]);
-    pw[13] = swap32(w3[0]);
-    pw[12] = swap32(w3[1]);
-    pw[15] = swap32(w3[2]);
-    pw[14] = swap32(w3[3]);
-
     u64x digest[8];
+    u64x m[16];
+    u64x v[16];
 
-    digest[0] = 0;
+    u64x h[8];
-    digest[1] = 0;
+    u64x t[2];
-    digest[2] = 0;
+    u64x f[2];
-    digest[3] = 0;
-    digest[4] = 0;
-    digest[5] = 0;
-    digest[6] = 0;
-    digest[7] = 0;
 
-    u64 p_salt[2] = { 0, 0 };
+    h[0] = esalt_bufs->h[0];
-    u64 p_key[16] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
+    h[1] = esalt_bufs->h[1];
-    blake2b_transform(pw, out_len, p_salt, p_key, 0, BLAKE2B_OUTBYTES, digest);
+    h[2] = esalt_bufs->h[2];
+    h[3] = esalt_bufs->h[3];
+    h[4] = esalt_bufs->h[4];
+    h[5] = esalt_bufs->h[5];
+    h[6] = esalt_bufs->h[6];
+    h[7] = esalt_bufs->h[7];
+
+    t[0] = esalt_bufs->t[0];
+    t[1] = esalt_bufs->t[1];
+    f[0] = esalt_bufs->f[0];
+    f[1] = esalt_bufs->f[1];
+
+    blake2b_compress(h, t, f, m, v, w0, w1, w2, w3, out_len, BLAKE2B_FINAL);
+
+    digest[0] = h[0];
+    digest[1] = h[1];
+    digest[2] = h[2];
+    digest[3] = h[3];
+    digest[4] = h[4];
+    digest[5] = h[5];
+    digest[6] = h[6];
+    digest[7] = h[7];
 
     const u32x r0 = h32_from_64(digest[0]);
     const u32x r1 = l32_from_64(digest[0]);
@@ -353,10 +301,10 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   }  
 }
 
-__kernel void m00600_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m00600_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const blake2_state_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
 }
 
-__kernel void m00600_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+__kernel void m00600_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const blake2_state_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 {
 }

b2test.dict

@@ -0,0 +1 @@
+123456

b2test.hash

@@ -0,0 +1 @@
+$BLAKE2$b3910b0f4b6f1aede44da90bb7705a868b265861b36e6f7f29dba7223f6f1ce7b10e0dd25e47deb70bd7f3b24f7da653409cd9014f8715e4013c15fee38ab418

b2test.sh

@@ -0,0 +1 @@
+./hashcat -m 600 -a 0 b2test.hash b2test.dict --opencl-vector-width=1 -u1 -n1 --force --weak-hash-threshold=0 --potfile-disable

include/cpu_blake2.h

@@ -0,0 +1,19 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+#ifndef _CPU_BLAKE2_H
+#define _CPU_BLAKE2_H
+
+#include <string.h>
+
+const u64 blake2b_IV[8] =
+{
+  0x6a09e667f3bcc908, 0xbb67ae8584caa73b,
+  0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1,
+  0x510e527fade682d1, 0x9b05688c2b3e6c1f,
+  0x1f83d9abfb41bd6b, 0x5be0cd19137e2179
+};
+
+#endif // _CPU_BLAKE2_H

include/types.h

@@ -663,6 +663,34 @@ typedef enum user_options_map
  * structs
  */
 
+typedef struct
+{
+  u8  digest_length;
+  u8  key_length;
+  u8  fanout;
+  u8  depth;
+  u32 leaf_length;
+  u32 node_offset;
+  u32 xof_length;
+  u8  node_depth;
+  u8  inner_length;
+  u8  reserved[14];
+  u8  salt[16];
+  u8  personnel[16];
+
+} blake2_params_t;
+
+typedef struct
+{
+  u64 h[8];
+  u64 t[2];
+  u64 f[2];
+  u32 buflen;
+  u32 outlen;
+  u8  last_node;
+
+} blake2_state_t;
+
 typedef struct
 {
   u32 salt_buf[16];

src/Makefile

@@ -4,7 +4,7 @@
 ##
 
 SHARED                  := 0
-DEBUG                   := 0
+DEBUG                   := 1
 PRODUCTION              := 0
 
 ##

src/cpu_blake.c

@@ -0,0 +1,31 @@
+
+
+
+typedef struct
+{
+  u8  digest_length;
+  u8  key_length;
+  u8  fanout;
+  u8  depth;
+  u32 leaf_length;
+  u32 node_offset;
+  u32 xof_length;
+  u8  node_depth;
+  u8  inner_length;
+  u8  reserved[14];
+  u8  salt[16];
+  u8  personnel[16];
+
+} blake2_params_t;
+
+typedef struct
+{
+  u64 h[8];
+  u64 t[2];
+  u64 f[2];
+  u32 buflen;
+  u32 outlen;
+  u8  last_node;
+
+} blake2_state_t;
+

src/interface.c

@@ -16,6 +16,7 @@
 #include "cpu_md5.h"
 #include "cpu_sha1.h"
 #include "cpu_sha256.h"
+#include "cpu_blake2.h"
 #include "interface.h"
 #include "ext_lzma.h"
 
@@ -5284,6 +5285,31 @@ int blake2b_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
   digest[6] = hex_to_u64 ((const u8 *) &input_hash_buf[ 96]);
   digest[7] = hex_to_u64 ((const u8 *) &input_hash_buf[112]);
 
+  // Initialize Blake2 Params and State
+
+  blake2_params_t P[1];
+  blake2_state_t  *S = (blake2_state_t *) hash_buf->esalt;
+
+  memset(P,  0, sizeof(blake2_params_t));
+  memset(S,  0, sizeof(blake2_state_t));
+
+  P->digest_length = 64;
+  P->fanout = 1;
+  P->depth  = 1;
+
+  S->h[0] = blake2b_IV[0];
+  S->h[1] = blake2b_IV[1];
+  S->h[2] = blake2b_IV[2];
+  S->h[3] = blake2b_IV[3];
+  S->h[4] = blake2b_IV[4];
+  S->h[5] = blake2b_IV[5];
+  S->h[6] = blake2b_IV[6];
+  S->h[7] = blake2b_IV[7];
+
+  // XOR P with S
+  for (uint i = 0; i < sizeof(blake2_params_t); i++)
+    ((u8 *)S)[i] ^= ((u8 *)P)[i];
+
   return (PARSER_OK);
 }
 
@@ -19364,15 +19390,13 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
                  break;
 
     case   600:  hashconfig->hash_type      = HASH_TYPE_BLAKE2B;
-                 hashconfig->salt_type      = SALT_TYPE_NONE;
+                 hashconfig->salt_type      = SALT_TYPE_EMBEDDED;
                  hashconfig->attack_exec    = ATTACK_EXEC_INSIDE_KERNEL;
                  hashconfig->opts_type      = OPTS_TYPE_PT_GENERATE_LE;
                  hashconfig->kern_type      = KERN_TYPE_BLAKE2B;
                  hashconfig->dgst_size      = DGST_SIZE_8_8;
                  hashconfig->parse_func     = blake2b_parse_hash;
                  hashconfig->opti_type      = OPTI_TYPE_ZERO_BYTE
-                                            | OPTI_TYPE_NOT_ITERATED
-                                            | OPTI_TYPE_NOT_SALTED
                                             | OPTI_TYPE_USES_BITS_64
                                             | OPTI_TYPE_RAW_HASH;
                  hashconfig->dgst_pos0      = 1;
@@ -22667,6 +22691,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
 
   switch (hashconfig->hash_mode)
   {
+    case   600: hashconfig->esalt_size = sizeof (blake2_state_t);   break;
     case  2500: hashconfig->esalt_size = sizeof (wpa_t);            break;
     case  5300: hashconfig->esalt_size = sizeof (ikepsk_t);         break;
     case  5400: hashconfig->esalt_size = sizeof (ikepsk_t);         break;