diff --git a/OpenCL/inc_hash_constants.h b/OpenCL/inc_hash_constants.h index 654a8159f..bef93efb7 100644 --- a/OpenCL/inc_hash_constants.h +++ b/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; + diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index 0bd649982..954f42e57 100644 --- a/OpenCL/inc_types.cl +++ b/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; - diff --git a/OpenCL/m00600_a0.cl b/OpenCL/m00600_a0.cl index 8a6bb16eb..d9d6e538d 100644 --- a/OpenCL/m00600_a0.cl +++ b/OpenCL/m00600_a0.cl @@ -14,19 +14,8 @@ #include "inc_rp.cl" #include "inc_simd.cl" - -inline u64 hl8_to_32(const u8 a, const u8 b, const u8 c, const u8 d) -{ - return as_uint((uchar4) (a, b, c, d)); -} - -__constant u64a blake2b_IV[8] = -{ - 0x6a09e667f3bcc908, 0xbb67ae8584caa73b, - 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, - 0x510e527fade682d1, 0x9b05688c2b3e6c1f, - 0x1f83d9abfb41bd6b, 0x5be0cd19137e2179 -}; +#define BLAKE2B_FINAL 1 +#define BLAKE2B_UPDATE 0 __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) { - /* - * Blake2b Init Param - */ + if (isFinal) + f[0] = -1; - u8 p_digest_length = digest_length; - 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]; + t[0] += hl32_to_64(0, out_len); - /* - * Blake2b Init State - */ - - u64x s_h[8]; - u64x s_t[2]; - u64x s_f[2]; - u32x s_buflen; - 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[0] = hl32_to_64(w0[1], w0[0]); + 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]); + m[4] = hl32_to_64(w2[1], w2[0]); + m[5] = hl32_to_64(w2[3], w2[2]); + m[6] = hl32_to_64(w3[1], w3[0]); + m[7] = hl32_to_64(w3[3], w3[2]); 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[ 1] = s_h[1]; - v[ 2] = s_h[2]; - v[ 3] = s_h[3]; - v[ 4] = s_h[4]; - v[ 5] = s_h[5]; - v[ 6] = s_h[6]; - v[ 7] = s_h[7]; - v[ 8] = blake2b_IV[0]; - v[ 9] = blake2b_IV[1]; - v[10] = blake2b_IV[2]; - v[11] = blake2b_IV[3]; - v[12] = blake2b_IV[4] ^ s_t[0]; - v[13] = blake2b_IV[5] ^ s_t[1]; - v[14] = blake2b_IV[6] ^ s_f[0]; - v[15] = blake2b_IV[7] ^ s_f[1]; + v[ 0] = h[0]; + v[ 1] = h[1]; + v[ 2] = h[2]; + v[ 3] = h[3]; + v[ 4] = h[4]; + v[ 5] = h[5]; + v[ 6] = h[6]; + v[ 7] = h[7]; + v[ 8] = BLAKE2B_IV_00; + v[ 9] = BLAKE2B_IV_01; + v[10] = BLAKE2B_IV_02; + v[11] = BLAKE2B_IV_03; + v[12] = BLAKE2B_IV_04 ^ t[0]; + v[13] = BLAKE2B_IV_05 ^ t[1]; + v[14] = BLAKE2B_IV_06 ^ f[0]; + 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) { - s_h[i] = s_h[i] ^ v[i] ^ v[i + 8]; - digest[i] = s_h[i]; - } + h[0] = h[0] ^ v[0] ^ v[ 8]; + h[1] = h[1] ^ v[1] ^ v[ 9]; + 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; - digest[1] = 0; - digest[2] = 0; - digest[3] = 0; - digest[4] = 0; - digest[5] = 0; - digest[6] = 0; - digest[7] = 0; + u64x m[16]; + u64x v[16]; - u64 p_salt[2] = { 0, 0 }; - u64 p_key[16] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 }; - blake2b_transform(pw, out_len, p_salt, p_key, 0, BLAKE2B_OUTBYTES, digest); + u64x h[8]; + u64x t[2]; + 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; - digest[1] = 0; - digest[2] = 0; - digest[3] = 0; - digest[4] = 0; - digest[5] = 0; - digest[6] = 0; - digest[7] = 0; + u64x h[8]; + u64x t[2]; + u64x f[2]; - u64 p_salt[2] = { 0, 0 }; - u64 p_key[16] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 }; - blake2b_transform(pw, out_len, p_salt, p_key, 0, BLAKE2B_OUTBYTES, digest); + 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]); @@ -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) { } diff --git a/b2test.dict b/b2test.dict new file mode 100644 index 000000000..9f358a4ad --- /dev/null +++ b/b2test.dict @@ -0,0 +1 @@ +123456 diff --git a/b2test.hash b/b2test.hash new file mode 100644 index 000000000..8e367750a --- /dev/null +++ b/b2test.hash @@ -0,0 +1 @@ +$BLAKE2$b3910b0f4b6f1aede44da90bb7705a868b265861b36e6f7f29dba7223f6f1ce7b10e0dd25e47deb70bd7f3b24f7da653409cd9014f8715e4013c15fee38ab418 diff --git a/b2test.sh b/b2test.sh new file mode 100755 index 000000000..1aeeb01d6 --- /dev/null +++ b/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 diff --git a/include/cpu_blake2.h b/include/cpu_blake2.h new file mode 100644 index 000000000..07a4d8e32 --- /dev/null +++ b/include/cpu_blake2.h @@ -0,0 +1,19 @@ +/** + * Author......: See docs/credits.txt + * License.....: MIT + */ + +#ifndef _CPU_BLAKE2_H +#define _CPU_BLAKE2_H + +#include + +const u64 blake2b_IV[8] = +{ + 0x6a09e667f3bcc908, 0xbb67ae8584caa73b, + 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, + 0x510e527fade682d1, 0x9b05688c2b3e6c1f, + 0x1f83d9abfb41bd6b, 0x5be0cd19137e2179 +}; + +#endif // _CPU_BLAKE2_H diff --git a/include/types.h b/include/types.h index 437e3bad0..c63ace75a 100644 --- a/include/types.h +++ b/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]; diff --git a/src/Makefile b/src/Makefile index 2b64c798e..5f61ee2c6 100644 --- a/src/Makefile +++ b/src/Makefile @@ -4,7 +4,7 @@ ## SHARED := 0 -DEBUG := 0 +DEBUG := 1 PRODUCTION := 0 ## diff --git a/src/cpu_blake.c b/src/cpu_blake.c new file mode 100644 index 000000000..7cada99b3 --- /dev/null +++ b/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; + diff --git a/src/interface.c b/src/interface.c index 2872b9959..5ddc05add 100644 --- a/src/interface.c +++ b/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;