mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-02 05:30:19 +00:00
Merge branch 'master' of https://github.com/hashcat/hashcat
This commit is contained in:
commit
282022406c
@ -12,7 +12,7 @@
|
||||
* *
|
||||
*/
|
||||
|
||||
__constant u32a k_sbox[256] =
|
||||
__constant const u32a k_sbox[256] =
|
||||
{
|
||||
0xfc, 0xee, 0xdd, 0x11, 0xcf, 0x6e, 0x31, 0x16,
|
||||
0xfb, 0xc4, 0xfa, 0xda, 0x23, 0xc5, 0x04, 0x4d,
|
||||
@ -48,7 +48,7 @@ __constant u32a k_sbox[256] =
|
||||
0xd1, 0x66, 0xaf, 0xc2, 0x39, 0x4b, 0x63, 0xb6
|
||||
};
|
||||
|
||||
__constant u32a k_sbox_inv[256] =
|
||||
__constant const u32a k_sbox_inv[256] =
|
||||
{
|
||||
0xa5, 0x2d, 0x32, 0x8f, 0x0e, 0x30, 0x38, 0xc0,
|
||||
0x54, 0xe6, 0x9e, 0x39, 0x55, 0x7e, 0x52, 0x91,
|
||||
@ -84,31 +84,26 @@ __constant u32a k_sbox_inv[256] =
|
||||
0xd6, 0x20, 0x0a, 0x08, 0x00, 0x4c, 0xd7, 0x74
|
||||
};
|
||||
|
||||
__constant int k_vec[16] =
|
||||
{
|
||||
0x94, 0x20, 0x85, 0x10, 0xc2, 0xc0, 0x01, 0xfb,
|
||||
0x01, 0xc0, 0xc2, 0x10, 0x85, 0x20, 0x94, 0x01
|
||||
};
|
||||
|
||||
#define extract_byte(x,n) (((x) >> (8 * (n))) & 0xff)
|
||||
|
||||
#define k_lookup(w,sbox) \
|
||||
for (int i = 0; i < 4; i++) \
|
||||
w[i] = (sbox[extract_byte(w[i],0)] << 0) \
|
||||
| (sbox[extract_byte(w[i],1)] << 8) \
|
||||
| (sbox[extract_byte(w[i],2)] << 16) \
|
||||
| (sbox[extract_byte(w[i],3)] << 24)
|
||||
w[i] = sbox[extract_byte (w[i], 0)] << 0 \
|
||||
| sbox[extract_byte (w[i], 1)] << 8 \
|
||||
| sbox[extract_byte (w[i], 2)] << 16 \
|
||||
| sbox[extract_byte (w[i], 3)] << 24
|
||||
|
||||
#define k_vec_xor(n) \
|
||||
for (int y = k_vec[(n)]; y > 0; y >>= 1) \
|
||||
{ \
|
||||
z ^= x * (y & 1); \
|
||||
x = ((x << 1) ^ ((x >> 7) * 0xc3)) & 0xff; \
|
||||
#define k_xor(n) \
|
||||
for (int i = (n); i > 0; i /= 2) \
|
||||
{ \
|
||||
z ^= x * (i % 2); \
|
||||
x = (x << 1) ^ ((x >> 7) * 0xc3); \
|
||||
x &= 0xff; \
|
||||
}
|
||||
|
||||
DECLSPEC void kuznyechik_linear (u32 *w)
|
||||
{
|
||||
// used inside k_vec_xor macro
|
||||
// used in k_xor macro
|
||||
u32 x;
|
||||
u32 z;
|
||||
|
||||
@ -116,22 +111,23 @@ DECLSPEC void kuznyechik_linear (u32 *w)
|
||||
{
|
||||
z = 0;
|
||||
|
||||
x = extract_byte (w[3], 3); k_vec_xor (15);
|
||||
x = extract_byte (w[3], 2); k_vec_xor (14);
|
||||
x = extract_byte (w[3], 1); k_vec_xor (13);
|
||||
x = extract_byte (w[3], 0); k_vec_xor (12);
|
||||
x = extract_byte (w[2], 3); k_vec_xor (11);
|
||||
x = extract_byte (w[2], 2); k_vec_xor (10);
|
||||
x = extract_byte (w[2], 1); k_vec_xor ( 9);
|
||||
x = extract_byte (w[2], 0); k_vec_xor ( 8);
|
||||
x = extract_byte (w[1], 3); k_vec_xor ( 7);
|
||||
x = extract_byte (w[1], 2); k_vec_xor ( 6);
|
||||
x = extract_byte (w[1], 1); k_vec_xor ( 5);
|
||||
x = extract_byte (w[1], 0); k_vec_xor ( 4);
|
||||
x = extract_byte (w[0], 3); k_vec_xor ( 3);
|
||||
x = extract_byte (w[0], 2); k_vec_xor ( 2);
|
||||
x = extract_byte (w[0], 1); k_vec_xor ( 1);
|
||||
x = extract_byte (w[0], 0); k_vec_xor ( 0);
|
||||
// k_xor (1) yields the same result as a simple xor
|
||||
x = extract_byte (w[3], 3); z ^= x;
|
||||
x = extract_byte (w[3], 2); k_xor (148);
|
||||
x = extract_byte (w[3], 1); k_xor (32);
|
||||
x = extract_byte (w[3], 0); k_xor (133);
|
||||
x = extract_byte (w[2], 3); k_xor (16);
|
||||
x = extract_byte (w[2], 2); k_xor (194);
|
||||
x = extract_byte (w[2], 1); k_xor (192);
|
||||
x = extract_byte (w[2], 0); z ^= x;
|
||||
x = extract_byte (w[1], 3); k_xor (251);
|
||||
x = extract_byte (w[1], 2); z ^= x;
|
||||
x = extract_byte (w[1], 1); k_xor (192);
|
||||
x = extract_byte (w[1], 0); k_xor (194);
|
||||
x = extract_byte (w[0], 3); k_xor (16);
|
||||
x = extract_byte (w[0], 2); k_xor (133);
|
||||
x = extract_byte (w[0], 1); k_xor (32);
|
||||
x = extract_byte (w[0], 0); k_xor (148);
|
||||
|
||||
// right-shift data block, prepend calculated byte
|
||||
w[3] = (w[3] << 8) | (w[2] >> 24);
|
||||
@ -143,7 +139,7 @@ DECLSPEC void kuznyechik_linear (u32 *w)
|
||||
|
||||
DECLSPEC void kuznyechik_linear_inv (u32 *w)
|
||||
{
|
||||
// used inside k_vec_xor macro
|
||||
// used in k_xor macro
|
||||
u32 x;
|
||||
u32 z;
|
||||
|
||||
@ -157,21 +153,21 @@ DECLSPEC void kuznyechik_linear_inv (u32 *w)
|
||||
w[2] = (w[2] >> 8) | (w[3] << 24);
|
||||
w[3] = (w[3] >> 8);
|
||||
|
||||
x = extract_byte (w[0], 0); k_vec_xor ( 0);
|
||||
x = extract_byte (w[0], 1); k_vec_xor ( 1);
|
||||
x = extract_byte (w[0], 2); k_vec_xor ( 2);
|
||||
x = extract_byte (w[0], 3); k_vec_xor ( 3);
|
||||
x = extract_byte (w[1], 0); k_vec_xor ( 4);
|
||||
x = extract_byte (w[1], 1); k_vec_xor ( 5);
|
||||
x = extract_byte (w[1], 2); k_vec_xor ( 6);
|
||||
x = extract_byte (w[1], 3); k_vec_xor ( 7);
|
||||
x = extract_byte (w[2], 0); k_vec_xor ( 8);
|
||||
x = extract_byte (w[2], 1); k_vec_xor ( 9);
|
||||
x = extract_byte (w[2], 2); k_vec_xor (10);
|
||||
x = extract_byte (w[2], 3); k_vec_xor (11);
|
||||
x = extract_byte (w[3], 0); k_vec_xor (12);
|
||||
x = extract_byte (w[3], 1); k_vec_xor (13);
|
||||
x = extract_byte (w[3], 2); k_vec_xor (14);
|
||||
x = extract_byte (w[0], 0); k_xor (148);
|
||||
x = extract_byte (w[0], 1); k_xor (32);
|
||||
x = extract_byte (w[0], 2); k_xor (133);
|
||||
x = extract_byte (w[0], 3); k_xor (16);
|
||||
x = extract_byte (w[1], 0); k_xor (194);
|
||||
x = extract_byte (w[1], 1); k_xor (192);
|
||||
x = extract_byte (w[1], 2); z ^= x;
|
||||
x = extract_byte (w[1], 3); k_xor (251);
|
||||
x = extract_byte (w[2], 0); z ^= x;
|
||||
x = extract_byte (w[2], 1); k_xor (192);
|
||||
x = extract_byte (w[2], 2); k_xor (194);
|
||||
x = extract_byte (w[2], 3); k_xor (16);
|
||||
x = extract_byte (w[3], 0); k_xor (133);
|
||||
x = extract_byte (w[3], 1); k_xor (32);
|
||||
x = extract_byte (w[3], 2); k_xor (148);
|
||||
|
||||
//append calculated byte
|
||||
w[3] |= (z << 24);
|
||||
|
@ -791,7 +791,7 @@ DECLSPEC void streebog256_transform (streebog256_ctx_t *ctx, const u32 *w0, cons
|
||||
streebog256_g (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
|
||||
|
||||
u64 counterbuf[8] = { 0 };
|
||||
counterbuf[7] = swap64_S ((u64) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
streebog256_add (ctx->n, counterbuf);
|
||||
|
||||
streebog256_add (ctx->s, m);
|
||||
@ -1479,7 +1479,7 @@ DECLSPEC void streebog256_transform_vector (streebog256_ctx_vector_t *ctx, const
|
||||
streebog256_g_vector (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
|
||||
|
||||
u64x counterbuf[8] = { 0 };
|
||||
counterbuf[7] = swap64 ((u64x) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
streebog256_add_vector (ctx->n, counterbuf);
|
||||
|
||||
streebog256_add_vector (ctx->s, m);
|
||||
|
@ -791,7 +791,7 @@ DECLSPEC void streebog512_transform (streebog512_ctx_t *ctx, const u32 *w0, cons
|
||||
streebog512_g (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
|
||||
|
||||
u64 counterbuf[8] = { 0 };
|
||||
counterbuf[7] = swap64_S ((u64) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
streebog512_add (ctx->n, counterbuf);
|
||||
|
||||
streebog512_add (ctx->s, m);
|
||||
@ -1479,7 +1479,7 @@ DECLSPEC void streebog512_transform_vector (streebog512_ctx_vector_t *ctx, const
|
||||
streebog512_g_vector (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
|
||||
|
||||
u64x counterbuf[8] = { 0 };
|
||||
counterbuf[7] = swap64 ((u64x) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
streebog512_add_vector (ctx->n, counterbuf);
|
||||
|
||||
streebog512_add_vector (ctx->s, m);
|
||||
|
@ -441,7 +441,7 @@ __kernel void m06212_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -449,7 +449,7 @@ __kernel void m06212_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -441,7 +441,7 @@ __kernel void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -449,7 +449,7 @@ __kernel void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -511,7 +511,7 @@ __kernel void m06213_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -585,7 +585,7 @@ __kernel void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -593,7 +593,7 @@ __kernel void m06222_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -585,7 +585,7 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -593,7 +593,7 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -655,7 +655,7 @@ __kernel void m06223_comp (KERN_ATTR_TMPS_ESALT (tc64_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -699,7 +699,7 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -707,7 +707,7 @@ __kernel void m06232_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -699,7 +699,7 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -707,7 +707,7 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -769,7 +769,7 @@ __kernel void m06233_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -411,7 +411,7 @@ __kernel void m13751_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -419,7 +419,7 @@ __kernel void m13751_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -427,7 +427,7 @@ __kernel void m13751_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -411,7 +411,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -419,7 +419,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -427,7 +427,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -473,7 +473,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -481,7 +481,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -489,7 +489,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -497,7 +497,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -505,7 +505,7 @@ __kernel void m13752_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -411,7 +411,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_aes (esalt_bufs, ukey1, ukey2, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -419,7 +419,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_serpent (esalt_bufs, ukey1, ukey2) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -427,7 +427,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_twofish (esalt_bufs, ukey1, ukey2) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -473,7 +473,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_aes_twofish (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -481,7 +481,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_serpent_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -489,7 +489,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -497,7 +497,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -505,7 +505,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -551,7 +551,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_aes_twofish_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -559,7 +559,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_serpent_twofish_aes (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -567,7 +567,7 @@ __kernel void m13753_comp (KERN_ATTR_TMPS_ESALT (tc_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -31,12 +31,12 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
u64x padding[8] = { 0 };
|
||||
u64x message[8];
|
||||
|
||||
padding[7] = swap64 ((u64x) 0x01);
|
||||
padding[7] = 0x0100000000000000;
|
||||
|
||||
//inner HMAC: ipad + message
|
||||
|
||||
//first transform: precalculated ipad hash
|
||||
counterbuf[7] = swap64 ((u64x) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
|
||||
//second transform: message = previous HMAC digest
|
||||
message[7] = hl32_to_64 (w3[2], w3[3]);
|
||||
@ -59,7 +59,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
|
||||
|
||||
counterbuf[7] = swap64 ((u64x) 0x400);
|
||||
counterbuf[7] = 0x0004000000000000;
|
||||
|
||||
//final: padding byte
|
||||
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64);
|
||||
@ -74,7 +74,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
//outer HMAC: opad + digest
|
||||
|
||||
//first transform: precalculated opad hash
|
||||
counterbuf[7] = swap64 ((u64x) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
|
||||
//second transform: message = inner HMAC digest
|
||||
message[0] = digest[0];
|
||||
@ -97,7 +97,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
|
||||
|
||||
counterbuf[7] = swap64 ((u64x) 0x400);
|
||||
counterbuf[7] = 0x0004000000000000;
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64);
|
||||
|
||||
|
@ -31,12 +31,12 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
u64x padding[8] = { 0 };
|
||||
u64x message[8];
|
||||
|
||||
padding[7] = swap64 ((u64x) 0x01);
|
||||
padding[7] = 0x0100000000000000;
|
||||
|
||||
//inner HMAC: ipad + message
|
||||
|
||||
//first transform: precalculated ipad hash
|
||||
counterbuf[7] = swap64 ((u64x) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
|
||||
//second transform: message = previous HMAC digest
|
||||
message[7] = hl32_to_64 (w3[2], w3[3]);
|
||||
@ -59,7 +59,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
|
||||
|
||||
counterbuf[7] = swap64 ((u64x) 0x400);
|
||||
counterbuf[7] = 0x0004000000000000;
|
||||
|
||||
//final: padding byte
|
||||
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64);
|
||||
@ -74,7 +74,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
//outer HMAC: opad + digest
|
||||
|
||||
//first transform: precalculated opad hash
|
||||
counterbuf[7] = swap64 ((u64x) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
|
||||
//second transform: message = inner HMAC digest
|
||||
message[0] = digest[0];
|
||||
@ -97,7 +97,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
|
||||
|
||||
counterbuf[7] = swap64 ((u64x) 0x400);
|
||||
counterbuf[7] = 0x0004000000000000;
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64);
|
||||
|
||||
@ -627,7 +627,7 @@ __kernel void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -635,7 +635,7 @@ __kernel void m13772_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
@ -31,12 +31,12 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
u64x padding[8] = { 0 };
|
||||
u64x message[8];
|
||||
|
||||
padding[7] = swap64 ((u64x) 0x01);
|
||||
padding[7] = 0x0100000000000000;
|
||||
|
||||
//inner HMAC: ipad + message
|
||||
|
||||
//first transform: precalculated ipad hash
|
||||
counterbuf[7] = swap64 ((u64x) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
|
||||
//second transform: message = previous HMAC digest
|
||||
message[7] = hl32_to_64 (w3[2], w3[3]);
|
||||
@ -59,7 +59,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
|
||||
|
||||
counterbuf[7] = swap64 ((u64x) 0x400);
|
||||
counterbuf[7] = 0x0004000000000000;
|
||||
|
||||
//final: padding byte
|
||||
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64);
|
||||
@ -74,7 +74,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
//outer HMAC: opad + digest
|
||||
|
||||
//first transform: precalculated opad hash
|
||||
counterbuf[7] = swap64 ((u64x) 0x200);
|
||||
counterbuf[7] = 0x0002000000000000;
|
||||
|
||||
//second transform: message = inner HMAC digest
|
||||
message[0] = digest[0];
|
||||
@ -97,7 +97,7 @@ DECLSPEC void hmac_streebog512_run_V (u32x *w0, u32x *w1, u32x *w2, u32x *w3, u6
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
|
||||
|
||||
counterbuf[7] = swap64 ((u64x) 0x400);
|
||||
counterbuf[7] = 0x0004000000000000;
|
||||
|
||||
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64);
|
||||
|
||||
@ -627,7 +627,7 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_kuznyechik (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -635,7 +635,7 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_camellia_serpent (esalt_bufs, ukey1, ukey2, ukey3, ukey4) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
@ -697,7 +697,7 @@ __kernel void m13773_comp (KERN_ATTR_TMPS_ESALT (vc64_sbog_tmp_t, tc_t))
|
||||
|
||||
if (verify_header_kuznyechik_serpent_camellia (esalt_bufs, ukey1, ukey2, ukey3, ukey4, ukey5, ukey6) == 1)
|
||||
{
|
||||
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
|
||||
if (atomic_inc (&hashes_shown[0]) == 0)
|
||||
{
|
||||
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
|
||||
}
|
||||
|
Loading…
Reference in New Issue
Block a user