github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-12-01 20:18:12 +01:00
Code Releases Activity

Merge pull request #1798 from Naufragous/cleanup-n-hardcode

Browse Source 
Cleanup VeraCrypt related code
This commit is contained in:
Jens Steube 2018-11-29 12:59:56 +01:00 committed by GitHub
commit 774b9bb3b6
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
15 changed files with 107 additions and 111 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

94
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, 0xfc, 0xee, 0xdd, 0x11, 0xcf, 0x6e, 0x31, 0x16,
0xfb, 0xc4, 0xfa, 0xda, 0x23, 0xc5, 0x04, 0x4d, 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 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, 0xa5, 0x2d, 0x32, 0x8f, 0x0e, 0x30, 0x38, 0xc0,
0x54, 0xe6, 0x9e, 0x39, 0x55, 0x7e, 0x52, 0x91, 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 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 extract_byte(x,n) (((x) >> (8 * (n))) & 0xff)
#define k_lookup(w,sbox) \ #define k_lookup(w,sbox) \
for (int i = 0; i < 4; i++) \ for (int i = 0; i < 4; i++) \
w[i] = (sbox[extract_byte(w[i],0)] << 0) \ w[i] = sbox[extract_byte (w[i], 0)] << 0 \
| (sbox[extract_byte(w[i],1)] << 8) \ | sbox[extract_byte (w[i], 1)] << 8 \
| (sbox[extract_byte(w[i],2)] << 16) \ | sbox[extract_byte (w[i], 2)] << 16 \
| (sbox[extract_byte(w[i],3)] << 24) | sbox[extract_byte (w[i], 3)] << 24
#define k_vec_xor(n) \ #define k_xor(n) \
for (int y = k_vec[(n)]; y > 0; y >>= 1) \ for (int i = (n); i > 0; i /= 2) \
{ \ { \
z ^= x * (y & 1); \ z ^= x * (i % 2); \
x = ((x << 1) ^ ((x >> 7) * 0xc3)) & 0xff; \ x = (x << 1) ^ ((x >> 7) * 0xc3); \
x &= 0xff; \
} }
DECLSPEC void kuznyechik_linear (u32 *w) DECLSPEC void kuznyechik_linear (u32 *w)
{ {
// used inside k_vec_xor macro // used in k_xor macro
u32 x; u32 x;
u32 z; u32 z;
@ -116,22 +111,23 @@ DECLSPEC void kuznyechik_linear (u32 *w)
{ {
z = 0; z = 0;
x = extract_byte (w[3], 3); k_vec_xor (15); // k_xor (1) yields the same result as a simple xor
x = extract_byte (w[3], 2); k_vec_xor (14); x = extract_byte (w[3], 3); z ^= x;
x = extract_byte (w[3], 1); k_vec_xor (13); x = extract_byte (w[3], 2); k_xor (148);
x = extract_byte (w[3], 0); k_vec_xor (12); x = extract_byte (w[3], 1); k_xor (32);
x = extract_byte (w[2], 3); k_vec_xor (11); x = extract_byte (w[3], 0); k_xor (133);
x = extract_byte (w[2], 2); k_vec_xor (10); x = extract_byte (w[2], 3); k_xor (16);
x = extract_byte (w[2], 1); k_vec_xor ( 9); x = extract_byte (w[2], 2); k_xor (194);
x = extract_byte (w[2], 0); k_vec_xor ( 8); x = extract_byte (w[2], 1); k_xor (192);
x = extract_byte (w[1], 3); k_vec_xor ( 7); x = extract_byte (w[2], 0); z ^= x;
x = extract_byte (w[1], 2); k_vec_xor ( 6); x = extract_byte (w[1], 3); k_xor (251);
x = extract_byte (w[1], 1); k_vec_xor ( 5); x = extract_byte (w[1], 2); z ^= x;
x = extract_byte (w[1], 0); k_vec_xor ( 4); x = extract_byte (w[1], 1); k_xor (192);
x = extract_byte (w[0], 3); k_vec_xor ( 3); x = extract_byte (w[1], 0); k_xor (194);
x = extract_byte (w[0], 2); k_vec_xor ( 2); x = extract_byte (w[0], 3); k_xor (16);
x = extract_byte (w[0], 1); k_vec_xor ( 1); x = extract_byte (w[0], 2); k_xor (133);
x = extract_byte (w[0], 0); k_vec_xor ( 0); 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 // right-shift data block, prepend calculated byte
w[3] = (w[3] << 8) | (w[2] >> 24); 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) DECLSPEC void kuznyechik_linear_inv (u32 *w)
{ {
// used inside k_vec_xor macro // used in k_xor macro
u32 x; u32 x;
u32 z; u32 z;
@ -157,21 +153,21 @@ DECLSPEC void kuznyechik_linear_inv (u32 *w)
w[2] = (w[2] >> 8) | (w[3] << 24); w[2] = (w[2] >> 8) | (w[3] << 24);
w[3] = (w[3] >> 8); w[3] = (w[3] >> 8);
x = extract_byte (w[0], 0); k_vec_xor ( 0); x = extract_byte (w[0], 0); k_xor (148);
x = extract_byte (w[0], 1); k_vec_xor ( 1); x = extract_byte (w[0], 1); k_xor (32);
x = extract_byte (w[0], 2); k_vec_xor ( 2); x = extract_byte (w[0], 2); k_xor (133);
x = extract_byte (w[0], 3); k_vec_xor ( 3); x = extract_byte (w[0], 3); k_xor (16);
x = extract_byte (w[1], 0); k_vec_xor ( 4); x = extract_byte (w[1], 0); k_xor (194);
x = extract_byte (w[1], 1); k_vec_xor ( 5); x = extract_byte (w[1], 1); k_xor (192);
x = extract_byte (w[1], 2); k_vec_xor ( 6); x = extract_byte (w[1], 2); z ^= x;
x = extract_byte (w[1], 3); k_vec_xor ( 7); x = extract_byte (w[1], 3); k_xor (251);
x = extract_byte (w[2], 0); k_vec_xor ( 8); x = extract_byte (w[2], 0); z ^= x;
x = extract_byte (w[2], 1); k_vec_xor ( 9); x = extract_byte (w[2], 1); k_xor (192);
x = extract_byte (w[2], 2); k_vec_xor (10); x = extract_byte (w[2], 2); k_xor (194);
x = extract_byte (w[2], 3); k_vec_xor (11); x = extract_byte (w[2], 3); k_xor (16);
x = extract_byte (w[3], 0); k_vec_xor (12); x = extract_byte (w[3], 0); k_xor (133);
x = extract_byte (w[3], 1); k_vec_xor (13); x = extract_byte (w[3], 1); k_xor (32);
x = extract_byte (w[3], 2); k_vec_xor (14); x = extract_byte (w[3], 2); k_xor (148);
//append calculated byte //append calculated byte
w[3] |= (z << 24); 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); streebog256_g (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
u64 counterbuf[8] = { 0 }; u64 counterbuf[8] = { 0 };
counterbuf[7] = swap64_S ((u64) 0x200); counterbuf[7] = 0x0002000000000000;
streebog256_add (ctx->n, counterbuf); streebog256_add (ctx->n, counterbuf);
streebog256_add (ctx->s, m); 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); streebog256_g_vector (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
u64x counterbuf[8] = { 0 }; u64x counterbuf[8] = { 0 };
counterbuf[7] = swap64 ((u64x) 0x200); counterbuf[7] = 0x0002000000000000;
streebog256_add_vector (ctx->n, counterbuf); streebog256_add_vector (ctx->n, counterbuf);
streebog256_add_vector (ctx->s, m); 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); streebog512_g (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
u64 counterbuf[8] = { 0 }; u64 counterbuf[8] = { 0 };
counterbuf[7] = swap64_S ((u64) 0x200); counterbuf[7] = 0x0002000000000000;
streebog512_add (ctx->n, counterbuf); streebog512_add (ctx->n, counterbuf);
streebog512_add (ctx->s, m); 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); streebog512_g_vector (ctx->h, ctx->n, m, ctx->s_sbob_sl64);
u64x counterbuf[8] = { 0 }; u64x counterbuf[8] = { 0 };
counterbuf[7] = swap64 ((u64x) 0x200); counterbuf[7] = 0x0002000000000000;
streebog512_add_vector (ctx->n, counterbuf); streebog512_add_vector (ctx->n, counterbuf);
streebog512_add_vector (ctx->s, m); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 (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); 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 padding[8] = { 0 };
u64x message[8]; u64x message[8];
padding[7] = swap64 ((u64x) 0x01); padding[7] = 0x0100000000000000;
//inner HMAC: ipad + message //inner HMAC: ipad + message
//first transform: precalculated ipad hash //first transform: precalculated ipad hash
counterbuf[7] = swap64 ((u64x) 0x200); counterbuf[7] = 0x0002000000000000;
//second transform: message = previous HMAC digest //second transform: message = previous HMAC digest
message[7] = hl32_to_64 (w3[2], w3[3]); 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); streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
counterbuf[7] = swap64 ((u64x) 0x400); counterbuf[7] = 0x0004000000000000;
//final: padding byte //final: padding byte
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); 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 //outer HMAC: opad + digest
//first transform: precalculated opad hash //first transform: precalculated opad hash
counterbuf[7] = swap64 ((u64x) 0x200); counterbuf[7] = 0x0002000000000000;
//second transform: message = inner HMAC digest //second transform: message = inner HMAC digest
message[0] = digest[0]; 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); 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); 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 padding[8] = { 0 };
u64x message[8]; u64x message[8];
padding[7] = swap64 ((u64x) 0x01); padding[7] = 0x0100000000000000;
//inner HMAC: ipad + message //inner HMAC: ipad + message
//first transform: precalculated ipad hash //first transform: precalculated ipad hash
counterbuf[7] = swap64 ((u64x) 0x200); counterbuf[7] = 0x0002000000000000;
//second transform: message = previous HMAC digest //second transform: message = previous HMAC digest
message[7] = hl32_to_64 (w3[2], w3[3]); 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); streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
counterbuf[7] = swap64 ((u64x) 0x400); counterbuf[7] = 0x0004000000000000;
//final: padding byte //final: padding byte
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); 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 //outer HMAC: opad + digest
//first transform: precalculated opad hash //first transform: precalculated opad hash
counterbuf[7] = swap64 ((u64x) 0x200); counterbuf[7] = 0x0002000000000000;
//second transform: message = inner HMAC digest //second transform: message = inner HMAC digest
message[0] = digest[0]; 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); 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); 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 (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); 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 (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); 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 padding[8] = { 0 };
u64x message[8]; u64x message[8];
padding[7] = swap64 ((u64x) 0x01); padding[7] = 0x0100000000000000;
//inner HMAC: ipad + message //inner HMAC: ipad + message
//first transform: precalculated ipad hash //first transform: precalculated ipad hash
counterbuf[7] = swap64 ((u64x) 0x200); counterbuf[7] = 0x0002000000000000;
//second transform: message = previous HMAC digest //second transform: message = previous HMAC digest
message[7] = hl32_to_64 (w3[2], w3[3]); 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); streebog512_g_vector (digest, counterbuf, message, s_sbob_sl64);
counterbuf[7] = swap64 ((u64x) 0x400); counterbuf[7] = 0x0004000000000000;
//final: padding byte //final: padding byte
streebog512_g_vector (digest, counterbuf, padding, s_sbob_sl64); 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 //outer HMAC: opad + digest
//first transform: precalculated opad hash //first transform: precalculated opad hash
counterbuf[7] = swap64 ((u64x) 0x200); counterbuf[7] = 0x0002000000000000;
//second transform: message = inner HMAC digest //second transform: message = inner HMAC digest
message[0] = digest[0]; 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); 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); 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 (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); 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 (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); 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 (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); mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, 0, gid, 0);
} }