arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-22 08:08:10 +00:00
Code Issues Releases Wiki Activity

Cleanup VeraCrypt related code

Browse Source 
Remove unnecessary constant variables by hardcoring values instead of
looking up. Precalculate swaps that are known at compile time. Hardcode
hashes_shown offset as zero for all TC / VC kernels.
This commit is contained in:
R. Yushaev 2018-11-29 12:49:03 +01:00
parent 01879c6395
commit 5efebb7b48
15 changed files with 107 additions and 111 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

96
OpenCL/inc_cipher_kuznyechik.cl
View File

@ -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);

4
OpenCL/inc_hash_streebog256.cl
View File

@ -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);

4
OpenCL/inc_hash_streebog512.cl
View File

@ -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);

4
OpenCL/m06212-pure.cl
View File

@ -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);
}

6
OpenCL/m06213-pure.cl
View File

@ -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);
}

4
OpenCL/m06222-pure.cl
View File

@ -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);
}

6
OpenCL/m06223-pure.cl
View File

@ -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);
}

4
OpenCL/m06232-pure.cl
View File

@ -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);
}

6
OpenCL/m06233-pure.cl
View File

@ -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);
}

6
OpenCL/m13751-pure.cl
View File

@ -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);
}

16
OpenCL/m13752-pure.cl
View File

@ -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);
}

22
OpenCL/m13753-pure.cl
View File

@ -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);
}

10
OpenCL/m13771-pure.cl
View File

@ -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);

14
OpenCL/m13772-pure.cl
View File

@ -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);
}

16
OpenCL/m13773-pure.cl
View File

@ -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);
}