diff --git a/OpenCL/m00600_a1.cl b/OpenCL/m00600_a1.cl index ce19e2543..101e30f3c 100644 --- a/OpenCL/m00600_a1.cl +++ b/OpenCL/m00600_a1.cl @@ -14,40 +14,8 @@ #include "inc_rp.cl" #include "inc_simd.cl" -typedef struct -{ - u8 digest_length; /* 1 */ - u8 key_length; /* 2 */ - u8 fanout; /* 3 */ - u8 depth; /* 4 */ - u32 leaf_length; /* 8 */ - u32 node_offset; /* 12 */ - u32 xof_length; /* 16 */ - u8 node_depth; /* 17 */ - u8 inner_length; /* 18 */ - u8 reserved[14]; /* 32 */ - u8 salt[BLAKE2B_SALTBYTES]; /* 48 */ - u8 personal[BLAKE2B_PERSONALBYTES]; /* 64 */ -} blake2b_param; - -typedef struct -{ - u64 h[8]; - u64 t[2]; - u64 f[2]; - u8 buf[BLAKE2B_BLOCKBYTES]; - u32 buflen; - u32 outlen; - u8 last_node; -} blake2b_state; - -__constant u64a blake2b_IV[8] = -{ - 0x6a09e667f3bcc908, 0xbb67ae8584caa73b, - 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, - 0x510e527fade682d1, 0x9b05688c2b3e6c1f, - 0x1f83d9abfb41bd6b, 0x5be0cd19137e2179 -}; +#define BLAKE2B_FINAL 1 +#define BLAKE2B_UPDATE 0 __constant u8 blake2b_sigma[12][16] = { @@ -89,71 +57,21 @@ __constant u8 blake2b_sigma[12][16] = BLAKE2B_G(r,7,v[ 3],v[ 4],v[ 9],v[14]); \ } while(0) -void blake2b_compress (const u32x pw[16], const u64 pw_len, 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 - */ - - u32 i; - blake2b_param P[1]; - - P->digest_length = BLAKE2B_OUTBYTES; - P->key_length = 0; - P->fanout = 1; - P->depth = 1; - P->leaf_length = 0; - P->node_offset = 0; - P->xof_length = 0; - P->node_depth = 0; - P->inner_length = 0; - - for (i = 0; i < 14; i++) - P->reserved[i] = 0; - for (i = 0; i < BLAKE2B_SALTBYTES; i++) - P->salt[i] = 0; - for (i = 0; i < BLAKE2B_PERSONALBYTES; i++) - P->personal[i] = 0; - - /* - * Blake2b Init State - */ - - blake2b_state S[1]; - - for (i = 0; i < 8; ++i) - S->h[i] = blake2b_IV[i]; - - S->t[0] = pw_len; - S->t[1] = 0; - S->f[0] = -1; - S->f[1] = 0; - S->buflen = 0; - S->outlen = 0; - S->last_node = 0; - - const u8 *p = (const u8 *)(P); - - /* IV XOR ParamBlock */ - for (i = 0; i < 8; ++i) - S->h[i] ^= *((u64*)(p + sizeof(S->h[i]) * i)); - - S->outlen = P->digest_length; - - /* - * Compress - */ - - u64 v[16]; - u64 m[16]; - - for (i = 0; i < 8; ++i) { - m[i] = swap32(pw[i * 2]); - m[i] <<= 32; - m[i] |= swap32(pw[i * 2 + 1]); - } - + if (isFinal) + f[0] = -1; + + t[0] += hl32_to_64(0, out_len); + + 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; @@ -163,17 +81,22 @@ void blake2b_compress (const u32x pw[16], const u64 pw_len, u64x digest[8]) m[14] = 0; m[15] = 0; - for (i = 0; i < 8; ++i) - v[i] = S->h[i]; - - 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); @@ -188,13 +111,17 @@ void blake2b_compress (const u32x pw[16], const u64 pw_len, u64x digest[8]) BLAKE2B_ROUND(10); BLAKE2B_ROUND(11); - for (i = 0; i < 8; ++i) { - S->h[i] = S->h[i] ^ v[i] ^ v[i + 8]; - digest[i] = swap64(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 @@ -225,7 +152,7 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule { const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); - const u32x pw_len = pw_l_len + pw_r_len; + const u32x out_len = pw_l_len + pw_r_len; /** * concat password candidate @@ -290,37 +217,38 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule w3[2] = wordl3[2] | wordr3[2]; w3[3] = wordl3[3] | wordr3[3]; - 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; - - blake2b_compress(pw, pw_len, digest); + u64x m[16]; + u64x v[16]; + + 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]); @@ -331,15 +259,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 @@ -385,7 +313,7 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule { const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); - const u32x pw_len = pw_l_len + pw_r_len; + const u32x out_len = pw_l_len + pw_r_len; /** * concat password candidate @@ -450,37 +378,38 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule w3[2] = wordl3[2] | wordr3[2]; w3[3] = wordl3[3] | wordr3[3]; - 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; - - blake2b_compress(pw, pw_len, digest); + u64x m[16]; + u64x v[16]; + + 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]); @@ -491,10 +420,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/OpenCL/m00600_a3.cl b/OpenCL/m00600_a3.cl index 3944dcf26..140b3cc66 100644 --- a/OpenCL/m00600_a3.cl +++ b/OpenCL/m00600_a3.cl @@ -12,40 +12,8 @@ #include "inc_common.cl" #include "inc_simd.cl" -typedef struct -{ - u8 digest_length; /* 1 */ - u8 key_length; /* 2 */ - u8 fanout; /* 3 */ - u8 depth; /* 4 */ - u32 leaf_length; /* 8 */ - u32 node_offset; /* 12 */ - u32 xof_length; /* 16 */ - u8 node_depth; /* 17 */ - u8 inner_length; /* 18 */ - u8 reserved[14]; /* 32 */ - u8 salt[BLAKE2B_SALTBYTES]; /* 48 */ - u8 personal[BLAKE2B_PERSONALBYTES]; /* 64 */ -} blake2b_param; - -typedef struct -{ - u64 h[8]; - u64 t[2]; - u64 f[2]; - u8 buf[BLAKE2B_BLOCKBYTES]; - u32 buflen; - u32 outlen; - u8 last_node; -} blake2b_state; - -__constant u64a blake2b_IV[8] = -{ - 0x6a09e667f3bcc908, 0xbb67ae8584caa73b, - 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, - 0x510e527fade682d1, 0x9b05688c2b3e6c1f, - 0x1f83d9abfb41bd6b, 0x5be0cd19137e2179 -}; +#define BLAKE2B_FINAL 1 +#define BLAKE2B_UPDATE 0 __constant u8 blake2b_sigma[12][16] = { @@ -87,71 +55,21 @@ __constant u8 blake2b_sigma[12][16] = BLAKE2B_G(r,7,v[ 3],v[ 4],v[ 9],v[14]); \ } while(0) -void blake2b_compress (const u32x pw[16], const u64 pw_len, 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 - */ - - u32 i; - blake2b_param P[1]; - - P->digest_length = BLAKE2B_OUTBYTES; - P->key_length = 0; - P->fanout = 1; - P->depth = 1; - P->leaf_length = 0; - P->node_offset = 0; - P->xof_length = 0; - P->node_depth = 0; - P->inner_length = 0; - - for (i = 0; i < 14; i++) - P->reserved[i] = 0; - for (i = 0; i < BLAKE2B_SALTBYTES; i++) - P->salt[i] = 0; - for (i = 0; i < BLAKE2B_PERSONALBYTES; i++) - P->personal[i] = 0; - - /* - * Blake2b Init State - */ - - blake2b_state S[1]; - - for (i = 0; i < 8; ++i) - S->h[i] = blake2b_IV[i]; - - S->t[0] = pw_len; - S->t[1] = 0; - S->f[0] = -1; - S->f[1] = 0; - S->buflen = 0; - S->outlen = 0; - S->last_node = 0; - - const u8 *p = (const u8 *)(P); - - /* IV XOR ParamBlock */ - for (i = 0; i < 8; ++i) - S->h[i] ^= *((u64*)(p + sizeof(S->h[i]) * i)); - - S->outlen = P->digest_length; - - /* - * Compress - */ - - u64 v[16]; - u64 m[16]; - - for (i = 0; i < 8; ++i) { - m[i] = swap32(pw[i * 2]); - m[i] <<= 32; - m[i] |= swap32(pw[i * 2 + 1]); - } - + if (isFinal) + f[0] = -1; + + t[0] += hl32_to_64(0, out_len); + + 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; @@ -161,17 +79,22 @@ void blake2b_compress (const u32x pw[16], const u64 pw_len, u64x digest[8]) m[14] = 0; m[15] = 0; - for (i = 0; i < 8; ++i) - v[i] = S->h[i]; - - 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); @@ -186,13 +109,17 @@ void blake2b_compress (const u32x pw[16], const u64 pw_len, u64x digest[8]) BLAKE2B_ROUND(10); BLAKE2B_ROUND(11); - for (i = 0; i < 8; ++i) { - S->h[i] = S->h[i] ^ v[i] ^ v[i + 8]; - digest[i] = swap64(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 u32x *words_buf_r, __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 u32x *words_buf_r, __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 @@ -201,7 +128,6 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); - /** * loop */ @@ -211,40 +137,64 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; - const u32x w0 = w0l | w0r; - u32x pw[16]; - - pw[ 1] = w0; - pw[ 0] = pws[gid].i[ 1]; - pw[ 3] = pws[gid].i[ 2]; - pw[ 2] = pws[gid].i[ 3]; - pw[ 5] = pws[gid].i[ 4]; - pw[ 4] = pws[gid].i[ 5]; - pw[ 7] = pws[gid].i[ 6]; - pw[ 6] = pws[gid].i[ 7]; - pw[ 9] = pws[gid].i[ 8]; - pw[ 8] = pws[gid].i[ 9]; - pw[11] = pws[gid].i[10]; - pw[10] = pws[gid].i[11]; - pw[13] = pws[gid].i[12]; - pw[12] = pws[gid].i[13]; - pw[15] = pws[gid].i[14]; - pw[15] = pws[gid].i[15]; - - u32 pw_len = pws[gid].pw_len; + const u32x w0x = w0l | w0r; + + u32x w0[4]; + u32x w1[4]; + u32x w2[4]; + u32x w3[4]; + + w0[0] = w0x; + w0[1] = pws[gid].i[ 1]; + w0[2] = pws[gid].i[ 2]; + w0[3] = pws[gid].i[ 3]; + w1[0] = pws[gid].i[ 4]; + w1[1] = pws[gid].i[ 5]; + w1[2] = pws[gid].i[ 6]; + w1[3] = pws[gid].i[ 7]; + w2[0] = pws[gid].i[ 8]; + w2[1] = pws[gid].i[ 9]; + w2[2] = pws[gid].i[10]; + w2[3] = pws[gid].i[11]; + w3[0] = pws[gid].i[12]; + w3[1] = pws[gid].i[13]; + w3[2] = pws[gid].i[14]; + w3[3] = pws[gid].i[15]; + + u32x out_len = pws[gid].pw_len; 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; - - blake2b_compress(pw, pw_len, digest); + u64x m[16]; + u64x v[16]; + + 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]); @@ -255,15 +205,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 u32x *words_buf_r, __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 u32x *words_buf_r, __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 @@ -293,40 +243,64 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; - const u32x w0 = w0l | w0r; - u32x pw[16]; - - pw[ 1] = w0; - pw[ 0] = pws[gid].i[ 1]; - pw[ 3] = pws[gid].i[ 2]; - pw[ 2] = pws[gid].i[ 3]; - pw[ 5] = pws[gid].i[ 4]; - pw[ 4] = pws[gid].i[ 5]; - pw[ 7] = pws[gid].i[ 6]; - pw[ 6] = pws[gid].i[ 7]; - pw[ 9] = pws[gid].i[ 8]; - pw[ 8] = pws[gid].i[ 9]; - pw[11] = pws[gid].i[10]; - pw[10] = pws[gid].i[11]; - pw[13] = pws[gid].i[12]; - pw[12] = pws[gid].i[13]; - pw[15] = pws[gid].i[14]; - pw[15] = pws[gid].i[15]; - - u32 pw_len = pws[gid].pw_len; + const u32x w0x = w0l | w0r; + + u32x w0[4]; + u32x w1[4]; + u32x w2[4]; + u32x w3[4]; + + w0[0] = w0x; + w0[1] = pws[gid].i[ 1]; + w0[2] = pws[gid].i[ 2]; + w0[3] = pws[gid].i[ 3]; + w1[0] = pws[gid].i[ 4]; + w1[1] = pws[gid].i[ 5]; + w1[2] = pws[gid].i[ 6]; + w1[3] = pws[gid].i[ 7]; + w2[0] = pws[gid].i[ 8]; + w2[1] = pws[gid].i[ 9]; + w2[2] = pws[gid].i[10]; + w2[3] = pws[gid].i[11]; + w3[0] = pws[gid].i[12]; + w3[1] = pws[gid].i[13]; + w3[2] = pws[gid].i[14]; + w3[3] = pws[gid].i[15]; + + u32x out_len = pws[gid].pw_len; 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; - - blake2b_compress(pw, pw_len, digest); + u64x m[16]; + u64x v[16]; + + 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]); @@ -337,10 +311,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 deleted file mode 100644 index 9f358a4ad..000000000 --- a/b2test.dict +++ /dev/null @@ -1 +0,0 @@ -123456 diff --git a/b2test.hash b/b2test.hash deleted file mode 100644 index 8e367750a..000000000 --- a/b2test.hash +++ /dev/null @@ -1 +0,0 @@ -$BLAKE2$b3910b0f4b6f1aede44da90bb7705a868b265861b36e6f7f29dba7223f6f1ce7b10e0dd25e47deb70bd7f3b24f7da653409cd9014f8715e4013c15fee38ab418 diff --git a/b2test.sh b/b2test.sh deleted file mode 100755 index 1aeeb01d6..000000000 --- a/b2test.sh +++ /dev/null @@ -1 +0,0 @@ -./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/src/Makefile b/src/Makefile index 5f61ee2c6..2b64c798e 100644 --- a/src/Makefile +++ b/src/Makefile @@ -4,7 +4,7 @@ ## SHARED := 0 -DEBUG := 1 +DEBUG := 0 PRODUCTION := 0 ##